Thesis：Engram 是 DeepEP V2 引入的 0 SM RDMA 远端内存访问原语，通过 NCCL Gin (GPU-Initiated Networking) 让 GPU 直接发起 RDMA Get 操作，绕过 CPU 和 NIC 软件栈。其核心设计在于将远端内存抽象为"对称内存窗口"中的全局可寻址表，通过 QP 并行化 + 路由分解 + 原子 flush 控制实现高并发查表，且与模型推理/训练共部署于同一 GPU 进程内。

概述

Engram（Experimental）是 DeepEP V2 中最具创新性的实验特性之一。它实现了 0 SM (Streaming Multiprocessor) 占用的远端内存访问——GPU 通过 RDMA 网络直接读取其他节点的 GPU/CPU 内存，无需占用任何 SM 计算资源进行数据搬运。

核心定位：Engram 面向的场景是 远程 KV Cache 访问——在 Disaggregated Inference（分离式推理）架构中，Prefill 节点的 KV Cache 需要被 Decode 节点远程读取。传统方案需要 CPU 参与（GPU → CPU → NIC → 网络 → NIC → CPU → GPU），而 Engram 允许 GPU 直接通过 RDMA Get 远程读取，实现真正的 GPU-to-GPU 零拷贝。

本报告基于 DeepEP v2.1.0 源码，从以下维度进行深度剖析：

源码架构：Python API → C++ Runtime → CUDA Kernel → JIT 编译四层结构
Gin 交互：NCCL Gin (GPU-Initiated Networking) 的 handle 构造、RDMA Get/Flush 原语
RoCE/NVLink 路径：跨节点 RoCE RDMA 与节点内 NVLink 的分流处理
查表并发：多 QP 并行、路由分解、原子计数与 flush 控制
对称内存：HybridElasticSymmetricMemory 的 CPU 段注册与跨进程映射
SM 调用 RoCE："0 SM" 的本质与实际 SM 占用分析
部署模型：与模型共部署 vs 独立部署的可行性

维度	传统 RDMA	Engram (0 SM)
数据路径	GPU→CPU→NIC→网络	GPU RDMA Get 直读
SM 占用	CPU 参与 + GPU 等待	0 SM（实际 kernel launch 极短）
延迟	高（CPU 中转）	低（GPU-direct）
编程模型	ibv_post_send (CPU)	gin.get (GPU kernel 内)
内存注册	固定 MR	对称内存窗口自动映射

1. 源码架构分层

Engram 的实现贯穿 DeepEP 的四层架构。下图展示完整调用链：

Python API (elastic.py)
    │
    ├─ ElasticBuffer.engram_write(storage, sf)
    │     写入本地 KV 到 CPU 段 (对称内存)
    │
    └─ ElasticBuffer.engram_fetch(indices) → hook
          返回 callable，调用时等待 RDMA 完成
              │
              ▼
C++ Runtime (buffer.hpp)
    │
    ├─ engram_write(): cudaMemcpyAsync → barrier
    │
    └─ engram_fetch(): launch_engram_fetch() → lambda
                          │
                          ▼
JIT Runtime (engram.hpp)
    │
    ├─ EngramFetchRuntime::generate() → 生成 CUDA 代码
    ├─ jit::compiler->build() → 编译
    ├─ EngramFetchRuntime::launch() → 启动 kernel
    │
    └─ EngramFetchWaitRuntime (lambda 被调用时)
           │
           ▼
CUDA Kernel (engram_fetch.cuh)
    │
    ├─ engram_fetch_impl: 每个 QP 一个 block，warp 协作发起 RDMA Get
    └─ engram_fetch_wait_impl: 等待所有 QP 的 RDMA 完成

1.1 Python API 层

Python 层定义在 deep_ep/buffers/elastic.py 的 ElasticBuffer 类中：

class ElasticBuffer:
    @staticmethod
    def get_engram_storage_size_hint(num_entries, hidden, num_max_tokens_per_rank, dtype):
        """计算 Engram 存储所需的最小 buffer 大小"""
        num_bytes_per_entry = align(hidden * dtype.itemsize, 32)  # LDG.256 对齐
        num_gpu_bytes = align(num_bytes_per_entry * num_max_tokens_per_rank, buffer_alignment)
        num_cpu_bytes = align(num_bytes_per_entry * num_entries, buffer_alignment)
        return num_gpu_bytes, num_cpu_bytes

    def engram_write(self, storage, sf=None):
        """写入 Engram 数据到 buffer（含前后 barrier）"""
        self.runtime.engram_write(storage, sf)

    def engram_fetch(self, indices, num_qps=0, use_tma_aligned_col_major_sf=False):
        """发起远端获取，返回 hook callable"""
        return self.runtime.engram_fetch(indices, num_qps, use_tma_aligned_col_major_sf)

关键设计：

get_engram_storage_size_hint：分离 GPU 和 CPU 存储需求。GPU 端只需接收缓冲区（num_max_tokens_per_rank），CPU 端存储全量 entries
engram_write：包含 barrier 确保写入对所有 rank 可见
engram_fetch 返回 callable：延迟执行模式——fetch 只发起 RDMA，返回的 hook 被调用时才 wait 完成

1.2 C++ Runtime 层

C++ 层在 csrc/elastic/buffer.hpp 的 ElasticBuffer 类中实现核心逻辑：

void ElasticBuffer::engram_write(const torch::Tensor& storage,
                                  const std::optional<torch::Tensor>& sf) {
    barrier(false, true);  // 确保之前 fetch 完成

    // 写入到 CPU 段 (buffer + num_gpu_buffer_bytes + offset)
    const auto cpu_write_offset = allow_hybrid_mode
        ? nccl_context->scaleup_rank_idx * num_cpu_buffer_bytes : 0;
    cudaMemcpyAsync(
        advance_ptr(buffer, num_gpu_buffer_bytes + cpu_write_offset),
        storage.data_ptr(), storage.nbytes(),
        cudaMemcpyDeviceToDevice, compute_stream);

    barrier(false, true);  // 确保写入可见
}

engram_fetch 的 C++ 实现：

auto ElasticBuffer::engram_fetch(const torch::Tensor& indices, int num_qps,
                                  const bool& use_tma_aligned_col_major_sf) const {
    // 1. 准备接收 tensor (zero-copy from GPU buffer)
    auto fetched = torch::from_blob(buffer,
        {num_tokens, engram_hidden * num_entries_per_token}, ...);

    // 2. 准备 Gin request 追踪区
    auto last_gin_requests = torch::empty(
        {num_gin_ranks * num_qps, sizeof(ncclGinRequest_t)}, ...);

    // 3. 发起 RDMA Get kernel
    launch_engram_fetch(
        nccl_context->dev_comm, nccl_context->window,
        advance_ptr(buffer, num_gpu_buffer_bytes),  // CPU 段基址
        fetched.data_ptr(),                          // GPU 接收区
        indices.data_ptr<int>(),
        last_gin_requests, ...);

    // 4. 返回 lambda：调用时执行 wait kernel
    return [=, this]() {
        launch_engram_fetch_wait(last_gin_requests, ...);
        return {fetched, fetched_sf};
    };
}

