 #include <stdio.h>
 #include "../debug.h"
 
-#define N ( 1 << 26 )
-#define THREADS_PER_BLOCK 256
+#define N ( 1 << 27 )
+#define THREADS_PER_BLOCK 128
 
-typedef double floatType_t;
+#define FLOATTYPE_T float
 
-__global__ void sumReduction(int n, floatType_t *in, floatType_t *out)
+__global__ void sumReduction(int n, FLOATTYPE_T *in, FLOATTYPE_T *out)
 {
 
-  __shared__ floatType_t sArray[THREADS_PER_BLOCK];
+  __shared__ FLOATTYPE_T sArray[THREADS_PER_BLOCK];
 
 /* calculate global index in the array */
   int globalIndex = blockIdx.x * blockDim.x + threadIdx.x;

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

  h_in = (floatType_t *)malloc( memBytes );

  for( int i = 0; i < size; i++ )
  {
   h_in[i] = floatType_t( rand() ) / ( floatType_t (RAND_MAX) + 1.0 );
  }

  h_out      = 0.0;
  good_out   = 0.0;

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

/* calculate block and grid sizes */



  sumReduction<<< blocks, threads1 >>>( size, d_in,  d_tempArray );
  checkKERNEL()
  sumReduction<<<      1, threads2 >>>( blocks.x, d_tempArray, d_sum );
  checkKERNEL()

/* stop the timers */
