 			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) {


		switch (m) {
		case DIFF:
			denoiseKernelDiff<<<blocksPerGrid,
				threadsPerBlock>>>(ctx->dbacklogs, i, dout, w);
			break;
		case TEMPORAL_AVG:
			denoiseKernelTemporalAvg<<<blocksPerGrid,
				threadsPerBlock>>>(ctx->dbacklogs, dout, w);
			break;
		case KNN:
			denoiseKernelKNN<<<blocksPerGrid,
				threadsPerBlock>>>(ctx->dbacklogs, i, dout, h, w);
			break;
		case AKNN:
			denoiseKernelAKNN<<<blocksPerGrid,
				threadsPerBlock>>>(ctx->dbacklogs, ctx->dmotions, i, dout, h, w);
			break;
		case MOTION:
			denoiseKernelMotion<<<blocksPerGrid,
				threadsPerBlock>>>(ctx->dmotions, dout, h, w);
			break;
		case ADAPTIVE_TEMPORAL_AVG:
			denoiseKernelAdaptiveTemporalAvg<<<blocksPerGrid,
				threadsPerBlock>>>(ctx->dbacklogs, i, dout, w);
			break;
		}
		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);
		}

