Skip to content

Flash RNN-T Loss: 2x Faster Loss and up to 50x VRAM Reduction - #16040

Open
MahmoudAshraf97 wants to merge 38 commits into
NVIDIA-NeMo:mainfrom
MahmoudAshraf97:flash-rnnt
Open

Flash RNN-T Loss: 2x Faster Loss and up to 50x VRAM Reduction#16040
MahmoudAshraf97 wants to merge 38 commits into
NVIDIA-NeMo:mainfrom
MahmoudAshraf97:flash-rnnt

Conversation

@MahmoudAshraf97

@MahmoudAshraf97 MahmoudAshraf97 commented Aug 6, 2026

Copy link
Copy Markdown
Contributor

What does this PR do ?

This is an exact implementation of RNNT loss using time-tiling and kernel fusion to reduce memory usage

flash_rnnt_packed_v4

We can safely delete all intermediate tensors since we recalculate the activations in the backward step

Results

All tests were using H100 SXM
Below are the VRAM consumption and Throughput in samples/s, the semi-transparent part of the bars is the result of tuning fused batch size between 8 to 64 in a batch of 64 across different buckets, batch sizes 1-4 were barely different in memory consumption than 8 but much slower, so they are excluded from the chart
Note that the current main branch does not support Numba BF16 and #15996 is used to enable it

mpl-throughput-bfloat16 mpl-memory-bfloat16

mpl-throughput-float32 mpl-memory-float32

The tunable memory-throughput knob is max_joint_rows, it decides how much lattice elements per tile, 200k was choosen as the default as it is the sweet spot IMO
image

The loss is almost identical and the difference can be attributed to different GEMM reduction order and different dropout implementation, the mask is drawn per call and is not what F.dropout produces for the same torch seed.

loss wer

Flash run finished 20k steps in 166 minutes and Numba finished in 207 minutes, so 20% e2e time saved, Flash also used 22% less total memory when using the same batch sizes for all buckets. The model used was a fast conformer medium with 1023+1 Vocab size, 8xH100 were used.

time and gpu util vram consumption

Usage

Add this to your rnnt training config

  model:
    joint:
      fuse_loss_wer: true
      fused_batch_size: 4  # required by the fused joint step; this loss does not read it
    loss:
      loss_name: "flash_rnnt"
      flash_rnnt_kwargs:
        fastemit_lambda: 0.0
        clamp: -1.0
        max_joint_rows: 200000  # rows of one tile

Notes for Reviewers

  1. the numpy implementation of warp loss still imports numba, this prevents some triton tests from running in any environment with missing numba although it is not needed and only needs a reference implementation.
  2. wer computation was moved into a separate util for usage in both regular fused path and flash path, note that this causes a slight difference in WER calculations since flash computes the WER for the whole batch while the other losses compute the mean of the subbatches.
  3. This is completely opt-in, the default pathes are not touched

Before your PR is "Ready for review"

Pre checks:

  • Make sure you read and followed Contributor guidelines
  • Did you write any new necessary tests?
  • Did you add or update any necessary documentation?
  • Does the PR affect components that are optional to install? (Ex: Numba, Pynini, Apex etc)
    • Reviewer: Does the PR have correct import guards for all optional libraries?

PR Type:

  • New Feature
  • Bugfix
  • Documentation

Who can review?

@artbataev @GNroy @titu1994 @hainan-xv

@copy-pr-bot

copy-pr-bot Bot commented Aug 6, 2026

Copy link
Copy Markdown

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

if TRITON_AVAILABLE:

@triton.jit
def _activate_fwd(value, activation: tl.constexpr):

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

This code path is not reachable since we raise an error in the joint if the required activation isn't implemented, so the return type is never None

return (2.0 * tl.sigmoid(2.0 * value)) - 1.0

@triton.jit
def _activate_bwd(pre_activation, activation: tl.constexpr):

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

This code path is not reachable since we raise an error in the joint if the required activation isn't implemented, so the return type is never None

