OpenACC:对数组进行归约操作
OpenACC: reduction operation on arrays
是否可以对静态数组使用“或”归约运算flag[nx3_tot][nx2_tot][nx1_tot]
?
这里 FLAG_HLL=4
和 FLAG_MINMOD=1
。如果没有减少,这个函数的结果 运行 或没有 OpenACC 是不同的,因为行:
flag[k][j][i+1] |= FLAG_MINMOD;
而且我不明白为什么:使用“或”运算符 |
即使不同的线程访问相同的内存地址,我也希望得到正确的结果。但由于是这种情况,我想使用减少条款,但我收到消息:
此变量数据类型不支持缩减类型 - 标记 flag_shock.c
需要无量纲数组的大小flag_shock.c
int flag[nx3_tot][nx2_tot][nx1_tot];
#pragma acc enter data create(flag[:nx3_tot][:nx2_tot][:nx1_tot])
#pragma acc parallel loop collapse(3) \
present(d, grid, pt[:nx3_tot][:nx2_tot][:nx1_tot]) reduction(||:flag)
for (k = INCLUDE_KDIR; k < nx3_tot-INCLUDE_KDIR; k++){
for (j = INCLUDE_JDIR; j < nx2_tot-INCLUDE_JDIR; j++){
for (i = INCLUDE_IDIR; i < nx1_tot-INCLUDE_IDIR; i++){
double divv, gradp, pt_min;
double dpx1, pt_min1, dvx1;
double dpx2, pt_min2, dvx2;
double dpx3, pt_min3, dvx3;
pt_min = pt[k][j][i];
DIM_EXPAND(pt_min1 = MIN(pt[k][j][i+1], pt[k][j][i-1]); ,
pt_min2 = MIN(pt[k][j+1][i], pt[k][j-1][i]); ,
pt_min3 = MIN(pt[k+1][j][i], pt[k-1][j][i]); )
DIM_EXPAND(pt_min = MIN(pt_min, pt_min1); ,
pt_min = MIN(pt_min, pt_min2); ,
pt_min = MIN(pt_min, pt_min3);)
DIM_EXPAND(dpx1 = fabs(pt[k][j][i+1] - pt[k][j][i-1]); ,
dpx2 = fabs(pt[k][j+1][i] - pt[k][j-1][i]); ,
dpx3 = fabs(pt[k+1][j][i] - pt[k-1][j][i]);)
gradp = DIM_EXPAND(dpx1, + dpx2, + dpx3);
if (gradp > EPS_PSHOCK_FLATTEN*pt_min)
{
flag[k][j][i] |= FLAG_HLL;
flag[k][j][i] |= FLAG_MINMOD;
flag[k][j][i+1] |= FLAG_MINMOD;
flag[k][j][i-1] |= FLAG_MINMOD;
flag[k][j-1][i] |= FLAG_MINMOD;
flag[k][j+1][i] |= FLAG_MINMOD;
}
}}}
我不建议在这种情况下使用缩减。每个线程都需要它自己的数组私有副本,考虑到它是 3D,这将需要相当多的内存,此外还有执行最终归约的开销。
相反,我建议使用原子操作。由于元素上的碰撞很少,因此开销会很低。
if (gradp > EPS_PSHOCK_FLATTEN*pt_min)
{
#pragma acc atomic update
flag[k][j][i] |= FLAG_HLL;
#pragma acc atomic update
flag[k][j][i] |= FLAG_MINMOD;
#pragma acc atomic update
flag[k][j][i+1] |= FLAG_MINMOD;
.. etc ...
您没有提及您使用的是哪个编译器或编译器版本。如果使用 NVHPC,最近添加了数组缩减支持,因此如果使用较旧的编译器,这就是为什么您会收到不支持此数据类型的消息。
是否可以对静态数组使用“或”归约运算flag[nx3_tot][nx2_tot][nx1_tot]
?
这里 FLAG_HLL=4
和 FLAG_MINMOD=1
。如果没有减少,这个函数的结果 运行 或没有 OpenACC 是不同的,因为行:
flag[k][j][i+1] |= FLAG_MINMOD;
而且我不明白为什么:使用“或”运算符 |
即使不同的线程访问相同的内存地址,我也希望得到正确的结果。但由于是这种情况,我想使用减少条款,但我收到消息:
此变量数据类型不支持缩减类型 - 标记 flag_shock.c
需要无量纲数组的大小flag_shock.c
int flag[nx3_tot][nx2_tot][nx1_tot];
#pragma acc enter data create(flag[:nx3_tot][:nx2_tot][:nx1_tot])
#pragma acc parallel loop collapse(3) \
present(d, grid, pt[:nx3_tot][:nx2_tot][:nx1_tot]) reduction(||:flag)
for (k = INCLUDE_KDIR; k < nx3_tot-INCLUDE_KDIR; k++){
for (j = INCLUDE_JDIR; j < nx2_tot-INCLUDE_JDIR; j++){
for (i = INCLUDE_IDIR; i < nx1_tot-INCLUDE_IDIR; i++){
double divv, gradp, pt_min;
double dpx1, pt_min1, dvx1;
double dpx2, pt_min2, dvx2;
double dpx3, pt_min3, dvx3;
pt_min = pt[k][j][i];
DIM_EXPAND(pt_min1 = MIN(pt[k][j][i+1], pt[k][j][i-1]); ,
pt_min2 = MIN(pt[k][j+1][i], pt[k][j-1][i]); ,
pt_min3 = MIN(pt[k+1][j][i], pt[k-1][j][i]); )
DIM_EXPAND(pt_min = MIN(pt_min, pt_min1); ,
pt_min = MIN(pt_min, pt_min2); ,
pt_min = MIN(pt_min, pt_min3);)
DIM_EXPAND(dpx1 = fabs(pt[k][j][i+1] - pt[k][j][i-1]); ,
dpx2 = fabs(pt[k][j+1][i] - pt[k][j-1][i]); ,
dpx3 = fabs(pt[k+1][j][i] - pt[k-1][j][i]);)
gradp = DIM_EXPAND(dpx1, + dpx2, + dpx3);
if (gradp > EPS_PSHOCK_FLATTEN*pt_min)
{
flag[k][j][i] |= FLAG_HLL;
flag[k][j][i] |= FLAG_MINMOD;
flag[k][j][i+1] |= FLAG_MINMOD;
flag[k][j][i-1] |= FLAG_MINMOD;
flag[k][j-1][i] |= FLAG_MINMOD;
flag[k][j+1][i] |= FLAG_MINMOD;
}
}}}
我不建议在这种情况下使用缩减。每个线程都需要它自己的数组私有副本,考虑到它是 3D,这将需要相当多的内存,此外还有执行最终归约的开销。
相反,我建议使用原子操作。由于元素上的碰撞很少,因此开销会很低。
if (gradp > EPS_PSHOCK_FLATTEN*pt_min)
{
#pragma acc atomic update
flag[k][j][i] |= FLAG_HLL;
#pragma acc atomic update
flag[k][j][i] |= FLAG_MINMOD;
#pragma acc atomic update
flag[k][j][i+1] |= FLAG_MINMOD;
.. etc ...
您没有提及您使用的是哪个编译器或编译器版本。如果使用 NVHPC,最近添加了数组缩减支持,因此如果使用较旧的编译器,这就是为什么您会收到不支持此数据类型的消息。