From 81bcdf9264935195988e171d1339d9a47df139f5 Mon Sep 17 00:00:00 2001
From: Stan Seibert <stan@mtrr.org>
Date: Wed, 15 Feb 2012 11:07:52 -0500
Subject: Bugfixes to BVH traversal and generation code.

---
 chroma/cuda/bvh.cu    | 5 ++++-
 chroma/cuda/render.cu | 2 +-
 2 files changed, 5 insertions(+), 2 deletions(-)

diff --git a/chroma/cuda/bvh.cu b/chroma/cuda/bvh.cu
index 5b40bf7..8bc5f04 100644
--- a/chroma/cuda/bvh.cu
+++ b/chroma/cuda/bvh.cu
@@ -288,6 +288,9 @@ extern "C"
     // Scan remaining children
     unsigned int real_children = 1;
     for (unsigned int i=1; i < n_children_per_node; i++) {
+      if (first_child + i >= num_children)
+	break;
+
       uint4 child_node = child_nodes[first_child + i];
       
       if (child_node.x == 0)
@@ -330,7 +333,7 @@ extern "C"
     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;    
+    dest_nodes[node_id] = src_node;
   }
 
   __global__ void distance_to_prev(unsigned int first_node, unsigned int threads_this_round,
diff --git a/chroma/cuda/render.cu b/chroma/cuda/render.cu
index eda9028..6c282d6 100644
--- a/chroma/cuda/render.cu
+++ b/chroma/cuda/render.cu
@@ -84,7 +84,7 @@ 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];
+	unsigned int nchild = nchild_ptr_stack[curr];
 	curr--;
 
 	for (unsigned int i=first_child; i < first_child + nchild; i++) {
-- 
cgit