From 04d952ea102d0bd922073f6d7a0b97b71e21ca0a Mon Sep 17 00:00:00 2001 From: Thomas Wang Date: Mon, 22 Jun 2026 22:32:43 +0800 Subject: [PATCH] [AMD] deepseek-v4 clean env vars (#28920) --- .../deepseek_v4_backend_hip_radix.py | 4 +--- .../dsv4/unified_kv_kernels/env_gate.py | 7 ++----- .../srt/layers/attention/hip_flash_mla.py | 5 ++--- python/sglang/srt/server_args.py | 12 ++++++++++++ .../amd/test_deepseek_v4_flash_fp4.py | 18 ++---------------- .../amd/test_deepseek_v4_flash_fp8.py | 18 ++---------------- .../amd/test_deepseek_v4_pro_fp4.py | 18 ++---------------- .../amd/test_deepseek_v4_pro_fp4_cp.py | 18 ++---------------- .../amd/test_deepseek_v4_pro_fp4_mtp.py | 19 ++----------------- .../amd/test_deepseek_v4_pro_fp8.py | 18 ++---------------- 10 files changed, 29 insertions(+), 108 deletions(-) diff --git a/python/sglang/srt/layers/attention/deepseek_v4_backend_hip_radix.py b/python/sglang/srt/layers/attention/deepseek_v4_backend_hip_radix.py index a5e55570b..c2d29b100 100644 --- a/python/sglang/srt/layers/attention/deepseek_v4_backend_hip_radix.py +++ b/python/sglang/srt/layers/attention/deepseek_v4_backend_hip_radix.py @@ -1472,13 +1472,11 @@ class DeepseekV4HipRadixBackend( extra_indices.shape[-1] % 64 == 0 ), f"{extra_indices.shape=}'s last dimension is not aligned to 64" - import os - from sglang.srt.layers.attention.hip_flash_mla import ( flash_mla_with_kvcache_entrypoint, ) - backend = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "kernel") + backend = envs.SGLANG_HACK_FLASHMLA_BACKEND.get() input_dict = dict( q=q, k_cache=swa_k_cache, diff --git a/python/sglang/srt/layers/attention/dsv4/unified_kv_kernels/env_gate.py b/python/sglang/srt/layers/attention/dsv4/unified_kv_kernels/env_gate.py index c6636f9dc..c55ba905e 100644 --- a/python/sglang/srt/layers/attention/dsv4/unified_kv_kernels/env_gate.py +++ b/python/sglang/srt/layers/attention/dsv4/unified_kv_kernels/env_gate.py @@ -1,15 +1,12 @@ from __future__ import annotations import functools -import os +from sglang.srt.environ import envs from sglang.srt.utils import is_hip @functools.lru_cache(maxsize=1) def is_unified_kv_triton() -> bool: # unified_kv_triton is only implemented on HIP (ROCm) - return ( - is_hip() - and os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "") == "unified_kv_triton" - ) + return is_hip() and envs.SGLANG_HACK_FLASHMLA_BACKEND.get() == "unified_kv_triton" diff --git a/python/sglang/srt/layers/attention/hip_flash_mla.py b/python/sglang/srt/layers/attention/hip_flash_mla.py index c26705c0c..460a8ebf6 100644 --- a/python/sglang/srt/layers/attention/hip_flash_mla.py +++ b/python/sglang/srt/layers/attention/hip_flash_mla.py @@ -2,6 +2,7 @@ from typing import Any, Optional import torch +from sglang.srt.environ import envs from sglang.srt.layers.quantization.fp8_kernel import is_fp8_fnuz from sglang.srt.utils import is_hip @@ -10,9 +11,7 @@ FP8_DTYPE = torch.float8_e4m3fnuz if is_fp8_fnuz() else torch.float8_e4m3fn def flash_mla_with_kvcache_entrypoint(backend: str, **kwargs): if is_hip(): - import os - - backend = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "tilelang") + backend = envs.SGLANG_HACK_FLASHMLA_BACKEND.get() else: import sgl_kernel.flash_mla as flash_mla diff --git a/python/sglang/srt/server_args.py b/python/sglang/srt/server_args.py index 91072b501..bbc2905d8 100644 --- a/python/sglang/srt/server_args.py +++ b/python/sglang/srt/server_args.py @@ -2626,6 +2626,18 @@ class ServerArgs: envs.SGLANG_OPT_USE_TILELANG_MHC_PRE.set(False) envs.SGLANG_OPT_DEEPGEMM_HC_PRENORM.set(False) envs.SGLANG_FP8_PAGED_MQA_LOGITS_TORCH.set(True) + elif is_hip(): + envs.SGLANG_OPT_DEEPGEMM_HC_PRENORM.set(False) + envs.SGLANG_OPT_USE_FUSED_COMPRESS.set(True) + envs.SGLANG_OPT_FP8_WO_A_GEMM.set(False) + envs.SGLANG_OPT_USE_JIT_INDEXER_METADATA.set(False) + envs.SGLANG_OPT_USE_TOPK_V2.set(False) + envs.SGLANG_OPT_USE_AITER_INDEXER.set(True) + envs.SGLANG_OPT_USE_TILELANG_MHC_PRE.set(False) + envs.SGLANG_OPT_USE_TILELANG_MHC_POST.set(False) + envs.SGLANG_FP8_PAGED_MQA_LOGITS_TORCH.set(True) + envs.SGLANG_OPT_USE_MULTI_STREAM_OVERLAP.set(False) + envs.SGLANG_EAGER_INPUT_NO_COPY.set(True) elif model_arch in ["GptOssForCausalLM"]: # Set attention backend for GPT-OSS diff --git a/test/registered/amd/test_deepseek_v4_flash_fp4.py b/test/registered/amd/test_deepseek_v4_flash_fp4.py index a0842bd0e..3ff281d12 100644 --- a/test/registered/amd/test_deepseek_v4_flash_fp4.py +++ b/test/registered/amd/test_deepseek_v4_flash_fp4.py @@ -36,26 +36,12 @@ DEEPSEEK_V4_FP4_MODEL_PATH = os.environ.get( SERVER_LAUNCH_TIMEOUT = 3600 FLASHMLA_BACKEND = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "unified_kv_triton") -# Common DeepSeek-V4 env vars (AMD ROCm 7.2 path: AITER indexer + triton attn + ROCm700A). COMMON_ENV_VARS = { "SGLANG_DEFAULT_THINKING": "1", "SGLANG_DSV4_REASONING_EFFORT": "max", - "SGLANG_OPT_DEEPGEMM_HC_PRENORM": "false", - "SGLANG_USE_AITER": "1", - "SGLANG_USE_ROCM700A": "1", - "SGLANG_OPT_USE_FUSED_COMPRESS": "true", - "SGLANG_OPT_USE_FUSED_COMPRESS_TRITON": "true", + "SGLANG_USE_ROCM700A": "0", + "SGLANG_DP_USE_GATHERV": "1", "SGLANG_HACK_FLASHMLA_BACKEND": FLASHMLA_BACKEND, - "SGLANG_OPT_FP8_WO_A_GEMM": "false", - "SGLANG_OPT_USE_JIT_INDEXER_METADATA": "false", - "SGLANG_OPT_USE_TOPK_V2": "false", - "SGLANG_OPT_USE_AITER_INDEXER": "true", - "SGLANG_OPT_USE_TILELANG_INDEXER": "false", - "SGLANG_OPT_USE_TILELANG_MHC_PRE": "false", - "SGLANG_OPT_USE_TILELANG_MHC_POST": "false", - "SGLANG_FP8_PAGED_MQA_LOGITS_TORCH": "1", - "SGLANG_OPT_USE_MULTI_STREAM_OVERLAP": "false", - "SGLANG_ROCM_USE_MULTI_STREAM": "false", "AITER_BF16_FP8_MOE_BOUND": "0", } diff --git a/test/registered/amd/test_deepseek_v4_flash_fp8.py b/test/registered/amd/test_deepseek_v4_flash_fp8.py index d8187c03e..f6d6d3859 100644 --- a/test/registered/amd/test_deepseek_v4_flash_fp8.py +++ b/test/registered/amd/test_deepseek_v4_flash_fp8.py @@ -36,26 +36,12 @@ DEEPSEEK_V4_FP8_MODEL_PATH = os.environ.get( SERVER_LAUNCH_TIMEOUT = 3600 FLASHMLA_BACKEND = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "unified_kv_triton") -# Common DeepSeek-V4 env vars (AMD ROCm 7.2 path: AITER indexer + triton attn + ROCm700A). COMMON_ENV_VARS = { "SGLANG_DEFAULT_THINKING": "1", "SGLANG_DSV4_REASONING_EFFORT": "max", - "SGLANG_OPT_DEEPGEMM_HC_PRENORM": "false", - "SGLANG_USE_AITER": "1", - "SGLANG_USE_ROCM700A": "1", - "SGLANG_OPT_USE_FUSED_COMPRESS": "true", - "SGLANG_OPT_USE_FUSED_COMPRESS_TRITON": "true", + "SGLANG_USE_ROCM700A": "0", + "SGLANG_DP_USE_GATHERV": "1", "SGLANG_HACK_FLASHMLA_BACKEND": FLASHMLA_BACKEND, - "SGLANG_OPT_FP8_WO_A_GEMM": "false", - "SGLANG_OPT_USE_JIT_INDEXER_METADATA": "false", - "SGLANG_OPT_USE_TOPK_V2": "false", - "SGLANG_OPT_USE_AITER_INDEXER": "true", - "SGLANG_OPT_USE_TILELANG_INDEXER": "false", - "SGLANG_OPT_USE_TILELANG_MHC_PRE": "false", - "SGLANG_OPT_USE_TILELANG_MHC_POST": "false", - "SGLANG_FP8_PAGED_MQA_LOGITS_TORCH": "1", - "SGLANG_OPT_USE_MULTI_STREAM_OVERLAP": "false", - "SGLANG_ROCM_USE_MULTI_STREAM": "false", "AITER_BF16_FP8_MOE_BOUND": "0", } diff --git a/test/registered/amd/test_deepseek_v4_pro_fp4.py b/test/registered/amd/test_deepseek_v4_pro_fp4.py index ca0560d4f..b1880ddd2 100644 --- a/test/registered/amd/test_deepseek_v4_pro_fp4.py +++ b/test/registered/amd/test_deepseek_v4_pro_fp4.py @@ -38,26 +38,12 @@ DEEPSEEK_V4_PRO_FP4_MODEL_PATH = os.environ.get( SERVER_LAUNCH_TIMEOUT = 5400 FLASHMLA_BACKEND = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "unified_kv_triton") -# Common DeepSeek-V4 env vars (AMD ROCm 7.2 path: AITER indexer + triton attn + ROCm700A). COMMON_ENV_VARS = { "SGLANG_DEFAULT_THINKING": "1", "SGLANG_DSV4_REASONING_EFFORT": "max", - "SGLANG_OPT_DEEPGEMM_HC_PRENORM": "false", - "SGLANG_USE_AITER": "1", - "SGLANG_USE_ROCM700A": "1", - "SGLANG_OPT_USE_FUSED_COMPRESS": "true", - "SGLANG_OPT_USE_FUSED_COMPRESS_TRITON": "true", + "SGLANG_USE_ROCM700A": "0", + "SGLANG_DP_USE_GATHERV": "1", "SGLANG_HACK_FLASHMLA_BACKEND": FLASHMLA_BACKEND, - "SGLANG_OPT_FP8_WO_A_GEMM": "false", - "SGLANG_OPT_USE_JIT_INDEXER_METADATA": "false", - "SGLANG_OPT_USE_TOPK_V2": "false", - "SGLANG_OPT_USE_AITER_INDEXER": "true", - "SGLANG_OPT_USE_TILELANG_INDEXER": "false", - "SGLANG_OPT_USE_TILELANG_MHC_PRE": "false", - "SGLANG_OPT_USE_TILELANG_MHC_POST": "false", - "SGLANG_FP8_PAGED_MQA_LOGITS_TORCH": "1", - "SGLANG_OPT_USE_MULTI_STREAM_OVERLAP": "false", - "SGLANG_ROCM_USE_MULTI_STREAM": "false", "AITER_BF16_FP8_MOE_BOUND": "0", } diff --git a/test/registered/amd/test_deepseek_v4_pro_fp4_cp.py b/test/registered/amd/test_deepseek_v4_pro_fp4_cp.py index 5a2732f71..cf3f198c5 100644 --- a/test/registered/amd/test_deepseek_v4_pro_fp4_cp.py +++ b/test/registered/amd/test_deepseek_v4_pro_fp4_cp.py @@ -37,24 +37,10 @@ SERVER_LAUNCH_TIMEOUT = 5400 COMMON_ENV_VARS = { "SGLANG_DEFAULT_THINKING": "1", "SGLANG_DSV4_REASONING_EFFORT": "max", - "SGLANG_OPT_DEEPGEMM_HC_PRENORM": "false", - "SGLANG_USE_AITER": "1", - "SGLANG_USE_ROCM700A": "1", - "SGLANG_OPT_USE_FUSED_COMPRESS": "true", - "SGLANG_OPT_USE_FUSED_COMPRESS_TRITON": "true", + "SGLANG_USE_ROCM700A": "0", + "SGLANG_DP_USE_GATHERV": "1", "SGLANG_HACK_FLASHMLA_BACKEND": "unified_kv_triton", - "SGLANG_OPT_FP8_WO_A_GEMM": "false", - "SGLANG_OPT_USE_JIT_INDEXER_METADATA": "false", - "SGLANG_OPT_USE_TOPK_V2": "false", - "SGLANG_OPT_USE_AITER_INDEXER": "true", - "SGLANG_OPT_USE_TILELANG_INDEXER": "false", - "SGLANG_OPT_USE_TILELANG_MHC_PRE": "false", - "SGLANG_OPT_USE_TILELANG_MHC_POST": "false", - "SGLANG_FP8_PAGED_MQA_LOGITS_TORCH": "1", - "SGLANG_OPT_USE_MULTI_STREAM_OVERLAP": "false", - "SGLANG_ROCM_USE_MULTI_STREAM": "false", "AITER_BF16_FP8_MOE_BOUND": "0", - "SGLANG_EAGER_INPUT_NO_COPY": "false", } # FP4 variant (matches test_deepseek_v4_pro_fp4.py; V4-Pro also auto-detects it). diff --git a/test/registered/amd/test_deepseek_v4_pro_fp4_mtp.py b/test/registered/amd/test_deepseek_v4_pro_fp4_mtp.py index 39613e512..97f5e94ed 100644 --- a/test/registered/amd/test_deepseek_v4_pro_fp4_mtp.py +++ b/test/registered/amd/test_deepseek_v4_pro_fp4_mtp.py @@ -38,28 +38,13 @@ FLASHMLA_BACKEND = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "unified_kv_tr GSM8K_ACCURACY_THRESHOLD = 0.92 AVG_SPEC_ACCEPT_LENGTH_THRESHOLD = 2.8 -# Common DeepSeek-V4 env vars (AMD ROCm 7.2 path: AITER indexer + triton attn + ROCm700A). COMMON_ENV_VARS = { "SGLANG_DEFAULT_THINKING": "1", "SGLANG_DSV4_REASONING_EFFORT": "max", - "SGLANG_OPT_DEEPGEMM_HC_PRENORM": "false", - "SGLANG_USE_AITER": "1", - "SGLANG_USE_ROCM700A": "1", - "SGLANG_OPT_USE_FUSED_COMPRESS": "true", - "SGLANG_OPT_USE_FUSED_COMPRESS_TRITON": "true", + "SGLANG_USE_ROCM700A": "0", + "SGLANG_DP_USE_GATHERV": "1", "SGLANG_HACK_FLASHMLA_BACKEND": FLASHMLA_BACKEND, - "SGLANG_OPT_FP8_WO_A_GEMM": "false", - "SGLANG_OPT_USE_JIT_INDEXER_METADATA": "false", - "SGLANG_OPT_USE_TOPK_V2": "false", - "SGLANG_OPT_USE_AITER_INDEXER": "true", - "SGLANG_OPT_USE_TILELANG_INDEXER": "false", - "SGLANG_OPT_USE_TILELANG_MHC_PRE": "false", - "SGLANG_OPT_USE_TILELANG_MHC_POST": "false", - "SGLANG_FP8_PAGED_MQA_LOGITS_TORCH": "1", - "SGLANG_OPT_USE_MULTI_STREAM_OVERLAP": "false", - "SGLANG_ROCM_USE_MULTI_STREAM": "false", "AITER_BF16_FP8_MOE_BOUND": "0", - "SGLANG_EAGER_INPUT_NO_COPY": "1", } FP4_ENV_VARS = { diff --git a/test/registered/amd/test_deepseek_v4_pro_fp8.py b/test/registered/amd/test_deepseek_v4_pro_fp8.py index f1f9618cd..ea74723ad 100644 --- a/test/registered/amd/test_deepseek_v4_pro_fp8.py +++ b/test/registered/amd/test_deepseek_v4_pro_fp8.py @@ -38,26 +38,12 @@ DEEPSEEK_V4_PRO_FP8_MODEL_PATH = os.environ.get( SERVER_LAUNCH_TIMEOUT = 5400 FLASHMLA_BACKEND = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "unified_kv_triton") -# Common DeepSeek-V4 env vars (AMD ROCm 7.2 path: AITER indexer + triton attn + ROCm700A). COMMON_ENV_VARS = { "SGLANG_DEFAULT_THINKING": "1", "SGLANG_DSV4_REASONING_EFFORT": "max", - "SGLANG_OPT_DEEPGEMM_HC_PRENORM": "false", - "SGLANG_USE_AITER": "1", - "SGLANG_USE_ROCM700A": "1", - "SGLANG_OPT_USE_FUSED_COMPRESS": "true", - "SGLANG_OPT_USE_FUSED_COMPRESS_TRITON": "true", + "SGLANG_USE_ROCM700A": "0", + "SGLANG_DP_USE_GATHERV": "1", "SGLANG_HACK_FLASHMLA_BACKEND": FLASHMLA_BACKEND, - "SGLANG_OPT_FP8_WO_A_GEMM": "false", - "SGLANG_OPT_USE_JIT_INDEXER_METADATA": "false", - "SGLANG_OPT_USE_TOPK_V2": "false", - "SGLANG_OPT_USE_AITER_INDEXER": "true", - "SGLANG_OPT_USE_TILELANG_INDEXER": "false", - "SGLANG_OPT_USE_TILELANG_MHC_PRE": "false", - "SGLANG_OPT_USE_TILELANG_MHC_POST": "false", - "SGLANG_FP8_PAGED_MQA_LOGITS_TORCH": "1", - "SGLANG_OPT_USE_MULTI_STREAM_OVERLAP": "false", - "SGLANG_ROCM_USE_MULTI_STREAM": "false", "AITER_BF16_FP8_MOE_BOUND": "0", }