尧图网站建设 免费方案咨询
ARTICLE DETAIL

资讯详情

深耕网站建设、视觉设计与SEO优化的一线实战洞察。

首页/资讯中心/资讯详情

SGLang 适配 AMD 平台的开发工作全面梳理

发布时间：2026/11/4 7:59:26 尧图网站建设
SGLang 适配 AMD 平台的开发工作全面梳理

SGLang 的核心分层是：纯 Python 的调度/前端层（硬件无关）→ 算子与后端层（sgl-kernel、attention backend、MoE、量化、通信）→ 构建/运行环境层。从 NVIDIA 迁到 AMD，绝大部分工作集中在后两层。下面按模块逐项梳理。

一、构建系统与运行环境

NVIDIA 路径：官方 pip wheel 开箱即用，CUDA toolkit + flashinfer 等预编译依赖直接装。

AMD 路径需要额外开发：

  • sgl-kernel 需用专门的 ROCm 构建脚本源码编译：cd sgl-kernel && python setup_rocm.py install；Python 包也要切换pyproject_other.toml并执行pip install -e "python[all_hip]"
  • 维护独立的 ROCm Dockerfile 与镜像（rocm/sgl-dev、lmsysorg/sglang:*-rocm*-mi30x），并按 GPU 架构分镜像：gfx942（MI300X/MI325X）、gfx950（MI350/MI355X）、gfx1151（Strix Halo，wave32 架构，需专门镜像且关闭 AITER）
  • 容器与宿主适配：暴露/dev/kfd、/dev/dri，加入 video/render 组，设置HIP_FORCE_DEV_KERNARG、RAY_EXPERIMENTAL_NOSET_ROCR_VISIBLE_DEVICES等环境变量
  • 条件依赖管理：pyproject.toml需要按平台分支（CUDA 版依赖 flashinfer/flash-attn，ROCm 版依赖 aiter/triton）

二、底层算子库（sgl-kernel / sglang-kernel）移植

这是工作量最大的部分。sgl-kernel 覆盖 MLA decode、GEMM（FP8/INT8/AWQ/GPTQ）、MoE topk、RMSNorm、RoPE、采样、投机解码等大量自定义 CUDA kernel。

  • CUDA→HIP 转换：hipify 工具可自动处理约 95% 的语法，但必须手工处理的关键点包括：
    • warp size 32→64（AMD wavefront 为 64 线程，所有依赖 warp 内 shuffle/ballot 语义的 kernel 都要改）
    • 删除/重写 PTX 内联汇编（如 PDL programmatic dependent launch）
    • NCCL→RCCL API 替换
  • 双后端条件编译：C++ 头文件用#ifdef USE_ROCM选择 CUDA/ROCm 两套实现，Python 侧按torch.version.hip分支导入不同算子
  • NVIDIA 专有库的替代：flashinfer 的 silu_and_mul / rmsnorm / RoPE / sampling 等算子在 AMD 上要么用 PyTorch 回退实现，要么接入 AITER/CK 等价算子；架构检测从is_sm90_supported换成is_gfx942/is_gfx950，NVTX 标记换成 roctx
  • CUTLASS/cuBLAS 类依赖替换：换成 Composable Kernel（CK）、hipBLASLt，并需要按 (M, N, K) shape 做 CSV 驱动的 GEMM tuning（AITER 提供 shape 采集和调优工作流）

三、Attention 后端开发

NVIDIA默认有 flashinfer、FlashAttention 3、trtllm_mha/gen、FlashMLA 等多套后端；这些全部不可用（FA3 是 Hopper+ 专属、TRT-LLM 无 ROCm 移植计划）。

AMD 需要：

  • 开发 Triton attention backend（HIP 版 Triton）和AITER backend（新增aiter_backend.py，遵循BaseAttnBackend接口封装 AITER 的 attention API，约 300–400 行）
  • 处理硬件约束差异：如 CK batch-prefill 路径 head_dim 上限 256，head_dim 512 的模型（如 Gemma-4）prefill 要路由到 Triton unified attention；某些路径禁止 fp8_e4m3 KV cache（batched decode 输出错误），须强制 bf16 KV
  • MLA 专项：NVIDIA 用 FlashMLA/cutlass_mla_decode；AMD 用 AITER 的 MLA decode kernel（gfx950 上走 gluon sparse kernel，其他架构走 tilelang partial+combine），还要做 fused decode MLA（SGLANG_ROCM_FUSED_DECODE_MLA）等开关
  • 长上下文 decode 的分区策略：如SGLANG_AITER_DECODE_PARTITION_BUCKETS按序列长度桶生成多个 decode graph