@github-actions github-actions Bot added the ASR label Aug 6, 2026
@svcnvidia-nemo-ci svcnvidia-nemo-ci added the waiting-on-maintainers Waiting on maintainers to respond label Aug 8, 2026
@MahmoudAshraf97
MahmoudAshraf97 force-pushed the flash-rnnt branch 2 times, most recently from bcb238f to d3a09dd Compare August 10, 2026 09:44
parts/triton/rnnt_loss.py scans alpha and beta in a log semiring with a single
associative scan per time step, and returns the occupancies the backward pass needs
so gradients never require a second traversal.

parts/triton/rnnt_joint.py fuses the joint broadcast-add with its activation. It is a
kernel rather than a torch.compile region because Dynamo forks a graph per shape,
dtype and grad mode, and silently falls back to eager once its recompile limit is
reached, which changes reduced-precision rounding part-way through a run. Backward
saves only the two projected operands and rebuilds the pre-activation inside the
reduction, so no [B, T, U + 1, H] tensor stays live.

The log-probability extractor gains optional gradient clamping and an opt-in mode that
overwrites its private logits with their gradient. Both default off, so the existing
graph-transducer caller keeps its behaviour. Its Triton import is now guarded like the
new modules, so importing it no longer requires Triton to be installed.

Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Standard RNN-T needs log probabilities for two transitions per lattice node, but a
dense joint materializes the whole [B, T, U + 1, V] tensor to get them. Flash projects
the encoder and prediction network once, then walks the batch in chunks and each chunk
in source-time tiles, reducing every tile to its blank and target scores under
activation checkpointing. Only those [B, T, U + 1] scores survive into the backward
pass, so the vocabulary axis is bounded by max_joint_rows instead of by the batch.

Sorting by target length before chunking keeps each chunk's padding close to its own
extents rather than the batch maximum.

Gradient clamping recovers each sample's upstream scale from the global final blank
transition, which only exists in a tile spanning all of source time. Splitting source
time under clamping corrupts gradients while leaving the loss untouched, so the tiling
path refuses that combination rather than producing silently wrong gradients.

Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
RNNTJoint's fused path assumes a loss that consumes dense logits, so it builds them
per sub-batch. A loss that reports requires_factorized_joint instead receives the
projected encoder and prediction network directly and does its own chunking, so the
sub-batch loop now runs only when WER is requested.

EncDecRNNTModel rejects the combination of such a loss with fuse_loss_wer disabled,
since the dense entry point cannot serve it.

Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Equivalence is checked against a dense-logit reference built from the same kernels, so
a mismatch points at the chunking rather than at the dynamic programming: across
activations, dtypes, chunk sizes and workspace budgets, and against warprnnt_numba for
the fastemit and clamp variants where numba is the authority.

Clamping is compared to numba end to end, with the batch split across chunks and under
both loss reduction and an AMP scale, because its upstream-scale recovery is the one
part that spans the chunk boundary. The joint activation is asserted to invoke no
Dynamo frames, since a silent fall back to eager would change rounding mid-run.

Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Measures the joint and loss region end to end -- whole batch, forward and backward --
across backends, chunk sizes and length distributions.

--dtype exists because warprnnt_numba runs its dynamic programming in FP32 whatever the
joint uses, so a BF16 comparison against it measures precision as much as scheduling;
--min-length-fraction because a batch of equal-length utterances hides the sorting and
trimming entirely, which is most of what the chunking does.

Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
The module depends on torch and Triton alone, and now has two callers rather than one,
so parts/k2 is no longer the right home for it. Relocate it beside the other Triton
kernels and repoint the four import sites.

Pure relocation: file contents and behaviour are unchanged.

Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
Signed-off-by: MahmoudAshraf97 <hassouna97.ma@gmail.com>
@nithinraok

Copy link
Copy Markdown
Member

@artbataev could you help review this PR

@svcnvidia-nemo-ci svcnvidia-nemo-ci removed the waiting-on-maintainers Waiting on maintainers to respond label Aug 24, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants