moe_combine_normal_zb/
├── op_kernel/          # 设备端 kernel 代码
│   ├── moe_combine_normal_zero_buffer.cpp    # kernel 入口 + 实现
│   ├── moe_combine_normal_zero_buffer.h      # 类定义
│   └── moe_combine_normal_zero_buffer_tiling.h  # tiling 数据结构
├── op_host/            # 主机端 tiling 代码
│   ├── moe_combine_normal_zero_buffer.cpp    # tiling 入口
│   └── moe_combine_normal_zero_buffer_tiling.cc  # tiling 实现
└── op_api/             # aclnn API 注册
    ├── aclnn_moe_combine_normal_zero_buffer.cpp
    └── aclnn_moe_combine_normal_zero_buffer.h

moe_combine_lowlatency_zb/
├── op_kernel/
│   ├── moe_combine_lowlatency_zero_buffer.cpp
│   ├── moe_combine_lowlatency_zero_buffer.h
│   └── moe_combine_lowlatency_zero_buffer_tiling.h
├── op_host/
│   ├── moe_combine_lowlatency_zero_buffer.cpp
│   └── moe_combine_lowlatency_zero_buffer_tiling.cpp
└── op_api/
    ├── aclnn_moe_combine_lowlatency_zero_buffer.cpp
    └── aclnn_moe_combine_lowlatency_zero_buffer.h

utils/op_kernel/        # 公共工具
├── zero_buffer_api.h         # SHMEM API 封装 (include shmem.h)
├── zero_buffer_sync_flag.h   # 跨卡 flag 同步机制
├── comm_args_zb.h            # ZB 版通信参数
├── sync_collectives_zb.h     # 同步原语
├── moe_lowlatency_zero_buffer_base.h  # lowlatency 共享工具
└── data_copy_zb.h            # 数据拷贝辅助

extern "C" __global__ __aicore__ void moe_combine_normal_zero_buffer(
    GM_ADDR recvX,           // [bs*k, h] 本 rank 收到的所有 token（各 expert 输出）
    GM_ADDR epRecvCount,     // [moeExpertNum, epWorldSize] 每个 expert 从各 rank 收到的 token 数
    GM_ADDR topkWeights,     // [bs, k] top-k 权重
    GM_ADDR topkIdx,         // [bs, k] top-k expert ID
    GM_ADDR sendTokenIdx,    // [bs, k] 在远端 recvX 中的偏移索引
    GM_ADDR probGrad,        // [bs, h] 概率梯度输入（可选）
    GM_ADDR XOut,            // [bs, h] 输出
    GM_ADDR sendCostStatsOut,// [epWorldSize] 诊断统计（可选）
    GM_ADDR gradOut,         // [bs, k] 梯度输出（可选）
    GM_ADDR workspaceGM,
    GM_ADDR tilingGM)

Process():
│
├── syncFlag_.IncrementMagic()     // Step 1: 获取 magic M，递增
├── syncFlag_.BarrierAll()         // Step 2: 全 rank 屏障，确保所有 GMM 完成
├── ReadBufferFromRemote()         // Step 3: 从远端 SHMEM 读取并加权求和
└── SyncAll<true>()                // Step 4: 全核同步

__aicore__ inline void IncrementMagic() {
    GM_ADDR magicAddr = flagBaseGM_ + coreIdx_ * MAGIC_SLOT_ALIGN;  // 512B 对齐
    // DataCacheCleanAndInvalid → 读 magic
    magic_ = static_cast<int32_t>(magicGT(0));
    magicGT(0) = magic_ + 1;  // 递增
    // DataCacheCleanAndInvalid → 写回
    barrierSeq_ = 0;
}

__aicore__ inline void BarrierAll() {
    SyncAll<true>();                              // 核内同步
    SplitCoreCal(epWorldSize_, ...);              // 分配远端 rank
    // Write: 向所有分配的远端 rank 写 barrier flag
    for (r = startRank; r < endRank; r++) {
        DataCopy(remoteSlot, ubFlag, ...);        // 写 flag 到远端 SHMEM
    }
    // Wait: 轮询本地 slot 等待所有远端 rank 的 flag
    for (r = startRank; r < endRank; r++) {
        do { observed = localSlot; } while (observed != barrierFlag);
    }
    SyncAll<true>();
}

