Skip to content

Add TSDF / ESDF / Occupancy / Decay APIs and a fast marching-cubes variant - #656

Draft
fwilliams wants to merge 1 commit into
feat/nanovdb-allocator-overridesfrom
feat/tsdf-esdf-stack
Draft

Add TSDF / ESDF / Occupancy / Decay APIs and a fast marching-cubes variant#656
fwilliams wants to merge 1 commit into
feat/nanovdb-allocator-overridesfrom
feat/tsdf-esdf-stack

Conversation

@fwilliams

Copy link
Copy Markdown
Collaborator

Summary

Native CUDA kernels and Python wrappers for TSDF and ESDF reconstruction, occupancy mapping, dynamic-scene decay, and a fast sparse-compact marching-cubes variant. These features share a common PersistentTSDFState + BuildPointTruncationShell substrate.

Stacked on top of #655 (nanoVDB allocator overrides). That PR should land first; the multi-frame integrator paths here depend on the allocator routing it sets up.

Topology and state primitives

  • BuildPointTruncationShell.{cu,h} — shared primitive that turns (points, base_grid, truncation_margin) into the set of voxels within the truncation shell. Used by both the depth and LiDAR TSDF integrators.

  • PersistentTSDFState.{cu,h} — grow-on-touch state holder for incremental integration. Wraps a monotonically-growing live grid with fixed-shape tsdf / weights / optional features sidecar tensors and exposes a grow method that expands the grid + sidecars atomically while preserving values at already-live voxels.

  • PersistentTSDFStateBinding.cpp — pybind11 binding for the above.

Integrators

  • IntegrateTSDF.{cu,h} (modified) — depth TSDF integrator now uses BuildPointTruncationShell and PersistentTSDFState, and exposes a new N-frame batched entry point integrateTSDFBatch that grows the union grid one frame at a time and copy-forwards sidecars through the persistent-state object. Bit-identical to the per-frame loop (pinned by test_integrate_tsdf_frames_matches_sequential).

  • IntegrateTSDFFromPoints.{cu,h} — native LiDAR / range-sensor TSDF integrator. Per-point thread HDDA-walks the union grid and atomicAdds a running-sum into (sum_w_sdf, sum_w, sum_w_feat) accumulators within the truncation (and optionally free-space) band. Single-frame, with-features, and N-frames-batched variants.

  • IntegrateOccupancyFromPoints.{cu,h} — LiDAR occupancy mapping with free-space carving and log-odds updates. Single-frame and N-frames-batched variants. Same ray-walk structure as the LiDAR TSDF integrator.

ESDF

  • ComputeESDF.{cu,h} — Euclidean Signed Distance Field from an integrated narrow-band TSDF. Composition pattern is dilateGrid -> esdfSeed -> N sweeps of 26-neighbour min-propagation, reusing the existing topology-op primitives.

  • DirtyMaskFromSidecars.{cu,h} — per-voxel dirty-mask primitive that lets the incremental ESDF variant scope work to just the voxels whose sidecars changed.

