OpenAI Triton From Zero • Thibault Castells

OpenAI Triton is a Python-based language and compiler for writing high-performance GPU kernels. It is useful when you want more control than ordinary PyTorch code gives you, but you do not want to write CUDA C++ directly. OpenAI introduced Triton 1.0 as an open-source Python-like language for efficient GPU code, and the current Triton repository describes it as a language and compiler for writing efficient custom deep-learning primitives.

Important naming warning: this tutorial is about Triton the GPU kernel programming language, not NVIDIA Triton Inference Server, which is a model-serving system for deploying models through HTTP/gRPC and backends such as Python, PyTorch, ONNX Runtime, TensorRT, TensorRT-LLM, and vLLM.

The goal is to build the mental model first, then make it concrete with progressively richer kernels: vector addition, fused elementwise operations, row-wise softmax, matrix multiplication, autotuning, integration with PyTorch, autograd, and debugging.

0. The core idea

Triton lets you write custom GPU kernels in Python while thinking in terms of blocks of tensors rather than individual CUDA threads. Instead of this CUDA-style mental model:

one thread computes one element
many threads form a block
many blocks form a grid

Triton encourages this one:

one Triton program computes one block of elements
many Triton programs form a grid
inside each program, operations happen on small tensor blocks

Most of the tutorial is just a gradual unpacking of that idea.

1. Why Triton exists

1.1 The PyTorch baseline and the fusion problem

Let us start with normal PyTorch. Suppose you write:

import torch

x = torch.randn(1024, 1024, device="cuda")
y = torch.randn(1024, 1024, device="cuda")

z = torch.relu(x * 0.5 + y)

This is easy. But under the hood, eager PyTorch may execute several operations:

tmp1 = x * 0.5
tmp2 = tmp1 + y
z = relu(tmp2)

That can mean:

read x
write tmp1

read tmp1
read y
write tmp2

read tmp2
write z

For many GPU workloads, especially elementwise and reduction-heavy workloads, the problem is not arithmetic. The problem is memory traffic. A GPU is very fast at math, but moving tensors back and forth through memory is expensive. The ideal kernel would do this:

read x once
read y once
compute x * 0.5 + y
apply relu
write z once

That is called kernel fusion. Triton exists because sometimes you want to write exactly the kernel you need:

one custom kernel
no unnecessary temporary tensors
better memory access pattern
more control over tiling
more control over precision
more control over fusion

OpenAI’s original Triton announcement framed the motivation similarly: deep-learning research is often implemented by combining framework operators, which is convenient but can create temporary tensors and extra memory movement; custom kernels can solve this, but CUDA programming is hard.

1.2 Triton vs CUDA vs PyTorch vs torch.compile

Triton is easiest to understand by comparing it with the tools around it.

Tool	What you write	Strength	Weakness
PyTorch eager	Normal Python tensor code	Easiest to use	May launch many separate kernels
torch.compile	Normal PyTorch code	Can automatically fuse and optimize many graphs	Less explicit control
Triton	Python GPU kernels	Explicit custom kernels without CUDA C++	You must reason about memory and blocks
CUDA C++	Low-level GPU code	Maximum control	Harder and more verbose
Vendor libraries	Calls like cuBLAS/cuDNN	Extremely optimized standard ops	Hard to customize

A practical rule:

Use PyTorch eager when performance is already good enough. Use torch.compile when you want automatic optimization with minimal code changes. Use Triton when you need a custom fused kernel, custom memory layout, custom reduction, custom attention variant, custom quantization kernel, or custom matmul-like operation. Use CUDA C++ when Triton is not expressive enough or when you need very low-level control.

PyTorch also supports user-defined Triton kernels inside torch.compile, so Triton is not separate from the PyTorch ecosystem; it can be integrated into compiled PyTorch programs.

1.3 The core mental model

Triton is five things glued together:

The Python launcher allocates tensors and calls a Triton kernel.
The Triton JIT function is decorated with @triton.jit.
The grid decides how many independent program instances run.
Each program instance works on a block or tile of data.
The compiler lowers the Triton program to GPU code.

In pseudocode:

# Python side.
x = torch.randn(..., device="cuda")
y = torch.randn(..., device="cuda")
out = torch.empty_like(x)

# Launch many Triton program instances.
kernel[grid](x, y, out, ...)

Inside the kernel:

# GPU side.
program_id = tl.program_id(0)
offsets = program_id * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)

x_block = tl.load(x_ptr + offsets)
y_block = tl.load(y_ptr + offsets)

out_block = x_block + y_block

tl.store(out_ptr + offsets, out_block)

The key difference with CUDA is that Triton wants you to write operations over blocks. In practice, you write a blocked program and let the compiler map that block-level work to GPU execution.

1.4 From Python to GPU

When you call a Triton kernel, roughly this happens:

Python code
  |
  v
Triton kernel launch
  |
  v
JIT compilation, if this shape/config was not compiled before
  |
  v
GPU launch
  |
  v
many program instances run in parallel
  |
  v
each program loads a block of memory
  |
  v
each program computes on that block
  |
  v
each program stores a block of output

The first call may be slower because Triton compiles the kernel. Later calls with the same specialization are faster because Triton can reuse compiled code. That means you should not benchmark a Triton kernel by timing only the first call.

2. Setup and the first kernel

2.1 Installation

The official Triton docs show the package installation as:

pip install triton

The current docs say binary wheels are available for CPython 3.10 through 3.14, and the GitHub compatibility section currently lists Linux support, NVIDIA GPUs with compute capability 8.0 or newer, AMD GPUs with ROCm 6.2 or newer, and CPU support as under development. Using uv pip:

python3 -m venv .venv
source .venv/bin/activate

uv pip install --upgrade pip setuptools wheel
uv pip install torch triton

Check the installation:

python - <<'PY'
import torch
import triton

print("torch:", torch.__version__)
print("triton:", triton.__version__)
print("cuda available:", torch.cuda.is_available())

if torch.cuda.is_available():
    print("gpu:", torch.cuda.get_device_name(0))
PY

For this tutorial, assume:

Python works.
PyTorch works.
A supported GPU is visible.
Triton imports successfully.

2.2 First PyTorch baseline

Create torch_add.py:

#!/usr/bin/env python3
import argparse

import torch


def parse_args() -> argparse.Namespace:
    parser = argparse.ArgumentParser(description="Simple PyTorch vector add.")
    parser.add_argument("--size", type=int, default=1_000_000)
    parser.add_argument("--seed", type=int, default=0)
    return parser.parse_args()


