OpenACC: nested loops with private variable

Question

I want to accelerate these nested loops. Because of the dimension of v (NMAX=MAX(NX1, NX2, NX3)), I understand that can be a conflict in the parallelization of the two external loops. I tried to use the private clause:

  static double **v;

  if (v == NULL) {
    v = ARRAY_2D(NMAX_POINT, NVAR, double);
  }
  
  #pragma acc parallel loop present(V, U) private(v[:NMAX_POINT][:NVAR])
  for (k = kbeg; k <= kend; k++){ g_k = k;
  #pragma acc loop
  for (j = jbeg; j <= jend; j++){ g_j = j;

    #pragma acc loop collapse(2)
    for (i = ibeg; i <= iend; i++) {  
    for (nv = 0; nv < NVAR; nv++){ 
      v[i][nv] = V[nv][k][j][i];
    }}

    #pragma acc routine(PrimToCons) seq
    PrimToCons (v, U[k][j], ibeg, iend);
  }}

I get these errors:

Generating present(V[:][:][:][:],U[:][:][:][:])

     Generating Tesla code

    144, #pragma acc loop seq

    146, #pragma acc loop seq

    151, #pragma acc loop gang, vector(128) collapse(2) /* blockIdx.x threadIdx.x */

    154,   /* blockIdx.x threadIdx.x collapsed */

144, Accelerator restriction: induction variable live-out from loop: g_k

     Complex loop carried dependence of v->-> prevents parallelization

146, Accelerator restriction: induction variable live-out from loop: g_j

     Loop carried dependence due to exposed use of v prevents parallelization

     Complex loop carried dependence of V->->->->,v->-> prevents parallelization

g_k and g_j are extern int. I've never seen the message "induction variable live-out from loop" before.

EDIT: I modified the loop as suggested but it sill doesn't work

#pragma acc parallel loop collapse(2) present(U, V) private(v[:NMAX_POINT][:NVAR])
  for (k = kbeg; k <= kend; k++){
  for (j = jbeg; j <= jend; j++){

    #pragma acc loop collapse(2)
    for (i = ibeg; i <= iend; i++) {  
    for (nv = 0; nv < NVAR; nv++){ 
      v[i][nv] = V[nv][k][j][i];
    }}
    PrimToCons (v, U[k][j], ibeg, iend, g_gamma);
  }}

I get this error:

Failing in Thread:1
call to cuStreamSynchronize returned error 700: Illegal address during kernel execution

It's as if the compiler cannot find v, U or V but in the main function I use these directives:

#pragma acc enter data copyin(data)
 #pragma acc enter data copyin(data.Vc[:NVAR][:NX3_TOT][:NX2_TOT][NX1_TOT], data.Uc[:NX3_TOT][:NX2_TOT][NX1_TOT][:NVAR])

data.Vc and data.Uc are V and U in this routine I want to parallelize.

Mat Colgrove Mat Colgrove · Accepted Answer · 2020-07-24T14:45:30

g_k and g_j are extern int. I've never seen the message "induction variable live-out from loop" before.

When run in parallel, the order in which the loop iterations are executed is non-deterministic. Hence the values of g_k and g_j once exiting the loop would be what ever iteration that happens to be last. This creates a dependency since in order to get correct answers (i.e. answers that would agree with those when running serially), the "k" and "j" loops must be run sequentially.

If "g_k" and "g_j" were local variables, then the compiler would implicitly privatize them in order to remove this dependency. However since they are global variables, it must assume other portions of the code uses the results and hence can't assume they can be made private. If you know the variables aren't used elsewhere, then you can fix this issue by adding them to your "private" clause. Note, it doesn't appear that these variables are used the loop itself so could be removed and just assigned the values "kend" and "jend" outside of the loop.

Unless "g_k" and "g_j" are used in the "PrimToCons" subroutine? In that case, you have a bigger problem in that this would cause a race condition in that the variables values may be updated by other threads and no longer be the value expected by the subroutine. In this case, the fix would be to pass "k" and "j" as arguments to "PrimToCons" and not use "g_k" and "g_j".

As for "v", it should be private to the "j" loop as well, not just the "k" loop. To fix, the I'd recommend adding a "collapse(2)" clause to the "k" loop's pragma and remove the loop directive about the "j" loop.

OpenACC: nested loops with private variable

1 Answers