diff options
| -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++) { |