HipKittens AMD Kernels: MI355X Performance in 2026

AMD's MI355X has had the memory bandwidth and the FLOPs on paper for a while. What it didn't have, until this year, was a kernel layer that got close to using them without an assembly expert on staff. HipKittens AMD kernels change that math: HipKittens shipped as an official AITER backend in March 2026, and the numbers behind that change are the actual argument for why AMD's cost-per-token story has started looking different from its spec sheet.

This post covers what changed, why the software stack (not the silicon) was the real reason AMD lost inference throughput comparisons for so long, where HipKittens' published gains land today, and what that does to the MI355X-versus-H200 rental decision. For the hardware specs behind CDNA4, see AMD MI350X vs NVIDIA B200.

What Changed: HipKittens AMD Kernels Land in AITER, MLPerf v6.0 Results

Two things happened close together this year, and only one of them got real attention. We'll take them in the order that matters for a buyer, not the order they made headlines.

The first is the MLPerf result already making the rounds: AMD submitted MI355X results to MLPerf Inference v6.0, the round MLCommons published on April 1, 2026. In the closed division, MI355X landed within 20% of H200 on LLaMA 2 70B and GPT-OSS 120B, numbers that are directly comparable since both ran under the same model and precision constraints. Some of that gap-closing is raw capacity: MI355X's 288 GB of HBM3e lets it run larger batch sizes than B200's 192 GB on workloads that fit inside it. But a real share of it is the kernel layer this post is about, since those MLPerf numbers depend on the same GEMM and GQA/MHA attention kernels HipKittens now ships inside AITER.

AITER, AMD's production kernel library, is the default inference backend both vLLM and SGLang use on AMD GPUs, spanning Triton, Composable Kernel, and hand-tuned assembly implementations across MHA, MLA, paged attention, fused MoE, and GEMM operators. Until this integration, the fastest path through AITER for a given operator was, more often than not, someone's hand-written assembly. Now some of that fastest path is HipKittens-authored tile code instead, meaning the framework's gains aren't a research paper sitting next to the production stack, they're inside it.

The MLPerf number is what a buyer sees first. The AITER integration is why that number is starting to be reproducible outside AMD's own submission runs.

Why the Software Stack Was AMD's Real Bottleneck, Not the Silicon

The HipKittens paper states the underlying problem plainly: "peak performance AMD kernels are written in raw assembly." Compiler-based alternatives haven't closed that gap on their own. Triton, according to the Hazy Research team, struggles with register lifetime tracking and lowering memory accesses to AMD's most performant intrinsics, and Mojo's MHA kernel lands around half of peak due to bank conflicts. That leaves a narrow set of engineers, the ones who can hand-write CDNA assembly, as the people standing between AMD's silicon specs and what a production inference stack actually gets to use. A GPU cloud buyer never sees that constraint on a spec sheet. They see it as MI355X quietly underperforming its paper numbers in whatever framework they happened to deploy with, and the H100 or H200 down the hall not doing that.

ROCm vs CUDA: Where the Gap Actually Lived

Framed as ROCm versus CUDA, this is the same story from a different angle. We covered the fuller framework-compatibility picture in our ROCm vs CUDA guide: standard PyTorch and vLLM inference on ROCm already reaches 90-95% of H100 throughput on MI300X and MI355X, so for the common path, the framework gap was mostly closed well before HipKittens shipped. Where ROCm kept losing ground was the tail: custom attention variants, anything that needed a hand-fused kernel, anything where the only fast implementation on CUDA was a library like TensorRT-LLM or FlashAttention 3 with no ROCm equivalent. That tail is exactly where "write it yourself in assembly, or don't get peak performance" bit hardest, and it's exactly the tail HipKittens targets. It doesn't replace ROCm. It replaces the assembly step for the specific operators, GEMM and attention chief among them, where assembly was previously the only way to get AMD's hardware to do what it's rated for.

Where HipKittens Wins Today: GEMM, GQA/MHA Attention, RoPE

HipKittens is a C++-embedded, tile-based programming framework, built as the AMD counterpart to Hazy Research's earlier NVIDIA framework, ThunderKittens. It targets CDNA3 (MI300X, MI325X) and CDNA4 (MI350X, MI355X) today. The high-level tile abstraction carries over from ThunderKittens largely unchanged, but the memory access patterns, the scheduling that overlaps compute with memory, and the way work gets ordered across processors were rebuilt for AMD's chiplet layout and wavefront model, which behave differently from NVIDIA's SM-based architecture at the level HipKittens operates.

