 					m_hitbools,
 					&hitcount);
 
-				AT_PRINTF("Hit Count %d\n", hitcount);
-
-				shade << <grid, block >> > (
+				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(),
-					width, height,
 					depth, rrDepth,
 					shapeparam.ptr(), shapeparam.num(),
 					mtrlparam.ptr(),

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

		ctxt.matrices = matrices;
	}

	idx = hitindices[idx];

	auto& shadowRay = shadowRays[idx];

