NanoVDB: remove multi-GPU diagnostics now that the race is fixed

Drop the per-device stripe/voxel-count fprintf in DistributedPointsToGrid and
the "--repeat until-fail:20" mgpu loop in the CodeBuild buildspec. Both were
temporary scaffolding for diagnosing the shared-tile race, which is now fixed
and covered by dedicated regression tests.

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

.github/workflows/nanovdb.yml

@@ -17,7 +17,7 @@ on:
       - 'fvdb/**'
       - 'pendingchanges/**'
       - '**.md'
-  pull_request:
+  pull_request_target:
     # Add 'ready_for_review' to run CI when a draft PR is marked ready.
     types: [opened, synchronize, reopened, ready_for_review]
     branches:
@@ -49,26 +49,30 @@ jobs:
     # push OR non-draft PR OR matching manual run
     if: |
       (github.event_name == 'push' ||
-       (github.event_name == 'pull_request' && github.event.pull_request.draft == false) ||
+       (github.event_name == 'pull_request_target' && github.event.pull_request.draft == false) ||
        (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' }}
     name: >
       linux-nanovdb:cxx:${{ matrix.config.cxx }}-${{ matrix.config.build }}
+    # Build in the same image used by the AWS CodeBuild GPU runner so the
+    # uploaded build tree is binary/runtime compatible when executed there.
     container:
-      image: aswf/ci-openvdb:${{ matrix.config.image }}
+      image: aswf/ci-openvdb:2024-clang17.2
     env:
       CXX: ${{ matrix.config.cxx }}
     strategy:
       matrix:
         config:
-          - { cxx: g++,     image: '2024-clang17.2', build: 'Release' }
-          - { cxx: g++,     image: '2024-clang17.2', build: 'Debug' }
-          - { cxx: clang++, image: '2024-clang17.2', build: 'Release' }
-          - { cxx: clang++, image: '2024-clang17.2', build: 'Debug' }
+          - { cxx: g++,     label: gcc,   build: 'Release' }
+          - { cxx: g++,     label: gcc,   build: 'Debug' }
+          - { cxx: clang++, label: clang, build: 'Release' }
+          - { cxx: clang++, label: clang, build: 'Debug' }
       fail-fast: false
     steps:
-      - uses: actions/checkout@v3
+      - uses: actions/checkout@v7
+        with:
+          ref: ${{ github.event_name == 'pull_request_target' && github.event.pull_request.head.sha || github.sha }}
       - name: setup_cuda_12
         run: |
           echo "/usr/local/cuda-12/bin" >> $GITHUB_PATH
@@ -78,26 +82,103 @@ jobs:
       - name: build
       # NOTE: CMAKE_POSITION_INDEPENDENT_CODE set to fix default behaviour change in clang 14
       #   https://github.com/AcademySoftwareFoundation/aswf-docker/issues/228
+      # NOTE: CUDA architecture 86 targets the AWS CodeBuild A10G GPUs used by
+      #   the linux-nanovdb-gpu test stage below.
         run: >
           ./ci/build.sh -v
           --build-type=${{ matrix.config.build }}
           --components=core,nano,nanotest,nanoexam,nanobench,nanotool,nanopython,nanopytest
           --cargs=\'
           -DUSE_EXPLICIT_INSTANTIATION=OFF
+          -DOPENVDB_CORE_STATIC=OFF
           -DNANOVDB_USE_CUDA=ON
-          -DCMAKE_CUDA_ARCHITECTURES="80"
+          -DCMAKE_CUDA_ARCHITECTURES="86"
           -DNANOVDB_USE_OPENVDB=ON
           -DCMAKE_INSTALL_PREFIX=`pwd`
           -DUSE_BLOSC=ON
           -DCMAKE_POSITION_INDEPENDENT_CODE=ON
           \'
+      # Run the non-GPU tests here. GPU-only (cuda/mgpu) tests are executed in
+      # the linux-nanovdb-gpu job. This step also generates the shared test
+      # data files (e.g. data/*.nvdb) that the GPU tests consume, so it must
+      # run before the build tree is uploaded.
       - name: test
-        run: cd build && sudo ctest -V -E ".*cuda.*|.*mgpu.*"
+        run: cd build && ctest -V -E ".*cuda.*|.*mgpu.*"
+      - name: package_gpu_build
+        # Record the absolute workspace path so the GPU runner can restore the
+        # build tree to the same location, keeping RPATHs and CTest-generated
+        # absolute paths valid.
+        run: |
+          META="$GITHUB_WORKSPACE/nanovdb-gpu-meta.env"
+          echo "BUILD_WORKSPACE=$GITHUB_WORKSPACE" > "$META"
+          echo "GITHUB_SHA=$GITHUB_SHA" >> "$META"
+          echo "CONFIG=${{ matrix.config.label }}-${{ matrix.config.build }}" >> "$META"
+          cat "$META"
+      - name: upload_gpu_build
+        uses: actions/upload-artifact@v7
+        with:
+          name: nanovdb-gpu-build-${{ matrix.config.label }}-${{ matrix.config.build }}
+          path: |
+            build
+            nanovdb/nanovdb/python/test/TestNanoVDB.py
+            nanovdb-gpu-meta.env
+          retention-days: 1
+          if-no-files-found: error
+
+  linux-nanovdb-gpu:
+    # Execute the CUDA / multi-GPU tests on AWS CodeBuild (4x A10G) using the
+    # prebuilt artifacts from linux-nanovdb. Repository secrets are not shared
+    # with fork PRs, so restrict this job to same-repo events.
+    if: |
+      github.repository == 'AcademySoftwareFoundation/openvdb' &&
+      (github.event_name == 'push' ||
+       (github.event_name == 'pull_request_target' &&
+        github.event.pull_request.draft == false &&
+        github.event.pull_request.head.repo.full_name == github.repository) ||
+       (github.event_name == 'workflow_dispatch' &&
+        (github.event.inputs.type == 'all' || github.event.inputs.type == 'linux')))
+    needs: linux-nanovdb
+    runs-on: ubuntu-latest
+    name: >
+      linux-nanovdb-gpu:${{ matrix.config.label }}-${{ matrix.config.build }}
+    # GITHUB_TOKEN needs actions:read to fetch the workflow-run artifacts.
+    permissions:
+      contents: read
+      actions: read
+    strategy:
+      matrix:
+        config:
+          - { label: gcc,   build: 'Release' }
+          - { label: gcc,   build: 'Debug' }
+          - { label: clang, build: 'Release' }
+          - { label: clang, build: 'Debug' }
+      fail-fast: false
+    env:
+      # Consumed by buildspec.yml on the CodeBuild side.
+      ARTIFACT_NAME: nanovdb-gpu-build-${{ matrix.config.label }}-${{ matrix.config.build }}
+      GH_TOKEN: ${{ github.token }}
+    steps:
+      - name: Configure AWS Credentials
+        uses: aws-actions/configure-aws-credentials@v6
+        with:
+          aws-access-key-id: ${{ secrets.CODEBUILD_ID }}
+          aws-secret-access-key: ${{ secrets.CODEBUILD_SECRET }}
+          aws-region: us-west-2
+      - name: Run CodeBuild GPU tests
+        uses: aws-actions/aws-codebuild-run-build@v1
+        with:
+          project-name: OpenVDB
+          buildspec-override: ci/buildspec_nanovdb.yml
+          image-override: aswf/ci-openvdb:2024-clang17.2
+          # GITHUB_* variables (GITHUB_RUN_ID, GITHUB_REPOSITORY, ...) are
+          # forwarded automatically. The job-scoped Actions runtime values let
+          # CodeBuild upload report.tgz directly to this workflow run.
+          env-vars-for-codebuild: ARTIFACT_NAME,GH_TOKEN,ACTIONS_RUNTIME_TOKEN,ACTIONS_RESULTS_URL
 
   windows-nanovdb:
     if: |
       (github.event_name == 'push' ||
-       (github.event_name == 'pull_request' && github.event.pull_request.draft == false) ||
+       (github.event_name == 'pull_request_target' && github.event.pull_request.draft == false) ||
        (github.event_name == 'workflow_dispatch' &&
         (github.event.inputs.type == 'all' || github.event.inputs.type == 'win')))
     runs-on: ${{ (github.repository_owner == 'AcademySoftwareFoundation' && 'windows-2022-8c-32g-300h') || 'windows-latest' }}
@@ -108,7 +189,9 @@ jobs:
     strategy:
       fail-fast: false
     steps:
-    - uses: actions/checkout@v3
+    - uses: actions/checkout@v7
+      with:
+        ref: ${{ github.event_name == 'pull_request_target' && github.event.pull_request.head.sha || github.sha }}
     - name: path
       run: |
         # note: system path must be modified in a previous step to it's use
@@ -145,7 +228,7 @@ jobs:
   macos-nanovdb:
     if: |
       (github.event_name == 'push' ||
-       (github.event_name == 'pull_request' && github.event.pull_request.draft == false) ||
+       (github.event_name == 'pull_request_target' && github.event.pull_request.draft == false) ||
        (github.event_name == 'workflow_dispatch' &&
         (github.event.inputs.type == 'all' || github.event.inputs.type == 'mac')))
     runs-on: ${{ matrix.config.runner }}
@@ -158,7 +241,9 @@ jobs:
           - { runner: 'macos-14', cxx: 'clang++', build: 'Debug'   }
       fail-fast: false
     steps:
-      - uses: actions/checkout@v3
+      - uses: actions/checkout@v7
+        with:
+          ref: ${{ github.event_name == 'pull_request_target' && github.event.pull_request.head.sha || github.sha }}
       - name: install_deps
         run: |
           ./ci/install_macos.sh
@@ -175,21 +260,23 @@ jobs:
   nanovdb-lite:
     if: |
       (github.event_name == 'push' ||
-       (github.event_name == 'pull_request' && github.event.pull_request.draft == false) ||
+       (github.event_name == 'pull_request_target' && github.event.pull_request.draft == false) ||
        (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' }}
     container:
       image: aswf/ci-openvdb:2024-clang17.2
     steps:
-      - uses: actions/checkout@v3
+      - uses: actions/checkout@v7
+        with:
+          ref: ${{ github.event_name == 'pull_request_target' && github.event.pull_request.head.sha || github.sha }}
       - name: install_gtest
         run: ./ci/install_gtest.sh latest
       - name: build_and_test
         run: |
           cd nanovdb/nanovdb
-          sudo mkdir .build
+          mkdir .build
           cd .build
-          sudo cmake -DUSE_EXPLICIT_INSTANTIATION=OFF -DNANOVDB_BUILD_UNITTESTS=ON -DNANOVDB_USE_OPENVDB=OFF -DNANOVDB_USE_CUDA=OFF ../
-          sudo make -j8 install
-          sudo ctest -V
+          cmake -DUSE_EXPLICIT_INSTANTIATION=OFF -DNANOVDB_BUILD_UNITTESTS=ON -DNANOVDB_USE_OPENVDB=OFF -DNANOVDB_USE_CUDA=OFF -DCMAKE_INSTALL_PREFIX="$(pwd)/install" ../
+          make -j8 install
+          ctest -V

ci/buildspec_nanovdb.yml

@@ -0,0 +1,208 @@
+version: 0.2
+
+# Runs the NanoVDB CUDA / multi-GPU unit tests on AWS CodeBuild (4x A10G).
+# This spec does NOT build the project. Instead it downloads the prebuilt
+# build tree produced by the matching GitHub Actions run (see the
+# linux-nanovdb / linux-nanovdb-gpu jobs in .github/workflows/nanovdb.yml),
+# restores it at its original absolute path, and executes the GPU and Python tests.
+#
+# Required environment variables (provided by aws-codebuild-run-build):
+#   ARTIFACT_NAME     - name of the GitHub Actions artifact to fetch
+#   GH_TOKEN          - ephemeral GITHUB_TOKEN with actions:read
+#   GITHUB_REPOSITORY - owner/repo (auto-forwarded GITHUB_* variable)
+#   GITHUB_RUN_ID     - workflow run that produced the artifact (auto-forwarded)
+#   ACTIONS_RUNTIME_TOKEN - job-scoped token used to upload the report
+#   ACTIONS_RESULTS_URL   - GitHub Actions artifact service endpoint
+
+env:
+  variables:
+    CUDA_HOME: /usr/local/cuda-12
+    DIAGNOSTIC_GIST_ID: c082e01bf80cd88d045e1e61bd78d43c
+    DIAGNOSTIC_GIST_REVISION: dcd46d4e3a65c5d36e1cde96ade0c102464d8a45
+
+phases:
+  install:
+    commands:
+      # Rocky Linux 8 image: tools required to fetch and unpack artifacts and
+      # clone the multi-GPU diagnostic gist.
+      - dnf install -y jq unzip tree 'dnf-command(config-manager)'
+      - dnf config-manager --add-repo https://cli.github.com/packages/rpm/gh-cli.repo
+      - dnf install -y gh
+      - dnf module reset -y nodejs
+      - dnf module enable -y nodejs:20
+      - dnf install -y nodejs
+      - npm install --prefix /tmp/actions-artifact-uploader --omit=dev --no-audit --no-fund @actions/artifact@6.2.1
+
+  pre_build:
+    commands:
+      # Confirm the GPUs are visible (expect 4x A10G).
+      - nvidia-smi
+      - nvidia-smi topo -m
+      - |
+        echo "=== IOMMU diagnostics ==="
+
+        kernel_cmdline=$(cat /proc/cmdline)
+        echo "Kernel command line:"
+        echo "${kernel_cmdline}"
+
+        iommu_args=$(echo "${kernel_cmdline}" |
+          grep -Eo '(^|[[:space:]])(iommu|amd_iommu|intel_iommu)=[^[:space:]]+' |
+          xargs || true)
+        if [ -n "${iommu_args}" ]; then
+          echo "Explicit IOMMU kernel arguments: ${iommu_args}"
+        else
+          echo "No explicit IOMMU kernel arguments found"
+        fi
+
+        if dmesg_output=$(dmesg 2>&1); then
+          echo "IOMMU kernel messages:"
+          echo "${dmesg_output}" |
+            grep -Ei 'DMAR|IOMMU|AMD-Vi|amd_iommu|intel_iommu' ||
+            echo "No IOMMU kernel messages found"
+
+          echo "IOMMU passthrough messages:"
+          echo "${dmesg_output}" |
+            grep -Ei 'iommu.*(pt|passthrough)|default domain type.*passthrough' ||
+            echo "No IOMMU passthrough messages found"
+        else
+          echo "dmesg unavailable in this container: ${dmesg_output}"
+        fi
+
+        if compgen -G '/sys/kernel/iommu_groups/*/devices/*' >/dev/null; then
+          echo "Populated IOMMU groups are visible"
+          group_count=0
+          typed_group_count=0
+          for group in /sys/kernel/iommu_groups/*; do
+            [ -d "${group}" ] || continue
+            group_count=$((group_count + 1))
+            if [ -r "${group}/type" ]; then
+              typed_group_count=$((typed_group_count + 1))
+              echo "IOMMU group ${group##*/} type: $(cat "${group}/type")"
+            fi
+          done
+          echo "Visible IOMMU group count: ${group_count}"
+          if [ "${typed_group_count}" -eq 0 ]; then
+            echo "Per-group IOMMU domain types are not exposed by this kernel"
+          fi
+        else
+          echo "No populated IOMMU groups are visible"
+        fi
+
+        if compgen -G '/sys/class/iommu/*' >/dev/null; then
+          echo "IOMMU devices visible in /sys/class/iommu:"
+          for iommu in /sys/class/iommu/*; do
+            echo "  ${iommu##*/}"
+          done
+        else
+          echo "No IOMMU devices are visible in /sys/class/iommu"
+        fi
+      # Locate, download and restore the prebuilt build tree.
+      - 'echo "Fetching artifact ${ARTIFACT_NAME} from run ${GITHUB_RUN_ID}"'
+      - |
+        set -euo pipefail
+        api="https://api.github.com/repos/${GITHUB_REPOSITORY}"
+        artifacts_json=$(curl -fsSL \
+          -H "Authorization: Bearer ${GH_TOKEN}" \
+          -H "Accept: application/vnd.github+json" \
+          -H "X-GitHub-Api-Version: 2022-11-28" \
+          "${api}/actions/runs/${GITHUB_RUN_ID}/artifacts?per_page=100")
+        count=$(echo "${artifacts_json}" | jq --arg n "${ARTIFACT_NAME}" \
+          '[.artifacts[] | select(.name==$n)] | length')
+        if [ "${count}" != "1" ]; then
+          echo "ERROR: expected exactly 1 artifact named '${ARTIFACT_NAME}', found ${count}"
+          exit 1
+        fi
+        artifact_id=$(echo "${artifacts_json}" | jq -r --arg n "${ARTIFACT_NAME}" \
+          '.artifacts[] | select(.name==$n) | .id')
+        echo "Resolved artifact id ${artifact_id}"
+        # The API returns a short-lived (60s) redirect to blob storage. Resolve
+        # the redirect first, then download without the Authorization header so
+        # it is not forwarded to the storage host.
+        download_url=$(curl -sS -o /dev/null -w '%{redirect_url}' \
+          -H "Authorization: Bearer ${GH_TOKEN}" \
+          -H "Accept: application/vnd.github+json" \
+          "${api}/actions/artifacts/${artifact_id}/zip")
+        curl -fsSL -o /tmp/nanovdb-gpu-build.zip "${download_url}"
+      - rm -rf /tmp/nanovdb-gpu-artifact
+      - mkdir -p /tmp/nanovdb-gpu-artifact
+      - unzip -q /tmp/nanovdb-gpu-build.zip -d /tmp/nanovdb-gpu-artifact
+      - |
+        set -euo pipefail
+        set -a
+        . /tmp/nanovdb-gpu-artifact/nanovdb-gpu-meta.env
+        set +a
+        echo "Restoring build tree and Python tests to ${BUILD_WORKSPACE}"
+        mkdir -p "${BUILD_WORKSPACE}"
+        rm -rf "${BUILD_WORKSPACE}/build"
+        mv /tmp/nanovdb-gpu-artifact/build "${BUILD_WORKSPACE}/build"
+        python_test_dir="${BUILD_WORKSPACE}/nanovdb/nanovdb/python/test"
+        mkdir -p "${python_test_dir}"
+        mv /tmp/nanovdb-gpu-artifact/nanovdb/nanovdb/python/test/TestNanoVDB.py \
+          "${python_test_dir}/TestNanoVDB.py"
+
+  build:
+    commands:
+      - |
+        set -euo pipefail
+        set -a
+        . /tmp/nanovdb-gpu-artifact/nanovdb-gpu-meta.env
+        set +a
+        export PATH="${CUDA_HOME}/bin:${PATH}"
+        export LD_LIBRARY_PATH="${CUDA_HOME}/lib64:${BUILD_WORKSPACE}/build/openvdb/openvdb:${BUILD_WORKSPACE}/build/nanovdb/nanovdb:${LD_LIBRARY_PATH:-}"
+        cd "${BUILD_WORKSPACE}/build"
+        ctest -V -R ".*cuda.*|.*mgpu.*|pytest_nanovdb" --output-on-failure
+
+  post_build:
+    commands:
+      - |
+        set -euo pipefail
+        diagnostic_dir="${CODEBUILD_SRC_DIR}/${DIAGNOSTIC_GIST_ID}"
+        rm -rf "${diagnostic_dir}" "${CODEBUILD_SRC_DIR}/report.tgz"
+        cd "${CODEBUILD_SRC_DIR}"
+        gh gist clone "${DIAGNOSTIC_GIST_ID}"
+        # Pin the executable diagnostic code rather than trusting the mutable
+        # head of the public gist on every CI run.
+        git -C "${diagnostic_dir}" checkout --detach "${DIAGNOSTIC_GIST_REVISION}"
+
+        diagnostic_status=0
+        (
+          cd "${diagnostic_dir}"
+          bash ./run_on_runner.sh
+        ) || diagnostic_status=$?
+
+        if [ -d "${diagnostic_dir}/report" ]; then
+          tar czf "${CODEBUILD_SRC_DIR}/report.tgz" \
+            -C "${diagnostic_dir}" report/
+        else
+          echo "ERROR: the diagnostic did not produce report/"
+          exit 1
+        fi
+
+        # Upload directly into the current GitHub Actions job. These values are
+        # short-lived and scoped to this job's artifact service.
+        : "${ACTIONS_RUNTIME_TOKEN:?ACTIONS_RUNTIME_TOKEN was not forwarded}"
+        : "${ACTIONS_RESULTS_URL:?ACTIONS_RESULTS_URL was not forwarded}"
+        cat > /tmp/actions-artifact-uploader/upload.mjs <<'EOF'
+        import {DefaultArtifactClient} from '@actions/artifact'
+
+        const reportPath = `${process.env.CODEBUILD_SRC_DIR}/report.tgz`
+        const artifactName = process.env.ARTIFACT_NAME.replace(
+          /^nanovdb-gpu-build-/,
+          'nanovdb-gpu-report-'
+        )
+        const client = new DefaultArtifactClient()
+        const result = await client.uploadArtifact(
+          artifactName,
+          [reportPath],
+          process.env.CODEBUILD_SRC_DIR,
+          {
+            retentionDays: 7,
+            compressionLevel: 0
+          }
+        )
+        console.log(
+          `Uploaded ${artifactName}: id=${result.id}, size=${result.size}, digest=${result.digest}`
+        )
+        EOF
+        node /tmp/actions-artifact-uploader/upload.mjs
+        exit "${diagnostic_status}"

nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh

@@ -37,21 +37,6 @@ private:
     uint32_t* mNodeCounts;
 };
 
-/// @brief Indicator functor that returns 1 if the input value matches the member value, 0 otherwise
-template <typename T, typename std::enable_if<std::is_integral<T>::value>::type* = nullptr>
-struct EqualityIndicator
-{
-    EqualityIndicator(const T* value) : mValue(value) {}
-
-    __hostdev__
-    T operator()(const T& x) const
-    {
-        return x == (*mValue);
-    }
-private:
-    const T* mValue;
-};
-
 /// @brief Implements the merge path binary search algorithm in order to find the median across two sorted input key arrays
 template<typename KeyIteratorIn>
 __device__
@@ -118,34 +103,6 @@ void mergePathKernel(KeyIteratorIn keys1, size_t keys1Count, KeyIteratorIn keys2
     mergePath(keys1, keys1Count, keys2, keys2Count, key1Intervals, key2Intervals, intervalIndex);
 }
 
-/// @brief Extends or shortens the left end of an array interval
-template<typename DistanceIteratorIn, typename CountIteratorOut, typename OffsetIteratorOut>
-__global__
-void leftRebalanceKernel(DistanceIteratorIn leftDistance, DistanceIteratorIn rightDistance, CountIteratorOut leftCount, OffsetIteratorOut leftOffset)
-{
-    if (*leftDistance < *rightDistance) {
-        *leftCount -= *leftDistance;
-    }
-    else {
-        *leftCount += *rightDistance;
-    }
-}
-
-/// @brief Extends or shortens the right end of an array interval
-template<typename DistanceIteratorIn, typename CountIteratorOut, typename OffsetIteratorOut>
-__global__
-void rightRebalanceKernel(DistanceIteratorIn leftDistance, DistanceIteratorIn rightDistance, CountIteratorOut rightCount, OffsetIteratorOut rightOffset)
-{
-    if (*leftDistance < *rightDistance) {
-        *rightCount += *leftDistance;
-        *rightOffset -= *leftDistance;
-    }
-    else {
-        *rightCount -= *rightDistance;
-        *rightOffset += *rightDistance;
-    }
-}
-
 } // namespace kernels
 
 // Define utility macro used to call cub functions that use dynamic temporary storage