关键洞察：engram_fetch 使用 torch::from_blob 在 GPU buffer 上零拷贝创建输出 tensor。RDMA Get 直接写入 GPU buffer 的前端，无需额外拷贝。

1.2 CUDA Kernel 层

CUDA kernel 在 deep_ep/include/deep_ep/impls/engram_fetch.cuh 中实现。这是 Engram 的核心：

template <int kNumQPs, int kNumEntriesPerRank, int kNumHiddenBytes,
          int kNumSFPacks, int kNumEntriesPerToken, int kNumRDMAPeers,
          int kNumRanksPerRDMAPeer, size_t kIntraRankStorageStride,
          int kNumThreads, typename team_t, int kNumWarps>
__global__ void __launch_bounds__(kNumThreads, 1)
engram_fetch_impl(const ncclDevComm_t nccl_dev_comm, const ncclWindow_t nccl_window,
                  void* storage, void* fetched, int* indices,
                  ncclGinRequest_t* last_gin_requests,
                  const sf_pack_t* sf_table, sf_pack_t* fetched_sf,
                  const int sf_token_stride, const int sf_hidden_stride,
                  const int num_tokens) {
    const auto qp_idx = blockIdx.x;          // 每个 block = 1 个 QP
    const auto warp_idx = ptx::get_warp_idx();
    const auto global_warp_idx = qp_idx * kNumWarps + warp_idx;

    // 构造 Gin handle
    const auto gin = handle::NCCLGin(nccl_dev_comm, nccl_window, qp_idx,
                                      NCCL_GIN_RESOURCE_SHARING_CTA);

    // 每个 warp 处理一个 token 的 entry fetch
    if (ptx::elect_one_sync()) {
        for (int i = global_warp_idx; i < num_tokens * kNumEntriesPerToken;
             i += kNumQPs * kNumWarps) {
            const auto global_idx = __ldg(indices + i);
            const auto owner_rank_idx = global_idx / kNumEntriesPerRank;
            const auto local_entry_idx = global_idx % kNumEntriesPerRank;

            // 路由到 RDMA peer
            const auto peer_idx = owner_rank_idx / kNumRanksPerRDMAPeer;
            const auto intra_peer_rank_idx = owner_rank_idx % kNumRanksPerRDMAPeer;

            // 计算源地址偏移
            const auto src_byte_offset = int64_t(intra_peer_rank_idx) * kIntraRankStorageStride
                                       + int64_t(local_entry_idx) * kNumHiddenBytes;

            // 发起 RDMA Get
            issue_rdma_get(i, peer_idx, src_byte_offset, ...);
        }
    }
}

Wait kernel 在 deep_ep/include/deep_ep/impls/engram_fetch_wait.cuh：

__global__ void __launch_bounds__(kNumThreads, 1)
engram_fetch_wait_impl(const ncclDevComm_t nccl_dev_comm,
                       const ncclWindow_t nccl_window,
                       ncclGinRequest_t* last_gin_requests) {
    const auto qp_idx = blockIdx.x;
    const auto gin = handle::NCCLGin(nccl_dev_comm, nccl_window, qp_idx, ...);

    // 遍历所有 RDMA peer，等待非零 request 完成
    for (int i = thread_idx; i < kNumRDMAPeers; i += kNumThreads) {
        auto last_gin_req_int4 = __ldg(reinterpret_cast<int4*>(
            last_gin_requests + qp_idx * kNumRDMAPeers + i));
        if (last_gin_req_int4.x != 0 || ...) {
            gin.wait(last_gin_req);  // 等待该 peer 的最后一个请求
        }
    }
}

1.4 JIT 编译层

JIT 层在 csrc/kernels/elastic/engram.hpp 中，使用 CRTP 模式：

class EngramFetchRuntime final : public jit::LaunchRuntime<EngramFetchRuntime> {
    struct Args {
        // Templated arguments (编译期常量)
        int num_entries_per_rank, num_hidden_bytes, num_sf_packs;
        int num_entries_per_token, num_scaleout_ranks, num_scaleup_ranks;
        int64_t num_cpu_bytes_per_rank;
        int num_qps;
        bool allow_hybrid_mode;
        // ...
    };

    static std::string generate_impl(const Args& args) {
        // 根据 hybrid 模式选择不同的 team 和 peer 计算
        if (args.allow_hybrid_mode) {
            num_rdma_peers = args.num_scaleout_ranks;
            num_ranks_per_rdma_peer = args.num_scaleup_ranks;
            team_tag = "ncclTeamTagRail";     // scale-out 使用 rail team
        } else {
            num_rdma_peers = args.num_scaleout_ranks * args.num_scaleup_ranks;
            num_ranks_per_rdma_peer = 1;
            team_tag = "ncclTeamTagWorld";    // 全局 world team
        }
        // 生成模板实例化代码
        return fmt::format(R"(
            #include <deep_ep/impls/engram_fetch.cuh>
            static void __instantiate_kernel() {{
                auto ptr = reinterpret_cast<void*>(&{});
            }}
        )", func_name);
    }
};

JIT 设计要点：

QP 数量、hidden size、entries per token 等作为模板参数传入，编译器可针对性优化
Hybrid 模式 vs Direct 模式生成不同的 team_tag 和 peer 计算逻辑
每个 kernel 使用 __launch_bounds__(kNumThreads, 1) 强制每 block 1 个线程块

2. 使用方式

2.1 Engram Write

engram_write 将本地 KV 数据写入对称内存的 CPU 端：

# 计算存储需求
num_gpu_bytes, num_cpu_bytes = deep_ep.ElasticBuffer.get_engram_storage_size_hint(
    num_entries=524288, hidden=128,
    num_max_tokens_per_rank=512, dtype=torch.bfloat16)

# 创建 buffer（CPU 段用于 Engram 存储）
buffer = deep_ep.ElasticBuffer(
    group,
    num_bytes=num_gpu_bytes + num_cpu_bytes,
    num_cpu_bytes=num_cpu_bytes,
    num_allocated_qps=num_qps,
    allow_hybrid_mode=True)

# 写入本地 KV 到 CPU 段
local_kv = torch.randn((num_entries, hidden), dtype=torch.bfloat16, device='cuda')
buffer.engram_write(local_kv)

注意：engram_write 内部调用 cudaMemcpyAsync 从 GPU tensor 拷贝到 CPU 段的对称内存区域。虽然源数据在 GPU 上，但目标是通过 CUDA Unified Memory 映射到 CPU 地址空间的物理内存，NIC 可以直接通过 RDMA 读取。

