Q: FreeTokenでサービングする際のQwen3.5/Qwen3.6 RMSNorm重みにおけるbf16の精度損失を修正する方法

バグ: FreeTokenの`models/qwen3_5_moe`ローダーは、読み込み時にGemma式の`+1`をbf16チェックポイントテンソルに畳み込み(`tensor + 1.0`)、`bf16(1 + w)`として保存します。一方、Transformersはゼロ中心化された重みをそのまま保持し、`(1 + w)`をfp32で計算するため、この畳み込みによって小さな重みが丸められてしまいます — `nvidia/Qwen3.6-35B-A3B-NVFP4`では、171,008個のノルム乗数のうち77.4%が変化し、平均相対誤差は1.4e-3、最大は3.9e-3(bf16 ulpの半分)に達し、すべてのトークンで`input_layernorm`、`post_attention_layernorm`、`self_attn.q_norm`、`self_attn.k_norm`、および最終の`model.norm`に影響します。修正方法: 生のゼロ中心化されたbf16重み`w`をそのまま保存し(ローダーで1を加算しない)、実行時に`+1`をfp32で適用します — 既存の`GemmaPlusOneRMSNorm`レイヤー(`layers/norm.py`)を使用するか、`sgl_kernel.rmsnorm(x, bf16(1+w))`の代わりに`sgl_kernel.gemma_rmsnorm(x, w)`を呼び出してください。これは`qwen4_exp`ローダーですでに採用されている手法と一致しており、fp64リファレンスとのbf16一致率を99.999%に回復します(残余誤差は出力の丸めのみで、rel L2は修正前の3.0e-3に対し約1.7e-3です)。

Question 1

FreeToken が --allowed-local-media-path 許可リストで 'file://localhost' の画像 URI を拒否するのはなぜですか?

Accepted Answer

これは `_load_local_media`（`python/freetoken/mm/media.py`）における URI 解析のバグです。この関数は URI の `netloc` と `path` を連結してファイルパスを構築するため、`file://localhost/path/image.png` が `/path/image.png` ではなく相対パス `localhost/path/image.png` になり、その後、解決済みパスの許可リストチェックに失敗して `ValueError` が発生します。修正方法: RFC 8089 に従い、空のオーソリティと `localhost` のオーソリティ（大文字小文字を区別しない）をローカルとして扱い、それ以外の非ローカルなオーソリティはすべて拒否し、パスを `url2pathname` でデコードし、既存の `Path.resolve()` による厳密なルート許可リストチェックは維持します。パッチが適用されるまでの回避策: `file://localhost/...` の代わりに、正規形式の `file:///path/image.png`（空のオーソリティ）を使用してください。

Question 2

なぜ FreeToken は RTX 5090 / 5070 Ti (sm_120) で最初の TMA カーネル起動後に OOM になったり、1.3-3.3 GiB の VRAM を失ったりするのか？

Accepted Answer

sm_120 GPUでは、3.8未満のTritonバージョンは、TMAロード（`TensorDescriptor`）をドライバのシステムコールにコンパイルし、最初のカーネル起動時にすべての常駐スレッドのスレッドあたりスタック上限を1024 Bから約14.4 KBに引き上げます。これにより、70-SMのRTX 5070 Tiでは約1.34 GiB、170-SMのRTX 5090では約3.3 GiBが永続的に予約されます。この予約はKVプールとエキスパートキャッシュが空き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 トークンのリクエストは 8192 でも問題なく処理され、単により多くのチャンクに分割されるだけです。警告が発生するのは、大きなチャンクが DSA/MLA アテンションパスで大きな一時的なアクティベーションピークを生じさせる（複数の生存中の [T, 64, 512] bf16 テンソル、12288 トークンのチャンクで約 3.8 GiB）ため、torch の expandable-segments アロケーターがメモリマッピングを再試行するからです。各再試行バーストはデバイス全体の同期ストール（1.5〜5 秒）を引き起こしますが、自然に解消されるため、リクエストは正確性への影響なく完了します。注意: overlap モードに対して MoE エキスパートキャッシュが小さすぎる場合、チャンクごとのフルエキスパート PCIe コピーがより多くのトークンに分散して償却されるため、大きなチャンクの方が全体的に高速になることがあります。そのため、ストールのないログとピーク 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ホイールはcu130ビルドであり、`cuda-toolkit==13.0.3`を必要とするため、インストーラー/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では一部のカーネルに性能回帰があります（例：NVFP4 MoEデコード用の`_decode_nvfp4_marlin_kernel`は40-130%低速化）。8件の回帰はフォローアップ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はゼロ中心化された重みをそのまま保持し、`(1 + w)`をfp32で計算するため、この畳み込みによって小さな重みが丸められてしまいます — `nvidia/Qwen3.6-35B-A3B-NVFP4`では、171,008個のノルム乗数のうち77.4%が変化し、平均相対誤差は1.4e-3、最大は3.9e-3(bf16 ulpの半分)に達し、すべてのトークンで`input_layernorm`、`post_attention_layernorm`、`self_attn.q_norm`、`self_attn.k_norm`、および最終の`model.norm`に影響します。修正方法: 生のゼロ中心化されたbf16重み`w`をそのまま保存し(ローダーで1を加算しない)、実行時に`+1`をfp32で適用します — 既存の`GemmaPlusOneRMSNorm`レイヤー(`layers/norm.py`)を使用するか、`sgl_kernel.rmsnorm(x, bf16(1+w))`の代わりに`sgl_kernel.gemma_rmsnorm(x, w)`を呼び出してください。これは`qwen4_exp`ローダーですでに採用されている手法と一致しており、fp64リファレンスとのbf16一致率を99.999%に回復します(残余誤差は出力の丸めのみで、rel L2は修正前の3.0e-3に対し約1.7e-3です)。

FreeToken

主な機能

できないこと

使用例

タグ

はじめ方

類似プロジェクト