[AMD] Prefer HIP Top-K for GLM-5.x on ROCm (#39631)
This commit is contained in:
@@ -340,24 +340,9 @@ def handle_model_specific_adjustments(server_args: Any):
|
||||
|
||||
run_post_process_pass(server_args, _deepseek_moe_quant_resolution)
|
||||
if get_platform().is_hip:
|
||||
if is_deepseek_dsa(hf_config):
|
||||
# The fused top-k v2 kernel (topk_transform_paged_v2) is a
|
||||
# CUDA/Hopper-only path: its JIT source includes
|
||||
# <cooperative_groups.h> and uses cg::this_cluster()
|
||||
# (thread-block clusters), neither of which exists on ROCm,
|
||||
# so it fails to JIT-compile on gfx9xx during CUDA-graph
|
||||
# capture. DeepSeek-V4 already disables it on HIP; mirror that
|
||||
# here for the rest of the DSA family (DeepSeek-V3.2 /
|
||||
# GLM-5.x) that shares the same decode top-k path.
|
||||
if is_deepseek_dsa(hf_config) and not envs.SGLANG_OPT_USE_TOPK_V2.is_set():
|
||||
# Prefer HIP top-k by default while honoring an explicit selection.
|
||||
envs.SGLANG_OPT_USE_TOPK_V2.set(False)
|
||||
if model_arch == "GlmMoeDsaForCausalLM":
|
||||
# Open the fused top-k v2 kernel for the GLM-5.x DSA
|
||||
# family on ROCm: it shares this decode top-k path, and
|
||||
# the kernel's ROCm build compiles the streaming levels
|
||||
# on gfx9xx. Order is load-bearing: the blanket disable
|
||||
# above `set`s the variable unconditionally, so this has
|
||||
# to follow it.
|
||||
envs.SGLANG_OPT_USE_TOPK_V2.set(True)
|
||||
if not resolved_view(server_args).enable_dp_attention and cfg.nnodes == 1:
|
||||
# TODO (Hubert): Put this back later
|
||||
# server_args.enable_aiter_allreduce_fusion = True
|
||||
|
||||
Reference in New Issue
Block a user