Skip to content

Major rework of two-level BVH, including both builder and traversal #18

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
iwald-nvidia merged 9 commits into from
Jan 18, 2026
Merged

Major rework of two-level BVH, including both builder and traversal #18

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
dfb5173
starting on openmp build
ingowald Jan 16, 2026
b48290f
added default constructor for cuBQL::ray_t, so it can be used in structs
ingowald Jan 17, 2026
4968a23
added shorthand cuBQL::bvh3d typedef
ingowald Jan 17, 2026
08515af
bugfix in two-level ray traversal
ingowald Jan 18, 2026
3e2d052
Merge branch 'main' into iw/openmp
ingowald Jan 18, 2026
fe24fbc
various warning fixes
ingowald Jan 18, 2026
3d6a85b
removed openmp experiments
ingowald Jan 18, 2026
a46d0ab
Merge branch 'iw/openmp'
ingowald Jan 18, 2026
b2b50fd
bugfix: sah builder now respects buildConfig::maxAllowedLeafSize
ingowald Jan 18, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
50 changes: 33 additions & 17 deletions cuBQL/builder/cuda/sah_builder.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,8 +28,14 @@ namespace cuBQL {
struct {
AtomicBox<box_t> centBounds;
uint32_t offset;
#if 1
int8_t dim;
int8_t bin;
int32_t atomicSideDeciderBits;
#else
int8_t dim;
int8_t bin;
#endif
} openNode;
struct {
uint32_t offset;
Expand All @@ -54,7 +60,8 @@ namespace cuBQL {
inline __device__
void evaluateSAH(int &splitDim,
int &splitBin,
const SAHBins<box_t> &sah)
const SAHBins<box_t> &sah,
int maxAllowedLeafSize)
{
float bestCost = CUBQL_INF;

Expand All @@ -69,7 +76,7 @@ namespace cuBQL {
rAreas[b] = surfaceArea(box);
}
const float leafCost = rAreas[0] * rCount;
if (leafCost < bestCost) {
if (leafCost < bestCost && rCount <= maxAllowedLeafSize) {
bestCost = leafCost;
splitDim = -1;
}
Expand Down Expand Up @@ -241,12 +248,17 @@ namespace cuBQL {
int splitDim = -1;
int splitBin;
if (in.count > buildConfig.makeLeafThreshold) {
evaluateSAH(splitDim,splitBin,sah);
evaluateSAH(splitDim,splitBin,sah,buildConfig.maxAllowedLeafSize);
}
if (splitDim < 0) {
bool makeLeaf
= (splitDim < 0)
&& (in.count <= buildConfig.maxAllowedLeafSize);

if (makeLeaf) {
nodeState = DONE_NODE;
auto &done = nodes[nodeID].doneNode;
done.count = in.count;

// set this to max-value, so the prims can later do atomicMin
// with their position ion the leaf list; this value is
// greater than any prim position.
Expand Down Expand Up @@ -295,22 +307,26 @@ namespace cuBQL {

auto open = nodes[me.nodeID].openNode;
const box_t<T,D> primBox = primBoxes[me.primID];
int side = -1;
if (open.dim == -1) {
// // use top bit to atomically 'sort' prims left vs right
// side = atomicAdd(&nodes[me.nodeID].openNode.atomicSideDeciderBits,1u<<31)>>31;
side = atomicAdd(&nodes[me.nodeID].openNode.atomicSideDeciderBits,1u)&1;
} else {

const int d = open.dim;
T lo = open.centBounds.get_lower(d);
T hi = open.centBounds.get_upper(d);
const int d = open.dim;
T lo = open.centBounds.get_lower(d);
T hi = open.centBounds.get_upper(d);

T prim_d = T(0.5)*(primBox.get_lower(d)+primBox.get_upper(d));
T rel
= (prim_d - lo)
/ (hi - lo + 1e-20f);
int prim_bin = int(rel*(int)SAHBins<box_t<T,D>>::numBins);
prim_bin = max(0,min((int)SAHBins<box_t<T,D>>::numBins-1,prim_bin));
T prim_d = T(0.5)*(primBox.get_lower(d)+primBox.get_upper(d));
T rel
= (prim_d - lo)
/ (hi - lo + 1e-20f);
int prim_bin = int(rel*(int)SAHBins<box_t<T,D>>::numBins);
prim_bin = max(0,min((int)SAHBins<box_t<T,D>>::numBins-1,prim_bin));

int side = (prim_bin >= open.bin);
// printf("updateprim %i node %i state %i dim %i bin %i -> prim bin %i -> side %i\n",
// primID,me.nodeID,ns,open.dim,open.bin,
// prim_bin,side);
side = (prim_bin >= open.bin);
}
int newNodeID = open.offset+side;
auto &myBranch = nodes[newNodeID].openBranch;
atomicAdd(&myBranch.count,1);
Expand Down
1 change: 1 addition & 0 deletions cuBQL/bvh.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -140,6 +140,7 @@ namespace cuBQL {
// easy short-hand - though cubql also supports other types of bvhs,
// scalars, etc, this will likely be the most commonly used one.
using bvh3f = BinaryBVH<float,3>;
using bvh3d = BinaryBVH<double,3>;

#ifdef __CUDACC__
typedef BinaryBVH<float,2> bvh_float2;
Expand Down
2 changes: 2 additions & 0 deletions cuBQL/math/Ray.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,8 @@ namespace cuBQL {
struct ray_t {
using vec3 = vec_t<T,3>;

inline ray_t() = default;
inline ray_t(const ray_t &) = default;
__cubql_both ray_t(vec3 org, vec3 dir, T tMin, T tMax);
__cubql_both ray_t(vec3 org, vec3 dir);
vec3 origin;
Expand Down
5 changes: 3 additions & 2 deletions cuBQL/math/affine.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -164,6 +164,7 @@ namespace cuBQL {
typename AffineSpaceT<L>::vector_t xfmPoint(const AffineSpaceT<L>& m,
const typename AffineSpaceT<L>::vector_t &p)
{
using vector_t = typename AffineSpaceT<L>::vector_t;
return madd(vector_t(p.x),m.l.vx,
madd(vector_t(p.y),m.l.vy,
madd(vector_t(p.z),m.l.vz,
Expand All @@ -185,11 +186,11 @@ namespace cuBQL {
/// Comparison Operators
////////////////////////////////////////////////////////////////////////////////

template<typename L> inline
template<typename L> inline __cubql_both
bool operator ==( const AffineSpaceT<L>& a, const AffineSpaceT<L>& b )
{ return a.l == b.l && a.p == b.p; }

template<typename L> inline
template<typename L> inline __cubql_both
bool operator !=( const AffineSpaceT<L>& a, const AffineSpaceT<L>& b )
{ return a.l != b.l || a.p != b.p; }

Expand Down
18 changes: 14 additions & 4 deletions cuBQL/math/linear.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -132,8 +132,13 @@ namespace cuBQL {
/// Comparison Operators
////////////////////////////////////////////////////////////////////////////////

template<typename T> inline __cubql_both bool operator ==( const LinearSpace2<T>& a, const LinearSpace2<T>& b ) { return a.vx == b.vx && a.vy == b.vy; }
template<typename T> inline __cubql_both bool operator !=( const LinearSpace2<T>& a, const LinearSpace2<T>& b ) { return a.vx != b.vx || a.vy != b.vy; }
template<typename T> inline __cubql_both
bool operator ==( const LinearSpace2<T>& a, const LinearSpace2<T>& b )
{ return a.vx == b.vx && a.vy == b.vy; }

template<typename T> inline __cubql_both
bool operator !=( const LinearSpace2<T>& a, const LinearSpace2<T>& b )
{ return a.vx != b.vx || a.vy != b.vy; }

////////////////////////////////////////////////////////////////////////////////
/// Output Operators
Expand Down Expand Up @@ -307,8 +312,13 @@ namespace cuBQL {
/// Comparison Operators
////////////////////////////////////////////////////////////////////////////////

template<typename T> inline bool operator ==( const LinearSpace3<T>& a, const LinearSpace3<T>& b ) { return a.vx == b.vx && a.vy == b.vy && a.vz == b.vz; }
template<typename T> inline bool operator !=( const LinearSpace3<T>& a, const LinearSpace3<T>& b ) { return a.vx != b.vx || a.vy != b.vy || a.vz != b.vz; }
template<typename T> inline __cubql_both
bool operator ==( const LinearSpace3<T>& a, const LinearSpace3<T>& b )
{ return a.vx == b.vx && a.vy == b.vy && a.vz == b.vz; }

template<typename T> inline __cubql_both
bool operator !=( const LinearSpace3<T>& a, const LinearSpace3<T>& b )
{ return a.vx != b.vx || a.vy != b.vy || a.vz != b.vz; }

////////////////////////////////////////////////////////////////////////////////
/// Output Operators
Expand Down
4 changes: 3 additions & 1 deletion cuBQL/math/vec.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -169,7 +169,9 @@ namespace cuBQL {
using vec_t_data<T,3>::y;
using vec_t_data<T,3>::z;

inline __cubql_both vec_t() {}
inline vec_t() = default;
inline vec_t(const vec_t &) = default;
// inline __cubql_both vec_t() {}
inline __cubql_both vec_t(const T &t) { x = y = z = t; }
inline __cubql_both vec_t(T x, T y, T z)
{ this->x = x; this->y = y; this->z = z; }
Expand Down
7 changes: 6 additions & 1 deletion cuBQL/queries/triangleData/Triangle.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,7 +34,12 @@ namespace cuBQL {
{
using vec3 = vec_t<T,3>;
using box3 = box_t<T,3>;


inline triangle_t() = default;
inline triangle_t(const triangle_t &) = default;
inline __cubql_both triangle_t(vec3 a, vec3 b, vec3 c)
: a(a), b(b), c(c)
{}
inline __cubql_both box3 bounds() const;
inline __cubql_both vec3 sample(float u, float v) const;
inline __cubql_both vec3 normal() const;
Expand Down
35 changes: 24 additions & 11 deletions cuBQL/traversal/rayQueries.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -168,6 +168,21 @@ namespace cuBQL {
return tin <= tout;
}

template<typename T>
inline __cubql_both
void rayBoxTest(T &tin, T &tout,
ray_t<T> ray, box_t<T,3> box)
{
using vec3 = vec_t<T,3>;
vec3 inv = rcp(ray.direction);
vec3 lo = (box.lower - ray.origin) * inv;
vec3 hi = (box.upper - ray.origin) * inv;
vec3 nr = min(lo,hi);
vec3 fr = max(lo,hi);
tin = max(ray.tMin,reduce_max(nr));
tout = min(ray.tMax,reduce_min(fr));
}

template<typename T>
inline __cubql_both
bool rayIntersectsBox(float &ret_t0,
Expand Down Expand Up @@ -462,7 +477,7 @@ namespace cuBQL {

node_t *tlasSavedNodePtr = 0;
uint32_t *tlasSavedPrimIDs = 0;
vec3f saved_dir, saved_org;
vec_t<T,3> saved_dir, saved_org;

if (ray.direction.x == (T)0) ray.direction.x = T(1e-20);
if (ray.direction.y == (T)0) ray.direction.y = T(1e-20);
Expand Down Expand Up @@ -512,6 +527,9 @@ namespace cuBQL {
enterBlas(transformed_ray,blas,instID);
ray.origin = transformed_ray.origin;
ray.direction = transformed_ray.direction;
if (ray.direction.x == (T)0) ray.direction.x = T(1e-20);
if (ray.direction.y == (T)0) ray.direction.y = T(1e-20);
if (ray.direction.z == (T)0) ray.direction.z = T(1e-20);
rcp_dir = rcp(ray.direction);
bvh.nodes = blas.nodes;
bvh.primIDs = blas.primIDs;
Expand All @@ -531,10 +549,13 @@ namespace cuBQL {
bool o0 = rayIntersectsBox(node_t0,ray,rcp_dir,n0.bounds);
bool o1 = rayIntersectsBox(node_t1,ray,rcp_dir,n1.bounds);

if (dbg)
if (dbg) {
// dout << " node L " << n0.bounds << "\n";
// dout << " node R " << n1.bounds << "\n";
printf("children L hit %i dist %f R hit %i dist %f\n",
int(o0),node_t0,
int(o1),node_t1);
}
if (o0) {
if (o1) {
if (stackPtr-traversalStack >= STACK_DEPTH) {
Expand Down Expand Up @@ -620,15 +641,7 @@ namespace cuBQL {
for (int i=0;i<count;i++) {
int primIdx = offset+i;
if (primIDs) primIdx = primIDs[primIdx];
ray.tMax = min(ray.tMax,intersectPrim(offset+i));

// if (primIDs == nullptr) {
// if (dbg) printf("leaf %p offset %i\n",bvh.primIDs,(int)offset);
// for (int i=0;i<count;i++)
// ray.tMax = min(ray.tMax,intersectPrim(offset+i));
// } else {
// for (int i=0;i<count;i++)
// ray.tMax = min(ray.tMax,intersectPrim(primIDs[offset+i]));
ray.tMax = min(ray.tMax,intersectPrim(primIdx));
}
if (dbg) printf("LEAVING LEAF! t = %f\n",ray.tMax);
return ray.tMax;
Expand Down