NVIDIA CUDA 高度并行处理器编程（二）：数据并行执行模型

内置变量：内置变量的值一般都由运行时系统预初始化，例如 CUDA 的 kernel 函数中， gradDim，blockDim，blockIdx，threadIdx都是内置变量，它们的值由 CUDA 运行时系统预初始化，可以在 kernel 函数中引用。其他地方应避免使用这些变量。

1. CUDA的线程组织

上一节向量加法中的线程被组织成二级的层次结构：一个网络包含一个或更多的线程块，每块包含一个或更多的线程。一个块中所有线程的 blockId 相同，每个块中又可通过唯一的 threadIdx 访问唯一的线程。
网络一般是由线程块组成的三维数组，线程块又是线程组成的三维数组。不需要的维度可以设为 1 。网络确切组织形式是由 kernel 函数启动语句中的配备参数决定的 (<<< 配备参数 1， 配备参数 2>>>)，参数 1 确定了网格的维度和线程块的数量，参数 2 确定了线程块的维度和每个线程块内线程的数量。每个参数都是 dim3 类型，dim3 类型为含有 x, y 和 z 三个无符号整形的结构体，分别对应三个维度。
下面的主机端代码启动 vecAddKernel() 函数，生成了包含 32 个一维线程块的一维网格，每个线程块包含 128 个线程：

dim3 dimGrid(32, 1, 1);
dim3 dimBlock(128, 1, 1);
vecAddKernel<<<dimGrid, dimBlock>>>(...);

对于向量加法，可以写成如下形式：

dim3 dimGrid(ceil(n/256.0), 1, 1);
dim3 dimBlock(256, 1, 1);
vecAddKernel<<<dimGrid, dimBlock>>>(...);

CUDA C提供了用一维网络和线程启动 kernel 的简单方式，就跟上一节的一样：

vecAddKernel<<<ceil(n/256.0), 256>>>(...);

在 CUDA C 中 gridDim.x，gridDim.y 和 gridDim.z 的取值范围是 1 ~ 65536。在所有的线程块中，blockIdx.x 的取值范围从 0 到 gridDim.x - 1，blockIdx.y 的取值范围从 0 到 gridDim.y - 1，blockIdx.z 的取值范围从 0 到 gridDim.z - 1。
在线程块的配置上，将 z 字段设为 1 可得到二维线程块，将 y, z 设为 1 可得到一维线程块。计算能力低于 2.0 的设备一个线程块最多包含 512 个线程，大于等于 2.0 的设备一个线程块最多包含 1024 个线程。（512， 1， 1）、（8， 16， 4）都是允许的 blockDim值。网格可以比线程块有更高的维度，反之亦然。可以使用下面的主机代码生成下图的网络：

dim3 dimGrid(2, 2, 1);
dim3 dimBlock(4, 2, 2);
KernelFunction<<<dimGrid, dimBlock>>>(...);

这个网络包含 4 个线程块，每个线程块可以通过（blockIdx.y, blockIdx.x）标记。例如，（1，0）位置上线程块的 blockIdx.y=1，blockIdx.x=0。每个线程块由 4x2x2 的线程数组组成。例如线程(1, 0, 2)的 threadIdx.z=1, threadIdx.y=0, threadIdx.z=2。注意标签是按照 z, y, x 的顺序来的。

2. 线程与多维数组的映射

线程组织形式的设置要依据数据的特点，灰度图片为二维数组，那就把网络和线程块组织成二维形式处理灰度图片。
以下图为例，对于一张 72x62 的图片用 16x16 的线程块处理，则在 x，y 方向分别需要 5 个和 4 个线程块。但 x，y 方向上线程块分别有 4 个线程和 2 个线程多余，需要像 vecAddKernel 函数一样用 if 语句防止多余的线程进行运算。