1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
|
// 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 + 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);
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]);
{ 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;
}
|
