Understanding CUDA shfl instruction
<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>