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

Question

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

并行构造

定义应该为并行编译的程序区域 在加速器设备上执行。

内核构造

定义程序的区域应该被编译成一个 用于在加速器设备上执行的内核序列。

我不明白“在加速器设备上并行执行”和“编译成一系列内核以在加速器设备上执行”之间的区别。如果加速器设备是 GPU，那么所有代码​​都被编译成某种大小的 CUDA 内核（我试图指的是 CUDA 网格和块），并且这些 CUDA 内核在 GPU 上的 CUDA 线程中执行，不是吗？什么是内核的“序列”？ “并行”指令生成 1 个内核，而“内核”可以从同一段代码生成一系列内核？

我也只在所有地方使用“并行”循环指令。例如，为了在 GPU 上并行执行 for 循环，我写了

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

正确吗？什么时候应该使用“内核”？或者它是“并行”的同义词，现在已被弃用且未使用？

Answer 1

已经发布了我的答案here，但又来了。

并行构造

1. 定义应编译为在加速器设备上并行执行的程序区域。

2. 并行循环指令是程序员的断言，即并行化受影响的循环既安全又可取。这依赖于程序员正确识别代码中的并行性并删除代码中可能对并行化不安全的任何内容。如果程序员错误地断言循环可能被并行化，那么生成的应用程序可能会产生错误的结果。

3. 并行结构允许更细粒度地控制编译器将如何尝试在加速器上构建工作。所以它不会严重依赖编译器自动并行化代码的能力。

4. 在访问相同数据的两个后续循环中使用并行循环时，编译器可能会也可能不会在两个循环之间的主机和设备之间来回复制数据。

5. 更有经验的并行程序员可能已经在他们的代码中识别出并行循环，他们可能会发现并行循环方法更可取。

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

内核构造

1. 定义应编译为内核序列以在加速器设备上执行的程序区域。

2. 关于内核构造需要注意的重要一点是，编译器将分析代码并仅在确定这样做安全时才进行并行化。在某些情况下，编译器在编译时可能没有足够的信息来确定循环是否安全并行化，在这种情况下它不会并行化循环，即使程序员可以清楚地看到循环是安全并行的。

3. kernels 结构为编译器提供了最大的余地来并行化和优化它认为适合目标加速器的代码，但它也很大程度上依赖于编译器自动并行化代码的能力。

4. 内核构造提供的一个更显着的好处是，如果多个循环访问相同的数据，它只会被复制到加速器一次，这可能会导致更少的数据移动。

5. 并行编程经验较少或代码包含大量需要分析的循环的程序员可能会发现内核方法要简单得多，因为它给编译器带来了更多负担。

例如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];
}

 生成两个内核

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

Answer 2

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

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

详情请见：https://www.pgroup.com/lit/articles/insider/v4n2a1.htm

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