CUDA编程：其二、网格与线程块

前文提到了GPU的强大之处是处理并行，那自然是要开大量线程。本文主要描述GPU如何组织大量线程。

一、GPU线程的组织结构

通俗的讲，GPU可以划分为多个网格（grid），每个网格下面有多个线程块（block），每个线程块上面可以运行多个线程（thread）。
但是注意，通常情况下，一个核函数只能运行在一个网格上。
一个网格之下，按照3维结构组织线程块，直观地看，按照一个正方体去组织block，长宽高分别在xyz三个方向扩展。如果在Z方向上的长度是1，那么grid中block的维度退化为2维；如果YZ方向上的长度都是1，那么grid中block的维度退化为1维。

block之下，thread的组织方式以此类推。

如题所示，假如这是一个grid，蓝色表示block，绿色表示thread，block维度是3×3×3，thread维度是4×4×4。
二、threadIdx、blockIdx、blockDim和gridDim
这一章重点介绍四个重要名词。
（1）gridDim
gridDim表示grid中block的维度大小，grid创建好之后可以看作常量。
gridDim.x、gridDim.y、gridDim.z分别表示线程格各个维度的大小，所以在上图中均为3。

（2）blockDim
blockDim表示block中thread的维度大小，block创建好之后可以看作常量。
blockDim.x、blockDim.y、blockDim.z分别表示线程格各个维度的大小，所以在上图中均为4。

（3）blockIdx
blockIdx表示在一个grid中该block所处的坐标，是变量。(blockIdx.x, blockIdx.y, blockIdx.z)就是block坐标。

（4）threadIdx
threadIdx表示在一个block中该线程所处的坐标，是变量。(threadId.x, threadIdx.y, threadIdx.z)就是线程坐标。

个人感觉这个名词起的有点歧义，记住一点：gridDim和blockIdx相对应；blockDim和threadIdx相对应。

三、线程ID。

我们设想，如果将grid中的线程全部铺展开，线性排列，如何定义线程ID呢？
（1）首先要找到线程所在block在grid中的id：

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

这个公式解释起来相当麻烦，可以参考这篇博客：

https://zhuanlan.zhihu.com/p/544864997

（2）要确定线程在block中的局部id:

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

（3）计算一个block中有多少个线程M:

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

（4）求得全局线程ID

idx = threadId + M*blockId;

线程ID的作用在并行计算中有很多体现。