DeepEP V2 ElasticBuffer 架构深度分析

// csrc/jit/launch_runtime.hpp:32-46
template <typename Derived>
class LaunchRuntime {
    template <typename Args>
    static std::string generate(const Args& args) {
        auto code = Derived::generate_impl(args);
        // 静态缓存 include hash（只计算一次）
        static std::string include_hash;
        if (include_hash.empty())
            include_hash = include_parser->get_hash_value(code);
        code = fmt::format("// Includes' hash value: {}\n{}", include_hash, code);
        return code;
    }

static std::string generate_impl(const Args& args) {
    // ...选择 header 和函数名
    return fmt::format(R"(
#include <deep_ep/impls/{}.cuh>
using namespace deep_ep::elastic;
static void __instantiate_kernel() {{
    auto ptr = reinterpret_cast<void*>(&{});
}}
)", header_name, func_name);
}

std::shared_ptr<KernelRuntime> build(const std::string& name, const std::string& code) const {
    // 签名 = name + compiler_signature + flags + code
    const auto kernel_signature = fmt::format("{}$${}$${}$${}", name, signature, flags, code);
    const auto dir_path = cache_dir_path / "cache" / fmt::format("kernel.{}.{}", name, get_hex_digest(kernel_signature));

    // 命中缓存
    if (const auto runtime = kernel_runtime_cache->get(dir_path); runtime != nullptr)
        return runtime;

    // 编译到临时目录后原子 rename
    const auto tmp_dir_path = make_tmp_dir() / get_uuid();
    // ... compile ...
    std::filesystem::rename(tmp_dir_path, dir_path);  // 原子操作
}

// csrc/jit/handle.hpp:23-81
#if CUDART_VERSION >= 12080 and defined(EP_JIT_USE_RUNTIME_API)
    using LibraryHandle = cudaLibrary_t;
    using KernelHandle = cudaKernel_t;
    // 使用 cudaLibraryLoadFromFile / cudaLibraryGetKernel
#else
    using LibraryHandle = CUmodule;
    using KernelHandle = CUfunction;
    // 使用 cuModuleLoad / cuModuleGetFunction
#endif

std::string get_hash_value(const std::string& code, const bool& exclude_code = true) {
    std::stringstream ss;
    for (const auto& i: get_includes(code))
        ss << get_hash_value_by_path(library_include_path / i) << "$";
    if (not exclude_code)
        ss << "#" << get_hex_digest(code);
    return get_hex_digest(ss.str());
}

// csrc/kernels/backend/nccl.cu:62-108
NCCLSymmetricMemoryContext::NCCLSymmetricMemoryContext(...) {
    // 查询 NCCL communicator 属性
    ncclCommProperties props = NCCL_COMM_PROPERTIES_INITIALIZER;
    NCCL_CHECK(ncclCommQueryProperties(comm, &props));

    // 配置 Gin 设备通信器
    ncclDevCommRequirements_t reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
    if (num_ranks > 1 and get_env("EP_DISABLE_GIN", 0) == 0) {
        reqs.ginContextCount = num_allocated_qps;          // QP 数量
        reqs.ginExclusiveContexts = true;                   // 独占 QP
        reqs.ginQueueDepth = kGinQPDepth;                  // QP 深度
        reqs.ginTrafficClass = sl_idx;                      // RDMA 服务等级
        reqs.ginSignalCount = num_ranks + 2 * 2;           // 信号量数
        reqs.ginConnectionType = allow_hybrid_mode ? 
            NCCL_GIN_CONNECTION_RAIL : NCCL_GIN_CONNECTION_FULL;
    }

    // 创建 NCCL 设备通信器
    NCCL_CHECK(ncclDevCommCreate(comm, &reqs, dev_comm.ptr));

    // 注册对称内存窗口
    NCCL_CHECK(ncclCommWindowRegister(comm, raw_window_ptr, num_bytes, &window, NCCL_WIN_STRICT_ORDERING));

    // 获取 LSA（Link-Symmetric Access）指针
    NCCL_CHECK(ncclGetLsaDevicePointer(window, 0, nvl_rank_idx, &mapped_window_ptr));
}

// csrc/kernels/elastic/handle.cuh — 实际在 deep_ep/include/deep_ep/common/handle.cuh
struct NCCLGin {
    const ncclDevComm_t& nccl_dev_comm;
    const ncclWindow_t& nccl_window;
    ncclGin gin;
    ncclTeam team_world, team_lsa, team_rail;

