 #include "cuda/cudamemory.h"
 
 #include "aten4idaten.h"
+
+struct ShadowRay : public aten::ray {
+	aten::vec3 lightPos;
+	int targetLightId;
+
+	struct {
+		uint32_t isActive : 1;
+	};
+};
 
 struct Path {
 	aten::ray ray;


		dirToLight.normalize();

		aten::ray shadowRay(path.rec.p, dirToLight);

		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, 
			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);
	}
}

namespace idaten {
	void RayTracing::prepare()

			//checkCudaErrors(cudaDeviceSynchronize());

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