cuda - 内核构造和并行构造之间的区别

Question

我研究了很多文章和 OpenACC 的手册，但我仍然不明白这两种结构的主要区别。

Answer 1

kernels指令是更一般的情况，如果您以前编写过 GPU（例如 CUDA）内核，您可能会想到这种情况。 kernels简单地指示编译器处理一段代码，并生成任意数量的“内核”，任意“维度”，按顺序执行，以将特定代码部分并行化/卸载到加速器。该parallel构造允许更细粒度地控制编译器将如何尝试在加速器上构建工作，例如通过指定并行化的特定维度。例如，作为parallel指令的一部分，工人和帮派的数量通常是恒定的（因为通常只暗示一个潜在的“内核”），kernels

这篇 PGI 文章中包含对这个特定问题的良好处理。

引用文章摘要：“OpenACC 内核和并行结构都试图解决相同的问题，识别循环并行性并将其映射到机器并行性。内核结构更加隐含，让编译器更自由地查找和映射并行性到目标加速器的要求。并行构造更明确，需要程序员进行更多分析以确定何时合法和合适。”

Answer 2

OpenACC 指令和 GPU 内核只是表示同一事物的两种方式——一段可以并行运行的代码。

当改造现有应用程序以利用 GPU 和/或希望让编译器处理与内存管理等问题相关的更多细节时，OpenACC 可能是最佳选择。这可以加快编写应用程序的速度，但可能会降低性能。

从头开始编写 GPU 应用程序和/或需要更细粒度的控制时，内核可能是最好的。这可能会使应用程序花费更长的时间来编写，但可能会提高性能。

我认为刚接触 GPU 的人可能会倾向于使用 OpenACC，因为它看起来更熟悉。但我认为实际上最好另辟蹊径，从编写内核开始，然后可能转向 OpenACC 以节省一些项目的时间。原因是 OpenACC 是一个泄漏的抽象。因此，虽然 OpenACC 可能会让 GPU 细节看起来像是抽象出来的，但它们仍然存在。因此，在不了解后台发生的情况下使用 OpenACC 编写 GPU 代码可能会令人沮丧，在尝试编译时会出现奇怪的错误消息，并导致应用程序性能低下。

Answer 3

并行构造

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

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

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

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

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

例如参考

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

例如参考

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

 生成两个内核

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