diff options
| author | Stan Seibert <stan@mtrr.org> | 2012-02-15 11:07:52 -0500 |
|---|---|---|
| committer | tlatorre <tlatorre@uchicago.edu> | 2021-05-09 08:42:38 -0700 |
| commit | 81bcdf9264935195988e171d1339d9a47df139f5 (patch) | |
| tree | 08a516ff8ebd86f25be17340d87bcf599e837d00 | |
| parent | 9e16751fcd575c268fd1df03c31b762475c5b480 (diff) | |
| download | chroma-81bcdf9264935195988e171d1339d9a47df139f5.tar.gz chroma-81bcdf9264935195988e171d1339d9a47df139f5.tar.bz2 chroma-81bcdf9264935195988e171d1339d9a47df139f5.zip | |
Bugfixes to BVH traversal and generation code.
| -rw-r--r-- | chroma/cuda/bvh.cu | 5 | ||||
| -rw-r--r-- | 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++) { |