SGLang W8A8 Block FP8 量化内核配置文件解析与调优指南【免费下载链接】sglangSGLang is a high-performance serving framework for large language models and multimodal models.项目地址: https://gitcode.com/GitHub_Trending/sg/sglang导读SGLang 的 W8A8 Block FP8 矩阵乘法MatMul内核使用一种按矩阵块block共享缩放因子的量化方案在保持精度的同时显著降低显存占用与带宽压力。为了让这颗内核在不同 GPU 型号、不同矩阵形状、不同 batch 大小下都能获得接近最优的性能SGLang 在仓库中内置了一套基于 JSON 的调优配置目录。本文以 configs 目录的官方说明文档 为核心骨架结合 fp8_kernel.py 与 int8_kernel.py 的源码实现完整讲解配置文件的命名规范、JSON 内部格式、内核运行时如何查找并应用这些配置以及当你需要为新的设备或矩阵形状添加调优条目时的具体方法。W8A8 Block FP8 内核的背景W8A8 表示权重Weight与激活Activation都以 8 bit 表示。FP8 使用 E4M3 格式指数 4 bit、尾数 3 bit最大约 448.0。Block FP8 量化的核心思想是不再为整张矩阵使用一个缩放因子而是把矩阵按[BLOCK_N, BLOCK_K]的粒度切成小块每个块独立计算绝对值最大值并生成一个缩放因子。相比 per-tensor 量化block 量化能更好地刻画权重矩阵中不同区域量级差异较大的情况例如 MoE 模型中各 expert 的权重分布差异从而降低量化误差。在 SGLang 中该内核的入口定义在 fp8_kernel.pyw8a8_block_fp8_matmul_triton(A, B, As, Bs, block_size, output_dtype)Triton 实现的 Block FP8 MatMul源码位置w8a8_block_fp8_matmul_deepgemm(...)在启用 DeepGEMM如 DeepSeek 系列模型时调用 deep_gemm 的快速路径源码位置w8a8_block_fp8_matmul(...)统一的测试入口。其中A是激活张量B是权重张量As是激活的逐 token-group 量化缩放矩阵Bs是权重的逐 block 量化缩放矩阵block_size形如[128, 128]。这两个内核同样被上层量化层 fp8_utils.py 引用见 调用点用于如 FP8 权重的线性层前向计算。配置目录与文件命名规则配置文件统一存放在仓库的python/sglang/kernels/ops/quantization/configs/目录下。根据 官方说明文档每个文件名的模式为N{N},K{K},device_name{DEVICE_NAME},dtypefp8_w8a8,block_shape[{BLOCK_N},{BLOCK_K}].json各字段的含义如下字段含义说明N输出维度权重矩阵的列数即B张量的shape[0]K输入收缩维度激活矩阵的列数即B张量的shape[1]需满足A.shape[-1] B.shape[-1]DEVICE_NAMEGPU 设备名空格替换为下划线例如NVIDIA_H100_80GB_HBM3、AMD_Instinct_MI300Xdtype量化数据类型fp8_w8a8或int8_w8a8见下文BLOCK_N/BLOCK_K逐块量化的块粒度目录中绝大多数条目为[128, 128]从仓库实际文件看配置目录命名与文档完全一致例如N1536,K7168,device_nameNVIDIA_H100_80GB_HBM3,dtypefp8_w8a8,block_shape[128, 128].jsonN1536,K7168,device_nameAMD_Instinct_MI300X,dtypefp8_w8a8,block_shape[128, 128].jsonN24576,K7168,device_nameNVIDIA_B200,dtypefp8_w8a8,block_shape[128, 128].json这些 N、K 取值并非随意挑选而是对应真实模型的线性层权重形状。例如K7168是 DeepSeek-V3 / Qwen 系列等大模型常见隐藏维度hidden size或其倍数N7168/24576对应各类 MoE 层、FFN 层与注意力投影层的输出维度。当推理某个模型时内核会按当前权重形状精确命中对应的配置文件。设备覆盖范围从目录清单可以统计出配置已覆盖的主流硬件包括NVIDIA 数据中心 GPUA100-SXM4-80GB、A800-SXM4-80GB、H100 80GB HBM3、H200、H20、B200、L20Y、L40、L40SAMD InstinctMI300X、MI325X文件名中亦有AMD_Radeon_Graphics条目用于 ROCm 环境下的兼容识别。这意味着同一套配置文件机制同时服务于 CUDA 与 ROCm 两条后端内核在运行时根据get_device_name()返回的当前设备名动态选择文件。JSON 配置文件的内部格式配置文件的内容是一个以 Mbatch / token 数为键的字典每个键对应一组 Triton 内核的编译期常量constexpr参数。以一个真实文件为例H100 上 N1536,K7168 的配置{ 1: { BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 64, BLOCK_SIZE_K: 128, GROUP_SIZE_M: 32, num_warps: 4, num_stages: 4 }, 2: { BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 32, BLOCK_SIZE_K: 128, GROUP_SIZE_M: 1, num_warps: 4, num_stages: 5 }, ... 4096: { BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 128, GROUP_SIZE_M: 16, num_warps: 4, num_stages: 3 } }各参数的含义如下参数类型作用顶层键字符串整数该配置适用的 batch 大小M即 token 数量如1、2、4…4096BLOCK_SIZE_MintTriton 线程块CTA在 M 维度上每次处理的 tile 行数BLOCK_SIZE_Nint线程块在 N 维度上每次处理的 tile 列数BLOCK_SIZE_Kint线程块在 K 维度上每次循环处理的 tile 深度需要能被block_k块量化粒度整除GROUP_SIZE_Mint控制程序 ID 在 M/N 维度上的分组调度策略grouped scheduling影响 L2 缓存局部性num_warpsint每个线程块使用的 warp 数Triton 编译参数num_stagesint软件流水线pipelining级数即同时预取的 K 循环阶段数可以看到同一张矩阵在不同 batch 下最优 tile 形状往往不同小 batchM 较小时倾向于较小的BLOCK_SIZE_M/BLOCK_SIZE_N以减少浪费大 batch 则倾向于更大的 tile 与更多 warp。这正是该配置表以 M 为键、按 batch 分段调优的原因。channelwise 变体对于 per-token / per-channel 的 W8A8 FP8 GEMM走triton_scaled_mm路径配置目录中还提供了一类dtypefp8_w8a8_channelwise的文件例如 L40S 上 N4096,K4096 的 channelwise 配置。其格式与 block 版类似字段为BLOCK_SIZE_M/N/K、num_warps、num_stages不含GROUP_SIZE_M对应源码中的get_w8a8_channelwise_fp8_configs与get_w8a8_channelwise_fp8_config源码位置。内核如何查找与应用配置查找流程配置查找逻辑集中在 get_w8a8_block_fp8_configs其工作流程如下跳过 torch.compile 场景当torch._dynamo.is_compiling()为真时直接返回None迫使调用方走默认配置路径避免在编译期执行设备名查询、文件 I/O 等不可追踪non-tensor操作构造文件名调用get_device_name()获取当前设备名将空格替换为下划线再拼接出N{N},K{K},device_name{device_name},dtypefp8_w8a8,block_shape[{block_n}, {block_k}].json定位文件文件路径为os.path.dirname(os.path.realpath(__file__)) /configs/ json_file_name即与内核源码同级的configs目录读取并解析若文件存在则以{int(key): value}的形式把 JSON 解析为Dict[int, Dict]并在日志中打印 Using configuration from ... for W8A8 Block FP8 kernel.兼容性矫正sanitize若某条配置的BLOCK_SIZE_K小于块量化粒度block_k且不满足block_k % BLOCK_SIZE_K 0的整除条件则把BLOCK_SIZE_K钳制clamp到block_k并给出 warning——因为 scale 的步进逻辑要求BLOCK_SIZE_K block_k未命中时回退若文件不存在则记录 warning Using default W8A8 Block FP8 kernel config. Performance might be sub-optimal! 并返回None。int8 变体的查找逻辑几乎一致实现在 int8_kernel.py 的 get_w8a8_block_int8_configs 中仅把dtype替换为int8_w8a8。两类查找函数都带有functools.lru_cache装饰同一 (N, K, block_n, block_k) 组合只做一次文件读取。M 维网格上的最邻近查找得到配置字典后内核需要在不规则的 M 网格上为当前 batch 选择配置。内核调用configs[min(configs.keys(), keylambda x: abs(x - M))]源码位置即选择与当前 M绝对值距离最近的键。例如当配置表键为{1, 2, 4, 8, 16, ..., 4096}时M3 会命中键 2M12 会命中键 16。这种分段近似策略既避免了为每个可能的 batch 单独调优的巨大开销又比单一默认配置更能贴近实际最优。默认配置回退当目录中没有对应文件例如尚未覆盖的 GPU 型号或矩阵形状时内核使用内置默认配置config { BLOCK_SIZE_M: 64, BLOCK_SIZE_N: block_size[0], # 例如 128 BLOCK_SIZE_K: block_size[1], # 例如 128 GROUP_SIZE_M: 32, num_warps: 4, num_stages: 3, }该默认值保证功能正确BLOCK_SIZE_K取块量化粒度block_k满足整除要求但文档与日志均明确提示性能可能次优。对于追求性能的生产环境应当为目标设备与形状补齐调优配置。内核调度细节在 w8a8_block_fp8_matmul_triton 中选中的配置最终以**config的方式作为 Triton 内核的 constexpr 元参数传入对于 gfx1250AMD 新一代架构强制num_stages1并切换到专用内核_w8a8_block_fp8_matmul_gfx1250对于其他 HIP 平台select_w8a8_block_fp8_matmul_kernel会根据网格大小cdiv(M, BLOCK_SIZE_M) * cdiv(N, BLOCK_SIZE_N)与设备计算单元数量对比决定使用常规内核还是手动 unroll x4 的内核_w8a8_block_fp8_matmul_unrolledx4当网格小于设备计算单元数时使用后者源码位置needs_masking K % config[BLOCK_SIZE_K] ! 0决定是否需要对 K 维度做边界掩码masking。此外prepare_block_fp8_matmul_inputs 会严格校验量化 scale 张量的形状例如要求triton.cdiv(N, block_n) Bs.shape[0]、triton.cdiv(K, block_k) Bs.shape[1]float scale 时保证传入的 block 量化缩放矩阵与量化粒度自洽。如何为你的设备/模型添加调优配置结合上面的机制为新的 GPU 或新权重形状添加调优配置的步骤是确认形状确定目标权重矩阵的N输出列数与K收缩维度以及量化块粒度block_n, block_k通常沿用[128, 128]确认设备名在目标机器上打印torch.cuda.get_device_name(0)并把空格替换为下划线例如NVIDIA_H100_80GB_HBM3生成配置表对一组代表性的 M 值如1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048, 4096通过 Triton autotune 或离线基准脚本在BLOCK_SIZE_M ∈ {16, 32, 64, 128, 256}、BLOCK_SIZE_N ∈ {16, 32, 64, 128, 256}、BLOCK_SIZE_K ∈ {64, 128, 256}、不同GROUP_SIZE_M、num_warps ∈ {2, 4, 8}、num_stages ∈ {1..5}的组合中搜索最优参数写入文件按命名规范创建 JSON 文件放入python/sglang/kernels/ops/quantization/configs/目录文件名中block_shape需与block_size参数保持一致验证生效运行推理并在日志中确认出现 Using configuration from ... 的信息若出现 Using default ... Performance might be sub-optimal! 则说明文件名未命中设备名、N、K 或 block_shape 不匹配。注意仓库是只读的上述添加配置属于你本地 fork 或自建部署目录中的操作而非对当前仓库的修改。小结SGLang 的 W8A8 Block FP8 内核配置文件体系是一个数据驱动调优的典型实现以N,K,device_name,dtype,block_shape五元组唯一定位配置文件以 M 为键存储分段最优的 Triton 元参数运行时通过最邻近查找快速选取并在缺失时优雅回退到默认配置。理解这套机制既能帮助你判断线上日志中的配置命中情况也能指导你为自己的硬件与模型形状补充调优数据从而在 FP8 量化的精度收益之外进一步榨取 GEMM 性能。【免费下载链接】sglangSGLang is a high-performance serving framework for large language models and multimodal models.项目地址: https://gitcode.com/GitHub_Trending/sg/sglang创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 
长河君 · 2026/9/10 16:18:57 · 阅读 1.2k SGLang W8A8 Block FP8 量化内核配置文件解析与调优指南
SGLang W8A8 Block FP8 量化内核配置文件解析与调优指南【免费下载链接】sglangSGLang is a high-performance serving framework for large language models and multimodal models.项目地址: https://gitcode.com/GitHub_Trending/sg/sglang导读SGLang 的 W8A8 Block FP8 矩阵乘法MatMul内核使用一种按矩阵块block共享缩放因子的量化方案在保持精度的同时显著降低显存占用与带宽压力。为了让这颗内核在不同 GPU 型号、不同矩阵形状、不同 batch 大小下都能获得接近最优的性能SGLang 在仓库中内置了一套基于 JSON 的调优配置目录。本文以 configs 目录的官方说明文档 为核心骨架结合 fp8_kernel.py 与 int8_kernel.py 的源码实现完整讲解配置文件的命名规范、JSON 内部格式、内核运行时如何查找并应用这些配置以及当你需要为新的设备或矩阵形状添加调优条目时的具体方法。W8A8 Block FP8 内核的背景W8A8 表示权重Weight与激活Activation都以 8 bit 表示。FP8 使用 E4M3 格式指数 4 bit、尾数 3 bit最大约 448.0。Block FP8 量化的核心思想是不再为整张矩阵使用一个缩放因子而是把矩阵按[BLOCK_N, BLOCK_K]的粒度切成小块每个块独立计算绝对值最大值并生成一个缩放因子。相比 per-tensor 量化block 量化能更好地刻画权重矩阵中不同区域量级差异较大的情况例如 MoE 模型中各 expert 的权重分布差异从而降低量化误差。在 SGLang 中该内核的入口定义在 fp8_kernel.pyw8a8_block_fp8_matmul_triton(A, B, As, Bs, block_size, output_dtype)Triton 实现的 Block FP8 MatMul源码位置w8a8_block_fp8_matmul_deepgemm(...)在启用 DeepGEMM如 DeepSeek 系列模型时调用 deep_gemm 的快速路径源码位置w8a8_block_fp8_matmul(...)统一的测试入口。其中A是激活张量B是权重张量As是激活的逐 token-group 量化缩放矩阵Bs是权重的逐 block 量化缩放矩阵block_size形如[128, 128]。这两个内核同样被上层量化层 fp8_utils.py 引用见 调用点用于如 FP8 权重的线性层前向计算。配置目录与文件命名规则配置文件统一存放在仓库的python/sglang/kernels/ops/quantization/configs/目录下。根据 官方说明文档每个文件名的模式为N{N},K{K},device_name{DEVICE_NAME},dtypefp8_w8a8,block_shape[{BLOCK_N},{BLOCK_K}].json各字段的含义如下字段含义说明N输出维度权重矩阵的列数即B张量的shape[0]K输入收缩维度激活矩阵的列数即B张量的shape[1]需满足A.shape[-1] B.shape[-1]DEVICE_NAMEGPU 设备名空格替换为下划线例如NVIDIA_H100_80GB_HBM3、AMD_Instinct_MI300Xdtype量化数据类型fp8_w8a8或int8_w8a8见下文BLOCK_N/BLOCK_K逐块量化的块粒度目录中绝大多数条目为[128, 128]从仓库实际文件看配置目录命名与文档完全一致例如N1536,K7168,device_nameNVIDIA_H100_80GB_HBM3,dtypefp8_w8a8,block_shape[128, 128].jsonN1536,K7168,device_nameAMD_Instinct_MI300X,dtypefp8_w8a8,block_shape[128, 128].jsonN24576,K7168,device_nameNVIDIA_B200,dtypefp8_w8a8,block_shape[128, 128].json这些 N、K 取值并非随意挑选而是对应真实模型的线性层权重形状。例如K7168是 DeepSeek-V3 / Qwen 系列等大模型常见隐藏维度hidden size或其倍数N7168/24576对应各类 MoE 层、FFN 层与注意力投影层的输出维度。当推理某个模型时内核会按当前权重形状精确命中对应的配置文件。设备覆盖范围从目录清单可以统计出配置已覆盖的主流硬件包括NVIDIA 数据中心 GPUA100-SXM4-80GB、A800-SXM4-80GB、H100 80GB HBM3、H200、H20、B200、L20Y、L40、L40SAMD InstinctMI300X、MI325X文件名中亦有AMD_Radeon_Graphics条目用于 ROCm 环境下的兼容识别。这意味着同一套配置文件机制同时服务于 CUDA 与 ROCm 两条后端内核在运行时根据get_device_name()返回的当前设备名动态选择文件。JSON 配置文件的内部格式配置文件的内容是一个以 Mbatch / token 数为键的字典每个键对应一组 Triton 内核的编译期常量constexpr参数。以一个真实文件为例H100 上 N1536,K7168 的配置{ 1: { BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 64, BLOCK_SIZE_K: 128, GROUP_SIZE_M: 32, num_warps: 4, num_stages: 4 }, 2: { BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 32, BLOCK_SIZE_K: 128, GROUP_SIZE_M: 1, num_warps: 4, num_stages: 5 }, ... 4096: { BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 128, GROUP_SIZE_M: 16, num_warps: 4, num_stages: 3 } }各参数的含义如下参数类型作用顶层键字符串整数该配置适用的 batch 大小M即 token 数量如1、2、4…4096BLOCK_SIZE_MintTriton 线程块CTA在 M 维度上每次处理的 tile 行数BLOCK_SIZE_Nint线程块在 N 维度上每次处理的 tile 列数BLOCK_SIZE_Kint线程块在 K 维度上每次循环处理的 tile 深度需要能被block_k块量化粒度整除GROUP_SIZE_Mint控制程序 ID 在 M/N 维度上的分组调度策略grouped scheduling影响 L2 缓存局部性num_warpsint每个线程块使用的 warp 数Triton 编译参数num_stagesint软件流水线pipelining级数即同时预取的 K 循环阶段数可以看到同一张矩阵在不同 batch 下最优 tile 形状往往不同小 batchM 较小时倾向于较小的BLOCK_SIZE_M/BLOCK_SIZE_N以减少浪费大 batch 则倾向于更大的 tile 与更多 warp。这正是该配置表以 M 为键、按 batch 分段调优的原因。channelwise 变体对于 per-token / per-channel 的 W8A8 FP8 GEMM走triton_scaled_mm路径配置目录中还提供了一类dtypefp8_w8a8_channelwise的文件例如 L40S 上 N4096,K4096 的 channelwise 配置。其格式与 block 版类似字段为BLOCK_SIZE_M/N/K、num_warps、num_stages不含GROUP_SIZE_M对应源码中的get_w8a8_channelwise_fp8_configs与get_w8a8_channelwise_fp8_config源码位置。内核如何查找与应用配置查找流程配置查找逻辑集中在 get_w8a8_block_fp8_configs其工作流程如下跳过 torch.compile 场景当torch._dynamo.is_compiling()为真时直接返回None迫使调用方走默认配置路径避免在编译期执行设备名查询、文件 I/O 等不可追踪non-tensor操作构造文件名调用get_device_name()获取当前设备名将空格替换为下划线再拼接出N{N},K{K},device_name{device_name},dtypefp8_w8a8,block_shape[{block_n}, {block_k}].json定位文件文件路径为os.path.dirname(os.path.realpath(__file__)) /configs/ json_file_name即与内核源码同级的configs目录读取并解析若文件存在则以{int(key): value}的形式把 JSON 解析为Dict[int, Dict]并在日志中打印 Using configuration from ... for W8A8 Block FP8 kernel.兼容性矫正sanitize若某条配置的BLOCK_SIZE_K小于块量化粒度block_k且不满足block_k % BLOCK_SIZE_K 0的整除条件则把BLOCK_SIZE_K钳制clamp到block_k并给出 warning——因为 scale 的步进逻辑要求BLOCK_SIZE_K block_k未命中时回退若文件不存在则记录 warning Using default W8A8 Block FP8 kernel config. Performance might be sub-optimal! 并返回None。int8 变体的查找逻辑几乎一致实现在 int8_kernel.py 的 get_w8a8_block_int8_configs 中仅把dtype替换为int8_w8a8。两类查找函数都带有functools.lru_cache装饰同一 (N, K, block_n, block_k) 组合只做一次文件读取。M 维网格上的最邻近查找得到配置字典后内核需要在不规则的 M 网格上为当前 batch 选择配置。内核调用configs[min(configs.keys(), keylambda x: abs(x - M))]源码位置即选择与当前 M绝对值距离最近的键。例如当配置表键为{1, 2, 4, 8, 16, ..., 4096}时M3 会命中键 2M12 会命中键 16。这种分段近似策略既避免了为每个可能的 batch 单独调优的巨大开销又比单一默认配置更能贴近实际最优。默认配置回退当目录中没有对应文件例如尚未覆盖的 GPU 型号或矩阵形状时内核使用内置默认配置config { BLOCK_SIZE_M: 64, BLOCK_SIZE_N: block_size[0], # 例如 128 BLOCK_SIZE_K: block_size[1], # 例如 128 GROUP_SIZE_M: 32, num_warps: 4, num_stages: 3, }该默认值保证功能正确BLOCK_SIZE_K取块量化粒度block_k满足整除要求但文档与日志均明确提示性能可能次优。对于追求性能的生产环境应当为目标设备与形状补齐调优配置。内核调度细节在 w8a8_block_fp8_matmul_triton 中选中的配置最终以**config的方式作为 Triton 内核的 constexpr 元参数传入对于 gfx1250AMD 新一代架构强制num_stages1并切换到专用内核_w8a8_block_fp8_matmul_gfx1250对于其他 HIP 平台select_w8a8_block_fp8_matmul_kernel会根据网格大小cdiv(M, BLOCK_SIZE_M) * cdiv(N, BLOCK_SIZE_N)与设备计算单元数量对比决定使用常规内核还是手动 unroll x4 的内核_w8a8_block_fp8_matmul_unrolledx4当网格小于设备计算单元数时使用后者源码位置needs_masking K % config[BLOCK_SIZE_K] ! 0决定是否需要对 K 维度做边界掩码masking。此外prepare_block_fp8_matmul_inputs 会严格校验量化 scale 张量的形状例如要求triton.cdiv(N, block_n) Bs.shape[0]、triton.cdiv(K, block_k) Bs.shape[1]float scale 时保证传入的 block 量化缩放矩阵与量化粒度自洽。如何为你的设备/模型添加调优配置结合上面的机制为新的 GPU 或新权重形状添加调优配置的步骤是确认形状确定目标权重矩阵的N输出列数与K收缩维度以及量化块粒度block_n, block_k通常沿用[128, 128]确认设备名在目标机器上打印torch.cuda.get_device_name(0)并把空格替换为下划线例如NVIDIA_H100_80GB_HBM3生成配置表对一组代表性的 M 值如1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048, 4096通过 Triton autotune 或离线基准脚本在BLOCK_SIZE_M ∈ {16, 32, 64, 128, 256}、BLOCK_SIZE_N ∈ {16, 32, 64, 128, 256}、BLOCK_SIZE_K ∈ {64, 128, 256}、不同GROUP_SIZE_M、num_warps ∈ {2, 4, 8}、num_stages ∈ {1..5}的组合中搜索最优参数写入文件按命名规范创建 JSON 文件放入python/sglang/kernels/ops/quantization/configs/目录文件名中block_shape需与block_size参数保持一致验证生效运行推理并在日志中确认出现 Using configuration from ... 的信息若出现 Using default ... Performance might be sub-optimal! 则说明文件名未命中设备名、N、K 或 block_shape 不匹配。注意仓库是只读的上述添加配置属于你本地 fork 或自建部署目录中的操作而非对当前仓库的修改。小结SGLang 的 W8A8 Block FP8 内核配置文件体系是一个数据驱动调优的典型实现以N,K,device_name,dtype,block_shape五元组唯一定位配置文件以 M 为键存储分段最优的 Triton 元参数运行时通过最邻近查找快速选取并在缺失时优雅回退到默认配置。理解这套机制既能帮助你判断线上日志中的配置命中情况也能指导你为自己的硬件与模型形状补充调优数据从而在 FP8 量化的精度收益之外进一步榨取 GEMM 性能。【免费下载链接】sglangSGLang is a high-performance serving framework for large language models and multimodal models.项目地址: https://gitcode.com/GitHub_Trending/sg/sglang创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考