Spaces:
Running on Zero
Running on Zero
Initial ABot-World interactive rollout demo
Browse filesThis view is limited to 50 files because it contains too many changes. See raw diff
- .gitattributes +4 -0
- README.md +21 -7
- app.py +328 -0
- configs/default_config.yaml +6 -0
- configs/long_forcing_dmd.yaml +36 -0
- examples/desert_valley.png +3 -0
- examples/example.png +3 -0
- examples/forest_stream.png +3 -0
- examples/mountain_meadow.png +3 -0
- pipeline/__init__.py +5 -0
- pipeline/causal_inference.py +822 -0
- quantizer/__init__.py +49 -0
- quantizer/kernels/__init__.py +13 -0
- quantizer/kernels/python/__init__.py +13 -0
- quantizer/kernels/python/gemm/__init__.py +34 -0
- quantizer/kernels/python/gemm/fp8_gemm.py +112 -0
- quantizer/kernels/python/gemm/fp8_gemm_torch.py +217 -0
- quantizer/kernels/python/mm/__init__.py +0 -0
- quantizer/kernels/python/mm/triton_kernels.py +743 -0
- quantizer/kernels/python/quantizers/__init__.py +46 -0
- quantizer/kernels/python/quantizers/fp8_per_block.py +78 -0
- quantizer/kernels/python/quantizers/fp8_per_block_torch.py +155 -0
- quantizer/kernels/python/quantizers/fp8_per_token_group.py +170 -0
- quantizer/kernels/python/quantizers/fp8_per_token_group_torch.py +163 -0
- quantizer/kernels/python/sgl/__init__.py +0 -0
- quantizer/kernels/python/sgl/int8_kernel.py +465 -0
- quantizer/quant/__init__.py +19 -0
- quantizer/quant/modules/__init__.py +17 -0
- quantizer/quant/modules/linear.py +589 -0
- quantizer/quant/ptq.py +272 -0
- quantizer/quant/quant_func.py +520 -0
- quantizer/quant/utils/__init__.py +37 -0
- quantizer/quant/utils/quant_io.py +149 -0
- quantizer/quant/utils/utils.py +219 -0
- requirements.txt +16 -0
- utils/cuda_kernel_buckets.py +227 -0
- utils/memory.py +135 -0
- utils/misc.py +39 -0
- utils/scheduler.py +194 -0
- utils/wan_wrapper.py +871 -0
- wan/__init__.py +2 -0
- wan/modules/__init__.py +17 -0
- wan/modules/attention.py +476 -0
- wan/modules/causal_model.py +0 -0
- wan/modules/helios_kernels/__init__.py +5 -0
- wan/modules/helios_kernels/fp32_rmsnorm.py +48 -0
- wan/modules/helios_kernels/tiled_linear.py +399 -0
- wan/modules/helios_kernels/triton_norm.py +413 -0
- wan/modules/helios_kernels/triton_rope.py +708 -0
- wan/modules/helios_kernels/utils.py +70 -0
.gitattributes
CHANGED
|
@@ -33,3 +33,7 @@ saved_model/**/* filter=lfs diff=lfs merge=lfs -text
|
|
| 33 |
*.zip filter=lfs diff=lfs merge=lfs -text
|
| 34 |
*.zst filter=lfs diff=lfs merge=lfs -text
|
| 35 |
*tfevents* filter=lfs diff=lfs merge=lfs -text
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 33 |
*.zip filter=lfs diff=lfs merge=lfs -text
|
| 34 |
*.zst filter=lfs diff=lfs merge=lfs -text
|
| 35 |
*tfevents* filter=lfs diff=lfs merge=lfs -text
|
| 36 |
+
examples/desert_valley.png filter=lfs diff=lfs merge=lfs -text
|
| 37 |
+
examples/example.png filter=lfs diff=lfs merge=lfs -text
|
| 38 |
+
examples/forest_stream.png filter=lfs diff=lfs merge=lfs -text
|
| 39 |
+
examples/mountain_meadow.png filter=lfs diff=lfs merge=lfs -text
|
README.md
CHANGED
|
@@ -1,13 +1,27 @@
|
|
| 1 |
---
|
| 2 |
-
title:
|
| 3 |
-
emoji:
|
| 4 |
-
colorFrom:
|
| 5 |
-
colorTo:
|
| 6 |
sdk: gradio
|
| 7 |
-
sdk_version: 6.20.0
|
| 8 |
-
python_version: '3.12'
|
| 9 |
app_file: app.py
|
| 10 |
pinned: false
|
|
|
|
|
|
|
|
|
|
| 11 |
---
|
| 12 |
|
| 13 |
-
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
---
|
| 2 |
+
title: ABot-World Interactive
|
| 3 |
+
emoji: 🌍
|
| 4 |
+
colorFrom: gray
|
| 5 |
+
colorTo: red
|
| 6 |
sdk: gradio
|
|
|
|
|
|
|
| 7 |
app_file: app.py
|
| 8 |
pinned: false
|
| 9 |
+
short_description: Interactive action-conditioned world rollout from one image
|
| 10 |
+
python_version: "3.10"
|
| 11 |
+
startup_duration_timeout: 40m
|
| 12 |
---
|
| 13 |
|
| 14 |
+
# 🌍 ABot-World — Interactive World Rollout
|
| 15 |
+
|
| 16 |
+
Turn a single image into a navigable world. Provide a first-frame image, describe
|
| 17 |
+
the scene, and choose an action (WASD to move / IJKL to look). The model
|
| 18 |
+
autoregressively rolls out an action-conditioned video of the world.
|
| 19 |
+
|
| 20 |
+
- **Model:** [`acvlab/ABot-World-0-5B-LF`](https://huggingface.co/acvlab/ABot-World-0-5B-LF) (built on Wan2.2-TI2V-5B)
|
| 21 |
+
- **Code:** [amap-cvlab/ABot-World](https://github.com/amap-cvlab/ABot-World)
|
| 22 |
+
- **Project:** [ABot-World](https://amap-cvlab.github.io/ABot-World/)
|
| 23 |
+
|
| 24 |
+
This Space adapts the upstream real-time interactive web client into a request/response
|
| 25 |
+
rollout suitable for ZeroGPU: instead of live keyboard streaming, you pick one action
|
| 26 |
+
to hold and a rollout length, and receive a video. Inference is causal / block-wise,
|
| 27 |
+
matching the upstream `CausalInferencePipeline`.
|
app.py
ADDED
|
@@ -0,0 +1,328 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
"""
|
| 2 |
+
ABot-World — Interactive Action-Conditioned World Rollout (ZeroGPU Gradio demo)
|
| 3 |
+
|
| 4 |
+
Given a first-frame image, a scene prompt, and an action (WASD / IJKL), the model
|
| 5 |
+
autoregressively rolls out an action-conditioned video of a navigable world.
|
| 6 |
+
|
| 7 |
+
Upstream: https://github.com/amap-cvlab/ABot-World
|
| 8 |
+
Model: https://huggingface.co/acvlab/ABot-World-0-5B-LF (built on Wan2.2-TI2V-5B)
|
| 9 |
+
"""
|
| 10 |
+
import os
|
| 11 |
+
os.environ.setdefault("PYTORCH_CUDA_ALLOC_CONF", "expandable_segments:True")
|
| 12 |
+
|
| 13 |
+
import spaces # must precede torch / CUDA-touching imports
|
| 14 |
+
|
| 15 |
+
import sys
|
| 16 |
+
import time
|
| 17 |
+
import tempfile
|
| 18 |
+
import threading
|
| 19 |
+
from pathlib import Path
|
| 20 |
+
|
| 21 |
+
import numpy as np
|
| 22 |
+
import torch
|
| 23 |
+
import imageio
|
| 24 |
+
import gradio as gr
|
| 25 |
+
from PIL import Image
|
| 26 |
+
from omegaconf import OmegaConf
|
| 27 |
+
from huggingface_hub import snapshot_download
|
| 28 |
+
|
| 29 |
+
# ── Repo paths ───────────────────────────────────────────────────────────────
|
| 30 |
+
APP_DIR = Path(__file__).resolve().parent
|
| 31 |
+
if str(APP_DIR) not in sys.path:
|
| 32 |
+
sys.path.insert(0, str(APP_DIR))
|
| 33 |
+
|
| 34 |
+
MODEL_ID = "acvlab/ABot-World-0-5B-LF"
|
| 35 |
+
CKPT_DIR = APP_DIR / "checkpoints" / "ABot-World-0-5B-LF"
|
| 36 |
+
|
| 37 |
+
# ── Stream / rollout configuration ───────────────────────────────────────────
|
| 38 |
+
# 704x1280 is the native training resolution used by the upstream web client.
|
| 39 |
+
# We keep it identical so the model behaves as intended.
|
| 40 |
+
STREAM_HEIGHT = 704
|
| 41 |
+
STREAM_WIDTH = 1280
|
| 42 |
+
VIDEO_FPS = 12
|
| 43 |
+
|
| 44 |
+
# Actions map to the 8-key one-hot the model was trained on (W A S D I J K L).
|
| 45 |
+
ACTIONS = {
|
| 46 |
+
"Forward (W)": {"W": True},
|
| 47 |
+
"Backward (S)": {"S": True},
|
| 48 |
+
"Turn left (A)": {"A": True},
|
| 49 |
+
"Turn right (D)": {"D": True},
|
| 50 |
+
"Look up (I)": {"I": True},
|
| 51 |
+
"Look down (K)": {"K": True},
|
| 52 |
+
"Pan left (J)": {"J": True},
|
| 53 |
+
"Pan right (L)": {"L": True},
|
| 54 |
+
"Forward + turn left": {"W": True, "A": True},
|
| 55 |
+
"Forward + turn right": {"W": True, "D": True},
|
| 56 |
+
"Idle (no action)": {},
|
| 57 |
+
}
|
| 58 |
+
|
| 59 |
+
DEFAULT_PROMPT = (
|
| 60 |
+
"A realistic outdoor world scene with a navigable path, natural lighting, "
|
| 61 |
+
"detailed ground texture, and stable forward motion."
|
| 62 |
+
)
|
| 63 |
+
|
| 64 |
+
# ── Download weights (once, at startup) ──────────────────────────────────────
|
| 65 |
+
print(f"[startup] downloading {MODEL_ID} weights ...", flush=True)
|
| 66 |
+
snapshot_download(
|
| 67 |
+
repo_id=MODEL_ID,
|
| 68 |
+
repo_type="model",
|
| 69 |
+
local_dir=str(CKPT_DIR),
|
| 70 |
+
)
|
| 71 |
+
print("[startup] weights downloaded.", flush=True)
|
| 72 |
+
|
| 73 |
+
# Point the taehv VAE at the downloaded checkpoints via env (config uses these).
|
| 74 |
+
os.environ["TAEW2_2_CHECKPOINT"] = str(CKPT_DIR / "taew2_2.pth")
|
| 75 |
+
|
| 76 |
+
# ── Build the causal streaming pipeline (module scope, eager to CUDA) ─────────
|
| 77 |
+
from pipeline import CausalInferencePipeline
|
| 78 |
+
from utils.misc import set_seed
|
| 79 |
+
from utils.wan_wrapper import create_vae_from_config
|
| 80 |
+
|
| 81 |
+
_CONFIG_PATH = APP_DIR / "configs" / "long_forcing_dmd.yaml"
|
| 82 |
+
_DEFAULT_CFG_PATH = APP_DIR / "configs" / "default_config.yaml"
|
| 83 |
+
|
| 84 |
+
|
| 85 |
+
def _build_config():
|
| 86 |
+
cfg = OmegaConf.merge(
|
| 87 |
+
OmegaConf.load(str(_DEFAULT_CFG_PATH)),
|
| 88 |
+
OmegaConf.load(str(_CONFIG_PATH)),
|
| 89 |
+
)
|
| 90 |
+
# Rewrite every checkpoint path to the absolute downloaded location.
|
| 91 |
+
cfg.taew2_2_checkpoint = str(CKPT_DIR / "taew2_2.pth")
|
| 92 |
+
cfg.lightvae_encoder_checkpoint = str(CKPT_DIR / "Wan2.2_VAE.pth")
|
| 93 |
+
cfg.model_kwargs.model_name = str(CKPT_DIR)
|
| 94 |
+
cfg.text_encoder_kwargs.tokenizer_path = str(CKPT_DIR / "google" / "umt5-xxl") + "/"
|
| 95 |
+
cfg.text_encoder_kwargs.encoder_pth_path = str(CKPT_DIR / "models_t5_umt5-xxl-enc-bf16.pth")
|
| 96 |
+
cfg.vae_kwargs.pretrained_path = str(CKPT_DIR / "Wan2.2_VAE.pth")
|
| 97 |
+
cfg.vae_type = "taew2_2"
|
| 98 |
+
# FP8 GEMM relies on Triton kernels tuned for RTX 5090; disable for portability.
|
| 99 |
+
cfg.use_fp8_gemm = False
|
| 100 |
+
return cfg
|
| 101 |
+
|
| 102 |
+
|
| 103 |
+
print("[startup] building pipeline ...", flush=True)
|
| 104 |
+
set_seed(42)
|
| 105 |
+
torch.set_grad_enabled(False)
|
| 106 |
+
|
| 107 |
+
CONFIG = _build_config()
|
| 108 |
+
_vae = create_vae_from_config(CONFIG)
|
| 109 |
+
pipeline = CausalInferencePipeline(CONFIG, device=torch.device("cuda"), vae=_vae)
|
| 110 |
+
|
| 111 |
+
# Optional Triton norm/rope speedups from the upstream repo. They work on the
|
| 112 |
+
# training GPU; wrap defensively so a kernel incompatibility can't break boot.
|
| 113 |
+
try:
|
| 114 |
+
from wan.modules.helios_kernels import (
|
| 115 |
+
replace_all_norms_with_flash_norms,
|
| 116 |
+
replace_rope_with_flash_rope,
|
| 117 |
+
)
|
| 118 |
+
replace_all_norms_with_flash_norms(pipeline.generator.model)
|
| 119 |
+
replace_rope_with_flash_rope()
|
| 120 |
+
print("[startup] helios Triton kernels enabled.", flush=True)
|
| 121 |
+
except Exception as e: # pragma: no cover
|
| 122 |
+
print(f"[startup] helios kernels disabled ({e!r}); using eager norms/rope.", flush=True)
|
| 123 |
+
|
| 124 |
+
pipeline = pipeline.to(dtype=torch.bfloat16)
|
| 125 |
+
# Full model on GPU — we have the VRAM budget, no CPU offload.
|
| 126 |
+
pipeline.text_encoder.to(device="cuda")
|
| 127 |
+
pipeline.generator.to(device="cuda")
|
| 128 |
+
pipeline.vae.to(device="cuda")
|
| 129 |
+
if pipeline.encoder is not None:
|
| 130 |
+
pipeline.encoder.to(device="cuda")
|
| 131 |
+
pipeline.torch_dtype = torch.bfloat16
|
| 132 |
+
print("[startup] pipeline ready.", flush=True)
|
| 133 |
+
|
| 134 |
+
NUM_FPB = int(getattr(CONFIG, "num_frame_per_block", 3))
|
| 135 |
+
_vae_for_shape = pipeline.encoder if pipeline.encoder is not None else pipeline.vae
|
| 136 |
+
LATENT_CHANNELS = _vae_for_shape.z_dim
|
| 137 |
+
UPSAMPLE = getattr(_vae_for_shape, "upsampling_factor", 16)
|
| 138 |
+
LATENT_H = STREAM_HEIGHT // UPSAMPLE
|
| 139 |
+
LATENT_W = STREAM_WIDTH // UPSAMPLE
|
| 140 |
+
|
| 141 |
+
_infer_lock = threading.Lock()
|
| 142 |
+
|
| 143 |
+
|
| 144 |
+
def _decode_block_to_frames(lat_block):
|
| 145 |
+
frames = []
|
| 146 |
+
|
| 147 |
+
class _Cap:
|
| 148 |
+
def append_data(self, f):
|
| 149 |
+
frames.append(f)
|
| 150 |
+
|
| 151 |
+
pipeline.decode_block_and_write(lat_block, _Cap())
|
| 152 |
+
return frames
|
| 153 |
+
|
| 154 |
+
|
| 155 |
+
def _estimate_duration(image, prompt, action_label, num_blocks, seed, *args, **kwargs):
|
| 156 |
+
try:
|
| 157 |
+
n = int(num_blocks)
|
| 158 |
+
except Exception:
|
| 159 |
+
n = 8
|
| 160 |
+
return min(230, 40 + int(n) * 12)
|
| 161 |
+
|
| 162 |
+
|
| 163 |
+
@spaces.GPU(duration=_estimate_duration)
|
| 164 |
+
def rollout(image, prompt, action_label, num_blocks, seed,
|
| 165 |
+
progress=gr.Progress(track_tqdm=False)):
|
| 166 |
+
"""Roll out an action-conditioned world video from a first-frame image.
|
| 167 |
+
|
| 168 |
+
Args:
|
| 169 |
+
image: first-frame image (PIL / filepath) to start the world from.
|
| 170 |
+
prompt: text description of the scene and desired dynamics.
|
| 171 |
+
action_label: which navigation action to hold during the rollout.
|
| 172 |
+
num_blocks: number of autoregressive latent blocks to generate (video length).
|
| 173 |
+
seed: RNG seed for reproducibility.
|
| 174 |
+
|
| 175 |
+
Returns:
|
| 176 |
+
Path to the generated MP4 video of the world rollout.
|
| 177 |
+
"""
|
| 178 |
+
if image is None:
|
| 179 |
+
raise gr.Error("Please provide a first-frame image to start the world.")
|
| 180 |
+
|
| 181 |
+
prompt = (prompt or DEFAULT_PROMPT).strip() or DEFAULT_PROMPT
|
| 182 |
+
keys = ACTIONS.get(action_label, {})
|
| 183 |
+
num_blocks = max(1, min(40, int(num_blocks)))
|
| 184 |
+
seed = int(seed)
|
| 185 |
+
|
| 186 |
+
device = torch.device("cuda")
|
| 187 |
+
set_seed(seed)
|
| 188 |
+
|
| 189 |
+
# Persist the first-frame image to a path (pipeline reads a file path).
|
| 190 |
+
if isinstance(image, str):
|
| 191 |
+
first_frame_path = image
|
| 192 |
+
else:
|
| 193 |
+
if not isinstance(image, Image.Image):
|
| 194 |
+
image = Image.fromarray(np.asarray(image))
|
| 195 |
+
tmp_img = tempfile.NamedTemporaryFile(suffix=".png", delete=False)
|
| 196 |
+
image.convert("RGB").save(tmp_img.name)
|
| 197 |
+
first_frame_path = tmp_img.name
|
| 198 |
+
|
| 199 |
+
with _infer_lock:
|
| 200 |
+
noise = torch.randn(
|
| 201 |
+
[1, NUM_FPB, LATENT_CHANNELS, LATENT_H, LATENT_W],
|
| 202 |
+
device=device, dtype=torch.bfloat16,
|
| 203 |
+
)
|
| 204 |
+
|
| 205 |
+
pipeline.set_prompts([prompt], device=device)
|
| 206 |
+
# No pre-generated head/left/right/front/back refs -> zero ref_latents
|
| 207 |
+
# (model ignores them; this matches the upstream fallback path).
|
| 208 |
+
pipeline.set_ref_latent_mask_from_exists_paths(
|
| 209 |
+
ref_dir=str(APP_DIR / "__no_such_ref_dir__"), device=device,
|
| 210 |
+
)
|
| 211 |
+
pipeline.reset_stream(batch_size=1, dtype=torch.bfloat16,
|
| 212 |
+
device=device, initial_latent=None)
|
| 213 |
+
pipeline.set_first_frame_latent(
|
| 214 |
+
first_frame_path, height=STREAM_HEIGHT, width=STREAM_WIDTH, device=device,
|
| 215 |
+
)
|
| 216 |
+
pipeline.set_act(keys, height=STREAM_HEIGHT, width=STREAM_WIDTH,
|
| 217 |
+
num_frames=NUM_FPB, device=device)
|
| 218 |
+
|
| 219 |
+
out_path = tempfile.NamedTemporaryFile(suffix=".mp4", delete=False).name
|
| 220 |
+
writer = imageio.get_writer(
|
| 221 |
+
out_path, fps=VIDEO_FPS, format="FFMPEG", codec="libx264",
|
| 222 |
+
ffmpeg_params=["-crf", "18", "-preset", "fast", "-pix_fmt", "yuv420p"],
|
| 223 |
+
)
|
| 224 |
+
try:
|
| 225 |
+
t0 = time.perf_counter()
|
| 226 |
+
for b in range(num_blocks):
|
| 227 |
+
pipeline.set_act(keys, height=STREAM_HEIGHT, width=STREAM_WIDTH,
|
| 228 |
+
num_frames=NUM_FPB, device=device)
|
| 229 |
+
lat_block = pipeline.generate_next_block(noise)
|
| 230 |
+
noise = torch.randn_like(noise)
|
| 231 |
+
frames = _decode_block_to_frames(lat_block)
|
| 232 |
+
for f in frames:
|
| 233 |
+
writer.append_data(np.asarray(f))
|
| 234 |
+
progress((b + 1) / num_blocks, desc=f"block {b + 1}/{num_blocks}")
|
| 235 |
+
elapsed = time.perf_counter() - t0
|
| 236 |
+
print(f"[rollout] {num_blocks} blocks in {elapsed:.1f}s "
|
| 237 |
+
f"(action={keys}, seed={seed})", flush=True)
|
| 238 |
+
finally:
|
| 239 |
+
writer.close()
|
| 240 |
+
# Reset streaming caches for the next request.
|
| 241 |
+
try:
|
| 242 |
+
pipeline.reset_stream(batch_size=1, dtype=torch.bfloat16,
|
| 243 |
+
device=device, initial_latent=None)
|
| 244 |
+
except Exception:
|
| 245 |
+
pass
|
| 246 |
+
try:
|
| 247 |
+
if hasattr(pipeline.vae, "model") and hasattr(pipeline.vae.model, "clear_cache"):
|
| 248 |
+
pipeline.vae.model.clear_cache()
|
| 249 |
+
except Exception:
|
| 250 |
+
pass
|
| 251 |
+
|
| 252 |
+
return out_path
|
| 253 |
+
|
| 254 |
+
|
| 255 |
+
# ── UI ───────���────────────────────────────────────────────────────────────────
|
| 256 |
+
CSS = """
|
| 257 |
+
#col-container { max-width: 1180px; margin: 0 auto; }
|
| 258 |
+
.dark .gradio-container { color: var(--body-text-color); }
|
| 259 |
+
"""
|
| 260 |
+
|
| 261 |
+
INTRO = """
|
| 262 |
+
# 🌍 ABot-World — Interactive World Rollout
|
| 263 |
+
Turn a single image into a navigable world. Pick a **first frame**, describe the
|
| 264 |
+
**scene**, choose an **action** to hold, and the model autoregressively rolls out
|
| 265 |
+
an action-conditioned video.
|
| 266 |
+
|
| 267 |
+
Built on [`acvlab/ABot-World-0-5B-LF`](https://huggingface.co/acvlab/ABot-World-0-5B-LF)
|
| 268 |
+
(Wan2.2-TI2V-5B) · [project page](https://amap-cvlab.github.io/ABot-World/) ·
|
| 269 |
+
[code](https://github.com/amap-cvlab/ABot-World)
|
| 270 |
+
"""
|
| 271 |
+
|
| 272 |
+
with gr.Blocks(theme=gr.themes.Citrus(), css=CSS) as demo:
|
| 273 |
+
with gr.Column(elem_id="col-container"):
|
| 274 |
+
gr.Markdown(INTRO)
|
| 275 |
+
with gr.Row():
|
| 276 |
+
with gr.Column(scale=1):
|
| 277 |
+
image_in = gr.Image(label="First frame", type="filepath", height=300)
|
| 278 |
+
prompt_in = gr.Textbox(
|
| 279 |
+
label="Scene prompt", value=DEFAULT_PROMPT, lines=3,
|
| 280 |
+
)
|
| 281 |
+
action_in = gr.Dropdown(
|
| 282 |
+
label="Action (held during rollout)",
|
| 283 |
+
choices=list(ACTIONS.keys()),
|
| 284 |
+
value="Forward (W)",
|
| 285 |
+
)
|
| 286 |
+
run_btn = gr.Button("Explore the world", variant="primary")
|
| 287 |
+
with gr.Accordion("Advanced settings", open=False):
|
| 288 |
+
blocks_in = gr.Slider(
|
| 289 |
+
label="Rollout length (latent blocks)",
|
| 290 |
+
minimum=2, maximum=24, step=1, value=8,
|
| 291 |
+
info="More blocks = longer video, longer runtime.",
|
| 292 |
+
)
|
| 293 |
+
seed_in = gr.Number(label="Seed", value=42, precision=0)
|
| 294 |
+
with gr.Column(scale=1):
|
| 295 |
+
video_out = gr.Video(label="World rollout", autoplay=True)
|
| 296 |
+
|
| 297 |
+
gr.Examples(
|
| 298 |
+
examples=[
|
| 299 |
+
["examples/example.png", DEFAULT_PROMPT, "Forward (W)", 8, 42],
|
| 300 |
+
["examples/desert_valley.png",
|
| 301 |
+
"A third-person view walking through a desert river valley, "
|
| 302 |
+
"sunlit rocky terrain, stable forward motion.",
|
| 303 |
+
"Forward (W)", 8, 42],
|
| 304 |
+
["examples/mountain_meadow.png",
|
| 305 |
+
"A first-person hike along an alpine meadow trail, lush grass, "
|
| 306 |
+
"distant peaks, gentle forward motion.",
|
| 307 |
+
"Forward + turn right", 8, 7],
|
| 308 |
+
["examples/forest_stream.png",
|
| 309 |
+
"A first-person walk stepping across stones in a forest stream, "
|
| 310 |
+
"dappled light, moss-covered rocks.",
|
| 311 |
+
"Turn left (A)", 8, 21],
|
| 312 |
+
],
|
| 313 |
+
inputs=[image_in, prompt_in, action_in, blocks_in, seed_in],
|
| 314 |
+
outputs=video_out,
|
| 315 |
+
fn=rollout,
|
| 316 |
+
cache_examples=True,
|
| 317 |
+
cache_mode="lazy",
|
| 318 |
+
)
|
| 319 |
+
|
| 320 |
+
run_btn.click(
|
| 321 |
+
fn=rollout,
|
| 322 |
+
inputs=[image_in, prompt_in, action_in, blocks_in, seed_in],
|
| 323 |
+
outputs=video_out,
|
| 324 |
+
api_name="rollout",
|
| 325 |
+
)
|
| 326 |
+
|
| 327 |
+
if __name__ == "__main__":
|
| 328 |
+
demo.queue().launch(mcp_server=True)
|
configs/default_config.yaml
ADDED
|
@@ -0,0 +1,6 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
vae_type: "taew2_2"
|
| 2 |
+
use_fp8_gemm: true
|
| 3 |
+
quant_type: "fp8-per-token"
|
| 4 |
+
|
| 5 |
+
taew2_2_checkpoint: "checkpoints/ABot-World-0-5B-LF/taew2_2.pth"
|
| 6 |
+
lightvae_encoder_checkpoint: "checkpoints/ABot-World-0-5B-LF/Wan2.2_VAE.pth"
|
configs/long_forcing_dmd.yaml
ADDED
|
@@ -0,0 +1,36 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
model_type: ci2v
|
| 2 |
+
|
| 3 |
+
denoising_step_list:
|
| 4 |
+
- 1000
|
| 5 |
+
- 750
|
| 6 |
+
- 500
|
| 7 |
+
- 250
|
| 8 |
+
warp_denoising_step: true
|
| 9 |
+
context_noise: 0
|
| 10 |
+
|
| 11 |
+
image_or_video_shape:
|
| 12 |
+
- 1
|
| 13 |
+
- 31
|
| 14 |
+
- 48
|
| 15 |
+
- 44
|
| 16 |
+
- 80
|
| 17 |
+
independent_first_frame: true
|
| 18 |
+
num_frame_per_block: 3
|
| 19 |
+
|
| 20 |
+
ref_num_slots: 5
|
| 21 |
+
ref_resolution: 512
|
| 22 |
+
|
| 23 |
+
model_kwargs:
|
| 24 |
+
model_name: checkpoints/ABot-World-0-5B-LF
|
| 25 |
+
timestep_shift: 5.0
|
| 26 |
+
local_attn_size: 21
|
| 27 |
+
use_relative_rope: true
|
| 28 |
+
downscale_factor_control_adapter: 16
|
| 29 |
+
|
| 30 |
+
text_encoder_kwargs:
|
| 31 |
+
tokenizer_path: checkpoints/ABot-World-0-5B-LF/google/umt5-xxl/
|
| 32 |
+
encoder_pth_path: checkpoints/ABot-World-0-5B-LF/models_t5_umt5-xxl-enc-bf16.pth
|
| 33 |
+
|
| 34 |
+
vae_kwargs:
|
| 35 |
+
vae_type: Wan2.2_VAE
|
| 36 |
+
pretrained_path: checkpoints/ABot-World-0-5B-LF/Wan2.2_VAE.pth
|
examples/desert_valley.png
ADDED
|
Git LFS Details
|
examples/example.png
ADDED
|
Git LFS Details
|
examples/forest_stream.png
ADDED
|
Git LFS Details
|
examples/mountain_meadow.png
ADDED
|
Git LFS Details
|
pipeline/__init__.py
ADDED
|
@@ -0,0 +1,5 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
from .causal_inference import CausalInferencePipeline
|
| 2 |
+
|
| 3 |
+
__all__ = [
|
| 4 |
+
"CausalInferencePipeline",
|
| 5 |
+
]
|
pipeline/causal_inference.py
ADDED
|
@@ -0,0 +1,822 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
from typing import List, Optional
|
| 2 |
+
from pathlib import Path
|
| 3 |
+
import math
|
| 4 |
+
import time
|
| 5 |
+
import numpy as np
|
| 6 |
+
import torch
|
| 7 |
+
from torch.nn import functional as Functional
|
| 8 |
+
from PIL import Image, ImageOps
|
| 9 |
+
from einops import repeat
|
| 10 |
+
|
| 11 |
+
from utils.wan_wrapper import (
|
| 12 |
+
WanDiffusionWrapper,
|
| 13 |
+
WanTextEncoder,
|
| 14 |
+
WanVAEWrapper,
|
| 15 |
+
model_kwargs_with_relative_rope,
|
| 16 |
+
)
|
| 17 |
+
|
| 18 |
+
from utils.memory import gpu, get_cuda_free_memory_gb, DynamicSwapInstaller, \
|
| 19 |
+
move_model_to_device_with_memory_preservation
|
| 20 |
+
import tqdm
|
| 21 |
+
|
| 22 |
+
# Supported image extensions for reference image loading (case-insensitive)
|
| 23 |
+
SUPPORTED_IMAGE_EXTS = (".jpg", ".jpeg", ".png", ".webp", ".bmp", ".tiff", ".tif")
|
| 24 |
+
|
| 25 |
+
|
| 26 |
+
def _resolve_ref_image_path(directory: Path, name: str) -> Optional[Path]:
|
| 27 |
+
"""Resolve a reference image name to an actual file with any supported extension."""
|
| 28 |
+
for ext in SUPPORTED_IMAGE_EXTS:
|
| 29 |
+
for variant in (ext, ext.upper()):
|
| 30 |
+
candidate = directory / f"{name}{variant}"
|
| 31 |
+
if candidate.is_file():
|
| 32 |
+
return candidate
|
| 33 |
+
return None
|
| 34 |
+
|
| 35 |
+
|
| 36 |
+
class CausalInferencePipeline(torch.nn.Module):
|
| 37 |
+
def __init__(
|
| 38 |
+
self,
|
| 39 |
+
args,
|
| 40 |
+
device,
|
| 41 |
+
generator=None,
|
| 42 |
+
text_encoder=None,
|
| 43 |
+
vae=None,
|
| 44 |
+
need_vae=True
|
| 45 |
+
):
|
| 46 |
+
super().__init__()
|
| 47 |
+
# Step 1: Initialize all models
|
| 48 |
+
model_kwargs = model_kwargs_with_relative_rope(args)
|
| 49 |
+
model_kwargs["local_attn_size"] = model_kwargs.get("local_attn_size", getattr(args, "local_attn_size", -1))
|
| 50 |
+
self.generator = WanDiffusionWrapper(
|
| 51 |
+
**model_kwargs, is_causal=True, model_type=args.model_type,
|
| 52 |
+
) if generator is None else generator
|
| 53 |
+
self.text_encoder = WanTextEncoder(tokenizer_path=args.text_encoder_kwargs.tokenizer_path,
|
| 54 |
+
encoder_pth_path=args.text_encoder_kwargs.encoder_pth_path) if text_encoder is None else text_encoder
|
| 55 |
+
if vae is not None and getattr(args, "model_type", "t2v") in ["ci2v", "ti2v"]:
|
| 56 |
+
# self.encoder = WanVAEWrapper()
|
| 57 |
+
self.encoder = WanVAEWrapper(**getattr(args, "vae_kwargs", {}))
|
| 58 |
+
else:
|
| 59 |
+
self.encoder = None
|
| 60 |
+
# self.vae = WanVAEWrapper() if vae is None else vae
|
| 61 |
+
self.vae = WanVAEWrapper(**getattr(args, "vae_kwargs", {})) if vae is None else vae
|
| 62 |
+
|
| 63 |
+
# Step 2: Initialize all causal hyperparmeters
|
| 64 |
+
self.scheduler = self.generator.get_scheduler()
|
| 65 |
+
self.denoising_step_list = torch.tensor(
|
| 66 |
+
args.denoising_step_list, dtype=torch.long)
|
| 67 |
+
if args.warp_denoising_step:
|
| 68 |
+
timesteps = torch.cat((self.scheduler.timesteps.cpu(), torch.tensor([0], dtype=torch.float32)))
|
| 69 |
+
self.denoising_step_list = timesteps[1000 - self.denoising_step_list]
|
| 70 |
+
|
| 71 |
+
self.num_transformer_blocks = 30
|
| 72 |
+
self.frame_seq_length = (
|
| 73 |
+
args.image_or_video_shape[3] * args.image_or_video_shape[4]
|
| 74 |
+
) // (
|
| 75 |
+
self.generator.model.patch_size[1] * self.generator.model.patch_size[2]
|
| 76 |
+
)
|
| 77 |
+
|
| 78 |
+
self.kv_cache1 = None
|
| 79 |
+
self.args = args
|
| 80 |
+
self.device = device
|
| 81 |
+
self.torch_dtype = torch.bfloat16
|
| 82 |
+
self.num_frame_per_block = getattr(args, "num_frame_per_block", 1)
|
| 83 |
+
self.independent_first_frame = args.independent_first_frame
|
| 84 |
+
self.local_attn_size = self.generator.model.local_attn_size
|
| 85 |
+
self.pyramid_sample_ratio = getattr(args, "pyramid_sample_ratio", None)
|
| 86 |
+
|
| 87 |
+
self.conditional_dict = None
|
| 88 |
+
self.crossattn_cache = None
|
| 89 |
+
|
| 90 |
+
# Streaming state
|
| 91 |
+
self.current_start_frame = 0
|
| 92 |
+
self.num_input_frames = 0
|
| 93 |
+
self._stream_block_diffusion_times: List[float] = []
|
| 94 |
+
self._stream_block_decode_times: List[float] = []
|
| 95 |
+
|
| 96 |
+
print(f"KV inference with {self.num_frame_per_block} frames per block")
|
| 97 |
+
|
| 98 |
+
if self.num_frame_per_block > 1:
|
| 99 |
+
self.generator.model.num_frame_per_block = self.num_frame_per_block
|
| 100 |
+
|
| 101 |
+
def _clear_generation_caches(self):
|
| 102 |
+
self.kv_cache1 = None
|
| 103 |
+
self.crossattn_cache = None
|
| 104 |
+
model = getattr(self.generator, "model", None)
|
| 105 |
+
if model is not None and hasattr(model, "clear_cache"):
|
| 106 |
+
model.clear_cache()
|
| 107 |
+
|
| 108 |
+
def _decode_output_to_video(self, output: torch.Tensor) -> torch.Tensor:
|
| 109 |
+
self._clear_generation_caches()
|
| 110 |
+
offload = bool(getattr(self.args, "offload_generator_for_decode", False))
|
| 111 |
+
generator_device = None
|
| 112 |
+
if offload:
|
| 113 |
+
try:
|
| 114 |
+
generator_device = next(self.generator.parameters()).device
|
| 115 |
+
except StopIteration:
|
| 116 |
+
generator_device = None
|
| 117 |
+
if generator_device is not None and generator_device.type == "cuda":
|
| 118 |
+
self.generator.to("cpu")
|
| 119 |
+
torch.cuda.empty_cache()
|
| 120 |
+
|
| 121 |
+
try:
|
| 122 |
+
return self.vae.decode_to_pixel(output, use_cache=False)
|
| 123 |
+
finally:
|
| 124 |
+
if (
|
| 125 |
+
offload
|
| 126 |
+
and generator_device is not None
|
| 127 |
+
and generator_device.type == "cuda"
|
| 128 |
+
and bool(getattr(self.args, "restore_generator_after_decode", False))
|
| 129 |
+
):
|
| 130 |
+
self.generator.to(generator_device)
|
| 131 |
+
|
| 132 |
+
def get_condition_split(self, conditional_dict, start, end):
|
| 133 |
+
# print(f"Getting condition split for frames {start} to {end}",conditional_dict.keys())
|
| 134 |
+
device = conditional_dict["prompt_embeds"].device
|
| 135 |
+
act_context = conditional_dict.get("act_context")
|
| 136 |
+
y = conditional_dict.get("y")
|
| 137 |
+
clip_fea = conditional_dict.get("clip_fea")
|
| 138 |
+
first_frame_latents = conditional_dict.get("first_frame_latents")
|
| 139 |
+
split_conditional_dict = {
|
| 140 |
+
"prompt_embeds": conditional_dict["prompt_embeds"],
|
| 141 |
+
"act_context": act_context[:, :, start:end].to(device) if act_context is not None else None,
|
| 142 |
+
"act_context_scale": conditional_dict.get("act_context_scale", 0.0),
|
| 143 |
+
"y": y[:, :, start:end].to(device) if y is not None else None,
|
| 144 |
+
"clip_fea": clip_fea,
|
| 145 |
+
"first_frame_latents": first_frame_latents,
|
| 146 |
+
"ref_latents": conditional_dict.get("ref_latents", None),
|
| 147 |
+
"ref_mask": conditional_dict.get("ref_mask", None),
|
| 148 |
+
}
|
| 149 |
+
return split_conditional_dict
|
| 150 |
+
|
| 151 |
+
|
| 152 |
+
def set_prompts(self, text_prompts: List[str], device: torch.device):
|
| 153 |
+
"""
|
| 154 |
+
刷新你的世界:重新算 conditional_dict,并让 crossattn_cache 在下次 generator 调用时重建。
|
| 155 |
+
"""
|
| 156 |
+
self.conditional_dict = self.text_encoder(text_prompts=text_prompts, device=device)
|
| 157 |
+
|
| 158 |
+
# 强制 cross-attn cache 重新 init(因为 prompt 变了)
|
| 159 |
+
if hasattr(self, "crossattn_cache") and self.crossattn_cache is not None:
|
| 160 |
+
for b in range(self.num_transformer_blocks):
|
| 161 |
+
self.crossattn_cache[b]["is_init"] = False
|
| 162 |
+
|
| 163 |
+
def set_ref_latent_mask_from_exists_paths(
|
| 164 |
+
self,
|
| 165 |
+
ref_dir: str,
|
| 166 |
+
device: Optional[torch.device] = None,
|
| 167 |
+
):
|
| 168 |
+
"""Set ref_latents and ref_mask from a directory of pre-generated reference images.
|
| 169 |
+
|
| 170 |
+
``ref_dir`` is expected to contain the 5 output images (head, left,
|
| 171 |
+
right, front, back) in any common image format (.jpg, .jpeg, .png,
|
| 172 |
+
.webp, .bmp, .tiff). The method only performs the VAE encoding step.
|
| 173 |
+
|
| 174 |
+
If any expected file is missing, falls back to zero ref_latents with zero
|
| 175 |
+
ref_mask (model ignores them).
|
| 176 |
+
Should be called after set_prompts().
|
| 177 |
+
"""
|
| 178 |
+
if self.conditional_dict is None:
|
| 179 |
+
raise RuntimeError("call set_prompts first")
|
| 180 |
+
dev = device or self.device
|
| 181 |
+
num_slots = getattr(self.args, "ref_num_slots", 5)
|
| 182 |
+
ref_resolution = getattr(self.args, "ref_resolution", 512)
|
| 183 |
+
# Get z_dim and upsampling_factor from encoder (preferred) or vae
|
| 184 |
+
vae_for_shape = self.encoder if self.encoder is not None else self.vae
|
| 185 |
+
z_dim = getattr(vae_for_shape, "z_dim", 48)
|
| 186 |
+
upsampling = getattr(vae_for_shape, "upsampling_factor", 16)
|
| 187 |
+
ref_h = ref_w = ref_resolution // upsampling
|
| 188 |
+
|
| 189 |
+
ref_names = ["head", "left", "right", "front", "back"]
|
| 190 |
+
ref_dir_path = Path(ref_dir)
|
| 191 |
+
|
| 192 |
+
# Resolve each reference name to an actual file with any supported extension
|
| 193 |
+
resolved_paths = []
|
| 194 |
+
missing = []
|
| 195 |
+
for name in ref_names[:num_slots]:
|
| 196 |
+
found = _resolve_ref_image_path(ref_dir_path, name)
|
| 197 |
+
if found is not None:
|
| 198 |
+
resolved_paths.append(found)
|
| 199 |
+
else:
|
| 200 |
+
missing.append(name)
|
| 201 |
+
|
| 202 |
+
if missing:
|
| 203 |
+
print(f"[set_ref_latent_mask_from_exists_paths] Missing ref images {missing} "
|
| 204 |
+
f"in {ref_dir}, falling back to zero ref_latents.")
|
| 205 |
+
ref_latents = torch.zeros(
|
| 206 |
+
[1, num_slots, z_dim, 1, ref_h, ref_w],
|
| 207 |
+
dtype=torch.bfloat16, device=dev,
|
| 208 |
+
)
|
| 209 |
+
ref_mask = torch.zeros([1, num_slots], dtype=torch.float32, device=dev)
|
| 210 |
+
source = "zeros_missing"
|
| 211 |
+
else:
|
| 212 |
+
vae = self.encoder if self.encoder is not None else self.vae
|
| 213 |
+
vae_device = next(vae.parameters()).device
|
| 214 |
+
|
| 215 |
+
latents_list = []
|
| 216 |
+
for img_path in resolved_paths:
|
| 217 |
+
img = Image.open(img_path).convert("RGB")
|
| 218 |
+
img = img.resize((ref_resolution, ref_resolution))
|
| 219 |
+
image = self.preprocess_image(img, device=dev, torch_dtype=torch.bfloat16)
|
| 220 |
+
ref_pixel = image.unsqueeze(2) # [1, C, 1, H, W]
|
| 221 |
+
with torch.no_grad():
|
| 222 |
+
latent = vae.encode_to_latent(ref_pixel.to(vae_device))
|
| 223 |
+
latent_slot = latent[0, 0].unsqueeze(1).to(device=dev, dtype=torch.bfloat16)
|
| 224 |
+
latents_list.append(latent_slot)
|
| 225 |
+
|
| 226 |
+
ref_latents = torch.stack(latents_list, dim=0).unsqueeze(0) # [1, K, C_lat, 1, H_lat, W_lat]
|
| 227 |
+
ref_mask = torch.ones([1, num_slots], dtype=torch.float32, device=dev)
|
| 228 |
+
source = "exists_paths"
|
| 229 |
+
|
| 230 |
+
self.conditional_dict["ref_latents"] = ref_latents
|
| 231 |
+
self.conditional_dict["ref_mask"] = ref_mask
|
| 232 |
+
print(f"[set_ref_latent_mask_from_exists_paths] ref_latents={ref_latents.shape}, "
|
| 233 |
+
f"ref_mask={ref_mask.shape}, "
|
| 234 |
+
f"source={source}")
|
| 235 |
+
|
| 236 |
+
def preprocess_image(self, image, torch_dtype=torch.bfloat16, device=None, pattern="B C H W", min_value=-1, max_value=1):
|
| 237 |
+
# Transform a PIL.Image to torch.Tensor
|
| 238 |
+
image = torch.Tensor(np.array(image, dtype=np.float32))
|
| 239 |
+
image = image.to(dtype=torch_dtype, device=device)
|
| 240 |
+
image = image * ((max_value - min_value) / 255) + min_value
|
| 241 |
+
image = repeat(image, f"H W C -> {pattern}", **({"B": 1} if "B" in pattern else {}))
|
| 242 |
+
return image
|
| 243 |
+
|
| 244 |
+
@torch.no_grad()
|
| 245 |
+
def set_first_frame_latent(
|
| 246 |
+
self,
|
| 247 |
+
ref_image_path: str,
|
| 248 |
+
height: int = 480,
|
| 249 |
+
width: int = 832,
|
| 250 |
+
device: Optional[torch.device] = None,
|
| 251 |
+
):
|
| 252 |
+
"""
|
| 253 |
+
设置首帧 clean latent:从路径加载图片,通过 vae 编码得到 latent,
|
| 254 |
+
存入 self.conditional_dict["first_frame_latents"]。
|
| 255 |
+
在 generate_next_block 的第一个 block 中会用它替换噪声的首帧。
|
| 256 |
+
"""
|
| 257 |
+
if self.conditional_dict is None:
|
| 258 |
+
raise RuntimeError("call set_prompts first")
|
| 259 |
+
dev = device or self.device
|
| 260 |
+
if self.encoder is not None:
|
| 261 |
+
vae = self.encoder
|
| 262 |
+
else:
|
| 263 |
+
vae = self.vae
|
| 264 |
+
|
| 265 |
+
if not ref_image_path:
|
| 266 |
+
self.conditional_dict["first_frame_latents"] = None
|
| 267 |
+
else:
|
| 268 |
+
input_image = Image.open(ref_image_path).convert("RGB")
|
| 269 |
+
# Preserve the source aspect ratio: scale to cover the target canvas
|
| 270 |
+
# then center-crop, instead of non-uniform stretching to (width, height).
|
| 271 |
+
input_image = ImageOps.fit(
|
| 272 |
+
input_image,
|
| 273 |
+
(width, height),
|
| 274 |
+
method=Image.LANCZOS,
|
| 275 |
+
centering=(0.5, 0.5),
|
| 276 |
+
)
|
| 277 |
+
image = self.preprocess_image(input_image, device=dev, torch_dtype=torch.bfloat16)
|
| 278 |
+
ref_pixel = image.unsqueeze(2) # [B, C, 1, H, W]
|
| 279 |
+
# encode_to_latent returns [B, 1, C_lat, H_lat, W_lat]
|
| 280 |
+
first_frame_latents = vae.encode_to_latent(ref_pixel.to(next(vae.parameters()).device))
|
| 281 |
+
self.conditional_dict["first_frame_latents"] = first_frame_latents.to(device=dev, dtype=torch.bfloat16)
|
| 282 |
+
print(f"set_first_frame_latent: shape={first_frame_latents.shape}")
|
| 283 |
+
|
| 284 |
+
def set_act(self, keys_dict, height: int = 480, width: int = 832, num_frames: int = 81, device: Optional[torch.device] = None):
|
| 285 |
+
"""
|
| 286 |
+
设置 act_context [B, 8]
|
| 287 |
+
device: 指定设备,默认 self.device
|
| 288 |
+
"""
|
| 289 |
+
if self.conditional_dict is None:
|
| 290 |
+
raise RuntimeError("call set_prompts first")
|
| 291 |
+
key_order = ['W', 'A', 'S', 'D', 'I', 'J', 'K', 'L']
|
| 292 |
+
action_list = [1 if keys_dict.get(k, False) else 0 for k in key_order]
|
| 293 |
+
action = torch.tensor(action_list, dtype=torch.float32, device=device).unsqueeze(0)
|
| 294 |
+
control_action_latents = action[:, None, None, :].repeat(1, height, width, 1).permute([3, 0, 1, 2]).unsqueeze(0).to(device=device, dtype=torch.bfloat16) #f,h,w,8 --> 1,8,f,h,w, ([1, 8, 1, 480, 832])
|
| 295 |
+
# repeat control_action_latents to num_frames
|
| 296 |
+
control_action_latents = control_action_latents.repeat_interleave(4, dim=1) # [1, 8*4, 1, 480, 832]
|
| 297 |
+
control_action_latents = control_action_latents.repeat(1, 1, num_frames, 1, 1) # [1, 8*4, F, 480, 832]
|
| 298 |
+
self.conditional_dict["act_context"] = control_action_latents
|
| 299 |
+
print(f'action_list, {action_list}')
|
| 300 |
+
|
| 301 |
+
|
| 302 |
+
def reset_stream(self, batch_size: int, dtype, device, initial_latent=None):
|
| 303 |
+
ref_cache_token_len = self._ref_cache_token_len(getattr(self, 'conditional_dict', None))
|
| 304 |
+
if self._kv_cache_needs_reinit(ref_cache_token_len):
|
| 305 |
+
self._initialize_kv_cache(
|
| 306 |
+
batch_size=batch_size, dtype=dtype, device=device,
|
| 307 |
+
extra_tokens=ref_cache_token_len,
|
| 308 |
+
)
|
| 309 |
+
self._initialize_crossattn_cache(batch_size=batch_size, dtype=dtype, device=device)
|
| 310 |
+
else:
|
| 311 |
+
# reset cross-attn cache flags
|
| 312 |
+
for b in range(self.num_transformer_blocks):
|
| 313 |
+
self.crossattn_cache[b]["is_init"] = False
|
| 314 |
+
# reset kv cache indices
|
| 315 |
+
for b in range(len(self.kv_cache1)):
|
| 316 |
+
self.kv_cache1[b]["global_end_index"].fill_(0)
|
| 317 |
+
self.kv_cache1[b]["local_end_index"].fill_(0)
|
| 318 |
+
|
| 319 |
+
self.current_start_frame = 0
|
| 320 |
+
self.num_input_frames = initial_latent.shape[1] if initial_latent is not None else 0
|
| 321 |
+
self._stream_block_diffusion_times = []
|
| 322 |
+
self._stream_block_decode_times = []
|
| 323 |
+
|
| 324 |
+
|
| 325 |
+
@torch.no_grad()
|
| 326 |
+
def generate_next_block(self, noise_block: torch.Tensor):
|
| 327 |
+
"""
|
| 328 |
+
noise_block: [B, F, C, H, W],其中 F=1 或 num_frame_per_block
|
| 329 |
+
return denoised_pred: [B, F, C, H, W] on GPU
|
| 330 |
+
"""
|
| 331 |
+
t0 = time.perf_counter()
|
| 332 |
+
B, F, C, H, W = noise_block.shape
|
| 333 |
+
noisy_input = noise_block
|
| 334 |
+
|
| 335 |
+
y = self.conditional_dict.get("y", None)
|
| 336 |
+
if y is not None:
|
| 337 |
+
y = y.clone()
|
| 338 |
+
|
| 339 |
+
action_context = self.conditional_dict.get("act_context", None)
|
| 340 |
+
if action_context is not None:
|
| 341 |
+
action_context = action_context.clone()
|
| 342 |
+
_, Ca, _, Ha, Wa = action_context.shape
|
| 343 |
+
|
| 344 |
+
first_frame_latents = self.conditional_dict.get("first_frame_latents", None)
|
| 345 |
+
replace_first = (self.current_start_frame - self.num_input_frames == 0) and (first_frame_latents is not None)
|
| 346 |
+
if replace_first:
|
| 347 |
+
print(f"replace the first timestep with the clean context. latent shape: {noisy_input.shape}, first_frame_latents shape: {first_frame_latents.shape}")
|
| 348 |
+
noisy_input[:, 0:1] = first_frame_latents
|
| 349 |
+
|
| 350 |
+
#Pyramid denoising
|
| 351 |
+
if self.pyramid_sample_ratio is not None:
|
| 352 |
+
assert len(self.pyramid_sample_ratio) == len(self.denoising_step_list)
|
| 353 |
+
noisy_input = noisy_input.reshape(B, F*C, H, W)
|
| 354 |
+
noisy_input = Functional.interpolate(noisy_input, scale_factor=self.pyramid_sample_ratio[0], mode='nearest')
|
| 355 |
+
noisy_input = noisy_input.reshape(B, F, C, int(H*self.pyramid_sample_ratio[0]), int(W*self.pyramid_sample_ratio[0]))
|
| 356 |
+
|
| 357 |
+
if y is not None:
|
| 358 |
+
cur_y = y.reshape(B, C*F, H, W)
|
| 359 |
+
cur_y = Functional.interpolate(cur_y, scale_factor=self.pyramid_sample_ratio[0], mode='nearest')
|
| 360 |
+
cur_y = cur_y.reshape(B, C, F, int(H*self.pyramid_sample_ratio[0]), int(W*self.pyramid_sample_ratio[0]))
|
| 361 |
+
self.conditional_dict["y"] = cur_y.clone()
|
| 362 |
+
|
| 363 |
+
if action_context is not None:
|
| 364 |
+
cur_action_context = action_context.reshape(B, Ca*F, Ha, Wa)
|
| 365 |
+
cur_action_context = Functional.interpolate(cur_action_context, scale_factor=self.pyramid_sample_ratio[0], mode='nearest')
|
| 366 |
+
cur_action_context = cur_action_context.reshape(B, Ca, F, int(Ha*self.pyramid_sample_ratio[0]), int(Wa*self.pyramid_sample_ratio[0]))
|
| 367 |
+
self.conditional_dict["act_context"] = cur_action_context.clone()
|
| 368 |
+
|
| 369 |
+
|
| 370 |
+
for index, current_timestep in enumerate(self.denoising_step_list):
|
| 371 |
+
timestep = torch.ones([B, F], device=noise_block.device, dtype=torch.int64) * current_timestep
|
| 372 |
+
if replace_first:
|
| 373 |
+
timestep[:, 0] = 0
|
| 374 |
+
|
| 375 |
+
if self.pyramid_sample_ratio is not None:
|
| 376 |
+
_, _, _, H_s, W_s = noisy_input.shape
|
| 377 |
+
|
| 378 |
+
_, denoised_pred = self.generator(
|
| 379 |
+
noisy_image_or_video=noisy_input,
|
| 380 |
+
conditional_dict=self.conditional_dict,
|
| 381 |
+
timestep=timestep,
|
| 382 |
+
kv_cache=self.kv_cache1,
|
| 383 |
+
crossattn_cache=self.crossattn_cache,
|
| 384 |
+
current_start=self.current_start_frame * self.frame_seq_length,
|
| 385 |
+
replace_first_timestep_and_noise_latents=replace_first,
|
| 386 |
+
)
|
| 387 |
+
|
| 388 |
+
if index < len(self.denoising_step_list) - 1:
|
| 389 |
+
next_timestep = self.denoising_step_list[index + 1]
|
| 390 |
+
if self.pyramid_sample_ratio is not None:
|
| 391 |
+
denoised_pred = denoised_pred.reshape(B, F*C, H_s, W_s)
|
| 392 |
+
denoised_pred = Functional.interpolate(denoised_pred, scale_factor=self.pyramid_sample_ratio[index+1], mode='nearest')
|
| 393 |
+
denoised_pred = denoised_pred.reshape(B, F, C, int(H_s*self.pyramid_sample_ratio[index+1]), int(W_s*self.pyramid_sample_ratio[index+1]))
|
| 394 |
+
|
| 395 |
+
if y is not None:
|
| 396 |
+
cur_y = y.reshape(B, C*F, H, W)
|
| 397 |
+
scale = math.prod(self.pyramid_sample_ratio[:index+2])
|
| 398 |
+
cur_y = Functional.interpolate(cur_y, scale_factor=scale, mode='nearest')
|
| 399 |
+
cur_y = cur_y.reshape(B, C, F, int(H*scale), int(W*scale))
|
| 400 |
+
self.conditional_dict["y"] = cur_y
|
| 401 |
+
|
| 402 |
+
if action_context is not None:
|
| 403 |
+
cur_action_context = action_context.reshape(B, Ca*F, Ha, Wa)
|
| 404 |
+
scale = math.prod(self.pyramid_sample_ratio[:index+2])
|
| 405 |
+
cur_action_context = Functional.interpolate(cur_action_context, scale_factor=scale, mode='nearest')
|
| 406 |
+
cur_action_context = cur_action_context.reshape(B, Ca, F, int(Ha*scale), int(Wa*scale))
|
| 407 |
+
self.conditional_dict["act_context"] = cur_action_context.clone()
|
| 408 |
+
|
| 409 |
+
noisy_input = self.scheduler.add_noise(
|
| 410 |
+
denoised_pred.flatten(0, 1),
|
| 411 |
+
torch.randn_like(denoised_pred.flatten(0, 1)),
|
| 412 |
+
next_timestep * torch.ones([B * F], device=noise_block.device, dtype=torch.long),
|
| 413 |
+
).unflatten(0, denoised_pred.shape[:2])
|
| 414 |
+
else:
|
| 415 |
+
noisy_input = denoised_pred
|
| 416 |
+
|
| 417 |
+
if replace_first:
|
| 418 |
+
noisy_input[:, 0:1] = first_frame_latents
|
| 419 |
+
|
| 420 |
+
# 用干净 context 更新 KV cache
|
| 421 |
+
context_timestep = torch.ones([B, F], device=noise_block.device, dtype=torch.int64) * self.args.context_noise
|
| 422 |
+
if replace_first:
|
| 423 |
+
context_timestep = context_timestep.clone()
|
| 424 |
+
context_timestep[:, 0] = 0
|
| 425 |
+
noisy_input[:, 0:1] = first_frame_latents
|
| 426 |
+
|
| 427 |
+
self.generator(
|
| 428 |
+
noisy_image_or_video=noisy_input,
|
| 429 |
+
conditional_dict=self.conditional_dict,
|
| 430 |
+
timestep=context_timestep,
|
| 431 |
+
kv_cache=self.kv_cache1,
|
| 432 |
+
crossattn_cache=self.crossattn_cache,
|
| 433 |
+
current_start=self.current_start_frame * self.frame_seq_length,
|
| 434 |
+
replace_first_timestep_and_noise_latents=replace_first,
|
| 435 |
+
)
|
| 436 |
+
|
| 437 |
+
self.current_start_frame += F
|
| 438 |
+
|
| 439 |
+
torch.cuda.synchronize()
|
| 440 |
+
block_diffusion_s = time.perf_counter() - t0
|
| 441 |
+
self._stream_block_diffusion_times.append(block_diffusion_s)
|
| 442 |
+
|
| 443 |
+
return noisy_input
|
| 444 |
+
|
| 445 |
+
@torch.no_grad()
|
| 446 |
+
def decode_block_and_write(self, latents_block: torch.Tensor, writer):
|
| 447 |
+
# latents_block: [B,F,C,H,W] on GPU
|
| 448 |
+
t0 = time.perf_counter()
|
| 449 |
+
vid = self.vae.decode_to_pixel(latents_block, use_cache=True, return_in_cpu=True)
|
| 450 |
+
torch.cuda.synchronize()
|
| 451 |
+
block_decode_s = time.perf_counter() - t0
|
| 452 |
+
self._stream_block_decode_times.append(block_decode_s)
|
| 453 |
+
|
| 454 |
+
vid = (vid * 0.5 + 0.5).clamp(0, 1)
|
| 455 |
+
# print('decode_block_and_write: ', latents_block.shape, vid.shape) # [1, 3, 16, 60, 104]) torch.Size([1, 12, 3, 480, 832])
|
| 456 |
+
frames = (vid[0].permute(0,2,3,1) * 255.0).clamp(0,255).to(torch.uint8).numpy()
|
| 457 |
+
for f in frames:
|
| 458 |
+
writer.append_data(f)
|
| 459 |
+
|
| 460 |
+
def inference(
|
| 461 |
+
self,
|
| 462 |
+
noise: torch.Tensor,
|
| 463 |
+
text_prompts: List[str],
|
| 464 |
+
conditional_dict=None,
|
| 465 |
+
initial_latent: Optional[torch.Tensor] = None,
|
| 466 |
+
return_latents: bool = False,
|
| 467 |
+
profile: bool = False,
|
| 468 |
+
low_memory: bool = False,
|
| 469 |
+
rectified_tf=False
|
| 470 |
+
) -> torch.Tensor:
|
| 471 |
+
"""
|
| 472 |
+
Perform inference on the given noise and text prompts.
|
| 473 |
+
Inputs:
|
| 474 |
+
noise (torch.Tensor): The input noise tensor of shape
|
| 475 |
+
(batch_size, num_output_frames, num_channels, height, width).
|
| 476 |
+
text_prompts (List[str]): The list of text prompts.
|
| 477 |
+
initial_latent (torch.Tensor): The initial latent tensor of shape
|
| 478 |
+
(batch_size, num_input_frames, num_channels, height, width).
|
| 479 |
+
If num_input_frames is 1, perform image to video.
|
| 480 |
+
If num_input_frames is greater than 1, perform video extension.
|
| 481 |
+
return_latents (bool): Whether to return the latents.
|
| 482 |
+
Outputs:
|
| 483 |
+
video (torch.Tensor): The generated video tensor of shape
|
| 484 |
+
(batch_size, num_output_frames, num_channels, height, width).
|
| 485 |
+
It is normalized to be in the range [0, 1].
|
| 486 |
+
"""
|
| 487 |
+
batch_size, num_frames, num_channels, height, width = noise.shape
|
| 488 |
+
if not self.independent_first_frame:
|
| 489 |
+
# If the first frame is independent and the first frame is provided, then the number of frames in the
|
| 490 |
+
# noise should still be a multiple of num_frame_per_block
|
| 491 |
+
# default here
|
| 492 |
+
# self.independent_first_frame: False
|
| 493 |
+
assert num_frames % self.num_frame_per_block == 0
|
| 494 |
+
num_blocks = num_frames // self.num_frame_per_block
|
| 495 |
+
elif self.independent_first_frame and initial_latent is not None:
|
| 496 |
+
assert (num_frames - 1) % self.num_frame_per_block == 0
|
| 497 |
+
num_blocks = (num_frames - 1) // self.num_frame_per_block
|
| 498 |
+
else:
|
| 499 |
+
# Using a [1, 4, 4, 4, 4, 4, ...] model to generate a video without image conditioning
|
| 500 |
+
assert (num_frames - 1) % self.num_frame_per_block == 0
|
| 501 |
+
num_blocks = (num_frames - 1) // self.num_frame_per_block
|
| 502 |
+
num_input_frames = initial_latent.shape[1] if initial_latent is not None else 0
|
| 503 |
+
num_output_frames = num_frames if self.independent_first_frame else num_frames + num_input_frames
|
| 504 |
+
# conditional_dict = self.text_encoder(
|
| 505 |
+
# text_prompts=text_prompts
|
| 506 |
+
# )
|
| 507 |
+
|
| 508 |
+
if low_memory:
|
| 509 |
+
gpu_memory_preservation = get_cuda_free_memory_gb(gpu) + 5
|
| 510 |
+
move_model_to_device_with_memory_preservation(self.text_encoder, target_device=gpu,
|
| 511 |
+
preserved_memory_gb=gpu_memory_preservation)
|
| 512 |
+
|
| 513 |
+
output = torch.zeros(
|
| 514 |
+
[batch_size, num_output_frames, num_channels, height, width],
|
| 515 |
+
device=noise.device,
|
| 516 |
+
dtype=noise.dtype
|
| 517 |
+
)
|
| 518 |
+
|
| 519 |
+
# Set up profiling if requested
|
| 520 |
+
if profile:
|
| 521 |
+
init_start = torch.cuda.Event(enable_timing=True)
|
| 522 |
+
init_end = torch.cuda.Event(enable_timing=True)
|
| 523 |
+
diffusion_start = torch.cuda.Event(enable_timing=True)
|
| 524 |
+
diffusion_end = torch.cuda.Event(enable_timing=True)
|
| 525 |
+
vae_start = torch.cuda.Event(enable_timing=True)
|
| 526 |
+
vae_end = torch.cuda.Event(enable_timing=True)
|
| 527 |
+
block_times = []
|
| 528 |
+
block_start = torch.cuda.Event(enable_timing=True)
|
| 529 |
+
block_end = torch.cuda.Event(enable_timing=True)
|
| 530 |
+
init_start.record()
|
| 531 |
+
|
| 532 |
+
# Step 1: Initialize KV cache to all zeros
|
| 533 |
+
ref_cache_token_len = self._ref_cache_token_len(conditional_dict)
|
| 534 |
+
if self._kv_cache_needs_reinit(ref_cache_token_len):
|
| 535 |
+
self._initialize_kv_cache(
|
| 536 |
+
batch_size=batch_size,
|
| 537 |
+
dtype=noise.dtype,
|
| 538 |
+
device=noise.device,
|
| 539 |
+
extra_tokens=ref_cache_token_len,
|
| 540 |
+
)
|
| 541 |
+
self._initialize_crossattn_cache(
|
| 542 |
+
batch_size=batch_size,
|
| 543 |
+
dtype=noise.dtype,
|
| 544 |
+
device=noise.device
|
| 545 |
+
)
|
| 546 |
+
else:
|
| 547 |
+
# reset cross attn cache
|
| 548 |
+
for block_index in range(self.num_transformer_blocks):
|
| 549 |
+
self.crossattn_cache[block_index]["is_init"] = False
|
| 550 |
+
# reset kv cache
|
| 551 |
+
for block_index in range(len(self.kv_cache1)):
|
| 552 |
+
self.kv_cache1[block_index]["global_end_index"] = torch.tensor(
|
| 553 |
+
[0], dtype=torch.long, device=noise.device)
|
| 554 |
+
self.kv_cache1[block_index]["local_end_index"] = torch.tensor(
|
| 555 |
+
[0], dtype=torch.long, device=noise.device)
|
| 556 |
+
self.kv_cache1[block_index]["ref_token_len"] = torch.tensor(
|
| 557 |
+
[0], dtype=torch.long, device=noise.device)
|
| 558 |
+
|
| 559 |
+
# Step 2: Cache context feature
|
| 560 |
+
current_start_frame = 0
|
| 561 |
+
if initial_latent is not None:
|
| 562 |
+
timestep = torch.ones([batch_size, 1], device=noise.device, dtype=torch.int64) * 0
|
| 563 |
+
if self.independent_first_frame:
|
| 564 |
+
# Assume num_input_frames is 1 + self.num_frame_per_block * num_input_blocks
|
| 565 |
+
assert (num_input_frames - 1) % self.num_frame_per_block == 0
|
| 566 |
+
num_input_blocks = (num_input_frames - 1) // self.num_frame_per_block
|
| 567 |
+
output[:, :1] = initial_latent[:, :1]
|
| 568 |
+
tmp_conditional_dict = self.get_condition_split(conditional_dict, 0, 1)
|
| 569 |
+
self.generator(
|
| 570 |
+
noisy_image_or_video=initial_latent[:, :1],
|
| 571 |
+
conditional_dict=tmp_conditional_dict,
|
| 572 |
+
timestep=timestep * 0,
|
| 573 |
+
kv_cache=self.kv_cache1,
|
| 574 |
+
crossattn_cache=self.crossattn_cache,
|
| 575 |
+
current_start=current_start_frame * self.frame_seq_length,
|
| 576 |
+
)
|
| 577 |
+
current_start_frame += 1
|
| 578 |
+
else:
|
| 579 |
+
# Assume num_input_frames is self.num_frame_per_block * num_input_blocks
|
| 580 |
+
assert num_input_frames % self.num_frame_per_block == 0
|
| 581 |
+
num_input_blocks = num_input_frames // self.num_frame_per_block
|
| 582 |
+
|
| 583 |
+
for _ in range(num_input_blocks):
|
| 584 |
+
current_ref_latents = \
|
| 585 |
+
initial_latent[:, current_start_frame:current_start_frame + self.num_frame_per_block]
|
| 586 |
+
output[:, current_start_frame:current_start_frame + self.num_frame_per_block] = current_ref_latents
|
| 587 |
+
tmp_conditional_dict = self.get_condition_split(conditional_dict, current_start_frame,
|
| 588 |
+
current_start_frame + self.num_frame_per_block)
|
| 589 |
+
|
| 590 |
+
self.generator(
|
| 591 |
+
noisy_image_or_video=current_ref_latents,
|
| 592 |
+
conditional_dict=tmp_conditional_dict,
|
| 593 |
+
timestep=timestep * 0,
|
| 594 |
+
kv_cache=self.kv_cache1,
|
| 595 |
+
crossattn_cache=self.crossattn_cache,
|
| 596 |
+
current_start=current_start_frame * self.frame_seq_length,
|
| 597 |
+
)
|
| 598 |
+
current_start_frame += self.num_frame_per_block
|
| 599 |
+
|
| 600 |
+
if profile:
|
| 601 |
+
init_end.record()
|
| 602 |
+
torch.cuda.synchronize()
|
| 603 |
+
diffusion_start.record()
|
| 604 |
+
|
| 605 |
+
# Step 3: Temporal denoising loop
|
| 606 |
+
all_num_frames = [self.num_frame_per_block] * num_blocks
|
| 607 |
+
if self.independent_first_frame and initial_latent is None:
|
| 608 |
+
all_num_frames = [1] + all_num_frames
|
| 609 |
+
for current_num_frames in tqdm.tqdm(all_num_frames):
|
| 610 |
+
if profile:
|
| 611 |
+
block_start.record()
|
| 612 |
+
|
| 613 |
+
if self.independent_first_frame:
|
| 614 |
+
noise_start_frame = current_start_frame
|
| 615 |
+
noise_end_frame = current_start_frame + current_num_frames
|
| 616 |
+
else:
|
| 617 |
+
noise_start_frame = current_start_frame - num_input_frames
|
| 618 |
+
noise_end_frame = current_start_frame + current_num_frames - num_input_frames
|
| 619 |
+
noisy_input = noise[:, noise_start_frame:noise_end_frame]
|
| 620 |
+
condition_start_frame = current_start_frame if self.independent_first_frame else noise_start_frame
|
| 621 |
+
condition_end_frame = (
|
| 622 |
+
current_start_frame + current_num_frames
|
| 623 |
+
if self.independent_first_frame
|
| 624 |
+
else noise_end_frame
|
| 625 |
+
)
|
| 626 |
+
tmp_conditional_dict = self.get_condition_split(
|
| 627 |
+
conditional_dict,
|
| 628 |
+
condition_start_frame,
|
| 629 |
+
condition_end_frame,
|
| 630 |
+
)
|
| 631 |
+
latents = noisy_input
|
| 632 |
+
replace_first_timestep_and_noise_latents = (
|
| 633 |
+
current_start_frame == 0
|
| 634 |
+
and num_input_frames == 0
|
| 635 |
+
and conditional_dict.get("first_frame_latents") is not None
|
| 636 |
+
)
|
| 637 |
+
if replace_first_timestep_and_noise_latents:
|
| 638 |
+
print(f"replace the first timestep with the clean context. latent shape: {noisy_input.shape}, conditional_dict['first_frame_latents'] shape: {conditional_dict['first_frame_latents'].shape}" )
|
| 639 |
+
noisy_input[:, 0:1] = conditional_dict["first_frame_latents"] # replace the first timestep with the clean context
|
| 640 |
+
|
| 641 |
+
# Step 3.1: Spatial denoising loop
|
| 642 |
+
for index, current_timestep in enumerate(self.denoising_step_list):
|
| 643 |
+
# print(f"current_timestep: {current_timestep}")
|
| 644 |
+
# set current timestep
|
| 645 |
+
timestep = torch.ones(
|
| 646 |
+
[batch_size, current_num_frames],
|
| 647 |
+
device=noise.device,
|
| 648 |
+
dtype=torch.int64) * current_timestep
|
| 649 |
+
|
| 650 |
+
if index < len(self.denoising_step_list) - 1:
|
| 651 |
+
_, denoised_pred = self.generator(
|
| 652 |
+
noisy_image_or_video=noisy_input,
|
| 653 |
+
conditional_dict=tmp_conditional_dict,
|
| 654 |
+
timestep=timestep,
|
| 655 |
+
kv_cache=self.kv_cache1,
|
| 656 |
+
crossattn_cache=self.crossattn_cache,
|
| 657 |
+
current_start=current_start_frame * self.frame_seq_length,
|
| 658 |
+
replace_first_timestep_and_noise_latents=replace_first_timestep_and_noise_latents,
|
| 659 |
+
)
|
| 660 |
+
next_timestep = self.denoising_step_list[index + 1]
|
| 661 |
+
noisy_input = self.scheduler.add_noise(
|
| 662 |
+
denoised_pred.flatten(0, 1),
|
| 663 |
+
torch.randn_like(denoised_pred.flatten(0, 1)),
|
| 664 |
+
next_timestep * torch.ones(
|
| 665 |
+
[batch_size * current_num_frames], device=noise.device, dtype=torch.long)
|
| 666 |
+
).unflatten(0, denoised_pred.shape[:2])
|
| 667 |
+
else:
|
| 668 |
+
# for getting real output
|
| 669 |
+
_, denoised_pred = self.generator(
|
| 670 |
+
noisy_image_or_video=noisy_input,
|
| 671 |
+
conditional_dict=tmp_conditional_dict,
|
| 672 |
+
timestep=timestep,
|
| 673 |
+
kv_cache=self.kv_cache1,
|
| 674 |
+
crossattn_cache=self.crossattn_cache,
|
| 675 |
+
current_start=current_start_frame * self.frame_seq_length,
|
| 676 |
+
replace_first_timestep_and_noise_latents=replace_first_timestep_and_noise_latents,
|
| 677 |
+
)
|
| 678 |
+
noisy_input = denoised_pred
|
| 679 |
+
if replace_first_timestep_and_noise_latents:
|
| 680 |
+
noisy_input[:, 0:1] = conditional_dict["first_frame_latents"]
|
| 681 |
+
# Step 3.2: record the model's output
|
| 682 |
+
output[:, current_start_frame:current_start_frame + current_num_frames] = noisy_input
|
| 683 |
+
|
| 684 |
+
# Step 3.3: rerun with timestep zero to update KV cache using clean context
|
| 685 |
+
context_timestep = torch.ones_like(timestep) * self.args.context_noise
|
| 686 |
+
if replace_first_timestep_and_noise_latents:
|
| 687 |
+
context_timestep = context_timestep.clone()
|
| 688 |
+
context_timestep[:, 0] = 0
|
| 689 |
+
noisy_input[:, 0:1] = conditional_dict["first_frame_latents"]
|
| 690 |
+
|
| 691 |
+
self.generator(
|
| 692 |
+
noisy_image_or_video=noisy_input,
|
| 693 |
+
conditional_dict=tmp_conditional_dict,
|
| 694 |
+
timestep=context_timestep,
|
| 695 |
+
kv_cache=self.kv_cache1,
|
| 696 |
+
crossattn_cache=self.crossattn_cache,
|
| 697 |
+
current_start=current_start_frame * self.frame_seq_length,
|
| 698 |
+
replace_first_timestep_and_noise_latents=replace_first_timestep_and_noise_latents,
|
| 699 |
+
)
|
| 700 |
+
|
| 701 |
+
if profile:
|
| 702 |
+
block_end.record()
|
| 703 |
+
torch.cuda.synchronize()
|
| 704 |
+
block_time = block_start.elapsed_time(block_end)
|
| 705 |
+
block_times.append(block_time)
|
| 706 |
+
|
| 707 |
+
# Step 3.4: update the start and end frame indices
|
| 708 |
+
current_start_frame += current_num_frames
|
| 709 |
+
|
| 710 |
+
if profile:
|
| 711 |
+
# End diffusion timing and synchronize CUDA
|
| 712 |
+
diffusion_end.record()
|
| 713 |
+
torch.cuda.synchronize()
|
| 714 |
+
diffusion_time = diffusion_start.elapsed_time(diffusion_end)
|
| 715 |
+
init_time = init_start.elapsed_time(init_end)
|
| 716 |
+
vae_start.record()
|
| 717 |
+
if rectified_tf:
|
| 718 |
+
mean = torch.load('laboratory/mean.pt').to(output.device)
|
| 719 |
+
std = torch.load('laboratory/std.pt').to(output.device)
|
| 720 |
+
noise = torch.randn_like(output).to(output.device)
|
| 721 |
+
output -= mean
|
| 722 |
+
# Step 4: Decode the output
|
| 723 |
+
video = self._decode_output_to_video(output)
|
| 724 |
+
video = (video * 0.5 + 0.5).clamp(0, 1)
|
| 725 |
+
|
| 726 |
+
if profile:
|
| 727 |
+
# End VAE timing and synchronize CUDA
|
| 728 |
+
vae_end.record()
|
| 729 |
+
torch.cuda.synchronize()
|
| 730 |
+
vae_time = vae_start.elapsed_time(vae_end)
|
| 731 |
+
total_time = init_time + diffusion_time + vae_time
|
| 732 |
+
|
| 733 |
+
print("Profiling results:")
|
| 734 |
+
print(f" - Initialization/caching time: {init_time:.2f} ms ({100 * init_time / total_time:.2f}%)")
|
| 735 |
+
print(f" - Diffusion generation time: {diffusion_time:.2f} ms ({100 * diffusion_time / total_time:.2f}%)")
|
| 736 |
+
for i, block_time in enumerate(block_times):
|
| 737 |
+
print(
|
| 738 |
+
f" - Block {i} generation time: {block_time:.2f} ms ({100 * block_time / diffusion_time:.2f}% of diffusion)")
|
| 739 |
+
print(f" - VAE decoding time: {vae_time:.2f} ms ({100 * vae_time / total_time:.2f}%)")
|
| 740 |
+
print(f" - Total time: {total_time:.2f} ms")
|
| 741 |
+
|
| 742 |
+
if return_latents:
|
| 743 |
+
return video, output
|
| 744 |
+
else:
|
| 745 |
+
return video
|
| 746 |
+
|
| 747 |
+
def _ref_cache_token_len(self, conditional_dict):
|
| 748 |
+
if conditional_dict is None:
|
| 749 |
+
return 0
|
| 750 |
+
ref_latents = conditional_dict.get("ref_latents", None)
|
| 751 |
+
if ref_latents is None:
|
| 752 |
+
return 0
|
| 753 |
+
in_ch_need = self.generator.model.patch_embedding.in_channels
|
| 754 |
+
if ref_latents.ndim == 4:
|
| 755 |
+
ref_latents = ref_latents.unsqueeze(0).unsqueeze(3)
|
| 756 |
+
elif ref_latents.ndim == 5 and ref_latents.shape[2] == in_ch_need:
|
| 757 |
+
ref_latents = ref_latents.unsqueeze(3)
|
| 758 |
+
elif ref_latents.ndim == 5:
|
| 759 |
+
ref_latents = ref_latents.unsqueeze(0)
|
| 760 |
+
if ref_latents.ndim != 6:
|
| 761 |
+
return 0
|
| 762 |
+
_, num_slots, _, ref_t, ref_h, ref_w = ref_latents.shape
|
| 763 |
+
patch_t, patch_h, patch_w = self.generator.model.patch_size
|
| 764 |
+
return (
|
| 765 |
+
int(num_slots)
|
| 766 |
+
* (int(ref_t) // int(patch_t))
|
| 767 |
+
* (int(ref_h) // int(patch_h))
|
| 768 |
+
* (int(ref_w) // int(patch_w))
|
| 769 |
+
)
|
| 770 |
+
|
| 771 |
+
def _kv_cache_needs_reinit(self, ref_cache_token_len):
|
| 772 |
+
if self.kv_cache1 is None:
|
| 773 |
+
return True
|
| 774 |
+
if getattr(self, "kv_cache_extra_tokens", 0) != int(ref_cache_token_len):
|
| 775 |
+
return True
|
| 776 |
+
return any("ref_token_len" not in cache for cache in self.kv_cache1)
|
| 777 |
+
|
| 778 |
+
def _initialize_kv_cache(self, batch_size, dtype, device, extra_tokens=0):
|
| 779 |
+
"""
|
| 780 |
+
Initialize a Per-GPU KV cache for the Wan model.
|
| 781 |
+
"""
|
| 782 |
+
kv_cache1 = []
|
| 783 |
+
if self.local_attn_size != -1:
|
| 784 |
+
# Use the local attention size to compute the KV cache size
|
| 785 |
+
print("self.local_attn_size:", self.local_attn_size)
|
| 786 |
+
kv_cache_size = self.local_attn_size * self.frame_seq_length
|
| 787 |
+
else:
|
| 788 |
+
# Use the default KV cache size
|
| 789 |
+
kv_cache_size = self.args.image_or_video_shape[1] * self.frame_seq_length
|
| 790 |
+
kv_cache_size = kv_cache_size + int(extra_tokens)
|
| 791 |
+
|
| 792 |
+
num_heads = self.generator.model.num_heads
|
| 793 |
+
head_dim = self.generator.model.dim // num_heads
|
| 794 |
+
|
| 795 |
+
for _ in range(self.num_transformer_blocks):
|
| 796 |
+
kv_cache1.append({
|
| 797 |
+
"k": torch.zeros([batch_size, kv_cache_size, num_heads, head_dim], dtype=dtype, device=device),
|
| 798 |
+
"v": torch.zeros([batch_size, kv_cache_size, num_heads, head_dim], dtype=dtype, device=device),
|
| 799 |
+
"global_end_index": torch.tensor([0], dtype=torch.long, device=device),
|
| 800 |
+
"local_end_index": torch.tensor([0], dtype=torch.long, device=device),
|
| 801 |
+
"ref_token_len": torch.tensor([0], dtype=torch.long, device=device),
|
| 802 |
+
})
|
| 803 |
+
|
| 804 |
+
self.kv_cache1 = kv_cache1 # always store the clean cache
|
| 805 |
+
self.kv_cache_extra_tokens = int(extra_tokens)
|
| 806 |
+
|
| 807 |
+
def _initialize_crossattn_cache(self, batch_size, dtype, device):
|
| 808 |
+
"""
|
| 809 |
+
Initialize a Per-GPU cross-attention cache for the Wan model.
|
| 810 |
+
"""
|
| 811 |
+
crossattn_cache = []
|
| 812 |
+
|
| 813 |
+
num_heads = self.generator.model.num_heads
|
| 814 |
+
head_dim = self.generator.model.dim // num_heads
|
| 815 |
+
|
| 816 |
+
for _ in range(self.num_transformer_blocks):
|
| 817 |
+
crossattn_cache.append({
|
| 818 |
+
"k": torch.zeros([batch_size, 512, num_heads, head_dim], dtype=dtype, device=device),
|
| 819 |
+
"v": torch.zeros([batch_size, 512, num_heads, head_dim], dtype=dtype, device=device),
|
| 820 |
+
"is_init": False
|
| 821 |
+
})
|
| 822 |
+
self.crossattn_cache = crossattn_cache
|
quantizer/__init__.py
ADDED
|
@@ -0,0 +1,49 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
from typing import Dict, List, Optional
|
| 16 |
+
|
| 17 |
+
from .quant import * # noqa: F401 F403
|
| 18 |
+
from .quant import DynamicDiTQuantizer
|
| 19 |
+
|
| 20 |
+
DEFAULT_FP8_INCLUDE_PATTERNS = ["blocks"]
|
| 21 |
+
DEFAULT_FP8_EXCLUDE_PATTERNS = []
|
| 22 |
+
|
| 23 |
+
|
| 24 |
+
def apply_fp8_quantization(
|
| 25 |
+
model,
|
| 26 |
+
quant_type: str = "fp8-per-token",
|
| 27 |
+
include_patterns: Optional[List[str]] = None,
|
| 28 |
+
exclude_patterns: Optional[List[str]] = None,
|
| 29 |
+
) -> Dict[str, object]:
|
| 30 |
+
"""Apply DynamicDiTQuantizer to the provided DiT model."""
|
| 31 |
+
final_include_patterns = (
|
| 32 |
+
include_patterns if include_patterns is not None else DEFAULT_FP8_INCLUDE_PATTERNS
|
| 33 |
+
)
|
| 34 |
+
final_exclude_patterns = (
|
| 35 |
+
exclude_patterns if exclude_patterns is not None else DEFAULT_FP8_EXCLUDE_PATTERNS
|
| 36 |
+
)
|
| 37 |
+
|
| 38 |
+
quantizer = DynamicDiTQuantizer(
|
| 39 |
+
quant_type=quant_type,
|
| 40 |
+
include_patterns=final_include_patterns,
|
| 41 |
+
exclude_patterns=final_exclude_patterns,
|
| 42 |
+
)
|
| 43 |
+
quantizer.convert_linear(model)
|
| 44 |
+
return {
|
| 45 |
+
"quantizer": quantizer,
|
| 46 |
+
"quant_type": quant_type,
|
| 47 |
+
"include_patterns": final_include_patterns,
|
| 48 |
+
"exclude_patterns": final_exclude_patterns,
|
| 49 |
+
}
|
quantizer/kernels/__init__.py
ADDED
|
@@ -0,0 +1,13 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
quantizer/kernels/python/__init__.py
ADDED
|
@@ -0,0 +1,13 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
quantizer/kernels/python/gemm/__init__.py
ADDED
|
@@ -0,0 +1,34 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
"""
|
| 16 |
+
FP8 GEMM kernels with automatic backend selection.
|
| 17 |
+
|
| 18 |
+
This module automatically selects between Triton (for Linux/CUDA) and
|
| 19 |
+
PyTorch (for Windows/CPU) implementations based on the runtime environment.
|
| 20 |
+
"""
|
| 21 |
+
|
| 22 |
+
# from angelslim.compressor._platform import use_triton
|
| 23 |
+
|
| 24 |
+
# Conditional imports based on platform/backend availability
|
| 25 |
+
# if use_triton():
|
| 26 |
+
from .fp8_gemm import fp8_gemm_triton_block
|
| 27 |
+
# else:
|
| 28 |
+
# # PyTorch fallback implementation
|
| 29 |
+
# from .fp8_gemm_torch import fp8_gemm_torch_block as fp8_gemm_triton_block
|
| 30 |
+
|
| 31 |
+
# Also export PyTorch version directly for explicit use
|
| 32 |
+
from .fp8_gemm_torch import fp8_gemm_torch_block
|
| 33 |
+
|
| 34 |
+
__all__ = ["fp8_gemm_triton_block", "fp8_gemm_torch_block"]
|
quantizer/kernels/python/gemm/fp8_gemm.py
ADDED
|
@@ -0,0 +1,112 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
import torch
|
| 16 |
+
import triton
|
| 17 |
+
import triton.language as tl
|
| 18 |
+
|
| 19 |
+
# modified from https://github.com/deepseek-ai/DeepSeek-V3/blob/main/inference/kernel.py
|
| 20 |
+
fp8_gemm_configs = [
|
| 21 |
+
triton.Config(
|
| 22 |
+
{"BLOCK_SIZE_M": block_m, "BLOCK_SIZE_N": block_n, "BLOCK_SIZE_K": 128},
|
| 23 |
+
num_stages=num_stages,
|
| 24 |
+
num_warps=8,
|
| 25 |
+
)
|
| 26 |
+
for block_m in [16, 32, 64]
|
| 27 |
+
for block_n in [32, 64, 128]
|
| 28 |
+
for num_stages in [3, 4, 5, 6]
|
| 29 |
+
]
|
| 30 |
+
|
| 31 |
+
|
| 32 |
+
@triton.autotune(configs=fp8_gemm_configs, key=["N", "K"])
|
| 33 |
+
@triton.jit
|
| 34 |
+
def _fp8_gemm_triton_block_kernel(
|
| 35 |
+
a_ptr,
|
| 36 |
+
b_ptr,
|
| 37 |
+
c_ptr,
|
| 38 |
+
a_s_ptr,
|
| 39 |
+
b_s_ptr,
|
| 40 |
+
M,
|
| 41 |
+
N: tl.constexpr,
|
| 42 |
+
K: tl.constexpr,
|
| 43 |
+
BLOCK_SIZE_M: tl.constexpr,
|
| 44 |
+
BLOCK_SIZE_N: tl.constexpr,
|
| 45 |
+
BLOCK_SIZE_K: tl.constexpr,
|
| 46 |
+
):
|
| 47 |
+
"""
|
| 48 |
+
Performs a matrix multiplication operation on FP8 matrices with scaling factors.
|
| 49 |
+
"""
|
| 50 |
+
pid_m = tl.program_id(axis=0)
|
| 51 |
+
pid_n = tl.program_id(axis=1)
|
| 52 |
+
k = tl.cdiv(K, BLOCK_SIZE_K)
|
| 53 |
+
offs_m = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M
|
| 54 |
+
offs_n = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
|
| 55 |
+
offs_k = tl.arange(0, BLOCK_SIZE_K)
|
| 56 |
+
a_ptrs = a_ptr + offs_m[:, None] * K + offs_k[None, :]
|
| 57 |
+
b_ptrs = b_ptr + offs_n[None, :] * K + offs_k[:, None]
|
| 58 |
+
a_s_ptrs = a_s_ptr + offs_m * k
|
| 59 |
+
b_s_ptrs = b_s_ptr + (offs_n // BLOCK_SIZE_K) * k
|
| 60 |
+
|
| 61 |
+
accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
|
| 62 |
+
for i in range(k):
|
| 63 |
+
a = tl.load(a_ptrs, mask=offs_k[None, :] < K - i * BLOCK_SIZE_K, other=0.0)
|
| 64 |
+
b = tl.load(b_ptrs, mask=offs_k[:, None] < K - i * BLOCK_SIZE_K, other=0.0)
|
| 65 |
+
a_s = tl.load(a_s_ptrs)
|
| 66 |
+
b_s = tl.load(b_s_ptrs)
|
| 67 |
+
|
| 68 |
+
accumulator += tl.dot(a, b) * a_s[:, None] * b_s[None, :]
|
| 69 |
+
a_ptrs += BLOCK_SIZE_K
|
| 70 |
+
b_ptrs += BLOCK_SIZE_K
|
| 71 |
+
a_s_ptrs += 1
|
| 72 |
+
b_s_ptrs += 1
|
| 73 |
+
c = accumulator.to(c_ptr.dtype.element_ty)
|
| 74 |
+
offs_m = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
|
| 75 |
+
offs_n = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
|
| 76 |
+
c_ptrs = c_ptr + offs_m[:, None] * N + offs_n[None, :]
|
| 77 |
+
mask = (offs_m[:, None] < M) & (offs_n[None, :] < N)
|
| 78 |
+
tl.store(c_ptrs, c, mask=mask)
|
| 79 |
+
|
| 80 |
+
|
| 81 |
+
# triton fp8 gemm for fp8 per-block weight & fp8 per-group activation
|
| 82 |
+
# modified from https://github.com/deepseek-ai/DeepSeek-V3/blob/main/inference/kernel.py
|
| 83 |
+
def fp8_gemm_triton_block(
|
| 84 |
+
a: torch.Tensor,
|
| 85 |
+
a_s: torch.Tensor,
|
| 86 |
+
b: torch.Tensor,
|
| 87 |
+
b_s: torch.Tensor,
|
| 88 |
+
out_dtype=torch.bfloat16,
|
| 89 |
+
bias=None,
|
| 90 |
+
) -> torch.Tensor:
|
| 91 |
+
"""
|
| 92 |
+
Perform a matrix multiplication using FP8 precision.
|
| 93 |
+
"""
|
| 94 |
+
assert a.is_contiguous() and b.is_contiguous()
|
| 95 |
+
assert a_s.is_contiguous() and b_s.is_contiguous()
|
| 96 |
+
K = a.size(-1)
|
| 97 |
+
M = a.numel() // K
|
| 98 |
+
N = b.size(0)
|
| 99 |
+
c = a.new_empty(*a.size()[:-1], N, dtype=out_dtype)
|
| 100 |
+
|
| 101 |
+
def grid(meta):
|
| 102 |
+
return (
|
| 103 |
+
triton.cdiv(M, meta["BLOCK_SIZE_M"]),
|
| 104 |
+
triton.cdiv(N, meta["BLOCK_SIZE_N"]),
|
| 105 |
+
)
|
| 106 |
+
|
| 107 |
+
_fp8_gemm_triton_block_kernel[grid](a, b, c, a_s, b_s, M, N, K)
|
| 108 |
+
|
| 109 |
+
if bias is not None:
|
| 110 |
+
c += bias
|
| 111 |
+
|
| 112 |
+
return c
|
quantizer/kernels/python/gemm/fp8_gemm_torch.py
ADDED
|
@@ -0,0 +1,217 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
"""
|
| 16 |
+
Pure PyTorch implementation of FP8 block-wise GEMM.
|
| 17 |
+
|
| 18 |
+
This module provides CPU/Windows-compatible implementations that mirror
|
| 19 |
+
the Triton kernel for FP8 GEMM with block-wise quantization.
|
| 20 |
+
"""
|
| 21 |
+
|
| 22 |
+
from typing import Optional
|
| 23 |
+
|
| 24 |
+
import torch
|
| 25 |
+
|
| 26 |
+
|
| 27 |
+
def fp8_gemm_torch_block(
|
| 28 |
+
a: torch.Tensor,
|
| 29 |
+
a_s: torch.Tensor,
|
| 30 |
+
b: torch.Tensor,
|
| 31 |
+
b_s: torch.Tensor,
|
| 32 |
+
out_dtype: torch.dtype = torch.bfloat16,
|
| 33 |
+
bias: Optional[torch.Tensor] = None,
|
| 34 |
+
block_size: int = 128,
|
| 35 |
+
) -> torch.Tensor:
|
| 36 |
+
"""
|
| 37 |
+
Pure PyTorch implementation of FP8 GEMM with block-wise quantization.
|
| 38 |
+
|
| 39 |
+
Performs a matrix multiplication using FP8 precision with per-block scaling.
|
| 40 |
+
This implementation dequantizes the inputs and performs standard matmul.
|
| 41 |
+
|
| 42 |
+
C = (A * A_scale) @ (B * B_scale).T + bias
|
| 43 |
+
|
| 44 |
+
Args:
|
| 45 |
+
a: Input activation tensor in FP8 format, shape [..., K]
|
| 46 |
+
a_s: Scale tensor for A, shape [..., K // block_size] or [..., num_k_blocks]
|
| 47 |
+
b: Weight tensor in FP8 format, shape [N, K]
|
| 48 |
+
b_s: Scale tensor for B, shape [N // block_size, K // block_size]
|
| 49 |
+
out_dtype: Output data type (default: bfloat16)
|
| 50 |
+
bias: Optional bias tensor, shape [N]
|
| 51 |
+
block_size: Block size used for quantization (default: 128)
|
| 52 |
+
|
| 53 |
+
Returns:
|
| 54 |
+
Output tensor of shape [..., N]
|
| 55 |
+
"""
|
| 56 |
+
assert a.is_contiguous() and b.is_contiguous()
|
| 57 |
+
assert a_s.is_contiguous() and b_s.is_contiguous()
|
| 58 |
+
|
| 59 |
+
K = a.size(-1)
|
| 60 |
+
orig_shape = a.shape[:-1]
|
| 61 |
+
M = a.numel() // K
|
| 62 |
+
N = b.size(0)
|
| 63 |
+
|
| 64 |
+
# Reshape for computation
|
| 65 |
+
a_2d = a.view(M, K) # [M, K]
|
| 66 |
+
|
| 67 |
+
# Dequantize A: expand scales to match tensor dimensions
|
| 68 |
+
# a_s shape is typically [M, K//block_size]
|
| 69 |
+
a_s_2d = a_s.view(M, -1) # [M, num_k_blocks]
|
| 70 |
+
|
| 71 |
+
# Dequantize by expanding scales
|
| 72 |
+
a_dq = _dequantize_per_group(a_2d, a_s_2d, block_size, K)
|
| 73 |
+
|
| 74 |
+
# Dequantize B: b_s is [N//block_size, K//block_size]
|
| 75 |
+
b_dq = _dequantize_blockwise_2d(b, b_s, block_size)
|
| 76 |
+
|
| 77 |
+
# Perform matmul: [M, K] @ [K, N] -> [M, N]
|
| 78 |
+
c = torch.matmul(a_dq.to(out_dtype), b_dq.to(out_dtype).t())
|
| 79 |
+
|
| 80 |
+
# Reshape output
|
| 81 |
+
c = c.view(*orig_shape, N)
|
| 82 |
+
|
| 83 |
+
if bias is not None:
|
| 84 |
+
c = c + bias
|
| 85 |
+
|
| 86 |
+
return c
|
| 87 |
+
|
| 88 |
+
|
| 89 |
+
def _dequantize_per_group(
|
| 90 |
+
x: torch.Tensor,
|
| 91 |
+
s: torch.Tensor,
|
| 92 |
+
group_size: int,
|
| 93 |
+
K: int,
|
| 94 |
+
) -> torch.Tensor:
|
| 95 |
+
"""
|
| 96 |
+
Dequantize tensor with per-group scales.
|
| 97 |
+
|
| 98 |
+
Args:
|
| 99 |
+
x: Quantized tensor [M, K]
|
| 100 |
+
s: Scale tensor [M, num_groups]
|
| 101 |
+
group_size: Size of each group
|
| 102 |
+
K: Total size of last dimension
|
| 103 |
+
|
| 104 |
+
Returns:
|
| 105 |
+
Dequantized tensor [M, K]
|
| 106 |
+
"""
|
| 107 |
+
M = x.shape[0]
|
| 108 |
+
num_groups = s.shape[1]
|
| 109 |
+
|
| 110 |
+
x_float = x.to(torch.float32)
|
| 111 |
+
|
| 112 |
+
# Expand scales to match K dimension
|
| 113 |
+
# s: [M, num_groups] -> [M, K]
|
| 114 |
+
s_expanded = s.unsqueeze(-1).expand(M, num_groups, group_size)
|
| 115 |
+
s_expanded = s_expanded.reshape(M, num_groups * group_size)
|
| 116 |
+
|
| 117 |
+
# Handle case where K is not exactly num_groups * group_size
|
| 118 |
+
if s_expanded.shape[1] > K:
|
| 119 |
+
s_expanded = s_expanded[:, :K]
|
| 120 |
+
elif s_expanded.shape[1] < K:
|
| 121 |
+
# Pad with last scale value
|
| 122 |
+
pad_size = K - s_expanded.shape[1]
|
| 123 |
+
s_expanded = torch.nn.functional.pad(s_expanded, (0, pad_size), mode="replicate")
|
| 124 |
+
|
| 125 |
+
return x_float * s_expanded
|
| 126 |
+
|
| 127 |
+
|
| 128 |
+
def _dequantize_blockwise_2d(
|
| 129 |
+
x: torch.Tensor,
|
| 130 |
+
s: torch.Tensor,
|
| 131 |
+
block_size: int,
|
| 132 |
+
) -> torch.Tensor:
|
| 133 |
+
"""
|
| 134 |
+
Dequantize 2D tensor with block-wise scales.
|
| 135 |
+
|
| 136 |
+
Args:
|
| 137 |
+
x: Quantized tensor [N, K]
|
| 138 |
+
s: Scale tensor [n_blocks, k_blocks]
|
| 139 |
+
block_size: Block size
|
| 140 |
+
|
| 141 |
+
Returns:
|
| 142 |
+
Dequantized tensor [N, K]
|
| 143 |
+
"""
|
| 144 |
+
N, K = x.shape
|
| 145 |
+
n_blocks, k_blocks = s.shape
|
| 146 |
+
|
| 147 |
+
x_float = x.to(torch.float32)
|
| 148 |
+
y = torch.empty_like(x_float)
|
| 149 |
+
|
| 150 |
+
for nb in range(n_blocks):
|
| 151 |
+
n_start = nb * block_size
|
| 152 |
+
n_end = min(n_start + block_size, N)
|
| 153 |
+
for kb in range(k_blocks):
|
| 154 |
+
k_start = kb * block_size
|
| 155 |
+
k_end = min(k_start + block_size, K)
|
| 156 |
+
|
| 157 |
+
scale = s[nb, kb]
|
| 158 |
+
y[n_start:n_end, k_start:k_end] = x_float[n_start:n_end, k_start:k_end] * scale
|
| 159 |
+
|
| 160 |
+
return y
|
| 161 |
+
|
| 162 |
+
|
| 163 |
+
def fp8_gemm_torch_simple(
|
| 164 |
+
a: torch.Tensor,
|
| 165 |
+
a_s: torch.Tensor,
|
| 166 |
+
b: torch.Tensor,
|
| 167 |
+
b_s: torch.Tensor,
|
| 168 |
+
out_dtype: torch.dtype = torch.bfloat16,
|
| 169 |
+
bias: Optional[torch.Tensor] = None,
|
| 170 |
+
) -> torch.Tensor:
|
| 171 |
+
"""
|
| 172 |
+
Simplified PyTorch FP8 GEMM using full dequantization.
|
| 173 |
+
|
| 174 |
+
This version is simpler but may use more memory for large tensors.
|
| 175 |
+
|
| 176 |
+
Args:
|
| 177 |
+
a: Input activation tensor in FP8 format
|
| 178 |
+
a_s: Scale tensor for A
|
| 179 |
+
b: Weight tensor in FP8 format
|
| 180 |
+
b_s: Scale tensor for B
|
| 181 |
+
out_dtype: Output data type
|
| 182 |
+
bias: Optional bias tensor
|
| 183 |
+
|
| 184 |
+
Returns:
|
| 185 |
+
Output tensor
|
| 186 |
+
"""
|
| 187 |
+
K = a.size(-1)
|
| 188 |
+
orig_shape = a.shape[:-1]
|
| 189 |
+
M = a.numel() // K
|
| 190 |
+
N = b.size(0)
|
| 191 |
+
|
| 192 |
+
# Reshape
|
| 193 |
+
a_2d = a.view(M, K)
|
| 194 |
+
a_s_2d = a_s.view(M, -1)
|
| 195 |
+
|
| 196 |
+
# Simple dequantization: repeat scales to match dimensions
|
| 197 |
+
block_size = K // a_s_2d.shape[1] if a_s_2d.shape[1] > 0 else K
|
| 198 |
+
|
| 199 |
+
# Dequantize A
|
| 200 |
+
a_dq = a_2d.to(torch.float32)
|
| 201 |
+
if a_s_2d.shape[1] > 1:
|
| 202 |
+
a_s_expanded = a_s_2d.repeat_interleave(block_size, dim=1)[:, :K]
|
| 203 |
+
a_dq = a_dq * a_s_expanded
|
| 204 |
+
else:
|
| 205 |
+
a_dq = a_dq * a_s_2d
|
| 206 |
+
|
| 207 |
+
# Dequantize B
|
| 208 |
+
b_dq = _dequantize_blockwise_2d(b, b_s, block_size)
|
| 209 |
+
|
| 210 |
+
# Matmul
|
| 211 |
+
c = torch.matmul(a_dq.to(out_dtype), b_dq.to(out_dtype).t())
|
| 212 |
+
c = c.view(*orig_shape, N)
|
| 213 |
+
|
| 214 |
+
if bias is not None:
|
| 215 |
+
c = c + bias
|
| 216 |
+
|
| 217 |
+
return c
|
quantizer/kernels/python/mm/__init__.py
ADDED
|
File without changes
|
quantizer/kernels/python/mm/triton_kernels.py
ADDED
|
@@ -0,0 +1,743 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
import torch
|
| 2 |
+
from triton import Config, autotune, cdiv, jit, next_power_of_2
|
| 3 |
+
from triton import language as tl
|
| 4 |
+
|
| 5 |
+
_ordered_datatypes = [torch.int8, torch.float16, torch.bfloat16, torch.float32]
|
| 6 |
+
|
| 7 |
+
|
| 8 |
+
@jit
|
| 9 |
+
def gelu(x):
|
| 10 |
+
return x * tl.sigmoid(x * 1.702)
|
| 11 |
+
|
| 12 |
+
|
| 13 |
+
@jit
|
| 14 |
+
def int8_quantize_kernel(X, OUT, SCALES, HDIM, BLOCK_SIZE: tl.constexpr):
|
| 15 |
+
row_idx = tl.program_id(0)
|
| 16 |
+
x_ptr = X + row_idx * HDIM
|
| 17 |
+
out_ptr = OUT + row_idx * HDIM
|
| 18 |
+
h_offset = tl.arange(0, BLOCK_SIZE)
|
| 19 |
+
x = tl.load(x_ptr + h_offset, mask=h_offset < HDIM).to(tl.float32)
|
| 20 |
+
x_scale = 127.0 / tl.max(tl.abs(x))
|
| 21 |
+
x_scaled = x * x_scale
|
| 22 |
+
x_scaled += (0.5 * tl.where(x_scaled >= 0, 1, -1)).to(tl.int8)
|
| 23 |
+
tl.store(out_ptr + h_offset, x_scaled, mask=h_offset < HDIM)
|
| 24 |
+
tl.store(SCALES + row_idx, 1 / x_scale)
|
| 25 |
+
|
| 26 |
+
|
| 27 |
+
def int8_quantize_triton(x):
|
| 28 |
+
x_shape_orig = x.shape
|
| 29 |
+
x = x.view(-1, x_shape_orig[-1])
|
| 30 |
+
out = torch.empty(x_shape_orig, dtype=torch.int8, device=x.device)
|
| 31 |
+
scales = torch.empty(x.shape[0], dtype=torch.float32, device=x.device)
|
| 32 |
+
BLOCK_SIZE = next_power_of_2(x_shape_orig[-1])
|
| 33 |
+
grid = (x.shape[0],)
|
| 34 |
+
int8_quantize_kernel[grid](x, out, scales, x_shape_orig[-1], BLOCK_SIZE, num_warps=8)
|
| 35 |
+
return out.view(x_shape_orig), scales.view(x_shape_orig[:-1])
|
| 36 |
+
|
| 37 |
+
|
| 38 |
+
@jit
|
| 39 |
+
def fp8_quantize_kernel(X, OUT, SCALES, HDIM, BLOCK_SIZE: tl.constexpr, FP8_MAX_VAL: tl.constexpr):
|
| 40 |
+
row_idx = tl.program_id(0)
|
| 41 |
+
x_ptr = X + row_idx * HDIM
|
| 42 |
+
out_ptr = OUT + row_idx * HDIM
|
| 43 |
+
h_offset = tl.arange(0, BLOCK_SIZE)
|
| 44 |
+
x = tl.load(x_ptr + h_offset, mask=h_offset < HDIM).to(tl.float32)
|
| 45 |
+
absmax = tl.max(tl.abs(x))
|
| 46 |
+
eps = 1e-8
|
| 47 |
+
absmax = tl.maximum(absmax, eps)
|
| 48 |
+
x_scale = absmax / FP8_MAX_VAL
|
| 49 |
+
x_scaled = x / x_scale
|
| 50 |
+
x_scaled = tl.clamp(x_scaled, -FP8_MAX_VAL, FP8_MAX_VAL)
|
| 51 |
+
tl.store(out_ptr + h_offset, x_scaled, mask=h_offset < HDIM)
|
| 52 |
+
tl.store(SCALES + row_idx, x_scale)
|
| 53 |
+
|
| 54 |
+
|
| 55 |
+
def fp8_quantize_triton(x):
|
| 56 |
+
x_shape_orig = x.shape
|
| 57 |
+
x = x.view(-1, x_shape_orig[-1])
|
| 58 |
+
out_scaled = torch.empty(x_shape_orig, dtype=torch.float32, device=x.device)
|
| 59 |
+
scales = torch.empty(x.shape[0], dtype=torch.bfloat16, device=x.device)
|
| 60 |
+
BLOCK_SIZE = next_power_of_2(x_shape_orig[-1])
|
| 61 |
+
grid = (x.shape[0],)
|
| 62 |
+
FP8_MAX = 448.0
|
| 63 |
+
fp8_quantize_kernel[grid](x, out_scaled, scales, x_shape_orig[-1], BLOCK_SIZE, FP8_MAX_VAL=FP8_MAX, num_warps=8)
|
| 64 |
+
quantized = out_scaled.to(torch.float8_e4m3fn)
|
| 65 |
+
return quantized.view(x_shape_orig), scales.view(x_shape_orig[:-1])
|
| 66 |
+
|
| 67 |
+
|
| 68 |
+
def upcast_if_fp8(a):
|
| 69 |
+
if "fp8" in str(a):
|
| 70 |
+
return torch.float16
|
| 71 |
+
return a
|
| 72 |
+
|
| 73 |
+
|
| 74 |
+
def get_higher_dtype(a, b):
|
| 75 |
+
a = upcast_if_fp8(a)
|
| 76 |
+
b = upcast_if_fp8(b)
|
| 77 |
+
if a is b:
|
| 78 |
+
return a
|
| 79 |
+
|
| 80 |
+
assert a in _ordered_datatypes
|
| 81 |
+
assert b in _ordered_datatypes
|
| 82 |
+
|
| 83 |
+
for d in _ordered_datatypes:
|
| 84 |
+
if a is d:
|
| 85 |
+
return b
|
| 86 |
+
if b is d:
|
| 87 |
+
return a
|
| 88 |
+
|
| 89 |
+
|
| 90 |
+
@autotune(
|
| 91 |
+
configs=[
|
| 92 |
+
Config({"BLOCK_M": 256, "BLOCK_N": 128, "BLOCK_K": 128, "SPLIT_K": 1}, num_stages=3, num_warps=8),
|
| 93 |
+
Config({"BLOCK_M": 128, "BLOCK_N": 128, "BLOCK_K": 128, "SPLIT_K": 1}, num_stages=4, num_warps=8),
|
| 94 |
+
Config({"BLOCK_M": 256, "BLOCK_N": 128, "BLOCK_K": 64, "SPLIT_K": 1}, num_stages=3, num_warps=8),
|
| 95 |
+
Config({"BLOCK_M": 256, "BLOCK_N": 128, "BLOCK_K": 64, "SPLIT_K": 1}, num_stages=4, num_warps=8),
|
| 96 |
+
],
|
| 97 |
+
key=["M", "N", "K"],
|
| 98 |
+
)
|
| 99 |
+
@jit
|
| 100 |
+
def int8_gemm_bias_kernel(
|
| 101 |
+
A,
|
| 102 |
+
B,
|
| 103 |
+
BIAS,
|
| 104 |
+
A_SCALES,
|
| 105 |
+
B_SCALES,
|
| 106 |
+
C,
|
| 107 |
+
M,
|
| 108 |
+
N,
|
| 109 |
+
K, #
|
| 110 |
+
stride_am,
|
| 111 |
+
stride_ak, #
|
| 112 |
+
stride_bk,
|
| 113 |
+
stride_bn, #
|
| 114 |
+
stride_cm,
|
| 115 |
+
stride_cn, #
|
| 116 |
+
acc_dtype: tl.constexpr, #
|
| 117 |
+
fuse_gelu: tl.constexpr,
|
| 118 |
+
BLOCK_M: tl.constexpr,
|
| 119 |
+
BLOCK_N: tl.constexpr,
|
| 120 |
+
BLOCK_K: tl.constexpr, #
|
| 121 |
+
GROUP_M: tl.constexpr,
|
| 122 |
+
SPLIT_K: tl.constexpr,
|
| 123 |
+
EVEN_K: tl.constexpr,
|
| 124 |
+
AB_DTYPE: tl.constexpr, #
|
| 125 |
+
):
|
| 126 |
+
# matrix multiplication
|
| 127 |
+
pid = tl.program_id(0)
|
| 128 |
+
pid_z = tl.program_id(1)
|
| 129 |
+
grid_m = tl.cdiv(M, BLOCK_M)
|
| 130 |
+
grid_n = tl.cdiv(N, BLOCK_N)
|
| 131 |
+
# re-order program ID for better L2 performance
|
| 132 |
+
width = GROUP_M * grid_n
|
| 133 |
+
group_id = pid // width
|
| 134 |
+
group_size = min(grid_m - group_id * GROUP_M, GROUP_M)
|
| 135 |
+
pid_m = group_id * GROUP_M + (pid % group_size)
|
| 136 |
+
pid_n = (pid % width) // (group_size)
|
| 137 |
+
# do matrix multiplication
|
| 138 |
+
rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
|
| 139 |
+
rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 140 |
+
ram = tl.max_contiguous(tl.multiple_of(rm % M, BLOCK_M), BLOCK_M)
|
| 141 |
+
rbn = tl.max_contiguous(tl.multiple_of(rn % N, BLOCK_N), BLOCK_N)
|
| 142 |
+
rk = pid_z * BLOCK_K + tl.arange(0, BLOCK_K)
|
| 143 |
+
# pointers
|
| 144 |
+
A = A + (ram[:, None] * stride_am + rk[None, :] * stride_ak)
|
| 145 |
+
B = B + (rk[:, None] * stride_bk + rbn[None, :] * stride_bn)
|
| 146 |
+
acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=acc_dtype)
|
| 147 |
+
for k in range(0, tl.cdiv(K, BLOCK_K * SPLIT_K)):
|
| 148 |
+
if EVEN_K:
|
| 149 |
+
a = tl.load(A)
|
| 150 |
+
b = tl.load(B)
|
| 151 |
+
else:
|
| 152 |
+
k_remaining = K - k * (BLOCK_K * SPLIT_K)
|
| 153 |
+
_0 = tl.zeros((1, 1), dtype=C.dtype.element_ty)
|
| 154 |
+
a = tl.load(A, mask=rk[None, :] < k_remaining, other=_0)
|
| 155 |
+
b = tl.load(B, mask=rk[:, None] < k_remaining, other=_0)
|
| 156 |
+
if AB_DTYPE is not None:
|
| 157 |
+
a = a.to(AB_DTYPE)
|
| 158 |
+
b = b.to(AB_DTYPE)
|
| 159 |
+
|
| 160 |
+
acc = tl.dot(a, b, acc, out_dtype=acc_dtype, input_precision=None)
|
| 161 |
+
A += BLOCK_K * SPLIT_K * stride_ak
|
| 162 |
+
B += BLOCK_K * SPLIT_K * stride_bk
|
| 163 |
+
|
| 164 |
+
acc = acc.to(tl.float32)
|
| 165 |
+
a_scales_ptr = A_SCALES + pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
|
| 166 |
+
b_scales_ptr = B_SCALES + pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 167 |
+
a_scales = tl.load(a_scales_ptr) # [BM]
|
| 168 |
+
b_scales = tl.load(b_scales_ptr) # [BN]
|
| 169 |
+
# [BM, BN] * [BM, 1] * [1, BN]
|
| 170 |
+
|
| 171 |
+
bias_ptr = BIAS + pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 172 |
+
bias = tl.load(bias_ptr)
|
| 173 |
+
if fuse_gelu:
|
| 174 |
+
acc = gelu(((acc * a_scales[:, None]) * b_scales[None, :]) + bias[None, :])
|
| 175 |
+
else:
|
| 176 |
+
acc = ((acc * a_scales[:, None]) * b_scales[None, :]) + bias[None, :]
|
| 177 |
+
|
| 178 |
+
acc = acc.to(C.dtype.element_ty)
|
| 179 |
+
|
| 180 |
+
# rematerialize rm and rn to save registers
|
| 181 |
+
rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
|
| 182 |
+
rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 183 |
+
C = C + (rm[:, None] * stride_cm + rn[None, :] * stride_cn)
|
| 184 |
+
mask = (rm < M)[:, None] & (rn < N)[None, :]
|
| 185 |
+
# handles write-back with reduction-splitting
|
| 186 |
+
if SPLIT_K == 1:
|
| 187 |
+
tl.store(C, acc, mask=mask)
|
| 188 |
+
else:
|
| 189 |
+
tl.atomic_add(C, acc, mask=mask)
|
| 190 |
+
|
| 191 |
+
|
| 192 |
+
# @torch.compiler.disable()
|
| 193 |
+
def int8_gemm_bias_triton(a, b, bias, a_scales, b_scales, fuse_gelu=False, output_dtype=None):
|
| 194 |
+
device = a.device
|
| 195 |
+
# handle non-contiguous inputs if necessary
|
| 196 |
+
a_orig_shape = a.shape
|
| 197 |
+
a = a.view(-1, a.shape[-1])
|
| 198 |
+
b = b.t()
|
| 199 |
+
|
| 200 |
+
if a.stride(0) > 1 and a.stride(1) > 1:
|
| 201 |
+
a = a.contiguous()
|
| 202 |
+
if b.stride(0) > 1 and b.stride(1) > 1:
|
| 203 |
+
b = b.contiguous()
|
| 204 |
+
# checks constraints
|
| 205 |
+
assert a.shape[1] == b.shape[0], f"incompatible dimensions {a.shape} and {b.shape}"
|
| 206 |
+
M, K = a.shape
|
| 207 |
+
_, N = b.shape
|
| 208 |
+
out_shape = a_orig_shape[:-1] + (N,)
|
| 209 |
+
# common type between a and b
|
| 210 |
+
ab_dtype = get_higher_dtype(a.dtype, b.dtype)
|
| 211 |
+
|
| 212 |
+
# allocates output
|
| 213 |
+
if output_dtype is None:
|
| 214 |
+
output_dtype = ab_dtype
|
| 215 |
+
|
| 216 |
+
c = torch.empty((M, N), device=device, dtype=output_dtype)
|
| 217 |
+
|
| 218 |
+
# Allowed types for acc_type given the types of a and b.
|
| 219 |
+
supported_acc_dtypes = {
|
| 220 |
+
torch.float16: (torch.float32, torch.float16),
|
| 221 |
+
torch.bfloat16: (torch.float32, torch.bfloat16),
|
| 222 |
+
torch.float32: (torch.float32,),
|
| 223 |
+
torch.int8: (torch.int32,),
|
| 224 |
+
}
|
| 225 |
+
|
| 226 |
+
acc_dtype = supported_acc_dtypes[ab_dtype][0]
|
| 227 |
+
|
| 228 |
+
def to_tl_type(ty):
|
| 229 |
+
return getattr(tl, str(ty).split(".")[-1])
|
| 230 |
+
|
| 231 |
+
acc_dtype = to_tl_type(acc_dtype)
|
| 232 |
+
ab_dtype = to_tl_type(ab_dtype)
|
| 233 |
+
output_dtype = to_tl_type(output_dtype)
|
| 234 |
+
|
| 235 |
+
# Tensor cores support input with mixed float8 types.
|
| 236 |
+
if a.dtype in [tl.float8e4nv, tl.float8e5] and b.dtype in [
|
| 237 |
+
tl.float8e4nv,
|
| 238 |
+
tl.float8e5,
|
| 239 |
+
]:
|
| 240 |
+
ab_dtype = None
|
| 241 |
+
# launch kernel
|
| 242 |
+
grid = lambda META: ( # noqa E731
|
| 243 |
+
cdiv(M, META["BLOCK_M"]) * cdiv(N, META["BLOCK_N"]),
|
| 244 |
+
META["SPLIT_K"],
|
| 245 |
+
) # noqa E731
|
| 246 |
+
int8_gemm_bias_kernel[grid](
|
| 247 |
+
a,
|
| 248 |
+
b,
|
| 249 |
+
bias,
|
| 250 |
+
a_scales,
|
| 251 |
+
b_scales,
|
| 252 |
+
c,
|
| 253 |
+
M,
|
| 254 |
+
N,
|
| 255 |
+
K, #
|
| 256 |
+
a.stride(0),
|
| 257 |
+
a.stride(1), #
|
| 258 |
+
b.stride(0),
|
| 259 |
+
b.stride(1), #
|
| 260 |
+
c.stride(0),
|
| 261 |
+
c.stride(1), #
|
| 262 |
+
acc_dtype=acc_dtype, #
|
| 263 |
+
fuse_gelu=fuse_gelu,
|
| 264 |
+
GROUP_M=8,
|
| 265 |
+
EVEN_K=True,
|
| 266 |
+
AB_DTYPE=ab_dtype,
|
| 267 |
+
)
|
| 268 |
+
return c.view(*out_shape)
|
| 269 |
+
|
| 270 |
+
|
| 271 |
+
@autotune(
|
| 272 |
+
configs=[
|
| 273 |
+
Config({"BLOCK_M": 256, "BLOCK_N": 128, "BLOCK_K": 128, "SPLIT_K": 1}, num_stages=3, num_warps=8),
|
| 274 |
+
Config({"BLOCK_M": 128, "BLOCK_N": 128, "BLOCK_K": 128, "SPLIT_K": 1}, num_stages=4, num_warps=8),
|
| 275 |
+
Config({"BLOCK_M": 256, "BLOCK_N": 128, "BLOCK_K": 64, "SPLIT_K": 1}, num_stages=3, num_warps=8),
|
| 276 |
+
Config({"BLOCK_M": 256, "BLOCK_N": 128, "BLOCK_K": 64, "SPLIT_K": 1}, num_stages=4, num_warps=8),
|
| 277 |
+
],
|
| 278 |
+
key=["M", "N", "K"],
|
| 279 |
+
)
|
| 280 |
+
@jit
|
| 281 |
+
def int8_gemm_kernel(
|
| 282 |
+
A,
|
| 283 |
+
B,
|
| 284 |
+
A_SCALES,
|
| 285 |
+
B_SCALES,
|
| 286 |
+
C,
|
| 287 |
+
M,
|
| 288 |
+
N,
|
| 289 |
+
K, #
|
| 290 |
+
stride_am,
|
| 291 |
+
stride_ak, #
|
| 292 |
+
stride_bk,
|
| 293 |
+
stride_bn, #
|
| 294 |
+
stride_cm,
|
| 295 |
+
stride_cn, #
|
| 296 |
+
acc_dtype: tl.constexpr, #
|
| 297 |
+
fuse_gelu: tl.constexpr,
|
| 298 |
+
BLOCK_M: tl.constexpr,
|
| 299 |
+
BLOCK_N: tl.constexpr,
|
| 300 |
+
BLOCK_K: tl.constexpr, #
|
| 301 |
+
GROUP_M: tl.constexpr,
|
| 302 |
+
SPLIT_K: tl.constexpr,
|
| 303 |
+
EVEN_K: tl.constexpr,
|
| 304 |
+
AB_DTYPE: tl.constexpr, #
|
| 305 |
+
):
|
| 306 |
+
# matrix multiplication
|
| 307 |
+
pid = tl.program_id(0)
|
| 308 |
+
pid_z = tl.program_id(1)
|
| 309 |
+
grid_m = tl.cdiv(M, BLOCK_M)
|
| 310 |
+
grid_n = tl.cdiv(N, BLOCK_N)
|
| 311 |
+
# re-order program ID for better L2 performance
|
| 312 |
+
width = GROUP_M * grid_n
|
| 313 |
+
group_id = pid // width
|
| 314 |
+
group_size = min(grid_m - group_id * GROUP_M, GROUP_M)
|
| 315 |
+
pid_m = group_id * GROUP_M + (pid % group_size)
|
| 316 |
+
pid_n = (pid % width) // (group_size)
|
| 317 |
+
# do matrix multiplication
|
| 318 |
+
rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
|
| 319 |
+
rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 320 |
+
ram = tl.max_contiguous(tl.multiple_of(rm % M, BLOCK_M), BLOCK_M)
|
| 321 |
+
rbn = tl.max_contiguous(tl.multiple_of(rn % N, BLOCK_N), BLOCK_N)
|
| 322 |
+
rk = pid_z * BLOCK_K + tl.arange(0, BLOCK_K)
|
| 323 |
+
# pointers
|
| 324 |
+
A = A + (ram[:, None] * stride_am + rk[None, :] * stride_ak)
|
| 325 |
+
B = B + (rk[:, None] * stride_bk + rbn[None, :] * stride_bn)
|
| 326 |
+
acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=acc_dtype)
|
| 327 |
+
for k in range(0, tl.cdiv(K, BLOCK_K * SPLIT_K)):
|
| 328 |
+
if EVEN_K:
|
| 329 |
+
a = tl.load(A)
|
| 330 |
+
b = tl.load(B)
|
| 331 |
+
else:
|
| 332 |
+
k_remaining = K - k * (BLOCK_K * SPLIT_K)
|
| 333 |
+
_0 = tl.zeros((1, 1), dtype=C.dtype.element_ty)
|
| 334 |
+
a = tl.load(A, mask=rk[None, :] < k_remaining, other=_0)
|
| 335 |
+
b = tl.load(B, mask=rk[:, None] < k_remaining, other=_0)
|
| 336 |
+
if AB_DTYPE is not None:
|
| 337 |
+
a = a.to(AB_DTYPE)
|
| 338 |
+
b = b.to(AB_DTYPE)
|
| 339 |
+
|
| 340 |
+
acc = tl.dot(a, b, acc, out_dtype=acc_dtype, input_precision=None)
|
| 341 |
+
A += BLOCK_K * SPLIT_K * stride_ak
|
| 342 |
+
B += BLOCK_K * SPLIT_K * stride_bk
|
| 343 |
+
|
| 344 |
+
acc = acc.to(tl.float32)
|
| 345 |
+
a_scales_ptr = A_SCALES + pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
|
| 346 |
+
b_scales_ptr = B_SCALES + pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 347 |
+
a_scales = tl.load(a_scales_ptr) # [BM]
|
| 348 |
+
b_scales = tl.load(b_scales_ptr) # [BN]
|
| 349 |
+
# [BM, BN] * [BM, 1] * [1, BN]
|
| 350 |
+
if fuse_gelu:
|
| 351 |
+
acc = gelu((acc * a_scales[:, None]) * b_scales[None, :])
|
| 352 |
+
else:
|
| 353 |
+
acc = (acc * a_scales[:, None]) * b_scales[None, :]
|
| 354 |
+
|
| 355 |
+
acc = acc.to(C.dtype.element_ty)
|
| 356 |
+
|
| 357 |
+
# rematerialize rm and rn to save registers
|
| 358 |
+
rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
|
| 359 |
+
rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 360 |
+
C = C + (rm[:, None] * stride_cm + rn[None, :] * stride_cn)
|
| 361 |
+
mask = (rm < M)[:, None] & (rn < N)[None, :]
|
| 362 |
+
# handles write-back with reduction-splitting
|
| 363 |
+
if SPLIT_K == 1:
|
| 364 |
+
tl.store(C, acc, mask=mask)
|
| 365 |
+
else:
|
| 366 |
+
tl.atomic_add(C, acc, mask=mask)
|
| 367 |
+
|
| 368 |
+
|
| 369 |
+
# @torch.compiler.disable()
|
| 370 |
+
def int8_gemm_triton(a, b, a_scales, b_scales, fuse_gelu=False, output_dtype=None):
|
| 371 |
+
device = a.device
|
| 372 |
+
# handle non-contiguous inputs if necessary
|
| 373 |
+
# USE ONLY IN linear layer. NOT GENERAL MATRIX MULTIPLY
|
| 374 |
+
a_orig_shape = a.shape
|
| 375 |
+
a = a.view(-1, a.shape[-1])
|
| 376 |
+
b = b.t()
|
| 377 |
+
|
| 378 |
+
if a.stride(0) > 1 and a.stride(1) > 1:
|
| 379 |
+
a = a.contiguous()
|
| 380 |
+
if b.stride(0) > 1 and b.stride(1) > 1:
|
| 381 |
+
b = b.contiguous()
|
| 382 |
+
# checks constraints
|
| 383 |
+
assert a.shape[1] == b.shape[0], f"incompatible dimensions {a.shape} and {b.shape}"
|
| 384 |
+
M, K = a.shape
|
| 385 |
+
_, N = b.shape
|
| 386 |
+
out_shape = a_orig_shape[:-1] + (N,)
|
| 387 |
+
# common type between a and b
|
| 388 |
+
ab_dtype = get_higher_dtype(a.dtype, b.dtype)
|
| 389 |
+
|
| 390 |
+
# allocates output
|
| 391 |
+
if output_dtype is None:
|
| 392 |
+
output_dtype = ab_dtype
|
| 393 |
+
|
| 394 |
+
c = torch.empty((M, N), device=device, dtype=output_dtype)
|
| 395 |
+
|
| 396 |
+
# Allowed types for acc_type given the types of a and b.
|
| 397 |
+
supported_acc_dtypes = {
|
| 398 |
+
torch.float16: (torch.float32, torch.float16),
|
| 399 |
+
torch.bfloat16: (torch.float32, torch.bfloat16),
|
| 400 |
+
torch.float32: (torch.float32,),
|
| 401 |
+
torch.int8: (torch.int32,),
|
| 402 |
+
}
|
| 403 |
+
|
| 404 |
+
acc_dtype = supported_acc_dtypes[ab_dtype][0]
|
| 405 |
+
|
| 406 |
+
def to_tl_type(ty):
|
| 407 |
+
return getattr(tl, str(ty).split(".")[-1])
|
| 408 |
+
|
| 409 |
+
acc_dtype = to_tl_type(acc_dtype)
|
| 410 |
+
ab_dtype = to_tl_type(ab_dtype)
|
| 411 |
+
output_dtype = to_tl_type(output_dtype)
|
| 412 |
+
|
| 413 |
+
# Tensor cores support input with mixed float8 types.
|
| 414 |
+
if a.dtype in [tl.float8e4nv, tl.float8e5] and b.dtype in [
|
| 415 |
+
tl.float8e4nv,
|
| 416 |
+
tl.float8e5,
|
| 417 |
+
]:
|
| 418 |
+
ab_dtype = None
|
| 419 |
+
# launch kernel
|
| 420 |
+
grid = lambda META: ( # noqa E731
|
| 421 |
+
cdiv(M, META["BLOCK_M"]) * cdiv(N, META["BLOCK_N"]),
|
| 422 |
+
META["SPLIT_K"],
|
| 423 |
+
) # noqa E731
|
| 424 |
+
int8_gemm_kernel[grid](
|
| 425 |
+
a,
|
| 426 |
+
b,
|
| 427 |
+
a_scales,
|
| 428 |
+
b_scales,
|
| 429 |
+
c,
|
| 430 |
+
M,
|
| 431 |
+
N,
|
| 432 |
+
K, #
|
| 433 |
+
a.stride(0),
|
| 434 |
+
a.stride(1), #
|
| 435 |
+
b.stride(0),
|
| 436 |
+
b.stride(1), #
|
| 437 |
+
c.stride(0),
|
| 438 |
+
c.stride(1), #
|
| 439 |
+
acc_dtype=acc_dtype, #
|
| 440 |
+
fuse_gelu=fuse_gelu,
|
| 441 |
+
EVEN_K=True,
|
| 442 |
+
GROUP_M=8,
|
| 443 |
+
AB_DTYPE=ab_dtype,
|
| 444 |
+
)
|
| 445 |
+
return c.view(*out_shape)
|
| 446 |
+
|
| 447 |
+
|
| 448 |
+
@autotune(
|
| 449 |
+
configs=[
|
| 450 |
+
Config({"BLOCK_M": 256, "BLOCK_N": 128, "BLOCK_K": 128, "SPLIT_K": 1}, num_stages=3, num_warps=8),
|
| 451 |
+
Config({"BLOCK_M": 128, "BLOCK_N": 128, "BLOCK_K": 128, "SPLIT_K": 1}, num_stages=4, num_warps=8),
|
| 452 |
+
Config({"BLOCK_M": 256, "BLOCK_N": 128, "BLOCK_K": 64, "SPLIT_K": 1}, num_stages=3, num_warps=8),
|
| 453 |
+
Config({"BLOCK_M": 256, "BLOCK_N": 128, "BLOCK_K": 64, "SPLIT_K": 1}, num_stages=4, num_warps=8),
|
| 454 |
+
],
|
| 455 |
+
key=["M", "N", "K"],
|
| 456 |
+
)
|
| 457 |
+
@jit
|
| 458 |
+
def fp8_gemm_bias_kernel(
|
| 459 |
+
A,
|
| 460 |
+
B,
|
| 461 |
+
BIAS,
|
| 462 |
+
A_SCALES,
|
| 463 |
+
B_SCALES,
|
| 464 |
+
C,
|
| 465 |
+
M,
|
| 466 |
+
N,
|
| 467 |
+
K,
|
| 468 |
+
stride_am,
|
| 469 |
+
stride_ak,
|
| 470 |
+
stride_bk,
|
| 471 |
+
stride_bn,
|
| 472 |
+
stride_cm,
|
| 473 |
+
stride_cn,
|
| 474 |
+
fuse_gelu: tl.constexpr,
|
| 475 |
+
BLOCK_M: tl.constexpr,
|
| 476 |
+
BLOCK_N: tl.constexpr,
|
| 477 |
+
BLOCK_K: tl.constexpr,
|
| 478 |
+
GROUP_M: tl.constexpr,
|
| 479 |
+
SPLIT_K: tl.constexpr,
|
| 480 |
+
EVEN_K: tl.constexpr,
|
| 481 |
+
):
|
| 482 |
+
pid = tl.program_id(0)
|
| 483 |
+
pid_z = tl.program_id(1)
|
| 484 |
+
|
| 485 |
+
grid_m = tl.cdiv(M, BLOCK_M)
|
| 486 |
+
grid_n = tl.cdiv(N, BLOCK_N)
|
| 487 |
+
|
| 488 |
+
width = GROUP_M * grid_n
|
| 489 |
+
group_id = pid // width
|
| 490 |
+
group_size = min(grid_m - group_id * GROUP_M, GROUP_M)
|
| 491 |
+
pid_m = group_id * GROUP_M + (pid % group_size)
|
| 492 |
+
pid_n = (pid % width) // group_size
|
| 493 |
+
|
| 494 |
+
rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
|
| 495 |
+
rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 496 |
+
ram = tl.max_contiguous(tl.multiple_of(rm % M, BLOCK_M), BLOCK_M)
|
| 497 |
+
rbn = tl.max_contiguous(tl.multiple_of(rn % N, BLOCK_N), BLOCK_N)
|
| 498 |
+
|
| 499 |
+
rk = pid_z * BLOCK_K + tl.arange(0, BLOCK_K)
|
| 500 |
+
|
| 501 |
+
A_ptr = A + (ram[:, None] * stride_am + rk[None, :] * stride_ak)
|
| 502 |
+
B_ptr = B + (rk[:, None] * stride_bk + rbn[None, :] * stride_bn)
|
| 503 |
+
|
| 504 |
+
acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
|
| 505 |
+
|
| 506 |
+
for k in range(0, tl.cdiv(K, BLOCK_K * SPLIT_K)):
|
| 507 |
+
if EVEN_K:
|
| 508 |
+
a = tl.load(A_ptr)
|
| 509 |
+
b = tl.load(B_ptr)
|
| 510 |
+
else:
|
| 511 |
+
k_remaining = K - k * (BLOCK_K * SPLIT_K)
|
| 512 |
+
a = tl.load(A_ptr, mask=rk[None, :] < k_remaining, other=0.0)
|
| 513 |
+
b = tl.load(B_ptr, mask=rk[:, None] < k_remaining, other=0.0)
|
| 514 |
+
|
| 515 |
+
acc = tl.dot(a, b, acc, out_dtype=tl.float32, input_precision=None)
|
| 516 |
+
|
| 517 |
+
A_ptr += BLOCK_K * SPLIT_K * stride_ak
|
| 518 |
+
B_ptr += BLOCK_K * SPLIT_K * stride_bk
|
| 519 |
+
|
| 520 |
+
a_scales_ptr = A_SCALES + pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
|
| 521 |
+
b_scales_ptr = B_SCALES + pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 522 |
+
a_scales = tl.load(a_scales_ptr).to(tl.float32) # [BM]
|
| 523 |
+
b_scales = tl.load(b_scales_ptr).to(tl.float32) # [BN]
|
| 524 |
+
|
| 525 |
+
bias_ptr = BIAS + pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 526 |
+
bias = tl.load(bias_ptr).to(tl.float32) # [BN]
|
| 527 |
+
|
| 528 |
+
out = (acc * a_scales[:, None]) * b_scales[None, :] + bias[None, :]
|
| 529 |
+
if fuse_gelu:
|
| 530 |
+
out = gelu(out)
|
| 531 |
+
|
| 532 |
+
out = out.to(C.dtype.element_ty)
|
| 533 |
+
|
| 534 |
+
rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
|
| 535 |
+
rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 536 |
+
C_ptr = C + (rm[:, None] * stride_cm + rn[None, :] * stride_cn)
|
| 537 |
+
mask = (rm < M)[:, None] & (rn < N)[None, :]
|
| 538 |
+
|
| 539 |
+
if SPLIT_K == 1:
|
| 540 |
+
tl.store(C_ptr, out, mask=mask)
|
| 541 |
+
else:
|
| 542 |
+
tl.atomic_add(C_ptr, out, mask=mask)
|
| 543 |
+
|
| 544 |
+
|
| 545 |
+
def fp8_gemm_bias_triton(a, b, bias, a_scales, b_scales, fuse_gelu=False, output_dtype=None):
|
| 546 |
+
assert a.is_cuda and b.is_cuda, "This kernel is for CUDA"
|
| 547 |
+
assert a.dtype in (getattr(torch, "float8_e4m3fn", None), getattr(torch, "float8_e4m3fnuz", None)), f"a.dtype={a.dtype} is not FP8 E4M3"
|
| 548 |
+
assert b.dtype in (getattr(torch, "float8_e4m3fn", None), getattr(torch, "float8_e4m3fnuz", None)), f"b.dtype={b.dtype} is not FP8 E4M3"
|
| 549 |
+
|
| 550 |
+
a_orig_shape = a.shape
|
| 551 |
+
a2 = a.view(-1, a.shape[-1])
|
| 552 |
+
b2 = b.t()
|
| 553 |
+
|
| 554 |
+
if a2.stride(0) > 1 and a2.stride(1) > 1:
|
| 555 |
+
a2 = a2.contiguous()
|
| 556 |
+
if b2.stride(0) > 1 and b2.stride(1) > 1:
|
| 557 |
+
b2 = b2.contiguous()
|
| 558 |
+
|
| 559 |
+
M, K = a2.shape
|
| 560 |
+
_, N = b2.shape
|
| 561 |
+
out_shape = a_orig_shape[:-1] + (N,)
|
| 562 |
+
|
| 563 |
+
if output_dtype is None:
|
| 564 |
+
output_dtype = torch.float16
|
| 565 |
+
|
| 566 |
+
c = torch.empty((M, N), device=a.device, dtype=output_dtype)
|
| 567 |
+
|
| 568 |
+
grid = lambda META: (cdiv(M, META["BLOCK_M"]) * cdiv(N, META["BLOCK_N"]), META["SPLIT_K"]) # noqa E731
|
| 569 |
+
even_k = K % 128 == 0
|
| 570 |
+
|
| 571 |
+
fp8_gemm_bias_kernel[grid](
|
| 572 |
+
a2,
|
| 573 |
+
b2,
|
| 574 |
+
bias,
|
| 575 |
+
a_scales,
|
| 576 |
+
b_scales,
|
| 577 |
+
c,
|
| 578 |
+
M,
|
| 579 |
+
N,
|
| 580 |
+
K,
|
| 581 |
+
a2.stride(0),
|
| 582 |
+
a2.stride(1),
|
| 583 |
+
b2.stride(0),
|
| 584 |
+
b2.stride(1),
|
| 585 |
+
c.stride(0),
|
| 586 |
+
c.stride(1),
|
| 587 |
+
fuse_gelu=fuse_gelu,
|
| 588 |
+
GROUP_M=8,
|
| 589 |
+
EVEN_K=even_k,
|
| 590 |
+
)
|
| 591 |
+
return c.view(*out_shape)
|
| 592 |
+
|
| 593 |
+
|
| 594 |
+
@autotune(
|
| 595 |
+
configs=[
|
| 596 |
+
Config({"BLOCK_M": 256, "BLOCK_N": 128, "BLOCK_K": 128, "SPLIT_K": 1}, num_stages=3, num_warps=8),
|
| 597 |
+
Config({"BLOCK_M": 128, "BLOCK_N": 128, "BLOCK_K": 128, "SPLIT_K": 1}, num_stages=4, num_warps=8),
|
| 598 |
+
Config({"BLOCK_M": 256, "BLOCK_N": 128, "BLOCK_K": 64, "SPLIT_K": 1}, num_stages=3, num_warps=8),
|
| 599 |
+
Config({"BLOCK_M": 256, "BLOCK_N": 128, "BLOCK_K": 64, "SPLIT_K": 1}, num_stages=4, num_warps=8),
|
| 600 |
+
],
|
| 601 |
+
key=["M", "N", "K"],
|
| 602 |
+
)
|
| 603 |
+
@jit
|
| 604 |
+
def fp8_gemm_kernel(
|
| 605 |
+
A,
|
| 606 |
+
B,
|
| 607 |
+
A_SCALES,
|
| 608 |
+
B_SCALES,
|
| 609 |
+
C,
|
| 610 |
+
M,
|
| 611 |
+
N,
|
| 612 |
+
K,
|
| 613 |
+
stride_am,
|
| 614 |
+
stride_ak,
|
| 615 |
+
stride_bk,
|
| 616 |
+
stride_bn,
|
| 617 |
+
stride_cm,
|
| 618 |
+
stride_cn,
|
| 619 |
+
fuse_gelu: tl.constexpr,
|
| 620 |
+
BLOCK_M: tl.constexpr,
|
| 621 |
+
BLOCK_N: tl.constexpr,
|
| 622 |
+
BLOCK_K: tl.constexpr,
|
| 623 |
+
GROUP_M: tl.constexpr,
|
| 624 |
+
SPLIT_K: tl.constexpr,
|
| 625 |
+
EVEN_K: tl.constexpr,
|
| 626 |
+
):
|
| 627 |
+
pid = tl.program_id(0)
|
| 628 |
+
pid_z = tl.program_id(1)
|
| 629 |
+
|
| 630 |
+
grid_m = tl.cdiv(M, BLOCK_M)
|
| 631 |
+
grid_n = tl.cdiv(N, BLOCK_N)
|
| 632 |
+
|
| 633 |
+
width = GROUP_M * grid_n
|
| 634 |
+
group_id = pid // width
|
| 635 |
+
group_size = min(grid_m - group_id * GROUP_M, GROUP_M)
|
| 636 |
+
pid_m = group_id * GROUP_M + (pid % group_size)
|
| 637 |
+
pid_n = (pid % width) // group_size
|
| 638 |
+
|
| 639 |
+
rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
|
| 640 |
+
rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 641 |
+
|
| 642 |
+
ram = tl.max_contiguous(tl.multiple_of(rm % M, BLOCK_M), BLOCK_M)
|
| 643 |
+
rbn = tl.max_contiguous(tl.multiple_of(rn % N, BLOCK_N), BLOCK_N)
|
| 644 |
+
|
| 645 |
+
rk = pid_z * BLOCK_K + tl.arange(0, BLOCK_K)
|
| 646 |
+
|
| 647 |
+
A_ptr = A + (ram[:, None] * stride_am + rk[None, :] * stride_ak)
|
| 648 |
+
B_ptr = B + (rk[:, None] * stride_bk + rbn[None, :] * stride_bn)
|
| 649 |
+
|
| 650 |
+
acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
|
| 651 |
+
|
| 652 |
+
for k in range(0, tl.cdiv(K, BLOCK_K * SPLIT_K)):
|
| 653 |
+
if EVEN_K:
|
| 654 |
+
a = tl.load(A_ptr)
|
| 655 |
+
b = tl.load(B_ptr)
|
| 656 |
+
else:
|
| 657 |
+
k_remaining = K - k * (BLOCK_K * SPLIT_K)
|
| 658 |
+
a = tl.load(A_ptr, mask=rk[None, :] < k_remaining, other=0.0)
|
| 659 |
+
b = tl.load(B_ptr, mask=rk[:, None] < k_remaining, other=0.0)
|
| 660 |
+
|
| 661 |
+
acc = tl.dot(a, b, acc, out_dtype=tl.float32, input_precision=None)
|
| 662 |
+
|
| 663 |
+
A_ptr += BLOCK_K * SPLIT_K * stride_ak
|
| 664 |
+
B_ptr += BLOCK_K * SPLIT_K * stride_bk
|
| 665 |
+
|
| 666 |
+
a_scales_ptr = A_SCALES + pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
|
| 667 |
+
b_scales_ptr = B_SCALES + pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 668 |
+
a_scales = tl.load(a_scales_ptr).to(tl.float32) # [BM]
|
| 669 |
+
b_scales = tl.load(b_scales_ptr).to(tl.float32) # [BN]
|
| 670 |
+
|
| 671 |
+
out = (acc * a_scales[:, None]) * b_scales[None, :]
|
| 672 |
+
if fuse_gelu:
|
| 673 |
+
out = gelu(out)
|
| 674 |
+
|
| 675 |
+
out = out.to(C.dtype.element_ty)
|
| 676 |
+
|
| 677 |
+
rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
|
| 678 |
+
rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
|
| 679 |
+
C_ptr = C + (rm[:, None] * stride_cm + rn[None, :] * stride_cn)
|
| 680 |
+
mask = (rm < M)[:, None] & (rn < N)[None, :]
|
| 681 |
+
|
| 682 |
+
if SPLIT_K == 1:
|
| 683 |
+
tl.store(C_ptr, out, mask=mask)
|
| 684 |
+
else:
|
| 685 |
+
tl.atomic_add(C_ptr, out, mask=mask)
|
| 686 |
+
|
| 687 |
+
|
| 688 |
+
def fp8_gemm_triton(a, b, a_scales, b_scales, fuse_gelu=False, output_dtype=None):
|
| 689 |
+
assert a.is_cuda and b.is_cuda
|
| 690 |
+
e4m3_ok = []
|
| 691 |
+
if hasattr(torch, "float8_e4m3fn"):
|
| 692 |
+
e4m3_ok.append(torch.float8_e4m3fn)
|
| 693 |
+
if hasattr(torch, "float8_e4m3fnuz"):
|
| 694 |
+
e4m3_ok.append(torch.float8_e4m3fnuz)
|
| 695 |
+
e4m3_ok = tuple(e4m3_ok)
|
| 696 |
+
|
| 697 |
+
assert a.dtype in e4m3_ok, f"a.dtype={a.dtype} is not FP8 E4M3"
|
| 698 |
+
assert b.dtype in e4m3_ok, f"b.dtype={b.dtype} is not FP8 E4M3"
|
| 699 |
+
|
| 700 |
+
a_orig_shape = a.shape
|
| 701 |
+
a2 = a.view(-1, a.shape[-1])
|
| 702 |
+
b2 = b.t()
|
| 703 |
+
|
| 704 |
+
if a2.stride(0) > 1 and a2.stride(1) > 1:
|
| 705 |
+
a2 = a2.contiguous()
|
| 706 |
+
if b2.stride(0) > 1 and b2.stride(1) > 1:
|
| 707 |
+
b2 = b2.contiguous()
|
| 708 |
+
|
| 709 |
+
M, K = a2.shape
|
| 710 |
+
_, N = b2.shape
|
| 711 |
+
out_shape = a_orig_shape[:-1] + (N,)
|
| 712 |
+
|
| 713 |
+
if output_dtype is None:
|
| 714 |
+
output_dtype = torch.float16
|
| 715 |
+
|
| 716 |
+
c = torch.empty((M, N), device=a.device, dtype=output_dtype)
|
| 717 |
+
|
| 718 |
+
grid = lambda META: ( # noqa E731
|
| 719 |
+
cdiv(M, META["BLOCK_M"]) * cdiv(N, META["BLOCK_N"]),
|
| 720 |
+
META["SPLIT_K"],
|
| 721 |
+
) # noqa E731
|
| 722 |
+
even_k = K % 128 == 0
|
| 723 |
+
|
| 724 |
+
fp8_gemm_kernel[grid](
|
| 725 |
+
a2,
|
| 726 |
+
b2,
|
| 727 |
+
a_scales,
|
| 728 |
+
b_scales,
|
| 729 |
+
c,
|
| 730 |
+
M,
|
| 731 |
+
N,
|
| 732 |
+
K,
|
| 733 |
+
a2.stride(0),
|
| 734 |
+
a2.stride(1),
|
| 735 |
+
b2.stride(0),
|
| 736 |
+
b2.stride(1),
|
| 737 |
+
c.stride(0),
|
| 738 |
+
c.stride(1),
|
| 739 |
+
fuse_gelu=fuse_gelu,
|
| 740 |
+
GROUP_M=8,
|
| 741 |
+
EVEN_K=even_k,
|
| 742 |
+
)
|
| 743 |
+
return c.view(*out_shape)
|
quantizer/kernels/python/quantizers/__init__.py
ADDED
|
@@ -0,0 +1,46 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
"""
|
| 16 |
+
FP8 quantization kernels with automatic backend selection.
|
| 17 |
+
|
| 18 |
+
This module automatically selects between Triton (for Linux/CUDA) and
|
| 19 |
+
PyTorch (for Windows/CPU) implementations based on the runtime environment.
|
| 20 |
+
"""
|
| 21 |
+
|
| 22 |
+
# from angelslim.compressor._platform import use_triton
|
| 23 |
+
|
| 24 |
+
# Conditional imports based on platform/backend availability
|
| 25 |
+
# if use_triton():
|
| 26 |
+
from .fp8_per_block import fp8_per_block_quant_triton
|
| 27 |
+
from .fp8_per_token_group import fp8_per_token_group_quant_triton
|
| 28 |
+
# else:
|
| 29 |
+
# # PyTorch fallback implementations
|
| 30 |
+
# from .fp8_per_block_torch import (
|
| 31 |
+
# fp8_per_block_quant_torch as fp8_per_block_quant_triton,
|
| 32 |
+
# )
|
| 33 |
+
# from .fp8_per_token_group_torch import (
|
| 34 |
+
# fp8_per_token_group_quant_torch as fp8_per_token_group_quant_triton,
|
| 35 |
+
# )
|
| 36 |
+
|
| 37 |
+
# Also export PyTorch versions directly for explicit use
|
| 38 |
+
from .fp8_per_block_torch import fp8_per_block_quant_torch
|
| 39 |
+
from .fp8_per_token_group_torch import fp8_per_token_group_quant_torch
|
| 40 |
+
|
| 41 |
+
__all__ = [
|
| 42 |
+
"fp8_per_token_group_quant_triton",
|
| 43 |
+
"fp8_per_block_quant_triton",
|
| 44 |
+
"fp8_per_block_quant_torch",
|
| 45 |
+
"fp8_per_token_group_quant_torch",
|
| 46 |
+
]
|
quantizer/kernels/python/quantizers/fp8_per_block.py
ADDED
|
@@ -0,0 +1,78 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
from typing import Tuple
|
| 16 |
+
|
| 17 |
+
import torch
|
| 18 |
+
import triton
|
| 19 |
+
import triton.language as tl
|
| 20 |
+
|
| 21 |
+
|
| 22 |
+
# https://github.com/deepseek-ai/DeepSeek-V3/blob/main/inference/kernel.py
|
| 23 |
+
@triton.jit
|
| 24 |
+
def _fp8_per_block_quant_kernel(x_ptr, y_ptr, s_ptr, M, N, BLOCK_SIZE: tl.constexpr):
|
| 25 |
+
"""Quantizes FP32 tensor to FP8 format using block-wise quantization."""
|
| 26 |
+
pid_m = tl.program_id(axis=0)
|
| 27 |
+
pid_n = tl.program_id(axis=1)
|
| 28 |
+
n = tl.cdiv(N, BLOCK_SIZE)
|
| 29 |
+
|
| 30 |
+
offs_m = pid_m * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
|
| 31 |
+
offs_n = pid_n * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
|
| 32 |
+
offs = offs_m[:, None] * N + offs_n[None, :]
|
| 33 |
+
|
| 34 |
+
mask = (offs_m[:, None] < M) & (offs_n[None, :] < N)
|
| 35 |
+
x = tl.load(x_ptr + offs, mask=mask).to(tl.float32)
|
| 36 |
+
max_val = tl.max(tl.abs(x))
|
| 37 |
+
scale = max_val / 448.0
|
| 38 |
+
scale = tl.where(max_val == 0.0, 1.0, scale)
|
| 39 |
+
y = x / scale
|
| 40 |
+
y = y.to(y_ptr.dtype.element_ty)
|
| 41 |
+
|
| 42 |
+
tl.store(y_ptr + offs, y, mask=mask)
|
| 43 |
+
tl.store(s_ptr + pid_m * n + pid_n, scale)
|
| 44 |
+
|
| 45 |
+
|
| 46 |
+
# triton implementation
|
| 47 |
+
# for weight quantization on gpu
|
| 48 |
+
def fp8_per_block_quant_triton(
|
| 49 |
+
x: torch.Tensor, block_size: int = 128
|
| 50 |
+
) -> Tuple[torch.Tensor, torch.Tensor]:
|
| 51 |
+
"""
|
| 52 |
+
Quantizes a FP32 2D tensor to FP8 (E4M3FN) using block-wise quantization.
|
| 53 |
+
For each (block_size x block_size) block:
|
| 54 |
+
- scale = max(abs(block)) / 448.0 (FP8 E4M3FN max magnitude)
|
| 55 |
+
- if block is all zeros, use scale = 1.0 to avoid div-by-zero
|
| 56 |
+
- scale, clamp and cast to FP8
|
| 57 |
+
Returns:
|
| 58 |
+
y: Quantized FP8 tensor, same shape as input
|
| 59 |
+
s: Per-block scales, shape (num_blocks_M, num_blocks_N)
|
| 60 |
+
"""
|
| 61 |
+
assert x.is_contiguous()
|
| 62 |
+
assert x.dim() == 2
|
| 63 |
+
|
| 64 |
+
M, N = x.size()
|
| 65 |
+
y = torch.empty_like(x, dtype=torch.float8_e4m3fn)
|
| 66 |
+
m_blocks = triton.cdiv(M, block_size)
|
| 67 |
+
n_blocks = triton.cdiv(N, block_size)
|
| 68 |
+
s = torch.empty((m_blocks, n_blocks), dtype=torch.float32, device=x.device)
|
| 69 |
+
|
| 70 |
+
def grid(meta):
|
| 71 |
+
return (
|
| 72 |
+
triton.cdiv(M, meta["BLOCK_SIZE"]),
|
| 73 |
+
triton.cdiv(N, meta["BLOCK_SIZE"]),
|
| 74 |
+
)
|
| 75 |
+
|
| 76 |
+
_fp8_per_block_quant_kernel[grid](x, y, s, M, N, BLOCK_SIZE=block_size)
|
| 77 |
+
|
| 78 |
+
return y, s
|
quantizer/kernels/python/quantizers/fp8_per_block_torch.py
ADDED
|
@@ -0,0 +1,155 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
"""
|
| 16 |
+
Pure PyTorch implementation of per-block FP8 quantization.
|
| 17 |
+
|
| 18 |
+
This module provides CPU/Windows-compatible implementations that mirror
|
| 19 |
+
the Triton kernel for FP8 per-block quantization.
|
| 20 |
+
"""
|
| 21 |
+
|
| 22 |
+
from typing import Tuple
|
| 23 |
+
|
| 24 |
+
import torch
|
| 25 |
+
|
| 26 |
+
# FP8 E4M3 max value
|
| 27 |
+
FP8_MAX = 448.0
|
| 28 |
+
|
| 29 |
+
|
| 30 |
+
def fp8_per_block_quant_torch(
|
| 31 |
+
x: torch.Tensor, block_size: int = 128
|
| 32 |
+
) -> Tuple[torch.Tensor, torch.Tensor]:
|
| 33 |
+
"""
|
| 34 |
+
Pure PyTorch implementation of block-wise FP8 quantization.
|
| 35 |
+
|
| 36 |
+
Quantizes a FP32 2D tensor to FP8 (E4M3FN) using block-wise quantization.
|
| 37 |
+
For each (block_size x block_size) block:
|
| 38 |
+
- scale = max(abs(block)) / 448.0 (FP8 E4M3FN max magnitude)
|
| 39 |
+
- if block is all zeros, use scale = 1.0 to avoid div-by-zero
|
| 40 |
+
- scale, clamp and cast to FP8
|
| 41 |
+
|
| 42 |
+
Args:
|
| 43 |
+
x: Input tensor of shape (M, N), must be contiguous
|
| 44 |
+
block_size: Size of each quantization block (default: 128)
|
| 45 |
+
|
| 46 |
+
Returns:
|
| 47 |
+
Tuple of (quantized_tensor, scale_tensor):
|
| 48 |
+
- y: Quantized FP8 tensor, same shape as input
|
| 49 |
+
- s: Per-block scales, shape (num_blocks_M, num_blocks_N)
|
| 50 |
+
"""
|
| 51 |
+
assert x.is_contiguous(), "Input tensor must be contiguous"
|
| 52 |
+
assert x.dim() == 2, "Input tensor must be 2D"
|
| 53 |
+
|
| 54 |
+
M, N = x.size()
|
| 55 |
+
device = x.device
|
| 56 |
+
|
| 57 |
+
# Calculate number of blocks
|
| 58 |
+
m_blocks = (M + block_size - 1) // block_size
|
| 59 |
+
n_blocks = (N + block_size - 1) // block_size
|
| 60 |
+
|
| 61 |
+
# Output tensors
|
| 62 |
+
y = torch.empty_like(x, dtype=torch.float8_e4m3fn)
|
| 63 |
+
s = torch.empty((m_blocks, n_blocks), dtype=torch.float32, device=device)
|
| 64 |
+
|
| 65 |
+
# Convert to float32 for computation
|
| 66 |
+
x_float = x.to(torch.float32)
|
| 67 |
+
|
| 68 |
+
# Process each block
|
| 69 |
+
for mb in range(m_blocks):
|
| 70 |
+
m_start = mb * block_size
|
| 71 |
+
m_end = min(m_start + block_size, M)
|
| 72 |
+
for nb in range(n_blocks):
|
| 73 |
+
n_start = nb * block_size
|
| 74 |
+
n_end = min(n_start + block_size, N)
|
| 75 |
+
|
| 76 |
+
block = x_float[m_start:m_end, n_start:n_end]
|
| 77 |
+
max_val = block.abs().amax()
|
| 78 |
+
|
| 79 |
+
# Compute scale (guard against zero)
|
| 80 |
+
scale = max_val / FP8_MAX
|
| 81 |
+
if scale.item() == 0.0:
|
| 82 |
+
scale = torch.tensor(1.0, dtype=torch.float32, device=device)
|
| 83 |
+
|
| 84 |
+
# Quantize block
|
| 85 |
+
y_block = (block / scale).to(torch.float8_e4m3fn)
|
| 86 |
+
|
| 87 |
+
# Store results
|
| 88 |
+
y[m_start:m_end, n_start:n_end] = y_block
|
| 89 |
+
s[mb, nb] = scale
|
| 90 |
+
|
| 91 |
+
return y, s
|
| 92 |
+
|
| 93 |
+
|
| 94 |
+
def fp8_per_block_quant_torch_fast(
|
| 95 |
+
x: torch.Tensor, block_size: int = 128
|
| 96 |
+
) -> Tuple[torch.Tensor, torch.Tensor]:
|
| 97 |
+
"""
|
| 98 |
+
Optimized PyTorch implementation using tensor operations.
|
| 99 |
+
|
| 100 |
+
This is faster than the loop-based version for large tensors on GPU.
|
| 101 |
+
|
| 102 |
+
Args:
|
| 103 |
+
x: Input tensor of shape (M, N), must be contiguous
|
| 104 |
+
block_size: Size of each quantization block (default: 128)
|
| 105 |
+
|
| 106 |
+
Returns:
|
| 107 |
+
Tuple of (quantized_tensor, scale_tensor)
|
| 108 |
+
"""
|
| 109 |
+
assert x.is_contiguous(), "Input tensor must be contiguous"
|
| 110 |
+
assert x.dim() == 2, "Input tensor must be 2D"
|
| 111 |
+
|
| 112 |
+
M, N = x.size()
|
| 113 |
+
|
| 114 |
+
# Pad tensor to be divisible by block_size
|
| 115 |
+
pad_m = (block_size - M % block_size) % block_size
|
| 116 |
+
pad_n = (block_size - N % block_size) % block_size
|
| 117 |
+
|
| 118 |
+
if pad_m > 0 or pad_n > 0:
|
| 119 |
+
x_padded = torch.nn.functional.pad(x, (0, pad_n, 0, pad_m), value=0.0)
|
| 120 |
+
else:
|
| 121 |
+
x_padded = x
|
| 122 |
+
|
| 123 |
+
M_padded, N_padded = x_padded.size()
|
| 124 |
+
m_blocks = M_padded // block_size
|
| 125 |
+
n_blocks = N_padded // block_size
|
| 126 |
+
|
| 127 |
+
# Reshape to blocks: [m_blocks, block_size, n_blocks, block_size]
|
| 128 |
+
x_blocks = x_padded.view(m_blocks, block_size, n_blocks, block_size)
|
| 129 |
+
x_blocks = x_blocks.permute(0, 2, 1, 3).contiguous()
|
| 130 |
+
|
| 131 |
+
# Compute max absolute value per block
|
| 132 |
+
x_float = x_blocks.to(torch.float32)
|
| 133 |
+
max_vals = x_float.abs().amax(dim=(2, 3)) # [m_blocks, n_blocks]
|
| 134 |
+
|
| 135 |
+
# Compute scales
|
| 136 |
+
s = max_vals / FP8_MAX
|
| 137 |
+
s = torch.where(s == 0.0, torch.ones_like(s), s)
|
| 138 |
+
|
| 139 |
+
# Quantize
|
| 140 |
+
s_expanded = s[:, :, None, None]
|
| 141 |
+
y_blocks = (x_float / s_expanded).to(torch.float8_e4m3fn)
|
| 142 |
+
|
| 143 |
+
# Reshape back
|
| 144 |
+
y_blocks = y_blocks.permute(0, 2, 1, 3).contiguous()
|
| 145 |
+
y_padded = y_blocks.view(M_padded, N_padded)
|
| 146 |
+
|
| 147 |
+
# Remove padding
|
| 148 |
+
y = y_padded[:M, :N].contiguous()
|
| 149 |
+
|
| 150 |
+
# Adjust scale tensor size if needed
|
| 151 |
+
actual_m_blocks = (M + block_size - 1) // block_size
|
| 152 |
+
actual_n_blocks = (N + block_size - 1) // block_size
|
| 153 |
+
s = s[:actual_m_blocks, :actual_n_blocks].contiguous()
|
| 154 |
+
|
| 155 |
+
return y, s
|
quantizer/kernels/python/quantizers/fp8_per_token_group.py
ADDED
|
@@ -0,0 +1,170 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
|
| 16 |
+
from typing import Tuple
|
| 17 |
+
|
| 18 |
+
import torch
|
| 19 |
+
import triton
|
| 20 |
+
import triton.language as tl
|
| 21 |
+
|
| 22 |
+
|
| 23 |
+
# quant function for per-group fp8 activation
|
| 24 |
+
# https://github.com/sgl-project/sglang/
|
| 25 |
+
# blob/a167fd0bcb9ef4b0f4331a109e40c8cdc770b026/python/sglang/srt/layers/
|
| 26 |
+
# quantization/fp8_kernel.py#L116
|
| 27 |
+
@triton.jit
|
| 28 |
+
def _per_token_group_quant_fp8(
|
| 29 |
+
y_ptr,
|
| 30 |
+
y_q_ptr,
|
| 31 |
+
y_s_ptr,
|
| 32 |
+
y_stride,
|
| 33 |
+
N,
|
| 34 |
+
eps,
|
| 35 |
+
fp8_min,
|
| 36 |
+
fp8_max,
|
| 37 |
+
BLOCK: tl.constexpr,
|
| 38 |
+
):
|
| 39 |
+
"""A Triton-accelerated function for per-token-group quantization."""
|
| 40 |
+
g_id = tl.program_id(0)
|
| 41 |
+
y_ptr += g_id * y_stride
|
| 42 |
+
y_q_ptr += g_id * y_stride
|
| 43 |
+
y_s_ptr += g_id
|
| 44 |
+
|
| 45 |
+
cols = tl.arange(0, BLOCK)
|
| 46 |
+
mask = cols < N
|
| 47 |
+
|
| 48 |
+
y = tl.load(y_ptr + cols, mask=mask, other=0.0).to(tl.float32)
|
| 49 |
+
_absmax = tl.maximum(tl.max(tl.abs(y)), eps)
|
| 50 |
+
y_s = _absmax / fp8_max
|
| 51 |
+
y_s_inv = 1.0 / y_s
|
| 52 |
+
y_q = tl.clamp(y * y_s_inv, fp8_min, fp8_max).to(y_q_ptr.dtype.element_ty)
|
| 53 |
+
|
| 54 |
+
tl.store(y_q_ptr + cols, y_q, mask=mask)
|
| 55 |
+
tl.store(y_s_ptr, y_s)
|
| 56 |
+
|
| 57 |
+
|
| 58 |
+
@triton.jit
|
| 59 |
+
def _per_token_group_quant_fp8_colmajor(
|
| 60 |
+
y_ptr,
|
| 61 |
+
y_q_ptr,
|
| 62 |
+
y_s_ptr,
|
| 63 |
+
group_size,
|
| 64 |
+
y_num_columns,
|
| 65 |
+
y_s_col_stride,
|
| 66 |
+
eps,
|
| 67 |
+
fp8_min,
|
| 68 |
+
fp8_max,
|
| 69 |
+
BLOCK: tl.constexpr,
|
| 70 |
+
):
|
| 71 |
+
"""A Triton-accelerated function for per-token-group quantization."""
|
| 72 |
+
g_id = tl.program_id(0)
|
| 73 |
+
y_ptr += g_id * group_size
|
| 74 |
+
y_q_ptr += g_id * group_size
|
| 75 |
+
|
| 76 |
+
blocks_per_row = y_num_columns // group_size
|
| 77 |
+
scale_col = g_id % blocks_per_row
|
| 78 |
+
scale_row = g_id // blocks_per_row
|
| 79 |
+
y_s_ptr += scale_col * y_s_col_stride + scale_row
|
| 80 |
+
|
| 81 |
+
cols = tl.arange(0, BLOCK)
|
| 82 |
+
mask = cols < group_size
|
| 83 |
+
|
| 84 |
+
y = tl.load(y_ptr + cols, mask=mask, other=0.0).to(tl.float32)
|
| 85 |
+
_absmax = tl.maximum(tl.max(tl.abs(y)), eps)
|
| 86 |
+
y_s = _absmax / fp8_max
|
| 87 |
+
y_q = tl.clamp(y / y_s, fp8_min, fp8_max).to(y_q_ptr.dtype.element_ty)
|
| 88 |
+
|
| 89 |
+
tl.store(y_q_ptr + cols, y_q, mask=mask)
|
| 90 |
+
tl.store(y_s_ptr, y_s)
|
| 91 |
+
|
| 92 |
+
|
| 93 |
+
def fp8_per_token_group_quant_triton(
|
| 94 |
+
x: torch.Tensor,
|
| 95 |
+
group_size: int,
|
| 96 |
+
eps: float = 1e-10,
|
| 97 |
+
dtype: torch.dtype = torch.float8_e4m3fn,
|
| 98 |
+
column_major_scales: bool = False,
|
| 99 |
+
scale_tma_aligned: bool = False,
|
| 100 |
+
) -> Tuple[torch.Tensor, torch.Tensor]:
|
| 101 |
+
"""Function to perform per-token-group quantization on an input tensor `x`."""
|
| 102 |
+
assert (
|
| 103 |
+
x.shape[-1] % group_size == 0
|
| 104 |
+
), "the last dimension of `x` cannot be divisible by `group_size`"
|
| 105 |
+
assert x.is_contiguous(), "`x` is not contiguous"
|
| 106 |
+
|
| 107 |
+
finfo = torch.finfo(dtype)
|
| 108 |
+
fp8_max = finfo.max
|
| 109 |
+
fp8_min = -fp8_max
|
| 110 |
+
|
| 111 |
+
x_q = torch.empty_like(x, device=x.device, dtype=dtype)
|
| 112 |
+
M = x.numel() // group_size
|
| 113 |
+
N = group_size
|
| 114 |
+
|
| 115 |
+
if column_major_scales:
|
| 116 |
+
if scale_tma_aligned:
|
| 117 |
+
aligned_size = (x.shape[-2] + 3) // 4 * 4
|
| 118 |
+
x_s = torch.empty(
|
| 119 |
+
x.shape[:-2] + (x.shape[-1] // group_size, aligned_size),
|
| 120 |
+
device=x.device,
|
| 121 |
+
dtype=torch.float32,
|
| 122 |
+
).permute(-1, -2)[: x.shape[-2], :]
|
| 123 |
+
else:
|
| 124 |
+
x_s = torch.empty(
|
| 125 |
+
(x.shape[-1] // group_size,) + x.shape[:-1],
|
| 126 |
+
device=x.device,
|
| 127 |
+
dtype=torch.float32,
|
| 128 |
+
).permute(-1, -2)
|
| 129 |
+
else:
|
| 130 |
+
x_s = torch.empty(
|
| 131 |
+
x.shape[:-1] + (x.shape[-1] // group_size,),
|
| 132 |
+
device=x.device,
|
| 133 |
+
dtype=torch.float32,
|
| 134 |
+
)
|
| 135 |
+
|
| 136 |
+
BLOCK = triton.next_power_of_2(N)
|
| 137 |
+
num_warps = min(max(BLOCK // 256, 1), 8)
|
| 138 |
+
num_stages = 1
|
| 139 |
+
|
| 140 |
+
if column_major_scales:
|
| 141 |
+
_per_token_group_quant_fp8_colmajor[(M,)](
|
| 142 |
+
x,
|
| 143 |
+
x_q,
|
| 144 |
+
x_s,
|
| 145 |
+
group_size,
|
| 146 |
+
x.shape[1],
|
| 147 |
+
x_s.stride(1),
|
| 148 |
+
eps,
|
| 149 |
+
fp8_min=fp8_min,
|
| 150 |
+
fp8_max=fp8_max,
|
| 151 |
+
BLOCK=BLOCK,
|
| 152 |
+
num_warps=num_warps,
|
| 153 |
+
num_stages=num_stages,
|
| 154 |
+
)
|
| 155 |
+
else:
|
| 156 |
+
_per_token_group_quant_fp8[(M,)](
|
| 157 |
+
x,
|
| 158 |
+
x_q,
|
| 159 |
+
x_s,
|
| 160 |
+
group_size,
|
| 161 |
+
N,
|
| 162 |
+
eps,
|
| 163 |
+
fp8_min=fp8_min,
|
| 164 |
+
fp8_max=fp8_max,
|
| 165 |
+
BLOCK=BLOCK,
|
| 166 |
+
num_warps=num_warps,
|
| 167 |
+
num_stages=num_stages,
|
| 168 |
+
)
|
| 169 |
+
|
| 170 |
+
return x_q, x_s
|
quantizer/kernels/python/quantizers/fp8_per_token_group_torch.py
ADDED
|
@@ -0,0 +1,163 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
"""
|
| 16 |
+
Pure PyTorch implementation of per-token-group FP8 quantization.
|
| 17 |
+
|
| 18 |
+
This module provides CPU/Windows-compatible implementations that mirror
|
| 19 |
+
the Triton kernels for FP8 per-token-group quantization.
|
| 20 |
+
"""
|
| 21 |
+
|
| 22 |
+
from typing import Tuple
|
| 23 |
+
|
| 24 |
+
import torch
|
| 25 |
+
|
| 26 |
+
|
| 27 |
+
def fp8_per_token_group_quant_torch(
|
| 28 |
+
x: torch.Tensor,
|
| 29 |
+
group_size: int,
|
| 30 |
+
eps: float = 1e-10,
|
| 31 |
+
dtype: torch.dtype = torch.float8_e4m3fn,
|
| 32 |
+
column_major_scales: bool = False,
|
| 33 |
+
scale_tma_aligned: bool = False,
|
| 34 |
+
) -> Tuple[torch.Tensor, torch.Tensor]:
|
| 35 |
+
"""
|
| 36 |
+
Pure PyTorch implementation of per-token-group FP8 quantization.
|
| 37 |
+
|
| 38 |
+
Function to perform per-token-group quantization on an input tensor `x`.
|
| 39 |
+
Each group of `group_size` elements along the last dimension is quantized
|
| 40 |
+
with its own scale factor.
|
| 41 |
+
|
| 42 |
+
Args:
|
| 43 |
+
x: Input tensor, must be contiguous
|
| 44 |
+
group_size: Size of each quantization group
|
| 45 |
+
eps: Small value to avoid division by zero (default: 1e-10)
|
| 46 |
+
dtype: Target FP8 dtype (default: float8_e4m3fn)
|
| 47 |
+
column_major_scales: Whether to use column-major scale layout (default: False)
|
| 48 |
+
scale_tma_aligned: Whether to use TMA-aligned scales (default: False)
|
| 49 |
+
|
| 50 |
+
Returns:
|
| 51 |
+
Tuple of (quantized_tensor, scale_tensor)
|
| 52 |
+
"""
|
| 53 |
+
assert (
|
| 54 |
+
x.shape[-1] % group_size == 0
|
| 55 |
+
), f"the last dimension ({x.shape[-1]}) cannot be divisible by group_size ({group_size})"
|
| 56 |
+
assert x.is_contiguous(), "`x` is not contiguous"
|
| 57 |
+
|
| 58 |
+
finfo = torch.finfo(dtype)
|
| 59 |
+
fp8_max = finfo.max
|
| 60 |
+
fp8_min = -fp8_max
|
| 61 |
+
|
| 62 |
+
device = x.device
|
| 63 |
+
orig_shape = x.shape
|
| 64 |
+
num_groups = orig_shape[-1] // group_size
|
| 65 |
+
|
| 66 |
+
# Reshape for group-wise processing: [..., N] -> [..., num_groups, group_size]
|
| 67 |
+
x_grouped = x.view(*orig_shape[:-1], num_groups, group_size)
|
| 68 |
+
|
| 69 |
+
# Convert to float32 for computation
|
| 70 |
+
x_float = x_grouped.to(torch.float32)
|
| 71 |
+
|
| 72 |
+
# Compute per-group max absolute values
|
| 73 |
+
absmax = x_float.abs().amax(dim=-1, keepdim=True) # [..., num_groups, 1]
|
| 74 |
+
absmax = torch.maximum(absmax, torch.tensor(eps, device=device, dtype=torch.float32))
|
| 75 |
+
|
| 76 |
+
# Compute scales
|
| 77 |
+
scales = absmax / fp8_max # [..., num_groups, 1]
|
| 78 |
+
|
| 79 |
+
# Quantize
|
| 80 |
+
x_scaled = x_float / scales
|
| 81 |
+
x_q = x_scaled.clamp(fp8_min, fp8_max).to(dtype)
|
| 82 |
+
|
| 83 |
+
# Reshape back to original shape
|
| 84 |
+
x_q = x_q.view(orig_shape)
|
| 85 |
+
|
| 86 |
+
# Prepare scale tensor based on output format
|
| 87 |
+
scales = scales.squeeze(-1) # [..., num_groups]
|
| 88 |
+
|
| 89 |
+
if column_major_scales:
|
| 90 |
+
if scale_tma_aligned:
|
| 91 |
+
# TMA-aligned column-major scales
|
| 92 |
+
aligned_size = (orig_shape[-2] + 3) // 4 * 4
|
| 93 |
+
x_s = torch.empty(
|
| 94 |
+
orig_shape[:-2] + (num_groups, aligned_size),
|
| 95 |
+
device=device,
|
| 96 |
+
dtype=torch.float32,
|
| 97 |
+
)
|
| 98 |
+
# Transpose and fill
|
| 99 |
+
scales_transposed = scales.permute(*range(len(orig_shape) - 2), -1, -2).contiguous()
|
| 100 |
+
x_s[..., : orig_shape[-2]] = scales_transposed
|
| 101 |
+
x_s = x_s.permute(-1, -2)[: orig_shape[-2], :]
|
| 102 |
+
else:
|
| 103 |
+
# Standard column-major scales
|
| 104 |
+
x_s = scales.permute(*range(len(orig_shape) - 2), -1, -2).contiguous()
|
| 105 |
+
else:
|
| 106 |
+
# Row-major scales (default)
|
| 107 |
+
x_s = scales
|
| 108 |
+
|
| 109 |
+
return x_q, x_s.float()
|
| 110 |
+
|
| 111 |
+
|
| 112 |
+
def fp8_per_token_group_quant_torch_simple(
|
| 113 |
+
x: torch.Tensor,
|
| 114 |
+
group_size: int,
|
| 115 |
+
eps: float = 1e-10,
|
| 116 |
+
dtype: torch.dtype = torch.float8_e4m3fn,
|
| 117 |
+
) -> Tuple[torch.Tensor, torch.Tensor]:
|
| 118 |
+
"""
|
| 119 |
+
Simplified PyTorch implementation for row-major scales only.
|
| 120 |
+
|
| 121 |
+
This version is more straightforward and easier to understand.
|
| 122 |
+
|
| 123 |
+
Args:
|
| 124 |
+
x: Input tensor, must be contiguous
|
| 125 |
+
group_size: Size of each quantization group
|
| 126 |
+
eps: Small value to avoid division by zero
|
| 127 |
+
dtype: Target FP8 dtype
|
| 128 |
+
|
| 129 |
+
Returns:
|
| 130 |
+
Tuple of (quantized_tensor, scale_tensor)
|
| 131 |
+
"""
|
| 132 |
+
assert x.shape[-1] % group_size == 0
|
| 133 |
+
assert x.is_contiguous()
|
| 134 |
+
|
| 135 |
+
finfo = torch.finfo(dtype)
|
| 136 |
+
fp8_max = finfo.max
|
| 137 |
+
fp8_min = -fp8_max
|
| 138 |
+
|
| 139 |
+
orig_shape = x.shape
|
| 140 |
+
num_groups = orig_shape[-1] // group_size
|
| 141 |
+
|
| 142 |
+
# Flatten all dimensions except last, then reshape for groups
|
| 143 |
+
x_flat = x.view(-1, orig_shape[-1]) # [M, N]
|
| 144 |
+
M = x_flat.shape[0]
|
| 145 |
+
|
| 146 |
+
# Reshape to [M, num_groups, group_size]
|
| 147 |
+
x_grouped = x_flat.view(M, num_groups, group_size).float()
|
| 148 |
+
|
| 149 |
+
# Compute per-group absmax
|
| 150 |
+
absmax = x_grouped.abs().amax(dim=-1, keepdim=True) # [M, num_groups, 1]
|
| 151 |
+
absmax = torch.clamp(absmax, min=eps)
|
| 152 |
+
|
| 153 |
+
# Compute scales
|
| 154 |
+
scales = absmax / fp8_max
|
| 155 |
+
|
| 156 |
+
# Quantize
|
| 157 |
+
x_q = (x_grouped / scales).clamp(fp8_min, fp8_max).to(dtype)
|
| 158 |
+
|
| 159 |
+
# Reshape back
|
| 160 |
+
x_q = x_q.view(orig_shape)
|
| 161 |
+
scales = scales.squeeze(-1).view(*orig_shape[:-1], num_groups)
|
| 162 |
+
|
| 163 |
+
return x_q, scales.float()
|
quantizer/kernels/python/sgl/__init__.py
ADDED
|
File without changes
|
quantizer/kernels/python/sgl/int8_kernel.py
ADDED
|
@@ -0,0 +1,465 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
import functools
|
| 2 |
+
import json
|
| 3 |
+
import logging
|
| 4 |
+
import os
|
| 5 |
+
from typing import Any, Dict, List, Optional, Tuple
|
| 6 |
+
from functools import lru_cache, partial
|
| 7 |
+
|
| 8 |
+
import torch
|
| 9 |
+
import triton
|
| 10 |
+
import triton.language as tl
|
| 11 |
+
|
| 12 |
+
@lru_cache(maxsize=1)
|
| 13 |
+
def is_musa() -> bool:
|
| 14 |
+
try:
|
| 15 |
+
import torchada # noqa: F401
|
| 16 |
+
except ImportError:
|
| 17 |
+
return False
|
| 18 |
+
return hasattr(torch.version, "musa") and torch.version.musa is not None
|
| 19 |
+
|
| 20 |
+
def get_device_name(device_id: int = 0) -> str:
|
| 21 |
+
if (hasattr(torch, "cuda") and torch.cuda.is_available()) or is_musa():
|
| 22 |
+
return torch.cuda.get_device_name(device_id)
|
| 23 |
+
|
| 24 |
+
if hasattr(torch, "xpu") and torch.xpu.is_available():
|
| 25 |
+
return torch.xpu.get_device_name(device_id)
|
| 26 |
+
|
| 27 |
+
if hasattr(torch, "hpu") and torch.hpu.is_available():
|
| 28 |
+
return torch.hpu.get_device_name(device_id)
|
| 29 |
+
|
| 30 |
+
if hasattr(torch, "npu") and torch.npu.is_available():
|
| 31 |
+
return torch.npu.get_device_name(device_id)
|
| 32 |
+
|
| 33 |
+
@lru_cache(maxsize=1)
|
| 34 |
+
def is_cuda():
|
| 35 |
+
return torch.cuda.is_available() and torch.version.cuda is not None
|
| 36 |
+
|
| 37 |
+
_is_cuda = is_cuda()
|
| 38 |
+
if _is_cuda:
|
| 39 |
+
# Temporary
|
| 40 |
+
try:
|
| 41 |
+
from sgl_kernel import sgl_per_token_group_quant_8bit
|
| 42 |
+
|
| 43 |
+
enable_sgl_per_token_group_quant_8bit = True
|
| 44 |
+
except ImportError:
|
| 45 |
+
from sgl_kernel import sgl_per_token_group_quant_int8
|
| 46 |
+
|
| 47 |
+
enable_sgl_per_token_group_quant_8bit = False
|
| 48 |
+
|
| 49 |
+
logger = logging.getLogger(__name__)
|
| 50 |
+
|
| 51 |
+
|
| 52 |
+
@triton.jit
|
| 53 |
+
def _per_token_quant_int8(
|
| 54 |
+
x_ptr,
|
| 55 |
+
xq_ptr,
|
| 56 |
+
scale_ptr,
|
| 57 |
+
x_sum_ptr,
|
| 58 |
+
stride_x,
|
| 59 |
+
stride_xq,
|
| 60 |
+
N,
|
| 61 |
+
CAL_SUM: tl.constexpr,
|
| 62 |
+
BLOCK: tl.constexpr,
|
| 63 |
+
):
|
| 64 |
+
# Adapted from https://github.com/InternLM/lmdeploy/blob/086481ed84b59bee3b8e4274e5fc69620040c048/lmdeploy/pytorch/kernels/cuda/w8a8_triton_kernels.py#L282
|
| 65 |
+
row_id = tl.program_id(0)
|
| 66 |
+
|
| 67 |
+
cols = tl.arange(0, BLOCK)
|
| 68 |
+
mask = cols < N
|
| 69 |
+
|
| 70 |
+
x = tl.load(x_ptr + row_id * stride_x + cols, mask=mask, other=0.0).to(tl.float32)
|
| 71 |
+
absmax = tl.maximum(tl.max(tl.abs(x)), 1e-10)
|
| 72 |
+
scale_x = absmax / 127
|
| 73 |
+
x_q = x * (127 / absmax)
|
| 74 |
+
x_q = tl.extra.cuda.libdevice.round(x_q).to(tl.int8)
|
| 75 |
+
if CAL_SUM:
|
| 76 |
+
x_sum = tl.sum(x, axis=0)
|
| 77 |
+
tl.store(x_sum_ptr + row_id, x_sum.to(x_sum_ptr.dtype.element_ty))
|
| 78 |
+
|
| 79 |
+
tl.store(xq_ptr + row_id * stride_xq + cols, x_q, mask=mask)
|
| 80 |
+
tl.store(scale_ptr + row_id, scale_x.to(scale_ptr.dtype.element_ty))
|
| 81 |
+
|
| 82 |
+
|
| 83 |
+
def per_token_quant_int8(x, scale_dtype=torch.float32, cal_sum=False):
|
| 84 |
+
M = x.numel() // x.shape[-1]
|
| 85 |
+
N = x.shape[-1]
|
| 86 |
+
x_q = torch.empty_like(x, device=x.device, dtype=torch.int8)
|
| 87 |
+
scales = torch.empty(x.shape[:-1] + (1,), device=x.device, dtype=scale_dtype)
|
| 88 |
+
if cal_sum:
|
| 89 |
+
x_sum = torch.empty(x.shape[:-1], device=x.device, dtype=x.dtype)
|
| 90 |
+
else:
|
| 91 |
+
x_sum = None
|
| 92 |
+
BLOCK = triton.next_power_of_2(N)
|
| 93 |
+
# heuristics for number of warps
|
| 94 |
+
num_warps = min(max(BLOCK // 256, 1), 8)
|
| 95 |
+
|
| 96 |
+
assert x.is_contiguous()
|
| 97 |
+
_per_token_quant_int8[(M,)](
|
| 98 |
+
x,
|
| 99 |
+
x_q,
|
| 100 |
+
scales,
|
| 101 |
+
x_sum,
|
| 102 |
+
stride_x=x.stride(-2),
|
| 103 |
+
stride_xq=x_q.stride(-2),
|
| 104 |
+
N=N,
|
| 105 |
+
CAL_SUM=cal_sum,
|
| 106 |
+
BLOCK=BLOCK,
|
| 107 |
+
num_warps=num_warps,
|
| 108 |
+
num_stages=1,
|
| 109 |
+
)
|
| 110 |
+
if cal_sum:
|
| 111 |
+
return x_q, scales, x_sum
|
| 112 |
+
else:
|
| 113 |
+
return x_q, scales
|
| 114 |
+
|
| 115 |
+
|
| 116 |
+
@triton.jit
|
| 117 |
+
def _per_token_group_quant_int8(
|
| 118 |
+
# Pointers to inputs and output
|
| 119 |
+
y_ptr,
|
| 120 |
+
y_q_ptr,
|
| 121 |
+
y_s_ptr,
|
| 122 |
+
# Stride of input
|
| 123 |
+
y_stride,
|
| 124 |
+
# Columns of input
|
| 125 |
+
N,
|
| 126 |
+
# Avoid to divide zero
|
| 127 |
+
eps,
|
| 128 |
+
# Information for int8
|
| 129 |
+
int8_min,
|
| 130 |
+
int8_max,
|
| 131 |
+
# Meta-parameters
|
| 132 |
+
BLOCK: tl.constexpr,
|
| 133 |
+
):
|
| 134 |
+
"""A Triton-accelerated function to perform per-token-group quantization on a
|
| 135 |
+
tensor.
|
| 136 |
+
|
| 137 |
+
This function converts the tensor values into int8 values.
|
| 138 |
+
"""
|
| 139 |
+
# Map the program id to the row of X and Y it should compute.
|
| 140 |
+
g_id = tl.program_id(0)
|
| 141 |
+
y_ptr += g_id * y_stride
|
| 142 |
+
y_q_ptr += g_id * y_stride
|
| 143 |
+
y_s_ptr += g_id
|
| 144 |
+
|
| 145 |
+
cols = tl.arange(0, BLOCK) # N <= BLOCK
|
| 146 |
+
mask = cols < N
|
| 147 |
+
|
| 148 |
+
y = tl.load(y_ptr + cols, mask=mask, other=0.0).to(tl.float32)
|
| 149 |
+
# Quant
|
| 150 |
+
_absmax = tl.maximum(tl.max(tl.abs(y)), eps)
|
| 151 |
+
y_s = _absmax / int8_max
|
| 152 |
+
y_q = tl.clamp(y / y_s, int8_min, int8_max).to(y_q_ptr.dtype.element_ty)
|
| 153 |
+
|
| 154 |
+
tl.store(y_q_ptr + cols, y_q, mask=mask)
|
| 155 |
+
tl.store(y_s_ptr, y_s)
|
| 156 |
+
|
| 157 |
+
|
| 158 |
+
def per_token_group_quant_int8(
|
| 159 |
+
x: torch.Tensor,
|
| 160 |
+
group_size: int,
|
| 161 |
+
eps: float = 1e-10,
|
| 162 |
+
dtype: torch.dtype = torch.int8,
|
| 163 |
+
) -> Tuple[torch.Tensor, torch.Tensor]:
|
| 164 |
+
"""Function to perform per-token-group quantization on an input tensor `x`.
|
| 165 |
+
|
| 166 |
+
It converts the tensor values into signed int8 values and returns the
|
| 167 |
+
quantized tensor along with the scaling factor used for quantization.
|
| 168 |
+
|
| 169 |
+
Args:
|
| 170 |
+
x: The input tensor with ndim >= 2.
|
| 171 |
+
group_size: The group size used for quantization.
|
| 172 |
+
eps: The minimum to avoid dividing zero.
|
| 173 |
+
dtype: The dype of output tensor. Note that only `torch.int8` is supported for now.
|
| 174 |
+
|
| 175 |
+
Returns:
|
| 176 |
+
Tuple[torch.Tensor, torch.Tensor]: The quantized tensor and the scaling factor for quantization.
|
| 177 |
+
"""
|
| 178 |
+
assert (
|
| 179 |
+
x.shape[-1] % group_size == 0
|
| 180 |
+
), "the last dimension of `x` cannot be divisible by `group_size`"
|
| 181 |
+
assert x.is_contiguous(), "`x` is not contiguous"
|
| 182 |
+
|
| 183 |
+
iinfo = torch.iinfo(dtype)
|
| 184 |
+
int8_max = iinfo.max
|
| 185 |
+
int8_min = iinfo.min
|
| 186 |
+
|
| 187 |
+
x_q = torch.empty_like(x, device=x.device, dtype=dtype)
|
| 188 |
+
M = x.numel() // group_size
|
| 189 |
+
N = group_size
|
| 190 |
+
x_s = torch.empty(
|
| 191 |
+
x.shape[:-1] + (x.shape[-1] // group_size,),
|
| 192 |
+
device=x.device,
|
| 193 |
+
dtype=torch.float32,
|
| 194 |
+
)
|
| 195 |
+
|
| 196 |
+
BLOCK = triton.next_power_of_2(N)
|
| 197 |
+
# heuristics for number of warps
|
| 198 |
+
num_warps = min(max(BLOCK // 256, 1), 8)
|
| 199 |
+
num_stages = 1
|
| 200 |
+
_per_token_group_quant_int8[(M,)](
|
| 201 |
+
x,
|
| 202 |
+
x_q,
|
| 203 |
+
x_s,
|
| 204 |
+
group_size,
|
| 205 |
+
N,
|
| 206 |
+
eps,
|
| 207 |
+
int8_min=int8_min,
|
| 208 |
+
int8_max=int8_max,
|
| 209 |
+
BLOCK=BLOCK,
|
| 210 |
+
num_warps=num_warps,
|
| 211 |
+
num_stages=num_stages,
|
| 212 |
+
)
|
| 213 |
+
|
| 214 |
+
return x_q, x_s
|
| 215 |
+
|
| 216 |
+
|
| 217 |
+
def sglang_per_token_group_quant_int8(
|
| 218 |
+
x: torch.Tensor,
|
| 219 |
+
group_size: int,
|
| 220 |
+
eps: float = 1e-10,
|
| 221 |
+
dtype: torch.dtype = torch.int8,
|
| 222 |
+
enable_v2: Optional[bool] = None,
|
| 223 |
+
):
|
| 224 |
+
assert (
|
| 225 |
+
x.shape[-1] % group_size == 0
|
| 226 |
+
), "the last dimension of `x` cannot be divisible by `group_size`"
|
| 227 |
+
assert x.is_contiguous(), "`x` is not contiguous"
|
| 228 |
+
|
| 229 |
+
iinfo = torch.iinfo(dtype)
|
| 230 |
+
int8_max = iinfo.max
|
| 231 |
+
int8_min = iinfo.min
|
| 232 |
+
|
| 233 |
+
x_q = torch.empty_like(x, device=x.device, dtype=dtype)
|
| 234 |
+
x_s = torch.empty(
|
| 235 |
+
x.shape[:-1] + (x.shape[-1] // group_size,),
|
| 236 |
+
device=x.device,
|
| 237 |
+
dtype=torch.float32,
|
| 238 |
+
)
|
| 239 |
+
|
| 240 |
+
# Temporary
|
| 241 |
+
if enable_sgl_per_token_group_quant_8bit:
|
| 242 |
+
sgl_per_token_group_quant_8bit(
|
| 243 |
+
x, x_q, x_s, group_size, eps, int8_min, int8_max, enable_v2=enable_v2
|
| 244 |
+
)
|
| 245 |
+
else:
|
| 246 |
+
assert not enable_v2
|
| 247 |
+
sgl_per_token_group_quant_int8(x, x_q, x_s, group_size, eps, int8_min, int8_max)
|
| 248 |
+
|
| 249 |
+
return x_q, x_s
|
| 250 |
+
|
| 251 |
+
|
| 252 |
+
@triton.jit
|
| 253 |
+
def _w8a8_block_int8_matmul(
|
| 254 |
+
# Pointers to inputs and output
|
| 255 |
+
A,
|
| 256 |
+
B,
|
| 257 |
+
C,
|
| 258 |
+
As,
|
| 259 |
+
Bs,
|
| 260 |
+
# Shape for matmul
|
| 261 |
+
M,
|
| 262 |
+
N,
|
| 263 |
+
K,
|
| 264 |
+
# Block size for block-wise quantization
|
| 265 |
+
group_n,
|
| 266 |
+
group_k,
|
| 267 |
+
# Stride for inputs and output
|
| 268 |
+
stride_am,
|
| 269 |
+
stride_ak,
|
| 270 |
+
stride_bk,
|
| 271 |
+
stride_bn,
|
| 272 |
+
stride_cm,
|
| 273 |
+
stride_cn,
|
| 274 |
+
stride_As_m,
|
| 275 |
+
stride_As_k,
|
| 276 |
+
stride_Bs_k,
|
| 277 |
+
stride_Bs_n,
|
| 278 |
+
# Meta-parameters
|
| 279 |
+
BLOCK_SIZE_M: tl.constexpr,
|
| 280 |
+
BLOCK_SIZE_N: tl.constexpr,
|
| 281 |
+
BLOCK_SIZE_K: tl.constexpr,
|
| 282 |
+
GROUP_SIZE_M: tl.constexpr,
|
| 283 |
+
):
|
| 284 |
+
"""Triton-accelerated function used to perform linear operations (dot
|
| 285 |
+
product) on input tensors `A` and `B` with block-wise quantization, and store the result in output
|
| 286 |
+
tensor `C`.
|
| 287 |
+
"""
|
| 288 |
+
|
| 289 |
+
pid = tl.program_id(axis=0)
|
| 290 |
+
num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)
|
| 291 |
+
num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)
|
| 292 |
+
num_pid_in_group = GROUP_SIZE_M * num_pid_n
|
| 293 |
+
group_id = pid // num_pid_in_group
|
| 294 |
+
first_pid_m = group_id * GROUP_SIZE_M
|
| 295 |
+
group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)
|
| 296 |
+
pid_m = first_pid_m + (pid % group_size_m)
|
| 297 |
+
pid_n = (pid % num_pid_in_group) // group_size_m
|
| 298 |
+
|
| 299 |
+
offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M
|
| 300 |
+
offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
|
| 301 |
+
offs_k = tl.arange(0, BLOCK_SIZE_K)
|
| 302 |
+
a_ptrs = A + (offs_am[:, None] * stride_am + offs_k[None, :] * stride_ak)
|
| 303 |
+
b_ptrs = B + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)
|
| 304 |
+
|
| 305 |
+
As_ptrs = As + offs_am * stride_As_m
|
| 306 |
+
offs_bsn = offs_bn // group_n
|
| 307 |
+
Bs_ptrs = Bs + offs_bsn * stride_Bs_n
|
| 308 |
+
|
| 309 |
+
accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
|
| 310 |
+
for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):
|
| 311 |
+
a = tl.load(a_ptrs, mask=offs_k[None, :] < K - k * BLOCK_SIZE_K, other=0.0)
|
| 312 |
+
b = tl.load(b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_SIZE_K, other=0.0)
|
| 313 |
+
|
| 314 |
+
k_start = k * BLOCK_SIZE_K
|
| 315 |
+
offs_ks = k_start // group_k
|
| 316 |
+
a_s = tl.load(As_ptrs + offs_ks * stride_As_k)
|
| 317 |
+
b_s = tl.load(Bs_ptrs + offs_ks * stride_Bs_k)
|
| 318 |
+
|
| 319 |
+
accumulator += tl.dot(a, b).to(tl.float32) * a_s[:, None] * b_s[None, :]
|
| 320 |
+
a_ptrs += BLOCK_SIZE_K * stride_ak
|
| 321 |
+
b_ptrs += BLOCK_SIZE_K * stride_bk
|
| 322 |
+
|
| 323 |
+
if C.dtype.element_ty == tl.bfloat16:
|
| 324 |
+
c = accumulator.to(tl.bfloat16)
|
| 325 |
+
elif C.dtype.element_ty == tl.float16:
|
| 326 |
+
c = accumulator.to(tl.float16)
|
| 327 |
+
else:
|
| 328 |
+
c = accumulator.to(tl.float32)
|
| 329 |
+
|
| 330 |
+
offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
|
| 331 |
+
offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
|
| 332 |
+
c_ptrs = C + stride_cm * offs_cm[:, None] + stride_cn * offs_cn[None, :]
|
| 333 |
+
c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)
|
| 334 |
+
tl.store(c_ptrs, c, mask=c_mask)
|
| 335 |
+
|
| 336 |
+
|
| 337 |
+
@functools.lru_cache
|
| 338 |
+
def get_w8a8_block_int8_configs(
|
| 339 |
+
N: int, K: int, block_n: int, block_k: int
|
| 340 |
+
) -> Optional[Dict[int, Any]]:
|
| 341 |
+
"""
|
| 342 |
+
Return optimized configurations for the w8a8 block fp8 kernel.
|
| 343 |
+
|
| 344 |
+
The return value will be a dictionary that maps an irregular grid of
|
| 345 |
+
batch sizes to configurations of the w8a8 block fp8 kernel. To evaluate the
|
| 346 |
+
kernel on a given batch size bs, the closest batch size in the grid should
|
| 347 |
+
be picked and the associated configuration chosen to invoke the kernel.
|
| 348 |
+
"""
|
| 349 |
+
|
| 350 |
+
# First look up if an optimized configuration is available in the configs
|
| 351 |
+
# directory
|
| 352 |
+
device_name = get_device_name().replace(" ", "_")
|
| 353 |
+
json_file_name = f"N={N},K={K},device_name={device_name},dtype=int8_w8a8,block_shape=[{block_n}, {block_k}].json"
|
| 354 |
+
|
| 355 |
+
config_file_path = os.path.join(
|
| 356 |
+
os.path.dirname(os.path.realpath(__file__)), "configs", json_file_name
|
| 357 |
+
)
|
| 358 |
+
if os.path.exists(config_file_path):
|
| 359 |
+
with open(config_file_path) as f:
|
| 360 |
+
logger.info(
|
| 361 |
+
"Using configuration from %s for W8A8 Block INT8 kernel.",
|
| 362 |
+
config_file_path,
|
| 363 |
+
)
|
| 364 |
+
# If a configuration has been found, return it
|
| 365 |
+
return {int(key): val for key, val in json.load(f).items()}
|
| 366 |
+
|
| 367 |
+
# If no optimized configuration is available, we will use the default
|
| 368 |
+
# configuration
|
| 369 |
+
logger.warning(
|
| 370 |
+
(
|
| 371 |
+
"Using default W8A8 Block INT8 kernel config. Performance might be sub-optimal! "
|
| 372 |
+
"Config file not found at %s"
|
| 373 |
+
),
|
| 374 |
+
config_file_path,
|
| 375 |
+
)
|
| 376 |
+
return None
|
| 377 |
+
|
| 378 |
+
|
| 379 |
+
def w8a8_block_int8_matmul(
|
| 380 |
+
A: torch.Tensor,
|
| 381 |
+
B: torch.Tensor,
|
| 382 |
+
As: torch.Tensor,
|
| 383 |
+
Bs: torch.Tensor,
|
| 384 |
+
block_size: List[int],
|
| 385 |
+
output_dtype: torch.dtype = torch.float16,
|
| 386 |
+
) -> torch.Tensor:
|
| 387 |
+
"""This function performs matrix multiplication with block-wise quantization.
|
| 388 |
+
|
| 389 |
+
It takes two input tensors `A` and `B` with scales `As` and `Bs`.
|
| 390 |
+
The output is returned in the specified `output_dtype`.
|
| 391 |
+
|
| 392 |
+
Args:
|
| 393 |
+
A: The input tensor, e.g., activation.
|
| 394 |
+
B: The input tensor, e.g., weight.
|
| 395 |
+
As: The per-token-group quantization scale for `A`.
|
| 396 |
+
Bs: The per-block quantization scale for `B`.
|
| 397 |
+
block_size: The block size for per-block quantization. It should be 2-dim, e.g., [128, 128].
|
| 398 |
+
output_dytpe: The dtype of the returned tensor.
|
| 399 |
+
|
| 400 |
+
Returns:
|
| 401 |
+
torch.Tensor: The result of matmul.
|
| 402 |
+
"""
|
| 403 |
+
assert len(block_size) == 2
|
| 404 |
+
block_n, block_k = block_size[0], block_size[1]
|
| 405 |
+
|
| 406 |
+
assert A.shape[-1] == B.shape[-1]
|
| 407 |
+
assert A.shape[:-1] == As.shape[:-1] and A.is_contiguous()
|
| 408 |
+
assert triton.cdiv(A.shape[-1], block_k) == As.shape[-1]
|
| 409 |
+
M = A.numel() // A.shape[-1]
|
| 410 |
+
|
| 411 |
+
assert B.ndim == 2 and B.is_contiguous() and Bs.ndim == 2
|
| 412 |
+
N, K = B.shape
|
| 413 |
+
assert triton.cdiv(N, block_n) == Bs.shape[0]
|
| 414 |
+
assert triton.cdiv(K, block_k) == Bs.shape[1]
|
| 415 |
+
|
| 416 |
+
C_shape = A.shape[:-1] + (N,)
|
| 417 |
+
C = A.new_empty(C_shape, dtype=output_dtype)
|
| 418 |
+
|
| 419 |
+
configs = get_w8a8_block_int8_configs(N, K, block_size[0], block_size[1])
|
| 420 |
+
if configs:
|
| 421 |
+
# If an optimal configuration map has been found, look up the
|
| 422 |
+
# optimal config
|
| 423 |
+
config = configs[min(configs.keys(), key=lambda x: abs(x - M))]
|
| 424 |
+
else:
|
| 425 |
+
# Default config
|
| 426 |
+
# Block-wise quant: BLOCK_SIZE_K must be divisible by block_size[1]
|
| 427 |
+
config = {
|
| 428 |
+
"BLOCK_SIZE_M": 64,
|
| 429 |
+
"BLOCK_SIZE_N": block_size[0],
|
| 430 |
+
"BLOCK_SIZE_K": block_size[1],
|
| 431 |
+
"GROUP_SIZE_M": 32,
|
| 432 |
+
"num_warps": 4,
|
| 433 |
+
"num_stages": 3,
|
| 434 |
+
}
|
| 435 |
+
|
| 436 |
+
def grid(META):
|
| 437 |
+
return (
|
| 438 |
+
triton.cdiv(M, META["BLOCK_SIZE_M"]) * triton.cdiv(N, META["BLOCK_SIZE_N"]),
|
| 439 |
+
)
|
| 440 |
+
|
| 441 |
+
_w8a8_block_int8_matmul[grid](
|
| 442 |
+
A,
|
| 443 |
+
B,
|
| 444 |
+
C,
|
| 445 |
+
As,
|
| 446 |
+
Bs,
|
| 447 |
+
M,
|
| 448 |
+
N,
|
| 449 |
+
K,
|
| 450 |
+
block_n,
|
| 451 |
+
block_k,
|
| 452 |
+
A.stride(-2),
|
| 453 |
+
A.stride(-1),
|
| 454 |
+
B.stride(1),
|
| 455 |
+
B.stride(0),
|
| 456 |
+
C.stride(-2),
|
| 457 |
+
C.stride(-1),
|
| 458 |
+
As.stride(-2),
|
| 459 |
+
As.stride(-1),
|
| 460 |
+
Bs.stride(1),
|
| 461 |
+
Bs.stride(0),
|
| 462 |
+
**config,
|
| 463 |
+
)
|
| 464 |
+
|
| 465 |
+
return C
|
quantizer/quant/__init__.py
ADDED
|
@@ -0,0 +1,19 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
from .ptq import DynamicDiTQuantizer
|
| 16 |
+
|
| 17 |
+
__all__ = [
|
| 18 |
+
"DynamicDiTQuantizer",
|
| 19 |
+
]
|
quantizer/quant/modules/__init__.py
ADDED
|
@@ -0,0 +1,17 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
from .linear import FP4DynamicLinear, FP8DynamicLinear, FP8WeightOnlyLinear, INT8DynamicLinear
|
| 16 |
+
|
| 17 |
+
__all__ = ["FP8DynamicLinear", "FP8WeightOnlyLinear", "FP4DynamicLinear", "INT8DynamicLinear"]
|
quantizer/quant/modules/linear.py
ADDED
|
@@ -0,0 +1,589 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
import os
|
| 16 |
+
import time
|
| 17 |
+
|
| 18 |
+
import torch
|
| 19 |
+
from lightx2v_kernel.gemm import (
|
| 20 |
+
cutlass_scaled_mxfp4_mm,
|
| 21 |
+
cutlass_scaled_mxfp6_mxfp8_mm,
|
| 22 |
+
cutlass_scaled_mxfp8_mm,
|
| 23 |
+
cutlass_scaled_nvfp4_mm,
|
| 24 |
+
)
|
| 25 |
+
|
| 26 |
+
try:
|
| 27 |
+
from torchao.quantization.utils import quant_int8_per_token_matmul as torchao_int8_gemm
|
| 28 |
+
from torchao.quantization.utils import quantize_activation_per_token_absmax as torchao_int8_quant
|
| 29 |
+
except ImportError:
|
| 30 |
+
try:
|
| 31 |
+
from torchao.quantization.utils import _quant_int8_per_token_matmul as torchao_int8_gemm
|
| 32 |
+
from torchao.quantization.utils import _quantize_activation_per_token_absmax as torchao_int8_quant
|
| 33 |
+
except ImportError:
|
| 34 |
+
torchao_int8_gemm, torchao_int8_quant = None, None
|
| 35 |
+
|
| 36 |
+
try:
|
| 37 |
+
from vllm import _custom_ops as vllm_ops
|
| 38 |
+
except ImportError:
|
| 39 |
+
vllm_ops = None
|
| 40 |
+
|
| 41 |
+
try:
|
| 42 |
+
from ...kernels.python.sgl.int8_kernel import per_token_quant_int8 as sglang_int8_act_quant
|
| 43 |
+
except ImportError:
|
| 44 |
+
sglang_int8_act_quant = None
|
| 45 |
+
|
| 46 |
+
try:
|
| 47 |
+
import sgl_kernel
|
| 48 |
+
except ImportError:
|
| 49 |
+
sgl_kernel = None
|
| 50 |
+
|
| 51 |
+
try:
|
| 52 |
+
from q8_kernels.functional.linear import q8_linear
|
| 53 |
+
except ImportError:
|
| 54 |
+
q8_linear = None
|
| 55 |
+
|
| 56 |
+
try:
|
| 57 |
+
from ...kernels.python.mm.triton_kernels import (
|
| 58 |
+
int8_gemm_bias_triton,
|
| 59 |
+
int8_gemm_triton,
|
| 60 |
+
int8_quantize_triton,
|
| 61 |
+
)
|
| 62 |
+
except ImportError:
|
| 63 |
+
int8_gemm_bias_triton, int8_gemm_triton, int8_quantize_triton = None, None, None
|
| 64 |
+
|
| 65 |
+
from ..quant_func import (
|
| 66 |
+
fp8_gemm,
|
| 67 |
+
fp8_per_block_quant,
|
| 68 |
+
fp8_per_tensor_quant,
|
| 69 |
+
fp8_per_token_group_quant,
|
| 70 |
+
fp8_per_token_quant_sgl,
|
| 71 |
+
fp8_weight_only_gemm,
|
| 72 |
+
mxfp4_per_tensor_quant,
|
| 73 |
+
mxfp6_per_tensor_quant,
|
| 74 |
+
mxfp8_per_tensor_quant,
|
| 75 |
+
nvfp4_per_tensor_quant,
|
| 76 |
+
)
|
| 77 |
+
|
| 78 |
+
|
| 79 |
+
# modified from https://github.com/neuralmagic/AutoFP8/blob/main/auto_fp8/quantize.py
|
| 80 |
+
class FP8DynamicLinear(torch.nn.Module):
|
| 81 |
+
def __init__(
|
| 82 |
+
self,
|
| 83 |
+
weight: torch.Tensor,
|
| 84 |
+
weight_scale: torch.Tensor,
|
| 85 |
+
bias: torch.nn.Parameter,
|
| 86 |
+
native_fp8_support: bool = False,
|
| 87 |
+
quant_type: str = "fp8-per-tensor",
|
| 88 |
+
block_size: int = 128,
|
| 89 |
+
):
|
| 90 |
+
super().__init__()
|
| 91 |
+
self.weight = torch.nn.Parameter(weight, requires_grad=False)
|
| 92 |
+
self.weight_scale = torch.nn.Parameter(weight_scale, requires_grad=False)
|
| 93 |
+
self.bias = bias
|
| 94 |
+
self.native_fp8_support = native_fp8_support
|
| 95 |
+
self.quant_type = quant_type
|
| 96 |
+
self.block_size = block_size
|
| 97 |
+
self.profile_enabled = os.environ.get("ANGELSLIM_FP8_PROFILE", "0") == "1"
|
| 98 |
+
|
| 99 |
+
@torch.compiler.disable(recursive=True)
|
| 100 |
+
def forward(self, x):
|
| 101 |
+
ori_dtype = x.dtype
|
| 102 |
+
assert ori_dtype in [
|
| 103 |
+
torch.float32,
|
| 104 |
+
torch.bfloat16,
|
| 105 |
+
torch.float16,
|
| 106 |
+
], "x.dtype must be float32, bfloat16, or float16"
|
| 107 |
+
|
| 108 |
+
if ori_dtype == torch.float32:
|
| 109 |
+
x = x.to(torch.bfloat16)
|
| 110 |
+
|
| 111 |
+
if self.profile_enabled and x.is_cuda:
|
| 112 |
+
torch.cuda.synchronize(x.device)
|
| 113 |
+
t0 = time.perf_counter()
|
| 114 |
+
|
| 115 |
+
if self.quant_type == "fp8-per-tensor":
|
| 116 |
+
origin_shape = None
|
| 117 |
+
qinput, x_scale = fp8_per_tensor_quant(x)
|
| 118 |
+
elif self.quant_type == "fp8-per-token":
|
| 119 |
+
origin_shape = None
|
| 120 |
+
x_2d = x.view(-1, x.shape[-1])
|
| 121 |
+
qinput, x_scale = fp8_per_token_group_quant(x_2d, x_2d.shape[-1])
|
| 122 |
+
elif self.quant_type == "fp8-per-token-sgl" and self.native_fp8_support:
|
| 123 |
+
origin_shape = x.shape
|
| 124 |
+
x_2d = x.view(-1, x.shape[-1])
|
| 125 |
+
qinput, x_scale = fp8_per_token_quant_sgl(x_2d)
|
| 126 |
+
elif self.quant_type == "fp8-per-block" and self.native_fp8_support:
|
| 127 |
+
origin_shape = x.shape
|
| 128 |
+
x = x.view(-1, x.shape[-1])
|
| 129 |
+
qinput, x_scale = fp8_per_token_group_quant(
|
| 130 |
+
x, group_size=128, column_major_scales=True, scale_tma_aligned=True
|
| 131 |
+
)
|
| 132 |
+
elif self.quant_type == "fp8-per-block" and not self.native_fp8_support:
|
| 133 |
+
origin_shape = x.shape
|
| 134 |
+
x_2d = x.view(-1, x.shape[-1])
|
| 135 |
+
qinput, x_scale = fp8_per_block_quant(x_2d, block_size=128)
|
| 136 |
+
elif self.quant_type == "fp8-per-channel-vllm":
|
| 137 |
+
if vllm_ops is None:
|
| 138 |
+
raise ImportError(
|
| 139 |
+
"quant_type='fp8-per-channel-vllm' requires vllm._custom_ops, but vllm is not installed"
|
| 140 |
+
)
|
| 141 |
+
origin_shape = x.shape if x.dim() == 3 else None
|
| 142 |
+
x_2d = x.view(-1, x.shape[-1]) if x.dim() == 3 else x
|
| 143 |
+
qinput, x_scale = vllm_ops.scaled_fp8_quant(
|
| 144 |
+
x_2d, None, scale_ub=None, use_per_token_if_dynamic=True
|
| 145 |
+
)
|
| 146 |
+
else:
|
| 147 |
+
raise ValueError(f"Invalid quant_type: {self.quant_type}")
|
| 148 |
+
|
| 149 |
+
if self.profile_enabled and qinput.is_cuda:
|
| 150 |
+
torch.cuda.synchronize(qinput.device)
|
| 151 |
+
t1 = time.perf_counter()
|
| 152 |
+
|
| 153 |
+
output = fp8_gemm(
|
| 154 |
+
A=qinput,
|
| 155 |
+
A_scale=x_scale,
|
| 156 |
+
B=self.weight,
|
| 157 |
+
B_scale=self.weight_scale,
|
| 158 |
+
bias=self.bias,
|
| 159 |
+
out_dtype=x.dtype,
|
| 160 |
+
native_fp8_support=self.native_fp8_support,
|
| 161 |
+
quant_type=self.quant_type,
|
| 162 |
+
origin_shape=origin_shape,
|
| 163 |
+
)
|
| 164 |
+
|
| 165 |
+
if self.profile_enabled and output.is_cuda:
|
| 166 |
+
torch.cuda.synchronize(output.device)
|
| 167 |
+
t2 = time.perf_counter()
|
| 168 |
+
|
| 169 |
+
if self.profile_enabled:
|
| 170 |
+
qshape = tuple(qinput.shape)
|
| 171 |
+
print(
|
| 172 |
+
f"[FP8Linear:{self.quant_type}] quant_ms={(t1 - t0) * 1000:.3f}, "
|
| 173 |
+
f"gemm_ms={(t2 - t1) * 1000:.3f}, qshape={qshape}"
|
| 174 |
+
)
|
| 175 |
+
|
| 176 |
+
if (
|
| 177 |
+
self.quant_type in ["fp8-per-token", "fp8-per-token-sgl"]
|
| 178 |
+
and x.dim() == 3
|
| 179 |
+
and output.dim() == 2
|
| 180 |
+
):
|
| 181 |
+
output = output.unsqueeze(0)
|
| 182 |
+
|
| 183 |
+
# Restore original shape for fp8-per-block with native_fp8_support=False
|
| 184 |
+
# (native_fp8_support=True case is handled in fp8_gemm_deepgemm_block)
|
| 185 |
+
if (
|
| 186 |
+
(
|
| 187 |
+
(self.quant_type == "fp8-per-block" and not self.native_fp8_support)
|
| 188 |
+
or self.quant_type == "fp8-per-channel-vllm"
|
| 189 |
+
)
|
| 190 |
+
and origin_shape is not None
|
| 191 |
+
and len(origin_shape) == 3
|
| 192 |
+
and output.dim() == 2
|
| 193 |
+
):
|
| 194 |
+
output = output.view(origin_shape[0], origin_shape[1], -1)
|
| 195 |
+
|
| 196 |
+
return output
|
| 197 |
+
|
| 198 |
+
|
| 199 |
+
class FP8WeightOnlyLinear(torch.nn.Module):
|
| 200 |
+
"""
|
| 201 |
+
FP8 Weight-Only Quantized Linear Layer.
|
| 202 |
+
|
| 203 |
+
This layer quantizes only the weights to FP8 while keeping activations
|
| 204 |
+
in higher precision (bfloat16/float16). This provides a good balance
|
| 205 |
+
between memory savings and accuracy.
|
| 206 |
+
"""
|
| 207 |
+
|
| 208 |
+
def __init__(
|
| 209 |
+
self,
|
| 210 |
+
weight: torch.Tensor,
|
| 211 |
+
weight_scale: torch.Tensor,
|
| 212 |
+
bias: torch.nn.Parameter,
|
| 213 |
+
native_fp8_support: bool = False, # not used
|
| 214 |
+
quant_type: str = "fp8-per-tensor-weight-only",
|
| 215 |
+
):
|
| 216 |
+
super().__init__()
|
| 217 |
+
self.weight = torch.nn.Parameter(weight, requires_grad=False)
|
| 218 |
+
self.weight_scale = torch.nn.Parameter(weight_scale, requires_grad=False)
|
| 219 |
+
self.bias = bias
|
| 220 |
+
self.native_fp8_support = native_fp8_support # not used
|
| 221 |
+
self.quant_type = quant_type
|
| 222 |
+
|
| 223 |
+
@torch.compiler.disable(recursive=True)
|
| 224 |
+
def forward(self, x):
|
| 225 |
+
ori_dtype = x.dtype
|
| 226 |
+
assert ori_dtype in [
|
| 227 |
+
torch.float32,
|
| 228 |
+
torch.bfloat16,
|
| 229 |
+
torch.float16,
|
| 230 |
+
], "x.dtype must be float32, bfloat16, or float16"
|
| 231 |
+
|
| 232 |
+
if ori_dtype == torch.float32:
|
| 233 |
+
x = x.to(torch.bfloat16)
|
| 234 |
+
|
| 235 |
+
# For weight-only quantization, we don't quantize activations
|
| 236 |
+
# Just use the original activations with quantized weights
|
| 237 |
+
output = fp8_weight_only_gemm(
|
| 238 |
+
A=x, # Keep activations in original precision
|
| 239 |
+
B=self.weight,
|
| 240 |
+
B_scale=self.weight_scale,
|
| 241 |
+
bias=self.bias,
|
| 242 |
+
out_dtype=x.dtype,
|
| 243 |
+
)
|
| 244 |
+
|
| 245 |
+
return output
|
| 246 |
+
|
| 247 |
+
|
| 248 |
+
class INT8DynamicLinear(torch.nn.Module):
|
| 249 |
+
"""
|
| 250 |
+
INT8 weight-only linear layer with per-channel scales.
|
| 251 |
+
"""
|
| 252 |
+
|
| 253 |
+
def __init__(
|
| 254 |
+
self,
|
| 255 |
+
weight: torch.Tensor,
|
| 256 |
+
weight_scale: torch.Tensor,
|
| 257 |
+
bias: torch.nn.Parameter,
|
| 258 |
+
native_fp8_support: bool = False, # not used
|
| 259 |
+
quant_type: str = "int8",
|
| 260 |
+
):
|
| 261 |
+
super().__init__()
|
| 262 |
+
self.weight = torch.nn.Parameter(weight, requires_grad=False)
|
| 263 |
+
self.weight_scale = torch.nn.Parameter(weight_scale, requires_grad=False)
|
| 264 |
+
self.bias = bias
|
| 265 |
+
self.native_fp8_support = native_fp8_support # not used
|
| 266 |
+
self.quant_type = quant_type
|
| 267 |
+
|
| 268 |
+
@staticmethod
|
| 269 |
+
def _is_backend_available(backend: str) -> bool:
|
| 270 |
+
if backend == "torchao":
|
| 271 |
+
return torchao_int8_quant is not None and torchao_int8_gemm is not None
|
| 272 |
+
if backend == "vllm":
|
| 273 |
+
return vllm_ops is not None and hasattr(torch.ops, "_C")
|
| 274 |
+
if backend == "triton":
|
| 275 |
+
return (
|
| 276 |
+
int8_quantize_triton is not None
|
| 277 |
+
and int8_gemm_triton is not None
|
| 278 |
+
and int8_gemm_bias_triton is not None
|
| 279 |
+
)
|
| 280 |
+
if backend == "sgl":
|
| 281 |
+
has_act = (
|
| 282 |
+
sglang_int8_act_quant is not None
|
| 283 |
+
or vllm_ops is not None
|
| 284 |
+
or torchao_int8_quant is not None
|
| 285 |
+
or int8_quantize_triton is not None
|
| 286 |
+
)
|
| 287 |
+
return sgl_kernel is not None and has_act
|
| 288 |
+
if backend == "q8f":
|
| 289 |
+
has_act = (
|
| 290 |
+
vllm_ops is not None
|
| 291 |
+
or torchao_int8_quant is not None
|
| 292 |
+
or int8_quantize_triton is not None
|
| 293 |
+
)
|
| 294 |
+
return q8_linear is not None and has_act
|
| 295 |
+
return False
|
| 296 |
+
|
| 297 |
+
def _resolve_int8_backend(self) -> str:
|
| 298 |
+
explicit_backend = {
|
| 299 |
+
"int8-torchao": "torchao",
|
| 300 |
+
"int8-vllm": "vllm",
|
| 301 |
+
"int8-triton": "triton",
|
| 302 |
+
"int8-sgl": "sgl",
|
| 303 |
+
"int8-q8f": "q8f",
|
| 304 |
+
}
|
| 305 |
+
if self.quant_type in explicit_backend:
|
| 306 |
+
backend = explicit_backend[self.quant_type]
|
| 307 |
+
if not self._is_backend_available(backend):
|
| 308 |
+
raise ImportError(
|
| 309 |
+
f"quant_type='{self.quant_type}' requires '{backend}' backend dependencies"
|
| 310 |
+
)
|
| 311 |
+
return backend
|
| 312 |
+
|
| 313 |
+
# quant_type='int8' uses auto priority: sgl > vllm > torchao > triton
|
| 314 |
+
for backend in ("sgl", "vllm", "torchao", "triton"):
|
| 315 |
+
if self._is_backend_available(backend):
|
| 316 |
+
return backend
|
| 317 |
+
raise ImportError(
|
| 318 |
+
"quant_type='int8' requires one of backends [sgl, vllm, torchao, triton], but none is available"
|
| 319 |
+
)
|
| 320 |
+
|
| 321 |
+
def _act_quant_int8_torchao(self, x_2d: torch.Tensor):
|
| 322 |
+
input_tensor_quant, input_tensor_scale = torchao_int8_quant(x_2d)
|
| 323 |
+
return input_tensor_quant, input_tensor_scale.float()
|
| 324 |
+
|
| 325 |
+
def _act_quant_int8_vllm(self, x_2d: torch.Tensor):
|
| 326 |
+
input_tensor_quant, input_tensor_scale, _ = vllm_ops.scaled_int8_quant(
|
| 327 |
+
x_2d, scale=None, azp=None, symmetric=True
|
| 328 |
+
)
|
| 329 |
+
return input_tensor_quant, input_tensor_scale.float()
|
| 330 |
+
|
| 331 |
+
def _act_quant_int8_triton(self, x_2d: torch.Tensor):
|
| 332 |
+
input_tensor_quant, input_tensor_scale = int8_quantize_triton(x_2d)
|
| 333 |
+
return input_tensor_quant, input_tensor_scale.float()
|
| 334 |
+
|
| 335 |
+
def _act_quant_int8_sgl(self, x_2d: torch.Tensor):
|
| 336 |
+
if sglang_int8_act_quant is not None:
|
| 337 |
+
input_tensor_quant, input_tensor_scale = sglang_int8_act_quant(x_2d)
|
| 338 |
+
return input_tensor_quant, input_tensor_scale.float()
|
| 339 |
+
if vllm_ops is not None:
|
| 340 |
+
return self._act_quant_int8_vllm(x_2d)
|
| 341 |
+
if torchao_int8_quant is not None:
|
| 342 |
+
return self._act_quant_int8_torchao(x_2d)
|
| 343 |
+
if int8_quantize_triton is not None:
|
| 344 |
+
return self._act_quant_int8_triton(x_2d)
|
| 345 |
+
raise ImportError("int8-sgl activation quantization requires sglang/vllm/torchao/triton")
|
| 346 |
+
|
| 347 |
+
def _act_quant_by_backend(self, x_2d: torch.Tensor, backend: str):
|
| 348 |
+
if backend == "torchao":
|
| 349 |
+
return self._act_quant_int8_torchao(x_2d)
|
| 350 |
+
if backend == "vllm":
|
| 351 |
+
return self._act_quant_int8_vllm(x_2d)
|
| 352 |
+
if backend == "triton":
|
| 353 |
+
return self._act_quant_int8_triton(x_2d)
|
| 354 |
+
if backend == "sgl":
|
| 355 |
+
return self._act_quant_int8_sgl(x_2d)
|
| 356 |
+
if backend == "q8f":
|
| 357 |
+
if vllm_ops is not None:
|
| 358 |
+
return self._act_quant_int8_vllm(x_2d)
|
| 359 |
+
if torchao_int8_quant is not None:
|
| 360 |
+
return self._act_quant_int8_torchao(x_2d)
|
| 361 |
+
return self._act_quant_int8_triton(x_2d)
|
| 362 |
+
raise ValueError(f"Unsupported int8 backend: {backend}")
|
| 363 |
+
|
| 364 |
+
def _gemm_int8_torchao(self, qinput, x_scale, out_dtype):
|
| 365 |
+
output = torchao_int8_gemm(
|
| 366 |
+
qinput,
|
| 367 |
+
x_scale,
|
| 368 |
+
self.weight.t(),
|
| 369 |
+
self.weight_scale.t().float(),
|
| 370 |
+
output_dtype=out_dtype,
|
| 371 |
+
)
|
| 372 |
+
if self.bias is not None:
|
| 373 |
+
output.add_(self.bias.to(output.dtype))
|
| 374 |
+
return output
|
| 375 |
+
|
| 376 |
+
def _gemm_int8_vllm(self, qinput, x_scale, out_dtype):
|
| 377 |
+
shape = (qinput.shape[0], self.weight.shape[0])
|
| 378 |
+
output = torch.empty(shape, dtype=out_dtype, device=qinput.device, requires_grad=False)
|
| 379 |
+
torch.ops._C.cutlass_scaled_mm(
|
| 380 |
+
output,
|
| 381 |
+
qinput,
|
| 382 |
+
self.weight.t(),
|
| 383 |
+
x_scale,
|
| 384 |
+
self.weight_scale.t(),
|
| 385 |
+
self.bias,
|
| 386 |
+
)
|
| 387 |
+
return output
|
| 388 |
+
|
| 389 |
+
def _gemm_int8_triton(self, qinput, x_scale, out_dtype):
|
| 390 |
+
if self.bias is not None:
|
| 391 |
+
return int8_gemm_bias_triton(
|
| 392 |
+
qinput,
|
| 393 |
+
self.weight,
|
| 394 |
+
self.bias,
|
| 395 |
+
x_scale,
|
| 396 |
+
self.weight_scale,
|
| 397 |
+
output_dtype=out_dtype,
|
| 398 |
+
)
|
| 399 |
+
return int8_gemm_triton(
|
| 400 |
+
qinput,
|
| 401 |
+
self.weight,
|
| 402 |
+
x_scale,
|
| 403 |
+
self.weight_scale,
|
| 404 |
+
output_dtype=out_dtype,
|
| 405 |
+
)
|
| 406 |
+
|
| 407 |
+
def _gemm_int8_sgl(self, qinput, x_scale, out_dtype):
|
| 408 |
+
return sgl_kernel.int8_scaled_mm(
|
| 409 |
+
qinput,
|
| 410 |
+
self.weight.t(),
|
| 411 |
+
x_scale,
|
| 412 |
+
self.weight_scale.t(),
|
| 413 |
+
out_dtype,
|
| 414 |
+
self.bias,
|
| 415 |
+
)
|
| 416 |
+
|
| 417 |
+
def _gemm_int8_q8f(self, qinput, x_scale, out_dtype):
|
| 418 |
+
bias_fp32 = self.bias.float() if self.bias is not None else None
|
| 419 |
+
return q8_linear(
|
| 420 |
+
qinput,
|
| 421 |
+
self.weight,
|
| 422 |
+
bias_fp32,
|
| 423 |
+
x_scale.float(),
|
| 424 |
+
self.weight_scale,
|
| 425 |
+
fuse_gelu=False,
|
| 426 |
+
out_dtype=out_dtype,
|
| 427 |
+
)
|
| 428 |
+
|
| 429 |
+
def _gemm_by_backend(self, qinput, x_scale, out_dtype, backend: str):
|
| 430 |
+
if backend == "torchao":
|
| 431 |
+
return self._gemm_int8_torchao(qinput, x_scale, out_dtype)
|
| 432 |
+
if backend == "vllm":
|
| 433 |
+
return self._gemm_int8_vllm(qinput, x_scale, out_dtype)
|
| 434 |
+
if backend == "triton":
|
| 435 |
+
return self._gemm_int8_triton(qinput, x_scale, out_dtype)
|
| 436 |
+
if backend == "sgl":
|
| 437 |
+
return self._gemm_int8_sgl(qinput, x_scale, out_dtype)
|
| 438 |
+
if backend == "q8f":
|
| 439 |
+
return self._gemm_int8_q8f(qinput, x_scale, out_dtype)
|
| 440 |
+
raise ValueError(f"Unsupported int8 backend: {backend}")
|
| 441 |
+
|
| 442 |
+
@torch.compiler.disable(recursive=True)
|
| 443 |
+
def forward(self, x):
|
| 444 |
+
ori_dtype = x.dtype
|
| 445 |
+
assert ori_dtype in [
|
| 446 |
+
torch.float32,
|
| 447 |
+
torch.bfloat16,
|
| 448 |
+
torch.float16,
|
| 449 |
+
], "x.dtype must be float32, bfloat16, or float16"
|
| 450 |
+
|
| 451 |
+
if ori_dtype == torch.float32:
|
| 452 |
+
x = x.to(torch.bfloat16)
|
| 453 |
+
|
| 454 |
+
need_reshape = x.dim() == 3
|
| 455 |
+
if need_reshape:
|
| 456 |
+
origin_shape = x.shape
|
| 457 |
+
x_2d = x.view(-1, x.shape[-1])
|
| 458 |
+
else:
|
| 459 |
+
origin_shape = None
|
| 460 |
+
x_2d = x
|
| 461 |
+
|
| 462 |
+
backend = self._resolve_int8_backend()
|
| 463 |
+
qinput, x_scale = self._act_quant_by_backend(x_2d, backend)
|
| 464 |
+
output = self._gemm_by_backend(qinput, x_scale, x.dtype, backend)
|
| 465 |
+
|
| 466 |
+
if need_reshape and output.dim() == 2:
|
| 467 |
+
output = output.view(origin_shape[0], origin_shape[1], -1)
|
| 468 |
+
return output.to(ori_dtype)
|
| 469 |
+
|
| 470 |
+
|
| 471 |
+
class FP4DynamicLinear(torch.nn.Module):
|
| 472 |
+
def __init__(
|
| 473 |
+
self,
|
| 474 |
+
weight: torch.Tensor,
|
| 475 |
+
weight_scale: torch.Tensor,
|
| 476 |
+
bias: torch.nn.Parameter,
|
| 477 |
+
weight_global_scale: torch.Tensor = None,
|
| 478 |
+
native_fp8_support: bool = False,
|
| 479 |
+
quant_type: str = "nvfp4",
|
| 480 |
+
block_size: int = 16,
|
| 481 |
+
):
|
| 482 |
+
super().__init__()
|
| 483 |
+
self.weight = torch.nn.Parameter(weight, requires_grad=False)
|
| 484 |
+
self.weight_scale = torch.nn.Parameter(weight_scale, requires_grad=False)
|
| 485 |
+
self.bias = bias
|
| 486 |
+
self.native_fp8_support = native_fp8_support
|
| 487 |
+
self.quant_type = quant_type
|
| 488 |
+
self.block_size = block_size
|
| 489 |
+
self.profile_enabled = os.environ.get("ANGELSLIM_NVFP4_PROFILE", "0") == "1"
|
| 490 |
+
if weight_global_scale is None:
|
| 491 |
+
weight_global_scale = torch.tensor(1.0, dtype=torch.float32, device=weight.device)
|
| 492 |
+
self.weight_global_scale = torch.nn.Parameter(
|
| 493 |
+
weight_global_scale.to(dtype=torch.float32), requires_grad=False
|
| 494 |
+
)
|
| 495 |
+
self.calibrate_x_absmax()
|
| 496 |
+
|
| 497 |
+
def calibrate_x_absmax(self):
|
| 498 |
+
if self.quant_type in ("mxfp4", "mxfp6", "mxfp8"):
|
| 499 |
+
self.x_absmax = torch.tensor(1.0, dtype=torch.float32, device=self.weight.device)
|
| 500 |
+
self.input_global_scale = torch.tensor(
|
| 501 |
+
1.0, dtype=torch.float32, device=self.weight.device
|
| 502 |
+
)
|
| 503 |
+
else:
|
| 504 |
+
self.x_absmax = torch.tensor(5.0, dtype=torch.float32, device=self.weight.device)
|
| 505 |
+
self.input_global_scale = (2688.0 / self.x_absmax).to(torch.float32)
|
| 506 |
+
self.alpha = 1.0 / (self.input_global_scale * self.weight_global_scale)
|
| 507 |
+
|
| 508 |
+
@torch.compiler.disable(recursive=True)
|
| 509 |
+
def forward(self, x):
|
| 510 |
+
ori_dtype = x.dtype
|
| 511 |
+
assert ori_dtype in [
|
| 512 |
+
torch.float32,
|
| 513 |
+
torch.bfloat16,
|
| 514 |
+
torch.float16,
|
| 515 |
+
], "x.dtype must be float32, bfloat16, or float16"
|
| 516 |
+
|
| 517 |
+
if ori_dtype == torch.float32:
|
| 518 |
+
x = x.to(torch.bfloat16)
|
| 519 |
+
|
| 520 |
+
need_reshape = x.dim() == 3
|
| 521 |
+
if need_reshape:
|
| 522 |
+
origin_shape = x.shape
|
| 523 |
+
x_2d = x.view(-1, x.shape[-1])
|
| 524 |
+
else:
|
| 525 |
+
x_2d = x
|
| 526 |
+
|
| 527 |
+
if self.profile_enabled and x_2d.is_cuda:
|
| 528 |
+
torch.cuda.synchronize(x_2d.device)
|
| 529 |
+
t0 = time.perf_counter()
|
| 530 |
+
if self.quant_type == "nvfp4":
|
| 531 |
+
qinput, x_scale, _ = nvfp4_per_tensor_quant(x_2d, self.input_global_scale)
|
| 532 |
+
output = cutlass_scaled_nvfp4_mm(
|
| 533 |
+
qinput,
|
| 534 |
+
self.weight,
|
| 535 |
+
x_scale,
|
| 536 |
+
self.weight_scale,
|
| 537 |
+
self.alpha,
|
| 538 |
+
bias=self.bias,
|
| 539 |
+
)
|
| 540 |
+
elif self.quant_type == "mxfp4":
|
| 541 |
+
qinput, x_scale, _ = mxfp4_per_tensor_quant(x_2d)
|
| 542 |
+
output = cutlass_scaled_mxfp4_mm(
|
| 543 |
+
qinput,
|
| 544 |
+
self.weight,
|
| 545 |
+
x_scale,
|
| 546 |
+
self.weight_scale,
|
| 547 |
+
self.alpha,
|
| 548 |
+
bias=self.bias,
|
| 549 |
+
)
|
| 550 |
+
elif self.quant_type == "mxfp8":
|
| 551 |
+
qinput, x_scale, _ = mxfp8_per_tensor_quant(x_2d)
|
| 552 |
+
output = cutlass_scaled_mxfp8_mm(
|
| 553 |
+
qinput,
|
| 554 |
+
self.weight,
|
| 555 |
+
x_scale,
|
| 556 |
+
self.weight_scale,
|
| 557 |
+
self.alpha,
|
| 558 |
+
bias=self.bias,
|
| 559 |
+
)
|
| 560 |
+
elif self.quant_type == "mxfp6":
|
| 561 |
+
qinput, x_scale, _ = mxfp8_per_tensor_quant(x_2d)
|
| 562 |
+
output = cutlass_scaled_mxfp6_mxfp8_mm(
|
| 563 |
+
qinput,
|
| 564 |
+
self.weight,
|
| 565 |
+
x_scale,
|
| 566 |
+
self.weight_scale,
|
| 567 |
+
self.alpha,
|
| 568 |
+
bias=self.bias,
|
| 569 |
+
)
|
| 570 |
+
else:
|
| 571 |
+
raise ValueError(f"Invalid quant_type for FP4DynamicLinear: {self.quant_type}")
|
| 572 |
+
if self.profile_enabled and x_2d.is_cuda:
|
| 573 |
+
torch.cuda.synchronize(x_2d.device)
|
| 574 |
+
t1 = time.perf_counter()
|
| 575 |
+
if self.profile_enabled and x_2d.is_cuda:
|
| 576 |
+
torch.cuda.synchronize(x_2d.device)
|
| 577 |
+
t2 = time.perf_counter()
|
| 578 |
+
|
| 579 |
+
if self.profile_enabled:
|
| 580 |
+
print(
|
| 581 |
+
f"[NVFP4Linear] quant_ms={(t1 - t0) * 1000:.3f}, "
|
| 582 |
+
f"gemm_ms={(t2 - t1) * 1000:.3f}, "
|
| 583 |
+
f"shape=({x_2d.shape[0]}, {x_2d.shape[1]})"
|
| 584 |
+
)
|
| 585 |
+
|
| 586 |
+
if need_reshape:
|
| 587 |
+
output = output.view(origin_shape[0], origin_shape[1], -1)
|
| 588 |
+
|
| 589 |
+
return output.to(ori_dtype)
|
quantizer/quant/ptq.py
ADDED
|
@@ -0,0 +1,272 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
import copy
|
| 16 |
+
import logging
|
| 17 |
+
import re
|
| 18 |
+
from typing import Callable, List, Optional, Tuple, Union
|
| 19 |
+
|
| 20 |
+
import torch
|
| 21 |
+
import tqdm
|
| 22 |
+
|
| 23 |
+
from .modules import FP4DynamicLinear, FP8DynamicLinear, FP8WeightOnlyLinear, INT8DynamicLinear
|
| 24 |
+
from .quant_func import (
|
| 25 |
+
fp8_per_block_quant,
|
| 26 |
+
fp8_per_channel_quant,
|
| 27 |
+
fp8_per_tensor_quant,
|
| 28 |
+
fp8_per_token_group_quant,
|
| 29 |
+
int8_per_channel_quant,
|
| 30 |
+
mxfp4_per_tensor_quant,
|
| 31 |
+
mxfp6_per_tensor_quant,
|
| 32 |
+
mxfp8_per_tensor_quant,
|
| 33 |
+
nvfp4_per_tensor_quant,
|
| 34 |
+
)
|
| 35 |
+
from .utils import (
|
| 36 |
+
QuantType,
|
| 37 |
+
_ensure_deep_gemm,
|
| 38 |
+
_ensure_sgl_kernel,
|
| 39 |
+
cleanup_memory,
|
| 40 |
+
load_fp8_scales,
|
| 41 |
+
load_quantized_model,
|
| 42 |
+
replace_module,
|
| 43 |
+
save_quantized_model,
|
| 44 |
+
should_quantize_layer,
|
| 45 |
+
)
|
| 46 |
+
|
| 47 |
+
__all__ = ["DynamicDiTQuantizer"]
|
| 48 |
+
|
| 49 |
+
logger = logging.getLogger(__name__)
|
| 50 |
+
|
| 51 |
+
|
| 52 |
+
class DynamicDiTQuantizer:
|
| 53 |
+
"""
|
| 54 |
+
Quantizer for DiT that supports various FP8 quantization strategies.
|
| 55 |
+
Optimized for efficient initialization and conversion.
|
| 56 |
+
"""
|
| 57 |
+
|
| 58 |
+
def __init__(
|
| 59 |
+
self,
|
| 60 |
+
quant_type: str = QuantType.FP8_PER_TENSOR,
|
| 61 |
+
layer_filter: Optional[Callable[[str], bool]] = None,
|
| 62 |
+
include_patterns: Optional[List[Union[str, re.Pattern]]] = None,
|
| 63 |
+
exclude_patterns: Optional[List[Union[str, re.Pattern]]] = None,
|
| 64 |
+
native_fp8_support: Optional[bool] = None,
|
| 65 |
+
):
|
| 66 |
+
QuantType.validate(quant_type)
|
| 67 |
+
self.fp8_scales_map = {}
|
| 68 |
+
self.quant_type = quant_type
|
| 69 |
+
self.include_patterns = include_patterns or [
|
| 70 |
+
"wrapped_module",
|
| 71 |
+
"block",
|
| 72 |
+
"lin",
|
| 73 |
+
"img",
|
| 74 |
+
"txt",
|
| 75 |
+
]
|
| 76 |
+
self.exclude_patterns = exclude_patterns or ["embed"]
|
| 77 |
+
|
| 78 |
+
# Configure layer filter function
|
| 79 |
+
self.layer_filter = (
|
| 80 |
+
layer_filter
|
| 81 |
+
if layer_filter is not None
|
| 82 |
+
else lambda name: should_quantize_layer(
|
| 83 |
+
name, self.include_patterns, self.exclude_patterns
|
| 84 |
+
)
|
| 85 |
+
)
|
| 86 |
+
|
| 87 |
+
# Auto-detect FP8 native support, fallback to False if not present
|
| 88 |
+
if native_fp8_support is not None:
|
| 89 |
+
self.native_fp8_support = native_fp8_support
|
| 90 |
+
else:
|
| 91 |
+
self.native_fp8_support = (
|
| 92 |
+
torch.cuda.is_available() and torch.cuda.get_device_capability() >= (8, 9)
|
| 93 |
+
)
|
| 94 |
+
|
| 95 |
+
self.quantize_linear_module = self._set_quantize_linear_module()
|
| 96 |
+
|
| 97 |
+
def _set_quantize_linear_module(self) -> torch.nn.Module:
|
| 98 |
+
if self.quant_type in (
|
| 99 |
+
QuantType.INT8,
|
| 100 |
+
QuantType.INT8_TORCHAO,
|
| 101 |
+
QuantType.INT8_VLLM,
|
| 102 |
+
QuantType.INT8_TRITON,
|
| 103 |
+
QuantType.INT8_SGL,
|
| 104 |
+
QuantType.INT8_Q8F,
|
| 105 |
+
):
|
| 106 |
+
return INT8DynamicLinear
|
| 107 |
+
if self.quant_type in (QuantType.NVFP4, QuantType.MXFP4, QuantType.MXFP6, QuantType.MXFP8):
|
| 108 |
+
return FP4DynamicLinear
|
| 109 |
+
if "fp8" in self.quant_type:
|
| 110 |
+
if self.quant_type == QuantType.FP8_PER_TENSOR_WEIGHT_ONLY:
|
| 111 |
+
return FP8WeightOnlyLinear
|
| 112 |
+
else:
|
| 113 |
+
return FP8DynamicLinear
|
| 114 |
+
raise ValueError(f"Invalid quant_type: {self.quant_type}")
|
| 115 |
+
|
| 116 |
+
def _quantize_linear_weight(
|
| 117 |
+
self, linear: torch.nn.Linear
|
| 118 |
+
) -> Union[Tuple[torch.Tensor, torch.Tensor], Tuple[torch.Tensor, torch.Tensor, torch.Tensor]]:
|
| 119 |
+
if self.quant_type == QuantType.FP8_PER_TENSOR:
|
| 120 |
+
quant_weight, weight_scale = fp8_per_tensor_quant(linear.weight)
|
| 121 |
+
elif self.quant_type == QuantType.FP8_PER_TOKEN:
|
| 122 |
+
quant_weight, weight_scale = fp8_per_token_group_quant(
|
| 123 |
+
linear.weight, linear.weight.shape[-1]
|
| 124 |
+
)
|
| 125 |
+
weight_scale = weight_scale.t()
|
| 126 |
+
elif self.quant_type == QuantType.FP8_PER_TOKEN_SGL:
|
| 127 |
+
if self.native_fp8_support:
|
| 128 |
+
_ensure_sgl_kernel()
|
| 129 |
+
quant_weight, weight_scale = fp8_per_channel_quant(linear.weight)
|
| 130 |
+
elif self.quant_type == QuantType.FP8_PER_BLOCK:
|
| 131 |
+
if self.native_fp8_support:
|
| 132 |
+
_ensure_deep_gemm()
|
| 133 |
+
quant_weight, weight_scale = fp8_per_block_quant(linear.weight)
|
| 134 |
+
elif self.quant_type == QuantType.FP8_PER_CHANNEL_VLLM:
|
| 135 |
+
quant_weight, weight_scale = fp8_per_channel_quant(linear.weight)
|
| 136 |
+
elif self.quant_type == QuantType.FP8_PER_TENSOR_WEIGHT_ONLY:
|
| 137 |
+
quant_weight, weight_scale = fp8_per_tensor_quant(linear.weight)
|
| 138 |
+
elif self.quant_type in (
|
| 139 |
+
QuantType.INT8,
|
| 140 |
+
QuantType.INT8_TORCHAO,
|
| 141 |
+
QuantType.INT8_VLLM,
|
| 142 |
+
QuantType.INT8_TRITON,
|
| 143 |
+
QuantType.INT8_SGL,
|
| 144 |
+
QuantType.INT8_Q8F,
|
| 145 |
+
):
|
| 146 |
+
quant_weight, weight_scale = int8_per_channel_quant(linear.weight)
|
| 147 |
+
elif self.quant_type == QuantType.NVFP4:
|
| 148 |
+
quant_weight, weight_scale, weight_global_scale = nvfp4_per_tensor_quant(linear.weight)
|
| 149 |
+
return quant_weight, weight_scale, weight_global_scale
|
| 150 |
+
elif self.quant_type == QuantType.MXFP4:
|
| 151 |
+
quant_weight, weight_scale, weight_global_scale = mxfp4_per_tensor_quant(linear.weight)
|
| 152 |
+
return quant_weight, weight_scale, weight_global_scale
|
| 153 |
+
elif self.quant_type == QuantType.MXFP8:
|
| 154 |
+
quant_weight, weight_scale, weight_global_scale = mxfp8_per_tensor_quant(linear.weight)
|
| 155 |
+
return quant_weight, weight_scale, weight_global_scale
|
| 156 |
+
elif self.quant_type == QuantType.MXFP6:
|
| 157 |
+
# Mixed kernel route expects MXFP8-formatted weights.
|
| 158 |
+
quant_weight, weight_scale, weight_global_scale = mxfp6_per_tensor_quant(linear.weight)
|
| 159 |
+
return quant_weight, weight_scale, weight_global_scale
|
| 160 |
+
else:
|
| 161 |
+
raise ValueError(f"Invalid quant_type: {self.quant_type}")
|
| 162 |
+
return quant_weight, weight_scale
|
| 163 |
+
|
| 164 |
+
def _convert_linear_with_scale(
|
| 165 |
+
self, model: torch.nn.Module, scale: Union[torch.Tensor, float]
|
| 166 |
+
):
|
| 167 |
+
model.to(torch.bfloat16)
|
| 168 |
+
assert scale is not None, "scale is required"
|
| 169 |
+
self.fp8_scales_map = load_fp8_scales(scale)
|
| 170 |
+
converted_count = 0
|
| 171 |
+
for name, module in tqdm.tqdm(list(model.named_modules()), desc="converting linear"):
|
| 172 |
+
if isinstance(module, torch.nn.Linear) and self.layer_filter(name):
|
| 173 |
+
# Prefer $name.weight_scale, fallback to "$name" key if needed
|
| 174 |
+
s = self.fp8_scales_map.get(f"{name}.weight_scale") or self.fp8_scales_map.get(
|
| 175 |
+
name
|
| 176 |
+
)
|
| 177 |
+
if s is None:
|
| 178 |
+
continue
|
| 179 |
+
# import pdb; pdb.set_trace()
|
| 180 |
+
weight = module.weight.to(torch.float8_e4m3fn)
|
| 181 |
+
bias = copy.deepcopy(module.bias) if module.bias is not None else None
|
| 182 |
+
quant_linear = self.quantize_linear_module(
|
| 183 |
+
weight=weight,
|
| 184 |
+
weight_scale=s.float(),
|
| 185 |
+
bias=bias,
|
| 186 |
+
native_fp8_support=self.native_fp8_support,
|
| 187 |
+
quant_type=self.quant_type,
|
| 188 |
+
)
|
| 189 |
+
replace_module(model, name, quant_linear)
|
| 190 |
+
del module.weight, module.bias, module
|
| 191 |
+
converted_count += 1
|
| 192 |
+
print(f"[INIT] quantized linear layers: {converted_count}")
|
| 193 |
+
cleanup_memory()
|
| 194 |
+
|
| 195 |
+
def _convert_linear(self, model: torch.nn.Module):
|
| 196 |
+
model.to(torch.bfloat16)
|
| 197 |
+
named_modules = list(model.named_modules())
|
| 198 |
+
converted_count = 0
|
| 199 |
+
for name, module in tqdm.tqdm(named_modules, desc="converting linear"):
|
| 200 |
+
if isinstance(module, torch.nn.Linear) and self.layer_filter(name):
|
| 201 |
+
quantized = self._quantize_linear_weight(module)
|
| 202 |
+
if self.quant_type in (
|
| 203 |
+
QuantType.NVFP4,
|
| 204 |
+
QuantType.MXFP4,
|
| 205 |
+
QuantType.MXFP6,
|
| 206 |
+
QuantType.MXFP8,
|
| 207 |
+
):
|
| 208 |
+
quant_weight, weight_scale, weight_global_scale = quantized
|
| 209 |
+
else:
|
| 210 |
+
quant_weight, weight_scale = quantized
|
| 211 |
+
self.fp8_scales_map[f"{name}.weight_scale"] = weight_scale
|
| 212 |
+
bias = copy.deepcopy(module.bias) if module.bias is not None else None
|
| 213 |
+
if self.quant_type in (
|
| 214 |
+
QuantType.NVFP4,
|
| 215 |
+
QuantType.MXFP4,
|
| 216 |
+
QuantType.MXFP6,
|
| 217 |
+
QuantType.MXFP8,
|
| 218 |
+
):
|
| 219 |
+
self.fp8_scales_map[f"{name}.weight_global_scale"] = weight_global_scale
|
| 220 |
+
quant_linear = self.quantize_linear_module(
|
| 221 |
+
weight=quant_weight,
|
| 222 |
+
weight_scale=weight_scale,
|
| 223 |
+
weight_global_scale=weight_global_scale,
|
| 224 |
+
bias=bias,
|
| 225 |
+
native_fp8_support=self.native_fp8_support,
|
| 226 |
+
quant_type=self.quant_type,
|
| 227 |
+
)
|
| 228 |
+
else:
|
| 229 |
+
quant_linear = self.quantize_linear_module(
|
| 230 |
+
weight=quant_weight,
|
| 231 |
+
weight_scale=weight_scale,
|
| 232 |
+
bias=bias,
|
| 233 |
+
native_fp8_support=self.native_fp8_support,
|
| 234 |
+
quant_type=self.quant_type,
|
| 235 |
+
)
|
| 236 |
+
replace_module(model, name, quant_linear)
|
| 237 |
+
del module.weight, module.bias, module
|
| 238 |
+
converted_count += 1
|
| 239 |
+
print(f"[INIT] quantized linear layers: {converted_count}")
|
| 240 |
+
cleanup_memory()
|
| 241 |
+
|
| 242 |
+
def convert_linear(
|
| 243 |
+
self, model: torch.nn.Module, scale: Optional[Union[torch.Tensor, float]] = None
|
| 244 |
+
):
|
| 245 |
+
if scale is not None:
|
| 246 |
+
self._convert_linear_with_scale(model, scale)
|
| 247 |
+
else:
|
| 248 |
+
self._convert_linear(model)
|
| 249 |
+
|
| 250 |
+
def export_quantized_weight(self, model: torch.nn.Module, save_path: str):
|
| 251 |
+
assert (
|
| 252 |
+
self.quant_type == QuantType.FP8_PER_TENSOR
|
| 253 |
+
), "Currently only FP8_PER_TENSOR is supported for export"
|
| 254 |
+
self.convert_linear(model)
|
| 255 |
+
save_quantized_model(model, save_path, self.fp8_scales_map)
|
| 256 |
+
logger.info(f"Quantized model saved to {save_path}")
|
| 257 |
+
logger.info(f"Quantized scales saved to {save_path}/fp8_scales.safetensors")
|
| 258 |
+
|
| 259 |
+
@staticmethod
|
| 260 |
+
def load_quantized_model(model_class, save_path: str, device: str = "cpu"):
|
| 261 |
+
"""
|
| 262 |
+
Load a quantized model from save_path.
|
| 263 |
+
|
| 264 |
+
Args:
|
| 265 |
+
model_class: The model class to instantiate
|
| 266 |
+
save_path: Path to the saved model directory
|
| 267 |
+
device: Device to load the model on
|
| 268 |
+
|
| 269 |
+
Returns:
|
| 270 |
+
Loaded model with quantized weights
|
| 271 |
+
"""
|
| 272 |
+
return load_quantized_model(model_class, save_path, device)
|
quantizer/quant/quant_func.py
ADDED
|
@@ -0,0 +1,520 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
from typing import Optional, Tuple
|
| 16 |
+
|
| 17 |
+
import torch
|
| 18 |
+
|
| 19 |
+
from ..kernels.python.gemm import fp8_gemm_triton_block
|
| 20 |
+
from lightx2v_kernel.gemm import (
|
| 21 |
+
cutlass_scaled_mxfp4_mm,
|
| 22 |
+
cutlass_scaled_mxfp6_mxfp8_mm,
|
| 23 |
+
cutlass_scaled_mxfp8_mm,
|
| 24 |
+
cutlass_scaled_nvfp4_mm,
|
| 25 |
+
scaled_mxfp4_quant,
|
| 26 |
+
scaled_mxfp6_quant,
|
| 27 |
+
scaled_mxfp8_quant,
|
| 28 |
+
scaled_nvfp4_quant,
|
| 29 |
+
)
|
| 30 |
+
from ..kernels.python.quantizers import (
|
| 31 |
+
fp8_per_block_quant_triton,
|
| 32 |
+
fp8_per_token_group_quant_triton,
|
| 33 |
+
)
|
| 34 |
+
from .utils import QuantType, _ensure_deep_gemm, _ensure_sgl_kernel
|
| 35 |
+
|
| 36 |
+
FP8_MAX = float(torch.finfo(torch.float8_e4m3fn).max)
|
| 37 |
+
FP8_MIN = float(torch.finfo(torch.float8_e4m3fn).min)
|
| 38 |
+
|
| 39 |
+
|
| 40 |
+
# quant function for per-tensor fp8
|
| 41 |
+
# modified from https://github.com/neuralmagic/AutoFP8/blob/main/auto_fp8/quantize.py
|
| 42 |
+
def fp8_per_tensor_quant(x: torch.Tensor) -> Tuple[torch.Tensor, float]:
|
| 43 |
+
"""Dynamically Quantize a tensor using per-tensor static scaling factor.
|
| 44 |
+
Args:
|
| 45 |
+
x: The input tensor.
|
| 46 |
+
"""
|
| 47 |
+
if x.numel() == 0:
|
| 48 |
+
# handle empty tensor for empty MoE experts
|
| 49 |
+
min_val, max_val = (
|
| 50 |
+
torch.tensor(-16.0, dtype=x.dtype),
|
| 51 |
+
torch.tensor(16.0, dtype=x.dtype),
|
| 52 |
+
)
|
| 53 |
+
else:
|
| 54 |
+
min_val, max_val = x.aminmax()
|
| 55 |
+
amax = torch.maximum(min_val.abs(), max_val.abs())
|
| 56 |
+
scale = FP8_MAX / amax.clamp(min=1e-12)
|
| 57 |
+
qx = (x * scale).clamp(min=FP8_MIN, max=FP8_MAX)
|
| 58 |
+
qx = qx.to(torch.float8_e4m3fn)
|
| 59 |
+
scale = scale.float().reciprocal()
|
| 60 |
+
return qx, scale
|
| 61 |
+
|
| 62 |
+
|
| 63 |
+
# quant function for per-token and per-group fp8
|
| 64 |
+
def fp8_per_token_group_quant(
|
| 65 |
+
x: torch.Tensor,
|
| 66 |
+
group_size: int,
|
| 67 |
+
eps: float = 1e-10,
|
| 68 |
+
dtype: torch.dtype = torch.float8_e4m3fn,
|
| 69 |
+
column_major_scales: bool = False,
|
| 70 |
+
scale_tma_aligned: bool = False,
|
| 71 |
+
) -> Tuple[torch.Tensor, torch.Tensor]:
|
| 72 |
+
"""
|
| 73 |
+
Per-token-group FP8 quantization with automatic backend selection.
|
| 74 |
+
|
| 75 |
+
Uses Triton kernel on GPU when available, otherwise falls back to PyTorch.
|
| 76 |
+
|
| 77 |
+
Args:
|
| 78 |
+
x: Input tensor
|
| 79 |
+
group_size: Size of each quantization group
|
| 80 |
+
eps: Small value to avoid division by zero
|
| 81 |
+
dtype: Target FP8 data type
|
| 82 |
+
column_major_scales: Whether to use column-major scale layout
|
| 83 |
+
scale_tma_aligned: Whether to use TMA-aligned scales
|
| 84 |
+
|
| 85 |
+
Returns:
|
| 86 |
+
Tuple of (quantized_tensor, scale_tensor)
|
| 87 |
+
"""
|
| 88 |
+
# fp8_per_token_group_quant_triton is automatically selected based on
|
| 89 |
+
# backend availability (Triton vs PyTorch) via __init__.py
|
| 90 |
+
return fp8_per_token_group_quant_triton(
|
| 91 |
+
x,
|
| 92 |
+
group_size,
|
| 93 |
+
eps,
|
| 94 |
+
dtype,
|
| 95 |
+
column_major_scales,
|
| 96 |
+
scale_tma_aligned,
|
| 97 |
+
)
|
| 98 |
+
|
| 99 |
+
|
| 100 |
+
def fp8_per_channel_quant(weight: torch.Tensor) -> Tuple[torch.Tensor, torch.Tensor]:
|
| 101 |
+
"""
|
| 102 |
+
Per-channel FP8 weight quantization (E4M3 format)
|
| 103 |
+
|
| 104 |
+
Args:
|
| 105 |
+
weight: Original weight tensor with shape [out_features, in_features]
|
| 106 |
+
|
| 107 |
+
Returns:
|
| 108 |
+
weight_quant: Quantized weight [out_features, in_features], dtype=float8_e4m3fn
|
| 109 |
+
weight_scale: Scale factors [out_features, 1], dtype=float32
|
| 110 |
+
"""
|
| 111 |
+
abs_max = torch.abs(weight).amax(dim=1, keepdim=True) # [out_features, 1]
|
| 112 |
+
|
| 113 |
+
weight_scale = abs_max / FP8_MAX
|
| 114 |
+
weight_scale = torch.clamp(weight_scale, min=1e-12)
|
| 115 |
+
|
| 116 |
+
weight_scaled = (weight / weight_scale).clamp(min=FP8_MIN, max=FP8_MAX)
|
| 117 |
+
weight_quant = weight_scaled.to(torch.float8_e4m3fn)
|
| 118 |
+
|
| 119 |
+
return weight_quant, weight_scale.float()
|
| 120 |
+
|
| 121 |
+
|
| 122 |
+
def int8_per_channel_quant(weight: torch.Tensor) -> Tuple[torch.Tensor, torch.Tensor]:
|
| 123 |
+
"""
|
| 124 |
+
Per-channel symmetric INT8 quantization for linear weights.
|
| 125 |
+
|
| 126 |
+
Args:
|
| 127 |
+
weight: Weight tensor with shape [out_features, in_features]
|
| 128 |
+
|
| 129 |
+
Returns:
|
| 130 |
+
weight_quant: Quantized INT8 weight [out_features, in_features]
|
| 131 |
+
weight_scale: Per-channel scales [out_features, 1], float32
|
| 132 |
+
"""
|
| 133 |
+
abs_max = torch.abs(weight).amax(dim=1, keepdim=True).clamp(min=1e-5)
|
| 134 |
+
qmin, qmax = -128, 127
|
| 135 |
+
weight_scale = (abs_max / qmax).to(torch.float32)
|
| 136 |
+
weight_quant = torch.clamp(torch.round(weight / weight_scale), qmin, qmax).to(torch.int8)
|
| 137 |
+
return weight_quant, weight_scale
|
| 138 |
+
|
| 139 |
+
|
| 140 |
+
def nvfp4_per_tensor_quant(
|
| 141 |
+
x: torch.Tensor, input_global_scale: Optional[torch.Tensor] = None
|
| 142 |
+
) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
|
| 143 |
+
"""
|
| 144 |
+
Per-tensor NVFP4 quantization using LightX2V kernel.
|
| 145 |
+
|
| 146 |
+
Returns packed uint8 values and per-block quantized scales.
|
| 147 |
+
"""
|
| 148 |
+
if x.dim() != 2:
|
| 149 |
+
raise ValueError(f"nvfp4_per_tensor_quant expects 2D tensor, but got shape={tuple(x.shape)}")
|
| 150 |
+
|
| 151 |
+
if input_global_scale is None:
|
| 152 |
+
x_absmax = torch.abs(x).amax().float().clamp(min=1e-12)
|
| 153 |
+
input_global_scale = (2688.0 / x_absmax).to(torch.float32)
|
| 154 |
+
else:
|
| 155 |
+
input_global_scale = input_global_scale.to(device=x.device, dtype=torch.float32)
|
| 156 |
+
|
| 157 |
+
qx, qscale = scaled_nvfp4_quant(x, input_global_scale)
|
| 158 |
+
return qx, qscale, input_global_scale
|
| 159 |
+
|
| 160 |
+
|
| 161 |
+
def mxfp4_per_tensor_quant(
|
| 162 |
+
x: torch.Tensor,
|
| 163 |
+
) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
|
| 164 |
+
"""
|
| 165 |
+
Per-tensor MXFP4 quantization using LightX2V kernel.
|
| 166 |
+
|
| 167 |
+
Returns packed uint8 values and per-block quantized scales.
|
| 168 |
+
"""
|
| 169 |
+
if x.dim() != 2:
|
| 170 |
+
raise ValueError(f"mxfp4_per_tensor_quant expects 2D tensor, but got shape={tuple(x.shape)}")
|
| 171 |
+
|
| 172 |
+
qx, qscale = scaled_mxfp4_quant(x)
|
| 173 |
+
# Keep a compatible triplet return signature with NVFP4 paths.
|
| 174 |
+
input_global_scale = torch.tensor(1.0, device=x.device, dtype=torch.float32)
|
| 175 |
+
return qx, qscale, input_global_scale
|
| 176 |
+
|
| 177 |
+
|
| 178 |
+
def mxfp6_per_tensor_quant(
|
| 179 |
+
x: torch.Tensor,
|
| 180 |
+
) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
|
| 181 |
+
"""
|
| 182 |
+
Per-tensor MXFP6 quantization using LightX2V kernel.
|
| 183 |
+
|
| 184 |
+
Returns packed uint8 values and per-block quantized scales.
|
| 185 |
+
"""
|
| 186 |
+
if x.dim() != 2:
|
| 187 |
+
raise ValueError(f"mxfp6_per_tensor_quant expects 2D tensor, but got shape={tuple(x.shape)}")
|
| 188 |
+
|
| 189 |
+
qx, qscale = scaled_mxfp6_quant(x)
|
| 190 |
+
input_global_scale = torch.tensor(1.0, device=x.device, dtype=torch.float32)
|
| 191 |
+
return qx, qscale, input_global_scale
|
| 192 |
+
|
| 193 |
+
|
| 194 |
+
def mxfp8_per_tensor_quant(
|
| 195 |
+
x: torch.Tensor,
|
| 196 |
+
) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
|
| 197 |
+
"""
|
| 198 |
+
Per-tensor MXFP8 quantization using LightX2V kernel.
|
| 199 |
+
|
| 200 |
+
Returns packed uint8 values and per-block quantized scales.
|
| 201 |
+
"""
|
| 202 |
+
if x.dim() != 2:
|
| 203 |
+
raise ValueError(f"mxfp8_per_tensor_quant expects 2D tensor, but got shape={tuple(x.shape)}")
|
| 204 |
+
|
| 205 |
+
qx, qscale = scaled_mxfp8_quant(x)
|
| 206 |
+
input_global_scale = torch.tensor(1.0, device=x.device, dtype=torch.float32)
|
| 207 |
+
return qx, qscale, input_global_scale
|
| 208 |
+
|
| 209 |
+
|
| 210 |
+
def fp8_per_token_quant_sgl(x: torch.Tensor) -> Tuple[torch.Tensor, torch.Tensor]:
|
| 211 |
+
m, k = x.shape
|
| 212 |
+
input_tensor_quant = torch.empty(
|
| 213 |
+
(m, k), dtype=torch.float8_e4m3fn, device="cuda", requires_grad=False
|
| 214 |
+
)
|
| 215 |
+
input_tensor_scale = torch.empty(
|
| 216 |
+
(m, 1), dtype=torch.float32, device="cuda", requires_grad=False
|
| 217 |
+
)
|
| 218 |
+
_sgl_kernel = _ensure_sgl_kernel()
|
| 219 |
+
_sgl_kernel.sgl_per_token_quant_fp8(x, input_tensor_quant, input_tensor_scale)
|
| 220 |
+
return input_tensor_quant, input_tensor_scale
|
| 221 |
+
|
| 222 |
+
|
| 223 |
+
# pure torch implementation of block-wise FP8 quantization on cpu
|
| 224 |
+
def fp8_per_block_quant_torch(
|
| 225 |
+
x: torch.Tensor, block_size: int = 128
|
| 226 |
+
) -> Tuple[torch.Tensor, torch.Tensor]:
|
| 227 |
+
"""
|
| 228 |
+
Pure Torch implementation of block-wise FP8 (e4m3fn) quantization.
|
| 229 |
+
|
| 230 |
+
Args:
|
| 231 |
+
x (torch.Tensor): 2D tensor (M, N) in float types.
|
| 232 |
+
block_size (int): Block size for both M and N dimensions.
|
| 233 |
+
|
| 234 |
+
Returns:
|
| 235 |
+
Tuple[torch.Tensor, torch.Tensor]:
|
| 236 |
+
- y: quantized tensor with dtype=torch.float8_e4m3fn and same
|
| 237 |
+
shape/device as x
|
| 238 |
+
- s: per-block scale tensor with shape (ceil(M/bs), ceil(N/bs)) on
|
| 239 |
+
x.device
|
| 240 |
+
"""
|
| 241 |
+
assert x.is_contiguous()
|
| 242 |
+
assert x.dim() == 2
|
| 243 |
+
|
| 244 |
+
M, N = x.size()
|
| 245 |
+
device = x.device
|
| 246 |
+
|
| 247 |
+
# Output tensors
|
| 248 |
+
y = torch.empty_like(x, dtype=torch.float8_e4m3fn)
|
| 249 |
+
m_blocks = (M + block_size - 1) // block_size
|
| 250 |
+
n_blocks = (N + block_size - 1) // block_size
|
| 251 |
+
s = torch.empty((m_blocks, n_blocks), dtype=torch.float32, device=device)
|
| 252 |
+
|
| 253 |
+
# Iterate over blocks
|
| 254 |
+
for mb in range(m_blocks):
|
| 255 |
+
m_start = mb * block_size
|
| 256 |
+
m_end = min(m_start + block_size, M)
|
| 257 |
+
for nb in range(n_blocks):
|
| 258 |
+
n_start = nb * block_size
|
| 259 |
+
n_end = min(n_start + block_size, N)
|
| 260 |
+
|
| 261 |
+
x_block = x[m_start:m_end, n_start:n_end].to(torch.float32)
|
| 262 |
+
# Compute per-block scale: max(abs(x))/448.0, guard zero to 1.0
|
| 263 |
+
max_val = x_block.abs().amax().to(torch.float32)
|
| 264 |
+
scale = max_val / FP8_MAX
|
| 265 |
+
if scale.item() == 0.0:
|
| 266 |
+
scale = torch.tensor(1.0, dtype=torch.float32, device=device)
|
| 267 |
+
|
| 268 |
+
# Quantize block
|
| 269 |
+
y_block = (x_block / scale).to(torch.float8_e4m3fn)
|
| 270 |
+
|
| 271 |
+
# Store results
|
| 272 |
+
y[m_start:m_end, n_start:n_end] = y_block
|
| 273 |
+
s[mb, nb] = scale
|
| 274 |
+
|
| 275 |
+
return y, s
|
| 276 |
+
|
| 277 |
+
|
| 278 |
+
# quant function for per-block fp8
|
| 279 |
+
def fp8_per_block_quant(
|
| 280 |
+
x: torch.Tensor, block_size: int = 128
|
| 281 |
+
) -> Tuple[torch.Tensor, torch.Tensor]:
|
| 282 |
+
"""
|
| 283 |
+
Per-block FP8 quantization with automatic GPU/CPU dispatch.
|
| 284 |
+
|
| 285 |
+
Quantizes a FP32 2D tensor to FP8 (E4M3FN) using block-wise quantization.
|
| 286 |
+
For each (block_size x block_size) block:
|
| 287 |
+
- scale = max(abs(block)) / 448.0 (FP8 E4M3FN max magnitude)
|
| 288 |
+
- if block is all zeros, use scale = 1.0 to avoid div-by-zero
|
| 289 |
+
- scale, clamp and cast to FP8
|
| 290 |
+
|
| 291 |
+
Args:
|
| 292 |
+
x: Input tensor (2D)
|
| 293 |
+
block_size: Block size for quantization
|
| 294 |
+
|
| 295 |
+
Returns:
|
| 296 |
+
Tuple of (quantized_tensor, scale_tensor):
|
| 297 |
+
- y: Quantized FP8 tensor, same shape as input
|
| 298 |
+
- s: Per-block scales, shape (num_blocks_M, num_blocks_N)
|
| 299 |
+
"""
|
| 300 |
+
if x.is_cuda:
|
| 301 |
+
return fp8_per_block_quant_triton(x, block_size)
|
| 302 |
+
else:
|
| 303 |
+
return fp8_per_block_quant_torch(x, block_size)
|
| 304 |
+
|
| 305 |
+
|
| 306 |
+
# gemm function for per-block quantization fp8 using deepgemm
|
| 307 |
+
def fp8_gemm_deepgemm_block(
|
| 308 |
+
a: torch.Tensor,
|
| 309 |
+
a_s: torch.Tensor,
|
| 310 |
+
b: torch.Tensor,
|
| 311 |
+
b_s: torch.Tensor,
|
| 312 |
+
out_dtype: torch.dtype = torch.bfloat16,
|
| 313 |
+
bias: Optional[torch.Tensor] = None,
|
| 314 |
+
origin_shape: Optional[Tuple[int, ...]] = None,
|
| 315 |
+
) -> torch.Tensor:
|
| 316 |
+
a_fp8 = (a, a_s)
|
| 317 |
+
b_fp8 = (b, b_s)
|
| 318 |
+
out = torch.empty((a.shape[0], b.shape[0]), device=a.device, dtype=torch.bfloat16)
|
| 319 |
+
|
| 320 |
+
_deep_gemm = _ensure_deep_gemm()
|
| 321 |
+
_deep_gemm.fp8_gemm_nt(a_fp8, b_fp8, out)
|
| 322 |
+
|
| 323 |
+
if origin_shape is not None:
|
| 324 |
+
out = out.reshape([*origin_shape[:-1], b.shape[0]])
|
| 325 |
+
if bias is not None:
|
| 326 |
+
out += bias
|
| 327 |
+
|
| 328 |
+
return out.to(out_dtype)
|
| 329 |
+
|
| 330 |
+
|
| 331 |
+
# gemm function for per-token and per-group fp8 quantization using torch native api
|
| 332 |
+
def fp8_gemm_torch_tensor_token(
|
| 333 |
+
a: torch.Tensor,
|
| 334 |
+
a_s: torch.Tensor,
|
| 335 |
+
b: torch.Tensor,
|
| 336 |
+
b_s: torch.Tensor,
|
| 337 |
+
out_dtype: torch.dtype = torch.bfloat16,
|
| 338 |
+
bias: Optional[torch.Tensor] = None,
|
| 339 |
+
) -> torch.Tensor:
|
| 340 |
+
need_reshape = a.dim() == 3
|
| 341 |
+
if need_reshape:
|
| 342 |
+
batch_size = a.shape[0]
|
| 343 |
+
A_input = a.reshape(-1, a.shape[-1])
|
| 344 |
+
else:
|
| 345 |
+
batch_size = None
|
| 346 |
+
A_input = a
|
| 347 |
+
|
| 348 |
+
output = torch._scaled_mm(
|
| 349 |
+
A_input,
|
| 350 |
+
b.t(),
|
| 351 |
+
out_dtype=out_dtype,
|
| 352 |
+
scale_a=a_s,
|
| 353 |
+
scale_b=b_s,
|
| 354 |
+
bias=bias,
|
| 355 |
+
)
|
| 356 |
+
# If output is a tuple, take the first element
|
| 357 |
+
if isinstance(output, tuple):
|
| 358 |
+
output = output[0]
|
| 359 |
+
|
| 360 |
+
if need_reshape:
|
| 361 |
+
output = output.reshape(batch_size, output.shape[0] // batch_size, output.shape[1])
|
| 362 |
+
|
| 363 |
+
return output
|
| 364 |
+
|
| 365 |
+
|
| 366 |
+
def fp8_weight_only_gemm(A, B, B_scale, bias, out_dtype):
|
| 367 |
+
"""Perform FP8 GEMM operation with fallback to standard linear.
|
| 368 |
+
|
| 369 |
+
Args:
|
| 370 |
+
A: Input tensor A.
|
| 371 |
+
B: Input tensor B.
|
| 372 |
+
B_scale: Scale factor for tensor B.
|
| 373 |
+
bias: Optional bias tensor.
|
| 374 |
+
out_dtype: Output data type.
|
| 375 |
+
native_fp8_support: Whether to use native FP8 support.
|
| 376 |
+
quant_type: Quantization type.
|
| 377 |
+
origin_shape: Original shape for reshaping.
|
| 378 |
+
|
| 379 |
+
Returns:
|
| 380 |
+
torch.Tensor: Result of the GEMM operation.
|
| 381 |
+
"""
|
| 382 |
+
if A.numel() == 0:
|
| 383 |
+
return torch.empty(size=(0, B.shape[0]), dtype=out_dtype, device=A.device)
|
| 384 |
+
|
| 385 |
+
output = torch.nn.functional.linear(
|
| 386 |
+
A.to(out_dtype),
|
| 387 |
+
B.to(out_dtype) * B_scale.to(out_dtype),
|
| 388 |
+
bias=bias,
|
| 389 |
+
)
|
| 390 |
+
|
| 391 |
+
return output
|
| 392 |
+
|
| 393 |
+
|
| 394 |
+
def fp8_gemm_sgl_token(A, A_scale, B, B_scale, out_dtype, bias):
|
| 395 |
+
"""GEMM function for FP8 per-token-sgl quantization using sgl-kernel.
|
| 396 |
+
|
| 397 |
+
Args:
|
| 398 |
+
A: Input activation tensor
|
| 399 |
+
A_scale: Scale tensor for input activations
|
| 400 |
+
B: Weight tensor
|
| 401 |
+
B_scale: Scale tensor for weights
|
| 402 |
+
out_dtype: Output data type.
|
| 403 |
+
bias: Optional bias tensor
|
| 404 |
+
|
| 405 |
+
Returns:
|
| 406 |
+
torch.Tensor: Result of the GEMM operation.
|
| 407 |
+
"""
|
| 408 |
+
_sgl_kernel = _ensure_sgl_kernel()
|
| 409 |
+
shape = (A.shape[0], B.shape[0])
|
| 410 |
+
output = torch.empty(shape, dtype=out_dtype, device=A.device, requires_grad=False)
|
| 411 |
+
output = _sgl_kernel.fp8_scaled_mm(
|
| 412 |
+
A,
|
| 413 |
+
B.t(),
|
| 414 |
+
A_scale,
|
| 415 |
+
B_scale.float(),
|
| 416 |
+
out_dtype,
|
| 417 |
+
bias=bias,
|
| 418 |
+
)
|
| 419 |
+
|
| 420 |
+
return output
|
| 421 |
+
|
| 422 |
+
|
| 423 |
+
def fp8_gemm(
|
| 424 |
+
A: torch.Tensor,
|
| 425 |
+
A_scale: torch.Tensor,
|
| 426 |
+
B: torch.Tensor,
|
| 427 |
+
B_scale: torch.Tensor,
|
| 428 |
+
bias: Optional[torch.Tensor],
|
| 429 |
+
out_dtype: torch.dtype,
|
| 430 |
+
native_fp8_support: bool = False,
|
| 431 |
+
quant_type: str = QuantType.FP8_PER_TENSOR,
|
| 432 |
+
origin_shape: Optional[Tuple[int, ...]] = None,
|
| 433 |
+
) -> torch.Tensor:
|
| 434 |
+
"""
|
| 435 |
+
Unified GEMM function for FP8 quantization.
|
| 436 |
+
|
| 437 |
+
Args:
|
| 438 |
+
A: Input activation tensor
|
| 439 |
+
A_scale: Scale tensor for input activations
|
| 440 |
+
B: Weight tensor
|
| 441 |
+
B_scale: Scale tensor for weights
|
| 442 |
+
bias: Optional bias tensor
|
| 443 |
+
out_dtype: Output data type
|
| 444 |
+
native_fp8_support: Whether to use native FP8 kernels
|
| 445 |
+
quant_type: Quantization type (QuantType enum values)
|
| 446 |
+
origin_shape: Original shape for reshaping output
|
| 447 |
+
|
| 448 |
+
Returns:
|
| 449 |
+
Output tensor after GEMM operation
|
| 450 |
+
|
| 451 |
+
Raises:
|
| 452 |
+
ValueError: If unsupported combination of parameters is provided
|
| 453 |
+
"""
|
| 454 |
+
# Return empty tensor if the input is empty
|
| 455 |
+
if A.numel() == 0:
|
| 456 |
+
return torch.empty(size=(0, B.shape[0]), dtype=out_dtype, device=A.device)
|
| 457 |
+
|
| 458 |
+
if quant_type == QuantType.FP8_PER_CHANNEL_VLLM:
|
| 459 |
+
if not hasattr(torch.ops, "_C") or not hasattr(torch.ops._C, "cutlass_scaled_mm"):
|
| 460 |
+
raise RuntimeError(
|
| 461 |
+
"quant_type='fp8-per-channel-vllm' requires torch.ops._C.cutlass_scaled_mm"
|
| 462 |
+
)
|
| 463 |
+
output = torch.empty(
|
| 464 |
+
(A.shape[0], B.shape[0]),
|
| 465 |
+
dtype=out_dtype,
|
| 466 |
+
device=A.device,
|
| 467 |
+
requires_grad=False,
|
| 468 |
+
)
|
| 469 |
+
weight_scale = B_scale
|
| 470 |
+
if (
|
| 471 |
+
weight_scale.dim() == 2
|
| 472 |
+
and weight_scale.shape[0] == B.shape[0]
|
| 473 |
+
and weight_scale.shape[1] == 1
|
| 474 |
+
):
|
| 475 |
+
weight_scale = weight_scale.t()
|
| 476 |
+
torch.ops._C.cutlass_scaled_mm(
|
| 477 |
+
output,
|
| 478 |
+
A,
|
| 479 |
+
B.t(),
|
| 480 |
+
A_scale,
|
| 481 |
+
weight_scale,
|
| 482 |
+
bias,
|
| 483 |
+
)
|
| 484 |
+
return output
|
| 485 |
+
|
| 486 |
+
# Prioritize native fp8 support if available
|
| 487 |
+
if native_fp8_support:
|
| 488 |
+
if quant_type in (QuantType.FP8_PER_TENSOR, QuantType.FP8_PER_TOKEN):
|
| 489 |
+
# Use torch native fp8 GEMM for per-tensor and per-token fp8 quantization
|
| 490 |
+
return fp8_gemm_torch_tensor_token(A, A_scale, B, B_scale, out_dtype, bias)
|
| 491 |
+
elif quant_type == QuantType.FP8_PER_TOKEN_SGL:
|
| 492 |
+
# Use sgl-kernel for per-token-sgl fp8 quantization
|
| 493 |
+
return fp8_gemm_sgl_token(A, A_scale, B, B_scale, out_dtype, bias)
|
| 494 |
+
elif quant_type == QuantType.FP8_PER_BLOCK:
|
| 495 |
+
# Use deepgemm accelerated blockwise fp8 GEMM
|
| 496 |
+
return fp8_gemm_deepgemm_block(A, A_scale, B, B_scale, out_dtype, bias, origin_shape)
|
| 497 |
+
else:
|
| 498 |
+
if quant_type == QuantType.FP8_PER_BLOCK:
|
| 499 |
+
# Use triton kernel for blockwise fp8 quantization
|
| 500 |
+
return fp8_gemm_triton_block(A, A_scale, B, B_scale, out_dtype, bias)
|
| 501 |
+
elif quant_type == QuantType.FP8_PER_TENSOR:
|
| 502 |
+
# Fall back to scaled linear for per-tensor fp8 without native fp8 kernel
|
| 503 |
+
return torch.nn.functional.linear(
|
| 504 |
+
A.to(out_dtype) * A_scale,
|
| 505 |
+
B.to(out_dtype) * B_scale.to(out_dtype),
|
| 506 |
+
bias=bias,
|
| 507 |
+
)
|
| 508 |
+
|
| 509 |
+
# Raise error for unsupported combination of quant_type and native_fp8_support
|
| 510 |
+
raise ValueError(
|
| 511 |
+
f"Unsupported combination: "
|
| 512 |
+
f"\n quant_type={quant_type},"
|
| 513 |
+
f"\n native_fp8_support={native_fp8_support}.\n"
|
| 514 |
+
"Supported combinations:\n"
|
| 515 |
+
" - native_fp8_support=True, "
|
| 516 |
+
"quant_type in [fp8-per-tensor, fp8-per-token,"
|
| 517 |
+
" fp8-per-block, fp8-per-token-sgl]\n"
|
| 518 |
+
" - native_fp8_support=False, "
|
| 519 |
+
"quant_type in [fp8-per-tensor, fp8-per-block]"
|
| 520 |
+
)
|
quantizer/quant/utils/__init__.py
ADDED
|
@@ -0,0 +1,37 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
from .quant_io import load_fp8_scales, load_quantized_model, save_quantized_model
|
| 16 |
+
from .utils import (
|
| 17 |
+
QuantType,
|
| 18 |
+
_compile_pattern,
|
| 19 |
+
_ensure_deep_gemm,
|
| 20 |
+
_ensure_sgl_kernel,
|
| 21 |
+
cleanup_memory,
|
| 22 |
+
replace_module,
|
| 23 |
+
should_quantize_layer,
|
| 24 |
+
)
|
| 25 |
+
|
| 26 |
+
__all__ = [
|
| 27 |
+
"load_fp8_scales",
|
| 28 |
+
"load_quantized_model",
|
| 29 |
+
"save_quantized_model",
|
| 30 |
+
"QuantType",
|
| 31 |
+
"_compile_pattern",
|
| 32 |
+
"_ensure_deep_gemm",
|
| 33 |
+
"cleanup_memory",
|
| 34 |
+
"replace_module",
|
| 35 |
+
"should_quantize_layer",
|
| 36 |
+
"_ensure_sgl_kernel",
|
| 37 |
+
]
|
quantizer/quant/utils/quant_io.py
ADDED
|
@@ -0,0 +1,149 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
import os
|
| 16 |
+
from typing import Dict, Optional, Union
|
| 17 |
+
|
| 18 |
+
import torch
|
| 19 |
+
|
| 20 |
+
__all__ = ["load_fp8_scales", "load_quantized_model", "save_quantized_model"]
|
| 21 |
+
|
| 22 |
+
|
| 23 |
+
def load_fp8_scales(
|
| 24 |
+
quant_scales: Optional[Union[str, Dict[str, torch.Tensor]]]
|
| 25 |
+
) -> Dict[str, torch.Tensor]:
|
| 26 |
+
"""Load FP8 quant scales from dict, file, or dir. Prefer .safetensors."""
|
| 27 |
+
if quant_scales is None:
|
| 28 |
+
raise ValueError("quant_scales is required")
|
| 29 |
+
|
| 30 |
+
if isinstance(quant_scales, dict) and len(quant_scales) > 0:
|
| 31 |
+
# Use provided dict
|
| 32 |
+
return quant_scales
|
| 33 |
+
|
| 34 |
+
if isinstance(quant_scales, str):
|
| 35 |
+
# Check if path is file
|
| 36 |
+
if os.path.isfile(quant_scales):
|
| 37 |
+
if quant_scales.endswith(".safetensors"):
|
| 38 |
+
import safetensors.torch
|
| 39 |
+
|
| 40 |
+
print(f"Loaded scale map from {quant_scales}")
|
| 41 |
+
return safetensors.torch.load_file(quant_scales)
|
| 42 |
+
else:
|
| 43 |
+
print(f"Loaded scale map from {quant_scales}")
|
| 44 |
+
return torch.load(quant_scales)
|
| 45 |
+
# Check if path is directory
|
| 46 |
+
if os.path.isdir(quant_scales):
|
| 47 |
+
safetensors_path = os.path.join(quant_scales, "fp8_scales.safetensors")
|
| 48 |
+
pth_path = os.path.join(quant_scales, "fp8_scales.pth")
|
| 49 |
+
if os.path.isfile(safetensors_path):
|
| 50 |
+
import safetensors.torch
|
| 51 |
+
|
| 52 |
+
print(f"Loaded scale map from {safetensors_path}")
|
| 53 |
+
return safetensors.torch.load_file(safetensors_path)
|
| 54 |
+
if os.path.isfile(pth_path):
|
| 55 |
+
print(f"Loaded scale map from {pth_path}")
|
| 56 |
+
return torch.load(pth_path)
|
| 57 |
+
raise FileNotFoundError(
|
| 58 |
+
f"Quant scale file not found: {pth_path} or {safetensors_path}"
|
| 59 |
+
)
|
| 60 |
+
raise FileNotFoundError(f"quant_scales path does not exist: {quant_scales}")
|
| 61 |
+
|
| 62 |
+
raise ValueError(f"Invalid quant_scales type: {type(quant_scales)}. Only str (path) or dict.")
|
| 63 |
+
|
| 64 |
+
|
| 65 |
+
def save_quantized_model(model: torch.nn.Module, save_path: str, fp8_scales_map: Dict):
|
| 66 |
+
"""
|
| 67 |
+
Save quantized model and scale dict to directory.
|
| 68 |
+
"""
|
| 69 |
+
import logging
|
| 70 |
+
|
| 71 |
+
logger = logging.getLogger(__name__)
|
| 72 |
+
|
| 73 |
+
if not os.path.exists(save_path):
|
| 74 |
+
try:
|
| 75 |
+
os.makedirs(save_path, exist_ok=True)
|
| 76 |
+
except Exception as e:
|
| 77 |
+
raise RuntimeError(f"Cannot create directory for save_path: {save_path}. Error: {e}")
|
| 78 |
+
|
| 79 |
+
try:
|
| 80 |
+
# If Hugging Face style, use save_pretrained
|
| 81 |
+
if hasattr(model, "save_pretrained"):
|
| 82 |
+
model.save_pretrained(save_path)
|
| 83 |
+
logger.info(f"Saved quantized model to {save_path} via save_pretrained")
|
| 84 |
+
else:
|
| 85 |
+
# Otherwise, save state_dict with safetensors
|
| 86 |
+
from safetensors.torch import save_file as safe_save
|
| 87 |
+
|
| 88 |
+
model_path = os.path.join(save_path, "model.safetensors")
|
| 89 |
+
safe_save(model.state_dict(), model_path)
|
| 90 |
+
logger.info(f"Saved state_dict to {model_path}")
|
| 91 |
+
|
| 92 |
+
# Always save scales dict
|
| 93 |
+
from safetensors.torch import save_file as safe_save
|
| 94 |
+
|
| 95 |
+
scale_save_path = os.path.join(save_path, "fp8_scales.safetensors")
|
| 96 |
+
safe_save(fp8_scales_map, scale_save_path)
|
| 97 |
+
logger.info(f"Saved scales map to {scale_save_path}")
|
| 98 |
+
|
| 99 |
+
except Exception as e:
|
| 100 |
+
raise RuntimeError(f"Failed to save model and scales map to {save_path}. Error: {e}")
|
| 101 |
+
|
| 102 |
+
|
| 103 |
+
def load_quantized_model(model_class, save_path: str, device: str = "cpu"):
|
| 104 |
+
"""
|
| 105 |
+
Load quantized model from directory.
|
| 106 |
+
"""
|
| 107 |
+
import logging
|
| 108 |
+
|
| 109 |
+
logger = logging.getLogger(__name__)
|
| 110 |
+
|
| 111 |
+
try:
|
| 112 |
+
# Try Hugging Face style first
|
| 113 |
+
if hasattr(model_class, "from_pretrained"):
|
| 114 |
+
model = model_class.from_pretrained(save_path)
|
| 115 |
+
logger.info(f"Loaded Hugging Face model from {save_path}")
|
| 116 |
+
return model
|
| 117 |
+
except Exception as e:
|
| 118 |
+
logger.warning(f"Failed to load as Hugging Face model: {e}")
|
| 119 |
+
|
| 120 |
+
try:
|
| 121 |
+
# Try safetensors file first
|
| 122 |
+
model_path = os.path.join(save_path, "model.safetensors")
|
| 123 |
+
if os.path.exists(model_path):
|
| 124 |
+
from safetensors.torch import load_file as safe_load
|
| 125 |
+
|
| 126 |
+
state_dict = safe_load(model_path, device=device)
|
| 127 |
+
model = model_class()
|
| 128 |
+
model.load_state_dict(state_dict)
|
| 129 |
+
model.to(device)
|
| 130 |
+
logger.info(f"Loaded model from {model_path} (safetensors)")
|
| 131 |
+
return model
|
| 132 |
+
|
| 133 |
+
# Try pytorch .bin next
|
| 134 |
+
model_path = os.path.join(save_path, "pytorch_model.bin")
|
| 135 |
+
if os.path.exists(model_path):
|
| 136 |
+
state_dict = torch.load(model_path, map_location=device)
|
| 137 |
+
model = model_class()
|
| 138 |
+
model.load_state_dict(state_dict)
|
| 139 |
+
model.to(device)
|
| 140 |
+
logger.info(f"Loaded model from {model_path} (pytorch)")
|
| 141 |
+
return model
|
| 142 |
+
else:
|
| 143 |
+
raise FileNotFoundError(
|
| 144 |
+
f"Model file not found at {save_path}. "
|
| 145 |
+
"Expected 'model.safetensors' or 'pytorch_model.bin'"
|
| 146 |
+
)
|
| 147 |
+
|
| 148 |
+
except Exception as e:
|
| 149 |
+
raise RuntimeError(f"Failed to load model from {save_path}. Error: {e}")
|
quantizer/quant/utils/utils.py
ADDED
|
@@ -0,0 +1,219 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2025 Tencent Inc. All Rights Reserved.
|
| 2 |
+
#
|
| 3 |
+
# Licensed under the Apache License, Version 2.0 (the "License");
|
| 4 |
+
# you may not use this file except in compliance with the License.
|
| 5 |
+
# You may obtain a copy of the License at
|
| 6 |
+
#
|
| 7 |
+
# http://www.apache.org/licenses/LICENSE-2.0
|
| 8 |
+
#
|
| 9 |
+
# Unless required by applicable law or agreed to in writing, software
|
| 10 |
+
# distributed under the License is distributed on an "AS IS" BASIS,
|
| 11 |
+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
| 12 |
+
# See the License for the specific language governing permissions and
|
| 13 |
+
# limitations under the License.
|
| 14 |
+
|
| 15 |
+
|
| 16 |
+
import gc
|
| 17 |
+
import re
|
| 18 |
+
from typing import List, Optional, Union
|
| 19 |
+
|
| 20 |
+
import torch
|
| 21 |
+
|
| 22 |
+
__all__ = [
|
| 23 |
+
"replace_module",
|
| 24 |
+
"cleanup_memory",
|
| 25 |
+
"should_quantize_layer",
|
| 26 |
+
"_compile_pattern",
|
| 27 |
+
"_ensure_deep_gemm",
|
| 28 |
+
"_ensure_sgl_kernel",
|
| 29 |
+
"QuantType",
|
| 30 |
+
]
|
| 31 |
+
|
| 32 |
+
|
| 33 |
+
class QuantType:
|
| 34 |
+
FP8_PER_TENSOR = "fp8-per-tensor"
|
| 35 |
+
FP8_PER_TOKEN = "fp8-per-token"
|
| 36 |
+
FP8_PER_BLOCK = "fp8-per-block"
|
| 37 |
+
FP8_PER_CHANNEL_VLLM = "fp8-per-channel-vllm"
|
| 38 |
+
FP8_PER_TENSOR_WEIGHT_ONLY = "fp8-per-tensor-weight-only"
|
| 39 |
+
FP8_PER_TOKEN_SGL = "fp8-per-token-sgl"
|
| 40 |
+
INT8 = "int8"
|
| 41 |
+
INT8_TORCHAO = "int8-torchao"
|
| 42 |
+
INT8_VLLM = "int8-vllm"
|
| 43 |
+
INT8_TRITON = "int8-triton"
|
| 44 |
+
INT8_SGL = "int8-sgl"
|
| 45 |
+
INT8_Q8F = "int8-q8f"
|
| 46 |
+
NVFP4 = "nvfp4"
|
| 47 |
+
MXFP4 = "mxfp4"
|
| 48 |
+
MXFP6 = "mxfp6"
|
| 49 |
+
MXFP8 = "mxfp8"
|
| 50 |
+
VALID_TYPES = [
|
| 51 |
+
FP8_PER_TENSOR,
|
| 52 |
+
FP8_PER_TOKEN,
|
| 53 |
+
FP8_PER_BLOCK,
|
| 54 |
+
FP8_PER_CHANNEL_VLLM,
|
| 55 |
+
FP8_PER_TENSOR_WEIGHT_ONLY,
|
| 56 |
+
FP8_PER_TOKEN_SGL,
|
| 57 |
+
INT8,
|
| 58 |
+
INT8_TORCHAO,
|
| 59 |
+
INT8_VLLM,
|
| 60 |
+
INT8_TRITON,
|
| 61 |
+
INT8_SGL,
|
| 62 |
+
INT8_Q8F,
|
| 63 |
+
NVFP4,
|
| 64 |
+
MXFP4,
|
| 65 |
+
MXFP6,
|
| 66 |
+
MXFP8,
|
| 67 |
+
]
|
| 68 |
+
|
| 69 |
+
@classmethod
|
| 70 |
+
def validate(cls, quant_type: str):
|
| 71 |
+
if quant_type not in cls.VALID_TYPES:
|
| 72 |
+
raise ValueError(f"Invalid quant_type: {quant_type}. Valid types: {cls.VALID_TYPES}")
|
| 73 |
+
|
| 74 |
+
|
| 75 |
+
def replace_module(model: torch.nn.Module, name: str, new_module: torch.nn.Module):
|
| 76 |
+
"""
|
| 77 |
+
Replace a submodule in the model with a new module by name.
|
| 78 |
+
"""
|
| 79 |
+
if "." in name:
|
| 80 |
+
parent_name = name.rsplit(".", 1)[0]
|
| 81 |
+
child_name = name[len(parent_name) + 1 :]
|
| 82 |
+
parent = model.get_submodule(parent_name)
|
| 83 |
+
else:
|
| 84 |
+
parent_name = ""
|
| 85 |
+
parent = model
|
| 86 |
+
child_name = name
|
| 87 |
+
|
| 88 |
+
setattr(parent, child_name, new_module)
|
| 89 |
+
|
| 90 |
+
|
| 91 |
+
def cleanup_memory():
|
| 92 |
+
"""
|
| 93 |
+
Run garbage collection and clear CUDA memory cache.
|
| 94 |
+
"""
|
| 95 |
+
gc.collect()
|
| 96 |
+
torch.cuda.empty_cache()
|
| 97 |
+
|
| 98 |
+
|
| 99 |
+
def _compile_pattern(pattern: Union[str, re.Pattern], case_sensitive: bool = False) -> re.Pattern:
|
| 100 |
+
"""
|
| 101 |
+
Compile a pattern (string or pre-compiled pattern) into a regex pattern object.
|
| 102 |
+
|
| 103 |
+
Args:
|
| 104 |
+
pattern: String pattern or already-compiled regex pattern.
|
| 105 |
+
case_sensitive: Whether the match is case sensitive.
|
| 106 |
+
|
| 107 |
+
Returns:
|
| 108 |
+
Compiled regex pattern object.
|
| 109 |
+
"""
|
| 110 |
+
if isinstance(pattern, str):
|
| 111 |
+
# If the string contains special regex characters, treat as regex.
|
| 112 |
+
if any(char in pattern for char in ".*+?^${}[]|()\\"):
|
| 113 |
+
flags = 0 if case_sensitive else re.IGNORECASE
|
| 114 |
+
return re.compile(pattern, flags)
|
| 115 |
+
else:
|
| 116 |
+
# Escape regular string so it matches literally.
|
| 117 |
+
escaped = re.escape(pattern)
|
| 118 |
+
flags = 0 if case_sensitive else re.IGNORECASE
|
| 119 |
+
return re.compile(escaped, flags)
|
| 120 |
+
else:
|
| 121 |
+
# Already a compiled pattern.
|
| 122 |
+
return pattern
|
| 123 |
+
|
| 124 |
+
|
| 125 |
+
def should_quantize_layer(
|
| 126 |
+
layer_name: str,
|
| 127 |
+
include_patterns: Optional[List[Union[str, re.Pattern]]] = None,
|
| 128 |
+
exclude_patterns: Optional[List[Union[str, re.Pattern]]] = None,
|
| 129 |
+
case_sensitive: bool = False,
|
| 130 |
+
) -> bool:
|
| 131 |
+
"""
|
| 132 |
+
Decide whether a layer should be quantized based on inclusion/exclusion patterns.
|
| 133 |
+
|
| 134 |
+
Args:
|
| 135 |
+
layer_name: Name of the layer.
|
| 136 |
+
include_patterns: List of patterns (str or regex) to include.
|
| 137 |
+
exclude_patterns: List of patterns (str or regex) to exclude.
|
| 138 |
+
case_sensitive: Whether patterns are matched case sensitively.
|
| 139 |
+
|
| 140 |
+
Returns:
|
| 141 |
+
bool: Whether this layer should be quantized.
|
| 142 |
+
|
| 143 |
+
Note:
|
| 144 |
+
String patterns are auto-detected as regex if they include special chars:
|
| 145 |
+
- If contains any . * + ? ^ $ { } [ ] | ( ) \\ -> treated as regex
|
| 146 |
+
- Otherwise, escaped and matched literally
|
| 147 |
+
"""
|
| 148 |
+
if include_patterns is None:
|
| 149 |
+
include_patterns = []
|
| 150 |
+
if exclude_patterns is None:
|
| 151 |
+
exclude_patterns = []
|
| 152 |
+
|
| 153 |
+
# Check exclusion patterns
|
| 154 |
+
for pattern in exclude_patterns:
|
| 155 |
+
compiled_pattern = _compile_pattern(pattern, case_sensitive)
|
| 156 |
+
if compiled_pattern.search(layer_name):
|
| 157 |
+
return False
|
| 158 |
+
|
| 159 |
+
# If no include patterns, default is to include all layers
|
| 160 |
+
if not include_patterns:
|
| 161 |
+
return True
|
| 162 |
+
|
| 163 |
+
# Check inclusion patterns
|
| 164 |
+
for pattern in include_patterns:
|
| 165 |
+
compiled_pattern = _compile_pattern(pattern, case_sensitive)
|
| 166 |
+
if compiled_pattern.search(layer_name):
|
| 167 |
+
return True
|
| 168 |
+
|
| 169 |
+
return False
|
| 170 |
+
|
| 171 |
+
|
| 172 |
+
_deep_gemm_cached = None
|
| 173 |
+
|
| 174 |
+
|
| 175 |
+
def _ensure_deep_gemm():
|
| 176 |
+
"""
|
| 177 |
+
Lazy, safe import of deep_gemm with process-level caching. Returns the module
|
| 178 |
+
if available, otherwise raises a clear error.
|
| 179 |
+
"""
|
| 180 |
+
global _deep_gemm_cached
|
| 181 |
+
if _deep_gemm_cached is not None:
|
| 182 |
+
return _deep_gemm_cached
|
| 183 |
+
try:
|
| 184 |
+
import deep_gemm
|
| 185 |
+
|
| 186 |
+
_deep_gemm_cached = deep_gemm
|
| 187 |
+
return _deep_gemm_cached
|
| 188 |
+
except ImportError as e:
|
| 189 |
+
raise ImportError(
|
| 190 |
+
(
|
| 191 |
+
"deep_gemm is required for 'fp8-per-block' quantization with "
|
| 192 |
+
"native_fp8_support, but was not found. Please install deep_gemm first."
|
| 193 |
+
)
|
| 194 |
+
) from e
|
| 195 |
+
|
| 196 |
+
|
| 197 |
+
_sgl_kernel_cached = None
|
| 198 |
+
|
| 199 |
+
|
| 200 |
+
def _ensure_sgl_kernel():
|
| 201 |
+
"""
|
| 202 |
+
Lazy, safe import of sgl_kernel with process-level caching. Returns the module
|
| 203 |
+
if available, otherwise raises a clear error.
|
| 204 |
+
"""
|
| 205 |
+
global _sgl_kernel_cached
|
| 206 |
+
if _sgl_kernel_cached is not None:
|
| 207 |
+
return _sgl_kernel_cached
|
| 208 |
+
try:
|
| 209 |
+
import sgl_kernel
|
| 210 |
+
|
| 211 |
+
_sgl_kernel_cached = sgl_kernel
|
| 212 |
+
return _sgl_kernel_cached
|
| 213 |
+
except ImportError as e:
|
| 214 |
+
raise ImportError(
|
| 215 |
+
(
|
| 216 |
+
"sgl_kernel is required for 'fp8-per-token-sgl' quantization with "
|
| 217 |
+
"native_fp8_support, but was not found. Please install sgl_kernel first"
|
| 218 |
+
)
|
| 219 |
+
) from e
|
requirements.txt
ADDED
|
@@ -0,0 +1,16 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
diffusers==0.31.0
|
| 2 |
+
transformers==4.57.3
|
| 3 |
+
tokenizers==0.22.1
|
| 4 |
+
accelerate==1.12.0
|
| 5 |
+
safetensors
|
| 6 |
+
opencv-python-headless
|
| 7 |
+
imageio==2.37.0
|
| 8 |
+
imageio-ffmpeg==0.6.0
|
| 9 |
+
easydict==1.13
|
| 10 |
+
ftfy==6.3.1
|
| 11 |
+
numpy==2.2.6
|
| 12 |
+
omegaconf==2.3.0
|
| 13 |
+
einops==0.8.1
|
| 14 |
+
av==13.1.0
|
| 15 |
+
sentencepiece==0.2.1
|
| 16 |
+
torchvision
|
utils/cuda_kernel_buckets.py
ADDED
|
@@ -0,0 +1,227 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
"""
|
| 2 |
+
将 PyTorch CUDA profiler 事件按内核/算子名称粗分为若干类,便于估算
|
| 3 |
+
memcpy、GEMM/Linear、Attention、Norm 等在 GPU 时间中的占比。
|
| 4 |
+
|
| 5 |
+
说明:
|
| 6 |
+
- 分类基于名称子串启发式,不同 CUDA / Triton / cuBLAS 版本下内核名会有差异;
|
| 7 |
+
- 若需更细粒度,请配合 prof.export_chrome_trace() 在 chrome://tracing 或 Perfetto 中查看。
|
| 8 |
+
"""
|
| 9 |
+
|
| 10 |
+
from __future__ import annotations
|
| 11 |
+
|
| 12 |
+
from collections import defaultdict
|
| 13 |
+
from typing import DefaultDict, Dict, Iterable, List, Tuple
|
| 14 |
+
|
| 15 |
+
# 顺序有意义:先匹配更具体的类别
|
| 16 |
+
_BUCKET_ORDER: List[str] = [
|
| 17 |
+
"memcpy_sync",
|
| 18 |
+
"attention",
|
| 19 |
+
"gemm_linear",
|
| 20 |
+
"conv",
|
| 21 |
+
"norm",
|
| 22 |
+
"elementwise_reduce",
|
| 23 |
+
"other",
|
| 24 |
+
]
|
| 25 |
+
|
| 26 |
+
|
| 27 |
+
def classify_cuda_name(name: str) -> str:
|
| 28 |
+
"""根据内核或 aten 算子名称归入粗分类。"""
|
| 29 |
+
n = (name or "").lower()
|
| 30 |
+
|
| 31 |
+
# 设备同步、拷贝、显存设置(含部分 launch 开销)
|
| 32 |
+
if any(
|
| 33 |
+
s in n
|
| 34 |
+
for s in (
|
| 35 |
+
"memcpy",
|
| 36 |
+
"memset",
|
| 37 |
+
"memgetinfo",
|
| 38 |
+
"cudaget",
|
| 39 |
+
"cudaevent",
|
| 40 |
+
"cudaoccupancy",
|
| 41 |
+
"cuda stream",
|
| 42 |
+
"cudastream",
|
| 43 |
+
"cudadevicesynchronize",
|
| 44 |
+
"cuda devicesynchronize",
|
| 45 |
+
"device synchronize",
|
| 46 |
+
"eventrecord",
|
| 47 |
+
"eventsynchronize",
|
| 48 |
+
)
|
| 49 |
+
):
|
| 50 |
+
return "memcpy_sync"
|
| 51 |
+
|
| 52 |
+
# FlashAttention / SDPA / 各类 fused attention
|
| 53 |
+
if any(
|
| 54 |
+
s in n
|
| 55 |
+
for s in (
|
| 56 |
+
"flash_attn",
|
| 57 |
+
"flash-attn",
|
| 58 |
+
"scaled_dot_product",
|
| 59 |
+
"sdpa",
|
| 60 |
+
"efficient_attention",
|
| 61 |
+
"mem_eff_attention",
|
| 62 |
+
"fused_attention",
|
| 63 |
+
"fmha",
|
| 64 |
+
"mha_default",
|
| 65 |
+
"multi_head_attention",
|
| 66 |
+
"contrib_attn",
|
| 67 |
+
"attention_forward",
|
| 68 |
+
"attention_backward",
|
| 69 |
+
"softmax",
|
| 70 |
+
)
|
| 71 |
+
):
|
| 72 |
+
# 独立 softmax 小核也可能与 attention 同桶;若需拆开可把 softmax 挪到 elementwise
|
| 73 |
+
return "attention"
|
| 74 |
+
|
| 75 |
+
# MatMul / Linear / GEMM(含 cutlass、cublas、triton、fp8)
|
| 76 |
+
if any(
|
| 77 |
+
s in n
|
| 78 |
+
for s in (
|
| 79 |
+
"gemm",
|
| 80 |
+
"cublas",
|
| 81 |
+
"cutlass",
|
| 82 |
+
"matmul",
|
| 83 |
+
"mat_mul",
|
| 84 |
+
"aten::linear",
|
| 85 |
+
"aten::mm",
|
| 86 |
+
"aten::bmm",
|
| 87 |
+
"aten::addmm",
|
| 88 |
+
"aten::matmul",
|
| 89 |
+
"scaled_mm",
|
| 90 |
+
"fp8",
|
| 91 |
+
"wmma",
|
| 92 |
+
"mma_sync",
|
| 93 |
+
"triton",
|
| 94 |
+
"dot",
|
| 95 |
+
)
|
| 96 |
+
):
|
| 97 |
+
return "gemm_linear"
|
| 98 |
+
|
| 99 |
+
if any(s in n for s in ("conv", "cudnn", "depthwise", "convolution")):
|
| 100 |
+
return "conv"
|
| 101 |
+
|
| 102 |
+
if any(
|
| 103 |
+
s in n
|
| 104 |
+
for s in (
|
| 105 |
+
"layernorm",
|
| 106 |
+
"layer_norm",
|
| 107 |
+
"rms_norm",
|
| 108 |
+
"group_norm",
|
| 109 |
+
"flash_norm",
|
| 110 |
+
"aten::layer_norm",
|
| 111 |
+
"aten::group_norm",
|
| 112 |
+
"aten::native_layer_norm",
|
| 113 |
+
)
|
| 114 |
+
):
|
| 115 |
+
return "norm"
|
| 116 |
+
|
| 117 |
+
if any(
|
| 118 |
+
s in n
|
| 119 |
+
for s in (
|
| 120 |
+
"elementwise",
|
| 121 |
+
"vectorized",
|
| 122 |
+
"unary",
|
| 123 |
+
"binary",
|
| 124 |
+
"reduce",
|
| 125 |
+
"reduction",
|
| 126 |
+
"activation",
|
| 127 |
+
"silu",
|
| 128 |
+
"gelu",
|
| 129 |
+
"swiglu",
|
| 130 |
+
"relu",
|
| 131 |
+
"aten::add",
|
| 132 |
+
"aten::mul",
|
| 133 |
+
"aten::div",
|
| 134 |
+
"aten::pow",
|
| 135 |
+
"aten::sqrt",
|
| 136 |
+
"aten::rsqrt",
|
| 137 |
+
)
|
| 138 |
+
):
|
| 139 |
+
return "elementwise_reduce"
|
| 140 |
+
|
| 141 |
+
return "other"
|
| 142 |
+
|
| 143 |
+
|
| 144 |
+
def _event_cuda_time_us(event) -> float:
|
| 145 |
+
for attr in (
|
| 146 |
+
"cuda_time_total",
|
| 147 |
+
"self_cuda_time_total",
|
| 148 |
+
"self_cuda_time_total_us",
|
| 149 |
+
):
|
| 150 |
+
v = getattr(event, attr, None)
|
| 151 |
+
if v is not None:
|
| 152 |
+
return float(v)
|
| 153 |
+
return 0.0
|
| 154 |
+
|
| 155 |
+
|
| 156 |
+
def aggregate_from_profiler(prof) -> Tuple[Dict[str, float], List[Tuple[str, float]]]:
|
| 157 |
+
"""
|
| 158 |
+
从 torch.profiler.profile 实例聚合:
|
| 159 |
+
- 返回 (bucket -> 微秒总和, 按耗时排序的 (name, us) 列表)
|
| 160 |
+
使用 prof.events() 以尽量接近 CUDA 内核级名称。
|
| 161 |
+
"""
|
| 162 |
+
bucket_us: DefaultDict[str, float] = defaultdict(float)
|
| 163 |
+
per_name: DefaultDict[str, float] = defaultdict(float)
|
| 164 |
+
|
| 165 |
+
try:
|
| 166 |
+
events = prof.events()
|
| 167 |
+
except Exception:
|
| 168 |
+
events = []
|
| 169 |
+
|
| 170 |
+
for e in events:
|
| 171 |
+
us = _event_cuda_time_us(e)
|
| 172 |
+
if us <= 0:
|
| 173 |
+
continue
|
| 174 |
+
name = getattr(e, "name", "") or ""
|
| 175 |
+
b = classify_cuda_name(name)
|
| 176 |
+
bucket_us[b] += us
|
| 177 |
+
per_name[name] += us
|
| 178 |
+
|
| 179 |
+
# 若 events() 为空,回退到 key_averages(多为 aten 级)
|
| 180 |
+
if not bucket_us and not per_name:
|
| 181 |
+
try:
|
| 182 |
+
for avg in prof.key_averages():
|
| 183 |
+
us = float(getattr(avg, "cuda_time_total", 0) or getattr(avg, "self_cuda_time_total", 0) or 0)
|
| 184 |
+
if us <= 0:
|
| 185 |
+
continue
|
| 186 |
+
name = getattr(avg, "key", "") or ""
|
| 187 |
+
b = classify_cuda_name(name)
|
| 188 |
+
bucket_us[b] += us
|
| 189 |
+
per_name[name] += us
|
| 190 |
+
except Exception:
|
| 191 |
+
pass
|
| 192 |
+
|
| 193 |
+
sorted_names = sorted(per_name.items(), key=lambda x: -x[1])
|
| 194 |
+
out = {k: bucket_us.get(k, 0.0) for k in _BUCKET_ORDER}
|
| 195 |
+
for k, v in bucket_us.items():
|
| 196 |
+
if k not in out:
|
| 197 |
+
out[k] = v
|
| 198 |
+
return out, sorted_names
|
| 199 |
+
|
| 200 |
+
|
| 201 |
+
def format_bucket_report(
|
| 202 |
+
bucket_us: Dict[str, float],
|
| 203 |
+
top_names: Iterable[Tuple[str, float]],
|
| 204 |
+
top_k: int = 25,
|
| 205 |
+
) -> str:
|
| 206 |
+
total = sum(bucket_us.values()) or 1e-9
|
| 207 |
+
lines: List[str] = []
|
| 208 |
+
lines.append("=== CUDA 时间按粗分类(微秒 / 占比)===")
|
| 209 |
+
for k in _BUCKET_ORDER:
|
| 210 |
+
if k in bucket_us and bucket_us[k] > 0:
|
| 211 |
+
us = bucket_us[k]
|
| 212 |
+
lines.append(f" {k:22s} {us:12.1f} us ({100.0 * us / total:5.1f}%)")
|
| 213 |
+
# 其它未在顺序表中的桶
|
| 214 |
+
for k, us in sorted(bucket_us.items(), key=lambda x: -x[1]):
|
| 215 |
+
if k in _BUCKET_ORDER:
|
| 216 |
+
continue
|
| 217 |
+
if us > 0:
|
| 218 |
+
lines.append(f" {k:22s} {us:12.1f} us ({100.0 * us / total:5.1f}%)")
|
| 219 |
+
lines.append(f" {'TOTAL':22s} {total:12.1f} us")
|
| 220 |
+
lines.append("")
|
| 221 |
+
lines.append(f"=== 耗时最高的 {top_k} 个事件名(微秒)===")
|
| 222 |
+
for i, (name, us) in enumerate(top_names):
|
| 223 |
+
if i >= top_k:
|
| 224 |
+
break
|
| 225 |
+
short = name if len(name) <= 120 else name[:117] + "..."
|
| 226 |
+
lines.append(f" {us:10.1f} {short}")
|
| 227 |
+
return "\n".join(lines)
|
utils/memory.py
ADDED
|
@@ -0,0 +1,135 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copied from https://github.com/lllyasviel/FramePack/tree/main/demo_utils
|
| 2 |
+
# Apache-2.0 License
|
| 3 |
+
# By lllyasviel
|
| 4 |
+
|
| 5 |
+
import torch
|
| 6 |
+
|
| 7 |
+
|
| 8 |
+
cpu = torch.device('cpu')
|
| 9 |
+
gpu = torch.device(f'cuda:{torch.cuda.current_device()}')
|
| 10 |
+
gpu_complete_modules = []
|
| 11 |
+
|
| 12 |
+
|
| 13 |
+
class DynamicSwapInstaller:
|
| 14 |
+
@staticmethod
|
| 15 |
+
def _install_module(module: torch.nn.Module, **kwargs):
|
| 16 |
+
original_class = module.__class__
|
| 17 |
+
module.__dict__['forge_backup_original_class'] = original_class
|
| 18 |
+
|
| 19 |
+
def hacked_get_attr(self, name: str):
|
| 20 |
+
if '_parameters' in self.__dict__:
|
| 21 |
+
_parameters = self.__dict__['_parameters']
|
| 22 |
+
if name in _parameters:
|
| 23 |
+
p = _parameters[name]
|
| 24 |
+
if p is None:
|
| 25 |
+
return None
|
| 26 |
+
if p.__class__ == torch.nn.Parameter:
|
| 27 |
+
return torch.nn.Parameter(p.to(**kwargs), requires_grad=p.requires_grad)
|
| 28 |
+
else:
|
| 29 |
+
return p.to(**kwargs)
|
| 30 |
+
if '_buffers' in self.__dict__:
|
| 31 |
+
_buffers = self.__dict__['_buffers']
|
| 32 |
+
if name in _buffers:
|
| 33 |
+
return _buffers[name].to(**kwargs)
|
| 34 |
+
return super(original_class, self).__getattr__(name)
|
| 35 |
+
|
| 36 |
+
module.__class__ = type('DynamicSwap_' + original_class.__name__, (original_class,), {
|
| 37 |
+
'__getattr__': hacked_get_attr,
|
| 38 |
+
})
|
| 39 |
+
|
| 40 |
+
return
|
| 41 |
+
|
| 42 |
+
@staticmethod
|
| 43 |
+
def _uninstall_module(module: torch.nn.Module):
|
| 44 |
+
if 'forge_backup_original_class' in module.__dict__:
|
| 45 |
+
module.__class__ = module.__dict__.pop('forge_backup_original_class')
|
| 46 |
+
return
|
| 47 |
+
|
| 48 |
+
@staticmethod
|
| 49 |
+
def install_model(model: torch.nn.Module, **kwargs):
|
| 50 |
+
for m in model.modules():
|
| 51 |
+
DynamicSwapInstaller._install_module(m, **kwargs)
|
| 52 |
+
return
|
| 53 |
+
|
| 54 |
+
@staticmethod
|
| 55 |
+
def uninstall_model(model: torch.nn.Module):
|
| 56 |
+
for m in model.modules():
|
| 57 |
+
DynamicSwapInstaller._uninstall_module(m)
|
| 58 |
+
return
|
| 59 |
+
|
| 60 |
+
|
| 61 |
+
def fake_diffusers_current_device(model: torch.nn.Module, target_device: torch.device):
|
| 62 |
+
if hasattr(model, 'scale_shift_table'):
|
| 63 |
+
model.scale_shift_table.data = model.scale_shift_table.data.to(target_device)
|
| 64 |
+
return
|
| 65 |
+
|
| 66 |
+
for k, p in model.named_modules():
|
| 67 |
+
if hasattr(p, 'weight'):
|
| 68 |
+
p.to(target_device)
|
| 69 |
+
return
|
| 70 |
+
|
| 71 |
+
|
| 72 |
+
def get_cuda_free_memory_gb(device=None):
|
| 73 |
+
if device is None:
|
| 74 |
+
device = gpu
|
| 75 |
+
|
| 76 |
+
memory_stats = torch.cuda.memory_stats(device)
|
| 77 |
+
bytes_active = memory_stats['active_bytes.all.current']
|
| 78 |
+
bytes_reserved = memory_stats['reserved_bytes.all.current']
|
| 79 |
+
bytes_free_cuda, _ = torch.cuda.mem_get_info(device)
|
| 80 |
+
bytes_inactive_reserved = bytes_reserved - bytes_active
|
| 81 |
+
bytes_total_available = bytes_free_cuda + bytes_inactive_reserved
|
| 82 |
+
return bytes_total_available / (1024 ** 3)
|
| 83 |
+
|
| 84 |
+
|
| 85 |
+
def move_model_to_device_with_memory_preservation(model, target_device, preserved_memory_gb=0):
|
| 86 |
+
print(f'Moving {model.__class__.__name__} to {target_device} with preserved memory: {preserved_memory_gb} GB')
|
| 87 |
+
|
| 88 |
+
for m in model.modules():
|
| 89 |
+
if get_cuda_free_memory_gb(target_device) <= preserved_memory_gb:
|
| 90 |
+
torch.cuda.empty_cache()
|
| 91 |
+
return
|
| 92 |
+
|
| 93 |
+
if hasattr(m, 'weight'):
|
| 94 |
+
m.to(device=target_device)
|
| 95 |
+
|
| 96 |
+
model.to(device=target_device)
|
| 97 |
+
torch.cuda.empty_cache()
|
| 98 |
+
return
|
| 99 |
+
|
| 100 |
+
|
| 101 |
+
def offload_model_from_device_for_memory_preservation(model, target_device, preserved_memory_gb=0):
|
| 102 |
+
print(f'Offloading {model.__class__.__name__} from {target_device} to preserve memory: {preserved_memory_gb} GB')
|
| 103 |
+
|
| 104 |
+
for m in model.modules():
|
| 105 |
+
if get_cuda_free_memory_gb(target_device) >= preserved_memory_gb:
|
| 106 |
+
torch.cuda.empty_cache()
|
| 107 |
+
return
|
| 108 |
+
|
| 109 |
+
if hasattr(m, 'weight'):
|
| 110 |
+
m.to(device=cpu)
|
| 111 |
+
|
| 112 |
+
model.to(device=cpu)
|
| 113 |
+
torch.cuda.empty_cache()
|
| 114 |
+
return
|
| 115 |
+
|
| 116 |
+
|
| 117 |
+
def unload_complete_models(*args):
|
| 118 |
+
for m in gpu_complete_modules + list(args):
|
| 119 |
+
m.to(device=cpu)
|
| 120 |
+
print(f'Unloaded {m.__class__.__name__} as complete.')
|
| 121 |
+
|
| 122 |
+
gpu_complete_modules.clear()
|
| 123 |
+
torch.cuda.empty_cache()
|
| 124 |
+
return
|
| 125 |
+
|
| 126 |
+
|
| 127 |
+
def load_model_as_complete(model, target_device, unload=True):
|
| 128 |
+
if unload:
|
| 129 |
+
unload_complete_models()
|
| 130 |
+
|
| 131 |
+
model.to(device=target_device)
|
| 132 |
+
print(f'Loaded {model.__class__.__name__} to {target_device} as complete.')
|
| 133 |
+
|
| 134 |
+
gpu_complete_modules.append(model)
|
| 135 |
+
return
|
utils/misc.py
ADDED
|
@@ -0,0 +1,39 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
import numpy as np
|
| 2 |
+
import random
|
| 3 |
+
import torch
|
| 4 |
+
|
| 5 |
+
|
| 6 |
+
def set_seed(seed: int, deterministic: bool = False):
|
| 7 |
+
"""
|
| 8 |
+
Helper function for reproducible behavior to set the seed in `random`, `numpy`, `torch`.
|
| 9 |
+
|
| 10 |
+
Args:
|
| 11 |
+
seed (`int`):
|
| 12 |
+
The seed to set.
|
| 13 |
+
deterministic (`bool`, *optional*, defaults to `False`):
|
| 14 |
+
Whether to use deterministic algorithms where available. Can slow down training.
|
| 15 |
+
"""
|
| 16 |
+
random.seed(seed)
|
| 17 |
+
np.random.seed(seed)
|
| 18 |
+
torch.manual_seed(seed)
|
| 19 |
+
torch.cuda.manual_seed_all(seed)
|
| 20 |
+
|
| 21 |
+
if deterministic:
|
| 22 |
+
torch.use_deterministic_algorithms(True)
|
| 23 |
+
|
| 24 |
+
|
| 25 |
+
def merge_dict_list(dict_list):
|
| 26 |
+
if len(dict_list) == 1:
|
| 27 |
+
return dict_list[0]
|
| 28 |
+
|
| 29 |
+
merged_dict = {}
|
| 30 |
+
for k, v in dict_list[0].items():
|
| 31 |
+
if isinstance(v, torch.Tensor):
|
| 32 |
+
if v.ndim == 0:
|
| 33 |
+
merged_dict[k] = torch.stack([d[k] for d in dict_list], dim=0)
|
| 34 |
+
else:
|
| 35 |
+
merged_dict[k] = torch.cat([d[k] for d in dict_list], dim=0)
|
| 36 |
+
else:
|
| 37 |
+
# for non-tensor values, we just copy the value from the first item
|
| 38 |
+
merged_dict[k] = v
|
| 39 |
+
return merged_dict
|
utils/scheduler.py
ADDED
|
@@ -0,0 +1,194 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
from abc import abstractmethod, ABC
|
| 2 |
+
import torch
|
| 3 |
+
|
| 4 |
+
|
| 5 |
+
class SchedulerInterface(ABC):
|
| 6 |
+
"""
|
| 7 |
+
Base class for diffusion noise schedule.
|
| 8 |
+
"""
|
| 9 |
+
alphas_cumprod: torch.Tensor # [T], alphas for defining the noise schedule
|
| 10 |
+
|
| 11 |
+
@abstractmethod
|
| 12 |
+
def add_noise(
|
| 13 |
+
self, clean_latent: torch.Tensor,
|
| 14 |
+
noise: torch.Tensor, timestep: torch.Tensor
|
| 15 |
+
):
|
| 16 |
+
"""
|
| 17 |
+
Diffusion forward corruption process.
|
| 18 |
+
Input:
|
| 19 |
+
- clean_latent: the clean latent with shape [B, C, H, W]
|
| 20 |
+
- noise: the noise with shape [B, C, H, W]
|
| 21 |
+
- timestep: the timestep with shape [B]
|
| 22 |
+
Output: the corrupted latent with shape [B, C, H, W]
|
| 23 |
+
"""
|
| 24 |
+
pass
|
| 25 |
+
|
| 26 |
+
def convert_x0_to_noise(
|
| 27 |
+
self, x0: torch.Tensor, xt: torch.Tensor,
|
| 28 |
+
timestep: torch.Tensor
|
| 29 |
+
) -> torch.Tensor:
|
| 30 |
+
"""
|
| 31 |
+
Convert the diffusion network's x0 prediction to noise predidction.
|
| 32 |
+
x0: the predicted clean data with shape [B, C, H, W]
|
| 33 |
+
xt: the input noisy data with shape [B, C, H, W]
|
| 34 |
+
timestep: the timestep with shape [B]
|
| 35 |
+
|
| 36 |
+
noise = (xt-sqrt(alpha_t)*x0) / sqrt(beta_t) (eq 11 in https://arxiv.org/abs/2311.18828)
|
| 37 |
+
"""
|
| 38 |
+
# use higher precision for calculations
|
| 39 |
+
original_dtype = x0.dtype
|
| 40 |
+
x0, xt, alphas_cumprod = map(
|
| 41 |
+
lambda x: x.double().to(x0.device), [x0, xt,
|
| 42 |
+
self.alphas_cumprod]
|
| 43 |
+
)
|
| 44 |
+
|
| 45 |
+
alpha_prod_t = alphas_cumprod[timestep].reshape(-1, 1, 1, 1)
|
| 46 |
+
beta_prod_t = 1 - alpha_prod_t
|
| 47 |
+
|
| 48 |
+
noise_pred = (xt - alpha_prod_t **
|
| 49 |
+
(0.5) * x0) / beta_prod_t ** (0.5)
|
| 50 |
+
return noise_pred.to(original_dtype)
|
| 51 |
+
|
| 52 |
+
def convert_noise_to_x0(
|
| 53 |
+
self, noise: torch.Tensor, xt: torch.Tensor,
|
| 54 |
+
timestep: torch.Tensor
|
| 55 |
+
) -> torch.Tensor:
|
| 56 |
+
"""
|
| 57 |
+
Convert the diffusion network's noise prediction to x0 predidction.
|
| 58 |
+
noise: the predicted noise with shape [B, C, H, W]
|
| 59 |
+
xt: the input noisy data with shape [B, C, H, W]
|
| 60 |
+
timestep: the timestep with shape [B]
|
| 61 |
+
|
| 62 |
+
x0 = (x_t - sqrt(beta_t) * noise) / sqrt(alpha_t) (eq 11 in https://arxiv.org/abs/2311.18828)
|
| 63 |
+
"""
|
| 64 |
+
# use higher precision for calculations
|
| 65 |
+
original_dtype = noise.dtype
|
| 66 |
+
noise, xt, alphas_cumprod = map(
|
| 67 |
+
lambda x: x.double().to(noise.device), [noise, xt,
|
| 68 |
+
self.alphas_cumprod]
|
| 69 |
+
)
|
| 70 |
+
alpha_prod_t = alphas_cumprod[timestep].reshape(-1, 1, 1, 1)
|
| 71 |
+
beta_prod_t = 1 - alpha_prod_t
|
| 72 |
+
|
| 73 |
+
x0_pred = (xt - beta_prod_t **
|
| 74 |
+
(0.5) * noise) / alpha_prod_t ** (0.5)
|
| 75 |
+
return x0_pred.to(original_dtype)
|
| 76 |
+
|
| 77 |
+
def convert_velocity_to_x0(
|
| 78 |
+
self, velocity: torch.Tensor, xt: torch.Tensor,
|
| 79 |
+
timestep: torch.Tensor
|
| 80 |
+
) -> torch.Tensor:
|
| 81 |
+
"""
|
| 82 |
+
Convert the diffusion network's velocity prediction to x0 predidction.
|
| 83 |
+
velocity: the predicted noise with shape [B, C, H, W]
|
| 84 |
+
xt: the input noisy data with shape [B, C, H, W]
|
| 85 |
+
timestep: the timestep with shape [B]
|
| 86 |
+
|
| 87 |
+
v = sqrt(alpha_t) * noise - sqrt(beta_t) x0
|
| 88 |
+
noise = (xt-sqrt(alpha_t)*x0) / sqrt(beta_t)
|
| 89 |
+
given v, x_t, we have
|
| 90 |
+
x0 = sqrt(alpha_t) * x_t - sqrt(beta_t) * v
|
| 91 |
+
see derivations https://chatgpt.com/share/679fb6c8-3a30-8008-9b0e-d1ae892dac56
|
| 92 |
+
"""
|
| 93 |
+
# use higher precision for calculations
|
| 94 |
+
original_dtype = velocity.dtype
|
| 95 |
+
velocity, xt, alphas_cumprod = map(
|
| 96 |
+
lambda x: x.double().to(velocity.device), [velocity, xt,
|
| 97 |
+
self.alphas_cumprod]
|
| 98 |
+
)
|
| 99 |
+
alpha_prod_t = alphas_cumprod[timestep].reshape(-1, 1, 1, 1)
|
| 100 |
+
beta_prod_t = 1 - alpha_prod_t
|
| 101 |
+
|
| 102 |
+
x0_pred = (alpha_prod_t ** 0.5) * xt - (beta_prod_t ** 0.5) * velocity
|
| 103 |
+
return x0_pred.to(original_dtype)
|
| 104 |
+
|
| 105 |
+
|
| 106 |
+
class FlowMatchScheduler():
|
| 107 |
+
|
| 108 |
+
def __init__(self, num_inference_steps=100, num_train_timesteps=1000, shift=3.0, sigma_max=1.0, sigma_min=0.003 / 1.002, inverse_timesteps=False, extra_one_step=False, reverse_sigmas=False):
|
| 109 |
+
self.num_train_timesteps = num_train_timesteps
|
| 110 |
+
self.shift = shift
|
| 111 |
+
self.sigma_max = sigma_max
|
| 112 |
+
self.sigma_min = sigma_min
|
| 113 |
+
self.inverse_timesteps = inverse_timesteps
|
| 114 |
+
self.extra_one_step = extra_one_step
|
| 115 |
+
self.reverse_sigmas = reverse_sigmas
|
| 116 |
+
self.set_timesteps(num_inference_steps)
|
| 117 |
+
|
| 118 |
+
def set_timesteps(self, num_inference_steps=100, denoising_strength=1.0, training=False):
|
| 119 |
+
sigma_start = self.sigma_min + \
|
| 120 |
+
(self.sigma_max - self.sigma_min) * denoising_strength
|
| 121 |
+
if self.extra_one_step:
|
| 122 |
+
self.sigmas = torch.linspace(
|
| 123 |
+
sigma_start, self.sigma_min, num_inference_steps + 1)[:-1]
|
| 124 |
+
else:
|
| 125 |
+
self.sigmas = torch.linspace(
|
| 126 |
+
sigma_start, self.sigma_min, num_inference_steps)
|
| 127 |
+
if self.inverse_timesteps:
|
| 128 |
+
self.sigmas = torch.flip(self.sigmas, dims=[0])
|
| 129 |
+
self.sigmas = self.shift * self.sigmas / \
|
| 130 |
+
(1 + (self.shift - 1) * self.sigmas)
|
| 131 |
+
if self.reverse_sigmas:
|
| 132 |
+
self.sigmas = 1 - self.sigmas
|
| 133 |
+
self.timesteps = self.sigmas * self.num_train_timesteps
|
| 134 |
+
if training:
|
| 135 |
+
x = self.timesteps
|
| 136 |
+
y = torch.exp(-2 * ((x - num_inference_steps / 2) /
|
| 137 |
+
num_inference_steps) ** 2)
|
| 138 |
+
y_shifted = y - y.min()
|
| 139 |
+
bsmntw_weighing = y_shifted * \
|
| 140 |
+
(num_inference_steps / y_shifted.sum())
|
| 141 |
+
self.linear_timesteps_weights = bsmntw_weighing
|
| 142 |
+
|
| 143 |
+
def step(self, model_output, timestep, sample, to_final=False):
|
| 144 |
+
if timestep.ndim == 2:
|
| 145 |
+
timestep = timestep.flatten(0, 1)
|
| 146 |
+
self.sigmas = self.sigmas.to(model_output.device)
|
| 147 |
+
self.timesteps = self.timesteps.to(model_output.device)
|
| 148 |
+
timestep_id = torch.argmin(
|
| 149 |
+
(self.timesteps.unsqueeze(0) - timestep.unsqueeze(1)).abs(), dim=1)
|
| 150 |
+
sigma = self.sigmas[timestep_id].reshape(-1, 1, 1, 1)
|
| 151 |
+
if to_final or (timestep_id + 1 >= len(self.timesteps)).any():
|
| 152 |
+
sigma_ = 1 if (
|
| 153 |
+
self.inverse_timesteps or self.reverse_sigmas) else 0
|
| 154 |
+
else:
|
| 155 |
+
sigma_ = self.sigmas[timestep_id + 1].reshape(-1, 1, 1, 1)
|
| 156 |
+
prev_sample = sample + model_output * (sigma_ - sigma)
|
| 157 |
+
return prev_sample
|
| 158 |
+
|
| 159 |
+
def add_noise(self, original_samples, noise, timestep):
|
| 160 |
+
"""
|
| 161 |
+
Diffusion forward corruption process.
|
| 162 |
+
Input:
|
| 163 |
+
- clean_latent: the clean latent with shape [B*T, C, H, W]
|
| 164 |
+
- noise: the noise with shape [B*T, C, H, W]
|
| 165 |
+
- timestep: the timestep with shape [B*T]
|
| 166 |
+
Output: the corrupted latent with shape [B*T, C, H, W]
|
| 167 |
+
"""
|
| 168 |
+
if timestep.ndim == 2:
|
| 169 |
+
timestep = timestep.flatten(0, 1)
|
| 170 |
+
self.sigmas = self.sigmas.to(noise.device)
|
| 171 |
+
self.timesteps = self.timesteps.to(noise.device)
|
| 172 |
+
timestep_id = torch.argmin(
|
| 173 |
+
(self.timesteps.unsqueeze(0) - timestep.unsqueeze(1)).abs(), dim=1)
|
| 174 |
+
sigma = self.sigmas[timestep_id].reshape(-1, 1, 1, 1)
|
| 175 |
+
sample = (1 - sigma) * original_samples + sigma * noise
|
| 176 |
+
return sample.type_as(noise)
|
| 177 |
+
|
| 178 |
+
def training_target(self, sample, noise, timestep):
|
| 179 |
+
target = noise - sample
|
| 180 |
+
return target
|
| 181 |
+
|
| 182 |
+
def training_weight(self, timestep):
|
| 183 |
+
"""
|
| 184 |
+
Input:
|
| 185 |
+
- timestep: the timestep with shape [B*T]
|
| 186 |
+
Output: the corresponding weighting [B*T]
|
| 187 |
+
"""
|
| 188 |
+
if timestep.ndim == 2:
|
| 189 |
+
timestep = timestep.flatten(0, 1)
|
| 190 |
+
self.linear_timesteps_weights = self.linear_timesteps_weights.to(timestep.device)
|
| 191 |
+
timestep_id = torch.argmin(
|
| 192 |
+
(self.timesteps.unsqueeze(1) - timestep.unsqueeze(0)).abs(), dim=0)
|
| 193 |
+
weights = self.linear_timesteps_weights[timestep_id]
|
| 194 |
+
return weights
|
utils/wan_wrapper.py
ADDED
|
@@ -0,0 +1,871 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
import types
|
| 2 |
+
from pathlib import Path
|
| 3 |
+
from typing import List, Optional
|
| 4 |
+
import torch
|
| 5 |
+
from torch import nn
|
| 6 |
+
from safetensors.torch import load_file as load_safetensors_file
|
| 7 |
+
|
| 8 |
+
from utils.scheduler import SchedulerInterface, FlowMatchScheduler
|
| 9 |
+
import os
|
| 10 |
+
|
| 11 |
+
|
| 12 |
+
def _wan_models_path(*parts) -> str:
|
| 13 |
+
"""Resolve wan_models path relative to project root (works with symlink and any cwd)."""
|
| 14 |
+
root = Path(__file__).resolve().parent.parent
|
| 15 |
+
return str((root / "wan_models").joinpath(Path(*parts)))
|
| 16 |
+
|
| 17 |
+
|
| 18 |
+
def _resolve_wan_path(path: str) -> str:
|
| 19 |
+
"""If path starts with wan_models/, resolve to absolute path (project root); else return as-is."""
|
| 20 |
+
if path and path.startswith("wan_models/"):
|
| 21 |
+
return _wan_models_path(path[len("wan_models/"):])
|
| 22 |
+
return path
|
| 23 |
+
|
| 24 |
+
|
| 25 |
+
def _resolve_wan_path_with_dir(path: str, wan_models_dir: Optional[str] = None) -> str:
|
| 26 |
+
"""Resolve path: if wan_models_dir is set and path starts with wan_models/, use wan_models_dir as base; else _resolve_wan_path."""
|
| 27 |
+
if not path:
|
| 28 |
+
return path
|
| 29 |
+
if wan_models_dir and path.startswith("wan_models/"):
|
| 30 |
+
return os.path.join(wan_models_dir, path[len("wan_models/"):])
|
| 31 |
+
return _resolve_wan_path(path)
|
| 32 |
+
|
| 33 |
+
|
| 34 |
+
def model_kwargs_with_relative_rope(args, default: bool = False) -> dict:
|
| 35 |
+
"""Merge top-level use_relative_rope into model_kwargs with a stable default."""
|
| 36 |
+
raw_model_kwargs = getattr(args, "model_kwargs", {}) or {}
|
| 37 |
+
model_kwargs = dict(raw_model_kwargs)
|
| 38 |
+
if "use_relative_rope" not in model_kwargs:
|
| 39 |
+
try:
|
| 40 |
+
model_kwargs["use_relative_rope"] = bool(getattr(args, "use_relative_rope"))
|
| 41 |
+
except Exception:
|
| 42 |
+
model_kwargs["use_relative_rope"] = bool(default)
|
| 43 |
+
return model_kwargs
|
| 44 |
+
|
| 45 |
+
from wan.modules.tokenizers import HuggingfaceTokenizer
|
| 46 |
+
from wan.modules.model import WanModel
|
| 47 |
+
from wan.modules.t5 import umt5_xxl
|
| 48 |
+
from wan.modules.causal_model import CausalWanModel
|
| 49 |
+
|
| 50 |
+
class WanTextEncoder(torch.nn.Module):
|
| 51 |
+
def __init__(
|
| 52 |
+
self,
|
| 53 |
+
tokenizer_path="wan_models/Wan2.1-T2V-1.3B/google/umt5-xxl/",
|
| 54 |
+
encoder_pth_path="wan_models/Wan2.1-T2V-1.3B/models_t5_umt5-xxl-enc-bf16.pth",
|
| 55 |
+
) -> None:
|
| 56 |
+
super().__init__()
|
| 57 |
+
# tokenizer_path = _resolve_wan_path_with_dir(tokenizer_path, wan_models_dir)
|
| 58 |
+
# encoder_pth_path = _resolve_wan_path_with_dir(encoder_pth_path, wan_models_dir)
|
| 59 |
+
|
| 60 |
+
self.text_encoder = umt5_xxl(
|
| 61 |
+
encoder_only=True,
|
| 62 |
+
return_tokenizer=False,
|
| 63 |
+
dtype=torch.bfloat16,
|
| 64 |
+
device=torch.device('cpu')
|
| 65 |
+
).eval().requires_grad_(False)
|
| 66 |
+
state_dict = torch.load(encoder_pth_path,
|
| 67 |
+
map_location='cpu', weights_only=False)
|
| 68 |
+
self.text_encoder.load_state_dict(state_dict)
|
| 69 |
+
del state_dict
|
| 70 |
+
|
| 71 |
+
self.tokenizer = HuggingfaceTokenizer(
|
| 72 |
+
name=tokenizer_path, seq_len=512, clean='whitespace')
|
| 73 |
+
|
| 74 |
+
@property
|
| 75 |
+
def device(self):
|
| 76 |
+
return next(self.text_encoder.parameters()).device
|
| 77 |
+
|
| 78 |
+
def forward(self, text_prompts: List[str], device: torch.device = None) -> dict:
|
| 79 |
+
ids, mask = self.tokenizer(
|
| 80 |
+
text_prompts, return_mask=True, add_special_tokens=True)
|
| 81 |
+
# When DynamicSwapInstaller is active, self.device returns cpu because
|
| 82 |
+
# parameters are swapped to GPU only during forward. Use the explicitly
|
| 83 |
+
# passed device (the intended execution device) when available.
|
| 84 |
+
target_device = device if device is not None else self.device
|
| 85 |
+
ids = ids.to(target_device)
|
| 86 |
+
mask = mask.to(target_device)
|
| 87 |
+
seq_lens = mask.gt(0).sum(dim=1).long()
|
| 88 |
+
context = self.text_encoder(ids, mask)
|
| 89 |
+
|
| 90 |
+
for u, v in zip(context, seq_lens):
|
| 91 |
+
u[v:] = 0.0 # set padding to 0.0
|
| 92 |
+
|
| 93 |
+
return {
|
| 94 |
+
"prompt_embeds": context
|
| 95 |
+
}
|
| 96 |
+
|
| 97 |
+
|
| 98 |
+
class WanVAEWrapper(torch.nn.Module):
|
| 99 |
+
def __init__(
|
| 100 |
+
self,
|
| 101 |
+
pretrained_path=None,
|
| 102 |
+
z_dim=48,
|
| 103 |
+
vae_type="Wan2.2_VAE",
|
| 104 |
+
wan_models_dir=None,
|
| 105 |
+
):
|
| 106 |
+
super().__init__()
|
| 107 |
+
if vae_type != "Wan2.2_VAE":
|
| 108 |
+
raise ValueError(f"Unsupported vae_type={vae_type!r}; only 'Wan2.2_VAE' is supported.")
|
| 109 |
+
from wan.modules.vae2_2 import _video_vae
|
| 110 |
+
self.mean = torch.tensor([
|
| 111 |
+
-0.2289, -0.0052, -0.1323, -0.2339, -0.2799, 0.0174, 0.1838, 0.1557,
|
| 112 |
+
-0.1382, 0.0542, 0.2813, 0.0891, 0.1570, -0.0098, 0.0375, -0.1825,
|
| 113 |
+
-0.2246, -0.1207, -0.0698, 0.5109, 0.2665, -0.2108, -0.2158, 0.2502,
|
| 114 |
+
-0.2055, -0.0322, 0.1109, 0.1567, -0.0729, 0.0899, -0.2799, -0.1230,
|
| 115 |
+
-0.0313, -0.1649, 0.0117, 0.0723, -0.2839, -0.2083, -0.0520, 0.3748,
|
| 116 |
+
0.0152, 0.1957, 0.1433, -0.2944, 0.3573, -0.0548, -0.1681, -0.0667,
|
| 117 |
+
], dtype=torch.float32)
|
| 118 |
+
|
| 119 |
+
self.std = torch.tensor([
|
| 120 |
+
0.4765, 1.0364, 0.4514, 1.1677, 0.5313, 0.4990, 0.4818, 0.5013,
|
| 121 |
+
0.8158, 1.0344, 0.5894, 1.0901, 0.6885, 0.6165, 0.8454, 0.4978,
|
| 122 |
+
0.5759, 0.3523, 0.7135, 0.6804, 0.5833, 1.4146, 0.8986, 0.5659,
|
| 123 |
+
0.7069, 0.5338, 0.4889, 0.4917, 0.4069, 0.4999, 0.6866, 0.4093,
|
| 124 |
+
0.5709, 0.6065, 0.6415, 0.4944, 0.5726, 1.2042, 0.5458, 1.6887,
|
| 125 |
+
0.3971, 1.0600, 0.3943, 0.5537, 0.5444, 0.4089, 0.7468, 0.7744,
|
| 126 |
+
], dtype=torch.float32)
|
| 127 |
+
self.scale = [self.mean, 1.0 / self.std]
|
| 128 |
+
self.upsampling_factor = 16
|
| 129 |
+
|
| 130 |
+
z_dim = 48
|
| 131 |
+
self.z_dim = z_dim
|
| 132 |
+
self.model = _video_vae(pretrained_path=pretrained_path,
|
| 133 |
+
z_dim=z_dim,).eval().requires_grad_(False)
|
| 134 |
+
|
| 135 |
+
def generate_noise(self, shape, seed=None, rand_device="cpu", rand_torch_dtype=torch.float32, device=None, torch_dtype=None):
|
| 136 |
+
# Initialize Gaussian noise
|
| 137 |
+
generator = None if seed is None else torch.Generator(rand_device).manual_seed(seed)
|
| 138 |
+
noise = torch.randn(shape, generator=generator, device=rand_device, dtype=rand_torch_dtype)
|
| 139 |
+
noise = noise.to(dtype=torch_dtype, device=device)
|
| 140 |
+
return noise
|
| 141 |
+
|
| 142 |
+
def encode_to_latent(self, pixel: torch.Tensor) -> torch.Tensor:
|
| 143 |
+
# pixel: [batch_size, num_channels, num_frames, height, width]
|
| 144 |
+
device, dtype = pixel.device, pixel.dtype
|
| 145 |
+
scale = [self.mean.to(device=device, dtype=dtype),
|
| 146 |
+
1.0 / self.std.to(device=device, dtype=dtype)]
|
| 147 |
+
|
| 148 |
+
output = [
|
| 149 |
+
self.model.encode(u.unsqueeze(0), scale).float().squeeze(0)
|
| 150 |
+
for u in pixel
|
| 151 |
+
]
|
| 152 |
+
output = torch.stack(output, dim=0)
|
| 153 |
+
# from [batch_size, num_channels, num_frames, height, width]
|
| 154 |
+
# to [batch_size, num_frames, num_channels, height, width]
|
| 155 |
+
output = output.permute(0, 2, 1, 3, 4)
|
| 156 |
+
return output
|
| 157 |
+
|
| 158 |
+
def decode_to_pixel(self, latent: torch.Tensor, use_cache: bool = False, return_in_cpu: bool = False) -> torch.Tensor:
|
| 159 |
+
# from [batch_size, num_frames, num_channels, height, width]
|
| 160 |
+
# to [batch_size, num_channels, num_frames, height, width]
|
| 161 |
+
zs = latent.permute(0, 2, 1, 3, 4)
|
| 162 |
+
if use_cache:
|
| 163 |
+
assert latent.shape[0] == 1, "Batch size must be 1 when using cache"
|
| 164 |
+
|
| 165 |
+
device, dtype = latent.device, latent.dtype
|
| 166 |
+
scale = [self.mean.to(device=device, dtype=dtype),
|
| 167 |
+
1.0 / self.std.to(device=device, dtype=dtype)]
|
| 168 |
+
|
| 169 |
+
if use_cache:
|
| 170 |
+
decode_function = self.model.cached_decode
|
| 171 |
+
else:
|
| 172 |
+
decode_function = self.model.decode
|
| 173 |
+
|
| 174 |
+
output = []
|
| 175 |
+
for u in zs:
|
| 176 |
+
decoded = decode_function(u.unsqueeze(0), scale).float().clamp_(-1, 1).squeeze(0)
|
| 177 |
+
if return_in_cpu:
|
| 178 |
+
decoded = decoded.cpu()
|
| 179 |
+
output.append(decoded)
|
| 180 |
+
output = torch.stack(output, dim=0)
|
| 181 |
+
# from [batch_size, num_channels, num_frames, height, width]
|
| 182 |
+
# to [batch_size, num_frames, num_channels, height, width]
|
| 183 |
+
output = output.permute(0, 2, 1, 3, 4)
|
| 184 |
+
return output
|
| 185 |
+
|
| 186 |
+
|
| 187 |
+
class TAEW2_2VAEWrapper(torch.nn.Module):
|
| 188 |
+
"""
|
| 189 |
+
VAE wrapper using TAEHV (TAEW2.2) for faster decoding.
|
| 190 |
+
Requires: pip install taehv (or install from https://github.com/madebyollin/taehv)
|
| 191 |
+
Checkpoint: taew2_2.pth (download from taehv releases)
|
| 192 |
+
"""
|
| 193 |
+
def __init__(self, checkpoint_path: str = "taew2_2.pth", dtype=torch.float16):
|
| 194 |
+
super().__init__()
|
| 195 |
+
try:
|
| 196 |
+
from wan.modules.taehv import TAEHV, StreamingTAEHV
|
| 197 |
+
except ImportError as e:
|
| 198 |
+
raise ImportError(
|
| 199 |
+
"taehv is required for TAEW2.2 VAE. Install with: pip install taehv"
|
| 200 |
+
) from e
|
| 201 |
+
self.taehv = TAEHV(checkpoint_path).to(dtype).eval().requires_grad_(False)
|
| 202 |
+
self.taehv = StreamingTAEHV(self.taehv)
|
| 203 |
+
self.dtype = dtype
|
| 204 |
+
# For compatibility with pipeline.vae.model.clear_cache()
|
| 205 |
+
self.model = _TAEW2_2ModelRef(self)
|
| 206 |
+
|
| 207 |
+
def warmup_first_frame(self, first_frame_latent: torch.Tensor):
|
| 208 |
+
"""Warm up the streaming decoder's MemBlock memory with the first-frame latent.
|
| 209 |
+
|
| 210 |
+
The TAeW2.2 MemBlocks use zero-initialized past context for the first frame,
|
| 211 |
+
causing blur. By feeding the first-frame latent as a warmup pass (output
|
| 212 |
+
discarded), subsequent decodes benefit from real temporal context.
|
| 213 |
+
|
| 214 |
+
Args:
|
| 215 |
+
first_frame_latent: [B, 1, C, H, W] latent of the first frame.
|
| 216 |
+
"""
|
| 217 |
+
if first_frame_latent is None:
|
| 218 |
+
return
|
| 219 |
+
# Reset decoder state, then feed first frame as warmup
|
| 220 |
+
self.taehv.reset()
|
| 221 |
+
with torch.no_grad(), torch.autocast(device_type="cuda", dtype=self.dtype):
|
| 222 |
+
# Feed the first-frame latent to populate MemBlock memory;
|
| 223 |
+
# the output (startup frames) is discarded.
|
| 224 |
+
_ = self.taehv.decode(first_frame_latent)
|
| 225 |
+
|
| 226 |
+
def decode_to_pixel(
|
| 227 |
+
self,
|
| 228 |
+
latent: torch.Tensor,
|
| 229 |
+
use_cache: bool = False,
|
| 230 |
+
return_in_cpu: bool = False
|
| 231 |
+
) -> torch.Tensor:
|
| 232 |
+
# latent: [B, F, C, H, W] = [B, T, C, H, W] (same as TAEHV's NTCHW)
|
| 233 |
+
# use_cache=True -> parallel=False for lower memory (streaming)
|
| 234 |
+
parallel = not use_cache
|
| 235 |
+
with torch.autocast(device_type="cuda", dtype=self.dtype):
|
| 236 |
+
# out = self.taehv.decode_video(
|
| 237 |
+
# latent, parallel=parallel, show_progress_bar=False
|
| 238 |
+
# )
|
| 239 |
+
out = self.taehv.decode(latent)
|
| 240 |
+
# TAEHV returns [0, 1], convert to [-1, 1] to match WanVAEWrapper
|
| 241 |
+
out = out.mul(2).sub(1).clamp(-1, 1).float()
|
| 242 |
+
if return_in_cpu:
|
| 243 |
+
out = out.cpu()
|
| 244 |
+
return out
|
| 245 |
+
|
| 246 |
+
|
| 247 |
+
class _TAEW2_2ModelRef:
|
| 248 |
+
"""Dummy ref for clear_cache compatibility; delegates to StreamingTAEHV.reset()."""
|
| 249 |
+
|
| 250 |
+
def __init__(self, parent):
|
| 251 |
+
self._parent = parent
|
| 252 |
+
|
| 253 |
+
def clear_cache(self):
|
| 254 |
+
self._parent.taehv.reset()
|
| 255 |
+
|
| 256 |
+
|
| 257 |
+
class MGLightVAEWrapper(torch.nn.Module):
|
| 258 |
+
"""VAE wrapper using MG-LightVAE (pruned Wan2.2 VAE) for faster decoding.
|
| 259 |
+
|
| 260 |
+
Wraps the ``Wan2_2_VAE`` class which supports different pruning rates.
|
| 261 |
+
The encoder uses the full (unpruned) Wan2.2 VAE teacher, while the decoder
|
| 262 |
+
uses the pruned student model.
|
| 263 |
+
|
| 264 |
+
Args:
|
| 265 |
+
vae_pth: Path to the pruned LightVAE checkpoint (student decoder).
|
| 266 |
+
lightvae_pruning_rate: Pruning rate for the decoder (e.g. 0.5, 0.75).
|
| 267 |
+
lightvae_encoder_vae_pth: Path to the full Wan2.2 VAE checkpoint
|
| 268 |
+
(teacher encoder). Required for mg_lightvae.
|
| 269 |
+
dtype: Data type for the VAE model.
|
| 270 |
+
device: Device to load the VAE on.
|
| 271 |
+
"""
|
| 272 |
+
|
| 273 |
+
def __init__(
|
| 274 |
+
self,
|
| 275 |
+
vae_pth: str,
|
| 276 |
+
lightvae_pruning_rate: float = 0.75,
|
| 277 |
+
lightvae_encoder_vae_pth: str | None = None,
|
| 278 |
+
dtype=torch.float,
|
| 279 |
+
device="cpu",
|
| 280 |
+
):
|
| 281 |
+
super().__init__()
|
| 282 |
+
from wan.modules.vae2_2 import Wan2_2_VAE
|
| 283 |
+
|
| 284 |
+
self._vae = Wan2_2_VAE(
|
| 285 |
+
z_dim=48,
|
| 286 |
+
c_dim=160,
|
| 287 |
+
vae_pth=vae_pth,
|
| 288 |
+
dtype=dtype,
|
| 289 |
+
device=device,
|
| 290 |
+
vae_type="mg_lightvae",
|
| 291 |
+
lightvae_pruning_rate=lightvae_pruning_rate,
|
| 292 |
+
lightvae_encoder_vae_pth=lightvae_encoder_vae_pth,
|
| 293 |
+
)
|
| 294 |
+
# Register model (pruned decoder) and encoder_model (teacher encoder)
|
| 295 |
+
# as submodules so that .to(), .eval(), .requires_grad_() propagate.
|
| 296 |
+
self.model = self._vae.model
|
| 297 |
+
if self._vae.encoder_model is not None:
|
| 298 |
+
self.encoder_model = self._vae.encoder_model
|
| 299 |
+
else:
|
| 300 |
+
self.encoder_model = None
|
| 301 |
+
|
| 302 |
+
self.z_dim = 48
|
| 303 |
+
self.upsampling_factor = 16
|
| 304 |
+
self.mean = self._vae.scale[0]
|
| 305 |
+
self.std = 1.0 / self._vae.scale[1]
|
| 306 |
+
|
| 307 |
+
# Initialize streaming cache attributes (_feat_map, _conv_idx, etc.)
|
| 308 |
+
# so cached_decode() can be called before any explicit clear_cache().
|
| 309 |
+
self.model.clear_cache()
|
| 310 |
+
if self.encoder_model is not None:
|
| 311 |
+
self.encoder_model.clear_cache()
|
| 312 |
+
|
| 313 |
+
def generate_noise(self, shape, seed=None, rand_device="cpu",
|
| 314 |
+
rand_torch_dtype=torch.float32, device=None, torch_dtype=None):
|
| 315 |
+
generator = None if seed is None else torch.Generator(rand_device).manual_seed(seed)
|
| 316 |
+
noise = torch.randn(shape, generator=generator, device=rand_device, dtype=rand_torch_dtype)
|
| 317 |
+
noise = noise.to(dtype=torch_dtype, device=device)
|
| 318 |
+
return noise
|
| 319 |
+
|
| 320 |
+
def encode_to_latent(self, pixel: torch.Tensor) -> torch.Tensor:
|
| 321 |
+
# pixel: [batch_size, num_channels, num_frames, height, width]
|
| 322 |
+
device, dtype = pixel.device, pixel.dtype
|
| 323 |
+
scale = [self.mean.to(device=device, dtype=dtype),
|
| 324 |
+
1.0 / self.std.to(device=device, dtype=dtype)]
|
| 325 |
+
|
| 326 |
+
encode_model = self.encoder_model if self.encoder_model is not None else self.model
|
| 327 |
+
output = [
|
| 328 |
+
encode_model.encode(u.unsqueeze(0), scale).float().squeeze(0)
|
| 329 |
+
for u in pixel
|
| 330 |
+
]
|
| 331 |
+
output = torch.stack(output, dim=0)
|
| 332 |
+
# from [B, C, F, H, W] to [B, F, C, H, W]
|
| 333 |
+
output = output.permute(0, 2, 1, 3, 4)
|
| 334 |
+
return output
|
| 335 |
+
|
| 336 |
+
def decode_to_pixel(self, latent: torch.Tensor, use_cache: bool = False,
|
| 337 |
+
return_in_cpu: bool = False) -> torch.Tensor:
|
| 338 |
+
# from [B, F, C, H, W] to [B, C, F, H, W]
|
| 339 |
+
zs = latent.permute(0, 2, 1, 3, 4)
|
| 340 |
+
if use_cache:
|
| 341 |
+
assert latent.shape[0] == 1, "Batch size must be 1 when using cache"
|
| 342 |
+
|
| 343 |
+
device, dtype = latent.device, latent.dtype
|
| 344 |
+
scale = [self.mean.to(device=device, dtype=dtype),
|
| 345 |
+
1.0 / self.std.to(device=device, dtype=dtype)]
|
| 346 |
+
|
| 347 |
+
if use_cache:
|
| 348 |
+
decode_function = self.model.cached_decode
|
| 349 |
+
else:
|
| 350 |
+
decode_function = self.model.decode
|
| 351 |
+
|
| 352 |
+
output = []
|
| 353 |
+
for u in zs:
|
| 354 |
+
decoded = decode_function(u.unsqueeze(0), scale).float().clamp_(-1, 1).squeeze(0)
|
| 355 |
+
if return_in_cpu:
|
| 356 |
+
decoded = decoded.cpu()
|
| 357 |
+
output.append(decoded)
|
| 358 |
+
output = torch.stack(output, dim=0)
|
| 359 |
+
# from [B, C, F, H, W] to [B, F, C, H, W]
|
| 360 |
+
output = output.permute(0, 2, 1, 3, 4)
|
| 361 |
+
return output
|
| 362 |
+
|
| 363 |
+
|
| 364 |
+
def create_vae_from_config(config) -> Optional[torch.nn.Module]:
|
| 365 |
+
"""Create a VAE wrapper based on the unified ``vae_type`` config field.
|
| 366 |
+
|
| 367 |
+
Supported vae_type values:
|
| 368 |
+
- "wan2.2": Standard Wan2.2 VAE (returns None, pipeline creates WanVAEWrapper)
|
| 369 |
+
- "taew2_2": TAeW2.2/taehv fast streaming decoder
|
| 370 |
+
- "mg_lightvae": MG-LightVAE pruned decoder (pruning rate 0.5)
|
| 371 |
+
- "mg_lightvae_v2": MG-LightVAE v2 pruned decoder (pruning rate 0.75)
|
| 372 |
+
|
| 373 |
+
If ``vae_type`` is not set, defaults to "taew2_2".
|
| 374 |
+
|
| 375 |
+
Returns:
|
| 376 |
+
A VAE wrapper instance, or None for wan2.2 (let pipeline create
|
| 377 |
+
WanVAEWrapper from vae_kwargs).
|
| 378 |
+
"""
|
| 379 |
+
vae_type = getattr(config, "vae_type", None)
|
| 380 |
+
|
| 381 |
+
if vae_type is None:
|
| 382 |
+
vae_type = "taew2_2"
|
| 383 |
+
else:
|
| 384 |
+
vae_type = str(vae_type).strip().lower()
|
| 385 |
+
|
| 386 |
+
if vae_type == "wan2.2":
|
| 387 |
+
return None # pipeline creates WanVAEWrapper(**vae_kwargs) internally
|
| 388 |
+
|
| 389 |
+
if vae_type == "taew2_2":
|
| 390 |
+
ckpt = os.environ.get("TAEW2_2_CHECKPOINT") or getattr(
|
| 391 |
+
config, "taew2_2_checkpoint", "taew2_2.pth"
|
| 392 |
+
)
|
| 393 |
+
return TAEW2_2VAEWrapper(checkpoint_path=ckpt).eval()
|
| 394 |
+
|
| 395 |
+
if vae_type in ("mg_lightvae", "mg_lightvae_v2"):
|
| 396 |
+
pruning_map = {"mg_lightvae": 0.5, "mg_lightvae_v2": 0.75}
|
| 397 |
+
# Explicit pruning rate overrides the default mapping
|
| 398 |
+
explicit_rate = getattr(config, "lightvae_pruning_rate", None)
|
| 399 |
+
if explicit_rate is not None:
|
| 400 |
+
pruning_rate = float(explicit_rate)
|
| 401 |
+
else:
|
| 402 |
+
pruning_rate = pruning_map[vae_type]
|
| 403 |
+
|
| 404 |
+
# Select checkpoint based on vae_type
|
| 405 |
+
ckpt_map = {
|
| 406 |
+
"mg_lightvae": "lightvae_checkpoint",
|
| 407 |
+
"mg_lightvae_v2": "lightvae_v2_checkpoint",
|
| 408 |
+
}
|
| 409 |
+
vae_ckpt = getattr(config, ckpt_map[vae_type], None)
|
| 410 |
+
if vae_ckpt is None:
|
| 411 |
+
raise ValueError(
|
| 412 |
+
f"vae_type={vae_type!r} requires '{ckpt_map[vae_type]}' config field "
|
| 413 |
+
f"(path to MG-LightVAE .pth file)."
|
| 414 |
+
)
|
| 415 |
+
|
| 416 |
+
# Encoder checkpoint: explicit config, or fall back to vae_kwargs.pretrained_path
|
| 417 |
+
encoder_ckpt = getattr(config, "lightvae_encoder_checkpoint", None)
|
| 418 |
+
if encoder_ckpt is None:
|
| 419 |
+
vae_kwargs = getattr(config, "vae_kwargs", {}) or {}
|
| 420 |
+
if isinstance(vae_kwargs, dict):
|
| 421 |
+
encoder_ckpt = vae_kwargs.get("pretrained_path")
|
| 422 |
+
else:
|
| 423 |
+
encoder_ckpt = getattr(vae_kwargs, "pretrained_path", None)
|
| 424 |
+
if encoder_ckpt is None:
|
| 425 |
+
raise ValueError(
|
| 426 |
+
f"vae_type={vae_type!r} requires 'lightvae_encoder_checkpoint' config field "
|
| 427 |
+
f"(path to full Wan2.2_VAE.pth for teacher encoder), "
|
| 428 |
+
f"or 'vae_kwargs.pretrained_path' must be set."
|
| 429 |
+
)
|
| 430 |
+
|
| 431 |
+
return MGLightVAEWrapper(
|
| 432 |
+
vae_pth=vae_ckpt,
|
| 433 |
+
lightvae_pruning_rate=pruning_rate,
|
| 434 |
+
lightvae_encoder_vae_pth=encoder_ckpt,
|
| 435 |
+
)
|
| 436 |
+
|
| 437 |
+
raise ValueError(
|
| 438 |
+
f"Unsupported vae_type={vae_type!r}. "
|
| 439 |
+
f"Choose from: wan2.2, taew2_2, mg_lightvae, mg_lightvae_v2."
|
| 440 |
+
)
|
| 441 |
+
|
| 442 |
+
|
| 443 |
+
class WanDiffusionWrapper(torch.nn.Module):
|
| 444 |
+
@staticmethod
|
| 445 |
+
def _materialize_meta_tensors(module: torch.nn.Module, device: torch.device = torch.device("cpu")):
|
| 446 |
+
materialized_names = []
|
| 447 |
+
|
| 448 |
+
def _materialize_recursive(mod: torch.nn.Module, prefix: str = ""):
|
| 449 |
+
for name, param in list(mod.named_parameters(recurse=False)):
|
| 450 |
+
if getattr(param, "is_meta", False):
|
| 451 |
+
new_param = torch.nn.Parameter(
|
| 452 |
+
torch.empty(tuple(param.shape), dtype=param.dtype, device=device),
|
| 453 |
+
requires_grad=param.requires_grad,
|
| 454 |
+
)
|
| 455 |
+
setattr(mod, name, new_param)
|
| 456 |
+
materialized_names.append(prefix + name)
|
| 457 |
+
|
| 458 |
+
for name, buf in list(mod.named_buffers(recurse=False)):
|
| 459 |
+
if getattr(buf, "is_meta", False):
|
| 460 |
+
setattr(mod, name, torch.empty(tuple(buf.shape), dtype=buf.dtype, device=device))
|
| 461 |
+
materialized_names.append(prefix + name)
|
| 462 |
+
|
| 463 |
+
for child_name, child in mod.named_children():
|
| 464 |
+
_materialize_recursive(child, prefix + child_name + ".")
|
| 465 |
+
|
| 466 |
+
_materialize_recursive(module)
|
| 467 |
+
return materialized_names
|
| 468 |
+
|
| 469 |
+
@staticmethod
|
| 470 |
+
def _normalize_model_state_dict_keys(state_dict: dict[str, torch.Tensor]) -> dict[str, torch.Tensor]:
|
| 471 |
+
key_prefixes = (
|
| 472 |
+
"generator.model._fsdp_wrapped_module.",
|
| 473 |
+
"generator.model.",
|
| 474 |
+
"model._fsdp_wrapped_module.",
|
| 475 |
+
"model.",
|
| 476 |
+
"_fsdp_wrapped_module.",
|
| 477 |
+
"module.",
|
| 478 |
+
)
|
| 479 |
+
keys = list(state_dict.keys())
|
| 480 |
+
for prefix in key_prefixes:
|
| 481 |
+
if keys and all(k.startswith(prefix) for k in keys):
|
| 482 |
+
return {k[len(prefix):]: v for k, v in state_dict.items()}
|
| 483 |
+
return state_dict
|
| 484 |
+
|
| 485 |
+
def _load_model_safetensors(self, model_safetensors_path: str) -> None:
|
| 486 |
+
if model_safetensors_path.startswith("oss://"):
|
| 487 |
+
raise ValueError(
|
| 488 |
+
"model_safetensors_path must be a local mounted path on AI-Hub, "
|
| 489 |
+
f"got {model_safetensors_path}"
|
| 490 |
+
)
|
| 491 |
+
|
| 492 |
+
model_safetensors_path = _resolve_wan_path(model_safetensors_path)
|
| 493 |
+
print(f"[WanDiffusionWrapper] Loading model safetensors from {model_safetensors_path}")
|
| 494 |
+
state_dict = load_safetensors_file(model_safetensors_path, device="cpu")
|
| 495 |
+
state_dict = self._normalize_model_state_dict_keys(state_dict)
|
| 496 |
+
|
| 497 |
+
model_keys = set(self.model.state_dict().keys())
|
| 498 |
+
matched_keys = model_keys.intersection(state_dict.keys())
|
| 499 |
+
if not matched_keys:
|
| 500 |
+
sample_keys = list(state_dict.keys())[:10]
|
| 501 |
+
raise ValueError(
|
| 502 |
+
"No safetensors keys matched the Wan model state_dict. "
|
| 503 |
+
f"First loaded keys: {sample_keys}"
|
| 504 |
+
)
|
| 505 |
+
|
| 506 |
+
match_ratio = len(matched_keys) / max(1, len(state_dict))
|
| 507 |
+
if match_ratio < 0.5:
|
| 508 |
+
sample_unexpected = [k for k in state_dict.keys() if k not in model_keys][:10]
|
| 509 |
+
raise ValueError(
|
| 510 |
+
f"Only {len(matched_keys)}/{len(state_dict)} safetensors keys match the Wan model "
|
| 511 |
+
f"state_dict after prefix normalization. Sample unexpected keys: {sample_unexpected}"
|
| 512 |
+
)
|
| 513 |
+
|
| 514 |
+
missing, unexpected = self.model.load_state_dict(state_dict, strict=False)
|
| 515 |
+
if missing:
|
| 516 |
+
print(
|
| 517 |
+
f"[WanDiffusionWrapper] model_safetensors missing {len(missing)} keys "
|
| 518 |
+
f"(showing first 20): {missing[:20]}"
|
| 519 |
+
)
|
| 520 |
+
if unexpected:
|
| 521 |
+
print(
|
| 522 |
+
f"[WanDiffusionWrapper] model_safetensors unexpected {len(unexpected)} keys "
|
| 523 |
+
f"(showing first 20): {unexpected[:20]}"
|
| 524 |
+
)
|
| 525 |
+
print(
|
| 526 |
+
f"[WanDiffusionWrapper] Loaded model safetensors with {len(matched_keys)} "
|
| 527 |
+
f"matched keys from {model_safetensors_path}"
|
| 528 |
+
)
|
| 529 |
+
|
| 530 |
+
def __init__(
|
| 531 |
+
self,
|
| 532 |
+
model_name="Wan2.1-T2V-1.3B",
|
| 533 |
+
timestep_shift=8.0,
|
| 534 |
+
is_causal=False,
|
| 535 |
+
local_attn_size=-1,
|
| 536 |
+
sink_size=0,
|
| 537 |
+
subfolder=None,
|
| 538 |
+
model_type='t2v',
|
| 539 |
+
num_frame_per_block=3,
|
| 540 |
+
model_safetensors_path: Optional[str] = None,
|
| 541 |
+
**model_init_kwargs,
|
| 542 |
+
):
|
| 543 |
+
super().__init__()
|
| 544 |
+
|
| 545 |
+
self.model_type = model_type
|
| 546 |
+
use_relative_rope = bool(model_init_kwargs.pop("use_relative_rope", False))
|
| 547 |
+
if is_causal:
|
| 548 |
+
model_init_kwargs["use_relative_rope"] = use_relative_rope
|
| 549 |
+
self.model = CausalWanModel.from_pretrained(
|
| 550 |
+
model_name, local_attn_size=local_attn_size, sink_size=sink_size, model_type=model_type, num_frame_per_block=num_frame_per_block,
|
| 551 |
+
**model_init_kwargs)
|
| 552 |
+
else:
|
| 553 |
+
if use_relative_rope:
|
| 554 |
+
print("[WanDiffusionWrapper] use_relative_rope is ignored for non-causal WanModel.")
|
| 555 |
+
self.model = WanModel.from_pretrained(model_name, model_type=model_type, **model_init_kwargs)
|
| 556 |
+
materialized = self._materialize_meta_tensors(self.model, device=torch.device("cpu"))
|
| 557 |
+
if materialized:
|
| 558 |
+
print(f"[WanDiffusionWrapper] Materialized {len(materialized)} meta tensors on CPU.")
|
| 559 |
+
if model_safetensors_path:
|
| 560 |
+
self._load_model_safetensors(model_safetensors_path)
|
| 561 |
+
self.model.eval()
|
| 562 |
+
|
| 563 |
+
# For non-causal diffusion, all frames share the same timestep
|
| 564 |
+
self.uniform_timestep = not is_causal
|
| 565 |
+
|
| 566 |
+
self.scheduler = FlowMatchScheduler(
|
| 567 |
+
shift=timestep_shift, sigma_min=0.0, extra_one_step=True
|
| 568 |
+
)
|
| 569 |
+
self.scheduler.set_timesteps(1000, training=True)
|
| 570 |
+
|
| 571 |
+
self.seq_len = None # [1, 21, 16, 60, 104]
|
| 572 |
+
self.post_init()
|
| 573 |
+
|
| 574 |
+
def enable_gradient_checkpointing(self) -> None:
|
| 575 |
+
self.model.enable_gradient_checkpointing()
|
| 576 |
+
|
| 577 |
+
def adding_cls_branch(self, atten_dim=1536, num_class=4, time_embed_dim=0) -> None:
|
| 578 |
+
# NOTE: This is hard coded for WAN2.1-T2V-1.3B for now!!!!!!!!!!!!!!!!!!!!
|
| 579 |
+
self._cls_pred_branch = nn.Sequential(
|
| 580 |
+
# Input: [B, 384, 21, 60, 104]
|
| 581 |
+
nn.LayerNorm(atten_dim * 3 + time_embed_dim),
|
| 582 |
+
nn.Linear(atten_dim * 3 + time_embed_dim, 1536),
|
| 583 |
+
nn.SiLU(),
|
| 584 |
+
nn.Linear(atten_dim, num_class)
|
| 585 |
+
)
|
| 586 |
+
self._cls_pred_branch.requires_grad_(True)
|
| 587 |
+
num_registers = 3
|
| 588 |
+
self._register_tokens = RegisterTokens(num_registers=num_registers, dim=atten_dim)
|
| 589 |
+
self._register_tokens.requires_grad_(True)
|
| 590 |
+
|
| 591 |
+
gan_ca_blocks = []
|
| 592 |
+
for _ in range(num_registers):
|
| 593 |
+
block = GanAttentionBlock()
|
| 594 |
+
gan_ca_blocks.append(block)
|
| 595 |
+
self._gan_ca_blocks = nn.ModuleList(gan_ca_blocks)
|
| 596 |
+
self._gan_ca_blocks.requires_grad_(True)
|
| 597 |
+
# self.has_cls_branch = True
|
| 598 |
+
|
| 599 |
+
def _convert_flow_pred_to_x0(self, flow_pred: torch.Tensor, xt: torch.Tensor, timestep: torch.Tensor) -> torch.Tensor:
|
| 600 |
+
"""
|
| 601 |
+
Convert flow matching's prediction to x0 prediction.
|
| 602 |
+
flow_pred: the prediction with shape [B, C, H, W]
|
| 603 |
+
xt: the input noisy data with shape [B, C, H, W]
|
| 604 |
+
timestep: the timestep with shape [B]
|
| 605 |
+
|
| 606 |
+
pred = noise - x0
|
| 607 |
+
x_t = (1-sigma_t) * x0 + sigma_t * noise
|
| 608 |
+
we have x0 = x_t - sigma_t * pred
|
| 609 |
+
see derivations https://chatgpt.com/share/67bf8589-3d04-8008-bc6e-4cf1a24e2d0e
|
| 610 |
+
"""
|
| 611 |
+
# use higher precision for calculations
|
| 612 |
+
original_dtype = flow_pred.dtype
|
| 613 |
+
flow_pred, xt, sigmas, timesteps = map(
|
| 614 |
+
lambda x: x.double().to(flow_pred.device), [flow_pred, xt,
|
| 615 |
+
self.scheduler.sigmas,
|
| 616 |
+
self.scheduler.timesteps]
|
| 617 |
+
)
|
| 618 |
+
|
| 619 |
+
timestep_id = torch.argmin(
|
| 620 |
+
(timesteps.unsqueeze(0) - timestep.unsqueeze(1)).abs(), dim=1)
|
| 621 |
+
sigma_t = sigmas[timestep_id].reshape(-1, 1, 1, 1)
|
| 622 |
+
x0_pred = xt - sigma_t * flow_pred
|
| 623 |
+
return x0_pred.to(original_dtype)
|
| 624 |
+
|
| 625 |
+
@staticmethod
|
| 626 |
+
def _convert_x0_to_flow_pred(scheduler, x0_pred: torch.Tensor, xt: torch.Tensor, timestep: torch.Tensor) -> torch.Tensor:
|
| 627 |
+
"""
|
| 628 |
+
Convert x0 prediction to flow matching's prediction.
|
| 629 |
+
x0_pred: the x0 prediction with shape [B, C, H, W]
|
| 630 |
+
xt: the input noisy data with shape [B, C, H, W]
|
| 631 |
+
timestep: the timestep with shape [B]
|
| 632 |
+
|
| 633 |
+
pred = (x_t - x_0) / sigma_t
|
| 634 |
+
"""
|
| 635 |
+
# use higher precision for calculations
|
| 636 |
+
original_dtype = x0_pred.dtype
|
| 637 |
+
x0_pred, xt, sigmas, timesteps = map(
|
| 638 |
+
lambda x: x.double().to(x0_pred.device), [x0_pred, xt,
|
| 639 |
+
scheduler.sigmas,
|
| 640 |
+
scheduler.timesteps]
|
| 641 |
+
)
|
| 642 |
+
timestep_id = torch.argmin(
|
| 643 |
+
(timesteps.unsqueeze(0) - timestep.unsqueeze(1)).abs(), dim=1)
|
| 644 |
+
sigma_t = sigmas[timestep_id].reshape(-1, 1, 1, 1)
|
| 645 |
+
flow_pred = (xt - x0_pred) / sigma_t
|
| 646 |
+
return flow_pred.to(original_dtype)
|
| 647 |
+
|
| 648 |
+
@staticmethod
|
| 649 |
+
def _history_x_to_model_format(history_x):
|
| 650 |
+
if history_x is None:
|
| 651 |
+
return None
|
| 652 |
+
|
| 653 |
+
if torch.is_tensor(history_x):
|
| 654 |
+
if history_x.ndim != 5:
|
| 655 |
+
raise ValueError(
|
| 656 |
+
f"history_x must be [B,F,C,H,W] when passed as a tensor, got {history_x.shape}"
|
| 657 |
+
)
|
| 658 |
+
return [u.permute(1, 0, 2, 3).contiguous() for u in history_x]
|
| 659 |
+
|
| 660 |
+
return history_x
|
| 661 |
+
|
| 662 |
+
@staticmethod
|
| 663 |
+
def _history_condition_to_model_format(value):
|
| 664 |
+
if value is None:
|
| 665 |
+
return None
|
| 666 |
+
|
| 667 |
+
if torch.is_tensor(value):
|
| 668 |
+
if value.ndim < 3:
|
| 669 |
+
return value
|
| 670 |
+
return [u.contiguous() for u in value]
|
| 671 |
+
|
| 672 |
+
return value
|
| 673 |
+
|
| 674 |
+
def forward(
|
| 675 |
+
self,
|
| 676 |
+
noisy_image_or_video: torch.Tensor, conditional_dict: dict,
|
| 677 |
+
timestep: torch.Tensor, kv_cache: Optional[List[dict]] = None,
|
| 678 |
+
crossattn_cache: Optional[List[dict]] = None,
|
| 679 |
+
current_start: Optional[int] = None,
|
| 680 |
+
|
| 681 |
+
classify_mode: Optional[bool] = False, # DF
|
| 682 |
+
concat_time_embeddings: Optional[bool] = False, #DF
|
| 683 |
+
clean_x: Optional[torch.Tensor] = None, # TF
|
| 684 |
+
aug_t: Optional[torch.Tensor] = None, # for TF clean GT, if it's also noisy and needs denoising by the model, aug_t is its timestep
|
| 685 |
+
|
| 686 |
+
cache_start: Optional[int] = None,
|
| 687 |
+
updating_cache: Optional[bool] = False,
|
| 688 |
+
replace_first_timestep_and_noise_latents: Optional[bool] = False,
|
| 689 |
+
history_x: Optional[torch.Tensor] = None,
|
| 690 |
+
history_y: Optional[torch.Tensor] = None,
|
| 691 |
+
history_act_context: Optional[torch.Tensor] = None,
|
| 692 |
+
history_y_action: Optional[torch.Tensor] = None,
|
| 693 |
+
noisy_start_frame: int = 0,
|
| 694 |
+
|
| 695 |
+
) -> torch.Tensor:
|
| 696 |
+
prompt_embeds = conditional_dict["prompt_embeds"]
|
| 697 |
+
act_context = conditional_dict.get("act_context", None)
|
| 698 |
+
act_context_scale = conditional_dict.get("act_context_scale", 1.0)
|
| 699 |
+
clip_fea = conditional_dict.get("clip_fea", None)
|
| 700 |
+
y = conditional_dict.get("y", None)
|
| 701 |
+
y_action = conditional_dict.get("y_action", None)
|
| 702 |
+
ref_latents = conditional_dict.get("ref_latents", None)
|
| 703 |
+
ref_mask = conditional_dict.get("ref_mask", None)
|
| 704 |
+
# first_frame_latents = conditional_dict.get("first_frame_latents", None)
|
| 705 |
+
|
| 706 |
+
raw_timestep = timestep
|
| 707 |
+
b, f, c, h, w = noisy_image_or_video.shape
|
| 708 |
+
|
| 709 |
+
if replace_first_timestep_and_noise_latents:
|
| 710 |
+
# Wan2.2 5B uses the first latent frame as a clean condition. Keep
|
| 711 |
+
# per-frame timesteps for score models so only frame 0 is forced to t=0.
|
| 712 |
+
if raw_timestep.dim() == 2:
|
| 713 |
+
input_timestep = raw_timestep.clone()
|
| 714 |
+
input_timestep[:, 0] = 0
|
| 715 |
+
elif raw_timestep.dim() == 1 and raw_timestep.shape[0] == f:
|
| 716 |
+
input_timestep = raw_timestep.unsqueeze(0).repeat(b, 1)
|
| 717 |
+
input_timestep[:, 0] = 0
|
| 718 |
+
elif raw_timestep.dim() == 1 and raw_timestep.shape[0] == b:
|
| 719 |
+
input_timestep = raw_timestep[:, None].repeat(1, f)
|
| 720 |
+
input_timestep[:, 0] = 0
|
| 721 |
+
else:
|
| 722 |
+
input_timestep = raw_timestep.reshape(-1)[0].view(1, 1).repeat(b, f)
|
| 723 |
+
input_timestep[:, 0] = 0
|
| 724 |
+
elif self.uniform_timestep:
|
| 725 |
+
# [B, F] -> [B] for legacy non-causal uniform score models.
|
| 726 |
+
input_timestep = raw_timestep[:, 0]
|
| 727 |
+
else:
|
| 728 |
+
input_timestep = raw_timestep
|
| 729 |
+
|
| 730 |
+
logits = None
|
| 731 |
+
if history_x is not None:
|
| 732 |
+
history_kwargs = {
|
| 733 |
+
"history_x": self._history_x_to_model_format(history_x),
|
| 734 |
+
"noisy_start_frame": int(noisy_start_frame),
|
| 735 |
+
}
|
| 736 |
+
history_y = self._history_condition_to_model_format(history_y)
|
| 737 |
+
history_act_context = self._history_condition_to_model_format(history_act_context)
|
| 738 |
+
history_y_action = self._history_condition_to_model_format(history_y_action)
|
| 739 |
+
|
| 740 |
+
if history_y is not None:
|
| 741 |
+
history_kwargs["history_y"] = history_y
|
| 742 |
+
if history_act_context is not None:
|
| 743 |
+
history_kwargs["history_act_context"] = history_act_context
|
| 744 |
+
if history_y_action is not None:
|
| 745 |
+
history_kwargs["history_y_action"] = history_y_action
|
| 746 |
+
|
| 747 |
+
model_out = self.model(
|
| 748 |
+
noisy_image_or_video.permute(0, 2, 1, 3, 4),
|
| 749 |
+
t=input_timestep,
|
| 750 |
+
context=prompt_embeds,
|
| 751 |
+
seq_len=self.seq_len,
|
| 752 |
+
kv_cache=None,
|
| 753 |
+
crossattn_cache=crossattn_cache,
|
| 754 |
+
current_start=0 if current_start is None else current_start,
|
| 755 |
+
cache_start=0 if cache_start is None else cache_start,
|
| 756 |
+
act_context=act_context,
|
| 757 |
+
y_action=y_action,
|
| 758 |
+
act_context_scale=act_context_scale,
|
| 759 |
+
clip_fea=clip_fea,
|
| 760 |
+
y=y,
|
| 761 |
+
ref_latents=ref_latents,
|
| 762 |
+
ref_mask=ref_mask,
|
| 763 |
+
**history_kwargs,
|
| 764 |
+
)
|
| 765 |
+
if isinstance(model_out, tuple):
|
| 766 |
+
flow_pred = model_out[0]
|
| 767 |
+
else:
|
| 768 |
+
flow_pred = model_out
|
| 769 |
+
flow_pred = flow_pred.permute(0, 2, 1, 3, 4)
|
| 770 |
+
# X0 prediction
|
| 771 |
+
elif kv_cache is not None:
|
| 772 |
+
kwargs = {}
|
| 773 |
+
if updating_cache:
|
| 774 |
+
kwargs["updating_cache"] = updating_cache
|
| 775 |
+
flow_pred = self.model(
|
| 776 |
+
noisy_image_or_video.permute(0, 2, 1, 3, 4), # => [B, C, F, H, W],
|
| 777 |
+
t=input_timestep, context=prompt_embeds,
|
| 778 |
+
seq_len=self.seq_len,
|
| 779 |
+
kv_cache=kv_cache,
|
| 780 |
+
crossattn_cache=crossattn_cache,
|
| 781 |
+
current_start=current_start,
|
| 782 |
+
cache_start=cache_start,
|
| 783 |
+
act_context=act_context,
|
| 784 |
+
act_context_scale=act_context_scale,
|
| 785 |
+
clip_fea=clip_fea,
|
| 786 |
+
y=y,
|
| 787 |
+
ref_latents=ref_latents,
|
| 788 |
+
ref_mask=ref_mask,
|
| 789 |
+
**kwargs,
|
| 790 |
+
).permute(0, 2, 1, 3, 4)
|
| 791 |
+
else:
|
| 792 |
+
if clean_x is not None:
|
| 793 |
+
# teacher forcing
|
| 794 |
+
flow_pred = self.model(
|
| 795 |
+
noisy_image_or_video.permute(0, 2, 1, 3, 4), # => [B, C, F, H, W]
|
| 796 |
+
t=input_timestep, context=prompt_embeds,
|
| 797 |
+
seq_len=self.seq_len,
|
| 798 |
+
clean_x=clean_x.permute(0, 2, 1, 3, 4), # => [B, C, F, H, W]
|
| 799 |
+
aug_t=aug_t,
|
| 800 |
+
act_context=act_context,
|
| 801 |
+
act_context_scale=act_context_scale,
|
| 802 |
+
clip_fea=clip_fea,
|
| 803 |
+
y=y,
|
| 804 |
+
ref_latents=ref_latents,
|
| 805 |
+
ref_mask=ref_mask,
|
| 806 |
+
).permute(0, 2, 1, 3, 4)
|
| 807 |
+
else:
|
| 808 |
+
# diffusion forcing or bidirectional
|
| 809 |
+
if classify_mode:
|
| 810 |
+
flow_pred, logits = self.model(
|
| 811 |
+
noisy_image_or_video.permute(0, 2, 1, 3, 4),
|
| 812 |
+
t=input_timestep, context=prompt_embeds,
|
| 813 |
+
seq_len=self.seq_len,
|
| 814 |
+
classify_mode=True,
|
| 815 |
+
register_tokens=self._register_tokens,
|
| 816 |
+
cls_pred_branch=self._cls_pred_branch,
|
| 817 |
+
gan_ca_blocks=self._gan_ca_blocks,
|
| 818 |
+
concat_time_embeddings=concat_time_embeddings,
|
| 819 |
+
act_context=act_context,
|
| 820 |
+
act_context_scale=act_context_scale,
|
| 821 |
+
clip_fea=clip_fea,
|
| 822 |
+
y=y,
|
| 823 |
+
ref_latents=ref_latents,
|
| 824 |
+
ref_mask=ref_mask,
|
| 825 |
+
)
|
| 826 |
+
flow_pred = flow_pred.permute(0, 2, 1, 3, 4)
|
| 827 |
+
else:
|
| 828 |
+
flow_pred = self.model(
|
| 829 |
+
noisy_image_or_video.permute(0, 2, 1, 3, 4),
|
| 830 |
+
t=input_timestep, context=prompt_embeds,
|
| 831 |
+
seq_len=self.seq_len,
|
| 832 |
+
act_context=act_context,
|
| 833 |
+
act_context_scale=act_context_scale,
|
| 834 |
+
clip_fea=clip_fea,
|
| 835 |
+
y=y,
|
| 836 |
+
ref_latents=ref_latents,
|
| 837 |
+
ref_mask=ref_mask,
|
| 838 |
+
).permute(0, 2, 1, 3, 4)
|
| 839 |
+
|
| 840 |
+
pred_x0 = self._convert_flow_pred_to_x0(
|
| 841 |
+
flow_pred=flow_pred.flatten(0, 1),
|
| 842 |
+
xt=noisy_image_or_video.flatten(0, 1),
|
| 843 |
+
timestep=timestep.flatten(0, 1)
|
| 844 |
+
).unflatten(0, flow_pred.shape[:2])
|
| 845 |
+
|
| 846 |
+
if logits is not None:
|
| 847 |
+
return flow_pred, pred_x0, logits
|
| 848 |
+
|
| 849 |
+
return flow_pred, pred_x0
|
| 850 |
+
|
| 851 |
+
def get_scheduler(self) -> SchedulerInterface:
|
| 852 |
+
"""
|
| 853 |
+
Update the current scheduler with the interface's static method
|
| 854 |
+
"""
|
| 855 |
+
scheduler = self.scheduler
|
| 856 |
+
scheduler.convert_x0_to_noise = types.MethodType(
|
| 857 |
+
SchedulerInterface.convert_x0_to_noise, scheduler)
|
| 858 |
+
scheduler.convert_noise_to_x0 = types.MethodType(
|
| 859 |
+
SchedulerInterface.convert_noise_to_x0, scheduler)
|
| 860 |
+
scheduler.convert_velocity_to_x0 = types.MethodType(
|
| 861 |
+
SchedulerInterface.convert_velocity_to_x0, scheduler)
|
| 862 |
+
self.scheduler = scheduler
|
| 863 |
+
return scheduler
|
| 864 |
+
|
| 865 |
+
def post_init(self):
|
| 866 |
+
"""
|
| 867 |
+
A few custom initialization steps that should be called after the object is created.
|
| 868 |
+
Currently, the only one we have is to bind a few methods to scheduler.
|
| 869 |
+
We can gradually add more methods here if needed.
|
| 870 |
+
"""
|
| 871 |
+
self.get_scheduler()
|
wan/__init__.py
ADDED
|
@@ -0,0 +1,2 @@
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2024-2025 The Alibaba Wan Team Authors. All rights reserved.
|
| 2 |
+
from . import modules
|
wan/modules/__init__.py
ADDED
|
@@ -0,0 +1,17 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2024-2025 The Alibaba Wan Team Authors. All rights reserved.
|
| 2 |
+
from .attention import flash_attention
|
| 3 |
+
from .model import WanModel
|
| 4 |
+
from .t5 import T5Decoder, T5Encoder, T5EncoderModel, T5Model
|
| 5 |
+
from .tokenizers import HuggingfaceTokenizer
|
| 6 |
+
from .vae2_2 import Wan2_2_VAE
|
| 7 |
+
|
| 8 |
+
__all__ = [
|
| 9 |
+
'Wan2_2_VAE',
|
| 10 |
+
'WanModel',
|
| 11 |
+
'T5Model',
|
| 12 |
+
'T5Encoder',
|
| 13 |
+
'T5Decoder',
|
| 14 |
+
'T5EncoderModel',
|
| 15 |
+
'HuggingfaceTokenizer',
|
| 16 |
+
'flash_attention',
|
| 17 |
+
]
|
wan/modules/attention.py
ADDED
|
@@ -0,0 +1,476 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
# Copyright 2024-2025 The Alibaba Wan Team Authors. All rights reserved.
|
| 2 |
+
import torch
|
| 3 |
+
|
| 4 |
+
try:
|
| 5 |
+
import flash_attn_interface
|
| 6 |
+
|
| 7 |
+
def is_hopper_gpu():
|
| 8 |
+
"""flashattn-hopper 仅支持 Hopper (H100),不支持 Blackwell,故只检测 Hopper。"""
|
| 9 |
+
if not torch.cuda.is_available():
|
| 10 |
+
return False
|
| 11 |
+
device_name = torch.cuda.get_device_name(0).lower()
|
| 12 |
+
return "h100" in device_name or "hopper" in device_name
|
| 13 |
+
FLASH_ATTN_3_AVAILABLE = is_hopper_gpu()
|
| 14 |
+
except ModuleNotFoundError:
|
| 15 |
+
FLASH_ATTN_3_AVAILABLE = False
|
| 16 |
+
|
| 17 |
+
try:
|
| 18 |
+
import flash_attn
|
| 19 |
+
FLASH_ATTN_2_AVAILABLE = True
|
| 20 |
+
except ModuleNotFoundError:
|
| 21 |
+
FLASH_ATTN_2_AVAILABLE = False
|
| 22 |
+
|
| 23 |
+
try:
|
| 24 |
+
from sageattention import sageattn
|
| 25 |
+
SAGE_ATTN_AVAILABLE = True
|
| 26 |
+
except ModuleNotFoundError:
|
| 27 |
+
SAGE_ATTN_AVAILABLE = False
|
| 28 |
+
|
| 29 |
+
try:
|
| 30 |
+
from sageattn3 import sageattn3_blackwell
|
| 31 |
+
SAGE_ATTN_3_BLACKWELL_AVAILABLE = True
|
| 32 |
+
except ModuleNotFoundError:
|
| 33 |
+
SAGE_ATTN_3_BLACKWELL_AVAILABLE = False
|
| 34 |
+
|
| 35 |
+
from .sla_attn import get_block_map
|
| 36 |
+
from .sla_kernel import _attention
|
| 37 |
+
|
| 38 |
+
|
| 39 |
+
# FLASH_ATTN_3_AVAILABLE = False
|
| 40 |
+
|
| 41 |
+
import warnings
|
| 42 |
+
|
| 43 |
+
# 优先级从高到低,模块加载时按此顺序选取第一个可用的作为默认后端
|
| 44 |
+
# ATTN_BACKEND_PRIORITY = ['flash_attn_3', 'flash_attn_2', 'sdpa']
|
| 45 |
+
ATTN_BACKEND_PRIORITY = ['sla_triton', 'sageattn', 'sageattn3', 'flash_attn_3', 'flash_attn_2', 'sdpa']
|
| 46 |
+
# ATTN_BACKEND_PRIORITY = ['flash_attn_3', 'flash_attn_2', 'sdpa']
|
| 47 |
+
print(f"Attention backend priority: {ATTN_BACKEND_PRIORITY}")
|
| 48 |
+
print(f"SAGE_ATTN_3_BLACKWELL_AVAILABLE: {SAGE_ATTN_3_BLACKWELL_AVAILABLE}")
|
| 49 |
+
print(f"SAGE_ATTN_AVAILABLE: {SAGE_ATTN_AVAILABLE}")
|
| 50 |
+
print(f"FLASH_ATTN_3_AVAILABLE: {FLASH_ATTN_3_AVAILABLE}")
|
| 51 |
+
print(f"FLASH_ATTN_2_AVAILABLE: {FLASH_ATTN_2_AVAILABLE}")
|
| 52 |
+
|
| 53 |
+
|
| 54 |
+
def _resolve_default_attn_backend():
|
| 55 |
+
"""根据优先级数组和当前环境可用性,解析出默认的 attention 后端。"""
|
| 56 |
+
availability = {
|
| 57 |
+
'sageattn3': SAGE_ATTN_3_BLACKWELL_AVAILABLE,
|
| 58 |
+
'sageattn': SAGE_ATTN_AVAILABLE,
|
| 59 |
+
'flash_attn_3': FLASH_ATTN_3_AVAILABLE,
|
| 60 |
+
'flash_attn_2': FLASH_ATTN_2_AVAILABLE,
|
| 61 |
+
'sla_triton': False,
|
| 62 |
+
'sdpa': True, # 始终可用
|
| 63 |
+
}
|
| 64 |
+
for name in ATTN_BACKEND_PRIORITY:
|
| 65 |
+
if availability.get(name, False):
|
| 66 |
+
return name
|
| 67 |
+
return 'sdpa'
|
| 68 |
+
|
| 69 |
+
|
| 70 |
+
def _is_power_of_two(x):
|
| 71 |
+
return x > 0 and (x & (x - 1)) == 0
|
| 72 |
+
|
| 73 |
+
|
| 74 |
+
def _is_backend_available(backend):
|
| 75 |
+
availability = {
|
| 76 |
+
'sageattn3': SAGE_ATTN_3_BLACKWELL_AVAILABLE,
|
| 77 |
+
'sageattn': SAGE_ATTN_AVAILABLE,
|
| 78 |
+
'flash_attn_3': FLASH_ATTN_3_AVAILABLE,
|
| 79 |
+
'flash_attn_2': FLASH_ATTN_2_AVAILABLE,
|
| 80 |
+
'sla_triton': True,
|
| 81 |
+
'sdpa': True,
|
| 82 |
+
}
|
| 83 |
+
return availability.get(backend, False)
|
| 84 |
+
|
| 85 |
+
|
| 86 |
+
def _can_use_backend(
|
| 87 |
+
backend,
|
| 88 |
+
q,
|
| 89 |
+
q_lens=None,
|
| 90 |
+
k_lens=None,
|
| 91 |
+
dropout_p=0.0,
|
| 92 |
+
causal=False,
|
| 93 |
+
window_size=(-1, -1),
|
| 94 |
+
):
|
| 95 |
+
# Layout in this module: [B, L, H, D]
|
| 96 |
+
b, _, _, d = q.shape
|
| 97 |
+
has_varlen = q_lens is not None or k_lens is not None
|
| 98 |
+
has_window = window_size != (-1, -1)
|
| 99 |
+
has_dropout = dropout_p > 0
|
| 100 |
+
|
| 101 |
+
if backend == 'sla_triton':
|
| 102 |
+
# Current SLA Triton path does not consume varlen/causal/window/dropout args,
|
| 103 |
+
# and Triton kernel in sla_attn.py requires power-of-two D.
|
| 104 |
+
return (
|
| 105 |
+
not has_varlen
|
| 106 |
+
and not has_window
|
| 107 |
+
and not causal
|
| 108 |
+
and not has_dropout
|
| 109 |
+
and b == 1
|
| 110 |
+
and _is_power_of_two(d)
|
| 111 |
+
)
|
| 112 |
+
if backend in ('sageattn', 'sageattn3'):
|
| 113 |
+
# Sage kernels in this file do not support varlen/window/dropout controls.
|
| 114 |
+
return not has_varlen and not has_window and not has_dropout
|
| 115 |
+
if backend == 'flash_attn_3':
|
| 116 |
+
# flash-attn3 path here does not support dropout/window_size.
|
| 117 |
+
return not has_window and not has_dropout
|
| 118 |
+
if backend == 'flash_attn_2':
|
| 119 |
+
return True
|
| 120 |
+
if backend == 'sdpa':
|
| 121 |
+
return True
|
| 122 |
+
return False
|
| 123 |
+
|
| 124 |
+
|
| 125 |
+
def _resolve_auto_backend(
|
| 126 |
+
q,
|
| 127 |
+
q_lens=None,
|
| 128 |
+
k_lens=None,
|
| 129 |
+
dropout_p=0.0,
|
| 130 |
+
causal=False,
|
| 131 |
+
window_size=(-1, -1),
|
| 132 |
+
):
|
| 133 |
+
candidates = [DEFAULT_ATTN_BACKEND] + [b for b in ATTN_BACKEND_PRIORITY if b != DEFAULT_ATTN_BACKEND]
|
| 134 |
+
for backend in candidates:
|
| 135 |
+
if not _is_backend_available(backend):
|
| 136 |
+
continue
|
| 137 |
+
if _can_use_backend(
|
| 138 |
+
backend=backend,
|
| 139 |
+
q=q,
|
| 140 |
+
q_lens=q_lens,
|
| 141 |
+
k_lens=k_lens,
|
| 142 |
+
dropout_p=dropout_p,
|
| 143 |
+
causal=causal,
|
| 144 |
+
window_size=window_size,
|
| 145 |
+
):
|
| 146 |
+
return backend
|
| 147 |
+
return 'sdpa'
|
| 148 |
+
|
| 149 |
+
|
| 150 |