企业官网建设流程全解析

SGLang MoE Triton Kernel 调优实战指南：TP/EP 并行模式、FP8 量化与配置落地

【免费下载链接】sglangSGLang is a high-performance serving framework for large language models and multimodal models.项目地址: https://gitcode.com/GitHub_Trending/sg/sglang

本篇技术指南以 SGLang 仓库中的 MoE（Mixture of Experts，混合专家）内核基准测试与调优工具集为核心，系统讲解fused_moe_triton内核在张量并行（TP）、专家并行（EP）以及多模态大模型（MLLM）场景下的自动调优流程、分离内核（Separate Kernel）调优方法、量化数据类型支持、以及调优结果配置文件的生成与落地部署。读完本文，你将掌握如何针对自己的 GPU 与模型组合，用仓库内置的调优脚本搜索最优 Triton 内核配置，并将其应用到 SGLang 服务中，同时学会使用对比基准脚本评估 SGLang 与 vLLM 的 MoE 内核性能差异。

一、工具集概览：一个针对 MoE 内核的完整调优闭环

调优脚本位于 benchmark/kernels/fused_moe_triton/ 目录，包含五个核心脚本：

该目录是 SGLang 仓库中"调优-产出配置-回填配置"闭环的入口：调优脚本搜索最佳内核配置，产出 JSON 配置文件，这些文件最终被 fused_moe_triton_config.py 在服务启动时按模型与设备名自动加载。

支持的并行模式

调优工具同时支持两种主流分布式并行模式：

• TP 模式（张量并行）：MoE 中间层权重按维度切分到多张 GPU 上，每张卡持有部分专家权重，所有专家在所有 GPU 上都有分片。
• EP 模式（专家并行）：专家本身被分布到不同 GPU 上，每张卡只持有部分完整专家。EP 可与 TP 组合使用（例如--tp-size 8 --ep-size 2），组合时要求tp_size 能被 ep_size 整除——这一点在 common_utils.py 的calculate_shard_intermediate_size中通过assert tp_size % ep_size == 0强制校验。

EP 模式对内核形状的影响体现在两个维度：专家数量E按E = num_experts // ep_size缩减（每卡只负责本地专家）；中间层大小则按shard_intermediate_size = 2 * intermediate_size / (tp_size / ep_size)计算，其中moe_tp_size = tp_size // ep_size表示 MoE 层实际参与切分的 GPU 数量。

MLLM 支持

调优工具支持带文本编码器的多模态大模型（如 Llama4-Vision、Qwen3-VL）。在 common_utils.py 的get_model_config中，当检测到编码器-解码器结构（hasattr(config, "text_config")）时，会自动替换为text_config再读取 MoE 参数，从而正确提取专家数、topk、中间层大小等调优所需形状信息。

二、统一内核调优：tuning_fused_moe_triton.py

tuning_fused_moe_triton.py是主调优工具，其实现思路改编自 vLLM 的benchmark_moe.py，但扩展了 EP 模式与更多模型架构支持。

2.1 基础 TP 模式调优

# 默认 TP 设置下调优 Mixtral-8x7B python benchmark/kernels/fused_moe_triton/tuning_fused_moe_triton.py \ --model mistralai/Mixtral-8x7B-Instruct-v0.1 \ --tune # 以 FP8 与 TP=4 调优 Qwen2-57B python benchmark/kernels/fused_moe_triton/tuning_fused_moe_triton.py \ --model Qwen/Qwen2-57B-A14B-Instruct \ --tp-size 4 \ --dtype fp8_w8a8 \ --tune # 以 FP8 与 TP=8 调优 DeepSeek-V3 python benchmark/kernels/fused_moe_triton/tuning_fused_moe_triton.py \ --model deepseek-ai/DeepSeek-V3-0324 \ --tp-size 8 \ --dtype fp8_w8a8 \ --tune

2.2 EP 模式调优

EP 模式可单独使用，也可与 TP 组合。需要注意：组合使用时tp_size必须能被ep_size整除。

# 仅 EP=2（此时 tp-size 2 等价于 2 卡做专家切分） python benchmark/kernels/fused_moe_triton/tuning_fused_moe_triton.py \ --model mistralai/Mixtral-8x7B-Instruct-v0.1 \ --tp-size 2 \ --ep-size 2 \ --tune # TP=8 与 EP=4 组合模式 python benchmark/kernels/fused_moe_triton/tuning_fused_moe_triton.py \ --model Qwen/Qwen2-57B-A14B-Instruct \ --tp-size 8 \ --ep-size 4 \ --dtype fp8_w8a8 \ --tune

