Q: 如何修复使用 FreeToken 提供服务时 Qwen3.5/Qwen3.6 RMSNorm 权重中的 bf16 精度损失问题？

问题：FreeToken 的 `models/qwen3_5_moe` 加载器在加载时将 Gemma 风格的 `+1` 折叠进 bf16 检查点张量（`tensor + 1.0`），存储的是 `bf16(1 + w)`。而 Transformers 则保留以零为中心的权重，并在 fp32 中计算 `(1 + w)`，因此这种折叠会把较小的权重舍入掉——在 `nvidia/Qwen3.6-35B-A3B-NVFP4` 上，171,008 个归一化乘数中有 77.4% 发生变化，平均相对误差为 1.4e-3，最大为 3.9e-3（半个 bf16 ulp），影响到每个 token 上的 `input_layernorm`、`post_attention_layernorm`、`self_attn.q_norm`、`self_attn.k_norm` 以及最终的 `model.norm`。修复方法：存储原始的零中心 bf16 权重 `w`（不要在加载器中加 1），并在运行时以 fp32 应用 `+1`——使用现有的 `GemmaPlusOneRMSNorm` 层（`layers/norm.py`），或调用 `sgl_kernel.gemma_rmsnorm(x, w)` 而不是 `sgl_kernel.rmsnorm(x, bf16(1+w))`。这与 `qwen4_exp` 加载器已采用的方法一致，并恢复了与 fp64 参考结果 99.999% 的 bf16 一致性（残余误差仅为输出舍入，rel L2 约为 1.7e-3，而修复前为 3.0e-3）。

Question 1

为什么 FreeToken 在使用 --allowed-local-media-path 白名单时拒绝 'file://localhost' 图像 URI？

Accepted Answer

这是 `_load_local_media`（`python/freetoken/mm/media.py`）中的一个 URI 解析 bug。该函数通过拼接 URI 的 `netloc` 和 `path` 来构建文件路径，因此 `file://localhost/path/image.png` 会变成相对路径 `localhost/path/image.png` 而不是 `/path/image.png`，这随后导致解析后路径的白名单检查失败并抛出 `ValueError`。修复方法：按照 RFC 8089 将空 authority 和 `localhost` authority（不区分大小写）视为本地，拒绝任何其他非本地 authority，使用 `url2pathname` 解码路径，并保留现有的 `Path.resolve()` 严格根目录白名单检查。在补丁发布前的临时解决方案：使用规范的 `file:///path/image.png` 形式（空 authority），而不是 `file://localhost/...`。

Question 2

为什么 FreeToken 在 RTX 5090 / 5070 Ti (sm_120) 上首次启动 TMA kernel 后会 OOM 或丢失 1.3-3.3 GiB 的 VRAM？

Accepted Answer

在 sm_120 GPU 上，低于 3.8 版本的 Triton 会将 TMA 加载（`TensorDescriptor`）编译为一个驱动程序系统调用（syscall），该调用会在首次 kernel 启动时将每个驻留线程的每线程栈上限从 1024 B 提升至约 14.4 KB。这会在 70-SM 的 RTX 5070 Ti 上永久预留约 1.34 GiB，在 170-SM 的 RTX 5090 上预留约 3.3 GiB。由于该预留发生在 KV pool 和 expert cache 根据空闲 VRAM 确定大小之后，服务器会发生 OOM。修复方法：升级到 torch 2.14 + Triton 3.8 或更高版本。Triton 3.8 在 sm_120 上会生成原生 `cp.async.bulk.tensor.2d.shared::cta` 指令，该指令无需系统调用，也不会造成栈预留。引擎本身已经要求这些版本；如果你遇到此问题，请检查是否存在仍在使用旧版锁定依赖的过时 venv 或源码构建，并使用更新后的 requirements 重新安装。

Question 3

如何修复分块预填充（chunked prefill）期间出现的 'expandable_segments: memory mapping failed with OOM' CUDACachingAllocator 警告？

Accepted Answer

