 				AT_PRINTF("Cuda Kernel Err(raytracing) [%s]\n", cudaGetErrorString(err));
 			}
 
+			hitShadowRay << <grid, block >> > (
+				//hitShadowRay << <1, 1 >> > (
+				outputSurf,
+				paths.ptr(),
+				shadowRays.ptr(),
+				width, height,
+				shapeparam.ptr(), shapeparam.num(),
+				mtrlparam.ptr(),
+				lightparam.ptr(), lightparam.num(),
+				nodetex.ptr(),
+				primparams.ptr(),
+				vtxTex,
+				mtxparams.ptr());
+
+			err = cudaGetLastError();
+			if (err != cudaSuccess) {
+				AT_PRINTF("Cuda Kernel Err(hitShadowRay) [%s]\n", cudaGetErrorString(err));
+			}
+
 			//checkCudaErrors(cudaDeviceSynchronize());
 
 			depth++;

#include "cuda/cudamemory.h"

#include "aten4idaten.h"

struct Path {
	aten::ray ray;


		dirToLight.normalize();

		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;

		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,
			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()
