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

  float elapsedTime;
  checkCUDA( cudaEventElapsedTime( &elapsedTime, start, stop ) );

  printf("Total elements is %d, %f GB\n", size, sizeof(floatType_t)*
    (double)size * 1.e-9 );
  printf("GPU total time is %f ms, bandwidth %f GB/s\n", elapsedTime,
    sizeof(floatType_t) * (double) size /
    ( (double) elapsedTime / 1000.0 ) * 1.e-9);

/* copy result back to host */

  checkCUDA( cudaMemcpy( &h_out, d_out, sizeof(floatType_t), 
    cudaMemcpyDeviceToHost ) );

  checkCUDA( cudaEventRecord( start, 0 ) );

  checkCUDA( cudaFree( d_in ) );
  checkCUDA( cudaFree( d_out ) );
  checkCUDA( cudaFree( d_tempArray ) );
  checkCUDA( cudaFree( d_temp_storage ) );

  checkCUDA( cudaDeviceReset() );
	
