OpenACC 复制指令如何工作？

Question

根据 OpenACC 文档：

copyin - 为设备上列出的变量创建空间，通过复制初始化变量 数据到设备的区域开始，并释放设备上的空间时 无需将数据复制回主机即可完成。

我已经创建了一个测试示例程序

int main(int argc, char** argv)
{
    int teste[] = { -15 };


    #pragma acc data copyin(teste[0:1])
    {

        #pragma acc parallel loop
        for (int p = 0; p < 5000; p++) {
            teste[0] = p;
        }
    }
    printf("%d", teste[0]);
    return 0;
}

根据文档，程序应该输出-15，因为数据在设备上被修改并且结果不会被复制回主机。但是一旦我编译并运行这段代码，输出是4999

我的编译器是gcc (tdm64-1) 10.3.0，我在一台具有独立设备和主机内存的计算机上运行程序

我想知道为什么这不起作用，我可以做些什么来防止从设备复制回主机。

这是在 windows 上使用 git bash 运行的程序：

$ cat test.c && echo "" &&gcc -fopenacc test.c && ./a.exe

#include <stdio.h>

int main(int argc, char** argv)
{
    int teste[] = { -15 };


    #pragma acc data copyin(teste[0:1])
    {

        #pragma acc parallel loop
        for (int p = 0; p < 5000; p++) {
            teste[0] = p;
        }
    }
    printf("%d\n", teste[0]);
    return 0;
}
4999

我还可以访问 Linux 机器，即使使用 nvc 也无法获得正确的结果

cat test.c && echo "" && /opt/nvidia/hpc_sdk/Linux_x86_64/2021/compilers/bin/nvc -acc -Minfo=accel test.c && ./a.out
#include <stdio.h>

int main(int argc, char** argv)
{
    int teste[] = { -15 };


    #pragma acc data copyin(teste[0:1])
    {

        #pragma acc parallel loop
        for (int p = 0; p < 5000; p++) {
            teste[0] = p;
        }
    }
    printf("%d\n", teste[0]);
    return 0;
}

main:
    9, Generating copyin(teste[:]) [if not already present]
       Generating NVIDIA GPU code
       12, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
4999

Answer 1

程序应该打印 -15，因为主机上的值没有改变。因此，这要么是 gcc 中的错误，要么您实际上并未启用 OpenACC。你使用什么编译器标志？

这是针对 NVIDIA A100 使用 nvc 的输出：

% cat test.c
#include <stdio.h>

int main(int argc, char** argv)
{
    int teste[] = { -15 };


    #pragma acc data copyin(teste[0:1])
    {

        #pragma acc parallel loop
        for (int p = 0; p < 5000; p++) {
            teste[0] = p;
        }
    }
    printf("%d\n", teste[0]);
    return 0;
}
% nvc test.c -acc -Minfo=accel ; a.out
main:
     10, Generating copyin(teste[:]) [if not already present]
         Generating NVIDIA GPU code
         13, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
-15