《CUDA C编程权威指南》——2.3节组织并行线程

本节书摘来自华章社区《CUDA C编程权威指南》一书中的第2章，第2.3节组织并行线程，作者[美]　马克斯·格罗斯曼（Max Grossman）　，更多章节内容可以访问云栖社区“华章社区”公众号查看

2.3　组织并行线程
从前面的例子可以看出，如果使用了合适的网格和块大小来正确地组织线程，那么可以对内核性能产生很大的影响。在向量加法的例子中，为了实现最佳性能我们调整了块的大小，并基于块大小和向量数据大小计算出了网格大小。
现在通过一个矩阵加法的例子来进一步说明这一点。对于矩阵运算，传统的方法是在内核中使用一个包含二维网格与二维块的布局来组织线程。但是，这种传统的方法无法获得最佳性能。在矩阵加法中使用以下布局将有助于了解更多关于网格和块的启发性的用法：
由二维线程块构成的二维网格
由一维线程块构成的一维网格
由一维线程块构成的二维网格
2.3.1　使用块和线程建立矩阵索引
通常情况下，一个矩阵用行优先的方法在全局内存中进行线性存储。图2-9所示的是一个8×6矩阵的小例子。
在一个矩阵加法核函数中，一个线程通常被分配一个数据元素来处理。首先要完成的任务是使用块和线程索引从全局内存中访问指定的数据。通常情况下，对一个二维示例来说，需要管理3种索引：

printThreadInfo函数被用于输出关于每个线程的以下信息：
线程索引
块索引
矩阵坐标
线性全局内存偏移量
相应元素的值
用以下命令编译并运行该程序：

2.3.2　使用二维网格和二维块对矩阵求和
在本节中，我们将使用一个二维网格和二维块来编写一个矩阵加法核函数。首先，应编写一个校验主函数以验证矩阵加法核函数是否能得出正确的结果：

然后，使用一个二维网格和二维块按如下方法设置核函数的执行配置：

接下来，调整块的尺寸为32×16并重新编译和运行该代码。核函数的执行速度几乎快了两倍：

你可能好奇为什么只是改变了执行配置，内核性能就几乎翻了一倍。直观地说，你可能会觉得这是因为第二次配置的线程块数是第一次配置块数的两倍，所以并行性也是两倍。你的直觉是正确的，但是，如果进一步减小块的大小变为16×16，相比第一次配置你已经将块的数量翻了四倍。如下所示，这种配置的结果比第一个好但是不如第二个。

表2-3总结了不同执行配置的性能。结果显示，增加块的数量不一定能提升内核性能。在第3章中，你将会学习到为什么不同的执行配置会影响核函数的性能。

由于在新的核函数中每个线程都要处理ny个元素，与使用二维网格和二维块的矩阵求和的核函数相比，从线程和块索引到全局线性内存索引的映射都将会有很大不同。由于在这个核函数启动中使用了一个一维块布局，因此只有threadIdx.x是有用的，并且使用内核中的一个循环来处理每个线程中的ny个元素。

一维网格和块的配置如下：

重新编译并运行，可以看出核函数运行得更快了。

2.3.4　使用二维网格和一维块对矩阵求和
当使用一个包含一维块的二维网格时，每个线程都只关注一个数据元素并且网格的第二个维数等于ny，如图2-14所示。
这可以看作是含有一个二维块的二维网格的特殊情况，其中块的第二个维数是1。因此，从块和线程索引到矩阵坐标的映射就变成：

从矩阵加法的例子中可以看出：
改变执行配置对内核性能有影响
传统的核函数实现一般不能获得最佳性能
对于一个给定的核函数，尝试使用不同的网格和线程块大小可以获得更好的性能
在第3章，将会从硬件的角度学习产生这些问题的原因。