NanoVDB PointsToGrid::countNodes: Use segmented radix sort for higher tile counts

[swahtz]

This pull request introduces a performance optimization to the voxel key sorting process in the implementation of PointsToGrid. The main improvement is the addition of a bulk segmented sort path for cases with many tiles instead of a serial loop of kernel launches per-tile, which significantly speeds up sorting in large datasets. I found that creating a grid for the Stanford dragon at a voxel size that produced 200 tiles, the speedup to sorting was 19x and for a voxel size producing 6000 tiles, the speedup was 73x. The end-to-end PointsToGrid improvements were 17% for the case with 6,000 tiles. For low tile counts, I found that the performance of segmented radix sort was worse than the original so I include a fallback when tile counts are low.

Performance and algorithm improvements:

• Added a new BulkVoxelKeyFunctor struct and associated kernel launch to efficiently compute voxel keys for all points in a single pass (instead of multiple kernel launches) when the number of tiles exceeds a threshold. This enables a bulk segmented sort path for large tile counts, improving performance for large datasets.
• Modified the sorting logic to choose between bulk segmented sort and serial per-tile sort based on the number of tiles (SEGMENTED_SORT_TILE_THRESHOLD). Bulk sort is used for large tile counts, while the original per-tile sort is retained for small tile counts where the original per-tile sort was faster.

Minor fixes:

• Fixed the setVerbose method to only set the local verbosity variable, removing the flag manipulation for clarity and correctness.

[harrism]

💡 suggestion: ‏Use thrust::lower_bound here with a thrust::seq execution policy to do the same thing more literately and robustly.

[swahtz]

Good idea, thanks. I actually think is upper_bound but the same gist.

[Copilot]

Pull request overview

This PR optimizes nanovdb::tools::cuda::PointsToGrid::countNodes by adding a bulk segmented radix-sort path for voxel-key sorting when the number of tiles is high, while keeping the existing per-tile sorting path for low tile counts to avoid overhead regressions.

Changes:

• Added a bulk voxel-key generation kernel (BulkVoxelKeyFunctor) to compute voxel keys for all points in one launch when tile counts are high.
• Switched sorting to cub::DeviceSegmentedRadixSort::SortPairs for high tile counts using computed per-tile segment offsets, with a threshold-based fallback to the original per-tile loop.
• Simplified setVerbose to only update the local verbosity level.

---

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

[harrism]

Looks like a great optimization. Well done.

[harrism]

👏 praise: ‏Nice!

[matthewdcong]

Is the difference in absolute time for small numbers of tiles large enough to warrant keeping the older path?

[swahtz]

Is the difference in absolute time for small numbers of tiles large enough to warrant keeping the older path?

I was unsure if we should keep the older path too. For small numbers of tiles, the overhead for the segmented sort makes it more expensive. For 100k points and 4 tiles, the new segmented sort is 0.371ms and the old serial per-tile sort is .278ms, a 34% regression (running on my Ada RTX 6000). For small tile counts, do you think it's reasonable to trade that off for code complexity?