def main() -> None:
    args = parse_args()

    if not torch.cuda.is_available():
        raise RuntimeError("CUDA is not available.")

    torch.manual_seed(args.seed)

    x = torch.randn(args.size, device="cuda")
    y = torch.randn(args.size, device="cuda")

    out = x + y

    print(out[:8])


if __name__ == "__main__":
    main()

Run:

python torch_add.py

This is the operation we will rewrite in Triton. Not because vector addition is hard. Because vector addition contains the whole Triton mental model in miniature.

2.3 First Triton kernel: vector addition

Create triton_add.py:

#!/usr/bin/env python3
import argparse

import torch
import triton
import triton.language as tl


@triton.jit
def add_kernel(
    x_ptr,
    y_ptr,
    out_ptr,
    n_elements,
    BLOCK_SIZE: tl.constexpr,
):
    # Each program instance owns one contiguous block of vector elements.
    pid = tl.program_id(axis=0)

    block_start = pid * BLOCK_SIZE
    offsets = block_start + tl.arange(0, BLOCK_SIZE)

    # The last block may run past n_elements, so every memory access is masked.
    mask = offsets < n_elements

    x = tl.load(x_ptr + offsets, mask=mask, other=0.0)
    y = tl.load(y_ptr + offsets, mask=mask, other=0.0)

    out = x + y

    tl.store(out_ptr + offsets, out, mask=mask)


def add_triton(
    x: torch.Tensor,
    y: torch.Tensor,
    block_size: int,
) -> torch.Tensor:
    if x.shape != y.shape:
        raise ValueError(f"Shape mismatch: {x.shape} vs {y.shape}")
    if x.device != y.device:
        raise ValueError(f"Device mismatch: {x.device} vs {y.device}")
    if not x.is_cuda:
        raise ValueError("This demo expects CUDA tensors.")

    out = torch.empty_like(x)
    n_elements = out.numel()

    # Launch enough programs to cover the whole vector.
    grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),)

    add_kernel[grid](
        x,
        y,
        out,
        n_elements,
        BLOCK_SIZE=block_size,
    )

    return out


def parse_args() -> argparse.Namespace:
    parser = argparse.ArgumentParser(description="Triton vector add.")
    parser.add_argument("--size", type=int, default=1_000_000)
    parser.add_argument("--block-size", type=int, default=1024)
    parser.add_argument("--seed", type=int, default=0)
    return parser.parse_args()


def main() -> None:
    args = parse_args()

    if not torch.cuda.is_available():
        raise RuntimeError("CUDA is not available.")

    torch.manual_seed(args.seed)

    x = torch.randn(args.size, device="cuda")
    y = torch.randn(args.size, device="cuda")

    out_torch = x + y
    out_triton = add_triton(x, y, block_size=args.block_size)

    torch.testing.assert_close(out_triton, out_torch)

    print(out_triton[:8])
    print("OK")


if __name__ == "__main__":
    main()

Run:

python triton_add.py

The official Triton vector-add tutorial uses the same essential ingredients: @triton.jit, tl.program_id, tl.arange, masked tl.load, masked tl.store, a launch grid, and a Python wrapper that allocates the output.

2.4 Understanding the kernel line by line

The kernel starts like this:

@triton.jit
def add_kernel(...):

@triton.jit means this function is not normal Python. It is compiled by the Triton compiler and runs on the GPU. A JIT function only has access to Python primitives, Triton builtins, function arguments, and other JIT functions. Then:

pid = tl.program_id(axis=0)

This asks:

Which program instance am I?

If the grid has 1000 programs, then pid ranges from:

0, 1, 2, ..., 999

Then:

block_start = pid * BLOCK_SIZE
offsets = block_start + tl.arange(0, BLOCK_SIZE)

If BLOCK_SIZE = 8, then:

pid = 0 -> offsets = [0, 1, 2, 3, 4, 5, 6, 7]
pid = 1 -> offsets = [8, 9, 10, 11, 12, 13, 14, 15]
pid = 2 -> offsets = [16, 17, 18, 19, 20, 21, 22, 23]

Then:

mask = offsets < n_elements

This protects the final block. If the vector length is 18 and BLOCK_SIZE = 8:

pid = 0 -> offsets 0..7   -> all valid
pid = 1 -> offsets 8..15  -> all valid
pid = 2 -> offsets 16..23 -> only 16 and 17 are valid

Without the mask, the kernel would read and write out of bounds. Then:

x = tl.load(x_ptr + offsets, mask=mask, other=0.0)
y = tl.load(y_ptr + offsets, mask=mask, other=0.0)

This loads a block of values. Then:

out = x + y

This adds the whole block. Then:

tl.store(out_ptr + offsets, out, mask=mask)

This writes the block back. That is Triton.

2.5 A Triton program is not a CUDA thread

This is the most common beginner confusion. In CUDA, you often think:

one thread handles one element

In Triton, you usually think:

one program handles one block of elements

For vector addition with BLOCK_SIZE = 1024:

Program 0 handles elements 0..1023
Program 1 handles elements 1024..2047
Program 2 handles elements 2048..3071
...

Inside one program:

offsets = block_start + tl.arange(0, BLOCK_SIZE)

offsets is not a Python list. It is a Triton block tensor. When you write:

x = tl.load(x_ptr + offsets)

you are loading many elements. When you write:

out = x + y

you are doing many additions. When you write:

tl.store(out_ptr + offsets, out)

you are storing many elements. So the unit you write is:

block in, block compute, block out

2.6 BLOCK_SIZE, tl.constexpr, and the grid

BLOCK_SIZE controls how much work each Triton program does. Small BLOCK_SIZE:

+ more programs
+ less work per program
- more overhead
- less reuse per program

Large BLOCK_SIZE:

+ fewer programs
+ more work per program
+ better amortization
- can use more registers
- can reduce occupancy
- may be too large for the compiler or hardware

For vector addition, 1024 is a reasonable beginner value. For real kernels, BLOCK_SIZE is a tuning parameter.

This line matters:

BLOCK_SIZE: tl.constexpr

It means BLOCK_SIZE is a compile-time constant. That matters because Triton needs to know the shape of:

tl.arange(0, BLOCK_SIZE)

at compile time. A useful mental rule:

Runtime arguments describe data.
constexpr arguments describe code shape.

Examples of runtime arguments:

x_ptr
y_ptr
out_ptr
n_elements

Examples of compile-time arguments:

BLOCK_SIZE
BLOCK_M
BLOCK_N
BLOCK_K
whether to fuse an activation

The launch looks like this:

grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),)

add_kernel[grid](...)

The grid says how many program instances to launch. For n_elements = 10_000 and BLOCK_SIZE = 1024:

ceil(10_000 / 1024) = 10 programs

So Triton launches:

program 0
program 1
program 2
...
program 9

