SDMA HIP device management fixes

examples/sdma/sdma_bw.cpp

@@ -176,20 +176,30 @@ void runExperiment(int srcDeviceId, const ExperimentParams& params) {
 
   size_t totalNumQueues = params.numOfQueues * params.numDestinations;
 
+  // Resolve KFD node ids once: anvil queues are keyed on node ids, not HIP
+  // device ordinals, so both connect() and getSdmaQueue() below must use
+  // the SAME converted id (calling them separately with raw device ids would
+  // silently miss the queue lookup whenever a device ordinal != its node id).
+  int srcNode = anvil::anvil.nodeForHipDevice(srcDeviceId);
+  std::vector<int> dstNodes;
+  dstNodes.reserve(dstDeviceIds.size());
   for (auto& dstDeviceId : dstDeviceIds) {
+    dstNodes.push_back(anvil::anvil.nodeForHipDevice(dstDeviceId));
+  }
+
+  for (auto& dstNode : dstNodes) {
     // Better performance if allocating all 8 queues
-    anvil::anvil.connect(srcDeviceId, dstDeviceId, 8);  // params.numOfQueues);
+    anvil::anvil.connect(srcNode, dstNode, 8);  // params.numOfQueues);
   }
 
   anvil::SdmaQueueDeviceHandle** deviceHandles_d = nullptr;
   CHECK_HIP_ERROR(
       hipMalloc(&deviceHandles_d, totalNumQueues * sizeof(anvil::SdmaQueueDeviceHandle*)));
 
   size_t queueIdx = 0;
-  for (auto& dstDeviceId : dstDeviceIds) {
+  for (auto& dstNode : dstNodes) {
     for (size_t q = 0; q < params.numOfQueues; q++) {
-      deviceHandles_d[queueIdx] =
-          anvil::anvil.getSdmaQueue(srcDeviceId, dstDeviceId, q)->deviceHandle();
+      deviceHandles_d[queueIdx] = anvil::anvil.getSdmaQueue(srcNode, dstNode, q)->deviceHandle();
       queueIdx++;
     }
   }

examples/sdma/sdma_bw_allgather.cpp

@@ -411,8 +411,17 @@ void runExperimentMPI(int srcDeviceId, int mpiRank, int mpiSize, const Experimen
   // =============================
   std::cout << "Process " << mpiRank << ": Setting up Anvil SDMA queues..." << std::endl;
 
+  // Resolve KFD node ids once: anvil queues are keyed on node ids, not HIP
+  // device ordinals, so both connect() and getSdmaQueue() below must use
+  // the SAME converted id. Since this example keeps HIP_VISIBLE_DEVICES
+  // unsliced/identical across all MPI ranks (each rank just binds device ==
+  // its own rank), resolving a peer's node id from THIS process's own device
+  // enumeration is valid -- no cross-process node-id exchange needed here.
+  int srcNode = anvil::anvil.nodeForHipDevice(srcDeviceId);
+
   // Establish connection for each destination GPU
   for (int dstDeviceId : dstDeviceIds) {
+    int dstNode = anvil::anvil.nodeForHipDevice(dstDeviceId);
     try {
       // Ensure correct device context
       err = hipSetDevice(srcDeviceId);
@@ -433,7 +442,7 @@ void runExperimentMPI(int srcDeviceId, int mpiRank, int mpiSize, const Experimen
       std::cout << "Process " << mpiRank << ": Connecting GPU" << srcDeviceId << " -> GPU"
                 << dstDeviceId << " ..." << std::endl;
 
-      anvil::anvil.connect(srcDeviceId, dstDeviceId, params.numOfQueues);
+      anvil::anvil.connect(srcNode, dstNode, params.numOfQueues);
 
       std::cout << "Process " << mpiRank << ": GPU" << srcDeviceId << " -> GPU" << dstDeviceId
                 << " Anvil connection successful" << std::endl;
@@ -444,7 +453,7 @@ void runExperimentMPI(int srcDeviceId, int mpiRank, int mpiSize, const Experimen
       // Try fallback to using 1 queue
       std::cerr << "Process " << mpiRank << ": Trying to reconnect with 1 queue..." << std::endl;
       try {
-        anvil::anvil.connect(srcDeviceId, dstDeviceId, 1);
+        anvil::anvil.connect(srcNode, dstNode, 1);
         std::cout << "Process " << mpiRank << ": GPU" << srcDeviceId << " -> GPU" << dstDeviceId
                   << " Fallback connection successful" << std::endl;
       } catch (const std::exception& e2) {
@@ -479,9 +488,10 @@ void runExperimentMPI(int srcDeviceId, int mpiRank, int mpiSize, const Experimen
 
   for (size_t dstIdx = 0; dstIdx < dstDeviceIds.size(); dstIdx++) {
     int dstDeviceId = dstDeviceIds[dstIdx];
+    int dstNode = anvil::anvil.nodeForHipDevice(dstDeviceId);
     for (size_t q = 0; q < params.numOfQueues; q++) {
       try {
-        auto queue = anvil::anvil.getSdmaQueue(srcDeviceId, dstDeviceId, q);
+        auto queue = anvil::anvil.getSdmaQueue(srcNode, dstNode, q);
         if (queue && queue->deviceHandle()) {
           host_device_handles[queueIdx] = queue->deviceHandle();
           std::cout << "Process " << mpiRank << ": Destination GPU" << dstDeviceId << " queue[" << q

examples/sdma/sdma_latency.cpp

@@ -148,11 +148,16 @@ void runExperiment(int srcDeviceId, int dstDeviceId, const ExperimentParams& par
   // 3. Queue Setup
   // ======================
 
-  anvil::anvil.connect(srcDeviceId, dstDeviceId);
+  // Resolve KFD node ids once: anvil queues are keyed on node ids, not HIP
+  // device ordinals, so both connect() and getSdmaQueue() below must use
+  // the SAME converted id.
+  int srcNode = anvil::anvil.nodeForHipDevice(srcDeviceId);
+  int dstNode = anvil::anvil.nodeForHipDevice(dstDeviceId);
+  anvil::anvil.connect(srcNode, dstNode);
 
   anvil::SdmaQueueDeviceHandle* deviceHandle_d = nullptr;
 
-  deviceHandle_d = anvil::anvil.getSdmaQueue(srcDeviceId, dstDeviceId)->deviceHandle();
+  deviceHandle_d = anvil::anvil.getSdmaQueue(srcNode, dstNode)->deviceHandle();
 
   // if (params.verbose)
   // {

examples/sdma/sdma_rate.cpp

@@ -137,16 +137,20 @@ void runExperiment(int srcDeviceId, int dstDeviceId, const ExperimentParams& par
   // ======================
 
   size_t totalNumQueues = params.numOfQueues;
-  anvil::anvil.connect(srcDeviceId, dstDeviceId, params.numOfQueues);
+  // Resolve KFD node ids once: anvil queues are keyed on node ids, not HIP
+  // device ordinals, so both connect() and getSdmaQueue() below must use
+  // the SAME converted id.
+  int srcNode = anvil::anvil.nodeForHipDevice(srcDeviceId);
+  int dstNode = anvil::anvil.nodeForHipDevice(dstDeviceId);
+  anvil::anvil.connect(srcNode, dstNode, params.numOfQueues);
 
   anvil::SdmaQueueDeviceHandle** deviceHandles_d = nullptr;
   CHECK_HIP_ERROR(
       hipMalloc(&deviceHandles_d, totalNumQueues * sizeof(anvil::SdmaQueueDeviceHandle*)));
 
   size_t queueIdx = 0;
   for (size_t q = 0; q < params.numOfQueues; q++) {
-    deviceHandles_d[queueIdx] =
-        anvil::anvil.getSdmaQueue(srcDeviceId, dstDeviceId, q)->deviceHandle();
+    deviceHandles_d[queueIdx] = anvil::anvil.getSdmaQueue(srcNode, dstNode, q)->deviceHandle();
 
     // if (params.verbose)
     // {

include/mori/application/context/context.hpp

@@ -78,7 +78,6 @@ class Context {
 
   int LocalRank() const { return bootNet.GetLocalRank(); }
   int WorldSize() const { return bootNet.GetWorldSize(); }
-  int LocalRankInNode() const { return rankInNode; }
   const std::string& HostName() const { return myHostname; }
 
   // Single-value transport selection driven by Context's default policy
@@ -95,6 +94,19 @@ class Context {
   const PeerCapabilities& GetPeerCapabilities(int destRank) const { return peerCaps[destRank]; }
   const std::vector<PeerCapabilities>& GetAllPeerCapabilities() const { return peerCaps; }
 
+  // KFD topology node id (== HSA_AGENT_INFO_NODE) of a peer's GPU. This is a
+  // process-independent, host-global identity: it does not depend on
+  // HIP_VISIBLE_DEVICES, so it is the correct key for wiring SDMA queues to a
+  // peer even when the peer GPU is not in this process's HIP device list.
+  // -1 if the peer's node id could not be resolved. Populated in
+  // CollectHostNames() via an allgather.
+  int KfdNodeId(int destRank) const { return peerInfos[destRank].kfdNodeId; }
+  int LocalKfdNode() const { return peerInfos[LocalRank()].kfdNodeId; }
+
+  // Local rank within this host (0-based index among same-host peers).
+  // Derived from peerInfos rather than a separately-tracked member.
+  int LocalRankInNode() const { return peerInfos[LocalRank()].rankInNode; }
+
   RdmaContext* GetRdmaContext() const { return rdmaContext.get(); }
   RdmaDeviceContext* GetRdmaDeviceContext() const { return rdmaDeviceContext.get(); }
   const std::vector<std::unique_ptr<RdmaDeviceContext>>& GetAllRdmaDeviceContexts() const {
@@ -184,16 +196,13 @@ class Context {
     bool sameHost{false};
     bool sameProcess{false};
     int rankInNode{-1};
+    int kfdNodeId{-1};  // KFD topology node id of this rank's GPU (host-global)
   };
 
   bool RailOnlyEligible() const;
 
  private:
-  // Number of same-host peers with rank < `rank` (a peer's within-node device id).
-  int SameHostPeersBefore(int rank) const;
-
   BootstrapNetwork& bootNet;
-  int rankInNode{-1};
   int numQpPerPe{4};
   bool sdmaEnabled{false};
   bool p2pDisabled{false};

include/mori/application/transport/sdma/anvil.hpp

@@ -47,13 +47,12 @@ namespace anvil {
 
 class SdmaQueue {
  public:
-  SdmaQueue(int localDeviceId, int remoteDeviceId, hsa_agent_t& localAgent, uint32_t engineId);
+  SdmaQueue(uint32_t localNodeId, uint32_t engineId);
   ~SdmaQueue();
 
   SdmaQueueDeviceHandle* deviceHandle() const;
 
  private:
-  int remoteDeviceId_;  // TODO unused
   uint64_t* cachedWptr_;
   uint64_t* committedWptr_;
   void* queueBuffer_;
@@ -76,8 +75,25 @@ class AnvilLib {
 
  public:
   void init();
-  bool connect(int srcDeviceId, int dstDeviceId, int numChannels = 1);
-  SdmaQueue* getSdmaQueue(int srcDeviceId, int dstDeviceId, int channelIdx = 0);
+  // srcNode/dstNode are KFD topology node ids (== HSA_AGENT_INFO_NODE), a
+  // host-global GPU identity that does NOT depend on HIP_VISIBLE_DEVICES. This
+  // is the correct key for a peer even when the peer GPU is not in this
+  // process's HIP device list. Channels for a pair are shared process-wide.
+  bool connect(int srcNode, int dstNode, int numChannels = 1);
+
+  // Get the SDMA queue for a given src/dst node pair and channel index.
+  SdmaQueue* getSdmaQueue(int srcNode, int dstNode, int channelIdx = 0);
+
+  // Map a HIP device ordinal to its KFD node id. For single-process callers
+  // (e.g. the examples) that only have HIP device ids; multi-process collectives
+  // should exchange node ids out-of-band instead (see Context::KfdNodeId).
+  static uint32_t nodeForHipDevice(int hipDev);
+
+  // Resolve the KFD topology node id of the given HIP device WITHOUT initializing
+  // HSA. The KFD node id (the directory index under /sys/class/kfd/kfd/topology/nodes)
+  // is identical to what HSA_AGENT_INFO_NODE / hsaKmtGetNodeProperties report,
+  // and is a host-global identity independent of HIP_VISIBLE_DEVICES.
+  static int kfdNodeIdForHipDevice(int hipDev);
 
  private:
   /*
@@ -101,25 +117,23 @@ class AnvilLib {
                                                      {5, 3, 2, 4, 6, 1, 0, 7},
                                                      {3, 6, 4, 2, 1, 5, 7, 0}}};
 
-  int getOamId(int deviceId);
+  // xGMI physical (OAM) id for a KFD node, read from the GPU's PCI sysfs.
+  int getOamId(int node);
 
-  int getSdmaEngineId(int srcDeviceId, int dstDeviceId);
-
-  // KFD topology node id for a HIP device id.
-  uint32_t getNodeId(int deviceId);
+  int getSdmaEngineId(int srcNode, int dstNode);
 
   // Bitmask of SDMA engine ids KFD recommends for the src->dst xGMI link to
   // reach maximum bandwidth (sysfs recommended_sdma_engine_id_mask). Returns 0
   // if the link or property is unavailable, in which case callers fall back to
   // the static OAM map.
-  uint32_t getRecommendedEngineMask(int srcDeviceId, int dstDeviceId);
+  uint32_t getRecommendedEngineMask(int srcNode, int dstNode);
 
   // Bitmask of the general (CPU-link, non-xGMI) SDMA engines. Zero if the node
   // reports no CPU link. Used to spread loopback channels on gfx1250.
-  uint32_t getHostLinkEngineMask(int srcDeviceId);
+  uint32_t getHostLinkEngineMask(int srcNode);
 
   // True for gfx1250 (gfx12.5), the only arch that spreads loopback channels.
-  bool isGfx1250(int deviceId);
+  bool isGfx1250(int node);
 
   struct PairHash {
     std::size_t operator()(const std::pair<int, int>& p) const {
@@ -144,22 +158,6 @@ inline void checkHipError(hipError_t err, const char* msg, const char* file, int
 }
 
 #define CHECK_HIP_ERROR(cmd) anvil::checkHipError((cmd), #cmd, __FILE__, __LINE__)
-// Allow access to peerDeviceId from deviceId
-inline void EnablePeerAccess(int const deviceId, int const peerDeviceId) {
-  int canAccess;
-  CHECK_HIP_ERROR(hipDeviceCanAccessPeer(&canAccess, deviceId, peerDeviceId));
-  if (!canAccess) {
-    std::cerr << "Unable to enable peer access from GPU devices " << deviceId << " to "
-              << peerDeviceId << "\n";
-  }
-
-  CHECK_HIP_ERROR(hipSetDevice(deviceId));
-  hipError_t error = hipDeviceEnablePeerAccess(peerDeviceId, 0);
-  if (error != hipSuccess && error != hipErrorPeerAccessAlreadyEnabled) {
-    std::cerr << "Unable to enable peer to peer access from " << deviceId << "  to " << peerDeviceId
-              << " (" << hipGetErrorString(error) << ")\n";
-  }
-}
 // Hardware cap on SDMA channels per GPU pair on CDNA (4 queues/engine ×
 // 2 recommended engines). Requests above this are clamped, not failed.
 inline constexpr int kMaxSdmaChannelsPerPair = 8;