  */
 
 #include <stdio.h>
-#include <cub/cub.cuh>
+#include <cub/block/block_reduce.cuh>
 #include "../debug.h"
 
 #define N ( 1 << 27 )
+#define THREADS_PER_BLOCK 128
+
 #define FLOATTYPE_T float
+
+__global__ void sumReduction(int n, FLOATTYPE_T *in, FLOATTYPE_T *out)
+{
+
+/* Allocate shared memory for cub::BlockReduce */
+   __shared__ typename cub::BlockReduce<FLOATTYPE_T,
+                THREADS_PER_BLOCK>::TempStorage sArray;
+
+/* calculate global index in the array */
+  int globalIndex = blockIdx.x * blockDim.x + threadIdx.x;
+
+/* grid stride handling case where array is larger than number of threads
+ * launched
+ * Loop over the grid stride so that each thread adds up its relevant 
+ * elements of the array and saves them to a register.
+ */
+
+  FLOATTYPE_T tempResult = 0.0;
+
+  for( int i = globalIndex; i < n; i += blockDim.x * gridDim.x )
+  {
+    tempResult += in[i];
+  } /* end for */
+
+/* Compute the block-wide sum for thread0 */
+   FLOATTYPE_T blockSum = cub::BlockReduce<FLOATTYPE_T,
+                 THREADS_PER_BLOCK>(sArray).Sum(tempResult);
+
+/* write the result back to global memory */
+  if( threadIdx.x == 0 ) out[blockIdx.x] = blockSum;
+
+  return;
+
+}
 
 int main()
 {
-  FLOATTYPE_T *h_in, h_out, good_out;
-  FLOATTYPE_T *d_in, *d_out, *d_tempArray;
+  FLOATTYPE_T *h_in, h_sum, cpu_sum;
+  FLOATTYPE_T *d_in, *d_sum, *d_tempArray;
   int size = N;
   int memBytes = size * sizeof( FLOATTYPE_T );
   int tempArraySize = 32768;

/* allocate space for device copies of in, out */

  checkCUDA( cudaMalloc( &d_in, memBytes ) );
  checkCUDA( cudaMalloc( &d_out, sizeof(FLOATTYPE_T) ) );
  checkCUDA( cudaMalloc( &d_tempArray, tempArraySize * sizeof(FLOATTYPE_T) ) );

/* allocate space for host copies of in, out and setup input values */


/* copy result back to host */

  checkCUDA( cudaMemcpy( &h_sum, d_sum, sizeof(FLOATTYPE_T), 
    cudaMemcpyDeviceToHost ) );

  checkCUDA( cudaEventRecord( start, 0 ) );

  for( int i = 0; i < size; i++ )
  {
    cpu_sum += h_in[i];
  } /* end for */

  checkCUDA( cudaEventRecord( stop, 0 ) );
