diff options
| -rw-r--r-- | tools/cuda-ab/blk512.cu | 18 |
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; } |