From 821027113bb804cf36f4df29c27f92608926edf6 Mon Sep 17 00:00:00 2001 From: cloudforge1 Date: Mon, 30 Mar 2026 12:52:56 +0200 Subject: [PATCH] =?UTF-8?q?=E3=80=90Hackathon=2010th=20Spring=20No.47?= =?UTF-8?q?=E3=80=91=E8=A1=A5=E5=85=85=20H10=20=E5=AE=9E=E7=8E=B0=E6=9B=B4?= =?UTF-8?q?=E6=96=B0=E8=87=B3=E5=B7=B2=E5=90=88=E5=B9=B6=20RFC?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .../20250916_add_minimax_m1_for_fastdeploy.md | 415 +++++++++--------- 1 file changed, 207 insertions(+), 208 deletions(-) diff --git a/rfcs/FastDeploy/20250916_add_minimax_m1_for_fastdeploy.md b/rfcs/FastDeploy/20250916_add_minimax_m1_for_fastdeploy.md index 484fcd4de..7683711b5 100644 --- a/rfcs/FastDeploy/20250916_add_minimax_m1_for_fastdeploy.md +++ b/rfcs/FastDeploy/20250916_add_minimax_m1_for_fastdeploy.md @@ -1,222 +1,221 @@ -## **项目：在 FastDeploy 中原生支持 MiniMax-M1** +# FastDeploy 新增 MiniMax-M1 模型 -**目标**: 实现一个高性能、功能完整的 MiniMax-M1 推理后端，支持其混合注意力、MoE、DeepNorm 及超长上下文特性。 +| 任务名称 | 【Hackathon 10th Spring No.47】为 FastDeploy 新增 MiniMax-M1 模型 | +|------|------| +| 提交作者 | cloudforge1（基于 ZhijunLStudio 原始 RFC 更新） | +| 提交时间 | 2025-09-16（初版）/ 2026-07-08（模板重组） | +| 版本号 | V2.0 | +| 文件名 | 20250916_add_minimax_m1_for_fastdeploy.md | -**核心技术路径**: -1. **复用**: 最大化复用 GLM-4.5 PR 中已有的 Partial RoPE 和标准 GQA Attention 组件。 -2. **翻译与开发**: 将 vLLM 的 `lightning_attn.py` (Triton) 翻译为高性能的 CUDA C++ 算子，以支持 MiniMax-M1 的线性注意力层。 -3. **集成**: 在 Python 层构建完整的 MiniMax-M1 模型结构，并正确调度新开发的算子。 +# 一、概述 -**核心策略**: 并行推进 CUDA 开发与 Python 集成，以 5 层模型快速验证核心算子，然后无缝扩展到 80 层全量模型。 ---- +## 1、相关背景 -### **Phase 0: 项目设置与配置 (1-2 天)** +MiniMax-M1 是 MiniMax 发布的 4560 亿参数混合注意力大语言模型，采用 80 层混合架构：标准 GQA（全注意力层）和线性注意力/Mamba 层交替排列，结合 MoE（Mixture of Experts）+ SharedExpert 和 DeepNorm 预归一化策略。该模型在长上下文推理、多轮对话和复杂推理任务上有优异表现。 -在写任何代码之前，先搭建好项目的骨架。 +FastDeploy 已通过 DeepSeek-v3 的 MoE 实现和 GLM-4.5 的 GQA 注意力积累了丰富的可复用组件，但**线性注意力/Mamba 后端**和**混合注意力调度**是全新模块，需要核心开发。 -**任务 1: 创建目录结构** -在 FastDeploy 代码库中，创建以下新文件和目录的占位符： -```bash -# 1. 创建 Mamba 算子的 C++/CUDA 目录 -mkdir -p custom_ops/gpu_ops/mamba_attn +## 2、功能目标 -# 2. 创建 Mamba 算子的 Python 包装器目录和文件 -touch fastdeploy/model_executor/layers/attention/ops/mamba_attention.py +1. 实现 MiniMax-M1 模型组网代码，提交至 `FastDeploy/fastdeploy/model_executor/models/` 目录 +2. 实现 Lightning Attention CUDA 内核（从 vLLM Triton → CUDA C++ 翻译），提交至 `FastDeploy/custom_ops/gpu_ops/` 目录 +3. 实现 `MambaBackend(AttentionBackend)` Python 后端，集成至注意力后端框架 +4. 适配 FastDeploy 现有的低bit量化推理能力 +5. 提交模型使用说明文档 -# 3. 创建 Mamba 算子的 Python 后端文件 -touch fastdeploy/model_executor/layers/attention/mamba_backend.py +## 3、意义 + +- 填补 FastDeploy 在混合注意力（标准 + 线性/Mamba）架构的空白 +- `MambaBackend` 后端可复用于后续 Mamba 系列模型（如 Jamba、Zamba 等） +- 验证 FastDeploy 注意力后端框架对非 Transformer-only 架构的扩展性 + +# 二、飞桨现状 + +## 1、已有基础设施 + +| 组件 | 文件 | 与 MiniMax-M1 的关系 | +|------|------|---------------------| +| GQA 注意力 | `layers/attention/flash_attn_backend.py` | 直接复用（标准注意力层） | +| MoE 分派 | `layers/moe/fused_moe_*.py` | 直接复用（Top-2 MoE 路由） | +| SharedExpert | `layers/moe/fused_moe_*.py` | 需确认是否支持 shared 路径 | +| Partial RoPE | `layers/rotary_embedding.py` | 直接复用 | +| RMSNorm / DeepNorm | `layers/normalization.py` | RMSNorm 可复用，DeepNorm alpha/beta 缩放需新增 | +| 并行线性层 | `layers/linear.py` | 直接复用 | +| 模型注册 | `models/__init__.py` | 直接复用 | + +## 2、缺失组件 + +| 组件 | 说明 | 复杂度 | +|------|------|--------| +| Lightning Attention CUDA 内核 | 5 个 Triton kernel → CUDA C++ 翻译 | **高** — 核心开发 | +| MambaBackend | 新增注意力后端 | **中** — 遵循 AttentionBackend 接口 | +| 混合注意力调度 | 逐层选择 GQA 或 Mamba | **低** — 配置驱动 | +| Mamba SSM 状态管理 | Block Manager 中管理 Mamba 状态缓存 | **中** — 需扩展 | + +# 三、业内方案调研 + +## 1、vLLM + +vLLM 的 MiniMax-M1 支持（`vllm/model_executor/models/minimax_text_01.py`）包含： +- `MiniMaxText01LinearAttention`：使用 Triton 实现的 Lightning Attention（5 个 `@triton.jit` kernel） +- `MiniMaxText01Attention`：标准 GQA 注意力 +- 混合层选择通过 `config.attn_type_list` 逐层配置 +- MoE + SharedExpert 集成 +- DeepNorm 通过 alpha/beta 参数缩放 residual + +**关键代码引用**：`vllm/model_executor/layers/lightning_attn.py` 中定义了 5 个核心 Triton kernel： +1. `_fwd_diag_kernel` — 对角块注意力计算 +2. `_fwd_kv_parallel` — KV 并行前缀计算 +3. `_fwd_kv_reduce` — KV 归约 +4. `_fwd_none_diag_kernel` — 非对角块注意力 +5. `_linear_attn_decode_kernel` — 解码阶段线性注意力 + +## 2、对比分析 + +| 方案 | 优点 | 缺点 | +|------|------|------| +| vLLM (Triton) | 成熟实现，社区验证 | Triton kernel 不可直接用于 Paddle | +| **FastDeploy（本方案）** | 原生 CUDA C++，与 Paddle 深度集成 | CUDA 翻译工作量大，需逐核验证 | + +# 四、设计思路与实现方案 + +## 1、主体设计 + +**核心策略**：将 vLLM 的 5 个 Triton kernel 翻译为 CUDA C++ 算子，封装为 `MambaBackend(AttentionBackend)`，然后在 MiniMax-M1 模型文件中通过 `attn_type_list` 配置逐层选择 GQA 或 Mamba 后端。 + +### 架构总览 -# 4. 创建 MiniMax-M1 的模型定义文件 -touch fastdeploy/model_executor/models/minimax_m1.py ``` +minimax_m1.py (模型定义) +├── MiniMaxM1ForCausalLM(ModelForCasualLM) +│ └── MiniMaxM1Model +│ └── MiniMaxM1DecoderLayer × 80 +│ ├── [attn_type=1] StandardAttention → FlashAttnBackend (已有) +│ ├── [attn_type=0] LinearAttention → MambaBackend (新增) +│ ├── MoE + SharedExpert (已有) +│ └── DeepNorm residual scaling (新增逻辑) +│ +mamba_backend.py (新增后端) +├── MambaBackend(AttentionBackend) +│ ├── init_attention_metadata() +│ ├── forward_extend() — Prefill 阶段 +│ └── forward_decode() — Decode 阶段 +│ +custom_ops/gpu_ops/mamba_attn/ (新增 CUDA 算子) +├── mamba_impl.cuh — 5 个翻译后的 CUDA kernel +├── mamba.cu — Host 启动器 + 算子注册 +└── 注册至 cpp_extensions.cc +``` + +## 2、关键技术点 + +### 2.1 Lightning Attention CUDA 翻译 + +5 个 Triton kernel 的 CUDA C++ 翻译规则： + +| Triton 概念 | CUDA 对应 | +|-------------|-----------| +| `tl.program_id(0/1)` | `blockIdx.x / blockIdx.y` | +| `tl.load(ptr + offsets, mask)` | `if (idx < N) data[idx]` | +| `tl.dot(a, b)` | 寄存器级矩阵乘或 `wmma` 指令 | +| `tl.sum(x, axis)` | Warp/Block 归约 | +| `tl.exp(x)` | `expf(x)` | + +**验证策略**：每个 kernel 翻译后独立对比 Triton 版本输出，要求 `max_abs_diff < 1e-3`（FP16）。 + +### 2.2 Mamba SSM 状态管理 + +与标准 KV Cache 不同，Mamba 的线性注意力维护一个**累积 SSM 状态**（形状 `[batch, heads, head_dim, head_dim]`），在时序步之间传递。 + +管理方案： +- 在 `ForwardMeta.caches` 中为 Mamba 层分配独立的状态缓存槽位 +- Prefill 阶段：从零初始化 SSM 状态，算子内部逐块累积 +- Decode 阶段：从缓存读取上一步状态，更新后写回 + +### 2.3 混合注意力调度 + +```python +# minimax_m1.py +class MiniMaxM1DecoderLayer: + def __init__(self, fd_config, layer_id): + attn_type = fd_config.model_config.attn_type_list[layer_id] + if attn_type == 0: # 线性注意力 + self.self_attn = MiniMaxM1LinearAttention(fd_config, layer_id) + else: # 标准全注意力 + self.self_attn = MiniMaxM1StandardAttention(fd_config, layer_id) +``` + +### 2.4 DeepNorm + +MiniMax-M1 使用 DeepNorm 变体，对注意力和 MLP 的输出分别应用不同的 alpha/beta 缩放： -**任务 2: 更新 `FDConfig`** -编辑 `fastdeploy/config.py`，在 `ModelConfig` 类中添加 MiniMax-M1 特有的配置项。 ```python -# fastdeploy/config.py -> class ModelConfig - -class ModelConfig: - def __init__(self, model: str, **args): - # ... - self.partial_rotary_factor: float = 1.0 - - # === 在这里添加新配置 === - self.attn_type_list: list = [] - self.layernorm_full_attention_alpha: float = 1.0 - self.layernorm_full_attention_beta: float = 1.0 - self.layernorm_linear_attention_alpha: float = 1.0 - self.layernorm_linear_attention_beta: float = 1.0 - self.layernorm_mlp_alpha: float = 1.0 - self.layernorm_mlp_beta: float = 1.0 - self.postnorm: bool = False - # ======================= - - for key, value in args.items(): - # ... +# 注意力 residual +hidden_states = residual * alpha_attn + attn_output * beta_attn +# MLP residual +hidden_states = residual * alpha_mlp + mlp_output * beta_mlp ``` ---- - -### **Phase 1: [核心开发] 实现 Mamba/线性注意力 CUDA 算子 (2-4 周)** - -这是整个项目中技术含量最高、最耗时的部分。 - -**任务 1.1: 翻译 Triton Kernels 为 CUDA C++ (`mamba_impl.cuh`)** -* **目标**: 将 `vllm/model_executor/layers/lightning_attn.py` 中的所有 `@triton.jit` 函数翻译成 `__global__` CUDA kernel。 -* **创建文件**: `custom_ops/gpu_ops/mamba_attn/mamba_impl.cuh` -* **翻译指南**: - 1. 将 `_fwd_diag_kernel` 翻译为 `MambaFwdDiagKernel`。 - 2. 将 `_fwd_kv_parallel` 翻译为 `MambaFwdKVParallelKernel`。 - 3. 将 `_fwd_kv_reduce` 翻译为 `MambaFwdKVReduceKernel`。 - 4. 将 `_fwd_none_diag_kernel` 翻译为 `MambaFwdNonDiagKernel`。 - 5. 将 `_linear_attn_decode_kernel` 翻译为 `MambaDecodeKernel`。 - * **关键替换**: `tl.program_id` -> `blockIdx`, `tl.load/store` -> `if-guarded memory access`, `tl.dot` -> 手写寄存器级矩阵乘法，`tl.sum` -> Warp/Block 级归约。 - -**任务 1.2: 编写 C++ Host 启动器 (`mamba.cu`)** -* **目标**: 编写一个 C++ Host 函数，模仿 `lightning_attn.py` 中的 `_attention.forward` 和 `linear_decode_forward_triton` 的调度逻辑。 -* **创建文件**: `custom_ops/gpu_ops/mamba_attn/mamba.cu` -* **代码框架**: - ```cpp - #include "mamba_impl.cuh" - #include "paddle/extension.h" - - void MambaAttentionForwardKernel( - const paddle::Tensor& q, const paddle::Tensor& k, const paddle::Tensor& v, - const paddle::Tensor& slope_rate, paddle::Tensor& mamba_state, - paddle::Tensor& out, /* ... 其他 ForwardMeta 参数 ... */ - ) { - // 1. 从 ForwardMeta 解析出 is_prefill, b, h, n, d 等信息 - bool is_prefill = ...; // 根据 seq_lens 判断 - - // 2. 模仿 _attention.forward 和 linear_decode_forward_triton - if (is_prefill) { - // 计算 grid/block 维度 - // 依次启动 MambaFwd... 系列的 CUDA Kernel - } else { // Decode - // 计算 grid/block 维度 - // 启动 MambaDecodeKernel - } - } - ``` - -**任务 1.3: 暴露自定义算子 (`mamba.cu` & `cpp_extensions.cc`)** -1. 在 `mamba.cu` 文件末尾添加算子注册代码： - ```cpp - PD_BUILD_STATIC_OP(mamba_attention_forward) - .Inputs({ ... }) - .Outputs({ "Out", "MambaStateOut" }) - .SetKernelFn(PD_KERNEL(MambaAttentionForwardKernel)); - ``` -2. 编辑 `custom_ops/gpu_ops/cpp_extensions.cc`，添加 `m.def("mamba_attention_forward", &MambaAttentionForward);`。 -3. 更新 `custom_ops/setup_ops.py` 或 `CMakeLists.txt`，将 `mamba.cu` 加入编译列表。 - -**里程碑**: 完成并编译通过后，拥有了一个底层的、可被调用的 Mamba/线性注意力算子。 - ---- - -### **Phase 2: 构建 Python 接口与后端 (1 周)** - -**任务 2.1: 创建底层 Python 包装器** -* **文件**: `fastdeploy/model_executor/layers/attention/ops/mamba_attention.py` -* **代码**: - ```python - from paddle.fluid import core - - def mamba_attention_forward(q, k, v, slope_rate, mamba_state, ...): - # 实际调用 C++ 扩展 - out, mamba_state_out = core.eager._run_custom_op( - "mamba_attention_forward", q, k, v, slope_rate, mamba_state, ... - ) - return out, mamba_state_out - ``` - -**任务 2.2: 创建高级 MambaBackend** -* **文件**: `fastdeploy/model_executor/layers/attention/mamba_backend.py` -* **代码框架**: - ```python - from .ops.mamba_attention import mamba_attention_forward - from fastdeploy.model_executor.layers.attention.base_attention_backend import AttentionBackend - - class MambaBackend(AttentionBackend): - def __init__(self, fd_config, ...): - # 初始化 Mamba 需要的参数 - pass - - def init_attention_metadata(self, forward_meta: ForwardMeta): - # Mamba 可能需要一些独特的元数据准备 - pass - - def forward(self, q, k, v, qkv, ..., layer: Attention, forward_meta: ForwardMeta): - # 1. 这里是核心业务逻辑，模仿 vLLM 的 MiniMaxText01LinearAttention._forward - # 它会从 forward_meta 中解析调度信息 - - # 2. 调用底层算子 - out, ssm_states_out = mamba_attention_forward(...) - - # 3. 更新 mamba_state (可能通过 forward_meta.caches) - forward_meta.caches[layer.layer_id].copy_(ssm_states_out, False) - - return out - ``` - - ---- - -### **Phase 3: 全量模型支持与性能验证 (1 周)** - -**目标**: 将已在 5 层模型上验证通过的核心组件无缝扩展至 80 层全量模型。 - -修改 `minimax_m1.py` 中的 `load_weights` 函数，使其能够加载并正确映射**全部 80 层**的权重。在启动脚本中，将 `num_hidden_layers` 设置为 80，并且对齐vllm精度。 - ---- - -### **Phase 4: 模型集成与测试 (1 周)** - -**任务 4.1: 编写 `minimax_m1.py`** -* **文件**: `fastdeploy/model_executor/models/minimax_m1.py` -* **核心逻辑**: - ```python - from fastdeploy.model_executor.layers.attention.mamba_backend import MambaBackend - from fastdeploy.model_executor.layers.attention.mla_attention_backend import MLAAttentionBackend - - class MiniMaxM1DecoderLayer(nn.Layer): - def __init__(self, fd_config, layer_id): - attn_type = fd_config.model_config.attn_type_list[layer_id] - if attn_type == 0: # 线性注意力 - self.self_attn = MiniMaxM1LinearAttention(fd_config, layer_id) # 这是一个新的 nn.Layer - else: # 标准注意力 - self.self_attn = MiniMaxM1StandardAttention(fd_config, layer_id) # 复用/包装标准 Attention - - # DeepNorm 参数 - self.layernorm_attention_alpha = ... - - def forward(self, hidden_states, ...): - # 实现 DeepNorm 和 MoE+SharedExpert 逻辑 - attn_output = self.self_attn(...) - layernorm_input = residual * self.layernorm_attention_alpha + attn_output * self.layernorm_attention_beta - ... - - # 建议为两种 Attention 创建独立的 nn.Layer 包装类 - class MiniMaxM1LinearAttention(nn.Layer): - def __init__(...): - self.backend = MambaBackend(...) - self.qkv_proj = ... - self.out_proj = ... - self.norm = ... - # ... - def forward(...): - qkv = self.qkv_proj(...) - # ... 准备输入 ... - attn_out = self.backend.forward(...) - # ... 后处理 ... - return final_out - - # MiniMaxM1StandardAttention 类似 - ``` - -**任务 4.2: 编写测试用例** -* **单元测试**: 编写一个 Python 测试脚本，直接调用 `ops.mamba_attention_forward`，并与一个纯 Python/Paddle 实现的、功能等价的 Mamba 算法进行对比，验证数值精度。 -* **端到端测试**: 创建一个完整的 MiniMax-M1 模型实例，加载权重，并运行推理。将输出与 Hugging Face `transformers` 的参考输出进行对比。 - -**任务 4.3: 性能基准测试** -* 在与 vLLM 相同的硬件和环境下，测试 FastDeploy 实现的吞吐量和时延，确保性能达标。 \ No newline at end of file +alpha/beta 值从 HuggingFace 配置的 `layernorm_full_attention_alpha`、`layernorm_linear_attention_alpha` 等字段读取。 + +### 2.5 文件结构 + +| 路径 | 说明 | +|------|------| +| `custom_ops/gpu_ops/mamba_attn/mamba_impl.cuh` | CUDA kernel 实现 | +| `custom_ops/gpu_ops/mamba_attn/mamba.cu` | Host 启动器 + 算子注册 | +| `custom_ops/gpu_ops/cpp_extensions.cc` | 新增 mamba 算子绑定 | +| `fastdeploy/model_executor/layers/attention/mamba_backend.py` | MambaBackend 后端 | +| `fastdeploy/model_executor/layers/attention/ops/mamba_attention.py` | Python 算子包装 | +| `fastdeploy/model_executor/models/minimax_m1.py` | 模型组网代码 | + +# 五、测试和验收的考量 + +1. **CUDA 内核精度**：每个翻译后的 kernel 与 vLLM Triton 版本对比，FP16 最大绝对差 < 1e-3 +2. **端到端推理**：对比 HuggingFace transformers 输出，token-level 一致性 +3. **量化精度**：WINT8 精度损失 <1%，WINT4 精度损失 <3% +4. **长上下文**：验证 4K/8K/16K 上下文长度下的推理正确性 +5. **性能基线**：多卡 A100 首 token 延迟和吞吐量，对标 vLLM + +# 六、影响面 + +## 对用户的影响 +- 新增模型选项，无 breaking changes + +## 对框架架构的影响 +- 新增 `MambaBackend` 注意力后端（扩展点，不修改已有后端） +- 新增 CUDA 算子 `mamba_attention_forward` +- `config.py` 新增 Mamba/DeepNorm 相关配置项 + +## 对性能的影响 +- 不影响现有模型性能 +- Mamba 线性注意力的 O(n) 复杂度可显著降低长文本推理延迟 + +# 七、排期规划 + +| 阶段 | 任务 | 预计时间 | +|------|------|---------| +| Phase 0 | 项目搭建 + 配置解析 + 目录结构 | 1-2天 | +| Phase 1 | Lightning Attention CUDA 内核翻译（5 个 kernel） | 2-3周 | +| Phase 2 | MambaBackend Python 后端 + 算子包装 | 1周 | +| Phase 3 | 模型组网 + 权重加载 + 混合调度 | 1周 | +| Phase 4 | 端到端测试 + 性能调优 + 文档 | 1周 | +| **合计** | | **~5-6周** | + +# 名词解释 + +| 名词 | 说明 | +|------|------| +| Lightning Attention | MiniMax 提出的线性注意力机制，O(n) 时间复杂度 | +| Mamba/SSM | 选择性状态空间模型，一种替代 Transformer 的序列建模方法 | +| DeepNorm | 深层残差连接的归一化策略，用不同 alpha/beta 缩放防止梯度消失 | +| MoE | 混合专家模型，每个 token 路由到 Top-K 专家处理 | + +# 附件及参考资料 + +1. MiniMax-M1 HuggingFace: https://huggingface.co/MiniMaxAI/MiniMax-M1-80B +2. vLLM MiniMax-M1 实现: `vllm/model_executor/models/minimax_text_01.py` +3. vLLM Lightning Attention: `vllm/model_executor/layers/lightning_attn.py` +4. FastDeploy 注意力后端基类: `fastdeploy/model_executor/layers/attention/base_attention_backend.py` +5. FastDeploy DeepSeek-v3 MoE 参考: `fastdeploy/model_executor/models/deepseek_v3.py` +6. 原始 H9 RFC: https://github.com/PaddlePaddle/community/pull/1156