title: Why is there a warp-level synchronization primitive in CUDA?
body: <p>I have two questions regarding <code>__syncwarp()</code> in CUDA:</p>
  
  <ol>
  <li>If I understand correctly, a warp in CUDA is executed in an SIMD fasion. Does that not imply that all threads in a warp are always synchronized? If so, what exactly does <code>__syncwarp()</code> do, and why is it necessary?</li>
  <li>Say we have a kernel launched with a block size of 1024, where the threads within a block are divided into groups of 32 threads each. Each thread communicates with other threads in it's group via shared memory, but does not communicate with any thread outside it's group. In such a kernel, I can see how a more granular synchronization than <code>__syncthreads()</code> may be useful, but since the warps the block is split into may not match with the groups, how would one guarantee correctness when using <code>__syncwarp()</code>?</li>
  </ol>