From f26716f810b744e78782854405d4ea730aac5835 Mon Sep 17 00:00:00 2001 From: helloJetBase-tech <178346048+marktech0813@users.noreply.github.com> Date: Wed, 5 Nov 2025 11:39:22 +0200 Subject: [PATCH] Fix--cdpQuadtree does not work correctly on a large number of points #389 Previously, children and child_offset were derived using global IDs which caused the base pointer to overshoot by multiples of 12 times the parent index at deeper levels, leading to out-of-bounds writes when setting child node metadata and ranges. The new logic makes children the base of the next-level slice for the current 4-node group and uses the local child index (node.id() & 3) to select the correct 4-node sub-slice. This keeps all accesses within the proper subtree memory. --- Samples/3_CUDA_Features/cdpQuadtree/cdpQuadtree.cu | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/Samples/3_CUDA_Features/cdpQuadtree/cdpQuadtree.cu b/Samples/3_CUDA_Features/cdpQuadtree/cdpQuadtree.cu index 60112fdae..c96f2dd94 100644 --- a/Samples/3_CUDA_Features/cdpQuadtree/cdpQuadtree.cu +++ b/Samples/3_CUDA_Features/cdpQuadtree/cdpQuadtree.cu @@ -502,11 +502,12 @@ __global__ void build_quadtree_kernel(Quadtree_node *nodes, Points *points, Para if (!(params.depth >= params.max_depth || num_points <= params.min_points_per_node)) { // The last thread launches new blocks. if (threadIdx.x == NUM_THREADS_PER_BLOCK - 1) { - // The children. - Quadtree_node *children = &nodes[params.num_nodes_at_this_level - (node.id() & ~3)]; + // The children. Move to the next-level slice relative to this 4-node group, + // and select the 4-children group of this node (local index within its parent). + Quadtree_node *children = &nodes[params.num_nodes_at_this_level]; - // The offsets of the children at their level. - int child_offset = 4 * node.id(); + // The offset of this node's 4 children within the next-level slice (local to group). + int child_offset = 4 * (node.id() & 3); // Set IDs. children[child_offset + 0].set_id(4 * node.id() + 0);