Each program handles one block.

2.7 A tiny CPU mental simulator

Here is a CPU-only toy version of what our Triton vector-add kernel is doing. This is not Triton code. It is a teaching sketch.

#!/usr/bin/env python3
import argparse


def parse_args() -> argparse.Namespace:
    parser = argparse.ArgumentParser(description="Toy Triton mental model.")
    parser.add_argument("--n-elements", type=int, default=20)
    parser.add_argument("--block-size", type=int, default=8)
    return parser.parse_args()


def main() -> None:
    args = parse_args()

    n_programs = (args.n_elements + args.block_size - 1) // args.block_size

    for pid in range(n_programs):
        block_start = pid * args.block_size
        offsets = list(range(block_start, block_start + args.block_size))
        valid_offsets = [offset for offset in offsets if offset < args.n_elements]

        print(f"program {pid}")
        print(f"  offsets:       {offsets}")
        print(f"  valid offsets: {valid_offsets}")


if __name__ == "__main__":
    main()

Run:

python toy_triton_grid.py

This is what tl.program_id, tl.arange, and mask are doing conceptually.

3. Benchmarking and fusion

3.1 Benchmarking correctly

Do not benchmark GPU kernels with plain time.time() unless you know exactly where you synchronize. GPU work is asynchronous. This is wrong or at least incomplete:

start = time.perf_counter()
out = add_triton(x, y, block_size=1024)
end = time.perf_counter()

The CPU may measure only launch overhead. Triton provides benchmarking utilities such as triton.testing.do_bench, which handles warmup and repeated timing. Create benchmark_add.py:

#!/usr/bin/env python3
import argparse

import torch
import triton
import triton.language as tl


@triton.jit
def add_kernel(
    x_ptr,
    y_ptr,
    out_ptr,
    n_elements,
    BLOCK_SIZE: tl.constexpr,
):
    pid = tl.program_id(axis=0)

    # One program handles BLOCK_SIZE consecutive elements.
    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements

    # Masked loads keep the final partial block safe.
    x = tl.load(x_ptr + offsets, mask=mask, other=0.0)
    y = tl.load(y_ptr + offsets, mask=mask, other=0.0)

    tl.store(out_ptr + offsets, x + y, mask=mask)


def parse_args() -> argparse.Namespace:
    parser = argparse.ArgumentParser(description="Benchmark Triton vector add.")
    parser.add_argument("--size", type=int, default=16_777_216)
    parser.add_argument("--block-size", type=int, default=1024)
    parser.add_argument("--dtype", choices=["float16", "float32"], default="float32")
    return parser.parse_args()


def main() -> None:
    args = parse_args()

    if not torch.cuda.is_available():
        raise RuntimeError("CUDA is not available.")

    dtype = {
        "float16": torch.float16,
        "float32": torch.float32,
    }[args.dtype]

    x = torch.randn(args.size, device="cuda", dtype=dtype)
    y = torch.randn(args.size, device="cuda", dtype=dtype)
    out = torch.empty_like(x)

    n_elements = out.numel()
    grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),)

    def run_triton() -> None:
        # do_bench expects a callable, so the kernel launch is wrapped here.
        add_kernel[grid](
            x,
            y,
            out,
            n_elements,
            BLOCK_SIZE=args.block_size,
        )

    def run_torch() -> torch.Tensor:
        return x + y

    ms_triton = triton.testing.do_bench(run_triton, warmup=25, rep=100)
    ms_torch = triton.testing.do_bench(run_torch, warmup=25, rep=100)

    # Vector add reads x and y and writes out: three tensor-sized transfers.
    bytes_moved = 3 * x.numel() * x.element_size()
    triton_gbps = bytes_moved / (ms_triton * 1e-3) / 1e9
    torch_gbps = bytes_moved / (ms_torch * 1e-3) / 1e9

    print(f"Triton: {ms_triton:.4f} ms, {triton_gbps:.2f} GB/s")
    print(f"Torch:  {ms_torch:.4f} ms, {torch_gbps:.2f} GB/s")


if __name__ == "__main__":
    main()

Run:

python benchmark_add.py

For vector addition, do not expect magic. PyTorch is already very good at simple elementwise operations. The point is to learn the structure. Triton becomes more interesting when you fuse operations or write custom memory patterns.

3.2 Kernel fusion: where Triton starts to matter

Suppose we want:

out = relu(x * scale + bias)

A simple PyTorch version is:

out = torch.relu(x * scale + bias)

A fused Triton version does:

load x
multiply by scale
add bias
apply relu
store out

in one kernel. Create fused_affine_relu.py:

#!/usr/bin/env python3
import argparse

import torch
import triton
import triton.language as tl


@triton.jit
def affine_relu_kernel(
    x_ptr,
    out_ptr,
    scale,
    bias,
    n_elements,
    BLOCK_SIZE: tl.constexpr,
):
    pid = tl.program_id(axis=0)

    # Each program handles one block of x.
    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements

    x = tl.load(x_ptr + offsets, mask=mask, other=0.0)

    # Fuse affine transform and ReLU before writing the result once.
    y = x * scale + bias
    y = tl.maximum(y, 0.0)

    tl.store(out_ptr + offsets, y, mask=mask)


def affine_relu_triton(
    x: torch.Tensor,
    scale: float,
    bias: float,
    block_size: int,
) -> torch.Tensor:
    if not x.is_cuda:
        raise ValueError("This demo expects a CUDA tensor.")

    out = torch.empty_like(x)
    n_elements = x.numel()
    grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),)

    affine_relu_kernel[grid](
        x,
        out,
        scale,
        bias,
        n_elements,
        BLOCK_SIZE=block_size,
    )

    return out


def parse_args() -> argparse.Namespace:
    parser = argparse.ArgumentParser(description="Fused affine + ReLU in Triton.")
    parser.add_argument("--size", type=int, default=1_000_000)
    parser.add_argument("--scale", type=float, default=0.5)
    parser.add_argument("--bias", type=float, default=1.0)
    parser.add_argument("--block-size", type=int, default=1024)
    return parser.parse_args()


def main() -> None:
    args = parse_args()

    if not torch.cuda.is_available():
        raise RuntimeError("CUDA is not available.")

    x = torch.randn(args.size, device="cuda")

    out_torch = torch.relu(x * args.scale + args.bias)
    out_triton = affine_relu_triton(
        x,
        scale=args.scale,
        bias=args.bias,
        block_size=args.block_size,
    )

    torch.testing.assert_close(out_triton, out_torch)

    print(out_triton[:8])
    print("OK")


if __name__ == "__main__":
    main()

The important idea:

PyTorch expression:
  many conceptual operations

