 	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;

	path.isHit = isHit;

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

template <bool isFirstBounce>

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

// NOTE
// persistent thread.
// https://gist.github.com/guozhou/b972bb42bbc5cba1f062#file-persistent-cpp-L15

// NOTE
// compute capability 6.0
// http://homepages.math.uic.edu/~jan/mcs572/performance_considerations.pdf
// p3

#define NUM_SM				64	// no. of streaming multiprocessors
#define NUM_WARP_PER_SM		64	// maximum no. of resident warps per SM
#define NUM_BLOCK_PER_SM	32	// maximum no. of resident blocks per SM
#define NUM_BLOCK			(NUM_SM * NUM_BLOCK_PER_SM)
#define NUM_WARP_PER_BLOCK	(NUM_WARP_PER_SM / NUM_BLOCK_PER_SM)
#define WARP_SIZE			32

__device__ unsigned int headDev = 0;

#define ENABLE_PERSISTENT_THREAD

__global__ void hitTest(
	Path* paths,
