┌─────────────────────────────────────────────────────────────────────┐
│                    FusedDeepMoE 统一执行流程                         │
│                                                                     │
│  ┌──────────┐   ┌──────────┐   ┌─────────┐   ┌────────┐   ┌──────┐ │
│  │ Dispatch │──▶│   GMM1   │──▶│ SwiGLU  │──▶│  GMM2  │──▶│Combine│ │
│  │ (AIV)    │   │ (AIC+AIV)│   │(Epilogue│   │(AIC+   │   │(AIV)  │ │
│  │          │   │          │   │ fused)  │   │ AIV)   │   │       │ │
│  └──────────┘   └──────────┘   └─────────┘   └────────┘   └──────┘ │
│       │              │              │              │           │     │
│  AlltoAll      PerTokenDeq   PerTokenDeq    PerTokenDeq  AlltoAll  │
│  Send          + Swiglu      + Quant        + Combine    Recv     │
│  + Quant       + Quant                                               │
│                                                                     │
│  ─────── 中间数据全部在片上 workspace 流转，不访问 GM ───────         │
└─────────────────────────────────────────────────────────────────────┘

// 完整的 6 参数模板，支持灵活的 Scale 类型
template <typename ExpandXType, typename W1ScaleType, typename W2ScaleType,
          typename ExpandIdxType, bool IsNeedReduceScatter, uint32_t EXEC_FLAG>
class FusedDeepMoe { ... };
// 实例化: FusedDeepMoe<DTYPE_X, DTYPE_GMM1_WEIGHT_SCALE, DTYPE_GMM2_WEIGHT_SCALE, int32_t, false, TILING_KEY_VAR>

// 简化模板：去掉 Scale 类型参数，EXEC_FLAG 提升为第一参数
template <typename ExpandXType, typename ExpandIdxType,
          bool IsNeedReduceScatter, uint32_t EXEC_FLAG>
class FusedDeepMoeFwk { ... };
// 实例化: FusedDeepMoeFwk<DTYPE_X, int32_t, false, TILING_KEY_VAR>

fused_deep_moe/
├── op_kernel/
│   ├── fused_deep_moe.{h,cpp}          ← 主 kernel 入口 + FusedDeepMoe 类
│   ├── fused_deep_moe_base.h           ← 模板宏定义 + 通信组结构
│   ├── fused_deep_moe_tiling.{h,cpp}   ← TilingData 定义 + 编译期常量
│   ├── fused_deep_moe/                 ← 子模块目录
│   │   ├── gemm/                       ← GEMM block_mmad, kernel
│   │   ├── epilogue/                   ← Epilogue block, tile, policy
│   │   └── raw_distributed/            ← dispatch/combine 通信原语
│   └── cam_moe_distribute_*.h          ← 兼容性别名（实际指向子模块）
├── op_host/
│   └── fused_deep_moe_infer.cpp        ← 仅 InferShape/InferDataType
└── op_api/
    └── aclnn_fused_deep_moe.{h,cpp}    ← aclnn 接口封装

// fused_deep_moe.cpp
extern "C" __global__ __aicore__ void fused_deep_moe(...) {
    icache_preload(8);
    REGISTER_TILING_DEFAULT(FusedDeepMoeTilingData);
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2);  // 1C2V 混合调度
    GET_TILING_DATA(tiling_data, tiling);
    // 64 个 tiling key 分支（编译期展开）
    CALL_FUSED_DEEP_MOE_IF_TILINGKEY(0)
    CALL_FUSED_DEEP_MOE_ELIF_TILINGKEY(1)
    // ... 直至 key 63
}

GemmCoord gmm1ProblemShape{maxTokenNum_, gmm1OutputDim_, tokenHiddenSize_};  // [M, N, K]
GemmCoord gmm2ProblemShape{maxTokenNum_, gmm2OutputDim_, gmm2InputDim_};

shmemWorkspaceGM_ (Base 独有，SHMEM 双缓冲):
├── gmShareX1 / gmX1          — GMM1 输入 token (int8)
├── gmShareX2 / gmX2          — GMM2 输入 token (int8)
├── gmTokenFlag               — token 到达标志
├── gmAllExpertTokenNums      — 全局专家 token 计数
└── gmCombineSend             — Combine 发送缓冲

workspaceGM_ (普通 GM workspace):
├── gmCVSwap                  — Cube 计算 workspace (blockDim * L1M*L1N * STAGES * int32)
├── gmShareSwigluOut/gmSwigluOut — SwiGLU 输出 (float)
├── gmGroupList               — 专家分组列表
├── gmExpandIdx               — 展开索引
├── gmEpSendCount             — EP 发送计数
├── gmResvered                — 保留空间 (256KB)
└── gmAllEpRecvCount          — 全局 EP 接收计数

if constexpr ((EXEC_FLAG & EXEC_FLAG_DEEP_FUSE) == 0) {
    if constexpr (g_coreType == AscendC::AIV) {
        // AIV 上执行 Dispatch：量化 + AlltoAll 发送
        CamMoeDistributeDispatch<...> dispatcher;
        dispatcher.Init(gmX_, gmexpertIds_, ...);
        dispatcher.Process();
        icache_preload(8);  // 预取 GMM 指令到 icache
    }
    // AIV 完成 → CrossCoreFlag 通知 AIC
    PipeBarrier<PIPE_ALL>();
    Arch::CrossCoreFlag gmm1AivFinished{0};
    if constexpr (g_coreType == AscendC::AIV) {
        Arch::CrossCoreSetFlag<0x2, PIPE_MTE3>(gmm1AivFinished);
    } else {
        Arch::CrossCoreWaitFlag(gmm1AivFinished);
    }
}

