Skip to content

Log when thd with dropout falls to the composite cuDNN engine - #3313

Open
bzantium wants to merge 1 commit into
NVIDIA:mainfrom
bzantium:log-thd-dropout-composite-engine
Open

Log when thd with dropout falls to the composite cuDNN engine#3313
bzantium wants to merge 1 commit into
NVIDIA:mainfrom
bzantium:log-thd-dropout-composite-engine

Conversation

@bzantium

@bzantium bzantium commented Aug 4, 2026

Copy link
Copy Markdown

What does this PR do?

Adds a logger.debug line for the case where qkv_format="thd" and attention_dropout > 0, which is served by the composite cuDNN engine rather than the unified one.

Related to #3312.

Why

fused_attn_f16_arbitrary_seqlen.cu already notes that dropout and stats generation cannot be combined on the unified engine, so a thd request with dropout is routed to the composite engine, which does not support cu_seqlens:

      // This extra restriction is needed because cuDNN frontend doesn't yet allow
      // the combination of dropout and stats generation for the fprop unified engine,
      // so any such request would always get routed to the old composite SDPA engine
      // (which doesn't support cu_seqlens). Remove this restriction when possible.
      !is_dropout;

The routing is correct, but it is expensive and silent. Forward + backward through one DotProductAttention, 4096 tokens, 16 heads, head_dim 128, bf16, median of 50 iterations after 20 warmup:

GPU layout p=0.0 p=0.1 dropout cost
B300 (sm103, TE 2.14.1) thd 0.585 ms 3.060 ms +2.476 ms (5.24x)
B300 (sm103, TE 2.14.1) sbhd 0.480 ms 0.525 ms +0.045 ms (1.09x)
H200 (sm90, TE 2.15.0) sbhd 0.594 ms 0.614 ms +0.020 ms (1.03x)

A profile attributes it to cudnn::fusion::gen_dropout_mask_4bit and its transpose variant, which together take 43% of CUDA time in an 8-layer training step on B300.

Backend selection already logs every case where a backend is disabled, so a user reading NVTE_DEBUG_LEVEL=2 output sees why a backend was not chosen. This case is different: the backend is chosen and quietly costs several times more. Most configurations set attention_dropout to 0 and never see it, but Megatron-Core's TransformerConfig defaults it to 0.1, so a packed run that does not set it explicitly inherits the slow path with nothing in the log to suggest it.

This does not change behaviour — it only makes the situation visible. The underlying fix belongs to the cuDNN frontend restriction quoted above.

Testing

black (repo settings) and pylint --rcfile=pylintrc clean on the changed file. No behavioural change, so no new tests; the line appears in existing NVTE_DEBUG_LEVEL=2 output when the condition holds.

Dropout keeps a thd request off cuDNN's unified engine, and the composite
engine it lands on instead generates the dropout mask in separate kernels.
On sm103 that is 5x the no-dropout cost for the same attention, and nothing
in the backend selection log says so.

Measured at 4096 tokens, 16 heads, head_dim 128, bf16, forward+backward:
thd 0.585 ms at p=0 against 3.060 ms at p=0.1, while sbhd goes 0.480 ms to
0.525 ms for the same dropout.

Signed-off-by: Minho Ryu <ryumin93@gmail.com>
@github-actions github-actions Bot added the community-contribution PRs from external contributor outside the core maintainers, representing community-driven work. label Aug 4, 2026
@bzantium
bzantium marked this pull request as ready for review August 4, 2026 12:03
@bzantium
bzantium requested a review from cyanguwa as a code owner August 4, 2026 12:03
@greptile-apps

greptile-apps Bot commented Aug 4, 2026

Copy link
Copy Markdown
Contributor

Greptile Summary

This PR adds a debug advisory for thd FusedAttention requests with dropout, intended to expose their slower composite-cuDNN execution path.

  • Logs the composite-engine performance warning during dropout filtering.
  • Does not alter backend selection or attention behavior.

Confidence Score: 4/5

The PR is safe to merge, though the new advisory should be moved or gated on final FusedAttention selection to avoid misleading diagnostics.

Later backend filters can reject FusedAttention after the new message claims its composite cuDNN engine is in use, while final execution proceeds through another backend.

Files Needing Attention: transformer_engine/pytorch/attention/dot_product_attention/utils.py

Important Files Changed

Filename Overview
transformer_engine/pytorch/attention/dot_product_attention/utils.py Adds the intended diagnostic, but emits it before later filters establish that FusedAttention will actually be selected.

Reviews (1): Last reviewed commit: "Log when thd with dropout falls to the c..." | Re-trigger Greptile

Comment on lines +1085 to +1091
if use_fused_attention and qkv_format == "thd":
# Dropout keeps thd off cuDNN's unified engine, so it falls to the much slower
# composite one. Nothing else reports this.
logger.debug(
"FusedAttention with dropout and qkv_format = thd uses the composite cuDNN"
" engine, which is much slower than the unified engine"
)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P2 Advisory precedes backend selection

