OpenAcc 标准中内核和并行指令之间的区别

Question

我已经使用支持 OpenAcc 的 PGI 编译器在 GPU 上启动代码大约 3 年了，但到目前为止我无法理解术语 "kernels" 和 "parallel" 之间的区别。我在 OpenAcc 入门指南中阅读：

Parallel Construct

Defines the region of the program that should be compiled for parallel execution on the accelerator device.

Kernels Construct

Defines the region of the program that should be compiled into a sequence of kernels for execution on the accelerator device.

我不明白 "parallel execution on the accelerator device" 和 "compiled into a sequence of kernels for execution on the accelerator device" 这两个词有什么区别。如果加速器设备是 GPU，那么所有代码都被编译成一定大小的 CUDA 内核（我试图指的是 CUDA 网格和块）并且这些 CUDA 内核在 GPU 上的 CUDA 线程中执行，不是吗？什么是 "sequence" 内核？ "parallel" 指令生成 1 个内核，而 "kernels" 可以从同一段代码生成一系列内核？

此外，我在任何地方都只使用 "parallel" 循环指令。例如，为了在 GPU 上并行执行 for 循环，我写

#pragma acc parallel loop gang vector copy(...) present(...)
  for(int i=0; i<N; ++i)
  {
    ...
  }

是否正确？什么时候应该使用 "kernels"？或者它是 "parallel" 的同义词，现在已弃用且不再使用？

Answer 1

考虑差异的最佳方式是，使用 "parallel"，程序员正在定义并行化哪些循环以及如何并行化。基本上，您是在告诉编译器并行化特定循环。使用 "kernels"，您正在定义一个可以并行化的代码区域，但编译器的工作是确定要并行化哪些循环以及如何并行化。

对于 "parallel"，该区域内的所有代码都作为一个 CUDA 内核卸载。如果在 "parallel" 区域内有多个外部循环，它们仍将在一个 CUDA 内核中卸载。由于编译器可以发现 "kernels" 的并行化，因此该区域内的多个循环可能会拆分为一系列单独的 CUDA 内核启动。

可在以下位置找到完整的详细信息：https://www.pgroup.com/lit/articles/insider/v4n2a1.htm

请注意，访问该文章需要您拥有 PGI Web 用户帐户。

Answer 2

已经发布了我的答案 here，但这里又出现了。

平行构造

定义应编译以在加速器设备上并行执行的程序区域。
并行循环指令是程序员断言并行化受影响的循环既安全又可取。这依赖于程序员正确识别代码中的并行性并删除代码中可能无法安全并行化的任何内容。如果程序员错误地断言循环可以并行化，那么最终的应用程序可能会产生错误的结果。
并行构造允许更细粒度地控制编译器将如何尝试在加速器上构建工作。所以它并不严重依赖编译器自动并行化代码的能力。
当并行循环用于访问相同数据的两个后续循环时，编译器可能会也可能不会在两个循环之间的主机和设备之间来回复制数据。
更有经验的并行程序员可能已经在他们的代码中识别出并行循环，他们可能会发现并行循环方法更可取。

例如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 不同）。

内核构造

定义应编译成内核序列以在加速器设备上执行的程序区域。
关于内核构造需要注意的一件重要事情是，编译器将分析代码并且仅在确定这样做是安全的情况下才进行并行化。在某些情况下，编译器在编译时可能没有足够的信息来确定循环是否是安全的并行化，在这种情况下它不会并行化循环，即使程序员可以清楚地看到循环是安全并行的。
内核构造为编译器提供了最大的余地来并行化和优化代码，使其适合目标加速器，但也最依赖于编译器自动并行化代码的能力。
内核构造提供的一个更显着的好处是，如果多个循环访问相同的数据，它只会被复制到加速器一次，这可能会导致更少的数据移动。
并行编程经验较少或代码包含大量需要分析的循环的程序员可能会发现内核方法更简单，因为它给编译器带来了更多负担。

例如refer

#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];
}

 生成两个内核

 两者之间存在隐性障碍循环：第二个循环将在第一个循环之后开始循环结束。

OpenAcc 标准中内核和并行指令之间的区别

difference between kernels and parallel directives in OpenAcc standard

c++

gpu

openacc

pgi