GmmDeqSwigluQuant<TemplateMC2TypeFunc, Gmm1L1TileShape, Gmm1L0TileShape, ...>(
    gmm1ProblemShape, groupCount_, gmGroupList, gmX1, layoutX1, ...);

// 先初始化 Combiner（AIV 上）
CamMoeDistributeCombine<...> combiner;
combiner.Init(gmGmm2DepOut, gmexpertIds_, gmExpandIdx, ...);

// 再执行 GMM2 kernel（AIC+AIV），Combiner 作为回调注入
GmmDeq<TemplateMC2TypeFunc, Gmm2L1TileShape, ...>(
    gmm2ProblemShape, ..., &combiner);

// GMM1 的 GEMM 类型定义
using AType = Gemm::GemmType<int8_t, layout::RowMajor>;     // 激活 (量化后)
using BType = Gemm::GemmType<int8_t, layout::zN>;           // 权重 (zN 布局)
using CType = Gemm::GemmType<int32_t, layout::RowMajor>;    // 累加器

// Epilogue 类型
using ScaleType = Gemm::GemmType<W1ScaleType, layout::VectorLayout>;      // 权重 scale
using PerTokenScaleType = Gemm::GemmType<float, layout::VectorLayout>;    // per-token scale
using DType = Gemm::GemmType<float, layout::RowMajor>;                    // SwiGLU 输出

using GemmKernel = typename std::conditional<
    (EXEC_FLAG & EXEC_FLAG_DEEP_FUSE),
    // 深度融合：Dispatch 在 GMM kernel 内部完成（AIV 先做量化+发送，AIC 同时计算）
    Gemm::Kernel::GroupedMatmulSliceMPerTokenDequantSwigluQuantMultiStageWorkspace<...>,
    // 非深度融合：独立的 Dispatch + GMM
    Gemm::Kernel::GroupedMatmulSliceMPerTokenDequantSwigluQuantMultiStageWorkspaceWithShallowDispatch<...>
>::type;

// GMM1 Tile Shape
constexpr uint32_t GMM1_L1M = 256;   // L1 tile M 维
constexpr uint32_t GMM1_L1N = 128;   // L1 tile N 维
constexpr uint32_t GMM1_L1K = 512;   // L1 tile K 维
constexpr uint32_t GMM1_L0K = 128;   // L0 tile K 维 (L1→L0 切分)
constexpr uint32_t GMM1_EPIM = 64;   // Epilogue interleaving M

// GMM2 Tile Shape
constexpr uint32_t GMM2_L1M = 128;
constexpr uint32_t GMM2_L1N = 256;
constexpr uint32_t GMM2_L1K = 512;
constexpr uint32_t GMM2_L0K = 128;
constexpr uint32_t GMM2_EPIM = 32;

// Pipeline 深度
constexpr uint32_t WORKSPACE_STAGES = 4;       // workspace 4 级流水线
constexpr uint32_t CUSTOM_PRELOAD_STAGES = 1;  // preload 1 级
constexpr uint32_t CUSTOM_L1_STAGES = 2;       // L1 双缓冲
constexpr uint32_t CUSTOM_L0A_STAGES = 2;      // L0A 双缓冲
constexpr uint32_t CUSTOM_L0B_STAGES = 2;      // L0B 双缓冲
constexpr uint32_t GMM2_L1A_STAGES = 4;        // GMM2 L1A 4 级 (ResidentA)
constexpr uint32_t GMM2_L0A_STAGES = 4;        // GMM2 L0A 4 级

extern "C" __global__ __aicore__ void fused_deep_moe(
    // ... 同 Base 的 14 个输入 ...
    GM_ADDR gmm1_bias, GM_ADDR gmm2_bias,           // ← 新增
    GM_ADDR share_gmm1_bias, GM_ADDR share_gmm2_bias, // ← 新增
    // ... 输出 + workspace + tiling ...
)

gmBias1_ = gmm1_bias;
gmBias2_ = gmm2_bias;
gmShareBias1_ = share_gmm1_bias;
gmShareBias2_ = share_gmm2_bias;

// W4A8: 权重和激活均为 INT4
using AType = Gemm::GemmType<AscendC::int4b_t, layout::RowMajor>;
using BType = Gemm::GemmType<AscendC::int4b_t, layout::zN>;
using CType = Gemm::GemmType<int32_t, layout::RowMajor>;

// 布局使用 int4b_t
layout::zN layoutWeight1 = layout::zN::template MakeLayout<AscendC::int4b_t>(tokenHiddenSize_, gmm1OutputDim_);

constexpr int CONSTANT_INT4_PACKING_RATIO = 2;  // 2 个 INT4 打包进 1 个 INT8

// L1/L0 tile 大小计算（注意除以打包比）
static constexpr uint32_t L1A_TILE_SIZE =
    L1TileShape::M * L1TileShape::K * sizeof(ElementA) / CONSTANT_INT4_PACKING_RATIO;
