 				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,

	out[k] = r | (g << 8) | (b << 16);
}

enum {
	Z_THRM = 20 * 0x10000,
	Z_THRD = 10,
};

__global__ void denoiseKernelMotion(
		uint32_t **diffs, uint32_t *out, int h, int w) {
	int i = blockDim.x * blockIdx.x + threadIdx.x;
	int j = blockDim.y * blockIdx.y + threadIdx.y;
	int k = i * w + j;

	int i2, j2;
	int t;
	// Motion detection
	for (t = 0; t < NDIFFS && diffs[t]; t++) {
		if (diffs[t][k] < Z_THRM) {
			continue;
		}
		int excess = 0;
		FOR_WINDOW(i2, j2, i, j, h, w, 1) {
			int k2 = i2 * w + j2;
			if (k2 != k && diffs[t][k2] >= Z_THRM) {
				excess++;
				if (excess == 2) {
					goto after_md;
				}
			}
		}
	}
after_md:
	int v = 0xff * (NDIFFS - t) / NDIFFS;
	out[k] = v | (v << 8) | (v << 16);
}

__global__ void denoiseKernelAKNN(
		uint32_t **backlogs, uint32_t **diffs, int nbacklogs, uint32_t *out,
		int h, int w) {
	int i = blockDim.x * blockIdx.x + threadIdx.x;
	int j = blockDim.y * blockIdx.y + threadIdx.y;

	for (i2 = max(0, i - wd); min(h, i2 < i + wd); i2++) \
		for (j2 = max(0, j - wd); min(w, j2 < j + wd); j2++)

__global__ static void updateMotion(
		bool *m, uint32_t *d, int h, int w) {
	int i = blockDim.x * blockIdx.x + threadIdx.x;
	int j = blockDim.y * blockIdx.y + threadIdx.y;
	int k = i * w + j;

	if (d[k] < Z_THRM) {
		m[k] = false;
	}

	int i2, j2;
	int excess = 0;
	FOR_WINDOW(i2, j2, i, j, h, w, 1) {
		int k2 = i2 * w + j2;
		if (k2 != k && d[k2] >= Z_THRM) {
			excess++;
			if (excess == 2) {
				m[k] = true;
				return;
			}
		}
	}
}

__global__ void denoiseKernelSpatial(
		const uint32_t *in, uint32_t *out, int h, int w) {
	int i = blockDim.x * blockIdx.x + threadIdx.x;