    // 支持三种 team 模式：
    // - ncclTeamTagWorld: 全局所有 rank
    // - ncclTeamTagLsa: NVLink 域内 rank
    // - ncclTeamTagRail: 同 rail 的 rank

    template <typename team_t>
    void get(void* src_ptr, void* dst_ptr, const int& num_bytes, const int& src_rank_idx) {
        gin.get(TEAM_WORLD_RAIL(), src_rank_idx, 
                nccl_window, offset(src_ptr), nccl_window, offset(dst_ptr), num_bytes, ...);
    }

    template <typename team_t>
    void put(void* recv_sym_ptr, void* send_sym_ptr, const int& num_bytes, const int& dst_rank_idx) {
        gin.put(TEAM_WORLD_RAIL(), dst_rank_idx, ...);
    }
};

template <int kNumSMs, int kNumQPs, int kNumChannelsPerSM, bool kWithNotifyWarps>
__device__ __forceinline__ std::pair<int, ncclGinResourceSharingMode> get_qp_mode(
    const int& sm_idx, const int& channel_in_sm_idx, const bool& is_notify_warp) {

    // Notify warp 独占 QP 0
    if (is_notify_warp) return {0, kSharingCTA};

    if constexpr (kNumSMs <= kNumAvailableQPs) {
        // SM 数 ≤ QP 数：每个 SM 至少一个独占 QP
        // e.g., 3 SMs, 10 QPs:
        // SM 0: 0 3 6 9
        // SM 1: 1 4 7
        // SM 2: 2 5 8
        return {kQPStartIdx + sm_idx + (channel_in_sm_idx % num_qps_in_sm) * kNumSMs, kSharingCTA};
    } else {
        // SM 数 > QP 数：所有 SM 共享所有 QP
        const auto global_channel_idx = sm_idx * kNumChannelsPerSM + channel_in_sm_idx;
        return {kQPStartIdx + (global_channel_idx % kNumAvailableQPs), kSharingGrid};
    }
}

[ GPU VRAM | CPU rank0 | CPU rank1 | ... | CPU rank(N-1) ]

static constexpr int64_t kNumAlignmentBytes = 2097152;  // 2 MB 对齐

template <int kNumRanks, int kNumSMs, int kNumThreads, int kNumQPs, int64_t kNumTimeoutCycles, typename team_t>
__device__ void gin_barrier_wo_local_sync(...) {
    // 1. 所有 SM 刷新所有 QPs（保证 store 可见性）
    for (int i = global_warp_idx; i < num_qps; i += kNumSMs * kNumWarps)
        ncclGin(nccl_dev_comm, i, NCCL_GIN_RESOURCE_SHARING_CTA).flush(ncclCoopWarp());
    cooperative_groups::this_grid().sync();

    // 2. SM 0 执行 barrier（只用 QP 0）
    if (sm_idx == 0) {
        for (int i = thread_idx; i < kNumRanks; i += kNumThreads)
            gin.signal(team, i, ncclGin_SignalInc{rank_idx});

        for (int i = thread_idx; i < kNumRanks; i += kNumThreads) {
            timeout_while<kNumTimeoutCycles>([=](const bool& is_last_check) {
                const auto signal = ptx::ld_acquire_sys<uint64_t>(signal_ptr);
                if (signal >= target) return true;
                // ... timeout 处理
            });
        }
    }
}

deep_ep/include/deep_ep/
├── common/           # 基础工具（header-only）
│   ├── comm.cuh      # Gin 通信原语
│   ├── handle.cuh    # NCCLGin handle 封装
│   ├── layout.cuh    # Token/Buffer/Workspace 布局
│   ├── math.cuh      # 数学工具
│   ├── ptx.cuh       # PTX 内联汇编
│   ├── compiled.cuh  # 编译期开关
│   └── exception.cuh # 异常处理
└── impls/            # Kernel 实现（header-only）
    ├── dispatch.cuh
    ├── dispatch_copy_epilogue.cuh
    ├── combine.cuh
    ├── combine_reduce_epilogue.cuh
    ├── hybrid_dispatch.cuh
    ├── hybrid_combine.cuh
    ├── engram_fetch.cuh
    ├── engram_fetch_wait.cuh
    ├── pp_send_recv.cuh
    └── barrier.cuh

// 通过 DISABLE_SM90_FEATURES 控制 FP8/TMA 支持
#ifndef DISABLE_SM90_FEATURES
    static constexpr bool kEnableSM90Features = true;
#else
    static constexpr bool kEnableSM90Features = false;
#endif

// csrc/python_api.cpp:22-39
PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("is_sm90_compiled", []() { return deep_ep::kEnableSM90Features; });
    m.attr("topk_idx_t") = py::cast(c10::CppTypeToScalarType<deep_ep::topk_idx_t>::value);
    deep_ep::jit::register_apis(m);        // JIT
    deep_ep::legacy::register_apis(m);     // Legacy
    deep_ep::elastic::register_apis(m);    // Elastic V2
}

