← All questions
Inference & ServingHard· Triton· Kernel Optimization· Softmax

Sketch a fused softmax kernel in Triton.

To sketch a fused softmax in Triton, you define a row-per-block grid where each program ID (pid = tl.program_id(0)) handles a specific row. You compute memory offsets using tl.arange and perform masked loads with other=-inf to handle padding. Stability is ensured by subtracting the row-wise maximum in-register before applying tl.exp and tl.sum. Finally, the values are normalized and written back via a masked store. The Triton compiler manages thread-level parallelism and memory coalescing automatically.

Code playground (Python)