Q: How to fix bf16 precision loss in Qwen3.5/Qwen3.6 RMSNorm weights when serving with FreeToken?

The bug: FreeToken's `models/qwen3_5_moe` loader folds the Gemma-style `+1` into the bf16 checkpoint tensor at load time (`tensor + 1.0`), storing `bf16(1 + w)`. Transformers instead keeps zero-centred weights and computes `(1 + w)` in fp32, so the fold rounds away small weights — on `nvidia/Qwen3.6-35B-A3B-NVFP4`, 77.4% of the 171,008 norm multipliers change, with mean relative error 1.4e-3 and max 3.9e-3 (half a bf16 ulp), affecting `input_layernorm`, `post_attention_layernorm`, `self_attn.q_norm`, `self_attn.k_norm` and the final `model.norm` on every token. Fix: store the raw zero-centred bf16 weight `w` (do NOT add 1 in the loader) and apply the `+1` in fp32 at runtime — use the existing `GemmaPlusOneRMSNorm` layer (`layers/norm.py`) or call `sgl_kernel.gemma_rmsnorm(x, w)` instead of `sgl_kernel.rmsnorm(x, bf16(1+w))`. This matches the approach already used by the `qwen4_exp` loader and restores 99.999% bf16 agreement with the fp64 reference (residual error is just output rounding, rel L2 ~1.7e-3 vs 3.0e-3 before).

Question 1

Why does FreeToken reject 'file://localhost' image URIs with the --allowed-local-media-path allowlist?

Accepted Answer

This is a URI parsing bug in `_load_local_media` (`python/freetoken/mm/media.py`). The function builds the file path by concatenating the URI's `netloc` and `path`, so `file://localhost/path/image.png` becomes the relative path `localhost/path/image.png` instead of `/path/image.png`, which then fails the resolved-path allowlist check and raises `ValueError`. Fix: treat empty and `localhost` authorities (case-insensitive) as local per RFC 8089, reject any other non-local authority, decode the path with `url2pathname`, and keep the existing `Path.resolve()` strict-root allowlist check. Workaround until patched: use the canonical `file:///path/image.png` form (empty authority) instead of `file://localhost/...`.

Question 2

Why does FreeToken OOM or lose 1.3-3.3 GiB of VRAM on RTX 5090 / 5070 Ti (sm_120) after the first TMA kernel launch?

Accepted Answer

On sm_120 GPUs, Triton versions below 3.8 compile TMA loads (`TensorDescriptor`) into a driver syscall that raises the per-thread stack limit from 1024 B to ~14.4 KB for every resident thread on the first kernel launch. This permanently reserves about 1.34 GiB on a 70-SM RTX 5070 Ti and about 3.3 GiB on a 170-SM RTX 5090. Since the reservation happens after the KV pool and expert cache are sized from free VRAM, the server OOMs. Fix: upgrade to torch 2.14 + Triton 3.8 or newer. Triton 3.8 emits the native `cp.async.bulk.tensor.2d.shared::cta` instruction on sm_120, which requires no syscall and causes no stack reservation. The engine already requires these versions; if you see the issue, check for a stale venv or source build with old pinned dependencies and reinstall with the updated requirements.

Question 3

How to fix 'expandable_segments: memory mapping failed with OOM' CUDACachingAllocator warnings during chunked prefill?

Accepted Answer

Lower `--max-prefill-length` (e.g. back to the default 8192). This flag sets the chunked-prefill chunk size, NOT the maximum prompt length the server can handle — a 30k-token request is still served fine at 8192, just split into more chunks. The warnings occur because large chunks create a big transient activation peak in the DSA/MLA attention path (several live [T, 64, 512] bf16 tensors, ~3.8 GiB at a 12288-token chunk), causing torch's expandable-segments allocator to retry memory mappings. Each retry burst costs a device-wide sync stall (1.5-5 s) but self-resolves, so requests still complete with no correctness impact. Note: with MoE expert caching too small for overlap mode, larger chunks can still be faster overall since the per-chunk full-expert PCIe copy is amortized over more tokens — so there is a trade-off between stall-free logs and peak prefill throughput.

Question 4

How to fix sglang-kernel load errors and 'C++20 or later compatible compiler is required' when upgrading to torch 2.14.1?

Accepted Answer

sglang-kernel links libtorch C++ symbols and loads only against the torch it was built for, so versions must be pinned together: use `sglang-kernel==0.4.9` with `torch>=2.14,<2.15` (0.4.5 is the torch 2.11 build). torch 2.14 requires Triton 3.8, so set `triton>=3.6,<3.9` and `torchvision>=0.29,<0.30`. The PyPI wheel of torch 2.14.1 is the cu130 build and requires `cuda-toolkit==13.0.3`, so keep your installer/uv index on cu130. If you compile C++ extensions, change `-std=c++17` to `-std=c++20` in setup.py, otherwise torch 2.14 headers fail with `#error C++20 or later compatible compiler is required`. Note: Triton 3.8 regresses some kernels (e.g. `_decode_nvfp4_marlin_kernel` for NVFP4 MoE decode is 40-130% slower); eight regressions were fixed in follow-up PRs and the rest have under ~4% end-to-end impact.

Question 5

How to fix bf16 precision loss in Qwen3.5/Qwen3.6 RMSNorm weights when serving with FreeToken?

Accepted Answer

The bug: FreeToken's `models/qwen3_5_moe` loader folds the Gemma-style `+1` into the bf16 checkpoint tensor at load time (`tensor + 1.0`), storing `bf16(1 + w)`. Transformers instead keeps zero-centred weights and computes `(1 + w)` in fp32, so the fold rounds away small weights — on `nvidia/Qwen3.6-35B-A3B-NVFP4`, 77.4% of the 171,008 norm multipliers change, with mean relative error 1.4e-3 and max 3.9e-3 (half a bf16 ulp), affecting `input_layernorm`, `post_attention_layernorm`, `self_attn.q_norm`, `self_attn.k_norm` and the final `model.norm` on every token. Fix: store the raw zero-centred bf16 weight `w` (do NOT add 1 in the loader) and apply the `+1` in fp32 at runtime — use the existing `GemmaPlusOneRMSNorm` layer (`layers/norm.py`) or call `sgl_kernel.gemma_rmsnorm(x, w)` instead of `sgl_kernel.rmsnorm(x, bf16(1+w))`. This matches the approach already used by the `qwen4_exp` loader and restores 99.999% bf16 agreement with the fp64 reference (residual error is just output rounding, rel L2 ~1.7e-3 vs 3.0e-3 before).

FreeToken

Core Features

What It Can't Do

Use Cases

Tags

Getting Started

Similar Projects