 	cudaUnbindTexture( ClassifyKeyholeTexture );
 	cudaFreeArray( PrimaryTextureArray );
 	cudaFreeArray( KeyholeTextureArray );
-	cudaStreamSynchronize(*stream);
 	cudaFree( dev_OutputData );
 
 }
	classification = (inClipping && inKeyhole) ? - tex2D(ClassifyKeyholeTexture, value.x, value.y) : classification;

	//output the final classification
	if( index.z < info.VolumeSize[2] ) outputVolume[inIndex] = classification;
}

void CUDAalgo_classifyVoxels( float* inputData, short* inputPrimaryTexture, short* inputKeyholeTexture, int textureSize,


	//allocate working memory for the output
	short* dev_OutputData = 0;
	cudaMalloc( (void**) &dev_OutputData, sizeof(short)*VolumeSize );
	
	//classify the volume
	dim3 grid( (VolumeSize-1) / NUM_THREADS+1,1,1);
	if( grid.x > MAX_GRID_SIZE ) grid.y = grid.x = (int) sqrt((double)VolumeSize/(double)NUM_THREADS) + 1;
	if( grid.y > MAX_GRID_SIZE ){ grid.y = grid.x = MAX_GRID_SIZE; grid.z = (VolumeSize/NUM_THREADS-1)/(grid.x*grid.y)+1; }
	dim3 threads(NUM_THREADS,1,1);
	ClassifyVolume<<< grid, threads, 0, *stream >>>((float2*)dev_InputData, dev_OutputData);
	
	cudaThreadSynchronize();
	printf( "Classify: " );
	printf( cudaGetErrorString( cudaGetLastError() ) );
	printf( "\n" );

	//retrieve classified output
	cudaMemcpyAsync( outputData, dev_OutputData, sizeof(short)*VolumeSize,
						cudaMemcpyDeviceToHost, *stream);
	cudaStreamSynchronize(*stream);

	cudaThreadSynchronize();
	printf( "Memcpy: " );
	printf( cudaGetErrorString( cudaGetLastError() ) );
	printf( "\n" );

	//deallocate textures and image memory
	cudaFree( dev_InputData );
