3dc35694b3
Signed-off-by: NumberWan <wantszkin2003@gmail.com> Signed-off-by: SamitHuang <285365963@qq.com> Signed-off-by: Samit <285365963@qq.com> Co-authored-by: Isotr0py <Isotr0py@outlook.com> Co-authored-by: Tianyu Guo <guoty9@mail2.sysu.edu.cn> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: SamitHuang <285365963@qq.com> Co-authored-by: Gao Han <hgaoaf@connect.ust.hk>
240 lines
12 KiB
Markdown
240 lines
12 KiB
Markdown
# FP8 Quantization
|
||
|
||
## Overview
|
||
|
||
FP8 quantization converts BF16/FP16 weights to FP8 at model load time. Online
|
||
activation scaling is the default and does not require calibration. Static
|
||
activation scaling is supported when calibrated scale information is available.
|
||
For ModelOpt-produced pre-quantized checkpoints, see
|
||
[ModelOpt Quantization](modelopt.md).
|
||
|
||
Some architectures can quantize all linear layers. Others have
|
||
quality-sensitive layers that should stay in BF16 through `ignored_layers`.
|
||
Image-stream MLPs (`img_mlp`) are a common sensitive target because denoising
|
||
latent ranges shift across timesteps and small per-layer errors can compound
|
||
in deep DiT blocks.
|
||
|
||
## Hardware Support
|
||
|
||
| Device | Support |
|
||
| -------- | --------- |
|
||
| NVIDIA Blackwell GPU (SM 100+) | ✅ |
|
||
| NVIDIA Ada/Hopper GPU (SM 89+) | ✅ |
|
||
| NVIDIA Ampere GPU (SM 80+) | ✅ |
|
||
| AMD ROCm | ⭕ |
|
||
| Intel XPU | ⭕ |
|
||
| Ascend NPU | ❌ |
|
||
|
||
Legend: `✅` supported, `❌` unsupported, `⭕` not verified in this
|
||
guide. FP8 on Ampere may use a weight-only path where available.
|
||
|
||
### Faster FP8 GEMM on Blackwell (quack)
|
||
|
||
On Blackwell (SM 100+), vLLM runs FP8 linears through the FlashInfer kernel, which
|
||
applies the bias as a separate kernel after the GEMM. On the small GEMMs in video
|
||
DiTs this bias add is a significant overhead. Installing the optional `quack` kernel
|
||
lets vLLM-Omni fuse `alpha * (A @ B) + bias` into a single CuteDSL GEMM, recovering
|
||
that overhead (e.g. HunyuanVideo-1.5 FP8 goes from slower-than-BF16 to faster).
|
||
|
||
```bash
|
||
# CUDA 12.9
|
||
pip install vllm-omni[quack]
|
||
|
||
# CUDA 13.x
|
||
pip install 'quack-kernels[cu13]' --extra-index-url https://download.pytorch.org/whl/cu130
|
||
```
|
||
|
||
It is enabled automatically once installed (no flag needed) and is **datacenter
|
||
Blackwell only** (`sm_100` / `sm_101` / `sm_103`, compute capability `10.x`, e.g.
|
||
B200): quack's CuteDSL GEMM uses the 5th-gen `tcgen05` tensor-core MMA, which exists
|
||
only on those parts. On Hopper/Ada the CUTLASS FP8 kernel already fuses bias, and on
|
||
workstation/consumer Blackwell (`sm_120` / `sm_121`, compute capability `12.x`, e.g.
|
||
RTX PRO 6000 / RTX 50-series) `tcgen05` is absent — so quack is **not** auto-enabled
|
||
there and FlashInfer's native FP8 path is used instead. Set
|
||
`VLLM_OMNI_USE_QUACK_FP8=1` to force quack on, or `VLLM_OMNI_USE_QUACK_FP8=0` to force
|
||
the FlashInfer path. If `quack-kernels` is not installed, FP8 still works — it just
|
||
keeps the unfused FlashInfer path.
|
||
|
||
#### Compile cache and warmup
|
||
|
||
quack JIT-compiles its kernel once per distinct GEMM shape (tens of seconds, longer
|
||
the first time across all autotuned configs). The compiled `.o` files are cached on
|
||
disk and reused on later runs, so this is a one-time cost — **not** per request.
|
||
|
||
vLLM-Omni points that cache at `~/.cache/vllm_omni/quack` (override with
|
||
`QUACK_CACHE_DIR`) instead of quack's default under `/tmp`, so it survives restarts.
|
||
In containers, set `QUACK_CACHE_DIR` to a mounted/persistent path — or bake it into
|
||
the image — so the first cold start does not recompile. The engine's startup dummy
|
||
run exercises the kernels, but new shapes, layouts, dtypes, or bias settings may
|
||
still need compilation or tuning. The warmup helper uses inference mode and
|
||
transposed weights without bias. Daemon workers compile candidates in-process
|
||
while retaining autotuning and caching.
|
||
|
||
Scale validation and Quack/FlashInfer dispatch run inside a PyTorch custom op.
|
||
This keeps layer-specific scale addresses and validation-cache updates out of
|
||
Dynamo tracing without introducing a graph break. Unpopulated scales and Quack
|
||
failures still fall back to FlashInfer at runtime. CUDA graph capture additionally
|
||
requires warming the dispatch with populated scales; Python validation and
|
||
dispatch do not rerun during CUDA graph replay.
|
||
|
||
To pre-warm specific shapes (e.g. at image build time):
|
||
|
||
```python
|
||
from vllm_omni.quantization.quack_fp8 import warmup_quack_fp8
|
||
# (M, K, N) per linear; M = number of tokens for your resolution/frame count
|
||
warmup_quack_fp8([(14040, 2048, 6144), (14040, 2048, 2048)])
|
||
```
|
||
|
||
> The PyPI package is `quack-kernels` (imported as `quack`); plain `pip install
|
||
> quack` is an unrelated statistics library. Requires CUDA 12.9+ and Python 3.12.
|
||
|
||
## Model Type Support
|
||
|
||
### Diffusion Models
|
||
|
||
| Model | HF models | Online | Pre-calibrated | Recommendation | `ignored_layers` | Text-Encoder quantization |
|
||
| ------- | ----------- | :-------: | :------: | ---------------- | ------------------ | ------------------ |
|
||
| Qwen-Image | `Qwen/Qwen-Image`, `Qwen/Qwen-Image-2512` | Yes | Yes | Skip sensitive image-stream MLPs when quality regresses | `img_mlp` | |
|
||
| Qwen-Image-2.1 | `Qwen/Qwen-Image-2.1` | Yes | Not validated | Skip image-stream MLPs for best fidelity; all-layer FP8 is loadable. On H200 (1024×1024, teapot prompt) DiT FP8-all is ~30.4 GB / 6.06 s vs BF16 ~36.9 GB / 7.34 s (~33 dB vs BF16). Text encoder FP8 covers the language model only — the vision tower and LM head stay BF16 | `img_mlp` | ✅︎ (language model only) |
|
||
| Wan2.2 | Wan2.2 diffusion pipelines | Not validated | Not validated | Validate against BF16 before documenting as supported | TBD | |
|
||
| LTX-2 | `Lightricks/LTX-2`, `rootonchair/LTX-2-19b-distilled` | Yes | Not validated | Transformer only; use dynamic phase LoRA for ordinary two-stage | None | |
|
||
| LTX-2.3 | `diffusers/LTX-2.3-Diffusers`, `diffusers/LTX-2.3-Distilled-Diffusers` | Yes | Not validated | Transformer only; use dynamic phase LoRA for ordinary two-stage | None | |
|
||
| Z-Image | `Tongyi-MAI/Z-Image-Turbo` | Yes | Yes | All layers | None | ✅︎ |
|
||
| FLUX.1 | `black-forest-labs/FLUX.1-dev`, `black-forest-labs/FLUX.1-schnell` | Yes | Yes | All layers | None | |
|
||
| FLUX.2-dev | `black-forest-labs/FLUX.2-dev` | Yes | Not validated | All layers | None | ✅︎ |
|
||
| FLUX.2-klein | `black-forest-labs/FLUX.2-klein-4B` | Yes | Yes | All layers | None | |
|
||
| HunyuanImage-3.0 | `tencent/HunyuanImage-3.0`, `tencent/HunyuanImage-3.0-Instruct` | Yes | Yes | All layers; use the Hunyuan deploy config for multi-stage runs | None | |
|
||
| HunyuanVideo-1.5 | `hunyuanvideo-community/HunyuanVideo-1.5-Diffusers-480p_t2v`, `720p_t2v`, `480p_i2v` | Yes | Yes | All layers | None | |
|
||
| Cosmos3 | `nvidia/Cosmos3-Nano`, `nvidia/Cosmos3-Super` | Yes | Not validated | All layers | None | |
|
||
| MiniMax-H3 | `MiniMaxAI/MiniMax-H3` (`FL2VA` / `Ref2VA`) | Yes | Not validated | `quantization="fp8"` quantizes eligible DiT and text-encoder linears; mixed-precision input/output heads stay FP32 | None | ✅︎ |
|
||
| SenseNova-U1.5 | `sensenova/SenseNova-U1.5-8B-MoT` | Yes | Not validated | UND/GEN language-model linears only | None | |
|
||
|
||
### Multi-Stage Omni/TTS Model (Qwen3-Omni, Qwen3-TTS)
|
||
|
||
| Model | Scope | Format | Status |
|
||
| ---------- | ---------------------------- | ---------------------------------------- | ----------------------------------- |
|
||
| Qwen3-Omni | Thinker language-model stage | [ModelOpt](modelopt.md) `quant_algo=FP8` | Tested for thinker memory reduction |
|
||
| Qwen3-TTS | TTS language-model stage | Checkpoint config | Not validated |
|
||
|
||
Audio encoder, vision encoder, talker, and code2wav stay in BF16 unless a
|
||
model-specific guide says otherwise.
|
||
|
||
### Multi-Stage Diffusion Model (BAGEL, GLM-Image)
|
||
|
||
| Model | Scope | Status | Notes |
|
||
| --------- | ---------------------------------------- | ------------- | ----------------------------------------------- |
|
||
| BAGEL | Stage-specific transformer or DiT module | Not validated | Route FP8 to the intended stage before enabling |
|
||
| GLM-Image | Stage-specific transformer or DiT module | Not validated | Validate quality against BF16 baseline |
|
||
|
||
## Configuration
|
||
|
||
Python API:
|
||
|
||
```python
|
||
from vllm_omni import Omni
|
||
from vllm_omni.inputs.data import OmniDiffusionSamplingParams
|
||
|
||
omni = Omni(model="<your-model>", quantization="fp8")
|
||
|
||
omni_with_skips = Omni(
|
||
model="<your-model>",
|
||
quantization_config={
|
||
"method": "fp8",
|
||
"ignored_layers": ["img_mlp"],
|
||
},
|
||
)
|
||
|
||
outputs = omni.generate(
|
||
"A cat sitting on a windowsill",
|
||
OmniDiffusionSamplingParams(num_inference_steps=50),
|
||
)
|
||
```
|
||
|
||
### Global and per-component scope
|
||
|
||
A plain, global quantization configuration is passed unchanged to every
|
||
quantization-aware component constructed by a pipeline. This includes an
|
||
eligible encoder when that encoder is implemented with vLLM quantizable
|
||
layers; it does not rewrite arbitrary `torch.nn` modules. A structured
|
||
component map is the only way to narrow that scope. Pipeline integrations that
|
||
do not yet expose an encoder through the quantization factory remain DiT-only.
|
||
|
||
For a pipeline that exposes both a transformer and a quantization-aware text
|
||
encoder, the scope is:
|
||
|
||
| Configuration | Transformer | Text encoder | Components without supported quantizable layers |
|
||
| --------------- | ------------- | -------------- | ------------------------------------------------- |
|
||
| `quantization="fp8"` | FP8 | FP8 | checkpoint precision |
|
||
| `{"transformer": {"method": "fp8"}}` | FP8 | checkpoint precision | checkpoint precision |
|
||
| `{"text_encoder": {"method": "fp8"}}` | checkpoint precision | FP8 | checkpoint precision |
|
||
|
||
Use `quantization_config` for component-selective Python configuration. For
|
||
example, quantize only the text decoder:
|
||
|
||
```python
|
||
omni = Omni(
|
||
model="<your-model>",
|
||
quantization_config={"text_encoder": {"method": "fp8"}},
|
||
)
|
||
```
|
||
|
||
Component keys are runtime prefixes exposed by the pipeline integration; common
|
||
keys include `transformer` and `text_encoder`. Entries may be combined, and
|
||
`ignored_layers` can keep named eligible layers in checkpoint precision. Check
|
||
the model recipe for supported components and their runtime prefixes.
|
||
|
||
CLI:
|
||
|
||
```bash
|
||
python text_to_image.py --model <your-model> --quantization fp8
|
||
python text_to_image.py --model <your-model> --quantization fp8 --ignored-layers "img_mlp"
|
||
vllm serve <your-model> --omni --quantization fp8
|
||
```
|
||
|
||
### Distributed layerwise offload
|
||
|
||
Per-tensor online FP8 linears can be combined with DLO's default AllGather path:
|
||
|
||
```bash
|
||
vllm serve <your-model> \
|
||
--omni \
|
||
--quantization fp8 \
|
||
--enable-distributed-layerwise-offload
|
||
```
|
||
|
||
Direct checkpoint mmap is not used for this combination. Each rank runs the
|
||
ordinary loader to create the final FP8 weights and scales before DLO shards
|
||
them, so startup temporarily materializes the complete FP8 model in host memory
|
||
on every rank. Runtime DLO host residency is sharded normally. Use
|
||
`--dlo-no-use-allgather` when the synchronized AllGather request-wave contract
|
||
does not fit the workload.
|
||
|
||
## Parameters
|
||
|
||
| Parameter | Type | Default | Description |
|
||
| ----------- | ------ | --------- | ------------- |
|
||
| `method` | str | - | Quantization method (`"fp8"`) |
|
||
| `ignored_layers` | list[str] | `[]` | Layer name patterns to keep in BF16 |
|
||
| `activation_scheme` | str | `"dynamic"` | `"dynamic"` selects online activation scaling, or `"static"` when scales are available |
|
||
| `weight_block_size` | list[int] \| None | `None` | Block size for block-wise weight quantization |
|
||
|
||
The available `ignored_layers` names depend on the model architecture, for
|
||
example `to_qkv`, `to_out`, `img_mlp`, or `txt_mlp`.
|
||
|
||
## Validation and Notes
|
||
|
||
FP8 quantization can be combined with cache acceleration:
|
||
|
||
```python
|
||
omni = Omni(
|
||
model="<your-model>",
|
||
quantization="fp8",
|
||
cache_backend="tea_cache",
|
||
cache_config={"rel_l1_thresh": 0.2},
|
||
)
|
||
```
|
||
|
||
Compare generated outputs with a BF16 baseline before adding a new model to the
|
||
supported table. GLM-Image and Helios are not listed as FP8-supported diffusion
|
||
models until they have method-specific validation.
|