From 8f953e86e73a621b1d28fd39d4517ba2de538f65 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Mon, 19 Oct 2026 09:27:19 -0400 Subject: 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. --- tools/cuda-ab/blk512.cu | 18 ++++++++++++++++-- 1 file changed, 16 insertions(+), 2 deletions(-) (limited to 'tools/cuda-ab') 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