我已经观看了录制的 openacc 概述课程视频，直到第 3 课，其中谈到了表达数据移动。如何使用copy_in() 将gsl_matrix* 从 cpu 移动到 gpu。例如在 CPU 上我可以做类似的事情，
gsl_matrix *Z = gsl_matrix_calloc(100, 100),
这会给我一个 100x100 的零矩阵。现在 Z 是一个指向 gsl_matrix 结构的指针，看起来像，
typedef struct{
size_t size1;
size_t size2;
size_t tda;
double * data;
gsl_block * block;
int owner;
} gsl_matrix;
如何使用copyin() 表达 Z（它是一个指针）从 CPU 到 GPU 的数据移动？
【问题讨论】：
我不能直接谈论在 OpenACC 数据和计算区域中使用 GSL，但可以为您提供有关具有动态数据成员的聚合类型的一般性答案。
假设您使用 PGI 编译器和较新的 NVIDIA 设备，首先要尝试的是 CUDA 统一内存 (UVM)。使用标志“-ta=tesla:managed”编译，所有动态分配的数据都将由 CUDA 运行时管理，因此您无需自己管理数据移动。涉及开销和警告，但它使事情更容易开始。请注意，与 PGI 16.9 或更高版本一起提供的 CUDA 8.0 提高了 UVM 性能。
如果没有 UVM，您需要对数据执行手动深度复制。下面是您首先在设备上创建父结构并执行浅拷贝的基本思想。接下来在设备上创建动态数组“数据”，将初始值复制到数组，然后将数据的设备指针附加到设备结构的数据指针。由于“块”本身是一个具有动态数据成员的结构数组，因此您需要遍历该数组，在设备上创建它的数据数组。
matrix * mat = (matrix*) malloc(sizeof(matrix));
#pragma acc enter data copyin(mat)
// Change this to the correct size of "data" and blocks
#pragma acc enter data copyin(mat.data[0:dataSize]);
#pragma acc enter data copyin(mat.block[0:blockSize]);
for (i=0; i < blockSize; ++i) {
#pragma acc enter data copyin(mat.block[i].data[0:mat.block[i].size])
}
要删除，再次遍历结构，自下而上删除
for (i=0; i < blockSize; ++i) {
#pragma acc exit data delete(mat.block[i].data)
}
#pragma acc exit data delete(mat.block);
#pragma acc exit data delete(mat.data);
#pragma acc exit data delete(mat);
更新时，请确保只更新基本数据类型的标量或数组。即，更新“数据”而不是“块”。更新会进行浅拷贝，因此更新“块”会更新主机或设备指针，从而导致非法地址。
最后，在计算区域中使用矩阵变量时，请务必将其放在“present”子句中。
#pragma acc parallel loop present(mat)
【讨论】：