【综述解析·III】Optimization Techniques for GPU Programming (推荐)(下)

本文是一篇 CSUR 长文，80 页，≈450 文献整理而成。作者来自阿姆斯特丹自由大学，QS209.
本文接续上一部分（正文+附录），这篇为附录部分。值得一提的是，附录也写的非常全面且很精彩。
Appendix A：不光解释了一些概念，还从实践而非理论的角度教读者如何进行操作，还有正文放不下的那些优化，这里都进一步进行了讲解。
Appendix B：告诉读者本综述筛选的视角和办法。（值得学习）
Appendix C：讲述了关于近些年 GPU 的发展和他们的论文数量，这个内容不多。
Appendix D：权衡因素，基本也能辅助读者快速入门 GPU 的性能考量了。

附录A：优化手段

本附录汇集了文献中提取的优化技术。为了便于参考与深入了解，这些优化技术按照第6节中的主题进行归类与组织，形成了系统化的参考资料。以下是所涵盖优化方法的概览：

A.1 内存访问

A.1.1 片上内存 —— 使用专用内存

对于大多数应用而言，GPU 通常具备比其内存系统更强的计算资源。因此，许多优化技术均围绕提高内存访问效率展开。与 CPU 类似，GPU 也提供缓存机制，尽管其层级较少（通常为两级，而 CPU 为三级）。此外，GPU 还提供映射到高速内存中的专用内存空间，包括常量内存、纹理或表面内存，以及共享内存。

