[AMD] Use tilelang as default NSA attention backend dispatch on AMD Instinct (#18319)
This commit is contained in:
@@ -1477,7 +1477,9 @@ class NativeSparseAttnBackend(
|
|||||||
page_size=1,
|
page_size=1,
|
||||||
)
|
)
|
||||||
else:
|
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(
|
def forward_decode(
|
||||||
self,
|
self,
|
||||||
|
|||||||
@@ -1150,7 +1150,10 @@ class ServerArgs:
|
|||||||
user_set_prefill = self.nsa_prefill_backend is not None
|
user_set_prefill = self.nsa_prefill_backend is not None
|
||||||
user_set_decode = self.nsa_decode_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
|
# flashmla_auto dispatches to flashmla_sparse/flashmla_kv based on hardware and heuristics
|
||||||
if not user_set_prefill:
|
if not user_set_prefill:
|
||||||
self.nsa_prefill_backend = "flashmla_auto"
|
self.nsa_prefill_backend = "flashmla_auto"
|
||||||
|
|||||||
Reference in New Issue
Block a user