线程管理

当内核函数开始执行，如何组织GPU的线程就变成了最主要的问题了，我们必须明确，一个核函数只能有一个grid，一个grid可以有很多个块，每个块可以有很多的线程，这种分层的组织结构使得我们的并行过程更加自如灵活：
2021-10-11 21-51-19 的屏幕截图.png
接下来就是给每个线程一个编号了，我们知道每个线程都执行同样的一段串行代码，那么怎么让这段相同的代码对应不同的数据呢？首先第一步就是让这些线程彼此区分开，才能对应到相应从线程，使得这些线程也能区分自己的数据。如果线程本身没有任何标记，那么没办法确认其行为。
依靠下面两个内置结构体确定线程标号：

  • blockIdx（线程块在线程网格内的位置索引）
  • threadIdx（线程在线程块内的位置索引）

两个内置结构体基于 uint3 定义，包含三个无符号整数的结构，通过三个字段来指定：

  • blockIdx.x
  • blockIdx.y
  • blockIdx.z
  • threadIdx.x
  • threadIdx.y
  • threadIdx.z

他们是dim3类型(基于uint3定义的数据结构)的变量，也包含三个字段x,y,z.
网格和块的维度一般是二维和三维的，也就是说一个网格通常被分成二维的块，而每个块常被分成三维的线程。
注意：dim3是手工定义的，主机端可见。uint3是设备端在执行的时候可见的，不可以在核函数运行时修改，初始化完成后uint3值就不变了。他们是有区别的！这一点必须要注意。

示例：

  1. /*
  2. *1_check_dimension
  3. */
  4. #include <cuda_runtime.h>
  5. #include <stdio.h>
  6. __global__ void checkIndex(void)
  7. {
  8.   printf("threadIdx:(%d,%d,%d) blockIdx:(%d,%d,%d) blockDim:(%d,%d,%d)\
  9.   gridDim(%d,%d,%d)\n",threadIdx.x,threadIdx.y,threadIdx.z,
  10.   blockIdx.x,blockIdx.y,blockIdx.z,blockDim.x,blockDim.y,blockDim.z,
  11.   gridDim.x,gridDim.y,gridDim.z);
  12. }
  13. int main(int argc,char **argv)
  14. {
  15.   int nElem=6;
  16.   dim3 block(3);
  17.   dim3 grid((nElem+block.x-1)/block.x);
  18.   printf("grid.x %d grid.y %d grid.z %d\n",grid.x,grid.y,grid.z);
  19.   printf("block.x %d block.y %d block.z %d\n",block.x,block.y,block.z);
  20.   checkIndex<<<grid,block>>>();
  21.   cudaDeviceReset();
  22.   return 0;
  23. }

接下来这段代码是检查网格和块的大小的：

  1. /*
  2. *2_grid_block
  3. */
  4. #include <cuda_runtime.h>
  5. #include <stdio.h>
  6. int main(int argc,char ** argv)
  7. {
  8.   int nElem=1024;
  9.   dim3 block(1024);
  10.   dim3 grid((nElem-1)/block.x+1);
  11.   printf("grid.x %d block.x %d\n",grid.x,block.x);
  13.   block.x=512;
  14.   grid.x=(nElem-1)/block.x+1;
  15.   printf("grid.x %d block.x %d\n",grid.x,block.x);
  17.   block.x=256;
  18.   grid.x=(nElem-1)/block.x+1;
  19.   printf("grid.x %d block.x %d\n",grid.x,block.x);
  21.   block.x=128;
  22.   grid.x=(nElem-1)/block.x+1;
  23.   printf("grid.x %d block.x %d\n",grid.x,block.x);
  25.   cudaDeviceReset();
  26.   return 0;
  27. }

模板

  1. __global__ void kernel_name(argument list);

