title: Block reduction in CUDA
body: <p>I am trying to do reduction in CUDA and I am really a newbie. I am currently studying a sample code from NVIDIA.</p>
  
  <p>I guess I am really not sure how to set up the block size and grid size, especially when my input array is larger (<code>512 X 512</code>) than a single block size. </p>
  
  <p>Here is the code.</p>
  
  <pre><code>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) 
      {
          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];
      }
  
      if (tid == 0) g_odata[blockIdx.x] = sdata[0];
  }
  </code></pre>
  
  <p>However, it seems to me the <code>g_odata[blockIdx.x]</code> saves the partial sums from all blocks, and, if I want to get the final result, I need to sum all the terms within the <code>g_odata[blockIdx.x]</code> array. </p>
  
  <p>I am wondering: is there a kernel to do the whole summation? or am I misunderstanding things here? I would really appreciate if anyone can educate me with this. Thanks very much.</p>