Skip to content

NanoVDB Python: GPU/device bindings (nanovdb.cuda) - #2225

Open
swahtz wants to merge 54 commits into
AcademySoftwareFoundation:masterfrom
swahtz:feature/nanovdb_python_gpu
Open

NanoVDB Python: GPU/device bindings (nanovdb.cuda)#2225
swahtz wants to merge 54 commits into
AcademySoftwareFoundation:masterfrom
swahtz:feature/nanovdb_python_gpu

Conversation

@swahtz

@swahtz swahtz commented Jun 2, 2026

Copy link
Copy Markdown
Contributor

Adds the GPU/device half of the NanoVDB Python bindings, on top of the CPU restructure that landed in #2219: a real nanovdb.cuda namespace, zero-copy interop, streams, device-side managers, and the tools.cuda algorithm surface.

Example (CuPy)

Build a grid on the device, view its buffer zero-copy, then run your own JIT-compiled CUDA kernel over it — compiled against the headers the wheel ships:

import cupy as cp
import nanovdb

handle = nanovdb.tools.cuda.createLevelSetSphere(nanovdb.GridType.Float, 20)
handle.deviceUpload(0, True)            # raw int stream handle; 0 == default stream

buf = cp.asarray(handle)                # zero-copy via __cuda_array_interface__
assert int(buf.data.ptr) == handle.device_ptr()

device_grid = handle.deviceGrid(0)      # raw NanoGrid<float>* in GPU memory
kernel = cp.RawKernel(r'''
#include <nanovdb/NanoVDB.h>
extern "C" __global__
void inspect(const nanovdb::NanoGrid<float>* grid, float* out) {
    auto acc = grid->getAccessor();
    out[0] = acc.getValue(nanovdb::Coord(0, 0, 0));
    out[1] = float(grid->activeVoxelCount());
}
''', "inspect", options=nanovdb.cuda.compile_options("-std=c++17"), backend="nvrtc")

out = cp.zeros(2, dtype=cp.float32)
kernel((1,), (1,), (device_grid.data_ptr(), out.data.ptr))
cp.cuda.runtime.deviceSynchronize()
print(out)   # -> [-3.0, 30254.0]

nanovdb.io.deviceReadGrid("grid.nvdb") reads straight onto the device the same way. See examples/ and its README for the full GPU walkthrough.

What's new

nanovdb.cudaDeviceBuffer/DeviceGridHandle move here from the root module (breaking) and gain __cuda_array_interface__ v3, __dlpack__, device_ptr/host_ptr, from_external/from_buffer, stream-aware upload/download, and recordUse(stream) to order the buffer's device free after kernels you enqueue on external non-blocking streams. Plus UnifiedBuffer/UnifiedGridHandle, DeviceMesh, DeviceStreamMap, DeviceResource, TempDevicePool, and device createNodeManager/buildVoxelBlockManager.

nanovdb.tools.cuda — all take an optional stream:

  • Rasterizers: generalized pointsToGrid, voxelsTo{OnIndex,Index,RGBA8}Grid, pointsToRGBA8Grid, and multi-GPU DistributedPointsToGrid.
  • Topology: dilateGrid / coarsenGrid / refineGrid / pruneGrid / mergeGrids.
  • Mesh → SDF: meshToGrid (device triangle mesh → narrow-band UDF + sidecar), chaining into indexToGrid to bake a float distance grid.
  • Index/QC: indexToGrid, addBlindData, updateGridStats, evalChecksum/validateChecksum/updateChecksum, isValid, signedFloodFill, sampleFromVoxels.
  • Sidecar injection (inject, injectPredicateToMask, injectGridMask) and dense VBM decode (gatherBoxStencil(valueCount, 27), activeVoxelCoords) — so dense array/tile frameworks (CuPy, cuTile) can run VDB stencils and bake results back without a hand-written decode kernel.

Custom-kernel interopgrid.data_ptr() + nanovdb.cuda.compile_options() (NVRTC include flags for the bundled headers), as in the example above.

Naming — bound names use camelCase consistently (includeStats, dGrid, log2BlockWidth, …); the zero-copy interop surface intentionally keeps ecosystem spellings (data_ptr, from_external, …) matching CuPy/PyTorch conventions.

Tests / examplesTestGpuInterop.py (ctest-wired; self-skips without CUDA/GPU/CuPy) and ~20 runnable GPU examples: device-tools ports of the C++ ex_* examples, cupy.RawKernel ray tracers, mesh→SDF, and a full GPU level-set filter (levelset_filter.py: Laplacian deform + Godunov reinit + narrow-band retrack via the VBM) with three interchangeable compute backends — rawkernel (fused CUDA), cupy (array ops), cutile (NVIDIA cuTile) — producing identical results.

Build/CI — nanobind 2.12.0 scoped to the NanoVDB CI job + wheel (>=2.12.0,<3); the binding code stays ≥2.5.0 source-compatible and the shared build.yml co-compiles it against 2.5.0.

Also included

tools/cuda/IndexToGrid.cuh: the output value grid now gets GridClass::Unknown instead of inheriting the source's Index class, so indexToGrid output passes isValid/checkGrid.

Validation

On a Blackwell (sm_120) GPU with CuPy: TestNanoVDB (166), TestGpuInterop (54), and TestExamples (14) all pass, and every GPU example runs and self-skips cleanly without CuPy. CI does not exercise GPU pathsctest excludes cuda/mgpu and the GPU tests self-skip via isCudaAvailable()/isGpuAvailable().

swahtz added 22 commits May 20, 2026 21:16
…eFoundation#2209)

* nanovdb python: Phase 0 foundation for C++ API mirror

