diff --git a/.github/workflows/nanovdb.yml b/.github/workflows/nanovdb.yml index ababa6c2ba..2f9ab09ec3 100644 --- a/.github/workflows/nanovdb.yml +++ b/.github/workflows/nanovdb.yml @@ -74,7 +74,11 @@ jobs: echo "/usr/local/cuda-12/bin" >> $GITHUB_PATH echo "LD_LIBRARY_PATH=/usr/local/cuda-12/lib64:$LD_LIBRARY_PATH" >> $GITHUB_ENV - name: nanobind - run: ./ci/install_nanobind.sh 2.5.0 + # The NanoVDB Python bindings (CUDA interop, DLPack) track a newer + # nanobind than the repo-wide 2.5.0 floor. Bump here only; the shared + # build.yml still co-compiles the bindings against 2.5.0, so the + # binding code must stay >=2.5.0 source-compatible. + run: ./ci/install_nanobind.sh 2.12.0 - name: build # NOTE: CMAKE_POSITION_INDEPENDENT_CODE set to fix default behaviour change in clang 14 # https://github.com/AcademySoftwareFoundation/aswf-docker/issues/228 diff --git a/nanovdb/nanovdb/python/CMakeLists.txt b/nanovdb/nanovdb/python/CMakeLists.txt index e7b0bab3ae..899d5cb095 100644 --- a/nanovdb/nanovdb/python/CMakeLists.txt +++ b/nanovdb/nanovdb/python/CMakeLists.txt @@ -32,13 +32,36 @@ nanobind_add_module(nanovdb_python NB_STATIC PyTree.cc PyVoxelBlockManager.cc cuda/PyDeviceBuffer.cc + cuda/PyUnifiedBuffer.cu + cuda/PyUnifiedGridHandle.cu + cuda/PyDeviceMesh.cu + cuda/PyDeviceStreamMap.cu + cuda/PyTempPool.cu + cuda/PyDistributedPointsToGrid.cu cuda/PyDeviceGridHandle.cu + cuda/PyDeviceNodeManager.cu + cuda/PyDeviceVoxelBlockManager.cu cuda/PyPointsToGrid.cu cuda/PySampleFromVoxels.cu cuda/PySignedFloodFill.cu + cuda/PyDilateGrid.cu + cuda/PyCoarsenGrid.cu + cuda/PyRefineGrid.cu + cuda/PyPruneGrid.cu + cuda/PyMergeGrids.cu + cuda/PyInjectData.cu + cuda/PyIndexToGrid.cu + cuda/PyMeshToGrid.cu + cuda/PyAddBlindData.cu + cuda/PyDeviceGridStats.cu + cuda/PyDeviceGridValidator.cu + cuda/PyDeviceGridChecksum.cu ) -target_include_directories(nanovdb_python PRIVATE ${CUDA_INCLUDE_DIRECTORY}) +# CMAKE_CURRENT_SOURCE_DIR lets sources in cuda/ include the shared binding +# headers by their plain names ("PyGridHandle.h", "BuildTypes.def") instead of +# "../"-relative paths, matching the include style used across the codebase. +target_include_directories(nanovdb_python PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${CUDA_INCLUDE_DIRECTORY}) target_link_libraries(nanovdb_python PRIVATE nanovdb ${CUDA_LIBRARIES} ${NANOVDB_BLOSC} ${NANOVDB_ZLIB} ${NANOVDB_OPENVDB} ${NANOVDB_TBB}) target_compile_definitions(nanovdb_python PRIVATE ${NANOVDB_USE_CUDA_FLAG} ${NANOVDB_USE_BLOSC_FLAG} ${NANOVDB_USE_ZLIB_FLAG} ${NANOVDB_USE_OPENVDB_FLAG} ${NANOVDB_USE_TBB_FLAG}) set_target_properties(nanovdb_python PROPERTIES OUTPUT_NAME "nanovdb") @@ -99,6 +122,14 @@ if(NANOVDB_BUILD_PYTHON_UNITTESTS) COMMAND ${Python_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/test/TestNanoVDB.py -v WORKING_DIRECTORY "${NANOVDB_PYTHON_WORKING_DIR}") + # GPU interop suite. Self-skips when the build has no CUDA or no GPU is + # present (so it is harmless to register unconditionally), but is a + # separate test file that the pytest_nanovdb command does not discover, + # so register it as its own ctest test mirroring the block above. + add_test(NAME pytest_nanovdb_gpu_interop + COMMAND ${Python_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/test/TestGpuInterop.py -v + WORKING_DIRECTORY "${NANOVDB_PYTHON_WORKING_DIR}") + # Smoke-run every example script (they self-skip when optional # dependencies such as NumPy or OpenVDB are unavailable). add_test(NAME pytest_nanovdb_examples @@ -110,8 +141,10 @@ if(NANOVDB_BUILD_PYTHON_UNITTESTS) set(PYTHONPATH "$ENV{PYTHONPATH};${NANOVDB_PYTHON_WORKING_DIR}") string(REPLACE "\\;" ";" PYTHONPATH "${PYTHONPATH}") string(REPLACE ";" "\\;" PYTHONPATH "${PYTHONPATH}") - set_tests_properties(pytest_nanovdb pytest_nanovdb_examples PROPERTIES ENVIRONMENT "PYTHONPATH=${PYTHONPATH}") + set_tests_properties(pytest_nanovdb pytest_nanovdb_gpu_interop pytest_nanovdb_examples + PROPERTIES ENVIRONMENT "PYTHONPATH=${PYTHONPATH}") else() - set_tests_properties(pytest_nanovdb pytest_nanovdb_examples PROPERTIES ENVIRONMENT "PYTHONPATH=$ENV{PYTHONPATH}:${NANOVDB_PYTHON_WORKING_DIR}") + set_tests_properties(pytest_nanovdb pytest_nanovdb_gpu_interop pytest_nanovdb_examples + PROPERTIES ENVIRONMENT "PYTHONPATH=$ENV{PYTHONPATH}:${NANOVDB_PYTHON_WORKING_DIR}") endif() endif() diff --git a/nanovdb/nanovdb/python/NanoVDBModule.cc b/nanovdb/nanovdb/python/NanoVDBModule.cc index 38b69f2db8..010d4259da 100644 --- a/nanovdb/nanovdb/python/NanoVDBModule.cc +++ b/nanovdb/nanovdb/python/NanoVDBModule.cc @@ -8,15 +8,25 @@ #include #include // for __repr__ +#include // host updateChecksum for the header setters #ifdef NANOVDB_USE_CUDA #include #endif +#include #include #include "cuda/PyDeviceBuffer.h" +#ifdef NANOVDB_USE_CUDA +#include "cuda/PyUnifiedBuffer.h" +#include "cuda/PyUnifiedGridHandle.h" +#include "cuda/PyDeviceMesh.h" +#include "cuda/PyDeviceStreamMap.h" +#include "cuda/PyTempPool.h" +#endif #include "PyBuildGrid.h" #include "PyGridHandle.h" +#include "PyValidate.h" #include "PyHostBuffer.h" #include "PyIO.h" #include "PyMath.h" @@ -236,21 +246,63 @@ void defineGrid(nb::module_& m) "the affine transform, grid flags and blind-data accessors. Concrete " "BuildT-typed subclasses (FloatGrid, Vec3fGrid, ...) add tree access " "and per-voxel queries.") - // Validation and flag mutators (already member functions on GridData). + // Validation and header mutators. Every mutator refreshes the grid's + // checksum afterwards (preserving its mode; a no-op when the checksum is + // disabled), so editing a header field never silently desyncs a + // checksummed grid -- a validating reader would otherwise reject it. .def("isValid", &GridData::isValid, "True iff the grid header looks consistent (magic / version / class tags).") - .def("setMinMaxOn", &GridData::setMinMaxOn, "on"_a = true, - "Toggle the HasMinMax grid flag.") - .def("setBBoxOn", &GridData::setBBoxOn, "on"_a = true, - "Toggle the HasBBox grid flag.") - .def("setLongGridNameOn", &GridData::setLongGridNameOn, "on"_a = true, - "Toggle the HasLongGridName grid flag.") - .def("setAverageOn", &GridData::setAverageOn, "on"_a = true, - "Toggle the HasAverage grid flag.") - .def("setStdDeviationOn", &GridData::setStdDeviationOn, "on"_a = true, - "Toggle the HasStdDeviation grid flag.") - .def("setGridName", &GridData::setGridName, "src"_a, - "Overwrite the grid's short name (truncated to the in-header buffer).") + .def("setMinMaxOn", + [](GridData& g, bool on) { g.setMinMaxOn(on); nanovdb::tools::updateChecksum(&g); }, + "on"_a = true, "Toggle the HasMinMax grid flag (refreshes the checksum).") + .def("setBBoxOn", + [](GridData& g, bool on) { g.setBBoxOn(on); nanovdb::tools::updateChecksum(&g); }, + "on"_a = true, "Toggle the HasBBox grid flag (refreshes the checksum).") + .def("setLongGridNameOn", + [](GridData& g, bool on) { g.setLongGridNameOn(on); nanovdb::tools::updateChecksum(&g); }, + "on"_a = true, "Toggle the HasLongGridName grid flag (refreshes the checksum).") + .def("setAverageOn", + [](GridData& g, bool on) { g.setAverageOn(on); nanovdb::tools::updateChecksum(&g); }, + "on"_a = true, "Toggle the HasAverage grid flag (refreshes the checksum).") + .def("setStdDeviationOn", + [](GridData& g, bool on) { g.setStdDeviationOn(on); nanovdb::tools::updateChecksum(&g); }, + "on"_a = true, "Toggle the HasStdDeviation grid flag (refreshes the checksum).") + .def("setGridName", + [](GridData& g, const char* src) { + const bool ok = g.setGridName(src); + nanovdb::tools::updateChecksum(&g); + return ok; + }, + "src"_a, + "Overwrite the grid's short name (truncated to the in-header buffer; " + "refreshes the checksum). Returns False if the name had to be truncated.") + .def("setGridClass", + [](GridData& g, GridClass gridClass) { + g.mGridClass = gridClass; + nanovdb::tools::updateChecksum(&g); + }, + "gridClass"_a, + "Set the grid's GridClass (e.g. GridClass.LevelSet) in place and refresh " + "the checksum. Host twin of tools.cuda.setGridClass for device grids.") + .def("setTransform", + [](GridData& g, double voxelSize, const Vec3d& translation) { + // Map::set only debug-asserts positivity, so validate here to + // keep a singular / non-finite transform out of the header. + requirePositiveFinite(voxelSize, "setTransform", "voxelSize"); + // Uniform scale + translation index->world map, then recompute the + // world-space AABB from the index bbox under the new map (mirrors + // how GridStats sets mWorldBBox). + g.mMap.set(voxelSize, translation, 1.0); + g.mVoxelSize = g.mMap.getVoxelSize(); + const CoordBBox& ib = g.indexBBox(); + g.mWorldBBox = CoordBBox(ib[0], ib[1].offsetBy(1)).transform(g.mMap); + nanovdb::tools::updateChecksum(&g); + }, + "voxelSize"_a, "translation"_a = Vec3d(0.0, 0.0, 0.0), + "Set a uniform-scale + translation index->world transform (voxelSize is " + "world units per voxel; translation is in world space), recompute the " + "world bounding box, and refresh the checksum. For non-uniform or rotated " + "maps, build the grid with the desired transform instead.") // Affine transforms (already member functions on GridData). .def("applyMap", nb::overload_cast(&GridData::template applyMap, nb::const_), "xyz"_a, "Transform an index-space point to world space using this grid's Map.") @@ -464,7 +516,16 @@ template void defineNanoGrid(nb::module_& m, const char* name) nb::overload_cast<>(&NanoGrid::tree, nb::const_), nb::rv_policy::reference_internal, "Return the tree associated with this grid. Lifetime is " - "anchored to the grid (and therefore to the GridHandle)."); + "anchored to the grid (and therefore to the GridHandle).") + .def("data_ptr", + [](const NanoGrid& grid) { + return reinterpret_cast(&grid); + }, + "Raw pointer to this grid as a Python int. The address is a HOST " + "pointer when the grid came from handle.grid(n) and a DEVICE " + "pointer when it came from handle.deviceGrid(n) — provenance is " + "the caller's responsibility, the grid object itself cannot tell " + "host from device."); // Grid::valueCount / pointCount are SFINAE-gated in C++ to the // index and Point BuildTs respectively — mirror that gating here. if constexpr (BuildTraits::is_index) { @@ -1218,8 +1279,25 @@ NB_MODULE(nanovdb, m) defineHostGridHandle(m); #ifdef NANOVDB_USE_CUDA - defineDeviceBuffer(m); - defineDeviceGridHandle(m); + nb::module_ cudaModule = m.def_submodule("cuda"); + cudaModule.doc() = "CUDA device buffers, the device GridHandle, and device infrastructure (mirrors nanovdb::cuda)."; + defineDeviceBuffer(cudaModule); + // UnifiedBuffer (CUDA managed memory) + its GridHandle. The unified handle + // is the type returned by nanovdb.tools.cuda.DistributedPointsToGrid. + defineUnifiedBuffer(cudaModule); + defineUnifiedGridHandle(cudaModule); + defineDeviceGridHandle(cudaModule); + // Device NodeManager (DeviceNodeManagerHandle + createDeviceNodeManager) + // on nanovdb.cuda, alongside the device GridHandle. + defineDeviceNodeManager(cudaModule); + // Multi-GPU / device infrastructure (mirrors nanovdb::cuda): the device + // mesh, the device->stream map, and the CUB temp-storage pool primitives. + defineDeviceMesh(cudaModule); + defineDeviceStreamMap(cudaModule); + defineTempPool(cudaModule); + // Device VoxelBlockManager (mirrors nanovdb::tools::cuda) is registered on + // the existing nanovdb.tools.cuda submodule in PyTools.cc — NOT here — so + // the Python layout matches the C++ namespaces. #endif nb::module_ toolsModule = m.def_submodule("tools"); diff --git a/nanovdb/nanovdb/python/PyMath.cc b/nanovdb/nanovdb/python/PyMath.cc index 2e212f09b9..d759da8a34 100644 --- a/nanovdb/nanovdb/python/PyMath.cc +++ b/nanovdb/nanovdb/python/PyMath.cc @@ -486,7 +486,7 @@ template void defineBBoxInteger(nb::module_& m, const char* nam "Construct a cube of side dim voxels anchored at the min Coord.") .def_static("createCube", nb::overload_cast(&math::BBox::createCube), "min"_a, "max"_a, "Construct a cube spanning [min, max] in every axis.") - .def("is_divisible", &math::BBox::is_divisible, + .def("isDivisible", &math::BBox::is_divisible, "True iff this CoordBBox has more than one voxel in every axis.") .def("empty", &math::BBox::empty, "True iff this CoordBBox is empty (any min component > the matching max).") diff --git a/nanovdb/nanovdb/python/PyTools.cc b/nanovdb/nanovdb/python/PyTools.cc index 9aa8f93452..fbc0e5fd99 100644 --- a/nanovdb/nanovdb/python/PyTools.cc +++ b/nanovdb/nanovdb/python/PyTools.cc @@ -14,10 +14,24 @@ #include "PyGridChecksum.h" #include "PyGridValidator.h" #include "PyNanoToOpenVDB.h" +#include "PyVoxelBlockManager.h" // for defineDeviceVoxelBlockManager (CUDA) #ifdef NANOVDB_USE_CUDA #include "cuda/PyPointsToGrid.h" +#include "cuda/PyDistributedPointsToGrid.h" #include "cuda/PySampleFromVoxels.h" #include "cuda/PySignedFloodFill.h" +#include "cuda/PyDilateGrid.h" +#include "cuda/PyCoarsenGrid.h" +#include "cuda/PyRefineGrid.h" +#include "cuda/PyPruneGrid.h" +#include "cuda/PyMergeGrids.h" +#include "cuda/PyInjectData.h" +#include "cuda/PyIndexToGrid.h" +#include "cuda/PyMeshToGrid.h" +#include "cuda/PyAddBlindData.h" +#include "cuda/PyDeviceGridStats.h" +#include "cuda/PyDeviceGridValidator.h" +#include "cuda/PyDeviceGridChecksum.h" #endif namespace nb = nanobind; @@ -60,10 +74,183 @@ void defineToolsModule(nb::module_& m) defineSignedFloodFill(cudaModule, "signedFloodFill"); defineSignedFloodFill(cudaModule, "signedFloodFill"); - definePointsToGrid(cudaModule, "pointsToRGBA8Grid"); + // Coordinate-input (index-space int32 (N,3)) -> grid. The legacy Rgba8 + // entry keeps its original Python name; OnIndex/Index get descriptive + // names. (Point is excluded here -- see PyPointsToGrid.cu -- and is built + // from world positions via pointsToGrid below.) + defineVoxelsToGrid(cudaModule, "pointsToRGBA8Grid"); + defineVoxelsToGrid(cudaModule, "voxelsToRGBA8Grid"); + defineVoxelsToGrid(cudaModule, "voxelsToOnIndexGrid"); + defineVoxelsToGrid(cudaModule, "voxelsToIndexGrid"); + + // World-position-input ((N,3) float OR double) -> NanoGrid. Both + // scalar precisions are bound under the same Python name; nanobind picks + // the overload that matches the input tensor dtype. + definePointsToGrid(cudaModule, "pointsToGrid"); + definePointsToGrid(cudaModule, "pointsToGrid"); + + // Triangle mesh -> narrow-band UDF (nanovdb::tools::cuda::MeshToGrid). + // Returns (ValueOnIndex device handle, per-value float UDF sidecar buffer). + defineMeshToGrid(cudaModule, "meshToGrid"); + + // Multi-GPU voxel-coordinate -> grid builder (nanovdb::tools::cuda:: + // DistributedPointsToGrid). Distributes an (N, 3) int32 unified-memory + // array of index-space voxel coordinates over a DeviceMesh. Bound for the + // index / Rgba8 build types (matching the voxelsTo*Grid set); Point is + // excluded (its coords must be world-space Vec3, not int32 Coord). + defineDistributedPointsToGrid(cudaModule, "DistributedPointsToGrid"); + defineDistributedPointsToGrid(cudaModule, "DistributedIndexPointsToGrid"); + defineDistributedPointsToGrid(cudaModule, "DistributedRGBA8PointsToGrid"); defineSampleFromVoxels(cudaModule, "sampleFromVoxels"); defineSampleFromVoxels(cudaModule, "sampleFromVoxels"); + + // Topological/morphological ops on OnIndex grids (nanovdb::tools::cuda). + // Each is constrained to OnIndex build types by TopologyBuilder's + // is_onindex static_assert, so only ValueOnIndex is instantiated. + defineDilateGrid(cudaModule, "dilateGrid"); + defineCoarsenGrid(cudaModule, "coarsenGrid"); + defineRefineGrid(cudaModule, "refineGrid"); + definePruneGrid(cudaModule, "pruneGrid"); + defineMergeGrids(cudaModule, "mergeGrids"); + + // Sidecar value transfer across a topology change, and the predicate->mask + // helper that feeds pruneGrid (nanovdb::util::cuda::Inject* functors). + defineInject(cudaModule, "inject"); + defineInject(cudaModule, "inject"); + defineInject(cudaModule, "inject"); + defineInject(cudaModule, "inject"); + defineInjectFeatures(cudaModule, "inject"); + defineInjectFeatures(cudaModule, "inject"); + defineInjectFeatures(cudaModule, "inject"); + defineInjectFeatures(cudaModule, "inject"); + defineInjectPredicateToMask(cudaModule, "injectPredicateToMask"); + defineInjectGridMask(cudaModule, "injectGridMask"); + + // Device VoxelBlockManager (nanovdb::tools::cuda) on nanovdb.tools.cuda. + defineDeviceVoxelBlockManager(cudaModule); + + // IndexGrid -> regular Grid (nanovdb::tools::cuda::indexToGrid). Source is + // an index grid (ValueIndex / ValueOnIndex); destination value type is a + // non-special type (float / double scalar, or Vec3f / Vec3d). All register + // under "indexToGrid"; nanobind disambiguates on the source grid class and + // the values ndarray dtype/shape. + defineIndexToGridScalar(cudaModule, "indexToGrid"); + defineIndexToGridScalar(cudaModule, "indexToGrid"); + defineIndexToGridScalar(cudaModule, "indexToGrid"); + defineIndexToGridScalar(cudaModule, "indexToGrid"); + defineIndexToGridScalar(cudaModule, "indexToGrid"); + defineIndexToGridScalar(cudaModule, "indexToGrid"); + defineIndexToGridVec3(cudaModule, "indexToGrid"); + defineIndexToGridVec3(cudaModule, "indexToGrid"); + defineIndexToGridVec3(cudaModule, "indexToGrid"); + defineIndexToGridVec3(cudaModule, "indexToGrid"); + + // Append blind data to a device grid (nanovdb::tools::cuda::addBlindData). + // Registered under "addBlindData" for a set of (grid BuildT, blind-data + // element type) combinations; nanobind disambiguates on the grid class and + // the blindData ndarray dtype. + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + // Signed-integer blind payloads (Int32 / Int64), e.g. integer labels / ids. + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + defineAddBlindData(cudaModule, "addBlindData"); + + // Device quality-control tools (mirror the host tools.* names on + // tools.cuda). updateGridStats covers scalar/vector/bool grids; isValid + // and the checksum entries cover the full callNanoGrid BuildT set. + defineDeviceUpdateGridStats(cudaModule, "updateGridStats"); + defineDeviceUpdateGridStats(cudaModule, "updateGridStats"); + defineDeviceUpdateGridStats(cudaModule, "updateGridStats"); + defineDeviceUpdateGridStats(cudaModule, "updateGridStats"); + defineDeviceUpdateGridStats(cudaModule, "updateGridStats"); + defineDeviceUpdateGridStats(cudaModule, "updateGridStats"); + defineDeviceUpdateGridStats(cudaModule, "updateGridStats"); + defineDeviceUpdateGridStats(cudaModule, "updateGridStats"); + defineDeviceUpdateGridStats(cudaModule, "updateGridStats"); + defineDeviceUpdateGridStats(cudaModule, "updateGridStats"); + defineDeviceUpdateGridStats(cudaModule, "updateGridStats"); + defineDeviceUpdateGridStats(cudaModule, "updateGridStats"); + + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + defineDeviceIsValid(cudaModule, "isValid"); + + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + defineDeviceGridChecksum(cudaModule); + + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); + defineDeviceGridMetadata(cudaModule); #endif } diff --git a/nanovdb/nanovdb/python/PyTree.h b/nanovdb/nanovdb/python/PyTree.h index f6029317e2..5aa7c4e97f 100644 --- a/nanovdb/nanovdb/python/PyTree.h +++ b/nanovdb/nanovdb/python/PyTree.h @@ -413,6 +413,11 @@ template void defineNodeManager(nb::module_& m, const char* nam void defineNodeManagerHandle(nb::module_& m); void defineCreateNodeManager(nb::module_& m); +#ifdef NANOVDB_USE_CUDA +// Device-side NodeManagerHandle + createDeviceNodeManager, registered on the +// nanovdb.cuda submodule (defined in cuda/PyDeviceNodeManager.cu). +void defineDeviceNodeManager(nb::module_& m); +#endif // -------------------- grid.leaf_values() bulk extractor -------------------- // @@ -439,7 +444,7 @@ struct PyLeafValuesBinder; using LeafT = nanovdb::NanoLeaf; using ValueT = typename LeafT::ValueType; - cls.def("leaf_values", + cls.def("leafValues", [](nb::handle py_self) -> nb::object { auto& grid = nb::cast(py_self); const auto& tree = grid.tree(); diff --git a/nanovdb/nanovdb/python/PyValidate.h b/nanovdb/nanovdb/python/PyValidate.h new file mode 100644 index 0000000000..0efbdfa377 --- /dev/null +++ b/nanovdb/nanovdb/python/PyValidate.h @@ -0,0 +1,33 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_PYVALIDATE_HAS_BEEN_INCLUDED +#define NANOVDB_PYVALIDATE_HAS_BEEN_INCLUDED + +#include + +#include +#include + +namespace pynanovdb { + +/// @brief Raise a Python ValueError unless @a value is a finite, strictly +/// positive number. Used to validate geometric parameters (voxelSize, +/// narrow-band halfWidth) before they reach nanovdb::Map::set and the +/// grid builders, which only debug-assert positivity — release builds +/// would otherwise persist a singular / non-finite transform in the +/// grid header. +inline void requirePositiveFinite(double value, const char* fnName, const char* paramName) +{ + if (!(std::isfinite(value) && value > 0.0)) { + std::string msg(fnName); + msg += ": "; + msg += paramName; + msg += " must be a finite, strictly positive number; got "; + msg += std::to_string(value); + throw nanobind::value_error(msg.c_str()); + } +} + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/PyVoxelBlockManager.cc b/nanovdb/nanovdb/python/PyVoxelBlockManager.cc index 5e61913050..3fa5de928e 100644 --- a/nanovdb/nanovdb/python/PyVoxelBlockManager.cc +++ b/nanovdb/nanovdb/python/PyVoxelBlockManager.cc @@ -44,15 +44,15 @@ static auto dispatchLog2BlockWidth(int log2BlockWidth, F&& fn) case 9: return fn(std::integral_constant{}); default: throw nb::value_error( - "VoxelBlockManager: log2_block_width must be 6, 7, 8, or 9 " + "VoxelBlockManager: log2BlockWidth must be 6, 7, 8, or 9 " "(BlockWidth = 64, 128, 256, or 512). Larger widths are not " "bound in Python by default."); } } // PyVBMHandle wraps the C++ VoxelBlockManagerHandle and carries the -// log2_block_width the handle was built with. The C++ handle does NOT store -// log2_block_width itself, so without this wrapper the Python binding would +// log2BlockWidth the handle was built with. The C++ handle does NOT store +// log2BlockWidth itself, so without this wrapper the Python binding would // have to ask the caller every time — which the user can lie about and // trigger out-of-bounds reads of the metadata buffers. Storing it once at // build time and consulting it in every accessor closes that hole. @@ -188,12 +188,12 @@ static void defineHandle(nb::module_& toolsModule) "by this handle.") .def("reset", &PyVBMHandle::reset, "Release this handle's buffers and reset it to the empty state.") - .def_prop_ro("log2_block_width", [](const PyVBMHandle& h) { return h.log2BlockWidth; }, - "The log2_block_width this handle was built with. The jumpMap " + .def_prop_ro("log2BlockWidth", [](const PyVBMHandle& h) { return h.log2BlockWidth; }, + "The log2BlockWidth this handle was built with. The jumpMap " "and decodeBlock outputs derive their shapes from this value.") - .def_prop_ro("block_width", &PyVBMHandle::blockWidth, - "BlockWidth = 1 << log2_block_width (64, 128, 256, or 512).") - .def_prop_ro("jump_map_length", &PyVBMHandle::jumpMapLength, + .def_prop_ro("blockWidth", &PyVBMHandle::blockWidth, + "BlockWidth = 1 << log2BlockWidth (64, 128, 256, or 512).") + .def_prop_ro("jumpMapLength", &PyVBMHandle::jumpMapLength, "JumpMapLength = BlockWidth / 64 (1, 2, 4, or 8).") .def( "__bool__", @@ -224,7 +224,7 @@ static void defineHandle(nb::module_& toolsModule) "default-constructed or reset() handle. The view keeps this " "handle alive.") // jumpMap is uint64_t[blockCount * JumpMapLength]. JumpMapLength is - // determined by the log2_block_width recorded on the handle, not by + // determined by the log2BlockWidth recorded on the handle, not by // the caller — that way the returned view always covers exactly the // allocated buffer, with no risk of OOB reads. .def("jumpMap", @@ -248,22 +248,22 @@ static void defineHandle(nb::module_& toolsModule) nb::keep_alive<0, 1>(), "Return a zero-copy (blockCount, jump_map_length) uint64 NumPy " "view of the jumpMap. The shape is determined by the " - "log2_block_width the handle was built with. Returns an empty " + "log2BlockWidth the handle was built with. Returns an empty " "(0, jump_map_length) array on a default-constructed or reset() " "handle. The view keeps this handle alive.") // Decode the inverse maps for a single block of this VBM. The - // log2_block_width is taken from the handle, so the caller cannot + // log2BlockWidth is taken from the handle, so the caller cannot // request a width that doesn't match what was built. .def("decodeBlock", [](PyVBMHandle& self, nb::handle py_grid, - uint64_t block_index) -> nb::object { + uint64_t blockIndex) -> nb::object { const auto* grid = castOnIndexGrid(py_grid, "VoxelBlockManagerHandle.decodeBlock"); - if (block_index >= self.blockCount()) { + if (blockIndex >= self.blockCount()) { throw nb::index_error( - "VoxelBlockManagerHandle.decodeBlock(block_index): " - "block_index out of range [0, blockCount)."); + "VoxelBlockManagerHandle.decodeBlock(blockIndex): " + "blockIndex out of range [0, blockCount)."); } // Defensive: NanoVDB's buildVoxelBlockManager doesn't always // initialize firstLeafID for blocks where no leaf starts at @@ -275,7 +275,7 @@ static void defineHandle(nb::module_& toolsModule) // read of tree.getFirstNode<0>()[garbage]. Catch the case // and raise rather than segfault. const uint32_t firstLeafID = - self.handle.hostFirstLeafID()[block_index]; + self.handle.hostFirstLeafID()[blockIndex]; const uint32_t nLeaves = grid->tree().nodeCount(0); if (firstLeafID >= nLeaves) { throw nb::value_error( @@ -295,19 +295,19 @@ static void defineHandle(nb::module_& toolsModule) constexpr int JumpMapLength = VoxelBlockManagerBase::JumpMapLength; const uint64_t blockFirstOffset = - self.firstOffset() + block_index * BlockWidth; + self.firstOffset() + blockIndex * BlockWidth; return pyDecodeInverseMapsImpl( *grid, firstLeafID, - self.handle.hostJumpMap() + block_index * JumpMapLength, + self.handle.hostJumpMap() + blockIndex * JumpMapLength, blockFirstOffset); }); }, - "grid"_a, "block_index"_a, - "Decode the inverse maps for the block_index-th block of this " + "grid"_a, "blockIndex"_a, + "Decode the inverse maps for the blockIndex-th block of this " "VBM. Returns (leaf_index, voxel_offset) uint32 / uint16 NumPy " - "arrays of length BlockWidth = 1< PyVBMHandle { + int log2BlockWidth, + uint64_t firstOffset, + uint64_t lastOffset, + uint64_t nBlocks) -> PyVBMHandle { const auto* grid = castOnIndexGrid(py_grid, "buildVoxelBlockManager"); // The C++ implementation only NANOVDB_ASSERTs these preconditions, // which makes them no-ops in release builds. Validate them here @@ -331,25 +331,25 @@ static void defineBuild(nb::module_& toolsModule) "layout). NanoVDB grids constructed via " "tools.createOnIndexGrid satisfy this by default."); } - return dispatchLog2BlockWidth(log2_block_width, [&](auto W) { + return dispatchLog2BlockWidth(log2BlockWidth, [&](auto W) { constexpr int LBW = decltype(W)::value; using Base = VoxelBlockManagerBase; constexpr uint64_t BlockWidth = Base::BlockWidth; constexpr uint64_t JumpMapLength = Base::JumpMapLength; - // first_offset must be 1 (mod BlockWidth). The single-arg + // firstOffset must be 1 (mod BlockWidth). The single-arg // C++ helper would normalize a zero input to 1; we do the // same here so the in-place builder below sees a valid // value. Validate the nonzero case ourselves. - if (first_offset != 0 && - ((first_offset - 1) & (BlockWidth - 1)) != 0) { + if (firstOffset != 0 && + ((firstOffset - 1) & (BlockWidth - 1)) != 0) { throw nb::value_error( - "buildVoxelBlockManager: first_offset must satisfy " - "first_offset == 1 (mod BlockWidth). Pass 0 (the " + "buildVoxelBlockManager: firstOffset must satisfy " + "firstOffset == 1 (mod BlockWidth). Pass 0 (the " "default) to let the implementation use 1."); } - if (first_offset == 0) first_offset = 1; - if (last_offset == 0) last_offset = grid->activeVoxelCount(); - if (last_offset < first_offset) return PyVBMHandle(); + if (firstOffset == 0) firstOffset = 1; + if (lastOffset == 0) lastOffset = grid->activeVoxelCount(); + if (lastOffset < firstOffset) return PyVBMHandle(); // Capacity must hold at least ceil((last - first + 1) / // BlockWidth) blocks; otherwise the handle's lastOffset // would advertise more coverage than blockCount allows @@ -357,18 +357,18 @@ static void defineBuild(nb::module_& toolsModule) // below equals the ceil() above when BlockWidth is a // power of two. const uint64_t minBlocks = - (last_offset - first_offset + BlockWidth) >> LBW; - if (n_blocks != 0 && n_blocks < minBlocks) { + (lastOffset - firstOffset + BlockWidth) >> LBW; + if (nBlocks != 0 && nBlocks < minBlocks) { std::string msg( - "buildVoxelBlockManager: n_blocks must be at " - "least ceil((last_offset - first_offset + 1) / " + "buildVoxelBlockManager: nBlocks must be at " + "least ceil((lastOffset - firstOffset + 1) / " "BlockWidth) = "); msg += std::to_string(minBlocks); msg += ". Pass 0 (the default) to use the minimum " "required capacity."; throw nb::value_error(msg.c_str()); } - if (n_blocks == 0) n_blocks = minBlocks; + if (nBlocks == 0) nBlocks = minBlocks; // Allocate the metadata buffers ourselves so we can // pre-initialize firstLeafID with a sentinel value before // the in-place builder runs. The C++ allocating overload @@ -380,20 +380,20 @@ static void defineBuild(nb::module_& toolsModule) // up front, every untouched slot deterministically trips // the guard. auto firstLeafIDBuf = HostBuffer::create( - n_blocks * sizeof(uint32_t)); + nBlocks * sizeof(uint32_t)); auto jumpMapBuf = HostBuffer::create( - n_blocks * JumpMapLength * sizeof(uint64_t)); + nBlocks * JumpMapLength * sizeof(uint64_t)); const uint32_t nLeaves = grid->tree().nodeCount(0); { uint32_t* slots = static_cast( firstLeafIDBuf.data()); - for (uint64_t i = 0; i < n_blocks; ++i) { + for (uint64_t i = 0; i < nBlocks; ++i) { slots[i] = nLeaves; } } VoxelBlockManagerHandle handle( std::move(firstLeafIDBuf), std::move(jumpMapBuf), - n_blocks, first_offset, last_offset); + nBlocks, firstOffset, lastOffset); // In-place builder zeros the jumpMap itself and only // touches firstLeafID slots it actually visits. Release the // GIL around it — it's pure C++ (touches no Python objects) @@ -406,16 +406,16 @@ static void defineBuild(nb::module_& toolsModule) }); }, "grid"_a, - "log2_block_width"_a = 6, - "first_offset"_a = 0, - "last_offset"_a = 0, - "n_blocks"_a = 0, + "log2BlockWidth"_a = 6, + "firstOffset"_a = 0, + "lastOffset"_a = 0, + "nBlocks"_a = 0, "Build a host-side VoxelBlockManager from an OnIndexGrid. " - "log2_block_width selects the per-block active-voxel count " - "(6=64, 7=128, 8=256, 9=512). Pass 0 for first_offset / " - "last_offset / n_blocks to use the full grid (first active " + "log2BlockWidth selects the per-block active-voxel count " + "(6=64, 7=128, 8=256, 9=512). Pass 0 for firstOffset / " + "lastOffset / nBlocks to use the full grid (first active " "voxel through grid.activeVoxelCount(), minimum block count). " - "first_offset, if nonzero, must satisfy first_offset == 1 " + "firstOffset, if nonzero, must satisfy firstOffset == 1 " "(mod BlockWidth)."); } @@ -425,14 +425,14 @@ static void defineDecode(nb::module_& toolsModule) { toolsModule.def("decodeInverseMaps", [](nb::handle py_grid, - uint32_t first_leaf_id, + uint32_t firstLeafId, nb::ndarray, - nb::c_contig, nb::device::cpu> jump_map, - uint64_t block_first_offset, - int log2_block_width) -> nb::object { + nb::c_contig, nb::device::cpu> jumpMap, + uint64_t blockFirstOffset, + int log2BlockWidth) -> nb::object { const auto* grid = castOnIndexGrid(py_grid, "decodeInverseMaps"); - // The C++ helper indexes tree.getFirstNode<0>()[first_leaf_id] - // without a bounds check, so a stray first_leaf_id leads to an + // The C++ helper indexes tree.getFirstNode<0>()[firstLeafId] + // without a bounds check, so a stray firstLeafId leads to an // OOB read. Validate up front. (We also require isSequential(); // getFirstNode only makes sense on a sequential tree.) if (!grid->isSequential()) { @@ -441,37 +441,37 @@ static void defineDecode(nb::module_& toolsModule) "grid.isSequential()."); } const uint32_t nLeaves = grid->tree().nodeCount(0); - if (first_leaf_id >= nLeaves) { + if (firstLeafId >= nLeaves) { throw nb::index_error( - "decodeInverseMaps: first_leaf_id out of range " + "decodeInverseMaps: firstLeafId out of range " "[0, grid.tree().nodeCount(0))."); } - return dispatchLog2BlockWidth(log2_block_width, [&](auto W) { + return dispatchLog2BlockWidth(log2BlockWidth, [&](auto W) { constexpr int LBW = decltype(W)::value; constexpr int JumpMapLength = VoxelBlockManagerBase::JumpMapLength; - if (jump_map.shape(0) != JumpMapLength) { - std::string msg("decodeInverseMaps: jump_map must have " + if (jumpMap.shape(0) != JumpMapLength) { + std::string msg("decodeInverseMaps: jumpMap must have " "length JumpMapLength = "); msg += std::to_string(JumpMapLength); - msg += " for log2_block_width="; + msg += " for log2BlockWidth="; msg += std::to_string(LBW); throw nb::value_error(msg.c_str()); } return pyDecodeInverseMapsImpl( - *grid, first_leaf_id, jump_map.data(), - block_first_offset); + *grid, firstLeafId, jumpMap.data(), + blockFirstOffset); }); }, "grid"_a, - "first_leaf_id"_a, - "jump_map"_a, - "block_first_offset"_a, - "log2_block_width"_a = 6, + "firstLeafId"_a, + "jumpMap"_a, + "blockFirstOffset"_a, + "log2BlockWidth"_a = 6, "Decode the inverse maps for a single voxel block of an OnIndexGrid. " "Returns a (leaf_index, voxel_offset) tuple of fresh NumPy arrays of " - "length BlockWidth = 1< static nb::object tryCreateOnIndexGrid(nb::handle py_grid, uint32_t channels, - bool include_stats, - bool include_tiles, + bool includeStats, + bool includeTiles, int verbose) { using SrcGridT = NanoGrid; @@ -494,7 +494,7 @@ static nb::object tryCreateOnIndexGrid(nb::handle py_grid, const SrcGridT& src = nb::cast(py_grid); return nb::cast( tools::createNanoGrid( - src, channels, include_stats, include_tiles, verbose)); + src, channels, includeStats, includeTiles, verbose)); } static void defineCreateOnIndexGrid(nb::module_& toolsModule) @@ -502,31 +502,31 @@ static void defineCreateOnIndexGrid(nb::module_& toolsModule) toolsModule.def("createOnIndexGrid", [](nb::handle py_grid, uint32_t channels, - bool include_stats, - bool include_tiles, + bool includeStats, + bool includeTiles, int verbose) -> nb::object { // Try every source BuildT we accept. if (auto r = tryCreateOnIndexGrid( - py_grid, channels, include_stats, include_tiles, verbose); + py_grid, channels, includeStats, includeTiles, verbose); r.is_valid()) return r; if (auto r = tryCreateOnIndexGrid( - py_grid, channels, include_stats, include_tiles, verbose); + py_grid, channels, includeStats, includeTiles, verbose); r.is_valid()) return r; if (auto r = tryCreateOnIndexGrid( - py_grid, channels, include_stats, include_tiles, verbose); + py_grid, channels, includeStats, includeTiles, verbose); r.is_valid()) return r; if (auto r = tryCreateOnIndexGrid( - py_grid, channels, include_stats, include_tiles, verbose); + py_grid, channels, includeStats, includeTiles, verbose); r.is_valid()) return r; throw nb::type_error( "createOnIndexGrid: source grid must be a FloatGrid, " "DoubleGrid, Int32Grid, or Vec3fGrid (other source BuildTs " "are not yet bound)."); }, - "src_grid"_a, + "srcGrid"_a, "channels"_a = 0u, - "include_stats"_a = true, - "include_tiles"_a = true, + "includeStats"_a = true, + "includeTiles"_a = true, "verbose"_a = 0, "Convert a source grid into a NanoGrid " "(OnIndexGrid). Accepts FloatGrid / DoubleGrid / Int32Grid / " diff --git a/nanovdb/nanovdb/python/PyVoxelBlockManager.h b/nanovdb/nanovdb/python/PyVoxelBlockManager.h index 6fe29570c6..fa79909370 100644 --- a/nanovdb/nanovdb/python/PyVoxelBlockManager.h +++ b/nanovdb/nanovdb/python/PyVoxelBlockManager.h @@ -15,6 +15,13 @@ namespace pynanovdb { /// Python submodule (expected to be the existing nanovdb.tools). void defineVoxelBlockManagerModule(nb::module_& toolsModule); +#ifdef NANOVDB_USE_CUDA +/// @brief Bind the device VoxelBlockManagerHandle wrapper and the device +/// buildVoxelBlockManager onto the given submodule (expected to be the +/// existing nanovdb.tools.cuda). Defined in cuda/PyDeviceVoxelBlockManager.cu. +void defineDeviceVoxelBlockManager(nb::module_& cudaToolsModule); +#endif + } // namespace pynanovdb #endif diff --git a/nanovdb/nanovdb/python/__init__.py b/nanovdb/nanovdb/python/__init__.py index 9ff10b6c28..2ce52cc294 100644 --- a/nanovdb/nanovdb/python/__init__.py +++ b/nanovdb/nanovdb/python/__init__.py @@ -24,3 +24,26 @@ def get_include(): from .lib.nanovdb import * # noqa: E402,F401,F403 + +# The compiled extension only exposes a `cuda` submodule when built with CUDA +# support. Attach `nanovdb.cuda.compile_options(*extra)` there so downstream +# code can obtain the NanoVDB include flag (plus any extra flags) for runtime +# CUDA compilation. Guarded so `import nanovdb` still succeeds in non-CUDA +# builds where the `cuda` submodule does not exist. +from .lib import nanovdb as _ext # noqa: E402 + +if hasattr(_ext, "cuda"): + + def _cuda_compile_options(*extra): + """Return the NanoVDB include flag followed by any extra flags. + + Suitable for passing to a runtime CUDA compiler (e.g. NVRTC) so kernels + compile against the same NanoVDB headers as the installed wheel:: + + opts = nanovdb.cuda.compile_options("-std=c++17") + """ + return (f"-I{get_include()}",) + tuple(extra) + + _ext.cuda.compile_options = _cuda_compile_options + +del _ext diff --git a/nanovdb/nanovdb/python/cuda/PyAddBlindData.cu b/nanovdb/nanovdb/python/cuda/PyAddBlindData.cu new file mode 100644 index 0000000000..bf5d2ad088 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyAddBlindData.cu @@ -0,0 +1,88 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#include "PyAddBlindData.h" + +#include +#include + +#include +#include + +#include + +namespace nb = nanobind; +using namespace nb::literals; +// NOTE: deliberately NOT `using namespace nanovdb;`. These tools instantiate +// CUB DeviceScan, whose nvcc-generated host stub references unqualified +// `cuda::std::...`; with `nanovdb::cuda` in scope that becomes ambiguous and +// fails to compile. Fully qualify nanovdb:: instead (matches PyPointsToGrid.cu). + +namespace pynanovdb { + +template +void defineAddBlindData(nb::module_& m, const char* name) +{ + m.def( + name, + [](nanovdb::NanoGrid* dGrid, + nb::ndarray, nb::c_contig, nb::device::cuda> blindData, + nanovdb::GridBlindDataClass blindClass, + nanovdb::GridBlindDataSemantic semantics, + const std::string& dataName, + uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + const BlindDataT* d_blindData = blindData.data(); + const uint64_t valueCount = static_cast(blindData.size()); + // addBlindData copies the grid into a fresh device buffer with the + // blind data appended and launches kernels on the stream; pure + // CUDA touching no Python objects, so release the GIL. dataName is + // already a C++-owned std::string (nanobind's stl/string caster + // materialized it from the Python str), so its c_str() stays valid + // across the GIL release for the duration of this call frame. + nb::gil_scoped_release release; + return nanovdb::tools::cuda::addBlindData( + dGrid, d_blindData, valueCount, blindClass, semantics, + dataName.c_str(), nanovdb::cuda::DeviceBuffer(), s); + }, + "dGrid"_a, + "blindData"_a, + "blindClass"_a = nanovdb::GridBlindDataClass::Unknown, + "semantics"_a = nanovdb::GridBlindDataSemantic::Unknown, + "name"_a = "", + "stream"_a = 0, + "Append a flat 1-D device array of blind data to a copy of a device " + "grid and return a fresh device GridHandle with the blind data " + "attached. valueCount is taken from the array length; blindClass and " + "semantics tag the new GridBlindMetaData entry (a RuntimeError is " + "raised on an invalid combination). stream is a raw CUDA stream " + "handle (Python int; 0 = default stream)."); +} + +// Grid BuildTs x blind-data element types. addBlindData itself has no BuildT +// restriction beyond BufferTraits::hasDeviceDual (satisfied by +// DeviceBuffer); we expose the common combinations. The same overload name is +// reused so nanobind dispatches on the grid class plus the array dtype. +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +// Signed-integer blind payloads (e.g. integer labels / ids), for the same grid +// BuildTs. +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); +template void defineAddBlindData(nb::module_&, const char*); + +} // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyAddBlindData.h b/nanovdb/nanovdb/python/cuda/PyAddBlindData.h new file mode 100644 index 0000000000..91ec75bbc0 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyAddBlindData.h @@ -0,0 +1,21 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYADDBLINDDATA_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYADDBLINDDATA_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +// Bind nanovdb::tools::cuda::addBlindData for a (grid BuildT, blind-data +// element type) pair. The blind data is a flat 1-D device array; all +// instantiations register under the same Python name and are disambiguated by +// nanobind on the grid class and the blind-data ndarray dtype. +template +void defineAddBlindData(nb::module_& m, const char* name); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyCoarsenGrid.cu b/nanovdb/nanovdb/python/cuda/PyCoarsenGrid.cu new file mode 100644 index 0000000000..aa4e5200df --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyCoarsenGrid.cu @@ -0,0 +1,39 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#include "PyCoarsenGrid.h" + +#include + +#include + +namespace nb = nanobind; +using namespace nb::literals; +// NOTE: deliberately NOT `using namespace nanovdb;`. These tools instantiate +// CUB DeviceScan, whose nvcc-generated host stub references unqualified +// `cuda::std::...`; with `nanovdb::cuda` in scope that becomes ambiguous and +// fails to compile. Fully qualify nanovdb:: instead (matches PyPointsToGrid.cu). + +namespace pynanovdb { + +template void defineCoarsenGrid(nb::module_& m, const char* name) +{ + m.def( + name, + [](nanovdb::NanoGrid* dGrid, uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + // CoarsenGrid::getHandle launches kernels and synchronizes the + // stream; pure CUDA touching no Python objects, so release the GIL. + nb::gil_scoped_release release; + nanovdb::tools::cuda::CoarsenGrid coarsener(dGrid, s); + return coarsener.getHandle(); + }, + "dGrid"_a, + "stream"_a = 0, + "Topologically coarsen (2x downsample) a device OnIndex grid and return " + "a fresh device GridHandle of the coarsened grid. stream is a raw CUDA " + "stream handle (Python int; 0 = default stream)."); +} + +template void defineCoarsenGrid(nb::module_&, const char*); + +} // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyCoarsenGrid.h b/nanovdb/nanovdb/python/cuda/PyCoarsenGrid.h new file mode 100644 index 0000000000..419e264120 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyCoarsenGrid.h @@ -0,0 +1,16 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYCOARSENGRID_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYCOARSENGRID_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +template void defineCoarsenGrid(nb::module_& m, const char* name); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceBuffer.cc b/nanovdb/nanovdb/python/cuda/PyDeviceBuffer.cc index e024d63939..b3d967e364 100644 --- a/nanovdb/nanovdb/python/cuda/PyDeviceBuffer.cc +++ b/nanovdb/nanovdb/python/cuda/PyDeviceBuffer.cc @@ -4,19 +4,51 @@ #include "PyDeviceBuffer.h" +#include + #include namespace nb = nanobind; +using namespace nb::literals; using namespace nanovdb; namespace pynanovdb { void defineDeviceBuffer(nb::module_& m) { - nb::class_(m, "DeviceBuffer", - "CUDA device-side buffer used to back a DeviceGridHandle. Holds a " - "host mirror and a device pointer; deviceUpload / deviceDownload on " - "the handle move bytes between the two."); + using BufferT = nanovdb::cuda::DeviceBuffer; + defineDeviceBufferLike(m, "DeviceBuffer") + .def_static( + "from_external", + [](uint64_t size, uintptr_t gpuPtr, uintptr_t cpuPtr) { + // Wrap externally-managed host + device memory in a NON-OWNING + // DeviceBuffer (mManaged == 0). The buffer will NOT free either + // pointer on destruction, upload, or download — the caller + // retains ownership of both allocations. + if (gpuPtr == 0) + throw nb::value_error( + "from_external: gpuPtr must be a non-null device pointer."); + if (cpuPtr == 0) + throw nb::value_error( + "from_external: cpuPtr must be a non-null host pointer; the " + "externally-managed DeviceBuffer constructor requires both a " + "host and a device pointer."); + return BufferT::create(size, + reinterpret_cast(cpuPtr), + reinterpret_cast(gpuPtr)); + }, + "size"_a, + "gpuPtr"_a, + "cpuPtr"_a, + "Wrap externally-managed host and device memory in a NON-OWNING " + "DeviceBuffer. size is the byte size of both allocations; gpuPtr " + "and cpuPtr are raw pointers (Python ints). The returned buffer " + "does NOT take ownership: it will never free either pointer, so " + "the caller must keep both allocations alive for the buffer's " + "lifetime. The device pointer is associated with the current CUDA " + "device.") + .def("recordUse", &recordUseChecked, "stream"_a, "device"_a = -1, + kRecordUseDoc); } } // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceBuffer.h b/nanovdb/nanovdb/python/cuda/PyDeviceBuffer.h index 87a081f638..c727c0b008 100644 --- a/nanovdb/nanovdb/python/cuda/PyDeviceBuffer.h +++ b/nanovdb/nanovdb/python/cuda/PyDeviceBuffer.h @@ -5,11 +5,217 @@ #include +#ifdef NANOVDB_USE_CUDA +#include + +#include +#include + +#include + +#include +#endif + namespace nb = nanobind; namespace pynanovdb { #ifdef NANOVDB_USE_CUDA + +/// @brief Bounds-checked forwarder for DeviceBuffer::recordUse, shared by the +/// DeviceBuffer and DeviceGridHandle bindings. device == -1 selects the +/// current CUDA device. recordUse itself indexes per-device tracking +/// state unchecked, so validate here where we can raise a Python +/// exception instead. +inline void recordUseChecked(nanovdb::cuda::DeviceBuffer& buf, uintptr_t stream, int device) +{ + int count = 0; + cudaCheck(cudaGetDeviceCount(&count)); + if (device < 0) cudaCheck(cudaGetDevice(&device)); + if (device >= count) { + const std::string msg = "recordUse: device id " + std::to_string(device) + + " out of range [0, " + std::to_string(count) + ")."; + throw nb::index_error(msg.c_str()); + } + nb::gil_scoped_release release; + buf.recordUse(device, reinterpret_cast(stream)); +} + +/// @brief Order the buffer's tracked prior uses (uploads / downloads / +/// recordUse'd kernels) before work subsequently issued on @a stream, +/// when the buffer type exposes use tracking. No-op overload for buffer +/// types without it (UnifiedBuffer). Call with a trailing 0 so the +/// tracking overload is preferred when available. +/// @note DeviceBuffer::orderAfterPriorUses is still private upstream, so the +/// SFINAE currently selects the no-op for DeviceBuffer too — the +/// CAI/DLPack exports below only gain real event ordering once the +/// upstream change making it public (and chaining recordUse) lands and +/// is merged in. The call sites are written against the final +/// semantics so no binding change is needed at that point. +template +inline auto orderPriorUsesBefore(const BufferT& buf, cudaStream_t stream, int) + -> decltype(buf.orderAfterPriorUses(0, stream)) +{ + int device = 0; + cudaCheck(cudaGetDevice(&device)); + buf.orderAfterPriorUses(device, stream); +} +template +inline void orderPriorUsesBefore(const BufferT&, cudaStream_t, long) {} + +/// @brief Resolve the DLPack-protocol `stream` argument to the cudaStream_t +/// the export must be ordered on. Per the protocol: None/0/1 = the +/// legacy default stream, 2 = the per-thread default stream, -1 = the +/// consumer does its own synchronization (no ordering requested; return +/// false), any other integer = a raw cudaStream_t handle. +inline bool resolveDlpackStream(nb::handle stream, cudaStream_t& out) +{ + if (stream.is_none()) { + out = nullptr; // legacy default stream + return true; + } + const long long v = nb::cast(stream); + if (v == -1) return false; + if (v == 2) { + out = cudaStreamPerThread; + return true; + } + if (v == 0 || v == 1) { + out = nullptr; // legacy default stream + return true; + } + out = reinterpret_cast(static_cast(v)); + return true; +} + +/// @brief Docstring shared by the DeviceBuffer and DeviceGridHandle recordUse +/// bindings (the two forward to the same underlying buffer method). +inline constexpr char kRecordUseDoc[] = + "Record that this buffer's device data was just used on `stream` (a raw " + "cudaStream_t as a Python int, e.g. cupy.cuda.Stream.ptr), so the device " + "free issued when the buffer is cleared or destroyed is ordered after " + "that work. Uploads/downloads record themselves automatically; call this " + "after enqueuing your own kernels or copies against device_ptr() / " + "__cuda_array_interface__ / __dlpack__ on a non-blocking stream — " + "without it such work is only safe if you synchronize before dropping " + "the buffer. device selects which device's buffer was used (-1 = the " + "current CUDA device). No-op on non-owning (from_external) buffers, " + "which never free their pointers."; + +/// @brief Bind the device-interop surface (CUDA Array Interface / DLPack, raw +/// device/host pointers, streams) onto a device-buffer-like class. +/// +/// @details This is the Phase-B interop hook. It is templated only on the +/// duck-typed buffer surface (size()/data()/deviceData()) so it can +/// be reused for any DeviceBuffer-like type; it must NOT reference +/// DeviceBuffer-specific members. It exposes the whole device buffer +/// as 1-D bytes through the CUDA Array Interface and DLPack so it can +/// be consumed zero-copy by CuPy / PyTorch / Numba, plus the raw +/// device/host pointers as Python ints. +template +void addDeviceInterop(nb::class_& cls) +{ + cls.def("size", &BufferT::size, "Total number of bytes managed by this buffer."); + + cls.def( + "device_ptr", + [](const BufferT& buf) { + return reinterpret_cast(buf.deviceData()); + }, + "Raw device pointer to the current device's buffer as a Python int " + "(0 if no device allocation exists yet). Work you enqueue against " + "this pointer on a non-blocking stream is invisible to the buffer's " + "lifetime tracking: record it afterwards where the buffer supports " + "it (DeviceBuffer.recordUse / DeviceGridHandle.recordUse), or " + "synchronize before the buffer is destroyed."); + + cls.def( + "host_ptr", + [](const BufferT& buf) { + return reinterpret_cast(buf.data()); + }, + "Raw host pointer to the buffer's host mirror as a Python int " + "(0 if no host allocation exists)."); + + cls.def_prop_ro( + "__cuda_array_interface__", + [](const BufferT& buf) { + // Hand-rolled CUDA Array Interface (version 3) describing the whole + // device buffer as a 1-D contiguous uint8 array. stream=1 selects + // the legacy default stream per the CAI v3 spec — make that claim + // true by ordering the legacy default stream after the buffer's + // tracked prior uses (async uploads, recordUse'd kernels), so a + // consumer that synchronizes on it per the spec sees complete data. + orderPriorUsesBefore(buf, cudaStream_t(0), 0); + nb::dict iface; + iface["shape"] = nb::make_tuple(buf.size()); + iface["typestr"] = "|u1"; + iface["data"] = + nb::make_tuple(reinterpret_cast(buf.deviceData()), false); + iface["version"] = 3; + iface["strides"] = nb::none(); + iface["stream"] = 1; + return iface; + }, + "CUDA Array Interface (v3) view of the whole device buffer as 1-D " + "uint8. Lets CuPy / Numba / PyTorch consume the device bytes " + "zero-copy. Returns a null data pointer until the buffer is populated " + "on the device."); + + cls.def( + "__dlpack_device__", + [](const BufferT&) { + int device = 0; + cudaGetDevice(&device); + // 2 == kDLCUDA. + return nb::make_tuple(2, device); + }, + "DLPack device tuple (kDLCUDA, device_id) for the device buffer."); + + cls.def( + "__dlpack__", + [](nb::handle self, nb::handle stream) { + const BufferT& buf = nb::cast(self); + // Honor the consumer-provided stream per the DLPack protocol: + // order it after the buffer's tracked prior uses so the consumer + // cannot read a partially-written buffer (e.g. after an async + // deviceUpload on another stream). + cudaStream_t consumer; + if (resolveDlpackStream(stream, consumer)) + orderPriorUsesBefore(buf, consumer, 0); + size_t shape[1] = {static_cast(buf.size())}; + // Delegate the capsule construction to nanobind: build a device + // ndarray view parented to this buffer (keep_alive via owner) and + // forward to its own __dlpack__ producer. + nb::ndarray> arr( + buf.deviceData(), 1, shape, self); + // nb::cast of a no-framework device ndarray IS the "dltensor" + // PyCapsule (nanobind ndarray_export), which is exactly what + // __dlpack__ must return — so return it directly (do NOT call + // .attr("__dlpack__") on it; a capsule has no such attribute). + // ndarray_inc_ref + owner=self keep the device memory alive. + return nb::cast(arr, nb::rv_policy::reference); + }, + nb::arg("stream") = nb::none(), + "DLPack capsule exporting the whole device buffer as 1-D uint8. " + "Delegates to a nanobind device ndarray view parented to this " + "buffer."); +} + +/// @brief Create an @c nb::class_ for a device-buffer-like type, attach the +/// shared device-interop surface via addDeviceInterop, and return it so +/// callers may chain additional @c .def() bindings. +template +nb::class_ defineDeviceBufferLike(nb::module_& m, const char* name) +{ + nb::class_ cls(m, name, + "CUDA device-side buffer used to back a DeviceGridHandle. Holds a " + "host mirror and a device pointer; deviceUpload / deviceDownload on " + "the handle move bytes between the two."); + addDeviceInterop(cls); + return cls; +} + void defineDeviceBuffer(nb::module_& m); #endif diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceGridChecksum.cu b/nanovdb/nanovdb/python/cuda/PyDeviceGridChecksum.cu new file mode 100644 index 0000000000..bb64e87da0 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDeviceGridChecksum.cu @@ -0,0 +1,167 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#include "PyDeviceGridChecksum.h" + +#include + +// Pull in cuda/GridHandle.cuh first: it includes GridChecksum.cuh and then +// uses nanovdb::tools::cuda::updateChecksum, so it must be parsed before any +// translation unit sets GridChecksum.cuh's include guard (otherwise the guard +// would skip GridChecksum.cuh's body and leave updateChecksum undeclared when +// GridHandle.cuh is parsed). Matches the unittest's include ordering. +#include +#include + +namespace nb = nanobind; +using namespace nb::literals; +// NOTE: deliberately NOT `using namespace nanovdb;`. These tools instantiate +// CUB DeviceScan, whose nvcc-generated host stub references unqualified +// `cuda::std::...`; with `nanovdb::cuda` in scope that becomes ambiguous and +// fails to compile. Fully qualify nanovdb:: instead (matches PyPointsToGrid.cu). + +namespace pynanovdb { + +// All three entries take the device NanoGrid* and use the GridData* +// device overloads in GridChecksum.cuh, which copy just the header host-side +// and run CRC32 on device — they never dereference the device pointer on the +// host, so reinterpreting the typed device grid pointer to GridData* is safe. +template +void defineDeviceGridChecksum(nb::module_& m) +{ + m.def( + "evalChecksum", + [](const nanovdb::NanoGrid* dGrid, nanovdb::CheckMode mode, + uintptr_t stream) -> nanovdb::Checksum { + cudaStream_t s = reinterpret_cast(stream); + const nanovdb::GridData* d_gridData = + reinterpret_cast(dGrid); + // Pure CUDA (CRC32 on device, header copied D2H); release the GIL. + nb::gil_scoped_release release; + return nanovdb::tools::cuda::evalChecksum(d_gridData, mode, s); + }, + "dGrid"_a, + "mode"_a = nanovdb::CheckMode::Default, + "stream"_a = 0, + "Compute and return the Checksum of the device grid for the given " + "CheckMode without modifying it. stream is a raw CUDA stream handle " + "(Python int; 0 = default stream)."); + + m.def( + "validateChecksum", + [](const nanovdb::NanoGrid* dGrid, nanovdb::CheckMode mode, + uintptr_t stream) -> bool { + cudaStream_t s = reinterpret_cast(stream); + const nanovdb::GridData* d_gridData = + reinterpret_cast(dGrid); + nb::gil_scoped_release release; + return nanovdb::tools::cuda::validateChecksum(d_gridData, mode, s); + }, + "dGrid"_a, + "mode"_a = nanovdb::CheckMode::Default, + "stream"_a = 0, + "Return True iff the device grid's stored checksum matches a freshly " + "computed one for the given CheckMode. A grid with an empty stored " + "checksum is considered valid. stream is a raw CUDA stream handle " + "(Python int; 0 = default stream)."); + + m.def( + "updateChecksum", + [](nanovdb::NanoGrid* dGrid, nanovdb::CheckMode mode, + uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + nanovdb::GridData* d_gridData = + reinterpret_cast(dGrid); + nb::gil_scoped_release release; + nanovdb::tools::cuda::updateChecksum(d_gridData, mode, s); + }, + "dGrid"_a, + "mode"_a = nanovdb::CheckMode::Default, + "stream"_a = 0, + "Recompute and write the checksum of the device grid in place using " + "the given CheckMode (returns None). stream is a raw CUDA stream " + "handle (Python int; 0 = default stream)."); +} + +// One-thread kernel that overwrites the grid's GridClass field in place. Grid +// publicly derives GridData, so mGridClass is reachable on the device pointer. +template +__global__ void setGridClassKernel(nanovdb::NanoGrid* dGrid, + nanovdb::GridClass gridClass) +{ + if (blockIdx.x == 0 && threadIdx.x == 0) dGrid->mGridClass = gridClass; +} + +// Device-side mutable grid-header metadata (the read-only counterparts already +// exist as getters on the grid objects) +template +void defineDeviceGridMetadata(nb::module_& m) +{ + m.def( + "setGridClass", + [](nanovdb::NanoGrid* dGrid, nanovdb::GridClass gridClass, + uintptr_t stream) { + if (!dGrid) throw nb::value_error("setGridClass: dGrid is None."); + cudaStream_t s = reinterpret_cast(stream); + nb::gil_scoped_release release; + setGridClassKernel<<<1, 1, 0, s>>>(dGrid, gridClass); + // Refresh the checksum (preserving its existing mode) so that + // checksum-validating readers still accept the grid after the + // class field changed; a no-op for grids with checksum disabled. + nanovdb::tools::cuda::updateChecksum( + reinterpret_cast(dGrid), s); + cudaStreamSynchronize(s); + }, + "dGrid"_a, + "gridClass"_a, + "stream"_a = 0, + "Set the device grid's GridClass (e.g. GridClass.LevelSet) in place and " + "refresh its checksum, preserving the checksum mode (returns None). " + "stream is a raw CUDA stream handle (Python int; 0 = default stream)."); +} + +// No BuildT restriction on the GridChecksum device entries; instantiate for +// the same set the host checksum dispatch covers via callNanoGrid. +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); +template void defineDeviceGridChecksum(nb::module_&); + +// Same BuildT set as the checksum entries above. +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); +template void defineDeviceGridMetadata(nb::module_&); + +} // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceGridChecksum.h b/nanovdb/nanovdb/python/cuda/PyDeviceGridChecksum.h new file mode 100644 index 0000000000..633b110cb6 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDeviceGridChecksum.h @@ -0,0 +1,27 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYDEVICEGRIDCHECKSUM_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYDEVICEGRIDCHECKSUM_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +// Bind the device checksum entry points for one grid BuildT. Each registers +// nanovdb.tools.cuda.evalChecksum / validateChecksum / updateChecksum as an +// overload taking a (device) NanoGrid* reinterpreted as GridData*. +// nanobind disambiguates the overloads on the device grid class. +template +void defineDeviceGridChecksum(nb::module_& m); + +// Bind mutable grid-header metadata setters for one grid BuildT. Currently +// registers nanovdb.tools.cuda.setGridClass(dGrid, gridClass, stream), which +// overwrites the device grid's GridClass in place and refreshes its checksum. +template +void defineDeviceGridMetadata(nb::module_& m); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceGridHandle.cu b/nanovdb/nanovdb/python/cuda/PyDeviceGridHandle.cu index 0214474420..a4bf3fa73f 100644 --- a/nanovdb/nanovdb/python/cuda/PyDeviceGridHandle.cu +++ b/nanovdb/nanovdb/python/cuda/PyDeviceGridHandle.cu @@ -2,9 +2,14 @@ // SPDX-License-Identifier: Apache-2.0 #ifdef NANOVDB_USE_CUDA -#include "../PyGridHandle.h" +#include "PyGridHandle.h" +#include "PyDeviceBuffer.h" // for recordUseChecked / kRecordUseDoc #include +#include + +#include + #include #include @@ -50,7 +55,7 @@ static nb::object pyDeviceGrid(nb::handle py_handle, uint32_t n) return grid ? nb::cast(grid, nb::rv_policy::reference, py_handle) \ : nb::none(); \ } -#include "../BuildTypes.def" +#include "BuildTypes.def" default: return nb::none(); } @@ -63,14 +68,14 @@ void defineDeviceGridHandle(nb::module_& m) .def( "__init__", [](GridHandle& handle, - nb::ndarray, nb::device::cpu> cpu_t, - nb::ndarray, nb::device::cuda> cuda_t) { - assert(cpu_t.size() == cuda_t.size()); - BufferT buffer(cpu_t.size() * sizeof(uint32_t), cpu_t.data(), cuda_t.data()); + nb::ndarray, nb::device::cpu> cpuT, + nb::ndarray, nb::device::cuda> cudaT) { + assert(cpuT.size() == cudaT.size()); + BufferT buffer(cpuT.size() * sizeof(uint32_t), cpuT.data(), cudaT.data()); new (&handle) GridHandle(std::move(buffer)); }, - "cpu_t"_a.noconvert(), - "cuda_t"_a.noconvert(), + "cpuT"_a.noconvert(), + "cudaT"_a.noconvert(), "Construct a DeviceGridHandle that wraps an existing pair of " "host and device uint32 arrays of equal length.") .def("deviceGrid", &pyDeviceGrid, "n"_a = 0, @@ -80,13 +85,115 @@ void defineDeviceGridHandle(nb::module_& m) "Python or the device copy has not been uploaded yet. The " "returned grid keeps this handle alive.") .def( - "deviceUpload", [](GridHandle& handle, bool sync) { handle.deviceUpload(nullptr, sync); }, "sync"_a = true, - "Copy the host-side buffer to the device. If sync is True the " - "call blocks until the transfer completes.") + "deviceUpload", + [](GridHandle& handle, uintptr_t stream, bool sync) { + cudaStream_t s = reinterpret_cast(stream); + // Use the current-device overload (void*, bool) — NOT the + // (int device, void*, bool) form — so the targeted device + // matches deviceData()/CAI/__dlpack__ (which use cudaGetDevice). + handle.deviceUpload(reinterpret_cast(s), sync); + }, + "stream"_a = 0, "sync"_a = true, + "Copy the host-side buffer to the device. stream is a raw CUDA " + "stream handle (Python int; 0 = default stream). If sync is True " + "the call blocks until the transfer completes.") + .def( + "deviceDownload", + [](GridHandle& handle, uintptr_t stream, bool sync) { + cudaStream_t s = reinterpret_cast(stream); + // Current-device overload, matching deviceData() (see deviceUpload). + handle.deviceDownload(reinterpret_cast(s), sync); + }, + "stream"_a = 0, "sync"_a = true, + "Copy the device-side buffer back to the host. stream is a raw " + "CUDA stream handle (Python int; 0 = default stream). If sync is " + "True the call blocks until the transfer completes.") + .def( + "device_ptr", + [](GridHandle& handle) { + return reinterpret_cast(handle.buffer().deviceData()); + }, + "Raw device pointer to the base of the whole device buffer as a " + "Python int (0 if the handle has not been uploaded to the device " + "yet). Work you enqueue against this pointer on a non-blocking " + "stream is invisible to the buffer's lifetime tracking: call " + "recordUse(stream) afterwards, or synchronize before the handle " + "is destroyed.") + .def( + "recordUse", + [](GridHandle& handle, uintptr_t stream, int device) { + recordUseChecked(handle.buffer(), stream, device); + }, + "stream"_a, "device"_a = -1, + kRecordUseDoc) + .def_prop_ro( + "__cuda_array_interface__", + [](GridHandle& handle) { + // CUDA Array Interface (v3) over the whole device buffer as a + // 1-D contiguous uint8 array. stream=1 selects the legacy + // default stream per the CAI v3 spec — make that claim true by + // ordering the legacy default stream after the buffer's tracked + // prior uses (async uploads, recordUse'd kernels). + orderPriorUsesBefore(handle.buffer(), cudaStream_t(0), 0); + nb::dict iface; + iface["shape"] = nb::make_tuple(handle.buffer().size()); + iface["typestr"] = "|u1"; + iface["data"] = nb::make_tuple( + reinterpret_cast(handle.buffer().deviceData()), false); + iface["version"] = 3; + iface["strides"] = nb::none(); + iface["stream"] = 1; + return iface; + }, + "CUDA Array Interface (v3) view of the whole device buffer as 1-D " + "uint8 — lets CuPy / Numba / PyTorch consume the serialized grid " + "bytes zero-copy. Returns a null data pointer until deviceUpload. " + "After enqueuing work on this view from a non-blocking stream, " + "call recordUse(stream) so the buffer's device free is ordered " + "after it.") .def( - "deviceDownload", [](GridHandle& handle, bool sync) { handle.deviceDownload(nullptr, sync); }, "sync"_a = true, - "Copy the device-side buffer back to the host. If sync is True " - "the call blocks until the transfer completes."); + "__dlpack_device__", + [](GridHandle&) { + int device = 0; + cudaGetDevice(&device); + return nb::make_tuple(2, device); // 2 == kDLCUDA + }, + "DLPack device tuple (kDLCUDA, device_id) for the device buffer.") + .def( + "__dlpack__", + [](nb::handle self, nb::handle stream) { + auto& handle = nb::cast&>(self); + // Honor the consumer-provided stream per the DLPack protocol: + // order it after the buffer's tracked prior uses so the + // consumer cannot read a partially-uploaded buffer. + cudaStream_t consumer; + if (resolveDlpackStream(stream, consumer)) + orderPriorUsesBefore(handle.buffer(), consumer, 0); + size_t shape[1] = {static_cast(handle.buffer().size())}; + nb::ndarray> arr( + handle.buffer().deviceData(), 1, shape, self); + // nb::cast of a no-framework device ndarray IS the "dltensor" + // capsule (what __dlpack__ must return); return it directly. + return nb::cast(arr, nb::rv_policy::reference); + }, + "stream"_a = nb::none(), + "DLPack capsule exporting the whole device buffer as 1-D uint8, " + "parented to this handle.") + .def_static( + "from_buffer", + [](BufferT& buffer) { + // Consumes (moves from) the buffer; the GridHandle ctor peeks + // the GridData header (host side if present, else a D2H copy of + // the device side) and throws std::runtime_error if it is not a + // valid grid. + return GridHandle(std::move(buffer)); + }, + "buffer"_a, + "Build a DeviceGridHandle that takes ownership of a DeviceBuffer. " + "The buffer is MOVED FROM (left empty), and its first GridData " + "header is validated — a RuntimeError is raised if it does not hold " + "a valid NanoVDB grid. Pair with DeviceBuffer.from_external to wrap " + "externally-managed device/host memory zero-copy."); // NOTE: defineGridHandleUtilities intentionally NOT called for // DeviceBuffer. Registering nanovdb.splitGrids / nanovdb.mergeGrids as a // second overload taking a DeviceGridHandle list conflicts with the host diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceGridStats.cu b/nanovdb/nanovdb/python/cuda/PyDeviceGridStats.cu new file mode 100644 index 0000000000..aaa577d46f --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDeviceGridStats.cu @@ -0,0 +1,66 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#include "PyDeviceGridStats.h" + +#include + +// GridStats.cuh uses NodeManager and (via GridHandle) the device +// checksum path but is not self-contained for either. Pull in +// cuda/GridHandle.cuh first: it includes cuda/NodeManager.cuh (for +// NodeManager) and GridChecksum.cuh, and must be parsed before any TU sets +// GridChecksum.cuh's include guard so tools::cuda::updateChecksum is declared +// when GridHandle.cuh is parsed. Matches the unittest's include ordering. +#include +#include + +namespace nb = nanobind; +using namespace nb::literals; +// NOTE: deliberately NOT `using namespace nanovdb;`. These tools instantiate +// CUB DeviceScan, whose nvcc-generated host stub references unqualified +// `cuda::std::...`; with `nanovdb::cuda` in scope that becomes ambiguous and +// fails to compile. Fully qualify nanovdb:: instead (matches PyPointsToGrid.cu). + +namespace pynanovdb { + +template +void defineDeviceUpdateGridStats(nb::module_& m, const char* name) +{ + m.def( + name, + [](nanovdb::NanoGrid* dGrid, nanovdb::tools::StatsMode mode, uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + // updateGridStats launches kernels and synchronizes the stream; + // pure CUDA touching no Python objects, so release the GIL. The + // operation mutates the device grid in place. + nb::gil_scoped_release release; + nanovdb::tools::cuda::updateGridStats(dGrid, mode, s); + }, + "dGrid"_a, + "mode"_a = nanovdb::tools::StatsMode::Default, + "stream"_a = 0, + "Recompute and write per-node statistics into the given device grid " + "in place (returns None). Does NOT recompute the grid checksum — call " + "updateChecksum afterward if the checksum must stay valid. stream is a " + "raw CUDA stream handle (Python int; 0 = default stream)."); +} + +// Scalar + vector + bool BuildTs. updateGridStats's MinMax / All branches +// instantiate Extrema / Stats, which are only meaningful for +// arithmetic value types, so the quantized / index / mask special BuildTs are +// intentionally NOT instantiated here (they would fail GridStats's ValueT +// static_assert / lack a usable Stats specialization). bool routes to the +// NoopStats path internally regardless of mode. +template void defineDeviceUpdateGridStats(nb::module_&, const char*); +template void defineDeviceUpdateGridStats(nb::module_&, const char*); +template void defineDeviceUpdateGridStats(nb::module_&, const char*); +template void defineDeviceUpdateGridStats(nb::module_&, const char*); +template void defineDeviceUpdateGridStats(nb::module_&, const char*); +template void defineDeviceUpdateGridStats(nb::module_&, const char*); +template void defineDeviceUpdateGridStats(nb::module_&, const char*); +template void defineDeviceUpdateGridStats(nb::module_&, const char*); +template void defineDeviceUpdateGridStats(nb::module_&, const char*); +template void defineDeviceUpdateGridStats(nb::module_&, const char*); +template void defineDeviceUpdateGridStats(nb::module_&, const char*); +template void defineDeviceUpdateGridStats(nb::module_&, const char*); + +} // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceGridStats.h b/nanovdb/nanovdb/python/cuda/PyDeviceGridStats.h new file mode 100644 index 0000000000..bf26314a14 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDeviceGridStats.h @@ -0,0 +1,20 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYDEVICEGRIDSTATS_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYDEVICEGRIDSTATS_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +// Bind nanovdb::tools::cuda::updateGridStats for one grid BuildT. All +// instantiations register under the same Python name ("updateGridStats") and +// are disambiguated by nanobind on the (device) grid class. +template +void defineDeviceUpdateGridStats(nb::module_& m, const char* name); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceGridValidator.cu b/nanovdb/nanovdb/python/cuda/PyDeviceGridValidator.cu new file mode 100644 index 0000000000..8a734b6926 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDeviceGridValidator.cu @@ -0,0 +1,73 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#include "PyDeviceGridValidator.h" + +#include + +// Pull in cuda/GridHandle.cuh first: it includes GridChecksum.cuh (which +// GridValidator.cuh also pulls in) and then uses +// nanovdb::tools::cuda::updateChecksum, so it must be parsed before any +// translation unit sets GridChecksum.cuh's include guard. Matches the +// unittest's include ordering. +#include +#include + +namespace nb = nanobind; +using namespace nb::literals; +// NOTE: deliberately NOT `using namespace nanovdb;`. These tools instantiate +// CUB DeviceScan, whose nvcc-generated host stub references unqualified +// `cuda::std::...`; with `nanovdb::cuda` in scope that becomes ambiguous and +// fails to compile. Fully qualify nanovdb:: instead (matches PyPointsToGrid.cu). + +namespace pynanovdb { + +template +void defineDeviceIsValid(nb::module_& m, const char* name) +{ + m.def( + name, + [](const nanovdb::NanoGrid* dGrid, nanovdb::CheckMode mode, + bool verbose, uintptr_t stream) -> bool { + cudaStream_t s = reinterpret_cast(stream); + // isValid runs structural checks in a device kernel plus a device + // checksum validation; pure CUDA touching no Python objects (the + // optional verbose diagnostic goes to std::cerr), so release the + // GIL. + nb::gil_scoped_release release; + return nanovdb::tools::cuda::isValid(dGrid, mode, verbose, s); + }, + "dGrid"_a, + "mode"_a = nanovdb::CheckMode::Default, + "verbose"_a = false, + "stream"_a = 0, + "Return True iff the device grid passes structural validation for the " + "given CheckMode AND its stored checksum matches a freshly computed " + "one. If verbose, the first failure is printed to stderr. stream is a " + "raw CUDA stream handle (Python int; 0 = default stream)."); +} + +// Instantiate for the same BuildT set the host isValid covers via +// callNanoGrid (NanoVDB.h). tools::checkGrid compiles for all of these. +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); +template void defineDeviceIsValid(nb::module_&, const char*); + +} // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceGridValidator.h b/nanovdb/nanovdb/python/cuda/PyDeviceGridValidator.h new file mode 100644 index 0000000000..28d01cf8b1 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDeviceGridValidator.h @@ -0,0 +1,20 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYDEVICEGRIDVALIDATOR_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYDEVICEGRIDVALIDATOR_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +// Bind nanovdb::tools::cuda::isValid for one grid BuildT. All instantiations +// register under the same Python name ("isValid") and are disambiguated by +// nanobind on the (device) grid class. +template +void defineDeviceIsValid(nb::module_& m, const char* name); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceMesh.cu b/nanovdb/nanovdb/python/cuda/PyDeviceMesh.cu new file mode 100644 index 0000000000..8aa9d8c59d --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDeviceMesh.cu @@ -0,0 +1,79 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifdef NANOVDB_USE_CUDA + +#include "PyDeviceMesh.h" + +#include + +#include + +#include + +namespace nb = nanobind; +using namespace nb::literals; + +namespace pynanovdb { + +void defineDeviceMesh(nb::module_& m) +{ + using DeviceNode = nanovdb::cuda::DeviceNode; + using DeviceMesh = nanovdb::cuda::DeviceMesh; + + nb::class_(m, "DeviceNode", + "A device id paired with a CUDA stream created on that device.") + .def_ro("id", &DeviceNode::id, + "CUDA device id this node refers to (-1 if unset).") + .def_prop_ro( + "stream", + [](const DeviceNode& node) { + return reinterpret_cast(node.stream); + }, + "Raw CUDA stream handle for this device as a Python int " + "(0 = default stream)."); + + nb::class_(m, "DeviceMesh", + "Multi-GPU context: enumerates every CUDA device on the host, creates " + "a stream per device, and caches peer-to-peer connectivity. Pass it to " + "nanovdb.tools.cuda.DistributedPointsToGrid. Move-only / not copyable.") + .def( + "__init__", + [](DeviceMesh* self) { + // The ctor touches every device (stream creation, P2P probe), + // so release the GIL while it runs. + nb::gil_scoped_release release; + new (self) DeviceMesh(); + }, + "Construct a DeviceMesh spanning every CUDA device on the host. " + "Each device must support managed memory.") + .def( + "deviceCount", + [](const DeviceMesh& mesh) { return static_cast(mesh.deviceCount()); }, + "Number of devices in this mesh.") + .def( + "__len__", + [](const DeviceMesh& mesh) { return static_cast(mesh.deviceCount()); }, + "Number of devices in this mesh (same as deviceCount()).") + .def( + "__getitem__", + [](const DeviceMesh& mesh, int deviceId) -> const DeviceNode& { + if (deviceId < 0 || static_cast(deviceId) >= mesh.deviceCount()) + throw nb::index_error("DeviceMesh index out of range [0, deviceCount())."); + return mesh[deviceId]; + }, + "deviceId"_a, + nb::rv_policy::reference_internal, + "Return the DeviceNode (id + stream) for the given device index.") + .def( + "canAccessPeer", + [](const DeviceMesh& mesh, int deviceId, int peerId) { + return mesh.canAccessPeer(deviceId, peerId); + }, + "deviceId"_a, "peerId"_a, + "True iff `deviceId` can directly access memory on `peerId` " + "(peer-to-peer support)."); +} + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceMesh.h b/nanovdb/nanovdb/python/cuda/PyDeviceMesh.h new file mode 100644 index 0000000000..05d0e6928a --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDeviceMesh.h @@ -0,0 +1,23 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYDEVICEMESH_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYDEVICEMESH_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +#ifdef NANOVDB_USE_CUDA +/// @brief Register nanovdb::cuda::DeviceNode and nanovdb::cuda::DeviceMesh on +/// the nanovdb.cuda submodule. DeviceMesh enumerates every CUDA device +/// on the host, creates a per-device stream, and caches P2P +/// connectivity; it is the multi-GPU context object consumed by +/// nanovdb.tools.cuda.DistributedPointsToGrid. +void defineDeviceMesh(nb::module_& m); +#endif + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceNodeManager.cu b/nanovdb/nanovdb/python/cuda/PyDeviceNodeManager.cu new file mode 100644 index 0000000000..495a87e903 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDeviceNodeManager.cu @@ -0,0 +1,172 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifdef NANOVDB_USE_CUDA + +#include "PyTree.h" + +#include + +#include + +#include +#include +#include +#include + +namespace nb = nanobind; +using namespace nb::literals; +using namespace nanovdb; + +namespace pynanovdb { + +// Device-side polymorphic mgr() — same dispatch shape as pyNodeMgr in +// PyTree.cc, but probes the DEVICE-resident NodeManager pointer +// (handle.deviceMgr()) instead of the host one. The per-BuildT +// NanoVDB NodeManager classes registered on the root module are reused +// as-is; nanobind does not distinguish a host vs device `this` pointer — the +// returned object is a NodeManager whose underlying address lives in device +// memory, so its accessors (leaf(i)/lower(i)/upper(i)/...) must only be used +// from CUDA kernels, never dereferenced on the host. +static nb::object pyDeviceNodeMgr(nb::handle py_self) +{ + using BufferT = nanovdb::cuda::DeviceBuffer; + using HandleT = NodeManagerHandle; + auto& handle = nb::cast(py_self); + // cuda::createNodeManager builds a DEVICE-only NodeManager: the buffer's + // host mirror (handle.data()) stays null while the device side is + // populated. Gate on size() (the allocated NodeManagerData byte count), + // NOT data(), so the device handle is observable. + if (handle.size() == 0) return nb::none(); + // The stored gridType is private; deviceMgr() returns NULL on a + // type mismatch, so iterate by BuildT (first non-null wins). The X-macro + // produces one case per bound BuildT. +#define NANOVDB_PY_FOR_EACH_SCALAR_BUILDT(T, Suffix, GridTypeEnum) \ + if (auto* m = handle.template deviceMgr()) { \ + return nb::cast(m, nb::rv_policy::reference, py_self); \ + } +#define NANOVDB_PY_FOR_EACH_VECTOR_BUILDT(T, Suffix, AccessorName, GridTypeEnum) \ + if (auto* m = handle.template deviceMgr()) { \ + return nb::cast(m, nb::rv_policy::reference, py_self); \ + } +#define NANOVDB_PY_FOR_EACH_POINT_BUILDT(T, Suffix, GridTypeEnum) \ + if (auto* m = handle.template deviceMgr()) { \ + return nb::cast(m, nb::rv_policy::reference, py_self); \ + } +#define NANOVDB_PY_FOR_EACH_READONLY_BUILDT(T, Suffix, GridTypeEnum) \ + if (auto* m = handle.template deviceMgr()) { \ + return nb::cast(m, nb::rv_policy::reference, py_self); \ + } +#include "BuildTypes.def" + return nb::none(); +} + +static void defineDeviceNodeManagerHandle(nb::module_& m) +{ + using BufferT = nanovdb::cuda::DeviceBuffer; + using HandleT = NodeManagerHandle; + // Distinct name from the host "NodeManagerHandle" because both live in the + // same nanobind type registry; a clashing name would collide. The device + // variant lives on the nanovdb.cuda submodule. + // + // No deviceUpload/deviceDownload here: cuda::createNodeManager builds the + // NodeManager directly on the device (no host mirror), so there is nothing + // to upload; NodeManagerHandle::deviceUpload would null-deref its host + // NodeManagerData. The handle is created device-resident and ready for + // kernel use. + nb::class_(m, "DeviceNodeManagerHandle", + "Owns the device memory backing a device-resident NodeManager. " + "Move-only. Obtain via nanovdb.cuda.createDeviceNodeManager(deviceGrid). " + "The NodeManager returned by mgr() is a device pointer: its node " + "accessors must only be used from CUDA kernels, never dereferenced on " + "the host.") + .def("size", + [](const HandleT& h) { return h.size(); }, + "Byte size of the device buffer backing this handle.") + .def( + "__bool__", + [](const HandleT& h) { return h.size() != 0; }, + nb::is_operator(), + "True iff this handle owns a non-empty (device-resident) buffer.") + .def("mgr", &pyDeviceNodeMgr, + nb::keep_alive<0, 1>(), + "Return the typed device NodeManager for the grid this handle was " + "built from, or None if the BuildT is not Python-visible. The " + "returned NodeManager's `this` is a DEVICE pointer — use it only " + "from CUDA kernels. It keeps this handle alive."); +} + +// cuda::createNodeManager has one template instantiation per BuildT. We expose +// a single polymorphic createDeviceNodeManager(deviceGrid, stream) that picks +// the right one based on the runtime type of `deviceGrid` (any bound +// NanoGrid whose underlying pointer is a device pointer, e.g. from +// DeviceGridHandle.deviceGrid(n)). The created handle stores a raw pointer back +// to the device grid, so the handle must keep the grid alive. +template +static nb::object tryCreateDeviceNodeManager(nb::handle py_grid, cudaStream_t stream) +{ + using GridT = NanoGrid; + if (!nb::isinstance(py_grid)) { + return nb::object(); // sentinel: "not this BuildT, try next" + } + // &grid is the device pointer (the NanoGrid object wraps a device this). + auto* dGrid = &nb::cast(py_grid); + NodeManagerHandle handle; + { + nb::gil_scoped_release release; + handle = nanovdb::cuda::createNodeManager( + dGrid, nanovdb::cuda::DeviceBuffer(), stream); + } + return nb::cast(std::move(handle)); +} + +static void defineCreateDeviceNodeManager(nb::module_& m) +{ + m.def("createDeviceNodeManager", + [](nb::handle py_grid, uintptr_t stream) -> nb::object { + cudaStream_t s = reinterpret_cast(stream); + // Try every bound BuildT; first matching runtime type wins. +#define NANOVDB_PY_FOR_EACH_SCALAR_BUILDT(T, Suffix, GridTypeEnum) \ + if (auto obj = tryCreateDeviceNodeManager(py_grid, s); obj.is_valid()) { \ + return obj; \ + } +#define NANOVDB_PY_FOR_EACH_VECTOR_BUILDT(T, Suffix, AccessorName, GridTypeEnum) \ + if (auto obj = tryCreateDeviceNodeManager(py_grid, s); obj.is_valid()) { \ + return obj; \ + } +#define NANOVDB_PY_FOR_EACH_POINT_BUILDT(T, Suffix, GridTypeEnum) \ + if (auto obj = tryCreateDeviceNodeManager(py_grid, s); obj.is_valid()) { \ + return obj; \ + } +#define NANOVDB_PY_FOR_EACH_READONLY_BUILDT(T, Suffix, GridTypeEnum) \ + if (auto obj = tryCreateDeviceNodeManager(py_grid, s); obj.is_valid()) { \ + return obj; \ + } +#include "BuildTypes.def" + throw nb::type_error( + "createDeviceNodeManager: argument is not a NanoVDB device " + "grid of any bound BuildT. Pass a device grid obtained from " + "DeviceGridHandle.deviceGrid(n)."); + }, + "deviceGrid"_a, "stream"_a = 0, + // The constructed NodeManager stores a raw pointer back to the device + // grid; the handle must therefore keep the grid (and transitively the + // DeviceGridHandle that owns the grid's device buffer) alive. + nb::keep_alive<0, 1>(), + "Build a device-resident NodeManager for the given DEVICE grid, " + "returning a DeviceNodeManagerHandle that owns the underlying device " + "buffer. deviceGrid MUST be a device grid (from " + "DeviceGridHandle.deviceGrid(n)); passing a host grid is a usage " + "error. stream is a raw CUDA stream handle (Python int; 0 = default " + "stream). The handle's mgr() returns the typed device NodeManager and " + "keeps the source grid alive for as long as it lives."); +} + +void defineDeviceNodeManager(nb::module_& m) +{ + defineDeviceNodeManagerHandle(m); + defineCreateDeviceNodeManager(m); +} + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceStreamMap.cu b/nanovdb/nanovdb/python/cuda/PyDeviceStreamMap.cu new file mode 100644 index 0000000000..c61a6188b3 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDeviceStreamMap.cu @@ -0,0 +1,128 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifdef NANOVDB_USE_CUDA + +#include "PyDeviceStreamMap.h" + +#include + +#include +#include + +#include +#include + +#include + +namespace nb = nanobind; +using namespace nb::literals; + +namespace pynanovdb { + +void defineDeviceStreamMap(nb::module_& m) +{ + using DeviceStreamMap = nanovdb::cuda::DeviceStreamMap; + + nb::class_ cls(m, "DeviceStreamMap", + "Maps each suitable CUDA device id to a freshly-created stream. The " + "constructor filters devices by a DeviceType policy. Iterates as " + "(device_id -> raw stream handle) pairs."); + + nb::enum_(cls, "DeviceType", + "Device-inclusion policy for DeviceStreamMap.") + .value("Any", DeviceStreamMap::DeviceType::Any, + "Include every available device (no filtering).") + .value("PeerToPeer", DeviceStreamMap::DeviceType::PeerToPeer, + "Include only devices that can peer-access every already-added " + "device.") + .value("Unified", DeviceStreamMap::DeviceType::Unified, + "Include only devices supporting unified addressing + concurrent " + "managed access AND peer access to the already-added devices " + "(the default)."); + + cls + .def( + "__init__", + [](DeviceStreamMap* self, DeviceStreamMap::DeviceType type, + std::vector exclude, int verbose) { + // Touches every device (attribute queries, stream creation), + // so release the GIL. + nb::gil_scoped_release release; + new (self) DeviceStreamMap(type, std::move(exclude), verbose); + }, + "type"_a = DeviceStreamMap::DeviceType::Unified, + "exclude"_a = std::vector{}, + "verbose"_a = 0, + "Build a map of device id -> CUDA stream over the devices that " + "satisfy `type`. `exclude` is a list of device ids to skip; " + "`verbose` is 0 (quiet), 1 (print ignored devices) or 2 (print " + "included devices).") + .def("deviceCount", &DeviceStreamMap::deviceCount, + "Number of devices in this map.") + .def("__len__", &DeviceStreamMap::deviceCount, + "Number of devices in this map (same as deviceCount()).") + .def( + "getMinPageSize", + [](const DeviceStreamMap& map) { + nb::gil_scoped_release release; + return map.getMinPageSize(); + }, + "Minimum CUDA allocation granularity (in bytes) across all devices " + "in this map.") + .def( + "printDevInfo", + [](const DeviceStreamMap& map) { map.printDevInfo(); }, + "Print device information for every device in this map to stdout.") + .def( + "items", + [](const DeviceStreamMap& map) { + // Expose the underlying std::map as a Python + // dict {device_id -> raw stream handle (int)}. + nb::dict out; + for (const auto& kv : map) + out[nb::int_(kv.first)] = + nb::int_(reinterpret_cast(kv.second)); + return out; + }, + "Return a dict mapping each device id to its raw CUDA stream handle " + "(Python int).") + .def( + "stream", + [](const DeviceStreamMap& map, int deviceId) { + auto it = map.find(deviceId); + if (it == map.end()) + throw nb::key_error("DeviceStreamMap has no stream for that device id."); + return reinterpret_cast(it->second); + }, + "deviceId"_a, + "Raw CUDA stream handle (Python int) for the given device id.") + .def( + "__contains__", + [](const DeviceStreamMap& map, int deviceId) { + return map.find(deviceId) != map.end(); + }, + "deviceId"_a, + "True iff this map holds a stream for the given device id.") + .def( + "__getitem__", + [](const DeviceStreamMap& map, int deviceId) { + auto it = map.find(deviceId); + if (it == map.end()) + throw nb::key_error("DeviceStreamMap has no stream for that device id."); + return reinterpret_cast(it->second); + }, + "deviceId"_a, + "Raw CUDA stream handle (Python int) for the given device id.") + .def( + "__iter__", + [](const DeviceStreamMap& map) { + nb::list keys; + for (const auto& kv : map) keys.append(kv.first); + return nb::iter(keys); + }, + "Iterate over the device ids in this map."); +} + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceStreamMap.h b/nanovdb/nanovdb/python/cuda/PyDeviceStreamMap.h new file mode 100644 index 0000000000..d8a79a643d --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDeviceStreamMap.h @@ -0,0 +1,25 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYDEVICESTREAMMAP_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYDEVICESTREAMMAP_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +#ifdef NANOVDB_USE_CUDA +/// @brief Register nanovdb::cuda::DeviceStreamMap and its DeviceType enum on +/// the nanovdb.cuda submodule. DeviceStreamMap is a std::map that, on construction, filters the available devices by +/// a DeviceType policy and creates one stream per surviving device. +/// @note DeviceStreamMap.h defines its (non-template) ctor/dtor in the header +/// without `inline`, so it must be included in exactly ONE translation +/// unit — this one. +void defineDeviceStreamMap(nb::module_& m); +#endif + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceVoxelBlockManager.cu b/nanovdb/nanovdb/python/cuda/PyDeviceVoxelBlockManager.cu new file mode 100644 index 0000000000..dfb243fdd2 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDeviceVoxelBlockManager.cu @@ -0,0 +1,625 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifdef NANOVDB_USE_CUDA + +#include +#include + +#include +#include +#include + +#include + +#include +#include +#include +#include +#include + +namespace nb = nanobind; +using namespace nb::literals; +using namespace nanovdb; +using nanovdb::tools::VoxelBlockManagerBase; +using nanovdb::tools::VoxelBlockManagerHandle; + +namespace pynanovdb { + +// ----------------------- Log2BlockWidth dispatch -------------------------- +// +// Mirrors the host dispatchLog2BlockWidth (PyVoxelBlockManager.cc): turn the +// runtime log2BlockWidth into one of the four compile-time widths the +// device builder is instantiated for (BlockWidth = 64, 128, 256, 512). +template +static auto dispatchLog2BlockWidth(int log2BlockWidth, F&& fn) +{ + switch (log2BlockWidth) { + case 6: return fn(std::integral_constant{}); + case 7: return fn(std::integral_constant{}); + case 8: return fn(std::integral_constant{}); + case 9: return fn(std::integral_constant{}); + default: + throw nb::value_error( + "VoxelBlockManager: log2BlockWidth must be 6, 7, 8, or 9 " + "(BlockWidth = 64, 128, 256, or 512). Larger widths are not " + "bound in Python by default."); + } +} + +// PyDeviceVBMHandle wraps the device VoxelBlockManagerHandle and carries the +// log2BlockWidth it was built with (parallel to the host PyVBMHandle). The +// C++ handle does NOT store log2BlockWidth itself, so recording it once at +// build time keeps the jumpMap view shape derivable from the handle rather +// than from a caller who could spoof it. +struct PyDeviceVBMHandle +{ + VoxelBlockManagerHandle handle; + int log2BlockWidth = 6; + + PyDeviceVBMHandle() = default; + PyDeviceVBMHandle(VoxelBlockManagerHandle&& h, + int lbw) noexcept + : handle(std::move(h)), log2BlockWidth(lbw) {} + + PyDeviceVBMHandle(const PyDeviceVBMHandle&) = delete; + PyDeviceVBMHandle& operator=(const PyDeviceVBMHandle&) = delete; + PyDeviceVBMHandle(PyDeviceVBMHandle&&) = default; + PyDeviceVBMHandle& operator=(PyDeviceVBMHandle&&) = default; + + uint64_t blockCount() const { return handle.blockCount(); } + uint64_t firstOffset() const { return handle.firstOffset(); } + uint64_t lastOffset() const { return handle.lastOffset(); } + void reset() { handle.reset(); } + int blockWidth() const { return 1 << log2BlockWidth; } + int jumpMapLength() const { return 1 << (log2BlockWidth - 6); } +}; + +// ------------------- OnIndex device grid cast helper ---------------------- + +static NanoGrid* castOnIndexDeviceGrid(nb::handle py_grid, + const char* fn_name) +{ + if (!nb::isinstance>(py_grid)) { + std::string msg(fn_name); + msg += ": deviceGrid must be a NanoVDB device grid of build type " + "ValueOnIndex (OnIndexGrid), obtained from " + "DeviceGridHandle.deviceGrid(n)"; + throw nb::type_error(msg.c_str()); + } + // The device builder takes a non-const NanoGrid*; the + // returned object's underlying address IS the device pointer. It must NOT + // be dereferenced on the host — it is only passed to device kernels. + return &nb::cast&>(py_grid); +} + +// gatherBoxStencil / activeVoxelCoords write out[valueIndex] for each active +// voxel, the caller having sized `out` to activeVoxelCount + 1. That is valid +// only when active-voxel indexing is CONTIGUOUS (value index 0 = background, +// 1..N = the N active voxels) -- the VoxelBlockManager invariant. A grid built +// with per-node statistics / tile values (the createOnIndexGrid defaults) has +// value indices beyond activeVoxelCount, so those writes run out of bounds (an +// illegal memory access). Detect it cheaply (two D2H header reads) and raise a +// clear error instead of crashing. +// Returns the grid's value count so callers can also bounds-check their +// value-indexed arrays without a second D2H read. +static uint64_t requireContiguousIndexing(const NanoGrid* dGrid, + const char* fn_name) +{ + using Traits = nanovdb::util::cuda::DeviceGridTraits; + const uint64_t valueCount = Traits::getValueCount(dGrid); + const uint64_t activeCount = Traits::getActiveVoxelCount(dGrid); + if (valueCount != activeCount + 1) { + std::string msg(fn_name); + msg += ": requires an OnIndex grid with contiguous active-voxel indexing " + "(value index 0 = background, 1..N = the N active voxels), but this " + "grid has " + std::to_string(valueCount) + " values for " + + std::to_string(activeCount) + " active voxels -- it carries " + "per-node statistics and/or tile values. Rebuild it with " + "createOnIndexGrid(grid, includeStats=False, includeTiles=False) " + "or voxelsToOnIndexGrid."; + throw nb::value_error(msg.c_str()); + } + return valueCount; +} + +// The gather / decode kernels index their value-keyed arrays by every value +// index in [0, valueCount), so a shorter array is an out-of-bounds device +// access; reject it with a Python exception instead. +static void requireValueCountRows(size_t rows, uint64_t valueCount, + const char* fn_name, const char* array_name) +{ + if (rows < valueCount) { + std::string msg(fn_name); + msg += ": "; + msg += array_name; + msg += " covers " + std::to_string(rows) + " rows but the grid stores " + + std::to_string(valueCount) + + " value indices (activeVoxelCount + 1); it must cover all of them."; + throw nb::value_error(msg.c_str()); + } +} + +// ------------------- DeviceVoxelBlockManagerHandle binding ----------------- + +static void defineHandle(nb::module_& m) +{ + nb::class_(m, "DeviceVoxelBlockManagerHandle", + "Owns the device-resident firstLeafID / jumpMap metadata buffers " + "backing a device VoxelBlockManager. Constructed by " + "nanovdb.tools.cuda.buildVoxelBlockManager. The firstLeafID / jumpMap " + "buffers are exposed zero-copy to CuPy / PyTorch / Numba via the CUDA " + "Array Interface and DLPack.") + .def("blockCount", &PyDeviceVBMHandle::blockCount, + "Number of voxel blocks managed by this handle.") + .def("firstOffset", &PyDeviceVBMHandle::firstOffset, + "Sequential voxel index of the first active voxel covered " + "by this handle (1 by default when the handle covers the " + "full grid).") + .def("lastOffset", &PyDeviceVBMHandle::lastOffset, + "Sequential voxel index of the last active voxel covered " + "by this handle.") + .def("reset", &PyDeviceVBMHandle::reset, + "Release this handle's device buffers and reset it to the empty state.") + .def_prop_ro("log2BlockWidth", + [](const PyDeviceVBMHandle& h) { return h.log2BlockWidth; }, + "The log2BlockWidth this handle was built with. The jumpMap " + "view derives its shape from this value.") + .def_prop_ro("blockWidth", &PyDeviceVBMHandle::blockWidth, + "BlockWidth = 1 << log2BlockWidth (64, 128, 256, or 512).") + .def_prop_ro("jumpMapLength", &PyDeviceVBMHandle::jumpMapLength, + "JumpMapLength = BlockWidth / 64 (1, 2, 4, or 8).") + .def( + "__bool__", + [](const PyDeviceVBMHandle& h) { return h.blockCount() > 0; }, + nb::is_operator()) + // ------------------- raw device pointers ------------------- + .def( + "first_leaf_id_ptr", + [](PyDeviceVBMHandle& h) { + return reinterpret_cast(h.handle.deviceFirstLeafID()); + }, + "Raw device pointer to the firstLeafID array (uint32 x blockCount) " + "as a Python int (0 if the handle is empty).") + .def( + "jump_map_ptr", + [](PyDeviceVBMHandle& h) { + return reinterpret_cast(h.handle.deviceJumpMap()); + }, + "Raw device pointer to the jumpMap array " + "(uint64 x blockCount x jump_map_length) as a Python int (0 if the " + "handle is empty).") + // ------------------- firstLeafID device view ------------------- + // Zero-copy DEVICE view of the (blockCount,) uint32 firstLeafID array, + // mirroring the host firstLeafID() but using the Phase-B device + // interop (nb::device::cuda ndarray) instead of a host numpy view. + .def( + "firstLeafID", + [](nb::handle py_self) -> nb::object { + auto& h = nb::cast(py_self); + size_t shape[1] = {static_cast(h.blockCount())}; + // A default-constructed / reset() handle has a null + // deviceFirstLeafID(); still return an empty (0,) array so + // callers don't branch on a None sentinel. The dummy non-null + // pointer (the handle itself) keeps nanobind happy; nothing is + // read since the leading shape is 0. + uint32_t* raw = h.handle.deviceFirstLeafID(); + void* data = (raw != nullptr) ? static_cast(raw) + : static_cast(&h); + nb::ndarray> arr( + data, size_t(1), shape, py_self); + return nb::cast(arr, nb::rv_policy::reference); + }, + // No nb::keep_alive<0,1> here: the returned no-framework device + // ndarray is exported as a DLPack capsule (not weak-referenceable), + // so keep_alive would throw "could not create a weak reference". + // Lifetime is already anchored by the ndarray's py_self owner arg. + "Return a zero-copy (blockCount,) uint32 DEVICE array view of the " + "firstLeafID array, consumable by CuPy / PyTorch / Numba. Returns " + "an empty (0,) array on a default-constructed or reset() handle. " + "The view keeps this handle alive.") + // ------------------- jumpMap device view ------------------- + // Zero-copy DEVICE view of the (blockCount, jumpMapLength) uint64 + // jumpMap. jumpMapLength derives from the handle's recorded + // log2BlockWidth, never the caller, so the view exactly covers the + // allocated buffer. + .def( + "jumpMap", + [](nb::handle py_self) -> nb::object { + auto& h = nb::cast(py_self); + size_t shape[2] = {static_cast(h.blockCount()), + static_cast(h.jumpMapLength())}; + uint64_t* raw = h.handle.deviceJumpMap(); + void* data = (raw != nullptr) ? static_cast(raw) + : static_cast(&h); + nb::ndarray> arr( + data, size_t(2), shape, py_self); + return nb::cast(arr, nb::rv_policy::reference); + }, + // No keep_alive (see firstLeafID): the device-ndarray capsule is + // not weak-referenceable; the py_self owner arg anchors lifetime. + "Return a zero-copy (blockCount, jump_map_length) uint64 DEVICE " + "array view of the jumpMap, consumable by CuPy / PyTorch / Numba. " + "The shape is determined by the log2BlockWidth the handle was " + "built with. Returns an empty (0, jump_map_length) array on a " + "default-constructed or reset() handle. The view keeps this handle " + "alive.") + // ------------------- CUDA Array Interface (v3) ------------------- + // CAI describing the firstLeafID array as 1-D uint32. (CuPy / Numba + // consume the __cuda_array_interface__ attribute directly; the jumpMap + // is reachable as a 2-D device ndarray via jumpMap() above.) + .def_prop_ro( + "__cuda_array_interface__", + [](PyDeviceVBMHandle& h) { + nb::dict iface; + iface["shape"] = nb::make_tuple(h.blockCount()); + iface["typestr"] = "(h.handle.deviceFirstLeafID()), + false); + iface["version"] = 3; + iface["strides"] = nb::none(); + iface["stream"] = 1; + return iface; + }, + "CUDA Array Interface (v3) view of the firstLeafID array as 1-D " + "uint32 — lets CuPy / Numba / PyTorch consume it zero-copy. The " + "jumpMap is available as a 2-D device array via jumpMap().") + .def( + "__dlpack_device__", + [](const PyDeviceVBMHandle&) { + int device = 0; + cudaGetDevice(&device); + return nb::make_tuple(2, device); // 2 == kDLCUDA + }, + "DLPack device tuple (kDLCUDA, device_id) for the firstLeafID " + "device buffer.") + .def( + "__dlpack__", + [](nb::handle self, nb::handle /*stream*/) { + auto& h = nb::cast(self); + size_t shape[1] = {static_cast(h.blockCount())}; + uint32_t* raw = h.handle.deviceFirstLeafID(); + void* data = (raw != nullptr) ? static_cast(raw) + : static_cast(&h); + // nb::cast of a no-framework device ndarray IS the "dltensor" + // capsule (what __dlpack__ must return); return it directly. + nb::ndarray> arr( + data, size_t(1), shape, self); + return nb::cast(arr, nb::rv_policy::reference); + }, + "stream"_a = nb::none(), + "DLPack capsule exporting the firstLeafID array as 1-D uint32, " + "parented to this handle. The jumpMap is available as a 2-D device " + "array via jumpMap()."); +} + +// ------------------- buildVoxelBlockManager (device) binding --------------- + +static void defineBuild(nb::module_& m) +{ + m.def("buildVoxelBlockManager", + [](nb::handle py_grid, + int log2BlockWidth, + uint64_t firstOffset, + uint64_t lastOffset, + uint64_t nBlocks, + uintptr_t stream) -> PyDeviceVBMHandle { + auto* dGrid = castOnIndexDeviceGrid(py_grid, "buildVoxelBlockManager"); + cudaStream_t s = reinterpret_cast(stream); + return dispatchLog2BlockWidth(log2BlockWidth, [&](auto W) { + constexpr int LBW = decltype(W)::value; + using Base = VoxelBlockManagerBase; + constexpr uint64_t BlockWidth = Base::BlockWidth; + // firstOffset, if nonzero, must satisfy firstOffset == 1 + // (mod BlockWidth); the C++ builder only NANOVDB_ASSERTs this + // (a no-op in release), so validate it here for a clear error. + if (firstOffset != 0 && + ((firstOffset - 1) & (BlockWidth - 1)) != 0) { + throw nb::value_error( + "buildVoxelBlockManager: firstOffset must satisfy " + "firstOffset == 1 (mod BlockWidth). Pass 0 (the " + "default) to let the implementation use 1."); + } + VoxelBlockManagerHandle handle; + { + // The device builder reads activeVoxelCount / lowerCount + // from device memory and launches kernels; pure C++/CUDA + // touching no Python objects, so release the GIL. + nb::gil_scoped_release release; + handle = nanovdb::tools::cuda::buildVoxelBlockManager< + LBW, nanovdb::cuda::DeviceBuffer>( + dGrid, firstOffset, lastOffset, nBlocks, s); + } + return PyDeviceVBMHandle(std::move(handle), LBW); + }); + }, + "deviceGrid"_a, + "log2BlockWidth"_a = 6, + "firstOffset"_a = 0, + "lastOffset"_a = 0, + "nBlocks"_a = 0, + "stream"_a = 0, + // The device builder reads the grid from device memory; keep the + // device grid (and its owning DeviceGridHandle) alive for the duration + // of the build. The returned handle owns its own device metadata + // buffers, so it does NOT need to keep the grid alive afterwards. + "Build a device-side VoxelBlockManager from an OnIndex DEVICE grid. " + "deviceGrid MUST be a device grid (from " + "DeviceGridHandle.deviceGrid(n)); passing a host grid is a usage " + "error. log2BlockWidth selects the per-block active-voxel count " + "(6=64, 7=128, 8=256, 9=512). Pass 0 for firstOffset / lastOffset / " + "nBlocks to use the full grid (first active voxel through " + "activeVoxelCount, minimum block count); these are read from device " + "memory. firstOffset, if nonzero, must satisfy firstOffset == 1 " + "(mod BlockWidth). stream is a raw CUDA stream handle (Python int; 0 = " + "default stream)."); +} + +// NOTE: decodeInverseMaps is intentionally NOT bound on the device. The device +// decode (VoxelBlockManager::decodeInverseMaps) is a __device__ +// function: it uses threadIdx / __syncthreads / shared-memory output arrays and +// is callable only from within a CUDA kernel, never from host code. Users who +// want device-side decode should call it from their own kernels via the shipped +// header , feeding it the firstLeafID +// / jumpMap device pointers (first_leaf_id_ptr / jump_map_ptr above) and the +// device grid pointer. + +// ------------------- gatherBoxStencil (VBM box-stencil gather) ------------- +// +// Materialise, for every active voxel, the values of its 3x3x3 neighbourhood +// into a dense (valueCount, 27) array -- the "dense-ise the sparse stencil" +// bridge that lets tile / array frameworks (CuPy, cuTile, ...) run VDB stencils +// without pointer-chasing the tree. Column j is the 3x3x3 spoke +// (di+1)*9 + (dj+1)*3 + (dk+1): the centre is column 13 and the six faces are +// columns 4, 10, 12, 14, 16, 22. Inactive neighbours read the sidecar's +// background slot (value index 0). +template +__global__ void gatherBoxStencilKernel( + const NanoGrid* grid, + const uint32_t* firstLeafID, const uint64_t* jumpMap, uint64_t firstOffset, + const T* values, T* out) +{ + constexpr int BW = 1 << Log2BlockWidth; + constexpr int JML = BW / 64; + using VBM = nanovdb::tools::cuda::VoxelBlockManager; + const uint64_t blockFirstOffset = firstOffset + uint64_t(blockIdx.x) * BW; + const int tID = threadIdx.x; + uint32_t leafIndex; + uint16_t voxelOffset; + VBM::template decodeInverseMap( + grid, firstLeafID[blockIdx.x], &jumpMap[uint64_t(blockIdx.x) * JML], + blockFirstOffset, tID, leafIndex, voxelOffset); + if (leafIndex == VBM::UnusedLeafIndex) return; + uint64_t st[27]; + VBM::template computeBoxStencil( + grid, leafIndex, voxelOffset, st); + const uint64_t c = st[13]; // centre value index + #pragma unroll + for (int j = 0; j < 27; ++j) out[c * 27 + j] = values[st[j]]; // 0 -> background +} + +template void defineGatherBoxStencil(nb::module_& m, const char* name) +{ + m.def( + name, + [](nb::handle py_grid, + nb::ndarray, nb::c_contig, nb::device::cuda> values, + nb::ndarray, nb::c_contig, nb::device::cuda> out, + int log2BlockWidth, uintptr_t stream) { + auto* dGrid = castOnIndexDeviceGrid(py_grid, "gatherBoxStencil"); + const uint64_t vc = requireContiguousIndexing(dGrid, "gatherBoxStencil"); + requireValueCountRows(values.size(), vc, "gatherBoxStencil", "values"); + requireValueCountRows(out.shape(0), vc, "gatherBoxStencil", "out"); + cudaStream_t s = reinterpret_cast(stream); + const T* dVals = values.data(); + T* dOut = out.data(); + // Build a transient VBM, then one block per VBM block decodes and + // gathers the 27 neighbour values; pure CUDA, so release the GIL. + nb::gil_scoped_release release; + dispatchLog2BlockWidth(log2BlockWidth, [&](auto W) { + constexpr int LBW = decltype(W)::value; + auto handle = nanovdb::tools::cuda::buildVoxelBlockManager< + LBW, nanovdb::cuda::DeviceBuffer>(dGrid, 0, 0, 0, s); + const uint32_t bc = static_cast(handle.blockCount()); + if (bc) + gatherBoxStencilKernel<<>>( + dGrid, handle.deviceFirstLeafID(), handle.deviceJumpMap(), + handle.firstOffset(), dVals, dOut); + cudaCheck(cudaStreamSynchronize(s)); + return 0; + }); + }, + "deviceGrid"_a, "values"_a, "out"_a, "log2BlockWidth"_a = 9, "stream"_a = 0, + "Gather the 3x3x3 box-stencil neighbour values of every active voxel " + "into a dense (valueCount, 27) array -- the bridge that lets tile / " + "array frameworks (CuPy, cuTile, ...) run VDB stencils without pointer-" + "chasing the tree. deviceGrid is an OnIndex device grid from " + "DeviceGridHandle.deviceGrid(n); values is a 1-D device array indexed by " + "value index (the per-voxel sidecar; entry 0 is the background slot); out " + "is a 2-D device array of shape (rows, 27) with rows >= valueCount, " + "filled so out[k, j] is voxel k's neighbour value at 3x3x3 spoke " + "j = (di+1)*9+(dj+1)*3+(dk+1) (centre j=13; the six faces are " + "j = 4, 10, 12, 14, 16, 22). Inactive neighbours read values[0]. A " + "transient VoxelBlockManager is built internally at log2BlockWidth " + "(6/7/8/9). stream is a raw CUDA stream handle (Python int; 0 = default " + "stream)."); +} + +// ------------------- gatherBoxStencilColumns (subset of the 27 spokes) ----- +// +// Like gatherBoxStencil, but writes only a chosen SUBSET of the 27 box-stencil +// spokes into an (valueCount, K) array. The VBM decode still computes all 27 +// neighbour indices in registers (that work is cheap and shared), so the only +// saving is the output table -- which matters when a kernel needs just a handful +// of the 27 neighbours (e.g. an SDF mesher's 8 corners + 6 faces): a (N, K) table +// instead of (N, 27). The K (<=27) spoke indices pass by value, no device scratch. +struct ColumnSpokes { int s[27]; }; + +template +__global__ void gatherBoxStencilColumnsKernel( + const NanoGrid* grid, + const uint32_t* firstLeafID, const uint64_t* jumpMap, uint64_t firstOffset, + const T* values, T* out, ColumnSpokes spokes, int K) +{ + constexpr int BW = 1 << Log2BlockWidth; + constexpr int JML = BW / 64; + using VBM = nanovdb::tools::cuda::VoxelBlockManager; + const uint64_t blockFirstOffset = firstOffset + uint64_t(blockIdx.x) * BW; + const int tID = threadIdx.x; + uint32_t leafIndex; + uint16_t voxelOffset; + VBM::template decodeInverseMap( + grid, firstLeafID[blockIdx.x], &jumpMap[uint64_t(blockIdx.x) * JML], + blockFirstOffset, tID, leafIndex, voxelOffset); + if (leafIndex == VBM::UnusedLeafIndex) return; + uint64_t st[27]; + VBM::template computeBoxStencil( + grid, leafIndex, voxelOffset, st); + const uint64_t c = st[13]; // centre value index + for (int col = 0; col < K; ++col) out[c * K + col] = values[st[spokes.s[col]]]; +} + +template void defineGatherBoxStencilColumns(nb::module_& m, const char* name) +{ + m.def( + name, + [](nb::handle py_grid, + nb::ndarray, nb::c_contig, nb::device::cuda> values, + nb::ndarray, nb::c_contig, nb::device::cuda> out, + nb::ndarray, nb::c_contig> spokes, + int log2BlockWidth, uintptr_t stream) { + auto* dGrid = castOnIndexDeviceGrid(py_grid, "gatherBoxStencilColumns"); + const uint64_t vc = + requireContiguousIndexing(dGrid, "gatherBoxStencilColumns"); + requireValueCountRows(values.size(), vc, "gatherBoxStencilColumns", "values"); + requireValueCountRows(out.shape(0), vc, "gatherBoxStencilColumns", "out"); + const int K = static_cast(spokes.shape(0)); + if (K < 1 || K > 27) + throw nb::value_error("gatherBoxStencilColumns: len(spokes) must be in [1, 27]."); + if (static_cast(out.shape(1)) != K) + throw nb::value_error("gatherBoxStencilColumns: out.shape[1] must equal len(spokes)."); + ColumnSpokes sp{}; // copy + validate spokes host-side + const int32_t* hSpokes = spokes.data(); + for (int i = 0; i < K; ++i) { + if (hSpokes[i] < 0 || hSpokes[i] >= 27) + throw nb::value_error("gatherBoxStencilColumns: each spoke must be in [0, 27)."); + sp.s[i] = static_cast(hSpokes[i]); + } + cudaStream_t s = reinterpret_cast(stream); + const T* dVals = values.data(); + T* dOut = out.data(); + nb::gil_scoped_release release; + dispatchLog2BlockWidth(log2BlockWidth, [&](auto W) { + constexpr int LBW = decltype(W)::value; + auto handle = nanovdb::tools::cuda::buildVoxelBlockManager< + LBW, nanovdb::cuda::DeviceBuffer>(dGrid, 0, 0, 0, s); + const uint32_t bc = static_cast(handle.blockCount()); + if (bc) + gatherBoxStencilColumnsKernel<<>>( + dGrid, handle.deviceFirstLeafID(), handle.deviceJumpMap(), + handle.firstOffset(), dVals, dOut, sp, K); + cudaCheck(cudaStreamSynchronize(s)); + return 0; + }); + }, + "deviceGrid"_a, "values"_a, "out"_a, "spokes"_a, "log2BlockWidth"_a = 9, "stream"_a = 0, + "Like gatherBoxStencil, but gathers only a chosen SUBSET of the 27 box-" + "stencil spokes into a dense (valueCount, K) array: out[k, col] is voxel " + "k's neighbour value at spoke spokes[col], where spoke " + "j = (di+1)*9+(dj+1)*3+(dk+1) (centre 13). `spokes` is a 1-D HOST int32 " + "array of K (1..27) spoke indices in [0,27); out has shape (rows, K) with " + "rows >= valueCount. Halves/shrinks the table when only a handful of the " + "27 neighbours are needed (e.g. an SDF mesher's corners + faces). Inactive " + "neighbours read values[0]; the grid must use contiguous active-voxel " + "indexing. stream is a raw CUDA stream handle (Python int; 0 = default)."); +} + +// ------------------- activeVoxelCoords (VBM coordinate decode) ------------- +// +// Write each active voxel's index-space coordinate into a dense (valueCount, 3) +// int32 array, keyed by value index -- the decode companion to +// gatherBoxStencil. Lets callers recover "where is value index k" (e.g. to bake +// a result back into a grid, or to scatter to a dense field) without a +// hand-written decode kernel. +template +__global__ void activeVoxelCoordsKernel( + const NanoGrid* grid, + const uint32_t* firstLeafID, const uint64_t* jumpMap, uint64_t firstOffset, + int32_t* out) +{ + constexpr int BW = 1 << Log2BlockWidth; + constexpr int JML = BW / 64; + using VBM = nanovdb::tools::cuda::VoxelBlockManager; + __shared__ uint32_t smem_leafIndex[BW]; + __shared__ uint16_t smem_voxelOffset[BW]; + const uint64_t blockFirstOffset = firstOffset + uint64_t(blockIdx.x) * BW; + VBM::template decodeInverseMaps( + grid, firstLeafID[blockIdx.x], &jumpMap[uint64_t(blockIdx.x) * JML], + blockFirstOffset, smem_leafIndex, smem_voxelOffset); + const int tID = threadIdx.x; + if (smem_leafIndex[tID] == VBM::UnusedLeafIndex) return; + const auto& leaf = grid->tree().getFirstNode<0>()[smem_leafIndex[tID]]; + const Coord c = leaf.offsetToGlobalCoord(smem_voxelOffset[tID]); + const uint64_t idx = leaf.getValue(smem_voxelOffset[tID]); + out[idx * 3 + 0] = c[0]; + out[idx * 3 + 1] = c[1]; + out[idx * 3 + 2] = c[2]; +} + +void defineActiveVoxelCoords(nb::module_& m, const char* name) +{ + m.def( + name, + [](nb::handle py_grid, + nb::ndarray, nb::c_contig, nb::device::cuda> out, + int log2BlockWidth, uintptr_t stream) { + auto* dGrid = castOnIndexDeviceGrid(py_grid, "activeVoxelCoords"); + const uint64_t vc = requireContiguousIndexing(dGrid, "activeVoxelCoords"); + requireValueCountRows(out.shape(0), vc, "activeVoxelCoords", "out"); + cudaStream_t s = reinterpret_cast(stream); + int32_t* dOut = out.data(); + nb::gil_scoped_release release; + dispatchLog2BlockWidth(log2BlockWidth, [&](auto W) { + constexpr int LBW = decltype(W)::value; + auto handle = nanovdb::tools::cuda::buildVoxelBlockManager< + LBW, nanovdb::cuda::DeviceBuffer>(dGrid, 0, 0, 0, s); + const uint32_t bc = static_cast(handle.blockCount()); + if (bc) + activeVoxelCoordsKernel<<>>( + dGrid, handle.deviceFirstLeafID(), handle.deviceJumpMap(), + handle.firstOffset(), dOut); + cudaCheck(cudaStreamSynchronize(s)); + return 0; + }); + }, + "deviceGrid"_a, "out"_a, "log2BlockWidth"_a = 9, "stream"_a = 0, + "Write each active voxel's index-space coordinate into a dense " + "(rows, 3) int32 device array keyed by value index (rows >= valueCount); " + "out[k] is the (i, j, k) coordinate of value index k (row 0, the " + "background slot, is left untouched). The decode companion to " + "gatherBoxStencil -- recovers per-voxel coordinates without a " + "hand-written decode kernel (e.g. to bake a sidecar result back into a " + "grid). deviceGrid is an OnIndex device grid from " + "DeviceGridHandle.deviceGrid(n); a transient VoxelBlockManager is built " + "internally at log2BlockWidth (6/7/8/9). stream is a raw CUDA stream " + "handle (Python int; 0 = default stream)."); +} + +void defineDeviceVoxelBlockManager(nb::module_& m) +{ + defineHandle(m); + defineBuild(m); + defineGatherBoxStencil(m, "gatherBoxStencil"); + defineGatherBoxStencil(m, "gatherBoxStencil"); + defineGatherBoxStencil(m, "gatherBoxStencil"); + defineGatherBoxStencil(m, "gatherBoxStencil"); + defineGatherBoxStencilColumns(m, "gatherBoxStencilColumns"); + defineGatherBoxStencilColumns(m, "gatherBoxStencilColumns"); + defineGatherBoxStencilColumns(m, "gatherBoxStencilColumns"); + defineGatherBoxStencilColumns(m, "gatherBoxStencilColumns"); + defineActiveVoxelCoords(m, "activeVoxelCoords"); +} + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyDilateGrid.cu b/nanovdb/nanovdb/python/cuda/PyDilateGrid.cu new file mode 100644 index 0000000000..b62ec517b4 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDilateGrid.cu @@ -0,0 +1,44 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#include "PyDilateGrid.h" + +#include + +#include +#include + +namespace nb = nanobind; +using namespace nb::literals; +// NOTE: deliberately NOT `using namespace nanovdb;`. These tools instantiate +// CUB DeviceScan, whose nvcc-generated host stub references unqualified +// `cuda::std::...`; with `nanovdb::cuda` in scope that becomes ambiguous and +// fails to compile. Fully qualify nanovdb:: instead (matches PyPointsToGrid.cu). + +namespace pynanovdb { + +template void defineDilateGrid(nb::module_& m, const char* name) +{ + m.def( + name, + [](nanovdb::NanoGrid* dGrid, int op, uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + // DilateGrid::getHandle launches kernels and synchronizes the + // stream; pure CUDA touching no Python objects, so release the GIL. + nb::gil_scoped_release release; + nanovdb::tools::cuda::DilateGrid dilator(dGrid, s); + dilator.setOperation(static_cast(op)); + return dilator.getHandle(); + }, + "dGrid"_a, + "op"_a = static_cast(nanovdb::tools::morphology::NN_FACE_EDGE_VERTEX), + "stream"_a = 0, + "Morphologically dilate a device OnIndex grid and return a fresh device " + "GridHandle of the dilated grid. op is a nearest-neighbor stencil: 6 " + "(NN_FACE) or 26 (NN_FACE_EDGE_VERTEX). NN_FACE_EDGE (18) is accepted by " + "the C++ setter but is not implemented and raises at getHandle time. " + "stream is a raw CUDA stream handle (Python int; 0 = default stream)."); +} + +template void defineDilateGrid(nb::module_&, const char*); + +} // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyDilateGrid.h b/nanovdb/nanovdb/python/cuda/PyDilateGrid.h new file mode 100644 index 0000000000..6e7eff9887 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDilateGrid.h @@ -0,0 +1,16 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYDILATEGRID_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYDILATEGRID_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +template void defineDilateGrid(nb::module_& m, const char* name); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyDistributedPointsToGrid.cu b/nanovdb/nanovdb/python/cuda/PyDistributedPointsToGrid.cu new file mode 100644 index 0000000000..1a8b507fd7 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDistributedPointsToGrid.cu @@ -0,0 +1,82 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifdef NANOVDB_USE_CUDA + +#include "PyDistributedPointsToGrid.h" + +#include + +#include + +#include +#include +#include + +namespace nb = nanobind; +using namespace nb::literals; + +namespace pynanovdb { + +template void defineDistributedPointsToGrid(nb::module_& m, const char* name) +{ + using ConverterT = nanovdb::tools::cuda::DistributedPointsToGrid; + using BufferT = nanovdb::cuda::UnifiedBuffer; + + nb::class_(m, name, + "Multi-GPU builder of a NanoVDB grid from an array of index-space voxel " + "coordinates, distributed over a DeviceMesh. Construct with a DeviceMesh " + "(which must outlive this object) plus a voxel scale and translation, " + "then call getHandle(voxels, count). On a single-GPU mesh this still " + "runs the trivial single-device path.") + .def( + "__init__", + [](ConverterT* self, const nanovdb::cuda::DeviceMesh& mesh, + double scale, nb::tuple trans) { + nanovdb::Vec3d t(0.0); + if (trans.size() == 3) + t = nanovdb::Vec3d(nb::cast(trans[0]), + nb::cast(trans[1]), + nb::cast(trans[2])); + else if (trans.size() != 0) + throw nb::value_error("translation must be a 3-tuple or empty."); + nb::gil_scoped_release release; + new (self) ConverterT(mesh, scale, t); + }, + "mesh"_a, "scale"_a = 1.0, "translation"_a = nb::make_tuple(0.0, 0.0, 0.0), + nb::keep_alive<1, 2>(), // keep the DeviceMesh alive for our lifetime + "Construct a converter over the given DeviceMesh. scale is the " + "uniform voxel size and translation is a 3-tuple world offset used " + "to build the output grid's index-to-world map. The DeviceMesh is " + "held by reference and MUST outlive this converter.") + .def( + "getHandle", + [](ConverterT& self, + nb::ndarray, nb::c_contig, nb::device::cuda_managed> voxels) { + // The (N, 3) c-contiguous int32 array is bit-compatible with a + // dense nanovdb::Coord[N] (Coord == 3 x int32). The pipeline + // issues cudaMemAdvise / cudaMemPrefetchAsync on this pointer, + // so the memory MUST be CUDA managed (unified) memory — hence the + // nb::device::cuda_managed constraint (DLPack kDLCUDAManaged=13). + // Plain device memory (cuda=2) is rejected: prefetch to a device + // ordinal is invalid on non-managed memory. + auto* coords = reinterpret_cast(voxels.data()); + const size_t count = voxels.shape(0); + nb::gil_scoped_release release; + return self.template getHandle( + coords, count, BufferT()); + }, + "voxels"_a, + "Rasterize the given (N, 3) int32 array of index-space voxel " + "coordinates into a fresh UnifiedGridHandle of type " + "NanoGrid. The array MUST be backed by CUDA managed " + "(unified) memory — the multi-GPU pipeline applies memory advise " + "and prefetch directly to its pointer."); +} + +template void defineDistributedPointsToGrid(nb::module_&, const char*); +template void defineDistributedPointsToGrid(nb::module_&, const char*); +template void defineDistributedPointsToGrid(nb::module_&, const char*); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyDistributedPointsToGrid.h b/nanovdb/nanovdb/python/cuda/PyDistributedPointsToGrid.h new file mode 100644 index 0000000000..a77161c9a3 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyDistributedPointsToGrid.h @@ -0,0 +1,24 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYDISTRIBUTEDPOINTSTOGRID_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYDISTRIBUTEDPOINTSTOGRID_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +#ifdef NANOVDB_USE_CUDA +/// @brief Register the multi-GPU nanovdb::tools::cuda::DistributedPointsToGrid< +/// BuildT> on the nanovdb.tools.cuda submodule. The Python class wraps a +/// nanovdb::cuda::DeviceMesh and a scale/translation (or unit map), and +/// getHandle(voxels, count) rasterizes an (N, 3) int32 unified-memory +/// array of index-space voxel coordinates into a UnifiedGridHandle. +/// @note The DeviceMesh passed to the constructor must outlive the converter. +template void defineDistributedPointsToGrid(nb::module_& m, const char* name); +#endif + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyIndexToGrid.cu b/nanovdb/nanovdb/python/cuda/PyIndexToGrid.cu new file mode 100644 index 0000000000..96fe5087fc --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyIndexToGrid.cu @@ -0,0 +1,131 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#include "PyIndexToGrid.h" + +#include + +#include +#include + +#include +#include // value count of a device grid + +namespace nb = nanobind; +using namespace nb::literals; +// NOTE: deliberately NOT `using namespace nanovdb;`. These tools instantiate +// CUB DeviceScan, whose nvcc-generated host stub references unqualified +// `cuda::std::...`; with `nanovdb::cuda` in scope that becomes ambiguous and +// fails to compile. Fully qualify nanovdb:: instead (matches PyPointsToGrid.cu). + +namespace pynanovdb { + +namespace { + +// indexToGrid reads d_srcValues[i] for every value index stored in the source +// grid (voxels, plus any per-node stats and tile slots), so the values array +// must cover the grid's full value count. Verify with a D2H header read and +// raise a Python exception instead of an out-of-bounds device read. +template +void requireValueCountRows(const nanovdb::NanoGrid* d_srcGrid, + size_t rows, const char* fnName) +{ + using Traits = nanovdb::util::cuda::DeviceGridTraits; + const uint64_t valueCount = Traits::getValueCount(d_srcGrid); + if (rows < valueCount) { + std::string msg(fnName); + msg += ": values array covers " + std::to_string(rows) + + " entries but the source IndexGrid stores " + + std::to_string(valueCount) + + " value indices (grid.valueCount()); the array must cover all " + "of them."; + throw nb::value_error(msg.c_str()); + } +} + +} // anonymous namespace + +// Scalar destination value type (e.g. float, double): d_srcValues is a flat, +// 1-D device array indexed by the IndexGrid's per-element uint64 indices. +template +void defineIndexToGridScalar(nb::module_& m, const char* name) +{ + m.def( + name, + [](nanovdb::NanoGrid* d_srcGrid, + nb::ndarray, nb::c_contig, nb::device::cuda> values, + uintptr_t stream) { + requireValueCountRows(d_srcGrid, values.size(), "indexToGrid"); + cudaStream_t s = reinterpret_cast(stream); + const DstBuildT* d_values = values.data(); + // indexToGrid launches kernels and synchronizes the stream; pure + // CUDA touching no Python objects, so release the GIL. + nb::gil_scoped_release release; + return nanovdb::tools::cuda::indexToGrid( + d_srcGrid, d_values, nanovdb::cuda::DeviceBuffer(), s); + }, + "d_srcGrid"_a, + "values"_a, + "stream"_a = 0, + "Combine a device IndexGrid (ValueIndex / ValueOnIndex) with a flat " + "1-D device array of destination values into a fresh device " + "GridHandle of the destination value type. values is indexed by the " + "IndexGrid's per-voxel / per-node uint64 indices, so it must be sized " + "to cover the grid's value count. stream is a raw CUDA stream handle " + "(Python int; 0 = default stream)."); +} + +// Vec3 destination value type (Vec3f / Vec3d): d_srcValues is an (N, 3) device +// array of the matching scalar; Vec3 is exactly three contiguous scalars so +// the c_contig tensor reinterprets element-for-element as Vec3*. +template +void defineIndexToGridVec3(nb::module_& m, const char* name) +{ + using ScalarT = typename DstBuildT::ValueType; + m.def( + name, + [](nanovdb::NanoGrid* d_srcGrid, + nb::ndarray, nb::c_contig, nb::device::cuda> values, + uintptr_t stream) { + requireValueCountRows(d_srcGrid, values.shape(0), "indexToGrid"); + cudaStream_t s = reinterpret_cast(stream); + const DstBuildT* d_values = reinterpret_cast(values.data()); + // indexToGrid launches kernels and synchronizes the stream; pure + // CUDA touching no Python objects, so release the GIL. + nb::gil_scoped_release release; + return nanovdb::tools::cuda::indexToGrid( + d_srcGrid, d_values, nanovdb::cuda::DeviceBuffer(), s); + }, + "d_srcGrid"_a, + "values"_a, + "stream"_a = 0, + "Combine a device IndexGrid (ValueIndex / ValueOnIndex) with an (N, 3) " + "device array of destination vector values into a fresh device " + "GridHandle of the destination Vec3 value type. The array is indexed " + "by the IndexGrid's per-voxel / per-node uint64 indices, so N must " + "cover the grid's value count. stream is a raw CUDA stream handle " + "(Python int; 0 = default stream)."); +} + +// Destination types: float / double / int32 (scalar) and Vec3f / Vec3d +// (vector), each for both index source types (ValueIndex and ValueOnIndex). +// DstBuildT is restricted by IndexToGrid's processLeafsKernel static_assert to +// non-special types (is_special = index || Fp || Point/bool/Mask), so the +// quantized / index / mask BuildTs are intentionally NOT instantiated; Int32 is +// a regular type, so it materialises an Int32Grid (e.g. an integer label grid). +template void defineIndexToGridScalar(nb::module_&, const char*); +template void defineIndexToGridScalar(nb::module_&, const char*); +template void defineIndexToGridScalar(nb::module_&, const char*); +template void defineIndexToGridScalar(nb::module_&, const char*); +template void defineIndexToGridScalar(nb::module_&, const char*); +template void defineIndexToGridScalar(nb::module_&, const char*); + +template void +defineIndexToGridVec3(nb::module_&, const char*); +template void +defineIndexToGridVec3(nb::module_&, const char*); +template void +defineIndexToGridVec3(nb::module_&, const char*); +template void +defineIndexToGridVec3(nb::module_&, const char*); + +} // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyIndexToGrid.h b/nanovdb/nanovdb/python/cuda/PyIndexToGrid.h new file mode 100644 index 0000000000..6a71216aa9 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyIndexToGrid.h @@ -0,0 +1,25 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYINDEXTOGRID_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYINDEXTOGRID_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +// Bind nanovdb::tools::cuda::indexToGrid for a (DstBuildT, SrcBuildT) pair. +// SrcBuildT must be an index build type (ValueIndex / ValueOnIndex); DstBuildT +// must be a non-special value type (float / double / Vec3f / ...). All +// instantiations are registered under the same Python name and disambiguated +// by nanobind on the source grid class and the value ndarray dtype/shape. +template +void defineIndexToGridScalar(nb::module_& m, const char* name); + +template +void defineIndexToGridVec3(nb::module_& m, const char* name); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyInjectData.cu b/nanovdb/nanovdb/python/cuda/PyInjectData.cu new file mode 100644 index 0000000000..bc2b23cd1b --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyInjectData.cu @@ -0,0 +1,247 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifdef NANOVDB_USE_CUDA + +#include "PyInjectData.h" + +#include + +#include +#include + +#include + +#include +#include // cudaCheck, operatorKernel +#include // leaf count of a device grid +#include // Inject*Functor + +namespace nb = nanobind; +using namespace nb::literals; +// Deliberately NOT `using namespace nanovdb;`: keep the device-grid type names +// fully qualified, matching the sibling tools.cuda bindings. + +namespace pynanovdb { + +namespace { + +// Cast a Python device-grid object to NanoGrid*. The returned +// object's underlying address IS the device pointer; it must NOT be +// dereferenced on the host -- it is only passed to device kernels. +nanovdb::NanoGrid* +castOnIndexDeviceGrid(nb::handle py_grid, const char* fn_name) +{ + if (!nb::isinstance>(py_grid)) { + std::string msg(fn_name); + msg += ": expected a NanoVDB device grid of build type ValueOnIndex " + "(OnIndexGrid), obtained from DeviceGridHandle.deviceGrid(n)"; + throw nb::type_error(msg.c_str()); + } + return &nb::cast&>(py_grid); +} + +// Leaf-node count read from device memory (one D2H copy of the tree header). +uint32_t leafCountOf(const nanovdb::NanoGrid* dGrid) +{ + using Traits = nanovdb::util::cuda::DeviceGridTraits; + return Traits::getTreeData(dGrid).mNodeCount[0]; +} + +// The injection kernels index the sidecar / predicate arrays by each grid's +// value indices, so an array shorter than the grid's value count is an +// out-of-bounds device access. Verify (one D2H header read per grid) and +// raise a Python exception instead. +void requireValueCountRows(const nanovdb::NanoGrid* dGrid, + size_t rows, const char* fnName, const char* arrayName) +{ + using Traits = nanovdb::util::cuda::DeviceGridTraits; + const uint64_t valueCount = Traits::getValueCount(dGrid); + if (rows < valueCount) { + std::string msg(fnName); + msg += ": "; + msg += arrayName; + msg += " covers " + std::to_string(rows) + " entries but its grid stores " + + std::to_string(valueCount) + + " value indices (grid.valueCount()); the array must cover all of them."; + throw nb::value_error(msg.c_str()); + } +} + +} // anonymous namespace + +template void defineInject(nb::module_& m, const char* name) +{ + m.def( + name, + [](nb::handle srcGrid, nb::handle dstGrid, + nb::ndarray, nb::c_contig, nb::device::cuda> srcSidecar, + nb::ndarray, nb::c_contig, nb::device::cuda> dstSidecar, + uintptr_t stream) { + auto* src = castOnIndexDeviceGrid(srcGrid, "inject"); + auto* dst = castOnIndexDeviceGrid(dstGrid, "inject"); + requireValueCountRows(src, srcSidecar.size(), "inject", "srcSidecar"); + requireValueCountRows(dst, dstSidecar.size(), "inject", "dstSidecar"); + cudaStream_t s = reinterpret_cast(stream); + const T* dSrc = srcSidecar.data(); + T* dDst = dstSidecar.data(); + const uint32_t srcLeafCount = leafCountOf(src); + using Op = nanovdb::util::cuda::InjectGridDataFunctor; + // operatorKernel launches one block per SOURCE leaf and copies the + // src/dst intersection bit-parallel per warp; pure CUDA, no Python. + nb::gil_scoped_release release; + if (srcLeafCount) + nanovdb::util::cuda::operatorKernel + <<>>(src, dst, dSrc, dDst); + cudaCheck(cudaStreamSynchronize(s)); + }, + "srcGrid"_a, "dstGrid"_a, "srcSidecar"_a, "dstSidecar"_a, "stream"_a = 0, + "Inject sidecar values from a source OnIndex device grid onto a " + "destination OnIndex device grid (injectData; NanoVDB 2.0 paper, " + "section 3.4). For every voxel present in BOTH grids the source sidecar " + "value is copied to that voxel's slot in the destination sidecar; " + "destination voxels with no source counterpart are left unchanged " + "(so the source need not be a subset -- the copy is over the " + "intersection). srcGrid / dstGrid are device grids from " + "DeviceGridHandle.deviceGrid(n); srcSidecar / dstSidecar are 1-D " + "device arrays indexed by each grid's value index (entry 0 is the " + "background slot, untouched). Wraps " + "nanovdb::util::cuda::InjectGridDataFunctor. stream is a raw CUDA " + "stream handle (Python int; 0 = default stream)."); +} + +template void defineInjectFeatures(nb::module_& m, const char* name) +{ + m.def( + name, + [](nb::handle srcGrid, nb::handle dstGrid, + nb::ndarray, nb::c_contig, nb::device::cuda> srcSidecar, + nb::ndarray, nb::c_contig, nb::device::cuda> dstSidecar, + uintptr_t stream) { + auto* src = castOnIndexDeviceGrid(srcGrid, "inject"); + auto* dst = castOnIndexDeviceGrid(dstGrid, "inject"); + if (srcSidecar.shape(1) != dstSidecar.shape(1)) + throw nb::value_error( + "inject: src and dst feature sidecars must share the same " + "feature dimension (shape[1])."); + requireValueCountRows(src, srcSidecar.shape(0), "inject", "srcSidecar"); + requireValueCountRows(dst, dstSidecar.shape(0), "inject", "dstSidecar"); + cudaStream_t s = reinterpret_cast(stream); + const T* dSrc = srcSidecar.data(); + T* dDst = dstSidecar.data(); + const size_t dim = srcSidecar.shape(1); + const uint32_t srcLeafCount = leafCountOf(src); + using Op = nanovdb::util::cuda::InjectGridFeatureFunctor; + nb::gil_scoped_release release; + if (srcLeafCount) + nanovdb::util::cuda::operatorKernel + <<>>(src, dst, dSrc, dDst, dim); + cudaCheck(cudaStreamSynchronize(s)); + }, + "srcGrid"_a, "dstGrid"_a, "srcSidecar"_a, "dstSidecar"_a, "stream"_a = 0, + "Inject vector-valued (feature) sidecar data across OnIndex device " + "grids -- the multi-channel form of inject. srcSidecar / dstSidecar " + "are 2-D device arrays of shape (value count, dim), row-major per voxel " + "(row 0 is the background slot); the feature dimension dim is taken " + "from shape[1] and must match. Values are copied for the src/dst voxel " + "intersection; the rest of the destination is left unchanged. Wraps " + "nanovdb::util::cuda::InjectGridFeatureFunctor. stream is a raw CUDA " + "stream handle (Python int; 0 = default stream)."); +} + +void defineInjectPredicateToMask(nb::module_& m, const char* name) +{ + m.def( + name, + [](nb::handle grid, + nb::ndarray, nb::c_contig, nb::device::cuda> predicate, + nb::ndarray, nb::c_contig, nb::device::cuda> leafMasks, + uintptr_t stream) { + auto* dGrid = castOnIndexDeviceGrid(grid, "injectPredicateToMask"); + requireValueCountRows(dGrid, predicate.size(), + "injectPredicateToMask", "predicate"); + cudaStream_t s = reinterpret_cast(stream); + const uint32_t leafCount = leafCountOf(dGrid); + constexpr size_t W = nanovdb::Mask<3>::WORD_COUNT; // 8 uint64 / leaf + if (leafMasks.size() < static_cast(leafCount) * W) + throw nb::value_error( + "injectPredicateToMask: leafMasks length must be at least " + "(leaf count) * 8 uint64 (one Mask<3> per leaf). A safe " + "upper bound is activeVoxelCount * 8, since every leaf " + "holds at least one active voxel."); + const bool* dPred = predicate.data(); + nanovdb::Mask<3>* dMask = + reinterpret_cast*>(leafMasks.data()); + using Op = nanovdb::util::cuda::InjectPredicateToMaskFunctor; + // One block per leaf; the functor zeroes each leaf mask, then sets + // the bit of every active voxel whose predicate slot is true. + nb::gil_scoped_release release; + if (leafCount) + nanovdb::util::cuda::operatorKernel + <<>>(dGrid, dPred, dMask); + cudaCheck(cudaStreamSynchronize(s)); + }, + "grid"_a, "predicate"_a, "leafMasks"_a, "stream"_a = 0, + "Build a per-leaf retain mask for pruneGrid from a boolean predicate " + "over an OnIndex device grid's value indices. grid is a device grid " + "from DeviceGridHandle.deviceGrid(n); predicate is a 1-D device bool " + "array indexed by value index (entry n true => keep that voxel); " + "leafMasks is a 1-D device uint64 output of length at least " + "(leaf count) * 8 (one nanovdb::Mask<3> per leaf, in leaf order), " + "ready to pass straight to pruneGrid; activeVoxelCount * 8 is a safe " + "size since every leaf holds at least one active voxel. Wraps " + "nanovdb::util::cuda::InjectPredicateToMaskFunctor. stream is a raw " + "CUDA stream handle (Python int; 0 = default stream)."); +} + +void defineInjectGridMask(nb::module_& m, const char* name) +{ + m.def( + name, + [](nb::handle srcGrid, nb::handle dstGrid, + nb::ndarray, nb::c_contig, nb::device::cuda> leafMasks, + uintptr_t stream) { + auto* src = castOnIndexDeviceGrid(srcGrid, "injectGridMask"); + auto* dst = castOnIndexDeviceGrid(dstGrid, "injectGridMask"); + cudaStream_t s = reinterpret_cast(stream); + const uint32_t dstLeafCount = leafCountOf(dst); + constexpr size_t W = nanovdb::Mask<3>::WORD_COUNT; // 8 uint64 / leaf + if (leafMasks.size() < static_cast(dstLeafCount) * W) + throw nb::value_error( + "injectGridMask: leafMasks length must be at least " + "(dst leaf count) * 8 uint64 (one Mask<3> per leaf). A safe " + "upper bound is the destination grid's activeVoxelCount * 8."); + nanovdb::Mask<3>* dMask = + reinterpret_cast*>(leafMasks.data()); + using Op = nanovdb::util::cuda::InjectGridMaskFunctor; + constexpr unsigned threads = 128; + nb::gil_scoped_release release; + if (dstLeafCount) + nanovdb::util::cuda::lambdaKernel + <<>>(dstLeafCount, Op{}, src, dst, dMask); + cudaCheck(cudaStreamSynchronize(s)); + }, + "srcGrid"_a, "dstGrid"_a, "leafMasks"_a, "stream"_a = 0, + "Build a per-leaf mask over the DESTINATION grid marking the voxels " + "that are ALSO active in the source grid (the src/dst intersection). " + "grid args are device grids from DeviceGridHandle.deviceGrid(n); " + "leafMasks is a 1-D device uint64 output of length at least " + "(dst leaf count) * 8 (one nanovdb::Mask<3> per leaf, in leaf order; " + "the destination activeVoxelCount * 8 is a safe size). Pass it to " + "pruneGrid to keep only the intersection. Wraps " + "nanovdb::util::cuda::InjectGridMaskFunctor. stream is a raw CUDA " + "stream handle (Python int; 0 = default stream)."); +} + +template void defineInject(nb::module_&, const char*); +template void defineInject(nb::module_&, const char*); +template void defineInject(nb::module_&, const char*); +template void defineInject(nb::module_&, const char*); +template void defineInjectFeatures(nb::module_&, const char*); +template void defineInjectFeatures(nb::module_&, const char*); +template void defineInjectFeatures(nb::module_&, const char*); +template void defineInjectFeatures(nb::module_&, const char*); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyInjectData.h b/nanovdb/nanovdb/python/cuda/PyInjectData.h new file mode 100644 index 0000000000..b2171b682b --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyInjectData.h @@ -0,0 +1,19 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYINJECTDATA_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYINJECTDATA_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +template void defineInject(nb::module_& m, const char* name); +template void defineInjectFeatures(nb::module_& m, const char* name); +void defineInjectPredicateToMask(nb::module_& m, const char* name); +void defineInjectGridMask(nb::module_& m, const char* name); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyMergeGrids.cu b/nanovdb/nanovdb/python/cuda/PyMergeGrids.cu new file mode 100644 index 0000000000..32c9dde8a1 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyMergeGrids.cu @@ -0,0 +1,71 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#include "PyMergeGrids.h" + +#include +#include +#include + +#include + +namespace nb = nanobind; +using namespace nb::literals; +// NOTE: deliberately NOT `using namespace nanovdb;`. These tools instantiate +// CUB DeviceScan, whose nvcc-generated host stub references unqualified +// `cuda::std::...`; with `nanovdb::cuda` in scope that becomes ambiguous and +// fails to compile. Fully qualify nanovdb:: instead (matches PyPointsToGrid.cu). + +namespace pynanovdb { + +template void defineMergeGrids(nb::module_& m, const char* name) +{ + m.def( + name, + [](nanovdb::NanoGrid* dGrid1, + nanovdb::NanoGrid* dGrid2, + uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + // MergeGrids::getHandle launches kernels and synchronizes the + // stream; pure CUDA touching no Python objects, so release the GIL. + nb::gil_scoped_release release; + nanovdb::tools::cuda::MergeGrids merger(dGrid1, dGrid2, s); + return merger.getHandle(); + }, + "dGrid1"_a, + "dGrid2"_a, + "stream"_a = 0, + "Topologically merge (active-mask union) two device OnIndex grids and " + "return a fresh device GridHandle of the union. The operation is " + "strictly binary; chain calls to union more than two grids. Output " + "metadata is taken from dGrid1. stream is a raw CUDA stream handle " + "(Python int; 0 = default stream)."); + + // List overload: N-ary merge + using GridT = nanovdb::NanoGrid; + m.def( + name, + [](nb::sequence grids_seq, uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + // Collect device-grid pointers (touches Python -> GIL held). + std::vector grids; + const size_t n = nb::len(grids_seq); + grids.reserve(n); + for (size_t i = 0; i < n; ++i) + grids.push_back(nb::cast(grids_seq[i])); + if (grids.empty()) + throw nb::value_error("mergeGrids: empty grid list"); + nb::gil_scoped_release release; + return nanovdb::tools::cuda::MergeGrids(grids, s).getHandle(); + }, + "grids"_a, + "stream"_a = 0, + "Topologically merge (active-mask union) a list of device OnIndex grids " + "into one fresh device GridHandle in a single N-ary pass. Pass device " + "grids, e.g. [h.deviceGrid(0) for h in handles]. Output metadata is taken " + "from the first grid. stream is a raw CUDA stream handle (Python int; " + "0 = default stream)."); +} + +template void defineMergeGrids(nb::module_&, const char*); + +} // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyMergeGrids.h b/nanovdb/nanovdb/python/cuda/PyMergeGrids.h new file mode 100644 index 0000000000..9ddb8958c3 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyMergeGrids.h @@ -0,0 +1,16 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYMERGEGRIDS_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYMERGEGRIDS_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +template void defineMergeGrids(nb::module_& m, const char* name); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyMeshToGrid.cu b/nanovdb/nanovdb/python/cuda/PyMeshToGrid.cu new file mode 100644 index 0000000000..4b8fdca1a0 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyMeshToGrid.cu @@ -0,0 +1,79 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#include "PyMeshToGrid.h" +#include "PyValidate.h" + +#include +#include +#include + +#include +#include + +#include + +namespace nb = nanobind; +using namespace nb::literals; +// NOTE: deliberately NOT `using namespace nanovdb;`. These tools instantiate +// CUB DeviceScan, whose nvcc-generated host stub references unqualified +// `cuda::std::...`; with `nanovdb::cuda` in scope that becomes ambiguous and +// fails to compile. Fully qualify nanovdb:: instead (matches PyPointsToGrid.cu). + +namespace pynanovdb { + +void defineMeshToGrid(nb::module_& m, const char* name) +{ + m.def( + name, + [](nb::ndarray, nb::c_contig, nb::device::cuda> points, + nb::ndarray, nb::c_contig, nb::device::cuda> triangles, + double voxelSize, float halfWidth, const std::string& gridName, + uintptr_t stream) { + // Map::set only debug-asserts a positive voxel size, and a + // non-positive halfWidth reverses the raster bounds; validate both + // before any CUDA work. + requirePositiveFinite(voxelSize, "meshToGrid", "voxelSize"); + requirePositiveFinite(halfWidth, "meshToGrid", "halfWidth"); + cudaStream_t s = reinterpret_cast(stream); + // Vec3f / Vec3i are three contiguous scalars, so the c_contig + // (N, 3) tensors reinterpret element-for-element. + const nanovdb::Vec3f* d_points = + reinterpret_cast(points.data()); + const nanovdb::Vec3i* d_triangles = + reinterpret_cast(triangles.data()); + const uint32_t pointCount = static_cast(points.shape(0)); + const uint32_t triangleCount = static_cast(triangles.shape(0)); + nanovdb::Map map; + map.set(voxelSize, nanovdb::Vec3d(0.0, 0.0, 0.0), 1.0); + // MeshToGrid launches kernels and synchronizes the stream; pure + // CUDA touching no Python objects, so release the GIL. + nb::gil_scoped_release release; + nanovdb::tools::cuda::MeshToGrid converter( + d_points, pointCount, d_triangles, triangleCount, map, s); + converter.setNarrowBandWidth(halfWidth); + if (!gridName.empty()) + converter.setGridName(gridName); + // Compute a checksum during the build so the result validates. + converter.setChecksum(nanovdb::CheckMode::Full); + return converter.getHandleAndUDF(); + }, + "points"_a, + "triangles"_a, + "voxelSize"_a = 1.0, + "halfWidth"_a = 3.0, + "gridName"_a = "", + "stream"_a = 0, + "Rasterize a triangle mesh into a narrow-band unsigned distance field on " + "the device. points is an (N, 3) float32 CUDA array of vertex world " + "positions; triangles is an (M, 3) int32 CUDA array of vertex indices. " + "Returns a tuple (handle, udf): handle is a device ValueOnIndex " + "GridHandle holding the narrow-band topology, and udf is a " + "nanovdb.cuda.DeviceBuffer of (valueCount) float32 unsigned distances " + "(in voxel units) indexed by the grid's per-voxel value index -- feed it " + "straight to indexToGrid to bake a Float distance grid. voxelSize is the " + "world size of a voxel; halfWidth is the narrow-band half-width in voxel " + "units; stream is a raw CUDA stream handle (Python int; 0 = default " + "stream)."); +} + +} // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyMeshToGrid.h b/nanovdb/nanovdb/python/cuda/PyMeshToGrid.h new file mode 100644 index 0000000000..a3d4f11775 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyMeshToGrid.h @@ -0,0 +1,19 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYMESHTOGRID_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYMESHTOGRID_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +// Bind nanovdb::tools::cuda::MeshToGrid: rasterize a triangle mesh (device +// vertex + triangle-index arrays) into a device ValueOnIndex GridHandle plus a +// per-value unsigned-distance-field sidecar buffer. Registered under `name`. +void defineMeshToGrid(nb::module_& m, const char* name); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyPointsToGrid.cu b/nanovdb/nanovdb/python/cuda/PyPointsToGrid.cu index 7ec30e4df3..1265ccd795 100644 --- a/nanovdb/nanovdb/python/cuda/PyPointsToGrid.cu +++ b/nanovdb/nanovdb/python/cuda/PyPointsToGrid.cu @@ -1,9 +1,12 @@ // Copyright Contributors to the OpenVDB Project // SPDX-License-Identifier: Apache-2.0 #include "PyPointsToGrid.h" +#include "PyValidate.h" #include +#include + #include namespace nb = nanobind; @@ -11,6 +14,9 @@ using namespace nb::literals; namespace pynanovdb { +// Fancy pointer adapter over an (N, 3) int32 device tensor that dereferences to +// a nanovdb::Coord, i.e. INDEX-space voxel coordinates. Used by the +// voxelsToGrid path (BuildT != Point, plus the Point/index build types). class NdArrayCoordPtr { const int32_t* data; @@ -40,20 +46,112 @@ public: } }; +// Fancy pointer adapter over an (N, 3) float/double device tensor that +// dereferences to a nanovdb::math::Vec3, i.e. WORLD-space point positions. +// Used by the pointsToGrid path which builds NanoGrid. The world +// pointsToGrid requires the dereferenced element to be Vec3f or Vec3d. +template class NdArrayVec3Ptr +{ + const ScalarT* data; + int64_t stride0, stride1; + +public: + using element_type = nanovdb::math::Vec3; + + __hostdev__ NdArrayVec3Ptr(const ScalarT* data, int64_t stride0, int64_t stride1) + : data(data) + , stride0(stride0) + , stride1(stride1) + { + } + __hostdev__ inline element_type operator[](size_t i) const + { + return element_type(data[i * stride0 + 0 * stride1], + data[i * stride0 + 1 * stride1], + data[i * stride0 + 2 * stride1]); + } + __hostdev__ inline element_type operator*() const + { + return element_type(data[0 * stride1], data[1 * stride1], data[2 * stride1]); + } +}; + +template void defineVoxelsToGrid(nb::module_& m, const char* name) +{ + m.def( + name, + [](nb::ndarray, nb::c_contig, nb::device::cuda> tensor, + double voxelSize, + uintptr_t stream) { + // Map::set only debug-asserts positivity; reject a singular / + // non-finite transform before any CUDA work. + requirePositiveFinite(voxelSize, "voxelsToGrid", "voxelSize"); + cudaStream_t s = reinterpret_cast(stream); + NdArrayCoordPtr points(tensor.data(), tensor.stride(0), tensor.stride(1)); + const size_t count = tensor.shape(0); + // voxelsToGrid only builds the grid topology (no blind data); pure + // CUDA touching no Python objects, so release the GIL. + nb::gil_scoped_release release; + return nanovdb::tools::cuda::voxelsToGrid( + points, count, voxelSize, nanovdb::cuda::DeviceBuffer(), s); + }, + "tensor"_a, + "voxelSize"_a = 1.0, + "stream"_a = 0, + "Rasterize the given (N, 3) int32 device tensor of index-space voxel " + "coordinates into a fresh device GridHandle. voxelSize is the world " + "size of a voxel; stream is a raw CUDA stream handle (Python int; " + "0 = default stream)."); +} + template void definePointsToGrid(nb::module_& m, const char* name) { m.def( name, - [](nb::ndarray, nb::c_contig, nb::device::cuda> tensor) { - NdArrayCoordPtr points(tensor.data(), tensor.stride(0), tensor.stride(1)); - nanovdb::tools::cuda::PointsToGrid converter; - auto handle = converter.getHandle(points, tensor.shape(0)); - return handle; + [](nb::ndarray, nb::c_contig, nb::device::cuda> tensor, + double voxelSize, + uintptr_t stream) { + // Map::set only debug-asserts positivity; reject a singular / + // non-finite transform before any CUDA work. + requirePositiveFinite(voxelSize, "pointsToGrid", "voxelSize"); + cudaStream_t s = reinterpret_cast(stream); + NdArrayVec3Ptr points(tensor.data(), tensor.stride(0), tensor.stride(1)); + const size_t count = tensor.shape(0); + // Build a NanoGrid at the requested fixed voxel size, + // encoding the world points as blind data. Mirror the converter + // usage of the legacy Rgba8 path (the fixed-voxelSize free function + // has no out-of-line definition in the header), but set the Point + // build type and the world->index map scale. + nb::gil_scoped_release release; + nanovdb::tools::cuda::PointsToGrid converter( + voxelSize, nanovdb::Vec3d(0.0), s); + converter.setPointType(nanovdb::PointType::Default); + return converter.getHandle(points, count); }, "tensor"_a, - "Rasterize the given (N, 3) int32 device tensor of points into a fresh GridHandle."); + "voxelSize"_a = 1.0, + "stream"_a = 0, + "Rasterize the given (N, 3) float or double device tensor of " + "WORLD-space point positions into a fresh device GridHandle of type " + "NanoGrid, encoding the point coordinates as blind data. " + "voxelSize is the world size of a voxel; stream is a raw CUDA stream " + "handle (Python int; 0 = default stream)."); } -template void definePointsToGrid(nb::module_&, const char*); +// voxelsToGrid (index-space int32 Coord input). The legacy Rgba8 binding maps +// onto defineVoxelsToGrid; additional build types build matching grids. +// +// NOTE: voxelsToGrid is intentionally NOT instantiated. The C++ +// PointsToGrid::countNodes static_asserts that Point coordinates be +// Vec3f or Vec3d (PointsToGrid.cuh:590), so the int32-Coord voxelsToGrid input +// is rejected at compile time for BuildT == Point. Build Point grids from +// world-space float/double positions via definePointsToGrid. +template void defineVoxelsToGrid(nb::module_&, const char*); +template void defineVoxelsToGrid(nb::module_&, const char*); +template void defineVoxelsToGrid(nb::module_&, const char*); + +// pointsToGrid (world-space float/double input -> NanoGrid). +template void definePointsToGrid(nb::module_&, const char*); +template void definePointsToGrid(nb::module_&, const char*); } // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyPointsToGrid.h b/nanovdb/nanovdb/python/cuda/PyPointsToGrid.h index 4af07f5657..166ede420f 100644 --- a/nanovdb/nanovdb/python/cuda/PyPointsToGrid.h +++ b/nanovdb/nanovdb/python/cuda/PyPointsToGrid.h @@ -9,6 +9,15 @@ namespace nb = nanobind; namespace pynanovdb { +// Rasterize an (N, 3) int32 device tensor of voxel (index-space) coordinates +// into a fresh device GridHandle of type NanoGrid. Used for the +// legacy pointsToRGBA8Grid binding and the new voxelsTo*Grid bindings. +template void defineVoxelsToGrid(nb::module_& m, const char* name); + +// Rasterize an (N, 3) float OR double device tensor of WORLD-space point +// positions into a fresh device GridHandle of type NanoGrid, encoding +// the point coordinates as blind data. BuildT selects the world coordinate +// scalar type (float or double). template void definePointsToGrid(nb::module_& m, const char* name); } // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyPruneGrid.cu b/nanovdb/nanovdb/python/cuda/PyPruneGrid.cu new file mode 100644 index 0000000000..e1ac8c631e --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyPruneGrid.cu @@ -0,0 +1,75 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#include "PyPruneGrid.h" + +#include + +#include +#include +#include + +#include +#include // leaf count of a device grid + +namespace nb = nanobind; +using namespace nb::literals; +// NOTE: deliberately NOT `using namespace nanovdb;`. These tools instantiate +// CUB DeviceScan, whose nvcc-generated host stub references unqualified +// `cuda::std::...`; with `nanovdb::cuda` in scope that becomes ambiguous and +// fails to compile. Fully qualify nanovdb:: instead (matches PyPointsToGrid.cu). + +namespace pynanovdb { + +template void definePruneGrid(nb::module_& m, const char* name) +{ + m.def( + name, + [](nanovdb::NanoGrid* dGrid, + nb::ndarray leafMask, + uintptr_t stream) { + // The sidecar is a device array of nanovdb::Mask<3> (one 512-bit / + // 8^3 leaf mask per leaf node, voxels to RETAIN), passed as raw + // uint64 words. Mask<3> is exactly WORD_COUNT 64-bit words, so the + // total word count must be a whole multiple of that, and should + // equal (leaf count) * Mask<3>::WORD_COUNT for a well-formed call. + constexpr size_t wordsPerMask = nanovdb::Mask<3>::WORD_COUNT; + const size_t totalWords = leafMask.size(); + if (totalWords % wordsPerMask != 0) + throw std::invalid_argument( + "pruneGrid: leafMask uint64 word count must be a multiple of " + "Mask<3>::WORD_COUNT (8); supply one 512-bit mask per leaf node"); + // PruneGrid indexes one Mask<3> per SOURCE leaf, so the sidecar + // must cover every leaf; a shorter multiple of 8 would pass the + // divisibility check but read out of bounds on the device. + using Traits = nanovdb::util::cuda::DeviceGridTraits; + const uint64_t leafCount = Traits::getTreeData(dGrid).mNodeCount[0]; + if (totalWords < leafCount * wordsPerMask) + throw nb::value_error( + ("pruneGrid: leafMask holds " + + std::to_string(totalWords / wordsPerMask) + + " masks but the grid has " + std::to_string(leafCount) + + " leaf nodes; supply one Mask<3> (8 uint64) per leaf.") + .c_str()); + cudaStream_t s = reinterpret_cast(stream); + const nanovdb::Mask<3>* d_mask = + reinterpret_cast*>(leafMask.data()); + // PruneGrid::getHandle launches kernels and synchronizes the + // stream; pure CUDA touching no Python objects, so release the GIL. + nb::gil_scoped_release release; + nanovdb::tools::cuda::PruneGrid pruner(dGrid, d_mask, s); + return pruner.getHandle(); + }, + "dGrid"_a, + "leafMask"_a, + "stream"_a = 0, + "Morphologically prune a device OnIndex grid against a per-leaf retain " + "mask and return a fresh device GridHandle of the pruned grid. leafMask " + "is a device uint64 array holding one nanovdb::Mask<3> (8 x uint64 = " + "512 bits) per leaf node, in leaf order, marking voxels to RETAIN; its " + "length must be (leaf count) * 8. stream is a raw CUDA stream handle " + "(Python int; 0 = default stream)."); +} + +template void definePruneGrid(nb::module_&, const char*); + +} // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyPruneGrid.h b/nanovdb/nanovdb/python/cuda/PyPruneGrid.h new file mode 100644 index 0000000000..c015958951 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyPruneGrid.h @@ -0,0 +1,16 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYPRUNEGRID_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYPRUNEGRID_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +template void definePruneGrid(nb::module_& m, const char* name); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyRefineGrid.cu b/nanovdb/nanovdb/python/cuda/PyRefineGrid.cu new file mode 100644 index 0000000000..c53fbae2b3 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyRefineGrid.cu @@ -0,0 +1,39 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#include "PyRefineGrid.h" + +#include + +#include + +namespace nb = nanobind; +using namespace nb::literals; +// NOTE: deliberately NOT `using namespace nanovdb;`. These tools instantiate +// CUB DeviceScan, whose nvcc-generated host stub references unqualified +// `cuda::std::...`; with `nanovdb::cuda` in scope that becomes ambiguous and +// fails to compile. Fully qualify nanovdb:: instead (matches PyPointsToGrid.cu). + +namespace pynanovdb { + +template void defineRefineGrid(nb::module_& m, const char* name) +{ + m.def( + name, + [](nanovdb::NanoGrid* dGrid, uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + // RefineGrid::getHandle launches kernels and synchronizes the + // stream; pure CUDA touching no Python objects, so release the GIL. + nb::gil_scoped_release release; + nanovdb::tools::cuda::RefineGrid refiner(dGrid, s); + return refiner.getHandle(); + }, + "dGrid"_a, + "stream"_a = 0, + "Topologically refine (2x upsample) a device OnIndex grid and return a " + "fresh device GridHandle of the refined grid. stream is a raw CUDA " + "stream handle (Python int; 0 = default stream)."); +} + +template void defineRefineGrid(nb::module_&, const char*); + +} // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/cuda/PyRefineGrid.h b/nanovdb/nanovdb/python/cuda/PyRefineGrid.h new file mode 100644 index 0000000000..9af34dcc1c --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyRefineGrid.h @@ -0,0 +1,16 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYREFINEGRID_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYREFINEGRID_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +template void defineRefineGrid(nb::module_& m, const char* name); + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PySampleFromVoxels.cu b/nanovdb/nanovdb/python/cuda/PySampleFromVoxels.cu index aaebf66772..9c13cdcff8 100644 --- a/nanovdb/nanovdb/python/cuda/PySampleFromVoxels.cu +++ b/nanovdb/nanovdb/python/cuda/PySampleFromVoxels.cu @@ -4,6 +4,8 @@ #include +#include + #include #include @@ -13,34 +15,34 @@ using namespace nanovdb; namespace { -template __global__ void sampleFromVoxels(unsigned int numPoints, const BuildT* points, const NanoGrid* d_grid, BuildT* values) +template __global__ void sampleFromVoxels(unsigned int numPoints, const BuildT* points, const NanoGrid* dGrid, BuildT* values) { using TreeT = NanoTree; using Vec3T = math::Vec3; for (unsigned int i = threadIdx.x + blockIdx.x * blockDim.x; i < numPoints; i += blockDim.x * gridDim.x) { Vec3T worldPos(points[3 * i], points[3 * i + 1], points[3 * i + 2]); - Vec3T indexPos = d_grid->worldToIndex(worldPos); + Vec3T indexPos = dGrid->worldToIndex(worldPos); - math::SampleFromVoxels sampler(d_grid->tree()); + math::SampleFromVoxels sampler(dGrid->tree()); values[i] = sampler(indexPos); } } template -__global__ void sampleFromVoxels(unsigned int numPoints, const BuildT* points, const NanoGrid* d_grid, BuildT* values, BuildT* gradients) +__global__ void sampleFromVoxels(unsigned int numPoints, const BuildT* points, const NanoGrid* dGrid, BuildT* values, BuildT* gradients) { using TreeT = NanoTree; using Vec3T = math::Vec3; for (unsigned int i = threadIdx.x + blockIdx.x * blockDim.x; i < numPoints; i += blockDim.x * gridDim.x) { Vec3T worldPos(points[3 * i], points[3 * i + 1], points[3 * i + 2]); - Vec3T indexPos = d_grid->worldToIndex(worldPos); + Vec3T indexPos = dGrid->worldToIndex(worldPos); - math::SampleFromVoxels sampler(d_grid->tree()); + math::SampleFromVoxels sampler(dGrid->tree()); values[i] = sampler(indexPos); - Vec3T inv2Dx = (BuildT).5 / d_grid->voxelSize(); + Vec3T inv2Dx = (BuildT).5 / dGrid->voxelSize(); Vec3T gradient = Vec3T(sampler(indexPos + Vec3T(1, 0, 0)) - sampler(indexPos - Vec3T(1, 0, 0)), sampler(indexPos + Vec3T(0, 1, 0)) - sampler(indexPos - Vec3T(0, 1, 0)), sampler(indexPos + Vec3T(0, 0, 1)) - sampler(indexPos - Vec3T(0, 0, 1))) * @@ -60,29 +62,41 @@ template void defineSampleFromVoxels(nb::module_& m, const char m.def( name, [](nb::ndarray, nb::c_contig, nb::device::cuda> points, - NanoGrid* d_grid, - nb::ndarray, nb::device::cuda> values) { + NanoGrid* dGrid, + nb::ndarray, nb::device::cuda> values, + uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); constexpr unsigned int numThreads = 128; unsigned int numBlocks = (points.shape(0) + numThreads - 1) / numThreads; - sampleFromVoxels<<>>(points.shape(0), points.data(), d_grid, values.data()); + // Raw kernel launch on the supplied stream; touches no Python + // objects, so release the GIL. + nb::gil_scoped_release release; + sampleFromVoxels<<>>(points.shape(0), points.data(), dGrid, values.data()); }, "points"_a, - "d_grid"_a, - "values"_a); + "dGrid"_a, + "values"_a, + "stream"_a = 0); m.def( name, [](nb::ndarray, nb::c_contig, nb::device::cuda> points, - NanoGrid* d_grid, + NanoGrid* dGrid, nb::ndarray, nb::device::cuda> values, - nb::ndarray, nb::device::cuda> gradients) { + nb::ndarray, nb::device::cuda> gradients, + uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); constexpr unsigned int numThreads = 128; unsigned int numBlocks = (points.shape(0) + numThreads - 1) / numThreads; - sampleFromVoxels<<>>(points.shape(0), points.data(), d_grid, values.data(), gradients.data()); + // Raw kernel launch on the supplied stream; touches no Python + // objects, so release the GIL. + nb::gil_scoped_release release; + sampleFromVoxels<<>>(points.shape(0), points.data(), dGrid, values.data(), gradients.data()); }, "points"_a, - "d_grid"_a, + "dGrid"_a, "values"_a, - "gradients"_a); + "gradients"_a, + "stream"_a = 0); } template void defineSampleFromVoxels(nb::module_&, const char*); diff --git a/nanovdb/nanovdb/python/cuda/PySignedFloodFill.cu b/nanovdb/nanovdb/python/cuda/PySignedFloodFill.cu index ddbd5c67cb..16efa1b583 100644 --- a/nanovdb/nanovdb/python/cuda/PySignedFloodFill.cu +++ b/nanovdb/nanovdb/python/cuda/PySignedFloodFill.cu @@ -2,6 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 #include "PySignedFloodFill.h" +#include + #include namespace nb = nanobind; @@ -13,7 +15,19 @@ namespace pynanovdb { template void defineSignedFloodFill(nb::module_& m, const char* name) { m.def( - name, [](NanoGrid* d_grid, bool verbose) { return tools::cuda::signedFloodFill(d_grid, verbose); }, "d_grid"_a, "verbose"_a = false); + name, + [](NanoGrid* dGrid, bool verbose, uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + // signedFloodFill launches kernels and synchronizes the stream; + // pure CUDA touching no Python objects, so release the GIL. + nb::gil_scoped_release release; + tools::cuda::signedFloodFill(dGrid, verbose, s); + }, + "dGrid"_a, + "verbose"_a = false, + "stream"_a = 0, + "Perform a signed flood fill on a device float/double grid in place. " + "stream is a raw CUDA stream handle (Python int; 0 = default stream)."); } template void defineSignedFloodFill(nb::module_&, const char*); diff --git a/nanovdb/nanovdb/python/cuda/PyTempPool.cu b/nanovdb/nanovdb/python/cuda/PyTempPool.cu new file mode 100644 index 0000000000..e193ab399b --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyTempPool.cu @@ -0,0 +1,106 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifdef NANOVDB_USE_CUDA + +#include "PyTempPool.h" + +#include + +#include + +#include +#include + +namespace nb = nanobind; +using namespace nb::literals; + +namespace pynanovdb { + +void defineTempPool(nb::module_& m) +{ + using DeviceResource = nanovdb::cuda::DeviceResource; + using TempDevicePool = nanovdb::cuda::TempDevicePool; + + // DeviceResource: a stateless async allocator over the current CUDA device. + // The C++ side is instance-based (the static allocateAsync / deallocateAsync + // are deprecated), but since it carries no state the binding stays static on + // the Python side. Raw pointers are Python ints. + nb::class_(m, "DeviceResource", + "Stateless CUDA async allocator: allocateAsync / deallocateAsync issue " + "cudaMallocAsync / cudaFreeAsync on a stream. Pointers are raw Python " + "ints. Backs TempDevicePool.") + .def_ro_static("DEFAULT_ALIGNMENT", &DeviceResource::DEFAULT_ALIGNMENT, + "Default allocation alignment in bytes (256).") + .def_static( + "allocateAsync", + [](size_t bytes, size_t alignment, uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + nb::gil_scoped_release release; + DeviceResource resource; + void* p = resource.allocate_async(bytes, alignment, s); + return reinterpret_cast(p); + }, + "bytes"_a, "alignment"_a = DeviceResource::DEFAULT_ALIGNMENT, "stream"_a = 0, + "Asynchronously allocate `bytes` of device memory on `stream` and " + "return the raw device pointer as a Python int. The alignment " + "argument is accepted for API parity but ignored by cudaMallocAsync.") + .def_static( + "deallocateAsync", + [](uintptr_t ptr, size_t bytes, size_t alignment, uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + nb::gil_scoped_release release; + DeviceResource resource; + resource.deallocate_async(reinterpret_cast(ptr), + bytes, alignment, s); + }, + "ptr"_a, "bytes"_a = 0, "alignment"_a = DeviceResource::DEFAULT_ALIGNMENT, + "stream"_a = 0, + "Asynchronously free a device pointer (Python int) on `stream`. The " + "bytes / alignment arguments are accepted for API parity but ignored " + "by cudaFreeAsync."); + + // TempDevicePool: a thin pool for CUB temporary storage. It owns a raw + // device pointer, so it is intentionally non-copyable (the default copy + // would double-free). Exposed in its current shape only. + nb::class_(m, "TempDevicePool", + "Thin pool of CUB temporary device storage backed by DeviceResource. " + "Owns a raw device pointer (non-copyable). reallocate(stream) grows it " + "to requestedSize when needed.") + .def(nb::init<>(), + "Construct an empty pool (no device allocation yet).") + .def( + "data", + [](TempDevicePool& pool) { + return reinterpret_cast(pool.data()); + }, + "Raw device pointer to the pooled storage as a Python int " + "(0 if nothing has been allocated yet).") + .def( + "size", + [](TempDevicePool& pool) { return pool.size(); }, + "Currently allocated size of the pool in bytes.") + .def( + "requestedSize", + [](TempDevicePool& pool) { return pool.requestedSize(); }, + "Size in bytes most recently requested for the pool. reallocate() " + "grows the allocation to this value when it exceeds size().") + .def( + "setRequestedSize", + [](TempDevicePool& pool, size_t value) { pool.requestedSize() = value; }, + "value"_a, + "Set the requested size (in bytes) used by the next reallocate().") + .def( + "reallocate", + [](TempDevicePool& pool, uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + nb::gil_scoped_release release; + pool.reallocate(s); + }, + "stream"_a = 0, + "Reallocate the pool on `stream` (a raw CUDA stream handle, Python " + "int) if it is empty or requestedSize() exceeds the current size()."); +} + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyTempPool.h b/nanovdb/nanovdb/python/cuda/PyTempPool.h new file mode 100644 index 0000000000..6522a8c34e --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyTempPool.h @@ -0,0 +1,23 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYTEMPPOOL_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYTEMPPOOL_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +#ifdef NANOVDB_USE_CUDA +/// @brief Register nanovdb::cuda::DeviceResource (static async allocator) and +/// nanovdb::cuda::TempDevicePool (= TempPool) on the +/// nanovdb.cuda submodule, in their current shape only. These are the +/// low-level temp-storage primitives used by the CUDA tools; they are +/// NOT reshaped into any resource-concept abstraction. +void defineTempPool(nb::module_& m); +#endif + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyUnifiedBuffer.cu b/nanovdb/nanovdb/python/cuda/PyUnifiedBuffer.cu new file mode 100644 index 0000000000..a6bb8ab8a7 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyUnifiedBuffer.cu @@ -0,0 +1,152 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifdef NANOVDB_USE_CUDA + +#include "PyUnifiedBuffer.h" +#include "PyDeviceBuffer.h" // for defineDeviceBufferLike (shared interop) + +#include + +#include + +#include + +namespace nb = nanobind; +using namespace nb::literals; + +namespace pynanovdb { + +void defineUnifiedBuffer(nb::module_& m) +{ + using BufferT = nanovdb::cuda::UnifiedBuffer; + + // defineDeviceBufferLike gives size() / device_ptr / host_ptr / + // __cuda_array_interface__ / __dlpack_device__ / __dlpack__ for free. For + // UnifiedBuffer the managed pointer is valid on host AND device, so + // host_ptr and device_ptr report the SAME address. + defineDeviceBufferLike(m, "UnifiedBuffer") + .def( + "__init__", + [](BufferT* self, size_t size) { + // Allocate a managed page table of `size` bytes (size == capacity). + nb::gil_scoped_release release; + new (self) BufferT(size); + }, + "size"_a, + "Construct a UnifiedBuffer backed by `size` bytes of CUDA managed " + "(unified) memory. The same pointer is valid on the host and on " + "every device.") + .def( + "__init__", + [](BufferT* self, size_t size, int device, uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + // The (size, device, stream) ctor allocates and applies a + // preferred-location advise + prefetch to `device`. + nb::gil_scoped_release release; + new (self) BufferT(static_cast(size), device, s); + }, + "size"_a, + "device"_a, + "stream"_a = 0, + "Construct a UnifiedBuffer of `size` bytes and set the preferred " + "location advise plus an async prefetch to `device`. stream is a " + "raw CUDA stream handle (Python int; 0 = default stream).") + .def("capacity", &BufferT::capacity, + "Number of bytes reserved in the managed page table (room for " + "growth; may exceed size()).") + .def("empty", &BufferT::empty, + "True iff this buffer manages no memory.") + .def("isEmpty", &BufferT::isEmpty, + "Same as empty(). Retained for parity with the C++ " + "UnifiedBuffer::isEmpty() member.") + .def( + "clear", + [](BufferT& buf) { + nb::gil_scoped_release release; + buf.clear(); + }, + "Free all managed memory and reset this buffer to empty.") + .def( + "resize", + [](BufferT& buf, size_t size, int device) { + nb::gil_scoped_release release; + buf.resize(size, device); + }, + "size"_a, "device"_a = cudaCpuDeviceId, + "Resize the managed memory block. If the new size fits inside the " + "current capacity this only redefines size(); otherwise a new page " + "table is allocated (with a preferred-location advise on `device`) " + "and the old contents are copied over. device defaults to the host " + "(cudaCpuDeviceId == -1).") + .def( + "advise", + [](const BufferT& buf, ptrdiff_t byteOffset, size_t size, int device, + int adv) { + nb::gil_scoped_release release; + buf.advise(byteOffset, size, device, + static_cast(adv)); + }, + "byteOffset"_a, "size"_a, "device"_a, "advise"_a, + "Apply a single cudaMemoryAdvise (passed as its integer enumerator, " + "e.g. cudaMemAdviseSetPreferredLocation == 3) to the [byteOffset, " + "byteOffset + size) range for `device` (cudaCpuDeviceId == -1 " + "selects the host).") + .def( + "prefetch", + [](const BufferT& buf, ptrdiff_t byteOffset, size_t size, int device, + uintptr_t stream) { + cudaStream_t s = reinterpret_cast(stream); + nb::gil_scoped_release release; + buf.prefetch(byteOffset, size, device, s); + }, + "byteOffset"_a = 0, "size"_a = 0, "device"_a = cudaCpuDeviceId, + "stream"_a = 0, + "Prefetch the [byteOffset, byteOffset + size) range to `device` " + "(cudaCpuDeviceId == -1 selects the host). size == 0 prefetches all " + "size() bytes. stream is a raw CUDA stream handle (Python int).") + .def( + "deviceUpload", + [](const BufferT& buf, int device, uintptr_t stream, bool sync) { + cudaStream_t s = reinterpret_cast(stream); + nb::gil_scoped_release release; + buf.deviceUpload(device, s, sync); + }, + "device"_a = 0, "stream"_a = 0, "sync"_a = false, + "Prefetch all managed bytes to `device` (legacy DeviceBuffer-compat; " + "internally a memPrefetchAsync). stream is a raw CUDA stream handle " + "(Python int); if sync is True the call blocks until the prefetch " + "completes.") + .def( + "deviceDownload", + [](const BufferT& buf, uintptr_t stream, bool sync) { + cudaStream_t s = reinterpret_cast(stream); + nb::gil_scoped_release release; + buf.deviceDownload(s, sync); + }, + "stream"_a = 0, "sync"_a = false, + "Prefetch all managed bytes back to the host (legacy " + "DeviceBuffer-compat). stream is a raw CUDA stream handle (Python " + "int); if sync is True the call blocks until the prefetch " + "completes.") + .def_static( + "create", + [](size_t size) { + nb::gil_scoped_release release; + return BufferT::create(size); + }, + "size"_a, + "Create a UnifiedBuffer of `size` bytes (size == capacity).") + .def_static( + "create", + [](size_t size, size_t capacity) { + nb::gil_scoped_release release; + return BufferT::create(size, capacity); + }, + "size"_a, "capacity"_a, + "Create a UnifiedBuffer with `size` used bytes and a managed page " + "table of `capacity` bytes for future growth."); +} + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyUnifiedBuffer.h b/nanovdb/nanovdb/python/cuda/PyUnifiedBuffer.h new file mode 100644 index 0000000000..2e0eafec4f --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyUnifiedBuffer.h @@ -0,0 +1,24 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYUNIFIEDBUFFER_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYUNIFIEDBUFFER_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +#ifdef NANOVDB_USE_CUDA +/// @brief Register nanovdb::cuda::UnifiedBuffer as "UnifiedBuffer" on the +/// nanovdb.cuda submodule. UnifiedBuffer is CUDA Unified (managed) +/// memory: a single pointer valid on the host AND every device, so it +/// gets the shared device-interop surface (CAI / DLPack / device_ptr / +/// host_ptr / size) for free — note host_ptr and device_ptr are the +/// same managed pointer. +void defineUnifiedBuffer(nb::module_& m); +#endif + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyUnifiedGridHandle.cu b/nanovdb/nanovdb/python/cuda/PyUnifiedGridHandle.cu new file mode 100644 index 0000000000..93402974cd --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyUnifiedGridHandle.cu @@ -0,0 +1,103 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifdef NANOVDB_USE_CUDA + +#include "PyUnifiedGridHandle.h" +#include "PyGridHandle.h" + +#include + +#include + +#include +#include + +namespace nb = nanobind; +using namespace nb::literals; + +namespace pynanovdb { + +// Polymorphic deviceGrid(n) for a unified GridHandle — same dispatch shape as +// pyHostGrid / the DeviceBuffer pyDeviceGrid, returning the grid via +// the unified (managed) pointer. With UnifiedBuffer host and device pointers +// coincide, so this is valid as soon as the handle holds a grid. +static nb::object pyUnifiedDeviceGrid(nb::handle py_handle, uint32_t n) +{ + using BufferT = nanovdb::cuda::UnifiedBuffer; + auto& handle = nb::cast&>(py_handle); + if (n >= handle.gridCount()) return nb::none(); + switch (handle.gridType(n)) { +#define NANOVDB_PY_FOR_EACH_SCALAR_BUILDT(T, Suffix, GridTypeEnum) \ + case nanovdb::GridType::GridTypeEnum: { \ + auto* grid = handle.template deviceGrid(n); \ + return grid ? nb::cast(grid, nb::rv_policy::reference, py_handle) \ + : nb::none(); \ + } +#define NANOVDB_PY_FOR_EACH_VECTOR_BUILDT(T, Suffix, AccessorName, GridTypeEnum) \ + case nanovdb::GridType::GridTypeEnum: { \ + auto* grid = handle.template deviceGrid(n); \ + return grid ? nb::cast(grid, nb::rv_policy::reference, py_handle) \ + : nb::none(); \ + } +#define NANOVDB_PY_FOR_EACH_POINT_BUILDT(T, Suffix, GridTypeEnum) \ + case nanovdb::GridType::GridTypeEnum: { \ + auto* grid = handle.template deviceGrid(n); \ + return grid ? nb::cast(grid, nb::rv_policy::reference, py_handle) \ + : nb::none(); \ + } +#define NANOVDB_PY_FOR_EACH_READONLY_BUILDT(T, Suffix, GridTypeEnum) \ + case nanovdb::GridType::GridTypeEnum: { \ + auto* grid = handle.template deviceGrid(n); \ + return grid ? nb::cast(grid, nb::rv_policy::reference, py_handle) \ + : nb::none(); \ + } +#include "BuildTypes.def" + default: + return nb::none(); + } +} + +void defineUnifiedGridHandle(nb::module_& m) +{ + using BufferT = nanovdb::cuda::UnifiedBuffer; + defineGridHandle(m, "UnifiedGridHandle") + .def("deviceGrid", &pyUnifiedDeviceGrid, "n"_a = 0, + nb::keep_alive<0, 1>(), + "Return the n-th grid as a typed Grid subclass selected by " + "gridType(n), accessed through the unified (managed) pointer, or " + "None if the BuildT is not bound in Python. The returned grid " + "keeps this handle alive.") + .def( + "deviceUpload", + [](nanovdb::GridHandle& handle, uintptr_t stream, bool sync) { + cudaStream_t s = reinterpret_cast(stream); + nb::gil_scoped_release release; + handle.deviceUpload(reinterpret_cast(s), sync); + }, + "stream"_a = 0, "sync"_a = true, + "Prefetch the unified buffer to the current device. stream is a raw " + "CUDA stream handle (Python int; 0 = default stream). If sync is " + "True the call blocks until the prefetch completes.") + .def( + "deviceDownload", + [](nanovdb::GridHandle& handle, uintptr_t stream, bool sync) { + cudaStream_t s = reinterpret_cast(stream); + nb::gil_scoped_release release; + handle.deviceDownload(reinterpret_cast(s), sync); + }, + "stream"_a = 0, "sync"_a = true, + "Prefetch the unified buffer back to the host. stream is a raw CUDA " + "stream handle (Python int; 0 = default stream). If sync is True " + "the call blocks until the prefetch completes.") + .def( + "device_ptr", + [](nanovdb::GridHandle& handle) { + return reinterpret_cast(handle.buffer().deviceData()); + }, + "Raw device (managed) pointer to the base of the whole buffer as a " + "Python int. For unified memory this equals the host pointer."); +} + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/cuda/PyUnifiedGridHandle.h b/nanovdb/nanovdb/python/cuda/PyUnifiedGridHandle.h new file mode 100644 index 0000000000..ec5a669616 --- /dev/null +++ b/nanovdb/nanovdb/python/cuda/PyUnifiedGridHandle.h @@ -0,0 +1,23 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +#ifndef NANOVDB_CUDA_PYUNIFIEDGRIDHANDLE_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_PYUNIFIEDGRIDHANDLE_HAS_BEEN_INCLUDED + +#include + +namespace nb = nanobind; + +namespace pynanovdb { + +#ifdef NANOVDB_USE_CUDA +/// @brief Register GridHandle as +/// "UnifiedGridHandle" on the nanovdb.cuda submodule. This is the handle +/// type returned by nanovdb.tools.cuda.DistributedPointsToGrid.getHandle +/// (its default BufferT is UnifiedBuffer), so the class must be +/// registered for nanobind to cast the result. +void defineUnifiedGridHandle(nb::module_& m); +#endif + +} // namespace pynanovdb + +#endif diff --git a/nanovdb/nanovdb/python/examples/README.md b/nanovdb/nanovdb/python/examples/README.md index 6cf82117ec..d1eab6d576 100644 --- a/nanovdb/nanovdb/python/examples/README.md +++ b/nanovdb/nanovdb/python/examples/README.md @@ -43,6 +43,230 @@ smoke-tested by [`../test/TestExamples.py`](../test/TestExamples.py) (ctest name `pytest_nanovdb_examples`) when the module is configured with `NANOVDB_BUILD_PYTHON_UNITTESTS=ON`. +### GPU examples + +These require a CUDA build of `nanovdb` and a CUDA-capable GPU; each +self-skips with a printed message when the build, GPU, or the optional +GPU-array framework it uses is unavailable. + +| Script | What it shows | +| --- | --- | +| [`gpu_load_inspect.py`](gpu_load_inspect.py) | `nanovdb.io.deviceReadGrid` → `handle.deviceUpload()` → zero-copy `cupy.asarray(handle)` view of the device buffer; the host `grid(n)` vs device `deviceGrid(n)` `data_ptr()` split. Requires CuPy. | +| [`cupy_rawkernel.py`](cupy_rawkernel.py) | A `cupy.RawKernel` that `#include ` (compiled with `nanovdb.cuda.compile_options()`) and reads a `const nanovdb::NanoGrid*` straight from `grid.data_ptr()`. Documents the device-pointer ABI. Requires CuPy. | +| [`numba_cuda.py`](numba_cuda.py) | Adopting a NanoVDB device buffer as a zero-copy `numba.cuda` array (Numba can't parse the C++ ABI, so it operates on the raw buffer). Requires Numba. | +| [`triton_kernel.py`](triton_kernel.py) | Handing a NanoVDB device buffer to a Triton kernel via a zero-copy `torch.from_dlpack` tensor. Requires Triton + PyTorch. | +| [`voxels_to_grid_cuda.py`](voxels_to_grid_cuda.py) | Build a grid on the GPU from a CuPy coordinate array with `tools.cuda.voxelsToOnIndexGrid` and `tools.cuda.pointsToGrid` — construction entirely on the device (no host tree). GPU counterpart to `build_grid.py`; port of `ex_voxels_to_grid_cuda`. Requires CuPy. | +| [`device_topology_ops.py`](device_topology_ops.py) | `tools.cuda.dilateGrid` / `coarsenGrid` / `refineGrid` / `mergeGrids` / `pruneGrid` on a device OnIndex grid. Ports the device path of `ex_{dilate,coarsen,refine,merge}_nanovdb_cuda` (OpenVDB-free: the source grid is built on the device). Requires CuPy. | +| [`sample_from_voxels_cuda.py`](sample_from_voxels_cuda.py) | `tools.cuda.sampleFromVoxels` — batched trilinear samples plus analytic gradients at device query points; the device analog of `createTrilinearSampler`, cross-checked against the host sampler. Requires CuPy. | +| [`index_to_grid_cuda.py`](index_to_grid_cuda.py) | `tools.cuda.indexToGrid` bakes a CuPy-computed value array onto an index grid on the device, with `tools.cuda.activeVoxelCoords` recovering each value slot's coordinate. Device half of `ex_index_grid_cuda`. Requires CuPy. | +| [`signed_flood_fill_cuda.py`](signed_flood_fill_cuda.py) | `tools.cuda.signedFloodFill` propagates interior/exterior sign of a level set in place on the device; verified with `tools.cuda.sampleFromVoxels`. Requires CuPy. | +| [`validate_cuda.py`](validate_cuda.py) | Device quality control: `tools.cuda.isValid`, `evalChecksum` / `updateChecksum` / `validateChecksum`, and `updateGridStats`, all in place on a device grid. GPU counterpart to `validate.py`. Requires CuPy. | +| [`collide_level_set_cuda.py`](collide_level_set_cuda.py) | GPU counterpart to `collide_level_set.py`: a whole particle set collided against a level set via one `tools.cuda.sampleFromVoxels` call (distances + gradients), with the reflection response as vectorized CuPy. Port of the device path of `ex_collide_level_set`. Requires CuPy. | +| [`raytrace_level_set_cuda.py`](raytrace_level_set_cuda.py) | GPU counterpart to `raytrace_level_set.py`: one CUDA thread per pixel in a `cupy.RawKernel`, using the faithful `nanovdb::math::Ray` + HDDA `ZeroCrossing` surface search (bound only in C++). Port of the device path of `ex_raytrace_level_set`. Requires CuPy + NanoVDB headers. | +| [`raytrace_fog_volume_cuda.py`](raytrace_fog_volume_cuda.py) | GPU counterpart to `raytrace_fog_volume.py`: per-pixel transmittance ray-march in a `cupy.RawKernel` over the device grid accessor. Port of the device path of `ex_raytrace_fog_volume`. Requires CuPy + NanoVDB headers. | +| [`mesh_to_grid_cuda.py`](mesh_to_grid_cuda.py) | `tools.cuda.meshToGrid` rasterizes a triangle mesh into a narrow-band unsigned distance field on the device (returns an OnIndex grid + per-value UDF sidecar), then `indexToGrid` bakes it into a Float distance grid. Port of `ex_mesh_to_grid_cuda` (OpenVDB-free: the mesh is generated with NumPy). Requires CuPy. | + +#### One level-set filter, three compute backends + +[`levelset_filter.py`](levelset_filter.py) is the runnable driver for **one** GPU +level-set filter — `tools::LevelSetFilter`-style Laplacian diffusion + Godunov +renormalisation + narrow-band retrack on a `.nvdb` file, driven by the device +`VoxelBlockManager`, reading/writing either a `FloatGrid` or an `OnIndexGrid` +(SDF in a blind channel) and preserving the input style — with the per-voxel +stencil math supplied by one of **three interchangeable backends**: + +``` +python levelset_filter.py {rawkernel|cupy|cutile} in.nvdb out.nvdb [outer_iters] +python levelset_filter.py {rawkernel|cupy|cutile} # self-test +``` + +The driver owns everything that isn't backend-specific (the style-detecting read, +the style-preserving write, the `dilateGrid` / `inject` / `injectPredicateToMask` +/ `pruneGrid` retrack, the outer loop). Each backend is a small `Backend` object +implementing only the three per-voxel stencils, and differs **only in how the +dense per-voxel math is computed** — a side-by-side of three GPU styles. Each +backend file also runs standalone (`python levelset_filter_.py`) as a +stencil-only smoke test (deform + renorm on a sphere, no I/O, no retrack). + +| Backend | Per-voxel compute | +| --- | --- | +| [`levelset_filter_rawkernel.py`](levelset_filter_rawkernel.py) | A hand-written CUDA kernel: `decodeInverseMaps` / `computeBoxStencil` fused with the update in a `cupy.RawModule` (nvcc backend), no dense gather. Fastest, most control. Requires CuPy + nvcc. | +| [`levelset_filter_cupy.py`](levelset_filter_cupy.py) | **Kernel-free**: `gatherBoxStencil` → dense `(N, 6)` face values, then all per-voxel math as plain **CuPy** array ops. No `RawModule` / CUDA C++ / nvcc. Requires only CuPy. | +| [`levelset_filter_cutile.py`](levelset_filter_cutile.py) | The same dense arrays, with the per-voxel stencils as **NVIDIA cuTile** (`cuda.tile`) tile kernels (`ct.load` / `ct.store` over `(TILE,)` tiles). Requires CuPy + `cuda-tile`. | + For full API signatures and per-argument docstrings, use Python's `help()` on any symbol — e.g. `help(nanovdb.tools.createNanoGridFpN)`. The bindings ship `.pyi` type stubs for IDE / type-checker support. + +## Migration notes + +The CUDA device buffer and device grid handle classes now live under the +`nanovdb.cuda` submodule (mirroring the C++ `nanovdb::cuda` namespace): + +- `nanovdb.DeviceBuffer` → `nanovdb.cuda.DeviceBuffer` +- `nanovdb.DeviceGridHandle` → `nanovdb.cuda.DeviceGridHandle` + +Factory functions that return a device handle are unaffected — they never name +the class — so `nanovdb.io.deviceReadGrid(s)`, `nanovdb.tools.cuda.create*`, +`nanovdb.tools.cuda.pointsToRGBA8Grid`, and the `deviceUpload` / +`deviceDownload` / `deviceGrid` handle methods continue to work unchanged. + +## GPU / CUDA + +The CUDA surface is only present when `nanovdb` is built with CUDA. Gate any +GPU code on both checks (compiled-with-CUDA and a GPU actually present): + +```python +import nanovdb +if nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable(): + ... # nanovdb.cuda / nanovdb.tools.cuda are safe to use +``` + +In a non-CUDA build the `nanovdb.cuda` submodule does not exist, so guard with +`hasattr(nanovdb, "cuda")` (or the checks above) before touching it. + +`nanovdb.cuda` and `nanovdb.tools.cuda` are **attributes** of the package, not +importable submodules. Use them via attribute access after `import nanovdb` +(`import nanovdb.tools.cuda` raises `ModuleNotFoundError`): + +```python +import nanovdb +TC = nanovdb.tools.cuda # ok +# import nanovdb.tools.cuda # NOT ok +``` + +### Namespace layout: `nanovdb.cuda` vs `nanovdb.tools.cuda` + +- **`nanovdb.cuda`** — CUDA *infrastructure* (mirrors the C++ `nanovdb::cuda` + namespace): the buffer / handle types and device plumbing. + `DeviceBuffer`, `DeviceGridHandle`, `UnifiedBuffer`, `UnifiedGridHandle`, + `DeviceMesh`, `DeviceStreamMap`, `DeviceResource`, `TempDevicePool`, + `DeviceNodeManagerHandle`, `createDeviceNodeManager`, and + `compile_options`. +- **`nanovdb.tools.cuda`** — device *algorithms* (mirrors + `nanovdb::tools::cuda`): point / voxel rasterizers (`pointsToGrid`, + `voxelsTo{OnIndex,Index,RGBA8}Grid`, `pointsToRGBA8Grid`), morphology / + topology (`dilateGrid`, `coarsenGrid`, `refineGrid`, `pruneGrid`, + `mergeGrids`), mesh rasterization (`meshToGrid`), index utilities + (`indexToGrid`, `addBlindData`), in-place + device QC (`updateGridStats`, `updateChecksum`, `evalChecksum`, + `validateChecksum`, `isValid`), `signedFloodFill`, `sampleFromVoxels`, + `buildVoxelBlockManager`, and the multi-GPU `Distributed*PointsToGrid` + converters. + +### Building a device grid + +```python +h = nanovdb.tools.createLevelSetSphere(radius=100.0, voxelSize=1.0) # host +fg = h.grid(0) # host FloatGrid +onh = nanovdb.tools.createOnIndexGrid(fg) # host OnIndexGrid handle +nanovdb.io.writeGrid("sphere.nvdb", onh) +dh = nanovdb.io.deviceReadGrid("sphere.nvdb") # nanovdb.cuda.DeviceGridHandle +dh.deviceUpload(0, True) # stream=0, sync=True +dg = dh.deviceGrid(0) # DEVICE grid; feed to tools.cuda.* +``` + +### Zero-copy interop: CAI, DLPack, and `data_ptr()` + +`DeviceGridHandle`, `DeviceBuffer`, and `UnifiedBuffer` expose the whole +device buffer as a 1-D `uint8` array through both the **CUDA Array Interface +(v3)** (`__cuda_array_interface__`) and **DLPack** (`__dlpack__` / +`__dlpack_device__`). So they bridge to CuPy / PyTorch / Numba with no copy: + +```python +import cupy as cp +dh.deviceUpload(0, True) +buf = cp.asarray(dh) # zero-copy; buf.data.ptr == dh.device_ptr() +buf = cp.from_dlpack(dh) # same, via DLPack +assert buf.nbytes == dh.size() +``` + +The data pointer is null and the array empty until `deviceUpload` runs. +`UnifiedGridHandle` intentionally does **not** expose the CAI / DLPack bridges +(it is the managed-memory handle returned by `Distributed*PointsToGrid`); read +its grid through `deviceGrid(n)` and the `tools.cuda.*` ops instead. + +Every typed grid exposes `grid.data_ptr() -> int`. **It is a HOST pointer when +the grid came from `handle.grid(n)` and a DEVICE pointer when it came from +`handle.deviceGrid(n)`** — the grid object itself cannot tell host from device, +so provenance is the caller's responsibility. The device pointer is the base +of a `nanovdb::NanoGrid` in GPU memory (the same C++ ABI), which is +exactly what you pass to a custom CUDA kernel (see `cupy_rawkernel.py`). + +> **Caveat — host accessors segfault on device grids.** A grid from +> `deviceGrid(n)` holds a DEVICE pointer; calling a host-side accessor on it +> (e.g. `dg.getAccessor().getValue(...)`) dereferences GPU memory on the CPU +> and **crashes the process** (SIGSEGV). Host reads are only legal on +> `grid(n)` / host grids. Feed `deviceGrid(n)` **only** to `tools.cuda.*` +> device entry points, or read its bytes via `data_ptr()` / the CAI / DLPack +> buffer from a device kernel. + +### Streams + +Stream arguments everywhere are **raw CUDA stream handles passed as Python +ints** (`0` == the default stream). `deviceUpload` / `deviceDownload` and the +`tools.cuda.*` ops all take a trailing `stream` int. From CuPy, pass +`stream.ptr`: + +```python +s = cp.cuda.Stream(non_blocking=True) +h = nanovdb.tools.cuda.pointsToGrid(points, 1.0, s.ptr) +``` + +### Wrapping external memory (`from_external`) + +`nanovdb.cuda.DeviceBuffer.from_external(size, gpu_ptr, cpu_ptr)` builds a +**non-owning** wrapper around memory you already allocated (e.g. a CuPy / +PyTorch buffer); it never frees, and the caller must keep the source alive. A +null `gpu_ptr` is rejected. Pair it with +`DeviceGridHandle.from_buffer(buffer)` (which **moves** the buffer, leaving it +empty, and validates the NanoVDB header) for a fully zero-copy adopt of grid +bytes you produced elsewhere. + +### Managed memory for the distributed pipeline + +`Distributed{Points,IndexPoints,RGBA8Points}ToGrid.getHandle(voxels)` requires +the `(N, 3)` int32 coordinate array to live in **CUDA managed (unified)** +memory and returns a `nanovdb.cuda.UnifiedGridHandle`. With CuPy, route the +allocation through the managed allocator: + +```python +prev = cp.cuda.get_allocator() +cp.cuda.set_allocator(cp.cuda.malloc_managed) +try: + voxels = cp.asarray(coords_int32) # managed (N,3) int32 + mesh = nanovdb.cuda.DeviceMesh() # must outlive the converter + conv = nanovdb.tools.cuda.DistributedPointsToGrid(mesh, 1.0, (0, 0, 0)) + uh = conv.getHandle(voxels) # UnifiedGridHandle (OnIndex) +finally: + cp.cuda.set_allocator(prev) +``` + +A plain (device-pool) array does not satisfy the `cuda_managed` constraint. + +### Compiling custom kernels against the bundled headers + +`nanovdb.cuda.compile_options(*extra)` returns the NanoVDB include flag +(`-I`) followed by any extra flags, for feeding NVRTC / a runtime CUDA +compiler so your kernel compiles against the same headers as the wheel: + +```python +opts = nanovdb.cuda.compile_options("-std=c++17") +kernel = cp.RawKernel(src, "my_kernel", options=opts, backend="nvrtc") +``` + +The include dir is only physically present in an installed wheel; in an +in-source dev build tree the path resolves but does not exist (see +`cupy_rawkernel.py` for a `NANOVDB_INCLUDE` fallback). + +### Future: pluggable memory resources (roadmap) + +The device buffers above own their allocations directly. A planned evolution is +a **pluggable memory-resource** model (à la a polymorphic allocator), so device +buffers can draw from a caller-supplied pool / async / pinned allocator — +e.g. RAPIDS Memory Manager–style resources. See Mark Harris's design sketch: +. The +following names are **reserved** for that work and should not be used today: +`Resource`, `MemoryResource`, `AsyncResource`, `PinnedResource`, +`ResourceDeviceBuffer`, `DeviceBuffer2`, `default_resource`, +`set_default_resource`. diff --git a/nanovdb/nanovdb/python/examples/build_grid.py b/nanovdb/nanovdb/python/examples/build_grid.py index d4c3f27bed..2e63c9a0bd 100644 --- a/nanovdb/nanovdb/python/examples/build_grid.py +++ b/nanovdb/nanovdb/python/examples/build_grid.py @@ -7,7 +7,7 @@ three construction loops you'll typically reach for (setValue directly, the cached ValueAccessor, and the thread-safe WriteAccessor), then bakes each build grid into a host NanoGrid via -.to_nanovdb() and reads it back through the regular polymorphic +.toNanoVDB() and reads it back through the regular polymorphic handle.grid() API. Run with: python build_grid.py diff --git a/nanovdb/nanovdb/python/examples/bulk_leaf_numpy.py b/nanovdb/nanovdb/python/examples/bulk_leaf_numpy.py index f967a2ae4a..6b34487910 100644 --- a/nanovdb/nanovdb/python/examples/bulk_leaf_numpy.py +++ b/nanovdb/nanovdb/python/examples/bulk_leaf_numpy.py @@ -2,7 +2,7 @@ # SPDX-License-Identifier: Apache-2.0 """Bulk per-leaf value access as a zero-copy NumPy array. -grid.leaf_values() is the highest-bandwidth path from NanoVDB into +grid.leafValues() is the highest-bandwidth path from NanoVDB into NumPy. It returns an (N_leaves, 512) view of every leaf's mValues without copying — modify it, slice it, feed it into a PyTorch tensor, hash it for cache lookup, whatever you need. @@ -21,18 +21,18 @@ def main(): # Build a fog volume sphere with stats so the leaves have meaningful # min/max attached (just for the printing below — not required by - # leaf_values itself). + # leafValues itself). handle = nanovdb.tools.createFogVolumeSphere( radius=20.0, name="bulk_demo") grid = handle.grid() print(f"Grid: {grid.gridType()}, active voxels = {grid.activeVoxelCount()}, " f"leaves = {grid.tree().nodeCount(0)}") - # leaf_values() is the zero-copy view. Modifying it modifies the grid. - bulk = grid.leaf_values() + # leafValues() is the zero-copy view. Modifying it modifies the grid. + bulk = grid.leafValues() # np.asarray adds a NumPy wrapper but doesn't copy. arr = np.asarray(bulk) - print(f"leaf_values: shape={arr.shape}, dtype={arr.dtype}, " + print(f"leafValues: shape={arr.shape}, dtype={arr.dtype}, " f"backed by grid memory (no copy).") # Global statistics across every voxel in every leaf, computed in C. diff --git a/nanovdb/nanovdb/python/examples/collide_level_set_cuda.py b/nanovdb/nanovdb/python/examples/collide_level_set_cuda.py new file mode 100644 index 0000000000..7d9fb28920 --- /dev/null +++ b/nanovdb/nanovdb/python/examples/collide_level_set_cuda.py @@ -0,0 +1,109 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Collide falling particles against a level set on the GPU. + +GPU counterpart to ``collide_level_set.py`` and a Python port of the +device path of ``ex_collide_level_set``. The CPU version loops over +particles one at a time, querying the SDF and its gradient per particle; +here the entire particle set is a pair of ``(N, 3)`` CuPy arrays and a +single ``nanovdb.tools.cuda.sampleFromVoxels`` call returns the signed +distance AND the analytic gradient for every particle at once — exactly +what a collision response needs (distance to push out along, normal to +reflect about). The per-step reflection is then plain vectorized CuPy. + +Requires CuPy and a CUDA-capable GPU. + +Run with: python collide_level_set_cuda.py +""" +import os +import tempfile + +import nanovdb + +NUM_PARTICLES = 4000 +NUM_STEPS = 60 +DT = 0.1 +GRAVITY = -9.8 +RADIUS = 100.0 +BAND = 3.0 # narrow-band half-width in world units (default 3 voxels) + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + import cupy as cp + except ImportError: + print("This example requires CuPy. Install it with: pip install cupy") + return + + host_handle = nanovdb.tools.createLevelSetSphere( + radius=RADIUS, voxelSize=1.0, name="sphere") + tmp = tempfile.NamedTemporaryFile(suffix=".nvdb", delete=False) + tmp.close() + nanovdb.io.writeGrid(tmp.name, host_handle) + try: + handle = nanovdb.io.deviceReadGrid(tmp.name) + finally: + os.unlink(tmp.name) + handle.deviceUpload(0, True) + deviceGrid = handle.deviceGrid(0) + + # Seed particles above the north pole of the sphere, falling down. + rng = cp.random.RandomState(42) + p = cp.empty((NUM_PARTICLES, 3), dtype=cp.float32) + p[:, 0] = rng.uniform(-30.0, 30.0, NUM_PARTICLES) + p[:, 1] = rng.uniform(RADIUS + 15.0, RADIUS + 40.0, NUM_PARTICLES) + p[:, 2] = rng.uniform(-30.0, 30.0, NUM_PARTICLES) + v = cp.zeros((NUM_PARTICLES, 3), dtype=cp.float32) + v[:, 1] = -20.0 + + values = cp.empty(NUM_PARTICLES, dtype=cp.float32) + grads = cp.empty((NUM_PARTICLES, 3), dtype=cp.float32) + total_collisions = 0 + + for step in range(NUM_STEPS): + v[:, 1] += GRAVITY * DT + next_p = cp.ascontiguousarray(p + v * DT) + + nanovdb.tools.cuda.sampleFromVoxels(next_p, deviceGrid, values, grads, 0) + cp.cuda.Stream.null.synchronize() + + # A collision is a point inside the surface but still within the + # meaningful narrow band (deep interior clamps to -background). + hit = (values <= 0.0) & (values > -BAND) + collisions = int(hit.sum()) + total_collisions += collisions + + # Normalized surface normals from the sampled gradients. + norm = cp.linalg.norm(grads, axis=1, keepdims=True) + n = grads / cp.maximum(norm, 1e-8) + mask = hit[:, None] + # Project penetrating particles back onto the surface... + next_p = cp.where(mask, next_p - values[:, None] * n, next_p) + # ...and reflect their velocity about the surface normal. + v_dot_n = (v * n).sum(axis=1, keepdims=True) + v = cp.where(mask, v - 2.0 * v_dot_n * n, v) + p = next_p + + if collisions: + print(f"step {step:2d}: {collisions} collisions") + + print(f"total collisions over {NUM_STEPS} steps: {total_collisions}") + assert total_collisions > 0 + + # No particle should end up deep inside the surface. + nanovdb.tools.cuda.sampleFromVoxels(cp.ascontiguousarray(p), deviceGrid, values, 0) + cp.cuda.Stream.null.synchronize() + in_band = (values > -BAND) & (values < BAND) + if bool(in_band.any()): + worst = float(values[in_band].min()) + print(f"min final SDF among in-band particles: {worst:.3f} world units " + f"(> 0 means all bounced clear of the surface)") + assert worst > -2.0 + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/examples/cupy_rawkernel.py b/nanovdb/nanovdb/python/examples/cupy_rawkernel.py new file mode 100644 index 0000000000..e092f84d94 --- /dev/null +++ b/nanovdb/nanovdb/python/examples/cupy_rawkernel.py @@ -0,0 +1,103 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Run a custom CUDA kernel over a NanoVDB device grid with CuPy. + +This shows the full "bring your own kernel" path: compile a CUDA kernel +that ``#include `` against the SAME headers the nanovdb +wheel was built with (via ``nanovdb.cuda.compile_options()``), then launch +it on the device grid pointer obtained from ``grid.data_ptr()``. + +Device-pointer ABI (read this before writing a kernel): + +* ``handle.deviceGrid(n).data_ptr()`` is a raw DEVICE pointer (a Python + int) to the base of a ``nanovdb::NanoGrid`` in GPU memory. For a + float grid that is ``const nanovdb::NanoGrid*``. +* Pass it to a kernel as a plain pointer argument. With CuPy RawKernel the + argument tuple takes the int directly (CuPy forwards it as a pointer). +* The grid layout in device memory is exactly the C++ ``NanoGrid`` + ABI, so the same accessor / sampler code you would write in C++ works + inside the kernel. +* ``grid.data_ptr()`` cannot tell host from device — only pass a pointer + from ``deviceGrid(n)`` to a device kernel. A host pointer from + ``grid(n)`` would dereference host memory on the GPU (garbage / fault). + +``nanovdb.cuda.compile_options(*extra)`` returns ``-I`` followed +by any extra flags. The header dir is only physically present in an +installed wheel; in an in-source dev build tree it resolves but does not +exist, so this example falls back to a ``NANOVDB_INCLUDE`` env var hint. + +Requires CuPy and a CUDA-capable GPU. + +Run with: python cupy_rawkernel.py +""" +import os + +import nanovdb + +KERNEL_SRC = r""" +#include + +// dGrid is the raw device pointer from FloatGrid.data_ptr(); out is a +// 2-float device buffer that receives [value@origin, activeVoxelCount]. +extern "C" __global__ +void inspect_float_grid(const nanovdb::NanoGrid* dGrid, float* out) +{ + auto acc = dGrid->getAccessor(); + out[0] = acc.getValue(nanovdb::Coord(0, 0, 0)); + out[1] = static_cast(dGrid->activeVoxelCount()); +} +""" + + +def _include_options(): + """compile_options(), falling back to $NANOVDB_INCLUDE in a dev tree.""" + opts = list(nanovdb.cuda.compile_options("-std=c++17")) + inc_dir = opts[0][2:] # strip the leading -I + if not os.path.isdir(inc_dir): + env_inc = os.environ.get("NANOVDB_INCLUDE") + if env_inc and os.path.isdir(env_inc): + opts[0] = f"-I{env_inc}" + else: + print(f"NanoVDB headers not found at {inc_dir!r} (expected in an " + "installed wheel). Set NANOVDB_INCLUDE to the dir that " + "contains nanovdb/NanoVDB.h to run from a source tree.") + return None + return tuple(opts) + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + import cupy as cp + except ImportError: + print("This example requires CuPy. Install it with: pip install cupy") + return + + options = _include_options() + if options is None: + return + + # Build a float level-set sphere directly on the device. + handle = nanovdb.tools.cuda.createLevelSetSphere(nanovdb.GridType.Float, 20) + handle.deviceUpload(0, True) + deviceGrid = handle.deviceGrid(0) + print(f"Device FloatGrid at {hex(deviceGrid.data_ptr())}") + + kernel = cp.RawKernel( + KERNEL_SRC, "inspect_float_grid", options=options, backend="nvrtc") + + out = cp.zeros(2, dtype=cp.float32) + # Launch with the raw device-grid pointer as the first argument. + kernel((1,), (1,), (deviceGrid.data_ptr(), out.data.ptr)) + cp.cuda.runtime.deviceSynchronize() + + value, active = cp.asnumpy(out) + print(f" kernel read value@origin = {value}") + print(f" kernel read activeVoxelCount = {int(active)}") + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/examples/device_topology_ops.py b/nanovdb/nanovdb/python/examples/device_topology_ops.py new file mode 100644 index 0000000000..18381e323f --- /dev/null +++ b/nanovdb/nanovdb/python/examples/device_topology_ops.py @@ -0,0 +1,93 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Morphology and topology operators on device OnIndex grids. + +Python port of the device path shared by ``ex_dilate_nanovdb_cuda``, +``ex_coarsen_nanovdb_cuda``, ``ex_refine_nanovdb_cuda`` and +``ex_merge_nanovdb_cuda``. The C++ examples read a ``.vdb`` through +OpenVDB to source the grid; here the source OnIndex grid is built +on the GPU with ``voxelsToOnIndexGrid``, so the example needs no +OpenVDB and everything stays on the device. + +Each operator returns a fresh device ``GridHandle``: + +* ``dilateGrid(g, op)`` — grow active topology (op 6 = faces, + 26 = faces+edges+vertices). +* ``coarsenGrid(g)`` — 2x topological downsample. +* ``refineGrid(g)`` — 2x topological upsample. +* ``mergeGrids(a, b)`` — active-mask union (strictly binary). +* ``pruneGrid(g, mask)`` — keep only voxels flagged in a per-leaf + retain mask (one ``nanovdb::Mask<3>`` = 8 x uint64 per leaf). + +Requires CuPy and a CUDA-capable GPU. + +Run with: python device_topology_ops.py +""" +import nanovdb + + +def _active(handle): + """Download a device handle and return its active voxel count.""" + handle.deviceDownload(0, True) + return handle.grid(0).activeVoxelCount() + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + import cupy as cp + except ImportError: + print("This example requires CuPy. Install it with: pip install cupy") + return + + # Source grid: a solid 8x8x8 block of voxels, built on the GPU. + lin = cp.arange(8, dtype=cp.int32) + i, j, k = cp.meshgrid(lin, lin, lin, indexing="ij") + coords = cp.ascontiguousarray( + cp.stack([i.ravel(), j.ravel(), k.ravel()], axis=1)) + src = nanovdb.tools.cuda.voxelsToOnIndexGrid(coords, 1.0, 0) + srcGrid = src.deviceGrid(0) + src_active = _active(src) + print(f"source block: {src_active} active voxels") + + dil6 = nanovdb.tools.cuda.dilateGrid(srcGrid, 6, 0) + dil26 = nanovdb.tools.cuda.dilateGrid(srcGrid, 26, 0) + coarse = nanovdb.tools.cuda.coarsenGrid(srcGrid, 0) + fine = nanovdb.tools.cuda.refineGrid(srcGrid, 0) + a6, a26 = _active(dil6), _active(dil26) + ac, af = _active(coarse), _active(fine) + print(f"dilate(faces) -> {a6} active") + print(f"dilate(faces+e+v) -> {a26} active") + print(f"coarsen (2x down) -> {ac} active") + print(f"refine (2x up) -> {af} active") + assert a6 > src_active and a26 >= a6 + assert ac < src_active < af + + # Union with a shifted copy of the block; the merged topology must + # cover at least the larger of the two inputs. + shifted = nanovdb.tools.cuda.voxelsToOnIndexGrid( + cp.ascontiguousarray(coords + cp.asarray([4, 0, 0], dtype=cp.int32)), + 1.0, 0) + merged = nanovdb.tools.cuda.mergeGrids(srcGrid, shifted.deviceGrid(0), 0) + am = _active(merged) + print(f"merge(block, block+4x) -> {am} active") + assert am > src_active + + # Prune with a retain-all mask (8 uint64 per leaf); topology is + # unchanged, demonstrating the mask-driven prune entry point. + src.deviceDownload(0, True) + leaf_count = src.grid(0).tree().nodeCount(0) + retain_all = cp.full(leaf_count * 8, 0xFFFFFFFFFFFFFFFF, dtype=cp.uint64) + pruned = nanovdb.tools.cuda.pruneGrid(srcGrid, retain_all, 0) + ap = _active(pruned) + print(f"prune(retain-all, {leaf_count} leaves) -> {ap} active") + assert ap == src_active + + print("OK: dilate / coarsen / refine / merge / prune on the device") + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/examples/gpu_load_inspect.py b/nanovdb/nanovdb/python/examples/gpu_load_inspect.py new file mode 100644 index 0000000000..19517e9278 --- /dev/null +++ b/nanovdb/nanovdb/python/examples/gpu_load_inspect.py @@ -0,0 +1,101 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Load a NanoVDB grid onto the GPU and inspect its device buffer zero-copy. + +This is the GPU counterpart to ``load_inspect.py``. It builds a grid, +writes it to disk, reads it straight onto the device with +``nanovdb.io.deviceReadGrid``, uploads it with ``deviceUpload``, and then +views the whole device buffer as a CuPy array WITHOUT copying via the +CUDA Array Interface (``cupy.asarray(handle)``). + +Key facts demonstrated: + +* ``handle.deviceUpload(stream, sync)`` takes a RAW CUDA stream handle as + a Python int (0 == the default stream). +* ``handle.__cuda_array_interface__`` (and ``__dlpack__``) expose the + whole device buffer as a 1-D ``uint8`` array; the data pointer is null + until ``deviceUpload`` runs. +* ``cupy.asarray(handle)`` aliases ``handle.device_ptr()`` with no copy; + ``arr.nbytes == handle.size()``. +* A single handle exposes BOTH a host grid via ``handle.grid(n)`` and a + device grid via ``handle.deviceGrid(n)``. They have different + ``data_ptr()`` values and the grid object cannot tell host from device + (see the docstring on ``Grid.data_ptr``) — feed ``deviceGrid(n)`` only + to ``nanovdb.tools.cuda.*`` device entry points. + +Requires CuPy and a CUDA-capable GPU. + +Run with: python gpu_load_inspect.py +""" +import os +import tempfile + +import nanovdb + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + import cupy as cp + except ImportError: + print("This example requires CuPy. Install it with: pip install cupy") + return + + # Build a small level-set sphere on the host, then write it so we can + # read it back straight onto the device. + host_handle = nanovdb.tools.createLevelSetSphere( + radius=20.0, voxelSize=1.0, name="gpu_sphere") + tmp = tempfile.NamedTemporaryFile(suffix=".nvdb", delete=False) + tmp.close() + nanovdb.io.writeGrid(tmp.name, host_handle) + + try: + # deviceReadGrid returns a nanovdb.cuda.DeviceGridHandle. + handle = nanovdb.io.deviceReadGrid(tmp.name) + finally: + os.unlink(tmp.name) + + print(f"Read device handle: gridCount={handle.gridCount()}, " + f"size={handle.size()} bytes") + print(f" device_ptr before upload = {handle.device_ptr()} " + f"(0 == not uploaded yet)") + print(f" deviceGrid(0) before upload = {handle.deviceGrid(0)}") + + # Upload to the device on the default stream, synchronously. + handle.deviceUpload(0, True) + print(f" device_ptr after upload = {hex(handle.device_ptr())}") + + # Zero-copy view of the whole device buffer as a CuPy uint8 array. + buf = cp.asarray(handle) + print(f" cupy.asarray(handle): shape={buf.shape}, dtype={buf.dtype}, " + f"nbytes={buf.nbytes}") + print(f" aliases device_ptr (no copy): " + f"{int(buf.data.ptr) == handle.device_ptr()}") + + # The host grid and the device grid share the handle but differ in + # provenance: the host grid's data_ptr is a host address, the device + # grid's data_ptr is a device address. The grid object itself cannot + # tell them apart. + host_grid = handle.grid(0) + deviceGrid = handle.deviceGrid(0) + print(f" host grid.data_ptr() = {hex(host_grid.data_ptr())} (CPU)") + print(f" device grid.data_ptr() = {hex(deviceGrid.data_ptr())} (GPU)") + print(f" device grid.data_ptr() == handle.device_ptr(): " + f"{deviceGrid.data_ptr() == handle.device_ptr()}") + + # The device grid is the input to nanovdb.tools.cuda.* ops. Validate + # it entirely on the device. + print(f" tools.cuda.isValid(deviceGrid) = " + f"{nanovdb.tools.cuda.isValid(deviceGrid)}") + + print("WARNING: calling a host-side accessor (e.g. " + "deviceGrid.getAccessor().getValue(...)) on a DEVICE grid " + "dereferences GPU memory on the CPU and SEGFAULTS. Use host_grid " + "for host reads.") + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/examples/index_grid_channels.py b/nanovdb/nanovdb/python/examples/index_grid_channels.py index 91e421d08f..dc54833867 100644 --- a/nanovdb/nanovdb/python/examples/index_grid_channels.py +++ b/nanovdb/nanovdb/python/examples/index_grid_channels.py @@ -17,15 +17,15 @@ def index_grid_with_channel(): src = nanovdb.tools.createLevelSetSphere(radius=50.0, name="sphere") - src_grid = src.grid() + srcGrid = src.grid() # channels=1 copies the source values into blind-data channel 0, # indexed by the per-voxel uint64 indices the OnIndex grid stores. - handle = nanovdb.tools.createNanoGridOnIndex(src_grid, channels=1) + handle = nanovdb.tools.createNanoGridOnIndex(srcGrid, channels=1) grid = handle.grid() print(f"OnIndex grid: valueCount={grid.valueCount()}, " - f"source activeVoxelCount={src_grid.activeVoxelCount()}") - assert grid.valueCount() >= src_grid.activeVoxelCount() + f"source activeVoxelCount={srcGrid.activeVoxelCount()}") + assert grid.valueCount() >= srcGrid.activeVoxelCount() # createChannelAccessor inspects the channel's recorded dataType and # returns the matching typed accessor (here: OnIndexFloat...). @@ -33,7 +33,7 @@ def index_grid_with_channel(): print(f"channel accessor: {type(channel).__name__}, " f"valueCount={channel.valueCount()}") - src_acc = src_grid.getAccessor() + src_acc = srcGrid.getAccessor() for ijk in (nanovdb.math.Coord(48, 0, 0), nanovdb.math.Coord(0, 50, 0), nanovdb.math.Coord(0, 0, 52)): via_channel = channel.getValue(ijk) diff --git a/nanovdb/nanovdb/python/examples/index_to_grid_cuda.py b/nanovdb/nanovdb/python/examples/index_to_grid_cuda.py new file mode 100644 index 0000000000..423ecced38 --- /dev/null +++ b/nanovdb/nanovdb/python/examples/index_to_grid_cuda.py @@ -0,0 +1,82 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Bake a per-voxel value array onto an index grid, on the device. + +Python port of the device path of ``ex_index_grid_cuda`` (the GPU half +of ``index_grid_channels.py``). An OnIndex grid stores only topology; +the per-voxel payload lives in a separate linear array indexed by +``grid.valueCount()``. ``nanovdb.tools.cuda.indexToGrid`` fuses the two +on the device — an index grid plus a value array in, a fully-typed +value grid out: + + value_handle = tools.cuda.indexToGrid(d_index_grid, values, stream) + +To fill the value array meaningfully we use +``tools.cuda.activeVoxelCoords`` to recover each value slot's +index-space coordinate (the decode companion to the index grid), then +compute a value per voxel with plain CuPy — here the distance from the +origin, turning the shell into a scalar field. + +Requires CuPy and a CUDA-capable GPU. + +Run with: python index_to_grid_cuda.py +""" +import nanovdb + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + import cupy as cp + except ImportError: + print("This example requires CuPy. Install it with: pip install cupy") + return + + # A hollow sphere shell of voxels, built on the GPU. + radius = 24.0 + lin = cp.arange(-30, 31, dtype=cp.float32) + x, y, z = cp.meshgrid(lin, lin, lin, indexing="ij") + r = cp.sqrt(x * x + y * y + z * z) + shell = cp.abs(r - radius) < 1.0 + coords = cp.ascontiguousarray( + cp.stack([x[shell], y[shell], z[shell]], axis=1).astype(cp.int32)) + index_handle = nanovdb.tools.cuda.voxelsToOnIndexGrid(coords, 1.0, 0) + index_grid = index_handle.deviceGrid(0) + + index_handle.deviceDownload(0, True) + value_count = index_handle.grid(0).valueCount() + print(f"OnIndex grid: {index_handle.grid(0).activeVoxelCount()} voxels, " + f"valueCount={value_count}") + + # Recover the coordinate of every value slot, then compute a value + # (distance from origin) per voxel with CuPy. Row 0 is the background + # slot; leave it at zero. + voxel_coords = cp.zeros((value_count, 3), dtype=cp.int32) + nanovdb.tools.cuda.activeVoxelCoords(index_grid, voxel_coords, 9, 0) + cp.cuda.Stream.null.synchronize() + fc = voxel_coords.astype(cp.float32) + values = cp.sqrt((fc * fc).sum(axis=1)).astype(cp.float32) + values[0] = 0.0 # background slot + values = cp.ascontiguousarray(values) + + value_handle = nanovdb.tools.cuda.indexToGrid(index_grid, values, 0) + value_grid = value_handle.deviceGrid(0) + print(f"indexToGrid -> {value_handle.gridType(0)} grid") + assert nanovdb.tools.cuda.isValid(value_grid, nanovdb.CheckMode.Full) + + # Confirm the CuPy-computed values actually landed in the grid by + # reading one back on the host: a shell voxel holds its radius. + value_handle.deviceDownload(0, True) + sampler = nanovdb.math.createTrilinearSampler(value_handle.grid(0)) + baked = sampler(nanovdb.math.Vec3f(radius, 0.0, 0.0)) + print(f" baked value at (r,0,0) = {baked:.2f} (shell radius {radius:.0f})") + assert abs(baked - radius) < 1.0 + + print("OK: baked a CuPy-computed value array onto an index grid on device") + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/examples/levelset_filter.py b/nanovdb/nanovdb/python/examples/levelset_filter.py new file mode 100644 index 0000000000..eae92fc9e7 --- /dev/null +++ b/nanovdb/nanovdb/python/examples/levelset_filter.py @@ -0,0 +1,330 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""GPU LevelSetFilter on NanoVDB .nvdb files -- driver for three compute backends. + +This is the runnable driver for ONE GPU level-set filter -- a +tools::LevelSetFilter-style Laplacian deform + first-order Godunov reinit + +narrow-band retrack, driven by the device VoxelBlockManager -- with the +per-voxel stencil math supplied by one of three interchangeable backends: + + python levelset_filter.py rawkernel in.nvdb out.nvdb [outer_iters] + python levelset_filter.py cupy in.nvdb out.nvdb [outer_iters] + python levelset_filter.py cutile in.nvdb out.nvdb [outer_iters] + +(no input/output files => a self-test for that backend over both input styles) + +The backends live in sibling files and differ ONLY in how the dense per-voxel +stencils are computed: + * levelset_filter_rawkernel.py -- a hand-written CUDA kernel (cupy.RawModule) + that decodes the VBM + gathers in-kernel + * levelset_filter_cupy.py -- pure CuPy array ops over a dense gather + * levelset_filter_cutile.py -- NVIDIA cuTile tile kernels over the gather +Each also runs standalone: `python levelset_filter_.py` executes a +stencil-only smoke test (deform + renorm on a sphere, no file I/O / no retrack) +of just that backend's math. + +Everything that is NOT backend-specific lives here and is shared by all three: +the .nvdb read (FloatGrid OR OnIndexGrid+blind-SDF, auto-detected), the +style-preserving write, the narrow-band retrack (dilateGrid -> inject -> +extrapolate -> injectPredicateToMask -> pruneGrid -> inject), the surface-radius +probe, and the outer filter loop. A backend is a small object implementing: + + class Backend: + NAME # "rawkernel" | "cupy" | "cutile" + cp # the cupy module + setup(handle) -> g # {"grid", "n", ...}; per topology change + active_coords(g) -> (n+1, 3) int32 # value-indexed voxel coords + laplacian(g, phi, half_width) -> phi' + godunov(g, phi, vx, half_width) -> phi' + extrapolate(g, phi, vx) -> phi' + +and a module-level `make_backend()` that returns a `Backend` (or None, with a +printed reason, if that backend's requirements -- a CUDA build, a GPU, CuPy, +nvcc, or cuda.tile -- are absent). The opaque `g` context is produced by the +backend and handed back to it, so each backend stashes whatever bookkeeping it +needs (the rawkernel backend caches VBM pointers + decoded coords; the gather +backends keep it minimal). `g["grid"]` (the device OnIndex grid) and `g["n"]` +(the active-voxel count, = valueCount - 1) are the only fields the driver reads. + +The SDF is carried as a value-indexed device array `phi` of length n+1: slot 0 +is the background, slots 1..n the active voxels (the order the VBM decode and +`getBlindData` use). `phi[0]` is held at a sentinel so inactive neighbours are +detectable in a dense gather. Scope: first-order Godunov reinit, no advection or +alpha mask; the output's inactive interior carries +background (no signed +flood-fill). +""" +import importlib +import os +import sys +import tempfile + +import numpy as np + +import nanovdb + + +BACKENDS = ("rawkernel", "cupy", "cutile") + +LOG2_BLOCK_WIDTH = 9 +SENTINEL = 1.0e30 # phi[0]: marks inactive neighbours in a dense gather +NN_FACE = 6 # nanovdb::tools::morphology::NN_FACE (6-face dilation) +# 3x3x3 box-stencil spoke columns for the six faces, in -/+ x, y, z order +# (spoke = (di+1)*9 + (dj+1)*3 + (dk+1); centre = 13). +FACES = [4, 22, 10, 16, 12, 14] +BAND = 3 # narrow-band half width, in voxels +DEFORM_ITERS = 4 # Laplacian deform sub-iterations per outer iteration +NORMALIZE_ITERS = 5 # Godunov reinit sub-iterations per outer iteration + + +class GatherBackend: + """Base for backends that read each voxel's 3x3x3 neighbourhood as a dense + (n+1, 27) array via the bound `gatherBoxStencil` -- the CuPy and cuTile + backends. Subclasses implement only the per-voxel stencils + (`laplacian`/`godunov`/`extrapolate`); `setup`/`active_coords` are generic.""" + + def __init__(self, cp): + self.cp = cp + + def setup(self, handle): + cp = self.cp + grid = handle.deviceGrid(0) + if grid is None or grid.data_ptr() == 0: + handle.deviceUpload(0, True) + grid = handle.deviceGrid(0) + n = int(nanovdb.tools.cuda.buildVoxelBlockManager( + grid, log2BlockWidth=LOG2_BLOCK_WIDTH).lastOffset()) + return {"handle": handle, "grid": grid, "n": n} + + def active_coords(self, g): + cp = self.cp + coords = cp.empty((g["n"] + 1, 3), dtype=cp.int32) + nanovdb.tools.cuda.activeVoxelCoords(g["grid"], coords) + return coords + + +def load_backend(name): + """Import levelset_filter_ and construct its Backend, or return None + (the backend prints why) if its requirements are unavailable.""" + if name not in BACKENDS: + raise SystemExit(f"unknown backend {name!r}; choose one of {', '.join(BACKENDS)}") + return importlib.import_module(f"levelset_filter_{name}").make_backend() + + +def gather_faces(cp, g, phi, background, clamp=True): + """`gatherBoxStencil` -> (centre values c, (n, 6) face values f). + + With clamp=True the sign-clamped background BC (`copysign(background, c)`) + replaces inactive (sentinel) spokes -- the boundary condition the deform and + reinit stencils use. With clamp=False the raw sentinel is preserved, so the + extrapolation stencil can tell which face neighbours are actually active.""" + n = g["n"] + nbrs = cp.empty((n + 1, 27), dtype=cp.float32) + nanovdb.tools.cuda.gatherBoxStencil(g["grid"], phi, nbrs) + c = phi[1:n + 1] + f = nbrs[1:n + 1][:, FACES] + if clamp: + f = cp.where(f == SENTINEL, cp.copysign(cp.float32(background), c)[:, None], f) + return c, f + + +def read_to_device(backend, path, band): + """Read a .nvdb (FloatGrid OR OnIndex+SDF) -> (g, phi, vx, half_width, style). + + A FloatGrid is baked to OnIndex+blind-SDF first (the representation the VBM + operates on); an OnIndexGrid is read directly. `style` (the input grid type) + is returned so the output can be written back in the same form.""" + cp = backend.cp + io, T = nanovdb.io, nanovdb.tools + host = io.readGrid(path) + gtype = host.gridType(0) + vx = float(host.grid(0).voxelSize()[0]) + tmp = None + if gtype == nanovdb.GridType.Float: + onh = T.createOnIndexGrid(host.grid(0), channels=1, + includeStats=False, includeTiles=False) + sdf = np.array(onh.grid(0).getBlindData(0), dtype=np.float32) + tmp = tempfile.NamedTemporaryFile(suffix=".nvdb", delete=False); tmp.close() + io.writeGrid(tmp.name, onh) + dh = io.deviceReadGrid(tmp.name) + elif gtype == nanovdb.GridType.OnIndex: + if host.grid(0).blindDataCount() == 0: + raise SystemExit(f"{path}: OnIndex grid has no blind-data SDF channel.") + sdf = np.array(host.grid(0).getBlindData(0), dtype=np.float32) + dh = io.deviceReadGrid(path) + else: + raise SystemExit(f"{path}: unsupported grid type {gtype} (expected Float or OnIndex).") + g = backend.setup(dh) + if sdf.shape[0] != g["n"] + 1: + raise SystemExit(f"{path}: SDF channel length {sdf.shape[0]} != activeVoxelCount+1 " + f"({g['n'] + 1}); the OnIndex grid must use contiguous voxel indexing " + "(built with includeStats=False, includeTiles=False).") + phi = cp.asarray(sdf) + phi[0] = SENTINEL # inactive-neighbour marker + if tmp is not None: + os.unlink(tmp.name) + return g, phi, vx, band * vx, gtype + + +def sphere_on_device(backend, radius, voxel_size=1.0, band=BAND, name="sphere"): + """Build a level-set sphere as OnIndex+SDF on the device -> (g, phi, vx, + half_width). Used by the backends' standalone stencil smoke tests.""" + cp = backend.cp + io, T = nanovdb.io, nanovdb.tools + fg = T.createLevelSetSphere(radius=radius, voxelSize=voxel_size, name=name) + onh = T.createOnIndexGrid(fg.grid(0), channels=1, + includeStats=False, includeTiles=False) + sdf = np.array(onh.grid(0).getBlindData(0), dtype=np.float32) + tmp = tempfile.NamedTemporaryFile(suffix=".nvdb", delete=False); tmp.close() + io.writeGrid(tmp.name, onh) + dh = io.deviceReadGrid(tmp.name) + os.unlink(tmp.name) + g = backend.setup(dh) + phi = cp.asarray(sdf) + phi[0] = SENTINEL + return g, phi, voxel_size, band * voxel_size + + +def rebuild(backend, g, phi, vx, half_width): + """Narrow-band retrack: dilate -> inject old phi -> extrapolate the new ring + (backend) -> |phi| <= halfWidth predicate -> injectPredicateToMask -> prune + -> inject. Returns the pruned grid context + its value-indexed sidecar.""" + cp, TC = backend.cp, nanovdb.tools.cuda + gd = backend.setup(TC.dilateGrid(g["grid"], op=NN_FACE)) + phi_d = cp.full(gd["n"] + 1, SENTINEL, dtype=cp.float32) + TC.inject(g["grid"], gd["grid"], phi, phi_d) # carry old phi (intersection) + phi_d = backend.extrapolate(gd, phi_d, vx) # fill the freshly-dilated ring + predicate = cp.abs(phi_d) <= half_width # phi_d[0]=SENTINEL -> False + leafMasks = cp.zeros(gd["n"] * 8, dtype=cp.uint64) + TC.injectPredicateToMask(gd["grid"], predicate, leafMasks) + gp = backend.setup(TC.pruneGrid(gd["grid"], leafMasks)) + phi_p = cp.full(gp["n"] + 1, SENTINEL, dtype=cp.float32) + TC.inject(gd["grid"], gp["grid"], phi_d, phi_p) + return gp, phi_p + + +def surface_radius(backend, g, phi, vx): + """Mean world radius of the zero-crossing voxels (|phi| < dx/2).""" + cp = backend.cp + coords = backend.active_coords(g) + v = phi[1:g["n"] + 1] + near = cp.abs(v) < 0.5 * vx + c = coords[1:g["n"] + 1][near].astype(cp.float64) * vx + return float(cp.mean(cp.linalg.norm(c, axis=1))) if int(near.sum()) else float("nan") + + +def write_output(backend, g, phi, vx, path, style, band, name="filtered"): + """Bake (coords, phi) into a host FloatGrid and write it in `style` + (GridType.Float -> a FloatGrid; GridType.OnIndex -> an OnIndexGrid with the + SDF in blind channel 0).""" + cp = backend.cp + T, io = nanovdb.tools, nanovdb.io + coords = cp.asnumpy(backend.active_coords(g)) + v = cp.asnumpy(phi) + builder = T.build.FloatGrid(float(band * vx)) + builder.setName(name) + builder.setTransform(vx) + acc = builder.getAccessor() + Coord = nanovdb.math.Coord + for k in range(1, g["n"] + 1): + i, j, kk = coords[k] + acc.setValue(Coord(int(i), int(j), int(kk)), float(v[k])) + fh = builder.toNanoVDB() + if style == nanovdb.GridType.OnIndex: + io.writeGrid(path, T.createOnIndexGrid(fh.grid(0), channels=1, + includeStats=False, includeTiles=False)) + else: + io.writeGrid(path, fh) + + +def run_filter(backend, in_path, out_path, outer_iters=6, band=BAND, + deform_iters=DEFORM_ITERS, normalize_iters=NORMALIZE_ITERS): + """Read -> N outer iterations (deform x k, renorm x k, retrack) -> write.""" + g, phi, vx, half_width, gtype = read_to_device(backend, in_path, band) + print(f"read {in_path}: {gtype}, {g['n']} active voxels, voxelSize {vx:g} " + f"[backend: {backend.NAME}]") + r0 = r = surface_radius(backend, g, phi, vx) + for it in range(outer_iters): + for _ in range(deform_iters): + phi = backend.laplacian(g, phi, half_width) + for _ in range(normalize_iters): + phi = backend.godunov(g, phi, vx, half_width) + g, phi = rebuild(backend, g, phi, vx, half_width) + r = surface_radius(backend, g, phi, vx) + print(f" iter {it + 1}: {g['n']:7d} active, surface radius = {r:.4f}") + write_output(backend, g, phi, vx, out_path, gtype, band) + style = "OnIndex+SDF" if gtype == nanovdb.GridType.OnIndex else "FloatGrid" + print(f"wrote {out_path}: {style} (same style as input), {g['n']} active voxels " + f"(surface radius {r0:.4f} -> {r:.4f}) [backend: {backend.NAME}]") + return r0, r, gtype + + +def stencil_demo(backend, radius=20.0, deform_iters=DEFORM_ITERS, + normalize_iters=NORMALIZE_ITERS): + """Run ONLY the per-voxel stencils (deform + renorm) on a sphere -- no file + I/O and no narrow-band retrack -- as a focused check of one backend's dense + math. Invoked by each backend file's `__main__`.""" + g, phi, vx, half_width = sphere_on_device(backend, radius) + r0 = surface_radius(backend, g, phi, vx) + for _ in range(deform_iters): + phi = backend.laplacian(g, phi, half_width) + for _ in range(normalize_iters): + phi = backend.godunov(g, phi, vx, half_width) + r = surface_radius(backend, g, phi, vx) + print(f"[{backend.NAME}] stencil-only smoke test on a sphere (radius {radius:g}, " + f"{g['n']} voxels, no retrack / no I/O): surface radius {r0:.4f} -> {r:.4f}") + assert r < r0, "Laplacian deform should shrink the sphere" + print(f"OK [{backend.NAME}]: deform + renorm stencils run and move the surface.") + + +def self_test(backend): + """Filter sphere .nvdb files of both styles; assert the style round-trips and + the sphere shrinks under curvature flow.""" + io, T, GT = nanovdb.io, nanovdb.tools, nanovdb.GridType + tmps = [tempfile.NamedTemporaryFile(suffix=".nvdb", delete=False) for _ in range(4)] + for t in tmps: + t.close() + f_in, f_out, o_in, o_out = (t.name for t in tmps) + + io.writeGrid(f_in, T.createLevelSetSphere(radius=20.0, voxelSize=1.0, name="sphere")) + print(f"self-test 1 [{backend.NAME}]: FloatGrid sphere -> filter -> FloatGrid") + r0, r, style = run_filter(backend, f_in, f_out, outer_iters=6) + assert style == GT.Float and io.readGrid(f_out).gridType(0) == GT.Float, \ + "output style is not FloatGrid" + assert r < r0 - 0.05, "sphere did not shrink under curvature flow" + + sphere = T.createLevelSetSphere(radius=18.0, voxelSize=1.0, name="sphere") + io.writeGrid(o_in, T.createOnIndexGrid(sphere.grid(0), channels=1, + includeStats=False, includeTiles=False)) + print(f"self-test 2 [{backend.NAME}]: OnIndex+SDF sphere -> filter -> OnIndex+SDF") + r0b, rb, style2 = run_filter(backend, o_in, o_out, outer_iters=4) + ro = io.readGrid(o_out) + assert style2 == GT.OnIndex and ro.gridType(0) == GT.OnIndex, "output style is not OnIndex" + assert ro.grid(0).blindDataCount() >= 1, "OnIndex output has no SDF blind channel" + assert rb < r0b - 0.05, "sphere did not shrink under curvature flow" + + print(f"OK [{backend.NAME}]: both input styles filter correctly and the output " + "style matches the input.") + for n in (f_in, f_out, o_in, o_out): + os.unlink(n) + + +def main(argv): + if len(argv) < 2 or argv[1] not in BACKENDS: + print(__doc__) + raise SystemExit(f"usage: levelset_filter.py {{{'|'.join(BACKENDS)}}} " + "[input.nvdb output.nvdb [outer_iterations]] (no files = self-test)") + backend = load_backend(argv[1]) + if backend is None: + return + rest = argv[2:] + if len(rest) >= 2: + outer = int(rest[2]) if len(rest) >= 3 else 6 + run_filter(backend, rest[0], rest[1], outer) + elif len(rest) == 0: + self_test(backend) + else: + raise SystemExit("provide BOTH input.nvdb and output.nvdb, or neither (self-test).") + + +if __name__ == "__main__": + main(sys.argv) diff --git a/nanovdb/nanovdb/python/examples/levelset_filter_cupy.py b/nanovdb/nanovdb/python/examples/levelset_filter_cupy.py new file mode 100644 index 0000000000..520e0ca24c --- /dev/null +++ b/nanovdb/nanovdb/python/examples/levelset_filter_cupy.py @@ -0,0 +1,93 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""levelset_filter pure-CuPy backend -- per-voxel stencils as plain CuPy arrays. + +The kernel-free backend for levelset_filter.py: every stage runs as plain +**CuPy** array math on the dense `(n, 6)` face values returned by the bound +`gatherBoxStencil` -- no `cupy.RawModule`, no CUDA C++, no nvcc. Run the full +file->file filter through the driver: + + python levelset_filter.py cupy input.nvdb output.nvdb [outer_iterations] + +Running this file directly executes a stencil-only smoke test (the deform + +renorm steps on a sphere, no file I/O and no narrow-band retrack): + + python levelset_filter_cupy.py + +The dense `(n, 6)` arrays this operates on are exactly the shape a tile +framework consumes -- see levelset_filter_cutile.py for the same math as cuTile +kernels, and levelset_filter_rawkernel.py for a fused hand-written CUDA kernel. +Requires only CuPy and a CUDA-capable GPU (no nvcc / NanoVDB headers). +""" +import nanovdb + +import levelset_filter as lsf + + +def make_backend(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This backend requires a CUDA build of nanovdb and a GPU. Skipping.") + return None + try: + import cupy as cp + except ImportError: + print("This backend requires CuPy. Skipping.") + return None + return Backend(cp) + + +class Backend(lsf.GatherBackend): + """Per-voxel stencils as CuPy array ops; `setup`/`active_coords` come from the + shared GatherBackend (gatherBoxStencil / activeVoxelCoords).""" + + NAME = "cupy" + + def laplacian(self, g, phi, half_width): + """phi += (sum6 - 6 phi)/6.""" + cp = self.cp + c, f = lsf.gather_faces(cp, g, phi, half_width, clamp=True) + out = phi.copy() + out[1:g["n"] + 1] = c + (f.sum(axis=1) - 6.0 * c) / 6.0 + out[0] = lsf.SENTINEL + return out + + def godunov(self, g, phi, vx, half_width): + """phi -= dt*S(phi)*(|grad phi| - 1) -- first-order Godunov reinit.""" + cp = self.cp + c, f = lsf.gather_faces(cp, g, phi, half_width, clamp=True) + xm, xp, ym, yp, zm, zp = (f[:, i] for i in range(6)) + dt = 0.3 * vx + s = c / cp.sqrt(c * c + vx * vx) + + def gd(dm, dp): + return cp.where(s > 0, + cp.maximum(cp.maximum(dm, 0.0) ** 2, cp.minimum(dp, 0.0) ** 2), + cp.maximum(cp.minimum(dm, 0.0) ** 2, cp.maximum(dp, 0.0) ** 2)) + + grad = cp.sqrt(gd((c - xm) / vx, (xp - c) / vx) + + gd((c - ym) / vx, (yp - c) / vx) + + gd((c - zm) / vx, (zp - c) / vx)) + out = phi.copy() + out[1:g["n"] + 1] = c - dt * s * (grad - 1.0) + out[0] = lsf.SENTINEL + return out + + def extrapolate(self, g, phi, vx): + """Fill freshly-dilated (sentinel) voxels from the nearest in-band face + neighbour: phi = phi_nbr + sign(phi_nbr)*dx.""" + cp = self.cp + c, f = lsf.gather_faces(cp, g, phi, 0.0, clamp=False) # raw spokes + known = f != lsf.SENTINEL + best = cp.take_along_axis( + f, cp.argmin(cp.where(known, cp.abs(f), cp.inf), axis=1)[:, None], axis=1)[:, 0] + out = phi.copy() + fill = (c == lsf.SENTINEL) & known.any(axis=1) + out[1:g["n"] + 1] = cp.where(fill, best + cp.copysign(cp.float32(vx), best), c) + out[0] = lsf.SENTINEL + return out + + +if __name__ == "__main__": + backend = make_backend() + if backend is not None: + lsf.stencil_demo(backend) diff --git a/nanovdb/nanovdb/python/examples/levelset_filter_cutile.py b/nanovdb/nanovdb/python/examples/levelset_filter_cutile.py new file mode 100644 index 0000000000..cde4172e5e --- /dev/null +++ b/nanovdb/nanovdb/python/examples/levelset_filter_cutile.py @@ -0,0 +1,155 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""levelset_filter NVIDIA cuTile backend -- per-voxel stencils as tile kernels. + +The cuTile backend for levelset_filter.py: this is levelset_filter_cupy.py with +the per-voxel stencil math (deform / renorm / extrapolate) moved from CuPy array +ops into NVIDIA cuTile (`cuda.tile`) kernels that `ct.load`/`ct.store` over +`(TILE,)` tiles. The SPARSE half (neighbour gather, coord decode, topology) +stays in bound NanoVDB ops; only the DENSE per-voxel compute is cuTile. Run the +full file->file filter through the driver: + + python levelset_filter.py cutile input.nvdb output.nvdb [outer_iterations] + +Running this file directly executes a stencil-only smoke test (the deform + +renorm kernels on a sphere, no file I/O and no narrow-band retrack): + + python levelset_filter_cutile.py + +The kernels load six 1-D face arrays per tile (cuTile tile dims must be +compile-time powers of two, so a `(TILE, 6)` tile of the gather columns isn't +allowed); the sign-clamped background BC is applied in CuPy before launch +(`lsf.gather_faces`), and the extrapolation BC is done inside its kernel. +Validated against levelset_filter_cupy.py (same spheres shrink the same amount). +Requires CuPy + cuda-tile (`cuda.tile`) and a CUDA-capable GPU. +""" +import nanovdb + +import levelset_filter as lsf + +try: + import cupy as cp + import cuda.tile as ct + HAVE_CUTILE = True +except ImportError: + HAVE_CUTILE = False + + +def _kernel(fn): + """Apply @cuda.tile.kernel when available; else a no-op so the module still + imports (the kernels are never called when cuTile / a GPU is absent).""" + return ct.kernel(fn) if HAVE_CUTILE else fn + + +def make_backend(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This backend requires a CUDA build of nanovdb and a GPU. Skipping.") + return None + if not HAVE_CUTILE: + print("This backend requires CuPy + cuda-tile (cuda.tile). Skipping.") + return None + return Backend(cp) + + +TILE = 256 +SENTINEL = 1.0e30 # must match levelset_filter.SENTINEL (read inside a kernel) + + +# ----------------------------- cuTile kernels ----------------------------- +@_kernel +def laplacian_kernel(c, xm, xp, ym, yp, zm, zp, out): + b = ct.bid(0) + cc = ct.load(c, index=(b,), shape=(TILE,)) + s6 = (ct.load(xm, index=(b,), shape=(TILE,)) + ct.load(xp, index=(b,), shape=(TILE,)) + + ct.load(ym, index=(b,), shape=(TILE,)) + ct.load(yp, index=(b,), shape=(TILE,)) + + ct.load(zm, index=(b,), shape=(TILE,)) + ct.load(zp, index=(b,), shape=(TILE,))) + ct.store(out, index=(b,), tile=cc + (s6 - 6.0 * cc) / 6.0) + + +@_kernel +def godunov_kernel(c, xm, xp, ym, yp, zm, zp, dx, dt, out): + b = ct.bid(0) + cc = ct.load(c, index=(b,), shape=(TILE,)) + xm_ = ct.load(xm, index=(b,), shape=(TILE,)); xp_ = ct.load(xp, index=(b,), shape=(TILE,)) + ym_ = ct.load(ym, index=(b,), shape=(TILE,)); yp_ = ct.load(yp, index=(b,), shape=(TILE,)) + zm_ = ct.load(zm, index=(b,), shape=(TILE,)); zp_ = ct.load(zp, index=(b,), shape=(TILE,)) + s = cc / ct.sqrt(cc * cc + dx * dx) + pos = s > 0.0 + + def axis(dm, dp): + return ct.where(pos, + ct.maximum(ct.maximum(dm, 0.0) * ct.maximum(dm, 0.0), + ct.minimum(dp, 0.0) * ct.minimum(dp, 0.0)), + ct.maximum(ct.minimum(dm, 0.0) * ct.minimum(dm, 0.0), + ct.maximum(dp, 0.0) * ct.maximum(dp, 0.0))) + + grad = ct.sqrt(axis((cc - xm_) / dx, (xp_ - cc) / dx) + + axis((cc - ym_) / dx, (yp_ - cc) / dx) + + axis((cc - zm_) / dx, (zp_ - cc) / dx)) + ct.store(out, index=(b,), tile=cc - dt * s * (grad - 1.0)) + + +@_kernel +def extrapolate_kernel(c, xm, xp, ym, yp, zm, zp, dx, out): + b = ct.bid(0) + cc = ct.load(c, index=(b,), shape=(TILE,)) + f0 = ct.load(xm, index=(b,), shape=(TILE,)); f1 = ct.load(xp, index=(b,), shape=(TILE,)) + f2 = ct.load(ym, index=(b,), shape=(TILE,)); f3 = ct.load(yp, index=(b,), shape=(TILE,)) + f4 = ct.load(zm, index=(b,), shape=(TILE,)); f5 = ct.load(zp, index=(b,), shape=(TILE,)) + best = ct.full((TILE,), SENTINEL, ct.float32) # min-|value| active neighbour + best = ct.where(ct.abs(f0) < ct.abs(best), f0, best) + best = ct.where(ct.abs(f1) < ct.abs(best), f1, best) + best = ct.where(ct.abs(f2) < ct.abs(best), f2, best) + best = ct.where(ct.abs(f3) < ct.abs(best), f3, best) + best = ct.where(ct.abs(f4) < ct.abs(best), f4, best) + best = ct.where(ct.abs(f5) < ct.abs(best), f5, best) + filled = best + ct.where(best >= 0.0, dx, -dx) # phi_nbr + sign*dx + is_new = cc == SENTINEL + has = ct.abs(best) < SENTINEL + ct.store(out, index=(b,), tile=ct.where(is_new, ct.where(has, filled, cc), cc)) + + +def _pad(cp, a, m): + out = cp.zeros(m, dtype=a.dtype) + out[:a.shape[0]] = a + return out + + +class Backend(lsf.GatherBackend): + """Per-voxel stencils as cuTile kernels; `setup`/`active_coords` come from the + shared GatherBackend (gatherBoxStencil / activeVoxelCoords).""" + + NAME = "cutile" + + def _apply(self, kernel, g, phi, clamp, background, extra): + """gather (driver) -> contiguous columns -> pad -> cuTile launch -> new + value-indexed sidecar.""" + cp = self.cp + n = g["n"] + c, f = lsf.gather_faces(cp, g, phi, background, clamp=clamp) + cols = [cp.ascontiguousarray(f[:, i]) for i in range(6)] + m = ct.cdiv(n, TILE) * TILE + args = [_pad(cp, a, m) for a in (cp.ascontiguousarray(c), *cols)] + out = cp.zeros(m, dtype=cp.float32) + ct.launch(cp.cuda.get_current_stream(), (ct.cdiv(m, TILE), 1, 1), + kernel, (*args, *extra, out)) + cp.cuda.runtime.deviceSynchronize() + new = cp.full(n + 1, SENTINEL, dtype=cp.float32) + new[1:n + 1] = out[:n] + return new + + def laplacian(self, g, phi, half_width): + return self._apply(laplacian_kernel, g, phi, True, half_width, ()) + + def godunov(self, g, phi, vx, half_width): + return self._apply(godunov_kernel, g, phi, True, half_width, + (float(vx), float(0.3 * vx))) + + def extrapolate(self, g, phi, vx): + return self._apply(extrapolate_kernel, g, phi, False, 0.0, (float(vx),)) + + +if __name__ == "__main__": + backend = make_backend() + if backend is not None: + lsf.stencil_demo(backend) diff --git a/nanovdb/nanovdb/python/examples/levelset_filter_rawkernel.py b/nanovdb/nanovdb/python/examples/levelset_filter_rawkernel.py new file mode 100644 index 0000000000..ec0e6432d5 --- /dev/null +++ b/nanovdb/nanovdb/python/examples/levelset_filter_rawkernel.py @@ -0,0 +1,300 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""levelset_filter compiled-CUDA-kernel backend -- one fused cupy.RawModule. + +The hand-written-CUDA-kernel backend for levelset_filter.py: each stencil is a +`cupy.RawModule` (nvcc) kernel that fuses the VoxelBlockManager decode +(`decodeInverseMaps`) + 3x3x3 neighbour gather (`computeBoxStencil`) + the +per-voxel update into a single launch -- no dense `gatherBoxStencil` array. Run +the full file->file filter through the driver: + + python levelset_filter.py rawkernel input.nvdb output.nvdb [outer_iterations] + +Running this file directly executes a stencil-only smoke test (the deform + +renorm kernels on a sphere, no file I/O and no narrow-band retrack): + + python levelset_filter_rawkernel.py + +This is the fastest of the three backends (the gather is fused into the compute, +no `(n, 27)` array is materialised) and the most code -- compare the kernel-free +levelset_filter_cupy.py / cuTile levelset_filter_cutile.py for the same math. +Requires CuPy, a CUDA-capable GPU, and nvcc (CuPy honours the NVCC env var). In +a dev/source tree, set NANOVDB_INCLUDE to the dir containing nanovdb/NanoVDB.h. +""" +import os + +import numpy as np + +import nanovdb + +import levelset_filter as lsf + + +LOG2_BLOCK_WIDTH = lsf.LOG2_BLOCK_WIDTH # 9 +BLOCK_WIDTH = 1 << LOG2_BLOCK_WIDTH # 512 +NN_FACE = lsf.NN_FACE # 6-face dilation +SENTINEL = lsf.SENTINEL # "value not yet known" marker + +# computeBoxStencil spoke ids: spoke = (di+1)*9 + (dj+1)*3 + (dk+1). +# centre=13; +x=22 -x=4; +y=16 -y=10; +z=14 -z=12 +KERNEL_SRC = r""" +#include + +using namespace nanovdb; +using VBM = nanovdb::tools::cuda::VoxelBlockManager<9>; // BlockWidth = 512 +static constexpr int BLOCK_WIDTH = 512; +static constexpr int JUMP_MAP_LENGTH = BLOCK_WIDTH / 64; // = 8 +static constexpr float SENTINEL = 1.0e30f; + +__device__ static unsigned long long +decodeBlock(const NanoGrid* grid, + const uint32_t* firstLeafID, const uint64_t* jumpMap, + uint64_t firstOffset, + uint32_t* smem_leafIndex, uint16_t* smem_voxelOffset) +{ + const int blockID = blockIdx.x; + const uint64_t blockFirstOffset = + firstOffset + (uint64_t)blockID * BLOCK_WIDTH; + VBM::decodeInverseMaps( + grid, firstLeafID[blockID], + &jumpMap[(uint64_t)blockID * JUMP_MAP_LENGTH], + blockFirstOffset, smem_leafIndex, smem_voxelOffset); + return blockFirstOffset; // decodeInverseMaps ends in __syncthreads() +} + +__device__ static inline float +readNbr(const float* v, uint64_t spoke, float vc, float background) +{ + return (spoke == 0) ? copysignf(background, vc) : v[spoke]; +} + +extern "C" __global__ +void decode_coords(const NanoGrid* grid, + const uint32_t* firstLeafID, const uint64_t* jumpMap, + uint64_t firstOffset, int* coords) +{ + __shared__ uint32_t smem_leafIndex[BLOCK_WIDTH]; + __shared__ uint16_t smem_voxelOffset[BLOCK_WIDTH]; + decodeBlock(grid, firstLeafID, jumpMap, firstOffset, + smem_leafIndex, smem_voxelOffset); + const int tID = threadIdx.x; + if (smem_leafIndex[tID] == VBM::UnusedLeafIndex) return; + const auto& leaf = grid->tree().getFirstNode<0>()[smem_leafIndex[tID]]; + const uint16_t off = smem_voxelOffset[tID]; + const Coord c = leaf.offsetToGlobalCoord(off); + const uint64_t idx = leaf.getValue(off); + coords[idx * 3 + 0] = c[0]; + coords[idx * 3 + 1] = c[1]; + coords[idx * 3 + 2] = c[2]; +} + +// One Laplacian-flow iteration (OpenVDB LevelSetFilter::laplacianImpl). +extern "C" __global__ +void laplacian_step(const NanoGrid* grid, + const uint32_t* firstLeafID, const uint64_t* jumpMap, + uint64_t firstOffset, float background, + const float* vin, float* vout) +{ + __shared__ uint32_t smem_leafIndex[BLOCK_WIDTH]; + __shared__ uint16_t smem_voxelOffset[BLOCK_WIDTH]; + decodeBlock(grid, firstLeafID, jumpMap, firstOffset, + smem_leafIndex, smem_voxelOffset); + const int tID = threadIdx.x; + if (smem_leafIndex[tID] == VBM::UnusedLeafIndex) return; + + uint64_t st[27]; + VBM::computeBoxStencil( + grid, smem_leafIndex[tID], smem_voxelOffset[tID], st); + + const uint64_t c = st[13]; + const float vc = vin[c]; + const float sum6 = readNbr(vin, st[22], vc, background) + + readNbr(vin, st[4], vc, background) + + readNbr(vin, st[16], vc, background) + + readNbr(vin, st[10], vc, background) + + readNbr(vin, st[14], vc, background) + + readNbr(vin, st[12], vc, background); + vout[c] = vc + (sum6 - 6.0f * vc) / 6.0f; +} + +// One Godunov reinitialisation iteration: phi -= dt*S(phi)*(|grad phi| - 1). +extern "C" __global__ +void godunov_step(const NanoGrid* grid, + const uint32_t* firstLeafID, const uint64_t* jumpMap, + uint64_t firstOffset, float dx, float dt, float background, + const float* vin, float* vout) +{ + __shared__ uint32_t smem_leafIndex[BLOCK_WIDTH]; + __shared__ uint16_t smem_voxelOffset[BLOCK_WIDTH]; + decodeBlock(grid, firstLeafID, jumpMap, firstOffset, + smem_leafIndex, smem_voxelOffset); + const int tID = threadIdx.x; + if (smem_leafIndex[tID] == VBM::UnusedLeafIndex) return; + + uint64_t st[27]; + VBM::computeBoxStencil( + grid, smem_leafIndex[tID], smem_voxelOffset[tID], st); + + const uint64_t c = st[13]; + const float vc = vin[c]; + const float xm = readNbr(vin, st[4], vc, background); + const float xp = readNbr(vin, st[22], vc, background); + const float ym = readNbr(vin, st[10], vc, background); + const float yp = readNbr(vin, st[16], vc, background); + const float zm = readNbr(vin, st[12], vc, background); + const float zp = readNbr(vin, st[14], vc, background); + + const float Dxm = (vc - xm) / dx, Dxp = (xp - vc) / dx; + const float Dym = (vc - ym) / dx, Dyp = (yp - vc) / dx; + const float Dzm = (vc - zm) / dx, Dzp = (zp - vc) / dx; + + const float s = vc / sqrtf(vc * vc + dx * dx); // smoothed sign + float gx, gy, gz; + if (s > 0.0f) { + gx = fmaxf(powf(fmaxf(Dxm, 0.0f), 2), powf(fminf(Dxp, 0.0f), 2)); + gy = fmaxf(powf(fmaxf(Dym, 0.0f), 2), powf(fminf(Dyp, 0.0f), 2)); + gz = fmaxf(powf(fmaxf(Dzm, 0.0f), 2), powf(fminf(Dzp, 0.0f), 2)); + } else { + gx = fmaxf(powf(fminf(Dxm, 0.0f), 2), powf(fmaxf(Dxp, 0.0f), 2)); + gy = fmaxf(powf(fminf(Dym, 0.0f), 2), powf(fmaxf(Dyp, 0.0f), 2)); + gz = fmaxf(powf(fminf(Dzm, 0.0f), 2), powf(fmaxf(Dzp, 0.0f), 2)); + } + const float grad = sqrtf(gx + gy + gz); + vout[c] = vc - dt * s * (grad - 1.0f); +} + +// Fill freshly-dilated (sentinel) voxels by SDF extrapolation from their +// nearest in-band face neighbour: phi = phi_nbr + sign(phi_nbr)*dx. +extern "C" __global__ +void extrapolate(const NanoGrid* grid, + const uint32_t* firstLeafID, const uint64_t* jumpMap, + uint64_t firstOffset, float dx, const float* vin, float* vout) +{ + __shared__ uint32_t smem_leafIndex[BLOCK_WIDTH]; + __shared__ uint16_t smem_voxelOffset[BLOCK_WIDTH]; + decodeBlock(grid, firstLeafID, jumpMap, firstOffset, + smem_leafIndex, smem_voxelOffset); + const int tID = threadIdx.x; + if (smem_leafIndex[tID] == VBM::UnusedLeafIndex) return; + + uint64_t st[27]; + VBM::computeBoxStencil( + grid, smem_leafIndex[tID], smem_voxelOffset[tID], st); + + const uint64_t c = st[13]; + const float vc = vin[c]; + if (vc != SENTINEL) { vout[c] = vc; return; } // already known + const int faces[6] = {22, 4, 16, 10, 14, 12}; + float best = SENTINEL; + for (int f = 0; f < 6; f++) { + const uint64_t ni = st[faces[f]]; + if (ni != 0) { // active neighbour + const float v = vin[ni]; + if (v != SENTINEL && fabsf(v) < fabsf(best)) best = v; + } + } + vout[c] = (best != SENTINEL) ? best + copysignf(dx, best) : vc; +} +""" + + +def _include_options(): + opts = list(nanovdb.cuda.compile_options("-std=c++17")) + inc = opts[0][2:] + if not os.path.isdir(inc): + env = os.environ.get("NANOVDB_INCLUDE") + if env and os.path.isdir(os.path.join(env, "nanovdb")): + opts[0] = f"-I{env}" + else: + print(f"NanoVDB headers not found at {inc!r}. In a source tree set " + "NANOVDB_INCLUDE to the dir containing nanovdb/NanoVDB.h.") + return None + return tuple(opts) + + +def make_backend(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This backend requires a CUDA build of nanovdb and a GPU. Skipping.") + return None + try: + import cupy as cp + except ImportError: + print("This backend requires CuPy (plus nvcc on PATH or $NVCC). Skipping.") + return None + options = _include_options() + if options is None: + return None + return Backend(cp, options) + + +class Backend: + """Per-voxel stencils as fused cupy.RawModule CUDA kernels that decode the VBM + and gather the 3x3x3 neighbourhood in-kernel (no dense gatherBoxStencil), so + this backend keeps its own VBM bookkeeping in the `g` context.""" + + NAME = "rawkernel" + + def __init__(self, cp, options): + self.cp = cp + self.tc = nanovdb.tools.cuda + m = cp.RawModule(code=KERNEL_SRC, backend="nvcc", options=options) + self.k_decode = m.get_function("decode_coords") + self.k_laplacian = m.get_function("laplacian_step") + self.k_godunov = m.get_function("godunov_step") + self.k_extrapolate = m.get_function("extrapolate") + + def setup(self, handle): + """Device OnIndex grid + VBM pointers; decode value-indexed coords once.""" + cp = self.cp + grid = handle.deviceGrid(0) + if grid is None or grid.data_ptr() == 0: + handle.deviceUpload(0, True) + grid = handle.deviceGrid(0) + vbm = self.tc.buildVoxelBlockManager(grid, log2BlockWidth=LOG2_BLOCK_WIDTH) + n, bc = int(vbm.lastOffset()), int(vbm.blockCount()) + coords = cp.zeros((n + 1, 3), dtype=cp.int32) + self.k_decode((bc,), (BLOCK_WIDTH,), + (grid.data_ptr(), vbm.first_leaf_id_ptr(), vbm.jump_map_ptr(), + np.uint64(vbm.firstOffset()), coords)) + cp.cuda.runtime.deviceSynchronize() + return dict(handle=handle, grid=grid, vbm=vbm, n=n, bc=bc, + fo=np.uint64(vbm.firstOffset()), + fid=vbm.first_leaf_id_ptr(), jmp=vbm.jump_map_ptr(), + gptr=grid.data_ptr(), coords=coords) + + def active_coords(self, g): + return g["coords"] # decoded in setup() + + def _vbm(self, g): + return (g["gptr"], g["fid"], g["jmp"], g["fo"]) + + def laplacian(self, g, phi, half_width): + cp = self.cp + out = cp.empty_like(phi) + self.k_laplacian((g["bc"],), (BLOCK_WIDTH,), + (*self._vbm(g), np.float32(half_width), phi, out)) + out[0] = SENTINEL + return out + + def godunov(self, g, phi, vx, half_width): + cp = self.cp + out = cp.empty_like(phi) + self.k_godunov((g["bc"],), (BLOCK_WIDTH,), + (*self._vbm(g), np.float32(vx), np.float32(0.3 * vx), + np.float32(half_width), phi, out)) + out[0] = SENTINEL + return out + + def extrapolate(self, g, phi, vx): + cp = self.cp + out = cp.empty_like(phi) + self.k_extrapolate((g["bc"],), (BLOCK_WIDTH,), + (*self._vbm(g), np.float32(vx), phi, out)) + out[0] = SENTINEL + return out + + +if __name__ == "__main__": + backend = make_backend() + if backend is not None: + lsf.stencil_demo(backend) diff --git a/nanovdb/nanovdb/python/examples/mesh_to_grid_cuda.py b/nanovdb/nanovdb/python/examples/mesh_to_grid_cuda.py new file mode 100644 index 0000000000..80b285309e --- /dev/null +++ b/nanovdb/nanovdb/python/examples/mesh_to_grid_cuda.py @@ -0,0 +1,94 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Rasterize a triangle mesh into a narrow-band SDF on the GPU. + +Python port of the device path of ``ex_mesh_to_grid_cuda``. +``nanovdb.tools.cuda.meshToGrid`` takes device arrays of mesh vertices and +triangle indices and builds, entirely on the GPU, a narrow-band ``OnIndex`` +grid plus a per-value unsigned-distance-field (UDF) sidecar. The C++ example +sources its mesh through OpenVDB; here the mesh (a closed box) is generated +with NumPy, so the example needs no OpenVDB. + +The returned pieces chain naturally: + +* ``handle`` — a device ``OnIndex`` grid holding the narrow-band topology. +* ``udf`` — a ``nanovdb.cuda.DeviceBuffer`` of ``valueCount`` float32 unsigned + distances (voxel units), indexed by the grid's per-voxel value index. That + is exactly the value array ``tools.cuda.indexToGrid`` consumes, so we bake + the UDF into a Float distance grid and sample it back to verify. + +Requires CuPy and a CUDA-capable GPU. + +Run with: python mesh_to_grid_cuda.py +""" +import nanovdb + + +def _box_mesh(cp, half): + """Return (vertices (8,3) float32, triangles (12,3) int32) for a cube.""" + corners = cp.asarray([ + [-1, -1, -1], [1, -1, -1], [1, 1, -1], [-1, 1, -1], + [-1, -1, 1], [1, -1, 1], [1, 1, 1], [-1, 1, 1], + ], dtype=cp.float32) * half + # 12 triangles (2 per face), outward winding. + tris = cp.asarray([ + [0, 3, 2], [0, 2, 1], # -z + [4, 5, 6], [4, 6, 7], # +z + [0, 1, 5], [0, 5, 4], # -y + [3, 7, 6], [3, 6, 2], # +y + [0, 4, 7], [0, 7, 3], # -x + [1, 2, 6], [1, 6, 5], # +x + ], dtype=cp.int32) + return cp.ascontiguousarray(corners), cp.ascontiguousarray(tris) + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + import cupy as cp + except ImportError: + print("This example requires CuPy. Install it with: pip install cupy") + return + + half = 40.0 + voxel_size = 1.0 + vertices, triangles = _box_mesh(cp, half) + print(f"Box mesh: {vertices.shape[0]} vertices, {triangles.shape[0]} " + f"triangles, half-extent {half:.0f}") + + # Build the narrow-band UDF on the device in one call. + handle, udf = nanovdb.tools.cuda.meshToGrid( + vertices, triangles, voxel_size, 3.0, "box_udf", 0) + index_grid = handle.deviceGrid(0) + assert nanovdb.tools.cuda.isValid(index_grid) + + handle.deviceDownload(0, True) + value_count = handle.grid(0).valueCount() + # The sidecar is a device buffer of `value_count` float32 UDF values. + udf_values = cp.asarray(udf).view(cp.float32)[:value_count] + udf_values = cp.ascontiguousarray(udf_values) + print(f"meshToGrid -> {handle.gridType(0)} grid, " + f"{handle.grid(0).activeVoxelCount()} narrow-band voxels, " + f"UDF range [{float(udf_values[1:].min()):.2f}, " + f"{float(udf_values[1:].max()):.2f}] voxels") + + # Bake the UDF into a Float distance grid on the device and sample it. + dist_handle = nanovdb.tools.cuda.indexToGrid(index_grid, udf_values, 0) + dist_grid = dist_handle.deviceGrid(0) + assert nanovdb.tools.cuda.isValid(dist_grid, nanovdb.CheckMode.Full) + + # A point right on a face should read a near-zero unsigned distance. + dist_handle.deviceDownload(0, True) + sampler = nanovdb.math.createTrilinearSampler(dist_handle.grid(0)) + on_face = sampler(nanovdb.math.Vec3f(half, 0.0, 0.0)) + print(f" baked UDF at a box face = {on_face:.3f} voxels (expect near 0)") + assert on_face < 2.0 + + print("OK: built a narrow-band SDF from a triangle mesh on the device") + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/examples/node_manager.py b/nanovdb/nanovdb/python/examples/node_manager.py index 3ca3feffbe..3e15de87e4 100644 --- a/nanovdb/nanovdb/python/examples/node_manager.py +++ b/nanovdb/nanovdb/python/examples/node_manager.py @@ -7,7 +7,7 @@ they can be visited by index instead of by tree traversal. Each node exposes its origin, per-node stats, and (on leaves) the raw 512-value buffer. For bulk NumPy analytics over every leaf at once, see -bulk_leaf_numpy.py's grid.leaf_values() instead. +bulk_leaf_numpy.py's grid.leafValues() instead. Run with: python node_manager.py """ diff --git a/nanovdb/nanovdb/python/examples/numba_cuda.py b/nanovdb/nanovdb/python/examples/numba_cuda.py new file mode 100644 index 0000000000..4a0d8832ac --- /dev/null +++ b/nanovdb/nanovdb/python/examples/numba_cuda.py @@ -0,0 +1,86 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Consume a NanoVDB device grid from a Numba CUDA kernel. + +Numba does not parse the C++ NanoVDB header, so it cannot use the +``NanoGrid`` accessor types directly the way the CuPy RawKernel example +(``cupy_rawkernel.py``) does. The realistic Numba pattern is therefore to +operate on the zero-copy *device arrays* NanoVDB hands back — the whole +grid buffer (CAI / DLPack on the handle) or the typed per-leaf / blind-data +buffers — rather than to dereference the ``NanoGrid`` ABI from JIT'd code. + +Device-pointer ABI recap (see cupy_rawkernel.py for the C++ ABI route): + +* ``handle.deviceGrid(n).data_ptr()`` is a raw DEVICE pointer (Python int) + to a ``nanovdb::NanoGrid``. Decoding the NanoGrid layout from a + Numba kernel means re-implementing the offset math by hand — out of + scope for this example. +* The supported, zero-copy Numba path: wrap a NanoVDB buffer that already + exposes ``__cuda_array_interface__`` (the DeviceGridHandle, a + UnifiedBuffer, or the VoxelBlockManager firstLeafID / jumpMap DLPack + capsules) with ``numba.cuda.as_cuda_array`` and process the raw bytes / + indices in a kernel. +* Provenance is the caller's responsibility: only feed device pointers / + device CAI buffers to device kernels. A host pointer from ``grid(n)`` + would fault on the GPU. + +This example wraps the whole device grid buffer as a uint8 CUDA array and +runs a trivial reduction over it to show the interop plumbing. Requires +Numba (with CUDA support) and a CUDA-capable GPU; it skips with a message +when Numba is not installed. + +Run with: python numba_cuda.py +""" +import nanovdb + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + from numba import cuda + except ImportError: + print("This example requires Numba (with CUDA support). Install it " + "with: pip install numba") + return + if not cuda.is_available(): + print("Numba reports no CUDA device available. Skipping.") + return + + # Build a float level-set sphere on the device. + handle = nanovdb.tools.cuda.createLevelSetSphere(nanovdb.GridType.Float, 20) + handle.deviceUpload(0, True) + print(f"Device grid buffer: {handle.size()} bytes at " + f"{hex(handle.device_ptr())}") + + # Zero-copy: Numba adopts the handle's __cuda_array_interface__ (the + # whole device buffer as 1-D uint8) with no copy. + dev_bytes = cuda.as_cuda_array(handle) + print(f" numba CUDA array: shape={dev_bytes.shape}, " + f"dtype={dev_bytes.dtype}") + + # Trivial per-element kernel over the raw buffer (XOR-fold into a + # device scalar) just to demonstrate launching JIT'd Numba code on a + # NanoVDB-owned device buffer. + @cuda.jit + def fold_xor(buf, out): + i = cuda.grid(1) + if i < buf.size: + cuda.atomic.xor(out, 0, buf[i]) + + import numpy as np + out = cuda.to_device(np.zeros(1, dtype=np.uint8)) + threads = 256 + blocks = (dev_bytes.size + threads - 1) // threads + fold_xor[blocks, threads](dev_bytes, out) + cuda.synchronize() + print(f" XOR-fold of device buffer = {int(out.copy_to_host()[0])}") + print(" (Decoding the NanoGrid ABI itself from Numba requires hand-" + "rolled offset math; use the CuPy RawKernel route for C++ " + "accessor access — see cupy_rawkernel.py.)") + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/examples/quantize.py b/nanovdb/nanovdb/python/examples/quantize.py index 83935f3b68..229d4ce211 100644 --- a/nanovdb/nanovdb/python/examples/quantize.py +++ b/nanovdb/nanovdb/python/examples/quantize.py @@ -20,9 +20,9 @@ def gridSize_in_kb(handle): def main(): # Source: a 50-radius sphere fog volume at full float precision. src_handle = nanovdb.tools.createFogVolumeSphere(radius=50.0) - src_grid = src_handle.grid() + srcGrid = src_handle.grid() print(f"Source FloatGrid: {gridSize_in_kb(src_handle):.1f} KB, " - f"active voxels = {src_grid.activeVoxelCount()}") + f"active voxels = {srcGrid.activeVoxelCount()}") # Fixed-width quantization. Each subsequent format roughly halves # the per-voxel storage cost; dithering optional. @@ -31,7 +31,7 @@ def main(): ("createNanoGridFp8", "Fp8 (8-bit fixed)"), ("createNanoGridFp4", "Fp4 (4-bit fixed)"), ]: - h = getattr(nanovdb.tools, fn_name)(src_grid, ditherOn=True) + h = getattr(nanovdb.tools, fn_name)(srcGrid, ditherOn=True) print(f" {label}: {gridSize_in_kb(h):.1f} KB, " f"active voxels = {h.grid().activeVoxelCount()}") @@ -40,11 +40,11 @@ def main(): # relative error. -1 tolerance means "uninitialized" so we pass # an explicit value. abs_oracle = nanovdb.tools.AbsDiff(0.05) # ±0.05 per voxel - h_fpn_abs = nanovdb.tools.createNanoGridFpN(src_grid, abs_oracle) + h_fpn_abs = nanovdb.tools.createNanoGridFpN(srcGrid, abs_oracle) print(f" FpN (AbsDiff 0.05): {gridSize_in_kb(h_fpn_abs):.1f} KB") rel_oracle = nanovdb.tools.RelDiff(0.1) # 10% relative error - h_fpn_rel = nanovdb.tools.createNanoGridFpN(src_grid, rel_oracle) + h_fpn_rel = nanovdb.tools.createNanoGridFpN(srcGrid, rel_oracle) print(f" FpN (RelDiff 0.10): {gridSize_in_kb(h_fpn_rel):.1f} KB") # The output is a regular NanoGrid — read-only, but exposes diff --git a/nanovdb/nanovdb/python/examples/raytrace_fog_volume_cuda.py b/nanovdb/nanovdb/python/examples/raytrace_fog_volume_cuda.py new file mode 100644 index 0000000000..e9800d8cdb --- /dev/null +++ b/nanovdb/nanovdb/python/examples/raytrace_fog_volume_cuda.py @@ -0,0 +1,139 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Ray-march a fog volume on the GPU with a custom CUDA kernel. + +GPU counterpart to ``raytrace_fog_volume.py`` and a Python port of the +device path of ``ex_raytrace_fog_volume``. The host port integrates +transmittance one pixel at a time in Python; here a ``cupy.RawKernel`` +does one thread per pixel with the full C++ NanoVDB API: a +``nanovdb::math::Ray`` clipped to the grid, then a fixed-step march +accumulating optical depth from ``acc.getValue(Coord::Floor(...))`` — the +accessor-based sampling idiom, on the device. The whole image renders in +one launch. + +The kernel compiles against the bundled NanoVDB headers via +``nanovdb.cuda.compile_options()`` (with the ``NANOVDB_INCLUDE`` dev-tree +fallback from ``cupy_rawkernel.py``). + +Requires CuPy and a CUDA-capable GPU. + +Run with: python raytrace_fog_volume_cuda.py +""" +import os +import tempfile + +import nanovdb + +RES = 256 +FOV_DEG = 45.0 +STEP = 0.5 # index-space march step +DENSITY_SCALE = 40.0 # maps fog density to optical depth + +KERNEL_SRC = r""" +#include +#include + +extern "C" __global__ +void render_fog(const nanovdb::NanoGrid* dGrid, + unsigned char* image, int res, float tan_fov, + float eye_x, float eye_y, float eye_z, + float step, float density_scale) +{ + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= res || y >= res) return; + + using Vec3T = nanovdb::math::Vec3f; + using RayT = nanovdb::math::Ray; + + auto acc = dGrid->tree().getAccessor(); + const float px = (2.0f * (x + 0.5f) / res - 1.0f) * tan_fov; + const float py = (2.0f * (y + 0.5f) / res - 1.0f) * tan_fov; + const float inv = 1.0f / sqrtf(px * px + py * py + 1.0f); + + RayT ray(Vec3T(eye_x, eye_y, eye_z), + Vec3T(px * inv, py * inv, -inv)); + float transmittance = 1.0f; + if (ray.clip(dGrid->indexBBox())) { + for (float t = ray.t0(); t < ray.t1(); t += step) { + const Vec3T pos = ray(t); + const float sigma = acc.getValue(nanovdb::Coord::Floor(pos)); + if (sigma > 0.0f) + transmittance *= expf(-sigma * step * density_scale); + } + } + // Opacity = 1 - transmittance, mapped to an 8-bit gray value. + image[y * res + x] = (unsigned char)(255.0f * (1.0f - transmittance)); +} +""" + + +def _include_options(): + """compile_options(), falling back to $NANOVDB_INCLUDE in a dev tree.""" + opts = list(nanovdb.cuda.compile_options("-std=c++17")) + inc_dir = opts[0][2:] + if not os.path.isdir(inc_dir): + env_inc = os.environ.get("NANOVDB_INCLUDE") + if env_inc and os.path.isdir(env_inc): + opts[0] = f"-I{env_inc}" + else: + print(f"NanoVDB headers not found at {inc_dir!r} (expected in an " + "installed wheel). Set NANOVDB_INCLUDE to the dir that " + "contains nanovdb/NanoVDB.h to run from a source tree.") + return None + return tuple(opts) + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + import cupy as cp + except ImportError: + print("This example requires CuPy. Install it with: pip install cupy") + return + + options = _include_options() + if options is None: + return + + handle = nanovdb.tools.cuda.createFogVolumeSphere( + nanovdb.GridType.Float, radius=100.0) + handle.deviceUpload(0, True) + handle.deviceDownload(0, True) + deviceGrid = handle.deviceGrid(0) + bbox = handle.grid(0).indexBBox() + + import math + dim = [bbox.max[i] + 1 - bbox.min[i] for i in range(3)] + center = [bbox.min[i] + 0.5 * dim[i] for i in range(3)] + eye = (center[0], center[1], center[2] + 2.0 * dim[2]) + tan_fov = math.tan(math.radians(FOV_DEG) * 0.5) + + kernel = cp.RawKernel( + KERNEL_SRC, "render_fog", options=options, backend="nvrtc") + image = cp.zeros(RES * RES, dtype=cp.uint8) + block = (16, 16) + grid = ((RES + 15) // 16, (RES + 15) // 16) + kernel(grid, block, + (deviceGrid.data_ptr(), image, RES, cp.float32(tan_fov), + cp.float32(eye[0]), cp.float32(eye[1]), cp.float32(eye[2]), + cp.float32(STEP), cp.float32(DENSITY_SCALE))) + cp.cuda.runtime.deviceSynchronize() + + host_image = cp.asnumpy(image) + lit = int((host_image > 0).sum()) + out_dir = tempfile.mkdtemp(prefix="nanovdb_") + path = os.path.join(out_dir, "raytrace_fog_volume_cuda.pgm") + with open(path, "wb") as f: + f.write(f"P5\n{RES} {RES}\n255\n".encode("ascii")) + f.write(host_image.tobytes()) + print(f"Rendered {RES}x{RES} fog volume on the GPU in one launch, " + f"{lit} lit pixels -> {path}") + assert lit > 0 + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/examples/raytrace_level_set_cuda.py b/nanovdb/nanovdb/python/examples/raytrace_level_set_cuda.py new file mode 100644 index 0000000000..9a5bda953c --- /dev/null +++ b/nanovdb/nanovdb/python/examples/raytrace_level_set_cuda.py @@ -0,0 +1,151 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Ray-trace a level set on the GPU with a custom CUDA kernel. + +GPU counterpart to ``raytrace_level_set.py`` and a Python port of the +device path of ``ex_raytrace_level_set``. The host port loops over +pixels in Python and had to approximate the C++ HDDA surface search with +sphere tracing (``math::Ray`` / ``ZeroCrossing`` are not bound in +Python). Inside a ``cupy.RawKernel`` we have the FULL C++ NanoVDB API, +so this version is the faithful original: one CUDA thread per pixel, +``nanovdb::math::Ray`` clipped to the grid, ``nanovdb::math::ZeroCrossing`` +(HDDA) for the exact surface hit, and a central-difference gradient for +Lambert shading. The whole image renders in one launch. + +The kernel is compiled against the bundled NanoVDB headers via +``nanovdb.cuda.compile_options()`` (with the ``NANOVDB_INCLUDE`` dev-tree +fallback from ``cupy_rawkernel.py``) and reads the device grid straight +from ``deviceGrid(0).data_ptr()``. + +Requires CuPy and a CUDA-capable GPU. + +Run with: python raytrace_level_set_cuda.py +""" +import os +import tempfile + +import nanovdb + +RES = 256 +FOV_DEG = 45.0 +LIGHT = (0.577, 0.577, 0.577) + +KERNEL_SRC = r""" +#include +#include +#include + +extern "C" __global__ +void render_level_set(const nanovdb::NanoGrid* dGrid, + unsigned char* image, int res, float tan_fov, + float eye_x, float eye_y, float eye_z, + float lx, float ly, float lz) +{ + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= res || y >= res) return; + + using Vec3T = nanovdb::math::Vec3f; + using RayT = nanovdb::math::Ray; + + auto acc = dGrid->tree().getAccessor(); + const float px = (2.0f * (x + 0.5f) / res - 1.0f) * tan_fov; + const float py = (2.0f * (y + 0.5f) / res - 1.0f) * tan_fov; + const float inv = 1.0f / sqrtf(px * px + py * py + 1.0f); + + RayT ray(Vec3T(eye_x, eye_y, eye_z), + Vec3T(px * inv, py * inv, -inv)); + unsigned char shade = 0; + if (ray.clip(dGrid->indexBBox())) { + nanovdb::Coord ijk; + float v = 0.0f, t = 0.0f; + if (nanovdb::math::ZeroCrossing(ray, acc, ijk, v, t)) { + float gx = acc.getValue(ijk.offsetBy(1, 0, 0)) + - acc.getValue(ijk.offsetBy(-1, 0, 0)); + float gy = acc.getValue(ijk.offsetBy(0, 1, 0)) + - acc.getValue(ijk.offsetBy(0, -1, 0)); + float gz = acc.getValue(ijk.offsetBy(0, 0, 1)) + - acc.getValue(ijk.offsetBy(0, 0, -1)); + const float gl = sqrtf(gx * gx + gy * gy + gz * gz); + if (gl > 0.0f) { gx /= gl; gy /= gl; gz /= gl; } + float s = gx * lx + gy * ly + gz * lz; + if (s < 0.0f) s = 0.0f; + shade = (unsigned char)(255.0f * s); + } + } + image[y * res + x] = shade; +} +""" + + +def _include_options(): + """compile_options(), falling back to $NANOVDB_INCLUDE in a dev tree.""" + opts = list(nanovdb.cuda.compile_options("-std=c++17")) + inc_dir = opts[0][2:] + if not os.path.isdir(inc_dir): + env_inc = os.environ.get("NANOVDB_INCLUDE") + if env_inc and os.path.isdir(env_inc): + opts[0] = f"-I{env_inc}" + else: + print(f"NanoVDB headers not found at {inc_dir!r} (expected in an " + "installed wheel). Set NANOVDB_INCLUDE to the dir that " + "contains nanovdb/NanoVDB.h to run from a source tree.") + return None + return tuple(opts) + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + import cupy as cp + except ImportError: + print("This example requires CuPy. Install it with: pip install cupy") + return + + options = _include_options() + if options is None: + return + + # Level-set sphere built on the device; download once for the host-side + # camera setup (index bounding box). + handle = nanovdb.tools.cuda.createLevelSetSphere( + nanovdb.GridType.Float, radius=100.0) + handle.deviceUpload(0, True) + handle.deviceDownload(0, True) + deviceGrid = handle.deviceGrid(0) + bbox = handle.grid(0).indexBBox() + + import math + dim = [bbox.max[i] + 1 - bbox.min[i] for i in range(3)] + center = [bbox.min[i] + 0.5 * dim[i] for i in range(3)] + eye = (center[0], center[1], center[2] + 2.0 * dim[2]) + tan_fov = math.tan(math.radians(FOV_DEG) * 0.5) + + kernel = cp.RawKernel( + KERNEL_SRC, "render_level_set", options=options, backend="nvrtc") + image = cp.zeros(RES * RES, dtype=cp.uint8) + block = (16, 16) + grid = ((RES + 15) // 16, (RES + 15) // 16) + kernel(grid, block, + (deviceGrid.data_ptr(), image, RES, cp.float32(tan_fov), + cp.float32(eye[0]), cp.float32(eye[1]), cp.float32(eye[2]), + cp.float32(LIGHT[0]), cp.float32(LIGHT[1]), cp.float32(LIGHT[2]))) + cp.cuda.runtime.deviceSynchronize() + + host_image = cp.asnumpy(image) + lit = int((host_image > 0).sum()) + out_dir = tempfile.mkdtemp(prefix="nanovdb_") + path = os.path.join(out_dir, "raytrace_level_set_cuda.pgm") + with open(path, "wb") as f: + f.write(f"P5\n{RES} {RES}\n255\n".encode("ascii")) + f.write(host_image.tobytes()) + print(f"Rendered {RES}x{RES} level set on the GPU in one launch, " + f"{lit} lit pixels -> {path}") + assert lit > 0 + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/examples/sample_from_voxels_cuda.py b/nanovdb/nanovdb/python/examples/sample_from_voxels_cuda.py new file mode 100644 index 0000000000..2603f58188 --- /dev/null +++ b/nanovdb/nanovdb/python/examples/sample_from_voxels_cuda.py @@ -0,0 +1,89 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Trilinearly sample a device grid at arbitrary points, with gradients. + +Demonstrates ``nanovdb.tools.cuda.sampleFromVoxels`` — the device +equivalent of the host ``createTrilinearSampler`` used by +``raytrace_level_set.py`` and ``collide_level_set.py``. Given a device +``FloatGrid`` and an ``(N, 3)`` CuPy array of WORLD-space query points, +it writes trilinear samples into an ``(N,)`` array and (optionally) +analytic gradients into an ``(N, 3)`` array, all on the GPU: + + tools.cuda.sampleFromVoxels(points, dGrid, values, stream) + tools.cuda.sampleFromVoxels(points, dGrid, values, gradients, stream) + +Here we sample a level-set sphere along a ray crossing its surface, +recover the surface normal from the normalized gradient, and cross-check +a few device samples against the host ``createTrilinearSampler``. + +Requires CuPy and a CUDA-capable GPU. + +Run with: python sample_from_voxels_cuda.py +""" +import nanovdb + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + import cupy as cp + except ImportError: + print("This example requires CuPy. Install it with: pip install cupy") + return + + # Level-set sphere (radius 50), read straight onto the device. + radius = 50.0 + host_handle = nanovdb.tools.createLevelSetSphere( + radius=radius, voxelSize=1.0, name="sphere") + import os + import tempfile + tmp = tempfile.NamedTemporaryFile(suffix=".nvdb", delete=False) + tmp.close() + nanovdb.io.writeGrid(tmp.name, host_handle) + try: + handle = nanovdb.io.deviceReadGrid(tmp.name) + finally: + os.unlink(tmp.name) + handle.deviceUpload(0, True) + deviceGrid = handle.deviceGrid(0) + + # Query points marching along +x across the surface (within the + # narrow band, where the SDF is meaningful). + xs = cp.linspace(radius - 3.0, radius + 3.0, 7, dtype=cp.float32) + points = cp.zeros((xs.size, 3), dtype=cp.float32) + points[:, 0] = xs + points = cp.ascontiguousarray(points) + + values = cp.empty(xs.size, dtype=cp.float32) + gradients = cp.empty((xs.size, 3), dtype=cp.float32) + nanovdb.tools.cuda.sampleFromVoxels(points, deviceGrid, values, gradients, 0) + cp.cuda.Stream.null.synchronize() + + print(f"Sampling a radius-{radius:.0f} level set along +x:") + v_host = cp.asnumpy(values) + g_host = cp.asnumpy(gradients) + for xi, vi, gi in zip(cp.asnumpy(xs), v_host, g_host): + gn = (gi[0] ** 2 + gi[1] ** 2 + gi[2] ** 2) ** 0.5 + print(f" x={xi:6.2f} sdf={vi:+.3f} |grad|={gn:.3f}") + + # SDF must increase monotonically as we move outward through the band. + assert (v_host[1:] >= v_host[:-1]).all() + # The sample straddling the surface should be ~0. + assert abs(v_host[xs.size // 2]) < 1.0 + + # Cross-check the device sampler against the host sampler. + host_sampler = nanovdb.math.createTrilinearSampler(host_handle.grid(0)) + mid = float(cp.asnumpy(xs)[xs.size // 2]) + host_val = host_sampler(nanovdb.math.Vec3f(mid, 0.0, 0.0)) + print(f"device vs host at x={mid:.2f}: " + f"{v_host[xs.size // 2]:+.4f} vs {host_val:+.4f}") + assert abs(v_host[xs.size // 2] - host_val) < 1e-3 + + print("OK: device trilinear sampling + gradients match the host sampler") + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/examples/signed_flood_fill_cuda.py b/nanovdb/nanovdb/python/examples/signed_flood_fill_cuda.py new file mode 100644 index 0000000000..86f8e34e2f --- /dev/null +++ b/nanovdb/nanovdb/python/examples/signed_flood_fill_cuda.py @@ -0,0 +1,74 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Signed flood fill of a level set on the device. + +Demonstrates ``nanovdb.tools.cuda.signedFloodFill`` — the device tool +that propagates interior/exterior sign out of a narrow band so the +inactive tiles inside the surface read as negative background and those +outside read as positive background. It runs in place on a device +``FloatGrid`` (or ``DoubleGrid``): + + tools.cuda.signedFloodFill(dGrid, verbose, stream) + +We upload a level-set sphere, run the flood fill on the device, then use +``tools.cuda.sampleFromVoxels`` to confirm the field is consistently +signed (negative inside, positive outside) and that the grid still +validates. + +Requires CuPy and a CUDA-capable GPU. + +Run with: python signed_flood_fill_cuda.py +""" +import os +import tempfile + +import nanovdb + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + import cupy as cp + except ImportError: + print("This example requires CuPy. Install it with: pip install cupy") + return + + radius = 40.0 + host_handle = nanovdb.tools.createLevelSetSphere( + radius=radius, voxelSize=1.0, name="sphere") + tmp = tempfile.NamedTemporaryFile(suffix=".nvdb", delete=False) + tmp.close() + nanovdb.io.writeGrid(tmp.name, host_handle) + try: + handle = nanovdb.io.deviceReadGrid(tmp.name) + finally: + os.unlink(tmp.name) + handle.deviceUpload(0, True) + deviceGrid = handle.deviceGrid(0) + + # Propagate signs across the whole grid on the device, in place. + nanovdb.tools.cuda.signedFloodFill(deviceGrid, False, 0) + print("signedFloodFill: done on the device") + + # The interior sign is only meaningful right at the band; sample just + # inside and just outside the surface and confirm the signs. + points = cp.asarray([[radius - 2.0, 0.0, 0.0], + [radius + 2.0, 0.0, 0.0]], dtype=cp.float32) + values = cp.empty(2, dtype=cp.float32) + nanovdb.tools.cuda.sampleFromVoxels(cp.ascontiguousarray(points), deviceGrid, values, 0) + cp.cuda.Stream.null.synchronize() + inside, outside = (float(v) for v in cp.asnumpy(values)) + print(f" sdf just inside surface = {inside:+.3f} (expect < 0)") + print(f" sdf just outside surface = {outside:+.3f} (expect > 0)") + assert inside < 0.0 < outside + + # The flooded grid must still be structurally valid. + assert nanovdb.tools.cuda.isValid(deviceGrid, nanovdb.CheckMode.Full) + print("OK: device signed flood fill produced a consistent, valid SDF") + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/examples/triton_kernel.py b/nanovdb/nanovdb/python/examples/triton_kernel.py new file mode 100644 index 0000000000..bfc519dbe4 --- /dev/null +++ b/nanovdb/nanovdb/python/examples/triton_kernel.py @@ -0,0 +1,88 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Consume a NanoVDB device buffer from a Triton kernel. + +Like Numba (``numba_cuda.py``), Triton does not parse the C++ NanoVDB +header and has no notion of the ``NanoGrid`` ABI. Triton kernels operate +on flat tensors addressed by pointer + offset, so the natural NanoVDB +interop is over the zero-copy *device buffers* NanoVDB exposes, not the +structured grid. + +Device-pointer ABI recap (see cupy_rawkernel.py for the C++ ABI route): + +* ``handle.deviceGrid(n).data_ptr()`` is a raw DEVICE pointer (Python int) + to a ``nanovdb::NanoGrid``. Triton cannot decode that layout, so + this example does not pass the grid pointer to the kernel directly. +* Triton works through a framework tensor: take the handle's zero-copy + ``__cuda_array_interface__`` / ``__dlpack__`` view (e.g. + ``torch.from_dlpack(handle)``) and hand the resulting tensor to the + kernel. Triton then sees a normal contiguous device tensor. +* Provenance is the caller's responsibility: only device buffers go to + device kernels. + +This example wraps the device grid buffer as a Torch CUDA tensor and runs +a trivial Triton element-wise kernel over it. Requires Triton, PyTorch +(with CUDA), and a CUDA-capable GPU; it skips with a message if any are +missing. + +Run with: python triton_kernel.py +""" +import nanovdb + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + import triton + import triton.language as tl + except ImportError: + print("This example requires Triton. Install it with: " + "pip install triton") + return + try: + import torch + except ImportError: + print("This example requires PyTorch (Triton kernels are launched " + "over Torch tensors). Install it with: pip install torch") + return + if not torch.cuda.is_available(): + print("PyTorch reports no CUDA device available. Skipping.") + return + + # Build a float level-set sphere on the device. + handle = nanovdb.tools.cuda.createLevelSetSphere(nanovdb.GridType.Float, 20) + handle.deviceUpload(0, True) + print(f"Device grid buffer: {handle.size()} bytes at " + f"{hex(handle.device_ptr())}") + + # Zero-copy: Torch adopts the handle's DLPack device buffer (1-D uint8). + buf = torch.from_dlpack(handle) + print(f" torch tensor: shape={tuple(buf.shape)}, dtype={buf.dtype}, " + f"is_cuda={buf.is_cuda}") + + @triton.jit + def count_nonzero(x_ptr, out_ptr, n, BLOCK: tl.constexpr): + pid = tl.program_id(0) + offs = pid * BLOCK + tl.arange(0, BLOCK) + mask = offs < n + x = tl.load(x_ptr + offs, mask=mask, other=0) + tl.atomic_add(out_ptr, tl.sum((x != 0).to(tl.int32))) + + n = buf.numel() + out = torch.zeros(1, dtype=torch.int32, device="cuda") + BLOCK = 1024 + grid = (triton.cdiv(n, BLOCK),) + count_nonzero[grid](buf, out, n, BLOCK=BLOCK) + torch.cuda.synchronize() + print(f" Triton counted {int(out.item())} non-zero bytes in the device " + "grid buffer") + print(" (Decoding the NanoGrid ABI itself is out of scope for Triton; " + "use the CuPy RawKernel route for C++ accessor access — see " + "cupy_rawkernel.py.)") + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/examples/validate_cuda.py b/nanovdb/nanovdb/python/examples/validate_cuda.py new file mode 100644 index 0000000000..ea2dfdc9a9 --- /dev/null +++ b/nanovdb/nanovdb/python/examples/validate_cuda.py @@ -0,0 +1,73 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Validate a grid, and its checksum and stats, entirely on the device. + +GPU counterpart to ``validate.py``. Once a grid is resident on the GPU +(``deviceReadGrid`` + ``deviceUpload``), the quality-control tools run +in place on the device grid — no download required: + +* ``tools.cuda.isValid(g, mode)`` — structural + checksum check. +* ``tools.cuda.evalChecksum(g, mode)`` — compute the CRC checksum. +* ``tools.cuda.updateChecksum(g, mode)`` — recompute it in place. +* ``tools.cuda.validateChecksum(g, mode)`` — verify it matches. +* ``tools.cuda.updateGridStats(g, mode)`` — recompute min/max/avg/bbox. + +``CheckMode`` trades coverage for speed (``Partial`` vs ``Full``); +``StatsMode`` selects how much of the per-node stats to refresh. + +Requires CuPy and a CUDA-capable GPU. + +Run with: python validate_cuda.py +""" +import os +import tempfile + +import nanovdb + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + import cupy # noqa: F401 (only needed to confirm a usable GPU stack) + except ImportError: + print("This example requires CuPy. Install it with: pip install cupy") + return + + host_handle = nanovdb.tools.createLevelSetSphere( + radius=40.0, voxelSize=1.0, name="sphere") + tmp = tempfile.NamedTemporaryFile(suffix=".nvdb", delete=False) + tmp.close() + nanovdb.io.writeGrid(tmp.name, host_handle) + try: + handle = nanovdb.io.deviceReadGrid(tmp.name) + finally: + os.unlink(tmp.name) + handle.deviceUpload(0, True) + deviceGrid = handle.deviceGrid(0) + + # Structural validation on the device, partial and full. + partial = nanovdb.tools.cuda.isValid(deviceGrid, nanovdb.CheckMode.Partial) + full = nanovdb.tools.cuda.isValid(deviceGrid, nanovdb.CheckMode.Full) + print(f"isValid: Partial={partial}, Full={full}") + assert partial and full + + # Checksum round-trip: recompute in place, then verify it matches. + nanovdb.tools.cuda.updateChecksum(deviceGrid, nanovdb.CheckMode.Full) + checksum = nanovdb.tools.cuda.evalChecksum(deviceGrid, nanovdb.CheckMode.Full) + ok = nanovdb.tools.cuda.validateChecksum(deviceGrid, nanovdb.CheckMode.Full) + print(f"checksum: {checksum} validateChecksum(Full)={ok}") + assert ok + + # Recompute per-node statistics on the device. + nanovdb.tools.cuda.updateGridStats(deviceGrid) + print("updateGridStats: OK") + assert nanovdb.tools.cuda.isValid(deviceGrid, nanovdb.CheckMode.Full) + + print("OK: device-side validation, checksum, and stats all pass") + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/examples/voxels_to_grid_cuda.py b/nanovdb/nanovdb/python/examples/voxels_to_grid_cuda.py new file mode 100644 index 0000000000..4c9c579d0b --- /dev/null +++ b/nanovdb/nanovdb/python/examples/voxels_to_grid_cuda.py @@ -0,0 +1,90 @@ +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 +"""Build a NanoVDB grid directly on the GPU from a CuPy coordinate array. + +GPU counterpart to ``build_grid.py`` and a Python port of +``ex_voxels_to_grid_cuda`` (and the device half of +``ex_make_custom_nanovdb_cuda``). Where the host builders write voxels +one at a time through a tree, the device rasterizers take a whole +``(N, 3)`` array already resident on the GPU and construct the grid in +one call, entirely on the device: + +* ``nanovdb.tools.cuda.voxelsToOnIndexGrid(coords, voxelSize, stream)`` + rasterizes ``(N, 3)`` int32 INDEX-space voxel coordinates into a + device ``OnIndexGrid`` handle. +* ``nanovdb.tools.cuda.pointsToGrid(points, voxelSize, stream)`` + rasterizes ``(N, 3)`` float WORLD-space positions into a device + ``NanoGrid`` (the point coordinates are stored as blind data). + +Key facts demonstrated: + +* The returned handle is already DEVICE-resident: ``deviceGrid(0)`` is + valid immediately with no ``deviceUpload``. ``grid(0)`` (the host + grid) is ``None`` until you call ``deviceDownload``. +* The device grid feeds straight into other ``nanovdb.tools.cuda.*`` + ops (here, device validation). + +Requires CuPy and a CUDA-capable GPU. + +Run with: python voxels_to_grid_cuda.py +""" +import nanovdb + + +def main(): + if not (nanovdb.isCudaAvailable() and nanovdb.isGpuAvailable()): + print("This example requires a CUDA build of nanovdb and a GPU. " + "Skipping.") + return + try: + import cupy as cp + except ImportError: + print("This example requires CuPy. Install it with: pip install cupy") + return + + # Build the voxel coordinates of a hollow sphere shell directly on the + # GPU with CuPy, so nothing round-trips through the host. + radius = 32.0 + lin = cp.arange(-40, 41, dtype=cp.float32) + x, y, z = cp.meshgrid(lin, lin, lin, indexing="ij") + r = cp.sqrt(x * x + y * y + z * z) + shell = (cp.abs(r - radius) < 1.0) + coords = cp.stack([x[shell], y[shell], z[shell]], axis=1).astype(cp.int32) + coords = cp.ascontiguousarray(coords) + print(f"Generated {coords.shape[0]} shell voxel coords on the GPU") + + # Rasterize into a device OnIndex grid in one call. + handle = nanovdb.tools.cuda.voxelsToOnIndexGrid(coords, 1.0, 0) + deviceGrid = handle.deviceGrid(0) + print(f"voxelsToOnIndexGrid -> {handle.gridType(0)} handle, " + f"deviceGrid={type(deviceGrid).__name__}") + print(f" grid(0) before download = {handle.grid(0)} (device-built)") + + # Validate entirely on the device, then download for host-side metadata. + print(f" tools.cuda.isValid(deviceGrid) = {nanovdb.tools.cuda.isValid(deviceGrid)}") + handle.deviceDownload(0, True) + host_grid = handle.grid(0) + active = host_grid.activeVoxelCount() + print(f" after deviceDownload: activeVoxelCount={active}, " + f"valueCount={host_grid.valueCount()}") + # OnIndex value 0 is the background slot, so valueCount == active + 1. + assert host_grid.valueCount() == active + 1 + assert active == int(coords.shape[0]) + + # pointsToGrid takes WORLD-space float positions and builds a Point grid. + world_points = coords.astype(cp.float32) * 0.5 # arbitrary world scale + world_points = cp.ascontiguousarray(world_points) + point_handle = nanovdb.tools.cuda.pointsToGrid(world_points, 0.5, 0) + print(f"pointsToGrid -> {point_handle.gridType(0)} handle " + f"(positions stored as blind data)") + assert nanovdb.tools.cuda.isValid(point_handle.deviceGrid(0)) + point_handle.deviceDownload(0, True) + point_active = point_handle.grid(0).activeVoxelCount() + print(f" point grid activeVoxelCount={point_active} (isValid on device)") + assert point_active > 0 + + print("OK: built OnIndex and Point grids on the device from CuPy arrays") + + +if __name__ == "__main__": + main() diff --git a/nanovdb/nanovdb/python/test/TestGpuInterop.py b/nanovdb/nanovdb/python/test/TestGpuInterop.py new file mode 100644 index 0000000000..2299dd6ba1 --- /dev/null +++ b/nanovdb/nanovdb/python/test/TestGpuInterop.py @@ -0,0 +1,1107 @@ +#!/usr/bin/env python +# Copyright Contributors to the OpenVDB Project +# SPDX-License-Identifier: Apache-2.0 + +"""GPU interop unit tests for the NanoVDB Python bindings. + +These exercise the ``nanovdb.cuda`` and ``nanovdb.tools.cuda`` surface: +the CUDA-Array-Interface (CAI v3) and DLPack zero-copy bridges on +``DeviceBuffer`` / ``DeviceGridHandle`` / ``UnifiedBuffer``, the +``grid.data_ptr()`` host/device pointer ABI, raw-stream arguments, the +device-grid morphology / index / QC ops, the device NodeManager and +VoxelBlockManager, and the multi-GPU ``DistributedPointsToGrid`` +pipeline over managed (unified) memory. + +The whole module self-skips when the extension was built without CUDA +or when no CUDA-capable GPU is present, mirroring the device-test +gating in TestNanoVDB.py. Individual tests additionally skip when CuPy +(the only GPU-array framework assumed present) or PyTorch is missing. + +Run directly with: python TestGpuInterop.py -v +""" + +import os +import subprocess +import sys +import tempfile +import unittest + +# If on Windows, add required dll directories from our binary build tree +# (mirrors the bootstrap in TestNanoVDB.py). +if 'add_dll_directory' in dir(os): + for p in os.environ.get('PATH', '').split(os.pathsep): + if os.path.isdir(p): + try: + os.add_dll_directory(p) + except OSError: + pass + +import nanovdb + + +def _require_cupy(test): + """Skip ``test`` (returns the cupy module) unless CuPy is importable.""" + try: + import cupy as cp + except ImportError: + test.skipTest("CuPy not installed") + return cp + + +def _build_device_onindex_grid(radius=20.0): + """Build, write, read-back, and upload a device OnIndex grid. + + Returns (handle, deviceGrid) where handle is a + nanovdb.cuda.DeviceGridHandle that has been deviceUpload()ed and + deviceGrid is the DEVICE OnIndexGrid (data_ptr is a device pointer). + The handle must be kept alive for as long as deviceGrid is used. + """ + h = nanovdb.tools.createLevelSetSphere(radius=radius, voxelSize=1.0) + fg = h.grid(0) + onh = nanovdb.tools.createOnIndexGrid(fg) + tmp = tempfile.NamedTemporaryFile(suffix=".nvdb", delete=False) + tmp.close() + nanovdb.io.writeGrid(tmp.name, onh) + try: + dh = nanovdb.io.deviceReadGrid(tmp.name) + finally: + os.unlink(tmp.name) + dh.deviceUpload(0, True) + return dh, dh.deviceGrid(0) + + +def _device_onindex_from_coords(cp, coords_np): + """Build a device OnIndex grid directly from (N,3) int32 index coords (no + host round-trip; voxelsToOnIndexGrid returns a device grid). Returns + (handle, deviceGrid, activeVoxelCount, coordsByValueIndex) where the last is + a CuPy (count+1, 3) int32 array mapping value index -> coord (row 0 unused).""" + import numpy as np + coords = cp.asarray(np.ascontiguousarray(coords_np, dtype=np.int32)) + dh = nanovdb.tools.cuda.voxelsToOnIndexGrid(coords, 1.0) + dg = dh.deviceGrid(0) + n = int(nanovdb.tools.cuda.buildVoxelBlockManager(dg, 9, 0, 0, 0, 0).lastOffset()) + by_index = cp.empty((n + 1, 3), dtype=cp.int32) + nanovdb.tools.cuda.activeVoxelCoords(dg, by_index) + return dh, dg, n, by_index + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestCompileOptions(unittest.TestCase): + """nanovdb.cuda.compile_options for feeding a runtime CUDA compiler.""" + + def test_include_flag_first(self): + opts = nanovdb.cuda.compile_options() + self.assertIsInstance(opts, tuple) + self.assertEqual(len(opts), 1) + self.assertTrue(opts[0].startswith("-I")) + # The include flag points at the bundled NanoVDB header dir. + self.assertTrue(opts[0].endswith(os.path.join("nanovdb", "include"))) + + def test_extra_flags_appended_in_order(self): + opts = nanovdb.cuda.compile_options("-std=c++17", "-O3") + self.assertEqual(opts[0][:2], "-I") + self.assertEqual(opts[1:], ("-std=c++17", "-O3")) + + def test_top_level_alias_matches(self): + # nanovdb._cuda_compile_options is the underlying closure. + self.assertEqual( + nanovdb.cuda.compile_options(), nanovdb._cuda_compile_options()) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestDeviceBufferInterop(unittest.TestCase): + """DeviceBuffer CAI v3 / DLPack and from_external wrapping.""" + + def test_from_external_wraps_managed_memory(self): + cp = _require_cupy(self) + prev = cp.cuda.get_allocator() + cp.cuda.set_allocator(cp.cuda.malloc_managed) + try: + buf = cp.zeros(256, dtype=cp.uint8) + gpu_ptr = int(buf.data.ptr) + # Managed memory: a single pointer is valid on both host and + # device, so host_ptr == device_ptr is legal here. + ext = nanovdb.cuda.DeviceBuffer.from_external(256, gpu_ptr, gpu_ptr) + self.assertEqual(ext.size(), 256) + self.assertEqual(ext.device_ptr(), gpu_ptr) + self.assertEqual(ext.host_ptr(), gpu_ptr) + finally: + cp.cuda.set_allocator(prev) + + def test_from_external_rejects_null_device_pointer(self): + # gpu_ptr == 0 is a usage error: there is no device memory to wrap. + with self.assertRaises(ValueError): + nanovdb.cuda.DeviceBuffer.from_external(256, 0, 0) + with self.assertRaises(ValueError): + nanovdb.cuda.DeviceBuffer.from_external(256, 0, 12345) + + def test_cuda_array_interface_v3(self): + cp = _require_cupy(self) + prev = cp.cuda.get_allocator() + cp.cuda.set_allocator(cp.cuda.malloc_managed) + try: + buf = cp.zeros(256, dtype=cp.uint8) + gpu_ptr = int(buf.data.ptr) + ext = nanovdb.cuda.DeviceBuffer.from_external(256, gpu_ptr, gpu_ptr) + cai = ext.__cuda_array_interface__ + self.assertEqual(cai["version"], 3) + self.assertEqual(cai["typestr"], "|u1") + self.assertEqual(cai["shape"], (256,)) + self.assertEqual(cai["data"][0], gpu_ptr) + # cupy zero-copy view aliases the same device pointer. + arr = cp.asarray(ext) + self.assertEqual(arr.nbytes, ext.size()) + self.assertEqual(int(arr.data.ptr), ext.device_ptr()) + finally: + cp.cuda.set_allocator(prev) + + def test_dlpack_round_trip(self): + cp = _require_cupy(self) + prev = cp.cuda.get_allocator() + cp.cuda.set_allocator(cp.cuda.malloc_managed) + try: + buf = cp.zeros(256, dtype=cp.uint8) + gpu_ptr = int(buf.data.ptr) + ext = nanovdb.cuda.DeviceBuffer.from_external(256, gpu_ptr, gpu_ptr) + # kDLCUDA == 2 + self.assertEqual(ext.__dlpack_device__()[0], 2) + arr = cp.from_dlpack(ext) + self.assertEqual(arr.nbytes, ext.size()) + self.assertEqual(int(arr.data.ptr), ext.device_ptr()) + finally: + cp.cuda.set_allocator(prev) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestDeviceGridHandleInterop(unittest.TestCase): + """DeviceGridHandle CAI/DLPack, the host/device data_ptr ABI, and + from_buffer zero-copy adoption.""" + + def test_cai_exposes_device_buffer_after_upload(self): + dh, dg = _build_device_onindex_grid(20.0) + # After upload the whole device buffer is exposed as 1-D uint8. + cai = dh.__cuda_array_interface__ + self.assertEqual(cai["version"], 3) + self.assertEqual(cai["typestr"], "|u1") + self.assertEqual(cai["shape"], (dh.size(),)) + self.assertNotEqual(dh.device_ptr(), 0) + self.assertEqual(cai["data"][0], dh.device_ptr()) + + def test_cupy_zero_copy_aliases_device_ptr(self): + cp = _require_cupy(self) + dh, dg = _build_device_onindex_grid(20.0) + arr = cp.asarray(dh) + self.assertEqual(arr.nbytes, dh.size()) + self.assertEqual(int(arr.data.ptr), dh.device_ptr()) + # The typed device grid's data_ptr() points at the same buffer base. + self.assertEqual(dg.data_ptr(), dh.device_ptr()) + + def test_dlpack_zero_copy(self): + cp = _require_cupy(self) + dh, dg = _build_device_onindex_grid(20.0) + arr = cp.from_dlpack(dh) + self.assertEqual(arr.nbytes, dh.size()) + self.assertEqual(int(arr.data.ptr), dh.device_ptr()) + + def test_host_vs_device_data_ptr_ambiguity(self): + # A single handle exposes BOTH a host grid() and a device + # deviceGrid() after upload; the data_ptr() values differ and the + # grid object cannot tell host from device — provenance is the + # caller's responsibility (see the docstring on data_ptr()). + dh, dg = _build_device_onindex_grid(20.0) + hg = dh.grid(0) + self.assertIsNotNone(hg) + self.assertIsNotNone(dg) + self.assertNotEqual(hg.data_ptr(), dg.data_ptr()) + self.assertEqual(dg.data_ptr(), dh.device_ptr()) + + def test_host_accessor_on_device_grid_segfaults_in_subprocess(self): + # Calling a HOST-side accessor on a grid from deviceGrid(n) + # dereferences device memory on the CPU and crashes the process. + # Run it in a subprocess so the SIGSEGV does not abort the test + # session, and assert the documented crash contract. + code = ( + "import nanovdb, tempfile, os\n" + "h = nanovdb.tools.createLevelSetSphere(radius=20.0, voxelSize=1.0)\n" + "fg = h.grid(0)\n" + "onh = nanovdb.tools.createOnIndexGrid(fg)\n" + "t = tempfile.NamedTemporaryFile(suffix='.nvdb', delete=False)\n" + "t.close()\n" + "nanovdb.io.writeGrid(t.name, onh)\n" + "dh = nanovdb.io.deviceReadGrid(t.name)\n" + "os.unlink(t.name)\n" + "dh.deviceUpload(0, True)\n" + "dg = dh.deviceGrid(0)\n" + "acc = dg.getAccessor()\n" + "acc.getValue(nanovdb.math.Coord(0, 0, 0))\n" + ) + r = subprocess.run([sys.executable, "-c", code], capture_output=True) + # Killed by a signal => negative return code; -11 is SIGSEGV. Any + # crash (negative) satisfies the "this is illegal" contract. + self.assertLess( + r.returncode, 0, + "host accessor on a device grid was expected to crash, " + f"got returncode {r.returncode}") + + def test_from_buffer_adopts_managed_buffer(self): + cp = _require_cupy(self) + dh, dg = _build_device_onindex_grid(20.0) + src = cp.asarray(dh) # zero-copy device view of a valid grid + prev = cp.cuda.get_allocator() + cp.cuda.set_allocator(cp.cuda.malloc_managed) + try: + # Managed buffer holding a copy of the valid grid bytes; a + # single pointer serves as both host and device pointer. + mbuf = cp.empty(int(dh.size()), dtype=cp.uint8) + mbuf[:] = src + cp.cuda.runtime.deviceSynchronize() + ptr = int(mbuf.data.ptr) + ext = nanovdb.cuda.DeviceBuffer.from_external( + int(dh.size()), ptr, ptr) + adopted = nanovdb.cuda.DeviceGridHandle.from_buffer(ext) + self.assertEqual(adopted.gridType(0), nanovdb.GridType.OnIndex) + self.assertEqual(adopted.size(), dh.size()) + # from_buffer MOVES the buffer; the source is left empty. + self.assertEqual(ext.size(), 0) + finally: + cp.cuda.set_allocator(prev) + + def test_from_buffer_rejects_garbage(self): + cp = _require_cupy(self) + prev = cp.cuda.get_allocator() + cp.cuda.set_allocator(cp.cuda.malloc_managed) + try: + mbuf = cp.zeros(512, dtype=cp.uint8) # not a NanoVDB grid header + ptr = int(mbuf.data.ptr) + ext = nanovdb.cuda.DeviceBuffer.from_external(512, ptr, ptr) + with self.assertRaises(RuntimeError): + nanovdb.cuda.DeviceGridHandle.from_buffer(ext) + finally: + cp.cuda.set_allocator(prev) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestRecordUse(unittest.TestCase): + """DeviceBuffer/DeviceGridHandle.recordUse: order the buffer's device free + after work enqueued on external non-blocking streams (the CAI/DLPack + interop pattern, where kernels launched against device_ptr() are invisible + to the buffer's automatic upload/download tracking).""" + + def test_free_ordered_after_external_stream_work(self): + cp = _require_cupy(self) + dh, _ = _build_device_onindex_grid(20.0) + # Baseline: reduce the serialized grid bytes while the handle is alive. + view = cp.asarray(dh) # zero-copy CAI view of the device buffer + expected = int(view.sum()) + # Enqueue the same reduction on a NON-BLOCKING stream, record the use, + # and drop the handle while the reduction may still be in flight. The + # buffer's cudaFreeAsync must be ordered after the recorded event; if + # it were not, the reduction would race the free and read freed memory. + s = cp.cuda.Stream(non_blocking=True) + with s: + pending = view.sum() # device scalar; no host sync yet + dh.recordUse(s.ptr) + del dh, view + s.synchronize() + self.assertEqual(int(pending), expected) + + def test_default_stream_and_explicit_device(self): + dh, _ = _build_device_onindex_grid(10.0) + # Default stream (0) and the current device, both explicit and implied. + dh.recordUse(0) + dh.recordUse(0, device=0) + + def test_rejects_out_of_range_device(self): + dh, _ = _build_device_onindex_grid(10.0) + with self.assertRaises(IndexError): + dh.recordUse(0, device=1_000_000) + + def test_noop_on_non_owning_buffer(self): + cp = _require_cupy(self) + prev = cp.cuda.get_allocator() + cp.cuda.set_allocator(cp.cuda.malloc_managed) + try: + mbuf = cp.zeros(256, dtype=cp.uint8) + ptr = int(mbuf.data.ptr) + ext = nanovdb.cuda.DeviceBuffer.from_external(256, ptr, ptr) + # Non-owning buffers never free their pointers, so recordUse has + # nothing to order — it must be accepted and be a no-op. + ext.recordUse(0) + finally: + cp.cuda.set_allocator(prev) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestBindingValidation(unittest.TestCase): + """Negative tests: undersized value-indexed arrays and invalid geometric + parameters must raise Python exceptions instead of out-of-bounds device + accesses or silently-invalid grid transforms.""" + + BAD_SIZES = (0.0, -1.0, float("nan"), float("inf")) + + def _grid(self, cp): + import numpy as np + ijk = np.stack(np.meshgrid(*([np.arange(8)] * 3), indexing="ij"), + axis=-1).reshape(-1, 3).astype(np.int32) + return _device_onindex_from_coords(cp, ijk) # (dh, dg, n, coords) + + def test_voxels_to_grid_rejects_invalid_voxel_size(self): + cp = _require_cupy(self) + coords = cp.zeros((4, 3), dtype=cp.int32) + for bad in self.BAD_SIZES: + with self.assertRaises(ValueError): + nanovdb.tools.cuda.voxelsToOnIndexGrid(coords, bad) + + def test_mesh_to_grid_rejects_invalid_geometry(self): + cp = _require_cupy(self) + pts = cp.asarray( + [[0.0, 0.0, 0.0], [8.0, 0.0, 0.0], [0.0, 8.0, 0.0]], dtype=cp.float32) + tris = cp.asarray([[0, 1, 2]], dtype=cp.int32) + for bad in self.BAD_SIZES: + with self.assertRaises(ValueError): + nanovdb.tools.cuda.meshToGrid(pts, tris, voxelSize=bad) + with self.assertRaises(ValueError): + nanovdb.tools.cuda.meshToGrid(pts, tris, halfWidth=bad) + + def test_index_to_grid_rejects_short_values(self): + cp = _require_cupy(self) + dh, dg, n, _ = self._grid(cp) + with self.assertRaises(ValueError): + nanovdb.tools.cuda.indexToGrid(dg, cp.zeros(n, dtype=cp.float32)) + # Exact valueCount (= n + 1) is accepted. + out = nanovdb.tools.cuda.indexToGrid(dg, cp.zeros(n + 1, dtype=cp.float32)) + self.assertEqual(out.gridCount(), 1) + + def test_inject_rejects_short_sidecars(self): + cp = _require_cupy(self) + dh, dg, n, _ = self._grid(cp) + good = cp.zeros(n + 1, dtype=cp.float32) + short = cp.zeros(n, dtype=cp.float32) + with self.assertRaises(ValueError): + nanovdb.tools.cuda.inject(dg, dg, short, good) + with self.assertRaises(ValueError): + nanovdb.tools.cuda.inject(dg, dg, good, short) + # 2-D (feature) overload: row counts are checked the same way. + good2 = cp.zeros((n + 1, 2), dtype=cp.float32) + short2 = cp.zeros((n, 2), dtype=cp.float32) + with self.assertRaises(ValueError): + nanovdb.tools.cuda.inject(dg, dg, short2, good2) + + def test_predicate_to_mask_rejects_short_predicate(self): + cp = _require_cupy(self) + dh, dg, n, _ = self._grid(cp) + masks = cp.zeros((n + 1) * 8, dtype=cp.uint64) + with self.assertRaises(ValueError): + nanovdb.tools.cuda.injectPredicateToMask( + dg, cp.zeros(n, dtype=cp.bool_), masks) + + def test_prune_grid_rejects_short_mask(self): + cp = _require_cupy(self) + import numpy as np + # 16^3 coords span 8 leaf nodes (2 per axis), so one 8-word Mask<3> is + # a valid multiple of 8 but covers only 1 of the 8 leaves — previously + # an out-of-bounds device read rather than a Python error. + ijk = np.stack(np.meshgrid(*([np.arange(16)] * 3), indexing="ij"), + axis=-1).reshape(-1, 3).astype(np.int32) + dh, dg, n, _ = _device_onindex_from_coords(cp, ijk) + with self.assertRaises(ValueError): + nanovdb.tools.cuda.pruneGrid(dg, cp.zeros(8, dtype=cp.uint64)) + + def test_gather_and_coords_reject_short_arrays(self): + cp = _require_cupy(self) + dh, dg, n, _ = self._grid(cp) + good_vals = cp.zeros(n + 1, dtype=cp.float32) + with self.assertRaises(ValueError): + nanovdb.tools.cuda.gatherBoxStencil( + dg, cp.zeros(n, dtype=cp.float32), cp.zeros((n + 1, 27), dtype=cp.float32)) + with self.assertRaises(ValueError): + nanovdb.tools.cuda.gatherBoxStencil( + dg, good_vals, cp.zeros((n, 27), dtype=cp.float32)) + with self.assertRaises(ValueError): + nanovdb.tools.cuda.gatherBoxStencilColumns( + dg, good_vals, cp.zeros((n, 2), dtype=cp.float32), + cp.asnumpy(cp.asarray([13, 4], dtype=cp.int32))) + with self.assertRaises(ValueError): + nanovdb.tools.cuda.activeVoxelCoords(dg, cp.zeros((n, 3), dtype=cp.int32)) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestExportStreamOrdering(unittest.TestCase): + """CAI / DLPack exports must order the consumer against the buffer's + tracked prior uses (async uploads on non-blocking streams).""" + + def test_cai_after_async_upload(self): + cp = _require_cupy(self) + h = nanovdb.tools.createLevelSetSphere(radius=20.0, voxelSize=1.0) + tmp = tempfile.NamedTemporaryFile(suffix=".nvdb", delete=False) + tmp.close() + try: + nanovdb.io.writeGrid(tmp.name, h) + dh = nanovdb.io.deviceReadGrid(tmp.name) + finally: + os.unlink(tmp.name) + s = cp.cuda.Stream(non_blocking=True) + # Async upload on a non-blocking stream, then consume via CAI without + # ever synchronizing s explicitly: the export's stream ordering (CAI + # stream=1 backed by the tracked upload event) must make this safe. + dh.deviceUpload(s.ptr, False) + view = cp.asarray(dh) + host = bytes(cp.asnumpy(view)) + self.assertEqual(len(host), dh.size()) + self.assertNotEqual(host.count(0), len(host)) # real grid bytes arrived + + def test_dlpack_accepts_protocol_stream_values(self): + cp = _require_cupy(self) + h = nanovdb.tools.createLevelSetSphere(radius=10.0, voxelSize=1.0) + tmp = tempfile.NamedTemporaryFile(suffix=".nvdb", delete=False) + tmp.close() + try: + nanovdb.io.writeGrid(tmp.name, h) + dh = nanovdb.io.deviceReadGrid(tmp.name) + finally: + os.unlink(tmp.name) + dh.deviceUpload(0, True) + s = cp.cuda.Stream(non_blocking=True) + # None / legacy default / per-thread / no-sync / raw handle must all + # be accepted per the DLPack protocol. + for stream in (None, 1, 2, -1, s.ptr): + capsule = dh.__dlpack__(stream=stream) + self.assertIsNotNone(capsule) + arr = cp.from_dlpack(dh) + self.assertEqual(arr.nbytes, dh.size()) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestUnifiedBufferInterop(unittest.TestCase): + """UnifiedBuffer create / capacity / zero-copy.""" + + def test_create_size_equals_capacity(self): + ub = nanovdb.cuda.UnifiedBuffer.create(1024) + self.assertEqual(ub.size(), 1024) + self.assertEqual(ub.capacity(), 1024) + self.assertFalse(ub.isEmpty()) + + def test_create_with_reserved_capacity(self): + ub = nanovdb.cuda.UnifiedBuffer.create(1024, 4096) + self.assertEqual(ub.size(), 1024) + self.assertEqual(ub.capacity(), 4096) + + def test_resize_within_capacity_keeps_page_table(self): + ub = nanovdb.cuda.UnifiedBuffer.create(1024, 4096) + ub.resize(2048) + self.assertEqual(ub.size(), 2048) + self.assertEqual(ub.capacity(), 4096) + + def test_cupy_zero_copy(self): + cp = _require_cupy(self) + ub = nanovdb.cuda.UnifiedBuffer.create(2048) + cai = ub.__cuda_array_interface__ + self.assertEqual(cai["version"], 3) + self.assertEqual(cai["shape"], (2048,)) + arr = cp.asarray(ub) + self.assertEqual(arr.nbytes, ub.size()) + arr2 = cp.from_dlpack(ub) + self.assertEqual(arr2.nbytes, ub.size()) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestPointAndVoxelRasterizers(unittest.TestCase): + """tools.cuda point / voxel rasterizers, including raw-stream args.""" + + def test_points_to_grid_world_float(self): + cp = _require_cupy(self) + import numpy as np + pts = cp.asarray(np.array([[0, 0, 0], [1, 1, 1], [5, 5, 5]], + dtype=np.float32)) + h = nanovdb.tools.cuda.pointsToGrid(pts, 1.0, 0) + self.assertEqual(h.gridType(0), nanovdb.GridType.PointIndex) + + def test_points_to_grid_world_double(self): + cp = _require_cupy(self) + import numpy as np + pts = cp.asarray(np.array([[0, 0, 0], [1, 1, 1]], dtype=np.float64)) + h = nanovdb.tools.cuda.pointsToGrid(pts, 1.0, 0) + self.assertEqual(h.gridType(0), nanovdb.GridType.PointIndex) + + def test_voxels_to_index_family(self): + cp = _require_cupy(self) + import numpy as np + vox = cp.asarray(np.array([[0, 0, 0], [1, 0, 0], [0, 1, 0]], + dtype=np.int32)) + self.assertEqual( + nanovdb.tools.cuda.voxelsToOnIndexGrid(vox, 1.0, 0).gridType(0), + nanovdb.GridType.OnIndex) + self.assertEqual( + nanovdb.tools.cuda.voxelsToIndexGrid(vox, 1.0, 0).gridType(0), + nanovdb.GridType.Index) + self.assertEqual( + nanovdb.tools.cuda.voxelsToRGBA8Grid(vox, 1.0, 0).gridType(0), + nanovdb.GridType.RGBA8) + self.assertEqual( + nanovdb.tools.cuda.pointsToRGBA8Grid(vox, 1.0, 0).gridType(0), + nanovdb.GridType.RGBA8) + + def test_non_default_stream_accepted(self): + cp = _require_cupy(self) + import numpy as np + pts = cp.asarray(np.array([[0, 0, 0], [2, 2, 2]], dtype=np.float32)) + stream = cp.cuda.Stream(non_blocking=True) + # Stream args everywhere are raw CUDA stream handles as Python ints. + h = nanovdb.tools.cuda.pointsToGrid(pts, 1.0, stream.ptr) + self.assertEqual(h.gridType(0), nanovdb.GridType.PointIndex) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestDeviceGridOps(unittest.TestCase): + """tools.cuda morphology / QC / sampling on device grids.""" + + def test_morphology_returns_onindex(self): + dh, dg = _build_device_onindex_grid(20.0) + # op 26 == NN_FACE_EDGE_VERTEX (verified). 18 is unimplemented. + self.assertEqual( + nanovdb.tools.cuda.dilateGrid(dg, 26, 0).gridType(0), + nanovdb.GridType.OnIndex) + self.assertEqual( + nanovdb.tools.cuda.coarsenGrid(dg, 0).gridType(0), + nanovdb.GridType.OnIndex) + self.assertEqual( + nanovdb.tools.cuda.refineGrid(dg, 0).gridType(0), + nanovdb.GridType.OnIndex) + + def test_is_valid_on_device_grid(self): + dh, dg = _build_device_onindex_grid(20.0) + self.assertTrue(nanovdb.tools.cuda.isValid(dg)) + + def test_signed_flood_fill_in_place(self): + h = nanovdb.tools.cuda.createLevelSetSphere(nanovdb.GridType.Float, 20) + h.deviceUpload(0, True) + dg = h.deviceGrid(0) + self.assertIsNotNone(dg) + # In place on the device grid; no return value. + self.assertIsNone(nanovdb.tools.cuda.signedFloodFill(dg)) + + def test_sample_from_voxels_float(self): + cp = _require_cupy(self) + import numpy as np + h = nanovdb.tools.cuda.createLevelSetSphere(nanovdb.GridType.Float, 20) + h.deviceUpload(0, True) + dg = h.deviceGrid(0) + pts = cp.asarray(np.array([[0, 0, 0], [20, 0, 0], [10, 0, 0]], + dtype=np.float32)) + vals = cp.empty(3, dtype=cp.float32) + nanovdb.tools.cuda.sampleFromVoxels(pts, dg, vals, 0) + out = cp.asnumpy(vals) + self.assertEqual(out.shape, (3,)) + + def test_sample_from_voxels_with_gradients(self): + cp = _require_cupy(self) + import numpy as np + h = nanovdb.tools.cuda.createLevelSetSphere(nanovdb.GridType.Float, 20) + h.deviceUpload(0, True) + dg = h.deviceGrid(0) + pts = cp.asarray(np.array([[10, 0, 0]], dtype=np.float32)) + vals = cp.empty(1, dtype=cp.float32) + grads = cp.empty((1, 3), dtype=cp.float32) + nanovdb.tools.cuda.sampleFromVoxels(pts, dg, vals, grads, 0) + self.assertEqual(cp.asnumpy(grads).shape, (1, 3)) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestDeviceNodeManager(unittest.TestCase): + """createDeviceNodeManager over a device OnIndex grid.""" + + def test_node_manager_handle(self): + dh, dg = _build_device_onindex_grid(20.0) + nmh = nanovdb.cuda.createDeviceNodeManager(dg, 0) + self.assertIsNotNone(nmh) + self.assertGreater(nmh.size(), 0) + mgr = nmh.mgr() + # mgr() returns the TYPED device NodeManager, not a raw int. + self.assertIsNotNone(mgr) + self.assertNotIsInstance(mgr, int) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestDeviceVoxelBlockManager(unittest.TestCase): + """buildVoxelBlockManager on a device OnIndex grid + zero-copy buffers.""" + + def test_block_manager_properties(self): + dh, dg = _build_device_onindex_grid(100.0) + vbm = nanovdb.tools.cuda.buildVoxelBlockManager(dg, 6, 0, 0, 0, 0) + self.assertGreater(vbm.blockCount(), 0) + # blockWidth / log2BlockWidth / jumpMapLength are PROPERTIES. + self.assertEqual(vbm.blockWidth, 64) + self.assertEqual(vbm.log2BlockWidth, 6) + # firstOffset == 1 (mod BlockWidth) for a full grid. + self.assertEqual(vbm.firstOffset(), 1) + self.assertGreater(vbm.lastOffset(), 0) + + def test_first_leaf_id_zero_copy(self): + cp = _require_cupy(self) + dh, dg = _build_device_onindex_grid(100.0) + vbm = nanovdb.tools.cuda.buildVoxelBlockManager(dg, 6, 0, 0, 0, 0) + # firstLeafID() / jumpMap() return DLPack capsules directly. + fl = cp.from_dlpack(vbm.firstLeafID()) + self.assertEqual(fl.dtype, cp.uint32) + self.assertEqual(fl.shape[0], vbm.blockCount()) + self.assertEqual(int(fl.data.ptr), vbm.first_leaf_id_ptr()) + + def test_jump_map_zero_copy(self): + cp = _require_cupy(self) + dh, dg = _build_device_onindex_grid(100.0) + vbm = nanovdb.tools.cuda.buildVoxelBlockManager(dg, 6, 0, 0, 0, 0) + jm = cp.from_dlpack(vbm.jumpMap()) + self.assertEqual(jm.dtype, cp.uint64) + self.assertEqual(jm.shape[0], vbm.blockCount()) + self.assertEqual(int(jm.data.ptr), vbm.jump_map_ptr()) + + def test_gather_box_stencil_dtypes(self): + """gatherBoxStencil works for float32 AND the integer payload dtypes + (int32/uint32). The centre spoke (column 13) is each voxel's own value, + inactive neighbours read values[0] (=0 here), and gathering arange yields + the SAME neighbour-index pattern in every dtype -- the property that lets + callers gather a neighbour-INDEX table directly in int32. + + Built from explicit coords (voxelsToOnIndexGrid) so active-voxel indexing + is contiguous -- the invariant the VoxelBlockManager gather relies on.""" + cp = _require_cupy(self) + import numpy as np + block = np.array([(i, j, k) for i in range(5) for j in range(5) for k in range(5)], + dtype=np.int32) + _h, dg, n, _co = _device_onindex_from_coords(cp, block) + N = n + 1 + ref = None + for dt in (cp.float32, cp.int32, cp.uint32): + idx = cp.arange(N, dtype=dt) + out = cp.empty((N, 27), dtype=dt) + nanovdb.tools.cuda.gatherBoxStencil(dg, idx, out) + self.assertEqual(out.dtype, cp.dtype(dt)) + self.assertTrue(bool((out[1:N, 13] == idx[1:N]).all())) # centre == self + self.assertTrue(bool((out[1:N] == 0).any())) # inactive -> values[0] + # row 0 (background slot) is never written by the kernel, so compare + # only the written rows across dtypes. + as_i64 = out[1:N].astype(cp.int64) + if ref is None: + ref = as_i64 + else: + self.assertTrue(bool((as_i64 == ref).all())) # dtype-independent + + def test_gather_rejects_noncontiguous_grid(self): + """gatherBoxStencil / activeVoxelCoords raise a clear ValueError (rather + than an illegal memory access) on a grid whose active-voxel indexing is + NOT contiguous -- e.g. createOnIndexGrid's default includeStats / + includeTiles, which add value slots past activeVoxelCount.""" + cp = _require_cupy(self) + dh, dg = _build_device_onindex_grid(20.0) # createOnIndexGrid defaults: stats+tiles + N = int(nanovdb.tools.cuda.buildVoxelBlockManager(dg, 9, 0, 0, 0, 0).lastOffset()) + 1 + with self.assertRaises(ValueError): + nanovdb.tools.cuda.gatherBoxStencil(dg, cp.arange(N, dtype=cp.int32), + cp.empty((N, 27), dtype=cp.int32)) + with self.assertRaises(ValueError): + nanovdb.tools.cuda.activeVoxelCoords(dg, cp.empty((N, 3), dtype=cp.int32)) + + def test_gather_box_stencil_columns(self): + """gatherBoxStencilColumns gathers a CHOSEN subset of the 27 spokes into + an (N,K) table -- equal to the matching columns of the full gather -- and + validates its (host) spoke list + the out shape.""" + cp = _require_cupy(self) + import numpy as np + block = np.array([(i, j, k) for i in range(5) for j in range(5) for k in range(5)], + dtype=np.int32) + _h, dg, n, _co = _device_onindex_from_coords(cp, block) + N = n + 1 + idx = cp.arange(N, dtype=cp.int32) + full = cp.empty((N, 27), dtype=cp.int32) + nanovdb.tools.cuda.gatherBoxStencil(dg, idx, full) + spokes = np.array([13, 22, 16, 14, 4, 10, 12], dtype=np.int32) # centre + 6 faces + sub = cp.empty((N, spokes.shape[0]), dtype=cp.int32) + nanovdb.tools.cuda.gatherBoxStencilColumns(dg, idx, sub, spokes) + self.assertTrue(bool((sub[1:N] == full[1:N][:, cp.asarray(spokes)]).all())) + with self.assertRaises(ValueError): # spoke out of [0, 27) + nanovdb.tools.cuda.gatherBoxStencilColumns( + dg, idx, cp.empty((N, 2), dtype=cp.int32), np.array([13, 99], np.int32)) + with self.assertRaises(ValueError): # out.shape[1] != len(spokes) + nanovdb.tools.cuda.gatherBoxStencilColumns(dg, idx, cp.empty((N, 3), dtype=cp.int32), spokes) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestDeviceInject(unittest.TestCase): + """inject / injectFeatures copy a sidecar over the src/dst voxel + intersection, for float AND the integer payloads (int32/uint32) -- the copy + is via the assignment operator, so it is type-generic.""" + + # src = a 4^3 block; dst = a 6^3 block (a superset), so the intersection is + # exactly src. enc(coord) gives each voxel a distinct in-range value. + SRC = [(i, j, k) for i in range(4) for j in range(4) for k in range(4)] + DST = [(i, j, k) for i in range(6) for j in range(6) for k in range(6)] + SENT = 255 # distinct from every enc value; exact in f32/i32/u32 + + @staticmethod + def _enc(c): # (...,3) int -> (...) int + return (c[..., 0] * 6 + c[..., 1]) * 6 + c[..., 2] + + def test_inject_scalar_dtypes(self): + cp = _require_cupy(self) + import numpy as np + _sh, sdg, ns, sco = _device_onindex_from_coords(cp, np.array(self.SRC, np.int32)) + _dh, ddg, nd, dco = _device_onindex_from_coords(cp, np.array(self.DST, np.int32)) + sco_h = cp.asnumpy(sco[1:]).astype(np.int64) + dco_h = cp.asnumpy(dco[1:]).astype(np.int64) + in_src = (dco_h < 4).all(axis=1) + expect = np.where(in_src, self._enc(dco_h), self.SENT) + for dt in (cp.float32, cp.int32, cp.uint32): + with self.subTest(dtype=np.dtype(dt).name): + src_vals = cp.zeros(ns + 1, dtype=dt) + src_vals[1:] = cp.asarray(self._enc(sco_h), dtype=dt) + dst_vals = cp.full(nd + 1, self.SENT, dtype=dt) + nanovdb.tools.cuda.inject(sdg, ddg, src_vals, dst_vals) + self.assertEqual(dst_vals.dtype, cp.dtype(dt)) + got = cp.asnumpy(dst_vals[1:]).astype(np.int64) + self.assertTrue(bool((got == expect).all())) + + def test_inject_features_dtypes(self): + """injectFeatures: the 2-D (value count, dim) form, integer + float.""" + cp = _require_cupy(self) + import numpy as np + _sh, sdg, ns, sco = _device_onindex_from_coords(cp, np.array(self.SRC, np.int32)) + _dh, ddg, nd, dco = _device_onindex_from_coords(cp, np.array(self.DST, np.int32)) + sco_h = cp.asnumpy(sco[1:]).astype(np.int64) + dco_h = cp.asnumpy(dco[1:]).astype(np.int64) + in_src = (dco_h < 4).all(axis=1) + enc_d = self._enc(dco_h) + for dt in (cp.float32, cp.int32): + with self.subTest(dtype=np.dtype(dt).name): + src_vals = cp.zeros((ns + 1, 2), dtype=dt) + e = self._enc(sco_h) + src_vals[1:, 0] = cp.asarray(e, dtype=dt) + src_vals[1:, 1] = cp.asarray(e + 1000, dtype=dt) # second channel + dst_vals = cp.full((nd + 1, 2), self.SENT, dtype=dt) + nanovdb.tools.cuda.inject(sdg, ddg, src_vals, dst_vals) + got = cp.asnumpy(dst_vals[1:]).astype(np.int64) + exp0 = np.where(in_src, enc_d, self.SENT) + exp1 = np.where(in_src, enc_d + 1000, self.SENT) + self.assertTrue(bool((got[:, 0] == exp0).all())) + self.assertTrue(bool((got[:, 1] == exp1).all())) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestDeviceTypedSidecars(unittest.TestCase): + """Integer payloads / build types: addBlindData with int32/int64 blind data, + and indexToGrid materialising an Int32 destination grid.""" + + BLOCK = [(i, j, k) for i in range(4) for j in range(4) for k in range(4)] + + def test_add_blind_data_int_dtypes(self): + cp = _require_cupy(self) + import numpy as np + _h, dg, n, _co = _device_onindex_from_coords(cp, np.array(self.BLOCK, np.int32)) + for dt in (cp.int32, cp.int64, cp.float32): # float32 = regression + with self.subTest(dtype=np.dtype(dt).name): + blind = (cp.arange(n + 1, dtype=dt) * 3) + out_dh = nanovdb.tools.cuda.addBlindData( + dg, blind, nanovdb.GridBlindDataClass.ChannelArray, + nanovdb.GridBlindDataSemantic.Unknown, "labels") + out_dh.deviceDownload(0, True) + g = out_dh.grid(0) + self.assertEqual(g.blindDataCount(), 1) + bd = np.asarray(g.getBlindData(0)) + self.assertEqual(bd.dtype, np.dtype(dt)) + self.assertTrue(bool((bd == cp.asnumpy(blind)).all())) + + def test_index_to_grid_int32(self): + cp = _require_cupy(self) + import numpy as np + _h, dg, n, _co = _device_onindex_from_coords(cp, np.array(self.BLOCK, np.int32)) + vals = cp.arange(n + 1, dtype=cp.int32) * 7 # int32 sidecar -> Int32Grid + out_dh = nanovdb.tools.cuda.indexToGrid(dg, vals) + self.assertEqual(out_dh.gridType(0), nanovdb.GridType.Int32) + out_dh.deviceDownload(0, True) + self.assertEqual(out_dh.grid(0).activeVoxelCount(), n) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestDeviceGridMetadata(unittest.TestCase): + """tools.cuda.setGridClass: retag a device grid's GridClass in place, + refreshing its checksum, with the change surviving download and file I/O.""" + + BLOCK = [(i, j, k) for i in range(4) for j in range(4) for k in range(4)] + + def _float_grid_from_block(self, cp): + """OnIndex(BLOCK) + per-voxel SDF -> device FloatGrid handle. The float + grid inherits the source's IndexGrid class (what setGridClass overrides).""" + import numpy as np + _h, dg, n, _co = _device_onindex_from_coords(cp, np.array(self.BLOCK, np.int32)) + sdf = cp.arange(n + 1, dtype=cp.float32) * 0.25 + float_dh = nanovdb.tools.cuda.indexToGrid(dg, sdf) + self.assertEqual(float_dh.gridType(0), nanovdb.GridType.Float) + return float_dh + + def test_set_grid_class_refreshes_checksum(self): + cp = _require_cupy(self) + float_dh = self._float_grid_from_block(cp) + float_dg = float_dh.deviceGrid(0) + # Populate a Full checksum so the retag has something to invalidate: if + # setGridClass failed to refresh it, the post-change Full validation + # below (computed over the new class field) would no longer match. + nanovdb.tools.cuda.updateChecksum(float_dg, nanovdb.CheckMode.Full) + self.assertTrue(nanovdb.tools.cuda.validateChecksum(float_dg, nanovdb.CheckMode.Full)) + + nanovdb.tools.cuda.setGridClass(float_dg, nanovdb.GridClass.LevelSet) + # Class changed AND the Full checksum still validates -> it was refreshed. + self.assertTrue(nanovdb.tools.cuda.validateChecksum(float_dg, nanovdb.CheckMode.Full)) + + float_dh.deviceDownload(0, True) + g = float_dh.grid(0) + self.assertEqual(g.gridClass(), nanovdb.GridClass.LevelSet) + self.assertTrue(g.isLevelSet()) + self.assertTrue(nanovdb.tools.validateChecksum(g, nanovdb.CheckMode.Full)) + + def test_set_grid_class_survives_file_roundtrip(self): + cp = _require_cupy(self) + float_dh = self._float_grid_from_block(cp) + nanovdb.tools.cuda.setGridClass(float_dh.deviceGrid(0), nanovdb.GridClass.FogVolume) + float_dh.deviceDownload(0, True) + tmp = tempfile.NamedTemporaryFile(suffix=".nvdb", delete=False) + tmp.close() + try: + float_dh.write(tmp.name) + g = nanovdb.io.readGrid(tmp.name).grid(0) + self.assertEqual(g.gridClass(), nanovdb.GridClass.FogVolume) + self.assertTrue(g.isFogVolume()) + finally: + os.unlink(tmp.name) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestDeviceInfra(unittest.TestCase): + """DeviceMesh / DeviceStreamMap / DeviceResource / TempDevicePool.""" + + def test_device_mesh(self): + mesh = nanovdb.cuda.DeviceMesh() + self.assertGreaterEqual(mesh.deviceCount(), 1) + # canAccessPeer is well-defined (self-peer is typically False). + self.assertIsInstance(mesh.canAccessPeer(0, 0), bool) + + def test_device_stream_map(self): + DSM = nanovdb.cuda.DeviceStreamMap + sm = DSM(DSM.DeviceType.Unified, [], 0) + self.assertGreaterEqual(sm.deviceCount(), 1) + # stream(deviceId) is a raw CUDA stream handle as a Python int. + self.assertIsInstance(sm.stream(0), int) + + def test_device_resource_alloc_dealloc(self): + DR = nanovdb.cuda.DeviceResource + self.assertEqual(DR.DEFAULT_ALIGNMENT, 256) + ptr = DR.allocateAsync(1024, 256, 0) + self.assertIsInstance(ptr, int) + self.assertNotEqual(ptr, 0) + DR.deallocateAsync(ptr, 1024, 256, 0) + + def test_temp_device_pool(self): + tp = nanovdb.cuda.TempDevicePool() + tp.setRequestedSize(2048) + tp.reallocate(0) + self.assertEqual(tp.requestedSize(), 2048) + self.assertGreaterEqual(tp.size(), 2048) + self.assertNotEqual(tp.data(), 0) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestDistributedPointsToGrid(unittest.TestCase): + """Multi-GPU DistributedPointsToGrid over MANAGED coordinate arrays. + + With a single device this exercises the trivial single-GPU path, but + still validates the cuda_managed memory constraint and the + UnifiedGridHandle result type. + """ + + def test_managed_coords_produce_unified_handle(self): + cp = _require_cupy(self) + import numpy as np + prev = cp.cuda.get_allocator() + cp.cuda.set_allocator(cp.cuda.malloc_managed) + try: + pts = np.array([[0, 0, 0], [1, 0, 0], [0, 1, 0], [5, 5, 5]], + dtype=np.int32) + voxels = cp.asarray(pts) # MANAGED (unified) memory + mesh = nanovdb.cuda.DeviceMesh() + # The mesh must outlive the converter (held by reference). + conv = nanovdb.tools.cuda.DistributedPointsToGrid( + mesh, 1.0, (0.0, 0.0, 0.0)) + uh = conv.getHandle(voxels) + self.assertEqual(uh.gridType(0), nanovdb.GridType.OnIndex) + # UnifiedGridHandle intentionally does NOT expose the + # CAI / DLPack bridges that DeviceGridHandle does. + self.assertFalse(hasattr(uh, "__cuda_array_interface__")) + finally: + cp.cuda.set_allocator(prev) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestTorchInterop(unittest.TestCase): + """Optional PyTorch interop (skips when torch is not installed).""" + + def test_torch_cuda_array_interface(self): + try: + import torch + except ImportError: + self.skipTest("PyTorch not installed") + if not torch.cuda.is_available(): + self.skipTest("PyTorch built without CUDA / no GPU") + dh, dg = _build_device_onindex_grid(20.0) + # torch can adopt the DeviceGridHandle device buffer via DLPack. + t = torch.from_dlpack(dh) + self.assertEqual(t.numel(), dh.size()) + self.assertTrue(t.is_cuda) + + +@unittest.skipIf( + not nanovdb.isCudaAvailable(), "nanovdb module was compiled without CUDA support" +) +@unittest.skipIf( + not nanovdb.isGpuAvailable(), "No CUDA-capable GPU available" +) +class TestMergeGridsList(unittest.TestCase): + """nanovdb.tools.cuda.mergeGrids -- topological (active-mask) union of device + OnIndex grids. The binary form merges exactly two; the list overload folds an + arbitrary number in a single native N-ary pass and must agree with chaining + the binary form.""" + + def _grid(self, cp, pts): + import numpy as np + coords = cp.asarray(np.ascontiguousarray(pts, dtype=np.int32)) + return nanovdb.tools.cuda.voxelsToOnIndexGrid(coords, 1.0) + + def _stats(self, h): + """(active voxels, (leaf, lower, upper) node counts, grid bytes).""" + h.deviceDownload(0, True) + t = h.grid(0).tree() + return (int(t.activeVoxelCount()), + (t.nodeCount(0), t.nodeCount(1), t.nodeCount(2)), + int(h.gridSize(0))) + + def test_disjoint_union_is_sum(self): + cp = _require_cupy(self) + sets = [[[0, 0, 0], [1, 0, 0]], [[5, 0, 0]], + [[0, 9, 0], [0, 10, 0]], [[4096, 0, 0]]] # last spans a new upper + hs = [self._grid(cp, s) for s in sets] + res = nanovdb.tools.cuda.mergeGrids([h.deviceGrid(0) for h in hs]) + self.assertEqual(self._stats(res)[0], sum(len(s) for s in sets)) + + def test_overlap_dedups_and_matches_pairwise(self): + cp = _require_cupy(self) + # B shares (2,0,0) with A; C lives in a separate upper region. + A = [[0, 0, 0], [1, 0, 0], [2, 0, 0]] + B = [[2, 0, 0], [3, 0, 0]] + C = [[0, 9, 0], [4096, 0, 0]] + hs = [self._grid(cp, s) for s in (A, B, C)] + truth = len({tuple(p) for p in A + B + C}) + self.assertLess(truth, len(A) + len(B) + len(C)) # overlap really exists + + nary = nanovdb.tools.cuda.mergeGrids([h.deviceGrid(0) for h in hs]) + ab = nanovdb.tools.cuda.mergeGrids(hs[0].deviceGrid(0), hs[1].deviceGrid(0)) + abc = nanovdb.tools.cuda.mergeGrids(ab.deviceGrid(0), hs[2].deviceGrid(0)) + + s_nary, s_chain = self._stats(nary), self._stats(abc) + self.assertEqual(s_nary[0], truth) # union dedups the shared voxel + self.assertEqual(s_nary, s_chain) # node counts + bytes identical to chaining + + def test_single_element_is_copy(self): + cp = _require_cupy(self) + h = self._grid(cp, [[0, 0, 0], [0, 1, 0], [0, 2, 0]]) + res = nanovdb.tools.cuda.mergeGrids([h.deviceGrid(0)]) + self.assertEqual(self._stats(res)[0], 3) + + def test_empty_list_raises(self): + _require_cupy(self) + with self.assertRaises(Exception): + nanovdb.tools.cuda.mergeGrids([]) + + +if __name__ == "__main__": + unittest.main() diff --git a/nanovdb/nanovdb/python/test/TestNanoVDB.py b/nanovdb/nanovdb/python/test/TestNanoVDB.py index f8dbaf92db..e05ce69e59 100644 --- a/nanovdb/nanovdb/python/test/TestNanoVDB.py +++ b/nanovdb/nanovdb/python/test/TestNanoVDB.py @@ -431,6 +431,94 @@ def test_constructed_from_grid(self): self.assertTrue(nanovdb.GridMetaData.safeCast(h.grid())) +class TestGridHeaderSetters(unittest.TestCase): + """Header mutators on the Grid base (setGridClass / setTransform / setGridName + / the flag toggles) edit the grid in place and refresh its checksum, so a + validating reader still accepts the grid after the edit.""" + + def _sphere(self): + # voxelSize 1.0 so setTransform's effect on voxelSize() is unambiguous. + h = nanovdb.tools.createLevelSetSphere(radius=10.0, voxelSize=1.0) + return h, h.grid() + + def test_set_grid_class_refreshes_checksum(self): + h, g = self._sphere() + self.assertTrue(g.isLevelSet()) + # Populate a Full checksum so a missing refresh would leave it stale. + nanovdb.tools.updateChecksum(g, nanovdb.CheckMode.Full) + self.assertTrue(nanovdb.tools.validateChecksum(g, nanovdb.CheckMode.Full)) + + g.setGridClass(nanovdb.GridClass.FogVolume) + self.assertEqual(g.gridClass(), nanovdb.GridClass.FogVolume) + self.assertTrue(g.isFogVolume()) + self.assertFalse(g.isLevelSet()) + # Still validates -> the class change refreshed the checksum. + self.assertTrue(nanovdb.tools.validateChecksum(g, nanovdb.CheckMode.Full)) + self.assertTrue(g.isValid()) + + def test_set_transform(self): + h, g = self._sphere() + index_bbox_before = str(g.indexBBox()) + nanovdb.tools.updateChecksum(g, nanovdb.CheckMode.Full) + + # Uniform scale 2 + world translation (5,0,0). + g.setTransform(2.0, nanovdb.math.Vec3d(5.0, 0.0, 0.0)) + self.assertEqual(g.voxelSize()[0], 2.0) + # The transform relabels world coordinates; it does not move voxels. + self.assertEqual(str(g.indexBBox()), index_bbox_before) + # world = scale * index + translation + self.assertEqual(g.applyMap(nanovdb.math.Vec3d(0, 0, 0)), nanovdb.math.Vec3d(5, 0, 0)) + self.assertEqual(g.applyMap(nanovdb.math.Vec3d(1, 0, 0)), nanovdb.math.Vec3d(7, 0, 0)) + # World bbox was recomputed under the new map (no longer the unit-scale box). + self.assertEqual(g.voxelSize(), nanovdb.math.Vec3d(2.0)) + self.assertTrue(nanovdb.tools.validateChecksum(g, nanovdb.CheckMode.Full)) + + # Default translation is the origin. + g.setTransform(0.5) + self.assertEqual(g.voxelSize()[0], 0.5) + self.assertEqual(g.applyMap(nanovdb.math.Vec3d(0, 0, 0)), nanovdb.math.Vec3d(0, 0, 0)) + + def test_set_transform_rejects_invalid_voxel_size(self): + # Map::set only debug-asserts positivity, so the binding must reject a + # singular / non-finite transform before it reaches the grid header. + h, g = self._sphere() + for bad in (0.0, -1.0, float("nan"), float("inf")): + with self.assertRaises(ValueError): + g.setTransform(bad) + # The transform is unchanged after the rejected calls. + self.assertEqual(g.voxelSize()[0], 1.0) + + def test_flag_and_name_setters_refresh_checksum(self): + h, g = self._sphere() + nanovdb.tools.updateChecksum(g, nanovdb.CheckMode.Full) + for toggle in (lambda: g.setBBoxOn(True), lambda: g.setMinMaxOn(False), + lambda: g.setAverageOn(True), lambda: g.setStdDeviationOn(False), + lambda: g.setLongGridNameOn(False)): + toggle() + self.assertTrue(nanovdb.tools.validateChecksum(g, nanovdb.CheckMode.Full)) + self.assertTrue(g.setGridName("renamed_grid")) + self.assertEqual(g.shortGridName(), "renamed_grid") + self.assertTrue(nanovdb.tools.validateChecksum(g, nanovdb.CheckMode.Full)) + + def test_setters_survive_file_roundtrip(self): + h, g = self._sphere() + g.setGridClass(nanovdb.GridClass.FogVolume) + g.setTransform(4.0, nanovdb.math.Vec3d(1.0, 2.0, 3.0)) + g.setGridName("retagged") + tmp = tempfile.NamedTemporaryFile(suffix=".nvdb", delete=False) + tmp.close() + try: + nanovdb.io.writeGrid(tmp.name, h) + g2 = nanovdb.io.readGrid(tmp.name).grid() + self.assertEqual(g2.gridClass(), nanovdb.GridClass.FogVolume) + self.assertEqual(g2.voxelSize()[0], 4.0) + self.assertEqual(g2.shortGridName(), "retagged") + self.assertEqual(g2.applyMap(nanovdb.math.Vec3d(0, 0, 0)), + nanovdb.math.Vec3d(1, 2, 3)) + finally: + os.unlink(tmp.name) + + class TestBlindDataEmpty(unittest.TestCase): """Blind data API (blindDataCount, blindMetaData, findBlindData, findBlindDataForSemantic, getBlindData) returns sensible None/-1 @@ -563,7 +651,7 @@ def test_readonly_accessors_have_neither_setvoxel_nor_nodeinfo(self): class TestTreeNodeWalking(unittest.TestCase): """Walk a grid's tree from Python: Grid.tree(), Root/Upper/Lower/Leaf node access and metadata, per-leaf zero-copy values() and bulk - grid.leaf_values() NumPy views, NodeManager + createNodeManager.""" + grid.leafValues() NumPy views, NodeManager + createNodeManager.""" @classmethod def setUpClass(cls): @@ -630,7 +718,7 @@ def test_bulk_leaf_values(self): import numpy as np except ImportError: self.skipTest("numpy not installed") - bulk = self.g.leaf_values() + bulk = self.g.leafValues() self.assertEqual(bulk.shape, (self.tree.nodeCount(0), 512)) self.assertEqual(bulk.dtype, np.float32) # First row should match per-leaf values(). @@ -649,7 +737,7 @@ def test_bulk_leaf_values_empty_grid_returns_empty_array(self): empty_h = nanovdb.tools.createFloatGrid( 0.0, "empty", nanovdb.GridClass.Unknown, lambda ijk: 0.0, empty_bbox) - bulk = empty_h.grid().leaf_values() + bulk = empty_h.grid().leafValues() self.assertEqual(bulk.shape, (0, 512)) self.assertEqual(bulk.dtype, np.float32) @@ -771,7 +859,7 @@ def test_grid_leaf_values_temporary(self): import numpy # noqa: F401 except ImportError: self.skipTest("numpy not installed") - bulk = nanovdb.tools.createFogVolumeSphere().grid().leaf_values() + bulk = nanovdb.tools.createFogVolumeSphere().grid().leafValues() self._force_gc() self.assertEqual(bulk.shape[1], 512) _ = float(bulk[0, 0]) @@ -821,7 +909,7 @@ def _make_cube_on_index_grid(self): 0.0, "cube", nanovdb.GridClass.Unknown, lambda ijk: 1.0, bbox) return nanovdb.tools.createOnIndexGrid( - float_h.grid(), include_stats=False, include_tiles=False) + float_h.grid(), includeStats=False, includeTiles=False) def test_create_on_index_grid(self): h = self._make_cube_on_index_grid() @@ -841,7 +929,7 @@ def test_create_on_index_grid_rejects_unsupported_source(self): def test_build_voxel_block_manager_handle(self): h = self._make_cube_on_index_grid() g = h.grid() - vbm = nanovdb.tools.buildVoxelBlockManager(g, log2_block_width=6) + vbm = nanovdb.tools.buildVoxelBlockManager(g, log2BlockWidth=6) self.assertGreater(vbm.blockCount(), 0) self.assertEqual(vbm.firstOffset(), 1) self.assertEqual(vbm.lastOffset(), g.activeVoxelCount()) @@ -853,21 +941,21 @@ def test_buffers_zero_copy_shape_and_dtype(self): except ImportError: self.skipTest("numpy not installed") h = self._make_cube_on_index_grid() - # log2_block_width=6 -> JumpMapLength=1. - vbm6 = nanovdb.tools.buildVoxelBlockManager(h.grid(), log2_block_width=6) - self.assertEqual(vbm6.log2_block_width, 6) - self.assertEqual(vbm6.block_width, 64) - self.assertEqual(vbm6.jump_map_length, 1) + # log2BlockWidth=6 -> JumpMapLength=1. + vbm6 = nanovdb.tools.buildVoxelBlockManager(h.grid(), log2BlockWidth=6) + self.assertEqual(vbm6.log2BlockWidth, 6) + self.assertEqual(vbm6.blockWidth, 64) + self.assertEqual(vbm6.jumpMapLength, 1) fl = vbm6.firstLeafID() self.assertEqual(fl.shape, (vbm6.blockCount(),)) self.assertEqual(fl.dtype, np.uint32) jm = vbm6.jumpMap() self.assertEqual(jm.shape, (vbm6.blockCount(), 1)) self.assertEqual(jm.dtype, np.uint64) - # log2_block_width=7 -> JumpMapLength=2 (independent build, separate + # log2BlockWidth=7 -> JumpMapLength=2 (independent build, separate # allocation; the jumpMap shape comes from the handle, not the caller). - vbm7 = nanovdb.tools.buildVoxelBlockManager(h.grid(), log2_block_width=7) - self.assertEqual(vbm7.jump_map_length, 2) + vbm7 = nanovdb.tools.buildVoxelBlockManager(h.grid(), log2BlockWidth=7) + self.assertEqual(vbm7.jumpMapLength, 2) self.assertEqual(vbm7.jumpMap().shape, (vbm7.blockCount(), 2)) def test_decode_block_zero(self): @@ -877,7 +965,7 @@ def test_decode_block_zero(self): self.skipTest("numpy not installed") h = self._make_cube_on_index_grid() g = h.grid() - vbm = nanovdb.tools.buildVoxelBlockManager(g, log2_block_width=6) + vbm = nanovdb.tools.buildVoxelBlockManager(g, log2BlockWidth=6) leaf_index, voxel_offset = vbm.decodeBlock(g, 0) self.assertEqual(leaf_index.shape, (64,)) self.assertEqual(leaf_index.dtype, np.uint32) @@ -887,7 +975,7 @@ def test_decode_block_zero(self): fl = np.asarray(vbm.firstLeafID()) jm = np.asarray(vbm.jumpMap()) li_free, vo_free = nanovdb.tools.decodeInverseMaps( - g, int(fl[0]), jm[0], vbm.firstOffset(), log2_block_width=6) + g, int(fl[0]), jm[0], vbm.firstOffset(), log2BlockWidth=6) self.assertTrue(np.array_equal(leaf_index, li_free)) self.assertTrue(np.array_equal(voxel_offset, vo_free)) @@ -902,24 +990,24 @@ def test_log2_block_width_out_of_range(self): h = self._make_cube_on_index_grid() g = h.grid() with self.assertRaises(ValueError): - nanovdb.tools.buildVoxelBlockManager(g, log2_block_width=5) + nanovdb.tools.buildVoxelBlockManager(g, log2BlockWidth=5) with self.assertRaises(ValueError): - nanovdb.tools.buildVoxelBlockManager(g, log2_block_width=10) + nanovdb.tools.buildVoxelBlockManager(g, log2BlockWidth=10) def test_build_voxel_block_manager_rejects_misaligned_first_offset(self): - # first_offset must satisfy first_offset == 1 (mod BlockWidth). - # For log2_block_width=6, BlockWidth=64, so 1, 65, 129, ... are valid + # firstOffset must satisfy firstOffset == 1 (mod BlockWidth). + # For log2BlockWidth=6, BlockWidth=64, so 1, 65, 129, ... are valid # but 2 is not. h = self._make_cube_on_index_grid() g = h.grid() with self.assertRaises(ValueError): nanovdb.tools.buildVoxelBlockManager( - g, log2_block_width=6, first_offset=2) - # And the wider-block case: log2_block_width=7 -> BlockWidth=128, - # first_offset=65 is valid for width=6 but misaligned for width=7. + g, log2BlockWidth=6, firstOffset=2) + # And the wider-block case: log2BlockWidth=7 -> BlockWidth=128, + # firstOffset=65 is valid for width=6 but misaligned for width=7. with self.assertRaises(ValueError): nanovdb.tools.buildVoxelBlockManager( - g, log2_block_width=7, first_offset=65) + g, log2BlockWidth=7, firstOffset=65) def test_decode_inverse_maps_rejects_bad_first_leaf_id(self): try: @@ -928,25 +1016,25 @@ def test_decode_inverse_maps_rejects_bad_first_leaf_id(self): self.skipTest("numpy not installed") h = self._make_cube_on_index_grid() g = h.grid() - vbm = nanovdb.tools.buildVoxelBlockManager(g, log2_block_width=6) + vbm = nanovdb.tools.buildVoxelBlockManager(g, log2BlockWidth=6) jm0 = np.asarray(vbm.jumpMap())[0] n_leaves = g.tree().nodeCount(0) with self.assertRaises(IndexError): nanovdb.tools.decodeInverseMaps( - g, n_leaves, jm0, vbm.firstOffset(), log2_block_width=6) + g, n_leaves, jm0, vbm.firstOffset(), log2BlockWidth=6) def test_build_voxel_block_manager_rejects_undersized_n_blocks(self): - # Caller-supplied n_blocks must hold at least - # ceil((last_offset - first_offset + 1) / BlockWidth) blocks; + # Caller-supplied nBlocks must hold at least + # ceil((lastOffset - firstOffset + 1) / BlockWidth) blocks; # smaller values would silently truncate coverage. h = self._make_cube_on_index_grid() g = h.grid() - # The cube grid has ~9261 active voxels, so at log2_block_width=6 + # The cube grid has ~9261 active voxels, so at log2BlockWidth=6 # (BlockWidth=64) the minimum is roughly 145 blocks. Passing 1 # must be rejected. with self.assertRaises(ValueError): nanovdb.tools.buildVoxelBlockManager( - g, log2_block_width=6, n_blocks=1) + g, log2BlockWidth=6, nBlocks=1) def test_build_voxel_block_manager_rejects_non_on_index_grid(self): # FloatGrid is not an OnIndexGrid. @@ -962,7 +1050,7 @@ def test_untouched_blocks_trip_sentinel_guard(self): # decodeBlock (raising ValueError). The sweep must therefore see # only two outcomes per block: a successful decode or a sentinel # ValueError — no segfaults, no IndexError (those would indicate - # block_index out of range, not sentinel), and no silent wrong + # blockIndex out of range, not sentinel), and no silent wrong # decode. try: import numpy as np # noqa: F401 @@ -970,7 +1058,7 @@ def test_untouched_blocks_trip_sentinel_guard(self): self.skipTest("numpy not installed") h = self._make_cube_on_index_grid() g = h.grid() - vbm = nanovdb.tools.buildVoxelBlockManager(g, log2_block_width=6) + vbm = nanovdb.tools.buildVoxelBlockManager(g, log2BlockWidth=6) n_leaves = g.tree().nodeCount(0) fl = vbm.firstLeafID() for b in range(vbm.blockCount()): @@ -1004,7 +1092,7 @@ def test_default_constructed_handle_returns_empty_arrays(self): self.assertEqual(fl.shape, (0,)) self.assertEqual(fl.dtype, np.uint32) jm = np.asarray(vbm.jumpMap()) - # Default-constructed handle uses log2_block_width=6 -> JumpMapLength=1. + # Default-constructed handle uses log2BlockWidth=6 -> JumpMapLength=1. self.assertEqual(jm.shape, (0, 1)) self.assertEqual(jm.dtype, np.uint64) @@ -1015,7 +1103,7 @@ def test_reset_handle_returns_empty_arrays(self): except ImportError: self.skipTest("numpy not installed") h = self._make_cube_on_index_grid() - vbm = nanovdb.tools.buildVoxelBlockManager(h.grid(), log2_block_width=6) + vbm = nanovdb.tools.buildVoxelBlockManager(h.grid(), log2BlockWidth=6) self.assertGreater(vbm.blockCount(), 0) vbm.reset() self.assertEqual(vbm.blockCount(), 0) diff --git a/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh b/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh index 2bf1b44e0f..27cf1b725c 100644 --- a/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh @@ -162,7 +162,15 @@ __global__ void processGridTreeRootKernel(IndexToGridNodeAccessor *no // process Grid *dstGrid.data() = *srcGrid.data(); dstGrid.mGridType = toGridType(); + // NOTE: The source is an index grid (GridClass::IndexGrid); a plain value + // grid must not inherit that class or it forms an invalid GridType/GridClass + // combination (e.g. Float + IndexGrid) that fails grid validation. + dstGrid.mGridClass = GridClass::Unknown; dstGrid.mData1 = 0u; + + dstGrid.mGridSize = nodeAcc->size; + dstGrid.mBlindMetadataOffset = nodeAcc->size; + dstGrid.mBlindMetadataCount = 0u; dstGrid.mGridIndex = 0u; // Possibly overwriting input; returned grid has batch size 1 dstGrid.mGridCount = 1u; // we will recompute GridData::mChecksum later diff --git a/pyproject.toml b/pyproject.toml index b42bcd5c80..6823cd8f66 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -1,5 +1,5 @@ [build-system] -requires = ["scikit_build_core", "nanobind"] +requires = ["scikit_build_core", "nanobind>=2.12.0,<3"] build-backend = "scikit_build_core.build" [project]