Skip to content

Add GPU-initiated and Host-initiated SDMA executors via Anvil - #343

Open
nileshnegi wants to merge 41 commits into
candidate-1.70from
users/nileshnegi/feat/rad-sdma-anvil
Open

Add GPU-initiated and Host-initiated SDMA executors via Anvil#343
nileshnegi wants to merge 41 commits into
candidate-1.70from
users/nileshnegi/feat/rad-sdma-anvil

Conversation

@nileshnegi

@nileshnegi nileshnegi commented Jul 21, 2026 •
edited
Loading

Copy link
Copy Markdown
Collaborator

Motivation

TransferBench can drive SDMA copies from the host (DMA/BMA executors), but has no way to benchmark GPU-initiated SDMA, where a GPU kernel itself builds SDMA packets and rings the KFD queue doorbell.

This PR adds two new SDMA executors inspired by the rocm-xio "anvil" SDMA queue layer:

  • GMA (EXE_GPU_INITIATED_DMA, char code S) => a GPU kernel builds the SDMA packet and rings the KFD doorbell (GPU-initiated).
  • HMA (EXE_HOST_INITIATED_DMA, char code H) => a CPU-driven variant of the same anvil path, where the host builds packets, rings the doorbell, and polls for completion (a correctness/perf reference for the kernel path, and useful in its own right).

Both are measured with the same size sweeps, concurrency, and validation as the existing executors.

The PR also adds the supporting pieces needed to make these paths correct and usable for large transfers (copy chunking, an optional fused copy+signal packet on gfx1250), extends the AllToAll and PodAllToAll presets to select all executors, and addresses reviewer feedback on BMA batching and the vendored anvil layer.

