 void denoise(void *p, int m, const uint32_t *in, uint32_t *out, int h, int w) {
 	static int frame = 0;
 	KernelContext *ctx = (KernelContext*) p;
-	dim3 threadsPerBlock(8, 8);
-	dim3 blocksPerGrid(h / 8, w / 8);
+	dim3 blockDim(8, 8);
+	dim3 gridDim(h / blockDim.x, w / blockDim.y);
 
 	int size = h * w * sizeof(uint32_t);
 	uint32_t *dout;

			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->m0,
						size / sizeof(uint32_t) * sizeof(bool));
			}
			updateMotion<<<gridDim, blockDim>>>(
					ctx->m0, ctx->diffs[0], h, w);
			dilateMotion<<<gridDim, blockDim>>>(
					ctx->motions[0], ctx->m0, h, w);
		}

