title: Using OpenACC to parallelize nested loops
body: <p>I am very new to openacc and have just high-level knowledge so any help and explanation of what I am doing wrong would be appreciated.</p>
  
  <p>I am trying to accelerate(parallelize) a not so straightforward nested loop that updates a flattened (3D to 1D) array using openacc directives. I have posted a simplified sample code below that when compiled using </p>
  
  <p><code>pgcc -acc -Minfo=accel test.c</code></p>
  
  <p>gives the following error:</p>
  
  <p><code>call to cuStreamSynchronize returned error 700: Illegal address during kernel execution</code></p>
  
  <p>Code:</p>
  
  <pre><code>#include &lt;stdio.h&gt;
  #include &lt;stdlib.h&gt;
  
  #define min(a,b) (a &gt; b) ? b : a
  #define max(a,b) (a &lt; b) ? b : a
  
  #define NX 10
  #define NY 10
  #define NZ 10
  
  struct phiType {
    double dx, dy, dz;
    double * distance;
  };
  
  typedef struct phiType Phi;
  
  #pragma acc routine seq
  double solve(Phi *p, int index) {
    // for simplicity just returning a value
    return 2;
  }
  
  void fast_sweep(Phi *p) {
  
    // removing boundaries
    int x = NX - 2; 
    int y = NY - 2;
    int z = NZ - 2;
  
    int startLevel = 3;
    int endLevel   = x + y + z;
  
    #pragma acc data copy(p-&gt;distance[0:NX*NY*NZ])
    for(int level = startLevel; level &lt;= endLevel; level++){
      int ks = max(1, level-(y + z));
      int ke = min(x, level-2);
  
      int js = max(1, level-(x + z));
      int je = min(y, level-2);
  
      #pragma acc region
      {
        #pragma acc loop independent
        for(int k = ks; k &lt;= ke; k++){
          #pragma acc loop independent
          for(int j = js; j &lt;= je; j++){
            int i = level - (k + j);
            if(i &gt; 0 &amp;&amp; i &lt;= z){
              int index = i * NX * NY + j * NX + k;
              p-&gt;distance[index] = solve(p, index);
            }
          }
        }
      }
    }
  }
  
  
  void create_phi(Phi *p){
  
    p-&gt;dx = 1;
    p-&gt;dy = 1;
    p-&gt;dz = 1;
  
    p-&gt;distance = (double *) malloc(sizeof(double) * NX * NY * NZ);
    for(int i = 0; i &lt; NZ; i++){
      for(int j = 0; j &lt; NY; j++){
        for(int k = 0; k &lt; NX; k++){
          int index = i * NX * NY + j * NX + k;
          p-&gt;distance[index] = (i*j*k == 0) ? 0 : 1;
        }
      }
    }
  
  }
  
  
  int main()
  {
    printf("start \n");
  
    Phi *p = (Phi *) malloc(sizeof(Phi));
    create_phi(p);
  
    printf("calling fast sweep \n");
    fast_sweep(p);
  
    printf(" print the results \n");
    for(int i = 0; i &lt; NZ; i++){
      for(int j = 0; j &lt; NY; j++){
        for(int k = 0; k &lt; NX; k++){
          int index = i * NX * NY + j * NX + k;
          printf("%f ", p-&gt;distance[index]);
        }
        printf("\n");
      }
      printf("\n");
    }
  
    return 0;
  }
  </code></pre>
  
  <p>Instead of using the <code>region</code> and <code>loop</code> directives, using the </p>
  
  <pre><code>#pragma acc kernels
  </code></pre>
  
  <p>produces the following error:</p>
  
  <pre><code>solve:
       19, Generating acc routine seq
  fast_sweep:
       34, Generating copy(p-&gt;distance[:1000])
       42, Generating copy(p[:1])
       45, Loop carried dependence due to exposed use of p[:1] prevents parallelization
           Accelerator scalar kernel generated
       47, Loop carried dependence due to exposed use of p[:i1+1] prevents parallelization
  </code></pre>
  
  <p>I am running this code on</p>
  
  <pre><code>GNU/Linux
  CentOS release 6.7 (Final)
  GeForce GTX Titan
  pgcc 15.7-0 64-bit target on x86-64 Linux -tp sandybridge 
  </code></pre>