triton

name	triton
description	Comprehensive reference documentation and skill for OpenAI Triton - a language and compiler for writing highly efficient custom Deep-Learning primitives on GPUs. Covers Python API (triton.language, triton.runtime, triton.compiler), MLIR dialects and passes, backends (NVIDIA CUDA, AMD ROCm/HIP), experimental features (Gluon, GSAN, triton_kernels), Proton profiler, tutorials, debugging, and build system.
version	3.7.0

Triton - GPU Programming Language & Compiler

Overview

Triton is a language and compiler for writing highly efficient custom Deep-Learning primitives. It provides an open-source environment to write fast code at higher productivity than CUDA, but also with higher flexibility than other existing DSLs.

Supported Hardware:

• NVIDIA GPUs (Compute Capability 8.0+)
• AMD GPUs (ROCm 6.2+)
• CPUs (under development)

Supported Platforms: Linux

Python Versions: CPython 3.10-3.14

Quick Reference

Installation

pip install triton

Minimal Kernel Example

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)
    output = x + y
    tl.store(output_ptr + offsets, output, mask=mask)

Documentation Structure

Core API

• 01-overview-and-architecture - Project architecture, compilation pipeline overview
• 02-getting-started - Installation, building from source, quick start
• 03-language-core - triton.language core: tensors, pointers, load/store, dtypes, dot, atomic ops
• 04-language-math - triton.language.math: exp, log, sin, cos, sqrt, etc.
• 05-language-random - triton.language.random: rand, randn, philox RNG
• 06-language-standard - triton.language.standard: reductions, sorting, softmax, cumsum
• 07-language-semantic - Semantic analysis layer

Runtime System

• 08-runtime-jit - @triton.jit decorator, JITFunction, specialization, caching
• 09-runtime-autotuner - @triton.autotune, Config, Heuristics, benchmarking
• 10-runtime-cache - Cache management, file cache, remote cache
• 11-runtime-driver - GPU driver abstraction, memory allocation
• 12-runtime-interpreter - TRITON_INTERPRET=1 interpreter mode

Compiler

• 13-compiler-pipeline - Full compilation pipeline: AST → TTIR → TTGIR → LLVMIR → PTX
• 14-compiler-codegen - AST-to-MLIR code generation, CodeGenerator class
• 15-compiler-errors - CompilationError, error handling

Backend System

• 16-backends - Backend architecture, BaseBackend, GPUTarget, driver/compiler split
• 17-nvidia-backend - NVIDIA CUDA backend specifics
• 18-amd-backend - AMD ROCm/HIP backend specifics

MLIR Dialects & Passes

• 19-mlir-dialects - Triton, TritonGPU, Gluon, TritonInstrument dialects
• 20-compilation-passes - All compilation passes: coalesce, pipeline, accelerate-matmul, etc.
• 21-memory-layouts - LinearLayout, blocked layouts, distributed/shared layouts

C++ Extension

• 22-cpp-extension - Python C++ bindings, pybind11 modules, ir.cc, passes.cc

Tutorials

• 23-tutorials - All 11 tutorials with detailed explanations

Advanced / Experimental

• 24-triton-kernels - triton_kernels library: matmul, reduce, topk, swiglu, etc.
• 25-gluon - Gluon experimental language and compiler
• 26-proton-profiler - Proton profiling system
• 27-gsan - GPU Sanitizer (GSAN) for symmetric memory

Operations

• 28-debugging - Environment variables, knobs, IR dumping, reproducer
• 29-building - Build system, CMake, custom LLVM, cross-compilation
• 30-plugin-system - Out-of-tree plugin dialects and passes
• 31-testing - Test framework, lit tests, pytest

Code Examples

Official tutorials from the Triton repository, providing verified reference implementations for common kernel patterns:

Example	Key Concepts
01-vector-add.py	Basic kernel, tl.load/tl.store, block programming model, autotune
02-fused-softmax.py	Fused kernel, online softmax, numerical stability, masking
03-matrix-multiplication.py	Tiled matmul, shared memory via tl.dot, block pointers, tl.autotune
04-low-memory-dropout.py	Random number generation, memory-efficient dropout, tl.rand
05-layer-norm.py	Reduction patterns, layer normalization, variance computation
06-fused-attention.py	Flash Attention, tiled softmax backward, complex multi-loop kernel
07-extern-functions.py	Calling external C/CUDA functions from Triton, tl.extern
08-grouped-gemm.py	Batched/variable-size grouped GEMM, problem splitting
09-persistent-matmul.py	Persistent kernels, software pipelining, long-lived kernel launch
10-block-scaled-matmul.py	Block-scaled (MxFP4/FP8) matmul, quantized GEMM patterns
11-programmatic-dependent-launch.py	Programmatic dependent launch, kernel synchronization via launch configs

Key Concepts

Programming Model

Triton kernels execute on a grid of programs (thread blocks). Each program:

1. Computes its ID via tl.program_id(axis)
2. Computes memory offsets from the program ID
3. Loads data using tl.load() with pointers and masks
4. Computes results using tensor operations
5. Stores results using tl.store()

Compilation Pipeline

Python AST → TTIR (Triton IR) → TTGIR (Triton GPU IR) → LLVM IR → PTX/AMDGPU → CUBIN/HSACO

Type System

• Scalar types: int1, int8, int16, int32, int64, uint8, uint16, uint32, uint64, float16, bfloat16, float32, float64, float8e4nv, float8e5, etc.
• Pointer types: tl.pointer_type(element_ty, const=False)
• Tensor types: block-level tensors with implicit parallelism
• constexpr: Compile-time constants annotated with tl.constexpr

Key Decorators

• @triton.jit - Mark a function for JIT compilation
• @triton.autotune(configs, key) - Auto-tune kernel configurations
• @triton.heuristics(values) - Compute meta-parameters heuristically
• @triton.experimental.constexpr_function - Compile-time function evaluation