Skip to content

Fix one-sided remote-read bandwidth inflation via source-slice rotation (SRC_ROTATE_BYTES) - #365

Draft
mustafabar wants to merge 1 commit into
developfrom
users/muabdulj/fix/src-slice-rotation
Draft

mustafabar wants to merge 1 commit into
developfrom
users/muabdulj/fix/src-slice-rotation

Conversation

@mustafabar

Copy link
Copy Markdown
Contributor

Fix one-sided remote-read bandwidth inflation via source-slice rotation (SRC_ROTATE_BYTES)

Motivation

The source buffer is written once (hipMemcpy into srcMem, prep phase) and never modified. The
timing loop reuses that allocation with no invalidation or rotation, so every iteration after the
first re-reads identical addresses and can be served from cache instead of the link under test.

The bias is one-sided: a remote write must be published regardless, but a remote read can
be satisfied out of a cache level shared with the peer. On an 8-partition gfx1250 system in DPX,
pull is inflated up to 4.4×, peaking at 4763 GB/s against a measured steady state of
~1.27 TB/s. That is the side of the push-vs-pull comparison this benchmark is used to decide.

Technical Details

Replicate each GPU source into K identical slices; iteration i reads slice i % K. All slices
hold the same reference data, so validation is unaffected by which slice is read. K is auto-sized:

K = ceil(SRC_ROTATE_BYTES / numBytes)      # default 256MB, 0 disables

  4MB transfer -> K=64 -> 256MB source     (artifact killed)
256MB transfer -> K=1  -> 256MB source     (unchanged, zero cost)

Memory is bounded at ~SRC_ROTATE_BYTES per source, and this is a no-op for transfers already
large enough to defeat the cache.

  • DataOptions::srcRotateBytes — default 256MB, 0 disables.
  • TransferResources — carries numSrcSlices, srcSliceFloats.
  • Allocation — K * numBytes + byteOffset; K decided once per Transfer, since all sources of
    a Transfer share one allocation shape.
  • Prep fill — slices 1..K-1 filled device-to-device from slice 0, so prep H2D stays at
    numBytes.
  • Param assembly — un-rotated src[] snapshotted into subExecSrcBase, keyed off
    rss.subExecIdx rather than assembly order, so it is correct for every GFX_BLOCK_ORDER.
  • RotateSourceSlices() — recomputes src[] from that base and re-uploads the param array
    (a few KB H2D); iteration folded into [0, K) to handle negative warmup indices.
  • Iteration loop — called before cpuStart and the barrier, so rotation is untimed.
  • ExecuteDmaTransfer — DMA reads srcMem directly and already has iteration, so it
    rotates inline.
  • EnvVars.hpp — SRC_ROTATE_BYTES plumbing, help text, banner line.

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; ping-pong (numLaps != 0) is excluded.

15 hunks, +127/−4. Applies cleanly to f8a8d83 (v1.71.01) and 6e5feaa (v1.71.00).

Test Plan

On an 8-partition gfx1250 system in DPX (96 CU / 3 XCC per partition), ROCm 10.2.0:

  • Apply to current upstream, build clean.
  • Confirm mechanism: a 4MB transfer allocates 268,435,456 B (K=64) with prep H2D still 4MB.
  • ALWAYS_VALIDATE=1 with rotation on.
  • Single-pair pull/push sweep 8M–512M, cold vs warm, median of 3 reps, 96 CUs.
  • Same sweep with SRC_ROTATE_BYTES=0 as a control.
  • Repeat on a scale-up (UALoE) pair, i.e. two different physical GPUs.
  • 8-GPU a2a (NUM_SUB_EXEC=8, 64M and 256M), on and off.
  • Cross-check on a 4-partition gfx1250 system in SPX (256 CU / 8 XCC).

Test Result

check result
applies to current upstream f8a8d83 ✅ 15/15 hunks
builds clean ✅
mechanism active ✅ 4MB transfer allocates 268,435,456 B (K=64); prep H2D still 4MB
correctness preserved ✅ ALWAYS_VALIDATE=1, 0 errors with rotation on
removes the artifact ✅ see table
control: OFF still reproduces it ✅ see table
no cost at a2a scale ✅ wall 1.21s both ways; aggregate BW within noise

