 			}
 		}
 
-		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;
 	}
 


void *kernelInit() {
	KernelContext *ctx = (KernelContext*) calloc(1, sizeof(KernelContext));
	return ctx;
}


		switch (m) {
		case SPATIAL:
			denoiseKernelSpatial<<<blocksPerGrid, threadsPerBlock>>>(
					din, dout, h, w);
			break;
		case SOBEL:
			double *dsobel = 0;
			denoiseKernelSobel<<<blocksPerGrid, threadsPerBlock>>>(
					din, frame, dsobel, dout, h, w);
			break;
		}
		cudaFree(din);