常量内存：常量内存的典型最大容量为 64 KB。常量内存在核函数执行期间为只读数据，必须在核函数启动前写入，并声明为常量内存。对于需要广播到多个线程的数据，常量内存特别有用，因为这类访问只会产生一次内存事务。它对于需要在一系列核函数调用中保持可用的数据也非常有用。使用常量缓存时需考虑缓存冲突的问题，Lee 等人测量了常量缓存的命中率。Jang 等人通过在 Tesla GPU 上使用常量内存获得了 3 ~ 5×的加速。Xiao 等人在 Fermi 架构下的基因测序应用中，将查询序列和评分矩阵存储在常量内存中，报告了 12% 的加速比。`

注意事项：①常量内存不能用于写操作，只能由主机端写入；②如果不同线程访问不同常量地址，则缓存命中率下降；③不能用于动态数据结构或需要线程写入的数据；④跨多个 kernel 使用时，需要特别注意数据一致性和管理。
实际使用举例：①广播参数（如模型参数、卷积核）；②查找表（如打分矩阵、常量系数）…

纹理内存：纹理内存也是一种只读内存，其典型大小为 12 KB 到 128 KB。尽管纹理内存常常仅因其访问具有缓存机制而被使用，纹理内存针对具有二维空间局部性的访问进行了优化，多个实现利用了这种空间局部性，或者将其用于非合并访问。此外，还可以利用纹理读取自动处理边界值的特性，从而可能减少分支发散。图 10(a) 展示了一个二维的 9 元素纹理，使用 “wrap” 寻址模式，意味着越界读取时将返回一个值，仿佛纹理是环绕的；图中透明元素表示越界访问。图 10(b) 展示了相同纹理在 “clamp” 寻址模式下的行为，此时越界读取将重复边界元素。

一次纹理读取可以获取 R、G、B 和 A 四个通道的值，这一特性启发 Liu 等人采用相同的技术来打包字符串数据，并通过一次纹理读取来提取这些数据，用于基因序列比对。纹理读取通过特殊的 API 进行，寻址通常由专用硬件完成。纹理单元还可以自动执行整数到浮点数的转换（带归一化），这一功能在图形应用中非常有用，但也可被通用计算应用所利用。

在视频压缩相关的核函数中，Ryoo 等人报告了一个典型的性能提升案例，其中边界计算由纹理单元完成，在 Tesla GPU 上实现了 2.8 倍的加速。

一句话解释：纹理内存是 GPU 上专为“类似图像”数据访问优化的只读内存，它能让读取二维数据更快、更安全，还能一次读取多个值，甚至自动做边界处理和数据变换，不仅在图像处理，很多通用计算任务中也被利用。
举例：假设你有一张地图（二维数组），你想查某一块的地形类型，如果直接查数组，每次都要手动检查是否越界、是否要处理边界。但如果你把这张地图放进纹理内存的话，它会自动帮你处理边界；邻近区域的访问也会变快（一目了然）；一次可以查多个值（比如颜色的 R/G/B/A）；特别适合那种“滑动窗口”扫描的任务，比如图像处理或基因序列比对。

共享内存：要使用共享内存，可以声明一个变量（通常是数组）为“shared”。共享数组驻留在流处理器的片上内存中，该片上内存的典型大小为 16、48、64 或 96 KB。这些数据可以在线程块内部的线程之间共享。由于这种共享允许对内存的并行访问，因此需要同步，这通过每个线程块中的 warp 到达屏障（barriers）来实现，以便继续执行。由于屏障造成的 warp 停顿通常对性能不利，一个常用策略是减少单个线程块使用的共享内存量，从而允许多个线程块在一个流处理器上同时运行。但由于线程块之间是独立的，不需要互相同步，所以当一个线程块在某个屏障上停顿时，另一个线程块可以继续执行。

片上内存通常以与一个 warp 中线程数量相等的 bank 数量组织。当一个 warp 中的多个线程访问同一个元素时，这些访问会被串行化，从而降低性能，这被称为“共享内存 bank 冲突”。一个常用的策略是减少 bank 冲突的数量。

例子（bank 冲突）：假设有 32 个学生（线程）同时去图书馆借书（访问共享内存），图书馆有 32 个窗口（memory banks），如果每个学生去不同窗口，就可以同时完成，不排队。但如果很多学生都跑去 同一个窗口，就会排队处理，这就叫做 bank 冲突。
所以：

如果线程访问的是 相同 bank 的不同地址，也会排队；
- 如果访问 同一地址，有的架构能做“广播”，不一定慢；
- 如果访问模式乱、冲突严重，就会导致 执行变慢。

共享内存可以用于多种目的。首先，共享内存用于在线程间共享数据。另一个用途是减少昂贵的全局内存操作，例如允许对全局内存访问进行合并并将其存储在共享内存中，以便允许在共享内存中进行非合并访问，或完全减少不规则访问的影响。Satish 等人报告称，在 Tesla 架构（2008）上，通过在共享内存中对数据进行排序以替代对全局内存的分散写入，可实现两个数量级的加速。共享内存也可以用于减少同步，例如为每个线程块使用一个哈希表，稍后再将其合并为一个全局哈希表。

哈希表目的是保证每个线程块之间互不干扰，最后合并时只需要少量同步。

如果能够利用数据重用或时间局部性，共享内存的效果会更加显著，但对于空间局部性，共享内存同样有用。通过共享内存常用的一种具体操作模式是归约。一旦核函数被转化为使用共享内存，优化共享内存的使用就变得重要。典型的变换包括使用共享内存进行的分块，或融合多个算子以增加数据重用的机会（见 A.1.8 节）。通常共享内存、纹理内存和常量内存被视为数据重用中最快的内存，这在空间局部性方面是正确的，但对于时间局部性而言，寄存器可能是一种更快的内存。

Volkov 和 Demmel 探讨了带宽、指令级并行性与内存带宽之间的关系，并得出结论：在指令级并行性足够的情况下，例如通过激进的循环展开获得，使用较少数量的线程可能更有利，从而使寄存器而不是共享内存成为主要的片上内存。Swirydowicz 等人延续了这一思路，并基于共享内存构建了一个 Roofline 模型。这类实现通常具有较低的占用率，但能实现极高的性能。

为以这种方式使用寄存器，通常共享内存被用于为寄存器提供数据暂存，或通过共享内存在寄存器之间传递数据。相比之下，对于其他实现，有时需要使用共享内存来减少寄存器溢出。在另一些情况下，寄存器被用作共享内存的溢出位置。

共享内存的使用量限制了每个处理器上线程块的数量，因此通常需要减少共享内存的使用，以允许更多的独立线程从而减少屏障同步的代价。例如，这可以通过循环分裂、重计算值或使用寄存器来实现。

另一项提高共享内存性能的重要技术是减少 bank 冲突。为此有多种技术可用。常用的方法是通过填充偏移线程的访问位置，使其访问不同的 bank。例如，Man 等人展示了图 11，其中他们为一个 32 × 32 的数组的每一行增加一个额外元素的填充，从而当线程块的 32 个线程访问数组的某一列（元素 0、32、64 等）时，线程不会访问相同的 bank，而是访问相邻的 bank，从而实现并行访问。

重新排序数据是另一种技术，Govindaraju 等人通过将结构体数组转换为数组结构的方式实现了这一点。也可以在保持相同数据布局的情况下重新映射线程到数据，或者执行相关的循环置换。Reddy 等人利用操作的交换性来重排数据访问。Bernstein 等人设计了一个搜索工具，为线程分配数据，从而几乎完全消除 bank 冲突。

所以要在线程数量（活跃度） & 单线程资源使用（如寄存器和共享内存） 中做出权衡，
即因为 共享内存、纹理内存、常量内存适合空间局部性、寄存器则更适合时间局部性；
如果每个线程使用更多的寄存器，通过高指令并行度优化执行效率，即便线程少也能获得更好性能

A.1.2 片上内存 —— 使用 Warp 级函数

一般而言，并且从历史上看，GPU 并不支持线程块之间的同步（参见第 A.3.10 节获取更详细的讨论），而是通过共享内存和同步屏障支持线程块内部线程之间的通信与同步。然而，NVIDIA 在 Fermi 架构（2010）中引入了 warp 投票函数（warp-voting functions），允许同一个 warp 中的线程无需显式同步即可达成一致。从 Kepler 架构（2012）开始，NVIDIA 对这些 warp 投票函数进行了扩展，加入了 warp shuffle 函数，从而使得 warp 级函数更为强大。由于能够带来性能提升，使用这些 warp 函数被视为一种优化手段。

shuffle 函数允许同一个 warp 中的线程直接在寄存器中交换数据，无需先将其存入共享内存，也无需在线程块级别进行显式同步。在图 12 中，Wang 等人展示了 Kepler 架构（2012）中引入的 shuffle 函数的若干变体。使用这一优化的优势是多方面的：首先，访问寄存器的延迟低于访问共享内存，带宽也更高；其次，所需分配的共享内存更少；第三，由于无需显式同步，通信开销也更低。基于这些优势，文献中最常见的 warp shuffle 函数使用场景是实现 intra-warp 的数据共享而无需借助共享内存。

尤其值得关注的是 Barnat 等人对这些函数的使用，他们利用这些函数实现了一种 warp 级别的缓存机制，从而减少对全局内存的访问次数并节省内存带宽。也有一些研究尝试将 warp shuffle 函数与基于共享内存的传统 intra-warp 通信机制进行比较。Wang 等人开发了微基准测试用于比较这两种方法，并得出结论：使用该优化可在两类序列比对算法（Kepler 2012）中带来显著性能提升，介于 20% ~200% 之间。Mawson 等人也表明，在其 Boltzmann 求解器中使用 warp shuffle 函数的速度略快于使用共享内存。尽管常常如此，但也存在此优化无法带来性能收益的情况，例如 Abdelrahman 等人使用其密码学微基准测试时所记录的结果。

warp shuffle 函数的另一个特殊用途是在实现协作算法，例如归约和前缀和。Wang 等人在图 13 中展示了如何使用 shuffle 指令实现归约操作。这类算法中该优化的主要优势是同步开销更小。warp 投票函数则用于在 warp 级别对谓词进行求值，并将求值结果广播给 warp 中的所有线程。通过这种方式，开发者可以测试 warp 中是否所有、是否有任一、或者具体哪些线程满足给定谓词，而无需它们之间进行显式通信。例如，Pai 等人在图处理任务中使用该优化判断哪个线程应执行算法中的 push 操作；Ashkiani 等人也使用相同的优化来计算 warp 级别的直方图。

①warp shuffle 是由硬件直接支持的功能，直接使用即可
②谓词就是一个布尔表达式，也就是判断某个条件是否成立的逻辑语句
③所谓的投票函数就是 __all (predicate), __any(predicate), __ballot(predicate) 等等，用来表示哪些线程满足某些条件

A.1.3 片上内存 —— 寄存器分块

寄存器是 GPU 上可用的最快速的内存，因此，通过将常用值存储在寄存器中，可以极大地提升性能。由于寄存器在线程之间不可共享，因此它们主要适用于具有时间局部性的重用数据，而不太适用于空间局部性，尽管有些方法使用 warp-shuffle 在同一 warp 的线程间分发值。由于这些限制，寄存器分块（register blocking，有时也称为时间分块 temporal blocking）与循环展开密切相关，其解释为“寄存器分块本质上是在 X、Y 或 Z 维度上的 unroll and jam”。它还与“每线程不同的工作量”（见 A.3.4 节）相关，因为有时同一线程会对多个元素重用相同的数据。Hong 等人确实指出，每线程不同的工作量可以实现寄存器分块。

寄存器分块通常与空间分块（见 A.1.7 节）结合使用，并有时被称为“双层分块”。从文献来看，寄存器分块在三维模板计算（3D stencil computations）中尤为常见。一种常用方法是：在 XY 平面使用共享内存，而在 Z 维度上串行计算，将 Z 方向上重用的数据保留在寄存器中。Mo 和 Li 采用了类似的技术，但他们重用的是邻近点的部分和，而不是 Z 方向上的元素。Matsumura 等人提供了关于时间分块及其与空间分块关系的全面综述，并提出了一个框架，用于在模板计算中实现高度的时间分块。

寄存器分块（register blocking）就是让每个线程多做一点事，把需要反复用的数据放在最快的寄存器里，而不是每次都去慢速内存里取。这样程序能跑得更快。

Daga 等人指出，由于寄存器中的值不易与其他线程共享，因此寄存器特别适合用于在 GPU 上存放累加器值，Hagedorn 等人也有类似的观察。这种用法在许多进行矩阵乘法的研究中尤为明显，其中输出矩阵的元素在寄存器中进行累加，输入矩阵的块可能也在寄存器中被分块存储。Remmelg 等人在图 14 中比较了寄存器分块（图 14(b)）与空间分块（图 14(a)，见 A.1.7 节）。在左侧，输出元素是根据 A 和 B 矩阵的行和列中的 tiles 计算得出的；在右侧，B 矩阵的列存储在寄存器中，并在 A 矩阵的一行中为每个高亮列重用，以累加输出元素，这些输出元素也存储在寄存器中。

Mukonoki 等人进行了矩阵/向量乘法，并利用向量指令实现了寄存器分块。在相同应用中，Xu 等人精心地将寄存器分块与合并内存访问结合，并考虑了为其他 warp 加载到寄存器中的值通过缓存进行重用的情况。

由于寄存器是 GPU 上最快的内存，因此它们可以提供显著的加速效果。Tang 等人展示了在 Volta 架构（2017）上，其图比较算法中根据 block 大小可获得 2 ~ 7 倍的加速比。Tan 等人指出，寄存器无法通过索引访问，因此必须使用宏展开或模板等技术，以便能够控制寄存器数组。Bernstein 等人使用了一种有趣的技术，将寄存器替代常量内存的功能。此外，他们还通过打包四个值进一个寄存器并利用位移操作进行计算，从而以更快的内存访问速度换取一定的计算开销。

A.1.4 片上内存 —— 减少寄存器使用

在编译 GPU 核函数时，编译器为每个线程块分配一定数量的寄存器用于保存中间值。由于每个流处理器（SM）中的寄存器数量是固定的，使用过多的寄存器可能会限制每个 SM 上可并发执行的线程块数量（即占用率）。此外，寄存器使用量的增加可能导致寄存器溢出（register spilling），即线程的局部变量被存储到较慢的片外设备内存中，尽管这些溢出变量通常会被缓存。因此，在某些情况下，减少寄存器使用量可以减少寄存器溢出，从而带来更高的并发性和性能提升。

一种显而易见的减少寄存器使用的方法是重写核函数的源代码。例如，有研究在一个三维模板应用中实验发现，在 CUDA 1.1 中，每次对多维数组的索引操作都会消耗一个寄存器。在后续的工作中，他们通过使用指针运算手动编写索引操作、最小化临时变量和重写算术操作，将寄存器使用量从 100 降低到 32。

另有研究针对 Lattice Boltzmann 模拟，使用多种技术将寄存器使用量从 70 减少到 40，包括将某些临时变量移入共享内存、将小类型数据打包为大类型（例如将 4 字节打包为一个整数）以及避免存储可重新计算的值。通过这些技术，在 Kepler GPU（2012）上获得了 7% 到 9% 的性能提升。

将小数据类型打包成较大的数据类型（如 4 字节打包为一个整数）的例子：4 和 char 需要四个寄存器，而合并为一个 int 只需要一个寄存器

小循环的展开也可以节省寄存器，因为它消除了对归纳变量的需求（见 A.2.1 节）。另一个节省寄存器的方法是更改算法。例如，在 GPU 上执行 SHA1 计算时，有研究使用的算法需要将 80 个预计算值存储在寄存器中，切换到另一种计算方案后节省了 64 个寄存器。还有研究在模板操作中从多对一方案切换到一对多方案，从而将寄存器使用量减少了 40%。

**Stencil（模板计算）** 是 GPU 并行编程中一个非常常见、重要的概念。简单例子就是假设你有一个一维数组：`A = [1, 2, 3, 4, 5]`，你想生成一个新数组 B，使得每个元素是它**左边、自己、右边**的平均值（边界先不考虑）：`B[1] = (A[0] + A[1] + A[2]) / 3 = (1+2+3)/3 = 2`, `B[2] = (A[1] + A[2] + A[3]) / 3 = (2+3+4)/3 = 3`，这就是一个**典型的 stencil 操作**，每个输出值依赖一个“局部模板”，即：当前点 ±1 的邻域。
那么**2D 和 3D stencil 更常见**，例如对一个像素 `[i][j]` 作如下操作 `B[i][j] = (A[i-1][j] + A[i+1][j] + A[i][j-1] + A[i][j+1] + A[i][j]) / 5`，在图像处理（模糊滤波）、流体模拟、热扩散模拟中很常见。

最后，还可以强制编译器使用更少的寄存器。有研究使用 CUDA 编译器的 maxrregcount 选项强制进行寄存器溢出，并发现这对某些基准程序是有利的。另有研究采取更为激进的方法：由于对 NVIDIA 的寄存器分配器不满意，使用反向工程得到的 CUDA 汇编器来执行自定义的寄存器分配。然而，需要注意的是，这项工作是在 2010 年进行的，之后的寄存器分配器可能已经有所改进。

GPU 上寄存器分配的论文确实挺罕见的，依稀只记得一篇 Shobaki 的论文。

A.1.5 片上内存 —— 重计算（Recompute）

重计算（recompute）优化指的是重新计算一个或多个先前已被计算出的值，这些值可能是由其他计算单元计算出来的。通常在应用该优化时，会进行一些本不必要的计算工作，以此来避免通信（若这些计算原先是由其他线程或线程块完成的），或者减少内存操作或内存占用（若这些计算原先是由相同线程或线程块完成的）。这种优化是一种资源权衡策略，在现代 GPU 上通常是有利的，因为其计算能力远超内存和通信带宽。

例子：在训练大型神经网络时，前向传播时每一层的激活值都会被缓存，以便在反向传播中计算梯度时使用。但这些激活值如果都保存在显存中，内存占用会非常高，尤其是在模型很深或 batch 很大时。与其保存每一层的激活值，不如在需要它们的时候重新计算它们。

文献中存在多个使用重计算以减少通信的示例。特别地，Holewinski 等人选择在其模板计算中为每个 tile 的 halo 重新计算，从而执行了额外的冗余计算工作，以避免线程块之间复杂的同步机制。该方法在不同架构（Tesla GPU（2008）、Fermi（2010）GPU，以及 AMD BeaverCreek（2011）GPU）上，根据计算指令与额外指令的比例，最高可获得 2.5 倍的加速。在高度计算密集的模板计算中，冗余计算无法与其他操作重叠执行，因此性能与原始核函数相当。

Li 等人的案例有所不同，其使用重计算的目标是减少主机与 GPU 之间的通信量。在该研究中，作者并未为确保所有值正确而进行主机与 GPU 之间的通信，而是接受 GPU 上因冲突而产生部分错误值的事实，并选择在主机端对这些值重新计算。

Lefebvre 等人的工作尤其值得注意，他们将算子融合（见 A.1.8 节）与重计算结合，减少了通过全局内存进行的昂贵同步操作，从而避免了中间结果的存储。使用重计算的一个经典示例是减少对全局内存的访问次数和读取数据的总量。Domínguez 等人选择仅存储那些无法轻易重计算的值，而不是存储所有的值。该优化也可用于减少共享内存需求，如 Lee 等人所展示。

此外，还有一些研究测试了该优化的效果。Lee 等人（与上述 Lee 等人为不同作者）实现了一个三维图像配准核函数（一个计算密集型算法），通过将该方法与存储中间结果以便重用的版本进行对比，以测试重计算优化的效果；他们的实验结果表明，重计算并不总是性能上的最优解。类似地，Střelák 等人对其 GPU 核函数的不同实现进行了比较，采用此优化在运行时重计算部分权重而不是存储这些权重；他们的结论是，对于具有更高浮点计算能力而内存带宽较低的架构，重计算可能更为有利。

A.1.6 片外内存 —— 合并访问（Coalesced Access）

合并访问或许是应用最广泛的一种优化。当一个 warp（32 个线程）的联合内存访问满足一组合并规则时，数据可以通过少于 32 次的设备内存事务（transaction）完成访问。这些合并规则随体系结构而异，并且随着时间推移逐渐变得不那么严格。一个典型的合并访问示例是：一个 warp 访问连续的、对齐到 128 字节的 32 位内存位置（128 字节是典型的缓存行大小）。如果这 128 字节已被缓存，则只需一次内存事务；否则，则需要四次 32 字节的设备内存事务。

图 15 展示了一个源自 CUDA Programming Guide 的示例：上半部分显示线程 0–31 以合并方式访问地址 128 到 255，如果数据未被缓存，将产生 4 次设备内存事务；如果数据已缓存，则只需 1 次缓存事务。下半部分，线程 0–31 访问地址 129 到 256，如果数据未缓存，将导致 5 次设备内存事务；如果数据已缓存，则为 2 次缓存事务。

有多种方法可以实现合并访问，其实现复杂度取决于多个因素，例如数据布局、线程组织方式以及共享内存的可用性。如果数据布局允许，一种简单的合并访问方法是将线程组织成每个线程对齐其所操作的数据元素。其他实现线程与数据对齐的方式包括选择合适的线程块大小、使用 tiling，或选择不同的并行化策略。

对于复杂的数据布局，也可以利用 GPU 的计算能力执行复杂的索引操作（例如空间填充曲线），以实现线程与数据的对齐。如果某个 kernel 表现出非合并的内存访问，一个常用方法是先以合并方式从全局内存加载数据到共享内存，再由共享内存进行非合并访问。这种方式还可以与 tiling 结合，将一个 tile 加载到共享内存中。另一种方式不是将数据加载到共享内存，而是将数据先暂存在共享内存中，然后再写回设备内存。有些研究同时采用了这两种方法。

上述技术保持了全局内存中的数据布局。然而，如果应用允许，另一种常用方法是重新组织数据，使线程以合并方式访问数据。除了以自定义方式重组数据之外，还有一些常见的策略。例如，从结构体数组（array of structs）转为数组结构体（struct of arrays），并可能结合 tiling 使用。另一种简单的转换是从行主序数组切换到列主序数组，或采用其他形式的转置。最后一种常见的技术是在数据中添加填充（padding）以实现对齐。Gómez-Luna 等人在其早期工作的基础上提出了一种最小填充策略，应用于就地矩阵转置 kernel。该新方法的吞吐量更高，并包括一个填充和去填充的 kernel。

Ito 与 Nakano 针对“最优多边形三角剖分”问题提出了两种不同的数据布局方式，分别在算法的不同时期表现出最优特性。他们指出，在适当时机将数据从一种布局转换为另一种布局（包括转换开销在内）相比于始终使用单一布局具有更高的整体性能。

在某些硬件架构中，合并读取与合并写入的吞吐能力存在差异，因此在设计数据布局时需在读取与写入的合并效率之间做出权衡。Lin 等人提出一种新的数据布局方式，与现有工作相比，能够同时实现读取与写入操作的合并访问。需要指出的是，在具有高度不规则访问模式的应用中（如稀疏矩阵-向量乘法与图处理），实现合并访问尤为困难。为应对此问题，可采用对代价高昂的随机访问进行缓存的方法，但如 A.2.3 节所示，也有大量稀疏格式的研究工作旨在提高访问模式的规律性。例如，Bell 与 Garland 展示了多种稀疏矩阵-向量乘法格式，这些格式均旨在提升访问合并性。Zhong 等人提出一种新型图数据结构，通过布局重组以支持合并访问。

对于稀疏数据结构，填充数据是一种常见手段，用于增强对齐性与合并性。Kreutzer 等人则提出了一种新的稀疏矩阵——向量格式，在降低填充开销的同时仍能实现高效的合并访问。

最后，另一类优化策略是直接重新设计数据结构以提升合并效率。例如，Ashkiani 等人提出一种支持更新操作的哈希表，其在设计上充分考虑合并访问的需求。该哈希表采用链式结构，每个链表节点（slab）包含多个键值对与一个指针，支持使用 warp 指令实现操作，并确保内存访问的合并性。

A.1.7 片外内存 —— 循环分块或空间分块（Loop Tiling or Spatial Blocking）

分块（blocking）或循环分块（loop tiling）是文献中广泛使用的优化术语，二者虽在表述上有所不同，但本质均指将计算或数据划分为多个小块，逐块处理的优化策略。

PS：现有翻译工具容易把 blocking 单独翻译为“阻塞”

Deftu 与 Murarasu 将循环分块定义为“通过将迭代空间划分为 tile，从而提升缓存复用率的循环重排转换”；而 Nath 等人将分块视为一种密集线性代数中的优化技术，其通过在原始矩阵的子矩阵上执行计算来提升效率。为消除术语混淆，本文统一使用“空间分块”（spatial blocking）一词，以区别于时间分块（详见 A.1.8 节）。图16左侧显示了以规则方式访问的数据。右侧则应用了空间分块技术，将数据划分为四个块。

在 GPU 编程中，空间分块尤为重要，它使程序员能够控制数据在何处、如何被使用与重用。通过该策略，GPU 核心可更有效地利用 L1/L2 缓存、纹理缓存、寄存器与共享内存，从而提高数据局部性与整体性能。空间分块还可带来其他优化效益。例如，Střelák 等人利用分块技术降低了 warp 发散程度（参见 A.2.2 节）；Stratton 等人报告称，对于规则应用，可获得 3–6 倍加速；即便在无缓存的架构中（如 Tesla 2008），其性能提升亦达到 12%。随着现代 GPU 普遍配备缓存机制，该优化在新架构中同样重要。

在应用空间分块前，可能需先对特定变量的数据布局进行调整。Nath 等人提出一种递归分块策略，适用于稠密矩阵-向量乘法核函数，并结合线程重组以实现局部计算。Rojek 等人引入“2.5D 分块”术语，用于描述在对流输运算法中应用的一类特殊分块形式。

分块维度通常与线程块的维度相关联，此外还受到其他因素限制，如向量数据类型的向量宽度。当线程块维度与 tile/block 尺寸解耦时，程序员可灵活调整每个线程及线程块所承担的计算量，这被视为另一种优化策略（参见 A.3.4 节）。在确定线程块与 tile 大小时，自动调优技术（参见 A.3.6 节）被广泛应用，以寻找最优参数组合，从而提升程序性能。

Auto-tuning（自动调优）是一种根据目标硬件和任务自动选择最佳参数配置的技术，用于优化程序性能。它通常通过在多个候选配置中自动试验和评估，选择出在运行时间、内存占用或吞吐量等方面最优的方案。在深度学习、编译器优化（如 TVM、TensorRT）或数值计算中非常常见，尤其适用于手动调优困难或搜索空间复杂的场景。

A.1.8 片外内存 —— 算子融合（Kernel Fusion）

GPU 核函数（kernel）是以特定线程与线程块配置在设备上启动执行的函数。算子融合（kernel fusion），有时亦称为算子合并（kernel merge）或算子统一（kernel unification），在概念上类似于循环融合（loop fusion），后者被定义为“合并连续且具有依赖关系的循环”。然而，核函数与循环之间存在关键差异，即核函数可在设备内存层面支持同步机制。若同一线程块内的线程需要同步，可通过共享内存与同步屏障实现；若是不同线程块之间的同步，仅能通过多次核函数启动来实现，尽管已有技术尝试实现无需重启核函数的全局同步机制（详见 A.3.10 节）。

尽管许多优化策略侧重于核函数级别的单独优化，研究表明，仅优化每个 Kernel并不总能达到全局最优配置，而通过算子融合可获得更优的执行时间。算子融合带来性能提升的原因多样，尽管部分文献未明确列出原因。其中广泛认可的一点是：若两个核函数共享相同数据，通过融合可避免冗余的全局内存访问，即仅需一次数据加载。此类中间数据的传递可通过共享内存或寄存器实现。以寄存器为媒介传递数据不仅减少内存访问次数，还可提升数据局部性与缓存命中率，从而增强重用效率。在图17中，Hou 等人展示了他们的融合核函数（右侧）如何节省了 10n 次访问。他们通过寄存器传递数据。

除减少全局内存访问外，算子融合还可带来以下益处：提升缓存效益、增强数据局部性与重用性、减少核函数启动开销、提升并发度与指令级并行性。后者意味着每个线程执行更多（可能彼此独立的）指令，从而在隐藏内存访问延迟方面表现更优。

值得注意的是，算子融合亦可能引入负面影响。例如，融合后的核函数可能因需要更多共享内存或寄存器而导致占用率（occupancy）下降。此外，两个核函数之间的并行度若不匹配，可能对整体性能产生抑制。针对这一问题，研究提出“算子可融合性（kernel fusibility）”的概念，指出：若中间数据通过寄存器传递，则需保持线程映射一致；若通过共享内存传递，则需线程块映射一致并能在块级完成同步。若融合的两个算子需全局同步，则通常不具备融合条件。

如果两个算子之间需要所有线程协同完成一次计算再继续（即全局同步），你没法把它们安全地融合进一个 kernel 中。

算子融合还可减少主机与设备之间的数据传输。此外，融合能够提升并发性与指令级并行性，并可进一步隐藏内存访问延迟。然而，若融合核函数处理的数据体积过大，可能导致较差的数据局部性，因为数据可能分布于不同的内存区域。

关于算子融合的具体实现方法，有三种主要策略：

线程内融合（Inner Thread Fusion）：每个线程合并多个算子的操作，是最常见的融合形式，要求两个核函数的线程块与网格尺寸一致。
线程块内融合（Inner Thread Block Fusion）：在融合后的核函数中，依据线程索引决定执行哪一个原始核函数，适用于彼此独立、无线程块同步的核函数。
线程块间融合（Inter Thread Block Fusion）：类似前者，但依据的是线程块索引进行区分。

此外，还有研究提出基于 warp 的融合选择策略，并强调融合过程中寄存器使用的重要性。然而，融合两个独立算子并不一定导致寄存器使用数量的简单相加。某些情况下，对于独立核函数而言，现代 GPU 所支持的并发核函数执行（concurrent kernel execution）可能较融合策略更具性能优势。

“融合两个独立算子并不一定导致寄存器使用数量的简单相加” 的根本原因在于寄存器生命周期不重叠、融合后代码优化机会更多、融合降低了中间值 materialization 的需要

Carabaño 等人根据核函数间共享数据的类型对算子融合进行了分类。他们指出，若多个核函数共享相同的输入数据，则可进行所谓的“扁平融合（flat fusion）”；若多个核函数之间存在生产者—消费者（producer/consumer）关系，则可采用“管道融合（pipe fusion）”。Yi 等人也提出了类似的区分方法，尽管未使用上述命名。

注：扁平融合可以理解为如下：

Input --> Kernel A
      --> Kernel B
      --> Kernel C
优化为：
for each i:
    tmp = input[i]
    out1[i] = f(tmp)     // Kernel A
    out2[i] = g(tmp)     // Kernel B
    out3[i] = h(tmp)     // Kernel C

管道融合则是指前后有依赖关系的：

Input --> Kernel A --> Kernel B --> Kernel C
优化为：
for each i:
    tmp1 = f(input[i])     // Kernel A
    tmp2 = g(tmp1)         // Kernel B
    out[i] = h(tmp2)       // Kernel C

如同多数优化策略，算子融合亦涉及性能权衡。例如，Lefebvre 等人将融合与重计算相结合以减少同步开销，但同时引入了冗余计算，若不加控制可能导致性能下降。Chen 等人借鉴 Xiao 等人的全局同步机制以实现融合，但因此不得不降低每个线程块中的线程数量，从而影响并发度。融合过程中的关键资源包括寄存器与共享内存，二者使用之间需谨慎权衡。Biferale 等人针对 stencil 计算进行融合优化，在边界区域执行未融合版本，而在非边界区域使用融合实现，以兼顾性能与正确性。Yan 等人针对前缀和（prefix sum）算法进行依赖分析，指出传统上需三次核函数调用的实现可被单核函数实现所取代。该优化通过将部分步骤串行化换取在另一维度上的更高并行性，为此他们提出采用自动调优策略以平衡并发度与串行计算开销。在图算法的上下文中，Liu 等人报告称，在 Kepler 架构（2012）下，算子融合可带来 10% 至 74% 的性能提升，具体取决于应用类型与融合策略。然而，他们也指出，在个别情况下融合可能导致性能下降（如某一案例中下降了 13%）。总体而言，对于内存密集型应用，算子融合表现出一致的性能提升趋势。

A.1.9 片外内存 —— 软件预取（Software Prefetching）

数据预取是一种通用的优化策略，其基本思想是在数据真正被使用之前提前加载。GPU 在架构设计上已具备一定的预取能力，具体表现为当某一 warp 因内存操作而停滞时，调度器会立即切换至其他活跃 warp。然而，为了更细致地控制全局内存访问延迟的隐藏，本文探讨通过软件手段实现的预取技术。该技术的核心在于：通过重写核函数代码，使得全局内存读取操作被提早调度，并将其结果暂存于临时寄存器中，即便这些数据尚未立即被使用。这样，调度器便可利用该时间窗口将内存访问延迟隐藏于后续计算之中。此策略的代价在于寄存器使用量的增加，从而可能影响线程占用率（occupancy）。

软件预取已被广泛应用于多种稠密线性代数核函数中，如矩阵——向量乘法、矩阵——矩阵乘法与 stencil 操作。在这类核函数中，输入数据通常以 tile 方式划分，线程块在循环中逐个 tile 处理数据。预取的实现方式包括：在处理当前 tile 的同时提前加载下一 tile 的数据。图 18 展示了Ryoo 等人在他们的矩阵乘法核函数中使用了预取技术。

这种机制有时被称为“双缓冲”（double buffering），因其需使用两个缓冲区分别存储当前 tile 与下一 tile 的数据；也有文献称之为“软件流水线”（software pipelining）。

预取（Prefetching）技术也被应用于其他类型的应用中。 例如，Yang 等人提出了一个框架，可以通过代码变换自动优化 GPU 核函数，其中一种变换方式就是预取变换。Bauer 等人提出了一种编程方法，在该方法中，不同的 warp 执行不同的角色（即 warp 专职化），其中一个 warp 的角色是为其他 warp 预取数据。Wu 等人探讨了使用 inter-block barrier 同步机制相比于多次调用同一个 kernel 的有效性（见附录 A.3.10）。他们在 inter-block barrier 之前预取 barrier 之后需要的数据，而这种预取方式在单独 kernel 调用的情况下是无法实现的。利用这一技术，他们成功提升了应用性能。

然而，预取带来的性能提升是有差异的：一些研究者声称看到了性能提升，而另一些则观察到性能下降。这种差异的原因在于，预取需要额外的寄存器来保存预取的数据，这反过来可能会降低线程占用率（occupancy）。例如，Anh 等人在 Tesla K20Xm 上对稀疏矩阵乘法的测试中看到了 20%–40% 的性能提升。Ma 等人在 GTX280 上对稠密矩阵乘法测试中，预取两个值可提升 8.5% 的性能，但当预取更多值时，性能反而下降了 62%。Yang 等人在用 10 个 kernel 测试他们的优化框架时发现，预取变换在整个测试集上并没有带来任何额外的性能提升，因此他们在后续的框架版本中甚至禁用了这项变换。

A.1.10 片外内存 —— 数据压缩（Compress Data）

主机中的内存比 GPU 上更为充裕，因此有许多优化技术致力于减少应用的内存占用，其中一种优化策略就是数据压缩。压缩可用于减小输入、输出，甚至中间结果的大小，这种尺寸上的缩减有助于降低存储需求和数据传输开销。

Samadi 等人在 Figure 19 中展示了他们如何通过压缩输入数据来减少内存传输次数。该优化广泛应用于稀疏矩阵，研究人员提出了多种压缩方案，以减小非零元素索引所占的空间。这类压缩不仅有助于降低通信开销和内存占用，还可以减少预处理时间，如 Neelima 等人所示。

此外，减少变量所使用的位数的做法并不仅限于稀疏矩阵。例如在序列比对中的最长公共子序列问题以及正则表达式匹配中，也使用了位宽压缩策略，以提升整体效率和节省内存。

在稀疏矩阵处理中，数据压缩不仅仅体现在减少索引表示所需的比特数。例如，Liu 与 Vinter 提出的工作中通过去除重复索引来节省内存；Nisa 等人则在某些切片仅包含单个非零值的情况下避免使用指针，从而降低存储开销。此外，Liu 等人（与前述 Liu 无关）将压缩技术应用于非确定有限状态机中的状态表示。由于该应用属于不规则计算，其压缩性能取决于状态的可压缩性及压缩带来的性能提升是否能够抵消其开销。在某些情况下，压缩所引入的开销未能带来净收益，表现为最多 20% 的性能下降；而在另一些情况下，压缩带来了 27% 至接近 5 倍的加速效果（在 Pascal 架构 GPU 上）。

A.1.11 片外内存 —— 预计算（Precompute）

预计算是一种优化策略，其核心思想是在主计算阶段开始之前，预先完成部分计算任务，并在后续过程中多次复用这些预先生成的值。该策略常被视为时间与空间之间的权衡：通过重用已存储的中间结果可节省运行时间，但同时也带来了额外的内存压力。在此意义上，预计算可被视为与 A.1.5 节所述的重计算策略的对偶操作。然而，预计算的应用场景并不仅限于提升时间效率。在某些情况下，预计算也用于将计算任务迁移至更适合的硬件设备上，即便这些预计算值仅被使用一次。例如，Fortin 等人将部分计算移至 CPU 执行，以实现资源的高效利用。预计算还可用于改变计算结构。例如，Greathouse 等人通过预先计算能够装入共享内存的矩阵行数，优化了后续数据布局与访问策略。

Zhou 等人将预计算与纹理内存结合使用：在解决旅行商问题时，预先计算城市之间的距离矩阵并将其存入纹理内存中以提升访问性能。相关研究还展示了预计算与异构计算策略（参见 A.4.2 节）的协同作用，其中预计算阶段在主机（CPU）上完成，而主要计算过程在 GPU 上执行。这种计算划分不仅提高了资源利用率，也带来了显著的性能提升。例如，在某些密码恢复应用中，预计算所有线程共用的数据可实现高达 49% 的加速效果（基于 AMD Tahiti 架构 GPU）。总体而言，预计算还可被视为实现“减少冗余计算”（详见 A.2.5 节）的一种具体手段，通过提前完成共性或重复性计算，有效提升整体系统性能。

A.2 不规则性（Irregularity）

A.2.1 循环展开（Loop Unrolling）

循环展开是一种将循环体显式重复若干次的优化技术。展开可通过手动方式完成，即显式复制循环体的内容，或通过自动化方式实现，如使用宏、C++ 模板或编译器指令等手段。图 20 展示了一个简单代码片段的手动循环展开示例。

循环展开对性能提升具有多方面的积极影响。首先，它可减少与循环控制相关的指令开销，例如分支跳转与地址计算。其次，通过增加可同时执行的独立指令数，循环展开提升了指令级并行性（ILP），从而更有效地隐藏延迟。此外，循环展开还能够触发编译器进一步优化，例如将数组索引替换为编译期常量，以便将数组元素存储于寄存器中；移除依赖于循环变量的分支；以及实现指令的向量化处理。

循环展开在多个应用场景中被证实具有显著效益。例如，在对 tile 化矩阵乘法进行优化时，完全展开最内层循环在某些 tile 大小（如 16 × 16）下可获得约 25% 的性能提升，尽管 tile 尺寸的选择对性能的影响更为显著。有趣的是，对于某些 tile 尺寸（如 8 × 8），循环展开不再带来性能增益，可能由于寄存器使用量增加导致线程数受限。其他研究也展示了循环展开在不同应用中的效果：在卷积操作中提高 GPU 执行效率；在哈希函数中移除大型 switch 语句；在某些算法中通过展开减少条件判断次数；在归约树中通过完全展开降低线程间同步开销；在小数组访问中消除动态索引，使编译器得以将数组元素寄存化，从而实现显著的性能提升。

尽管循环展开通常能提升性能，但展开因子若设置过大亦可能导致性能下降。主要原因包括：代码体积膨胀可能导致超出指令缓存容量，进而降低指令访问效率；寄存器压力增大可能导致线程占用率下降，尤其在寄存器资源紧张的情况下。尽管部分研究表明展开可降低寄存器使用，但在实际应用中其效果往往依赖于具体核函数与硬件平台。因此，如何确定最佳展开因子成为研究重点之一，并被广泛认为需针对不同核函数与设备进行调优以实现最优性能。

A.2.2 减少分支发散

分支发散（Branch divergence，有时也称为路径发散）是 SIMD 或基于 warp 的架构（如 GPU）中的一个难题，因为一个 warp 内的所有线程都是锁步（lock-step）执行指令的。当某个分支条件依赖于线程索引时，warp 内不同线程可能需要走不同的分支路径。在 GPU 上，这会导致所有线程都必须执行两条分支——只是那些本不该执行某条分支的线程虽然执行了指令，却不会写回结果。这会严重影响性能，因为两条分支被串行地执行。若分支内的指令足够简单，就可以用谓词化指令（predicated instructions）来替代分支——只有谓词条件为真的线程才会将计算结果写回，从而避免发散带来的性能损失。

移除分支。
移除分支。第一种技术是完全消除分支，因为分支会带来额外的开销。在某些情况下，可以使用最大值和最小值函数来消除分支，或使用其他算术运算。图 21 展示了 Wang 等人的一个示例：顶部是带分支的原始代码，底部则用算术运算替代了分支。阴影框表示在 warp 中处于非活动状态的线程（条件为假），它们仍执行该指令，但结果不会写回。

为了说明减少分支发散的潜力与局限，Wang 等人在 2008 年的一块 AMD GPU 上，使用 Brook+ 语言对一个分支密集型核函数进行了优化，观察到 3.45× 的加速，但对其它核函数则无明显提升。Tran 等人在 2012 年的 Kepler GPU 上，将一个三维模板（stencil）计算的边界处分支去除后，对 Lattice Boltzmann 求解器实现了约 10% 的加速。

Chakroun 等人用正弦、余弦和指数函数（采用快速数学函数，见 A.3.2 节）替代分支，这种做法称为“分支重构”（branch refactoring）。Vespa 等人在所谓的“算法扁平化”（algorithm flattening）中将此技术发挥到极致，将每个分支都表达为可进一步化简和优化的算术计算。

Jiang 等人执行“指令流归一化”（instruction stream normalization），同样消除了分支：他们将条件编码到一个 4 位宽的特征向量中，用作常量内存中查找表的索引，为线程提供正确的值。Daga 等人则通过核函数裂变（kernel-fission，见 A.2.4 节）在 CPU 启动核函数之前，就根据可预计算的条件拆分核函数，从而移除分支。他们指出，在他们的测试中，这项优化对 AMD GPU 的效果更明显，因为其每个处理核心上的分支单元比 NVIDIA GPU 要少。

减少分支（Reducing Branches）
如果无法完全移除分支，那么次优的策略是尽量减少分支的数量。Cecilia 等人将蚁群优化（ACO）应用于旅行商问题，通过共享内存通知同一个线程块内的其他线程某个城市是否已被访问。他们之后进一步使用了如 shuffle 等 warp 指令来优化这一过程。还有一些研究采用循环展开技术，来减少核函数中的控制流。

在循环结构中，还可以采用不同的策略。Han 和 Abdelrahman 提出了一种称为“延迟迭代”的方法，该方法会延迟带有分支的循环迭代，使每次迭代仅执行大多数线程会采取的路径；当另一条路径成为大多数线程所需路径时，才在后续迭代中执行。他们使用 warp 投票函数来判断线程的多数路径。作者指出这种方法可能导致少数线程长时间被阻塞。Novak 的方法对此进行了改进，允许只有被阻塞的线程参与投票，从而牺牲部分效率换取线程进度。Zhang 等人使用了类似方法，在循环末尾判断条件，从而允许无任务的线程跳过某些迭代。Han 等人还提出一种基于循环的技术，通过合并嵌套循环，使 warp 中的某些线程能开始下一次迭代，而其他线程则继续执行当前内层循环代码。

分支的折中策略（Trading Off Branching）
另一种减少分支的方法是将分支发散与其他性能指标进行权衡。例如，Okada 等人在基因序列比对中，采用了让 warp 中线程执行相似路径的方式，虽然带来了冗余计算，但由于其代价小于分支开销，整体性能反而更优。Sartori 和 Kumar 提出了一种称为“分支牧集”（branch-herding）的技术，应用于对错误容忍度高的核函数中。他们使用 warp 投票函数确定大多数线程所选路径，并强制整个 warp 执行该路径。用户可以通过性能计数器控制允许的错误率，从而设定强制分支的阈值。他们也将类似策略用于内存合并优化。

Hong 等人针对动态应用（如图处理）提出一种新的编程方法，可区分 SIMD 执行和传统的串行 SISD 执行。在后者中，每个 warp 执行相同的指令和数据，虽然牺牲了一些并行性，但消除了分支发散。通过将 warp 拆分为更小的虚拟 warp，可以在分支发散与 warp 内串行执行之间取得平衡。

削弱分支影响（Reducing the Effect of Branches）
除了减少分支数量外，也可以尝试降低分支带来的性能影响。Han 和 Abdelrahman 提出了一种称为“分支分布”的技术，尽量缩减分支内的代码，将不必要放入分支的代码移到外部，代价是可能引入更多寄存器来存储临时结果。如果分支内的指令足够简单，编译器可能会将其转换为谓词指令，即仅对满足条件的线程写回结果。Lalami 和 El-Baz 也使用了类似的方法，仅在分支中使用寄存器，从而减少其他资源开销。

线程/数据重映射（Thread/Data Remapping）
上述技术主要是直接作用于分支本身，但也可以通过间接手段影响分支行为。本部分将讨论并行化策略和数据布局，首先聚焦于静态性较强的应用。

在模板计算（stencil computation）中，常见做法是将核心数据及其边界一起加载到共享内存中。与其通过条件判断来加载边界数据，一些方案选择将线程直接分配给边界区域，这些线程在后续阶段不参与计算，从而避免了分支。另一些研究中，采用不同策略，不将左右边界加载入共享内存，而是直接从全局内存读取对应值。还有一种方式是将所有潜在引起发散的数据加载操作集中到一个 warp 中执行，从而加速其余 warp 的执行。

在 GPU 上的归约操作中，避免分支发散常常是首要优化手段之一。可以通过间隔式访问（strided access）实现该目标，并可在不同代际的 GPU 上评估其效果。结果显示，较新的硬件（如 GTX 970）从中受益更大。

例如如下代码在warp 中引入了多条分支（if），会导致 warp 发散：

if (tid < 16) data[tid] += data[tid + 16];
if (tid < 8)  data[tid] += data[tid + 8];
if (tid < 4)  data[tid] += data[tid + 4];
if (tid < 2)  data[tid] += data[tid + 2];
if (tid < 1)  data[tid] += data[tid + 1];

但是使用间隔式访问 + warp shuffle

float sum = data[tid];

// 32 个线程的 warp 归约，无分支
for (int offset = 16; offset > 0; offset /= 2)
    sum += __shfl_down_sync(0xFFFFFFFF, sum, offset);

也有研究通过修改并行化方案以减少分支发散。例如，有研究提出用于高性能三对角矩阵求解器的动态 tiling 方法；也有工作在执行高精度整数乘法（1024 位）时采用 warp 同步编程，合理分配线程避免发散；还有研究在多边形最优三角剖分中重新设计线程到数据的映射方式，从而完全消除分支发散。在图计算场景中，通过每个 warp 处理一条边（而不是每个线程处理一条边）也可以有效减少发散。

在高度动态、结构不规则的应用中，静态分析失效，因此有研究提出线程——数据重映射（thread-data remapping）技术以减少分支发散。这一策略主要包含两方面机制：数据布局转换和访问重定向。前者通过复制原始数据以减小线程间发散；后者则为内存访问添加一层间接索引，从而引导线程访问更加一致的数据区域。考虑到数据布局转换本身代价较高，尤其在 CPU 上进行，因此引入流水线方式，允许在 CPU 修改数据布局的同时，GPU 执行已有核函数。此外，还提出算法以用引用重定向部分替代数据布局转换，从而进一步降低开销。在后续工作中，该策略被拓展至 G-Streamline 框架，专门处理不规则内存访问，并引入了“任务交换”（job swapping）技术，其核心思想是对线程索引进行重排。

除了 CPU 侧实现，也有研究提出基于源到源编译器的自动 GPU 端线程——数据重映射方法，后续工作还进一步支持了嵌套分支的情况。与 warp 发散紧密相关的一个术语是线程块压缩（thread block compaction），这是由 Fung 和 Aamodt 提出的微架构级优化概念。该术语或更广义的“流压缩”（stream compaction）虽然也用于 GPU 并行过滤算法，但在此上下文中，指的是将相同行为的线程聚集到一起执行，从而避免发散。

除 Zhang 等人的研究外，Khorasani 等人也属于该类方案的代表。他们提出了“协同上下文收集”（Collaborative-Context-Collection）方法，用于处理含有分支的循环。warp 内的线程利用共享内存存储其将采取的分支路径，当某一分支路径收集到足够多线程上下文后，warp 决定执行该路径，从而消除发散。

Gmys 等人采用多种技术减少分支代码。他们首先分析哪些线程需要执行相似任务，并据此进行流压缩，只启动必要数量的线程执行操作；而对于无法避免高度分支的代码，他们采取每个 warp 仅启动一个线程的策略，类似于虚拟 warp 的概念。

与线程压缩相关的还有一种数据排序策略，即将工作负载相近的数据分配给同一 warp 或线程块。例如，有研究在 Barnes-Hut n-body 应用中，通过排序将空间上相邻的物体分组，从而减少因受力计算路径不同而引起的发散。在 Delaunay 网格细化中也采用了类似思路，对三角形进行排序以降低发散程度。其他研究也使用了类似的数据预排序或任务预测技术，如根据线程需执行的迭代次数排序任务，或依据上一轮迭代的行为预测线程分配。还有工作根据任务规模对输入进行排序，并均匀分配到线程块中。为避免该分配策略仍可能导致发散，一些研究还引入了统计方法来进一步降低发散概率。

更改数据布局（Change the Data Layout）
从某种意义上说，对数据进行排序本质上是对数据布局的一种改变，但它高度依赖于并行性。另一种更“纯粹”地改变数据布局的方式是填充（padding）。例如，Sitardi 等人在进行并行字符串匹配时，首先按字符串长度对其进行分组，然后在多个阶段内将字符串划分，使得 warp 中每个线程获得长度相同的字符串，从而实现负载均衡。

对于稀疏数据结构，如稀疏矩阵——向量乘法或图算法，也存在多种用于减少分支发散的技术。一些研究者提出了专门的稀疏格式，以降低分支发散的发生。这些格式中的一部分是以减少分支发散为设计初衷的。例如，Ashari 等人提出了 Blocked-Row Column 格式，Khorasani 等人则提出了一种图表示格式 G-shards，还有其他多种图格式也专门面向减少分支发散进行优化。Sha 等人通过将图的邻接表解码过程拆分为两个阶段，从而有效减少了分支发散。

算法级变更（Algorithmic Changes）
最后，还有一些高度依赖具体算法本身的技术。Zhang 等人在 AutoRopes 基础上提出了名为 RegTT 的框架，以支持在遍历树结构时并行执行多个查询操作。RegTT 相较于 AutoRopes 的改进在于：在遍历初期采用广度优先搜索（BFS），在达到某一深度后切换为深度优先搜索（DFS），并对查询请求进行重排序，使得具有相同路径前缀的查询能被集中处理。这一设计显著提高了同一个 warp 中线程执行相似路径的可能性，从而有效减少了分支发散。

A.2.3 稀疏矩阵格式（Sparse Matrix Format）

GPU 在执行稀疏线性代数计算（特别是稀疏矩阵-向量乘法，SpMV）方面被广泛采用。用于表示稀疏矩阵的内存格式对性能有显著影响，因此已有大量不同的格式被提出。本节讨论了一些专为 GPU 优化的主流稀疏格式。关于基础知识的综述可参见 Bell 和 Garland 的工作；而关于近期的全面回顾，则可参考 Muhammed 等人的研究。图 22 展示了三种由 Bell 等人提出的常见稀疏格式。其中，ELL 格式采用填充策略，其填充元素以 * 标示。

由于 SpMV 属于高度不规则的应用，其性能表现极度依赖于输入矩阵，因此难以对不同格式的性能优势给出统一定量。然而，Muhammed 等人提供了一份包含多个输入样本上平均与总吞吐率的综合评估。例如，在其提出的 SURAA 稀疏格式上，在 Kepler 架构的 GPU（2012）上，相对于不同的对比格式，其加速比从 1.1 ~ 40 倍不等。在某些特定输入上，加速比可达 1.94 至 562 倍，但也存在在个别输入下性能反而下降的情况。

ELL 格式（也称为 ELLPACK）通常适用于每行非零元数量大致一致的稀疏矩阵，这类矩阵常见于结构化网格或半结构化网格生成的离散系统中。在此格式中，稀疏矩阵被压缩成一个稠密矩阵：通过去除零元素、将非零元素向左对齐、并在每行末尾填充零元素，确保每行恰好包含 K 个元素。该格式的优点在于结构规则，适合 GPU 并行计算；但缺点是会引入额外的填充开销。为了克服这一问题，研究者提出了多种改进方案。例如，ELL-R 格式在原始 ELL 结构基础上额外维护一个数组，用以记录每行的真实长度，从而允许线程跳过尾部的填充值，避免无效计算，但代价是需要额外的存储。Sliced-ELL 格式将矩阵划分为多个连续的 C 行条带，并对每个条带独立采用 ELL 存储，每个条带可使用不同的 K 值，从而减小整体填充冗余。ELLR-T 格式则结合了上述两者的优点：条带划分与每行长度记录同时存在。SELL-C-σ 格式进一步在 Sliced-ELL 基础上进行优化：其一，强制 C 为 SIMD 宽度的整数倍，以支持向量化执行；其二，将每个连续的 σ 行按照行长度排序，使长度相近的行聚集在同一个条带中。Kreutzer 等人对参数 C 与 σ 进行了系统性的调优，以适配不同架构下的最优性能。

除对单一格式进行改进外，另一研究方向是混合稀疏格式的组合使用。例如，Bell 和 Garland 提出了 HYB（Hybrid）格式，其将每行的前 K 个非零元存储为 ELL 格式，超出部分使用 COO 格式存储，从而在保留结构规则性的同时减小填充冗余。Yang 等人提出将矩阵按行分组，并基于启发式方法选择每组使用 ELL 或 CSR 格式。Su 等人则提出了更严格的 cocktail 格式：将矩阵划分为多个子块，并为每个子块选择 9 种现有格式之一进行存储；该方案通过离线训练构建推荐系统，根据体系结构特性为任意稀疏矩阵选择最优的格式组合。

稀疏矩阵运算通常具有较低的计算密度，属于内存带宽受限型任务。因此，对格式进行压缩可以提升吞吐率，尽管可能引入额外的解码计算。例如，Xu 等人采用索引压缩方法，用 16 位代替 32 位存储列索引，仅适用于特定矩阵。Tang 等人提出了位表示优化（Bit-Representation Optimized, BRO）压缩策略，结合增量编码与位打包技术，应用于三种常见稀疏格式，并实现了平均 1.5 倍的加速。Yan 等人提出了 BCCOO（Blocked Compressed Common Coordinate）格式，通过位标志对行索引数组进行压缩，实现从 32 位整数压缩为每索引 1 位的理论比率，但代价是在运行时需要进行索引解码。

上述格式大多假设稀疏矩阵为静态结构，但在图处理等应用中，稀疏矩阵通常作为图的表示，会随时动态变化，因此也催生了支持边的添加与删除的动态稀疏格式。Méndez-Lojo 提出了一种基于稀疏位向量的动态格式，用于并发地应用图重写规则。Busato 等人则提出 Hornet 格式，其通过分层数据结构表示稀疏矩阵，依赖向量化位树实现对图结构的高效更新和并行操作。

A.2.4 Kernel Fission（核函数裂变）

核函数裂变（Kernel Fission）指的是将一个核函数划分为多个较小核函数的过程，或者将单个核函数的迭代过程拆分为多个迭代过程。从优化意图上看，它可以被视为与算子融合（见第 A.1.8 节）相对的技术。该优化能够提升性能的核心机制在于：通过简化和结构规则化核函数，改善硬件资源的利用效率。

核函数裂变的经典应用可见于 Reddy 等人的工作中。他们在并行归约操作中，将整个任务划分为两个连续执行的核函数：第一个核函数并行地计算多个部分归约结果，第二个核函数再将这些中间结果合并，以得到最终值。类似地，Satish 等人利用此技术来应对 GPU 上缺乏线程块间同步机制的问题（见第 A.3.10 节）。

在密码学应用场景中，核函数裂变常被用于简化复杂核函数的结构。值得注意的是，An 等人甚至将裂变后的某个核函数迁移至 CPU 执行，从而将核函数裂变与异构计算（CPU/GPU 协同计算，见第 A.4.2 节）结合使用。

此外，核函数裂变还被用作减少分支发散的一种手段（参见第 A.2.2 节）。其基本思想是：为每条可能的分支路径生成一个独立的核函数，运行时由主机程序根据条件选择适当的核函数执行。然而，尽管该方案在理论上有效，实际应用中却可能导致工程复杂度显著增加。例如，Daga 等人通过该技术将一个核函数拆分为 16 个，分别对应 16 条可能的执行路径。尽管如此，该方法在 2009 年发布的 AMD Cypress GPU（该架构分支单元数量较少）上实现了约 70% 的加速，在分子建模类应用中取得了显著效果。

提升计算规则性也是核函数裂变的重要应用方向。诸如稀疏矩阵-向量乘法等场景中，有大量工作表明将一个复杂核函数划分为多个子核函数可以更好地利用局部结构的规律性。即使在处理稠密矩阵-向量乘法时，核函数裂变也能被用于支持任意尺寸矩阵的处理。例如，Abdelfattah 等人在其工作中，为矩阵底部的不规则行单独设计了一个核函数，从而避免在主核函数中实现复杂的边界处理逻辑。

最后还值得强调的是核函数裂变与自动调优（见第 A.3.6 节）之间的紧密关联。Alimirzazadeh 等人将原始核函数划分为更小、更简单的核函数，部分原因正是这些更细粒度的核函数在运行时配置上更易调优，也更利于利用编译器或运行时系统进行高效优化。

A.2.5 减少冗余计算（Reduce Redundant Work）

顾名思义，减少冗余计算是一种避免执行被视为冗余的计算工作的优化方式。它与重新计算值（参见 A.1.5 节）相对立。虽然此优化的主要效果似乎仅是减少执行的操作数量，例如 Qiu 等人和 Střelák 等人的研究所展示的那样，但一般而言，减少冗余工作量还可以带来更少的内存操作和通信开销。该优化常用于处理不规则数据结构，例如图结构、动态规划问题，或线性代数问题。在非确定性有限状态机领域，Liu 等人通过一种不存储冗余信息的实现，在 Pascal GPU（2016）上相较于朴素实现达到了 70% 到 6 倍的性能提升。尽管减少冗余计算是一种通用优化策略，其具体实现通常具有算法依赖性。

A.3 平衡(权衡)

A.3.1 指令流平衡——向量化

向量化指的是使用向量代替标量变量，并通过对向量所有元素统一施加的指令来操作这些向量。这种优化通常通过在 CUDA 和 OpenCL 中使用向量数据类型来实现，尽管这两种语言也支持专门的向量化加载和存储操作。然而，尽管两种编程语言都提供了向量化的支持，但并非所有 GPU 在硬件层面上都支持向量指令，因此这些指令可能会在内部被实现为对向量元素依次施加的一系列标量指令。甚至存在一些情形，其中向量化是通过数据重排与编译器优化的组合方式来实现的。

在 GPU 上，这种优化最常见的使用场景是使用向量数据类型来访问内存。通过使用向量，确实可以减少为加载或存储一定数量的内存而需要发出的指令数，从而实现更高的内存带宽。特别地，对内存操作进行向量化有助于改进合并访问。在代码可移植性方面，Yang 等人展示了 AMD 和 NVIDIA GPU 都能从向量化中受益，尽管使用的向量大小不同，并且 NVIDIA GPU 主要在使用纹理内存时受益于向量化。具体而言，在一款 Fermi（2011）GPU 上，他们达到了 2.3 倍的加速；而在一款 Cypress AMD（2009）GPU 上，他们在所有基准测试中平均测得了 6.1 倍的加速。

向量化的另一个常见使用场景是控制某次计算中线程的数量以及每个线程所执行的工作量。例如，Jang 等人使用该优化以每个线程计算更多的值，从而减少解决问题所需的线程数量。Reis 等人也使用向量化来增加每个线程的工作量，同时利用向量操作来并行化顺序循环。类似地，Ragan-Kelley 等人使用单个向量指令来替代小型循环，而 Nisa 等人则使用向量来减少内层循环的迭代次数。另一个值得一提的例子是将线程合并与向量化结合起来，使用向量重新引入因线程合并而减少的并行性。

向量化在 GPU 上还被用于实现其他目标，从访问稀疏矩阵，到改进数据存储方式，再到优化数值运算，以及减少在小数据类型上进行的条件语句和赋值操作。与其他任何优化一样，向量化的使用也并非没有代价，正如 Volkov 和 Demmel 在他们的微基准测试中所展示的那样，在某些情况下，使用向量甚至可能导致性能下降，正如 Doerksen 等人所指出的那样。

A.3.2 指令流平衡——快速数学函数

快速数学函数优化是指使用近似的数学函数，这些函数比标准函数显著更快，并且通常在硬件的特殊功能单元中实现。该优化可以由程序员手动应用，使用内建函数（intrinsics）代替标准函数，或者通过编译器启用，对所有数学函数生效。在我们选取的文献中，作者通常选择性地应用快速数学函数，使用内建函数（intrinsic）手动选择所需数学函数的快速版本，仅有一项研究是通过编译器标志全局启用该优化。所有情况下，作者都表示使用快速数学函数所提供的精度已足以满足他们的需求。值得注意的是，在该优化中，快速函数与另一项优化“将分支改写为算术操作”（见 A.2.2 节）结合使用，在代码中引入的用于替代条件语句的等式中使用了近似函数。

为了展示使用快速数学函数对性能的影响，Cecilia 等人报告在 Fermi 架构 GPU（2010）上实现了 3.3 到 4.8 倍的加速。该加速依赖于旅行商问题的规模，优化通过将代价昂贵的 powf() 函数替换为内建函数实现。在近期架构（Volta、Turing、Ampere）中，NVIDIA 引入了张量核心，这些是专门用于张量操作的功能单元，广泛用于机器学习应用。这些张量核心提供了具有不同精度的矩阵乘加（MMA）指令，例如半精度。Yan 等人详细研究了该操作，分析了其吞吐量和延迟与架构瓶颈的关系。通常，这些张量核心通过如 cuBLAS 或 cuDNN 等库来调用，但也可以直接使用。Hagedorn 等人支持在调度语言中使用这些指令，Tang 等人将其用于稀疏矩阵乘法，Yan 等人用于优化的半精度矩阵乘法，Zhai 等人则将其用于小矩阵批量乘法。

A.3.3 指令流平衡——基于 Warp 的编程

该优化背后的思想是将 warp 不仅作为程序执行模型的核心组成部分，也作为编程模型的一部分，因此编写的代码应围绕将 warp 作为计算单元来组织。尤其值得关注的是此优化与“减少同步”之间的关系（参见 A.3.8 节），因为以 warp 为中心的编程方式所带来的一个效果正是同步开销的降低。实现以 warp 为中心的编程的一种方式是重组代码，使得工作不再分配给线程或线程块，而是直接分配给 warp。

Hong 等人使用 Tesla GPU（2008）给出了该技术潜在加速效果的一个示例：他们在四个不同的图上应用广度优先搜索，具体加速效果依图而异，分别为 4%、2.5 倍、8 倍。但在某一图上，他们也观察到 18% 的性能下降。

Cecilia 等人将该概念扩展为所谓的“超级 warp”，即虚拟 warp，其大小超过硬件所支持的 warp，并将任务分配给这些虚拟 warp。然而，他们必须通过编写 PTX 代码手动为这些 warp 提供同步原语。尽管如此，文献中更常见的是将 warp 分解为更小单元的做法，而不是使用超级 warp。

一个使用任意大小虚拟 warp 的有趣示例由 Busato 等人提供，他们的虚拟 warp 是动态的，其大小基于输入图的一些特性。采用以 warp 为中心的编程的常见原因包括实现某种形式的负载均衡、隐藏延迟，以及实现嵌套并行性。该优化的另一个应用是 Bauer 等人所描述的“warp 专用化”。在这项工作中，一个线程块被划分为执行不同工作的多个 warp，这些 warp 被专用化为两类：一类执行内存操作，另一类执行计算操作。

A.3.4 与并行性相关的平衡 —— 每线程工作量的变化

改变每个线程和线程块分配的工作量 是 GPU 编程中最重要且普适的代码优化策略之一。该优化的主要思想是改变线程与问题域的映射方式，从而增加或减少每个线程执行的工作量。这可以通过多种方式实现，例如，在给定的核函数中，可以通过减少每个线程块中的线程数或减少线程块的总数来增加每个线程的工作量。图 23 展示了一个二维问题域的划分方式：左侧为每个线程执行一个工作单元，右侧通过 work_per_thread_x 增加了三倍的工作量，通过 work_per_thread_y 增加了两倍的工作量。

最早将这一策略作为 GPU 编程上下文中一种优化技术的文章之一是 Ryoo 等人，他们使用“1x2 矩形 tiling”一词来表示使用一个线程块计算两个 tile 的工作。Volkov 和 Demmel 也采用了这一优化，并将其称为 strip mining，这一术语与部分循环展开或循环划分同义。增加每个线程的工作量可以被视为对已在 GPU 核函数中并行化的外层循环进行 strip mining。Yang 等人将该优化区分为线程块合并（threadblock merge）和线程合并（thread merge），其中“线程块合并决定每个线程块的工作量，而线程合并决定每个线程的工作量”。

Hsu 等人指出，线程合并也被称为工作项合并（work-item merge）。Magni 等人对这一优化进行了广泛评估，并称之为线程粗化（thread coarsening）。他们展示了在所用的一半基准测试中该优化有效，且其效果在不同架构间有所差异。例如，该优化在他们使用的 AMD GPU（Cypress 2009，Tahiti 2011）上比在 NVIDIA GPU（Fermi 2011，Kepler 2012）上更有效。在 Fermi 上获得的最大加速比为 3.95，在 Kepler 上为 2，在 Cypress 上为 3.78，在 Tahiti 上为 12.01。Hong 等人也使用了“块粗化（block coarsening）”这一术语，与线程粗化类似，但作用于线程块层面。他们还指出，“线程粗化是一种实现寄存器 tiling 的方法”，见 A.1.3 节。Garvey 等人也将其称为块合并（block merge），并指出这是一种 tiling 形式。

一般而言，增加每个线程的工作量有助于在具备数据复用特征的核函数中更好地利用数据复用，但代价是寄存器和共享内存等资源使用量的增加。增加每个线程的工作量还可以通过消除先前分布在多个线程间的冗余指令来压缩指令流，例如，多个索引计算、循环计数器和分支指令可以用来支持多个元素的计算，而不只是一个。增加每个线程的工作量有时也用于为使用向量数据类型创造条件，后者反过来又可以提升内存吞吐量（见 A.3.1 节）。通常，找到每个线程最优的工作量并非易事，这可能取决于输入规模和具体的硬件架构。因此，该参数通常被纳入自动调优优化中（见 A.3.6 节）。

A.3.5 与并行性相关的平衡 —— 调整线程块大小

一般而言，在 GPU 编程中，程序员在选择每个线程块的线程数量和线程块数量方面具有相当大的自由度。例如，对于一个包含 $2^{20}$ 个元素的简单向量加法操作，其中每个线程处理一个元素，可以使用每个线程块 1,024 个线程和 1,024 个线程块，或 512 个线程和 2,048 个线程块，或 256 个线程和 4,096 个线程块等等。在某些情况下，每个线程块中的线程数并不会对性能产生显著影响，但在更复杂的核函数中，性能差异可能非常显著。具体而言，有报告指出，在一个二维结构网格计算流体力学求解器中，在 Tesla（2008）GPU 上，最优线程块配置与非最优配置之间的加速比达到了 30 倍。

通常，编译器会决定每个线程使用的寄存器数量，因此，线程块中线程数量的增加会影响寄存器使用量，并可能导致寄存器溢出。此外，线程数量还决定了每个 SM 上可以运行多少个线程块（除其他因素外），而每个 SM 上更多的（相互独立的）线程块可以在遇到同步屏障时保持计算单元的繁忙状态。许多作者明确地比较了内存操作与 warp 计算操作的延迟。线程数量通常也影响线程块使用的共享内存数量，而共享内存的使用也会影响每个 SM 上可运行的线程块数量。

每个线程块的线程数量是自动调优中使用最广泛的参数之一（见 A.3.6 节），原因有二：首先，线程数量对上述所有因素都有显著影响，性能往往依赖于这些因素之间的良好协同，而这种正确的平衡通常难以分析和找到；其次，线程数量相对容易调整，并且可以与线程块总数配合变化，因此一些作者利用这一点来自动调优其核函数。

根据具体核函数的不同，调整线程块大小也可能影响其他方面。有些研究根据是否使用常规或宽幅（例如 128 位）加载操作来调整线程数量。Lee 等人希望增加线程数量，但受限于寄存器的使用数量。因此，他们使用共享内存代替寄存器，同时增加数据重用。Krotkiewski 和 Dabrowski 调整线程块大小，以在一个模板应用中平衡内部元素与边界元素的数量。Sugimoto 等人认为，如果一个核函数在线程块之间具有高度局部性，则在一个 SM 上保持大量驻留线程块对于良好地利用缓存非常重要。相比之下，他们的另一个应用局部性较差，因此通过增加每个线程块中的线程数量来掩盖内存延迟。Wu 等人根据其动态规划应用中不同阶段的需求自适应地调整线程数量。

尽管调整线程块大小的动机常常与改变每个线程的工作量相同（见 A.3.4 节），但改变线程数量并不一定会改变每个线程的实际工作量。然而，也可以将核函数组织成一种形式，使得每个线程的工作量依赖于线程数量。

A.3.6 与并行性相关的平衡 —— 自动调优（Auto-tuning）

算法常常需要设置某些参数以保证其按预期运行，同时也存在一些性能优化方法需要设置特定参数才能有效发挥作用。由于配置空间的组合性质，这些参数的最优值往往难以确定，并且搜索过程可能因输入规模或执行平台的依赖性而进一步复杂化。以自动方式或仅需最少用户交互探索配置空间，以找到最优配置的过程称为自动调优（auto-tuning）。自动调优在 CPU 上早已被使用，早在 2008 年，Govindaraju 等人就已将该优化技术用于确定在 GPU 上使用哪种 FFT 实现。

尽管自动调优已被用于不同领域的多种 GPU 应用参数的调整，但在某些领域，这种优化技术的应用更为广泛。其中最主要的应用领域是稀疏矩阵相关的计算（见 A.2.3 节）。在这些研究中，Yoshizawa 等人的工作尤为值得关注，他们不仅对算法的某些参数进行了调优，还将自动调优应用于稀疏格式本身的参数。此外，Ashari 等人提出了基于矩阵稀疏特性的调优方法，Choi 等人提出的技术则结合了建模与基准测试。

另外两个广泛应用自动调优的领域是模板（stencil）操作和稠密矩阵乘法。尤其值得一提的是，Garvey 等人使用基于机器学习的调优器来加速其模板计算，以及 Christen 等人提出的模板计算自动调优框架。另一个相关的应用领域是线性代数，尤其是 Cholesky 分解。

自动调优的一个经典应用是编译器参数的调优（在传统编译器中：例如内联自动化、自动选择循环展开因子），这在 GPU 上同样适用。其他更特定于 GPU 的应用包括线程数量的调优、线程粗化因子的调优，以及执行流的数量调优。调度也可从调优中受益。

GPU 上的自动调优研究还包括一些更通用的调优器和调优框架的使用，如 CLTune、PADL、Kernel Tuner、OpenTuner 以及 Kernel Tuning Toolkit，该工具包提供了用于运行和调优核函数的 API。我们也发现 Hong 等人提出的 GPU 通用优化方法中包含了调优功能。

尽管自动调优的主要目标是提升性能，但性能可移植性也常被提及作为该优化方法的另一目标。如果能找到正确的配置，自动调优可以显著提升性能。Magni 等人提供了一个很好的示例，他们找到了每个线程的最优工作量（见 A.3.4 节），在 Fermi（2010）、Kepler（2012）、AMD Cypress（2009）和 AMD Tahiti GPU（2011）上运行不同基准测试时获得了 2–12 倍的加速比。

A.3.7 与并行性相关的平衡 —— 负载均衡（Load Balancing）

GPU 提供了多个层次的并行性，而在每个层次上，负载均衡对性能都至关重要。我们将在多个层面上讨论负载均衡问题。

Warp 内部
warp 发散与 warp 内负载均衡是相似但又不同的问题。Warp 发散指的是同一 warp 内的线程在遇到分支时必须串行执行两个路径，导致工作重复。而负载均衡则指的是部分线程没有执行有用的工作，而其他线程执行了工作，因此并无工作重复的现象。在许多情况下，减少分支发散同时也能提升负载均衡，已有多项工作同时以这两个目标为优化方向。下文我们将讨论专门为负载均衡而设计的技术，关于减少分支发散的方法见 A.2.2 节。

GPU 上的图处理通常会导致负载不均衡。Hong 等人采用虚拟 warp 的方式解决了 warp 内部的发散问题。然而，一个 warp 内的某个线程仍可能承担较大的工作负载，从而造成 warp 内部负载不均衡。他们采用了一种名为 “defer outliers” 的技术，将大任务放入一个全局工作列表中，由另一个 kernel 执行该任务，并将其并行化分配给 warp 内的线程。

Merrill 等人比较了多种用于广度优先搜索的策略。他们得出结论，一种称为 “CTA+Warp+Scan” 的混合策略在负载均衡方面表现最佳。该方法依据图中邻接表的大小，将工作分配给线程块、warp，或使用基于 scan 的策略在 warp 内分配线程。其他研究者也采用了这一技术或类似的设计。

Davidson 等人比较了包括 Merrill 的 “CTA+Warp+Scan” 在内的多种负载均衡机制，并提出了 Load-Balanced Partitioning 方法，该方法依据图中边的数量（使用 CSR 格式）对工作进行分组，实现在线程块之间以及线程块内部的负载均衡，应用于单源最短路径问题。

Gunrock 提出了一种通用图遍历 API，并在其框架中结合了 Merrill 和 Davidson 的成果，形成了自己的负载均衡策略。Brahmakshatriya 等人对图处理框架中的负载均衡方案进行了较为全面的综述，并引入了 “CTA+Warp+Scan” 的基于边的变体。他们的框架 GraphIt 支持七种不同的负载均衡方案。

稀疏矩阵领域与图处理密切相关。AdELL 稀疏矩阵格式专为高度倾斜的稀疏矩阵设计，提供了一种启发式方法，以实现在 warp 内线程间的负载均衡。Ashari 等人提出了用于稀疏向量——矩阵乘法的自适应 CSR 格式 ACSR。他们将矩阵行按非零元数量划分为不同的 bin，并依据非零元数量启动不同的 kernel。Muhammed 等人提出了 SURAA 格式，基于 CSR 增加了一个排列数组，用于对具有相同非零元数量的行进行排序和分组。

高度倾斜（highly skewed）的稀疏矩阵：某些行拥有大量非零元素，而其他大多数行的非零元素极少甚至为零。

Khorasani 等人跳出传统稀疏格式和算法的范畴，将其推广为 “嵌套并行模式”，即一组粗粒度任务，每个任务包含一组细粒度子任务。他们提出 Collaborative-Task-Engagement 方法，旨在提升执行可拆分粗粒度任务的线程的性能。该方法使用共享内存登记细粒度任务，并对任务进行调度以避免 warp 效率低下。其他研究也采用了类似技术。

与减少分支发散类似，还有一些方法通过对数据排序，使线程更有可能拥有相同的工作量。其他方案包括 Plauth 等人针对 N 皇后问题的研究，每个线程在皇后排列树中搜索解。由于在较深层次，正确解数量减少，他们通过调整搜索深度来补偿线程的负载。

Zhang 等人同样聚焦于 N 皇后问题，但采用每个线程块一个工作队列的方式，由线程从中获取任务。他们将此方法与单一工作队列进行比较，得出结论：单一队列存在过多竞争，而每个线程块一个队列性能更好。Nasre 等人则为每个线程分配一个工作列表，并将其存储在共享内存中。

线程块内部的负载均衡
在一个线程块内部，warp 之间的工作负载可能不同，导致资源利用率不足。前述的 Hong 等人研究了 GPU 上的图算法，并针对线程块内的负载不均，将任务存储在一个全局工作队列中，warp 从中提取任务。Zhu 等人使用了一种类似的 warp-centric 技术（见 A.3.3 节）。在图 24 左侧，他们展示了一种方式：在广度优先搜索程序中，将图的顶点分配给线程，例如，将顶点 1、2、5 和 9 的边分配给线程 1。在右侧，他们将顶点分配给一个 warp，从而使该 warp 中的线程能够并行处理来自同一顶点的不同边，从而提高了负载平衡性和性能。

Lee 等人事先分析图的结构，并通过分裂和合并线程块来重新组织线程块的分配。Nisa 等人在进行稀疏张量计算时，将分配给线程的 fiber 拆分以平衡负载。Nasre 等人使用共享内存进行工作捐赠，即从一个线程将工作捐赠给另一个线程。Zhang 等人通过对序列进行排序以对齐，从而在一个线程块中平衡线程间的工作负载。Lee 等人展示了线程块内部负载均衡可带来显著的加速效果。根据不同架构（Pascal 2016、Volta 2017 和 Turing 2018）在各种图算法中的平均结果显示，其性能提升介于 40% 至 66% 之间。

线程块之间
通常线程块之间的负载均衡并不是问题，因为 GPU 编程模型假定线程块会超额分配到流处理器（Streaming Multiprocessors）上。然而，对于不规则负载来说，这种编程方式可能并非最佳。持久线程模型（persistent threads model）不同于默认的最大化启动的超额分配线程块的方式，该模型中程序员分配的是可以在 GPU 上物理并发运行的最大线程和线程块数。这些线程在整个核函数运行期间始终处于活跃状态，并通过队列交换工作。Gupta 等人对该模型进行了深入探讨，并将其与传统的 GPU 编程方式进行了比较。他们指出，对于某些应用，该模型确实能带来性能提升，但并非适用于所有情况，这一点也由 Nasre 等人指出。该模型的主要优势在于负载均衡，并使得线程块之间的同步更加高效，尽管这需要采用 A.3.10 节中所讨论的技术。

多位作者在处理不规则负载时使用了持久线程模型。Tzengy 等人将光线追踪应用中的工作划分为 bin，并采用工作窃取和工作捐赠技术，即当某个 bin 的任务溢出时，对应的线程块以轮转方式将任务捐赠给另一个流处理器。另一个线程块之间的负载均衡技术可见于稀疏张量计算领域。Nisa 等人可能会遇到某个线程块负载远高于其他线程块的情况。他们扩展了 Ashari 的 binning 技术（在“warp 内”一节中已讨论）以实现线程块的负载均衡。

CPU/GPU 之间的负载均衡
在同时使用 GPU 和 CPU 的应用中（见 A.4.2 节），可能需要在这两类计算设备之间平衡负载。作者们通常采用静态划分的方式，即预先确定一个固定比例的总工作量被分配给 GPU，其余部分则留给 CPU。

然而，这种解决方案不容易扩展，因此也有一些更通用的调度问题解决方案。Wan 等人基于执行时间的反馈动态地确定任务的最优比例。这种方法在执行时间可预测的情况下有效，但对于如稀疏矩阵乘法这类不规则应用则可能无效。Matam 等人展示了几种启发式方法，用于预测 CPU/GPU 计算比例的良好平衡。Wang 等人使用支持向量机算法学习工作负载之间的最优比例。

A.3.8 与同步相关的平衡——减少同步操作

并行算法因多种原因依赖于同步原语。其中常见的原因包括：确保所有线程都已到达某段代码，或者确保内存操作已完成，从而使对内存的更新对所有线程可见。在图 25 中，Phuong 等人可视化了一个使用共享内存的简单归约操作。线程块中的线程从共享内存中读取两个值，对这两个值求和，并将结果写回共享内存。由于线程是并行执行的，因此在各个步骤之间需要加入屏障。然而，同步可能会对性能产生显著影响，尤其是在线程数量不足、工作量不够的情况下，线程因空闲等待而导致的延迟难以被隐藏。因此，减少同步可被视为一种性能优化，尽管这一优化通常依赖于算法级的改动以在不影响正确性的前提下规避同步。

一个有趣的例子是 Petre 等人所采用的 Cooley-Tukey FFT 算法的变体，该变体无需同步。另一个值得注意的工作是 Lin 等人提出的方案，用自定义的同步机制替代共享内存屏障，以改善嵌套分支中的分支发散问题。在 CUDA 程序中，曾经较为常见的做法是通过使用隐式的 warp 级同步来减少显式的线程块级同步。然而，自 CUDA 9.0 起，这种做法已被弃用，程序员被建议使用显式的 warp 级同步原语。Agarwal 等人也指出，如果没有隐式同步，他们的性能将显著下降。OpenCL 中亦可通过使用大向量实现类似的隐式同步机制。一种相似的方法是，通过多次核函数执行所提供的隐式同步机制来消除核函数内的同步，如 Mawson 等人所示。

另一种减少同步的方法是增加每个线程的工作量，从而减少线程间同步的需求（见 A.3.4 节）。为实现这一目标，还可采用其他优化策略，例如使用共享内存和分块。具体来说，Reddy 等人通过完全展开一个基于树的归约算法，成功地规避了同步操作。

值得指出的是，在减少同步方面存在应用上的冲突现象。例如，一方面 Bauer 等人通过用更细粒度的原语替代屏障来减少同步开销；另一方面，Nasre 等人则通过将细粒度的同步替换为屏障来实现优化。Nasre 等人还提供了通过减少原子指令带来潜在加速的量化指标：在各种图算法中，依据架构（Fermi 2011 和 Kepler 2012）、图算法种类和输入数据，他们测得的加速比在 3% 到 50% 之间。

A.3.9 与同步相关的平衡——减少原子操作

原子操作允许程序员在进行内存更新时获得无冲突结果的保证。这种保证对于 GPU 尤为重要，因为在 GPU 上可能有数千个线程并行运行，并可能同时操作同一内存位置。然而，尽管原子操作对许多算法的正确性至关重要，但它们会引入计算开销，这是由于串行化所致：如果两个或多个原子操作访问相同位置，则一次只能执行一个，其余操作必须等待。因此，可以明确地看到，冲突次数越多，性能损失就越大。因此，GPU 编程文献中存在两种常见优化：避免原子操作（avoid atomics）和减少原子操作（reduce atomics）。虽然这两种优化类似，但“避免原子操作”指的是将代码中所有原子操作用其他方法替代的目标，而“减少原子操作”则是指仅替代其中的一部分，同时保留那些不可或缺的原子操作。

避免原子操作。
并没有一种通用的方法可以避免使用原子操作，大多数避免原子操作的技术都是特定于应用的。例如，Leist 等人通过将计数器替换为布尔标志来实现该目标，而 Vetter 等人则需要执行一些预处理来实现相同的效果。Cecka 等人提出了类似的方法，即应用特定于应用的数据划分方式，从而避免冲突，使作者能够避免使用原子操作。甚至在某些情况下，避免原子操作的方法是放宽某些约束条件，并维护图状态的近似视图，或者需要放宽应用的内存模型。通过使用屏障对线程何时看到内存操作结果施加顺序性，Burtscher 等人提出了该优化的另一种实现方式。Cecilia 等人通过应用另一种优化——scatter-to-gather 转换——来避免原子操作，而 Soman 等人依赖其应用的内在属性来实现相同的目标。Střelák 等人通过使用细粒度并行化策略来避免原子操作。最后，Xian 等人提出的线程块间同步策略的无锁实现是不依赖原子操作的广为人知的例子。

减少原子操作。
与前述优化类似，减少原子操作也有多种实现方式，其中大多数也是特定于应用的。例如，Pai 等人通过对图中的推送操作进行聚合，从而减少所需更新的总次数，并因此减少了原子操作的数量；Leist 等人则通过在共享内存中引入原子操作来减少全局内存上的原子操作数量。让线程在块级或 warp 级协作以减少对全局内存的更新次数是另一种减少原子操作的技术。Gaihre 等人将每个 warp 的原子操作数量减少到一个，代价是需要更多的 warp 内协作以及使用 shuffle 指令（见 A.1.2 节）。Anzt 等人采用了类似的技术。Gaihre 等人利用 Pascal GPU（2016）在其广度优先搜索实现中，在多个数据集上获得了平均 35% 的性能提升。其他研究表明，通过使用第三方库也可以减少原子操作。Samadi 等人选择性地减少原子操作，仅移除那些潜在冲突程度最高的原子操作实例，而将其余部分保留在代码中。在某些情况下，代码中原子操作的数量还可能受到另一个独立参数的影响，例如 Ashari 等人通过控制每个线程块的工作量来调节原子操作的数量。Anzt 等人通过对数据进行排序，降低了线程同时执行原子更新的可能性，从而减少了线程之间的原子操作冲突。

A.3.10 同步相关的负载均衡优化——线程块间同步（Inter-Block Synchronization）

CUDA 和 OpenCL 均提供了一种使用屏障（barrier）轻松同步同一线程块内线程的方法。然而，从历史上看，在线程块之外进行同步的唯一方式是通过多次核函数调用。为了克服多次调用所带来的开销，研究人员提出了多种实现线程块间同步的方法。由于该技术可以用于提升性能，因此我们将其视为一种优化方法。

例如，Wu 等人使用线程块间同步在全局屏障之后预取共享内存中的数据（见第 A.1.9 节），而这在仅使用核函数调用的情况下是不可能实现的。在图 26 的左侧，他们展示了由于全局内存是通过第二次核函数调用隐式同步的，因而无法将数据预取到共享内存。而在右侧，他们展示了通过线程块间的屏障，可以将数据预取到共享内存中，以供第二阶段使用。

该优化技术最早的引用可见于一个用于研究线性代数核函数性能的微基准测试背景下的工作中，Volkov 和 Demmel 使用全局内存中的原子操作实现了一个全局屏障。然而，Xiao 等人提出的两种全局屏障实现方式——一种使用原子操作，另一种为无锁实现——在文献中是该优化方法最常见的用法。Xiao 等人展示了其无锁技术在 Tesla GPU（2008）上针对不同基准测试的性能提升介于 11% 至 66% 之间。

总体而言，所有这些线程块间同步的实现（包括类似方法）都基于在全局内存中拥有一个所有线程块均可访问和修改的同步数据结构。NVIDIA 认识到线程块间同步及更细粒度同步的实用性，并在 2017 年发布了支持协作组（cooperative groups）的新版本 CUDA（9.0）。协作组允许程序员定义用于同步、shuffle 和投票函数的一组线程（见第 A.1.2 节）。它支持定义小于 warp 的线程组以用于屏障同步，也支持定义线程块级的线程组，相当于传统的线程块屏障，并进一步支持定义大于线程块的线程组，潜在地包括线程块网格甚至跨多个 GPU 的多个网格。

这种高级的线程块间同步方式需要架构支持，由 Pascal（2016）、Volta（2017）及之后的架构提供。

A.4 主机交互

A.4.1 主机通信

GPU 计算的主流模型是以 CPU 为主机、GPU 作为 PCI-Express 卡形式的加速器。这就要求在主机内存与 GPU 内存（即设备内存）之间通过 PCI-Express 总线来回复制数据。与设备内存和计算单元之间的带宽相比，PCI-Express 总线的带宽较低，因此在许多情况下必须优化主机通信以提升整体应用性能。主机与设备之间的通信主要有两种方式：隐式传输和显式传输。

隐式传输发生在主机上分配了页锁定（或称“固定”）内存，但这些内存已被映射到设备的地址空间中。虽然仍然通过 PCIe 总线进行传输，但可以与核函数执行重叠进行。这种技术有时被称为零拷贝（zero-copy）。更为复杂的隐式传输方式是统一内存（Unified Memory，也称为托管内存），其中设备（GPU 或主机）按需获取并交换页面。这种方式既简化了代码，又具有加速潜力，尽管有研究未观察到性能提升，这很可能是由于内存访问模式造成的。统一内存在指针密集型代码中可能带来性能优势，因为通过页面交换，可以直接通过指针访问数据元素；而使用零拷贝访问时，每次通过指针访问数据元素都会导致 PCIe 传输。

也可以在主机程序中显式地将数据从主机传输到设备。使用页锁定内存可以使这些传输与核函数执行重叠。这通常通过 CUDA 的“流”（streams）抽象或 OpenCL 的“命令队列”（command queues）实现。核函数执行和数据传输是与流相关联的命令，或被放入 OpenCL 队列中，其中这些流或队列彼此独立。每个命令返回一个事件，并可通过这些事件在流或队列之间实现同步。在 OpenCL 中，队列可以是顺序的（等同于 CUDA 流），也可以是无序的，后者需要使用事件进行同步。

有研究展示了额外使用流带来的性能提升。传输可以与 GPU 计算重叠，从而提高性能。最有效的优化方式是完全消除通信，即将整个算法都在 GPU 上执行，或尽可能多地在 GPU 上执行，或尽可能多地将数据保留在 GPU 上。另一种方法是压缩需要通过 PCI-e 总线传输的数据。有研究指出了这种优化可能带来的加速效果。在 Tesla（2009）和 Fermi（2010）GPU 上，他们在一个无网格计算流体力学应用上报告了分别为 28.9% 和 18.9% 的性能提升。

另有一种有趣的方法是使用动态并行性（即线程也可以启动核函数）将 CPU 的计算迁移到 GPU。在主机端启动一个单线程的核函数，该线程会调用其他核函数，实质上将控制循环移到 GPU 上负责核函数调用。这种方案需要调整同步开销。

还有一种有趣的方法是使用一种称为 GMAC 的非对称分布式共享内存（ADSM）实现，它基于 OpenCL 简化并加速了数据传输。它在 OpenCL 中为用户提供了 CPU 和 GPU 内存的统一视图，CUDA 也提供了类似功能。

许多针对主机通信的优化聚焦于传输的组织方式。许多方法简单地使用页锁定内存来加速传输。其他方法则使用页锁定和/或映射内存来实现传输与计算的重叠。更高级的重叠方案可通过使用流或命令队列实现，这种方法被称为“流化”。流被用于多种方案，例如：每个线程一个流，一个流用于传输、一个用于核函数执行，或者分别设置用于磁盘到主机、主机到设备传输以及核函数执行的流。

页锁定内存 是指一段不会被操作系统分页管理机制移动或回收的主机内存区域，通常用于 CPU 与 GPU 间高效的数据传输。

在这些更复杂的方案中，也必须妥善管理缓冲区。一些方法采用双缓冲，甚至采用三缓冲，使用三个线程和三个流分别用于主机到设备传输、设备到主机传输和核函数执行。有研究使用调优过的缓冲区，以实现加载与计算的完美重叠。通过这些技术可以实现流水线化，这对带宽受限的核函数尤其有益。

例如，有研究使用双缓冲和三缓冲流水线，包含两个命令队列，一个用于数据传输，一个用于核函数执行，通过事件进行同步。另有研究使用流并将任务拆分为更小的部分以重叠传输。还有更多方法使用流水线，但也有研究对其提出了有趣的应用。他们设计了一种流水线方案以隐藏“线程-数据重映射”的延迟，该内容在 A.2.2 节中进行了讨论。

A.4.2 CPU/GPU 协同计算

GPU 几乎总是作为包含一个或多个 CPU、主机内存、与其他系统的互联以及 I/O 的系统的一部分使用。因此，开发者需要决定应用程序的每一部分应在何种计算设备上运行。尽管通常的选择是在 CPU 或 GPU 上运行，但并非总是如此。这种优化技术包括将计算在 CPU 和 GPU 之间划分，使用两个设备同时执行有用的工作。

Che 等人提供了关于性能提升的一个示例，他们在一块 Fermi（2010）GPU 上测量了 K-means 应用在不同数据布局下获得的 15% 和 20% 的性能提升。该优化技术的一个常见使用场景是，当应用程序可以被划分为一组独立或半独立的任务时。如果这些任务之间存在依赖关系，那么优化 CPU 与 GPU 之间的数据传输就变得对性能至关重要，正如 Fan 等人所示。该技术的另一个示例是将非性能关键的代码段保留在 CPU 上执行。

当关于如何在 CPU 和 GPU 之间划分计算的选择并不明显时，一个关键任务是在两个设备之间找到某种平衡，正如在负载均衡章节（A.3.7 节）中所讨论的那样。最后还值得指出的是，这种优化技术并不总能带来性能提升。

附录 B：方法论

本附录对第 3 章中的内容进行了扩展，详细说明了文献筛选过程及处理方式。

B.1 阶段一：基于标题、发表会议和关键词的筛选

在第一阶段中，我们选取了一批作者熟悉的 GPU 优化相关文献，从中收集作者提供的关键词，并据此构建了用于 Scopus 数据库检索的关键词集合，以启动优化相关文献的查询。我们在查询中使用了下列关键词。必要时，我们通过添加 AND KEY(gpu) 或其他 GPU 相关术语来补充关键词上下文。每一行表示一个独立的查询语句，并与以下固定查询模板组合：

{} AND SUBJAREA(COMP OR ENGI) AND DOCTYPE(ar OR cp) AND (PUBYEAR > 2005)

其中，{} 代表关键词查询语句。我们通过该方式跟踪各个关键词的检索效果。这些关键词包括：

KEY(gpu) AND KEY(optimization*)
KEY(graphics processing unit) AND KEY(optimization*)
KEY(gpgpu) AND KEY(optimization*)
KEY(cuda) AND KEY(optimization*)
KEY(opencl) AND KEY(optimization*)
KEY(many*core) AND KEY(optimization*)
KEY(control flow divergence)
KEY(warp-synchronous programming)
KEY(branch divergence)
KEY(warp execution)
KEY(thread-data remapping)
KEY(memory coalescing)
KEY(coalesced memory accesses)
KEY(divergence) AND KEY(gpu)
KEY(divergence) AND KEY(graphics processing unit)
KEY(divergence) AND KEY(gpgpu)
KEY(divergence) AND KEY(cuda)
KEY(divergence) AND KEY(opencl)
KEY(divergence) AND KEY(many*core)
KEY(kernel fusion) AND KEY(gpu)
KEY(kernel fusion) AND KEY(graphics processing unit)
KEY(kernel fusion) AND KEY(gpgpu)
KEY(kernel fusion) AND KEY(cuda)
KEY(kernel fusion) AND KEY(opencl)
KEY(kernel fusion) AND KEY(many*core)
KEY(fusion) AND KEY(gpu)
KEY(fusion) AND KEY(graphics processing unit)
KEY(fusion) AND KEY(gpgpu)
KEY(fusion) AND KEY(cuda)
KEY(fusion) AND KEY(opencl)
KEY(fusion) AND KEY(many*core)
KEY(irregular algorithms) AND KEY(gpu)
KEY(irregular algorithms) AND KEY(graphics processing unit)
KEY(irregular algorithms) AND KEY(gpgpu)
KEY(irregular algorithms) AND KEY(cuda)
KEY(irregular algorithms) AND KEY(opencl)
KEY(irregular algorithms) AND KEY(many*core)
KEY(irregular computations) AND KEY(gpu)
KEY(irregular computations) AND KEY(graphics processing unit)
KEY(irregular computations) AND KEY(gpgpu)
KEY(irregular computations) AND KEY(cuda)
KEY(irregular computations) AND KEY(opencl)
KEY(irregular computations) AND KEY(many*core)
KEY(data transformation) AND KEY(gpu)
KEY(data transformation) AND KEY(graphics processing unit)
KEY(data transformation) AND KEY(gpgpu)
KEY(data transformation) AND KEY(cuda)
KEY(data transformation) AND KEY(opencl)
KEY(data transformation) AND KEY(many*core)
KEY(graph representation) AND KEY(gpu)
KEY(graph representation) AND KEY(graphics processing unit)
KEY(graph representation) AND KEY(gpgpu)
KEY(graph representation) AND KEY(cuda)
KEY(graph representation) AND KEY(opencl)
KEY(graph representation) AND KEY(many*core)

该检索最初于2019年11月29日进行，并于2021年5月27日重复执行，以纳入最新发表的文献。检索共获得3,973篇文章，随后按发表年份降序排序，并在同一年内依据引用次数降序排列。此排序方式旨在初步判断文献的重要性，前提是假设同一年内的引用次数具有可比性，且引用次数可作为衡量文章影响力的合理指标。文献列表中包含标题、引用次数、年份、发表会议与关键词。当某篇文章明确涉及GPU优化时予以保留，例如，仅将应用程序移植至GPU的文章则被剔除。根据上述标准，共剔除2,853篇文章，保留1,120篇。

B.2 第二阶段：基于摘要与作者信息的筛选

在第二阶段中，延续第一阶段的排序与展示方式，并补充了作者名单及文章摘要。基于这些信息，对文献进行保留、剔除或标记为辅助文献。辅助文献的典型例子包括综述文章或关于GPU安全属性的研究文献。要使文章被正式纳入，需提供实施GPU优化的明确证据。在该阶段，我们剔除386篇文章，保留532篇，并将202篇标记为辅助文献。对辅助文献进行重新分类后，最终选出其中10篇，从而共纳入542篇文献。

B.3 第三阶段：文献扫描

在此阶段，我们对文献进行扫描，即阅读标题、作者、会议、关键词与摘要，并检查图表、图像以及各章节标题。根据获取的信息，依据以下纳入与排除标准进行筛选：
纳入标准：

提出对单一GPU核函数的优化；
提出对多个GPU核函数的优化（如算子融合）；
提出用于异构CPU-GPU系统的优化；
提出CPU-GPU数据通信相关优化；
在GPU核函数上应用自动调优；
直接在GPU上应用优化；
衡量优化效果的文献。

排除标准：
8. 针对多GPU集群的优化；
9. 针对硬件的优化（如warp调度器）；
10. 编译器执行、开发者无法自行实现的优化；
11. 单纯的算法层级优化。

纳入标准聚焦于GPU核函数优化及多核函数优化（如融合）；第3与第4条亦涉及主机侧优化。后三项标准则覆盖未必提出新优化方案的文献，例如只应用现有优化技术并结合自动调优的研究。排除标准则过滤掉对GPU集群、硬件层面或编译器不可控优化的研究，最后一条则排除仅为算法调整的文章。

据此标准，最终筛选出401篇文献，剔除141篇。

B.4 高频被引文献

我们使用Scopus数据库查询被上述401篇文章引用的所有文献，从中筛选出那些未被纳入但被引用次数超过5次的文献。因其在现有文献集中被频繁引用，认为其足够重要，应纳入整体分析流程。经此操作额外纳入149篇文献，并依据第二阶段的标准筛选，选出63篇；再根据第三阶段标准进一步筛选，最终纳入49篇。

B.5 第四阶段：优化技术提取

在此阶段，我们对总计450篇文献（包括第三阶段的401篇和高频引用的49篇）进行分析，提取其中明确提出的优化技术、性能瓶颈与使用的GPU信息。若文章中明确出现某种优化技术，我们将其归类记录，如“blocking”或“coalesced-access”，并关联文献标识以及相应的引文语句作为佐证。同样地，文中提到的GPU也按关键词（如“titanx”）归档，并关联文献标识。性能瓶颈信息亦采用相同方式记录。

B.6 第五阶段：优化技术的描述

在本阶段中，我们对已提取的优化技术、GPU信息以及性能瓶颈的数据库进行了清理，并按照文献数量对优化技术与性能瓶颈进行了排序。由于所选优化技术中存在大量仅在少数文献中提及的“长尾”项，考虑到时间与篇幅的限制，我们将分析范围限制在出现频次超过20篇文献的优化技术上。

我们将所归纳的优化技术按主题组织，并编写了相应章节进行描述。例如，“使用专用内存”这一章节涵盖了多个优化策略，如“使用共享内存”、“使用常量内存”以及“使用纹理内存”。这些章节在附录 A 中列出，旨在为希望深入了解特定优化技术的读者提供系统的参考资料。

附录 C：数据分析

本附录在第四节的基础上，进一步补充了图表和细节说明。Scopus 提供了文献的发表来源，如图 28 所示，展示了各发表源的文章数量。IPDPS、SC 和 PPoPP 是 GPU 优化研究中最受欢迎的三个会议，而 CPE、TPDS 和 JPDC 是最常见的三本期刊。整体分布呈现长尾特征：共有 188 个不同的发表源，其中许多源仅包含一至两篇本文数据集中的文章。这表明 GPU 优化研究广泛分布于多个研究群体与发表平台，而并不局限于特定社区或期刊。

我们对每篇文章中用于实验评估的 GPU 进行了人工提取，因为这可作为作者所针对硬件架构类型的有效指标。图 29 显示了所有文章中每种 GPU 被提及的频次（注意，一篇文章可能提及多个 GPU）。图中可以明显看出，NVIDIA 是最常被使用的 GPU 厂商。图 30 展示了各 GPU 架构所对应的文章数量。NVIDIA 架构构成了绝大多数评估使用的 GPU，而使用 AMD GPU 的文章略高于 10%。

图 4（第 5 页）进一步展示了各 GPU 架构在不同时期被文章采用的频率变化。我们可以观察到各个架构随时间的兴衰过程。例如，Fermi 架构发布于 2010 年，当年仅有一篇文章以其为研究对象，随后该数字逐年上升，并在 2013 年达到峰值，此时 73% 的文章使用 Fermi 架构，之后逐年下降。GPU 架构具有较长的生命周期。例如，尽管 Kepler 架构于 2012 年发布，但即便到 2019 年，它依然是使用最广泛的架构。至 2021 年，最常见的架构分别是 Pascal、Volta 和 Maxwell，尽管它们的发布年份分别为 2016、2017 和 2014。

图 3（第 4 页）显示了每种优化技术在文献中被提及的比例。“内存访问合并”（coalesced access）和“使用专用内存”（use dedicated memories）是最常见的优化技术，这与其对 GPU 性能的关键影响相符。“减少分支发散”（reduce branch divergence）和“循环展开”（loop unrolling）也非常常见，这是因为控制流结构在性能分析中同样起着重要作用。此外，“自动调优”（auto-tuning）也较为普遍，超过八分之一的文章采用了某种形式的参数自动调节策略。

附录 D：性能参数的平衡

本附录扩展了第 5.4 节内容，介绍了 Volkov 提出的一个简化模型，用于刻画 GPU 核函数的合理利用程度。该模型以单个流式多处理器（SM）为视角，并假设线程块数量足够多，从而保证整体 GPU 的性能。通常，“占用率”（occupancy）一词用来描述 GPU 的利用率，指的是活动 warp 数与 SM 所支持的最大 warp 数的比值。人们普遍认为，如果在选择寄存器数量、共享内存或线程块大小时使得占用率较低，则无法实现高性能。

然而，Volkov 指出，占用率仅衡量线程级并行性（TLP），而高性能还依赖于指令级并行性（ILP）、带宽级并行性（BLP）、向量指令和数据类型的利用。此外，对于某些核函数而言，低占用率反而是实现高性能的必要条件。后续研究中，Volkov 强调了指令流中指令类型的重要性，并指出许多分析模型忽略了 ILP，因而提出了一个更为复杂但能涵盖 ILP 的分析性能模型。

本节介绍 Volkov 提出的简化利用模型，该模型已足以阐明参数平衡的重要性。模型使用以下参数：

latency：内存或算术指令的延迟（以时钟周期为单位）
throughput：指令吞吐率（以 SM 中核心数衡量）
parallelism：并行度（以每个 SM 的操作数表示）
nrThreads：每个线程块的线程数（由程序员定义）
activeTBs：每个 SM 上的活动线程块数
ILP：指令流中的独立指令数
BLP：向量宽度

首先，确定指令流中最关键的指令类型（例如单精度浮点运算），并通过文档或基准测试获得该指令的延迟值。接着，根据指令所使用的计算单元数量，确定其吞吐率（例如某 GPU 每个 SM 有 64 个核心，则单精度指令的吞吐率为 64）。

接下来，依据 Little 定律确定所需的最小并行度：

$\text{parallelism}_{min}=\text{latency}×\text{throughput}$

其中，latency 以周期计，throughput 为每个 SM 的核心数，parallelism_min 单位为每个 SM 的操作数。
线程级并行度（TLP）由 SM 上活动线程块数决定，公式如下：

$\text{TLP} = \text{activeTBs} \times \text{nrThreads}$

在该公式中，每个活动线程执行一次操作，因此 TLP 单位也为每 SM 的操作数。

最终，利用率定义为：

$\text{utilization} = \min\left(1,\ \frac{\text{TLP} \times \text{ILP} \times \text{BLP}}{\text{parallelism}_{\text{min}}}\right)$

利用率范围为 0 到 1。若上述比值大于 1，意味着核函数的并行能力已超过延迟的掩盖需求，从而可以充分利用资源。

以 Volkov 举的 GTX480 为例，单精度操作的延迟约为 18 个周期，吞吐率为 48 核心/SM，因此最小并行度为 864 ops/SM。在无指令级并行与位级并行的情况下，需要 864 个活动线程才能实现 100% 利用率；若指令流中有两个独立指令，则只需 432 个线程；若有四个，则仅需 216 个，以此类推。同样的道理也适用于带宽，使用更宽的数据类型可降低线程数的需求，例如加载 float2 元素时，仅需加载 float 元素时线程数的一半。

综上所述，实现高性能往往是一种参数平衡行为，其中 TLP 和 ILP 类似于连通容器的两个腔体，但中间设有“阀门”，一旦资源限制（如寄存器使用量）被触发，就会影响线程块活跃数，进而影响 TLP。因此，只有在明确限制条件的情况下，才能准确判断哪个因素主导性能。图 8（第 24 页）展示了这些参数之间的相互影响关系。

说明

①文章中，算子、Kernel、核函数、核函数均是一个英文单词 kernel，只是在不同语境下翻译方式换了