CUDA编程笔记003. 线程索引

1、线程索引

CUDA提供了两层的层次线程模型：

• Grid：
  • 定义：执行同一个kernel代码的所有thread集合称为一个grid。
  • 位于同一grid中的thread共享相同的global memory。
  • 一个grid包含多个block。
• Block：
  • 定义：一种thread集合，同一block的thread可通过block内同步和block内显存贡献来相互协作。

对应地，在kernel中可使用两层坐标来索引thread：

• 上层：blockIdx（在一个grid中的block索引）；
• 下层：threadIdx（在某一个block中的thread索引）。

通俗地来说，程序要执行某一个kernel时，在kernel内部想知道当前所处的线程，可以先通过blockIdx找到当前线程位于哪个block，然后在根据threadIdx找到block中具体的thread位置。

kernel，grid，block和thread的关系如下：

2、索引细节

• Host侧

在host侧，由程序员自行指定CUDA kernel的grid和block维度。

kernel_name <<<grid, block>>>(argument list);

CUDA提供了对应的数据类型是dim3。3维大小不必全部设置，可按需设置，未设置的维度默认为1，取各维度值时，可用.x, .y, .z三个成员变量。

比如6000x80的矩阵，想要划分为grid=10x20x30，block=1x20x4的话。写法如下：

dim3 grid(10, 20, 30); // 指定grid中3维block的各维度大小
dim3 block(1, 20, 4); // 指定block中3维thread的各维度大小

testKernel<<<grid, block>>>()；// 传给grid, block。

疑问1：grid和block对于kernel来说，是变量还是常量？是否在程序启动时可调。

• Device侧

CUDA对kernel，预定义了两套变量，分别是：

（1）获取维度：

• gridDim(uint3类型):和host侧grid值一致。
• blockDim(uint3类型)：和host侧block值一致。

（2）获取各维度索引：

• blockIdx(uint3类型)：grid中block的索引值；
• threadIdx(uint3类型)：指定block中thread的索引值。

3、动手测试

Demo代码（GitHub）：

__global__ void CheckIndex(void) {
  printf("Device: 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() {
  printf("CUDA Demo: illustrate block and thread index.\n\n");

  int n_elem = 6;
  // 计算block和grid数量：1个grid包含多个block，1个block包含多个thread.
  dim3 block(3); // block内的thread分布
  dim3 grid((n_elem + block.x - 1) / block.x); // grid中的block分布

  printf("Host: block:(%d, %d, %d), grid:(%d, %d, %d)\n\n",
         block.x, block.y, block.z, grid.x, grid.y, grid.z);

  CheckIndex<<<grid, block>>>();
  cudaDeviceReset();
}

打印结果：

Host: block:(3, 1, 1), grid:(2, 1, 1)

Device: threadIdx:(0, 0, 0), blockIdx:(0, 0, 0), blockDim:(3, 1, 1), gridDim:(2, 1, 1)
Device: threadIdx:(1, 0, 0), blockIdx:(0, 0, 0), blockDim:(3, 1, 1), gridDim:(2, 1, 1)
Device: threadIdx:(2, 0, 0), blockIdx:(0, 0, 0), blockDim:(3, 1, 1), gridDim:(2, 1, 1)
Device: threadIdx:(0, 0, 0), blockIdx:(1, 0, 0), blockDim:(3, 1, 1), gridDim:(2, 1, 1)
Device: threadIdx:(1, 0, 0), blockIdx:(1, 0, 0), blockDim:(3, 1, 1), gridDim:(2, 1, 1)
Device: threadIdx:(2, 0, 0), blockIdx:(1, 0, 0), blockDim:(3, 1, 1), gridDim:(2, 1, 1)

参考

[1] “Professional Cuda C Programming” Chapter 2.