Triton GPU Kernel Programming: Tiling, Fusion and Choosing CUDA, CUTLASS or CuTe

Most GPU kernels you will ever need to write are not compute-bound. They are bandwidth-bound: a chain of small operations that each drag a tensor out of high-bandwidth memory (HBM) and push it back. Triton GPU kernel programming attacks that problem directly. You write one Python function that describes what a block of data should do, and the compiler handles thread mapping, shared-memory staging, vectorized access and software pipelining for you.

That trade matters in 2026 because the practical GPU toolchain has fractured into layers. PyTorch’s compiler emits Triton by default for fused pointwise and reduction code. NVIDIA now ships a Python DSL (CuTe DSL) inside CUTLASS 4. Newer entrants such as Helion, TileLang and Pallas all argue that tile-level abstractions beat thread-level ones. Knowing where each one fits saves weeks of effort.

You will leave this post able to write a masked vector add, a fused softmax and an autotuned tiled matmul in Triton, read the compiler pipeline, profile what you wrote, and decide when to drop down to CUDA, CUTLASS or CuTe DSL.

What this covers: the Triton programming model, three runnable kernels, the compiler pipeline, profiling, pitfalls, and a decision framework against CUDA, CUTLASS, CuTe DSL, Pallas, TileLang and Helion.

Context and Background: Why Tile-Level Programming Won

A modern GPU is a throughput machine with a steep memory hierarchy. Registers are fastest, shared memory (SRAM) next, then L2 cache, then HBM. Arithmetic is cheap relative to data movement, so a kernel’s speed usually depends on how few times it touches HBM. If you want the hardware picture first, our explainer on the GPU thread, warp and streaming-multiprocessor hierarchy covers the execution model, and DRAM versus SRAM and kernel fusion covers why fusion pays.

Classic CUDA exposes this hierarchy at thread granularity. You decide what each thread does, how threads cooperate through shared memory, where to place __syncthreads(), how to avoid bank conflicts, and how to coalesce global loads. That control is powerful, and it is also why a competitive CUDA matmul takes expert effort. Our overview of CUDA kernel execution and data starvation shows how much bookkeeping hides behind a “simple” kernel.

Triton, which grew out of a 2019 research paper on tile-based neural network compilation (the project README cites the 2019 MAPL publication), moves the unit of programming up one level. According to the project’s own README, the goal is “faster code than CUDA at higher productivity,” with more flexibility than other DSLs. You reason about blocks of elements, and the compiler maps blocks onto threads, warps and memory spaces. The official Triton repository lists Linux, NVIDIA GPUs with compute capability 8.0 or higher, and AMD GPUs with ROCm 6.2 or higher as supported targets, with CPU support still under development. Those requirements come from the README as retrieved while writing this post; check the repository for the current matrix before you plan a deployment.

Two things make Triton strategically important beyond its own users. First, PyTorch’s compiler stack leans on it: TorchInductor lowers fused graphs of pointwise, reduction and many matmul-adjacent operations to Triton code on GPUs, so many engineers run Triton kernels without ever writing one. Second, it is the substrate for higher-level layers. PyTorch’s Helion DSL, for example, is documented as compiling down to Triton rather than replacing it.

The honest framing: Triton is not a universal CUDA replacement. It is the right tool when your kernel is naturally expressed as operations on rectangular tiles and the compiler’s heuristics are good enough. When they are not, you have to know what to reach for next, and that is the second half of this post.

The Triton Programming Model: Program Instances, Tiles and Masks

In one paragraph: a Triton kernel is a Python function decorated with @triton.jit. Launching it creates a grid of independent program instances, each identified by tl.program_id. Each instance operates on a block of data expressed as tensors of offsets, loads that block with masks, computes on it in registers, and stores the result. You never write per-thread code.

Program instances instead of threads

In CUDA a kernel body describes one thread, and you manually derive a global index from blockIdx, blockDim and threadIdx. In Triton the body describes one program instance, and the compiler decides how many threads and warps cooperate to execute it. The num_warps launch parameter is a hint about parallelism inside an instance, not a thread index you can read.

The official vector-add tutorial shows the whole model in a handful of lines:

import torch
import triton
import triton.language as tl

