Triton Basics: Vector Add

Triton is a language for writing GPU kernels in Python syntax. Compared to CUDA, Triton handles many low-level details (shared memory, synchronization), letting you focus on the algorithm.

This chapter uses the simplest example—vector addition—to learn Triton’s core model.

SPMD Programming Model

Before writing code, understand Triton’s core idea: SPMD (Single Program, Multiple Data).

In short: the same kernel code runs in many “programs” in parallel, each handling a different chunk of data.

Suppose we add two vectors of length 256 with BLOCK_SIZE = 64. Triton launches 4 programs:

Input vector (N=256, BLOCK_SIZE=64):

┌────────────┬────────────┬────────────┬────────────┐
│  0 ... 63  │ 64 ... 127 │ 128 .. 191 │ 192 .. 255 │
├────────────┼────────────┼────────────┼────────────┤
│ Program 0  │ Program 1  │ Program 2  │ Program 3  │
└────────────┴────────────┴────────────┴────────────┘

Each program handles its own block. How does a program know which block it owns? Use tl.program_id().