我学习了很多OpenACC的文章和手册,但仍然不理解这两个结构之间的主要区别。
kernels
指令是更常见的情况,如果您之前编写过GPU(例如CUDA)内核,则可能会考虑这种情况。 kernels
指令简单地指示编译器处理代码段,并生成任意数量的“内核”,以任意“维度”顺序执行,从而将特定部分的代码并行化/卸载到加速器。 parallel
构造允许更精细地控制编译器尝试在加速器上组织工作的方式,例如通过指定特定的并行化维度。例如,工作者和帮派的数量通常是parallel
指令的一部分(因为通常只暗示一个基础“内核”),但可能不适用于kernels
指令(因为它可以翻译成多个基础“内核”)。OpenACC指令和GPU内核只是表示同一件事情的两种方式——可以并行运行的代码部分。
当改装现有应用程序以利用GPU和/或希望让编译器处理与内存管理等问题相关的更多细节时,OpenACC可能是最好的选择。这可以使编写应用程序更快,但可能会降低性能。
当从头开始编写GPU应用程序和/或需要更精细的控制时,内核可能是最好的选择。这可能会使编写应用程序的时间更长,但可能会提高性能。
我认为,对于新手来说,他们可能会倾向于选择OpenACC,因为它看起来更熟悉。但我认为更好的方法是先编写内核,然后在某些项目中转向OpenACC以节省时间。原因是OpenACC是一个泄漏的抽象。因此,使用OpenACC编写GPU代码而不了解背景中发生的事情很可能会令人沮丧,并导致在尝试编译时出现奇怪的错误消息,并且结果是应用程序性能较低。
kernels
是由OpenACCæ ‡å‡†å®šä¹‰çš„æŒ‡ä»¤ã€‚å®ƒè¿˜åŒ…æ‹¬parallels
指令。 - Mark Ebersole并行构造
定义程序应该编译为在加速设备上并行执行的区域。
并行循环指令是程序员断言影响循环可以安全和可取并行化的声明。这依赖于程序员正确地识别代码中的并行性,并删除可能不安全并行化的任何内容。如果程序员错误地断言循环可以并行化,则生成的应用程序可能会产生不正确的结果。
并行构造允许更细粒度地控制编译器如何尝试在加速器上结构化工作,因此不过度依赖编译器自动并行化代码的能力。
当在两个相邻访问相同数据的循环上使用并行循环时,编译器可能会在两个循环之间在主机和设备之间来回复制数据。
有经验的并行程序员可能已经在其代码中确定了并行循环,他们可能会发现并行循环方法更可取。
例如 refer
#pragma acc parallel
{
#pragma acc loop
for (i=0; i<n; i++)
a[i] = 3.0f*(float)(i+1);
#pragma acc loop
for (i=0; i<n; i++)
b[i] = 2.0f*a[i];
}
生成一个内核
两个循环之间没有障碍:第二个循环可以在第一个循环结束之前开始。(这与OpenMP不同)。
内核构造
定义程序区域,应编译为一系列内核以在加速器设备上执行。
关于内核结构的重要事项是,编译器将分析代码,并仅在确定安全并行化时才进行并行化。在某些情况下,编译器可能无法在编译时获得足够的信息来确定循环是否安全并行化,在这种情况下,即使程序员清楚地看到循环是安全的并行化,它也不会并行化循环。
内核结构为编译器提供了最大的自由度,以便根据目标加速器自行并行化和优化代码,但同时也最大程度地依赖于编译器自动并行化代码的能力。
内核结构提供的另一个显着好处是,如果多个循环访问相同的数据,则该数据仅会被复制到加速器一次,这可能导致更少的数据移动。
具有较少并行编程经验或其代码包含需要分析的大量循环的程序员可能会发现内核方法更简单,因为它将更多的负担放在编译器上。
#pragma acc kernels
{
for (i=0; i<n; i++)
a[i] = 3.0f*(float)(i+1);
for (i=0; i<n; i++)
b[i] = 2.0f*a[i];
}
生成两个内核
两个循环之间存在隐式屏障:第二个循环将在第一个循环结束后开始。
parallel
的实现要好得多。据我所知,在GCC中不支持kernel
的reduction
。 - Z boson