Triton Gluon: The Language That Hand-Writes the Layouts the Compiler Used to Hide

• Introduction — Where Triton Gets Stuck Isn't "Too Slow"
• What Gluon Is — Same Compiler Stack, Inverted Contract
• The First Kernel Disappoints — 666 GB/s
• Layout — What Gluon Hands Back to You
• Layouts, in Numbers
• Get the Layout Wrong and Performance Collapses
• Where It Beats PyTorch, and What It Costs
• Release Timeline — How Far It's Come
• Not NVIDIA-Only — AMD gfx950 and MI355
• Under the Layout — Linear Layouts
• The Honest Tradeoffs — When Not to Use It
• Closing
• References

Introduction — Where Triton Gets Stuck Isn't "Too Slow"

Triton's appeal is obvious. Write a tile-level kernel in Python, and the compiler handles layout, memory allocation, data movement, and synchronization for you. Most of the time this is a great trade — you get usable performance without dropping down to CUDA C++.

The problem is the moment that trade breaks down. The Gluon tutorial's introduction describes this exactly: the Triton compiler generates efficient code for a wide range of kernels, but it can lose to hand-tuned low-level code. And when that happens — in the tutorial's own words — "there is very little a user can do to meaningfully improve performance, since all the details are hidden."

That is the crux of it. Triton's limitation isn't that it's slow — it's that there is no door to walk through once you're stuck. The abstraction is a blessing when it works, and a wall the moment it doesn't. Until now, there was exactly one option in front of that wall: rewrite from scratch in CUDA C++.

Gluon is an attempt to create a third option.

What Gluon Is — Same Compiler Stack, Inverted Contract

The official docs define it in one line. Gluon is Triton's low-level GPU programming model, one that "directly exposes layouts, shared memory, warp specialization, and target-specific features so advanced kernels can trade convenience for control."

The important part is that Gluon is not a separate project. As the tutorial makes clear, Gluon and Triton run on the same compiler stack. Both implement a tile-based SPMD programming model, and both are Python DSLs sharing the same frontend and JIT infrastructure. The host-side code is nearly identical too — how you call kernels, how you specify the grid, even triton.autotune carries over unchanged.

What changes is how you write device code. To paraphrase the tutorial's own summary: Gluon changes how you write device code, while host code changes only in that the kernel picks up more hyperparameters.

The simplest possible Gluon kernel is indistinguishable from Triton.

import torch
import triton
from triton.experimental import gluon
from triton.experimental.gluon import language as gl


@gluon.jit
def copy_scalar_kernel(in_ptr, out_ptr):
    value = gl.load(in_ptr)
    gl.store(out_ptr, value)

@gluon.jit instead of @triton.jit, gl instead of tl. So far it looks like a rename.

And take note of that import path — triton.experimental. We'll come back to it at the end of this post.

The First Kernel Disappoints — 666 GB/s

The tutorial is honest enough to start with a failure. It writes a memcpy kernel that copies scalars one at a time, autotunes XBLOCK, and copies 8 GB on a GB200. The result, recorded as a comment in the tutorial source, is this.

Time:        24.00 ms
Throughput: 666.24 GB/s

The autotuner's chosen XBLOCK was 2048, and the result is about 666 GB/s. The tutorial notes this is "far below the GPU's peak bandwidth of 8 TB/s" — roughly 8% of peak.

Look at the kernel and the reason is immediately visible.

@gluon.jit
def memcpy_kernel(in_ptr, out_ptr, xnumel, XBLOCK: gl.constexpr):
    pid = gl.program_id(0)
    start = pid * XBLOCK
    end = min(start + XBLOCK, xnumel)
    for i in range(start, end):
        value = gl.load(in_ptr + i)
        gl.store(out_ptr + i, value)

Each program (a CTA on the GPU — a thread block) copies one element at a time. It uses none of the GPU's parallelism. Moving multiple elements at once means loading a tile instead of a scalar, and loading a tile means — you have to pick a layout.

This is exactly where Gluon parts ways with Triton. What the compiler would have chosen for you in Triton, you write yourself in Gluon.

Layout — What Gluon Hands Back to You