The advisory runs before later filters can reject FusedAttention, so configurations such as thd with dropout and an arbitrary mask log that the composite cuDNN engine is in use even though another backend is ultimately selected. This sends users toward the wrong performance diagnosis; emit the advisory only after final FusedAttention selection.

Knowledge Base Used: PyTorch Attention Stack

Note: If this suggestion doesn't match your team's coding style, reply to this and let me know. I'll remember it for next time!

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Remove the current early debug advisory from the dropout-filtering section. It runs before final backend selection and can therefore report FusedAttention even when a later filter selects another backend.

Suggested change
if use_fused_attention and qkv_format == "thd":
# Dropout keeps thd off cuDNN's unified engine, so it falls to the much slower
# composite one. Nothing else reports this.
logger.debug(
"FusedAttention with dropout and qkv_format = thd uses the composite cuDNN"
" engine, which is much slower than the unified engine"
)

@cyanguwa
cyanguwa requested a review from KshitijLakhani August 4, 2026 16:02
@cyanguwa

cyanguwa commented Aug 4, 2026

Copy link
Copy Markdown
Collaborator

Thanks for the detailed reporting in #3312!

I agree that this affects users' performance and without warning, but I don't think it's TE's responsibility to cache these default-value issues coming from Megatron. Given that dropout is going out of fashion, Megatron can probably change their default to 0 but in a gradual deprecation code cycle. This way, it wouldn't incorrectly and silently call FlashAttention for dropout either.

With cuDNN, we can check with them to see if they can speed up the dropout implementation on Blackwell for THD, to match FlashAttention's performance. For now, we can disable FusedAttention for THD + dropout + Blackwell, and allow users to use FlashAttention instead. @KshitijLakhani, could you please lead the discussion with cuDNN, file a bug if necessary, and guide @bzantium to disable FusedAttention in this PR? Thanks!

@KshitijLakhani

Copy link
Copy Markdown
Collaborator

Thanks for the detailed reporting in #3312!

I agree that this affects users' performance and without warning, but I don't think it's TE's responsibility to cache these default-value issues coming from Megatron. Given that dropout is going out of fashion, Megatron can probably change their default to 0 but in a gradual deprecation code cycle. This way, it wouldn't incorrectly and silently call FlashAttention for dropout either.

With cuDNN, we can check with them to see if they can speed up the dropout implementation on Blackwell for THD, to match FlashAttention's performance. For now, we can disable FusedAttention for THD + dropout + Blackwell, and allow users to use FlashAttention instead. @KshitijLakhani, could you please lead the discussion with cuDNN, file a bug if necessary, and guide @bzantium to disable FusedAttention in this PR? Thanks!

@bzantium thanks for reporting this.
I definitely am able to reproduce your observations with TE 2.17 as well, so let me take up this discussion with the cuDNN team and report back once I have a good solution.
In the mean time, let me recommend a simple, no code changes WAR.

Please set NVTE_FUSED_ATTN=0 NVTE_FLASH_ATTN=1 (hopefully this unblocks you temporarily)
I independently validated this on GB200/SM100 with TE 2.17, cuDNN 9.24, and FA2 2.7.4:

  • THD FusedAttention, dropout 0.0: 0.626 ms
  • THD FusedAttention, dropout 0.1: 4.407 ms — 7.04× slower
  • THD FA2, dropout 0.0: 0.695 ms
  • THD FA2, dropout 0.1: 0.728 ms — only 1.05× slower

For THD with dropout 0.1, switching to FA2 reduced time from 4.407 ms to approximately 0.728 ms, about a 6.1× speedup.
This agrees with your B300 results:

  • FusedAttention: 0.585 → 3.060 ms, or 5.24×
  • FA2 dropout overhead: only approximately +0.024 ms

To verify that FA2 is selected, temporarily use:
NVTE_FUSED_ATTN=0 NVTE_FLASH_ATTN=1 NVTE_UNFUSED_ATTN=0 NVTE_DEBUG=1 NVTE_DEBUG_LEVEL=2

The log should report something like this:
Selected backend = FlashAttention (2.7.4.post1)

NOTE: I am suggesting NVTE_UNFUSED_ATTN=0 only as a validation guard to prevent silent fallback.

@KshitijLakhani KshitijLakhani self-assigned this Aug 6, 2026
Comment on lines +1085 to +1091
if use_fused_attention and qkv_format == "thd":
# Dropout keeps thd off cuDNN's unified engine, so it falls to the much slower
# composite one. Nothing else reports this.
logger.debug(
"FusedAttention with dropout and qkv_format = thd uses the composite cuDNN"
" engine, which is much slower than the unified engine"
)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Remove the current early debug advisory from the dropout-filtering section. It runs before final backend selection and can therefore report FusedAttention even when a later filter selects another backend.

Suggested change
if use_fused_attention and qkv_format == "thd":
# Dropout keeps thd off cuDNN's unified engine, so it falls to the much slower
# composite one. Nothing else reports this.
logger.debug(
"FusedAttention with dropout and qkv_format = thd uses the composite cuDNN"
" engine, which is much slower than the unified engine"
)