2.2 Engram Fetch

engram_fetch 发起远端查表：

# indices: [num_tokens, num_entries_per_token] 全局 entry 索引
indices = torch.randint(0, num_ranks * num_entries,
                        (num_tokens, num_entries_per_token), dtype=torch.int32)

# 发起 RDMA Get（非阻塞）
hook = buffer.engram_fetch(indices)

# ... 可在此插入计算 kernel 以 overlap ...

# 等待 RDMA 完成并获取数据
data, sf = hook()  # data: [num_tokens, num_entries_per_token * hidden]

延迟执行（deferred execution）模式允许用户在 issue 和 wait 之间插入计算 kernel，实现通信与计算的 overlap。

2.3 FP8 模式

Engram 支持 FP8 存储 + 全局复制的 scaling factor：

# FP8 模式：数据以 FP8 存储，SF 全局复制
local_fp8, local_sf = per_token_cast_to_fp8(local_bf16)

# SF 全局复制到所有 rank（all-gather）
sf = torch.empty((num_ranks * num_entries, num_sf_packs), ...)
dist.all_gather_into_tensor(sf, local_sf, group)

# 写入：FP8 数据走 RDMA，SF 本地查表
buffer.engram_write(local_fp8, sf=sf)

# 获取：FP8 数据通过 RDMA，SF 通过本地 gather
data, fetched_sf = buffer.engram_fetch(indices)()

设计智慧：FP8 scaling factor 不通过 RDMA 传输，而是在所有 rank 本地冗余存储。这样：

RDMA 只传输 FP8 数据（节省带宽）
SF 查表通过本地 memory read（延迟极低）
避免 RDMA 传输小数据（SF 通常只有 4-16 bytes/entry）的 overhead

3. NCCL Gin 交互机制

3.1 Gin Handle 构造

每个 QP 对应一个 NCCL Gin handle：

const auto gin = handle::NCCLGin(nccl_dev_comm, nccl_window, qp_idx,
                                  NCCL_GIN_RESOURCE_SHARING_CTA);

参数含义：

nccl_dev_comm：NCCL 设备端 communicator
nccl_window：对称内存窗口（所有 rank 可访问的共享内存区域）
qp_idx：QP 索引（每个 block 一个 QP）
NCCL_GIN_RESOURCE_SHARING_CTA：资源分享模式（CTA = Cooperative Thread Array，即 block 内共享）

3.2 RDMA Get 原语

核心 RDMA 操作通过 gin.get 发起：

gin.get<team_t, ncclCoopThread, ncclGin_SegmentMixed>(
    math::advance_ptr(storage, src_byte_offset),  // 源地址（远端对称内存）
    math::advance_ptr(fetched, token_idx * kNumHiddenBytes),  // 目的地址（本地 GPU）
    kNumHiddenBytes,    // 传输字节数
    peer_idx,           // 目标 RDMA peer
    extra_options);     // 聚合选项

模板参数：

team_t：通信组类型（ncclTeamTagWorld 或 ncclTeamTagRail）
ncclCoopThread：协作线程模式（单线程发起，非 warp 协作）
ncclGin_SegmentMixed：混合段模式（源在 CPU 段，目的在 GPU 段）

3.3 Flush 与 Doorbell

RDMA 请求默认聚合（aggregate），达到一定数量后 flush：

// 每 kGinQPFlushDepth 个请求 flush 一次
(request_idx % kGinQPFlushDepth == (kGinQPFlushDepth - 1))
    ? 0   // flush（不聚合）
    : ncclGinOptFlagsAggregateRequests  // 聚合请求

// 最终 flush：ring doorbell 通知 NIC
gin.flush_async<team_t, ncclCoopThread>(i, request_ptr);

Flush 机制：ncclGinOptFlagsAggregateRequests 告诉 Gin 不要立即发送，而是将多个 get 请求聚合后一次性 flush。这减少了 doorbell ringing 次数，提高吞吐。flush_async 将最后一个请求的 handle 写入 last_gin_requests，供 wait kernel 使用。

3.4 gin.get<...> 模板参数

模板参数	值	含义
team_t	ncclTeamTagWorld / ncclTeamTagRail	全局组 vs rail-local 组
CoopLevel	ncclCoopThread	单线程发起操作
Segment	ncclGin_SegmentMixed	源/目的在不同内存段
Options	ncclGinOptFlagsAggregateRequests	请求聚合优化

4. RoCE 与 NVLink 路径处理

4.1 RoCE (RDMA) 路径

Engram 的远程访问完全基于 RDMA（RoCE 或 InfiniBand）：

// 路由计算：owner rank → RDMA peer
const auto peer_idx = owner_rank_idx / kNumRanksPerRDMAPeer;
const auto intra_peer_rank_idx = owner_rank_idx % kNumRanksPerRDMAPeer;

// 源地址在 CPU 段（对称内存）
const auto src_byte_offset = int64_t(intra_peer_rank_idx) * kIntraRankStorageStride
                           + int64_t(local_entry_idx) * kNumHiddenBytes;

// 发起 RDMA Get（跨节点）
gin.get<team_t, ncclCoopThread, ncclGin_SegmentMixed>(
    advance_ptr(storage, src_byte_offset),  // 远端 CPU 段
    advance_ptr(fetched, token_idx * kNumHiddenBytes),  // 本地 GPU
    kNumHiddenBytes, peer_idx, ...);

RoCE 路径特点：

数据从远端节点的 CPU 物理内存（通过 CUDA Unified Memory 映射）直接读取
经过 RDMA NIC（RoCE v2 或 IB）传输
写入本地 GPU VRAM
全程无需 CPU 参与（GPU-Initiated）

4.2 NVLink 路径

Engram 的 engram_fetch kernel 本身不直接处理 NVLink——它只发起 RDMA Get。但 DeepEP 的 barrier 机制区分了 NVLink 和 RDMA：

// comm.cuh: scaleup_barrier 根据拓扑选择路径
if constexpr (kIsScaleupNVLink) {
    // NVLink 路径：使用 LSA (Load-Store Atomic) 屏障
    nvlink_barrier_wo_local_sync<...>(gin, workspace, rank_idx, sm_idx, thread_idx);
} else {
    // RDMA 路径：使用 Gin signal 屏障
    gin_barrier_wo_local_sync<...>(gin.nccl_dev_comm, 1, rank_idx, sm_idx, thread_idx);
}

对于 Engram 的 NVLink 场景（同一节点内 GPU 访问另一 GPU 的 Engram 存储）：

如果 scale-up 域使用 NVLink，engram_write 后的 barrier 使用 NVLink 原子操作
但 engram_fetch 的 RDMA Get 仍然走 RDMA 路径（即使目标在同一节点）
这是因为 Engram 的存储位于 CPU 段，GPU 间 NVLink 无法直接访问 CPU 内存

