 	for(; i > NUMTHREADS; i = i/2){
 		dim3 tempGrid( i>NUMTHREADS ? i/NUMTHREADS : 1, 1, 1);
 		ReduceBySummation<<<tempGrid, threads, 0, *stream>>>(devDataTermBuffer, i, size );
+
+		#ifdef DEBUG_VTKCUDAHMFD
+			cudaThreadSynchronize();
+			printf( "Reduce: " );
+			printf( cudaGetErrorString( cudaGetLastError() ) );
+			printf( "\n" );
+		#endif
+
 	}
 	reduceData( min(NUMTHREADS,size), min(NUMTHREADS,size), 1, devDataTermBuffer, stream );
 
 	//return result to CPU
-	float retVal = 0.0f;
 	cudaMemcpyAsync( &retVal, devDataTermBuffer,  sizeof(float), cudaMemcpyDeviceToHost, *stream );
 	cudaStreamSynchronize(*stream);
 

	cudaMalloc( &devLabelBuffer, sizeof(float)*size );
	cudaMemcpyAsync( devDataTermBuffer, data,  sizeof(float)*size, cudaMemcpyHostToDevice, *stream );
	cudaMemcpyAsync( devLabelBuffer,    label, sizeof(float)*size, cudaMemcpyHostToDevice, *stream );

	//multiply buffers
	dim3 threads(NUMTHREADS,1,1);

#include "cuda.h"

#define NUMTHREADS 512
//#define DEBUG_VTKCUDAHMFD

__global__ void kern_GHMFD_MultiplyBuffers(float* storage, float* multiplicand, int size){
	int idx = threadIdx.x + blockIdx.x * blockDim.x;
