 					AT_PRINTF("Cuda Kernel Err(shade) [%s]\n", cudaGetErrorString(err));
 				}
 
-				hitShadowRay << <grid, block >> > (
+				hitShadowRay << <blockPerGrid, threadPerBlock >> > (
 				//hitShadowRay << <1, 1 >> > (
 					paths.ptr(),
+					m_hitidx.ptr(), hitcount,
 					shadowRays.ptr(),
-					width, height,
 					shapeparam.ptr(), shapeparam.num(),
 					mtrlparam.ptr(),
 					lightparam.ptr(), lightparam.num(),

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


		ctxt.matrices = matrices;
	}

	idx = hitindices[idx];

	auto& shadowRay = shadowRays[idx];

