Technical Walkthrough

通过 GPU 内存访问调整提高应用程序性能

 

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_arraysmall_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 倍。

Three line graphs showing a drop in time it takes to perform kernel launches when using memory prefetch.
图 1 :。由于负载较宽,内核持续时间减少

第二步:寄存器使用

我们可能想到此为止并宣布成功,但使用 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 行,指令的执行次数突然跳到其最大值。

Screen capture from Nsight Compute tool showing inline hexadecimal encoding of assembly language instructions indicating GPU time taken to execute each instruction.
图 2 :。演示内部循环开始的 SASS 代码段(第 166 行)

LDG 。 E 、 64 是我们所追求的指令。它从全局内存( DRAM )加载一个具有扩展地址的 64 位字。宽单词的负载对应于 LDG 。 E 、 128 。加载指令名称后的第一个参数(图 2 中的 R34 )是接收该值的寄存器。由于双精度值占用两个相邻寄存器,因此加载指令中隐含 R35 。接下来,我们比较三个版本的代码( 1.基线, 2.宽负载的small_array, 3.宽负载的small_arraybig_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 提供了在本博客讨论的示例中分析寄存器使用的专业知识。