Merged experimental lighttree code.

apps/benchmarkapp/camera.xml

@@ -1,12 +1,12 @@
 <camera>
-    <transform m00="0.98938847" m01="-0.022192042" m02="-0.14359" m03="-2.8366079" m10="0" m11="0.98826677" m12="-0.15273803" m13="0.45841199" m20="-0.1452948" m21="-0.15111725" m22="-0.97777963" m23="30.572617" m30="0" m31="0" m32="0" m33="1"/>
+    <transform m00="-0.58767647" m01="-0.25613666" m02="-0.76748312" m03="9.9049749" m10="0" m11="0.9485687" m12="-0.31657138" m13="10.969443" m20="-0.80909604" m21="0.18604155" m22="0.55745149" m23="-17.508549" m30="0" m31="0" m32="0" m33="1"/>
     <FOV>40</FOV>
     <brightness>0</brightness>
     <contrast>0</contrast>
     <gamma>2.2</gamma>
-    <aperture>0.0089872945</aperture>
+    <aperture>0.0089999996</aperture>
     <distortion>0.050000001</distortion>
-    <focalDistance>5.0255618</focalDistance>
+    <focalDistance>21.969528</focalDistance>
     <clampValue>2.5</clampValue>
     <tonemapper>3</tonemapper>
 </camera>

lib/CUDA/shared_kernel_code/lights_shared.h

@@ -22,6 +22,7 @@
 #include "noerrors.h"
 
 #define ISLIGHTS
+// #define LIGHTTREE
 #define MAXISLIGHTS	64
 
 #define TRILIGHTCOUNT			(lightCounts.x & 0xffff)
@@ -112,6 +113,59 @@ LH2_DEVFUNC float CalculateLightPDF( const float3& D, const float t, const float
 	return (t * t) / (abs( dot( D, lightNormal ) ) * lightArea);
 }
 
+//  +-----------------------------------------------------------------------------+
+//  |  CalculateChildNodeWeights                                                  |
+//  |  Helper to compute child node weights.                                LH2'20|
+//  +-----------------------------------------------------------------------------+
+LH2_DEVFUNC float CalculateChildNodeWeights( const int node, const float3& I, const float3& N, uint& seed, const bool debug = false )
+{
+	const int left = lightTree[node].left;
+	const int right = lightTree[node].right;
+	const float3 b1j = make_float3( lightTree[left].bmin ), b1k = make_float3( lightTree[right].bmin );
+	const float3 b2j = make_float3( lightTree[left].bmax ), b2k = make_float3( lightTree[right].bmax );
+	const float3 diag_j = b2j - b1j;
+	const float3 diag_k = b2k - b1k;
+	const float3 LN = make_float3( lightTree[node].N );
+	// calculate (squared) minimum and maximum distance from I to aabb
+	// method: GPU-Accelerated Minimum Distance and Clearance Queries, Krishnamurthy et al., 2011
+	const float3 Bj = 0.5f * diag_j;
+	const float3 Bk = 0.5f * diag_k;
+	const float3 Cj = (b1j + b2j) * 0.5f;
+	const float3 Ck = (b1k + b2k) * 0.5f;
+	const float3 Dj = Cj - I;
+	const float3 Dk = Ck - I;
+	const float3 min_j = make_float3( max( Dj.x - Bj.x, 0.0f ), max( Dj.y - Bj.y, 0.0f ), max( Dj.z - Bj.z, 0.0f ) );
+	const float3 min_k = make_float3( max( Dk.x - Bk.x, 0.0f ), max( Dk.y - Bk.y, 0.0f ), max( Dk.z - Bk.z, 0.0f ) );
+	const float dist2j = dot( min_j, min_j );
+	const float dist2k = dot( min_k, min_k );
+	const float3 max_j = Dj + Bj;
+	const float3 max_k = Dk + Bk;
+	const float dist2j_max = dot( max_j, max_j );
+	const float dist2k_max = dot( max_k, max_k );
+	// get the left and right cluster intensities
+	const float Ij = lightTree[left].intensity;
+	const float Ik = lightTree[right].intensity;
+	// get a reasonable value for F using the normals at I and the light
+	const float3 Rj = b1j + (b2j - b1j) * make_float3( RandomFloat( seed ), RandomFloat( seed ), RandomFloat( seed ) );
+	const float3 Rk = b1k + (b2k - b1k) * make_float3( RandomFloat( seed ), RandomFloat( seed ), RandomFloat( seed ) );
+	const float3 Lj = normalize( Rj - I );
+	const float3 Lk = normalize( Rk - I );
+	float Fj = max( 0.001f, dot( N, Lj ) );
+	float Fk = max( 0.001f, dot( N, Lk ) );
+	if (dot( LN, LN ) > 0.001f)
+		Fj *= max( 0.001f, dot( LN, Lj * -1.0f ) ),
+		Fk *= max( 0.001f, dot( LN, Lk * -1.0f ) );
+	// calculate final probabilities according to the realtime stochastic lightcuts paper
+	const bool insideBoth = dist2j == 0 && dist2k == 0;
+	const float wmin_j = (Fj * Ij) / (insideBoth ? 1 : max( 0.0001f, dist2j) );
+	const float wmin_k = (Fk * Ik) / (insideBoth ? 1 : max( 0.0001f, dist2k) );
+	const float wmax_j = (Fj * Ij) / max( 0.0001f, dist2j_max );
+	const float wmax_k = (Fj * Ij) / max( 0.0001f, dist2k_max );
+	const float pmin_j = wmin_j / (wmin_j + wmin_k);
+	const float pmax_j = wmax_j / (wmax_j + wmax_k);
+	return 0.5f * (pmin_j + pmax_j);
+}
+
 //  +-----------------------------------------------------------------------------+
 //  |  LightPickProb                                                              |
 //  |  Calculates the probability with which the specified light woukd be picked  |