2.3 MLLM 多模态模型调优

python benchmark/kernels/fused_moe_triton/tuning_fused_moe_triton.py \ --model Qwen/Qwen3-VL-30B-A3B-Instruct \ --tp-size 2 \ --tune

2.4 高级选项

# 通道级量化（per-channel quantization），适用于 meituan/DeepSeek-R1-Channel-INT8 python benchmark/kernels/fused_moe_triton/tuning_fused_moe_triton.py \ --model meituan/DeepSeek-R1-Channel-INT8 \ --tp-size 16 \ --dtype int8_w8a8 \ --per-channel-quant \ --tune # 针对特定 batch size 调优 python benchmark/kernels/fused_moe_triton/tuning_fused_moe_triton.py \ --model mistralai/Mixtral-8x7B-Instruct-v0.1 \ --batch-size 2048 \ --tune

2.5 底层实现细节

从源码看，tuning_fused_moe_triton.py的调优流程有几个值得注意的设计（见 tuning_fused_moe_triton.py）：

• Ray 分布式分发：main中通过ray.init()与BenchmarkWorker.remote(...)为每张可见 GPU 创建 worker，多个 batch size 的调优任务通过轮询（round-robin）方式分发到所有 GPU，实现并行调优。
• CUDA Graph 捕获：benchmark_config将 10 次内核调用捕获进 CUDA Graph 并 replay 计时，最终平均延迟avg = sum(latencies) / (num_iters * 10) * 1000以微秒为单位输出。
• L2 Cache 冲刷：每次迭代前用 256 MB 数据冲刷 L2 缓存，确保测得的延迟反映真实计算时间而非缓存命中带来的偏差。
• 无效配置容错：调优循环中捕获triton.runtime.autotuner.OutOfResources与RuntimeError——部分配置可能因共享内存超限等原因是非法组合，直接跳过继续搜索。
• 默认 batch size 网格：默认调优一批 token 数[1, 2, 4, 8, 16, 24, 32, 48, 64, 96, 128, 256, 512, 1024, 1536, 2048, 3072, 4096]，每个 batch size 产出最优配置，最终 JSON 中每个 batch size 对应一组BLOCK_SIZE_M/N/K、GROUP_SIZE_M、num_warps、num_stages。
• 搜索空间：CUDA 环境下get_configs_compute_bound()枚举num_stages ∈ {2,3,4,5}、BLOCK_SIZE_M ∈ {16,32,64,128,256}、BLOCK_SIZE_K ∈ {64,128,256}、BLOCK_SIZE_N ∈ {32,64,128,256}、num_warps ∈ {4,8}、GROUP_SIZE_M ∈ {1,16,32,64}；ROCm 环境则额外包含waves_per_eu维度，且参数范围不同（见 common_utils.py）。

三、分离内核调优：tuning_fused_moe_triton_sep.py

tuning_fused_moe_triton_sep.py是面向精细化优化的专用工具：它把 MoE 计算拆成第一个内核（up 投影，含 silu_and_mul 激活）与第二个内核（down 投影），分别搜索最优配置，并支持 TMA（Tensor Memory Accelerator）加速变体。

3.1 第一步：生成 topk_ids 数据

与统一调优不同，分离内核调优需要预先收集真实推理中的 topk_ids 路由结果，以保证调优时的 token 分布贴近真实负载。操作步骤如下：

1. 在 Python site-packages 中的模型代码（如srt/models/deepseek_v2.py）中，于DeepseekV2MoE::forward_normal内加入保存逻辑：

# import get_tensor_model_parallel_rank # DeepseekV2MoE::forward_normal if hidden_states.shape[0] >= 4096 and get_tensor_model_parallel_rank() == 0: topk_ids_dir = xxxx # 替换为你的保存目录 if not hasattr(self, "save_idx"): self.save_idx = 0 if self.save_idx <= 1: torch.save(topk_output.topk_ids, f"{topk_ids_dir}/topk_ids_layer{self.layer_id}_idx{self.save_idx}.pt") self.save_idx += 1

从源码（load_topk_ids，见 tuning_fused_moe_triton_sep.py）可以看出，文件命名约定为topk_ids_layer{layer}_idx{idx}.pt，其中 layer 编号从 dense 层之后开始（DeepSeek 结构默认num_layers=61、dense_layers=3，即 MoE 层从第 3 层起算），一共会读取 100 份 topk_ids 用于多轮迭代。