# deep_ep/buffers/elastic.py:301
self.nccl_comm_handle = get_nccl_comm_handle(group, force_new_comm=num_cpu_bytes > 0)

// csrc/kernels/backend/nccl.cu:76
comm = reinterpret_cast<ncclComm_t>(nccl_comm_comm);

// csrc/kernels/backend/nccl.cu:49-54
std::tuple<int, int> get_physical_domain_size(const int64_t& nccl_comm) {
    const int num_ranks = ncclTeamWorld(comm).nRanks;
    const int num_nvl_ranks = ncclTeamLsa(comm).nRanks;
    return {num_ranks / num_nvl_ranks, num_nvl_ranks};
}

std::tuple<int, int> get_logical_domain_size(const int64_t& nccl_comm, const bool& allow_hybrid_mode) {
    const auto [num_rdma_ranks, num_nvl_ranks] = get_physical_domain_size(nccl_comm);
    return {allow_hybrid_mode ? num_rdma_ranks : 1,
            allow_hybrid_mode ? num_nvl_ranks : num_rdma_ranks * num_nvl_ranks};
}

Symmetric Memory Layout:
[[[Workspace] GPU Buffer] CPU Buffer]
 ↑                       ↑          ↑
 mapped_window_ptr       workspace  cpu_segment
                         + gpu_buf  (for Engram)

// layout.cuh:43-80
static int64_t get_num_bytes() {
    // 16 bytes NVLink barrier counter
    // (1024 + 2048) * 8 bytes notify reduction workspace
    // 1024 * 8 * 2 scaleup rank count (send/recv)
    // 2048 * 8 * 2 scaleup expert count (send/recv)
    // 1024 * 4 scaleup atomic sender counter
    // 1024 * 4 * 2 scaleout rank count (send/recv)
    // 2048 * 4 * 2 scaleout expert count (send/recv)
    // 1024 * 256 * 8 scaleout channel metadata
    // 1024 * 256 * 4 channel scaleup tail
    // 2 * 2 * 8 PP send/recv count
    // (32 + 1) * 1024 * 4 AGRS signals
}

struct TokenLayout {
    int num_hidden_bytes, num_sf_bytes;
    int num_topk, num_metadata_bytes;
    bool with_metadata;

    // 每个 token 的内存布局:
    // [hidden (TMA-aligned)] [sf (TMA-aligned)] [metadata (TMA-aligned)] [mbarrier (optional)]
    // metadata = topk_idx (num_topk * 4) + topk_weights (num_topk * 4) + src_token_global_idx (4) + linked_list_idx (num_topk * 4)

