From 594eb703d941f55a4dcde6809401e7e03ab952cc Mon Sep 17 00:00:00 2001 From: gilbertlee-amd <44450918+gilbertlee-amd@users.noreply.github.com> Date: Mon, 29 Jun 2026 10:10:46 -0600 Subject: [PATCH 1/2] Fixing GFX Executor missing threadfence_system from v1.67 (#327) --- CHANGELOG.md | 6 ++++ src/header/TransferBench.hpp | 67 ++++++++++++++++++++---------------- 2 files changed, 44 insertions(+), 29 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 84847a37..e5c22af8 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -3,6 +3,12 @@ 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 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 - 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..d72f4f88 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 @@ -5036,21 +5036,25 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } } } - // Allows for numSubiterations == 0 to run infinitely - if (++subIterations == numSubIterations) break; - } - // Wait for all threads to finish - if (seType == 1) { - // For warp-level, sync within warp only + // Drain every thread's writes out to system scope + __threadfence_system(); + + // 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) { @@ -5213,22 +5217,25 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } } } - // Allows for numSubiterations == 0 to run infinitely - if (++subIterations == numSubIterations) break; - } - // Wait for all threads to finish - if (seType == 1) { - // For warp-level, sync within warp only + // Drain every thread's writes out to system scope + __threadfence_system(); + + // 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) { @@ -5459,11 +5466,13 @@ 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))); + } } } From 69464481a1163fc8407968c632b379b6ba05dc6b Mon Sep 17 00:00:00 2001 From: gilbertlee-amd <44450918+gilbertlee-amd@users.noreply.github.com> Date: Mon, 29 Jun 2026 15:05:09 -0600 Subject: [PATCH 2/2] Adjusting location of the fence (#329) --- src/header/TransferBench.hpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index d72f4f88..b2f3ffdc 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -5037,9 +5037,6 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } } - // Drain every thread's writes out to system scope - __threadfence_system(); - // Wait for all threads to finish this subiteration if (seType == 1) { // For warp-level, sync within warp only @@ -5058,6 +5055,9 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } if (shouldRecordTiming) { + // Previous sync ensures data from all threads in this subexecutor has landed in cache + // so only one thread needs to do system level threadfence. + __threadfence_system(); p.stopCycle = GetTimestamp(); p.startCycle = startCycle; GetHwId(p.hwId); @@ -5218,9 +5218,6 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } } - // Drain every thread's writes out to system scope - __threadfence_system(); - // Wait for all threads to finish this subiteration if (seType == 1) { // For warp-level, sync within warp only @@ -5239,6 +5236,9 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } if (shouldRecordTiming) { + // Previous sync ensures data from all threads in this subexecutor has landed in cache + // so only one thread needs to do system level threadfence. + __threadfence_system(); p.stopCycle = GetTimestamp(); p.startCycle = startCycle; GetHwId(p.hwId);