Refactor the saving and loading of packed BVH nodes to fully abstract

away the split storage.

Also fixes a bug discovered by Mike Jewell that crashed BVH creation after
the last commit.

2 files changed, 21 insertions, 10 deletions

diff --git a/chroma/cuda/bvh.cu b/chroma/cuda/bvh.cu
index 053eb2b..6ee8404 100644
--- a/chroma/cuda/bvh.cu
+++ b/chroma/cuda/bvh.cu
@@ -548,14 +548,12 @@ extern "C"
      unsigned int thread_id = blockDim.x * blockIdx.x + threadIdx.x;
      unsigned int stride = gridDim.x * blockDim.x;
 
-     uint4 *node = geometry->nodes;
-
      const int MAX_CHILD = 1 << (32 - CHILD_BITS);
      float distance[MAX_CHILD];
      uint4 children[MAX_CHILD];
 
      for (unsigned int i=start+thread_id; i < end; i += stride) {
-       uint4 this_node = node[i];
+       uint4 this_node = get_packed_node(geometry, i);
        unsigned int nchild = this_node.w >> CHILD_BITS;
        unsigned int child_id = this_node.w &  ~NCHILD_MASK;
 
@@ -563,7 +561,7 @@ extern "C"
 	 continue;
 
        for (unsigned int i=0; i < nchild; i++) {
-	 children[i] = node[child_id + i];
+	 children[i] = get_packed_node(geometry, child_id+i);
 	 Node unpacked = get_node(geometry, child_id+i);
 	 float3 delta = unpacked.upper - unpacked.lower;
 	 distance[i] = -(delta.x * delta.y + delta.y * delta.z + delta.z * delta.x);
@@ -572,7 +570,7 @@ extern "C"
        piksrt2(nchild, distance, children);
 
        for (unsigned int i=0; i < nchild; i++)
-	 node[child_id + i] = children[i];
+	 put_packed_node(geometry, child_id + i, children[i]);
      }
   }
 
diff --git a/chroma/cuda/geometry.h b/chroma/cuda/geometry.h
index 4d21779..6a930e7 100644
--- a/chroma/cuda/geometry.h
+++ b/chroma/cuda/geometry.h
@@ -19,14 +19,27 @@ __device__ uint4 read_skip_l1(uint4 *ptr)
   return val;
 }
 
-__device__ Node
-get_node(Geometry *geometry, const unsigned int &i)
+__device__ uint4
+get_packed_node(Geometry *geometry, const unsigned int &i)
+{
+    if (i < geometry->nprimary_nodes)
+	return geometry->primary_nodes[i];
+    else
+	return geometry->extra_nodes[i - geometry->nprimary_nodes];
+}
+__device__ void
+put_packed_node(Geometry *geometry, const unsigned int &i, const uint4 &node)
 {
-    uint4 node; 
     if (i < geometry->nprimary_nodes)
-	node = geometry->primary_nodes[i];
+	geometry->primary_nodes[i] = node;
     else
-	node = geometry->extra_nodes[i - geometry->nprimary_nodes];
+        geometry->extra_nodes[i - geometry->nprimary_nodes] = node;
+}
+
+__device__ Node
+get_node(Geometry *geometry, const unsigned int &i)
+{
+    uint4 node = get_packed_node(geometry, i); 
 	
     Node node_struct;