OpenACC:数组的约简操作

Question

是否可以使用缩减操作"或"对静态阵列flag[nx3_tot][nx2_tot][nx1_tot]?这里是FLAG_HLL=4和FLAG_MINMOD=1。如果没有减少，这个函数使用或不使用OpenACC运行的结果是不同的，因为以下行:

flag[k][j][i+1] |= FLAG_MINMOD;

和我不明白为什么:有"或"操作符|我希望有正确的结果，即使不同的线程访问相同的内存地址。但由于这种情况，我想使用还原子句，但我得到的消息:

减少类型不支持这个变量datatype - flag flag_shock.c

Sizeof dimensionless array required 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;
}
}}}

Answer 1

我不建议在这种情况下使用缩减。每个线程都需要它自己的数组私有副本，这将占用相当多的内存，因为它是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，最近添加了数组缩减支持，所以如果使用旧的编译器，这就是为什么你会得到不支持此数据类型的消息。