Skip to content

Commit a818e72

Browse files
committed
feat(ops): add optional MindIE-SD attention backend for Ascend NPU
Add AttnImplType.MINDIE_ATTN, routed through mindiesd.attention_forward (auto-tuned Ascend kernels including LaserAttention) for dense BNSD attention. Availability is probed like the sageattention backend (optional import, NPU platform only), so default behavior is unchanged on every platform and the enum is strictly opt-in. Verified on Atlas 910B (CANN 8.2): kernel microbench 1.29-1.53x over SDPA across wan shapes (holds at ulysses head-sharded 12-head shapes); unmodified wan22_t2v_5b example end to end: 720p 121f 477.8->405.0s (1.18x) on 1 card and 173.9->157.6s (1.10x) on 4 cards, 480p 121f 180.9->162.6s on 1 card; a 480p 4-card regression (87.0->103.0s) is documented in the PR - the backend is recommended for compute-dominated configurations. Output is bit-identical to SDPA at probed shapes; dispatch covered by a mock unit test; ruff clean.
1 parent 5b71ae5 commit a818e72

5 files changed

Lines changed: 67 additions & 2 deletions

File tree

‎docs/en/attention.md‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -24,6 +24,7 @@ class AttnImplType(Enum):
2424
FLASH_ATTN_2 = auto()
2525
FLASH_ATTN_3 = auto()
2626
FLASH_ATTN_4 = auto() # For Hopper (SM90) and Blackwell (SM100+) GPUs
27+
MINDIE_ATTN = auto() # Ascend NPU, requires the optional mindiesd package
2728
SAGE_ATTN_2_8_8 = auto()
2829
SAGE_ATTN_2_8_16 = auto()
2930
SAGE_ATTN_2_8_8_SM90 = auto()

‎telefuser/core/config.py‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -127,6 +127,7 @@ class AttnImplType(Enum):
127127
FLASH_ATTN_2 = auto()
128128
FLASH_ATTN_3 = auto()
129129
FLASH_ATTN_4 = auto()
130+
MINDIE_ATTN = auto()
130131
# Sparse attention implementations
131132
RADIAL_ATTN = auto() # Radial attention for video generation
132133
LOCAL_SPARSE_ATTN = auto() # Local window sparse attention