降低 `--max-prefill-length`（例如恢复到默认值 8192）。该标志设置的是 chunked-prefill 的分块大小，而不是服务器可处理的最大提示词长度——30k-token 的请求在 8192 下仍能正常处理，只是会被拆分为更多分块。出现这些警告是因为较大的分块会在 DSA/MLA 注意力路径中造成较大的瞬时激活峰值（多个存活的 [T, 64, 512] bf16 张量，在 12288-token 分块时约占用 3.8 GiB），从而导致 torch 的 expandable-segments 分配器重试内存映射。每轮重试突发都会带来一次设备级同步停顿（1.5-5 秒），但会自行恢复，因此请求仍会正常完成，不影响正确性。注意：当 MoE 专家缓存对 overlap 模式而言过小时，较大的分块总体上仍可能更快，因为每个分块的完整专家 PCIe 拷贝开销可分摊到更多 token 上——因此需要在无停顿的日志与峰值 prefill 吞吐量之间进行权衡。

Question 4

升级到 torch 2.14.1 时如何修复 sglang-kernel 加载错误和 "C++20 or later compatible compiler is required" 错误？

Accepted Answer

sglang-kernel 链接了 libtorch 的 C++ 符号，并且只能针对其构建时所用的 torch 加载，因此两者的版本必须一起锁定：使用 `sglang-kernel==0.4.9` 搭配 `torch>=2.14,<2.15`（0.4.5 是针对 torch 2.11 构建的版本）。torch 2.14 需要 Triton 3.8，因此请设置 `triton>=3.6,<3.9` 和 `torchvision>=0.29,<0.30`。torch 2.14.1 的 PyPI wheel 是 cu130 构建版本，需要 `cuda-toolkit==13.0.3`，因此请让你的 installer/uv 索引保持在 cu130 上。如果你要编译 C++ 扩展，请将 setup.py 中的 `-std=c++17` 改为 `-std=c++20`，否则 torch 2.14 的头文件会报错失败，提示 `#error C++20 or later compatible compiler is required`。注意：Triton 3.8 导致部分 kernel 出现性能回退（例如，用于 NVFP4 MoE 解码的 `_decode_nvfp4_marlin_kernel` 慢了 40-130%）；其中八处回退已在后续 PR 中修复，其余回退对端到端性能的影响低于约 4%。

Question 5

如何修复使用 FreeToken 提供服务时 Qwen3.5/Qwen3.6 RMSNorm 权重中的 bf16 精度损失问题？

Accepted Answer

问题：FreeToken 的 `models/qwen3_5_moe` 加载器在加载时将 Gemma 风格的 `+1` 折叠进 bf16 检查点张量（`tensor + 1.0`），存储的是 `bf16(1 + w)`。而 Transformers 则保留以零为中心的权重，并在 fp32 中计算 `(1 + w)`，因此这种折叠会把较小的权重舍入掉——在 `nvidia/Qwen3.6-35B-A3B-NVFP4` 上，171,008 个归一化乘数中有 77.4% 发生变化，平均相对误差为 1.4e-3，最大为 3.9e-3（半个 bf16 ulp），影响到每个 token 上的 `input_layernorm`、`post_attention_layernorm`、`self_attn.q_norm`、`self_attn.k_norm` 以及最终的 `model.norm`。修复方法：存储原始的零中心 bf16 权重 `w`（不要在加载器中加 1），并在运行时以 fp32 应用 `+1`——使用现有的 `GemmaPlusOneRMSNorm` 层（`layers/norm.py`），或调用 `sgl_kernel.gemma_rmsnorm(x, w)` 而不是 `sgl_kernel.rmsnorm(x, bf16(1+w))`。这与 `qwen4_exp` 加载器已采用的方法一致，并恢复了与 fp64 参考结果 99.999% 的 bf16 一致性（残余误差仅为输出舍入，rel L2 约为 1.7e-3，而修复前为 3.0e-3）。

FreeToken

核心功能

避坑指南

适用场景

标签

快速上手

相似推荐