OpenACC: reduction operation on arrays

Is it possible to use the reduction operation "or" on the static array flag[nx3_tot][nx2_tot][nx1_tot] ? here FLAG_HLL=4 and FLAG_MINMOD=1. Without the reduction the results of this function running with or without OpenACC are different due to the lines:

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

and I don't understand why: having the "or" operator | I would expect to have correct results even if different threads access the same memory address. But since this is the case I would like to use the reduction clause but I get the message:

Reduction type not supported for this variable 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;
  }

}}}

I wouldn't recommend using reductions in this case. Each thread would need it's own private copy of the array which would be quite a bit of memory given it's 3D, plus there's the overhead of performing the final reduction.

Instead, I'd recommend using atomic operations. Since there will very few collisions on the elements, the overhead would be very low.

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 ...

You don't mention which compiler or compiler version you're using. If using NVHPC, array reduction support was added recently so if using an older compiler, this is why you'd be getting the message that this data type is not supported.