[Refactor] Remove dead code for quantization (#52221)
Signed-off-by: yewentao256 <zhyanwentao@126.com>
Wentao Ye committed
Aug 16, 2026 at 13:10 UTC
7d7b6f26f4120b5db10d0a697a878ad605f09579
14 files changed
+1
-395
tests/kernels/quantization/test_block_int8.py
deleted
-75
@@ -1,75 +0,0 @@
1
-# SPDX-License-Identifier: Apache-2.0
2
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
3
-
4
-# Adapted from https://github.com/sgl-project/sglang/blob/main/test/srt/test_block_int8.py
5
-import itertools
6
-
7
-import pytest
8
-import torch
9
-
10
-from tests.kernels.quant_utils import native_w8a8_block_matmul
11
-from vllm.config import VllmConfig
12
-from vllm.model_executor.layers.quantization.utils.int8_utils import (
13
- w8a8_block_int8_matmul,
14
-)
15
-from vllm.platforms import current_platform
16
-
17
-if not (current_platform.is_cuda_alike() or current_platform.is_xpu()):
18
- pytest.skip(
19
- "INT8 Triton kernels require a CUDA-alike or XPU device",
20
- allow_module_level=True,
21
- )
22
-
23
-if current_platform.is_cuda_alike() and not current_platform.has_device_capability(
24
- (7, 0)
25
-):
26
- pytest.skip("INT8 Triton requires CUDA 7.0 or higher", allow_module_level=True)
27
-
28
-vllm_config = VllmConfig()
29
-
30
-DTYPES = [torch.half, torch.bfloat16]
31
-M = [1, 33, 64, 222]
32
-N = [128, 1024]
33
-K = [256, 4096]
34
-# BLOCK_SIZE = [[64, 64], [64, 128], [128, 64], [128, 128]]
35
-BLOCK_SIZE = [[128, 128]]
36
-SEEDS = [0]
37
-
38
-
39
-@pytest.mark.parametrize(
40
- "M,N,K,block_size,out_dtype,seed",
41
- itertools.product(M, N, K, BLOCK_SIZE, DTYPES, SEEDS),
42
-)
43
-@torch.inference_mode()
44
-def test_w8a8_block_int8_matmul(M, N, K, block_size, out_dtype, seed):
45
- torch.manual_seed(seed)
46
- device = current_platform.device_type
47
- factor_for_scale = 1e-2
48
- int8_info = torch.iinfo(torch.int8)
49
- int8_max, int8_min = int8_info.max, int8_info.min
50
-
51
- A_fp32 = torch.rand(M, K, dtype=torch.float32, device=device)
52
- A_fp32 = (A_fp32 - 0.5) * 2 * int8_max
53
- A_fp8 = A_fp32.clamp(min=int8_min, max=int8_max).to(torch.float8_e4m3fn)
54
-
55
- B_fp32 = torch.rand(N, K, dtype=torch.float32, device=device)
56
- B_fp32 = (B_fp32 - 0.5) * 2 * int8_max
57
- B_fp8 = B_fp32.clamp(min=int8_min, max=int8_max).to(torch.float8_e4m3fn)
58
-
59
- block_n, block_k = block_size[0], block_size[1]
60
- n_tiles = (N + block_n - 1) // block_n
61
- k_tiles = (K + block_k - 1) // block_k
62
-
63
- As = torch.rand(M, k_tiles, dtype=torch.float32, device=device) * factor_for_scale
64
- Bs = (
65
- torch.rand(n_tiles, k_tiles, dtype=torch.float32, device=device)
66
- * factor_for_scale
67
- )
68
-
69
- ref_out = native_w8a8_block_matmul(A_fp8, B_fp8, As, Bs, block_size, out_dtype)
70
- out = w8a8_block_int8_matmul(A_fp8, B_fp8, As, Bs, block_size, out_dtype)
71
-
72
- rel_diff = torch.mean(
73
- torch.abs(out.to(torch.float32) - ref_out.to(torch.float32))
74
- ) / torch.mean(torch.abs(ref_out.to(torch.float32)))
75
- assert rel_diff < 0.001
vllm/model_executor/kernels/linear/base.py
+1
-23
@@ -89,28 +89,6 @@ class FP8Params(Params):
89
)
90
91
92
-@dataclass
93
-class Int8Params(Params):
94
- """Int8 layer parameters with typed fields"""
95
-
96
- input_zero_point: torch.Tensor | None
97
- azp_adj: torch.Tensor | None
98
-
99
- INPUT_ZERO_POINT: ClassVar[str] = "input_zero_point"
100
- AZP_ADJ: ClassVar[str] = "azp_adj"
101
-
102
- @classmethod
103
- def from_layer(cls, layer: torch.nn.Module) -> "Int8Params":
104
- """Extract parameters from layer"""
105
- return cls(
106
- weight=getattr(layer, cls.WEIGHT),
107
- weight_scale=getattr(layer, cls.WEIGHT_SCALE),
108
- input_scale=getattr(layer, cls.INPUT_SCALE, None),
109
- input_zero_point=getattr(layer, cls.INPUT_ZERO_POINT, None),
110
- azp_adj=getattr(layer, cls.AZP_ADJ, None),
111
- )
112
-
113
-
92
_ParamsT = TypeVar("_ParamsT", bound=Params)
93
_ConfigT = TypeVar("_ConfigT", bound=MMLinearLayerConfig)
94
@@ -130,7 +108,7 @@ class MMLinearKernel(ABC, Generic[_ConfigT, _ParamsT]):
108
109
Typical Usage:
110
1. Define a config dataclass inheriting from MMLinearLayerConfig
133
- 2. Define a params dataclass inheriting from Params (or FP8Params/Int8Params)
111
+ 2. Define a params dataclass inheriting from Params
112
3. Subclass MMLinearKernel with your config and params types
113
4. Implement all abstract methods
114
5. Register the kernel with the quantization method
vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
-8
@@ -130,14 +130,6 @@ class Mxfp4MoeBackend(Enum):
130
HUMMING = "HUMMING"
131
132
133
-# AITER backends group
134
-AITER_BACKENDS = (
135
- Mxfp4MoeBackend.AITER_MXFP4_BF16,
136
- Mxfp4MoeBackend.AITER_MXFP4_FP8,
137
- Mxfp4MoeBackend.AITER_MXFP4_MXFP4,
138
-)
139
-
140
-
133
# Backends that share the same TRTLLM weight format
134
TRTLLM_BACKENDS = (
135
Mxfp4MoeBackend.FLASHINFER_TRTLLM_MXFP4_BF16,
vllm/model_executor/layers/quantization/auto_awq.py
-25
@@ -358,31 +358,6 @@ class AutoAWQConfig(QuantizationConfig):
358
359
return None
360
361
- @classmethod
362
- def is_awq_marlin_compatible(cls, quant_config: dict[str, Any]):
363
- # Extract data from quant config.
364
- quant_method = quant_config.get("quant_method", "").lower()
365
- num_bits = quant_config.get("bits")
366
- group_size = quant_config.get("group_size")
367
- zero_point = quant_config.get("zero_point")
368
-
369
- if not (current_platform.is_cuda_alike() or current_platform.is_cpu()):
370
- return False
371
-
372
- if quant_method != "awq":
373
- return False
374
-
375
- # If we cannot find the info needed in the config, cannot convert.
376
- if num_bits is None or group_size is None or zero_point is None:
377
- return False
378
-
379
- if num_bits not in cls.TYPE_MAP:
380
- return False
381
-
382
- return check_marlin_supported(
383
- quant_type=cls.TYPE_MAP[num_bits], group_size=group_size, has_zp=zero_point
384
- )
385
-
361
def apply_vllm_mapper(self, hf_to_vllm_mapper: "WeightsMapper"):
362
if self.modules_to_not_convert:
363
self.modules_to_not_convert = hf_to_vllm_mapper.apply_list(
vllm/model_executor/layers/quantization/compressed_tensors/schemes/compressed_tensors_w4a8_fp8.py
-1
@@ -32,7 +32,6 @@ __all__ = ["CompressedTensorsW4A8Fp8"]
32
W4A8_SUPPORTED_TYPES_MAP = {
33
4: scalar_types.int4,
34
}
35
-W4A8_SUPPORTED_BITS = list(W4A8_SUPPORTED_TYPES_MAP.keys())
35
36
37
class CompressedTensorsW4A8Fp8(CompressedTensorsScheme):
vllm/model_executor/layers/quantization/compressed_tensors/schemes/compressed_tensors_w4a8_int.py
-1
@@ -27,7 +27,6 @@ __all__ = ["CompressedTensorsW4A8Int"]
27
W4A8_SUPPORTED_TYPES_MAP = {
28
4: scalar_types.int4,
29
}
30
-W4A8_SUPPORTED_BITS = list(W4A8_SUPPORTED_TYPES_MAP.keys())
30
31
32
class CompressedTensorsW4A8Int(CompressedTensorsScheme):
vllm/model_executor/layers/quantization/humming.py
-22
@@ -1,6 +1,5 @@
1
# SPDX-License-Identifier: Apache-2.0
2
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
3
-import math
3
from collections.abc import Callable
4
from typing import TYPE_CHECKING, Any
5
@@ -61,11 +60,6 @@ if TYPE_CHECKING:
60
)
61
62
64
-def prepare_padded_shape(shape, x):
65
- padded_shape = math.ceil(shape / x) * x
66
- return padded_shape, padded_shape - shape
67
-
68
-
63
def prepare_param(tensor, name, extra_attrs):
64
extra_attrs = extra_attrs.copy()
65
scale_type = extra_attrs.pop("scale_type", None)
@@ -129,22 +123,6 @@ def prepare_moe_param(tensor: torch.Tensor, name: str, extra_attrs: dict[str, An
123
return param
124
125
132
-def may_pad_loaded_weight(param, loaded_weight):
133
- pad_shape = getattr(param, "pad_shape", None)
134
- if pad_shape is None:
135
- return loaded_weight
136
- value = 1 if loaded_weight.dtype == torch.float8_e8m0fnu else 0
137
- padding = []
138
- for x in pad_shape[::-1][: loaded_weight.ndim]:
139
- padding += [0, x]
140
- loaded_weight = torch.nn.functional.pad(
141
- input=loaded_weight,
142
- pad=padding,
143
- value=value,
144
- )
145
- return loaded_weight
146
-
147
-
126
def compressed_tensors_get_config(config: dict[str, Any], key: str):
127
assert key in ["weights", "input_activations"]
128
target_group_config = None
vllm/model_executor/layers/quantization/inc/inc.py
-1
@@ -53,7 +53,6 @@ class INCConfig(QuantizationConfig):
53
MXFP8_GROUP_SIZE = 32
54
MXFP8_DATA_TYPE = "mx_fp"
55
MXFP8_PACKING_FORMAT = "auto_round:llm_compressor"
56
- MXFP8_SUPPORTED_ACT_DTYPES = {"mx_fp", "mx_fp_rceil"}
56
57
def __init__(
58
self,
vllm/model_executor/layers/quantization/inc/schemes/inc_wna16_linear.py
-4
@@ -458,7 +458,3 @@ class INCARKLinearMethod(INCXPULinearBase):
458
layer.ark_scale_type,
459
not self.sym,
460
)
461
-
462
-
463
-class INCXPUW4A16LinearScheme(INCXPULinearMethod):
464
- pass
vllm/model_executor/layers/quantization/modelopt.py
-1
@@ -112,7 +112,6 @@ QUANT_ALGOS = [
112
# MIXED_PRECISION,
113
"MIXED_PRECISION",
114
]
115
-KV_CACHE_QUANT_ALGOS = ["FP8", "NVFP4"]
115
116
117
class ModelOptKVCacheMethod(BaseKVCacheMethod):
vllm/model_executor/layers/quantization/torchao.py
-3
@@ -283,9 +283,6 @@ class TorchAOConfig(QuantizationConfig):
283
284
return TorchAOLinearMethod(self)
285
286
- def get_scaled_act_names(self) -> list[str]:
287
- return []
288
-
286
287
def torchao_quantize_param_data(
288
param: torch.Tensor, torchao_config: Any
vllm/model_executor/layers/quantization/utils/int8_utils.py
-223
@@ -2,17 +2,12 @@
2
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
3
4
# Adapted from https://github.com/sgl-project/sglang/blob/4cb53ecd0cffceb6dee5c011a58f65997a86f151/python/sglang/srt/layers/quantization/int8_kernel.py
5
-import functools
6
-import json
5
import logging
8
-import os
9
-from typing import Any
6
7
import torch
8
9
from vllm.platforms import current_platform
10
from vllm.triton_utils import tl, triton
15
-from vllm.utils.platform_utils import get_device_name_as_file_name
11
12
logger = logging.getLogger(__name__)
13
@@ -228,221 +223,3 @@ def per_token_group_quant_int8(
223
)
224
225
return x_q, x_s
231
-
232
-
233
-@triton.jit
234
-def _w8a8_block_int8_matmul(
235
- # Pointers to inputs and output
236
- A,
237
- B,
238
- C,
239
- As,
240
- Bs,
241
- # Shape for matmul
242
- M,
243
- N,
244
- K,
245
- # Block size for block-wise quantization
246
- group_n,
247
- group_k,
248
- # Stride for inputs and output
249
- stride_am,
250
- stride_ak,
251
- stride_bk,
252
- stride_bn,
253
- stride_cm,
254
- stride_cn,
255
- stride_As_m,
256
- stride_As_k,
257
- stride_Bs_k,
258
- stride_Bs_n,
259
- # Meta-parameters
260
- BLOCK_SIZE_M: tl.constexpr,
261
- BLOCK_SIZE_N: tl.constexpr,
262
- BLOCK_SIZE_K: tl.constexpr,
263
- GROUP_SIZE_M: tl.constexpr,
264
-):
265
- """Triton-accelerated function used to perform linear operations (dot
266
- product) on input tensors `A` and `B` with block-wise quantization, and
267
- store the result in output tensor `C`.
268
- """
269
-
270
- pid = tl.program_id(axis=0)
271
- num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)
272
- num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)
273
- num_pid_in_group = GROUP_SIZE_M * num_pid_n
274
- group_id = pid // num_pid_in_group
275
- first_pid_m = group_id * GROUP_SIZE_M
276
- group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)
277
- pid_m = first_pid_m + (pid % group_size_m)
278
- pid_n = (pid % num_pid_in_group) // group_size_m
279
-
280
- offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M
281
- offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
282
- offs_k = tl.arange(0, BLOCK_SIZE_K)
283
- a_ptrs = A + (offs_am[:, None] * stride_am + offs_k[None, :] * stride_ak)
284
- b_ptrs = B + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)
285
-
286
- As_ptrs = As + offs_am * stride_As_m
287
- offs_bsn = offs_bn // group_n
288
- Bs_ptrs = Bs + offs_bsn * stride_Bs_n
289
-
290
- accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
291
- for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):
292
- a = tl.load(a_ptrs, mask=offs_k[None, :] < K - k * BLOCK_SIZE_K, other=0.0)
293
- b = tl.load(b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_SIZE_K, other=0.0)
294
-
295
- k_start = k * BLOCK_SIZE_K
296
- offs_ks = k_start // group_k
297
- a_s = tl.load(As_ptrs + offs_ks * stride_As_k)
298
- b_s = tl.load(Bs_ptrs + offs_ks * stride_Bs_k)
299
-
300
- accumulator += tl.dot(a, b).to(tl.float32) * a_s[:, None] * b_s[None, :]
301
- a_ptrs += BLOCK_SIZE_K * stride_ak
302
- b_ptrs += BLOCK_SIZE_K * stride_bk
303
-
304
- if C.dtype.element_ty == tl.bfloat16:
305
- c = accumulator.to(tl.bfloat16)
306
- elif C.dtype.element_ty == tl.float16:
307
- c = accumulator.to(tl.float16)
308
- else:
309
- c = accumulator.to(tl.float32)
310
-
311
- offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
312
- offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
313
- c_ptrs = C + stride_cm * offs_cm[:, None] + stride_cn * offs_cn[None, :]
314
- c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)
315
- tl.store(c_ptrs, c, mask=c_mask)
316
-
317
-
318
-@functools.lru_cache
319
-def get_w8a8_block_int8_configs(
320
- N: int, K: int, block_n: int, block_k: int
321
-) -> dict[int, Any] | None:
322
- """
323
- Return optimized configurations for the w8a8 block fp8 kernel.
324
-
325
- The return value will be a dictionary that maps an irregular grid of
326
- batch sizes to configurations of the w8a8 block fp8 kernel. To evaluate the
327
- kernel on a given batch size bs, the closest batch size in the grid should
328
- be picked and the associated configuration chosen to invoke the kernel.
329
- """
330
-
331
- # First look up if an optimized configuration is available in the configs
332
- # directory
333
- device_name = get_device_name_as_file_name()
334
- json_file_name = f"N={N},K={K},device_name={device_name},dtype=int8_w8a8,block_shape=[{block_n}, {block_k}].json" # noqa: E501
335
-
336
- config_file_path = os.path.join(
337
- os.path.dirname(os.path.realpath(__file__)), "configs", json_file_name
338
- )
339
- if os.path.exists(config_file_path):
340
- with open(config_file_path) as f:
341
- logger.info(
342
- "Using configuration from %s for W8A8 Block INT8 kernel.",
343
- config_file_path,
344
- )
345
- # If a configuration has been found, return it
346
- return {int(key): val for key, val in json.load(f).items()}
347
-
348
- # If no optimized configuration is available, we will use the default
349
- # configuration
350
- logger.warning(
351
- (
352
- "Using default W8A8 Block INT8 kernel config. Performance might "
353
- "be sub-optimal! Config file not found at %s"
354
- ),
355
- config_file_path,
356
- )
357
- return None
358
-
359
-
360
-def w8a8_block_int8_matmul(
361
- A: torch.Tensor,
362
- B: torch.Tensor,
363
- As: torch.Tensor,
364
- Bs: torch.Tensor,
365
- block_size: list[int],
366
- output_dtype: torch.dtype = torch.float16,
367
-) -> torch.Tensor:
368
- """This function performs matrix multiplication with block-wise
369
- quantization.
370
-
371
- It takes two input tensors `A` and `B` with scales `As` and `Bs`.
372
- The output is returned in the specified `output_dtype`.
373
-
374
- Args:
375
- A: The input tensor, e.g., activation.
376
- B: The input tensor, e.g., weight.
377
- As: The per-token-group quantization scale for `A`.
378
- Bs: The per-block quantization scale for `B`.
379
- block_size: The block size for per-block quantization. It should be
380
- 2-dim, e.g., [128, 128].
381
- output_dtype: The dtype of the returned tensor.
382
-
383
- Returns:
384
- torch.Tensor: The result of matmul.
385
- """
386
- assert len(block_size) == 2
387
- block_n, block_k = block_size[0], block_size[1]
388
-
389
- assert A.shape[-1] == B.shape[-1]
390
- assert A.shape[:-1] == As.shape[:-1] and A.is_contiguous()
391
- assert triton.cdiv(A.shape[-1], block_k) == As.shape[-1]
392
- M = A.numel() // A.shape[-1]
393
-
394
- assert B.ndim == 2 and B.is_contiguous() and Bs.ndim == 2
395
- N, K = B.shape
396
- assert triton.cdiv(N, block_n) == Bs.shape[0]
397
- assert triton.cdiv(K, block_k) == Bs.shape[1]
398
-
399
- C_shape = A.shape[:-1] + (N,)
400
- C = A.new_empty(C_shape, dtype=output_dtype)
401
-
402
- configs = get_w8a8_block_int8_configs(N, K, block_size[0], block_size[1])
403
- if configs:
404
- # If an optimal configuration map has been found, look up the
405
- # optimal config
406
- config = configs[min(configs.keys(), key=lambda x: abs(x - M))]
407
- else:
408
- # Default config
409
- # Block-wise quant: BLOCK_SIZE_K must be divisible by block_size[1]
410
- config = {
411
- "BLOCK_SIZE_M": 64,
412
- "BLOCK_SIZE_N": block_size[0],
413
- "BLOCK_SIZE_K": block_size[1],
414
- "GROUP_SIZE_M": 32,
415
- "num_warps": 4,
416
- "num_stages": 3,
417
- }
418
-
419
- def grid(META):
420
- return (
421
- triton.cdiv(M, META["BLOCK_SIZE_M"]) * triton.cdiv(N, META["BLOCK_SIZE_N"]),
422
- )
423
-
424
- _w8a8_block_int8_matmul[grid](
425
- A,
426
- B,
427
- C,
428
- As,
429
- Bs,
430
- M,
431
- N,
432
- K,
433
- block_n,
434
- block_k,
435
- A.stride(-2),
436
- A.stride(-1),
437
- B.stride(1),
438
- B.stride(0),
439
- C.stride(-2),
440
- C.stride(-1),
441
- As.stride(-2),
442
- As.stride(-1),
443
- Bs.stride(1),
444
- Bs.stride(0),
445
- **config,
446
- )
447
-
448
- return C
vllm/model_executor/layers/quantization/utils/mxfp4_utils.py
-7
@@ -12,13 +12,6 @@ from vllm.utils.torch_utils import direct_register_custom_op, is_torch_equal_or_
12
13
logger = init_logger(__name__)
14
15
-# CK's pre-compiled MXFP4 MoE GEMM kernel instances require the
16
-# intermediate_size (after TP split) to be a multiple of this value.
17
-# This arises from FP4 packing (2 values per byte) combined with CK
18
-# tile size constraints. When violated, AITER raises:
19
-# "device_gemm ... does not support this GEMM problem".
20
-CK_MXFP4_MOE_DIM_ALIGNMENT = 256
21
-
15
16
def should_use_cdna4_mx_scale_swizzle() -> bool:
17
"""Whether to use the CDNA4 swizzled scale layout for mxfp4 on gfx950.
vllm/model_executor/layers/quantization/utils/ocp_mx_utils.py
-1
@@ -16,7 +16,6 @@ OCP_MX_DTYPES = {
16
"mxfp8_e5m2",
17
"mxint8",
18
}
19
-SUPPORTED_OCP_MX_DTYPES = {"mxfp4", "mxfp6_e3m2", "mxfp6_e2m3"}
19
20
21
class OCP_MX_Scheme(str, Enum):