OpenACC:对数组进行归约操作

OpenACC: reduction operation on arrays

是否可以对静态数组使用“或”归约运算flag[nx3_tot][nx2_tot][nx1_tot]? 这里 FLAG_HLL=4FLAG_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,最近添加了数组缩减支持,因此如果使用较旧的编译器,这就是为什么您会收到不支持此数据类型的消息。