From f8064daccf39c230b29cfd5d84cd57c8ab1f38fb Mon Sep 17 00:00:00 2001 From: gilbertlee-amd <44450918+gilbertlee-amd@users.noreply.github.com> Date: Mon, 6 Jul 2026 13:36:23 -0600 Subject: [PATCH 1/3] Fix for using byte offsets with DMA executors (#330) --- CHANGELOG.md | 4 ++++ src/client/Client.cpp | 2 +- src/header/TransferBench.hpp | 20 ++++++++++++-------- 3 files changed, 17 insertions(+), 9 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index e5c22af8..5d745b00 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -3,6 +3,10 @@ Documentation for TransferBench is available at [https://rocm.docs.amd.com/projects/TransferBench](https://rocm.docs.amd.com/projects/TransferBench). +## v1.69.00 +### Fixed +- Fix for non-zero byte offsets used with DMA executor + ## v1.68.00 ### Fixed - Improper draining of writes that could artificially inflate transfer timing diff --git a/src/client/Client.cpp b/src/client/Client.cpp index 3074f3b1..95a0f9a5 100644 --- a/src/client/Client.cpp +++ b/src/client/Client.cpp @@ -259,7 +259,7 @@ void DisplayVersion() #ifndef TB_GIT_COMMIT #define TB_GIT_COMMIT "unknown" #endif - Print("TransferBench v%s.%s (%s:%s)%s%s\n", VERSION, CLIENT_VERSION, TB_GIT_BRANCH, TB_GIT_COMMIT, support.c_str(), multiNodeMode.c_str()); + Print("TransferBench v%s.%sC (%s:%s)%s%s\n", VERSION, CLIENT_VERSION, TB_GIT_BRANCH, TB_GIT_COMMIT, support.c_str(), multiNodeMode.c_str()); Print("=============================================================================================================\n"); } diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index b2f3ffdc..5f22f52c 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.68"; + constexpr char VERSION[] = "1.69"; /** * Enumeration of supported Executor types @@ -5507,6 +5507,8 @@ static bool IsConfiguredGid(union ibv_gid const& gid) int numDsts = (int)resources.dstMem.size(); ERR_CHECK(hipSetDevice(exeIndex)); int subIterations = 0; + size_t const initOffset = cfg.data.byteOffset / sizeof(float); + float* const src = resources.srcMem[0] + initOffset; if (!useSubIndices && !cfg.dma.useHsaCopy) { if (cfg.dma.useHipEvents) ERR_CHECK(hipEventRecord(startEvent, stream)); @@ -5520,12 +5522,13 @@ static bool IsConfiguredGid(union ibv_gid const& gid) do { // Queue for each output location for (int dstIdx = 0; dstIdx < numDsts; dstIdx++) { + float* const dst = resources.dstMem[dstIdx] + initOffset; #if defined(CUMEM_ENABLED) - ERR_CHECK(cuMemcpyAsync((CUdeviceptr)resources.dstMem[dstIdx], - (CUdeviceptr)resources.srcMem[0], + ERR_CHECK(cuMemcpyAsync((CUdeviceptr)dst, + (CUdeviceptr)src, resources.numBytes, stream)); #else - ERR_CHECK(hipMemcpyAsync(resources.dstMem[dstIdx], resources.srcMem[0], resources.numBytes, + ERR_CHECK(hipMemcpyAsync(dst, src, resources.numBytes, memcpyKind, stream)); #endif } @@ -5542,14 +5545,15 @@ static bool IsConfiguredGid(union ibv_gid const& gid) do { hsa_signal_store_screlease(resources.signal, numDsts); for (int dstIdx = 0; dstIdx < numDsts; dstIdx++) { + float* const dst = resources.dstMem[dstIdx] + initOffset; if (!useSubIndices) { - ERR_CHECK(hsa_amd_memory_async_copy(resources.dstMem[dstIdx], resources.dstAgent[dstIdx], - resources.srcMem[0], resources.srcAgent, + ERR_CHECK(hsa_amd_memory_async_copy(dst, resources.dstAgent[dstIdx], + src, resources.srcAgent, resources.numBytes, 0, NULL, resources.signal)); } else { - HSA_CALL(hsa_amd_memory_async_copy_on_engine(resources.dstMem[dstIdx], resources.dstAgent[dstIdx], - resources.srcMem[0], resources.srcAgent, + HSA_CALL(hsa_amd_memory_async_copy_on_engine(dst, resources.dstAgent[dstIdx], + src, resources.srcAgent, resources.numBytes, 0, NULL, resources.signal, resources.sdmaEngineId, true)); From 6e52f749427493aa01832f7e1c29cf39a0465ae6 Mon Sep 17 00:00:00 2001 From: nileshnegi Date: Sun, 10 May 2026 09:01:16 -0500 Subject: [PATCH 2/3] Add DMA memset support (no-src EXE_GPU_DMA transfer) Allow EXE_GPU_DMA transfers with zero sources to perform a memset using hsa_amd_memory_fill, which enqueues a LINEAR_FILL operation on the SDMA engines. Fill value: uint32_t fillVal = bit_cast(MEMSET_VAL); // 0x4B4B4B4B count = numBytes / sizeof(uint32_t); // count is in uint32_t units 0x4B4B4B4B matches both memset(MEMSET_CHAR) (used by dstReference[0]) and MemsetVal() used by the GFX no-src kernel, so existing correctness validation passes without changes. Validation changes (AMD only, gated on !__NVCC__): - DMA no-src is now valid; rejected only on NVIDIA builds - DMA no-src with a specific SDMA engine (e.g. "n d0.2 g1") is rejected because hsa_amd_memory_fill has no engine-selection parameter - Copy-agent-selection warnings guarded by !t.srcs.empty() to avoid out-of-bounds access when no source is specified Execution changes (ExecuteDmaTransfer): - no-src hoisted before hipMemcpy/HSA-async-copy branches - Copy paths (hipMemcpy and HSA async copy) unchanged HSA resource setup: - srcMem pointer-info query guarded by !rss.srcMem.empty() Co-authored-by: Claude --- src/header/TransferBench.hpp | 73 +++++++++++++++++++++++++----------- 1 file changed, 52 insertions(+), 21 deletions(-) diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 5f22f52c..c60f0705 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -2327,12 +2327,20 @@ namespace { } break; case EXE_GPU_DMA: - if (t.srcs.size() != 1) { + if (t.srcs.size() > 1) { + errors.push_back({ERR_FATAL, + "Transfer %d: DMA executor must have 0 or 1 sources", i}); + hasFatalError = true; + break; + } +#if defined(__NVCC__) + if (t.srcs.empty()) { errors.push_back({ERR_FATAL, - "Transfer %d: DMA executor must have exactly 1 source", i}); + "Transfer %d: DMA memset (0 sources) not supported on NVIDIA hardware", i}); hasFatalError = true; break; } +#endif if (t.dsts.size() < 1) { errors.push_back({ERR_FATAL, "Transfer %d: DMA executor must have at least 1 destination", i}); @@ -2355,6 +2363,12 @@ namespace { hasFatalError = true; break; #else + if (t.srcs.empty()) { + errors.push_back({ERR_FATAL, + "Transfer %d: DMA memset (0 sources) does not support engine selection", i}); + hasFatalError = true; + break; + } useSubIndexCount[t.exeDevice]++; int numSubIndices = GetNumExecutorSubIndices(t.exeDevice); if (t.exeSubIndex >= numSubIndices) { @@ -2375,7 +2389,6 @@ namespace { hasFatalError = true; break; } - } int numDsts = (int)t.dsts.size(); @@ -2414,22 +2427,24 @@ namespace { #endif } - if (!IsGpuMemType(t.srcs[0].memType) && !IsGpuMemType(t.dsts[0].memType)) { - errors.push_back({ERR_WARN, - "Transfer %d: No GPU memory for source or destination. Copy might not execute on DMA %d", - i, t.exeDevice.exeIndex}); - } else { - // Currently HIP will use src agent if source memory is GPU, otherwise dst agent - if (IsGpuMemType(t.srcs[0].memType)) { - if (t.srcs[0].memIndex != t.exeDevice.exeIndex) { + if (!t.srcs.empty()) { + if (!IsGpuMemType(t.srcs[0].memType) && !IsGpuMemType(t.dsts[0].memType)) { + errors.push_back({ERR_WARN, + "Transfer %d: No GPU memory for source or destination. Copy might not execute on DMA %d", + i, t.exeDevice.exeIndex}); + } else { + // Currently HIP will use src agent if source memory is GPU, otherwise dst agent + if (IsGpuMemType(t.srcs[0].memType)) { + if (t.srcs[0].memIndex != t.exeDevice.exeIndex) { + errors.push_back({ERR_WARN, + "Transfer %d: DMA executor may automatically switch to using the source memory device (%d) not (%d)", + i, t.srcs[0].memIndex, t.exeDevice.exeIndex}); + } + } else if (t.dsts[0].memIndex != t.exeDevice.exeIndex) { errors.push_back({ERR_WARN, - "Transfer %d: DMA executor may automatically switch to using the source memory device (%d) not (%d)", - i, t.srcs[0].memIndex, t.exeDevice.exeIndex}); + "Transfer %d: DMA executor may automatically switch to using the destination memory device (%d) not (%d)", + i, t.dsts[0].memIndex, t.exeDevice.exeIndex}); } - } else if (t.dsts[0].memIndex != t.exeDevice.exeIndex) { - errors.push_back({ERR_WARN, - "Transfer %d: DMA executor may automatically switch to using the destination memory device (%d) not (%d)", - i, t.dsts[0].memIndex, t.exeDevice.exeIndex}); } } break; @@ -4311,8 +4326,10 @@ static bool IsConfiguredGid(union ibv_gid const& gid) 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)); @@ -5508,8 +5525,21 @@ static bool IsConfiguredGid(union ibv_gid const& gid) ERR_CHECK(hipSetDevice(exeIndex)); int subIterations = 0; size_t const initOffset = cfg.data.byteOffset / sizeof(float); - float* const src = resources.srcMem[0] + initOffset; - if (!useSubIndices && !cfg.dma.useHsaCopy) { + if (resources.srcMem.empty()) { + // DMA memset: fill each destination via LINEAR_FILL on SDMA DACC BE. + // count is in uint32_t units; value matches MEMSET_VAL byte pattern (0x4B4B4B4B). +#if !defined(__NVCC__) + uint32_t fillVal; + float const f = MEMSET_VAL; + memcpy(&fillVal, &f, sizeof(fillVal)); + size_t const count = resources.numBytes / sizeof(uint32_t); + do { + for (int dstIdx = 0; dstIdx < numDsts; dstIdx++) + ERR_CHECK(hsa_amd_memory_fill(resources.dstMem[dstIdx], fillVal, count)); + } while (++subIterations != cfg.general.numSubIterations); +#endif + } else if (!useSubIndices && !cfg.dma.useHsaCopy) { + float* const src = resources.srcMem[0] + initOffset; if (cfg.dma.useHipEvents) ERR_CHECK(hipEventRecord(startEvent, stream)); @@ -5541,6 +5571,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) #if defined(__NVCC__) return {ERR_FATAL, "HSA copy not supported on NVIDIA hardware"}; #else + float* const src = resources.srcMem[0] + initOffset; // Use HSA async copy do { hsa_signal_store_screlease(resources.signal, numDsts); From d3db0a330398de91a6ba759e340c017aa7c1710c Mon Sep 17 00:00:00 2001 From: nileshnegi Date: Sun, 10 May 2026 09:38:51 -0500 Subject: [PATCH 3/3] Fix DMA memset timing: skip HIP event elapsed time for 0-src transfers hsa_amd_memory_fill does not record HIP events, so querying hipEventElapsedTime after a 0-src DMA transfer produced an "invalid resource handle" error. Guard the HIP event timing path with !resources.srcMem.empty(); the fill path falls back to CPU wall-clock time, which is accurate since hsa_amd_memory_fill is synchronous. Co-authored-by: Claude --- src/header/TransferBench.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index c60f0705..99f43cd6 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -5602,7 +5602,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) if (iteration >= 0) { double deltaMsec = cpuDeltaMsec; - if (!useSubIndices && !cfg.dma.useHsaCopy && cfg.dma.useHipEvents) { + if (!resources.srcMem.empty() && !useSubIndices && !cfg.dma.useHsaCopy && cfg.dma.useHipEvents) { float gpuDeltaMsec; ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent)); deltaMsec = gpuDeltaMsec / cfg.general.numSubIterations;