Repository navigation
[rocshmem] SDMA enablement - #5268
Merged
Merged
Conversation
Import SDMA (System DMA) communication channel from Anvil into rocshmem. SDMA offloads memory copy operations to dedicated DMA engines, bypassing GPU shader execution for improved performance on larger transfers. Key features: - Integrate SDMA at IpcOnImpl policy level following rocshmem patterns - Runtime routing based on transfer size (default threshold: 8KB) - Signal-based completion tracking with quiet() synchronization - OAM topology mapping for optimal SDMA engine selection on MI300X New files: - src/sdma/anvil.hpp, anvil.cpp: Host-side queue creation via HSA KMT - src/sdma/anvil_device.hpp: GPU-side queue handle with device functions - src/sdma/sdma_pkt_struct.h: SDMA packet structures - src/sdma_policy.hpp, cpp: SdmaOnImpl/SdmaOffImpl policy classes Environment variables: - ROCSHMEM_SDMA_THRESHOLD: Size threshold in bytes (default: 8192) - ROCSHMEM_SDMA_NUM_CHANNELS: Channels per destination (default: 2) Build with: cmake -DUSE_IPC=ON -DUSE_SDMA=ON .. Original Anvil implementation by Sidler, David Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
Enable concurrent SDMA packet preparation across multiple channels when using wave-level or workgroup-level copy operations. Key changes: - sdmaCopy_wave: Each participating lane handles a different channel, splitting the transfer across up to numChannels queues - sdmaCopy_wg: Multiple threads prepare and submit packets concurrently to different channels - Device handle and signal arrays now 2D: [shm_size * numChannels] - Added minChunkPerChannel threshold to avoid over-parallelization for small transfers New environment variable: - ROCSHMEM_SDMA_MIN_CHUNK_PER_CHANNEL: Minimum bytes per channel for multi-channel transfers (default: 4096) Channel selection logic: - channels_to_use = min(numChannels, size / minChunkPerChannel) - For wave: also limited by number of active lanes - Transfer is split evenly, last thread handles remainder Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
optimization point
- Fix waitForSignal: use >= instead of == to avoid infinite spin if the SDMA engine signal counter overshoots the expected value - Increase SDMA_QUEUE_SIZE from 256KB to 1MB to reduce wraparound stalls under multi-channel load - Fix NOP padding: first DWORD now encodes NOP count per SDMA spec instead of writing all-zeros - Set BREAK_ON_RETRIES=false so hangs are debuggable rather than silently exiting after 2^30 retries - Default numChannels to 1 (envvar + policy in sync) Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com> Signed-off-by: Aurelien Bouteiller <abouteil@amd.com>
- anvil_device.hpp: add maxWritePtr field to SdmaQueueDeviceHandle; submitPacket() now tracks maxWritePtr for rptr-based quiet - Add flushTo()/quietAll() methods to SdmaQueueDeviceHandle - Add proper SdmaQueueSingleProducerDeviceHandle with CAS-free ReserveQueueSpace and simplified submitPacket - Replace putWithSignal with put_signal_counter_impl<> template; add putSignal (rename), putCounter, putSignalCounter, signalCounter free functions; keep putWithSignal as backward-compat alias - Add waitSignal(), waitCounter(), flush(), quiet() free functions using poll_until_ge<> template - anvil.cpp: initialize maxWritePtr in device handle constructor - sdma_policy.hpp/cpp: switch completion tracking from signal/counter arrays to rptr-based anvil::quiet(); remove signalPtrs and expectedSignals allocations Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
- anvil.hpp: add dump() declaration to SdmaQueue; add createSdmaQueue(src, dst, engineId, channelIdx*) to AnvilLib; add initEndpoint() / shutdownEndpoint() free function declarations; add <fstream> include for dump() - anvil.cpp: implement SdmaQueue::dump() — logs wptr/rptr/doorbell/ cachedWptr/committedWptr and decodes COPY/ATOMIC/FENCE packets in the ring buffer with hex addresses - anvil.cpp: implement AnvilLib::createSdmaQueue() — creates a queue with an explicit engineId, bypassing the OAM map lookup; connect() now delegates to createSdmaQueue() internally - anvil.cpp: add s_kfd_opened guard so ~AnvilLib only calls CloseKFD / hsa_shut_down when SetUpKFD actually ran (matches rocm-xio pattern) - anvil.cpp: implement initEndpoint() / shutdownEndpoint() free functions Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
New CMake option USE_SDMA_OSS7 (default OFF) enables CDNA4/MI350X+ packet format support, gated by the preprocessor define USE_SDMA_OSS7. sdma_opcodes.h (new): Consolidates SDMA opcode constants previously inlined in sdma_pkt_struct.h. Adds MI4-specific sub-opcodes and signal/wait operation codes under USE_SDMA_OSS7 guard. sdma_pkt_struct.h: Replace inline opcode definitions with #include "sdma_opcodes.h". sdma_pkt_struct_mi4.h (new): OSS7 packet structs: SDMA_PKT_COPY_LINEAR_WAIT_SIGNAL_MI4 (19 DW, fused copy+wait+signal), SDMA_PKT_COPY_LINEAR_PHY_MI4 (8 DW), SDMA_PKT_FENCE_MI4 (4 DW), SDMA_PKT_FENCE_64B_MI4 (5 DW). All gated on USE_SDMA_OSS7. anvil_device.hpp: Under USE_SDMA_OSS7: add CreateCopyWaitSignalPacketMI4(), CreateFencePacketMI4(), CreateFence64BPacketMI4() device helpers. put_signal_counter_impl<PUT,SIGNAL,COUNTER>: OSS7 fast path fuses copy + signal (or counter) into one COPY_LINEAR_WAIT_SIGNAL_MI4 packet, eliminating the extra ATOMIC packet needed on MI300X. When both signal and counter are active, signal is fused and counter falls back to a separate ATOMIC. Falls through to pre-OSS7 path when USE_SDMA_OSS7 is not set. anvil.cpp: dump() decodes COPY_LINEAR_WAIT_SIGNAL_MI4, FENCE_MI4, and FENCE_64B_MI4 packets under USE_SDMA_OSS7; falls back to pre-OSS7 COPY_LINEAR / FENCE decoders otherwise. CMakeLists.txt: Add option(USE_SDMA_OSS7 ...) under USE_SDMA block; propagate as target compile definition when both USE_SDMA and USE_SDMA_OSS7 are ON. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
…etection OSS7 packet format (MI350X / CDNA4) is now detected per compilation unit via the HIP predefined macro __gfx950__ rather than a CMake flag: - sdma_opcodes.h: define SDMA_IS_OSS7 = 1 when __gfx950__ is set, 0 otherwise; MI4 sub-opcode constants are always defined (no gate needed) - sdma_pkt_struct_mi4.h: MI4 struct definitions are unconditional — safe for all arches, used only when SDMA_IS_OSS7 selects the fast path - anvil_device.hpp: replace #if USE_SDMA_OSS7 with #if SDMA_IS_OSS7 - anvil.cpp dump(): remove preprocessor gates; MI4 packet decoding is now data-driven (checks subop byte from ring buffer), no compile-time arch guard - CMakeLists.txt: remove USE_SDMA_OSS7 option and GPU_TARGETS detection loop Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
…ffImpl) Replace scattered #if defined(USE_SDMA) guards with a clean class hierarchy: - IpcOnImpl: IPC with load/store memcpy, no SDMA awareness - IpcSdmaImpl: extends IpcOnImpl with SDMA-routed copy and quiet - IpcOffImpl: no-op stubs (unchanged) IpcImpl typedef selects: USE_SDMA -> IpcSdmaImpl, USE_IPC -> IpcOnImpl, else -> IpcOffImpl. All #if USE_SDMA removed from context_ipc_device.cpp and IpcOnImpl; unified ipcPut/ipcQuiet/initFrom methods handle dispatch. SdmaOffImpl removed (no longer needed); SdmaOnImpl renamed to SdmaImpl. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
remain correct in mixed use
config. Removed currently useless isSdmaAvailable
abouteiller
changed the title
syncthread in fillbuffer
Without this flush, maxWritePtr is only visible in the GL1 of the CU that submitted the SDMA packet. A thread on a different CU (e.g. the count- sending thread calling sdmaQuietAll in rocshmem_fence) reads maxWritePtr as a plain load from GL2 and may see a stale value, causing quietAll to return prematurely and the signal to fire before SDMA data has arrived.
Fine-grain memory on AMD CDNA is cache-coherent (CC), cached in GL2. The SDMA engine reads source data from GL2, but __syncthreads() in the caller only drains stores to GL0/GL1 without flushing to GL2. Without this agent-scope release fence, SDMA may read stale source data from GL2 instead of the values written by GPU threads. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
…ention The ctx-level sdmaChannel (ctx_id % numChannels) distributes blocks across channels, but kNumWarpsPerGroup=8 warps within each block all call put_nbi_warp to the same dest PE using the same sdmaChannel — causing 8-way intra-block CAS contention on the same SDMA ring buffer. Across 4 blocks per channel this gives 32 concurrent CAS attempts per (channel, dest_PE) ring, producing the observed ~1ms median submit latency. Fix: effective_channel = (sdmaChannel + warp_id_in_wg) % numChannels where warp_id_in_wg = threadIdx.x / 64 (AMD wavefront width). With 8 warps and 8 channels each warp in a block gets a unique ring, reducing per-ring concurrency from 32 to 4 (one warp per block, 4 blocks per channel). sdmaDirtyPECh tracks (pe, channel) pairs so sdmaQuiet correctly drains all channels used for a given dest PE regardless of which warp submitted. No DeepEP changes required. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
ipcFence and ipcQuiet unconditionally called sdmaQuietAll even for pure load/store operations (e.g. getmem_wave/wg below SDMA threshold), paying an agent-scope atomic exchange on every call. Add a relaxed agent-scope atomic load pre-check on sdmaDirtyPECh before entering sdmaQuietAll or sdmaQuiet. The relaxed agent-scope load is coherent through GL2, so it correctly observes dirty bits set by any thread via fetch_or on the same GPU. When dirty=0 (common in pure-IPC workloads), this replaces an unconditional agent-scope RMW with a cheaper agent-scope read and skips the drain entirely.
a per-ctx channle allocation or a spread over all channel policy, the default context always spread on all channels
github-actions
Bot
added
the
documentation
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
pending for future quiet/fences to capture
Yiltan
reviewed
May 14, 2026
edgargabriel
approved these changes
May 18, 2026
miss the pod-like initialization
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
SDMA enablement can improve bandwidth dramatically for low-thread count deployments.
Technical Details
JIRA ID
AIROCSHMEM-322
Test Plan
Submission Checklist