 /* 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) ) );
+  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 );
+  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_in[i] = FLOATTYPE_T( rand() ) / ( FLOATTYPE_T (RAND_MAX) + 1.0 );
+    if( i % 2 == 0 ) h_in[i] = -h_in[i];
   }
 
-  h_out      = 0.0;
-  good_out   = 0.0;
+  h_sum      = 0.0;
+  cpu_sum   = 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_sum, 0, sizeof(FLOATTYPE_T) ) );
   checkCUDA( cudaMemset( d_tempArray, 0, 
-    tempArraySize * sizeof(floatType_t) ) );
+    tempArraySize * sizeof(FLOATTYPE_T) ) );
 
 /* calculate block and grid sizes */
 

  float elapsedTime;
  checkCUDA( cudaEventElapsedTime( &elapsedTime, start, stop ) );

  printf("Total elements is %d, %f GB\n", size, sizeof(floatType_t)*
    (double)size * 1.e-9 );
  printf("GPU total time is %f ms, bandwidth %f GB/s\n", elapsedTime,
    sizeof(floatType_t) * (double) size /
    ( (double) elapsedTime / 1000.0 ) * 1.e-9);

/* copy result back to host */

  checkCUDA( cudaMemcpy( &h_out, d_out, sizeof(floatType_t), 
    cudaMemcpyDeviceToHost ) );

  checkCUDA( cudaEventRecord( start, 0 ) );

  for( int i = 0; i < size; i++ )
  {
    good_out += h_in[i];
  } /* end for */

  checkCUDA( cudaEventRecord( stop, 0 ) );
  checkCUDA( cudaEventSynchronize( stop ) );
  checkCUDA( cudaEventElapsedTime( &elapsedTime, start, stop ) );
  printf("CPU total time is %f ms, bandwidth %f GB/s\n", elapsedTime,
    sizeof(floatType_t) * (double) size /
    ( (double) elapsedTime / 1000.0 ) * 1.e-9);


  floatType_t diff = abs( good_out - h_out );

  if( diff / h_out < 0.001 ) printf("PASS\n");
  else
  {                       
    printf("FAIL\n");
    printf("Error is %f\n", diff / h_out );
    printf("GPU result is %f, CPU result is %f\n",h_out, good_out );
  } /* end else */

/* clean up */

  free(h_in);
  checkCUDA( cudaFree( d_in ) );
  checkCUDA( cudaFree( d_out ) );
  checkCUDA( cudaFree( d_tempArray ) );

  checkCUDA( cudaDeviceReset() );


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

/* stop the timers */
