NVIDIA GPU 具有强大的计算能力，通常需要高速传输数据才能部署这种能力。原则上，这是可能的，因为 GPU 也有很高的内存带宽，但有时他们需要程序员的帮助来饱和带宽。在这篇博文中，我们研究了一种实现这一点的方法，并将其应用于金融计算中的一个示例。我们将解释在什么情况下这种方法可以很好地工作，以及如何找出这些情况是否适用于您的工作负载。
上下文
NVIDIA GPU 的力量来自大规模并行。可以将 32 个线程的许多扭曲放置在流式多处理器（ SM ）上，等待轮到它们执行。当一个 warp 因任何原因暂停时， warp 调度程序将切换到另一个，开销为零，确保 SM 始终有工作要做。在高性能 NVIDIA Ampere 100 （ A100 ） GPU 上，多达 64 个活动经线可以共享一个 SM ，每个都有自己的资源。除此之外， A100 还有许多 SMs-108 ，它们都可以同时执行 warp 指令。大多数指令都必须对数据进行操作，而这些数据几乎总是源自连接到 GPU 的设备内存（ DRAM ）。 SM 上大量的翘曲也可能无法工作的一个主要原因是，它们正在等待来自内存的数据。如果发生这种情况，并且内存带宽没有得到充分利用，则可以重新组织程序以改进内存访问并减少扭曲暂停，从而使程序更快完成。
第一步：宽负载
在之前的 博客文章 中，我们检查了一个工作负载，该工作负载没有充分利用 GPU 的可用计算和内存带宽资源。我们确定，在需要之前从内存中预取数据可以大大减少内存暂停并提高性能。当预取不适用时，需要确定哪些其他因素可能会限制内存子系统的性能。一种可能性是，向该子系统发出请求的速率太高。直观地说，我们可以通过在每个加载指令中提取多个单词来降低请求速率。最好用一个例子来说明这一点。
在本文的所有代码示例中，大写变量都是编译时常量。 BLOCKDIMX 采用预定义变量 blockDim 的值。 x 、 出于某些目的，它必须是编译时已知的常量，而出于其他目的，它有助于避免在运行时进行计算。
原始代码如下所示，index是计算数组索引的辅助函数。它隐式地假设只使用了一个一维线程块，而派生它的激励应用程序则不是这样。但是，它减少了代码混乱，并且不会更改参数。
for (pt = threadIdx.x; pt < ptmax ; pt += BLOCKDIMX ) { double best = 0.0; #pragma unroll for (int k = 0; k < kmax; ++k) { double c = big_array[index(pt, k)]; c += small_array[k] ; best = max(c, best); } final[pt] = best;
}请注意，每个线程从建议命名的small_array中加载kmax个连续值。此阵列足够小，完全适合一级缓存，但要求它以非常高的速率返回数据可能会出现问题。下面的更改表明，如果我们稍微重新构造代码并引入 double2 数据类型，则每个线程可以在同一条指令中发出两个双精度字的请求，这在 NVIDIA GPU 上本机支持；它将两个双精度字存储在相邻的内存位置，可以使用字段选择器“ x ”和“ y ”访问这些位置。之所以这样做，是因为每个线程都访问small_array的连续元素。我们称这种技术为 VZX28 。请注意，索引“k”上的内部循环现在增加了 2 ，而不是 1 。
for (pt = threadIdx.x; pt < ptmax ; pt += BLOCKDIMX ) { double best = 0.0; #pragma unroll for (int k = 0; k < kmax; k+=2) { double c = big_array[index(pt, k)]; double2 val = *(double2 *) &small_array[k]; c += val.x; best = max(c, best); c = big_array[index(pt, k+1)]; c += val.y; best = max(c, best); } final[pt] = best;
}有几个注意事项。首先，我们没有检查kmax是否为偶数。如果没有，修改后的k循环将执行额外的迭代，我们需要编写一些特殊代码来防止这种情况发生。其次，我们没有确认small_array是否在 16 字节边界上正确对齐。否则，宽荷载将失效。如果它是使用cudaMalloc分配的，它将自动在 256 字节的边界上对齐。但是，如果使用指针算法将其传递给内核，则需要执行一些检查。
接下来，我们检查辅助函数指数，发现它在 pt 中与系数 1 呈线性关系。因此，通过在一条指令中请求两个双精度值，我们可以对从 big \ U 数组获取的值应用类似的宽负载方法。对big_array和small_array的访问之间的区别在于，现在 warp 中的连续线程访问相邻的数组元素。下面重构的代码将数组元素上的循环增量加倍big_array，现在每个线程在每次迭代中处理两个数组元素。
for (pt = 2*threadIdx.x; pt < ptmax ; pt += 2*BLOCKDIMX ) { double best1 = 0.0, best2 = 0.0; #pragma unroll for (int k = 0; k < kmax; k+=2) { double2 c1 = *(double2 *) &big_array[index(pt, k)]; double2 c2 = *(double2 *) &big_array[index(pt, k+1)]; double2 val = *(double2 *) &small_array[k]; c1.x += val.x; best1 = max(c1.x, best1); c2.x += val.y; best1 = max(c2.x, best1); c1.y += val.x; best2 = max(c1.y, best2); c2.y += val.y; best2 = max(c2.y, best2); } final[pt] = best1; final[pt+1] = best2;
}与之前相同的注意事项也适用，现在应该扩展到ptmax的奇偶校验和big_array的对齐。幸运的是，从中派生此示例的应用程序满足所有要求。下图显示了在应用程序中重复多次的一组内核的持续时间（以纳秒为单位）。对于宽负载组合，内核的平均加速比为 1.63 倍。
第二步：寄存器使用
我们可能想到此为止并宣布成功，但使用 NVIDIA Nsight Compute 对程序执行的深入分析表明，即使我们将加载指令的数量减少了一半，我们也没有从根本上改变对内存子系统的请求速率。原因是一条扭曲加载指令（即 32 个线程同时发出加载指令）会导致一个或多个扇区请求，这是硬件处理的实际内存访问单元。每个扇区是 32 字节，因此每个线程一条 8 字节双精度字的扭曲加载指令会导致 8 个扇区请求（访问以单位跨距进行），而一条双精度字的扭曲加载指令会导致 16 个扇区请求。普通负载和宽负载的扇区请求总数相同。那么，是什么导致了性能的提高呢？
为了理解代码行为，我们需要考虑一个尚未讨论的资源，即寄存器。这些用于存储从内存加载的数据，并用作算术指令的输入。寄存器是一种有限的资源。如果流式多处理器（ SM ）在 A100 GPU 上承载尽可能多的扭曲，则每个线程可以使用 32 个 4 字节寄存器，这些寄存器总共可以容纳 16 个双精度字。将代码翻译成机器语言的编译器知道这一点，并将限制每个线程的寄存器数量。我们如何确定代码的寄存器使用及其在性能中所起的作用？我们使用 Nsight Compute 中的“ source ”视图来并排查看汇编代码（“ SASS ”）和 C 源代码。
代码的最内层循环是执行次数最多的循环，因此，如果我们在导航菜单中选择“已执行的指令”，然后要求转到 SASS 代码中数量最多的那一行，我们会自动进入内部循环。如果不确定，可以将 SASS 与突出显示的相应源代码进行比较以确认。接下来，我们在内环的 SASS 代码中识别从内存（ LDG ）加载数据的所有指令。图 2 显示了 SASS 的一个片段，我们在其中搜索以找到内部循环的开始；在第 166 行，指令的执行次数突然跳到其最大值。
LDG 。 E 、 64 是我们所追求的指令。它从全局内存（ DRAM ）加载一个具有扩展地址的 64 位字。宽单词的负载对应于 LDG 。 E 、 128 。加载指令名称后的第一个参数（图 2 中的 R34 ）是接收该值的寄存器。由于双精度值占用两个相邻寄存器，因此加载指令中隐含 R35 。接下来，我们比较三个版本的代码（ 1.基线， 2.宽负载的small_array， 3.宽负载的small_array和big_array）在内部循环中使用寄存器的方式。回想一下，编译器试图保持在限制范围内，有时需要对寄存器进行处理。也就是说，如果没有足够的寄存器可用于从内存接收每个唯一值，它将重用以前在内部循环中使用的寄存器。
这样做的结果是，算术指令需要使用以前的值，以便新值可以覆盖它。此时，从内存加载需要等待该指令完成：内存延迟暴露。在所有现代计算机体系结构上，此延迟构成了一个显著的延迟。在 GPU 上，可以通过切换到另一个扭曲来隐藏部分扭曲，但通常不是全部扭曲。因此，寄存器在内环中被重用的次数可以表示代码的速度变慢。
有了这一见解，我们分析了代码的三个版本，发现它们在每个内部循环中分别经历了 8 、 6 和 3 个内存延迟，这解释了图 1 所示的性能差异。不同寄存器重用模式背后的主要原因是，当两个普通加载融合为单个宽加载时，通常需要更少的地址计算，并且地址计算的结果也会进入寄存器。随着持有地址的寄存器越来越多，剩下来充当从内存中提取的值的“着陆区”的地址越来越少，我们在 Music chairs 游戏中失去了席位；寄存器压力增大。
第三步：启动边界
我们还没有完成。现在我们知道了寄存器在程序性能中所起的关键作用，我们将查看三个版本的代码使用的寄存器总数。最简单的方法是再次检查 Nsight Compute 报告。我们发现使用的寄存器数量分别为 40 、 36 和 44 。
编译器确定这些数字的方法是使用复杂的启发式算法，该算法考虑了大量因素，包括 SM 上可能存在多少活动扭曲、在忙循环中加载的唯一值的数量以及每个操作所需的寄存器数量。如果编译器不知道 SM 上可能存在的扭曲数，它将尝试将每个线程的寄存器数限制为 32 ，因为如果存在硬件允许的绝对最大同时扭曲数（ 64 ），那么这就是可用的数字。在我们的例子中，我们没有告诉编译器期望的是什么，所以它尽了最大努力，但显然确定仅使用 32 个寄存器生成的代码效率太低。
然而，内核的 launch 语句中指定的线程块的实际大小是 1024 个线程，因此有 32 个扭曲。这意味着，如果 SM 上只存在一个线程块，则每个线程最多可以使用 64 个线程。在实际使用的每个线程中有 40 、 36 和 44 个寄存器时，没有足够的寄存器可用于支持每个 SM 的两个或多个线程块，因此将只启动一个，每个线程分别保留 24 、 28 和 20 个未使用的寄存器。
通过使用 launch bounds 将我们的意图告知编译器，我们可以做得更好。通过告诉编译器一个线程块中的最大线程数（ 1024 ）和同时支持的最小块数（ 1 ），编译器可以放松，并且很高兴每个线程分别使用 63 、 56 和 64 个寄存器。
有趣的是，最快的代码版本现在是基线版本，没有任何广泛的负载。虽然组合宽负载 without 启动边界的加速比为 1.64 倍，但宽负载 with 启动边界的加速比为 1.76 倍，而基线代码的加速比为 1.77 倍。这意味着我们不必费心修改内核定义；在这种情况下，仅提供启动边界就足以获得这种特定线程块大小的最佳性能。
通过对 SM 上的线程块大小和预期的最小线程块数进行更多的实验，我们在每个 SM 有 512 个线程的 2 个线程块的情况下达到了 1.79 倍的加速，对于没有宽负载的基线版本也是如此。
结论
寄存器的有效使用对于获得良好的 GPU 内核性能至关重要。有时，一种称为“宽负载”的技术可以带来显著的好处。它减少了计算并需要存储在寄存器中的内存地址的数量，留下更多的寄存器来接收来自内存的数据。然而，向编译器提示在应用程序中启动内核的方式可能会带来同样的好处，而无需更改内核本身。
确认书
作者要感谢 NVIDIA 的 Mark Gebhart 和 Jerry Zheng 提供了在本博客讨论的示例中分析寄存器使用的专业知识。
