 #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;


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

/* stop the timers */

  {
    sArray[threadIdx.x] += in[i];
  } /* end for */
  
  __syncthreads();

/* do the final reduction in SMEM */
  for( int i = blockDim.x/2; i > 0; i = i / 2 )
  {
    if( threadIdx.x < i )
    {
      sArray[threadIdx.x] += sArray[threadIdx.x + i];
    } /* end if */
    __syncthreads();
  } /* end for */

  if( threadIdx.x == 0 ) out[blockIdx.x] = sArray[0]; 

  return;

}

int main()
{
  FLOATTYPE_T *h_in, h_sum, cpu_sum;
  FLOATTYPE_T *d_in, *d_sum, *d_tempArray;
  int size = N;
  int memBytes = size * sizeof( FLOATTYPE_T );
  int tempArraySize = 32768;

/* get GPU device number and name */
