 
 		dirToLight.normalize();
 
-		aten::ray shadowRay(path.rec.p, dirToLight);
+		shadowRays[idx].isActive = true;
+		shadowRays[idx].org = path.rec.p;
+		shadowRays[idx].dir = dirToLight;
+		shadowRays[idx].lightPos = sampleres.pos;
+		shadowRays[idx].targetLightId = lightidx;
 
 		aten::hitrecord tmpRec;
 
-		bool isHit = intersectBVH(&ctxt, shadowRay, AT_MATH_EPSILON, AT_MATH_INF, &tmpRec);
-
-		if (AT_NAME::scene::hitLight(
-			isHit, 
-			&light, 
-			sampleres.pos, 
-			shadowRay, 
+		if (light.attrib.isInfinite) {
+			len = 1.0f;
+		}
+
+		const auto c0 = max(0.0f, dot(orienting_normal, dirToLight));
+		float c1 = 1.0f;
+
+		if (!light.attrib.isSingular) {
+			c1 = max(0.0f, dot(sampleres.nml, -dirToLight));
+		}
+
+		auto G = c0 * c1 / (len * len);
+
+		path.throughput = path.throughput * (mtrl->baseColor * sampleres.finalColor) * G;
+	}
+}
+
+__global__ void hitShadowRay(
+	cudaSurfaceObject_t outSurface,
+	Path* paths,
+	ShadowRay* shadowRays,
+	int width, int height,
+	aten::ShapeParameter* shapes, int geomnum,
+	aten::MaterialParameter* mtrls,
+	aten::LightParameter* lights, int lightnum,
+	cudaTextureObject_t* nodes,
+	aten::PrimitiveParamter* prims,
+	cudaTextureObject_t vertices,
+	aten::mat4* matrices)
+{
+	const auto ix = blockIdx.x * blockDim.x + threadIdx.x;
+	const auto iy = blockIdx.y * blockDim.y + threadIdx.y;
+
+	if (ix >= width && iy >= height) {
+		return;
+	}
+
+	const auto idx = iy * width + ix;
+
+	auto& path = paths[idx];
+
+	if (path.isTerminate) {
+		return;
+	}
+
+	Context ctxt;
+	{
+		ctxt.geomnum = geomnum;
+		ctxt.shapes = shapes;
+		ctxt.mtrls = mtrls;
+		ctxt.lightnum = lightnum;
+		ctxt.lights = lights;
+		ctxt.nodes = nodes;
+		ctxt.prims = prims;
+		ctxt.vertices = vertices;
+		ctxt.matrices = matrices;
+	}
+
+	auto& shadowRay = shadowRays[idx];
+
+	if (shadowRay.isActive) {
+		aten::hitrecord rec;
+		bool isHit = intersectBVH(&ctxt, shadowRay, AT_MATH_EPSILON, AT_MATH_INF, &rec);
+
+		auto light = ctxt.lights[shadowRay.targetLightId];
+		if (light.object.idx >= 0) {
+			light.object.ptr = &ctxt.shapes[light.object.idx];
+		}
+
+		shadowRay.isActive = AT_NAME::scene::hitLight(
+			isHit,
+			&light,
+			shadowRay.lightPos,
+			shadowRay,
 			AT_MATH_EPSILON, AT_MATH_INF,
-			tmpRec.p,
-			tmpRec.t,
-			tmpRec.obj))
-		{
-			if (light.attrib.isInfinite) {
-				len = 1.0f;
-			}
-
-			const auto c0 = max(0.0f, dot(orienting_normal, dirToLight));
-			float c1 = 1.0f;
-
-			if (!light.attrib.isSingular) {
-				c1 = max(0.0f, dot(sampleres.nml, -dirToLight));
-			}
-
-			auto G = c0 * c1 / (len * len);
-
-			contrib += path.throughput * (mtrl->baseColor * sampleres.finalColor) * G;
-		}
-
-		path.isTerminate = true;
-		//p[idx] = make_float4(contrib.x, contrib.y, contrib.z, 1);
-		surf2Dwrite(make_float4(contrib.x, contrib.y, contrib.z, 1), outSurface, ix * sizeof(float4), iy, cudaBoundaryModeTrap);
+			rec.p,
+			rec.t,
+			rec.obj);
+
+		if (shadowRay.isActive) {
+			path.isTerminate = true;
+
+			auto contrib = path.throughput;
+
+			//p[idx] = make_float4(contrib.x, contrib.y, contrib.z, 1);
+			surf2Dwrite(
+				make_float4(contrib.x, contrib.y, contrib.z, 1), outSurface, ix * sizeof(float4), iy, cudaBoundaryModeTrap);
+		}
 	}
 }
+
 
 namespace idaten {
 	void RayTracing::prepare()


	auto& path = paths[idx];

	if (!path.isHit) {
		return;
	}

			//checkCudaErrors(cudaDeviceSynchronize());

			raytracing << <grid, block >> > (
				outputSurf,
				paths.ptr(),
				shadowRays.ptr(),
				width, height,
				shapeparam.ptr(), shapeparam.num(),
				mtrlparam.ptr(),
