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))); + } } }