__aicore__ inline void ReadBufferFromRemote() {
    SplitCoreCal(axisBS_, ...);  // 将 bs 个 token 均匀分配到各 AIV core
    // 初始化所有 buffer
    tpipe_->InitBuffer(xOutBuf_, h32AlignRecvXLen_);         // 14KB (h=7168, bf16)
    tpipe_->InitBuffer(tokenFloatBuf_, h32AlignFloatLen_);   // 28KB (fp32)
    tpipe_->InitBuffer(weightedMulBuf_, h256AlignFloatLen_); // 28KB
    tpipe_->InitBuffer(sumFloatBuf_, h32AlignFloatLen_);     // 28KB
    tpipe_->InitBuffer(weightedSumQueue_, DOUBLE_BUFFER, h32AlignRecvXLen_); // 28KB
    // ... topkWeightsBuf_, sendTokenIdxBuf_, topkIdxBuf_, allRecvCountBuf_

    // 读取全局 recvCount 到本地
    DataCopyPad(allRecvCountLocal, epRecvCountGT_, ...);

    for (tokenIndex = startTokenIndex; tokenIndex < endTokenIndex; tokenIndex++) {
        // 读取当前 token 的 topkWeights, topkIdx, sendTokenIdx
        DataCopyPad(topkWeightsLocal, topkWeightsGT_[tokenIndex * axisK_], ...);
        DataCopyPad(topkIdxLocal, topkIdxGT_[tokenIndex * axisK_], ...);
        DataCopyPad(sendTokenIdxLocal, sendTokenIdxGT_[tokenIndex * axisK_], ...);

        ReadBufferAndWeightedSum(tokenIndex, startTokenIndex);
    }
}

__aicore__ inline void ReadBufferAndWeightedSum(uint32_t tokenIndex, uint32_t startTokenIndex) {
    Duplicate(sumFloatBufLocal, 0.0f, axisH_);  // 清零累加器

    for (uint32_t topkId = 0U; topkId < axisK_; topkId++) {
        float scale = topkWeightsLocal.GetValue(topkId);
        int32_t expertId = topkIdxLocal.GetValue(topkId);
        int32_t remoteReadOffset = sendTokenIdxLocal(topkId);
        int32_t remoteReadBase = allRecvCountLocal(expertId * epWorldSize_ + epRankId_);
        uint64_t remoteReadAddr = (remoteReadBase + remoteReadOffset) * hRecvXTypeLen_;

        int32_t dstRankId = expertId / moeExpertPerRankNum_;
        auto ptr = reinterpret_cast<__gm__ uint8_t *>(shmem_ptr(recvXGM_, dstRankId));
        dstGT.SetGlobalBuffer((__gm__ XType *)(ptr + remoteReadAddr));

        // 从远端 SHMEM 读取 token
        LocalTensor<XType> tmpToken = weightedSumQueue_.AllocTensor<XType>();
        DataCopyPad(tmpToken, dstGT, xOutCopyParams, ...);
        weightedSumQueue_.EnQue(tmpToken);
        tmpToken = weightedSumQueue_.DeQue<XType>();

        // bf16 → fp32 → 乘权重 → 累加
        Cast(tokenFloatLocal, tmpToken, CAST_NONE, axisH_);      // bf16 → fp32
        Muls(weightedMulBufLocal, tokenFloatLocal, scale, axisH_); // × weight
        Add(sumFloatBufLocal, sumFloatBufLocal, weightedMulBufLocal, axisH_); // 累加

        weightedSumQueue_.FreeTensor<XType>(tmpToken);
    }

    // fp32 → bf16 → 写输出
    Cast(xOutLocal, sumFloatBufLocal, CAST_RINT, axisH_);
    SyncFunc<V_MTE3>();
    DataCopyPad(xOutGlobal_[tokenIndex * axisH_], xOutLocal, ...);
}

expertId = topkIdx[i][j]                    // 目标 expert
dstRankId = expertId / moeExpertPerRankNum_  // expert 所在的 rank
remoteReadBase = epRecvCount[expertId][epRankId_]  // 本 rank 在该 expert 上的基址偏移
remoteReadOffset = sendTokenIdx[i][j]        // 在该 expert 接收列表中的位置
remoteReadAddr = (remoteReadBase + remoteReadOffset) * hRecvXTypeLen_