Technical Details

  • New GMA executor (GPU-initiated SDMA)
    • Adds EXE_GPU_INITIATED_DMA (char S) to the executor enum / ExeTypeStr ("CGDINBTSH") / IsGpuExeType dispatch, ExeTypeToStr ("GMA"), and HSA-agent resolution.
    • PrepareAnvilExecutor / RunAnvilExecutor / teardown: one KFD SDMA queue and completion signal per transfer, launched with hipExtLaunchKernelGGL for event-bracketed timing; transfers run concurrently on their own streams.
    • The SDMA kernel is split into four non-branching variants (signal/legacy × wait/no-wait); runtime toggles: ANVIL_WAIT_SIGNAL, ANVIL_LEGACY, ANVIL_REMOTE_DOORBELL, opt-in round-robin SDMA engine selection (ANVIL_ENGINE_ROUND_ROBIN), and an XGMI-optimal SDMA affinity topology table.
  • New HMA executor (host-initiated SDMA)
    • Adds EXE_HOST_INITIATED_DMA (char H), driven by RunAnvilHostExecutor => promoted from the earlier ANVIL_HOST_QUEUE=1 env toggle into a first-class executor.
    • SdmaQueueHostHandle: the host reserves/wrap-pads/places packets into the same host-accessible KFD ring, advances the wptr, rings the doorbell, and polls quiet(). Channels are keyed by {src,dst}; sdma_packets.hpp provides the host-side packet builders. A queue is driven by either the host handle or the device kernel, never both at once.
  • anvil library (src/anvil/)
    • inspired by rocm-xio sdma-ep; gated behind ANVIL_EXEC_ENABLED (CMake -DENABLE_ANVIL_EXEC, Makefile ENABLE_ANVIL_EXEC), linking libhsakmt + libdrm transitively.
  • Chunking (>1 GiB transfers)
    • the SDMA COPY_LINEAR count field is 30-bit (stores size-1), capping one packet at 1 GiB.
    • put_chunked / put_signal_chunked split larger transfers into ≤1 GiB packets, with the signal on the final chunk only.
    • size is controlled by ANVIL_CHUNK_SIZE (raw bytes or K/M/G, clamped to (0, 1 GiB]).
  • OSS7 fused copy+signal (gfx1250)
    • an optional single fused COPY_LINEAR_WAIT_SIGNAL_MI4 packet replacing the separate COPY_LINEAR + ATOMIC pair, selected via ANVIL_FUSED_SIGNAL=1.
    • Two-tier gating: XIO_SDMA_OSS7 (build-level, ABI-portable structs/host code) and XIO_SDMA_OSS7_ENABLED (arch-gates the fused device code to gfx1250/gfx950), so all-arch package builds are safe; the fused path is selected at runtime only on gfx1250.
    • On by default; -DENABLE_XIO_SDMA_OSS7=OFF forces the separate path.
  • AllToAll / PodAllToAll executor selection
    • USE_DMA_EXEC now selects the executor in both presets: 0=GFX, 1=DMA, 2=BMA, 3=GMA, 4=HMA.
    • GMA/HMA (3/4) require -DENABLE_ANVIL_EXEC=ON and are rejected on multi-rank runs (anvil uses process-local virtual addresses, so it cannot reach another rank's memory).
  • BMA batching (addresses review feedback)
    • BMA (USE_DMA_EXEC=2) now batches every copy from a given source GPU into a single hipMemcpyBatchAsync launch by emitting one multi-destination Transfer per source GPU, rather than launching each (src,dst) copy independently.
    • As noted in review, per-link timing can no longer be measured individually; each (src,dst) cell instead shows its equal share of the measured per-source aggregate (rows sum to the source total, diagonal stays N/A). USE_REMOTE_READ is ignored for BMA (the source GPU is always the executor).
  • Correct byte accounting for multi-destination transfers
    • An executor's byte total (and hence CPU-timed bandwidth) now scales by the number of destinations, so a merged BMA batch that writes numBytes to each of K destinations is counted correctly instead of being under-counted by ~K×.
  • Vendored anvil review fixes
    • Fix CHECK_HSAKMT_SUCCESS double-evaluation of its call argument (capture the result once).
    • Cast std::tolower arguments to unsigned char (avoid UB on negative char), add <cctype>.
    • Fix AnvilLib& anvil = anvil.getInstance(); self-reference → AnvilLib::getInstance().
    • Add <iosfwd> for the dump(std::ofstream&) declaration.
  • Diagnostics/tuning
    • ANVIL_CPU_TIMING=1 reports CPU wall time next to the hipEvent time to quantify launch+sync overhead.

Test Plan

  • Build with the anvil executor enabled for gfx1250 (-DENABLE_ANVIL_EXEC=ON -DGPU_TARGETS=gfx1250), default XIO_SDMA_OSS7=ON.
  • Run GMA G→G transfers across a size sweep, including sizes >1 GiB, to exercise the chunking path.
  • Force small ANVIL_CHUNK_SIZE values to validate multi-chunk copies produce a single completion signal and correct data.
  • Validate correctness with FILL_PATTERN + always-validate.
  • Exercise the HMA (host-initiated) executor (USE_DMA_EXEC=4) and the host-queue API unit harness (tools/anvil_harness.cpp).
  • Exercise the fused path (ANVIL_FUSED_SIGNAL=1) on gfx1250 and confirm the non-gfx1250 runtime fallback.
  • Run a2a and poda2a across all executors (USE_DMA_EXEC=0..4) and confirm BMA merged batching produces correct per-source aggregates and CPU-timed bandwidth in line with GPU-timed.

Test Result

  • Clean build for gfx1250 with ENABLE_ANVIL_EXEC=ON and XIO_SDMA_OSS7=ON.
  • GiB transfers complete correctly via chunking (single completion signal preserved); small chunk sizes validated for correctness.
  • Fused copy+signal validated on gfx1250 (64 KiB - 1 GiB), perf-neutral vs the separate path; falls back with a log line on non-gfx1250 runtimes.
  • a2a (4×GPU, 128 MiB) verified across all executors with no GPU faults; representative aggregate bandwidth (GPU-timed / CPU-timed)

Submission Checklist

@nileshnegi
nileshnegi force-pushed the users/nileshnegi/feat/rad-sdma-anvil branch from fc1c16d to ebbb4e2 Compare July 21, 2026 20:14
Copilot AI reviewed Jul 21, 2026
View reviewed changes

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.

Pull request overview

Adds a new GPU-initiated SDMA execution path to TransferBench by vendoring an “anvil”-style SDMA queue layer and wiring it into the existing executor/config infrastructure, enabling benchmarking of kernel-driven SDMA (plus a host-driven reference mode), including large-transfer chunking and optional OSS7 fused copy+signal on supported runtimes.

Changes:

  • Introduces EXE_GPU_INITIATED_DMA (“GISDMA”) executor with prepare/run/teardown integration and env-var runtime toggles (wait/legacy/host-queue/fused/chunking).
  • Vendors src/anvil/ SDMA queue + packet building code (device + host submission paths, OSS7 packet structs, chunking helpers).
  • Adds diagnostics/utilities: SDMA affinity table in topology output and a standalone tools/anvil_harness.cpp for API/packet emission validation; updates Makefile/CMake gating for optional build support.

Reviewed changes

Copilot reviewed 15 out of 16 changed files in this pull request and generated 9 comments.

Show a summary per file
File Description
tools/anvil_harness.cpp New standalone harness to validate anvil host APIs and packet emission (incl. fused-path inspection).
src/header/TransferBench.hpp Adds GISDMA executor type, integrates anvil prepare/run/teardown, and adds SDMA affinity query helper.
src/client/Utilities.hpp Maps new executor type to "GISDMA" string.
src/client/Topology.hpp Prints an SDMA engine affinity table (preferred/available masks per GPU pair).
src/client/Presets/AllToAll.hpp Extends USE_DMA_EXEC to select GISDMA and adds validation/build-flag gating.
src/client/Client.cpp Modifies version printing format string.
src/anvil/sdma-ep.h Adds SDMA endpoint public API and device-side SDMA queue/ops implementation.
src/anvil/sdma_pkt_struct.h Adds pre-OSS7 SDMA packet struct definitions used by anvil paths.
src/anvil/sdma_pkt_struct_mi4.h Adds OSS7/MI4 SDMA packet structs gated by XIO_SDMA_OSS7.
src/anvil/sdma_packets.hpp Host/device packet builders for host-queue submission.
src/anvil/sdma_opcodes.h Shared SDMA opcode/sub-opcode constants (incl. OSS7-gated constants).
src/anvil/anvil.hpp Declares anvil queue types, singleton library, and host-queue handle API.
src/anvil/anvil.cpp Implements KFD/HSA queue creation, engine selection, host-queue submission, and singleton lifecycle.
src/anvil/anvil_device.hpp Device-side compatibility shim + chunking helpers + OSS7 fused-path helpers.
Makefile Adds optional ENABLE_ANVIL_EXEC build path and conditionally links hsakmt/drm deps.
CMakeLists.txt Adds optional ENABLE_ANVIL_EXEC build path and hsakmt/drm detection/linking.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Comment thread src/anvil/anvil.cpp
Comment thread src/anvil/anvil.cpp
Comment thread src/anvil/anvil.cpp Outdated
Comment thread src/anvil/anvil.cpp
Comment thread src/anvil/anvil.hpp
Comment thread src/client/Client.cpp Outdated
Comment thread Makefile
ANVIL_SRCS = ./src/anvil/anvil.cpp
endif

TransferBench: ./src/client/Client.cpp $(ANVIL_SRCS) $(shell find -regex ".*\.\hpp")
Comment thread Makefile
TransferBench: ./src/client/Client.cpp $(ANVIL_SRCS) $(shell find -regex ".*\.\hpp")
$(CXX) $(CXXFLAGS) $(HIPFLAGS) $(COMMON_FLAGS) ./src/client/Client.cpp $(ANVIL_SRCS) -o $@ $(HIPLDFLAGS) $(LDFLAGS)

TransferBenchCuda: ./src/client/Client.cpp $(shell find -regex ".*\.\hpp")
Comment thread CMakeLists.txt
option(ENABLE_AMD_SMI "Enable AMD-SMI pod membership queries" OFF)
option(ENABLE_POD_COMM "Enable pod communication" OFF)
option(BUILD_RELOCATABLE_PACKAGE "Build with RVS-style relocatable RPATH and amdrocm<MAJOR>-transferbench package naming" OFF)
option(ENABLE_ANVIL_EXEC "Enable GPU-initiated SDMA executor via anvil (AMD MI300X only)" OFF)
gilbertlee-amd
gilbertlee-amd reviewed Jul 21, 2026
View reviewed changes
Comment thread src/client/Presets/AllToAll.hpp Outdated
// Collect the number of GPU devices to use
ExeType exeType = useDmaExec ? EXE_GPU_DMA : EXE_GPU_GFX;
ExeType exeType = (useDmaExec == 3) ? EXE_GPU_INITIATED_DMA :
(useDmaExec == 2) ? EXE_GPU_BDMA :

@gilbertlee-amd gilbertlee-amd Jul 21, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

BDMA is a good add, however the correct way for this to be handled will be to actually "batch" together multiple "B" Transfers into a single launch, (which unfortunately will make us lose per Transfer information). We will need to expose this as a new config option for BMA. This is ok as a place holder for now

@nileshnegi nileshnegi Aug 14, 2026

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

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

made some mods in newer commits. pls check

gilbertlee-amd
gilbertlee-amd reviewed Jul 21, 2026
View reviewed changes
Comment thread src/client/Client.cpp Outdated
@nileshnegi
nileshnegi force-pushed the users/nileshnegi/feat/rad-sdma-anvil branch from ebbb4e2 to 85b06c1 Compare July 22, 2026 22:12
@nileshnegi
nileshnegi force-pushed the users/nileshnegi/feat/rad-sdma-anvil branch from 7c48528 to 95f39bb Compare July 31, 2026 04:28
gilbertlee-amd
gilbertlee-amd approved these changes Aug 9, 2026
View reviewed changes
Comment thread src/header/TransferBench.hpp
nileshnegi and others added 19 commits August 13, 2026 20:08
@nileshnegi @claude
feat: add anvil KFD/SDMA queue library headers
0a435c8 
Import sdma_ep.h, anvil.hpp, anvil_device.hpp, and packet struct headers
from rocm-xio verbatim. These provide the GPU-initiated SDMA queue
abstraction (SdmaQueueHandle, AnvilLib, SdmaQueue) that will back the
new EXE_GPU_INITIATED_DMA executor.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
feat: add anvil.cpp — xio-free adaptation of anvil.hip
cb8a576 
Remove the single xio framework dependency from anvil.hip:
replace xio::allocDeviceMemory (XIO_DEVICE_MEM_HIP) with hipMalloc,
and xio::allocDeviceMemory (XIO_DEVICE_MEM_UNCACHED) with
hipExtMallocWithFlags(hipDeviceMallocUncached). Free paths analogously.
No logic changes.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
build: add ENABLE_ANVIL_EXEC cmake option for GPU-initiated SDMA exec…
1b1ec70 
…utor

Add optional ENABLE_ANVIL_EXEC flag (default OFF) that finds libhsakmt,
compiles src/anvil/anvil.cpp, and defines ANVIL_EXEC_ENABLED. Follows
the same pattern as ENABLE_NIC_EXEC and BMA_EXEC_ENABLED.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
feat: add EXE_GPU_INITIATED_DMA enum value (char code 'S')
324d0d9 
Extend ExeType enum with EXE_GPU_INITIATED_DMA=6, update ExeTypeStr to
"CGDINBS", include anvil headers under ANVIL_EXEC_ENABLED guard, and
extend IsGpuExeType to cover the new executor so memory allocation,
peer access, and stream/event setup paths all apply automatically.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
build: fix hsakmt IMPORTED target — use STATIC for .a, add drm transi…
6e4b475 
…tive deps

The found libhsakmt is a static archive (.a), so declare it STATIC IMPORTED
(not SHARED). libhsakmt.a references amdgpu_device_initialize and drmClose
from libdrm_amdgpu and libdrm respectively; add these as INTERFACE_LINK_LIBRARIES
so the linker resolves them transitively.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
feat: add anvilQueues field to ExeInfo for SDMA queue handles
6fd0022 
Store one sdma_ep::SdmaQueueInfo per transfer in ExeInfo, guarded by
ANVIL_EXEC_ENABLED. SdmaQueueInfo holds the host-side KFD queue
descriptor (deviceHandle pointer, src/dst device IDs, channel index)
returned by sdma_ep::createQueue().

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
feat: add TransfersHaveErrors validation for EXE_GPU_INITIATED_DMA
fd0502a 
Validate that the anvil executor has exactly 1 GPU source matching the
executor GPU, at least 1 destination, and same-rank src/dst. Without
ANVIL_EXEC_ENABLED, any 'S' executor in a config produces a clear build-
time error message pointing to the cmake flag.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
feat: implement PrepareAnvilExecutor and extend Teardown for anvil
e8232c7 
PrepareAnvilExecutor: initializes AnvilLib singleton (idempotent via
std::call_once), enables peer access, resolves XGMI-optimal SDMA engine
via getSdmaEngineId, creates one SdmaQueue per transfer via
createSdmaQueue, and stores SdmaQueueInfo in exeInfo.anvilQueues.
SdmaQueue objects are owned by AnvilLib singleton (process lifetime);
TeardownExecutor clears anvilQueues to drop references.

PrepareExecutor and TeardownExecutor extended: EXE_GPU_INITIATED_DMA
included in GPU stream/event create/destroy blocks; HIP events always
created for anvil executor; anvilQueues cleared on teardown.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
feat: implement AnvilTransferKernel and RunAnvilExecutor
26d5483 
AnvilTransferKernel: single-thread __global__ that writes an SDMA
linear-copy packet to the KFD ring buffer via anvil::put(), then
spin-polls for completion via anvil::quiet(). The kernel body is
a no-op for gfx1250 (RDNA4) which uses a different s_waitcnt encoding
incompatible with __builtin_amdgcn_s_waitcnt at ROCm 7.2.

RunAnvilExecutor: launches one kernel per transfer on a shared HIP
stream, brackets with HIP events for wall-clock timing, accumulates
totalDurationMsec. Wired into RunExecutor switch under
ANVIL_EXEC_ENABLED guard.

Also fixes -Wswitch for EXE_GPU_INITIATED_DMA in two additional
switch statements: RecursiveWildcardTransferExpansion (line ~6554)
and GetHsaAgent (line ~7918), adding it alongside the existing GPU
executor cases in each.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
fix: add gfx1250 s_waitcnt compat in sdma-ep.h; remove kernel no-op g…
8218dcd 
…uard

gfx1250 uses s_wait_loadcnt/s_wait_storecnt instead of the legacy s_waitcnt
encoding. Add SDMA_EP_S_WAITCNT() macro that emits inline asm on gfx1250 and
falls back to the builtin on older targets.

Remove the #ifndef __gfx1250__ no-op guard added in Task 8 — AnvilTransfer
Kernel now compiles and runs correctly on all three targets: gfx942,
gfx950, gfx1250.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
fix: make IsGpuExeType unconditional to prevent segfault on gfx950
cfc5fc3 
EXE_GPU_INITIATED_DMA (enum value 6) is always parseable from char 'S'
via CharToExeType regardless of ANVIL_EXEC_ENABLED, because ExeTypeStr
is "CGDINBS" unconditionally. When the define was absent, IsGpuExeType
returned false for value 6, so RecursiveWildcardTransferExpansion skipped
the subindex-assignment block, leaving exeSubIndices empty. Line 6636
then accessed [0] on an empty vector → SIGSEGV.

Fix: include EXE_GPU_INITIATED_DMA in IsGpuExeType unconditionally.
Simplify PrepareExecutor and TeardownExecutor to use IsGpuExeType()
instead of the explicit enum list with ANVIL_EXEC_ENABLED guards.
The event-creation condition is also de-guarded; EXE_GPU_INITIATED_DMA
in the condition is a no-op without the define because that type will
never reach PrepareExecutor's GPU branch on an unsupported build.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
build: add ENABLE_ANVIL_EXEC flag to Makefile
03f4638 
Mirrors the CMake ENABLE_ANVIL_EXEC option. Disabled by default.
When set to 1:
- Detects hsakmt/hsakmt.h and libhsakmt (static or shared)
- Adds -DANVIL_EXEC_ENABLED and -I./src/anvil to COMMON_FLAGS
- Appends src/anvil/anvil.cpp to the build sources
- Links libhsakmt.a with transitive drm deps (or -lhsakmt for .so)

Skipped silently for TransferBenchCuda targets (AMD-only feature).

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
build: pass libhsakmt.a via -Wl, to avoid -x hip parsing it as source
1419b3c 
-x hip applies to all positional arguments on the amdclang++ command
line.  When libhsakmt.a was appended to LDFLAGS without a prefix, the
compiler treated it as HIP source and emitted:

  /opt/rocm/lib/libhsakmt.a:1:1: error: expected unqualified-id
  /opt/rocm/lib/libhsakmt.a:3:4: error: source file is not valid UTF-8

Using -Wl,<path> passes the archive directly to the linker driver,
bypassing the compiler frontend entirely.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
feat: add TB_VERBOSE=1 logging for anvil alloc/dealloc lifecycle
dad5aa7 
anvil.cpp (SdmaQueue ctor/dtor):
- Log entry/exit with GPU src/dst, engine ID, KFD node ID
- Log each hsaKmtAllocMemory/MapMemoryToGPU call with pointer + size
- Log hsaKmtCreateQueueExt result (QueueId, rptr, wptr, doorbell ptrs)
- Log hipMalloc (deviceHandle_) and hipExtMalloc (cachedWptr_,
  committedWptr_) with returned pointers
- Log each dtor step: hsaKmtDestroyQueue, hipFree x3,
  hsaKmtUnmapMemoryToGPU, hsaKmtFreeMemory with pointer/size

TransferBench.hpp (PrepareAnvilExecutor / TeardownExecutor):
- Log resource count and src GPU at prepare entry
- Log per-resource dst GPU, SDMA engine ID, channel index, and
  device handle pointer after queue creation
- Log queue reference release at teardown with per-entry details

All logging is gated on TB_VERBOSE=1; zero overhead when unset.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
fix: add EXE_GPU_INITIATED_DMA to ExeTypeToStr as "GISDMA"
716c1a2 
Missing case caused ExeTypeToStr to return "N/A" for Anvil executors,
showing "Executor: N/A 00" in results output.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
feat: extend USE_DMA_EXEC in AllToAll preset to support BMA and GISDMA
5d10c46 
USE_DMA_EXEC now accepts four values:
  0 = GFX (default, shader-based)
  1 = DMA (CPU-initiated hipMemcpy/HSA)
  2 = BMA (GPU batched DMA)
  3 = GISDMA (GPU-initiated SDMA via anvil/KFD)

All non-GFX modes are restricted to A2A_MODE=0 (copy). Selecting
USE_DMA_EXEC=3 without ANVIL_EXEC_ENABLED at build time emits a clear
error at runtime rather than a confusing crash.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
fix: remove ANVIL_EXEC_ENABLED guard from RecursiveWildcardTransferEx…
062d1e4 
…pansion switch

EXE_GPU_INITIATED_DMA exists in the enum unconditionally, so its switch
case must also be unconditional — the same reasoning that fixed
IsGpuExeType. Without this, non-anvil builds emit:

  warning: enumeration value 'EXE_GPU_INITIATED_DMA' not handled in switch

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
fix: remove ANVIL_EXEC_ENABLED guard from GetHsaAgent switch case
90d287b 
Same pattern as the RecursiveWildcardTransferExpansion fix: the
EXE_GPU_INITIATED_DMA case in GetHsaAgent was guarded by
builds. The case falls naturally into the GPU agent lookup branch which
doesn't use any anvil-specific code, so the guard is unnecessary.

The TransfersHaveErrors switch (line 2545) is correct as-is: it has
both #ifdef and #else branches, so the case label is always present.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
anvil: suppress nodiscard warnings on hipFree in SdmaQueue dtor
a5fffe3 
hipFree is marked nodiscard but we intentionally ignore the return
value in the destructor where error propagation is not possible.

Co-authored-by: Claude <claude@anthropic.com>
nileshnegi and others added 20 commits August 13, 2026 20:10
@nileshnegi @claude
anvil: add HSA preferred-engine selection + SDMA affinity topology table
44068b9 
ANVIL_USE_HSA_ENGINE (default=1) switches getSdmaEngineId() to use
hsa_amd_memory_get_preferred_copy_engine() instead of the hardcoded
MI300X OAM lookup table. The HSA path extracts the lowest-set-bit
engine index from the returned mask; it falls back to the OAM table
if the query fails or returns 0 (no preference). Set
ANVIL_USE_HSA_ENGINE=0 to always use the OAM table.

Add GetSdmaAffinityMask() public API (AMD-only) that wraps both
hsa_amd_memory_copy_engine_status and
hsa_amd_memory_get_preferred_copy_engine for a given GPU pair.

Print a GPU-SDMA affinity matrix in DisplaySingleRankTopology() so
`./TransferBench` (no args) shows the preferred SDMA engine index for
each src->dst GPU pair alongside the available-engine mask.

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @claude
topology: fix SDMA affinity table column alignment
98fd410 
All data cells must be exactly 8 chars wide to match the header
column width (" GPU %02d " = 8 chars, "--------" = 8 chars).

  " %2d[%02x] " was 10 chars → "%2d[%04x]" is 8 chars
  "  [%04x] "  was 9 chars  → " [%04x] "  is 8 chars

Co-authored-by: Claude <claude@anthropic.com>
@nileshnegi @cursoragent
fix: honor byteOffset in DMA and Anvil writes
f174c05 
Apply data.byteOffset to DMA and Anvil source/destination addresses so transfer writes target the same offset window used by data prep and validation.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
anvil: run GISDMA transfers concurrently with per-transfer timing
9d4fc48 
Launch each GPU-initiated SDMA transfer on its own stream and time it
with a dedicated event pair, mirroring the DMA/BMA executors. Previously
all transfers were serialized on a single stream and each was credited
with total_time/N, which produced identical per-link values and a
physically impossible aggregate bandwidth. Executor wall time is now
bounded by the slowest concurrent transfer.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
anvil: add opt-in round-robin SDMA engine selection
d542d0b 
Spread concurrent GISDMA transfers on a source GPU across the set bits of
the HSA preferred-engine mask instead of funneling all onto the lowest
engine. getSdmaEngineId now takes a rotation index (resource index) and
selects the k-th preferred engine.

Gated behind ANVIL_ENGINE_ROUND_ROBIN (default off): neutral on gfx1250
where transfers are bandwidth/latency-bound, but kept as opt-in for
topologies where per-engine contention dominates.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
anvil: use hipExtLaunchKernelGGL for GISDMA kernel timing
151769f 
The Anvil executor is AMD-only (no NVCC path), so launch AnvilTransferKernel
via hipExtLaunchKernelGGL with start/stop events folded into the launch,
matching the GFX executor's timing path and dropping the separate
hipEventRecord calls.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
anvil: use engine-ordered signal for GISDMA completion
ebb9301 
Replace the GISDMA completion (put + quiet, which polls the source ring
read pointer) with put_signal + waitSignal. rptr only reflects ring-space
reclamation and can advance before the copy's writes retire, so quiet
could report completion early; the fused copy + atomic-increment signal is
ordered by the engine after the copy, giving a correct completion barrier.

Add a per-transfer uncached device signal buffer (allocated in
PrepareAnvilExecutor, freed in teardown), reset on the stream before each
launch. This is a fabric-independent ordering fix and does not include
HDP flush / GCR cache ops, which are only needed for non-coherent
host-data-path/aperture consumers (not this XGMI/coherent hardware).

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
anvil: add runtime toggles for GISDMA doorbell placement and completi…
09c87b2 
…on wait

Add ANVIL_REMOTE_DOORBELL (default 0) to place the per-transfer completion
signal in destination memory instead of the source, so the source kernel
remote-polls a doorbell that lands at the same endpoint as the copied data.
Add ANVIL_WAIT_SIGNAL (default 1) to gate the in-kernel waitSignal, allowing
submission-only timing to isolate the completion wait cost. Both are wired
through PrepareAnvilExecutor / AnvilTransferKernel / RunAnvilExecutor.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
anvil: add ANVIL_LEGACY toggle for legacy put/quiet GISDMA completion
0ce43f7 
Add ANVIL_LEGACY (default 0) to switch the GISDMA completion between the
default put_signal + waitSignal (engine-ordered atomic signal) and the
legacy put + quiet (ring rptr poll). Kept for A/B comparison; quiet can
report completion before the copy's writes retire, so the signal path
remains the default. Composes with ANVIL_WAIT_SIGNAL, which still gates
the completion wait in either mode.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
anvil: split GISDMA kernel into four non-branching variants
b61b3e1 
Replace the single AnvilTransferKernel (which branched on doWait/useLegacy
at runtime) with four specialized single-thread kernels selected on the
host: signal+wait, signal+nowait, legacy+wait, legacy+nowait. All share
one signature so the launch site stays uniform, and each combination is a
distinct kernel (no in-kernel branching; distinct entries in profiler traces).

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
anvil: chunk GISDMA copies to lift the 1 GiB per-transfer limit
6fa3874 
The SDMA COPY_LINEAR count field is 30-bit (stores size-1), capping a
single packet at 1 GiB. Split larger transfers into <=1 GiB packets:

- add anvil::put_chunked / put_signal_chunked device wrappers plus
  ANVIL_MAX_COPY_CHUNK and anvil_clamp_chunk in anvil_device.hpp; only
  the final chunk of put_signal_chunked carries the atomic signal so
  waitSignal semantics are unchanged.
- thread a chunkBytes argument through the four GISDMA kernels and the
  launch macro, driven by the ANVIL_CHUNK_SIZE env var (raw bytes or
  K/M/G suffix, clamped to (0, 1 GiB]).
- correct the CreateCopyPacket size assert/comment (1 GiB, not 4 GiB).
- add the ANVIL_CPU_TIMING diagnostic to report CPU wall time alongside
  the hipEvent time (quantifies launch+sync overhead; GB/s unchanged).

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
anvil: add host-initiated (CPU-driven) SDMA queue path
1aa6b4e 
Add a CPU-driven alternative to the GPU-initiated (kernel) SDMA path:
the host builds packets, copies them into the same host-accessible KFD
ring, advances the wptr, and rings the doorbell. A queue is driven by
either the host handle or the device kernel, never both at once.

- SdmaQueueHostHandle (anvil.hpp/anvil.cpp): reserve / wrap-pad / place
  / submit ring mechanics plus put, signal, put_signal, wait_flag_then_
  put, timestamp, and quiet; put/put_signal split into <=1 GiB chunks
  (ANVIL_CHUNK_SIZE) with only the final chunk carrying the flag.
- key SDMA channels by {src,dst} (ChannelKey) instead of dst alone, make
  connect() idempotent, and add connectHost()/getHostHandle() with a
  separate host_sdma_channels_ map.
- sdma_packets.hpp: host-side COPY_LINEAR / ATOMIC / POLL_REGMEM /
  TIMESTAMP packet builders, with the opcodes they need in sdma_opcodes.h.
- TransferBench: ANVIL_HOST_QUEUE=1 selects the host path (one host
  channel per transfer, CPU wall-clock timed).

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
anvil: add OSS7 fused copy+signal SDMA packet (gfx1250)
e533b6c 
Replace the separate COPY_LINEAR + ATOMIC pair with a single fused
copy+atomic-signal packet on hardware that can consume it.

- introduce the two-tier gate: XIO_SDMA_OSS7 (build-level, ABI-portable
  structs/host code) and XIO_SDMA_OSS7_ENABLED (arch-gates the fused
  *device* code to gfx1250/gfx950 so all-arch builds are safe).
- add the compact 12-DWORD COPY_LINEAR_WAIT_SIGNAL_MI4 layout (the
  fixed 19-DWORD form with wait=0 faults gfx1250) plus put_signal_fused
  / put_signal_fused_chunked in anvil_device.hpp, and remove the dead
  anvil:: duplicate of put_signal_counter_impl.
- add AnvilKernelFusedSignal{Wait,NoWait} (bodies arch-gated, symbols
  always emitted) and select them via ANVIL_FUSED_SIGNAL=1, guarded by a
  gfx1250 runtime check that falls back to the separate path otherwise.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
build: enable XIO_SDMA_OSS7 fused packets by default (all-arch safe)
45c36c1 
Add an ENABLE_XIO_SDMA_OSS7 option (CMake and Makefile, on by default)
that defines XIO_SDMA_OSS7=1 for the anvil build. The MI4 structs and
host code are ABI-portable and the fused device code is arch-gated via
XIO_SDMA_OSS7_ENABLED, so this is safe for an all-arch package build;
the fused path is still selected at runtime only on gfx1250. Set
-DENABLE_XIO_SDMA_OSS7=OFF (CMake) or ENABLE_XIO_SDMA_OSS7=0 (make) to
force the separate COPY_LINEAR + ATOMIC path.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
anvil: rename GISDMA executor to GMA and split host path into HMA
3c7f986 
Rename the user-facing GISDMA label to GMA across code and comments
(executor char stays 'S'). Promote the former ANVIL_HOST_QUEUE=1 mode
into a first-class host-initiated executor, HMA (EXE_HOST_INITIATED_DMA,
char 'H'), driven by RunAnvilHostExecutor instead of an env toggle.
AllToAll gains USE_DMA_EXEC=4 to select HMA.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
docs: document GMA/HMA anvil SDMA executors in CHANGELOG
8a36010 
Record the GPU-initiated (GMA) and host-initiated (HMA) anvil SDMA
executors, their runtime toggles, transfer chunking, and the AllToAll
USE_DMA_EXEC selection under v1.70.00.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
anvil: address PR #343 review feedback in vendored SDMA layer
49dae56 
- CHECK_HSAKMT_SUCCESS: evaluate the call argument once into a local so
  side-effecting KFD calls (open/alloc/map/free) are not re-invoked when
  the failure branch prints the status code.
- getBusId: include <cctype> for std::tolower and cast to unsigned char
  to avoid UB on negative char values.
- Fix self-referential global initializer (anvil.getInstance() ->
  AnvilLib::getInstance()).
- anvil.hpp: include <iosfwd> so the dump(std::ofstream&) declaration is
  self-contained.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
a2a: batch each source GPU's BMA copies into one launch
7e82b52 
For USE_DMA_EXEC=2 (BMA), AllToAll now emits one multi-destination
Transfer per source GPU instead of one Transfer per (src,dst) pair, so a
source's copies share a single hipMemcpyBatchAsync launch - the intended
use of the batched API. Because the launch is timed as a whole, per-link
bandwidth can't be resolved individually; each (src,dst) cell reports its
equal share of the measured per-source aggregate (STotal/RTotal still add
up), avoiding empty cells. USE_REMOTE_READ is incompatible (the single
executor must be the source) and is ignored with a warning. Also reject
GMA/HMA on multi-rank runs, since anvil uses process-local addresses.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
core: scale executor byte total by destination count
9fb2b4b 
exeInfo.totalBytes only added a Transfer's numBytes once, ignoring that a
multi-destination Transfer moves numBytes to each destination. This
undercounted totalBytesTransferred (and thus the CPU-timed aggregate
bandwidth) by the destination-count factor for such Transfers - most
visibly the merged BMA all-to-all path, where CPU-timed read ~3x low
against a correct GPU-timed table. Count every destination write so the
CPU-timed figure matches reality.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
poda2a: add BMA/GMA/HMA executor support
4587bd0 
Extend PodAllToAll's USE_DMA_EXEC to the full set (0=GFX, 1=DMA, 2=BMA,
3=GMA, 4=HMA), mirroring AllToAll: executor selection, display labels and
validation (range check plus the ANVIL_EXEC_ENABLED build guard). BMA uses
the merged per-source batched launch with equal-share cells. GMA/HMA are
rejected on multi-rank pods, since anvil's process-local addresses cannot
reach another rank's memory.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi
nileshnegi force-pushed the users/nileshnegi/feat/rad-sdma-anvil branch from 95f39bb to 4587bd0 Compare August 14, 2026 01:12
@nileshnegi
nileshnegi marked this pull request as ready for review August 14, 2026 01:12
@nileshnegi
nileshnegi requested review from a team as code owners August 14, 2026 01:12
nileshnegi and others added 2 commits August 15, 2026 01:13
@nileshnegi @cursoragent
anvil: reuse GMA SDMA queues across re-prepares to fix sweep queue ex…
e4a4222 
…haustion