四、MoE 与专家并行（EP）

  • All-to-All 通信后端替换：NVIDIA 用 DeepEP（基于 NVSHMEM/IBGDA）、pplx-kernels、NIXL；AMD 需要接入MORI-EP（--moe-a2a-backend mori），即 AMD 原生的 dispatch/combine 实现
  • MoE 计算后端替换：NVIDIA 的 DeepGEMM/CUTLASS/TRT-LLM runner 不可用，AMD 走 Triton runner 或 AITER fused MoE（CK 内核），还要处理 MoE padding、gate/up 权重 tile 布局（interleave vs separated）等差异
  • 计算通信重叠适配：TBO（two-batch overlap）在 AMD 上要基于 MORI 的异步 API 重做双流流水线，并利用 SDMA copy engine 做零算力开销的数据搬运
  • ROCm 特有的多流限制：HIP 默认硬件队列只有 4 个，共享专家与路由专家重叠需设置GPU_MAX_HW_QUEUES>=5并开SGLANG_ROCM_USE_MULTI_STREAM
  • 还需修多节点稳定性问题：MORI inter_node 模式曾出现 CUDA graph capture 挂起和高并发下 dispatch/combine RDMA 死锁

五、量化路径重建

量化是平台差异最明显的领域之一：

  • NVIDIA 专有路径全部失效：Marlin（awq_marlin、gptq_marlin）、modelopt_fp8/fp4、gguf 在 AMD 上不支持
  • AMD 替代路径：FP8 走 AITER 或 Triton（DeepSeek-V3/R1 预量化模型开箱即用）；AWQ 用 Triton 反量化 kernel（无 Marlin 加速）；MXFP4 需要 CDNA3/CDNA4 +SGLANG_USE_AITER=1；NVFP4 模型需通过 Petit kernel（petit_nvfp4）；另有 AMD 专属的在线量化quark_int4fp8_moe
  • 新量化格式的适配常留尾巴，例如 MiniMax-M3 的 MXFP8 需要 PR 把mxfp8加进 ROCm 支持列表并绕过 bf16 router GEMM 缺口
  • block-scale FP8 GEMM 还需在 CK 与 Triton 实现间做精度/速度权衡（gfx942 默认 CK，但 GSM8K 等评测有小幅精度差）

六、CUDA Graph 与投机解码

  • HIP graph 基本可用但坑更多：EAGLE 投机解码在 MI300X 上开启 CUDA graph 时，因tree speculative sampling kernel 在 HIP 上不可用而回退到 greedy verification，且该回退路径存在不确定性死锁
  • 需要移植/重写投机解码相关 kernel（verify、draft_extend），AITER 侧有 unified verify / draft_extend 的独立开关（SGLANG_AITER_UNIFIED_VERIFY等）
  • 内存策略差异：AITER 注意力在长上下文（>8K）下要把mem-fraction-static乘 0.85 预留非静态 workspace，否则 OOM——这类默认值都要为 AMD 单独调校

七、通信与 PD 分离

  • NCCL→RCCL：API 基本兼容，但环境变量、拓扑感知、allreduce 性能调优要重做；SGLang 的 custom all-reduce CUDA kernel 需要 HIP 移植，或改用 AITER 的 all-gather 实现（SGLANG_USE_AITER_AG）
  • PD 分离的 KV cache 传输：NVIDIA 走 Mooncake/NIXL；AMD 用 MORI-IO 做 p2p KV 传输