// 最终地址 = 远端 rank 的 SHMEM window + remoteReadAddr
finalAddr = shmem_ptr(recvXGM_, dstRankId) + remoteReadAddr

extern "C" __global__ __aicore__ void moe_combine_lowlatency_zero_buffer(
    GM_ADDR expandX,            // [A*tpWorldSize, h] 展开的 expert 输出
    GM_ADDR expertIds,          // [bs, k] top-k expert ID
    GM_ADDR assistInfoForCombine, // [A*128] 辅助信息
    GM_ADDR epSendCount,        // [moeExpertNum*epWorldSize] EP 发送计数
    GM_ADDR scales,             // [bs, k] top-k 权重
    GM_ADDR tpSendCount,        // [tpWorldSize] TP 发送计数
    GM_ADDR xActiveMask,        // [bs] 或 [bs, k] 活跃 mask（可选）
    GM_ADDR activationScale,    // 预留
    GM_ADDR weightScale,        // 预留
    GM_ADDR groupList,          // 预留
    GM_ADDR expandScales,       // 预留
    GM_ADDR sharedExpertX,      // [bs, h] 共享专家输出（可选）
    GM_ADDR elasticInfo,        // 弹性信息（可选）
    GM_ADDR oriX,               // [bs, h] 原始输入（copy expert 用）
    GM_ADDR constExpertAlpha1,  // [constExpertNum] 常量专家参数
    GM_ADDR constExpertAlpha2,  // [constExpertNum]
    GM_ADDR constExpertV,       // [constExpertNum, h]
    GM_ADDR XOut,               // [bs, h] 输出
    GM_ADDR workspaceGM, GM_ADDR tilingGM)

10000: tp=1, IsInt8Quant=0
10100: tp=2, IsInt8Quant=0
10020: tp=1, IsInt8Quant=1
10120: tp=2, IsInt8Quant=1

Process():
│
├── CopyValidExpandXToRemote()    // Step 1: 拷贝本地 expert 数据到 SHMEM
├── SyncAll<true>()               // Step 2: 全核同步
├── tpipe_->Reset()               // 重置 pipe
├── SetSyncFlag()                 // Step 3: 向远端 rank 写 flag
├── WaitSyncFlag()                 // Step 4: 轮询等待远端 flag
├── SyncAll<true>()               // Step 5: 全核同步
└── InputToDstOutput()            // Step 6: 从 SHMEM 读取并加权求和

__aicore__ inline void CopyValidExpandXToRemote() {
    if (!needCopyExpandX_) return;  // oriX==0 或 validTokenNum==0 时跳过

    SplitToCore(validTokenNum_, aivNum_, ...);  // 分配 token
    tpipe_->InitBuffer(copyInBuf_, COPY_BUFFER_NUM * GM2GM_COPY_BLOCK_SIZE);  // 2×64KB

    // 双缓冲 GM→GM 拷贝：oriX → expandX (SHMEM)
    for (i = 0; dataSizeRemain > 0; ++i) {
        copySize = min(dataSizeRemain, 64KB);
        WaitFlag<MTE3_MTE2>(eventId);           // 等待 DMA 通道
        DataCopyPad(copyLocal[bufferOffset], oriXByteGT[srcOffset], ...);  // GM→UB
        SetFlag<MTE2_MTE3>(eventId);
        WaitFlag<MTE2_MTE3>(eventId);
        DataCopyPad(expandXByteGT[dstOffset], copyLocal[bufferOffset], ...);  // UB→GM(SHMEM)
        SetFlag<MTE3_MTE2>(eventId);
    }
}

__aicore__ inline void SetSyncFlag() {
    // 构造 flag 值：0x3F800000 = 1.0f (IEEE 754)
    Duplicate(statusTensor, 0, rankNumPerBlock_ * 8);
    Duplicate(statusTensor, 0x3F800000, mask, ...);  // 按 mask 设置 1.0

    for (targetRankId = curBlockStartRankId_; targetRankId < curBlockEndRankId_; ++targetRankId) {
        auto ptr = GetSyncFlagAddrByRankId(targetRankId) + epRankId_ * COMBINE_SYNC_STATE_OFFSET;
        gmRemoteStatusGT.SetGlobalBuffer(ptr);
        DataCopy(gmRemoteStatusGT, statusTensor[...], 8UL);
    }
}

