diff options
| -rw-r--r-- | tools/cuda-ab/blk512.cu | 25 |
1 files changed, 25 insertions, 0 deletions
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; +} |