@@ -135,6 +189,33 @@ LH2_DEVFUNC float LightPickProb( int idx, const float3& O, const float3& N, cons
 #endif
 }
 
+//  +-----------------------------------------------------------------------------+
+//  |  LightPickProbLTree                                                         |
+//  |  Calculates the probability with which the specified light woukd be picked  |
+//  |  from the specified world space location and normal using the stochastic    |
+//  |  lightcuts approach.                                                  LH2'20|
+//  +-----------------------------------------------------------------------------+
+LH2_DEVFUNC float LightPickProbLTree( int idx, const float3& O, const float3& N, const float3& I, uint& seed )
+{
+#ifndef LIGHTTREE
+	return LightPickProb( idx, O, N, I );
+#else
+	LightCluster* tree = lightTree;
+	int node = idx + 1; // leaf for light i is at index i + 1, see UpdateLightTree in rendercore.cpp.
+	float pickProb = 1;
+	while (1)
+	{
+		if (node == 0) break; // we are the root node
+		// determine probability of selecting the current node over its sibling
+		int parent = __float_as_int( tree[node].N.w /* we abused N.w to store the parent node index */ );
+		const float p = CalculateChildNodeWeights( parent, I, N, seed );
+		if (tree[parent].left == node) pickProb *= p /* we are the left child */; else pickProb *= 1 - p;
+		node = parent;
+	}
+	return pickProb;
+#endif
+}
+
 //  +-----------------------------------------------------------------------------+
 //  |  RandomPointOnLight                                                         |
 //  |  Selects a random point on a random light. Returns a position, a normal on  |
@@ -231,6 +312,88 @@ LH2_DEVFUNC float3 RandomPointOnLight( float r0, float r1, const float3& I, cons
 	}
 }
 