    template <bool kWithMBarrier, typename dtype_t = int>
    dtype_t get_num_bytes() const {
        return align(num_hidden_bytes, 128) + align(num_sf_bytes, 128) + 
               align(num_metadata_bytes, 128) + align(kWithMBarrier ? sizeof(mbarrier) : 0, 128);
    }
};

class ElasticBuffer:
    def dispatch(self, x, topk_idx, topk_weights, ...) -> (recv_x, recv_topk_idx, recv_topk_weights, handle, event)
    def combine(self, x, handle, topk_weights, ...) -> (combined_x, combined_topk_weights, event)
    def engram_write(self, storage, sf) -> None
    def engram_fetch(self, indices, num_qps, use_tma_aligned_col_major_sf) -> Callable
    def pp_send(self, x, dst_rank_idx, num_sms) -> None
    def pp_recv(self, x, src_rank_idx, num_sms) -> None
    def all_gather(self, tensors) -> (gathered_tensors, handle)

# elastic.py:271-277
if use_tma_aligned_col_major_sf:
    // TMA-aligned column-major layout for the next GEMM input
    sf_token_stride = 1
    sf_hidden_stride = math.align(num_tokens, kNumAlignedSFPacks)

static int64_t calculate_buffer_size(const int64_t& nccl_comm,
                                     const int& num_max_tokens_per_rank, const int& hidden,
                                     int num_topk, const bool& use_fp8_dispatch,
                                     const bool& allow_hybrid_mode,
                                     const bool& allow_multiple_reduction) {
    // 自动推断 topology
    const auto [num_rdma_ranks, num_nvl_ranks] = nccl::get_physical_domain_size(nccl_comm);
    const auto [num_scaleout_ranks, num_scaleup_ranks] = nccl::get_logical_domain_size(nccl_comm, allow_hybrid_mode);

    // 计算 dispatch 和 combine 各自需要的 buffer size
    const auto num_dispatch_bytes = get_dispatch_buffer_size(...);
    const auto num_combine_bytes = get_combine_buffer_size(...);

    // 取最大值，对齐到 2 MB
    return math::align(std::max(num_dispatch_bytes, num_combine_bytes), symmetric::kNumAlignmentBytes);
}

static constexpr int kNumMaxRanks = 1024;       // 最大 rank 数
static constexpr int kNumMaxExperts = 2048;     // 最大 expert 数
static constexpr int kNumMaxExpertsPerRank = 256; // 每个 rank 最大 expert 数

if os.environ.get('NCCL_GIN_CROSS_NIC') == '0':
    # Multi-plane: all ranks share CPU segments, skip proxy re-export
    os.environ.setdefault('NCCL_SYM_REUSE_SYSMEM_HANDLES', '1')

if num_cpu_bytes > 0:
    num_gpu_bytes = num_bytes - num_cpu_bytes
    num_max_local_ranks = int(os.getenv('EP_NUM_MAX_LOCAL_RANKS', 16)) if allow_hybrid_mode else 1

    num_registered_bytes = num_gpu_bytes + num_cpu_bytes * num_max_local_ranks + (1 << 32)
    num_total_gpu_bytes = torch.cuda.get_device_properties('cuda').total_memory
    if num_registered_bytes > num_total_gpu_bytes:
        win_stride = align(num_registered_bytes, 1 << 32)
        os.environ['NCCL_WIN_STRIDE'] = str(win_stride)

