当前位置：首页 > news >正文

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

news 2025/5/3 3:46:59

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个线程，可以去捋一捋输出结果
在这里插入图片描述