title: reduction example using cuda and CUB
body: <p>I'm trying to get my head around CUB, and having a bit of trouble following the (rather incomplete) worked examples. CUB looks like it is a fantastic tool, I just can't make sense of the example code.</p>
  
  <p>I've built a simple proto-warp reduce example:</p>
  
  <pre><code>#include &lt;cub/cub.cuh&gt;
  #include &lt;cuda.h&gt;
  #include &lt;vector&gt;
  using std::vector;
  #include &lt;iostream&gt;
  using std::cout;
  using std::endl;
  
  const int N = 128;
  
  __global__ void sum(float *indata, float *outdata) {
      typedef cub::WarpReduce&lt;float,4&gt; WarpReduce;
      __shared__ typename WarpReduce::TempStorage temp_storage;
      int id = blockIdx.x*blockDim.x+threadIdx.x;
      if( id &lt; 128 ) {
          outdata[id] = WarpReduce(temp_storage).Sum(indata[id]);
      }
  }
  
  int main() {
      vector&lt;float&gt; y(N), sol(N);
      float *dev_y, *dev_sol;
      cudaMalloc((void**)&amp;dev_y,N*sizeof(float));
      cudaMalloc((void**)&amp;dev_sol,N*sizeof(float));
      for( int i = 0; i &lt; N; i++ ) {
          y[i] = (float)i;
      }
      cout &lt;&lt; "input: ";
      for( int i = 0; i &lt; N; i++ ) cout &lt;&lt; y[i] &lt;&lt; " ";
      cout &lt;&lt; endl;
      cudaMemcpy(&amp;y[0],dev_y,N*sizeof(float),cudaMemcpyHostToDevice);
      sum&lt;&lt;&lt;1,32&gt;&gt;&gt;(dev_y,dev_sol);
      cudaMemcpy(dev_sol,&amp;sol[0],N*sizeof(float),cudaMemcpyDeviceToHost);
      cout &lt;&lt; "output: ";
      for( int i = 0; i &lt; N; i++ ) cout &lt;&lt; sol[i] &lt;&lt; " ";
      cout &lt;&lt; endl;
      cudaFree(dev_y);
      cudaFree(dev_sol);
      return 0;
  }
  </code></pre>
  
  <p>which returns all zeros.</p>
  
  <p>I'm aware that this code would return a reduction that was banded with every 32nd element being the sum of a warp and the other elements being undefined - I just want to get a feel for how CUB works. Can someone point out what I'm doing wrong?</p>
  
  <p>(also, does CUB deserve its own tag yet?)</p>