__aicore__ inline void WaitSyncFlag() {
    // 轮询本地 SHMEM 中的 flag 区域
    compareTarget = 1.0f * rankNumPerBlock_;
    while ((sumOfFlag < minTarget) || (sumOfFlag > maxTarget)) {
        DataCopy(statusFp32Tensor, statusFp32TensorGT[curBlockStartRankId_ * ...], ...);
        ReduceSum(statusSumOutTensor, statusFp32Tensor, ...);
        sumOfFlag = statusSumOutTensor.GetValue(0);
    }
    // 清理 flag
    Duplicate(cleanStateTensor, 0, ...);
    DataCopy(statusFp32TensorGT[...], cleanStateTensor.ReinterpretCast<float>(), ...);
}

__aicore__ inline void InputToDstOutput() {
    SplitToCore(axisBS_, aivNum_, ...);  // 分配 token

    for (tokenIndex = startTokenId; tokenIndex < endTokenId; tokenIndex++) {
        DataCopyPad(expertIdsTensor, expertIdsGM_[tokenIndex * axisK_], ...);
        Duplicate(sumTokenTensor_, 0.0f, axisH_);

        for (topkId = 0U; topkId < axisK_; topkId++) {
            int32_t dstExpertId = expertIdsTensor.GetValue(topkId);
            int32_t srcRankId = dstExpertId / moeExpertPerRankNum_;
            uint32_t offsetIdx = dstExpertId * epWorldSize_ + epRankId_;
            uint32_t col = offsetIdx % moeExpertNum_;
            int32_t remoteBase = (col == 0) ? 0 : epSendCountGM_(offsetIdx - 1);
            int32_t remoteOffset = expandIdxLocal(localTokenId * axisK_ + topkId);
            float scale = expertScalesLocal_.GetValue(localTokenId * axisK_ + topkId);

            // 从远端 SHMEM 读取
            auto srcPtr = shmem_ptr(expandXGM_, srcRankId);
            srcWinGMTensor.SetGlobalBuffer(srcPtr + hExpandXTypeSize_ * (remoteBase + remoteOffset));

            // 加权求和（同 normal_zb）
            DataCopyPad(tmpToken, srcWinGMTensor, ...);
            Cast(tokenReduceTensor_, tmpToken, CAST_NONE, axisH_);
            Muls(weightedSumTensor_, tokenReduceTensor_, scale, axisH_);
            Add(sumTokenTensor_, sumTokenTensor_, weightedSumTensor_, axisH_);
        }

        Cast(xOutLocal, sumTokenTensor_, CAST_RINT, axisH_);
        DataCopyPad(XOutGT_[tokenIndex * axisH_], xOutLocal, ...);
    }
}

// 前缀和计算远端基址
offsetIdx = dstExpertId * epWorldSize_ + epRankId_
col = offsetIdx % moeExpertNum_
remoteBase = (col == 0) ? 0 : epSendCountGM_(offsetIdx - 1)
remoteOffset = expandIdxLocal(localTokenId * axisK_ + topkId)

// 最终地址
srcAddr = shmem_ptr(expandXGM_, srcRankId) + hExpandXTypeSize_ * (remoteBase + remoteOffset)

shmem.h (CANN 提供)
    ├── shmem_ptr(baseGM, rankId)              → 获取远端 rank 的 SHMEM window 基址
    ├── aclshmem_half_put_nbi(...)             → 非阻塞 half 写入远端
    ├── aclshmem_bfloat16_put_nbi(...)         → 非阻塞 bf16 写入远端
    ├── aclshmem_float_put_nbi(...)            → 非阻塞 float 写入远端
    ├── aclshmem_half_get_nbi(...)             → 非阻塞 half 读取远端
    └── aclshmem_bfloat16_get_nbi(...)         → 非阻塞 bf16 读取远端

zero_buffer_api.h (CAM 封装)
    └── #include "shmem.h"

zero_buffer_sync_flag.h (同步框架)
    └── ZeroBufferSyncFlag 类

// normal_zb
auto ptr = reinterpret_cast<__gm__ uint8_t *>(shmem_ptr(recvXGM_, dstRankId));
dstGT.SetGlobalBuffer((__gm__ XType *)(ptr + remoteReadAddr));

