Redo node format to include number of children, rather than just leaf bit.

5 files changed, 30 insertions, 16 deletions

diff --git a/chroma/cuda/bvh.cu b/chroma/cuda/bvh.cu
index 2ce0580..5b40bf7 100644
--- a/chroma/cuda/bvh.cu
+++ b/chroma/cuda/bvh.cu
@@ -1,6 +1,7 @@
 //-*-c++-*-
 #include <cuda.h>
 
+#include "geometry_types.h"
 #include "linalg.h"
 #include "physical_constants.h"
 
@@ -123,8 +124,6 @@ __device__ unsigned long long surface_half_area(const uint4 &node)
   return x*y + y*z + z*x;
 }
 
-const unsigned int LEAF_BIT = (1U << 31);
-
 extern "C"
 {
 
@@ -192,7 +191,7 @@ extern "C"
     leaf_node.x = q_lower.x | (q_upper.x << 16);
     leaf_node.y = q_lower.y | (q_upper.y << 16);
     leaf_node.z = q_lower.z | (q_upper.z << 16);
-    leaf_node.w = triangle_id | LEAF_BIT;
+    leaf_node.w = triangle_id;
 
     leaf_nodes[triangle_id] = leaf_node;
     morton_codes[triangle_id] = morton;
@@ -238,12 +237,15 @@ extern "C"
     
 
     // Scan remaining children
+    unsigned int real_children = 1;
     for (unsigned int i=1; i < n_children_per_node; i++) {
       uint4 child_node = nodes[first_child + i];
       
       if (child_node.x == 0)
 	break;  // Hit first padding node in list of children
 
+      real_children++;
+
       uint3 child_lower = make_uint3(child_node.x & 0xFFFF, child_node.y & 0xFFFF, child_node.z & 0xFFFF);
       uint3 child_upper = make_uint3(child_node.x >> 16, child_node.y >> 16, child_node.z >> 16);
 
@@ -251,7 +253,7 @@ extern "C"
       upper = max(upper, child_upper);
     }
 
-    parent_node.w = first_child;
+    parent_node.w = (real_children << CHILD_BITS) | first_child;
     parent_node.x = upper.x << 16 | lower.x;
     parent_node.y = upper.y << 16 | lower.y;
     parent_node.z = upper.z << 16 | lower.z;
@@ -284,12 +286,15 @@ extern "C"
     uint3 upper = make_uint3(parent_node.x >> 16, parent_node.y >> 16, parent_node.z >> 16);
     
     // Scan remaining children
+    unsigned int real_children = 1;
     for (unsigned int i=1; i < n_children_per_node; i++) {
       uint4 child_node = child_nodes[first_child + i];
       
       if (child_node.x == 0)
 	break;  // Hit first padding node in list of children
 
+      real_children++;
+
       uint3 child_lower = make_uint3(child_node.x & 0xFFFF, child_node.y & 0xFFFF, child_node.z & 0xFFFF);
       uint3 child_upper = make_uint3(child_node.x >> 16, child_node.y >> 16, child_node.z >> 16);
 
@@ -297,7 +302,8 @@ extern "C"
       upper = max(upper, child_upper);
     }
 
-    parent_node.w = first_child + child_id_offset;
+    parent_node.w = (real_children << CHILD_BITS)
+      | (first_child + child_id_offset);
     parent_node.x = upper.x << 16 | lower.x;
     parent_node.y = upper.y << 16 | lower.y;
     parent_node.z = upper.z << 16 | lower.z;
@@ -320,9 +326,9 @@ extern "C"
     unsigned int node_id = first_node + thread_id;
     uint4 src_node = src_nodes[node_id];
     
-    unsigned int leaf_flag = src_node.w & 0x80000000;
-    unsigned int child_id = src_node.w &  0x7FFFFFFF;
-    src_node.w = leaf_flag | (child_id + child_id_offset);
+    unsigned int nchild = src_node.w >> CHILD_BITS;
+    unsigned int child_id = src_node.w &  ~NCHILD_MASK;
+    src_node.w = (nchild << CHILD_BITS) | (child_id + child_id_offset);
 
     dest_nodes[node_id] = src_node;    
   }
diff --git a/chroma/cuda/geometry.h b/chroma/cuda/geometry.h
index 577da72..0735792 100644
--- a/chroma/cuda/geometry.h
+++ b/chroma/cuda/geometry.h
@@ -4,9 +4,6 @@
 #include "geometry_types.h"
 #include "linalg.h"
 
