From 26a32d70f2eb24d19312aefb55eeb3dbba569366 Mon Sep 17 00:00:00 2001 From: Mustafa Abduljabbar Date: Thu, 22 Oct 2026 21:48:03 -0400 Subject: [PATCH] Fix one-sided remote-read bandwidth inflation via source-slice rotation The source buffer is written once during prep and never modified, and the timing loop reuses that allocation with no invalidation or rotation. Every iteration after the first re-reads identical addresses and can be served from cache instead of the link under test. Remote writes must be published regardless, so the bias is one-sided and always flatters remote read -- measured up to 4.4x on an 8-partition gfx1250 system in DPX. Replicate each GPU source into K identical slices and have iteration i read slice i % K, where K = ceil(SRC_ROTATE_BYTES / numBytes). All slices hold the same reference data, so validation is unaffected. Memory is bounded at ~SRC_ROTATE_BYTES per source (default 256MB), and the change is a no-op for transfers already large enough to defeat the cache. SRC_ROTATE_BYTES=0 restores the previous behaviour. Scoped to GPU-resident sources on a GFX or DMA executor; the NIC executor registers memory regions against a fixed vaddr and must not be rotated, and ping-pong transfers are excluded. Co-Authored-By: Claude Opus 5 (1M context) --- src/client/EnvVars.hpp | 11 ++++ src/header/TransferBench.hpp | 120 ++++++++++++++++++++++++++++++++++- 2 files changed, 129 insertions(+), 2 deletions(-) diff --git a/src/client/EnvVars.hpp b/src/client/EnvVars.hpp index 122f4dd0..1cc7fd6b 100644 --- a/src/client/EnvVars.hpp +++ b/src/client/EnvVars.hpp @@ -85,6 +85,7 @@ class EnvVars int alwaysValidate; // <0 = disable validation, 0 = validate once after all iterations, >0 = validate after each iteration int blockBytes; // Each subexecutor, except the last, gets a multiple of this many bytes to copy int byteOffset; // Byte-offset for memory allocations + size_t srcRotateBytes; // Min source working-set bytes; replicate+rotate source slices to defeat cache reuse vector fillPattern; // Pattern of floats used to fill source data vector fillCompress; // Percentages of 64B lines to be filled by random/1B0/2B0/4B0/32B0 int sweepMaxPow2; // Maximum power of two to sweep up to when number of bytes to transfer is set to 0 @@ -170,6 +171,7 @@ class EnvVars alwaysValidate = GetEnvVar("ALWAYS_VALIDATE", 0); blockBytes = GetEnvVar("BLOCK_BYTES" , 256); byteOffset = GetEnvVar("BYTE_OFFSET" , 0); + srcRotateBytes = (size_t)GetEnvVar("SRC_ROTATE_BYTES", 256 * 1024 * 1024); fillCompress = GetEnvVarArray("FILL_COMPRESS" , {}); gfxBlockOrder = GetEnvVar("GFX_BLOCK_ORDER" , 0); gfxBlockSize = GetEnvVar("GFX_BLOCK_SIZE" , 256); @@ -393,6 +395,10 @@ class EnvVars printf(" GFX_KERNEL - -1=auto, 0=force GpuReduceKernel, 1=force GpuCopyKernel (may error if ineligible)\n"); printf(" GFX_SE_TYPE - SubExecutor granularity type (0=threadblock, 1=warp)\n"); printf(" GFX_TEMPORAL - Use of non-temporal loads or stores (0=none 1=loads 2=stores 3=both)\n"); + printf(" SRC_ROTATE_BYTES - Min source working-set bytes. GPU sources are replicated into\n"); + printf(" ceil(SRC_ROTATE_BYTES/N) identical slices and each iteration reads a\n"); + printf(" different slice, so repeat iterations cannot hit cache. Defeats the\n"); + printf(" one-sided inflation of remote-read bandwidth. 0 disables. (default 256MB)\n"); printf(" GFX_UNROLL - Unroll factor for GFX kernel\n"); printf(" GFX_SINGLE_TEAM - Have subexecutors work together on full array instead of working on disjoint subarrays\n"); printf(" GFX_WAVE_ORDER - Stride pattern for GFX kernel (0=UWC,1=UCW,2=WUC,3=WCU,4=CUW,5=CWU)\n"); @@ -498,6 +504,10 @@ class EnvVars "Each CU gets a mulitple of %d bytes to copy", blockBytes); Print("BYTE_OFFSET", byteOffset, "Using byte offset of %d", byteOffset); + Print("SRC_ROTATE_BYTES", (int32_t)srcRotateBytes, + "%s", srcRotateBytes == 0 + ? "Source slice rotation OFF (remote-read bandwidth may be cache-inflated)" + : "Replicating GPU sources to >= this many bytes and rotating slices per iteration"); Print("CU_MASK", getenv("CU_MASK") ? 1 : 0, "%s", (cuMask.size() ? GetCuMaskDesc().c_str() : "All")); Print("FILL_COMPRESS", getenv("FILL_COMPRESS") ? 1 : 0, @@ -771,6 +781,7 @@ class EnvVars cfg.data.alwaysValidate = alwaysValidate; cfg.data.blockBytes = blockBytes; cfg.data.byteOffset = byteOffset; + cfg.data.srcRotateBytes = srcRotateBytes; cfg.data.fillCompress = fillCompress; cfg.data.fillPattern = fillPattern; cfg.data.validateDirect = validateDirect; diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index cb991588..790f5822 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -63,6 +63,7 @@ THE SOFTWARE. #include #include #include +#include #include #include "IbvDynLoad.hpp" @@ -266,6 +267,14 @@ namespace TransferBench int validateDirect = 0; ///< Validate GPU results directly instead of copying to host int validateOnDevice = 0; ///< Validate GPU dst/src memory via an on-device kernel instead of host memcmp int validateSource = 0; ///< Validate src GPU memory immediately after preparation + size_t srcRotateBytes = (size_t)256 * 1024 * 1024; + ///< Minimum source working-set size, in bytes. GPU source buffers are + ///< replicated into ceil(srcRotateBytes / numBytes) identical slices and + ///< successive iterations read a different slice, so repeat iterations + ///< cannot be served from cache. Without this, a source buffer is written + ///< once and then re-read unchanged every iteration, which inflates + ///< *remote read* bandwidth by up to 4.4x (writes cannot benefit the same + ///< way, so the bias is one-sided). 0 disables replication. }; /** @@ -3310,6 +3319,8 @@ const auto& AmdSmiFabricInfoV1(const T& info) vector dstActualBytes; ///< Actual amount of dst memory allocated (after padding) vector srcMemHandle; ///< Memory handles for source memory vector dstMemHandle; ///< Memory handles for destination memory + int numSrcSlices = 1; ///< Number of identical source replicas (see DataOptions::srcRotateBytes) + size_t srcSliceFloats = 0;///< Stride in floats between consecutive source replicas vector subExecParamCpu; ///< Defines subarrays for each subexecutor PingpongParam pingpongParamCpu; ///< Pingpong parameter vector subExecIdx; ///< Indices into subExecParamGpu @@ -3494,6 +3505,12 @@ const auto& AmdSmiFabricInfoV1(const T& info) vector resources; ///< Per-transfer resources (normal and ping/pong halves) vector transferRssIdx; ///< Indices into resources of the normal Transfers + // Source-slice rotation (see DataOptions::srcRotateBytes), parallel to subExecParamCpu + vector> subExecSrcBase; ///< Un-rotated src pointers, snapshot after assembly + vector subExecSliceFloats; ///< Per-entry stride between source slices + vector subExecNumSlices; ///< Per-entry slice count (1 = rotation disabled) + bool rotatesSrcSlices = false; ///< True if any entry has >1 slice + // For GPU-Executors SubExecParam* subExecParamGpu; ///< GPU copy of subExecutor parameters PingpongParam* pingpongParamGpu; ///< GPU copy of pingpong parameters @@ -5241,9 +5258,29 @@ const auto& AmdSmiFabricInfoV1(const T& info) rss.srcMem.resize(srcs.size()); rss.srcActualBytes.resize(srcs.size()); rss.srcMemHandle.resize(srcs.size(), NULL); + // Replicate the source into N identical slices so that successive iterations read + // different addresses. The source is written once during prep and never modified, so + // without this every iteration after the first can be served from cache -- which + // inflates remote-read bandwidth only (remote writes must be published regardless). + // Decided once per Transfer, since every source shares one allocation shape. + // Scoped to GPU-resident sources driven by a GFX/DMA executor: the NIC executor + // registers memory regions against a fixed vaddr and must not be rotated. + rss.numSrcSlices = 1; + rss.srcSliceFloats = t.numBytes / sizeof(float); + { + bool allSrcsOnGpu = !srcs.empty(); + for (auto const& s : srcs) + if (!IsGpuMemType(s.memType)) allSrcsOnGpu = false; + if (cfg.data.srcRotateBytes > 0 && t.numBytes > 0 && rss.numLaps == 0 && allSrcsOnGpu && + (exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA)) { + size_t const slices = (cfg.data.srcRotateBytes + t.numBytes - 1) / t.numBytes; + rss.numSrcSlices = (int)std::max((size_t)1, slices); + } + } + for (int iSrc = 0; iSrc < srcs.size(); ++iSrc) { MemDevice const& srcMemDevice = srcs[iSrc]; - size_t srcAllocBytes = t.numBytes + cfg.data.byteOffset; + size_t srcAllocBytes = (size_t)rss.numSrcSlices * 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)); @@ -5516,6 +5553,29 @@ const auto& AmdSmiFabricInfoV1(const T& info) } } + // Snapshot un-rotated source pointers so RotateSourceSlices() can recompute from a fixed + // base. Driven off rss.subExecIdx rather than assembly order, so it is correct for every + // GFX_BLOCK_ORDER (sequential, interleaved and randomized). + { + size_t const numEntries = exeInfo.subExecParamCpu.size(); + exeInfo.subExecSrcBase.assign(numEntries, {}); + exeInfo.subExecSliceFloats.assign(numEntries, 0); + exeInfo.subExecNumSlices.assign(numEntries, 1); + exeInfo.rotatesSrcSlices = false; + for (auto& rss : exeInfo.resources) { + if (rss.numLaps != 0) continue; + for (int k = 0; k < (int)rss.subExecIdx.size(); ++k) { + int const idx = rss.subExecIdx[k]; + if (idx < 0 || idx >= (int)numEntries) continue; + for (int s = 0; s < MAX_SRCS; ++s) + exeInfo.subExecSrcBase[idx][s] = exeInfo.subExecParamCpu[idx].src[s]; + exeInfo.subExecSliceFloats[idx] = rss.srcSliceFloats; + exeInfo.subExecNumSlices[idx] = rss.numSrcSlices; + if (rss.numSrcSlices > 1) exeInfo.rotatesSrcSlices = true; + } + } + } + // Copy sub executor parameters to GPU ERR_CHECK(hipSetDevice(exeDevice.exeIndex)); if (verbose) { @@ -7043,7 +7103,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) ERR_CHECK(hipSetDevice(exeIndex)); int subIterations = 0; size_t const initOffset = cfg.data.byteOffset / sizeof(float); - float* const src = resources.srcMem[0] + initOffset; + // Select this iteration's source slice (see DataOptions::srcRotateBytes). The DMA executor + // reads srcMem directly rather than through SubExecParam, so it rotates here. + int const numSrcSlices = std::max(1, resources.numSrcSlices); + size_t const sliceOffset = (numSrcSlices > 1) + ? (size_t)(((iteration % numSrcSlices) + numSrcSlices) % numSrcSlices) + * resources.srcSliceFloats + : 0; + float* const src = resources.srcMem[0] + initOffset + sliceOffset; if (!useSubIndices && !cfg.dma.useHsaCopy) { if (cfg.general.useHipEvents) ERR_CHECK(hipEventRecord(startEvent, stream)); @@ -7538,6 +7605,39 @@ const auto& AmdSmiFabricInfoV1(const T& info) va_end(args); } + /** + * Point each subexecutor at the source slice for this iteration and re-upload the + * parameter array. Called outside the timed region, so the H2D copy (a few KB) is + * not attributed to transfer bandwidth. No-op unless some Transfer was replicated. + */ + static ErrResult RotateSourceSlices(int iteration, + ExeDevice const& exeDevice, + ExeInfo& exeInfo) + { + if (!exeInfo.rotatesSrcSlices || exeInfo.totalSubExecs == 0) return ERR_NONE; + + for (size_t i = 0; i < exeInfo.subExecParamCpu.size(); ++i) { + int const numSlices = exeInfo.subExecNumSlices[i]; + if (numSlices <= 1) continue; + // iteration is negative during warmup, so fold into [0, numSlices) + int const slice = ((iteration % numSlices) + numSlices) % numSlices; + size_t const offset = (size_t)slice * exeInfo.subExecSliceFloats[i]; + SubExecParam& p = exeInfo.subExecParamCpu[i]; + for (int s = 0; s < p.numSrcs; ++s) + p.src[s] = exeInfo.subExecSrcBase[i][s] + offset; + } + + size_t const numToUpload = std::min((size_t)exeInfo.totalSubExecs, + exeInfo.subExecParamCpu.size()); + if (numToUpload == 0) return ERR_NONE; + ERR_CHECK(hipSetDevice(exeDevice.exeIndex)); + ERR_CHECK(hipMemcpy(exeInfo.subExecParamGpu, + exeInfo.subExecParamCpu.data(), + numToUpload * sizeof(SubExecParam), + hipMemcpyHostToDevice)); + return ERR_NONE; + } + bool RunTransfers(ConfigOptions const& cfg, std::vector const& transfers, TestResults& results) @@ -7702,6 +7802,17 @@ const auto& AmdSmiFabricInfoV1(const T& info) hipMemcpyDefault), errResults); ERR_APPEND(hipDeviceSynchronize(), errResults); + // Fill the remaining source slices with the same reference data. Copying + // device-to-device from slice 0 keeps prep cost off the host interconnect. + // All slices are identical, so validation is unaffected by which slice is read. + for (int slice = 1; slice < resource->numSrcSlices; ++slice) { + ERR_APPEND(hipMemcpy(resource->srcMem[srcIdx] + initOffset + (size_t)slice * resource->srcSliceFloats, + resource->srcMem[srcIdx] + initOffset, + resource->numBytes, hipMemcpyDeviceToDevice), errResults); + } + if (resource->numSrcSlices > 1) + ERR_APPEND(hipDeviceSynchronize(), errResults); + // Optionally validate source memory right after preparation if (cfg.data.validateSource) { if (cfg.data.validateOnDevice && IsGpuMemType(t.srcs[srcIdx].memType)) { @@ -7881,6 +7992,11 @@ const auto& AmdSmiFabricInfoV1(const T& info) } // Wait for all ranks before starting any timing + // Advance to this iteration's source slice before timing starts, so repeat iterations + // read different addresses and cannot be served from cache. + for (size_t i = 0; i < localExecutors.size(); ++i) + ERR_APPEND(RotateSourceSlices(iteration, localExecutors[i], *localExeInfos[i]), errResults); + System::Get().Barrier(); // Start CPU timing for this iteration