Flash RNN-T Loss: 2x Faster Loss and up to 50x VRAM Reduction - #16040
Open
MahmoudAshraf97 wants to merge 38 commits into
Open
Flash RNN-T Loss: 2x Faster Loss and up to 50x VRAM Reduction#16040MahmoudAshraf97 wants to merge 38 commits into
MahmoudAshraf97 wants to merge 38 commits into
Conversation
MahmoudAshraf97
force-pushed
the
flash-rnnt
branch
from
August 6, 2026 17:24
f66c4ff to
431c114
Compare
| if TRITON_AVAILABLE: | ||
|
|
||
| @triton.jit | ||
| def _activate_fwd(value, activation: tl.constexpr): |
Contributor
Author
There was a problem hiding this comment.
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): |
Contributor
Author
There was a problem hiding this comment.
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
MahmoudAshraf97
force-pushed
the
flash-rnnt
branch
2 times, most recently
from
August 10, 2026 09:44
bcb238f to
d3a09dd
Compare
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>
MahmoudAshraf97
force-pushed
the
flash-rnnt
branch
from
August 22, 2026 08:55
d3a09dd to
f9be9a3
Compare
Member
|
@artbataev could you help review this PR |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What does this PR do ?
This is an exact implementation of RNNT loss using time-tiling and kernel fusion to reduce memory usage
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
mainbranch does not support Numba BF16 and #15996 is used to enable itThe 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 IMOThe 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.dropoutproduces for the same torch seed.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.
Usage
Add this to your rnnt training config
Notes for Reviewers
numba, this prevents some triton tests from running in any environment with missingnumbaalthough it is not needed and only needs a reference implementation.Before your PR is "Ready for review"
Pre checks:
PR Type:
Who can review?
@artbataev @GNroy @titu1994 @hainan-xv