@triton.jit
def add_kernel(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(axis=0)
    block_start = pid * BLOCK_SIZE
    offsets = block_start + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements
    x = tl.load(x_ptr + offsets, mask=mask)
    y = tl.load(y_ptr + offsets, mask=mask)
    tl.store(output_ptr + offsets, x + y, mask=mask)

def add(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
    output = torch.empty_like(x)
    n_elements = output.numel()
    grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']),)
    add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=1024)
    return output

Read it as a block-level program. tl.arange(0, BLOCK_SIZE) creates a vector of indices, and every subsequent operation applies to that whole vector. x_ptr + offsets is pointer arithmetic over a block, so tl.load fetches BLOCK_SIZE elements in one logical operation. Triton picks the vector width and thread mapping that make those loads coalesced.

Figure 1: The Triton execution model. A launch grid creates program instances; each computes offsets from its program id, masks out-of-range lanes, loads a block, computes in registers and stores.

Figure 1 summarizes the flow. The grid is launched with kernel[grid](...), where grid is either a tuple or a function of the kernel’s meta-parameters. Ceiling division (triton.cdiv) ensures the last instance covers the tail of the array.

Masks are not optional

Because the last block may extend past the end of the data, every load and store that can run out of bounds must carry a mask. The tutorial describes the mask as there to guard memory operations against out-of-bounds accesses. A masked-off load never touches memory, and other= supplies the value that lane sees instead.

That other value is a correctness decision, not a formality. In a sum, padding with 0.0 is neutral. In a max, -inf is neutral. In a product, 1.0 would be. A common bug is to mask the load but leave the default, so padded lanes silently contribute garbage to a reduction. We return to this in the pitfalls section.

constexpr and compile-time specialization

Parameters annotated tl.constexpr, such as BLOCK_SIZE, must be known at compile time. Triton compiles a separate specialization of the kernel for each distinct value. This is what lets tl.arange(0, BLOCK_SIZE) produce a statically shaped tensor, and it is also why block sizes are normally powers of two. The cost is that every new constexpr value triggers a fresh compile, which matters for autotuning and for dynamic shapes.

Block pointers and descriptors

Raw pointer arithmetic is flexible but leaves the compiler to infer access patterns. Triton also offers structured ways to describe tiles of multi-dimensional tensors: block pointers (tl.make_block_ptr) and, on newer NVIDIA hardware, tensor descriptors that can map to the Tensor Memory Accelerator (TMA). I did not re-verify the current stability or exact API names of these features for this post, so treat the names as pointers to the documentation and check the Triton release you target. Helion’s documentation does list pointer, block_ptr and tensor_descriptor as indexing strategies it can tune, and notes that tensor_descriptor uses TMAs on Hopper and Blackwell GPUs, which confirms that all three paths exist in the Triton ecosystem.

Fused Kernels: From Eager Softmax to One Pass Over Memory

In one paragraph: fused kernels combine several operations into one GPU kernel so intermediate tensors stay in registers or shared memory instead of round-tripping through HBM. Softmax is the canonical example: eager PyTorch executes it as several memory-bound steps, while a fused Triton kernel reads each row once and writes it once.

Counting the traffic

Take an M x N matrix and apply row-wise softmax. The official fused-softmax tutorial counts the naive PyTorch-style implementation as reading roughly 5MN + 2M elements and writing roughly 3MN + 2M, because each intermediate (row max, shifted values, exponentials, row sum) is materialized. The fused kernel reads MN and writes MN. The tutorial estimates a theoretical speedup of about 4x from the ratio (8MN + 4M) / 2MN. That is a ceiling set by bytes moved, not a measured result, and real speedups depend on hardware, dtype and shape.

Figure 3: Eager softmax versus a fused Triton kernel. The fused version keeps max, exp, sum and divide in registers between a single read and a single write.

The insight generalizes. Any chain of elementwise and reduction operations over data that fits in on-chip memory is a fusion candidate, and the saving scales with how many intermediates you eliminate. Kernel launch overhead is a secondary gain.

The fused softmax kernel

This is the structure of the official tutorial’s kernel, with the host-side occupancy heuristics trimmed. The kernel is persistent: a fixed number of programs loop over rows with a stride equal to the number of programs.

import torch
import triton
import triton.language as tl

@triton.jit
def softmax_kernel(output_ptr, input_ptr, input_row_stride, output_row_stride,
                   n_rows, n_cols, BLOCK_SIZE: tl.constexpr,
                   num_stages: tl.constexpr):
    row_start = tl.program_id(0)
    row_step = tl.num_programs(0)
    for row_idx in tl.range(row_start, n_rows, row_step, num_stages=num_stages):
        row_start_ptr = input_ptr + row_idx * input_row_stride
        col_offsets = tl.arange(0, BLOCK_SIZE)
        input_ptrs = row_start_ptr + col_offsets
        mask = col_offsets < n_cols
        row = tl.load(input_ptrs, mask=mask, other=-float('inf'))
        row_minus_max = row - tl.max(row, axis=0)
        numerator = tl.exp(row_minus_max)
        denominator = tl.sum(numerator, axis=0)
        softmax_output = numerator / denominator
        output_ptrs = output_ptr + row_idx * output_row_stride + col_offsets
        tl.store(output_ptrs, softmax_output, mask=mask)

def softmax(x: torch.Tensor) -> torch.Tensor:
    n_rows, n_cols = x.shape
    BLOCK_SIZE = triton.next_power_of_2(n_cols)
    y = torch.empty_like(x)
    # The tutorial sizes the grid from measured register and shared-memory
    # usage. A simpler, slightly less tuned launch is one program per row:
    softmax_kernel[(n_rows,)](y, x, x.stride(0), y.stride(0),
                              n_rows, n_cols, BLOCK_SIZE=BLOCK_SIZE,
                              num_stages=2, num_warps=8)
    return y

if __name__ == "__main__":
    x = torch.randn(4096, 1000, device="cuda")
    torch.testing.assert_close(softmax(x), torch.softmax(x, dim=1))

Check the pieces against what we said earlier. BLOCK_SIZE is the next power of two at or above n_cols, so one instance holds a whole row. The mask handles the padding, and other=-float('inf') makes padded lanes irrelevant to the max and, because exp(-inf) is zero, to the sum. Subtracting the row max is the standard numerical-stability trick: it keeps exp arguments at or below zero so they cannot overflow.

The tutorial itself notes that exponentiation in Triton is fast but approximate, comparable to __expf in CUDA. The final assert_close uses PyTorch’s default tolerances; tighten or loosen them deliberately rather than by accident.

One caveat on the simplified launch above: it is my reduction of the tutorial’s code, not the tutorial’s code verbatim, and I have not benchmarked it. The tutorial’s persistent launch exists precisely to keep every streaming multiprocessor busy without launching one instance per row.

Why persistent programs

The tutorial computes occupancy on the host by pre-compiling the kernel (warmup), reading the register count and shared-memory footprint, and then dividing the device’s register file and shared memory by those numbers. It then launches NUM_SM * occupancy programs, capped at the number of rows. Each program walks rows with tl.range(..., num_stages=...), which asks the compiler to software-pipeline the loop so the next row’s load overlaps the current row’s math.

This design pays off when rows are short relative to the number of SMs, because launching a CUDA-style program per row adds scheduling overhead. It also illustrates a theme: Triton gives you the high-level structure, but performance work still means understanding registers, shared memory and occupancy.

Fusion beyond softmax

The same pattern covers layer normalization and RMSNorm (reduce, then scale), bias-plus-activation epilogues, rotary position embeddings, dequantize-then-multiply, and the online-softmax recurrence inside attention. The last is worth its own treatment; see our deep dive on FlashAttention’s tiling and online softmax for the algorithm that turns one softmax into a streaming computation over tiles.

Tiled Matmul with Autotune: Where Triton Earns Its Keep

In one paragraph: a tiled matrix multiply assigns each program instance one BLOCK_M x BLOCK_N tile of the output, loops over the K dimension in BLOCK_K slices, accumulates in float32 with tl.dot (which maps to tensor cores), and writes the tile back. Autotuning searches block sizes, warps and pipeline stages for the best configuration per problem shape.

The kernel

The kernel below follows the official matrix-multiplication tutorial. It adds two ideas on top of the plain tiling loop: grouped program ordering for L2 cache reuse, and a config list for @triton.autotune.

import torch
import triton
import triton.language as tl

def get_autotune_config():
    return [
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 64,
                       'GROUP_SIZE_M': 8}, num_stages=3, num_warps=8),
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32,
                       'GROUP_SIZE_M': 8}, num_stages=4, num_warps=4),
        triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32,
                       'GROUP_SIZE_M': 8}, num_stages=4, num_warps=4),
        triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 32, 'BLOCK_SIZE_K': 32,
                       'GROUP_SIZE_M': 8}, num_stages=5, num_warps=2),
    ]

