如何将 OpenACC 数据区域用于结构化数据？答案

【问题标题】：How to use OpenACC data region for structured data?如何将 OpenACC 数据区域用于结构化数据？
【发布时间】：2020-08-08 17:28:03
【问题描述】：

我是 OpenACC 的新手。我有一个使用#pragma acc data 指令与结构化数据传输相关的查询。根据网站https://docs.computecanada.ca/wiki/OpenACC_Tutorial_-_Data_movement

data 指令定义了一个代码区域，其中 GPU 数组保留在 GPU 上并在该区域中的所有内核之间共享。

我明白copy 子句的用法。我想知道这个指令是否可以在没有任何子句的情况下使用？

我阅读了 OpenACC 2.7 规范。这部分不清楚子句是否强制。我的理解是，如果在未明确指定任何数据的情况下定义数据区域，则该区域内使用的所有数据将在整个数据区域内隐式保留在 GPU 上。

#pragma acc data
{
    #pragma acc kernels 
    // Kernel 1

    #pragma acc kernels
    // Kernel 2
}

这意味着，对于上述代码，内核 1 和内核 2 中使用的所有数据将在数据区域的整个持续时间内保留在 GPU 上。

如果我错了，请纠正我。

提前谢谢你。

【问题讨论】：

【解决方案1】：

作为并行构造的一部分存在隐式数据区域（即，作为“并行”或“内核”区域一部分的数据区域），编译器将尝试隐式地将数据复制到设备，假设大小和数据的形状是已知的。否则，您确实需要使用数据子句来定义形状和大小。

对于其他数据区域构造（结构化、非结构化和声明），您确实需要将设备上所需的变量包含在数据子句中，其中数据子句可能是复制、复制、复制、创建、 present 或 deviceptr（或 delete 用于退出数据指令）。编译器无法假定您想要设备上的数据，因此一般不会为您隐式复制它。

【讨论】：

谢谢你的回答。我正在寻找的是编译器可以通过任何方式隐式决定并保留 GPU 上的数据以供多个内核使用。我知道我可以通过使用 copyin 和 coyout 子句来实现这一点。我在想是否可以隐式完成。谢谢。
编译器无法查看整个程序，因此无法真正隐式决定何时传输以及传输多少数据。它可以用于单个计算区域（如果数据的大小和形状已知），但不能用于更大的区域。如果您使用带有 PGI 编译器的 NVIDIA 设备，您可以考虑使用 CUDA 统一内存（通过 -ta=tesla:managed 标志启用）。使用统一内存，CUDA 驱动程序将管理动态分配数据的传输，因此您根本不需要使用数据区域。虽然静态和堆栈数据仍然需要用户管理。