 		for (j2 = max(0, j - wd); min(w, j2 < j + wd); j2++)
 
 __global__ void denoiseKernelSpatial(
-		const uint32_t *in, 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;
-	}
-	for (k = lower; k < upper; k++) {
-		int i = k / w, j = k % w;
-		double dmin = INFINITY;
-		int i2, j2, i3, j3;
-		FOR_WINDOW(i2, j2, i, j, h, w, HALF_WINDOW) {
-			double d = 0;
-			FOR_WINDOW(i3, j3, i, j, h, w, HALF_WINDOW) {
-				d += distance(in(i2, j2), in(i3, j3));
-			}
-			if (dmin > d) {
-				dmin = d;
-				out(i, j) = in(i2, j2);
-			}
+		const uint32_t *in, 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;
+
+	double dmin = INFINITY;
+	int i2, j2, i3, j3;
+	FOR_WINDOW(i2, j2, i, j, h, w, HALF_WINDOW) {
+		double d = 0;
+		FOR_WINDOW(i3, j3, i, j, h, w, HALF_WINDOW) {
+			d += distance(in(i2, j2), in(i3, j3));
+		}
+		if (dmin > d) {
+			dmin = d;
+			out[k] = in(i2, j2);
 		}
 	}
 }
 
 __global__ void denoiseKernelSobel(
 		const uint32_t *in, int frame, double *sout, 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;
-	}
-	for (k = lower; k < upper; k++) {
-		int x = k / w, y = k % w;
-		if (x == 0 || x == h-1 || y == 0 || y == w-1) {
-			out[k] = 0;
-			if (sout != 0) {
-				sout[k] = 0;
-			}
-			continue;
-		}
-		int gx = lum(in(x-1, y+1)) + 2 * lum(in(x, y+1)) + lum(in(x+1, y+1))
-			   - lum(in(x-1, y-1)) - 2 * lum(in(x, y-1)) - lum(in(x+1, y-1));
-		int gy = lum(in(x-1, y-1)) + 2 * lum(in(x-1, y)) + lum(in(x-1, y+1))
-			   - lum(in(x+1, y-1)) - 2 * lum(in(x+1, y)) - lum(in(x+1, y+1));
-		/*
-		int g = sqrt((gx * gx + gy * gy) / 50.0);
-		*/
-		int g = ((gx * gx + gy * gy) > 2000) * 0xff;
-		/*
-		if (g > 0x5) {
-			g = 0xff;
-		} else {
-			g = 0;
-		}
-		*/
-		out[k] = g | (g << 8) | (g << 16);
+		int w) {
+	int x = blockDim.x * blockIdx.x + threadIdx.x;
+	int y = blockDim.y * blockIdx.y + threadIdx.y;
+	int k = x * w + y;
+
+	if (x == 0 || x == h-1 || y == 0 || y == w-1) {
+		out[k] = 0;
 		if (sout != 0) {
-			sout[k] = sqrt(double(gx * gx + gy * gy));
-		}
+			sout[k] = 0;
+		}
+		return;
+	}
+	int gx = lum(in(x-1, y+1)) + 2 * lum(in(x, y+1)) + lum(in(x+1, y+1))
+		   - lum(in(x-1, y-1)) - 2 * lum(in(x, y-1)) - lum(in(x+1, y-1));
+	int gy = lum(in(x-1, y-1)) + 2 * lum(in(x-1, y)) + lum(in(x-1, y+1))
+		   - lum(in(x+1, y-1)) - 2 * lum(in(x+1, y)) - lum(in(x+1, y+1));
+	/*
+	int g = sqrt((gx * gx + gy * gy) / 50.0);
+	*/
+	int g = ((gx * gx + gy * gy) > 2000) * 0xff;
+	/*
+	if (g > 0x5) {
+		g = 0xff;
+	} else {
+		g = 0;
+	}
+	*/
+	out[k] = g | (g << 8) | (g << 16);
+	if (sout != 0) {
+		sout[k] = sqrt(double(gx * gx + gy * gy));
 	}
 }
 
 __global__ void denoiseKernelTemporalAvg(
-		KernelContext *ctx, 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 i, r = 0, g = 0, b = 0;
-		for (i = 0; i < BACKLOG_SIZE && backlogs[i]; i++) {
-			uint32_t v = backlogs[i][k];
-			r += v & 0xff;
-			g += (v >> 8) & 0xff;
-			b += (v >> 16) & 0xff;
-		}
-		r /= i;
-		g /= i;
-		b /= i;
-		r &= 0xff;
-		g &= 0xff;
-		b &= 0xff;
-		out[k] = r | (g << 8) | (b << 16);
-	}
+		uint32_t **backlogs, uint32_t *out, int w) {
+	int k = (blockDim.x * blockIdx.x + threadIdx.x) * w +
+		blockDim.y * blockIdx.y + threadIdx.y;
+
+	int i, r = 0, g = 0, b = 0;
+	for (i = 0; i < BACKLOG_SIZE && backlogs[i]; i++) {
+		uint32_t v = backlogs[i][k];
+		r += v & 0xff;
+		g += (v >> 8) & 0xff;
+		b += (v >> 16) & 0xff;
+	}
+	r /= i;
+	g /= i;
+	b /= i;
+	r &= 0xff;
+	g &= 0xff;
+	b &= 0xff;
+	out[k] = r | (g << 8) | (b << 16);
 }
 
 enum {

}

__global__ static void updateDiff(
		KernelContext *ctx, 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 *d0 = ctx->diffs[0];
	uint32_t *b0 = ctx->backlogs[0];
	uint32_t *b1 = ctx->backlogs[1];
	if (b0 == 0 || b1 == 0) {
		return;
	}
	for (k = lower; k < upper; k++) {
		d0[k] = diff(b0[k], b1[k]);
	}
}

#define in(i, j) in[(i)*w + (j)]

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