OpenACC 2.0 例程：数据局部性答案

【问题标题】：OpenACC 2.0 routine: data localityOpenACC 2.0 例程：数据局部性
【发布时间】：2015-01-08 12:37:30
【问题描述】：

以下面的代码为例，它说明了在加速器上调用一个简单的例程，在设备上使用 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 的设备副本？我们如何强制使用设备副本？

【问题讨论】：

标签： gpgpu openacc

【解决方案1】：

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

如果您想强制使用更多的设备副本，您可以创建例程 nohost，这样从主机调用从技术上讲是无效的，但这并不能真正做到您的目的询问，即在 GPU 上检查该阵列是否真的是一个设备阵列。

请记住，虽然并行区域内的任何代码不在 loop 内，都将运行gang-redundantly，因此写入 out 可能是一种竞争条件，除非你碰巧和一个帮派一起跑，或者你用atomic给它写信。

【讨论】：

【解决方案2】：

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

【讨论】：

【解决方案3】：

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

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

【讨论】：