本节书摘来自华章计算机《CUDA C编程权威指南》一书中的第2章，第2.3节，作者 [美] 马克斯·格罗斯曼（Max Grossman），译 颜成钢 殷建 李亮，更多章节内容可以访问云栖社区“华章计算机”公众号查看。
2.3 组织并行线程
从前面的例子可以看出，如果使用了合适的网格和块大小来正确地组织线程，那么可以对内核性能产生很大的影响。在向量加法的例子中，为了实现最佳性能我们调整了块的大小，并基于块大小和向量数据大小计算出了网格大小。
现在通过一个矩阵加法的例子来进一步说明这一点。对于矩阵运算，传统的方法是在内核中使用一个包含二维网格与二维块的布局来组织线程。但是，这种传统的方法无法获得最佳性能。在矩阵加法中使用以下布局将有助于了解更多关于网格和块的启发性的用法：
- 由二维线程块构成的二维网格
- 由一维线程块构成的一维网格
- 由一维线程块构成的二维网格
2.3.1 使用块和线程建立矩阵索引
通常情况下，一个矩阵用行优先的方法在全局内存中进行线性存储。图2-9所示的是一个8×6矩阵的小例子。
在一个矩阵加法核函数中，一个线程通常被分配一个数据元素来处理。首先要完成的任务是使用块和线程索引从全局内存中访问指定的数据。通常情况下，对一个二维示例来说，需要管理3种索引：
- 线程和块索引
- 矩阵中给定点的坐标
- 全局线性内存中的偏移量
对于一个给定的线程，首先可以通过把线程和块索引映射到矩阵坐标上来获取线程块和线程索引的全局内存偏移量，然后将这些矩阵坐标映射到全局内存的存储单元中。
第一步，可以用以下公式把线程和块索引映射到矩阵坐标上：
第二步，可以用以下公式把矩阵坐标映射到全局内存中的索引/存储单元上：
图2-10说明了块和线程索引、矩阵坐标以及线性全局内存索引之间的对应关系。
printThreadInfo函数被用于输出关于每个线程的以下信息：
- 线程索引
- 块索引
- 矩阵坐标
- 线性全局内存偏移量
- 相应元素的值
用以下命令编译并运行该程序：
对于每个线程，你可以获取以下信息：
图2-11说明了这三项索引之间的关系。
2.3.2 使用二维网格和二维块对矩阵求和
在本节中，我们将使用一个二维网格和二维块来编写一个矩阵加法核函数。首先，应编写一个校验主函数以验证矩阵加法核函数是否能得出正确的结果：
然后，创建一个新的核函数，目的是采用一个二维线程块来进行矩阵求和：
这个核函数的关键步骤是将每个线程从它的线程索引映射到全局线性内存索引中，如图2-12所示。
接下来，每个维度下的矩阵大小可以按如下方法设置为16 384个元素：
然后，使用一个二维网格和二维块按如下方法设置核函数的执行配置：
把所有的代码整合到名为sumMatrixOnGPU-2D-grid-2D-block.cu的文件中。主函数代码如代码清单2-7所示。
用以下命令编译并运行该代码：
在Tesla M2070上运行的结果：
接下来，调整块的尺寸为32×16并重新编译和运行该代码。核函数的执行速度几乎快了两倍：
你可能好奇为什么只是改变了执行配置，内核性能就几乎翻了一倍。直观地说，你可能会觉得这是因为第二次配置的线程块数是第一次配置块数的两倍，所以并行性也是两倍。你的直觉是正确的，但是，如果进一步减小块的大小变为16×16，相比第一次配置你已经将块的数量翻了四倍。如下所示，这种配置的结果比第一个好但是不如第二个。
表2-3总结了不同执行配置的性能。结果显示，增加块的数量不一定能提升内核性能。在第3章中，你将会学习到为什么不同的执行配置会影响核函数的性能。
2.3.3 使用一维网格和一维块对矩阵求和
为了使用一维网格和一维块，你需要写一个新的核函数，其中每个线程处理ny个数据元素，如图2-13所示。
由于在新的核函数中每个线程都要处理ny个元素，与使用二维网格和二维块的矩阵求和的核函数相比，从线程和块索引到全局线性内存索引的映射都将会有很大不同。由于在这个核函数启动中使用了一个一维块布局，因此只有threadIdx.x是有用的，并且使用内核中的一个循环来处理每个线程中的ny个元素。
一维网格和块的配置如下：
使用以下配置调用核函数：
使用一维网格和一维块的更改替换代码清单2-7中的部分，并保存到文件sumMatrix-OnGPU-1D-grid-1D-block.cu中，使用以下命令编译并运行该程序：
结果显示，与使用一个二维网格和块（32×32）的配置结果相比，两者的性能基本相同。
接下来，按如下所示的方法增加块的大小：
重新编译并运行，可以看出核函数运行得更快了。
2.3.4 使用二维网格和一维块对矩阵求和
当使用一个包含一维块的二维网格时，每个线程都只关注一个数据元素并且网格的第二个维数等于ny，如图2-14所示。
这可以看作是含有一个二维块的二维网格的特殊情况，其中块的第二个维数是1。因此，从块和线程索引到矩阵坐标的映射就变成：
从矩阵坐标到全局线性内存偏移量的映射保持不变。新的核函数如下：
注意，二维核函数sumMatrixOnGPU2D也为这个执行配置工作。编写新内核的唯一优点是每个线程省去了一次整数乘法和一次整数加法的运算。
将块尺寸设置为32，并在此基础上计算网格大小：
如下所示调用内核：
对代码清单2-7进行更改替换，并将替换后的程序保存到名为sumMatrixOnGPU-2D-grid-1D-block.cu的文件中，然后使用以下命令编译并运行。
运行结果为：
如下所示，将线程块的大小增加到256：
然后重新编译运行，系统会表现出目前为止最佳的性能（见表2-4）：
从矩阵加法的例子中可以看出：
- 改变执行配置对内核性能有影响
- 传统的核函数实现一般不能获得最佳性能
- 对于一个给定的核函数，尝试使用不同的网格和线程块大小可以获得更好的性能
在第3章，将会从硬件的角度学习产生这些问题的原因。
