 #include "../debug.h"
 
 #define N ( 1 << 26 )
-#define THREADS_PER_BLOCK 128
-
-typedef double floatType_t;
-
-__global__ void sumReduction(int n, floatType_t *in, floatType_t *out)
-{
-
-  __shared__ floatType_t sArray[THREADS_PER_BLOCK];
-
-/* calculate global index in the array */
-  int globalIndex = blockIdx.x * blockDim.x + threadIdx.x;
-	
-/* zero out the smem array */
-  sArray[threadIdx.x] = 0.0;
-
-/* 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 SMEM
- */
-
-  for( int i = globalIndex; i < n; i += blockDim.x * gridDim.x )
-  {
-    sArray[threadIdx.x] += in[i];
-  } /* end for */
-  
-  __syncthreads();
-
-//  printf("tid %d blockId %d value %f\n",threadIdx.x, blockIdx.x, 
- //   sArray[threadIdx.x] );
-/* do the final reduction in SMEM */
-  //printf("blockDim is %d\n",blockDim.x);
-  for( int i = blockDim.x/2; i > 0; i = i / 2 )
-  {
-    if( threadIdx.x < i )
-    {
-      sArray[threadIdx.x] += sArray[threadIdx.x + i];
-   //   printf("tid %d, i %d, sArray %f\n",threadIdx.x, i,sArray[threadIdx.x] );
-    } /* end if */
-    __syncthreads();
-  } /* end for */
-
-  if( threadIdx.x == 0 ) out[blockIdx.x] = sArray[0]; 
-  //if( threadIdx.x == 0 ) printf("out is %f\n",out[blockIdx.x]);
-
-  return;
-
-}
+#define FLOATTYPE_T double
 
 int main()
 {
-  floatType_t *h_in, h_out, good_out;
-  floatType_t *d_in, *d_out, *d_tempArray;
+  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 memBytes = size * sizeof( FLOATTYPE_T );
   int tempArraySize = 32768;
 
 /* get GPU device number and name */

  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()

/* stop the timers */


/* copy inputs to device */

  checkCUDA( cudaMemcpy( d_in, h_in, memBytes, cudaMemcpyHostToDevice ) );
  checkCUDA( cudaMemset( d_out, 0, sizeof(FLOATTYPE_T) ) );
  checkCUDA( cudaMemset( d_tempArray, 0, 
    tempArraySize * sizeof(FLOATTYPE_T) ) );

  void *d_temp_storage = NULL;
  size_t temp_storage_bytes = 0;
  cub::DeviceReduce::Sum( d_temp_storage, temp_storage_bytes, d_in, d_out, 
    size );

  printf("temp storage is %ld\n", temp_storage_bytes );

  checkCUDA( cudaMalloc( &d_temp_storage, temp_storage_bytes ) ); 

/* start the timers */