@triton.autotune(configs=get_autotune_config(), key=['M', 'N', 'K'])
@triton.jit
def matmul_kernel(a_ptr, b_ptr, c_ptr, M, N, K,
                  stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn,
                  BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr,
                  BLOCK_SIZE_K: tl.constexpr, GROUP_SIZE_M: tl.constexpr):
    pid = tl.program_id(axis=0)
    num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)
    num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)
    num_pid_in_group = GROUP_SIZE_M * num_pid_n
    group_id = pid // num_pid_in_group
    first_pid_m = group_id * GROUP_SIZE_M
    group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)
    pid_m = first_pid_m + ((pid % num_pid_in_group) % group_size_m)
    pid_n = (pid % num_pid_in_group) // group_size_m

    offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M
    offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
    offs_k = tl.arange(0, BLOCK_SIZE_K)
    a_ptrs = a_ptr + (offs_am[:, None] * stride_am + offs_k[None, :] * stride_ak)
    b_ptrs = b_ptr + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)

    accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
    for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):
        a = tl.load(a_ptrs, mask=offs_k[None, :] < K - k * BLOCK_SIZE_K, other=0.0)
        b = tl.load(b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_SIZE_K, other=0.0)
        accumulator = tl.dot(a, b, accumulator)
        a_ptrs += BLOCK_SIZE_K * stride_ak
        b_ptrs += BLOCK_SIZE_K * stride_bk
    c = accumulator.to(tl.float16)

    offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
    offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
    c_ptrs = c_ptr + stride_cm * offs_cm[:, None] + stride_cn * offs_cn[None, :]
    c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)
    tl.store(c_ptrs, c, mask=c_mask)

