Skip to content

Glossary

Short definitions for words used across the lessons, labs and case studies. Each entry says where you first need it, so you can skip terms you have not reached yet. Lesson numbers refer to the beginner course.

Execution model

Term Meaning First needed
Host / device The CPU running Python / the GPU running kernels Lesson 0
Kernel A function compiled to run on the GPU (@triton.jit or CUDA __global__) Lesson 0
Launch Submitting a kernel with a grid: kernel[grid](args) in Triton, kernel<<<blocks, threads>>>(args) in CUDA Lesson 0
Grid How many kernel instances run: Triton programs or CUDA thread blocks Lesson 1
Program (Triton) One kernel instance. It processes a whole block of values; tl.program_id(0) says which one Lesson 1
Block (Triton) The vector of values one program handles, e.g. tl.arange(0, BLOCK). Not the same as a CUDA thread block Lesson 1
Thread / thread block (CUDA) One scalar worker / a group of threads that can share memory and use __syncthreads() Lesson 9
Warp / lane 32 NVIDIA threads that execute together / one thread's position (0-31) inside its warp Lesson 9
Asynchronous launch Python returns before the GPU finishes, so a plain stopwatch measures submission, not work Lesson 0
Stream An ordered GPU work queue. Kernels on one stream run in submission order Workshop A2

Memory and layout

Term Meaning First needed
Offset Distance in elements (not bytes) from a tensor's first element Lesson 1
Mask / tail Per-lane true/false saying whether a load or store may touch memory / the final, partly filled block Lesson 1
Stride Elements to skip to move one step along a dimension. Shape alone cannot locate a row Lesson 2
Contiguous Rows stored back to back with column stride 1, so the strides follow from the shape Lesson 2
View A tensor sharing another tensor's storage with different shape/strides, e.g. a slice or transpose Lesson 2
Global memory (DRAM) The GPU's main memory, 16 GiB on the RTX 4080 SUPER; about 736 GB/s peak Lesson 4
L2 cache 64 MiB on-chip cache in front of DRAM. Data used again soon may come from L2, much faster than DRAM Lesson 7
Shared memory Fast memory shared by the threads of one CUDA block, e.g. scratch in the CUDA softmax Lesson 9
Registers / spill Each thread's fastest private storage / values moved to slower memory when registers run out Lesson 9
Coalescing Neighboring threads reading neighboring addresses, so the hardware combines their accesses Lesson 2
Occupancy How many warps an SM (streaming multiprocessor) keeps resident. Higher is not automatically faster Lab 2

Numbers and precision

