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

> Discover how DeepEP leverages TMA instructions for minimal SM usage on Hopper GPUs. Offloads tensor movement to TMA, achieving high bandwidth with fewer SMs.

- Repository: [DeepSeek/DeepEP](https://github.com/deepseek-ai/DeepEP)
- Tags: deep-dive
- Published: 2026-04-25

---

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

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

```cpp
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:

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

```python
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:

```cpp
#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

```python
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.