def matmul(a: torch.Tensor, b: torch.Tensor) -> torch.Tensor:
    M, K = a.shape
    K, N = b.shape
    c = torch.empty((M, N), device=a.device, dtype=torch.float16)
    grid = lambda META: (triton.cdiv(M, META['BLOCK_SIZE_M']) *
                         triton.cdiv(N, META['BLOCK_SIZE_N']),)
    matmul_kernel[grid](a, b, c, M, N, K,
                        a.stride(0), a.stride(1), b.stride(0), b.stride(1),
                        c.stride(0), c.stride(1))
    return c

if __name__ == "__main__":
    a = torch.randn((512, 512), device="cuda", dtype=torch.float16)
    b = torch.randn((512, 512), device="cuda", dtype=torch.float16)
    torch.testing.assert_close(matmul(a, b), torch.matmul(a, b), atol=1e-2, rtol=0)

The config list is a four-entry subset of the tutorial’s much longer CUDA list, so it will not be as good across shapes. The tutorial also adds tl.assume hints and an optional fused leaky-ReLU activation applied while the accumulator is still float32; both are worth reading in the original.

What each part is doing

Accumulate in float32. Inputs are float16, but the running sum is float32, and only the final tile is cast down. Accumulating in half precision across a long K loop would lose bits quickly. This mirrors what cuBLAS-class kernels do in spirit, though the exact internal behavior of vendor libraries is theirs to define.

Modulo on the row and column offsets. % M and % N wrap out-of-range rows and columns onto valid ones so loads never go out of bounds. The wrapped values are garbage for those lanes, but the final store uses c_mask, which discards them. The K dimension is masked separately with other=0.0, because zeros are neutral for a dot product.

Epilogue fusion. Because the tile is still in registers after the K loop, you can apply an activation, bias, scale or dtype conversion before the single store. This is where Triton matmuls routinely beat “library matmul then separate elementwise kernel.”

Grouped ordering for L2 reuse

If program ids simply walked tiles of C in row-major order, consecutive programs would stream through an entire row of tiles, loading new columns of B each time while reusing one band of A. The tutorial groups tiles into bands of GROUP_SIZE_M rows and walks each band column by column. Its illustrative 9×9-tile example says row-major needs 90 blocks loaded into SRAM to compute the first 9 output blocks, versus 54 with grouping.

The tutorial reports that this ordering improved performance by more than 10% on some hardware, from roughly 220 to 245 TFLOPS on an A100. Treat that as one author’s measurement on one device. Whether grouping helps on your GPU depends on L2 size, problem shape and how many programs are resident at once.

Reading the autotuner

@triton.autotune benchmarks every config the first time it sees a new combination of the key arguments (M, N, K here), caches the winner, and reuses it afterwards. Three details matter in production:

Cost scales with configs times distinct shapes. Each config is a separate compile plus several timed runs. An inference server with ragged batch sizes can retune constantly unless you bucket shapes or choose the key carefully.
Cold-start latency. The first request for each new key pays the tuning cost. Warm up the shapes you care about at startup.
Persistence. Make sure your deployment caches compiled kernels and tuning results on disk, or every pod restart repeats the work. Check the cache settings for your Triton version.

What the tutorial’s numbers do and do not tell you

The tutorial states that the kernel reaches performance on par with cuBLAS or rocBLAS. In its FP16 square benchmark, Triton is close to cuBLAS at large sizes, with about 216.8 TFLOPS for Triton versus about 219.0 for cuBLAS at M=N=K=4096, slower at some sizes (512 and 1408) and faster at others (896 and 1792). Those are the docs’ own runs on their own hardware. They show that a 60-line Triton kernel can approach a vendor library on a friendly, dense, square case. They do not show that you will match cuBLAS on skinny shapes, on the newest architecture, or at FP8, where the tutorial notes there was no cuBLAS comparison available through torch.matmul at the time.

The Compiler Pipeline: From Python to PTX

