GPU上的强缩放

Question

我想研究一下我的并行GPU代码(用OpenACC编写)的强大扩展性。使用GPU进行强扩展的概念--至少据我所知--比CPU更模糊。我发现的关于GPU强扩展的唯一资源建议解决问题的大小和增加GPU的数量。然而，我相信在GPU中有一些强大的扩展，例如在流多处理器上(在中)。

OpenACC和CUDA的目的是明确地将硬件抽象给并行程序员，并将她限制在他们的三层编程模型上，包括帮派(线程块)、工作人员(翘曲)和向量(SIMT线程组)。我的理解是，CUDA模型的目标是为其线程块提供可伸缩性，这些线程块是独立的，并且映射到SMX。因此，我看到了两种利用GPU进行强缩放的方法：

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

我的问题是:我关于GPU强扩展的思路正确/相关吗？如果是这样的话，是否有办法在OpenACC中实现上面的第2步呢？

Answer 1

GPU具有很强的扩展性，但并不一定按照您的想法，这就是为什么您只能找到关于强扩展到多个GPU的信息。有了一个多核CPU，您就可以精确地决定要运行多少个CPU核心，这样您就可以修复工作并调整跨核线程的程度。使用GPU，跨SMs的分配是自动处理的，完全超出了您的控制范围。这是由设计，因为它意味着一个良好的GPU代码将强大的规模，以填补任何GPU (或GPU)，你抛向它，而没有任何程序员或用户的干预。

您可以运行一些少量的OpenACC帮派/CUDA线程块，并假设14个帮派将运行在14个不同的短信，但这有两个问题。首先，1组/螺纹块不会使一个开普勒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组中对循环进行分区。我所期待的是，当你加入帮派，你会看到更好的表现，直到一些多倍的短信，届时的性能将变得大致持平(当然，离群点)。正如我上面所说，在你看到最佳表现的时候固定帮派的数量并不一定是最好的主意，除非这是你唯一感兴趣的设备。相反，要么让编译器决定如何分解循环(这允许编译器根据您告诉它构建的体系结构做出智能决策)，要么公开尽可能多的帮派，这将使您具有更强的并行性，可以扩展到更大的GPU或多个GPU，这样就会有更多的可移植代码。

Answer 2

为了占用完整的SMX，我建议使用共享内存作为占用的限制资源。编写一个内核，它消耗了所有32 of的共享内存，这个块将占用整个SMX，因为SMX耗尽了另一个块的资源。您可以将块从1扩展到13 (对于K20c)，调度程序(希望如此)可以将每个块调度到不同的SMX。你可以先把每个区块的数据扩展到192，这样每个CUDA核心都会很忙，然后你就可以进一步提高翘曲调度器的效率了。GPU通过隐藏延迟来提供性能。所以你必须从一个街区占据一个SMX到N个区块。您可以通过使用较少的共享内存来做到这一点。再一次放大你的翘曲以掩盖隐藏的延迟。

我从未碰过OpenACC，如果您真的想完全控制您的实验代码，请使用CUDA而不是OpenACC。您无法看到OpenACC编译器内部以及它如何处理代码中使用的实用程序。