 	int bounce,
 	float hitDistLimit)
 {
-#if 0
-	int idx = blockIdx.x * blockDim.x + threadIdx.x;
-
-	if (idx >= width * height) {
-		return;
-	}
-#else
-	const auto ix = blockIdx.x * blockDim.x + threadIdx.x;
-	const auto iy = blockIdx.y * blockDim.y + threadIdx.y;
-
-	if (ix >= width && iy >= height) {
-		return;
-	}
-
-	const auto idx = getIdx(ix, iy, width);
-#endif
-
-	auto& path = paths[idx];
-	path.isHit = false;
-
-	hitbools[idx] = 0;
-
-	if (path.isTerminate) {
-		return;
+#ifdef ENABLE_PERSISTENT_THREAD
+	// warp-wise head index of tasks in a block
+	__shared__ volatile unsigned int headBlock[NUM_WARP_PER_BLOCK];
+
+	volatile unsigned int& headWarp = headBlock[threadIdx.y];
+
+	if (blockIdx.x == 0 && threadIdx.x == 0) {
+		g_headDev = 0;
 	}
 
 	Context ctxt;

	// Accumulate value, so do not reset.
	//path.contrib = aten::vec3(0);
}

__global__ void hitTest(
	idaten::SVGFPathTracing::Path* paths,

	path.isHit = isHit;

	hitbools[idx] = isHit ? 1 : 0;
#endif
}

template <bool isFirstBounce>
