title: When to use volatile with shared CUDA Memory
body: <p>Under what circumstances should you use the <code>volatile</code> keyword with a CUDA kernel's shared memory? I understand that <code>volatile</code> tells the compiler never to cache any values, but my question is about the behavior with a shared array:</p>
  
  <pre><code>__shared__ float products[THREADS_PER_ACTION];
  
  // some computation
  products[threadIdx.x] = localSum;
  
  // wait for everyone to finish their computation
  __syncthreads();
  
  // then a (basic, ugly) reduction:
  if (threadIdx.x == 0) {
      float globalSum = 0.0f;
      for (i = 0; i &lt; THREADS_PER_ACTION; i++)
          globalSum += products[i];
  }
  </code></pre>
  
  <p>Do I need <code>products</code> to be volatile in this case? Each array entry is only accessed by a single thread, except at the end, where everything is read by thread 0. Is it possible that the compiler could cache the entire array, and so I need it to be  <code>volatile</code>, or will it only cache elements?</p>
  
  <p>Thanks!</p>