 		ctxt.matrices = matrices;
 	}
 
-	const auto idx = getIdx(ix, iy, width);
+	idx = hitindices[idx];
 
 	auto& path = paths[idx];
 	const auto& ray = rays[idx];
 
 	shadowRays[idx].isActive = false;
-
-	if (!path.isHit) {
-		return;
-	}
-	if (path.isTerminate) {
-		return;
-	}
 
 	aten::hitrecord rec;
 

__global__ void shade(
	cudaSurfaceObject_t outSurface,
	Path* paths,
	aten::Intersection* isects,
	aten::ray* rays,
	ShadowRay* shadowRays,
	int width, int height,
	int depth, int rrDepth,
	aten::ShapeParameter* shapes, int geomnum,
	aten::MaterialParameter* mtrls,

					m_hitbools,
					&hitcount);

				dim3 blockPerGrid((hitcount + 256 - 1) / 256);
				dim3 threadPerBlock(256);

				shade << <blockPerGrid, threadPerBlock >> > (
				//shade << <1, 1 >> > (
					outputSurf,
					paths.ptr(),
					m_hitidx.ptr(), hitcount,
					isects.ptr(),
					rays.ptr(),
					shadowRays.ptr(),
					depth, rrDepth,
					shapeparam.ptr(), shapeparam.num(),
					mtrlparam.ptr(),