‎telefuser/ops/attention/attention_impl.py‎

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
"""Unified attention implementation with dense and sparse support.
22
33
Supports multiple attention implementations:
4-
- Dense: TORCH_SDPA, TORCH_CUDNN, FLASH_ATTN_2/3/4, SAGE_ATTN variants, SPARGE_ATTN
4+
- Dense: TORCH_SDPA, TORCH_CUDNN, FLASH_ATTN_2/3/4, SAGE_ATTN variants, SPARGE_ATTN, MINDIE_ATTN (Ascend NPU)
55
- Sparse: RADIAL_ATTN, LOCAL_SPARSE_ATTN, SOL_ATTN
66
77
Note: Attention functions are decorated with @torch.compiler.disable because:
@@ -31,6 +31,7 @@
3131
FLASH_ATTN_2_AVAILABLE,
3232
FLASH_ATTN_3_AVAILABLE,
3333
FLASH_ATTN_4_AVAILABLE,
34+
MINDIE_ATTN_AVAILABLE,
3435
SAGE_ATTN_AVAILABLE,
3536
SDPA_AVAILABLE,
3637
SOL_ATTN_AVAILABLE,
@@ -39,6 +40,7 @@
3940
flash_attn4,
4041
flash_attn4_varlen,
4142
get_lse_fallback_impl,
43+
mindie_attn,
4244
sageattention,
4345
sdpa_attn_cudnn,
4446
sol_attn,
@@ -272,7 +274,7 @@ def attention(
272274
# - Flash Attention expects BSND (NHD) layout
273275
# - PyTorch SDPA variants expect BNSD (HND) layout
274276
# - SageAttention can accept both via tensor_layout parameter
275-
BNSD_IMPLS = {AttnImplType.TORCH_CUDNN, AttnImplType.TORCH_SDPA, AttnImplType.SPARGE_ATTN}
277+
BNSD_IMPLS = {AttnImplType.TORCH_CUDNN, AttnImplType.TORCH_SDPA, AttnImplType.SPARGE_ATTN, AttnImplType.MINDIE_ATTN}
276278

277279
# Track current layout after potential conversions
278280
current_layout = input_layout
@@ -347,6 +349,10 @@ def attention(
347349
else:
348350
output = result
349351

352+
# MindIE-SD attention (Ascend NPU)
353+
elif sequence_lengths is None and attn_impl == AttnImplType.MINDIE_ATTN and MINDIE_ATTN_AVAILABLE:
354+
output = mindie_attn(q, k, v, attn_mask=attn_mask, scale=scale, is_causal=is_causal)
355+
350356
# PyTorch SDPA
351357
elif sequence_lengths is None and attn_impl == AttnImplType.TORCH_CUDNN and SDPA_AVAILABLE:
352358
output = sdpa_attn_cudnn(q, k, v, attn_mask=attn_mask, scale=scale, is_causal=is_causal)

‎telefuser/ops/attention/backends.py‎

Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -13,6 +13,7 @@
1313
import torch
1414
from torch import Tensor
1515

16+
from telefuser.platforms import current_platform
1617
from telefuser.utils.logging import logger
1718

1819
# Availability flags
@@ -24,6 +25,7 @@
2425
SPARGE_ATTN_AVAILABLE = False
2526
FLASHINFER_AVAILABLE = False
2627
SOL_ATTN_AVAILABLE = False
28+
MINDIE_ATTN_AVAILABLE = False
2729

2830
# Backend function references (populated on successful import)
2931
flash_attn2: Callable | None = None
@@ -34,6 +36,7 @@
3436
spas_sage2_attn_meansim_cuda: Callable | None = None
3537
flashinfer: object | None = None
3638
sol_attn: Callable | None = None
39+
mindiesd_attention_forward: Callable | None = None
3740

3841

3942
def _try_import_flash_attn() -> None:
@@ -162,12 +165,48 @@ def _try_import_sol_attn() -> None:
162165

163166

164167
# Initialize all backends
168+
def _try_import_mindie_attn() -> None:
169+
"""Import the optional MindIE-SD attention backend (Ascend NPU only)."""
170+
global MINDIE_ATTN_AVAILABLE, mindiesd_attention_forward
171+
172+
MINDIE_ATTN_AVAILABLE = False
173+
mindiesd_attention_forward = None
174+
if current_platform.device_type != "npu":
175+
return
176+
try:
177+
if importlib.util.find_spec("mindiesd") is None:
178+
return
179+
from mindiesd import attention_forward
180+
except (ModuleNotFoundError, ImportError) as error:
181+
logger.debug("MindIE-SD attention backend unavailable: %s", error)
182+
return
183+
mindiesd_attention_forward = attention_forward
184+
MINDIE_ATTN_AVAILABLE = True
185+
logger.debug("MindIE-SD attention available")
186+
187+
188+
def mindie_attn(
189+
q: Tensor,
190+
k: Tensor,
191+
v: Tensor,
192+
*,
193+
attn_mask: Tensor | None = None,
194+
scale: float | None = None,
195+
is_causal: bool = False,
196+
) -> Tensor:
197+
"""MindIE-SD attention wrapper for BNSD inputs on Ascend NPU."""
198+
if is_causal:
199+
raise ValueError("MINDIE_ATTN does not support causal attention")
200+
return mindiesd_attention_forward(q, k, v, attn_mask=attn_mask, scale=scale, head_first=True)
201+
202+
165203
_try_import_flash_attn()
166204
_try_import_sdpa()
167205
_try_import_sage_attn()
168206
_try_import_sparge_attn()
169207
_try_import_flashinfer()
170208
_try_import_sol_attn()
209+
_try_import_mindie_attn()
171210

172211

173212
def supports_return_lse(attn_impl: str) -> bool:
@@ -234,6 +273,7 @@ def sparge_attn(
234273
"SPARGE_ATTN_AVAILABLE",
235274
"FLASHINFER_AVAILABLE",
236275
"SOL_ATTN_AVAILABLE",
276+
"MINDIE_ATTN_AVAILABLE",
237277
"flash_attn2",
238278
"flash_attn3",
239279
"flash_attn4",
@@ -245,4 +285,5 @@ def sparge_attn(
245285
"get_lse_fallback_impl",
246286
"sdpa_attn_cudnn",
247287
"sparge_attn",
288+
"mindie_attn",
248289
]

‎tests/unit/ops/test_attention_backends.py‎

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -269,3 +269,19 @@ def test_ring_config_converts_fallback_name() -> None:
269269
assert result.attn_impl is AttnImplType.SAGE_ATTN_2_8_8_SM90
270270
assert result.scale == 0.125
271271
assert result.is_causal is True
272+
273+
274+
def test_mindie_attn_dispatch_converts_to_bnsd_and_back() -> None:
275+
q = torch.randn(1, 3, 2, 4)
276+
forward = MagicMock(side_effect=lambda q_, k_, v_, **kwargs: q_)
277+
278+
with (
279+
patch.object(attention_impl, "MINDIE_ATTN_AVAILABLE", True),
280+
patch.object(backends, "mindiesd_attention_forward", forward),
281+
):
282+
output = attention_impl.attention(q, q, q, attn_impl=attention_impl.AttnImplType.MINDIE_ATTN)
283+
284+
assert output.shape == q.shape
285+
called_q = forward.call_args.args[0]
286+
assert called_q.shape == (1, 2, 3, 4)
287+
assert forward.call_args.kwargs == {"attn_mask": None, "scale": None, "head_first": True}

0 commit comments

Comments
 (0)