 
     CUDA_SAFE_CALL( cudaMemcpy(h_keys, device_storage.d_keys, sizeof(unsigned int) * num_elements, cudaMemcpyDeviceToHost) );
 	CUDA_SAFE_CALL( cudaFree(device_storage.d_keys) );		// clean up keys
+	CUDA_SAFE_CALL( cudaFree(device_storage.d_values) );		// clean up values
 
     // Clean up events
 	CUDA_SAFE_CALL( cudaEventDestroy(start_event) );

	// Allocate device storage and create sorting enactor  
	//

	RadixSortStorage<unsigned int> device_storage;	
	CUDA_SAFE_CALL( cudaMalloc((void**) &device_storage.d_keys, sizeof(unsigned int) * num_elements) );

	RadixSortingEnactor<unsigned int> sorting_enactor(num_elements);


	

		
		Base::template DigitPlacePass<0, 4, 0,  MandatoryPassNopFunctor<ConvertedKeyType>, MandatoryPassNopFunctor<ConvertedKeyType> >(converted_storage);		// d_keys (orig) -> d_alt_keys (orig, inner)
		Base::template DigitPlacePass<1, 4, 4,  MandatoryPassNopFunctor<ConvertedKeyType>, MandatoryPassNopFunctor<ConvertedKeyType> >(inner_pass_storage); 	// d_alt_keys (orig, inner) -> d_keys (inner)
		Base::template DigitPlacePass<2, 3, 8,  MandatoryPassNopFunctor<ConvertedKeyType>, MandatoryPassNopFunctor<ConvertedKeyType> >(inner_pass_storage); 	// d_keys (inner) -> d_alt_keys  (orig, inner)
		Base::template DigitPlacePass<3, 3, 11, MandatoryPassNopFunctor<ConvertedKeyType>, MandatoryPassNopFunctor<ConvertedKeyType> >(inner_pass_storage); 	// d_alt_keys (orig, inner) -> d_keys (inner)
		Base::template DigitPlacePass<4, 3, 14, MandatoryPassNopFunctor<ConvertedKeyType>, MandatoryPassNopFunctor<ConvertedKeyType> >(final_pass_storage); 	// d_keys (inner) -> d_keys (orig)

		return cudaSuccess;
	}