@@ -560,8 +517,6 @@ void DistributedPointsToGrid<BuildT>::countNodes(const PtrT coords, size_t coord
     // to reduce overhead.
     std::vector<cudaEvent_t> sortEvents(mDeviceMesh.deviceCount());
     std::vector<cudaEvent_t> runLengthEncodeEvents(mDeviceMesh.deviceCount());
-    std::vector<cudaEvent_t> transformReduceEvents(mDeviceMesh.deviceCount());
-    std::vector<cudaEvent_t> rebalanceEvents(mDeviceMesh.deviceCount());
     std::vector<cudaEvent_t> tilePrefixSumEvents(mDeviceMesh.deviceCount());
     std::vector<cudaEvent_t> voxelCountEvents(mDeviceMesh.deviceCount());
     std::vector<cudaEvent_t> leafCountEvents(mDeviceMesh.deviceCount());
@@ -572,8 +527,6 @@ void DistributedPointsToGrid<BuildT>::countNodes(const PtrT coords, size_t coord
         cudaCheck(cudaSetDevice(deviceId));
         cudaEventCreateWithFlags(&sortEvents[deviceId], cudaEventDisableTiming);
         cudaEventCreateWithFlags(&runLengthEncodeEvents[deviceId], cudaEventDisableTiming);
-        cudaEventCreateWithFlags(&transformReduceEvents[deviceId], cudaEventDisableTiming);
-        cudaEventCreateWithFlags(&rebalanceEvents[deviceId], cudaEventDisableTiming);
         cudaEventCreateWithFlags(&tilePrefixSumEvents[deviceId], cudaEventDisableTiming);
         cudaEventCreateWithFlags(&voxelCountEvents[deviceId], cudaEventDisableTiming);
         cudaEventCreateWithFlags(&leafCountEvents[deviceId], cudaEventDisableTiming);
@@ -626,60 +579,68 @@ void DistributedPointsToGrid<BuildT>::countNodes(const PtrT coords, size_t coord
 
     radixSortAsync(mDeviceMesh, mTempDevicePools, mData->d_keys, mKeys, mData->d_indx, mIndices, coordCount, mIntervals, mStripeOffsets, mStripeCounts, sortEvents.data(), sortEvents.data());
 
-    // For each segment of sorted keys on each device, we count how many of the leftmost key occur past the left boundary of the segment. The same is done for the rightmost key with the right boundary of the segment.
-    auto leftIntervals = mIntervals;
-    auto rightIntervals = mIntervals + mDeviceMesh.deviceCount();
+    // Rebalance the device segments so that a device boundary always coincides
+    // with a change in key value. Because TileKeyFunctor assigns identical keys
+    // to every point that falls in the same upper-node "tile", this aligns the
+    // device ownership boundaries with tile boundaries. Downstream construction
+    // assumes each tile (and therefore each lower node, leaf node, and voxel) is
+    // owned by exactly one device; if a tile straddled a boundary, multiple
+    // devices would concurrently build the same leaf and race on its value mask.
+    //
+    // A single tile can span three or more devices (e.g. one dense leaf whose
+    // points are split evenly across the mesh). Adjusting only adjacent pairs of
+    // boundaries cannot consolidate such a tile because a fully-interior device
+    // lies entirely within it, so we compute the boundaries globally and
+    // monotonically. mKeys is globally sorted at this point, so a tile boundary
+    // is simply a position where mKeys changes. Snapping is performed on the host
+    // over the (small) set of device boundaries; fully-interior devices are left
+    // empty, which the rest of the pipeline already handles.
     for (const auto& [deviceId, stream] : mDeviceMesh) {
         cudaCheck(cudaSetDevice(deviceId));
-
-        auto deviceStripeCount = mStripeCounts[deviceId];
-        auto deviceStripeOffset = mStripeOffsets[deviceId];
-        uint64_t* deviceInputKeys = mKeys + deviceStripeOffset;
-
-        if (deviceId > 0) {
-            cudaStreamWaitEvent(stream, sortEvents[deviceId - 1]);
-            EqualityIndicator<uint64_t> indicator(deviceInputKeys - 1);
-            CUB_LAUNCH(DeviceReduce::TransformReduce, mTempDevicePools[deviceId], stream, deviceInputKeys, rightIntervals + deviceId, deviceStripeCount, ::cuda::std::plus(), indicator, 0);
-        }
-        else {
-            rightIntervals[deviceId] = 0;
-        }
-
-        if (deviceId < static_cast<int>(mDeviceMesh.deviceCount() - 1)) {
-            cudaStreamWaitEvent(stream, sortEvents[deviceId + 1]);
-            EqualityIndicator<uint64_t> indicator(deviceInputKeys + deviceStripeCount);
-            CUB_LAUNCH(DeviceReduce::TransformReduce, mTempDevicePools[deviceId], stream, deviceInputKeys, leftIntervals + deviceId, deviceStripeCount, ::cuda::std::plus(), indicator, 0);
-        }
-        else {
-            leftIntervals[deviceId] = 0;
-        }
-        cudaEventRecord(transformReduceEvents[deviceId], stream);
+        cudaCheck(cudaStreamSynchronize(stream));
     }
 
-    // Rebalance the segments so that a device segment boundary also corresponds to a change in key value. Effectively, this aligns upper node boundaries with device ownership boundaries.
-    for (const auto& [deviceId, stream] : mDeviceMesh) {
-        cudaCheck(cudaSetDevice(deviceId));
-
-        if (deviceId > 0)
-        {
-            cudaStreamWaitEvent(stream, transformReduceEvents[deviceId - 1]);
-            kernels::rightRebalanceKernel<<<1, 1, 0, stream>>>(leftIntervals + deviceId - 1, rightIntervals + deviceId, mStripeCounts + deviceId, mStripeOffsets + deviceId);
+    {
+        const int deviceCount = static_cast<int>(mDeviceMesh.deviceCount());
+        const ptrdiff_t keyCount = static_cast<ptrdiff_t>(coordCount);
+        ptrdiff_t previousBoundary = mStripeOffsets[0]; // device 0 always starts at 0
+        for (int deviceId = 1; deviceId < deviceCount; ++deviceId) {
+            ptrdiff_t boundary = mStripeOffsets[deviceId];
+            if (boundary >= keyCount) {
+                boundary = keyCount;
+            } else if (boundary > previousBoundary && mKeys[boundary] == mKeys[boundary - 1]) {
+                // The even-split boundary falls inside a tile run; find its extent
+                // and snap to whichever end keeps the boundary closest to the
+                // even split without crossing the previous boundary.
+                ptrdiff_t runStart = boundary;
+                while (runStart > previousBoundary && mKeys[runStart - 1] == mKeys[boundary]) --runStart;
+                ptrdiff_t runEnd = boundary;
+                while (runEnd < keyCount && mKeys[runEnd] == mKeys[boundary]) ++runEnd;
+                if (runStart <= previousBoundary) {
+                    boundary = runEnd; // the run reaches the previous device, give the whole tile away
+                } else {
+                    boundary = (boundary - runStart <= runEnd - boundary) ? runStart : runEnd;
+                }
+            }
+            if (boundary < previousBoundary) boundary = previousBoundary;
+            mStripeOffsets[deviceId] = boundary;
+            previousBoundary = boundary;
         }
 
-        if (deviceId < static_cast<int>(mDeviceMesh.deviceCount() - 1))
-        {
-            cudaStreamWaitEvent(stream, transformReduceEvents[deviceId + 1]);
-            kernels::leftRebalanceKernel<<<1, 1, 0, stream>>>(leftIntervals + deviceId, rightIntervals + deviceId + 1, mStripeCounts + deviceId, mStripeOffsets + deviceId);
+        // Recompute the per-device counts from the adjusted, monotonic offsets.
+        for (int deviceId = 0; deviceId < deviceCount; ++deviceId) {
+            const ptrdiff_t nextOffset = (deviceId + 1 < deviceCount)
+                ? mStripeOffsets[deviceId + 1]
+                : keyCount;
+            mStripeCounts[deviceId] = static_cast<size_t>(nextOffset - mStripeOffsets[deviceId]);
         }
-        cudaEventRecord(rebalanceEvents[deviceId], stream);
     }
 
-    // Parallel RLE in order to obtain tiles
+    // Parallel RLE in order to obtain tiles. The device boundaries were finalized
+    // synchronously on the host above, so no per-device rebalance event is needed.
     for (const auto& [deviceId, stream] : mDeviceMesh) {
         cudaCheck(cudaSetDevice(deviceId));
 
-        cudaCheck(cudaEventSynchronize(rebalanceEvents[deviceId]));
-
         auto deviceStripeCount = mStripeCounts[deviceId];
         auto deviceStripeOffset = mStripeOffsets[deviceId];
 
@@ -837,8 +798,6 @@ void DistributedPointsToGrid<BuildT>::countNodes(const PtrT coords, size_t coord
         cudaCheck(cudaSetDevice(deviceId));
         cudaEventDestroy(sortEvents[deviceId]);
         cudaEventDestroy(runLengthEncodeEvents[deviceId]);
-        cudaEventDestroy(transformReduceEvents[deviceId]);
-        cudaEventDestroy(rebalanceEvents[deviceId]);
         cudaEventDestroy(tilePrefixSumEvents[deviceId]);
         cudaEventDestroy(voxelCountEvents[deviceId]);
         cudaEventDestroy(leafCountEvents[deviceId]);
@@ -915,7 +874,7 @@ inline void DistributedPointsToGrid<BuildT>::processGridTreeRoot(const PtrT poin
     } else {
         cudaCheck(cudaMemsetAsync(dst, 0, GridData::MaxNameSize, stream));
     }
-    cudaEventRecord(processGridTreeRootEvent);
+    cudaCheck(cudaEventRecord(processGridTreeRootEvent, stream));
 
     for (const auto& [otherDeviceId, otherStream] : mDeviceMesh) {
         cudaSetDevice(otherDeviceId);

nanovdb/nanovdb/unittest/TestMultiGPU.cu

@@ -10,6 +10,9 @@
 #include <cuda_runtime_api.h>
 #include <gtest/gtest.h>
 #include <thread> // for std::thread
+#include <algorithm> // for std::sort, std::unique
+#include <cstdint>
+#include <vector>
 
 #include <thrust/fill.h>
 #include <thrust/universal_vector.h>
@@ -243,7 +246,7 @@ TEST(TestNanoVDBMultiGPU, DenseLeaf_DistributedCudaPointsToGrid_UnifiedBuffer)
     EXPECT_TRUE(data);
     grid = handle.grid<BuildT>();
     EXPECT_TRUE(grid);
-    EXPECT_TRUE(grid->activeVoxelCount() == 512);
+    EXPECT_EQ(voxelCount, grid->activeVoxelCount());
     EXPECT_EQ(nanovdb::Vec3d(1.0), grid->voxelSize());
 
     cudaCheck(cudaFree(voxels));
@@ -415,3 +418,118 @@ TEST(TestNanoVDBMultiGPU, ManyTiles_DistributedCudaPointsToGrid)
     cudaCheck(cudaFree(voxels));
     cudaSetDevice(current);
 }// ManyTiles_DistributedCudaPointsToGrid
+
+/// @brief Regression test for the multi-GPU shared-tile race. All coordinates
+///        fall inside a single upper-node tile (each upper node spans 4096
+///        voxels per axis, so ijk in [0, 4095] maps to tile (0,0,0)), yet they
+///        populate many lower and leaf nodes. With the default even initial
+///        split this one tile is shared across every GPU, so the builder must
+///        consolidate it onto a single device. A cross-device race in leaf
+///        construction would drop a device's contribution and undercount the
+///        active voxels, which the exact-count assertion below catches
+///        deterministically.
+TEST(TestNanoVDBMultiGPU, SingleUpperNode_DistributedCudaPointsToGrid_UnifiedBuffer)
+{
+    int current = 0;
+    cudaCheck(cudaGetDevice(&current));
+
+    using BufferT = nanovdb::cuda::UnifiedBuffer;
+    using BuildT = nanovdb::ValueOnIndex;
+
+    const size_t inputCount = 1 << 18;// 262144
+    nanovdb::Coord* voxels = nullptr;
+    cudaCheck(cudaMallocManaged(&voxels, inputCount * sizeof(nanovdb::Coord)));
+    std::srand(24680);
+    auto op = [](){ return rand() % 4096; };// stays within a single upper node
+    for (size_t i = 0; i < inputCount; ++i)
+        voxels[i] = nanovdb::Coord(op(), op(), op());
+
+    // Deterministic expected active-voxel count (the input may contain duplicates).
+    std::vector<uint64_t> packed(inputCount);
+    for (size_t i = 0; i < inputCount; ++i)
+        packed[i] = (uint64_t(voxels[i][0]) << 24) | (uint64_t(voxels[i][1]) << 12) | uint64_t(voxels[i][2]);
+    std::sort(packed.begin(), packed.end());
+    const size_t uniqueCount = static_cast<size_t>(std::unique(packed.begin(), packed.end()) - packed.begin());
+
+    nanovdb::cuda::DeviceMesh deviceMesh;
+    nanovdb::tools::cuda::DistributedPointsToGrid<BuildT> converter(deviceMesh);
+    auto handle = converter.getHandle(voxels, inputCount);
+
+    EXPECT_TRUE(handle.deviceData());
+    EXPECT_TRUE(handle.deviceGrid<BuildT>());
+    handle.deviceDownload();
+    auto *grid = handle.grid<BuildT>();
+    EXPECT_TRUE(grid);
+    EXPECT_EQ(nanovdb::Vec3d(1.0), grid->voxelSize());
+    // The input occupies exactly one upper-node tile, so the shared-tile
+    // consolidation path is exercised.
+    EXPECT_EQ(1u, grid->tree().nodeCount(2));
+    // Every unique input voxel must be active exactly once.
+    EXPECT_EQ(static_cast<uint64_t>(uniqueCount), grid->activeVoxelCount());
+
+    nanovdb::util::forEach(0, inputCount, 1, [&](const nanovdb::util::Range1D &r){
+        auto acc = grid->getAccessor();
+        for (size_t i=r.begin(); i!=r.end(); ++i) {
+            const nanovdb::Coord &ijk = voxels[i];
+            EXPECT_TRUE(acc.probeLeaf(ijk)!=nullptr);
+            EXPECT_TRUE(acc.isActive(ijk));
+            EXPECT_TRUE(acc.getValue(ijk) > 0u);
+        }
+    });
+
+    cudaCheck(cudaFree(voxels));
+    cudaSetDevice(current);
+}// SingleUpperNode_DistributedCudaPointsToGrid_UnifiedBuffer
+
+/// @brief Cross-checks the distributed builder against the trusted single-GPU
+///        PointsToGrid on an input that forces a single tile to be split across
+///        devices. Index assignment order may differ between the two builders,
+///        so we compare topology (node counts, active-voxel count) and voxel
+///        occupancy rather than the ValueOnIndex indices themselves.
+TEST(TestNanoVDBMultiGPU, MatchesSingleGpu_DistributedCudaPointsToGrid)
+{
+    int current = 0;
+    cudaCheck(cudaGetDevice(&current));
+
+    using BufferT = nanovdb::cuda::UnifiedBuffer;
+    using BuildT = nanovdb::ValueOnIndex;
+
+    const size_t inputCount = 1 << 17;// 131072, all within a single upper node
+    nanovdb::Coord* voxels = nullptr;
+    cudaCheck(cudaMallocManaged(&voxels, inputCount * sizeof(nanovdb::Coord)));
+    std::srand(1357);
+    auto op = [](){ return rand() % 4096; };
+    for (size_t i = 0; i < inputCount; ++i)
+        voxels[i] = nanovdb::Coord(op(), op(), op());
+
+    nanovdb::cuda::DeviceMesh deviceMesh;
+    nanovdb::tools::cuda::DistributedPointsToGrid<BuildT> converter(deviceMesh);
+    auto distributedHandle = converter.getHandle(voxels, inputCount);
+    distributedHandle.deviceDownload();
+    auto *distributedGrid = distributedHandle.grid<BuildT>();
+    EXPECT_TRUE(distributedGrid);
+
+    cudaSetDevice(current);
+    auto referenceHandle = nanovdb::tools::cuda::voxelsToGrid<BuildT, nanovdb::Coord*, BufferT>(voxels, inputCount);
+    referenceHandle.deviceDownload();
+    auto *referenceGrid = referenceHandle.grid<BuildT>();
+    EXPECT_TRUE(referenceGrid);
+
+    EXPECT_EQ(referenceGrid->activeVoxelCount(), distributedGrid->activeVoxelCount());
+    EXPECT_EQ(referenceGrid->tree().nodeCount(0), distributedGrid->tree().nodeCount(0));
+    EXPECT_EQ(referenceGrid->tree().nodeCount(1), distributedGrid->tree().nodeCount(1));
+    EXPECT_EQ(referenceGrid->tree().nodeCount(2), distributedGrid->tree().nodeCount(2));
+
+    nanovdb::util::forEach(0, inputCount, 1, [&](const nanovdb::util::Range1D &r){
+        auto distributedAcc = distributedGrid->getAccessor();
+        auto referenceAcc = referenceGrid->getAccessor();
+        for (size_t i=r.begin(); i!=r.end(); ++i) {
+            const nanovdb::Coord &ijk = voxels[i];
+            EXPECT_TRUE(referenceAcc.isActive(ijk));
+            EXPECT_TRUE(distributedAcc.isActive(ijk));
+        }
+    });
+
+    cudaCheck(cudaFree(voxels));
+    cudaSetDevice(current);
+}// MatchesSingleGpu_DistributedCudaPointsToGrid

nanovdb/nanovdb/unittest/TestNanoVDB.cu

@@ -791,20 +791,31 @@ TEST(TestNanoVDBCUDA, FewTiles_CudaPointsToGrid)
 }// FewTiles_CudaPointsToGrid
 
 /// @brief Exercises the segmented sort path (>= 32 tiles) in PointsToGrid.
-///        Coordinates in [-8000, 8000] produce 4^3 = 64 upper internal node tiles
+///        Coordinates are placed in 4^3 = 64 upper internal node tiles
 ///        (each upper node covers 4096 voxels per dimension), exceeding the threshold of 32.
+///        Each tile contains 4^3 voxels in one leaf, giving every segment multiple keys to sort
+///        without creating the very large number of lower nodes produced by sparse random voxels.
 TEST(TestNanoVDBCUDA, ManyTiles_CudaPointsToGrid)
 {
     using BuildT = nanovdb::ValueOnIndex;
-    const size_t voxelCount = 1 << 20;// 1048576
     std::vector<nanovdb::Coord> voxels;
-    {
-        voxels.reserve(voxelCount);
-        std::srand(54321);
-        const int max = 8000, min = -max;
-        auto op = [&](){return rand() % (max - min) + min;};
-        while (voxels.size() < voxelCount) voxels.push_back(nanovdb::Coord(op(), op(), op()));
+    voxels.reserve(1 << 12);// 64 upper tiles * 64 voxels per tile = 4096
+    for (int tx = -2; tx < 2; ++tx) {
+        for (int ty = -2; ty < 2; ++ty) {
+            for (int tz = -2; tz < 2; ++tz) {
+                const nanovdb::Coord base(tx * 4096 + 64, ty * 4096 + 64, tz * 4096 + 64);
+                for (int x = 0; x < 4; ++x) {
+                    for (int y = 0; y < 4; ++y) {
+                        for (int z = 0; z < 4; ++z) {
+                            voxels.emplace_back(base + nanovdb::Coord(x, y, z));
+                        }
+                    }
+                }
+            }
+        }
     }
+    const size_t voxelCount = voxels.size();
+    ASSERT_EQ(size_t(1 << 12), voxelCount);
 
     nanovdb::Coord* d_coords;
     cudaCheck(cudaMalloc(&d_coords, voxels.size() * sizeof(nanovdb::Coord)));
@@ -819,6 +830,7 @@ TEST(TestNanoVDBCUDA, ManyTiles_CudaPointsToGrid)
     EXPECT_TRUE(grid);
     EXPECT_TRUE(grid->valueCount() > 0);
     EXPECT_EQ(nanovdb::Vec3d(1.0), grid->voxelSize());
+    EXPECT_EQ(64u, grid->tree().nodeCount(2));
 
     nanovdb::util::forEach(voxels, [&](const nanovdb::util::Range1D &r){
         auto acc = grid->getAccessor();