   {
     sArray[threadIdx.x] += in[i];
   } /* end for */
-
   
   __syncthreads();
 
-#if 1
-  if( blockDim.x >= 512 ) 
-  {
-    if( threadIdx.x < 256 )
-    {
-      sArray[threadIdx.x] += sArray[threadIdx.x+256];
-    } /* end if */
-    __syncthreads();
-  } /* end if */
-  if( blockDim.x >= 256 ) 
-  {
-    if( threadIdx.x < 128 )
-    {
-      sArray[threadIdx.x] += sArray[threadIdx.x+128];
-    } /* end if */
-    __syncthreads();
-  } /* end if */
-  if( blockDim.x >= 128 ) 
-  {
-    if( threadIdx.x < 64 )
-    {
-      sArray[threadIdx.x] += sArray[threadIdx.x+64];
-    } /* end if */
-    __syncthreads();
-  } /* end if */
-  if( blockDim.x >= 64 ) 
-  {
-    if( threadIdx.x < 32 )
-    {
-      sArray[threadIdx.x] += sArray[threadIdx.x+32];
-    } /* end if */
-    __syncthreads();
-  } /* end if */
-  if( blockDim.x >= 32 ) 
-  {
-    if( threadIdx.x < 16 )
-    {
-      sArray[threadIdx.x] += sArray[threadIdx.x+16];
-    } /* end if */
-    __syncthreads();
-  } /* end if */
-  if( blockDim.x >= 16 ) 
-  {
-    if( threadIdx.x < 8 )
-    {
-      sArray[threadIdx.x] += sArray[threadIdx.x+8];
-    } /* end if */
-    __syncthreads();
-  } /* end if */
-  if( blockDim.x >= 8 ) 
-  {
-    if( threadIdx.x < 4 )
-    {
-      sArray[threadIdx.x] += sArray[threadIdx.x+4];
-    } /* end if */
-    __syncthreads();
-  } /* end if */
-  if( blockDim.x >= 4 ) 
-  {
-    if( threadIdx.x < 2 )
-    {
-      sArray[threadIdx.x] += sArray[threadIdx.x+2];
-    } /* end if */
-    __syncthreads();
-  } /* end if */
-  if( blockDim.x >= 2 ) 
-  {
-    if( threadIdx.x < 1 )
-    {
-      sArray[threadIdx.x] += sArray[threadIdx.x+1];
-    } /* end if */
-    __syncthreads();
-  } /* end if */
-
-#else
-//  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]);
+  if( threadIdx.x == 0 ) out[blockIdx.x] = sArray[0]; 
 
-#endif
-  if( threadIdx.x == 0 ) out[blockIdx.x] = sArray[0]; 
   return;
 
 }
 
 int main()
 {
-  floatType_t *h_in, h_out, good_out;
-  floatType_t *d_in, *d_out, *d_tempArray;
+  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 memBytes = size * sizeof( FLOATTYPE_T );
   int tempArraySize = 32768;
 
 /* get GPU device number and name */

#include <stdio.h>
#include "../debug.h"

#define N ( 1 << 26 )
#define THREADS_PER_BLOCK 256

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;

/* allocate space for device copies of in, out */

  checkCUDA( cudaMalloc( &d_in, memBytes ) );
  checkCUDA( cudaMalloc( &d_sum, 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 );
    if( i % 2 == 0 ) h_in[i] = -h_in[i];
  }

  h_sum      = 0.0;
  cpu_sum   = 0.0;

/* copy inputs to device */

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

/* calculate block and grid sizes */

