From 1a8cb909e6e1a6a4dd0e80a68f628703e2361a8f Mon Sep 17 00:00:00 2001 From: Ingo Wald Date: Thu, 19 Feb 2026 10:37:57 -0700 Subject: [PATCH 1/6] starting on nbody sample --- cuBQL/traversal/nBodyStyle.h | 0 samples/CMakeLists.txt | 13 +++ samples/s07_aggregateNBody/aggregateNBody.cu | 100 +++++++++++++++++++ 3 files changed, 113 insertions(+) create mode 100644 cuBQL/traversal/nBodyStyle.h create mode 100644 samples/s07_aggregateNBody/aggregateNBody.cu diff --git a/cuBQL/traversal/nBodyStyle.h b/cuBQL/traversal/nBodyStyle.h new file mode 100644 index 0000000..e69de29 diff --git a/samples/CMakeLists.txt b/samples/CMakeLists.txt index 5b23713..e478516 100644 --- a/samples/CMakeLists.txt +++ b/samples/CMakeLists.txt @@ -112,6 +112,19 @@ if (CUBQL_HAVE_CUDA) ) endif() +if (CUBQL_HAVE_CUDA) + add_executable(sample07_aggregateNBody + s07_aggregateNBody/aggregateNBody.cu + ) + target_link_libraries(sample07_aggregateNBody + # the cuda-side builders for float3 data + cuBQL_cuda_float3 + # common samples stuff + cuBQL_samples_common + stb_image + ) +endif() + diff --git a/samples/s07_aggregateNBody/aggregateNBody.cu b/samples/s07_aggregateNBody/aggregateNBody.cu new file mode 100644 index 0000000..f54146a --- /dev/null +++ b/samples/s07_aggregateNBody/aggregateNBody.cu @@ -0,0 +1,100 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA +// CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +/*! \file closestPointGPU.cu Implements a small demo-app that + generates a set of data points, another set of query points, and + then uses cuBQL to perform closest-point qeuries (ie, it finds, + for each query point, the respectively closest data point */ + +// cuBQL itself, and the BVH type(s) it defines +#include "cuBQL/bvh.h" +#include "cuBQL/builder/cuda.h" +#include "cuBQL/builder/cuda/aggregate_refit.h" +// some specialized query kernels for find-closest, on 'points' data +#include "cuBQL/queries/pointData/findClosest.h" +// helper class to generate various data distributions +#include "samples/common/Generator.h" + +using namespace cuBQL; + +__global__ +void computeBoxes(box3f *d_boxes, const vec3f *d_data, int numData) +{ + int tid = threadIdx.x+blockIdx.x*blockDim.x; + if (tid >= numData) return; + + d_boxes[tid] = box3f().including(d_data[tid]); +} + +__global__ +void runQueries(bvh3f bvh, + const vec3f *d_data, + const vec3f *d_queries, + int numQueries) +{ + int tid = threadIdx.x+blockIdx.x*blockDim.x; + if (tid >= numQueries) return; + + vec3f queryPoint = d_queries[tid]; + int closestID = cuBQL::points::findClosest(/* the cubql bvh we've built */ + bvh, + /* data that this bvh was built over*/ + d_data, + queryPoint); + vec3f closestPoint = d_data[closestID]; + printf("[%i] closest point to (%f %f %f) is point #%i, at (%f %f %f)\n", + tid, + queryPoint.x, + queryPoint.y, + queryPoint.z, + closestID, + closestPoint.x, + closestPoint.y, + closestPoint.z); +} + + +int main(int, char **) +{ + int numDataPoints = 10000; + int numQueryPoints = 20; + /*! generate 10,000 uniformly distributed data points */ + std::vector dataPoints + = cuBQL::samples::convert + (cuBQL::samples::UniformPointGenerator<3>() + .generate(numDataPoints,290374)); + std::cout << "#cubql: generated " << dataPoints.size() + << " data points" << std::endl; + std::vector queryPoints + = cuBQL::samples::convert + (cuBQL::samples::UniformPointGenerator<3>() + .generate(numQueryPoints,/*seed*/1234567)); + std::cout << "#cubql: generated " << queryPoints.size() + << " query points" << std::endl; + + vec3f *d_queryPoints = 0; + vec3f *d_dataPoints = 0; + box3f *d_primBounds = 0; + CUBQL_CUDA_CALL(Malloc((void **)&d_queryPoints,queryPoints.size()*sizeof(vec3f))); + CUBQL_CUDA_CALL(Memcpy(d_queryPoints,queryPoints.data(), + queryPoints.size()*sizeof(queryPoints[0]), + cudaMemcpyDefault)); + CUBQL_CUDA_CALL(Malloc((void **)&d_dataPoints,dataPoints.size()*sizeof(vec3f))); + CUBQL_CUDA_CALL(Memcpy(d_dataPoints,dataPoints.data(), + dataPoints.size()*sizeof(dataPoints[0]), + cudaMemcpyDefault)); + CUBQL_CUDA_CALL(Malloc((void **)&d_primBounds,dataPoints.size()*sizeof(box3f))); + computeBoxes<<>> + (d_primBounds,d_dataPoints,numDataPoints); + + // generate cuBQL bvh + bvh3f bvh; + cuBQL::gpuBuilder(bvh,d_primBounds,numDataPoints,BuildConfig()); + runQueries<<>> + (bvh,d_dataPoints,d_queryPoints,numQueryPoints); + + CUBQL_CUDA_SYNC_CHECK(); + return 0; +} + From 176777049d59ca8e734a1cb9c181180de9440c23 Mon Sep 17 00:00:00 2001 From: Ingo Wald Date: Thu, 19 Feb 2026 10:59:16 -0700 Subject: [PATCH 2/6] separated out refit into its own header files and cuda:: kernel --- cuBQL/builder/cuda.h | 16 +++ cuBQL/builder/cuda/aggregate_refit.h | 13 +++ cuBQL/builder/cuda/builder_common.h | 6 ++ cuBQL/builder/cuda/gpu_builder.h | 2 +- cuBQL/builder/cuda/profiling_helper.h | 55 ++++++++++ cuBQL/builder/cuda/radix.h | 2 +- cuBQL/builder/cuda/rebinMortonBuilder.h | 24 +++-- cuBQL/builder/cuda/refit.h | 92 ++++++++++++++++ cuBQL/builder/cuda/sah_builder.h | 2 +- cuBQL/builder/cuda/sm_builder.h | 133 +----------------------- 10 files changed, 202 insertions(+), 143 deletions(-) create mode 100644 cuBQL/builder/cuda/aggregate_refit.h create mode 100644 cuBQL/builder/cuda/profiling_helper.h create mode 100644 cuBQL/builder/cuda/refit.h diff --git a/cuBQL/builder/cuda.h b/cuBQL/builder/cuda.h index 6556da2..0b5fbe6 100644 --- a/cuBQL/builder/cuda.h +++ b/cuBQL/builder/cuda.h @@ -186,16 +186,32 @@ namespace cuBQL { cudaStream_t s=0, GpuMemoryResource &memResource=defaultGpuMemResource()); + // ------------------------------------------------------------------ + /*! refit a previously built boxes to a new set of bounding + boxes. The order of boxes in the array boxes[] has to + correspond to that used when building the tree. */ + // ------------------------------------------------------------------ + template + void refit(BinaryBVH &bvh, + const box_t *boxes, + cudaStream_t s=0, + GpuMemoryResource &memResource=defaultGpuMemResource()); + + // ------------------------------------------------------------------ /*! frees the bvh.nodes[] and bvh.primIDs[] memory allocated when building the BVH. this assumes that the 'memResource' provided here was the same that was used during building */ + // ------------------------------------------------------------------ template void free(BinaryBVH &bvh, cudaStream_t s=0, GpuMemoryResource& memResource=defaultGpuMemResource()); + + // ------------------------------------------------------------------ /*! frees the bvh.nodes[] and bvh.primIDs[] memory allocated when building the BVH. this assumes that the 'memResource' provided here was the same that was used during building */ + // ------------------------------------------------------------------ template void free(WideBVH &bvh, cudaStream_t s=0, diff --git a/cuBQL/builder/cuda/aggregate_refit.h b/cuBQL/builder/cuda/aggregate_refit.h new file mode 100644 index 0000000..557d1b2 --- /dev/null +++ b/cuBQL/builder/cuda/aggregate_refit.h @@ -0,0 +1,13 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA +// CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "cuBQL/builder/cuda/builder_common.h" + +namespace cuBQL { + namespace cuda { + + } +} diff --git a/cuBQL/builder/cuda/builder_common.h b/cuBQL/builder/cuda/builder_common.h index c11b9ed..47e0ea4 100644 --- a/cuBQL/builder/cuda/builder_common.h +++ b/cuBQL/builder/cuda/builder_common.h @@ -14,6 +14,12 @@ #include #include +#ifdef __HIPCC__ +namespace cub { + using namespace hipcub; +} +#endif + namespace cuBQL { namespace gpuBuilder_impl { diff --git a/cuBQL/builder/cuda/gpu_builder.h b/cuBQL/builder/cuda/gpu_builder.h index 6e1e45c..37aee9f 100644 --- a/cuBQL/builder/cuda/gpu_builder.h +++ b/cuBQL/builder/cuda/gpu_builder.h @@ -53,7 +53,7 @@ namespace cuBQL { buildConfig.makeLeafThreshold = 1; gpuBuilder_impl::build(bvh,boxes,numBoxes,buildConfig,s,memResource); } - gpuBuilder_impl::refit(bvh,boxes,s,memResource); + cuBQL::cuda::refit(bvh,boxes,s,memResource); } namespace cuda { diff --git a/cuBQL/builder/cuda/profiling_helper.h b/cuBQL/builder/cuda/profiling_helper.h new file mode 100644 index 0000000..b3a8747 --- /dev/null +++ b/cuBQL/builder/cuda/profiling_helper.h @@ -0,0 +1,55 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA +// CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +namespace cuBQL { + namespace gpuBuilder_impl { + + //#define CUBQL_PROFILE 1 + +#if CUBQL_PROFILE + struct Profile { + void setName(std::string name, int sub=-1) + { + if (sub >= 0) { + char suff[1000]; + sprintf(suff,"[%2i]",sub); + this->name = name+suff; + } else + this->name = name; + } + ~Profile() { ping(); } + + void start() { + t0 = getCurrentTime(); + } + void sync_start() { + CUBQL_CUDA_SYNC_CHECK(); + start(); + } + void sync_stop() { + CUBQL_CUDA_SYNC_CHECK(); + stop(); + } + void stop(bool do_ping = false) { + double t1 = getCurrentTime(); + t_sum += (t1-t0); + count ++; + if (do_ping) ping(); + } + void ping() + { + if (count) + std::cout << "#PROF " << name << " = " << prettyDouble(t_sum / count) << std::endl; + } + double t0 = 0.; + double t_sum = 0.; + int count = 0; + std::string name = ""; + }; +#endif + + } +} diff --git a/cuBQL/builder/cuda/radix.h b/cuBQL/builder/cuda/radix.h index 1475577..825af91 100644 --- a/cuBQL/builder/cuda/radix.h +++ b/cuBQL/builder/cuda/radix.h @@ -742,7 +742,7 @@ namespace cuBQL { // ================================================================== // done. all we need to do now is refit the bboxes // ================================================================== - gpuBuilder_impl::refit(bvh,boxes,s,memResource); + cuBQL::cuda::refit(bvh,boxes,s,memResource); } } diff --git a/cuBQL/builder/cuda/rebinMortonBuilder.h b/cuBQL/builder/cuda/rebinMortonBuilder.h index 877f7bc..00fc6c2 100644 --- a/cuBQL/builder/cuda/rebinMortonBuilder.h +++ b/cuBQL/builder/cuda/rebinMortonBuilder.h @@ -1459,20 +1459,24 @@ namespace cuBQL { // ================================================================== // done. all we need to do now is refit the bboxes // ================================================================== - gpuBuilder_impl::refit(bvh,boxes,s,memResource); + cuBQL::cuda::refit(bvh,boxes,s,memResource); } } - + namespace cuda { template void rebinRadixBuilder(BinaryBVH &bvh, - const box_t *boxes, - uint32_t numPrims, - BuildConfig buildConfig, - cudaStream_t s, - GpuMemoryResource &memResource) - { rebinRadixBuilder_impl::build(bvh,boxes,numPrims,buildConfig,s,memResource); } - } -} + const box_t *boxes, + uint32_t numPrims, + BuildConfig buildConfig, + cudaStream_t s, + GpuMemoryResource &memResource) + { + rebinRadixBuilder_impl::build + (bvh,boxes,numPrims,buildConfig,s,memResource); + } + + } // ::cuBQL::cuda +} // ::cuBQL #endif diff --git a/cuBQL/builder/cuda/refit.h b/cuBQL/builder/cuda/refit.h new file mode 100644 index 0000000..457b092 --- /dev/null +++ b/cuBQL/builder/cuda/refit.h @@ -0,0 +1,92 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA +// CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "cuBQL/builder/cuda/builder_common.h" + +namespace cuBQL { + namespace cuda { + + template + __global__ void + refit_init(const typename BinaryBVH::Node *nodes, + uint32_t *refitData, + int numNodes) + { + const int nodeID = threadIdx.x+blockIdx.x*blockDim.x; + if (nodeID == 1 || nodeID >= numNodes) return; + if (nodeID < 2) + refitData[0] = 0; + const auto &node = nodes[nodeID]; + if (node.admin.count) return; + + refitData[node.admin.offset+0] = nodeID << 1; + refitData[node.admin.offset+1] = nodeID << 1; + } + + template + __global__ + void refit_run(BinaryBVH bvh, + uint32_t *refitData, + const box_t *boxes) + { + int nodeID = threadIdx.x+blockIdx.x*blockDim.x; + if (nodeID == 1 || nodeID >= bvh.numNodes) return; + + typename BinaryBVH::Node *node = &bvh.nodes[nodeID]; + if (node->admin.count == 0) + // this is a inner node - exit + return; + + box_t bounds; bounds.set_empty(); + for (int i=0;iadmin.count;i++) { + const box_t primBox = boxes[bvh.primIDs[node->admin.offset+i]]; + bounds.lower = min(bounds.lower,primBox.lower); + bounds.upper = max(bounds.upper,primBox.upper); + } + + int parentID = (refitData[nodeID] >> 1); + while (true) { + node->bounds = bounds; + __threadfence(); + if (node == bvh.nodes) + break; + + uint32_t refitBits = atomicAdd(&refitData[parentID],1u); + if ((refitBits & 1) == 0) + // we're the first one - let other one do it + break; + + nodeID = parentID; + node = &bvh.nodes[parentID]; + parentID = (refitBits >> 1); + + typename BinaryBVH::Node l = bvh.nodes[node->admin.offset+0]; + typename BinaryBVH::Node r = bvh.nodes[node->admin.offset+1]; + bounds.lower = min(l.bounds.lower,r.bounds.lower); + bounds.upper = max(l.bounds.upper,r.bounds.upper); + } + } + + template + void refit(BinaryBVH &bvh, + const box_t *boxes, + cudaStream_t s, + GpuMemoryResource &memResource) + { + uint32_t *refitData = 0; + memResource.malloc((void**)&refitData,bvh.numNodes*sizeof(int),s); + + int numNodes = bvh.numNodes; + refit_init<<>> + (bvh.nodes,refitData,bvh.numNodes); + refit_run<<>> + (bvh,refitData,boxes); + memResource.free((void*)refitData,s); + // we're not syncing here - let APP do that + } + + } // ::cuBQL::gpuBuilder_impl +} // ::cuBQL diff --git a/cuBQL/builder/cuda/sah_builder.h b/cuBQL/builder/cuda/sah_builder.h index dafb440..5374615 100644 --- a/cuBQL/builder/cuda/sah_builder.h +++ b/cuBQL/builder/cuda/sah_builder.h @@ -542,7 +542,7 @@ namespace cuBQL { _FREE(buildState,s,memResource); _FREE(sahBins,s,memResource); - gpuBuilder_impl::refit(bvh,boxes,s,memResource); + cuBQL::cuda::refit(bvh,boxes,s,memResource); } template<> diff --git a/cuBQL/builder/cuda/sm_builder.h b/cuBQL/builder/cuda/sm_builder.h index 2442048..59395cd 100644 --- a/cuBQL/builder/cuda/sm_builder.h +++ b/cuBQL/builder/cuda/sm_builder.h @@ -1,63 +1,15 @@ -// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA +// CORPORATION & AFFILIATES. All rights reserved. // SPDX-License-Identifier: Apache-2.0 #pragma once #include "cuBQL/builder/cuda/builder_common.h" - -#ifdef __HIPCC__ -namespace cub { - using namespace hipcub; -} -#endif +#include "cuBQL/builder/cuda/refit.h" namespace cuBQL { namespace gpuBuilder_impl { - //#define CUBQL_PROFILE 1 - -#if CUBQL_PROFILE - struct Profile { - void setName(std::string name, int sub=-1) - { - if (sub >= 0) { - char suff[1000]; - sprintf(suff,"[%2i]",sub); - this->name = name+suff; - } else - this->name = name; - } - ~Profile() { ping(); } - - void start() { - t0 = getCurrentTime(); - } - void sync_start() { - CUBQL_CUDA_SYNC_CHECK(); - start(); - } - void sync_stop() { - CUBQL_CUDA_SYNC_CHECK(); - stop(); - } - void stop(bool do_ping = false) { - double t1 = getCurrentTime(); - t_sum += (t1-t0); - count ++; - if (do_ping) ping(); - } - void ping() - { - if (count) - std::cout << "#PROF " << name << " = " << prettyDouble(t_sum / count) << std::endl; - } - double t0 = 0.; - double t_sum = 0.; - int count = 0; - std::string name = ""; - }; -#endif - struct PrimState { union { /* careful with this order - this is intentionally chosen such @@ -609,85 +561,6 @@ namespace cuBQL { _FREE(buildState,s,memResource); } - template - __global__ void - refit_init(const typename BinaryBVH::Node *nodes, - uint32_t *refitData, - int numNodes) - { - const int nodeID = threadIdx.x+blockIdx.x*blockDim.x; - if (nodeID == 1 || nodeID >= numNodes) return; - if (nodeID < 2) - refitData[0] = 0; - const auto &node = nodes[nodeID]; - if (node.admin.count) return; - - refitData[node.admin.offset+0] = nodeID << 1; - refitData[node.admin.offset+1] = nodeID << 1; - } - - template - __global__ - void refit_run(BinaryBVH bvh, - uint32_t *refitData, - const box_t *boxes) - { - int nodeID = threadIdx.x+blockIdx.x*blockDim.x; - if (nodeID == 1 || nodeID >= bvh.numNodes) return; - - typename BinaryBVH::Node *node = &bvh.nodes[nodeID]; - if (node->admin.count == 0) - // this is a inner node - exit - return; - - box_t bounds; bounds.set_empty(); - for (int i=0;iadmin.count;i++) { - const box_t primBox = boxes[bvh.primIDs[node->admin.offset+i]]; - bounds.lower = min(bounds.lower,primBox.lower); - bounds.upper = max(bounds.upper,primBox.upper); - } - - int parentID = (refitData[nodeID] >> 1); - while (true) { - node->bounds = bounds; - __threadfence(); - if (node == bvh.nodes) - break; - - uint32_t refitBits = atomicAdd(&refitData[parentID],1u); - if ((refitBits & 1) == 0) - // we're the first one - let other one do it - break; - - nodeID = parentID; - node = &bvh.nodes[parentID]; - parentID = (refitBits >> 1); - - typename BinaryBVH::Node l = bvh.nodes[node->admin.offset+0]; - typename BinaryBVH::Node r = bvh.nodes[node->admin.offset+1]; - bounds.lower = min(l.bounds.lower,r.bounds.lower); - bounds.upper = max(l.bounds.upper,r.bounds.upper); - } - } - - template - void refit(BinaryBVH &bvh, - const box_t *boxes, - cudaStream_t s=0, - GpuMemoryResource &memResource=defaultGpuMemResource()) - { - uint32_t *refitData = 0; - memResource.malloc((void**)&refitData,bvh.numNodes*sizeof(int),s); - - int numNodes = bvh.numNodes; - refit_init<<>> - (bvh.nodes,refitData,bvh.numNodes); - refit_run<<>> - (bvh,refitData,boxes); - memResource.free((void*)refitData,s); - // we're not syncing here - let APP do that - } - } // ::cuBQL::gpuBuilder_impl } // ::cuBQL From bbeb306aca03630693ff7040efefee9092c385d8 Mon Sep 17 00:00:00 2001 From: Ingo Wald Date: Sat, 21 Feb 2026 12:45:43 -0700 Subject: [PATCH 3/6] first framework for aggreagete refit and query --- cuBQL/builder/cuda/aggregate_refit.h | 13 -- cuBQL/builder/cuda/refit_aggregate.h | 22 +++ cuBQL/traversal/aggregateApproximate.h | 81 +++++++++ cuBQL/traversal/nBodyStyle.h | 0 samples/s07_aggregateNBody/aggregateNBody.cu | 165 ++++++++++++++++--- 5 files changed, 245 insertions(+), 36 deletions(-) delete mode 100644 cuBQL/builder/cuda/aggregate_refit.h create mode 100644 cuBQL/builder/cuda/refit_aggregate.h create mode 100644 cuBQL/traversal/aggregateApproximate.h delete mode 100644 cuBQL/traversal/nBodyStyle.h diff --git a/cuBQL/builder/cuda/aggregate_refit.h b/cuBQL/builder/cuda/aggregate_refit.h deleted file mode 100644 index 557d1b2..0000000 --- a/cuBQL/builder/cuda/aggregate_refit.h +++ /dev/null @@ -1,13 +0,0 @@ -// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA -// CORPORATION & AFFILIATES. All rights reserved. -// SPDX-License-Identifier: Apache-2.0 - -#pragma once - -#include "cuBQL/builder/cuda/builder_common.h" - -namespace cuBQL { - namespace cuda { - - } -} diff --git a/cuBQL/builder/cuda/refit_aggregate.h b/cuBQL/builder/cuda/refit_aggregate.h new file mode 100644 index 0000000..0e7c03a --- /dev/null +++ b/cuBQL/builder/cuda/refit_aggregate.h @@ -0,0 +1,22 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA +// CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "cuBQL/builder/cuda/builder_common.h" + +namespace cuBQL { + namespace cuda { + + template< + typename T, + int D, + typename AggregateNodeData, + typename AggregateFct> + void refit_aggregate(bvh_t bvh, + AggregateNodeData *d_aggregateNodeData, + const AggregateFct &aggregateFct); + + } +} diff --git a/cuBQL/traversal/aggregateApproximate.h b/cuBQL/traversal/aggregateApproximate.h new file mode 100644 index 0000000..01fdc6e --- /dev/null +++ b/cuBQL/traversal/aggregateApproximate.h @@ -0,0 +1,81 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "cuBQL/bvh.h" +#include +#include +#include +#include + +/* Defines and implements cuBQL "approximate/aggregate" style + traversals that can, for example, be used for N-body style + problems. + + The core idea of these types of queries is that the user provides + three things: + + - one, some per-subtree 'aggregate data' (of the user's choosing, + and computed, for example, via refit_aggregate()). For an n-body + style problem this could, for example, be the sum of all + planets/bodies/masses in a subtree. + + - second, a callback function that checks if a given query can be + approximately fulfilled with the subtree's aggregate data; i.e., + _without_ having to traverse that subtree's children. If so, this + helper function can accumulate this partial result (in whatever + way it chooses - it's user code, after all), and returns 'true' + to tell cuBQL that this subtree is 'done' and does not require + further processing. Otherwise, it returns 'false' and cuBQL will + process the children + + - third, a second callback function that operates on individual + primitmives, and gets called by cuBQL if traversal reaches a leaf + without ever having decided to approximate in any of that child + dnoe's parent nodes + + Obviously both callback functions need additional data to do their + job: the bvh to be traversed (eg to get a node's bounding box), the + (tempalted) aggregate data (obviously), the (templated) query_t for + which the query is performed, and some (templated) result_t in + which both callbacks can accumulate their partial results (eg for + an n-body style, this could be the sum of all forces) + + Note that "approximate/aggregate" refers to the two key concepts + required to realize these kind of traversals: the idea to avoid a + "full" tree traversal by "approimating" certain subtrees (instead + of just traversing both children); and the idea that one needs some + sort of "aggregate data" for a subtree to even decide whether + that's possible or not. +*/ +namespace cuBQL { + namespace aggregateApproximate { + + /*! implements a approximate/aggregate traversal (see above for + the core idea). Note this function is heavily templated, so to + allow template matchign to do its magic the order of + parameters is pretty important. + + agg + */ + template + + inline __device__ + void traverse(bvh_t bvh, + aggregateNodeData_t aggregateData[], + primitive_t primitives[], + result_t result, + query_t queryPoint, + const approximateSubtreeFct_t &approximateSubtreeFct, + const perPrimFct_t perPrimFct); + } +} + diff --git a/cuBQL/traversal/nBodyStyle.h b/cuBQL/traversal/nBodyStyle.h deleted file mode 100644 index e69de29..0000000 diff --git a/samples/s07_aggregateNBody/aggregateNBody.cu b/samples/s07_aggregateNBody/aggregateNBody.cu index f54146a..a1b64b9 100644 --- a/samples/s07_aggregateNBody/aggregateNBody.cu +++ b/samples/s07_aggregateNBody/aggregateNBody.cu @@ -2,20 +2,23 @@ // CORPORATION & AFFILIATES. All rights reserved. // SPDX-License-Identifier: Apache-2.0 -/*! \file closestPointGPU.cu Implements a small demo-app that - generates a set of data points, another set of query points, and - then uses cuBQL to perform closest-point qeuries (ie, it finds, - for each query point, the respectively closest data point */ +/*! \file aggregateNBodya.cu Provides a mini-sample for how to use the + "aggregate refit" and "nbody-style traversal" concepts, in the + example of a simplified n-body problem */ // cuBQL itself, and the BVH type(s) it defines #include "cuBQL/bvh.h" #include "cuBQL/builder/cuda.h" -#include "cuBQL/builder/cuda/aggregate_refit.h" // some specialized query kernels for find-closest, on 'points' data #include "cuBQL/queries/pointData/findClosest.h" // helper class to generate various data distributions #include "samples/common/Generator.h" +// pull in ability to refit aggregate data onto a BVH: +#include "cuBQL/builder/cuda/refit_aggregate.h" +// pull in traversal that can cull based on that aggregate data: +#include "cuBQL/traversal/aggregateApproximate.h" + using namespace cuBQL; __global__ @@ -27,8 +30,105 @@ void computeBoxes(box3f *d_boxes, const vec3f *d_data, int numData) d_boxes[tid] = box3f().including(d_data[tid]); } +namespace nBody { + /*! the data we want cuBQL to store for us in each subtree. For this + sample we simply store the total number of points in each + subtree. A real n-body code might want to track more or other + data (like sum of all masses, etc) - if so, just change this */ + struct AggregateNodeData { + int numBodiesInSubtree; + }; + + /*! aggregation function that computes a node's aggregate data + during aggragate_refit */ + inline __device__ + void aggregate(int nodeID, + AggregateNodeData nodeAggregates[], + bvh3f bvh) + { + auto node = bvh.nodes[nodeID].admin; + if (node.count != 0) { + // this is a leaf - aggregate data is the leaf count itself + nodeAggregates[nodeID].numBodiesInSubtree = node.count; + } else { + // this is a inner node - aggregate data is the sum of both + // children. note that aggragate_refit() guarantees that both + // children have alrady been aggregated before this gets called + nodeAggregates[nodeID].numBodiesInSubtree + = nodeAggregates[node.offset+0].numBodiesInSubtree + + nodeAggregates[node.offset+1].numBodiesInSubtree; + } + } + + /*! the final result type that results from iterating over the + entire tree. FOr our simple n-body mock-up, it's simply a float + to track sum_i{1/sqrDistance(queryPoint,dataPoint[i])} */ + struct ResultType { + float sumOfForces; + }; + + /*! callback function that checks if a subtree can be approximated; + if so it accumulates the approximated result and returns true + ('yes, i could approximate this subtree); otherwise it returns + false and let's cuBQL traverse to the children */ + inline __device__ + bool approximateSubtree(/* param #1: the actual BVH */ + bvh3f bvh, + /* param #2: the pre-computed per-node + aggregate data */ + AggregateNodeData nodeAggregates[], + /*! param #3: the subtree we're supposed to + evaluate (specified by its node ID) */ + int nodeID, + /*! param #4: the user-supplied struct where + we can accumulate the approximated + subtree's partial result in (if we chose + to do so) */ + ResultType &result, + /*! param #5 the actual query that this + traversal is run on */ + const vec3f &queryPoint) + { + auto node = bvh.nodes[nodeID]; + // first, check if we can approximate this subtree + float sqrDist = sqrDistance(node.bounds,queryPoint); + float sqrDiag = sqrLength(node.bounds.size()); + const float approxThreshold = /* 1% */0.01f; + bool canApproximate + = sqrDist > 0.f + && (sqrDiag / sqrDist <= sqr(approxThreshold)); + + if (canApproximate) { + result.sumOfForces + += nodeAggregates[nodeID].numBodiesInSubtree + * (1.f/sqrDist); + /* yes we DID approximate the subtree (cuBQL doesn't need to + traverse any more*/ + return true; + } + + return false; + } + + inline __device__ + void processPrim(/*! the query's final result type */ + ResultType &result, + /*! the query's actual query point */ + const vec3f &queryPoint, + int primID, + const vec3f prims[]) + { + float sqrDist = sqrLength(prims[primID]-queryPoint); + if (sqrDist != 0.f) + result.sumOfForces += 1.f/sqrDist; + } +} + + __global__ -void runQueries(bvh3f bvh, +void runQueries(float *d_results, + bvh3f bvh, + nBody::AggregateNodeData nodeAggregates[], const vec3f *d_data, const vec3f *d_queries, int numQueries) @@ -37,21 +137,21 @@ void runQueries(bvh3f bvh, if (tid >= numQueries) return; vec3f queryPoint = d_queries[tid]; - int closestID = cuBQL::points::findClosest(/* the cubql bvh we've built */ - bvh, - /* data that this bvh was built over*/ - d_data, - queryPoint); - vec3f closestPoint = d_data[closestID]; - printf("[%i] closest point to (%f %f %f) is point #%i, at (%f %f %f)\n", - tid, - queryPoint.x, - queryPoint.y, - queryPoint.z, - closestID, - closestPoint.x, - closestPoint.y, - closestPoint.z); + nBody::ResultType result = { 0.f }; + /* run a cuBQL "approximate/aggregate" traversal, which requires two + callbacks: one that tries to approximate a subtree (using + pre-aggregated data), and one that processes individual + primitives (if traversal went all the way to leaves */ + cuBQL::aggregateApproximate::traverse + (/* all the per-tree input data: */ + bvh,nodeAggregates,d_data, + /* all the per-query state data */ + result,queryPoint, + /* and the two callbacks */ + nBody::approximateSubtree, + nBody::processPrim + ); + d_results[tid] = result.sumOfForces; } @@ -76,6 +176,8 @@ int main(int, char **) vec3f *d_queryPoints = 0; vec3f *d_dataPoints = 0; box3f *d_primBounds = 0; + float *d_results = 0; + CUBQL_CUDA_CALL(Malloc((void **)&d_results,queryPoints.size()*sizeof(*d_results))); CUBQL_CUDA_CALL(Malloc((void **)&d_queryPoints,queryPoints.size()*sizeof(vec3f))); CUBQL_CUDA_CALL(Memcpy(d_queryPoints,queryPoints.data(), queryPoints.size()*sizeof(queryPoints[0]), @@ -88,11 +190,28 @@ int main(int, char **) computeBoxes<<>> (d_primBounds,d_dataPoints,numDataPoints); - // generate cuBQL bvh + // ------------------------------------------------------------------ + // generate initial cuBQL bvh over the data points + // ------------------------------------------------------------------ bvh3f bvh; cuBQL::gpuBuilder(bvh,d_primBounds,numDataPoints,BuildConfig()); + + // ------------------------------------------------------------------ + // re-fit kernel-specific aggregate data on top of the bvh + // ------------------------------------------------------------------ + nBody::AggregateNodeData *d_nodeAggregates = 0; + CUBQL_CUDA_CALL(Malloc((void **)&d_nodeAggregates, + bvh.numNodes*sizeof(*d_nodeAggregates))); + cuBQL::cuda::refit_aggregate(bvh, + d_nodeAggregates, + nBody::aggregate); + + // ------------------------------------------------------------------ + // ready to run query + // ------------------------------------------------------------------ runQueries<<>> - (bvh,d_dataPoints,d_queryPoints,numQueryPoints); + (d_results,bvh,d_nodeAggregates, + d_dataPoints,d_queryPoints,numQueryPoints); CUBQL_CUDA_SYNC_CHECK(); return 0; From a326fada7b14fc2a6fa51be5fb1f8ae62b0d8f35 Mon Sep 17 00:00:00 2001 From: Ingo Wald Date: Sat, 21 Feb 2026 13:55:54 -0700 Subject: [PATCH 4/6] added n-body traversal and sample --- cuBQL/builder/cuda/refit.h | 7 +- cuBQL/builder/cuda/refit_aggregate.h | 88 ++++++++++++- cuBQL/traversal/aggregateApproximate.h | 130 ++++++++++++++++++- samples/s07_aggregateNBody/aggregateNBody.cu | 4 +- 4 files changed, 216 insertions(+), 13 deletions(-) diff --git a/cuBQL/builder/cuda/refit.h b/cuBQL/builder/cuda/refit.h index 457b092..eb51b6c 100644 --- a/cuBQL/builder/cuda/refit.h +++ b/cuBQL/builder/cuda/refit.h @@ -76,12 +76,13 @@ namespace cuBQL { cudaStream_t s, GpuMemoryResource &memResource) { + int numNodes = bvh.numNodes; + uint32_t *refitData = 0; - memResource.malloc((void**)&refitData,bvh.numNodes*sizeof(int),s); + memResource.malloc((void**)&refitData,numNodes*sizeof(*refitData),s); - int numNodes = bvh.numNodes; refit_init<<>> - (bvh.nodes,refitData,bvh.numNodes); + (bvh.nodes,refitData,numNodes); refit_run<<>> (bvh,refitData,boxes); memResource.free((void*)refitData,s); diff --git a/cuBQL/builder/cuda/refit_aggregate.h b/cuBQL/builder/cuda/refit_aggregate.h index 0e7c03a..c2039d6 100644 --- a/cuBQL/builder/cuda/refit_aggregate.h +++ b/cuBQL/builder/cuda/refit_aggregate.h @@ -5,18 +5,100 @@ #pragma once #include "cuBQL/builder/cuda/builder_common.h" +#include "cuBQL/builder/cuda/refit.h" namespace cuBQL { namespace cuda { + // ------------------------------------------------------------------ + // INTERFACE + // ------------------------------------------------------------------ template< typename T, int D, typename AggregateNodeData, typename AggregateFct> - void refit_aggregate(bvh_t bvh, + void refit_aggregate(BinaryBVH bvh, AggregateNodeData *d_aggregateNodeData, - const AggregateFct &aggregateFct); - + const AggregateFct &aggregateFct, + cudaStream_t s + =0, + GpuMemoryResource &memResource + =defaultGpuMemResource()); + + template + __global__ + void refit_aggregate_run(BinaryBVH bvh, + AggregateNodeData *aggregateNodeData, + const AggregateFct &aggregateFct, + uint32_t *refitData) + { + int nodeID = threadIdx.x+blockIdx.x*blockDim.x; + if (nodeID == 1 || nodeID >= bvh.numNodes) return; + + typename BinaryBVH::Node *node = &bvh.nodes[nodeID]; + if (node->admin.count == 0) + // this is a inner node - exit + return; + + // box_t bounds; bounds.set_empty(); + // for (int i=0;iadmin.count;i++) { + // const box_t primBox = boxes[bvh.primIDs[node->admin.offset+i]]; + // bounds.lower = min(bounds.lower,primBox.lower); + // bounds.upper = max(bounds.upper,primBox.upper); + // } + + int parentID = (refitData[nodeID] >> 1); + while (true) { + aggregateFct(bvh,aggregateNodeData,nodeID); + __threadfence(); + if (node == bvh.nodes) + break; + + uint32_t refitBits = atomicAdd(&refitData[parentID],1u); + if ((refitBits & 1) == 0) + // we're the first one - let other one do it + break; + + nodeID = parentID; + node = &bvh.nodes[parentID]; + parentID = (refitBits >> 1); + + // typename BinaryBVH::Node l = bvh.nodes[node->admin.offset+0]; + // typename BinaryBVH::Node r = bvh.nodes[node->admin.offset+1]; + // bounds.lower = min(l.bounds.lower,r.bounds.lower); + // bounds.upper = max(l.bounds.upper,r.bounds.upper); + } + } + + + + // ------------------------------------------------------------------ + // IMPLEMENTATION + // ------------------------------------------------------------------ + template< + typename T, + int D, + typename AggregateNodeData, + typename AggregateFct> + void refit_aggregate(BinaryBVH bvh, + AggregateNodeData *d_aggregateNodeData, + const AggregateFct &aggregateFct, + cudaStream_t s, + GpuMemoryResource &memResource) + { + int numNodes = bvh.numNodes; + + uint32_t *refitData = 0; + memResource.malloc((void**)&refitData,numNodes*sizeof(*refitData),s); + refit_init<<>> + (bvh.nodes,refitData,numNodes); + refit_aggregate_run<<>> + (bvh,d_aggregateNodeData,aggregateFct,refitData); + memResource.free((void*)refitData,s); + // we're not syncing here - let APP do that + } } } diff --git a/cuBQL/traversal/aggregateApproximate.h b/cuBQL/traversal/aggregateApproximate.h index 01fdc6e..a85e55c 100644 --- a/cuBQL/traversal/aggregateApproximate.h +++ b/cuBQL/traversal/aggregateApproximate.h @@ -52,21 +52,53 @@ namespace cuBQL { namespace aggregateApproximate { + // ------------------------------------------------------------------ + // INTERFACE + // ------------------------------------------------------------------ + /*! implements a approximate/aggregate traversal (see above for the core idea). Note this function is heavily templated, so to allow template matchign to do its magic the order of parameters is pretty important. - agg + `approximateSubtreeFct_t` is a lambda with signature + inline __device__ + bool approximateSubtree(bvh_t, + aggregateNodeData_t [], + int nodeID, + result_t &, + query_t) + + `perPrimFct_t` is a lambda with signature + inline __device__ + void processPrim(result_t &result, + const query_t &queryPoint, + int primID, + const primitive_t prims[]) + */ template - inline __device__ void traverse(bvh_t bvh, @@ -76,6 +108,94 @@ namespace cuBQL { query_t queryPoint, const approximateSubtreeFct_t &approximateSubtreeFct, const perPrimFct_t perPrimFct); - } -} - + + + // ------------------------------------------------------------------ + // IMPLEMENTATION + // ------------------------------------------------------------------ + template + + inline __device__ + void traverse(bvh_t bvh, + aggregateNodeData_t aggregateData[], + primitive_t primitives[], + result_t result, + query_t queryPrim, + const approximateSubtreeFct_t &approximateSubtreeFct, + const perPrimFct_t perPrimFct) + { + struct StackEntry { + uint32_t idx; + }; + bvh3f::node_t::Admin traversalStack[64], *stackPtr = traversalStack; + bvh3f::node_t::Admin node = bvh.nodes[0].admin; + // ------------------------------------------------------------------ + // traverse until there's nothing left to traverse: + // ------------------------------------------------------------------ + while (true) { + + // ------------------------------------------------------------------ + // traverse INNER nodes downward; breaking out if we either + // find a leaf, or a encounter subtrees that can be either + // approximated or culled with the approximateSubtreeFct() + // ------------------------------------------------------------------ + while (true) { + if (node.count != 0) + // it's a boy! - seriously: this is not a inner node, step + // out of down-travesal and let leaf code pop in. + break; + + uint32_t n0Idx = (uint32_t)node.offset+0; + uint32_t n1Idx = (uint32_t)node.offset+1; + bvh3f::node_t n0 = bvh.nodes[n0Idx]; + bvh3f::node_t n1 = bvh.nodes[n1Idx]; + bool done0 = approximateSubtreeFct(bvh,aggregateData,n0Idx, + result,queryPrim); + bool done1 = approximateSubtreeFct(bvh,aggregateData,n1Idx, + result,queryPrim); + bool o0 = !done0; + bool o1 = !done1; + if (o0) { + if (o1) { + *stackPtr++ = n1.admin; + } else { + } + node = n0.admin; + } else { + if (o1) { + node = n1.admin; + } else { + // both children are too far away; this is a dead end + node.count = 0; + break; + } + } + } + + if (node.count != 0) { + for (int i=0;i Date: Sat, 21 Feb 2026 20:42:43 -0700 Subject: [PATCH 5/6] nbody apparently working --- cuBQL/builder/cuda/refit_aggregate.h | 38 ++++++++--- cuBQL/bvh.h | 3 + samples/CMakeLists.txt | 10 +++ samples/s07_aggregateNBody/aggregateNBody.cu | 72 +++++++++++++------- 4 files changed, 89 insertions(+), 34 deletions(-) diff --git a/cuBQL/builder/cuda/refit_aggregate.h b/cuBQL/builder/cuda/refit_aggregate.h index c2039d6..1c621ea 100644 --- a/cuBQL/builder/cuda/refit_aggregate.h +++ b/cuBQL/builder/cuda/refit_aggregate.h @@ -16,23 +16,33 @@ namespace cuBQL { template< typename T, int D, - typename AggregateNodeData, - typename AggregateFct> + typename AggregateNodeData + // , + // typename AggregateFct + > void refit_aggregate(BinaryBVH bvh, AggregateNodeData *d_aggregateNodeData, - const AggregateFct &aggregateFct, + void (*aggregateFct)(bvh3f, + AggregateNodeData[], + int), + // const AggregateFct &aggregateFct, cudaStream_t s =0, GpuMemoryResource &memResource =defaultGpuMemResource()); template + typename AggregateNodeData + // , + // typename AggregateFct + > __global__ void refit_aggregate_run(BinaryBVH bvh, AggregateNodeData *aggregateNodeData, - const AggregateFct &aggregateFct, + void (*aggregateFct)(bvh3f, + AggregateNodeData[], + int), + // const AggregateFct &aggregateFct, uint32_t *refitData) { int nodeID = threadIdx.x+blockIdx.x*blockDim.x; @@ -81,11 +91,17 @@ namespace cuBQL { template< typename T, int D, - typename AggregateNodeData, - typename AggregateFct> + typename AggregateNodeData + // , + // typename AggregateFct + > void refit_aggregate(BinaryBVH bvh, AggregateNodeData *d_aggregateNodeData, - const AggregateFct &aggregateFct, + // const AggregateFct &aggregateFct, + // __device__ + void (*aggregateFct)(bvh3f, + AggregateNodeData[], + int), cudaStream_t s, GpuMemoryResource &memResource) { @@ -93,11 +109,15 @@ namespace cuBQL { uint32_t *refitData = 0; memResource.malloc((void**)&refitData,numNodes*sizeof(*refitData),s); + CUBQL_CUDA_SYNC_CHECK(); refit_init<<>> (bvh.nodes,refitData,numNodes); + CUBQL_CUDA_SYNC_CHECK(); refit_aggregate_run<<>> (bvh,d_aggregateNodeData,aggregateFct,refitData); + CUBQL_CUDA_SYNC_CHECK(); memResource.free((void*)refitData,s); + CUBQL_CUDA_SYNC_CHECK(); // we're not syncing here - let APP do that } } diff --git a/cuBQL/bvh.h b/cuBQL/bvh.h index aeda611..73fe0cd 100644 --- a/cuBQL/bvh.h +++ b/cuBQL/bvh.h @@ -16,6 +16,9 @@ namespace cuBQL { build the tree; in particular, at which threshold to make a leaf */ struct BuildConfig { + BuildConfig(int makeLeafThreshold=0) + : makeLeafThreshold(makeLeafThreshold) + {} inline BuildConfig &enableSAH() { buildMethod = SAH; return *this; } inline BuildConfig &enableELH() { buildMethod = ELH; return *this; } typedef enum diff --git a/samples/CMakeLists.txt b/samples/CMakeLists.txt index e478516..851ef15 100644 --- a/samples/CMakeLists.txt +++ b/samples/CMakeLists.txt @@ -123,6 +123,16 @@ if (CUBQL_HAVE_CUDA) cuBQL_samples_common stb_image ) + set_target_properties(sample07_aggregateNBody + PROPERTIES + CUDA_SEPARABLE_COMPILATION ON + POSITION_INDEPENDENT_CODE ON + CUDA_USE_STATIC_CUDA_RUNTIME ON + CUDA_RESOLVE_DEVICE_SYMBOLS ON + CXX_VISIBILITY_PRESET hidden + CUDA_VISIBILITY_PRESET hidden + VISIBILITY_INLINES_HIDDEN ON + ) endif() diff --git a/samples/s07_aggregateNBody/aggregateNBody.cu b/samples/s07_aggregateNBody/aggregateNBody.cu index bd02c1a..cca8c7a 100644 --- a/samples/s07_aggregateNBody/aggregateNBody.cu +++ b/samples/s07_aggregateNBody/aggregateNBody.cu @@ -41,7 +41,8 @@ namespace nBody { /*! aggregation function that computes a node's aggregate data during aggragate_refit */ - inline __device__ + // inline + __device__ void aggregate(bvh3f bvh, AggregateNodeData nodeAggregates[], int nodeID) @@ -60,6 +61,26 @@ namespace nBody { } } + typedef void (*AggregateNodeFctPtr)(bvh3f, AggregateNodeData *, int); + __global__ + void k_get_aggregate(AggregateNodeFctPtr *d_result) + { + if (threadIdx.x != 0) return; + *d_result = aggregate; + } + + AggregateNodeFctPtr get_aggregate() + { + AggregateNodeFctPtr result = 0; + AggregateNodeFctPtr *d_resultPtr = 0; + CUBQL_CUDA_CALL(Malloc((void**)&d_resultPtr, + sizeof(AggregateNodeFctPtr))); + k_get_aggregate<<<1,32>>>(d_resultPtr); + CUBQL_CUDA_CALL(Memcpy((void*)&result,(void*)d_resultPtr, + sizeof(void*), cudaMemcpyDefault)); + return result; + } + /*! the final result type that results from iterating over the entire tree. FOr our simple n-body mock-up, it's simply a float to track sum_i{1/sqrDistance(queryPoint,dataPoint[i])} */ @@ -93,10 +114,10 @@ namespace nBody { // first, check if we can approximate this subtree float sqrDist = sqrDistance(node.bounds,queryPoint); float sqrDiag = sqrLength(node.bounds.size()); - const float approxThreshold = /* 1% */0.01f; + const float approxThreshold = /* 1% of fource*/0.01f; bool canApproximate = sqrDist > 0.f - && (sqrDiag / sqrDist <= sqr(approxThreshold)); + && (sqrDiag / sqrDist <= approxThreshold); if (canApproximate) { result.sumOfForces @@ -157,8 +178,7 @@ void runQueries(float *d_results, int main(int, char **) { - int numDataPoints = 10000; - int numQueryPoints = 20; + int numDataPoints = 100000; /*! generate 10,000 uniformly distributed data points */ std::vector dataPoints = cuBQL::samples::convert @@ -166,27 +186,17 @@ int main(int, char **) .generate(numDataPoints,290374)); std::cout << "#cubql: generated " << dataPoints.size() << " data points" << std::endl; - std::vector queryPoints - = cuBQL::samples::convert - (cuBQL::samples::UniformPointGenerator<3>() - .generate(numQueryPoints,/*seed*/1234567)); - std::cout << "#cubql: generated " << queryPoints.size() - << " query points" << std::endl; - vec3f *d_queryPoints = 0; vec3f *d_dataPoints = 0; box3f *d_primBounds = 0; - float *d_results = 0; - CUBQL_CUDA_CALL(Malloc((void **)&d_results,queryPoints.size()*sizeof(*d_results))); - CUBQL_CUDA_CALL(Malloc((void **)&d_queryPoints,queryPoints.size()*sizeof(vec3f))); - CUBQL_CUDA_CALL(Memcpy(d_queryPoints,queryPoints.data(), - queryPoints.size()*sizeof(queryPoints[0]), + CUBQL_CUDA_CALL(Malloc((void **)&d_dataPoints, + numDataPoints*sizeof(*d_dataPoints))); + CUBQL_CUDA_CALL(Memcpy((void *)d_dataPoints,dataPoints.data(), + numDataPoints*sizeof(*d_dataPoints), cudaMemcpyDefault)); - CUBQL_CUDA_CALL(Malloc((void **)&d_dataPoints,dataPoints.size()*sizeof(vec3f))); - CUBQL_CUDA_CALL(Memcpy(d_dataPoints,dataPoints.data(), - dataPoints.size()*sizeof(dataPoints[0]), - cudaMemcpyDefault)); - CUBQL_CUDA_CALL(Malloc((void **)&d_primBounds,dataPoints.size()*sizeof(box3f))); + + CUBQL_CUDA_CALL(Malloc((void **)&d_primBounds, + numDataPoints*sizeof(box3f))); computeBoxes<<>> (d_primBounds,d_dataPoints,numDataPoints); @@ -194,21 +204,33 @@ int main(int, char **) // generate initial cuBQL bvh over the data points // ------------------------------------------------------------------ bvh3f bvh; - cuBQL::gpuBuilder(bvh,d_primBounds,numDataPoints,BuildConfig()); - + cuBQL::gpuBuilder(bvh,d_primBounds,numDataPoints,BuildConfig(8)); + CUBQL_CUDA_SYNC_CHECK(); + // ------------------------------------------------------------------ // re-fit kernel-specific aggregate data on top of the bvh // ------------------------------------------------------------------ nBody::AggregateNodeData *d_nodeAggregates = 0; CUBQL_CUDA_CALL(Malloc((void **)&d_nodeAggregates, bvh.numNodes*sizeof(*d_nodeAggregates))); + CUBQL_CUDA_SYNC_CHECK(); cuBQL::cuda::refit_aggregate(bvh, d_nodeAggregates, - nBody::aggregate); + nBody::get_aggregate() + // nBody::aggregate + ); + CUBQL_CUDA_SYNC_CHECK(); // ------------------------------------------------------------------ // ready to run query // ------------------------------------------------------------------ + float *d_results = 0; + int numQueryPoints = numDataPoints; + // int numQueryPoints = std::min(numDataPoints,16*1024); + CUBQL_CUDA_CALL(Malloc((void **)&d_results, + numQueryPoints*sizeof(*d_results))); + + auto d_queryPoints = d_dataPoints; runQueries<<>> (d_results,bvh,d_nodeAggregates, d_dataPoints,d_queryPoints,numQueryPoints); From 2ee25402a14fc965f7ec8b5c63db74c04990de0b Mon Sep 17 00:00:00 2001 From: Ingo Wald Date: Sun, 22 Feb 2026 13:27:08 -0700 Subject: [PATCH 6/6] various cleanups --- cuBQL/builder/cuda/refit_aggregate.h | 34 +++----------------- samples/s07_aggregateNBody/aggregateNBody.cu | 23 ++++++++----- 2 files changed, 19 insertions(+), 38 deletions(-) diff --git a/cuBQL/builder/cuda/refit_aggregate.h b/cuBQL/builder/cuda/refit_aggregate.h index 1c621ea..3971ba2 100644 --- a/cuBQL/builder/cuda/refit_aggregate.h +++ b/cuBQL/builder/cuda/refit_aggregate.h @@ -13,6 +13,7 @@ namespace cuBQL { // ------------------------------------------------------------------ // INTERFACE // ------------------------------------------------------------------ + template< typename T, int D, @@ -25,24 +26,18 @@ namespace cuBQL { void (*aggregateFct)(bvh3f, AggregateNodeData[], int), - // const AggregateFct &aggregateFct, - cudaStream_t s - =0, + cudaStream_t s =0, GpuMemoryResource &memResource =defaultGpuMemResource()); template + typename AggregateNodeData> __global__ void refit_aggregate_run(BinaryBVH bvh, AggregateNodeData *aggregateNodeData, void (*aggregateFct)(bvh3f, AggregateNodeData[], int), - // const AggregateFct &aggregateFct, uint32_t *refitData) { int nodeID = threadIdx.x+blockIdx.x*blockDim.x; @@ -53,13 +48,6 @@ namespace cuBQL { // this is a inner node - exit return; - // box_t bounds; bounds.set_empty(); - // for (int i=0;iadmin.count;i++) { - // const box_t primBox = boxes[bvh.primIDs[node->admin.offset+i]]; - // bounds.lower = min(bounds.lower,primBox.lower); - // bounds.upper = max(bounds.upper,primBox.upper); - // } - int parentID = (refitData[nodeID] >> 1); while (true) { aggregateFct(bvh,aggregateNodeData,nodeID); @@ -75,11 +63,6 @@ namespace cuBQL { nodeID = parentID; node = &bvh.nodes[parentID]; parentID = (refitBits >> 1); - - // typename BinaryBVH::Node l = bvh.nodes[node->admin.offset+0]; - // typename BinaryBVH::Node r = bvh.nodes[node->admin.offset+1]; - // bounds.lower = min(l.bounds.lower,r.bounds.lower); - // bounds.upper = max(l.bounds.upper,r.bounds.upper); } } @@ -91,14 +74,9 @@ namespace cuBQL { template< typename T, int D, - typename AggregateNodeData - // , - // typename AggregateFct - > + typename AggregateNodeData> void refit_aggregate(BinaryBVH bvh, AggregateNodeData *d_aggregateNodeData, - // const AggregateFct &aggregateFct, - // __device__ void (*aggregateFct)(bvh3f, AggregateNodeData[], int), @@ -109,15 +87,11 @@ namespace cuBQL { uint32_t *refitData = 0; memResource.malloc((void**)&refitData,numNodes*sizeof(*refitData),s); - CUBQL_CUDA_SYNC_CHECK(); refit_init<<>> (bvh.nodes,refitData,numNodes); - CUBQL_CUDA_SYNC_CHECK(); refit_aggregate_run<<>> (bvh,d_aggregateNodeData,aggregateFct,refitData); - CUBQL_CUDA_SYNC_CHECK(); memResource.free((void*)refitData,s); - CUBQL_CUDA_SYNC_CHECK(); // we're not syncing here - let APP do that } } diff --git a/samples/s07_aggregateNBody/aggregateNBody.cu b/samples/s07_aggregateNBody/aggregateNBody.cu index cca8c7a..034d92c 100644 --- a/samples/s07_aggregateNBody/aggregateNBody.cu +++ b/samples/s07_aggregateNBody/aggregateNBody.cu @@ -41,7 +41,6 @@ namespace nBody { /*! aggregation function that computes a node's aggregate data during aggragate_refit */ - // inline __device__ void aggregate(bvh3f bvh, AggregateNodeData nodeAggregates[], @@ -60,24 +59,32 @@ namespace nBody { + nodeAggregates[node.offset+1].numBodiesInSubtree; } } + typedef void (*AggregateNodeFctPtr)(bvh3f bvh, + AggregateNodeData nodeAggregates[], + int nodeID); + - typedef void (*AggregateNodeFctPtr)(bvh3f, AggregateNodeData *, int); - __global__ - void k_get_aggregate(AggregateNodeFctPtr *d_result) - { - if (threadIdx.x != 0) return; - *d_result = aggregate; - } + __device__ AggregateNodeFctPtr aggregate_funcPtr = nBody::aggregate; AggregateNodeFctPtr get_aggregate() { AggregateNodeFctPtr result = 0; +#if 1 + // CUBQL_CUDA_CALL(Memcpy((void*)&result,(void*)&funcPtr, + // sizeof(void*), cudaMemcpyDefault)); + cudaMemcpyFromSymbol((void*)&result, + nBody::aggregate_funcPtr, + // nBody::funcPtr, + sizeof(void*)); +#else + AggregateNodeFctPtr *d_resultPtr = 0; CUBQL_CUDA_CALL(Malloc((void**)&d_resultPtr, sizeof(AggregateNodeFctPtr))); k_get_aggregate<<<1,32>>>(d_resultPtr); CUBQL_CUDA_CALL(Memcpy((void*)&result,(void*)d_resultPtr, sizeof(void*), cudaMemcpyDefault)); +#endif return result; }