[Refactor] Set fp4-gemm-backend=auto on SM100 and rename fp4-gemm-backend with flashinfer_ prefix (#17309)

This commit is contained in:
b8zhong
2026-01-19 04:09:07 -08:00
committed by GitHub
parent 5c02217746
commit f374623fa9
7 changed files with 42 additions and 28 deletions

View File

@@ -98,7 +98,7 @@ class CompressedTensorsW4A4Fp4(CompressedTensorsScheme):
layer.weight_global_scale.max().to(torch.float32), requires_grad=False
)
if get_fp4_gemm_runner_backend().is_trtllm():
if get_fp4_gemm_runner_backend().is_flashinfer_trtllm():
# FlashInfer TRTLLM FP4 GEMM requires a different weight layout.
# FlashInfer provides nvfp4_quantize to quantize + shuffle the
# layout but we use our own quantization so we have to call

View File

@@ -16,21 +16,35 @@ class Fp4GemmRunnerBackend(Enum):
"""Enum for FP4 GEMM runner backend selection."""
AUTO = "auto"
CUDNN = "cudnn"
CUTLASS = "cutlass"
TRTLLM = "trtllm"
FLASHINFER_CUDNN = "flashinfer_cudnn"
FLASHINFER_CUTLASS = "flashinfer_cutlass"
FLASHINFER_TRTLLM = "flashinfer_trtllm"
def is_auto(self) -> bool:
return self == Fp4GemmRunnerBackend.AUTO
def is_cudnn(self) -> bool:
return self == Fp4GemmRunnerBackend.CUDNN
def is_flashinfer_cudnn(self) -> bool:
return self == Fp4GemmRunnerBackend.FLASHINFER_CUDNN
def is_cutlass(self) -> bool:
return self == Fp4GemmRunnerBackend.CUTLASS
def is_flashinfer_cutlass(self) -> bool:
return self == Fp4GemmRunnerBackend.FLASHINFER_CUTLASS
def is_trtllm(self) -> bool:
return self == Fp4GemmRunnerBackend.TRTLLM
def is_flashinfer_trtllm(self) -> bool:
return self == Fp4GemmRunnerBackend.FLASHINFER_TRTLLM
def get_flashinfer_backend(self) -> str:
"""Get the backend string to pass to FlashInfer's mm_fp4 API.
This remaps SGLang's user-facing backend names to FlashInfer's API names.
Examples:
'flashinfer_trtllm' -> 'trtllm'
'flashinfer_cutlass' -> 'cutlass'
'flashinfer_cudnn' -> 'cudnn'
"""
if self.value.startswith("flashinfer_"):
return self.value.removeprefix("flashinfer_")
else:
return self.value
FP4_GEMM_RUNNER_BACKEND: Fp4GemmRunnerBackend | None = None

View File

@@ -128,10 +128,9 @@ def fp4_gemm(
out_features: int,
) -> torch.Tensor:
fp4_backend = get_fp4_gemm_runner_backend()
# TODO(shuw@nvidia.com): Remove the "cutlass" default override after flashinfer 0.6.0
# and let flashinfer's auto backend selection handle it.
backend = fp4_backend.value if not fp4_backend.is_auto() else "cutlass"
if enable_flashinfer_fp4_gemm:
# Use the remapping logic to convert SGLang backend names to FlashInfer API names
backend = fp4_backend.get_flashinfer_backend()
return flashinfer_fp4_gemm(
input, weight, input_sf, weight_sf, alpha, out_dtype, backend=backend
)
@@ -1155,7 +1154,7 @@ class ModelOptFp4LinearMethod(LinearMethodBase):
layer.input_scale_inv = Parameter(
(1 / input_scale_2).to(torch.float32), requires_grad=False
)
if get_fp4_gemm_runner_backend().is_trtllm():
if get_fp4_gemm_runner_backend().is_flashinfer_trtllm():
# FlashInfer TRTLLM FP4 GEMM requires a different weight layout.
# FlashInfer provides nvfp4_quantize to quantize + shuffle the
# layout but we use our own quantization so we have to call

View File

@@ -194,9 +194,9 @@ FP8_GEMM_RUNNER_BACKEND_CHOICES = [
FP4_GEMM_RUNNER_BACKEND_CHOICES = [
"auto",
"cudnn",
"cutlass",
"trtllm",
"flashinfer_cudnn",
"flashinfer_cutlass",
"flashinfer_trtllm",
]
MAMBA_SSM_DTYPE_CHOICES = ["float32", "bfloat16"]
@@ -3564,10 +3564,10 @@ class ServerArgs:
default=ServerArgs.fp4_gemm_runner_backend,
dest="fp4_gemm_runner_backend",
help="Choose the runner backend for NVFP4 GEMM operations. "
"Options: 'auto' (default, selects between cudnn/cutlass based on CUDA/cuDNN version), "
"'cudnn' (cuDNN backend, optimal on CUDA 13+ with cuDNN 9.15+), "
"'cutlass' (CUTLASS backend, optimal on CUDA 12), "
"'trtllm' (TensorRT-LLM backend, requires different weight preparation with shuffling). "
"Options: 'auto' (default, selects between flashinfer_cudnn/flashinfer_cutlass based on CUDA/cuDNN version), "
"'flashinfer_cudnn' (FlashInfer cuDNN backend, optimal on CUDA 13+ with cuDNN 9.15+), "
"'flashinfer_cutlass' (FlashInfer CUTLASS backend, optimal on CUDA 12), "
"'flashinfer_trtllm' (FlashInfer TensorRT-LLM backend, requires different weight preparation with shuffling). "
"NOTE: This replaces the deprecated environment variable "
"SGLANG_FLASHINFER_FP4_GEMM_BACKEND.",
)