2021-10-11 21-19-24 的屏幕截图.png
Kernel核函数编写有以下限制

  1. 只能访问设备内存
  2. 必须有void返回类型
  3. 不支持可变数量的参数
  4. 不支持静态变量
  5. 显示异步行为

    并行化一个for

    串行：
    1. void sumArraysOnHost(float *A, float *B, float *C, const int N) {
    2. for (int i = 0; i < N; i++)
    3.  C[i] = A[i] + B[i];
    4. }
    并行：
    1. __global__ void sumArraysOnGPU(float *A, float *B, float *C) {
    2. int i = threadIdx.x;
    3. C[i] = A[i] + B[i];
    4. }
    验证： ```cpp /*

void sumArrays(float a,float b,float res,const int size) { for(int i=0;i<size;i+=4) { res[i]=a[i]+b[i]; res[i+1]=a[i+1]+b[i+1]; res[i+2]=a[i+2]+b[i+2]; res[i+3]=a[i+3]+b[i+3]; } } global void sumArraysGPU(floata,floatb,floatres) { int i=threadIdx.x; res[i]=a[i]+b[i]; }

int main(int argc,char **argv) { int dev = 0; cudaSetDevice(dev);

int nElem=32; printf(“Vector size:%d\n”,nElem); int nByte=sizeof(float)nElem; //host上面运行的东西表示为_h float a_h=(float)malloc(nByte); float b_h=(float)malloc(nByte); float res_h=(float)malloc(nByte); float res_from_gpu_h=(float*)malloc(nByte); memset(res_h,0,nByte); memset(res_from_gpu_h,0,nByte);

//device上面的就表示为_d float a_d,b_d,res_d; CHECK(cudaMalloc((float)&a_d,nByte)); CHECK(cudaMalloc((float)&b_d,nByte)); CHECK(cudaMalloc((float*)&res_d,nByte));

initialData(a_h,nElem); initialData(b_h,nElem);

//这一步表示的是从host上面拷贝内容？ CHECK(cudaMemcpy(a_d,a_h,nByte,cudaMemcpyHostToDevice)); CHECK(cudaMemcpy(b_d,b_h,nByte,cudaMemcpyHostToDevice));

//然后这他M的有是个啥，前面不是还1，2，3的写吗 dim3 block(nElem); dim3 grid(nElem/block.x); sumArraysGPU<<>>(a_d,b_d,res_d); printf(“Execution configuration<<<%d,%d>>>\n”,block.x,grid.x);

CHECK(cudaMemcpy(res_from_gpu_h,res_d,nByte,cudaMemcpyDeviceToHost)); sumArrays(a_h,b_h,res_h,nElem);

checkResult(res_h,res_from_gpu_h,nElem); cudaFree(a_d); cudaFree(b_d); cudaFree(res_d);

free(a_h); free(b_h); free(res_h); free(res_from_gpu_h);

return 0; }

  1. ![2021-10-11 21-43-16 的屏幕截图.png](https://cdn.nlark.com/yuque/0/2021/png/2771394/1633959810369-2fdfe7d0-fbb6-4a34-9a1e-ba7bc0d428ea.png#clientId=ub5da1d6c-e3e6-4&from=drop&id=ubd78af35&margin=%5Bobject%20Object%5D&name=2021-10-11%2021-43-16%20%E7%9A%84%E5%B1%8F%E5%B9%95%E6%88%AA%E5%9B%BE.png&originHeight=85&originWidth=721&originalType=binary&ratio=1&size=17851&status=done&style=none&taskId=u56adc0d0-b13b-4a9b-88f6-9389f324f35)<br />CUDA小技巧，当我们进行调试的时候可以把核函数配置成单线程的：
  2. ```cpp
  3. kernel_name<<<1,1>>>(argument list)

<<的表示，我真TM是什么都不会呀
2021-10-11 21-32-39 的屏幕截图.png
2021-10-11 21-37-55 的屏幕截图.png
日了，越来越不想学了

错误处理

我们就需要对错误进行防御性处理了，例如我们代码库头文件里面的这个宏：

  1. #define CHECK(call)\
  2. {\
  3.   const cudaError_t error=call;\
  4.   if(error!=cudaSuccess)\
  5.   {\
  6.       printf("ERROR: %s:%d,",__FILE__,__LINE__);\
  7.       printf("code:%d,reason:%s\n",error,cudaGetErrorString(error));\
  8.       exit(1);\
  9.   }\
  10. }

就是获得每个函数执行后的返回结果，然后对不成功的信息加以处理，CUDA C 的API每个调用都会返回一个错误代码，这个代码我们就可以好好利用了，当然在release版本中可以去除这部分，但是开发的时候一定要有的。学到了好像也没学到。。。。

编译执行

  1. nvcc xxxx.cu -o xxxx

懂了吗，我是没太懂。。。