static constexpr uint32_t L1B_TILE_SIZE =
    L1TileShape::N * L1TileShape::K * sizeof(ElementB) / CONSTANT_INT4_PACKING_RATIO;

// W4A8 tiling.h
constexpr uint32_t GMM1_L1K = 1024;  // Base 是 512
constexpr uint32_t GMM1_L0K = 256;   // Base 是 128
constexpr uint32_t GMM2_L1K = 1024;
constexpr uint32_t GMM2_L0K = 256;

// W4A8 workspace 计算
size_t CVSwapBufferSize = CeilUp(aicNum * L1_TILE_BYTE_SIZE * CUBE_WORKSPACE_STAGE * sizeof(int32_t), GM_ALIGN_SIZE);
workSpaces[0] = SYSTEM_NEED_WORKSPACE + usrSize;  // SYSTEM_NEED_WORKSPACE = 16MB

// W4A8 tiling.cpp
uint64_t tilingKey = 0;
if (xActiveMaskEnable) tilingKey |= EXEC_FLAG_X_ACTIVE_MASK;
if (calShareExpert) tilingKey |= EXEC_FLAG_SHARED_EXPERT;
if (moeExpertNumPerRank != 1) tilingKey |= EXEC_FLAG_DEEP_FUSE;
if (isTensorList) tilingKey |= EXEC_FLAG_TENSOR_LIST;
if (expertSmoothScalesExist) tilingKey |= EXEC_FLAG_SMOOTH_QUANT;
context.SetTilingKey(tilingKey);
context.SetBlockDim(aicNum);

// FWK kernel 入口 — 无 share_gmm1/2_weight，无 x_active_mask
extern "C" __global__ __aicore__ void fused_deep_moe_fwk(
    GM_ADDR x, GM_ADDR expert_ids,
    GM_ADDR gmm1_permuted_weight, GM_ADDR gmm1_permuted_weight_scale,  // ← permuted
    GM_ADDR gmm2_weight, GM_ADDR gmm2_weight_scale,
    GM_ADDR expert_smooth_scales, GM_ADDR expert_scales,               // ← 无 share_smooth
    GM_ADDR output, GM_ADDR expertTokenNums,                           // ← 无 share_output
    GM_ADDR workspace, GM_ADDR tiling)

sharedExpertNum_ = tilingData->...sharedExpertNum;
sharedExpertRankNum_ = tilingData->...sharedExpertRankNum;

bool isShareExpert = (epRankId_ < sharedExpertRankNum_);
if (isShareExpert) {
    m_ = maxBs_ * epRankSize_ / sharedExpertRankNum_;  // 共享专家处理全部 token
} else {
    m_ = maxBs_ * epRankSize_ * min(topK_, moeExpertNumPerRank_);  // 路由专家
}
groupCount_ = isShareExpert ? 1 : moeExpertNumPerRank_;  // 共享专家只有 1 组

lib_ops.h
├── 工具函数
│   ├── PipeSync<EVENT>()          — 管道同步
│   └── GetBufferByByte<T>()      — UB 缓冲区分配
├── LibOps namespace
│   ├── MoEContext                — MoE 全局上下文（窗口地址、状态管理）
│   ├── OpDispatchSendToken       — Dispatch 发送端（量化 + 窗口写入）
│   ├── OpDispatchRecvToken       — Dispatch 接收端（窗口读取 + 反量化）
│   ├── OpDequant                 — 反量化操作（INT32 → float）
│   ├── OpSwiglu                  — SwiGLU 激活
│   ├── OpSyncAllCleanAndUpdate   — 全局同步 + 状态清零
│   ├── OpBlockQuant              — 分块量化（float → INT8）
│   ├── OpCombineSend             — Combine 发送端（窗口写入）
│   ├── OpCombineSetFlag          — Combine 标志设置
│   ├── OpCombineRecv             — Combine 接收端（ReducePermute）
│   ├── GmScaleReader             — GM 标量读取器
│   ├── GmPerTokenScaleReader     — GM per-token scale 读取器
│   └── GmSwigluOutputWriter      — SwiGLU 输出写入器

struct MoEContext {
    __gm__ HcclOpResParam *winContext;   // HCCL 窗口上下文
    GM_ADDR gmStatusDataSpace;           // 状态数据空间
    uint32_t epRankSize, epRankId;
    uint32_t aicNum, aivNum;
    uint32_t moeExpertNumPerRank;
    int32_t dataState;      // 双缓冲状态 (0/1)
    int32_t tokenFlag;      // token 到达标志 (0x55555555 / 0x33333333)
    int32_t vToCFlag;       // AIV→AIC 通信标志
    float sumTarget;        // ReduceSum 目标值（用于 flag 检测）
    uint64_t expertPerSizeOnWin;  // 单个专家在窗口中的空间大小
    uint64_t winDataSizeOffset;   // 窗口数据偏移（双缓冲）
};

GM_ADDR GetWindowStateAddrByRankId(uint32_t rankId) {
    if (epRankId == rankId) return (GM_ADDR)winContext->localWindowsExp;
    return (GM_ADDR)((HcclRankRelationResV2 *)winContext->remoteRes[rankId].nextDevicePtr)->windowsExp
           + dataState * WIN_STATE_OFFSET;  // 双缓冲偏移
}

