From 81bcdf9264935195988e171d1339d9a47df139f5 Mon Sep 17 00:00:00 2001 From: Stan Seibert 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