@weak_lru(maxsize=None)
def get_theoretical_num_sms(self, num_experts, num_topk, num_scaleout_topk=0,
                            rdma_gbs=0, nvlink_gbs=0,
                            sm_read_gbs=200, sm_write_gbs=50) -> int:
    # 1. 自动获取带宽
    if rdma_gbs == 0 and self.num_rdma_ranks > 1:
        rdma_gbs = get_rdma_gbs()
    if nvlink_gbs == 0:
        nvlink_gbs = get_nvlink_gbs()

    # 2. 计算期望 topk（去重后）
    def get_expected_topk(num_groups):
        return num_groups * (1 - math.comb(num_experts - num_experts // num_groups, num_topk) / math.comb(num_experts, num_topk))

    # 3. 计算各类型流量
    sm_read, sm_write = 0, 0
    rdma_traffic, nvlink_traffic = 0, 0

    # ... 详细流量建模 ...

    # 4. 找到瓶颈
    if self.num_scaleout_ranks > 1 and (rdma_traffic / rdma_gbs) > (nvlink_traffic / nvlink_gbs):
        bounded_traffic, bounded_gbs = rdma_traffic, rdma_gbs
    else:
        bounded_traffic, bounded_gbs = nvlink_traffic, nvlink_gbs

    # 5. 计算 SM 数
    num_sms = max(
        bounded_gbs / bounded_traffic * sm_read / sm_read_gbs,
        bounded_gbs / bounded_traffic * sm_write / sm_write_gbs,
    )
    num_sms = align(max(4, math.ceil(num_sms * 1.25)), 2)  # 至少 4，向上取偶数，+25% 余量
    num_sms = num_sms if self.prefer_overlap_with_compute else max(num_sms, 64)
    num_sms = min(num_sms, num_device_sms)

def get_theoretical_num_qps(self, num_sms: int) -> int:
    # Direct mode: 鼓励较少 QPs（减少 DB ringing 开销）
    num_qps = min(num_sms, 8 + 1)

    # Hybrid mode: 每个 channel（和 notify）独立 QP
    if self.allow_hybrid_mode:
        num_qps = num_sms * 16 + 1

    return min(num_qps, self.num_allocated_qps)

# Automatic decide SM and QP count
num_topk = (handle.topk_idx if topk_idx is None else topk_idx).shape[1]
num_sms = self.get_theoretical_num_sms(num_experts, num_topk) if num_sms == 0 else num_sms
num_qps = self.get_theoretical_num_qps(num_sms) if num_qps == 0 else num_qps
assert num_qps <= self.num_allocated_qps

// 根据 allow_hybrid_mode 决定逻辑拓扑
if (allow_hybrid_mode) {
    num_scaleout_ranks = num_rdma_ranks, num_scaleup_ranks = num_nvl_ranks;
    scaleout_rank_idx = rdma_rank_idx, scaleup_rank_idx = nvl_rank_idx;
} else {
    num_scaleout_ranks = 1, num_scaleup_ranks = num_ranks;
    scaleout_rank_idx = 0, scaleup_rank_idx = rank_idx;
}

static std::string generate_impl(const Args& args) {
    if (args.num_scaleout_ranks == 1) {
        header_name = "dispatch";
        func_name = fmt::format("dispatch_impl<...>", ...);
    } else {
        header_name = "hybrid_dispatch";
        func_name = fmt::format("hybrid_dispatch_impl<...>", ...);
    }
}

# elastic.py:823
num_sms = align(max(4, math.ceil(num_sms * 1.25)), 2)
num_sms = num_sms if self.prefer_overlap_with_compute else max(num_sms, 64)

// dispatch.hpp:226
.launch_args = jit::LaunchArgs(num_sms, num_threads, num_smem_bytes, 2 - (num_sms % 2), true);
//                                                                          ^^^^^^^^^^^^^^^^
//                                                                          cluster_dim=2（偶数 SMs）或 1（奇数 SMs）

// csrc/kernels/elastic/engram.hpp:98-126
static void launch_engram_fetch(...) {
    constexpr int kNumEngramFetchThreads = 1024;  // 仅 1 个 block！

    const EngramFetchRuntime::Args args = {
        .launch_args = jit::LaunchArgs(num_qps, kNumEngramFetchThreads)
        //                                       ^^^^^^^^
        //                                       grid_dim = num_qps（通常 1-8）
    };
}

// 伪代码
template <int kNumQPs, ...>
__global__ void engram_fetch_impl(...) {
    const auto sm_idx = blockIdx.x;  // 每个 block 对应 1 个 QP
    const auto gin = handle::NCCLGin(nccl_dev_comm, nccl_window, sm_idx, ...);

    // 每个 token 的 entry 通过 RDMA get 拉取
    for (int i = thread_idx; i < num_tokens * num_entries_per_token; i += blockDim.x) {
        const auto token_idx = i / num_entries_per_token;
        const auto entry_idx = i % num_entries_per_token;
        const auto storage_idx = indices[token_idx * num_entries_per_token + entry_idx];

        // RDMA one-sided get — 不占用 SM 计算资源
        gin.get<team_t>(storage_ptr, fetched_ptr, num_bytes, src_rank_idx);
        last_gin_requests[sm_idx] = gin.last_request();
    }
}

// csrc/kernels/elastic/engram.hpp:169-189
static void launch_engram_fetch_wait(...) {
    constexpr int kNumEngramFetchWaitThreads = 1024;
    const EngramFetchWaitRuntime::Args args = {
        .launch_args = jit::LaunchArgs(num_qps, kNumEngramFetchWaitThreads)
    };
}

void engram_write(const torch::Tensor& storage, const std::optional<torch::Tensor>& sf) {
    // 写入 CPU segment（在 symmetric memory 的后部）
    const auto cpu_write_offset = allow_hybrid_mode
        ? static_cast<int64_t>(nccl_context->scale_up_rank_idx) * num_cpu_buffer_bytes : 0;
    CUDA_RUNTIME_CHECK(cudaMemcpyAsync(
        math::advance_ptr(buffer, num_gpu_buffer_bytes + cpu_write_offset),
        storage.data_ptr(), storage.nbytes(),
        cudaMemcpyDeviceToDevice, compute_stream));
}

# elastic.py:327-375
def pp_set_config(self, num_max_tensor_bytes, num_max_inflight_tensors):
    self.runtime.pp_set_config(num_max_tensor_bytes, num_max_inflight_tensors)

def pp_send(self, x, dst_rank_idx, num_sms=0):
    self.runtime.pp_send(x, dst_rank_idx, 
                         num_sms=num_sms if num_sms else jit.device_runtime.get_num_sms())

def pp_recv(self, x, src_rank_idx, num_sms=0):
    self.runtime.pp_recv(x, src_rank_idx, 
                         num_sms=num_sms if num_sms else jit.device_runtime.get_num_sms())

// csrc/kernels/elastic/pp_send_recv.hpp:64-95
static void launch_pp_send(...) {
    const PPSendRuntime::Args args = {
        .launch_args = jit::LaunchArgs(num_sms, 32, num_smem_bytes, 1, true)
        //                                       ^^^^^^^^
        //                                       默认使用所有 SMs
    };
}

void pp_set_config(const int64_t& num_max_tensor_bytes, const int& num_max_inflight_tensors) {
    EP_HOST_ASSERT(num_max_tensor_bytes * num_max_inflight_tensors * 2 * 2 <= num_buffer_bytes);
    //                                                                          ^^^^^^^
    //                                                                          send/recv × prev/next
    this->prev_rank_idx = (nccl_context->rank_idx + nccl_context->num_ranks - 1) % nccl_context->num_ranks;
    this->next_rank_idx = (nccl_context->rank_idx + 1) % nccl_context->num_ranks;
    this->num_max_pp_tensor_bytes = math::align<int64_t>(num_max_tensor_bytes, 32);
    this->num_max_pp_inflight_tensors = num_max_inflight_tensors;
}

# elastic.py:638-726
def create_agrs_session(self):
    self.runtime.create_agrs_session()

def destroy_agrs_session(self):
    self.runtime.destroy_agrs_session()

def all_gather(self, tensors):
    return self.runtime.all_gather(tensors)

std::pair<std::vector<torch::Tensor>, std::function<void()>>
all_gather(const std::vector<torch::Tensor>& tensors) {
    // 1. 计算需要多少 copies
    int num_copies = 0;
    for (int i = 0; i < nccl_context->num_ranks; ++ i)
        for (int j = 0; j < num_tensors; ++ j)
            num_copies += nccl_context->num_ranks - is_inplace;

    // 2. 使用 cudaMemcpyBatchAsync — 完全由 Copy Engine 执行
    cudaMemcpyAttributes attrs = {
        .srcAccessOrder = cudaMemcpySrcAccessOrderStream,
        .flags = cudaMemcpyFlagPreferOverlapWithCompute
    };
    CUDA_RUNTIME_CHECK(cudaMemcpyBatchAsync(dst_ptrs.data(), src_ptrs.data(), sizes.data(), num_copies, attrs, comm_stream));

    // 3. 信号通知其他 rank
    cuda_driver::batched_write_and_wait(comm_stream, write_ptrs, wait_ptrs, current_session);

    // 4. 返回 gathered tensors 和 wait handle
    return {out, handle};
}

// csrc/kernels/backend/cuda_driver.cu
void batched_write_and_wait(CUstream stream, const std::vector<void*>& write_ptrs, 
                            const std::vector<void*>& wait_ptrs, const int& value) {
    // 批量发起 writes
    for (auto& ptr : write_ptrs)
        cuStreamWriteValue32(stream, reinterpret_cast<CUdeviceptr>(ptr), value, CU_STREAM_WRITE_VALUE_DEFAULT);

    // 批量等待 waits
    for (auto& ptr : wait_ptrs)
        cuStreamWaitValue32(stream, reinterpret_cast<CUdeviceptr>(ptr), value, CU_STREAM_WAIT_VALUE_EQ);
}

DeepEP/
├── csrc/
│   ├── jit/                         # JIT 编译系统
│   │   ├── compiler.hpp             # NVCC 编译 + 缓存
│   │   ├── kernel_runtime.hpp       # CUBIN 加载
│   │   ├── launch_runtime.hpp       # Launch 配置
│   │   ├── include_parser.hpp       # 递归 hash
│   │   ├── handle.hpp               # Driver/Runtime API 双路径
│   │   ├── cache.hpp                # 双层缓存
│   │   ├── device_runtime.hpp       # GPU 属性查询
│   │   └── api.hpp                  # pybind11 注册
│   ├── kernels/
│   │   ├── backend/
│   │   │   ├── nccl.cu              # NCCL Gin 后端
│   │   │   ├── symmetric.hpp        # 对称内存三层架构
│   │   │   ├── nvshmem.cu           # NVSHMEM 后备
│   │   │   ├── cuda_driver.cu       # Driver API 工具
│   │   │   └── api.cuh              # Backend API 声明
│   │   └── elastic/
│   │       ├── dispatch.hpp         # Dispatch launch
│   │       ├── combine.hpp          # Combine launch
│   │       ├── engram.hpp           # Engram launch
│   │       ├── pp_send_recv.hpp     # PP launch
│   │       ├── barrier.hpp          # Barrier launch
│   │       └── api.hpp              # Kernel API 汇聚
│   ├── elastic/
│   │   ├── buffer.hpp               # ElasticBuffer 核心（1382 行）
│   │   └── utils.hpp                # Elastic 工具
│   └── python_api.cpp               # pybind11 模块（39 行）
├── deep_ep/
│   ├── buffers/
│   │   ├── elastic.py               # Python 入口（1107 行）
│   │   └── legacy.py                # Legacy API 兼容
│   └── include/deep_ep/
│       ├── common/                  # Header-only 基础
│       │   ├── comm.cuh             # Gin 通信原语
│       │   ├── handle.cuh           # NCCLGin handle
│       │   ├── layout.cuh           # Token/Buffer/Workspace 布局
│       │   ├── math.cuh             # 数学工具
│       │   ├── ptx.cuh              # PTX 内联汇编
│       │   ├── compiled.cuh         # 编译期开关
│       │   └── exception.cuh        # 异常处理
│       └── impls/                   # Header-only kernel 实现
│           ├── dispatch.cuh
│           ├── dispatch_copy_epilogue.cuh
│           ├── combine.cuh
│           ├── combine_reduce_epilogue.cuh
│           ├── hybrid_dispatch.cuh
│           ├── hybrid_combine.cuh
│           ├── engram_fetch.cuh
│           ├── engram_fetch_wait.cuh
│           ├── pp_send_recv.cuh
│           └── barrier.cuh
└── setup.py                         # 构建配置