 		cudaMemcpy(din, in, size, cudaMemcpyHostToDevice);
 		switch (m) {
 		case SPATIAL:
-			denoiseKernelSpatial<<<blocksPerGrid, threadsPerBlock>>>(
+			denoiseKernelSpatial<<<gridDim, blockDim>>>(
 					din, dout, h, w);
 			break;
 		case SPATIAL_KNN:
-			denoiseKernelSpatialKNN<<<blocksPerGrid, threadsPerBlock>>>(
+			denoiseKernelSpatialKNN<<<gridDim, blockDim>>>(
 					din, dout, h, w);
 			break;
 		case SOBEL:
 			double *dsobel = 0;
-			denoiseKernelSobel<<<blocksPerGrid, threadsPerBlock>>>(
+			denoiseKernelSobel<<<gridDim, blockDim>>>(
 					din, frame, dsobel, dout, h, w);
 			break;
 		}

				cudaMalloc(&ctx->m0,
						size / sizeof(uint32_t) * sizeof(bool));
			}
			updateMotion<<<blocksPerGrid, threadsPerBlock>>>(
					ctx->m0, ctx->diffs[0], h, w);
			dilateMotion<<<blocksPerGrid, threadsPerBlock>>>(
					ctx->motions[0], ctx->m0, h, w);
		}


void denoise(void *p, int m, const uint32_t *in, uint32_t *out, int h, int w) {
	static int frame = 0;
	KernelContext *ctx = (KernelContext*) p;
	dim3 blockDim(8, 8);
	dim3 gridDim(h / blockDim.x, w / blockDim.y);

	int size = h * w * sizeof(uint32_t);
	uint32_t *dout;