关键限制：Engram 的远程访问始终走 RDMA，即使目标 rank 在同一 NVLink 域内。这是因为 Engram 数据存储在 CPU 物理内存（对称内存的 CPU 段），而 NVLink 不支持 GPU 直接访问远端 CPU 内存。如果同一节点内 GPU 需要访问对方的 Engram 数据，仍然需要经过 RDMA loopback（RoCE/IB 回环）。

4.3 Hybrid 模式

Hybrid 模式（allow_hybrid_mode=True）改变了 RDMA peer 的路由方式：

if (args.allow_hybrid_mode) {
    // Hybrid: scale-out 每个 rail 一个 peer
    num_rdma_peers = args.num_scaleout_ranks;
    num_ranks_per_rdma_peer = args.num_scaleup_ranks;
    team_tag = "ncclTeamTagRail";  // 使用 rail team
} else {
    // Direct: 每个 rank 一个 peer
    num_rdma_peers = args.num_scaleout_ranks * args.num_scaleup_ranks;
    num_ranks_per_rdma_peer = 1;
    team_tag = "ncclTeamTagWorld";  // 使用 world team
}

模式	num_rdma_peers	num_ranks_per_rdma_peer	team	适用场景
Direct	scaleout × scaleup	1	ncclTeamTagWorld	扁平拓扑
Hybrid	scaleout	scaleup	ncclTeamTagRail	NVLink + RDMA 两层

Hybrid 模式下，一个 RDMA peer 对应一个 scale-out 域（rail），该 rail 内的多个 scale-up rank 通过 NVLink 共享数据。这减少了跨节点 RDMA 的 QP 数量。

5. 查表并发机制

5.1 QP 并行化

Engram 的并发核心是 多 QP (Queue Pair) 并行：

// 每个 QP 对应一个 CUDA block
const auto qp_idx = blockIdx.x;

// 每个 QP 内有多个 warp 并行处理不同 token
const auto global_warp_idx = qp_idx * kNumWarps + warp_idx;

// 循环处理：每个 warp 处理 num_tokens * kNumEntriesPerToken / (kNumQPs * kNumWarps) 个 entry
for (int i = global_widx; i < num_tokens * kNumEntriesPerToken; i += kNumQPs * kNumWarps) {
    // 查表 + RDMA Get
}

并行度分析：

num_qps：默认 = GPU SM 数量（如 H100 = 132 SM），每个 SM 一个 QP
kNumWarps：每个 QP 内的 warp 数（1024 threads / 32 = 32 warps）
总并行度：num_qps × 32 warps 同时发起 RDMA Get

5.2 路由分解

全局 entry index 被分解为 RDMA peer + 本地 entry：

const auto global_idx = __ldg(indices + i);           // 全局索引
const auto owner_rank_idx = global_idx / kNumEntriesPerRank;  // 所属 rank
const auto local_entry_idx = global_idx % kNumEntriesPerRank;  // rank 内偏移

// 路由到 RDMA peer
const auto peer_idx = owner_rank_idx / kNumRanksPerRDMAPeer;
const auto intra_peer_rank_idx = owner_rank_idx % kNumRanksPerRDMAPeer;

// 计算源地址
const auto src_byte_offset = int64_t(intra_peer_rank_idx) * kIntraRankStorageStride
                           + int64_t(local_entry_idx) * kNumHiddenBytes;

查表并发设计：每个 entry 的查表完全独立——通过 __ldg (load global 只读缓存) 读取 index，计算路由，发起 RDMA Get。不同 warp/QP 之间无数据依赖，天然支持最大并发。atomicAdd_block 仅用于统计每个 peer 的请求数以支持 flush 控制。

5.3 原子计数与 Flush 控制

__shared__ int num_requests_per_peer[kNumRDMAPeers];

// 原子计数：追踪每个 peer 的请求数
const auto request_idx = atomicAdd_block(num_requests_per_peer + peer_idx, 1);

// 每 kGinQPFlushDepth 个请求 flush 一次
issue_rdma_get(i, peer_idx, src_byte_offset,
    (request_idx % kGinQPFlushDepth == (kGinQPFlushDepth - 1))
        ? 0   // 不聚合 = flush
        : ncclGinOptFlagsAggregateRequests);

// 最终 flush：确保所有请求被提交
gin.flush_async<team_t, ncclCoopThread>(i, request_ptr);

Flush 策略：

聚合请求减少 doorbell ringing 开销
每 kGinQPFlushDepth 个请求强制 flush 一次，防止请求无限堆积
最终无条件 flush 确保所有请求被 NIC 处理

6. 对称内存定义

6.1 内存布局

Engram 使用 HybridElasticSymmetricMemory（定义在 csrc/kernels/backend/symmetric.hpp）：

Virtual Address Space (per rank):
┌──────────────────────────────────────────────────────────────┐
│  [Workspace]  │  GPU Buffer (front)  │  CPU Buffer (back)   │
│  (2MB align)  │  (dispatch/combine   │  (Engram storage)     │
│               │   + fetch recv)      │                       │
└──────────────────────────────────────────────────────────────┘
                                      │
                                      ▼
                            CPU Buffer 内部布局:
                            ┌─────────┬─────────┬─────┬─────────┐
                            │ CPU     │ CPU     │     │ CPU     │
                            │ rank 0  │ rank 1  │ ... │ rank N-1│
                            └─────────┴─────────┴─────┴─────────┘
                            (每个 scale-up rank 一段)

关键代码：

// HybridElasticSymmetricMemory 构造
this->num_bytes = num_gpu_bytes + num_cpu_bytes * num_scaleup_ranks;

// Reserve VA
cuMemAddressReserve(&addr, num_bytes, kNumAlignmentBytes, 0, 0);

// Map GPU segment (front)
cuMemMap(addr, num_gpu_bytes, 0, gpu_handle, 0);

// Import and map CPU segments (back) - 每个 scale-up rank 一段
for (int i = 0; i < num_scaleup_ranks; ++i) {
    const auto offset = num_gpu_bytes + i * num_cpu_bytes;
    cuMemImportFromShareableHandle(&cpu_handles[i], local_fd, ...);
    cuMemMap(addr + offset, num_cpu_bytes, 0, cpu_handles[i], 0);
}

6.2 CPU 段注册与跨进程映射

CPU 段的创建使用 CUDA Virtual Memory Management：

// 1. 创建 NUMA-local CPU 分配
static cpu_handle_t create_cpu_handle(int64_t num_cpu_bytes) {
    DeviceContext ctx;
    auto cpu_prop = ctx.cpu_alloc_prop();  // NUMA-local 属性

    CUmemGenericAllocationHandle handle;
    cumem_create_with_fallback(&handle, num_cpu_bytes, &cpu_prop);

    // 2. 导出为 POSIX FD（跨进程共享）
    int fd = -1;
    cuMemExportToShareableHandle(&fd, handle,
        CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR, 0);

    return {getpid(), fd};
}

