From 44132d7887222a300dcc155e0376f445dca41655 Mon Sep 17 00:00:00 2001
From: Arthur Hanson Rasmusson <arthur@vgpu.io>
Date: Mon, 19 Oct 2026 08:55:48 -0400
Subject: cuda-ab: blk512 - minimal 2-D-grid CTA band-bug reproducer

w512<<<8,512>>> writes 16384 ints out[i]=i perfectly (bad=0 across
repeated runs); w256<<<(16,4),(16,16)>>> with the same write does
bad=12288 deterministically: rows 0..15 correct, rows >= 16 wrong. This is
the transformer/row-band failure without any probe machinery, and the host
scan reports the engine wrote 31/31 tracked pages for the failing launch,
i.e. the by>=1 stores exist but do not reach the pages the client reads.
Build: nvcc -O2 --cudart=static -arch=sm_75 -o blk512 blk512.cu
---
 tools/cuda-ab/blk512.cu | 25 +++++++++++++++++++++++++
 1 file changed, 25 insertions(+)
 create mode 100644 tools/cuda-ab/blk512.cu

diff --git a/tools/cuda-ab/blk512.cu b/tools/cuda-ab/blk512.cu
new file mode 100644
index 0000000..9b91c54
--- /dev/null
+++ b/tools/cuda-ab/blk512.cu
@@ -0,0 +1,25 @@
+// SPDX-License-Identifier: BSD-3-Clause
+#include <cstdio>
+#include <cstdlib>
+#include <cuda_runtime.h>
+#define CHECK(x) do { cudaError_t e_=(x); if(e_!=cudaSuccess){printf("CUDA_ERROR %s L%d\n", cudaGetErrorString(e_), __LINE__); printf("BLK512 DONE FAIL\n"); return 1;} } while(0)
+__global__ void w512(int *out){ int i = (blockIdx.x*512 + threadIdx.x)*4;
+  out[i]=i; out[i+1]=i+1; out[i+2]=i+2; out[i+3]=i+3; }
+__global__ void w256(int *out){ int i = blockIdx.y*16*256 + blockIdx.x*1024 + threadIdx.y*256 + threadIdx.x*4;
+  out[i]=i; out[i+1]=i+1; out[i+2]=i+2; out[i+3]=i+3; }
+int main(void){
+  int *d2, *d1; const int n=16384; int *h=(int*)malloc(n*4); int bad1=0,bad2=0,i;
+  setvbuf(stdout,NULL,_IONBF,0);
+  CHECK(cudaMalloc(&d2, n*4)); CHECK(cudaMalloc(&d1, n*4));
+  CHECK(cudaMemset(d2,0,n*4)); CHECK(cudaMemset(d1,0,n*4));
+  w512<<<8,512>>>(d2); CHECK(cudaGetLastError()); CHECK(cudaDeviceSynchronize());
+  CHECK(cudaMemcpy(h,d2,n*4,cudaMemcpyDeviceToHost));
+  for(i=0;i<n;i++) if(h[i]!=i) bad1++;
+  printf("BLK512 w512 grid=(8,1) block=(512,1) bad=%d first8=%d %d %d %d %d %d %d %d\n",bad1,h[0],h[1],h[2],h[3],h[4],h[5],h[6],h[7]);
+  CHECK(cudaMemset(d2,0,n*4));
+  w256<<<dim3(16,4),dim3(16,16)>>>(d1); CHECK(cudaGetLastError()); CHECK(cudaDeviceSynchronize());
+  CHECK(cudaMemcpy(h,d1,n*4,cudaMemcpyDeviceToHost));
+  for(i=0;i<n;i++) if(h[i]!=i) bad2++;
+  printf("BLK512 w256 grid=(16,4) block=(16,16) bad=%d first8=%d %d %d %d %d %d %d %d\n",bad2,h[0],h[1],h[2],h[3],h[4],h[5],h[6],h[7]);
+  printf("BLK512 DONE\n"); return 0;
+}
-- 
cgit v1.2.3