SDMA HIP device management fixes - #623
Open
pemeliya wants to merge 2 commits into
Open
Conversation
There was a problem hiding this comment.
Pull request overview
Updates SDMA queue management to use host-global KFD node IDs, supporting sliced HIP_VISIBLE_DEVICES.
Changes:
- Exchanges KFD node IDs between peers and keys Anvil queues by node.
- Updates SDMA examples and symmetric-memory wiring.
- Adds sliced-device all-gather coverage and test-runner integration.
Reviewed changes
Copilot reviewed 12 out of 12 changed files in this pull request and generated 3 comments.
Show a summary per file
| File | Description |
|---|---|
tools/run_cco_tests.sh |
Runs the new device-assignment test. |
tests/cpp/cco/test_sdma_hip_dev_assign.cpp |
Tests sliced HIP device assignments. |
src/cco/cco_init.cpp |
Uses KFD nodes for CCO queue handles. |
src/application/transport/sdma/anvil.cpp |
Refactors queues around KFD node IDs. |
src/application/memory/symmetric_memory.cpp |
Maps symmetric-memory handles by KFD node. |
src/application/context/context.cpp |
Resolves and exchanges GPU KFD nodes. |
include/mori/application/transport/sdma/anvil.hpp |
Updates the Anvil API contract. |
include/mori/application/context/context.hpp |
Exposes peer KFD node identities. |
examples/sdma/sdma_rate.cpp |
Converts example HIP ordinals to nodes. |
examples/sdma/sdma_latency.cpp |
Converts example HIP ordinals to nodes. |
examples/sdma/sdma_bw.cpp |
Converts bandwidth example queue keys. |
examples/sdma/sdma_bw_allgather.cpp |
Converts all-gather example queue keys. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Comment on lines
+252
to
+258
| // cco_test_harness.hpp's HIP_CHECK macro reports failures via the | ||
| // process-wide g_rank global (not thread-local). Setting it here is | ||
| // best-effort for diagnostics only: under concurrent threads the printed | ||
| // rank in a HIP_CHECK failure message can race, but HIP_CHECK always | ||
| // _exit(1)s the whole process immediately regardless, so this never | ||
| // affects correctness — only which rank number an error message blames. | ||
| g_rank = rank; |
| if (!strcmp(argv[i], "--n_processes") && i + 1 < argc) { | ||
| nProcesses = atoi(argv[++i]); | ||
| } else if (!strcmp(argv[i], "--n_local_devices") && i + 1 < argc) { | ||
| g_devicesPerProcess = std::max(1, atoi(argv[++i])); |
pemeliya
force-pushed
the
pemeliya/sdma-hsa-node-id
branch
2 times, most recently
from
September 1, 2026 10:03
658db2d to
c32f50c
Compare
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 8 out of 8 changed files in this pull request and generated 3 comments.
Suppressed comments (1)
tests/cpp/cco/test_sdma_hip_dev_assign.cpp:253
- Concurrent local-rank workers write this process-wide non-atomic variable while their
HIP_CHECKexpansions also read it. That is a C++ data race, not merely a potentially inaccurate diagnostic. Keep the process index assigned byrun_test, or use a rank-explicit check likeHIP_CHECK_MTintest_sdma_put_mt.cpp.
g_rank = rank;
Comment on lines
+260
to
+263
| if (SdmaAllGatherSetup(rank, nranks, uid, &ctx) != 0) { | ||
| snprintf(result->detail, sizeof(result->detail), "setup failed"); | ||
| return; | ||
| } |
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 12 out of 12 changed files in this pull request and generated 2 comments.
Suppressed comments (2)
Previously missed (2) — in code that hasn't changed since the last review.
tools/run_cco_tests.sh:21
- This test is now skipped by the default sweep, but
.github/workflows/ci_cco.yml:82-83also omits it from the explicit SDMA test loop. As a result, CI never runs the new sliced-visibility regression with SDMA enabled (CTest only reaches its self-SKIP path). Add this binary to the SDMA CI step when excluding it here.
test_sdma_hip_dev_assign) continue ;;
src/application/context/context.cpp:192
- This contradicts the PR's stated “without HIP/HSA calls (pure sysfs)” initialization path:
CollectHostNamesnow callshipGetDevice, andkfdNodeIdForHipDevicecallshipDeviceGetPCIBusId. If avoiding runtime initialization is a requirement, resolve the active device from environment/sysfs data; otherwise update the PR and API claims to state that only direct HSA initialization is avoided.
int hipDev = 0;
HIP_RUNTIME_CHECK(hipGetDevice(&hipDev));
| } | ||
| setenv("HIP_VISIBLE_DEVICES", vis.c_str(), /*overwrite=*/1); | ||
| } | ||
| // Else: leave HIP_VISIBLE_DEVICES unset, so native full visibility applies. |
Comment on lines
+461
to
+462
| } else if (!g_visibleDevices.empty() && | ||
| static_cast<int>(g_visibleDevices.size()) != nProcesses * g_devicesPerProcess) { |
pemeliya
force-pushed
the
pemeliya/sdma-hsa-node-id
branch
from
September 1, 2026 15:13
9f72cb6 to
584de3d
Compare
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> test(cco/sdma): add torch- and jax-style SDMA all-gather tests 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> updated anvil refactoring and added tests changed the tests fixed clang some fixes clang fix fix clang refactor(context): dedup rankInNode; derive LocalRankInNode from peerInfos The Context::rankInNode member was a redundant recomputation of peerInfos[LocalRank()].rankInNode (the -1 init just cancelled the self-count in the loop). Drop the member and the counting loop: LocalRankInNode() now returns peerInfos[LocalRank()].rankInNode directly (mirroring LocalKfdNode()), and InitializeTopologyAndTransports takes a local from the same source for NIC selection. No behavior change to rail-only routing or NIC matching. Co-authored-by: Pavel Emeliyanenko <pemeliya@users.noreply.github.com> refactoring continued restored legacy tests that use hipDevID update fixing precommit updated clang refactor Breaking the loop on the failure Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> interface update
pemeliya
force-pushed
the
pemeliya/sdma-hsa-node-id
branch
from
September 1, 2026 15:29
584de3d to
91b832b
Compare
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
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.
Summary
Key intra-node SDMA queues on the host-global KFD topology node id (==
HSA_AGENT_INFO_NODE) instead of a HIP device ordinal / within-node index. HIP ordinals are only a per-process slice ofHIP_VISIBLE_DEVICES, and the previous within-node index (SameHostPeersBefore,hipDev/lsaused directly as anvil's device-pair key) only lined up with physical GPUs in the full-8-GPU-visibility case. Under slicedHIP_VISIBLE_DEVICES(the JAX-style multi-process launch pattern) those indices diverge from topology, so SDMA queues were wired to the wrong peer. KFD node ids are resolved without HIP/HSA calls (pure sysfs) and exchanged once per rank, so they stay correct regardless of what each process can see.What changed
Context:CollectHostNames()now also allgathers each rank's local KFD node id (resolved viaanvil::kfdNodeIdForHipDevice(), sysfs-only, no HSA init) alongside the existing pid/rail-flag/nodeId record, stored per-peer asPeerInfo::kfdNodeId. New accessorsKfdNodeId(rank)/LocalKfdNode(). RemovedSameHostPeersBefore()and therankInNodemember —rankInNodeis now read once frompeerInfos[LocalRank()].rankInNode(already computed during the allgather) instead of being recomputed by a separate loop inInitializeTopologyAndTransports().EnsureSdmaTransport(): connects peers viaanvil::anvil.connect(localNode, peerNode, ...)keyed on KFD node ids (aborting loudly if either side's node id failed to resolve), and drops theEnablePeerAccess/hipDeviceEnablePeerAccesscall entirely — the SDMA copy engine operates on already-mapped fabric/IPC VAs, so HIP peer access was never required and was the only remaining hard dependency on HIP-visibility-derived device ids.anvil(AnvilLib):sdma_channels_is now keyed on(srcNode, dstNode)node-id pairs rather than HIP device ids.connect(srcNode, dstNode, numChannels)andgetSdmaQueue(srcNode, dstNode, channelIdx)both take node ids now (the old device-id-takingconnect()overload is gone — there's a single, consistently node-id-keyed API).getSdmaEngineId,getRecommendedEngineMask,getHostLinkEngineMask,isGfx1250,getOamId) all now take node ids instead of device ids.SdmaQueue's constructor drops the unusedremoteDeviceId_member and the HSA-agent-lookup-by-device-id it did internally; it now takes the resolved node id directly.nodeForHipDevice()(via HSA, for single-process callers that only have HIP device ids) andkfdNodeIdForHipDevice()(sysfs-only, used byContext) statics.getBusId()now returns the packed BDF integer directly (previously a hex string that callers re-parsed);getOamId()andkfdNodeIdForHipDevice()share a newlocIdAndDomainForNode()sysfs-parsing helper instead of duplicating thelocation_id/domainfile-scan logic.kfdNodeIdForHipDevice(): the node-id scan didn't stop when it ran past the last real KFD node (sysfs read failure was silently treated as "keep scanning" rather than "no more nodes"); now breaks out oncelocIdAndDomainForNode()reports an unavailable node.cco_init.cpp(ccoSdmaSetupCommQueues) andsymmetric_memory.cpp(RegisterSymmMemObj): both switch theirgetSdmaQueue()calls from a within-node index (comm->hipDev/lsa, or a manually-recounted "SDMA peers before me" index) tocontext->LocalKfdNode()/context->KfdNodeId(pe). The kernel-facing device-handle arrays are unchanged — still indexed by logical rank (lsaRank, or globalpe), since only the anvil lookup key was wrong, not the on-device layout.examples/sdma/{sdma_bw,sdma_rate,sdma_latency,sdma_bw_allgather}.cpp: these single-process (and, forsdma_bw_allgather, full-visibility MPI) examples now resolve each HIP device id to its KFD node id once viaanvil::anvil.nodeForHipDevice(), and use the resulting node id consistently for bothconnect()and everygetSdmaQueue()lookup (previously they mixed the two, which happened to work only because HIP ordinal == KFD node id on this particular box).tools/run_cco_tests.sh: skipstest_sdma_hip_dev_assignin the default sweep (alongside the other SDMA tests already skipped there), since it takes its own--n_processes/--n_local_devices/--visible_devicesflags instead of a bare positional rank count and needsMORI_ENABLE_SDMA=1.New test coverage
Added
tests/cpp/cco/test_sdma_hip_dev_assign.cpp— a self-contained CCO SDMA all-gather test (auto-discovered by the existingtest_sdma_*CMake glob, built only underBUILD_CCO_SDMA, self-SKIPs whenMORI_ENABLE_SDMAis unset) that exercises both real-world multi-process GPU-visibility patterns the node-id keying is meant to unify:--n_processesforked processes slicesHIP_VISIBLE_DEVICESinto a disjoint--n_local_devices-sized chunk before its first HIP call, then spawns one thread per local device (SPMT), each binding local ordinal0..n_local_devices-1. This is the sliced-visibility case where a peer's HIP ordinal doesn't correspond to the same physical GPU across processes — exactly what broke with device-id-based keying.--torch_style): every process shares the same (unsliced) visible-device set and binds device ordinal == its own rank viahipSetDevice(i), mirroring a real PyTorch multi-process launch where all GPUs are visible to every process.--visible_devices D0,D1,...lets either mode run over a custom/non-contiguous physical device order (jax-style: sliced per process, exact length required; torch-style: full list shared unsliced, only needslen >= n_processes).Testing
test_sdma_hip_dev_assignand smoke-tested on an 8-GPU ROCm host (pemeliya-rocm7.14container):--n_processes/--n_local_devices/--visible_devicescombinations — all ranks PASSED.--torch_stylealone (auto-detected GPU count), with--n_processes, and with--visible_devices— all ranks PASSED.--torch_style --n_processes 99(exceeds available devices) fails loudly with a clear error instead of hanging or silently clamping.Risk / follow-ups
anvil,Context,cco_init,symmetric_memory); worth a careful look at the sysfs-based node-id resolution (kfdNodeIdForHipDevice/locIdAndDomainForNode) on hosts with unusual KFD topology numbering or multi-segment PCI domains.EnablePeerAccess/hipDeviceEnablePeerAccessremoval inEnsureSdmaTransportis intentional (SDMA doesn't need HIP peer access) but changes a previously "always-on" call, so worth flagging explicitly in review.EnsureSdmaTransportnowstd::abort()s if a local or peer KFD node id fails to resolve, rather than silently falling back — this is a stricter failure mode than before and should be validated doesn't fire spuriously on other topologies.