GM_ADDR GetWindowDataAddrByRankId(uint32_t rankId) {
    // ... 类似，但使用 localWindowsIn + winDataSizeOffset
}

template <typename XType, uint32_t BUFFER_NUM, bool ENABLE_X_ACTIVE_MASK>
class OpDispatchSendToken {
    // 1. TokenActiveMaskCal: 计算有效 token 数（可选）
    // 2. LoadExpertIds: 加载 expert_ids 到 UB
    // 3. CalAndSendTokenCount: 计算每个专家的 token 数并写入远端状态空间
    // 4. SendToMoeExpert: 量化 token 并写入远端窗口
    //    - DataCopy: GM → UB (读取 token)
    //    - QuantToken: Cast to float → Abs → ReduceMax → scale → Cast to int8
    //    - DataCopy: UB → GM (写入远端窗口)
};

void QuantToken(...) {
    Cast(xFp32TmpTensor, xInTensor, CAST_NONE, tokenLength);  // bf16/fp16 → float
    Abs(xFp32AbsTensor, xFp32TmpTensor, tokenLength);
    ReduceMax(xRowMaxTensor, xFp32AbsTensor, ..., tokenLength, false);
    float dynamicQuantScale = 127.0 / xRowMaxTensor.GetValue(0);
    // 写入 scale 到 token 尾部的 extra space
    yFp32Tensor.SetValue(tokenLength / sizeof(float), 1.0 / dynamicQuantScale);
    yInt32Tensor.SetValue(tokenLength / sizeof(int32_t) + 1, ctx.tokenFlag);  // 写入 flag
    Muls(xFp32TmpTensor, xFp32TmpTensor, dynamicQuantScale, tokenLength);
    Cast(ytmpInt32Tensor, xFp32TmpTensor, CAST_RINT, tokenLength);  // float → int32
    Cast(yHalfTensor, ytmpInt32Tensor, CAST_ROUND, tokenLength);    // int32 → fp16
    Cast(yInt8Tensor, yHalfTensor, CAST_TRUNC, tokenLength);        // fp16 → int8
}

class OpDispatchRecvToken {
    // 1. RecvTokenCountInfo: 轮询状态空间直到所有源 rank 的 flag 到达
    //    - GatherMask: 从 512B 状态块中提取 flag 和 count
    //    - ReduceSum: 检测 flag 总和是否达到目标值
    // 2. RecvToken: 从窗口读取 token 并写入本地 GM
    //    - 轮询 token flag 直到到达
    //    - DataCopy: GM → UB → GM (窗口 → 本地)
    // 3. NotifyNextStage: AtomicAdd 通知 AIC 该组就绪
};

// 使用 GatherMask 从 256B 状态块中提取 flag
// 每个 rank 的状态: [flag(0), count(1), pad(2), pad(3)] × 8 blocks = 256B
// src1Pattern=3: 每 4 个 int32 取第 1 个（flag）
// src1Pattern=4: 每 4 个 int32 取第 2 个（count）
do {
    DataCopy(statusFp32Tensor, windowInStatusFp32Tensor, dataCopyParams);
    GatherMask(gatherMaskOutFp32Tensor, statusFp32Tensor, src1Pattern=3, ...);
    ReduceSum(gatherMaskOutFp32Tensor, ..., recStatusNumPerCore);
} while (!(minTarget < sumOfFlag && sumOfFlag < maxTarget));

// OpDequant: INT32 → float (per-token dequant)
// 输入: cQue(int32), scaleQue(float), perTokenScaleQue(float)
// 输出: outQue(float)
// 计算: float = int32 * weightScale * perTokenScale

// OpSwiglu: SwiGLU 激活
// 输入: inQue(float)
// 输出: outQue(float)
// 计算: swiglu(x) = x[:, :N/2] * sigmoid(x[:, :N/2]) * x[:, N/2:]
//   = x_left * (1 / (1 + exp(-x))) * x_right

// OpBlockQuant: float → INT8 (per-token block quant)
// 输入: inQue(float)
// 输出: dequantScaleQue(float), outQue(int8)
// 计算: abs → Max(左右两半) → scale = Max/127 → quant

// OpCombineSend: 将 GMM2 输出写入远端窗口
// 使用 epSendCount 前缀和确定每个 rank 的 token 范围
// DataCopyPad: UB → GM (远端窗口)

// OpCombineRecv: 调用 combiner->ReducePermute()
// ReducePermute 完成 AlltoAll 接收 + 专家输出累加 + 写回输出

