cuda学习3: 全局线程id计算

最新推荐文章于 2025-11-14 13:03:27 发布

原创最新推荐文章于 2025-11-14 13:03:27 发布 · 1.2k 阅读

11 ·

CC 4.0 BY-SA版权

文章标签：

#学习

cuda 专栏收录该内容

3 篇文章

订阅专栏

部署运行你感兴趣的模型镜像

cuda全局线程id计算

cuda线程模型

在这里再次提到几个概念：

Thread，线程，并行的基本单位
Block，线程块，包含多个线程，线程块中所有线程在同一个SM上执行
Grid，网格，由一组Block组成

线程在CUDA中以三维组织的形势组织。
在这里插入图片描述

如上图，把网格和线程块都看作一个三维的矩阵。这里假设网格是一个3x3x3的三维矩阵，线程块是一个4x4x4的三维矩阵，线程就是最小的绿色小格子。

CUDA可以组织3维的grid和block。blockIdx表示线程块在线程格内的索引，threadIdx表示块内的线程索引；blockDim表示每个线程块中的线程数，gridDim表示网格中的线程块数。

CUDA内建变量

CUDA有几个内建变量，可以直接使用，运行时获得网格和块的尺寸及线程索引等信息。位于device_launch_parameters.h头文件。

gridDim：包含三个元素x, y, z的结构体，表示网格在x,y,z方向上的尺寸，对应于执行配置中的第一个参数。
blockDim：包含三个元素x, y, z的结构体，表示块在x,y,z方向上的尺寸，对应于执行配置的第二个参数
blockIdx：包含三个元素x, y, z的结构体，分别表示当前线程所在块在网格中x, y, z方向上的索引
threadIdx：包含三个元素x, y, z的结构体，分别表示当前线程在其所在块中x, y, z方向上的索引
warpSize：表明warp的尺寸，在计算能力1.0的设备中，这个值是24，在1.0以上的设备中，这个值是32。

gridDim和blockDim都是dim3结构体，定义如下

struct __device_builtin__ dim3
{
    unsigned int x, y, z;
};

blockIdx和threadIdx都是uint3结构体，定义如下

struct __device_builtin__ uint3
{
    unsigned int x, y, z;
};

虚假的一维

__global__ void checkDim(void) {
    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() {
    checkDim<<<2, 3>>>();
    cudaDeviceSynchronize();
    return 0;
}

如上示例，我们在执行核函数的网格大小设置为2，线程块大小设置为3，都是一维数据，实际在执行的时候看看结果。

$ ./calcId
threadIdx:(0, 0, 0) blockIdx:(1, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1)
threadIdx:(1, 0, 0) blockIdx:(1, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1)
threadIdx:(2, 0, 0) blockIdx:(1, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1)
threadIdx:(0, 0, 0) blockIdx:(0, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1)
threadIdx:(1, 0, 0) blockIdx:(0, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1)
threadIdx:(2, 0, 0) blockIdx:(0, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1)

blockDim 和 gridDim其实都是三维数据，只是在它们的y，z数据都默认为1了。
和如下指定是等价的，只是简略写法而已。

    dim3 grid_size(2, 1, 1);
    dim3 block_size(3, 1, 1);
    checkDim<<<grid_size, block_size>>>();

全局线程索引threadId的计算

cuda按照3维网格嵌套3维线程块的方式来组织线程，我们在编程的时候是需要知道线程的位置——也就是全局的索引，本来可以有很多种计算方式，但是我们还是遵循统一的原则来处理。

先计算线程块的索引blockId

先找到当前线程位于线程格中的哪一个线程块blockId

blockId = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;

这个公式第一次看到是比较抽象的。
先看一个二维的示例
在这里插入图片描述

你怎么求得(1,1)的全局索引？方式有很多种，但是我们选择其中一种。

第一步：先拿去y==1的时候x方向索引，也就是1，相当于blockIdx.x

第二部：再加上y方向覆盖的数量（y==0这一层）, 相当于blockIdx.y*gridDim.x，gridDim.x就是x方向的宽度，blockIdx.y就是覆盖的高度。

至此，二维id计算可以理解了，拓展到三维就是：再加上Z方向的覆盖blockIdx.z*gridDim.x*gridDim.y，其中gridDim.x*gridDim.y是覆盖面积，blockIdx.z是Z方向覆盖高度。

再计算当前线程在线程块中的索引threadId

原理同线程块id计算是一样的

threadId = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;

计算一个线程块中一共有多少个线程M

M = blockDim.x*blockDim.y*blockDim.z

求得当前的线程序列号idx

idx = threadId + M*blockId;

公式汇总就是

idx = (threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y) + (blockIdx.x + blockIdx.y * gridDim.x + blockIdx.z * gridDim.x * gridDim.y) * (blockDim.x * blockDim.y * blockDim.z)

来个例子

#include <cuda_runtime_api.h>
#include <iostream>

__global__ void checkIndex(void) {
    // gridDim表示grid的维度，blockDim表示block的维度，grid维度表示grid中block的数量，block维度表示block中thread的数量

    unsigned int blockId = blockIdx.x + blockIdx.y * gridDim.x +
                           blockIdx.z * gridDim.x * gridDim.y;

    unsigned int blockCount = blockDim.x * blockDim.y * blockDim.z;

    unsigned int threadId = threadIdx.x + threadIdx.y * blockDim.x +
                            threadIdx.z * blockDim.x * blockDim.y;

    unsigned int idx = threadId + blockCount * blockId;

    printf(
        "idx: %d, tid: %d, bid: %d, threadIdx:(%d, %d, %d) blockIdx:(%d, %d, "
        "%d) blockDim:(%d, %d, %d) gridDim:(%d, %d, %d)\n",
        idx, threadId, blockId, 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() {
    dim3 grid_size(2, 3, 4);
    dim3 block_size(2, 2, 3);
    checkIndex<<<grid_size, block_size>>>();

    cudaDeviceSynchronize();
    return 0;
}

总共有288个线程，可以去捋一捋输出结果
在这里插入图片描述