title: Sum reduction with CUDA: What is N?
body: <p>According to NVIDIA, <a href="http://developer.download.nvidia.com/compute/cuda/1_1/Website/projects/reduction/doc/reduction.pdf" rel="noreferrer">this</a> is the fastest sum reduction kernel:</p>
  
  <pre><code>template &lt;unsigned int blockSize&gt;
  __device__ void warpReduce(volatile int *sdata, unsigned 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];
  }
  template &lt;unsigned int blockSize&gt;
  __global__ void reduce6(int *g_idata, int *g_odata, unsigned int n) {
  extern __shared__ int sdata[];
  unsigned int tid = threadIdx.x;
  unsigned int i = blockIdx.x*(blockSize*2) + tid;
  unsigned int gridSize = blockSize*2*gridDim.x;
  sdata[tid] = 0;
  while (i &lt; n) { sdata[tid] += g_idata[i] + g_idata[i+blockSize];  i += gridSize;  }
  __syncthreads();
  if (blockSize &gt;= 512) { if (tid &lt; 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads(); }
  if (blockSize &gt;= 256) { if (tid &lt; 128) { sdata[tid] += sdata[tid + 128]; } __syncthreads(); }
  if (blockSize &gt;= 128) { if (tid &lt;   64) { sdata[tid] += sdata[tid +   64]; } __syncthreads(); }
  if (tid &lt; 32) warpReduce(sdata, tid);
  if (tid == 0) g_odata[blockIdx.x] = sdata[0];
  }
  </code></pre>
  
  <p>However, I don't understand the "n" parameter. Any clues? I don't think it's the size of the array to reduce, since in the while loop there would be a buffer overflow.</p>