Windows 原生编译 SGLang7/8·收尾架构裁剪与 LNK2019 链接收尾前四篇都是撞墙 → 排查 → 修的微观战斗。本篇要往上抽一层,讲一条贯穿整个攻关的主线策略:面对一个为数据中心级 GPU(Hopper / Blackwell)精心优化过的库,在一张消费级显卡(RTX 3090)上编译,最重要的判断不是这段代码报错了怎么修,而是——“这段代码,我的硬件到底用不用得上?用不上的,根本不该编。”这条主线讲透之后,我们处理裁剪之后冒出的最后一道坎(链接期的未解析符号 LNK2019),然后产出 wheel、完成安装验证——与第 1 篇开篇的三条铁证闭环。这是整个编译攻关的收官篇。一、一条主线:架构相关性裁剪为什么修不完,但裁得完sgl-kernel是 sglang 的高性能算子库,它为多种 NVIDIA 架构都写了专门优化的代码路径:Hopper(SM90):H100 / H200 等,FA3、WGMMA、TMA 等专属优化;Blackwell(SM100 系列):更新的数据中心卡,tcgen05、块缩放 MoE 等;以及 FlashMLA(DeepSeek MLA 专用)、各类 SM90a/SM100a 专属 kernel……这些代码路径,绝大多数是为数据中心级算力写的。而我们的目标硬件是RTX 3090(sm_86,Ampere 架构)——一张消费级显卡。如果硬要把这些 Hopper / Blackwell 专属代码也在 MSVC 下逐个编译通过,你会陷入一场打不完的仗:它们用了大量最新架构才有的指令和 CUTLASS 模板,其中一些(如上一篇的 C1001)根本无法在 MSVC 下编译。关键的认知转变是:这些代码,RTX 3090 运行时本来就用不到。sglang 在实际推理时,会根据当前 GPU 的算力自动选择对应的 kernel——在 sm_86 上,它走的是 Ampere 路径(以及通用的 FA2),根本不会调用那些 Hopper / Blackwell 专属实现。既然运行时用不上,编译期就没有理由去编它。“修不完是因为方向错了;换成裁”,就裁得完。裁了哪些,以及判断标准本次攻关里,按架构相关性裁掉的主要有这么几块:裁剪对象性质与 RTX 3090 的关系FlashMLADeepSeek MLA 专用 Hopper kernel无关,OCR/VLM 推理用不到SM100A(Blackwell)自动触发一串compute_100a/120a/103agencode无关,牵连进tcgen05等 Blackwell 头文件FA3(FlashAttention-3)Hopper 专属,几百个 sm_90a kernel无关,3090 走 FA2 即可fp8_blockwise_moe_kernel.cuBlackwell 块缩放分组 MoE无关,且触发 MSVC C1001 崩溃transfer.cuKV-cache 传输,依赖 POSIXdlfcn.h平台无关,Windows 上不可用判断标准很统一:这段代码是不是某个我不拥有的架构(Hopper / Blackwell)专属的?或者是不是依赖某个 Windows 没有的平台特性(POSIX)?是,就裁。裁剪手法:三种关掉的层次裁剪不是粗暴删代码,而是按它为什么该关选对应的关法,这一点很考究:① 用项目自带的开关关(最优雅)。比如 FlashMLA,项目本身就有SGL_KERNEL_BUILD_FLASHMLA这个开关,我们只需在构建命令里传-Ccmake.define.SGL_KERNEL_BUILD_FLASHMLAOFF,不动一行源码。② 修正自动触发的判断条件(治本)。这是最有代表性的一类。比如 SM100A(Blackwell)的开关SGL_KERNEL_ENABLE_SM100A默认是OFF的,但 CMakeLists 里的触发条件写成了:# 原始:只要 CUDA 版本够新就自动开,无视你有没有 Blackwell 卡 if(${CUDA_VERSION} VERSION_GREATER_EQUAL 12.8 OR SGL_KERNEL_ENABLE_SM100A)这个OR前半段是个工具链版本够新就自动开的兜底——它检测的是CUDA 13.1 知不知道 SM100 指令,而不是你这次要不要给 SM100 编译。结果:任何装了 CUDA ≥12.8 的环境,都会被迫编译 Blackwell 专属代码,把tcgen05_ld.h这类 CUDA 自带的 Blackwell 头文件也牵连进来。修法是把那个自动触发去掉,让开关回归它本来的语义:# 改后:只在显式要求时才开,尊重那个默认 OFF 的开关 if(SGL_KERNEL_ENABLE_SM100A)FA3 是同一类问题——它被一个CUDA_VERSION 12.4的条件在 Windows 上强制打开了。修法也一样:给条件加上NOT MSVC,让 FA3 在 MSVC 下不自动开(3090 走 FA2 即可,需要时可显式开启)。这一类版本号自动触发是个值得警惕的反模式:它把工具链支不支持某架构和本次要不要为某架构编译混为一谈。前者是工具链能力探测,后者该由用户根据自己的硬件决定。一旦用OR 版本号把两者绑在一起,就会出现我没有那张卡,却被迫编了它的代码的局面。③ 从源文件列表里整个排除(兜底)。当一个文件整体都是某架构专属、或会触发编译器崩溃时,直接把它从 CMakeLists 的源文件列表里移出去。比如fp8_blockwise_moe_kernel.cu(Blackwell MoE C1001 崩溃),用if(NOT MSVC)把它挡在 MSVC 构建之外;transfer.cu(POSIX 依赖)用if(NOT WIN32)排除。一个细节:排除条件用NOT MSVC还是NOT WIN32,要按为什么排除来选。fp8_blockwise_moe_kernel.cu是 MSVC 编译器崩溃,跟是不是 Windows无关,所以用NOT MSVC(理论上换 GCC 编同样代码不崩);transfer.cu是 POSIX 平台依赖,是 Windows 这个平台的问题,所以用NOT WIN32。条件选得精确,语义才站得住。二、裁剪的代价:最后一道坎 LNK2019架构裁剪解决了编译期的一大批问题,但它带来一个滞后显现的副作用——直到所有源文件都编译通过、进入链接阶段,才暴露出来。现象编译全部通过后,链接器报出 17 个未解析外部符号:common_extension.cc.obj : error LNK2019: unresolved external symbol mscclpp_generate_unique_id common_extension.cc.obj : error LNK2019: unresolved external symbol fp8_blockwise_scaled_grouped_mm common_extension.cc.obj : error LNK2019: unresolved external symbol transfer_kv_cache_cpu_to_gpu ... (共 17 个)成因:注册与实现脱钩了根源在算子注册入口csrc/common_extension.cc。这个文件无条件地注册了所有算子(把它们登记进 torch 的算子表),包括那些我们刚刚在 CMakeLists 里按架构/平台排除掉了实现文件的算子:mscclpp_*(3 个)→ 实现在mscclpp_allreduce.cu,已被if(NOT WIN32)排除;fp8_blockwise_scaled_grouped_mm→ 实现在fp8_blockwise_moe_kernel.cu,已被if(NOT MSVC)排除;transfer_kv_*(13 个)→ 实现在transfer.cu,已被if(NOT WIN32)排除。于是局面就是:注册声明还在(说有这么个算子),但实现没编进来(链接器找不到它的函数体)——注册与实现脱钩,链接器自然报未解析符号。解法:让注册与编译保持同步修法是在common_extension.cc里,给这些算子的注册代码也加上与 CMakeLists 里完全对应的平台守卫,让注册和实现是否编译同步开关:#ifndef_WIN32// 对应 CMakeLists 的 if(NOT WIN32)m.def(mscclpp_generate_unique_id,...);m.impl(mscclpp_init_context,...);// ... mscclpp 与 transfer_kv 系列的注册全包进来#endif#ifndef_MSC_VER// 对应 CMakeLists 的 if(NOT MSVC)m.def(fp8_blockwise_scaled_grouped_mm,...);m.impl(fp8_blockwise_scaled_grouped_mm,...);#endif这里的宏选择跟 CMakeLists 的条件一一对应:if(NOT WIN32)排除的,注册处用#ifndef _WIN32守卫;if(NOT MSVC)排除的,注册处用#ifndef _MSC_VER。(_WIN32在 MinGW 下也定义、_MSC_VER只在 MSVC 下定义,两者语义有细微差别,这里按与 CMakeLists 条件最贴合来选——本项目纯 MSVC,不涉及 MinGW 场景。)这道坎的普适教训:当你按条件排除了某些实现文件(无论是架构裁剪还是平台适配),务必检查算子/符号的注册入口,把注册也同步加上一样的条件。否则实现没了、注册还在,必然在链接期以 LNK2019 收场。注册与实现,要么一起在、要么一起不在。三、收官:wheel 产出与安装验证链接通过,wheel 产出:sgl-kernel/dist/sglang_kernel-0.4.3-cp310-abi3-win_amd64.whl但编出来了不等于能用。完整的验收要走完安装 导入 依赖闭合三步——这正是第 1 篇开篇亮出的三条铁证,这里补全它们的来历。第一步:装 flashinfer-windows(前置依赖)sglang 的 attention 计算后端依赖 flashinfer,但官方包不支持 Windows。需提前安装单独维护的 Windows 兼容 fork 所编译的 wheel:pip install K:\PythonProjects5\Unlimited-OCR\flashinfer-windows\dist\flashinfer_python-0.6.11.post3-py3-none-any.whl --no-depsSuccessfully installed flashinfer-python-0.6.11.post3py3-none-any表示这个 wheel 的 CUDA kernel 代码是 JIT 模式(运行时即时编译),wheel 本身是纯 Python 外壳。它必须在 sglang 主包之前安装到位。必须加--no-deps:flashinfer_python 的 wheel 元数据里声明了对 torch 的依赖。不加这个 flag 时,pip 会重新解析依赖并从默认 PyPI 索引拉 torch——而 PyPI 上裸的torch包是 CPU-only 构建,会悄悄把已经装好的torchcu130换掉,导致后续torch.cuda.is_available()返回False。--no-deps告诉 pip 只装这一个包,不动已经装好的依赖树。第二步:装 sgl-kernel 子包pip install sgl-kernel\dist\sglang_kernel-0.4.3-cp310-abi3-win_amd64.whlSuccessfully installed sglang-kernel-0.4.3第三步:补齐缺失依赖 kernelskernels是 sglang 依赖链里一个未被自动拉取的缺失包。PyPI 上有预编译版本,直接装即可:uv pip install --link-modecopy kernels0.11.7Successfully installed kernels-0.11.7--link-modecopy是 uv 在 Windows 上处理硬链接限制的标准写法。整个包几秒内装完。第四步:装项目自带的主包(必须--no-deps)如第 1 篇所述,项目自带的 sglang 主包元数据把依赖钉死在sglang-kernel0.4.1,跟我们编出的0.4.3对不上,直接装会去公网找不存在的0.4.1win 包而失败。前三步已把所有依赖就位,用--no-deps跳过依赖解析,让已装好的包顶上:pip install K:\...\wheel\sglang-0.0.0.dev11416g92e8bb79e-py3-none-any.whl --no-depsSuccessfully installed sglang-0.0.0.dev11416g92e8bb79e第五步:三项验证python -c import sgl_kernel; print(Kernel OK)Kernel OK这一步:我们逐行抠出来的 C/CUDA 扩展,能被 Python 真正加载——编译成功 ≠ 能加载,这一条证明了后者。python -c import torch; print(CUDA OK:, torch.cuda.is_available())CUDA OK: True这一步:CUDA 运行时可用。pip show sglang sglang-kernel关键看到sglang-kernel 0.4.3的Required-by: sglang,且sglang主包的Requires:列表里含sglang-kernel——这一步:我们编的子包不是孤立产物,而是严丝合缝补进了主包本来需要、Windows 上一直缺失的那个依赖位。python -c import sglang; print(sglang, sglang.__version__)sglang 0.0.0.dev11416g92e8bb79e三步走完,闭环成立:sglang 高性能后端,在原生 Windows 上,装上了、导入了、依赖闭合了。第 1 篇开篇承诺的做成了,到这里给齐了全部证据链。一个诚实的边界说明:本次验收以导入无异常 CUDA 可用 依赖闭合 版本符合预期为准,尚未跑 kernel 级单元测试或端到端推理基准。也就是说,能正确加载并具备运行前提已经坐实;在真实负载下的性能与数值正确性属于进一步的验收工作,可作为后续补充。把话说到这个份上,既不夸大,也不留含糊。小结与下一篇本篇把整个编译攻关收了官:架构裁剪是主线:面对为数据中心 GPU 优化的库,在消费级卡上编译,最重要的不是逐个修,而是判断这段代码该不该编——用不上的架构专属代码(Hopper / Blackwell / FlashMLA / FA3)和平台不支持的代码(POSIX),一律裁掉;裁剪有三种层次:项目开关 修正自动触发条件 整文件排除,按为什么该关选对应手法;警惕版本号自动触发反模式:把工具链支不支持和本次要不要编绑在一起,会逼你编用不上的代码;裁剪的代价是 LNK2019:实现排除了、注册没同步,就会在链接期暴露,解法是让注册与编译条件一一对应;验收要闭环:编出来 ≠ 能用,走完装载 导入 依赖闭合,才算真正做成。到这里,从环境、源码方言、MSVC 语义、到架构裁剪与链接,整条编译路径已经完整记录完毕。但还剩一个问题没回答:这样一场动辄上百轮、跨越多个工具、还时常自己走错路的攻关,是怎么做到不乱、可复现、最终收敛的?最后一篇(第 7 篇·机动篇)不讲具体 bug,讲方法论——台账驱动、幂等补丁脚本、改完即验证的铁律、以及多个 AI 工具协作时如何交接信息、避免互相覆盖。这部分经验,比任何单个补丁都更可迁移到你自己的硬仗里。系列导航全 14 篇编译移植篇怎么把 sglang 从源码编出来00 · 系列总览01 · EPGF 环境地基与岔路口02 · 结论与可行性三铁证 --no-deps03 · 编译篇·前置FlashInfer Windows 源码编译04 · 编译篇·环境关VS 版本、venv 顺序、CUDA 多版本、生成器缓存05 · 移植篇(上)GCC 方言与 MSVC 预处理器严格性06 · 移植篇(下)常量求值、重载决议与编译器崩溃07 · 编译篇·收尾架构裁剪与 LNK2019 链接收尾08 · 方法论台账、幂等补丁脚本与多 AI 协作部署运行篇怎么跑起来并排障09 · 正确启动 SGLang Unlimited-OCR10 · 排障①推理输出乱码/数值错误根因定位11 · 排障②环境变量块超限导致 spawn 子进程崩溃12 · 性能调优RTX 3090 MoE triton autotune config13 · 长文档验证 代理/端口冲突坑 使用指南编译移植篇讲能不能编出来、怎么编部署运行篇讲编出来之后怎么跑通、怎么排障、怎么调优。两篇之间最关键的交叉点本机实际编译用的是 第 07 篇 产物sglang_kernel-0.4.3-cp310-abi3-win_amd64.whl而 第 09 篇 的启动命令正是加载它 Unlimited-OCR 模型。参考资料与延伸阅读以下为本文涉及的官方仓库、文档与规格站建议发布前点一遍确认可达Unlimited-OCR 官方仓库模型与项目源码SGLang 官方仓库SGLang 官方文档启动参数 / OpenAI 兼容 APIflashinfer-windowsWindows 兼容 fork编译前置vllm-windows同作者可对照的 Windows 移植思路PyTorch Windows CUDA 预编译索引cu130NVIDIA CUDA Toolkit 下载uv 官方文档Python 环境治理MSVC /Zc:preprocessor 标准预处理器MSVC 致命错误 C1001编译器内部错误nvcc -Xcompiler 转发 host 编译器选项CMake 生成器Visual Studio / NinjaRTX 3090 规格GA102 / sm_86共享内存 100KBCUDA 共享内存上限与 dynamic_shared_memory 限制Windows 子进程环境变量块限制CreateProcess / ~32KBOpenAI 兼容 API 参考推理调用
Lamp架构(3))