八、工具链、调试与 CI

  • 性能分析：Nsight Systems/Compute → rocprof / roctx / rpd_tracer；torch profiler 的 kineto 在 ROCm 上的行为差异要适配
  • CI 基础设施：需要 MI300/MI350 的 GPU runner，为 ROCm 单独跑测试集（很多 kernel 路径依赖特定 gfx 架构，无法进常规 CI，只能人工覆盖精度与性能验证）
  • 单 kernel 正确性验证策略：warp 语义变化、原子操作、共享内存 bank conflict 模式在 wave64 下都不同，移植后需逐一跑数值对齐测试

九、平台检测与代码组织约定

整个 Python 层要建立统一的设备抽象：所有硬件分支收拢到if torch.version.hip:条件或独立文件中，保证 NVIDIA 路径零改动；工具函数层（is_cuda()/is_hip()、显存查询、设备属性）要封装好，因为torch.cuda.*在 ROCm 下虽是别名可用，但显存上报、SM/CU 数量、wave size 等语义不同。

工作量参考

以精简版 mini-sglang 的评估为例，完整 AMD 适配约1000–1200 行改动、涉及约 20 个文件，可拆成 4 个 PR：设备抽象与构建系统（~200 行）、HIP kernel 移植 + RCCL（~400 行）、AITER attention backend（~300–400 行）、Docker/文档/CI（~150 行）。完整版 SGLang 则是长期持续投入：量化矩阵、每个新模型的 per-arch tuning、MORI 多节点稳定性、投机解码 kernel 补齐都需要随版本迭代。

总结：Python 调度层几乎零成本；真正的开发工作在「kernel 移植（hipify + 手工修 warp/汇编）→ 用 AITER/Triton/CK 重建 attention、MoE、量化后端 → 用 MORI/RCCL 替换 NVSHMEM/NCCL 通信 → 处理 CUDA graph 与投机解码的 HIP 特有缺陷 → 按 gfx 架构做调优与 CI」这条链路上。回报是显著的——AITER 集成后 DeepSeek-R1 在 MI300X 上端到端吞吐从 6484 提升到 13704 tok/s（约 2.1 倍）。

参考1：工作内容梳理一

一、构建系统与安装（对应「一、构建系统与运行环境」）

  1. SGLang 官方文档 · AMD GPU 平台安装指南
    https://docs.sglang.ai/_sources/platforms/amd_gpu.md （中文版镜像：https://sglang.org/zh/platforms/amd_gpu ）
    —— 支撑：源码安装流程setup_rocm.py install、切换pyproject_other.toml+pip install -e "python[all_hip]"、rocm.Dockerfile、Docker 运行参数（/dev/kfd、/dev/dri、--group-add video等）、GRUB/NUMA 系统级调优。
  2. AMD ROCm 官方文档 · SGLang inference and serving on ROCm
    https://rocm.docs.amd.com/projects/ai-ecosystem/en/latest/inference/sglang.html
    —— 支撑：AMD 官方维护的 SGLang on ROCm 支持矩阵（ROCm 版本 × Instinct/Radeon × gfx 架构）。
  3. sgl-project/sglang PR #33939：[AMD] Add gfx1151 (Strix Halo) Docker image
    https://github.com/sgl-project/sglang/pull/33939 （镜像解析：https://prhub.com.cn/sgl-project/sglang/pull/33939 ）
    —— 支撑：gfx1151 专用镜像、setup_rocm.py仅接受 gfx942/gfx950/gfx1250 需打补丁、RDNA 上SGLANG_USE_AITER=0且默认--attention-backend triton等细节。

二、sgl-kernel 内核移植（对应「二、底层算子库移植」）

  1. sgl-project/mini-sglang Issue #90：Interest in adding AMD GPU (ROCm) support
    https://github.com/sgl-project/mini-sglang/issues/90
    —— 这是「工作量参考」的直接出处：~1000–1200 行 / ~20 个文件 / 4 个 PR 的拆分；hipify-perl 覆盖约 95% 语法，warp 32→64、PDL 内联汇编、NCCL→RCCL 需手工处理；flashinfer 算子替换与 NVTX→roctx；新增attention/aiter.py后端（~300–400 行）等。
  2. sglang-kernel（sgl-kernel）内核库完全指南（GitCode 博客）
    https://blog.gitcode.com/73c0a33a7f2453365058fcff9bc3fca1.html
    —— 支撑：sgl-kernel 按平台拆分pyproject_rocm.toml/setup_rocm.py、按torch.version.hip动态加载 ROCm 算子的机制。
  3. SGLang 源码阅读 · sgl-kernel “five wheels, one tree”
    https://www.jinnpan.com/sources/sglang.html
    —— 支撑：一份 C++/Triton 源码、五个平台 wheel（CUDA/ROCm/CPU/MUSA/Metal）的工程组织方式，ROCm 用 CMake + hipcc + Composable Kernel 构建。

