* improved the vdb option in the vdb_tool
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* vdb_tool: address review of the vdb/geo name-resolution option
Two issues raised in review of the name-resolution feature:
1. Resolution was written back into Option::value, destroying the parsed source
expression. Because a loop body revisits the same Action every iteration, this
pinned the option to whatever the first iteration resolved: the documented
pattern "-for v=0,3 -render vdb={$v} -end" rendered age 0 on every pass. It
would also have corrupted the source form that -write emits to a config file,
which reads Option::value directly.
Resolution now runs per-read on a copy instead of mutating the option:
Parser::beforeActionRun is replaced by Parser::onGetOption, invoked from
Parser::getStr() after expression expansion. Every accessor (get<T>, getVec<T>,
getVec3<T>) funnels through getStr(), so coverage is unchanged, but the stored
expression survives and is re-expanded and re-resolved each iteration.
2. Expanding a duplicated name to several ages was unsafe. Most actions read
vdb/geo as a single age via get<int>() and cannot parse a list at all (29 such
call sites versus 19 list consumers), and list consumers such as -clear erase
by index while iterating, so a multi-age value could delete the wrong entry.
A name must now identify exactly one entry; an ambiguous name is rejected with
an error listing the matching ages so an explicit index can be used. Because
this fires during option resolution, the action aborts before mutating anything.
Also exempts -soup2ls, whose "vdb" option selects an output resolution level of
the shrink-wrap hierarchy rather than an entry on the grid stack.
Tests: inverted the ambiguity case, added the requested loop regression using
vdb={$v}, a bare-name-in-loop case, a no-partial-deletion check, and a
singleton-consumer (-rename) ambiguity check.
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* vdb_tool: fix partial deletion in -clear with multiple ages
"-clear" erased entries one at a time inside its loop, but an age is a distance
from the back of the stack, so removing one shifted the ages of those still to be
processed. On a three-grid stack "-clear vdb=0,2" erased age 0, leaving two
grids, and then failed on age 2 -- having already destroyed a grid. An invalid
age behaved the same way, aborting only after earlier deletions had been applied.
Reachable with plain integer indices and unrelated to name resolution, though the
name-resolution work made it easy to hit.
Every requested age is now resolved to a list iterator, and duplicate ages
removed, before anything is erased. An out-of-range index therefore leaves both
stacks untouched, matching how "-copy" validates all of its indices up front, and
a repeated age such as "vdb=0,0" deletes a single entry instead of erasing the
same iterator twice (which is undefined behaviour). Resolving up front is safe
because mGrid/mGeom are std::lists, whose iterators remain valid when other
elements are erased.
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* vdb_tool: rename the -soup2ls "vdb" option to "levels"
"-soup2ls" used "vdb" for something no other action does: selecting which
resolution levels of the shrink-wrap hierarchy to output, rather than referring
to an entry on the VDB stack. That collision forced the grid-name resolution
introduced earlier in this branch to special-case the action, and it read badly
in help output, needing the caveat "unlike other actions, here vdb selects
outputs (not inputs)".
Renaming the option to "levels" removes both problems: the resolution hook no
longer needs an exemption, and every remaining "vdb"/"geo" option in the tool
now uniformly means "an entry on one of the two stacks". No other option on the
action begins with "l", so the name stays unambiguous under prefix matching.
Also drops the stale reference to "vdb" in the neighbouring "dim" description,
copy-pasted from actions where "vdb" supplies a transform; -soup2ls never read
it that way.
This option has not appeared in a release (it is still described only in
pendingchanges), so the rename is not a break for any published version and the
existing changelog entry is amended in place rather than adding a deprecation.
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* review comments
Signed-off-by: Ken Museth <ken.museth@gmail.com>
---------
Signed-off-by: Ken Museth <ken.museth@gmail.com>
Signed-off-by: Ken Museth <1495380+kmuseth@users.noreply.github.com>
* Add fused ReadAccessor::getDimAndActive API
Introduce a fused getDim + isActive traversal that shares the
Root->Internal->Leaf descent, the accessor cache update, and the
terminating-level mask read between the two queries. Adds a
DimAndActive POD, getDimAndActiveAndCache on Root/Internal/Leaf,
and getDimAndActive on all four ReadAccessor specializations.
Two policy tags control how `active` is resolved when an
InternalNode skip-flag fires along the descent:
- ActiveExact (default): preserves byte-for-byte semantics of
separate getDim + getAndCache<GetState>. `active` is correct
even when the skip-flag is set, at the cost of forcing the
child descent. Required by callers that read `active` at
upper levels (e.g. fvdb HDDA iterators).
- ActiveOnLeafOnly: cheap skip-flag short-circuit; `active` is
unspecified when the returned `dim` exceeds leaf-dim, and the
caller must gate on `dim` first
(`if (hdda.dim() > 1 || !active) continue;`).
HDDA.h ZeroCrossing is intentionally not migrated -- on the
register-tight sm_120 kernel, any fused shape grows reg/thread by
2-3 and drops occupancy from 4 to 3 blocks/SM. The kernel is
memory-latency bound (~14% achieved occupancy) and the L1 cache
absorbs the second descent at ~87% hit rate, so the occupancy hit
dominates the descent-fusion savings. Rationale recorded inline.
Also rework the ex_raytrace_level_set benchmark harness for honest
measurement: kernel-only cudaEvent timing on the CUDA path, 2
untimed warmup iters, min/median/mean/max reporting,
numIterations 50 -> 200, block size 512 -> 256.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb: pack DimAndActive into a single int32_t
The fused getDim+isActive traversal returns a {dim, active} pair. The
otherwise-natural {uint32_t, bool} struct layout passes through two
registers (or a packed 64-bit pair) at the function-return boundary.
NanoVDB's getDim return value is bounded by the largest internal-node
DIM in the default FloatTree hierarchy (4096 for an upper internal --
2^12, well below 2^31). The sign bit is unused, so we can pack:
low 31 bits = dim
sign bit = active
The struct keeps its existing API surface (aggregate-init via a 2-arg
constructor, accessors dim() and active()), so callers that used
`return {dim, active}` continue to work; consumers that previously
read `.dim` / `.active` as members switch to `.dim()` / `.active()`.
Effect on consumers:
- On the static ZeroCrossing kernel (which does not use the fused
API; rationale comment updated below): no change.
- On lean_hdda + fusion (a separate experiment) the renderOp kernel
drops from 62 regs (struct encoding) to 61 regs (packed encoding).
Doesn't cross an occupancy tier on sm_120 at b=256, so the wall-
time delta is within noise. Still useful: any consumer holding
DimAndActive live across other work gains 1 reg of headroom for
no semantic change.
- On fvdb's HDDA iterators (the productive consumer that reads
active unconditionally at upper levels), the packing reduces
return-boundary register pressure in a hot path.
Also expand the rationale comment in HDDA.h's ZeroCrossing. The
original justification framed the don't-migrate decision as a
register-cliff problem; testing against lean_hdda (which has 3 regs
of headroom for fusion) shows the same negative wall-time result
without an occupancy hit. The real reason is the second descent is
L1-warm at ~87 % hit rate, so eliminating it saves ~no cycles, and
the fused call's slightly heavier per-call overhead nets out slightly
worse. Verified both struct- and packed-encoded fusion across both
regimes.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb: drop ZeroCrossing rationale comment
The long NB block that explained why HDDA::ZeroCrossing keeps
separate getDim + isActive calls (instead of using the fused API
this branch adds) was useful during the iteration but doesn't earn
its keep in the source tree -- a future reader doesn't need to relive
the cliff/lean/descent-cost investigation that justified the decision.
The rationale is preserved in the PR description: ZeroCrossing's
`dim > 1 || !active` short-circuit makes the active read dead at upper
levels, and the L1 cache absorbs the leaf-level descent at ~87% hit
rate, so fusion saves essentially no cycles while adding a small per-
call overhead. The fused API is retained for callers that read active
unconditionally (fvdb's HDDA iterators -- -27% to -62% wall time win).
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* comment cleanup
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* comment update
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb: fix DimAndActive packing UB and add getDimAndActive test
The DimAndActive constructor packed active via 'int32_t(a) << 31',
which shifts 1 into the sign bit of a signed int -- undefined behavior
under C++17 (the repo's minimum standard). Do the OR in uint32_t so the
shift is well-defined, and mask dim to 31 bits so a stray high bit can't
clobber the active bit.
Add TestNanoVDB.GetDimAndActive, which asserts the fused
getDimAndActive matches the separate getDim + isActive it replaces:
- ActiveExact: {dim, active} equals getDim(ijk, ray) and isActive(ijk)
byte-for-byte across a box spanning interior, surface and exterior.
- ActiveOnLeafOnly: dim always matches getDim; active matches isActive
whenever dim == 1 (the contract gate).
The sphere's narrow band is too thin to fill an InternalNode, so a
skip-flagged internal node above an active voxel cannot arise from the
geometry. The test forces that case by setting the skip-flag on the
lower node over a known-active voxel, then confirms ActiveExact descends
past it (active == true) while ActiveOnLeafOnly short-circuits
(active == false), both reporting the skipped dim.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* comment cleanup
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb: address PR review on DimAndActive and raytrace example
- DimAndActive: store the packed field as uint32_t instead of int32_t.
The previous int32_t form converted a uint32_t with bit 31 set to a
signed int (implementation-defined) and left the defaulted ctor's
member uninitialized. uint32_t packs/decodes with no conversion, the
default ctor zero-initializes, and the size/ABI is unchanged.
- ex_raytrace_level_set reportStats: take samples by const reference,
early-return on an empty vector (front()/back() would otherwise be
UB if numIterations were 0), and sort a local copy.
- ex_raytrace_level_set common.h: check the cudaEventDestroy return
codes with NANOVDB_CUDA_CHECK_ERROR, consistent with the rest of the
CUDA timing path.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb: assert dim fits in 31 bits when packing DimAndActive
The packed encoding masks dim to 31 bits (bit 31 holds active). dim is
1u << TOTAL, bounded by the deepest node dim (4096 for the default tree),
so it always fits -- but a pathological tree config with TOTAL >= 31 would
be silently truncated. Add a debug assert documenting the invariant, in
the same style as the version-packing asserts above. Release builds are
unaffected (assert compiles out; the mask still protects the active bit).
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Add pending changes entry
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Fix compile warnings
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Reduce comment verbosity
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
cuda::copyTo is the explicit, stream-carrying deep copy between grid
handles in different address spaces: a single-space device handle copies to
a host-readable handle (HostBuffer or a cuda::Buffer over a host-accessible
resource) and back, and device-to-device across different resources. An
optional prototype buffer routes the destination allocation through a
caller-supplied resource, which non-default-constructible resources (e.g.
ResourceRef) require. A host-readable destination synchronizes the transfer
stream before returning; a device destination parses and validates its
metadata on the transferred bytes. Dual-space buffers are rejected with a
named error: they keep deviceUpload/deviceDownload until their retirement.
GridHandle now accepts buffers over host-accessible resources (e.g.
PinnedResource) as host-readable single-space handles instead of rejecting
them: metadata parses on the host, the host accessors and write() apply,
and read() and copy() allocate through the buffer's resource via a shared
host-storage factory (BufferT::create for buffers that provide it, the
resource otherwise). The byte-element requirement now covers host-single
buffers, and splitGrids/mergeGrids name their create() requirement instead
of failing inside the template body.
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* NanoVDB: chain DeviceBuffer::recordUse and expose orderAfterPriorUses (CUDA)
recordUse re-records the single per-device tracking event, and re-recording
moves an event: uses recorded on streams A then B left only B's coverage, so
the device free could run while A's work was in flight — violating the
invariant documented on mEvents ("every use waits on this event before
issuing work and re-records it afterwards"). Record now waits on the prior
capture first, so the event transitively covers every recorded use. Also make
orderAfterPriorUses public as the consume-side companion, so external
consumers (e.g. the Python bindings' zero-copy exports) can order their own
stream after the tracked uses. Adds a DeviceBufferChainedRecordUse regression
test that fails before the fix and passes after.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Address review: doc concurrent-reader trade-off, guard vacuous test pass, tidy access sections
- recordUse's doc note now states that concurrent READERS recording uses
serialize behind each other (the single per-device event cannot distinguish
read-read from write-read) and names the upgrade path (read/write-separated
or per-record events) should a profile ever surface it.
- DeviceBufferChainedRecordUse now GTEST_SKIPs when the allocator did not
recycle the freed block, making a vacuous run visible instead of green.
Verified on sm_120 that stream-ordered pools recycle WITH the dependency
attached even post-fix, so the skip does not fire on a working platform.
- orderAfterPriorUses moved into the public section next to recordUse instead
of a public/private island inside the private section.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* NanoVDB: assign each upper-node tile to a single GPU in DistributedPointsToGrid
The device segment boundaries produced by the initial even split were
rebalanced by adjusting only adjacent pairs of boundaries. That cannot
consolidate an upper-node tile spanning three or more GPUs, since a
fully-interior device lies entirely within the tile. Downstream
construction assumes each tile - and therefore each lower node, leaf node
and voxel - is owned by exactly one device, so multiple devices built the
same leaf concurrently and raced on its value mask, silently dropping
active voxels.
Replace the pairwise GPU rebalance with a global, monotonic snap performed
on the host over the (small) set of device boundaries. The keys are
globally sorted at that point, so a tile boundary is simply a position
where the key changes; fully-interior devices are left empty, which the
rest of the pipeline already handles.
Add two multi-GPU regression tests: SingleUpperNode, which places all
input in one upper-node tile and asserts the exact unique active-voxel
count, and MatchesSingleGpu, which cross-checks topology and occupancy
against the single-GPU builder on the same input.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Added a `snapBoundariesToRunsKernel` and running on device to avoid stream sync
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
* added build support on Ampere
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* vdb_tool: link vdb_tool_test against GTest::gtest only
The unit test provides its own main() and uses only gtest (no gmock), so it
needs just GTest::gtest. Listing gmock / gtest_main / gmock_main as well pulled
libgtest.a and libgmock.a onto the link line multiple times, which the macOS
(Xcode 15+) linker flags as "ignoring duplicate libraries". Linking only
GTest::gtest removes the warning and also works against gtest-only GoogleTest
installs (BUILD_GMOCK=OFF).
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* review comments
Signed-off-by: Ken Museth <ken.museth@gmail.com>
---------
Signed-off-by: Ken Museth <ken.museth@gmail.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
* added histogram action to vdb_tool
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* now -stats and -histogram handles vector grids
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* fixed various issues (see comment in PR)
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* addressed reviewer comments
Signed-off-by: Ken Museth <ken.museth@gmail.com>
---------
Signed-off-by: Ken Museth <ken.museth@gmail.com>
The handles (GridHandle, NodeManagerHandle, tools::VoxelBlockManager) now
release storage through the buffer's destroy() when it provides one and
fall back to the legacy clear() otherwise, detected by a BufferHasDestroy
trait companion. That lets the transitional cuda::Buffer::clear() and
cuda::BufferView::clear() finally carry their [[deprecated]] attribute
(BufferView gains destroy() with its detach semantics); the legacy
HostBuffer/DeviceBuffer clear() methods are unaffected.
cuda::createNodeManager gains a single-space overload taking a memory
resource: the returned handle owns a cuda::Buffer<std::byte,
ResourceRef<ResourceT>> whose storage and size scratch allocate through the
injected resource, which must outlive the handle -- the builders'
ownership contract. NodeManagerHandle supports single-space buffers
(deviceMgr maps onto the buffer, host accessors become compile-time
errors, constructors member-initialize so non-default-constructible
buffers work), and tools::cuda::GridStats builds its temporary NodeManager
through its injected resource, closing the last allocation that bypassed
ResourceT. The dual-space overload's raw size-scratch pair becomes an RAII
buffer, and both overloads now synchronize before the host reads the
device-computed size, which the old code consumed without ordering.
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* Add NanoVDB TEACHME interactive tutorial + docs-compile test harness
doc/nanovdb/TEACHME/ is an interactive, LLM-guided tutorial for the
NanoVDB user API: 10 modules (reading grids, ReadAccessor, math/HDDA
sampling, GPU kernels, GPU grid builders, topology operators + data
re-homing via Injection, IndexGrid/VoxelBlockManager, OpenVDB
conversion), a cheat sheet, and a GPU level-set ray-march capstone.
To keep the tutorial from drifting from the API, the test harness in
doc/nanovdb/TEACHME/test/ extracts every C++/CUDA code block from the
docs and compiles it against the real NanoVDB headers, and validates
that every #include and prose header reference resolves. ~64 blocks
compile green.
A CI workflow (.github/workflows/teachme.yml) runs the harness whenever
the NanoVDB headers or the tutorial change, so API drift is caught
automatically. It is compile-only and needs no GPU.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: add pending-changes note for the TEACHME tutorial
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: note the TEACHME tutorial in pendingchanges
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
---------
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
* NanoVDB: support single-space device buffers in GridHandle (CUDA)
GridHandle can now own a nanovdb::cuda::Buffer<std::byte, R>: a new
hasDeviceSingle buffer trait selects a constructor (implemented in
cuda/GridHandle.cuh) that parses the grid metadata through the device,
with every operation -- the validated header walk, metadata scratch
allocation through the buffer's resource, the copy kernel and the
readback -- ordered on the buffer's retained stream.
deviceData()/deviceGrid() work on such handles. The host accessors
(data, grid, gridData, gridMetaData) are SFINAE-removed for them,
since the handle owns no host-readable bytes; the read/write I/O
members and splitGrids/mergeGrids stay addressable (the python
bindings take write's address via overload_cast) and fail instead
with an explanatory static_assert when instantiated for a
single-space handle.
The device parse validates the whole grid chain (per-header bounds
checks against the allocation) before launching the metadata kernel,
so truncated buffers and forged mGridCount/mGridSize headers are
rejected with an exception instead of an out-of-bounds device read;
the pre-existing dual-space parse had the same hole and now shares the
validation helper. Its metadata scratch uses MallocResource so that
long-standing path keeps working on devices without memory-pool
support, and the dirty-flag scratch in splitGridHandles and
mergeGridHandles is a resource-aware buffer as well -- together
replacing the file's six raw cudaMalloc/mallocAsync sites, one of
which was missing its cudaCheck.
copy() dispatches on the traits: host buffers keep the memcpy path,
single-space buffers deep-copy device-to-device through their own
resource on the retained stream and reuse the host-resident metadata
instead of re-parsing. Buffer gains ElementType/ResourceType typedefs,
a resource() accessor, and a no-argument copy() for stream-ordered
resources that orders the copy on the retained stream.
Host accessibility is a property of the resource, not the element
type: the new is_host_accessible_resource trait detects a
HOST_ACCESSIBLE marker (declared by PinnedResource and forwarded by
ResourceRef and AsyncFromSync). A cuda::Buffer over a host-accessible
resource is rejected at handle scope with a named error for now --
GridHandle's host paths require an allocation interface cuda::Buffer
does not yet provide -- and non-byte element types fail loudly in the
single-space constructor. Both hasDeviceSingle and hasHostSingle are
detected rather than required, so pre-existing BufferTraits
specializations in and out of tree compile unchanged.
Part of #2232 (step 3).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: validate and parse the device grid chain in one kernel pass
Replace the per-header device-to-host walk in the GridHandle constructors
with a single host read of the head header (which sizes the metadata
scratch) followed by one kernel that validates every header and fills the
metadata scratch in the same pass, plus one readback and one
synchronization regardless of the grid count. The previous walk issued one
synchronizing copy per header, so constructing the handle that
cuda::mergeGridHandles returns cost O(N) stream drains for N grids; it now
costs the same two round trips as a single-grid buffer, and the separate
metadata kernel launch, its scratch copy and final synchronization are
gone. Bounds are still checked before every header read, on the host for
the head and on the device for the rest, so a truncated buffer or forged
header can never read out of bounds.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: normalize grid index/count in addBlindData and indexToGrid output
Both tools copy the source grid's header verbatim into a new single-grid
buffer, so a source grid taken from a multi-grid handle left a stale
mGridIndex/mGridCount in the output. The returned handle's metadata parse
then walked a grid chain that is not there: past the end of the buffer for
index 0 of a multi-grid source, or starting at a nonzero index. The output
header is now normalized to index 0 / count 1 in the same kernels that
already patch its other fields, before each tool's existing checksum
update. The other handle-returning GPU tools were audited: the topology
ops and MeshToGrid normalize in TopologyBuilder, and the point builders
initialize fresh headers via GridData::init.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: share the buffer trait detectors and deduplicate GridHandle lookups
BufferHasDeviceSingle/BufferHasHostSingle move next to the BufferTraits
primary in HostBuffer.h: they are companions to the trait protocol, not
GridHandle-specific, so other handle types can consume them without
including GridHandle.h. The three identical grid-lookup bodies behind
grid() and both deviceGrid() families collapse into one private gridAt()
helper, and the supported-combination guard for GridHandle::copy() now
lives solely in the no-arg overload, so the two spellings cannot drift
apart when cross-space transfers relax it. cuda::Buffer documents the
single-space device-buffer concept on its BufferTraits specialization --
naming exactly the interface a consumer of hasDeviceSingle may rely on --
and gains a rebind alias for generic code that constructs sibling buffers.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
---------
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* added vdb option to -soup2ls action and improved error handeling
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* improved the documentation
Signed-off-by: Ken Museth <ken.museth@gmail.com>
---------
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* NanoVDB CUDA: select-based VoxelBlockManager inverse-map decode
Replace the O(nLeaves x 512) cooperative sweep in decodeInverseMaps with an
O(1)-per-slot select: one thread per output slot ranks itself into its leaf via
the jumpMap popcount, then locates its voxel with the leaf's 9-bit mPrefixSum plus
an in-word __fns bit select. Same signature, byte-exact decode maps.
Verified byte-exact on the VBM goldens for all consumers (decode/box/lap/weno) at
widths 64/128/256/512, and faster for every consumer with no regression: decode-only
2.17x/1.97x/1.62x/1.42x (w6/w7/w8/w9), 7-pt Laplacian -8.5..-15.4%, box/weno -4..-9%.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* NanoVDB: add pending changes entry for the select-based VBM decode
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Additional explanatory comments and cleanup
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* NanoVDB CUDA: consolidate VBM decode to a thread-local per-slot API
Replace the cooperative shared-memory decodeInverseMaps with decodeInverseMap:
a per-slot decode into registers (no shared memory or synchronization; callable
from divergent threads), factored into jumpMapRank + selectVoxelInLeaf helpers.
computeBoxStencil now takes the decoded (leafIndex, voxelOffset) by value.
Block-level facts formerly read from the materialized maps derive directly from
the VBM metadata: the block's first leaf is firstLeafID, slot p starts a new
leaf iff jumpMap bit p is set, and the spanned-leaf count is 1 + the jumpMap
popcount.
Decode maps and all 27 box-stencil taps verified byte-exact against the prior
select decode at widths 64/128/512 across 16-86% leaf occupancy. Decode-only
+14..17% at width 128 and +8..13% at width 512 over the prior select, neutral
at width 64; vs master's sweep: 1.56-2.08x (w7) and 1.24-1.68x (w9) across the
sparsity range.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* NanoVDB CUDA: preserve decodeInverseMaps as a wrapper over the per-slot decode
Restore the cooperative decodeInverseMaps with its pre-existing signature as a
thin wrapper that materializes all BlockWidth slots via decodeInverseMap
(grid-striding over the slots, ending in __syncthreads), so existing callers
keep compiling and behaving unchanged. The per-slot decodeInverseMap remains
the preferred form for consumers that only read their own slot's result.
The unit test decodes through the wrapper (which exercises decodeInverseMap
internally) and resolves via the by-value computeBoxStencil; both
VoxelBlockManager tests pass.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* NanoVDB: own the small builders' device scratch with cuda::Buffer (CUDA)
Route the last raw mallocAsync/freeAsync pairs in the GPU tools through
RAII cuda::Buffer scratch over an injectable resource, completing the
audit that PointsToGrid, TopologyBuilder, MeshToGrid and the topology-op
consumers already went through: IndexToGrid (node accessor + grid name),
addBlindData (byte-size scratch), GridStats (per-node statistics
scratch) and SignedFloodFill (node counts). Each gains a defaulted
ResourceT template parameter following the established conventions --
classes take a must-outlive resource instance in the constructor, free
functions take the template parameter and route through the per-type
default instance. Every replaced free keeps its position in stream
order via destroy(stream) at the old freeAsync site, and GridStats'
mode-dependent allocation becomes a zero-element buffer when the mode
has no average, so the resource sees identical traffic.
IndexToGrid's nested NodeAccessor is hoisted to the namespace-scope
IndexToGridNodeAccessor so the kernels that name it do not acquire a
spurious ResourceT dependency; a public alias preserves the nested
spelling.
Seam tests with exact allocation counts cover all four (including the
grid-name path, the Extrema zero-allocation path, and free-function
forwarding to the default instance).
Part of #2232 (step 2).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: document GridStats' NodeManager allocation boundary
The temporary NodeManager in GridStats::update allocates through the
dual-space DeviceBuffer, which cannot yet take a resource; make that
visible on update() so users injecting a resource for accounting or
limits know this allocation bypasses it until the single-space handle
work lands.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
---------
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
TopologyBuilder.cuh and PointsToGrid.cuh both undef CALL_CUBS after
their last use so the macro doesn't leak into whatever gets included
next in the same translation unit. MeshToGrid.cuh was missing this.
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* NanoVDB: align cuda::Buffer member names with cuda::buffer
Name the free operation destroy, as cuda::buffer does, and add the
stream-taking overload it also provides. clear stays but only as a
transitional delegate, marked as such: it exists because
GridHandle::reset still calls it, and goes away with the legacy dual
buffers that cuda::Buffer replaces.
Rename setStream to set_stream. Member names cannot be aliased, so
matching the standard spelling is the whole reason the resource concept
kept allocate_async; the same argument applies here. Note in passing that
ours deliberately does not synchronize, matching cuda::buffer's
set_stream_unsynchronized rather than its set_stream, whose own
documentation and implementation disagree (NVIDIA/cccl#10649).
Add swap. The generic std::swap already does the right thing through the
move operations, but both std::vector and rmm::device_buffer provide one.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: allocate TopologyBuilder scratch from an injected resource
Eleven of the builder's buffers are device-only -- nothing reads them on
the host -- yet they were cuda::DeviceBuffer, whose host pointer and
per-device array they never use. Move them to the single-space
cuda::Buffer and give the builder a resource parameter, so its scratch is
injectable like PointsToGrid's already is, and freed by scope rather than
by hand. mProcessedRoot and mData are read on the host and stay dual.
This is not a speedup: DeviceBuffer::init allocates host or device memory
but never both, and these buffers always passed a real device id, so no
cudaMallocHost was on this path to begin with. Measured dilate on 1k/20k/
200k points, and the difference is within run-to-run noise.
Hoist the Data struct out of the class. It does not depend on the
resource, and leaving it nested would give every ResourceT its own
incompatible type for the device functors to name.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: add SyncFromAsync, and give MeshToGrid a resource seam
A stream-ordered resource must also model the synchronous concept, which
means writing four methods where two would do. The synchronous pair is
not a bare delegate -- memory from allocate must be usable on any stream
when it returns, so the null-stream allocation has to be synchronized
first -- and omitting that yields memory which satisfies the concept but
is not actually synchronous. Put it in one place rather than leaving each
author to rediscover it.
The two resources in TestMemoryResource are the first users, and were
already wrong in exactly that way: they provide only the async pair, so
they never modelled is_async_resource. TempPool duck-typed and never
checked, so nothing caught it.
MeshToGrid was the last builder allocating from a hard-wired
DeviceResource, through TempDevicePool. Give it a ResourceT parameter and
thread it into both its TopologyBuilder and its pool. As with Data in the
builder, BoxTrianglePair is hoisted out of the class: it does not depend
on the resource, and leaving it nested would give every ResourceT its own
incompatible type for the device functors to name.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: add ResourceRef and route TempPool's scratch through cuda::Buffer
Buffer holds its resource by value, matching cuda::buffer -- whose model
this completes: in CCCL the ownership semantics are selected by what is
placed in the by-value slot, an owning any_resource or a borrowing
resource_ref. We adopted the slot without the borrowing type, so a
container like TempPool, whose contract is a non-owning pointer to a
possibly stateful resource, had no way to hold a Buffer without copying
that resource and stranding its state. ResourceRef is the missing piece:
a non-owning reference that is itself a resource, so copying the ref
shares the underlying instance. Its async methods exist only when R
models AsyncResource, so a ref over a synchronous resource does not
misreport its tier, and two refs compare equal exactly when they
reference the same resource.
TempPool now keeps its bytes in a Buffer<std::byte, ResourceRef<R>>:
same resource contract, same stream retention, same discard-on-growth
reallocation, but the block is freed by ownership rather than by hand.
The TempPool unit tests, which assert traffic against the caller's own
resource instance, pass unchanged.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: note cuda::Buffer and cuda::BufferView in pendingchanges
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: express PointsToGrid's density search as a loop
The bisection search over voxel size was written as a backward goto,
which made the lifetimes of the buffers it retries over non-lexical.
Rewrite it as while(true) with continue on retry and break on
convergence; the six hand-written frees before the jump are unchanged.
The change is easiest to review with whitespace ignored, since the loop
body re-indents: git diff -w shows 27 changed lines.
d_keys and d_node_count carry results past the loop, so their
declarations move above it, as does the copy event, which was created
inside the retried region on every iteration but destroyed only once at
the end -- each retry leaked the previous handle.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: own PointsToGrid's device arrays with cuda::Buffer
Every device array PointsToGrid allocates is now owned by a
Buffer<T, ResourceRef<ResourceT>> borrowing the injected resource --
members where the pipeline frees them in a later member function than
the one that allocated them (countNodes allocates; processUpperNodes,
processLeafNodes, processPoints and processBBox release), locals where
the lifetime is contained. The raw pointers survive only as views: the
device-visible fields inside mData, and the working pointers the cub
and kernel calls take. Every owner event -- assignment, swap, destroy
-- immediately refreshes its view, and released views are nulled so a
stale use faults instead of reading a freed block. The index ping-pong
becomes a swap of owners across the member/local boundary, replacing
the bare pointer swap whose safety depended on nothing reading the
device copy of d_indx between the two uploads.
The density-search retry keeps its free-before-reallocate order via
explicit destroy calls, so peak device memory is unchanged. The
hand-matched byte sizes at every free site disappear, and the arrays
released one line before scope exit now just leave scope. One behavior
change worth naming: the too-many-points-per-leaf throw previously
leaked the reduction scratch; ownership now releases it during unwind.
Verified against the previous commit with a counting resource over
three shapes (bulk segmented-sort branch, serial per-tile branch, and
the bisection retry engaged): allocation count, free count, and total
bytes are identical. Full CUDA and memory-resource suites unchanged.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: free TopologyBuilder scratch on the caller's stream
Each release site sits in a function that receives the stream, so pass
it to destroy rather than relying on the retained stream matching --
they are the same on every current path, but the explicit form does not
depend on that staying true. Also drop the cudaGetDevice calls whose
result the Buffer conversion left unused.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: borrow TopologyBuilder scratch through ResourceRef
The scratch buffers held their resource by value, so each of the eight
carried its own copy -- fine for the stateless default, wrong for a
stateful resource, whose accounting would be split across copies while
the caller's instance saw nothing. Borrow through ResourceRef instead,
the same reconciliation TempPool uses.
Assert the stream-ordered requirement directly in TopologyBuilder and
MeshToGrid so a synchronous-only resource fails with a diagnostic that
names the builder, not just the pool inside it. Note SyncFromAsync's
synchronize cost on its allocate.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: guard PointsToGrid's copy event with a scope owner
The too-many-points-per-leaf throw unwinds past the event's manual
destroy, leaking the handle. Own it with a small guard so unwinding
releases it, consistent with the buffer ownership in this function.
Also assert the stream-ordered resource requirement on the class, so a
synchronous-only resource fails naming PointsToGrid rather than the
pool inside it.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: run the builders on a synchronous memory resource
The builders require a stream-ordered resource, which a device without
memory-pool support cannot provide through cudaMallocAsync. Two additions
close the gap. MallocResource is a synchronous cudaMalloc/cudaFree
resource that works on any device. AsyncFromSync presents a synchronous
resource as a stream-ordered one -- the mirror of SyncFromAsync, and the
analog of cuda::mr's synchronous_resource_adapter: allocate_async
forwards, since synchronously allocated memory is already valid on every
stream, and deallocate_async synchronizes the stream first, which makes
the synchronous contract's quiescence requirement hold. The serialization
cost of that synchronize is documented at the type, so a synchronous
backend is an explicit caller choice rather than a silent substitution.
The new test drives PointsToGrid end-to-end with an injected
AsyncFromSync over a stateful synchronous resource: every scratch
allocation routes through cudaMalloc and is freed through the caller's
instance, never touching stream-ordered allocation -- which is exactly
what a pool-less device requires of the scratch path.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: assert the scratch alignment TopologyBuilder relies on
The byte scratch is reinterpreted as word-sized types, which is valid
for every resource whose DEFAULT_ALIGNMENT is at least word alignment --
all CUDA allocation paths give 256 -- but nothing said so. Assert it, so
a custom resource with a weaker guarantee fails at compile time instead
of misaligning on the device.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: terminate a bare cudaCheckError with a semicolon
Legal without one -- the macro expands to a braced block -- but every
other use in the file spells it as a statement.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: make AsyncFromSync's null deallocation a no-op
Null-free is a no-op everywhere else, so there is no quiescence to
establish and no reason to synchronize the stream for it. Also correct
the sync-resource test's comment: the grid handle's output buffer is
allocated through BufferT, not the injected resource, so "every builder
allocation" overstated what the test demonstrates -- it is the scratch
allocations that never touch stream-ordered allocation.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: set_stream matches cuda::buffer's contract, not a divergence
cuda::buffer's set_stream is deliberately non-synchronizing -- its
documented synchronization is a stale note left behind when the
synchronization was removed -- so our non-synchronizing set_stream
matches it in both name and contract, and the comment claiming a
divergence was wrong.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: mark cuda::Buffer::clear deprecated in the documentation
Documentation-level only: the [[deprecated]] attribute would warn from
GridHandle::reset and NodeManager::reset, template members in our own
headers that must keep calling clear() until every buffer type provides
destroy() -- an unactionable diagnostic for callers of reset(), and a
build break under -Werror. The attribute lands when those callers
migrate.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: exercise the concept-required members of the test resources
The trait checks detect allocate/deallocate through unevaluated contexts,
which never odr-use them, so nvcc warned that every synchronous pair and
stub member was declared but never referenced (#177-D). The attribute
route is closed -- nvcc's front end ignores [[maybe_unused]] for this
diagnostic -- so reference them the honest way: a test that verifies the
synchronous halves of the counting and stream-recording doubles behave
like their stream-ordered halves, and one that pins the trait probes'
stub behavior. The TU now builds with no warnings at all.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
---------
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
* nanovdb python: Phase 0 foundation for C++ API mirror (#2209)
* nanovdb python: Phase 0 foundation for C++ API mirror
First slice of the Python bindings restructure tracked in #2208 and
laid out in nanovdb-python-plan.md. Phase 0 is the mechanical
groundwork the rest of the plan builds on:
- BuildTypes.def: single X-macro list of currently-bound BuildT types
(scalar / vector / point / sampleable). NanoVDBModule.cc, PyMath.cc,
PySampleFromVoxels.cc, PyCreateNanoGrid.cc, PyTools.cc, PyGridHandle.h
and cuda/PyDeviceGridHandle.cu now drive their per-type instantiations
from this one file. Adding a BuildT in Phase 2 becomes one line.
- CMakeLists.txt: under SKBUILD, ship the nanovdb/ headers inside the
Python wheel at nanovdb/include/nanovdb/ so downstream extension
authors can compile against the same headers the wheel was built with.
Also wire nanobind_add_stub so a nanovdb.pyi (and py.typed marker) are
emitted into the wheel for IDE / type-checker support; gated behind
NANOVDB_BUILD_PYTHON_STUBS and silently skipped on older nanobind.
- __init__.py: add nanovdb.get_include() returning the bundled include
dir. Also fix the Windows DLL shim, which referenced an undefined
`directory` variable instead of the local `openvdb_dll_directory`.
- Relocate the batched sampleFromVoxels CUDA kernel binding from the
phantom nanovdb.math.cuda submodule (which has no C++ counterpart) to
nanovdb.tools.cuda, alongside the existing signedFloodFill and
pointsToRGBA8Grid kernels. The math.cuda submodule is no longer
registered. TestNanoVDB.py updated to match.
- Pre-existing observable bug fixes:
* GridHandle.__bool__ returned None (lambda was missing return);
it now returns !handle.empty() as intended.
* Enable __repr__ on GridType, GridClass and io::Codec via the
nanovdb::toStr / strlen<> helpers — previously commented out.
Build + pytest verified locally (CPU-only): 28 tests pass; the single
pre-existing test_read_write_grid BLOSC failure is unrelated (test
doesn't wrap the optional codec call in try/except like its sibling
does).
Part of #2208.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: align Phase 0 X-macros with codingstyle.txt
Audit of the new BuildTypes.def + consumers against
nanovdb/nanovdb/docs/codingstyle.txt:
- Rename the X-macro family NVDB_PY_FOR_EACH_*_BUILDT to
NANOVDB_PY_FOR_EACH_*_BUILDT and the helper sentinels
NVDB_PY_LOCAL_DEFINED_* to NANOVDB_PY_LOCAL_DEFINED_* to match the
established NANOVDB_ macro prefix used everywhere else in the codebase
(NANOVDB_USE_CUDA, NANOVDB_BUILD_PYTHON_MODULE, NANOVDB_HOSTDEV, ...).
- Bring lines under the 100-column limit:
* Wrap the long VECTOR_BUILDT lines in BuildTypes.def across two
lines each (cleaner alignment, no behavior change).
* Rename the consumer macro parameter DeviceHandleMethod to
DeviceMethod so the #define line itself fits (was 102 cols).
* Reformat the GridHandle __bool__ lambda onto three lines instead
of one 119-col line (made worse by the `return` fix in the
previous commit).
- Add a top-of-file justification block in BuildTypes.def explaining why
this file is a deliberate exception to the codingstyle "avoid macro
functions" rule (templates can't emit top-level declarations and
explicit instantiations across translation units from a single
canonical list).
No functional change. Rebuild + pytest_nanovdb is identical to the
previous commit: 28 tests pass, 8 CUDA skips, 1 pre-existing BLOSC
failure.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: only generate .pyi stubs under SKBUILD by default
Fixes the macOS CI failure on #2209. The previous commit defaulted
NANOVDB_BUILD_PYTHON_STUBS to ON, which made the nanovdb_python_stub
target fire on every in-tree CI build. The macOS GitHub Actions runner
wraps the stubgen invocation with
cmake -E env DYLD_INSERT_LIBRARIES=.../libclang_rt.tsan_osx_dynamic.dylib:
.../libclang_rt.asan_osx_dynamic.dylib:
.../libclang_rt.ubsan_osx_dynamic.dylib
ASAN_OPTIONS=detect_leaks=0 python stubgen.py -m nanovdb ...
(this preamble is added by the runner / CMake env wrapper — nothing in
our nanobind_add_stub call sets it, and the .so itself is not built with
-fsanitize). ThreadSanitizer can only install its interceptors at
process start; Python loads first and then dlopen's the compiled .so,
which TSan considers "too late" and aborts:
==49907==ERROR: Interceptors are not working. This may be because
ThreadSanitizer is loaded too late (e.g. via dlopen).
So the macOS build target failed with exit 2 after the .so itself built
fine. The other matrix legs (linux-nanovdb Debug/Release for clang/gcc)
all passed.
Stubs are only useful to wheel consumers — they ship next to the .so in
the installed package layout under SKBUILD. The in-source OpenVDB CI
build never consumes them, so making stub generation default ON only
when SKBUILD is set keeps the wheel build path unchanged and stops the
macOS CI from invoking stubgen under the sanitizer wrapper. The user can
still force generation with -DNANOVDB_BUILD_PYTHON_STUBS=ON for local
dev builds where it's useful.
Verified locally:
- in-tree config (SKBUILD unset): no nanovdb_python_stub target
defined; `make nanovdb_python_stub` errors with "no rule". Matches
desired CI behavior.
- SKBUILD=ON config: stub target exists, builds, emits nanovdb.pyi
and py.typed into the install layout (28 kB stub file with all the
expected GridHandle/GridType/GridClass symbols).
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: Phase 1 — polymorphic Grid, GridMetaData, blind data, PointAccessor (#2210)
* nanovdb python: Phase 1 — polymorphic Grid, GridMetaData, blind data, PointAccessor
Second slice of the Python bindings restructure tracked in #2208 and
laid out in nanovdb-python-plan.md. Phase 1 bundles the three sub-todos
(1a polymorphic API, 1b type-erased introspection + blind data, 1c
PointAccessor + handle utilities) into a single change on top of the
Phase 0 X-macro foundation.
API changes (pre-1.0 breaks called out in the plan):
- Polymorphic accessors. handle.grid(n=0) and handle.deviceGrid(n=0)
replace the typed handle.floatGrid()/doubleGrid()/int32Grid()/
vec3fGrid()/rgba8Grid() and their device equivalents. Dispatch is
driven by gridType(n) through a switch generated from BuildTypes.def;
unbound BuildTs route to None rather than throwing.
- Grid base class rename. The Python class previously bound as
nanovdb.GridData is now nanovdb.Grid (matching the C++ user-facing
class name Grid<TreeT>). All typed grid classes (FloatGrid, ...)
inherit from Grid.
- Base-class method lift. version/gridSize/gridIndex/gridCount/voxelSize/
map/gridType/gridClass/checksum/isLevelSet/isFogVolume/.../hasMinMax/
hasBBox/.../isBreadthFirst/shortGridName move from per-BuildT
defineNanoGrid<T> up to defineGrid via lambdas that read GridData data
members directly. defineNanoGrid<T> now only binds getAccessor,
activeVoxelCount, and isSequential — the BuildT-dependent slice.
Additive surface:
- GridMetaData. Bound as nanovdb.GridMetaData with constructor from a
Grid and the full read-only accessor surface (gridType, gridClass,
shortGridName, gridSize/Index/Count, map, worldBBox, indexBBox,
voxelSize, blindDataCount, activeVoxelCount, activeTileCount(level),
nodeCount(level), checksum, version, isValid, isLevelSet/...,
hasMinMax/..., isBreadthFirst, rootTableSize, isEmpty). Type-erased
introspector — answer "what's in this buffer?" without knowing BuildT.
- Blind data API on Grid. blindDataCount, blindMetaData(n),
findBlindData(name), findBlindDataForSemantic(sem), getBlindData(n).
getBlindData returns a zero-copy NumPy view typed by mDataType (Float
-> 1D float32, Vec3f -> (N, 3) float32, RGBA8 -> (N, 4) uint8, etc.).
Out-of-range and unknown-type paths return None / fall back to a flat
uint8 byte view.
- Enums: GridBlindDataClass (Unknown/IndexArray/AttributeArray/GridName/
ChannelArray/End) and GridBlindDataSemantic (Unknown/PointPosition/
PointColor/PointNormal/PointRadius/PointVelocity/PointId/WorldCoords/
GridCoords/VoxelCoords/LevelSet/FogVolume/Staggered/End). Bound as
nb::enum_ with .export_values() so the names are also top-level
attributes of the module.
- GridBlindMetaData struct. Read-only fields valueCount/valueSize/
semantic/dataClass/dataType, name() accessor, isValid(),
blindDataSize().
- PointAccessor variants. nanovdb.PointIndexAccessor (uint32 indices,
used by PointIndex grids) and nanovdb.PointDataAccessor (Vec3f
positions, used by PointData grids). Methods gridPoints(),
leafPoints(ijk), voxelPoints(ijk) each return a zero-copy NumPy view
onto the underlying blind-data buffer, anchored to the accessor
lifetime via keep_alive.
- GridHandle utilities. handle.copy() does a deep copy into a freshly
allocated buffer of the same buffer type. Module-scope splitGrids(h)
-> list[GridHandle] and mergeGrids(handles) -> GridHandle are
registered for both host and device handles via the existing
defineGridHandleUtilities<BufferT> template (nanobind merges them as
an overload set).
Mechanical X-macro changes:
- BuildTypes.def gains a GridTypeEnum column on each row so the
polymorphic dispatch in pyHostGrid/pyDeviceGrid can `case
nanovdb::GridType::<GridTypeEnum>:` on it. The Point row maps to
GridType::PointIndex (there is no GridType::Point).
- HandleMethod/DeviceMethod columns dropped — the typed handle.fooGrid()
accessors no longer exist.
NB_MODULE bind order:
- defineCheckMode + defineChecksum now bind BEFORE defineGrid because
Grid.checksum() returns Checksum by value (registration must precede
use).
- defineGridBlindData binds BEFORE defineGrid for the same reason
(Grid.findBlindDataForSemantic / blindMetaData reference the new enum
and class in their signatures).
Tests (TestNanoVDB.py): all typed-accessor call sites rewritten to
handle.grid(i)/handle.deviceGrid(i). New test classes cover the new
surface — TestPolymorphicGridAccess, TestGridBase, TestGridMetaData,
TestBlindDataEmpty, TestSplitMergeCopy — 11 new tests; the full suite
is now 48 tests, 39 pass on a minimal CPU build (8 CUDA skip, 1
pre-existing test_read_write_grid BLOSC failure unrelated to this PR).
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: don't expose splitGrids/mergeGrids for DeviceGridHandle
Caught during a real-CUDA verification pass of #2210: calling
nanovdb.mergeGrids([device_h1, device_h2]) raised std::bad_cast on
sm_120 (Blackwell, CUDA 13.2). Both the host and device overloads of
splitGrids/mergeGrids take nb::list, so nanobind's overload resolution
can't disambiguate by element type — it picks the first match and the
inner nb::cast<HostHandle&&>(device_h) fails.
The host-only variant is what the Phase 1 plan calls for. A properly
typed device variant (with its own name, or strongly-typed
std::vector<HandleT> args via nanobind/stl/vector.h) can land later if
it's actually needed. handle.copy() on a DeviceGridHandle continues to
work because copy() is a regular method, no overload resolution
involved.
Full CUDA pytest now reports 46/48 (the 2 failures are the pre-existing
test_read_write_grid BLOSC bug, host + device variants — both call
writeGrid(..., Codec.BLOSC) without try/except, identical to master).
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: address Copilot review on #2210
Three concrete fixes from Copilot's review of the Phase 1 PR
(https://github.com/AcademySoftwareFoundation/openvdb/pull/2210):
1) mergeGrids no longer consumes its input handles.
The previous implementation built a std::vector<HandleT> via
nb::cast<HandleT&&>(h), which move-constructs the C++ GridHandle
out of the Python wrapper — leaving caller's h1/h2 silently
emptied (gridCount went 1 -> 0, size went non-zero -> 0).
Reproduced before fix; locked into a regression test
(TestSplitMergeCopy.test_merge_does_not_consume_inputs).
Rewrote the binding to read each handle by const reference and
inline the merge concat directly. The nanovdb::mergeGrids C++
helper's signature requires a std::vector<GridHandle> (a
move-only type), so reusing it from Python without moving from
the inputs would have meant deep-copying each handle twice; the
inlined version is ~15 lines and does one memcpy per source grid
with tools::updateGridCount fixing up the per-grid header.
2) getBlindData validates mValueSize against the implied dtype/shape
before building a typed NumPy view.
A blind-data channel with mDataType=Float but mValueSize != 4
(corruption, version mismatch, or an unknown variant of a known
tag) would previously be exposed as `count` float32 elements —
i.e. count*4 bytes — even though the underlying region is only
count*mValueSize bytes. That overruns the channel and returns a
view onto unrelated bytes.
Added a `valueSize == sizeof(...)` (or `dim*sizeof(scalar)` for
vector cases) check on every handled GridType. On mismatch the
binding falls back to a raw uint8 byte view of mValueCount *
mValueSize, which is by definition the actual byte extent and
therefore always safe.
3) GridMetaData ctor + safeCast guard against invalid grids before
calling into NanoVDB, where NANOVDB_ASSERT(gridData->isValid())
would abort debug builds and undefined-behave in release.
nanobind's type system already rejects Python None at the bind-
site (None can't bind to const GridData*), so the literal
"GridMetaData(None)" case Copilot called out is a TypeError
today — but the broader concern (an otherwise-valid Grid object
wrapping a corrupted buffer) is real.
__init__ now does an explicit `gd == nullptr || !gd->isValid()`
check and raises nb::value_error with a descriptive message
before calling into nanovdb::GridMetaData. safeCast does the
same and returns False on bad input, matching the spirit of
"is this safe to cast?".
New tests: TestGridMetaDataGuards covers the rejection paths and
the still-works happy path.
Build + test verified locally on both CPU (52 tests, 43 pass, 8 CUDA
skip, 1 pre-existing BLOSC) and full CUDA (52 tests, 50 pass, 2 pre-
existing BLOSC host+device).
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: Phase 2 — broaden BuildT coverage to 24 grid types (#2211)
* nanovdb python: Phase 2 — broaden BuildT coverage to 24 grid types
Third slice of the Python bindings restructure tracked in #2208 and
laid out in nanovdb-python-plan.md. Phase 2 expands the bound BuildT
list from the six surfaced in Phase 0/1 (float, double, int32_t, Vec3f,
Rgba8, Point) to seventeen new types, riding the Phase 0 X-macro so
each addition is a single row in BuildTypes.def plus a polymorphic
dispatch arm.
New BuildTs by category:
- SCALAR (+4): int16_t, int64_t, uint8_t, uint32_t.
Bound exactly like the existing float/double/int32_t — full
defineScalarAccessor (with setVoxel) plus defineNodeInfo.
Class names follow the GridType enum: Int16Grid, Int64Grid,
UInt8Grid, UInt32Grid.
- VECTOR (+5): Vec3d, Vec4f, Vec4d, Vec3u8, Vec3u16.
Same shape as the existing Vec3f / Rgba8 — defineVectorAccessor
with setVoxel, no NodeInfo. Accessor names use the consistent
"<Suffix>ReadAccessor" form (Vec3dReadAccessor, ...); only the
legacy Vec3fReadVectorAccessor / RGBA8ReadAccessor names from
Phase 0 stay as-is.
- READONLY (new category, +8): bool, Fp4, Fp8, Fp16, FpN,
ValueIndex, ValueOnIndex, ValueMask. These all have
nanovdb::BuildTraits<T>::is_special == true, which means the C++
SetVoxel<T> specialization static_asserts and won't compile.
They get a bare defineAccessor<T> binding — getValue() only,
no setVoxel, no NodeInfo. Class names: BooleanGrid, Fp4Grid /
Fp8Grid / Fp16Grid / FpNGrid (quantized — getValue returns
float), IndexGrid / OnIndexGrid (getValue returns uint64), and
MaskGrid (getValue returns bool).
The accessor's value type now resolves through
nanovdb::BuildToValueMap<BuildT>::Type rather than
DefaultReadAccessor<BuildT>::ValueType. For ordinary types they're
identical, but for Half / Fp* the accessor decodes to float on read,
for ValueIndex / OnIndex it returns uint64, and for ValueMask / bool
it returns bool — the bound probeValue() out-parameter and the
Python-side return type both want the decoded form. (Without this
change, probeValue<Half>'s instantiation chain mismatches its own
ProbeValue<Half>::ValueT = float.)
nanovdb::Half is intentionally NOT bound. The source declares it as
`class Half{};` (an empty placeholder for IEEE 754 half-precision)
and the C++ ProbeValue<Half> chain is inconsistent — leaf storage
carries Half but ProbeValue expects float, so the template doesn't
instantiate. When the upstream Half implementation lands we can add
it via the same X-macro path.
Polymorphic dispatch in pyHostGrid / pyDeviceGrid gains a fourth
arm (NANOVDB_PY_FOR_EACH_READONLY_BUILDT) covering all eight new
GridType enumerators (Boolean, Fp4/Fp8/Fp16/FpN, Index, OnIndex,
Mask). handle.grid(n) / handle.deviceGrid(n) now return the right
typed subclass for these too.
Tests: a single new TestPhase2BuildTCoverage class with 5 methods
verifies every new BuildT registered, every accessor surface matches
its category (scalars have setVoxel + getNodeInfo; vectors have
setVoxel; read-only have neither), and all the new typed grids
inherit from the polymorphic Grid base. We can't host-construct
Int16Grid / Fp4Grid / IndexGrid / etc. yet because the C++
create*Grid factories for those types land in Phase 5 — but the
registration and the dispatch surface are locked in.
Build + test verified locally:
- CPU (no CUDA, no OpenVDB, no BLOSC): 57 tests, 48 pass, 8 CUDA
skip, 1 pre-existing test_read_write_grid BLOSC failure unrelated
to this PR.
- CUDA sm_120 (Blackwell, CUDA 13.2): 57 tests, 55 pass, 2 pre-
existing BLOSC failures (host + device variants).
No new lines over 100 cols.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: address Copilot review on #2211
Three concrete fixes from Copilot's review of the Phase 2 PR
(https://github.com/AcademySoftwareFoundation/openvdb/pull/2211):
1) Bind GridType.UInt8 in the Python enum.
Real bug — a Phase 0 oversight that was harmless until Phase 2
surfaced UInt8Grid. The C++ enumerator nanovdb::GridType::UInt8
existed (value 26), and the polymorphic dispatch routes it
correctly internally, but the enum binding was missing the
`.value("UInt8", GridType::UInt8)` line. Python users therefore
couldn't write `handle.gridType(n) == nanovdb.GridType.UInt8` to
discriminate UInt8 grids.
New regression test TestPhase2BuildTCoverage.
test_all_grid_type_enums_reachable walks every BuildT we bind
and confirms its GridType enumerator is reachable from Python,
so the next oversight gets caught at test time.
2) Update pyHostGrid() docstring in PyGridHandle.h. Was claiming
"Boolean, Half, Fp16 land in Phase 2" as examples of types that
weren't yet Python-visible — those examples are now stale (Phase
2 binds Boolean and Fp16; Half stays unbound but for a different
reason). Reworded to point at BuildTypes.def as the source of
truth so the docstring can't go stale again as Phase 5+ adds
more types.
3) Disambiguate "bool for ValueMask and bool" wording in the
READONLY macro doc comment in BuildTypes.def. The second "bool"
referred to the literal BuildT=bool grid; the phrasing read as
redundant. Now: "bool for ValueMask and for the BuildT=bool
(Boolean) grid". Same fix also enumerates the quantized types
explicitly (Fp4/Fp8/Fp16/FpN) instead of writing "Fp*".
Build + test on CUDA sm_120 with BLOSC + ZLIB: 58/58 pass.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: Phase 3 — tree / nodes / NodeManager / leaf_values (#2212)
* nanovdb python: Phase 3 — tree / nodes / NodeManager / leaf_values
Fourth slice of the Python bindings restructure tracked in #2208 and
laid out in nanovdb-python-plan.md. Phase 3 makes the tree itself
walkable from Python: every BuildT now has a bound NanoTree, Root,
Upper / Lower internal node, Leaf, and a host-side NodeManager. Where
the leaf actually carries a contiguous T mValues[512] (regular
scalar BuildTs) we also expose zero-copy NumPy views, including a
high-level grid.leaf_values() bulk extractor.
Per BuildT we now register six new classes:
- <Suffix>Leaf — origin, bbox, dim, voxelCount, memUsage, flags,
isActive(ijk|n), getValue(offset|ijk), getFirstValue, getLastValue,
minimum / maximum / average / stdDeviation, valueMask, probeValue,
and (for arithmetic non-special ValueTs) values() returning a
zero-copy (512,) NumPy view of the leaf's mValues array.
- <Suffix>Upper, <Suffix>Lower — origin, bbox, dim, memUsage,
minimum / maximum / average / stdDeviation, valueMask, childMask,
getValue, getFirstValue, getLastValue, isActive, probeValue. The
two internal node levels share a single defineInternalNodeBase
helper since their C++ APIs are identical.
- <Suffix>Root — background, tileCount, getTableSize, isEmpty, bbox,
minimum / maximum / average / stdDeviation, memUsage, getValue,
isActive, probeValue.
- <Suffix>Tree — root, background, activeVoxelCount,
activeTileCount(level), nodeCount(level), totalNodeCount,
memUsage, getValue, isActive, probeValue, extrema() (returns
(min, max) tuple), getFirstLeaf / getFirstLower / getFirstUpper.
Grid.tree() is bound on NanoGrid<T> and returns the typed tree
as reference_internal.
- <Suffix>NodeManager — isLinear, memUsage, nodeCount(level),
leafCount, lowerCount, upperCount, leaf(i), lower(i), upper(i)
returning typed node refs. Constructed via the module-scope
nanovdb.createNodeManager(grid) which polymorphically picks the
right BuildT and returns a NodeManagerHandle. Handle exposes
size(), __bool__(), and mgr() — which itself dispatches by stored
gridType to return the right typed NodeManager.
- grid.leaf_values() bulk extractor — for arithmetic non-special
BuildTs (float, double, Int16/32/64, UInt8/UInt32), walks the
breadth-first leaf array and returns a strided zero-copy
(N_leaves, 512) NumPy view. Stride between leaves is sizeof(LeafT)
/ sizeof(ValueT), reflecting the leaf header between value blocks.
Throws ValueError on non-breadth-first grids (createNanoGrid
produces breadth-first by default).
Mechanical bits:
- All six bindings driven from BuildTypes.def via the existing
X-macro. New file PyTree.h holds the templated definitions; PyTree.cc
holds the non-templated NodeManagerHandle + createNodeManager bindings
(the latter dispatches over every BuildT via the same X-macro).
- Tree / node classes are registered BEFORE defineNanoGrid because
NanoGrid<T>.tree() returns NanoTree<T>& and nanobind needs the
return type registered first.
Skipped / deferred:
- LeafT::variance() is NOT bound. NanoVDB.h line 4388 reads
`Pow2(DataType::getDev())` unqualified, which fails ADL for
non-float ValueTs (ValueIndex / ValueMask). Same for InternalNode.
Users can compute variance from stdDeviation() in Python.
- VoxelBlockManager (OnIndexGrid-specific) deferred to a follow-up;
surface is sizeable enough to merit its own PR.
- Vector leaf values() (Vec3f / Vec3d / Vec4f / Vec4d / Vec3u8 /
Vec3u16 / Rgba8) deferred — these need a flattened
(count, dim) component view since nanobind ndarray<Vec3f> isn't
well-formed. A future PR can add float[N, 512, 3] views.
- Tree iterators (beginValueOn etc.) deferred — the bulk leaf_values
view + per-leaf valueMask covers the most common use case (mask
the bulk array and you have your active values).
Verified locally (BLOSC + ZLIB on so the optional-codec tests pass):
- CPU build: 58 tests + 8 new TestPhase3TreeNodes — **all 66 pass**.
- CUDA sm_120 build (RTX PRO 6000 Blackwell, CUDA 13.2):
**all 66 pass**.
- numpy-backed shape/dtype assertions in the new tests confirm the
per-leaf values() (512,) float32 and bulk leaf_values()
(N_leaves, 512) float32 views. NodeManager.leaf(i).values() is
byte-identical to tree.getFirstLeaf().values() for i==0.
No new lines over 100 cols.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: drop phase numbers + reviewer markers from tests and comments
User-facing strings in tests/comments that reference in-flight project
history (phase numbers, reviewer names) lose their meaning once the
project merges. Drop them and reword the surrounding text so each name
and comment self-describes the feature being tested.
- Renamed TestPhase2BuildTCoverage -> TestBuildTRegistrations.
- Renamed TestPhase3TreeNodes -> TestTreeNodeWalking.
- Reworded docstrings on TestPolymorphicGridAccess, TestGridBase,
TestGridMetaData, TestBlindDataEmpty, TestSplitMergeCopy,
TestGridMetaDataGuards to describe the feature instead of the phase.
- Inline comments referencing "Phase 0", "Phase 1a.3", "Copilot review",
etc. reworded into prose about the actual behavior being asserted or
the underlying bug being regression-tested.
- Same cleanup on the cuda/PyDeviceGridHandle.cu note about why we don't
register splitGrids/mergeGrids on DeviceBuffer.
No behavior change. 66/66 tests still pass with BLOSC+ZLIB on.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: fix Clang Tree.nodeCount overload + drop exception-driven createNodeManager dispatch
Two real issues caught in CI / review on #2212:
1) Tree.nodeCount(int) binding doesn't compile on Clang.
PyTree.h used `nb::overload_cast<int>(&TreeT::nodeCount, nb::const_)`
to select the non-templated `uint32_t nodeCount(int) const` overload
on nanovdb::Tree<RootT>. Tree<RootT> also has a templated
`template<typename NodeT> uint32_t nodeCount() const` overload.
GCC accepted the overload_cast under SFINAE rules; Clang (used by
the linux-nanovdb:cxx:clang++-Debug CI leg) rejects it with
`no matching function for call to object of type
'const detail::overload_cast_impl<int>'`, repeated once per
BuildT instantiation.
Replaced with a direct static_cast to the function pointer type,
which is unambiguous to both compilers:
static_cast<uint32_t (TreeT::*)(int) const>(&TreeT::nodeCount)
2) createNodeManager dispatch was exception-driven.
The previous binding tried nb::cast<NanoGrid<T>&>(py_grid) for every
BuildT and caught nb::cast_error on each mismatch — so a single call
to nanovdb.createNodeManager(grid) threw and caught 22 cast_error
exceptions before landing on the matching BuildT. Replaced with an
nb::isinstance<GridT>(py_grid) pre-check so the cast is only ever
attempted on the matching BuildT.
Both linux-nanovdb:cxx:clang++-Debug should now build, and
createNodeManager() no longer pays per-call exception overhead.
Verified locally on CPU and CUDA sm_120 builds: 66/66 tests pass with
BLOSC + ZLIB on.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: bounds-check + lifetime fixes on #2212
Addresses Copilot review notes on PR #2212. Two real classes of bug.
(1) Out-of-range arguments fell through into raw memory access.
Several entry points on Leaf, Tree, and NodeManager rely on
NANOVDB_ASSERT in the underlying C++ to catch invalid indices. That
assertion is a no-op in release builds, so passing an OOB index from
Python would read off the end of mValueMask / mValues / mNodeOffset[]
arrays. Wrapped each with an explicit range check that raises a
Python IndexError or ValueError:
- Leaf.isActive(n) n must be < voxelCount() (512)
- Leaf.getValue(offset) offset must be < voxelCount()
- Tree.activeTileCount(level) level must be 1, 2, or 3
- Tree.nodeCount(level) level must be 0, 1, or 2
- NodeManager.nodeCount(L) same as Tree
- NodeManager.leaf(i) i must be < leafCount()
- NodeManager.lower(i) i must be < lowerCount()
- NodeManager.upper(i) i must be < upperCount()
(2) Returned pointers / NumPy views did not actually keep their
backing buffers alive.
The pattern `nb::cast(value, nb::rv_policy::reference, parent)` was
used at multiple sites under the assumption that the third argument
established a Python-level keep_alive linkage between the returned
object and `parent`. It doesn't — rv_policy::reference is "no
ownership, no keep_alive" by definition; the parent argument is only
a hint to the cast machinery, not a lifetime guarantee.
As a result, expressions that drop the intermediate handle would
silently free the underlying buffer:
g = nanovdb.tools.createFogVolumeSphere(name='probe').grid()
g.gridName() # SEGFAULT — handle was GC'd
vals = (nanovdb.tools.createFogVolumeSphere()
.grid().tree().getFirstLeaf().values())
vals.sum() # SEGFAULT
nm = nanovdb.createNodeManager(
nanovdb.tools.createFogVolumeSphere().grid()).mgr()
nm.leaf(0) # SEGFAULT — NodeManager holds raw ptr to grid
Added explicit `nb::keep_alive<0, 1>()` to the .def for every
affected site so the returned value keeps its parent alive:
- GridHandle.grid(n) (PyGridHandle.h)
- DeviceGridHandle.deviceGrid(n) (cuda/PyDeviceGridHandle.cu)
- NodeManagerHandle.mgr() (PyTree.cc)
- Grid.getBlindData(n) (NanoVDBModule.cc)
- Leaf.values() (PyTree.h)
- Grid.leaf_values() (PyTree.h)
- PointAccessor.gridPoints / leafPoints / voxelPoints
(NanoVDBModule.cc)
- nanovdb.createNodeManager(grid) keep arg 1 (grid) alive as long
as the returned NodeManagerHandle lives — because the underlying
NodeManager stores a raw pointer to the grid.
These were pre-existing bugs introduced in Phase 1 (pyHostGrid) and
inherited by every subsequent zero-copy view; Phase 3 surfaced more
of them via the new Tree/Leaf/NodeManager bindings.
Two new test classes lock the fixes in:
- TestBoundsChecks asserts each guarded entry point raises the
correct exception type at every boundary and still accepts in-
range inputs.
- TestZeroCopyViewLifetimes invokes each fixed binding in the
chained-temporary form (handle.grid().tree().getFirstLeaf()
.values(), createNodeManager(temp).mgr().leaf(0).values(), ...),
runs gc.collect(), then touches the returned value. Pre-fix
these segfaulted; now they pass cleanly.
CPU + CUDA sm_120 (BLOSC + ZLIB on): 75/75 pass.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: address Copilot follow-up review on #2212
Three further Copilot notes on the bounds-check/lifetime commit
(f6735717), all valid.
(1) PyTree.h relied on transitive includes for std::is_arithmetic_v /
std::enable_if. Works on GCC + libstdc++ because <type_traits> is
pulled in by <nanobind/nanobind.h> -> standard headers, but MSVC /
libc++ may break the chain. Added an explicit
#include <type_traits>
near the top.
(2) grid.leaf_values() returned None for empty grids (nLeaves == 0 or
getFirstLeaf() == nullptr) while its docstring promised an
(N_leaves, 512) NumPy view. Callers had to special-case the None
sentinel before iterating. Now returns an empty (0, 512) ndarray
of the right dtype, so the contract reads cleanly:
for row in grid.leaf_values(): ...
works on every grid, empty or not. When nLeaves == 0 we pass a
dummy non-null aligned pointer (the grid itself) to nb::ndarray so
nanobind has a valid base for the empty array — no data is read
since the leading shape is 0. Docstring updated to call out the
empty-grid behavior explicitly.
(3) The ValueError raised on a non-breadth-first grid said
"rebuild via tools::createNanoGrid(...)"
which reads like a C++ symbol. Reworded to the Python-API form
"rebuild via nanovdb.tools.createNanoGrid(...)"
so Python users see a Python entry point.
New test method TestTreeNodeWalking.test_bulk_leaf_values_empty_grid_returns_empty_array
constructs a grid with an empty bbox (nLeaves == 0) and asserts
leaf_values() is an (0, 512) float32 NumPy array (not None).
CPU + CUDA sm_120 (BLOSC + ZLIB on): 76/76 pass.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: VoxelBlockManager (host) — Phase 3 follow-up (#2213)
* nanovdb python: VoxelBlockManager (host) — Phase 3 follow-up
Phase 3 deferred the VoxelBlockManager surface; this picks it up before
moving on to Phase 4. Adds host-side bindings for everything under
nanovdb::tools::VoxelBlockManager*, plus a minimal createOnIndexGrid
scaffold needed to build OnIndex grids from Python (the broader
createNanoGrid<SrcGridT, DstBuildT> binding lands in Phase 5).
New module surface, all under nanovdb.tools:
- VoxelBlockManagerHandle (host) — owns the firstLeafID + jumpMap
metadata buffers; exposes blockCount(), firstOffset(), lastOffset(),
reset(), __bool__. Buffers exposed as zero-copy NumPy views:
firstLeafID() -> (blockCount,) uint32
jumpMap(jump_map_length=1) -> (blockCount, jump_map_length) uint64
jumpMap takes a jump_map_length argument because the value depends on
log2_block_width (= 1 << (log2_block_width - 6)) and isn't stored on
the handle. handle.decodeBlock(grid, i, log2_block_width=6) is a
convenience method that slices firstLeafID / jumpMap for block i and
calls decodeInverseMaps internally.
- buildVoxelBlockManager(grid, log2_block_width=6, first_offset=0,
last_offset=0, n_blocks=0) — runtime switch over Log2BlockWidth ∈
{6, 7, 8, 9} (BlockWidth 64/128/256/512) dispatching to the right
template instantiation. Rejects non-OnIndex grids with TypeError;
rejects out-of-range log2_block_width with ValueError.
- decodeInverseMaps(grid, first_leaf_id, jump_map, block_first_offset,
log2_block_width=6) — free function. jump_map is a uint64 NumPy
array of length JumpMapLength (= 1 << (log2_block_width - 6));
returns (leaf_index, voxel_offset) freshly-allocated NumPy arrays
(uint32, uint16) of length BlockWidth.
- createOnIndexGrid(src_grid, channels=0, include_stats=True,
include_tiles=True, verbose=0) — minimal test-scaffold factory
that binds tools::createNanoGrid<SrcGridT, ValueOnIndex>. Accepts
FloatGrid, DoubleGrid, Int32Grid, Vec3fGrid sources. Required for
end-to-end VBM testing since no other path produces an OnIndex
grid from Python today. The full createNanoGrid<SrcGridT, DstBuildT>
surface remains scoped to Phase 5.
Defensive checks:
- decodeBlock validates firstLeafID[block_index] is in [0, leafCount)
before passing it into the C++ decodeInverseMaps. The underlying
NanoVDB algorithm doesn't always initialize firstLeafID — blocks
that no leaf's iteration sweep reaches (e.g. on tile-compressed
OnIndex grids where some sequential offsets correspond to tile
values rather than leaf voxels) are left with uninitialized memory.
Without the guard, decodeInverseMaps would read
tree.getFirstNode<0>()[garbage] and crash; the guard converts that
into a Python ValueError with a clear message pointing at the
workaround (build the source grid voxel-by-voxel via build::Grid
so it stays uncompressed).
- All entry points reject out-of-range indices, levels, and grid
build types up front (TypeError / ValueError / IndexError).
Tests under TestVoxelBlockManager exercise:
- createOnIndexGrid produces an OnIndex / IndexGrid / sequential grid
- Buffer shapes and dtypes (firstLeafID, jumpMap default + reshape)
- decodeBlock(0) returns (uint32, uint16) arrays of length BlockWidth
- decodeInverseMaps free function agrees with handle.decodeBlock
- Out-of-range block_index raises IndexError
- Out-of-range log2_block_width (5, 10) raises ValueError
- Non-OnIndex grid argument raises TypeError
- createOnIndexGrid(None) raises TypeError
End-to-end decode verification across every block is intentionally
deferred until Phase 4's build::Grid bindings land — that's the only
host-side path to construct a tile-free OnIndex grid where every
block's firstLeafID is reliably initialized by the current upstream
algorithm.
Verified locally:
- CPU build: 84 tests, 75 pass + 8 new VBM tests, 1 pre-existing
test_read_write_grid BLOSC env-dependent failure unrelated to this PR.
- CUDA sm_120 + BLOSC + ZLIB: 84/84 pass.
No new lines over 100 cols.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: skip ndarray tests cleanly when numpy is missing
Windows CI surfaced three test errors on #2213 from
TestZeroCopyViewLifetimes — the test methods invoked
.values() / .leaf_values() directly without an `import numpy` guard,
so on environments without numpy installed (e.g. the Windows runner)
the bindings raised:
TypeError: could not export nanobind::ndarray:
ModuleNotFoundError: No module named 'numpy'
The other ndarray-touching tests (test_leaf_values_zero_copy,
test_bulk_leaf_values, test_node_manager_round_trip, the VBM
test_decode_block_zero, etc.) already guard with `import numpy`
and call self.skipTest. The three lifetime tests just skipped
that guard.
Added the same try/except ImportError + skipTest pattern to:
- TestZeroCopyViewLifetimes.test_handle_grid_tree_leaf_values_chain
- TestZeroCopyViewLifetimes.test_grid_leaf_values_temporary
- TestZeroCopyViewLifetimes.test_node_manager_temporary_grid
While I was here, also dropped an unnecessary numpy guard on
TestVoxelBlockManager.test_create_on_index_grid_rejects_unsupported_source
— that test only checks TypeError on a non-grid input and doesn't
touch numpy, so it can run unconditionally.
Verified locally with both a numpy-enabled venv (84/84 pass minus the
pre-existing BLOSC env failure) and a numpy-less venv (84 ran, 17
skipped, no new errors).
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: install numpy on Windows CI so ndarray tests run
The Windows job for the NanoVDB workflow was missing the install_numpy
step that the main openvdb build.yml and weekly.yml workflows already
run after install_windows.ps1. As a result the vcpkg-supplied Python
on the Windows runner had no numpy, so every test that touches
nb::ndarray<nb::numpy, ...> either errored (before #2213's skip
guards) or now skips silently.
Add the install_numpy step using the existing
ci/install_windows_numpy.ps1 helper, matching the pattern already in
build.yml. The Windows runner will then actually execute the
ndarray-touching VoxelBlockManager / Tree / GridHandle tests instead
of skipping them.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: address Copilot review on #2213
Four issues raised by Copilot on the open VBM follow-up PR, all real:
1. VoxelBlockManagerHandle.jumpMap(jump_map_length) accepted any
caller-supplied length and used it to shape a zero-copy NumPy view
over hostJumpMap(). Since the underlying buffer was sized as
blockCount * JumpMapLength (where JumpMapLength is derived from
the log2_block_width the handle was BUILT with), a caller passing
a larger value produced an ndarray whose elements lived past the
end of the allocated buffer — an OOB read on access.
Fix: wrap VoxelBlockManagerHandle in a small PyVBMHandle struct
that also stores the log2_block_width. jumpMap() takes no
arguments and derives JumpMapLength from the stored width, so the
returned view always matches the buffer exactly. decodeBlock() no
longer accepts a log2_block_width either, removing the equivalent
mismatch hazard there. The new PyVBMHandle exposes
log2_block_width / block_width / jump_map_length as read-only
properties for introspection.
2. decodeInverseMaps() did not validate first_leaf_id against
grid.tree().nodeCount(0). VoxelBlockManager::decodeInverseMaps
indexes tree.getFirstNode<0>()[first_leaf_id] unconditionally, so
a stray ID produced an OOB read of the leaf array. Validate up
front and raise IndexError.
3. buildVoxelBlockManager() did not enforce grid.isSequential() (a
precondition guarded only by NANOVDB_ASSERT, which is a no-op in
release) nor that first_offset == 1 (mod BlockWidth). Validate
both at the Python boundary and raise ValueError; the zero
first_offset default still flows through unchanged because the
C++ helper normalizes it to 1 itself.
4. The zero-copy shape/dtype test exercised
vbm.jumpMap(jump_map_length=2) on a handle built with
log2_block_width=6 — exactly the OOB case (1) was about. Rewrite
the test to build a second handle with log2_block_width=7 and
verify its jumpMap shape is (blockCount, 2), confirming the
shape now follows the handle. Drop the obsolete
log2_block_width=6 kwarg from decodeBlock. Add two new tests:
misaligned first_offset is rejected, and an out-of-range
first_leaf_id on decodeInverseMaps raises IndexError.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: empty-handle guards + accurate ownership comment
Two follow-up items from Copilot on #2213:
1. firstLeafID() and jumpMap() called hostFirstLeafID() / hostJumpMap()
on the underlying VoxelBlockManagerHandle and fed the result straight
into nb::ndarray, but both accessors legally return nullptr on a
default-constructed or reset() handle (blockCount() == 0). Passing
nullptr into nb::ndarray is unsafe even with a zero leading shape.
Mirror the pattern PyTree.h uses for the empty-grid leaf_values()
case: when the buffer is null, substitute a non-null dummy pointer
(the handle itself) so nanobind has something to base the empty
ndarray on; nothing is read because the leading shape is 0. Add
tests for both default-constructed and reset() handles.
2. The pyDecodeInverseMapsImpl comment described nanobind as allocating
"fresh memory through numpy", which doesn't match the implementation
(which uses new[] + a capsule with delete[] as deleter). Rewrite the
comment to describe the actual ownership model so future maintainers
aren't misled.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: VBM exception-safety + firstLeafID sentinel prefill
Two more items from Copilot on #2213:
1. pyDecodeInverseMapsImpl allocated leafIndex / voxelOffset with raw
new[] and only wrapped them in nb::capsule after several intervening
operations (a second new[], the decodeInverseMaps call itself, and
the first capsule's own construction). If anything in that window
threw, the half-built state leaked. Hold the raw arrays in
std::unique_ptr until each capsule has been constructed, then
release() so ownership transfers cleanly; any throw during that
sequence now unwinds without leaking.
2. The C++ allocating overload of buildVoxelBlockManager uses
HostBuffer::create (uninitialized malloc) for firstLeafID, then
touches only the slots its iteration sweep reaches. Blocks the
algorithm doesn't visit retain arbitrary values; the existing
decodeBlock guard (firstLeafID >= nLeaves) catches values past the
leaf array but cannot tell garbage that happens to be < nLeaves
from a real leaf id, so a low garbage byte would silently decode
against the wrong leaf.
Switch the Python binding to allocate the metadata buffers itself,
prefill every firstLeafID slot with the sentinel value `nLeaves`,
then call the in-place buildVoxelBlockManager(grid, handle)
overload. The in-place builder zeros the jumpMap and only writes
firstLeafID slots it visits, so every untouched slot keeps the
sentinel and deterministically trips the decodeBlock guard.
Add test_untouched_blocks_trip_sentinel_guard, which sweeps every
block of the cube VBM and asserts each firstLeafID slot is either
a real leaf id (< nLeaves) or exactly the sentinel — never any
other value.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: VBM popcount upper-bound + n_blocks validation
Two more items from Copilot on #2213:
1. The decodeBlock guard only checked firstLeafID < nLeaves, but the
C++ decodeInverseMaps loops leafID = firstLeafID ..
firstLeafID + nExtraLeaves where nExtraLeaves is the popcount of
this block's jumpMap (each set bit marks an additional leaf
boundary crossed within the block). With a corrupt jumpMap or a
handle paired with a different grid, the loop could read past
tree.getFirstNode<0>() even with a valid firstLeafID.
Hoist a popcount-based upper-bound check into
pyDecodeInverseMapsImpl: compute nExtraLeaves locally and raise
ValueError if firstLeafID + nExtraLeaves >= grid.tree().nodeCount(0).
Because both the handle.decodeBlock() and the free-function
decodeInverseMaps() funnel through this impl, both paths are
covered without per-caller duplication.
2. buildVoxelBlockManager accepted an explicit n_blocks but didn't
validate it against the documented precondition
n_blocks >= ceil((last_offset - first_offset + 1) / BlockWidth).
A smaller value produced a handle whose blockCount() < the
coverage implied by lastOffset, silently truncating later
decodeBlock sweeps. Validate when nonzero and raise ValueError
with the minimum required value in the message.
New test_build_voxel_block_manager_rejects_undersized_n_blocks
covers the n_blocks=1 case on the cube VBM.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: explicit std headers in PyVoxelBlockManager.cc
Copilot noted PyVoxelBlockManager.cc uses std::integral_constant,
std::string/std::to_string, and std::move but relied on transitive
includes from nanobind / NanoVDB headers to bring those in. That works
on the toolchains we currently test but is fragile against stricter
ones.
Add explicit <string>, <type_traits>, and <utility> includes; drop
<cstring> and <stdexcept>, neither of which the translation unit
actually uses now.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: Phase 4a — tools.build.Grid<T> mutable CPU builder (#2214)
* nanovdb python: Phase 4a — tools.build.Grid<T> mutable CPU builder
Bind nanovdb::tools::build::Grid<BuildT>, its ValueAccessor<BuildT>, and
Tree<BuildT>::WriteAccessor under a new nanovdb.tools.build submodule.
One set of classes per writable BuildT in BuildTypes.def — every scalar
(float, double, int16, int32, int64, uint8, uint32) and every vector
(Vec3f, Vec3d, Vec4f, Vec4d, Vec3u8, Vec3u16, Rgba8). Naming mirrors
the C++ namespace: nanovdb.tools.build.FloatGrid is the mutable
counterpart of the read-only nanovdb.FloatGrid.
The binding exposes:
- Grid<T>(background, name='', gridClass=Unknown) — constructor
- getValue / setValue / setValueOn / isActive (the last two convenience-
wrap a fresh ValueAccessor under the hood, since C++ has no
setValueOn/isActive on Grid itself)
- nodeCount, gridType, gridClass, getName/setName, setTransform,
.background property
- getAccessor() / getWriteAccessor() returning the typed Value /
WriteAccessor proxies
- .to_nanovdb(sMode, cMode, verbose) — bakes the build grid into a
host NanoGrid<BuildT> handle by calling tools::createNanoGrid
ValueAccessor exposes getValue/setValue/setValueOn/isActive/isValueOn.
WriteAccessor exposes setValue/setValueOn/merge. Both are wired up so
the parent grid is kept alive by Python while the accessor lives.
Implementation notes:
- WriteAccessor's defaulted move constructor leaves its internal
ValueAccessor::mRoot reference dangling (the reference points into
the moved-from WriteAccessor's own mRoot field, which is per-object
state — not the parent Tree's mRoot). The Python binding bypasses
the move path by heap-allocating via nb::rv_policy::take_ownership
so the C++ object's address is stable for its entire lifetime.
- ValueAccessor doesn't have this hazard because its mRoot reference
points at the parent Tree's mRoot (a stable address external to the
accessor), so move construction is safe.
- Read-only special BuildTs (Boolean, Fp4/8/16/N, ValueIndex,
ValueOnIndex, ValueMask) and Point are deliberately excluded — they
have no SetValue<T> specialization and can't be built voxel-by-voxel.
Test coverage in new TestBuildGrid: constructor defaults, setValue
marks active, setValueOn preserves background, accessor parity with
grid, WriteAccessor merge-on-destruction, .to_nanovdb() round-trip
with metadata preserved, .to_nanovdb() doesn't consume the source,
Int32Grid + Vec3fGrid spot-checks, and setTransform propagation to
the baked grid's voxelSize / map.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: address Copilot review on #2214
Four items from the Phase 4a review, all valid:
1. Grid<T>.background property read self.mRoot.mBackground directly,
touching an internal field even though RootNode exposes a
background() accessor. Switch to self.mRoot.background() so the
binding doesn't depend on the underlying field name staying put.
2. .to_nanovdb() can be an expensive bake for large grids but held
the GIL the whole time. Add nb::call_guard<nb::gil_scoped_release>()
so other Python threads can run during the conversion (the lambda
only touches C++ state, no Python object handling).
3. nodeCount()'s docstring said the tuple was "internal node counts",
but the first element is the leaf (level-0) count — leaves are not
internal nodes. Reword to just "(leaf_count, lower_count,
upper_count)" and drop the "internal" mislabel.
4. The WriteAccessor merge test was named "merges_on_destruction" but
actually called wa.merge() explicitly and never forced the
destructor to run. Split into two cases:
- test_write_accessor_explicit_merge — calls .merge() explicitly,
matching what the original test actually checked
- test_write_accessor_merges_on_destruction — drops the only
reference (del wa) and runs gc.collect() so the C++ destructor
fires, then asserts the change is visible
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: use Vec3f.__eq__ in build::Vec3fGrid test
Copilot noted the component-by-component comparison in
test_vec3f_build_grid was justified by an out-of-date claim that
Vec3f equality isn't bound — it is (PyMath.cc defineVec3 wires
nb::self == nb::self). Use self.assertEqual(got, v) directly.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: Phase 4b + 4c — stats, validation, checksum (#2215)
* nanovdb python: Phase 4b + 4c — stats, validation, checksum
Round out Phase 4 by binding the GridStats, GridValidator, and
GridChecksum surfaces that didn't ship with Phase 4a.
Phase 4b — Stats
================
* Per-BuildT `tools.<Suffix>Extrema` and `tools.<Suffix>Stats` classes
for every scalar and vector BuildT in `BuildTypes.def`. Extrema
exposes `min` / `max` / `add(value)` and a truthy `bool()` for "has
at least one sample"; Stats inherits from Extrema and adds `size`,
`avg` / `mean`, `var` / `variance`, `std` / `stdDev`. Static
predicates `hasMinMax` / `hasAverage` / `hasStdDeviation` /
`hasStats` mirror the C++ trait queries.
* `tools.updateGridStats(grid, mode=Default)` — polymorphic dispatch
via `callNanoGrid` over an `UpdateGridStatsOp`. Scalar and vector
BuildTs route to `tools::updateGridStats<BuildT>`; the special /
quantized / index / mask types (`Fp4/8/16/N`, `ValueIndex`,
`ValueOnIndex`, `ValueMask`, `Point`) raise `ValueError` because
`Stats<ValueT>` isn't meaningful for them. `bool` falls through to
the C++ `NoopStats` arm.
* Per-BuildT `tools.getExtrema(grid, bbox)` returning the matching
`Extrema`. Restricted to scalar + vector BuildTs (the only ones
with an arithmetic ValueType).
Both `updateGridStats` and `getExtrema` release the GIL during the
traversal.
Phase 4c — Validation & checksum
================================
* `tools.validateGrid(handle, gridID, mode=Default, verbose=False)`
for `GridHandle<HostBuffer>` — single-grid complement of the
existing `validateGrids`. Returns `False` (without raising) when
the gridID is out of range.
* `tools.checkGrid(grid, mode=Full)` — polymorphic via `callNanoGrid`.
Returns `(ok, error_message)`. The 256-byte char buffer the C++
helper writes into is hidden inside the binding so Python callers
see a `(bool, str)` tuple.
* `tools.isValid(grid, mode=Default, verbose=False)` — polymorphic via
`callNanoGrid<IsValidOp>`, equivalent to `checkGrid` + checksum
verification rolled into a single bool.
* `tools.evalChecksum(grid, mode=Default)` and
`tools.validateChecksum(grid, mode=Default)` — accept any bound
NanoGrid via the existing `GridData` Python upcast (every
`NanoGrid<T>` Python class is registered with `GridData` as its
base, so nanobind handles the dispatch transparently). All three
release the GIL.
Test coverage in new `TestGridStats`, `TestGridValidate`, and
`TestGridChecksum` (11 cases total): Extrema/Stats default-and-add,
polymorphic `updateGridStats` on a float grid, `updateGridStats`
rejection on an OnIndex grid, `getExtrema` over a sub-bbox,
`checkGrid` / `isValid` / `validateGrid` happy paths plus
out-of-range `gridID` and `CheckMode.Disable` short-circuit, and
`evalChecksum` → `updateChecksum` → `validateChecksum` round-trip.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: address Copilot review on #2215
Four items from the Phase 4b/4c review, all valid:
1. UpdateGridStatsOp rejected every BuildTraits<BuildT>::is_special
type unconditionally (except bool), but tools::updateGridStats
actually supports StatsMode::BBox on any ValueT via the NoopStats
path. Drop directly into NoopStats for special BuildTs when mode
is Disable or BBox; only MinMax / All raise (because those would
instantiate Stats / Extrema over a non-arithmetic ValueT and the
semantics are ill-defined). Update the binding docstring to
describe the actual matrix.
2. The test_update_grid_stats_rejects_index_grid case was renamed to
test_update_grid_stats_on_index_grid and extended to cover all
four StatsMode arms on an OnIndexGrid: MinMax and All raise,
BBox and Disable now succeed.
3. validateGrid's docstring didn't mention the CheckMode.Disable
short-circuit (which returns True without inspecting the handle
or gridID). The Python tests already exercised the short-circuit
via test_validateGrid_disable_mode_always_true, so the behavior
was correct — just the docstring was missing it.
4. IsValidOp::unknown ignored verbose=True and silently returned
false, but the C++ callNanoGrid::unknown arm writes an
"Unsupported GridType" message to std::cerr when verbose is set.
Mirror that — pull in <iostream> in the binding TU and emit the
same diagnostic from C++.
5. test_get_extrema_over_active_bbox used a bbox that exactly equals
the root's active bbox, which triggers C++ getExtrema's
"bbox contains root.bbox()" short-circuit that unconditionally
folds the grid background into the extrema — so the min came
back as 0.0 (background), not 1.0 (the smallest active voxel).
The test passed but the semantics were muddy: it was really
asserting "background gets added in this branch" rather than
anything about getExtrema's bbox restriction. Replace with
test_get_extrema_strictly_inside_active_region, which uses a
bbox strictly inside the active region (so the recursive branch
runs) and asserts min/max are exactly the smallest/largest
sampled active values.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: bind validateGrid for device handles too
Copilot noted tools.validateGrids was registered for both
GridHandle<HostBuffer> and GridHandle<cuda::DeviceBuffer> (the
latter behind NANOVDB_USE_CUDA), but my Phase 4c addition of the
single-grid tools.validateGrid only covered HostBuffer — leaving
device handles able to validate the whole bundle but not an
individual grid.
Add the matching #ifdef NANOVDB_USE_CUDA overload binding
tools::validateGrid<GridHandle<cuda::DeviceBuffer>>. The C++
helper does host-side dispatch via callNanoGrid on the
host-resident gridData() pointer that DeviceGridHandle exposes, so
the same Python overload pair is appropriate.
New test_validateGrid_on_device_handle exercises the device path
end-to-end (build a CUDA level-set sphere, validate it, confirm
the same out-of-range and Disable-mode short-circuits as the host
overload). The test is gated on cuda module availability so it's
skipped cleanly on CPU-only builds.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: address more Copilot review on #2215
Two more items from the review, both valid:
1. CheckGridOp wrote tools::checkGrid's error message into a 256-byte
stack buffer. tools::checkGrid (and its util::sprint / util::strcpy
helpers) trust the caller's buffer is large enough — there is no
bounds-checked variant. The current error messages are at most ~80
characters once GridType and GridClass enumerator names are
stringified, so 256 wasn't actually overflowing today, but the
margin is uncomfortable and a future error string addition could
push past it. Bump to a 4096-byte buffer (named via a constexpr
kErrorBufSize) and zero-init the first byte before the call so the
ok-detection still works if the helper bails out before writing
anything.
2. updateGridStats's docstring claimed special (quantized / index /
mask) BuildTs only accept Disable / BBox, but Boolean grids are
special yet fall through the if-constexpr filter to the regular
tools::updateGridStats path (which routes any mode to NoopStats
for ValueT=bool inside the C++ helper). Reword the docstring to
call out Boolean as the exception that accepts every StatsMode
via the NoopStats internal path.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: Phase 5 — primitives + quantized/index createNanoGrid (#2216)
* nanovdb python: Phase 5 — primitives + quantized/index createNanoGrid
Closes out Phase 5 of the NanoVDB Python bindings restructure (#2208).
Lands all three sub-phases (5a primitives, 5b quantized createNanoGrid
overloads, 5c generic createNanoGrid from build::Grid) in one PR.
Phase 5a — Host primitives
==========================
Bind the nine primitives that didn't ship with Phase 0:
* tools.createLevelSetBox(gridType, width, height, depth, ...) — narrow-
band level set of a solid axis-aligned box.
* tools.createLevelSetBBox(gridType, width, height, depth, thickness, ...) —
narrow-band level set of a hollow box wireframe.
* tools.createLevelSetOctahedron(gridType, scale, ...) — narrow-band
level set of an octahedron.
* tools.createFogVolumeBox(...) and tools.createFogVolumeOctahedron(...) —
the fog-volume counterparts of the above.
* tools.createPointSphere / createPointTorus / createPointBox(gridType,
pointsPerVoxel, ...) — PointDataGrids scattered on the surface of each
primitive shape.
* tools.createPointScatter(srcGrid, pointsPerVoxel, ...) — scatter a
PointDataGrid into the active voxels of a NanoGrid<float> level set or
fog volume. The C++ template also accepts double sources; the binding
is float-only for simplicity (the runtime grid pointer carries the
source BuildT and adding a per-type dispatch is mechanical follow-up).
Each non-point primitive is instantiated for float and double via the
same runtime-GridType switch the existing createLevelSetSphere et al.
already use. The FpN overloads of these primitives are deliberately
not bound here — quantization is reachable via Phase 5b's generic
createNanoGrid* path with explicit oracle and dither parameters.
Phase 5b — Quantized createNanoGrid overloads
=============================================
* tools.AbsDiff(tolerance=-1.0) and tools.RelDiff(tolerance=-1.0) —
compression-oracle classes used by FpN. Both expose getTolerance /
setTolerance and a truthy __bool__ that returns True iff the
tolerance has been initialized (>= 0). The default tolerance of -1
matches the C++ "uninitialized — fill in via init()" sentinel.
* tools.createNanoGridFp4 / Fp8 / Fp16(src, sMode, cMode, ditherOn,
verbose) — quantize a float source into a fixed-bit-width grid.
ditherOn adds sub-quantum noise to break up banding.
* tools.createNanoGridFpN(src, oracle, sMode, cMode, ditherOn, verbose) —
variable-bit-width quantization. Two overloads: one accepting an
AbsDiff oracle (the default), one accepting a RelDiff oracle. Python
picks the right overload from the oracle argument's type.
The C++ Fp{4,8,16,N} preProcess templates static_assert SrcValueT ==
float, so the binding rejects double sources with a TypeError instead
of letting the compile-time assertion trip a hard abort.
Phase 5c — Generic createNanoGrid from build::Grid
==================================================
Each conversion entry accepts both NanoGrid<SrcBuildT> and
tools::build::Grid<SrcBuildT> as its source. Internally a small
template helper tries each source-side BuildT in turn (matching the
existing tryCreateOnIndexGrid pattern from the Phase 3 follow-up) and
dispatches via nb::isinstance.
* tools.createNanoGridIndex(src, channels=0, includeStats=True,
includeTiles=True, verbose=0) — NEW. Bake any supported source into
a NanoGrid<ValueIndex> with all voxels (active and inactive) given a
uint64 sequential index. Original values can be carried as blind
data when channels > 0.
* tools.createNanoGridOnIndex(src, ...) — same as Index but only the
active voxels get a sequential index. Supersedes the Phase 3
follow-up's tools.createOnIndexGrid test scaffold (which is kept
alive in PyVoxelBlockManager.cc for backwards compatibility with
the VBM tests; new code should prefer createNanoGridOnIndex).
Source types accepted by the index path are NanoGrid<float | double |
int32 | Vec3f> and tools::build::Grid<float | double | int32 | Vec3f>.
The same-type bake of build::Grid<T> -> NanoGrid<T> remains served by
the Phase 4a .to_nanovdb() method.
Test plan
=========
Three new TestCase classes covering: every primitive (TestNewPrimitives,
11 cases), every quantization path including double-source rejection
and oracle defaults (TestCreateNanoGridQuantized, 8 cases), and every
index/onindex source type including the build::Grid path
(TestCreateNanoGridIndex, 6 cases). All 25 new cases pass locally; the
full pytest_nanovdb suite stays green except for the two pre-existing
BLOSC-disabled errors in my local build.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: address Copilot review on #2216
Four items from the Phase 5 review, plus a follow-on segfault fix
caught while addressing the third item:
1. openToNanoVDB binding accidentally dropped the sMode argument and
gave `base` a default of tools::StatsMode::Default — nonsense both
semantically (base is an OpenVDB GridBase::Ptr) and structurally
(the C++ template expects 4 args, the binding declared 3 keyword
args). The CI build broke for openvdb-enabled configs with a
"number of nb::arg annotations must match the argument count"
static_assert from nanobind. Restore the correct signature: base
has no default, sMode/cMode/verbose carry their defaults.
2. Octahedron primitive default name strings carried the upstream C++
typo "octadedron" (sic). The C++ side keeps the misspelling for
binary compatibility with old grids; correct it on the Python
side since the default propagates into help() output and grid
metadata where end-users see it.
3. The point primitives took a `gridType` argument that was
misleading — it controls the intermediate level-set's precision
the scatter starts from, not the returned PointDataGrid's type
(which is always UInt32). While investigating the rename, I found
that the createPointSphere<double> / createPointTorus<double> /
createPointBox<double> code paths actually segfault during scatter
with the current C++ implementation (only the float path is
exercised by the C++ unit tests). Take Copilot's "remove it and
always use float" alternative: drop the argument entirely. The
binding now always uses the float intermediate level-set.
createPointSphere / Torus / Box no longer have the parameter;
their docstrings explicitly note "always returns a UInt32
PointDataGrid".
4. The defineCreateNanoGridConversions doc-comment claimed sources
were accepted for "every scalar/vector BuildT in BuildTypes.def
that tools::createNanoGrid supports", which was wishful — the
implementation actually tries a narrower set per destination kind.
Spell out the actual matrix: quantized Fp{4,8,16,N} paths accept
float only (the C++ preProcess static-asserts SrcValueT == float);
the createNanoGridIndex / OnIndex paths accept float / double /
int32 / Vec3f. Adding more source types is a one-line extension
of the explicit try-each-SrcBuildT chains.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: address more Copilot review on #2216 + CI GPU gating
Four items in this round, three from Copilot plus a CI failure caught
on the linux-clang++ runner.
1. test_validateGrid_on_device_handle was gated only on the cuda
submodule being importable, but the linux-clang++ CI runner has a
CUDA toolkit installed (so the submodule loads) but no GPU
driver. The test then crashed inside the C++ helper with
"CUDA error 35: CUDA driver version is insufficient for CUDA
runtime version". Switch the gate to the project's existing pair
of helpers — nanovdb.isCudaAvailable() (build-time CUDA support)
AND nanovdb.isGpuAvailable() (runtime device probe) — matching
how TestPointsToGrid / TestSignedFloodFill / TestSampleFromPoints
already gate. The test skips cleanly on CPU-only runners and
driverless CUDA runners alike.
2. The Phase 5b/5c conversion helpers (tryQuantizeFpX, tryQuantizeFpN,
tryIndexify) ran the heavy tools::createNanoGrid traversal while
holding the Python GIL. Restructure each helper so the GIL is
held only for the isinstance / cast dispatch, then released
around the conversion itself (the source data lives in stable
C++ storage anchored by the Python wrapper passed in via py_src,
so it's safe to read without the GIL). Other Python threads can
now make progress during large-grid quantization / indexing.
3. AbsDiff's class docstring said "pass an explicit positive value"
for the tolerance, but the C++ operator bool() actually treats
any non-negative value (including 0.0) as initialized. Reword to
match what the API actually does.
4. The createNanoGridIndex/OnIndex source-rejection test was named
test_index_rejects_unsupported_source with a comment about
"bool / Boolean grids" but really only exercised the None case.
Split into two cases with single, accurate purposes:
- test_index_rejects_none — passes None; matches neither the
NanoGrid nor the build::Grid isinstance arms.
- test_index_rejects_unsupported_buildt — builds a
tools.build.Vec3dGrid (structurally valid but a BuildT outside
the float/double/int32/Vec3f source set) and confirms the
try-each-SrcBuildT chain falls through to a TypeError.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: correct createPointScatter docs + add fog-rejection test
Copilot noted the createPointScatter binding docstring claimed the
source could be a "level set or fog volume", but the C++ implementation
explicitly checks srcGrid.isLevelSet() at line 1679 of
tools/CreatePrimitives.h and throws std::runtime_error("Expected a
level set grid") otherwise. Match the actual behavior: the docstring
now says the source must satisfy srcGrid.isLevelSet() and that
non-level-set sources (e.g. fog volumes) raise RuntimeError.
Add test_create_point_scatter_rejects_fog_volume to lock in the
behavior (build a fog volume sphere, confirm createPointScatter raises
RuntimeError).
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: fully-qualify nanovdb.tools.build.* in 5b/5c docs
Copilot noted the Phase 5b/5c TypeError messages and per-function
docstrings referred to the mutable build grids as "tools.build.FloatGrid",
but the actual Python import path is "nanovdb.tools.build.FloatGrid".
Replace every occurrence so help() / autocomplete / error messages
point at a path users can actually import.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: examples + docstring coverage sweep (#2218)
* nanovdb python: Phase 6 — docs, migration guide, examples
Closes Phase 6 of the NanoVDB Python bindings restructure (#2208).
Ships:
* doc/nanovdb/PythonMigration.md — porting guide for users on the
pre-Phase-1 typed-accessor API. Covers the polymorphic
handle.grid(n) replacement of floatGrid() / doubleGrid() / etc.,
the removal of the nanovdb.math.cuda submodule (sampleFromVoxels
is now in nanovdb.tools.cuda), the GridHandle.__bool__ + enum
__repr__ behavior changes, and a survey of every new surface that
used to require dropping to C++ (GridMetaData, blind data,
PointAccessor, tree/node walking, NodeManager, leaf_values,
VoxelBlockManager, tools.build.Grid, stats, validation,
primitives, quantized createNanoGrid). Closes with a mechanical
before/after replacement table.
* doc/nanovdb/PythonAPI.md — narrative API reference for the final
surface. One section per submodule (nanovdb root, math, io,
tools, tools.build, tools.cuda, cuda) listing the bound classes
and functions with one-line descriptions. Points readers at
help(x) for full signatures.
* nanovdb/nanovdb/python/examples/ — five runnable .py scripts:
- load_inspect.py: polymorphic handle.grid(n), GridMetaData
introspection, mixed-type handle via mergeGrids.
- build_grid.py: tools.build.FloatGrid setValue / ValueAccessor /
WriteAccessor / .to_nanovdb() round-trip.
- bulk_leaf_numpy.py: zero-copy (N_leaves, 512) NumPy view via
grid.leaf_values(), with global-stats reduction and in-place
mutation demonstrating the no-copy semantics.
- quantize.py: createNanoGridFp{4,8,16} fixed-width + createNanoGridFpN
with AbsDiff and RelDiff oracles. Prints the per-format size
so users see the compression tradeoff at a glance.
- validate.py: validateGrid / validateGrids / checkGrid / isValid
plus evalChecksum / updateChecksum / validateChecksum.
* nanovdb/nanovdb/python/examples/README.md — index of the above
with one-line summaries.
* nanovdb/nanovdb/Readme.md — added a "Python bindings" section
linking the API reference, migration guide, and examples.
Each example is self-contained (builds its own input data,
prints a small stdout summary) and was smoke-tested locally
against the post-Phase-5 binding. The bulk_leaf_numpy example is
the only one requiring NumPy; it prints a friendly skip message
otherwise.
API reference docs via Sphinx / mkdocs are deferred to a follow-up;
the existing docstrings already make help() useful and the .pyi
type stubs (from Phase 0) cover IDE / type-checker integration.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: drop end-user docs, scrub refactor framing from examples
Two changes in response to review feedback:
* Remove doc/nanovdb/PythonAPI.md and doc/nanovdb/PythonMigration.md.
End users don't need a separate narrative reference (help() and
the .pyi stubs cover the same ground) and don't need a migration
guide framed around the multi-phase refactor history.
* Rewrite the docstring blurb at the top of each example
(load_inspect.py, build_grid.py, bulk_leaf_numpy.py, quantize.py,
validate.py) so the prose describes what the example demonstrates
about the current API, with no mention of which phase of the
refactor added the feature. Same scrub on
python/examples/README.md (drop the "bindings restructure" framing
and the now-dead links to the deleted docs) and on
nanovdb/nanovdb/Readme.md (drop the new "Python bindings" section
header — just leave the single examples link alongside the existing
Examples link).
Examples themselves remain unchanged in behavior; all five still
run cleanly against the current binding.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: docstring coverage sweep across the binding surface
Add a short prose docstring to every nanobind .def() and nb::class_
call that didn't already carry one. The audit dropped from 374
distinct undocumented call sites (3278 raw, counting per-BuildT
template duplicates) to 0.
Style:
* One short imperative sentence per method, focused on what it does
in Python terms.
* For methods that exactly mirror a C++ member of the same name, a
"See nanovdb::<Class>::<method> in NanoVDB.h." tail so users can
find the canonical reference without us re-stating it.
* Class-level docstrings on every bound class (Grid, GridHandle,
GridMetaData, Tree, Root, Upper, Lower, Leaf, NodeManager, the
per-BuildT accessor subclasses, Map, Coord, BBox*, samplers,
build::Grid + ValueAccessor + WriteAccessor, Extrema / Stats,
oracle classes, Checksum, etc.).
* Two-to-three sentence blurbs on the entry-point classes that
genuinely need orientation.
Files touched: every Py*.cc / Py*.h binding file under
nanovdb/nanovdb/python/ plus the three cuda/ binding files. No code
changes — only the trailing const char* docstring argument on .def
calls and the third arg of nb::class_. The auto-generated .pyi stub
output picks the new prose up automatically; help() shows the same
text at runtime.
Also drop the few remaining "Phase N" / "from Phase N follow-up"
references that survived in internal C++ comments in
PyCreateNanoGrid.cc, PyPrimitives.cc, and PyVoxelBlockManager.cc,
keeping the source consistent with the no-refactor-history policy
already applied to the public docstrings, tests, and examples.
Test plan: full pytest_nanovdb green locally (138/140; the two
errors are pre-existing BLOSC-disabled errors). Module compiles
cleanly with CUDA + nanobind stub generation enabled.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: address Copilot review on #2218
Two items from the review:
1. The Readme link to the Python examples directory used the wrong
relative path. From nanovdb/nanovdb/Readme.md, the examples live
at python/examples/ (sibling), not nanovdb/python/examples/
(which would resolve to nanovdb/nanovdb/nanovdb/python/examples/).
2. The Map.applyMap / applyJacobian / applyInverseMap /
applyInverseJacobian docstrings used the words "double precision"
and "single precision" to describe the variant suffix. That
wording is misleading: the F-suffixed variants use 32-bit math
internally, the unsuffixed variants use 64-bit math, but BOTH
return a vector whose dtype matches the input (Vec3f -> Vec3f,
Vec3d -> Vec3d). Switching to "uses 64-bit math" / "uses 32-bit
math" with an explicit "returns a vector of the same dtype as
the input" tail makes the semantics unambiguous. Same fix on the
Jacobian variants (which previously didn't mention precision at
all on the 64-bit overloads, only on the 32-bit ones — also
confusing).
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: clarify Grid.gridName docstring
Copilot caught that the Grid.gridName docstring described it as "the
in-header buffer; truncated if very long" — which is actually the
semantics of shortGridName(). gridName() reads the LONG-form name
from blind data when the HasLongGridName flag is set, falling back
to the in-header buffer otherwise. Reword the docstring to match
and explicitly point at shortGridName() for the truncated form.
(The other comment in the same review pass flagged the Examples
table in python/examples/README.md as having "double leading pipes"
but the actual file uses single-pipe GitHub-flavored markdown
syntax — verified via per-character inspection. The table renders
correctly on GitHub; no change needed there.)
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: clarify VBM offset and Root.bbox docstrings
Two more Copilot review items:
* VoxelBlockManagerHandle.firstOffset / lastOffset docstrings said
"Linear leaf offset of the first/last block..." but these are
sequential VOXEL indices, not leaf IDs. Switch the wording to
"Sequential voxel index of the first/last active voxel covered
by this handle" so users aren't confused into thinking the value
is a leaf number.
* RootT.bbox docstring said "Index-space bounding box of every
active tile" but RootNode::bbox() returns the bounding box of
every active VALUE in the tree, not just the tile entries in the
root table. Drop the misleading "tile" wording.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: remove dead code in bulk_leaf_numpy example
Copilot caught that bulk_leaf_numpy.py computed n_active_after via a
per-voxel sum over leaf.isActive() but never used the result — dead
code that linters would flag. Replace with leaf.getFirstValue(),
which is actually informative for what the surrounding comment is
trying to show (the values changed in place via the zero-copy
NumPy write) and prints in the same line.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: bounds-check GridHandle grid accessors; fix stale docstring
Addresses review feedback on PR #2219.
GridHandle::gridSize(n) and gridType(n) index mMetaData[n] without a
bounds check (the C++ API documents "assumed to be less than
gridCount()"), so calling handle.gridSize()/gridType()/gridData() with
an out-of-range n — including any n on an empty handle — was undefined
behaviour reachable straight from Python. Wrap all three in
bounds-checked lambdas that raise IndexError. gridData() is guarded too:
gridData(n) itself returns nullptr for bad n, but the binding passes
gridSize(n) alongside it, and that call is the one that reads
mMetaData[n] out of bounds.
Also refresh the createOnIndexGrid docstring: it claimed the broader
createNanoGrid<SrcGridT, DstBuildT> surface "lands in a later phase",
but tools.createNanoGridOnIndex already ships in this PR
(PyCreateNanoGrid.cc). Point users at the canonical binding instead.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: harden mergeGrids for empty input and use gridData(n)
Addresses review feedback on PR #2219.
Empty input: mergeGrids([]) — or a sequence of only-empty handles —
produced totalGrids == 0 and called BufferT::create(0). For HostBuffer
that yields a buffer whose data() is non-null over a zero-byte region,
so the GridHandle(buffer&&) ctor then reads a full GridData header out
of it (heap-overflow read; in practice an opaque "invalid host buffer"
throw). Return an empty handle up front when there's nothing to merge.
Per-grid source pointer: copy from h->gridData(n) — the authoritative
start pointer that applies mMetaData[n].offset — instead of walking a
raw data() pointer advanced by gridSize(n). The two are equivalent for
the current tightly-packed layout (offsets are a running sum of grid
sizes), but the accessor form doesn't bake in that assumption and drops
the manual pointer arithmetic.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: release the GIL around VoxelBlockManager compute kernels
Addresses review feedback on PR #2219.
buildVoxelBlockManager, decodeBlock, and decodeInverseMaps all run
non-trivial C++ (the build may parallelize internally via
util::forEach; the decode sweeps multiple leaves) while holding the
GIL, blocking unrelated Python threads.
Release the GIL around the pure-C++ kernels only, matching the
established pattern in PyCreateNanoGrid.cc (tryIndexify / tryQuantizeFpX):
- pyDecodeInverseMapsImpl (shared by decodeBlock and the free
decodeInverseMaps): scope a gil_scoped_release around the
VBM::decodeInverseMaps call. The surrounding grid cast, bounds checks
(which throw Python exceptions), and the NumPy array / capsule / tuple
construction stay under the GIL.
- buildVoxelBlockManager: scope a gil_scoped_release around the
buildVoxelBlockManager<LBW> call only.
Deliberately not a blanket nb::call_guard<nb::gil_scoped_release>() as
suggested: these entry points take an nb::handle and call
castOnIndexGrid (isinstance/cast) plus allocate NumPy arrays inside the
function body, so dropping the GIL across the whole call would touch the
Python C API without the GIL. The scoped release covers exactly the
GIL-free compute.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: index-grid completeness, transform aliases, sampler gradients
Close the residual gaps between the Python bindings and the C++ host API:
- Bind Grid.valueCount() on Index/OnIndex grids and Grid.pointCount() on
PointGrid, gated per BuildT like the C++ SFINAE.
- Add tools.CreateNanoGrid mirroring the C++ converter class, exposing
addBlindData() so blind-data channels can be authored from Python and
filled through the writable getBlindData() NumPy view. Authored
channels are zero-filled (the C++ path leaves them uninitialized).
- Bind ChannelAccessor<ChannelT, IndexT> for Index/OnIndex x
float/double/int32/Vec3f, plus a polymorphic createChannelAccessor()
dispatching on the channel's recorded dataType.
- Add the worldToIndex/indexToWorld convenience names (Dir/Grad/F
variants included) as aliases on the type-erased Grid base.
- Bind gradient() and zeroCrossing() on the host samplers for the
orders/BuildTs where the C++ templates compile (trilinear gradient,
trilinear+triquadratic zeroCrossing, floating-point only).
- Fix a pre-existing bug: getBlindData() and the PointAccessor views
built ndarrays without the nb::numpy framework tag, so the def-site
keep_alive raised 'could not create a weak reference' whenever a
non-empty view was returned.
- Paper cuts: restore the empty error messages in the sphere/torus
primitive dispatchers, drop the unused baseName parameter from
defineStats, and bind Checksum.isEmpty/isHalf/isFull/mode.
TestNanoVDB.py grows 139 -> 161 cases, covering the new surface plus the
previously untested populated blind-data and channel-accessor paths.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* nanovdb python: port host C++ examples to Python, smoke-test all examples
Nine new scripts under python/examples/, each a Python equivalent of a
host-capable C++ example (or its host code path):
- io_roundtrip.py: ex_write_nanovdb_grids + ex_read_nanovdb_sphere(_accessor)
as a self-contained write/read round trip, plus splitGrids and zero-copy
point positions via getBlindData.
- make_funny_nanovdb.py: ex_make_funny_nanovdb via the functor-based
createFloatGrid factory (reduced domain — one Python call per voxel).
- make_typed_grids.py: ex_make_typed_grids across seven tools.build grid
types with polymorphic re-read.
- raytrace_level_set.py / raytrace_fog_volume.py: the host render paths,
re-expressed with the bound samplers (sphere tracing + zeroCrossing +
gradient shading) and accessor marching respectively; PGM output.
- collide_level_set.py: the host particle-collision path using
worldToIndexF, narrow-band isActive, and sampler.gradient normals.
- index_grid_channels.py: host half of ex_index_grid_cuda extended with
ChannelAccessor read-back and CreateNanoGrid blind-data authoring.
- node_manager.py: host half of ex_nodemanager_cuda.
- openvdb_interop.py: ex_openvdb_to_nanovdb(_accessor); self-skips unless
built with NANOVDB_USE_OPENVDB and openvdb imports.
Not ported: the pool-buffer examples (HostBuffer pool API unbound), the
MagicaVoxel converter (external parser/asset), device-only examples, and
a dedicated PointAccessor example (host point primitives bake UInt32
PointData grids; Point{Index,Data}Accessor require GridType.PointIndex).
test/TestExamples.py smoke-runs all 14 example scripts in subprocesses;
wired as ctest pytest_nanovdb_examples (~5 s) next to pytest_nanovdb,
gated on NANOVDB_BUILD_PYTHON_UNITTESTS.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* fix example DLL lookup on Windows
Run example smoke tests through a Windows-specific child launcher that
registers dependent DLL directories before importing nanovdb. Keep the
DLL directory handles alive while preserving subprocess isolation and
normal script execution behavior.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* fix Windows example test startup
Register dependent DLL directories through a test-only sitecustomize module so example smoke tests can run as normal Python scripts. This avoids the native access violations caused by executing them through runpy while preserving subprocess isolation.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* diagnose Windows example crashes
Enable Python's fault handler and unbuffered output for example subprocesses so Windows access violations report the active Python frame and preserve output leading up to the crash.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* fix Windows example DLL search
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* add CUDA DLL path to Windows examples
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Rename to_nanovdb->toNanoVDB to match convention
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Added comments on nb::object returned sentinel return in `try*` functions
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Fix `toNanoVDB` renames that were not caught.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* NanoVDB: add non-owning mergeGrids overload and use it from Python
Address review feedback on the Python bindings, which had reimplemented
the grid-merge logic instead of reusing nanovdb::mergeGrids. The
existing overload takes a const std::vector<GridHandle>&, which cannot
be built from Python-owned handles without moving from them (GridHandle
is move-only, and moving would silently empty the caller's objects) or
deep-copying every input buffer.
- Add mergeGrids(const GridHandle<BufferT>* const* handles, size_t
count, pool) to GridHandle.h, which borrows the inputs; the vector
overload now forwards to it, so there is a single implementation.
- While consolidating, fold in two fixes the Python binding had
accumulated: empty input returns an empty handle instead of
constructing from BufferT::create(0) (whose zero-byte region the
GridHandle constructor would read a header out of), and the source
walk uses gridData(n) — which applies the per-grid offset — instead
of assuming grids are contiguous from data(), so padded handles
merge correctly.
- Shrink the Python mergeGrids binding to collecting pointers and
calling the new overload; drop the now-unused <cstring> and
GridChecksum.h includes.
Verified with a standalone C++ test of both overloads (inputs left
intact, grid index/count rewritten, empty input, split round-trip) and
by running TestNanoVDB.py before/after: identical results, with all
merge/split tests passing.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* NanoVDB: use constructor-free shared-memory storage in GridStats kernels
CUDA does not support dynamic initialization of function-scope __shared__
variables, so nvcc skips the StatsT/CoordBBox constructors in processLeaf
and processInternal and emits warning #20054-D -- twelve copies in every
TU that instantiates the kernels, propagated to downstream builds since
GridStats.cuh is a header. The skipped initialization was harmless
(every shared slot is fully assigned before it is read), but the noise
is not.
Wrap the scratch arrays in cub::Uninitialized, whose backing store is a
trivially-constructible word array, and bind array references to the
aliased element type. The reduction code is unchanged.
Verified warning-free (0 vs 12) and performance-neutral: per-kernel SASS
comparison on sm_89 shows identical instruction counts, registers,
shared-memory bytes, and zero spills for all 14 kernels; timing 200
iterations of updateGridStats on an 8M-voxel grid shows matched-round
deltas within +/-0.1% with alternating sign.
Closes#2279
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: make GridChecksum's header-defined CRC32 kernels static
nvcc ignores the inline qualifier on __global__ functions and emits
warning #20050-D for crc32SlicedKernel and crc32CombineKernel. The
inline was doing ODR duty for a non-template kernel defined in a
header, so removing it outright would invite multiple-definition link
errors; static provides internal linkage per TU instead, matching the
effect of the anonymous namespace used by GridStats.cuh. Both launch
sites live in the same header, so nothing needs external linkage.
Note these warnings are invisible in CMake builds: CMake passes the
CUDA toolkit include directory via -isystem, and the diagnostic is
attributed through the __global__ macro machinery in those system
headers, so it is suppressed. Plain nvcc consumers of the headers see
it. With this change, TestNanoVDB.cu -- which includes essentially
every CUDA tool header -- compiles warning-free even in the
warning-visible (non -isystem) configuration.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: treat NVCC diagnostics as errors by default
Add the CMake option NANOVDB_CUDA_WERROR (default ON, also implied by
OPENVDB_CXX_STRICT), which prepends --Werror=all-warnings to
CMAKE_CUDA_FLAGS so every nvcc front-end diagnostic fails the build of
the NanoVDB tests, tools and examples. This locks in the warning-free
state: the #177-D, #20050-D and #20054-D classes fixed recently cannot
silently reappear. The option only affects NanoVDB's own build tree,
not projects that merely include the headers, and can be disabled if a
newer CUDA toolkit introduces new diagnostics.
Verified locally at CMAKE_CUDA_ARCHITECTURES=80 (matching CI) in both
Release and Debug with unit tests, tools and examples enabled: zero
diagnostics, full builds succeed, 7/7 ctests pass on a GPU.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: fix Windows fallout from NVCC warnings-as-errors
The first Windows CI run under --Werror=all-warnings surfaced two
issues:
Remove an unreferenced local in DistributedPointsToGrid's stripe
setup. deviceCoords became dead when the cudaMemAdvise call on the
external coords was removed; only the MSVC-host front end reports it
(#177-D), but the variable is dead on every platform.
Suppress the EDG dll-interface diagnostics (#1394, #1388) for CUDA
sources on Windows. This is the NVCC analog of the /wd4251 and /wd4275
suppressions in OpenVDBCXX.cmake: it's not possible to use STL types
in DLL interfaces in a portable and reliable way, so these fire on
every include of the OpenVDB core headers from a .cu file (404 + 14
instances in the CI log). Uses the top-level --diag-suppress option;
the -Xcudafe --diag_suppress spelling does not reach the CUDA
front-end phase that emits these.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: fix remaining Windows warnings-as-errors fallout
Clearing the dll-interface diagnostics unblocked the Windows projects
to compile further, exposing two more pre-existing warning classes that
NANOVDB_CUDA_WERROR now promotes to errors:
Guard the _USE_MATH_DEFINES definition in the four examples that
define it. The Windows build also passes -D_USE_MATH_DEFINES on the
command line (as 1), so the bare #define (empty) triggered MSVC C4005
macro-redefinition warnings, which the flag turns into C2220 errors.
Suppress EDG #177 for the CUDA gtest targets on Windows. gtest's TEST
macro declares each test's static test_info_ member with an
unused-attribute on GCC/Clang but not under MSVC, so for tests defined
inside an anonymous namespace (TestBuffer.cu, TestMemoryResource.cu,
TestUtilCuda.cu) the front end proves it unreferenced and reports #177
-- 76 errors, all from the macro machinery, none actionable in our
code. TestNanoVDB.cu is unaffected because its tests are declared at
namespace scope. The suppression is per-target and Windows-only, so
#177 stays enforced for all library headers, tools and examples on
every platform, and for everything on Linux.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: guard _USE_MATH_DEFINES in the example headers and CXX sources
The previous commit guarded only the nanovdb.cu files that the first
Windows log reported; each of these examples defines the macro in up to
three files, and common.h is included by the .cu sources, so the C4005
redefinition warnings persisted. Every definition site in the examples
is now guarded, verified by a tree-wide sweep for unguarded instances.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
---------
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Part of #2232 (B5), following #2268 which gave
tools::cuda::TopologyBuilder an injectable-resource seam and #2269 which
gave one to MeshToGrid.
The five classes that instantiate TopologyBuilder -- DilateGrid,
MergeGrids, PruneGrid, RefineGrid and CoarsenGrid, all in
nanovdb/tools/cuda/ -- were left behind as unaffected callers: each held
a TopologyBuilder<BuildT> constructed as mBuilder(stream), binding
default_resource<DeviceResource>() with no injection path. Downstream,
fvdb-core needs to route these operators' scratch through PyTorch's
c10::cuda::CUDACachingAllocator, and cannot delete its forked headers
until the seam exists.
Each of the five headers receives the conversion established by
MeshToGrid:
- the class template gains typename ResourceT =
nanovdb::cuda::DeviceResource, with an is_async_resource guard
static_assert first in the class body;
- the builder member becomes TopologyBuilder<BuildT, ResourceT>
mBuilder;
- the constructor(s) gain a trailing defaulted ResourceT& resource
argument, forwarded as mBuilder(stream, resource). Trailing position
keeps every existing call source-compatible. MergeGrids takes the
argument on both its N-ary and its binary-convenience constructor,
the delegating one forwarding;
- every out-of-class member definition moves to the two-parameter
form, including the dual-template getHandle.
No nested-type hoisting is needed: TopologyBuilder::Data was hoisted to
TopologyBuilderData in #2268, and none of the five classes declares its
own nested struct. The change is additive in behavior -- ResourceT
defaults everywhere, so existing callers compile unchanged with
identical allocation behavior.
Left out of scope: each operator's mBuilder.mProcessedRoot and
TopologyBuilder's internal mData are still host+device dual-space
nanovdb::cuda::DeviceBuffer, and the grid handle's output buffer still
goes through BufferT::create. None of these is observed by the injected
resource, which is expected and is noted in the new tests' comments; the
dual-space buffers are Step 3's problem.
Adds one injected-resource test per operator to TestMemoryResource.cu
(DilateGrid_InjectedResourceSeam, MergeGrids_InjectedResourceSeam,
PruneGrid_InjectedResourceSeam, RefineGrid_InjectedResourceSeam,
CoarsenGrid_InjectedResourceSeam), mirroring the existing PointsToGrid
seam tests: build a small ValueOnIndex source grid, run the operator
with a CountingResource passed as the new trailing constructor argument,
synchronize, and assert allocs > 0 && allocs == deallocs. MergeGrids
merges two disjoint grids, PruneGrid supplies a single all-on Mask<3>
sidecar over a one-leaf grid, and CoarsenGrid spreads voxels over a
2x2x2 leaf block so coarsening is non-degenerate.
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* NanoVDB: align cuda::Buffer member names with cuda::buffer
Name the free operation destroy, as cuda::buffer does, and add the
stream-taking overload it also provides. clear stays but only as a
transitional delegate, marked as such: it exists because
GridHandle::reset still calls it, and goes away with the legacy dual
buffers that cuda::Buffer replaces.
Rename setStream to set_stream. Member names cannot be aliased, so
matching the standard spelling is the whole reason the resource concept
kept allocate_async; the same argument applies here. Note in passing that
ours deliberately does not synchronize, matching cuda::buffer's
set_stream_unsynchronized rather than its set_stream, whose own
documentation and implementation disagree (NVIDIA/cccl#10649).
Add swap. The generic std::swap already does the right thing through the
move operations, but both std::vector and rmm::device_buffer provide one.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: allocate TopologyBuilder scratch from an injected resource
Eleven of the builder's buffers are device-only -- nothing reads them on
the host -- yet they were cuda::DeviceBuffer, whose host pointer and
per-device array they never use. Move them to the single-space
cuda::Buffer and give the builder a resource parameter, so its scratch is
injectable like PointsToGrid's already is, and freed by scope rather than
by hand. mProcessedRoot and mData are read on the host and stay dual.
This is not a speedup: DeviceBuffer::init allocates host or device memory
but never both, and these buffers always passed a real device id, so no
cudaMallocHost was on this path to begin with. Measured dilate on 1k/20k/
200k points, and the difference is within run-to-run noise.
Hoist the Data struct out of the class. It does not depend on the
resource, and leaving it nested would give every ResourceT its own
incompatible type for the device functors to name.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: add SyncFromAsync, and give MeshToGrid a resource seam
A stream-ordered resource must also model the synchronous concept, which
means writing four methods where two would do. The synchronous pair is
not a bare delegate -- memory from allocate must be usable on any stream
when it returns, so the null-stream allocation has to be synchronized
first -- and omitting that yields memory which satisfies the concept but
is not actually synchronous. Put it in one place rather than leaving each
author to rediscover it.
The two resources in TestMemoryResource are the first users, and were
already wrong in exactly that way: they provide only the async pair, so
they never modelled is_async_resource. TempPool duck-typed and never
checked, so nothing caught it.
MeshToGrid was the last builder allocating from a hard-wired
DeviceResource, through TempDevicePool. Give it a ResourceT parameter and
thread it into both its TopologyBuilder and its pool. As with Data in the
builder, BoxTrianglePair is hoisted out of the class: it does not depend
on the resource, and leaving it nested would give every ResourceT its own
incompatible type for the device functors to name.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: add ResourceRef and route TempPool's scratch through cuda::Buffer
Buffer holds its resource by value, matching cuda::buffer -- whose model
this completes: in CCCL the ownership semantics are selected by what is
placed in the by-value slot, an owning any_resource or a borrowing
resource_ref. We adopted the slot without the borrowing type, so a
container like TempPool, whose contract is a non-owning pointer to a
possibly stateful resource, had no way to hold a Buffer without copying
that resource and stranding its state. ResourceRef is the missing piece:
a non-owning reference that is itself a resource, so copying the ref
shares the underlying instance. Its async methods exist only when R
models AsyncResource, so a ref over a synchronous resource does not
misreport its tier, and two refs compare equal exactly when they
reference the same resource.
TempPool now keeps its bytes in a Buffer<std::byte, ResourceRef<R>>:
same resource contract, same stream retention, same discard-on-growth
reallocation, but the block is freed by ownership rather than by hand.
The TempPool unit tests, which assert traffic against the caller's own
resource instance, pass unchanged.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: note cuda::Buffer and cuda::BufferView in pendingchanges
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: express PointsToGrid's density search as a loop
The bisection search over voxel size was written as a backward goto,
which made the lifetimes of the buffers it retries over non-lexical.
Rewrite it as while(true) with continue on retry and break on
convergence; the six hand-written frees before the jump are unchanged.
The change is easiest to review with whitespace ignored, since the loop
body re-indents: git diff -w shows 27 changed lines.
d_keys and d_node_count carry results past the loop, so their
declarations move above it, as does the copy event, which was created
inside the retried region on every iteration but destroyed only once at
the end -- each retry leaked the previous handle.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: own PointsToGrid's device arrays with cuda::Buffer
Every device array PointsToGrid allocates is now owned by a
Buffer<T, ResourceRef<ResourceT>> borrowing the injected resource --
members where the pipeline frees them in a later member function than
the one that allocated them (countNodes allocates; processUpperNodes,
processLeafNodes, processPoints and processBBox release), locals where
the lifetime is contained. The raw pointers survive only as views: the
device-visible fields inside mData, and the working pointers the cub
and kernel calls take. Every owner event -- assignment, swap, destroy
-- immediately refreshes its view, and released views are nulled so a
stale use faults instead of reading a freed block. The index ping-pong
becomes a swap of owners across the member/local boundary, replacing
the bare pointer swap whose safety depended on nothing reading the
device copy of d_indx between the two uploads.
The density-search retry keeps its free-before-reallocate order via
explicit destroy calls, so peak device memory is unchanged. The
hand-matched byte sizes at every free site disappear, and the arrays
released one line before scope exit now just leave scope. One behavior
change worth naming: the too-many-points-per-leaf throw previously
leaked the reduction scratch; ownership now releases it during unwind.
Verified against the previous commit with a counting resource over
three shapes (bulk segmented-sort branch, serial per-tile branch, and
the bisection retry engaged): allocation count, free count, and total
bytes are identical. Full CUDA and memory-resource suites unchanged.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: free TopologyBuilder scratch on the caller's stream
Each release site sits in a function that receives the stream, so pass
it to destroy rather than relying on the retained stream matching --
they are the same on every current path, but the explicit form does not
depend on that staying true. Also drop the cudaGetDevice calls whose
result the Buffer conversion left unused.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: borrow TopologyBuilder scratch through ResourceRef
The scratch buffers held their resource by value, so each of the eight
carried its own copy -- fine for the stateless default, wrong for a
stateful resource, whose accounting would be split across copies while
the caller's instance saw nothing. Borrow through ResourceRef instead,
the same reconciliation TempPool uses.
Assert the stream-ordered requirement directly in TopologyBuilder and
MeshToGrid so a synchronous-only resource fails with a diagnostic that
names the builder, not just the pool inside it. Note SyncFromAsync's
synchronize cost on its allocate.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: guard PointsToGrid's copy event with a scope owner
The too-many-points-per-leaf throw unwinds past the event's manual
destroy, leaking the handle. Own it with a small guard so unwinding
releases it, consistent with the buffer ownership in this function.
Also assert the stream-ordered resource requirement on the class, so a
synchronous-only resource fails naming PointsToGrid rather than the
pool inside it.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: assert the scratch alignment TopologyBuilder relies on
The byte scratch is reinterpreted as word-sized types, which is valid
for every resource whose DEFAULT_ALIGNMENT is at least word alignment --
all CUDA allocation paths give 256 -- but nothing said so. Assert it, so
a custom resource with a weaker guarantee fails at compile time instead
of misaligning on the device.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: terminate a bare cudaCheckError with a semicolon
Legal without one -- the macro expands to a braced block -- but every
other use in the file spells it as a statement.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: set_stream matches cuda::buffer's contract, not a divergence
cuda::buffer's set_stream is deliberately non-synchronizing -- its
documented synchronization is a stale note left behind when the
synchronization was removed -- so our non-synchronizing set_stream
matches it in both name and contract, and the comment claiming a
divergence was wrong.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: mark cuda::Buffer::clear deprecated in the documentation
Documentation-level only: the [[deprecated]] attribute would warn from
GridHandle::reset and NodeManager::reset, template members in our own
headers that must keep calling clear() until every buffer type provides
destroy() -- an unactionable diagnostic for callers of reset(), and a
build break under -Werror. The attribute lands when those callers
migrate.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: exercise the concept-required members of the test resources
The trait checks detect allocate/deallocate through unevaluated contexts,
which never odr-use them, so nvcc warned that every synchronous pair and
stub member was declared but never referenced (#177-D). The attribute
route is closed -- nvcc's front end ignores [[maybe_unused]] for this
diagnostic -- so reference them the honest way: a test that verifies the
synchronous halves of the counting and stream-recording doubles behave
like their stream-ordered halves, and one that pins the trait probes'
stub behavior. The TU now builds with no warnings at all.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
---------
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
* NanoVDB: align cuda::Buffer member names with cuda::buffer
Name the free operation destroy, as cuda::buffer does, and add the
stream-taking overload it also provides. clear stays but only as a
transitional delegate, marked as such: it exists because
GridHandle::reset still calls it, and goes away with the legacy dual
buffers that cuda::Buffer replaces.
Rename setStream to set_stream. Member names cannot be aliased, so
matching the standard spelling is the whole reason the resource concept
kept allocate_async; the same argument applies here. Note in passing that
ours deliberately does not synchronize, matching cuda::buffer's
set_stream_unsynchronized rather than its set_stream, whose own
documentation and implementation disagree (NVIDIA/cccl#10649).
Add swap. The generic std::swap already does the right thing through the
move operations, but both std::vector and rmm::device_buffer provide one.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: allocate TopologyBuilder scratch from an injected resource
Eleven of the builder's buffers are device-only -- nothing reads them on
the host -- yet they were cuda::DeviceBuffer, whose host pointer and
per-device array they never use. Move them to the single-space
cuda::Buffer and give the builder a resource parameter, so its scratch is
injectable like PointsToGrid's already is, and freed by scope rather than
by hand. mProcessedRoot and mData are read on the host and stay dual.
This is not a speedup: DeviceBuffer::init allocates host or device memory
but never both, and these buffers always passed a real device id, so no
cudaMallocHost was on this path to begin with. Measured dilate on 1k/20k/
200k points, and the difference is within run-to-run noise.
Hoist the Data struct out of the class. It does not depend on the
resource, and leaving it nested would give every ResourceT its own
incompatible type for the device functors to name.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: add SyncFromAsync, and give MeshToGrid a resource seam
A stream-ordered resource must also model the synchronous concept, which
means writing four methods where two would do. The synchronous pair is
not a bare delegate -- memory from allocate must be usable on any stream
when it returns, so the null-stream allocation has to be synchronized
first -- and omitting that yields memory which satisfies the concept but
is not actually synchronous. Put it in one place rather than leaving each
author to rediscover it.
The two resources in TestMemoryResource are the first users, and were
already wrong in exactly that way: they provide only the async pair, so
they never modelled is_async_resource. TempPool duck-typed and never
checked, so nothing caught it.
MeshToGrid was the last builder allocating from a hard-wired
DeviceResource, through TempDevicePool. Give it a ResourceT parameter and
thread it into both its TopologyBuilder and its pool. As with Data in the
builder, BoxTrianglePair is hoisted out of the class: it does not depend
on the resource, and leaving it nested would give every ResourceT its own
incompatible type for the device functors to name.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: add ResourceRef and route TempPool's scratch through cuda::Buffer
Buffer holds its resource by value, matching cuda::buffer -- whose model
this completes: in CCCL the ownership semantics are selected by what is
placed in the by-value slot, an owning any_resource or a borrowing
resource_ref. We adopted the slot without the borrowing type, so a
container like TempPool, whose contract is a non-owning pointer to a
possibly stateful resource, had no way to hold a Buffer without copying
that resource and stranding its state. ResourceRef is the missing piece:
a non-owning reference that is itself a resource, so copying the ref
shares the underlying instance. Its async methods exist only when R
models AsyncResource, so a ref over a synchronous resource does not
misreport its tier, and two refs compare equal exactly when they
reference the same resource.
TempPool now keeps its bytes in a Buffer<std::byte, ResourceRef<R>>:
same resource contract, same stream retention, same discard-on-growth
reallocation, but the block is freed by ownership rather than by hand.
The TempPool unit tests, which assert traffic against the caller's own
resource instance, pass unchanged.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: note cuda::Buffer and cuda::BufferView in pendingchanges
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: free TopologyBuilder scratch on the caller's stream
Each release site sits in a function that receives the stream, so pass
it to destroy rather than relying on the retained stream matching --
they are the same on every current path, but the explicit form does not
depend on that staying true. Also drop the cudaGetDevice calls whose
result the Buffer conversion left unused.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: borrow TopologyBuilder scratch through ResourceRef
The scratch buffers held their resource by value, so each of the eight
carried its own copy -- fine for the stateless default, wrong for a
stateful resource, whose accounting would be split across copies while
the caller's instance saw nothing. Borrow through ResourceRef instead,
the same reconciliation TempPool uses.
Assert the stream-ordered requirement directly in TopologyBuilder and
MeshToGrid so a synchronous-only resource fails with a diagnostic that
names the builder, not just the pool inside it. Note SyncFromAsync's
synchronize cost on its allocate.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: assert the scratch alignment TopologyBuilder relies on
The byte scratch is reinterpreted as word-sized types, which is valid
for every resource whose DEFAULT_ALIGNMENT is at least word alignment --
all CUDA allocation paths give 256 -- but nothing said so. Assert it, so
a custom resource with a weaker guarantee fails at compile time instead
of misaligning on the device.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: set_stream matches cuda::buffer's contract, not a divergence
cuda::buffer's set_stream is deliberately non-synchronizing -- its
documented synchronization is a stale note left behind when the
synchronization was removed -- so our non-synchronizing set_stream
matches it in both name and contract, and the comment claiming a
divergence was wrong.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: mark cuda::Buffer::clear deprecated in the documentation
Documentation-level only: the [[deprecated]] attribute would warn from
GridHandle::reset and NodeManager::reset, template members in our own
headers that must keep calling clear() until every buffer type provides
destroy() -- an unactionable diagnostic for callers of reset(), and a
build break under -Werror. The attribute lands when those callers
migrate.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: exercise the concept-required members of the test resources
The trait checks detect allocate/deallocate through unevaluated contexts,
which never odr-use them, so nvcc warned that every synchronous pair and
stub member was declared but never referenced (#177-D). The attribute
route is closed -- nvcc's front end ignores [[maybe_unused]] for this
diagnostic -- so reference them the honest way: a test that verifies the
synchronous halves of the counting and stream-recording doubles behave
like their stream-ordered halves, and one that pins the trait probes'
stub behavior. The TU now builds with no warnings at all.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
---------
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
* Balance the multi-GPU merge tree across every device
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Generalize merge to non power-of-two GPU counts
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Remove prefetches causing performance regressions
Signed-off-by: Matthew Cong <mcong@nvidia.com>
---------
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* NanoVDB: handle empty partitions in the distributed merge path search
mergePath's bounds checks compare against (keys1Count - 1) and
(keys2Count - 1) on size_t counts, so a zero count underflows to
SIZE_MAX and the guards never fire; the binary search then reads the
empty array (and past either array) and produces wrong or garbage
intervals. An empty partition is reachable from radixSortAsync's
ceil-split whenever a device receives no items, e.g. sorting a single
element across two devices.
Return the trivial split before the search: with an empty side, every
element up to the diagonal comes from the non-empty side.
New single-GPU test MergePathEmptyPartition covers the empty-right,
empty-left, both-empty, and non-empty control cases; the empty cases
fail deterministically without the fix (inverted split / underflow
garbage) and run clean under compute-sanitizer with it.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: compute the empty-partition merge split on the host
An empty side needs no merge path search, so skip the kernel launch:
the trivial split (everything up to the diagonal comes from the
non-empty side) is computed on the host by mergePathTrivial at the
launch site. mergePath reverts to the unguarded search and documents
that both inputs must be non-empty.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
---------
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
* NanoVDB: align cuda::Buffer member names with cuda::buffer
Name the free operation destroy, as cuda::buffer does, and add the
stream-taking overload it also provides. clear stays but only as a
transitional delegate, marked as such: it exists because
GridHandle::reset still calls it, and goes away with the legacy dual
buffers that cuda::Buffer replaces.
Rename setStream to set_stream. Member names cannot be aliased, so
matching the standard spelling is the whole reason the resource concept
kept allocate_async; the same argument applies here. Note in passing that
ours deliberately does not synchronize, matching cuda::buffer's
set_stream_unsynchronized rather than its set_stream, whose own
documentation and implementation disagree (NVIDIA/cccl#10649).
Add swap. The generic std::swap already does the right thing through the
move operations, but both std::vector and rmm::device_buffer provide one.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: allocate TopologyBuilder scratch from an injected resource
Eleven of the builder's buffers are device-only -- nothing reads them on
the host -- yet they were cuda::DeviceBuffer, whose host pointer and
per-device array they never use. Move them to the single-space
cuda::Buffer and give the builder a resource parameter, so its scratch is
injectable like PointsToGrid's already is, and freed by scope rather than
by hand. mProcessedRoot and mData are read on the host and stay dual.
This is not a speedup: DeviceBuffer::init allocates host or device memory
but never both, and these buffers always passed a real device id, so no
cudaMallocHost was on this path to begin with. Measured dilate on 1k/20k/
200k points, and the difference is within run-to-run noise.
Hoist the Data struct out of the class. It does not depend on the
resource, and leaving it nested would give every ResourceT its own
incompatible type for the device functors to name.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: note cuda::Buffer and cuda::BufferView in pendingchanges
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: free TopologyBuilder scratch on the caller's stream
Each release site sits in a function that receives the stream, so pass
it to destroy rather than relying on the retained stream matching --
they are the same on every current path, but the explicit form does not
depend on that staying true. Also drop the cudaGetDevice calls whose
result the Buffer conversion left unused.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: set_stream matches cuda::buffer's contract, not a divergence
cuda::buffer's set_stream is deliberately non-synchronizing -- its
documented synchronization is a stale note left behind when the
synchronization was removed -- so our non-synchronizing set_stream
matches it in both name and contract, and the comment claiming a
divergence was wrong.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: mark cuda::Buffer::clear deprecated in the documentation
Documentation-level only: the [[deprecated]] attribute would warn from
GridHandle::reset and NodeManager::reset, template members in our own
headers that must keep calling clear() until every buffer type provides
destroy() -- an unactionable diagnostic for callers of reset(), and a
build break under -Werror. The attribute lands when those callers
migrate.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: exercise the concept-required members of the test resources
The trait checks detect allocate/deallocate through unevaluated contexts,
which never odr-use them, so nvcc warned that every synchronous pair and
stub member was declared but never referenced (#177-D). The attribute
route is closed -- nvcc's front end ignores [[maybe_unused]] for this
diagnostic -- so reference them the honest way: a test that verifies the
synchronous halves of the counting and stream-recording doubles behave
like their stream-ordered halves, and one that pins the trait probes'
stub behavior. The TU now builds with no warnings at all.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
---------
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
* Fix empty chunks in the multi-GPU points-to-grid pipeline
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Fix race condition with prefix update and incorrect stream in process root
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Remove unsupported cudaMemAdvise call on memory pool allocations
Signed-off-by: Matthew Cong <mcong@nvidia.com>
---------
Signed-off-by: Matthew Cong <mcong@nvidia.com>
Review follow-ups from the Copilot pass on this PR.
The lifetime event only covers uses issued through deviceUpload/deviceDownload; work enqueued
against the raw pointer from deviceData() is invisible to it, so on a non-blocking stream such
work could still outlive the buffer (as on master, where no API could express the dependency).
Make recordUse public so raw-pointer callers can register their stream, and document the
contract on deviceData().
freeDeviceBuffers issued every device's free on one stream, but a stream belongs to a single
device; free each allocation on its own device instead (caller's stream for the current device,
the owning device's default stream otherwise, switching devices as needed). Same pre-existing
shape as master's loop; correctness on multi-GPU is by inspection - this machine has one GPU.
Split the free-ordering regression test into the two scenarios that discriminate the two
historical bugs: a blocking stream with an unregistered write (fails if frees move off the
default stream: the reviewed revision) and a non-blocking stream with a recordUse-registered
write (fails without the event: master). Verified per revision: current passes both, the
reviewed revision fails both, master fails the non-blocking one.
Add NonBlockingStreamSignedFloodFill, mirroring the dilation regression: identical output on
the default stream and on a non-blocking stream while the default stream is occupied, with the
readback on the producing stream. This locks the cross-stream contract; it cannot discriminate
processRoot's internal ordering on constructible inputs (verified: the pre-fix header passes
it) because that path only does work when interior root-level tiles exist and pre-fix code was
accidentally host-synchronous otherwise - noted in the test.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Two non-blocking review points from harrism.
Self-move assignment freed the buffers and then read its own (now dangling) members. Verified:
without the guard 'buf = std::move(buf)' aborts on the deviceData assertion because mDeviceCount
has been zeroed; with it the buffer is left intact. Note UnifiedBuffer::operator=(&&) has the
same flaw and is left for a separate change.
Also document clear()'s stream argument. The review noted it had become dead, which was true of
the revision reviewed - the per-device stream overrode it. With frees now ordered by the tracking
event the argument is honoured again and simply selects where the free is enqueued, so any stream
is safe to pass regardless of where the buffer was used.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
The existing DeviceBufferNonBlockingStreamLifetime test only asserts that repeated
allocate/use/free returns cudaSuccess, so it passes with or without a correctly ordered free.
Its comment also described the superseded 'free on the stream it was allocated on' mechanism.
Reword it as the smoke test it is, and add a test that actually discriminates.
DeviceBufferMultiStreamFreeOrdering uses a device-only buffer on purpose - it owns no pinned
host memory, so clear()'s cudaFreeHost (which implicitly synchronizes) cannot mask the problem.
A second stream is parked behind a busy-wait with a write to the buffer queued behind it, the
buffer is destroyed, and the next allocation is checked for corruption. Verified to fail against
the pre-fix DeviceBuffer (64 MB clobbered) and pass with the fix.
A warm-up pass is required: on a cold context the first launches serialize and the race never
forms. The recycled/still-pending state is reported in the failure message but deliberately not
used as a preconditon - with a correctly ordered free the allocator cannot hand the block out
again until the other stream drains, so those observations are the fix working rather than a
reason to skip.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
The entry described the mechanism as freeing on the stream the buffer was last used on,
which was superseded: frees are now ordered after every stream the buffer was used on via a
per-device event.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Tracking only the stream a device buffer was LAST used on is not sufficient: two streams can
both have work in flight, and freeing on the later one leaves the earlier one's work running
against memory the allocator has already recycled. Reported in review by matthewdcong.
Reproduced: a device-only DeviceBuffer plus a user kernel on a second blocking stream, with a
warm-up pass so the launches actually overlap. The freed block is recycled and all 64 MB of the
next allocation are overwritten by the late kernel. Freeing on the default stream (the previous
behaviour) passes, because the legacy stream is implicitly ordered after every blocking stream;
freeing on the last-used stream fails.
Replace the per-device stream with a per-device event. Each use waits on the event before
issuing work and re-records it afterwards, so one event transitively covers every stream the
buffer has been used on, and the frees wait on it. This is correct for any mix of blocking and
non-blocking streams and for any number of them, and it removes the requirement that callers
keep their streams alive until the buffer is destroyed, since the event is owned by the buffer.
Verified: both the multi-blocking-stream case and the non-blocking case pass, and the VBM,
topology and maintenance goldens are unchanged (46/46 on dragon and emu).
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* NanoVDB: encode points for any resource in PointsToGrid (CUDA)
The point-encode step in tools::cuda::PointsToGrid was only reachable for
the default DeviceResource, because it lived in the BuildT=Point full
specialization of processPoints (i.e. <Point, DeviceResource>). A Point
grid built with a custom resource fell through to the generic no-op and
was silently left with unencoded point data.
Fold the encode into the generic processPoints and gate it with
if constexpr(is_same<BuildT, Point>). The encode kernels run on device
data and are independent of the resource; only the trailing d_indx
deallocation routes through ResourceT. Non-Point builds and default-
resource Point builds are unchanged.
Add a regression test that builds a NanoGrid<Point> through a custom
resource and asserts every input point is recoverable from the encoded
per-voxel data.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: move the point-encode kernels out of the if-constexpr block
nvcc rejects an extended lambda defined directly inside the block of an
if constexpr statement on some host compilers, so the nine encode
launches failed to compile against MSVC while building fine with gcc and
clang. Give them their own member function and have the branch call it.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
---------
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
* NanoVDB: fail loudly when CUDA memory pools are unavailable
Fractional vGPU configurations commonly expose no stream-ordered memory
pools (cudaDevAttrMemoryPoolsSupported == 0), so the first cudaMallocAsync
fails with cudaErrorNotSupported. Rather than silently substituting
synchronous cudaMalloc -- which would make an async resource misrepresent
its own semantics, since the choice of a synchronous backend belongs to
the caller -- the util::cuda::mallocAsync wrapper now detects the missing
capability and emits an actionable diagnostic naming
NANOVDB_USE_SYNC_CUDA_MALLOC before returning cudaErrorNotSupported, which
every caller already surfaces via cudaCheck. freeAsync mirrors the check.
The NANOVDB_USE_SYNC_CUDA_MALLOC macro remains the caller's explicit
opt-in to synchronous allocation, and must be defined identically in every
translation unit (the wrappers are inline). memoryPoolsSupported() exposes
the cached per-device capability query, and mallocSync/freeSync name the
synchronous pair the macro path maps to.
TestUtilCuda.cu is compiled into two binaries -- default and
macro-defined -- covering both modes; its async-path tests skip on a device
without memory pools, where those wrappers fail by design.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: drop redundant grid-build test from TestUtilCuda
GridBuildThroughWrappers built a device grid through the wrappers, but
PointsToGrid device builds are already covered by TestMemoryResource,
TestNanoVDB, and TestMultiGPU. Its only unique coverage was the
macro-forced-synchronous build, which is better established by the
MallocAsyncRoundTrip test running in the NANOVDB_USE_SYNC_CUDA_MALLOC
binary plus validation on pool-less hardware -- and it pulled the
PointsToGrid builder header into a util-wrapper unit test. Remove it and
the now-unused include.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: inline cudaMalloc/cudaFree instead of mallocSync/freeSync helpers
mallocSync/freeSync were the shared synchronous implementation for two
callers: the macro tier and a runtime fallback. With the fallback removed
in favor of failing loudly, only the macro tier remains, so the helpers
are one-line pass-throughs with a single caller each. Inline cudaMalloc /
cudaFree directly, restoring the macro-tier wrappers to their prior form
and leaving memoryPoolsSupported as the only added function.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: clarify the memory-pool diagnostic names a build-time macro
Say "Define NANOVDB_USE_SYNC_CUDA_MALLOC when building" so the message
reads as a compile-time option rather than a runtime environment variable
that would be respected.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: hoist the device query out of its if-statement condition
Declaring the error code in the if-init clause packed a query, a test
and an early return onto one line. Split them so each reads on its own.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
---------
Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Ken Museth <1495380+kmuseth@users.noreply.github.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Co-authored-by: Ken Museth <1495380+kmuseth@users.noreply.github.com>
* NanoVDB: add cuda::Buffer / cuda::BufferView and the synchronous resource trait (CUDA)
Adds nanovdb/cuda/Buffer.h with two typed containers:
- cuda::Buffer<T, R>: owning, move-only container of T elements allocated
from a memory resource R held by value. With a stream-ordered resource
(is_async_resource) it retains the stream of the most recent allocation
and orders deallocation, resize, and copy on it; with a synchronous
resource (is_resource) it carries no stream state and exposes no stream
API, enforced via an empty-base StreamHolder and enable_if-gated
overloads. Constructors value-initialize elements to zero unless the
NoInit tag is passed. Element counts whose byte size would overflow
size_t are rejected with an exception.
- cuda::BufferView<T>: non-owning, trivially copyable view with span
semantics plus a detaching clear(), which is what the GridHandle buffer
static interface requires, enabling zero-copy
GridHandle<cuda::BufferView<std::byte>> over externally owned grids.
Also adds the is_resource detection trait (synchronous allocate/deallocate)
to nanovdb/cuda/DeviceResource.h alongside is_async_resource, and a new
unit-test suite (TestBuffer.cu, ctest name nanovdb_cuda_buffer_unit_test)
covering traits, counting and stream-recording resources, stream
retention, resize, move semantics, deep copy, pinned-host buffers, and a
zero-copy GridHandle round trip.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: free the new block when a synchronous Buffer resize fails
The stream-ordered resize guarded its prefix copy so a failure frees the
new allocation and leaves the buffer untouched; the synchronous overload
had no such guard. Mirror it so both overloads share one invariant.
Also document parameters with @param rather than describing them in
@brief, split the stream-ordered c-tor's initializer list onto one
member per line, and note that copy construction and assignment are
explicitly disallowed.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
---------
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Per review feedback, the converted release-note entries now spell out the
complete historical signature (including parameters) rather than a bare
name, so readers still see exactly what each change referred to.
Uses <TT> rather than @c for these, since @c only formats the next
whitespace-delimited word and cannot span a signature with arguments.
Follows the existing <TT> precedent in this file.
Note that argument-qualified @vdblink/@link targets do not resolve under
the versioned inline namespace even when the API still exists, so plain
text is required here regardless of whether the symbol was removed.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* addressed bug in nanovdb::ReadAccesor<0,1,2>
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* fixed build of vdb_tool with OpenEXR on OSX
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* changed default to off
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* updated nanovdb.txt
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* improved openvdb_cmd/vdb_tool/CMakeLists.txt
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* Clean up vdb_tool OpenEXR/Imath CMake logic.
Signed-off-by: Andre Pradhana <andre.pradhana@gmail.com>
* Simplify Old ReadAccessor CMake option and ensures users who include NanoVDB.h directly get the new ReadAccessor behavior by default.
Signed-off-by: Andre Pradhana <andre.pradhana@gmail.com>
* reviewer feedback
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* improved writing
Signed-off-by: Ken Museth <ken.museth@gmail.com>
---------
Signed-off-by: Ken Museth <ken.museth@gmail.com>
Signed-off-by: Andre Pradhana <andre.pradhana@gmail.com>
Co-authored-by: Andre Pradhana <andre.pradhana@gmail.com>
- Removed APIs now use @c with fully qualified historical names, such as @c
openvdb::Grid::signedFloodFill().
- Restored names omitted entirely by the PR, including the Transform methods and
Filter::laplacian().
- Kept createLevelSetSphere() as a corrected @vdblink because it still exists.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Such as `TBB_INTERFACE_VERSION`.
Without it, the `#else` statement always triggers, breaking build for
newer TBB version (albeit rare at this point).
Signed-off-by: Maciej Małecki <me@mmalecki.com>
When either constexpr-if condition in the evalFirstCached lambda
evaluates to true, the remaining code including the return statement
is unreachable. MSVC emits warning C4702 (unreachable code) in this
case.
Add else keywords to form a proper constexpr-if / else-if / else
chain so that the unreachable code path is excluded from the
instantiated function body.
Fixes#2239
Signed-off-by: Mustafa Senoglu <mmustafasenoglu0@gmail.com>
crc32CombineKernel rebuilt its two GF(2) shift operators by binary
exponentiation on a single device thread every call - a fixed ~480 us that
dominated the checksum for small and medium grids (profiling showed it as large
as the entire tail-streaming pass). The operators depend only on the chunk
length, so they are now built on the host in microseconds and uploaded, and the
device kernel is reduced to the O(chunkCount) fold.
Bit-identical: all checksum goldens pass on dragon/emu/crawler/wdas_cloud.
A/B vs the parallel-checksum base (RTX PRO 6000, 3 rounds):
dragon +40.7% (1.69x) emu +15.3% crawler +9.8% wdas_cloud +1.9%
The gain shrinks with grid size because it removes a fixed cost; transient
memory unchanged.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* fixed cmake for vdb_tool to use OpenEXR v3
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* added min version to cmake for imath
Signed-off-by: Ken Museth <ken.museth@gmail.com>
---------
Signed-off-by: Ken Museth <ken.museth@gmail.com>
Addresses PR review on the cooperative GridStats:
- processLeaf is now launched only when nodeCount[0] > 0, matching the guards on
the internal-node launches (a leaf-less grid would otherwise be a <<<0,...>>>
launch).
- The average path now gives each node its own d_stats slot instead of reusing
one of its children's. A fully-tiled internal node (active value tiles, no
children) never claimed a child slot, so the old code wrote d_stats[-1] (out
of bounds) and propagated an invalid slot to its parent. d_stats is now sized
for all nodes (leaves, then lower, then upper) and each node writes slot
leafCount + levelOffset + nodeID; the sSlot/atomicMax machinery is removed.
stats_minmax is unaffected (it never touches d_stats) and stays byte-exact;
stats_all output is unchanged for grids without childless internal nodes and
correct for those that have them. compute-sanitizer memcheck clean.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Port of Agent A's exp/20. Consecutive threads now process consecutive table
entries and leaf values - the former mapping gave each thread a consecutive RUN
(a 32-entry stride across the warp for internal nodes, 8 for leaves) on every
iteration. The 4 KB upper-node value/child masks are copied cooperatively across
the block instead of member-wise by thread 0 while the rest idled.
Combines with the Tranche-1 scoped zero-init already on this file: the leaf
kernel's stats/padding-gap zeroing is preserved inside the thread-0 block, and
the coalesced value loop still writes every mValues[i], so output stays
bit-identical - all indextogrid goldens pass on dragon/emu/crawler/wdas_cloud.
A/B vs the Tranche-1 base (RTX PRO 6000, 3 rounds): dragon +23%, emu +21%,
crawler +18%, wdas_cloud +14%; transient memory unchanged. The remap gain is on
top of Tranche 1's scoped init, so absolute indextogrid is now faster than both
master and A's exp/20.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Port of Agent A's exp/13. Replaces the thread-per-node serial reductions:
- processLeaf: one warp per leaf - lanes stride the 512 voxel slots (mask-gated)
and merge partial statistics through shared memory; lane 0 keeps the bbox
update and the final store.
- processInternal: one 128-thread block per internal node striding the
4,096/32,768-entry child table with a shared-memory tree reduction of Stats +
bbox. The former kernel gave each thread up to 32,768 serial child visits and
ran an entire lower level of a 181M-voxel grid on ~9 warps of a 188-SM GPU.
min/max are order-independent, so stats_minmax output is bit-identical - all
stats_minmax goldens pass on dragon/emu/crawler/wdas_cloud. avg/stddev
(stats_all) use a Welford-correct parallel merge but are NOT bit-identical to
the serial result (floating-point reduction order differs); stats_all stays a
bench-only measurement, never a byte-exact claim.
A/B vs the Tranche-1 base (RTX PRO 6000, 3 rounds):
stats_minmax dragon +87.3% emu +82.5% crawler +80.3% wdas +83.5% (5-8x)
stats_all dragon +71.9% emu +51.7% crawler +29.8% wdas +35.1%
transient memory unchanged.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Port of Agent A's exp/14. Replaces the two serial anti-patterns in blockedCRC32
with parallel equivalents:
- Per-block CRC via slicing-by-4: the 256-entry base LUT is staged into shared
memory and three derived slice tables are built in place, so the dependent CRC
chain advances four bytes per step with four shared-memory lookups instead of
one byte per L2 lookup. (Not __constant__: divergent indices serialize there.)
- The final fold (a single thread CRCing the whole block-CRC array - megabytes
for multi-GB grids) becomes parallel per-chunk CRCs plus a GF(2) crc32_combine
(zlib construction: one shift operator per fixed chunk length, applied per
fold). Small inputs keep the single-thread path.
Checksum VALUES are bit-identical (slicing-by-4 and the GF(2) combine are both
value-preserving); all checksum_full goldens pass on dragon/emu/crawler/
wdas_cloud. The pre-v32.6.0 crc32TailOld path is untouched.
A/B vs the Tranche-1 base (RTX PRO 6000, 3 rounds): dragon +84.7%, emu +92.0%,
crawler +93.2%, wdas_cloud +94.1% (~6.5x / 12.5x / 15x / 17x); transient memory
unchanged.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Addresses PR review:
- PointsToGrid/MeshToGrid/DistributedPointsToGrid use std::min in the grid-name
copy but did not include <algorithm> (it compiled only via transitive
includes). Add the include so each header is self-sufficient.
- DeviceBuffer comments still described mStreams as the "allocation stream", but
it is now updated on every deviceUpload/deviceDownload; reword them to
"last-used stream" to match the actual behavior.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Addresses PR review: mStreams[device] was set only when the device allocation
was first created. If the buffer is later re-uploaded (or downloaded) on a
different stream, the destructor/move-assign/clear would still free on the old
stream, which does not necessarily order after the most recent work on the new
stream (notably for cudaStreamNonBlocking). Update mStreams[device] to the
current stream on every host<->device copy in deviceUpload/deviceDownload, not
just on first allocation. Guarded on mStreams (null for externally-managed
buffers, which are never freed here).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Addresses PR review: the test launches streamBusyWaitKernel on the default
stream to occupy it during the candidate dilation, but only synchronized the
non-blocking stream afterward. The busy-wait could remain queued on stream 0
and leak into subsequent tests, causing timing-dependent flakes. Drain it with
a cudaDeviceSynchronize() after the candidate run.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Addresses PR review: the Tranche-1 stream-ordered-free change stored a single
mStream, overwritten by every DeviceBuffer::deviceUpload. A DeviceBuffer can
hold an allocation on each device (mGpuData is per-device), and the multi-GPU
example (ex_make_mgpu_nanovdb) uploads one handle to every device with that
device's own stream. The destructor/move-assign then freed *every* device's
allocation on the last device's stream - a cross-device cudaFreeAsync(ptr@devA,
stream@devB), which is invalid (the free stream must belong to the allocation's
device). Pre-Tranche-1 code freed on stream 0 uniformly, so this was a
regression for the multi-GPU path.
Track the allocation stream per device in a cudaStream_t array parallel to
mGpuData, set in init()/deviceUpload() at mStreams[device], and free each
mGpuData[i] on its own mStreams[i] in clear()/move-assignment. The single-GPU
path is unchanged (one device, one stream); the multi-GPU path now frees each
device's allocation on its own device-bound stream.
Single-GPU verified (DeviceBuffer lifetime gtest + byte-exact validation); the
multi-GPU free path is correct by construction but not runtime-tested here
(single-GPU environment).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
The Tranche-1 stream-correctness fix added an unconditional
cudaStreamSynchronize(stream) before processRoot's synchronous default-stream
cudaMemcpy of the tree. That sync is only needed when the caller's node passes
ran on a non-default stream; on the default stream (0) the blocking copy already
orders after prior stream-0 work, so the extra sync there was pure overhead - a
measured ~4-11% regression on floodfill/dragon vs master. Guard it with
`if (stream != 0)`. Correctness for non-blocking-stream callers is unchanged.
floodfill is now NEUTRAL vs master on dragon/emu/crawler (0 regressions), and
byte-exact validation is unaffected.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Addresses PR review:
- The grid-name copy could leave GridData::mGridName without a null terminator
when the name length >= MaxNameSize: min(size+1, MaxNameSize) == MaxNameSize
overwrote the entire zeroed field, and gridName() (a C-string accessor) would
then read past it. Copy at most MaxNameSize-1 bytes and rely on the preceding
memset for termination, truncating over-long names instead. Fixed in
PointsToGrid, MeshToGrid and DistributedPointsToGrid.
- GridName_CudaPointsToGrid gains an over-long (>= MaxNameSize) case asserting
the result is truncated to MaxNameSize-1 chars and stays null-terminated.
- NonBlockingStreamDilate_ValueOnIndex queried device 0 for the clock rate;
query the current device via cudaGetDevice() instead.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
The Tranche-1 determinism fix zero-initialized the *entire* IndexToGrid output
buffer so that bytes the build kernels never write - stats fields absent from
the source grid, and struct alignment padding - are deterministic instead of
carrying recycled pool bytes (which otherwise leak into written .nvdb files).
Measured against master this cost 10-22% on indextogrid across dragon/emu/
crawler/wdas_cloud: the leaf array dominates the buffer and was memset to zero
and then immediately overwritten value-by-value - a redundant multi-GB pass.
Scope the initialization to the bytes that are genuinely left unwritten:
* getBuffer() memsets only the non-leaf region [0, node[0]) - grid, tree,
root, root tiles and the internal nodes - a small fraction of the buffer.
* processLeafsKernel zeroes each leaf's own gap [&mMinimum, mValues): the
stats fields (skipped when the source has no stats) and any padding before
the 32-aligned mValues array. Every mValues[i] is written anyway, so this
reproduces the former full zero-init byte-for-byte.
Output is bit-identical to the full zero-init - all indextogrid goldens pass on
dragon/emu/crawler/wdas_cloud, compute-sanitizer memcheck is clean and initcheck
shows no new reports - and the regression vs master is gone (NEUTRAL on all four
datasets, RTX PRO 6000 / sm_120).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Five correctness/determinism fixes to the NanoVDB CUDA tools:
* GridChecksum: evalChecksum copied sizeof(GridData)+sizeof(TreeData) (736 B)
from a device buffer holding a single uint32 CRC into the 4-byte head/tail of
a stack Checksum -- an out-of-bounds device read plus host stack corruption
(compute-sanitizer visible; a hard CUDA "invalid argument" on some configs).
Copy sizeof(uint32_t) at all four sites, matching validateChecksum.
* Stream propagation: the NN_FACE_EDGE_VERTEX leaf-dilation kernel and
SignedFloodFill's root pass ran on the default stream instead of the caller's
stream, silently corrupting output for callers using a non-blocking stream.
Launch and synchronize on the operation's stream.
* Bounded grid-name copy: the grid-name copy read a fixed MaxNameSize bytes from
a possibly shorter std::string (whose .data() is never null, making the memset
branch dead), over-reading up to 255 bytes of host heap into the grid and any
written .nvdb file. Zero the field, then copy only min(name.size()+1,
MaxNameSize) bytes. Applied to PointsToGrid, DistributedPointsToGrid, MeshToGrid.
* Deterministic output initialization: PointsToGrid and IndexToGrid left
alignment padding and stats fields carrying recycled pool bytes, making output
nondeterministic and leaking host heap into files. Zero the whole grid buffer
once after allocation. For PointsToGrid this makes the three dense
background-fill passes (upper/lower value tables, and the inactive-leaf-values
pass for all but the Point build type) redundant, so they are dropped.
* DeviceBuffer stream-ordered lifetime: the destructor, move-assignment and
clear() freed device allocations on the default stream regardless of the
stream the buffer was last used on -- a latent use-after-free for
non-blocking-stream callers, masked today by legacy-stream semantics. Track the
allocation stream and free on it, matching the stream-ordered free TempPool
now performs.
Adds gtest coverage in unittest/TestNanoVDB.cu: cross-stream dilation
determinism, grid-name round-trip (empty/short/max-length), PointsToGrid output
determinism, and DeviceBuffer non-blocking-stream lifetime. No performance claims.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
GL/glu.h uses APIENTRY without defining it. This previously compiled because
oneTBB (<= 2022.3) transitively included windows.h from its public headers;
oneTBB 2023.0 removed that include, so APIENTRY is no longer defined when
RenderModules.cc includes GL/glu.h. GL/glew.h cannot provide it either: it
undefines APIENTRY at the end of the header when it was the one defining it
(and the vcpkg glew package additionally patches out glew.h's internal
GL/glu.h include). Include windows.h explicitly before GL/glu.h instead of
relying on a transitive include.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* NanoVDB: add stream-ordered async memory-resource seam (CUDA)
Introduce an instance-based async memory-resource concept on NanoVDB's CUDA
path so downstream code can inject its own allocator without forking headers,
with zero change to existing call sites at the default.
- cuda::DeviceResource: add instance allocate_async/deallocate_async alongside
the existing static methods; add default_resource<R>() (a program-lifetime
singleton) and the is_async_resource<R> trait.
- cuda::PinnedResource: new host-pinned resource modelling the same concept
(default impl wraps cudaMallocHost/cudaFreeHost).
- cuda::TempPool: route allocations through a resource instance and free on the
stream of the most recent reallocate() at destruction, not the null stream.
- tools::cuda::PointsToGrid: route internal scratch and mDeviceData through an
injected resource instance via a defaulted ResourceT& constructor argument.
- unittest: add TestMemoryResource.cu (TestMemoryResource suite) with
resource-concept, pinned page-lock, stream-retained-free, no-leak, and
PointsToGrid instance-injection tests.
TopologyBuilder/DistributedPointsToGrid routing and the dual-buffer pinned-host
fix allocate via cuda::DeviceBuffer and are deferred to a follow-up that
introduces a resource-parameterized buffer.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
* NanoVDB: make cuda::PinnedResource synchronous; document is_async_resource
cudaMallocHost / cudaFreeHost are synchronous, so PinnedResource now
models the synchronous Resource concept -- allocate(bytes, alignment) /
deallocate(ptr, bytes, alignment), no stream -- instead of carrying the
stream-ordered AsyncResource shape with an ignored stream argument. The
static allocateAsync / deallocateAsync helpers are removed; the type is
allocator-stateless, so the instance methods call CUDA directly.
is_async_resource<PinnedResource> is now false; the unit test asserts
that and exercises the synchronous allocate / deallocate API.
Add a @code example to is_async_resource documenting the if-constexpr
dispatch between a stream-ordered allocate_async and a synchronous
allocate.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
---------
Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
tools::cuda::MergeGrids merged exactly two device OnIndex grids, but its
algorithm is already an accumulate-into-shared-masks union, so it extends
to any number of inputs without changing the kernels.
- Add an N-ary constructor taking std::vector<const GridT*>; the binary
(grid1, grid2) constructor now delegates to it, leaving current callers
unchanged. Loop over the inputs in mergeRoot / mergeInternalNodes /
mergeLeafNodes instead of handling two grids explicitly; guard the
empty list.
- A single N-ary pass uses one accumulator and reads each input once,
versus materializing log(N) intermediate grids when chaining binaries.
Add TestNanoVDBCUDA.MergeGridsNary_ValueOnIndex: verify a 3-grid merge
matches chaining the binary form (and edge cases: single-element copy,
empty throws).
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
- UT_CameraParms changes
- Assume Uniform option in Filter to prevent any anistropy being
introduced.
- format override to pretty print ut types
- Sync of the reference files
Signed-off-by: Jeff Lait <jlait@andorra.sidefx.com>
With the changes.txt, Half.h and ComputePrimitives.h warnings fixed, the
per-pattern suppression layer has served its purpose. Switch the CI docs build
to Doxygen's native FAIL_ON_WARNINGS so any net-new warning fails the PR
directly, and remove the filter scaffolding:
* .github/workflows/docs.yml passes -DOPENVDB_DOXYGEN_WARN_AS_ERROR=ON and
drops the check_doxygen_warnings step. The warning log is still uploaded as
an artifact (FAIL_ON_WARNINGS writes it before exiting non-zero).
* ci/check_doxygen_warnings.py and doc/.doxygen_warning_suppress.txt removed.
* doc/CMakeLists.txt option comment updated to note CI now uses the gate.
NOTE: the docs job stays red until the remaining openvdb_ax/ast/AST.h
INLINE_INHERITED_MEMB warning (handled separately) is merged into this branch
-- it is the only warning the build still emits.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
The cxx14 namespace was a C++11 forward-compat reimplementation of
std::index_sequence / std::make_index_sequence. NanoVDB requires C++17
(MINIMUM_CXX_STANDARD 17), so the standard versions in <utility> (already
included) are guaranteed present, and openvdb/Types.h already uses them.
Switching to std:: deletes the recursive-inheritance template
(make_index_sequence<N> : make_index_sequence<N-1, ...>) that Doxygen's
class-relation analyser flagged as a "potential recursive class relation".
Doxygen documents template entities rather than instantiations, so the
recursive base resolves back to the same class and looks like a self-loop;
it can't see that the <0, Is...> specialization terminates the recursion.
Removing the shim eliminates the warning at the source for every compiler,
not just the CI Doxygen, so its suppression entry is dropped too. The only
remaining backlog entry is the AST.h INLINE_INHERITED_MEMB interaction.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Whittles two more entries off doc/.doxygen_warning_suppress.txt by fixing
the underlying warnings instead of suppressing them.
doc/changes.txt (26 unresolved-link warnings -> 0):
* Historical Release Notes \link references to APIs that were since renamed
or removed are delinked to @c <name> (the file's existing inline-code
style), or dropped to plain prose where the link text was an English word
woven into a sentence. createLevelSetSphere() is requalified tools:: (it
only lacked the namespace, matching the working link elsewhere in the file).
doc/CMakeLists.txt:
* EXCLUDE_SYMBOLS no longer carries a capital-I "*Internal*" glob: it also
matched the public tree::InternalNode / nanovdb::InternalNode classes,
silently dropping them from the docs and breaking InternalNode::childCount
in changes.txt while its RootNode/LeafNode siblings resolved. Both classes
are now documented and the link resolves.
* OPENVDB_API is predefined empty so Doxygen stops folding the export macro
into the name of `enum OPENVDB_API FromBitsTag` in the vendored Half.h.
ext/imath/openvdb/math/Half.h (vendored, OpenVDB-adapted Imath fork):
* @file half.h -> @file Half.h (the file was renamed during vendoring).
* The internal half-float limit/config macros are wrapped in @cond/@endcond:
Doxygen auto-links the literal `#define` token on documented active macros,
and these are implementation details (the half type lives in the excluded
internal namespace).
* The file-level description now renders cleanly: its ASCII bit-layout
diagrams, formulas and example tables are wrapped in <pre>, and its
double-backtick code spans are reduced to single backticks (the double
form broke Doxygen 1.14's markdown paragraphing and collapsed the prose).
doc/.doxygen_warning_suppress.txt:
* Drop the changes.txt and Half.h entries; both are now warning-clean.
Remaining backlog: the ComputePrimitives.h recursive-template false
positive and the AST.h INLINE_INHERITED_MEMB interaction.
Verified against Doxygen 1.14 (the CI version): 0 surviving warnings.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Add MeshToGrid: CUDA triangle mesh to NanoVDB IndexGrid/UDF converter
Introduces nanovdb::tools::cuda::MeshToGrid<BuildT>, a GPU-accelerated
rasterizer that converts a triangle soup (vertex list + index list) into
a sparse NanoVDB grid on the device. Two output modes are implemented:
- getHandle() — produces a ValueOnIndex (topology-only) grid
- getHandleAndUDF() — produces the same grid plus a device sidecar
buffer of per-active-voxel unsigned distances
(world-space UDF, ordered by active voxel index)
SDF (signed distance field) output is planned future work.
The pipeline uses a hierarchical top-down subdivision approach:
1. transformTriangles() — vertices transformed to NanoVDB index space
via nanovdb::Map::applyInverseMap()
2. processRootTrianglePairs() — enumerates (root tile, triangle) pairs whose
padded AABBs overlap, via CUB prefix-sum +
scatter (no atomics)
3. processLeafTrianglePairs() — 3-level 8x subdivision (4096→512→64→8),
refining the pair list at each level using
triangle-AABB / SAT intersection tests
(evaluateAndCountSubBoxesKernel: 1 CTA per
pair, 512 threads, warp ballot + CUB reduce)
4. Topology assembly — unique leaf origins extracted, NanoVDB
ValueOnIndex tree built via TopologyBuilder,
empty leaves pruned via PruneGrid
5. UDF computation (optional) — ComputeUDFFunctor computes per-voxel
point-to-triangle distances via atomic-min,
finalized to world-space in a second pass
The example driver (ex_mesh_to_grid_cuda) reads an OBJ file, builds an
OpenVDB UDF as reference, and runs a thorough correctness validation:
- Active voxel set comparison (false negatives, false positive rate)
- Per-voxel UDF error histogram vs. OpenVDB reference (in voxel units)
- CPU brute-force ground-truth check on the worst-error voxels
New files:
nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh
nanovdb/nanovdb/math/Proximity.h — point/triangle proximity utils
nanovdb/nanovdb/util/cuda/Rasterization.cuh — ComputeUDFFunctor and helpers
Modified files:
nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh — adds InitGridTreeRootFunctor
for constructing grids from scratch (no source grid to copy metadata from)
nanovdb/nanovdb/util/cuda/Util.h — adds operatorKernelInstance for launching
pre-constructed device functors carrying data members
nanovdb/nanovdb/examples/CMakeLists.txt — registers ex_mesh_to_grid_cuda
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* Formatting changes
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* Remove trailing whitespace in MeshToGrid.cuh and mesh_to_grid_cuda.cpp
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* Formatting changes in Rasterization.cuh
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* Handle empty triangle list gracefully in MeshToGrid
Replace throws-on-zero with early returns throughout the pipeline so
that getHandle() and getHandleAndUDF() return a valid empty grid when
triangleCount == 0. Kernel launches over null sets are explicitly
guarded before any buffer pointer is dereferenced.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* Add unit tests for MeshToGrid: empty mesh and unit tetrahedron
MeshToGrid_EmptyMesh: verifies that getHandle() returns a valid grid
with zero active voxels when the triangle list is empty.
MeshToGrid_UnitTetrahedron: constructs a unit tetrahedron (4 vertices,
4 faces) with voxel size 0.1. Tests both getHandle() (topology-only)
and getHandleAndUDF() (with sidecar). Checks:
- Topology: checksum from getHandle() equals that from getHandleAndUDF(),
confirming both paths produce identical grids.
- Per-voxel UDF: sidecar values match CPU brute-force (index-space
pointToTriangleDistSqr * voxelSize) within 1e-3 voxels.
- No false positives: every active voxel lies within the narrow band.
- No false negatives: every voxel with CPU UDF strictly inside the band
is active (checked via leaf-iteration active set).
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* Slim down example driver: cleanup and leaner OpenVDB comparison
mesh_to_grid_cuda.cpp:
- Remove unused #include <openvdb/tools/Morphology.h>
- Fix cout label "openvdb_triangles" -> "triangles"
- Replace atof(argv[3]) with std::stof
- Remove spurious blank lines
mesh_to_grid_cuda_kernels.cu:
- Remove unused #include <array> (leftover from std::array<Vec3f,3> era)
- Collapse triple blank line after getHandleAndUDF() to single
- Replace the detailed per-voxel histogram, top-N worst-voxel heap,
and inline Ericson brute-force with a compact summary: active voxel
counts, false negative/positive counts, and a single UDF quality
line (max error + count exceeding 0.1 voxels vs OpenVDB reference)
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* MeshToGrid: remove verbose>=2 diagnostic prints
All mVerbose>=2 guarded timer restarts and printf diagnostics have been
removed. Verbosity level 1 (function-scale start/stop timing) is retained.
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* MeshToGrid: rename pool/gridPool params to buffer; remove diagnostic accessors
- getHandle: parameter renamed from pool to buffer, matching the
convention used throughout NanoVDB tools (PointsToGrid, DilateGrid, etc.)
- getHandleAndUDF: parameters renamed from gridPool/sidecarPool to
buffer/sidecarBuffer for the same reason
- Local raw-buffer variables renamed to gridBuffer to avoid shadowing
the (now renamed) buffer parameter
- Remove getPairCount(), getDevicePairs(), getRootTileCount(), and
getDeviceRootOrigins() — unused diagnostic-only accessors
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* ci: trigger CI
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* MeshToGrid: fix leaf-local X↔Z swap in CUDA rasterization kernels
Four kernels were decoding the per-thread leaf offset with the wrong bit
ordering (lx = vi & 7, lz = vi >> 6), transposing X and Z relative to
NanoVDB's Mask<3> convention. The kernels were self-consistent, so the
generated grid was a topologically valid X↔Z mirror of the input mesh's
UDF rather than a crash. Replace all four inline decodes with
NanoLeaf<BuildT>::OffsetToLocalCoord:
- Rasterization.cuh: RasterizeLeafNodesFunctor, ComputeUDFFunctor
- MeshToGrid.cuh: evaluateAndCountSubBoxesKernel, ScatterChildPairsFunctor
Also fix the matching inline decodes in the unit test and example
verifiers, which were silently cancelling the bug on the consumer side.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* Rasterization.cuh: qualify ::atomicOr to bypass namespace collision
The RasterizeLeafNodesFunctor lives in nanovdb::util::rasterization::cuda.
Unqualified atomicOr lookup from inside an innermost namespace named cuda
confuses nvcc's name resolution on CUDA 12.6 (clang/gcc on Linux): it
picks up an atomicOr(uint64_t*, uint64_t) overload from libcu++/cuda::
and then rejects the reinterpret_cast<unsigned long long*> argument with
"argument of type unsigned long long * is incompatible with parameter
of type uint64_t * (aka unsigned long *)". Qualifying as ::atomicOr
forces global-namespace lookup and resolves to the canonical CUDA
builtin atomicOr(unsigned long long*, unsigned long long).
Also pull the reinterpret_cast onto its own line into a typed pointer
variable, matching the pattern used by Morphology.cuh.
Mask::setOnAtomic in NanoVDB.h is unaffected because it lives directly
in nanovdb, not in a nested cuda namespace.
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
---------
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
Motivated by issue #2161 (docs drift away from code), this change wires
in a layered set of gates so the same class of rot fails on the PR that
introduces it.
CI gates (new files):
* ci/validate_docs_references.py runs on every PR via a new validate_docs
job. It scans doc/**/*.md and doc/**/*.txt and verifies that every
#include <...> inside a fenced code block (Markdown or Doxygen @code)
resolves to a real header in this checkout, and that every
github.com/AcademySoftwareFoundation/openvdb/blob/<ref>/<path> URL
points at a file on disk. Catches the exact pattern from #2161 (e.g.
nanovdb/util/CudaDeviceBuffer.h -> nanovdb/cuda/DeviceBuffer.h).
* ci/check_doxygen_warnings.py filters the Doxygen warning log against
doc/.doxygen_warning_suppress.txt and fails the build on any surviving
warning. Replaces Doxygen's all-or-nothing WARN_AS_ERROR with a
per-pattern gate that honours a documented backlog without blocking
new regressions. Reports suppression hit counts so dead patterns are
flagged automatically.
* doc/.doxygen_warning_suppress.txt is the explicit triage backlog: four
narrowly-scoped, heavily-commented regex patterns (doc/changes.txt
historical \link rot, vendored ext/imath/Half.h, a Doxygen recursive
variadic template false positive in ex_util/ComputePrimitives.h, and
the INLINE_INHERITED_MEMB + unnamed-override-param interaction in
openvdb_ax/ast/AST.h).
* .github/workflows/docs.yml runs the validator first, then doxygen,
then the warning check (which fails the job on any unsuppressed
warning), and always uploads the doxygen.warnings.log as an artifact.
Drops '**.md' from paths-ignore so doc-only PRs are still gated.
Single-source-of-truth fix for the bug in #2161:
* doc/nanovdb/HelloWorld.md and four backing example sources
(ex_openvdb_to_nanovdb_accessor, ex_read_nanovdb_sphere_accessor,
ex_read_nanovdb_sphere_accessor_cuda) now use Doxygen's \snippet
mechanism instead of duplicating code into markdown. The Hello World
rendered on openvdb.org will always match what is built and tested in
CI. The exact namespace drift from #2161 is now structurally
impossible to reintroduce.
* doc/CMakeLists.txt sets DOXYGEN_EXAMPLE_PATH/EXAMPLE_RECURSIVE so
\snippet paths resolve, captures DOXYGEN_WARN_LOGFILE, and adds an
OPENVDB_DOXYGEN_WARN_AS_ERROR option for local opt-in to Doxygen's
native all-or-nothing gate.
Doxygen warning backlog (44 of 45 surviving warnings fixed in this PR):
* nanovdb/GridHandle.h - bogus @param verbose removed; 3x
@param filename -> @param fileName; @param pool -> @param other.
* nanovdb/NanoVDB.h - 3x @param ijk -> @param xyz on
Map::apply*Jacobian* methods; fixed @node -> @note typo.
* nanovdb/util/PrefixSum.h - added missing @param op documentation.
* openvdb/TypeList.h - backtick-wrapped 'std::get<I>(mTuple)' so
Doxygen stops parsing <I> as an unterminated HTML italics tag.
* openvdb/tools/RayTracer.h - added explicit @brief docs to seven
shader-subclass operator() overrides so the inherited @param nml/dir
no longer conflict with the override signatures.
* openvdb/version.h.in - @file path now matches the actual filename.
* openvdb_houdini/ParmFactory.h - 2x \link references qualified with
the houdini_utils:: namespace so they resolve from class scope.
* doc/ax/doc.txt - deleted orphan </div> with no matching <div>.
* doc/nanovdb/SourceTree.md - closed the missing code fence; replaced
'foo@bar:~$' bash prompt (Doxygen parses @bar as a command) with
the unlabeled '$' style used elsewhere in doc/nanovdb/.
* doc/nanovdb/FAQ.md - corrected link to nanovdb/cuda/DeviceBuffer.h
(was util/CudaDeviceBuffer.h, a sibling of #2161).
LaTeX -> MathJax migration:
* DOXYGEN_USE_MATHJAX YES + MATHJAX_VERSION MathJax_3 in
doc/CMakeLists.txt switches formula rendering from latex/dvips/PNG
to MathJax. Removes the install_latex step from docs.yml and lets
contributors build the docs locally without any TeX install (the
conda-forge texlive-core package on osx-arm64 is incomplete and
has no working alternative on that channel). AX docs use @f$/@f[
math; output is now browser-rendered, selectable, and crisper.
Result: 73 raw Doxygen warnings, 73 suppressed by documented patterns,
0 surviving. The strict gate is live and enforces 'no new warnings'
from the next PR onward.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
The Python binding for nanovdb.io.writeGrids iterated the handle list
and called io::writeGrid(fileName, ...) per element. That file-level
overload opens the path with std::ios::trunc, so every iteration wiped
the previously written grid and the resulting .nvdb only ever contained
the last handle in the list (same story for deviceWriteGrids on the
CUDA side).
Mirror io::writeGrids(fileName, handles, codec, verbose) from
nanovdb/io/IO.h directly in the wrapper: open the output stream once
with trunc, then call the ostream overload of writeGrid for each handle
cast from the nb::list. We can't simply forward to io::writeGrids
because GridHandle<BufferT> is move-only (see GridHandle.h) and the
nb::list holds references we must not invalidate.
Tests (nanovdb/python/test/TestNanoVDB.py):
- TestGridHandleExchange.test_list_to_vector now reads the written
file back via readGridMetaData and readGrids and asserts the full
handle count is preserved.
- Added TestDeviceReadWriteGrids.test_device_write_grids_multi as a
CUDA-guarded regression for deviceWriteGrids with multiple handles.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* NanoVDB: add CPU port of VoxelBlockManager and supporting utilities
- util/MaskPrefixSum.h (new): bit-parallel prefix sum over NanoVDB leaf
masks via a three-pass (X/Y/Z) algorithm using SIMD-friendly byte
transpose and in-place shuffle-down scan
- tools/VoxelBlockManager.h (new): CPU port of VoxelBlockManager, sharing
VoxelBlockManagerBase with the CUDA implementation; buildVoxelBlockManager
uses nanovdb::util::forEach for portable multi-threaded execution
(TBB or std::thread fallback, no OpenMP runtime dependency)
- tools/cuda/VoxelBlockManager.cuh: refactor to inherit VoxelBlockManagerBase;
simplify buildVoxelBlockManager to a single-call API; rename template
parameter to Log2BlockWidth
- util/Util.h: add NANOVDB_RESTRICT macro and portable __hostdev__
util::atomicOr / util::atomicAnd, dispatching to CUDA intrinsics,
std::atomic_ref (C++20), __atomic_fetch_* (GCC/Clang), or
_Interlocked*64 (MSVC)
- NanoVDB.h: promote Mask::setOnAtomic/setOffAtomic/setAtomic from
__device__-only to __hostdev__ using util::atomicOr/util::atomicAnd;
bump patch version to 1
- util/MorphologyHelpers.h, util/cuda/Morphology.cuh: portability fixes
- unittest/TestNanoVDB.cc: add CPU VoxelBlockManager test
- unittest/TestNanoVDB.cu: simplify and improve CUDA VoxelBlockManager test
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* VoxelBlockManager: use initializer list in move constructor
Replace std::exchange calls in the move constructor body with proper
member initialization in the initializer list, zeroing out the source
members in the body.
The move assignment operator retains std::exchange, which remains the
right idiom there since assignment operators have no initializer list.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* VoxelBlockManager: document decodeInverseMaps arguments
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* VoxelBlockManager: document VoxelBlockManagerHandle constructor arguments
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
* MaskPrefixSum: introduce zRow reference in Z-pass prefix sum loop
Signed-off-by: Efty Sifakis <esifakis@gmail.com>
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
---------
Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
Signed-off-by: Efty Sifakis <esifakis@gmail.com>
Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
Improvements:
- CNanoVDB supports integer (I) grids
- CNanoVDB has setValue methods to write to voxels
- CNanoVDB has getLeaf methods to acquire raw leaf pointers
Bug Fixes:
- static used for local kernel methods in CNanoVDB
- Remove auto-computed reserved sizes in favour of hard coded sizes
in CNanoVDB. While less maintainble, C99 forbids zero sized
arrays so some compilers crashed when encountering this.
- Initialize key value when constructing an accessor, some compilers
entered bad states with this undefined.
Signed-off-by: Jeff Lait <jlait@andorra.sidefx.com>
CI: Enable NanoVDB Python Tests
Enable build of nanovdb python module and running python tests in NanoVDB CI action to make sure we catch any regressions.
* Using segmented radix sort for higher tile counts for optimized voxel key computation.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Switch handrolled binary search to thrust::upper_bound
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Implemented segmented radix sort in DistributedPointsToGrid
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Add tests for FewTiles and ManyTiles scenarios in PointsToGrid
- Implemented unit tests for the serial per-tile sort path with fewer than 32 tiles.
- Added tests for the segmented sort path handling 32 or more tiles.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* patch version update
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* pnano version bump
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
- Implemented isGpuAvailable function in NanoVDBModule.cc to determine if a CUDA-capable GPU is available at runtime.
- Updated TestNanoVDB.py to include tests that skip if no CUDA-capable GPU is detected.
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
- Added a new root directory, ext/ to house imported third-party code
- Moved Imath Half.h into ext/ such that it's clear it is third party.
- Added doxygen-awesome css/js extensions
- Fixes and improvements to doxygen generation. Leverage new doxygen-awesome styles
- Add a THIRD-PARTY.md license document to ext/ to detail all imported licenses (copied from OSL)
Signed-off-by: Nick Avramoussis <4256455+Idclip@users.noreply.github.com>
- Added customization of the output transformation attribute for points::pca. Can be output as a quaternion or as a combined affine transform
- Added support for quaternions attributes and affine transformations (that include scale) in PointRasterizeSDF for ellipsoidal rasterizations.
Signed-off-by: Nick Avramoussis <4256455+Idclip@users.noreply.github.com>
This adds the latest round of improvements to NanoVDB from the fVDB team, including:
Improvements and bugfixes to PNanoVDB
Add Mat classes and other math types to NanoVDB
Add support for morphology operators in NanoVDB
Improvements to mGPU grid building and introduction of a separate test executable for mGPU tests
Several fixes for compiler warnings
This PR fixes a few discrepancies between grid construction via (a) PointsToGrid, and (b) OpenToNanoVDB. The discrepancies are not typically inhibiting functionality of either construction, but result in grids that are not bit-identical, when they would otherwise be expected to be. The changes are:
Removal of the PointsToGridData::flags member. The actual functionality of this flag was to allow control of solely the leaf-level flags (interior node, and GridData flags were already set independently), and in that aspect it was even incorrectly setting the IsBreadthFirst bit (which only makes sense for GridData::mFlags, not the leaf node flags). The method includeBBox was having no effect as it was, it was never called in any of the unittests, and its effect was overwritten later in the code.
PointsToGrid has been corrected to flag HasBBox as active in every node (not just the leaf nodes, as was happening prior; the flags were previously being set to zero).
The world space bounding box in OpenToNanoVDB was incrementing the upper bound of the index-space bounding box, prior to applying the transform ( hence, a grid with a single voxel [0,0,0] would have a WS bounding box of [0,0,0]->[1,1,1]). PointsToGrid was not applying this increment, resulting in a smaller WS bounding box.
Signed-off-by: Matthew Cong <mcong@nvidia.com>
The CALL_CUBS macro has a slightly different definition in the two similar code files that use it, namely PointsToGrid.cuh and TopologyBuilder.cuh. This fix restricts the scope of each macro definition to just the code file where it is used.
Signed-off-by: Matthew Cong <mcong@nvidia.com>
Refactor the distributed radix sort code into a separate function so that it is more modular, testable, reusable. Furthermore, it fixes a small rounding error edge case in the merge code.
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Separate data from algorithm
* Add custom resource support
* Add docs
* Add support for cudax::async_resource
* Fix default alignment
* Add cleanup
* Split up DeviceResource into separate file
* Fix file
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Adding VoxelBlockManager with base functionality
* switching from deprecated CountOn implementation
* Unit test for basic VoxelBlockManager use
* fixing trailing whitespace
* fixing trailing whitespace
* Fixing bug in get() method, that prevents proper operation for other than leaf nodes
* Fix licenses
* Add stream parameter
* Verify linear access
* More cleanup
* Cleanup constants
* Remove header
* Add comment
* Expose NumThreads for VBM build
* Fix up comments
* Avoid reinterpret_cast and use cuda::atomic_ref
* Clean up params and add docs
* replaced or commented on cryptic constants
* Add named constants
* More cleanup
---------
Co-authored-by: Matthew Cong <mcong@nvidia.com>
Signed-off-by: Matthew Cong <mcong@nvidia.com>
- Added OPENVDB_HAS_INVOKABLE_MEMBER_FUNCTION and OPENVDB_HAS_MEMBER_FUNCTION
macros for compile time checking of member functions
- Updated points::rasterize signature to no longer accept interrupter or
filters and instead delegate these to the transfer scheme. Added transfer
modules to support them
- Added support for transfer schemes to provide a rasterizePoints which takes
a range of points
- Updated PCA to support target and source point per voxel counts
- Remove dependency on PCA in ellips rasterizer
- Added projected ellips distance and improved axis-aligned ellips calculation kernels
for better representations of SDF distance calculations.
- Made the pws value to the ellips rasterizer optional
- Removed templated interrupter type from all rasterize methods
- Removed old ABI<=9 support from PCA methods
- Move to new OPENVDB_ASSERT macro
- Use #if instead of #ifdefs to ensure they are present
Signed-off-by: Nick Avramoussis <4256455+Idclip@users.noreply.github.com>
Non trivial changes:
- Env flag to disable stream writing (backported to 19.50)
- H21 style camera options
- OP_InputIdx and OP_OutputIdx support.
Signed-off-by: Jeff Lait <jlait@andorra.sidefx.com>
There only appeared to be one obvious case left.
Code search process:
grep -C 3 delete .
... then look for if statements, in C-like code, conditioned on the same
variable, and no other obvious purpose that required the if statement.
Signed-off-by: Eric Powers <eric.powers@dreamworks.com>
* Fix memory leaks in PointsToGrid tests
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Fix leapfrogging across recursion levels
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Parallelize kernel dispatch for different levels
Signed-off-by: Matthew Cong <mcong@nvidia.com>
---------
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Start support for more than two GPUs in key merge
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Flip flop storage
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Generalize concurrent leftIntervals/rightIntervals usage
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Fix median search
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Fix rebalancing
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Fix recursive MGPU merge
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Update TODO
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Add guards for zero-tile GPUs
Signed-off-by: Matthew Cong <mcong@nvidia.com>
---------
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Adding /bigobj to TestNanoVDB.cu compilation options to fix Windows build failure
---------
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Add distributed implementation of PointsToGrid
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Fix race condition in merge and single GPU case
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Fix kernel/cudaFree race conditions
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Add DistributedPointsToGrid unittest
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Clean up example
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Use range-based for loops
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Use structured bindings inside for loop
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Fix Windows build and some warnings
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Update for refactored DeviceMesh
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Fix copyright/include
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Add parallelForEach helper
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Templatize kernels to align with CUB better
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Shorten TemporaryDevicePool to TempDevicePool
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Refactor to use TempDevicePool
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Switch to class
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Parallelize kernel dispatch in build
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Fix race condition wrt to root node processing
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Speed up synchronization
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Add comments and fix deprecated TransformInputIterator
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Address review comments
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Address more review comments
* Add more cudaCheck calls and restrict EqualityIndicator type
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Fix license identifiers and remove unused code
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Add missing include
Signed-off-by: Matthew Cong <mcong@nvidia.com>
---------
Signed-off-by: Matthew Cong <mcong@nvidia.com>
Co-authored-by: = <=>
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Add device mesh
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Restore current device in constructor and destructor
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Enable range-based for loops
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Use structured bindings inside for loop
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Remove uneeded accessors and modularize functions
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Statically initialize entry point
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Encapsulate
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Switch to lightweight wrapper class for DeviceNodes vector
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Cleanup
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Minor changes needed for DistributedPointsToGrid
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Add RAII DeviceGuard
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* More cleanup
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Clean up NCCL dependency
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Remove parallel_for to modularize
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Implement and test move constructor/assignment
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Add docs
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Expand test and documentation
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Move non-reinit comment to function description
Co-authored-by: Mark Harris <783069+harrism@users.noreply.github.com>
* Switch to using aliases and modern type names
Co-authored-by: Mark Harris <783069+harrism@users.noreply.github.com>
* Use size_type instead of int for return
Co-authored-by: Mark Harris <783069+harrism@users.noreply.github.com>
* Separate out DeviceGuard and clean up API
Signed-off-by: Matthew Cong <mcong@nvidia.com>
---------
Signed-off-by: Matthew Cong <mcong@nvidia.com>
Co-authored-by: Mark Harris <783069+harrism@users.noreply.github.com>
Signed-off-by: Matthew Cong <mcong@nvidia.com>
Adds Various level set constructors:
* tools::createLevelSetCapsule creates a grid containing a narrow-band level set representation of a capsule (tube with constant radius and sphere caps).
* tools::createLevelSetTaperedCapsule creates a grid containing a narrow-band level set representation of a tapered capsule (tube with sphere caps and different radii at both ends, or equivalently the convex hull of two spheres with possibly different centers and radii).
* tools::createLevelSetTubeComplex creates a grid containing a narrow-band level set representation of a tube complex (a collection of capsules or tapered capsules defined by endpoint coordinates and segment indices).
* tools::createLevelSetDilatedMesh creates a grid containing a narrow-band level set representation of a dilated surface mesh (dilated by a radius in all directions).
Base class that can be used to create other level sets:
* tools::ConvexVoxelizer, which can be inherited to implement a convex level set constructor.
Signed-off-by: ghurstunither <62885595+ghurstunither@users.noreply.github.com>
it doesn't do an in-place transpose.
Thus the line
mat.transpose(); // the 3x3 in nanovdb is transposed relative to openvdb's 3x3
is a no-op, and not actually transposing the matrix. Since we DO
transpose when we do an OpenToNano; it is important we transpose here
or we can't round-trip.
Unfortunately, it has been this way since the initial commit, so it is
quite likely tooling has already worked around this....
Signed-off-by: Jeff Lait <jlait@andorra.sidefx.com>
* Merge changes from latest NanoVDB
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* Fix comments and trailing include
Signed-off-by: Matthew Cong <mcong@nvidia.com>
---------
Signed-off-by: Matthew Cong <mcong@nvidia.com>
* added read/write support of OFF files to vdb_tool
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* added pendingchanges/vdb_tool.txt
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* added patch from Jonathan S
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* minor cleanup
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* addressed review comments
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* before major improvements to readPLY
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* cleanup
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* removed whitespace
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* added nanovdb::vdb_tool::swapBytes
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* added more documentation
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* patch by Jonathan S to fix CI
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* Fix gtest install for rpm. not deb
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Building gtest from source; newer version required than available on distro
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* permissions fix
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Attempting to fix doxygen node version issue
Missing iomanip include in TestNanoVDB
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* Bump container versions attempt to fix doxygen test
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* install_latex fix
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
* added ascii option to Geometry::writePLY
Signed-off-by: Ken Museth <ken.museth@gmail.com>
* fixing issue introduced in the previous commit
Signed-off-by: Ken Museth <ken.museth@gmail.com>
---------
Signed-off-by: Ken Museth <ken.museth@gmail.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Co-authored-by: Jonathan Swartz <jonathan@jswartz.info>
* nonLeafCount and treeDepth return Index64 in TreeBase.
* All ABI breakages safeguarded for versions < 12.
* Deprecated implementations of nodeCount for vector<Index32> arguments.
Signed-off-by: ghurstunither <62885595+ghurstunither@users.noreply.github.com>
leafCount, unallocatedLeafCount, nodeCount now work with Index64 to accommodate trees with more than 2^32-1 leaf nodes.
Signed-off-by: ghurstunither <62885595+ghurstunither@users.noreply.github.com>
* Uses any available pdal plugin to read input file
* Writes RGB into Cd channel of points (also extended to PTS implementation)
* Adds a requirement on PDAL, so hidden behind a feature VDB_TOOL_USE_PDAL - off by default
Signed-off-by: Tom Matterson <tommatterson@gmail.com>
# Fix error: Cannot prepare internal mirrorlist: No URLs in mirrorlist. CentOS 8 reached EOL means need to replace the official mirror to vault.centos.org
# Comment out mirrorlist and replace #baseurl=...mirror.centos.org with baseurl=...vault.centos.org in files starting with CentOS- in /etc/yum.repos.d folder
sed -i 's/mirrorlist/#mirrorlist/g' /etc/yum.repos.d/CentOS-*
sed -i 's|#baseurl=http://mirror.centos.org|baseurl=http://vault.centos.org|g' /etc/yum.repos.d/CentOS-*
# Fix error: Cannot prepare internal mirrorlist: No URLs in mirrorlist. CentOS 8 reached EOL means need to replace the official mirror to vault.centos.org
# Comment out mirrorlist and replace #baseurl=...mirror.centos.org with baseurl=...vault.centos.org in files starting with CentOS- in /etc/yum.repos.d folder
sed -i 's/mirrorlist/#mirrorlist/g' /etc/yum.repos.d/CentOS-*
sed -i 's|#baseurl=http://mirror.centos.org|baseurl=http://vault.centos.org|g' /etc/yum.repos.d/CentOS-*
- {name:'asan', build:'asan', components:'core,test,axcore,axtest', cmake:'-DNANOVDB_USE_OPENVDB=ON -DOPENVDB_AX_STATIC=OFF -DOPENVDB_CORE_STATIC=OFF -DUSE_BLOSC=OFF'}# We never called blosc_destroy(), so disable blosc to silence these errors
- {name:'asan', image:'2026', build:'asan', components:'core,test', cmake:'-DDISABLE_DEPENDENCY_VERSION_CHECKS=ON -DNANOVDB_USE_OPENVDB=ON -DOPENVDB_AX_STATIC=OFF -DOPENVDB_CORE_STATIC=OFF -DUSE_BLOSC=OFF'}# We never called blosc_destroy(), so disable blosc to silence these errors
@@ -22,7 +22,7 @@ open source software best practice policies of the Linux Foundation.
#### License
OpenVDB is licensed under the [Mozilla Public License, version 2.0](LICENSE.md)
OpenVDB is licensed under the [Apache License, version 2.0](LICENSE.md)
license. Contributions to OpenVDB should abide by that standard
license.
@@ -54,9 +54,13 @@ Every commit must be signed off. That is, every commit log message
must include a “`Signed-off-by`” line (generated, for example, with
“`git commit --signoff`”), indicating that the committer wrote the
code and has the right to release it under the
[Mozilla Public License, version 2.0](LICENSE.md)
[Apache License, version 2.0](LICENSE.md)
license. See the [TAC documentation on contribution sign off](https://github.com/AcademySoftwareFoundation/tac/blob/master/process/contributing.md#contribution-sign-off) for more information on this requirement.
#### Stale PRs
A PR that does not have a signed CLA or DCO sign-off will be given at least one week in which to resolve the issue and then will be closed. The TSC cannot review or merge PRs without these requirements having been met. Please re-open PRs that are now expected to pass once this criteria has been met.
## Maintainer
The maintainer role is the equivalent of the "Committer" role in the charter.
OpenVDB is an open source C++ library comprising a novel hierarchical data
@@ -28,14 +28,14 @@ generally less stable than the [production releases](https://github.com/AcademyS
### License
OpenVDB is released under the [Mozilla Public License Version 2.0](https://www.mozilla.org/MPL/2.0/),
OpenVDB is released under the [Apache License, Version 2.0](https://www.apache.org/licenses/LICENSE-2.0),
which is a free, open source software license developed and maintained by the
Mozilla Foundation.
Apache Software Foundation.
The trademarks of any contributor to this project may not be used in
association with the project without the contributor's express permission.
NOTE: OpenVDB is in the process of changing its license from [Mozilla Public License Version 2.0](https://www.mozilla.org/MPL/2.0/) to [Apache License, Version 2.0](https://www.apache.org/licenses/LICENSE-2.0)! Please see the file RE-LICENSE_NOTE.txt for more details.
NOTE: OpenVDB has completed re-licensing from [Mozilla Public License Version 2.0](https://www.mozilla.org/MPL/2.0/) to [Apache License, Version 2.0](https://www.apache.org/licenses/LICENSE-2.0). Please see the file RE-LICENSE_NOTE.txt for more details.
No, NanoVDB is by design agnostic and works with most graphics APIs, e.g. OpenGL, OpenCL, OptiX, CUDA, HLSL, GLSL, DirectX, and even the CPU! The data structure itself (defined in [NanoVDB.h](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/NanoVDB.h) and [PNanoVDB.h](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/PNanoVDB.h)) only depends on C++11 or C99. However, some of the auxiliary math classes are optionally optimized with intrinsic CUDA calls (e.g. fma in Vec3) and some of the [Allocator](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/util/CudaDeviceBuffer.h) uses CUDA for explicit memory management. It should be straightforward to replace (or remove) these CUDA dependencies if you so desire.
No, NanoVDB is by design agnostic and works with most graphics APIs, e.g. OpenGL, OpenCL, OptiX, CUDA, HLSL, GLSL, DirectX, and even the CPU! The data structure itself (defined in [NanoVDB.h](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/NanoVDB.h) and [PNanoVDB.h](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/PNanoVDB.h)) only depends on C++11 or C99. However, some of the auxiliary math classes are optionally optimized with intrinsic CUDA calls (e.g. fma in Vec3) and the [DeviceBuffer](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/cuda/DeviceBuffer.h) allocator uses CUDA for explicit memory management. It should be straightforward to replace (or remove) these CUDA dependencies if you so desire.
* **How does the data structure of NanoVDB compare to that of OpenVDB?**
# NanoVDB Hello World Examples {#NanoVDB_HelloWorld}
[The first example](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/examples/ex_openvdb_to_nanovdb_accessor/openvdb_to_nanovdb_accessor.cc) shows how to convert an OpenVDB level set sphere into a NanoVDB level set, use accessors to print out multiple values from both grids, and save the NanoVDB grid to file. Note that **this example depends on both OpenVDB and NanoVDB**.
The example snippets below are pulled directly from compiled example source
files via Doxygen's `\snippet` mechanism, so the code shown here is always
in sync with what is built and tested in CI. Enable
`-DNANOVDB_BUILD_EXAMPLES=ON` to build them locally.
```cpp
#include <openvdb/tools/LevelSetSphere.h> // replace with your own dependencies for generating the OpenVDB grid
#include <nanovdb/util/CreateNanoGrid.h> // converter from OpenVDB to NanoVDB (includes NanoVDB.h and GridManager.h)
#include <nanovdb/util/IO.h>
## Convert an OpenVDB grid to NanoVDB and use accessors
// Convert an openvdb level set sphere into a nanovdb, use accessors to print out multiple values from both
// grids and save the NanoVDB grid to file.
// Note, main depends on BOTH OpenVDB and NanoVDB.
int main()
{
try {
// Create an OpenVDB grid (here a level set surface but replace this with your own code)
auto srcGrid = openvdb::tools::createLevelSetSphere<openvdb::FloatGrid>(100.0f, openvdb::Vec3f(0.0f), 1.0f);
[The first example](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/examples/ex_openvdb_to_nanovdb_accessor/openvdb_to_nanovdb_accessor.cc)
shows how to convert an OpenVDB level set sphere into a NanoVDB level set, use
accessors to print out multiple values from both grids, and save the NanoVDB
grid to file. Note that **this example depends on both OpenVDB and NanoVDB**.
// Convert the OpenVDB grid, srcGrid, into a NanoVDB grid handle.
// Define a (raw) pointer to the NanoVDB grid on the host. Note we match the value type of the srcGrid!
auto* dstGrid = handle.grid<float>();
## Read a NanoVDB grid from a file (host only)
if (!dstGrid)
throw std::runtime_error("GridHandle does not contain a grid with value type float");
[The second example](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/examples/ex_read_nanovdb_sphere_accessor/read_nanovdb_sphere_accessor.cc)
reads a NanoVDB grid from a file (the one saved in the previous example) and
prints out multiple values. Note that **this example does not depend on
OpenVDB (nor CUDA), only NanoVDB**.
// Get accessors for the two grids. Note that accessors only accelerate repeated access!
[The third example](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/examples/ex_read_nanovdb_sphere_accessor_cuda/read_nanovdb_sphere_accessor_cuda.cu)
reads a NanoVDB grid from a file (the one saved in the first example) and
prints out multiple values on both the CPU and GPU. Note that **this example
does NOT depend on OpenVDB, only NanoVDB and CUDA**.
[The second example](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/examples/ex_read_nanovdb_sphere_accessor/read_nanovdb_sphere_accessor.cc) reads a NanoVDB grid from a file (the one saved in the previous example) and prints out multiple values. Note that **this example does not depend on OpenVDB (nor CUDA), only NanoVDB**.
#include <nanovdb/util/IO.h> // this is required to read (and write) NanoVDB files on the host
NVCC requires the CUDA kernel to be defined in
[a separate .cu file](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/examples/ex_read_nanovdb_sphere_accessor_cuda/read_nanovdb_sphere_accessor_cuda_kernel.cu):
/// @brief Read a NanoVDB grid from a file and print out multiple values.
///
/// @note Note This example does NOT depend on OpenVDB (nor CUDA), only NanoVDB.
int main()
{
try {
auto handle = nanovdb::io::readGrid("data/sphere.nvdb"); // reads first grid from file
auto* grid = handle.grid<float>(); // get a (raw) pointer to a NanoVDB grid of value type float
if (!grid)
throw std::runtime_error("File did not contain a grid with value type float");
auto acc = grid->getAccessor(); // create an accessor for fast access to multiple values
for (int i = 97; i < 104; ++i) {
printf("(%3i,0,0) NanoVDB cpu: % -4.2f\n", i, acc.getValue(nanovdb::Coord(i, 0, 0)));
[The third example](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/examples/ex_read_nanovdb_sphere_accessor_cuda/read_nanovdb_sphere_accessor_cuda.cc) reads a NanoVDB grid from a file (the one saved in the first example) and prints out multiple values on both the CPU and GPU. Note that **this example does NOT depend on OpenVDB, only NanoVDB and CUDA**.
```cpp
#include <nanovdb/util/IO.h> // this is required to read (and write) NanoVDB files on the host
#include <nanovdb/util/cuda/CudaDeviceBuffer.h> // required for CUDA memory management
NVCC requires the CUDA kernel to be defined in [a separate .cu file](https://github.com/AcademySoftwareFoundation/openvdb/blob/master/nanovdb/nanovdb/examples/ex_read_nanovdb_sphere_accessor_cuda/read_nanovdb_sphere_accessor_cuda.cu):
```cpp
#include <nanovdb/NanoVDB.h> // this defined the core tree data structure of NanoVDB accessible on both the host and device
This is a C++ class library for using the Single Instruction Multiple Data (SIMD) instructions to improve performance on modern microprocessors with the x86 or x86/64 instruction set on Windows, Linux, and Mac platforms. There are no plans to support ARM or other instruction sets.
[Add-on packages for particular applications](https://github.com/vectorclass/add-on)
[Getting-started video.](https://www.youtube.com/watch?v=TKjYdLIMTrI) Video blogger Christopher Rose has made this nice video telling how to get started with the Vector Class Library.
**Help:** You may ask for programming help on [StackOverflow](https://stackoverflow.com) using the tag vector-class-library.
Some files were not shown because too many files have changed in this diff
Show More
Reference in New Issue
Block a user
Blocking a user prevents them from interacting with repositories, such as opening or commenting on pull requests or issues. Learn more about blocking a user.
Ken Museth
