 		ctxt.matrices = matrices;
 	}
 
-	aten::Intersection isect;
-
-	bool isHit = intersectClosest(&ctxt, rays[idx], &isect);
-
-	//isects[idx].t = isect.t;
-	isects[idx].objid = isect.objid;
-	isects[idx].mtrlid = isect.mtrlid;
-	isects[idx].meshid = isect.meshid;
-	isects[idx].primid = isect.primid;
-	isects[idx].a = isect.a;
-	isects[idx].b = isect.b;
-
-	path.isHit = isHit;
-
-	hitbools[idx] = isHit ? 1 : 0;
+	do
+	{
+		// let lane 0 fetch [wh, wh + WARP_SIZE - 1] for a warp
+		if (threadIdx.x == 0) {
+			headWarp = atomicAdd(&_headDev, WARP_SIZE);
+		}
+
+		// task index per thread in a warp
+		unsigned int idx = headWarp + threadIdx.x;
+
+		if (idx >= testNum) {
+			return;
+		}
+
+		idx = notIntersectInScreenSpaceIds[idx];
+
+		int ix = idx % width;
+		int iy = idx / width;
+
+		idx = getIdx(ix, iy, width);
+
+		auto& path = paths[idx];
+		path.isHit = false;
+
+		hitbools[idx] = 0;
+
+		if (path.isTerminate) {
+			continue;
+		}
+
+		aten::Intersection isect;
+
+		bool isHit = intersectClosest(&ctxt, rays[idx], &isect);
+
+		//isects[idx].t = isect.t;
+		isects[idx].objid = isect.objid;
+		isects[idx].mtrlid = isect.mtrlid;
+		isects[idx].meshid = isect.meshid;
+		isects[idx].primid = isect.primid;
+		isects[idx].a = isect.a;
+		isects[idx].b = isect.b;
+
+		path.isHit = isHit;
+
+		hitbools[idx] = isHit ? 1 : 0;
+	} while (true);
 }
 
 template <bool isFirstBounce>

			dim3 blockPerGrid((hitTestCount + 64 - 1) / 64);
			dim3 threadPerBlock(64);

			hitTest << <blockPerGrid, threadPerBlock >> > (
				m_hitidx.ptr(), hitTestCount,
				m_paths.ptr(),
				m_isects.ptr(),

#endif
	}
}

#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;

__global__ void hitTest(
	const int* __restrict__ notIntersectInScreenSpaceIds,