2. 启动 SGLang 服务，并用tuning_client.py发送长文本请求触发推理：

python benchmark/kernels/fused_moe_triton/tuning_client.py --port 8000

tuning_client.py会读取同目录下的tuning_text.json长文本作为 prompt，以 OpenAI 协议流式请求服务（base_url=http://{ip}:{port}/v1），并打印 TTFT 与 TPOT 指标。生成的 topk_ids 文件保存在你指定的topk_ids_dir中。

3.2 第二步：分离内核调优

# TP 模式：TP=4 调优分离内核 python benchmark/kernels/fused_moe_triton/tuning_fused_moe_triton_sep.py \ --model Qwen/Qwen2-57B-A14B-Instruct \ --tp-size 4 \ --topk-ids-dir /path/to/topk_ids \ --tune # EP 模式：TP=4 与 EP=2 组合 python benchmark/kernels/fused_moe_triton/tuning_fused_moe_triton_sep.py \ --model mistralai/Mixtral-8x7B-Instruct-v0.1 \ --tp-size 4 \ --ep-size 2 \ --topk-ids-dir /path/to/topk_ids \ --tune # MLLM：DeepSeek-V3 分离内核，TP=8 与 EP=4 python benchmark/kernels/fused_moe_triton/tuning_fused_moe_triton_sep.py \ --model deepseek-ai/DeepSeek-V3-0324 \ --tp-size 8 \ --ep-size 4 \ --dtype fp8_w8a8 \ --topk-ids-dir /path/to/topk_ids \ --tune # 不调优，仅用指定配置基准测试 python benchmark/kernels/fused_moe_triton/tuning_fused_moe_triton_sep.py \ --model deepseek-ai/DeepSeek-V3-0324 \ --tp-size 4 \ --batch-size 1024 \ --dtype fp8_w8a8 \ --configs 128 256 128 16 8 4 \ --topk-ids-dir /path/to/topk_ids

--configs参数依次指定[BLOCK_M, BLOCK_N, BLOCK_K, GROUP_M, warps, stages]六个数值，对应benchmark_config中的BLOCK_SIZE_M、BLOCK_SIZE_N、BLOCK_SIZE_K、GROUP_SIZE_M、num_warps、num_stages。

3.3 TMA 与两轮耦合调优

分离内核工具的核心优势是 TMA 支持。TMA 是 Hopper 及后续架构上的硬件张量内存加速器，可降低数据搬运开销。从 tuning_fused_moe_triton_sep.py 源码可以确认以下几点：

• 默认（tune_round="both"）时，up 内核使用非 TMA 变体、down 内核同时评估 TMA 与非 TMA 变体，最终在配置文件的USE_TMA字段记录 down 投影是否启用 TMA（取两者中更快者）。
• 当启用--enable-tune-up-tma时，工具进入两轮耦合调优：第一轮先只调优 down 内核，确定每个BLOCK_SIZE_M下 down 是否使用 TMA；第二轮再以第一轮的c_sorted决策为前提调优 up 内核的 TMA 变体。这样保证 up 内核的排序假设与运行时行为一致（up 内核配置含USE_TMA字段）。
• EP 模式下，保存的 topk_ids 是全局专家索引，prepare 阶段会通过topk_ids // ep_size将其映射为本地专家索引后再喂给内核。
• 通过NCU_ENABLE=1环境变量可以进入单次迭代模式，便于配合 Nsight Compute（ncu）进行逐配置的 profiling。
• 分离调优产出两份配置：up 投影一份（save_configs_sep，down_moe=False）、down 投影一份（_down.json，down_moe=True），文件命名遵循E={E},N={N},device_name={...},dtype={...},block_shape={...}的格式，down 投影文件名额外带_down后缀（该命名规则定义于 fused_moe_triton_config.py 的get_config_file_name）。
• 另外支持--cmp-configs参数，可指定多份既有配置文件进行横向对比，验证不同配置在各 batch size 下的内核耗时。

四、调优结果与配置落地

4.1 生成的文件

调优完成后，脚本会按模型与设备特征自动生成 JSON 配置文件，文件名形如：

• 标准调优：E=64,N=640,device_name=NVIDIA_GeForce_RTX_4090,dtype=fp8_w8a8.json
• 分离内核调优：分别生成 up/down 两份文件，down 文件带_down后缀，并可能包含 TMA 优化标志（USE_TMA）。

