huggingface / huggingface/diffusers
cuDNN attention CP path consumes a bool attn_mask as an additive 0/1 bias
- Lingua principale
- Python
- Stelle
- 34.5k
- Fork
- 7.3k
- Merge medio
- 3g 3h
- PR unite (30g)
- 91
Descrizione
# cuDNN attention CP path consumes a bool `attn_mask` as an additive 0/1 bias
Split out of #14338 so that issue can be closed cleanly by #14341, which fixes the
K/V transpose but not this.
## What's wrong
`_cudnn_attention_forward_op` passes `attn_mask` straight into
`torch.ops.aten._scaled_dot_product_cudnn_attention`:
```python
# src/diffusers/models/attention_dispatch.py
out, lse, ... = torch.ops.aten._scaled_dot_product_cudnn_attention(
query=query, key=key, value=value,
attn_bias=attn_mask, # <-- additive bias slot
...
)
```
The public `F.scaled_dot_product_attention` converts a boolean mask to an additive
`-inf`/`0` bias before dispatching. The raw ATen op does not. So a `torch.bool` mask
is consumed as an additive **0/1**: "masked" positions get `+0.0` and "keep" positions
get `+1.0`, which is neither a mask nor a no-op — it perturbs the scores of exactly the
positions that were supposed to be kept, and does not suppress the ones that were
supposed to be dropped.
`_native_flash_attention_forward_op` takes the same `attn_mask` argument through the
same wrapper, so it is worth checking whether it has the same exposure.
## Repro
With the conversion removed (i.e. current `main`), comparing the op against
`F.scaled_dot_product_attention` given the same boolean mask, bf16, `B=1 H=2 S=16 D=64`,
one partially masked row and one fully masked row:
```
partial FAIL
fully_masked_row FAIL
```
With a `-inf`/`0` conversion in front of the call, both pass.
## Fix
Convert before the call, matching what `F.scaled_dot_product_attention` does:
```python
if attn_mask is not None and attn_mask.dtype == torch.bool:
attn_mask = torch.zeros_like(attn_mask, dtype=query.dtype).masked_fill_(
attn_mask.logical_not(), float("-inf")
)
```
One note on the fill value, since it is easy to get wrong: `-inf` and
`torch.finfo(dtype).min` are **not** interchangeable in general — on a fully masked row
the CPU math path gives zeros for `-inf` and `mean(V)` for `finfo.min`. Through the
cuDNN op specifically I measured no difference between the two, so this is not a
cuDNN-visible bug; `-inf` is simply the value that matches boolean semantics on every
backend, which is what this wrapper is emulating.
I have this fix plus a parametrized regression test (partial / fully-masked row) rebased
on top of #14341, and can open it as a PR whenever a maintainer would like — happy to
wait until #14341 lands to avoid conflicting with it.
Found during an audit of cuDNN SDPA integrations across the ecosystem, run by the NVIDIA
cuDNN team.
Guida per i contributori
Apri la guida per i contributori
Direzione di ricerca
Start in src/diffusers/models/attention_dispatch.py at _cudnn_attention_forward_op, then inspect _native_flash_attention_forward_op for the same boolean-mask handling. Compare both paths with torch.nn.functional.scaled_dot_product_attention using the partial and fully masked-row cases, and add the parametrized regression test described in the issue. Done means boolean masks match the public function without changing non-boolean masks.
Scritto dal modello di indicizzazione a partire dal testo della issue.
Valutazione
- Stack tecnologico
- python, pytorch
- Ambito
- machine-learning, performance
- Tipo di issue
- Bug
- Difficoltà
- 3/5
- Tempo stimato
- 1-2 giorni
- Stato di attività
- Tranquilla
- Chiarezza
- Specificata chiaramente
- Idoneità per principianti
- 74/100