Add GPU-initiated and Host-initiated SDMA executors via Anvil by nileshnegi · Pull Request #343 · ROCm/TransferBench

nileshnegi · 2026-07-21T20:07:23Z

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 a new GISDMA executor (EXE_GPU_INITIATED_DMA, char code S) inspired by rocm-xio "anvil" SDMA queue layer, so GPU-initiated SDMA can be measured with the same size sweeps, concurrency, and validation as the existing executors. It also adds the supporting pieces needed to make that path correct and usable for large transfers: copy chunking, a host-initiated comparison path, and an optional fused copy+signal packet on gfx1250.

Technical Details

New GISDMA executor
- Adds EXE_GPU_INITIATED_DMA to the executor enum/ExeTypeStr/IsGpuExeType dispatch, ExeTypeToStr ("GISDMA"), 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.
- GISDMA 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, and an
  XGMI-optimal SDMA affinity topology table.
anvil library (src/anvil/)
- inspired by rocm-xio sdma-ep
- anvil.cpp; 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]).
Host-initiated SDMA path
- SdmaQueueHostHandle (CPU builds packets, rings the doorbell, polls quiet()) selected via ANVIL_HOST_QUEUE=1
- used as a correctness/perf reference for the kernel path.
- Channels are keyed by {src,dst}; sdma_packets.hpp provides the host-side packet builders.
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.
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 GISDMA 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 host-queue path (ANVIL_HOST_QUEUE=1) 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.

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.
Host-queue path and anvil_harness API tests pass.
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.

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

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>

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>

…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>

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>

…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>

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>

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>

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>

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>

…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>

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>

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>

-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>

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>

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>

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>

…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>

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>

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>

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>

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>

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>

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>

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>

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>

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>

…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>

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>

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>

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>

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>

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>

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>

Copilot

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.

+#define CHECK_HSAKMT_SUCCESS(call, msg)                                        \
+  do {                                                                         \
+    if ((call) != HSAKMT_STATUS_SUCCESS) {                                     \
+      std::cout << "ERROR code: " << std::dec << call << " " << msg            \
+                << " (File: " << __FILE__ << ", Line: " << __LINE__ << ")"     \
+                << std::endl;                                                  \
+      exit(EXIT_FAILURE);                                                      \
+    }                                                                          \
+  } while (0)


+    uint32_t* queueBuf = static_cast<uint32_t*>(queue_->queueBuffer_);
+    uint64_t offset_dwords = wrapIntoRing(cur_index) / sizeof(uint32_t);
+    for (uint64_t i = 0; i < num_nops; i++) {
+      queueBuf[offset_dwords + i] = SDMA_OP_NOP;
+    }


+  return mi300xOamMap[srcOamId][dstOamId] * 2;
+}
+
+AnvilLib& anvil = anvil.getInstance();


+#include <atomic>
+#include <cstdlib>
+#include <cstring>


+#include <array>
+#include <atomic>
+#include <cstdint>
+#include <initializer_list>


 #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());


+  ANVIL_SRCS = ./src/anvil/anvil.cpp
+endif
+
+TransferBench: ./src/client/Client.cpp $(ANVIL_SRCS) $(shell find -regex ".*\.\hpp")


+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")


 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 · 2026-07-21T22:58:03Z

  // 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 :


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

gilbertlee-amd · 2026-07-21T23:11:27Z

 #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());


gilbertlee-amd and others added 30 commits July 21, 2026 00:51

Fix for using byte offsets with DMA executors (#330)

1e66b46

fix: add EXE_GPU_INITIATED_DMA to ExeTypeToStr as "GISDMA"

8077c00

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>

anvil: suppress nodiscard warnings on hipFree in SdmaQueue dtor

40b97cd

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>

fix: honor byteOffset in DMA and Anvil writes

2dfc295

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 requested review from AtlantaPepsi, Copilot and gilbertlee-amd July 21, 2026 20:09

Copilot started reviewing on behalf of nileshnegi July 21, 2026 20:10 View session

nileshnegi and others added 3 commits July 21, 2026 20:13

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

gilbertlee-amd reviewed Jul 21, 2026

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

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

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

nileshnegi commented Jul 21, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

gilbertlee-amd Jul 21, 2026

Uh oh!

gilbertlee-amd Jul 21, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Uh oh!

Conversation

nileshnegi commented Jul 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

gilbertlee-amd Jul 21, 2026

Choose a reason for hiding this comment

Uh oh!

gilbertlee-amd Jul 21, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

nileshnegi commented Jul 21, 2026 •

edited

Loading