Normalizes ray directions when raymarching. Fixes an artifact in g3DP…

…rint!

Previously, ray directions were normalized, and then multiplied by the
application-to-voxel matrix. This meant that SetSteps specified how far the ray
should march at each step in the application's coordinate space, instead of
in terms of voxels!

As a result, `SetSteps` now specifies how far the ray should march at each step
in terms of voxels. This commit updates all of the samples to fix this (which
fixes an artifact visible on g3DPrint when using a high voxel resolution plus
the surface visualization mode would result in undersampling), which improves
rendering performance.

Also:

- Fixes ray-depth buffer intersection (#48, #95) without adding a new variable
to the GVDB scene or info struct, by adding a new function, `getRayDepthBufferMax`.

- Hopefully correct ray transformation in volume shaders (#89).

- Fixes a bug where RotateTZYXS was multiplying on the left instead of on the
right by the scaling matrix.

- Renames `SCN_PSTEP`, `SCN_SSTEP`, and `SCN_FSTEP` to use more descriptive
names.

- Adds an `m_draw_topology` flag to `gInteractiveOptix` to choose whether to
draw topology at compile-time

- Adjusted volume translation in `gInteractiveOptix` and extinction to
account for scale change

- Includes minor formatting changes

source/g3DPrint/main_3dprint.cpp

@@ -222,7 +222,7 @@ bool Sample::init ()
 
 	// Set volume params
 	printf ( "Volume params.\n" );
-	gvdb1.getScene()->SetSteps ( 0.25f, 16.f, 0.25f );			// Set raycasting steps
+	gvdb1.getScene()->SetSteps ( 0.5f, 16.f, 0.5f );			// Set raycasting steps
 	gvdb1.getScene()->SetVolumeRange ( 0.25f, 0.0f, 1.0f );		// Set volume value range
 	gvdb1.getScene()->SetExtinct ( -1.0f, 1.1f, 0.f );			// Set volume extinction	
 	gvdb1.getScene()->SetCutoff ( 0.005f, 0.005f, 0.f );
@@ -233,7 +233,7 @@ bool Sample::init ()
 	gvdb1.getScene()->SetBackgroundClr ( 0.1f, 0.2f, 0.4f, 1.0f );	
 
 #ifdef USE_GVDB2
-	gvdb2.getScene()->SetSteps(0.25f, 16.f, 0.25f);			// Set raycasting steps
+	gvdb2.getScene()->SetSteps(0.5f, 16.f, 0.5f);		// Set raycasting steps
 	gvdb2.getScene()->SetVolumeRange(0.5f, 0.0f, 1.0f);	// Set volume value range	
 	gvdb2.getScene()->SetExtinct(-1.0f, 1.5f, 0.f);		// Set volume extinction	
 	gvdb2.getScene()->SetCutoff(0.005f, 0.01f, 0.f);

source/gDepthMap/main_depthmap.cpp

@@ -82,7 +82,7 @@ bool Sample::init()
 	gvdb.UpdateApron();
 
 	// Set volume params
-	gvdb.getScene()->SetSteps ( .1f, 16, .1f );				// Set raycasting steps
+	gvdb.getScene()->SetSteps ( .5f, 16, .5f );				// Set raycasting steps
 	gvdb.getScene()->SetExtinct ( -1.0f, 1.0f, 0.0f );		// Set volume extinction
 	gvdb.getScene()->SetVolumeRange ( 0.1f, 0.0f, 0.5f );	// Set volume value range
 	gvdb.getScene()->SetCutoff ( 0.005f, 0.005f, 0.0f );

source/gFluidSurface/main_fluid_surface.cpp

@@ -244,7 +244,7 @@ bool Sample::init()
 	gvdb.AddPath ( ASSET_PATH );
 
 	// Set volume params
-	gvdb.getScene()->SetSteps ( 0.2f, 16, 0.2f );			// Set raycasting steps
+	gvdb.getScene()->SetSteps ( 0.25f, 16, 0.25f );			// Set raycasting steps
 	gvdb.getScene()->SetExtinct ( -1.0f, 1.5f, 0.0f );		// Set volume extinction
 	gvdb.getScene()->SetVolumeRange ( 0.0f, 3.0f, -1.0f );	// Set volume value range
 	gvdb.getScene()->SetCutoff ( 0.005f, 0.01f, 0.0f );

source/gImportVDB/main_import_vdb.cpp

@@ -68,7 +68,7 @@ int main (int argc, char* argv)
 	// Remainder of this sample demonstrates level set rendering to file.
 
 	// Set volume params
-	gvdb.getScene()->SetSteps ( 0.25*0.05f, 16*0.05f, 0.25*0.05f );	// Set raycasting steps (note that this is in world-space!)
+	gvdb.getScene()->SetSteps ( 0.25f, 16, 0.25f );			// Set raycasting steps per voxel
 	gvdb.getScene()->SetExtinct ( -1.0f, 1.5f, 0.0f );		// Set volume extinction
 	gvdb.getScene()->SetVolumeRange ( 0.0f, 1.0f, -1.0f );	// Set volume value range (for a level set)
 	gvdb.getScene()->SetCutoff ( 0.005f, 0.01f, 0.0f );

source/gInteractiveGL/main_interactive_gl.cpp

@@ -56,7 +56,7 @@ bool Sample::init()
 
 	// Set volume params
 	gvdb.SetTransform(m_pretrans, m_scale, m_angs, m_trans);
-	gvdb.getScene()->SetSteps ( .25, 16, .25 );				// Set raycasting steps
+	gvdb.getScene()->SetSteps ( .25f, 16, .25f );			// Set raycasting steps
 	gvdb.getScene()->SetExtinct ( -1.0f, 1.0f, 0.0f );		// Set volume extinction
 	gvdb.getScene()->SetVolumeRange ( 0.1f, 0.0f, .5f );	// Set volume value range
 	gvdb.getScene()->SetCutoff ( 0.005f, 0.005f, 0.0f );

source/gInteractiveOptix/main_interactive_optix.cpp

@@ -63,6 +63,7 @@ class Sample : public NVPWindow {
 	int			m_shading;
 	bool		m_render_optix;
 	Vector3DF	m_translate;
+	bool		m_draw_topology = false;// Whether to draw topology
 
 	int			mat_surf1;				// material id for surface objects
 	int			mat_deep;				// material id for volumetric objects
@@ -142,7 +143,12 @@ bool Sample::init()
 	max_samples = 1024;
 	m_render_optix = true;
 	m_shading = SHADE_VOLUME;
-	m_translate.Set(150, 0, 100);
+	m_translate.Set(44, 0, 16);
+
+	// Initialize debug drawing if enabled
+	if (m_draw_topology) {
+		init2D("arial");
+	}
 
 	// Initialize Optix Scene
 	if (m_render_optix) 
@@ -182,8 +188,8 @@ bool Sample::init()
 	// Set volume params		
 	gvdb.SetTransform(Vector3DF(-125, -160, -125), Vector3DF(.25, .25, .25), Vector3DF(0, 0, 0), m_translate);
 	gvdb.SetEpsilon(0.001f, 256);
-	gvdb.getScene()->SetSteps ( 0.2f, 16, 0.2f );			// SCN_PSTEP, SCN_SSTEP, SCN_FSTEP - Raycasting steps
-	gvdb.getScene()->SetExtinct ( -1.0f, 1.0f, 0.0f );		// SCN_EXTINCT, SCN_ALBEDO - Volume extinction	
+	gvdb.getScene()->SetSteps ( 0.5f, 16, 0.5f );			// SCN_DIRECTSTEP, SCN_SHADOWSTEP, SCN_FINESTEP - Raycasting steps
+	gvdb.getScene()->SetExtinct ( -0.25f, 1.0f, 0.0f );		// SCN_EXTINCT, SCN_ALBEDO - Volume extinction	
 	gvdb.getScene()->SetVolumeRange(0.1f, 0.0f, 0.3f);		// Threshold: Isoval, Vmin, Vmax
 	gvdb.getScene()->SetCutoff(0.001f, 0.001f, 0.0f);		// SCN_MINVAL, SCN_ALPHACUT
 	gvdb.getScene()->SetBackgroundClr(0.1f, 0.2f, 0.4f, 1);
@@ -288,6 +294,11 @@ void Sample::display()
 	// renders an opengl 2D texture to the screen.
 	renderScreenQuadGL ( gl_screen_tex );
 
+	if (m_draw_topology) {
+		draw_topology();
+		draw3D();
+	}
+
 	postRedisplay();
 }
 

source/gPointCloud/main_point_cloud.cpp

@@ -466,7 +466,7 @@ bool Sample::init()
 	gvdb.getScene()->SetLight(0, lgt);
 
 	// Default volume params
-	gvdb.getScene()->SetSteps(0.1f, 16, 0.1f);			// Set raycasting steps
+	gvdb.getScene()->SetSteps(0.25f, 16, 0.25f);			// Set raycasting steps
 	gvdb.getScene()->SetExtinct(-1.0f, 1.1f, 0.0f);			// Set volume extinction
 	gvdb.getScene()->SetVolumeRange(0.0f, -1.0f, 3.0f);		// Set volume value range
 	gvdb.getScene()->SetCutoff(0.005f, 0.001f, 0.0f);

source/gRenderToFile/main_rendertofile.cpp

@@ -53,7 +53,7 @@ int main (int argc, char* argv)
 	gvdb.LoadVBX ( scnpath );							// Load VBX
 
 	// Set volume params
-	gvdb.getScene()->SetSteps ( .25, 16, .25 );				// Set raycasting steps
+	gvdb.getScene()->SetSteps ( .25f, 16, .25f );			// Set raycasting steps
 	gvdb.getScene()->SetExtinct ( -1.0f, 1.5f, 0.0f );		// Set volume extinction
 	gvdb.getScene()->SetVolumeRange ( 0.1f, 0.0f, .1f );	// Set volume value range
 	gvdb.getScene()->SetCutoff ( 0.005f, 0.01f, 0.0f );

source/gResample/main_resample.cpp

@@ -203,7 +203,7 @@ bool Sample::init()
 	Rebuild ( m_VolMax, m_sparse, m_halo );
 
 	// Set volume params
-	gvdb.getScene()->SetSteps ( .5, 16, .5 );				// Set raycasting steps
+	gvdb.getScene()->SetSteps ( .5f, 16, .5f );				// Set raycasting steps
 	gvdb.getScene()->SetExtinct ( -1.0f, 1.5f, 0.0f );		// Set volume extinction
 	gvdb.getScene()->SetVolumeRange ( 0.05f, 0.0f, 1.f );	// Set volume value range
 	gvdb.getScene()->SetCutoff ( 0.001f, 0.001f, 0.0f );

source/gvdb_library/kernels/cuda_gvdb.cuh

@@ -106,7 +106,7 @@ inline __device__ float3 getRayPoint ( float3 pos, float3 dir, float t )
 inline __device__ float3 getViewRay ( float x, float y )
 {
   float3 v = x*scn.camu + y*scn.camv + scn.cams;
-  return mmult(SCN_INVXFORM, normalize(v));
+  return normalize(mmult(SCN_INVXROT, v));
 }
 
 

source/gvdb_library/kernels/cuda_gvdb_geom.cuh

@@ -72,7 +72,7 @@ inline __device__ float3 getViewRay ( float x, float y )
   #else
 	float3 v = make_float3(0,0,0);
   #endif
-  return mmult(SCN_INVXROT, normalize(v));
+  return normalize(mmult(SCN_INVXROT, v));
 }
 
 

source/gvdb_library/kernels/cuda_gvdb_raycast.cuh

@@ -183,7 +183,7 @@ inline __device__ float3 getGradientTricubic ( VDBInfo* gvdb, uchar chan, float3
 	return g;
 }
 
-// Marches along the ray p + o + rdir*t in atlas space in steps of SCN_FSTEP, using default interpolation.
+// Marches along the ray p + o + rdir*t in atlas space in steps of SCN_FINESTEP, using default interpolation.
 // If it samples a point less than SCN_THRESH, halts and returns:
 //	 `p`: The atlas-space coordinate of the intersection relative to `o`
 //   returned value: The index-space coordinate of the intersection
@@ -198,7 +198,7 @@ inline __device__ float3 getGradientTricubic ( VDBInfo* gvdb, uchar chan, float3
 //   `vmin`: The minimum AABB vertex of the brick in index-space
 __device__ float3 rayLevelSet ( VDBInfo* gvdb, uchar chan, float3& p, float3 o, float3 rpos, float3 rdir, float3 vmin )
 {
-	float dt = SCN_FSTEP;
+	float dt = SCN_FINESTEP;
 	float3 pt = dt*rdir;
 
 	for ( int i=0; i < 512; i++ ) {
@@ -297,7 +297,7 @@ __device__ void raySurfaceTrilinearBrick ( VDBInfo* gvdb, uchar chan, int nodeid
 	VDBNode* node	= getNode ( gvdb, 0, nodeid, &vmin );			// Get the VDB leaf node	
 	float3  o = make_float3( node->mValue ) ;				// Atlas sub-volume to trace	
 	float3	p = pos + t.x*dir - vmin;					// sample point in index coords			
-	t.x = SCN_PSTEP * ceil ( t.x / SCN_PSTEP );
+	t.x = SCN_DIRECTSTEP * ceil ( t.x / SCN_DIRECTSTEP );
 
 	for (int iter=0; iter < MAX_ITER && p.x >=0 && p.y >=0 && p.z >=0 && p.x < gvdb->res[0] && p.y < gvdb->res[0] && p.z < gvdb->res[0]; iter++) 
 	{	
@@ -307,8 +307,8 @@ __device__ void raySurfaceTrilinearBrick ( VDBInfo* gvdb, uchar chan, int nodeid
 			if ( gvdb->clr_chan != CHAN_UNDEF ) hclr = getColorF ( gvdb, gvdb->clr_chan, p+o );
 			return;	
 		}	
-		p += SCN_PSTEP*dir;		
-		t.x += SCN_PSTEP;
+		p += SCN_DIRECTSTEP*dir;		
+		t.x += SCN_DIRECTSTEP;
 	}
 }
 
@@ -338,16 +338,16 @@ __device__ void raySurfaceTricubicBrick ( VDBInfo* gvdb, uchar chan, int nodeid,
 	{
 		v.z = getTricubic ( gvdb, chan, p, o );
 		if ( v.z >= SCN_THRESH) {
-			v.x = getTricubic ( gvdb, chan, p - SCN_FSTEP*dir, o );
+			v.x = getTricubic ( gvdb, chan, p - SCN_FINESTEP*dir, o );
 			v.y = (v.z - SCN_THRESH)/(v.z-v.x);
-			p += -v.y*SCN_FSTEP*dir;
+			p += -v.y*SCN_FINESTEP*dir;
 			hit = p + vmin;
 			norm = getGradientTricubic ( gvdb, chan, p, o );
 			if ( gvdb->clr_chan != CHAN_UNDEF ) hclr = getColorF ( gvdb, gvdb->clr_chan, p+o );
 			return;			
 		}	
-		p += SCN_PSTEP*dir;		
-		t.x += SCN_PSTEP;
+		p += SCN_DIRECTSTEP*dir;		
+		t.x += SCN_DIRECTSTEP;
 	}
 }
 
@@ -364,6 +364,24 @@ inline __device__ float getLinearDepth(float* depthBufFloat)
 	return (-n * f / (f - n)) / (z - (f / (f - n)));				// Return linear depth
 }
 
+// Get the value of t at which the ray starting at the camera position with direction `dir` intersects the depth buffer
+// at the current thread. INFINITY if there is no depth buffer.
+inline __device__ float getRayDepthBufferMax(const float3& rayDir) {
+	if (SCN_DBUF != 0x0) {
+		// Solve
+		//   t * (length of rayDir in app space) == getLinearDepth(SCN_DBUF)
+		// for t, where (length of rayDir in app space) is length((SCN_XFORM * float4(rayDir, 0)).xyz):
+		float3 rayInWorldSpace;
+		rayInWorldSpace.x = rayDir.x * SCN_XFORM[0] + rayDir.y * SCN_XFORM[4] + rayDir.z * SCN_XFORM[ 8];
+		rayInWorldSpace.y = rayDir.x * SCN_XFORM[1] + rayDir.y * SCN_XFORM[5] + rayDir.z * SCN_XFORM[ 9];
+		rayInWorldSpace.z = rayDir.x * SCN_XFORM[2] + rayDir.y * SCN_XFORM[6] + rayDir.z * SCN_XFORM[10];
+		return getLinearDepth(SCN_DBUF) / length(rayInWorldSpace);
+	}
+	else {
+		return INFINITY;
+	}
+}
+
 #define EPSTEST(a,b,c)	(a>b-c && a<b+c)
 #define VOXEL_EPS 0.0001
 
@@ -387,7 +405,7 @@ __device__ void rayLevelSetBrick ( VDBInfo* gvdb, uchar chan, int nodeid, float3
 	VDBNode* node	= getNode ( gvdb, 0, nodeid, &vmin );			// Get the VDB leaf node	
 	float3  o = make_float3( node->mValue ) ;				// Atlas sub-volume to trace	
 	float3	p = pos + t.x*dir - vmin;					// sample point in index coords			
-	t.x = SCN_PSTEP * ceil ( t.x / SCN_PSTEP );
+	t.x = SCN_DIRECTSTEP * ceil ( t.x / SCN_DIRECTSTEP );
 
 	for (int iter=0; iter < MAX_ITER && p.x >=0 && p.y >=0 && p.z >=0 && p.x <= gvdb->res[0] && p.y <= gvdb->res[0] && p.z <= gvdb->res[0]; iter++) {	
 
@@ -399,8 +417,8 @@ __device__ void rayLevelSetBrick ( VDBInfo* gvdb, uchar chan, int nodeid, float3
 				return;
 			}
 		}	
-		p += SCN_PSTEP*dir;
-		t.x += SCN_PSTEP;
+		p += SCN_DIRECTSTEP*dir;
+		t.x += SCN_DIRECTSTEP;
 	}
 }
 
@@ -423,9 +441,9 @@ __device__ void rayEmptySkipBrick ( VDBInfo* gvdb, uchar chan, int nodeid, float
 
 // Returns deep shadow accumulation along a ray. Each sample's value is mapped to a density value using the transfer
 // function. This sample density is then treated as an opaque layer with opacity
-//   exp( SCN_EXTINCT * density * SCN_SSTEP / (1.0 + t.x * 0.4) )
-// where t.x in the above equation increments in steps of SCN_SSTEP, while the parameter of the ray increments in steps
-// of SCN_PSTEP.
+//   exp( SCN_EXTINCT * density * SCN_SHADOWSTEP / (1.0 + t.x * 0.4) )
+// where t.x in the above equation increments in steps of SCN_SHADOWSTEP, while the parameter of the ray increments in steps
+// of SCN_DIRECTSTEP.
 // Inputs:
 //   `gvdb`: The volume's `VDBInfo` object
 //   `chan`: The channel to render
@@ -445,15 +463,15 @@ __device__ void rayShadowBrick ( VDBInfo* gvdb, uchar chan, int nodeid, float3 t
 	t.y -= gvdb->epsilon;												// make sure we end insidoke	
 	float3 o = make_float3( node->mValue );					// atlas sub-volume to trace	
 	float3 p = pos + t.x*dir - vmin;		// sample point in index coords
-	float3 pt = SCN_PSTEP * dir;								// index increment	
+	float3 pt = SCN_DIRECTSTEP * dir;								// index increment	
 	float val = 0;
 
 	// accumulate remaining voxels	
 	for (; clr.w < 1 && p.x >=0 && p.y >=0 && p.z >=0 && p.x < gvdb->res[0] && p.y < gvdb->res[0] && p.z < gvdb->res[0];) {		
-		val = exp ( SCN_EXTINCT * transfer( gvdb, tex3D<float> ( gvdb->volIn[chan], p.x+o.x, p.y+o.y, p.z+o.z )).w * SCN_SSTEP/(1.0 + t.x * 0.4) );		// 0.4 = shadow gain
+		val = exp ( SCN_EXTINCT * transfer( gvdb, tex3D<float> ( gvdb->volIn[chan], p.x+o.x, p.y+o.y, p.z+o.z )).w * SCN_SHADOWSTEP/(1.0 + t.x * 0.4) );		// 0.4 = shadow gain
 		clr.w = 1.0 - (1.0-clr.w) * val;
 		p += pt;	
-		t.x += SCN_SSTEP;
+		t.x += SCN_SHADOWSTEP;
 	}	
 }
 
@@ -462,7 +480,7 @@ __device__ void rayShadowBrick ( VDBInfo* gvdb, uchar chan, int nodeid, float3 t
 // This samples in increments of `SCN_PSTEP` in `t`.
 // Each sample's value is mapped to a density value using the transfer function. This sample density is then treated as
 // an opaque layer with opacity
-//    exp(SCN_EXTINCT * val.w * SCN_PSTEP).
+//    exp(SCN_EXTINCT * val.w * SCN_DIRECTSTEP).
 // Inputs:
 //   `gvdb`: The volume's `VDBInfo` object
 //   `chan`: The channel to render
@@ -482,49 +500,49 @@ __device__ void rayDeepBrick ( VDBInfo* gvdb, uchar chan, int nodeid, float3 t,
 	float3 vmin;
 	VDBNode* node = getNode ( gvdb, 0, nodeid, &vmin );			// Get the VDB leaf node		
 
-	//t.x = SCN_PSTEP * ceil( t.x / SCN_PSTEP );						// start on sampling wavefront	
+	//t.x = SCN_DIRECTSTEP * ceil( t.x / SCN_DIRECTSTEP );						// start on sampling wavefront	
 
 	float3 o = make_float3( node->mValue );					// atlas sub-volume to trace
 	float3 wp = pos + t.x*dir;	
 	float3 p = wp-vmin;					// sample point in index coords	
-	float3 wpt = SCN_PSTEP*dir;					// world increment
+	float3 wpt = SCN_DIRECTSTEP*dir;					// world increment
 	float4 val = make_float4(0,0,0,0);
 	float4 hclr;
 	int iter = 0;
-	float dt = length(SCN_PSTEP*dir);
+	float dt = length(SCN_DIRECTSTEP*dir);
+	const float tDepthIntersection = getRayDepthBufferMax(dir); // The t.x at which the ray intersects the depth buffer
 
 	// record front hit point at first significant voxel
 	if (hit.x == 0) hit.x = t.x; // length(wp - pos);
 
 	// skip empty voxels
 	for (iter=0; val.w < SCN_MINVAL && iter < MAX_ITER && p.x >= 0 && p.y >=0 && p.z >=0 && p.x < gvdb->res[0] && p.y < gvdb->res[0] && p.z < gvdb->res[0]; iter++) {		
 		val.w = transfer ( gvdb, tex3D<float> ( gvdb->volIn[chan], p.x+o.x, p.y+o.y, p.z+o.z ) ).w;
-		p += SCN_PSTEP*dir;
+		p += SCN_DIRECTSTEP*dir;
 		wp += wpt;
 		t.x += dt;
 	}	
 
 	// accumulate remaining voxels
 	for (; clr.w > SCN_ALPHACUT && iter < MAX_ITER && p.x >=0 && p.y >=0 && p.z >=0 && p.x < gvdb->res[0] && p.y < gvdb->res[0] && p.z < gvdb->res[0]; iter++) {			
 
-		// depth buffer test [optional]
-		if (SCN_DBUF != 0x0) {
-			if (t.x > getLinearDepth(SCN_DBUF) ) {
-				hit.y = length(wp - pos);
-				hit.z = 1;
-				clr = make_float4(fmin(clr.x, 1.f), fmin(clr.y, 1.f), fmin(clr.z, 1.f), fmax(clr.w, 0.f));
-				return;
-			}
+		// Test to see if we've intersected the depth buffer (if there is no depth buffer, then this will never happen):
+		if (t.x > tDepthIntersection) {
+			hit.y = length(wp - pos);
+			hit.z = 1;
+			clr = make_float4(fmin(clr.x, 1.f), fmin(clr.y, 1.f), fmin(clr.z, 1.f), fmax(clr.w, 0.f));
+			return;
 		}
+
 		val = transfer ( gvdb, tex3D<float> ( gvdb->volIn[chan], p.x+o.x, p.y+o.y, p.z+o.z ) );
-		val.w = exp ( SCN_EXTINCT * val.w * SCN_PSTEP );
+		val.w = exp ( SCN_EXTINCT * val.w * SCN_DIRECTSTEP );
 		hclr = (gvdb->clr_chan==CHAN_UNDEF) ? make_float4(1,1,1,1) : getColorF (gvdb, gvdb->clr_chan, p+o );
 		clr.x += val.x * clr.w * (1 - val.w) * SCN_ALBEDO * hclr.x;
 		clr.y += val.y * clr.w * (1 - val.w) * SCN_ALBEDO * hclr.y;
 		clr.z += val.z * clr.w * (1 - val.w) * SCN_ALBEDO * hclr.z;
 		clr.w *= val.w;		
 
-		p += SCN_PSTEP*dir;		
+		p += SCN_DIRECTSTEP*dir;		
 		wp += wpt;		
 		t.x += dt;
 	}			
@@ -562,17 +580,16 @@ __device__ void rayCast ( VDBInfo* gvdb, uchar chan, float3 pos, float3 dir, flo
 	HDDAState dda;
 	dda.SetFromRay(pos, dir, tStart);
 	dda.Prepare(vmin, gvdb->vdel[lev]);
+	const float tDepthIntersection = getRayDepthBufferMax(dir); // The t.x at which the ray intersects the depth buffer
 
 	for (iter=0; iter < MAX_ITER && lev > 0 && lev <= gvdb->top_lev && dda.p.x >=0 && dda.p.y >=0 && dda.p.z >=0 && dda.p.x <= gvdb->res[lev] && dda.p.y <= gvdb->res[lev] && dda.p.z <= gvdb->res[lev]; iter++ ) {
 
 		dda.Next();
 
-		// depth buffer test [optional]
-		if (SCN_DBUF != 0x0) {
-			if (dda.t.x > getLinearDepth(SCN_DBUF) ) {
-				hit.z = 0;
-				return;
-			}
+		// Test to see if we've intersected the depth buffer (if there is no depth buffer, then this will never happen):
+		if (dda.t.x > tDepthIntersection) {
+			hit.z = 0;
+			return;
 		}
 
 		// node active test

source/gvdb_library/kernels/cuda_gvdb_scene.cuh

@@ -103,9 +103,9 @@ struct ALIGN(16) ScnRay {
 #define SCN_SHADE			scn.shading
 #define SCN_EXTINCT			scn.extinct.x
 #define SCN_ALBEDO			scn.extinct.y
-#define SCN_PSTEP			scn.steps.x
-#define SCN_SSTEP			scn.steps.y
-#define SCN_FSTEP			scn.steps.z
+#define SCN_DIRECTSTEP		scn.steps.x
+#define SCN_SHADOWSTEP		scn.steps.y
+#define SCN_FINESTEP		scn.steps.z
 #define SCN_MINVAL			scn.cutoff.x
 #define SCN_ALPHACUT		scn.cutoff.y
 #define SCN_THRESH			scn.thresh.x