 	}
 }
 
+// population normalization of complex cell responses
+// note the 0.1, probably to avoid division by zero
 __global__ void dev_normalize(float *resp, float *pop, int nrXnrY, int nrZ) {
 	const int     tid = IMUL(blockDim.x, blockIdx.x) + threadIdx.x;
 	const int threadN = IMUL(blockDim.x, gridDim.x);

	dev_scaleHalfRect<<<iDivUp(nrX*nrY,128), 128>>>(d_color_tmp, 50.0, nrX*nrY);
	cutilSafeCall(cudaMemcpy(blue_yellow,d_color_tmp,sizeof(float)*nrX*nrY,GPUpointers?cudaMemcpyDeviceToDevice:cudaMemcpyDeviceToHost));

	// shift d_stimBuf in time by 1 frame
	for(int i=1;i<nrT;i++)
		cutilSafeCall(cudaMemcpy(&d_stimBuf[nrX*nrY*(i-1)],&d_stimBuf[nrX*nrY*i],sizeof(float)*nrX*nrY,cudaMemcpyDeviceToDevice));

	cutilSafeCall(cudaMemset (d_resp, 0, sizeof(float)*nrX*nrY*nrDirs*nrScales));

	//copy d_stimBuf to d_scalingStimBuf
	cutilSafeCall(cudaMemcpy(d_scalingStimBuf,d_stimBuf,sizeof(float)*nrX*nrY*nrT,cudaMemcpyDeviceToDevice));
	for (int scale=1; scale<=nrScales; scale++) {
		if (scale > 1) {
			float* tmp;
			cutilSafeCall(cudaMalloc((void**)&tmp, sizeof(float)*nrX*nrY*nrT));

float* d_diffV1GausBuf;
float* diffV1GausBufT;

unsigned char* d_stim; // the video input
float* d_pop;

// following 3 lines are needed for dev_split(), altho we only care about d_stim
float* d_red;
float* d_green;
float* d_blue;
