 	
 	bool isHit = intersectBVH(&ctxt, rays[idx], AT_MATH_EPSILON, AT_MATH_INF, &isect);
 
-	path.isect.t = isect.t;
-	path.isect.objid = isect.objid;
-	path.isect.mtrlid = isect.mtrlid;
-	path.isect.area = isect.area;
-	path.isect.primid = isect.primid;
-	path.isect.a = isect.a;
-	path.isect.b = isect.b;
+	isects[idx].t = isect.t;
+	isects[idx].objid = isect.objid;
+	isects[idx].mtrlid = isect.mtrlid;
+	isects[idx].area = isect.area;
+	isects[idx].primid = isect.primid;
+	isects[idx].a = isect.a;
+	isects[idx].b = isect.b;
 
 	path.isHit = isHit;
 }

__global__ void shade(
	cudaSurfaceObject_t outSurface,
	Path* paths,
	aten::ray* rays,
	ShadowRay* shadowRays,
	int width, int height,

		idaten::TypedCudaMemory<Path> paths;
		paths.init(width * height);

		idaten::TypedCudaMemory<aten::Intersection> isects;
		isects.init(width * height);

		idaten::TypedCudaMemory<aten::ray> rays;
		rays.init(width * height);

