本文参考https://blog.csdn.net/hujingshuang/article/details/53097222上的内容
不同在于,本文 将对原文进行补充,同时澄清一些易混淆的概念
CUDA中
grid、block、thread只要在最大值范围内,可以随意设置,最后系统会根据内部算法重新分配。因此grid、block、thread的设置其实只以你的“目的”有关。也就是说这几个值的索引如何方便如何来,要处理的矩阵就用二维,要处理的向量就用一维,要处理的张量就用三维
在设置上,变量是遵循向下设置的, 即
dim3 grid(1, 1, 1)表示的是一个grid由一个block组成
dim3 grid(4, 1, 1)表示的是一个grid由四个一维block组成,不要混淆为有4个grid
block(8, 1, 1)表示的是一个block由1维thread组成,1乘8的一维分布,不要混淆为有8个block
block(4, 2, 1)表示的是一个block由2维thread组成,4乘2的二维分布,不要混淆为有4个长度为2的block
这就是为什么没有threadDim.x,threadDim.y,threadDim.z的原因,如果你在window下配置CUDA,用VS编写的时候会发现没有threadDim变量
threadIdx.x,blockIdx.x,gridIdx.x下标都是从0开始
因为本质上是在网格索引,所以从x方向开始查找和从y方向开始查找结果是一样的。对于一个block(4, 2, 1)而言
int tid = threadIdx.x*blockDim.y + threadIdx.y;
int tid = threadIdx.y*blockDim.x + threadIdx.x;
以上两个索引结果等价。
在索引的时候要计算,不要自己套 公式,例如:
dim3 grid(8, 1, 1), block(2, 1, 1); 的情况,索引就是:
int tid = blockIdx.x * blockDim.x + threadIdx.x;而不是原文中的blockIdx.x * gridDim.x + threadIdx.x, 原文只是恰巧gridDim.x = blockDim.x = 4 而已
在dim3 grid(2, 2, 1), block(2, 2, 1); 这种block和thread都是二维分布的情况下:
int tid = (blockIdx.y * gridDim.x + blockIdx.x)*(blockDim.x * blockDim.y)
+ threadIdx.y*blockDim.x + threadIdx.x;
int tid = (blockIdx.x * gridDim.y + blockIdx.y)*(blockDim.x * blockDim.y)
+ threadIdx.x*blockDim.y + threadIdx.y;
以上是等价 的
在三维情况下,更是多种等价的查询,比如可以走列切片,或是走横切片,或是走高度切片,其实意义不大,只要按顺序索引就行。