三、Attention 后端与 AITER（对应「三、Attention 后端开发」）

  1. AMD ROCm 博客 · Accelerate DeepSeek-R1 Inference: Integrate AITER into SGLang
    https://rocm.blogs.amd.com/artificial-intelligence/aiter-intergration-s/README.html
    —— 支撑：AITER 替换 MoE / attention / GEMM / all-reduce 的集成点、AITER 的 JIT + ASM/CK/HIP 多级 kernel 层次，以及 prefill 延迟降 52%、decode 降 47%、吞吐翻倍的实测数据。
  2. AMD ROCm 博客 · AITER: AI Tensor Engine For ROCm
    https://rocm.blogs.amd.com/software-tools-optimization/aiter-ai-tensor-engine/README.html
    —— 支撑：结尾引用的性能数字来源（SGLang + DeepSeek V3/R1 在 MI300X 上 6484.76 → 13704.36 tok/s，2 倍以上提升）。
  3. AMD ROCm 博客 · Supercharge DeepSeek-R1 Inference on AMD Instinct MI300X
    https://rocm.blogs.amd.com/artificial-intelligence/DeepSeekR1-Part2/README.html
    —— 支撑：AITER kernel 级加速幅度（GEMM +2x、MoE +3x、MLA decode +17x、MHA prefill +14x）、MI300X vs H200 的对比、SGLang 服务端超参调优（chunked_prefill_size 等）。

四、MoE / EP 与通信（对应「四、MoE 与专家并行」「七、通信与 PD 分离」）

  1. SGLang 文档 · Expert Parallelism（Hardware-Specific Configuration）
    https://sgl-project-sglang-93.mintlify.app/distributed/expert-parallelism
    —— 支撑：AMD 上--moe-a2a-backend mori的配置方式及「MORI 目前只支持 normal mode」。
  2. GitHub · ROCm/mori（Modular RDMA Interface）
    https://github.com/rocm/mori
    —— 支撑：MORI-EP / MORI-IO 已集成进 SGLang 和 vLLM 用于 MoE EP 与 PD 分离、MORI-SHMEM、扩展到 64 GPU 等项目动态。
  3. SGLang 文档 · Environment Variables（MORI Configuration）
    https://sgl-project.github.io/references/environment_variables.html
    —— 支撑：SGLANG_MORI_DISPATCH_DTYPE、MORI_DISABLE_AUTO_XGMI等 MORI 调优环境变量。
  4. AMD 招聘 JD · AI Software Engineer (Shanghai)
    https://bebee.com/cn/jobs/ai-software-engineer-high-performance-amd-shanghai-shanghai–t7xk-751205891
    —— 支撑：AMD 官方对「把 MORI-EP/IO 落进 SGLang/vLLM」工作内容的描述（dispatch/combine、KVCache transfer、EPLB、8–64 GPU EP 维护），可作为「该方向是持续工程投入」的旁证。

