[Kernel] Set sgl_per_token_group_quant_8bit_v2 as default choice (#22467)
This commit is contained in:
@@ -122,7 +122,6 @@ SGLang supports various environment variables that can be used to configure its
|
|||||||
| Environment Variable | Description | Default Value |
|
| Environment Variable | Description | Default Value |
|
||||||
| --- | --- | --- |
|
| --- | --- | --- |
|
||||||
| `SGLANG_INT4_WEIGHT` | Enable INT4 weight quantization | `false` |
|
| `SGLANG_INT4_WEIGHT` | Enable INT4 weight quantization | `false` |
|
||||||
| `SGLANG_PER_TOKEN_GROUP_QUANT_8BIT_V2` | Apply per token group quantization kernel with fused silu and mul and masked m | `false` |
|
|
||||||
| `SGLANG_FORCE_FP8_MARLIN` | Force using FP8 MARLIN kernels even if other FP8 kernels are available | `false` |
|
| `SGLANG_FORCE_FP8_MARLIN` | Force using FP8 MARLIN kernels even if other FP8 kernels are available | `false` |
|
||||||
| `SGLANG_NVFP4_CKPT_FP8_GEMM_IN_ATTN` | Quantize q_b_proj from BF16 to FP8 when launching DeepSeek NVFP4 checkpoint | `false` |
|
| `SGLANG_NVFP4_CKPT_FP8_GEMM_IN_ATTN` | Quantize q_b_proj from BF16 to FP8 when launching DeepSeek NVFP4 checkpoint | `false` |
|
||||||
| `SGLANG_MOE_NVFP4_DISPATCH` | Use nvfp4 for moe dispatch (on flashinfer_cutlass or flashinfer_cutedsl moe runner backend) | `"false"` |
|
| `SGLANG_MOE_NVFP4_DISPATCH` | Use nvfp4 for moe dispatch (on flashinfer_cutlass or flashinfer_cutedsl moe runner backend) | `"false"` |
|
||||||
|
|||||||
@@ -3,7 +3,7 @@ import subprocess
|
|||||||
import warnings
|
import warnings
|
||||||
from contextlib import ExitStack, contextmanager
|
from contextlib import ExitStack, contextmanager
|
||||||
from enum import IntEnum
|
from enum import IntEnum
|
||||||
from typing import Any
|
from typing import Any, Optional
|
||||||
|
|
||||||
|
|
||||||
@contextmanager
|
@contextmanager
|
||||||
@@ -341,7 +341,6 @@ class Envs:
|
|||||||
SGLANG_FORCE_FP8_MARLIN = EnvBool(False)
|
SGLANG_FORCE_FP8_MARLIN = EnvBool(False)
|
||||||
SGLANG_MOE_NVFP4_DISPATCH = EnvBool(False)
|
SGLANG_MOE_NVFP4_DISPATCH = EnvBool(False)
|
||||||
SGLANG_NVFP4_CKPT_FP8_GEMM_IN_ATTN = EnvBool(False)
|
SGLANG_NVFP4_CKPT_FP8_GEMM_IN_ATTN = EnvBool(False)
|
||||||
SGLANG_PER_TOKEN_GROUP_QUANT_8BIT_V2 = EnvBool(False)
|
|
||||||
SGLANG_NVFP4_CKPT_FP8_NEXTN_MOE = EnvBool(False)
|
SGLANG_NVFP4_CKPT_FP8_NEXTN_MOE = EnvBool(False)
|
||||||
SGLANG_QUANT_ALLOW_DOWNCASTING = EnvBool(False)
|
SGLANG_QUANT_ALLOW_DOWNCASTING = EnvBool(False)
|
||||||
SGLANG_FP8_IGNORED_LAYERS = EnvStr("")
|
SGLANG_FP8_IGNORED_LAYERS = EnvStr("")
|
||||||
@@ -552,8 +551,11 @@ envs = Envs()
|
|||||||
EnvField._allow_set_name = False
|
EnvField._allow_set_name = False
|
||||||
|
|
||||||
|
|
||||||
def _print_deprecated_env(new_name: str, old_name: str):
|
def _print_deprecated_env(old_name: str, new_name: Optional[str] = None):
|
||||||
if old_name in os.environ:
|
if old_name in os.environ:
|
||||||
|
if new_name is None:
|
||||||
|
warnings.warn(f"Environment variable {old_name} has been deprecated.")
|
||||||
|
else:
|
||||||
warnings.warn(
|
warnings.warn(
|
||||||
f"Environment variable {old_name} will be deprecated, please use {new_name} instead"
|
f"Environment variable {old_name} will be deprecated, please use {new_name} instead"
|
||||||
)
|
)
|
||||||
@@ -570,14 +572,15 @@ def _warn_deprecated_env_to_cli_flag(env_name: str, suggestion: str):
|
|||||||
|
|
||||||
|
|
||||||
def _convert_SGL_to_SGLANG():
|
def _convert_SGL_to_SGLANG():
|
||||||
_print_deprecated_env("SGLANG_LOG_GC", "SGLANG_GC_LOG")
|
_print_deprecated_env("SGLANG_GC_LOG", "SGLANG_LOG_GC")
|
||||||
_print_deprecated_env(
|
_print_deprecated_env(
|
||||||
"SGLANG_MOE_NVFP4_DISPATCH", "SGLANG_CUTEDSL_MOE_NVFP4_DISPATCH"
|
"SGLANG_CUTEDSL_MOE_NVFP4_DISPATCH", "SGLANG_MOE_NVFP4_DISPATCH"
|
||||||
)
|
)
|
||||||
_print_deprecated_env(
|
_print_deprecated_env(
|
||||||
"SGLANG_ENABLE_TP_MEMORY_INBALANCE_CHECK",
|
|
||||||
"SGL_DISABLE_TP_MEMORY_INBALANCE_CHECK",
|
"SGL_DISABLE_TP_MEMORY_INBALANCE_CHECK",
|
||||||
|
"SGLANG_ENABLE_TP_MEMORY_INBALANCE_CHECK",
|
||||||
)
|
)
|
||||||
|
_print_deprecated_env("SGLANG_PER_TOKEN_GROUP_QUANT_8BIT_V2")
|
||||||
_deprecated_ms_to_s = {
|
_deprecated_ms_to_s = {
|
||||||
"SGLANG_QUEUED_TIMEOUT_MS": "SGLANG_REQ_WAITING_TIMEOUT",
|
"SGLANG_QUEUED_TIMEOUT_MS": "SGLANG_REQ_WAITING_TIMEOUT",
|
||||||
"SGLANG_FORWARD_TIMEOUT_MS": "SGLANG_REQ_RUNNING_TIMEOUT",
|
"SGLANG_FORWARD_TIMEOUT_MS": "SGLANG_REQ_RUNNING_TIMEOUT",
|
||||||
|
|||||||
@@ -324,10 +324,6 @@ def _per_token_group_quant_8bit_raw(
|
|||||||
return x_q, x_s
|
return x_q, x_s
|
||||||
|
|
||||||
|
|
||||||
# backward compatibility
|
|
||||||
per_token_group_quant_fp8 = _per_token_group_quant_8bit_raw
|
|
||||||
|
|
||||||
|
|
||||||
def _per_token_group_quant_8bit_fuse_silu_and_mul(
|
def _per_token_group_quant_8bit_fuse_silu_and_mul(
|
||||||
x: torch.Tensor,
|
x: torch.Tensor,
|
||||||
group_size: int,
|
group_size: int,
|
||||||
@@ -507,6 +503,11 @@ def sglang_per_token_group_quant_fp8(
|
|||||||
scale_ue8m0=scale_ue8m0,
|
scale_ue8m0=scale_ue8m0,
|
||||||
)
|
)
|
||||||
|
|
||||||
|
# Enable v2 kernel by default on supported group sizes
|
||||||
|
_V2_KERNEL_SUPPORTED_GROUP_SIZES = [16, 32, 64, 128]
|
||||||
|
if enable_v2 is None:
|
||||||
|
enable_v2 = group_size in _V2_KERNEL_SUPPORTED_GROUP_SIZES
|
||||||
|
|
||||||
if x.shape[0] > 0:
|
if x.shape[0] > 0:
|
||||||
# Temporary
|
# Temporary
|
||||||
if enable_sgl_per_token_group_quant_8bit:
|
if enable_sgl_per_token_group_quant_8bit:
|
||||||
@@ -606,6 +607,12 @@ def sglang_per_token_quant_fp8(
|
|||||||
return x_q, x_s
|
return x_q, x_s
|
||||||
|
|
||||||
|
|
||||||
|
if _is_cuda:
|
||||||
|
per_token_group_quant_fp8 = sglang_per_token_group_quant_fp8
|
||||||
|
else:
|
||||||
|
per_token_group_quant_fp8 = _per_token_group_quant_8bit_raw
|
||||||
|
|
||||||
|
|
||||||
@triton.jit
|
@triton.jit
|
||||||
def _static_quant_fp8(
|
def _static_quant_fp8(
|
||||||
# Pointers to inputs and output
|
# Pointers to inputs and output
|
||||||
|
|||||||
@@ -109,10 +109,9 @@ def sgl_per_token_group_quant_8bit(
|
|||||||
masked_m: Optional[torch.Tensor] = None,
|
masked_m: Optional[torch.Tensor] = None,
|
||||||
enable_v2: Optional[bool] = None,
|
enable_v2: Optional[bool] = None,
|
||||||
) -> None:
|
) -> None:
|
||||||
|
_V2_KERNEL_SUPPORTED_GROUP_SIZES = [16, 32, 64, 128]
|
||||||
if enable_v2 is None:
|
if enable_v2 is None:
|
||||||
from sglang.srt.utils import get_bool_env_var
|
enable_v2 = group_size in _V2_KERNEL_SUPPORTED_GROUP_SIZES
|
||||||
|
|
||||||
enable_v2 = get_bool_env_var("SGLANG_PER_TOKEN_GROUP_QUANT_8BIT_V2")
|
|
||||||
|
|
||||||
if enable_v2:
|
if enable_v2:
|
||||||
return torch.ops.sgl_kernel.sgl_per_token_group_quant_8bit_v2.default(
|
return torch.ops.sgl_kernel.sgl_per_token_group_quant_8bit_v2.default(
|
||||||
|
|||||||
@@ -96,7 +96,9 @@ class TestPerTokenGroupQuantFP8(TestFP8Base):
|
|||||||
A, A_quant_gt, scale_gt = self._make_A(
|
A, A_quant_gt, scale_gt = self._make_A(
|
||||||
M=self.M, K=self.K, group_size=self.group_size, out_dtype=self.quant_type
|
M=self.M, K=self.K, group_size=self.group_size, out_dtype=self.quant_type
|
||||||
)
|
)
|
||||||
A_quant, scale = per_token_group_quant_fp8(x=A, group_size=self.group_size)
|
A_quant, scale = per_token_group_quant_fp8(
|
||||||
|
x=A.to(torch.bfloat16), group_size=self.group_size
|
||||||
|
)
|
||||||
torch.testing.assert_close(scale, scale_gt)
|
torch.testing.assert_close(scale, scale_gt)
|
||||||
diff = (A_quant.to(torch.float16) - A_quant_gt.to(torch.float16)).abs()
|
diff = (A_quant.to(torch.float16) - A_quant_gt.to(torch.float16)).abs()
|
||||||
diff_count = (diff > 1e-5).count_nonzero()
|
diff_count = (diff > 1e-5).count_nonzero()
|
||||||
|
|||||||
Reference in New Issue
Block a user