minke.yu/sglang
1
0
Fork 0
Code Issues Pull Requests Actions 12 Packages Projects Releases Wiki Activity

[AMD] deepseek-v4 clean env vars (#28920)

Browse Source
This commit is contained in:
Thomas Wang
2026-06-22 07:32:43 -07:00
committed by GitHub
parent dd39ef6cec
commit 04d952ea10
10 changed files with 29 additions and 108 deletions
Download Patch File Download Diff File Expand all files Collapse all files
python/sglang/srt/layers/attention/deepseek_v4_backend_hip_radix.py
+1 -3
View File
@@ -1472,13 +1472,11 @@ class DeepseekV4HipRadixBackend(
extra_indices.shape[-1] % 64 == 0 extra_indices.shape[-1] % 64 == 0
), f"{extra_indices.shape=}'s last dimension is not aligned to 64" ), f"{extra_indices.shape=}'s last dimension is not aligned to 64"
import os
from sglang.srt.layers.attention.hip_flash_mla import ( from sglang.srt.layers.attention.hip_flash_mla import (
flash_mla_with_kvcache_entrypoint, flash_mla_with_kvcache_entrypoint,
) )
backend = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "kernel") backend = envs.SGLANG_HACK_FLASHMLA_BACKEND.get()
input_dict = dict( input_dict = dict(
q=q, q=q,
k_cache=swa_k_cache, k_cache=swa_k_cache,
python/sglang/srt/layers/attention/dsv4/unified_kv_kernels/env_gate.py
+2 -5
View File
@@ -1,15 +1,12 @@
from __future__ import annotations from __future__ import annotations
import functools import functools
import os
from sglang.srt.environ import envs
from sglang.srt.utils import is_hip from sglang.srt.utils import is_hip
@functools.lru_cache(maxsize=1) @functools.lru_cache(maxsize=1)
def is_unified_kv_triton() -> bool: def is_unified_kv_triton() -> bool:
# unified_kv_triton is only implemented on HIP (ROCm) # unified_kv_triton is only implemented on HIP (ROCm)
return ( return is_hip() and envs.SGLANG_HACK_FLASHMLA_BACKEND.get() == "unified_kv_triton"
is_hip()
and os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "") == "unified_kv_triton"
)
python/sglang/srt/layers/attention/hip_flash_mla.py
+2 -3
View File
@@ -2,6 +2,7 @@ from typing import Any, Optional
import torch import torch
from sglang.srt.environ import envs
from sglang.srt.layers.quantization.fp8_kernel import is_fp8_fnuz from sglang.srt.layers.quantization.fp8_kernel import is_fp8_fnuz
from sglang.srt.utils import is_hip 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): def flash_mla_with_kvcache_entrypoint(backend: str, **kwargs):
if is_hip(): if is_hip():
import os backend = envs.SGLANG_HACK_FLASHMLA_BACKEND.get()
backend = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "tilelang")
else: else:
import sgl_kernel.flash_mla as flash_mla import sgl_kernel.flash_mla as flash_mla
python/sglang/srt/server_args.py
+12
View File
@@ -2626,6 +2626,18 @@ class ServerArgs:
envs.SGLANG_OPT_USE_TILELANG_MHC_PRE.set(False) envs.SGLANG_OPT_USE_TILELANG_MHC_PRE.set(False)
envs.SGLANG_OPT_DEEPGEMM_HC_PRENORM.set(False) envs.SGLANG_OPT_DEEPGEMM_HC_PRENORM.set(False)
envs.SGLANG_FP8_PAGED_MQA_LOGITS_TORCH.set(True) 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"]: elif model_arch in ["GptOssForCausalLM"]:
# Set attention backend for GPT-OSS # Set attention backend for GPT-OSS
test/registered/amd/test_deepseek_v4_flash_fp4.py
+2 -16
View File
@@ -36,26 +36,12 @@ DEEPSEEK_V4_FP4_MODEL_PATH = os.environ.get(
SERVER_LAUNCH_TIMEOUT = 3600 SERVER_LAUNCH_TIMEOUT = 3600
FLASHMLA_BACKEND = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "unified_kv_triton") 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 = { COMMON_ENV_VARS = {
"SGLANG_DEFAULT_THINKING": "1", "SGLANG_DEFAULT_THINKING": "1",
"SGLANG_DSV4_REASONING_EFFORT": "max", "SGLANG_DSV4_REASONING_EFFORT": "max",
"SGLANG_OPT_DEEPGEMM_HC_PRENORM": "false", "SGLANG_USE_ROCM700A": "0",
"SGLANG_USE_AITER": "1", "SGLANG_DP_USE_GATHERV": "1",
"SGLANG_USE_ROCM700A": "1",
"SGLANG_OPT_USE_FUSED_COMPRESS": "true",
"SGLANG_OPT_USE_FUSED_COMPRESS_TRITON": "true",
"SGLANG_HACK_FLASHMLA_BACKEND": FLASHMLA_BACKEND, "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", "AITER_BF16_FP8_MOE_BOUND": "0",
} }
test/registered/amd/test_deepseek_v4_flash_fp8.py
+2 -16
View File
@@ -36,26 +36,12 @@ DEEPSEEK_V4_FP8_MODEL_PATH = os.environ.get(
SERVER_LAUNCH_TIMEOUT = 3600 SERVER_LAUNCH_TIMEOUT = 3600
FLASHMLA_BACKEND = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "unified_kv_triton") 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 = { COMMON_ENV_VARS = {
"SGLANG_DEFAULT_THINKING": "1", "SGLANG_DEFAULT_THINKING": "1",
"SGLANG_DSV4_REASONING_EFFORT": "max", "SGLANG_DSV4_REASONING_EFFORT": "max",
"SGLANG_OPT_DEEPGEMM_HC_PRENORM": "false", "SGLANG_USE_ROCM700A": "0",
"SGLANG_USE_AITER": "1", "SGLANG_DP_USE_GATHERV": "1",
"SGLANG_USE_ROCM700A": "1",
"SGLANG_OPT_USE_FUSED_COMPRESS": "true",
"SGLANG_OPT_USE_FUSED_COMPRESS_TRITON": "true",
"SGLANG_HACK_FLASHMLA_BACKEND": FLASHMLA_BACKEND, "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", "AITER_BF16_FP8_MOE_BOUND": "0",
} }
test/registered/amd/test_deepseek_v4_pro_fp4.py
+2 -16
View File
@@ -38,26 +38,12 @@ DEEPSEEK_V4_PRO_FP4_MODEL_PATH = os.environ.get(
SERVER_LAUNCH_TIMEOUT = 5400 SERVER_LAUNCH_TIMEOUT = 5400
FLASHMLA_BACKEND = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "unified_kv_triton") 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 = { COMMON_ENV_VARS = {
"SGLANG_DEFAULT_THINKING": "1", "SGLANG_DEFAULT_THINKING": "1",
"SGLANG_DSV4_REASONING_EFFORT": "max", "SGLANG_DSV4_REASONING_EFFORT": "max",
"SGLANG_OPT_DEEPGEMM_HC_PRENORM": "false", "SGLANG_USE_ROCM700A": "0",
"SGLANG_USE_AITER": "1", "SGLANG_DP_USE_GATHERV": "1",
"SGLANG_USE_ROCM700A": "1",
"SGLANG_OPT_USE_FUSED_COMPRESS": "true",
"SGLANG_OPT_USE_FUSED_COMPRESS_TRITON": "true",
"SGLANG_HACK_FLASHMLA_BACKEND": FLASHMLA_BACKEND, "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", "AITER_BF16_FP8_MOE_BOUND": "0",
} }
test/registered/amd/test_deepseek_v4_pro_fp4_cp.py
+2 -16
View File
@@ -37,24 +37,10 @@ SERVER_LAUNCH_TIMEOUT = 5400
COMMON_ENV_VARS = { COMMON_ENV_VARS = {
"SGLANG_DEFAULT_THINKING": "1", "SGLANG_DEFAULT_THINKING": "1",
"SGLANG_DSV4_REASONING_EFFORT": "max", "SGLANG_DSV4_REASONING_EFFORT": "max",
"SGLANG_OPT_DEEPGEMM_HC_PRENORM": "false", "SGLANG_USE_ROCM700A": "0",
"SGLANG_USE_AITER": "1", "SGLANG_DP_USE_GATHERV": "1",
"SGLANG_USE_ROCM700A": "1",
"SGLANG_OPT_USE_FUSED_COMPRESS": "true",
"SGLANG_OPT_USE_FUSED_COMPRESS_TRITON": "true",
"SGLANG_HACK_FLASHMLA_BACKEND": "unified_kv_triton", "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", "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). # FP4 variant (matches test_deepseek_v4_pro_fp4.py; V4-Pro also auto-detects it).
test/registered/amd/test_deepseek_v4_pro_fp4_mtp.py
+2 -17
View File
@@ -38,28 +38,13 @@ FLASHMLA_BACKEND = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "unified_kv_tr
GSM8K_ACCURACY_THRESHOLD = 0.92 GSM8K_ACCURACY_THRESHOLD = 0.92
AVG_SPEC_ACCEPT_LENGTH_THRESHOLD = 2.8 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 = { COMMON_ENV_VARS = {
"SGLANG_DEFAULT_THINKING": "1", "SGLANG_DEFAULT_THINKING": "1",
"SGLANG_DSV4_REASONING_EFFORT": "max", "SGLANG_DSV4_REASONING_EFFORT": "max",
"SGLANG_OPT_DEEPGEMM_HC_PRENORM": "false", "SGLANG_USE_ROCM700A": "0",
"SGLANG_USE_AITER": "1", "SGLANG_DP_USE_GATHERV": "1",
"SGLANG_USE_ROCM700A": "1",
"SGLANG_OPT_USE_FUSED_COMPRESS": "true",
"SGLANG_OPT_USE_FUSED_COMPRESS_TRITON": "true",
"SGLANG_HACK_FLASHMLA_BACKEND": FLASHMLA_BACKEND, "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", "AITER_BF16_FP8_MOE_BOUND": "0",
"SGLANG_EAGER_INPUT_NO_COPY": "1",
} }
FP4_ENV_VARS = { FP4_ENV_VARS = {
test/registered/amd/test_deepseek_v4_pro_fp8.py
+2 -16
View File
@@ -38,26 +38,12 @@ DEEPSEEK_V4_PRO_FP8_MODEL_PATH = os.environ.get(
SERVER_LAUNCH_TIMEOUT = 5400 SERVER_LAUNCH_TIMEOUT = 5400
FLASHMLA_BACKEND = os.environ.get("SGLANG_HACK_FLASHMLA_BACKEND", "unified_kv_triton") 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 = { COMMON_ENV_VARS = {
"SGLANG_DEFAULT_THINKING": "1", "SGLANG_DEFAULT_THINKING": "1",
"SGLANG_DSV4_REASONING_EFFORT": "max", "SGLANG_DSV4_REASONING_EFFORT": "max",
"SGLANG_OPT_DEEPGEMM_HC_PRENORM": "false", "SGLANG_USE_ROCM700A": "0",
"SGLANG_USE_AITER": "1", "SGLANG_DP_USE_GATHERV": "1",
"SGLANG_USE_ROCM700A": "1",
"SGLANG_OPT_USE_FUSED_COMPRESS": "true",
"SGLANG_OPT_USE_FUSED_COMPRESS_TRITON": "true",
"SGLANG_HACK_FLASHMLA_BACKEND": FLASHMLA_BACKEND, "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", "AITER_BF16_FP8_MOE_BOUND": "0",
} }