0

以下面的代码为例，它说明了在加速器上调用一个简单的例程，使用 OpenACC 2.0 的routine指令在设备上编译：

#include <iostream>

#pragma acc routine
int function(int *ARRAY,int multiplier){
        int sum=0;

        #pragma acc loop reduction(+:sum)
        for(int i=0; i<10; ++i){
                sum+=multiplier*ARRAY[i];
        }

        return sum;
}

int main(){
        int *ARRAY = new int[10];
        int multiplier = 5;
        int out;

        for(int i=0; i<10; i++){
                ARRAY[i] = 1;
        }

        #pragma acc enter data create(out) copyin(ARRAY[0:10],multiplier)

        #pragma acc parallel present(out,ARRAY[0:10],multiplier)
        if (function(ARRAY,multiplier) == 50){
                out = 1;
        }else{
                out = 0;
        }

        #pragma acc exit data copyout(out) delete(ARRAY[0:10],multiplier)

        std::cout << out << std::endl;
}

如何function知道使用设备副本ARRAY[0:10]以及multiplier何时从并行区域内调用它？我们如何强制使用设备副本？

4

3 回答 3

1

当您的例程在设备区域（parallel在您的代码中）被调用时，它被设备上的线程调用，这意味着这些线程将只能访问设备上的数组。编译器实际上可能选择内联该函数，或者它可能是设备端函数调用。这意味着您可以知道，当从设备调用该函数时，它将接收数据的设备副本，因为该函数本质上是present从并行区域继承数据子句。如果您仍然想说服自己在函数内部一次在设备上运行，您可以调用acc_on_device，但这只会告诉您您正在加速器上运行，而不是您收到了设备指针。

如果您想强制使用更多设备副本，您可以制作例程nohost，以便从主机调用从技术上讲是无效的，但这并不能真正满足您的要求，即在 GPU 上检查该阵列确实是一个设备阵列。

请记住，任何不在 a 内的并行区域内的代码loop都将运行gang-redundantly，因此写入out可能是一种竞争条件，除非您碰巧与一个 gang 一起运行，或者您使用atomic.

于 2015-08-10T20:41:15.233 回答
0

基本上，当您涉及“数据”子句时，设备将创建/复制数据到设备内存，然后使用“acc 例程”定义的代码块将在设备上执行。请注意，与多线程 (OpenMP) 不同，主机和设备之间的内存不共享。所以是的，只要“函数”在数据段下，它将使用 ARRAY 和乘法器的设备副本。希望这可以帮助！:)

于 2015-03-09T03:59:47.170 回答
0

您应该为函数分配一个并行级别，例如 gang/worker/vector。这是一种更准确的方法。

该例程将使用设备内存中的日期。

于 2015-04-07T05:50:55.130 回答