// lowlatency_zb
auto srcPtr = reinterpret_cast<__gm__ uint8_t *>(shmem_ptr(expandXGM_, srcRankId));
srcWinGMTensor.SetGlobalBuffer((__gm__ ExpandXType *)(srcPtr + hExpandXTypeSize_ * (remoteBase + remoteOffset)));

gva_gm + FLAG_AREA_BASE (100KB):
├── Magic Area (24KB = 48 cores × 512B)
│   ├── slot[0] @ 0×512B     ← core 0 的 magic
│   ├── slot[1] @ 1×512B     ← core 1 的 magic
│   └── ...
├── Barrier PingPong Buffer 0 (epWorldSize × 32B)
│   ├── slot[srcRank=0]       ← rank 0 写入的 barrier flag
│   ├── slot[srcRank=1]
│   └── ...
├── Barrier PingPong Buffer 1 (同上)
├── Flag PingPong Buffer 0 (epWorldSize × slotsPerRank × 32B)
│   ├── slot[srcRank=0][eventID=0]
│   ├── slot[srcRank=0][eventID=1]
│   └── ...
└── Flag PingPong Buffer 1 (同上)

1. 完成数据写入 (DMA 到远端 SHMEM)
2. SetFlag(destRank=B, eventID=i, phase)
   → MTE3 fence (flush 所有先前的写入)
   → 写 flag 到 B 的本地 SHMEM 的 slot [A][i]

1. WaitFlag(srcRank=A, eventID=i, phase)
   → 轮询本地 SHMEM 直到 flag 匹配 (magic, phase)
2. 返回后，rank A 的数据保证可见

DataWrite(A→B) ── MTE3 fence ──► SetFlag(A→B)
                                    │
                                    ▼
                              WaitFlag(B←A) ──► DataRead(B←A)

// 同步区域: metaDataGvaGM_ + COMBINE_SYNC_FLAG_BASE_OFFSET (100KB)
// 每个 rank 占 COMBINE_SYNC_STATE_OFFSET (32B)

// SetSyncFlag: 向远端 rank 的 flag 区域写入 1.0f
GetSyncFlagAddrByRankId(rankId) = shmem_ptr(metaDataGvaGM_, rankId) + 100KB + epRankId_ * 32B

// WaitSyncFlag: 轮询本 rank 的 flag 区域，直到所有远端 rank 都写入
// 使用 ReduceSum 求和，期望值 = rankNumPerBlock_

// master: moe_combine_normal.cpp
Process():
├── CopyBufferToShareAndSetStatus()    // 将本地 token 拷贝到远端 window
│   └── for each token:
│       ├── CopyBufferToShare()        // DataCopyPad 到远端 SHMEM
│       └── SetStatusBySrcInfo()       // 写 status flag
└── ReadBufferFromRemote()             // 读取远端数据
    └── for each token:
        ├── WaitBuffCopy()             // 轮询等待所有源 rank 数据到位
        └── ReadBufferAndWeightedSum() // 加权求和

// cam_feature: moe_combine_normal_zero_buffer.cpp
Process():
├── syncFlag_.IncrementMagic()         // 获取 magic
├── syncFlag_.BarrierAll()             // 全 rank 屏障
├── ReadBufferFromRemote()             // 直接读取远端 SHMEM
│   └── for each token:
│       └── ReadBufferAndWeightedSum() // 直接从远端 SHMEM 读取 + 加权
└── SyncAll<true>()

Rank A: CopyBufferToShare → 写入 Rank B 的 window
Rank B: WaitBuffCopy → 轮询等待 → 从本地 window 读取

Rank A: (dispatch 阶段已写入 SHMEM)
Rank B: shmem_ptr(recvXGM_, rankA) → 直接读取 Rank A 的 SHMEM

// 每个 (token, topk) 对有一个 32B status slot
GM_ADDR stateGM = GetStateAddrByRankId(srcRankId) + (srcTokenId * axisK_ + srcTopkId) * UB_32_ALIGN;
// 轮询直到所有 K 个 source 都到位
while (current != target) {
    DataCopy(stateTensorLocal, stateGMTensor, calCount);
    Sum(tempStateTensorLocal, stateTensorLocal, sumPerKParams);
    current = tempStateTensorLocal(0);
}