The paper's own framing of the goal is worth stating directly: to build "a single, tile-based software layer for high-performance AI kernels that translates across GPU vendors." HipKittens isn't positioned as an AMD-only tool competing with CUDA. It's a bet that one tile abstraction, expressed once, should be able to hit near-assembly performance on both NVIDIA and AMD silicon, which is a materially different ambition than "port a CUDA kernel to HIP and hope."

The paper puts a number on how far that bet actually pays off. Validated across CDNA3 and CDNA4, HipKittens kernels compete with AMD's hand-optimized assembly on GEMM and attention overall, and in a specific subset of cases, d=64 attention, GQA backwards, and memory-bound kernels among them, they outperform every available baseline, assembly included, by 1.2 to 2.4x. That's the detail worth sitting with: in those cases assembly isn't just matched, it's beaten, because a hand-written kernel is tuned for one shape while a tile abstraction can be re-scheduled across the breadth of shapes real AI workloads produce.

GEMM and Memory-Bound Kernels

GEMM sat in the same assembly-only category as attention before this integration. AITER's and HipBLASLt's GEMM kernels on AMD hardware are written in raw assembly, the same constraint the paper documents for attention. That's the whole shift in one sentence: a few dozen lines of tile code standing in for what used to require an assembly rewrite, and in the memory-bound cases the paper measured, that tile code came out ahead of the assembly it replaced.

Attention Forward and Backward

Attention is where the software gap showed up most starkly, and it's where HipKittens' published numbers do the most work. HipKittens' attention kernels close that distance substantially, competing with hand-optimized assembly on both the forward and backward passes for GQA and MHA attention rather than sitting in that same underperforming range. That's what actually matters for anyone running GQA-based models (which is most current open-weight LLMs) on AMD hardware: the backward pass, the workload where AMD's own tooling struggled hardest before this integration, is now a kernel that AITER ships by default instead of a gap teams had to route around. GQA backwards specifically is one of the named cases above where HipKittens didn't just close the gap to assembly, it moved past it.

RoPE and the other attention-adjacent operators the outline flags follow the same pattern as GEMM: they're kernels where a correctly-scheduled tile implementation was always achievable in principle, and the barrier was that achieving it required an assembly-level rewrite that most teams weren't going to do for a single operator. Folding RoPE handling into a tile framework that's already shipping in AITER removes that barrier without requiring anyone downstream to write AMD assembly themselves.

What This Means for Renting MI355X vs H200 Right Now

Here's how we'd frame the direct answer: if your inference stack calls AITER's default kernels, or you're willing to write HipKittens kernels for what it doesn't cover yet, MI355X's software disadvantage against NVIDIA just got smaller in the specific places (GEMM, GQA/MHA attention) that decide most inference throughput. If your stack depends on CUDA-only tooling that has no ROCm equivalent, that gap is still there, and no kernel framework closes it from the AMD side.

Run the numbers on a like-for-like inference workload and the decision comes down to what you're actually paying for a GPU-hour, not just what the chip can theoretically do. Most of the cheaper MI355X quotes in circulation are spot or preemptible, not on-demand, so they're not a fair comparison point. On the NVIDIA side, H200 currently rents from Spheron at $4.96/hr dedicated on-demand, or $2.57/hr spot (reclaimable anytime). Even before factoring in a single throughput number, that's roughly half the hourly cost for the on-demand tier. HipKittens closing AITER's kernel gap makes MI355X's tokens-per-dollar story more credible than it was a year ago; it doesn't yet make MI355X the cheaper on-demand hour.

Pricing fluctuates based on GPU availability and provider. The prices above are current as of 31 Aug 2026 and may have changed. Check current GPU pricing → for live NVIDIA rates, and confirm AMD on-demand rates directly with whichever provider you're evaluating.

Worth being direct about the limit here: Spheron's own catalog is NVIDIA-only today, spanning H100, H200, B200, B300, A100, and the rest of the current lineup. If HipKittens' AITER integration convinces you MI355X is worth evaluating for a GQA-heavy inference workload, that's a real, evidence-backed reason to test it, and you'll need to test it on a provider that actually carries AMD Instinct hardware. What Spheron can speak to directly is the NVIDIA side of that decision: an H200 instance is a reasonable place to run the H100/H200-side baseline you'd want before deciding whether MI355X's narrower software gap is worth the migration.

For the kernel-authoring approach on the NVIDIA side of this same comparison, see our Triton kernel development guide, and for the training-side ROCm story on this same hardware, see training LLMs on AMD GPUs with ROCm.

HipKittens narrowed AMD's kernel gap, but the buying decision still comes down to what your stack actually calls and what the hour costs. Test the NVIDIA side of that comparison on Spheron with no long-term contract.

Rent H200 GPU →