线程组织：内置变量

2024-6-3

CUDA

在上一篇文章中，我们已经了解了 CUDA 执行配置的三大组成：线程、线程块 block 和 网格 grid。

1. threadIdx

threadIdx 是一个内置的变量，用来表示当前线程在其线程块中的索引。

具体的索引规则，举个例子就能很清楚的了解。如图 1 所示，有两个 block，threadIdx 都是独立于所在的 block 进行依次“编码”的。

图1 threadIdx 举例

如图 2 所示，再举一个二维 block 的例子，“编码”规则也不难理解。比如其中的 X，在第 0 行，第 1 列，所以 threadIdx.x=1,threadIdx.y=0；再比如 Q，在第 1 行，第 2 列，所以 threadIdx.x=2,threadIdx.y=1。

图2 threadIdx 举例

我们写一段代码来了解如何使用。如代码清单 1 所示，grid 的布局是 2x2，即 4 组 block，每组 block 里有 8x8 个线程。在核函数 print_threadIds 中，我们打印每个执行线程的 threadIdx 信息。

我们通过打印信息“验证”我们的理解：总共打印 256 次。threadIdx.x 的范围在 0 到 7；threadIdx.y 的范围也在 0 到 7；threadIdx.z 都是 0。而且 (threadIdx.x，threadIdx.y) 的相同组合，会出现 4 次。

代码清单 1 threadIdx
  1. #include "cuda_runtime.h"
  2. #include "device_launch_parameters.h"
  3.  
  4. #include <stdio.h>
  5.  
  6. __global__ void print_threadIds()
  7. {
  8.     printf("threadIdx.x=%d, threadIdx.y=%d, threadIdx.z=%d\n",
  9.         threadIdx.x, threadIdx.y, threadIdx.z);
  10. }
  11.  
  12. int main()
  13. {
  14.     int nx = 16, ny = 16;
  15.  
  16.     dim3 block(8, 8);
  17.     dim3 grid(nx / block.x, ny / block.y);
  18.  
  19.     print_threadIds << <grid, block>> > ();
  20.     cudaDeviceSynchronize();
  21.  
  22.     cudaDeviceReset();
  23.     return 0;
  24. }

2. blockIdx

blockIdx 也是一个内置的变量，用于表示当前执行线程块的索引。

我们看到图 3 的示例，blockIdx 已经和其内部的 thread 排布无关了，我们只关系整体的 block 排布。比如 X 和 Y 一样，blockIdx.x=0,blockIdx.y=0；T 和 U 一样，blockIdx.x=1,blockIdx.y=1。

图3 blockIdx 举例

3. blockDim 和 gridDim

我们再来看两个内置变量。如图 4 所示，blockDim 提供了当前执行的线程块的维度信息。如图 5 所示，gridDim 提供整个网格的维度信息。

图4 blockDim 举例
图5 gridDim 举例

最后我们写程序“验证”一下 blockIdx 和 blockDim/gridDim 的概念。如代码清单 2 所示，执行配置和代码清单 1 是一样的，不过每个执行线程里打印 blockIdx、blockDim、gridDim。

从打印信息来看，blockIdx.x 的范围在 0 到 1；blockIdx.y 的范围也在 0 到 1。blockDim 都为 8x8，gridDim 都为 2x2。

blockDim 和 gridDim 在核函数启动时被设定，在整个核函数执行期间保持不变。有一丝“常量”的意思。

代码清单 2 blockIdx 和 blockDim/gridDim
  1. #include "cuda_runtime.h"
  2. #include "device_launch_parameters.h"
  3.  
  4. #include <stdio.h>
  5.  
  6. __global__ void print_details()
  7. {
  8.     printf("blockIdx.x=%d, blockIdx.y=%d, blockIdx.z=%d, "
  9.            "blockDim.x=%d, blockDim.y=%d, "
  10.            "gridDim.x=%d, gridDim.y=%d\n",
  11.         blockIdx.x, blockIdx.y, blockIdx.z,
  12.         blockDim.x, blockDim.y,
  13.         gridDim.x, gridDim.y);
  14. }
  15.  
  16. int main()
  17. {
  18.     int nx = 16, ny = 16;
  19.  
  20.     dim3 block(8, 8);
  21.     dim3 grid(nx / block.x, ny / block.y);
  22.  
  23.     print_details << <grid, block>> > ();
  24.     cudaDeviceSynchronize();
  25.  
  26.     cudaDeviceReset();
  27.     return 0;
  28. }