我的世界做头像的网站网站搜索优化方法
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个线程，可以去捋一捋输出结果
参考文章
https://zhuanlan.zhihu.com/p/544864997
https://blog.csdn.net/qq_43715171/article/details/121794135