PrepareAnvilExecutor's GMA path called createSdmaQueue(), which always
appends a new KFD SDMA queue for a {src,dst} pair. Since queues are owned
by the AnvilLib singleton (process lifetime) and teardown only drops the
ExeInfo references, every re-prepare (e.g. each size step of a SWEEP)
leaked another queue per pair. After a handful of steps this exhausted the
KFD SDMA queue slots and hsaKmtCreateQueueExt failed with NO_MEMORY.

Add AnvilLib::getOrCreateSdmaQueue(), an idempotent get-or-create that
reuses the existing channel for {src,dst} and only allocates when a new
channel index is genuinely needed. PrepareAnvilExecutor now tracks a
per-dst channel counter and uses it, mirroring the already-idempotent HMA
connectHost() path. Concurrent transfers to the same dst still get
distinct channels within a Test.

Verified: SWEEP_MAX_POW2=32 with two GMA transfers now runs all 23 size
steps (up to 4 GiB, exercising chunking) with no queue exhaustion.

Co-authored-by: Cursor <cursoragent@cursor.com>
@nileshnegi @cursoragent
anvil: tighten comments across the SDMA executor changes
dcac4d8 
Reduce comment verbosity in the anvil GMA/HMA and BMA-merged code without
changing any logic: condense the executor env-var blocks, kernel-selection
and doorbell notes in TransferBench.hpp, the OSS7 gating/HW-ABI notes in
anvil_device.hpp, the engine-selection and queue-reuse docs in anvil.*, the
BMA-merged notes in the AllToAll/PodAllToAll presets, and the XIO_SDMA_OSS7
option comments in the build files. Comment-only change.

Co-authored-by: Cursor <cursoragent@cursor.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

Copilot code review Copilot
Copilot review effort
Applies to this pull request for everyone.Learn more about Copilot code review.
Copilot left review comments
@gilbertlee-amd gilbertlee-amd gilbertlee-amd approved these changes
@AtlantaPepsi AtlantaPepsi Awaiting requested review from AtlantaPepsi

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

3 participants

@nileshnegi @gilbertlee-amd