Marching cubes

  • MarchingCubesFast.{cu,h} — sparse-compact, packed-key marching cubes for fp32 / fp16 CUDA. Surface-voxel compaction + 1-D packed-key dedup + an in-register fp16-to-fp32 cast on load. Bit-identical to the legacy at fp32 and numerically-identical at fp16 (kernels are templated on the input scalar type so there's no transient fp32 buffer).

  • MarchingCubes.{cu,h} (modified) — the public marchingCubes dispatcher routes to MarchingCubesFast for fp32 / fp16 CUDA inputs and to the previous implementation (now named marchingCubesLegacy, kept verbatim) for other dtype / device combinations.

Python surface

Test plan

New unit tests cover the persistent-state invariants, bit-identity of the batched-vs-sequential TSDF paths, atomic-noise tolerance for the LiDAR / occupancy variants, and fp16-vs-fp32 numerical agreement for the new marching-cubes fast path:

  • tests/unit/test_persistent_tsdf_state.py — 7 tests
  • tests/unit/test_compute_esdf.py — 17 tests
  • tests/unit/test_dirty_mask.py — 9 tests
  • tests/unit/test_integrate_occupancy.py — 7 tests
  • tests/unit/test_decay_and_prune.py — 9 tests
  • tests/unit/test_basic_ops.py — 16 new TSDF / MC fp16 / occupancy tests

All pass locally on the freshly-built install. Full pytest unit/ suite is green.

cd tests && pytest unit/

Followups

Cross-library benchmark drivers (Replica, KITTI, Mai City, 7-Scenes, fvdb vs nvblox / VDBFusion / Open3D) for these new APIs live in a companion PR against openvdb/fvdb-reality-capture.

@fwilliams
fwilliams force-pushed the feat/nanovdb-allocator-overrides branch from 14613ce to 4ec7afe Compare May 19, 2026 17:44
@fwilliams
fwilliams force-pushed the feat/tsdf-esdf-stack branch from 859f849 to 959d383 Compare May 19, 2026 17:47
@fwilliams
fwilliams force-pushed the feat/nanovdb-allocator-overrides branch from 4ec7afe to bedf87e Compare May 19, 2026 17:49
@fwilliams
fwilliams force-pushed the feat/tsdf-esdf-stack branch 2 times, most recently from 31dde5d to a0bd959 Compare May 19, 2026 19:09
@fwilliams
fwilliams force-pushed the feat/nanovdb-allocator-overrides branch from 010424d to c700a40 Compare May 19, 2026 19:41
@fwilliams
fwilliams force-pushed the feat/tsdf-esdf-stack branch from a0bd959 to 63b12a2 Compare May 19, 2026 19:43
Add native CUDA kernels and Python wrappers for TSDF and ESDF
reconstruction, occupancy mapping, dynamic-scene decay, and a fast
sparse-compact marching-cubes variant. These features sit on top of
the nanoVDB allocator-overrides change (parent PR) and share a
common `PersistentTSDFState` + `BuildPointTruncationShell` substrate.

Topology + state primitives

  src/fvdb/detail/ops/BuildPointTruncationShell.{cu,h}
      Shared primitive that turns `(points, base_grid, truncation_margin)`
      into the set of voxels within the truncation shell. Used by both
      depth and LiDAR TSDF integrators.

  src/fvdb/detail/ops/PersistentTSDFState.{cu,h}
      Grow-on-touch state holder for incremental integration: wraps a
      monotonically-growing live grid with fixed-shape tsdf / weights /
      optional feature sidecars and exposes a `grow` method that
      expands the grid + sidecars atomically while preserving values
      at already-live voxels.

  src/python/PersistentTSDFStateBinding.cpp
      Pybind11 binding for the above.

Integrators

  src/fvdb/detail/ops/IntegrateTSDF.{cu,h}  (modified)
      Depth TSDF integrator now uses `BuildPointTruncationShell` and
      `PersistentTSDFState`, and exposes a new N-frame batched entry
      point `integrateTSDFBatch` that grows the union grid one frame
      at a time and copy-forwards sidecars through the
      persistent-state object. Bit-identical to the per-frame loop
      (pinned by `test_integrate_tsdf_frames_matches_sequential`).

  src/fvdb/detail/ops/IntegrateTSDFFromPoints.{cu,h}
      Native LiDAR / range-sensor TSDF integrator: per-point thread
      HDDA-walks the union grid and `atomicAdd`s a running-sum into
      (sum_w_sdf, sum_w, sum_w_feat) accumulators within the
      truncation (and optionally free-space) band. Single-frame,
      with-features, and N-frames-batched variants.

  src/fvdb/detail/ops/IntegrateOccupancyFromPoints.{cu,h}
      LiDAR occupancy mapping with free-space carving and log-odds
      updates. Single-frame and N-frames-batched variants. Same
      ray-walk structure as the LiDAR TSDF integrator.

ESDF

  src/fvdb/detail/ops/ComputeESDF.{cu,h}
      Euclidean Signed Distance Field from an integrated narrow-band
      TSDF. Composition pattern is
      `dilateGrid -> esdfSeed -> N sweeps of 26-N min-propagation`,
      reusing the topology-op primitives.

  src/fvdb/detail/ops/DirtyMaskFromSidecars.{cu,h}
      Per-voxel dirty-mask primitive that lets the incremental ESDF
      variant scope work to just the voxels whose sidecars changed.

Marching cubes

  src/fvdb/detail/ops/MarchingCubesFast.{cu,h}
      Sparse-compact, packed-key marching cubes for fp32 / fp16 CUDA.
      `marchingCubes` now dispatches to this for eligible inputs and
      to `marchingCubesLegacy` (the previous default, kept verbatim)
      otherwise.

  src/fvdb/detail/ops/MarchingCubes.{cu,h}  (modified)
      Routes through to the new fast path.

Python surface

  fvdb/functional/_meshing.py
      Wrappers for the new N-frame + with-features + LiDAR variants
      of TSDF integration, occupancy mapping (single + frames), and
      ESDF (single + incremental).

  fvdb/functional/_topology.py
      Wrapper for `dirty_mask_from_sidecars_single`.

  fvdb/grid.py
      New methods on `Grid`: `decay_and_prune`,
      `integrate_tsdf_frames`, `integrate_tsdf_with_features`,
      `integrate_tsdf_from_points` (+ frames + with-features
      variants), `integrate_occupancy_from_points` (+ frames),
      `compute_esdf`, `compute_esdf_incremental`. `decay_and_prune`
      is implemented entirely in Python on top of existing fvdb
      sidecar + topology primitives.

  fvdb/functional/__init__.py
      Export the new functional names.

  src/python/Bindings.cpp, src/python/GridBatchOps.cpp
      Register the new C++ bindings.

Tests

  tests/unit/test_persistent_tsdf_state.py
  tests/unit/test_compute_esdf.py
  tests/unit/test_dirty_mask.py
  tests/unit/test_integrate_occupancy.py
  tests/unit/test_decay_and_prune.py
  tests/unit/test_basic_ops.py  (extended)

      Cover the new primitives, the persistent-state invariants
      (`grow` semantics, sidecar carry-forward), bit-identity of the
      batched-vs-sequential TSDF paths, atomic-noise tolerance for
      the LiDAR/occupancy variants, and fp16-vs-fp32 numerical
      agreement for the new marching-cubes fast path.

Signed-off-by: Francis Williams <francis@fwilliams.info>
@fwilliams
fwilliams force-pushed the feat/tsdf-esdf-stack branch from 63b12a2 to cd63cb3 Compare May 19, 2026 20:04
swahtz added a commit that referenced this pull request Aug 20, 2026
… via upstream memory-resource seams (#732)

## Summary

fvdb's grid builders allocate their device scratch from nanoVDB's
default `DeviceResource` — a second `cudaMallocAsync` pool that
partitions VRAM against PyTorch's. Large workloads (e.g. multi-frame
TSDF integration) then hit a clean OOM even when the GPU has free memory
in aggregate.

This routes that scratch — O(N-points) sort keys, CUB temp storage,
topology mask buffers — through PyTorch's CUDA allocator instead, so it
shares one pool with fvdb / PyTorch tensors. **22 sites across 13 `.cu`
files plus `PadGrid.cuh`.**

Note this is not hardcoded to Torch's *native* caching allocator:
`c10::cuda::CUDACachingAllocator` is a namespace, and its
`raw_alloc_with_stream` / `raw_delete` free functions dispatch through
`CUDACachingAllocator::get()` — the runtime-swappable allocator Torch
itself allocates tensors from. fvdb's scratch therefore follows whatever
allocator the user has installed: the native caching allocator
(including `PYTORCH_CUDA_ALLOC_CONF` knobs), the `cudaMallocAsync`
backend (`PYTORCH_CUDA_ALLOC_CONF=backend:cudaMallocAsync`), or a custom
allocator installed via
`torch.cuda.memory.change_current_allocator(CUDAPluggableAllocator(...))`.

**Supersedes #655**, which vendored modified nanoVDB headers into the
tree. This instead uses the injectable-memory-resource seams we
developed upstream (AcademySoftwareFoundation/openvdb#2232; PRs
[#2268](AcademySoftwareFoundation/openvdb#2268),
[#2269](AcademySoftwareFoundation/openvdb#2269),
[#2270](AcademySoftwareFoundation/openvdb#2270),
[#2272](AcademySoftwareFoundation/openvdb#2272),
[#2273](AcademySoftwareFoundation/openvdb#2273))
— now merged, so the pin is plain upstream master. No fork, no
include-path shadowing, no resync procedure.

## What's in this PR

1. **Pin nanovdb to upstream master.** `src/cmake/get_nanovdb.cmake` →
`AcademySoftwareFoundation/openvdb @ 7946f17e`, which includes the
small-builder `ResourceT` seams
([#2286](AcademySoftwareFoundation/openvdb#2286)),
the synchronous resource adapters
([#2272](AcademySoftwareFoundation/openvdb#2272)),
and the `MeshToGrid` `CALL_CUBS` `#undef` fix
([#2284](AcademySoftwareFoundation/openvdb#2284)).

2. **`fvdb::TorchResource`.** A ~40-line stateless resource
([`src/fvdb/TorchResource.h`](src/fvdb/TorchResource.h)) modeling
nanoVDB's stream-ordered `AsyncResource` concept over
`c10::cuda::CUDACachingAllocator::raw_alloc_with_stream` / `raw_delete`
— the dispatchers to Torch's currently active CUDA allocator (see
Summary). Passed as the `ResourceT` template parameter at all 13
upstream builder call sites — `voxelsToGrid`, `DilateGrid`,
`MergeGrids`, `PruneGrid`, `RefineGrid`, `CoarsenGrid` — always via the
`fvdb::BuilderResource` alias
([`src/fvdb/BuilderResource.h`](src/fvdb/BuilderResource.h)), never
named directly, so the allocator policy lives in a single line (a
non-torch build, e.g. the ONNX Runtime EP planned in #579, retargets the
alias there instead of touching every op). Being stateless, it binds
through each builder's defaulted constructor argument, so no instance is
plumbed through. Retains #655's `FVDB_NANOVDB_TRACE_ALLOCS` tracing
(`=1` traces ≥ 256 KiB, a value starting with `2` traces everything).

3. **`PadGrid` gains a `ResourceT` seam.** The conv builders used
`DilateGrid<..., TorchResource>` for odd kernels and `PadGrid` — on the
rival pool — for even ones: same loop, same grid, a different allocator
depending on kernel parity. fvdb's own `morphology::PadGrid` drives
nanoVDB's `TopologyBuilder` (internal mask buffers, `countNodes` CUB
scratch, `TempPool`) but hardcoded `DeviceResource`. It now takes a
`ResourceT` parameter mirroring the upstream `DilateGrid` signature and
forwards it, with `BuilderResource` passed at all 7 call sites. The
default keeps it source-compatible.

4. **CUB scratch in `BuildFineGridFromCoarse`.**
`cub::DeviceSegmentedReduce` temp storage used a bare `cudaMallocAsync`;
it now routes through `BuilderResource`. Both `cub` calls are also now
`C10_CUDA_CHECK`-wrapped — previously unchecked, as was the allocation.

5. **The `SaveNanoVDB` CUDA path.** The save path allocated its largest
device buffers from nanoVDB's default pool: the per-batch
`(N+1)`-element value staging buffer, the `indexToGrid` output grid
handle, and the defensive host-upload buffer. All three now use
`TorchDeviceBuffer`, and `indexToGrid`'s internal scratch routes through
`TorchResource` via the #2286 seam. Stream-ordering is preserved: the
replaced stream-ordered `DeviceBuffer` constructors become `raw_alloc`
on the same current stream the copies and kernels are queued on. The
host path (`indexToGridHost`) and the `HostBuffer` file-staging buffers
are unchanged.

## Not routed (no upstream seam yet; all off the hot paths)

- `DistributedPointsToGrid` multi-GPU scratch (deferred upstream behind
AcademySoftwareFoundation/openvdb#2248) — the most valuable remaining
seam
- `VoxelBlockManager` / `buildVoxelBlockManager` scratch in
`ReinitializeSdf.cu`
- the builders' small dual-space `mProcessedRoot` / `mData` buffers
(upstream roadmap Step 3)

`MeshToGrid` is the one merged seam fvdb does not use:
`BuildGridFromMesh.cu` does its own parametric surface sampling and goes
through `_createNanoGridFromIJK`, so there is nothing to route.

## Test plan

- [x] `./build.sh install` succeeds on a clean tree against the new pin
(full CUDA build, `-Werror`).
- [x] Injection verified live via `FVDB_NANOVDB_TRACE_ALLOCS`: a
500k-point `Grid.from_points` + `dilated_grid(2)` prints 42
`TorchResource` traces with correct results (484,631 → 8,461,871
voxels); `from_nearest_voxels_to_points` at 2M points shows `PadGrid`
scratch routed.
- [x] 584 tests passed — conv semantics + integration (203),
conv/conv-transpose default + prune + empty grids (103), basic ops (276,
1 skipped), sliced batch (2, covering the `BuildFineGridFromCoarse` CUB
path).
- [x] `test_io.py` — 622 passed against the `7946f17e` pin; a traced
`save_nanovdb` (`FVDB_NANOVDB_TRACE_ALLOCS=2`) shows the `indexToGrid`
scratch flowing through `TorchResource`.

## Followups

- Rebase the TSDF / ESDF / Occupancy stack (#656) onto this branch in
place of #655, threading `TorchResource` through the new ops it adds.

🤖 Generated with [Claude Code](https://claude.com/claude-code)

---------

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

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant