CUDA 编程（一）--grid、block、thread

1.计算机架构

正常程序执行在host（CPU）端，当需要到device（GPU）执行时，启动kernel（核函数，一个kernel对应一个线程）调用GPU。（K-G-B-T）

• 一个kernel对应一个grid，启动一个kernel的时候需要制定grid和block的维度
•  一个grid可以有多个block，一维~三维
•  一个block可以有多个thread，一维~三维

每个thread有独立的Register和Local Memory，每个block有独立的share memory，一个block中的所有thread共享share memory，每个grid 有独立的global、constant、texture memory，所有的block共享global memory、constant memory、texture memory。

2.线程管理

由一个内核启动所产生的所有线程统称为一个网格（grid）。同一网格中的所有线程共享相同的全局内存空间（相当于CPU中的系统内存）。

一个网格由多个线程块（block）构成，一个线程块包含一组线程，同一线程块内的线程协作通过同步、共享内存实现。不同线程块（block）内的线程不能协作。

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

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

网格（grid）和块（block）的维度由下列两个内置变量指定。

• blockDim（线程块的维度，每个块中的线程数目）
• gridDim（线程格的维度，启动块的数目）

网格和块从逻辑上代表了一个核函数的线程层次结构。CUDA的特点之一就是通过编程模型揭示了一个两层的线程层次结构。

3. thread遍历

3.1 一维block、grid

dim3 block(blockDim);  //一维block
这一行创建了一个三维向量block，用来定义每个block的大小。blockDim表示每个block中包含blockDim个线程。dim3数据类型是CUDA中的一个特殊数据类型，用于表示三维向量。在这个情况下，block的其余维度将被默认设置为1。这意味着你将有一个包含blockDim个线程的一维block。
dim3 grid(gridDim); // 一维grid
这一行创建了一个三维向量grid，这意味着整个grid将包含gridDim个block。与block一样，只传递一个整数值给grid，所以其余维度将被默认设置为1，得到一个一维grid。

3.2 多维

dim3 block(blockDimx, blockDimy) ; //二维

dim3 grid(gridDimx, gridDimy); // 二维grid

dim3 block(blockDimx, blockDimy，blockDimz) ; //三维

dim3 grid(gridDimx, gridDimy,，gridDimz); // 三维grid

3.3 遍历

block方向寻找threadId：

先遍历z方向（每一个z方向都有blockDim.x * blockDim.y的thread），再遍历y方向（每一个y方向都有blockDim.x的thread），最后遍历x方向。

如需在grid方向寻找threadId：

1. 首先计算blockSize：blockSize = blockDim.z * blockDim.y * blockDim.z ;（每个block中含有多少个thread）
2. 在grid方向查找block坐标前有多少个block，bIndex = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x;
3. 在block方向查找thread坐标前有多少个thread：tIndex = threadIdx.z * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
4. 最后计算整个index：index = bIndex * blockSize + tIndex ;

__global__ void print_thread_idx_kernel(){
int bSize = blockDim.z * blockDim.y * blockDim.x;
int bIndex = blockIdx.z * gridDim.y * gridDim.x + \
             blockIdx.y * gridDim.x + \
             blockIdx.x;
int tIndex = threadIdx.z * blockDim.y * blockDim.x + \
             threadIdx.y * blockDim.x + \
             threadIdx.x;
int index = bIndex * bSize + tIndex;
printf("block idx: %3d, thread idx in block: %3d, thread idx: %3d\n",
        bIndex, tIndex, index);
}