 	}
 }
 
+enum {
+	D0_THRESHOLD = 100,
+	DSUM_THRESHOLD = 250,
+};
+
 __global__ void denoiseKernelAdaptiveTemporalAvg(
-		KernelContext *ctx, uint32_t *out,
+		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;

				threadsPerBlock>>>(dctx, dout, h, w, perThread);
			break;
		case ADAPTIVE_TEMPORAL_AVG:
			denoiseKernelAdaptiveTemporalAvg<<<blocksPerGrid,
				threadsPerBlock>>>(dctx, dout, h, w, perThread);
			break;
		}
		cudaFree(dctx);

	}
}

__global__ void denoiseKernelTemporalAvg(
		KernelContext *ctx, uint32_t *out,
		int h, int w, int perThread) {
