OpenACC 减和输出在每次执行时递增总和

Question

为什么会出现以下代码：

#include <iostream>

int main(int argc, char const *argv[])
{
    int sum = 0;
    int *array;
    array = new int [100];

    #pragma acc enter data create(array[0:100],sum)

    #pragma acc parallel loop present(array[0:100])
    for (int i = 0; i < 100; ++i)
    {
        array[i] = 1;
    }

    #pragma acc parallel loop present(array[0:100],sum) reduction(+:sum)
    for (int i = 0; i < 100; ++i)
    {
        sum += array[i];
    }

    #pragma acc exit data delete(array[0:100]) copyout(sum)

    std::cout << sum << std::endl;

    return 0;
}

每次执行都返回不同的结果？

$ pgcpp -acc -Minfo main.cpp
main:
      7, Generating enter data create(sum)
         Generating enter data create(array[:100])
         Generating present(array[:100])
         Accelerator kernel generated
         12, #pragma acc loop gang, vector(256) /* blockIdx.x threadIdx.x */
      7, Generating Tesla code
     15, Generating present(array[:100])
         Generating present(sum)
         Accelerator kernel generated
         18, #pragma acc loop gang, vector(256) /* blockIdx.x threadIdx.x */
         20, Sum reduction generated for sum
     15, Generating Tesla code
     25, Generating exit data copyout(sum)
         Generating exit data delete(array[:100])
$ ./a.out
100
$ ./a.out
200
$ ./a.out
300
$ ./a.out
400

根据 OpenACC 标准：

在退出数据指令中，数据被复制回本地内存并 解除分配。

sum 似乎没有被释放，而是在程序的每次运行时重新使用（并增加）。此外，reduction 指令中的+ 运算符将归约变量初始化为0，因此即使在执行之间没有释放sum，也不应该发生这种情况。

我可以通过在 enter data 指令中使用 copyin 而不是 create 代替 sum 或在单个 gang、单个工作内核中设置 sum = 0 来避免这种行为：

#pragma acc parallel present(sum) num_gangs(1) num_workers(1)
sum = 0;

但这并不令人满意，因为它需要昂贵的主机到设备数据复制，或者内核启动。为什么我的程序会这样？

Answer 1

您误解了归约运算符初始化值的含义。参考openACC specification，第 20-21 页：

并行结构上允许使用归约子句。它指定一个归约运算符和一个或多个标量变量。对于每个变量，为每个并行组创建一个私有副本并为该操作员初始化。在该区域的末端，值为 每个gang用reduction算子组合，将结果与原变量的值组合保存在原变量中。

这意味着整个还原问题被分解成小块，每个小块都由一个帮派处理。 gang 处理的部分问题将使用减少变量的指示初始化值。但是，当最终结果创建时，每个帮派的单个结果将与原始变量的值（在您的情况下为 sum）相结合，这就是结果。

因此，您必须正确初始化 sum，也许使用您在问题中概述的方法之一。

此外，尽管这不是问题的症结所在，但请注意，无论是释放还是分配都不会对内存内容产生任何影响。在该位置分配的新变量，未经适当初始化，将获取该位置当前的值。