Sweep at 96 CUs, median of 3 reps. Pull = remote read, G0→G1 (on-die). Steady state ~1270 GB/s:

size OFF cold OFF warm w/c ON cold ON warm w/c
8M 256 637 2.48 294 331 1.13
16M 547 1159 2.12 550 622 1.13
32M 664 1824 (impossible) 2.75 686 712 1.04
64M 925 3226 (impossible) 3.49 964 1013 1.05
128M 1043 1101 1.06 1053 1075 1.02
256M 1164 1159 1.00 1152 1176 1.02
512M 1200 1199 1.00 1189 1201 1.01

With rotation on, warm (1013) ≈ cold (964) — what a correct steady-state measurement looks like.
SRC_ROTATE_BYTES=0 still reproduces the 3.49× inflation on the same build.

Remote write is never inflated and never exceeds that steady state, on or off (push w/c:
1.23 → 0.93 at 16M, 1.16 → 1.00 at 64M, 1.00 → 1.01 at 256M).

Scope: on-die peers only. G0→G1 is on-die — two partitions of the same physical GPU, sharing
PCIe bus 0001:04:00.* — so they share a cache level. G0→G2 crosses scale-up (UALoE) to a
different physical GPU, which is not reachable through that cache level: pull w/c there is
1.01–1.11 across 1M–512M with rotation both on and off, so there is nothing to remove.

Partition mode decides whether the artifact appears. The SPX system does not exhibit it —
stock pull w/c is 1.01–1.30× across 1M–512M. The DPX system shows up to 3.4–4.4×. Same chip
generation, same ROCm; reproduction must match the partition mode.

No cost at a2a scale:

a2a, 8 GPUs, NUM_SUB_EXEC=8 push pull pull/push
64M, rotation OFF 2560 815 0.318
64M, rotation ON 2615 816 0.312
256M, rotation OFF 2618 827 0.316
256M, rotation ON 2556 827 0.324

a2a was never contaminated — each GPU reads from 7 peers, so the working set is already 7 × buffer.
The artifact only distorted single-pair measurements in the 4M–64M range.

Known side effect. Push at 4M/8M drops below 1.0 (0.89, 0.83) with rotation on, because
rotation also denies the push path its local source-read cache reuse. Arguably more correct for
a link-bandwidth benchmark, but small-message push does fall slightly versus stock;
SRC_ROTATE_BYTES=0 restores previous behavior exactly.

Two pre-existing bugs found while validating (NOT caused by this patch)

Both confirmed against a stock build of the same commit. On record only; out of scope here.

  1. Remote write fails validation on the SPX system. (X0->G0->X1) reports ~2,055,873
    mismatched bytes out of 64MB at the destination, with every memory type (U, G, F). The
    mirror pull (X0->G1->X1) is always clean, so push bandwidth is unmeasurable there.
  2. Per-Transfer bandwidth row is wrong on the SPX system. Transfer 0 reports a constant ~4.47 ms
    regardless of size (64MB → 14.998 GB/s) while Executor for the same run reports a sane
    747 GB/s. Use the Executor / Aggregate rows on that build.

Submission Checklist

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) <noreply@anthropic.com>
Copilot AI balanced review requested due to automatic review settings October 23, 2026 01:48

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

Source alignment, cross-rank consistency, byte-count parsing, and subiteration handling need correction, and small-transfer initialization adds substantial overhead.

5 open issues
What changed in this PR

This PR aims to reduce cache-inflated remote-read bandwidth measurements by rotating GPU source buffers between iterations.

Changes:

  • Adds source replication and rotation for eligible GFX and DMA transfers, excluding NIC and ping-pong transfers.
  • Initializes replicas from the original source and updates GFX parameters before timing.
  • Exposes SRC_ROTATE_BYTES with a 256 MiB default, help text, and configuration reporting.
File Description
src/​header/​TransferBench.hpp Adds replica allocation, initialization, and per-iteration source selection.
src/​client/​EnvVars.hpp Adds environment parsing, help text, reporting, and configuration plumbing.

Review effort: Balanced

Comment thread src/client/EnvVars.hpp
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);
// 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);
Comment on lines +5274 to +5277
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);
Comment on lines +7808 to +7812
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);
}
Comment on lines +7995 to +7998
// 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);
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants