NVIDIA CUDA2023春训营（三）CUDA 线程层次结构与线程索引

Reference

• 我的CUDA学习之旅——启程
• CUDA线程索引计算

线程层次结构

核函数在 device 端执行时，会启动若干线程，一个核函数所启动的所有线程被称为一个线程网格（Thread Grid），同一个线程网格上的线程共享相同的全局内存空间，每个线程网格又可分为若干线程块（Thread Block），每个线程块中又包含若干线程

通过如下语句，可以定义一个 3*2 的 grid，每个 block 中含有 5*3 的 thread 的二维线程组织，其中 gird 和 block 是定义为 dim3 类型的变量，在定义时，缺省值初始化为 1，可以灵活的定义 1-dim、2-dim、3-dim 结构

dim3 grid(3, 2, 1);
dim3 block(5, 3, 1);
kernel_function<<<grid, block>>>(prams...);

相应的线程层次如下图所示

线程坐标

在 CUDA 中，一个线程需要两个内置的坐标变量 (blockIdx，threadIdx) 来标识，它们都是 dim3 类型变量，其中 blockIdx 指明线程所在的 block 在 grid 中的位置，而 threaIdx 指明线程在 block 中的位置

blockIdx 与 threadIdx 可通过内置变量来获得：

• threadIdx.[x y z]：执行当前核函数的线程在 block 中的索引值
• blockIdx.[x y z]：执行当前核函数的线程所在的 block 在 grid 中的索引值

例如，threadIdx.x 是执行当前核函数的线程在 block 中的 x 方向的序号，blockIdx.x 是执行当前核函数的线程所在 block 在 grid 中的 x 方向的序号

#include <stdio.h>

__global__ void hello_from_gpu() {
    const int bx = blockIdx.x;
    const int tx = threadIdx.x;
    printf("block: %d, thread: %d\n", bx, tx);
}

int main() {
    hello_from_gpu<<<4, 4>>>();
    cudaDeviceSynchronize();
    return 0;
}

线程唯一索引

唯一索引计算公式

一个线程块上的线程是放在同一个流式多处理器（SM）上的，但是单个 SM 的资源有限，这就使得 block 中的线程数是有限制的，现代 GPUs 的 block 可支持的线程数可达 1024 个，这就使得有时若想要知道一个线程在所有线程中的全局 ID，就必须要知道相应的组织结构，通过以下两个内置变量，可获得相应的组织结构信息

• blockDim.[x y z]：一个 block 中包含多少个线程
• gridDim.[x y z]：一个 grid 中包含多少个 block

对于线程唯一索引，需要知道以下三个信息：

• blockIndex：block 在整个 grid 中的索引（1 维到 3 维）
• blockSize ：block 的大小，描述其中含有多少个线程
• threadId ：线程在 block 中的索引（1 维到 3 维）

进而有唯一索引计算公式：index = blockIndex * blockSize + threadIndex

1D grid, 1D block

对于 1-dim 的 grid 与 1-dim 的 block，有：

// block在整个一维grid中的索引
int blockIndex = blockIdx.x;
// 一维block的大小
int blockSize = blockDim.x;
// 线程在一维block中的索引
int threadIndex = threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

1D grid, 2D block

对于 1-dim 的 grid 与 2-dim 的 block，有：

// block在整个一维grid中的索引
int blockIndex = blockIdx.x;
// 二维block的大小
int blockSize = blockDim.x * blockDim.y;
// 线程在二维block中的索引
int threadIndex = blockDim.x * threadIdx.y + threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

1D grid, 3D block

对于 1-dim 的 grid 与 3-dim 的 block，有：

// block在整个一维grid中的索引
int blockIndex = blockIdx.x;
// 三维block的大小
int blockSize = blockDim.x * blockDim.y * blockDim.z;
// 线程在三维block中的索引
int threadIndex = blockDim.x * blockDim.y * threadIdx.z + blockDim.x * threadIdx.y + threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

2D grid, 1D block

对于 2-dim 的 grid 与 1-dim 的 block，有：

// block在整个二维grid中的索引
int blockIndex = gridDim.x * blockIdx.y + blockIdx.x
// 一维block的大小
int blockSize = blockDim.x;
// 线程在一维block中的索引
int threadIndex = threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

2D grid, 2D block

对于 2-dim 的 grid 与 2-dim 的 block，有：

// block在整个二维grid中的索引
int blockIndex = gridDim.x * blockIdx.y + blockIdx.x
// 二维block的大小
int blockSize = blockDim.x * blockDim.y;
// 线程在二维block中的索引
int threadIndex = blockDim.x * threadIdx.y + threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

2D grid, 3D block

对于 2-dim 的 grid 与 3-dim 的 block，有：

// block在整个二维grid中的索引
int blockIndex = gridDim.x * blockIdx.y + blockIdx.x
// 三维block的大小
int blockSize = blockDim.x * blockDim.y * blockDim.z;
// 线程在三维block中的索引
int threadIndex = blockDim.x * blockDim.y * threadIdx.z + blockDim.x * threadIdx.y + threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

3D grid, 1D block

对于 3-dim 的 grid 与 1-dim 的 block，有：

// block在整个三维grid中的索引
int blockIndex = gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x
// 一维block的大小
int blockSize = blockDim.x;
// 线程在一维block中的索引
int threadIndex = threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

3D grid, 2D block

对于 3-dim 的 grid 与 2-dim 的 block，有：

// block在整个三维grid中的索引
int blockIndex = gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x
// 二维block的大小
int blockSize = blockDim.x * blockDim.y;
// 线程在二维block中的索引
int threadIndex = blockDim.x * threadIdx.y + threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

3D grid, 3D block

对于 3-dim 的 grid 与 3-dim 的 block，有：

// block在整个三维grid中的索引
int blockIndex = gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x
// 三维block的大小
int blockSize = blockDim.x * blockDim.y * blockDim.z;
// 线程在三维block中的索引
int threadIndex = blockDim.x * blockDim.y * threadIdx.z + blockDim.x * threadIdx.y + threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;