cuda-ab: blk512 w256 kernel index fix

The committed w256 kernel mixed element and block strides
(bx*1024 + tx*4) and could not cover the buffer. Replace with the
well-formed index i = (blockIdx.y*16 + threadIdx.y)*256 +
blockIdx.x*16 + threadIdx.x, writing out[i..i+3] = i..i+3.

Verified on the emulated vGPU: w512 bad=0; w256 bad=0 with all 16
blocks zero-bad=256/full-bad=0; span sweep firstbad=-1, incl. the
page-15 tail (16376..16383). Two consecutive boots green. This
isolates the earlier w256 bad=12288 to the test kernel itself and
retires blk512 as an index-math control; A-read (arow) failures
remain tracked via ctaid_probe PROBE2.

1 files changed, 16 insertions, 2 deletions

diff --git a/tools/cuda-ab/blk512.cu b/tools/cuda-ab/blk512.cu
index 9b91c54..f657d39 100644
--- a/tools/cuda-ab/blk512.cu
+++ b/tools/cuda-ab/blk512.cu
@@ -5,8 +5,8 @@
 #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; }
+__global__ void w256(int *out){ int i = (blockIdx.y*16 + threadIdx.y)*256 + blockIdx.x*16 + threadIdx.x;
+  out[i]=i; }
 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);
@@ -21,5 +21,19 @@ int main(void){
   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]);
+  { int z; int firstbad=-1,lastbad=-1; for(z=0;z<n;z++){ if(h[z]!=z){ if(firstbad<0) firstbad=z; lastbad=z; } }
+    printf("BLK512 w256 span firstbad=%d lastbad=%d | p0[1000..1023]=%d %d %d %d %d %d %d %d %d %d | p1[0..7]=%d %d %d %d %d %d %d %d | p3[1016..1023]=%d %d %d %d %d %d %d %d | p4[0..7]=%d %d %d %d %d %d %d %d | p15[1016..1023]=%d %d %d %d %d %d %d %d\n",
+      firstbad,lastbad,
+      h[1000],h[1001],h[1010],h[1015],h[1016],h[1017],h[1020],h[1021],h[1022],h[1023],
+      h[1024],h[1025],h[1026],h[1027],h[1028],h[1029],h[1030],h[1031],
+      h[4096-8],h[4096-7],h[4096-6],h[4096-5],h[4096-4],h[4096-3],h[4096-2],h[4096-1],
+      h[4096],h[4097],h[4098],h[4099],h[4100],h[4101],h[4102],h[4103],
+      h[16384-8],h[16384-7],h[16384-6],h[16384-5],h[16384-4],h[16384-3],h[16384-2],h[16384-1]); }
+  { int b,bb; int zc=0,fc=0; int prevoff=0;
+    for(b=0;b<256;b++){ int cb=0; for(bb=0;bb<64;bb++) if(h[b*64+bb]!=b*64+bb) cb++;
+      if(cb==0) zc++; if(cb==64) fc++;
+      if(b<16) printf("BLK512 w256 blk%d bad=%d\n",b,cb); }
+    printf("BLK512 w256 blocks zero-bad=%d full-bad=%d\n",zc,fc); }
+
   printf("BLK512 DONE\n"); return 0;
 }