[Perf] Enable Flashinfer autotune by default (#14357)
This commit is contained in:
@@ -246,6 +246,7 @@ Please consult the documentation below and [server_args.py](https://github.com/s
|
|||||||
| `--nsa-prefill` | Choose the NSA backend for the prefill stage (overrides `--attention-backend` when running DeepSeek NSA-style attention). | `flashmla_sparse` | `flashmla_sparse`, `flashmla_decode`, `fa3`, `tilelang`, `aiter` |
|
| `--nsa-prefill` | Choose the NSA backend for the prefill stage (overrides `--attention-backend` when running DeepSeek NSA-style attention). | `flashmla_sparse` | `flashmla_sparse`, `flashmla_decode`, `fa3`, `tilelang`, `aiter` |
|
||||||
| `--nsa-decode` | Choose the NSA backend for the decode stage when running DeepSeek NSA-style attention. Overrides `--attention-backend` for decoding. | `flashmla_kv` | `flashmla_prefill`, `flashmla_kv`, `fa3`, `tilelang`, `aiter` |
|
| `--nsa-decode` | Choose the NSA backend for the decode stage when running DeepSeek NSA-style attention. Overrides `--attention-backend` for decoding. | `flashmla_kv` | `flashmla_prefill`, `flashmla_kv`, `fa3`, `tilelang`, `aiter` |
|
||||||
| `--fp8-gemm-backend` | Choose the runner backend for Blockwise FP8 GEMM operations. Options: 'auto' (default, auto-selects based on hardware), 'deep_gemm' (JIT-compiled; enabled by default on NVIDIA Hopper (SM90) and Blackwell (SM100) when DeepGEMM is installed), 'flashinfer_trtllm' (optimal for Blackwell and low-latency), 'cutlass' (optimal for Hopper/Blackwell GPUs and high-throughput), 'triton' (fallback, widely compatible), 'aiter' (ROCm only). **NOTE**: This replaces the deprecated environment variables SGLANG_ENABLE_FLASHINFER_FP8_GEMM and SGLANG_SUPPORT_CUTLASS_BLOCK_FP8. | `auto` | `auto`, `deep_gemm`, `flashinfer_trtllm`, `cutlass`, `triton`, `aiter` |
|
| `--fp8-gemm-backend` | Choose the runner backend for Blockwise FP8 GEMM operations. Options: 'auto' (default, auto-selects based on hardware), 'deep_gemm' (JIT-compiled; enabled by default on NVIDIA Hopper (SM90) and Blackwell (SM100) when DeepGEMM is installed), 'flashinfer_trtllm' (optimal for Blackwell and low-latency), 'cutlass' (optimal for Hopper/Blackwell GPUs and high-throughput), 'triton' (fallback, widely compatible), 'aiter' (ROCm only). **NOTE**: This replaces the deprecated environment variables SGLANG_ENABLE_FLASHINFER_FP8_GEMM and SGLANG_SUPPORT_CUTLASS_BLOCK_FP8. | `auto` | `auto`, `deep_gemm`, `flashinfer_trtllm`, `cutlass`, `triton`, `aiter` |
|
||||||
|
| `--disable-flashinfer-autotune` | Flashinfer autotune is enabled by default. Set this flag to disable the autotune. | `False` | bool flag (set to enable) |
|
||||||
|
|
||||||
## Speculative decoding
|
## Speculative decoding
|
||||||
| Argument | Description | Defaults | Options |
|
| Argument | Description | Defaults | Options |
|
||||||
|
|||||||
@@ -54,6 +54,7 @@ from sglang.srt.utils import (
|
|||||||
is_cpu,
|
is_cpu,
|
||||||
is_flashinfer_available,
|
is_flashinfer_available,
|
||||||
is_hip,
|
is_hip,
|
||||||
|
next_power_of_2,
|
||||||
round_up,
|
round_up,
|
||||||
)
|
)
|
||||||
|
|
||||||
@@ -1087,6 +1088,7 @@ class FlashInferFusedMoE(FusedMoE):
|
|||||||
local_expert_offset=self.moe_ep_rank * self.num_local_experts,
|
local_expert_offset=self.moe_ep_rank * self.num_local_experts,
|
||||||
local_num_experts=self.num_local_experts,
|
local_num_experts=self.num_local_experts,
|
||||||
routing_method_type=self.routing_method_type,
|
routing_method_type=self.routing_method_type,
|
||||||
|
tune_max_num_tokens=next_power_of_2(hidden_states.shape[0]),
|
||||||
)
|
)
|
||||||
|
|
||||||
else:
|
else:
|
||||||
@@ -1228,6 +1230,7 @@ class FlashInferFP4MoE(FusedMoE):
|
|||||||
tile_tokens_dim=None,
|
tile_tokens_dim=None,
|
||||||
routing_method_type=routing_method_type,
|
routing_method_type=routing_method_type,
|
||||||
do_finalize=True,
|
do_finalize=True,
|
||||||
|
tune_max_num_tokens=next_power_of_2(hs_fp4.shape[0]),
|
||||||
output=symm_output,
|
output=symm_output,
|
||||||
)[0]
|
)[0]
|
||||||
|
|
||||||
|
|||||||
@@ -69,6 +69,7 @@ from sglang.srt.utils import (
|
|||||||
is_sm90_supported,
|
is_sm90_supported,
|
||||||
is_sm100_supported,
|
is_sm100_supported,
|
||||||
log_info_on_rank0,
|
log_info_on_rank0,
|
||||||
|
next_power_of_2,
|
||||||
print_warning_once,
|
print_warning_once,
|
||||||
set_weight_attrs,
|
set_weight_attrs,
|
||||||
use_intel_amx_backend,
|
use_intel_amx_backend,
|
||||||
@@ -1384,6 +1385,7 @@ class Fp8MoEMethod(FusedMoEMethodBase):
|
|||||||
tile_tokens_dim=None,
|
tile_tokens_dim=None,
|
||||||
routing_method_type=routing_method_type,
|
routing_method_type=routing_method_type,
|
||||||
use_shuffled_weight=False,
|
use_shuffled_weight=False,
|
||||||
|
tune_max_num_tokens=next_power_of_2(a_q.shape[0]),
|
||||||
)
|
)
|
||||||
else:
|
else:
|
||||||
routing_bias_cast = (
|
routing_bias_cast = (
|
||||||
@@ -1415,6 +1417,7 @@ class Fp8MoEMethod(FusedMoEMethodBase):
|
|||||||
),
|
),
|
||||||
use_routing_scales_on_input=False,
|
use_routing_scales_on_input=False,
|
||||||
routing_method_type=routing_method_type,
|
routing_method_type=routing_method_type,
|
||||||
|
tune_max_num_tokens=next_power_of_2(a_q.shape[0]),
|
||||||
)
|
)
|
||||||
|
|
||||||
def maybe_apply_hip_fused_experts(
|
def maybe_apply_hip_fused_experts(
|
||||||
|
|||||||
@@ -793,6 +793,7 @@ class ModelOptFp8MoEMethod(FusedMoEMethodBase):
|
|||||||
use_routing_scales_on_input=use_routing_scales_on_input,
|
use_routing_scales_on_input=use_routing_scales_on_input,
|
||||||
tile_tokens_dim=None,
|
tile_tokens_dim=None,
|
||||||
routing_method_type=routing_method_type,
|
routing_method_type=routing_method_type,
|
||||||
|
tune_max_num_tokens=next_power_of_2(x.shape[0]),
|
||||||
)
|
)
|
||||||
|
|
||||||
from sglang.srt.layers.moe.token_dispatcher import StandardCombineInput
|
from sglang.srt.layers.moe.token_dispatcher import StandardCombineInput
|
||||||
|
|||||||
@@ -2219,11 +2219,15 @@ class ModelRunner:
|
|||||||
|
|
||||||
def _should_run_flashinfer_autotune(self) -> bool:
|
def _should_run_flashinfer_autotune(self) -> bool:
|
||||||
"""Check if flashinfer autotune should be run."""
|
"""Check if flashinfer autotune should be run."""
|
||||||
if not self.server_args.enable_flashinfer_autotune:
|
if self.server_args.disable_flashinfer_autotune:
|
||||||
return False
|
return False
|
||||||
|
|
||||||
backend_str = self.server_args.attention_backend
|
backend_str = self.server_args.moe_runner_backend
|
||||||
if backend_str not in ["flashinfer", "trtllm_mla", "trtllm_mha"]:
|
if backend_str not in [
|
||||||
|
"flashinfer_trtllm",
|
||||||
|
"flashinfer_cutlass",
|
||||||
|
"flashinfer_mxfp4",
|
||||||
|
]:
|
||||||
return False
|
return False
|
||||||
|
|
||||||
major, _ = torch.cuda.get_device_capability()
|
major, _ = torch.cuda.get_device_capability()
|
||||||
|
|||||||
@@ -418,7 +418,7 @@ class ServerArgs:
|
|||||||
fp8_gemm_runner_backend: str = "auto"
|
fp8_gemm_runner_backend: str = "auto"
|
||||||
nsa_prefill_backend: str = "flashmla_sparse"
|
nsa_prefill_backend: str = "flashmla_sparse"
|
||||||
nsa_decode_backend: str = "fa3"
|
nsa_decode_backend: str = "fa3"
|
||||||
enable_flashinfer_autotune: bool = False
|
disable_flashinfer_autotune: bool = False
|
||||||
|
|
||||||
# Speculative decoding
|
# Speculative decoding
|
||||||
speculative_algorithm: Optional[str] = None
|
speculative_algorithm: Optional[str] = None
|
||||||
@@ -3313,10 +3313,10 @@ class ServerArgs:
|
|||||||
"SGLANG_ENABLE_FLASHINFER_FP8_GEMM and SGLANG_SUPPORT_CUTLASS_BLOCK_FP8.",
|
"SGLANG_ENABLE_FLASHINFER_FP8_GEMM and SGLANG_SUPPORT_CUTLASS_BLOCK_FP8.",
|
||||||
)
|
)
|
||||||
parser.add_argument(
|
parser.add_argument(
|
||||||
"--enable-flashinfer-autotune",
|
"--disable-flashinfer-autotune",
|
||||||
default=ServerArgs.enable_flashinfer_autotune,
|
default=ServerArgs.disable_flashinfer_autotune,
|
||||||
action="store_true",
|
action="store_true",
|
||||||
help="Enable FlashInfer autotuning for optimal kernel selection.",
|
help="Disable FlashInfer autotuning.",
|
||||||
)
|
)
|
||||||
|
|
||||||
# Speculative decoding
|
# Speculative decoding
|
||||||
|
|||||||
@@ -16,7 +16,7 @@ from sglang.test.test_utils import (
|
|||||||
)
|
)
|
||||||
|
|
||||||
FULL_DEEPSEEK_V3_FP4_MODEL_PATH = "nvidia/DeepSeek-V3-0324-FP4"
|
FULL_DEEPSEEK_V3_FP4_MODEL_PATH = "nvidia/DeepSeek-V3-0324-FP4"
|
||||||
SERVER_LAUNCH_TIMEOUT = 1000
|
SERVER_LAUNCH_TIMEOUT = 1200
|
||||||
|
|
||||||
|
|
||||||
class TestDeepseekV3FP4(CustomTestCase):
|
class TestDeepseekV3FP4(CustomTestCase):
|
||||||
|
|||||||
Reference in New Issue
Block a user