Triton kernel:
  one custom operation

Modern PyTorch compilers can often fuse simple cases automatically. Triton is what you reach for when you want explicit control.

3.3 Memory traffic is often the real enemy

For this operation:

out = torch.relu(x * scale + bias)

A naive implementation might do something like:

1. read x, write tmp1
2. read tmp1, write tmp2
3. read tmp2, write out

A fused implementation does:

1. read x, write out

That can be much faster because it reduces memory traffic. This same idea appears in the official Triton fused-softmax tutorial: a naive PyTorch softmax can require multiple reads and writes of the matrix, while a fused kernel can keep the row on-chip and reduce global-memory traffic.

4. Reductions and row-wise softmax

4.1 Row-wise softmax

Now let us write something more realistic. Softmax over each row:

out = torch.softmax(x, dim=1)

Mathematically:

softmax(x_i) = exp(x_i - max(x)) / sum_j exp(x_j - max(x))

The stable version subtracts the row maximum first. For one row:

load row
compute max
subtract max
compute exp
compute sum
divide
store row

This is a good Triton example because each program can handle one row. Create row_softmax.py:

#!/usr/bin/env python3
import argparse

import torch
import triton
import triton.language as tl


@triton.jit
def row_softmax_kernel(
    x_ptr,
    out_ptr,
    n_cols: tl.constexpr,
    x_row_stride,
    out_row_stride,
    BLOCK_SIZE: tl.constexpr,
):
    # One program processes one row.
    row_id = tl.program_id(axis=0)

    # BLOCK_SIZE is a power of two, so it may be larger than n_cols.
    col_offsets = tl.arange(0, BLOCK_SIZE)
    mask = col_offsets < n_cols

    x_ptrs = x_ptr + row_id * x_row_stride + col_offsets

    # Padding lanes use -inf so they do not affect the row maximum.
    x = tl.load(x_ptrs, mask=mask, other=-float("inf"))

    # Stable softmax: subtract the row maximum before exponentiating.
    x = x - tl.max(x, axis=0)
    numerator = tl.exp(x)
    denominator = tl.sum(numerator, axis=0)
    out = numerator / denominator

    out_ptrs = out_ptr + row_id * out_row_stride + col_offsets
    tl.store(out_ptrs, out, mask=mask)


def softmax_triton(
    x: torch.Tensor,
    num_warps: int,
) -> torch.Tensor:
    if x.ndim != 2:
        raise ValueError(f"Expected a 2D tensor, got shape {x.shape}")
    if not x.is_cuda:
        raise ValueError("This demo expects a CUDA tensor.")
    if not x.is_contiguous():
        x = x.contiguous()

    n_rows, n_cols = x.shape

    # Triton reductions work best with static block sizes.
    block_size = triton.next_power_of_2(n_cols)

    out = torch.empty_like(x)

    row_softmax_kernel[(n_rows,)](
        x,
        out,
        n_cols,
        x.stride(0),
        out.stride(0),
        BLOCK_SIZE=block_size,
        num_warps=num_warps,
    )

    return out


def parse_args() -> argparse.Namespace:
    parser = argparse.ArgumentParser(description="Row-wise softmax in Triton.")
    parser.add_argument("--n-rows", type=int, default=1024)
    parser.add_argument("--n-cols", type=int, default=1024)
    parser.add_argument("--num-warps", type=int, default=4)
    parser.add_argument("--seed", type=int, default=0)
    return parser.parse_args()


def main() -> None:
    args = parse_args()

    if not torch.cuda.is_available():
        raise RuntimeError("CUDA is not available.")

    torch.manual_seed(args.seed)

    x = torch.randn(args.n_rows, args.n_cols, device="cuda")
    out_torch = torch.softmax(x, dim=1)
    out_triton = softmax_triton(x, num_warps=args.num_warps)

    torch.testing.assert_close(out_triton, out_torch, rtol=1e-4, atol=1e-4)

    print(out_triton[0, :8])
    print("OK")


if __name__ == "__main__":
    main()

Run:

python row_softmax.py

The official Triton softmax tutorial uses the same high-level strategy: each program loads a row, masks padding, subtracts the maximum for numerical stability, computes exponentials, reduces the denominator, and stores the result. It also shows a practical Triton constraint: block sizes often need to be powers of two, so rows are padded internally and protected with masks.

4.2 Why softmax teaches more than vector add

Vector add:

load x
load y
add
store out

Softmax:

load row
reduce max
subtract max
exp
reduce sum
divide
store row

This teaches more Triton ideas:

block-level reductions
masking
power-of-two block sizes
numerical stability
working on rows
choosing num_warps

This is closer to real deep-learning kernels.

4.3 Reductions in Triton

In the softmax kernel, these two lines are reductions:

tl.max(x, axis=0)
tl.sum(numerator, axis=0)

Remember: x is a block tensor. If x has shape:

[BLOCK_SIZE]

then:

tl.max(x, axis=0)

returns one value: the maximum over that block. This is not a Python loop. It is compiled into GPU work. A useful mental model:

tl.load gives a block tensor.
tl.max reduces a block tensor.
tl.sum reduces a block tensor.
tl.store writes a block tensor.

4.4 num_warps and num_stages

When launching the softmax kernel, we used:

num_warps=num_warps

A warp is a group of GPU threads that execute together. In Triton, you usually do not directly write thread-level code, but you still choose how many warps participate in a program. Small num_warps:

+ less overhead
+ can be good for small blocks
- less parallelism inside a program

Large num_warps:

+ more parallelism inside a program
+ can help large reductions or matmuls
- more resource usage
- can reduce occupancy

Beginner rule:

Try 4 or 8.
Benchmark.
Do not guess forever.

You will often see:

num_stages=3
num_stages=4

This controls software pipelining depth. It matters especially for kernels like matrix multiplication where the kernel repeatedly loads the next tile while computing on the current tile. Beginner mental model:

num_warps  -> how much parallelism inside one program
num_stages -> how aggressively loads and compute are pipelined

Do not tune these by intuition alone. Benchmark.

5. Matrix multiplication and tuning

5.1 Matrix multiplication: the big example

Matrix multiplication is the classic Triton example. We want:

C = A @ B

where:

A has shape [M, K]
B has shape [K, N]
C has shape [M, N]

A naive mathematical loop is:

for m in range(M):
    for n in range(N):
        acc = 0
        for k in range(K):
            acc += A[m, k] * B[k, n]
        C[m, n] = acc

A GPU kernel should not compute one element at a time. It should compute tiles:

one Triton program computes one tile of C

C tile shape:
  [BLOCK_M, BLOCK_N]

A tile shape:
  [BLOCK_M, BLOCK_K]

