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/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..eb51b6c --- /dev/null +++ b/cuBQL/builder/cuda/refit.h @@ -0,0 +1,93 @@ +// 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) + { + int numNodes = bvh.numNodes; + + uint32_t *refitData = 0; + memResource.malloc((void**)&refitData,numNodes*sizeof(*refitData),s); + + refit_init<<>> + (bvh.nodes,refitData,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/refit_aggregate.h b/cuBQL/builder/cuda/refit_aggregate.h new file mode 100644 index 0000000..3971ba2 --- /dev/null +++ b/cuBQL/builder/cuda/refit_aggregate.h @@ -0,0 +1,98 @@ +// 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" +#include "cuBQL/builder/cuda/refit.h" + +namespace cuBQL { + namespace cuda { + + // ------------------------------------------------------------------ + // INTERFACE + // ------------------------------------------------------------------ + + template< + typename T, + int D, + typename AggregateNodeData + // , + // typename AggregateFct + > + void refit_aggregate(BinaryBVH bvh, + AggregateNodeData *d_aggregateNodeData, + void (*aggregateFct)(bvh3f, + AggregateNodeData[], + int), + cudaStream_t s =0, + GpuMemoryResource &memResource + =defaultGpuMemResource()); + + template + __global__ + void refit_aggregate_run(BinaryBVH bvh, + AggregateNodeData *aggregateNodeData, + void (*aggregateFct)(bvh3f, + AggregateNodeData[], + int), + 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; + + 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); + } + } + + + + // ------------------------------------------------------------------ + // IMPLEMENTATION + // ------------------------------------------------------------------ + template< + typename T, + int D, + typename AggregateNodeData> + void refit_aggregate(BinaryBVH bvh, + AggregateNodeData *d_aggregateNodeData, + void (*aggregateFct)(bvh3f, + AggregateNodeData[], + int), + 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/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 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/cuBQL/traversal/aggregateApproximate.h b/cuBQL/traversal/aggregateApproximate.h new file mode 100644 index 0000000..a85e55c --- /dev/null +++ b/cuBQL/traversal/aggregateApproximate.h @@ -0,0 +1,201 @@ +// 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 { + + // ------------------------------------------------------------------ + // 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. + + `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 + < /*! T/D that describes the BVH data/dimensionality */ + typename T, int D, + /*! the aggregate data we store per node (to base the user's + culling test on) */ + typename aggregateNodeData_t, + /* the type of primitmives stored in leaves, for when traversal + reaches leaves */ + typename primitive_t, + /*! the result type of the query, for when traversal decides to + either aprpxoimate a subtree (and has to udpate the result + with that subtree's approximateion), or when prim tests are + done */ + typename result_t, + /*! the type of query primitive that the traversal operates + on */ + typename query_t, + /*! the lambda function that tests a subtree on whether it can + be solved with an approximation */ + typename approximateSubtreeFct_t, + /*! the lambda function that gets executed for each prim (those + that did not get approximates higher up the tree */ + typename perPrimFct_t> + 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); + + + // ------------------------------------------------------------------ + // 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