First slice of the Python bindings restructure tracked in AcademySoftwareFoundation#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 AcademySoftwareFoundation#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 AcademySoftwareFoundation#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>
… PointAccessor (AcademySoftwareFoundation#2210)

* nanovdb python: Phase 1 — polymorphic Grid, GridMetaData, blind data, PointAccessor

Second slice of the Python bindings restructure tracked in AcademySoftwareFoundation#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 AcademySoftwareFoundation#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 AcademySoftwareFoundation#2210

Three concrete fixes from Copilot's review of the Phase 1 PR
(AcademySoftwareFoundation#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>
…cademySoftwareFoundation#2211)

* nanovdb python: Phase 2 — broaden BuildT coverage to 24 grid types

Third slice of the Python bindings restructure tracked in AcademySoftwareFoundation#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 AcademySoftwareFoundation#2211

Three concrete fixes from Copilot's review of the Phase 2 PR
(AcademySoftwareFoundation#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>
…cademySoftwareFoundation#2212)

* nanovdb python: Phase 3 — tree / nodes / NodeManager / leaf_values

Fourth slice of the Python bindings restructure tracked in AcademySoftwareFoundation#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 AcademySoftwareFoundation#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 AcademySoftwareFoundation#2212

Addresses Copilot review notes on PR AcademySoftwareFoundation#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 AcademySoftwareFoundation#2212

Three further Copilot notes on the bounds-check/lifetime commit
(f673571), 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>
…SoftwareFoundation#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 AcademySoftwareFoundation#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 AcademySoftwareFoundation#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 AcademySoftwareFoundation#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 AcademySoftwareFoundation#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 AcademySoftwareFoundation#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 AcademySoftwareFoundation#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>
…cademySoftwareFoundation#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 AcademySoftwareFoundation#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>
…oftwareFoundation#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 AcademySoftwareFoundation#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 AcademySoftwareFoundation#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>
…AcademySoftwareFoundation#2216)

* nanovdb python: Phase 5 — primitives + quantized/index createNanoGrid

Closes out Phase 5 of the NanoVDB Python bindings restructure (AcademySoftwareFoundation#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 AcademySoftwareFoundation#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 AcademySoftwareFoundation#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>
…oundation#2218)

* nanovdb python: Phase 6 — docs, migration guide, examples

Closes Phase 6 of the NanoVDB Python bindings restructure (AcademySoftwareFoundation#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 AcademySoftwareFoundation#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>
…string

Addresses review feedback on PR AcademySoftwareFoundation#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>
Addresses review feedback on PR AcademySoftwareFoundation#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>
Addresses review feedback on PR AcademySoftwareFoundation#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>
…andle seam

First step of the GPU/device bindings epoch (issue AcademySoftwareFoundation#2208 follow-on).

- Create a real root `nanovdb.cuda` submodule (mirrors C++ nanovdb::cuda)
  and MOVE DeviceBuffer + DeviceGridHandle into it. Breaking, pre-1.0
  relocation: `nanovdb.DeviceBuffer` -> `nanovdb.cuda.DeviceBuffer`,
  `nanovdb.DeviceGridHandle` -> `nanovdb.cuda.DeviceGridHandle`.
  `nanovdb.tools.cuda` is untouched. No in-repo call site named the moved
  classes (handles come from io.deviceReadGrid / tools.cuda.create*), so
  only a migration note in python/examples/README.md was needed.
- Add a BufferT-templated seam in PyDeviceBuffer.h: defineDeviceBufferLike
  <BufferT>() + addDeviceInterop<BufferT>() keyed only on the duck-typed
  buffer surface (size()/data()/deviceData()), so a future resource-backed
  buffer or UnifiedBuffer registers via one call. Phase A binds only the
  additive size() accessor; CAI/DLPack/pointers/streams land in Phase B.
- Expose nanovdb.cuda.compile_options(*extra) via __init__.py, hasattr-
  guarded so non-CUDA builds still import.

Compatible with nanobind >= 2.5.0 (no newer-only APIs). Verified by per-TU
g++/nvcc compile checks of the changed TUs; runtime import is exercised by
the shared build.yml pytest target.

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Fills the Phase-A addDeviceInterop<BufferT> seam and the device GridHandle
with the GPU interop surface:

- __cuda_array_interface__ v3 (read-only dict, whole device buffer as 1-D
  uint8) on DeviceBuffer and DeviceGridHandle.
- __dlpack__ / __dlpack_device__ on both: __dlpack__ builds an
  nb::ndarray<device::cuda,uint8> view parented to the owner (keep_alive)
  and returns nanobind's "dltensor" capsule directly — nanobind owns the
  capsule/deleter/alignment; no hand-built DLManagedTensor.
- Raw pointers: DeviceBuffer.device_ptr()/host_ptr(), DeviceGridHandle
  .device_ptr(), and a generic grid.data_ptr() (host or device per
  provenance — grid(n) vs deviceGrid(n)).
- stream:int=0 threaded through deviceUpload/deviceDownload (was hardcoded
  nullptr), via the current-device (void*,bool) overload so the targeted
  device matches deviceData().
- DeviceBuffer.from_external(size, gpu_ptr, cpu_ptr) (non-owning wrap of
  external host+device memory; mManaged==0 so neither ptr is freed) and
  DeviceGridHandle.from_buffer(buffer) (moves the buffer; ctor validates
  the grid header).

Device-only array protocols live ONLY on DeviceBuffer/DeviceGridHandle,
never the shared NanoGrid<T> class (which also wraps host grids).

Two correctness fixes from adversarial review folded in: __dlpack__ now
returns the cast capsule directly (calling .attr("__dlpack__") on it would
AttributeError, since nb::cast of a no-framework device ndarray already IS
the capsule); deviceUpload/deviceDownload use the current-device overload
instead of hardcoding device 0.

Known limitation: from_external requires BOTH a host and device pointer
because nanovdb::cuda::DeviceBuffer's external ctor asserts both non-null;
wrapping a device-only CuPy pointer would need a device-only external ctor
in the core (follow-up / resource-backed buffer). Compatible with nanobind
>= 2.5.0. Verified by per-TU g++/nvcc compile checks.

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
…BlockManager

- nanovdb.cuda.createDeviceNodeManager(device_grid, stream=0) + a
  DeviceNodeManagerHandle (over NodeManagerHandle<cuda::DeviceBuffer>):
  polymorphic over BuildTypes.def; takes a device grid (from
  DeviceGridHandle.deviceGrid(n)); mgr() returns the typed device
  NodeManager (kernel-only `this`). New cuda/PyDeviceNodeManager.cu.
- nanovdb.tools.cuda.buildVoxelBlockManager(device_onindex_grid,
  log2_block_width=6, ..., stream=0) + a DeviceVoxelBlockManagerHandle
  whose firstLeafID() (uint32) / jumpMap() (uint64) are zero-copy DEVICE
  arrays exposing the Phase-B CAI/DLPack interop. New
  cuda/PyDeviceVoxelBlockManager.cu. The __device__ decodeInverseMaps is
  intentionally not bound (kernel-only; use the shipped headers).

Four adversarial-review/runtime fixes folded in (all validated on the GPU):
1. Dropped nb::keep_alive<0,1> on the device firstLeafID()/jumpMap() views
   — the returned no-framework device ndarray is a DLPack capsule (not
   weak-referenceable), so keep_alive threw "could not create a weak
   reference"; the ndarray's py_self owner already anchors lifetime.
2. DeviceNodeManagerHandle mgr()/__bool__ now gate on size(), not the host
   data() pointer (which is null for a device-only NodeManager).
3. Dropped DeviceNodeManagerHandle.deviceUpload/deviceDownload — they
   null-deref on a device-built NodeManager (no host mirror); the handle
   is created device-resident, nothing to upload.
4. Device VoxelBlockManager registered on the existing nanovdb.tools.cuda
   (mirrors nanovdb::tools::cuda), not a bogus new nanovdb.cuda.tools.

Compatible with nanobind >= 2.5.0. Verified by per-TU compile + a real
incremental build + GPU runtime smoke (createDeviceNodeManager, device VBM
build, cupy CAI/DLPack round-trips).

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
…params

- Generalize the CUDA points->grid surface beyond Rgba8 (nanovdb.tools.cuda):
  - Coordinate input (int32 (N,3) voxel coords) -> grid via voxelsToGrid<T>
    for Rgba8 / ValueOnIndex / ValueIndex (voxelsToRGBA8Grid /
    voxelsToOnIndexGrid / voxelsToIndexGrid). Legacy pointsToRGBA8Grid kept
    (routes to voxelsToGrid<Rgba8>), so existing callers/tests are unchanged.
  - World-position input ((N,3) float OR double) -> NanoGrid<Point> via a new
    pointsToGrid (both precisions under one name; nanobind dtype dispatch).
  - Point is intentionally NOT offered on the coord path: PointsToGrid<Point>
    static_asserts Vec3f/Vec3d coords, so Point grids build from world
    positions only (documented).
- Add stream:int=0 (reinterpret_cast<cudaStream_t>) to every CUDA-launching
  binding — pointsToGrid/voxelsToGrid, signedFloodFill, both sampleFromVoxels
  overloads — with nb::gil_scoped_release around each kernel launch. Stream is
  a trailing default kwarg, so existing call sites are unaffected.
- World path constructs PointsToGrid<Point>(voxelSize, ...) directly rather
  than the fixed-voxelSize free function, which is declared-but-undefined in
  the header (would link-fail).

Returns DeviceGridHandle (Phase-A/B device interop). Compatible with
nanobind >= 2.5.0. Verified by per-TU compile + real incremental build + GPU
runtime smoke (world float32/float64 -> Point grids, coord -> OnIndex/Index/
RGBA8, non-default streams) and the full Python suite (140/140 pass).

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Bind 5 of the 6 tools::cuda morphology/topology grid operators into
nanovdb.tools.cuda, one Py*.cu + .h per header:

- dilateGrid(d_grid, op=26, stream=0)  (NN_FACE=6 / NN_FACE_EDGE_VERTEX=26)
- coarsenGrid(d_grid, stream=0)        (2x topological downsample)
- refineGrid(d_grid, stream=0)         (2x topological upsample)
- pruneGrid(d_grid, leafMask, stream=0) (leafMask: device uint64 array
  reinterpreted as Mask<3>* per leaf; word count validated as a multiple of 8)
- mergeGrids(d_grid1, d_grid2, stream=0) (binary active-mask union; chain for >2)

All take a device grid (NanoGrid<ValueOnIndex>* from deviceGrid(n)) + a
stream:int=0, run getHandle() under nb::gil_scoped_release, and return a
DeviceGridHandle (Phase-A/B CAI/DLPack). Instantiated for ValueOnIndex ONLY:
every op embeds TopologyBuilder<BuildT>, which static_asserts is_onindex.

TopologyBuilder itself is intentionally NOT bound — it is the low-level
multi-stage engine (ordered allocate/count/process* calls, no one-shot
getHandle), not a user-facing op.

Each .cu drops `using namespace nanovdb;` and fully-qualifies nanovdb:: to
avoid a CUB DeviceScan vs nanovdb::cuda 'reference to cuda is ambiguous'
error (same fix PyPointsToGrid.cu already uses).

Compatible with nanobind >= 2.5.0. Verified by per-TU compile + real
incremental build + GPU runtime smoke: all 5 ops run on a device OnIndex
grid with correct active-voxel deltas (dilate grows, coarsen ~8x shrink,
refine ~8x grow; prune + binary merge succeed).

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Both evalChecksum overloads (GridData* and NanoGrid<BuildT>*) copied the
computed head/tail CRC from device to host with size `headSize`
(sizeof(GridData)+sizeof(TreeData), ~672 bytes) into the 4-byte
Checksum::head()/tail() uint32_t destinations, reading ~672 bytes from
`d_lut.get()+256` — one uint32_t past the 1024-byte CRC LUT. On a real
device this trips "CUDA error 1: invalid argument" and cudaCheck aborts.
The sibling validateChecksum copies the CRC correctly with
sizeof(uint32_t); match it. The full-GridData header copies (which
legitimately use headSize) are unchanged.

Surfaced by the new nanovdb.tools.cuda.evalChecksum Python binding;
verified fixed on a Blackwell GPU (evalChecksum returns a Checksum for
CheckMode.Partial and Full instead of aborting).

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Bind into nanovdb.tools.cuda (one Py*.cu + .h per header; stream:int=0 +
nb::gil_scoped_release; fully-qualified nanovdb:: per the CUB ambiguity fix):

- indexToGrid(index_grid, values, stream=0) -> DeviceGridHandle: device
  ValueIndex/ValueOnIndex grid + a device values ndarray -> typed grid.
  Scalar (float/double, flat 1-D) and Vec3 (Vec3f/Vec3d, (N,3) c_contig)
  destinations. Dst restricted to non-special types (processLeafsKernel
  static_assert !is_special); Src restricted to index types (SFINAE).
- addBlindData(grid, blind_data, class, semantic, name, stream=0) ->
  DeviceGridHandle: appends a device blind-data channel (valueCount from
  the array length). Reuses existing GridBlindDataClass/Semantic enums.
- Device QC mirroring the host names on tools.cuda:
  - updateGridStats(grid, mode=StatsMode.Default, stream=0) -> None
  - isValid(grid, mode=CheckMode.Default, verbose=False, stream=0) -> bool
  - evalChecksum/validateChecksum/updateChecksum(grid, mode, stream=0)
    (device grid reinterpreted as GridData*; the GridData* overloads only
    copy the header host-side).
  updateGridStats bound for scalar/vector/bool BuildTs (special/index/mask
  trip Extrema/Stats static_asserts); isValid + checksum bound across the
  full 20-type set.

Include ordering: the QC .cu files include <nanovdb/cuda/GridHandle.cuh>
before the tool .cuh (the tool headers aren't self-contained for the
device GridHandle's updateChecksum/NodeManager uses), matching
unittest/TestNanoVDB.cu.

Compatible with nanobind >= 2.5.0. Verified by per-TU compile + real
incremental build + GPU runtime smoke (updateGridStats/isValid/evalChecksum/
validateChecksum/updateChecksum, indexToGrid float+OnIndex, addBlindData).

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
…PointsToGrid

On nanovdb.cuda (NanoVDBModule.cc):
- UnifiedBuffer: registered via the Phase-A defineDeviceBufferLike seam, so
  it gets CAI/DLPack/device_ptr/host_ptr/size for free (managed memory:
  host_ptr == device_ptr). Plus ctors, capacity/resize/clear, advise,
  prefetch, deviceUpload/deviceDownload.
- UnifiedGridHandle (GridHandle<UnifiedBuffer>): needed because
  DistributedPointsToGrid.getHandle returns one; full GridHandle surface +
  deviceGrid/deviceUpload/deviceDownload/device_ptr.
- DeviceMesh (+ DeviceNode), DeviceStreamMap (+ DeviceType enum): host-query
  surface for multi-device/stream coordination.
- DeviceResource (static allocateAsync/deallocateAsync + DEFAULT_ALIGNMENT)
  and TempDevicePool (TempPool<DeviceResource>) bound in their CURRENT shape
  only — not reshaped into the future resource concept.

On nanovdb.tools.cuda (PyTools.cc):
- DistributedPointsToGrid / DistributedIndexPointsToGrid /
  DistributedRGBA8PointsToGrid (ValueOnIndex / ValueIndex / Rgba8): ctor
  (DeviceMesh + scale/translation, mesh kept alive via keep_alive) +
  getHandle((N,3) int32 coords) -> UnifiedGridHandle.

getHandle requires nb::device::cuda_managed input: the C++ pipeline issues
cudaMemAdvise/memPrefetchAsync directly on the pointer, so the array MUST be
CUDA managed/unified memory. (Adversarial review caught that the initial
nb::device::cuda tag made it uncallable — managed input rejected, plain
device input crashed; fixed to cuda_managed.)

Forward-compat guardrail honored: none of the reserved future names
(Resource/MemoryResource/AsyncResource/PinnedResource/ResourceDeviceBuffer/
DeviceBuffer2/default_resource/set_default_resource) are bound. NCCL-gated
and __device__-only internals left unbound.

Compatible with nanobind >= 2.5.0. Verified by concurrent per-TU compile +
real incremental build + GPU runtime smoke: UnifiedBuffer zero-copy
CAI/DLPack round-trip, DeviceMesh/DeviceStreamMap construct, and all three
Distributed*PointsToGrid build OnIndex/Index/RGBA8 grids from managed coords
on a single GPU (plain-device input cleanly rejected).

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
…examples, and docs

Adds the test/example/documentation surface for the NanoVDB Python GPU
bindings landed in the preceding phases. No C++ binding changes.

- test/TestGpuInterop.py: stdlib-unittest suite (37 tests) mirroring
  TestNanoVDB.py gating — every case is class-guarded on isCudaAvailable()
  and isGpuAvailable(), CuPy is guarded per-test, and Torch via try-import +
  skipTest, so the suite self-skips cleanly on a non-CUDA/no-GPU build.
  Covers CAI v3 dict + DLPack round-trips aliasing device_ptr (nbytes==size),
  from_external (managed wrap + null rejection), from_buffer move/reject,
  non-default streams, the host-accessor-on-device-grid SIGSEGV contract
  (asserted in a subprocess so it cannot abort the runner), UnifiedBuffer,
  the generalized point/voxel rasterizers, device morphology/QC/flood-fill/
  sampling, device NodeManager and VoxelBlockManager (firstLeafID/jumpMap
  device views), the cuda infrastructure types, and DistributedPointsToGrid
  over managed memory.

- examples/{gpu_load_inspect,cupy_rawkernel,numba_cuda,triton_kernel}.py:
  runnable, import-guarded examples. cupy_rawkernel.py demonstrates the
  device-pointer ABI — a cp.RawKernel that #includes NanoVDB.h via
  compile_options() and reads a const NanoGrid<float>* from grid.data_ptr().
  numba/triton examples self-skip with an install hint when absent.

- examples/README.md: a GPU/CUDA section covering the nanovdb.cuda (infra)
  vs nanovdb.tools.cuda (algorithms) split, attribute-not-submodule access,
  the device-grid build recipe, zero-copy interop + device-pointer ABI, the
  host-accessor caveat, streams, from_external/from_buffer, the managed-memory
  distributed pipeline, compile_options for NVRTC, the relocation migration
  note, and a forward-looking pluggable-memory-resources roadmap.

- python/CMakeLists.txt: registers TestGpuInterop.py as its own ctest test
  (pytest_nanovdb_gpu_interop) mirroring pytest_nanovdb; self-skips without
  a GPU so it is harmless in non-CUDA CI.

Validated locally on a Blackwell GPU with CuPy: full suite green (1 torch
skip), both required examples run end-to-end, TestNanoVDB.py shows no
regression, and a CMake reconfigure confirms both ctest entries register.

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
…e wheel

The NanoVDB Python GPU bindings (CUDA array interface, DLPack export, device
interop) benefit from a newer nanobind than the repo-wide 2.5.0 floor. Scope
the bump to the two places that build the published NanoVDB surface:

- .github/workflows/nanovdb.yml: install nanobind 2.12.0 (latest 2.x) for the
  dedicated NanoVDB job.
- pyproject.toml: pin the wheel's nanobind to >=2.12.0,<3 (it was previously
  unpinned and silently floated to latest, diverging from CI).

The other seven install sites (build/houdini/docs/ax/weekly) and the global
FUTURE_MINIMUM_NANOBIND_VERSION stay at 2.5.0: nanobind is a shared dependency
and the co-compiled pyopenvdb module plus the shared build.yml still target
2.5.0, so the binding code remains >=2.5.0 source-compatible. The <3 ceiling
keeps the wheel on the 2.x series.

Verified locally on a Blackwell GPU: nanobind 2.12.0 (installed to a separate
prefix, CMake-resolved version confirmed 2.12.0) compiles both nanovdb_python
and openvdb_python with no errors and no deprecation warnings.

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Expose NanoVDB's sidecar-injection operators (util/cuda/Injection.cuh) on the
nanovdb.tools.cuda submodule. These carry a grid's per-voxel data across a
topology change (dilate/prune/merge), which the morphology ops alone do not do
— the OnIndex value-index space changes, so values must be re-injected.

- inject(src_grid, dst_grid, src_sidecar, dst_sidecar, stream=0): copies sidecar
  values for every voxel present in BOTH grids (the intersection); destination
  voxels with no source counterpart are left unchanged. Wraps
  InjectGridDataFunctor via util::cuda::operatorKernel (one block per source
  leaf, bit-parallel per warp). Bound for float and double sidecars.
- injectPredicateToMask(grid, predicate, leaf_masks, stream=0): turns a boolean
  predicate over a grid's value indices into the per-leaf Mask<3> retain mask
  that pruneGrid consumes, so a value-driven trim (e.g. |phi| <= halfWidth) is
  two device ops with no hand-rolled kernel. Wraps InjectPredicateToMaskFunctor.
  leaf_masks need only be >= (leaf count)*8 uint64; activeVoxelCount*8 is a safe
  size since every leaf holds at least one active voxel.

This is the GPU realization of the injectData abstraction in the NanoVDB 2.0
paper (section 3.4), and completes the dilate -> inject -> prune narrow-band
rebuild loop entirely with bound device ops.

Validated locally on a Blackwell GPU: inject copies exactly the src/dst
intersection (leaving the rest, including the background slot, untouched);
injectPredicateToMask + pruneGrid trims to exactly the predicate-true voxel set;
a second inject carries values onto the pruned grid; bad grid arguments raise.

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
swahtz added 12 commits June 9, 2026 04:32
The device VoxelBlockManager gather and the sidecar-transfer / index
conversion bindings were instantiated only for float/double, even though
their kernels are type-generic copies. Expose the integer types so an
integer label / id / occupancy sidecar can be carried through the GPU
pipeline without a float round-trip:

  * gatherBoxStencil: add int32/uint32 overloads. Lets callers gather a
    neighbour-VALUE-INDEX table directly in int32 instead of the old
    float64-gather + int64-cast, halving the (valueCount, 27) table.
  * inject / injectFeatures: add int32/uint32 (InjectGrid*Functor copies
    via the assignment operator, so any trivially-copyable T works).
  * addBlindData: add int32/int64 blind payloads for every grid BuildT.
  * indexToGrid: add an Int32 scalar destination (Int32 is non-special,
    so it satisfies processLeafsKernel's static_assert) -> an Int32Grid.

Also guard gatherBoxStencil / activeVoxelCoords. They write
out[valueIndex] for each active voxel with the caller sizing out to
activeVoxelCount + 1, which is valid only when active-voxel indexing is
contiguous (the VBM invariant). On a grid carrying per-node stats / tile
values (the createOnIndexGrid defaults) value indices run past
activeVoxelCount and the writes go out of bounds -- previously an illegal
memory access. Read valueCount / activeVoxelCount from the device grid
header (DeviceGridTraits) and raise a clear error pointing at the fix.

Add TestGpuInterop coverage -- the first tests for these device ops:
TestDeviceInject (inject / injectFeatures over float32/int32/uint32),
TestDeviceTypedSidecars (addBlindData int32/int64, indexToGrid -> Int32),
and gatherBoxStencil dtype + non-contiguous-grid-rejection tests.

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
gatherBoxStencil always materialises the full (valueCount, 27) neighbour
table, but a caller often needs only a handful of the 27 spokes (e.g. an
SDF mesher's 8 corners + 6 faces -- 14 of 27). Add a sibling that writes
only a chosen subset into an (valueCount, K) array.

The kernel is the gatherBoxStencil kernel with the write narrowed: the
VBM decode still computes all 27 neighbour
shared), then it writes out[k, col] = values[st[spokes[col]]] for the K
requested spokes. The K (<=27) spoke indice
struct, so there is no device scratch. spokes is a 1-D host int32 array,
validated in [0, 27) (and against out.shape
the same contiguous-active-indexing guard as gatherBoxStencil.
Instantiated for float / double / int32 / u

Cuts the dominant array roughly in proporti
contouring mesher this is the (N,14) vs (N,27) table difference, the
final step taking its full-resolution peak

Add TestDeviceVoxelBlockManager.test_gather
subset equals the matching columns of the full gather, and an out-of-
range spoke / mismatched out.shape[1] raise

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
The grid bindings exposed getters for gridClass and the transform but no
way to set them after a grid was built. indexToGrid, for one, copies the
source's GridClass onto its output, so an SDF FloatGrid materialised from
an IndexGrid came out tagged IndexGrid with no way to retag it LevelSet.

Add mutable header setters on both the host and device grid:

- tools.cuda.setGridClass(d_grid, gridClass, stream): new
  defineDeviceGridMetadata<BuildT> in PyDeviceGridChecksum.cu (a one-thread
  setGridClassKernel writes mGridClass on the device grid, then
  tools::cuda::updateChecksum refreshes the checksum preserving its mode),
  registered for the full checksum BuildT set in PyTools.cc.

- Grid.setGridClass(gridClass) and Grid.setTransform(voxelSize,
  translation=Vec3d(0,0,0)) on the host GridData type. setTransform builds
  a uniform-scale + translation Map, updates mVoxelSize, and recomputes
  mWorldBBox from the index bbox under the new map (mirroring GridStats).

Also fix a latent wart: the existing host header setters (setGridName and
the setMinMaxOn/setBBoxOn/setLongGridNameOn/setAverageOn/setStdDeviationOn
flag toggles) edited the header without refreshing the checksum, silently
desyncing a checksummed grid so a validating reader would reject it. They
now call updateChecksum after the edit, which preserves the checksum mode
and is a no-op when the checksum is disabled (so it never enables one).

Tests: TestNanoVDB.TestGridHeaderSetters (host setGridClass/setTransform,
the flag/name setters keeping a Full checksum valid, and all surviving a
file round-trip) and TestGpuInterop.TestDeviceGridMetadata (device
setGridClass proven to refresh a pre-populated Full checksum, and
surviving a round-trip as FogVolume). Full suites: TestNanoVDB 144 OK,
TestGpuInterop 46 OK (1 skip).

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
processGridTreeRootKernel builds the destination grid header by copying
the source index grid's header verbatim (*dstGrid.data() = *srcGrid.data())
and then only patches mGridType and mData1. It never updated mGridSize or
the blind-metadata fields, so the destination inherited the SOURCE grid's
values for them.

A destination value grid has a different memory footprint than its source
index grid -- e.g. a NanoLeaf<float> stores a full 512-value array where the
compact OnIndex leaf stores none -- so a dense float grid is roughly 2x the
bytes of the index grid it came from. With mGridSize left at the smaller
source value, GridHandle::write emits only mGridSize bytes and truncates the
grid: about half the leaves are lost and the surviving node offsets dangle
past the end of the file. The device buffer itself was allocated correctly
(getBuffer uses nodeAcc.size), so the grid was fine in memory and only the
written file was corrupt -- which then crashed downstream readers such as
nanoToOpenVDB / nanovdb_convert. The bug only bites when the destination is
larger than the source (large/dense grids); small or sparse grids, where the
inherited size still covers the data, were unaffected.

Set mGridSize to nodeAcc->size (the actual destination layout) and reset the
blind-metadata fields, since indexToGrid copies no blind data.

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
…geGrids

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
…radients

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>
Conflicts:
- nanovdb/nanovdb/python/NanoVDBModule.cc: defineNanoGrid gained data_ptr
  on this branch and the valueCount/pointCount bindings on
  feature/nanovdb_python — kept both.
- nanovdb/nanovdb/python/test/TestNanoVDB.py: test_list_to_vector's
  multi-grid round-trip check (added on this branch) used the removed
  typed doubleGrid() accessor — ported to the polymorphic grid().

TestNanoVDB.py: 166 cases green on a CUDA build.

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
…ython

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
…ples

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>
# Conflicts:
#	nanovdb/nanovdb/python/CMakeLists.txt
Add device/GPU Python examples exercising the nanovdb.cuda /
nanovdb.tools.cuda surface, bind tools::cuda::MeshToGrid, and fix two
device-tool issues surfaced while writing the examples.

New GPU examples (nanovdb/nanovdb/python/examples/):
- voxels_to_grid_cuda, device_topology_ops, sample_from_voxels_cuda,
  index_to_grid_cuda, signed_flood_fill_cuda, validate_cuda,
  collide_level_set_cuda: device grid construction, morphology/topology
  ops, trilinear sampling + gradients, index->value baking, signed flood
  fill, device QC, and vectorized particle collision.
- raytrace_level_set_cuda, raytrace_fog_volume_cuda: cupy.RawKernel
  renderers using the in-kernel HDDA / accessor C++ API.
- mesh_to_grid_cuda: triangle mesh -> narrow-band UDF via meshToGrid.

Bindings:
- Bind nanovdb::tools::cuda::MeshToGrid as tools.cuda.meshToGrid,
  returning a device OnIndex handle plus a per-value UDF sidecar buffer.
- Bind device isValid for Point grids.

Fixes:
- IndexToGrid.cuh: the destination value grid inherited the source index
  grid's GridClass (Index), forming an invalid GridType/GridClass
  combination that failed grid validation; set GridClass::Unknown on the
  output value grid.

Document the new GPU examples and meshToGrid in the examples README.

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
swahtz added 7 commits July 30, 2026 14:10
  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>
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>
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>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
@swahtz
swahtz changed the base branch from feature/nanovdb_python to master August 18, 2026 21:59
swahtz and others added 3 commits August 19, 2026 10:36
…ython_gpu

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
The Python bindings had drifted into two naming conventions. The method
and function surface was already 208 camelCase names to 6 snake_case, but
parameter names were split almost evenly (27 camelCase, 24 snake_case),
with the snake_case ones confined to the CUDA device bindings and
PyVoxelBlockManager.cc while the host bindings used camelCase throughout.
That put the same concept under two spellings in one namespace, e.g.
createNanoGridOnIndex(includeStats=) next to
createOnIndexGrid(include_stats=), and left three files internally mixed.

Rename the snake_case parameters (d_grid, log2_block_width, include_stats,
first_offset, src_grid, leaf_masks, ...) and the five snake_case
method/property outliers (block_width, log2_block_width, jump_map_length,
leaf_values, is_divisible) to camelCase, and update every call site in the
examples and tests along with the prose that referenced them.

The interop surface keeps its snake_case names on purpose: data_ptr,
device_ptr, host_ptr, from_external, from_buffer, jump_map_ptr and
first_leaf_id_ptr sit in the zero-copy CuPy / PyTorch / Numba path, where
data_ptr() is the spelling those libraries use. Default grid-name values
such as "sphere_ls" are data rather than identifiers and are untouched.

The bindings have never appeared in a release (they landed on master after
v13.0.0), so no deprecation aliases are needed.

Verified with TestNanoVDB.py (166), TestGpuInterop.py (50, 31 skipped for
absent torch/cupy) and TestExamples.py (14) — all passing.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
…rnal streams

The bindings' interop surface hands raw device pointers to CuPy / PyTorch /
Numba via device_ptr(), __cuda_array_interface__ and __dlpack__, and those
frameworks launch kernels on their own non-blocking streams. Such work is
invisible to the buffer's automatic upload/download use tracking, so the
cudaFreeAsync issued when the buffer is cleared or destroyed could race it —
the C++ side grew DeviceBuffer::recordUse(device, stream) for exactly this
caller, but Python had no way to reach it.

Bind recordUse(stream, device=-1) on DeviceBuffer and on DeviceGridHandle
(which owns its buffer internally and is where every real workflow's owning
buffer lives). Both forward to a shared recordUseChecked helper that
validates the device id — the C++ method indexes per-device tracking state
unchecked — raising IndexError on an out-of-range id, with -1 selecting the
current CUDA device. The device_ptr and __cuda_array_interface__ docstrings
now point callers at recordUse. Non-owning (from_external) buffers accept
the call as a no-op, matching the C++ behavior.

Add TestRecordUse to TestGpuInterop.py: a CuPy reduction enqueued on a
non-blocking stream against a zero-copy CAI view, recorded, and validated
after the handle is destroyed while the work may still be in flight; plus
default-stream/explicit-device acceptance, IndexError on a bad device id,
and the non-owning no-op.

Also normalize the cuda/ sources' includes of shared binding headers: add
the python source dir to the target's private include paths so they use
plain names ("PyGridHandle.h", "BuildTypes.def") instead of "../"-relative
paths, which the rest of the codebase does not use.

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

Adds comprehensive CUDA/device support to NanoVDB’s Python API, including zero-copy interoperability, GPU algorithms, infrastructure, tests, and examples.

Changes:

  • Introduces nanovdb.cuda device buffers, handles, streams, and multi-GPU infrastructure.
  • Exposes CUDA topology, rasterization, validation, checksum, and sidecar algorithms.
  • Adds GPU interoperability examples and test registration.

Reviewed changes

Copilot reviewed 82 out of 82 changed files in this pull request and generated 11 comments.

Show a summary per file
File Description
pyproject.toml Updates nanobind requirement.
.github/workflows/nanovdb.yml Updates NanoVDB CI nanobind version.
nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh Corrects output-grid metadata.
nanovdb/nanovdb/python/CMakeLists.txt Builds and tests CUDA bindings.
nanovdb/nanovdb/python/__init__.py Adds CUDA compile options.
nanovdb/nanovdb/python/NanoVDBModule.cc Registers CUDA infrastructure and metadata APIs.
nanovdb/nanovdb/python/PyMath.cc Renames the divisibility binding.
nanovdb/nanovdb/python/PyTools.cc Registers CUDA tool algorithms.
nanovdb/nanovdb/python/PyTree.h Adds device manager declarations and renames leaf access.
nanovdb/nanovdb/python/PyVoxelBlockManager.h Declares device VBM bindings.
nanovdb/nanovdb/python/cuda/PyAddBlindData.h Declares device blind-data binding.
nanovdb/nanovdb/python/cuda/PyAddBlindData.cu Implements device blind-data binding.
nanovdb/nanovdb/python/cuda/PyCoarsenGrid.h Declares coarsening binding.
nanovdb/nanovdb/python/cuda/PyCoarsenGrid.cu Implements coarsening binding.
nanovdb/nanovdb/python/cuda/PyDeviceBuffer.h Adds shared zero-copy interoperability.
nanovdb/nanovdb/python/cuda/PyDeviceBuffer.cc Adds external buffer wrapping.
nanovdb/nanovdb/python/cuda/PyDeviceGridChecksum.h Declares device checksum APIs.
nanovdb/nanovdb/python/cuda/PyDeviceGridChecksum.cu Implements checksum and metadata APIs.
nanovdb/nanovdb/python/cuda/PyDeviceGridHandle.cu Adds streams and zero-copy exports.
nanovdb/nanovdb/python/cuda/PyDeviceGridStats.h Declares device statistics binding.
nanovdb/nanovdb/python/cuda/PyDeviceGridStats.cu Implements device statistics binding.
nanovdb/nanovdb/python/cuda/PyDeviceGridValidator.h Declares device validation binding.
nanovdb/nanovdb/python/cuda/PyDeviceGridValidator.cu Implements device validation binding.
nanovdb/nanovdb/python/cuda/PyDeviceMesh.h Declares multi-GPU mesh bindings.
nanovdb/nanovdb/python/cuda/PyDeviceMesh.cu Implements multi-GPU mesh bindings.
nanovdb/nanovdb/python/cuda/PyDeviceNodeManager.cu Implements device node managers.
nanovdb/nanovdb/python/cuda/PyDeviceStreamMap.h Declares stream-map binding.
nanovdb/nanovdb/python/cuda/PyDeviceStreamMap.cu Implements stream-map binding.
nanovdb/nanovdb/python/cuda/PyDeviceVoxelBlockManager.cu Implements device VBM and dense decoding.
nanovdb/nanovdb/python/cuda/PyDilateGrid.h Declares dilation binding.
nanovdb/nanovdb/python/cuda/PyDilateGrid.cu Implements dilation binding.
nanovdb/nanovdb/python/cuda/PyDistributedPointsToGrid.h Declares distributed rasterization.
nanovdb/nanovdb/python/cuda/PyDistributedPointsToGrid.cu Implements distributed rasterization.
nanovdb/nanovdb/python/cuda/PyIndexToGrid.h Declares index-grid conversion.
nanovdb/nanovdb/python/cuda/PyIndexToGrid.cu Implements index-grid conversion.
nanovdb/nanovdb/python/cuda/PyInjectData.h Declares sidecar injection APIs.
nanovdb/nanovdb/python/cuda/PyInjectData.cu Implements sidecar injection APIs.
nanovdb/nanovdb/python/cuda/PyMergeGrids.h Declares device grid merging.
nanovdb/nanovdb/python/cuda/PyMergeGrids.cu Implements device grid merging.
nanovdb/nanovdb/python/cuda/PyMeshToGrid.h Declares mesh rasterization.
nanovdb/nanovdb/python/cuda/PyMeshToGrid.cu Implements mesh-to-UDF rasterization.
nanovdb/nanovdb/python/cuda/PyPointsToGrid.h Declares point and voxel rasterizers.
nanovdb/nanovdb/python/cuda/PyPointsToGrid.cu Implements point and voxel rasterizers.
nanovdb/nanovdb/python/cuda/PyPruneGrid.h Declares pruning binding.
nanovdb/nanovdb/python/cuda/PyPruneGrid.cu Implements mask-driven pruning.
nanovdb/nanovdb/python/cuda/PyRefineGrid.h Declares refinement binding.
nanovdb/nanovdb/python/cuda/PyRefineGrid.cu Implements refinement binding.
nanovdb/nanovdb/python/cuda/PySampleFromVoxels.cu Adds stream-aware device sampling.
nanovdb/nanovdb/python/cuda/PySignedFloodFill.cu Adds stream-aware flood filling.
nanovdb/nanovdb/python/cuda/PyTempPool.h Declares temporary allocation APIs.
nanovdb/nanovdb/python/cuda/PyTempPool.cu Implements temporary allocation APIs.
nanovdb/nanovdb/python/cuda/PyUnifiedBuffer.h Declares managed-memory buffers.
nanovdb/nanovdb/python/cuda/PyUnifiedBuffer.cu Implements managed-memory buffers.
nanovdb/nanovdb/python/cuda/PyUnifiedGridHandle.h Declares managed grid handles.
nanovdb/nanovdb/python/cuda/PyUnifiedGridHandle.cu Implements managed grid handles.
nanovdb/nanovdb/python/test/TestGpuInterop.py Exercises GPU interoperability.
nanovdb/nanovdb/python/examples/build_grid.py Updates API naming.
nanovdb/nanovdb/python/examples/bulk_leaf_numpy.py Updates leaf-value API naming.
nanovdb/nanovdb/python/examples/collide_level_set_cuda.py Demonstrates GPU collisions.
nanovdb/nanovdb/python/examples/cupy_rawkernel.py Demonstrates custom CuPy kernels.
nanovdb/nanovdb/python/examples/device_topology_ops.py Demonstrates topology operations.
nanovdb/nanovdb/python/examples/gpu_load_inspect.py Demonstrates GPU loading and inspection.
nanovdb/nanovdb/python/examples/index_grid_channels.py Updates API naming.
nanovdb/nanovdb/python/examples/index_to_grid_cuda.py Demonstrates device value baking.
nanovdb/nanovdb/python/examples/levelset_filter_cupy.py Adds CuPy filter backend.
nanovdb/nanovdb/python/examples/levelset_filter_cutile.py Adds cuTile filter backend.
nanovdb/nanovdb/python/examples/mesh_to_grid_cuda.py Demonstrates mesh-to-UDF conversion.
nanovdb/nanovdb/python/examples/node_manager.py Updates leaf-value documentation.
nanovdb/nanovdb/python/examples/numba_cuda.py Demonstrates Numba interoperability.
nanovdb/nanovdb/python/examples/quantize.py Updates API naming.
nanovdb/nanovdb/python/examples/raytrace_fog_volume_cuda.py Adds GPU fog rendering.
nanovdb/nanovdb/python/examples/raytrace_level_set_cuda.py Adds GPU level-set rendering.
nanovdb/nanovdb/python/examples/sample_from_voxels_cuda.py Demonstrates device sampling.
nanovdb/nanovdb/python/examples/signed_flood_fill_cuda.py Demonstrates device flood filling.
nanovdb/nanovdb/python/examples/triton_kernel.py Demonstrates Triton interoperability.
nanovdb/nanovdb/python/examples/validate_cuda.py Demonstrates device validation.
nanovdb/nanovdb/python/examples/voxels_to_grid_cuda.py Demonstrates device rasterization.
Suppressed comments (6)

nanovdb/nanovdb/python/cuda/PyDeviceBuffer.h:127

  • The DLPack stream argument is ignored. A consumer passes its CUDA stream here specifically so the producer can order that stream after outstanding writes; after deviceUpload(non_default_stream, sync=False), this returns a capsule immediately and the consumer may read partially uploaded bytes. Honor the requested stream using the buffer's tracked completion event, or synchronize before exporting.
        "__dlpack__",
        [](nb::handle self, nb::handle /*stream*/) {
            const BufferT& buf = nb::cast<const BufferT&>(self);

nanovdb/nanovdb/python/cuda/PyDeviceGridHandle.cu:162

  • This second DLPack producer also discards the consumer-provided stream. An asynchronous deviceUpload(..., sync=False) on another stream can therefore race torch.from_dlpack(handle)/cupy.from_dlpack(handle). Order the supplied consumer stream after the handle buffer's tracked upload event before returning the capsule.
            "__dlpack__",
            [](nb::handle self, nb::handle /*stream*/) {
                auto& handle = nb::cast<GridHandle<BufferT>&>(self);

nanovdb/nanovdb/python/cuda/PyDeviceVoxelBlockManager.cu:552

  • activeVoxelCoordsKernel writes one row for every value index, but the output row count is never checked despite the docstring requiring rows >= valueCount. An undersized output produces out-of-bounds GPU writes. Reject it before launching using the value count already queried during contiguous-index validation.
        [](nb::handle py_grid,
           nb::ndarray<int32_t, nb::shape<-1, 3>, nb::c_contig, nb::device::cuda> out,
           int log2BlockWidth, uintptr_t stream) {
            auto* dGrid = castOnIndexDeviceGrid(py_grid, "activeVoxelCoords");
            requireContiguousIndexing(dGrid, "activeVoxelCoords");

nanovdb/nanovdb/python/cuda/PyInjectData.cu:99

  • Only the feature dimensions are compared; the row counts are not checked against the source and destination grid value counts. The kernel indexes rows by grid value index, so undersized (rows, dim) arrays can read or write beyond their allocations. Validate both row counts before launch.
    nanovdb/nanovdb/python/cuda/PyDeviceBuffer.h:106
  • This hard-codes the producer stream to the legacy default stream even when the last upload was issued asynchronously on a caller-supplied non-blocking stream. CUDA Array Interface consumers trust this field for synchronization, so cp.asarray(buffer) can race an unfinished upload. Report/order against the actual tracked producer event, or synchronize before exposing the interface.
            iface["data"] =
                nb::make_tuple(reinterpret_cast<uintptr_t>(buf.deviceData()), false);
            iface["version"] = 3;
            iface["strides"] = nb::none();
            iface["stream"] = 1;

nanovdb/nanovdb/python/cuda/PyDeviceVoxelBlockManager.cu:469

  • As in the full-stencil overload, this kernel indexes values and out by grid value index but only checks the number of output columns. Validate values.size() and out.shape(0) against the grid's value count; otherwise a short array causes out-of-bounds device access.
           nb::ndarray<const T, nb::ndim<1>, nb::c_contig, nb::device::cuda> values,
           nb::ndarray<T, nb::shape<-1, -1>, nb::c_contig, nb::device::cuda>  out,
           nb::ndarray<const int32_t, nb::ndim<1>, nb::c_contig>             spokes,
           int log2BlockWidth, uintptr_t stream) {
            auto* dGrid = castOnIndexDeviceGrid(py_grid, "gatherBoxStencilColumns");

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

Comment thread nanovdb/nanovdb/python/cuda/PyDeviceBuffer.h
Comment thread nanovdb/nanovdb/python/NanoVDBModule.cc
Comment thread nanovdb/nanovdb/python/cuda/PyIndexToGrid.cu
Comment thread nanovdb/nanovdb/python/cuda/PyInjectData.cu
Comment thread nanovdb/nanovdb/python/cuda/PyInjectData.cu
Comment thread nanovdb/nanovdb/python/cuda/PyPointsToGrid.cu
Comment thread nanovdb/nanovdb/python/cuda/PyPointsToGrid.cu
Comment thread nanovdb/nanovdb/python/cuda/PyMeshToGrid.cu
Comment thread nanovdb/nanovdb/python/cuda/PyPruneGrid.cu
Comment thread nanovdb/nanovdb/python/cuda/PyDeviceVoxelBlockManager.cu Outdated
…/DLPack exports

Address the Copilot review on AcademySoftwareFoundation#2225: every kernel-facing binding that indexes
a user array by grid value index or leaf ID now checks the array covers the
grid's count (ValueError instead of an out-of-bounds device access), geometric
parameters (voxelSize, halfWidth) are rejected unless finite and positive
before reaching Map::set (which only debug-asserts), and the CAI / DLPack
exports order the consumer against the buffer's tracked prior uses — the CAI
export backs its stream=1 claim by ordering the legacy default stream, and
__dlpack__ honors the protocol's consumer-stream argument. The ordering
activates once DeviceBuffer::orderAfterPriorUses becomes public upstream; the
call sites are written against those semantics already.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants