From bb683b62c3adae82bee5154e992ba62de73d2386 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Mon, 5 Oct 2026 16:18:55 -0400 Subject: [PATCH 01/10] fix(rocshmem): make Anvil SDMA usable on MI300X CPX partitions Anvil SDMA could not initialize on a CPX partition. Two defects: getOamId read xgmi_physical_id from a sysfs node that does not exist for a partition. A CPX partition is a HIP alias at PCI function .1-.7 with no node of its own, so the open failed. Fall back to the physical function-0 BDF, and fold the PCI function into the engine id so same-device peers do not all land on engine 0. Engine selection indexed mi300xOamMap, which describes the 14 xGMI SDMA engines of an unpartitioned MI300X. A CPX partition owns 2 host engines and 0 xGMI engines, so every map id was out of range and hsaKmtCreateQueueExt returned INVALID_PARAMETER. Fold the undoubled map value plus the PCI function into the reported range, retry with a generic HSA_QUEUE_SDMA when the engine-pinned queue is still rejected, and report failure through SdmaQueue::valid() so connect() fails instead of aborting. KFD caps user SDMA queues per engine, so a partition also has a small total budget. Count usage across both engines and refuse a request that cannot fit before KFD returns NO_MEMORY part way through building it. Adds disconnectDevice() so a caller can release the queues it took for one destination without tearing down the process-global channel map. Verified on dell300x-ccs-aus-k13-09 (8x MI300X, CPX): SDMA queues create on a partition and both partition engines are used for same-device peers. JIRA ID: AICOMRCCL-2387 Co-authored-by: Cursor --- projects/rocshmem/src/sdma/anvil.cpp | 356 ++++++++++++++++++++++---- projects/rocshmem/src/sdma/anvil.hpp | 89 ++++++- projects/rocshmem/src/sdma_policy.cpp | 12 +- 3 files changed, 396 insertions(+), 61 deletions(-) diff --git a/projects/rocshmem/src/sdma/anvil.cpp b/projects/rocshmem/src/sdma/anvil.cpp index f9f181d0e174..87a4dce45256 100644 --- a/projects/rocshmem/src/sdma/anvil.cpp +++ b/projects/rocshmem/src/sdma/anvil.cpp @@ -36,6 +36,54 @@ namespace sdma_anvil { +static const char* hsakmtStatusName(HSAKMT_STATUS status) { + switch (status) { + case HSAKMT_STATUS_SUCCESS: + return "SUCCESS"; + case HSAKMT_STATUS_ERROR: + return "ERROR"; + case HSAKMT_STATUS_DRIVER_MISMATCH: + return "DRIVER_MISMATCH"; + case HSAKMT_STATUS_INVALID_PARAMETER: + return "INVALID_PARAMETER"; + case HSAKMT_STATUS_INVALID_HANDLE: + return "INVALID_HANDLE"; + case HSAKMT_STATUS_INVALID_NODE_UNIT: + return "INVALID_NODE_UNIT"; + case HSAKMT_STATUS_NO_MEMORY: + return "NO_MEMORY"; + case HSAKMT_STATUS_BUFFER_TOO_SMALL: + return "BUFFER_TOO_SMALL"; + case HSAKMT_STATUS_NOT_IMPLEMENTED: + return "NOT_IMPLEMENTED"; + case HSAKMT_STATUS_NOT_SUPPORTED: + return "NOT_SUPPORTED"; + case HSAKMT_STATUS_UNAVAILABLE: + return "UNAVAILABLE"; + case HSAKMT_STATUS_OUT_OF_RESOURCES: + return "OUT_OF_RESOURCES"; + case HSAKMT_STATUS_KERNEL_IO_CHANNEL_NOT_OPENED: + return "KERNEL_IO_CHANNEL_NOT_OPENED"; + case HSAKMT_STATUS_KERNEL_COMMUNICATION_ERROR: + return "KERNEL_COMMUNICATION_ERROR"; + case HSAKMT_STATUS_KERNEL_ALREADY_OPENED: + return "KERNEL_ALREADY_OPENED"; + case HSAKMT_STATUS_HSAMMU_UNAVAILABLE: + return "HSAMMU_UNAVAILABLE"; + case HSAKMT_STATUS_WAIT_FAILURE: + return "WAIT_FAILURE"; + case HSAKMT_STATUS_WAIT_TIMEOUT: + return "WAIT_TIMEOUT"; + case HSAKMT_STATUS_MEMORY_ALREADY_REGISTERED: + return "MEMORY_ALREADY_REGISTERED"; + case HSAKMT_STATUS_MEMORY_NOT_REGISTERED: + return "MEMORY_NOT_REGISTERED"; + case HSAKMT_STATUS_MEMORY_ALIGNMENT: + return "MEMORY_ALIGNMENT"; + default: + return "UNKNOWN"; + } +} #define CHECK_HSAKMT_SUCCESS(call, msg) do { \ if ((call) != HSAKMT_STATUS_SUCCESS) \ @@ -117,17 +165,19 @@ static const std::string getBusId(int deviceId) { return std::string(busIdChar); } -SdmaQueue::SdmaQueue([[maybe_unused]] int localDeviceId, int remoteDeviceId, - const hsa_agent_t& localAgent, uint32_t engineId) +SdmaQueue::SdmaQueue(int localDeviceId, int remoteDeviceId, const hsa_agent_t& localAgent, + uint32_t engineId, const EngineSelection& selection) : remoteDeviceId_(remoteDeviceId) { int originalDeviceId; ANVIL_CHECK_HIP_ERROR(hipGetDevice(&originalDeviceId)); // Save the current device - uint32_t localNodeId; + uint32_t localNodeId = 0; hsa_status_t status = hsa_agent_get_info(localAgent, HSA_AGENT_INFO_NODE, &localNodeId); if (status != HSA_STATUS_SUCCESS) { - LOG_TRACE("Failure to get device info: %#x", status); + LOG_ERROR("anvil: HSA_AGENT_INFO_NODE failed device=%d status=%#x", localDeviceId, status); + createStatus_ = HSAKMT_STATUS_ERROR; + return; } // Allocate SDMA queue buffer on device side, requires ExecuteAccess @@ -149,10 +199,53 @@ SdmaQueue::SdmaQueue([[maybe_unused]] int localDeviceId, int remoteDeviceId, // Create SDMA Queue memset(&queue_, 0, sizeof(HsaQueueResource)); - CHECK_HSAKMT_SUCCESS(hsaKmtCreateQueueExt(localNodeId, HSA_QUEUE_SDMA_BY_ENG_ID, - DEFAULT_QUEUE_PERCENTAGE, DEFAULT_PRIORITY, engineId, - queueBuffer_, SDMA_QUEUE_SIZE, nullptr, &queue_), - "hsaKmtCreateQueueExt failed"); + HSAKMT_STATUS pinnedStatus = hsaKmtCreateQueueExt( + localNodeId, HSA_QUEUE_SDMA_BY_ENG_ID, DEFAULT_QUEUE_PERCENTAGE, DEFAULT_PRIORITY, engineId, + queueBuffer_, SDMA_QUEUE_SIZE, nullptr, &queue_); + HSAKMT_STATUS queueStatus = pinnedStatus; + + if (queueStatus != HSAKMT_STATUS_SUCCESS) { + // An engine-pinned queue needs an engine id the KFD node actually owns. A CPX/DPX partition owns + // one XCD's SDMA engines and no xGMI engines at all, so an id that is valid on the unpartitioned + // device is rejected here. A generic HSA_QUEUE_SDMA lets KFD pick an engine it owns: that gives + // up the per-peer xGMI engine affinity but still reaches the peer over xGMI. + LOG_WARN("anvil: engine-pinned queue rejected (hsakmt=%d %s node=%u engineId=%u), retrying with " + "a generic SDMA queue", + static_cast(queueStatus), hsakmtStatusName(queueStatus), localNodeId, engineId); + memset(&queue_, 0, sizeof(queue_)); + queueStatus = + hsaKmtCreateQueueExt(localNodeId, HSA_QUEUE_SDMA, DEFAULT_QUEUE_PERCENTAGE, DEFAULT_PRIORITY, + 0, queueBuffer_, SDMA_QUEUE_SIZE, nullptr, &queue_); + } + + // A generic retry can return NO_MEMORY after the pinned create returned INVALID_PARAMETER. + // Keeping only the pinned status would hide the budget failure from connect(). + if (queueStatus == HSAKMT_STATUS_SUCCESS) { + createStatus_ = HSAKMT_STATUS_SUCCESS; + } else if (pinnedStatus == HSAKMT_STATUS_NO_MEMORY || queueStatus == HSAKMT_STATUS_NO_MEMORY) { + createStatus_ = HSAKMT_STATUS_NO_MEMORY; + } else { + createStatus_ = pinnedStatus; + } + + if (queueStatus != HSAKMT_STATUS_SUCCESS) { + const std::string srcBus = getBusId(localDeviceId); + const std::string dstBus = getBusId(remoteDeviceId); + LOG_ERROR( + "anvil: hsaKmtCreateQueueExt failed hsakmt=%d (%s) node=%u engineId=%u srcDev=%d (%s) " + "dstDev=%d (%s) usedPreferred=%d preferredStatus=%#x preferredMask=0x%x hostEng=%u " + "xgmiEng=%u total=%u", + static_cast(createStatus_), hsakmtStatusName(createStatus_), localNodeId, engineId, + localDeviceId, srcBus.c_str(), remoteDeviceId, dstBus.c_str(), + selection.usedPreferred ? 1 : 0, static_cast(selection.preferredStatus), + selection.preferredMask, selection.numSdmaEngines, selection.numSdmaXgmiEngines, + selection.numSdmaEnginesTotal); + // Leave valid_ false so the caller can drop the Anvil backend instead of killing the job. + hsaKmtUnmapMemoryToGPU(queueBuffer_); + hsaKmtFreeMemory(queueBuffer_, SDMA_QUEUE_SIZE); + queueBuffer_ = nullptr; + return; + } // Populate Device Handle ANVIL_CHECK_HIP_ERROR(hipMalloc(&deviceHandle_, sizeof(SdmaQueueDeviceHandle))); @@ -186,9 +279,12 @@ SdmaQueue::SdmaQueue([[maybe_unused]] int localDeviceId, int remoteDeviceId, ANVIL_CHECK_HIP_ERROR(hipMemcpy(cachedWptr_, &cachedWptr, sizeof(uint64_t), hipMemcpyHostToDevice)); ANVIL_CHECK_HIP_ERROR( hipMemcpy(committedWptr_, &committedWptr, sizeof(uint64_t), hipMemcpyHostToDevice)); + + valid_ = true; } SdmaQueue::~SdmaQueue() { + if (!valid_) return; CHECK_HSAKMT_SUCCESS(hsaKmtDestroyQueue(queue_.QueueId), "Failed to destroy queue."); ANVIL_CHECK_HIP_ERROR(hipFree(deviceHandle_)); if (singleProducerDeviceHandle_) ANVIL_CHECK_HIP_ERROR(hipFree(singleProducerDeviceHandle_)); @@ -198,6 +294,10 @@ SdmaQueue::~SdmaQueue() { CHECK_HSAKMT_SUCCESS(hsaKmtFreeMemory(queueBuffer_, SDMA_QUEUE_SIZE), "Failed"); } +bool SdmaQueue::valid() const { return valid_; } + +HSAKMT_STATUS SdmaQueue::createStatus() const { return createStatus_; } + SdmaQueueDeviceHandle* SdmaQueue::deviceHandle() const { return deviceHandle_; } SdmaQueueSingleProducerDeviceHandle* SdmaQueue::singleProducerDeviceHandle() const { @@ -385,8 +485,18 @@ void AnvilLib::querySdmaEngineCounts() { } numSdmaEnginesTotal_ = numSdmaEngines_ + numSdmaXgmiEngines_; - LOG_TRACE("anvil: SDMA engines host=%u xgmi=%u total=%u", numSdmaEngines_, numSdmaXgmiEngines_, - numSdmaEnginesTotal_); + + uint32_t node = 0; + HsaNodeProperties nodeProps{}; + if (hsa_agent_get_info(agent, HSA_AGENT_INFO_NODE, &node) == HSA_STATUS_SUCCESS && + hsaKmtGetNodeProperties(node, &nodeProps) == HSAKMT_STATUS_SUCCESS) { + numSdmaQueuesPerEngine_ = nodeProps.NumSdmaQueuesPerEngine; + } else { + LOG_WARN("anvil: SDMA queues-per-engine unknown for node %u; queue budget unchecked", node); + } + + LOG_TRACE("anvil: SDMA engines host=%u xgmi=%u total=%u queuesPerEngine=%u", numSdmaEngines_, + numSdmaXgmiEngines_, numSdmaEnginesTotal_, numSdmaQueuesPerEngine_); } void AnvilLib::init() { @@ -406,18 +516,23 @@ void AnvilLib::init() { } buildGpuAgentMap(); - querySdmaEngineCounts(); + // Before querySdmaEngineCounts: the per-engine queue budget comes from the KFD node properties. SetUpKFD(); s_kfd_opened = true; + + querySdmaEngineCounts(); }); } SdmaQueue* AnvilLib::createSdmaQueue(int srcDeviceId, int dstDeviceId, uint32_t engineId, - int* channelIdx) { + const EngineSelection& selection, int* channelIdx) { auto& vec = sdma_channels_[dstDeviceId]; - vec.emplace_back(std::make_unique(srcDeviceId, dstDeviceId, - getHipGpuAgent(srcDeviceId), engineId)); + auto queue = std::make_unique(srcDeviceId, dstDeviceId, getHipGpuAgent(srcDeviceId), + engineId, selection); + lastQueueStatus_ = queue->createStatus(); + if (!queue->valid()) return nullptr; + vec.emplace_back(std::move(queue)); if (channelIdx != nullptr) { *channelIdx = static_cast(vec.size() - 1); } @@ -425,11 +540,64 @@ SdmaQueue* AnvilLib::createSdmaQueue(int srcDeviceId, int dstDeviceId, uint32_t } bool AnvilLib::connect(int srcDeviceId, int dstDeviceId, int numChannels) { - uint32_t engineId = getSdmaEngineId(srcDeviceId, dstDeviceId); + const EngineSelection selection = getSdmaEngineId(srcDeviceId, dstDeviceId); + if (selection.engineId < 0) { + LOG_ERROR("anvil: no SDMA engine mapping for %d -> %d", srcDeviceId, dstDeviceId); + return false; + } + const uint32_t engineId = static_cast(selection.engineId); LOG_TRACE("SDMA: Connect from %d to %d with %d channels using engine %d", srcDeviceId, dstDeviceId, numChannels, engineId); + + // The queue budget is a property of the partition mode rather than of this peer: a CPX partition + // owns one XCD's engines, so the mesh a whole MI300X supports does not fit. KFD also does not + // report how many queues ROCr already holds, so an exhausted budget can only be predicted for + // gross over-subscription and must additionally be recognised when queue creation fails. + // The budget is counted across all engines, not per engine. A rejected engine-pinned create + // retries as a generic queue and reports engine 0, so on a partition every queue charges the + // same key; a per-engine cap would then refuse at numSdmaQueuesPerEngine_ while the partition + // really offers numSdmaEnginesTotal_ times that, which is the figure this message quotes. The + // total is an upper bound, so an uneven distribution is still caught by KFD's own NO_MEMORY. + const uint32_t queueBudget = numSdmaEnginesTotal_ * numSdmaQueuesPerEngine_; + auto reportBudget = [&](uint32_t used) { + LOG_ERROR( + "anvil: SDMA queue budget exhausted: %u queue(s) taken by this process + %d requested, " + "limit %u (%u engines: host=%u xgmi=%u, %u queues per engine). ROCm already holds some of " + "them, so this partition cannot cover this peer count at %d channel(s) per peer. Use fewer " + "ranks per node, fewer channels, or a coarser partition mode (DPX/SPX).", + used, numChannels, queueBudget, numSdmaEnginesTotal_, numSdmaEngines_, numSdmaXgmiEngines_, + numSdmaQueuesPerEngine_, numChannels); + }; + + // The whole request is known up front. Refusing here avoids creating a queue + // that rollback would destroy immediately. + const uint32_t used = queuesUsedTotal_; + if (numChannels > 0 && queueBudget > 0 && + used + static_cast(numChannels) > queueBudget) { + reportBudget(used); + return false; + } + + auto& vec = sdma_channels_[dstDeviceId]; + const size_t already = vec.size(); + auto rollback = [&]() { + while (vec.size() > already) { + SdmaQueue* q = vec.back().get(); + if (q != nullptr && q->valid() && queuesUsedTotal_ > 0) queuesUsedTotal_ -= 1; + vec.pop_back(); + } + }; + for (int c = 0; c < numChannels; ++c) { - createSdmaQueue(srcDeviceId, dstDeviceId, engineId); + SdmaQueue* queue = createSdmaQueue(srcDeviceId, dstDeviceId, engineId, selection); + if (queue == nullptr) { + if (queueBudget > 0 && lastQueueStatus_ == HSAKMT_STATUS_NO_MEMORY) { + reportBudget(queuesUsedTotal_); + } + rollback(); + return false; + } + queuesUsedTotal_ += 1; } return true; } @@ -437,9 +605,21 @@ bool AnvilLib::connect(int srcDeviceId, int dstDeviceId, int numChannels) { void AnvilLib::disconnect() { // Destroy all SDMA queues. SdmaQueue destructor calls hsaKmtDestroyQueue. sdma_channels_.clear(); + queuesUsedTotal_ = 0; LOG_TRACE("SDMA: Disconnected all queues"); } +void AnvilLib::disconnectDevice(int dstDeviceId) { + auto it = sdma_channels_.find(dstDeviceId); + if (it == sdma_channels_.end()) return; + // Only valid queues were charged to the budget in connect(). + for (const auto& q : it->second) { + if (q != nullptr && q->valid() && queuesUsedTotal_ > 0) queuesUsedTotal_ -= 1; + } + sdma_channels_.erase(it); + LOG_TRACE("SDMA: Disconnected queues for device %d", dstDeviceId); +} + SdmaQueue* AnvilLib::getSdmaQueue([[maybe_unused]] int srcDeviceId, int dstDeviceId, int channel_idx) { if (sdma_channels_.find(dstDeviceId) == sdma_channels_.end()) { @@ -461,66 +641,142 @@ AnvilLib& AnvilLib::getInstance() { return *instance; } +// PCI function and the function-0 BDF of the same device. function is -1 when the tail is not a +// function digit, so an unreadable id is not treated as function 0. The physical BDF is rewritten +// only in that case; a bad tail is left unchanged. +struct PciFunctionBus { + std::string busId; + std::string physBusId; + int function; +}; + +static PciFunctionBus pciFunctionBus(const std::string& busId) { + PciFunctionBus loc; + loc.busId = busId; + loc.physBusId = busId; + loc.function = -1; + if (busId.empty()) return loc; + const char c = busId.back(); + if (c >= '0' && c <= '7') { + loc.function = c - '0'; + loc.physBusId.back() = '0'; + } + return loc; +} + int AnvilLib::getOamId(int deviceId) { - std::string busId = getBusId(deviceId); - std::string file_str = "/sys/bus/pci/devices/" + busId + "/xgmi_physical_id"; - std::ifstream file(file_str); - int xgmi_physical_id; - if (file.is_open()) { - if (!(file >> xgmi_physical_id)) { - LOG_ERROR_EXIT("anvil: failed to read xGMI physical id from %s", file_str.c_str()); - } - } else { - LOG_ERROR_EXIT("anvil: failed to open file: %s", file_str.c_str()); + // xgmi_physical_id is a property of the physical GPU. A CPX/DPX partition is a HIP alias at PCI + // function .1-.7 of that GPU and has no sysfs node of its own, so read the physical function + // whenever the partition's own BDF is absent. Returns -1 when neither is readable. + const PciFunctionBus loc = pciFunctionBus(getBusId(deviceId)); + const std::string* candidates[2] = {&loc.busId, &loc.physBusId}; + const int nCandidates = loc.physBusId == loc.busId ? 1 : 2; + for (int i = 0; i < nCandidates; ++i) { + if (candidates[i]->empty()) continue; + std::ifstream file("/sys/bus/pci/devices/" + *candidates[i] + "/xgmi_physical_id"); + int xgmi_physical_id = 0; + if (file.is_open() && (file >> xgmi_physical_id)) return xgmi_physical_id; } - return xgmi_physical_id; + + const int preferredEngine = + selection_.preferredMask == 0 ? -1 : __builtin_ctz(selection_.preferredMask); + // LOG_ERROR so a default ROCSHMEM_DEBUG_LEVEL=ERROR banff repro still shows the preferred-engine + // values that explain why the OAM map was needed. + LOG_ERROR( + "anvil: no xGMI physical id for %s or %s device=%d pair=%d->%d preferredQueried=%d " + "preferredStatus=%#x preferredMask=0x%x preferredEngine=%d hostEng=%u xgmiEng=%u total=%u", + loc.busId.c_str(), loc.physBusId.c_str(), deviceId, selection_.srcDeviceId, + selection_.dstDeviceId, selection_.preferredQueried ? 1 : 0, + static_cast(selection_.preferredStatus), selection_.preferredMask, preferredEngine, + selection_.numSdmaEngines, selection_.numSdmaXgmiEngines, selection_.numSdmaEnginesTotal); + return -1; } int AnvilLib::getSdmaEngineIdFromOamMap(int srcDeviceId, int dstDeviceId) { - auto normalizeOamId = [](int oamId) { - if (oamId < 0) return 0; - return oamId % static_cast(AnvilLib::kMi300xOamMapDim); - }; + const int srcOamId = getOamId(srcDeviceId); + const int dstOamId = getOamId(dstDeviceId); + // Do not clamp an unreadable id to 0: that would make every failed pair select mi300xOamMap[0][0] + // and pin them all on engine 0. Propagate to connect() instead. + if (srcOamId < 0 || dstOamId < 0) return -1; - const int srcOamId = normalizeOamId(getOamId(srcDeviceId)); - const int dstOamId = normalizeOamId(getOamId(dstDeviceId)); + const int srcIdx = srcOamId % static_cast(AnvilLib::kMi300xOamMapDim); + const int dstIdx = dstOamId % static_cast(AnvilLib::kMi300xOamMapDim); + const int oamEngine = mi300xOamMap[static_cast(srcIdx)][static_cast(dstIdx)]; // Use even engines only (MI300X xGMI SDMA layout). - int engineId = - mi300xOamMap[static_cast(srcOamId)][static_cast(dstOamId)] * 2; - - if (numSdmaEnginesTotal_ > 0 && - static_cast(engineId) >= numSdmaEnginesTotal_) { - LOG_WARN("anvil: legacy OAM-map engine %d >= total %u", engineId, numSdmaEnginesTotal_); + int engineId = oamEngine * 2; + + // The map and the doubling assume the 14 xGMI SDMA engines of an unpartitioned MI300X. A + // partition owns fewer engines (CPX: 2 host, 0 xGMI). Same-device peers share an OAM id, so they + // hit the diagonal (oamEngine 0) and engineId 0 is in range. The fold has to run there too, not + // only when the doubled id is past the engine count. + // + // (srcFn + dstFn) rather than dstFn alone: with 2 engines the two split a mixed-parity mesh the + // same way, but once the node has more than 2 engines the sum still separates pairs that share a + // destination function. Two engines cannot give 8 partitions distinct ids. + const bool partition = numSdmaXgmiEngines_ == 0 && numSdmaEnginesTotal_ > 0; + if (oamMapEngineNeedsFold(numSdmaXgmiEngines_, numSdmaEnginesTotal_, engineId)) { + const PciFunctionBus srcPci = pciFunctionBus(getBusId(srcDeviceId)); + const PciFunctionBus dstPci = pciFunctionBus(getBusId(dstDeviceId)); + // An unreadable tail adds 0 and therefore collides with function 0. The warning logs the raw + // BDF so that case is not silent. The info line reports the values the fold actually added. + const int srcFn = srcPci.function < 0 ? 0 : srcPci.function; + const int dstFn = dstPci.function < 0 ? 0 : dstPci.function; + const int folded = foldOamMapEngine(oamEngine, srcFn, dstFn, numSdmaEnginesTotal_); + if (srcPci.function < 0 || dstPci.function < 0) { + LOG_WARN("anvil: PCI function unreadable src=%s dst=%s, using engine %d (oam=%d total=%u)", + srcPci.busId.c_str(), dstPci.busId.c_str(), folded, oamEngine, numSdmaEnginesTotal_); + } else if (partition) { + LOG_INFO("anvil: partition engine %d (oam=%d srcFn=%d dstFn=%d total=%u)", folded, oamEngine, + srcFn, dstFn, numSdmaEnginesTotal_); + } else { + LOG_WARN( + "anvil: legacy OAM-map engine %d >= total %u, using engine %d (oam=%d srcFn=%d dstFn=%d)", + engineId, numSdmaEnginesTotal_, folded, oamEngine, srcFn, dstFn); + } + engineId = folded; } return engineId; } -int AnvilLib::getSdmaEngineId(int srcDeviceId, int dstDeviceId) { +EngineSelection AnvilLib::getSdmaEngineId(int srcDeviceId, int dstDeviceId) { + // Published on the way in: getOamId reads it to report why the map was needed. + selection_ = EngineSelection{}; + selection_.srcDeviceId = srcDeviceId; + selection_.dstDeviceId = dstDeviceId; + selection_.numSdmaEngines = numSdmaEngines_; + selection_.numSdmaXgmiEngines = numSdmaXgmiEngines_; + selection_.numSdmaEnginesTotal = numSdmaEnginesTotal_; + if (srcDeviceId >= 0 && dstDeviceId >= 0 && srcDeviceId < static_cast(gpuAgentsByHipDev_.size()) && dstDeviceId < static_cast(gpuAgentsByHipDev_.size()) && hsaAgentIsValid(gpuAgentsByHipDev_[static_cast(srcDeviceId)]) && hsaAgentIsValid(gpuAgentsByHipDev_[static_cast(dstDeviceId)])) { - uint32_t engineMask = 0; const hsa_agent_t srcAgent = gpuAgentsByHipDev_[static_cast(srcDeviceId)]; const hsa_agent_t dstAgent = gpuAgentsByHipDev_[static_cast(dstDeviceId)]; - const hsa_status_t status = hsa_amd_memory_get_preferred_copy_engine( - dstAgent, srcAgent, &engineMask); - if (status == HSA_STATUS_SUCCESS && engineMask != 0) { - const int engineId = __builtin_ctz(engineMask); + selection_.preferredStatus = + hsa_amd_memory_get_preferred_copy_engine(dstAgent, srcAgent, &selection_.preferredMask); + selection_.preferredQueried = true; + if (selection_.preferredStatus == HSA_STATUS_SUCCESS && selection_.preferredMask != 0) { + const int engineId = __builtin_ctz(selection_.preferredMask); if (engineId >= 0 && (numSdmaEnginesTotal_ == 0 || static_cast(engineId) < numSdmaEnginesTotal_)) { + selection_.usedPreferred = true; + selection_.engineId = engineId; LOG_TRACE("SDMA: HSA preferred engine %d for %d -> %d (mask=0x%x)", engineId, srcDeviceId, - dstDeviceId, engineMask); - return engineId; + dstDeviceId, selection_.preferredMask); + return selection_; } - LOG_WARN("anvil: HSA preferred engine %d out of range (total=%u), using OAM map", engineId, - numSdmaEnginesTotal_); + LOG_WARN("anvil: HSA preferred engine %d out of range (total=%u status=%#x mask=0x%x), " + "using OAM map", + engineId, numSdmaEnginesTotal_, + static_cast(selection_.preferredStatus), selection_.preferredMask); } } - return getSdmaEngineIdFromOamMap(srcDeviceId, dstDeviceId); + selection_.engineId = getSdmaEngineIdFromOamMap(srcDeviceId, dstDeviceId); + return selection_; } AnvilLib& anvil = anvil.getInstance(); diff --git a/projects/rocshmem/src/sdma/anvil.hpp b/projects/rocshmem/src/sdma/anvil.hpp index e8af4289322e..404a8bdfac87 100644 --- a/projects/rocshmem/src/sdma/anvil.hpp +++ b/projects/rocshmem/src/sdma/anvil.hpp @@ -43,23 +43,74 @@ namespace sdma_anvil { +// True when the doubled OAM-map id cannot be used as-is. A partition (no xGMI +// engines) folds even when the doubled id is in range: same-device peers share +// the map diagonal and would otherwise all land on engine 0. An id past the +// engines this node reports folds too. numSdmaEnginesTotal == 0 never folds, +// so the modulo below is not asked to divide by zero. +inline bool oamMapEngineNeedsFold(uint32_t numSdmaXgmiEngines, uint32_t numSdmaEnginesTotal, + int doubledEngineId) { + const bool partition = numSdmaXgmiEngines == 0 && numSdmaEnginesTotal > 0; + const bool outOfRange = + numSdmaEnginesTotal > 0 && static_cast(doubledEngineId) >= numSdmaEnginesTotal; + return partition || outOfRange; +} + +// Fold the undoubled OAM-map value and both PCI functions into the engines this +// node reports. numEngines must be > 0. An unreadable function (-1) counts as 0 +// and therefore collides with function 0. +inline int foldOamMapEngine(int oamEngine, int srcFn, int dstFn, uint32_t numEngines) { + const int src = srcFn < 0 ? 0 : srcFn; + const int dst = dstFn < 0 ? 0 : dstFn; + return (oamEngine + src + dst) % static_cast(numEngines); +} + +// How an engine was chosen for one peer, and the values that explain the choice. The two failure +// logs that report it -- the queue-create error and the missing-xGMI-id error -- need the same +// fields, so they travel as one record: a second copy means the next diagnostic field has to be +// remembered in two places. +struct EngineSelection { + int engineId{-1}; // -1 when no engine could be mapped for this pair + int srcDeviceId{-1}; + int dstDeviceId{-1}; + // Result of hsa_amd_memory_get_preferred_copy_engine. queried stays false when the agents were + // not valid enough to ask, which is different from asking and being refused. + hsa_status_t preferredStatus{HSA_STATUS_ERROR}; + uint32_t preferredMask{0}; + bool preferredQueried{false}; + bool usedPreferred{false}; + // Engine counts this node reports, which bound every id above. + uint32_t numSdmaEngines{0}; + uint32_t numSdmaXgmiEngines{0}; + uint32_t numSdmaEnginesTotal{0}; +}; class SdmaQueue { public: - SdmaQueue(int localDeviceId, int remoteDeviceId, const hsa_agent_t& localAgent, uint32_t engineId); + SdmaQueue(int localDeviceId, int remoteDeviceId, const hsa_agent_t& localAgent, uint32_t engineId, + const EngineSelection& selection); ~SdmaQueue(); SdmaQueueDeviceHandle* deviceHandle() const; SdmaQueueSingleProducerDeviceHandle* singleProducerDeviceHandle() const; void dump(std::ofstream& logFile); + // False when the constructor could not create an SDMA queue. The constructor releases whatever it + // acquired before giving up, so an invalid queue must be dropped without being used or destroyed + // any further. + bool valid() const; + // Status of the queue-creation attempt, so the caller can tell an exhausted queue budget + // (NO_MEMORY) apart from other failures. + HSAKMT_STATUS createStatus() const; private: - int remoteDeviceId_; - uint64_t* cachedWptr_; - uint64_t* committedWptr_; - void* queueBuffer_; - HsaQueueResource queue_; - SdmaQueueDeviceHandle* deviceHandle_; + bool valid_{false}; + HSAKMT_STATUS createStatus_{HSAKMT_STATUS_ERROR}; + int remoteDeviceId_{-1}; + uint64_t* cachedWptr_{nullptr}; + uint64_t* committedWptr_{nullptr}; + void* queueBuffer_{nullptr}; + HsaQueueResource queue_{}; + SdmaQueueDeviceHandle* deviceHandle_{nullptr}; SdmaQueueSingleProducerDeviceHandle* singleProducerDeviceHandle_{nullptr}; }; @@ -80,9 +131,13 @@ class AnvilLib { void init(); bool connect(int srcDeviceId, int dstDeviceId, int numChannels = 1); void disconnect(); + // Drop only the queues held for one destination. The channel map is process-global and a + // second communicator reuses a previous create's queues, so a failed create must not take + // disconnect()'s whole-process path and destroy handles already live on the GPU. + void disconnectDevice(int dstDeviceId); SdmaQueue* getSdmaQueue(int srcDeviceId, int dstDeviceId, int channel_idx = 0); SdmaQueue* createSdmaQueue(int srcDeviceId, int dstDeviceId, uint32_t engineId, - int* channelIdx = nullptr); + const EngineSelection& selection, int* channelIdx = nullptr); private: /* @@ -111,13 +166,29 @@ class AnvilLib { uint32_t numSdmaEngines_{0}; uint32_t numSdmaXgmiEngines_{0}; uint32_t numSdmaEnginesTotal_{0}; + // KFD caps user SDMA queues per engine, so a partition with few engines also has a small total + // queue budget. Track usage to refuse a mesh that cannot fit before KFD returns NO_MEMORY part + // way through building it. + uint32_t numSdmaQueuesPerEngine_{0}; + // Queues already taken by this process, counted across all engines rather than per engine: a + // rejected engine-pinned create retries as a generic queue and reports engine 0, so on a + // partition every queue would charge the same key and a per-engine cap would refuse at a + // fraction of the real budget. connect() compares this against + // numSdmaEnginesTotal_ * numSdmaQueuesPerEngine_. + uint32_t queuesUsedTotal_{0}; + HSAKMT_STATUS lastQueueStatus_{HSAKMT_STATUS_SUCCESS}; + // Selection in progress, so getOamId can report why the map was consulted without the caller + // threading the same values back down. + EngineSelection selection_; void buildGpuAgentMap(); hsa_agent_t getHipGpuAgent(int hipDeviceId) const; void querySdmaEngineCounts(); int getOamId(int deviceId); int getSdmaEngineIdFromOamMap(int srcDeviceId, int dstDeviceId); - int getSdmaEngineId(int srcDeviceId, int dstDeviceId); + // Chooses an engine for the pair and records why, so the caller can hand the same record to + // every queue it then creates. + EngineSelection getSdmaEngineId(int srcDeviceId, int dstDeviceId); std::once_flag init_flag; std::vector gpuAgentsByHipDev_; diff --git a/projects/rocshmem/src/sdma_policy.cpp b/projects/rocshmem/src/sdma_policy.cpp index 0dd1e6ca05f2..7f7d2b7493aa 100644 --- a/projects/rocshmem/src/sdma_policy.cpp +++ b/projects/rocshmem/src/sdma_policy.cpp @@ -64,12 +64,20 @@ __host__ void SdmaImpl::sdmaHostInit(int pe, int num_pes, int rank) { int deviceId; CHECK_HIP(hipGetDevice(&deviceId)); - // Create SDMA connections to all local PEs including self + // Create SDMA connections to all local PEs including self. A false return leaves a partially + // wired mesh whose null handles would silently skip puts above the SDMA threshold, so tear + // everything down and fall back to IPC memcpy instead. for (int i = 0; i < shm_size; i++) { if (i != deviceId) { sdma_anvil::EnablePeerAccess(deviceId, i); } - sdma_anvil::anvil.connect(deviceId, i, numChannels); + if (!sdma_anvil::anvil.connect(deviceId, i, numChannels)) { + LOG_ERROR("SDMA: connect failed from device %d to %d with %d channel(s); disabling SDMA", + deviceId, i, numChannels); + sdma_anvil::anvil.disconnect(); + sdmaEnabled = false; + return; + } } // Total number of handles: shm_size * numChannels From 904f4ba88b26a1b831a3b4f718ed609f2291fa1b Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Mon, 5 Oct 2026 17:42:37 -0400 Subject: [PATCH 02/10] Expose pciFunctionBus and the queue budget for testing Both helpers decided behaviour that no test could reach. pciFunctionBus was a file-static in anvil.cpp. It is pure and it picks the same fallback as oamMapEngineNeedsFold and foldOamMapEngine -- a tail outside '0'-'7' leaves function at -1, which collapses getOamId's candidate list and skips the physical-BDF read -- so it belongs beside them in the header as inline rather than hidden in the translation unit. The per-process queue count was only observable through a log line, which left connect()'s refusal and disconnectDevice()'s teardown with nothing to assert against: a refused connect has to leave the count unchanged and a fully abandoned create has to return it to where it started. queuesUsed() makes that readable without widening anything else. No functional change. Co-authored-by: Cursor --- projects/rocshmem/src/sdma/anvil.cpp | 23 ------------------- projects/rocshmem/src/sdma/anvil.hpp | 34 ++++++++++++++++++++++++++++ 2 files changed, 34 insertions(+), 23 deletions(-) diff --git a/projects/rocshmem/src/sdma/anvil.cpp b/projects/rocshmem/src/sdma/anvil.cpp index 87a4dce45256..5fdece715066 100644 --- a/projects/rocshmem/src/sdma/anvil.cpp +++ b/projects/rocshmem/src/sdma/anvil.cpp @@ -641,29 +641,6 @@ AnvilLib& AnvilLib::getInstance() { return *instance; } -// PCI function and the function-0 BDF of the same device. function is -1 when the tail is not a -// function digit, so an unreadable id is not treated as function 0. The physical BDF is rewritten -// only in that case; a bad tail is left unchanged. -struct PciFunctionBus { - std::string busId; - std::string physBusId; - int function; -}; - -static PciFunctionBus pciFunctionBus(const std::string& busId) { - PciFunctionBus loc; - loc.busId = busId; - loc.physBusId = busId; - loc.function = -1; - if (busId.empty()) return loc; - const char c = busId.back(); - if (c >= '0' && c <= '7') { - loc.function = c - '0'; - loc.physBusId.back() = '0'; - } - return loc; -} - int AnvilLib::getOamId(int deviceId) { // xgmi_physical_id is a property of the physical GPU. A CPX/DPX partition is a HIP alias at PCI // function .1-.7 of that GPU and has no sysfs node of its own, so read the physical function diff --git a/projects/rocshmem/src/sdma/anvil.hpp b/projects/rocshmem/src/sdma/anvil.hpp index 404a8bdfac87..58a811626960 100644 --- a/projects/rocshmem/src/sdma/anvil.hpp +++ b/projects/rocshmem/src/sdma/anvil.hpp @@ -33,6 +33,7 @@ #include #include #include +#include #include #include @@ -65,6 +66,34 @@ inline int foldOamMapEngine(int oamEngine, int srcFn, int dstFn, uint32_t numEng return (oamEngine + src + dst) % static_cast(numEngines); } +// PCI function and the function-0 BDF of the same device. function is -1 when the tail is not a +// function digit, so an unreadable id is not treated as function 0. The physical BDF is rewritten +// only in that case; a bad tail is left unchanged. +struct PciFunctionBus { + std::string busId; + std::string physBusId; + int function; +}; + +// Split a BDF into its function digit and the function-0 BDF of the same device. Pure, and beside +// the two helpers above because it decides the same fallback they do: a tail outside '0'-'7' +// leaves function at -1 and physBusId equal to busId, which collapses getOamId's candidate list to +// one entry so the physical-BDF read never runs. PCI function numbers are three bits, so '0'-'7' +// is the whole range. +inline PciFunctionBus pciFunctionBus(const std::string& busId) { + PciFunctionBus loc; + loc.busId = busId; + loc.physBusId = busId; + loc.function = -1; + if (busId.empty()) return loc; + const char c = busId.back(); + if (c >= '0' && c <= '7') { + loc.function = c - '0'; + loc.physBusId.back() = '0'; + } + return loc; +} + // How an engine was chosen for one peer, and the values that explain the choice. The two failure // logs that report it -- the queue-create error and the missing-xGMI-id error -- need the same // fields, so they travel as one record: a second copy means the next diagnostic field has to be @@ -135,6 +164,11 @@ class AnvilLib { // second communicator reuses a previous create's queues, so a failed create must not take // disconnect()'s whole-process path and destroy handles already live on the GPU. void disconnectDevice(int dstDeviceId); + // Queues this process currently holds, counted across all engines. The budget is otherwise only + // visible in a log line, which leaves the refusal in connect() and the teardown in + // disconnectDevice() with nothing a test can assert: a refused connect must leave this unchanged + // and a fully abandoned create must bring it back to where it started. + uint32_t queuesUsed() const { return queuesUsedTotal_; } SdmaQueue* getSdmaQueue(int srcDeviceId, int dstDeviceId, int channel_idx = 0); SdmaQueue* createSdmaQueue(int srcDeviceId, int dstDeviceId, uint32_t engineId, const EngineSelection& selection, int* channelIdx = nullptr); From 25264383827c62c39d6d296ca2cbaa3bd208cc58 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Tue, 6 Oct 2026 22:09:56 -0400 Subject: [PATCH 03/10] ci: retrigger math-ci after a Docker Hub registry timeout The rocshmem precheckin Docker stage failed on hyd-7c-zt05-02 with 'auth.docker.io/token: context deadline exceeded' 1m47s into a 5m timeout. No code involved; the other node built rocSHMEM to 94% on the same commit. Co-authored-by: Cursor From 61c5ab98ceb8b5c90592c41217830354a8d16770 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Wed, 7 Oct 2026 22:30:03 -0400 Subject: [PATCH 04/10] Address Argus review on the Anvil CPX change - connect() refuses a source device with no mapped HSA agent instead of reaching getHipGpuAgent's process exit through createSdmaQueue. - The queue-create error logs the generic retry's own status. - The post-create budget hint no longer depends on a known budget. - sdmaHostInit still allocates the (all-null) handle array after a failed connect, since the USE_SDMA testers index it unguarded. - isSdmaPartition() is the single copy of the partition predicate. - Align oamMapEngineNeedsFold's continuation line. Co-authored-by: Cursor --- projects/rocshmem/src/sdma/anvil.cpp | 20 ++++++++++++++------ projects/rocshmem/src/sdma/anvil.hpp | 9 +++++++-- projects/rocshmem/src/sdma_policy.cpp | 5 +++-- 3 files changed, 24 insertions(+), 10 deletions(-) diff --git a/projects/rocshmem/src/sdma/anvil.cpp b/projects/rocshmem/src/sdma/anvil.cpp index 5fdece715066..bf1ed6c8afa3 100644 --- a/projects/rocshmem/src/sdma/anvil.cpp +++ b/projects/rocshmem/src/sdma/anvil.cpp @@ -232,10 +232,11 @@ SdmaQueue::SdmaQueue(int localDeviceId, int remoteDeviceId, const hsa_agent_t& l const std::string srcBus = getBusId(localDeviceId); const std::string dstBus = getBusId(remoteDeviceId); LOG_ERROR( - "anvil: hsaKmtCreateQueueExt failed hsakmt=%d (%s) node=%u engineId=%u srcDev=%d (%s) " - "dstDev=%d (%s) usedPreferred=%d preferredStatus=%#x preferredMask=0x%x hostEng=%u " - "xgmiEng=%u total=%u", - static_cast(createStatus_), hsakmtStatusName(createStatus_), localNodeId, engineId, + "anvil: hsaKmtCreateQueueExt failed hsakmt=%d (%s) genericRetry=%d (%s) node=%u " + "engineId=%u srcDev=%d (%s) dstDev=%d (%s) usedPreferred=%d preferredStatus=%#x " + "preferredMask=0x%x hostEng=%u xgmiEng=%u total=%u", + static_cast(createStatus_), hsakmtStatusName(createStatus_), + static_cast(queueStatus), hsakmtStatusName(queueStatus), localNodeId, engineId, localDeviceId, srcBus.c_str(), remoteDeviceId, dstBus.c_str(), selection.usedPreferred ? 1 : 0, static_cast(selection.preferredStatus), selection.preferredMask, selection.numSdmaEngines, selection.numSdmaXgmiEngines, @@ -540,6 +541,13 @@ SdmaQueue* AnvilLib::createSdmaQueue(int srcDeviceId, int dstDeviceId, uint32_t } bool AnvilLib::connect(int srcDeviceId, int dstDeviceId, int numChannels) { + // createSdmaQueue's getHipGpuAgent exits the process on an unmapped device; refuse here instead. + if (srcDeviceId < 0 || srcDeviceId >= static_cast(gpuAgentsByHipDev_.size()) || + !hsaAgentIsValid(gpuAgentsByHipDev_[static_cast(srcDeviceId)])) { + LOG_ERROR("anvil: no HSA agent mapped for HIP device %d, cannot connect to %d", srcDeviceId, + dstDeviceId); + return false; + } const EngineSelection selection = getSdmaEngineId(srcDeviceId, dstDeviceId); if (selection.engineId < 0) { LOG_ERROR("anvil: no SDMA engine mapping for %d -> %d", srcDeviceId, dstDeviceId); @@ -591,7 +599,7 @@ bool AnvilLib::connect(int srcDeviceId, int dstDeviceId, int numChannels) { for (int c = 0; c < numChannels; ++c) { SdmaQueue* queue = createSdmaQueue(srcDeviceId, dstDeviceId, engineId, selection); if (queue == nullptr) { - if (queueBudget > 0 && lastQueueStatus_ == HSAKMT_STATUS_NO_MEMORY) { + if (lastQueueStatus_ == HSAKMT_STATUS_NO_MEMORY) { reportBudget(queuesUsedTotal_); } rollback(); @@ -691,7 +699,7 @@ int AnvilLib::getSdmaEngineIdFromOamMap(int srcDeviceId, int dstDeviceId) { // (srcFn + dstFn) rather than dstFn alone: with 2 engines the two split a mixed-parity mesh the // same way, but once the node has more than 2 engines the sum still separates pairs that share a // destination function. Two engines cannot give 8 partitions distinct ids. - const bool partition = numSdmaXgmiEngines_ == 0 && numSdmaEnginesTotal_ > 0; + const bool partition = isSdmaPartition(numSdmaXgmiEngines_, numSdmaEnginesTotal_); if (oamMapEngineNeedsFold(numSdmaXgmiEngines_, numSdmaEnginesTotal_, engineId)) { const PciFunctionBus srcPci = pciFunctionBus(getBusId(srcDeviceId)); const PciFunctionBus dstPci = pciFunctionBus(getBusId(dstDeviceId)); diff --git a/projects/rocshmem/src/sdma/anvil.hpp b/projects/rocshmem/src/sdma/anvil.hpp index 58a811626960..5d3b3e25f0a0 100644 --- a/projects/rocshmem/src/sdma/anvil.hpp +++ b/projects/rocshmem/src/sdma/anvil.hpp @@ -44,14 +44,19 @@ namespace sdma_anvil { +// A CPX/DPX partition owns host SDMA engines but no xGMI engines. +inline bool isSdmaPartition(uint32_t numSdmaXgmiEngines, uint32_t numSdmaEnginesTotal) { + return numSdmaXgmiEngines == 0 && numSdmaEnginesTotal > 0; +} + // True when the doubled OAM-map id cannot be used as-is. A partition (no xGMI // engines) folds even when the doubled id is in range: same-device peers share // the map diagonal and would otherwise all land on engine 0. An id past the // engines this node reports folds too. numSdmaEnginesTotal == 0 never folds, // so the modulo below is not asked to divide by zero. inline bool oamMapEngineNeedsFold(uint32_t numSdmaXgmiEngines, uint32_t numSdmaEnginesTotal, - int doubledEngineId) { - const bool partition = numSdmaXgmiEngines == 0 && numSdmaEnginesTotal > 0; + int doubledEngineId) { + const bool partition = isSdmaPartition(numSdmaXgmiEngines, numSdmaEnginesTotal); const bool outOfRange = numSdmaEnginesTotal > 0 && static_cast(doubledEngineId) >= numSdmaEnginesTotal; return partition || outOfRange; diff --git a/projects/rocshmem/src/sdma_policy.cpp b/projects/rocshmem/src/sdma_policy.cpp index 7f7d2b7493aa..bf3aea59ebbd 100644 --- a/projects/rocshmem/src/sdma_policy.cpp +++ b/projects/rocshmem/src/sdma_policy.cpp @@ -66,7 +66,8 @@ __host__ void SdmaImpl::sdmaHostInit(int pe, int num_pes, int rank) { // Create SDMA connections to all local PEs including self. A false return leaves a partially // wired mesh whose null handles would silently skip puts above the SDMA threshold, so tear - // everything down and fall back to IPC memcpy instead. + // everything down and fall back to IPC memcpy instead. The handle array below is still allocated + // (all null after disconnect) because USE_SDMA testers index it without checking sdmaEnabled. for (int i = 0; i < shm_size; i++) { if (i != deviceId) { sdma_anvil::EnablePeerAccess(deviceId, i); @@ -76,7 +77,7 @@ __host__ void SdmaImpl::sdmaHostInit(int pe, int num_pes, int rank) { deviceId, i, numChannels); sdma_anvil::anvil.disconnect(); sdmaEnabled = false; - return; + break; } } From c87f3767b45e3a6e2684aece03ae812f2ecb7d0c Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Thu, 8 Oct 2026 15:09:12 -0400 Subject: [PATCH 05/10] Test the CPX engine fold and harden AnvilLib inputs - Move the pure engine-selection helpers into anvil_engine_map.hpp, which has no HSA/KFD dependency, and cover them with an unconditional rocSHMEM gtest so CI exercises the partition fold without USE_SDMA. - A failed SDMA engine-count query now reports all counts as unknown (total 0) at ERROR level instead of reading as a CPX partition. - createSdmaQueue charges the queue budget itself, so direct callers stay consistent with disconnectDevice. - connect() refuses an out-of-range destination device before getBusId can exit the process. Co-authored-by: Cursor --- projects/rocshmem/src/sdma/anvil.cpp | 34 +++--- projects/rocshmem/src/sdma/anvil.hpp | 56 +-------- .../rocshmem/src/sdma/anvil_engine_map.hpp | 93 ++++++++++++++ .../rocshmem/tests/unit_tests/CMakeLists.txt | 1 + .../unit_tests/anvil_engine_map_gtest.cpp | 113 ++++++++++++++++++ 5 files changed, 228 insertions(+), 69 deletions(-) create mode 100644 projects/rocshmem/src/sdma/anvil_engine_map.hpp create mode 100644 projects/rocshmem/tests/unit_tests/anvil_engine_map_gtest.cpp diff --git a/projects/rocshmem/src/sdma/anvil.cpp b/projects/rocshmem/src/sdma/anvil.cpp index bf1ed6c8afa3..6de34a6d7010 100644 --- a/projects/rocshmem/src/sdma/anvil.cpp +++ b/projects/rocshmem/src/sdma/anvil.cpp @@ -470,23 +470,23 @@ void AnvilLib::querySdmaEngineCounts() { LOG_WARN("anvil: no mapped HIP GPU agents; SDMA engine count unknown"); return; } - hsa_status_t status = hsa_agent_get_info( + const hsa_status_t hostStatus = hsa_agent_get_info( agent, static_cast(HSA_AMD_AGENT_INFO_NUM_SDMA_ENG), &numSdmaEngines_); - if (status != HSA_STATUS_SUCCESS) { - LOG_WARN("anvil: HSA_AMD_AGENT_INFO_NUM_SDMA_ENG query failed: %#x", status); - numSdmaEngines_ = 0; - } - - status = hsa_agent_get_info( + const hsa_status_t xgmiStatus = hsa_agent_get_info( agent, static_cast(HSA_AMD_AGENT_INFO_NUM_SDMA_XGMI_ENG), &numSdmaXgmiEngines_); - if (status != HSA_STATUS_SUCCESS) { - LOG_WARN("anvil: HSA_AMD_AGENT_INFO_NUM_SDMA_XGMI_ENG query failed: %#x", status); + if (hostStatus != HSA_STATUS_SUCCESS || xgmiStatus != HSA_STATUS_SUCCESS) { + // A zero count from a failed query would read as a CPX partition (no xGMI engines) and fold a + // full SPX node onto the host engines. A total of 0 means "unknown": no fold, no budget check. + LOG_ERROR("anvil: SDMA engine count query failed (host=%#x xgmi=%#x); engine counts unknown", + hostStatus, xgmiStatus); + numSdmaEngines_ = 0; numSdmaXgmiEngines_ = 0; + numSdmaEnginesTotal_ = 0; + } else { + numSdmaEnginesTotal_ = numSdmaEngines_ + numSdmaXgmiEngines_; } - numSdmaEnginesTotal_ = numSdmaEngines_ + numSdmaXgmiEngines_; - uint32_t node = 0; HsaNodeProperties nodeProps{}; if (hsa_agent_get_info(agent, HSA_AGENT_INFO_NODE, &node) == HSA_STATUS_SUCCESS && @@ -534,6 +534,7 @@ SdmaQueue* AnvilLib::createSdmaQueue(int srcDeviceId, int dstDeviceId, uint32_t lastQueueStatus_ = queue->createStatus(); if (!queue->valid()) return nullptr; vec.emplace_back(std::move(queue)); + queuesUsedTotal_ += 1; if (channelIdx != nullptr) { *channelIdx = static_cast(vec.size() - 1); } @@ -541,13 +542,19 @@ SdmaQueue* AnvilLib::createSdmaQueue(int srcDeviceId, int dstDeviceId, uint32_t } bool AnvilLib::connect(int srcDeviceId, int dstDeviceId, int numChannels) { - // createSdmaQueue's getHipGpuAgent exits the process on an unmapped device; refuse here instead. + // createSdmaQueue's getHipGpuAgent exits the process on an unmapped device, and getBusId exits on + // an out-of-range one; refuse here instead. if (srcDeviceId < 0 || srcDeviceId >= static_cast(gpuAgentsByHipDev_.size()) || !hsaAgentIsValid(gpuAgentsByHipDev_[static_cast(srcDeviceId)])) { LOG_ERROR("anvil: no HSA agent mapped for HIP device %d, cannot connect to %d", srcDeviceId, dstDeviceId); return false; } + if (dstDeviceId < 0 || dstDeviceId >= static_cast(gpuAgentsByHipDev_.size())) { + LOG_ERROR("anvil: destination HIP device %d out of range [0, %zu), cannot connect from %d", + dstDeviceId, gpuAgentsByHipDev_.size(), srcDeviceId); + return false; + } const EngineSelection selection = getSdmaEngineId(srcDeviceId, dstDeviceId); if (selection.engineId < 0) { LOG_ERROR("anvil: no SDMA engine mapping for %d -> %d", srcDeviceId, dstDeviceId); @@ -605,7 +612,6 @@ bool AnvilLib::connect(int srcDeviceId, int dstDeviceId, int numChannels) { rollback(); return false; } - queuesUsedTotal_ += 1; } return true; } @@ -620,7 +626,7 @@ void AnvilLib::disconnect() { void AnvilLib::disconnectDevice(int dstDeviceId) { auto it = sdma_channels_.find(dstDeviceId); if (it == sdma_channels_.end()) return; - // Only valid queues were charged to the budget in connect(). + // Only valid queues were charged to the budget in createSdmaQueue(). for (const auto& q : it->second) { if (q != nullptr && q->valid() && queuesUsedTotal_ > 0) queuesUsedTotal_ -= 1; } diff --git a/projects/rocshmem/src/sdma/anvil.hpp b/projects/rocshmem/src/sdma/anvil.hpp index 5d3b3e25f0a0..022b8c91390f 100644 --- a/projects/rocshmem/src/sdma/anvil.hpp +++ b/projects/rocshmem/src/sdma/anvil.hpp @@ -38,67 +38,13 @@ #include #include "anvil_device.hpp" +#include "anvil_engine_map.hpp" #include "hsa/hsa_ext_amd.h" #include "hsakmt/hsakmt.h" #include "hsakmt/hsakmttypes.h" namespace sdma_anvil { -// A CPX/DPX partition owns host SDMA engines but no xGMI engines. -inline bool isSdmaPartition(uint32_t numSdmaXgmiEngines, uint32_t numSdmaEnginesTotal) { - return numSdmaXgmiEngines == 0 && numSdmaEnginesTotal > 0; -} - -// True when the doubled OAM-map id cannot be used as-is. A partition (no xGMI -// engines) folds even when the doubled id is in range: same-device peers share -// the map diagonal and would otherwise all land on engine 0. An id past the -// engines this node reports folds too. numSdmaEnginesTotal == 0 never folds, -// so the modulo below is not asked to divide by zero. -inline bool oamMapEngineNeedsFold(uint32_t numSdmaXgmiEngines, uint32_t numSdmaEnginesTotal, - int doubledEngineId) { - const bool partition = isSdmaPartition(numSdmaXgmiEngines, numSdmaEnginesTotal); - const bool outOfRange = - numSdmaEnginesTotal > 0 && static_cast(doubledEngineId) >= numSdmaEnginesTotal; - return partition || outOfRange; -} - -// Fold the undoubled OAM-map value and both PCI functions into the engines this -// node reports. numEngines must be > 0. An unreadable function (-1) counts as 0 -// and therefore collides with function 0. -inline int foldOamMapEngine(int oamEngine, int srcFn, int dstFn, uint32_t numEngines) { - const int src = srcFn < 0 ? 0 : srcFn; - const int dst = dstFn < 0 ? 0 : dstFn; - return (oamEngine + src + dst) % static_cast(numEngines); -} - -// PCI function and the function-0 BDF of the same device. function is -1 when the tail is not a -// function digit, so an unreadable id is not treated as function 0. The physical BDF is rewritten -// only in that case; a bad tail is left unchanged. -struct PciFunctionBus { - std::string busId; - std::string physBusId; - int function; -}; - -// Split a BDF into its function digit and the function-0 BDF of the same device. Pure, and beside -// the two helpers above because it decides the same fallback they do: a tail outside '0'-'7' -// leaves function at -1 and physBusId equal to busId, which collapses getOamId's candidate list to -// one entry so the physical-BDF read never runs. PCI function numbers are three bits, so '0'-'7' -// is the whole range. -inline PciFunctionBus pciFunctionBus(const std::string& busId) { - PciFunctionBus loc; - loc.busId = busId; - loc.physBusId = busId; - loc.function = -1; - if (busId.empty()) return loc; - const char c = busId.back(); - if (c >= '0' && c <= '7') { - loc.function = c - '0'; - loc.physBusId.back() = '0'; - } - return loc; -} - // How an engine was chosen for one peer, and the values that explain the choice. The two failure // logs that report it -- the queue-create error and the missing-xGMI-id error -- need the same // fields, so they travel as one record: a second copy means the next diagnostic field has to be diff --git a/projects/rocshmem/src/sdma/anvil_engine_map.hpp b/projects/rocshmem/src/sdma/anvil_engine_map.hpp new file mode 100644 index 000000000000..a3b39318f73a --- /dev/null +++ b/projects/rocshmem/src/sdma/anvil_engine_map.hpp @@ -0,0 +1,93 @@ +/****************************************************************************** + * Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ +#ifndef LIBRARY_SRC_SDMA_ANVIL_ENGINE_MAP_HPP_ +#define LIBRARY_SRC_SDMA_ANVIL_ENGINE_MAP_HPP_ + +// Pure engine-selection helpers for AnvilLib. Kept free of HSA/KFD headers so the unit tests can +// build them without USE_SDMA. + +#include +#include + +namespace sdma_anvil { + +// A CPX/DPX partition owns host SDMA engines but no xGMI engines. Callers pass a total of 0 when +// the engine counts could not be queried, so a failed query is never read as a partition. +inline bool isSdmaPartition(uint32_t numSdmaXgmiEngines, uint32_t numSdmaEnginesTotal) { + return numSdmaXgmiEngines == 0 && numSdmaEnginesTotal > 0; +} + +// True when the doubled OAM-map id cannot be used as-is. A partition (no xGMI +// engines) folds even when the doubled id is in range: same-device peers share +// the map diagonal and would otherwise all land on engine 0. An id past the +// engines this node reports folds too. numSdmaEnginesTotal == 0 never folds, +// so the modulo below is not asked to divide by zero. +inline bool oamMapEngineNeedsFold(uint32_t numSdmaXgmiEngines, uint32_t numSdmaEnginesTotal, + int doubledEngineId) { + const bool partition = isSdmaPartition(numSdmaXgmiEngines, numSdmaEnginesTotal); + const bool outOfRange = + numSdmaEnginesTotal > 0 && static_cast(doubledEngineId) >= numSdmaEnginesTotal; + return partition || outOfRange; +} + +// Fold the undoubled OAM-map value and both PCI functions into the engines this +// node reports. numEngines must be > 0. An unreadable function (-1) counts as 0 +// and therefore collides with function 0. +inline int foldOamMapEngine(int oamEngine, int srcFn, int dstFn, uint32_t numEngines) { + const int src = srcFn < 0 ? 0 : srcFn; + const int dst = dstFn < 0 ? 0 : dstFn; + return (oamEngine + src + dst) % static_cast(numEngines); +} + +// PCI function and the function-0 BDF of the same device. function is -1 when the tail is not a +// function digit, so an unreadable id is not treated as function 0. The physical BDF is rewritten +// only in that case; a bad tail is left unchanged. +struct PciFunctionBus { + std::string busId; + std::string physBusId; + int function; +}; + +// Split a BDF into its function digit and the function-0 BDF of the same device. Pure, and beside +// the two helpers above because it decides the same fallback they do: a tail outside '0'-'7' +// leaves function at -1 and physBusId equal to busId, which collapses getOamId's candidate list to +// one entry so the physical-BDF read never runs. PCI function numbers are three bits, so '0'-'7' +// is the whole range. +inline PciFunctionBus pciFunctionBus(const std::string& busId) { + PciFunctionBus loc; + loc.busId = busId; + loc.physBusId = busId; + loc.function = -1; + if (busId.empty()) return loc; + const char c = busId.back(); + if (c >= '0' && c <= '7') { + loc.function = c - '0'; + loc.physBusId.back() = '0'; + } + return loc; +} + +} // namespace sdma_anvil + +#endif // LIBRARY_SRC_SDMA_ANVIL_ENGINE_MAP_HPP_ diff --git a/projects/rocshmem/tests/unit_tests/CMakeLists.txt b/projects/rocshmem/tests/unit_tests/CMakeLists.txt index b6f0b2902611..8d6ffc2ca7ce 100644 --- a/projects/rocshmem/tests/unit_tests/CMakeLists.txt +++ b/projects/rocshmem/tests/unit_tests/CMakeLists.txt @@ -64,6 +64,7 @@ target_sources( atomic_wf_queue_gtest.cpp envvar_gtest.cpp backend_bc_buffer_register_gtest.cpp + anvil_engine_map_gtest.cpp ) if (USE_IPC) diff --git a/projects/rocshmem/tests/unit_tests/anvil_engine_map_gtest.cpp b/projects/rocshmem/tests/unit_tests/anvil_engine_map_gtest.cpp new file mode 100644 index 000000000000..b4565beaf903 --- /dev/null +++ b/projects/rocshmem/tests/unit_tests/anvil_engine_map_gtest.cpp @@ -0,0 +1,113 @@ +/****************************************************************************** + * Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#include "gtest/gtest.h" +#include "../src/sdma/anvil_engine_map.hpp" + +using namespace sdma_anvil; + +namespace { + +// MI300X engine counts as ROCr reports them. +constexpr uint32_t kSpxXgmi = 14; +constexpr uint32_t kSpxTotal = 16; +constexpr uint32_t kCpxXgmi = 0; +constexpr uint32_t kCpxTotal = 2; + +} // namespace + +TEST(AnvilEngineMap, PartitionIsNoXgmiEnginesWithKnownTotal) { + EXPECT_TRUE(isSdmaPartition(kCpxXgmi, kCpxTotal)); + EXPECT_FALSE(isSdmaPartition(kSpxXgmi, kSpxTotal)); + // A failed engine-count query reports a total of 0 and must not be read as a partition. + EXPECT_FALSE(isSdmaPartition(0, 0)); +} + +TEST(AnvilEngineMap, PartitionFoldsEvenWhenDoubledIdIsInRange) { + // Same-device CPX peers share an OAM id and hit the map diagonal, doubled id 0. Without the + // partition fold they would all be pinned to engine 0. + EXPECT_TRUE(oamMapEngineNeedsFold(kCpxXgmi, kCpxTotal, 0)); + EXPECT_TRUE(oamMapEngineNeedsFold(kCpxXgmi, kCpxTotal, 1)); +} + +TEST(AnvilEngineMap, SpxUsesInRangeIdsAsIs) { + for (int id = 0; id < static_cast(kSpxTotal); id += 2) { + EXPECT_FALSE(oamMapEngineNeedsFold(kSpxXgmi, kSpxTotal, id)) << "id=" << id; + } +} + +TEST(AnvilEngineMap, OutOfRangeIdFolds) { + EXPECT_TRUE(oamMapEngineNeedsFold(kSpxXgmi, kSpxTotal, static_cast(kSpxTotal))); + EXPECT_TRUE(oamMapEngineNeedsFold(2, 4, 12)); +} + +TEST(AnvilEngineMap, UnknownTotalNeverFolds) { + EXPECT_FALSE(oamMapEngineNeedsFold(0, 0, 0)); + EXPECT_FALSE(oamMapEngineNeedsFold(0, 0, 12)); +} + +TEST(AnvilEngineMap, FoldStaysInRangeForEveryInput) { + for (int oam = 0; oam < 8; ++oam) { + for (int src = -1; src < 8; ++src) { + for (int dst = -1; dst < 8; ++dst) { + const int e = foldOamMapEngine(oam, src, dst, kCpxTotal); + EXPECT_GE(e, 0); + EXPECT_LT(e, static_cast(kCpxTotal)); + } + } + } +} + +TEST(AnvilEngineMap, FoldSeparatesPartitionsOfOneDevice) { + // Diagonal (oam 0) pairs from function 0 to functions 0 and 1 land on different engines. + EXPECT_NE(foldOamMapEngine(0, 0, 0, kCpxTotal), foldOamMapEngine(0, 0, 1, kCpxTotal)); +} + +TEST(AnvilEngineMap, FoldTreatsUnreadableFunctionAsZero) { + EXPECT_EQ(foldOamMapEngine(3, -1, -1, kCpxTotal), foldOamMapEngine(3, 0, 0, kCpxTotal)); +} + +TEST(AnvilEngineMap, PciFunctionBusRewritesPartitionToFunctionZero) { + const PciFunctionBus loc = pciFunctionBus("0000:05:00.1"); + EXPECT_EQ(loc.function, 1); + EXPECT_EQ(loc.busId, "0000:05:00.1"); + EXPECT_EQ(loc.physBusId, "0000:05:00.0"); +} + +TEST(AnvilEngineMap, PciFunctionBusAcceptsOnlyThreeBitFunctions) { + for (char c = '0'; c <= '7'; ++c) { + EXPECT_EQ(pciFunctionBus(std::string("0000:05:00.") + c).function, c - '0'); + } + for (const char* bad : {"0000:05:00.8", "0000:05:00.9", "0000:05:00.a"}) { + const PciFunctionBus loc = pciFunctionBus(bad); + EXPECT_EQ(loc.function, -1) << bad; + EXPECT_EQ(loc.physBusId, loc.busId) << bad; + } +} + +TEST(AnvilEngineMap, PciFunctionBusHandlesEmptyInput) { + const PciFunctionBus loc = pciFunctionBus(""); + EXPECT_EQ(loc.function, -1); + EXPECT_TRUE(loc.physBusId.empty()); +} From 97738bfd30a019db694c5cd3ae69716846e3db43 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Thu, 8 Oct 2026 16:44:39 -0400 Subject: [PATCH 06/10] Scope SDMA init teardown and test the queue budget - sdmaHostInit releases only the peers it wired via disconnectDevice() instead of clearing the process-global channel map. - Move the budget comparison into anvil_engine_map.hpp as queueBudgetExceeded() and cover it in the unit test. - Pass raw PCI functions to foldOamMapEngine so its clamp is the one that runs, and make the test sensitive to it. - One releaseQueues() helper credits the budget for both rollback and disconnectDevice(). - Both engine-selection failure logs share describeSelection(), so the queue-create error now also reports preferredQueried. - Trim the budget comment in connect() to what that code needs. Co-authored-by: Cursor --- projects/rocshmem/src/sdma/anvil.cpp | 76 +++++++++---------- projects/rocshmem/src/sdma/anvil.hpp | 2 + .../rocshmem/src/sdma/anvil_engine_map.hpp | 13 +++- projects/rocshmem/src/sdma_policy.cpp | 13 ++-- .../unit_tests/anvil_engine_map_gtest.cpp | 20 +++++ 5 files changed, 76 insertions(+), 48 deletions(-) diff --git a/projects/rocshmem/src/sdma/anvil.cpp b/projects/rocshmem/src/sdma/anvil.cpp index 6de34a6d7010..18b69bfa1fce 100644 --- a/projects/rocshmem/src/sdma/anvil.cpp +++ b/projects/rocshmem/src/sdma/anvil.cpp @@ -96,6 +96,19 @@ std::vector gpuAgents_; static bool hsaAgentIsValid(const hsa_agent_t& agent) { return agent.handle != 0; } +// The engine-choice fields shared by every failure log that reports an EngineSelection. +static std::string describeSelection(const EngineSelection& s) { + const int preferredEngine = s.preferredMask == 0 ? -1 : __builtin_ctz(s.preferredMask); + char buf[256]; + snprintf(buf, sizeof(buf), + "pair=%d->%d preferredQueried=%d preferredStatus=%#x preferredMask=0x%x " + "preferredEngine=%d usedPreferred=%d hostEng=%u xgmiEng=%u total=%u", + s.srcDeviceId, s.dstDeviceId, s.preferredQueried ? 1 : 0, + static_cast(s.preferredStatus), s.preferredMask, preferredEngine, + s.usedPreferred ? 1 : 0, s.numSdmaEngines, s.numSdmaXgmiEngines, s.numSdmaEnginesTotal); + return buf; +} + static std::string hsaAgentBusId(const hsa_agent_t& agent) { uint32_t domain = 0; uint32_t bdfid = 0; @@ -233,14 +246,11 @@ SdmaQueue::SdmaQueue(int localDeviceId, int remoteDeviceId, const hsa_agent_t& l const std::string dstBus = getBusId(remoteDeviceId); LOG_ERROR( "anvil: hsaKmtCreateQueueExt failed hsakmt=%d (%s) genericRetry=%d (%s) node=%u " - "engineId=%u srcDev=%d (%s) dstDev=%d (%s) usedPreferred=%d preferredStatus=%#x " - "preferredMask=0x%x hostEng=%u xgmiEng=%u total=%u", + "engineId=%u srcDev=%d (%s) dstDev=%d (%s) %s", static_cast(createStatus_), hsakmtStatusName(createStatus_), static_cast(queueStatus), hsakmtStatusName(queueStatus), localNodeId, engineId, localDeviceId, srcBus.c_str(), remoteDeviceId, dstBus.c_str(), - selection.usedPreferred ? 1 : 0, static_cast(selection.preferredStatus), - selection.preferredMask, selection.numSdmaEngines, selection.numSdmaXgmiEngines, - selection.numSdmaEnginesTotal); + describeSelection(selection).c_str()); // Leave valid_ false so the caller can drop the Anvil backend instead of killing the job. hsaKmtUnmapMemoryToGPU(queueBuffer_); hsaKmtFreeMemory(queueBuffer_, SDMA_QUEUE_SIZE); @@ -564,15 +574,8 @@ bool AnvilLib::connect(int srcDeviceId, int dstDeviceId, int numChannels) { LOG_TRACE("SDMA: Connect from %d to %d with %d channels using engine %d", srcDeviceId, dstDeviceId, numChannels, engineId); - // The queue budget is a property of the partition mode rather than of this peer: a CPX partition - // owns one XCD's engines, so the mesh a whole MI300X supports does not fit. KFD also does not - // report how many queues ROCr already holds, so an exhausted budget can only be predicted for - // gross over-subscription and must additionally be recognised when queue creation fails. - // The budget is counted across all engines, not per engine. A rejected engine-pinned create - // retries as a generic queue and reports engine 0, so on a partition every queue charges the - // same key; a per-engine cap would then refuse at numSdmaQueuesPerEngine_ while the partition - // really offers numSdmaEnginesTotal_ times that, which is the figure this message quotes. The - // total is an upper bound, so an uneven distribution is still caught by KFD's own NO_MEMORY. + // KFD does not report how many queues ROCr already holds, so this only predicts gross + // over-subscription; a NO_MEMORY from queue creation below is reported against the same budget. const uint32_t queueBudget = numSdmaEnginesTotal_ * numSdmaQueuesPerEngine_; auto reportBudget = [&](uint32_t used) { LOG_ERROR( @@ -586,36 +589,36 @@ bool AnvilLib::connect(int srcDeviceId, int dstDeviceId, int numChannels) { // The whole request is known up front. Refusing here avoids creating a queue // that rollback would destroy immediately. - const uint32_t used = queuesUsedTotal_; - if (numChannels > 0 && queueBudget > 0 && - used + static_cast(numChannels) > queueBudget) { - reportBudget(used); + if (queueBudgetExceeded(queuesUsedTotal_, numChannels, numSdmaEnginesTotal_, + numSdmaQueuesPerEngine_)) { + reportBudget(queuesUsedTotal_); return false; } auto& vec = sdma_channels_[dstDeviceId]; const size_t already = vec.size(); - auto rollback = [&]() { - while (vec.size() > already) { - SdmaQueue* q = vec.back().get(); - if (q != nullptr && q->valid() && queuesUsedTotal_ > 0) queuesUsedTotal_ -= 1; - vec.pop_back(); - } - }; - for (int c = 0; c < numChannels; ++c) { SdmaQueue* queue = createSdmaQueue(srcDeviceId, dstDeviceId, engineId, selection); if (queue == nullptr) { if (lastQueueStatus_ == HSAKMT_STATUS_NO_MEMORY) { reportBudget(queuesUsedTotal_); } - rollback(); + releaseQueues(vec, already); return false; } } return true; } +void AnvilLib::releaseQueues(std::vector>& vec, size_t keep) { + // createSdmaQueue stores only valid queues and charges each one, so every popped entry is + // credited back. + while (vec.size() > keep) { + if (queuesUsedTotal_ > 0) queuesUsedTotal_ -= 1; + vec.pop_back(); + } +} + void AnvilLib::disconnect() { // Destroy all SDMA queues. SdmaQueue destructor calls hsaKmtDestroyQueue. sdma_channels_.clear(); @@ -626,10 +629,7 @@ void AnvilLib::disconnect() { void AnvilLib::disconnectDevice(int dstDeviceId) { auto it = sdma_channels_.find(dstDeviceId); if (it == sdma_channels_.end()) return; - // Only valid queues were charged to the budget in createSdmaQueue(). - for (const auto& q : it->second) { - if (q != nullptr && q->valid() && queuesUsedTotal_ > 0) queuesUsedTotal_ -= 1; - } + releaseQueues(it->second, 0); sdma_channels_.erase(it); LOG_TRACE("SDMA: Disconnected queues for device %d", dstDeviceId); } @@ -669,17 +669,10 @@ int AnvilLib::getOamId(int deviceId) { if (file.is_open() && (file >> xgmi_physical_id)) return xgmi_physical_id; } - const int preferredEngine = - selection_.preferredMask == 0 ? -1 : __builtin_ctz(selection_.preferredMask); // LOG_ERROR so a default ROCSHMEM_DEBUG_LEVEL=ERROR banff repro still shows the preferred-engine // values that explain why the OAM map was needed. - LOG_ERROR( - "anvil: no xGMI physical id for %s or %s device=%d pair=%d->%d preferredQueried=%d " - "preferredStatus=%#x preferredMask=0x%x preferredEngine=%d hostEng=%u xgmiEng=%u total=%u", - loc.busId.c_str(), loc.physBusId.c_str(), deviceId, selection_.srcDeviceId, - selection_.dstDeviceId, selection_.preferredQueried ? 1 : 0, - static_cast(selection_.preferredStatus), selection_.preferredMask, preferredEngine, - selection_.numSdmaEngines, selection_.numSdmaXgmiEngines, selection_.numSdmaEnginesTotal); + LOG_ERROR("anvil: no xGMI physical id for %s or %s device=%d %s", loc.busId.c_str(), + loc.physBusId.c_str(), deviceId, describeSelection(selection_).c_str()); return -1; } @@ -713,7 +706,8 @@ int AnvilLib::getSdmaEngineIdFromOamMap(int srcDeviceId, int dstDeviceId) { // BDF so that case is not silent. The info line reports the values the fold actually added. const int srcFn = srcPci.function < 0 ? 0 : srcPci.function; const int dstFn = dstPci.function < 0 ? 0 : dstPci.function; - const int folded = foldOamMapEngine(oamEngine, srcFn, dstFn, numSdmaEnginesTotal_); + const int folded = + foldOamMapEngine(oamEngine, srcPci.function, dstPci.function, numSdmaEnginesTotal_); if (srcPci.function < 0 || dstPci.function < 0) { LOG_WARN("anvil: PCI function unreadable src=%s dst=%s, using engine %d (oam=%d total=%u)", srcPci.busId.c_str(), dstPci.busId.c_str(), folded, oamEngine, numSdmaEnginesTotal_); diff --git a/projects/rocshmem/src/sdma/anvil.hpp b/projects/rocshmem/src/sdma/anvil.hpp index 022b8c91390f..0a060908a272 100644 --- a/projects/rocshmem/src/sdma/anvil.hpp +++ b/projects/rocshmem/src/sdma/anvil.hpp @@ -169,6 +169,8 @@ class AnvilLib { void buildGpuAgentMap(); hsa_agent_t getHipGpuAgent(int hipDeviceId) const; void querySdmaEngineCounts(); + // Pop `vec` back to `keep` entries, crediting each popped queue to the budget. + void releaseQueues(std::vector>& vec, size_t keep); int getOamId(int deviceId); int getSdmaEngineIdFromOamMap(int srcDeviceId, int dstDeviceId); // Chooses an engine for the pair and records why, so the caller can hand the same record to diff --git a/projects/rocshmem/src/sdma/anvil_engine_map.hpp b/projects/rocshmem/src/sdma/anvil_engine_map.hpp index a3b39318f73a..d463d316c6cb 100644 --- a/projects/rocshmem/src/sdma/anvil_engine_map.hpp +++ b/projects/rocshmem/src/sdma/anvil_engine_map.hpp @@ -24,8 +24,8 @@ #ifndef LIBRARY_SRC_SDMA_ANVIL_ENGINE_MAP_HPP_ #define LIBRARY_SRC_SDMA_ANVIL_ENGINE_MAP_HPP_ -// Pure engine-selection helpers for AnvilLib. Kept free of HSA/KFD headers so the unit tests can -// build them without USE_SDMA. +// Pure engine-selection and queue-budget helpers for AnvilLib. Kept free of HSA/KFD headers so the +// unit tests can build them without USE_SDMA. #include #include @@ -88,6 +88,15 @@ inline PciFunctionBus pciFunctionBus(const std::string& busId) { return loc; } +// True when taking `requested` more queues would exceed the queue budget, counted across all +// engines as numEnginesTotal * queuesPerEngine. Either count being 0 means the budget is unknown, +// which is never reported as exhausted. +inline bool queueBudgetExceeded(uint32_t used, int requested, uint32_t numEnginesTotal, + uint32_t queuesPerEngine) { + const uint32_t budget = numEnginesTotal * queuesPerEngine; + return requested > 0 && budget > 0 && used + static_cast(requested) > budget; +} + } // namespace sdma_anvil #endif // LIBRARY_SRC_SDMA_ANVIL_ENGINE_MAP_HPP_ diff --git a/projects/rocshmem/src/sdma_policy.cpp b/projects/rocshmem/src/sdma_policy.cpp index bf3aea59ebbd..d8e88656b236 100644 --- a/projects/rocshmem/src/sdma_policy.cpp +++ b/projects/rocshmem/src/sdma_policy.cpp @@ -65,9 +65,11 @@ __host__ void SdmaImpl::sdmaHostInit(int pe, int num_pes, int rank) { CHECK_HIP(hipGetDevice(&deviceId)); // Create SDMA connections to all local PEs including self. A false return leaves a partially - // wired mesh whose null handles would silently skip puts above the SDMA threshold, so tear - // everything down and fall back to IPC memcpy instead. The handle array below is still allocated - // (all null after disconnect) because USE_SDMA testers index it without checking sdmaEnabled. + // wired mesh whose null handles would silently skip puts above the SDMA threshold, so release the + // peers this call wired and fall back to IPC memcpy instead. connect() already rolled back the + // failed peer, and disconnect() is not used because the channel map is process-global. The + // handle array below is still allocated because USE_SDMA testers index it without checking + // sdmaEnabled; it is filled with nulls once sdmaEnabled is false. for (int i = 0; i < shm_size; i++) { if (i != deviceId) { sdma_anvil::EnablePeerAccess(deviceId, i); @@ -75,7 +77,7 @@ __host__ void SdmaImpl::sdmaHostInit(int pe, int num_pes, int rank) { if (!sdma_anvil::anvil.connect(deviceId, i, numChannels)) { LOG_ERROR("SDMA: connect failed from device %d to %d with %d channel(s); disabling SDMA", deviceId, i, numChannels); - sdma_anvil::anvil.disconnect(); + for (int j = 0; j < i; j++) sdma_anvil::anvil.disconnectDevice(j); sdmaEnabled = false; break; } @@ -95,7 +97,8 @@ __host__ void SdmaImpl::sdmaHostInit(int pe, int num_pes, int rank) { for (int i = 0; i < shm_size; i++) { for (int ch = 0; ch < numChannels; ch++) { int idx = i * numChannels + ch; - sdma_anvil::SdmaQueue* queue = sdma_anvil::anvil.getSdmaQueue(deviceId, i, ch); + sdma_anvil::SdmaQueue* queue = + sdmaEnabled ? sdma_anvil::anvil.getSdmaQueue(deviceId, i, ch) : nullptr; handles_h[idx] = queue ? queue->deviceHandle() : nullptr; } } diff --git a/projects/rocshmem/tests/unit_tests/anvil_engine_map_gtest.cpp b/projects/rocshmem/tests/unit_tests/anvil_engine_map_gtest.cpp index b4565beaf903..2d0dfa185a27 100644 --- a/projects/rocshmem/tests/unit_tests/anvil_engine_map_gtest.cpp +++ b/projects/rocshmem/tests/unit_tests/anvil_engine_map_gtest.cpp @@ -85,6 +85,9 @@ TEST(AnvilEngineMap, FoldSeparatesPartitionsOfOneDevice) { } TEST(AnvilEngineMap, FoldTreatsUnreadableFunctionAsZero) { + // One unreadable side: without the clamp, (0 + -1 + 0) % 2 would be -1. + EXPECT_EQ(foldOamMapEngine(0, -1, 0, kCpxTotal), foldOamMapEngine(0, 0, 0, kCpxTotal)); + EXPECT_EQ(foldOamMapEngine(0, 0, -1, kCpxTotal), foldOamMapEngine(0, 0, 0, kCpxTotal)); EXPECT_EQ(foldOamMapEngine(3, -1, -1, kCpxTotal), foldOamMapEngine(3, 0, 0, kCpxTotal)); } @@ -106,6 +109,23 @@ TEST(AnvilEngineMap, PciFunctionBusAcceptsOnlyThreeBitFunctions) { } } +TEST(AnvilEngineMap, QueueBudgetRefusesOverSubscription) { + // CPX: 2 engines x 8 queues per engine. + EXPECT_FALSE(queueBudgetExceeded(9, 7, kCpxTotal, 8)); + EXPECT_TRUE(queueBudgetExceeded(10, 7, kCpxTotal, 8)); + EXPECT_TRUE(queueBudgetExceeded(0, 17, kCpxTotal, 8)); +} + +TEST(AnvilEngineMap, QueueBudgetUnknownNeverRefuses) { + EXPECT_FALSE(queueBudgetExceeded(100, 8, 0, 8)); + EXPECT_FALSE(queueBudgetExceeded(100, 8, kCpxTotal, 0)); +} + +TEST(AnvilEngineMap, QueueBudgetIgnoresEmptyRequest) { + EXPECT_FALSE(queueBudgetExceeded(16, 0, kCpxTotal, 8)); + EXPECT_FALSE(queueBudgetExceeded(16, -1, kCpxTotal, 8)); +} + TEST(AnvilEngineMap, PciFunctionBusHandlesEmptyInput) { const PciFunctionBus loc = pciFunctionBus(""); EXPECT_EQ(loc.function, -1); From 2a40260960c8ae32f153fea47a15b50c15854dd7 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Thu, 8 Oct 2026 16:51:13 -0400 Subject: [PATCH 07/10] Scope sdmaHostStop teardown to the peers it wired Like the failed-init path, normal shutdown now calls disconnectDevice() for each local peer instead of disconnect(), which would destroy queues another user of the process-global channel map still holds. Co-authored-by: Cursor --- projects/rocshmem/src/sdma_policy.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/projects/rocshmem/src/sdma_policy.cpp b/projects/rocshmem/src/sdma_policy.cpp index d8e88656b236..fa44bae58dfa 100644 --- a/projects/rocshmem/src/sdma_policy.cpp +++ b/projects/rocshmem/src/sdma_policy.cpp @@ -115,7 +115,11 @@ __host__ void SdmaImpl::sdmaHostStop() { CHECK_HIP(hipFree(deviceHandles_d)); deviceHandles_d = nullptr; } - sdma_anvil::anvil.disconnect(); + // Release only the peers sdmaHostInit wired: the channel map is process-global, so disconnect() + // would also destroy queues another user still holds. A disabled or failed init left none. + if (sdmaEnabled) { + for (int i = 0; i < shm_size; i++) sdma_anvil::anvil.disconnectDevice(i); + } } #endif // USE_SDMA From f1742deb4f7f52464eded5f2ea56f6c15ce99a5e Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Thu, 8 Oct 2026 17:51:49 -0400 Subject: [PATCH 08/10] Fail fast on an SDMA connect failure in rocSHMEM init sdmaEnabled also selects the alltoall algorithm, so clearing it on one rank while peers keep SDMA leads to incompatible synchronization and a hang. Abort instead; a cross-rank agreed fallback is tracked as AICOMRCCL-2795. Co-authored-by: Cursor --- projects/rocshmem/src/sdma_policy.cpp | 22 ++++++++-------------- 1 file changed, 8 insertions(+), 14 deletions(-) diff --git a/projects/rocshmem/src/sdma_policy.cpp b/projects/rocshmem/src/sdma_policy.cpp index fa44bae58dfa..a02a2f80679e 100644 --- a/projects/rocshmem/src/sdma_policy.cpp +++ b/projects/rocshmem/src/sdma_policy.cpp @@ -64,22 +64,17 @@ __host__ void SdmaImpl::sdmaHostInit(int pe, int num_pes, int rank) { int deviceId; CHECK_HIP(hipGetDevice(&deviceId)); - // Create SDMA connections to all local PEs including self. A false return leaves a partially - // wired mesh whose null handles would silently skip puts above the SDMA threshold, so release the - // peers this call wired and fall back to IPC memcpy instead. connect() already rolled back the - // failed peer, and disconnect() is not used because the channel map is process-global. The - // handle array below is still allocated because USE_SDMA testers index it without checking - // sdmaEnabled; it is filled with nulls once sdmaEnabled is false. + // Create SDMA connections to all local PEs including self. A failed connect aborts rather than + // clearing sdmaEnabled on this rank alone: sdmaEnabled also selects the alltoall algorithm, and + // ranks that disagree on it run incompatible synchronization and hang. for (int i = 0; i < shm_size; i++) { if (i != deviceId) { sdma_anvil::EnablePeerAccess(deviceId, i); } if (!sdma_anvil::anvil.connect(deviceId, i, numChannels)) { - LOG_ERROR("SDMA: connect failed from device %d to %d with %d channel(s); disabling SDMA", - deviceId, i, numChannels); - for (int j = 0; j < i; j++) sdma_anvil::anvil.disconnectDevice(j); - sdmaEnabled = false; - break; + LOG_ERROR_ABORT("SDMA: connect failed from device %d to %d with %d channel(s); set " + "ROCSHMEM_SDMA_ENABLED=0 or lower ROCSHMEM_SDMA_NUM_CHANNELS", + deviceId, i, numChannels); } } @@ -97,8 +92,7 @@ __host__ void SdmaImpl::sdmaHostInit(int pe, int num_pes, int rank) { for (int i = 0; i < shm_size; i++) { for (int ch = 0; ch < numChannels; ch++) { int idx = i * numChannels + ch; - sdma_anvil::SdmaQueue* queue = - sdmaEnabled ? sdma_anvil::anvil.getSdmaQueue(deviceId, i, ch) : nullptr; + sdma_anvil::SdmaQueue* queue = sdma_anvil::anvil.getSdmaQueue(deviceId, i, ch); handles_h[idx] = queue ? queue->deviceHandle() : nullptr; } } @@ -116,7 +110,7 @@ __host__ void SdmaImpl::sdmaHostStop() { deviceHandles_d = nullptr; } // Release only the peers sdmaHostInit wired: the channel map is process-global, so disconnect() - // would also destroy queues another user still holds. A disabled or failed init left none. + // would also destroy queues another user still holds. A disabled init wired none. if (sdmaEnabled) { for (int i = 0; i < shm_size; i++) sdma_anvil::anvil.disconnectDevice(i); } From aad8f6b8b82dbffedc8df2bcc72b3f211f1874d5 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Wed, 14 Oct 2026 20:52:38 -0400 Subject: [PATCH 09/10] Keep Anvil partition logs quiet and copy the engine map into RCCL The partition-engine line is diagnostic detail, and the missing-xGMI comment named a local machine. RCCL's --rocshmem-gin recipe copies a fixed source list, so anvil.hpp's new header has to be on that list or the copied translation unit cannot compile. Co-authored-by: Cursor --- projects/rccl/src/CMakeLists.txt | 2 +- projects/rocshmem/src/sdma/anvil.cpp | 6 ++---- 2 files changed, 3 insertions(+), 5 deletions(-) diff --git a/projects/rccl/src/CMakeLists.txt b/projects/rccl/src/CMakeLists.txt index cb250b9972e2..6e5f72c9c3e0 100644 --- a/projects/rccl/src/CMakeLists.txt +++ b/projects/rccl/src/CMakeLists.txt @@ -1195,7 +1195,7 @@ if (ENABLE_ROCSHMEM_GIN) # SDMA anvil host code: copied into build tree (no sed — host code needs # unqualified log.hpp to find the host shim via include path ordering). set(_GIN_SDMA_DIR "${PROJECT_BINARY_DIR}/gin/sdma") - foreach(_f anvil.cpp anvil.hpp) + foreach(_f anvil.cpp anvil.hpp anvil_engine_map.hpp) add_custom_command( OUTPUT "${_GIN_SDMA_DIR}/${_f}" COMMAND ${CMAKE_COMMAND} -E make_directory "${_GIN_SDMA_DIR}" diff --git a/projects/rocshmem/src/sdma/anvil.cpp b/projects/rocshmem/src/sdma/anvil.cpp index 18b69bfa1fce..e1ca828aacbc 100644 --- a/projects/rocshmem/src/sdma/anvil.cpp +++ b/projects/rocshmem/src/sdma/anvil.cpp @@ -669,8 +669,6 @@ int AnvilLib::getOamId(int deviceId) { if (file.is_open() && (file >> xgmi_physical_id)) return xgmi_physical_id; } - // LOG_ERROR so a default ROCSHMEM_DEBUG_LEVEL=ERROR banff repro still shows the preferred-engine - // values that explain why the OAM map was needed. LOG_ERROR("anvil: no xGMI physical id for %s or %s device=%d %s", loc.busId.c_str(), loc.physBusId.c_str(), deviceId, describeSelection(selection_).c_str()); return -1; @@ -703,7 +701,7 @@ int AnvilLib::getSdmaEngineIdFromOamMap(int srcDeviceId, int dstDeviceId) { const PciFunctionBus srcPci = pciFunctionBus(getBusId(srcDeviceId)); const PciFunctionBus dstPci = pciFunctionBus(getBusId(dstDeviceId)); // An unreadable tail adds 0 and therefore collides with function 0. The warning logs the raw - // BDF so that case is not silent. The info line reports the values the fold actually added. + // BDF so that case is not silent. The trace line reports the values the fold actually added. const int srcFn = srcPci.function < 0 ? 0 : srcPci.function; const int dstFn = dstPci.function < 0 ? 0 : dstPci.function; const int folded = @@ -712,7 +710,7 @@ int AnvilLib::getSdmaEngineIdFromOamMap(int srcDeviceId, int dstDeviceId) { LOG_WARN("anvil: PCI function unreadable src=%s dst=%s, using engine %d (oam=%d total=%u)", srcPci.busId.c_str(), dstPci.busId.c_str(), folded, oamEngine, numSdmaEnginesTotal_); } else if (partition) { - LOG_INFO("anvil: partition engine %d (oam=%d srcFn=%d dstFn=%d total=%u)", folded, oamEngine, + LOG_TRACE("anvil: partition engine %d (oam=%d srcFn=%d dstFn=%d total=%u)", folded, oamEngine, srcFn, dstFn, numSdmaEnginesTotal_); } else { LOG_WARN( From b47b4f7ccfafd2ab5cbc7c096688320d92defe7d Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Thu, 15 Oct 2026 14:41:58 -0400 Subject: [PATCH 10/10] fix(rocshmem): keep Anvil queues on safe SDMA engines Co-authored-by: Cursor --- projects/rocshmem/src/sdma/anvil.cpp | 70 +++++++------------ projects/rocshmem/src/sdma/anvil.hpp | 9 ++- .../rocshmem/src/sdma/anvil_engine_map.hpp | 36 +++++++--- .../unit_tests/anvil_engine_map_gtest.cpp | 66 +++++++++++++---- 4 files changed, 110 insertions(+), 71 deletions(-) diff --git a/projects/rocshmem/src/sdma/anvil.cpp b/projects/rocshmem/src/sdma/anvil.cpp index e1ca828aacbc..4ebb987dac90 100644 --- a/projects/rocshmem/src/sdma/anvil.cpp +++ b/projects/rocshmem/src/sdma/anvil.cpp @@ -212,43 +212,20 @@ SdmaQueue::SdmaQueue(int localDeviceId, int remoteDeviceId, const hsa_agent_t& l // Create SDMA Queue memset(&queue_, 0, sizeof(HsaQueueResource)); - HSAKMT_STATUS pinnedStatus = hsaKmtCreateQueueExt( + // A generic HSA_QUEUE_SDMA retry cannot constrain the engine or doorbell slot. KFD could place + // it on one of the odd-engine slots that a GPU-written doorbell cannot reach (ROCM-32598), so a + // rejected pinned create must fail the connect instead of risking a later hang. + createStatus_ = hsaKmtCreateQueueExt( localNodeId, HSA_QUEUE_SDMA_BY_ENG_ID, DEFAULT_QUEUE_PERCENTAGE, DEFAULT_PRIORITY, engineId, queueBuffer_, SDMA_QUEUE_SIZE, nullptr, &queue_); - HSAKMT_STATUS queueStatus = pinnedStatus; - - if (queueStatus != HSAKMT_STATUS_SUCCESS) { - // An engine-pinned queue needs an engine id the KFD node actually owns. A CPX/DPX partition owns - // one XCD's SDMA engines and no xGMI engines at all, so an id that is valid on the unpartitioned - // device is rejected here. A generic HSA_QUEUE_SDMA lets KFD pick an engine it owns: that gives - // up the per-peer xGMI engine affinity but still reaches the peer over xGMI. - LOG_WARN("anvil: engine-pinned queue rejected (hsakmt=%d %s node=%u engineId=%u), retrying with " - "a generic SDMA queue", - static_cast(queueStatus), hsakmtStatusName(queueStatus), localNodeId, engineId); - memset(&queue_, 0, sizeof(queue_)); - queueStatus = - hsaKmtCreateQueueExt(localNodeId, HSA_QUEUE_SDMA, DEFAULT_QUEUE_PERCENTAGE, DEFAULT_PRIORITY, - 0, queueBuffer_, SDMA_QUEUE_SIZE, nullptr, &queue_); - } - - // A generic retry can return NO_MEMORY after the pinned create returned INVALID_PARAMETER. - // Keeping only the pinned status would hide the budget failure from connect(). - if (queueStatus == HSAKMT_STATUS_SUCCESS) { - createStatus_ = HSAKMT_STATUS_SUCCESS; - } else if (pinnedStatus == HSAKMT_STATUS_NO_MEMORY || queueStatus == HSAKMT_STATUS_NO_MEMORY) { - createStatus_ = HSAKMT_STATUS_NO_MEMORY; - } else { - createStatus_ = pinnedStatus; - } - if (queueStatus != HSAKMT_STATUS_SUCCESS) { + if (createStatus_ != HSAKMT_STATUS_SUCCESS) { const std::string srcBus = getBusId(localDeviceId); const std::string dstBus = getBusId(remoteDeviceId); LOG_ERROR( - "anvil: hsaKmtCreateQueueExt failed hsakmt=%d (%s) genericRetry=%d (%s) node=%u " - "engineId=%u srcDev=%d (%s) dstDev=%d (%s) %s", - static_cast(createStatus_), hsakmtStatusName(createStatus_), - static_cast(queueStatus), hsakmtStatusName(queueStatus), localNodeId, engineId, + "anvil: engine-pinned hsaKmtCreateQueueExt failed hsakmt=%d (%s) node=%u engineId=%u " + "srcDev=%d (%s) dstDev=%d (%s) %s; not retrying as a generic queue", + static_cast(createStatus_), hsakmtStatusName(createStatus_), localNodeId, engineId, localDeviceId, srcBus.c_str(), remoteDeviceId, dstBus.c_str(), describeSelection(selection).c_str()); // Leave valid_ false so the caller can drop the Anvil backend instead of killing the job. @@ -566,8 +543,9 @@ bool AnvilLib::connect(int srcDeviceId, int dstDeviceId, int numChannels) { return false; } const EngineSelection selection = getSdmaEngineId(srcDeviceId, dstDeviceId); - if (selection.engineId < 0) { - LOG_ERROR("anvil: no SDMA engine mapping for %d -> %d", srcDeviceId, dstDeviceId); + if (!isUsableSdmaEngine(selection.engineId)) { + LOG_ERROR("anvil: no usable even SDMA engine for %d -> %d (selected %d)", srcDeviceId, + dstDeviceId, selection.engineId); return false; } const uint32_t engineId = static_cast(selection.engineId); @@ -576,20 +554,21 @@ bool AnvilLib::connect(int srcDeviceId, int dstDeviceId, int numChannels) { // KFD does not report how many queues ROCr already holds, so this only predicts gross // over-subscription; a NO_MEMORY from queue creation below is reported against the same budget. - const uint32_t queueBudget = numSdmaEnginesTotal_ * numSdmaQueuesPerEngine_; + const uint32_t usableEngines = usableSdmaEngines(numSdmaEnginesTotal_); + const uint32_t queueBudget = usableEngines * numSdmaQueuesPerEngine_; auto reportBudget = [&](uint32_t used) { LOG_ERROR( "anvil: SDMA queue budget exhausted: %u queue(s) taken by this process + %d requested, " - "limit %u (%u engines: host=%u xgmi=%u, %u queues per engine). ROCm already holds some of " - "them, so this partition cannot cover this peer count at %d channel(s) per peer. Use fewer " - "ranks per node, fewer channels, or a coarser partition mode (DPX/SPX).", - used, numChannels, queueBudget, numSdmaEnginesTotal_, numSdmaEngines_, numSdmaXgmiEngines_, - numSdmaQueuesPerEngine_, numChannels); + "limit %u (%u usable even engines of %u: host=%u xgmi=%u, %u queues per engine). ROCm " + "already holds some of them, so this GPU cannot cover this peer count at %d channel(s) per " + "peer. Use fewer ranks per node, fewer channels, or a coarser partition mode (DPX/SPX).", + used, numChannels, queueBudget, usableEngines, numSdmaEnginesTotal_, numSdmaEngines_, + numSdmaXgmiEngines_, numSdmaQueuesPerEngine_, numChannels); }; // The whole request is known up front. Refusing here avoids creating a queue // that rollback would destroy immediately. - if (queueBudgetExceeded(queuesUsedTotal_, numChannels, numSdmaEnginesTotal_, + if (queueBudgetExceeded(queuesUsedTotal_, numChannels, usableEngines, numSdmaQueuesPerEngine_)) { reportBudget(queuesUsedTotal_); return false; @@ -690,12 +669,11 @@ int AnvilLib::getSdmaEngineIdFromOamMap(int srcDeviceId, int dstDeviceId) { // The map and the doubling assume the 14 xGMI SDMA engines of an unpartitioned MI300X. A // partition owns fewer engines (CPX: 2 host, 0 xGMI). Same-device peers share an OAM id, so they - // hit the diagonal (oamEngine 0) and engineId 0 is in range. The fold has to run there too, not - // only when the doubled id is past the engine count. + // hit the diagonal (oamEngine 0) and engineId 0 is in range. The fold still runs there, but it + // now targets only the even engines whose every queue slot accepts GPU-written doorbells. // - // (srcFn + dstFn) rather than dstFn alone: with 2 engines the two split a mixed-parity mesh the - // same way, but once the node has more than 2 engines the sum still separates pairs that share a - // destination function. Two engines cannot give 8 partitions distinct ids. + // (srcFn + dstFn) rather than dstFn alone lets modes with multiple usable even engines separate + // pairs that share a destination function. CPX has only one such engine and therefore uses 0. const bool partition = isSdmaPartition(numSdmaXgmiEngines_, numSdmaEnginesTotal_); if (oamMapEngineNeedsFold(numSdmaXgmiEngines_, numSdmaEnginesTotal_, engineId)) { const PciFunctionBus srcPci = pciFunctionBus(getBusId(srcDeviceId)); @@ -742,7 +720,7 @@ EngineSelection AnvilLib::getSdmaEngineId(int srcDeviceId, int dstDeviceId) { hsa_amd_memory_get_preferred_copy_engine(dstAgent, srcAgent, &selection_.preferredMask); selection_.preferredQueried = true; if (selection_.preferredStatus == HSA_STATUS_SUCCESS && selection_.preferredMask != 0) { - const int engineId = __builtin_ctz(selection_.preferredMask); + const int engineId = preferredUsableEngine(selection_.preferredMask); if (engineId >= 0 && (numSdmaEnginesTotal_ == 0 || static_cast(engineId) < numSdmaEnginesTotal_)) { selection_.usedPreferred = true; diff --git a/projects/rocshmem/src/sdma/anvil.hpp b/projects/rocshmem/src/sdma/anvil.hpp index 0a060908a272..7f97d6dadf44 100644 --- a/projects/rocshmem/src/sdma/anvil.hpp +++ b/projects/rocshmem/src/sdma/anvil.hpp @@ -155,11 +155,10 @@ class AnvilLib { // queue budget. Track usage to refuse a mesh that cannot fit before KFD returns NO_MEMORY part // way through building it. uint32_t numSdmaQueuesPerEngine_{0}; - // Queues already taken by this process, counted across all engines rather than per engine: a - // rejected engine-pinned create retries as a generic queue and reports engine 0, so on a - // partition every queue would charge the same key and a per-engine cap would refuse at a - // fraction of the real budget. connect() compares this against - // numSdmaEnginesTotal_ * numSdmaQueuesPerEngine_. + // Queues already taken by this process, counted across the usable even engines rather than per + // engine. connect() compares this against usableSdmaEngines(numSdmaEnginesTotal_) * + // numSdmaQueuesPerEngine_; an engine-pinned rejection fails instead of retrying on an + // unconstrained engine. uint32_t queuesUsedTotal_{0}; HSAKMT_STATUS lastQueueStatus_{HSAKMT_STATUS_SUCCESS}; // Selection in progress, so getOamId can report why the map was consulted without the caller diff --git a/projects/rocshmem/src/sdma/anvil_engine_map.hpp b/projects/rocshmem/src/sdma/anvil_engine_map.hpp index d463d316c6cb..c994da1c3f3e 100644 --- a/projects/rocshmem/src/sdma/anvil_engine_map.hpp +++ b/projects/rocshmem/src/sdma/anvil_engine_map.hpp @@ -32,6 +32,25 @@ namespace sdma_anvil { +// Every queue slot on an even-numbered MI300 SDMA engine is reachable by a GPU-written doorbell. +// Slots 0-3 on odd engines are not (ROCM-32598), and this code cannot choose a slot, so use only +// even engines until the slot-aware workaround can inspect and park unusable queues. +inline bool isUsableSdmaEngine(int engineId) { return engineId >= 0 && engineId % 2 == 0; } + +// Number of even engines in a node-local engine range. A zero total means the count is unknown. +inline uint32_t usableSdmaEngines(uint32_t numSdmaEnginesTotal) { + return (numSdmaEnginesTotal + 1) / 2; +} + +// Select an even engine from an HSA preferred-copy-engine mask. Prefer an even bit from the mask; +// if it contains only odd engines, use the paired even engine immediately below the lowest one. +// Return -1 for an empty mask. +inline int preferredUsableEngine(uint32_t preferredMask) { + if (preferredMask == 0) return -1; + const uint32_t evenMask = preferredMask & 0x55555555u; + return evenMask != 0 ? __builtin_ctz(evenMask) : (__builtin_ctz(preferredMask) & ~1); +} + // A CPX/DPX partition owns host SDMA engines but no xGMI engines. Callers pass a total of 0 when // the engine counts could not be queried, so a failed query is never read as a partition. inline bool isSdmaPartition(uint32_t numSdmaXgmiEngines, uint32_t numSdmaEnginesTotal) { @@ -51,13 +70,14 @@ inline bool oamMapEngineNeedsFold(uint32_t numSdmaXgmiEngines, uint32_t numSdmaE return partition || outOfRange; } -// Fold the undoubled OAM-map value and both PCI functions into the engines this -// node reports. numEngines must be > 0. An unreadable function (-1) counts as 0 -// and therefore collides with function 0. -inline int foldOamMapEngine(int oamEngine, int srcFn, int dstFn, uint32_t numEngines) { +// Fold the undoubled OAM-map value and both PCI functions onto an even engine this node reports. +// numEnginesTotal must be > 0. An unreadable function (-1) counts as 0 and therefore collides with +// function 0. +inline int foldOamMapEngine(int oamEngine, int srcFn, int dstFn, uint32_t numEnginesTotal) { const int src = srcFn < 0 ? 0 : srcFn; const int dst = dstFn < 0 ? 0 : dstFn; - return (oamEngine + src + dst) % static_cast(numEngines); + return 2 * ((oamEngine + src + dst) % + static_cast(usableSdmaEngines(numEnginesTotal))); } // PCI function and the function-0 BDF of the same device. function is -1 when the tail is not a @@ -88,9 +108,9 @@ inline PciFunctionBus pciFunctionBus(const std::string& busId) { return loc; } -// True when taking `requested` more queues would exceed the queue budget, counted across all -// engines as numEnginesTotal * queuesPerEngine. Either count being 0 means the budget is unknown, -// which is never reported as exhausted. +// True when taking `requested` more queues would exceed the queue budget, counted across the +// engines the caller can safely use. Either count being 0 means the budget is unknown, which is +// never reported as exhausted. inline bool queueBudgetExceeded(uint32_t used, int requested, uint32_t numEnginesTotal, uint32_t queuesPerEngine) { const uint32_t budget = numEnginesTotal * queuesPerEngine; diff --git a/projects/rocshmem/tests/unit_tests/anvil_engine_map_gtest.cpp b/projects/rocshmem/tests/unit_tests/anvil_engine_map_gtest.cpp index 2d0dfa185a27..269ea0152295 100644 --- a/projects/rocshmem/tests/unit_tests/anvil_engine_map_gtest.cpp +++ b/projects/rocshmem/tests/unit_tests/anvil_engine_map_gtest.cpp @@ -32,6 +32,8 @@ namespace { // MI300X engine counts as ROCr reports them. constexpr uint32_t kSpxXgmi = 14; constexpr uint32_t kSpxTotal = 16; +constexpr uint32_t kDpxTotal = 8; +constexpr uint32_t kQpxTotal = 4; constexpr uint32_t kCpxXgmi = 0; constexpr uint32_t kCpxTotal = 2; @@ -44,9 +46,31 @@ TEST(AnvilEngineMap, PartitionIsNoXgmiEnginesWithKnownTotal) { EXPECT_FALSE(isSdmaPartition(0, 0)); } +TEST(AnvilEngineMap, OnlyEvenEnginesAreUsable) { + for (int engine = 0; engine < static_cast(kSpxTotal); ++engine) { + EXPECT_EQ(isUsableSdmaEngine(engine), engine % 2 == 0) << "engine=" << engine; + } + EXPECT_FALSE(isUsableSdmaEngine(-1)); +} + +TEST(AnvilEngineMap, UsableEngineCountIncludesOnlyEvenEngines) { + EXPECT_EQ(usableSdmaEngines(kSpxTotal), 8u); + EXPECT_EQ(usableSdmaEngines(kDpxTotal), 4u); + EXPECT_EQ(usableSdmaEngines(kQpxTotal), 2u); + EXPECT_EQ(usableSdmaEngines(kCpxTotal), 1u); + EXPECT_EQ(usableSdmaEngines(0), 0u); +} + +TEST(AnvilEngineMap, PreferredMaskSelectsAnEvenEngine) { + EXPECT_EQ(preferredUsableEngine(0x0004), 2); + EXPECT_EQ(preferredUsableEngine(0x0006), 2); + EXPECT_EQ(preferredUsableEngine(0x0002), 0); + EXPECT_EQ(preferredUsableEngine(0x8000), 14); + EXPECT_EQ(preferredUsableEngine(0), -1); +} + TEST(AnvilEngineMap, PartitionFoldsEvenWhenDoubledIdIsInRange) { - // Same-device CPX peers share an OAM id and hit the map diagonal, doubled id 0. Without the - // partition fold they would all be pinned to engine 0. + // The partition path must apply the safe-engine fold even when the legacy map id is in range. EXPECT_TRUE(oamMapEngineNeedsFold(kCpxXgmi, kCpxTotal, 0)); EXPECT_TRUE(oamMapEngineNeedsFold(kCpxXgmi, kCpxTotal, 1)); } @@ -67,21 +91,29 @@ TEST(AnvilEngineMap, UnknownTotalNeverFolds) { EXPECT_FALSE(oamMapEngineNeedsFold(0, 0, 12)); } -TEST(AnvilEngineMap, FoldStaysInRangeForEveryInput) { - for (int oam = 0; oam < 8; ++oam) { - for (int src = -1; src < 8; ++src) { - for (int dst = -1; dst < 8; ++dst) { - const int e = foldOamMapEngine(oam, src, dst, kCpxTotal); - EXPECT_GE(e, 0); - EXPECT_LT(e, static_cast(kCpxTotal)); +TEST(AnvilEngineMap, FoldStaysOnAnInRangeEvenEngine) { + for (uint32_t total : {kDpxTotal, kQpxTotal, kCpxTotal}) { + for (int oam = 0; oam < 8; ++oam) { + for (int src = -1; src < 8; ++src) { + for (int dst = -1; dst < 8; ++dst) { + const int engine = foldOamMapEngine(oam, src, dst, total); + EXPECT_TRUE(isUsableSdmaEngine(engine)) << "total=" << total << " engine=" << engine; + EXPECT_LT(engine, static_cast(total)) << "total=" << total; + } } } } } -TEST(AnvilEngineMap, FoldSeparatesPartitionsOfOneDevice) { - // Diagonal (oam 0) pairs from function 0 to functions 0 and 1 land on different engines. - EXPECT_NE(foldOamMapEngine(0, 0, 0, kCpxTotal), foldOamMapEngine(0, 0, 1, kCpxTotal)); +TEST(AnvilEngineMap, CpxFoldUsesOnlyEngineZero) { + for (int oam = 0; oam < 8; ++oam) { + EXPECT_EQ(foldOamMapEngine(oam, 3, 5, kCpxTotal), 0); + } +} + +TEST(AnvilEngineMap, FoldSeparatesPairsWhenMultipleEvenEnginesExist) { + EXPECT_NE(foldOamMapEngine(0, 0, 1, kDpxTotal), + foldOamMapEngine(0, 1, 1, kDpxTotal)); } TEST(AnvilEngineMap, FoldTreatsUnreadableFunctionAsZero) { @@ -116,6 +148,16 @@ TEST(AnvilEngineMap, QueueBudgetRefusesOverSubscription) { EXPECT_TRUE(queueBudgetExceeded(0, 17, kCpxTotal, 8)); } +TEST(AnvilEngineMap, QueueBudgetCountsEvenEnginesOnly) { + const uint32_t cpxUsable = usableSdmaEngines(kCpxTotal); + EXPECT_FALSE(queueBudgetExceeded(7, 1, cpxUsable, 8)); + EXPECT_TRUE(queueBudgetExceeded(8, 1, cpxUsable, 8)); + + const uint32_t qpxUsable = usableSdmaEngines(kQpxTotal); + EXPECT_FALSE(queueBudgetExceeded(15, 1, qpxUsable, 8)); + EXPECT_TRUE(queueBudgetExceeded(16, 1, qpxUsable, 8)); +} + TEST(AnvilEngineMap, QueueBudgetUnknownNeverRefuses) { EXPECT_FALSE(queueBudgetExceeded(100, 8, 0, 8)); EXPECT_FALSE(queueBudgetExceeded(100, 8, kCpxTotal, 0));