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

  checkCUDA( cudaEventSynchronize( stop ) );
  checkCUDA( cudaEventElapsedTime( &elapsedTime, start, stop ) );
  printf("CPU total time is %f ms, bandwidth %f GB/s\n", elapsedTime,
    sizeof(floatType_t) * (double) size /
    ( (double) elapsedTime / 1000.0 ) * 1.e-9);


  floatType_t diff = abs( good_out - h_out );

  if( diff / h_out < 0.001 ) printf("PASS\n");
  else

#include "../debug.h"

#define N ( 1 << 26 )
#define FLOATTYPE_T double

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;

/* get GPU device number and name */
