 	// of values seen). n is number of values corresponding
 	// to each M1 and M2 value.
 	mean_stddev_reduction_kernel1<<<grid,block,0,sa.getStream(stream)>>>(dPssIn, d_M1, d_M2, d_n);
+	SAFE_CALL(cudaStreamSynchronize(sa.getStream(stream)),"ZCA cudaStreamSynchronize failed");
 
 	// Second reduction generates mean and stddev values
 	// 12x12 fit in a single block so there's no

	// image has the mean and stddev computed individually
	// in the CUDA kernels above
	global_contrast_normalization_kernel<<<grid,block,0,sa.getStream(stream)>>>(dPssIn, dMean, dStddev, dFlattenedImages);

	// Multiply images by weights to get the ZCA-whitened output
	gemm(dFlattenedImages, weights, 1.0, buf, 0.0, zcaOut, 0, stream);

	// Copy to output buffer in the order expected by
	// neural net input

}
#define SAFE_CALL(call,msg) _safe_cuda_call((call),(msg),__FILE__,__LINE__)

#ifndef cublasSafeCall
#define cublasSafeCall(err) __cublasSafeCall(err, __FILE__, __LINE__)
#endif
inline void __cublasSafeCall(cublasStatus_t err, const char *file, const int line)
{
if( CUBLAS_STATUS_SUCCESS != err) {
fprintf(stderr, "CUBLAS error in file '%s', line %d\n \nerror %d \nterminating!\n",__FILE__, __LINE__,err);
cudaDeviceReset(); assert(0);
}
}
// Apply global contrast normalization to
// each input image.
// For each channel in each image, the mean and stddev has