B tile shape:
  [BLOCK_K, BLOCK_N]

Then:

C_tile += A_tile @ B_tile

The official Triton matrix multiplication tutorial teaches this blocked algorithm, block-level matrix multiplication, multidimensional pointer arithmetic, program ordering for better L2 cache hit rate, and autotuning.

5.2 A small educational matmul kernel

This is not meant to beat torch.matmul. It is meant to make the blocking idea concrete. Create matmul_basic.py:

#!/usr/bin/env python3
import argparse

import torch
import triton
import triton.language as tl


@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_M: tl.constexpr,
    BLOCK_N: tl.constexpr,
    BLOCK_K: tl.constexpr,
):
    pid = tl.program_id(axis=0)

    num_blocks_n = tl.cdiv(N, BLOCK_N)

    # Map the 1D program id to a 2D tile of C.
    block_m = pid // num_blocks_n
    block_n = pid % num_blocks_n

    # Row, column, and reduction offsets for this tile.
    offs_m = block_m * BLOCK_M + tl.arange(0, BLOCK_M)
    offs_n = block_n * BLOCK_N + tl.arange(0, BLOCK_N)
    offs_k = tl.arange(0, BLOCK_K)

    # Accumulate in float32 even though inputs and output are float16.
    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)

    for k_block in range(0, tl.cdiv(K, BLOCK_K)):
        k_offsets = k_block * BLOCK_K + offs_k

        # Build pointer matrices for the A and B tiles used in this K block.
        a_ptrs = a_ptr + offs_m[:, None] * stride_am + k_offsets[None, :] * stride_ak
        b_ptrs = b_ptr + k_offsets[:, None] * stride_bk + offs_n[None, :] * stride_bn

        # Boundary masks make non-multiple M/N/K sizes safe.
        a_mask = (offs_m[:, None] < M) & (k_offsets[None, :] < K)
        b_mask = (k_offsets[:, None] < K) & (offs_n[None, :] < N)

        a = tl.load(a_ptrs, mask=a_mask, other=0.0)
        b = tl.load(b_ptrs, mask=b_mask, other=0.0)

        # Compute one partial tile: [BLOCK_M, BLOCK_K] @ [BLOCK_K, BLOCK_N].
        acc = tl.dot(a, b, acc)

    c_ptrs = c_ptr + offs_m[:, None] * stride_cm + offs_n[None, :] * stride_cn
    c_mask = (offs_m[:, None] < M) & (offs_n[None, :] < N)

    # Store only the valid C elements for boundary tiles.
    tl.store(c_ptrs, acc, mask=c_mask)


def matmul_triton(
    a: torch.Tensor,
    b: torch.Tensor,
    block_m: int,
    block_n: int,
    block_k: int,
) -> torch.Tensor:
    if a.ndim != 2 or b.ndim != 2:
        raise ValueError("Expected two 2D tensors.")
    if a.shape[1] != b.shape[0]:
        raise ValueError(f"Shape mismatch: {a.shape} vs {b.shape}")
    if not a.is_cuda or not b.is_cuda:
        raise ValueError("This demo expects CUDA tensors.")
    if a.dtype != torch.float16 or b.dtype != torch.float16:
        raise ValueError("This demo expects float16 inputs.")

    a = a.contiguous()
    b = b.contiguous()

    M, K = a.shape
    _, N = b.shape

    c = torch.empty((M, N), device=a.device, dtype=torch.float16)

    # One program computes one [block_m, block_n] tile of C.
    grid = (
        triton.cdiv(M, block_m) * triton.cdiv(N, block_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),
        BLOCK_M=block_m,
        BLOCK_N=block_n,
        BLOCK_K=block_k,
        num_warps=4,
        num_stages=3,
    )

    return c


def parse_args() -> argparse.Namespace:
    parser = argparse.ArgumentParser(description="Basic educational matmul in Triton.")
    parser.add_argument("--m", type=int, default=512)
    parser.add_argument("--n", type=int, default=512)
    parser.add_argument("--k", type=int, default=512)
    parser.add_argument("--block-m", type=int, default=16)
    parser.add_argument("--block-n", type=int, default=16)
    parser.add_argument("--block-k", type=int, default=32)
    parser.add_argument("--seed", type=int, default=0)
    return parser.parse_args()


def main() -> None:
    args = parse_args()

    if not torch.cuda.is_available():
        raise RuntimeError("CUDA is not available.")

    torch.manual_seed(args.seed)

    a = torch.randn((args.m, args.k), device="cuda", dtype=torch.float16)
    b = torch.randn((args.k, args.n), device="cuda", dtype=torch.float16)

    c_torch = a @ b
    c_triton = matmul_triton(
        a,
        b,
        block_m=args.block_m,
        block_n=args.block_n,
        block_k=args.block_k,
    )

    torch.testing.assert_close(c_triton, c_torch, rtol=1e-2, atol=1e-2)

    print(c_triton[:2, :8])
    print("OK")


if __name__ == "__main__":
    main()

Run:

python matmul_basic.py

Again, this is educational. A serious matmul kernel needs:

better block sizes
autotuning
better program ordering
hardware-specific configs
possibly fused activation
possibly FP8 or quantized paths

The official matmul tutorial goes deeper, including grouped ordering to improve L2 cache behavior and autotuning over multiple configurations.

5.3 Understanding matmul pointer arithmetic

This line:

a_ptrs = a_ptr + offs_m[:, None] * stride_am + k_offsets[None, :] * stride_ak

builds a block of pointers. If:

offs_m has shape [BLOCK_M]
k_offsets has shape [BLOCK_K]

then:

offs_m[:, None]

has shape:

[BLOCK_M, 1]

and:

k_offsets[None, :]

has shape:

[1, BLOCK_K]

Broadcasting gives:

[BLOCK_M, BLOCK_K]

So a_ptrs points to a tile of A. Similarly:

b_ptrs = b_ptr + k_offsets[:, None] * stride_bk + offs_n[None, :] * stride_bn

points to a tile of B with shape:

[BLOCK_K, BLOCK_N]

Then:

acc = tl.dot(a, b, acc)

computes:

[BLOCK_M, BLOCK_K] @ [BLOCK_K, BLOCK_N]
=
[BLOCK_M, BLOCK_N]

This is the heart of tiled matmul.

5.4 What the compiler does for you

Your Triton code says:

a = tl.load(...)
b = tl.load(...)
acc = tl.dot(a, b, acc)
tl.store(...)

The compiler has to decide many low-level details:

how to map block operations to GPU instructions
how to use registers
how to use shared memory
how to coalesce memory accesses
how many warps execute one program
how to schedule loads and compute
how to use tensor cores for dot products

