2014-09-23 116 views
1

在此代码:外部函数使用的设备内存变量如何?

#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 

我不明白为什么指定valuepresent导致上述故障。我使用NVVP进行了检查,value仅在enter data指令中被复制一次,即它不会在intfunparallel指令中再次被复制。 OpenACC如何发挥其魔力?

回答

4

你已经被自己的语法和已经指出的内容再次弄糊涂了。

intfunvalue是不一样的value您在maincopyin(value)。该函数调用通过值value,意味着它使副本它。因此将其添加到present()子句是没有意义的,因为它在设备上的不存在。编译器必须将其复制。(并且,如果根本没有提及它,编译器会自动识别它是否需要并将其复制到您的目录中。)

存在的项目在设备上是main -scope变量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; 
} 
[[email protected] 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子句是合法和功能。

+0

然后我不明白为什么在我的起源例子中,如果我省略'main'中编译指令中的'copyin(value)'和'delete(value)','value'不会复制到设备上, 。设备从哪里检索29的值?这存储在主机内存中,因此必须在某个点上将其复制到设备内存中? – lodhb 2014-09-24 19:07:36

+1

我已经在我的回答中说过,如果你省略了任何提及'value'的东西,编译器会自动识别内核需要'value',它会自动为它生成所需的副本。只有当你指定'present'时才会覆盖这个行为,那么事情就会破裂。如果您有一个新问题,即您不明白的其他问题,它基于与问题中显示的代码不同的代码,那么您可能会提出一个新问题。 'main'中的'copyin(value)'与你的原始代码无关,你的原始代码在没有它的情况下工作得很好。 – 2014-09-24 19:20:59

+0

我已经用NVVP验证过,在'main'中省略''copyin''的值'导致没有主机到设备的数据传输。因此,我对设备从哪里获得“价值”的价值感到困惑。 – lodhb 2014-09-24 19:26:44

相关问题