下载掌阅APP，畅读海量书库
立即打开
当核函数在主机端启动后，它的执行会移动到设备端，此时，设备端会产生大量的工作项，也就是线程（Thread）。每个线程都执行核函数指定的语句。当核函数开始执行时，如何组织GPU线程是MXMACA编程的关键部分。MXMACA明确了一个两层的线程层次结构，其由线程块（Block）和线程网格（Grid）组成。
在一个核函数执行之前，需要指定一个 N 维的线程网格。一个线程网格是一个一维、二维或三维的索引空间。还需要指定全局工作线程的数目和线程块中工作线程的数目。MXMACA线程管理示意图如图3-6所示，以图中的线程网格为例，全局工作线程网格两个维度的线程范围为{ m × j ， n × k }，线程块的线程范围为{ j ， k }，总共有 j × k 个线程块。如果线程块的GPU线程范围为{ j ，2 k }，则该核函数的线程网格的范围就变为{ m × j ，（ n /2）×（2 k ）}。其中， m 、 n 、 j 和 k 是根据应用程序设计的需要进行定义的变量。定义线程块主要是为有些仅需在线程块内交换数据的程序提供方便，不过线程块内的线程数的多少要受到AP的资源限制。
由一个核函数启动所产生的所有线程统称为一个线程网格，同一线程网格的所有线程共享相同的全局内存空间。一个线程网格有多个线程块，一个线程块包含一组线程，同一线程块内的线程可以通过同步、共享内存的方式进行协作，不同线程块之间的线程不能协作。线程依靠以下两个坐标变量来区分彼此。
图3-6 MXMACA线程管理示意图
● blockIdx：线程块在线程网格内的索引。
● threadIdx：线程在线程块内的索引。
这两个坐标变量是基于unit3定义的MXMACA内置向量类型，是包含3个无符号整数的结构，可通过x、y、z三个字段指定。这些是核函数中需要预初始化的内置变量。当执行一个核函数时，MXMACA会为每个线程分配坐标变量blockIdx和threadIdx。基于这些坐标变量，程序员可以将不同的数据分配给不同的线程。
MXMACA可以组织三维的线程网格和线程块，它们的维度由以下两个内置变量来指定。
● blockDim：线程块的维度，用每个线程块中的线程数来表示。
● gridDim：线程网格的维度，用每个线程网格中的线程块数来表示。
它们是dim3类型的变量，也是整型向量，用来表示维度。当定义一个dim3类型变量时，所有未指定的元素都被初始化为1，dim3类型的每个元素也可以通过其x、y、z字段来获得。
图3-7是线程网格、线程块、线程示意图。图中，1个线程网格包含27个线程块（深灰色的格子），每个线程块又包含64个线程（浅灰色的格子），线程是最小的单位。
以图3-7为例，把线程网格和线程块都看作一个三维的矩阵。这里假设线程网格是一个3×3×3的三维矩阵，线程块是一个4×4×4的三维矩阵。举例说明各个变量的用法。
● gridDim：gridDim.x、gridDim.y、gridDim.z分别表示线程网格各个维度的大小，那么
● blockDim：blockDim.x、blockDim.y、blockDim.z分别表示线程块中各个维度的大小，那么
● blockIdx：blockIdx.x、blockIdx.y、blockIdx.z分别表示当前线程块所处的线程网格的坐标位置。
图3-7 MXMACA编程中的线程网格、线程块、线程示意图
● threadIdx：threadIdx.x、threadIdx.y、threadIdx.z分别表示当前线程所处的线程块的坐标位置。
因此，在MXMACA中有以下两组不同的线程网格和线程块变量。
● 手动定义的dim3数据类型：在主机端使用，作为核函数调用的一部分，定义一个线程网格和数据块的维度，仅在主机端可见。
● 预定义的unit3数据类型：在运行核函数时生成的线程网格、线程块和线程变量，可在核函数内被访问到，仅在设备端可见。
对于给定的数据大小，确定线程网格和线程块尺寸的一般步骤为：（1）确定线程块大小；（2）在已知数据大小和线程块大小的基础上计算线程网格的维度。
要确定线程块的大小，通常需要考虑核函数性能特征和GPU资源限制等因素。MXMACA的特点之一就是通过编程模型揭示了一个两层的线程层次结构。由于一个核函数启动的线程网格和线程块的维度数会影响性能，这一结构为程序员优化性能提供了额外的途径。线程网格和线程块从逻辑上代表了一个核函数的线程层次结构。
第5章会结合示例对线程管理进行更详尽的介绍。