2

我想研究我的并行 GPU 代码（用 OpenACC 编写）的强大扩展性。使用 GPU 进行强大扩展的概念——至少据我所知——比使用 CPU 更模糊。我发现的关于 GPU 强大扩展的唯一资源建议修复问题大小并增加 GPU 的数量。然而，我相信在GPU中存在一定程度的强大扩展，例如通过流式多处理器（在 Nvidia Kepler 架构中）进行扩展。

OpenACC 和 CUDA 的目的是明确地将硬件抽象给并行程序员，将她限制在他们的三级编程模型中，包括 gangs（线程块）、workers（warp）和vectors（SIMT 线程组）。据我了解，CUDA 模型旨在为其线程块提供可扩展性，这些线程块是独立的并映射到 SMX。因此，我看到了两种研究使用 GPU 进行强大扩展的方法：

  1. 修复问题大小，并将线程块大小和每个块的线程数设置为任意常数。缩放线程块的数量（网格大小）。
  2. 给定关于底层硬件的额外知识（例如 CUDA 计算能力、最大扭曲/多处理器、最大线程块/多处理器等），设置线程块大小和每个块的线程数，以便一个块占用整个和单个 SMX。因此，在线程块上进行缩放等同于在 SMX 上进行缩放。

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

4

2 回答 2

5

GPU 具有强大的扩展能力，但不一定以您所想的方式，这就是为什么您只能找到有关对多个 GPU 进行强大扩展的信息的原因。使用多核 CPU，您可以轻松确定要在多少 CPU 内核上运行，这样您就可以修复工作并调整内核之间的线程化程度。使用 GPU，SM 之间的分配是自动处理的，完全不受您的控制。这是设计使然，因为这意味着编写良好的 GPU 代码将具有强大的扩展能力，可以在没有任何程序员或用户干预的情况下填充你扔给它的任何 GPU（或 GPU）。

您可以在少量 OpenACC gangs/CUDA 线程块上运行，并假设 14 个 gangs 将在 14 个不同的 SM 上运行，但这有几个问题。首先，1 个帮派/线程块不会使单个 Kepler SMX 饱和。无论有多少线程，无论占用率如何，每个 SM 都需要更多的块才能充分利用硬件。其次，您并不能真正保证硬件会选择以这种方式调度块。最后，即使您在您拥有的设备上找到每个 SM 的最佳块数或帮派数，它也不会扩展到其他设备。GPU 的诀窍是尽可能多地公开并行性，以便您可以从具有 1 个 SM 的设备扩展到具有 100 个（如果它们存在的话）的设备，或者扩展到多个设备。

如果您想试验为固定工作量改变 OpenACC 帮派的数量如何影响性能，您可以使用num_gangs子句（如果使用parallel区域）或gang子句（如果使用kernels. 由于您试图强制循环的特定映射，因此您最好使用parallel，因为这是更具规范性的指令。您想要做的是如下所示：

#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 数量的倍数，此时性能将大致持平（当然有异常值）。正如我上面所说，在您看到最佳性能的点上固定 gang 的数量不一定是最好的主意，除非这是您唯一感兴趣的设备。相反，通过让编译器决定如何分解循环，它允许编译器根据你告诉它构建的架构做出明智的决定，或者通过暴露尽可能多的帮派，这为你提供了额外的并行性，可以强大地扩展到更大的 GPU 或多个 GPU，你

于 2014-12-18T17:01:06.763 回答
0

为了占用一个完整的 SMX，我建议使用共享内存作为占用的限制资源。编写一个内核，消耗所有 32kB 的共享内存，该块将占用整个 SMX，因为 SMX 的资源不足，无法容纳另一个块。您可以将块从 1 扩展到 13（对于 K20c），并且调度程序将（希望）将每个块调度到不同的 SMX。你可以先将每个块的 therads 扩大到 192 以使每个 CUDA 核心忙碌，然后你可以进一步让 warp 调度程序满意。GPU 通过延迟隐藏提供性能。因此，您必须从占用 SMX 的 1 个块移动到 N 个块。您可以通过使用更少的共享内存来做到这一点。并再次扩大你的扭曲以覆盖延迟隐藏。

我从未接触过 OpenACC，如果您真的想完全控制您的实验代码，请使用 CUDA 而不是 OpenACC。您看不到 OpenACC 编译器的内部以及它对代码中使用的 pragma 的作用。

于 2014-11-14T09:17:54.870 回答