template <uint32_t EXEC_FLAG, typename XType_, ...>
CATLASS_DEVICE void GmmDeqSwigluQuant(
    GemmCoord problemShape, uint32_t groupCount, GM_ADDR gmGroupList,
    GM_ADDR gmA, layout::RowMajor layoutA,
    GM_ADDR gmB, layout::zN layoutB,
    GM_ADDR gmScale, layout::VectorLayout layoutScale,
    GM_ADDR gmPerTokenScale, layout::VectorLayout layoutPerTokenScale,
    GM_ADDR gmD, layout::RowMajor layoutD,
    GM_ADDR gmDequantScale, layout::VectorLayout layoutDequantScale,
    GM_ADDR gmWorkspace,
    GM_ADDR gmX, GM_ADDR debugGm,                    // ← debugGm 替代 smoothScales
    GM_ADDR gmexpertIds, GM_ADDR gmExpandIdx,
    GM_ADDR gmEpSendCount, GM_ADDR gmResvered,
    GM_ADDR gmExpertTokenNums,
    uint32_t epRankSize, uint32_t epRankId,          // ← 标量参数替代 tiling 引用
    uint32_t moeExpertNum, uint32_t moeExpertNumPerRank,
    uint32_t sharedExpertNum, uint32_t sharedExpertRankNum,  // ← 新增
    uint32_t quantMode, uint32_t globalBs, uint32_t bs,
    uint32_t topK, uint32_t tokenLen)

using BType = Gemm::GemmType<int8_t, layout::nZ>;  // FWK: nZ
// vs Base:
using BType = Gemm::GemmType<int8_t, layout::zN>;  // Base: zN

AIV: Dispatch (量化 + AlltoAll 发送) ──────────────────────▶
AIC:  等待 ◀── CrossCoreFlag ── 等待结束 ──▶ GMM1+Swiglu+Quant ──▶ GMM2+Combine

AIV: [Dispatch 量化+发送] ─┐
                           ├─▶ 在 GMM Kernel 内部流水线化
AIC: [GMM1 计算] ──────────┘     [SwiGLU+Quant] ──▶ [GMM2+Combine]

路由专家 Rank:
  Dispatch → GMM1_route → SwiGLU → GMM2_route → Combine

共享专家 Rank (epRankId < sharedExpertRankNum):
  GMM1_share (全部 token) → SwiGLU → GMM2_share → 直接写回

bool isShareExpert = (epRankId_ < sharedExpertRankNum_);
if (isShareExpert) {
    m_ = maxBs_ * epRankSize_ / sharedExpertRankNum_;  // 处理全部 token
} else {
    m_ = maxBs_ * epRankSize_ * min(topK_, moeExpertNumPerRank_);
}
groupCount_ = isShareExpert ? 1 : moeExpertNumPerRank_;

constexpr uint32_t STATE_OFFSET = 512;        // 每个 rank 的状态空间偏移
constexpr uint32_t STATE_SIZE = 1024 * 1024;  // 1MB 状态空间
constexpr uint64_t WIN_STATE_OFFSET = 512 * 1024;  // 窗口状态双缓冲偏移
constexpr uint64_t STATE_WIN_OFFSET = 900 * 1024;  // 状态在窗口中的偏移
constexpr uint32_t SELF_STATE_OFFSET = 256 * 1024; // 自状态偏移

窗口状态空间 (每个 rank 1MB):
├── [0, 512KB):     数据状态空间 (dataState=0)
│   ├── rank0: [0, 512B)     — 每个 rank 512B 状态块
│   │   ├── [0, 256B):  flag + count 数组 (64 个 int32)
│   │   └── [256B, 512B): 保留
│   ├── rank1: [512B, 1024B)
│   └── ...
├── [512KB, 1024KB): 数据状态空间 (dataState=1) — 双缓冲
├── [900KB, 932KB):  窗口状态空间 (STATE_WIN_OFFSET)
├── [932KB, 964KB):  组 token 计数 (GROUP_TOKEN_NUM_OFFSET)
└── [964KB, 988KB):  软同步空间 (SOFT_SYNC_OFFSET)

GM_ADDR GetWindAddrByRankId(uint8_t ctxIdx, const int32_t rankId) {
    if (curRankId == rankId)
        return (GM_ADDR)(winContext_[ctxIdx]->localWindowsIn) + winDataSizeOffset_ + rankId * OPT_RANK_OFFSET;
    return (GM_ADDR)(((HcclRankRelationResV2 *)(winContext_[ctxIdx]->remoteRes[rankId].nextDevicePtr))->windowsIn)
           + winDataSizeOffset_ + rankId * OPT_RANK_OFFSET;
}

窗口数据空间 (每个 rank):
├── expert0: [0, expertPerSizeOnWin)
│   ├── token0: [0, hCommuSize)  — token 数据 + extra space (scale + flag)
│   ├── token1: [hCommuSize, 2*hCommuSize)
│   └── ...
├── expert1: [expertPerSizeOnWin, 2*expertPerSizeOnWin)
└── ...

hCommuSize = tokenLength * sizeof(int8_t) + TOKEN_EXTRA_SPACE
TOKEN_EXTRA_SPACE = 512B (存储 scale + flag)

struct CombineCalcInfo {
    uint64_t expertPerSizeOnWin_;    // 单个专家在窗口中的空间
    uint32_t epRankId_;              // 本 rank ID
    uint32_t epWorldSize_;           // EP 世界大小
    uint32_t moeExpertPerRankNum_;   // 每 rank 专家数
    uint32_t axisH_;                 // 隐藏维度
    uint32_t moeSendNum_;            // MoE 发送数
    GM_ADDR epSendCount_;            // EP 发送计数
    __gm__ HcclOpResParam *epWinContext_;  // EP 窗口上下文
    uint64_t winDataSizeOffset_;     // 窗口数据偏移
    uint32_t aivNum_;                // AIV 数量
    uint32_t moeExpertNum_;          // 总专家数
    GM_ADDR metaInfoGm_;             // 元信息 GM 地址
    GM_ADDR allExpertTokenNums_;     // 全局专家 token 数
    GM_ADDR allEpRecvCount_;         // 全局 EP 接收计数
    GM_ADDR combineSend_;            // Combine 发送缓冲
};

