[Nvidia][TritonGPU] Add phases to RemoveLayoutConversions#763
Merged
Conversation
zeroherolin
requested review from
menchunlei,
sunnycase and
zhzhcookie
as code owners
July 8, 2026 02:36
zeroherolin
force-pushed
the
zl_lyt_dev
branch
4 times, most recently
from
July 8, 2026 09:53
537e2b1 to
1343f09
Compare
Collaborator
Author
|
TLE performance is unaffected. A/B on Hopper (master switch ON vs OFF) shows the compute-bound TLE kernels (sparse-MLA TLE / FlashMLA-Prefill / Pipelined across all SKV sizes) within <0.5% (measurement noise), and the distributed cluster-GEMM within ~1% noise. This is expected: the phases either don't touch TLE anchors (dot/wgmma/TMA/pipeline) or explicitly skip TLE cluster remote-address chains, so TLE codegen is essentially unchanged. |
…ime macro Move the four RemoveLayoutConversions phases behind a runtime master switch FLAGTREE_RLC_ENHANCE (compiler.py, NVIDIA only, on by default) and wrap all enhancement code in the __FLAGTREE_RLC_ENHANCE__ compile-time macro (enabled via CMake for the default/NVIDIA build). Undefining the macro removes the enhancement and reverts the pass byte-for-byte to the upstream implementation. Revert the Passes.td pass options and the sunrise spec override back to upstream; the per-phase switches now live as file-local constexpr flags for internal debugging only.
zhzhcookie
reviewed
Jul 13, 2026
zhzhcookie
reviewed
Jul 13, 2026
zhzhcookie
reviewed
Jul 13, 2026
zhzhcookie
reviewed
Jul 13, 2026
zhzhcookie
reviewed
Jul 13, 2026
zhzhcookie
reviewed
Jul 13, 2026
zhzhcookie
reviewed
Jul 13, 2026
zhzhcookie
reviewed
Jul 13, 2026
zhzhcookie
reviewed
Jul 13, 2026
zhzhcookie
reviewed
Jul 13, 2026
Address review feedback on the RemoveLayoutConversions enhancement: - Drive the runtime master switch through the idiomatic knobs mechanism: add nvidia_knobs.rlc_enhance (FLAGTREE_RLC_ENHANCE, on by default) and read it in compiler.py, dropping the ad-hoc env-var helper. - Normalize the __FLAGTREE_RLC_ENHANCE__ / __TLE__ conditional blocks: annotate #endif/#else with their controlling macro, restore normal blank-line separation between functions, and fix include grouping. Comment/whitespace/structure only; the macro-on expansion is code-identical to the previous commit and macro-off still reverts to upstream behavior.
The conflict-resolution loop interleaved the __FLAGTREE_RLC_ENHANCE__ #ifdef with nested if/for bodies, so the shared closing braces meant different things in each branch and were hard to follow. Split the loop body into two fully self-contained #ifdef/#else branches (each with matched braces), duplicating only the two trivial common lines. No behavior change: macro-on expansion is code-identical and macro-off still matches upstream byte-for-byte in this region.
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.
Background
A
ttg.convert_layout(cvt) usually lowers to a cross-thread data reshuffle that round-trips through shared memory, making it one of the more expensive ops in TritonGPU. Yet tensor layouts are chosen independently by several upstream passes: Coalesce picks coalescing-friendly blocked layouts for memory ops, AccelerateMatmul picks MMA layouts for dots, and so on. These decisions are uncoordinated, so cvts are inserted at the layout boundaries to bridge them.The
tritongpu-remove-layout-conversions(RLC) pass, which is responsible for cleaning up redundant cvts, natively only performs forward propagation from anchors and resolves conflicts with a fixed heuristic ("prefer blocked for memory ops, mma otherwise"). This has two structural limitations:The result is that many kernels retain redundant cvts on the hot path — each a shared-memory round-trip — hurting bandwidth on memory-bound kernels and the epilogue on compute-bound ones.
Motivation: without changing RLC's existing semantics and with zero impact on other backends, add four capabilities — backward preference propagation, global cost-based selection, whole small-component solving, and MMA write-back rematerialization — to systematically remove these residual cvts.
Design
This PR non-intrusively enhances the
tritongpu-remove-layout-conversions(RLC) pass with four new optimization phases, eliminating redundantttg.convert_layout(cvt) ops produced by the independent layout decisions of upstream passes (Coalesce, AccelerateMatmul, etc.). The enhancement is enabled by default only on the Nvidia backend (transparent to users, no configuration required) and leaves other backends untouched. To turn it off at runtime, set the environment variableFLAGTREE_RLC_ENHANCE=0, which reverts to the original pass. In addition, all enhancement code is wrapped in the compile-time macro__FLAGTREE_RLC_ENHANCE__(enabled by default for Nvidia builds); undefining it in CMake removes the enhancement and reverts the pass to behavior identical to the upstream implementation, serving as an absolute fallback if issues arise. All four phases are TLE-aware: they skip TLE cluster remote-address (tle.remote_pointers) chains, leaving their layout intact.32 × byteCountand pick the minimum, replacing the hard rule of "prefer blocked for memory ops, mma otherwise"; on ties, fall back to the legacy rule to guarantee zero regression.scf.ifbranches are retagged together); otherwise the original IR is kept intact.tt.store/ result-unusedtt.atomic_rmw) whose stored value comes from a single-use cvt with an MMA-layout source, recompute the store's ptr/mask pure-index chains in the value's layout — without degrading coalescing — thereby removing the cvt from the hot data path. Includes attg.local_storecounterpart and access-preserving duplication of small-vector loads.Performance
Performance gains are observed for 100+ Ops. Representative gains (median):
replication_pad3dmm/mm_outcol2imbatch_normmax_pool2d_backwardfftmasked_selectrouter_gemmtopkrand/randn/uniform_/bernoulli_/dropout)amax/max/min/all/any/std/prodetc.)fp8_einsumsort/argsortNet cvt elimination is ~68% ~ 79% overall. A very small number of ops show minor regressions on individual shapes; all of them stem from the same source as larger gains within the same kernel (inseparable at compile time), have been individually attributed and verified, and are an acceptable trade-off. The overall net result is strongly positive.