 		int bounce,
 		cudaTextureObject_t texVtxPos)
 	{
-#if 0
-		dim3 blockPerGrid_HitTest((width * height + 128 - 1) / 128);
-		dim3 threadPerBlock_HitTest(128);
-
-		hitTest << <blockPerGrid_HitTest, threadPerBlock_HitTest >> > (
+#ifdef ENABLE_PERSISTENT_THREAD
+		hitTest << <NUM_BLOCK, dim3(WARP_SIZE, NUM_WARP_PER_BLOCK) >> > (
 #else
 		dim3 block(BLOCK_SIZE, BLOCK_SIZE);
 		dim3 grid(

	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 g_headDev = 0;

#define ENABLE_PERSISTENT_THREAD

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