文件名中的字段含义为：E为专家数（EP 下为本地专家数）、N为 MoE 中间层经silu_and_mul后的大小（即shard_intermediate_size / 2，INT4 时再除以 2）、device_name为torch.cuda.get_device_name()去掉空格后的设备名、dtype为量化数据类型、block_shape为块量化形状（如 DeepSeek-V3 常用的[128, 128]）。

配置内容是一个以 batch size 为键的映射，例如 python/sglang/srt/layers/moe/moe_runner/triton_utils/configs/triton_3_5_1/E=128,N=1024,device_name=NVIDIA_H100_80GB_HBM3,dtype=fp8_w8a8.json：

{ "1": { "BLOCK_SIZE_M": 16, "BLOCK_SIZE_N": 64, "BLOCK_SIZE_K": 128, "GROUP_SIZE_M": 1, "num_warps": 4, "num_stages": 5 }, "128": { "BLOCK_SIZE_M": 32, "BLOCK_SIZE_N": 64, "BLOCK_SIZE_K": 128, "GROUP_SIZE_M": 1, "num_warps": 4, "num_stages": 5 }, "2048":{ "BLOCK_SIZE_M": 64, "BLOCK_SIZE_N": 128, "BLOCK_SIZE_K": 128, "GROUP_SIZE_M": 1, "num_warps": 8, "num_stages": 4 }, "4096":{ "BLOCK_SIZE_M": 128, "BLOCK_SIZE_N": 256, "BLOCK_SIZE_K": 128, "GROUP_SIZE_M": 32, "num_warps": 8, "num_stages": 4 } }

可以看到，随着 batch size 增大，最优配置趋向更大的BLOCK_SIZE_M与BLOCK_SIZE_N、更多的 warp 数——这正是内核调优要捕捉的规律：小 batch 下过大的 tile 会浪费计算资源，大 batch 下则需要更大的 tile 来摊薄调度开销。

4.2 部署到 SGLang

将生成的 JSON 文件移动到 SGLang 的配置目录即可让服务自动加载：

python/sglang/srt/layers/moe/moe_runner/triton_utils/configs/triton_version/

其中triton_version为当前环境的 Triton 版本目录（如triton_3_5_1）。配置文件按 Triton 版本隔离是刻意设计：源码注释明确指出，Triton 3.1.0 的配置直接用于 3.2.0 会出现性能回退，因此 fused_moe_triton_config.py 在get_moe_configs中会先按当前 Triton 版本查找对应子目录，未命中时再按版本号从新到旧回退搜索。

加载时还遵循以下逻辑：

• 若环境变量SGLANG_MOE_CONFIG_DIR已设置，则优先从该目录读取配置（默认是triton_utils目录本身）；
• 命中配置文件后，运行时会以 batch size 为键，选取网格中最接近实际 batch size的配置调用内核；
• 若开启了确定性推理（deterministic.enable_deterministic_inference），会回退到默认内核配置；
• 源码中的提示也值得注意：由于调优环境与运行环境可能存在差异（如 Triton 版本升级），旧配置可能并非最优，建议在目标环境重新调优。

仓库 configs 目录 下已经预置了大量覆盖 A100/H100/H200/B200/MI300X/RTX 4090 等设备、多种 dtype 与 block_shape 的配置文件，可以作为对照参考。

五、参数参考

六、支持的模型

调优工具对模型架构的解析集中在 common_utils.py 的get_model_config中，不同架构的字段名（专家数、topk、中间层大小）差异较大，需要逐架构适配。README 列出的已支持模型包括：

• Mixtral：mistralai/Mixtral-8x7B-Instruct-v0.1、mixtral-8x22b
• Qwen：Qwen2-57B、Qwen3-235B、Qwen3VL（MLLM）
• DeepSeek：DeepSeek-V2、DeepSeek-V3、DeepSeek-R1
• Llama：Llama4-Vision（MLLM）
• DBRX：databricks/dbrx-instruct
• Jamba：ai21labs/AI21-Jamba
• Grok：xai-org/grok-1
• GLM：THUDM/glm-4-9b-chat
• Bailing：自定义 MoE 模型