In one paragraph: Triton parses the Python AST of a @triton.jit function into Triton IR (TTIR), lowers it to a GPU-specific form (TTGIR) that assigns data layouts and applies pipelining and tensor-core mapping, converts that to LLVM IR, emits PTX, and assembles a binary the driver loads. Each stage is inspectable.

Figure 2: The Triton compilation pipeline. The Python function becomes Triton IR, then GPU-layout IR, LLVM IR, PTX and a binary cached for later launches.

What the stages mean

The Triton README confirms the architectural bones: version 2.0 rewrote the backend to use MLIR, and LLVM generates code for GPUs. The stage names TTIR and TTGIR are the names commonly used in Triton’s compiler and debug output for the generic and GPU-specific dialect levels. I could not retrieve a page that enumerates the dialects, so the following description is my summary of widely documented behavior rather than a quoted specification.

TTIR is hardware-agnostic. It represents block-level tensors, pointer arithmetic, loads, stores and dot operations.
TTGIR attaches layouts that say how each tensor is distributed across threads, warps and shared memory. This is where most performance decisions happen: choosing a blocked layout for global loads, an MMA layout for tensor-core operands, inserting layout conversions through shared memory, and software-pipelining loops according to num_stages.
LLVM IR to PTX is conventional compilation. On AMD the final target differs, which is how one frontend serves two vendors.

How to look inside

The README documents environment knobs for debugging: MLIR_ENABLE_DUMP dumps IR before MLIR passes, LLVM_IR_ENABLE_DUMP dumps before LLVM passes, and TRITON_REPRODUCER_PATH writes an MLIR reproducer before each compiler stage. In a notebook, the compiled-kernel handle returned from a launch or warmup call also carries the generated IR and metadata, such as register counts and shared-memory use, which the softmax tutorial reads through kernel.n_regs and kernel.metadata.shared.

When a kernel is slower than expected, the useful questions are concrete. How many registers does it use? Did the compiler insert shared-memory layout conversions you did not expect? Is the dot operating on a tensor-core layout? Is the loop actually pipelined? You answer these from the IR and the profiler, not from reading source.

Why this matters for portability

Because the programming model is expressed above the thread level, the same source can target different hardware. The cost is that peak performance on a specific architecture’s newest features (for example, Hopper’s asynchronous warpgroup instructions, or Blackwell’s tensor memory) depends on the compiler’s support landing, not on you writing the instruction directly. That dependency is the root of the “when should I leave Triton” question.

torch.compile and Inductor: Triton You Did Not Write

PyTorch 2’s torch.compile captures a graph with TorchDynamo, hands it to AOTAutograd and then to TorchInductor. On CUDA devices Inductor generates Triton kernels for fused pointwise and reduction groups; on CPU it generates C++. That is Inductor’s documented design, and it is the main way most teams consume Triton.

import torch

def gelu_residual_norm(x, residual, weight):
    y = torch.nn.functional.gelu(x) + residual
    return torch.nn.functional.layer_norm(y, y.shape[-1:], weight)

compiled = torch.compile(gelu_residual_norm)
x = torch.randn(8192, 4096, device="cuda", dtype=torch.float16)
r = torch.randn_like(x)
w = torch.ones(4096, device="cuda", dtype=torch.float16)
out = compiled(x, r, w)

Run it with the environment variable TORCH_LOGS=output_code (or set TORCHINDUCTOR_DEBUG-style options per the PyTorch docs for your version) and you can read the Triton source Inductor produced. Reading these generated kernels is one of the fastest ways to learn idiomatic Triton, because they show exactly how a compiler fuses an elementwise chain around a reduction.

The practical rule: try torch.compile first. If Inductor already fuses your pattern, a handwritten kernel buys little and costs maintenance. Write your own Triton when the fusion crosses a boundary the compiler will not cross (custom data structures, unusual masking, a hand-built online algorithm), when you need a specific tiling the heuristics do not choose, or when you need a kernel that behaves predictably outside the compiler’s graph capture.

Profiling Triton Kernels: ncu, Proton and Roofline Thinking

In one paragraph: time the kernel first, then ask whether it is memory-bound or compute-bound, and only then profile at instruction level. Use Triton’s benchmarking helpers for wall-clock numbers, Nsight Compute (ncu) for hardware counters, and Triton’s Proton profiler for lightweight, Triton-aware tracing.

Start with a roofline estimate

Compute your kernel’s arithmetic intensity: floating-point operations divided by bytes moved to and from HBM. A softmax does a few operations per element and moves two elements, so it is firmly bandwidth-bound, and its best-case time is bytes divided by memory bandwidth. A large matmul does 2*M*N*K operations on (M*K + K*N + M*N) elements, so its intensity grows with size and it becomes compute-bound. Compare your measured time with the bound for the right resource. If a bandwidth-bound kernel runs at a fraction of peak bandwidth, look at coalescing, occupancy and masking waste. If a compute-bound kernel is far below peak tensor-core throughput, look at tile shape, pipelining and layout conversions.