// 3. 跨进程导入（通过 pidfd_open + pidfd_getfd）
int pidfd = syscall(SYS_pidfd_open, pid, 0);
int local_fd = syscall(SYS_pidfd_getfd, pidfd, fd, 0);
cuMemImportFromShareableHandle(&cpu_handles[i], (void*)(uintptr_t), local_fd, ...);

跨进程内存共享机制：每个 rank 创建自己的 CPU 段，通过 POSIX file descriptor 导出。其他 rank 通过 Linux pidfd API 获取 FD 并 import 到自己的 VA 空间。这允许 NIC 通过 RDMA 直接访问远端进程的 CPU 内存。

6.3 Engram 存储的对称内存语义

Engram 的"对称内存"含义：

地址对称：每个 rank 的 CPU 段在各自 VA 空间中的偏移相同（num_gpu_buffer_bytes + rank_idx * num_cpu_buffer_bytes）
内容不对称：每个 rank 的 CPU 段只存储自己的 local entries
全局可寻址：通过 ncclWindow_t，任何 rank 可以计算远端 rank 的 CPU 段地址并发起 RDMA Get

Engram 的存储布局（CPU 段内）：

CPU segment for rank r:
┌────────────────────────────────────────────────────────┐
│ entry 0  │ entry 1  │ ...  │ entry (num_entries-1)     │
│ hidden*B │ hidden*B │      │ hidden*B                  │
└────────────────────────────────────────────────────────┘

// 全局视角（所有 rank 的 CPU 段拼接）:
┌──────────────┬──────────────┬─────┬──────────────┐
│ rank 0       │ rank 1       │ ... │ rank N-1     │
│ entries      │ entries      │     │ entries      │
└──────────────┴──────────────┴─────┴──────────────┘

// 全局索引 → (rank, local_entry):
owner_rank = global_idx / num_entries_per_rank
local_entry = global_idx % num_entries_per_rank

7. HybridElasticSymmetricMemory 深度专题

专题核心问题：HybridElasticSymmetricMemory 为什么要设计成"对称"架构？这种对称性为 Engram 的 GPU 单向查表式内存访问带来什么本质好处？从 RoCE 相关代码、测试代码和初始化操作三个维度分析，对称内存到底能否跨越节点，还是仅限于节点内或 Engram 域内？

7.1 为什么要设计对称架构

"对称"的精确定义：在 DeepEP 中，"symmetric"不是指数据内容对称（各 rank 存储不同数据），而是指虚拟地址布局对称——每个 rank 的 NCCL Window 具有完全相同的 VA 结构：

// 所有 rank 的窗口布局完全相同（nccl.cu L140）
ncclCommWindowRegister(comm, raw_window_ptr, sym->num_bytes, &window, ...);

// 每个 rank 的 VA 布局:
┌─────────────────────────────────────────────────────────────────┐
│  [GPU VRAM (num_gpu_bytes)]  │  CPU Segments (num_cpu_bytes)   │
│  所有 rank 大小相同          │  所有 rank 大小相同              │
└─────────────────────────────────────────────────────────────────┘

对称架构的核心好处：O(1) 全局地址计算。

关键洞察：因为所有 rank 的窗口布局对称，任意 rank 可以在本地计算出任意其他 rank 的存储偏移量，无需查询、无需通信、无需地址转换：

// engram_fetch.cuh L57-66 — 本地计算远端地址
const auto owner_rank_idx = global_idx / kNumEntriesPerRank;
const auto local_entry_idx = global_idx % kNumEntriesPerRank;
const auto peer_idx = owner_rank_idx / kNumRanksPerRDMAPeer;
const auto intra_peer_rank_idx = owner_rank_idx % kNumRanksPerRDMAPeer;
const auto src_byte_offset = int64_t(intra_peer_rank_idx) * kIntraRankStorageStride
                           + int64_t(local_entry_idx) * kNumHiddenBytes;

这个偏移量在任何 rank 上计算结果都相同——对称性保证了偏移量的全局一致性。

对比非对称方案：

方案	地址计算	开销	RDMA 友好
非对称（各 rank 布局不同）	需要查询远端布局信息	O(通信) 或 O(查表)	否
对称（DeepEP 采用）	本地 O(1) 计算	零	是
全局统一 VA（CXL 风格）	直接指针访问	需要硬件支持	部分

7.2 对称性对 Engram 单向查表的意义

Engram 的 GPU 单向内存访问（GPU 主动读远端，远端 GPU 不参与）依赖对称性实现三个关键能力：

7.2.1 无状态路由

GPU kernel 计算出的 src_byte_offset 直接传递给 NCCL Gin：

gin.get<team_t, ncclCoopThread, ncclGin_SegmentMixed>(
    math::advance_ptr(storage, src_byte_offset),  // 本地 VA 计算的偏移
    math::advance_ptr(fetched, token_idx * kNumHiddenBytes),  // 本地 GPU 目的
    kNumHiddenBytes, peer_idx, ...);


关键机制：gin.get 内部将 src_ptr 转换为窗口偏移（src_ptr - lsa_base_ptr），然后告诉 NIC："从 peer_idx 的窗口偏移 XX 处读取数据"。因为所有 peer 的窗口布局对称，本地计算的偏移量在远端同样有效。


7.2.2 单向性（One-Sided）保证

远端 GPU 完全不参与数据传输：

Node A (发起端)                    Node B (远端)
┌─────────────────┐                ┌─────────────────┐
│ GPU Kernel      │                │ GPU 空闲        │
│ 计算 src_offset │                │ (不参与)        │
│     ↓           │                │                 │
│ NIC (A)         │══ RDMA Get ═══▶│ NIC (B)         │
│                 │                │ 从本地 CPU 段   │
│                 │                │ 读取物理内存    │
│     ↓           │                │                 │
│ GPU VRAM (A)    │◀═══════════════│                 │
└─────────────────┘                └─────────────────┘

对称性使得这种单向操作成为可能——发起端不需要远端的任何信息（地址、状态、同步），仅凭全局索引即可定位数据。

7.2.3 对称性消除了 CPU 查表需求

在传统 RDMA 编程中，CPU 需要维护地址翻译表（MR 查找、lkey/rkey 管理）。对称架构下：


偏移量 = 地址（无需 MR 查找）
窗口基址 = 常量（lsa_base_ptr）
远端 key = 隐式（peer_idx → NCCL 内部管理 rkey）


7.3 对称内存是否跨越节点：三维证据分析

这是最核心的问题。我从三个维度进行分析：

7.3.1 维度一：RoCE 相关代码证据

证据 1：Hybrid 模式的 RDMA peer 覆盖所有 scale-out 节点

