 			depth = 0;
 
 			while (depth < maxDepth) {
+#ifdef ENABLE_PERSISTENT_THREAD
+				hitTest << <NUM_BLOCK, dim3(WARP_SIZE, NUM_WARP_PER_BLOCK) >> > (
+#else
 				hitTest << <blockPerGrid_HitTest, threadPerBlock_HitTest >> > (
+#endif
 				//hitTest << <1, 1 >> > (
 					paths.ptr(),
 					isects.ptr(),

		ctxt.matrices = matrices;
	}

	aten::Intersection isect;
	
	bool isHit = intersectBVH(&ctxt, rays[idx], &isect);

	isects[idx].t = isect.t;

	// 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,
