title: Can the #pragma unroll statement be used to clean up compiler pre-evaluated code?
body: <p>In NVIDIA's well made <a href="http://developer.download.nvidia.com/assets/cuda/files/reduction.pdf" rel="nofollow">reduction optimization documentation</a>, they end up with a warpReduce that looks like:</p>
  
  <pre><code>Template &lt;unsigned int blockSize&gt;
  __device__ void warpReduce(volatile int* sdata, int tid) {
      if (blockSize &gt;= 64) sdata[tid] += sdata[tid + 32]; 
      if (blockSize &gt;= 32) sdata[tid] += sdata[tid + 16]; 
      if (blockSize &gt;= 16) sdata[tid] += sdata[tid + 8]; 
      if (blockSize &gt;= 8) sdata[tid] += sdata[tid + 4]; 
      if (blockSize &gt;= 4) sdata[tid] += sdata[tid + 2]; 
      if (blockSize &gt;= 2) sdata[tid] += sdata[tid + 1]; 
  }
  </code></pre>
  
  <p>This hurts my sensibility as a proponent of <a href="http://en.wikipedia.org/wiki/Don&#39;t_repeat_yourself" rel="nofollow">Don't Repeat Yourself</a>, why didn't they save programmer time and introduce constructs more like</p>
  
  <pre><code>Template &lt;unsigned int blockSize&gt;
  __device__ void warpReduce(volatile int* sdata, int tid) {
      #pragma unroll(6)
      for (int i = 64; i&gt;1; i&gt;&gt;=1){
          if (blockSize &gt;= i) sdata[tid] += sdata[tid + (i/2)]; 
      }
  }
  </code></pre>
  
  <p>Will this compile to the same code? Or does this confuse the compiler too much to optimize it away? I know the case statement suffers from a similar DRY issue, but I don't know that there's a way to get around that, since it maps runtime inputs to compiler time constant defined template kernels, but I don't see how making programmers need to write out the same thing a fair number of times is helpful.</p>
  
  <p>Not to mention, if the warp size ever changes, which it may not (but who knows) the code I introduced will be much easier to change, and could even use a constant to define the warp size, allowing for one change to change the optimization here, and anywhere else that relies on the warp size.</p>
  
  <p>Very related subquestion, I read using the built-in <code>warpSize</code> is questionable for a compiler optimization reason as well. Otherwise I would have included it in the code above. Going to the definition in <code>device_launch_parameters.h</code>, it calls a built-in device query <code>__cudaGet_warpSize()</code>. If this isn't an issue, NVIDIA could have provided in that document a complete, generalized optimization of <code>warpReduce</code> for reference, but didn't, and I would like to know why? </p>
  
  <p><strong>Tl;dr:</strong> Can I write this as</p>
  
  <pre><code>Template &lt;unsigned int blockSize&gt;
  __device__ void warpReduce(volatile int* sdata, int tid) {
      #pragma unroll(6)
      for (int i = warpSize * 2; i&gt;1; i&gt;&gt;=1){
          if (blockSize &gt;= i) sdata[tid] += sdata[tid + (i/2)]; 
      }
  }
  </code></pre>
  
  <p>and enjoy the same optimization as in their documentation?</p>