A layout specifies how a tensor's elements are distributed among the threads in a thread block. The distribution follows the GPU's hierarchy — thread block, warp, lane, and each lane's individual registers.

The most common kind is BlockedLayout.

gl.BlockedLayout(
    size_per_thread=[2, 4],
    threads_per_warp=[16, 2],
    warps_per_cta=[2, 2],
    order=[1, 0],
)

Here's how to read it.

• size_per_thread=[2, 4] — each thread owns a contiguous 2x4 sub-tile, held in that thread's registers.
• threads_per_warp=[16, 2] — thread tiles are stacked this way to form a warp tile. The product must match the number of threads per warp (32 on NVIDIA hardware).
• warps_per_cta=[2, 2] — warp tiles are stacked this way to form the whole block.
• order=[1, 0] — which dimension to tile first. This is row-major.

Multiplying the three element-wise gives the block shape — [64, 16] here.

Why order matters becomes obvious the moment you draw a single thread's tile. With order=[1, 0], registers increase along the inner dimension.

[[T:0, T:1, T:2, T:3],
 [T:4, T:5, T:6, T:7]]

With order=[0, 1], the same 2x4 tile turns into this.

[[T:0, T:2, T:4, T:6],
 [T:1, T:3, T:5, T:7]]

Same elements, same threads, different arrangement. And as we'll see, this single difference can split bandwidth by multiples.

Making the layout explicit gives you a side benefit too — the register budget becomes computable. Following the tutorial's example, using the layout above (block [64, 16]) on a 128x128 f32 tensor tiles the block into [2, 8], using 8 registers per thread per block, 8 * 16 = 128 registers total for the tensor. In Triton, that number was hidden behind the compiler.

There's a trap that comes with this too. If the tensor is smaller than the block, it gets broadcast. A 32x8 f32 tensor has only 256 elements, but since the block tiling doesn't change, each program stores it in 64 * 16 = 1024 physical registers. A small tensor eats 4x the registers — the kind of cost that's hard to notice without an explicit layout.

Layouts, in Numbers

From here on, these are measurements recorded as comments in the Triton repo's tutorial source (python/tutorials/gluon/02-layouts.py). Upstream developers measured them on a GB200 and put them directly into the source; I did not reproduce them myself.

The 1D memcpy is rewritten as a proper tile load, XBLOCK=2048 and num_warps=4 are held fixed, and only size_per_thread (R below) is varied.

gl.BlockedLayout(
    size_per_thread=[R],
    threads_per_warp=[32],
    warps_per_cta=[4],
    order=[0],
)

The recorded results.

R=1   6.574 TB/s
R=2   6.476 TB/s
R=4   6.474 TB/s
R=8   6.502 TB/s
R=16  6.214 TB/s

