 		switch (m) {
 		case AKNN:
 		case MOTION:
-			for (int i = NDIFFS; i > 0; i--) {
-				ctx->diffs[i] = ctx->diffs[i-1];
-			}
-			ctx->diffs[0] = ctx->diffs[NDIFFS];
+			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++)

	KNN_FRAMES = 5,
};

typedef struct {
	uint32_t v;
	double d;

		g = ((u >> 8) & 0xff) - ((v >> 8) & 0xff),
		b = ((u >> 16) & 0xff) - ((v >> 16) & 0xff);
	return sqrt(double(r * r + g * g + b * b));
}

__device__ uint32_t lum(uint32_t u) {