+//  +-----------------------------------------------------------------------------+
+//  |  RandomPointOnLightLTree                                                    |
+//  |  Selects a random point on a random light, using the stochastic lightcuts   |
+//  |  approach, via a binary light tree. Default method for the Optix7 core.     |
+//  |  Returns a position, a normal on the light source, the pick probability,    |
+//  |  and the importance of the explicit connection.                       LH2'20|
+//  +-----------------------------------------------------------------------------+
+LH2_DEVFUNC float3 RandomPointOnLightLTree( float r0, float r1, uint& seed, const float3& I, const float3& N, float& pickProb, float& lightPdf, float3& lightColor, const bool debug = false )
+{
+#ifndef LIGHTTREE
+	return RandomPointOnLight( r0, r1, I, N, pickProb, lightPdf, lightColor );
+#else
+	LightCluster* tree = lightTree;
+	int node = 0; // index of root node
+	int lightIdx = 0;
+	pickProb = 1;
+	while (1)
+	{
+		if (tree[node].left == -1)
+		{
+			// reached a leaf, use this light
+			lightIdx = tree[node].light;
+			break;
+		}
+		// interior node; randomly pick a child
+		const float p_j = CalculateChildNodeWeights( node, I, N, seed, debug );
+		// decide
+		if (r1 < p_j)
+			node = tree[node].left, r1 *= 1.0f / p_j, pickProb *= p_j;
+		else
+			node = tree[node].right, r1 = (r1 - p_j) / (1 - p_j), pickProb *= 1 - p_j;
+	}
+	if (lightIdx & (1 << 30))
+	{
+		// pick a pointlight
+		const CorePointLight4& light = (const CorePointLight4&)pointLights[lightIdx - (1 << 30)];
+		const float3 P = make_float3( light.data0 );	// position
+		const float3 L = P - I;
+		const float sqDist = dot( L, L );
+		lightColor = make_float3( light.data1 ) / sqDist;		// radiance
+		lightPdf = dot( L, N ) > 0 ? 1 : 0;
+		return P;
+	}
+	else if (lightIdx & (1 << 29))
+	{
+		// spotlight
+		const CoreSpotLight4& light = (const CoreSpotLight4&)spotLights[lightIdx - (1 << 29)];
+		const float4 V0 = light.data0;				// position + cos_inner
+		const float4 V1 = light.data1;				// radiance + cos_outer
+		const float4 D = light.data2;				// direction
+		const float3 P = make_float3( V0 );
+		float3 L = I - P;
+		const float sqDist = dot( L, L );
+		L = normalize( L );
+		float d = (max( 0.0f, L.x * D.x + L.y * D.y + L.z * D.z ) - V1.w) / (V0.w - V1.w);
+		const float LNdotL = min( 1.0f, d );
+		lightPdf = (LNdotL > 0 && dot( L, N ) < 0) ? (sqDist / LNdotL) : 0;
+		lightColor = make_float3( V1 );
+		return P;
+	}
+	else
+	{
+		// light triangle
+		float3 bary = RandomBarycentrics( r0 );
+		const CoreLightTri4& light = (const CoreLightTri4&)triLights[lightIdx];
+		const float4 V0 = light.data3;				// vertex0
+		const float4 V1 = light.data4;				// vertex1
+		const float4 V2 = light.data5;				// vertex2
+		lightColor = make_float3( light.data2 );	// radiance
+		const float4 LN = light.data1;				// N
+		const float3 P = make_float3( bary.x * V0 + bary.y * V1 + bary.z * V2 );
+		float3 L = I - P;							// reversed: from light to intersection point
+		const float sqDist = dot( L, L );
+		L = normalize( L );
+		const float LNdotL = L.x * LN.x + L.y * LN.y + L.z * LN.z;
+		const float reciSolidAngle = sqDist / (LN.w * LNdotL); // LN.w contains area
+		lightPdf = (LNdotL > 0 && dot( L, N ) < 0) ? reciSolidAngle : 0;
+		return P;
+	}
+#endif
+}
+
 //  +-----------------------------------------------------------------------------+
 //  |  Sample_Le                                                                  |
 //  |  Part of the BDPT core.                                               LH2'19|

lib/RenderCore_Optix7Adaptive/kernels/.cuda.cu

@@ -34,6 +34,7 @@ __constant__ int skywidth;
 __constant__ int skyheight;
 __constant__ PathState* pathStates;
 __constant__ float4* debugData;
+__constant__ LightCluster* lightTree;
 
 __constant__ mat4 worldToSky;
 
@@ -44,7 +45,7 @@ __constant__ __device__ float clampValue;
 // staging: copies will be batched and carried out after rendering completes, 
 // to allow the CPU to update the scene concurrently with GPU rendering.
 
-enum { INSTS = 0, MATS, TLGHTS, PLGHTS, SLGHTS, DLGHTS, LCNTS, RGB32, RGBH, NRMLS, SKYPIX, SKYW, SKYH, SMAT, DBGDAT, GEPS, CLMPV };
+enum { INSTS = 0, MATS, TLGHTS, PLGHTS, SLGHTS, DLGHTS, LCNTS, RGB32, RGBH, NRMLS, SKYPIX, SKYW, SKYH, SMAT, DBGDAT, GEPS, CLMPV, LTREE };
 
 // device pointers are not real pointers for nvcc, so we need a bit of a hack.
 
