cuda - GPU 上的强大扩展

Question

我想研究我的并行 GPU 代码(使用 OpenACC 编写)的强大扩展性。至少据我所知，GPU 强扩展的概念比 CPU 更加模糊。 The only resource I found regarding strong scaling on GPUs建议修复问题大小并增加 GPU 数量。然而，我相信 GPU 内部存在一定程度的强大扩展能力，例如流式多处理器上的扩展(在 Nvidia Kepler 架构中)。

OpenACC 和 CUDA 的目的是明确地将硬件抽象给并行程序员，将其限制在带有组(线程 block )、工作线程(扭曲)和向量(SIMT 线程组)的三级编程模型中。据我了解，CUDA 模型旨在提供与其线程 block 相关的可扩展性，这些线程 block 是独立的并映射到 SMX。因此，我看到了两种研究 GPU 强扩展的方法:

1. 固定问题大小，并将线程 block 大小和每个 block 的线程数设置为任意常数。缩放线程 block 的数量(网格大小)。
2. 考虑到有关底层硬件的其他知识(例如 CUDA 计算能力、最大扭曲/多处理器、最大线程 block /多处理器等)，设置线程 block 大小和每个 block 的线程数，以便 block 占据整个和单个 SMX。因此，在线程 block 上进行扩展相当于在 SMX 上进行扩展。

我的问题是:我关于 GPU 强扩展的思路是否正确/相关？如果是这样，有没有办法在 OpenACC 中执行上述#2 操作？

Answer 1

GPU 可以进行强大的扩展，但不一定以您所想的方式进行，这就是为什么您只能找到有关多个 GPU 的强大扩展的信息。使用多核 CPU，您可以轻松决定要在多少个 CPU 核心上运行，这样您就可以修复工作并调整跨核心的线程程度。使用 GPU，SM 之间的分配是自动处理的，完全不受您的控制。这是设计使然，因为这意味着编写良好的 GPU 代码将具有强大的扩展能力，可以填充您投入的任何 GPU(或多个 GPU)，而无需任何程序员或用户干预。

您可以在少量 OpenACC 组/CUDA 线程 block 上运行，并假设 14 个组将在 14 个不同的 SM 上运行，但这存在一些问题。首先，1 个组/线程 block 不会使单个 Kepler SMX 饱和。无论有多少线程，无论占用多少，每个 SM 都需要更多的 block 才能充分利用硬件。其次，您并不能真正保证硬件会选择以这种方式安排 block 。最后，即使您在自己的设备上找到了每个 SM 的最佳 block 或组数，它也无法扩展到其他设备。 GPU 的技巧是公开尽可能多的并行性，以便您可以从具有 1 个 SM 的设备扩展到具有 100 个 SM 的设备(如果存在)，或者扩展到多个设备。

如果您想试验固定工作量下 OpenACC gang 数量的变化如何影响性能，您可以使用 num_gangs 子句(如果您使用的是 parallel 区域，或 gang 子句(如果您使用的是kernels)。由于您试图强制循环的特定映射，因此最好使用并行，因为这是更具规范性的指令。您想要执行的操作如下所示:

#pragma acc parallel loop gang vector num_gangs(vary this number) vector_length(fix this number)
for(i=0; i<N; i++)
  do something

这告诉编译器使用一些提供的向量长度对循环进行向量化，然后在 OpenACC 组之间划分循环。我期望的是，当你添加帮派时，你会看到更好的性能，直到 SM 数量的某个倍数，此时性能将变得大致持平(当然也有异常值)。正如我上面所说，在您看到最佳性能时固定组的数量不一定是最好的主意，除非这是您感兴趣的唯一设备。相反，可以让编译器决定如何分解循环，这允许编译器根据您告诉它构建的架构做出明智的决策，或者通过暴露尽可能多的组，这为您提供了额外的并行性，可以强大地扩展到更大的 GPU 或多个 GPU，您将拥有更多可移植代码。