五、PD 分离 KV 传输（对应「七、通信与 PD 分离」）

  1. Infera 文档 · PD disaggregation（Mooncake vs MoRIIO 对比）
    https://rocm.docs.amd.com/projects/infera/en/main/features/pd_disaggregation.html
    —— 支撑：SGLang 侧--disaggregation-transfer-backend mooncake/mori的差异（write/push vs read/pull、bootstrap server vs ZMQ）。
  2. SGLang 文档 · DeepSeek-V4 Cookbook · PD Disaggregation on MI355X (MORI-IO)
    https://lmsysorg.mintlify.app/cookbook/autoregressive/DeepSeek/DeepSeek-V4
    —— 支撑：GPU_MAX_HW_QUEUES=5、MORI_SHMEM_HEAP_SIZE、MegaMoEv2、FP8 KV pool 拆分等 AMD 侧 PD 调优细节。
  3. sgl-project/sglang PR #26288：[PD][AMD] incremental KV transfer with decode radix cache
    https://github.com/sgl-project/sglang/pull/26288 （镜像解析：https://prhub.com.cn/sgl-project/sglang/pull/26288 ）
    —— 支撑：mori 后端的增量 KV 传输与 decode radix cache 适配工作（mori/conn.py、参数白名单等）。
  4. Prefill/decode disaggregation reaches production: state of play（lecompute.fr）
    https://lecompute.fr/en/runtimes/disaggregation-prefill-decode-production/
    —— 支撑：MORI-IO 作为 AMD 的 NIXL 对位、SGLang 同时支持 NIXL / Mooncake / MORI-IO 三种 KV 传输。

六、CUDA Graph 与投机解码（对应「六、CUDA Graph 与投机解码」）

  1. sgl-project/sglang Issue #29347：[Bug][AMD][ROCm/MI300X] EAGLE + CUDA graph 不确定性死锁
    https://github.com/sgl-project/sglang/issues/29347
    —— 支撑：tree speculative sampling kernel 在 HIP 上不可用 → 回退 greedy verification → 约 1/3 概率挂死；--disable-cuda-graph或去掉 EAGLE 作为 workaround，但均牺牲吞吐。

七、补充阅读

  1. SGLang 遇上 ROCm 首测体验（CSDN）
    https://blog.csdn.net/2600_96323197/article/details/162310474
    —— ROCm 7.x 下源码编译实践：PYTORCH_ROCM_ARCH="gfx942"、Triton ROCm 分支锁定等踩坑记录。
  2. AMD ROCm 文档 · SGLang inference performance testing（MAD 基准）
    https://rocm.docs.amd.com/en/docs-6.4.2/how-to/rocm-for-ai/inference/benchmark-docker/sglang.html
    —— 官方基准复现流程，可用于移植后的性能回归验证。

参考2: 工作内容梳理2

作为一家新的 GPGPU 公司，要让 SGLang 在你的硬件上高效运行，核心工作不是修改 SGLang 源码，而是通过其 Out-of-Tree（OOT）插件体系，以旁加载（side-installed）的方式补齐硬件特有的实现。SGLang 近年来的架构演进，正是为了系统性地降低新硬件厂商的接入成本。

适配工作可以归纳为四个核心层面：平台插件、算子库、通信库，以及可选的通用插件钩子。

一、硬件平台插件（Hardware Platform Plugin）：声明“你的硬件是什么”

这是最基础的一层。你需要发布一个 Python 包，在其中实现一个SRTPlatform的子类，并通过setuptools入口点sglang.srt.platforms注册它。SGLang 启动时，会通过activate()回调自动探测你的硬件是否可用，并加载对应的平台类。

平台类需要覆盖 SGLang 运行时依赖的设备侧核心抽象，主要包括：

需要实现的接口/能力作用说明
DeviceMixin设备操作设备身份识别（如is_gpgpu()的等价逻辑）、set_device()、get_device_name()、显存查询等基础操作
KV Cache 池通过build_kv_pool(request)或工厂方法，返回适配你显存布局的 KV Cache 池类；需覆盖 MHA、MLA、稀疏注意力等不同架构的池实现
Attention 后端提供 Prefill 和 Decode 阶段的 Attention 执行后端，使 SGLang 的注意力调度能调用你的硬件实现
Graph 捕获若你的硬件支持图执行（类似 CUDA Graph），需要提供图捕获与回放的后端；这对降低解码阶段 CPU 开销至关重要
编译后端若你的软件栈提供torch.compile的等效后端，需要在此声明；否则 SGLang 会回退到 eager 模式

这一层的设计原则是fail-fast：未实现的方法默认抛出NotImplementedError，迫使你在适配初期就暴露缺失项，而不是静默产生错误结果。

二、算子库：提供“在你的硬件上怎么算”

