From 9496bbd7b11c7ce7ac96b339c7a1fb23d0406545 Mon Sep 17 00:00:00 2001 From: fxmarty-amd Date: Fri, 27 Feb 2026 10:43:34 +0100 Subject: [PATCH] [AMD] Use `tilelang` as default NSA attention backend dispatch on AMD Instinct (#18319) --- python/sglang/srt/layers/attention/nsa_backend.py | 4 +++- python/sglang/srt/server_args.py | 5 ++++- 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/python/sglang/srt/layers/attention/nsa_backend.py b/python/sglang/srt/layers/attention/nsa_backend.py index 170f75e3e..649ca9e50 100644 --- a/python/sglang/srt/layers/attention/nsa_backend.py +++ b/python/sglang/srt/layers/attention/nsa_backend.py @@ -1477,7 +1477,9 @@ class NativeSparseAttnBackend( page_size=1, ) else: - raise ValueError(f"Unsupported {nsa_impl = }") + raise ValueError( + f"Unsupported {nsa_impl = } for forward_extend. Consider using an other attention backend." + ) def forward_decode( self, diff --git a/python/sglang/srt/server_args.py b/python/sglang/srt/server_args.py index a91e17d80..27196c549 100644 --- a/python/sglang/srt/server_args.py +++ b/python/sglang/srt/server_args.py @@ -1150,7 +1150,10 @@ class ServerArgs: user_set_prefill = self.nsa_prefill_backend is not None user_set_decode = self.nsa_decode_backend is not None - if kv_cache_dtype == "fp8_e4m3": + if not user_set_prefill and not user_set_decode and is_hip(): + self.nsa_prefill_backend = "tilelang" + self.nsa_decode_backend = "tilelang" + elif kv_cache_dtype == "fp8_e4m3": # flashmla_auto dispatches to flashmla_sparse/flashmla_kv based on hardware and heuristics if not user_set_prefill: self.nsa_prefill_backend = "flashmla_auto"