 		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>

	int bounce,
	float hitDistLimit)
{
#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;
