外部函数如何使用设备内存中的变量？

Question

在这段代码中：

#include <iostream>

void intfun(int * variable, int value){
    #pragma acc parallel present(variable[:1]) num_gangs(1) num_workers(1)
    {
        *variable = value;
    }
}

int main(){
    int var, value = 29;

    #pragma acc enter data create(var) copyin(value)
        intfun(&var,value);
    #pragma acc exit data copyout(var) delete(value)

    std::cout << var << std::endl;
}

如何将int value 识别为位于intfun 的设备内存中？如果我在 intfun 杂注中将 present(variable[:1]) 替换为 present(variable[:1],value)，我会收到以下运行时错误：

FATAL ERROR: data in PRESENT clause was not found on device 1: name=_43144_33_value
 file:/opt/pgi/linux86-64/14.9/include/CC/iostream intfun__FPii line:5
Present table dump for device[1]: NVIDIA Tesla GPU 1, compute capability 3.5
host:0x7fffc11faa28 device:0x2303f20200 size:4 presentcount:1 line:14 name:_43152_14_value
host:0x7fffc11faa34 device:0x2303f20000 size:4 presentcount:2 line:14 name:_43152_9_var

我不明白为什么指定 value 是 present 会导致上述失败。我用 NVVP 检查了value 在enter data 指令中只复制一次，即它不会在intfun 的parallel 指令中再次复制。 OpenACC 如何发挥它的魔力？

Answer 1

您再次对自己的语法和已经向您指出的内容感到困惑。

intfun 中的value 与main 中的value 中的copyin(value) 不同。函数调用按值传递value，这意味着它制作了它的副本。因此，将其添加到 present() 子句是没有意义的，因为它在设备上不存在。编译器必须把它拷贝进去。（当你根本不提它时，编译器会自动识别出它是需要的，并为你拷贝进去。）

设备上存在的项目是main-范围变量value，intfun未使用。为所有变量赋予相同的名称可能无法帮助您理解这一点。

为了证明这一点，让我们通过value 通过引用而不是通过值：

$ cat main8.cpp
#include <iostream>

void intfun(int * variable, int &value){
    #pragma acc parallel present(variable[:1],value) num_gangs(1) num_workers(1)
    {
        *variable = value;
    }
}

int main(){
    int var, value = 29;

    #pragma acc enter data create(var) copyin(value)
        intfun(&var,value);
    #pragma acc exit data copyout(var) delete(value)

    std::cout << var << std::endl;
}
[user2@dc12 misc]$ pgcpp -acc -Minfo main8.cpp
intfun(int *, int &):
      5, Generating present(variable[:1])
         Generating present(value[:])
         Accelerator kernel generated
         Generating Tesla code
main:
     14, Generating enter data copyin(value)
         Generating enter data create(var)
     17, Generating exit data delete(value)
         Generating exit data copyout(var)
$ ./a.out
29
$

现在主作用域变量value在设备上，编译器知道了，它会被intfun直接使用，并且把它加到present子句中是合法的和功能性的。