 				threadsPerBlock>>>(ctx->dbacklogs, i, dout, h, w);
 			break;
 		case AKNN:
-			updateDiff<<<blocksPerGrid, threadsPerBlock>>>(
-					ctx->diffs[0], ctx->backlogs[0], ctx->backlogs[1], w);
 			denoiseKernelAKNN<<<blocksPerGrid,
-				threadsPerBlock>>>(ctx->dbacklogs, ctx->ddiffs, i, dout, h, w);
+				threadsPerBlock>>>(ctx->dbacklogs, ctx->dmotions, i, dout, h, w);
 			break;
 		case MOTION:
-			updateDiff<<<blocksPerGrid, threadsPerBlock>>>(
-					ctx->diffs[0], ctx->backlogs[0], ctx->backlogs[1], w);
 			denoiseKernelMotion<<<blocksPerGrid,
-				threadsPerBlock>>>(ctx->ddiffs, dout, h, w);
+				threadsPerBlock>>>(ctx->dmotions, dout, h, w);
 			break;
 		case ADAPTIVE_TEMPORAL_AVG:
 			denoiseKernelAdaptiveTemporalAvg<<<blocksPerGrid,

	KNN_FRAMES = 5,
};

typedef struct {
	uint32_t v;
	double d;

		switch (m) {
		case AKNN:
		case MOTION:
			ror(ctx->diffs, NDIFFS);
			if (ctx->diffs[0] == 0) {
				cudaMalloc(&ctx->diffs[0], size);
			}
			updateDiff<<<blocksPerGrid, threadsPerBlock>>>(
					ctx->diffs[0], ctx->backlogs[0], ctx->backlogs[1], w);
			ror(ctx->motions, NDIFFS);
			if (ctx->motions[0] == 0) {
				cudaMalloc(&ctx->motions[0],
						size / sizeof(uint32_t) * sizeof(bool));
			}
			updateMotion<<<blocksPerGrid, threadsPerBlock>>>(
					ctx->motions[0], ctx->diffs[0], h, w);
		}

		cudaMemcpy(ctx->dbacklogs, ctx->backlogs, sizeof(ctx->backlogs),
				cudaMemcpyHostToDevice);
		cudaMemcpy(ctx->ddiffs, ctx->diffs, sizeof(ctx->diffs),
				cudaMemcpyHostToDevice);
		cudaMemcpy(ctx->dmotions, ctx->motions, sizeof(ctx->motions),
				cudaMemcpyHostToDevice);

		int i;
		for (i = 0; i < BACKLOG_SIZE && ctx->backlogs[i]; i++)
