[copilot mirror] B3: PointsToGrid onto cuda::Buffer

[harrism]

Copilot-review mirror of AcademySoftwareFoundation#2270, based on the B2 branch so only B3's payload shows.

[Copilot]

Pull request overview

Warning

Copilot couldn't run its full agentic review because it didn't start before the timeout. Make sure your repository has a runner available, or add a copilot-code-review.yml file specifying one with the runs-on attribute. See the docs for more details.

This PR migrates NanoVDB’s CUDA PointsToGrid builder scratch/intermediate device allocations to nanovdb::cuda::Buffer (backed by an injectable memory resource), reducing manual allocate/free management and aligning with other GPU builders.

Changes:

• Updated PointsToGrid to own device scratch + intermediate arrays via cuda::Buffer members and refresh raw-pointer views in mData.
• Reworked the countNodes “dx search” loop to reuse a single CUDA event and to free buffers via Buffer::destroy() instead of paired allocate_async/deallocate_async.
• Updated pending change notes to document the PointsToGrid migration.

Reviewed changes

Copilot reviewed 2 out of 2 changed files in this pull request and generated 1 comment.

File	Description
pendingchanges/nanovdb.txt	Documents the PointsToGrid scratch/intermediate allocation change to cuda::Buffer.
nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh	Replaces raw async allocations with cuda::Buffer ownership and updates lifecycle management across the pipeline.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[harrism]

Fixed in 4957582 — the event is now owned by a scope guard, so the too-many-points-per-leaf throw releases it during unwinding, consistent with the buffer ownership this PR introduces.

[Copilot]

Pull request overview

Copilot reviewed 2 out of 2 changed files in this pull request and generated 1 comment.

Suppressed comments (3)

nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh:653

• Calling cudaCheck(...) in a destructor is risky if cudaCheck can throw (it can lead to std::terminate during stack unwinding). Prefer making the destructor non-throwing: call cudaEventDestroy(event) directly (optionally ignoring errors), or adjust cudaCheck usage here to guarantee noexcept behavior.

    struct EventGuard {
        cudaEvent_t event;
        EventGuard() { cudaCheck(cudaEventCreate(&event)); }
        ~EventGuard() { cudaCheck(cudaEventDestroy(event)); }
        EventGuard(const EventGuard&) = delete;
        EventGuard& operator=(const EventGuard&) = delete;
    } eventGuard;

nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh:722

• This manual new[]/delete[] is not exception-safe if any of the CUDA/CUB checks throw between allocation and delete[]. Use an RAII container (e.g., std::vector<uint32_t> or std::unique_ptr<uint32_t[]>) so host memory is released automatically on early exits.

            uint32_t *points_per_tile = new uint32_t[mData.nodeCount[2]];
            cudaCheck(cudaMemcpyAsync(points_per_tile, d_points_per_tile, mData.nodeCount[2]*sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream));
            pointsPerTileScratch.destroy(); d_points_per_tile = nullptr;
            for (uint32_t id = 0, offset = 0; id < mData.nodeCount[2]; ++id) {
                const uint32_t count = points_per_tile[id];
                util::cuda::offsetLambdaKernel<<<numBlocks(count), mNumThreads, 0, mStream>>>(count, offset, VoxelKeyFunctor<BuildT, PtrT>(), mDeviceData, points, id, d_keys, d_indx);
                cudaCheckError();
                CALL_CUBS(DeviceRadixSort::SortPairs, d_keys + offset, mData.d_keys + offset, d_indx + offset, mData.d_indx + offset, count, 0, 36);
                offset += count;
            }
            delete [] points_per_tile;

nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh:408

• ref() does not modify object state and can be marked const (and potentially noexcept) to improve API correctness and allow calling it from const contexts.

    nanovdb::cuda::ResourceRef<ResourceT> ref() { return nanovdb::cuda::ResourceRef<ResourceT>(*mResource); }

[harrism]

Fixed in 2e1e237 — for consistency, not correctness: it compiled because cudaCheckError expands to a braced block, so no semicolon is required, but every other use in the file spells it as a statement. Pre-existing, incidentally (visible at line 1328 of the base).

[harrism]

Copilot review cycles converged (final cycle: no new findings). All accepted changes are on the shared branch and reflected in the upstream PR; threads here document the dismissals. Branch retained.