Term Meaning First needed
dtype Element format: FP32 (4 bytes), FP16 and BF16 (2 bytes; BF16 keeps FP32's range with less precision) Lesson 0
Accumulation dtype The precision used while combining values. Here reductions use FP32 even for FP16 inputs Lesson 3
Reduction / identity Combining many values into one (sum, max) / the padding value that changes nothing (0 for sum, -inf for max) Lesson 3
Stable softmax Subtracting the row maximum before exp so large inputs cannot overflow Lesson 4
Oracle / reference An independent, simpler computation (often in FP64) that a kernel's output is compared against Lesson 4
Tolerance (atol, rtol) Allowed absolute and relative difference from the oracle; rounding makes exact equality unrealistic Lesson 8
Epsilon (eps) Small positive number added before a square root so an all-zero row does not divide by zero Lesson 5

Compiling Triton

Term Meaning First needed
JIT compilation Triton compiles a kernel the first time it is called with new settings. That first call can take a second or more Lesson 8
tl.constexpr A value fixed at compile time. Each new value compiles a new kernel. Use it for tile sizes like BLOCK, and for sizes a model fixes, such as row width Lesson 8
Runtime argument An ordinary integer or pointer argument. It can change per call without recompiling (sizes, strides) Lesson 8
Specialization Compiling a version tuned to known values. Triton also specializes runtime ints equal to 1 or divisible by 16 Lesson 8
Autotuning Timing several tile configurations on the first call for a shape and keeping the fastest Lab 5
Fusion Doing several steps in one kernel so intermediates never go through global memory Lesson 4
Tensor Core / mma.sync Matrix-multiply hardware units / the PTX instruction that uses them Lesson 6
TTIR, TTGIR, LLVM IR, PTX, SASS Successive compiler stages, from Triton's own IR down to NVIDIA's final machine code Lab 6

Measuring

Term Meaning First needed
Speedup Baseline time / candidate time. Above 1 means the candidate is faster Lesson 7
Median Middle sample; less sensitive to one slow outlier than the mean Lesson 7
CUDA events GPU-side timestamps; the elapsed time between two events measures device work Lesson 7
CUDA graph / graph replay A recorded sequence of launches replayed with almost no Python cost. The default benchmark mode Lesson 7
Event timing mode --timing events: ordinary Python calls, so host dispatch gaps are included Lesson 7
Warm vs cold cache Warm: inputs likely already in L2 from the previous call. Cold (--cache cold): L2 flushed first, so data comes from DRAM Lesson 7
GPU clocks / boost The GPU lowers its clock speed when idle and raises it under load. The climb takes a few hundred milliseconds, so a timing taken right after a pause can run several times slower Lesson 7
GPU warm-up A burst of compute-bound work (large matrix multiplies) run before timing so the GPU is at boost clocks. Separate from JIT warm-up, which only compiles Lesson 7
Regret How much slower a rule's choice is than the best choice for the same shape (time / best time - 1). Used to compare tuning rules over a grid of shapes Case studies
Logical GB/s Minimum useful bytes / time. A model, not a hardware counter; above DRAM peak means the data came from cache Lesson 7
TFLOP/s 2*M*N*K / seconds for GEMM: trillions of floating-point operations per second Lab 5
Arithmetic intensity Operations per byte moved. Low intensity means memory-bound Lesson 6
Amdahl's law Speeding up part of a program only helps in proportion to that part's share of the runtime Lesson 7
Eager mode Normal PyTorch: each operation launches its own kernel(s) immediately Lesson 4
torch.compile / Inductor PyTorch's compiler; Inductor generates fused (often Triton) kernels. A strong baseline for fusion claims Lesson 5
cuBLAS / SDPA NVIDIA's tuned matrix-multiply library / PyTorch's scaled_dot_product_attention, which selects a fused backend Lesson 6 / Workshop A3

Algorithms in the workshops

Term Meaning First needed
Jacobian / VJP Matrix of all partial derivatives / vector-Jacobian product: the gradient a backward pass actually needs Workshop I2
Autograd PyTorch's automatic differentiation. Every portfolio operator registers a backward pass with it: Triton kernels for softmax, residual RMSNorm and GEMM, exact formulas for add and row sum Workshop I2
Epilogue Work applied to a GEMM result before storing it, e.g. bias and ReLU Workshop I3
Grouped tile ordering Launching output tiles so neighbors reuse cached inputs; GROUP in the fused GEMM Workshop I3
Online normalizer Mergeable (max, sum) state that lets softmax statistics be computed chunk by chunk Workshop A1
Scratch Temporary GPU buffers a multi-kernel operation allocates for intermediate results Workshop A2
Custom op A function registered with PyTorch (torch.library) so torch.compile can trace it; see kernel_portfolio/library.py Setup
Causal mask Attention rule that a query may only see keys at the same or earlier positions Workshop A3

Tools

Term Meaning First needed
WSL Windows Subsystem for Linux: Ubuntu running inside Windows. One of the two GPU options here Setup
triton-windows Community build of Triton for native Windows; the PowerShell GPU option uses it Setup
Virtual environment An isolated Python install with its own packages (.venv-win in PowerShell, .venv in WSL); activate it in each new terminal Setup
Triton interpreter TRITON_INTERPRET=1 runs Triton kernels on the CPU with NumPy, for debugging and CI Setup
Compute Sanitizer NVIDIA tool: memcheck finds out-of-bounds accesses, racecheck finds shared-memory races Lesson 9
Nsight Systems / Compute NVIDIA profilers: a timeline of launches / per-kernel hardware counters Profiling