这是性能差异的最大来源。SGLang 的核心热点算子（Attention、GEMM、归一化、激活、MoE 路由等）在 NVIDIA 生态中依赖flashinfer、sgl_kernel等 CUDA 专用库；在你的 GPGPU 上，需要用你的软件栈重新实现这些算子。

有两种并存的算子集成路径，建议同时使用：

  1. 自定义算子库（静态注册）：将高性能 Kernel 编译成.so或 Python 扩展，通过 PyTorch 的TORCH_LIBRARY_FRAGMENT机制注册为torch.ops.<你的命名空间>.<op_name>。当 SGLang 检测到你的平台激活时，自动加载该算子库，并在热路径中替换 CUDA 实现。昇腾 NPU 的sgl_kernel_npu就是这一模式的参考。

  2. OOT 算子分发（动态优先级替换）：SGLang 提供了MultiPlatformOp的 OOT 注册机制。你的插件可以通过register_oot_impl()为指定的算子（如SiluAndMul、RMSNorm）注册VENDOR类型的实现，并指定优先级。运行时，dispatch_forward()会优先查找 OOT 注册表；若你的算子可用，则调用你的实现；否则回退到内置或其他后端。可以通过SGLANG_OOT_PREFER=vendor全局设定你的算子优先，或通过SGLANG_OOT_OP_PREFER按算子精确控制。

需要优先覆盖的算子类别（参考摩尔线程 MATE 库的实践）：Attention（对接 FlashAttention / FlashMLA 等接口语义）、GEMM / DeepGEMM、归一化（RMSNorm、LayerNorm）、激活函数、MoE 相关路由与专家计算、采样（Sampling）算子。

三、分布式通信：让多卡跑起来

SGLang 的分布式推理依赖集合通信。你的 GPGPU 通常不会直接运行 NCCL，因此需要在 OOT 框架下注册通信器。

具体做法是：通过GroupCoordinator.register_oot_communicator(factory_fn)注册一个通信器工厂，并提供all_reduce、reduce_scatter_tensor、all_gather_into_tensor等集合操作的实现。你的通信库可以基于自研的集合通信库（类似 HCCL、FlagCX 的定位），或封装torch.distributed中你已支持的后端。通信器工厂的优先级高于 SGLang 内置的硬件通信逻辑，确保你的路径被选中。

四、通用插件钩子（可选，但强烈建议用于长尾兼容）

SGLang 还提供了sglang.srt.plugins入口点的通用插件机制，支持before / after / around / replace类型的钩子，可以拦截任意函数或方法。当某些模型或路径中存在平台接口无法优雅覆盖的硬件相关分支时，可以用通用插件做临时性兼容，避免 fork SGLang。但应优先通过平台接口和算子注册解决问题；通用钩子更适合作为模型支持初期的过渡手段。

参考路径：以摩尔线程 MUSA 为镜

摩尔线程的 MUSA 后端从“代码共建”走到 SGLang 主线“原生支持”，提供了一条可参照的路线图：在算子层开源了 MATE 算子库，对接 FlashAttention、FlashMLA、DeepGEMM 等主流接口，并通过多平台分发机制将 MUSA 实现路由到 SGLang 执行路径中；同时完成了从环境构建、算子覆盖到分布式推理的全链路打通。你的适配工作本质上是在重复这条已被验证的路径，只是硬件和软件栈不同。

建议的推进顺序

第一步：最小可运行闭环——实现SRTPlatform的设备操作与 KV 池，提供最基本的 Attention 和 GEMM 算子，确保单卡上能跑通一个主流模型（如 Qwen 或 Llama 系列）的generate请求。

第二步：性能关键路径——覆盖 Prefill / Decode Attention、MoE、采样等热点算子，接入图捕获，验证吞吐与延迟是否进入可用区间。

第三步：分布式与长尾模型——注册集合通信器，跑通 TP / PP / EP 配置；针对目标客户的重点模型（如 DeepSeek、Kimi 等）补齐模型特有的算子与缓存路径。

第四步：上游化（可选但有利于客户信心）——像摩尔线程那样将适配代码合入 SGLang 主线，获得“官方原生支持”的身份，会显著降低客户在选型时的采用顾虑。

网站建设 视觉设计 SEO优化
返回列表