From 972d1059a1a7da00ce64d14843dfef922b051c7a Mon Sep 17 00:00:00 2001 From: WanRui37 <1446145652@qq.com> Date: Thu, 16 Oct 2025 16:11:19 +0800 Subject: [PATCH 1/5] v1: Simply fill in the RFC --- ..._FastDeploy_add_moe_groupgemm_int8_int8.md | 48 +++++++++++++++++++ 1 file changed, 48 insertions(+) create mode 100644 rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md diff --git a/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md b/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md new file mode 100644 index 000000000..5a7b805cc --- /dev/null +++ b/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md @@ -0,0 +1,48 @@ +# FastDeploy中的MoE GroupGEMM支持INT8*INT8实现 + +| 方案名称 | FastDeploy中的MoE GroupGEMM支持INT8*INT8实现 | +|----------------------------------------------------------|-------------------------------------------| +| 提交作者 | WanRui37 | +| 提交时间 | 2025-10-16 | +| 版本号 | V1.1 | +| 依赖飞桨版本 | paddlepaddle-gpu==3.2.0 | +| 文件名 | 20251016_FastDeploy_add_moe_groupgemm_int8_int8.md
| + +# 一、概述 +## 1、相关背景 +大规模模型在自然语言处理、计算机视觉等领域取得了显著成果。其中，混合专家模型（Mixture of Experts，MoE）作为一种高效的模型架构，通过将输入数据分配给不同的专家子网络进行处理，能够有效提升模型的性能和计算效率。在MoE模型的训练和推理过程中，GroupGEMM（Group General Matrix Multiply）操作是核心计算步骤 + +## 2、功能目标 +为FastDeploy 开发高性能 MoE算子(INT8*INT8)，将上述算子集成到EB、Qwen等开源模型中。 + +## 3、意义 +支持INT8*INT8的MoE GroupGEMM实现能够充分利用硬件的整数计算单元，相较于高精度计算，大幅减少计算延迟，提高模型推理速度。 + +# 二、FastDeploy现状 +- 目前FastDeploy中MoE GroupGEMM没有支持INT8*INT8的实现 + +# 三、业内方案调研 +- 目前业内MoE GroupGEMM没有支持INT8*INT8的实现 + +# 四、设计思路与实现方案 +## 总体设计 + + + +# 五、测试和验收的考量 +- 增加算子测试 +- 在EB，Qwen开源模型上测试数据精度&性能 + +# 六、影响面 +为FastDeploy集成 SageAttn v2++，不影响其他部分 + +# 七、排期规划 +* 2025-10-16 ~ 2025-11-16：完成集成代码开发 +* 2025-11-16 ~ 2025-11-25：完成代码测试 +* 2025-11-25 ~ 2025-12-01： 完成部署示例及文档 + +# 八、参考资料 + +[Accelerating MoE's with a Triton Persistent Cache-Aware Grouped GEMM Kernel](https://pytorch.org/blog/accelerating-moes-with-a-triton-persistent-cache-aware-grouped-gemm-kernel/) + +[上述为 vllm 增加 BF16 Grouped GEMM Kernel 的 PR](https://github.com/vllm-project/vllm/pull/19443) \ No newline at end of file From 19a1c756eecd9b5e63a3615d3823e411a9167f18 Mon Sep 17 00:00:00 2001 From: WanRui37 <1446145652@qq.com> Date: Thu, 16 Oct 2025 21:41:38 +0800 Subject: [PATCH 2/5] v1: Simply fill in the RFC --- .../20251016_FastDeploy_add_moe_groupgemm_int8_int8.md | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md b/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md index 5a7b805cc..2ec2fe701 100644 --- a/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md +++ b/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md @@ -25,9 +25,7 @@ - 目前业内MoE GroupGEMM没有支持INT8*INT8的实现 # 四、设计思路与实现方案 -## 总体设计 - - +参考 FastDeploy 其余算子实现，添加好INT8*INT8的MoE GroupGEMM实现。 # 五、测试和验收的考量 - 增加算子测试 From 5a34a4efc715baffe135aae592e66880173cf0c6 Mon Sep 17 00:00:00 2001 From: WanRui37 <1446145652@qq.com> Date: Thu, 16 Oct 2025 21:47:35 +0800 Subject: [PATCH 3/5] v1: Simply fill in the RFC --- .../20251016_FastDeploy_add_moe_groupgemm_int8_int8.md | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md b/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md index 2ec2fe701..7214d97e4 100644 --- a/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md +++ b/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md @@ -19,20 +19,20 @@ 支持INT8*INT8的MoE GroupGEMM实现能够充分利用硬件的整数计算单元，相较于高精度计算，大幅减少计算延迟，提高模型推理速度。 # 二、FastDeploy现状 -- 目前FastDeploy中MoE GroupGEMM没有支持INT8*INT8的实现 +- 目前`FastDeploy`中`MoE GroupGEMM`没有支持`INT8*INT8`的实现 # 三、业内方案调研 -- 目前业内MoE GroupGEMM没有支持INT8*INT8的实现 +- 目前业内`MoE GroupGEMM`没有支持`INT8*INT8`的实现 # 四、设计思路与实现方案 -参考 FastDeploy 其余算子实现，添加好INT8*INT8的MoE GroupGEMM实现。 +参考`FastDeploy`其余算子实现，添加好`INT8*INT8`的`MoE GroupGEMM`实现。 # 五、测试和验收的考量 - 增加算子测试 - 在EB，Qwen开源模型上测试数据精度&性能 # 六、影响面 -为FastDeploy集成 SageAttn v2++，不影响其他部分 +为`FastDeploy`集成`MoE GroupGEMM`，不影响其他部分 # 七、排期规划 * 2025-10-16 ~ 2025-11-16：完成集成代码开发 From 09b6eb39cb74ecc5d0870a91bcecfa65272807af Mon Sep 17 00:00:00 2001 From: WanRui37 <1446145652@qq.com> Date: Wed, 22 Oct 2025 21:48:28 +0800 Subject: [PATCH 4/5] v2: Added some design ideas --- ..._FastDeploy_add_moe_groupgemm_int8_int8.md | 25 ++++++++++++++++++- 1 file changed, 24 insertions(+), 1 deletion(-) diff --git a/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md b/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md index 7214d97e4..767b405e2 100644 --- a/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md +++ b/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md @@ -25,7 +25,30 @@ - 目前业内`MoE GroupGEMM`没有支持`INT8*INT8`的实现 # 四、设计思路与实现方案 -参考`FastDeploy`其余算子实现，添加好`INT8*INT8`的`MoE GroupGEMM`实现。 +1. 一些参考的代码路径 + 1. `sm89`架构可以参考 + - `24_gemm_grouped`: `FastDeploy/custom_ops/third_party/cutlass/examples/24_gemm_grouped` + - `64_ada_fp8_gemm_grouped`: `FastDeploy/custom_ops/third_party/cutlass/examples/64_ada_fp8_gemm_grouped` + - 比较重要的就是`cutlass::gemm::device::GemmGrouped` + 1. `sm90`架构可以参考 + - `w4afp8_gemm`: `FastDeploy/custom_ops/gpu_ops/w4afp8_gemm` + - 比较重要的是`cute::gemm` + - `57_hopper_grouped_gemm`: `FastDeploy/custom_ops/third_party/cutlass/examples/57_hopper_grouped_gemm` + - 比较重要的是`cutlass::gemm::device::GemmUniversalAdapter` + +1. 需要修改的代码路径 + ```text + custom_ops/ + └── gpu_ops/ # GPU相关自定义算子 + ├── int8_gemm_with_cutlass/ # INT8*INT8 GEMM with Cutlass 算子实现 + │ ├── w8a8_group_gemm.cu # Cutlass Kernel实现 + │ └── w8a8_group_gemm.h # Cutlass Kernel头文件 + └── ... + test/ + └── operators/ + ├── test_w8a8_group_gemm.py # 测试INT8*INT8 GEMM with Cutlass 算子 + └── ... + ``` # 五、测试和验收的考量 - 增加算子测试 From 754593c96a3a8d836d034f7d451e4c2cc6dc0e9c Mon Sep 17 00:00:00 2001 From: WanRui37 <1446145652@qq.com> Date: Wed, 22 Oct 2025 21:54:49 +0800 Subject: [PATCH 5/5] v2: Added some design ideas --- ...6_FastDeploy_add_moe_groupgemm_int8_int8.md | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md b/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md index 767b405e2..2a53ff9b5 100644 --- a/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md +++ b/rfcs/FastDeploy/20251016_FastDeploy_add_moe_groupgemm_int8_int8.md @@ -1,12 +1,12 @@ # FastDeploy中的MoE GroupGEMM支持INT8*INT8实现 -| 方案名称 | FastDeploy中的MoE GroupGEMM支持INT8*INT8实现 | +| 方案名称 | FastDeploy中的MoE GroupGEMM支持INT8*INT8实现 | |----------------------------------------------------------|-------------------------------------------| -| 提交作者 | WanRui37 | -| 提交时间 | 2025-10-16 | -| 版本号 | V1.1 | -| 依赖飞桨版本 | paddlepaddle-gpu==3.2.0 | -| 文件名 | 20251016_FastDeploy_add_moe_groupgemm_int8_int8.md
| +| 提交作者 | WanRui37 | +| 提交时间 | 2025-10-16 | +| 版本号 | V1.1 | +| 依赖飞桨版本 | paddlepaddle-gpu==3.2.0 | +| 文件名 | 20251016_FastDeploy_add_moe_groupgemm_int8_int8.md
| # 一、概述 ## 1、相关背景 @@ -38,12 +38,12 @@ 1. 需要修改的代码路径 ```text - custom_ops/ + custom_ops/ └── gpu_ops/ # GPU相关自定义算子 ├── int8_gemm_with_cutlass/ # INT8*INT8 GEMM with Cutlass 算子实现 │ ├── w8a8_group_gemm.cu # Cutlass Kernel实现 │ └── w8a8_group_gemm.h # Cutlass Kernel头文件 - └── ... + └── ... test/ └── operators/ ├── test_w8a8_group_gemm.py # 测试INT8*INT8 GEMM with Cutlass 算子 @@ -66,4 +66,4 @@ [Accelerating MoE's with a Triton Persistent Cache-Aware Grouped GEMM Kernel](https://pytorch.org/blog/accelerating-moes-with-a-triton-persistent-cache-aware-grouped-gemm-kernel/) -[上述为 vllm 增加 BF16 Grouped GEMM Kernel 的 PR](https://github.com/vllm-project/vllm/pull/19443) \ No newline at end of file +[上述为 vllm 增加 BF16 Grouped GEMM Kernel 的 PR](https://github.com/vllm-project/vllm/pull/19443)