The compiler uses block-level data-flow analysis and applies optimizations such as coalescing, vectorization, tensor-core-aware instruction selection, shared-memory allocation/synchronization, and asynchronous copy scheduling. This is why Triton can feel much higher-level than CUDA while still being low-level enough to matter.

5.5 Autotuning

The annoying truth:

The best BLOCK_M, BLOCK_N, BLOCK_K, num_warps, and num_stages depend on the GPU and tensor shapes.

So Triton provides autotuning. The triton.autotune decorator tries multiple configurations and selects a good one for a given key. A small example:

import triton
import triton.language as tl


@triton.autotune(
    configs=[
        # Try a few static code shapes for the same logical kernel.
        triton.Config({"BLOCK_SIZE": 256}, num_warps=4),
        triton.Config({"BLOCK_SIZE": 512}, num_warps=4),
        triton.Config({"BLOCK_SIZE": 1024}, num_warps=8),
    ],
    # Retune when this runtime argument changes.
    key=["n_elements"],
)
@triton.jit
def add_kernel_autotuned(
    x_ptr,
    y_ptr,
    out_ptr,
    n_elements,
    BLOCK_SIZE: tl.constexpr,
):
    pid = tl.program_id(axis=0)
    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements

    x = tl.load(x_ptr + offsets, mask=mask, other=0.0)
    y = tl.load(y_ptr + offsets, mask=mask, other=0.0)

    tl.store(out_ptr + offsets, x + y, mask=mask)

The key idea:

For each n_elements value, try candidate configs.
Benchmark them.
Cache the best one.
Reuse it later.

Autotuning is powerful, but it can surprise you because the kernel may run multiple times during tuning. Be careful with kernels that mutate inputs or have side effects.

5.6 Heuristics

Sometimes autotuning is too expensive. You may just want:

BLOCK_SIZE = next power of two of n_cols

Triton has triton.heuristics for this. It computes meta-parameter values before compilation, which is useful when autotuning is too expensive or not applicable. Example sketch:

import triton
import triton.language as tl


@triton.heuristics(
    values={
        # Derive BLOCK_SIZE from the runtime n_cols argument.
        "BLOCK_SIZE": lambda args: triton.next_power_of_2(args["n_cols"]),
    }
)
@triton.jit
def kernel_with_heuristic(
    x_ptr,
    out_ptr,
    n_cols,
    BLOCK_SIZE: tl.constexpr,
):
    offsets = tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_cols
    x = tl.load(x_ptr + offsets, mask=mask, other=0.0)
    tl.store(out_ptr + offsets, x, mask=mask)

Beginner rule:

Use heuristics for obvious shape-derived choices.
Use autotuning for performance-sensitive choices.

6. Triton semantics and memory habits

6.1 Triton syntax looks like NumPy, but it is not NumPy

Triton code often looks like:

offsets = tl.arange(0, BLOCK_SIZE)
x = tl.load(x_ptr + offsets)
y = x * 2.0 + 1.0

This feels like NumPy. But it is compiled GPU code. Triton mostly follows NumPy semantics, with exceptions, including behavior around type promotion and broadcasting. A good mental model:

NumPy syntax style
GPU execution model
compiler-controlled lowering

So this:

x = tl.load(x_ptr + offsets)
y = x + 1

does not create a Python array. It creates GPU code operating on a block tensor.

6.2 Masks are not optional

Whenever your block can go out of bounds, use a mask. Example:

mask = offsets < n_elements
x = tl.load(x_ptr + offsets, mask=mask, other=0.0)
tl.store(out_ptr + offsets, x, mask=mask)

Without a mask, the final program may access invalid memory. This is one of the most important Triton habits. Bad:

x = tl.load(x_ptr + offsets)

Good:

x = tl.load(x_ptr + offsets, mask=offsets < n_elements, other=0.0)

For stores:

tl.store(out_ptr + offsets, y, mask=offsets < n_elements)

For loads, the other value is used for masked-out lanes. For softmax, we used:

other=-float("inf")

because masked-out values should not affect the maximum or denominator. For matmul, we used:

other=0.0

because masked-out values should not affect the sum.

6.3 Strides

PyTorch tensors are not always contiguous. A tensor element address is usually:

base_pointer + i * stride_0 + j * stride_1

For a row-major contiguous matrix:

shape = [M, N]
stride = [N, 1]

So:

x[i, j]

lives at:

x_ptr + i * N + j

In Triton, you often pass strides explicitly:

x.stride(0)
x.stride(1)

That is why the matmul kernel had:

stride_am
stride_ak
stride_bk
stride_bn
stride_cm
stride_cn

This makes kernels more general. But for beginner code, using .contiguous() is often simpler.

6.4 Memory coalescing

GPUs like consecutive memory access. Good:

offsets = block_start + tl.arange(0, BLOCK_SIZE)
x = tl.load(x_ptr + offsets)

This loads consecutive elements. Less good:

offsets = block_start + tl.arange(0, BLOCK_SIZE) * 17
x = tl.load(x_ptr + offsets)

This jumps through memory. A useful rule:

Consecutive offsets are friendly.
Strided or random offsets are harder.

Triton can optimize many patterns, but it cannot make arbitrary random memory access free.

6.5 When Triton is a good idea

Triton is a good fit for:

custom elementwise fusion
row-wise reductions
softmax-like kernels
layer norm and RMSNorm
custom matmul variants
matmul plus activation fusion
attention variants
quantization and dequantization kernels
custom indexing kernels
inference-specific kernels
research kernels that PyTorch does not provide

Triton is less likely to help when:

the operation is already a standard cuBLAS/cuDNN call
the tensor is tiny
the code is dominated by Python overhead outside the kernel
the operation is too irregular
you need full low-level CUDA control
you have not benchmarked the PyTorch baseline

A beginner mistake is to rewrite everything in Triton. A better approach:

Start with PyTorch.
Measure.
Try torch.compile.
Identify the bottleneck.
Write one Triton kernel for that bottleneck.
Measure again.

7. PyTorch integration, autograd, and debugging

7.1 Using Triton with torch.compile

Here is a tiny example of a user-defined Triton kernel inside a compiled PyTorch function. Create compile_triton_add.py:

#!/usr/bin/env python3
import argparse

import torch
import triton
import triton.language as tl


@triton.jit
def add_kernel(
    x_ptr,
    y_ptr,
    out_ptr,
    n_elements,
    BLOCK_SIZE: tl.constexpr,
):
    pid = tl.program_id(axis=0)

    # Same vector-add pattern: one program owns one block.
    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements

    x = tl.load(x_ptr + offsets, mask=mask, other=0.0)
    y = tl.load(y_ptr + offsets, mask=mask, other=0.0)

    tl.store(out_ptr + offsets, x + y, mask=mask)


