Attempting to fix secretes manager access error

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Improving buildspec logging
2025-06-25 10:35:02 +12:00 · 2025-06-25 10:23:53 +12:00 · 2025-06-25 10:18:03 +12:00 · 2025-06-25 09:52:32 +12:00 · 2025-06-25 09:31:01 +12:00 · 2025-06-25 09:28:46 +12:00
684 changed files with 144968 additions and 1444 deletions
@@ -0,0 +1 @@
+
@@ -55,7 +55,7 @@ jobs:
      github.event_name != 'workflow_dispatch' ||
      github.event.inputs.type == 'all' ||
      github.event.inputs.type == 'linux'
-    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
+    runs-on: ${{ (github.repository_owner == 'NVIDIA-Omniverse' && 'linux-amd64-cpu32') || (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
    name: >
      linux-ax:${{ matrix.config.image }}-cxx:${{ matrix.config.cxx }}-${{ matrix.config.build }}
    container:
@@ -152,7 +152,7 @@ jobs:
    if: |
      github.event_name == 'workflow_dispatch' &&
      github.event.inputs.type == 'grammar'
-    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
+    runs-on: ${{ (github.repository_owner == 'NVIDIA-Omniverse' && 'linux-amd64-cpu32') || (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
    container:
      image: aswf/ci-openvdb:2023-clang15
    steps:
@@ -160,7 +160,7 @@ jobs:
    - name: build
      run: ./ci/build.sh -v --components=axgr --target=openvdb_ax_grammar --cargs=\"-DOPENVDB_AX_GRAMMAR_NO_LINES=ON\"
    - name: upload grammar
-      uses: actions/upload-artifact@v4
+      uses: actions/upload-artifact@v2
      with:
        name: ax_grammar
        path: ./build/openvdb_ax/openvdb_ax/openvdb_ax/grammar
@@ -62,7 +62,7 @@ jobs:
      github.event_name != 'workflow_dispatch' ||
      github.event.inputs.type == 'all' ||
      github.event.inputs.type == 'linux'
-    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
+    runs-on: ${{ (github.repository_owner == 'NVIDIA-Omniverse' && 'linux-amd64-cpu32') || (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
    name: >
      linux-vfx:${{ matrix.config.image }}-
      abi:${{ matrix.config.abi }}-
@@ -0,0 +1,118 @@
+# Copyright Contributors to the OpenVDB Project
+# SPDX-License-Identifier: Apache-2.0
+#
+# GitHub Actions workflow file
+# https://help.github.com/en/actions/reference/workflow-syntax-for-github-actions
+
+name: fVDB Codebuild
+
+on:
+  pull_request:
+    paths-ignore:
+        - 'CHANGES'
+        - 'CODEOWNERS'
+        - 'doc/**'
+        - 'openvdb/**'
+        - 'openvdb_cmd/**'
+        - 'openvdb_ax/**'
+        - 'openvdb_maya/**'
+        - 'openvdb_houdini/**'
+        - 'nanovdb/**'
+        - 'pendingchanges/**'
+        - '**.md'
+        - 'fvdb/debug/**'
+        - 'fvdb/docs/**'
+        - 'fvdb/examples/**'
+        - 'fvdb/notebooks/**'
+        - 'fvdb/scripts/**'
+
+permissions:
+  contents: read
+  id-token: write
+
+jobs:
+  # ---------------------------------------------------------------------------
+  # Linux GPU (AWS CodeBuild)
+  # ---------------------------------------------------------------------------
+
+  linux_gpu:
+    name: 'ci-openvdb 2024'
+    # Don't run on forks
+    if: github.repository == 'AcademySoftwareFoundation/openvdb'
+    # GH-hosted VM. The build runs in CentOS 7 'container' hard-coded in
+    # AWS CodeBuild project.
+    # TODO: Add support for dynamic GH Action defined CodeBuild
+    #       container choice.
+    runs-on: ubuntu-latest
+    steps:
+      - name: Verify AWS Credentials Set
+        run: |
+          if [ -z "${{ secrets.CODEBUILD_ID }}" ] || [ -z "${{ secrets.CODEBUILD_SECRET }}" ]; then
+            echo "Error: AWS credentials secrets not properly configured"
+            echo "Please ensure CODEBUILD_ID and CODEBUILD_SECRET are set in organization secrets"
+            exit 1
+          fi
+          echo "AWS credentials secrets are configured"
+      - name: Configure AWS Credentials
+        uses: aws-actions/configure-aws-credentials@v4
+        with:
+          aws-access-key-id: ${{ secrets.CODEBUILD_ID }}
+          aws-secret-access-key: ${{ secrets.CODEBUILD_SECRET }}
+          aws-region: us-west-2
+      - name: Verify AWS Authentication
+        run: |
+          echo "Testing AWS authentication..."
+          aws sts get-caller-identity || {
+            echo "Error: AWS authentication failed"
+            exit 1
+          }
+          echo "AWS authentication successful"
+      - name: CodeBuild Monitoring URLs
+        run: |
+          echo "🚀 Starting CodeBuild job..."
+          echo "Project: OpenVDB"
+          echo ""
+          echo "📊 Monitor the build progress at:"
+
+          # Check if public builds are enabled by looking for the public project URL
+          PUBLIC_URL=$(aws codebuild batch-get-projects --names OpenVDB --query 'projects[0].publicProjectUrl' --output text 2>/dev/null || echo "None")
+
+          if [ "$PUBLIC_URL" != "None" ] && [ "$PUBLIC_URL" != "null" ]; then
+            echo "📖 Public Build Results (no sign-in required):"
+            echo "   $PUBLIC_URL"
+            echo ""
+          fi
+
+          echo "🔒 AWS Console (requires sign-in):"
+          echo "   Project: https://console.aws.amazon.com/codesuite/codebuild/projects/OpenVDB/history"
+          echo ""
+          echo "ℹ️  Individual build URLs will be available once the job starts..."
+      - name: Run CodeBuild
+        id: codebuild
+        uses: aws-actions/aws-codebuild-run-build@v1
+        with:
+          project-name: OpenVDB
+          buildspec-override: ci/buildspec.yml
+      - name: CodeBuild Results
+        if: always()
+        run: |
+          echo "📋 CodeBuild Results:"
+          echo "Build ID: ${{ steps.codebuild.outputs.aws-build-id }}"
+          echo "Build Status: ${{ steps.codebuild.outputs.aws-build-status }}"
+          echo ""
+          echo "🔗 Detailed Build Links:"
+          if [ -n "$CODEBUILD_PUBLIC_BUILD_URL" ]; then
+            echo "📖 Public Build Details: $CODEBUILD_PUBLIC_BUILD_URL"
+          fi
+          echo "🔒 AWS Console Build Details: https://console.aws.amazon.com/codesuite/codebuild/projects/OpenVDB/build/${{ steps.codebuild.outputs.aws-build-id }}/"
+          echo ""
+          if [ "${{ steps.codebuild.outputs.aws-build-status }}" != "SUCCEEDED" ]; then
+            echo "❌ Build failed! Check the build logs at the links above for detailed error information."
+            echo "Common issues to check:"
+            echo "- Build environment configuration"
+            echo "- Source code compilation errors"
+            echo "- Missing dependencies"
+            echo "- Buildspec.yml configuration"
+          else
+            echo "✅ Build succeeded!"
+          fi
@@ -0,0 +1,80 @@
+name: fVDB Code Style
+on:
+  push:
+    branches:
+      - "pull-request/[0-9]+"
+    paths-ignore:
+        - 'CHANGES'
+        - 'CODEOWNERS'
+        - 'doc/**'
+        - 'ci/**'
+        - 'openvdb/**'
+        - 'openvdb_cmd/**'
+        - 'openvdb_ax/**'
+        - 'openvdb_maya/**'
+        - 'openvdb_houdini/**'
+        - 'openvdb_wolfram/**'
+        - 'tsc/**'
+        - 'nanovdb/**'
+        - 'pendingchanges/**'
+        - '**.md'
+        - 'fvdb/docs/**'
+        - 'fvdb/env/**'
+        - 'fvdb/notebooks/**'
+
+
+
+# Allow subsequent pushes to the same PR or REF to cancel any previous jobs.
+concurrency:
+  group: ${{ github.workflow }}-${{ github.event.pull_request.number || github.ref }}
+  cancel-in-progress: true
+
+jobs:
+  test-python-black-lint:
+    name: Check Python code style with black
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/checkout@v4
+      - uses: psf/black@stable
+        with:
+          options: "--check --verbose --target-version=py311 --line-length=120"
+          src: "fvdb/"
+          version: "~= 24.0"
+
+  test-cpp-clang-format-lint:
+    name: Check C++ code style with clang-format
+    runs-on: ubuntu-latest
+    steps:
+    - uses: actions/checkout@v4
+    - uses: DoozyX/clang-format-lint-action@v0.18.2
+      with:
+        source: 'fvdb/src/'
+        extensions: 'h,cpp,cc,cu,cuh'
+        clangFormatVersion: 18
+        style: file
+
+  include-guards:
+    name: Check include guards
+    runs-on: ubuntu-latest
+    steps:
+    - uses: actions/checkout@v4
+    - uses: swahtz/include-guards-check-action@master
+      with:
+        path: 'fvdb/src/'
+
+  check-spdx-identifiers:
+    name: Check SPDX identifiers
+    runs-on: ubuntu-latest
+    steps:
+    - name: checkout
+      uses: actions/checkout@v4
+    - uses: swahtz/spdx@feature/ignore_paths
+      with:
+        licenses: |-
+          Apache-2.0
+        ignore-paths: |-
+            openvdb/openvdb/math/Half.cc
+            openvdb/openvdb/math/Half.h
+            openvdb_wolfram/OpenVDBLink
+            openvdb_ax/openvdb_ax/grammar/generated
+            fvdb/projects/panoptic_segmentation/garfield/garfield
@@ -65,7 +65,7 @@ jobs:
    if: >
      ${{ needs.checksecret.outputs.HOUDINI_SECRETS == 'true' ||
          github.repository_owner == 'AcademySoftwareFoundation' }}
-    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
+    runs-on: ${{ (github.repository_owner == 'NVIDIA-Omniverse' && 'linux-amd64-cpu32') || (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
    name: hou:${{ matrix.config.hou_hash }}-vfx:${{ matrix.config.image }}-cxx:${{ matrix.config.cxx }}
    container:
      image: aswf/ci-base:${{ matrix.config.image }}
@@ -51,7 +51,7 @@ jobs:
      github.event_name != 'workflow_dispatch' ||
      github.event.inputs.type == 'all' ||
      github.event.inputs.type == 'linux'
-    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
+    runs-on: ${{ (github.repository_owner == 'NVIDIA-Omniverse' && 'linux-amd64-cpu32') || (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
    name: >
      linux-nanovdb:cxx:${{ matrix.config.cxx }}-${{ matrix.config.build }}
    container:
@@ -171,7 +171,7 @@ jobs:
      github.event_name != 'workflow_dispatch' ||
      github.event.inputs.type == 'all' ||
      github.event.inputs.type == 'linux'
-    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
+    runs-on: ${{ (github.repository_owner == 'NVIDIA-Omniverse' && 'linux-amd64-cpu32') || (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
    container:
      image: aswf/ci-openvdb:2024
    steps:
@@ -61,7 +61,7 @@ jobs:
      (github.event_name != 'workflow_dispatch' ||
       github.event.inputs.type == 'all' ||
       github.event.inputs.type == 'houdini')
-    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
+    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-20.04-8c-32g-300h') || 'ubuntu-latest' }}
    name: linux-houdini:${{ matrix.config.hou_hash }}
    env:
      CXX: clang++
@@ -146,7 +146,7 @@ jobs:
      github.event_name != 'workflow_dispatch' ||
      github.event.inputs.type == 'all' ||
      github.event.inputs.type == 'extra'
-    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
+    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-20.04-8c-32g-300h') || 'ubuntu-latest' }}
    name: linux-extra:${{ matrix.config.name }}
    container:
      image: aswf/ci-openvdb:2024
@@ -203,7 +203,7 @@ jobs:
      - name: install_deps
        run: |
          if [ "$RUNNER_OS" == "Linux" ]; then
-            sudo apt-get -q install -y libboost-dev libboost-iostreams-dev libtbb-dev libblosc-dev llvm-15-dev libgtest-dev libgmock-dev libcppunit-dev
+            sudo apt-get -q install -y libboost-dev libboost-iostreams-dev libtbb-dev libblosc-dev llvm-dev libgtest-dev libcppunit-dev
            ./ci/install_nanobind.sh 2.0.0
          elif [ "$RUNNER_OS" == "macOS" ]; then
            ./ci/install_macos.sh 15
@@ -284,7 +284,7 @@ jobs:
      github.event_name != 'workflow_dispatch' ||
      github.event.inputs.type == 'all' ||
      github.event.inputs.type == 'ax'
-    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
+    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-20.04-8c-32g-300h') || 'ubuntu-latest' }}
    name: >
      linux-ax:${{ matrix.config.image }}-cxx:${{ matrix.config.cxx }}-${{ matrix.config.build }}
    container:
@@ -313,7 +313,6 @@ jobs:
          -DOPENVDB_AX_TEST_CMD_DOWNLOADS=ON
          -DUSE_EXPLICIT_INSTANTIATION=OFF
          -DOPENVDB_CXX_STRICT=ON
-          -DOPENVDB_BUILD_VDB_TOOL=OFF
          \"
      - name: clean
        if: matrix.config.components == 'core'
@@ -329,7 +328,6 @@ jobs:
          -DOPENVDB_AX_TEST_CMD_DOWNLOADS=ON
          -DUSE_EXPLICIT_INSTANTIATION=OFF
          -DOPENVDB_CXX_STRICT=ON
-          -DOPENVDB_BUILD_VDB_TOOL=OFF
          \"
      - name: test
        run: cd build && ctest -V
@@ -496,7 +494,7 @@ jobs:
      github.event_name != 'workflow_dispatch' ||
      github.event.inputs.type == 'all' ||
      github.event.inputs.type == 'blosc'
-    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-22.04-8c-32g-300h') || 'ubuntu-latest' }}
+    runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'ubuntu-20.04-8c-32g-300h') || 'ubuntu-latest' }}
    name: linux-blosc:${{ matrix.blosc }}
    container:
      image: aswf/ci-base:2023
@@ -592,7 +590,7 @@ jobs:
        -strict
        -extended
    - name: upload_report
-      uses: actions/upload-artifact@v4
+      uses: actions/upload-artifact@v3
      if: always()
      with:
        name: abi_report
@@ -1,2 +1,4 @@
 /build/*
+**/src/benchmarks/build/*
+nanovdb/nanovdb/build/*
 .DS_Store
@@ -41,3 +41,6 @@ CHANGES                                                     @kmuseth @apradhana
 /openvdb_houdini/openvdb_houdini/PointUtils.*               @danrbailey @idclip
 /openvdb_houdini/openvdb_houdini/*_OpenVDB_Points_*         @danrbailey @idclip
 /openvdb_houdini/openvdb_houdini/GR_PrimVDBPoints.*         @danrbailey @idclip
+
+# fVDB
+/fvdb/                                                      @NVIDIA-Omniverse/fvdb-dev
@@ -0,0 +1,84 @@
+version: 0.2
+
+env:
+  variables:
+    PYTHONPATH: ""
+  # secrets-manager:
+  #   GITHUB_TOKEN: "github-token:token"
+
+phases:
+  install:
+    runtime-versions:
+      python: 3.10
+    commands:
+      # Install Conda
+      - wget https://github.com/conda-forge/miniforge/releases/latest/download/Miniforge3-Linux-x86_64.sh
+      - bash Miniforge3-Linux-x86_64.sh -b -p $HOME/miniforge3
+      - eval "$($HOME/miniforge3/bin/conda shell.bash hook)"
+      - conda init
+      - source ~/.bashrc || true
+
+      # Create build environment
+      - conda env create -f fvdb/env/build_environment.yml -n fvdb_build
+      - echo "Build environment created successfully"
+
+      # Create test environment
+      - conda env create -f fvdb/env/test_environment.yml -n fvdb_test
+      - echo "Test environment created successfully"
+
+  pre_build:
+    commands:
+      - echo "Starting pre_build phase"
+      - git config --global --add safe.directory "$(pwd)"
+
+      # Verify environments exist
+      - conda env list
+      - echo "Pre-build phase completed"
+
+  build:
+    commands:
+      - echo "Starting build phase"
+
+      # Build fvdb wheel using conda run (this ensures proper environment activation)
+      - cd fvdb
+      - export TORCH_CUDA_ARCH_LIST="7.0;7.5;8.0;8.6+PTX"
+      - echo "Building wheel with conda run..."
+      - conda run --no-capture-output -n fvdb_build ./build.sh wheel verbose gtests benchmarks
+
+      # Verify wheel was created
+      - ls -la dist/
+      - echo "Wheel build completed"
+
+      # Move back to root directory
+      - cd ..
+
+      # Install package in test environment using conda run
+      - echo "Installing wheel in test environment..."
+      - conda run --no-capture-output -n fvdb_test pip install ./fvdb/dist/*.whl
+      - echo "Package installation completed"
+
+  post_build:
+    commands:
+      - echo "Starting post_build phase"
+
+      # Run unit tests using conda run
+      - echo "Running unit tests..."
+      - cd fvdb/tests
+      - conda run --no-capture-output -n fvdb_test pytest -v unit
+
+      # Run documentation tests using conda run
+      - echo "Running documentation tests..."
+      - cd ..
+      - conda run --no-capture-output -n fvdb_test pytest --markdown-docs docs
+
+      - echo "All tests completed successfully"
+
+artifacts:
+  files:
+    - fvdb/dist/*.whl
+  name: fvdb-artifacts
+
+cache:
+  paths:
+    - '/root/.cache/pip/**/*'
+    - '$HOME/miniforge3/**/*'
@@ -0,0 +1,24 @@
+# SPDX-License-Identifier: MIT
+#
+# SPDX-FileCopyrightText: Copyright (c) 2019-2023 Lars Melchior and contributors
+
+set(CPM_DOWNLOAD_VERSION 0.40.2)
+set(CPM_HASH_SUM "c8cdc32c03816538ce22781ed72964dc864b2a34a310d3b7104812a5ca2d835d")
+
+if(CPM_SOURCE_CACHE)
+  set(CPM_DOWNLOAD_LOCATION "${CPM_SOURCE_CACHE}/cpm/CPM_${CPM_DOWNLOAD_VERSION}.cmake")
+elseif(DEFINED ENV{CPM_SOURCE_CACHE})
+  set(CPM_DOWNLOAD_LOCATION "$ENV{CPM_SOURCE_CACHE}/cpm/CPM_${CPM_DOWNLOAD_VERSION}.cmake")
+else()
+  set(CPM_DOWNLOAD_LOCATION "${CMAKE_BINARY_DIR}/cmake/CPM_${CPM_DOWNLOAD_VERSION}.cmake")
+endif()
+
+# Expand relative path. This is important if the provided path contains a tilde (~)
+get_filename_component(CPM_DOWNLOAD_LOCATION ${CPM_DOWNLOAD_LOCATION} ABSOLUTE)
+
+file(DOWNLOAD
+     https://github.com/cpm-cmake/CPM.cmake/releases/download/v${CPM_DOWNLOAD_VERSION}/CPM.cmake
+     ${CPM_DOWNLOAD_LOCATION} EXPECTED_HASH SHA256=${CPM_HASH_SUM}
+)
+
+include(${CPM_DOWNLOAD_LOCATION})
@@ -0,0 +1,44 @@
+include (FindPackageHandleStandardArgs)
+
+if(NCCL_ROOT)
+  set(_NCCL_ROOT ${NCCL_ROOT})
+elseif(DEFINED ENV{NCCL_ROOT})
+  set(_NCCL_ROOT $ENV{NCCL_ROOT})
+endif()
+
+set(_NCCL_INCLUDE_SEARCH_DIRS "")
+list(APPEND _NCCL_INCLUDE_SEARCH_DIRS
+  ${_NCCL_ROOT}/include
+)
+find_path(NCCL_INCLUDE_DIR
+    NAMES nccl.h
+    PATHS ${_NCCL_INCLUDE_SEARCH_DIRS}
+)
+
+set(_NCCL_LIBRARY_SEARCH_DIRS "")
+list(APPEND _NCCL_LIBRARY_SEARCH_DIRS
+  ${_NCCL_ROOT}/bin
+)
+find_library(NCCL_LIBRARY
+    NAMES nccl
+    PATHS ${_NCCL_LIBRARY_SEARCH_DIRS}
+)
+
+find_package_handle_standard_args(NCCL REQUIRED_VARS NCCL_INCLUDE_DIR NCCL_LIBRARY)
+
+if (NCCL_FOUND)
+    mark_as_advanced(NCCL_INCLUDE_DIR)
+    mark_as_advanced(NCCL_LIBRARY)
+endif()
+
+if (NCCL_FOUND AND NOT TARGET NCCL::NCCL)
+    add_library(NCCL::NCCL SHARED IMPORTED)
+    if(MSVC)
+      string(REPLACE ".lib" ".dll" NCCL_DLL ${NCCL_LIBRARY})
+      set_property(TARGET NCCL::NCCL PROPERTY IMPORTED_LOCATION ${NCCL_DLL})
+      set_property(TARGET NCCL::NCCL PROPERTY IMPORTED_IMPLIB ${NCCL_LIBRARY})
+    else()
+      set_property(TARGET NCCL::NCCL PROPERTY IMPORTED_LOCATION ${NCCL_LIBRARY})
+    endif()
+    target_include_directories(NCCL::NCCL INTERFACE ${NCCL_INCLUDE_DIR})
+endif()
@@ -0,0 +1,139 @@
+---
+Language: Cpp
+Standard: c++20
+
+# Indentation
+TabWidth: 4
+IndentWidth: 4
+UseTab: Never
+IndentPPDirectives: None
+IndentWrappedFunctionNames: false
+NamespaceIndentation: None
+
+# Empty lines
+KeepEmptyLinesAtTheStartOfBlocks: false
+MaxEmptyLinesToKeep: 1
+
+# Line length
+ColumnLimit: 100
+
+# Line endings
+DeriveLineEnding: false
+LineEnding: LF
+
+# Breaking and Penalties
+BreakBeforeBinaryOperators: None
+BreakBeforeBraces: Attach
+BreakBeforeTernaryOperators: true
+BreakConstructorInitializersBeforeComma: false
+BreakInheritanceList: BeforeColon
+BreakStringLiterals: false
+
+PenaltyBreakAssignment: 2
+PenaltyBreakBeforeFirstCallParameter: 19
+PenaltyBreakComment: 300
+PenaltyBreakFirstLessLess: 120
+PenaltyBreakString: 1000
+PenaltyExcessCharacter: 1000000
+PenaltyReturnTypeOnItsOwnLine: 60
+
+# Spacing and padding
+SpaceAfterCStyleCast: false
+SpaceAfterTemplateKeyword: true
+SpaceAfterLogicalNot: false
+SpaceAfterControlStatementKeyword: true
+SpaceBeforeAssignmentOperators: true
+SpaceBeforeParens: ControlStatements
+SpaceBeforeRangeBasedForLoopColon: false
+SpaceInEmptyParentheses: false
+SpacesInAngles: false
+SpacesInContainerLiterals: false
+SpacesInCStyleCastParentheses: false
+SpacesInParentheses: false
+SpacesBeforeTrailingComments: 1
+SpaceBeforeCpp11BracedList: false
+SpaceBeforeCtorInitializerColon: true
+SpaceBeforeInheritanceColon: true
+SpaceBeforeSquareBrackets: false
+SpacesInSquareBrackets: false
+SpaceBeforeCaseColon: false
+
+# Brace placement
+BraceWrapping:
+  AfterClass: false
+  AfterControlStatement: false
+  AfterEnum: false
+  AfterFunction: false
+  AfterNamespace: false
+  AfterStruct: false
+  AfterUnion: false
+  BeforeCatch: false
+  BeforeElse: false
+  IndentBraces: false
+  SplitEmptyFunction: false
+  SplitEmptyRecord: false
+  SplitEmptyNamespace: false
+Cpp11BracedListStyle: true
+
+# Function definitions
+# BreakAfterReturnType: AllDefinitions
+AlwaysBreakAfterDefinitionReturnType: All
+
+AttributeMacros:
+  - __host__
+  - __device__
+  - __hostdev__
+  - __global__
+  - __forceinline__
+  - __shared__
+  - __launch_bounds__
+
+# Alignment
+AlignConsecutiveAssignments: true
+AlignConsecutiveBitFields: true
+AlignConsecutiveDeclarations: false
+AlignConsecutiveMacros: true
+AlignEscapedNewlines: Left
+AlignOperands: true
+AlignTrailingComments:
+  Kind: Always
+  OverEmptyLines: 2
+
+# Single line allowances
+BinPackParameters: false
+BinPackArguments: false
+AllowAllArgumentsOnNextLine: true
+AllowAllConstructorInitializersOnNextLine: true
+AllowAllParametersOfDeclarationOnNextLine: true
+AllowShortBlocksOnASingleLine: true
+AllowShortCaseLabelsOnASingleLine: true
+AllowShortEnumsOnASingleLine: true
+AllowShortFunctionsOnASingleLine: Inline
+AllowShortIfStatementsOnASingleLine: false
+AllowShortLambdasOnASingleLine: true
+AllowShortLoopsOnASingleLine: false
+
+# Sorting
+IncludeBlocks: Regroup
+IncludeIsMainRegex: "$"
+IncludeCategories:
+  - Regex:           '^"' # quoted includes
+    Priority:        1
+  - Regex:           '^<fvdb/' # fvdb includes
+    Priority:        2
+  - Regex:           '^<nanovdb/' # nanovdb includes
+    Priority:        3
+  - Regex:           '^<(torch|c10|THC|ATen)/' # torch includes
+    Priority:        4
+  - Regex:           '^<(thrust|cub|cuda)/' # CCCL includes
+    Priority:        5
+  - Regex:           '^<(cooperative_groups|cuco|cuda.h|cuda_|device_types|math_constants|nvtx3|cute)' # CUDA includes
+    Priority:        6
+  - Regex:           '^<.*\..*' # other system includes (e.g. with a '.')
+    Priority:        7
+  - Regex:           '^<[^.]+' # STL includes (no '.')
+    Priority:        8
+IncludeIsMainSourceRegex: "$"
+SortIncludes: true
+SortUsingDeclarations: true
+
@@ -0,0 +1,71 @@
+# https://clangd.llvm.org/config
+
+# Apply a config conditionally to all C files
+If:
+  PathMatch: .*\.(c|h)$
+
+---
+
+# Apply a config conditionally to all C++ files
+If:
+  PathMatch: .*\.(c|h)pp
+
+---
+
+# Apply a config conditionally to all CUDA files
+If:
+  PathMatch: .*\.cuh?
+CompileFlags:
+  Add:
+    - "-x"
+    - "cuda"
+    # No error on unknown CUDA versions
+    - "-Wno-unknown-cuda-version"
+    # Allow variadic CUDA functions
+    - "-Xclang=-fcuda-allow-variadic-functions"
+Diagnostics:
+  Suppress:
+    - "variadic_device_fn"
+    - "attributes_not_allowed"
+    - "unknown_cuda_version"
+    - "unknown_type_name"
+    - "implicit_return_from_non_void_function"
+    - "no_template_args_on_forwarding_function"
+
+---
+
+# Tweak the clangd parse settings for all files
+CompileFlags:
+  Add:
+    # report all errors
+    - "-ferror-limit=0"
+    - "-fmacro-backtrace-limit=0"
+    - "-ftemplate-backtrace-limit=0"
+    # Skip the CUDA version check
+    - "--no-cuda-version-check"
+  Remove:
+    # remove gcc's -fcoroutines
+    - -fcoroutines
+    # remove nvc++ flags unknown to clang
+    - "-gpu=*"
+    - "-stdpar*"
+    # remove nvcc flags unknown to clang
+    - "-arch*"
+    - "-gencode*"
+    - "--generate-code*"
+    - "-ccbin*"
+    - "-t=*"
+    - "--threads*"
+    - "-Xptxas*"
+    - "-Xcudafe*"
+    - "-Xfatbin*"
+    - "-Xcompiler*"
+    - "--diag-suppress*"
+    - "--diag_suppress*"
+    - "--compiler-options*"
+    - "--extended-lambda"
+    - "--expt-extended-lambda"
+    - "--expt-relaxed-constexpr"
+    - "-forward-unknown-to-host-compiler"
+    - "-Werror=cross-execution-space-call"
+    - "-Werror=all-warnings"
@@ -0,0 +1,41 @@
+// File location: openvdb/fvdb/.devcontainer/devcontainer.json
+{
+    // A friendly name for the container configuration shown in VS Code
+    "name": "FVDB Development Environment",
+
+    // Tells VS Code to use Docker Compose to define and run the environment.
+    // The path is relative to this devcontainer.json file.
+    "dockerComposeFile": [
+        "../docker-compose.yml" // Points to openvdb/fvdb/docker-compose.yml
+    ],
+
+    // The name of the specific service within the docker-compose.yml file
+    // that VS Code should connect to.
+    "service": "fvdb-dev",
+
+    // Specifies the default directory VS Code should open *inside* the container.
+    // Since the docker-compose.yml mounts the host's `openvdb/` directory to `/openvdb`
+    // in the container, the fvdb code resides at `/openvdb/fvdb`.
+    "workspaceFolder": "/openvdb/fvdb",
+
+    // Extensions to install inside the container's VS Code server
+    "customizations": {
+        "vscode": {
+            "extensions": [
+                "ms-python.python", // Python language support
+                "ms-vscode.cpptools", // C/C++ language support
+                "ms-vscode.cmake-tools", // CMake integration
+                "ms-azuretools.vscode-docker" // Docker integration (useful for managing containers)
+            ],
+            "settings": {
+                "python.defaultInterpreterPath": "/opt/conda/envs/fvdb/bin/python"
+            }
+        }
+    },
+
+    // Optional: What to do when the VS Code window is closed.
+    // "stopCompose" will stop the container(s) defined in the compose file.
+    // "none" would leave them running.
+    "shutdownAction": "stopCompose"
+}
+
@@ -0,0 +1,106 @@
+name: Building Release Wheels
+
+on:
+  release:
+    types: [published, edited]
+  workflow_dispatch:
+
+jobs:
+
+  wheel:
+    runs-on:
+      - self-hosted
+
+    strategy:
+      fail-fast: false
+      matrix:
+        # support version based on: https://download.pytorch.org/whl/torch/
+        python-version: ['3.8', '3.9', '3.10', '3.11', '3.12']
+        torch-version: [2.1.0, 2.1.1, 2.1.2, 2.2.0, 2.2.1, 2.2.2, 2.3.0]
+        cuda-version: ['12.1.1']
+        exclude:
+          - torch-version: 2.1.2
+            python-version: '3.12'
+          - torch-version: 2.1.1
+            python-version: '3.12'
+          - torch-version: 2.1.0
+            python-version: '3.12'
+
+    steps:
+      - uses: actions/checkout@v2
+      - name: Generate release_environment.yml
+        run: |
+          echo "cp env/release_base_environment.yml release_environment.yml"
+          cp env/release_base_environment.yml release_environment.yml
+          echo "  - python=${{ matrix.python-version }}" >> release_environment.yml
+          echo "  - pytorch=${{ matrix.torch-version }}" >> release_environment.yml
+          perl -E '"${{ matrix.cuda-version }}" =~ /^(\d+\.\d+)/; say "  - pytorch-cuda=$1"' >> release_environment.yml
+          perl -E '"${{ matrix.cuda-version }}" =~ /^(\d+)\.(\d+)/; say "  - cuda>=$1.$2.0, <$1." . ($2+1) . ".0"' >> release_environment.yml
+          perl -E '"${{ matrix.cuda-version }}" =~ /^(\d+)\.(\d+)/; say "  - cuda-nvcc>=$1.$2.0, <$1." . ($2+1) . ".0"' >> release_environment.yml
+
+      - name: Set up fvdb_release Conda env
+        uses: conda-incubator/setup-miniconda@v3
+        with:
+          miniconda-version: "latest"
+          mamba-version: "*"
+          channels: nvidia,pytorch,conda-forge,defaults
+          channel-priority: true
+          activate-environment: fvdb_release
+          environment-file: release_environment.yml
+
+      - name: Set version
+        run: |
+          VERSION=`sed -n "s/^__version__ = '\(.*\)'/\1/p" fvdb/__init__.py`
+          TORCH_VERSION=`echo "pt${{ matrix.torch-version }}" | sed "s/..$//" | sed "s/\.//g"`
+          CUDA_VERSION=`perl -E '"${{ matrix.cuda-version }}" =~ /^(\d+)\.(\d+)/; say "cu$1$2"'`
+          echo "New version name: $VERSION+$TORCH_VERSION$CUDA_VERSION"
+          sed -i "s/$VERSION/$VERSION+$TORCH_VERSION$CUDA_VERSION/" fvdb/__init__.py
+        shell:
+          bash -el {0}
+
+      - name: Build wheel
+        run: |
+          conda activate fvdb_release
+          mkdir dist
+          TORCH_CUDA_ARCH_LIST="6.1;7.0;7.5;8.0;8.6+PTX" python setup.py bdist_wheel --dist-dir=dist
+        shell:
+          bash -el {0}
+
+      - name: Rename wheel
+        run: |
+          python scripts/rename_wheels.py ${{ matrix.torch-version }} `perl -E '"${{ matrix.cuda-version }}" =~ /^(\d+)\.(\d+)/; say "cu$1$2"'`
+        shell:
+          bash -el {0}
+
+
+      - name: Release
+        uses: softprops/action-gh-release@v1
+        if: startsWith(github.ref, 'refs/tags/')
+        with:
+          token: ${{ secrets.RELEASE_ACTION_TOKEN }}
+          files: |
+            dist/*.whl
+
+      # - name: Upload to artifacts
+      #   uses: actions/upload-artifact@v3
+      #   with:
+      #     name: wheels
+      #     path: |
+      #       dist/*.whl
+
+      # - name: Configure AWS
+      #   uses: aws-actions/configure-aws-credentials@v1
+      #   with:
+      #     aws-access-key-id: ${{ secrets.AWS_ACCESS_KEY_ID }}
+      #     aws-secret-access-key: ${{ secrets.AWS_SECRET_ACCESS_KEY }}
+      #     aws-region: ap-northeast-1
+
+      # - name: Upload wheel
+      #   run: |
+      #     aws s3 sync dist s3://nksr/whl/torch-${{ matrix.torch-version }}+${{ matrix.cuda-version }} --grants read=uri=http://acs.amazonaws.com/groups/global/AllUsers
+
+      - name: Clean Conda
+        run: |
+          conda clean -pty
+        shell:
+          bash -el {0}
@@ -0,0 +1,221 @@
+name: Running Tests
+
+on:
+  pull_request:
+    branches: [main]
+    paths-ignore:
+      - 'debug/**'
+      - 'docs/**'
+      - 'examples/**'
+      - 'notebooks/**'
+      - 'scripts/**'
+
+# Allow subsequent pushes to the same PR or REF to cancel any previous jobs.
+concurrency:
+  group: ${{ github.workflow }}-${{ github.event.pull_request.number || github.ref }}
+  cancel-in-progress: true
+
+permissions:
+  contents: write
+  deployments: write
+
+jobs:
+  fvdb-build:
+    if: ${{ !startsWith(github.event.pull_request.title, 'Draft:') }}
+    name: fVDB Build
+    runs-on:
+      - self-hosted
+    container:
+      image: aswf/ci-openvdb:2024
+      env:
+        PYTHONPATH: ""
+      options: --rm
+    defaults:
+      run:
+        shell: bash -el {0}
+    steps:
+      - uses: actions/checkout@v4
+
+      - name: Set up fvdb_build Conda env
+        uses: conda-incubator/setup-miniconda@v3
+        with:
+          miniforge-version: latest
+          conda-remove-defaults: "true"
+          activate-environment: fvdb_build
+          environment-file: env/build_environment.yml
+
+
+      - name: Buid fvdb
+        run: |
+          TORCH_CUDA_ARCH_LIST="7.0;7.5;8.0;8.6+PTX" MAX_JOBS=$(($(nproc) < $(free -g | awk '/^Mem:/{jobs=int($4/2.5); if(jobs<1) jobs=1; print jobs}') ? $(nproc) : $(free -g | awk '/^Mem:/{jobs=int($4/2.5); if(jobs<1) jobs=1; print jobs}'))) conda run --no-capture-output -n fvdb_build python setup.py bdist_wheel --dist-dir=dist
+
+      - name: Upload package
+        uses: actions/upload-artifact@v4
+        with:
+            name: fvdb-test-package
+            path: dist/*.whl
+            retention-days: 2
+
+      - name: Cleanup
+        if: always()
+        run: |
+          echo "Cleaning up /__w/_temp directory"
+          sudo rm -rf /__w/_temp/*
+          echo "Cleanup completed"
+
+
+  fvdb-unit-test:
+    needs: [fvdb-build]
+    name: fVDB Unit Tests
+    runs-on:
+      - self-hosted
+    container:
+      image: aswf/ci-openvdb:2024
+      env:
+        PYTHONPATH: ""
+      options: --rm
+    defaults:
+      run:
+        shell: bash -el {0}
+    steps:
+      - uses: actions/checkout@v4
+      - name: Set up fvdb_test Conda env
+        uses: conda-incubator/setup-miniconda@v3
+        with:
+          miniforge-version: latest
+          conda-remove-defaults: "true"
+          activate-environment: fvdb_test
+          environment-file: env/test_environment.yml
+
+      - name: Download package
+        uses: actions/download-artifact@v4
+        with:
+            name: fvdb-test-package
+            path: ./dist
+
+      - name: Install package
+        run: |
+            conda activate fvdb_test
+            pip install ./dist/*.whl
+
+      - name: Run tests
+        run: |
+            cd tests;
+            pytest -v unit
+
+      - name: Cleanup
+        if: always()
+        run: |
+          echo "Cleaning up /__w/_temp directory"
+          sudo rm -rf /__w/_temp/*
+          echo "Cleanup completed"
+
+  fvdb-docs-test:
+    needs: [fvdb-build]
+    name: fVDB Documentation Tests
+    runs-on:
+      - self-hosted
+    container:
+      image: aswf/ci-openvdb:2024
+      env:
+        PYTHONPATH: ""
+      options: --rm
+    defaults:
+      run:
+        shell: bash -el {0}
+    steps:
+      - uses: actions/checkout@v4
+      - name: Set up fvdb_test Conda env
+        uses: conda-incubator/setup-miniconda@v3
+        with:
+          miniforge-version: latest
+          conda-remove-defaults: "true"
+          activate-environment: fvdb_test
+          environment-file: env/test_environment.yml
+
+      - name: Download package
+        uses: actions/download-artifact@v4
+        with:
+            name: fvdb-test-package
+            path: ./dist
+
+      - name: Install package
+        run: |
+            conda activate fvdb_test
+            pip install ./dist/*.whl
+
+      - name: Run tests
+        run: |
+            cd .. && pytest --markdown-docs fvdb/docs
+
+      - name: Cleanup
+        if: always()
+        run: |
+          echo "Cleaning up /__w/_temp directory"
+          sudo rm -rf /__w/_temp/*
+          echo "Cleanup completed"
+
+  fvdb-benchmarks:
+    needs: [fvdb-build]
+    name: fVDB Continuous Benchmarking
+    runs-on:
+      - self-hosted
+    container:
+      image: aswf/ci-openvdb:2024
+      env:
+        PYTHONPATH: ""
+      options: --rm
+    defaults:
+      run:
+        shell: bash -el {0}
+    steps:
+      - uses: actions/checkout@v4
+      - name: Set up fvdb_test Conda env
+        uses: conda-incubator/setup-miniconda@v3
+        with:
+          miniforge-version: latest
+          conda-remove-defaults: "true"
+          activate-environment: fvdb_test
+          environment-file: env/test_environment.yml
+
+      - name: Download package
+        uses: actions/download-artifact@v4
+        with:
+            name: fvdb-test-package
+            path: ./dist
+
+      - name: Install package
+        run: |
+            conda activate fvdb_test
+            pip install ./dist/*.whl
+
+      - name: Disable git ownership verification
+        run: |
+          git config --global --add safe.directory "$(pwd)"
+
+      - name: Run benchmarks
+        run: |
+          cd tests;
+          pytest benchmark --benchmark-json benchmark/output.json
+
+      - name: Store benchmark result
+        uses: benchmark-action/github-action-benchmark@v1
+        with:
+            name: Python Benchmark with pytest-benchmark
+            tool: 'pytest'
+            output-file-path: tests/benchmark/output.json
+            # Use personal access token instead of GITHUB_TOKEN due to https://github.community/t/github-action-not-triggering-gh-pages-upon-push/16096
+            github-token: ${{ secrets.GITHUB_TOKEN }}
+            auto-push: true
+            # Show alert with commit comment on detecting possible performance regression
+            alert-threshold: '200%'
+            comment-on-alert: true
+            fail-on-alert: true
+            alert-comment-cc-users: '@swahtz'
+
+      - name: Cleanup
+        if: always()
+        run: |
+          echo "Cleaning up /__w/_temp directory"
+          sudo rm -rf /__w/_temp/*
+          echo "Cleanup completed"
@@ -0,0 +1,60 @@
+# Generated version file
+fvdb/version.py
+
+# Generated compile_commands.json symlink
+compile_commands.json
+
+# Build files
+/build/*
+/src/benchmarks/build/*
+/src/tests/build/*
+/src/build/*
+
+*.creator
+*.includes
+*.files
+*.config
+*.obj
+*.off
+__pycache__
+*.pyc
+*.creator.user
+*.egg-info
+*.so
+dist
+.cache
+.idea
+.eggs
+cmake-build-debug
+external
+.polyscope.ini
+imgui.ini
+datasets/lib/external
+datasets/lib/build
+*.png
+!docs/imgs/**/*.png
+!notebooks/img/**/*.png
+*.ply
+apex/
+*.zip
+*.npz
+!kernel.npz
+*.npy
+*.o
+*.h5
+docs/_build
+fvdb/include/
+trash
+scratch
+_build
+docs/build
+*tmp*.py
+*.whl
+tests/benchmark/data
+releases/
+/data/
+/data/*
+!tests/data
+.vscode/launch.json
+lightning_logs
+tests/data/
@@ -0,0 +1,30 @@
+# Gitlab CI Pipeline for building and deploying the documentation
+default:
+  image: mambaorg/micromamba:focal-cuda-12.1.1
+  tags:
+    - os/linux
+    - type/docker
+
+before_script:
+    # - apt-get update
+    - export TORCH_CUDA_ARCH_LIST="Turing"
+    - export FORCE_CUDA=1
+    # install all required libraries using existing conda environment requirements file:
+    - micromamba env create -f env/build_environment.yml
+    - eval "$(micromamba shell hook --shell bash)"
+    - micromamba activate
+    - micromamba activate fvdb_build
+
+pages:
+  stage: deploy
+  script:
+    # build and publish automated documentation with Sphinx:
+    - export MAX_JOBS=$(free -g | awk '/^Mem:/{jobs=int($4/3); if(jobs<1) jobs=1; print jobs}')
+    - python setup.py build_ext --inplace
+    - sphinx-build -E -a docs/ public/
+
+  artifacts:
+    paths:
+    - public
+  only:
+  - main # Documentation will only be generated on the main branch
@@ -0,0 +1,51 @@
+{
+    "configurations": [
+        {
+            "name": "Linux",
+            "includePath": [
+                "${env:CONDA_PREFIX}/envs/fvdb/include",
+                "${env:CONDA_PREFIX}/envs/fvdb/targets/x86_64-linux/include/",
+                "${workspaceFolder}/src/tests/build/_deps/googletest-src/googletest/include",
+                "${workspaceFolder}/src/tests/build/_deps/googletest-src/googlemock/include",
+                "${workspaceFolder}/src/tests/build/_deps/nvtx3-src/include",
+                "${workspaceFolder}/src",
+                "${workspaceFolder}/../nanovdb",
+                "${workspaceFolder}/external/glm",
+                "${workspaceFolder}/external/cudnn_fe/include",
+                "${workspaceFolder}/external/cutlass/include",
+                "${workspaceFolder}/external/tinyply/source",
+                "${env:CONDA_PREFIX}/envs/fvdb/include/python3.11",
+                "${env:CONDA_PREFIX}/envs/fvdb/lib/python3.11/site-packages/torch/include",
+                "${env:CONDA_PREFIX}/envs/fvdb/lib/python3.11/site-packages/torch/include/torch/csrc/api/include",
+                "${env:CONDA_PREFIX}/envs/fvdb/lib/python3.11/site-packages/torch/include/THC",
+                "${env:CONDA_PREFIX}/envs/fvdb/lib/python3.11/site-packages/torch/include/c10",
+                "${default}"
+            ],
+            "macFrameworkPath": [
+                "${default}"
+            ],
+            "defines": [
+                "${default}",
+                "NDEVELOP_IDE_ONLY",
+                "__CUDA_ARCH__=800"
+            ],
+            "forcedInclude": [
+                "${default}"
+            ],
+            "compileCommands": "${default}",
+            "browse": {
+                "limitSymbolsToIncludedHeaders": true,
+                "databaseFilename": "${default}",
+                "path": [
+                    "${default}"
+                ]
+            },
+            "intelliSenseMode": "${default}",
+            "cStandard": "${default}",
+            "cppStandard": "c++20",
+            "compilerPath": "${default}",
+            "configurationProvider": "ms-vscode.cpptools"
+        }
+    ],
+    "version": 4
+}
@@ -0,0 +1,15 @@
+{
+    // See https://go.microsoft.com/fwlink/?LinkId=827846 to learn about workspace recommendations.
+    // Extension identifier format: ${publisher}.${name}. Example: vscode.csharp
+
+    // List of extensions which should be recommended for users of this workspace.
+    "recommendations": [
+        "llvm-vs-code-extensions.vscode-clangd",
+        "ms-vscode.cpptools-extension-pack",
+        "Seaube.clangformat"
+    ],
+    // List of extensions recommended by VS Code that should not be recommended for users of this workspace.
+    "unwantedRecommendations": [
+
+    ]
+}
@@ -0,0 +1,58 @@
+{
+    "files.associations": {
+        "*.cu": "cuda-cpp",
+        "*.cuh": "cuda-cpp"
+    },
+    "files.trimTrailingWhitespace": true,
+    "files.insertFinalNewline": true,
+    "editor.rulers": [ 100 ],
+    "C_Cpp.intelliSenseEngine": "disabled",
+    "C_Cpp.enhancedColorization": "enabled",
+    "C_Cpp.errorSquiggles": "enabledIfIncludesResolve",
+    "C_Cpp.autocomplete": "default",
+    "C_Cpp.default.compilerPath": "${env:CONDA_PREFIX}/envs/fvdb/bin/x86_64-conda-linux-gnu-c++",
+    "terminal.integrated.defaultLocation": "editor",
+    "cmake.environment": {
+        "PATH": "${env:CONDA_PREFIX}/envs/fvdb/bin:${env:PATH}"
+    },
+    "cmake.configureOnOpen": false,
+    "C_Cpp.formatting": "clangFormat",
+    "C_Cpp.clang_format_path": "${env:CONDA_PREFIX}/envs/fvdb/bin/clang-format-18",
+    "C_Cpp.autoAddFileAssociations": false,
+
+    "[cpp]": {
+        "editor.defaultFormatter": "seaube.clangformat",
+        "editor.formatOnSave": true
+    },
+    "[cuda-cpp]": {
+        "editor.defaultFormatter": "seaube.clangformat",
+        "editor.formatOnSave": true
+    },
+    "[python]": {
+        "editor.defaultFormatter": "ms-python.black-formatter",
+        "editor.formatOnSave": true,
+        "editor.codeActionsOnSave": {
+            "source.organizeImports": "explicit"
+        },
+    },
+    "python.analysis.typeCheckingMode": "basic",
+    "python.testing.unittestArgs": [
+        "-v", "-s", "."
+    ],
+    "python.testing.pytestEnabled": false,
+    "python.testing.unittestEnabled": true,
+    "black-formatter.args": [
+        "--target-version=py311",
+        "--line-length=120"
+    ],
+    "python.analysis.diagnosticSeverityOverrides": {
+        "reportPrivateImportUsage": "none",
+    },
+    "isort.args": [
+        "--profile", "black"
+    ],
+    "json.format.keepLines": true,
+
+    "clangd.path": "${env:CONDA_PREFIX}/envs/fvdb/bin/clangd",
+    "clangFormat.executable.linux": "${env.CONDA_PREFIX}/envs/fvdb/bin/clang-format",
+}
@@ -0,0 +1,98 @@
+# Copyright Contributors to the OpenVDB Project
+# SPDX-License-Identifier: Apache-2.0
+
+cmake_minimum_required(VERSION 3.25 FATAL_ERROR)
+
+# Set global download timeout to 600 seconds (10 minutes)
+set(CMAKE_DOWNLOAD_TIMEOUT 600 CACHE STRING "Time allowed for downloads in seconds")
+
+project(${SKBUILD_PROJECT_NAME} VERSION ${SKBUILD_PROJECT_VERSION} LANGUAGES CXX)
+
+find_package(Python3 COMPONENTS Development Interpreter)
+
+message(STATUS "CMAKE_BUILD_PARALLEL_LEVEL: ${CMAKE_BUILD_PARALLEL_LEVEL}")
+
+include(${CMAKE_CURRENT_SOURCE_DIR}/src/cmake/get_cpm.cmake)
+include(${CMAKE_CURRENT_SOURCE_DIR}/src/cmake/get_nvtx.cmake)
+include(${CMAKE_CURRENT_SOURCE_DIR}/src/cmake/get_torch.cmake)
+include(${CMAKE_CURRENT_SOURCE_DIR}/src/cmake/get_pybind11.cmake)
+include(${CMAKE_CURRENT_SOURCE_DIR}/src/cmake/configure_torch_pybind11.cmake)
+
+# Generate compile_commands.json
+set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
+
+# configure libfvdb
+add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/src/)
+
+# Source files
+set(FVDB_BINDINGS_CPP_FILES
+    src/python/Bindings.cpp
+    src/python/GaussianSplatBinding.cpp
+    src/python/GridBatchBinding.cpp
+    src/python/JaggedTensorBinding.cpp)
+
+# Build library
+Python3_add_library(_Cpp ${FVDB_BINDINGS_CPP_FILES})
+if (DEFINED TORCH_PYBIND11_INCLUDE_DIR)
+    target_link_libraries(_Cpp PUBLIC torch_pybind11_headers)
+else()
+    target_link_libraries(_Cpp PUBLIC pybind11::headers)
+endif()
+set_target_properties(_Cpp PROPERTIES
+                                INTERPROCEDURAL_OPTIMIZATION ON
+                                CXX_VISIBILITY_PRESET hidden
+                                VISIBILITY_INLINES_HIDDEN ON)
+
+target_link_options(_Cpp PRIVATE "-Wl,--no-as-needed")
+
+set_target_properties(_Cpp PROPERTIES
+    BUILD_RPATH "\$ORIGIN"
+    INSTALL_RPATH "\$ORIGIN"
+    INSTALL_RPATH_USE_LINK_PATH FALSE
+    BUILD_WITH_INSTALL_RPATH TRUE
+    CXX_STANDARD 20
+    CXX_STANDARD_REQUIRED ON
+    CXX_EXTENSIONS ON
+    POSITION_INDEPENDENT_CODE ON
+    INTERFACE_POSITION_INDEPENDENT_CODE ON)
+
+target_include_directories(_Cpp PUBLIC
+    ${CMAKE_CURRENT_SOURCE_DIR}/src
+    ${TORCH_INCLUDE_DIRS})
+
+target_link_libraries(_Cpp PRIVATE
+    ${TORCH_PYTHON_LIBRARY}
+    ${TORCH_LIBRARIES}
+    fvdb)
+
+target_compile_options(_Cpp PRIVATE
+    $<$<COMPILE_LANGUAGE:CXX>:
+    #"-Wno-unused-variable"
+    "-fvisibility=hidden"
+    "-fvisibility-inlines-hidden"
+    "-Wall"
+    "-Werror"
+    "${TORCH_CXX_FLAGS}"
+    >)
+
+target_compile_definitions(_Cpp PRIVATE
+    TORCH_EXTENSION_NAME=_Cpp
+    TORCH_API_INCLUDE_EXTENSION_H)
+
+# Configure the pybind11 module with PyTorch compatibility settings
+configure_torch_pybind11(_Cpp)
+
+install(TARGETS _Cpp fvdb
+    LIBRARY
+    DESTINATION fvdb
+    COMPONENT PythonModule
+    RUNTIME
+    DESTINATION fvdb
+    COMPONENT PythonModule)
+
+# Add at the beginning after project() call:
+set(VERSION_PY_PATH "${CMAKE_CURRENT_SOURCE_DIR}/fvdb/version.py")
+configure_file("${CMAKE_CURRENT_SOURCE_DIR}/fvdb/version.py.in" "${VERSION_PY_PATH}" @ONLY)
+
+# Make sure version.py is included in the installation
+install(FILES "${VERSION_PY_PATH}" DESTINATION fvdb)
@@ -0,0 +1,19 @@
+FROM condaforge/miniforge3:24.11.3-2
+
+ARG MODE=dev
+RUN echo "Building fVDB container in $MODE mode"
+
+# used for cross-compilation in docker build
+ENV FORCE_CUDA=1
+
+WORKDIR /openvdb
+
+# force this CUDA version to be used to build the docker container because `docker build` does not
+# expose the GPU to the docker build process for it to be detected
+ENV CONDA_OVERRIDE_CUDA=12.0
+# copy env/dev_environment.yml to /tmp/
+COPY fvdb/env/dev_environment.yml /tmp/
+RUN  conda env create -f /tmp/dev_environment.yml
+
+RUN conda init
+RUN echo "conda activate fvdb" >> ~/.bashrc
@@ -0,0 +1,244 @@
+# *ƒ*(VDB)
+
+
+This repository contains the code for *f*VDB, a data structure for encoding and operating on *sparse voxel hierarchies* of features in PyTorch. A sparse voxel hierarchy is a coarse-to-fine hierarchy of sparse voxel grids such that every fine voxel is contained within some coarse voxel. The image below illustrates an example. *f*VDB supports using PyTorch Tensors to represent features at the corners and centers of voxels in a hierarchy and enables a number of differentiable operations on these Tensors (*e.g.* trilinear interpolation, convolution, splatting, ray tracing).
+
+<p align="center">
+  <img src="docs/imgs/fvdb_teaser.png" style="width: 40%;"alt="fVDB Teaser">
+  <!-- <img src="docs/imgs/readme/av_screenshot.png" style="width: 100%;"alt="fVDB Teaser"> -->
+  <figcaption style="text-align: center; font-style: italic;">An example of a sparse voxel hierarchy with 3 levels. Each fine voxel is contained within exactly one coarse voxel.</figcaption>
+</p>
+
+## Learning to Use *f*VDB
+
+After [installing *f*VDB](#installing-fvdb), we recommend starting with our walk-through [notebooks](notebooks) which provide a gentle, illustrated introduction to the main concepts and operations in *f*VDB.
+
+Once familiar with the basics, [Usage Examples](#usage-examples) introduces a few of the practical python scripts that can be further explored in the [examples](examples) directory.
+
+Lastly, our [documentation](docs) provides deeper details on the concepts as well as an exhaustive set of illustrations of all the operations available in *f*VDB and an API reference. The [documentation can be built locally](#building-documentation) or can be accessed online at https://www.openvdb.org/documentation/fvdb/.
+
+## Installing *f*VDB
+
+During the project's initial development stages, it is necessary to [run the build steps](#building-fvdb-from-source) to install ƒVDB. Eventually, ƒVDB will be provided as a pre-built, installable package from anaconda.  We support building the latest ƒVDB version for the following dependent library configurations:
+
+|   PyTorch      | Python      | CUDA |
+| -------------- | ----------- | ------------ |
+|  2.4.0-2.4.1   | 3.10 - 3.12 | 12.1 - 12.4 |
+
+
+
+** Notes:**
+* Linux is the only platform currently supported (Ubuntu >= 20.04 recommended).
+* A CUDA-capable GPU with Ampere architecture or newer (i.e. compute capability >=8.0) is recommended to run the CUDA-accelerated operations in ƒVDB.  A GPU with compute capabililty >=7.0 (Volta architecture) is the minimum requirement but some operations and data types are not supported.
+
+
+## Building *f*VDB from Source
+
+### Environment Management
+ƒVDB is a Python library implemented as a C++ Pytorch extension.  Of course you can build ƒVDB in whatever environment suits you, but we provide three paths to constructing reliable environments for building and running ƒVDB:  using [docker](#setting-up-a-docker-container), [conda](#setting-up-a-conda-environment), or a Python virtual environment.
+
+`conda` tends to be more flexible since reconfiguring toolchains and modules to suit your larger project can be dynamic, but at the same time this can be a more brittle experience compared to using a virtualized `docker` container.  Using `conda` is generally recommended for development and testing, while using `docker` is recommended for CI/CD and deployment.
+
+#### Setting up a Docker Container
+
+Running a docker container is a great way to ensure that you have a consistent environment for building and running ƒVDB.
+
+Our provided [`Dockerfile`](Dockerfile) constructs a Docker image which is ready to build ƒVDB.  The docker image is configured to install miniforge and the `fvdb` conda environment with all the dependencies needed to build and run ƒVDB.
+
+Building and starting the docker image is done by running the following command from the fvdb directory:
+```shell
+docker compose run --rm fvdb-dev
+```
+
+
+When you are ready to build ƒVDB, run the following command within the docker container.  `TORCH_CUDA_ARCH_LIST` specifies which CUDA architectures to build for.
+```shell
+conda activate fvdb;
+cd /openvdb/fvdb;
+TORCH_CUDA_ARCH_LIST="7.0;7.5;8.0;8.6+PTX" \
+./build.sh install
+```
+
+Additional information about the ƒVDB docker setup, with troubleshooting for common errors, can be found
+here: [`ƒVDB Docker`](docs/markdown/docker_readme.md)
+
+#### Setting up a Conda Environment
+
+*f*VDB can be used with any Conda distribution installed on your system. Below is an installation guide using
+[miniforge](https://github.com/conda-forge/miniforge). You can skip steps 1-3 if you already have a Conda installation.
+
+1. Download and Run Install Script. Copy the command below to download and run the [miniforge install script](https://github.com/conda-forge/miniforge?tab=readme-ov-file#unix-like-platforms-macos--linux):
+
+```shell
+curl -L -O "https://github.com/conda-forge/miniforge/releases/latest/download/Miniforge3-$(uname)-$(uname -m).sh"
+bash Miniforge3-$(uname)-$(uname -m).sh
+```
+
+2. Follow the prompts to customize Conda and run the install. Note, we recommend saying yes to enable `conda-init`.
+
+3. Start Conda. Open a new terminal window, which should now show Conda initialized to the `(base)` environment.
+
+4. Create the `fvdb` conda environment. Run the following command from the directory containing this README:
+
+```shell
+conda env create -f env/dev_environment.yml
+```
+
+5. Activate the *f*VDB environment:
+
+```shell
+conda activate fvdb
+```
+
+##### Other available environments
+* `fvdb_build`: Use `env/build_environment.yml` for a minimum set of dependencies needed just to build/package *f*VDB (note this environment won't have all the runtime dependencies needed to `import fvdb`).
+* `fvdb_test`: Use `env/test_environment.yml` for a runtime environment which has only the packages required to run the unit tests after building ƒVDB. This is the environment used by the CI pipeline to run the tests after building ƒVDB in the `fvdb_build` environment.
+* `fvdb_learn`: Use `env/learn_environment.yml` for additional runtime requirements and packages needed to run the [notebooks](notebooks) or [examples](examples) and view their visualizations.
+
+#### Setting up a Python virtual environment
+
+After creating a Python virtual environment, proceed to install the exact version of PyTorch that corresponds to your CUDA version. Then, install the rest of the build requirements.
+
+```shell
+python -m venv fvdb
+source fvdb/bin/activate
+pip install torch==2.6.0 torchvision==0.21.0 torchaudio==2.6.0 --index-url https://download.pytorch.org/whl/cu126
+pip install -r env/build_requirements.txt
+```
+
+When you're ready to build fVDB, run the following command after activating the Python virtual environment
+```shell
+TORCH_CUDA_ARCH_LIST="7.0;7.5;8.0;8.6+PTX" ./build.sh install
+```
+
+### Building *f*VDB
+
+**:warning: Note:** Compilation can be very memory-consuming. As part of our build script, we set the `CMAKE_BUILD_PARALLEL_LEVEL` environment variable to control compilation job parallelism with a value that we find works well for most machines (allowing for one job every 2.5GB of memory) but this can be overridden by setting the `CMAKE_BUILD_PARALLEL_LEVEL` environment variable to a different value.
+
+
+You can either perform an install:
+```shell
+./build.sh
+```
+
+or if you would like to build a packaged wheel for installing in other environments, you can run the following command:
+```shell
+./build.sh wheel
+```
+
+
+### Running Tests
+
+To make sure that everything works by running tests:
+```shell
+cd tests
+pytest unit
+```
+
+### Building Documentation
+
+To build the documentation, simply run:
+```shell
+python setup.py build_ext --inplace
+sphinx-build -E -a docs/ build/sphinx
+# View the docs
+open build/sphinx/index.html
+```
+
+### Setting up Intellisense with clangd in Visual Studio Code
+
+Please see the guide [`Clangd for Intellisense in fVDB`](docs/markdown/clangd.md)
+
+
+## Usage Examples
+The [examples](examples) directory contains a number of useful illustrations using the `fvdb` Python package. The sections below show some notable examples and their outputs. Run all commands from the directory containing this README.
+
+### Trilinear sampling of grids
+```
+python examples/sample_trilinear.py
+```
+This script generates a grid with scalars at the corners of each voxel and samples this grid at points. The visualization below shows the points colored according to their sampled values as well as the values at grid corners.
+<p align="center">
+  <img src="docs/imgs/readme/trilerp.png" style="width: 40%;"alt="fVDB trilinear interpolation demo">
+  <figcaption style="text-align: center; font-style: italic;">Trilinearly interpolate the corner values at the points.</figcaption>
+</p>
+
+
+### Trilinear splatting into grids
+```
+python examples/splat_trilinear.py
+```
+This script splats normals of a point cloud onto grid centers. The green arrows represent the values of the normals splatted onto each grid center
+<p align="center">
+  <img src="docs/imgs/readme/splat.png" style="width: 40%;"alt="fVDB trilinear splatting demo">
+  <figcaption style="text-align: center; font-style: italic;">Splat the normals at the blue points into the center of each grid cell. The green arrows are the splatted normals</figcaption>
+</p>
+
+
+### Tracing voxels along rays (hierarchical DDA)
+```
+python examples/ray_voxel_marching.py
+```
+This script demonstrates finding the first `N` voxels which lie along a ray (returning thier index as well as their entry and exit points).
+<p align="center">
+  <img src="docs/imgs/readme/rayvox.png" style="width: 70%;"alt="fVDB ray voxel marching">
+  <figcaption style="text-align: center; font-style: italic;">Find the voxels (yellow) which intersect the pink rays eminating from the green dot.</figcaption>
+</p>
+
+
+### Tracing contiguous segments along rays
+```
+python examples/ray_segment_marching.py
+```
+This script demonstrates finding the first `N` continuous segments of voxels which lie along a ray (returning thier index as well as their entry and exit points).
+<p align="center">
+  <img src="docs/imgs/readme/rayseg.png" style="width: 70%;"alt="fVDB ray voxel marching">
+  <figcaption style="text-align: center; font-style: italic;">Find the contiguous segments of voxels (red and blue lines) which intersect the cyan rays eminating from the pink dot.</figcaption>
+</p>
+
+
+### Backpropagating through sampling and splatting
+```
+python examples/overfit_sdf.py
+```
+This scripts fits SDF values at a grid corner to the SDF of a mesh using gradient descent.
+<p align="center">
+  <img src="docs/imgs/readme/fitsdf.png" style="width: 70%;"alt="fVDB SDF fitting">
+  <figcaption style="text-align: center; font-style: italic;">SDF values at grid corners (colored dots) fitted using gradient descent to the SDF of a mesh.</figcaption>
+</p>
+
+The following scripts also show how to bakcprop through splatting and sampling with fVDB:
+```
+python scripts/debug_grad_trilerp.py
+```
+```
+python scripts/debug_grad_splat.py
+```
+
+## Code Structure
+The main source code for fVDB lives in the [src](src) directory. There are several important files here:
+* `src/python/Bindings.cpp` exposes functionality directly to Python. It is mainly a wrapper around the core classes such as `fvdb::GridBatch` and `fvdb::JaggedTensor`.
+* `src/GridBatch.h` contains the implementation of `fvdb::GridBatch` which is the core data structure on which fVDB is built. A `GridBatch` acts as a map between `(i, j, k)` integer coordinates and offsets in linear memory. This mapping can be used to perform a host of operations. The methods in this class are mostly lightweight wrappers around a set of CPU and CUDA *kernels*. The function prototypes for these kernels are defined in `src/detail/ops/Ops.h`.
+* `src/detail/ops/Ops.h` contains the function prototypes for the main kernels used by fVDB. Host and device kernel implementations are provided in the `src/detail/ops/*.cu` source files.
+* `src/detail/autograd` contains C++ implementations of PyTorch autograd functions for differentiable operations.  `#include <detail/autograd/Autograd.h>` includes all of the functions in this directory.
+* `src/detail/utils/nanovdb` contains a number of utilities which make it easier to use NanoVDB.
+
+
+## References
+
+Please consider citing this when using *f*VDB in a project. You can use the citation BibTeX:
+
+```bibtex
+@article{williams2024fvdb,
+  title={fVDB: A Deep-Learning Framework for Sparse, Large-Scale, and High-Performance Spatial Intelligence},
+  author={Williams, Francis and Huang, Jiahui and Swartz, Jonathan and Klar, Gergely and Thakkar, Vijay and Cong, Matthew and Ren, Xuanchi and Li, Ruilong and Fuji-Tsang, Clement and Fidler, Sanja and Sifakis, Eftychios and Museth, Ken},
+  journal={ACM Transactions on Graphics (TOG)},
+  volume={43},
+  number={4},
+  pages={133:1--133:15},
+  year={2024},
+  publisher={ACM New York, NY, USA}
+}
+```
+
+## Contact
@@ -0,0 +1,195 @@
+#!/bin/bash
+# Copyright Contributors to the OpenVDB Project
+# SPDX-License-Identifier: Apache-2.0
+
+usage() {
+  echo "Usage: $0 [-h|--help] [build_type] [options...]"
+  echo ""
+  echo "Builds or tests FVDB."
+  echo ""
+  echo "Arguments:"
+  echo "  build_type     Specifies the build operation. Can be one of:"
+  echo "                   install    - Build and install the package (default)."
+  echo "                   wheel      - Build the Python wheel."
+  echo "                   ctest      - Run tests (requires tests to be built)."
+  echo "                   debug      - Build in debug mode with full debug symbols and no optimizations."
+  echo ""
+  echo "Options:"
+  echo "  -h, --help     Display this help message and exit."
+  echo ""
+  echo "Build Modifiers (for 'install' and 'wheel' build types, typically passed after build_type):"
+  echo "  gtests         Enable building tests (sets FVDB_BUILD_TESTS=ON)."
+  echo "  benchmarks     Enable building benchmarks (sets FVDB_BUILD_BENCHMARKS=ON)."
+  echo "  verbose        Enable verbose build output for pip and CMake."
+  echo ""
+  echo "  Any modifier arguments not matching above are passed through to pip."
+  exit 0
+}
+
+setup_parallel_build_jobs() {
+  # Calculate the optimal number of parallel build jobs based on available RAM
+  RAM_GB=$(free -g | awk '/^Mem:/{print $4}')
+  if [ -z "$RAM_GB" ]; then
+      echo "Error: Unable to determine available RAM"
+      exit 1
+  fi
+  JOB_RAM_GB=2.5
+
+  # Calculate max jobs based on RAM
+  RAM_JOBS=$(awk -v ram="$RAM_GB" -v job_ram="$JOB_RAM_GB" 'BEGIN { print int(ram / job_ram) }')
+
+  # Get number of processors
+  NPROC=$(nproc)
+
+  # if CMAKE_BUILD_PARALLEL_LEVEL is set, use that
+  if [ -n "$CMAKE_BUILD_PARALLEL_LEVEL" ]; then
+    echo "Using CMAKE_BUILD_PARALLEL_LEVEL=$CMAKE_BUILD_PARALLEL_LEVEL"
+  else
+    # Determine the minimum of RAM-based jobs and processor count
+    if [ "$RAM_JOBS" -lt "$NPROC" ]; then
+      CMAKE_BUILD_PARALLEL_LEVEL="$RAM_JOBS"
+    else
+      CMAKE_BUILD_PARALLEL_LEVEL="$NPROC"
+    fi
+
+    # Ensure at least 1 job
+    if [ "$CMAKE_BUILD_PARALLEL_LEVEL" -lt 1 ]; then
+      CMAKE_BUILD_PARALLEL_LEVEL=1
+    fi
+
+    echo "Setting CMAKE_BUILD_PARALLEL_LEVEL to $CMAKE_BUILD_PARALLEL_LEVEL based on available RAM to target $JOB_RAM_GB GB per translation unit"
+    export CMAKE_BUILD_PARALLEL_LEVEL
+  fi
+}
+
+if [[ "$1" == "-h" || "$1" == "--help" ]]; then
+  usage
+fi
+
+# Determine BUILD_TYPE from the first positional argument, default to 'install'.
+# Handle shifting of arguments appropriately.
+_first_arg_val="$1"
+BUILD_TYPE="install" # Default build type
+
+if [[ -n "$_first_arg_val" ]]; then
+  if [[ "$_first_arg_val" == "install" || "$_first_arg_val" == "wheel" || "$_first_arg_val" == "ctest" ]]; then
+    BUILD_TYPE="$_first_arg_val"
+    shift # Consume the build_type argument
+  else
+    # _first_arg_val is not a recognized build type. Print usage and exit.
+    echo "Error: Argument '$_first_arg_val' is not a recognized build_type."
+    usage # This will also exit
+  fi
+fi
+
+CONFIG_SETTINGS=""
+PASS_THROUGH_ARGS=""
+
+while (( "$#" )); do
+  is_config_arg_handled=false
+  if [[ "$BUILD_TYPE" == "install" || "$BUILD_TYPE" == "wheel" ]]; then
+    if [[ "$1" == "gtests" ]]; then
+      echo "Detected 'gtests' flag for $BUILD_TYPE build. Enabling FVDB_BUILD_TESTS."
+      CONFIG_SETTINGS+=" --config-settings=cmake.define.FVDB_BUILD_TESTS=ON"
+      is_config_arg_handled=true
+    elif [[ "$1" == "benchmarks" ]]; then
+      echo "Detected 'benchmarks' flag for $BUILD_TYPE build. Enabling FVDB_BUILD_BENCHMARKS."
+      CONFIG_SETTINGS+=" --config-settings=cmake.define.FVDB_BUILD_BENCHMARKS=ON"
+      is_config_arg_handled=true
+    elif [[ "$1" == "verbose" ]]; then
+      echo "Enabling verbose build"
+      CONFIG_SETTINGS+=" -v -C build.verbose=true"
+      is_config_arg_handled=true
+    elif [[ "$1" == "debug" ]]; then
+      echo "Enabling debug build"
+      CONFIG_SETTINGS+=" --config-settings=cmake.build-type=debug"
+      is_config_arg_handled=true
+    fi
+  fi
+
+  if ! $is_config_arg_handled; then
+    # Append other arguments, handling potential spaces safely
+    PASS_THROUGH_ARGS+=" $(printf "%q" "$1")"
+  fi
+  shift
+done
+
+# Construct PIP_ARGS with potential CMake args and other pass-through args
+export PIP_ARGS="--no-build-isolation$CONFIG_SETTINGS$PASS_THROUGH_ARGS"
+
+if [ "$BUILD_TYPE" != "ctest" ]; then
+    setup_parallel_build_jobs
+fi
+
+# if the user specified 'wheel' as the build type, then we will build the wheel
+if [ "$BUILD_TYPE" == "wheel" ]; then
+    echo "Build wheel"
+    echo "pip wheel . --wheel-dir dist/ $PIP_ARGS"
+    pip wheel . --wheel-dir dist/ $PIP_ARGS
+elif [ "$BUILD_TYPE" == "install" ]; then
+    echo "Build and install package"
+    echo "pip install --force-reinstall $PIP_ARGS ."
+    pip install --force-reinstall $PIP_ARGS .
+# TODO: Fix editable install
+# else
+#     echo "Build and install editable package"
+#     echo "pip install $PIP_ARGS -e .  "
+#     pip install $PIP_ARGS -e .
+elif [ "$BUILD_TYPE" == "ctest" ]; then
+
+    # --- Ensure Test Data is Cached via CMake Configure Step ---
+    echo "Ensuring test data is available in CPM cache..."
+
+    if [ -z "$CPM_SOURCE_CACHE" ]; then
+         echo "CPM_SOURCE_CACHE is not set"
+    else
+        echo "Using CPM_SOURCE_CACHE: $CPM_SOURCE_CACHE"
+    fi
+
+    # Assume this script runs from the source root directory
+    SOURCE_DIR=$(pwd)
+    TEMP_BUILD_DIR="build_temp_download_data"
+
+    # Clean up previous temp dir and create anew
+    rm -rf "$TEMP_BUILD_DIR"
+    mkdir "$TEMP_BUILD_DIR"
+
+    echo "Running CMake configure in temporary directory ($TEMP_BUILD_DIR) to trigger data download..."
+    pushd "$TEMP_BUILD_DIR" > /dev/null
+    cmake "$SOURCE_DIR/src/cmake/download_test_data"
+    popd > /dev/null # Back to SOURCE_DIR
+
+    # Clean up temporary directory
+    rm -rf "$TEMP_BUILD_DIR"
+    echo "Test data caching step finished."
+    # --- End Test Data Caching ---
+
+    # --- Find and Run Tests ---
+    echo "Searching for test build directory..."
+    # Find the directory containing the compiled tests (adjust if needed)
+    # Using -print -quit to stop after the first match for efficiency
+    BUILD_DIR=$(find build -name tests -type d -print -quit)
+
+    if [ -z "$BUILD_DIR" ]; then
+        echo "Error: Could not find build directory with tests"
+        echo "Please enable tests by building with pip argument"
+        echo "-C cmake.define.FVDB_BUILD_TESTS=ON"
+        exit 1
+    fi
+    echo "Found test build directory: $BUILD_DIR"
+
+    # Run ctest within the test build directory
+    pushd "$BUILD_DIR" > /dev/null
+    echo "Running ctest..."
+    ctest --output-on-failure
+    CTEST_EXIT_CODE=$?
+    popd > /dev/null # Back to SOURCE_DIR
+
+    echo "ctest finished with exit code $CTEST_EXIT_CODE."
+    exit $CTEST_EXIT_CODE
+
+else
+    echo "Invalid build/run type: $BUILD_TYPE"
+    echo "Valid build/run types are: wheel, install, ctest"
+    exit 1
+fi
@@ -0,0 +1,25 @@
+version: '3'
+services:
+  fvdb-dev:
+    build:
+      context: ../
+      dockerfile: fvdb/Dockerfile
+      args:
+        MODE: dev
+    volumes:
+      - ../:/openvdb
+    # Uncomment and adjust the following as needed
+    # ports:
+    #   - "8888:8888"  # Example for Jupyter notebook
+    entrypoint: /bin/bash
+    stdin_open: true
+    tty: true
+    environment:
+      - NVIDIA_VISIBLE_DEVICES=all
+    deploy:
+      resources:
+        reservations:
+          devices:
+            - driver: nvidia
+              count: all
+              capabilities: [gpu]
@@ -0,0 +1,5 @@
+GridBatch
+===============
+
+.. autoclass:: fvdb.GridBatch
+   :members:
@@ -0,0 +1,5 @@
+JaggedTensor
+===============
+
+.. autoclass:: fvdb.JaggedTensor
+   :members:
@@ -0,0 +1,20 @@
+fvdb.nn
+=======
+
+`fvdb.nn` is a collection of neural network layers to build sparse neural networks.
+
+.. autoclass:: fvdb.nn.VDBTensor
+.. autoclass:: fvdb.nn.MaxPool
+.. autoclass:: fvdb.nn.AvgPool
+.. autoclass:: fvdb.nn.UpsamplingNearest
+.. autoclass:: fvdb.nn.FillFromGrid
+.. autoclass:: fvdb.nn.SparseConv3d
+.. autoclass:: fvdb.nn.GroupNorm
+.. autoclass:: fvdb.nn.Linear
+.. autoclass:: fvdb.nn.ReLU
+.. autoclass:: fvdb.nn.LeakyReLU
+.. autoclass:: fvdb.nn.SELU
+.. autoclass:: fvdb.nn.SiLU
+.. autoclass:: fvdb.nn.Tanh
+.. autoclass:: fvdb.nn.Sigmoid
+.. autoclass:: fvdb.nn.Dropout
@@ -0,0 +1,4 @@
+fvdb.utils
+==================
+
+.. autofunction:: fvdb.utils.build_ext.FVDBExtension
@@ -0,0 +1,91 @@
+# Copyright Contributors to the OpenVDB Project
+# SPDX-License-Identifier: Apache-2.0
+#
+# -- Path setup --------------------------------------------------------------
+
+# If extensions (or modules to document with autodoc) are in another directory,
+# add these directories to sys.path here. If the directory is relative to the
+# documentation root, use os.path.abspath to make it absolute, like shown here.
+
+import os
+import sys
+
+sys.path.insert(0, os.path.abspath(".."))
+
+
+# -- Project information -----------------------------------------------------
+
+project = "fVDB"
+copyright = "Contributors to the OpenVDB Project"
+author = "Contributors to the OpenVDB Project"
+
+
+# -- General configuration ---------------------------------------------------
+
+# Add any Sphinx extension module names here, as strings. They can be
+# extensions coming with Sphinx (named 'sphinx.ext.*') or your custom
+# ones.
+extensions = ["sphinx.ext.autodoc", "sphinx.ext.viewcode", "sphinx.ext.napoleon", "myst_parser"]
+
+myst_enable_extensions = [
+    "amsmath",
+    "attrs_inline",
+    "colon_fence",
+    "deflist",
+    "dollarmath",
+    "fieldlist",
+    "html_admonition",
+    "html_image",
+    "linkify",
+    "replacements",
+    "smartquotes",
+    "strikethrough",
+    "substitution",
+    "tasklist",
+]
+
+# Fix return-type in google-style docstrings
+napoleon_custom_sections = [("Returns", "params_style")]
+
+# Add any paths that contain templates here, relative to this directory.
+templates_path = ["_templates"]
+
+source_suffix = [".rst", ".md"]
+
+# List of patterns, relative to source directory, that match files and
+# directories to ignore when looking for source files.
+# This pattern also affects html_static_path and html_extra_path.
+exclude_patterns = ["_build", "Thumbs.db", ".DS_Store"]
+
+autodoc_default_options = {"undoc-members": "forward, extra_repr"}
+
+# -- Options for HTML output -------------------------------------------------
+
+# The theme to use for HTML and HTML Help pages.  See the documentation for
+# a list of builtin themes.
+#
+html_theme = "sphinx_rtd_theme"
+
+# Add any paths that contain custom static files (such as style sheets) here,
+# relative to this directory. They are copied after the builtin static files,
+# so a file named "default.css" will overwrite the builtin "default.css".
+# html_static_path = ['imgs']
+
+
+# -- Custom hooks ------------------------------------------------------------
+
+
+def process_signature(app, what, name, obj, options, signature, return_annotation):
+    if signature is not None:
+        signature = signature.replace("._Cpp", "")
+        signature = signature.replace("fvdb::", "fvdb.")
+
+    if return_annotation is not None:
+        return_annotation = return_annotation.replace("._Cpp", "")
+        return_annotation = return_annotation.replace("fvdb::", "fvdb.")
+
+    return signature, return_annotation
+
+
+def setup(app):
+    app.connect("autodoc-process-signature", process_signature)
@@ -0,0 +1,55 @@
+Welcome to fVDB!
+=================
+
+fVDB, inspired by function notation to resemble :math:`f(VDB)`, is a data structure for encoding and operating on *sparse* voxel hierarchies of features in PyTorch.
+A sparse voxel hierarchy is a coarse-to-fine hierarchy of sparse voxel grids such that every fine voxel is contained within some coarse voxel.
+fvdb supports storing PyTorch tensors at the corners and centers of voxels in a hierarchy and enables a number of differentiable operations on these tensors (e.g. trilinear interpolation, splatting, ray tracing).
+
+.. image:: imgs/fvdb_teaser.png
+   :align: center
+   :width: 400
+
+Please refer to the tutorials for examples of how to install and use fVDB to build your own sparse voxel pipelines.
+
+.. toctree::
+   :caption: Introduction
+   :hidden:
+
+   self
+
+.. toctree::
+   :maxdepth: 1
+   :caption: Tutorials
+
+   tutorials/installation
+   tutorials/basic_concepts
+   tutorials/building_grids
+   tutorials/basic_grid_ops
+   tutorials/ray_tracing
+   tutorials/simple_unet
+   tutorials/io
+   tutorials/volume_rendering
+
+.. toctree::
+   :maxdepth: 1
+   :caption: API References
+
+   api/grid_batch
+   api/jagged_tensor
+
+.. toctree::
+   :maxdepth: 2
+
+   api/nn
+   api/utils
+
+.. raw:: html
+
+   <hr>
+
+Indices and tables
+==================
+
+* :ref:`genindex`
+* :ref:`modindex`
+* :ref:`search`
@@ -0,0 +1,60 @@
+# Clangd for "Intellisense" in fVDB
+
+In fVDB we typically develop using Visual Studio Code (VSCode), both local, in containers, and
+remove via SSH. To aid development we use `clangd` for intellisense (code navigation) and to
+detect syntax and usage errors before compiling. With `clangd` active, you will see "red squiggles"
+under actual syntax / compilation errors in your code as you edit; you will get code completion
+suggestions; and you will get tooltips for definitions when you mouse over code.
+
+The fVDB VSCode workspace includes configuration for clangd but developers new to fVDB should follow
+these instructions.
+
+## Setting up clangd in your environment
+
+The fvdb Conda development environment installs the CUDA toolkit, but clangd runs outside the
+environment and so may not find the correct CUDA installation. Therefore you should create the file
+
+To fix this, you need to create a user clangd config file to point to the CUDA installation in the
+FVDB conda environment: Create ~/.config/clangd/config.yaml with these contents. This assumes you
+use the `dev_environment.yml` to create the `fvdb` Conda environment.
+
+```yaml
+# https://clangd.llvm.org/config
+
+# Apply a config conditionally to all CUDA files
+If:
+  PathMatch: .*\.cuh?
+CompileFlags:
+  Add:
+    - "--cuda-path=/home/mharris/miniforge3/envs/fvdb"
+```
+
+The first time you load the fVDB VSCode workspace, it should prompt you to install recommended
+extensions. Specifically: "llvm-vs-code-extensions.vscode-clangd", and
+"ms-vscode.cpptools-extension-pack", as defined in the `extensions.json` file. Install these
+extensions. With Clangd, the MS CPPTools is only used for syntax highlighting. Clangd is used for
+intellisense and (CUDA) C++ file formatting.
+
+## Troubleshooting Intellisense
+
+Clangd uses the file `compile_commands.json` to determine the compiler command line used so that it
+can build an index of the source code for "intellisense" navigation. This file is generated by
+CMake when you configure the build. This happens when you run fVDB's `build.sh`. Here are some
+tips for troubleshooting intellisense issues.
+
+- If you get a red squiggle on the first character of a file, it’s usually a major failure. For
+  example when clangd fails to find the CUDA installation. This is often a configuration problem.
+- A common problem is when you are editing a file that is not in `compile_commands.json`. if this is
+  the case, you may not have CMake-configured your build.
+    - A key indicator of this is seeing a message about ‘inferred’ commands like this in the clangd
+    output log `ASTWorker building file XXX with command inferred from YYY`
+- Switching conda environments: if you activate a conda environment, e.g. `fvdb_build` and build
+  (or cmake configure), then switch to another conda environment, your `compile_commands.json` will
+  be out of date with the wrong include paths. You need to rerun cmake configure or rebuild in the
+  active conda environment.
+- red squiggles in google test .cpp : did you build the tests? Are they in compile_commands.cpp?
+  Make sure you add the `gtests` argument to your `build.sh` command line:
+
+```sh
+./build.sh install gtests
+```
@@ -0,0 +1,51 @@
+# ƒVDB Docker Help
+
+Running a docker container is a great way to ensure that you have a consistent environment for building and running ƒVDB.
+
+Our provided [`Dockerfile`](../../Dockerfile) constructs a Docker image which is ready to build ƒVDB.  The docker image is configured to install miniforge and the `fvdb` conda environment with all the dependencies needed to build and run ƒVDB.
+
+## Setting up a Docker Container
+
+Building and starting the docker image is done by running the following command from the fvdb directory:
+```shell
+docker compose run --rm fvdb-dev
+```
+
+When you are ready to build ƒVDB, run the following command within the docker container.  `TORCH_CUDA_ARCH_LIST` specifies which CUDA architectures to build for.
+```shell
+conda activate fvdb;
+cd /openvdb/fvdb;
+TORCH_CUDA_ARCH_LIST="7.0;7.5;8.0;8.6+PTX" \
+./build.sh install
+```
+
+### Troubleshooting
+
+* **docker daemon runtime/driver errors**
+Errors like these:
+```shell
+docker: Error response from daemon: could not select device driver "" with capabilities: [[gpu]].
+```
+```shell
+docker: Error response from daemon: Unknown runtime specified nvidia.
+```
+most likely indicate that the [`NVIDIA Container Toolkit`](https://docs.nvidia.com/datacenter/cloud-native/container-toolkit/latest/index.html) is not installed.
+You can install it by following the [installation instructions here](https://docs.nvidia.com/datacenter/cloud-native/container-toolkit/latest/install-guide.html)
+
+* **GPU access blocked by the operating system**
+Errors like these:
+```shell
+Failed to initialize NVML: GPU access blocked by the operating system
+Failed to properly shut down NVML: GPU access blocked by the operating system
+```
+may be solved by making a change to the file `/etc/nvidia-container-toolkit/config.toml`, setting
+the property "no-cgroups" from "true" to "false":
+
+```
+# ...
+[nvidia-container-cli]
+# ...
+# Below changed from previous no-cgroups = true
+no-cgroups = false
+# ...
+```
@@ -0,0 +1,87 @@
+# Basic Concepts
+
+We give a high-level overview of the main concepts in fVDB. Namely, how fVDB encodes sparse voxel grids with attributes, as well as how fVDB efficiently manages batches with non-uniform numbers of elements.
+
+## `GridBatch`: Voxel Grids with Attributes
+
+At its core, fVDB provides means to efficiently encode mini-batches of sparse voxel grids where each voxel contains arbitrary vector or scalar attributes encoded as torch tensors. In practice, such an encoding is implemented via the `GridBatch` class.
+
+![Minibatch2.png](../imgs/fig/Minibatch2.png)
+
+A `GridBatch` is an indexing structure which maps 3D `ijk` coordinates to integer offsets which can be used to look up attributes in a tensor. The figure below illustrates this process for a GridBatch containing a single grid.
+
+![gridbatch.png](../imgs/fig/gridbatch.png)
+
+In the figure, the `GridBatch` acts as an acceleration structure which encodes the sparsity pattern of the grid (also known as the **********topology**********) and can translate grid coordinates into offsets into a data tensor.
+
+By separating the grid topology and data tensors, the same grid can be used to operate on many different attributes without rebuilding.
+
+Every operation in fVDB is built upon this kind of query (e.g. Sparse Convolution uses this query to look up features in the neighborhood of each voxel).
+
+## `JaggedTensor` and Batching
+
+Each grid in a `GridBatch` can have a different number of voxels (****e.g.**** in the mini batch of four cars above, each car has a different number of voxels). This means that unlike the dense case, fVDB needs to handle parallel operations over ***jagged batches***. I.e. batches containing different numbers of elements.
+
+To handle jagged batches, fVDB provides a `JaggedTensor` class. Conceptually, a `JaggedTensor` is a list of tensors with shapes $[N_0, *], [N_1, *], \ldots, [N_{B-1}, *]$ where $B$ is the number of elements in the batch, $N_i$ is the number of elements in the $i^\text{th}$ batch item and $*$ is an arbitrary numer of additional dimensions that all match between the tensors. The figure below illustrates such a list of tensors pictorially.
+
+![jaggedtensor1.png](../imgs/fig/jaggedtensor1.png)
+
+In practice, `JaggedTensor`s are represented in memory by concatenating each tensor in the list into a single `jdata` (for Jagged Data) tensor of shape $[N_0 + N_1 + \ldots + N_{B-1}, *]$. Additionally, each `JaggedTensor` stores an additional `jidx` tensor (for Jagged Indexes) of shape $[N_0 + N_1 + \ldots + N_{B-1}]$ containing one int per element in the jagged tensor. `jidx[i]` is the batch index of the $i^\text{th}$ element of `jdata`. Finally, a `JaggedTensor` contains a `joffsets` tensor (for Jagged Offsets) of shape $[B, 2]$ which indicates the start and end positions of the $i^\text{th}$ tensor in the batch.
+
+![jaggedtensor4.png](../imgs/fig/jaggedtensor4.png)
+
+Similarly, each `GridBatch` also has `jidx` and `joffsets` corresponding to the batch index of each voxel in the grid, and the start and end offsets of each voxel index in a batch.
+
+## A simple example
+
+To illustrate the use of `GridBatch`and `JaggedTensor`, consider a simple example where we build a grid from a point cloud, splat some values onto the voxels of that grid, and then sample them again using a different set of points.
+
+First, we construct a minibatch of grids using the input points. These input points have corresponding color attributes.
+
+```python
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+import torch
+import point_cloud_utils as pcu
+
+# We're going to create a minibatch of two point clouds each of which
+# has a different number of points
+pts1, clrs1 = load_car_1_mesh(mode = "vn")
+pts2, clrs2 = load_car_2_mesh(mode = "vn")
+
+# Creating JaggedTensors: one for points and one for colors
+points = fvdb.JaggedTensor([pts1, pts2])
+colors = fvdb.JaggedTensor([clrs1, clrs2])
+
+# Create a grid where the voxels each have unit sidelength
+grid = fvdb.gridbatch_from_points(points, voxel_sizes=1.0)
+
+# Indexing into a JaggedTensor returns a JaggedTensor
+print(points[0].jdata.shape)
+print(points[1].jdata.shape)
+```
+
+![Minibatch of grids constructed from the input points. These input points have corresponding color attributes.](../imgs/fig/screenshot_000000.png.trim.png)
+
+Next, we splat the colors at the points to the constructed grid, yielding per-voxel colors.
+
+```python continuation
+# Splat the colors into the grid with trilinear interpolation
+# vox_colors is a JaggedTensor of per-voxel normas
+vox_colors = grid.splat_trilinear(points, colors)
+```
+
+![Colors splat at the input points to grid, yielding per-voxel colors.](../imgs/fig/screenshot_000006.png.trim.png)
+
+Finally, we generate a new set of noisy points and sample the grid to recover colors at those new samples.
+
+```python continuation
+# Now let's generate some random points and sample the grid at those points
+sample_points = fvdb.JaggedTensor([torch.rand(10_000, 3), torch.rand(11_000, 3)]).cuda()
+
+# sampled_colors is a JaggedTensor with the same shape as sample_points with
+# one color sampled from the grid at each point
+sampled_colors = grid.sample_trilinear(sample_points, vox_colors)
+```
+
+![Colors resampled at random locations from the grid.](../imgs/fig/screenshot_000004.png.trim.png)
@@ -0,0 +1,868 @@
+# Basic GridBatch Operations
+
+Here we describe basic operations you can perform on a `GridBatch`. Generally, these operations involve a `GridBatch` and an associated `JaggedTensor` of data representing the attributes/features at each voxel within the grid.
+
+## Sampling grids
+
+You can differentiably sample data stored at the voxels of a grid using *trilinear* or *Bézier* interpolation as follows:
+
+```python
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+import torch
+import numpy as np
+
+# We're going to create a minibatch of two meshes
+v1, f1 = load_car_1_mesh(mode="vf")
+v2, f2 = load_car_2_mesh(mode="vf")
+
+# Build a GridBatch from two meshes
+mesh_v_jagged = fvdb.JaggedTensor([v1, v2])
+mesh_f_jagged = fvdb.JaggedTensor([f1, f2]).int()
+grid = fvdb.gridbatch_from_mesh(mesh_v_jagged, mesh_f_jagged, voxel_sizes=0.1)
+
+# Generate some sample points by adding random gaussian noise to the center of each voxel
+world_space_centers = grid.grid_to_world(grid.ijk.float())
+# 5 samples per voxel for the first grid and 7 samples per voxel for the second
+sample_pts = fvdb.JaggedTensor([
+    torch.cat([world_space_centers[0].jdata] * 5),
+    torch.cat([world_space_centers[1].jdata] * 7)])
+sample_pts += fvdb.JaggedTensor([
+    torch.randn(grid.num_voxels_at(0) * 5, 3).to(grid.device) * grid.voxel_sizes[0]*0.3,
+    torch.randn(grid.num_voxels_at(1) * 7, 3).to(grid.device) * grid.voxel_sizes[1]*0.3])
+
+# Generate RGB values per voxel as the normalized absolute grid coordinate of the voxel
+per_voxel_colors = world_space_centers.clone()
+per_voxel_colors.jdata = torch.abs(world_space_centers.jdata) / torch.norm(world_space_centers.jdata, dim=-1, keepdim=True)
+
+# Sample these RGB colors at each sample point with trilinear interpolation
+# NOTE: You can use grid.sample_bezier to sample using bezier interpolation
+sampled_colors = grid.sample_trilinear(sample_pts, per_voxel_colors)
+```
+
+Grid with color attributes            |  Points with sampled colors
+:------------------------------------:|:------------------------------------:
+![](../imgs/fig/sampling_1.png)       |  ![](../imgs/fig/sampling_2.png)
+
+## Splatting point data to a grid
+
+You can differentiably splat data at a set of points into voxels in a grid using *trilinear* or *Bézier* interpolation as follows:
+
+```python
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+import torch
+import point_cloud_utils as pcu
+
+# We're going to create a minibatch of two point clouds each of which
+# has a different number of points
+pts1, clrs1 = load_car_1_mesh(mode="vn")
+pts2, clrs2 = load_car_2_mesh(mode="vn")
+
+# JaggedTensors of points and normals
+points = fvdb.JaggedTensor([pts1, pts2])
+colors = fvdb.JaggedTensor([clrs1, clrs2])
+
+# Create a grid where the voxels each have unit sidelength
+grid = fvdb.gridbatch_from_points(points, voxel_sizes=1.0)
+
+# Splat the normals into the grid with trilinear interpolation
+# vox_normals is a JaggedTensor of per-voxel normas
+# NOTE: You can use grid.splat_bezier to splat using bezier interpolation
+vox_colors = grid.splat_trilinear(points, colors)
+```
+
+Input grid and points with colors                     |  Splat colors onto the grid
+:----------------------------------------------------:|:----------------------------------------------------:
+![](../imgs/fig/screenshot_000000.png.trim.png)       |  ![](../imgs/fig/screenshot_000006.png.trim.png)
+
+## Checking if points are in a grid
+
+You can query whether points lie in a grid as follows:
+
+```python
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+import torch
+import numpy as np
+import point_cloud_utils as pcu
+
+# We're going to create a minibatch of two meshes
+v1, f1 = load_car_1_mesh(mode="vf")
+v2, f2 = load_car_2_mesh(mode="vf")
+f1, f2 = f1.to(torch.int32), f2.to(torch.int32)
+
+# Build a GridBatch from two meshes
+mesh_v_jagged = fvdb.JaggedTensor([v1, v2]).cuda()
+mesh_f_jagged = fvdb.JaggedTensor([f1, f2]).cuda()
+grid = fvdb.gridbatch_from_mesh(mesh_v_jagged, mesh_f_jagged, voxel_sizes=0.1)
+
+# Generate some points and check if they lie within the grid
+bbox_sizes = grid.bbox[:, 1] - grid.bbox[:, 0]
+bbox_origins = grid.bbox[:, 0]
+pts = fvdb.JaggedTensor([
+    (torch.randn(10_000, 3, device='cuda') - bbox_origins[0]) * bbox_sizes[0],
+    (torch.randn(11_000, 3, device='cuda') - bbox_origins[0]) * bbox_sizes[0],
+])
+
+# Get a mask indicating which points lie in the grid
+mask = grid.points_in_active_voxel(pts)
+```
+
+We visualize the points which intersect the grid (yellow points intersect and purple points do not).
+![](../imgs/fig/pts_in_grid.png)
+
+## Checking if ijk coordinates are in a grid
+
+Similar to querying whether world space points lie in a grid, you can query whether the index space `ijk` integer coordinates lie in a grid as follows:
+
+```python
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+import torch
+import numpy as np
+import point_cloud_utils as pcu
+import polyscope as ps
+import os
+
+
+# We're going to create a minibatch of two meshes
+v1, f1 = load_car_1_mesh(mode="vf")
+v2, f2 = load_car_2_mesh(mode="vf")
+f1, f2 = f1.to(torch.int32), f2.to(torch.int32)
+
+# Build a GridBatch from two meshes
+mesh_v_jagged = fvdb.JaggedTensor([v1, v2]).cuda()
+mesh_f_jagged = fvdb.JaggedTensor([f1, f2]).cuda()
+grid = fvdb.gridbatch_from_mesh(mesh_v_jagged, mesh_f_jagged, voxel_sizes=0.025)
+
+rand_idx_pts = []
+for b, b_grid in enumerate(grid.bbox):
+    pts = [torch.randint(int(b_grid[0][i]), int(b_grid[1][i]), size=(2_000 * (b+1),), device='cuda') for i in range(3)]
+    rand_idx_pts.append(torch.stack(pts, dim=1))
+
+pts = fvdb.JaggedTensor(rand_idx_pts)
+
+coords_in_grid = grid.coords_in_active_voxel(pts)
+```
+
+We visualize the coordinates which intersect the grid (yellow coordinates intersect and purple coordinates do not).
+![](../imgs/fig/coords_in_grid.png)
+
+
+## Checking if axis aligned cubes intersect a grid
+
+There are methods of a grid to help ascertain whether a provided set of axis-aligned cubes (all of the same size) each are contained within voxels of the grid or intersect with any voxel of the grid.  These methods are called `cubes_in_grid` and `cubes_intersect_grid`, respectively, and return a `JaggedTensor` of boolean values which indicate the result.
+
+In this example we create some random points to represent the centers of cubes of size `0.03` units along-a-side. Then we use these two methods to discover which of those cubes would intersect a voxel of the grid and which would be entirely enclosed by a voxel in the grid.
+
+```python
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+import torch
+import numpy as np
+import point_cloud_utils as pcu
+import polyscope as ps
+import os
+
+# We're going to create a minibatch of two meshes
+v1, f1 = load_car_1_mesh(mode="vf")
+v2, f2 = load_car_2_mesh(mode="vf")
+f1, f2 = f1.to(torch.int32), f2.to(torch.int32)
+
+# Build a GridBatch from two meshes
+mesh_v_jagged = fvdb.JaggedTensor([v1, v2]).cuda()
+mesh_f_jagged = fvdb.JaggedTensor([f1, f2]).cuda()
+grid = fvdb.gridbatch_from_mesh(mesh_v_jagged, mesh_f_jagged, voxel_sizes=0.025)
+
+# Generate some points and check if they lie within the grid
+bbox_sizes = (grid.bbox[:, 1] - grid.bbox[:, 0]) * grid.voxel_sizes
+bbox_origins = (grid.bbox[:, 0] + grid.origins) * grid.voxel_sizes
+pts = fvdb.JaggedTensor(
+    [
+        (torch.rand(3000, 3, device="cuda")) * bbox_sizes[0] + bbox_origins[0],
+        (torch.rand(2000, 3, device="cuda")) * bbox_sizes[1] + bbox_origins[1],
+    ]
+)
+
+cube_size = 0.03
+
+# We can check if the axis-aligned cubes intersect any voxels of the grid...
+cubes_intersect_grid = grid.cubes_intersect_grid(pts, -cube_size / 2, cube_size / 2)
+# ... or if they are fully contained within the grid
+cubes_in_grid = grid.cubes_in_grid(pts, -cube_size / 2, cube_size / 2)
+```
+
+We visualize the cubes which intersect a voxel in the grid (yellow cubes intersect and purple cubes do not).
+![](../imgs/fig/cubes_intersect_grid.png)
+
+
+And which cubes lie entirely within a voxel of the grid
+![](../imgs/fig/cubes_inside_grid.png)
+
+
+## Converting ijk values to indexes
+
+There is a convenient method to convert ijk values, which are integer coordinates in the index space of a grid, to indexes, which are the linearized indices of the voxels in the GridBatch.  This method, `ijk_to_index`, returns a `JaggedTensor` of indexes.  If the provided ijk values do not correspond to any voxels in the grid, the returned index will be `-1`.
+
+```python
+
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+import torch
+import numpy as np
+import point_cloud_utils as pcu
+import polyscope as ps
+import os
+
+torch.random.manual_seed(0)
+
+def generate_random_points(bounding_box, num_points):
+    min_i, min_j, min_k = bounding_box[0]
+    max_i, max_j, max_k = bounding_box[1]
+
+    # Generate random integer points within the bounding box
+    random_i = torch.randint(min_i, max_i, size=(num_points,))
+    random_j = torch.randint(min_j, max_j, size=(num_points,))
+    random_k = torch.randint(min_k, max_k, size=(num_points,))
+
+    random_points = torch.stack([random_i, random_j, random_k], dim=1)
+
+    return random_points
+
+# We're going to create a minibatch of two meshes
+v1, f1 = load_car_1_mesh(mode="vf")
+v2, f2 = load_car_1_mesh(mode="vf")
+f1, f2 = f1.to(torch.int32), f2.to(torch.int32)
+
+# Build a GridBatch from two meshes
+mesh_v_jagged = fvdb.JaggedTensor([v1, v2]).cuda()
+mesh_f_jagged = fvdb.JaggedTensor([f1, f2]).cuda()
+grid = fvdb.gridbatch_from_mesh(mesh_v_jagged, mesh_f_jagged, voxel_sizes=0.025)
+
+rand_pts = fvdb.JaggedTensor([generate_random_points(bbox, 1000) for bbox in grid.bbox]).cuda()
+
+rand_pts_indices= grid.ijk_to_index(rand_pts)
+
+print(rand_pts_indices.jdata)
+```
+```bash
+tensor([  -1,   -1,  121,  ..., 4614, 5695,   -1], device='cuda:0')
+```
+
+## Converting indexes to ijk values
+
+If we have the value of an index into the feature data and want to obtain its corresponding ijk value, this is as simple as indexing into the `ijk` attribute of the grid which is itself just a `JaggedTensor` representing the `ijk` values of each index.  Here we get the `ijk` values for 1000 random indexes of a `GridBatch`.
+
+```python
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+import torch
+import numpy as np
+import point_cloud_utils as pcu
+import polyscope as ps
+import os
+
+# We're going to create a minibatch of two meshes
+v1, f1 = load_car_1_mesh(mode="vf")
+v2, f2 = load_car_2_mesh(mode="vf")
+f1, f2 = f1.to(torch.int32), f2.to(torch.int32)
+
+# Build a GridBatch from two meshes
+mesh_v_jagged = fvdb.JaggedTensor([v1, v2]).cuda()
+mesh_f_jagged = fvdb.JaggedTensor([f1, f2]).cuda()
+grid = fvdb.gridbatch_from_mesh(mesh_v_jagged, mesh_f_jagged, voxel_sizes=0.025)
+
+rand_indexes = torch.randint(0, grid.total_voxels, size=(1000,)).cuda()
+print(grid.ijk.jdata[rand_indexes])
+print(grid.ijk.jidx[rand_indexes])
+```
+```bash
+tensor([[-11,   5,   7],
+        [  8,  -4,  -5],
+        [-15,   3,  -2],
+        ...,
+        [ 18,   1,   6],
+        [ 19,  -3,   6],
+        [-14,  -7,   8]], device='cuda:0', dtype=torch.int32)
+```
+
+However, you may also want to obtain the batch index of the grid that the `ijk` coordinate belongs to.  This is just as simple by using the `jidx` attribute of the `JaggedTensor` which represents the batch index of each element in the `JaggedTensor`.  Using the same random indexes as above, we can get the batch index of the grid for each random index.
+
+```python continuation
+print(grid.ijk.jidx[rand_indexes])
+```
+```bash
+tensor([0, 1, 1, 0, 1, 1, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 0, 1, 1, 0, 1, 1,
+        1, 1, 0, 1, 1, 0, 1, 0, 0, 1, 1, 0, 1, 1, 1, 0, 1, 1, 0, 1, 1, 0, 1, 1,
+        0, 0, 1, 0, 0, 1, 1, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 1, 1, 0, 0, 1, 1, 1,
+        1, 0, 0, 0, 0, 1, 0, 1, 1, 0,…], device='cuda:0')
+```
+
+While this case of having a random set of `ijk` values is very contrived, it is more likely that we would have a `JaggedTensor` of features that correspond to each `ijk` value but are out-of-order in relation to another `grid.ijk`'s and we need to reorder these features to for this other grid.  For this, there is the `ijk_to_inv_index` function.
+
+Given a `JaggedTensor` of `ijk` values, `grid_batch.ijk_to_inv_index` will return a scalar integer `JaggedTensor` of size `[B, -1]`, where `B` is the number of grids in the batch and -1 represents the number of voxels in each grid.  This `JaggedTensor` of indexes can be used to permute the input to `ijk_to_inv_index` (or the feature `JaggedTensor` that correspond to those `ijk`'s) to match the ordering of the `grid_batch`.
+
+To concisely illustrate the properties of this function,
+* if `idx = grid.ijk_to_inv_index(misordered_ijk)`, then `grid.ijk == misordered_ijk[idx]`
+* if `idx = grid.ijk_to_index(misordered_ijk)`, then `grid.ijk[idx] == misordered_ijk`
+
+Note:  If any `ijk` values in the grid are not present in the input to `ijk_to_inv_index`, the returned index will be `-1` at that position.
+
+In this example, let's illustrate the more useful case where we have corresponding features and `ijk` values from a grid that are ordered differently from another reference grid and we want to re-order the features to match the order they should be in the reference grid.
+
+```python continuation
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+import torch
+import numpy as np
+import point_cloud_utils as pcu
+
+# We're going to create a minibatch of two meshes
+v1, f1 = load_car_1_mesh(mode="vf")
+v2, f2 = load_car_2_mesh(mode="vf")
+f1, f2 = f1.to(torch.int32), f2.to(torch.int32)
+
+# Build a GridBatch from two meshes
+mesh_v_jagged = fvdb.JaggedTensor([v1, v2]).cuda()
+mesh_f_jagged = fvdb.JaggedTensor([f1, f2]).cuda()
+reference_grid = fvdb.gridbatch_from_mesh(mesh_v_jagged, mesh_f_jagged, voxel_sizes=0.025)
+
+# 7 random feature values for each voxel in the grid
+features = reference_grid.jagged_like(torch.rand(grid.total_voxels, 7).to(grid.device))
+
+# Create a set of randomly shuffled corresponding ijk/features from our original grid/features
+shuffled_ijks = []
+shuffled_features = []
+
+for i in range(reference_grid.grid_count):
+    perm = torch.randperm(reference_grid.num_voxels_at(i))
+    shuffled_ijks.append(reference_grid.ijk.jdata[grid.ijk.jidx==i][perm])
+    shuffled_features.append(features.jdata[features.jidx==i][perm])
+
+shuffled_ijks = fvdb.JaggedTensor(shuffled_ijks)
+shuffled_features = fvdb.JaggedTensor(shuffled_features)
+
+# Get the indexes to reorder the shuffled features to match the grid's original ijk ordering
+idx = reference_grid.ijk_to_inv_index(shuffled_ijks)
+
+# Permute the shuffled features based on the ordering from `ijk_to_inv_index`
+unshuffled_features = shuffled_features.jdata[idx.jdata]
+print("Do the shuffled features that have been permuted based on `ijk_to_inv_index` match the original features? ", "Yes!" if torch.all(unshuffled_features == features.jdata) else "No!")
+```
+
+```bash
+Do the shuffled features that have been permuted based on `ijk_to_inv_index` match the original features?  Yes!
+```
+
+
+## Getting indexes of neighbors
+
+If we want to get the indexes of all the spatial neighbors of a set of ijk values, we can use the `neighbor_indexes` method of a `GridBatch`.  This method receives a set of `ijk` values and an `extent` (the number of voxels away from the ijk value to consider neighbors) and returns a `JaggedTensor` of indexes of neighbors for each ijk value.  The returned `JaggedTensor` has `jdata` of size `[N, extent*2+1, extent*2+1, extent*2+1]` where `N` is the number of requested ijk values.
+
+In this example we create a set of 24 random ijk values per grid and get the indexes of all neighbors 2 voxels away from each ijk.
+
+```python
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+import torch
+import numpy as np
+import point_cloud_utils as pcu
+import polyscope as ps
+import os
+
+# We're going to create a minibatch of two meshes
+v1, f1 = load_car_1_mesh(mode="vf")
+v2, f2 = load_car_2_mesh(mode="vf")
+f1, f2 = f1.to(torch.int32), f2.to(torch.int32)
+
+# Build a GridBatch from two meshes
+mesh_v_jagged = fvdb.JaggedTensor([v1, v2]).cuda()
+mesh_f_jagged = fvdb.JaggedTensor([f1, f2]).cuda()
+grid = fvdb.gridbatch_from_mesh(mesh_v_jagged, mesh_f_jagged, voxel_sizes=0.025)
+
+# Build a set of 24 randomly selected ijk values per grid
+rand_ijks = fvdb.JaggedTensor(
+    [
+        grid.ijk.jdata[
+            torch.randint(int(grid.ijk.joffsets[b]), int(grid.ijk.joffsets[b + 1]), (24,))
+        ]
+        for b in range(grid.grid_count - 1)
+    ],
+)
+
+# Get the indexes of all neighbors 2 voxels away from each ijk
+# Returns a JaggedTensor with jdata of size [48, 5, 5, 5] ([ N x (extent*2+1)^3 ])
+#   where each [5, 5, 5] describes the layout of neighbouring indexes of each ijk
+neighbor_idxs = grid.neighbor_indexes(rand_ijks, 2)
+```
+
+We visualize the voxels we selected in red and their neighbors in blue.
+![](../imgs/fig/neighbor_indexes.png)
+
+
+## Clipping a grid to a bounding box
+
+If we want to 'clip' a grid, meaning remove all the voxels that are outside of a bounding box, we can use the `clipped_grid` method of a `GridBatch`.  The minimum and maximum `ijk` extents of the bounding box used for this clipping are provided as arguments and can be specified per-grid of the batch.
+
+If we want to both clip the batch of grids and the accompanying `JaggedTensor` of batched features, we can use the `clip` method of a `GridBatch` which can be provided the `features` to be clipped along with the grid.
+
+In this example we clip a batch of grids to independent bounding boxes and visualize the result.
+
+```python
+import os
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh, load_car_3_mesh, load_car_4_mesh
+import torch
+import numpy as np
+import point_cloud_utils as pcu
+
+device = 'cuda'
+mesh_vs = []
+mesh_fs = []
+mesh_load_funcs = [load_car_1_mesh, load_car_2_mesh, load_car_3_mesh, load_car_4_mesh]
+
+for func in mesh_load_funcs:
+    v, f = func(mode="vf")
+    mesh_vs.append(v)
+    mesh_fs.append(f.to(torch.int32))
+
+mesh_vs = fvdb.JaggedTensor(mesh_vs)
+mesh_fs = fvdb.JaggedTensor(mesh_fs)
+
+vox_size = 0.01
+
+# Build GridBatch from meshes
+grid = fvdb.gridbatch_from_mesh(mesh_vs, mesh_fs, vox_size)
+
+# Use `clipped_grid` to clip the grids outside of the specified minimum and maximum bounding boxes for each grid…
+clipped_grid = grid.clipped_grid(
+                         ijk_min=[[-200, -200, -200], [0, -200, -200],  [-200, -200, -200], [-200, 0, -200]],
+                         ijk_max=[[0, 200, 200],      [200, 200, 200],  [200, 0, 200],       [200, 200, 200]])
+
+# Use `clip` to both clip the grids and the accompanying JaggedTensor of features
+features = grid.jagged_like(torch.rand(grid.total_voxels, 7).to(device))
+clipped_features, clipped_grid = grid.clip(
+                         features=features,
+                         ijk_min=[[-200, -200, -200], [0, -200, -200],  [-200, -200, -200], [-200, 0, -200]],
+                         ijk_max=[[0, 200, 200],      [200, 200, 200],  [200, 0, 200],       [200, 200, 200]])
+
+```
+
+We visualize these grids, each clipped to a different bounding box.
+
+![](../imgs/fig/clip.png)
+
+
+## Max/Mean Pooling
+
+'Pooling' a grid has the effect of coarsening the resolution of the grid by a constant factor (and can be performed anisotropically for different factors in each `xyz` spatial dimension).  When pooling a grid, the values of the new, coarser grid can be determined by the maximum or mean of the values of the voxels in the original grid that are covered by each voxel in the new grid.  Maximum and mean pooling can be accomplished by the `max_pool` and `avg_pool` operators.
+
+In this example, we create a grid from a mesh and then perform max and mean pooling on it to illustrate the difference between the two.
+
+```python
+import os
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh
+import torch
+import numpy as np
+import point_cloud_utils as pcu
+
+vox_size = 0.02
+num_pts = 10_000
+
+mesh_load_funcs = [load_car_1_mesh]
+
+points = []
+normals = []
+
+for func in mesh_load_funcs:
+    pts, nms = func(mode="vn")
+    pmt = torch.randperm(pts.shape[0])[:num_pts]
+    pts, nms = pts[pmt], nms[pmt]
+    points.append(pts)
+    normals.append(nms)
+
+# JaggedTensors of points and normals
+points = fvdb.JaggedTensor(points)
+normals = fvdb.JaggedTensor(normals)
+
+# Create a grid
+grid = fvdb.gridbatch_from_points(points, voxel_sizes=vox_size)
+
+# Splat the normals into the grid with trilinear interpolation
+vox_normals = grid.splat_trilinear(points, normals)
+
+# Mean Pooling of normals features
+avg_normals, avg_grid = grid.avg_pool(4, vox_normals)
+# Max Pooling of normals features
+max_normals, max_grid = grid.max_pool(4, vox_normals)
+```
+
+We visualize the original grid with values of the mesh normals as features, the features/grid after mean pooling, and the features/grid after max pooling to illustrate how these pooling modes affect the resulting features of the grid.
+
+![](../imgs/fig/max_avg_pool.png)
+
+
+## Subdividing a grid
+
+Subdividing a grid has the effect of increasing the resolution of the grid by a constant factor (and can be performed anisotropically for different factors in each `xyz` spatial dimension).
+
+The most straightforward way to use the `subdivide` method is to provide an integer value for the subdivision factor.  This will result in a grid that is `subdiv_factor` times the resolution of the original grid in each spatial dimension.
+
+```python
+import os
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh
+import torch
+import numpy as np
+import point_cloud_utils as pcu
+
+vox_size = 0.02
+num_pts = 10_000
+
+mesh_load_funcs = [load_car_1_mesh]
+
+points = []
+normals = []
+
+for func in mesh_load_funcs:
+    pts, nms = func(mode="vn")
+    pmt = torch.randperm(pts.shape[0])[:num_pts]
+    pts, nms = pts[pmt], nms[pmt]
+    points.append(pts)
+    normals.append(nms)
+
+# JaggedTensors of points and normals
+points = fvdb.JaggedTensor(points)
+normals = fvdb.JaggedTensor(normals)
+
+# Create a grid
+grid = fvdb.gridbatch_from_points(points, voxel_sizes=vox_size)
+
+# Splat the normals into the grid with trilinear interpolation
+vox_normals = grid.splat_trilinear(points, normals)
+
+# Subdivide by a constant factor of 2
+subdiv_normals, subdiv_grid = grid.subdivide(2, vox_normals)
+```
+
+Here we visualize the original grid on the right and the grid after subdivision by a factor of 2 on the left.
+
+![](../imgs/fig/subdivide.png)
+
+
+In practice, in a deep neural network like a U-Net architecture, the resolution of the grid can be decreased early in the network and then the grid's features need to be concatenated with features of the grid after its resolution is increased again.  When working with traditional, dense 2D or 3D data, cropping any mismatched outputs is straightforward to be able to concatenate these features.  However, in a sparse 3D grid, this is not straightforward and it's entirely unclear how to align the features.
+
+Take this simple example to illustrate the difficulties created by changing the grid topology in this way.  We take the grid created from our mesh, perform Pooling by a factor of 2 and then try to Subdivide by an equal factor of 2 to invert the Pooling operation.
+
+```python continuation
+max_normals, max_grid = grid.max_pool(2, vox_normals)
+
+subdiv_normals, subdiv_grid = max_grid.subdivide(2, max_normals)
+```
+
+
+Let's visualize these results from left to right of the original grid, the grid after max pooling, and the grid after max pooling and then subdividing again by the same factor.
+
+![](../imgs/fig/subdiv_pool_incorrect.png)
+
+
+Notice how we have not obtained the topology of the original grid and have obtained a grid with many more voxels than the original.
+
+To correctly invert the Pooling operation, we can provide the `subdivide` function with a `fine_grid` optional argument which describes the topology we want the grid to have after the subdivision.  The original grid before the Pooling operation can be used as this `fine_grid`.
+
+```python continuation
+max_normals, max_grid = grid.max_pool(2, vox_normals)
+
+# Providing the original grid as our fine_grid target
+subdiv_normals, subdiv_grid = max_grid.subdivide(2, max_normals, fine_grid=grid)
+```
+
+![](../imgs/fig/subdiv_pool_correct.png)
+
+
+Note the matching topology of the grid after the Subdivision operation to the original grid (though the features are different due to the Pooling operation).
+
+One other useful feature of the `subdivide` operator is that this operation can be *masked* so that subdivision is only performed on a subset of the voxels in the grid.
+
+The optional `mask` argument to `subdivide` is a `JaggedTensor` of boolean values that indicates which voxels should be subdivided.  Given this `mask` is simply a `JaggedTensor`, this operation can be made differentiable in a neural network and the `subdivide` operator can be learned.
+
+Let's illustrate a very simple example of how to use the `mask` argument to only subdivide the voxels which have a feature value greater than a certain threshold.
+
+```python continuation
+# Mask the grid with the normals where only the normals with a value greater than 0.5 on the x-axis are subdivided
+mask = vox_normals.jdata[:, 0] > 0.5
+
+subdiv_normals, subdiv_grid = grid.subdivide(2, vox_normals, mask=mask)
+```
+
+![](../imgs/fig/subdiv_mask.png)
+
+
+## Getting the number of voxels per grid in a batch
+
+Getting the number of voxels in the grids of a `GridBatch` can be easily accomplished with `num_voxels`:
+
+```python
+import fvdb
+import torch
+
+# Create a GridBatch of random points
+batch_size = 4
+pts = fvdb.JaggedTensor([torch.rand(10_000*(i+1),3) for i in range(batch_size)])
+grid = fvdb.GridBatch()
+grid.set_from_points(pts, voxel_sizes=0.02)
+
+# Get the number of voxels per grid in the batch
+for batch, num_voxels in enumerate(grid.num_voxels):
+    print(f"Grid {batch} has {num_voxels} voxels")
+```
+```bash
+Grid 0 has 9645 voxels
+Grid 1 has 18503 voxels
+Grid 2 has 26749 voxels
+Grid 3 has 34378 voxels
+```
+
+## Converting between ijk (grid) coordinates and world coordinates
+
+A `GridBatch` can contain multiple grids, each with its own coordinate system that relate the grids' `ijk` integer index space coordinates to `xyz` floating-point world-space coordinates.  These axis-aligned coordinate systems are defined by specifying a list of three-dimensional world-space origin and scale values when we define the topology of the grids in the `GridBatch` like so:
+
+```python
+import fvdb
+import torch
+
+# Create a GridBatch of random points
+batch_size = 4
+pts = fvdb.JaggedTensor([torch.rand(10_000*(i+1),3) for i in range(batch_size)])
+grid = fvdb.GridBatch()
+grid.set_from_points(pts,
+                    voxel_sizes=[[0.02, 0.02, 0.02], [0.03, 0.03, 0.03], [0.04, 0.04, 0.04], [0.05, 0.05, 0.05]],
+                    origins=[[-.1,-.1,-.1], [0,0,0], [.1,.1,.1], [.2,-.2,.2]])
+```
+
+When `voxel_sizes` and `origins` are not defined, it is assumed all grids have a unit scale voxel_size and an origin at `[0.0, 0.0, 0.0]`.
+
+We can use `GridBatch`'s `grid_to_world` function to convert between `ijk` index coordinates and their corresponding world-space `xyz` coordinates.  In this example, let's obtain the world-space position that would lie at the index-space `[1, 1, 1]` point of each grid:
+
+```python continuation
+# Convert ijk coordinates to world coordinates
+ijk = fvdb.JaggedTensor([torch.ones(1,3, dtype=torch.float) for _ in range(batch_size)])
+world_coords = grid.grid_to_world(ijk)
+for i in range(grid.grid_count):
+    print(f"World-space point that lies at index [1,1,1] for Grid {i} is positioned at {world_coords.jdata[world_coords.jidx==i].tolist()}")
+```
+```bash
+World-space point that lies at index [1,1,1] for Grid 0 is positioned at [[-0.07999999821186066, -0.07999999821186066, -0.07999999821186066]]
+World-space point that lies at index [1,1,1] for Grid 1 is positioned at [[0.030000001192092896, 0.030000001192092896, 0.030000001192092896]]
+World-space point that lies at index [1,1,1] for Grid 2 is positioned at [[0.14000000059604645, 0.14000000059604645, 0.14000000059604645]]
+World-space point that lies at index [1,1,1] for Grid 3 is positioned at [[0.25, -0.15000000596046448, 0.25]]
+```
+
+We can also do the inverse operation and convert world-space `xyz` coordinates to their corresponding `ijk` index-space coordinates using `world_to_grid`.  In this example, let's find the `ijk` index coordinates of the voxel which would contain the world-space point located at `[1.0, 1.0, 1.0]` for each grid in our `GridBatch`:
+
+```python continuation
+# Convert world coordinates to ijk coordinates
+xyz = fvdb.JaggedTensor([torch.ones(1,3, dtype=torch.float) for _ in range(batch_size)])
+ijk_coords = grid.world_to_grid(xyz)
+for i in range(grid.grid_count):
+    print(f"Index-space voxel that contains the point [1.0, 1.0, 1.0] for Grid {i} {ijk_coords.jdata[ijk_coords.jidx==i].int().tolist()}")
+```
+```bash
+Index-space voxel that contains the point [1.0, 1.0, 1.0] for Grid 0 [[55, 55, 55]]
+Index-space voxel that contains the point [1.0, 1.0, 1.0] for Grid 1 [[33, 33, 33]]
+Index-space voxel that contains the point [1.0, 1.0, 1.0] for Grid 2 [[22, 22, 22]]
+Index-space voxel that contains the point [1.0, 1.0, 1.0] for Grid 3 [[16, 24, 16]]
+```
+
+While `grid_to_world` and `world_to_grid` are convenient functions for these purposes, the row-major transformation matrices used for these calculations can be obtained directly from a `GridBatch` for use in your own logic:
+
+```python continuation
+print(f"Grid to world matrices:\n{grid.grid_to_world_matrices}")
+print(f"World to grid matrices:\n{grid.grid_to_world_matrices}")
+```
+```bash
+Grid to world matrices:
+tensor([[[ 0.0200,  0.0000,  0.0000,  0.0000],
+         [ 0.0000,  0.0200,  0.0000,  0.0000],
+         [ 0.0000,  0.0000,  0.0200,  0.0000],
+         [-0.1000, -0.1000, -0.1000,  1.0000]],
+
+        [[ 0.0300,  0.0000,  0.0000,  0.0000],
+         [ 0.0000,  0.0300,  0.0000,  0.0000],
+         [ 0.0000,  0.0000,  0.0300,  0.0000],
+         [ 0.0000,  0.0000,  0.0000,  1.0000]],
+
+        [[ 0.0400,  0.0000,  0.0000,  0.0000],
+         [ 0.0000,  0.0400,  0.0000,  0.0000],
+         [ 0.0000,  0.0000,  0.0400,  0.0000],
+         [ 0.1000,  0.1000,  0.1000,  1.0000]],
+
+        [[ 0.0500,  0.0000,  0.0000,  0.0000],
+         [ 0.0000,  0.0500,  0.0000,  0.0000],
+         [ 0.0000,  0.0000,  0.0500,  0.0000],
+         [ 0.2000, -0.2000,  0.2000,  1.0000]]])
+World to grid matrices:
+tensor([[[ 0.0200,  0.0000,  0.0000,  0.0000],
+         [ 0.0000,  0.0200,  0.0000,  0.0000],
+         [ 0.0000,  0.0000,  0.0200,  0.0000],
+         [-0.1000, -0.1000, -0.1000,  1.0000]],
+
+        [[ 0.0300,  0.0000,  0.0000,  0.0000],
+         [ 0.0000,  0.0300,  0.0000,  0.0000],
+         [ 0.0000,  0.0000,  0.0300,  0.0000],
+         [ 0.0000,  0.0000,  0.0000,  1.0000]],
+
+        [[ 0.0400,  0.0000,  0.0000,  0.0000],
+         [ 0.0000,  0.0400,  0.0000,  0.0000],
+         [ 0.0000,  0.0000,  0.0400,  0.0000],
+         [ 0.1000,  0.1000,  0.1000,  1.0000]],
+
+        [[ 0.0500,  0.0000,  0.0000,  0.0000],
+         [ 0.0000,  0.0500,  0.0000,  0.0000],
+         [ 0.0000,  0.0000,  0.0500,  0.0000],
+         [ 0.2000, -0.2000,  0.2000,  1.0000]]])
+```
+
+
+## Convolution
+
+Convolving the features of a `GridBatch` can be accomplished with either a high-level `torch.nn.Module` derived class provided by `fvdb.nn` or with more low-level methods available with `GridBatch`, we will illustrate both techniques.
+
+### High-level Usage with `fvdb.nn`
+
+`fvdb.nn.SparseConv3d` provides a high-level `torch.nn.Module` class for convolution on `fvdb` classes that is an analogue to the use of `torch.nn.Conv3d`.  Using this module is the recommended functionality for performing convolution with `fvdb` because it not only manages functionality such as initializing the weights of the convolution and calling appropriate backend implementation functions but it also provides certain backend optimizations which will be illustrated in the [Low-level usage](#low-level-usage-with-gridbatch) section.
+
+One thing to note is `fvdb.nn.SparseConv3d` operates on a class that wraps a `GridBatch` and `JaggedTensor` together into a convenience object, `fvdb.VDBTensor`, which is used by all the `fvdb.nn` modules.
+
+A simple example of using `fvdb.nn.SparseConv3d` is as follows:
+
+```python
+import fvdb
+import fvdb.nn as fvdbnn
+from fvdb.utils.examples import load_car_1_mesh
+import torch
+import numpy as np
+import point_cloud_utils as pcu
+
+num_pts = 10_000
+vox_size = 0.02
+
+mesh_load_funcs = [load_car_1_mesh]
+
+points = []
+normals = []
+
+for func in mesh_load_funcs:
+    pts, nms = func(mode="vn")
+    pmt = torch.randperm(pts.shape[0])[:num_pts]
+    pts, nms = pts[pmt], nms[pmt]
+    points.append(pts)
+    normals.append(nms)
+
+# JaggedTensors of points and normals
+points = fvdb.JaggedTensor(points)
+normals = fvdb.JaggedTensor(normals)
+
+# Create a grid
+grid = fvdb.gridbatch_from_points(points, voxel_sizes=vox_size)
+
+# Splat the normals into the grid with trilinear interpolation
+vox_normals = grid.splat_trilinear(points, normals)
+
+# VDBTensor is a simple wrapper of a grid and a feature tensor
+vdbtensor = fvdbnn.VDBTensor(grid, vox_normals)
+
+# fvdb.nn.SparseConv3d is a convenient torch.nn.Module implementing the fVDB convolution
+conv = fvdbnn.SparseConv3d(in_channels=3, out_channels=3, kernel_size=3, stride=1, bias=False).to(vdbtensor.device)
+
+output = conv(vdbtensor)
+```
+Let's visualize the original grid with normals visualized as colours alongside the result of these features after a convolution initialized with random weights:
+![](../imgs/fig/simple_conv.png)
+
+For stride values greater than 1, the output of the convolution will be a grid with a smaller resolution than the input grid (similar in topological effect to the output of a Pooling operator).  Let's illustrate this:
+
+```python continuation
+# We would expect for stride=2 that the output grid would have half the resolution (or twice the world-space size) of the input grid
+conv = fvdbnn.SparseConv3d(in_channels=3, out_channels=3, kernel_size=3, stride=2, bias=False).to(vdbtensor.device)
+
+output = conv(vdbtensor)
+```
+
+![](../imgs/fig/stride_conv.png)
+
+
+Transposed convolution can be performed with `fvdb.nn.SparseConv3d` which can increase the resolution of the grid.  It only really makes sense to perform transposed sparse convolution with a target grid topology we wish to produce with this operation (see the [Pooling Operators](#maxmean-pooling) for an explanation).  Therefore, an `out_grid` argument must be provided in this case to specify the target grid topology:
+
+```python continuation
+# Tranposed convolution operator, stride=2
+transposed_conv = fvdbnn.SparseConv3d(in_channels=3, out_channels=3, kernel_size=3, stride=2, bias=False, transposed=True).to(vdbtensor.device)
+
+# Note the use of the `out_grid` argument to specify the target grid topology
+transposed_output = transposed_conv(output, out_grid=vdbtensor.grid)
+```
+
+Here we visuzlie the original grid, the grid after strided convolution and the grid after transposed convolution inverts the topological operation of the strided convolution to produce the same topology as the original grid with the features convolved by our two layers:
+
+![](../imgs/fig/transposed_stride_conv.png)
+
+
+### Low-level Usage with `GridBatch`
+
+The [high-level `fvdb.nn.SparseConv3d` class](#high-level-convolution-with-fvdbnn) wraps several pieces of `GridBatch` functionality to provide a convenient `torch.nn.Module` for convolution.  However, for a more low-level approach that accomplishes the same outcome, the `GridBatch` class itself can be the starting point for performing convolution on the grid and its features.  We will illustrate this approach for completeness, though we do recommend the use of the `fvdb.nn.SparseConv3d` Module for most use-cases.
+
+Using the `GridBatch` convolution functions directly requires a little more knowledge about the implementation under-the-hood.  Due to the nature of a sparse grid, in order to make convolution performant, it is useful to pre-compute a mapping of which features in the input grid will contribute to the values of the output grid when convolved by a kernel of a particular dimension and stride.  This mapping structure is called a 'kernel map'.
+
+The kernel map, as well as the functionality for using it to compute the convolution, is contained within a `fvdb.SparseConvPackInfo` class which can be constructed by `GridBatch.sparse_conv_kernel_map`.  A `fvdb.SparseConvPackInfo` must perform a pre-computation of the kernel map based on the style expected by the backend implementation of the convolution utilized by `fvdb.SparseConvPackInfo.sparse_conv_3d`.  Here is an example of how to construct a `fvdb.SparseConvPackInfo` and use it to perform a convolution:
+
+```python
+import fvdb
+import fvdb.nn as fvdbnn
+from fvdb.utils.examples import load_car_1_mesh
+import torch
+import numpy as np
+import point_cloud_utils as pcu
+
+num_pts = 10_000
+vox_size = 0.02
+
+mesh_load_funcs = [load_car_1_mesh]
+
+points = []
+normals = []
+
+for func in mesh_load_funcs:
+    pts, nms = func(mode="vn")
+    pmt = torch.randperm(pts.shape[0])[:num_pts]
+    pts, nms = pts[pmt], nms[pmt]
+    points.append(pts)
+    normals.append(nms)
+
+# JaggedTensors of points and normals
+points = fvdb.JaggedTensor(points)
+normals = fvdb.JaggedTensor(normals)
+
+# Create a grid
+grid = fvdb.gridbatch_from_points(points, voxel_sizes=vox_size)
+
+# Splat the normals into the grid with trilinear interpolation
+vox_normals = grid.splat_trilinear(points, normals)\
+
+# Create the kernel map (housed in a SparseConvPackInfo) and the output grid's topology based on the kernel parameters
+sparse_conv_packinfo, out_grid = grid.sparse_conv_kernel_map(kernel_size=3, stride=1)
+
+# The kernel map must be pre-computed based on the backend implementation we plan on using, here we use gather/scatter, the default implementation
+sparse_conv_packinfo.build_gather_scatter()
+
+# Create random weights for our convolution kernel of size 3x3x3 that takes 3 input channels and produces 3 output channels
+kernel_weights = torch.randn(3, 3, 3, 3, 3, device=grid.device)
+
+# Perform convolution on the normals colours.  Gather/scatter is used as the backend, it is the default
+conv_vox_normals = sparse_conv_packinfo.sparse_conv_3d(vox_normals, weights=kernel_weights, backend=fvdb.ConvPackBackend.GATHER_SCATTER)
+```
+Here we visualize the output of our convolution alongside the original grid with normals visualized as colours:
+![](../imgs/fig/gridbatch_conv.png)
+
+The kernel map can potentially be expensive to compute, so it is often useful to re-use the `SparseConvPackInfo` in the same network to perform a convolution on other features or with different weights.  This optimization is something `fvdb.nn.SparseConv3d` attempts to do where appropriate and is one reason we recommend using `fvdb.nn.SparseConv3d` over this low-level approach.
@@ -0,0 +1,186 @@
+# Building Sparse Grids
+
+We introduce several ways to construct sparse voxel grids from various data sources including point clouds, coordinate lists, triangle meshes, and deriving from other grids.
+All the examples below could be found in full version (including visualization) at `examples/grid_building.py` and `examples/grid_subdivide_coarsen.py`.
+
+## From Coordinate Lists
+
+If you already have an integer list of the `ijk` coordinates for each voxel in the grid, you could directly build grids from the list.
+Additionally, you will have to specify the voxel sizes and voxel origins for the grid.
+An example is as follows:
+
+```python
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+
+coords_1, _ = load_car_1_mesh()
+coords_2, _ = load_car_2_mesh()
+coords_jagged = fvdb.JaggedTensor([
+    coords_1.long().cuda(),
+    coords_2.long().cuda()
+])
+voxel_sizes = [[0.1, 0.1, 0.1], [0.15, 0.15, 0.15]]
+
+grid = fvdb.gridbatch_from_ijk(coords_jagged, voxel_sizes=voxel_sizes, origins=[0.0] * 3)
+```
+
+The above code assumes that you want to build a grid with two batch elements, one with voxel size `[0.1, 0.1, 0.1]`, and the other with voxel size `[0.15, 0.15, 0.15]` (although usually you just want all the elements in your batch to have the same size, in which case you could just pass in `voxel_sizes=[0.1, 0.1, 0.1]`).
+The same logic applies for `origins` that specifies the world coordinates of voxel `(0, 0, 0)`, which we set to the origin here.
+The grid will be constructed on the same device as `coords_jagged`, which is a `JaggedTensor`. The JaggedTensor is built from two numpy `int64` arrays with size `[*, 3]` called `coords_1` and `coords_2`.
+
+![build_from_coordinates.png](../imgs/fig/build_from_coordinates.png)
+
+## From Point Clouds
+
+You could either choose to quantize the coordinates of your point cloud into `ijk` coordinates yourself (e.g. using `np.unique((xyz / voxel_size).floor(), axis=0)`), or let `fvdb` handle this logic. Specifically, you could do:
+
+```python
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+
+coords_1, _ = load_car_1_mesh()
+coords_2, _ = load_car_2_mesh()
+
+# Assemble point clouds into JaggedTensor
+pcd_jagged = fvdb.JaggedTensor([
+    coords_1.cuda(),
+    coords_2.cuda()
+])
+voxel_sizes = [[0.1, 0.1, 0.1], [0.15, 0.15, 0.15]]
+
+# Method 1:
+grid_a1 = fvdb.gridbatch_from_points(pcd_jagged, voxel_sizes=voxel_sizes, origins=[0.0] * 3)
+
+# Method 2:
+grid_a2 = fvdb.GridBatch(device=pcd_jagged.device)
+grid_a2.set_from_points(pcd_jagged, voxel_sizes=voxel_sizes, origins=[0.0] * 3)
+```
+
+Above we show two methods of building grids from points. Similar functions exist for other grid building approaches. The built grids are shown as following:
+
+![build_from_points.png](../imgs/fig/build_from_points.png)
+
+In some applications, you may want to build a dilated version of the grid by ensuring that all $2\times 2 \times 2$ voxels around each point are included in the built grid. That said, you could build the grid by:
+
+```python
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+
+coords_1, _ = load_car_1_mesh()
+coords_2, _ = load_car_2_mesh()
+
+# Assemble point clouds into JaggedTensor
+pcd_jagged = fvdb.JaggedTensor([
+    coords_1.cuda(),
+    coords_2.cuda()
+])
+voxel_sizes = [[0.1, 0.1, 0.1], [0.15, 0.15, 0.15]]
+
+# Build grid from containing nearest voxels to the points
+grid_b = fvdb.gridbatch_from_nearest_voxels_to_points(pcd_jagged, voxel_sizes=voxel_sizes, origins=[0.0] * 3)
+```
+
+![build_from_points_nn.png](../imgs/fig/build_from_points_nn.png)
+
+
+## From Meshes
+
+We allow building grids enclosing a triangle mesh easily. The given triangle mesh does not have to be manifold nor watertight and it will be treated as a triangle soup internally.
+An example to build grids from meshes is shown as follows:
+
+```python
+import fvdb
+from fvdb.utils.examples import load_car_1_mesh, load_car_2_mesh
+
+mesh_1_v, mesh_1_f = load_car_1_mesh(mode='vf')
+mesh_2_v, mesh_2_f = load_car_2_mesh(mode='vf')
+
+mesh_v_jagged = fvdb.JaggedTensor([
+    mesh_1_v.float().cuda(),
+    mesh_2_v.float().cuda()
+])
+mesh_f_jagged = fvdb.JaggedTensor([
+    mesh_1_f.long().cuda(),
+    mesh_2_f.long().cuda()
+])
+
+voxel_sizes = [[0.1, 0.1, 0.1], [0.15, 0.15, 0.15]]
+grid = fvdb.gridbatch_from_mesh(mesh_v_jagged, mesh_f_jagged, voxel_sizes=voxel_sizes, origins=[0.0] * 3)
+```
+
+Here `mesh_1_v` and `mesh_1_f` are the vertex array and triangle array of the mesh to build grid from, with the shape of $(V, 3)$ and $(F, 3)$. The triangle array is an integer array that indexes into the vertex array (starting from 0 for each element in the batch). Same for another `mesh_2_v` and `mesh_2_f`.
+
+![build_from_mesh.png](../imgs/fig/build_from_mesh.png)
+
+## From Dense
+
+We have APIs for you to build dense grids of shape $(D, H, W)$ containing the full $D\times H \times W$ voxels.
+
+```python
+import fvdb
+
+grid = fvdb.gridbatch_from_dense(num_grids=1, dense_dims=[32, 32, 32], device="cuda")
+```
+
+![build_from_dense.png](../imgs/fig/build_from_dense.png)
+
+If you are comparing the performance of dense pytorch 3D tensors vs sparse grids, it is usually very helpful to build the exact same input (including grid and features). In `fvdb.nn`, we provide a thin wrapper class `VDBTensor` that works like a `torch.Tensor`, yet enclosing the grid topology. To convert data back and forth from dense PyTorch `Tensor`s, we could do:
+
+```python
+import torch
+import fvdb
+from fvdb.nn import VDBTensor
+
+# Easy way to initialize a VDBTensor from a torch 3D tensor [B, D, H, W, C]
+dense_data = torch.ones(2, 32, 32, 32, 16).cuda()
+sparse_data = fvdb.nn.vdbtensor_from_dense(dense_data, voxel_sizes=[0.1] * 3)
+dense_data_back = sparse_data.to_dense()
+assert torch.all(dense_data == dense_data_back)
+```
+
+Here `sparse_data` will be a `fvdb.nn.VDBTensor` class, containing both `feature` and `grid` attribute.
+Such a class could be fed into all the neural network components available in `fvdb.nn`.
+
+## Deriving from other grids
+
+### Dual grid
+
+A dual grid (of a primal grid) is also a grid, with its voxel *centers* covering the *corners* of the primal grid, and *corners* taking the positions of the *centers* of the primal grid.
+A dual grid shares the same voxel sizes as its primal grid, and is just a simple shift of half the voxel size in translation.
+In the picture below, green grid is the primal grid while purple grid is its dual.
+
+![build_dual.png](../imgs/fig/build_dual.png)
+
+To create a dual grid from a given primal grid, use `GridBatch.dual_grid()`:
+
+```python
+import fvdb
+from fvdb.utils.examples import load_dragon_mesh
+
+coords_1, _ = load_dragon_mesh()
+
+grid_primal = fvdb.gridbatch_from_points(fvdb.JaggedTensor([coords_1]))
+
+grid_dual = grid_primal.dual_grid()
+```
+
+### Subdividing and Coarsening a grid
+
+The grid could be subdivided or coarsened with a subdivision/coarsening factor provided:
+
+```python
+import fvdb
+from fvdb.utils.examples import load_happy_mesh
+
+coords_1, _ = load_happy_mesh(mode='vf')
+
+grid = fvdb.gridbatch_from_points(fvdb.JaggedTensor([coords_1]))
+
+grid_subdivided = grid.subdivided_grid(2)
+grid_coarsened = grid.coarsened_grid(2)
+```
+
+![build_coarse_subdivide.png](../imgs/fig/build_coarse_subdivide.png)
+
+The features associated with the grids could be processed via the examples in `Grid Operations`.
+Refer to the corresponding section for more details.
@@ -0,0 +1,150 @@
+# Sparse Grid I/O
+
+We give an overview of ways to save and load sparse grids including how fVDB's serialized format relates to other libraries such as OpenVDB and NanoVDB.  All of the examples in this tutorial are available in the `examples/io.py` file of the fVDB repository.
+
+In these examples we will be using tools which are part of the `NanoVDB` project such as `nanovdb_convert`.  It is assumed that the `NanoVDB` tools are available to call (i.e. findable on the system's `$PATH`).  If you have not already installed these tools, you can find instructions on building NanoVDB with OpenVDB from its documentation:
+
+https://www.openvdb.org/documentation/doxygen/NanoVDB_HowToBuild.html
+
+While not necessary for fVDB's functionality, they are useful utilities to have available for inspecting and manipulating sparse grids.
+
+## Python Serialization
+
+Batches of sparse grids can be serialized to a NanoVDB file using the `fvdb.save` method.  Here, we create two grids of different sizes and numbers of points and save them to a compressed NanoVDB file with specified names.  The names are optional.
+
+```python
+import torch
+import fvdb
+import tempfile
+import os
+import subprocess
+
+p = fvdb.JaggedTensor(
+    [
+        torch.randn(10, 3),
+        torch.randn(100, 3),
+    ]
+)
+grid = fvdb.gridbatch_from_points(
+    p, voxel_sizes=[[0.1, 0.1, 0.1], [0.15, 0.15, 0.15]], origins=[0.0] * 3
+)
+
+# save the grid and features to a compressed nvdb file
+path = os.path.join(tempfile.gettempdir(), "two_random_grids.nvdb")
+fvdb.save(path, grid, names=["taco1", "taco2"], compressed=True)
+```
+
+We can use the `nanovdb_print` command line tool to show information about our saved file.  Note how our grids have the `INDEX` class since they have no features and only store the voxel indices.
+
+```bash
+
+The file "/tmp/tmpwnu7qc_7/two_random_grids.nvdb" contains the following 2 grids:
+#  Name   Type     Class  Version  Codec  Size      File      Scale             # Voxels  Resolution
+1  taco1  OnIndex  INDEX  32.6.0   BLOSC  1.453 MB  7.181 KB  (0.1,0.1,0.1)     10        22 x 43 x 37
+2  taco2  OnIndex  INDEX  32.6.0   BLOSC  2.326 MB  12.7 KB   (0.15,0.15,0.15)  100       33 x 39 x 38
+```
+
+We can include N-dimensional features by passing a JaggedTensor as the second argument (or `data` kwarg) to `fvdb.save`.  Here, we create a grid with a single, float feature channel for our grids and save it to a `nvdb` file.
+
+```python continuation
+# a single, scalar float feature per grid
+feats = fvdb.JaggedTensor([torch.randn(x, 1) for x in grid.num_voxels])
+
+# save the grid and features to a compressed nvdb file
+path = os.path.join(tempfile.gettempdir(), "two_random_grids.nvdb")
+fvdb.save(path, grid, feats, names=["taco1", "taco2"], compressed=True)
+```
+
+Again, we can use the `nanovdb_print` command line tool to show information about our saved file.
+
+```bash
+The file "/tmp/tmpg9ol5841/two_random_grids.nvdb" contains the following 2 grids:
+#  Name   Type   Class  Version  Codec  Size      File      Scale    # Voxels  Resolution
+1  taco1  float  ?      32.6.0   BLOSC  1.47 MB   7.959 KB  (1,1,1)  10        26 x 30 x 36
+2  taco2  float  ?      32.6.0   BLOSC  2.411 MB  16.48 KB  (1,1,1)  100       32 x 33 x 34
+```
+
+Note how our serialized NanoVDB grids are now of type `float`.  fVDB will automatically map N-dimensional features to appropriate NanoVDB types.  For feature sizes that don't naturally map to any NanoVDB data types, fVDB will save the feature data as NanoVDB blind-data which will be appropriately read back as N-dimensional features by fVDB.
+
+Let's try to save the same two grids with a `Vec3d` type by creating a JaggedTensor of 3-dimensional double-precision features.
+
+```python continuation
+# a 3-vector double feature per grid
+feats = fvdb.JaggedTensor([torch.randn(x, 3, dtype=torch.float64) for x in grid.num_voxels])
+
+# save the grid and features to a compressed nvdb file
+saved_nvdb = os.path.join(tempfile.gettempdir(), "two_random_vec3d_grids.nvdb")
+fvdb.save(saved_nvdb, grid, feats, names=["taco1", "taco2"], compressed=True)
+```
+
+```bash
+The file "/tmp/tmpwnu7qc_7/two_random_grids.nvdb" contains the following 2 grids:
+#  Name   Type   Class  Version  Codec  Size      File      Scale    # Voxels  Resolution
+1  taco1  Vec3d  ?      32.6.0   BLOSC  6.077 MB  28.18 KB  (1,1,1)  10        23 x 36 x 35
+2  taco2  Vec3d  ?      32.6.0   BLOSC  7.346 MB  41.37 KB  (1,1,1)  100       37 x 40 x 34
+```
+
+## Loading NanoVDB Files
+
+Loading NanoVDB files is as simple as calling `fvdb.load`.  You can optionally supply a PyTorch device you'd like the grids and features loaded onto.  Here, we load the two grids we saved in the previous section onto our GPU.
+
+```python continuation
+# Load the grid and features from the compressed nvdb file
+grid_batch, features, names = fvdb.load(saved_nvdb, device=torch.device("cuda:0"))
+print("Loaded grid batch total number of voxels: ", grid_batch.total_voxels)
+print("Loaded grid batch data type: %s, device: %s" % (features.dtype, features.device))
+```
+
+```bash
+Loaded grid batch total number of voxels:  110
+Loaded grid batch data type: torch.float64, device: cuda:0
+```
+
+## Saving/Loading OpenVDB Files
+
+While saving and loading from OpenVDB files is not directly supported by fVDB, it is possible to easily convert between NanoVDB and OpenVDB files using the `nanovdb_convert` command line tool.  Here, we convert our previously saved NanoVDB file to an OpenVDB file.
+
+```python notest
+vdb_path = os.path.join(tempfile.gettempdir(), "two_random_grids.vdb")
+convert_cmd = "nanovdb_convert -v %s %s" % (saved_nvdb, vdb_path)
+print("nanovdb_convert our nvdb to vdb: ", convert_cmd)
+print(subprocess.check_output(convert_cmd.split()).decode("utf-8"))
+```
+
+```bash
+nanovdb_convert our nvdb to vdb:  nanovdb_convert -v /tmp/tmpnr7gk0tf/two_random_vec3d_grids.nvdb /tmp/tmpnr7gk0tf/two_random_grids.vdb
+Opening NanoVDB file named "/tmp/tmpnr7gk0tf/two_random_vec3d_grids.nvdb"
+Read 1 NanoGrid(s) from the file named "/tmp/tmpnr7gk0tf/two_random_vec3d_grids.nvdb"
+Converting NanoVDB grid named "taco1" to OpenVDB
+Converting NanoVDB grid named "taco2" to OpenVDB
+```
+
+From here, our grid can be loaded by OpenVDB tools.  Roundtripping our converted cache back to NanoVDB is possible with the `nanovdb_convert` tool as well.
+
+Loading the converted OpenVDB file into fVDB shows our familiar grids and features as we expect:
+
+
+```python notest
+convert_cmd = "nanovdb_convert -v -f %s %s" % (  # -f flag forces overwriting existing file
+    vdb_path,
+    saved_nvdb,
+)
+print("nanovdb_convert roundtrip the vdb to nvdb: ", convert_cmd)
+print(subprocess.check_output(convert_cmd.split()).decode("utf-8"))
+
+# Load the nvdb file of the converted vdb
+grid_batch, features, names = fvdb.load(saved_nvdb, device=torch.device("cuda:0"))
+print("Loaded grid batch total number of voxels: ", grid_batch.total_voxels)
+print("Loaded grid batch data type: %s, device: %s" % (features.dtype, features.device))
+print("\n")
+```
+
+```bash
+nanovdb_convert roundtrip the vdb to nvdb:  nanovdb_convert -v -f /tmp/tmpuwq2r5mx/two_random_grids.vdb /tmp/tmpuwq2r5mx/two_random_vec3d_grids.nvdb
+Opening OpenVDB file named "/tmp/tmpuwq2r5mx/two_random_grids.vdb"
+Converting OpenVDB grid named "taco1" to NanoVDB
+Converting OpenVDB grid named "taco2" to NanoVDB
+
+Loaded grid batch total number of voxels:  110
+Loaded grid batch data type: torch.float64, device: cuda:0
+```
@@ -0,0 +1 @@
+# Ray Tracing
@@ -0,0 +1,272 @@
+# A Simple Convolutional U-Net
+
+In this tutorial, you will be guided on how to build a simple sparse convolutional neural network using fVDB.
+If you were using MinkowskiEngine to tackle sparse 3D data previously, we will also guide you step-by-step to help you smoothly transfer from it and enjoy speed-ups and memory-savings.
+
+In our simplistic U-Net case, we want to build a Res-UNet with four layers, and each layer contains several blocks.
+First, we import basic `fvdb` libraries:
+
+```python
+import fvdb
+import fvdb.nn as fvnn
+from fvdb.nn import VDBTensor
+import torch
+```
+
+Here `fvdb.nn` is a namespace similar to `torch.nn`, containing a broad definition of different neural layers.
+`VDBTensor` is a very thin wrapper around a grid (with type `GridBatch`) and its corresponding feature (with type `JaggedTensor`), and internally makes sure that the two members align.
+It also overloads a bunch of operators such as arithmetic computations. Please refer to our API docs to learn more.
+
+We could then build a basic block as follows:
+
+```python continuation
+class BasicBlock(torch.nn.Module):
+    expansion = 1
+
+    def __init__(self, in_channels: int, out_channels: int, downsample=None, bn_momentum: float = 0.1):
+        super().__init__()
+        self.conv1 = fvnn.SparseConv3d(in_channels, out_channels, kernel_size=3, stride=1)
+        self.norm1 = fvnn.BatchNorm(out_channels, momentum=bn_momentum)
+        self.conv2 = fvnn.SparseConv3d(out_channels, out_channels, kernel_size=3, stride=1)
+        self.norm2 = fvnn.BatchNorm(out_channels, momentum=bn_momentum)
+        self.relu = fvnn.ReLU(inplace=True)
+        self.downsample = downsample
+
+    def forward(self, x: VDBTensor):
+        residual = x
+
+        out = self.conv1(x)
+        out = self.norm1(out)
+        out = self.relu(out)
+
+        out = self.conv2(out)
+        out = self.norm2(out)
+
+        if self.downsample is not None:
+          residual = self.downsample(x)
+
+        out += residual
+        out = self.relu(out)
+
+        return out
+```
+
+This defines a similar block as `MinkowskiEngine`:
+
+```python notest
+import MinkowskiEngine as ME
+
+
+class BasicBlock(torch.nn.Module):
+    expansion = 1
+
+    def __init__(self, in_channels: int, out_channels: int, downsample=None, bn_momentum: float = 0.1):
+        super().__init__()
+        self.conv1 = ME.MinkowskiConvolution(
+            in_channels, out_channels, kernel_size=3, stride=1, dilation=1, dimension=3)
+        self.norm1 = ME.MinkowskiBatchNorm(out_channels, momentum=bn_momentum)
+        self.conv2 = ME.MinkowskiConvolution(
+            out_channels, out_channels, kernel_size=3, stride=1, dilation=1, dimension=3)
+        self.norm2 = ME.MinkowskiBatchNorm(out_channels, momentum=bn_momentum)
+        self.relu = ME.MinkowskiReLU(inplace=True)
+        self.downsample = downsample
+
+    def forward(self, x):
+        residual = x
+
+        out = self.conv1(x)
+        out = self.norm1(out)
+        out = self.relu(out)
+
+        out = self.conv2(out)
+        out = self.norm2(out)
+
+        if self.downsample is not None:
+          residual = self.downsample(x)
+
+        out += residual
+        out = self.relu(out)
+
+        return out
+```
+
+All the network layers are fully compatible with `torch.nn`. The only difference is that they take `VDBTensor` as input and return a `VDBTensor`.
+A full network definition could then be built as:
+
+```python continuation
+class FVDBUNetBase(torch.nn.Module):
+    LAYERS = (2, 2, 2, 2, 2, 2, 2, 2)
+    CHANNELS = (32, 64, 128, 256, 256, 128, 96, 96)
+    INIT_DIM = 32
+    OUT_TENSOR_STRIDE = 1
+
+    def __init__(self, in_channels, out_channels, D=3):
+        super().__init__()
+
+        # Output of the first conv concated to conv6
+        self.inplanes = self.INIT_DIM
+        self.conv0p1s1 = fvnn.SparseConv3d(in_channels, self.inplanes, kernel_size=5, stride=1, bias=False)
+        self.bn0 = fvnn.BatchNorm(self.inplanes)
+
+        self.conv1p1s2 = fvnn.SparseConv3d(self.inplanes, self.inplanes, kernel_size=2, stride=2, bias=False)
+        self.bn1 = fvnn.BatchNorm(self.inplanes)
+
+        self.block1 = self._make_layer(BasicBlock, self.CHANNELS[0], self.LAYERS[0])
+
+        self.conv2p2s2 = fvnn.SparseConv3d(
+            self.inplanes, self.inplanes, kernel_size=2, stride=2, bias=False)
+        self.bn2 = fvnn.BatchNorm(self.inplanes)
+
+        self.block2 = self._make_layer(BasicBlock, self.CHANNELS[1], self.LAYERS[1])
+
+        self.conv3p4s2 = fvnn.SparseConv3d(
+            self.inplanes, self.inplanes, kernel_size=2, stride=2, bias=False)
+
+        self.bn3 = fvnn.BatchNorm(self.inplanes)
+        self.block3 = self._make_layer(BasicBlock, self.CHANNELS[2], self.LAYERS[2])
+
+        self.conv4p8s2 = fvnn.SparseConv3d(
+            self.inplanes, self.inplanes, kernel_size=2, stride=2, bias=False)
+        self.bn4 = fvnn.BatchNorm(self.inplanes)
+        self.block4 = self._make_layer(BasicBlock, self.CHANNELS[3], self.LAYERS[3])
+
+        self.convtr4p16s2 = fvnn.SparseConv3d(
+            self.inplanes, self.CHANNELS[4], kernel_size=2, stride=2, transposed=True, bias=False)
+        self.bntr4 = fvnn.BatchNorm(self.CHANNELS[4])
+
+        self.inplanes = self.CHANNELS[4] + self.CHANNELS[2]
+        self.block5 = self._make_layer(BasicBlock, self.CHANNELS[4], self.LAYERS[4])
+        self.convtr5p8s2 = fvnn.SparseConv3d(
+            self.inplanes, self.CHANNELS[5], kernel_size=2, stride=2, transposed=True, bias=False)
+        self.bntr5 = fvnn.BatchNorm(self.CHANNELS[5])
+
+        self.inplanes = self.CHANNELS[5] + self.CHANNELS[1]
+        self.block6 = self._make_layer(BasicBlock, self.CHANNELS[5], self.LAYERS[5])
+        self.convtr6p4s2 = fvnn.SparseConv3d(
+            self.inplanes, self.CHANNELS[6], kernel_size=2, stride=2, transposed=True, bias=False)
+        self.bntr6 = fvnn.BatchNorm(self.CHANNELS[6])
+
+        self.inplanes = self.CHANNELS[6] + self.CHANNELS[0]
+        self.block7 = self._make_layer(BasicBlock, self.CHANNELS[6], self.LAYERS[6])
+        self.convtr7p2s2 = fvnn.SparseConv3d(
+            self.inplanes, self.CHANNELS[7], kernel_size=2, stride=2, transposed=True, bias=False)
+        self.bntr7 = fvnn.BatchNorm(self.CHANNELS[7])
+
+        self.inplanes = self.CHANNELS[7] + self.INIT_DIM
+        self.block8 = self._make_layer(BasicBlock, self.CHANNELS[7], self.LAYERS[7])
+
+        self.final = fvnn.SparseConv3d(self.CHANNELS[7], out_channels, kernel_size=1)
+        self.relu = fvnn.ReLU(inplace=True)
+
+    def _make_layer(self, block, planes, blocks):
+        downsample = None
+        if self.inplanes != planes * block.expansion:
+            downsample = torch.nn.Sequential(
+                fvnn.SparseConv3d(
+                    self.inplanes,
+                    planes * block.expansion,
+                    kernel_size=1,
+                    stride=1
+                ),
+                fvnn.BatchNorm(planes * block.expansion),
+            )
+        layers = []
+        layers.append(
+            BasicBlock(
+                self.inplanes, planes,
+                downsample=downsample
+            )
+        )
+        self.inplanes = planes * block.expansion
+        for _ in range(1, blocks):
+            layers.append(BasicBlock(self.inplanes, planes))
+
+        return torch.nn.Sequential(*layers)
+
+    def forward(self, x):
+        out = self.conv0p1s1(x)
+        out = self.bn0(out)
+        out_p1 = self.relu(out)
+        grid1 = out_p1.grid
+
+        out = self.conv1p1s2(out_p1)
+        out = self.bn1(out)
+        out = self.relu(out)
+        out_b1p2 = self.block1(out)
+        grid2 = out_b1p2.grid
+
+        out = self.conv2p2s2(out_b1p2)
+        out = self.bn2(out)
+        out = self.relu(out)
+        out_b2p4 = self.block2(out)
+        grid4 = out_b2p4.grid
+
+        out = self.conv3p4s2(out_b2p4)
+        out = self.bn3(out)
+        out = self.relu(out)
+        out_b3p8 = self.block3(out)
+        grid8 = out_b3p8.grid
+
+        # tensor_stride=16
+        out = self.conv4p8s2(out_b3p8)
+        out = self.bn4(out)
+        out = self.relu(out)
+        out = self.block4(out)
+
+        # tensor_stride=8
+        out = self.convtr4p16s2(out, out_grid=grid8)
+        out = self.bntr4(out)
+        out = self.relu(out)
+
+        out = fvdb.jcat([out, out_b3p8], dim=1)
+        out = self.block5(out)
+
+        # tensor_stride=4
+        out = self.convtr5p8s2(out, out_grid=grid4)
+        out = self.bntr5(out)
+        out = self.relu(out)
+
+        out = fvdb.jcat([out, out_b2p4], dim=1)
+        out = self.block6(out)
+
+        # tensor_stride=2
+        out = self.convtr6p4s2(out, out_grid=grid2)
+        out = self.bntr6(out)
+        out = self.relu(out)
+
+        out = fvdb.jcat([out, out_b1p2], dim=1)
+        out = self.block7(out)
+
+        # tensor_stride=1
+        out = self.convtr7p2s2(out, out_grid=grid1)
+        out = self.bntr7(out)
+        out = self.relu(out)
+
+        out = fvdb.jcat([out, out_p1], dim=1)
+        out = self.block8(out)
+
+        return self.final(out)
+```
+
+Please note that here, when we apply transposed convolution layers, we additionally introduce the `out_grid` keyword arguments.
+This is needed to guide the output domain of the network, because for perception networks, the output grid topology should align with the input topology.
+Note that fVDB will NOT cache the grids to maintain maximum flexibility.
+
+To perform inference with the network, you could simply create a VDBTensor and feed it into the model:
+
+```python continuation
+coords = fvdb.JaggedTensor([
+    (torch.randn(10_000, 3, device='cuda')),
+    (torch.randn(11_000, 3, device='cuda')),
+])
+
+grid = fvdb.gridbatch_from_points(coords)
+features = grid.jagged_like(torch.randn(grid.total_voxels, 32, device='cuda'))
+sinput = fvnn.VDBTensor(grid, features)
+
+model = FVDBUNetBase(32, 1).to('cuda')
+soutput = model(sinput)
+```
+
+The output `soutput` will carry gradients during training, and you could train the sparse network accordingly.
+Please find a fully working example at `examples/perception_example.py`. The same network is implemented using `MinkowskiEngine` for reference.
@@ -0,0 +1,641 @@
+# Volume Rendering
+
+In this example we replace `nerfacc`'s acceleration structure with fVDB and hence scale to unbounded scenes:
+
+```python
+import json
+import math
+import os
+from typing import Optional, Tuple, Union
+
+import imageio.v2 as imageio
+import matplotlib.pyplot as plt
+import numpy as np
+import polyscope as ps
+import torch
+import tqdm
+
+from fvdb import GridBatch
+from fvdb import volume_render
+
+TensorPair = Tuple[torch.Tensor, torch.Tensor]
+TensorTriple = Tuple[torch.Tensor, torch.Tensor, torch.Tensor]
+
+class _TruncExp(torch.autograd.Function):
+    @staticmethod
+    def forward(ctx, x):
+        ctx.save_for_backward(x)
+        return torch.exp(x)
+
+    @staticmethod
+    def backward(ctx, dL_dout):
+        x = ctx.saved_tensors[0]
+        return dL_dout * torch.exp(x.clamp(-15, 15))
+
+# SH
+MAX_SH_BASIS = 10
+def eval_sh_bases(basis_dim : int, dirs : torch.Tensor):
+    """
+    Evaluate spherical harmonics bases at unit directions,
+    without taking linear combination.
+    At each point, the final result may the be
+    obtained through simple multiplication.
+
+    :param basis_dim: int SH basis dim. Currently, 1-25 square numbers supported
+    :param dirs: torch.Tensor (..., 3) unit directions
+
+    :return: torch.Tensor (..., basis_dim)
+    """
+    SH_C0 = 0.28209479177387814
+    SH_C1 = 0.4886025119029199
+    SH_C2 = [
+        1.0925484305920792,
+        -1.0925484305920792,
+        0.31539156525252005,
+        -1.0925484305920792,
+        0.5462742152960396
+    ]
+    SH_C3 = [
+        -0.5900435899266435,
+        2.890611442640554,
+        -0.4570457994644658,
+        0.3731763325901154,
+        -0.4570457994644658,
+        1.445305721320277,
+        -0.5900435899266435
+    ]
+    SH_C4 = [
+        2.5033429417967046,
+        -1.7701307697799304,
+        0.9461746957575601,
+        -0.6690465435572892,
+        0.10578554691520431,
+        -0.6690465435572892,
+        0.47308734787878004,
+        -1.7701307697799304,
+        0.6258357354491761,
+    ]
+    result = torch.empty((*dirs.shape[:-1], basis_dim), dtype=dirs.dtype, device=dirs.device)
+    result[..., 0] = SH_C0
+    if basis_dim > 1:
+        x, y, z = dirs.unbind(-1)
+        result[..., 1] = -SH_C1 * y;
+        result[..., 2] = SH_C1 * z;
+        result[..., 3] = -SH_C1 * x;
+        if basis_dim > 4:
+            xx, yy, zz = x * x, y * y, z * z
+            xy, yz, xz = x * y, y * z, x * z
+            result[..., 4] = SH_C2[0] * xy;
+            result[..., 5] = SH_C2[1] * yz;
+            result[..., 6] = SH_C2[2] * (2.0 * zz - xx - yy);
+            result[..., 7] = SH_C2[3] * xz;
+            result[..., 8] = SH_C2[4] * (xx - yy);
+
+            if basis_dim > 9:
+                result[..., 9] = SH_C3[0] * y * (3 * xx - yy);
+                result[..., 10] = SH_C3[1] * xy * z;
+                result[..., 11] = SH_C3[2] * y * (4 * zz - xx - yy);
+                result[..., 12] = SH_C3[3] * z * (2 * zz - 3 * xx - 3 * yy);
+                result[..., 13] = SH_C3[4] * x * (4 * zz - xx - yy);
+                result[..., 14] = SH_C3[5] * z * (xx - yy);
+                result[..., 15] = SH_C3[6] * x * (xx - 3 * yy);
+
+                if basis_dim > 16:
+                    result[..., 16] = SH_C4[0] * xy * (xx - yy);
+                    result[..., 17] = SH_C4[1] * yz * (3 * xx - yy);
+                    result[..., 18] = SH_C4[2] * xy * (7 * zz - 1);
+                    result[..., 19] = SH_C4[3] * yz * (7 * zz - 3);
+                    result[..., 20] = SH_C4[4] * (zz * (35 * zz - 30) + 3);
+                    result[..., 21] = SH_C4[5] * xz * (7 * zz - 3);
+                    result[..., 22] = SH_C4[6] * (xx - yy) * (7 * zz - 1);
+                    result[..., 23] = SH_C4[7] * xz * (xx - 3 * yy);
+                    result[..., 24] = SH_C4[8] * (xx * (xx - 3 * yy) - yy * (3 * xx - yy));
+    return result
+
+def speherical_harmonics(deg: int, sh: torch.Tensor, dirs: torch.Tensor):
+    C0 = 0.28209479177387814
+    C1 = 0.4886025119029199
+    C2 = [
+        1.0925484305920792,
+        -1.0925484305920792,
+        0.31539156525252005,
+        -1.0925484305920792,
+        0.5462742152960396
+    ]
+    C3 = [
+        -0.5900435899266435,
+        2.890611442640554,
+        -0.4570457994644658,
+        0.3731763325901154,
+        -0.4570457994644658,
+        1.445305721320277,
+        -0.5900435899266435
+    ]
+    C4 = [
+        2.5033429417967046,
+        -1.7701307697799304,
+        0.9461746957575601,
+        -0.6690465435572892,
+        0.10578554691520431,
+        -0.6690465435572892,
+        0.47308734787878004,
+        -1.7701307697799304,
+        0.6258357354491761,
+    ]
+
+    # sh is a tensor of shape [N, C, (deg+1)**2]
+    # dirs is a tensor of shape [N, 3]
+    assert 0 <= deg <= 4
+    assert (deg + 1) ** 2 == sh.shape[-1]
+    # C = sh.shape[-2]
+
+    result = C0 * sh[..., 0]
+    if deg > 0:
+        x, y, z = dirs[..., 0:1], dirs[..., 1:2], dirs[..., 2:3]
+        result = (result -
+                C1 * y * sh[..., 1] +
+                C1 * z * sh[..., 2] -
+                C1 * x * sh[..., 3])
+        if deg > 1:
+            xx, yy, zz = x * x, y * y, z * z
+            xy, yz, xz = x * y, y * z, x * z
+            result = (result +
+                    C2[0] * xy * sh[..., 4] +
+                    C2[1] * yz * sh[..., 5] +
+                    C2[2] * (2.0 * zz - xx - yy) * sh[..., 6] +
+                    C2[3] * xz * sh[..., 7] +
+                    C2[4] * (xx - yy) * sh[..., 8])
+
+            if deg > 2:
+                result = (result +
+                        C3[0] * y * (3 * xx - yy) * sh[..., 9] +
+                        C3[1] * xy * z * sh[..., 10] +
+                        C3[2] * y * (4 * zz - xx - yy)* sh[..., 11] +
+                        C3[3] * z * (2 * zz - 3 * xx - 3 * yy) * sh[..., 12] +
+                        C3[4] * x * (4 * zz - xx - yy) * sh[..., 13] +
+                        C3[5] * z * (xx - yy) * sh[..., 14] +
+                        C3[6] * x * (xx - 3 * yy) * sh[..., 15])
+                if deg > 3:
+                    result = (result + C4[0] * xy * (xx - yy) * sh[..., 16] +
+                            C4[1] * yz * (3 * xx - yy) * sh[..., 17] +
+                            C4[2] * xy * (7 * zz - 1) * sh[..., 18] +
+                            C4[3] * yz * (7 * zz - 3) * sh[..., 19] +
+                            C4[4] * (zz * (35 * zz - 30) + 3) * sh[..., 20] +
+                            C4[5] * xz * (7 * zz - 3) * sh[..., 21] +
+                            C4[6] * (xx - yy) * (7 * zz - 1) * sh[..., 22] +
+                            C4[7] * xz * (xx - 3 * yy) * sh[..., 23] +
+                            C4[8] * (xx * (xx - 3 * yy) - yy * (3 * xx - yy)) * sh[..., 24])
+    return result
+
+
+
+def compute_psnr(rgb_gt: torch.Tensor, rgb_est: torch.Tensor) ->  torch.Tensor:
+    x = torch.mean((rgb_gt - rgb_est)**2)
+    return -10. * torch.log10(x)
+
+
+def nerf_matrix_to_ngp(pose: np.ndarray, scale: float = 0.33, offset: Union[tuple, list, torch.Tensor] = (0, 0, 0)) -> np.ndarray:
+    new_pose = np.array([
+        [pose[1, 0], -pose[1, 1], -pose[1, 2], pose[1, 3] * scale + offset[0]],
+        [pose[2, 0], -pose[2, 1], -pose[2, 2], pose[2, 3] * scale + offset[1]],
+        [pose[0, 0], -pose[0, 1], -pose[0, 2], pose[0, 3] * scale + offset[2]],
+        [0, 0, 0, 1],], dtype=np.float32)
+
+    return new_pose
+
+
+def get_rays(pose: torch.Tensor, intrinsic: torch.Tensor, H: int, W: int, depth: Optional[torch.Tensor] = None) -> TensorTriple:
+    fx, fy, cx, cy = intrinsic
+
+    i, j = torch.meshgrid(torch.linspace(0, W-1, W, device='cpu'), torch.linspace(0, H-1, H, device='cpu'), indexing='ij')
+    i = i.t().reshape([1, H*W]).expand([1, H*W]) + 0.5
+    j = j.t().reshape([1, H*W]).expand([1, H*W]) + 0.5
+    zs = torch.ones_like(i)
+    xs = (i - cx) / fx * zs
+    ys = (j - cy) / fy * zs
+    directions = torch.cat((xs.reshape(-1,1), ys.reshape(-1,1), zs.reshape(-1,1)), dim=-1)
+
+    # compute distances
+    if depth is not None:
+        dist = torch.norm(directions * depth[:,None], dim=-1, keepdim=True)
+    else:
+        dist = torch.empty([])
+
+    directions = directions / torch.norm(directions, dim=-1, keepdim=True)
+    rays_d = (pose[:3,:3] @ directions.transpose(0,1)).transpose(0,1)
+
+    rays_o = pose[:3, 3] # [3]
+    rays_o = rays_o[None, :].expand_as(rays_d) # [N, 3]
+
+    return rays_o.squeeze(), rays_d.squeeze(), dist.squeeze()
+
+
+class NeRFDataset:
+    def __init__(self, root_path: str = 'data/lego/', scale: float = 1.0, num_rays: int = 4096, mode: str = 'train'):
+        super().__init__()
+
+        self.root_path  = root_path
+        self.scale = scale
+        self.num_rays = num_rays
+        self.mode = mode
+
+        with open(os.path.join(self.root_path, f'transforms_{self.mode}.json'), 'r', encoding='utf-8') as f:
+            transform = json.load(f)
+
+        # read images
+        frames = transform["frames"]
+        self.n_frames = len(frames)
+
+        # Read the intrinsics
+        image = imageio.imread(os.path.join(self.root_path, frames[0]['file_path'] + '.png')) # [H, W, 3] o [H, W, 4]
+        self.H, self.W  = image.shape[:2]
+        fl_x = fl_y = self.W / (2 * np.tan(transform['camera_angle_x'] / 2)) if 'camera_angle_x' in transform else None
+        cx = (transform['cx']) if 'cx' in transform else (self.W / 2)
+        cy = (transform['cy']) if 'cy' in transform else (self.H / 2)
+        self.intrinsics = np.array([fl_x, fl_y, cx, cy])
+
+        self.rays = []
+        self.rgbs = []
+        self.depths = []
+        self.poses = []
+        self.pc = []
+        self.pc_rgbs = []
+
+        for f in tqdm.tqdm(frames, desc=f'Loading {self.mode} data'):
+            f_path = os.path.join(self.root_path, f['file_path'] + '.png')
+            pose = nerf_matrix_to_ngp(np.array(f['transform_matrix'], dtype=np.float32), scale=self.scale)
+            image = imageio.imread(f_path) / 255.0 # [H, W, 3] o [H, W, 4]
+            depth = None
+
+            if self.mode == 'train':
+                f_path_depth = os.path.join(self.root_path, f['file_path'] + '_depth.npy')
+                depth = np.load(f_path_depth).reshape(-1)
+
+            ray_o, ray_d, depth = get_rays(torch.from_numpy(pose), torch.from_numpy(self.intrinsics), self.H, self.W, depth)
+
+            # Scale the depth
+            depth_mask = depth < 1000
+            depth *= scale
+
+            rgbs = torch.from_numpy(image).reshape(self.H * self.W, -1)
+            self.poses.append(pose)
+            self.rays.append(torch.cat([ray_o, ray_d], 1))
+            self.rgbs.append(rgbs)
+
+            if self.mode == 'train':
+                self.depths.append(depth)
+                self.pc.append(ray_o[depth_mask, :3] + ray_d[depth_mask, :3] * depth[depth_mask,None])
+                self.pc_rgbs.append(rgbs[depth_mask,:3])
+
+        self.rays = torch.vstack(self.rays)
+        self.rgbs = torch.vstack(self.rgbs)
+
+        if self.mode == 'train':
+            self.depths = torch.cat(self.depths) # Note that depth denotes the distance along the ray
+            self.pc = torch.vstack(self.pc)
+            self.pc_rgbs = torch.vstack(self.pc_rgbs)
+
+    def get_point_cloud(self, downsample_ratio: float = 1.0, return_color: bool = False) -> Union[torch.Tensor, TensorPair]:
+        if self.mode == 'train':
+            assert isinstance(self.pc, torch.Tensor)
+            if return_color:
+                assert isinstance(self.pc_rgbs, torch.Tensor)
+                dri = int(1 / downsample_ratio)
+                pts = self.pc[::dri, :]
+                rgb = self.pc_rgbs[::dri, :]
+                return pts, rgb
+            return self.pc[::int(1/downsample_ratio),:]
+        else:
+            raise ValueError('Only training data has depth information!')
+
+    def __len__(self):
+        if self.mode == 'train':
+            return 1000
+        else:
+            return self.n_frames
+
+
+    def __getitem__(self, idx):
+        # raise an error to now iterate in infinity
+        if idx >= len(self): raise IndexError
+
+        if self.mode == 'train':
+            assert isinstance(self.rays, torch.Tensor)
+            idxs = np.random.choice(self.rays.shape[0], self.num_rays)
+            return  {'rays_o': self.rays[idxs,:3],
+                     'rays_d': self.rays[idxs,3:6],
+                     'rgba': self.rgbs[idxs],
+                     'depth': self.depths[idxs],
+                     'idxs': idxs}
+
+        else:
+            # raise an error to now iterate in infinity
+            if idx >= len(self): raise IndexError
+            assert isinstance(self.rays, torch.Tensor)
+            assert isinstance(self.rgbs, torch.Tensor)
+            return  {'rays_o': self.rays[idx * self.W * self.H : (idx + 1) * self.W * self.H, :3],
+                     'rays_d': self.rays[idx * self.W * self.H : (idx + 1) * self.W * self.H, 3:6],
+                     'rgba': self.rgbs[idx * self.W * self.H : (idx + 1) * self.W * self.H,],
+                     'depth': None}
+
+
+def make_ray_grid(origin, nrays, minb=(-0.45, -0.45), maxb=(0.45, 0.45),
+                  device: Union[str, torch.device]='cpu', dtype=torch.float32):
+
+    ray_o = torch.tensor([origin] * nrays**2) #+ p.mean(0, keepdim=True)
+    ray_d = torch.from_numpy(
+        np.stack([a.ravel() for a in
+                  np.mgrid[minb[0]:maxb[0]:nrays*1j,
+                           minb[1]:maxb[1]:nrays*1j]] +
+                  [np.ones(nrays**2)], axis=-1).astype(np.float32))
+    ray_d /= torch.norm(ray_d, dim=-1, keepdim=True)
+
+    ray_o, ray_d = ray_o.to(device).to(dtype), ray_d.to(device).to(dtype)
+
+    return ray_o, ray_d
+
+
+def evaluate_density_and_color(dual_grid: GridBatch, sh_features: torch.Tensor, o_features: torch.Tensor,
+                               ray_d: torch.Tensor, pts: torch.Tensor) -> TensorPair:
+
+    pt_features = dual_grid.sample_trilinear(pts, sh_features.view(sh_features.shape[0], -1)).jdata.view(pts.shape[0], 3, 9)
+    pt_o_features = dual_grid.sample_trilinear(pts, o_features.unsqueeze(-1)).jdata.squeeze(-1)
+    return _TruncExp.apply(pt_o_features), torch.sigmoid(speherical_harmonics(2, pt_features, ray_d))
+
+
+def render(primal_grid: GridBatch, dual_grid: GridBatch,
+           sh_features: torch.Tensor, o_features: torch.Tensor,
+           ray_o: torch.Tensor, ray_d: torch.Tensor,
+           tmin: torch.Tensor, tmax: torch.Tensor, step_size: float,
+           t_threshold: float = 0.0, chunk: bool = False) -> TensorTriple:
+
+    pack_info, ray_idx, ray_intervals = \
+        primal_grid.uniform_ray_samples(ray_o, ray_d, tmin, tmax, step_size)
+
+    ray_t = ray_intervals.jdata.mean(1)
+    ray_delta_t = (ray_intervals.jdata[:, 1] - ray_intervals.jdata[:, 0]).contiguous()
+    ray_pts = ray_o[ray_idx.jdata] + ray_t[:, None] * ray_d[ray_idx.jdata]
+
+    if chunk:
+        ray_density = []
+        ray_color = []
+        ray_d = ray_d[ray_idx.jdata]
+        chunk_size = 400000
+
+        for i in range(ray_d.shape[0]//chunk_size + 1):
+            ray_density_chunk, ray_color_chunk = evaluate_density_and_color(dual_grid, sh_features, o_features,
+                                                                            ray_d[i*chunk_size:(i+1)*chunk_size, :],
+                                                                            ray_pts[i*chunk_size:(i+1)*chunk_size, :])
+
+            ray_density.append(ray_density_chunk)
+            ray_color.append(ray_color_chunk)
+
+        ray_density = torch.cat(ray_density, 0)
+        ray_color = torch.vstack(ray_color)
+    else:
+        ray_density, ray_color = evaluate_density_and_color(dual_grid, sh_features, o_features,
+                                                            ray_d[ray_idx.jdata], ray_pts)
+
+    # Do the volume rendering
+    # print(ray_density.shape, ray_color.shape, ray_delta_t.shape, ray_t.shape, pack_info.jdata.shape)
+    rgb, depth, opacity, _, _ = volume_render(ray_density, ray_color, ray_delta_t,
+                                                         ray_t, pack_info.jdata, t_threshold)
+
+    return rgb, depth, opacity[:, None]
+
+
+def tv_loss(dual_grid: GridBatch, ijk: torch.Tensor, sh_features: torch.Tensor, o_features: torch.Tensor, res) -> TensorPair:
+    nhood = dual_grid.neighbor_indexes(ijk, 1).jdata.view(-1, 3, 3, 3)
+    n_up = nhood[:, 1, 0, 0]
+    n_right = nhood[:, 0, 1, 0]
+    n_front = nhood[:, 0, 0, 1]
+    n_center = nhood[:, 0, 0, 0]
+
+    mask = torch.logical_and(torch.logical_and(n_center != -1, n_up != -1), n_front != -1)
+    fmask = mask.float()
+    n_up_mask, n_right_mask, n_center_mask, n_front_mask = n_up[mask], n_right[mask], n_center[mask], n_front[mask]
+
+    diff_up_sh = (sh_features[n_up_mask] - sh_features[n_center_mask]) / (256.0 / res)
+    diff_right_sh = (sh_features[n_right_mask] - sh_features[n_center_mask]) / (256.0 / res)
+    diff_front_sh = (sh_features[n_front_mask] - sh_features[n_center_mask]) / (256.0 / res)
+
+    diff_up_o = (o_features[n_up] * fmask - o_features[n_center]) / (256.0 / res)
+    diff_right_o = (o_features[n_right] * fmask- o_features[n_center]) / (256.0 / res)
+    diff_front_o = (o_features[n_front] * fmask - o_features[n_center]) / (256.0 / res)
+
+    tv_reg_sh = (diff_up_sh ** 2.0 + diff_right_sh ** 2.0 + diff_front_sh ** 2.0).sum(-1).sum(-1)
+    tv_reg_o = (diff_up_o ** 2.0 + diff_right_o ** 2.0 + diff_front_o ** 2.0)
+    return tv_reg_sh.mean(), tv_reg_o.mean()
+
+
+def main():
+    # Configuration parameters
+    device = torch.device('cuda')
+    dtype = torch.float32
+    # scene_aabb = 1.0
+    starting_resolution = 256
+    resolution = starting_resolution
+    vox_size = (1.0 / resolution, 1.0 / resolution, 1.0 / resolution)
+    vox_origin = (vox_size[0]/2, vox_size[1]/2, vox_size[2]/2)
+    ray_step_size = math.sqrt(3) / 512
+    rays_per_batch = 4096
+    lr_o = 1e-1
+    lr_sh = 1e-2
+
+    plot_every = 2
+    num_epochs = 30
+    bg_color = (0.0, 0.0, 0.0)
+    t_threshold = 1e-5
+
+    # Create the dataset. Assumes there is a file <repository_root>/data/lego_test.h5
+    data_path = os.path.join(os.path.dirname(__file__), "..", "data/lego/")
+
+    if not os.path.exists(data_path):
+        data_url = "https://drive.google.com/drive/folders/1i6qMn-mnPwPEioiNIFMO8QJlTjU0dS1b?usp=share_link"
+        raise RuntimeError(f"You need to download the data at {data_url} "
+                           "into <repository_root>/data "
+                           "in order to run this script")
+
+    train_dataset = NeRFDataset(data_path, scale=0.33, num_rays=rays_per_batch, mode='train')
+    test_dataset = NeRFDataset(data_path, scale=0.33, num_rays=rays_per_batch, mode='test')
+
+
+    # Create a sparse grid used to support features and do ray queries
+    print("Building grid...")
+    primal_grid = GridBatch(device=device)
+    primal_grid.set_from_dense_grid(1, [resolution]*3, [-resolution//2]*3, voxel_sizes=vox_size, voxel_origins=vox_origin)
+    dual_grid = primal_grid # primal_grid.dual_grid()
+
+    print("Done bulding the grid!")
+
+    # Initialize features at the voxel centers
+    xyz = primal_grid.ijk.jdata / (resolution / 2.0)
+    print(xyz.min(0)[0], xyz.max(0)[0])
+
+    sh_features = torch.stack([eval_sh_bases(9, xyz)]*3, dim=1)
+    print(sh_features.shape)
+    sh_features = sh_features.to(device=device, dtype=dtype)
+    o_features = torch.rand(dual_grid.total_voxels)
+    o_features = o_features.to(device=device, dtype=dtype)
+    o_features.requires_grad = True
+    sh_features.requires_grad = True
+
+    # Init optimizer
+    param_group = []
+    param_group.append({'params': o_features, 'lr': lr_o })
+    param_group.append({'params': sh_features, 'lr': lr_sh })
+    # optimizer = torch.optim.Adam(param_group)
+    optimizer = torch.optim.RMSprop(param_group)
+    # scheduler = torch.optim.lr_scheduler.CosineAnnealingLR(optimizer, T_max=30, eta_min=lr/30)
+
+
+    print("Starting training!")
+    ps.init()  # Initialize 3d plotting
+    for epoch in tqdm.trange(num_epochs):
+        if resolution <= starting_resolution:
+            all_ijk = dual_grid.ijk.jdata
+        else:
+            all_ijk = None
+        pbar = tqdm.tqdm(enumerate(train_dataset))  # type: ignore
+        for _, batch in pbar: # type: ignore
+            optimizer.zero_grad()
+            ray_o, ray_d = batch['rays_o'].to(device=device, dtype=dtype), \
+                           batch['rays_d'].to(device=device, dtype=dtype)
+            tmin = torch.zeros(ray_o.shape[0]).to(ray_o)
+            tmax = torch.full_like(tmin, 1e10)
+
+            # Render color and depth along rays
+            rgb, depth, opacity = render(primal_grid, dual_grid, sh_features, o_features, ray_o, ray_d,
+                                         tmin, tmax, ray_step_size, t_threshold=t_threshold)
+
+            rgb = opacity * rgb + (1.0 - opacity) * torch.tensor(bg_color).to(rgb)[None, :]
+
+            # RGB loss
+            rgb_gt = batch['rgba'].to(rgb)[:,:3]
+            loss_rgb = torch.nn.functional.mse_loss(rgb, rgb_gt) # torch.nn.functional.huber_loss(rgb, rgb_gt) / 5.
+
+            # Depth loss
+            # depth_gt = batch['depth'].to(rgb)
+            # depth_mask = depth_gt < np.sqrt(3) # Mask out rays that miss the object
+            # loss_depth = torch.nn.functional.l1_loss(depth[depth_mask], depth_gt[depth_mask]) / 100
+
+            if resolution <= starting_resolution:
+                assert all_ijk is not None
+                random_ijk = all_ijk[torch.randperm(all_ijk.shape[0])[:int(.1*dual_grid.total_voxels)]]
+                tv_reg_sh, tv_reg_o = tv_loss(dual_grid, random_ijk, sh_features, o_features, resolution)
+                tv_reg = 1e-1 * tv_reg_sh + 1e-2 * tv_reg_o
+            else:
+                tv_reg_sh = torch.tensor([0.0]).to(device)
+                tv_reg_o = torch.tensor([0.0]).to(device)
+                tv_reg = torch.tensor([0.0]).to(device)
+            # Total loss to minimize
+            loss = loss_rgb + tv_reg #+ 0.1 * loss_depth
+
+            loss.backward()
+            optimizer.step()
+
+            # Compute current PSNR
+            psnr = compute_psnr(rgb, rgb_gt)
+
+            # Log losses in tqdm progress bar
+            pbar.set_postfix({"Loss": f"{loss.item():.4f}",
+                              "Loss RGB": f"{loss_rgb.item():.4f}",
+                            #   "Loss Depth": f"{loss_depth.item():.4f}",
+                              "Loss TV (sh)": f"{tv_reg_sh.item():.4f}",
+                              "Loss TV (o)": f"{tv_reg_o.item():.4f}",
+                              "PSNR": f"{psnr.item():.2f}"})
+
+        # scheduler.step()
+
+        if epoch % plot_every == 0:
+            with torch.no_grad():
+                torch.cuda.empty_cache()
+
+                grid_res = 512
+                ray_o, ray_d = make_ray_grid((0., 0.15, 1.2), grid_res, device=device, dtype=dtype)
+                ray_d = - ray_d
+                tmin = torch.zeros(ray_o.shape[0]).to(ray_o)
+                tmax = torch.full_like(tmin, 1e10)
+                rgb, depth, opacity = render(primal_grid,dual_grid, sh_features, o_features, ray_o, ray_d,
+                                             tmin, tmax, ray_step_size, t_threshold=t_threshold, chunk=True)
+                rgb = opacity * rgb + (1.0 - opacity) * torch.tensor(bg_color).to(rgb)[None, :]
+
+                rgb_img = rgb.clip(0.0, 1.0).detach().cpu().numpy().reshape([grid_res, grid_res, 3])
+                depth_img = depth.detach().cpu().numpy().reshape([grid_res, grid_res])
+
+                plt.figure()
+                plt.imshow(rgb_img)
+                plt.figure()
+                plt.imshow(depth_img)
+                plt.show()
+
+                ray_v = torch.cat([ray_o, ray_o + ray_d*0.33]).cpu().numpy()
+                ray_e = np.array([[i, i + ray_o.shape[0]] for i in range(ray_o.shape[0])])
+
+                ps.register_curve_network("rays", ray_v, ray_e, radius=0.00002)
+                vox_ijk = primal_grid.ijk.jdata
+                vox_ctrs = primal_grid.grid_to_world(vox_ijk.to(dtype)).jdata
+                vox_density, vox_color = evaluate_density_and_color(dual_grid, sh_features, o_features,
+                                                                    torch.ones_like(vox_ctrs), vox_ctrs)
+
+                # Subdivide
+                if epoch > 0:
+                    sh_features, sub_grid = dual_grid.subdivide(2, sh_features.view(sh_features.shape[0], -1), mask=vox_density > 0.25)
+                    o_features, sub_grid = dual_grid.subdivide(2, o_features.unsqueeze(-1), mask=vox_density > 0.25)
+                    o_features = o_features.jdata.squeeze(-1)
+                    sh_features = sh_features.jdata.reshape(sh_features.rshape[0], 3, -1)
+                    sh_features.requires_grad = True
+                    o_features.requires_grad = True
+                    resolution *= 2.0
+                    ray_step_size /= 2.0
+
+                    print(f"Subdivided grid with {dual_grid.total_voxels} to {sub_grid.total_voxels}")
+                    dual_grid = sub_grid
+                    primal_grid = sub_grid
+
+                camera_origins = []
+                for pose in test_dataset.poses:
+                    camera_origins.append(pose @ np.array([0.0, 0.0, 0.0, 1.0]))
+                camera_origins = np.stack(camera_origins)[:, :3]
+                ps.register_point_cloud("camera origins", camera_origins)
+
+                v, e = primal_grid.viz_edge_network
+                v, e = v.jdata, e.jdata
+                ps.register_curve_network("grid", v.cpu(), e.cpu(), radius=0.0001)
+                pc = ps.register_point_cloud("vox centers", vox_ctrs.cpu(),
+                                            point_render_mode='quad')
+                pc.add_scalar_quantity("density", vox_density.cpu(), enabled=True)
+                pc.add_scalar_quantity("density thresh", (vox_density.cpu() < .25).float(), enabled=True)
+                pc.add_color_quantity("rgb", vox_color.cpu(), enabled=False)
+                ps.show()
+
+
+    print("Starting testing!")
+    pbar = tqdm.tqdm(enumerate(test_dataset))  # type: ignore
+    psnr_test = []
+    for _, batch in pbar: # type: ignore
+        with torch.no_grad():
+            ray_o, ray_d = batch['rays_o'].to(device=device, dtype=dtype), \
+                           batch['rays_d'].to(device=device, dtype=dtype)
+            tmin = torch.zeros(ray_o.shape[0]).to(ray_o)
+            tmax = torch.full_like(tmin, 1e10)
+            rgb_gt = batch['rgba'].to(rgb)[:,:3]  # type: ignore
+
+            # Render color and depth along rays
+            rgb, depth, opacity = render(primal_grid, dual_grid, sh_features, o_features, ray_o, ray_d,
+                                         tmin, tmax, ray_step_size, t_threshold=t_threshold, chunk=True)
+
+            rgb = opacity * rgb + (1.0 - opacity) * torch.tensor(bg_color).to(rgb)[None, :]
+
+            # Compute current PSNR
+            psnr = compute_psnr(rgb, rgb_gt)
+            psnr_test.append(psnr.item())
+            # Log losses in tqdm progress bar
+            pbar.set_postfix({"PSNR": f"{psnr.item():.2f}"})
+
+    print(f"Mean PSNR on the test set across {len(test_dataset)} images: {torch.tensor(psnr_test).mean().item()} ")
+
+if __name__ == "__main__":
+    main()
+
+```
@@ -0,0 +1,23 @@
+diff --git a/CMakeLists.txt b/CMakeLists.txt
+index f726a0a..78dde8b 100644
+--- a/CMakeLists.txt
+++ b/CMakeLists.txt
+@@ -72,7 +72,7 @@
+ #    DEV: static includes blosc.a and blosc.h
+ 
+ 
+-cmake_minimum_required(VERSION 2.8.12)
+cmake_minimum_required(VERSION 3.5)
+ if(NOT CMAKE_VERSION VERSION_LESS 3.3)
+     cmake_policy(SET CMP0063 NEW)
+ endif()
+diff --git a/internal-complibs/zlib-1.2.13/CMakeLists.txt b/internal-complibs/zlib-1.2.13/CMakeLists.txt
+index b412dc7..3efb674 100644
+--- a/internal-complibs/zlib-1.2.13/CMakeLists.txt
+++ b/internal-complibs/zlib-1.2.13/CMakeLists.txt
+@@ -1,4 +1,4 @@
+-cmake_minimum_required(VERSION 2.4.4)
+cmake_minimum_required(VERSION 3.5)
+ set(CMAKE_ALLOW_LOOSE_LOOP_CONSTRUCTS ON)
+ 
+ project(zlib C)
@@ -0,0 +1,38 @@
+name: fvdb_build
+channels:
+  - conda-forge
+  - nodefaults
+dependencies:
+  - blas=*=mkl
+  - ca-certificates
+  - certifi
+  - cmake
+  - cuda-command-line-tools
+  - cuda-compiler
+  - cuda-nvrtc-dev
+  - cuda-nvtx-dev
+  - cuda-version=12.0
+  - cxx-compiler
+  - gcc_linux-64=11
+  - git
+  - gitpython
+  - gxx_linux-64=11
+  - libcublas-dev
+  - libcusolver-dev
+  - libcusparse-dev
+  - libpng
+  - linkify-it-py
+  - make
+  - myst-parser
+  - ninja
+  - numpy
+  - openssl
+  - psutil
+  - python=3.11
+  - pytorch-gpu=2.4.1[build=cuda120*]
+  - scikit-build-core
+  - setuptools>=68.2.2
+  - sphinx_rtd_theme
+  - sphinx>=7.0.0
+  - tqdm
+  - zlib
@@ -0,0 +1,9 @@
+pip >= 23.3.1
+scikit_build_core
+wheel
+GitPython
+ninja
+nvidia-cudnn-cu12
+tqdm
+numpy
+torch
--- a/Show More
+++ b/Show More