// engram.hpp L44-47 — Hybrid 模式
if (args.allow_hybrid_mode) {
    num_rdma_peers = args.num_scaleout_ranks;  // 所有跨节点 peer
    num_ranks_per_rdma_peer = args.num_scaleup_ranks;
    team_tag = "ncclTeamTagRail";  // Rail team 覆盖所有节点
}

证据 2：NCCL 窗口注册是全局 collective

// nccl.cu L135-140 — 窗口注册跨越所有 rank
// NOTES: `ncclCommWindowRegister` is collective: it internally calls
// bootstrapBarrier across all ranks, so no explicit barrier is needed.
NCCL_CHECK(ncclCommWindowRegister(comm, raw_window_ptr,
    this->symmetric_memory->num_bytes, &window, NCCL_WIN_STRICT_ORDERING));

证据 3：Gin get 的 peer_idx 可以是跨节点 peer

// engram_fetch.cuh L61-62 — peer_idx 路由到任意 RDMA peer
const auto peer_idx = owner_rank_idx / kNumRanksPerRDMAPeer;
// peer_idx 可以是 0 到 num_scaleout_ranks-1（跨节点）

证据 4：CPU handle 通过 all_gather 全局分发

# elastic.py L338-342 — 全局分发（跨节点）
cpu_comm = []
if allow_hybrid_mode and num_cpu_bytes > 0:
    pid, fd = _C.create_cpu_handle(num_cpu_bytes)
    cpu_comm = [None] * self.num_ranks
    dist.all_gather_object(cpu_comm, (pid, fd), self.group)  # 全局 all_gather

7.3.2 维度二：测试代码证据

测试使用全局 group（跨节点）：

# test_engram.py L15,124
rank, num_ranks, group = init_dist(local_rank, num_local_ranks)
# init_dist 返回的 group 是全局 group:
# dist.new_group(list(range(num_local_ranks * num_nodes)))

# 测试默认 4 个进程（可跨节点）
torch.multiprocessing.spawn(test, args=(num_processes, args), nprocs=num_processes)

# indices 范围覆盖所有 rank 的 entries:
indices = torch.randint(0, num_ranks * args.num_entries, ...)  # 全局索引

测试的正确性验证：通过 dist.all_gather_into_tensor 收集所有 rank 的数据作为 reference，验证 Engram fetch 结果。这意味着测试设计本身就假设了跨节点数据访问。

7.3.3 维度三：初始化操作证据

证据 1：VA 空间分配只包含本地节点

// symmetric.hpp L211 — 只分配本地节点的 CPU 段
this->num_bytes = num_gpu_bytes + num_cpu_bytes * num_scaleup_ranks;
// 注意：不是 num_ranks，而是 num_scaleup_ranks（本地节点 rank 数）

证据 2：只 import 本地节点的 CPU segments

// symmetric.hpp L223-226 — 仅映射本地节点
// Import and map all intra-node CPU segments
for (int i = 0; i < num_scaleup_ranks; ++i) {
    auto [pid, fd] = cpu_comm[num_scaleup_ranks * scaleout_rank_idx + i];
    // scaleout_rank_idx 选择当前节点的 segments
}

证据 3：set_access 只设置本地设备访问权限

// symmetric.hpp L104-109 — 仅授予本地 GPU 访问
for (int i = 0; i < num_devices; ++i) {
    int can_access_peer = 0;
    if (i == device_idx ||
        (cudaDeviceCanAccessPeer(&can_access_peer, i, device_idx) && can_access_peer)) {
        desc[0].location.id = i;
        cuMemSetAccess(addr, num_bytes, desc, ...);
    }
}
// 跨节点 GPU 无法 peer access，不授予权限

证据 4：LSA pointer 只获取节点内 peer

// nccl.cu L141-147 — 仅 NVLink peer 的 LSA pointer
ncclGetLsaDevicePointer(window, 0, nvl_rank_idx, &mapped_window_ptr);
for (int i = 0; i < num_nvl_ranks; ++i)  // 仅节点内 NVLink peers
    ncclGetLsaDevicePointer(window, 0, i, &nvl_window_ptrs[i]);

7.3.4 综合结论：物理本地 + 逻辑全局


结论：HybridElasticSymmetricMemory 是"物理本地、逻辑全局"的架构：

物理层：CPU 段位于各节点本地（NUMA-local），仅本地节点 GPU 可直接访问
逻辑层：通过对称布局，所有节点使用相同的 VA 方案，偏移量全局有效
访问层：跨节点访问通过 RDMA 实现——本地 NIC 读取远端节点的本地 CPU 段



完整的数据流（跨节点场景）：

Node 0 (发起 fetch)                           Node 1 (远端)
┌────────────────────────┐                    ┌────────────────────────┐
│ GPU Kernel (Node 0)    │                    │ GPU 空闲               │
│  计算: peer_idx=1      │                    │ (完全不参与)           │
│  src_off = 3*stride+off│                    │                        │
│         ↓              │                    │                        │
│ NIC (Node 0)           │                    │ NIC (Node 1)           │
│  发起 RDMA Get         │═══ RoCE/IB ═══════▶│  接收 Get 请求         │
│  peer=1, offset=XXX    │                    │  从本地 CPU 段读取     │
│                        │                    │  (物理内存 DMA)        │
│                        │◀═══════════════════│                        │
│  接收数据              │                    │                        │
│         ↓              │                    │                        │
│ GPU VRAM (Node 0)      │                    │                        │
│  fetched[...] 就绪     │                    │                        │
└────────────────────────┘                    └────────────────────────┘

关键: NIC 读取的是 Node 1 的本地 CPU 物理内存
     Node 1 的 GPU 完全不参与
     对称性保证: Node 0 计算的 offset 在 Node 1 同样有效

"Engram 域"的定义：


层级 范围 访问方式 延迟
GPU 段 本地节点所有 GPU NVLink peer access ~10 µs
CPU 段（本地） 本地节点所有 rank 本地 GPU 不可直接访问，RDMA loopback ~1-2 µs
CPU 段（远端） 所有 scale-out 节点的 CPU 段 RDMA Get (RoCE/IB) ~5-15 µs



重要澄清：Engram 的"对称内存"不是"跨节点共享内存"（不是 CXL 内存池），而是"对称地址布局的分布式本地内存"。每个节点只存储自己的数据，但通过 RDMA 网络让其他节点可以读取。对称性体现在地址计算层面，而非数据物理位置。


7.4 设计权衡总结


设计决策 选择 原因
对称布局 所有 rank VA 相同 O(1) 地址计算，无需查表
物理分布 各节点存各节点数据 避免一致性开销，最大化本地带宽
跨节点访问 RDMA one-sided read 远端 GPU 不参与，真正的 0 SM
CPU 段注册 NUMA-local + GDR RDMA NIC 可直接 DMA 读取
全局窗口 ncclCommWindowRegister NCCL 内部管理跨节点 MR