Always use the device’s published specifications for the bounds, and note that sustained bandwidth is lower than the datasheet peak.

Timing correctly

GPU launches are asynchronous, so wall-clock timing without synchronization measures nothing useful. Triton ships triton.testing.do_bench, which handles warmup, repetition and cache flushing, and triton.testing.perf_report for sweeping sizes into a plot, as the tutorials do. Report medians or percentiles over many runs and fix clocks where you can; run-to-run variance on a shared GPU can exceed the gain you are chasing.

Nsight Compute

NVIDIA’s ncu collects per-kernel counters: achieved occupancy, DRAM throughput, L2 hit rates, warp stall reasons, tensor-core utilization and register usage. Triton kernels appear under generated names (typically derived from your function name), so filter with --kernel-name. Build the habit of checking three numbers first: DRAM throughput as a percentage of peak, achieved occupancy, and the dominant warp-stall reason. Add -lineinfo-style source correlation where your toolchain supports it, though mapping PTX back to Python lines is not always precise.

Proton

Proton is the profiler that lives in the Triton repository. I attempted to fetch its documentation while preparing this post and the page was blocked, so I am not quoting its command-line flags or output format here. Consult the third_party/proton directory of the Triton repository for the current interface before you rely on it. In general terms it is intended for Triton-aware, low-overhead profiling and can sit alongside ncu rather than replace it.

A repeatable workflow

Verify correctness against a reference in float32.
Measure with do_bench across the shapes you ship.
Compute the roofline bound and decide which resource limits you.
Inspect registers and shared memory from the compiled handle; look for spills.
Use ncu to confirm the bottleneck before changing the kernel.
Change one thing, such as block size, num_warps, num_stages or access pattern, and remeasure.

Triton vs CUDA, CUTLASS and CuTe DSL: A Decision Framework

In one paragraph: choose by the unit of control you need. Use torch.compile for fusion you do not want to write, Triton or Helion for custom block-level kernels, CuTe DSL or CUTLASS C++ when you need explicit control of tensor-core pipelines and data layouts, and raw CUDA C++ when you need thread-level behavior no tile abstraction exposes.

Figure 4: A decision tree for picking a kernel tool. Start from the job: fuse existing ops, write a custom block kernel, chase peak tensor-core control, or target non-NVIDIA hardware.

The landscape, with what is verified

CUTLASS and CuTe DSL. NVIDIA’s CUTLASS documentation describes CUTLASS 4 as adding Python-based DSLs to the long-standing C++ template library, and CuTe DSL as the first of them. According to that page, CuTe DSL is a low-level programming model that mirrors the CuTe C++ abstractions (layouts, tensors, hardware atoms and control over the thread and data hierarchy), targets tensor cores on Ampere, Hopper and Blackwell, and was in public beta, with the page expecting it to leave beta by the end of summer 2026. The same page says CUTLASS C++ supports Volta through Blackwell and that release 4.8 adds initial Rubin support. It also claims the DSL gives a smoother learning curve, much faster compile times than C++ templates, and native framework integration without performance compromises. Those are NVIDIA’s claims; I did not independently benchmark them, and beta status may have changed since the page was retrieved.

The conceptual difference from Triton is where the abstraction sits. CuTe’s central idea is the layout, a function from logical coordinates to memory offsets, composed algebraically. You decide how data is partitioned across threads and how it moves between global memory, shared memory and registers. Triton infers those decisions from block-level code. CuTe DSL gives you the explicit control, in Python, without C++ template metaprogramming.

Helion. PyTorch’s Helion, announced in beta on October 22, 2025 per the PyTorch blog, is a Python-embedded DSL described as “PyTorch with Tiles.” A hl.tile loop defines a tile iteration space, and Helion compiles the device code to a Triton kernel, using TorchInductor to lower PyTorch ops inside it. Its pitch is an implicit autotuning space: one tile call stands for thousands of Triton configurations covering block sizes, loop orders, L2 grouping and indexing strategy. The post says a first autotune typically takes about ten minutes, and gives an example of 1,520 configurations searched in 586.6 seconds. It also reports geometric-mean speedups over eager PyTorch on NVIDIA B200 (Helion 3.27x, torch.compile max-autotune 2.7x, handwritten Triton 1.76x) and AMD MI350X (2.37x, 2.26x and 1.65x). These are vendor-reported figures from the project’s own benchmark suite, and the blog’s own caveat applies: they depend on the kernel set and hardware.

