CUDA编程模型#2

线程管理

由一个内核启动所产生的所有线程统称为一个网格。同一网格中的所有线程共享相同的全局内存空间。一个网格由多个线程块构成，一个线程块包含一组线程，同一线程块内的线程协作可以通过以下方式来实现。

• 同步
• 共享内存

不同块内的线程不能协作。

线程依靠以下两个坐标变量来区分彼此。

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

执行一个核函数时，CUDA运行时为每个线程分配坐标变量blockIdx和threadIdx。基于这些坐标，可以将部分数据分配给不同的线程。

该坐标变量是基于uint3定义的cuda内置的向量类型，是一个包含3个无符号整数的结构，可以通过x,y,z三个字段来指定。

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

CUDA可以组织三维的网格和块。网格和块的维度由下列两个内置变量指定

• blockDim（线程块的维度，用每个线程块中的线程数来表示）
• gridDim（线程格的维度，用每个线程格中的线程数来表示）

它们是dim3类型的变量，基于uint3定义的整数型向量，用来表示维度。

在CUDA程序中，有两组不同的网格和块变量：手动定义的dim3数据类型和预定义的uint3数据类型。在主机端，作为内核调用的一部分，可以使用dim3数据类型定义一个网格和块的维度。当执行核函数时，CUDA运行时会生成相应的内置预初始化的网格、块和线程变量，它们在核函数内均可被访问到且为uint3类型。

举个栗子

#include <cuda_runtime.h>
#include <stdio.h>
#include <cstdlib>

__global__ void checkIndex(){
        printf("threadIdx:(%d, %d, %d) blockIdx:(%d, %d, %d) blockDim:(%d, %d, %d)"
                        "gridDim:(%d, %d, %d)\n",threadIdx.x,threadIdx.y,threadIdx.z,
                        blockIdx.x,blockIdx.y,blockIdx.z,blockDim.x,blockDim.y,blockDim.z,
                        gridDim.x,gridDim.y,gridDim.z);
}

int main(int argc, char **argv){
        int nElem = std::atoi(argv[1]);

        dim3 block (3);
        dim3 grid ((nElem+block.x-1)/block.x);

        fprintf(stdout,"grid.x %d, grid.y %d, grid.z %d\n",grid.x, grid.y, grid.z);
        fprintf(stdout,"block.x %d, block.y %d, block.z %d\n",block.x, block.y, block.z);

        checkIndex<<<grid, block>>>();

        cudaDeviceReset();

        return 0;
}

运行命令

nvcc -o check checkDimension.cu

./check 3

运行结果