summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorStan Seibert <stan@mtrr.org>2012-02-15 11:07:52 -0500
committertlatorre <tlatorre@uchicago.edu>2021-05-09 08:42:38 -0700
commit81bcdf9264935195988e171d1339d9a47df139f5 (patch)
tree08a516ff8ebd86f25be17340d87bcf599e837d00
parent9e16751fcd575c268fd1df03c31b762475c5b480 (diff)
downloadchroma-81bcdf9264935195988e171d1339d9a47df139f5.tar.gz
chroma-81bcdf9264935195988e171d1339d9a47df139f5.tar.bz2
chroma-81bcdf9264935195988e171d1339d9a47df139f5.zip
Bugfixes to BVH traversal and generation code.
-rw-r--r--chroma/cuda/bvh.cu5
-rw-r--r--chroma/cuda/render.cu2
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++) {