__aicore__ GM_ADDR GetShmemAddrByRankId(GM_ADDR baseAddr, int64_t rankId) {
    if (epRankId_ == rankId) return baseAddr;
    return (GM_ADDR)aclshmem_ptr(baseAddr, rankId);  // SHMEM 指针转换
}

__aicore__ GM_ADDR GetShmemStateAddrByRankId(int64_t rankId) {
    return GetShmemAddrByRankId(metaInfoGm_ + statusDataSpaceOffset_ + dataState_ * WIN_STATE_OFFSET, rankId);
}

void ShmemCleanUp() {
    if (coreIdx_ != 0) return;  // 仅 core 0 执行
    // 清零本地 notify 数据
    Duplicate(cleanTempLt, 0, moeExpertNum_);
    DataCopyPad(localNotifyDataTensor, cleanTempLt, ...);
}

// 1. AllToAllSend: 将 GMM2 输出发送到远端 rank
void AllToAllSend() {
    // 读取 epSendCount 确定每个 rank 的 token 范围
    // DataCopyPad: UB → GM (远端窗口)
}

// 2. ReducePermute: 接收远端 rank 的结果并累加
void ReducePermute() {
    // 从窗口读取 token
    // 按 expert 累加 (CustomAdd)
    // 写回输出 GM
}

// 3. ProcessCombine: 主流程
void ProcessCombine() {
    AllToAllSend();
    ReducePermute();
}

fused_deep_moe/op_kernel/fused_deep_moe/
├── gemm/
│   ├── block/
│   │   ├── block_mmad.h                              ← BlockMmad 通用实现
│   │   ├── block_mmad_preload_async_with_callback_resident_a.h  ← GMM2 ResidentA
│   │   └── (W4A8) block_mmad_preload_async_with_callback_w4a8.hpp
│   ├── kernel/
│   │   ├── grouped_matmul_slice_m_per_token_dequant_multistage_workspace.h      ← GMM2 kernel
│   │   └── grouped_matmul_slice_m_per_token_dequant_swiglu_quant_multistage_workspace.h  ← GMM1 kernel
│   └── dispatch_policy.h
├── epilogue/
│   ├── block/
│   │   ├── block_epilogue.h                          ← BlockEpilogue 通用实现
│   │   └── (W4A8) block_epilogue_per_token_dequant.hpp / _swiglu.h
│   ├── tile/
│   │   ├── tile_broadcast_mul.hpp                    ← 行广播乘
│   │   ├── tile_broadcast_one_blk.hpp                ← 单块广播
│   │   ├── tile_swizzle.hpp                          ← Tile 调度
│   │   └── (FWK) tile_stride_muls.h / tile_stride_binary.h
│   └── dispatch_policy.h
└── raw_distributed/
    ├── cam_moe_distribute_dispatch.h
    └── cam_moe_distribute_combine.h

template <class DispatchPolicy, class L1TileShape, class L0TileShape,
          class AType, class BType, class CType, class BiasType,
          class TileCopy, class TileMmad>
struct BlockMmad {
    // 1. InitL1: 初始化 L1 双缓冲
    // 2. InitL0A/InitL0B/InitL0C: 初始化 L0 缓冲
    // 3. operator(): 主计算循环
    //    - K 轴分 tile → GM→L1 预取 → L1→L0 搬运 → TileMmad
    // 4. SynchronizeBlock: 排空流水线
};

K 轴循环:
  for kTileIdx in [0, kTileCount):
    1. GM→L1: copyGmToL1A/L1B (MTE2)
    2. 如果 preloadCount == PRELOAD_STAGES:
       L1TileMmad(preload 最早的 tile)  ← 计算与预取重叠
    3. 存储当前参数到 l1TileMmadParamsList
    4. 更新 preloadCount / l1ListId

L1TileMmad (内部):
  for mPartIdx, kPartIdx, nPartIdx:
    L1→L0A: copyL1ToL0A (MTE1)
    L1→L0B: copyL1ToL0B (MTE1)
    TileMmad: 硬件乘累加 (FIX)
    if isKLoopLast:
      callbackBeforeFixpipe()
      copyL0CToGm (MTE3)
      callbackAfterFixpipe()  ← 触发 Epilogue

template <class EpilogueDispatchPolicy, class CType, class ScaleType,
          class PerTokenScaleType, class DType, class TileRowBroadcastMul,
          class TileBroadcastOneBlk, class TileOneBlkColumnBroadcastMul,
          class TileCopy, class TileScheduler>
struct BlockEpilogue {
    // 在 L0C→GM 写入前/后执行:
    // 1. PerTokenDequant: INT32 * weightScale * perTokenScale → float
    // 2. SwiGLU: x * sigmoid(x) * x_right (half-glu)
    // 3. BlockQuant: float → INT8 + dequantScale
};

