为什么下面的代码:
#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
没有在执行之间释放,也不应该发生这种情况。
我可以通过在指令中使用copyin
而不是create
for来避免这种行为,或者在单个 gang、单个工作内核中设置:sum
enter data
sum = 0
#pragma acc parallel present(sum) num_gangs(1) num_workers(1)
sum = 0;
但这并不令人满意,因为它需要昂贵的主机到设备数据复制,或者内核启动。为什么我的程序会这样?