GPU上的强大扩展

Question

我想调查我的并行GPU代码(使用OpenACC编写)的强大扩展.使用GPU进行强扩展的概念 - 至少据我所知 - 比使用CPU更加模糊.我发现有关GPU的强大扩展的唯一资源建议修复问题大小并增加GPU的数量.不过,我相信有很强的比例一定量的范围内,例如缩放过流多处理器(在NVIDIA开普勒架构)的GPU.

OpenACC和CUDA的目的是明确地将硬件抽象给并行程序员,将其限制为使用帮派(线程块),工作者(warps)和向量(SIMT线程组)的三级编程模型.据我所知,CUDA模型旨在提供与其线程块相关的可伸缩性,这些线程块是独立的并映射到SMX.因此,我看到了两种方法来研究GPU的强缩放:

1. 修复问题大小,并将每个块的线程块大小和线程数设置为任意常数.缩放线程块数(网格大小).
2. 有了底层硬件的额外知识(例如CUDA计算能力,最大扭曲/多处理器,最大线程块/多处理器等),设置线程块大小和每个块的线程数,使得块占用整个单个SMX.因此,扩展线程块相当于在SMX上进行扩展.

我的问题是:关于GPU上的强缩放是否正确/相关,我的思路是什么？如果是这样,有没有办法在OpenACC中做到#2？

Answer 1

GPU的规模很大,但不一定与您的思维方式一致,这就是为什么您只能找到有关强扩展到多个GPU的信息的原因.使用多核CPU,您可以轻松确定要运行的CPU核心数,因此您可以修复工作并调整核心的线程程度.使用GPU,SM之间的分配将自动处理,完全不受您的控制.这是设计使然,因为它意味着编写良好的GPU代码可以强大地扩展以填充您在其中投入的任何GPU(或GPU)而无需任何程序员或用户干预.

您可以运行一些少数OpenACC帮派/ CUDA线程块,并假设14个团伙将在14个不同的SM上运行,但这有几个问题.首先,1个gang/threadblock不会使单个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的数量的倍数,此时性能将变得大致平坦(当然有异常值).正如我上面所说的那样,在你看到最佳性能的时候修复帮派的数量并不一定是最好的主意,除非这是你唯一感兴趣的设备.相反,要么让编译器决定如何分解循环,它允许编译器根据您告诉它构建的体系结构做出明智的决策,或者通过尽可能多地暴露团队,这为您提供额外的并行性,可以扩展到更大的GPU或多个GPU,您需要更多便携式代码.