Comment on lines 1632 to 1638
# Select FusedAttention for performance
if use_flash_attention and use_fused_attention and device_compute_capability >= (9, 0):
logger.debug(
"Disabling FlashAttention to give FusedAttention preference on Hopper+ "
"for performance reasons"
)
use_flash_attention = False

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Replace the existing Hopper+ FusedAttention preference block with the following narrow selection rule. This safely prefers FA2 only when it is confirmed installed, enabled, and eligible. Otherwise, FusedAttention remains enabled.

      # Prefer FA2 for THD training with dropout on SM100/103, where FusedAttention has a known
      # performance issue. At this point use_flash_attention_2 confirms a usable installation.
      if (
          is_training
          and qkv_format == "thd"
          and attention_dropout != 0.0
          and device_compute_capability in ((10, 0), (10, 3))
          and use_flash_attention_2
          and use_fused_attention
      ):
          logger.debug(
              "Disabling FusedAttention to give FlashAttention 2 preference for THD with dropout on SM100/103"
          )
          use_fused_attention = False
          fused_attention_backend = None
      # Select FusedAttention for performance in all other Hopper+ configurations.
      elif use_flash_attention and use_fused_attention and device_compute_capability >= (9, 0):
          logger.debug(
              "Disabling FlashAttention to give FusedAttention preference on Hopper+ "
              "for performance reasons"
          )
          use_flash_attention = False

Please place this after unavailable FlashAttention installations have been filtered and use_flash_attention_2 has been finalized. This placement is important because it ensures that FusedAttention is disabled only when FA2 is actually usable.

@KshitijLakhani

KshitijLakhani commented Aug 7, 2026

Copy link
Copy Markdown
Collaborator

@bzantium Here's a quick update: I discussed with the cuDNN team and it's a known shortcoming. The issue is not unified v/s composite but it is the slow dropout kernels used for this very specific combination being discussed here. The fact that the slow dropout kernels are being triggered in a composite setup worsens the timing. cuDNN does not have this on their roadmap for the foreseeable future as it is low priority and the suggested "fix" is for TE users to use FA2 instead if performance is of importance.

If you'd like to contribute to TE, which we highly encourage, I've reviewed your PR and added suggested changes. Please review these changes and commit them. Once you've done this I can help you launch CI on this PR and review/approve the PR to have it merged in to main.

On the contrary, if you'd rather have me add in these changes, I'm happy to do so in a separate PR. Let me know.

Notes:

  • Please make your branch up-to-date with main
  • Please change the title of the PR to something more suitable now that we know it isn't a composite engine problem per se. A suggestion is Prefer Flash Attn 2 instead of Fused Attn for THD dropout on Blackwell but feel free to use your own

@KshitijLakhani

Copy link
Copy Markdown
Collaborator

Lastly, here's a small mock up test I came up with using my agent. I just asked it to write a test to confirm the backend choice. This is juts a guidance but it does the trick. Feel free to be creative :)

  @pytest.mark.skipif(
      device_compute_capability not in ((10, 0), (10, 3)),
      reason="This backend preference applies only to SM100/SM103.",
  )
  @pytest.mark.skipif(
      not FlashAttentionUtils.is_installed,
      reason="A supported FlashAttention 2 installation is required.",
  )
  @pytest.mark.skipif(
      get_cudnn_version() < (9, 0, 0),
      reason="cuDNN 9.0.0+ is required for THD FusedAttention.",
  )
  def test_thd_dropout_prefers_flash_attention_2(monkeypatch):
      # Ensure both candidate backends are enabled for this test. Pytest restores
      # the previous environment automatically after the test completes.
      monkeypatch.setenv("NVTE_FLASH_ATTN", "1")
      monkeypatch.setenv("NVTE_FLASH_ATTN_V2", "1")
      monkeypatch.setenv("NVTE_FUSED_ATTN", "1")
      attention_params = AttentionParams(
          qkv_dtype=torch.bfloat16,
          qkv_layout="thd_thd_thd",
          batch_size=4,
          num_heads=16,
          num_gqa_groups=16,
          max_seqlen_q=1024,
          max_seqlen_kv=1024,
          head_dim_qk=128,
          head_dim_v=128,
          attn_mask_type="padding_causal",
          core_attention_bias_shape=None,
          attention_dropout=0.1,
          is_training=True,
      )
      (
          use_flash_attention,
          flash_attention_backend,
          use_fused_attention,
          fused_attention_backend,
          _,
          available_backends,
      ) = get_attention_backend(attention_params)
      assert available_backends[0]  # Real FA2 support check
      assert available_backends[1]  # Real cuDNN FusedAttention support check
      assert use_flash_attention
      assert flash_attention_backend == FlashAttentionUtils.version
      assert not use_fused_attention
      assert fused_attention_backend is None

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

community-contribution PRs from external contributor outside the core maintainers, representing community-driven work.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants