 
 /* launch the kernel on the GPU */
 
-  sumReduction<<< 1, 1 >>>( size, d_in, d_out );
+  sumReduction<<< 1, 1 >>>( size, d_in, d_sum );
   checkKERNEL()
 
 /* stop the timers */

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

  checkCUDA( cudaMalloc( &d_in, memBytes ) );
  checkCUDA( cudaMalloc( &d_out, 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) ) );

/* calculate block and grid sizes */


#include "../debug.h"

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

#define FLOATTYPE_T double

__global__ void sumReduction(int n, FLOATTYPE_T *in, FLOATTYPE_T *sum)
{
/* calculate global index in the array */
  int globalIndex = blockIdx.x * blockDim.x + threadIdx.x;
