 	// Thread index within block - used for addressing smem below
 	const unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x;
 
+	// Shared memory per channel per thread = 1 long, 2 floats.
+	// So a 3 channel image needs 3 longs and 6 floats
+	// Thread blocks are 16x16 threads
 	// TODO : fixme for variable sized thread blocks
 	__shared__ float M1[16*16*3];
 	__shared__ float M2[16*16*3];

	}
}

__device__ void combine_running_totals(float &M1_1, float M1_2, float &M2_1, float M2_2, unsigned int &n_1, unsigned int n_2)
{
	unsigned int combined_n = n_1 + n_2;


	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");
}

