 }
 
 
+// Take the output of the ZCA matrix mul - that will
+// be a matrix. Each image is a row, each row is the pixels
+// in BGRBGRBGR.. order
+// Convert that to a flat 1-D array as expected by the neural
+// net input stages
 __global__ void unflatten_kernel(const cv::gpu::PtrStepSz<float> input,
 								 const size_t rows,
 								 const size_t cols,

	// each GPU mat in input. Copy it to device memory
	cv::gpu::PtrStepSz<float> hPssIn[input.size()];
	for (size_t i = 0; i < input.size(); ++i)
	{
		hPssIn[i] = input[i];
		//cout << "hPssIn[i] r=" << hPssIn[i].rows << " cols=" << hPssIn[i].cols << " step=" << hPssIn[i].step << endl;
	}
	cudaMemcpy(dPssIn, hPssIn, input.size() * sizeof(cv::gpu::PtrStepSz<float>), cudaMemcpyHostToDevice);

	// Specify a reasonable block size


	gemm(dFlattenedImages, weights, 1.0, buf, 0.0, zcaOut);

	unflatten_kernel<<<grid,block>>>(zcaOut, input[0].rows, input[0].cols, output);
	SAFE_CALL(cudaDeviceSynchronize(),"GCN kernel launch failed");

	SAFE_CALL(cudaFree(d_M1),"CUDA Free Failed");
	SAFE_CALL(cudaFree(d_M2),"CUDA Free Failed");
	SAFE_CALL(cudaFree(d_n),"CUDA Free Failed");
}

