huggingface / huggingface/diffusers
cuDNN attention CP path consumes a bool attn_mask as an additive 0/1 bias
- Lenguaje dominante
- Python
- Estrellas
- 34.5k
- Forks
- 7.3k
- Merge medio
- 3 d 3 h
- PR fusionados (30 d)
- 91
Descripción
# 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.
Guía de contribución
Línea de trabajo
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.
Escrito por el modelo de indexación a partir del texto del issue.
Evaluación
- Stack tecnológico
- python, pytorch
- Área
- machine-learning, performance
- Tipo de issue
- Error
- Dificultad
- 3/5
- Tiempo estimado
- 1-2 días
- Estado de actividad
- Tranquilo
- Claridad
- Bien especificado
- Aptitud para principiantes
- 74/100