TileLang. TileLang is an academic and open-source tile-programming language with a paper titled “TileLang: A Composable Tiled Programming Model for AI Systems” on arXiv (2504.17577) and an ICLR 2026 conference listing. I have not verified its backend list or current maturity in this post. The Helion blog cites a case where Helion beat TileLang on a Mamba-2 chunk-scan kernel on H100, which is a vendor comparison, not a neutral one.

Pallas. Pallas is JAX’s kernel extension, aimed at writing custom kernels for TPUs and GPUs from JAX. The exact set of current backends and lowering paths for GPUs has changed over time, so I have not stated them here; check the JAX documentation before choosing it. For teams already on JAX it is the natural counterpart to Triton.

Raw CUDA C++. Still the escape hatch. Only CUDA gives you every instruction, every memory space and every synchronization primitive, at the price of everything else.

A comparison matrix

Need	torch.compile	Triton	Helion	CuTe DSL / CUTLASS	Raw CUDA
Fuse existing PyTorch ops	Best fit	Overkill	Overkill	Overkill	Overkill
Custom fused block kernel	Limited	Strong fit	Strong fit	Possible	Possible
Control over tensor-core pipeline	None	Indirect	Indirect	Direct	Direct
Learning curve	Low	Medium	Low to medium	High	Highest
Portability across vendors	Broad	NVIDIA and AMD	Triton targets	NVIDIA only	NVIDIA only
Autotuning burden	Automatic	You configure	Automatic search	You design	You design
Newest-architecture features	Follows Triton	Follows compiler	Follows Triton	Earliest	Earliest

The ratings are my qualitative judgments from the documented design of each tool, not measured results. The portability row reflects that Triton lists NVIDIA and AMD targets while CuTe DSL and CUTLASS target NVIDIA tensor cores.

The reasoning behind the matrix

Productivity versus ceiling. Every layer you move up trades control for effort. A fused softmax in Triton is about twenty lines. The equivalent hand-tuned CUDA kernel, with vectorized loads, warp shuffles for the reductions and a shared-memory staging plan, is several times that and needs retuning for each architecture. In exchange, CUDA lets you do things the tile model cannot express.

Where Triton falls short. The weak spots are workloads that need fine-grained, warp-specialized pipelines: producer warps issuing asynchronous copies while consumer warps run tensor-core instructions, explicit use of cluster-level shared memory, or bespoke scheduling between SMs. Whether the Triton compiler exposes or automatically derives a given pattern depends on the version and the architecture. This is the situation where teams reach for CUTLASS or CuTe, which are designed around exactly those primitives.

Where CuTe DSL fits. It targets the same practitioners who used CUTLASS C++ but want Python-native iteration and shorter compile cycles. If you are writing a GEMM variant with an unusual epilogue or a grouped GEMM for mixture-of-experts layers, and you need to pin down the pipeline structure, this is the intended zone. If your kernel is a fused elementwise or reduction chain, you are over-engineering.

Where Helion fits. Between torch.compile and Triton. You write code that looks like PyTorch with explicit tiles, and the system explores the Triton configuration space for you. Its tradeoff is autotune time and its beta maturity, plus dependence on Triton’s code generation as its only backend per the project’s own description.

Fusion in inference systems

Triton kernels show up throughout inference stacks: fused attention variants, dequantize-and-multiply for weight-only quantization, rotary embeddings, sampling and top-k kernels. Techniques like speculative decoding and LLM prompt caching change the shapes and reuse patterns those kernels see (more small, verification-sized batches; reused KV blocks), which is another reason shape-aware autotuning keys matter in serving.

Trade-offs, Gotchas, and What Goes Wrong

Triton’s abstraction is a loan, and these are the usual repayments.

Masking bugs that pass tests. Forgetting a mask is invisible on sizes that are multiples of the block size, and it produces out-of-bounds reads or silent corruption on the ragged tail. Test with awkward sizes such as 1, 127, 1001 and prime numbers. Use the other= value that is neutral for the reduction you perform: -inf for max and softmax, 0 for sums and dots.

Numerical stability. Always subtract the row max before exp. Accumulate reductions and dot products in float32 even if inputs are float16 or bfloat16. Remember the tutorial’s warning that Triton’s exp is fast but approximate; do not compare against a float64 reference with tight tolerances and then conclude the kernel is broken. Fused online algorithms (rescaling a running max and sum) are more delicate than the one-pass-per-row case.

Register pressure and spills. Large tiles and many live intermediates raise per-thread register use, which lowers occupancy and can spill to local memory, a slow path through the cache hierarchy. The occupancy formula in the softmax tutorial (registers per SM divided by registers per thread, threads per warp and warps) makes the relationship explicit. If increasing BLOCK_SIZE or num_warps makes things slower, suspect spills before suspecting the compiler.

