SDMA HIP device management fixes - #623
Draft
pemeliya wants to merge 9 commits into
Draft
Conversation
Wire intra-node SDMA queues by the host-global KFD topology node id (== HSA_AGENT_INFO_NODE) rather than a HIP device ordinal. HIP ordinals are a per-process slice of HIP_VISIBLE_DEVICES, so the old within-node derivations (Context::SameHostPeersBefore, cco's dstDeviceId=lsa, and the symmetric-memory SDMA-peer counter) plus the BDF band-aid only agreed in the full-8-GPU-visibility case and broke under sliced visibility. - Context now resolves each rank's local GPU KFD node id HSA-free from the bound device BDF via sysfs and allgathers it in CollectHostNames; adds KfdNodeId()/LocalKfdNode() accessors. - anvil re-keys sdma_channels_ on (srcNode,dstNode); connect/getSdmaQueue and the engine-mask helpers take node ids; adds agentForNode() to pick the local queue's HSA agent and nodeForHipDevice() for single-process callers. - EnsureSdmaTransport uses node ids and drops hipDeviceEnablePeerAccess: the SDMA copy engine targets already-mapped fabric/IPC VAs, so HIP peer access is unnecessary and was the only hard HIP-visibility dependency. - cco_init and symmetric_memory pass node ids to getSdmaQueue; kernel-facing handle arrays stay indexed by logical rank (lsaRank / global pe). - examples/sdma convert HIP ordinals via nodeForHipDevice(). Cannot be built/tested in this environment (no ROCm/HIP); WIP for hardware validation, including a sliced HIP_VISIBLE_DEVICES run and a full-visibility regression. Co-authored-by: Pavel Emeliyanenko <pemeliya@users.noreply.github.com>
Two multi-rank CCO SDMA all-gather tests exercising the KFD-node-id queue keying across both launch topologies: - test_sdma_allgather_torch: full HIP visibility, one distinct GPU per rank (hipSetDevice(rank % numDevices)) -- HIP ordinals match physical GPUs. - test_sdma_allgather_jax: each rank slices HIP_VISIBLE_DEVICES=<rank> before its first HIP call and binds device 0, so HIP ordinal 0 maps to a different physical GPU per process (ROCR_VISIBLE_DEVICES left unset so HSA still sees all GPUs). This is the sliced-visibility case that HIP-ordinal keying could not handle and that the node-id keying fixes. Shared kernel body + comm setup/verify live in sdma_allgather_common.hpp; each test keeps its own __global__ wrapper so the CMake auto-discovery flags it HIP. Auto-registered as cco_sdma_allgather_torch / cco_sdma_allgather_jax under BUILD_CCO_SDMA; both SKIP when MORI_ENABLE_SDMA is unset. Not built here (no ROCm/HIP); WIP for hardware validation. Co-authored-by: Pavel Emeliyanenko <pemeliya@users.noreply.github.com>
pemeliya
marked this pull request as draft
August 31, 2026 11:14
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.
Motivation
Technical Details
Test Plan
Test Result
Submission Checklist