 			if (ctx->diffs[0] == 0) {
 				cudaMalloc(&ctx->diffs[0], size);
 			}
-			updateDiff<<<blocksPerGrid, threadsPerBlock>>>(
+			updateDiff<<<gridDim, blockDim>>>(
 					ctx->diffs[0], ctx->backlogs[0], ctx->backlogs[1], w);
 			ror(ctx->motions, NDIFFS);
 			if (ctx->motions[0] == 0) {

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

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