Autotune cost and instability. Tuning is compile-heavy, results are shape-specific, and a noisy GPU can pick the wrong winner. Pin tuning to representative shapes, bucket dynamic dimensions, persist caches, and re-tune when you upgrade the driver, Triton version or GPU.

Version coupling. Triton, PyTorch and CUDA toolkit versions move together, and a kernel that is fast on one release can regress on another. Pin versions in production and include a performance smoke test in CI.

Hardware coupling. A config tuned for an A100 will not be best on an H100, B200 or an AMD part; the tutorial itself uses a different config list for HIP. Treat tuned parameters as data keyed by device, not constants in the source.

Debuggability. Printing inside kernels is awkward and stepping through compiled code is limited. The README’s IR dump and reproducer variables help for compiler issues, but for logic bugs your best tool is a small reference implementation and property tests across shapes.

Writing a kernel you did not need. The most common failure is organizational. A custom kernel that saves 5% of a layer that is 3% of end-to-end time is not worth owning. Profile the whole model first, and let the numbers justify the maintenance.

Practical Recommendations

Start from the application, not the kernel. Profile a real workload, find the operations dominated by memory traffic, and see whether torch.compile already fuses them. If it does, stop. If it does not, write a minimal Triton kernel that reproduces one fusion, validate it against a float32 PyTorch reference on awkward shapes, and measure it end to end.

Keep your first kernels simple: one program per row or per tile, no persistence tricks, explicit masks, float32 accumulation. Move to persistent scheduling, grouped ordering and exotic pipelining only when a profile says the simple version leaves bandwidth or compute on the table. Read the code Inductor generates for your pattern, since it encodes many of the idioms Triton’s authors consider normal.

Adopt CuTe DSL or CUTLASS when you have evidence that the tile abstraction is the limit: tensor-core utilization stuck well below what your hardware vendor’s library reaches, or a required hardware feature the compiler does not expose to you. Given that CuTe DSL was in public beta at the time of the documentation I read, budget for API changes.

A short checklist before merging a custom kernel:

Passes correctness tests at sizes 1, one less than a block, exactly one block, one more than a block, and a large prime.
Neutral other= values chosen and commented for every masked load.
Float32 accumulation, with the tolerance justified.
Benchmarked with do_bench against the eager and torch.compile baselines on production shapes.
Roofline bound computed and the gap explained.
Autotune keys, config list and cache location documented.
Triton, PyTorch and driver versions pinned, with a CI performance guard.

Frequently Asked Questions

What is Triton GPU kernel programming?

Triton GPU kernel programming means writing GPU kernels in a Python-embedded language where the unit of work is a block of data rather than a single thread. You decorate a function with @triton.jit, compute offsets for a tile, load with masks, compute, and store. The compiler maps that onto threads, warps, shared memory and tensor cores, so you get CUDA-class performance on many workloads with far less low-level code.

Is Triton faster than CUDA?

Not inherently. Triton compiles to the same hardware, so its ceiling is at most what a perfect CUDA kernel can reach. Its value is reaching good performance with much less effort, and sometimes matching vendor libraries: the official matmul tutorial reports FP16 results close to cuBLAS at large square sizes on its test hardware. Hand-tuned CUDA, CUTLASS or CuTe can still win when you need explicit pipeline control or the newest hardware features.

Do I need to know CUDA to learn Triton?

You do not need to write CUDA, but you do need the mental model it teaches: the memory hierarchy, coalesced access, occupancy and why fusing kernels saves bandwidth. Without those, Triton’s tuning knobs such as num_warps and num_stages look arbitrary. Learn the hierarchy first, then use Triton to apply it with far less boilerplate than thread-level code requires.

How does torch.compile relate to Triton?

On NVIDIA GPUs, torch.compile uses TorchInductor, which generates Triton kernels for fused pointwise and reduction code. That means many users run Triton kernels without writing any. You can inspect the generated code through PyTorch’s logging options. Writing your own Triton kernel makes sense when Inductor cannot fuse your pattern or you need a specific algorithm or tiling it does not choose.

When should I choose CUTLASS or CuTe DSL over Triton?

Choose them when you need direct control over tensor-core pipelines, data layouts and memory movement, for example a custom GEMM or grouped GEMM where you must dictate the structure, or when you need the earliest access to new NVIDIA architecture features. CuTe DSL offers that control in Python, while CUTLASS C++ is the mature option. If your kernel is a fusion of elementwise and reduction operations, Triton is usually simpler.

How long does Triton autotuning take?

It depends on the number of configurations, how long each compiles, and how many distinct shapes you tune. Each config means a separate compile plus timed runs, so a long list across many shapes can take minutes or more. Helion’s documentation cites roughly ten minutes for a typical first autotune of one kernel. Reduce the cost by pruning configs, bucketing shapes and persisting caches.