 		cam.ptr(),
 		spheres.ptr(), AT_COUNTOF(g_spheres),
 		materials.ptr(),
-		lightmem.ptr());
+		lightmem.ptr(), lightmem.num());
 
 	checkCudaErrors(cudaDeviceSynchronize());
 

	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;
