title: OpenACC sum reduction outputs increments sum at each execution
body: <p>Why does the following code:</p>
  
  <pre><code>#include &lt;iostream&gt;
  
  int main(int argc, char const *argv[])
  {
      int sum = 0;
      int *array;
      array = new int [100];
  
      #pragma acc enter data create(array[0:100],sum)
  
      #pragma acc parallel loop present(array[0:100])
      for (int i = 0; i &lt; 100; ++i)
      {
          array[i] = 1;
      }
  
      #pragma acc parallel loop present(array[0:100],sum) reduction(+:sum)
      for (int i = 0; i &lt; 100; ++i)
      {
          sum += array[i];
      }
  
      #pragma acc exit data delete(array[0:100]) copyout(sum)
  
      std::cout &lt;&lt; sum &lt;&lt; std::endl;
  
      return 0;
  }
  </code></pre>
  
  <p>Return different results at every execution?</p>
  
  <pre><code>$ pgcpp -acc -Minfo main.cpp
  main:
        7, Generating enter data create(sum)
           Generating enter data create(array[:100])
           Generating present(array[:100])
           Accelerator kernel generated
           12, #pragma acc loop gang, vector(256) /* blockIdx.x threadIdx.x */
        7, Generating Tesla code
       15, Generating present(array[:100])
           Generating present(sum)
           Accelerator kernel generated
           18, #pragma acc loop gang, vector(256) /* blockIdx.x threadIdx.x */
           20, Sum reduction generated for sum
       15, Generating Tesla code
       25, Generating exit data copyout(sum)
           Generating exit data delete(array[:100])
  $ ./a.out
  100
  $ ./a.out
  200
  $ ./a.out
  300
  $ ./a.out
  400
  </code></pre>
  
  <p>According to the OpenACC standard: </p>
  
  <blockquote>
    <p>On an exit data directive, the data is copied back to the local memory and 
    deallocated. </p>
  </blockquote>
  
  <p>It would seem that <code>sum</code> is <strong>not</strong> deallocated and instead re-used (and incremented) at every run of the program. Furthermore, the <code>+</code> operator in the <code>reduction</code> directive initializes the reduction variable to <code>0</code>, so this shouldn't happen even if <code>sum</code> weren't deallocated between executions. </p>
  
  <p>I can avoid this behavior by either using <code>copyin</code> instead of <code>create</code> for <code>sum</code> in the <code>enter data</code> directive, or setting <code>sum = 0</code> in a single gang, single worker kernel:</p>
  
  <pre><code>#pragma acc parallel present(sum) num_gangs(1) num_workers(1)
  sum = 0;
  </code></pre>
  
  <p>But this is not satisfactory as it requires either a costly host to device data copy, respectively a kernel launch. Why is my program behaving thusly?</p>