From 91a2d105b08ca26f1afeadb56cebd75f1e005ec2 Mon Sep 17 00:00:00 2001 From: Gilbert Lee Date: Sun, 28 Jun 2026 06:59:38 +0000 Subject: [PATCH 1/2] Fixing GFX Executor missing threadfence from v1.67 + warp-subexec timing bug --- CHANGELOG.md | 5 +++++ src/header/TransferBench.hpp | 24 ++++++++++++++++-------- 2 files changed, 21 insertions(+), 8 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 84847a37..bea60804 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -3,6 +3,11 @@ Documentation for TransferBench is available at [https://rocm.docs.amd.com/projects/TransferBench](https://rocm.docs.amd.com/projects/TransferBench). +## v1.68.00 +### Fixed +- Improper draining of writes introduced in v1.67 that impacts GFX executor timing +- Potential timing bug when running GFX Executor in warp-subexecutor mode with small data sizes + ## v1.67.00 ### Added - Added NIC_TRAFFIC_CLASS to set the DSCP/traffic class byte in the RoCE GRH for QPs (RoCE only) diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index fa536155..255af546 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -92,7 +92,7 @@ namespace TransferBench using std::set; using std::vector; - constexpr char VERSION[] = "1.67"; + constexpr char VERSION[] = "1.68"; /** * Enumeration of supported Executor types @@ -5040,6 +5040,9 @@ static bool IsConfiguredGid(union ibv_gid const& gid) if (++subIterations == numSubIterations) break; } + // Drain every thread's writes out to system scope + __threadfence_system(); + // Wait for all threads to finish if (seType == 1) { // For warp-level, sync within warp only @@ -5217,6 +5220,9 @@ static bool IsConfiguredGid(union ibv_gid const& gid) if (++subIterations == numSubIterations) break; } + // Drain every thread's writes out to system scope + __threadfence_system(); + // Wait for all threads to finish if (seType == 1) { // For warp-level, sync within warp only @@ -5400,8 +5406,8 @@ static bool IsConfiguredGid(union ibv_gid const& gid) int const exeIndex, ExeInfo& exeInfo) { - auto cpuStart = std::chrono::high_resolution_clock::now(); ERR_CHECK(hipSetDevice(exeIndex)); + auto cpuStart = std::chrono::high_resolution_clock::now(); int xccDim = exeInfo.useSubIndices ? exeInfo.numSubIndices : 1; @@ -5459,12 +5465,14 @@ static bool IsConfiguredGid(union ibv_gid const& gid) std::set> CUs; for (auto subExecIdx : rss.subExecIdx) { - minStartCycle = std::min(minStartCycle, exeInfo.subExecParamGpu[subExecIdx].startCycle); - maxStopCycle = std::max(maxStopCycle, exeInfo.subExecParamGpu[subExecIdx].stopCycle); - if (cfg.general.recordPerIteration) { - CUs.insert(std::make_pair(exeInfo.subExecParamGpu[subExecIdx].xccId, - GetId(exeInfo.subExecParamGpu[subExecIdx].hwId))); - } + if (exeInfo.subExecParamCpu[subExecIdx].N != 0) { + minStartCycle = std::min(minStartCycle, exeInfo.subExecParamGpu[subExecIdx].startCycle); + maxStopCycle = std::max(maxStopCycle, exeInfo.subExecParamGpu[subExecIdx].stopCycle); + if (cfg.general.recordPerIteration) { + CUs.insert(std::make_pair(exeInfo.subExecParamGpu[subExecIdx].xccId, + GetId(exeInfo.subExecParamGpu[subExecIdx].hwId))); + } + } } double deltaMsec = (maxStopCycle - minStartCycle) / (double)(exeInfo.wallClockRate); From ddc7ba350fa61ed15c8717ccd452208ab5531480 Mon Sep 17 00:00:00 2001 From: Gilbert Lee Date: Sun, 28 Jun 2026 08:54:55 +0000 Subject: [PATCH 2/2] Keeping subiterations in sync --- CHANGELOG.md | 3 +- src/header/TransferBench.hpp | 73 ++++++++++++++++++------------------ 2 files changed, 39 insertions(+), 37 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index bea60804..e5c22af8 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -5,8 +5,9 @@ Documentation for TransferBench is available at ## v1.68.00 ### Fixed -- Improper draining of writes introduced in v1.67 that impacts GFX executor timing +- Improper draining of writes that could artificially inflate transfer timing - Potential timing bug when running GFX Executor in warp-subexecutor mode with small data sizes +- Keeping subiteration threads in sync per subiteration ## v1.67.00 ### Added diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 255af546..d72f4f88 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -5036,24 +5036,25 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } } } - // Allows for numSubiterations == 0 to run infinitely - if (++subIterations == numSubIterations) break; - } - // Drain every thread's writes out to system scope - __threadfence_system(); + // Drain every thread's writes out to system scope + __threadfence_system(); - // Wait for all threads to finish - if (seType == 1) { - // For warp-level, sync within warp only + // Wait for all threads to finish this subiteration + if (seType == 1) { + // For warp-level, sync within warp only #if defined(__HIP_PLATFORM_AMD__) && (HIP_VERSION_MAJOR < 7) - __builtin_amdgcn_wave_barrier(); + __builtin_amdgcn_wave_barrier(); #else - __syncwarp(); + __syncwarp(); #endif - } else { - // For threadblock-level, sync all threads - __syncthreads(); + } else { + // For threadblock-level, sync all threads + __syncthreads(); + } + + // Allows for numSubiterations == 0 to run infinitely + if (++subIterations == numSubIterations) break; } if (shouldRecordTiming) { @@ -5216,25 +5217,25 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } } } - // Allows for numSubiterations == 0 to run infinitely - if (++subIterations == numSubIterations) break; - } - // Drain every thread's writes out to system scope - __threadfence_system(); + // Drain every thread's writes out to system scope + __threadfence_system(); - // Wait for all threads to finish - if (seType == 1) { - // For warp-level, sync within warp only + // Wait for all threads to finish this subiteration + if (seType == 1) { + // For warp-level, sync within warp only #if defined(__HIP_PLATFORM_AMD__) && (HIP_VERSION_MAJOR < 7) - __builtin_amdgcn_wave_barrier(); + __builtin_amdgcn_wave_barrier(); #else - - __syncwarp(); + __syncwarp(); #endif - } else { - // For threadblock-level, sync all threads - __syncthreads(); + } else { + // For threadblock-level, sync all threads + __syncthreads(); + } + + // Allows for numSubiterations == 0 to run infinitely + if (++subIterations == numSubIterations) break; } if (shouldRecordTiming) { @@ -5406,8 +5407,8 @@ static bool IsConfiguredGid(union ibv_gid const& gid) int const exeIndex, ExeInfo& exeInfo) { - ERR_CHECK(hipSetDevice(exeIndex)); auto cpuStart = std::chrono::high_resolution_clock::now(); + ERR_CHECK(hipSetDevice(exeIndex)); int xccDim = exeInfo.useSubIndices ? exeInfo.numSubIndices : 1; @@ -5465,14 +5466,14 @@ static bool IsConfiguredGid(union ibv_gid const& gid) std::set> CUs; for (auto subExecIdx : rss.subExecIdx) { - if (exeInfo.subExecParamCpu[subExecIdx].N != 0) { - minStartCycle = std::min(minStartCycle, exeInfo.subExecParamGpu[subExecIdx].startCycle); - maxStopCycle = std::max(maxStopCycle, exeInfo.subExecParamGpu[subExecIdx].stopCycle); - if (cfg.general.recordPerIteration) { - CUs.insert(std::make_pair(exeInfo.subExecParamGpu[subExecIdx].xccId, - GetId(exeInfo.subExecParamGpu[subExecIdx].hwId))); - } - } + if (exeInfo.subExecParamCpu[subExecIdx].N != 0) { + minStartCycle = std::min(minStartCycle, exeInfo.subExecParamGpu[subExecIdx].startCycle); + maxStopCycle = std::max(maxStopCycle, exeInfo.subExecParamGpu[subExecIdx].stopCycle); + if (cfg.general.recordPerIteration) { + CUs.insert(std::make_pair(exeInfo.subExecParamGpu[subExecIdx].xccId, + GetId(exeInfo.subExecParamGpu[subExecIdx].hwId))); + } + } } double deltaMsec = (maxStopCycle - minStartCycle) / (double)(exeInfo.wallClockRate);