L0C (int32) → Cast → float
  → RowBroadcastMul(weightScale)        // 乘权重 scale
  → BroadcastOneBlk(perTokenScale)      // 乘 per-token scale
  → SwiGLU:                             // 激活函数
      left = x[:, :N/2]
      right = x[:, N/2:]
      silu = left * sigmoid(left)
      out = silu * right
  → BlockQuant:                         // 量化为 INT8
      abs → rowMax → scale = rowMax/127
      quant = round(x * 127/rowMax)
  → GM (int8) + dequantScale (float)

// 深度融合: Dispatch 在 GMM kernel 内部
template <TemplateMC2TypeClass, class BlockMmad, class BlockEpilogue,
          class BlockScheduler, uint32_t WORKSPACE_STAGES, class ElementGroupList>
struct GroupedMatmulSliceMPerTokenDequantSwigluQuantMultiStageWorkspace {
    // AIV 在 GMM1 计算前完成 Dispatch (量化 + AlltoAll 发送)
    // AIC 同时计算 GMM1
    // 通过 workspace 流水线实现计算-通信重叠
};

// 非深度融合: 独立的 Dispatch + GMM
template <TemplateMC2TypeClass, class BlockMmad, class BlockEpilogue,
          class BlockScheduler, uint32_t WORKSPACE_STAGES, class ElementGroupList>
struct GroupedMatmulSliceMPerTokenDequantSwigluQuantMultiStageWorkspaceWithShallowDispatch {
    // Dispatch 已在外部完成，直接执行 GMM1
};

template <class ArchTag>
class BlockQuant {
    // 输入: SwiGLU 输出 (float)
    // 输出: INT8 token + dequantScale
    // 计算:
    //   1. Abs + ReduceMax → rowMax
    //   2. scale = rowMax / 127
    //   3. quant = Cast(Cast(Cast(x * 127/scale, RINT), ROUND), TRUNC)
};

using Gmm2DispatchPolicy =
    Gemm::MmadAtlasA2PreloadAsyncWithCallbackResidentA<
        CUSTOM_PRELOAD_STAGES, GMM2_L1A_STAGES, GMM2_L1B_STAGES,
        GMM2_L0A_STAGES, GMM2_L0B_STAGES, CUSTOM_L0C_STAGES,
        CUSTOM_ENABLE_UNIT_FLAG, CUSTOM_ENABLE_SHUFFLE_K>;

// fused_deep_moe.h Init()
if constexpr (EXEC_FLAG & EXEC_FLAG_ZERO_BUFFER) {
    shmemWorkspaceGM_ = (GM_ADDR)(tilingData->disGmmDeqSwigluQuantGmmDeqComInfo.shmemWorkspacePtr);
}

shmemWorkspaceGM_:
├── gmShareX1 / gmX1          — GMM1 输入 token (int8)
├── gmShareX2 / gmX2          — GMM2 输入 token (int8)
├── gmShareX1Scale / gmX1Scale — GMM1 输入 scale
├── gmShareX2Scale / gmX2Scale — GMM2 输入 scale
├── gmTokenFlag               — token 到达标志
├── gmAllExpertTokenNums      — 全局专家 token 计数
└── gmCombineSend             — Combine 发送缓冲

// combine 中的 SHMEM 地址计算
__aicore__ GM_ADDR GetShmemAddrByRankId(GM_ADDR baseAddr, int64_t rankId) {
    if (epRankId_ == rankId) return baseAddr;
    return (GM_ADDR)aclshmem_ptr(baseAddr, rankId);  // SHMEM 指针转换
}

// dataState 在 {0, 1} 之间交替
winDataSizeOffset_ = dataState * epRankSize * expertPerSizeOnWin * moeExpertNumPerRank;

// 窗口地址 = base + winDataSizeOffset_ + rankId * OPT_RANK_OFFSET
GM_ADDR GetWindowDataAddrByRankId(uint32_t rankId) {
    return baseAddr + winDataSizeOffset_ + rankId * WIN_OPT_RANK_OFFSET;
}

void ShmemCleanUp() {
    if (coreIdx_ != 0) return;  // 仅 core 0 执行
    // 清零本地 notify 数据
    AscendC::LocalTensor<int32_t> cleanTempLt = cleanUpBuf_.Get<int32_t>();
    AscendC::Duplicate(cleanTempLt, (int32_t)0, moeExpertNum_);
    // DataCopyPad 到远端窗口的 notify 区域
    DataCopyPad(localNotifyDataTensor, cleanTempLt, ...);
}

struct FusedDeepMoeInfo {
    uint32_t epRankSize;           // EP rank 数量
    uint32_t epRankId;             // 本 rank ID
    uint32_t moeExpertNum;         // 总 MoE 专家数
    uint32_t moeExpertNumPerRank;  // 每 rank 专家数
    uint32_t quantMode;            // 量化模式
    uint32_t globalBs;             // 全局 batch size
    uint32_t bs;                   // 本 rank batch size
    uint32_t k;                    // topK
    uint32_t h;                    // 隐藏维度
    uint32_t aicNum;               // AIC 数量
    uint32_t aivNum;               // AIV 数量
    uint64_t totalUbSize;          // UB 总大小
    uint64_t totalWinSize;         // 窗口总大小
    uint64_t gmm1HLen;             // GMM1 隐藏维度
    uint64_t shareGmm1HLen;        // 共享专家 GMM1 隐藏维度
    uint64_t shmemWorkspacePtr;    // SHMEM workspace 地址 (Base 独有)
    uint64_t metaInfoPtr;          // 元信息地址 (Base 独有)
    bool isTensorList;             // 权重是否为 TensorList
};

