From 3d321ae915fd448af27a3224dfe3ca6f88ed89cb Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Tue, 22 Sep 2026 09:39:14 -0500 Subject: [PATCH 01/13] ibv port maxMsgSize check --- src/client/Topology.hpp | 19 ++++-- src/client/Utilities.hpp | 7 +- src/header/TransferBench.hpp | 123 ++++++++++++++++++++++++++++++----- 3 files changed, 128 insertions(+), 21 deletions(-) diff --git a/src/client/Topology.hpp b/src/client/Topology.hpp index 98cb717f..e9eca347 100644 --- a/src/client/Topology.hpp +++ b/src/client/Topology.hpp @@ -35,9 +35,9 @@ static int RemappedCpuIndex(int origIdx) static void PrintNicToGPUTopo(bool outputToCsv) { if (!IsIbvSymbolsReady()) return; - printf(" NIC | Device Name | Active | PCIe Bus ID | NUMA | Closest GPU(s) | GID Index | GID Descriptor\n"); + printf(" NIC | Device Name | Active | PCIe Bus ID | NUMA | Closest GPU(s) | MaxMsgSz | GID Index | GID Descriptor\n"); if(!outputToCsv) - printf("-----+-------------+--------+--------------+------+----------------+-----------+-------------------\n"); + printf("-----+-------------+--------+--------------+------+----------------+------------+-----------+-------------------\n"); int numGpus = TransferBench::GetNumExecutors(EXE_GPU_GFX); auto const& ibvDeviceList = GetIbvDeviceList(); @@ -56,12 +56,13 @@ static void PrintNicToGPUTopo(bool outputToCsv) for (int i = 0; i < ibvDeviceList.size(); i++) { std::string closestGpusStr = closestGpusForNic[i]; - printf(" %-3d | %-11s | %-6s | %-12s | %-4d | %-14s | %-9s | %-20s\n", + printf(" %-3d | %-11s | %-6s | %-12s | %-4d | %-14s | %-10s | %-9s | %-20s\n", i, ibvDeviceList[i].name.c_str(), ibvDeviceList[i].hasActivePort ? "Yes" : "No", ibvDeviceList[i].busId.c_str(), ibvDeviceList[i].numaNode, closestGpusStr.c_str(), + ibvDeviceList[i].maxMsgSize ? std::to_string(ibvDeviceList[i].maxMsgSize).c_str() : "-", ibvDeviceList[i].isRoce && ibvDeviceList[i].hasActivePort ? std::to_string(ibvDeviceList[i].gidIndex).c_str() : "N/A", ibvDeviceList[i].isRoce && ibvDeviceList[i].hasActivePort ? ibvDeviceList[i].gidDescriptor.c_str() : "N/A" ); @@ -239,6 +240,7 @@ void DisplayMultiRankTopology(bool outputToCsv, bool showBorders) std::vector nicClosestCpu = std::get<7>(key); std::vector nicClosestGpu = std::get<8>(key); std::vector nicIsActive = std::get<9>(key); + std::vector nicMaxMsgSize = std::get<10>(key); int numRanks = hosts.size(); int numCpus = cpuNames.size(); @@ -253,7 +255,7 @@ void DisplayMultiRankTopology(bool outputToCsv, bool showBorders) groupNum++, numRanks, numCpus, numGpus, numNics, numActiveNics); // Determine size of table - int numCols = 6; + int numCols = 7; int numRows = 1 + std::max(numRanks, numExecutors); TransferBench::Utils::TableHelper table(numRows, numCols); @@ -273,6 +275,7 @@ void DisplayMultiRankTopology(bool outputToCsv, bool showBorders) table.Set(0, 3, " Executor "); table.Set(0, 4, " Executor Name "); table.Set(0, 5, " #SE "); + table.Set(0, 6, " MaxMsgSz "); // Fill in ranks / hosts for (int i = 0; i < numRanks; i++) { @@ -306,6 +309,10 @@ void DisplayMultiRankTopology(bool outputToCsv, bool showBorders) table.Set(rowIdx, 3, " - NIC %02d ", nicIndex); table.Set(rowIdx, 4, " - %s", nicNames[nicIndex].c_str()); table.Set(rowIdx, 5, " %s ", nicIsActive[nicIndex] ? "ON" : "OFF"); + if (nicMaxMsgSize[nicIndex]) + table.Set(rowIdx, 6, " %u ", nicMaxMsgSize[nicIndex]); + else + table.Set(rowIdx, 6, " - "); rowIdx++; } } @@ -316,6 +323,10 @@ void DisplayMultiRankTopology(bool outputToCsv, bool showBorders) table.Set(rowIdx, 3, " - NIC %02d ", nicIndex); table.Set(rowIdx, 4, " - %s ", nicNames[nicIndex].c_str()); table.Set(rowIdx, 5, " %s ", nicIsActive[nicIndex] ? "ON" : "OFF"); + if (nicMaxMsgSize[nicIndex]) + table.Set(rowIdx, 6, " %u ", nicMaxMsgSize[nicIndex]); + else + table.Set(rowIdx, 6, " - "); rowIdx++; } } diff --git a/src/client/Utilities.hpp b/src/client/Utilities.hpp index ad3a5641..6d69f29a 100644 --- a/src/client/Utilities.hpp +++ b/src/client/Utilities.hpp @@ -116,7 +116,8 @@ namespace TransferBench::Utils std::vector, // NIC Names std::vector, // NIC Closest NUMA std::vector, // NIC Closest GPU - std::vector // NIC is active + std::vector, // NIC is active + std::vector // NIC max_msg_sz > GroupKey; typedef std::map> RankGroupMap; @@ -393,17 +394,19 @@ namespace TransferBench::Utils std::vector nicNames; std::vector nicClosestCpu; std::vector nicIsActive; + std::vector nicMaxMsgSize; for (int exeIndex = 0; exeIndex < numNics; exeIndex++) { ExeDevice exeDevice = {EXE_NIC, exeIndex, rank}; nicNames.push_back(TransferBench::GetExecutorName(exeDevice)); nicClosestCpu.push_back(TransferBench::GetClosestCpuNumaToNic(exeIndex, rank)); nicIsActive.push_back(TransferBench::NicIsActive(exeIndex, rank)); + nicMaxMsgSize.push_back(TransferBench::GetNicMaxMsgSize(exeIndex, rank)); } GroupKey key(podId, cpuNames, cpuNumSubExecs, gpuNames, gpuNumSubExecs, gpuClosestCpu, - nicNames, nicClosestCpu, nicClosestGpu, nicIsActive); + nicNames, nicClosestCpu, nicClosestGpu, nicIsActive, nicMaxMsgSize); groups[key].push_back(rank); } diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index d2a59f4b..2b0b47d8 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -610,6 +610,13 @@ namespace TransferBench */ int NicIsActive(int nicIndex, int targetRank = -1); + /** + * @param[in] nicIndex The NIC index to query + * @param[in] targetRank Rank to query (-1 for local rank) + * @returns Smallest max_msg_sz across the NIC's active ports in bytes, or 0 if unknown + */ + uint32_t GetNicMaxMsgSize(int nicIndex, int targetRank = -1); + /** * Helper function to parse a line containing Transfers into a vector of Transfers * @@ -1110,6 +1117,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) bool IsSamePod(int targetRank, int sourceRank) const; std::string GetExecutorName(ExeDevice exeDevice) const; int NicIsActive(int nicIndex, int targetRank) const; + uint32_t GetNicMaxMsgSize(int nicIndex, int targetRank) const; // Translate a logical CPU NUMA index (as exposed to users) into the physical // NUMA node id used by libnuma / HSA. Returns the index unchanged if out of range. @@ -1183,6 +1191,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) std::map closestCpuNumaToGpu; std::map closestCpuNumaToNic; std::map nicIsActive; + std::map nicMaxMsgSize; std::map> closestNicsToGpu; std::map> closestGpusToNic; std::map, std::string> executorName; @@ -2757,6 +2766,41 @@ const auto& AmdSmiFabricInfoV1(const T& info) hasFatalError = true; break; } + + // Each posted RDMA WR is min(chunkBytes, bytes assigned to that QP). + // That size must not exceed the port's max_msg_sz on either NIC. + // Split matches PrepareSubExecParams: leftover blocks go to later QPs, so the + // last active QP can be the largest. + // This mimics PrepareSubExecParams. + size_t const N = t.numBytes / sizeof(float); + int const targetMultiple = cfg.data.blockBytes / sizeof(float); + int const qpCount = t.numSubExecs; + int const maxSubExecToUse = std::min((size_t)(N + targetMultiple - 1) / targetMultiple, + (size_t)qpCount); + size_t assigned = 0; + size_t maxQpBytes = 0; + for (int qp = 0; qp < qpCount; ++qp) { + int const subExecLeft = std::max(0, maxSubExecToUse - qp); + size_t const leftover = N - assigned; + size_t const roundedN = (leftover + targetMultiple - 1) / targetMultiple; + size_t const qpN = subExecLeft + ? std::min(leftover, (roundedN / subExecLeft) * (size_t)targetMultiple) + : 0; + maxQpBytes = std::max(maxQpBytes, qpN * sizeof(float)); + assigned += qpN; + } + size_t const maxWrBytes = maxQpBytes ? std::min(cfg.nic.chunkBytes, maxQpBytes) : 0; + uint32_t const maxMsgSz = std::min(GetNicMaxMsgSize(srcExeDevice.exeIndex, srcExeDevice.exeRank), + GetNicMaxMsgSize(dstExeDevice.exeIndex, dstExeDevice.exeRank)); + if (maxMsgSz && maxWrBytes > maxMsgSz) { + errors.push_back({ERR_FATAL, + "Transfer %d: NIC work request size (%lu bytes) exceeds IBV max_msg_sz (%u bytes) " + "with %d queue pair(s) and chunkBytes %lu. Reduce NIC_CHUNK_BYTES / transfer size, " + "or increase the queue pair count", + i, maxWrBytes, maxMsgSz, qpCount, cfg.nic.chunkBytes}); + hasFatalError = true; + break; + } } else { errors.push_back({ERR_FATAL, "Transfer %d: NIC executor is requested but is not available.", i}); hasFatalError = true; @@ -3204,6 +3248,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) int gidIndex; std::string gidDescriptor; bool isRoce; + uint32_t maxMsgSize; ///< Smallest max_msg_sz across active ports (largest RDMA work request) }; // Function to collect information about IBV devices @@ -3343,29 +3388,38 @@ const auto& AmdSmiFabricInfoV1(const T& info) ibvDevice.devicePtr = deviceList[i]; ibvDevice.name = deviceList[i]->name; ibvDevice.hasActivePort = false; + ibvDevice.maxMsgSize = 0; + ibvDevice.isRoce = false; + ibvDevice.gidIndex = -1; { struct ibv_context *context = ibv_open_device(ibvDevice.devicePtr); if (context) { struct ibv_device_attr deviceAttr; if (!ibv_query_device(context, &deviceAttr)) { int activePort; - ibvDevice.gidIndex = -1; for (int port = 1; port <= deviceAttr.phys_port_cnt; ++port) { struct ibv_port_attr portAttr; if (ibv_query_port(context, port, &portAttr)) continue; - if (portAttr.state == IBV_PORT_ACTIVE) { - activePort = port; - ibvDevice.hasActivePort = true; - if(portAttr.link_layer == IBV_LINK_LAYER_ETHERNET) { - ibvDevice.isRoce = true; - std::pair gidInfo (-1, ""); - auto res = GetGidIndex(context, portAttr.gid_tbl_len, activePort, gidInfo); - if (res.errType == ERR_NONE) { - ibvDevice.gidIndex = gidInfo.first; - ibvDevice.gidDescriptor = gidInfo.second; - } + if (portAttr.state != IBV_PORT_ACTIVE) continue; + + // Any active port may be selected for QP setup via cfg.nic.ibPort, which is + // not visible here, so keep the most restrictive limit across all of them. + if (portAttr.max_msg_sz && + (!ibvDevice.maxMsgSize || portAttr.max_msg_sz < ibvDevice.maxMsgSize)) + ibvDevice.maxMsgSize = portAttr.max_msg_sz; + + // The first active port supplies the RoCE / GID information + if (ibvDevice.hasActivePort) continue; + activePort = port; + ibvDevice.hasActivePort = true; + if (portAttr.link_layer == IBV_LINK_LAYER_ETHERNET) { + ibvDevice.isRoce = true; + std::pair gidInfo (-1, ""); + auto res = GetGidIndex(context, portAttr.gid_tbl_len, activePort, gidInfo); + if (res.errType == ERR_NONE) { + ibvDevice.gidIndex = gidInfo.first; + ibvDevice.gidDescriptor = gidInfo.second; } - break; } } } @@ -3805,6 +3859,13 @@ const auto& AmdSmiFabricInfoV1(const T& info) int const port = cfg.nic.ibPort; + size_t maxWrBytes = 0; + for (int i = 0; i < rss.qpCount; i++) { + size_t qpBytes = rss.subExecParamCpu[i].N * sizeof(float); + if (qpBytes) + maxWrBytes = std::max(maxWrBytes, std::min(cfg.nic.chunkBytes, qpBytes)); + } + // Prepare NIC on SRC mem rank int srcGidIndex = cfg.nic.ibGidIndex; bool srcIsRoCE = false; @@ -3846,6 +3907,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) IBV_PTR_CALL(rss.srcCompQueue, ibv_create_cq, rss.srcContext, srcCQSize, NULL, NULL, 0); // Get SRC port attributes IBV_CALL(ibv_query_port, rss.srcContext, port, &rss.srcPortAttr); + if (rss.srcPortAttr.max_msg_sz && maxWrBytes > rss.srcPortAttr.max_msg_sz) { + return {ERR_FATAL, + "Transfer %d: NIC work request size (%lu bytes) exceeds SRC NIC %d IBV max_msg_sz (%u bytes) " + "with %u queue pair(s) and chunkBytes %lu. Reduce NIC_CHUNK_BYTES / transfer size, " + "or increase the queue pair count", + rss.transferIdx, maxWrBytes, rss.srcNicIndex, rss.srcPortAttr.max_msg_sz, + rss.qpCount, cfg.nic.chunkBytes}; + } // Check for RDMA over Converged Ethernet (RoCE) and update GID index appropriately srcIsRoCE = (rss.srcPortAttr.link_layer == IBV_LINK_LAYER_ETHERNET); if (srcIsRoCE) { @@ -3907,6 +3976,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) IBV_PTR_CALL(rss.dstCompQueue, ibv_create_cq, rss.dstContext, dstCQSize, NULL, NULL, 0); // Get DST port attributes IBV_CALL(ibv_query_port, rss.dstContext, port, &rss.dstPortAttr); + if (rss.dstPortAttr.max_msg_sz && maxWrBytes > rss.dstPortAttr.max_msg_sz) { + return {ERR_FATAL, + "Transfer %d: NIC work request size (%lu bytes) exceeds DST NIC %d IBV max_msg_sz (%u bytes) " + "with %u queue pair(s) and chunkBytes %lu. Reduce NIC_CHUNK_BYTES / transfer size, " + "or increase the queue pair count", + rss.transferIdx, maxWrBytes, rss.dstNicIndex, rss.dstPortAttr.max_msg_sz, + rss.qpCount, cfg.nic.chunkBytes}; + } // Check for RDMA over Converged Ethernet (RoCE) and update GID index appropriately dstIsRoCE = (rss.dstPortAttr.link_layer == IBV_LINK_LAYER_ETHERNET); if (dstIsRoCE) { @@ -8325,13 +8402,15 @@ const auto& AmdSmiFabricInfoV1(const T& info) topo.closestCpuNumaToNic[exeIndex] = GetClosestLogicalCpu(nicPhysNode); topo.executorName[{EXE_NIC, exeIndex}] = GetIbvDeviceList()[exeIndex].name; topo.nicIsActive[exeIndex] = GetIbvDeviceList()[exeIndex].hasActivePort; + topo.nicMaxMsgSize[exeIndex] = GetIbvDeviceList()[exeIndex].maxMsgSize; if (verbose) { auto const& nic = GetIbvDeviceList()[exeIndex]; - Log("[INFO] Rank %03d: NIC [%02d/%02d] %s BDF %s NUMA %d active=%s\n", + Log("[INFO] Rank %03d: NIC [%02d/%02d] %s BDF %s NUMA %d active=%s max_msg_sz=%u\n", rank, exeIndex, numNics, nic.name.c_str(), nic.busId.empty() ? "?" : nic.busId.c_str(), topo.closestCpuNumaToNic[exeIndex], - nic.hasActivePort ? "yes" : "no"); + nic.hasActivePort ? "yes" : "no", + nic.maxMsgSize); } } } @@ -8572,6 +8651,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) SendMap(peerRank, topo.closestCpuNumaToGpu); SendMap(peerRank, topo.closestCpuNumaToNic); SendMap(peerRank, topo.nicIsActive); + SendMap(peerRank, topo.nicMaxMsgSize); SendMap(peerRank, topo.closestNicsToGpu); SendMap(peerRank, topo.closestGpusToNic); SendMap(peerRank, topo.executorName); @@ -8588,6 +8668,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) RecvMap(peerRank, topo.closestCpuNumaToGpu); RecvMap(peerRank, topo.closestCpuNumaToNic); RecvMap(peerRank, topo.nicIsActive); + RecvMap(peerRank, topo.nicMaxMsgSize); RecvMap(peerRank, topo.closestNicsToGpu); RecvMap(peerRank, topo.closestGpusToNic); RecvMap(peerRank, topo.executorName); @@ -8957,6 +9038,13 @@ const auto& AmdSmiFabricInfoV1(const T& info) return rankInfo[targetRank].nicIsActive.at(nicIndex); } + uint32_t System::GetNicMaxMsgSize(int nicIndex, int targetRank) const + { + if (targetRank < 0 || targetRank >= numRanks) targetRank = rank; + if (rankInfo[targetRank].nicMaxMsgSize.count(nicIndex) == 0) return 0; + return rankInfo[targetRank].nicMaxMsgSize.at(nicIndex); + } + int GetNumExecutors(ExeType exeType, int targetRank) { return System::Get().GetNumExecutors(exeType, targetRank); @@ -9071,6 +9159,11 @@ const auto& AmdSmiFabricInfoV1(const T& info) return System::Get().NicIsActive(nicIndex, targetRank); } + uint32_t GetNicMaxMsgSize(int nicIndex, int targetRank) + { + return System::Get().GetNicMaxMsgSize(nicIndex, targetRank); + } + // Undefine CUDA compatibility macros #if defined(__NVCC__) From 98a43f5dd0dcb182244e1c6ff96bc91af3aafdf4 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Tue, 22 Sep 2026 13:29:14 -0500 Subject: [PATCH 02/13] fixing logical CU ID report --- src/header/TransferBench.hpp | 64 +++++++++++++++++++++++++++--------- 1 file changed, 49 insertions(+), 15 deletions(-) diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 2b0b47d8..533ad736 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -734,9 +734,17 @@ namespace TransferBench // Helper macro functions //========================================================================================== -// Returns xccId and a unified cuId for the current wavefront. -// cuId encoding is arch-specific (see comments) but is always a dense index -// suitable for CU-set tracking. +// Returns xccId and the logical CU_MASK slot for the current wavefront. +// +// Logical CU ID: KFD's mqd_symmetrically_map_cu_mask() maps mask bit to (CU, SH, SE) +// by walking CU outermost and SE innermost, so the slot is the same mixed-radix +// expression on every architecture, only the radices differ per arch: +// slot = (cuIdx * numSHs + sh) * numSEs + se +// A wave launched under CU_MASK=N therefore reports cuId == N, on every XCC. +// +// Exact only while every (SE,SH) has the same number of active CUs. +// KFD skips visits in the final, ragged CU row (e.g. 38 CUs on MI300X), +// which shifts the highest few slots per XCC; slots below min(cu_per_sh) * numSEs are unaffected. __device__ __forceinline__ void GetXccHwId(uint32_t& xccId, uint32_t& cuId) { #if defined(__gfx942__) || defined(__gfx950__) @@ -745,10 +753,15 @@ __device__ __forceinline__ void GetXccHwId(uint32_t& xccId, uint32_t& cuId) uint32_t hwId = 0, xccReg = 0; asm volatile("s_getreg_b32 %0, hwreg(HW_REG_HW_ID)" : "=s"(hwId)); asm volatile("s_getreg_b32 %0, hwreg(HW_REG_XCC_ID)" : "=s"(xccReg)); + // constexpr uint32_t numSEs = 4; // SEs per XCC + // constexpr uint32_t numSHs = 1; // SHs per SE (SH_ID is therefore always 0) + // uint32_t const se = (hwId >> 13) & 0x3; // SE_ID [15:13] + // uint32_t const sh = (hwId >> 12) & 0x1; // SH_ID [12] + // uint32_t const cu = (hwId >> 8) & 0xF; // CU_ID [11:8] + // cuId = (cu * numSHs + sh) * numSEs + se; xccId = xccReg & 0xF; - cuId = (((hwId >> 12) & 1) << 5) // SH_ID - | (((hwId >> 8) & 15) << 2) // CU_ID - | ((hwId >> 13) & 3); // SE_ID + cuId = (((hwId >> 8) & 15) << 2) // CU_ID [11:8] → bits [5:2] + | ((hwId >> 13) & 3); // SE_ID [15:13] → bits [1:0] #elif defined(__gfx1250__) // CDNA5: HW_ID1 (code 23) + RTN_GET_SE_HW_ID (0x87) @@ -756,20 +769,41 @@ __device__ __forceinline__ void GetXccHwId(uint32_t& xccId, uint32_t& cuId) uint32_t hwId = 0, seHwId = 0; asm volatile("s_getreg_b32 %0, hwreg(HW_REG_HW_ID1)" : "=s"(hwId)); asm volatile("s_sendmsg_rtn_b32 %0, 0x87\ns_wait_kmcnt 0" : "=s"(seHwId)); - xccId = (seHwId >> 16) & 0xF; // Virtual_XCC_ID [19:16] - cuId = ((seHwId & 0xF) << 4) // SE_ID [3:0] (gfx1250: 2 SEs → 1 bit) - | (((hwId >> 16) & 0x1) << 3) // SA_ID [16] (gfx1250: 2 SAs → 1 bit) - | ((hwId >> 10) & 0x7); // WGP_ID [12:10] (gfx1250: 8 WGPs per SA → 3 bits) + // Masking is at WGP granularity here, so cuIdx is the WGP index and the SA + // takes the role of KFD's SH. 32 WGPs per XCC gives slots 0-31. + // constexpr uint32_t numSEs = 2; // SEs per XCC + // constexpr uint32_t numSAs = 2; // SAs per SE + // uint32_t const se = seHwId & 0x1; // SE_ID [3:0] + // uint32_t const sa = (hwId >> 16) & 0x1; // SA_ID [16] + // uint32_t const wgp = (hwId >> 10) & 0x7; // WGP_ID [12:10] (8 WGPs per SA) + // cuId = (wgp * numSAs + sa) * numSEs + se; + xccId = (seHwId >> 16) & 0xF; // Virtual_XCC_ID [19:16] + cuId = (((hwId >> 10) & 0x7) << 2) // WGP_ID [12:10] → bits [4:2] + | (((hwId >> 16) & 0x1) << 1) // SA_ID [16] → bit 1 + | (seHwId & 0x1); // SE_ID [0] → bit 0 + +#elif defined(__gfx90a__) + // CDNA2 / MI200 (MI210, MI250, MI250X): one XCC per HIP device (each GCD is + // its own GPU). slot = CU * 8 + SE (SH is always 0). + // 8 SEs, 1 SH per SE. KFD xcc_mask = 1, so CU_MASK=N is HIP bit N. + // MI250X with 110 CUs over 8 SEs are susceptible to the ragged row issue as noted above. + // CDNA2 ISA §3.12 Table 6. + uint32_t hwId = 0; + asm volatile("s_getreg_b32 %0, hwreg(HW_REG_HW_ID)" : "=s"(hwId)); + xccId = 0; + cuId = (((hwId >> 8) & 15) << 3) // CU_ID [11:8] → bits [6:3] + | ((hwId >> 13) & 7); // SE_ID [15:13] → bits [2:0] (3 bits) #elif defined(__GFX9__) - // Other GFX9 (gfx90a/MI200, gfx908, gfx906) — single die, no XCC - // CDNA2 ISA §3.12 Table 6 + // Remaining GFX9 (gfx908/MI100, gfx906/MI50): SE counts differ, so this is a + // packed (SH,CU,SE) coordinate, not a CU_MASK slot. gfx908 also has 8 SEs; + // its slot formula would match gfx90a but is left as a coordinate for now. uint32_t hwId = 0; asm volatile("s_getreg_b32 %0, hwreg(HW_REG_HW_ID)" : "=s"(hwId)); xccId = 0; - cuId = (((hwId >> 12) & 1) << 5) - | (((hwId >> 8) & 15) << 2) - | ((hwId >> 13) & 3); + cuId = (((hwId >> 12) & 1) << 6) // SH_ID + | (((hwId >> 8) & 15) << 2) // CU_ID + | ((hwId >> 13) & 3); // SE_ID #elif defined(__GFX10__) || defined(__GFX11__) || defined(__GFX12__) // RDNA2/3/4 (non-CDNA5) — HW_ID1 present, no XCC From bed30a5d47c4661e5e4eb03a9da62b7feb20df45 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Tue, 22 Sep 2026 16:49:06 -0500 Subject: [PATCH 03/13] gfx1250strict --- CMakeLists.txt | 3 ++- docs/install/build_from_source.rst | 2 +- src/header/TransferBench.hpp | 4 ++-- src/header/tdmCopy.h | 2 +- 4 files changed, 6 insertions(+), 5 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 8252b5f9..cc11a742 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -220,7 +220,8 @@ set(DEFAULT_GPUS gfx1151 gfx1200 gfx1201 - gfx1250) + gfx1250 + gfx1250-strict) if(BUILD_LOCAL_GPU_TARGET_ONLY) message(STATUS "Building only for local GPU target") diff --git a/docs/install/build_from_source.rst b/docs/install/build_from_source.rst index f80d6570..d5750849 100644 --- a/docs/install/build_from_source.rst +++ b/docs/install/build_from_source.rst @@ -317,7 +317,7 @@ To modify the CMake behavior, use the following environment variables: CMake cache variables GPU_TARGETS Semicolon-separated GPU architectures. Overridden if BUILD_LOCAL_GPU_TARGET_ONLY is ON - gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1200;gfx1201;gfx1250 + gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1200;gfx1201;gfx1250;gfx1250-strict AMD_SMI_EXECUTABLE diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 533ad736..a2e207ac 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -763,7 +763,7 @@ __device__ __forceinline__ void GetXccHwId(uint32_t& xccId, uint32_t& cuId) cuId = (((hwId >> 8) & 15) << 2) // CU_ID [11:8] → bits [5:2] | ((hwId >> 13) & 3); // SE_ID [15:13] → bits [1:0] -#elif defined(__gfx1250__) +#elif defined(__gfx1250__) || defined(__gfx1250_strict__) // CDNA5: HW_ID1 (code 23) + RTN_GET_SE_HW_ID (0x87) // CDNA5 ISA §3.4.9, §5.4 Table 19 uint32_t hwId = 0, seHwId = 0; @@ -829,7 +829,7 @@ __device__ __forceinline__ uint32_t GetXccId() uint32_t xccReg = 0; asm volatile("s_getreg_b32 %0, hwreg(HW_REG_XCC_ID)" : "=s"(xccReg)); return xccReg & 0xF; -#elif defined(__gfx1250__) +#elif defined(__GFX12__) uint32_t seHwId = 0; asm volatile("s_sendmsg_rtn_b32 %0, 0x87\ns_wait_kmcnt 0" : "=s"(seHwId)); return (seHwId >> 16) & 0xF; // Virtual_XCC_ID [19:16] diff --git a/src/header/tdmCopy.h b/src/header/tdmCopy.h index bdaef85c..954bef26 100644 --- a/src/header/tdmCopy.h +++ b/src/header/tdmCopy.h @@ -112,7 +112,7 @@ THE SOFTWARE. # else # define TDM_TOOLCHAIN_AVAILABLE 0 # endif -# if defined(__gfx1250__) && \ +# if defined(__GFX12__) && \ __has_builtin(__builtin_amdgcn_tensor_load_to_lds) && \ TDM_TOOLCHAIN_AVAILABLE /* extend: || (defined(__gfxNNNN__) && ...) */ From 85d92c19c1a077890bb393e38d472fce70894814 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Tue, 8 Sep 2026 01:37:22 -0500 Subject: [PATCH 04/13] pingpong latency test --- src/client/Client.cpp | 24 +- src/client/EnvVars.hpp | 108 +-- src/client/Utilities.hpp | 286 +++++--- src/header/TransferBench.hpp | 1194 +++++++++++++++++++++++++++++----- 4 files changed, 1296 insertions(+), 316 deletions(-) diff --git a/src/client/Client.cpp b/src/client/Client.cpp index 93e39954..2a653330 100644 --- a/src/client/Client.cpp +++ b/src/client/Client.cpp @@ -107,16 +107,24 @@ int main(int argc, char **argv) if (transfers.empty()) { Print("\n"); } else { - bool isMultiNode = GetNumRanks() > 1; for (size_t i = 0; i < transfers.size(); i++) { Transfer const& t = transfers[i]; - Print("Transfer %5lu: (%s->", i, MemDevicesToStr(t.srcs).c_str()); - if (isMultiNode) Print("R%d", t.exeDevice.exeRank); - Print("%c%d", ExeTypeStr[t.exeDevice.exeType], t.exeDevice.exeIndex); - if (t.exeDevice.exeSlot) Print("%c", 'A' + t.exeDevice.exeSlot); - if (t.exeSubIndex != -1) Print(".%d", t.exeSubIndex); - if (t.exeSubSlot != 0) Print("%c", 'A' + t.exeSubSlot); - Print("->%s)\n", MemDevicesToStr(t.dsts).c_str()); + if (t.numLaps > 0) { + Print("Transfer %5lu: PingPong +%d: (%s->%s->%s) <+> (%s->%s->%s)\n", + i, t.numLaps, + MemDeviceToStr(t.srcs[0]).c_str(), + ExeDeviceToStr(t.exeDevice, t.exeSubIndex, t.exeSubSlot).c_str(), + MemDeviceToStr(t.dsts[0]).c_str(), + MemDeviceToStr(t.srcs[1]).c_str(), + ExeDeviceToStr(t.exeDevicePong, t.exeSubIndexPong, t.exeSubSlotPong).c_str(), + MemDeviceToStr(t.dsts[1]).c_str()); + } else { + Print("Transfer %5lu: (%s->%s->%s)\n", + i, + MemDevicesToStr(t.srcs).c_str(), + ExeDeviceToStr(t.exeDevice, t.exeSubIndex, t.exeSubSlot).c_str(), + MemDevicesToStr(t.dsts).c_str()); + } } } return 0; diff --git a/src/client/EnvVars.hpp b/src/client/EnvVars.hpp index 575f6a67..dc74d6fa 100644 --- a/src/client/EnvVars.hpp +++ b/src/client/EnvVars.hpp @@ -74,6 +74,8 @@ class EnvVars int numIterations; // Number of timed iterations to perform. If negative, run for -numIterations seconds instead int numSubIterations; // Number of subiterations to perform int numWarmups; // Number of un-timed warmup iterations to perform + int pingpongFlagBuffer; // Size in bytes of the pingpong flag buffer (must be positive) + int pingpongStride; // Stride in bytes between flag slots for pingpong laps (wraps within the flag buffer) int showBorders; // Show ASCII box-drawing characaters in tables int showIterations; // Show per-iteration timing info int useInteractive; // Pause for user-input before starting transfer loop @@ -158,55 +160,57 @@ class EnvVars else if (archName == "gfx950") defaultGfxUnroll = 4; else if (archName == "gfx1250") defaultGfxUnroll = 32; - alwaysValidate = GetEnvVar("ALWAYS_VALIDATE", 0); - blockBytes = GetEnvVar("BLOCK_BYTES" , 256); - byteOffset = GetEnvVar("BYTE_OFFSET" , 0); - fillCompress = GetEnvVarArray("FILL_COMPRESS" , {}); - gfxBlockOrder = GetEnvVar("GFX_BLOCK_ORDER" , 0); - gfxBlockSize = GetEnvVar("GFX_BLOCK_SIZE" , 256); - gfxKernel = GetEnvVar("GFX_KERNEL" , 0); - gfxSeType = GetEnvVar("GFX_SE_TYPE" , 0); - gfxSingleTeam = GetEnvVar("GFX_SINGLE_TEAM" , 0); - gfxTemporal = GetEnvVar("GFX_TEMPORAL" , 0); - gfxUnroll = GetEnvVar("GFX_UNROLL" , defaultGfxUnroll); - gfxWaveOrder = GetEnvVar("GFX_WAVE_ORDER" , 0); - gfxWordSize = GetEnvVar("GFX_WORD_SIZE" , 4); - hideEnv = GetEnvVar("HIDE_ENV" , 0); - minNumVarSubExec = GetEnvVar("MIN_VAR_SUBEXEC" , 1); - maxNumVarSubExec = GetEnvVar("MAX_VAR_SUBEXEC" , 0); - numIterations = GetEnvVar("NUM_ITERATIONS" , 10); - numSubIterations = GetEnvVar("NUM_SUBITERATIONS" , 1); - numWarmups = GetEnvVar("NUM_WARMUPS" , 3); - outputToCsv = GetEnvVar("OUTPUT_TO_CSV" , 0); - samplingFactor = GetEnvVar("SAMPLING_FACTOR" , 1); - showBorders = GetEnvVar("SHOW_BORDERS" , 1); - showIterations = GetEnvVar("SHOW_ITERATIONS" , 0); - showPercentiles = GetEnvVarArray("SHOW_PERCENTILES", {}); - sweepMaxPow2 = GetEnvVar("SWEEP_MAX_POW2" , 29); - sweepMinPow2 = GetEnvVar("SWEEP_MIN_POW2" , 10); - sweepMinPow2 = std::clamp(sweepMinPow2, 0, 62); - sweepMaxPow2 = std::clamp(sweepMaxPow2, 0, 62); + alwaysValidate = GetEnvVar("ALWAYS_VALIDATE", 0); + blockBytes = GetEnvVar("BLOCK_BYTES" , 256); + byteOffset = GetEnvVar("BYTE_OFFSET" , 0); + fillCompress = GetEnvVarArray("FILL_COMPRESS" , {}); + gfxBlockOrder = GetEnvVar("GFX_BLOCK_ORDER" , 0); + gfxBlockSize = GetEnvVar("GFX_BLOCK_SIZE" , 256); + gfxKernel = GetEnvVar("GFX_KERNEL" , 0); + gfxSeType = GetEnvVar("GFX_SE_TYPE" , 0); + gfxSingleTeam = GetEnvVar("GFX_SINGLE_TEAM" , 0); + gfxTemporal = GetEnvVar("GFX_TEMPORAL" , 0); + gfxUnroll = GetEnvVar("GFX_UNROLL" , defaultGfxUnroll); + gfxWaveOrder = GetEnvVar("GFX_WAVE_ORDER" , 0); + gfxWordSize = GetEnvVar("GFX_WORD_SIZE" , 4); + hideEnv = GetEnvVar("HIDE_ENV" , 0); + minNumVarSubExec = GetEnvVar("MIN_VAR_SUBEXEC" , 1); + maxNumVarSubExec = GetEnvVar("MAX_VAR_SUBEXEC" , 0); + numIterations = GetEnvVar("NUM_ITERATIONS" , 10); + numSubIterations = GetEnvVar("NUM_SUBITERATIONS" , 1); + numWarmups = GetEnvVar("NUM_WARMUPS" , 3); + pingpongFlagBuffer = GetEnvVar("PINGPONG_FLAG_BUFFER", 1); + pingpongStride = GetEnvVar("PINGPONG_STRIDE" , 1); + outputToCsv = GetEnvVar("OUTPUT_TO_CSV" , 0); + samplingFactor = GetEnvVar("SAMPLING_FACTOR" , 1); + showBorders = GetEnvVar("SHOW_BORDERS" , 1); + showIterations = GetEnvVar("SHOW_ITERATIONS" , 0); + showPercentiles = GetEnvVarArray("SHOW_PERCENTILES", {}); + sweepMaxPow2 = GetEnvVar("SWEEP_MAX_POW2" , 29); + sweepMinPow2 = GetEnvVar("SWEEP_MIN_POW2" , 10); + sweepMinPow2 = std::clamp(sweepMinPow2, 0, 62); + sweepMaxPow2 = std::clamp(sweepMaxPow2, 0, 62); if (sweepMinPow2 > sweepMaxPow2) std::swap(sweepMinPow2, sweepMaxPow2); - tdmBlockOrder = GetEnvVar("TDM_BLOCK_ORDER" , 0); - tdmBlockSize = GetEnvVar("TDM_BLOCK_SIZE" , 256); - tdmLdsBytes = GetEnvVar("TDM_LDS_BYTES" , 0); - useHipEvents = GetEnvVar("USE_HIP_EVENTS" , 1); - useHsaDma = GetEnvVar("USE_HSA_DMA" , 0); - useInteractive = GetEnvVar("USE_INTERACTIVE" , 0); - useSingleStream = GetEnvVar("USE_SINGLE_STREAM" , 1); - validateDirect = GetEnvVar("VALIDATE_DIRECT" , 0); - validateOnDevice = GetEnvVar("VALIDATE_ON_DEVICE" , 0); - validateSource = GetEnvVar("VALIDATE_SOURCE" , 0); - - ibGidIndex = GetEnvVar("IB_GID_INDEX" ,-1); - ibPort = GetEnvVar("IB_PORT_NUMBER" , 1); - roceVersion = GetEnvVar("ROCE_VERSION" , 2); - ipAddressFamily = GetEnvVar("IP_ADDRESS_FAMILY" , 4); - nicChunkBytes = GetEnvVar("NIC_CHUNK_BYTES" , 1073741824); - nicCqPollBatch = GetEnvVar("NIC_CQ_POLL_BATCH" , 4); - nicRelaxedOrder = GetEnvVar("NIC_RELAX_ORDER" , 1); - nicServiceLevel = GetEnvVar("NIC_SERVICE_LEVEL" , 0); - nicTrafficClass = GetEnvVar("NIC_TRAFFIC_CLASS" , 0); + tdmBlockOrder = GetEnvVar("TDM_BLOCK_ORDER" , 0); + tdmBlockSize = GetEnvVar("TDM_BLOCK_SIZE" , 256); + tdmLdsBytes = GetEnvVar("TDM_LDS_BYTES" , 0); + useHipEvents = GetEnvVar("USE_HIP_EVENTS" , 1); + useHsaDma = GetEnvVar("USE_HSA_DMA" , 0); + useInteractive = GetEnvVar("USE_INTERACTIVE" , 0); + useSingleStream = GetEnvVar("USE_SINGLE_STREAM" , 1); + validateDirect = GetEnvVar("VALIDATE_DIRECT" , 0); + validateOnDevice = GetEnvVar("VALIDATE_ON_DEVICE" , 0); + validateSource = GetEnvVar("VALIDATE_SOURCE" , 0); + + ibGidIndex = GetEnvVar("IB_GID_INDEX" ,-1); + ibPort = GetEnvVar("IB_PORT_NUMBER" , 1); + roceVersion = GetEnvVar("ROCE_VERSION" , 2); + ipAddressFamily = GetEnvVar("IP_ADDRESS_FAMILY" , 4); + nicChunkBytes = GetEnvVar("NIC_CHUNK_BYTES" , 1073741824); + nicCqPollBatch = GetEnvVar("NIC_CQ_POLL_BATCH" , 4); + nicRelaxedOrder = GetEnvVar("NIC_RELAX_ORDER" , 1); + nicServiceLevel = GetEnvVar("NIC_SERVICE_LEVEL" , 0); + nicTrafficClass = GetEnvVar("NIC_TRAFFIC_CLASS" , 0); // Check that NIC service level and traffic class are in valid ranges if (nicServiceLevel < 0 || nicServiceLevel > 15) { @@ -404,6 +408,8 @@ class EnvVars printf(" NUM_SUBITERATIONS - # of sub-iterations to run per iteration. Must be non-negative\n"); printf(" NUM_WARMUPS - # of untimed warmup iterations per test\n"); printf(" OUTPUT_TO_CSV - Outputs to CSV format if set\n"); + printf(" PINGPONG_FLAG_BUFFER - Size in bytes of the pingpong flag buffer (default 1, must be positive). Each flag slot is 1 byte\n"); + printf(" PINGPONG_STRIDE - Stride in bytes between flag slots for pingpong laps (default 1, must be positive unless the flag buffer is 1 byte, wraps within the flag buffer)\n"); printf(" SAMPLING_FACTOR - Add this many samples (when possible) between powers of 2 when auto-generating data sizes\n"); printf(" SHOW_BORDERS - Show ASCII box-drawing characters in tables\n"); printf(" SHOW_ITERATIONS - Show per-iteration timing info\n"); @@ -555,6 +561,10 @@ class EnvVars "Running %s subiterations", (numSubIterations == 0 ? "infinite" : std::to_string(numSubIterations)).c_str()); Print("NUM_WARMUPS", numWarmups, "Running %d warmup iteration(s) per Test", numWarmups); + Print("PINGPONG_FLAG_BUFFER", pingpongFlagBuffer, + "Pingpong flag buffer of %d bytes", pingpongFlagBuffer); + Print("PINGPONG_STRIDE", pingpongStride, + "Pingpong flag stride %d bytes per lap", pingpongStride); Print("SHOW_BORDERS", showBorders, "%s ASCII box-drawing characaters in tables", showBorders ? "Showing" : "Hiding"); Print("SHOW_ITERATIONS", showIterations, "%s per-iteration timing", showIterations ? "Showing" : "Hiding"); @@ -738,6 +748,8 @@ class EnvVars cfg.general.numIterations = numIterations; cfg.general.numSubIterations = numSubIterations; cfg.general.numWarmups = numWarmups; + cfg.general.pingpongFlagBuffer = pingpongFlagBuffer; + cfg.general.pingpongStride = pingpongStride; cfg.general.recordPerIteration = ((showIterations != 0) || !showPercentiles.empty()) ? 1 : 0; cfg.general.useHipEvents = useHipEvents; cfg.general.useInteractive = useInteractive; diff --git a/src/client/Utilities.hpp b/src/client/Utilities.hpp index 6d69f29a..64406454 100644 --- a/src/client/Utilities.hpp +++ b/src/client/Utilities.hpp @@ -141,6 +141,12 @@ namespace TransferBench::Utils // Helper function that converts MemDevices to a string std::string MemDevicesToStr(std::vector const& memDevices); + // Helper function that converts a single MemDevice to a string ("N" for MEM_NULL) + std::string MemDeviceToStr(MemDevice const& memDevice); + + // Helper function that converts an ExeDevice (with subindex/subslot) to a string + std::string ExeDeviceToStr(ExeDevice const& exeDevice, int32_t subIndex = -1, int32_t subSlot = 0); + // Helper function to determine if current rank does output bool RankDoesOutput(); @@ -474,6 +480,33 @@ namespace TransferBench::Utils return ss.str(); } + std::string MemDeviceToStr(MemDevice const& memDevice) + { + if (memDevice.memType == TransferBench::MEM_NULL) return "N"; + bool isMultiNode = TransferBench::GetNumRanks() > 1; + std::stringstream ss; + if (isMultiNode) + ss << "R" << memDevice.memRank; + ss << TransferBench::MemTypeStr[memDevice.memType] << memDevice.memIndex; + return ss.str(); + } + + std::string ExeDeviceToStr(ExeDevice const& exeDevice, int32_t subIndex, int32_t subSlot) + { + bool isMultiNode = TransferBench::GetNumRanks() > 1; + std::stringstream ss; + if (isMultiNode) + ss << "R" << exeDevice.exeRank; + ss << TransferBench::ExeTypeStr[exeDevice.exeType] << exeDevice.exeIndex; + if (exeDevice.exeSlot) + ss << char('A' + exeDevice.exeSlot); + if (subIndex != -1) + ss << "." << subIndex; + if (subSlot != 0) + ss << char('A' + subSlot); + return ss.str(); + } + template struct is_std_vector : std::false_type {}; @@ -574,12 +607,20 @@ namespace TransferBench::Utils bool isMultiRank = TransferBench::GetNumRanks() > 1; + // The pong half owns no result row, so its executor shows up with no transfers beneath it + std::set pongExeDevices; + for (auto const& t : transfers) + if (t.numLaps > 0) pongExeDevices.insert(t.exeDevicePong); + // Figure out table dimensions int numCols = 5, numRows = 1; size_t numTimedIterations = results.numTimedIterations; for (auto const& exeInfoPair : results.exeResults) { ExeResult const& exeResult = exeInfoPair.second; - numRows += 1 + exeResult.transferIdx.size(); + int displayCount = 0; + for (int idx : exeResult.transferIdx) + if (transfers[idx].numLaps >= 0) displayCount++; + numRows += 1 + displayCount; if (!ev.showPercentiles.empty()) { numRows += static_cast(ev.showPercentiles.size()) * static_cast(exeResult.transferIdx.size()); } @@ -588,6 +629,7 @@ namespace TransferBench::Utils } if (ev.showIterations || !ev.showPercentiles.empty()) { for (int idx : exeResult.transferIdx) { + if (transfers[idx].numLaps > 0) continue; // pingpong iteration rows handled below TransferResult const& r = results.tfrResults[idx]; if (r.perIterMsec.size() != numTimedIterations) { Print("[ERROR] Per iteration timing data unavailable: Expected %lu data points, but have %lu\n", @@ -612,18 +654,33 @@ namespace TransferBench::Utils ExeType const exeType = exeDevice.exeType; int32_t const exeIndex = exeDevice.exeIndex; + // Executors running only pingpong halves move no payload, so bytes/bandwidth are meaningless + bool const isPingpongExe = (exeResult.numBytes == 0); + // Display Executor results table.DrawRowBorder(rowIdx); - if (isMultiRank) { + if (isMultiRank) table.Set(rowIdx, 0, " Executor: Rank %d %3s %02d ", exeDevice.exeRank, ExeTypeToStr(exeType).c_str(), exeIndex); - table.Set(rowIdx, 4, " %7.3f GB/s (sum) [%s]", exeResult.sumBandwidthGbPerSec, GetHostname(exeDevice.exeRank).c_str()); - } else { + else table.Set(rowIdx, 0, " Executor: %3s %02d ", ExeTypeToStr(exeType).c_str(), exeIndex); - table.Set(rowIdx, 4, " %7.3f GB/s (sum)", exeResult.sumBandwidthGbPerSec); + + std::string exeSummary; + if (isPingpongExe) { + exeSummary = pongExeDevices.count(exeDevice) && exeResult.transferIdx.empty() + ? " pingpong (pong half)" : " pingpong"; + table.Set(rowIdx, 1, " "); + table.Set(rowIdx, 3, " "); + } else { + char buf[64]; + snprintf(buf, sizeof(buf), " %7.3f GB/s (sum)", exeResult.sumBandwidthGbPerSec); + exeSummary = buf; + table.Set(rowIdx, 1, "%8.3f GB/s " , exeResult.avgBandwidthGbPerSec); + table.Set(rowIdx, 3, "%12lu bytes ", exeResult.numBytes); } - table.Set(rowIdx, 1, "%8.3f GB/s " , exeResult.avgBandwidthGbPerSec); - table.Set(rowIdx, 2, "%8.3f ms " , exeResult.avgDurationMsec); - table.Set(rowIdx, 3, "%12lu bytes ", exeResult.numBytes); + if (isMultiRank) exeSummary += " [" + GetHostname(exeDevice.exeRank) + "]"; + + table.Set(rowIdx, 2, "%8.3f ms ", exeResult.avgDurationMsec); + table.Set(rowIdx, 4, "%s", exeSummary.c_str()); table.SetCellAlignment(rowIdx, 4, TableHelper::ALIGN_LEFT); rowIdx++; table.DrawRowBorder(rowIdx); @@ -633,87 +690,138 @@ namespace TransferBench::Utils Transfer const& t = transfers[idx]; TransferResult const& r = results.tfrResults[idx]; - table.Set(rowIdx, 0, "Transfer %-4d ", idx); - table.Set(rowIdx, 1, "%8.3f GB/s " , r.avgBandwidthGbPerSec); - table.Set(rowIdx, 2, "%8.3f ms " , r.avgDurationMsec); - table.Set(rowIdx, 3, "%12lu bytes " , r.numBytes); - - char exeSubIndexStr[32] = ""; - if (t.exeSubIndex != -1) - sprintf(exeSubIndexStr, ".%d", t.exeSubIndex); - - if (isMultiRank) { - table.Set(rowIdx, 4, " %s -> R%d%c%d%s:%d -> %s", - MemDevicesToStr(t.srcs).c_str(), - exeDevice.exeRank, ExeTypeStr[t.exeDevice.exeType], t.exeDevice.exeIndex, - exeSubIndexStr, t.numSubExecs, - MemDevicesToStr(t.dsts).c_str()); + if (t.numLaps > 0) { + // Pingpong row: show latency using ping's round-trip delta + double latencyUs = r.avgDurationMsec * 1000.0; + table.Set(rowIdx, 0, "PingPong %-4d ", idx); + table.Set(rowIdx, 1, "%8.3f us " , latencyUs); + table.Set(rowIdx, 2, "%8.3f ms " , r.avgDurationMsec); + table.Set(rowIdx, 3, "%8d laps " , t.numLaps); + + if (isMultiRank) { + table.Set(rowIdx, 4, " %s->R%d%c%d->%s <+> %s->R%d%c%d->%s", + MemDeviceToStr(t.srcs[0]).c_str(), + t.exeDevice.exeRank, ExeTypeStr[t.exeDevice.exeType], t.exeDevice.exeIndex, + MemDeviceToStr(t.dsts[0]).c_str(), + MemDeviceToStr(t.srcs[1]).c_str(), + t.exeDevicePong.exeRank, ExeTypeStr[t.exeDevicePong.exeType], t.exeDevicePong.exeIndex, + MemDeviceToStr(t.dsts[1]).c_str()); + } else { + table.Set(rowIdx, 4, " %s->%c%d->%s <+> %s->%c%d->%s", + MemDeviceToStr(t.srcs[0]).c_str(), + ExeTypeStr[t.exeDevice.exeType], t.exeDevice.exeIndex, + MemDeviceToStr(t.dsts[0]).c_str(), + MemDeviceToStr(t.srcs[1]).c_str(), + ExeTypeStr[t.exeDevicePong.exeType], t.exeDevicePong.exeIndex, + MemDeviceToStr(t.dsts[1]).c_str()); + } + table.SetCellAlignment(rowIdx, 4, TableHelper::ALIGN_LEFT); + rowIdx++; + + if (ev.showIterations) { + std::set> times; + double stdDevTime = 0; + for (size_t i = 0; i < numTimedIterations; i++) { + times.insert(std::make_pair(r.perIterMsec[i], i+1)); + double const varTime = fabs(r.avgDurationMsec - r.perIterMsec[i]); + stdDevTime += varTime * varTime; + } + stdDevTime = sqrt(stdDevTime / numTimedIterations); + + for (auto& time : times) { + double iterUs = time.first * 1000.0; + table.Set(rowIdx, 0, "Iter %03d ", time.second); + table.Set(rowIdx, 1, "%8.3f us ", iterUs); + table.Set(rowIdx, 2, "%8.3f ms ", time.first); + rowIdx++; + } + + table.Set(rowIdx, 0, "StandardDev "); + table.Set(rowIdx, 1, "%8.3f us ", stdDevTime * 1000.0); + table.Set(rowIdx, 2, "%8.3f ms ", stdDevTime); + rowIdx++; + table.DrawRowBorder(rowIdx); + } } else { - table.Set(rowIdx, 4, " %s -> %c%d%s:%d -> %s", - MemDevicesToStr(t.srcs).c_str(), - ExeTypeStr[t.exeDevice.exeType], t.exeDevice.exeIndex, - exeSubIndexStr, t.numSubExecs, - MemDevicesToStr(t.dsts).c_str()); - } - table.SetCellAlignment(rowIdx, 4, TableHelper::ALIGN_LEFT); - rowIdx++; - - // Show per-iteration timing information - if (ev.showIterations) { - - // Compute standard deviation and track iterations by speed - std::set> times; - double stdDevTime = 0; - double stdDevBw = 0; - for (int i = 0; i < numTimedIterations; i++) { - times.insert(std::make_pair(r.perIterMsec[i], i+1)); - double const varTime = fabs(r.avgDurationMsec - r.perIterMsec[i]); - stdDevTime += varTime * varTime; - - double iterBandwidthGbs = (t.numBytes / 1.0E9) / r.perIterMsec[i] * 1000.0f; - double const varBw = fabs(iterBandwidthGbs - r.avgBandwidthGbPerSec); - stdDevBw += varBw * varBw; + // Regular transfer row (numLaps == 0) + table.Set(rowIdx, 0, "Transfer %-4d ", idx); + table.Set(rowIdx, 1, "%8.3f GB/s " , r.avgBandwidthGbPerSec); + table.Set(rowIdx, 2, "%8.3f ms " , r.avgDurationMsec); + table.Set(rowIdx, 3, "%12lu bytes " , r.numBytes); + + char exeSubIndexStr[32] = ""; + if (t.exeSubIndex != -1) + sprintf(exeSubIndexStr, ".%d", t.exeSubIndex); + + if (isMultiRank) { + table.Set(rowIdx, 4, " %s -> R%d%c%d%s:%d -> %s", + MemDevicesToStr(t.srcs).c_str(), + exeDevice.exeRank, ExeTypeStr[t.exeDevice.exeType], t.exeDevice.exeIndex, + exeSubIndexStr, t.numSubExecs, + MemDevicesToStr(t.dsts).c_str()); + } else { + table.Set(rowIdx, 4, " %s -> %c%d%s:%d -> %s", + MemDevicesToStr(t.srcs).c_str(), + ExeTypeStr[t.exeDevice.exeType], t.exeDevice.exeIndex, + exeSubIndexStr, t.numSubExecs, + MemDevicesToStr(t.dsts).c_str()); } - stdDevTime = sqrt(stdDevTime / numTimedIterations); - stdDevBw = sqrt(stdDevBw / numTimedIterations); - - // Loop over iterations (fastest to slowest) - for (auto& time : times) { - double iterDurationMsec = time.first; - double iterBandwidthGbs = (t.numBytes / 1.0E9) / iterDurationMsec * 1000.0f; - - std::set usedXccs; - std::stringstream ss1; - if (exeDevice.exeType == EXE_GPU_GFX) { - if (time.second - 1 < r.perIterCUs.size()) { - ss1 << " CUs: "; - for (auto x : r.perIterCUs[time.second - 1]) { - ss1 << x.first << ":" << std::setfill('0') << std::setw(2) << x.second << " "; - usedXccs.insert(x.first); + table.SetCellAlignment(rowIdx, 4, TableHelper::ALIGN_LEFT); + rowIdx++; + + if (ev.showIterations) { + std::set> times; + double stdDevTime = 0; + double stdDevBw = 0; + for (size_t i = 0; i < numTimedIterations; i++) { + times.insert(std::make_pair(r.perIterMsec[i], i+1)); + double const varTime = fabs(r.avgDurationMsec - r.perIterMsec[i]); + stdDevTime += varTime * varTime; + + double iterBandwidthGbs = (t.numBytes / 1.0E9) / r.perIterMsec[i] * 1000.0f; + double const varBw = fabs(iterBandwidthGbs - r.avgBandwidthGbPerSec); + stdDevBw += varBw * varBw; + } + stdDevTime = sqrt(stdDevTime / numTimedIterations); + stdDevBw = sqrt(stdDevBw / numTimedIterations); + + for (auto& time : times) { + double iterDurationMsec = time.first; + double iterBandwidthGbs = (t.numBytes / 1.0E9) / iterDurationMsec * 1000.0f; + + std::set usedXccs; + std::stringstream ss1; + if (exeDevice.exeType == EXE_GPU_GFX) { + if (time.second - 1 < r.perIterCUs.size()) { + ss1 << " CUs: "; + for (auto x : r.perIterCUs[time.second - 1]) { + ss1 << x.first << ":" << std::setfill('0') << std::setw(2) << x.second << " "; + usedXccs.insert(x.first); + } } } - } - std::stringstream ss2; - if (!usedXccs.empty()) { - ss2 << " XCCs:"; - for (auto x : usedXccs) - ss2 << " " << x; + std::stringstream ss2; + if (!usedXccs.empty()) { + ss2 << " XCCs:"; + for (auto x : usedXccs) + ss2 << " " << x; + } + + table.Set(rowIdx, 0, "Iter %03d ", time.second); + table.Set(rowIdx, 1, "%8.3f GB/s ", iterBandwidthGbs); + table.Set(rowIdx, 2, "%8.3f ms ", iterDurationMsec); + table.Set(rowIdx, 3, ss1.str()); + table.Set(rowIdx, 4, ss2.str()); + rowIdx++; } - table.Set(rowIdx, 0, "Iter %03d ", time.second); - table.Set(rowIdx, 1, "%8.3f GB/s ", iterBandwidthGbs); - table.Set(rowIdx, 2, "%8.3f ms ", iterDurationMsec); - table.Set(rowIdx, 3, ss1.str()); - table.Set(rowIdx, 4, ss2.str()); + table.Set(rowIdx, 0, "StandardDev "); + table.Set(rowIdx, 1, "%8.3f GB/s ", stdDevBw); + table.Set(rowIdx, 2, "%8.3f ms ", stdDevTime); rowIdx++; + table.DrawRowBorder(rowIdx); } - - table.Set(rowIdx, 0, "StandardDev "); - table.Set(rowIdx, 1, "%8.3f GB/s ", stdDevBw); - table.Set(rowIdx, 2, "%8.3f ms ", stdDevTime); - rowIdx++; - table.DrawRowBorder(rowIdx); } // Show percentiles @@ -722,9 +830,13 @@ namespace TransferBench::Utils std::sort(sortedDur.begin(), sortedDur.end()); for (int pct : ev.showPercentiles) { double dur = PercentileDurationMsecFromSorted(sortedDur, pct); - double bwGbs = dur > 0.0 ? (t.numBytes / 1.0E9) / dur * 1000.0 : 0.0; table.Set(rowIdx, 0, "p%d ", pct); - table.Set(rowIdx, 1, "%8.3f GB/s ", bwGbs); + if (t.numLaps > 0) { + table.Set(rowIdx, 1, "%8.3f us ", dur * 1000.0); + } else { + double bwGbs = dur > 0.0 ? (t.numBytes / 1.0E9) / dur * 1000.0 : 0.0; + table.Set(rowIdx, 1, "%8.3f GB/s ", bwGbs); + } table.Set(rowIdx, 2, "%8.3f ms ", dur); table.Set(rowIdx, 3, " "); table.Set(rowIdx, 4, " "); @@ -737,9 +849,15 @@ namespace TransferBench::Utils } table.DrawRowBorder(rowIdx); table.Set(rowIdx, 0, "Aggregate (CPU) "); - table.Set(rowIdx, 1, "%8.3f GB/s " , results.avgTotalBandwidthGbPerSec); + if (results.totalBytesTransferred == 0) { + // Pingpong-only run: no payload was moved, so leave the bandwidth/byte cells empty + table.Set(rowIdx, 1, " "); + table.Set(rowIdx, 3, " "); + } else { + table.Set(rowIdx, 1, "%8.3f GB/s " , results.avgTotalBandwidthGbPerSec); + table.Set(rowIdx, 3, "%12lu bytes " , results.totalBytesTransferred); + } table.Set(rowIdx, 2, "%8.3f ms " , results.avgTotalDurationMsec); - table.Set(rowIdx, 3, "%12lu bytes " , results.totalBytesTransferred); table.Set(rowIdx, 4, " Overhead %.3f ms", results.overheadMsec); table.SetCellAlignment(rowIdx, 4, TableHelper::ALIGN_LEFT); table.DrawRowBorder(rowIdx+1); diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index a2e207ac..48b46960 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -45,6 +45,7 @@ THE SOFTWARE. #include #include // If not found, try installing libnuma-dev (e.g apt-get install libnuma-dev) #include +#include #include #include #include @@ -198,16 +199,39 @@ namespace TransferBench /** * A Transfer adds together data from zero or more sources then writes the sum to zero or more desintations + * + * Normal transfer (numLaps == 0): + * srcs/dsts are variable-length lists as today. + * + * Pingpong transfer (numLaps > 0): + * srcs and dsts are each exactly 2 entries: + * [0] = ping half, [1] = pong half + * dsts[0] and dsts[1] are required (non-NULL); srcs[0] and srcs[1] are optional (MEM_NULL if absent) + * Ping/pong MemDevices are not restricted — any supported MemType is allowed per slot. + * exeDevice / exeSubIndex / exeSubSlot = ping executor (any ExeType) + * exeDevicePong / exeSubIndexPong / exeSubSlotPong = pong executor (any ExeType) + * numLaps = number of pingpong laps (must be > 0); specify as "+N" after ping half ("+" alone defaults to 1) + * + * Ping/pong executors may differ in type (e.g. ping on GFX, pong on DMA). The struct is + * executor-agnostic; additional executor types are enabled by implementing their dispatch + * paths. As of now, pingpong execution is implemented for EXE_GPU_GFX only. + * + * On a GFX executor, all ping/pong halves assigned to that executor are launched on one + * dedicated HIP stream, with one threadblock per half (transfer.numSubExecs is ignored). */ struct Transfer { size_t numBytes = 0; ///< Number of bytes to Transfer - vector srcs = {}; ///< List of source memory devices - vector dsts = {}; ///< List of destination memory devices - ExeDevice exeDevice = {}; ///< Executor to use - int32_t exeSubIndex = -1; ///< Executor subindex - int32_t exeSubSlot = 0; ///< Executor subslot + vector srcs = {}; ///< Source memory (pingpong: [ping, pong]) + vector dsts = {}; ///< Destination memory (pingpong: [ping, pong]) + ExeDevice exeDevice = {}; ///< (Transfer or Ping) Executor to use + int32_t exeSubIndex = -1; ///< (Transfer or Ping) Executor subindex + int32_t exeSubSlot = 0; ///< (Transfer or Ping) Executor subslot int numSubExecs = 0; ///< Number of subExecutors to use for this Transfer + int numLaps = 0; ///< 0 = normal transfer; >0 = pingpong lap count + ExeDevice exeDevicePong = {}; ///< Pong executor (pingpong only) + int32_t exeSubIndexPong = -1; ///< Pong executor subindex (pingpong only) + int32_t exeSubSlotPong = 0; ///< Pong executor subslot (pingpong only) }; /** @@ -222,6 +246,8 @@ namespace TransferBench int useHipEvents = 1; ///< Use HIP events for timing Executors that support it int useInteractive = 0; ///< Pause for user-input before starting transfer loop int useMultiStream = 0; ///< Split GFX/TDM Transfers into separate kernel launches in separate stream + int pingpongStride = 1; ///< Stride in bytes between flag slots for pingpong laps (positive, or 0 when pingpongFlagBuffer is 1) + int pingpongFlagBuffer = 1; ///< Size of the pingpong flag buffer in bytes (must be positive) }; /** @@ -2105,6 +2131,8 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (general.numIterations != cfg.general.numIterations) ADD_ERROR("cfg.general.numIterations"); if (general.numSubIterations != cfg.general.numSubIterations) ADD_ERROR("cfg.general.numSubIterations"); if (general.numWarmups != cfg.general.numWarmups) ADD_ERROR("cfg.general.numWarmups"); + if (general.pingpongStride != cfg.general.pingpongStride) ADD_ERROR("cfg.general.pingpongStride"); + if (general.pingpongFlagBuffer != cfg.general.pingpongFlagBuffer) ADD_ERROR("cfg.general.pingpongFlagBuffer"); if (general.recordPerIteration != cfg.general.recordPerIteration) ADD_ERROR("cfg.general.recordPerIteration"); if (general.useHipEvents != cfg.general.useHipEvents) ADD_ERROR("cfg.general.useHipEvents"); if (general.useInteractive != cfg.general.useInteractive) ADD_ERROR("cfg.general.useInteractive"); @@ -2216,7 +2244,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) #undef ADD_ERROR } - // Forward declaration + // Forward declarations int GetGpuKernelUnrollIdx(int unroll); // Validate configuration options - return trues if and only if an fatal error is detected @@ -2226,6 +2254,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Check general options if (cfg.general.numWarmups < 0) errors.push_back({ERR_FATAL, "[general.numWarmups] must be a non-negative number"}); + if (cfg.general.pingpongFlagBuffer < 1) + errors.push_back({ERR_FATAL, "[general.pingpongFlagBuffer] must be a positive number of bytes"}); + // A 0 stride keeps every lap on the same flag slot, which only alternates safely when the + // buffer is a single byte and there is no other slot to rotate through + if (cfg.general.pingpongStride < 0 || + (cfg.general.pingpongStride == 0 && cfg.general.pingpongFlagBuffer != 1)) + errors.push_back({ERR_FATAL, "[general.pingpongStride] must be a positive number of bytes " + "(0 is only allowed when [general.pingpongFlagBuffer] is 1)"}); // Check that config options are consistent (where necessary) across all ranks CheckMultiNodeConfigConsistency(cfg, errors); @@ -2412,6 +2448,10 @@ const auto& AmdSmiFabricInfoV1(const T& info) System::Get().Broadcast(root, sizeof(t.exeSubIndex), &t.exeSubIndex); System::Get().Broadcast(root, sizeof(t.exeSubSlot), &t.exeSubSlot); System::Get().Broadcast(root, sizeof(t.numSubExecs), &t.numSubExecs); + System::Get().Broadcast(root, sizeof(t.numLaps), &t.numLaps); + System::Get().Broadcast(root, sizeof(t.exeDevicePong), &t.exeDevicePong); + System::Get().Broadcast(root, sizeof(t.exeSubIndexPong), &t.exeSubIndexPong); + System::Get().Broadcast(root, sizeof(t.exeSubSlotPong), &t.exeSubSlotPong); if (t.numBytes != transfers[i].numBytes) ADD_ERROR("numBytes"); if (t.srcs != transfers[i].srcs) ADD_ERROR("Source memory locations"); @@ -2421,6 +2461,11 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (t.exeSubIndex != transfers[i].exeSubIndex) ADD_ERROR("Executor subindex"); if (t.exeSubSlot != transfers[i].exeSubSlot) ADD_ERROR("Executor dst slot"); if (t.numSubExecs != transfers[i].numSubExecs) ADD_ERROR("Num SubExecutors"); + if (t.numLaps != transfers[i].numLaps) ADD_ERROR("numLaps"); + if (t.exeDevicePong < transfers[i].exeDevicePong || + transfers[i].exeDevicePong < t.exeDevicePong) ADD_ERROR("Pong executor device"); + if (t.exeSubIndexPong != transfers[i].exeSubIndexPong) ADD_ERROR("Pong executor subindex"); + if (t.exeSubSlotPong != transfers[i].exeSubSlotPong) ADD_ERROR("Pong executor subslot"); } if (isInconsistent && !System::Get().IsVerbose()) { @@ -2433,6 +2478,15 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Returns true if the given Transfer requires pod communication static bool IsPodTransfer(Transfer const& t) { + if (t.numLaps > 0) { + for (int half = 0; half < 2; half++) { + ExeDevice const& exe = half == 0 ? t.exeDevice : t.exeDevicePong; + if (t.srcs[half].memType != MEM_NULL && t.srcs[half].memRank != exe.exeRank) return true; + if (t.dsts[half].memType != MEM_NULL && t.dsts[half].memRank != exe.exeRank) return true; + } + return false; + } + if (IsCpuExeType(t.exeDevice.exeType) || IsGpuExeType(t.exeDevice.exeType)) { for (auto const& src : t.srcs) if (src.memRank != t.exeDevice.exeRank) return true; @@ -2451,6 +2505,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) std::map transferCount; std::map useSubIndexCount; std::map totalSubExecs; + std::map totalPingpong; // Check that the set of requested transfers is consistent across all ranks CheckMultiNodeTransferConsistency(transfers, errors); @@ -2486,8 +2541,29 @@ const auto& AmdSmiFabricInfoV1(const T& info) i, maxSubExecToUse, t.numSubExecs, cfg.data.blockBytes}); } + bool const isPingpong = t.numLaps > 0; + if (t.numLaps < 0) { + errors.push_back({ERR_FATAL, + "Transfer %zu: numLaps must be > 0 for pingpong transfers (negative values belong on resources)", i}); + hasFatalError = true; + break; + } + // Check sources and destinations - if (t.srcs.empty() && t.dsts.empty()) { + if (isPingpong) { + if (t.srcs.size() != 2 || t.dsts.size() != 2) { + errors.push_back({ERR_FATAL, + "Transfer %zu: Pingpong transfer requires srcs and dsts each have exactly 2 entries ([ping, pong])", i}); + hasFatalError = true; + break; + } + if (t.dsts[0].memType == MEM_NULL || t.dsts[1].memType == MEM_NULL) { + errors.push_back({ERR_FATAL, + "Transfer %zu: Pingpong transfer requires a non-NULL dst for both ping and pong halves", i}); + hasFatalError = true; + break; + } + } else if (t.srcs.empty() && t.dsts.empty()) { errors.push_back({ERR_FATAL, "Transfer %d: Must have at least one source or destination", i}); break; } @@ -2512,6 +2588,56 @@ const auto& AmdSmiFabricInfoV1(const T& info) } if (hasFatalError) break; + if (isPingpong) { + if (t.exeDevice.exeType != EXE_GPU_GFX || t.exeDevicePong.exeType != EXE_GPU_GFX) { + errors.push_back({ERR_FATAL, + "Transfer %zu: Pingpong is currently supported on GFX executors only (ping type %c, pong type %c)", + i, ExeTypeStr[t.exeDevice.exeType], ExeTypeStr[t.exeDevicePong.exeType]}); + hasFatalError = true; + break; + } + + auto validateGfxExecutor = [&](ExeDevice const& exe, int32_t subIndex, char const* role) { + if (exe.exeRank < 0 || exe.exeRank >= GetNumRanks()) { + errors.push_back({ERR_FATAL, + "Transfer %zu: %s executor rank must be between 0 and %d (instead of %d)", + i, role, GetNumRanks() - 1, exe.exeRank}); + return true; + } + executors.insert(exe); + transferCount[exe]++; + int numExecutors = GetNumExecutors(EXE_GPU_GFX, exe.exeRank); + if (exe.exeIndex < 0 || exe.exeIndex >= numExecutors) { + errors.push_back({ERR_FATAL, + "Transfer %zu: %s GFX index must be between 0 and %d (instead of %d) for rank %d", + i, role, numExecutors - 1, exe.exeIndex, exe.exeRank}); + return true; + } + if (subIndex != -1) { +#if defined(__NVCC__) + errors.push_back({ERR_FATAL, + "Transfer %zu: %s GFX executor subindex not supported on NVIDIA hardware", i, role}); + return true; +#else + useSubIndexCount[exe]++; + int numSubIndices = GetNumExecutorSubIndices(exe); + if (subIndex >= numSubIndices) { + errors.push_back({ERR_FATAL, + "Transfer %zu: %s GFX subIndex (XCC) must be between 0 and %d for rank %d", + i, role, numSubIndices - 1, exe.exeRank}); + return true; + } +#endif + } + return false; + }; + + if (validateGfxExecutor(t.exeDevice, t.exeSubIndex, "Ping") || + validateGfxExecutor(t.exeDevicePong, t.exeSubIndexPong, "Pong")) { + hasFatalError = true; + break; + } + } else { // Check executor rank if (t.exeDevice.exeRank < 0 || t.exeDevice.exeRank >= GetNumRanks()) { errors.push_back({ERR_FATAL, @@ -2842,6 +2968,8 @@ const auto& AmdSmiFabricInfoV1(const T& info) break; } + } // !isPingpong + // Skip further tests if fatal error detected if (hasFatalError) break; @@ -2854,18 +2982,31 @@ const auto& AmdSmiFabricInfoV1(const T& info) break; #endif // In order to support pod communication, the participanting ranks need to be members of the same pod - int exeRank = t.exeDevice.exeRank; bool samePod = true; - for (auto const& src : t.srcs) { - if (!(samePod = IsSamePod(src.memRank, exeRank))) - break; - } - if (samePod) { - for (auto const& dst : t.dsts) { - if (!(samePod = IsSamePod(dst.memRank, exeRank))) + if (isPingpong) { + for (int half = 0; half < 2 && samePod; half++) { + ExeDevice const& exe = half == 0 ? t.exeDevice : t.exeDevicePong; + int const exeRank = exe.exeRank; + if (t.srcs[half].memType != MEM_NULL && + !(samePod = IsSamePod(t.srcs[half].memRank, exeRank))) + break; + if (t.dsts[half].memType != MEM_NULL && + !(samePod = IsSamePod(t.dsts[half].memRank, exeRank))) break; } + } else { + int exeRank = t.exeDevice.exeRank; + for (auto const& src : t.srcs) { + if (!(samePod = IsSamePod(src.memRank, exeRank))) + break; + } + if (samePod) { + for (auto const& dst : t.dsts) { + if (!(samePod = IsSamePod(dst.memRank, exeRank))) + break; + } + } } if (!samePod || IsCpuExeType(t.exeDevice.exeType)) { @@ -2877,7 +3018,35 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Pod (cross-rank) transfers with a GPU executor are exchanged via CUDA/HIP fabric handles, // which, for current version, only support device backed allocations. // Reject host memory allocations up front. - if (IsGpuExeType(t.exeDevice.exeType)) { + if (isPingpong) { + for (int half = 0; half < 2 && !hasFatalError; half++) { + ExeDevice const& exe = half == 0 ? t.exeDevice : t.exeDevicePong; + if (!IsGpuExeType(exe.exeType)) continue; + if (t.srcs[half].memType != MEM_NULL && + t.srcs[half].memRank != exe.exeRank && IsCpuMemType(t.srcs[half].memType)) { + errors.push_back({ERR_FATAL, + "Transfer %d: Cross-rank GPU executor (R%d%c%d) cannot access remote host memory " + "(%s on rank %d is %s). Fabric-handle sharing only supports GPU memory for 1.67; use a NIC " + "executor (e.g. R%dN..) for cross-rank transfers involving host memory.", + i, exe.exeRank, ExeTypeStr[exe.exeType], exe.exeIndex, + "SRC", t.srcs[half].memRank, GetMemTypeName(t.srcs[half].memType), exe.exeRank}); + hasFatalError = true; + break; + } + if (t.dsts[half].memType != MEM_NULL && + t.dsts[half].memRank != exe.exeRank && IsCpuMemType(t.dsts[half].memType)) { + errors.push_back({ERR_FATAL, + "Transfer %d: Cross-rank GPU executor (R%d%c%d) cannot access remote host memory " + "(%s on rank %d is %s). Fabric-handle sharing only supports GPU memory for 1.67; use a NIC " + "executor (e.g. R%dN..) for cross-rank transfers involving host memory.", + i, exe.exeRank, ExeTypeStr[exe.exeType], exe.exeIndex, + "DST", t.dsts[half].memRank, GetMemTypeName(t.dsts[half].memType), exe.exeRank}); + hasFatalError = true; + break; + } + } + if (hasFatalError) break; + } else if (IsGpuExeType(t.exeDevice.exeType)) { bool hasRemoteCpuMem = false; MemDevice offender = {}; char const* role = nullptr; @@ -2909,8 +3078,16 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Check subexecutors if (t.numSubExecs <= 0) errors.push_back({ERR_FATAL, "Transfer %d: # of subexecutors must be positive", i}); - else + else if (isPingpong) { + if (t.numSubExecs != 1) + errors.push_back({ERR_WARN, + "Transfer %d: pingpong uses one threadblock per half; numSubExecs (%d) is ignored", + i, t.numSubExecs}); + totalPingpong[t.exeDevice]++; + totalPingpong[t.exeDevicePong]++; + } else { totalSubExecs[t.exeDevice] += t.numSubExecs; + } } @@ -2941,11 +3118,12 @@ const auto& AmdSmiFabricInfoV1(const T& info) int warpsPerBlock = cfg.gfx.blockSize / GetWarpSize(&errors); numGpuSubExec *= warpsPerBlock; } - if (totalSubExecs[exeDevice] > numGpuSubExec) + int totalTotalSubExecs = totalSubExecs[exeDevice] + totalPingpong[exeDevice]; + if (totalTotalSubExecs > numGpuSubExec) errors.push_back({ERR_WARN, "GPU %d requests %d total %s however only %d available. " "Serialization will occur", - exeDevice.exeIndex, totalSubExecs[exeDevice], + exeDevice.exeIndex, totalTotalSubExecs, cfg.gfx.seType == 0 ? "CUs" : "warps", numGpuSubExec}); // Check that if executor subindices are used, all Transfers specify executor subindices if (useSubIndexCount[exeDevice] > 0 && useSubIndexCount[exeDevice] != transferCount[exeDevice]) { @@ -2956,10 +3134,22 @@ const auto& AmdSmiFabricInfoV1(const T& info) break; } - if (cfg.general.useMultiStream && transferCount[exeDevice] > gpuMaxHwQueues) { - errors.push_back({ERR_WARN, - "GPU %d attempting %d parallel transfers, however GPU_MAX_HW_QUEUES only set to %d", - exeDevice.exeIndex, transferCount[exeDevice], gpuMaxHwQueues}); + // Parallel HIP stream count: one stream per normal transfer in multistream mode + // (or one shared stream for all normal transfers otherwise), plus one shared + // pingpong stream when this executor has any ping/pong halves. + { + int const normalTransferCount = transferCount[exeDevice] - totalPingpong[exeDevice]; + int streamCount = cfg.general.useMultiStream ? normalTransferCount + : (normalTransferCount > 0 ? 1 : 0); + if (totalPingpong[exeDevice] > 0) + streamCount++; + if (streamCount > gpuMaxHwQueues) { + errors.push_back({ERR_WARN, + "GPU %d attempting %d parallel streams (%d normal transfer stream(s)" + " + %d pingpong stream), however GPU_MAX_HW_QUEUES only set to %d", + exeDevice.exeIndex, streamCount, normalTransferCount, + totalPingpong[exeDevice] > 0 ? 1 : 0, gpuMaxHwQueues}); + } } break; } @@ -3047,6 +3237,23 @@ const auto& AmdSmiFabricInfoV1(const T& info) uint32_t xccId; ///< XCC ID }; + // Pingpong parameters (parallel to SubExecParam; one entry per Ping/Pong) + struct PingpongParam + { + volatile uint8_t* srcMem[2]; ///< Device pointers to uint8_t values {0, 1} (even/odd laps) + volatile uint8_t* localFlagMem; ///< Partner half's dst; poll here for signal arrival + volatile uint8_t* flagMem; ///< Own dst; write here to signal partner + int numLaps; ///< 0 = normal, >0 = ping, <0 = pong + int flagStride; ///< Stride in bytes between flag slots per lap + int flagAllocBytes; ///< Total flag allocation size in bytes (for wrap-around) + int hopPeriod; ///< Laps between extra stride hops (0 = never hop) + int32_t preferredXccId; ///< XCC ID to execute on (GFX only) + + // Outputs (ping half only) + int64_t startCycle; ///< Start timestamp for in-kernel timing + int64_t stopCycle; ///< Stop timestamp for in-kernel timing + }; + // Internal resources allocated per Transfer typedef hipMemGenericAllocationHandle_t memHandle_t; struct TransferResources @@ -3060,11 +3267,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) vector srcMemHandle; ///< Memory handles for source memory vector dstMemHandle; ///< Memory handles for destination memory vector subExecParamCpu; ///< Defines subarrays for each subexecutor + PingpongParam pingpongParamCpu; ///< Pingpong parameter vector subExecIdx; ///< Indices into subExecParamGpu + int pingpongParamIdx = -1; ///< Index into exeInfo.pingpongParamCpu/Gpu int numaNode; ///< NUMA node to use for this Transfer // For GFX executor SubExecParam* subExecParamGpuPtr; + PingpongParam* pingpongParamGpuPtr; // For on-device validation (VALIDATE_ON_DEVICE) vector dstExpectedMem; ///< Per-dst device copy of expected values @@ -3115,6 +3325,9 @@ const auto& AmdSmiFabricInfoV1(const T& info) vector batchBytes; ///< Bytes to copy (per batch item) #endif + // Pingpong role/lap count on this resource half (0 = normal, >0 = ping, <0 = pong) + int numLaps = 0; + // Counters double totalDurationMsec; ///< Total duration for all iterations for this Transfer vector perIterMsec; ///< Duration for each individual iteration @@ -3223,22 +3436,26 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Internal resources allocated per Executor struct ExeInfo { - size_t totalBytes; ///< Total bytes this executor transfers - double totalDurationMsec; ///< Total duration for all iterations for this Executor - int totalSubExecs; ///< Total number of subExecutors to use - bool useSubIndices; ///< Use subexecutor indicies - int numSubIndices; ///< Number of subindices this ExeDevice has - vector subExecParamCpu; ///< Subexecutor parameters for this executor - vector resources; ///< Per-Transfer resources + size_t totalBytes; ///< Total bytes this executor transfers + double totalDurationMsec; ///< Total duration for all iterations for this Executor + int totalSubExecs; ///< Total normal-transfer threadblocks across this executor + int totalPingpong; ///< Total pingpong threadblocks (one per ping/pong resource half) + bool useSubIndices; ///< Use subexecutor indicies + int numSubIndices; ///< Number of subindices this ExeDevice has + vector subExecParamCpu; ///< Subexecutor parameters for this executor + vector pingpongParamCpu; ///< Pingpong parameters for this executor + vector resources; ///< Per-transfer resources (normal and ping/pong halves) // For GPU-Executors - SubExecParam* subExecParamGpu; ///< GPU copy of subExecutor parameters + SubExecParam* subExecParamGpu; ///< GPU copy of subExecutor parameters + PingpongParam* pingpongParamGpu; ///< GPU copy of pingpong parameters bool subExecParamHostAccessible; ///< Host can directly read subExecParamGpu - vector streams; ///< HIP streams to launch on - vector startEvents; ///< HIP start timing event - vector stopEvents; ///< HIP stop timing event - int wallClockRate; ///< (GFX-only) Device wall clock rate - int gfxKernelToUse; ///< (GFX-only) Which GFX kernel to use + bool pingpongParamHostAccessible; ///< Host can directly read pingpongParamGpu + vector streams; ///< HIP streams (normal transfers, then pingpong) + vector startEvents; ///< HIP start timing event + vector stopEvents; ///< HIP stop timing event + int wallClockRate; ///< (GFX-only) Device wall clock rate + int gfxKernelToUse; ///< (GFX-only) Which GFX kernel to use // For TDM-Executors uint32_t ldsBytesActual; ///< Actual number of LDS bytes to use as buffer @@ -4458,6 +4675,10 @@ const auto& AmdSmiFabricInfoV1(const T& info) int transferIdx = rss->transferIdx; Transfer const& t = transfers[transferIdx]; + // Pingpong halves exchange handshake flags instead of data, so their destinations + // hold lap flags that no dstReference entry describes + if (t.numLaps != 0) continue; + float const* expected = dstReference[t.srcs.size()].data(); for (int dstIdx = 0; dstIdx < (int)rss->dstMem.size(); dstIdx++) { // Validation is only done on the rank the destination memory is on @@ -4542,7 +4763,17 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (exeInfo.resources.empty()) return false; for (auto const& rss : exeInfo.resources) { Transfer const& t = transfers[rss.transferIdx]; - if (t.srcs.size() > 1 || t.dsts.size() > 1) return false; + vector srcs, dsts; + if (t.numLaps > 0) { + bool const isPong = (rss.numLaps < 0); + int const h = isPong ? 1 : 0; + if (t.srcs[h].memType != MEM_NULL) srcs.push_back(t.srcs[h]); + dsts.push_back(t.dsts[h]); + } else { + srcs = t.srcs; + dsts = t.dsts; + } + if (srcs.size() > 1 || dsts.size() > 1) return false; if (cfg.gfx.useSingleTeam && t.numSubExecs > 1) return false; } return true; @@ -4572,6 +4803,28 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Preparation-related functions //======================================================================================== + // Resolve src/dst MemDevices and subIndex for a TransferResources entry (normal or ping/pong half). + static void ResolveTransferResourceMem(Transfer const& transfer, + TransferResources const& rss, + vector& srcs, + vector& dsts, + int32_t& subIndex) + { + srcs.clear(); + dsts.clear(); + if (transfer.numLaps > 0) { + bool const isPong = (rss.numLaps < 0); + int const halfIdx = isPong ? 1 : 0; + if (transfer.srcs[halfIdx].memType != MEM_NULL) srcs.push_back(transfer.srcs[halfIdx]); + dsts.push_back(transfer.dsts[halfIdx]); + subIndex = isPong ? transfer.exeSubIndexPong : transfer.exeSubIndex; + } else { + srcs = transfer.srcs; + dsts = transfer.dsts; + subIndex = transfer.exeSubIndex; + } + } + // Prepares input parameters for each subexecutor // Determines how sub-executors will split up the work // Initializes counters @@ -4664,6 +4917,78 @@ const auto& AmdSmiFabricInfoV1(const T& info) return ERR_NONE; } + // Returns how many distinct flag slots the lap offsets cycle through, which is also the lap + // distance between successive reuses of any one slot, since + // offset(lap) = (lap * stride) % allocBytes + // repeats with period allocBytes / gcd(stride, allocBytes). + static int PingpongFlagBufferPeriod(int stride, int allocBytes) + { + if (allocBytes <= 0) return 1; + return allocBytes / (int)std::gcd(std::max(0, stride), allocBytes); + } + + // Prepares pingpong parameters for a ping or pong resource half. + // Assumes PrepareExecutor has already allocated rss.srcMem / rss.dstMem for this half. + // Flag cross-linking (localFlagMem) is deferred to PingpongPostPrep. + static ErrResult PreparePingpongParam(ConfigOptions const& cfg, + Transfer const& transfer, + TransferResources& rss) + { + int const initOffset = cfg.data.byteOffset / sizeof(float); + bool const isPong = (rss.numLaps < 0); + int const halfIdx = isPong ? 1 : 0; + + ExeDevice const& exeDevice = isPong ? transfer.exeDevicePong : transfer.exeDevice; + int32_t const subIndex = isPong ? transfer.exeSubIndexPong : transfer.exeSubIndex; + MemDevice const& dstMemDevice = transfer.dsts[halfIdx]; + + PingpongParam& p = rss.pingpongParamCpu; + p.srcMem[0] = nullptr; + p.srcMem[1] = nullptr; + p.localFlagMem = nullptr; + p.flagMem = static_cast(static_cast(rss.dstMem[0])); + p.numLaps = rss.numLaps; + p.flagStride = cfg.general.pingpongStride; + p.flagAllocBytes = cfg.general.pingpongFlagBuffer; + int flagPeriod = PingpongFlagBufferPeriod(p.flagStride, p.flagAllocBytes); + p.hopPeriod = (flagPeriod % 2 == 0) ? flagPeriod : 0; + p.preferredXccId = subIndex; + p.startCycle = 0; + p.stopCycle = 0; + + // Device-resident uint8_t values {0, 1} at rss.srcMem[0] + byteOffset (even/odd lap signaling) + if (exeDevice.exeType == EXE_GPU_GFX && !rss.srcMem.empty()) { + volatile uint8_t* base = static_cast( + static_cast(rss.srcMem[0] + initOffset)); + p.srcMem[0] = base; + p.srcMem[1] = base + 1; + } + + // Override if XCC table has been specified + vector> const& table = cfg.gfx.prefXccTable; + if (exeDevice.exeType == EXE_GPU_GFX && subIndex == -1 && !table.empty() && + IsGpuMemType(dstMemDevice.memType)) { + if (table.size() <= exeDevice.exeIndex || + table[exeDevice.exeIndex].size() <= dstMemDevice.memIndex) { + return {ERR_FATAL, "[gfx.xccPrefTable] is too small"}; + } + p.preferredXccId = table[exeDevice.exeIndex][dstMemDevice.memIndex]; + if (p.preferredXccId < 0 || p.preferredXccId >= GetNumExecutorSubIndices(exeDevice)) { + return {ERR_FATAL, "[gfx.xccPrefTable] defines out-of-bound XCC index %d", p.preferredXccId}; + } + } + + if (System::Get().IsVerbose()) { + System::Get().Log("[INFO] Pingpong flags (%s): %d laps stride %d B buffer %d B " + "%d slot(s) hop %s\n", + isPong ? "pong" : "ping", abs(rss.numLaps), p.flagStride, p.flagAllocBytes, + flagPeriod, p.hopPeriod ? "on" : "off"); + } + + rss.totalDurationMsec = 0.0; + return ERR_NONE; + } + static ErrResult ExchangeMemory(MemDevice const& memDevice, ExeDevice const& exeDevice, size_t* pActualBytes, float** memPtr, hipMemGenericAllocationHandle_t* memHandle) { @@ -4759,9 +5084,13 @@ const auto& AmdSmiFabricInfoV1(const T& info) Transfer const& t = transfers[rss.transferIdx]; rss.numBytes = t.numBytes; + vector srcs, dsts; + int32_t subIndex; + ResolveTransferResourceMem(t, rss, srcs, dsts, subIndex); + if (verbose) { System::Get().Log("[INFO] Rank %d preparing transfer %d (%lu SRC %lu DST) %zu bytes\n", - localRank, rss.transferIdx, t.srcs.size(), t.dsts.size(), t.numBytes); + localRank, rss.transferIdx, srcs.size(), dsts.size(), t.numBytes); System::Get().Log("[INFO] EXE: R%d%c%d NUMA %d%s%s\n", exeDevice.exeRank, ExeTypeStr[exeDevice.exeType], exeDevice.exeIndex, exeNuma, @@ -4770,11 +5099,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) } // Allocate source memory - rss.srcMem.resize(t.srcs.size()); - rss.srcActualBytes.resize(t.srcs.size()); - rss.srcMemHandle.resize(t.srcs.size(), NULL); - for (int iSrc = 0; iSrc < t.srcs.size(); ++iSrc) { - MemDevice const& srcMemDevice = t.srcs[iSrc]; + rss.srcMem.resize(srcs.size()); + rss.srcActualBytes.resize(srcs.size()); + rss.srcMemHandle.resize(srcs.size(), NULL); + for (int iSrc = 0; iSrc < srcs.size(); ++iSrc) { + MemDevice const& srcMemDevice = srcs[iSrc]; + size_t srcAllocBytes = t.numBytes + cfg.data.byteOffset; + if (rss.numLaps != 0) // room for the two 1-byte flag values {0, 1}, rounded up to a float + srcAllocBytes = std::max(srcAllocBytes, cfg.data.byteOffset + sizeof(float)); // Ensure executing GPU can access source memory // This only applies to memory being accessed by a local GPU executor @@ -4799,10 +5131,10 @@ const auto& AmdSmiFabricInfoV1(const T& info) iSrc, GetMemTypeName(srcMemDevice.memType), srcMemDevice.memIndex, GetMemDeviceNuma(srcMemDevice).c_str(), bdf.empty() ? "" : " BDF ", bdf.c_str(), - srcMemDevice.memRank, t.numBytes + cfg.data.byteOffset, + srcMemDevice.memRank, srcAllocBytes, requiresFabricHandle ? " [fabric-exportable]" : ""); } - ERR_CHECK(AllocateMemory(srcMemDevice, t.numBytes + cfg.data.byteOffset, (void**)&rss.srcMem[iSrc], + ERR_CHECK(AllocateMemory(srcMemDevice, srcAllocBytes, (void**)&rss.srcMem[iSrc], &rss.srcActualBytes[iSrc], requiresFabricHandle ? &rss.srcMemHandle[iSrc] : nullptr)); if (verbose) { System::Get().Log("[INFO] SRC[%d]: allocated at %p\n", iSrc, rss.srcMem[iSrc]); @@ -4812,14 +5144,27 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Exchange memory pointer across ranks ERR_CHECK(ExchangeMemory(srcMemDevice, exeDevice, &rss.srcActualBytes[iSrc], &rss.srcMem[iSrc], &rss.srcMemHandle[iSrc])); + + // Pingpong: seed uint8_t lap values {0, 1} into an existing src buffer + // (NULL src skips this and the kernel stores the lap bit directly) + if (rss.numLaps != 0 && rss.srcMem[iSrc] && srcMemDevice.memRank == localRank) { + int const initOffset = cfg.data.byteOffset / sizeof(float); + uint8_t const vals[2] = {0, 1}; + if (IsGpuMemType(srcMemDevice.memType)) { + ERR_CHECK(hipSetDevice(srcMemDevice.memIndex)); + ERR_CHECK(hipMemcpy(rss.srcMem[iSrc] + initOffset, vals, sizeof(vals), hipMemcpyHostToDevice)); + } else if (IsCpuMemType(srcMemDevice.memType)) { + memcpy(static_cast(rss.srcMem[iSrc] + initOffset), vals, sizeof(vals)); + } + } } // Allocate destination memory - rss.dstMem.resize(t.dsts.size()); - rss.dstActualBytes.resize(t.dsts.size()); - rss.dstMemHandle.resize(t.dsts.size(), NULL); - for (int iDst = 0; iDst < t.dsts.size(); ++iDst) { - MemDevice const& dstMemDevice = t.dsts[iDst]; + rss.dstMem.resize(dsts.size()); + rss.dstActualBytes.resize(dsts.size()); + rss.dstMemHandle.resize(dsts.size(), NULL); + for (int iDst = 0; iDst < dsts.size(); ++iDst) { + MemDevice const& dstMemDevice = dsts[iDst]; // Ensure executing GPU can access destination memory if (IsGpuExeType(exeDevice.exeType) && @@ -4835,6 +5180,10 @@ const auto& AmdSmiFabricInfoV1(const T& info) } // Allocate destination memory (on the correct rank) + // For pingpong transfers, allocate a larger buffer to hold multiple flag slots for UALoE station rotation + size_t dstAllocBytes = t.numBytes + cfg.data.byteOffset; + if (rss.numLaps != 0) + dstAllocBytes = std::max(dstAllocBytes, (size_t)cfg.general.pingpongFlagBuffer); bool requiresFabricHandle = (dstMemDevice.memRank != exeDevice.exeRank) && IsGpuExeType(exeDevice.exeType); if (dstMemDevice.memRank == localRank) { if (verbose) { @@ -4846,7 +5195,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) dstMemDevice.memRank, t.numBytes + cfg.data.byteOffset, requiresFabricHandle ? " [fabric-exportable]" : ""); } - ERR_CHECK(AllocateMemory(dstMemDevice, t.numBytes + cfg.data.byteOffset, (void**)&rss.dstMem[iDst], + ERR_CHECK(AllocateMemory(dstMemDevice, dstAllocBytes, (void**)&rss.dstMem[iDst], &rss.dstActualBytes[iDst], requiresFabricHandle ? &rss.dstMemHandle[iDst] : NULL)); if (verbose) { System::Get().Log("[INFO] DST[%d]: allocated at %p\n", iDst, rss.dstMem[iDst]); @@ -4859,7 +5208,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) } // Prepare HSA DMA copy specific resources - if (exeDevice.exeType == EXE_GPU_DMA && (t.exeSubIndex != -1 || cfg.dma.useHsaCopy) && exeDevice.exeRank == localRank) { + if (exeDevice.exeType == EXE_GPU_DMA && (subIndex != -1 || cfg.dma.useHsaCopy) && exeDevice.exeRank == localRank) { #if !defined(__NVCC__) // Collect HSA agent information hsa_amd_pointer_info_t info; @@ -4871,30 +5220,40 @@ const auto& AmdSmiFabricInfoV1(const T& info) rss.dstAgent[dstIdx] = info.agentOwner; } - ERR_CHECK(hsa_amd_pointer_info(rss.srcMem[0], &info, NULL, NULL, NULL)); - rss.srcAgent = info.agentOwner; + if (!rss.srcMem.empty()) { + ERR_CHECK(hsa_amd_pointer_info(rss.srcMem[0], &info, NULL, NULL, NULL)); + rss.srcAgent = info.agentOwner; + } // Create HSA completion signal ERR_CHECK(hsa_signal_create(1, 0, NULL, &rss.signal)); - if (t.exeSubIndex != -1) - rss.sdmaEngineId = (hsa_amd_sdma_engine_id_t)(1U << t.exeSubIndex); + if (subIndex != -1) + rss.sdmaEngineId = (hsa_amd_sdma_engine_id_t)(1U << subIndex); #endif } // Prepare subexecutor parameters (on all ranks) - ERR_CHECK(PrepareSubExecParams(cfg, t, rss)); + if (rss.numLaps == 0) { + ERR_CHECK(PrepareSubExecParams(cfg, t, rss)); + } else { + ERR_CHECK(PreparePingpongParam(cfg, t, rss)); + } } // Prepare additional requirements for GPU-based executors if (IsGpuExeType(exeDevice.exeType) && exeDevice.exeRank == localRank) { ERR_CHECK(hipSetDevice(exeDevice.exeIndex)); - // Determine how many streams to use - int const numStreamsToUse = (exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA || - (cfg.general.useMultiStream && (exeDevice.exeType == EXE_GPU_GFX || - exeDevice.exeType == EXE_GPU_TDM))) - ? exeInfo.resources.size() : 1; + // Determine how many streams to use. + // Pingpong halves do not get their own stream, + // only normal Transfers are counted here, followed by one shared pingpong stream + int const numTransfers = (int)exeInfo.resources.size() - exeInfo.totalPingpong; + bool const multistream = (exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA || + (cfg.general.useMultiStream && (exeDevice.exeType == EXE_GPU_GFX || + exeDevice.exeType == EXE_GPU_TDM))); + int numStreamsToUse = multistream ? numTransfers : (numTransfers > 0 ? 1 : 0); + if (exeInfo.totalPingpong) numStreamsToUse++; exeInfo.streams.resize(numStreamsToUse); // Create streams @@ -4943,12 +5302,20 @@ const auto& AmdSmiFabricInfoV1(const T& info) #else MemType memType = MEM_MANAGED; // NVIDIA hardware requires managed memory to access from host #endif - ERR_CHECK(AllocateMemory({memType, exeDevice.exeIndex}, exeInfo.totalSubExecs * sizeof(SubExecParam), - (void**)&exeInfo.subExecParamGpu)); - ERR_CHECK(GetMemHostAccessibility({memType, exeDevice.exeIndex}, exeInfo.subExecParamHostAccessible)); + if (exeInfo.totalSubExecs > 0) { + ERR_CHECK(AllocateMemory({memType, exeDevice.exeIndex}, exeInfo.totalSubExecs * sizeof(SubExecParam), + (void**)&exeInfo.subExecParamGpu)); + ERR_CHECK(GetMemHostAccessibility({memType, exeDevice.exeIndex}, exeInfo.subExecParamHostAccessible)); + } + if (exeInfo.totalPingpong > 0) { + ERR_CHECK(AllocateMemory({memType, exeDevice.exeIndex}, exeInfo.totalPingpong * sizeof(PingpongParam), + (void**)&exeInfo.pingpongParamGpu)); + ERR_CHECK(GetMemHostAccessibility({memType, exeDevice.exeIndex}, exeInfo.pingpongParamHostAccessible)); + } // Create subexecutor parameter array for entire executor exeInfo.subExecParamCpu.clear(); + exeInfo.pingpongParamCpu.clear(); exeInfo.numSubIndices = GetNumExecutorSubIndices(exeDevice); #if defined(__NVCC__) exeInfo.wallClockRate = 1000000; @@ -4957,9 +5324,11 @@ const auto& AmdSmiFabricInfoV1(const T& info) exeDevice.exeIndex)); #endif int transferOffset = 0; + int pingpongOffset = 0; if (cfg.general.useMultiStream || cfg.gfx.blockOrder == 0) { // Threadblocks are ordered sequentially one transfer at a time for (auto& rss : exeInfo.resources) { + if (rss.numLaps != 0) continue; rss.subExecParamGpuPtr = exeInfo.subExecParamGpu + transferOffset; for (auto p : rss.subExecParamCpu) { rss.subExecIdx.push_back(exeInfo.subExecParamCpu.size()); @@ -4971,6 +5340,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Interleave threadblocks of different Transfers for (int subExecIdx = 0; exeInfo.subExecParamCpu.size() < exeInfo.totalSubExecs; ++subExecIdx) { for (auto& rss : exeInfo.resources) { + if (rss.numLaps != 0) continue; Transfer const& t = transfers[rss.transferIdx]; if (subExecIdx < t.numSubExecs) { rss.subExecIdx.push_back(exeInfo.subExecParamCpu.size()); @@ -4983,6 +5353,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) std::vector> indices; for (int i = 0; i < exeInfo.resources.size(); i++) { auto const& rss = exeInfo.resources[i]; + if (rss.numLaps != 0) continue; Transfer const& t = transfers[rss.transferIdx]; for (int j = 0; j < t.numSubExecs; j++) indices.push_back(std::make_pair(i,j)); @@ -5010,12 +5381,25 @@ const auto& AmdSmiFabricInfoV1(const T& info) exeInfo.subExecParamCpu.data(), exeInfo.subExecParamGpu); } - ERR_CHECK(hipMemcpy(exeInfo.subExecParamGpu, - exeInfo.subExecParamCpu.data(), - exeInfo.totalSubExecs * sizeof(SubExecParam), - hipMemcpyHostToDevice)); - ERR_CHECK(hipDeviceSynchronize()); + if (exeInfo.totalSubExecs > 0) { + ERR_CHECK(hipMemcpy(exeInfo.subExecParamGpu, + exeInfo.subExecParamCpu.data(), + exeInfo.totalSubExecs * sizeof(SubExecParam), + hipMemcpyHostToDevice)); + ERR_CHECK(hipDeviceSynchronize()); + } + + // Pingpong resources are always single stream + for (auto& rss : exeInfo.resources) { + if (rss.numLaps == 0) continue; + rss.pingpongParamIdx = exeInfo.pingpongParamCpu.size(); + rss.pingpongParamGpuPtr = exeInfo.pingpongParamGpu + pingpongOffset; + exeInfo.pingpongParamCpu.push_back(rss.pingpongParamCpu); + pingpongOffset++; + } + // PingpongParam upload deferred to PingpongPostPrep (after flag cross-linking) } + // Prepare for NIC-based executors if (IsNicExeType(exeDevice.exeType)) { @@ -5069,6 +5453,82 @@ const auto& AmdSmiFabricInfoV1(const T& info) return ERR_NONE; } + // Pingpong post-prep: cross-link flag slots, then propagate into PingpongParam. + static ErrResult PingpongPostPrep(ConfigOptions const& cfg, + int const localRank, + vector const& transfers, + vector const& transferResources, + std::map& executorMap) + { + // Cross-link ping and pong flag memory to their partners + std::map pongByTransferIdx; + for (auto* rss : transferResources) { + if (rss->numLaps < 0) pongByTransferIdx[rss->transferIdx] = rss; + } + for (auto* pingRss : transferResources) { + if (pingRss->numLaps <= 0) continue; + + auto it = pongByTransferIdx.find(pingRss->transferIdx); + if (it == pongByTransferIdx.end()) continue; + TransferResources* pongRss = it->second; + + for (auto* rss : {pingRss, pongRss}) { + TransferResources* partnerRss = (rss == pingRss ? pongRss : pingRss); + volatile uint8_t* partnerFlag = static_cast(static_cast( + partnerRss->dstMem[0])); + PingpongParam& pp = rss->pingpongParamCpu; + pp.localFlagMem = partnerFlag; + + // Each half polls the partner half's flag buffer, which PrepareExecutor did not + // cover since the partner's memory belongs to the other half's src/dst list + Transfer const& t = transfers[rss->transferIdx]; + bool const isPong = (rss->numLaps < 0); + MemDevice const& partnerMem = t.dsts[isPong ? 0 : 1]; + ExeDevice exeDevice; + ERR_CHECK(GetActualExecutor(isPong ? t.exeDevicePong : t.exeDevice, exeDevice)); + if (IsGpuExeType(exeDevice.exeType) && IsGpuMemType(partnerMem.memType) && + exeDevice.exeRank == localRank && partnerMem.memRank == localRank && + partnerMem.memIndex != exeDevice.exeIndex) { + if (System::Get().IsVerbose()) { + System::Get().Log("[INFO] Enabling pingpong peer access: GPU %d -> GPU %d\n", + exeDevice.exeIndex, partnerMem.memIndex); + } + ERR_CHECK(EnablePeerAccess(exeDevice.exeIndex, partnerMem.memIndex)); + } + } + } + + // Upload updated PingpongParam to GPU (one bulk copy per executor) + bool const verbose = System::Get().IsVerbose(); + for (auto& exeInfoPair : executorMap) { + ExeDevice const& exeDevice = exeInfoPair.first; + ExeInfo& exeInfo = exeInfoPair.second; + if (exeDevice.exeRank != localRank) continue; + if (!exeInfo.pingpongParamGpu || exeInfo.totalPingpong == 0) continue; + + for (auto& rss : exeInfo.resources) { + if (rss.numLaps == 0 || rss.pingpongParamIdx < 0) continue; + exeInfo.pingpongParamCpu[rss.pingpongParamIdx] = rss.pingpongParamCpu; + } + + ERR_CHECK(hipSetDevice(exeDevice.exeIndex)); + ERR_CHECK(hipMemcpy(exeInfo.pingpongParamGpu, + exeInfo.pingpongParamCpu.data(), + exeInfo.totalPingpong * sizeof(PingpongParam), + hipMemcpyHostToDevice)); + + if (verbose) { + System::Get().Log("[INFO] PingpongParam upload: GPU%d %zu bytes (%zu params) host=%p dev=%p\n", + exeDevice.exeIndex, + exeInfo.totalPingpong * sizeof(PingpongParam), + exeInfo.totalPingpong, + exeInfo.pingpongParamCpu.data(), + exeInfo.pingpongParamGpu); + } + } + return ERR_NONE; + } + // Teardown-related functions //======================================================================================== @@ -5087,26 +5547,31 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Loop over each transfer this executor is involved in for (auto& rss : exeInfo.resources) { Transfer const& t = transfers[rss.transferIdx]; + vector srcs, dsts; + int32_t subIndex; + ResolveTransferResourceMem(t, rss, srcs, dsts, subIndex); if (verbose) { - System::Get().Log("[INFO] Rank %d tearing down transfer %d\n", localRank, rss.transferIdx); + System::Get().Log("[INFO] Rank %d tearing down transfer %d%s\n", + localRank, rss.transferIdx, + rss.numLaps > 0 ? " (ping)" : rss.numLaps < 0 ? " (pong)" : ""); } // Deallocate source memory - for (int iSrc = 0; iSrc < t.srcs.size(); ++iSrc) { - if (t.srcs[iSrc].memRank == localRank) { + for (int iSrc = 0; iSrc < (int)srcs.size(); ++iSrc) { + if (srcs[iSrc].memRank == localRank) { if (verbose) { System::Get().Log("[INFO] Free SRC[%d]: %s idx=%d %p (%zu bytes)\n", - iSrc, GetMemTypeName(t.srcs[iSrc].memType), - t.srcs[iSrc].memIndex, rss.srcMem[iSrc], rss.srcActualBytes[iSrc]); + iSrc, GetMemTypeName(srcs[iSrc].memType), + srcs[iSrc].memIndex, rss.srcMem[iSrc], rss.srcActualBytes[iSrc]); } - ERR_CHECK(DeallocateMemory(t.srcs[iSrc].memType, rss.srcMem[iSrc], + ERR_CHECK(DeallocateMemory(srcs[iSrc].memType, rss.srcMem[iSrc], rss.srcActualBytes[iSrc], &rss.srcMemHandle[iSrc])); } else if (exeDevice.exeRank == localRank && rss.srcMemHandle[iSrc] != 0) { if (verbose) { System::Get().Log("[INFO] Unmap remote SRC[%d]: %p (%zu bytes) from Rank %d\n", - iSrc, rss.srcMem[iSrc], rss.srcActualBytes[iSrc], t.srcs[iSrc].memRank); + iSrc, rss.srcMem[iSrc], rss.srcActualBytes[iSrc], srcs[iSrc].memRank); } #ifdef POD_COMM_ENABLED ERR_CHECK(hipMemUnmap((gpu_device_ptr)rss.srcMem[iSrc], rss.srcActualBytes[iSrc])); @@ -5117,20 +5582,20 @@ const auto& AmdSmiFabricInfoV1(const T& info) } // Deallocate destination memory - for (int iDst = 0; iDst < t.dsts.size(); ++iDst) { - if (t.dsts[iDst].memRank == localRank) { + for (int iDst = 0; iDst < (int)dsts.size(); ++iDst) { + if (dsts[iDst].memRank == localRank) { if (verbose) { System::Get().Log("[INFO] Free DST[%d]: %s idx=%d %p (%zu bytes)\n", - iDst, GetMemTypeName(t.dsts[iDst].memType), - t.dsts[iDst].memIndex, rss.dstMem[iDst], rss.dstActualBytes[iDst]); + iDst, GetMemTypeName(dsts[iDst].memType), + dsts[iDst].memIndex, rss.dstMem[iDst], rss.dstActualBytes[iDst]); } - ERR_CHECK(DeallocateMemory(t.dsts[iDst].memType, rss.dstMem[iDst], + ERR_CHECK(DeallocateMemory(dsts[iDst].memType, rss.dstMem[iDst], rss.dstActualBytes[iDst], &rss.dstMemHandle[iDst])); } else if (exeDevice.exeRank == localRank && rss.dstMemHandle[iDst] != 0) { if (verbose) { System::Get().Log("[INFO] Unmap remote DST[%d]: %p (%zu bytes) from Rank %d\n", - iDst, rss.dstMem[iDst], rss.dstActualBytes[iDst], t.dsts[iDst].memRank); + iDst, rss.dstMem[iDst], rss.dstActualBytes[iDst], dsts[iDst].memRank); } #ifdef POD_COMM_ENABLED ERR_CHECK(hipMemUnmap((gpu_device_ptr)rss.dstMem[iDst], rss.dstActualBytes[iDst])); @@ -5156,7 +5621,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Destroy HSA signal for DMA executor #if !defined(__NVCC__) - if (exeDevice.exeType == EXE_GPU_DMA && (t.exeSubIndex != -1 || cfg.dma.useHsaCopy) && exeDevice.exeRank == localRank) { + if (exeDevice.exeType == EXE_GPU_DMA && (subIndex != -1 || cfg.dma.useHsaCopy) && exeDevice.exeRank == localRank) { ERR_CHECK(hsa_signal_destroy(rss.signal)); } #endif @@ -5184,12 +5649,34 @@ const auto& AmdSmiFabricInfoV1(const T& info) #else MemType memType = MEM_MANAGED; #endif - ERR_CHECK(DeallocateMemory(memType, exeInfo.subExecParamGpu, exeInfo.totalSubExecs * sizeof(SubExecParam))); + if (exeInfo.totalSubExecs > 0) + ERR_CHECK(DeallocateMemory(memType, exeInfo.subExecParamGpu, exeInfo.totalSubExecs * sizeof(SubExecParam))); + if (exeInfo.pingpongParamGpu) + ERR_CHECK(DeallocateMemory(memType, exeInfo.pingpongParamGpu, + exeInfo.totalPingpong * sizeof(PingpongParam))); } return ERR_NONE; } +// PingPong Wait primitives +//======================================================================================== + + // GPU-side spin wait: polls a flag until it equals the expected value. + // Used by GPU-GFX executors (called inline from the transfer kernel). + __device__ void GpuWait(volatile uint8_t* flag, uint8_t val) + { +#if defined(__NVCC__) + // CUDA has no 1-byte atomic, so poll through volatile and fence to order subsequent loads + while (*flag != val) + ; + __threadfence_system(); +#else + while (__hip_atomic_load(flag, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_SYSTEM) != val) + ; +#endif + } + // CPU Executor-related functions //======================================================================================== @@ -5915,6 +6402,81 @@ const auto& AmdSmiFabricInfoV1(const T& info) return (blocksize + 255) / 256 - 1; } + __global__ void GpuPingpongKernel(PingpongParam* params /*int seType*/) + { + int const pingpongIdx = blockIdx.y; + PingpongParam& p = params[pingpongIdx]; + +#if !defined(__NVCC__) + int32_t const xccId = (int32_t)GetXccId(); + if (p.preferredXccId != -1 && xccId != p.preferredXccId) return; +#endif + + if (threadIdx.x != 0) return; + + bool const isPing = p.numLaps > 0; + int const laps = isPing ? p.numLaps : -p.numLaps; + + // Hoist all parameters into registers so that the lap loop only touches the flag slots + int const y = p.flagAllocBytes; + int const sx = y > 0 ? p.flagStride % y : 0; + // hopPeriod 0 disables hopping; a period past the last lap keeps the loop body branch-identical + int const hp = p.hopPeriod > 0 ? p.hopPeriod : laps + 1; + + volatile uint8_t* const localBase = p.localFlagMem; + volatile uint8_t* const remoteBase = p.flagMem; + // Kept as scalars rather than an array so that the lap-parity select stays in registers + uint8_t* const srcVal0 = const_cast(p.srcMem[0]); + uint8_t* const srcVal1 = const_cast(p.srcMem[1]); + bool const useSrcMem = (srcVal0 != nullptr); + + int off = 0; + int hopCnt = hp; + uint8_t val = 0; + + int64_t startCycle = GetTimestamp(); + + for (int lap = 0; lap < laps; lap++) { + volatile uint8_t* localFlag = localBase + off; + volatile uint8_t* remoteFlag = remoteBase + off; + // TODO: replace with hip_atomic_store + if (!useSrcMem) { + if (isPing) { + __atomic_store_n((uint8_t*)remoteFlag, val, __ATOMIC_RELEASE); + GpuWait(localFlag, val); + } else { + GpuWait(localFlag, val); + __atomic_store_n((uint8_t*)remoteFlag, val, __ATOMIC_RELEASE); + } + } else{ + uint8_t* const srcPtr = val ? srcVal1 : srcVal0; + if (isPing) { + __atomic_store((uint8_t*)remoteFlag, srcPtr, __ATOMIC_RELEASE); + GpuWait(localFlag, val); + } else { + GpuWait(localFlag, val); + __atomic_store((uint8_t*)remoteFlag, srcPtr, __ATOMIC_RELEASE); + } + + } + + // Advance one stride, plus an extra stride every hp laps so that a slot is never + // revisited an even number of laps later (which would leave a stale matching value) + off += sx; if (off >= y) off -= y; + if (--hopCnt == 0) { + hopCnt = hp; + off += sx; if (off >= y) off -= y; + } + val ^= 1; + } + + if (isPing) { + __threadfence_system(); + p.stopCycle = GetTimestamp(); + p.startCycle = startCycle; + } + } + // Execute a single GPU Transfer (when using 1 stream per Transfer) static ErrResult ExecuteGpuTransfer(int const iteration, int const exeTotalSubExecs, @@ -5963,6 +6525,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) params, cfg.gfx.seType, cfg.gfx.waveOrder, cfg.general.numSubIterations); #endif + ERR_CHECK(hipGetLastError()); ERR_CHECK(hipStreamSynchronize(stream)); // Record this timing if this Transfer is being run in multistream mode @@ -6001,6 +6564,39 @@ const auto& AmdSmiFabricInfoV1(const T& info) return ERR_NONE; } + // Launch all pingpong halves on the executor's shared pingpong stream (one threadblock per half). + static ErrResult ExecuteGpuPingpong(int const iteration, + int const numPingpong, + PingpongParam* params, + hipStream_t const stream, + hipEvent_t const startEvent, + hipEvent_t const stopEvent, + int const xccDim, + ConfigOptions const& cfg) + { + (void)iteration; + + dim3 const gridSize(xccDim, numPingpong, 1); + dim3 const blockSize(1); + +#if defined(__NVCC__) + if (cfg.general.useHipEvents && startEvent) + ERR_CHECK(hipEventRecord(startEvent, stream)); + GpuPingpongKernel<<>>(params); + if (cfg.general.useHipEvents && stopEvent) + ERR_CHECK(hipEventRecord(stopEvent, stream)); +#else + hipExtLaunchKernelGGL(GpuPingpongKernel, gridSize, blockSize, 0, stream, + cfg.general.useHipEvents ? startEvent : NULL, + cfg.general.useHipEvents ? stopEvent : NULL, 0, + params); +#endif + + ERR_CHECK(hipGetLastError()); + ERR_CHECK(hipStreamSynchronize(stream)); + return ERR_NONE; + } + // Execute a single GPU executor static ErrResult RunGpuExecutor(int const iteration, ConfigOptions const& cfg, @@ -6013,30 +6609,49 @@ const auto& AmdSmiFabricInfoV1(const T& info) int xccDim = exeInfo.useSubIndices ? exeInfo.numSubIndices : 1; if (cfg.general.useMultiStream) { - // Launch one task per Transfer in separate streams on the persistent worker pool - int const numStreams = (int)exeInfo.streams.size(); - std::vector tfrErr(numStreams); - exeInfo.pool->ParallelFor(numStreams, [&](int i) { + std::vector normalIdx; + for (int r = 0; r < (int)exeInfo.resources.size(); r++) + if (exeInfo.resources[r].numLaps == 0) + normalIdx.push_back(r); + + int const numNormal = (int)normalIdx.size(); + std::vector tfrErr(numNormal); + exeInfo.pool->ParallelFor(numNormal, [&](int i) { tfrErr[i] = ExecuteGpuTransfer(iteration, - exeInfo.totalSubExecs, - exeInfo.subExecParamGpu, - exeInfo.streams[i], - cfg.general.useHipEvents ? exeInfo.startEvents[i] : NULL, - cfg.general.useHipEvents ? exeInfo.stopEvents[i] : NULL, - xccDim, - cfg, - exeInfo.gfxKernelToUse, - exeInfo.subExecParamHostAccessible, - exeInfo.resources[i]); + exeInfo.totalSubExecs, + exeInfo.subExecParamGpu, + exeInfo.streams[i], + cfg.general.useHipEvents ? exeInfo.startEvents[i] : NULL, + cfg.general.useHipEvents ? exeInfo.stopEvents[i] : NULL, + xccDim, + cfg, + exeInfo.gfxKernelToUse, + exeInfo.subExecParamHostAccessible, + exeInfo.resources[normalIdx[i]]); }); for (auto& e : tfrErr) ERR_CHECK(e); - } else { - // Launch all Transfers in one kernel launch (avoid extra thread creation) - ExecuteGpuTransfer(iteration, exeInfo.totalSubExecs, exeInfo.subExecParamGpu, exeInfo.streams[0], - cfg.general.useHipEvents ? exeInfo.startEvents[0] : NULL, - cfg.general.useHipEvents ? exeInfo.stopEvents[0] : NULL, - xccDim, cfg, exeInfo.gfxKernelToUse, - exeInfo.subExecParamHostAccessible, exeInfo.resources[0]); + } else if (exeInfo.totalSubExecs > 0) { + TransferResources* normalRss = nullptr; + for (auto& rss : exeInfo.resources) { + if (rss.numLaps == 0) { normalRss = &rss; break; } + } + ERR_CHECK(ExecuteGpuTransfer(iteration, exeInfo.totalSubExecs, exeInfo.subExecParamGpu, exeInfo.streams[0], + cfg.general.useHipEvents ? exeInfo.startEvents[0] : NULL, + cfg.general.useHipEvents ? exeInfo.stopEvents[0] : NULL, + xccDim, cfg, exeInfo.gfxKernelToUse, + exeInfo.subExecParamHostAccessible, *normalRss)); + } + + if (exeInfo.totalPingpong > 0) { + int const ppIdx = (int)exeInfo.streams.size() - 1; + ERR_CHECK(ExecuteGpuPingpong(iteration, + exeInfo.totalPingpong, + exeInfo.pingpongParamGpu, + exeInfo.streams[ppIdx], + cfg.general.useHipEvents ? exeInfo.startEvents[ppIdx] : NULL, + cfg.general.useHipEvents ? exeInfo.stopEvents[ppIdx] : NULL, + xccDim, + cfg)); } auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; @@ -6047,7 +6662,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) // - Otherwise, Use CPU timing if (cfg.general.useHipEvents && !cfg.general.useMultiStream) { float gpuDeltaMsec; - ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, exeInfo.startEvents[0], exeInfo.stopEvents[0])); + if (exeInfo.totalSubExecs > 0) { + ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, exeInfo.startEvents[0], exeInfo.stopEvents[0])); + } else if (exeInfo.totalPingpong > 0) { + int const ppIdx = exeInfo.streams.size() - 1; + ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, exeInfo.startEvents[ppIdx], exeInfo.stopEvents[ppIdx])); + } else { + gpuDeltaMsec = 0.0f; + } gpuDeltaMsec /= cfg.general.numSubIterations; exeInfo.totalDurationMsec += gpuDeltaMsec; } else { @@ -6070,6 +6692,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) for (int i = 0; i < exeInfo.resources.size(); i++) { TransferResources& rss = exeInfo.resources[i]; + if (rss.numLaps != 0) continue; int64_t minStartCycle = std::numeric_limits::max(); int64_t maxStopCycle = std::numeric_limits::min(); std::set> CUs; @@ -6094,6 +6717,31 @@ const auto& AmdSmiFabricInfoV1(const T& info) } } } + + // Pingpong timing is reported per lap, from the in-kernel timestamps that the + // ping half wrote into its PingpongParam (pingpong always uses its own stream) + if (exeInfo.totalPingpong > 0) { + std::vector pingpongParamHost; + PingpongParam const* pingpongParam = exeInfo.pingpongParamGpu; + if (!exeInfo.pingpongParamHostAccessible) { + pingpongParamHost.resize(exeInfo.totalPingpong); + ERR_CHECK(hipMemcpy(pingpongParamHost.data(), exeInfo.pingpongParamGpu, + exeInfo.totalPingpong * sizeof(PingpongParam), hipMemcpyDefault)); + pingpongParam = pingpongParamHost.data(); + } + + for (TransferResources& rss : exeInfo.resources) { + // Only the ping half timestamps the exchange, and it owns the reported row + if (rss.numLaps <= 0 || rss.pingpongParamIdx < 0) continue; + + PingpongParam const& p = pingpongParam[rss.pingpongParamIdx]; + double deltaMsec = (p.stopCycle - p.startCycle) / (double)(exeInfo.wallClockRate); + deltaMsec /= rss.numLaps; + rss.totalDurationMsec += deltaMsec; + if (cfg.general.recordPerIteration) + rss.perIterMsec.push_back(deltaMsec); + } + } } return ERR_NONE; } @@ -6647,15 +7295,39 @@ const auto& AmdSmiFabricInfoV1(const T& info) TransferResources resource = {}; resource.transferIdx = i; + resource.numLaps = t.numLaps; + + bool isPingpong = t.numLaps != 0; ExeInfo& exeInfo = executorMap[exeDevice]; - exeInfo.totalBytes += t.numBytes; - exeInfo.totalSubExecs += t.numSubExecs; + if (!isPingpong) { + // Pingpong exchanges 1-byte flags, not payload, so t.numBytes (a placeholder that + // only sizes the flag allocation) is left out of the reported byte totals + exeInfo.totalBytes += t.numBytes; + exeInfo.totalSubExecs += t.numSubExecs; + } else { + exeInfo.totalPingpong ++; + } exeInfo.useSubIndices |= (t.exeSubIndex != -1 || (t.exeDevice.exeType == EXE_GPU_GFX && !cfg.gfx.prefXccTable.empty())); exeInfo.resources.push_back(resource); minNumSrcs = std::min(minNumSrcs, (int)t.srcs.size()); maxNumSrcs = std::max(maxNumSrcs, (int)t.srcs.size()); maxNumBytes = std::max(maxNumBytes, t.numBytes); + + // proceed to check latter half of a Transfer (pong) + if (isPingpong) { + ExeDevice pongExe; + ERR_APPEND(GetActualExecutor(t.exeDevicePong, pongExe), errResults); + + TransferResources pong = {}; + pong.transferIdx = i; + pong.numLaps = -t.numLaps; + + ExeInfo& pongInfo = executorMap[pongExe]; + pongInfo.totalPingpong ++; + pongInfo.useSubIndices |= (t.exeSubIndexPong != -1 || (t.exeDevicePong.exeType == EXE_GPU_GFX && !cfg.gfx.prefXccTable.empty())); + pongInfo.resources.push_back(pong); + } } // Empty transfer list leaves minNumSrcs at its sentinel (MAX_SRCS + 1); @@ -6687,6 +7359,9 @@ const auto& AmdSmiFabricInfoV1(const T& info) } } + // After all Executors are prepared and subExecParam is set up, we need to link ping and pong flag memory. + ERR_APPEND(PingpongPostPrep(cfg, localRank, transfers, transferResources, executorMap), errResults); + // Prepare reference src/dst arrays - only once for largest size. // dstReference (expected results) is only needed when validation is enabled. bool const validateEnabled = (cfg.data.alwaysValidate >= 0); @@ -6716,6 +7391,8 @@ const auto& AmdSmiFabricInfoV1(const T& info) bool const verbose = System::Get().IsVerbose(); for (auto resource : transferResources) { Transfer const& t = transfers[resource->transferIdx]; + // Ping and Pong will start with src value of 0 + if (t.numLaps != 0) continue; for (int srcIdx = 0; srcIdx < resource->srcMem.size(); srcIdx++) { if (t.srcs[srcIdx].memRank == localRank) { if (IsGpuMemType(t.srcs[srcIdx].memType)) { @@ -6801,37 +7478,50 @@ const auto& AmdSmiFabricInfoV1(const T& info) System::Get().Log("Memory prepared:\n"); for (int i = 0; i < transfers.size(); i++) { - Transfer const& t = transfers[i]; - ExeDevice const& exe = t.exeDevice; - - // Executor info - std::string exeBdf = IsGpuExeType(exe.exeType) ? GetGpuBdf(exe.exeIndex) : ""; - int exeNuma = IsGpuExeType(exe.exeType) - ? System::Get().GetClosestCpuNumaToGpu(exe.exeIndex, exe.exeRank) - : IsNicExeType(exe.exeType) - ? System::Get().GetClosestCpuNumaToNic(exe.exeIndex, exe.exeRank) - : exe.exeIndex; - System::Get().Log("Transfer %03d: EXE=R%d%c%d NUMA=%d%s%s %zu bytes\n", - i, exe.exeRank, ExeTypeStr[exe.exeType], exe.exeIndex, exeNuma, - exeBdf.empty() ? "" : " BDF=", exeBdf.c_str(), t.numBytes); - - for (int iSrc = 0; iSrc < t.srcs.size(); ++iSrc) { - MemDevice const& md = t.srcs[iSrc]; - std::string bdf = IsGpuMemType(md.memType) ? GetGpuBdf(md.memIndex) : ""; - System::Get().Log(" SRC[%d]: %p type=%-18s idx=%d NUMA=%s Rank=%d%s%s\n", - iSrc, transferResources[i]->srcMem[iSrc], - GetMemTypeName(md.memType), md.memIndex, - GetMemDeviceNuma(md).c_str(), md.memRank, - bdf.empty() ? "" : " BDF=", bdf.c_str()); - } - for (int iDst = 0; iDst < t.dsts.size(); ++iDst) { - MemDevice const& md = t.dsts[iDst]; - std::string bdf = IsGpuMemType(md.memType) ? GetGpuBdf(md.memIndex) : ""; - System::Get().Log(" DST[%d]: %p type=%-18s idx=%d NUMA=%s Rank=%d%s%s\n", - iDst, transferResources[i]->dstMem[iDst], - GetMemTypeName(md.memType), md.memIndex, - GetMemDeviceNuma(md).c_str(), md.memRank, - bdf.empty() ? "" : " BDF=", bdf.c_str()); + Transfer const& t = transfers[i]; + + // A pingpong Transfer owns two resources (one per half), each holding only its own + // half's src/dst, so walk the resources and resolve the matching mem lists + for (auto rss : transferResources) { + if (rss->transferIdx != i) continue; + + bool const isPong = (rss->numLaps < 0); + vector srcs, dsts; + int32_t subIndex; + ResolveTransferResourceMem(t, *rss, srcs, dsts, subIndex); + + ExeDevice const& exe = isPong ? t.exeDevicePong : t.exeDevice; + char const* halfStr = rss->numLaps == 0 ? "" : isPong ? " (pong)" : " (ping)"; + + // Executor info + std::string exeBdf = IsGpuExeType(exe.exeType) ? GetGpuBdf(exe.exeIndex) : ""; + int exeNuma = IsGpuExeType(exe.exeType) + ? System::Get().GetClosestCpuNumaToGpu(exe.exeIndex, exe.exeRank) + : IsNicExeType(exe.exeType) + ? System::Get().GetClosestCpuNumaToNic(exe.exeIndex, exe.exeRank) + : exe.exeIndex; + System::Get().Log("Transfer %03d%s: EXE=R%d%c%d NUMA=%d%s%s %zu bytes\n", + i, halfStr, exe.exeRank, ExeTypeStr[exe.exeType], exe.exeIndex, + exeNuma, exeBdf.empty() ? "" : " BDF=", exeBdf.c_str(), t.numBytes); + + for (int iSrc = 0; iSrc < (int)srcs.size(); ++iSrc) { + MemDevice const& md = srcs[iSrc]; + std::string bdf = IsGpuMemType(md.memType) ? GetGpuBdf(md.memIndex) : ""; + System::Get().Log(" SRC[%d]: %p type=%-18s idx=%d NUMA=%s Rank=%d%s%s\n", + iSrc, rss->srcMem[iSrc], + GetMemTypeName(md.memType), md.memIndex, + GetMemDeviceNuma(md).c_str(), md.memRank, + bdf.empty() ? "" : " BDF=", bdf.c_str()); + } + for (int iDst = 0; iDst < (int)dsts.size(); ++iDst) { + MemDevice const& md = dsts[iDst]; + std::string bdf = IsGpuMemType(md.memType) ? GetGpuBdf(md.memIndex) : ""; + System::Get().Log(" DST[%d]: %p type=%-18s idx=%d NUMA=%s Rank=%d%s%s\n", + iDst, rss->dstMem[iDst], + GetMemTypeName(md.memType), md.memIndex, + GetMemDeviceNuma(md).c_str(), md.memRank, + bdf.empty() ? "" : " BDF=", bdf.c_str()); + } } } System::Get().Log("Hit to continue: "); @@ -6869,6 +7559,35 @@ const auto& AmdSmiFabricInfoV1(const T& info) System::Get().Broadcast(0, sizeof(shouldStop), &shouldStop); if (shouldStop) break; + // Reset pingpong flag memory to idle (-1) before each iteration + for (auto* rss : transferResources) { + if (rss->numLaps == 0) continue; + + Transfer const& t = transfers[rss->transferIdx]; + bool const isPong = (rss->numLaps < 0); + MemDevice const& dstMem = t.dsts[isPong ? 1 : 0]; + + // Flag buffers are physically allocated on the owning rank only + if (dstMem.memRank != localRank) continue; + + void* dst = rss->dstMem[0]; + if (!dst) { + // Defensive against future executors + // TODO: improve exit path to prevent hang + System::Get().Log("[ERROR] dstMem[0] is NULL for ping/pong transfer %d on rank %d\n", + rss->transferIdx, localRank); + exit(1); + } + + size_t allocBytes = (size_t)cfg.general.pingpongFlagBuffer; + if (IsCpuMemType(dstMem.memType)) { + memset(dst, -1, allocBytes); + } else if (IsGpuMemType(dstMem.memType)) { + ERR_APPEND(ErrResult(hipSetDevice(dstMem.memIndex)), errResults); + ERR_APPEND(ErrResult(hipMemset(dst, -1, allocBytes)), errResults); + } + } + // Wait for all ranks before starting any timing System::Get().Barrier(); @@ -6951,6 +7670,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) for (auto rss : transferResources) { int transferIdx = rss->transferIdx; Transfer const& t = transfers[transferIdx]; + + // Pingpong halves carry lap flags rather than data (report once, from the ping half) + if (t.numLaps != 0) { + if (rss->numLaps > 0) + System::Get().Log(" Transfer %03d: SKIP(pingpong)\n", transferIdx); + continue; + } + float const* expected = dstReference[t.srcs.size()].data(); bool transferOk = true; bool anyLocalDst = false; @@ -7032,21 +7759,28 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Local executor collects results exeResult.numBytes = exeInfo.totalBytes; exeResult.avgDurationMsec = exeInfo.totalDurationMsec / numTimedIterations; - exeResult.avgBandwidthGbPerSec = (exeResult.numBytes / 1.0e6) / exeResult.avgDurationMsec; + exeResult.avgBandwidthGbPerSec = exeResult.numBytes ? (exeResult.numBytes / 1.0e6) / exeResult.avgDurationMsec + : 0.0; exeResult.sumBandwidthGbPerSec = 0.0; exeResult.transferIdx.clear(); - // Copy over transfer results + // Copy over transfer results (ping half owns pingpong latency) for (auto const& rss : exeInfo.resources) { + if (rss.numLaps < 0) continue; // skip pong half — ping owns the result row int const transferIdx = rss.transferIdx; exeResult.transferIdx.push_back(transferIdx); + // A pingpong half moves a 1-byte flag per lap, so it reports latency only + bool const isPingpong = (rss.numLaps > 0); + size_t const reportedBytes = isPingpong ? 0 : rss.numBytes; + TransferResult& tfrResult = results.tfrResults[transferIdx]; tfrResult.exeDevice = exeDevice; tfrResult.exeDstDevice = {exeDevice.exeType, rss.dstNicIndex}; - tfrResult.numBytes = rss.numBytes; + tfrResult.numBytes = reportedBytes; tfrResult.avgDurationMsec = rss.totalDurationMsec / numTimedIterations; - tfrResult.avgBandwidthGbPerSec = (rss.numBytes / 1.0e6) / tfrResult.avgDurationMsec; + tfrResult.avgBandwidthGbPerSec = reportedBytes ? (reportedBytes / 1.0e6) / tfrResult.avgDurationMsec + : 0.0; if (cfg.general.recordPerIteration) { tfrResult.perIterMsec = rss.perIterMsec; tfrResult.perIterCUs = rss.perIterCUs; @@ -7065,7 +7799,9 @@ const auto& AmdSmiFabricInfoV1(const T& info) results.overheadMsec = std::min(results.overheadMsec, (results.avgTotalDurationMsec - exeResult.avgDurationMsec)); } - results.avgTotalBandwidthGbPerSec = (results.totalBytesTransferred / 1.0e6) / results.avgTotalDurationMsec; + results.avgTotalBandwidthGbPerSec = results.totalBytesTransferred + ? (results.totalBytesTransferred / 1.0e6) / results.avgTotalDurationMsec + : 0.0; // Teardown executors for (auto& exeInfoPair : executorMap) { @@ -7420,7 +8156,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) ErrResult ParseTransfers(std::string line, std::vector& transfers) { - // Replace any round brackets or '->' with spaces, + // Replace round brackets, '->', and ':' with spaces, but preserve '+' for (int i = 1; line[i]; i++) if (line[i] == '(' || line[i] == ')' || line[i] == '-' || line[i] == ':' || line[i] == '>' ) line[i] = ' '; @@ -7482,11 +8218,89 @@ const auto& AmdSmiFabricInfoV1(const T& info) ERR_CHECK(ParseMemType(dstStr, wct.mem[1])); ERR_CHECK(ParseExeType(exeStr, wct.exe)); - // Perform wildcard expansion - int numRanks = GetNumRanks(); - for (int localRankIndex = 0; localRankIndex < numRanks; localRankIndex++) { - bool localRankModified = RecursiveWildcardTransferExpansion(wct, localRankIndex, numBytes, numSubExecs, transfers); - if (!localRankModified) break; + // Check for '+' to detect pingpong; optional lap count immediately follows '+' + // e.g. "+500" or "+" (default numLaps) + std::string nextToken; + auto pos = iss.tellg(); + int numLaps = 1; + bool isPingpong = false; + if (iss >> nextToken && !nextToken.empty() && nextToken[0] == '+') { + isPingpong = true; + if (nextToken.size() > 1) { + numLaps = atoi(nextToken.c_str() + 1); + if (numLaps < 1) + return {ERR_FATAL, "Parsing error: Pingpong %d lap count must be positive (got %d)", i+1, numLaps}; + } + } else { + iss.clear(); + iss.seekg(pos); + } + + if (isPingpong) { + // Parse the pong triplet + std::string pongSrcStr, pongExeStr, pongDstStr; + iss >> pongSrcStr >> pongExeStr >> pongDstStr; + if (iss.fail()) + return {ERR_FATAL, + "Parsing error: Incomplete pong triplet for Pingpong %d", i+1}; + + WildcardTransfer pongWct; + ERR_CHECK(ParseMemType(pongSrcStr, pongWct.mem[0])); + ERR_CHECK(ParseMemType(pongDstStr, pongWct.mem[1])); + ERR_CHECK(ParseExeType(pongExeStr, pongWct.exe)); + + // Temporary transfers to store ping and pong halves + // Expand ping half + std::vector pingTransfers; + int numRanks = GetNumRanks(); + for (int r = 0; r < numRanks; r++) { + if (!RecursiveWildcardTransferExpansion(wct, r, numBytes, numSubExecs, pingTransfers)) + break; + } + + // Expand pong half + std::vector pongTransfers; + for (int r = 0; r < numRanks; r++) { + if (!RecursiveWildcardTransferExpansion(pongWct, r, numBytes, numSubExecs, pongTransfers)) + break; + } + + // Cartesian product: pair every ping with every pong into one Transfer + for (size_t p = 0; p < pingTransfers.size(); p++) { + for (size_t q = 0; q < pongTransfers.size(); q++) { + Transfer const& pingHalf = pingTransfers[p]; + Transfer const& pongHalf = pongTransfers[q]; + + auto singleMemOrNull = [](vector const& mems) { + if (mems.empty()) return MemDevice{MEM_NULL, 0, 0}; + return mems[0]; + }; + + Transfer t; + t.numLaps = numLaps; + // Pingpong exchanges a 1-byte flag per lap; numBytes only has to be a non-zero + // multiple of 4 large enough for the two seed flag values, and is not reported as traffic + t.numBytes = 8; + t.numSubExecs = 1; + t.srcs = {singleMemOrNull(pingHalf.srcs), singleMemOrNull(pongHalf.srcs)}; + t.dsts = {singleMemOrNull(pingHalf.dsts), singleMemOrNull(pongHalf.dsts)}; + t.exeDevice = pingHalf.exeDevice; + t.exeSubIndex = pingHalf.exeSubIndex; + t.exeSubSlot = pingHalf.exeSubSlot; + t.exeDevicePong = pongHalf.exeDevice; + t.exeSubIndexPong = pongHalf.exeSubIndex; + t.exeSubSlotPong = pongHalf.exeSubSlot; + transfers.push_back(t); + } + } + + } else { + // Normal transfer -- expand into transfers (existing behavior) + int numRanks = GetNumRanks(); + for (int localRankIndex = 0; localRankIndex < numRanks; localRankIndex++) { + bool localRankModified = RecursiveWildcardTransferExpansion(wct, localRankIndex, numBytes, numSubExecs, transfers); + if (!localRankModified) break; + } } } @@ -7918,40 +8732,67 @@ const auto& AmdSmiFabricInfoV1(const T& info) { if (!dumpCfgFile || !rankDoesOutput) return; + auto printMem = [&](MemDevice const& m) { + if (m.memType == MEM_NULL) + fprintf(dumpCfgFile, "N"); + else + fprintf(dumpCfgFile, "R%d%c%d", m.memRank, MemTypeStr[m.memType], m.memIndex); + }; + + auto printExe = [&](ExeDevice const& exe, int32_t subIndex, int32_t subSlot) { + fprintf(dumpCfgFile, "R%d%c%d", exe.exeRank, ExeTypeStr[exe.exeType], exe.exeIndex); + if (exe.exeSlot != 0) + fprintf(dumpCfgFile, "%c", 'A' + exe.exeSlot); + if (subIndex != -1) + fprintf(dumpCfgFile, ".%d", subIndex); + if (subSlot != 0) + fprintf(dumpCfgFile, "%c", 'A' + subSlot); + }; + fprintf(dumpCfgFile, "-%lu ", transfers.size()); for (auto const& t : transfers) { fprintf(dumpCfgFile, "("); - // Print SRCs - for (auto const& src : t.srcs) { - fprintf(dumpCfgFile, "R%d%c%d", src.memRank, MemTypeStr[src.memType], src.memIndex); - } - if (t.srcs.empty()) - fprintf(dumpCfgFile, "N"); + if (t.numLaps > 0) { + printMem(t.srcs[0]); + fprintf(dumpCfgFile, "->"); + printExe(t.exeDevice, t.exeSubIndex, t.exeSubSlot); + fprintf(dumpCfgFile, "->"); + printMem(t.dsts[0]); + fprintf(dumpCfgFile, " +"); + if (t.numLaps != 1) + fprintf(dumpCfgFile, "%d", t.numLaps); + fprintf(dumpCfgFile, " "); + printMem(t.srcs[1]); + fprintf(dumpCfgFile, "->"); + printExe(t.exeDevicePong, t.exeSubIndexPong, t.exeSubSlotPong); + fprintf(dumpCfgFile, "->"); + printMem(t.dsts[1]); + fprintf(dumpCfgFile, " %d %lu)", t.numSubExecs, t.numBytes); + } else { + // Print SRCs + for (auto const& src : t.srcs) { + fprintf(dumpCfgFile, "R%d%c%d", src.memRank, MemTypeStr[src.memType], src.memIndex); + } + if (t.srcs.empty()) + fprintf(dumpCfgFile, "N"); - fprintf(dumpCfgFile, "->"); + fprintf(dumpCfgFile, "->"); - // Print Executor - fprintf(dumpCfgFile, "R%d%c%d", t.exeDevice.exeRank, ExeTypeStr[t.exeDevice.exeType], t.exeDevice.exeIndex); - if (t.exeDevice.exeSlot != 0) - fprintf(dumpCfgFile, "%c", 'A' + t.exeDevice.exeSlot); - if (t.exeSubIndex != -1) { - fprintf(dumpCfgFile, ".%d", t.exeSubIndex); - } - if (t.exeSubSlot != 0) { - fprintf(dumpCfgFile, "%c", 'A' + t.exeSubSlot); - } + // Print Executor + printExe(t.exeDevice, t.exeSubIndex, t.exeSubSlot); - fprintf(dumpCfgFile, "->"); + fprintf(dumpCfgFile, "->"); - // Print DSTs - for (auto const& dst : t.dsts) { - fprintf(dumpCfgFile, "R%d%c%d", dst.memRank, MemTypeStr[dst.memType], dst.memIndex); - } - if (t.dsts.empty()) - fprintf(dumpCfgFile, "N"); + // Print DSTs + for (auto const& dst : t.dsts) { + fprintf(dumpCfgFile, "R%d%c%d", dst.memRank, MemTypeStr[dst.memType], dst.memIndex); + } + if (t.dsts.empty()) + fprintf(dumpCfgFile, "N"); - fprintf(dumpCfgFile, " %d %lu)", t.numSubExecs, t.numBytes); + fprintf(dumpCfgFile, " %d %lu)", t.numSubExecs, t.numBytes); + } fflush(dumpCfgFile); } fprintf(dumpCfgFile, "\n"); @@ -9278,3 +10119,4 @@ const auto& AmdSmiFabricInfoV1(const T& info) #undef ERR_CHECK #undef ERR_APPEND } + From 6f775df0470be1601e7d3b5b4042a820e8660d0e Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Tue, 8 Sep 2026 20:30:53 +0000 Subject: [PATCH 05/13] adjusting output format --- src/client/Utilities.hpp | 21 ++++++++++----------- src/header/TransferBench.hpp | 11 +++++------ 2 files changed, 15 insertions(+), 17 deletions(-) diff --git a/src/client/Utilities.hpp b/src/client/Utilities.hpp index 64406454..c1bb222c 100644 --- a/src/client/Utilities.hpp +++ b/src/client/Utilities.hpp @@ -607,16 +607,13 @@ namespace TransferBench::Utils bool isMultiRank = TransferBench::GetNumRanks() > 1; - // The pong half owns no result row, so its executor shows up with no transfers beneath it - std::set pongExeDevices; - for (auto const& t : transfers) - if (t.numLaps > 0) pongExeDevices.insert(t.exeDevicePong); - // Figure out table dimensions int numCols = 5, numRows = 1; size_t numTimedIterations = results.numTimedIterations; for (auto const& exeInfoPair : results.exeResults) { ExeResult const& exeResult = exeInfoPair.second; + // Pong-only executors own no result rows (ping reports the pingpong) + if (exeResult.transferIdx.empty()) continue; int displayCount = 0; for (int idx : exeResult.transferIdx) if (transfers[idx].numLaps >= 0) displayCount++; @@ -654,6 +651,8 @@ namespace TransferBench::Utils ExeType const exeType = exeDevice.exeType; int32_t const exeIndex = exeDevice.exeIndex; + if (exeResult.transferIdx.empty()) continue; // pong-only executor: ping owns the report + // Executors running only pingpong halves move no payload, so bytes/bandwidth are meaningless bool const isPingpongExe = (exeResult.numBytes == 0); @@ -666,8 +665,7 @@ namespace TransferBench::Utils std::string exeSummary; if (isPingpongExe) { - exeSummary = pongExeDevices.count(exeDevice) && exeResult.transferIdx.empty() - ? " pingpong (pong half)" : " pingpong"; + exeSummary = " pingpong"; table.Set(rowIdx, 1, " "); table.Set(rowIdx, 3, " "); } else { @@ -695,7 +693,7 @@ namespace TransferBench::Utils double latencyUs = r.avgDurationMsec * 1000.0; table.Set(rowIdx, 0, "PingPong %-4d ", idx); table.Set(rowIdx, 1, "%8.3f us " , latencyUs); - table.Set(rowIdx, 2, "%8.3f ms " , r.avgDurationMsec); + table.Set(rowIdx, 2, "%8.3f ms " , r.avgDurationMsec * t.numLaps); table.Set(rowIdx, 3, "%8d laps " , t.numLaps); if (isMultiRank) { @@ -732,13 +730,13 @@ namespace TransferBench::Utils double iterUs = time.first * 1000.0; table.Set(rowIdx, 0, "Iter %03d ", time.second); table.Set(rowIdx, 1, "%8.3f us ", iterUs); - table.Set(rowIdx, 2, "%8.3f ms ", time.first); + table.Set(rowIdx, 2, "%8.3f ms ", time.first * t.numLaps); rowIdx++; } table.Set(rowIdx, 0, "StandardDev "); table.Set(rowIdx, 1, "%8.3f us ", stdDevTime * 1000.0); - table.Set(rowIdx, 2, "%8.3f ms ", stdDevTime); + table.Set(rowIdx, 2, "%8.3f ms ", stdDevTime * t.numLaps); rowIdx++; table.DrawRowBorder(rowIdx); } @@ -833,11 +831,12 @@ namespace TransferBench::Utils table.Set(rowIdx, 0, "p%d ", pct); if (t.numLaps > 0) { table.Set(rowIdx, 1, "%8.3f us ", dur * 1000.0); + table.Set(rowIdx, 2, "%8.3f ms ", dur * t.numLaps); } else { double bwGbs = dur > 0.0 ? (t.numBytes / 1.0E9) / dur * 1000.0 : 0.0; table.Set(rowIdx, 1, "%8.3f GB/s ", bwGbs); + table.Set(rowIdx, 2, "%8.3f ms ", dur); } - table.Set(rowIdx, 2, "%8.3f ms ", dur); table.Set(rowIdx, 3, " "); table.Set(rowIdx, 4, " "); table.SetCellAlignment(rowIdx, 4, TableHelper::ALIGN_LEFT); diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 48b46960..915fce3d 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -6658,15 +6658,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (iteration >= 0) { // Determine executor timing - // - Use HIP event timing if enabled and not using multi-stream - // - Otherwise, Use CPU timing - if (cfg.general.useHipEvents && !cfg.general.useMultiStream) { + // - HIP events cover only the combined copy launch on stream 0, so they miss a + // later pingpong stream. If this GPU has pingpong (alone or mixed with copies), + // use the CPU clock around this function, same as multi-stream. + // - Otherwise HIP events when enabled and not multi-stream. + if (cfg.general.useHipEvents && !cfg.general.useMultiStream && exeInfo.totalPingpong == 0) { float gpuDeltaMsec; if (exeInfo.totalSubExecs > 0) { ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, exeInfo.startEvents[0], exeInfo.stopEvents[0])); - } else if (exeInfo.totalPingpong > 0) { - int const ppIdx = exeInfo.streams.size() - 1; - ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, exeInfo.startEvents[ppIdx], exeInfo.stopEvents[ppIdx])); } else { gpuDeltaMsec = 0.0f; } From 93227a0e8dbc0c8f6fe7ad71f49117e8a28101f4 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Wed, 9 Sep 2026 16:04:47 -0500 Subject: [PATCH 06/13] relaxed atomics for pingpong kernel --- src/header/TransferBench.hpp | 38 +++++++++++++++++------------------- 1 file changed, 18 insertions(+), 20 deletions(-) diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 915fce3d..3d1e5b14 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -5672,11 +5672,20 @@ const auto& AmdSmiFabricInfoV1(const T& info) ; __threadfence_system(); #else - while (__hip_atomic_load(flag, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_SYSTEM) != val) + while (__hip_atomic_load(flag, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM) != val) ; #endif } + __device__ void GpuStore(volatile uint8_t* flag, uint8_t val) + { +#if defined(__NVCC__) + __atomic_store_n((uint8_t*)flag, val, __ATOMIC_RELAXED); +#else + __hip_atomic_store((uint8_t*)flag, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); +#endif + } + // CPU Executor-related functions //======================================================================================== @@ -6439,25 +6448,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) for (int lap = 0; lap < laps; lap++) { volatile uint8_t* localFlag = localBase + off; volatile uint8_t* remoteFlag = remoteBase + off; - // TODO: replace with hip_atomic_store - if (!useSrcMem) { - if (isPing) { - __atomic_store_n((uint8_t*)remoteFlag, val, __ATOMIC_RELEASE); - GpuWait(localFlag, val); - } else { - GpuWait(localFlag, val); - __atomic_store_n((uint8_t*)remoteFlag, val, __ATOMIC_RELEASE); - } - } else{ - uint8_t* const srcPtr = val ? srcVal1 : srcVal0; - if (isPing) { - __atomic_store((uint8_t*)remoteFlag, srcPtr, __ATOMIC_RELEASE); - GpuWait(localFlag, val); - } else { - GpuWait(localFlag, val); - __atomic_store((uint8_t*)remoteFlag, srcPtr, __ATOMIC_RELEASE); - } - + if (isPing) { + uint8_t const storeVal = useSrcMem ? (val ? *srcVal1 : *srcVal0) : val; + GpuStore(remoteFlag, storeVal); + GpuWait(localFlag, val); + } else { + GpuWait(localFlag, val); + uint8_t const storeVal = useSrcMem ? (val ? *srcVal1 : *srcVal0) : val; + GpuStore(remoteFlag, storeVal); } // Advance one stride, plus an extra stride every hp laps so that a slot is never From 6954cdd45d3d79a1e6e5c4c3bee6185196ed0e85 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Wed, 23 Sep 2026 01:00:20 -0500 Subject: [PATCH 07/13] refactor GFX executor timing --- src/header/TransferBench.hpp | 135 ++++++++++++++++++++--------------- 1 file changed, 78 insertions(+), 57 deletions(-) diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 3d1e5b14..f4b68502 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -3445,6 +3445,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) vector subExecParamCpu; ///< Subexecutor parameters for this executor vector pingpongParamCpu; ///< Pingpong parameters for this executor vector resources; ///< Per-transfer resources (normal and ping/pong halves) + vector transferRssIdx; ///< Indices into resources of the normal Transfers // For GPU-Executors SubExecParam* subExecParamGpu; ///< GPU copy of subExecutor parameters @@ -5067,6 +5068,13 @@ const auto& AmdSmiFabricInfoV1(const T& info) exeInfo.totalDurationMsec = 0.0; int const localRank = GetRank(); bool const verbose = System::Get().IsVerbose(); + + // Pick out the normal Transfers among resources + exeInfo.transferRssIdx.clear(); + for (int r = 0; r < exeInfo.resources.size(); r++) + if (exeInfo.resources[r].numLaps == 0) + exeInfo.transferRssIdx.push_back(r); + // Executor BDF and NUMA are constant across all transfers — compute once std::string exeBdf = IsGpuExeType(exeDevice.exeType) ? GetGpuBdf(exeDevice.exeIndex) : ""; int exeNuma = IsGpuExeType(exeDevice.exeType) @@ -6601,78 +6609,91 @@ const auto& AmdSmiFabricInfoV1(const T& info) int const exeIndex, ExeInfo& exeInfo) { - auto cpuStart = std::chrono::high_resolution_clock::now(); + using Clock = std::chrono::high_resolution_clock; + + auto cpuStart = Clock::now(); ERR_CHECK(hipSetDevice(exeIndex)); int xccDim = exeInfo.useSubIndices ? exeInfo.numSubIndices : 1; - if (cfg.general.useMultiStream) { - std::vector normalIdx; - for (int r = 0; r < (int)exeInfo.resources.size(); r++) - if (exeInfo.resources[r].numLaps == 0) - normalIdx.push_back(r); - - int const numNormal = (int)normalIdx.size(); - std::vector tfrErr(numNormal); - exeInfo.pool->ParallelFor(numNormal, [&](int i) { - tfrErr[i] = ExecuteGpuTransfer(iteration, + std::vector const& normalIdx = exeInfo.transferRssIdx; + + // Normal Transfers run as one task per stream in multi-stream mode, otherwise all of + // them are combined into a single launch on stream 0 + bool const hasPingpong = (exeInfo.totalPingpong > 0); + int const numCopyTasks = cfg.general.useMultiStream ? (int)normalIdx.size() + : ((exeInfo.totalSubExecs > 0 && !normalIdx.empty()) ? 1 : 0); + int const numTasks = numCopyTasks + (hasPingpong ? 1 : 0); + + // Pingpong shares the pool with the copies so that both run concurrently on their own + // streams (a second ParallelFor would serialize behind this one). It is dispatched + // first since it is typically the longer running of the two. + int const ppTask = hasPingpong ? 0 : -1; + int const copyTaskBase = hasPingpong ? 1 : 0; + + std::vector taskErr(numTasks); + std::vector copyStarts(numCopyTasks), copyStops(numCopyTasks); + + exeInfo.pool->ParallelFor(numTasks, [&](int i) { + if (i == ppTask) { + int const ppIdx = (int)exeInfo.streams.size() - 1; + taskErr[i] = ExecuteGpuPingpong(iteration, + exeInfo.totalPingpong, + exeInfo.pingpongParamGpu, + exeInfo.streams[ppIdx], + cfg.general.useHipEvents ? exeInfo.startEvents[ppIdx] : NULL, + cfg.general.useHipEvents ? exeInfo.stopEvents[ppIdx] : NULL, + xccDim, + cfg); + return; + } + + // Each copy task brackets itself so that executor timing can be derived from the + // copies alone, no matter how long the pingpong exchange keeps running + int const c = i - copyTaskBase; + int const streamIdx = cfg.general.useMultiStream ? c : 0; + TransferResources& rss = exeInfo.resources[cfg.general.useMultiStream ? normalIdx[c] : normalIdx[0]]; + + copyStarts[c] = Clock::now(); + taskErr[i] = ExecuteGpuTransfer(iteration, exeInfo.totalSubExecs, exeInfo.subExecParamGpu, - exeInfo.streams[i], - cfg.general.useHipEvents ? exeInfo.startEvents[i] : NULL, - cfg.general.useHipEvents ? exeInfo.stopEvents[i] : NULL, + exeInfo.streams[streamIdx], + cfg.general.useHipEvents ? exeInfo.startEvents[streamIdx] : NULL, + cfg.general.useHipEvents ? exeInfo.stopEvents[streamIdx] : NULL, xccDim, cfg, exeInfo.gfxKernelToUse, exeInfo.subExecParamHostAccessible, - exeInfo.resources[normalIdx[i]]); - }); - for (auto& e : tfrErr) ERR_CHECK(e); - } else if (exeInfo.totalSubExecs > 0) { - TransferResources* normalRss = nullptr; - for (auto& rss : exeInfo.resources) { - if (rss.numLaps == 0) { normalRss = &rss; break; } - } - ERR_CHECK(ExecuteGpuTransfer(iteration, exeInfo.totalSubExecs, exeInfo.subExecParamGpu, exeInfo.streams[0], - cfg.general.useHipEvents ? exeInfo.startEvents[0] : NULL, - cfg.general.useHipEvents ? exeInfo.stopEvents[0] : NULL, - xccDim, cfg, exeInfo.gfxKernelToUse, - exeInfo.subExecParamHostAccessible, *normalRss)); - } - - if (exeInfo.totalPingpong > 0) { - int const ppIdx = (int)exeInfo.streams.size() - 1; - ERR_CHECK(ExecuteGpuPingpong(iteration, - exeInfo.totalPingpong, - exeInfo.pingpongParamGpu, - exeInfo.streams[ppIdx], - cfg.general.useHipEvents ? exeInfo.startEvents[ppIdx] : NULL, - cfg.general.useHipEvents ? exeInfo.stopEvents[ppIdx] : NULL, - xccDim, - cfg)); - } + rss); + copyStops[c] = Clock::now(); + }); + for (auto& e : taskErr) ERR_CHECK(e); - auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; + auto cpuDelta = Clock::now() - cpuStart; if (iteration >= 0) { - // Determine executor timing - // - HIP events cover only the combined copy launch on stream 0, so they miss a - // later pingpong stream. If this GPU has pingpong (alone or mixed with copies), - // use the CPU clock around this function, same as multi-stream. - // - Otherwise HIP events when enabled and not multi-stream. - if (cfg.general.useHipEvents && !cfg.general.useMultiStream && exeInfo.totalPingpong == 0) { + // Determine executor timing. Pingpong moves flags rather than payload and does not + // contribute to exeInfo.totalBytes, so it is kept out of the executor duration that + // backs the reported bandwidth - otherwise a lap count that outlasts the copies would + // stretch the denominator. + // - HIP events on stream 0 cover exactly the combined copy launch (non multi-stream) + // - Otherwise the span from the first copy task starting to the last one finishing + // - Pingpong-only executors transfer no bytes, so report the exchange's wall time + double const scale = 1000.0 / cfg.general.numSubIterations; + if (numCopyTasks == 0) { + exeInfo.totalDurationMsec += hasPingpong + ? std::chrono::duration_cast>(cpuDelta).count() * scale + : 0.0; + } else if (cfg.general.useHipEvents && !cfg.general.useMultiStream) { float gpuDeltaMsec; - if (exeInfo.totalSubExecs > 0) { - ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, exeInfo.startEvents[0], exeInfo.stopEvents[0])); - } else { - gpuDeltaMsec = 0.0f; - } - gpuDeltaMsec /= cfg.general.numSubIterations; - exeInfo.totalDurationMsec += gpuDeltaMsec; + ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, exeInfo.startEvents[0], exeInfo.stopEvents[0])); + exeInfo.totalDurationMsec += gpuDeltaMsec / cfg.general.numSubIterations; } else { - double cpuDeltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 - / cfg.general.numSubIterations; - exeInfo.totalDurationMsec += cpuDeltaMsec; + auto const minStart = *std::min_element(copyStarts.begin(), copyStarts.end()); + auto const maxStop = *std::max_element(copyStops.begin(), copyStops.end()); + exeInfo.totalDurationMsec += std::chrono::duration_cast>(maxStop - minStart).count() + * scale; } // If Transfers were combined into a single launch, figure out per-Transfer timing From 81865faadaf5f33594cdecacc1eeee4779c8e87b Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Wed, 23 Sep 2026 10:14:13 -0500 Subject: [PATCH 08/13] cuda fix and cross pod fix --- src/header/TransferBench.hpp | 104 ++++++++++++++++++++++++++++++----- 1 file changed, 91 insertions(+), 13 deletions(-) diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index f4b68502..523c2809 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -2480,10 +2480,13 @@ const auto& AmdSmiFabricInfoV1(const T& info) { if (t.numLaps > 0) { for (int half = 0; half < 2; half++) { - ExeDevice const& exe = half == 0 ? t.exeDevice : t.exeDevicePong; + ExeDevice const& exe = !half ? t.exeDevice : t.exeDevicePong; if (t.srcs[half].memType != MEM_NULL && t.srcs[half].memRank != exe.exeRank) return true; if (t.dsts[half].memType != MEM_NULL && t.dsts[half].memRank != exe.exeRank) return true; } + // Partner flag buffers are polled by the opposite half's executor + if (t.dsts[0].memType != MEM_NULL && t.dsts[0].memRank != t.exeDevicePong.exeRank) return true; + if (t.dsts[1].memType != MEM_NULL && t.dsts[1].memRank != t.exeDevice.exeRank) return true; return false; } @@ -2995,6 +2998,11 @@ const auto& AmdSmiFabricInfoV1(const T& info) !(samePod = IsSamePod(t.dsts[half].memRank, exeRank))) break; } + // Partner flag buffers must be reachable by the opposite half's executor + if (samePod && t.dsts[0].memType != MEM_NULL) + samePod = IsSamePod(t.dsts[0].memRank, t.exeDevicePong.exeRank); + if (samePod && t.dsts[1].memType != MEM_NULL) + samePod = IsSamePod(t.dsts[1].memRank, t.exeDevice.exeRank); } else { int exeRank = t.exeDevice.exeRank; for (auto const& src : t.srcs) { @@ -3010,8 +3018,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) } if (!samePod || IsCpuExeType(t.exeDevice.exeType)) { - errors.push_back({ERR_FATAL, "Transfer %d: Executor on rank %d can not access memory across ranks\n", - i, t.exeDevice.exeRank}); + if (isPingpong) { + errors.push_back({ERR_FATAL, + "Transfer %d: Ping/Pong Executor on rank %d/%d cannot access memory (src or partner flag) across ranks", + i, t.exeDevice.exeRank, t.exeDevicePong.exeRank}); + } else { + errors.push_back({ERR_FATAL, "Transfer %d: Executor on rank %d cannot access memory across ranks\n", + i, t.exeDevice.exeRank}); + } break; } @@ -3045,6 +3059,24 @@ const auto& AmdSmiFabricInfoV1(const T& info) break; } } + if (!hasFatalError && t.dsts[0].memType != MEM_NULL && + t.dsts[0].memRank != t.exeDevicePong.exeRank && IsCpuMemType(t.dsts[0].memType)) { + errors.push_back({ERR_FATAL, + "Transfer %d: Cross-rank GPU executor (R%d%c%d) cannot access remote host memory " + "(ping DST on rank %d is %s) for partner flag polling.", + i, t.exeDevicePong.exeRank, ExeTypeStr[t.exeDevicePong.exeType], t.exeDevicePong.exeIndex, + t.dsts[0].memRank, GetMemTypeName(t.dsts[0].memType)}); + hasFatalError = true; + } + if (!hasFatalError && t.dsts[1].memType != MEM_NULL && + t.dsts[1].memRank != t.exeDevice.exeRank && IsCpuMemType(t.dsts[1].memType)) { + errors.push_back({ERR_FATAL, + "Transfer %d: Cross-rank GPU executor (R%d%c%d) cannot access remote host memory " + "(pong DST on rank %d is %s) for partner flag polling.", + i, t.exeDevice.exeRank, ExeTypeStr[t.exeDevice.exeType], t.exeDevice.exeIndex, + t.dsts[1].memRank, GetMemTypeName(t.dsts[1].memType)}); + hasFatalError = true; + } if (hasFatalError) break; } else if (IsGpuExeType(t.exeDevice.exeType)) { bool hasRemoteCpuMem = false; @@ -3327,6 +3359,9 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Pingpong role/lap count on this resource half (0 = normal, >0 = ping, <0 = pong) int numLaps = 0; + float* partnerFlagMem = nullptr; ///< Partner dst imported onto this executor (cross-rank/pod) + size_t partnerFlagActualBytes = 0; ///< Size of partnerFlagMem mapping + memHandle_t partnerFlagMemHandle = 0; ///< Import handle for partnerFlagMem teardown // Counters double totalDurationMsec; ///< Total duration for all iterations for this Transfer @@ -5193,6 +5228,12 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (rss.numLaps != 0) dstAllocBytes = std::max(dstAllocBytes, (size_t)cfg.general.pingpongFlagBuffer); bool requiresFabricHandle = (dstMemDevice.memRank != exeDevice.exeRank) && IsGpuExeType(exeDevice.exeType); + if (rss.numLaps != 0 && IsGpuMemType(dstMemDevice.memType)) { + // Partner half's executor may fabric-import this flag buffer for cross-rank pingpong + ExeDevice const& partnerExe = (rss.numLaps > 0) ? t.exeDevicePong : t.exeDevice; + if (IsGpuExeType(partnerExe.exeType) && partnerExe.exeRank != dstMemDevice.memRank) + requiresFabricHandle = true; + } if (dstMemDevice.memRank == localRank) { if (verbose) { std::string bdf = IsGpuMemType(dstMemDevice.memType) ? GetGpuBdf(dstMemDevice.memIndex) : ""; @@ -5494,15 +5535,33 @@ const auto& AmdSmiFabricInfoV1(const T& info) MemDevice const& partnerMem = t.dsts[isPong ? 0 : 1]; ExeDevice exeDevice; ERR_CHECK(GetActualExecutor(isPong ? t.exeDevicePong : t.exeDevice, exeDevice)); - if (IsGpuExeType(exeDevice.exeType) && IsGpuMemType(partnerMem.memType) && - exeDevice.exeRank == localRank && partnerMem.memRank == localRank && - partnerMem.memIndex != exeDevice.exeIndex) { - if (System::Get().IsVerbose()) { - System::Get().Log("[INFO] Enabling pingpong peer access: GPU %d -> GPU %d\n", - exeDevice.exeIndex, partnerMem.memIndex); + + float* partnerFlagPtr = partnerRss->dstMem[0]; + if (IsGpuExeType(exeDevice.exeType) && IsGpuMemType(partnerMem.memType)) { + if (partnerMem.memRank != exeDevice.exeRank) { + // Partner dst must be fabric-imported onto this half's executor (pingDst<->pongExe, etc.) + rss->partnerFlagActualBytes = partnerRss->dstActualBytes.empty() ? 0 : partnerRss->dstActualBytes[0]; + memHandle_t* exchangeHandle = &rss->partnerFlagMemHandle; + if (localRank == partnerMem.memRank) + exchangeHandle = &partnerRss->dstMemHandle[0]; + if (System::Get().IsVerbose()) { + System::Get().Log("[INFO] Exchanging pingpong partner flag: %s rank %d -> executor R%d%c%d\n", + isPong ? "ping DST" : "pong DST", partnerMem.memRank, + exeDevice.exeRank, ExeTypeStr[exeDevice.exeType], exeDevice.exeIndex); + } + ERR_CHECK(ExchangeMemory(partnerMem, exeDevice, &rss->partnerFlagActualBytes, + &partnerFlagPtr, exchangeHandle)); + rss->partnerFlagMem = partnerFlagPtr; + } else if (partnerMem.memIndex != exeDevice.exeIndex) { + if (System::Get().IsVerbose()) { + System::Get().Log("[INFO] Enabling pingpong peer access: GPU %d -> GPU %d\n", + exeDevice.exeIndex, partnerMem.memIndex); + } + ERR_CHECK(EnablePeerAccess(exeDevice.exeIndex, partnerMem.memIndex)); } - ERR_CHECK(EnablePeerAccess(exeDevice.exeIndex, partnerMem.memIndex)); } + + pp.localFlagMem = static_cast(static_cast(partnerFlagPtr)); } } @@ -5638,6 +5697,22 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (IsIbvSymbolsReady() && IsNicExeType(exeDevice.exeType)) { ERR_CHECK(TeardownNicTransferResources(rss, t)); } + + // Unmap fabric-imported partner flag buffer (cross-rank pingpong) + if (rss.numLaps != 0 && exeDevice.exeRank == localRank && rss.partnerFlagMemHandle) { + if (verbose) { + System::Get().Log("[INFO] Unmap pingpong partner flag: %p (%zu bytes)\n", + rss.partnerFlagMem, rss.partnerFlagActualBytes); + } +#ifdef POD_COMM_ENABLED + ERR_CHECK(hipMemUnmap((gpu_device_ptr)rss.partnerFlagMem, rss.partnerFlagActualBytes)); + ERR_CHECK(hipMemRelease(rss.partnerFlagMemHandle)); + ERR_CHECK(hipMemAddressFree((gpu_device_ptr)rss.partnerFlagMem, rss.partnerFlagActualBytes)); +#endif + rss.partnerFlagMem = nullptr; + rss.partnerFlagMemHandle = 0; + rss.partnerFlagActualBytes = 0; + } } // Teardown additional requirements for GPU-based executors @@ -5675,10 +5750,9 @@ const auto& AmdSmiFabricInfoV1(const T& info) __device__ void GpuWait(volatile uint8_t* flag, uint8_t val) { #if defined(__NVCC__) - // CUDA has no 1-byte atomic, so poll through volatile and fence to order subsequent loads + // ld.volatile has the same semantics as ld.relaxed.sys on sm_70+, matching the HIP path below while (*flag != val) ; - __threadfence_system(); #else while (__hip_atomic_load(flag, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM) != val) ; @@ -5688,7 +5762,8 @@ const auto& AmdSmiFabricInfoV1(const T& info) __device__ void GpuStore(volatile uint8_t* flag, uint8_t val) { #if defined(__NVCC__) - __atomic_store_n((uint8_t*)flag, val, __ATOMIC_RELAXED); + // st.volatile has the same semantics as st.relaxed.sys on sm_70+ + *flag = val; #else __hip_atomic_store((uint8_t*)flag, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); #endif @@ -7603,6 +7678,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) } else if (IsGpuMemType(dstMem.memType)) { ERR_APPEND(ErrResult(hipSetDevice(dstMem.memIndex)), errResults); ERR_APPEND(ErrResult(hipMemset(dst, -1, allocBytes)), errResults); + ERR_APPEND(ErrResult(hipDeviceSynchronize()), errResults); } } @@ -7636,6 +7712,8 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Clear destination memory after validation so each iteration starts from a known-zero state size_t const initOffset = cfg.data.byteOffset / sizeof(float); for (auto rss : transferResources) { + // Pingpong flag buffers are reset at the start of every iteration instead + if (rss->numLaps != 0) continue; Transfer const& t = transfers[rss->transferIdx]; for (int dstIdx = 0; dstIdx < (int)rss->dstMem.size(); dstIdx++) { if (t.dsts[dstIdx].memRank != localRank) continue; From e5166b63899605d819876b679ddcdb81c8fd761d Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Wed, 23 Sep 2026 16:05:11 -0500 Subject: [PATCH 09/13] Latency Presets --- src/client/Presets/Latency.hpp | 234 +++++++++++++++++++++++++++++++++ src/client/Presets/Presets.hpp | 58 ++++---- 2 files changed, 265 insertions(+), 27 deletions(-) create mode 100644 src/client/Presets/Latency.hpp diff --git a/src/client/Presets/Latency.hpp b/src/client/Presets/Latency.hpp new file mode 100644 index 00000000..564415e2 --- /dev/null +++ b/src/client/Presets/Latency.hpp @@ -0,0 +1,234 @@ +/* +Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + +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. +*/ + +// Pingpong latency presets, which differ only in which pairs of GPU executors run together: +// p2p_latency - Every pair runs by itself, one pair at a time +// one2all_latency - One PING executor runs against all PONG executors at once, one PING executor at a time +// a2a_latency - Every pair runs at once +int LatencyPreset(EnvVars& ev, + size_t const /*numBytesPerTransfer*/, + std::string const presetName, + [[maybe_unused]] bool const bytesSpecified) +{ + bool const isP2p = (presetName == "p2p_latency"); + bool const isOne2All = (presetName == "one2all_latency"); + bool const isA2a = (presetName == "a2a_latency"); + + if (!Utils::AllRanksHaveSameGpuCount()) { + Utils::Print("[ERROR] %s preset requires all ranks to have the same number of GPUs\n", presetName.c_str()); + Utils::Print("[ERROR] Run ./TransferBench without any args to display topology information\n"); + return ERR_FATAL; + } + + int const numRanks = TransferBench::GetNumRanks(); + int const numDetectedGpus = TransferBench::GetNumExecutors(EXE_GPU_GFX); + + // Collect env vars for this preset + int a2aLocal = isA2a ? EnvVars::GetEnvVar("A2A_LOCAL", 0) : 0; + int memTypeIdx = EnvVars::GetEnvVar("GPU_MEM_TYPE" , 0); + int numGpuDevices = EnvVars::GetEnvVar("NUM_GPU_DEVICES", numDetectedGpus); + int numLaps = EnvVars::GetEnvVar("NUM_LAPS" , 1000); + int useRemoteRead = EnvVars::GetEnvVar("USE_REMOTE_READ", 0); + + MemType const memType = Utils::GetGpuMemType(memTypeIdx); + std::string const memTypeStr = Utils::GetGpuMemTypeStr(memTypeIdx); + + // Display environment variables + if (Utils::RankDoesOutput()) { + ev.DisplayEnvVars(); + if (!ev.hideEnv) { + if (!ev.outputToCsv) printf("[Latency Related]\n"); + if (isA2a) + ev.Print("A2A_LOCAL", a2aLocal, "%s pairs with PING and PONG on the same GPU", a2aLocal ? "Include" : "Exclude"); + ev.Print("GPU_MEM_TYPE", memTypeIdx, "Using %s memory for flags (%s)", memTypeStr.c_str(), Utils::GetAllGpuMemTypeStr().c_str()); + ev.Print("NUM_GPU_DEVICES", numGpuDevices, "Using %d GPUs%s", numGpuDevices, numRanks > 1 ? " per rank" : ""); + ev.Print("NUM_LAPS", numLaps, "Timing %d round trips per iteration", numLaps); + ev.Print("USE_REMOTE_READ", useRemoteRead, "%s", useRemoteRead ? "Executors write to their own memory and poll their partner's" + : "Executors write to their partner's memory and poll their own"); + printf("\n"); + } + } + + // Check that input parameters are uniform across all ranks + IS_UNIFORM(a2aLocal, "A2A_LOCAL"); + IS_UNIFORM(memTypeIdx, "GPU_MEM_TYPE"); + IS_UNIFORM(numGpuDevices, "NUM_GPU_DEVICES"); + IS_UNIFORM(numLaps, "NUM_LAPS"); + IS_UNIFORM(useRemoteRead, "USE_REMOTE_READ"); + + // Validate env vars + if (numGpuDevices < 1 || numGpuDevices > numDetectedGpus) { + Utils::Print("[ERROR] NUM_GPU_DEVICES must be between 1 and %d (got %d)\n", numDetectedGpus, numGpuDevices); + return ERR_FATAL; + } + if (numLaps < 1) { + Utils::Print("[ERROR] NUM_LAPS must be positive (got %d)\n", numLaps); + return ERR_FATAL; + } + + // Executors are ordered rank-major so that each rank's GPUs stay together in the matrix + std::vector exes; + for (int rank = 0; rank < numRanks; rank++) + for (int gpu = 0; gpu < numGpuDevices; gpu++) + exes.push_back({EXE_GPU_GFX, gpu, rank}); + int const numExes = (int)exes.size(); + + // Select the (PING, PONG) pairs to measure. Cross-rank pairs exchange flags through fabric + // handles, which requires pod support and both ranks to be in the same pod + std::vector> isMeasured(numExes, std::vector(numExes, false)); + int numPairs = 0, numSkipped = 0; + for (int i = 0; i < numExes; i++) { + for (int j = 0; j < numExes; j++) { + if (i == j && !isP2p && !a2aLocal) continue; + bool canPair = (exes[i].exeRank == exes[j].exeRank); +#ifdef POD_COMM_ENABLED + canPair |= TransferBench::IsSamePod(exes[j].exeRank, exes[i].exeRank); +#endif + if (!canPair) { + numSkipped++; + continue; + } + isMeasured[i][j] = true; + numPairs++; + } + } + + if (numSkipped > 0) + Utils::Print("[WARN] %d cross-rank pair(s) are shown as N/A: this requires pod communication support and ranks in the same pod\n", + numSkipped); + if (numPairs == 0) { + Utils::Print("[WARN] No pairs to measure. %s requires at least 2 GPUs%s\n", + presetName.c_str(), isA2a ? " (or A2A_LOCAL=1)" : ""); + return ERR_NONE; + } + + TransferBench::ConfigOptions cfg = ev.ToConfigOptions(); + TransferBench::TestResults results; + + // Round-trip latency per lap in microseconds, negative when not measured + std::vector> latencyUs(numExes, std::vector(numExes, -1.0)); + + // Runs the given (PING, PONG) pairs in parallel + auto runPairs = [&](std::vector> const& pairs) { + std::vector transfers; + for (auto const& [i, j] : pairs) { + MemDevice const pingMem = {memType, exes[i].exeIndex, exes[i].exeRank}; + MemDevice const pongMem = {memType, exes[j].exeIndex, exes[j].exeRank}; + + Transfer t; + t.numBytes = 8; // Only sizes the flag seed allocation; pingpong always exchanges 1-byte flags + t.numSubExecs = 1; + t.numLaps = numLaps; + t.srcs = {{MEM_NULL, 0}, {MEM_NULL, 0}}; + t.exeDevice = exes[i]; + t.exeDevicePong = exes[j]; + // Each half writes to its own DST, which the partner half polls + if (useRemoteRead) t.dsts = {pingMem, pongMem}; + else t.dsts = {pongMem, pingMem}; + transfers.push_back(t); + } + + if (!TransferBench::RunTransfers(cfg, transfers, results)) + Utils::PrintErrors(results.errResults); + + for (size_t k = 0; k < pairs.size(); k++) + latencyUs[pairs[k].first][pairs[k].second] = results.tfrResults[k].avgDurationMsec * 1000.0; + }; + + char const sep = ev.outputToCsv ? ',' : ' '; + int const width = numRanks > 1 ? 12 : 10; + auto exeStr = [&](ExeDevice const& exe) { + char buf[32]; + if (numRanks > 1) snprintf(buf, sizeof(buf), "R%d GPU %02d", exe.exeRank, exe.exeIndex); + else snprintf(buf, sizeof(buf), "GPU %02d", exe.exeIndex); + return std::string(buf); + }; + + Utils::Print("Pingpong round-trip latency per lap (us), %s\n", + isP2p ? "each pair run by itself" : + isOne2All ? "each PING executor run against all PONG executors in parallel" + : "all pairs run in parallel"); + Utils::Print("[%d laps] [%s memory flags] [%s]\n", numLaps, memTypeStr.c_str(), + useRemoteRead ? "local write / remote poll" : "remote write / local poll"); + + // PING executors are rows, PONG executors are columns + Utils::Print("%*s", width, "PING\\PONG"); + for (int j = 0; j < numExes; j++) + Utils::Print("%c%*s", sep, width, exeStr(exes[j]).c_str()); + Utils::Print("\n"); + + if (isA2a) { + std::vector> pairs; + for (int i = 0; i < numExes; i++) + for (int j = 0; j < numExes; j++) + if (isMeasured[i][j]) pairs.push_back({i, j}); + runPairs(pairs); + } + + for (int i = 0; i < numExes; i++) { + if (isP2p) { + for (int j = 0; j < numExes; j++) + if (isMeasured[i][j]) runPairs({{i, j}}); + } else if (isOne2All) { + std::vector> pairs; + for (int j = 0; j < numExes; j++) + if (isMeasured[i][j]) pairs.push_back({i, j}); + if (!pairs.empty()) runPairs(pairs); + } + + Utils::Print("%*s", width, exeStr(exes[i]).c_str()); + for (int j = 0; j < numExes; j++) { + if (latencyUs[i][j] < 0) + Utils::Print("%c%*s", sep, width, "N/A"); + else + Utils::Print("%c%*.3f", sep, width, latencyUs[i][j]); + } + Utils::Print("\n"); + } + + // Summarize latencies between distinct executors + int count = 0; + double sumUs = 0.0; + std::pair minPair, maxPair; + for (int i = 0; i < numExes; i++) { + for (int j = 0; j < numExes; j++) { + if (i == j || latencyUs[i][j] < 0) continue; + if (count == 0 || latencyUs[i][j] < latencyUs[minPair.first][minPair.second]) minPair = {i, j}; + if (count == 0 || latencyUs[i][j] > latencyUs[maxPair.first][maxPair.second]) maxPair = {i, j}; + sumUs += latencyUs[i][j]; + count++; + } + } + if (count > 0) { + Utils::Print("\n"); + Utils::Print("Average latency: %8.3f us over %d pairs\n", sumUs / count, count); + Utils::Print("Minimum latency: %8.3f us (PING %s, PONG %s)\n", latencyUs[minPair.first][minPair.second], + exeStr(exes[minPair.first]).c_str(), exeStr(exes[minPair.second]).c_str()); + Utils::Print("Maximum latency: %8.3f us (PING %s, PONG %s)\n", latencyUs[maxPair.first][maxPair.second], + exeStr(exes[maxPair.first]).c_str(), exeStr(exes[maxPair.second]).c_str()); + } + + if (numRanks > 1 && Utils::HasDuplicateHostname()) { + Utils::Print("[WARN] It is recommended to run TransferBench with one rank per host to avoid potential aliasing of executors\n"); + } + return ERR_NONE; +} diff --git a/src/client/Presets/Presets.hpp b/src/client/Presets/Presets.hpp index 8465e389..d6e62ac6 100644 --- a/src/client/Presets/Presets.hpp +++ b/src/client/Presets/Presets.hpp @@ -38,6 +38,7 @@ THE SOFTWARE. #include "HbmBandwidth.hpp" #include "HealthCheck.hpp" #include "Help.hpp" +#include "Latency.hpp" #include "NicAllToAll.hpp" #include "NicRings.hpp" #include "NicPeerToPeer.hpp" @@ -66,40 +67,43 @@ struct PresetInfo std::map presetFuncMap = { - {"a2a", {AllToAllPreset, "Tests parallel transfers between all pairs of GPU devices"}}, - {"a2a_n", {AllToAllRdmaPreset, "Tests parallel transfers between all pairs of GPU devices using Nearest NIC RDMA transfers"}}, - {"a2asweep", {AllToAllSweepPreset, "Test GFX-based all-to-all transfers swept across different CU and GFX unroll counts"}}, - {"bmasweep", {BmaSweepPreset, "Test and compare batched DMA executor for multi destination copies"}}, - {"empty", {EmptyKernelPreset, "Empty GFX kernel launch latency"}}, - {"envvars", {EnvVarsPreset, "Show list of environment variables that can be used to modify behavior"}}, - {"gfxsweep", {GfxSweepPreset, "Sweep over various GFX kernel options for a given GFX Transfer"}}, - {"hbm", {HbmBandwidthPreset, "Tests HBM bandwidth"}}, - {"healthcheck", {HealthCheckPreset, "Simple bandwidth health check (MI300X series only)"}}, - {"help", {HelpPreset, "Shows example usage details"}}, - {"nica2a", {NicAllToAllPreset, "All-to-all GPU traffic over NIC transfers using each NIC's closest GPU/CPU endpoint"}}, - {"nicp2p", {NicPeerToPeerPreset, "Multi-node peer-to-peer RDMA transfer test between all NICs"}}, - {"nicrings", {NicRingsPreset, "Tests NIC rings created across identical NIC indices across ranks"}}, - {"one2all", {OneToAllPreset, "Test all subsets of parallel transfers from one GPU to all others"}}, - {"p2p" , {PeerToPeerPreset, "Peer-to-peer device memory bandwidth test"}}, - {"poda2a", {PodAllToAllPreset, "All-to-all transfers between subgroups of ranks within a pod"}}, - {"podp2p", {PodPeerToPeerPreset, "Peer-to-peer transfers test among ranks within a pod"}}, - {"rings", {RingsPreset, "Ring transfers within subgroups of ranks in a pod"}}, - {"rsweep", {SweepPreset, "Randomly sweep through sets of Transfers"}}, - {"scaling", {ScalingPreset, "Run scaling test from one GPU to other devices"}}, - {"schmoo", {SchmooPreset, "Scaling tests for local/remote read/write/copy"}}, - {"smoketest", {SmokeTestPreset, "Simple correctness smoke-test"}}, - {"sweep", {SweepPreset, "Ordered sweep through sets of Transfers"}}, - {"tdmsweep", {TdmSweepPreset, "Sweep over TDM executor options (block size / LDS / order / subExecs) for a given TDM Transfer"}}, - {"wallclock", {WallClockPreset, "Tests wallclock consistency across XCCs within a GPU"}}, + {"a2a", {AllToAllPreset, "Tests parallel transfers between all pairs of GPU devices"}}, + {"a2a_latency", {LatencyPreset, "Latency values between all pairs when all are run in parallel"}}, + {"a2a_n", {AllToAllRdmaPreset, "Tests parallel transfers between all pairs of GPU devices using Nearest NIC RDMA transfers"}}, + {"a2asweep", {AllToAllSweepPreset, "Test GFX-based all-to-all transfers swept across different CU and GFX unroll counts"}}, + {"bmasweep", {BmaSweepPreset, "Test and compare batched DMA executor for multi destination copies"}}, + {"empty", {EmptyKernelPreset, "Empty GFX kernel launch latency"}}, + {"envvars", {EnvVarsPreset, "Show list of environment variables that can be used to modify behavior"}}, + {"gfxsweep", {GfxSweepPreset, "Sweep over various GFX kernel options for a given GFX Transfer"}}, + {"hbm", {HbmBandwidthPreset, "Tests HBM bandwidth"}}, + {"healthcheck", {HealthCheckPreset, "Simple bandwidth health check (MI300X series only)"}}, + {"help", {HelpPreset, "Shows example usage details"}}, + {"nica2a", {NicAllToAllPreset, "All-to-all GPU traffic over NIC transfers using each NIC's closest GPU/CPU endpoint"}}, + {"nicp2p", {NicPeerToPeerPreset, "Multi-node peer-to-peer RDMA transfer test between all NICs"}}, + {"nicrings", {NicRingsPreset, "Tests NIC rings created across identical NIC indices across ranks"}}, + {"one2all", {OneToAllPreset, "Test all subsets of parallel transfers from one GPU to all others"}}, + {"one2all_latency", {LatencyPreset, "Latency values from one Executor to many, run in parallel"}}, + {"p2p", {PeerToPeerPreset, "Peer-to-peer device memory bandwidth test"}}, + {"p2p_latency", {LatencyPreset, "Latency values between pairs of Executors, run serially"}}, + {"poda2a", {PodAllToAllPreset, "All-to-all transfers between subgroups of ranks within a pod"}}, + {"podp2p", {PodPeerToPeerPreset, "Peer-to-peer transfers test among ranks within a pod"}}, + {"rings", {RingsPreset, "Ring transfers within subgroups of ranks in a pod"}}, + {"rsweep", {SweepPreset, "Randomly sweep through sets of Transfers"}}, + {"scaling", {ScalingPreset, "Run scaling test from one GPU to other devices"}}, + {"schmoo", {SchmooPreset, "Scaling tests for local/remote read/write/copy"}}, + {"smoketest", {SmokeTestPreset, "Simple correctness smoke-test"}}, + {"sweep", {SweepPreset, "Ordered sweep through sets of Transfers"}}, + {"tdmsweep", {TdmSweepPreset, "Sweep over TDM executor options (block size / LDS / order / subExecs) for a given TDM Transfer"}}, + {"wallclock", {WallClockPreset, "Tests wallclock consistency across XCCs within a GPU"}}, }; void DisplayPresets() { if (!Utils::RankDoesOutput()) return; - printf(" %-12s | %-56s\n", "Preset", "Description"); + printf(" %-15s | %-56s\n", "Preset", "Description"); printf("=============================================================================================================\n"); for (auto const& x : presetFuncMap) { - printf(" %-12s | %-56s\n", + printf(" %-15s | %-56s\n", x.first.c_str(), x.second.description.c_str()); } From f75f9c1d9c79f28379e457258109c0945550e60b Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Thu, 24 Sep 2026 20:20:06 -0500 Subject: [PATCH 10/13] xcc subindices reporting + advanced mode fix --- src/header/TransferBench.hpp | 86 ++++++++++++++++++++++++------------ 1 file changed, 58 insertions(+), 28 deletions(-) diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 523c2809..a946a68e 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -211,6 +211,7 @@ namespace TransferBench * exeDevice / exeSubIndex / exeSubSlot = ping executor (any ExeType) * exeDevicePong / exeSubIndexPong / exeSubSlotPong = pong executor (any ExeType) * numLaps = number of pingpong laps (must be > 0); specify as "+N" after ping half ("+" alone defaults to 1) + * Both halves are plain (SRC EXE DST) triplets in simple and advanced mode alike, e.g. "(N G0 F1 +1000 N G1 F0)" * * Ping/pong executors may differ in type (e.g. ping on GFX, pong on DMA). The struct is * executor-agnostic; additional executor types are enabled by implementing their dispatch @@ -3284,6 +3285,10 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Outputs (ping half only) int64_t startCycle; ///< Start timestamp for in-kernel timing int64_t stopCycle; ///< Stop timestamp for in-kernel timing + + // Outputs (both halves) + uint32_t hwId; ///< Hardware ID of the CU that ran this half + uint32_t xccId; ///< XCC ID that ran this half }; // Internal resources allocated per Transfer @@ -4991,6 +4996,8 @@ const auto& AmdSmiFabricInfoV1(const T& info) p.preferredXccId = subIndex; p.startCycle = 0; p.stopCycle = 0; + p.hwId = ~0u; + p.xccId = ~0u; // Device-resident uint8_t values {0, 1} at rss.srcMem[0] + byteOffset (even/odd lap signaling) if (exeDevice.exeType == EXE_GPU_GFX && !rss.srcMem.empty()) { @@ -6506,6 +6513,9 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (threadIdx.x != 0) return; + uint32_t hwXccId, hwCuId; + GetXccHwId(hwXccId, hwCuId); + bool const isPing = p.numLaps > 0; int const laps = isPing ? p.numLaps : -p.numLaps; @@ -6556,6 +6566,8 @@ const auto& AmdSmiFabricInfoV1(const T& info) p.stopCycle = GetTimestamp(); p.startCycle = startCycle; } + p.xccId = hwXccId; + p.hwId = hwCuId; } // Execute a single GPU Transfer (when using 1 stream per Transfer) @@ -6823,6 +6835,16 @@ const auto& AmdSmiFabricInfoV1(const T& info) pingpongParam = pingpongParamHost.data(); } + if (iteration == 0 && System::Get().IsVerbose()) { + for (TransferResources const& rss : exeInfo.resources) { + if (rss.numLaps == 0 || rss.pingpongParamIdx < 0) continue; + PingpongParam const& p = pingpongParam[rss.pingpongParamIdx]; + System::Get().Log("[INFO] Pingpong %d (%s) on GPU %d: requested XCC %d, ran on XCC %u CU %u\n", + rss.transferIdx, rss.numLaps > 0 ? "ping" : "pong", exeIndex, + p.preferredXccId, p.xccId, p.hwId); + } + } + for (TransferResources& rss : exeInfo.resources) { // Only the ping half timestamps the exchange, and it owns the reported row if (rss.numLaps <= 0 || rss.pingpongParamIdx < 0) continue; @@ -8267,6 +8289,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) // If numTransfers < 0, read 5-tuple (srcMem, exeMem, dstMem, #CUs, #Bytes) // otherwise read triples (srcMem, exeMem, dstMem) + // Pingpongs are always two triples joined by "+N", as #CUs and #Bytes do not apply to them bool const advancedMode = (numTransfers < 0); numTransfers = abs(numTransfers); @@ -8282,19 +8305,36 @@ const auto& AmdSmiFabricInfoV1(const T& info) } for (int i = 0; i < numTransfers; i++) { - size_t numBytes; - if (!advancedMode) { - iss >> srcStr >> exeStr >> dstStr; - if (iss.fail()) { - return {ERR_FATAL, - "Parsing error: Unable to read valid Transfer %d (SRC EXE DST) triplet", i+1}; + size_t numBytes = 0; + iss >> srcStr >> exeStr >> dstStr; + if (iss.fail()) { + return {ERR_FATAL, + "Parsing error: Unable to read valid Transfer %d (SRC EXE DST) triplet", i+1}; + } + + // Check for '+' to detect pingpong; optional lap count immediately follows '+' + // e.g. "+500" or "+" (default numLaps) + std::string nextToken; + auto pos = iss.tellg(); + int numLaps = 1; + bool isPingpong = false; + if (iss >> nextToken && !nextToken.empty() && nextToken[0] == '+') { + isPingpong = true; + if (nextToken.size() > 1) { + numLaps = atoi(nextToken.c_str() + 1); + if (numLaps < 1) + return {ERR_FATAL, "Parsing error: Pingpong %d lap count must be positive (got %d)", i+1, numLaps}; } - numBytes = 0; } else { - iss >> srcStr >> exeStr >> dstStr >> numSubExecs >> numBytesToken; + iss.clear(); + iss.seekg(pos); + } + + if (advancedMode && !isPingpong) { + iss >> numSubExecs >> numBytesToken; if (iss.fail()) { return {ERR_FATAL, - "Parsing error: Unable to read valid Transfer %d (SRC EXE DST $CU #Bytes) tuple", i+1}; + "Parsing error: Unable to read valid Transfer %d (SRC EXE DST #CU #Bytes) tuple", i+1}; } if (sscanf(numBytesToken.c_str(), "%lu", &numBytes) != 1) { return {ERR_FATAL, @@ -8307,6 +8347,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) case 'M': numBytes *= 1024; case 'K': numBytes *= 1024; } + + pos = iss.tellg(); + if (iss >> nextToken && !nextToken.empty() && nextToken[0] == '+') { + return {ERR_FATAL, + "Parsing error: Pingpong %d must not specify #CU / #Bytes, use (SRC EXE DST)+N(SRC EXE DST)", i+1}; + } + iss.clear(); + iss.seekg(pos); } WildcardTransfer wct; @@ -8314,24 +8362,6 @@ const auto& AmdSmiFabricInfoV1(const T& info) ERR_CHECK(ParseMemType(dstStr, wct.mem[1])); ERR_CHECK(ParseExeType(exeStr, wct.exe)); - // Check for '+' to detect pingpong; optional lap count immediately follows '+' - // e.g. "+500" or "+" (default numLaps) - std::string nextToken; - auto pos = iss.tellg(); - int numLaps = 1; - bool isPingpong = false; - if (iss >> nextToken && !nextToken.empty() && nextToken[0] == '+') { - isPingpong = true; - if (nextToken.size() > 1) { - numLaps = atoi(nextToken.c_str() + 1); - if (numLaps < 1) - return {ERR_FATAL, "Parsing error: Pingpong %d lap count must be positive (got %d)", i+1, numLaps}; - } - } else { - iss.clear(); - iss.seekg(pos); - } - if (isPingpong) { // Parse the pong triplet std::string pongSrcStr, pongExeStr, pongDstStr; @@ -8864,7 +8894,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) printExe(t.exeDevicePong, t.exeSubIndexPong, t.exeSubSlotPong); fprintf(dumpCfgFile, "->"); printMem(t.dsts[1]); - fprintf(dumpCfgFile, " %d %lu)", t.numSubExecs, t.numBytes); + fprintf(dumpCfgFile, ")"); } else { // Print SRCs for (auto const& src : t.srcs) { From 352cd1dfa1ccebd89b564baf0be2c8baf93d7677 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Mon, 28 Sep 2026 10:23:35 -0500 Subject: [PATCH 11/13] fix pingpong timing wrt subiterations --- src/header/TransferBench.hpp | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index a946a68e..ae8a561c 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -7379,6 +7379,19 @@ const auto& AmdSmiFabricInfoV1(const T& info) std::vector const& transfers, TestResults& results) { + // Pingpong runs its laps once per launch rather than once per subiteration, + // a test made up only of pingpongs is timed with a single subiteration + if (cfg.general.numSubIterations != 1 && !transfers.empty() && + std::all_of(transfers.begin(), transfers.end(), [](Transfer const& t) { return t.numLaps != 0; })) { + ConfigOptions pingpongCfg = cfg; + pingpongCfg.general.numSubIterations = 1; + bool const success = RunTransfers(pingpongCfg, transfers, results); + results.errResults.push_back({ERR_WARN, + "[general.numSubIterations] (%d) is ignored when all Transfers are pingpongs", + cfg.general.numSubIterations}); + return success; + } + // Clear all errors; auto& errResults = results.errResults; errResults.clear(); From 60c7d5c8fe6419d96d5bd20d2b3aee4ac4efe837 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Mon, 28 Sep 2026 10:24:20 -0500 Subject: [PATCH 12/13] addition of latency ring; skip failing latency tests --- src/client/Presets/Latency.hpp | 117 +++++++++++++++++++++++++++------ src/client/Presets/Presets.hpp | 1 + 2 files changed, 99 insertions(+), 19 deletions(-) diff --git a/src/client/Presets/Latency.hpp b/src/client/Presets/Latency.hpp index 564415e2..d85a47c4 100644 --- a/src/client/Presets/Latency.hpp +++ b/src/client/Presets/Latency.hpp @@ -24,6 +24,7 @@ THE SOFTWARE. // p2p_latency - Every pair runs by itself, one pair at a time // one2all_latency - One PING executor runs against all PONG executors at once, one PING executor at a time // a2a_latency - Every pair runs at once +// ring_latency - Each GPU pings its neighbor; every hop in the ring(s) runs at once int LatencyPreset(EnvVars& ev, size_t const /*numBytesPerTransfer*/, std::string const presetName, @@ -32,6 +33,7 @@ int LatencyPreset(EnvVars& ev, bool const isP2p = (presetName == "p2p_latency"); bool const isOne2All = (presetName == "one2all_latency"); bool const isA2a = (presetName == "a2a_latency"); + bool const isRing = (presetName == "ring_latency"); if (!Utils::AllRanksHaveSameGpuCount()) { Utils::Print("[ERROR] %s preset requires all ranks to have the same number of GPUs\n", presetName.c_str()); @@ -48,6 +50,8 @@ int LatencyPreset(EnvVars& ev, int numGpuDevices = EnvVars::GetEnvVar("NUM_GPU_DEVICES", numDetectedGpus); int numLaps = EnvVars::GetEnvVar("NUM_LAPS" , 1000); int useRemoteRead = EnvVars::GetEnvVar("USE_REMOTE_READ", 0); + int stride = isRing ? EnvVars::GetEnvVar("STRIDE", 1) : 1; + int ringSize = isRing ? EnvVars::GetEnvVar("RING_SIZE", numRanks * numGpuDevices) : 0; MemType const memType = Utils::GetGpuMemType(memTypeIdx); std::string const memTypeStr = Utils::GetGpuMemTypeStr(memTypeIdx); @@ -62,6 +66,10 @@ int LatencyPreset(EnvVars& ev, ev.Print("GPU_MEM_TYPE", memTypeIdx, "Using %s memory for flags (%s)", memTypeStr.c_str(), Utils::GetAllGpuMemTypeStr().c_str()); ev.Print("NUM_GPU_DEVICES", numGpuDevices, "Using %d GPUs%s", numGpuDevices, numRanks > 1 ? " per rank" : ""); ev.Print("NUM_LAPS", numLaps, "Timing %d round trips per iteration", numLaps); + if (isRing) { + ev.Print("RING_SIZE", ringSize, "Building rings of size %d", ringSize); + ev.Print("STRIDE", stride, "Reordering devices by taking %d steps", stride); + } ev.Print("USE_REMOTE_READ", useRemoteRead, "%s", useRemoteRead ? "Executors write to their own memory and poll their partner's" : "Executors write to their partner's memory and poll their own"); printf("\n"); @@ -73,6 +81,10 @@ int LatencyPreset(EnvVars& ev, IS_UNIFORM(memTypeIdx, "GPU_MEM_TYPE"); IS_UNIFORM(numGpuDevices, "NUM_GPU_DEVICES"); IS_UNIFORM(numLaps, "NUM_LAPS"); + if (isRing) { + IS_UNIFORM(ringSize, "RING_SIZE"); + IS_UNIFORM(stride, "STRIDE"); + } IS_UNIFORM(useRemoteRead, "USE_REMOTE_READ"); // Validate env vars @@ -92,23 +104,64 @@ int LatencyPreset(EnvVars& ev, exes.push_back({EXE_GPU_GFX, gpu, rank}); int const numExes = (int)exes.size(); + // Ring hops follow the same ordering as the rings preset: devices are listed rank-major, + // reordered by STRIDE, then split into rings of RING_SIZE. GPU i pings the next GPU in its ring. + std::vector ringOrder; + if (isRing) { + if (ringSize <= 1) { + if (numExes < 2) + Utils::Print("[ERROR] %s requires at least 2 GPUs\n", presetName.c_str()); + else + Utils::Print("[ERROR] RING_SIZE must be greater than 1 (got %d)\n", ringSize); + return ERR_FATAL; + } + if (numExes % ringSize) { + Utils::Print("[ERROR] Ring size %d must evenly divide the total number of GPUs %d\n", ringSize, numExes); + return ERR_FATAL; + } + ringOrder.resize(numExes); + for (int i = 0; i < numExes; i++) ringOrder[i] = i; + Utils::StrideGenerate(ringOrder, stride); + } + // Select the (PING, PONG) pairs to measure. Cross-rank pairs exchange flags through fabric // handles, which requires pod support and both ranks to be in the same pod - std::vector> isMeasured(numExes, std::vector(numExes, false)); - int numPairs = 0, numSkipped = 0; - for (int i = 0; i < numExes; i++) { - for (int j = 0; j < numExes; j++) { - if (i == j && !isP2p && !a2aLocal) continue; - bool canPair = (exes[i].exeRank == exes[j].exeRank); + auto canPair = [&](int i, int j) { + bool ok = (exes[i].exeRank == exes[j].exeRank); #ifdef POD_COMM_ENABLED - canPair |= TransferBench::IsSamePod(exes[j].exeRank, exes[i].exeRank); + ok |= TransferBench::IsSamePod(exes[j].exeRank, exes[i].exeRank); #endif - if (!canPair) { - numSkipped++; - continue; + return ok; + }; + + std::vector> isMeasured(numExes, std::vector(numExes, false)); + int numPairs = 0, numSkipped = 0; + if (isRing) { + int const numRings = numExes / ringSize; + for (int ringIdx = 0; ringIdx < numRings; ringIdx++) { + int const ringBase = ringIdx * ringSize; + for (int hop = 0; hop < ringSize; hop++) { + int const src = ringOrder[ringBase + hop]; + int const dst = ringOrder[ringBase + (hop + 1) % ringSize]; + if (!canPair(src, dst)) { + numSkipped++; + continue; + } + isMeasured[src][dst] = true; + numPairs++; + } + } + } else { + for (int i = 0; i < numExes; i++) { + for (int j = 0; j < numExes; j++) { + if (i == j && !isP2p && !a2aLocal) continue; + if (!canPair(i, j)) { + numSkipped++; + continue; + } + isMeasured[i][j] = true; + numPairs++; } - isMeasured[i][j] = true; - numPairs++; } } @@ -147,8 +200,11 @@ int LatencyPreset(EnvVars& ev, transfers.push_back(t); } - if (!TransferBench::RunTransfers(cfg, transfers, results)) - Utils::PrintErrors(results.errResults); + if (!TransferBench::RunTransfers(cfg, transfers, results)) { + for (auto const& err : results.errResults) + Utils::Print("[%s] %s\n", err.errType == ERR_FATAL ? "ERROR" : "WARN", err.errMsg.c_str()); + return; + } for (size_t k = 0; k < pairs.size(); k++) latencyUs[pairs[k].first][pairs[k].second] = results.tfrResults[k].avgDurationMsec * 1000.0; @@ -163,10 +219,33 @@ int LatencyPreset(EnvVars& ev, return std::string(buf); }; - Utils::Print("Pingpong round-trip latency per lap (us), %s\n", - isP2p ? "each pair run by itself" : - isOne2All ? "each PING executor run against all PONG executors in parallel" - : "all pairs run in parallel"); + char modeBuf[192]; + if (isRing) { + int const numRings = numExes / ringSize; + if (stride == 1) + snprintf(modeBuf, sizeof(modeBuf), "%d parallel ring%s of %d, each GPU pinging its neighbor", + numRings, numRings == 1 ? "" : "s", ringSize); + else + snprintf(modeBuf, sizeof(modeBuf), "%d parallel ring%s of %d (stride %d), each GPU pinging its neighbor", + numRings, numRings == 1 ? "" : "s", ringSize, stride); + } else if (isP2p) { + snprintf(modeBuf, sizeof(modeBuf), "each pair run by itself"); + } else if (isOne2All) { + snprintf(modeBuf, sizeof(modeBuf), "each PING executor run against all PONG executors in parallel"); + } else { + snprintf(modeBuf, sizeof(modeBuf), "all pairs run in parallel"); + } + Utils::Print("Pingpong round-trip latency per lap (us), %s\n", modeBuf); + if (isRing) { + int const numRings = numExes / ringSize; + for (int ringIdx = 0; ringIdx < numRings; ringIdx++) { + int const ringBase = ringIdx * ringSize; + Utils::Print("Ring %02d:", ringIdx); + for (int hop = 0; hop < ringSize; hop++) + Utils::Print(" %s ->", exeStr(exes[ringOrder[ringBase + hop]]).c_str()); + Utils::Print(" %s\n", exeStr(exes[ringOrder[ringBase]]).c_str()); + } + } Utils::Print("[%d laps] [%s memory flags] [%s]\n", numLaps, memTypeStr.c_str(), useRemoteRead ? "local write / remote poll" : "remote write / local poll"); @@ -176,7 +255,7 @@ int LatencyPreset(EnvVars& ev, Utils::Print("%c%*s", sep, width, exeStr(exes[j]).c_str()); Utils::Print("\n"); - if (isA2a) { + if (isA2a || isRing) { std::vector> pairs; for (int i = 0; i < numExes; i++) for (int j = 0; j < numExes; j++) diff --git a/src/client/Presets/Presets.hpp b/src/client/Presets/Presets.hpp index d6e62ac6..86410a94 100644 --- a/src/client/Presets/Presets.hpp +++ b/src/client/Presets/Presets.hpp @@ -87,6 +87,7 @@ std::map presetFuncMap = {"p2p_latency", {LatencyPreset, "Latency values between pairs of Executors, run serially"}}, {"poda2a", {PodAllToAllPreset, "All-to-all transfers between subgroups of ranks within a pod"}}, {"podp2p", {PodPeerToPeerPreset, "Peer-to-peer transfers test among ranks within a pod"}}, + {"ring_latency", {LatencyPreset, "Each GPU pingpongs its neighbor, all hops in parallel"}}, {"rings", {RingsPreset, "Ring transfers within subgroups of ranks in a pod"}}, {"rsweep", {SweepPreset, "Randomly sweep through sets of Transfers"}}, {"scaling", {ScalingPreset, "Run scaling test from one GPU to other devices"}}, From 91f99bddb2b55619f8634078ddaf28f794c1fc98 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Mon, 28 Sep 2026 10:26:53 -0500 Subject: [PATCH 13/13] TransferBench v1.71.00 --- .github/workflows/README_BUILD_PACKAGES.md | 6 ++--- CHANGELOG.md | 28 ++++++++++++++++++++++ CMakeLists.txt | 4 ++-- src/client/EnvVars.hpp | 2 +- src/header/TransferBench.hpp | 2 +- 5 files changed, 35 insertions(+), 7 deletions(-) diff --git a/.github/workflows/README_BUILD_PACKAGES.md b/.github/workflows/README_BUILD_PACKAGES.md index bbaecd83..c17cfd09 100644 --- a/.github/workflows/README_BUILD_PACKAGES.md +++ b/.github/workflows/README_BUILD_PACKAGES.md @@ -70,9 +70,9 @@ sudo -E BUILD_TYPE=Debug ./build_packages_local.sh After the script completes, packages live under `build/`: ``` -build/amdrocm7-transferbench_1.66.02-_amd64.deb -build/amdrocm7-transferbench-1.66.02-.x86_64.rpm -build/amdrocm7-transferbench-1.66.02-Linux.tar.gz +build/amdrocm7-transferbench_1.71.00-_amd64.deb +build/amdrocm7-transferbench-1.71.00-.x86_64.rpm +build/amdrocm7-transferbench-1.71.00-Linux.tar.gz ``` ## Installing built packages diff --git a/CHANGELOG.md b/CHANGELOG.md index 22c87b02..92cf0e58 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -3,6 +3,34 @@ Documentation for TransferBench is available at [https://rocm.docs.amd.com/projects/TransferBench](https://rocm.docs.amd.com/projects/TransferBench). +## v1.71.00 +### Added +- Added PingPong operations to measure latency between two GPU Executors. A PingPong joins a Ping and a Pong triplet + with `+` and an optional lap count, e.g. `(N->G0->F1)+1000(N->G1->F0)`. Each lap, the Ping Executor writes a + 1-byte flag to its DST and waits for the Pong Executor to write one back to the Pong DST. Results report the + GPU-timed round-trip latency per lap + - PingPongs can run in parallel with regular Transfers in the same Test, and do not count toward Executor bandwidth + - Each Ping and Pong uses a single threadblock; the number of SubExecutors and bytes to transfer are ignored + - Currently supported on GFX Executors only + - Ping and Pong may run on different ranks within the same pod (requires pod communication support) +- Added `PINGPONG_FLAG_BUFFER` and `PINGPONG_STRIDE` to spread PingPong flags across a buffer, moving by the stride each lap +- Added latency presets built on PingPong operations: + - "p2p_latency": latency between pairs of GPU Executors, run serially + - "one2all_latency": latency from one GPU Executor to all others, run in parallel + - "a2a_latency": latency between all pairs of GPU Executors, run in parallel + - Behavior can be modified with `NUM_LAPS`, `GPU_MEM_TYPE`, `USE_REMOTE_READ`, `NUM_GPU_DEVICES` and, for a2a_latency, `A2A_LOCAL` +- Added gfx1250-strict to the default GPU targets for CMake builds +- NIC topology output now shows each NIC's IBV max_msg_sz +### Modified +- When GFX Executor timing does not use HIP events (multi-stream mode or `USE_HIP_EVENTS=0`), it now spans from the first + Transfer launch to the last Transfer completion, excluding executor dispatch overhead +### Fixed +- CU IDs reported with `SHOW_ITERATIONS` now match `CU_MASK` indices on gfx90a, gfx942, gfx950 and gfx1250 +- GFX kernel launch failures are now reported instead of being ignored +- RoCE / GID index fields are now initialized for NICs that are not RoCE or have no active port +- NIC Transfers are now rejected during validation if a queue pair's work request (the smaller of `NIC_CHUNK_BYTES` and + the bytes assigned to that queue pair) exceeds the IBV max_msg_sz of either NIC + ## v1.70.02 ### Modified - rings preset defaults `NUM_SUB_EXEC` to 0, which uses all available subexecutors per Transfer diff --git a/CMakeLists.txt b/CMakeLists.txt index cc11a742..c6703d75 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -103,8 +103,8 @@ set(ENV{ROCM_PATH} "${ROCM_PATH}") # TransferBench project definitions #================================================================================================== set(TRANSFERBENCH_VERSION_MAJOR 1) -set(TRANSFERBENCH_VERSION_MINOR 70) -set(TRANSFERBENCH_VERSION_PATCH_FALLBACK "02") +set(TRANSFERBENCH_VERSION_MINOR 71) +set(TRANSFERBENCH_VERSION_PATCH_FALLBACK "00") # Auto-compute patch from git: count commits since the last v..* tag. # Falls back to TRANSFERBENCH_VERSION_PATCH_FALLBACK when git is unavailable, diff --git a/src/client/EnvVars.hpp b/src/client/EnvVars.hpp index dc74d6fa..c81811e5 100644 --- a/src/client/EnvVars.hpp +++ b/src/client/EnvVars.hpp @@ -44,7 +44,7 @@ THE SOFTWARE. #include #include -#define CLIENT_VERSION "02" +#define CLIENT_VERSION "00" #include "TransferBench.hpp" using namespace TransferBench; diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index ae8a561c..32992252 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -102,7 +102,7 @@ namespace TransferBench using std::set; using std::vector; - constexpr char VERSION[] = "1.70"; + constexpr char VERSION[] = "1.71"; /** * Enumeration of supported Executor types