共享变量和 OpenACC

Shared variables and OpenACC

在 OpenMP 中,可以通过

在循环中使用共享变量
#pragma omp parallel for shared(foo) private(bar)

在 OpenACC 中我们有一个 private 子句,但没有 shared 子句。另一方面有copycopyincopyout等数据子句。

有时,我们使用加速器,它们有自己的私有内存,但也可以访问公共内存。

在这种情况下,我们可能希望加速器避免将数据复制到自己的私有内存中,而是对公共内存中的实例进行操作。

我们如何告诉 OpenACC 不要复制数据?

请注意,默认情况下数组是共享的。

"create" 数据子句将在设备上创建数据但不执行复制。

如果你想使用已经在设备上创建的数据,例如通过调用 cudaMalloc 或 acc_malloc,你可以使用 "deviceptr" 数据子句告诉编译器使用指针的设备代码中的地址,而不是使用主机地址在当前 table 中查找设备指针。

如果要将设备变量与主机变量相关联,可以使用 API 调用 "acc_map_data"。

听起来你有更大的内存池,然后你想重新使用。在这种情况下,您可以使用 OpenACC 数据子句、cudaMalloc 或 acc_malloc 创建内存池。如果使用数据子句,则调用 "acc_deviceptr" 以获取设备指针地址。接下来您可以使用 "acc_map_data" 将主机指针与设备指针相关联。请注意,映射数据可以是较大设备池的子集,您可以映射到偏移量,即 "devptr+offset"。

有关使用 "acc_map_data" 的示例,请参阅:https://github.com/rmfarber/ParallelProgrammingWithOpenACC/blob/master/Chapter05/acc_map.c