Kernel portability across NVIDIA and AMD

In one paragraph

Most AI kernels are written for NVIDIA GPUs first. AMD's ROCm stack and its HIP language mirror much of CUDA, and several newer languages aim to make kernels portable across vendors. Portable source does not mean portable performance: the hardware differs enough that a kernel tuned for one vendor usually needs retuning for the other.

How it works

• HIP. Mirrors much of the CUDA API; tools exist to translate CUDA source. Not a drop-in replacement.
• Hardware differences. AMD data-centre GPUs schedule 64-thread wavefronts rather than 32-thread warps, and differ in memory hierarchy and matrix instructions.
• Portable layers. ThunderKittens (tile primitives) and its AMD port HipKittens; TileLang (tile-level language with several backends); Mojo (a language aimed at portable high-performance code); Triton also has an AMD backend.

Design choices and trade-offs

• Engine support first. For serving, check whether the chosen engine and its attention kernels support the target AMD GPU before considering custom kernels.
• Portability vs peak speed. Portable code is cheaper to maintain across vendors; vendor-specific code reaches peak performance.

When not to use it

When the deployment will only ever run on one vendor's hardware.

How we evaluate it

The same correctness tests and benchmarks on each target, with results reported per vendor and never mixed.

Sources

• AMD: what is HIP (opens in a new tab)
• ThunderKittens (opens in a new tab), HipKittens (opens in a new tab), TileLang (opens in a new tab)