// 每个 rank 有一个 64B flag slot
// BarrierAll 一次等待所有 rank
syncFlag_.BarrierAll();

auto contextGM0 = AscendC::GetHcclContext<HCCL_GROUP_ID_0>();
epWinContext_ = (__gm__ HcclOpResParam *)contextGM0;
selfMagicTensor.SetGlobalBuffer(epWinContext_->localWindowsExp + MAGIC_WIN_OFFSET + coreIdx_ * WIN_512_ALIGN);
magic_ = selfMagicTensor(0);  // 0 或 1 翻转
selfMagicTensor(0) = ((magic_ == 0) ? 1 : 0);

magic_ = static_cast<int32_t>(magicGT(0));
magicGT(0) = magic_ + 1;  // 单调递增

// 1. 获取属性
GetAttrAndSetTilingData(context, *tilingData, nodeName);
// → 读取 meta_data_ptr, ep_world_size, ep_rank_id, moe_expert_num, global_bs

// 2. 校验输入输出
TilingCheckMoeCombineNormalZeroBuffer(context, nodeName, isEnableDiagnose);
// → 检查 dim/format/dataType

// 3. 校验属性
CheckAttrs(context, *tilingData, nodeName, localMoeExpertNum);
// → moeExpertNum % epWorldSize == 0
// → tp=2 时 localMoeExpertNum == 1
// → globalBs == 0 || globalBs >= bs * epWorldSize

// 4. 校验 shape
CheckTensorShape(context, *tilingData, nodeName, localMoeExpertNum);
// → k ∈ (0, min(16, moeExpertNum)]
// → h ∈ [1024, 7168]

// 5. 校验 win 区大小
tokenNeedSizeCombine = Ceil(h * 2 + 511, 512) * 512;  // 512 对齐
actualSize = (maxBs * k * tokenNeedSizeCombine + COMBINE_STATE_WIN_OFFSET + NOTIFY_DISPATCH_WIN_OFFSET) * 2;

// 6. 设置 blockDim 和 schedule mode
blockDim = ascendcPlatform.CalcTschBlockDim(aivNum, 0, aivNum);
context->SetScheduleMode(1);  // batch mode，所有核同时启动

// 1. 获取属性（更复杂）
GetAttrAndSetTilingData(context, *tilingData, nodeName, groupEp, groupTp, commQuantMode);
// → 额外读取 sharedExpertNum, sharedExpertRankNum, zeroExpertNum, copyExpertNum, constExpertNum

// 2. 计算 A（expandX 的 dim0）
if (isShared) {
    A = maxBs * maxSharedGroupNum;  // 共享专家卡
} else {
    A = globalBs * min(localMoeExpertNum, expertIdsDim1);  // MoE 专家卡
}

// 3. 校验 win 区大小
tokenNeedSizeCombine = Ceil(h * 2 + 511, 512) * 512;
tokenActualLen = Ceil(h * 2 + 31, 32) * 32 + 44;  // + scale(32B) + expandIdx(3*4B)
tokenNeedSizeDispatch = Ceil(tokenActualLen + 511, 512) * 512;
actualSize = ((maxBs * tokenNeedSizeDispatch * epWorldSize * localMoeExpertNum) +
              (maxBs * tokenNeedSizeCombine * (k + sharedExpertNum))) * 2;

// 4. 计算 TilingKey
tilingKey = 10000;
if (tpWorldSize == 2) tilingKey += 100;
if (commQuantMode == 2) tilingKey += 20;

// 5. 设置 Mc2CcTiling（HCCL 通信配置）
SetHCommCfg(context, tilingData, groupEp, groupTp);
// → AlltoAll: level0:fullmesh;level1:pairwise
// → ReduceScatter: level0:ring

actualSize = (maxBs × k × tokenNeedSizeCombine + COMBINE_STATE_WIN_OFFSET + NOTIFY_DISPATCH_WIN_OFFSET) × 2
           = (maxBs × k × Ceil(h×2, 512) + 5MB + 204KB) × 2

actualSize = ((maxBs × tokenNeedSizeDispatch × epWorldSize × localMoeExpertNum) +
              (maxBs × tokenNeedSizeCombine × (k + sharedExpertNum))) × 2

