 	cudaTextureObject_t vtxPos,
 	aten::mat4* matrices)
 {
-	int idx = blockIdx.x * blockDim.x + threadIdx.x;
-
-	if (idx >= width * height) {
-		return;
-	}
-
-	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) {
+		headDev = 0;
 	}
 
 	Context ctxt;

			depth = 0;

			while (depth < maxDepth) {
				hitTest << <blockPerGrid_HitTest, threadPerBlock_HitTest >> > (
				//hitTest << <1, 1 >> > (
					paths.ptr(),
					isects.ptr(),

	path.isHit = isHit;

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

template <bool isFirstBounce>