对称性带来的根本优势：它将"远端地址计算"从 O(通信) 或 O(查表) 降为 O(1) 算术运算。GPU kernel 只需一次除法和一次取模即可定位任意 rank 的任意 entry——这是 GPU-Initiated RDMA 能够高效工作的前提。

8. SM 如何调用 RoCE 能力

7.1 "0 SM" 的本质

Engram 被称为"0 SM"通信，但这需要精确理解：


"0 SM" 的准确含义：Engram 的 RDMA Get 操作不长期占用 SM。kernel 启动后，SM 只执行少量指令（计算地址、发起 RDMA doorbell），然后 kernel 结束。RDMA 数据传输由 NIC 异步完成，无需 SM 参与。这与传统通信 kernel（如 NCCL）需要 SM 持续参与数据搬运完全不同。


7.2 NCCL Gin GPU-Initiated API

SM 通过 NCCL Gin 的 GPU 端 API 直接操作 RDMA：

// 1. 构造 Gin handle（绑定到特定 QP）
const auto gin = handle::NCCLGin(nccl_dev_comm, nccl_window, qp_idx,
                                  NCCL_GIN_RESOURCE_SHARING_CTA);

// 2. 发起 RDMA Get（GPU 指令，直接写 NIC doorbell）
gin.get<team_t, ncclCoopThread, ncclGin_SegmentMixed>(
    src_ptr, dst_ptr, size, peer_idx, options);

// 3. Flush（ring doorbell）
gin.flush_async<team_t, ncclCoopThread>(peer_idx, request_ptr);

// 4. Wait（轮询 completion）
gin.wait(last_gin_req);

这些 API 在 GPU kernel 内执行，直接操作 NIC 的 doorbell register，绕过 CPU 和 NIC driver。

7.3 实际 SM 占用分析

虽然称为"0 SM"，但实际仍有短暂占用：


阶段 SM 占用 持续时间 说明
engram_fetch kernel num_qps 个 block ~μs 级 计算地址 + 发起 RDMA
engram_fetch_wait kernel num_qps 个 block 等待 RDMA 完成 轮询 completion
数据传输 0 SM 取决于数据量 NIC DMA 异步完成



实际占用：engram_fetch kernel 使用 __launch_bounds__(1024, 1)，每个 QP 一个 block（1024 threads）。对于 H100（132 SM），如果 num_qps=132，则短暂占用全部 SM，但 kernel 执行时间极短（微秒级）。wait kernel 类似。真正的数据传输期间 SM 完全空闲。


9. Engram 部署模型

8.1 与模型共部署

当前 Engram 的实现是与模型共部署的：

# 模型进程内创建 ElasticBuffer（包含 Engram 支持）
buffer = deep_ep.ElasticBuffer(
    group,
    num_bytes=num_gpu_bytes + num_cpu_bytes,
    num_cpu_bytes=num_cpu_bytes,  # CPU 段用于 Engram
    ...)

# 模型推理/训练中使用 Engram
# Prefill 节点：写入 KV Cache 到 Engram
buffer.engram_write(kv_cache)

# Decode 节点：从 Engram 读取远端 KV
data, _ = buffer.engram_fetch(indices)()

共部署的特点：

ElasticBuffer 与模型在同一进程、同一 GPU
对称内存窗口由 NCCL communicator 管理
Engram 的 CPU 段与模型的 GPU buffer 共享同一 VA 空间
无需跨进程通信，直接函数调用


8.2 ElasticBuffer 生命周期


初始化:
  1. 创建 NCCL communicator (ProcessGroup)
  2. 计算 buffer 大小 (GPU + CPU)
  3. 创建 HybridElasticSymmetricMemory
     - Reserve VA space
     - Map GPU segment (front)
     - Import + map CPU segments from all scale-up ranks
  4. 初始化 workspace (barrier 等)

使用:
  - engram_write(): GPU tensor → CPU 段 (cudaMemcpyAsync)
  - engram_fetch(): 发起 RDMA Get → 返回 hook
  - hook(): 等待 RDMA 完成 → 返回数据

销毁:
  - barrier() 确保所有操作完成
  - 释放 host workspace
  - 销毁 NCCL symmetric memory context
  - Unmap + release GPU/CPU segments
  - Free VA space

8.3 独立部署可行性分析

Engram 能否作为独立服务部署（类似 KV Cache 池化）？


部署模式 当前状态 可行性 挑战
与模型共部署 ✅ 当前实现 完全支持 无
独立进程（同 GPU） ❌ 未实现 可行但复杂 需要 IPC 机制传递 indices
独立 GPU（分离式） ❌ 未实现 理论上可行 需要独立 NCCL group
CPU-only 节点 ❌ 不支持 需要重新设计 Engram 需要 GPU 发起 RDMA



独立部署的关键障碍：

NCCL Communicator 绑定：Engram 依赖 NCCL communicator 管理对称内存窗口，独立部署需要独立的 NCCL group
地址空间隔离：engram_fetch 直接在 GPU kernel 内计算远端地址，独立进程需要跨进程 VA 映射
同步机制：当前使用 dist.barrier() 进行 CPU 同步，独立部署需要替代同步方案
内存注册：CPU 段的跨进程共享依赖 POSIX FD + pidfd，独立部署需要额外的服务发现机制



10. 测试方法

9.1 正确性验证

测试代码在 tests/elastic/test_engram.py：

@torch.inference_mode()
def test(local_rank, num_local_ranks, args):
    rank, num_ranks, group = init_dist(local_rank, num_local_ranks)

    # 1. 创建 buffer
    buffer = deep_ep.ElasticBuffer(group, ...)

    # 2. 写入本地数据
    local_bf16 = torch.randn((num_entries, hidden), dtype=torch.bfloat16)
    buffer.engram_write(local_bf16)

    # 3. 生成随机 indices
    indices = torch.randint(0, num_ranks * num_entries, (num_tokens, num_entries_per_token))

    # 4. 正确性检查：与 all-gather 结果对比
    global_storage = torch.empty((num_ranks * num_entries, hidden), ...)
    dist.all_gather_into_tensor(global_storage, local_storage, group)
    ref_data = global_storage[indices.view(-1)].view(num_tokens, -1)

    data, _ = buffer.engram_fetch(indices)()
    assert torch.equal(ref_data, data), 'data mismatch'

正确性验证策略：

使用 dist.all_gather_into_tensor 获取全局数据作为 reference
对比 Engram fetch 结果与 reference 是否 bitwise identical
FP8 模式下额外验证 scaling factor 一致性


9.2 性能基准

# 性能测量
issue_t, wait_t = bench_kineto(
    fetch_and_wait,
    kernel_names=('engram_fetch_impl', 'engram_fetch_wait_impl'),
    barrier_comm_profiling=True,
    barrier=buffer.barrier)

