 			}
 		}
 
-		KernelContext *dctx;
-		cudaMalloc(&dctx, sizeof(KernelContext));
-		cudaMemcpy(dctx, ctx, sizeof(KernelContext), cudaMemcpyHostToDevice);
+		cudaMemcpy(ctx->dbacklogs, ctx->backlogs, sizeof(ctx->backlogs),
+				cudaMemcpyHostToDevice);
+		cudaMemcpy(ctx->ddiffs, ctx->diffs, sizeof(ctx->diffs),
+				cudaMemcpyHostToDevice);
+
 		int i;
 		for (i = 0; i < BACKLOG_SIZE && ctx->backlogs[i]; i++)
 			;
+
 		switch (m) {
 		case DIFF:
 			denoiseKernelDiff<<<blocksPerGrid,
-				threadsPerBlock>>>(dctx, i, dout, h, w, perThread);
+				threadsPerBlock>>>(ctx->dbacklogs, i, dout, w);
 			break;
 		case TEMPORAL_AVG:
 			denoiseKernelTemporalAvg<<<blocksPerGrid,
-				threadsPerBlock>>>(dctx, dout, h, w, perThread);
+				threadsPerBlock>>>(ctx->dbacklogs, dout, w);
 			break;
 		case KNN:
 			denoiseKernelKNN<<<blocksPerGrid,
-				threadsPerBlock>>>(dctx, i, dout, h, w, perThread);
+				threadsPerBlock>>>(ctx->dbacklogs, i, dout, h, w);
 			break;
 		case ZLOKOLICA:
-			updateDiff<<<blocksPerGrid,
-				threadsPerBlock>>>(dctx, h, w, perThread);
+			updateDiff<<<blocksPerGrid, threadsPerBlock>>>(
+					ctx->diffs[0], ctx->backlogs[0], ctx->backlogs[1], w);
 			denoiseKernelZlokolica<<<blocksPerGrid,
-				threadsPerBlock>>>(dctx, i, dout, h, w, perThread);
+				threadsPerBlock>>>(ctx->dbacklogs, ctx->ddiffs, i, dout, h, w);
 			break;
 		case MOTION:
-			updateDiff<<<blocksPerGrid,
-				threadsPerBlock>>>(dctx, h, w, perThread);
+			updateDiff<<<blocksPerGrid, threadsPerBlock>>>(
+					ctx->diffs[0], ctx->backlogs[0], ctx->backlogs[1], w);
 			denoiseKernelMotion<<<blocksPerGrid,
-				threadsPerBlock>>>(dctx, i, dout, h, w, perThread);
+				threadsPerBlock>>>(ctx->ddiffs, dout, h, w);
 			break;
 		case ADAPTIVE_TEMPORAL_AVG:
 			denoiseKernelAdaptiveTemporalAvg<<<blocksPerGrid,
-				threadsPerBlock>>>(dctx, i, dout, h, w, perThread);
-			break;
-		}
-		cudaFree(dctx);
+				threadsPerBlock>>>(ctx->dbacklogs, i, dout, w);
+			break;
+		}
 		break;
 	}
 

};

__global__ void denoiseKernelAdaptiveTemporalAvg(
		KernelContext *ctx, int nbacklogs, uint32_t *out,
		int h, int w, int perThread) {
	int k = blockDim.x * blockIdx.x + threadIdx.x;
	int lower = k * perThread;
	int upper = (k + 1) * perThread;
	if (k == 256 * 32 - 1) {
		upper = h * w;
	}
	uint32_t **backlogs = ctx->backlogs;
	for (k = lower; k < upper; k++) {
		int c = nbacklogs / 2, i, j, r = 0, g = 0, b = 0;
		int dsum = 0;
		uint32_t vc = backlogs[c][k];
		/*
		for (i = c; i >= 0; i--) {
			uint32_t v = backlogs[i][k];
			dsum += distance(v, lastv);
			if (distance(vc, v) > D0_THRESHOLD || dsum > DSUM_THRESHOLD) {
				break;
			}
			r += v & 0xff;
			g += (v >> 8) & 0xff;
			b += (v >> 16) & 0xff;
			lastv = v;
		}
		j = i;
		lastv = vc;
		for (i = c + 1; i < BACKLOG_SIZE && backlogs[i]; i++) {
			uint32_t v = backlogs[i][k];
			dsum += distance(v, lastv);
			if (distance(vc, v) > D0_THRESHOLD || dsum > DSUM_THRESHOLD) {
				break;
			}
			r += v & 0xff;
			g += (v >> 8) & 0xff;
			b += (v >> 16) & 0xff;
			lastv = v;
		}
		*/
		uint32_t vi_last = vc, vj_last = vc;
		for (i = j = c; i < nbacklogs && j >= 0; i++, j--) {
			uint32_t vi = backlogs[i][k], vj = backlogs[j][k];
			dsum += distance(vi, vi_last) + distance(vj, vj_last);
			if (distance(vc, vi) > D0_THRESHOLD ||
				distance(vc, vj) > D0_THRESHOLD ||
				dsum > DSUM_THRESHOLD) {
				break;
			}
			r += vi & 0xff;
			g += (vi >> 8) & 0xff;
			b += (vi >> 16) & 0xff;
			r += vj & 0xff;
			g += (vj >> 8) & 0xff;
			b += (vj >> 16) & 0xff;
			vi_last = vi;
			vj_last = vj;
		}
		r /= i - j;
		g /= i - j;
		b /= i - j;
		r &= 0xff;
		g &= 0xff;
		b &= 0xff;

		//r = g = b = 0xff * (i - j) / nbacklogs;
		out[k] = r | (g << 8) | (b << 16);
	}
}

enum {

	for (int i = 0; i < NDIFFS; i++) {
		cudaFree(ctx->diffs[i]);
	}
	cudaFree(ctx->dbacklogs);
	cudaFree(ctx->ddiffs);
	free(ctx);
}