@torch.compile(fullgraph=True)
def compiled_add(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
    out = torch.empty_like(x)
    n_elements = out.numel()

    # The compiled PyTorch graph contains a Triton kernel launch.
    grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),)

    add_kernel[grid](
        x,
        y,
        out,
        n_elements,
        BLOCK_SIZE=1024,
    )

    return out


def parse_args() -> argparse.Namespace:
    parser = argparse.ArgumentParser(description="Use a Triton kernel with torch.compile.")
    parser.add_argument("--size", type=int, default=1_000_000)
    return parser.parse_args()


def main() -> None:
    args = parse_args()

    if not torch.cuda.is_available():
        raise RuntimeError("CUDA is not available.")

    x = torch.randn(args.size, device="cuda")
    y = torch.randn(args.size, device="cuda")

    out = compiled_add(x, y)

    torch.testing.assert_close(out, x + y)

    print(out[:8])
    print("OK")


if __name__ == "__main__":
    main()

Run:

python compile_triton_add.py

The PyTorch tutorial on user-defined Triton kernels shows this same broad pattern and notes that custom Triton kernels can optimize specific parts of a model’s computation. It also discusses more structured integration through torch.library.triton_op for better composability with PyTorch subsystems.

7.2 Autograd: Triton does not automatically give gradients

If you write:

out = my_triton_kernel(x)
loss = out.sum()
loss.backward()

PyTorch does not magically know the derivative of your custom Triton kernel. You usually need one of these:

a torch.autograd.Function
a torch.library custom op with autograd registration
a backward implemented in normal PyTorch
a backward implemented in Triton

A small educational example:

#!/usr/bin/env python3
import argparse

import torch
import triton
import triton.language as tl


@triton.jit
def square_forward_kernel(
    x_ptr,
    out_ptr,
    n_elements,
    BLOCK_SIZE: tl.constexpr,
):
    pid = tl.program_id(axis=0)

    # Forward pass: y = x^2 for one block.
    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements

    x = tl.load(x_ptr + offsets, mask=mask, other=0.0)
    tl.store(out_ptr + offsets, x * x, mask=mask)


@triton.jit
def square_backward_kernel(
    x_ptr,
    grad_out_ptr,
    grad_x_ptr,
    n_elements,
    BLOCK_SIZE: tl.constexpr,
):
    pid = tl.program_id(axis=0)

    # Backward pass: d(x^2)/dx = 2x, multiplied by upstream grad.
    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements

    x = tl.load(x_ptr + offsets, mask=mask, other=0.0)
    grad_out = tl.load(grad_out_ptr + offsets, mask=mask, other=0.0)

    grad_x = 2.0 * x * grad_out

    tl.store(grad_x_ptr + offsets, grad_x, mask=mask)


def square_forward_triton(x: torch.Tensor, block_size: int) -> torch.Tensor:
    out = torch.empty_like(x)
    n_elements = x.numel()
    grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),)

    square_forward_kernel[grid](
        x,
        out,
        n_elements,
        BLOCK_SIZE=block_size,
    )

    return out


def square_backward_triton(
    x: torch.Tensor,
    grad_out: torch.Tensor,
    block_size: int,
) -> torch.Tensor:
    grad_x = torch.empty_like(x)
    n_elements = x.numel()
    grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),)

    square_backward_kernel[grid](
        x,
        grad_out,
        grad_x,
        n_elements,
        BLOCK_SIZE=block_size,
    )

    return grad_x


class TritonSquare(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x: torch.Tensor) -> torch.Tensor:
        # Save x so the backward pass can compute 2 * x * grad_out.
        ctx.save_for_backward(x)
        return square_forward_triton(x, block_size=1024)

    @staticmethod
    def backward(ctx, grad_out: torch.Tensor) -> tuple[torch.Tensor]:
        (x,) = ctx.saved_tensors
        grad_x = square_backward_triton(
            x,
            grad_out.contiguous(),
            block_size=1024,
        )
        return (grad_x,)


def parse_args() -> argparse.Namespace:
    parser = argparse.ArgumentParser(description="Tiny Triton autograd demo.")
    parser.add_argument("--size", type=int, default=1024)
    return parser.parse_args()


def main() -> None:
    args = parse_args()

    if not torch.cuda.is_available():
        raise RuntimeError("CUDA is not available.")

    x = torch.randn(args.size, device="cuda", requires_grad=True)

    y = TritonSquare.apply(x)
    loss = y.sum()
    loss.backward()

    torch.testing.assert_close(y, x.detach() * x.detach())
    torch.testing.assert_close(x.grad, 2.0 * x.detach())

    print(y[:8])
    print(x.grad[:8])
    print("OK")


if __name__ == "__main__":
    main()

This example is intentionally simple. The important point:

Forward kernel and backward kernel are separate pieces of work.

7.3 Debugging Triton

Triton kernels can fail in confusing ways:

wrong output
out-of-bounds memory access
compilation error
invalid constexpr use
dtype mismatch
mask bug
bad stride
unexpected non-contiguous tensor

Useful debugging tools:

torch.testing.assert_close
small tensor sizes
print intermediate values in Python wrappers
TRITON_INTERPRET=1
tl.static_print
tl.static_assert
tl.device_print
tl.device_assert
NVIDIA compute-sanitizer

Triton has compile-time debugging operators such as static_print and static_assert, runtime debugging operators such as device_print and device_assert, and an interpreter mode enabled with TRITON_INTERPRET=1, which simulates Triton programs on CPU using NumPy equivalents. Example:

TRITON_INTERPRET=1 python triton_add.py --size 16 --block-size 8

Another useful pattern:

torch.testing.assert_close(out_triton, out_torch, rtol=1e-4, atol=1e-4)

Do this before benchmarking. Correctness first. Performance second.

8. Practical workflow and final intuition

8.1 Common beginner mistakes

36.1 Forgetting masks

Bad:

x = tl.load(x_ptr + offsets)

Good:

mask = offsets < n_elements
x = tl.load(x_ptr + offsets, mask=mask, other=0.0)

36.2 Benchmarking first-call compilation

Bad:

time one call
declare Triton slow

Good:

warm up
benchmark repeated calls
ignore compilation time unless compilation time is the thing you care about

36.3 Comparing against the wrong PyTorch baseline

Bad:

torch_baseline = slow_python_loop(x)

Good:

torch_baseline = real_pytorch_operator(x)

Also test:

torch.compile(...)

before writing a custom kernel.

36.4 Using Triton for tiny tensors

