 
 	Context ctx;
 	{
-		ctx.geomnum = num;
+		ctx.geomnum = geomnum;
 		ctx.shapes = shapes;
 		ctx.mtrls = mtrls;
+		ctx.lightnum = lightnum;
+		ctx.lights = lights;
 	}
 
 	const auto idx = iy * camera->width + ix;

	aten::ShapeParameter* shapes;

	aten::MaterialParameter* mtrls;
};

struct CameraSampleResult {

	float4* p,
	int width, int height,
	Camera* camera,
	aten::ShapeParameter* shapes, int geomnum,
	aten::MaterialParameter* mtrls,
	aten::LightParameter* lights, int lightnum)
{
	const auto ix = blockIdx.x * blockDim.x + threadIdx.x;
	const auto iy = blockIdx.y * blockDim.y + threadIdx.y;
