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

					AT_PRINTF("Cuda Kernel Err(shade) [%s]\n", cudaGetErrorString(err));
				}

				hitShadowRay << <grid, block >> > (
				//hitShadowRay << <1, 1 >> > (
					paths.ptr(),
					shadowRays.ptr(),
					width, height,
					shapeparam.ptr(), shapeparam.num(),
					mtrlparam.ptr(),
					lightparam.ptr(), lightparam.num(),

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