How DeepEP Utilizes TMA Instructions for Minimal SM Usage on Hopper GPUs

DeepEP offloads bulk tensor movement from Streaming Multiprocessors (SMs) to the Tensor Memory Access (TMA) engine using cp.async.bulk PTX instructions, allowing communication kernels to run on as few as 16-24 SMs while saturating NVLink and RDRAM bandwidth.

DeepEP, DeepSeek-AI's open-source communication library for Mixture-of-Experts (MoE) models, leverages Hopper architecture features to minimize SM occupancy during inter-GPU data transfers. By utilizing Tensor Memory Access (TMA) instructions, the library moves gigabytes of tensor data between global and shared memory without consuming SM compute cycles, enabling the "SM-free" execution mode that leaves the majority of GPU resources available for model computation.

TMA Architecture and SM Offloading

Traditional CUDA kernels perform global memory access using ld.global and st.global instructions that execute directly on SM load/store units. These operations occupy warps and consume valuable instruction issue cycles.

On Hopper (SM90) GPUs, DeepEP employs the Tensor Memory Access (TMA) engine—a hardware unit that performs asynchronous bulk copies between global and shared memory. This engine executes independently of the SM compute pipeline, allowing warps to continue with arithmetic operations (such as GEMM computations) while data movement proceeds in parallel.

Dynamic Shared Memory Configuration

DeepEP configures kernel resources using a compile-time macro that allocates exactly the shared memory required for TMA operations. The SET_SHARED_MEMORY_FOR_TMA macro in csrc/kernels/launch.cuh sets the cudaFuncAttributeMaxDynamicSharedMemorySize attribute before kernel launch:

// From csrc/kernels/launch.cuh
#define SET_SHARED_MEMORY_FOR_TMA(kernel)                                                \
    EP_HOST_ASSERT(cudaFuncSetAttribute(kernel,                                         \
        cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size) == cudaSuccess);        \
    cfg.dynamicSmemBytes = smem_size;

The smem_size value is calculated per-kernel as kNumTMABytesPerWarp * (kNumThreads/32), ensuring each warp receives a dedicated buffer slice without wasting SM memory resources.

Per-Warp TMA Buffer Layout

Each warp operates on a private 8 KB TMA buffer slice allocated within dynamic shared memory. In csrc/kernels/intranode.cu, the layout is defined:

constexpr int kNumThreads = 768;
constexpr int kNumTMABytesPerWarp = 8192;  // 8 KB per warp
constexpr int smem_size = kNumTMABytesPerWarp * (kNumThreads/32);  // 192 KB total

With 24 warps active (kNumThreads/32), the kernel reserves 192 KB of shared memory. This slice holds the source and destination fragments for TMA operations, eliminating register pressure while maintaining coalesced access patterns.

Asynchronous Copy Primitives

DeepEP wraps PTX cp.async.bulk instructions in utility functions defined in csrc/kernels/utils.cuh. The receive loop in csrc/kernels/intranode.cu demonstrates the complete pattern:

// Simplified from intranode.cu receive loop
for (int i = 0; i < 2; ++i) {
    tma_store_wait<0>();                                   // Flush previous copies
    
    if (elect_one_sync()) {                                // Single lane per warp
        // Global → Shared via TMA
        tma_load_1d(tma_buffer,
                    shifted_buffer_x_int4 + i*half_hidden_int4,
                    tma_mbarrier,
                    half_hidden_bytes);
        
        mbarrier_arrive_and_expect_tx(tma_mbarrier, half_hidden_bytes);
        mbarrier_wait(tma_mbarrier, tma_phase);
        
        // Shared → Global via TMA
        tma_store_1d(tma_buffer,
                     shifted_recv_x_int4 + i*half_hidden_int4,
                     half_hidden_bytes,
                     false);
    }
}
__syncwarp();

Key mechanisms include:

• elect_one_sync(): Guarantees exactly one lane per warp issues the TMA instruction, preventing duplicate traffic (implemented in csrc/kernels/utils.cuh, lines 517-528).
• tma_load_1d/tma_store_1d: Thin wrappers emitting cp.async.bulk PTX (lines 387-404 in utils.cuh).
• mbarrier synchronization: Cluster-wide memory barriers track asynchronous completion without SM polling.

Minimal SM Occupancy Configuration

Because TMA handles data movement asynchronously, DeepEP can saturate NVLink/RDMA bandwidth while occupying minimal SM resources. The library exposes this through the Python API:

from deep_ep import Buffer

# Configure only 24 SMs for communication kernels

Buffer.set_num_sms(24)

According to the DeepEP source code, this "SM-free" mode leaves the remaining SMs (up to 132 on H100) available for forward and backward computation, eliminating the resource contention typical of communication-heavy MoE workloads.

TMA Implementation Across Kernel Types

DeepEP applies TMA optimizations consistently across its kernel suite, varying buffer sizes based on data transfer requirements:

Kernel	TMA Buffer Size	Key Source Locations
intranode.cu (NVLink)	8,192 bytes per warp	tma_load_1d at line 470
internode.cu (RDMA)	16,384 bytes per warp	tma_load_1d at line 982, tma_store_1d at line 1137
internode_ll.cu (Low-latency)	sizeof(int4)*32*kNumSendUnrolls	tma_load_1d at line 828

All kernels rely on the SET_SHARED_MEMORY_FOR_TMA macro from csrc/kernels/launch.cuh and the utility functions in csrc/kernels/utils.cuh.

Practical Code Examples

Implementing Direct TMA Copies

The following pattern from csrc/kernels/intranode.cu shows how to implement custom TMA operations:

#include "csrc/kernels/utils.cuh"

__global__ void custom_tma_copy(const float* __restrict__ src,
                               float* __restrict__ dst,
                               int bytes_per_warp,
                               uint64_t* mbar) {
    extern __shared__ uint8_t smem[];
    void* tma_buf = smem + (threadIdx.x / 32) * 8192;

    // Single lane issues TMA load
    if (threadIdx.x % 32 == 0) {
        tma_load_1d(tma_buf, 
                    src + (threadIdx.x / 32) * (bytes_per_warp/4),
                    mbar, bytes_per_warp);
        mbarrier_arrive_and_expect_tx(mbar, bytes_per_warp);
    }
    mbarrier_wait(mbar, 0);

    // Single lane issues TMA store
    if (threadIdx.x % 32 == 0) {
        tma_store_1d(tma_buf,
                    dst + (threadIdx.x / 32) * (bytes_per_warp/4),
                    bytes_per_warp);
    }
}

Compile with -arch=sm_90 and configure shared memory using the SET_SHARED_MEMORY_FOR_TMA pattern before launching.

High-Level DeepEP API Usage

import torch
import torch.distributed as dist
from deep_ep import Buffer

# Initialize with minimal SM footprint

Buffer.set_num_sms(16)
buffer = Buffer(process_group, 0, rdma_bytes=1048576, low_latency_mode=False)

# Dispatch uses TMA internally for transfers

output,idx, weight, _, handle, event = buffer.dispatch(
    hidden_states, 
    topk_idx, 
    topk_weights,
    num_experts=64
)

Summary

• TMA offloading: DeepEP uses Hopper's cp.async.bulk instructions via tma_load_1d and tma_store_1d wrappers to perform asynchronous copies without SM compute cycles.
• Shared memory layout: Each warp receives an 8-16 KB dedicated buffer slice allocated through the SET_SHARED_MEMORY_FOR_TMA macro, calculated as kNumTMABytesPerWarp * (kNumThreads/32).
• Single-lane issuance: The elect_one_sync() primitive ensures exactly one thread per warp issues TMA commands, preventing duplicate transfers.
• Minimal occupancy: The Buffer.set_num_sms() API allows restricting communication kernels to 16-24 SMs, leaving remaining resources for model computation.
• Hardware requirement: TMA optimizations require SM90 (Hopper) architecture; pre-Hopper GPUs fall back to traditional UNROLLED_WARP_COPY pathways.

Frequently Asked Questions

What are TMA instructions in CUDA?

Tensor Memory Access (TMA) instructions are hardware-accelerated copy operations introduced with the NVIDIA Hopper architecture (SM90). They allow asynchronous bulk transfers between global memory and shared memory using dedicated hardware units rather than SM load/store units. In DeepEP, these map to PTX cp.async.bulk instructions wrapped by tma_load_1d and tma_store_1d functions in csrc/kernels/utils.cuh.

Why does DeepEP use only one lane per warp for TMA operations?

DeepEP uses the elect_one_sync() utility to select exactly one active lane per warp because TMA instructions are warp-level operations that do not benefit from multiple threads participating in the same copy. Electing a single lane prevents redundant memory traffic and simplifies barrier synchronization, as only one thread needs to signal the mbarrier arrival while all threads wait on completion.

How much shared memory does DeepEP allocate for TMA buffers?

The allocation scales with thread count and kernel type. For intranode NVLink kernels in csrc/kernels/intranode.cu, DeepEP allocates 8,192 bytes per warp (8 KB), while internode RDMA kernels in csrc/kernels/internode.cu use 16,384 bytes per warp (16 KB). With 768 threads (24 warps), this results in 192 KB or 384 KB of dynamic shared memory per kernel, configured via the SET_SHARED_MEMORY_FOR_TMA macro.

Can DeepEP's TMA optimization work on GPUs older than Hopper (SM90)?

No. The TMA instructions and mbarrier synchronization primitives require SM90 architecture (Hopper) or newer. DeepEP handles this through conditional compilation using the DISABLE_SM90_FEATURES macro; on pre-Hopper GPUs, the library automatically falls back to traditional ld.global/st.global copy loops implemented in UNROLLED_WARP_COPY pathways, though these consume significantly more SM resources and do not support the minimal-SM optimization mode.

{
  "mcpServers": {
    "instagit": {
      "command": "npx",
      "args": ["-y", "instagit@latest"]
    }
  }
}