# 计算吞吐
mpps = num_tokens * num_entries_per_token / (issue_t + wait_t) / 1e6  # MPPS
gb_s = num_fetched_bytes / (issue_t + wait_t) / 1e9  # GB/s

性能指标：

issue time：从发起 RDMA 到 kernel 结束的时间
wait time：从 wait 开始到所有 RDMA 完成的时间
MPPS: Million Packets Per Second（百万次查表/秒）
GB/s：有效带宽


9.3 参数空间


参数 默认值 说明
num_processes 4 测试进程数
num_qps 0 (= max) QP 数量（0 = 全部 SM）
num_entries 524288 每个 rank 的 entry 数
hidden 128 每个 entry 的 hidden dim
num_tokens 512 每次 fetch 的 token 数
num_entries_per_token 24 每个 token 查表的 entry 数
use_fp8 False 是否使用 FP8 存储
allow_hybrid_mode False 是否启用 hybrid 模式


11. 总结与设计洞察

核心设计原则


GPU-Initiated Everything：从地址计算到 RDMA doorbell，全部在 GPU kernel 内完成，消除 CPU 瓶颈
对称内存抽象：通过 CUDA Virtual Memory Management 将远端 CPU 内存映射到本地 VA，实现位置透明的查表
QP 并行化：每个 SM 一个 QP，warp 级协作，最大化 RDMA 并发
延迟执行：fetch 返回 callable，支持通信-计算 overlap
FP8 优化：数据走 RDMA，scaling factor 本地冗余，避免小数据传输


与 DeepEP 其他原语的关系


原语 SM 占用 通信方向 使用场景
Dispatch/Combine 4-64 SM All-to-all MoE EP 通信
PP Send/Recv 0 SM Point-to-point Pipeline Parallelism
AGRS 0 SM (Copy Engine) All-gather 权重同步
Engram 0 SM (瞬时) Remote read KV Cache 远程访问


局限与展望


只支持 read：当前 Engram 只支持远程读取（Get），不支持远程写入
CPU 段依赖：Engram 存储必须在 CPU 物理内存，无法直接存储在远端 GPU VRAM
RDMA 路径唯一：即使目标在同一 NVLink 域，仍然走 RDMA（无法利用 NVLink 的 GPU-to-GPU 带宽）
实验性质：API 标记为 Experimental，接口可能变化
未来方向：支持 GPU VRAM 直接注册（绕过 CPU 段）、支持 write 原语、与 Disaggregated Inference 框架集成


References


文件 路径 说明
Python API deep_ep/buffers/elastic.py (L409-604) Engram Python 接口
C++ Runtime csrc/elastic/buffer.hpp (L210-325) engram_write / engram_fetch 实现
CUDA Kernel (Fetch) deep_ep/include/deep_ep/impls/engram_fetch.cuh RDMA Get kernel
CUDA Kernel (Wait) deep_ep/include/deep_ep/impls/engram_fetch_wait.cuh RDMA wait kernel
JIT Runtime csrc/kernels/elastic/engram.hpp 模板代码生成与 launch
Symmetric Memory csrc/kernels/backend/symmetric.hpp (L191-309) HybridElasticSymmetricMemory
Comm Utils deep_ep/include/deep_ep/common/comm.cuh Gin barrier / QP mode
Test tests/elastic/test_engram.py 正确性 + 性能测试
Indexing csrc/indexing/main.cu Kernel 实例化入口


— Engram 架构深潜 · DeepEP v2.1.0 源码分析 · backyes.github.io/deepepv2 —

层级	范围	访问方式	延迟
GPU 段	本地节点所有 GPU	NVLink peer access	~10 µs
CPU 段（本地）	本地节点所有 rank	本地 GPU 不可直接访问，RDMA loopback	~1-2 µs
CPU 段（远端）	所有 scale-out 节点的 CPU 段	RDMA Get (RoCE/IB)	~5-15 µs

设计决策	选择	原因
对称布局	所有 rank VA 相同	O(1) 地址计算，无需查表
物理分布	各节点存各节点数据	避免一致性开销，最大化本地带宽
跨节点访问	RDMA one-sided read	远端 GPU 不参与，真正的 0 SM
CPU 段注册	NUMA-local + GDR	RDMA NIC 可直接 DMA 读取
全局窗口	ncclCommWindowRegister	NCCL 内部管理跨节点 MR

阶段	SM 占用	持续时间	说明
engram_fetch kernel	num_qps 个 block	~μs 级	计算地址 + 发起 RDMA
engram_fetch_wait kernel	num_qps 个 block	等待 RDMA 完成	轮询 completion
数据传输	0 SM	取决于数据量	NIC DMA 异步完成

部署模式	当前状态	可行性	挑战
与模型共部署	✅ 当前实现	完全支持	无
独立进程（同 GPU）	❌ 未实现	可行但复杂	需要 IPC 机制传递 indices
独立 GPU（分离式）	❌ 未实现	理论上可行	需要独立 NCCL group
CPU-only 节点	❌ 不支持	需要重新设计	Engram 需要 GPU 发起 RDMA

参数	默认值	说明
num_processes	4	测试进程数
num_qps	0 (= max)	QP 数量（0 = 全部 SM）
num_entries	524288	每个 rank 的 entry 数
hidden	128	每个 entry 的 hidden dim
num_tokens	512	每次 fetch 的 token 数
num_entries_per_token	24	每个 token 查表的 entry 数
use_fp8	False	是否使用 FP8 存储
allow_hybrid_mode	False	是否启用 hybrid 模式

原语	SM 占用	通信方向	使用场景
Dispatch/Combine	4-64 SM	All-to-all	MoE EP 通信
PP Send/Recv	0 SM	Point-to-point	Pipeline Parallelism
AGRS	0 SM (Copy Engine)	All-gather	权重同步
Engram	0 SM (瞬时)	Remote read	KV Cache 远程访问

文件	路径	说明
Python API	`deep_ep/buffers/elastic.py` (L409-604)	Engram Python 接口
C++ Runtime	`csrc/elastic/buffer.hpp` (L210-325)	engram_write / engram_fetch 实现
CUDA Kernel (Fetch)	`deep_ep/include/deep_ep/impls/engram_fetch.cuh`	RDMA Get kernel
CUDA Kernel (Wait)	`deep_ep/include/deep_ep/impls/engram_fetch_wait.cuh`	RDMA wait kernel
JIT Runtime	`csrc/kernels/elastic/engram.hpp`	模板代码生成与 launch
Symmetric Memory	`csrc/kernels/backend/symmetric.hpp` (L191-309)	HybridElasticSymmetricMemory
Comm Utils	`deep_ep/include/deep_ep/common/comm.cuh`	Gin barrier / QP mode
Test	`tests/elastic/test_engram.py`	正确性 + 性能测试
Indexing	`csrc/indexing/main.cu`	Kernel 实例化入口