title: Understanding CUDA shfl instruction
body: <p>I've read the <a href="http://on-demand.gputechconf.com/gtc/2013/presentations/S3174-Kepler-Shuffle-Tips-Tricks.pdf" rel="nofollow">Shuffle Tips and Tricks</a> paper, but I'm not sure exactly how to apply it to some dodgy code I have inherited:</p>
  
  <pre><code>extern __shared__ unsigned int lpSharedMem[];
  int tid = threadIdx.x;
  lpSharedMem[tid] = startValue;
  volatile unsigned int *srt = lpSharedMem;
  
  // ...various stuff
  srt[tid] = min( srt[tid], srt[tid+32] );
  srt[tid] = min( srt[tid], srt[tid+16] );
  srt[tid] = min( srt[tid], srt[tid+8] );
  srt[tid] = min( srt[tid], srt[tid+4] );
  srt[tid] = min( srt[tid], srt[tid+2] );
  srt[tid] = min( srt[tid], srt[tid+1] );
  __syncthreads();
  </code></pre>
  
  <p>Even without CUDA, this code is rather dodgy, but looking at <a href="https://github.com/fixstars/libSGM/blob/master/src/winner_takes_all.cu" rel="nofollow">this implementation</a> I see:</p>
  
  <pre><code>__device__ inline int min_warp(int val) {
      val = min(val, __shfl_xor(val, 16));
      val = min(val, __shfl_xor(val, 8));
      val = min(val, __shfl_xor(val, 4));
      val = min(val, __shfl_xor(val, 2));
      val = min(val, __shfl_xor(val, 1));
      return __shfl(val, 0);
  }
  </code></pre>
  
  <p>This code might be callable with:</p>
  
  <pre><code>int minVal = min_warp(startValue);
  </code></pre>
  
  <p>Thus I can replace my rather dodgy <code>volatile</code> with the above code. However, I cannot really understand what is going on; can someone explain if I am correct, and what exactly is happening in the <code>min_warp()</code> function.</p>