Triton vs CUDA — where each one wins
Triton buys you productivity and portability; hand-written CUDA buys you the last mile. A pragmatic take on when to reach for which.
The Triton-vs-CUDA debate usually generates more heat than light, because the honest answer is “it depends on how close to the metal you need to be.” Here’s a working heuristic.
Reach for Triton when
- You’re writing fused elementwise, reduction, or attention-style kernels.
- You want autotuning over block sizes and pipelining without writing it yourself.
- You value one kernel that stays fast across architectures.
Triton’s block-pointer model and software pipelining will emit cp.async and, on Hopper/Blackwell, TMA under the hood — so you get a lot of the modern memory-movement machinery for free.
Reach for CUDA C++ when
- You need explicit control over shared-memory layout, warp specialization, or tensor-core fragment scheduling.
- You’re squeezing the last 10–15% out of a kernel that already dominates your runtime.
- You need a memory-movement pattern Triton won’t express.
The part both share
Whichever you pick, the same trap applies: an “optimized” kernel that’s subtly wrong is worse than the slow one you started with. Kernalin verifies both paths through the same gauntlet, and its technique recommendations are backend-aware — vectorized_loads and make_block_ptr for Triton, vectorized_ldg_128bit and TMA staging for CUDA.
Pick the tool that fits the kernel. Verify the result the same way regardless.
This is sample content published to illustrate Kernalin’s engineering perspective.