从源码看，架构分发还覆盖了 Qwen2Moe/Qwen3Moe/Qwen3Next、DeepseekV4、Glm4MoeLite/GlmMoeDsa、KimiVL、MistralLarge3、NemotronH、Gemma4、MiniMax-M3、InternS2、Lfm2Moe、HYV3、UnlimitedOCR 等更新的架构。其中 DeepSeek-V3/GLM4-MoE/Llama4/MiniMax-M3 等模型在默认情况下会将共享专家（shared expert）融合进路由专家张量（E = n_routed_experts // ep_size + 1），可通过--disable-shared-experts-fusion关闭该融合行为。

另外需要注意的是，调优时的 block_shape 也会从模型量化配置中自动提取：若quantization_config含weight_block_size则直接使用；若含config_groups则取首组的group_size作为 block_k；针对 AMD MI300X 的 MXFP8（[1, 32]）还会在非 gfx95 平台上重映射为[128, 128]以匹配加载时的行为（源码注释对此有明确说明）。

七、性能对比工具

7.1 SGLang vs vLLM 对比

benchmark_vllm_vs_sglang_fused_moe_triton.py在相同输入与权重下分别调用 vLLM 与 SGLang 的 fused MoE 内核，基于triton.testing.perf_report绘制随 batch size 变化的耗时曲线：

# 默认设置（Mixtral 模型）对比 python benchmark/kernels/fused_moe_triton/benchmark_vllm_vs_sglang_fused_moe_triton.py # Qwen2-57B 的 FP8 模式对比 python benchmark/kernels/fused_moe_triton/benchmark_vllm_vs_sglang_fused_moe_triton.py \ --model Qwen/Qwen2-57B-A14B-Instruct \ --use-fp8-w8a8 # DeepSeek-V3 自定义 TP 大小 python benchmark/kernels/fused_moe_triton/benchmark_vllm_vs_sglang_fused_moe_triton.py \ --model deepseek-ai/DeepSeek-V3-0324 \ --tp-size 8

基准结果以折线图与数据文件形式保存到输出目录（默认./configs/benchmark_ops/vllm_sglang_fused_moe/）。从源码看，对比通过@triton.testing.perf_report装饰器驱动，batch size 覆盖 1～512，两个实现分别以vllm_fused_moe_triton与sglang_fused_moe_triton两条曲线呈现，并复用与调优工具相同的get_model_config解析模型形状，支持 FP8（--use-fp8-w8a8）与自定义--tp-size。

7.2 torch.compile 对比

benchmark_torch_compile_fused_moe.py用于对比torch.compile编译后的内核与原始 fused MoE 内核的性能，用法与 vLLM 对比工具类似。需要特别注意的是：torch.compile不支持fp8_w8a8与int8_w8a8的 fused_moe_kernel，因此对比这两种量化类型时只能测原始内核。两个对比工具目前都支持通过--ep-size参数启用 EP 模式。

八、调优流程速查与注意事项

完整的调优落地流程总结如下：

1. 确认并行配置：根据 GPU 数量确定--tp-size与--ep-size，组合时确保 tp 能被 ep 整除。
2. 选择数据类型：对照模型权重的实际量化格式选择--dtype（fp8_w8a8 / int8_w8a8 / int8_w8a16 / int4_w4a16 / auto）；INT4 与 INT8 系列需要对应的权重张量格式与 scale 参数，调优脚本会据此构造测试张量。
3. 标准调优：对多数场景直接运行tuning_fused_moe_triton.py --tune即可；若追求极致性能或需要 TMA 支持，则按"收集 topk_ids → 启动服务 → 分离调优"的流程使用tuning_fused_moe_triton_sep.py。
4. 部署配置：将生成的 JSON 放入python/sglang/srt/layers/moe/moe_runner/triton_utils/configs/triton_{version}/目录（或通过SGLANG_MOE_CONFIG_DIR指定目录），重启服务后由get_moe_configs自动按模型形状、设备名与 Triton 版本匹配加载。
5. 验证：使用对比工具（vLLM 对比或 torch.compile 对比）或tuning_fused_moe_triton_sep.py --configs ...手动指定配置复核，确认配置在目标 batch size 上的表现。

最后需要强调的是，调优结果与硬件、Triton 版本、驱动环境强相关：仓库源码在加载配置文件时也明确提示"调优环境可能与当前环境存在差异，旧配置可能并非最优"。因此在新环境或升级 Triton 后，建议使用本文的调优脚本重新生成本机配置，以获得与运行环境匹配的最佳内核性能。

【免费下载链接】sglangSGLang is a high-performance serving framework for large language models and multimodal models.项目地址: https://gitcode.com/GitHub_Trending/sg/sglang