For tiny tensors, kernel launch overhead can dominate. Triton is often most useful when each launch does enough work.

36.5 Ignoring contiguity and strides

This can break correctness:

x = x.transpose(0, 1)
out = my_kernel_that_assumes_contiguous(x)

Either call:

x = x.contiguous()

or pass and use strides correctly.

36.6 Thinking `tl.arange` is a Python list

It is not. It is a Triton block tensor used by the compiler.

36.7 Forgetting that `tl.constexpr` changes specialization

If you call the same kernel with many different BLOCK_SIZE values, Triton may compile many specializations.

8.2 A practical workflow for writing a Triton kernel

Use this loop:

1. Write the PyTorch reference.
2. Write a tiny Triton kernel.
3. Test on small shapes.
4. Test weird shapes.
5. Test non-multiple-of-block-size shapes.
6. Test dtype behavior.
7. Compare with torch.testing.assert_close.
8. Benchmark with warmup.
9. Tune BLOCK_SIZE, num_warps, num_stages.
10. Try autotune.
11. Compare against torch.compile.
12. Keep the Triton version only if it actually helps.

Example shape tests:

size = 1
size = 7
size = 1024
size = 1025
size = 1_000_000

For matrix kernels:

M, N, K all multiples of block sizes
M not multiple
N not multiple
K not multiple
skinny matrix
wide matrix
square matrix

8.3 How Triton works, in one final diagram

Python program
  |
  v
PyTorch tensors on GPU
  |
  v
Triton kernel launch
  |
  v
@triton.jit function
  |
  v
grid of program instances
  |
  v
program 0: block/tile of data
program 1: block/tile of data
program 2: block/tile of data
...
  |
  v
tl.load
  |
  v
block tensor math
  |
  v
tl.store
  |
  v
output tensor

The key loop in your head should be:

Which program am I?
Which block of data do I own?
Which offsets do I load?
Which masks protect the boundary?
What computation do I do?
Where do I store?

For vector add:

program owns a slice of a vector

For softmax:

program owns a row

For matmul:

program owns a tile of C

For attention:

program owns a tile of the attention computation

8.4 The core intuition to remember

Triton is not “Python that magically runs fast on GPU.” It is a way to write GPU kernels with this abstraction:

one program instance handles one block of data

Your job is to choose:

how big the block is
how many programs run
which memory each program loads
which masks protect the edges
which computation happens on-chip
which result gets stored

The compiler’s job is to turn that block-level program into efficient GPU code. So the real Triton optimization problem is:

Given:
  - a tensor operation,
  - limited memory bandwidth,
  - GPU parallelism,
  - registers,
  - shared memory,
  - tensor cores,
  - launch overhead,
  - shape variability,

choose:
  - block sizes,
  - program grid,
  - memory layout,
  - masks,
  - reductions,
  - num_warps,
  - num_stages,
  - autotune configs,

so that:
  - memory traffic is low,
  - data access is friendly,
  - the GPU stays busy,
  - correctness holds for edge shapes,
  - performance beats the baseline.

That is Triton.

References

OpenAI Triton announcement: https://openai.com/index/triton/
NVIDIA Triton Inference Server documentation: https://docs.nvidia.com/deeplearning/triton-inference-server/user-guide/docs/python_backend/README.html
Triton programming guide: https://triton-lang.org/main/programming-guide/chapter-1/introduction.html
Triton installation docs: https://triton-lang.org/main/getting-started/installation.html
Triton vector-add tutorial: https://triton-lang.org/main/getting-started/tutorials/01-vector-add.html
Triton fused-softmax tutorial: https://triton-lang.org/main/getting-started/tutorials/02-fused-softmax.html
Triton matrix multiplication tutorial: https://triton-lang.org/main/getting-started/tutorials/03-matrix-multiplication.html
Triton JIT API docs: https://triton-lang.org/main/python-api/generated/triton.jit.html
Triton benchmarking API docs: https://triton-lang.org/main/python-api/generated/triton.testing.do_bench.html
Triton autotune API docs: https://triton-lang.org/main/python-api/generated/triton.autotune.html
Triton heuristics API docs: https://triton-lang.org/main/python-api/generated/triton.heuristics.html
Triton semantics docs: https://triton-lang.org/main/python-api/triton-semantics.html
Triton debugging guide: https://triton-lang.org/main/programming-guide/chapter-3/debugging.html
PyTorch tutorial on user-defined Triton kernels: https://docs.pytorch.org/tutorials/recipes/torch_compile_user_defined_triton_kernel_tutorial.html

0. The core idea

1. Why Triton exists

1.1 The PyTorch baseline and the fusion problem

1.2 Triton vs CUDA vs PyTorch vs torch.compile

1.3 The core mental model

1.4 From Python to GPU

2. Setup and the first kernel

2.1 Installation

2.2 First PyTorch baseline

2.3 First Triton kernel: vector addition

2.4 Understanding the kernel line by line

2.5 A Triton program is not a CUDA thread

2.6 BLOCK_SIZE, tl.constexpr, and the grid

2.7 A tiny CPU mental simulator

3. Benchmarking and fusion

3.1 Benchmarking correctly

3.2 Kernel fusion: where Triton starts to matter

3.3 Memory traffic is often the real enemy

4. Reductions and row-wise softmax

4.1 Row-wise softmax

4.2 Why softmax teaches more than vector add

4.3 Reductions in Triton

4.4 num_warps and num_stages

5. Matrix multiplication and tuning

5.1 Matrix multiplication: the big example

5.2 A small educational matmul kernel

5.3 Understanding matmul pointer arithmetic

5.4 What the compiler does for you

5.5 Autotuning

5.6 Heuristics

6. Triton semantics and memory habits

6.1 Triton syntax looks like NumPy, but it is not NumPy

6.2 Masks are not optional

6.3 Strides

6.4 Memory coalescing

6.5 When Triton is a good idea

7. PyTorch integration, autograd, and debugging

7.1 Using Triton with torch.compile

7.2 Autograd: Triton does not automatically give gradients

7.3 Debugging Triton

8. Practical workflow and final intuition

8.1 Common beginner mistakes

36.1 Forgetting masks

36.2 Benchmarking first-call compilation

36.3 Comparing against the wrong PyTorch baseline

36.4 Using Triton for tiny tensors

36.5 Ignoring contiguity and strides

36.6 Thinking tl.arange is a Python list

36.7 Forgetting that tl.constexpr changes specialization

8.2 A practical workflow for writing a Triton kernel

8.3 How Triton works, in one final diagram

8.4 The core intuition to remember

References

36.6 Thinking `tl.arange` is a Python list

36.7 Forgetting that `tl.constexpr` changes specialization