struct FusedDeepMoeFwkInfo {
    uint32_t epRankSize;
    uint32_t epRankId;
    uint32_t moeExpertNum;
    uint32_t moeExpertNumPerRank;
    uint32_t sharedExpertNum;      // ← FWK 独有：共享专家数
    uint32_t sharedExpertRankNum;  // ← FWK 独有：共享专家 rank 数
    uint32_t quantMode;
    uint32_t globalBs;
    uint32_t bs;
    uint32_t k;
    uint32_t h;
    uint32_t aicNum;
    uint32_t aivNum;
    uint64_t totalUbSize;
    uint64_t totalWinSize;
    uint64_t gmm1HLen;
    bool isTensorList;
    // 注意: 无 shmemWorkspacePtr / metaInfoPtr
};

// Base/W4A8
struct FusedDeepMoeTilingData {
    Mc2InitTiling mc2InitTiling;    // MC2 初始化 tiling
    Mc2CcTiling mc2CcTiling;        // MC2 通信 tiling
    FusedDeepMoeInfo disGmmDeqSwigluQuantGmmDeqComInfo;  // 算子信息
};

// FWK
struct FusedDeepMoeFwkTilingData {
    Mc2InitTiling mc2InitTiling;
    Mc2CcTiling mc2CcTiling;
    FusedDeepMoeFwkInfo disGmmDeqSwigluQuantGmmDeqComInfo;
};

HCCL Window (每个 rank):
├── DataSpace (数据空间)
│   ├── rank0: [0, expertPerSizeOnWin * moeExpertNumPerRank)
│   │   ├── expert0: [0, expertPerSizeOnWin)
│   │   │   ├── token0: [0, hCommuSize)  — int8 token + 512B extra
│   │   │   └── ...
│   │   └── ...
│   └── ...
├── ExpSpace (状态空间, 1MB)
│   ├── [0, 512KB):     dataState=0 状态
│   │   ├── rank0: [0, 512B)
│   │   │   ├── [0, 256B):  [flag0, count0, pad, pad, flag1, count1, ...]
│   │   │   └── [256B, 512B): 保留
│   │   └── ...
│   ├── [512KB, 1024KB): dataState=1 状态 (双缓冲)
│   ├── [900KB, 932KB):  窗口状态 (STATE_WIN_OFFSET)
│   ├── [932KB, 964KB):  组 token 计数 (GROUP_TOKEN_NUM_OFFSET)
│   └── [964KB, 988KB):  软同步空间 (SOFT_SYNC_OFFSET)

fused_deep_moe()
  └── FusedDeepMoe::Process()
       ├── [if !DEEP_FUSE] CamMoeDistributeDispatch::Init/Process()  [AIV]
       │    ├── SendToMoeExpert()        — 量化 + AlltoAll 发送
       │    ├── AlltoAllDispatch()        — 窗口数据搬运
       │    └── ResetStatus()            — 状态清零
       ├── PipeBarrier<PIPE_ALL>()
       ├── CrossCoreFlag: AIV→AIC 同步
       ├── GmmDeqSwigluQuant()           — GMM1 + SwiGLU + Quant
       │    └── GroupedMatmulSliceMPerTokenDequantSwigluQuantMultiStageWorkspace
       │         ├── BlockMmad::operator()     — GM→L1→L0→TileMmad
       │         ├── BlockEpilogue::Process()  — Dequant + SwiGLU
       │         └── BlockQuant::operator()    — float→INT8
       ├── PipeBarrier<PIPE_ALL>()
       ├── CrossCoreFlag: AIV→AIC 同步
       ├── CamMoeDistributeCombine::Init()  [AIV]
       └── GmmDeq()                      — GMM2 + Combine
            └── GroupedMatmulSliceMPerTokenDequantMultiStageWorkspace
                 ├── BlockMmad::operator()     — ResidentA 策略
                 └── callbackAfterFixpipe → combiner.AllToAllSend()

fused_deep_moe_fwk()
  └── FusedDeepMoeFwk::Process()
       ├── [if !DEEP_FUSE] CamMoeDistributeDispatch::Init/Process()  [AIV]
       ├── PipeBarrier<PIPE_ALL>()
       ├── CrossCoreFlag: AIV→AIC 同步
       ├── GmmDeqSwigluQuant<EXEC_FLAG, ExpandXType, ...>()
       │    └── GroupedMatmulSliceMPerTokenDequantSwigluQuantMultiStageWorkspace
       │         └── (内部使用 LibOps 的 Op 类)
       ├── PipeBarrier<PIPE_ALL>()
       ├── CrossCoreFlag: AIV→AIC 同步
       ├── CamMoeDistributeCombine::Init()  [AIV]
       └── GmmDeq()                      — GMM2 + Combine
            └── callbackAfterFixpipe → combiner.AllToAllSend()
                 └── OpCombineSend/OpCombineRecv (LibOps)