CUDA 线程索引计算

终于搞清楚了 thread 索引的计算方式，简单来说很像小学学的除法公式
被除数 = 除数 * 商 + 余数
*用公式表示：最终的线程Id = blockId blockSize + threadId**

blockId ：当前 block 在 grid 中的坐标（可能是1维到3维）
blockSize ：block 的大小，描述其中含有多少个 thread
threadId ：当前 thread 在 block 中的坐标（同样从1维到3维）

下面先理清几个关键点：
grid 中含有若干个 blocks，其中 blocks 的数量由 gridDim.x/y/z 来描述。某个 block 在此 grid 中的坐标由 blockIdx.x/y/z 描述。
blocks 中含有若干个 threads，其中 threads 的数量由 blockDim.x/y/z 来描述。某个 thread 在此 block 中的坐标由 threadIdx.x/y/z 描述。
接着一个多维的坐标如何用一维数据表达呢？这里大家想一想两位数和三位数，就是很好的例子。数字 = 百位数字 100 + 十位数字 10 + 个位数字。

当我们得知每个维度上的大小时，就可以利用这样的进制将三维坐标转换为1维坐标 。
一般来说坐标（x, y, z）分别所在的维度大小是（Dx, Dy, Dz），一般会把 z 看成高纬度，接着是 y ，最后是 x。
*高维度坐标转一维坐标公式 id = Dx Dy z + Dx y + x**

搞清楚了这些，我们找几个例子开始计算

1D grid, 1D block

blockSize = blockDim.x
blockId = blockIdx.x
threadId = threadIdx.x

Id = blockIdx.x * blockDim.x + threadIdx.x

3D grid, 1D block

blockSize = blockDim.x（一维 block 的大小）
blockId = Dx Dy z + Dx y + x （三维 grid 中 block 的 id，用公式）
= gridDim.x gridDim.y blockIdx.z + gridDim.x blockIdx.y + blockIdx.x
threadId = threadIdx.x （一维 block 中 thread 的 id）

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

1D grid, 2D block

blockSize = blockDim.x * blockDim.y（二维 block 的大小）
blockId = blockIdx.x（一维 grid 中 block id）
threadId = Dx y + x （二维 block 中 thread 的 id）
= blockDim.x threadIdx.y + threadIdx.x

Id = blockIdx.x (blockDim.x blockDim.y) + blockDim.x * threadIdx.y + threadIdx.x

3D grid, 3D block

blockSize = blockDim.x blockDim.y blockDim.z（三维 block 的大小）
blockId = Dx Dy z + Dx y + x（三维 grid 中 block 的 id，用公式）
= gridDim.x gridDim.y blockIdx.z + gridDim.x blockIdx.y + blockIdx.x
threadId = Dx Dy z + Dx y + x（三维 block 中 thread 的 id，用公式）
= blockDim.x blockDim.y threadIdx. z + blockDim.x threadIdx.y + threadIdx.x

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