@@ -74,6 +75,7 @@ __host__ static void pushPtrCpy( int id, void* p )
 	if (id == NRMLS) cudaMemcpyToSymbol( nrm32, &p, sizeof( void* ) );
 	if (id == SKYPIX) cudaMemcpyToSymbol( skyPixels, &p, sizeof( void* ) );
 	if (id == DBGDAT) cudaMemcpyToSymbol( debugData, &p, sizeof( void* ) );
+	if (id == LTREE) cudaMemcpyToSymbol( lightTree, &p, sizeof( void* ) );
 }
 __host__ static void pushIntCpy( int id, const int v )
 {
@@ -151,6 +153,7 @@ __host__ void stageWorldToSky( const mat4& worldToLight ) { stageMatCpy( SMAT /*
 __host__ void stageDebugData( float4* p ) { stagePtrCpy( DBGDAT /* debugData */, p ); }
 __host__ void stageGeometryEpsilon( float e ) { stageF32Cpy( GEPS /* geometryEpsilon */, e ); }
 __host__ void stageClampValue( float c ) { stageF32Cpy( CLMPV /* clampValue */, c ); }
+__host__ void stageLightTree( LightCluster* t ) { stagePtrCpy( LTREE /* light tree */, t ); }
 __host__ void stageLightCounts( int tri, int point, int spot, int directional )
 {
 	const int4 counts = make_int4( tri, point, spot, directional );

lib/RenderCore_Optix7Filter/kernels/.cuda.cu

@@ -34,6 +34,7 @@ __constant__ int skywidth;
 __constant__ int skyheight;
 __constant__ PathState* pathStates;
 __constant__ float4* debugData;
+__constant__ LightCluster* lightTree;
 
 __constant__ mat4 worldToSky;
 
@@ -44,7 +45,7 @@ __constant__ __device__ float clampValue;
 // staging: copies will be batched and carried out after rendering completes, 
 // to allow the CPU to update the scene concurrently with GPU rendering.
 
-enum { INSTS = 0, MATS, TLGHTS, PLGHTS, SLGHTS, DLGHTS, LCNTS, RGB32, RGBH, NRMLS, SKYPIX, SKYW, SKYH, SMAT, DBGDAT, GEPS, CLMPV };
+enum { INSTS = 0, MATS, TLGHTS, PLGHTS, SLGHTS, DLGHTS, LCNTS, RGB32, RGBH, NRMLS, SKYPIX, SKYW, SKYH, SMAT, DBGDAT, GEPS, CLMPV, LTREE };
 
 // device pointers are not real pointers for nvcc, so we need a bit of a hack.
 
@@ -74,6 +75,7 @@ __host__ static void pushPtrCpy( int id, void* p )
 	if (id == NRMLS) cudaMemcpyToSymbol( nrm32, &p, sizeof( void* ) );
 	if (id == SKYPIX) cudaMemcpyToSymbol( skyPixels, &p, sizeof( void* ) );
 	if (id == DBGDAT) cudaMemcpyToSymbol( debugData, &p, sizeof( void* ) );
+	if (id == LTREE) cudaMemcpyToSymbol( lightTree, &p, sizeof( void* ) );
 }
 __host__ static void pushIntCpy( int id, const int v )
 {
@@ -117,6 +119,7 @@ __host__ void stageWorldToSky( const mat4& worldToLight ) { stageMatCpy( SMAT /*
 __host__ void stageDebugData( float4* p ) { stagePtrCpy( DBGDAT /* debugData */, p ); }
 __host__ void stageGeometryEpsilon( float e ) { stageF32Cpy( GEPS /* geometryEpsilon */, e ); }
 __host__ void stageClampValue( float c ) { stageF32Cpy( CLMPV /* clampValue */, c ); }
+__host__ void stageLightTree( LightCluster* t ) { stagePtrCpy( LTREE /* light tree */, t ); }
 __host__ void stageLightCounts( int tri, int point, int spot, int directional )
 {
 	const int4 counts = make_int4( tri, point, spot, directional );

lib/RenderCore_Optix7Guiding/kernels/.cuda.cu

@@ -36,6 +36,7 @@ __constant__ int skywidth;
 __constant__ int skyheight;
 __constant__ PathState* pathStates;
 __constant__ float4* debugData;
+__constant__ LightCluster* lightTree;
 
 __constant__ mat4 worldToSky;
 
@@ -46,7 +47,7 @@ __constant__ __device__ float clampValue;
 // staging: copies will be batched and carried out after rendering completes, 
 // to allow the CPU to update the scene concurrently with GPU rendering.
 
-enum { INSTS = 0, MATS, TLGHTS, PLGHTS, SLGHTS, DLGHTS, LCNTS, RGB32, RGBH, NRMLS, SKYPIX, SKYW, SKYH, SMAT, DBGDAT, GEPS, CLMPV, GUIDANCE, PMIN, PEXT };
+enum { INSTS = 0, MATS, TLGHTS, PLGHTS, SLGHTS, DLGHTS, LCNTS, RGB32, RGBH, NRMLS, SKYPIX, SKYW, SKYH, SMAT, DBGDAT, GEPS, CLMPV, GUIDANCE, PMIN, PEXT, LTREE };
 
 // device pointers are not real pointers for nvcc, so we need a bit of a hack.
 
@@ -78,6 +79,7 @@ __host__ static void pushPtrCpy( int id, void* p )
 	if (id == NRMLS) cudaMemcpyToSymbol( nrm32, &p, sizeof( void* ) );
 	if (id == SKYPIX) cudaMemcpyToSymbol( skyPixels, &p, sizeof( void* ) );
 	if (id == DBGDAT) cudaMemcpyToSymbol( debugData, &p, sizeof( void* ) );
+	if (id == LTREE) cudaMemcpyToSymbol( lightTree, &p, sizeof( void* ) );
 	if (id == GUIDANCE) cudaMemcpyToSymbol( guidance, &p, sizeof( void* ) );
 }
 __host__ static void pushIntCpy( int id, const int v )
@@ -170,6 +172,7 @@ __host__ void stageWorldToSky( const mat4& worldToLight ) { stageMatCpy( SMAT /*
 __host__ void stageDebugData( float4* p ) { stagePtrCpy( DBGDAT /* debugData */, p ); }
 __host__ void stageGeometryEpsilon( float e ) { stageF32Cpy( GEPS /* geometryEpsilon */, e ); }
 __host__ void stageClampValue( float c ) { stageF32Cpy( CLMPV /* clampValue */, c ); }
+__host__ void stageLightTree( LightCluster* t ) { stagePtrCpy( LTREE /* light tree */, t ); }
 __host__ void stageLightCounts( int tri, int point, int spot, int directional )
 {
 	const int4 counts = make_int4( tri, point, spot, directional );

lib/RenderCore_Optix7Guiding/kernels/pathtracer.h

@@ -65,7 +65,7 @@ LH2_DEVFUNC float3 RandomPointOnLightPNEE( float r0, float r1, const float3& I,
 	if (g1.x == 0)
 	{
 		// nothing here; sample lights uniformly
-		lightIdx = (int)(r1 * (TRILIGHTCOUNT + POINTLIGHTCOUNT));
+		lightIdx = min( (int)(r1 * (TRILIGHTCOUNT + POINTLIGHTCOUNT)), TRILIGHTCOUNT + POINTLIGHTCOUNT - 1 );
 		pickProb = 1.0f / (TRILIGHTCOUNT + POINTLIGHTCOUNT);
 	}
 	else
@@ -374,9 +374,9 @@ void shadeKernel( float4* accumulator, const uint stride,
 	// next event estimation: connect eye path to light
 	if ((FLAGS & S_SPECULAR) == 0 && connections != 0) // skip for specular vertices
 	{
-		float pickProb, lightPdf = 0;
+		float pickProb, lightPdf = (pixelIdx == probePixelIdx && pathLength == 1) ? -999 : 0;
 		float3 lightColor, L;
-		if (pixelIdx % SCRWIDTH < SCRWIDTH / 2)
+		if (1) // pixelIdx % SCRWIDTH < SCRWIDTH / 2)
 		{
 			float3 jitter = make_float3( 1.0f / PNEErext.x, 1.0f / PNEErext.y, 1.0f / PNEErext.z );
 			float3 II = make_float3(
@@ -410,7 +410,7 @@ void shadeKernel( float4* accumulator, const uint stride,
 				const uint shadowRayIdx = atomicAdd( &counters->shadowRays, 1 ); // compaction
 				connections[shadowRayIdx] = make_float4( SafeOrigin( I, L, N, geometryEpsilon ), 0 ); // O4
 				connections[shadowRayIdx + stride * 2] = make_float4( L, dist - 2 * geometryEpsilon ); // D4
-				connections[shadowRayIdx + stride * 2 * 2] = make_float4( contribution, __int_as_float( pixelIdx ) ); // E4
+				connections[shadowRayIdx + stride * 2 * 2] = make_float4( contribution, __int_as_float( 0 ) ); // E4
 			}
 		}
 	}

lib/RenderCore_OptixPrime_B/kernels/.cuda.cu

@@ -33,6 +33,7 @@ __constant__ float4* skyPixels;
 __constant__ int skywidth;
 __constant__ int skyheight;
 __constant__ float4* debugData;
+__constant__ LightCluster* lightTree;
 
 __constant__ mat4 worldToSky;
 
@@ -43,7 +44,7 @@ __constant__ __device__ float clampValue;
 // staging: copies will be batched and carried out after rendering completes, 
 // to allow the CPU to update the scene concurrently with GPU rendering.
 
-enum { INSTS = 0, MATS, TLGHTS, PLGHTS, SLGHTS, DLGHTS, LCNTS, RGB32, RGBH, NRMLS, SKYPIX, SKYW, SKYH, SMAT, DBGDAT, GEPS, CLMPV };
+enum { INSTS = 0, MATS, TLGHTS, PLGHTS, SLGHTS, DLGHTS, LCNTS, RGB32, RGBH, NRMLS, SKYPIX, SKYW, SKYH, SMAT, DBGDAT, GEPS, CLMPV, LTREE };
 
 // device pointers are not real pointers for nvcc, so we need a bit of a hack.
 
@@ -73,6 +74,7 @@ __host__ static void pushPtrCpy( int id, void* p )
 	if (id == NRMLS) cudaMemcpyToSymbol( nrm32, &p, sizeof( void* ) );
 	if (id == SKYPIX) cudaMemcpyToSymbol( skyPixels, &p, sizeof( void* ) );
 	if (id == DBGDAT) cudaMemcpyToSymbol( debugData, &p, sizeof( void* ) );
+	if (id == LTREE) cudaMemcpyToSymbol( lightTree, &p, sizeof( void* ) );
 }
 __host__ static void pushIntCpy( int id, const int v )
 {
@@ -150,6 +152,7 @@ __host__ void stageWorldToSky( const mat4& worldToLight ) { stageMatCpy( SMAT /*
 __host__ void stageDebugData( float4* p ) { stagePtrCpy( DBGDAT /* debugData */, p ); }
 __host__ void stageGeometryEpsilon( float e ) { stageF32Cpy( GEPS /* geometryEpsilon */, e ); }
 __host__ void stageClampValue( float c ) { stageF32Cpy( CLMPV /* clampValue */, c ); }
+__host__ void stageLightTree( LightCluster* t ) { stagePtrCpy( LTREE /* light tree */, t ); }
 __host__ void stageLightCounts( int tri, int point, int spot, int directional )
 {
 	const int4 counts = make_int4( tri, point, spot, directional );

lib/RenderCore_OptixPrime_B/kernels/pathtracer.h

@@ -1,4 +1,4 @@
-/* pathtracer.cu - Copyright 2019/2020 Utrecht University
+/* pathtracer.h - Copyright 2019/2020 Utrecht University
 
    Licensed under the Apache License, Version 2.0 (the "License");
    you may not use this file except in compliance with the License.

lib/RenderCore_OptixPrime_BDPT/kernels/.cuda.cu

@@ -33,6 +33,7 @@ __constant__ float4* skyPixels;
 __constant__ int skywidth;
 __constant__ int skyheight;
 __constant__ float4* debugData;
+__constant__ LightCluster* lightTree;
 
 __constant__ mat4 worldToSky;
 
@@ -50,6 +51,7 @@ __host__ void stageTriLights( CoreLightTri* p ) { stagedcpy( triLights, p ); }
 __host__ void stagePointLights( CorePointLight* p ) { stagedcpy( pointLights, p ); }
 __host__ void stageSpotLights( CoreSpotLight* p ) { stagedcpy( spotLights, p ); }
 __host__ void stageDirectionalLights( CoreDirectionalLight* p ) { stagedcpy( directionalLights, p ); }
+__host__ void stageLightTree( LightCluster* t ) { stagedcpy( lightTree, t ); }
 __host__ void stageLightCounts( int tri, int point, int spot, int directional )
 {
 	const int4 counts = make_int4( tri, point, spot, directional );