struct MoeCombineNormalZeroBufferInfo {
    uint32_t epWorldSize;        // EP 通信域大小
    uint32_t tpWorldSize;        // TP 通信域大小 (0 or 1 or 2)
    uint32_t epRankId;           // 本 rank ID
    uint32_t tpRankId;           // 本 TP rank ID
    uint32_t expertShardType;    // 专家分片类型
    uint32_t moeExpertNum;       // MoE 专家总数
    uint32_t moeExpertPerRankNum; // 每 rank 专家数
    uint32_t globalBs;           // 全局 batch size
    uint32_t bs;                 // 本 rank batch size
    uint32_t k;                  // top-K
    uint32_t h;                  // hidden size
    uint32_t aivNum;             // AIV 核数
    uint64_t totalUbSize;        // UB 总大小
    uint64_t totalWinSize;       // Window 总大小
    float armAvgFactor;          // ARM 平均因子
    float epsilon;               // 数值精度
    bool isEnableDiagnose;       // 是否启用诊断
    bool isGetProb;              // 是否获取概率梯度
};
struct MoeCombineNormalZeroBufferTilingData {
    MoeCombineNormalZeroBufferInfo moeCombineNormalInfo;
    uint64_t zeroBufferPtr;      // Zero-Buffer 对称地址
};

struct MoeCombineLowlatencyZeroBufferInfo {
    // ... (同上，额外字段)
    uint32_t sharedExpertNum;    // 共享专家数
    uint32_t sharedExpertRankNum; // 共享专家所在 rank 数
    uint32_t zeroExpertNum;      // 零专家数
    uint32_t copyExpertNum;      // 拷贝专家数
    uint32_t constExpertNum;     // 常量专家数
    bool isTokenMask;            // token 级 mask
    bool isExpertMask;           // expert 级 mask
    bool hasSharedExpertX;       // 是否有共享专家输入
    bool hasElasticInfo;         // 是否有弹性信息
    uint64_t zeroBufferPtr;      // Zero-Buffer 地址
};
struct MoeCombineLowlatencyZeroBufferTilingData {
    Mc2InitTiling mc2InitTiling;     // MC2 初始化 tiling
    Mc2CcTiling mc2CcTiling1;        // MC2 CC tiling 1 (AlltoAll)
    Mc2CcTiling mc2CcTiling2;        // MC2 CC tiling 2 (ReduceScatter)
    MoeCombineLowlatencyZeroBufferInfo moeCombineLowlatencyZeroBufferInfo;
};

// normal_zb 内存布局常量
constexpr uint64_t NOTIFY_MAGIC_OFFSET = 50UL * 1024UL;           // 50KB
constexpr uint64_t WIN_MAGIC_OFFSET = 100UL * 1024UL;             // 100KB
constexpr uint64_t HALF_WIN_STATE_OFFSET = 8UL * 1024UL * 1024UL; // 8MB
constexpr uint64_t COMBINE_WIN_STATE_OFFSET = 5UL * 1024UL * 1024UL; // 5MB
constexpr uint64_t WIN_ADDR_ALIGN = 512UL;

// lowlatency_zb 同步常量
constexpr uint64_t COMBINE_SYNC_FLAG_BASE_OFFSET = 100UL * 1024UL;  // 100KB
constexpr uint64_t COMBINE_SYNC_STATE_OFFSET = 32U;                  // 32B per rank

// 通用常量
constexpr uint32_t UB_32_ALIGN = 32U;
constexpr uint32_t MUL_256_ALIGN = 256U;
constexpr uint8_t DOUBLE_BUFFER = 2;

MoeCombineNormalZeroBuffer<RecvXType, XType, SrcInfoType>
├── 使用 ZeroBufferSyncFlag (通过 TBuf)
├── 使用 shmem_ptr() 访问远端 SHMEM
└── 使用 TQue/TBuf 管理 UB

MoeCombineLowlatencyZeroBuffer<ExpandXType, XType, ExpandIdxType, IsNeedReduceScatter, IsInt8Quant>
├── 自实现 SetSyncFlag/WaitSyncFlag
├── 使用 shmem_ptr() 访问远端 SHMEM
├── 使用 moe_lowlatency_zero_buffer_base.h 工具
└── 支持 TP=2, 共享专家, 弹性缩容, 常量专家