如何在动态结构中使用#pragma acc shape及-ta=tesla:deepcopy?
问题分析与解决方案
报错原因
你遇到的PRESENT clause was not found错误,核心原因是:
#pragma acc enter data copyin(data)仅拷贝了Data结构体本身(也就是两个指针Vc和Uc的主机地址值)到设备,但指针指向的多维数组并未被拷贝到设备内存。- 当
PrimToCons3D中的present(V[:...])检查设备内存时,找不到V对应的设备端数组,因此触发错误。 - 而
-ta=tesla:managed使用统一内存(Unified Memory),主机和设备共享内存空间,指针可直接访问,因此不会出现这个问题。
修复代码:显式拷贝多维数组并关联结构体指针
要解决这个问题,需要先将Vc、Uc指向的多维数组拷贝到设备,再确保设备端结构体的指针指向设备上的数组。具体步骤如下:
1. 修改Initialize函数,显式拷贝数组并关联结构体
int Initialize(Data *data){ // 分配主机端多维数组 data->Uc = ARRAY_4D(ntot[KDIR], ntot[JDIR], ntot[IDIR], NVAR, double); data->Vc = ARRAY_4D(NVAR, ntot[KDIR], ntot[JDIR], ntot[IDIR], double); // 第一步:将Vc和Uc指向的多维数组完整拷贝到设备内存 // 注意维度要和ARRAY_4D的分配顺序完全匹配 #pragma acc enter data copyin(data->Vc[0:NVAR][0:ntot[KDIR]][0:ntot[JDIR]][0:ntot[IDIR]]) #pragma acc enter data copyin(data->Uc[0:ntot[KDIR]][0:ntot[JDIR]][0:ntot[IDIR]][0:NVAR]) // 第二步:拷贝结构体,并将结构体的指针成员关联到设备端数组 // attach子句告诉编译器:Vc/Uc已经在设备上存在对应的内存 #pragma acc enter data copyin(data) attach(data->Vc, data->Uc) PrimToCons3D(data->Vc, data->Uc, grid, NULL); }
2. 确保PrimToCons3D中的present范围与实际分配一致
检查nx3_tot、nx2_tot、nx1_tot是否与分配时的ntot[KDIR]、ntot[JDIR]、ntot[IDIR]完全匹配。如果这些变量是运行时确定的,需要将它们也拷贝到设备:
void PrimToCons3D(double ****V, double ****U, Grid *grid, RBox *box){ // 先拷贝运行时维度变量到设备 #pragma acc enter data copyin(nx3_tot, nx2_tot, nx1_tot, kbeg, kend, jbeg, jend, ibeg, iend) #pragma acc parallel loop collapse(3) present(V[:NVAR][:nx3_tot][:nx2_tot][:nx1_tot]) for (k = kbeg; k <= kend; k++){ for (j = jbeg; j <= jend; j++){ for (i = ibeg; i <= iend; i++){ double v[NVAR]; #pragma acc loop for (nv = 0; nv < NVAR; nv++) v[nv] = V[nv][k][j][i]; } } } }
关于#pragma acc shape的用法
#pragma acc shape的作用是向编译器提示数组的维度信息,帮助编译器生成正确的设备拷贝代码,尤其适用于结构体中的指针成员。对于你这种运行时动态确定维度的场景:
1. 配合运行时变量指定形状
在结构体定义后,使用shape指令关联成员与运行时维度变量(注意变量需要在使用时可见):
typedef struct Data_{ double ****Vc; double ****Uc; } Data; // 指定Vc和Uc的动态维度 #pragma acc shape(Vc: NVAR, ntot[KDIR], ntot[JDIR], ntot[IDIR]) #pragma acc shape(Uc: ntot[KDIR], ntot[JDIR], ntot[IDIR], NVAR)
2. 结合enter data使用
有了shape指令后,拷贝结构体时可以简化数组拷贝的写法,但仍需确保维度变量已被拷贝到设备:
int Initialize(Data *data){ data->Uc = ARRAY_4D(ntot[KDIR], ntot[JDIR], ntot[IDIR], NVAR, double); data->Vc = ARRAY_4D(NVAR, ntot[KDIR], ntot[JDIR], ntot[IDIR], double); // 先拷贝维度变量到设备 #pragma acc enter data copyin(ntot) // 利用shape指令的信息,自动拷贝结构体成员指向的数组 #pragma acc enter data copyin(data) copyin(data->Vc, data->Uc) PrimToCons3D(data->Vc, data->Uc, grid, NULL); }
不过需要注意:shape指令依赖于编译器对维度变量的可见性,若维度变量是局部变量或动态变化的,显式指定数组范围的方式(如前面的修复方案)会更可靠。
内容的提问来源于stack exchange,提问作者Steve
相关产品推荐
相关产品推荐

