NanoVDB: assign each upper-node tile to a single GPU in DistributedPointsToGrid - #2283
Open
swahtz wants to merge 2 commits into
Open
NanoVDB: assign each upper-node tile to a single GPU in DistributedPointsToGrid#2283swahtz wants to merge 2 commits into
swahtz wants to merge 2 commits into
Conversation
…intsToGrid The device segment boundaries produced by the initial even split were rebalanced by adjusting only adjacent pairs of boundaries. That cannot consolidate an upper-node tile spanning three or more GPUs, since a fully-interior device lies entirely within the tile. Downstream construction assumes each tile - and therefore each lower node, leaf node and voxel - is owned by exactly one device, so multiple devices built the same leaf concurrently and raced on its value mask, silently dropping active voxels. Replace the pairwise GPU rebalance with a global, monotonic snap performed on the host over the (small) set of device boundaries. The keys are globally sorted at that point, so a tile boundary is simply a position where the key changes; fully-interior devices are left empty, which the rest of the pipeline already handles. Add two multi-GPU regression tests: SingleUpperNode, which places all input in one upper-node tile and asserts the exact unique active-voxel count, and MatchesSingleGpu, which cross-checks topology and occupancy against the single-GPU builder on the same input. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
| leftIntervals[deviceId] = 0; | ||
| } | ||
| cudaCheck(cudaEventRecord(transformReduceEvents[deviceId], stream)); | ||
| cudaCheck(cudaStreamSynchronize(stream)); |
Contributor
There was a problem hiding this comment.
Is it possible to do this on the device instead of the host in order to avoid the stream sync? Alternatively, would it be possible to run a benchmark on an analytic example (e.g. a sampled torus) to show that performance isn't affected?
Contributor
Author
There was a problem hiding this comment.
Sure, I added a snapBoundariesToRunsKernel so we run the snapping computation on-device.
…tream sync Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
The device segment boundaries produced by the initial even split were rebalanced by adjusting only adjacent pairs of boundaries. That cannot consolidate an upper-node tile spanning three or more GPUs, since a fully-interior device lies entirely within the tile. Downstream construction assumes each tile - and therefore each lower node, leaf node and voxel - is owned by exactly one device, so multiple devices built the same leaf concurrently and raced on its value mask, silently dropping active voxels.
Replace the pairwise GPU rebalance with a global, monotonic snap performed on the host over the (small) set of device boundaries. The keys are globally sorted at that point, so a tile boundary is simply a position where the key changes; fully-interior devices are left empty, which the rest of the pipeline already handles.
Add two multi-GPU regression tests: SingleUpperNode, which places all input in one upper-node tile and asserts the exact unique active-voxel count, and MatchesSingleGpu, which cross-checks topology and occupancy against the single-GPU builder on the same input.