   checkCUDA( cudaEventRecord( start, 0 ) );
 
 /* launch the kernel on the GPU */
-  cub::DeviceReduce::Sum( d_temp_storage, temp_storage_bytes, d_in, d_out, 
-    size );
+
+  sumReduction<<< blocks, threads1 >>>( size, d_in,  d_tempArray );
+  checkKERNEL()
+  sumReduction<<<      1, threads2 >>>( blocks.x, d_tempArray, d_sum );
+  checkKERNEL()
 
 /* stop the timers */
 

 */

#include <stdio.h>
#include <cub/cub.cuh>
#include "../debug.h"

#define N ( 1 << 27 )
#define FLOATTYPE_T float

int main()
{
  FLOATTYPE_T *h_in, h_out, good_out;
  FLOATTYPE_T *d_in, *d_out, *d_tempArray;
  int size = N;
  int memBytes = size * sizeof( FLOATTYPE_T );
  int tempArraySize = 32768;


/* 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 ) );
