小伙伴们看了是不是有点晕呢其实LZ现在对于CUDA理解也是很浅显算是看书的一个笔记吧首先threadIdx, blockIdx, blockDim, gridDimwarpsize这五个变量都是CUDA的内建变量gridDim指定grid维度类型为dim3blockDIm指定block维度类型为dim3blockIdx:指定grid内block索引号类型为uint3threadIdx指定block内thread索引号类型为uint3warpsize指定warp内thread数量类型为intdim3类型它的定义在头文件vector_types.h中看定义等同于uint3用于指定线程或线程块的维度在没有赋值时初始化值为1.struct__device_builtin__dim3{unsignedintx,y,z;#ifdefined(__cplusplus)#if__cplusplus201103L__host__ __device__ constexprdim3(unsignedintvx1,unsignedintvy1,unsignedintvz1):x(vx),y(vy),z(vz){}#else__host__ __device__dim3(unsignedintvx1,unsignedintvy1,unsignedintvz1):x(vx),y(vy),z(vz){}#endif__host__ __device__dim3(uint3 v):x(v.x),y(v.y),z(v.z){}__host__ __device__ operatoruint3(void){uint3 t;t.xx;t.yy;t.zz;returnt;}#endif/* __cplusplus */};具体的定义类型有很多详细的定义可以在头文件中自行查阅typedef__device_builtin__structchar1char1;typedef__device_builtin__structuchar1uchar1;typedef__device_builtin__structchar2char2;typedef__device_builtin__structuchar2uchar2;typedef__device_builtin__structchar3char3;typedef__device_builtin__structuchar3uchar3;typedef__device_builtin__structchar4char4;typedef__device_builtin__structuchar4uchar4;typedef__device_builtin__structshort1short1;typedef__device_builtin__structushort1ushort1;typedef__device_builtin__structshort2short2;typedef__device_builtin__structushort2ushort2;typedef__device_builtin__structshort3short3;typedef__device_builtin__structushort3ushort3;typedef__device_builtin__structshort4short4;typedef__device_builtin__structushort4ushort4;typedef__device_builtin__structint1int1;typedef__device_builtin__structuint1uint1;typedef__device_builtin__structint2int2;typedef__device_builtin__structuint2uint2;typedef__device_builtin__structint3int3;typedef__device_builtin__structuint3uint3;typedef__device_builtin__structint4int4;typedef__device_builtin__structuint4uint4;typedef__device_builtin__structlong1long1;typedef__device_builtin__structulong1ulong1;typedef__device_builtin__structlong2long2;typedef__device_builtin__structulong2ulong2;typedef__device_builtin__structlong3long3;typedef__device_builtin__structulong3ulong3;typedef__device_builtin__structlong4long4;typedef__device_builtin__structulong4ulong4;typedef__device_builtin__structfloat1float1;typedef__device_builtin__structfloat2float2;typedef__device_builtin__structfloat3float3;typedef__device_builtin__structfloat4float4;typedef__device_builtin__structlonglong1longlong1;typedef__device_builtin__structulonglong1ulonglong1;typedef__device_builtin__structlonglong2longlong2;typedef__device_builtin__structulonglong2ulonglong2;typedef__device_builtin__structlonglong3longlong3;typedef__device_builtin__structulonglong3ulonglong3;typedef__device_builtin__structlonglong4longlong4;typedef__device_builtin__structulonglong4ulonglong4;typedef__device_builtin__structdouble1double1;typedef__device_builtin__structdouble2double2;typedef__device_builtin__structdouble3double3;typedef__device_builtin__structdouble4double4;每个grid中有多个Block每个block有多个线程束线程每个线程束包含多个线程.线程束的概念用的比较少但是要明白GPU不是一个线程一个线程处理的而是以线程束的方式来处理的。下图为书《 GPU编程与优化》中截图但是在nvidia官方文件中可以看到可以看到两者的坐标并不一致所以参考多方资料应该是按照这种方式来进行索引的可以看到一个grid有6个block每个block有8个线程。可以按照如下的方式计算对应的索引号。不同维度block和thread索引方式#includecuda_runtime.h#includedevice_launch_parameters.h#includestdio.h#includestdlib.h#includeiostreamusing namespace std;//thread 1D__global__voidtestThread1(int*c,constint*a,constint*b){intithreadIdx.x;c[i]b[i]-a[i];}//thread 2D__global__voidtestThread2(int*c,constint*a,constint*b){intithreadIdx.xthreadIdx.y*blockDim.x;c[i]b[i]-a[i];}//thread 3D__global__voidtestThread3(int*c,constint*a,constint*b){intithreadIdx.xthreadIdx.y*blockDim.xthreadIdx.z*blockDim.x*blockDim.y;c[i]b[i]-a[i];}//block 1D__global__voidtestBlock1(int*c,constint*a,constint*b){intiblockIdx.x;c[i]b[i]-a[i];}//block 2D__global__voidtestBlock2(int*c,constint*a,constint*b){intiblockIdx.xblockIdx.y*gridDim.x;c[i]b[i]-a[i];}//block 3D__global__voidtestBlock3(int*c,constint*a,constint*b){intiblockIdx.xblockIdx.y*gridDim.xblockIdx.z*gridDim.x*gridDim.y;c[i]b[i]-a[i];}//block-thread 1D-1D__global__voidtestBlockThread1(int*c,constint*a,constint*b){intithreadIdx.xblockDim.x*blockIdx.x;c[i]b[i]-a[i];}//block-thread 1D-2D__global__voidtestBlockThread2(int*c,constint*a,constint*b){intthreadId_2DthreadIdx.xthreadIdx.y*blockDim.x;intithreadId_2D(blockDim.x*blockDim.y)*blockIdx.x;c[i]b[i]-a[i];}//block-thread 1D-3D__global__voidtestBlockThread3(int*c,constint*a,constint*b){intthreadId_3DthreadIdx.xthreadIdx.y*blockDim.xthreadIdx.z*blockDim.x*blockDim.y;intithreadId_3D(blockDim.x*blockDim.y*blockDim.z)*blockIdx.x;c[i]b[i]-a[i];}//block-thread 2D-1D__global__voidtestBlockThread4(int*c,constint*a,constint*b){intblockId_2DblockIdx.xblockIdx.y*gridDim.x;intithreadIdx.xblockDim.x*blockId_2D;c[i]b[i]-a[i];}//block-thread 3D-1D__global__voidtestBlockThread5(int*c,constint*a,constint*b){intblockId_3DblockIdx.xblockIdx.y*gridDim.xblockIdx.z*gridDim.x*gridDim.y;intithreadIdx.xblockDim.x*blockId_3D;c[i]b[i]-a[i];}//block-thread 2D-2D__global__voidtestBlockThread6(int*c,constint*a,constint*b){intthreadId_2DthreadIdx.xthreadIdx.y*blockDim.x;intblockId_2DblockIdx.xblockIdx.y*gridDim.x;intithreadId_2D(blockDim.x*blockDim.y)*blockId_2D;c[i]b[i]-a[i];}//block-thread 2D-3D__global__voidtestBlockThread7(int*c,constint*a,constint*b){intthreadId_3DthreadIdx.xthreadIdx.y*blockDim.xthreadIdx.z*blockDim.x*blockDim.y;intblockId_2DblockIdx.xblockIdx.y*gridDim.x;intithreadId_3D(blockDim.x*blockDim.y*blockDim.z)*blockId_2D;c[i]b[i]-a[i];}//block-thread 3D-2D__global__voidtestBlockThread8(int*c,constint*a,constint*b){intthreadId_2DthreadIdx.xthreadIdx.y*blockDim.x;intblockId_3DblockIdx.xblockIdx.y*gridDim.xblockIdx.z*gridDim.x*gridDim.y;intithreadId_2D(blockDim.x*blockDim.y)*blockId_3D;c[i]b[i]-a[i];}//block-thread 3D-3D__global__voidtestBlockThread9(int*c,constint*a,constint*b){intthreadId_3DthreadIdx.xthreadIdx.y*blockDim.xthreadIdx.z*blockDim.x*blockDim.y;intblockId_3DblockIdx.xblockIdx.y*gridDim.xblockIdx.z*gridDim.x*gridDim.y;intithreadId_3D(blockDim.x*blockDim.y*blockDim.z)*blockId_3D;c[i]b[i]-a[i];}voidaddWithCuda(int*c,constint*a,constint*b,unsignedintsize){int*dev_a0;int*dev_b0;int*dev_c0;cudaSetDevice(0);cudaMalloc((void**)dev_c,size*sizeof(int));cudaMalloc((void**)dev_a,size*sizeof(int));cudaMalloc((void**)dev_b,size*sizeof(int));cudaMemcpy(dev_a,a,size*sizeof(int),cudaMemcpyHostToDevice);cudaMemcpy(dev_b,b,size*sizeof(int),cudaMemcpyHostToDevice);//testThread11, size(dev_c, dev_a, dev_b);//uint3 s;s.x size/5;s.y 5;s.z 1;//testThread2 1,s(dev_c, dev_a, dev_b);//uint3 s; s.x size / 10; s.y 5; s.z 2;//testThread31, s (dev_c, dev_a, dev_b);//testBlock1size,1 (dev_c, dev_a, dev_b);//uint3 s; s.x size / 5; s.y 5; s.z 1;//testBlock2s, 1 (dev_c, dev_a, dev_b);//uint3 s; s.x size / 10; s.y 5; s.z 2;//testBlock3s, 1 (dev_c, dev_a, dev_b);//testBlockThread1size/10, 10(dev_c, dev_a, dev_b);//uint3 s1; s1.x size / 100; s1.y 1; s1.z 1;//uint3 s2; s2.x 10; s2.y 10; s2.z 1;//testBlockThread2 s1, s2 (dev_c, dev_a, dev_b);//uint3 s1; s1.x size / 100; s1.y 1; s1.z 1;//uint3 s2; s2.x 10; s2.y 5; s2.z 2;//testBlockThread3 s1, s2 (dev_c, dev_a, dev_b);//uint3 s1; s1.x 10; s1.y 10; s1.z 1;//uint3 s2; s2.x size / 100; s2.y 1; s2.z 1;//testBlockThread4 s1, s2 (dev_c, dev_a, dev_b);//uint3 s1; s1.x 10; s1.y 5; s1.z 2;//uint3 s2; s2.x size / 100; s2.y 1; s2.z 1;//testBlockThread5 s1, s2 (dev_c, dev_a, dev_b);//uint3 s1; s1.x size / 100; s1.y 10; s1.z 1;//uint3 s2; s2.x 5; s2.y 2; s2.z 1;//testBlockThread6 s1, s2 (dev_c, dev_a, dev_b);//uint3 s1; s1.x size / 100; s1.y 5; s1.z 1;//uint3 s2; s2.x 5; s2.y 2; s2.z 2;//testBlockThread7 s1, s2 (dev_c, dev_a, dev_b);//uint3 s1; s1.x 5; s1.y 2; s1.z 2;//uint3 s2; s2.x size / 100; s2.y 5; s2.z 1;//testBlockThread8 s1, s2 (dev_c, dev_a, dev_b);uint3 s1;s1.x5;s1.y2;s1.z2;uint3 s2;s2.xsize/200;s2.y5;s2.z2;testBlockThread9s1,s2(dev_c,dev_a,dev_b);cudaMemcpy(c,dev_c,size*sizeof(int),cudaMemcpyDeviceToHost);cudaFree(dev_a);cudaFree(dev_b);cudaFree(dev_c);cudaGetLastError();}intmain(){constintn1000;int*anewint[n];int*bnewint[n];int*cnewint[n];int*ccnewint[n];for(inti0;in;i){a[i]rand()%100;b[i]rand()%100;c[i]b[i]-a[i];}addWithCuda(cc,a,b,n);FILE*fpfopen(out.txt,w);for(inti0;in;i)fprintf(fp,%d %d\n,c[i],cc[i]);fclose(fp);bool flagtrue;for(inti0;in;i){if(c[i]!cc[i]){flagfalse;break;}}if(flagfalse)printf(no pass);elseprintf(pass);cudaDeviceReset();delete[]a;delete[]b;delete[]c;delete[]cc;getchar();return0;}参考地址和书籍https://www.cnblogs.com/rainbow70626/p/6498738.html?utm_sourceitdadaoutm_mediumreferralhttps://www.cnblogs.com/tiandsp/p/9458734.htmlGPU编程与优化https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#creating-a-graph-using-stream-capturehttps://zhuanlan.zhihu.com/p/99947605?from_voters_pagetrue