-const unsigned int LEAF_BIT = (1U << 31);
-
-
 __device__ float3 
 to_float3(const uint3 &a)
 {
@@ -39,8 +36,9 @@ get_node(Geometry *geometry, const unsigned int &i)
 
     node_struct.lower = geometry->world_origin + to_float3(lower_int) * geometry->world_scale;
     node_struct.upper = geometry->world_origin + to_float3(upper_int) * geometry->world_scale;
-    node_struct.child = node.w & ~LEAF_BIT; // Mask off leaf bit
-    node_struct.kind = node.w & LEAF_BIT ? LEAF_NODE : INTERNAL_NODE;
+    node_struct.child = node.w & ~NCHILD_MASK;
+    node_struct.nchild = node.w >> CHILD_BITS;
+    node_struct.kind = node_struct.nchild == 0 ? LEAF_NODE : INTERNAL_NODE;
     
     return node_struct;
 }
diff --git a/chroma/cuda/geometry_types.h b/chroma/cuda/geometry_types.h
index 89c53bc..b2b5947 100644
--- a/chroma/cuda/geometry_types.h
+++ b/chroma/cuda/geometry_types.h
@@ -28,12 +28,15 @@ struct Triangle
 };
 
 enum { INTERNAL_NODE, LEAF_NODE, PADDING_NODE };
+const unsigned int CHILD_BITS = 26;
+const unsigned int NCHILD_MASK = (0xFFFFu << CHILD_BITS);
 
 struct Node
 {
     float3 lower;
     float3 upper;
     unsigned int child;
+    unsigned int nchild;
     unsigned int kind;
 };
 
@@ -48,7 +51,6 @@ struct Geometry
     Surface **surfaces;
     float3 world_origin;
     float world_scale;
-    unsigned int branch_degree;
 };
 
 #endif
diff --git a/chroma/cuda/mesh.h b/chroma/cuda/mesh.h
index 70ffc23..c52a241 100644
--- a/chroma/cuda/mesh.h
+++ b/chroma/cuda/mesh.h
@@ -56,7 +56,9 @@ intersect_mesh(const float3 &origin, const float3& direction, Geometry *g,
 	return -1;
 
     unsigned int child_ptr_stack[STACK_SIZE];
+    unsigned int nchild_ptr_stack[STACK_SIZE];
     child_ptr_stack[0] = root.child;
+    nchild_ptr_stack[0] = root.nchild;
 
     int curr = 0;
 
@@ -65,9 +67,10 @@ intersect_mesh(const float3 &origin, const float3& direction, Geometry *g,
 
     while (curr >= 0) {
 	unsigned int first_child = child_ptr_stack[curr];
+	unsigned int nchild = nchild_ptr_stack[curr];
 	curr--;
 
-	for (unsigned int i=first_child; i < first_child + g->branch_degree; i++) {
+	for (unsigned int i=first_child; i < first_child + nchild; i++) {
 	    Node node = get_node(g, i);
 	    count++;
 
@@ -97,6 +100,7 @@ intersect_mesh(const float3 &origin, const float3& direction, Geometry *g,
 		} else {
 		    curr++;
 		    child_ptr_stack[curr] = node.child;
+		    nchild_ptr_stack[curr] = node.nchild;
 		} // leaf or internal node?
 	    } // hit node?
 	    
diff --git a/chroma/cuda/render.cu b/chroma/cuda/render.cu
index bbb7383..eda9028 100644
--- a/chroma/cuda/render.cu
+++ b/chroma/cuda/render.cu
@@ -70,7 +70,9 @@ render(int nthreads, float3 *_origin, float3 *_direction, Geometry *g,
     }
 
     unsigned int child_ptr_stack[STACK_SIZE];
+    unsigned int nchild_ptr_stack[STACK_SIZE];
     child_ptr_stack[0] = root.child;
+    nchild_ptr_stack[0] = root.nchild;
 
     int curr = 0;
 
@@ -82,9 +84,10 @@ render(int nthreads, float3 *_origin, float3 *_direction, Geometry *g,
 
     while (curr >= 0) {
 	unsigned int first_child = child_ptr_stack[curr];
+	unsigned int nchild = child_ptr_stack[curr];
 	curr--;
 
-	for (unsigned int i=first_child; i < first_child + g->branch_degree; i++) {
+	for (unsigned int i=first_child; i < first_child + nchild; i++) {
 	    Node node = get_node(g, i);
 	    count++;
 
@@ -128,6 +131,7 @@ render(int nthreads, float3 *_origin, float3 *_direction, Geometry *g,
 	      } else {
 		curr++;
 		child_ptr_stack[curr] = node.child;
+		nchild_ptr_stack[curr] = node.nchild;
 	      } // leaf or internal node?
 	    } // hit node?