Why was my kernel slower than an identical kernel?
The honest version of how this piece started is that I was deeply annoyed. I had a Triton kernel that was slower than a kernel that someone else had written, and the two files differed by nothing I could point at.
Same block sizes. Same num_stages. Same num_warps. I read both of them several times, in the way you do when you are convinced you are missing something obvious and you are not yet willing to consider that the thing you are missing might not be in the file.
I didn’t have a GPU that weekend. So I did the thing I have been doing for a while now with NVIDIA’s toolchain, which essentially consiosts to install the compiler and take it apart on the host, and it turned out that the answer was not in either file.
The setup works better than it has any right to. Triton’s entire compilation pipeline runs on the CPU. Source becomes TTIR, then TTIR becomes TTGIR, which becomes LLVM IR; this piece of code gets tanslated into PTX, and PTX becomes a cubin through ptxas, which is an x86 binary shipped inside the wheel.
Nothing in that chain touches a device. You need a GPU to run a Triton kernel. You do not need one to find out what the compiler decided, and what the compiler decided is most of the story.
# 197 MB from PyPI. No CUDA install, no driver, no device.
pip install triton==3.7.1 --no-deps --target ./t
# the full pipeline, on one CPU core
ck = triton.compile(src, target=GPUTarget(“cuda”, 90, 32),
options={“num_warps”: 8, “num_stages”: 3})
list(ck.asm.keys())
# [’source’, ‘ttir’, ‘ttgir’, ‘llir’, ‘ptx’, ‘cubin’]
# and, from the same wheel, the same source, a different vendor
ck = triton.compile(src, target=GPUTarget(“hip”, “gfx942”, 64), ...)
# [’source’, ‘ttir’, ‘ttgir’, ‘llir’, ‘amdgcn’, ‘hsaco’]What follows is a census of the decisions Triton makes for you, measured from Triton, on the theory that a decision you can’t see is more interesting than one you can.
It’s not an introduction, neither a benchmark. As said before, we have no hardware here and we make no claims about wall clock speed. Everything is what the compiler emits.
Three questions organise it, and we found answers to all three that we did not expect:
What decides whether your loop is pipelined? Not
num_stages. An alignment attribute the JIT infers from your pointers at launch, with a threshold at four bytes and a ramp to sixteen.What decides which matrix instruction you get? On Hopper,
num_warps. Below four warps Triton silently drops off the warpgroup path and back onto an Ampere-era instruction.Where is the compiler’s real work? Not in the arithmetic. Eighty eight percent of the passes Triton contributes, over and above what it inherits from MLIR, are about where data lives and when it moves.
On method. We wrote the verification harness before the prose, a rule we adopted after the last piece. It is 117 assertions against a live Triton install and it exits non-zero if any of them fail.
Building it caught a number we had already written down wrong, for the second time in three articles. Every figure below is tagged in the dossier, and you can reproduce some of the outputs too.
Thirty thousand lines of Python on 180 megabytes of C++
We keep seeing that Triton is described by everyone, almost universally, as a Python DSL. That is both true and wrong. It is correct for the part you write, but it’s a poor description of the whole artifact.
The Python frontend in 3.7.1 is 115 files and 29,984 lines. It sits on libtriton.so, which is 461,559,216 bytes as shipped and 180,128,520 stripped of debug information.
What is interesting to compare actually, is with the closed compiler underneath, because the usual framing has Triton as a thin open layer over an opaque assembler.
Stripped Triton is 4.35 times the size of the Blackwell ptxas it feeds. This is a huge number.
libtriton.so is 461,559,216 bytes before stripping.The ratio of 80/1 between the two backend directories will get quoted by everyone without its caveat, but in this case the caveat we are about to explore is not a measure of engineering effort.
As you may know, AMD’s code generation happens inside LLVM, which is already linked into libtriton.so; on the other side, NVIDIA’s requires shipping proprietary binaries that cannot be linked.
Thedirectories are said to be described as asymmetric, just because the licensing is asymmetric.
The narrow version of the claim survives and is still worth something: a portable compiler carries 219 MB of one vendor’s closed tooling inside its own wheel, and carries none of the other vendor’s because there is nothing closed to carry.
Which half of this stack is slow?
I assumed multiple times, before measuring, that ptxas was the expensive part, and for a good reason. It is the closed one, it does register allocation and instruction scheduling, and it is the thing people complain about when Triton compile times come up.
Spoiler alert: it is not the expensive part. Compiling the 128 by 128 by 64 matmul for sm_90a, five cold runs with the cache cleared between them, taking medians:
ptxas -O3 invocation.This matters for two concrete reasons.
The first is more practical: if you are waiting on Triton compile times, tuning
ptxasflags is not where the win is.The second is more interesting. The open part of this stack is where the complexity is, by a factor of four in time and a factor of four in binary size.
How much of “portable” survives contact with the backend?
Take just a single one matmul. Compile it for Ampere, Hopper, Blackwell datacenter, Blackwell consumer, and three AMD targets. The TTIR is byte for byte identical across all seven, at 165 lines.
Everything the language expresses is architecture neutral. Every architectural decision happens below it. What comes out the other end is not neutral in any sense.
sm_100a additionally allocates 128 columns of tensor memory, an address space the other six targets do not have.We get four observations, in ascending order of how much they should bother you.
The accumulator moves. Registers on Ampere, fed from shared memory by a warpgroup instruction on Hopper, in tensor memory on Blackwell datacenter, a fifth address space that did not exist two years ago. The programmer wrote
tl.dot(a, b, acc)in every case.Blackwell consumer does not use wgmma. It falls back to
mma.sync, the same instruction Ampere uses, and allocates no tensor memory. We reached the same conclusion from the opposite direction in Issue 09 by disassembling cubins. Here it falls out of a Triton compile:sm_120aemits zerowgmmaand zerotcgen05. The warpgroup instruction has no successor on that part. Onlymma.syncspans Hopper and both Blackwells.The AMD path is a real path, and it is target-specialised too. Same 165-line TTIR, three CDNA generations, three different TTGIRs. gfx950 issues 24
v_mfmawhere gfx942 issues 48, because the instruction is wider. And gfx942 gets a layout that no other target in the entire compiler uses:// layout kinds present in TTGIR, per target gfx90a blocked, dot_op, shared_memory, slice, swizzled_shared, amd_mfma gfx942 blocked, dot_op, shared_memory, slice, swizzled_shared, amd_mfma, linear, amd_rotating_shared gfx950 blocked, dot_op, shared_memory, slice, swizzled_shared, amd_mfma, linearThe rotating shared layout cuts LDS writes from 54 to 26 on the same kernel. It exists for one hardware generation. This is the concrete form of what portability costs: the IR is shared, the language is shared, and the layouts are per-silicon.
Register pressure is a function of pipelining, not of architecture. The figures above are from a build that pipelines. The same kernel compiled without the pipeline uses 226, 227, 244 and 234 registers on the four NVIDIA targets. On Hopper that is 227 against 128. Staging operands through shared memory buys back nearly half the register file. So pipelining is not simply a cost paid in shared memory, and the tradeoff is two-sided.
Which brings us to the question I could not answer about my own kernel: what makes the compiler pipeline the loop in the first place?
An integer you never write decides your vector width
The short answer is tt.divisibility, an attribute attached to the kernel’s pointer and integer arguments. You don’t write it. Triton’s JIT computes it at launch by looking at the runtime value of each argument and checking what it divides by.
If your tensor happens to be sixteen-byte aligned, your loads vectorize and your loop pipelines. If it does not, the same Python file compiles to a kernel that is correct, slower, and gives no indication that anything was declined.
In the first draft of this piece we wrote that the threshold was sixteen. That was wrong, or rather it was one point on a curve we had not measured. So we swept it. The structure is much better than “sixteen or nothing”, and it is exactly what the hardware would predict if you knew where to look.
The cliff is at four bytes. That is the minimum granularity of
cp.async, which moves 4, 8 or 16 bytes per thread and nothing else. Below four, the compiler cannot express the copy as an asynchronous one, so it cannot pipeline, sonum_stagesbecomes inert at every value. The threshold is not a heuristic. It is an instruction encoding.Between four and sixteen it is a ramp, not a switch. Vector width tracks alignment one for one: 4 bytes gives
[1,2], 8 gives[1,4], 16 gives[1,8]. The count ofcp.asyncinstructions falls correspondingly, 70 to 38 to 22, because the same bytes move in fewer, wider transactions.Above sixteen it saturates. A divisibility of 32 buys nothing, because 16 bytes is the widest vector load the ISA has. The compiler stops asking for more alignment than it can use.
Shared memory is identical for every pipelined width. 81,920 bytes at divisibility 4, 8, 16 and 32. Vector width changes how the bytes travel, not how many buffers exist. Those are two independent decisions the same attribute happens to gate.
cp.async has no 2-byte form. Shared memory is 81,920 bytes at every pipelined width: vector width and buffer depth are independent decisions the same attribute happens to gate.Without the attribute, num_stages is inert. It is accepted, recorded in the kernel metadata, and ignored. There is no warning.
So what does this mean for someone writing kernels? It means the shape of an allocation upstream of your kernel is a performance parameter of your kernel, and the kernel author has no lever, because the lever is not in the kernel.
A tensor produced by a slice, a view at an odd offset, or a buffer carved from a pool at an unaligned boundary takes the slow path silently. That is what had happened to me. The other person’s kernel was not better. Their tensors were.
Is this a bug? We do not think so, and it is worth being clear about why. Specialize on what you can prove, fall back to what is always safe, is the correct design for a JIT that has to be correct on every input. The complaint is narrower: the fallback is unobservable from the language.
There is no warning, no metadata field the user reads, no assert_aligned to make the assumption explicit and fail loudly when it breaks. The compiler knows it declined. It just does not say so.
What does a stage actually cost?
The folklore answer is that stages multiply your operand footprint, so shared memory goes as num_stages × (A + B). That is not what happens inthe real world. We swept five block shapes with deliberately asymmetric operands to decompose it.
Every measured value matches one closed form, across all fifteen combinations:
shared = max( num_stages × A_tile + B_tile , C_tile )
BLOCK_K = 32 the epilogue scratch dominates the loop allocation, so two stage counts produce byte-identical kernels.The marginal cost of a stage is the A tile alone. B is not multi-buffered. The TTGIR says so directly:
// 128x256x64, num_stages = 4, sm_90a
%a = ttg.local_alloc : () -> !ttg.memdesc<4x128x64xf16, #shared, #smem, mutable>
%b = ttg.local_alloc %b_reg : (...) -> !ttg.memdesc<64x256xf16, #shared, #smem>A is allocated once, outside the loop, four deep, mutable. B is allocated inside the loop from a register tensor, one deep. So A travels global to shared asynchronously and B travels global to register to shared synchronously, in a Hopper kernel where the matrix instruction reads both operands from shared memory.
We do not have a confident explanation. Both loads carry the same divisibility attributes, both index expressions have the same shape, and both operands feed the same warp_group_dot.
It is stable across shapes and stage counts, which argues against an accident of one configuration, but we are reporting it as measured rather than as a rule of the pipeliner. If you know why, we would like to hear it, and we will publish the correction.
The max term matters more than it looks. At BLOCK_K = 32 the epilogue scratch for storing C dominates the loop allocation entirely, so num_stages of 2 and 3 produce byte identical kernels: 32,768 either way.
Anyone autotuning that shape is paying compile time to distinguish two configurations that are the same kernel.
What happens below four warps?
This one we found by accident, sweeping the other option people sweep blindly, and it is the second answer that is not in the file.
warpsPerCTA and instrShape also move: [1,1] and [16,8] at one warp, [8,1] and [16,128,16] at eight, [8,2] and [16,64,16] at sixteen. Three tensor core configurations across five values of one integer.At num_warps of 1 or 2, Triton does not emit wgmma. It emits MMA version 2, which is mma.sync with an instrShape of [16, 8], the Ampere-era instruction. The reason is structural and unavoidable: wgmma is a warpgroup instruction and a warpgroup is four warps.
Below four warps there is no warpgroup, so there is no warpgroup instruction, so you are on the previous generation of tensor core path on brand new silicon.
On Hopper, num_warps is not a parallelism knob: it’s an instruction selection knob with a hard threshold at four, and nothing in the language says so.
The register column is the other half of it. Total register file consumption, which is registers per thread times threads per block, is 2,048 at one and two warps and 32,640 at four. A sixteen-fold jump for a doubling of the warp count, because crossing the threshold changes which instruction runs and therefore how much state has to be live.
And at exactly four warps the compiler lands on 255 registers per thread, which is the architectural ceiling, with zero spills. It is sitting on the edge. Eight warps halves it to 128.
Then at sixteen warps the instrShape narrows from [16, 128, 16] to [16, 64, 16] and warpsPerCTA becomes [8, 2]: the compiler splits the N dimension across two warp columns and each warpgroup gets a narrower instruction. Three different tensor core configurations across five values of one integer.
So if you autotune num_warps over [1, 2, 4, 8], which is a common default, half your search space is not testing parallelism. It is testing a different instruction. Does that matter for the result? Not necessarily, since the autotuner measures wall clock and does not care why one config is faster.
It matters for the interpretation, and it matters for anyone reasoning about the space rather than searching it exhaustively, which is everyone with a compile budget.
If it is not arithmetic, what is it?
Triton 3.7.1 registers 82 distinct passes across its core and its two backends. We classified every one by hand, into four buckets, and shipped the classification in the repository so it can be argued with line by line.
The headline number depends entirely on the classification, so it should be possible to check the classification.
m2_pipeline.py so the classification can be disputed line by line, which matters because the headline number depends on it.Strip out the borrowed machinery, which tells you nothing about Triton because every MLIR project has it, and 60 passes remain. Of those, 53 are about data placement and movement. Five are about arithmetic. Two are checks.
Triton is a layout compiler with an arithmetic language attached, not an arithmetic compiler with a layout system attached.
This is consistent with the project’s own bug data. The linear layouts paper from the Triton team reports that that 12 percent of issues filed against the Triton repository are layout related, and that the pre-linear-layout system suffered a quadratic blow-up in the number of layout-to-layout conversions that had to be implemented by hand.
The architecture specific branching goes further than lowering. The TTGIR pipeline in the NVIDIA backend is a three-way conditional on compute capability, and the branches are not the same length. The Ampere and Hopper branch makes 9 pass registration calls.
The Blackwell branch makes 14, adding accumulator initialization, tensor memory hoisting twice, promotion of the left operand into tensor memory, automatic warp specialization, partition warp optimization, and tensor memory token removal.
A “portable” compiler with a per-architecture pipeline is portable in the sense that it will produce working code everywhere, not in the sense that it does the same thing everywhere.
Eleven bits describe a two thousand element tile
The reason the placement machinery is tractable at all is a change that landed in Triton over the last two years and that we think is the most interesting idea in the project.
A layout is a map from hardware coordinates, which is to say the triple of register index, lane index and warp index, to positions in a logical tensor.
The old way to represent such a map was a family of hand-written attributes, one per pattern, plus a conversion routine for each ordered pair of patterns. That is quadratic in the number of patterns and it is where the bugs lived.
The new representation treats the hardware coordinate as a vector of bits and the layout as a linear map over the field with two elements. Because it is linear, it is a matrix. Because it is a matrix, composition is multiplication and a conversion between two layouts is one composed with the inverse of the other. One function instead of a table.
The bindings are exposed to Python, so this is not a description, it is something you can execute:
from triton._C.libtriton.linear_layout import LinearLayout
# the #blocked layout the compiler chose for the A operand in section 04:
# 8 elements per thread on the fast axis, lanes split 4 by 8, 8 warps
reg = LinearLayout.identity_1d(8, “register”, “dim1”)
lane = LinearLayout.identity_1d(8, “lane”, “dim1”)
slow = LinearLayout.identity_1d(4, “lane”, “dim0”)
warp = LinearLayout.identity_1d(8, “warp”, “dim0”)
ll = (reg * lane) * (slow * warp)
ll.is_surjective(), ll.is_injective() # True, True
sum(len(b) for _, b in ll.bases) # 11Eleven basis vectors, each one telling you where a single bit of the hardware index lands in the tensor:
The scaling is the point. A layout over 2k hardware slots needs exactly k basis vectors, which we checked from 16 elements to 65,536: 4 vectors and 16 vectors respectively. The representation is logarithmic in the size of the thing it describes, and a table would be linear. A 65,536 element tile is sixteen integers.
Conversion falls out of the same algebra. Take an identity assignment of 32 lanes and a bit-reversed assignment of the same 32 lanes, which is the shape of a swizzle, and ask for the map between them:
ident = LinearLayout.identity_1d(32, “lane”, “dim0”)
rev = LinearLayout.from_bases([(“lane”, [[16],[8],[4],[2],[1]])], [“dim0”], [32])
conv = ident.invert_and_compose(rev)
# bases: lane[0]->16, lane[1]->8, lane[2]->4, lane[3]->2, lane[4]->1No case analysis, no table entry, no new pass. This is why the placement machinery can be 53 pasasses instead of 53 passes plus a combinatorial explosion of conversion special cases, and it is the piece of Triton that we would expect to outlive the language it currently serves.
The limitation is stated openly by the authors and shows up immediately in practice: the algebra is over powers of two. Non power of two shapes have to be padded and masked, and operations like slicing and flipping are affine rather than linear, y = Ax + b rather than y = Ax, so they sit outside the framework as originally formulated.
It is worth asking how general this is, because if it is general it outlives Triton. The signs are that it is.
Work published in January 2026 gives a categorical account of CuTe layouts and observes that Triton’s F2 layouts compose naturally with swizzles, which CuTe layouts generally cannot express, while being less expressive in the other direction because of the power-of-two constraint and the inability to scale by a non-power-of-two.
A separate 2026 proposal, Axe, argues for a single unified layout abstraction across ML compilers and cites the linear layouts work as prior art. Two independent groups converging on “layout is the object worth formalising” is the same conclusion the pass census reaches by counting.
Why does one wheel need two assemblers?
The wheel ships two copies of ptxas, at CUDA 12.8.93 and CUDA 13.1.80, and picks between them with one line:
def get_ptxas(arch: int) -> knobs.NvidiaTool:
return knobs.nvidia.ptxas_blackwell if arch >= 100 else knobs.nvidia.ptxasThe reason is that the two target sets are disjoint at both ends. CUDA 12.8 still accepts ten targets that 13.1 dropped, including all of Maxwell, Pascal and Volta plus sm_101 and sm_101a.
CUDA 13.1 adds twelve that 12.8 has never heard of, including sm_88, sm_103, sm_110, sm_121 and the entire family-compatible class. Eleven targets are common to both.
Supporting the range Triton claims to support requires carrying two closed assemblers, and there is no version of CUDA that covers it.
That last row is the one worth dwelling on. Here is how Triton picks the target string it hands to ptxas:
def sm_arch_from_capability(capability: int):
# TODO: Handle non-”a” sms
suffix = “a” if capability >= 90 else “”
return f“sm_{capability}{suffix}”Unconditionally. Every kernel Triton compiles for Hopper or later goes to an a suffixed target. The a targets are the ones that expose architecture-specific instructions and, in exchange, give up PTX forward compatibility: a cubin built for sm_90a will not load on a later architecture, and the embedded PTX will not JIT forward either.
We argued in the compiler moat piece that PTX forward compatibility does not cover the instructions that matter. Here is the same claim from the other side, in four lines of Python, with a live TODO on top of it.
CUDA 13 shipped the middle option. The f targets keep forward compatibility within an architecture family while still admitting most of the family’s instructions. Triton ships the assembler that accepts them and asks for none of them.
Whether that is a deliberate choice about capability or simply work nobody has done, the effect on users is the same: your Triton kernels are locked to the exact architecture they were built for, and the lock is a hardcoded string with a comment saying somebody should look at it.
What is the compiler deciding for you, exactly?
Inside the same wheel, sharing the same compiler and the same IR, is a second language. Gluon is a lower-level DSL that hands the programmer the decisions Triton makes automatically.
Its own tutorial says the quiet part plainly: the Triton compiler generates efficient code across a wide range of kernels but can be beaten by hand-tuned low-level code, and when that happens there is little the user can do.
That gives us a way to size the automation without arguing about it. Whatever Gluon exposes and Triton does not is, by construction, a decision the Triton compiler is making on your behalf. So we counted.
triton.language exposes 122 public symbols. gluon.language exposes 147, sharing 76 with Triton and adding 71 of its own. Strip the Python plumbing and roughly forty are substantive.
Those forty are a precise inventory of the automation: choose a layout, place a tile in shared or tensor memory, pick a matrix instruction, insert a fence, decide whether to specialize warps, and check for bank conflicts.
A language that grows a second, lower-level language inside itself has told you where its ceiling is. The interesting part is that Gluon reuses the entire compiler, so the ceiling is in the inference, not in the representation.
That last distinction is what separates Gluon from the usual story of an abstraction failing. Triton’s IR can express everything Gluon can express, and Gluon compiles through the same passes to the same PTX.
What Gluon removes is the inference step: the guessing about which layout, how many buffers, which instruction. What is not always good enough is the search over it, which is a much more tractable problem and a much better place for a moat to be than in a representation.
Is the ceiling in the representation or in the search?
Triton’s ceiling is in the inference, not the representation, on the grounds that Gluon reuses the entire compiler and only removes the guessing. That is an argument from architecture.
Is there independent evidence?
There is, and it is unusually direct. In December 2025 a group from Stanford, Toronto and NVIDIA published Twill, which formulates software pipelining and warp specialization as a single joint optimization problem and solves it with an off-the-shelf constraint solver rather than with heuristics.
Their framing of the problem is worth reading against the measurements above: they describe the current state as a mix of brittle compilation heuristics and fallible human intuition, with little insight into the space of solutions, and they point at the year that elapsed between Hopper shipping and FlashAttention-3 arriving with a hand-designed schedule for it.
Two critical findings in that paper bear directly on this one.
The first concerns what Triton’s pipeliner is. The authors report that Triton’s Blackwell backend heuristically applies the FlashAttention-3 pipelining strategy. Not derives, applies. A schedule that a human designed for one kernel on one architecture is baked in as the backend’s default plan, which is a reasonable engineering decision and also exactly the thing sections 05 through 07 keep running into: the compiler is not searching, it is pattern matching against a small number of known-good plans, and whether you land on one is gated by conditions the language does not surface.
The second is harsher, and we quote its shape carefully because it is someone else’s measurement and not ours. When the Twill authors tried to have Triton compile the schedules their solver found, they report that Triton made incorrect decisions in memory allocation, layout conversion and synchronization placement, and that the result was either a compile failure or poorly performing code. Their workaround was to hand-translate their pipelined IR into CUDA C++. On Blackwell attention their solver completed in 19 seconds and found a strategy that beat Triton substantially and was competitive with cuDNN and FlashAttention-4, and the strategy it discovered was the same one the FlashAttention-4 authors had arrived at by hand.
A constraint solver rediscovered a hand-tuned schedule in nineteen seconds. The obstacle to using it was not the schedule. It was getting the compiler to lower it.
This is corroboration and it is also a caveat on our own framing. We said the ceiling is in the inference rather than the representation. Twill’s experience says the lowering has gaps too: a schedule that is expressible in principle was not compilable in practice. Those are different problems with different fixes.
The optimistic reading, which we lean toward but hold loosely, is that lowering gaps are the kind of thing that gets closed by ordinary engineering, whereas a representation that cannot express the schedule at all would be a structural problem.
Gluon existing, and shipping precisely the primitives that Twill needed to place by hand, is weak evidence for the optimistic reading.
There is a second literature pointing the same way from further out. He and Yoneki’s CuAsmRL, at CGO 2025, intercepts the cubin Triton produces, disassembles it, and has a reinforcement learning agent mutate the SASS schedule that ptxas -O3 already optimized.
They report up to 26 percent improvement and 9 percent on average, a geometric mean of 1.09 times, transparently, on kernels Triton had already compiled as well as it knows how.
Two caveats we would want if someone quoted this at us. Their evaluation used Triton 2.1.0 and ptxas 12.2 on an A100, which is several compiler generations behind everything measured in this piece, so the number should not be read as a current gap.
And the optimization happens below PTX, which is precisely where Triton has no visibility at all: it is a measurement of what the whole stack leaves on the table, not of what Triton specifically gets wrong. What survives both caveats is the direction. There was single-digit percent lying underneath an -O3 schedule, found by search.
And the automated-kernel-generation literature, TritonBench and AutoTriton among others, consistently finds generating good Triton harder than generating good CUDA.
That is a strange result for a higher-level language until you notice what is actually being generated: not a program, but a set of hints to a heuristic, several of which are the invisible switches of sections 05 and 07.
Six minutes of CPU before a single kernel runs
Because the compiler is a host program, the cost of an autotune space can be priced without a GPU.
We took a conventional matmul space, three block M by three block N by three block K by three stage counts by two warp counts, which is 162 configurations, sampled 24 of them and compiled each.
Two things follow. The first is that autotuning has a substantial fixed cost that is paid in host CPU and is invisible in any GPU-side measurement. Six minutes for one operator, on one shape, on one architecture, before measuring anything.
A serving stack with a dozen tuned operators and a per-architecture rebuild is spending real time on this, and it is time that does not show up in a kernel benchmark.
The second is that a meaningful fraction of the space is not launchable. Two of our 24 samples committed more than 232,448 bytes of shared memory, which is above what an H100 will give a single block. Those configurations compile successfully, cost their full compile time, and fail at launch.
Shared memory is known to the compiler at the end of compilation, so this is not information that has to be discovered on hardware. It could be a pre-filter. It is not.
Is 80 percent the right way to read a range?
The most useful recent number on Triton’s absolute performance comes from Yadav, Zhao and Kumar at Wisconsin-Milwaukee and Illinois Tech, who ran GEMM, fused multi-head attention and end-to-end LLM inference on an H100 NVL, a B200 and an RTX PRO 6000 Blackwell Server Edition, in BF16 and FP16.
Triton sustains 62 to 101 percent of cuBLAS across all three with no architecture-specific tuning. NVIDIA’s own CuTile reaches 52 to 79 percent of cuBLAS on GEMM in 22 lines of Python against 123 for WMMA, and on B200 its attention kernel hits 1,007 TFLOP/s, beating FlashAttention-2 by 2.5 times in 60 lines.
On the RTX PRO 6000 the same CuTile attention kernel gets 53 percent of FlashAttention-2. An earlier result from the Triton-distributed work puts Triton GEMM at roughly 95 percent of cuBLAS and CUTLASS on H800.
One caveat, which the authors raise themselves and which we would have raised anyway: their H100 ran PyTorch 2.7.1 with CUDA 12.6 while both Blackwell machines ran PyTorch 2.8.0 with CUDA 12.8. They flag it as a possible confound in cross-GPU comparison.
It does not affect the within-GPU comparisons between Triton, CuTile and cuBLAS, which is what we are using the paper for, but anyone quoting the 62 to 101 band as a clean cross-architecture result should know it is not one.
The temptation is to average that band and call Triton an eighty percent solution. We think the width of the band is the finding, not its centre. A compiler that ranges from 62 to 101 percent depending on shape and architecture is not delivering eighty percent of peak. It is delivering peak on some inputs and losing a third on others, and the variance is not something the programmer can see from the source.
Everything in the preceding sections is a mechanism for that variance, and this is the part where the compile-time census earns its keep. The divisibility hint moves the vector width by a factor of eight and gates pipelining altogether. num_warps below four takes you off the warpgroup instruction entirely.
The buffering rule means some stage counts are byte identical kernels while others triple the footprint. The max term means the epilogue can dominate a small-K shape. The architecture branch means Blackwell runs five more passes than Hopper, and one AMD generation gets a layout no other target has.
None of these are visible at the call site. A band from 62 to 101 percent is what you would expect from a compiler whose output is controlled by four or five invisible switches, some of which are set by your allocator rather than by you.
Where does the seam sit after Tile IR?
In January 2026 NVIDIA published a Triton backend that emits CUDA Tile IR instead of PTX, as an incubator repository under the triton-lang organisation, enabled with ENABLE_TILE=1, requiring CUDA 13.1 and Blackwell. Helion has a backend for it too.
Read against the rest of this piece, that is not a portability feature. It is a proposal to move the boundary. Today the boundary between Triton and NVIDIA sits at PTX, and everything in sections 04 through 07, the layouts, the buffering, the pipelining, the placement passes, happens above it, inside Triton.
Tile IR sits above PTX and its type system already encodes tile semantics. A Triton that lowers to Tile IR hands a large part of its placement work to a closed lowering compiler, which is the exact machinery this piece has spent thirty pages measuring.
We wrote in July that Tile IR was the new PTX: publish the interface, keep the lowering, one level higher than before, as a response to Triton’s position inside PyTorch. A first-party Triton backend for it, six months later, is consistent with that reading.
NVIDIA’s own documentation for the backend notes that Tile IR in CUDA 13.1 does not support num_warps, replacing it with an occupancy attribute, and that tensor-of-pointer patterns, which is the ordinary way people write Triton, perform poorly and should be rewritten to use the TMA descriptor APIs.
Both are the interface asserting itself over the language.
Where we might be wrong
The B operand asymmetry might be our kernel, not the pipeliner. We observe that only A is multi-buffered, across five shapes and three stage counts. We do not know why, and we did not test enough distinct kernels to separate a property of the pipeliner from a property of our index expressions. If the modulo we use for bounds wrapping on the N axis defeats the analysis for B specifically, the closed form in section 06 is a fact about our kernel.
Everything here is compile-time observation presented next to performance claims. We measured what the compiler emits, not what a GPU does with it. It is possible, though we think unlikely given that the unpipelined kernel has no asynchronous copies at all, that the gap is smaller in wall clock than in shared memory. We have no hardware and we are not pretending otherwise. This is the single largest weakness of the piece and the obvious thing to fix: the same harness with a device attached would settle it.
The pass classification is a judgement call and the headline moves with it. We put add_fuse_nested_loops and add_triton_licm in the borrowed bucket even though both exist in Triton partly to enable pipelining, which biases 88.3 percent downward. Moving those plus two or three similar calls pushes it past 90. Someone arguing the other way could reclassify add_accelerate_matmul and add_optimize_dot_operands as arithmetic and pull it to 85. The full list is in the repository so the argument can be had with the names visible.
The compile-time split is a lower bound on the Triton share, not a precise ratio. The outer measurement includes Python AST walking and object construction; the inner one is a bare ptxas process. A fairer accounting would instrument the MLIR pass manager. We would expect that to move the number somewhat and not to change the direction.
The 80 to 1 backend size ratio is close to meaningless and we include it only because it is measured and will be quoted anyway. It measures licensing, not effort.
The Twill results are someone else’s measurements on hardware we do not have. We are relying on them for a claim central to section 12. Two of the seven authors work at NVIDIA, which cuts both ways: they have unusually good access, and they have an interest in the conclusion that heuristic compilers leave performance on the table.
We are one version deep. All of this is Triton 3.7.1. The pipeline has been rewritten more than once in two years, the AMD backend changed substantially in 3.7, and the numbers in sections 06, 08 and 11 should be assumed to drift. Rerun the harness rather than trusting the figures.
The four-byte cliff is an inference from one instruction’s encoding. We observe the threshold and we observe that cp.async has 4, 8 and 16 byte forms. The causal claim connecting them is ours, not something we read in the compiler. It is a very short inference and we could still be wrong about the mechanism while being right about the number.
Five dated predictions
By the end of 2027, Triton will request non-
atargets for at least one architecture family. The TODO is four lines long, the shipped assembler already accepts theftargets, and the pressure comes from anyone shipping precompiled kernels. 60 percent.Alignment will become assertable in the Triton language before the inference is removed. Some annotation or type-level marker letting an author declare alignment rather than having it guessed at launch, plus a diagnostic when the pipeliner declines. 50 percent by end of 2027, and we would rather be wrong in the direction of it happening sooner.
The Tile IR backend will not be merged into mainline Triton before 2028. It is an incubator repository, it is Blackwell only, and merging it puts a closed lowering path inside the project whose institutional value is being the open one. 70 percent.
A solver-based scheduler will ship inside a mainstream kernel compiler by the end of 2027. Twill demonstrated 19-second solve times for a problem currently handled by baked-in heuristics, and the gap it exposed is the kind that gets closed. 55 percent.
Gluon’s public surface will grow faster than Triton’s over the next eighteen months. Measured in public symbols. If the ceiling is in the inference rather than the representation, the escape hatch is where the work goes. 65 percent.
The July predictions from the compiler moat piece are on the record and two resolved early, one in our favour and one against, which we scored in the GPU software gap piece. These get the same treatment.
Confidence dossier