Two things stand out. First, this went from 0.666 TB/s in the scalar loop to 6.574 TB/s — roughly a 10x jump if you just divide the two recorded numbers (this is my arithmetic, not upstream's own phrasing). That is the scale of what a tile load plus a layout buys you.

Second, making R bigger does not help. R=1 is fastest and R=16 is slowest. From here the tutorial opens the SASS and counts LDG.E/STG.E instructions directly. This says a lot about what Gluon actually is — exposing the layout doesn't mean giving you "better defaults," it means giving you a choice, and the responsibility to verify that choice. The intuition that widening it makes it faster is simply wrong here.

Get the Layout Wrong and Performance Collapses

Move to 2D and the story sharpens.

An 8 GB, 32K x 64K tensor is copied with XBLOCK=1, YBLOCK=2048, using a layout designed to behave like the R=1 winner from 1D.

layout = gl.BlockedLayout([1, 1], [1, 32], [1, 4], [1, 0])

The result is 6.260 TB/s — 5% slower than 1D. The tutorial chalks this up to 2D arithmetic being more complex, and calls it acceptable.

But what happens if you transpose the input tensor? Same kernel, same layout.

0.774 TB/s

In the tutorial's own words, performance "craters." The reason is simple — the inner dimension is no longer contiguous, so coalescing disappears. 32 lanes need to read adjacent addresses to be bundled into one transaction, but on the transposed tensor, each lane points to a far-apart address.

The fix is just as simple. Change the block size and transpose the layout.

layout = gl.BlockedLayout([1, 1], [32, 1], [4, 1], [0, 1])
# XBLOCK=2048, YBLOCK=1

6.590 TB/s

From 0.774 to 6.590. The kernel's algorithm didn't change by a single character — only the layout descriptor changed.

There's an interesting observation the tutorial adds here. The fixed version (6.590 TB/s) is slightly faster than the 1D memcpy (6.574 TB/s). Even though the per-program memory access pattern is identical in both the transposed and non-transposed cases, there's still a gap, and the tutorial attributes it to where the programs get scheduled on the GPU — data locality effects like TLB virtual-address-translation caching, or the fact that on an H100 the L2 cache is split into partitions that talk to each other. At this level, even "the same access pattern" doesn't mean the same performance.

The practical lesson here: the optimal layout depends not on the kernel but on the stride of the tensor as it sits in global memory. Which is why the tutorial eventually converges on a helper like this.

def get_layout_for_gmem_access(tensor, num_warps):
    if len(tensor.shape) == 1:
        return gl.BlockedLayout([1], [32], [num_warps], [0])

    assert len(tensor.shape) == 2, "only 1D and 2D tensors are supported"
    assert 1 in tensor.stride(), "expected at least 1 contiguous dimension"
    if tensor.stride(1) == 1:
        return gl.BlockedLayout([1, 1], [1, 32], [1, num_warps], [1, 0])
    else:
        return gl.BlockedLayout([1, 1], [32, 1], [num_warps, 1], [0, 1])

Instead of baking the layout in as a constant, it's derived from the runtime stride. This is the kind of code Gluon kernels need if they're going to be flexible and fast at the same time.

Where It Beats PyTorch, and What It Costs

The tutorial also covers a case where Gluon's 2D memcpy actually earns its keep. Copying a row-skipping view (a non-contiguous tensor) into a contiguous tensor for an 8 GB tensor — the same job PyTorch's x.contiguous() does.

2D memcpy:                6.258 TB/s
torch.Tensor.contiguous:  2.946 TB/s
2D memcpy (transposed):   6.398 TB/s

The tutorial's own wording is "already more than 2x faster than the PyTorch implementation." This is a real example of Gluon earning its value — when the input or output has an exotic layout. Though there's a condition attached: the 1D memcpy only works when both input and output are views of contiguous blocks of memory.

What if the input and output layouts are opposite each other? The tutorial measures this case too. The input is contiguous along dim 1, and the output is contiguous along dim 0, on a 32K x 32K tensor.

2D memcpy (order=[1, 0]): 0.978 TB/s
2D memcpy (order=[0, 1]): 1.674 TB/s

Either layout choice is terrible — satisfying one breaks the other. The fix is to use different layouts for the load and the store, and convert between them.

@gluon.jit
def memcpy_2d_inout_kernel(in_ptr, out_ptr, ..., layout_in: gl.constexpr, layout_out: gl.constexpr, ...):
    ...
    value = gl.load(in_ptr + in_offsets, mask=mask_in)

    # Use `gl.convert_layout` to perform layout conversions.
    value = gl.convert_layout(value, layout_out)

    gl.store(out_ptr + out_offsets, value, mask=mask_out)

2D memcpy (in/out layouts): 4.814 TB/s

From 0.978 to 4.814 — but still short of the 6.5 TB/s range. And the tutorial doesn't hide why. convert_layout needs shared memory, and shared memory is a scarce resource on the GPU. Using it for a layout conversion cuts occupancy and reduces the maximum pipeline depth, which hurts performance. The tutorial's conclusion is that "in this case, the conversion cost is unavoidable, and it's much cheaper than the cost of inefficient global memory accesses."

In other words, 4.814 TB/s isn't a win — it's the least-bad choice. Making costs visible like this is Gluon's whole character.

Release Timeline — How Far It's Come

You can track how far along Gluon is from the release notes. The dates below are confirmed from GitHub's release API.

There's a legible arc here. The skeleton landed in mid-2025, and through 2026 it's been expanding into multi-CTA and the newest hardware features (Blackwell's tcgen05, cluster barriers). And the fact that v3.7.1 was a pure bug-fix release also means the Gluon surface has been quiet for about a month as of the latest stable release.

The 3.6.0 Triton → Gluon converter deserves a separate mention. It sounds like a tool that makes migration easy, but the description on the actual PR (#8417, merged 2025-10-11) pins expectations down precisely: "This is not intended for production use, and simply allows converting Triton kernels into a naive Gluon version, which will likely be quite a bit slower."

That's an honest sentence. And it summarizes what Gluon fundamentally is — moving to Gluon by itself buys you no performance at all. Performance comes afterward, when you get the layout right yourself.

Not NVIDIA-Only — AMD gfx950 and MI355

One thing that stands out in the table above is how often AMD entries show up. async_copy for gfx1250, AMD warp pipelining, Stream-K kernels for gfx1250, AMD cluster barriers. Gluon was not designed as an NVIDIA-only escape hatch.

AMD is pushing this too. On May 22, 2026, the ROCm blog published "From Naive to Near-Peak: Building High-Performance GEMM Kernels with Gluon" by Lixun Zhang, Jason Furmanek, Peng Sun, and Emad Barsoum. It's a tutorial that optimizes a GEMM kernel step by step from v0 to v9 using Gluon, on gfx950 (CDNA4)-based MI350/MI355 hardware.

Every number below is AMD's own vendor-measured figure. The blog states its measurement conditions as a single MI355, ROCm 7.0, and Triton built from the gfx950-tutorial-v0.1 tag.

There's a point that's easy to misread here, so let me spell it out. The roughly 3x from 520 → 1489 TFLOPS does not mean Gluon beats Triton by 3x. Both v0 and v9 are Gluon kernels — they're just the two ends of an educational optimization journey in the same repo (ROCm/gfx950-gluon-tutorials), from v0_naive to v9_beyond_hotloop. A naive Gluon kernel being slow is expected (exactly what the converter PR said above), and the 3x is a story about "this is how far you get if you do the layout and scheduling right" — not a cross-language comparison.

And the blog draws its own line — it explicitly states this is not a replacement for a production library like hipBLASLt. The body of the post, as far as I could confirm, contains no direct comparison figures against hipBLASLt. So the claim "a Gluon kernel beat a vendor library" does not come from this material.

There's also an honest methodological choice in this blog. It uses MFMA efficiency, not TFLOPS, as its primary metric, citing that MFMA efficiency is clock-independent and more reproducible across runs than raw TFLOPS. That's a step above vendor material that just waves around TFLOPS numbers that swing with boost clocks. (Though of course, 99% MFMA efficiency doesn't by itself mean the workload is fast — it means the matrix engine sits idle for almost none of the hot loop, and that's only meaningful if the kernel is running the right algorithm to begin with.)

Under the Layout — Linear Layouts

Go one level deeper and you find linear layouts. As the tutorial explains, Gluon has no canonical layout representation — multiple layouts can describe the same element mapping. For example, the following two are equivalent.

gl.BlockedLayout([1], [32], [4], [0])
gl.SliceLayout(1, gl.BlockedLayout([1, 1], [32, 1], [4, 1], [1, 0]))

If you know they're equivalent, you can have the compiler verify it with gl.convert_layout(x, layout, assert_trivial=True) — checking that the conversion really is free (nothing more than a register rearrangement). It's an API that fits an explicit language.

And every Gluon layout can be expressed as a linear layout. The tutorial calls this "the most expressive and powerful representation" while also noting it's "relatively rare and can be hard to understand." The underlying theory is published as a paper — "Linear Layouts: Robust Code Generation of Efficient Tensor Computation Using F2" by Keren Zhou, Mario Lezcano, Jeff Niu, Phil Tillet, Thomas Raoux, and others (arXiv:2505.23819, submitted May 28, 2025, revised March 6, 2026). The approach models layouts as binary matrices over GF(2), avoiding the quadratic complexity of prior systems.

For a practitioner, the takeaway is this — BlockedLayout and SliceLayout cover most cases, and you don't need to know linear layouts until you do. But the moment you need zero-cost split/join/reshape/permute, that's where it lives.

The Honest Tradeoffs — When Not to Use It

Time to be honest.

It is still experimental. The import path is triton.experimental.gluon, and in the repo it lives at python/triton/experimental/gluon. The official Gluon overview docs carry no separate warning about experimental status or stability, but the path itself already says API stability is not promised. The fact that the API kept getting added to and changed across every release from 3.4.0 through 3.7.0 is consistent with that reading. If you put a production kernel here, be ready for it to break every time you bump the Triton version.

It demands a different amount of knowledge. The tutorial's introduction is direct about it — writing Gluon kernels requires a deeper understanding of several aspects of GPU hardware and GPU programming. If you don't know that a warp has 32 threads, why coalescing depends on the inner dimension, what a shared-memory bank conflict is, or how register pressure affects occupancy, you cannot fill in BlockedLayout's four arguments. Triton hiding all this wasn't laziness — it was a service.

Moving to Gluon by itself is not free performance. As the converter PR said, a naively ported Gluon kernel is significantly slower. And as the R sweep above showed, the intuitive move (widening it) can actually make things worse. If you're taking over work the automated compiler used to do, you need to prove, with a benchmark every time, that you're doing it better.

The whole tutorial is memcpy. The GB200 numbers quoted in this post are values upstream developers measured and recorded in the source — I did not reproduce them myself. And the target is memory-bandwidth-bound memcpy specifically — if your kernel is compute-bound, you won't see a 10x from a layout sweep.

To sum it up, here's the decision criteria.

When it earns its cost

• You've already written it in Triton, and profiling has confirmed with measurements that the compiler-generated code falls well short of the hardware limit.
• The bottleneck is in layout, data movement, or synchronization, and there's no lever left to pull at the Triton level.
• You need to reach directly for a recent hardware feature (Blackwell tcgen05, Hopper WGMMA, 2CTA, TMA, CDNA4's MXFP4).
• The kernel is long-lived and runs often enough to pay back the tuning time and the per-version maintenance cost.

When it's overkill

• You haven't profiled it in Triton yet. (This is most cases.)
• What torch.compile produces is already good enough.
• The kernel is compute-bound and layout isn't the bottleneck.
• A vendor library (cuBLAS, hipBLASLt, cuDNN) already covers the operation — even the AMD blog is explicit that its own tutorial is not a replacement for hipBLASLt.
• Nobody on the team knows GPU microarchitecture.

Closing

Gluon is not "a better Triton." It's a different contract.

Triton's contract is "I'll handle the details, you just write the algorithm." That's a good trade most of the time, and most kernels should keep living in Triton going forward. Gluon's contract is "you handle the details, but in exchange there's a door instead of a wall when you get stuck."

That spectrum matters. There used to be a cliff between Triton and CUDA C++ — when the compiler couldn't produce what you needed, you had to switch the entire language and ecosystem. Gluon puts stairs on that cliff. You stay on the same compiler stack, the same frontend, the same JIT, the same autotuner, and only step down into the parts you actually need. And the fact that those stairs are being built on both NVIDIA and AMD — with AMD writing its own gfx950 tutorial, and an AMD entry showing up in nearly every release note — might be the more important piece of news here.

But keep the order straight. Write it in Triton, profile it, confirm the compiler actually lost, and only then reach for Gluon. The fact that what split 0.774 TB/s from 6.590 TB/s wasn't the algorithm but a single layout descriptor line cuts both ways — get it right and you gain 10x, get it wrong and you lose 10x. Triton let you skip that choice. Gluon hands it back to you. Whether that's a gift or a bill depends on whether you'd actually opened the profiler.

References

• Gluon Overview — official Triton docs
• Introduction to Gluon — tutorial part 1
• 01-intro.py — original tutorial source (includes GB200 memcpy measurements)
• 02-layouts.py — original layout tutorial source (includes TB/s measurements)
• triton-lang/triton release notes — 3.4.0 through 3.7.1
• PR #8417 — experimental Triton to Gluon converter
• From Naive to Near-Peak: Building High-Performance GEMM Kernels with Gluon — ROCm blog (AMD vendor-measured, 2026-05-22)
• ROCm/gfx950-gluon-tutorials — kernel source for the AMD blog
• Linear Layouts: Robust Code Generation of Efficient Tensor Computation Using F2 (arXiv:2505.23819)