 		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>

	printf("P0 (%f, %f), dP (%f, %f)\n", P0.x, P0.y, dP.x, dP.y);
#endif

	static const int maxSteps = 50;

	for (;
		((PQk.x * stepDir) <= end)	// I_ɓB.
		&& (stepCount < maxSteps)	// ő又ɓB.
		&& (sceneZMax != 0.0);	// ȂƂɓBĂȂ.
		++stepCount)
	{
		// OZ̍ől̍ŏlɂȂ.
		rayZMin = prevZMaxEstimate;

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

			//hitTest << <blockPerGrid, threadPerBlock >> > (
			hitTest << <NUM_BLOCK, dim3(WARP_SIZE, NUM_WARP_PER_BLOCK) >> > (
				m_hitidx.ptr(), hitTestCount,
				m_paths.ptr(),
				m_isects.ptr(),
