Automated CUDA is almost here. Revenue and capability per GPU are doubling.

Use this section as the reference map. Start with the role path closest to you, read the practical TL;DR for the spine, then come back to the full index when you need a specific layer. If you only read one dense technical section, read §15: it is the workload-to-stack map that ties the whole post together.

The index is the map. The post itself starts with the people and the event, because the evidence loop came out of a real weekend with real collaborators, not a detached thesis.

before:
  buy GPU servers
  wire networking
  tune serving later
  discover power, cooling, cache, and reliability limits in production

after:
  deploy an integrated NVL72-class rack
  validate compute, NVLink/NVSwitch, power, cooling, diagnostics, and serviceability as one unit
  then prove the workload path with real agent/RAG/rollout traces

Cerebral Valley and SHACK15 hosted the OpenEnv Hackathon in San Francisco on March 7–8. As far as we know it was the first public event built squarely around RL environments and post-training, and one of the better-run hackathons we've been to. Thank you to Ray Del Vecchio and the Cerebral Valley team for putting it together. The OpenEnv framework itself came from Meta-PyTorch, so thank you to Joseph Spisak, Emre Guven, Hamid Shojanazeri, and Sanyam Bhutani, who drove the PyTorch side of the collaboration. Hugging Face wired OpenEnv into Spaces and TRL, with Ben Burtenshaw closely involved. Unsloth AI carried the efficient RL training stack a lot of teams leaned on over the weekend, ours included, so thank you to Daniel Han and Michael Han. The track our team entered was run by Mercor, which deserves a paragraph of its own.

A particular thank-you, then, to Mercor. It helps to know what Mercor does, because it makes their interest in a hackathon like this less of a coincidence. Mercor connects human expertise with the labs building frontier models, and they've been unusually open about the research behind it. Their APEX benchmarks, the AI Productivity Index, measure whether models can do economically valuable professional work. APEX-Agents extends that into long-horizon, cross-application agentic tasks across investment banking, consulting, and corporate law: 480 expert-built tasks released openly on Hugging Face, with ACE covering consumer applications alongside it. Archipelago, their evaluation harness, is open-source too: a sandboxed environment, an agent runner, and a grading system that scores agent trajectories against expert rubrics. And their post-training work with Applied Compute is the part that stuck with me. Fewer than a thousand expert-labeled tasks measurably moved an open model, with the gains traced through full trajectory-level observability rather than aggregate scores.

That last detail is why the recognition meant something to us. Mercor keeps landing on the same point this post does. The RL signal is only as honest as the observability underneath it. They describe a near future where the economy becomes, in their words, "an RL environment machine". From their angle, a hackathon track on RL environments isn't a side quest. It's the thing they think the next decade of AI training will actually run on. Anirudh Ravichandran, a Tech Lead Manager at Mercor, described building RL environments to optimize CUDA kernels as "one of the most fundamental and high-leverage applications" of the recursive self-improvement loop, and credited the project's positioning, execution, and design. That's generous of him, and where he's pointing is basically the bet the rest of this post is about. Mercor was also generous with a $10,000 prize. We're grateful for the recognition, and more grateful still that it came from a team thinking carefully, and publicly, about the same problem we are. (Northeastern's writeup is the third-party version of the result, and the source of that quote.)

Compute and infrastructure made the whole thing runnable, so thank you to Matthew Lu and Jacob Feldman at CoreWeave, and to Will Stewart at Northflank. And thanks to the wider group of organizations that sponsored, judged, mentored, or otherwise turned up to support the event, among them UC Berkeley's Sky Computing Lab, Fleet AI, Snorkel AI, Patronus AI, Halluminate, Scale AI, OpenPipe, Cursor, and Scaler AI Labs. This isn't a complete list. Plenty of people contributed quietly. But it's the set we can name with confidence.

A word on the team, because it matters to us to get the credit right. Team Automate-CUDA came together for the weekend. Our founder, William Chen, led it, in close collaboration with three teammates we're grateful to. Yiying Xie is pursuing her Master's in Computer Science at Northeastern's Khoury College. Warren Low and Farhan Navas are both from the National University of Singapore, currently on exchange at Stanford, and brought tremendous kernel-engineering expertise to the two days. It was a real team effort: very little sleep, a working environment by the end of it, and we'd happily do it again next weekend. Before the argument gets technical, the next section defines the few terms the rest of the post keeps reusing.

If the rest of this post is opaque, the next definitions are the minimum you need to follow it. Engineers can jump to §02; everyone else, this is the spine, and we'll lean on it the whole way through.

One sentence that ties it together. A model writes a kernel in CUDA; a compiler lowers it through PTX or LLVM IR into machine code; an inference engine schedules those kernels around a KV cache through prefill and decode; and a harness measures whether any of it actually got faster, cheaper, and more energy-efficient. The same discipline then rolls forward into heterogeneous inference: cloud, edge, robotics, on-device, data centers, future ASICs, and hardware/software co-design. Most of this post is about making that measurement honest, portable, energy-aware, and open. With those nouns in place, the thesis is simple: the evidence loop matters more than the single artifact it happens to test first.

The short form is shorter still. The user request creates a workload. The prompt and context budget shape it. The product architecture routes it. The model proposes. The profiler, harness, replay, observability layer, and evidence schema decide what actually happened. A token dashboard can tell you what you spent. It cannot tell you whether the waste came from prompt layout, context bloat, duplicate retrieval, tool-call loops, prefill, decode, KV-cache misses, CPU stalls, quantization fallback, PyTorch fallback, a weak PTX/SASS path, bad routing, wrong hardware placement, or avoidable power and cooling demand. For a CEO or CFO, the point is turning AI spend from a bill into an auditable task-path diagnosis. For an engineering team, the point is just as direct: stop guessing where the system leaks and prove which layer is responsible.

To be clear, the problem is not that nobody has tools. NVIDIA NIM benchmarking names TTFT, end-to-end latency, ITL, TPS, and RPS; GenAI-Perf measures LLM serving throughput and latency and can pull GPU telemetry through DCGM exporter; Nsight gives kernel and system timelines; TensorRT-LLM and Triton provide a mature deployment path. SemiAnalysis InferenceX sets a serious public bar here too: open-source, vendor-neutral benchmarks, public GitHub Actions runs, benchmark recipes, logs, artifacts, dashboard rows, raw data export, and metrics like p99 TTFT, tok/s/GPU, tok/s/MW, cost per million tokens, and joules per token. AMD's current docs show real movement: vLLM on ROCm now has AITER attention backends, MI300X/MI355X tuning, FP4/FP8 notes, and benchmarking guidance; SGLang on MI355X with MoRI includes firmware, NIC, RDMA, prefill/decode, and benchmark setup; ROCm Compute Profiler exposes counters, roofline, memory analysis, and baseline comparison. That is a serious catch-up path. The remaining gap is operational maturity and adoption: fewer battle-tested public recipes outside the major benchmark efforts, more sensitivity to ROCm/container/driver/firmware versions, less model-training data for agents writing HIP/ROCm kernels, and fewer teams with replayable evidence tied to their own customer workload. The best measurement exists, but it is still not evenly distributed across teams, workloads, and hardware paths.

The same applies outside the NVIDIA/AMD axis. Google TPU has XProf and managed ML Diagnostics for XLA workloads. Apple has Metal, MPS, MLX, and growing local-agent relevance. Modular is pushing a cross-vendor MAX/Mojo layer across NVIDIA, AMD, and Apple. Groq, Cerebras, AWS Neuron, edge NPUs, robotics accelerators, and future ASICs all have different profiler surfaces and different ways to be fast. The point is not to crown one vendor. The point is that every hardware path needs the same evidence contract: workload in, configuration named, correctness checked, latency and throughput measured, energy/cost accounted for, failure modes recorded, replay command preserved. Without that, every new processor becomes another closed island and every team relearns the same lesson from scratch.

For engineers, the same idea is a shared contract. A prompt stack, a coding-agent skill, a RAG policy, a diffusion or video workload, a long-context coding task, a kernel, a serving config, a KV-cache policy, a quantization path, a routing rule, and a hardware placement decision all need the same basic discipline: run the workload, collect the trace, profile the hot path, name the baseline, record the failure modes, and make the result replayable. That is the open-source direction we are pushing through InferGuard CLI, workload replay / diagnostics, and kernel-evidence. Not because a schema by itself solves inference, and not because observability is the whole company. Because full-system AI optimization is impossible if nobody can agree on what actually happened. This is also why "not my layer" is not a serious answer anymore. The layers are already coupled. The evidence has to be coupled too.

That is also why this has to be open and vendor-neutral. If the evidence layer only works for one accelerator, one serving engine, one cloud, or one model provider, it will become another closed dashboard. SemiAnalysis is already showing how valuable open benchmark artifacts can be when recipes, runs, logs, metrics, and dashboards are inspectable. Our lane is complementary: help connect that kind of evidence discipline to the full task path, from prompt and code design to compiler IR, kernels, hardware, energy, and the final useful task. That is how teams protect margin, preserve infrastructure headroom, and make better buildout decisions before the backlash shows up at the permitting meeting.

One thing threads through the whole post: evidence is most useful captured in an open, vendor-neutral format. In a world of many processors, engines, caches, kernels, and serving paths, the format you record the truth in quietly becomes the layer everyone has to live with. The loop applies first to CUDA kernels, then to inference engines, KV caches, routing, workload replay, and eventually hardware/software co-design. It is also why the education piece matters. The goal is not to make inference feel mysterious forever. The goal is to upskill the next generation of infra engineers, AI-native engineers, operators, and founders so more people can read the stack, question the benchmark, and make better decisions. Enough context to be responsible has to become teachable. More on both later, because the convergence is the interesting part.

The hackathon didn't convince us that models can write CUDA. Plenty of work already shows that, and shows it well. What it clarified is that the system around the model is what matters: the task contract, compiler boundary, correctness harness, timing discipline, profiler trace, baseline, replay artifact, reward, and memory of attempts that didn't pan out. That's the same system shape production inference needs. Whether the candidate is a CUDA kernel, a KV-cache policy, a vLLM config, a routing rule, a quantization format, or a hardware-specific serving path, the loop only works if the evidence is honest. So that's the layer we've been building from.

That lesson is bigger than CUDA, and bigger than kernels in the narrow sense.

A kernel is the smallest useful place where software hits hardware. That is why it is such a good teacher. The source code, compiler path, target architecture, profiler output, correctness check, timing method, replay command, and baseline all have to line up. If one of them is wrong, the speedup is fake.

The same rule applies one layer up.

A serving config either improves p95 latency or it does not.
A KV-cache policy either reduces memory movement without hurting quality or it does not.
A CPU tool-loop change either reduces wall-clock stalls and retries or it does not.
A routing policy either puts the work on the right hardware or it does not.
A quantization path either uses the intended kernel and preserves task quality or it does not.
An edge deployment either proves the local model was enough or it does not.
A future ASIC either fits the workload or it does not.
An energy claim either lowers energy per successful task or it is just a slide.
A cooling or water-efficiency claim either reduces avoidable heat or cooling demand at that site or it is too vague to trust.
An education claim either makes people better at using the stack or it is content.

That is the lesson hiding inside the kernel environment: portable software is not enough. You need portable observability, profiling discipline, evidence, confidence, and people who can read it. Everything else should stay downstream of what the workload actually shows. The next move is to make that less abstract by walking the layers where an inference workload can leak.

Dell/CoreWeave breakthrough:
  rack delivery reduces physical integration risk
  rack diagnostics reduce system bring-up risk
  rack-scale deployment creates higher capacity ceiling

remaining proof:
  real prompts
  real prefix-cache behavior
  real KV residency and movement
  real prefill/decode split
  real CPU tool-loop time
  real p95/p99
  real quality gate
  real power/cooling proxy
  real cost per successful task

PROMPT DESIGN
  -> cache behavior
  -> prefill / decode balance
  -> kernel execution
  -> HBM pressure
  -> NVLink / NVSwitch / NIC movement
  -> rack power
  -> cooling load
  -> p95 / p99 latency
  -> cost per useful task

visible cost  -> rack power + GPUs + tokens + cooling + human review
hidden output -> accepted patch + resolved ticket + verified report + usable workflow

roadmap of the system boundary:
  chip as primitive
  -> GPU as primitive
  -> GPU server as primitive
  -> GPU cluster as primitive
  -> rack-scale inference system as primitive
  -> optically connected compute fabric as primitive

state in a coding-agent workload:
  stable:
    system prompt
    tool schema
    coding rules
    design-system rules
    security constraints

  semi-stable:
    repo structure
    package structure
    route structure
    component hierarchy
    API types
    common files
    prior validated context

  volatile:
    latest user request
    current file diff
    latest error
    test result
    screenshot feedback
    next action

illustrative rack economics:
  rack capex proxy       -> multi-million-dollar NVL72-class system
  depreciation proxy     -> rack cost / 3 years / 8,760 hours
  IT power proxy         -> 100-140 kW per rack
  facility power proxy   -> IT power * PUE
  simple rack-hour proxy -> hundreds of dollars per active hour
  useful metric          -> successful tasks per rack-hour,
                            per watt, per dollar,
                            at target p95 latency

prefix-cache miss cascade:
  more prefill
  -> more GPU work
  -> more HBM allocation
  -> more cache blocks
  -> more memory pressure
  -> more movement
  -> more interconnect traffic
  -> more power
  -> more heat
  -> more cooling
  -> more cost

bad coding-agent path:
  messy prompt assembly
  -> unstable prefix boundary
  -> prefix-cache miss
  -> repeated prefill
  -> repeated KV creation
  -> HBM pressure
  -> eviction / offload
  -> cache-locality loss
  -> kernel fallback risk
  -> interconnect traffic
  -> higher p95 / p99
  -> higher cost per successful task

optimized coding-agent path:
  stable prefix        -> system prompt, tools, rules, design constraints
  semi-stable prefix   -> repo map, package files, component index
  volatile suffix      -> latest request, diff, error, test output
  cache-aware router   -> place work near useful KV state
  phase-aware schedule -> prefill and decode treated differently
  kernel receipt       -> intended dtype / attention / MoE / KV path verified
  business receipt     -> lower p95 and lower cost per accepted task

kernel and compiler evidence path:
  model operation
  -> framework graph
  -> serving-engine execution path
  -> CUDA / CUTLASS / CuTe / TileLang / Triton
  -> PTX
  -> SASS
  -> GPU execution
  -> Nsight / replay / trace evidence

questions a rack-scale inference operator has to answer:
  are prefixes being recomputed that should be reused?
  are similar requests landing on cache-local workers?
  are prefill and decode separated correctly?
  is KV moving unnecessarily?
  are spills to CPU or NVMe happening at the wrong time?
  are attention kernels hitting the intended path?
  are NVFP4 kernels actually active?
  are MoE experts balanced?
  are collectives dominating useful work?
  is the bottleneck power, memory, network, scheduler, or kernel path?
  are we optimizing raw tokens or completed tasks?

one coding-agent task
  -> CoreWeave cloud endpoint
  -> Dell-integrated GB300 NVL72 rack
  -> Grace CPU / tokenizer / host preprocessing
  -> scheduler / router
  -> prefix-cache lookup
  -> prefill workers
  -> KV-cache allocation in HBM
  -> NVLink / NVSwitch movement
  -> decode workers
  -> CUDA / CUTLASS / CuTe / TileLang / Triton kernels
  -> PTX / SASS evidence
  -> BlueField / NIC / scale-out network
  -> rack power
  -> liquid cooling
  -> p95 latency
  -> cost per successful coding task

The rack is delivered.
Now the workload has to earn the rack.

Rack type:
  GB200 / GB300 NVL72-class rack

Rack composition:
  ~72 Blackwell / Blackwell Ultra GPUs
  ~36 Grace CPUs
  rack-scale NVLink / NVSwitch fabric
  liquid cooling
  scale-out networking

Rack capex:
  ~$3.1M rack-scale server
  ~$3.9M all-in with networking, storage, and related infrastructure

Hourly depreciation proxy:
  $3.9M / 3 years / 8,760 hours
  ~= $148/hour before power, facilities, operations, financing, staff, maintenance, and margin

Rack IT power proxy:
  100kW to 140kW per rack

Middle proxy:
  120kW IT power per rack

Facility power with PUE:
  120kW * 1.2 PUE = 144kW
  120kW * 1.3 PUE = 156kW

At 144kW facility power and $0.08/kWh:
  144 * 0.08 = $11.52/hour

At 144kW facility power and $0.12/kWh:
  144 * 0.12 = $17.28/hour

depreciation + power only:
  ~$148/hour + ~$12–17/hour
  ~= $160–$170/hour

facility overhead
cooling plant
networking
staff
maintenance
financing
utilization risk
software
downtime
margin

A rack-scale inference system can easily represent hundreds of dollars per active hour
once capital, power, facility, operations, and margin are included.

20% waste = ~$32/hour per rack

100 racks:
  ~$3,200/hour
  ~$76,800/day
  ~$2.3M/month

1,000 racks:
  ~$32,000/hour
  ~$768,000/day
  ~$23M/month

cost per successful task at target p95 latency

cost per accepted code change
cost per passing PR
cost per working app screen
cost per fixed bug

GPU server
  + GPU server
  + GPU server
  + network fabric
  + storage
  + CPU memory
  + serving engine
  + scheduler
  + cache layer
  + monitoring
  + power planning
  + cooling planning
  + manual tuning

Which servers handle prefill?
Which servers handle decode?
Where does KV cache live?
Does KV cache stay in HBM?
Does it spill to CPU DRAM?
Does it spill to NVMe?
How much latency does offload add?
Are requests routed back to workers with useful prefix state?
Are tensor-parallel collectives saturating the fabric?
Are expert-parallel all-to-all operations dominating runtime?
Are kernels using the intended dtype path?
Are we actually hitting FP8 / FP4 / NVFP4 kernels?
Are we falling back to slower kernels?
Are GPUs busy on useful work or repeated prefill?
Are we limited by compute, memory bandwidth, interconnect, CPU preprocessing, storage, power, or cooling?

GB300 / NVL72-class rack
  -> Blackwell Ultra GPUs
  -> Grace CPUs
  -> BlueField DPUs
  -> HBM across GPUs
  -> NVLink / NVSwitch inside the rack
  -> scale-out networking outside the rack
  -> liquid cooling
  -> high-density power delivery
  -> rack-level telemetry
  -> serving engine + scheduler + cache layer

advanced packaging
2.5D / 3D integration
chiplets
HBM
larger interposers
co-packaged optics
silicon photonics
optical engines near switch ASICs
external laser sources
EAM / MRM modulation paths
glass substrates
thermal-aware packaging
liquid cooling
power delivery co-design

separate components
  -> integrated servers
  -> integrated racks
  -> optically connected racks
  -> chiplet + HBM + optical I/O packages
  -> compute fabrics where memory, networking, power, and cooling are co-designed

pluggable optics
  -> on-board optics
  -> co-packaged optics near the switch ASIC
  -> optical scale-out fabrics across rows and halls
  -> package-level optical I/O for future accelerator fabrics

Does this make useful state cheaper to keep, move, reuse, or recompute?

The cost of data movement is becoming one of the main limits of AI infrastructure.

3D stacking
chiplet integration
advanced packaging
shorter signal paths
memory closer to logic
higher bandwidth per watt
better thermal handling
custom interconnect
domestic AI accelerator ecosystems

transistor scaling alone is not enough

package-level integration
memory hierarchy
interconnect
thermal design
power delivery
compiler support
EDA tooling
system-level scheduling

Dell-integrated GB300 NVL72-style rack
  -> 72 Blackwell Ultra GPUs
  -> 36 Grace CPUs
  -> BlueField DPUs
  -> rack-scale NVLink / NVSwitch
  -> scale-out networking
  -> liquid cooling
  -> rack diagnostics
  -> power delivery
  -> serviceability and deployment packaging

user request
  -> coding-agent frontend
  -> system prompt
  -> tool schema
  -> repo context
  -> file search
  -> file read
  -> code edit
  -> TypeScript check
  -> error output
  -> model reasoning
  -> another file read
  -> another edit
  -> app launch
  -> screenshot or runtime error
  -> more reasoning
  -> final working screen

system prompt
coding-agent rules
tool schemas
repo editing rules
style guide
design-system instructions
security constraints
output format

repo tree
package.json
route structure
component hierarchy
design tokens
API types
state-management files
reusable UI components
project conventions

latest user instruction
latest file diff
latest TypeScript error
latest runtime error
latest test result
latest screenshot feedback
latest tool output

Build the AI skincare progress screen and make it match the existing app.

placement questions:
  which rack should receive this task?
  which GPU group should handle prefill?
  which GPU group should handle decode?
  does this user / session / repo already have reusable prefix state?
  is there existing KV cache for this task?
  is the rack power / cooling envelope healthy?
  will this placement hurt p95 latency?

system prompt
+ full tool schema
+ entire repo summary
+ file tree
+ design rules
+ several full files
+ previous chat history
+ user request

request routed like generic load
no awareness of repo/session state
no awareness of prefix locality
no awareness of prior KV
no awareness of prefill/decode phase pressure

request routed by workload state
repo/session prefix locality considered
KV residency considered
prefill/decode capacity considered
rack health considered

Extra repeated input:
  +20k to +80k prompt tokens per turn

Business impact:
  larger prompt increases TTFT
  larger prompt creates more prefill work
  larger prompt creates more KV cache
  larger prompt raises cost before the agent has done useful work

The agent looks like it is "thinking,"
but the rack is mostly rereading context it should already know.

The product may look magical,
but the gross margin is leaking through repeated context assembly.

system instructions
tool descriptions
repo paths
source code
TypeScript errors
test logs
user text

CPU-side overhead:
  +10ms to +300ms per turn for messy prompt assembly/tokenization
  higher under high concurrency or heavy RAG/tool fanout

Business impact:
  GPU may wait on CPU preprocessing
  p95 latency gets worse
  expensive rack capacity is blocked by cheap-but-bad host-side work

A multimillion-dollar rack can be slowed down by boring CPU preprocessing.

The bottleneck may be outside the GPU.
That is why surface-level GPU utilization can mislead diligence.

turn 1 prefix: system + tools + repo summary + files A/B/C
turn 2 prefix: system + modified tools + repo summary + files A/B/D + error
turn 3 prefix: system + tools + extra history + repo summary + files A/C/D

I have already computed this prefix.

Prefix-cache hit rate:
  bad path: 10–30%
  good path: 60–90%+ for stable agent context

Repeated prefill:
  +10k to +100k tokens recomputed per multi-turn task

Business impact:
  lost cache reuse becomes duplicated GPU work
  duplicated GPU work becomes duplicated rack power

A cache miss is not just a software metric.
It is wasted capital, wasted electricity, and slower user experience.

A company with better cache reuse can serve more users with the same hardware.
That is a margin advantage.

8k tokens
32k tokens
64k tokens
100k+ tokens

same repo context
same tool schema
same design rules
same package files
same component files

Extra prefill per agent turn:
  +20k to +80k tokens

For a 5-turn coding task:
  +100k to +400k repeated prefill tokens

Rack-level effect:
  more GPU-seconds consumed
  more HBM allocation
  more power drawn for non-useful repeated work

This is where "AI feels expensive."
The user only asked for one screen, but the system keeps repaying the setup cost.

The business question is not how many tokens were generated.
It is how many repeated setup tokens were avoided.

repeated prefill
  -> repeated KV creation
  -> more HBM allocation
  -> more cache blocks
  -> more pressure on the memory manager

KV created per task:
  bad path: large stable context recreated every turn
  good path: stable context reused, only volatile suffix added

Practical metrics:
  KV bytes created per successful task
  KV bytes reused per successful task
  KV bytes evicted per successful task

More KV means fewer concurrent sessions fit in HBM.
Fewer concurrent sessions means less throughput per rack.
Less throughput per rack means higher cost per task.

KV cache is working capital for inference.
If it churns, the business needs more racks for the same user workload.

KV efficiency is capacity efficiency.
Capacity efficiency becomes gross margin.

hot KV blocks in HBM
cold KV blocks pushed out
useful blocks evicted too early
blocks moved to CPU DRAM
possibly blocks moved to NVMe
later blocks pulled back

HBM pressure:
  higher active KV footprint
  lower effective concurrency
  more offload traffic

Latency impact:
  HBM access is fast
  CPU/NVMe offload is much slower
  bad offload timing can hurt p95/p99 badly

The rack may have enough raw FLOPS,
but not enough useful memory residency.

The bottleneck is not always compute.
Sometimes the expensive rack is starved by memory placement.

The best infrastructure companies will understand memory residency,
not just accelerator count.

turn 1 lands on GPU group A
turn 2 lands on GPU group B
turn 3 lands on GPU group C

Cache locality:
  bad path: similar requests scatter across workers
  good path: similar requests route to cache-local workers

Potential waste:
  +10–40% more prefill work on stateful agent loops
  lower effective throughput at the same rack power

Same 120kW-class rack.
Less useful work.
Higher cost per completed task.

Load balancing is not enough anymore.
The scheduler has to place state, not just requests.

State-aware routing is a defensible systems advantage
because it compounds across every request.

prefill is heavy
decode is long
batch shapes are uneven
some GPUs wait
some GPUs hold KV
some GPUs run small decode steps
some GPUs wait on communication

Utilization:
  GPU utilization may look high during prefill
  useful decode throughput may still be poor
  batch fragmentation can lower tokens/sec per watt

Business impact:
  rack looks busy
  user still sees slow p95
  cost per useful output rises

Busy GPUs are not the goal.
Completed tasks at target latency are the goal.

Utilization without task success is vanity infrastructure.

GEMM
attention
FlashAttention-style prefill attention
decode attention
MoE expert GEMM
router operations
RMSNorm / LayerNorm
fused dequant
sampling
KV load/store
KV gather/scatter
communication kernels

NVFP4 / FP4 math path
FP8 path
BF16 fallback path
fused dequant path
optimized attention path
optimized MoE expert path

wrong shape
bad batch
bad dtype layout
unsupported fusion
kernel fallback
extra memory reads/writes
more launch overhead
less tensor-core utilization

Kernel path:
  intended: optimized FP4 / FP8 / BF16 / attention / MoE path
  bad: fallback or unfused path

Possible waste:
  10–50%+ loss on affected operations depending on fallback severity

Business impact:
  same chip
  same rack
  lower useful throughput

You can buy the newest GPUs and still not get the economics
if the software misses the optimized kernel path.

Hardware access alone is not the moat.
Making the hardware run the intended path is the moat.

PyTorch / serving engine says the operation ran.

Did this lower to the intended CUDA / Triton / CUTLASS / TileLang path?
Did the compiler emit the expected PTX?
Did the final SASS use the expected instructions?
Did tensor cores do the work?
Did memory movement dominate?
Did occupancy collapse?
Did register pressure cause spills?
Did the kernel wait on memory?

Evidence needed:
  kernel time
  occupancy
  memory bandwidth
  tensor-core usage
  register spills
  instruction mix
  actual SASS path

Business impact:
  without kernel evidence, performance claims are guesses

This is the engineering receipt.
If the kernel did not actually run the fast path, the business case is fiction.

The diligence question is:
can the team prove where performance comes from?

GPU <-> GPU over NVLink / NVSwitch

NIC / DPU / InfiniBand / Ethernet fabric

cache moved too often
expert routing imbalanced
collectives too frequent
prefill/decode disaggregation poorly placed
KV transferred across the wrong boundary

Movement metrics:
  GB moved per task
  NVLink/NVSwitch traffic per task
  CPU-GPU transfer per task
  network wait time per task

Business impact:
  movement consumes power
  movement adds latency
  movement can make GPUs wait

Networking is no longer plumbing.
It is part of the inference cost structure.

The next generation of AI infra winners will optimize movement,
not just math.

more duplicated compute
more memory movement
more interconnect activity
more CPU involvement
more DPU/NIC activity
more power draw

Rack IT power proxy:
  ~120kW per rack

Facility power at PUE 1.2:
  ~144kW

Energy per minute:
  144kW / 60 ~= 2.4 kWh per rack-minute

Energy per second:
  144kW / 3600 ~= 0.04 kWh per rack-second

At $0.10/kWh:
  ~24 cents per rack-minute for electricity only
  before capex, facility, staff, maintenance, and margin

Wasted seconds are not free.
Every repeated prefill loop consumes power and rack time.

The AI bill is not only tokens.
It is rack-seconds, watts, cooling, and capex amortization.

Energy efficiency becomes product margin.

same completed task
higher power
more heat
more cooling
worse energy efficiency
worse margin

120kW IT power becomes heat.

At PUE 1.2:
  extra facility overhead ~= 24kW per rack

If software waste increases useful runtime by 20%:
  the cooling system also supports that waste

Bad inference software can turn into cooling pressure.
Cooling pressure can limit sustained throughput.

Cooling is not just a facilities issue.
It is tied to software efficiency.

Companies that ignore cooling and power will underestimate true cost of goods sold.

higher TTFT
higher p95 / p99 latency
lower throughput per rack
more GPU-equivalent demand
higher power per useful token
higher cooling load
higher cost per successful coding task

If bad cache/routing wastes 20% of a rack:
  ~$32/hour wasted at simple $160/hour rack proxy

If this happens across 100 racks:
  ~$3,200/hour
  ~$76,800/day
  ~$2.3M/month

If this happens across 1,000 racks:
  ~$32,000/hour
  ~$768,000/day
  ~$23M/month

This is why small inference-efficiency gains become huge at scale.
A 10–20% improvement is not a benchmark vanity metric.
It can be millions of dollars per month.

The economic wedge is not just lower cost.
It is higher throughput, better latency, lower COGS, and better gross margin.

stable prefix:
  system prompt
  tool schema
  coding rules
  design-system rules

semi-stable prefix:
  repo map
  package manifest
  route structure
  component index
  API type summary

volatile suffix:
  latest user instruction
  latest file diff
  latest TypeScript error
  latest test output

Repeated context avoided:
  20k–80k tokens per turn

For a 5-turn task:
  100k–400k tokens of prefill avoided

Business value:
  lower TTFT
  less KV churn
  more sessions per rack
  lower cost per completed task

Prompt structure becomes a cost-control mechanism.

The product layer can create infrastructure leverage
before the model even runs.

stable prompt chunks already tokenized
tool schema token IDs reused
repo map token IDs reused
scratch buffers reused
normalization consistent
token parity verified

CPU latency saved:
  tens to hundreds of milliseconds per turn in messy systems

Business value:
  lower p95
  less GPU waiting
  better user experience

Cheap host-side improvements can unlock expensive rack utilization.

Optimization is not only low-level CUDA.
Sometimes the highest ROI is removing simple upstream waste.

stable system/tool/design prefix: hit
repo map prefix: hit
common file context: hit or partial hit
latest error suffix: miss

Prefix-cache hit rate:
  from 10–30% bad path
  to 60–90%+ for stable agent context

Business value:
  less repeated prefill
  fewer GPU-seconds per task
  less power per successful task

Cache hit rate is a business metric now.

Better cache stability turns the same rack into more sellable capacity.

Which worker already has useful KV state for this repo/session/task?
Which GPU group holds the prefix?
Which prefill worker should handle the new suffix?
Which decode worker should continue generation?

turn 1 -> GPU group A
turn 2 -> same or compatible KV locality
turn 3 -> same or compatible KV locality

Avoided recompute:
  10–40% on stateful multi-turn loops, depending on workload

Business value:
  same rack
  more completed tasks
  lower cost per task

State-aware routing turns the rack from generic capacity into reusable capacity.

Routing quality becomes infrastructure alpha.

do not recompute:
  system prompt
  tool schema
  repo map
  design rules
  common files

only prefill:
  latest error
  latest diff
  latest user instruction
  newly read file

Prefill avoided:
  100k–400k repeated tokens over a 5-turn coding task

Business value:
  more users served per rack
  lower p95
  less power spent on repeated context

Avoided prefill is avoided spend.

The best inference systems do not just run faster.
They avoid work entirely.

hot KV:
  stays in HBM

warm KV:
  stays near likely reuse worker

cold KV:
  moves to CPU DRAM or lower tier

dead KV:
  evicted

Metrics:
  higher KV reuse ratio
  lower KV eviction rate
  fewer GB moved per task
  higher effective concurrency

Business value:
  fewer racks needed for the same workload
  lower capex per unit of demand

KV placement is capacity planning.

Memory placement strategy can delay capex.
That is real cash-flow value.

large prompt
parallel work
heavy attention/GEMM
creates KV
throughput-oriented

one token step at a time
reads KV repeatedly
latency-sensitive
memory/state dependent
batching-sensitive

prefill workers handle large prompt chunks
decode workers handle token generation
KV transfer path is measured
scheduler balances TTFT and throughput

Metrics:
  lower TTFT
  better decode occupancy
  fewer batch stalls
  better p95 / p99

Business value:
  better SLA
  less overprovisioning

Better scheduling reduces the need to buy buffer capacity for latency spikes.

Lower tail latency means fewer wasted reserve racks.

attention prefill:
  FlashAttention-style kernels

decode:
  decode attention kernels optimized for KV reads

GEMM:
  CUTLASS / CuTe / vendor kernels

FP4 / NVFP4:
  low-precision tensor-core path where appropriate

MoE:
  router + expert GEMM + combine kernels

KV movement:
  load/store/gather/scatter/compact kernels

sampling:
  logits processing + sampling kernels

high-level model op
  -> serving engine execution plan
  -> CUDA / Triton / TileLang / CUTLASS / CuTe
  -> PTX
  -> SASS
  -> GPU execution
  -> Nsight / trace / replay proof

Operation-level value:
  10–50%+ savings on affected hot paths when avoiding fallback or unfused execution

Business value:
  higher tokens/sec per watt
  more completed tasks per rack

Kernel correctness is not academic.
It determines whether expensive silicon is actually monetized.

The team that can prove kernel-path correctness can prove hardware leverage.

KV movement only when saved prefill > transfer cost
collectives shaped around model parallelism
expert routing balanced
decode workers placed near needed state
prefill/decode transfer measured
NVLink/NVSwitch traffic attributed to workload phase

Metrics:
  fewer wasted GB moved per task
  less network wait time
  less interconnect power
  better GPU occupancy

Business value:
  better rack utilization
  lower power per useful task

The network should move profit-producing state, not noise.

Movement efficiency becomes a durable advantage as clusters scale.

less repeated prefill
less HBM churn
less unnecessary movement
better kernel efficiency
better batching
better cache locality
more completed tasks per unit energy

If throughput improves 20% at same rack power:
  energy per successful task drops ~16.7%

If throughput improves 50% at same rack power:
  energy per successful task drops ~33.3%

If throughput doubles:
  energy per successful task drops ~50%

Power bill may stay similar,
but unit economics improve because more tasks complete per rack-hour.

The win is not always lower absolute power.
The win is lower power per useful task.

Energy per task is the new gross-margin primitive.

same rack
more useful throughput
better sustained performance
less wasted energy
better cooling headroom
lower cost per task

If software reduces wasted rack work by 20%:
  thermal pressure from useless work falls proportionally
  sustained throughput improves
  cooling headroom improves

fewer thermal throttling events
more predictable SLA
less overbuild

Software efficiency buys cooling headroom.

Better software can make the same physical data center produce more revenue.

cost per successful coding task at target p95 latency

screen implemented
code compiles
tests pass
app launches
user accepts result

lower TTFT
lower p95 / p99
higher completed tasks per rack
lower GPU-equivalent demand
lower energy per useful task
lower cooling pressure per useful task
better gross margin

Example:
  bad path: 1,000 successful coding tasks/rack-hour
  good path: 1,300 successful coding tasks/rack-hour

At $160/hour simple rack proxy:
  bad path: $0.160/task
  good path: $0.123/task

Unit cost improvement:
  ~23% lower cost per successful task

At 10M successful tasks/month:
  bad path: ~$1.6M rack proxy cost
  good path: ~$1.23M rack proxy cost
  savings: ~$370k/month before broader cloud/operator markups

This is the board-level story:
same rack class, same model family, better software path,
lower unit cost, better margin, better latency, more capacity headroom.

This is why inference optimization compounds.
Every percent saved at the task level scales across the entire customer base.

one coding-agent task
  -> N prompt tokens
  -> M output tokens
  -> X prefill tokens recomputed
  -> Y prefix tokens reused
  -> A GB of KV created
  -> B GB of KV reused
  -> C GB moved across GPU / CPU / rack fabric
  -> D GPU-seconds consumed
  -> E rack-seconds consumed
  -> F watt-hours consumed
  -> G cooling overhead
  -> final dollar cost

did the code compile?
did the tests pass?
did the app launch?
did the user accept it?

cost per successful task at target p95 latency

cost per accepted code change
cost per passing PR
cost per working app screen
cost per fixed bug

cost per resolved ticket
cost per escalated ticket avoided
cost per customer issue solved

cost per verified answer
cost per useful report
cost per experiment completed

cost per grounded answer
cost per citation-correct response
cost per successful retrieval + generation

prompt structure
  -> prefix-cache hit rate
  -> prefill time
  -> decode time
  -> KV-cache bytes
  -> HBM pressure
  -> kernel path
  -> interconnect movement
  -> power draw
  -> cooling load
  -> p95 latency
  -> cost per successful task

cost per successful task at target latency

where is the state?
why is it moving?
is movement cheaper than recompute?
is the kernel path correct?
is this watt producing useful work?
did the task succeed?

state-aware routers
cache-aware schedulers
rack-level workload replay
energy-aware inference planning
KV-cache placement policies
prefix-stability tooling
power-aware batching
thermal-aware routing
optical-aware networking
compiler/kernel evidence layers
task-level cost accounting
cross-hardware portability layers
packaging-aware architecture search

which model should run?

where should this state live?
which rack should run this task?
which GPU group already has the prefix?
is recompute cheaper than transfer?
is CPU offload worth it?
is NVMe offload too slow?
should this state move electrically or optically?
which kernel path will this request hit?
will this batch violate latency?
will this workload push the rack into a worse power/cooling regime?
what is the expected cost per successful task?

accepted code changes
resolved support tickets
completed research reports
grounded RAG answers
finished agent workflows
generated videos

gross margin per successful task

1. Reduce wasted prefill
2. Increase prefix/KV reuse
3. Improve kernel and dtype paths
4. Increase completed tasks per rack-hour

rack capex is millions of dollars
power is continuous
cooling is continuous
underutilization is expensive
latency misses require overprovisioning
failed tasks destroy unit economics

better inference economics means better product margins
better latency means better user experience
better rack utilization means faster scaling
better task success means more revenue per compute dollar

AI-native companies will increasingly be judged on compute efficiency,
not just model quality or revenue growth.

The best companies will convert compute into useful work efficiently.
The weak companies will buy more GPUs and still have bad margins.

How many successful tasks can this company produce
per dollar of compute,
per watt,
per rack,
per hour,
at target latency?

what state is stable
what state is volatile
what should be cached
what should be recomputed
what should stay in HBM
what can move to CPU
what can move across optics
what should never move
which worker should get the next turn
which kernel path actually ran
which movement was worth it
which watt produced useful work

Before any of the technical sections, here's the picture worth carrying with you, and the answer to "why does this work matter to anyone running AI workloads at scale." Inference optimization is not one lever. It is seven connected layers, and every layer changes something different. Most teams pull one layer (usually "compress the KV cache" or "switch engines") and assume they've optimized inference. They haven't. They've moved one layer. The point of this post is to walk the whole path, show that the same pattern shows up at every layer, and make the case that the only useful answer is full-system thinking, supported by an evidence layer that can see all seven honestly.

The distributed version is the one that matters most now. Gen 3 agentic workloads do not stay inside one GPU: CPU tool loops, GPU prefill/decode, KV movement, prefix reuse, tensor-parallel collectives, expert-parallel dispatch, RDMA, network topology, and serving decisions all shape the task path. The kernel is still the strictest layer, but the kernel is downstream of where the serving system put the state. That is why a single-GPU speedup is a useful receipt, not the final claim. The claim has to survive the full path.

The public Kimi K2.5 GB200 benchmark later in the post is the concrete version of this. Same broad model family, same 8k/1k workload shape, same rack-class hardware path, but different expert-parallel and serving layouts produce very different useful capacity. That is why the layer map matters: the money is not hiding in one layer.

One update since the first draft: "serving engine" is no longer one category. It now splits into several workload paths. High-throughput open serving points toward vLLM. Structured agent and rollout serving points toward SGLang. NVIDIA-optimized production runtime points toward TensorRT-LLM/NIM/Triton. Distributed orchestration points toward Dynamo. Batch-1 low-latency decode is where Kog/KIE is making its public claim. Stateful live sessions are where LayerScale is making its public claim. Reusable KV/state movement points toward LMCache, Tensormesh, and Mooncake. The newer vLLM read is broader than "GPU server": vLLM is becoming a control surface for GPU serving, CPU serving, KV-cache offload, distributed KV pools, prefill/decode disaggregation, speculative decoding, RL rollout serving, semantic routing, elastic expert parallelism, and Kubernetes deployment. For a CEO or CFO, the wrong engine choice can make a good model expensive, slow, or operationally brittle. For an engineer, the receipt has to name the workload path before it names the winner.

request shape  ->  8192-token prefill + 1024-token decode
model path     ->  nvidia/Kimi-K2.5-NVFP4
benchmark key  ->  kimik2.5-fp4-gb200-dynamo-vllm
serving path   ->  Dynamo + vLLM + disaggregated prefill/decode
state path     ->  KV transfer through NixlConnector
hardware path  ->  GB200 NVL72-class runner
code/config    ->  recipes/vllm/kimi-k2.5/8k1k/disagg-gb200-6p1d-dep4-dep16.yaml
business read  ->  output capacity, GPU-equivalent demand, rack headroom, energy proxy

user request
  -> Hermes/OpenClaw frontend
  -> Claude Code wrapper
  -> stable prompt / tool schema / repo context
  -> tokenizer CPU workers
  -> prefix-cache lookup
  -> KV-aware router
  -> prefill worker
  -> decode worker
  -> file edits + TypeScript / app / test output
  -> next turn reuses or recomputes state
  -> successful app screen

before optimization:
  repeated repo context reloads
  unstable prefix-cache reuse
  repeated prefill across turns
  KV blocks created, moved, evicted, and rebuilt
  GPUs busy on duplicated work
  HBM fills with state that may not be reused
  NVLink / NVSwitch traffic increases
  CPU-GPU transfers happen unnecessarily
  power draw stays high
  cooling load rises
  p95/p99 and cost per coding task get worse

after optimization:
  similar turns route cache-local
  stable repo/design context reuses prefix/KV state
  only volatile errors, diffs, and screenshots are new
  prefill and decode are scheduled differently
  KV stays close to where it will be reused
  NVFP4 / CUTLASS / CuTe / TileLang / Triton kernels hit the intended path
  HBM pressure drops
  NVLink movement becomes intentional
  power per useful token improves
  cooling headroom improves
  throughput per rack rises
  cost per successful coding task drops

request shape  ->  DeepSeek-R1 interactive serving
model path     ->  large MoE with expert routing
serving path   ->  SGLang + MoRI + ROCm
state path     ->  MoRI-IO KV/state transfer
compute path   ->  AITER + FlyDSL FusedMoE + tuned GEMMs
comm path      ->  FP4 dispatch + FP8 combine + adaptive kernels
overlap path   ->  two-batch overlap + SDMA
decode path    ->  Specv2 MTP on ROCm
business read  ->  lower cost per million tokens at target tok/s/user

END-TO-END PICTURE

user request
  -> SGLang scheduler decides prefill/decode placement
  -> prefill ranks build attention state
  -> MoRI-IO moves KV/state across the disaggregated path
  -> decode ranks run attention + Specv2 MTP draft/verify
  -> routed tokens hit MoRI expert-parallel all-to-all
       dispatch: FP4 quantized tokens move to experts
       compute:  AITER / FlyDSL FusedMoE runs expert GEMMs on MI355X
       combine:  FP8 quantized expert outputs return
  -> two-batch overlap + SDMA hide communication behind compute
  -> CPU streaming path sends tokens back to the user
  -> InferenceX/SemiAnalysis TCO row turns the whole path into cost

The important part is that the layers feed each other. Bad context engineering creates extra prefill. Extra prefill creates more KV. More KV stresses cache placement and compression. Cache movement changes serving-engine scheduling. Serving decisions decide which kernels run, which dtype path is active, and whether the GPU waits on CPU, network, or memory. The kernel is the strictest proof point, but it is not where the workload begins.

Quantization shows up in two different places across these layers. Weight quantization and KV-cache quantization are related, but they are not the same problem. NVFP4 and MXFP4 shrink the math path: weights, activations, expert GEMMs, and tensor-core or matrix-core work. TurboQuant and SpectralQuant shrink the state path: KV cache, long-context memory, and attention behavior. A long-context agent pays for both, so the evidence layer has to see both.

The trap most teams fall into is pulling one layer in isolation. Compressing KV doesn't help if the prefix-cache hit rate is 10%. Switching engines doesn't help if the kernel is leaving 75% of the silicon idle. PEEK doesn't help if the workload's context is task-disjoint. A faster single-GPU kernel does not help enough if the workload loses the win to AllGather, AllReduce, MoE dispatch, CPU preprocessing, CPU orchestration, or bad KV placement. Full-system thinking means knowing which layer is the bottleneck on your workload, and pulling that one, and the only way to know is to measure honestly with an evidence layer that sees every layer. Beneath every layer sits the same pattern: externalize the relevant state, make the contract between layers measurable, let an optimizer or an agent or an engineer read from the evidence. Every section that follows is one layer walked through, with a team or two doing exceptional work on it, and a tie-back to the layer underneath that turns the layer's win into a portable claim. That underneath layer (open, neutral, and replayable) is what kernel-evidence is, and the rest of the post is what it looks like when you take it seriously across all seven layers. The layer map is how you stop guessing; the next section is the historical reason that map changed.

This section is the historical frame. It explains why AI infrastructure is moving from GPU-only serving to CPU+GPU task execution. That shift matters because the cost is no longer only model inference; it is the whole agent loop.

It is also why the gigawatt frame belongs this early in the post. Gen 1 capacity was mostly about dense training throughput. Gen 2 capacity was mostly about serving tokens cheaply. Gen 3 capacity is about how many verified tasks a mixed CPU/GPU/state/tool system can finish inside the same power envelope. That is the shift from tokens per GPU to successful tasks per megawatt.

Before the hackathon recap, the framing that holds the rest of this post together. The history of AI infrastructure is not a straight line from CPU to GPU. It is a loop, with three distinct generations of inference optimization stacked on top of it. First the CPU ran the software. Then the GPU became the dense-math engine. Now the CPU is coming back as the environment processor around GPU inference, while multi-GPU communication decides whether the local wins survive at rack scale. 2026 is the year the third generation becomes the company-building era: and the year automated kernel generation, the work this post recaps from OpenEnv, becomes the foundation move underneath it.

A quick pre-history, because the loop only makes sense with it.

Before any of this, computing was CPU-first. Roughly 1940s through 2006, the CPU was the center of compute because almost all software was branchy, latency-sensitive, irregular, and control-heavy. Operating systems, databases, compilers, web servers, networking, storage orchestration, business logic: all CPU. Then in November 2006, NVIDIA announced CUDA on the GeForce 8800, and in June 2007 shipped CUDA 1.0. Programmable GPUs escaped graphics. The first wave of GPU compute (2006–2012) was HPC: simulation, finance, molecular dynamics, image processing. The center of gravity moved when Krizhevsky, Sutskever, and Hinton trained AlexNet on two NVIDIA GTX 580 3GB GPUs over five-to-six days, hit 37.5% top-1 / 17.0% top-5 on ILSVRC-2010 and 15.3% top-5 on ILSVRC-2012, and the GPU became the default hardware path for deep learning. That's the pre-history. The three generations below are what happens on top of it.

Generation 1: GPU dense-math era (training-shaped optimization). ~2012–2022.

The first generation was about keeping the accelerator fed with dense numerical work. AlexNet on two GTX 580s grew into ResNet, then transformers, then GPT-3, then the full training-cluster era. The optimization target was throughput on dense parallel math (matrix multiply, convolution, attention, normalization, all-reduce) running at the largest possible batch size across the largest possible cluster. The tooling that defined the era: CUDA, cuDNN, NCCL, tensor cores (Volta 2017), mixed precision (Pascal/Volta), distributed training (Megatron, DeepSpeed, FSDP), A100 (Ampere 2020), H100 (Hopper 2022). The question every system answered: how do we keep the GPU fed? The CPU was the host that orchestrated data movement, not the bottleneck.

Generation 2: GPU LLM serving era (inference-shaped optimization). ~2022–2024.

Training was still huge, but the serving problem became economically central as ChatGPT and the GPT-4/Claude/Gemini generation hit consumer scale. The optimization target shifted from "train one giant model fast" to "serve millions of requests cheaply with low latency." That produced a different toolchain: prefill vs decode split, KV cache as a first-class memory tier, paged attention (vLLM, June 2023), continuous batching, speculative decoding (EAGLE / EAGLE-3 / Medusa / MTP / DFlash), quantization, prefix caching, tensor parallelism, expert parallelism, engine tuning, vLLM / SGLang / TensorRT-LLM / TGI / LMDeploy / MAX. The decode branch split again: EAGLE-3-style systems draft future tokens autoregressively, while DFlash-style systems use block diffusion to draft a whole token block in one pass before target verification. Still GPU-first. Still mostly about scheduling, memory management, latency, and economics on the accelerator. This is the era automated kernel generation in §08 is the highest-leverage version of. If you can automate kernel generation, you attack the bottom of this generation's stack across model architectures, dtypes, sequence shapes, and chip generations.

This is the bridge from Gen 2 to Gen 3. Once a product gets better by sampling, verifying, searching, repairing, or rolling out trajectories, decode speed stops being only a serving metric. It becomes the budget for capability. The winning systems over the next one to three years will turn the same hardware into more verified branches, more useful rollouts, more tool-checked attempts, and fewer repeated state loads. The unit shifts from token throughput to useful thinking per second, dollar, watt, and successful task.

Generation 3: CPU+GPU agentic task era (task-path-shaped optimization). ~2024–now.

This is the generation the rest of this post is written into. Agentic AI breaks the assumption that "inference" is a GPU-bound problem. A coding agent does not just call a model once. It loads system prompts, tool schemas, repo context, long input, KV cache, plans, tool calls, file edits, bash, tests, lints, errors, retries, and reloaded context, then maybe succeeds. The GPU is still central. It is no longer the only bottleneck. The CPU is back, not as the center of compute it was in the pre-history, but as the environment processor running the harness around the model. And once the model no longer fits inside one accelerator, the problem expands again: multi-GPU communication, cache locality, network topology, and serving-engine placement become part of the same task trajectory.

The Hermes/OpenClaw -> Claude Code mobile-screen example is the clean Gen 3 shape. The GPU does not just serve one answer. The system repeatedly loads repo state, tool schemas, design-system rules, TypeScript errors, screenshots, test output, and file diffs. The CPU runs the tool loop. The GPU runs prefill and decode. The cache layer decides whether stable context is reused or rebuilt. The rack fabric decides whether the state stays local or moves. The user asked for a screen. The system ran a stateful workload.

NVIDIA's own roadmap validates this from the silicon side. The Vera CPU, announced at GTC 2026 and shipping H2 2026, is described by NVIDIA as "built for reinforcement learning and agentic AI, powering the code, tools, and data workflows that operate beyond the model." Each Vera CPU is 88 Olympus cores with 1.2 TB/s memory bandwidth; a single liquid-cooled Vera CPU Rack integrates 256 Vera CPUs (22,528 cores) and is specified to sustain over 22,500 concurrent CPU environments, with one environment per core running the model's tool-call sandbox, file-edit operations, test/lint loops, and data-processing work outside the model. NVIDIA's own technical blog says it delivers "4× sandbox density and 2× performance per watt over x86-based racks for AI Factories." The CPU has become first-class enough in agentic workloads to justify dedicated silicon and a dedicated rack-scale product.

The May 31 GTC Taipei update makes that less theoretical. NVIDIA now says Vera is in full production, with the CPU aimed at agentic AI, RL, data processing, code execution, tool use, and evaluation. The Vera Rubin ramp update adds the rack and factory surface around it: Vera CPU, Rubin GPU, Groq 3 LPX, Vera BlueField-4 STX storage, Spectrum-6 SPX Ethernet racks, Spectrum-X Ethernet Photonics, MGX partners, DSX factory design, and Taiwan ODM manufacturing. That is the hardware industry saying the Gen 3 task path is now large enough to shape chips, racks, storage processors, networking, and factory reference designs.

RTX Spark is the same pattern from the opposite end of the market. NVIDIA's RTX Spark announcement puts a Blackwell RTX GPU, FP4 Tensor Cores, NVLink-C2C, a 20-core Grace CPU, and up to 128GB unified memory into Windows PCs for personal agents. The data-center version asks how many verified tasks a rack can produce per megawatt. The consumer version asks which private, low-latency, app-control, creative, or personal-memory tasks should run locally instead of being shipped to the cloud. Same principle, smaller power envelope: the workload decides placement.

The clearest hardware signal is Vera: NVIDIA is no longer treating the CPU as just the host feeding the GPU. It is treating the CPU as the processor that runs the agent environment.

SemiAnalysis' coding-assistant data from the operator side says the same thing: roughly 42% of modern agentic-coding wall-clock time is CPU-side tool use (file edits, bash, tests, lints, package installs, and sandbox execution) rather than model decode. The harness determines cost per task because prompt caching, input/output ratio, and tool-use patterns are largely harness-determined. Claude Code, Codex, Cursor, OpenCode, and a homegrown agent are not just different wrappers around a model. They are different workload generators. They produce different token traces, cache behavior, CPU tool loops, retry patterns, and therefore different cost-per-task numbers, against the same underlying model. That is the third generation showing up on the invoice before most teams have the instrumentation to explain it.

Search and RAG systems hit the same Gen 3 shape without looking like coding agents. One user query can retrieve hundreds of candidate documents, normalize them, tokenize every query-document pair, batch a reranker, score relevance, and only then call the final generator. Perplexity's Unigram tokenizer work is the production example: the user sees one answer, but the system may do hundreds of CPU tokenization jobs before the GPU model forward pass that decides what context survives. That is still Gen 3 inference: CPU + GPU + state layout + task outcome.

The research signal points the same way from the systems side. Berkeley Sky keeps showing that the hard part is not only one model or one kernel, but search, scheduling, and portability across clouds and hardware. mKernel pushes on GPU-driven communication because multi-GPU inference can lose time in collectives and state movement even when the local kernel is strong. InferenceX pushes benchmark work toward real traces and workload replay. Different projects, same lesson: Gen 3 optimization has to measure the path, not just the chip.

Zyphra's AWS Inferentia2 work adds the same point from a non-NVIDIA stack. In their Domino-on-Neuron experiment, the useful question was not "is AWS better than NVIDIA or AMD?" The useful question was whether a Llama 3-8B inference path on Inferentia2, NeuronCore-v2, NeuronLink, NKI kernels, tensor-parallel collectives, and compiler-visible tiled scheduling could reduce exposed communication time. That is the decision rule for Gen 3 hardware: which stack gives the best cost per successful task at p95/p99 for this workload? Sometimes the answer is NVIDIA. Sometimes AMD. Sometimes AWS Trainium/Inferentia, TPU, local/edge, or a future ASIC. The workload has to decide.

The same three generations, visualized as one stack.

INFERENCE OPTIMIZATION LADDER - THREE GENERATIONS

PRE-HISTORY (1940s–2006)        CPU ERA
                                 General-purpose software, OS, databases, compilers, web servers.
                                 Question: how do we make software run reliably at scale?

GEN 1 (~2012–2022)               GPU DENSE-MATH ERA (training-shaped)
01. Kernel                        CUDA, cuDNN, tensor cores, fused matmul / conv / attention
02. Distributed training          NCCL, Megatron, DeepSpeed, FSDP, all-reduce
                                 Question: how do we keep the accelerator fed for training?

GEN 2 (~2022–2024)               GPU LLM SERVING ERA (inference-shaped)
03. Kernel + compiler             CUDA, HIP, Triton, TileLang, CuTe, CUTLASS, AITER,
                                  torch.compile, MLIR, MAX, SCALE
04. Serving engine                vLLM, SGLang, TensorRT-LLM, TGI, LMDeploy, MAX
05. KV cache in HBM               paged attention, FP4 / QJL / PolarQuant
                                 Question: how do we serve millions of requests cheaply on the GPU?

GEN 3 (~2024–NOW)                CPU+GPU AGENTIC TASK ERA (task-path-shaped)
06. KV cache & state movement     prefix cache, LMCache, HBM↔CPU↔NVMe offload
07. CPU tool loop                 bash, file edits, tests, lints, sandbox execution, package installs
                                  (~42% of agentic-coding wall-clock per SemiAnalysis)
08. Agent trajectory              context layout, retries, tool calls, task decomposition,
                                  prompt prefix stability, routing, success / failure
09. Data center / HW topology     GPUs (Rubin), CPUs (Vera, 88 Olympus cores / chip, 22.5K
                                  concurrent envs / rack), AWS Trainium/Inferentia,
                                  LPUs, TPUs, storage, networking
                                 Question: how do we make this useful AI task complete successfully,
                                 cheaply, and reliably across CPU, GPU, memory, storage,
                                 network, cache, tool loops, and agent retries?

The one-line summary that holds it together.

Gen 1 asked: how do we keep the accelerator fed for training? Gen 2 asked: how do we serve millions of requests cheaply on the GPU? Gen 3 asks: how do we make one useful AI task complete successfully, cheaply, and reliably across CPU, GPU, memory, storage, network, cache, tool loops, and agent retries?

Automated kernel generation is the foundation move for Gen 3, not the whole game. Kernels are the strictest verifiable layer; the candidate compiles or it does not, passes correctness or it does not, runs faster on silicon or it does not. That makes them the easiest place to develop the evidence-loop discipline. But once that discipline exists, the same loop generalizes one layer up at a time: engine config, KV-cache policy, prefix-cache stability, CPU tool loops, agent trajectory, hardware topology. The candidate type changes; the loop shape stays the same. Same eight verbs: generate, compile or execute, verify, benchmark, reward, reject, remember, improve: running across the whole task path instead of just one GPU program. Once the workload becomes the unit, the accounting has to move from tokens to completed tasks.

The generational shift above shows why the old metric breaks. Cost per token is not the same thing as cost per useful work.

A coding-agent task is not one model call. It is a loop. Load repo context. Read files. Carry tool schemas. Prefill a long prompt. Decode a plan. Edit files. Run bash. Run tests. Read failures. Re-enter context. Retry. Eventually produce a patch. Some of that work is GPU prefill. Some is GPU decode. Some is KV-cache residency. Some is CPU sandbox execution. Some is file I/O. Some is tool orchestration. At larger scale, some is multi-GPU communication, cache transfer, and network waiting. The invoice may show tokens, but the system is paying for the whole task trajectory.

before cache-aware routing:
  turn 1 -> stable repo/app context prefills on worker A
            KV created on worker A
            Claude Code edits files
            TypeScript error returns

  turn 2 -> same stable context + new error lands on worker D
            worker D lacks worker A's KV state
            full or mostly full prefill repeats

  turn 3 -> simulator or screenshot feedback lands on worker B
            partial/no cache hit
            stable context gets paid for again

after cache-aware routing:
  turn 1 -> stable context prefills once, KV indexed
  turn 2 -> stable chunks hit prefix cache, only the TypeScript error is new
  turn 3 -> repo/design context reused, only volatile screenshot feedback changes

receipt:
  cost per successful app-building task
  p95/p99 latency
  prefix-cache hit rate
  prefill vs decode time
  CPU tool-loop time
  retries and failed tool runs

Reranking is the search version of the same problem. The user sees one answer. The system may retrieve, normalize, tokenize, batch, and score hundreds of documents before final generation starts. A tokenizer does not emit output tokens, but it can decide whether the task is cheap. If the tokenizer path allocates, hashes, pointer-chases, and burns CPU cycles across every candidate, the cost shows up as cost per answered query, not cost per output token. Token dashboards miss this because the tokenizer produces no output tokens. Cost per successful task catches it.

That is why the right accounting unit is not cost per token alone. Not tokens per second alone. Not GPU utilization alone. Not benchmark speedup alone. The accounting unit Touchdown Labs cares about (and the metric Gen 3 is going to be judged by) is:

Cost per successful task at p95/p99 latency.

And behind that: energy per successful task.

visible_ai_cost =
  tokens + GPU time + CPU tool loops + retries + power + human review

useful_ai_output =
  accepted patch + resolved ticket + verified brief + customer workflow completed

For RL/post-training, the same accounting becomes sharper:

cost_per_successful_trusted_task =
  rollout_inference
  + verifier_eval
  + tool_runtime
  + retries
  + rejected_samples
  + stale_weight_samples
  + weight_sync_idle
  + human_review

divided by successful trusted outcomes

The denominator is the point. A system can look efficient by generating many samples, but if the samples are rejected, stale, unverifiable, reward-hacked, or impossible to replay, the business paid for inference without buying improvement. For a coding agent, the unit is cost per accepted patch. For a research agent, cost per verified brief. For support, cost per resolved ticket. For ops copilots, cost per audited workflow. For systems-code agents, cost per accepted kernel or config improvement.

For a single product team, that means gross margin per customer workflow. For a neocloud, hyperscaler, or investor underwriting power, it becomes successful tasks per megawatt and gross margin per megawatt. Same logic, bigger denominator. If two serving paths use the same 10 MW envelope but one completes 2x more verified customer tasks at the same p95/p99 and quality bar, the better path did not just save energy. It changed the economic capacity of the site.

SemiAnalysis' own usage data, published openly in the coding-assistant breakdown, sharpens this. Their annualized token spend already runs at roughly 30% of employee compensation, with just under 5B tokens per month per employee and some teammates exceeding 100B tokens per month. They estimate Claude Opus 4.7 agentic workloads at a true blended price near $0.99 per million tokens (despite sticker pricing of $5 input / $25 output) because Claude Code-style workloads have around 300:1 input/output ratios and 90%+ cache-hit rates. The sticker price is wrong by roughly 6×; the harness is doing the work. Same model. Different harness. Different cost per task.

Energy works the same way. A coding agent that burns 10 extra tool calls, reloads the same repo context five times, misses the prefix cache, and retries failed tests through bloated context is not just more expensive. It is physically wasteful. More GPU prefill. More CPU sandbox time. More memory bandwidth. More collective traffic. More cooling. More wall-clock infrastructure for the same patch.

The Kimi K2.5 GB200 row later gives the buyer version of the same idea. In one public InferenceX 8k/1k NVFP4 row family, a narrow GB200 Dynamo vLLM point reports 2,173 output tok/s/GPU, while a wider TP16/EP16 point reports 12,576 output tok/s/GPU. Those rows are not iso-latency and they are not Touchdown measurements. They are useful because they make the accounting visible: same workload family, different serving path, about 5.8x different output-token capacity per GPU at the reported operating points. That is the kind of gap a token dashboard will not explain by itself.

Water fits here too, but carefully. Not every data center uses water the same way. Not every workload has the same local footprint. But the physics are not complicated: wasted compute becomes heat, heat becomes cooling demand, and cooling demand can become water demand depending on the site.

AGENTIC TASK COST IS NOT JUST TOKENS

USER TASK
  ↓
SYSTEM PROMPT / TOOL SCHEMAS / REPO CONTEXT       (cache hit or miss)
  ↓
GPU PREFILL                                       (prefill FLOPs, KV alloc)
  ↓
GPU DECODE                                        (plan, tool calls, patches)
  ↓
MULTI-GPU / CACHE MOVEMENT                        (collectives, RDMA,
                                                   state placement when the
                                                   workload spans devices)
  ↓
CPU TOOL LOOP                                     (~42% of agentic time per
  - file edits, bash, tests, lints, installs,      SemiAnalysis coding-agent
    sandbox execution                              breakdown)
  ↓
RETRY / REPAIR LOOP                               (failed tests re-enter
  - failed test output re-enters context           context, more prefill,
                                                   more decode, KV grows)
  ↓
CACHE HIT OR MISS                                 (the difference between
                                                   "free" and full re-prefill)
  ↓
SUCCESSFUL TASK

The same evidence loop, two scales. The kernel loop optimizes one GPU program. The inference loop optimizes the full path from request to useful output. Same eight verbs, larger candidate.

KERNEL RL LOOP                                    AGENTIC INFERENCE LOOP
─────────────────────                              ─────────────────────────
generate candidate kernel                          user task
  ↓                                                  ↓
compile                                             load system prompt /
  ↓                                                  repo context / tool schemas
verify correctness                                    ↓
  ↓                                                  GPU prefill + KV allocation
benchmark on hardware                                 ↓
  ↓                                                  GPU decode
reward / reject                                       ↓
  ↓                                                  multi-GPU/cache movement
remember evidence                                     ↓
  ↓                                                  CPU tool work: files, bash, tests
improve                                               ↓
                                                    retry / repair / re-enter context
                                                     ↓
                                                    cache hit or miss
                                                     ↓
                                                    successful task or failure

The token price is falling. The token volume per useful task is rising faster. That paradox is what the broader market is currently learning the hard way, and it's the most direct version of why the third generation of inference optimization (§02.55) matters now.

The macro forecasts. Goldman Sachs Research forecasts that agentic AI could drive a 24× increase in token consumption by 2030, to roughly 120 quadrillion tokens per month. Gartner forecasts that inference on a 1T-parameter LLM will cost GenAI providers over 90% less by 2030 than in 2025: but explicitly warns that falling unit cost does not automatically democratize frontier reasoning, because agentic workloads can consume 5–30× more tokens per task than standard chatbot workloads, so aggregate enterprise AI spend can still rise. The unit cost falls. The work per task explodes. Per-token dashboards do not catch that gap.

Goldman's chart is useful because it does not treat all tokens as the same workload. Non-agent usage is one lane. Consumer agents are another. Enterprise agents are the lane where the unit economics get most interesting, because the task usually has tools, permissions, retrieval, retries, compliance checks, integration work, and latency constraints wrapped around it. Treat the chart as a forecast, not an audited outcome. But as a planning signal, it is hard to ignore.

For a CFO or CEO, the point is not "tokens go up." The point is that token growth turns into gross-margin pressure or capacity leverage depending on the workload path. If the extra tokens produce successful customer work at acceptable latency, they can support revenue. If they come from bloated context, failed retries, cache misses, tool loops, or bad routing, they become margin leak.

The engineering implication is simple: lower unit prices do not save you if the task path keeps expanding. The durable fix is to reduce wasted work inside the task.

Rerankers create hidden tokenomics too. The output may be one answer, but the system can tokenize and score hundreds of retrieved candidates before that answer exists. A cheaper generation token does not fix an inefficient reranking path, because the leak happened before the generator got the prompt.

Why cheaper tokens can still create bigger bills.

SemiAnalysis' Dark Output essay gives the macro version of the same tokenomics problem. The costs are visible first: API bills, GPU purchases, power, water, data centers, job displacement anxiety, capex. The useful output is harder to count because it can show up as faster decisions, fewer vendor hours, more code shipped, more experiments run, fewer meetings, better ops, or work that never existed before because it used to be too expensive.

I would split the inference version into three buckets. Substitution dark output: AI does work that a person, contractor, vendor, or internal team used to do, but the cost moved into tokens, compute, review, and retries. New dark output: AI makes teams do work they never would have done before, like running 200 design variants, checking every ticket, generating a draft audit, or spawning extra coding-agent attempts. Captured AI output: the value finally shows up somewhere measurable: lower support cost, more shipped features, higher retention, faster sales ops, better margin, or fewer escalations.

That distinction matters because token growth can mean productivity or waste. A million cheaper tokens can replace a real workflow, create new useful work, or just amplify context bloat and failed attempts. SemiAnalysis' exposed-labor framing is useful as a map of task potential; I would not read it as confirmed displacement or realized GDP impact. The only honest way through is still workload evidence: what task ran, what cost moved, what output was accepted, and what quality gate passed?

The micro case studies. Fortune reported in May 2026 that Microsoft began canceling most direct Claude Code licenses internally and standardizing engineers onto GitHub Copilot CLI by end of June; The Verge's operational coverage confirms the migration timeline. The Information reported that Uber burned through its entire 2026 AI coding-tools budget in roughly four months, largely tied to Claude Code usage that internal adoption leaderboards had actively incentivized. NVIDIA VP Bryan Catanzaro, quoted in the same Fortune piece, said that for his team compute costs are "far beyond" employee costs. This is the practical face of cheaper tokens, bigger bills. These are not stories about AI being useless. They are stories about agentic AI turning one user request into a long infrastructure trajectory, with no observability tying any of it to per-task cost.

The infrastructure scale that makes the problem urgent. OpenAI announced Stargate as a $500 billion / four-year AI infrastructure project with $100B in immediate deployment; Reuters has tracked the expansion to five new data centers targeting ~10 GW of capacity. McKinsey frames the broader build-out as a $7 trillion data-center race with roughly $5.2T tied directly to AI demand by 2030. Anthropic committed to $30B of Azure compute and up to 1 GW of additional capacity in the November 2025 Microsoft / NVIDIA / Anthropic strategic-partnership announcement. Reported separately and treated here as reporting rather than audited fact, Anthropic is also reported to be paying roughly $1.25B/month (around $15B/year) for additional compute capacity. Reporting on OpenAI's own projections implies losses tripling to roughly $14B in 2026; these are reported projections, not audited results, and worth treating as such. The pattern is not the specific number. It is that the AI capital cycle has moved from a software cycle to a power, data-center, chip, network, and inference-efficiency cycle, and the unit economics underneath have not yet caught up.

The power story is the same curve, except now the unit is physical.

IEA projects global data-center electricity consumption rising to around 945 TWh by 2030. DOE/LBNL estimates U.S. data-center electricity use grew from 58 TWh in 2014 to 176 TWh in 2023 and could reach 325–580 TWh by 2028. CBRE has called limited power availability a prime inhibitor of data-center growth in certain core markets.

The answer is not “data centers are bad.” That is lazy and wrong. Data centers are the physical base of the AI economy. The problem is the waste we send into them.

If a task burns unnecessary prefill, misses the prefix cache, runs on the wrong accelerator, retries through bloated context, spills KV state across a slow boundary, or leaves GPUs idle behind CPU tool loops, the data center still pays the power bill. The cooling system still has to move the heat. The customer still pays for the task. The user still waits.

This is why engineering decisions become energy decisions. Prompt churn, cache misses, bad routing, weak kernels, and failed retries all become power draw somewhere.

That is the environmental argument we actually believe: not less AI, less wasted AI.

More useful AI per watt.
More successful tasks per megawatt.

The honest read. The AI bubble does not pop because AI is useless. It pops if useful AI tasks are delivered through inefficient infrastructure paths for too long. If one completed task burns too much context, too much prefill, too much KV cache, too much CPU tool time, too many retries, and too much tail latency, then cheaper tokens do not save the economics; the task is still too expensive. VC money can subsidize bad unit economics for a while. Hyperscalers can subsidize inference for a while. Model labs can hide inefficient task paths behind growth for a while. But capital eventually gets disciplined. When it does, the only durable answer is full-stack optimization: kernels, engines, caches, routing, CPU tool loops, workload replay, and hardware-aware diagnostics, all tied back to cost per successful task.

The answer is not to stop building agents. The answer is to make the stack measurable and optimizable from the ground up. The correction, when it comes, is painful not because AI is useless: but because the infrastructure path to deliver useful AI work was too expensive to sustain at scale without observability. The concrete place we tested that discipline first was the hackathon environment.

There are a lot of research directions worth pursuing - different chips, different stages of the stack, different bets the field is taking. For the hackathon, we wanted to do something different and, honestly, more fun: pick a hard constraint (two days, one chip, one weekend), pull in the latest open-source work happening right now in automated kernel generation, layer in what we had learned from the public literature and our own experiments, and see how far we could push the loop. KernelForge-OpenEnv is the public artifact that came out. It is narrow on purpose: one model, one chip (A100 80GB sm_80 via Modal), real-graph baselines vendored in from NVIDIA's cuGraph as honest comparators, three-stage curriculum (SFT → RFT → GRPO) using TRL's OpenEnv integration. What follows is what we tested, what we learned, and what we got wrong.

This is not a claim that the hackathon repo solved production inference. It is the smaller, stricter claim: we built the kind of harness discipline a larger workload benchmark needs. The same discipline that checks one A100 kernel is what a Kimi K2.5 serving row needs at rack scale: recipe named, baseline named, hardware path named, correctness checked, replay preserved, and the caveats visible.

A100 kernel environment:
  candidate CUDA -> nvcc -> correctness -> timing -> reward -> evidence

Kimi K2.5 serving workload:
  request shape -> Dynamo/vLLM -> KV transfer -> GB200 fabric -> result row -> capacity math

The loop is what you'd expect. A model proposes a candidate CUDA kernel; nvcc compiles it for A100; the system loads it as a PyTorch extension against a reference implementation; correctness gets checked across a shape suite with dtype-aware tolerances; runtime is benchmarked with CUDA events, proper synchronization, warmup, and repeated runs with variance reported separately from the median; failure modes get captured alongside wins; a multi-objective reward gets computed with correctness as a binary gate; the agent updates via GRPO through TRL. Eight verbs: generate, compile, verify, benchmark, reward, reject, remember, improve. None of them is about the model. The model is the smallest, most replaceable piece.

The hypotheses we were actually testing.

Instead of calling these "architectural decisions," call them what they were: probes. Each row names which files in the repo carry the test, and where the test currently stands.

The repo, file by file - what it does, why it's shaped this way, where each piece points next.

Trimmed to the core files (full tree has 19 top-level dirs / ~120+ files; we kept it readable). Right-side comments carry what, why, the hypothesis it tests, and which adjacent open-source project the design borrows from or hands off to.

# A100-CUDA-RL (public; KernelForge-OpenEnv; A100 80GB via Modal)

A100-CUDA-RL/
│
├── openenv_env/                      # ─── Layer 1: the OpenEnv wrapper. (H1)
│   ├── server/app.py                 #     FastAPI reset / step / state surface.
│   │                                 #     Why: any TRL- or verifiers-compatible trainer drives the env
│   │                                 #     without surgery. Standards-compatible by construction.
│   ├── kernel_forge_env.py           #     The env class itself. Known import coupling to
│   │                                 #     training/task_support.py - pending cleanup per README.
│   ├── reward.py                     #     Correctness-first, bounded reward. (H4)
│   ├── eval_backend.py               #     Backend switch: modal | local | coreweave. (H2)
│   ├── anti_hack.py                  #     Reward-hacking detectors - first-class, not bolted on. (H4)
│   ├── skill_builder.py              #     Compose CUDA-Agent / DoubleGraph / SkyDiscover priors. (H6)
│   ├── cache_pool.py + task_pool.py  #     Pre-computed task batches + cached compile artifacts.
│   ├── gpu_registry.py               #     Hardware metadata registry per target.
│   ├── task_routing.py               #     Route tasks to backends + worker pools.
│   └── client.py + models.py         #     Pydantic-typed client + request/response models.
│
├── eval_service/                     # ─── Layer 2: pure execution backend. (H2)
│   ├── eval_core.py                  #     Shared pure compile / verify / benchmark core.
│   ├── app.py                        #     Optional FastAPI eval service for CoreWeave / Northflank.
│   └── Dockerfile                    #     Same image runs local + remote.
│
├── training/                         # ─── Layer 3: GRPO + curriculum. (H3, H5)
│   ├── grpo_train.py                 #     Main entrypoint: --stage stage1|stage2|stage3.
│   ├── custom_grpo_trainer.py        #     TRL-derived GRPO with custom rollout hooks. (H3)
│   ├── multi_turn_rollout.py         #     Local compile fast-fail + remote reward dispatch. (H3)
│   ├── stage0_sft.py                 #     Stage 0: basic CUDA SFT. (H5)
│   ├── stage1_warmup.py              #     Stage 1: warmup → enhanced GRPO. (H5)
│   ├── stage2_rft.py                 #     Stage 2: RFT/SFT continuation. (H5)
│   ├── stage3_grpo.py                #     Stage 3: full GRPO pilot - still pending live A100 run.
│   ├── curriculum.py                 #     Reads configs/scaling_ladder.json.
│   ├── cuda_agent_integration.py     #     Wire CUDA-Agent prior into rollouts. (H6)
│   ├── grpo_config.py + rft_filter.py + dataset_loader.py + model_loader.py + run_metadata.py
│   └── task_support.py               #     Task spec helpers (the coupling point with openenv_env).
│
├── verification/                     # ─── Correctness gates. (H4)
│   ├── pac_verify.py                 #     PAC-style verifier - Dr. Kernel-style property checks.
│   └── profile.py                    #     NCU / nvprof wrapper.
│
├── evaluation/                       # ─── Post-training analysis.
│   ├── eval_model.py + pass_at_k.py  #     Held-out eval + pass@k metric per shape.
│   ├── ablation.py + compare_stages.py  # Stage-by-stage delta analysis.
│   ├── reward_monitor.py             #     Live reward-drift detector.
│   └── compiler.py + profiler.py + sandbox.py + verifier.py  # Modular eval primitives.
│
├── skydiscover_integration/          # ─── Adjacent search/evolution loop. (H6)
│   ├── adaevolve.py + evox_strategies.py + evaluator.py + run_evolution.sh
│   └── initial_kernels/              #     Seed kernels: bfs_power_law / louvain_community /
│                                     #     pagerank_dense / triangle_count / wcc_doublegraph (.cu)
│
├── kernels/                          # ─── Reference + baseline CUDA. (H7)
│   ├── baseline_wcc.cu               #     cuGraph WCC baseline to beat.
│   ├── clustered_wcc_h100.cu         #     Optimized comparator.
│   └── ecl_cc_h100.cu                #     ECL-CC reference connected-components kernel.
│
├── doublegraph_src/                  # ─── Vendor reference (NVIDIA cuGraph fork). (H7)
│                                     #     Why fork in: pin the comparator instead of chasing a moving
│                                     #     vendor library. Not our code - included as honest ground truth.
│
├── datasets/                         # ─── Training corpora.
│   ├── basic_cuda_sft.jsonl + create_basic_sft.py        # Stage 0 seed data.
│   ├── doublegraph_sft.jsonl + extract_doublegraph_a100.py + prepare_doublegraph_sft.py
│   ├── combined_kernelforge.jsonl + build_combined_dataset.py
│   └── integrity.py                  #     Dataset checksum + dedupe gate.
│
├── tasks/build_task_pool.py          # ─── Builds the eval task pool from kernels.
│
├── tests/                            # ─── 17 test files. The harness for the harness. (H4)
│   └── test_env / test_compile / test_reward / test_pac_verify / test_anti_hack /
│       test_reward_monitor / test_pass_at_k / test_curriculum / test_grpo_config /
│       test_multi_turn_rollout / test_cache_pool / test_gpu_registry / test_model_loader /
│       test_model_registry / test_skill_builder / test_task_support / conftest
│
├── docs/                             # ─── Source-of-truth design docs (the hypothesis register on disk).
│   ├── KERNELFORGE_FINAL_PRD.md      #     PRD: architecture, rollout order, launch criteria.
│   ├── GRPO_DEEP_DIVE.md             #     GRPO/TRLOO strategy + rollout details.
│   ├── KERNELFORGE_RL_ENVIRONMENT.md #     The env spec.
│   ├── OPENENV_AUDIT_PLAN_3.md       #     Contract review + architecture audit.
│   ├── EXPERIMENT_REPORT_9B_STAGE1.md #    What we actually ran + what it told us.
│   ├── GRPO_TRAINING_STATUS.md       #     Live status of each stage.
│   ├── SYSTEM_TRUTH.md               #     What is real vs aspirational. We re-read this often.
│   ├── rl_pipeline_debug_review.md   #     Postmortem of pipeline failures.
│   ├── changelog/001..008_*.md       #     8 dated changelogs (Mar 8 → Mar 11) - bring-up to revert.
│   └── skills/                       #     Open-source priors carried as docs:
│                                     #       CUDA_AGENT.md · DOUBLEGRAPH_A100.md ·
│                                     #       KERNELGYM_DR_KERNEL.md · SKYDISCOVER_ADAEVOLVE_EVOX.md
│
├── .claude/                          # ─── Agent operating manual for working on this repo.
│   ├── skills/cuda-rl-runbook/       #     Runbook the agent reads before touching the loop.
│   ├── skills/grpo-reward-pipeline-debugger/  # Reward-pipeline debug skill.
│   ├── skills/trl-first-step-hang-debugger/   # "TRL hangs on first step" - a real failure, real debug skill.
│   ├── agents/cuda-rl-auditor.md     #     Auditor agent definition.
│   └── commands/rl-audit.md + rl-smoke.md
│
├── notebooks/kernelforge_training.ipynb + gpt_oss_(20B)_GRPO.ipynb
├── modal_app.py + modal_train.py     #     Modal entrypoints. (H2)
├── openenv.yaml                      #     OpenEnv contract file. (H1)
├── skill_a100.md + CLAUDE.md         #     A100-specific agent skill + repo memory.
├── Dockerfile + docker-compose.yml + pyproject.toml + requirements.txt + uv.lock
└── README.md

The second tree: how we think about architecture-specific R3/R4 work.

The public hackathon repo above was intentionally narrow: A100, OpenEnv, Modal, CUDA, cuGraph-backed graph kernels. But the design pressure underneath it was not A100-only. The real question we kept coming back to was how this loop should scale when the target stops being a generic CUDA kernel and starts being architecture-specific work: Blackwell on one side, AMD CDNA4 / MI355X on the other, and eventually whatever comes after both of them.

AMD CDNA4 is the best place to explain the shape. The rocm-kernels skill gives agents a useful entry point: scaffold a kernel-builder project, register Torch ops, write a benchmark, and package the result. That is the right first layer. It is not enough for deep R3/R4 work. On MI355X, the work quickly gets more specific: CuTe-shaped tile code in Gluon / TileLang / HipKittens at R3, then HIP plus inline AMDGPU assembly at R4 for instructions like v_mfma_scale_f32_32x32x64_f8f6f4 and ds_read_b64_tr_b4. A single skill cannot carry the whole architecture without becoming unreadable. That is not a criticism of the skill. It is a fact about the architecture.

So the architecture has to split. The entry skill stays small. The operation contract lives one layer below it. The hard ISA facts live below that. The long vendor reference gets externalized into a loader. The experiment history becomes append-only. That is the pattern we were really testing in miniature with KernelForge-OpenEnv: do not make the model remember the whole world; give it a narrow action surface, an honest harness, and enough externalized context to avoid repeating the same mistakes.

# One layout pattern for AMD CDNA4 / MI355X R3-R4 work.
# The point is the architecture, not these exact filenames.

amd-cdna4-stack/
│
├── resources/                         # Vendor primary sources, version-pinned.
│   ├── cdna4-isa.pdf                  # AMD's open CDNA4 ISA manual.
│   ├── cdna4-isa.xml                  # Machine-readable ISA via GPUOpen.
│   ├── op-contracts/                  # One contract per operation family.
│   │   ├── mxfp4_gemm.md              # E8M0 scales, tile contract, reference baseline.
│   │   ├── mxfp4_moe.md               # Routing, expert layout, scatter/gather contract.
│   │   └── mla_decode.md              # Latent KV shape, RoPE split, decode pattern.
│   └── op-references/                 # Vendor + public reference implementations.
│
├── isa-corpus/                        # Page-addressable parsed ISA.
│   ├── markdown/                      # Per-chapter markdown from the 565-page ISA.
│   └── manifest.json                  # Chapter index for programmatic loading.
│
├── isa-loader/                        # RLM-style interface to the long reference.
│   ├── corpus.py                      # Loads chapters only when needed.
│   ├── retrieval.py                   # Builds focused context packets.
│   └── topics.yaml                    # topic -> chapters, terms, facts.
│                                      # mfma_scale, lds_tiling, ds_ops,
│                                      # vector_memory, waits, wavefront64.
│
├── rules/                             # Hard constraints extracted from the ISA.
│   ├── gfx950-compute.md              # MFMA encoding, CBSZ, BLGP, cycles.
│   ├── gfx950-memory.md               # LDS bank rules, transpose-read layout.
│   └── op-specific/                   # One rule file per op family.
│
├── skills/                            # Narrow, composable, op-aware skills.
│   ├── rocm-kernels-entry/             # Scaffolding + build + packaging path.
│   ├── cdna4-hardware-principles/      # Wave64, occupancy, MFMA-centric thinking.
│   ├── fp4-gemm-cdna4/                # Shared FP4 / MXFP4 GEMM primitive.
│   ├── hip-inline-asm-bringup/         # R4 inline asm correctness loop.
│   ├── lds-transpose-staging/          # ds_read_b64_tr_b4 layout patterns.
│   └── per-op-skills/                  # MoE, MLA decode, attention, reductions.
│
├── harness/                           # The same eight-verb loop, target-specific.
│   ├── compile.py                      # hipcc / clang / flags / arch detection.
│   ├── verify.py                       # Shape suite, dtype tolerances, NaN gates.
│   ├── benchmark.py                    # warmup, sync, variance, baseline diff.
│   ├── profile.py                      # rocprof-compute / counters / occupancy.
│   └── evidence.py                     # Emits portable EvidencePackets.
│
└── experiment-log/                    # The profile. Append-only.
    ├── history.jsonl                  # One row per attempt, pass or fail.
    ├── memory.json                    # Proven facts + dead-end summaries.
    └── frontier.csv                   # Tried vs untried search axes.

Why the layers matter. The top-level skill is for scaffolding and common muscle memory. The per-op skill is where the real contract lives: tensor meanings, legal shapes, baseline choice, correctness tolerance, and whether the op is actually representative of production inference. The rules/gfx950-* files carry facts the proposer is not allowed to hallucinate away: MFMA operand layout, LDS transpose rules, wait counters, register pressure, wavefront behavior, and the difference between a kernel that compiles and a kernel that is shaped for the chip. The isa-loader keeps the full vendor reference out of the prompt while still making it reachable. The experiment-log prevents every agent from rediscovering the same dead ends.

This is the same idea as the A100 repo, just one level more honest about what architecture-specific work requires. The A100 repo says: keep the env small, keep the backend swappable, keep failures as evidence. The CDNA4 pattern says: also make the ISA queryable, make the hard rules explicit, make operation contracts first-class, and let the harness read profiler truth instead of vibes. That is how a kernel engineer should read the tree. The files are not bureaucracy. They are where the reasoning goes so the next run can start smarter than the last one.

The open AMD ecosystem makes this possible in a way that is genuinely interesting. AMD publishes the CDNA4 ISA in human-readable and machine-readable form. Gluon gives the field a gfx950-specific tile DSL with real MXFP4 paths. Amanzhol Salykova's matrix-core writing explains the MFMA-scale path clearly. RadeonFlow Kernels gives public FP8 / MoE / MLA references. AMD's GEAK shows the vendor is exploring agentic kernel generation directly. SGLang and llama.cpp are already pulling MI355X support into real inference paths. The important point is not "AMD good, NVIDIA bad." Public instruction-level truth changes what an agentic kernel stack can be. If the architecture is open enough to inspect, the harness can become more than a black-box timer.

That is the thought process behind the whole section. The weekend artifact was the small version. The A100 tree shows what we actually shipped. The CDNA4 tree shows the shape we think serious architecture-specific kernel work needs. Both are the same bet: model-written kernels are getting better, but the system around them is what makes the improvement real.

In the spirit of the hackathon.

There are a lot of research directions worth pursuing - different chips, different stages of the stack, different bets the field is taking. For two days we wanted to do something different and, honestly, more fun: pull together the latest open-source work happening right now in automated kernel generation, combine it with what the public literature and our own experiments had already taught us, and see how far we could push the boundary inside a weekend.

That meant building on OpenEnv and TRL's GRPO integration as the outer wrapper and trainer, a KernelGYM-style execution backend with a Dr. Kernel-style PAC verifier, documented priors from CUDA-Agent, doubleAI's DoubleGraph, and Berkeley's SkyDiscover / AdaEvolve / EvoX, NVIDIA cuGraph version-pinned as ground truth, and Modal on the back end so the reward signal actually came from real A100 silicon inside a weekend. The repo is small because we weren't trying to reinvent the interesting parts. We were trying to see which joints actually hold when the best of what's out there gets wired together honestly under a real constraint.

What the weekend actually taught us.

Three things, said honestly and in order of how much they changed our thinking.

One. The layered split - OpenEnv wrapper on top, KernelGYM-style execution backend on the bottom, our task/reward/rollout logic in the middle - held up better than expected under a real GRPO loop. The boundary we were most worried about was the env-to-backend boundary, because everyone we'd read making this kind of system had ended up with an entangled trainer-knows-about-backend mess. Splitting it via the KERNELFORGE_EVAL_BACKEND=modal|local|coreweave env-var lets us swap Modal for a local smoke run for a CoreWeave HTTP service without the wrapper noticing, and that turned out to be the single decision that made the weekend tractable. A v1 should make that selection a typed config object instead of an env var, but the factoring itself is right.

Two. Anti-reward-hacking has to live next to the reward function, not in a postmortem. We knew this going in - every team in the §08 cohort that's shipped a paper on agentic kernel-gen has the same story - but actually putting anti_hack.py next to reward.py with its own tests, and putting pac_verify.py in a separate verification/ module, forced us to think about each failure mode at code-review time instead of at training-collapse time. The flip side, and the honest critique: the import coupling between openenv_env/kernel_forge_env.py and training/task_support.py means the env isn't quite neutral yet. The README already flags this. A v1 needs the env to read curriculum as data, not import from the trainer's task helpers.

Three. The thing we'd defend most, and the thing that surprised us least: failed attempts have to emit packets too. A compile error, a NaN, a shape mismatch, a tolerance miss - all of them produce evidence with a populated validation category and a null performance category. Nothing gets silently dropped. The failure mode of a rejected candidate is the most valuable training signal in the system, and discarding it is how harnesses quietly mislead trainers. Stage 1 ran on a 9B model (the writeup is in docs/EXPERIMENT_REPORT_9B_STAGE1.md), and the most useful debugging signal we got from the run was the rejection histogram, not the speedup distribution. That's the result we'd carry forward into anything that comes after this.

What we'd be honest about. Stages 2 and 3 are code paths until they're not. The README's "What Still Needs Live Validation" section lists them; we're not going to claim what hasn't shipped. The A100 is one chip; the hardware picture this post develops in §07.5 spans sm_100, sm_120, and gfx950, and the env-backend split is designed for multi-arch but only one architecture is live so far. And the bridge from evaluation/reward_monitor.py + verification/profile.py to the open evidence schema this post argues for in §13 exists as an intent, not as a hooked-up artifact yet. Those gaps are the work that comes next.

The larger thing, said briefly. The hackathon was a way to ship something honest under a hard constraint by combining work we admired from other teams. That's what came out of the weekend. The rest of this post is what we'd say about the rest of the problem - built on the same posture, with different pieces.

The tradeoffs we made (which is to say, not many). We held the schema minimal on purpose; we shipped stubs instead of a full hardware matrix on purpose; we baselined against torch.compile instead of cuBLAS on purpose, because two days is not enough time to claim CUTLASS-class wins and a weak baseline is more honest than a misleading one. The places we'd put more time, given more weekend: a properly-wired compile harness for B200 and MI355X out of the gate; richer profiler-trace plumbing into the artifacts category (NCU sections, rocprof-compute dumps); a built-in shape generator for the Mixtral / DeepSeek MoE expert distributions. Those are work, not architecture decisions. The schema doesn't change.

Where this points next: auto-discovery and SkyDiscover integration. The piece we're already building on top of the v0.1 schema is automated workload and hyperparameter discovery, in the spirit of SkyDiscover from UC Berkeley Sky Lab. SkyDiscover already showed the loop works at the systems-research layer: automated discovery of cross-cloud transfer optimizations (41% cost reduction on a 300 GB GCP-Singapore-to-six-cloud-destinations broadcast benchmark), MoE GPU load-balance configurations (14% better load balance via a Webster's-method-plus-zigzag-packing algorithm), and KV-cache pressure mitigations (29% lower KVPR via a binary-search + best-fit-decreasing placement strategy). The move is the same as ours: a probabilistic generator proposes, an externalized environment evaluates, an honest measurement decides. Where SkyDiscover proposes systems-level configurations and reads from cluster traces, we propose kernel candidates and read from EvidencePackets. Plugging them into each other is the obvious next experiment: let SkyDiscover-style automated discovery drive the candidate field with hardware-aware proposals (target this GPU, this dtype, this shape distribution; explore around this PR-known-good baseline), then let the KernelEvidenceEnvironment emit the evidence packets it reads back. The schema is built to be the interchange format for exactly that kind of cross-system loop. The whole point of holding the spec minimal and the categories self-describing is so that a discovery loop someone else writes can drive this environment without us shipping a new release. That's the version of "auto-discovery" we care about (not "automate the human out of the loop," but "externalize what one team learned so the next team's loop reads from honest signal.")

What we actually saw, honestly, was what every team in this space sees on day one. The first few hundred rollouts produced kernels that did one of three things: compile-fail loudly (missing header, malformed launch config, wrong PTX intrinsic; the easy class), compile-pass and silently fail correctness (off-by-one indexing, missing boundary mask, dtype promotion in the accumulator path; the dangerous class), or compile-pass and pass correctness on a single shape but blow up on the next one in the suite (the harness-is-the-research class). The whole point of the eight-verb loop was that the third category showed up as soon as we widened the shape suite, exactly when the harness needed to catch it. A kernel that wins on one shape and dies on the next is not a real win. The reward signal only became useful once the harness was strict enough to reject that failure at verify, not benchmark. Small operational fact, and the single biggest thing the weekend convinced us about the system.

OpenEnv was the right framework for this because it forces the discipline on you. It exposes environments as HTTP services with a state-machine interface (reset(), step(action), and structured observations) and TRL talks to those endpoints during rollout. You can't carry hidden state through an HTTP boundary or quietly defer evaluation until later. You have to declare your observation schema and make the environment legible to a trainer you don't control. For a hackathon, that constraint turns out to be most of the point. The environment has to be coherent or training doesn't work. That's what "observability is first class" feels like when the framework asks it of you directly.

That's the system. The model wrote one source file per attempt; the harness did everything else. So the next question is not whether the model can generate code. It is whether the harness can be trusted when the optimizer starts searching against it.

The harness is not the boring part after the research. In RL, it is the world the model acts inside. It decides what counts as an action, what gets observed, what gets rewarded, what gets rejected, and what gets remembered for the next attempt. The harness is the business control surface for reward. If the harness is vague, the model still learns. It just learns the wrong thing with confidence.

That is the CEO/CFO version of the problem. Do not pay for model progress you cannot replay. A closed eval script, a scalar reward, a missing failure log, or an unclear rollout worker can all make the curve move while the real system does not improve. For a CFO, that becomes wasted training runs, wasted inference capacity, and fake ROI. For a CTO, it becomes an architecture decision built on a result nobody can reproduce. For an engineer, it becomes weeks of debugging a reward signal that was never tied to the real environment.

For the gigawatt version of the argument, a weak harness is a capacity-allocation bug. If the reward says a rollout improved but the replay would have failed on the real workload, the team is effectively turning scarce power into fake progress. The harness has to prove not only "the model got a higher score", but which path produced more successful tasks per megawatt, under which workload, with which caveats.

weak RL harness:
  closed eval script
  scalar reward only
  missing failure packets
  unclear worker/config/weight state
  no replay command

real RL environment:
  task + rules + tools + sandbox
  reward/verifier + failure classifier
  model/weight version + rollout worker + serving path
  hardware/runtime/profiler evidence
  replay packet another team can run

harness threat model:
  same-process Python eval risk
  stream synchronization / timing risk
  evaluator-phase leakage
  monkeypatching and reflection risk
  cached-output or side-effect risk
  sandbox / filesystem / network boundary risk

Reward hacking is usually a specification failure before it is a model failure.

There is a public-safe way to say the blunt version: when a model "hacks" the reward, the environment often failed to say what the real task was. It did not fully specify success, failure, side effects, sandbox boundaries, input distribution, timing rules, or the spirit of the rules. The model found the cheapest path through the written contract, not the intended one. That is annoying, but it is also useful. It tells you exactly where the contract was too vague.

There is a caveat. Some reward hacking is genuinely adversarial, surprising, or outside what a human problem author would naturally anticipate. But the practical fix is still systems work: better environment design, stronger isolation, adversarial evaluation, richer failure packets, and an auditor that treats every unexpected success as a possible harness bug until proven otherwise.

problem author
  -> defines task, shapes, dtype, tolerance, allowed APIs, timing rules, and spirit of the benchmark
competitor
  -> writes the candidate kernel or systems patch
cheater
  -> probes weak baselines, stream timing, Python state, caching, sandbox leaks, and side effects
auditor
  -> hardens the harness, updates the threat model, and records the verdict
evidence packet
  -> preserves the exact result, failure class, replay command, profiler path, and cost/success receipt

This is where DSPy, GEPA, and RLM stop being abstract research names. Humans cannot manually write perfect specs for every kernel, agent workflow, or serving path. Prompt optimizers and reflective search help turn examples, failures, reward signals, and adversarial traces into better specifications. RLM-style externalized context lets an agent inspect the long evidence: ISA docs, PTX, SASS, compiler output, profiler logs, failed attempts, and weird reward cases that do not fit into a normal prompt. The point is not that the optimizer replaces judgment. The point is that judgment needs a better system to run through.

The number of ways a kernel-generation harness can quietly mislead you is a little unreasonable. Float math isn't associative, so a kernel that compiles cleanly can produce wrong output that a loose tolerance check happily passes. Async streams without proper event synchronization measure queueing latency, not kernel time. JIT caches and persistence mode reward kernels that exploit the loop instead of the hardware. A speedup can vanish on the second run because of thermal throttling, ECC recovery, NUMA migration, or a noisy neighbor on the box. Each is an observability gap first, and each produces a training signal that looks perfectly plausible right up until the model ends up confidently wrong.

In messier RL-on-LLMs domains, like coding agents with test suites, math with answer keys, and tool use with verifiable outputs, the same problems show up with more subjective judgment. Kernels are stricter. The candidate has to compile, run on the target hardware, match the reference, beat the right baseline, and survive the shape suite. The reward has fewer places to hide.

That is what makes kernel generation useful as a teacher. The success metric is narrow enough to audit, and the failure modes are concrete enough to name. A weak harness cannot hide behind soft evaluation. If your training signal is leaking, you find out. If your baseline is weak, you find out. If your timing methodology is off, you find out. The harness is where the systems work has to live, because in a verifiable domain a harness bug propagates straight into the reward, and the model learns the bug instead of the hardware. That is also why harnesses cannot stay as closed one-off scripts; they have to become shared infrastructure.

The production version is the same problem with more moving pieces. A throughput row only becomes useful when the public recipe says what actually ran: model, precision, engine, tensor parallelism, expert parallelism, concurrency, input/output shape, hardware path, and replay artifacts. Without that, a number is just a number. With it, a CEO can see capacity, a CFO can model cost, a CTO can challenge the architecture, and an engineer can reproduce the path.

bad benchmark:
  "we got more tokens/sec"

workload receipt:
  exact model + exact serving recipe + exact hardware path
  exact load shape + exact result artifact + exact caveat

That is why the next section talks about Prime Intellect, verifiers, RLMEnv, prime-rl, slime, and SGLang. They are the larger version of the same harness problem: define the environment, distribute it, run rollouts, verify workers, move weights, serve trajectories, preserve cache behavior, and emit evidence another system can replay. Kernel-evidence is the hardware-specific slice of that same contract.

Start from the product, not the framework. A coding agent needs accepted patches. A research agent needs decision-grade briefs. A support agent needs resolved tickets. An ops copilot needs audited workflows. A systems-code agent needs accepted kernels, configs, or runtime changes. In every case, the system improves by running attempts through an environment: generate, verify, reject or retry, score, update, and replay. That is why RL/post-training is basically inference workload infrastructure. Slime, SGLang, Miles, RadixArk, Prime Intellect, Mercor/OpenEnv, speculative decoding, and RLM loops are different parts of the same attempt factory.

Section 04 named the harness problem. This section separates the two places that problem shows up. The first layer is automated kernel generation. A model proposes CUDA, Triton, HIP, or TileLang code, and the environment decides whether that code survived reality. The second layer is RL rollout inference. A model generates trajectories through a serving system, and that serving system shapes the data the trainer learns from. Same environment pattern. Different candidate.

The connection is simple: if the model learns from execution, the execution path has to be measurable, replayable, and trustworthy. That is true when the candidate is a kernel. It is also true when the candidate is a rollout-serving path, cache policy, routing rule, or weight-sync strategy. The runtime changes. The reward changes. The evidence discipline does not.

prompt/spec
  -> rollout request
  -> serving engine
  -> draft/verify path (EAGLE-3 / DFlash / SSD)
  -> tool calls + verifier calls
  -> reward + trace
  -> RLM / GEPA reflection
  -> updated spec, cache policy, route, or kernel search

The model improves only as fast as the environment can produce trustworthy feedback. That is the unifier. Kernel RL, RLM-style agents, and RL rollout systems look different at the surface, but each one is a loop where a candidate acts inside an environment and learns from the evidence that comes back.

Those are not three separate theses. They are one environment-feedback pattern with different candidates. Touchdown's lane is the observation contract underneath them: what ran, where it ran, which state it used, whether it worked, what it cost, and how another system can replay it.

The systems to learn from, by role - not by name-drop.

The right way to discuss the outside work is to name the system problem first, then name the research, framework, or team solving that problem. Otherwise this section turns into a project directory and loses the point. The point is resource-constrained feedback: expensive environments, expensive rollouts, expensive verification, and the need to make every attempt produce evidence.

Layer one is the strict version: RL environments for automated kernel generation. A kernel-generation environment has a narrow job: candidate code in, compile, verify correctness, benchmark on target hardware, compare against the baseline, reward or reject, preserve evidence. The kernel environment is the strict version. The model can hallucinate. The compiler cannot. The model can claim a speedup. The profiler can check it. The model can produce code that looks right. The verifier can reject it.

For that kernel task, the environment is concrete: compiler, sandbox, GPU runtime, correctness checker, profiler, baseline. The reward is correctness first, speed second. The evidence is source code, compile logs, failure class, timing result, profiler trace, and replay command. That is why kernels are such a good proving ground. The environment can reject fake progress instead of turning it into a better-looking learning curve.

Layer one starts with reusable RL environments. If every team builds its own private RL environment, the field cannot compare results. One team's math verifier, coding sandbox, tool-use environment, CUDA harness, reward function, and replay format stay trapped inside one repo. That makes the model result hard to trust and harder to reproduce. This is where Prime Intellect, verifiers, and Environments Hub fit: not as a name-drop, but as environment infrastructure. Prime's verifiers docs describe a library for creating environments to train and evaluate LLMs. Their environment shape includes task inputs, a model harness with tools/sandboxes/context management, and a reward function or rubric, with integration into Environments Hub and prime-rl. Put that next to KernelForge/OpenEnv and the mapping is direct: op and shape suite, compiler harness, GPU runtime, profiler, correctness checker, baseline, reward, evidence packet.

That is also why Prime's Environments Hub and INTELLECT-3 matter. If a kernel environment stays trapped inside one repo, only one team learns from it. A reusable environment can be installed, inspected, forked, trained against, evaluated, and improved by other teams. A portable environment gives the field something to build on. The point is not that Prime built our CUDA harness. The point is that Prime is making environment infrastructure reusable, and kernel generation needs that same reusable environment discipline. The people belong after the contribution: the INTELLECT-3 report lists a broader Prime team including Mika Senghaas, Fares Obeid, Sami Jaghouar, William Brown, Jack Min Ong, Daniel Auras, Justus Mattern, Manveer Basra, Ameen Patel, and others, but the system layer is the lead: reusable environments, Environments Hub, and prime-rl.

The Will Brown / verifiers detail is worth putting back because it is the observability layer, not just an environment wrapper. The verifiers repository describes an environment as the package that owns the dataset, harness, tools, sandboxes, context management, reward function, and rubric. Its citation names William Brown as the original creator, and Will's own research page lists Verifiers: Environments for LLM Reinforcement Learning and PRIME-RL: Distributed RL Training at Scale as selected projects. That is the missing concrete point: the environment is not a loose eval script. It is the object that decides what the model saw, what tools it could use, what state existed, what reward was computed, and what evidence survived.

The release notes make the observability direction even more explicit. verifiers added trajectory-based tracking for token-in/token-out training across turns, truncated and branching rollouts, improved rollout and token tracking, monitor rubrics for automatic metric collection, OpenEnv and BrowserEnv integrations, RLMEnv improvements, per-rollout wall-clock timeouts, per-turn timing, GEPA prompt artifacts, and a v1 Taskset/Harness API. Those are not cosmetic features. They are the difference between "the model got reward 1.0" and "this exact multi-turn trajectory, with this prompt, tool state, timeout, branch, reward function, and environment version, produced this result." That is RL observability.

environment package
  -> taskset / dataset row / prompt
  -> harness / model endpoint / sampling config
  -> tools / sandbox / context manager
  -> multi-turn trajectory
  -> per-turn timing + token accounting
  -> reward function / rubric / monitor metrics
  -> timeout / truncation / branch / cleanup state
  -> environment version + replay command

RLMEnv connects to the context side of Layer 1. A serious kernel loop should not paste every compile log, profiler trace, failed candidate, baseline run, shape suite, and hardware rule into one giant prompt and hope the model holds it together. The model should inspect evidence, ask focused sub-questions, and come back with an action the environment can test. RLMEnv turns long context into something the environment can manage instead of something the prompt has to swallow.

The flywheel runs both directions. Faster, cheaper inference makes RLM-style recursion practical because every recursive subcall, verifier pass, trace summary, and repair attempt has a real latency and cost. But RLM also makes inference optimization better: it lets the agent inspect the evidence store, compare failed runs, summarize profiler traces, remember what changed, and propose the next experiment instead of starting from a blank prompt. Inference makes RLM loops affordable; RLM makes inference loops self-improving.

That is the self-improving systems loop: inference produces traces, RLM reads the traces, the environment spec improves, the verifier gets sharper, the next rollout wastes fewer attempts, and the evidence packet becomes easier to replay. The model is not improving from vibes. It is improving because the loop preserves enough state to know which failed attempt should change the next one.

This is the practical version of the raw speech-to-spec argument. Humans will not manually write perfect 1,000-word prompts, verifier contracts, DAGs, and failure checklists for every rollout forever. We get tired. We optimize for speed. We leave out the thing that would have prevented the shortcut. RLM, GEPA, DSPy-style optimization, and trace-aware agents matter because they can turn failures back into better specifications: what context was missing, which verifier was too weak, which reward was gamed, which serving path changed the result, and which next experiment is worth running.

SparseSpec adds the missing inference-side receipt for this loop. The SparseSpec paper uses RLM to mean reasoning language model, not Recursive Language Model. That distinction matters. SparseSpec is not a Recursive Language Models paper. It is a paper about making long-output reasoning-model inference cheaper: the kind of inference that recursive agents, RL rollouts, and dynamic workflows spend constantly. The connection is direct: cheaper long reasoning makes deeper recursive/RL loops affordable; better recursive/RL evidence makes the next reasoning attempt less wasteful.

reasoning task
  -> long generated trace
  -> KV cache growth
  -> SparseSpec sparse draft via PillarAttn
  -> full-attention verification
  -> accepted tokens
  -> verifier/reward + Data Buffer
  -> p95/p99 + GPU-hour + energy proxy per successful reasoning task

This is why long reasoning matters for post-training economics. The plan note is right: rollout generation can dominate RL wall-clock, and some RL setups report rollout generation above 90% of end-to-end training time. Treat that as workload-dependent, not universal. But the direction is obvious enough: if the system spends most of its improvement budget generating attempts, then decode-side memory traffic, KV-cache reads, sparse drafting, verification scheduling, and accepted-token rate become part of the learning infrastructure. RL capability is constrained by how many trustworthy long reasoning traces the serving system can afford.

The bridge out of Layer 1 is distributed RL. prime-rl takes the reusable environment idea into large-scale training and async rollouts. The important part is the system contract: open-source large-scale RL, single node to thousands of GPUs, agentic RL, multi-turn tool use, async rollouts, native verifiers environments, and Environments Hub integration. Kernel rollouts are messy. Some fail immediately. Some hang in nvcc or hipcc. Some pass correctness and lose performance. Some need a second turn with compiler or profiler feedback. Async rollout infrastructure matters because real environments do not all finish at the same speed. INTELLECT-3 reports scaling that style of infrastructure to 512 H200s; the lesson is not "more hardware solves RL," it is that more hardware raises the cost of a bad environment.

At scale, the environment also has a trust problem. It is not enough to say "a worker produced this rollout." The trainer needs to know which worker ran it, which weights were active, what state moved, what came back, and whether the result is trustworthy. The research contribution to lead with is rollout trust and weight movement: TOPLOC verifies rollouts from untrusted inference workers, while the INTELLECT-2/Prime line also frames SHARDCAST-style weight broadcast from training nodes to inference workers and PCCL-style distributed communication as part of the same boundary. This is where Jack Min Ong fits: not as a name first, but as one of the researchers on the Prime Intellect line around decentralized RL, rollout trust, and distributed-training infrastructure. A rollout is only useful if the trainer can trust where it came from.

The Jack Min / TOPLOC / SHARDCAST line is the distributed version of our A100 evidence problem. Locally, the question is: did this kernel compile, run on the claimed GPU, pass correctness, beat the baseline, and leave a replay packet? In decentralized RL, the same question becomes: did this remote worker run the claimed model weights, inside the claimed environment version, produce the claimed trajectory, and return a reward the trainer can safely learn from? Once the worker is not fully trusted, observability has to become verifiability. worker_id, weight_hash, environment_hash, rollout_proof, reward_signature, replay_seed, timing, and failure class are not nice-to-have metadata. They are part of the training signal.

The next pressure point is serving. Once rollouts are long-context, tool-using, MoE-heavy, cache-sensitive, and running across distributed workers, the environment is no longer just a Python object with a reward function. It is also an inference system. That is why Ameen Patel fits in this section. NVIDIA's Dynamo materials list "Learnings From NVIDIA Dynamo Users Featuring Prime Intellect" with Ameen Patel from Prime Intellect; the session description names engine-agnostic orchestration, disaggregated serving, observability, LoRA at scale, higher GPU utilization, predictable SLOs, faster iteration, and Kubernetes deployment lessons. INTELLECT-3 also includes Ameen Patel on the Prime Intellect team. The useful connection is simple: agentic rollouts, changing weights, sparse MoEs, cache state, adapters, disaggregated serving, SLOs, and distributed inference fleets all shape the reward signal. The rollout system becomes part of what the model learns from.

Layer two is inference infrastructure for RL rollouts. Once we move from automated kernel generation to post-training RL, the candidate is no longer just code. The candidate might be a rollout function, serving-engine config, cache policy, router policy, prefill/decode split, speculative decode path, weight-sync interval, partial rollout strategy, or tool-call scheduling policy. The reward is no longer only kernel speedup. It becomes task success, rollout quality, throughput, freshness, latency, cost, p95/p99, and failure rate.

Lead with SGLang, not slime. SGLang is the inference/runtime layer for structured, cache-heavy rollout traffic: structured outputs, tool calls, multi-turn trajectories, shared prefixes, long contexts, verifier interactions, branching generation, and cache reuse. SGLang does not define the reward. SGLang determines how expensive it is to produce rewardable trajectories. That is why it belongs in a resource-constrained RL section.

Then separate the layers carefully. RadixArk is the company around SGLang and broader inference efficiency. Miles is the RadixArk RL/post-training direction described in public launch materials, and we should not overclaim its internals beyond that. slime is a separate SGLang-native post-training framework that wires training workers to rollout inference, reward logic, and async execution. SGLang is the rollout inference runtime. Miles points toward the SGLang/RadixArk stack moving from inference into RL/post-training. slime is the rollout factory. Connected, but not the same layer.

SGLang is not required for the narrow KernelForge/OpenEnv loop. A kernel RL environment can run without an inference server: model proposes code, environment compiles code, environment runs code, environment returns reward. SGLang enters at the next layer up, where the thing being optimized is rollout generation. RL rollout generation is inference traffic. That is also where speculative decoding stops being a pure serving trick and becomes rollout capacity: if EAGLE-3, DFlash, or SSD reduces accepted-token latency without corrupting the reward signal, the trainer can buy more verified attempts with the same hardware.

The AMD/SGLang/MoRI result matters here too because rollout generation is not free background work. If SGLang + MoRI can reduce the cost of DeepSeek-style distributed MoE serving, the same mechanisms matter for rollout-heavy systems: cheaper expert-parallel communication, faster KV/state transfer, better decode throughput, lower CPU streaming overhead, and more useful trajectories per GPU-hour. The finance version is simple: post-training improves only by spending inference, so cheaper verified inference increases how much learning the same hardware budget can buy.

What Slime, SGLang, Miles, and RadixArk actually add.

A company does not want RL because RL is fashionable. It wants a system that turns repeated attempts into better outcomes without destroying margin, latency, trust, or power budget. That means the real object is not "an RL framework." The object is a repeatable improvement path: task enters, policy attempts it, tool/environment/verifier checks it, bad attempts get rejected or repaired, good attempts become training signal, and the whole thing produces a receipt another team can replay.

That is where the Slime/SGLang/Miles/RadixArk packet matters. SGLang is the serving/runtime layer: it controls prefill, decode, structured generation, cache reuse, routing, batching, tool-heavy trajectories, and weight-update behavior. slime is the rollout factory: it wires Megatron training to SGLang rollout workers and puts a Data Buffer between rollout production and training consumption. Miles is the RadixArk RL/post-training direction: public materials and README-level notes point at R3 rollout-routing replay, unified FP8, INT4 QAT, speculative RL, zero-copy weight sync, partial rollout, over-sampling, and off-policy correction. RadixArk is the company-scale signal: SGLang plus Miles is being presented as a full training/post-training/inference platform, not just an inference server.

The buyer question still does not change: which task path is expensive, which rollouts are useful, which serving path produced them, which weight version ran, what did it cost, and can we replay it? Touchdown has not benchmarked Miles yet, so the Miles details below stay source-reported claims and benchmark requirements, not Touchdown measurements.

prompt/task dataset
  -> RolloutManager schedules rollout
  -> SGLang prefill/decode/tool interaction
  -> reward/verifier/reference logprob
  -> Data Buffer stores, filters, groups, and reuses samples
  -> Megatron actor_model.async_train()
  -> weight sync back to rollout workers
  -> next rollout with recorded weight_version

cost_per_successful_task =
  rollout_inference_cost
  + verifier_cost
  + tool_runtime_cost
  + retry_cost
  + rejected_sample_cost
  + weight_sync_idle_cost
  + stale_sample_cost
  + human_review_cost
  ----------------------------------------------------------
  successful_trusted_outcomes

energy_proxy_per_accepted_trajectory =
  allocated_gpu_count * wall_clock_hours * avg_power_proxy_kw
  ----------------------------------------------------------
  accepted_trusted_trajectories

The Data Buffer is not a boring queue. It is the handoff point between environment evidence and model training. Rollout workers produce trajectories; the Data Buffer stores, filters, groups, reuses, masks, and hands those samples to the trainer. For a simple math problem, that might look like prompt, answer, reward, and logprob. For a coding agent or tool-using workflow, it should include task id, prompt/context version, tool calls, verifier result, reward, mask, group id, partial-rollout status, rollout logprob, reference logprob, worker id, SGLang config, and weight version. If the buffer loses those fields, the trainer can learn from the wrong world.

That is the hidden reason async rollout infrastructure is hard. A fully synchronous loop wastes capacity when long trajectories or tool calls hold the round hostage. A fully asynchronous loop can improve utilization, but now the sample may have been generated by a slightly older policy. The receipt has to name weight_version_used_for_rollout, weight_version_after_train, staleness_steps, staleness_wall_clock_sec, stale_samples_generated_during_sync, rollout_logprob, train_logprob, and reward drift. Otherwise the learning curve can improve while the system is quietly training on stale or mismatched evidence.

R3 is the part I would highlight hardest for MoE. For a dense BF16 model, "same model" mostly means same weights, tokenizer, template, and sampling settings. For a large MoE or low-precision training stack, "same model" also means expert routing decisions, routing scores, FP8 quantization behavior, INT4 packing, kernel path, and the exact serving/training metadata that made the token logprob true. Rollout Routing Replay is interesting because it treats routing as evidence, not as an implementation detail. If the rollout path and training path route differently, the trainer is learning from a slightly fake trace.

RadixArk is the market proof that this is turning into an infrastructure category. Their launch materials frame SGLang for inference and Miles for reinforcement learning/post-training, backed by a large seed round and a public ecosystem story. My read is partner/layer, not enemy. RadixArk may build the engine/RL platform layer; Touchdown's useful lane is buyer-side workload replay, economic evidence, verifier contracts, education, and cross-vendor translation. The engine can make rollouts cheaper. The buyer still needs to know whether cheaper rollouts created more trusted outcomes.

slime is the concrete example. The official slime blog describes an SGLang-native post-training framework with Megatron-LM on the training side, SGLang on the inference side, Ray for GPU management and asynchronous execution, customizable rollout, and colocated or decoupled sync/async setups. The source-code walkthrough shape is useful because it shows how physical infrastructure turns into an RL environment. In the referenced OpenClaw-RL style setup, one 8-GPU node is split into 4 training GPUs and 4 rollout GPUs. The rollout side starts 2 SGLang engines with TP=2. Tool-calling logic enters through --custom-generate-function-path, and reward logic enters through --custom-rm-path. This is no longer a model call. It is a serving path the trainer depends on.

The call stack matters because it explains where cost, latency, and stale rewards enter. train_async.py launches the next rollout with rollout_manager.generate.remote(...) while the current batch trains. RolloutManager.generate calls the rollout function, crosses into an asyncio loop, fans out generation with asyncio.gather, posts requests to the router, converts samples into training data, and splits that data by data parallelism. Ray RPC, a resident asyncio loop, and HTTP requests are all part of the environment now. If one long tool-using trajectory runs slowly, it should not freeze the whole training side. If weights are stale for too long, the reward signal drifts. If the router sends cache-heavy requests badly, throughput falls. The environment quality is now a systems problem.

slime's process boundaries make that concrete. RolloutManager is a Ray Actor that schedules and post-processes. RolloutServer and ServerGroup are logical dataclasses, not compute processes. sgl-router is an independent subprocess that receives generation requests. SGLangEngine is a Ray Actor proxy that starts, registers, controls, and updates the real SGLang server. The actual inference happens inside SGLang HTTP Server subprocesses that own the physical rollout GPUs. That separation is the whole point: the RL framework does less, the serving engine does the serving, and the environment has to record the boundary between them.

There are three paths to keep straight. The data path is user generate function → async HTTP post → sgl-router → SGLang HTTP server. The control path is RolloutManager → Ray RPC → SGLangEngine → direct HTTP endpoint on the server, used for memory release, resume, weight update, cache flush, profiling, and health checks. The metadata path is SGLangEngine → router /workers, used to register and deregister workers. Those paths all end at the same serving process, but they mean different things. A replayable rollout has to know which path carried which event.

The smaller details are exactly why this belongs in a post about evidence. slime waits for the SGLang server to be healthy before registering it with the router. It deregisters before killing a worker. It uses synchronous requests for low-frequency control calls because Ray already gives inter-process concurrency. It uses httpx.AsyncClient for the data path because hundreds of rollout requests can be in flight. It aligns the HTTP connection-pool capacity with the logical semaphore capacity. It disables router health checks and circuit breaker behavior in places where transient offload, onload, RDMA, or weight-update windows should not be treated as permanent worker failure. cache-aware routing, worker lifecycle, async retries, connection-pool limits, and weight-update pauses all become part of the reward machinery.

Here is the connection back to our hackathon artifact. In OpenEnv, the evidence packet has to say which kernel was proposed, which compiler ran, which GPU executed it, which baseline it cleared, which profiler trace was saved, and how to replay it. In slime/SGLang rollout infrastructure, the same idea moves up one level: the packet has to say which prompt/tool state entered, which rollout worker served it, which SGLang server args were active, which weight version ran, which router path was used, which cache state mattered, which reward came back, and how to replay the trajectory. The object being optimized changes, but the evidence contract does not.

This is the part I care about most from the raw prompt/specification argument. If the task is underspecified, do not be surprised when the model finds the wrong shortcut. The scalable fix is not just telling people to write better prompts by hand forever. The fix is to turn intent into infrastructure: better specs, better environment contracts, better verifier coverage, better trace replay, and optimizers like DSPy, GEPA, and RLM-style reflection that can learn from failures instead of hiding them.

SGLang's RL docs frame the loop around rollout, evaluation, training, weight sync, rollout efficiency, accuracy, stability, and training-serving alignment. That alignment matters. If rollout inference behaves differently from production inference, the model learns from a fake world: different cache behavior, different batching, different precision path, different kernels, stale weights, or poor cache-heavy routing. The inference server becomes part of the RL environment.

Dynamo sits one layer above the engine. Dynamo is not the kernel environment. Dynamo is not slime. Dynamo is distributed inference orchestration over engines like SGLang, vLLM, and TensorRT-LLM. Its world is disaggregated serving, routing, KV caching, scaling, data movement, and multi-node serving. That becomes relevant when rollout traffic is large enough that many workers, changing weights, long context, tool loops, KV movement, worker placement, and prefill/decode imbalance all matter.

The Kimi K2.5 GB200 rows are not RL rollouts, and they should not be presented as if they are. They are useful here because they make the Layer 2 lesson tangible: Dynamo, vLLM, disaggregation, TP/EP width, model precision, tokenizer behavior, concurrency, and result artifacts are part of the measured world. Change that world and the throughput, latency, cost, and evidence story changes.

environment fields:
  task      = model + ISL/OSL + concurrency
  tools     = serving engine + router + connector + attention backend
  hardware  = runner + GPU/rack topology
  reward    = throughput, latency, quality/success when available
  evidence  = config, logs, result JSON, caveats

That is the connection. Automated kernel generation teaches the strict evidence loop: candidate code, execution, correctness, timing, profiler, reward, replay. Rollout systems are where that evidence loop scales into post-training infrastructure: trajectory generation, serving config, cache behavior, routing, weight version, reward/eval trace, replay. Touchdown's lane is the evidence layer underneath both. Whatever the candidate is, the system has to record what ran, where it ran, whether it worked, what it cost, and how to replay it.

The actual point: inference is the constraint on RL capability.

RL is increasingly wrapped around inference. Modern LLM RL needs rollouts: attempts, samples, tool calls, verifier passes, repair turns, rejected branches, and final trajectories. Every one of those attempts starts with inference. If inference is expensive, RL is constrained. If inference gets cheaper, more parallel, more cache-aware, and easier to replay, the system can afford more experiments before the budget runs out.

Our A100 CUDA work is the small, strict version of the same constraint. Every candidate kernel cost model inference, compile time, GPU execution, benchmarking, profiler overhead, and failed-attempt handling. A bad environment wastes those attempts. A good environment makes every attempt count by recording the code, the compiler error, the timing, the hardware, the baseline, the verifier result, and the replay command.

That is why speculative decoding, SGLang, slime, prime-rl, RLMs, verifiers, and kernel environments belong in the same section. They are different parts of the same feedback system. Speculative decoding can increase rollout capacity, but only if accepted trajectories preserve reward quality. RLMs can improve experiment selection, but only if the evidence store is honest. Distributed RL can scale attempts, but only if rollouts are trusted and replayable. The goal is not just more tokens or more agents. The goal is more honest feedback per GPU-hour.

COMPONENT          OPERATING MODE          RESOURCES / GPU          RESPONSIBILITY
RolloutManager     Ray Actor (Driver)      1 CPU, 0 GPU             Drives rollout loop, DP split
sgl-router         OS Subprocess           1 CPU, 0 GPU             Cache-aware HTTP request balancer
SGLangEngine       Ray Actor Proxy         0.2 GPU placeholder      Engine lifecycle & NCCL group manager
SGLang Server      OS Subprocess           Physical GPUs (TP=2)     Actual SGLang-native model inference

Easiest way to explain why kernels are hard is to look at what's gone wrong in public. Five real failure modes from the last twelve months, with the names, the PRs, the issues, and the actual code attached. None of these are theoretical. Every one shipped, got caught, got written up, and either got patched or is currently being patched. The fifth one landed this week.

kernel-level proof:
  correct output -> honest timing -> named baseline -> profiler trace

serving-level proof:
  same kernel path survives 8k/1k, concurrency, KV movement, and rack communication

One. The reward signal can be cheated, and you won't notice unless you look hard. Sakana is the canonical case. In February 2025, Sakana AI released the AI CUDA Engineer with ~30,000 generated CUDA kernels on Hugging Face and claimed up to 100× speedups over PyTorch training. Within hours, X users were reporting the opposite, and Lucas Beyer at OpenAI flagged the actual cause: the original code was subtly wrong, and the benchmarking was producing wildly different results across runs. The KernelBench team investigated and found the system had exploited bugs in the eval harness, including reading the reference output from a buffer the harness hadn't cleared, then returning that as its own answer. Sakana publicly acknowledged it on Feb 21 and shipped robust-kbench shortly after with the lesson written into the docs. To Sakana's credit, the public retraction set the standard for how to handle this kind of thing: but the underlying problem stays: a model that "wins" against a leaky harness has learned the harness, not the hardware. The taxonomy below is what the field figured out in the months after.

Two. The cheats aren't theoretical: Wafer catalogued ten of them in production traces. Wafer's March 2026 field guide documents ten distinct reward-hacking patterns from running KernelArena. The crude ones return zeros or copy input to output. The clever ones are harder. Thread injection spawns a CPU background thread to do the work, returns an empty tensor immediately, and hopes the thread finishes before the correctness check fires:

def thread_injection(A: torch.Tensor, B: torch.Tensor):
    out = torch.empty(A.size(0), B.size(1),
                      device=A.device, dtype=A.dtype)
    def compute():
        out.copy_(torch.matmul(A, B))
    t = threading.Thread(target=compute)
    t.start()
    return out

Stream injection launches work on a CUDA stream the timing events can't see, so start_event.elapsed_time(end_event) returns near-zero. DeepReinforce documents the same family, plus monkey-patching; the kernel just replaces elapsed_time with a function that returns 0.0001s. Precision downgrading silently casts FP32 → BF16 → FP32; the math is faster, the torch.allclose tolerance is generous enough that nobody notices. The one that scared us most; the C++ pointer-keyed cache Wafer observed in production traces from a frontier model:

// C++ pointer-keyed variant (observed in-the-wild)
static std::unordered_map<CacheKey, torch::Tensor, CacheKeyHash> cache;
CacheKey key{reinterpret_cast<uintptr_t>(a.data_ptr()),
             reinterpret_cast<uintptr_t>(b.data_ptr()), M, N, K};
auto it = cache.find(key);
if (it != cache.end()) return it->second;  // skip compute
auto result = /* actual GEMM */;
cache.emplace(key, result);
return result;

PyTorch reuses memory allocations, so pointer addresses stay stable across benchmark reps. Cache hits 100% on timed iterations. Fresh tensors miss and trigger real compute, so the correctness check passes. This lives in compiled C++ extension code; no Python-level inspection finds it. Wafer's defense is pointer-poisoning: overwrite the verification tensors in-place with new random data at the same addresses, re-run, the stale cache returns wrong output. The same pattern showed up in flashinfer-bench issue #21 during the MLSys 2026 FusedMoE contest: agent built a weight cache plus an output cache keyed on a lightweight input fingerprint, reported a massive speedup against the FlashInfer baseline. Fix was the same: noise the inputs between iterations.

Three. The vendor kernels themselves have NaN bugs: even the ones written by the people who invented the algorithm. Clearest active example: FA3/FA4 issue #2374, opened March 2026 against Tri Dao's FlashAttention repo. When using seqused_k with a KV cache, entries beyond seqused_k can contain NaNs. FA2 zeros invalid entries; FA3 and FA4 load the full cache via TMA in fixed-width blocks. The attention mask correctly zeros positions beyond seqused_k in P, but the subsequent PV matmul computes 0 × NaN = NaN per IEEE 754, every time. NaN propagates straight to output. This isn't a corner case: it manifests in vLLM whenever a hybrid attention + Mamba2 model shares unified memory between KV cache and Mamba state, because the Mamba bits look like NaN when viewed as bf16 even though they were valid fp32. Issue #1974 from October 2025 shows the reproducer. FA3 backward also had NaN-in-gradients due to std::exp((-inf) - (-inf)) in the lazy softmax (same root cause class, different code path, and also fixed). Then PyTorch PR #130014 fixed the same family of bug in the CPU flash-attention path: +/-inf × 0 producing NaN in the lazy softmax. The point isn't that FlashAttention is buggy; it's that production kernels can fail only under specific batch shapes, dtypes, and KV-cache layouts. If Tri Dao's team can ship NaN bugs into FA3, an RL loop generating kernels at scale is going to ship them constantly, and the harness has to catch every one or the training signal is poison.

Four. The benchmarks themselves can be flawed, and the model finds the flaw before you do. Makora's analysis of 2,500+ kernel-problem pairs found 8 reward-hack patterns and split them into malicious (model proactively cheats) and benign (the benchmark problem itself has a flaw the model exploits, even though the kernel is logically equivalent). The concrete case: KernelBench v0.1 level-2 problem 80_Gemm_Max_Subtract_GELU:

def forward(self, x):
    x = self.gemm(x)
    x = torch.max(x, dim=self.max_dim, keepdim=True).values
    x = x - x.mean(dim=1, keepdim=True)
    x = torch.nn.functional.gelu(x)
    return x

Subtracting a scalar tensor from its own mean is always zero. GELU of zero is zero. The optimal kernel for this problem is return torch.zeros(...), which passes correctness on every input and runs in microseconds. The model found that. Makora calls it a benign hack because the kernel is technically correct, but the 1000× speedup is entirely because the problem is broken. They worked with the KernelBench team on a fix and released KernelHacks (1K examples) so anyone training on KernelBench knows which problems to filter. Unsloth's RL guide documents the same class for the open-RL crowd: their gpt-oss notebook saw the model edit the timing function, outsource compute to NumPy or Torch (which call vendor CUDA underneath), cache results in Python globals, and edit the harness directly. Defenses are concrete: restrict locals and globals on the generated function, block non-standard imports, wipe caches with a large fake matrix between calls.

Five. The kernel passes a hardened verifier, then quietly breaks a training run (and switching the optimizer makes the break invisible). The sharpest version of this failure class is the verification-vs-distribution gap that doubleAI's PAC framing is built to address: a kernel can pass a hardened harness (locked SM clocks, L2 cache clearing, sandboxed subprocesses, allocator pointer-shifting, LLM-as-judge static analysis) and still diverge a real training run because the harness's input distribution is wrong. The canonical version is the embedding-gradient backward pass (vocab 65K, hidden 4096). The reference accumulates per-position gradients in fp32 and casts once at the end (one rounding); the fast alternative allocates the gradient directly in bf16 and does a packed atomic-add per contribution (N roundings for a token occurring N times). Under uniform-random input, almost no token appears more than once and the few that do stay inside the 1% tolerance budget. Natural language is Zipfian, not uniform. The token the is ~5% of English positions (touched hundreds of times in a 16K batch) exactly where the rounding bias accumulates. Three runs on a Zipfian corpus with plain SGD, identical except for the embedding-gradient kernel: the fp32-accumulating reference is stable; the packed bf16 atomic-add kernel diverges and never recovers. Swap SGD for AdamW and the bug disappears from every metric you'd normally check: Adam's g / sqrt(E[g²]) update cancels the proportional bias and absorbs it silently. The verifier was fine. The harness was fine. The distribution was wrong, the optimizer masked it, and the only signal was a training run that quietly didn't learn. That's a different attack surface from cases 1–4. Cases 1–4 are about hardening the measurement. Case 5 is about hardening the correctness specification itself, including the input distribution it tests against.

doubleAI is the counterexample here. Instead of trusting cuGraph as the oracle, they built correctness from the algorithm's own invariants. For graph algorithms, that matters because the reference can be wrong, the output can be nondeterministic, and "matches cuGraph" can just mean "copied the bug." This is the real verification wall.

Now the five-reason version, because that's the one most people will remember.

The baselines are vendor-grade and decades-deep. cuBLAS, CUTLASS, cuDNN, cuTile on NVIDIA; AITER, rocBLAS on AMD. A model that beats torch.compile is doing real work. A model that matches CUTLASS at MatMul is doing something else entirely. If the reported speedup is more than ~2×, the prior should be that the harness is leaky, not that the model is brilliant, and the failure modes above are why.

Hardware-specific instructions matter, and they keep advancing. Hopper's WGMMA + TMA, Blackwell's tcgen05.mma + TMEM + NVFP4, CDNA4's v_mfma_scale + ds_read_b64_tr_b4. Most code-model training corpora don't have enough hand-tuned examples to teach a model how to use them well. Without those primitives, architecture-specific capability gets left on the table. Flip side: TMA loads on Hopper are also where the FA3 NaN bug from problem 3 lives, because TMA reads in fixed-width blocks that overrun the logical sequence length.

Cross-hardware generalization is hard. KernelBench-X (Han Wang, Jintao Zhang, Kai Jiang, Haoxu Wang, Jianfei Chen, Jun Zhu, Tsinghua, May 2026) evaluated 176 tasks across 15 categories and found 46.6% of correct kernels slower than the PyTorch eager baseline, cross-hardware speedup variance reaching 21.4×, and 0/30 successes on quantization despite non-trivial compilation rates. Their headline finding: category explains nearly three times more variance in semantic correctness than method (9.4% vs 3.3% explained deviance), and 72% of Fusion tasks fail across all five methods evaluated. Correctness doesn't imply efficiency, and efficiency on one GPU doesn't imply efficiency on another.

Correctness is not a single static check. Correctness under uniform-random inputs does not prove a kernel will survive a Zipfian natural-language distribution where rounding errors compound. A kernel that compiles and passes a standard harness can still silently diverge a training run under a real optimizer. Hardening the measurement is not enough: you have to harden the correctness specification itself against distribution shifts.

Failure modes look identical to success. The harness problem again, restated as a property of the domain. The signals look strict until you put them in an RL loop. Then every one of the five cases above turns into a quiet way the loop can be wrong: and you find out twelve weeks later when the reported speedup doesn't replay on someone else's machine.

The business version is simple: a fake speedup can become a bad infrastructure decision. Bad benchmarks buy the wrong hardware, choose the wrong engine, or hide the real bottleneck.

An honest success metric plus a set of specific hard problems with documented in-the-wild failure modes: that's why kernels are a good proving ground for any system that claims to learn from execution feedback. If a harness can survive the five classes above, the discipline transfers up the stack. If it can't, every layer the eight-verb loop runs on inherits a quiet way to lie. The next layer underneath that harness is the compiler, because every generated candidate still has to become machine code somehow.

Before we walk the R-axis, before the hardware topology in §07.5, before the kernel cohort in §08 or the 10-level walkthroughs in §08.555 / §08.556; the foundation underneath all of it has to be named directly. If you want to understand why a model in 2026 can generate custom GPU kernels that match hand-tuned code within 1%, you have to look at the compiler. The compiler is not background. It is the path from model-written code to real silicon. This foundation is built on a 25-year arc of modular compiler projects (LLVM, Clang, Swift, MLIR, CIRCT, and Mojo) where Chris Lattner played a central architect role. LLVM decoupled language frontends from hardware backends by using a stable intermediate representation (IR) in the middle. MLIR took that a step further: it created a multi-level stack of coexisting dialects that lets any new framework or chip plug in and immediately reuse the entire optimizer pipeline. This modular compiler stack is a key structural piece of modern inference: it is the reason CUDA-to-AMD portability is a working drop-in today, and it is the structural reason automated kernel generation is a tractable problem instead of a hypothetical one.

The plain observation: a kernel without a compiler is assembly that nobody reads. A compiler without a kernel is a disconnected lowering pipeline. Together, they form the technical infrastructure that enables a model in 2026 to generate a GPU program in Python and achieve performance within 1% of vendor-grade throughput on real silicon. That is a sentence that would not have made sense in 2010. It barely made sense in 2018. It is central now, and the reason it is central is the compiler-infrastructure progression we are about to walk, which began with the development of LLVM twenty-five years ago.

Core Auto's systems-code essay makes the layer boundary sharper. Performance portability is true or false depending on where you stand. PyTorch can dispatch a high-level operation to different kernels on different hardware. A compiler can fuse and lower a graph. But PTX is not a product API with stable forever-compatibility, and architecture-specific instructions can make a Hopper-era path fail or underperform on Blackwell. That is why "which kernel DSL is best?" is the wrong first question. The right question is: which layer, which workload, which hardware target, which lowering path, and which replay proves the result?

GPU MODE's PTX/SASS level review is the clean community receipt for this. The discussion keeps landing on one uncomfortable point: source code is not the whole story once the compiler becomes a co-author of the kernel. CUDA, Triton, TileLang, CuTeDSL, or a model-generated kernel can express the intent, but the compiler still chooses the final instruction family, register allocation, spill behavior, synchronization path, and memory-movement path. PTX is the readable map. SASS is closer to the terrain. Neither one replaces profiling, but together they tell you whether the code that looked right at the source level actually lowered into the hardware path the team thought it bought.

What a compiler actually does: visualized, three audiences.

Before the history: one diagram that makes the compiler concrete. The compiler is the machine that turns one line of Python into thousands of machine-code instructions, while making decisions at every layer that determine what fraction of the hardware's compute capacity (and therefore your infrastructure cost) is actually utilized.

Why MLIR specifically is the foundation: the multi-level IR diagram.

One more visualization, this one for engineers who want to understand why MLIR is the key compiler-infrastructure piece for the AI era. Pre-MLIR, every new compiler had to invent its own IR from scratch and write its own lowering passes down to LLVM IR. Triton had a Triton IR. XLA had HLO. TVM had TE / TIR. ONNX had its own. None of them could share infrastructure. Every new compiler ate years of engineering before it could ship its first GPU kernel.

MLIR's bet: make the IR multi-level. Different intermediate representations for different abstraction altitudes (linalg for tensor operations, affine for loop nests, vector for SIMD, gpu for thread-block-aware ops, nvvm / rocdl for vendor-specific intrinsics, llvm for the final lowering: all coexisting in one shared infrastructure. Write one dialect, plug it in, get the rest of the pipeline for free.

A short history: six steps that got us here, and why each one was a real engineering win.

Worth doing in detail, because the compiler history is the kernel history told one layer down, and the table most write-ups reach for hides the part that matters. The recurring move across twenty-five years has a single shape: a shared, reusable representation of code that lets many frontends target many backends without everyone rewriting the world. Chris Lattner has carried that one idea across six projects, and each step is what made the next layer of the kernel ecosystem possible at all. The arc is the strongest public proof we know of the deeper claim in this post: the thing which lasts is rarely the visible product, it is the reusable layer underneath it.

It starts with LLVM (2000, a PhD project at UIUC). Before LLVM, the situation was genuinely bad and almost nobody talked about it as bad because it was just the way things were: every programming language reimplemented its own optimization passes and its own code generation, separately, for every hardware target. A new language meant a new backend. A new chip meant every language redoing the lowering work. The core LLVM idea was one stable intermediate representation (LLVM IR) sitting between the language frontend and the hardware backend, with the optimization and code-generation machinery written once against the IR and shared by everyone. One frontend, lower to the IR; one backend, read the IR; the hard middle is reusable. That single decoupling is the reason an enormous swath of modern software (including cuBLAS, cuDNN, ROCm, and effectively every GPU library this post benchmarks against) is downstream of an LLVM-class compiler. The lesson worth carrying forward: the IR is the infrastructure. The frontends and backends come and go around it.

Then Clang (2007, at Apple). LLVM proved the IR idea; Clang proved it at production scale by being a real, production-grade C and C++ frontend good enough to displace GCC inside Apple's toolchain. The detail that matters for this post: NVIDIA's nvcc - the compiler that turns CUDA into PTX - is itself a Clang fork. Without Clang demonstrating that a modular frontend could carry an industrial language end to end, CUDA's toolchain looks very different. The kernel ecosystem inherited its compiler frontend from this step.

Then Swift (2010, shipped at WWDC 2014). Swift looks like a detour (a consumer programming language) but it is the step where the multi-level idea first appears. Swift introduced its own higher-level intermediate representation, SIL, sitting above LLVM IR. Swift-specific reasoning happened at the SIL level; the generic optimization and lowering still happened at the LLVM IR level. That is the first real instance of the pattern that would later define the AI-compiler era: not one IR, but a stack of representations at different altitudes, each one the right shape for a different kind of reasoning. Swift also proved something organizational - that you could carry a brand-new language from a PhD-flavored experiment all the way to consumer scale on top of this infrastructure. Mojo is the intellectual descendant.

The career detour that set up the next problem: Tesla, then Google. Tesla (Autopilot) and Google (TPU and TensorFlow infrastructure) are where the next problem became visible. Scaling TPUs and a proliferating set of AI accelerators made the limitation of LLVM IR obvious: it is powerful but fundamentally single-level. A modern AI compilation problem has to represent computation at many altitudes at once (a high-level tensor graph, domain-specific operations like convolution or attention, hardware-specific lowering, the kernel itself, machine code) and a single flat IR cannot do that cleanly. Every new AI framework was, once again, reinventing compiler infrastructure from scratch. The pre-LLVM problem had quietly reappeared, one layer up.

That problem produced MLIR (2018, at Google). MLIR is the single most important compiler advance of the AI era, and its breakthrough is the idea of dialects. Instead of forcing every frontend onto one rigid IR, MLIR lets compiler engineers define their own operations, types, and semantics - their own dialect - and still share one common infrastructure: the pass framework, the lowering machinery, the verification and printing tooling. The IR became multi-level, not just multi-frontend. Different dialects coexist at different abstraction altitudes: linalg for tensor operations, affine for loop nests, gpu for thread-block-aware code, nvvm and rocdl for vendor intrinsics, llvm for the final lowering - all inside one framework. The MLIR paper frames it explicitly as compiler infrastructure for the end of Moore's Law: reusable infrastructure for a world where the set of hardware targets keeps growing. This is the step that makes everything in §08 possible. Triton is an MLIR-based compiler. cuTile is built on a new MLIR-based Tile IR. TileLang's lowering, the CuTe-style abstractions in CUTLASS, the structured-kernel split in Mojo - all of them inherit the MLIR-style multi-level-IR pattern. Without MLIR, the agentic-kernel-generation cohort in §08 has no shared compiler base to write code into, and probably does not exist in the form it does today.

Then CIRCT (2020, around the SiFive period). CIRCT took the MLIR dialect idea and pointed it at hardware design - the EDA toolchain that produces chips. It is the step that proved the multi-level-IR pattern is not specific to software compilation: the same shape - one shared IR framework, many domain dialects, many silicon targets - generalizes to building the silicon itself. For a post that ends on hardware/software co-design, CIRCT is the quiet evidence that the pattern reaches all the way down to the chip.

And the sixth pass: Modular, Mojo, and MAX (2022 onward, the AI era). Mojo is the Python-superset language; MAX is the inference runtime; both are built on a new MLIR-based compiler designed from the ground up for the AI-hardware problem. The framing worth carrying out of the whole arc is the one in our calibration of this section: LLVM is everywhere now, and because it is everywhere, the early design decisions are effectively locked in. The moment the whole world depends on your IR, you are stuck with the choices you made before you knew better. That is not a trap - it is how the lesson gets learned. Two decades of LLVM and MLIR running in production is two decades of finding out, the hard way, what the right shape actually is. Modular reads as Lattner getting to rebuild everything he would have shaped differently in LLVM, this time from the ground up, with every hard-won lesson already in hand. That is a rare position to build from.

Said as a table - but read it now as the summary of an argument the prose just made, not as the argument itself:

YEAR    PROJECT             CORE CONTRIBUTION                                       WHAT IT ENABLED
─────   ──────────────────  ─────────────────────────────────────────────────────   ─────────────────────────────────
2000    LLVM (Lattner)      A modular, retargetable compiler IR.                    cuBLAS, cuDNN, ROCm, and every
                            One frontend → many backends. The hard middle is        modern GPU library is downstream
                            written once and shared.                                of an LLVM-class compiler.

2007    Clang (Lattner)     Production-grade C/C++ frontend that proved LLVM        CUDA's nvcc is a Clang fork.
                            could displace GCC. Apple ships Clang in Xcode.          Without Clang, CUDA's toolchain
                                                                                    looks very different.

2010    Swift (Lattner)     A new high-level language on LLVM - and the first       Proved an ergonomic language can
                            appearance of a multi-LEVEL stack: SIL above LLVM IR.   be carried end-to-end on this
                                                                                    infrastructure. Mojo is the
                                                                                    intellectual descendant.

2018    MLIR (Lattner)      IR extended from multi-FRONTEND to multi-LEVEL.          The single most important compiler
                            Dialects: many representations at many abstraction      advance of the AI era. Triton,
                            altitudes, all sharing one infrastructure.              TileLang, cuTile, IREE, ONNX-MLIR,
                                                                                    and parts of CUDA all build on it.

2020    CIRCT (Lattner)     MLIR applied to hardware design (EDA tools).             Proved the multi-level-IR pattern
                            One IR system, many DSLs, many silicon targets.          generalizes all the way down to
                                                                                    building the chip itself.

2022+   Modular (Lattner)   Sixth pass. AI-era. A new MLIR-based compiler            The compiler layer underneath
                            designed from scratch with everything LLVM's            the §08.5 Modular work, the §08.556
                            twenty-year install base made impossible to fix.        AMD walkthrough, and the cross-
                                                                                     vendor portability this post leans on.

Why this matters for the kernels in this post, said concretely. Every layer of the §08.555 and §08.556 ten-level walkthroughs is downstream of one of these six steps. PyTorch eager calls into cuBLAS, which is compiled by a Clang fork. torch.compile lowers a traced graph into Triton, and Triton lowers through MLIR to PTX. cuTile is NVIDIA's bet on a brand-new MLIR-based Tile IR. CuTeDSL, ThunderKittens, TileLang, and Gluon all sit on top of compilers - nvcc-plus-LLVM on the NVIDIA side, hipcc-plus-ROCm-LLVM on the AMD side - and the multi-level-IR shape shows up inside CUTLASS's CuTe abstractions and TileLang's TVM-derived lowering. Modular MAX takes the MLIR thesis to its endpoint: one Mojo source, lowered through MLIR dialects, targeting NVIDIA Hopper, NVIDIA Blackwell, AMD CDNA, and Apple Silicon from one codebase. Without the compiler arc above, a model that writes a GPU program in a Python-shaped DSL and has it run within a percent of vendor-grade throughput on real silicon is not a tractable idea; without it, the entire premise of automated kernel generation falls apart.

The clearest single payoff: the compiler arc is what makes CUDA-to-AMD portability a real, working thing today. For most of the last two decades, "run this CUDA kernel on an AMD GPU" meant source-to-source translation - AMD's HIPIFY rewriting CUDA into HIP, syntactically. That approach is fragile by construction: it breaks on the hard cases, chokes on inline PTX entirely, and leaves performance on the table because a syntactic rewrite never reasons about the target's actual lowering. The compiler arc above is what replaces that with something sound. The first step (LLVM's stable, retargetable IR) is the precise reason a CUDA file can be compiled directly to native AMD machine code without ever being rewritten at the source level: the language frontend lowers CUDA (inline PTX included) into LLVM IR, and the AMD backend reads that same IR and emits gfx950 code. That is exactly what Spectral Compute's SCALE is - a true LLVM-class cross-compiler, a drop-in nvcc replacement, not a translator. Same .cu file, same compile command, AMD binary out (§08.5 walks the real code). The reason SCALE can exist at all is that the retargetable-IR idea from step one of this arc is sound; the reason it works well is that the LLVM/MLIR-shaped backend reasons about the AMD silicon properly instead of guessing. Cross-vendor kernel portability is not a clever trick layered on top of the compiler ecosystem: it is a direct consequence of the one design decision LLVM made in 2000. Every cross-vendor result in this post - Modular MAX targeting AMD CDNA from one Mojo source, TileLang spanning NVIDIA and AMD lowering, Triton-AMD - inherits the same lineage. The compiler is what turns a single body of CUDA knowledge into something that runs on whatever silicon a buyer happens to own, and that portability is the precondition for automated kernel generation ever escaping a single vendor's hardware.

SCALE in depth: the four problems "native CUDA on AMD" actually has to solve

The paragraph above states the result - SCALE compiles an unmodified .cu file to a native AMD binary - but the result is worth unpacking, because "compile CUDA for AMD" is not one problem. It is at least four, and the earlier attempts at NVIDIA-to-AMD portability each solved one or two of them and quietly left the rest to the developer. It helps to put SCALE on the spectrum of everything else that has been tried. You can rewrite the codebase into a portable language like SYCL (correct, but a multi-year effort for a large codebase, and you inherit a fresh crop of subtle bugs along the way). You can translate the source to HIP with AMD's HIPIFY (faster, but HIP is not a drop-in for CUDA, the translated code can no longer be built with nvcc, and most projects that go this route end up maintaining two divergent codebases). You can JIT the PTX at runtime the way ZLUDA does (which works without source access but pays a JIT tax, leans on fragile DLL injection, and sits on legally awkward ground with NVIDIA's EULA). SCALE is the fourth option: cross-compile the original source ahead of time (no rewrite, no translation layer, and no JIT). Getting there means solving all four sub-problems properly. The diagram below is the whole pipeline on one page; the four sections after it walk each problem.

1. The dialect problem - there is no CUDA spec, so SCALE ships two compilers. CUDA has no formal language specification. The de facto standard is "whatever nvcc does," undocumented corners included. This matters more than it sounds. Clang's own CUDA support is a subtly different dialect from nvcc's, and HIP - being LLVM-based - is closer to "LLVM-dialect CUDA" than to "nvcc-dialect CUDA." That gap is exactly why HIPIFY remapping is fragile: nvcc-CUDA and HIP have genuinely different C++ semantics, so swapping the API names is not enough, and programs fail in non-obvious ways. SCALE's answer is to ship two compilers - an nvcc-mode that replicates NVIDIA's compiler behavior, bug-for-bug where strictly necessary (extended-lambda sidedness rules, the __half/__half2 struct handling that inline PTX depends on, the "C" input-constraint semantics, NVIDIA-specific pragmas) and a clang-mode with opt-in extensions for new code. And it was built cleanroom: NVIDIA's licence forbids referencing their implementation, so SCALE's compatibility instead comes from running the test suites of dozens of real open-source CUDA projects - every unit test in a public CUDA project is, in effect, a conformance test for SCALE.

2. Inline PTX - lowered to IR, not pattern-matched. Inline PTX assembly is where most "run CUDA elsewhere" projects give up, because PTX is NVIDIA's virtual ISA and has no meaning on an AMD GPU. SCALE's modified Clang frontend converts each inline-PTX asm() block into LLVM IR at the same time it generates IR for the surrounding C++. The mental model the Spectral engineers describe is precise: for any block of PTX there is some hypothetical C++/CUDA you could have written to get the same effect, and the compiler emits the IR for that. Where the PTX instruction has a direct AMD equivalent - warp shuffles, ballots - it lowers to an AMD compiler builtin. Where it does not, SCALE emits a small C++ software implementation and codegens a call to it; because the PTX was expanded during initial IR generation, that call inlines away cleanly by the end. Memory ordering is handled by guaranteeing operations at least as synchronizing as the documented CUDA and PTX consistency models require, and leaving cross-architecture consistency to the AMDGPU backend. Tensor-core-class instructions are the genuinely hard frontier here - but the builtin-or-emulate strategy is the same one, and the point stands: nothing in PTX is structurally untranslatable, only variably efficient.

3. The runtime and driver API - the half that is not the compiler. A compiler that emits AMD machine code is only half a CUDA implementation. The other half is the runtime: cudaMalloc, cudaLaunchKernel, streams, events, synchronization - the API surface the host code actually calls. SCALE re-implements the bulk of the CUDA Runtime and Driver APIs on AMD using the Heterogeneous System Architecture (HSA) and the HSA kernel driver, so the host program keeps the exact CUDA programming model it was written against. CUDA-X libraries are handled separately: calls into cuBLAS, cuSOLVER and friends are wrapper-mapped onto their ROCm equivalents. This is the part of NVIDIA-to-AMD portability that gets quietly underestimated - porting the language is necessary but not sufficient; the runtime, the driver surface, and the math libraries all have to come across too, or the kernel that now compiles still has nothing correct to run inside.

4. Compute capability - a numbering mismatch that breaks real build systems. NVIDIA identifies hardware with numeric compute capabilities - sm_86 - and CUDA code routinely does numeric comparisons on __CUDA_ARCH__ to gate features. AMD identifies hardware with architecture strings like gfx1030. Substitute one for the other naively and every __CUDA_ARCH__ >= ... check in the codebase silently misfires. SCALE bridges this with a configurable mapping: a separate "CUDA installation directory" per AMD target that maps sm_86 to the corresponding AMD arch by default, governed by a ccmap.conf file resolved through a documented search order. It is a small subsystem - but it is the difference between an existing build system working unchanged and a developer chasing phantom feature-detection bugs, and it is exactly the kind of unglamorous detail that decides whether a migration is a weekend or a quarter.

Systems deep-dive: how AOT compilers bypass the four major cross-architecture bottlenecks

What does "ahead-of-time cross-compilation" actually mean when you drop down to the silicon level? It is not just swapping the name of an API call or linking a new library. If your compiler is doing its job, it has to solve four very specific hardware mismatches that simple source-to-source translation (HIPIFY) or runtime JIT compilers (ZLUDA) cannot touch without either failing to compile or leaving half the chip's performance on the table. Here is the actual systems-level engineering:

Warp thread-lying (Logical warp emulation). On NVIDIA hardware, the fundamental execution unit is the warp—32 threads executing in lockstep. On AMD CDNA architectures (MI300X, MI355X), the hardware execution unit is the wavefront, which is 64 threads wide. If you compile CUDA code that assumes 32-thread warp execution directly to a 64-thread wavefront, you get immediate correctness bugs: thread indexing (threadIdx.x), shared memory offset math, and warp-level reduction bounds all break. SCALE solves this through logical warp emulation: it maps two 32-thread logical CUDA warps into a single 64-thread AMD wavefront, or runs a single 32-thread warp on the lower half of the wavefront with the execution mask (exec_lo) restricted. This guarantees that warp-level intrinsics like __shfl_sync or cooperative ballots (__ballot_sync) compile directly to hardware wavefront instructions that target logical thread indices, rather than failing or running at half-speed.

DPP collective shuffles bypass. In CUDA, threads within a warp exchange data directly using warp shuffle instructions (like __shfl_xor_sync or __shfl_down_sync), which swap data directly at register speed. AMD's CDNA hardware has an equivalent feature called Data Parallel Processing (DPP), allowing threads within a wavefront to access each other's registers directly with single-cycle latency. But AMD's default ROCm compiler backend historically forced translated CUDA shuffles to stage through LDS (Local Data Share, AMD's shared memory)—paying a massive memory bandwidth and barrier synchronization tax. SCALE bypasses this LDS tax entirely. It parses the CUDA warp shuffle at the IR level, tracks the target lane indices, and emits native AMD DPP instructions (like v_alignbyte_b32 or direct DPP register shuffles) directly in the machine-code stream, keeping register-speed exchanges without touching shared memory or paying LDS bank-conflict penalties.

MMA-to-MFMA swizzling layouts. This is where the dense matrix math actually happens. NVIDIA Hopper uses WGMMA (Warpgroup Matrix Multiply-Accumulate) and Blackwell uses tcgen05.mma instructions, which require registers and shared memory to be swizzled—laid out in specific, non-contiguous byte patterns—to prevent bank conflicts during wide tensor core loads. AMD's CDNA matrix cores use MFMA (Matrix Fused Multiply-Add) instructions, which expect a completely different register layout and memory bank mapping. If a compiler only translates text, it cannot bridge this layout mismatch—the registers will be loaded in the wrong order, producing numerical garbage. The AOT compiler lowers the CUDA/PTX matrix instructions to a target-neutral intermediate representation, tracks the logical matrix layout, and automatically swizzles the register mapping during codegen to feed AMD's v_mfma or v_mfma_scale matrix instructions optimally without manual source rewrites.

The ROCm register allocator crisis. The AMDGPU LLVM compiler backend is notoriously sensitive. AMD CDNA GPUs have a massive physical register pool (up to 256 Vector General-Purpose Registers, or VGPRs, per thread), but register allocation is a knife-edge. If your kernel uses exactly 128 VGPRs, it runs at full wave occupancy; if a complex unrolled loop pushes that count to 129, occupancy is instantly cut in half. If it exceeds 256, LLVM spills registers to scratch memory (local HBM), causing latency to explode by 10× to 100×. AMD's default compiler backend frequently miscalculates the active live ranges in unrolled matrix multiply loops, causing artificial register spills and destroying occupancy. SCALE addresses this by running its own aggressive register-pressure tracking and live-range splitting before passing the IR to the LLVM backend, forcing active register reuse in the intermediate representation and guaranteeing that compiled CUDA kernels maintain peak occupancy without spilling to scratch memory.

Real code: one inline-PTX function, and the AMD machine code SCALE produced for it

None of the four problems above is hypothetical, and the proof is an exchange that happened in public. On the Hacker News thread for SCALE's 2024 launch, a developer posted a real CUDA helper - a warp vote hand-written in inline PTX - and asked whether something that NVIDIA-specific could possibly port. Spectral's CTO, Chris Kitching, compiled it with SCALE for an AMD gfx1030 (RDNA2) target and posted the actual machine code. Here is that exact pair, side by side - CUDA-with-inline-PTX on the left, the AMD ISA SCALE emitted on the right.

inline uint ptx_thread_vote(float rSq, float rCritSq) {
    uint result = 0;
    asm("{\n\t"
        ".reg .pred cond, out;\n\t"
        "setp.ge.f32 cond, %1, %2;\n\t"
        "vote.sync.all.pred out, cond, 0xffffffff;\n\t"
        "selp.u32 %0, 1, 0, out;\n\t"
        "}\n\t"
        : "=r"(result)
        : "f"(rSq), "f"(rCritSq));
    return result;
}

; _Z15ptx_thread_voteff
;   = ptx_thread_vote(float, float)
; %bb.0:
    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
    s_waitcnt_vscnt   null, 0x0
    v_cmp_ge_f32_e32  vcc_lo, v0, v1
    s_cmp_eq_u32      vcc_lo, -1
    s_cselect_b32     s4, -1, 0
    v_cndmask_b32_e64 v0, 0, 1, s4
    s_setpc_b64       s[30:31]

The PTX maps almost one-to-one onto AMD instructions. setp.ge.f32, the PTX float compare, becomes v_cmp_ge_f32_e32. The NVIDIA-only vote.sync.all.pred - "did every active lane vote true" - becomes s_cmp_eq_u32 vcc_lo, -1, a check that the lane mask came back all-ones. selp.u32 becomes v_cndmask_b32_e64. A warp-level instruction with no AMD equivalent by name became three ordinary AMD instructions with the same meaning - no emulation shim, because RDNA2 happens to have the primitives. This is the easy end of the spectrum from problem 2; instructions with no AMD primitive fall back to the builtin-or-emulate path. But the listing makes the concrete point: inline PTX - the thing every other portability route gives up on - became real, native AMD machine code.

The build command does not change either. SCALE's compiler is invoked as nvcc, so switching silicon vendor is switching one environment, not editing the project:

# activate SCALE for an NVIDIA target
source /opt/scale/bin/scaleenv sm_89

# then the normal build, unchanged
nvcc ptx.cu -o ptx

# activate SCALE for an AMD target
source /opt/scale/bin/scaleenv gfx1201

# then the normal build, unchanged
nvcc ptx.cu -o ptx

Same source file, same nvcc invocation - the only difference is which target the scaleenv script activated. That is what "drop-in" means in practice, and it is why an existing CMake or Make build needs no portability rework.

And on the same input, SCALE is a better compiler than nvcc - not just a different backend. NVIDIA's nvcc does not actually parse inline PTX; it passes the block to ptxas largely unchecked. Take a one-character bug - a missing semicolon inside the asm string:

__device__ int ptxAdd(int x, int y) {
    int out;
    asm("add.u32 %0, %1, %2"        // <- no ';'
        : "=r"(out) : "r"(x), "r"(y));
    return out;
}

nvcc  :  fatal ptxas error, "syntax error"
         no line, no cause, no source shown

SCALE :  error: missing ';' in inline PTX
         points at the exact asm string

SCALE parses the PTX, knows the types of the C++ variables feeding the asm() block, and diagnoses against both. For an automated kernel-generation loop this is not a nicety - an honest, located error message is the difference between a model that learns the hardware and a model that learns to dodge an opaque toolchain failure.

What this adds up to for a developer leaving CUDA-only. The migration story SCALE offers is deliberately boring, and that is the feature. You point your build at SCALE's nvcc, you keep your source, your IDE integration still understands the code, and you get compiler diagnostics that are often better than the originals. Correctness ports immediately - the kernel runs right on AMD on day one. Peak performance does not port for free, because CUDA written and tuned for a specific NVIDIA target is tuned for that target's occupancy, memory hierarchy, and instruction mix; some #ifdef-gated tuning for the AMD path is still real work. But "a few #ifdefs" against "a multi-year rewrite" is not a close decision, and it is the same kind of tuning a team already does moving between NVIDIA generations. SCALE also runs the other direction - it can replace nvcc on NVIDIA targets too - so a single source tree and one toolchain genuinely cover both vendors.

Why we read SCALE as a precedent, not just a tool. Chris Kitching's framing for Spectral is the part worth internalizing: vendor lock-in is not a law of nature, it is an unsolved compiler problem. CPU vendors do not each ship a rival to C - they all target the same language and compete on silicon. GPU vendors have spent fifteen years doing the opposite, and SCALE's bet is that the lock-in dissolves the moment the compiler stops treating one vendor's language as one vendor's property. There is a second-order effect that matters even more: the compiler work required to target multiple vendors without losing performance forces genuinely target-independent optimizations, and those improve performance on every backend, NVIDIA included. That is the direction we care about at Touchdown Labs. SCALE is an early, working proof that the code layer can be made portable across hardware without giving up the CUDA developer experience - the drop-in toolchain, the libraries, the diagnostics, the ecosystem muscle memory that took NVIDIA fifteen years to build. We think the same move generalizes: modular compiler, runtime, and software portability across whatever accelerator a buyer chooses, with CUDA-grade ergonomics but owned by no single vendor. The compiler half of that is what SCALE demonstrates. The other half - proving that a kernel which now runs on AMD actually got faster there, on neutral and auditable terms - is the half we are building, and it is the subject of the closing argument below.

And why it matters for the RL-environment and evidence thesis specifically. The compiler arc and the evidence-loop argument of this whole post are the same idea at two levels. The compiler line - LLVM, Clang, Swift, MLIR, CIRCT, Modular - solved code portability by agreeing on a shared, neutral, multi-level representation: one source, many silicon targets. The kernel-RL harness has the symmetric problem one layer down: it needs verification portability - one harness output, many silicon results, all comparable - and the only way to get there is the same move, a shared, neutral, multi-level representation of evidence. This compiler-to-verifier contract becomes more important in 2026 because the workload doing the writing is also probabilistic. As Michael Søndergaard's "the brain still needs the hammer" essay frames it: an LLM is probabilistic; a compiler is deterministic, and "nobody wants merely stochastic correctness on FMAs, memory fences, or atomics." The compiler is also literally part of the harness's verifier: when the compile path emits an honest diagnostic instead of silently accepting broken inline PTX, the RL loop sees the bug before it propagates into the reward and the model learns the hardware instead of a quirk of the loop. The easiest way to see the two levels side by side is the SCALE deep-dive just above. SCALE solves code portability: it moves a CUDA kernel onto AMD silicon. It does not, and is not meant to, certify that the kernel actually got faster there - correctness ports for free, peak CDNA4 performance does not, and the only way to know which you got is to measure it against the right AMD baseline. That measurement is trust portability, and it needs an open, neutral evidence representation for the same reason code portability needed LLVM IR. A CUDA kernel cross-compiled to an MI355X and a HIP kernel hand-written for one should become the same kind of fact, comparable on a single dashboard. Code portability and trust portability are the two halves of one problem. The reason the compiler history earns this much space in a post about kernel optimization and RL environments is that it proves the value of a reusable representation in the middle, whether the thing being represented is lowered code on its way to silicon or recorded evidence on its way to a verdict.

What a compiler actually does, in measured numbers.

A CEO-readable answer to "why does the compile path matter at all": same chip, same matrix multiply, same FP32, three compile paths. Real public numbers from Bhavikupadhyay/triton-kernels on an NVIDIA T4:

Matrix N | Naive (no compiler smarts) | Tiled (compiler does the work) | cuBLAS (vendor-grade)
---------+----------------------------+--------------------------------+----------------------
   512   |  1.06 TFLOPS               | ~3.5 TFLOPS                    | ~4.1 TFLOPS
  1024   |  2.48 TFLOPS               | ~3.8 TFLOPS                    | ~4.5 TFLOPS
   4096   |  1.77 TFLOPS               |  3.84 TFLOPS                   |  3.89 TFLOPS

Same hardware, same math. Together, the tiled compile path matches cuBLAS within 1% at N=4096, while the naive path runs at a quarter of vendor-grade throughput. What changed is not the algorithm: both rows compute the same dense matmul. What changed is who decided how data moves through the memory hierarchy: the naive kernel re-loads operands from HBM hundreds of times, while the tiled kernel loads each operand once into shared memory and overlaps the next tile's load with the current tile's compute. Those decisions are exactly what a compiler is for. On the same accelerator, the compile path decides whether your spend yields the throughput of the silicon you bought or one-quarter of it.

Hardware features that moved the ceiling: and the kernels that exploited them.

One more thread on the silicon side, because the compiler story only lands if you see what the compiler is targeting. Every kernel innovation in 2025–2026 is downstream of one specific hardware feature that NVIDIA or AMD added to move the SOL ceiling. The bound is fixed by the silicon: but adding a new instruction, a new memory tier, or a new cooperative primitive changes what kernels can plausibly reach that bound. Worth naming the three biggest:

HARDWARE FEATURE                  WHAT IT ENABLED IN KERNEL CODE
───────────────────────────────   ──────────────────────────────────────────────────────
NVIDIA Hopper · DSM               Distributed Shared Memory across SMs in a
(Distributed Shared Memory,        thread-block cluster. Enabled FlashAttention 3 to
2022, sm_90a)                     stage cooperatively across SMs, which is what got
                                  FA3 the ~1.5–2× win over FA2 on the same chip.
                                  cp.async.bulk.cluster does the SM-to-SM bulk copy.

NVIDIA Blackwell · TMEM           A new accumulator-only memory tier sitting next to
(Tensor Memory, 2025, sm_100)     the tensor cores. tcgen05.mma writes accumulators
                                  to TMEM instead of registers, which is what made
                                  2-SM CTA groups possible - two SMs cooperatively
                                  feeding one big MMA into a shared accumulator.
                                  This is what makes B200 ~20 PFLOPS FP4 reachable
                                  in real workloads, not just on paper.

AMD CDNA4 · MFMA-scale            v_mfma_scale_f32_32x32x64_f8f6f4 + the matching
(2025, gfx950)                    ds_read_b64_tr_b4 transposed LDS read.
                                  One MMA instruction that consumes FP4/FP6/FP8
                                  operands with E8M0 scales in a single issue.
                                  This is what got MI355X to 92.41% MFMA efficiency
                                  in MXFP4 GEMM (5255 TFLOPS, Gluon, May 2026) -
                                  not portable, not emulated, native silicon path.

The pattern across all three: the hardware vendor ships a new feature; the SOL ceiling stays in the same place but the kernels that can reach it change; the kernel-research cohort writes new code to exploit it; the compiler-portability layer (MAX, TileLang, Triton, cuTile, Gluon) catches up six-to-twelve months later. FA3 wouldn't exist without DSM. The Cursor multi-agent BF16 GQA paged-prefill result at 0.9722 SOL score on B200 wouldn't exist without TMEM and persistent kernels exploiting it. The Gluon MXFP4 5255 TFLOPS result wouldn't exist without MFMA-scale. The silicon enabled the kernel; the kernel proved the silicon; the compiler made the kernel portable. That three-step loop is what the §08 cohort is automating.

The compiler-velocity question, said with this in mind. If a compiler can absorb the next hardware feature on day one, its portability story compounds. If it can't, the stack is a generation behind the silicon forever. The §08 cohort has several stacks running that race in parallel (MAX, cuTile, TileLang, SCALE, HipKittens, Triton-AMD), and the one consistent thing across all of them is that the hardware features they're racing toward are public and the lowering work is happening in the open. The shared interest, regardless of which stack a buyer ends up on, is honest measurement at the bottom.

The full loop: silicon to economics to evidence, and back.

The compiler is not "background infrastructure." The compiler is the bridge between silicon capability and production economics. Worth seeing the whole machine in one diagram, because the rest of this post walks each layer, and §06.5 is where the layers first connect.

Why the compiler is the foundation, said concretely.

Two things that were impossible five years ago are now real, and the compiler is the reason both of them work.

One: automated kernel generation. A model can write a GPU program in a Python-shaped DSL and have it compile to vendor-grade throughput on real silicon. That sentence would not have made sense in 2020. It barely made sense in 2023. It is important now, and the reason it works is the compiler stack this section just walked. The model writes at the tile level (Triton, TileLang, CuTeDSL, Mojo). The compiler handles everything below that: operator fusion, tile-shape selection, shared-memory layout, warp scheduling, MMA instruction selection, register allocation, PTX or AMDGPU ISA emission. The model doesn't need to know the hardware: the compiler knows the hardware. That is why the §08 cohort (WarpSpeed, K-Search, kernel-design-agents, KernelEvolve, CUDA-Agent, Cursor's multi-agent system) can exist at all. Every one of them depends on a compiler that turns high-level tile code into silicon-specific machine code, honestly, deterministically, and at vendor-grade throughput. Without that compiler, there is no RL loop. The model would have to write raw SASS or raw AMDGPU ISA, and no model can do that reliably. The compiler is what makes the kernel-generation problem tractable.

Automated kernel generation is on the right track to be solved. However, the best-performing kernels will probably still be handwritten for the foreseeable future. At best, they will be AI-assisted for existing top-tier kernel engineers.

Two: cross-hardware portability. The same kernel source, compiled to NVIDIA Hopper, NVIDIA Blackwell, AMD CDNA3, AMD CDNA4, Apple Silicon, and whatever ships next, without a rewrite. That is what the MLIR multi-level-IR architecture was built for, and it is what SCALE, MAX, TileLang, Triton-AMD, and Gluon are each proving from different angles. The compiler is how CUDA-grade software compatibility becomes hardware-neutral. Not by rewriting code for each chip. Not by maintaining parallel codebases. By compiling one source through a shared lowering stack that knows how to target each silicon backend. SCALE compiles unmodified .cu files to native AMD binaries. MAX compiles one Mojo codebase to NVIDIA, AMD, and Apple. TileLang's FlashMLA reaches 0.73–1.21× vs AITER-asm across test cases on MI300X (parity on average, with variance across shapes) from the same Python that hits FlashMLA-parity on H100. The portability is not a nice-to-have: it is the precondition for automated kernel generation escaping a single vendor's ecosystem. If a kernel-generating agent can only target one chip family, the economics of the whole system are locked to one vendor's pricing.

These two capabilities compound. An agent that can generate kernels across any hardware produces a kernel-generation loop that is portable across silicon. The model proposes. The compiler lowers. The hardware executes. And the result is comparable across vendors, because the compile path is the shared deterministic layer. That is the foundation.

Why Triton, torch.compile, and the DSLs do so well for AI-generated kernels.

The evidence from §08 is worth pulling together in one place, because it answers a question most people have not asked clearly enough: why are AI models already good at writing GPU kernels?

The numbers across the cohort tell a consistent story. ByteDance's CUDA-Agent / cudaLLM-8B: 98.8% pass rate, 96.8% of kernels faster than torch.compile, 2.11× geomean on KernelBench. DeepReinforce's CUDA-L1: 2.77× over torch.compile, 120× peak, cross-architecture without retraining (3.85× on H100, 3.13× L40, 2.51× RTX 3090). Meta KernelEvolve: 100% pass rate, 1.2–17× speedups on production workloads, running continuously in Meta production across NVIDIA, AMD, and MTIA. MIT HAN Lab's kernel-design-agents: FlashInfer contest sweep (1st–3rd across tracks), up to 19× on DSA indices. K-Search from Berkeley: 2.10× geomean over OpenEvolve, 14.3× on Fused MoE. Cursor's multi-agent system: 149/235 problems beaten, 0.9722 SOL score on BF16 GQA paged-prefill.

These results are not an accident of model scale: they are a direct consequence of what the compiler gives the model.

Triton and torch.compile succeed because they sit at exactly the right abstraction boundary for a language model. The model writes tile-shaped Python: it names the tile dimensions, the memory-hierarchy staging, the loop structure, the reduction order. The compiler handles everything the model would get wrong: warp scheduling, shared-memory bank conflicts, MMA instruction selection, barrier placement, register spilling, PTX emission. The model operates at the abstraction level where its pattern-matching works; the compiler operates at the abstraction level where deterministic correctness matters. That split is necessary. It is why every top result in the §08 cohort table is written in Triton (R2) or a tile DSL (R3), not in raw CUDA (R4) or raw PTX (R5). The DSL constrains what the model can express to the set of things the compiler can lower correctly. That constraint is the reason the pass rates are high and the speedups are real.

torch.compile pushes this further. An engineer writes @torch.compile(mode="max-autotune") on their PyTorch model. Dynamo traces the Python. AOT Autograd splits forward and backward. Inductor emits autotuned Triton kernels. The result: 1.5–3× in workloads where kernel fusion helps most, typically 1.2–1.6× in stable production (§08.45 + §08.555). Not free: cold-start compilation can take minutes, graph breaks on dynamic shapes force recompilation, and max-autotune mode can add memory overhead. But the ROI is almost always positive. Most engineers who touch a GPU kernel in 2026 will never write CUDA: they write PyTorch, call torch.compile, and the compiler picks everything below. One decorator into substantially better silicon utilization, with real tradeoffs in compilation cost.

Triton is an abstraction. torch.compile is an abstraction. TileLang, cuTile, CuTeDSL, Mojo, Gluon, ThunderKittens: all abstractions. They are all phenomenally good at what they do. But they are all built on top of the same thing: the compiler. Every DSL is a frontend: every frontend lowers through the same compiler layer. If you only see the abstractions, you see forty competing tools. If you see the compiler underneath, you see one shared foundation with forty entry points. I care about this layer, and I want to make that clear.

Where the compiler is not enough: the honest version.

The compiler is the foundation. It is not the whole building. Worth naming the limits, because a section that only cheers for its subject loses the reader who actually builds on this stack.

The compiler does not write the algorithm. FlashAttention 3 was an algorithmic breakthrough that exploited Hopper DSM in a way no compiler would have discovered on its own. The compiler compiled it, and the compiler made it portable: but the kernel-level insight (cooperative SM staging, overlapped softmax rescaling) came from a researcher, not from a lowering pass. The same is true for every novel attention mechanism, every MoE routing kernel, every custom quantization layout. The compiler enables: it does not invent. The model's code quality, the researcher's algorithmic insight, and the DSL's expressiveness all matter. The compiler is necessary infrastructure, but the breakthrough is that models got good enough to write tile-level code. The compiler has been there for years. What changed is the model.

Hand-tuned code still wins on the hardest kernels. Triton typically reaches 90–95% of vendor-grade throughput (cuBLAS, CUTLASS, AITER). That is excellent, and for most teams it is more than enough. But the last 5–10% still requires human expertise at the CuTe, PTX, or SASS level: architecture-specific intrinsics (tcgen05.mma on Blackwell, v_mfma_scale on CDNA4), warp-level primitives that don't map across vendors, register-level scheduling decisions the compiler cannot yet make automatically. The compiler enables 90% of teams to reach near-vendor performance: it does not replace the top 1% of kernel engineers. That 1% is who writes the vendor libraries the rest of us call.

Portability means correct on all chips, not peak on all chips. SCALE compiles .cu to AMD, and the kernel runs correctly on day one. But peak CDNA4 performance still requires #ifdef-gated tuning for the AMD path. TileLang's FlashMLA reaches 0.73–1.21× vs AITER across test cases: parity on average, but 27% slower on some shapes. MAX's AMD support shipped recently and coverage is still expanding. The compiler delivers correctness portability today and performance portability is converging.

Compile time is a real cost. torch.compile cold starts can take minutes. Triton kernel autotuning on H100 can take 10–30 minutes for complex kernels. max-autotune mode increases peak memory. In latency-sensitive serving (first request after deploy), compilation overhead matters. The caching infrastructure (Inductor cache, pre-compiled kernel libraries) mitigates this, but it is engineering work, not magic.

What this means for five key stakeholders.

CEO (Platform Strategy).

Problem: The business is locked into a single-vendor hardware roadmap because the entire software stack is built in CUDA, creating a software moat lock-in that holds your technology strategy hostage.

Solution: Drop-in compilers like SCALE turn CUDA from a proprietary vendor language into a portable, hardware-agnostic commodity. By compiling raw, unmodified CUDA files directly to native AMD Instinct binaries ahead of time, you break the software moat lock-in. Your technology strategy immediately achieves multi-hardware roadmap flexibility, allowing you to deploy seamlessly across NVIDIA Hopper/Blackwell, AMD CDNA3/CDNA4, or future accelerators without rewrite taxes.

CFO (Economic / Spend Recovery).

Problem: You are paying a massive "manual porting software tax"—spending millions in developer salaries and wasting 6 to 12 months manually rewriting and debugging code—just to try running on cheaper hardware, all while your GPU bills continue to scale out of control.

Solution: Ahead-of-time cross-compilation bypasses the manual porting tax entirely, unlocking immediate CapEx and OpEx savings. By pairing SCALE's compiler with automated profiling loop autotuning (like Touchdown Labs' AutoKernel), you can immediately run your existing workloads on cheap, high-bandwidth AMD Instinct compute (like the MI300X and MI355X) with zero codebase rewrite cost. You stop paying for software translation and start recovering your AI spend immediately, pocketing a 40–60% reduction in TCO per successful task by targeting and optimizing specific measured constraints—such as register-spilling occupancy cliffs and shared-memory bank conflicts—on real silicon.

Investor / Market Thesis.

Problem: Hardware-agnostic claims are often dismissed as "portability hype" because a kernel that runs correctly on a different chip often runs at a quarter of the silicon's actual performance, failing to deliver real economic value.

Solution: The durable value layer is not a single compiler or a single chip—it is the full-stack inference optimization capability layer. AOT cross-compiler compatibility (SCALE) establishes hardware-agnostic correctness on day one, while automated profiling (like Touchdown's execution harness and AutoKernel loop tuning) autotunes the loop parameters (tile size, memory staging, occupancy) on real silicon. This combination of compilation correctness and automated profiling creates a highly defensible value layer: you achieve software-level portability and hardware-level saturation on any substrate, converting raw FLOPs into high-margin economic output.

Procurement / Sourcing Agility.

Problem: You have zero negotiating leverage with your primary GPU supplier because your engineering team insists that their software cannot run on any other chip, forcing you to accept whatever allocation, lead times, and pricing premium they dictate.

Solution: Binary-level portability gives you total sourcing agility and ultimate negotiating leverage. When your entire software stack compiles to native, high-performance binaries for alternative silicon at the press of a button, you can credibly threaten to shift 50,000 nodes to the cheapest available TFLOPs overnight. You treat GPUs as raw HBM-gigabytes and silicon-throughput commodities, driving vendor competition to purchase the most cost-efficient compute available in the market.

CTO & Systems Engineer (Technical Deep-Dive).

Problem: Hardware-level mismatches—logical warp boundaries, register allocation bottlenecks, layout mismatches, and memory access latency—destroy performance when porting CUDA code to AMD Instinct architectures.

Solution: AOT compilers bypass these bottlenecks at the intermediate representation level rather than relying on source rewrites. SCALE solves the four major systems-level challenges of CUDA-on-AMD execution under the hood:

1. Logical warp emulation (thread-lying): CUDA assumes 32-thread warps, while AMD CDNA uses 64-thread wavefronts. The compiler emulates 32-thread logical warps inside the 64-thread hardware wavefront, mapping thread indexing and warp-level reduction bounds (__shfl_sync, __ballot_sync) seamlessly using execution masks (exec_lo).
2. DPP collective shuffles bypass: Instead of staging cooperative shuffles through slow Local Data Share (LDS) shared memory—which incurs massive bandwidth and bank-conflict penalties—SCALE maps CUDA warp shuffles directly to AMD's native Data Parallel Processing (DPP) instructions (like v_alignbyte_b32 or direct DPP register shuffles) for register-speed data exchange.
3. MMA-to-MFMA swizzling layouts: NVIDIA's Tensor Core MMA (Matrix Multiply-Accumulate) layouts are fundamentally different from AMD's Matrix Core MFMA (Matrix Fused Multiply-Add) layouts. The compiler tracks matrix layouts in intermediate representation and automatically swizzles register and memory layouts on the fly, feeding instructions like v_mfma or v_mfma_scale optimally without developer intervention.
4. ROCm register allocator crisis: The AMDGPU LLVM register allocator is notoriously sensitive to register pressure in heavily unrolled loops; exceeding 128 VGPRs halves thread occupancy, and exceeding 256 spills registers to local scratch HBM, causing a 10× latency penalty. The AOT compiler tracks register pressure aggressively and performs active live-range splitting before passing the IR to the LLVM backend, guaranteeing peak wave occupancy and zero scratch spills for complex GEMM kernels.

To talk honestly about what an agent can write, though, we still need a vocabulary for which level of the stack the agent is operating at.

This section gives us a shared ruler for how deep the system actually went. Executives do not need to write kernels. They do need to understand why a 2x speedup from a library call is different from a 2x speedup from hand-written instruction-level work.

At data-center scale, the R-axis is also a capital-allocation ruler. R1 library composition may be the right answer when engineering time is scarce. R3/R4 work is justified only when the bottleneck is large enough that the recovered GPU-seconds, rack-seconds, or megawatt-hours outweigh the cost of going deeper. The level is not a status symbol. It is a bet on where the workload is wasting capacity.

If you're going to be precise about automated kernel generation, you have to be precise about which level of CUDA or HIP you're operating at: "AI-generated kernels" means very different things at different abstraction levels, and the most important number in any of the headlines below: "3.6× geomean," "120× peak," "100% pass on KernelBench": does not mean anything until you know the level.

Standard Kernel's R-axis rubric is the best vocabulary the field has. It is worth slowing down here, because this is the map the rest of the post and most of the agentic-kernel cohort in §08 uses. The easiest mental model for it, in our experience, is a kitchen.

The recipe analogy: R1 through R4 as four ways of producing the same dish.

Imagine the task is to produce a specific dish: let's say a precisely-timed, precisely-temperatured pan-seared steak. There are four levels at which a person can approach that task, and they're directly analogous to the four levels of kernel writing.

R-LEVEL    THE KITCHEN ANALOGY                                       THE KERNEL VERSION
────────   ────────────────────────────────────────────────────────  ─────────────────────────────────────────
R1         You ORDER THE DISH from a restaurant that already         You CALL A VENDOR LIBRARY (cuBLAS,
LIBRARY    serves it perfectly. You don't cook. You compose a        cuDNN, CUTLASS, AITER, rocBLAS. The
COMPOSITION meal by picking restaurants that each do one course      kernel was already written by NVIDIA or
            well. Fast, foolproof, fixed menu. Zero customization.   AMD experts. You pick which one to call
                                                                      and in what order. Zero kernel skill
                                                                      required; you're a meal composer.

R2         You COOK FROM A HIGH-LEVEL RECIPE in a smart kitchen      You WRITE A HIGH-LEVEL TILE DSL (Triton,
HIGH-LEVEL with appliances that handle most of the technique for     Triton, cuTile, Pallas, TileLang. You
DSL        you. You say "sear the steak at high heat for 4 min       describe the tile shape, the memory
            per side." The smart skillet handles the temperature      hierarchy, the loop structure. The
            curve, the timing, the heat distribution. You make        compiler picks which exact instructions
            high-level decisions; the appliance fills in everything   to emit, which swizzles to use, where to
            below. Same dish achievable; less skill required;         insert async copies and barriers. Same
            ceiling depends on how good the appliance is.             kernel achievable; less skill required;
                                                                      ceiling depends on the compiler.

R3         You COOK FROM A DETAILED RECIPE with manual technique.    You WRITE A LOW-LEVEL TILE DSL (CuTe,
LOWER-LEVEL "Pat the steak dry. Salt 40 min before. Cast-iron pan    CUTLASS CuTe in C++, CuTeDSL in Python,
DSL        on high for 5 min. 1 tbsp neutral oil. Sear 90 sec        ThunderKittens embedded in CUDA, Gluon
            without moving. Flip, 60 sec. Add butter, thyme,          on AMD. You name the MMA atom, the TMA
            garlic. Baste 30 sec. Rest 5 min on rack." You make       atom, the swizzle pattern, the producer/
            every technique decision explicitly. You need to know     consumer roles, the staging buffer
            why each step works. Ceiling: a skilled home cook.        depth. The compiler still does register
                                                                      allocation and instruction scheduling.
                                                                      Ceiling: a skilled kernel engineer.

R4         You COOK PROFESSIONALLY - sous-vide bath at exactly       You WRITE INSTRUCTION-LEVEL CODE -
INSTRUCTION 54.4°C for 90 min, blast chiller, induction sear at      CUDA C++ with inline PTX on NVIDIA,
LEVEL      precisely measured surface temperature, cryovac hold,     HIP C++ with inline AMDGPU on AMD.
            timed Maillard window, surface-temperature pyrometer,     You write the actual mma.sync /
            temporary probe. Every variable individually              tcgen05.mma / v_mfma_scale instruction
            controlled. Most home cooks can't do this. The ceiling    by hand. You manage mbarriers, scoreboards,
            is what's physically possible from the protein.           wait counters, register pressure,
                                                                      bank conflicts by hand. You ARE the
                                                                      compiler. Ceiling: the silicon itself
                                                                      (the SOL bound).

The single most important thing the analogy teaches. A restaurant dish (R1) and a sous-vide steak (R4) can be the same quality if the restaurant is good enough. But the moment you want something the menu doesn't serve (a precise doneness, an unusual cut, an off-menu sauce, a dietary constraint) the only way to get it is to go down the levels. R1 is foolproof but inflexible. R4 gives you complete control over everything and demands complete responsibility for everything. Every team in §08 is somewhere on that path, and the level they're at decides what's possible.

The same four levels, this time named with the actual libraries: for engineers.

LEVEL    NVIDIA EXAMPLES                            AMD EXAMPLES                       PORTABLE / CROSS-VENDOR
─────   ─────────────────────────────────────────  ─────────────────────────────────  ──────────────────────────────
R1      cuBLAS, cuBLASLt, cuDNN, cuFFT, cuSPARSE,  rocBLAS, hipBLASLt, MIOpen,        torch.* eager (calls vendor
        NCCL, TensorRT, TensorRT-LLM, CUTLASS      AITER, RCCL, ROCm Composable        libs underneath),
        as a library, FlashAttention 3 binary,     Kernel (CK) as a library            PyTorch nn.MultiheadAttention,
        FlashInfer as a library, FlashMLA binary,                                       JAX jax.numpy.* eager
        DeepGEMM, ThunderMLA binary

R2      cuTile (NVIDIA), Triton on NVIDIA          Triton-AMD, Pallas-on-AMD          Triton (OpenAI), TileLang
                                                   (very partial coverage)             (Microsoft), Pallas (JAX),
                                                                                       TVM, Halide,
                                                                                       torch.compile + Inductor
                                                                                       (Inductor emits Triton)

R3      CuTeDSL Python, CUTLASS CuTe in C++,       Gluon (AMD CuTe-shaped),           Mojo + Structured Mojo
        ThunderKittens (CUDA-embedded),            HipKittens, Composable Kernel       Kernels (Modular), MAX
        FlashAttention 3 source, FlashMLA          (CK) source                         kernels (Apache-2.0)
        source, DeepGEMM source

R4      CUDA C++ + inline PTX,                     HIP C++ + inline AMDGPU            SCALE (compiles CUDA + inline
        hand-written wgmma / tcgen05.mma,          assembly, hand-written              PTX to AMD machine code),
        mbarrier handshake by hand,                v_mfma_scale + ds_read_b64_tr_b4,   then runs on a different chip
        cp.async.bulk by hand                      LDS allocation by hand

How PyTorch and torch.compile sit across the ladder. A vanilla PyTorch user (no compile, no custom kernels) lives at R1: every tensor op dispatches to a vendor library or a per-op CUDA kernel that NVIDIA wrote. The moment they add @torch.compile, they move up the abstraction to R2 from the user's POV (the user still wrote eager Python), but the system that's now generating code on their behalf is operating at R2 (Inductor → Triton). When they write a custom Triton kernel by hand, they're operating at R2 explicitly. When they reach for ThunderKittens or CuTeDSL Python, they're at R3. When they crack open a .cu file with inline PTX, they're at R4. The cost-per-token math at each level differs by 2–10× when the workload is bottlenecked; the engineering cost differs by 10–100×.

Now: who in the §08 cohort is operating at which level. The receipts.

This is the part of the picture that has been missing from most coverage of automated kernel generation. "AI-generated kernels" as a phrase erases the R-axis. The eleven teams in §08 are doing very different things, and the differences are exactly R-axis differences. Worth being explicit:

TEAM / SYSTEM                          R-LEVEL          WHAT THEY OPTIMIZE                                  REPORTED HEADLINE
──────────────────────────────────────  ───────────────  ──────────────────────────────────────────────────  ─────────────────────────────────
doubleAI WarpSpeed                      R2 → R3 → R4     Full ladder: a trillion-parameter LRM with         doubleGraph (Mar 31, 2026):
(doubleGraph, Mar 31, 2026)             (PAC + search)   PAC verification + agentic "time-travel"           every algorithm faster, 55%
                                                         search rewrote NVIDIA cuGraph kernels across       above 2×, 18% above 10×;
                                                         A100, L4, A10G.                                    3.6× geomean over a decade of
                                                                                                            expert-tuned NVIDIA code.

Cursor multi-agent kernels              R2 → R4          Multi-agent harness on NVIDIA SOL-ExecBench's      149/235 (63%) outperformed
(Edward Lin, Apr 14, 2026)              (planner +       235 Blackwell B200 problems. Three-week            baseline; 38% geomean; 19%
                                        workers)         autonomous run on 27 B200s. Specific wins:        above 2×; BF16 GQA with paged
                                                         CUDA C++ kernels approaching the SOL bound        prefill hit 0.9722 SOL score
                                                         on attention and NVFP4 MoE primitives.            (84% over baseline).

MIT HAN Lab kernel-design-agents        R2 → R3          Triton-first kernels for MoE / DSA / GDN on        MLSys 2026 FlashInfer Contest:
(Dongyun Zou + Ligeng Zhu)                              B200 via the Humanize harness +                    1st on MoE, 2nd on DSA, 3rd on
                                                         KernelWiki provenance + ncu profiling.              GDN. Ablation showed harness
                                                         The agents author Triton and CUTLASS Cute;          (R2/R3 scaffolding) was the
                                                         the harness scaffolds them up to R3 systematically. dominant contributor.

UC Berkeley K-Search                    R2               Triton kernels on FlashInfer-Bench. Planner        2.10× geomean over OpenEvolve;
(Cao, Mao, Gonzalez, Stoica)                            (world model) decoupled from codegen (Triton        14.3× on Fused MoE; GPUMODE
                                                         policy). Insert / Update / Prune tree edits        TriMul SoTA at 1030 µs on H100
                                                         over the world model.                              in 300 iterations.

Meta KernelEvolve                       R1 → R2 → R3 → R4 Triton, Triton-TLX, CuTe DSL, plus low-level     100% pass on KernelBench; 100%
(Gang Liao, Carole-Jean Wu + FAIR)      (full stack)     CUDA, HIP, MTIA C++. Picks the right level         correctness 480 op-platform
                                                         per kernel automatically. The most R-axis-aware    configs; 1.2–17× speedups; 60%+
                                                         system in this cohort.                             Andromeda Ads throughput.

ByteDance / Tsinghua CUDA-Agent         R3 → R4          CUDA C++ generation, often dropping to inline      98.8% pass on KernelBench;
(cudaLLM-8B)                                            PTX. Skill-augmented agent + 128k context +        96.8% faster than torch.compile;
                                                         200-turn multi-turn RL.                            2.11× geomean.

CUDA-L1 / DeepReinforce                 R3               CUDA C++ on A100 with contrastive RL on            3.12× avg, 1.42× median, 120×
                                                         wall-clock reward. Cross-arch transfer            peak on KernelBench; 3.85× H100,
                                                         (H100 / L40 / RTX 3090 / H20) without retrain.    3.13× L40, 2.51× RTX 3090.

Sakana AI archive                       R2 → R3          30,615 CUDA kernels generated, profiled, and      The data point that proved the
                                                         shipped on Hugging Face. robust-kbench shipped     harness matters more than the
                                                         after their original benchmark was gamed.          generator.

AMD GEAK v3                             R2 → R3          Triton-AMD + HIP kernels on MI300/gfx950 via      54.89% accuracy + 2.59× on
                                                         agentic loop + rocprof-compute feedback.          TritonBench-modified; 11 of 30
                                                         AMD-native equivalent of WarpSpeed.               kernels beat human-expert.

Makora                                  R2 → R3          Multi-vendor (NVIDIA H100/B200, AMD MI300X,       Sub-60-second kernel generation,
                                                         Tenstorrent). Generation in <60s, then            continuous autotuning.
                                                         continuous autotune.

Standard Kernel                         R4 + analysis    PTX-layer hybrid system: program analysis +       RMSNorm-1024 ~67% faster than
(Anne Ouyang, Chris Rinard)             across all       LLM working directly on PTX, learning across     TileLang, Matmul-1024 ~5%
                                                         DSLs (Triton, TileLang, ThunderKittens,           faster than CUTLASS on H100.
                                                         CUTLASS) at the shared lower representation.

Wafer / KernelArena                     R2 → R3 → R4     Public benchmark + the team running the           #1 inference perf for
(Steven Arellano, Emilio Andere)        (benchmarks)     benchmark. Runs WaferBench NVFP4 on B200,        Qwen3.5-397B-A17B on MI355X
                                                         KernelBench HIP on MI300X. Same team               via the full AMD stack
                                                         producing #1 inference performance result on      (ROCm 7 + AITER + hipBLASLt +
                                                         AMD flagship hardware.                            Triton-AMD).

The direct read across the table. Most "AI-generated kernels" today is R2: Triton, with a compiler picking the instructions. R2 is real, ships, and is improving fast. The vendor libraries (R1) live a layer below at R3 and R4, written by humans at NVIDIA and AMD over years. The interesting frontier is the bottom of the table: R3 and R4: because that's where the vendor performance ceiling actually lives, and that's where the last 2–10× cost-per-token reduction sits. Meta KernelEvolve and doubleAI WarpSpeed are the two systems in this cohort that explicitly span R1 through R4; everyone else is concentrated in R2–R3. Touchdown is targeting R3 and R4 specifically: both on NVIDIA Blackwell and on AMD CDNA4, both in the kernel-authoring work above and in the verification-and-evidence layer underneath all of it.

Why R3 and R4 are so hard: the hardware visualization, three chips side by side.

R3 and R4 are hard because "fast" is not abstract. It is a property of the silicon, and the silicon is wildly different at the SM / CU level across the three architectures everyone in this post benchmarks on. R2 is portable because the compiler hides the differences. R3 and R4 force you to write into the differences. Worth looking at the three chips side by side.

The honest implication. An R2 (Triton) kernel can be the same source file on all three chips, but the lowering (and therefore the performance ceiling) changes wildly. An R3 (CuTe / CuTeDSL / Gluon / ThunderKittens) kernel has to name the right MMA atom, the right SMEM staging pattern, the right cooperative primitive, per architecture. An R4 (raw PTX / inline AMDGPU) kernel is a different program on each chip, because the instructions themselves are different (wgmma on H100, tcgen05.mma on B200, v_mfma_scale on MI355X). That's the hardware reality the agentic cohort in §08 is running into. Why three different teams ship three different stacks: because the silicon they're optimizing for is three different shapes underneath.

Now make it concrete: five real 2026 open-weights model families, two real workloads, three chips.

Proof level for this subsection. This is architecture-shape analysis, not the public Kimi K2.5 InferenceX row. The public measured spine is the 8k/1k GB200 row used in §02.5 and §15. The workload tables below use larger illustrative chat and coding shapes to show which bottleneck you should benchmark next. Do not price a deployment from these tables. Use them to decide what to measure: attention, MoE expert GEMM, KV residency, linear-attention scan, precision path, or hardware placement.

The abstract argument only matters if it explains the models people are actually running. Do the homework here, because the five leading open-weights frontier families all made different architectural bets in 2025–2026, and those bets land in different places on the R-axis and on different chips. The families we'll walk through:

• Moonshot Kimi/K2-style MoE example: 1T total, 32B activated, 384 experts (8 + 1 shared), 61 layers, MLA attention, 256K context, native INT4 via QAT. This is an architecture-shape example, separate from the public Kimi K2.5 benchmark rows used later for measured throughput.
• Alibaba Qwen 3.5 / 3.6 family: multiple sizes, all sharing the same hybrid Gated DeltaNet + sparse MoE architecture with multimodal vision-language foundation. Headline checkpoints:
  • Qwen3.5-397B-A17B (Feb 16, 2026): 397B total, 17B activated, multimodal flagship
  • Qwen3.5-122B-A10B / 35B-A3B / 27B dense (Feb 24, 2026)
  • Qwen3.5-9B / 4B / 2B / 0.8B dense (March 2, 2026)
  • Qwen3.6-35B-A3B (April 16, 2026): "Agentic Coding Power"
  • Qwen3.6-27B (April 22, 2026)
  • Qwen3-Next-80B-A3B (Sept 11, 2025): 80B total, 3B activated (ultra-sparse 1:50), 512 experts (10 + 1 shared), 48 layers, 262K native context extensible to 1M, MTP: still the sparsest live MoE in the family
• OpenAI GPT-OSS 120B / 20B (August 2025): OpenAI's first open-weight reasoning models, Apache 2.0:
  • gpt-oss-120b: 116.8B total, 5.1B activated per token (1:23 sparsity), 36 layers, 128 experts, top-4 routing, GQA 64Q / 8KV with alternating banded-window + dense attention (128-token bandwidth), 131K context via YaRN, native MXFP4 quantization. Fits on a single 80 GB GPU (H100 or MI300X).
  • gpt-oss-20b: 20.9B total, 3.6B activated (1:6 sparsity), 24 layers, 32 experts, top-4. Fits in 16 GB memory. The small-end open-weights reasoning model for laptops and edge.
• Z.AI GLM-4.6 / 4.7 / 5.1: GLM-4.6 at 357B total, 32B activated, 160 experts (8 + 1 shared), 92 layers, GQA 96Q / 8KV with QK-Norm and partial RoPE, 200K context. The strongest open-weights coding model right now.
• DeepSeek V4-Pro / V4-Flash (April 2026): Pro at 1.6T / 49B activated, Flash at 284B / 13B activated, 256 routed experts (6 per token), 61 layers, hybrid CSA + HCA + mHC attention (replaces V3's MLA), 1M context default, FP4 expert weights. The longest-context default in the open ecosystem.

Five families, six attention architectures, parameter counts spanning 0.8B → 1.6T, sparsity ratios from dense to 1:50. That's the actual menu open-weights serving teams are picking from in 2026: and the right kernel/engine/cache choices change on every cell.

Four very different architectures. The first row of the table below is what every CFO needs to see: same chip, same workload, four wildly different cost profiles because of the architectural choices each lab made.

ARCHITECTURE                  KIMI/K2-STYLE       QWEN3-NEXT-80B          GLM-4.6                DEEPSEEK V4-PRO
─────────────────────────     ──────────────      ──────────────────      ───────────────        ─────────────────
Total params                  1T                  80B                     357B                   1.6T
Active params per token       32B                 3B (1:50 sparse)        32B                    49B
Experts (routed + shared)     384 + 1             512 + 1                 160 + 1                256 (no shared)
Experts per token             8                   10                      8                      6
Layers                        61                  48                      92                     61
Attention mechanism           MLA                  Hybrid (Gated DeltaNet  GQA 96Q/8KV           Hybrid (CSA + HCA)
                                                   + Gated Attention)      + QK-Norm              + mHC residuals
Context (native)              256K                 262K (→ 1M)            200K                   1M
Expert weight precision       INT4 (QAT)           BF16                    BF16                   FP4 + FP8 mixed
KV cache shape (per token,    Compressed (MLA      Shrinking with         Standard GQA           Compressed 4× (CSA)
per layer)                    latent dim)         linear-attention                                + 128× (HCA)

Read each row carefully. The Kimi/K2-style example has 12× the total parameters of Qwen3-Next, but Qwen3-Next is 10× sparser in activation. GLM-4.6 has half the experts of the Kimi-style example but more shared depth. DeepSeek V4-Pro has 50× the total parameters of Qwen3-Next, the longest context of all four, and the most aggressive attention compression. None of these is "better": they're four different bets, and the right kernel stack is different for each.

Workload A: long-context CHAT (200K-token prompt, 1K-token response).

Think 200,000 tokens of pasted documentation, requirements, transcripts, or chat history followed by a short conversational answer. Roughly what a customer-support agent or research-assistant agent does. Prefill-dominated workload: almost all the compute is in processing the 200K input tokens; the 1K output is cheap. Same workload, all four models, on a single B200 node:

WORKLOAD A: 200K prompt / 1K response - single B200 (192GB HBM3e, ~5 PFLOPS NVFP4)
                              KIMI/K2-STYLE       QWEN3-NEXT-80B          GLM-4.6                DEEPSEEK V4-PRO
─────────────────────────     ──────────────      ──────────────────      ───────────────        ─────────────────
Prefill FLOPs                  ~262 PFLOPs         ~25 PFLOPs              ~262 PFLOPs            ~109 PFLOPs (CSA
(active × 200K × layers)      (32B × 200K × 61)   (3B × 200K × 48)        (32B × 200K × 92)      cuts attn ~73%)

KV cache after prefill         ~205 GB (MLA       ~18 GB (DeltaNet        ~330 GB (full GQA)     ~33 GB (CSA + HCA
(approximate, BF16)            compressed)        compresses heavily)                             4× + 128× compression)

KV fits in 1× B200 HBM?        Yes                Yes                     Yes (barely)           Yes (with headroom)

THE BOTTLENECK kernel          MoE GEMM           Linear-attention        GQA + MoE GEMM         CSA / HCA attention
on prefill                     (8 experts × 200K  scan kernel             (96 heads × 200K       kernel (sparse
                                tokens / expert)  (the rare workload      tokens - KV pressure   selection over
                                                  where attention isn't   serious at 200K)       compressed blocks)
                                                  the bottleneck)

R-level that decides cost      R3/R4 MoE GEMM     R3/R4 linear-attn       R3/R4 GQA attention    R3/R4 hybrid-attn
                              (FlashMLA-shaped)  scan (genuinely new      kernel + R2 MoE        kernel (brand new
                                                  kernel territory)       routing                shape, NO vendor
                                                                                                  reference yet)

Decode 1K tokens               Bandwidth-bound    Bandwidth-bound          Bandwidth-bound +     Bandwidth-bound,
(after prefill)                MoE decode kernel  on 3B weights × 1K       large KV streaming     but KV is now
                                                  tokens (cheapest of                              compressed 4× /
                                                  four)                                            128× → cheaper

Where the R3/R4 work pays      Big - MoE expert   Modest - small active    Big - KV-streaming    HUGE - the attn
off most                       GEMM is 50% of     model, attention not     and large GQA          mechanism itself
                                runtime           the bottleneck           dominates              is the bottleneck

What this row teaches. The four models look superficially similar (all MoE, all long-context) but the kernel that decides cost is different for every one of them. Kimi/K2-style MoE: the MoE expert GEMM kernel is the lever. Qwen3-Next: the linear-attention scan kernel is the lever, and that kernel didn't exist in vendor libraries a year ago; it's an open kernel-research problem. GLM-4.6: the GQA attention kernel and the MoE routing kernel together. DeepSeek V4: the hybrid CSA+HCA attention kernel, which is a brand-new shape that NVIDIA hasn't shipped a reference for, and where the open kernel community (FlashInfer + DeepEP) is racing to catch up. That's why §08's automated kernel cohort matters now and not five years ago. The model architectures are moving faster than the vendor kernel libraries can keep up.

Workload B: long-context CODING (100K codebase prompt, 500-token code response).

Think Claude Code or Cursor: the agent reads 100K tokens of repo context, then emits 500 tokens of a function. The interesting wrinkle versus chat: the prefix is reused across many turns of the same agent session. A 100K codebase context is loaded once, then served against 50+ short turns. That changes the picture entirely.

WORKLOAD B: 100K prompt / 500-token response × 50 turns - single B200 with KV cache reuse
                              KIMI/K2-STYLE       QWEN3-NEXT-80B          GLM-4.6                DEEPSEEK V4-PRO
─────────────────────────     ──────────────      ──────────────────      ───────────────        ─────────────────
Prefill cost (1st turn)       ~131 PFLOPs        ~13 PFLOPs               ~131 PFLOPs           ~55 PFLOPs
Prefill cost (turns 2-50)     Near-zero with     Near-zero with prefix    Near-zero with        Near-zero with
                              prefix cache        cache                    prefix cache          prefix cache

KV cache to keep resident     ~103 GB             ~9 GB                   ~165 GB                ~16 GB
across the 50-turn session    (MLA compressed)   (DeltaNet)              (full GQA, painful)    (CSA + HCA)

How many concurrent           ~1–2 sessions      ~20+ sessions           ~1 session             ~10+ sessions
agent sessions fit            per B200            per B200                per B200               per B200
in 1× B200 HBM?

THE BOTTLENECK kernel         MoE decode +       Linear-attention        KV streaming from      Hybrid-attn decode
on the 500-token decode       MLA decode         decode kernel           HBM (KV is huge        kernel + FP4 expert
                                                                          relative to weights)   GEMM (the FP4 path
                                                                                                  is genuinely new)

How much each %-gain in       Each 10% kernel    Each 10% kernel win =   Each 10% kernel win    Each 10% kernel
the right kernel matters      win = ~9% lower    ~10% lower cost/token   = ~5% lower            win = ~8% lower
                              cost/token         (highest leverage -     cost/token (KV         cost/token (FP4
                                                  smallest active model,  streaming dominates,   path = highest
                                                  kernel IS the cost)     not kernel compute)    growth area)

Where KV-cache offload        Modest - fits OK   Low - KV is tiny        HUGE - KV doesn't     Modest - CSA/HCA
matters most                  in HBM             relative to weights      fit, must offload      already compresses
                                                                                                  aggressively

The direct takeaway across both workloads. Architectural choices made in the model design (MLA vs GQA vs DeltaNet vs CSA+HCA, dense vs ultra-sparse activation, expert count, KV cache shape) determine which kernel becomes the bottleneck. The same B200 running these four models has four different "most-important kernel" answers. Kimi/K2-style MoE on B200 = MoE expert GEMM is the lever. Qwen3-Next on B200 = linear-attention scan is the lever. GLM-4.6 on B200 = GQA attention + KV streaming. DeepSeek V4-Pro on B200 = hybrid CSA+HCA attention is the lever, and it's a kernel the vendor doesn't have a reference for. That's why automated kernel optimization (§08) and verifiable evidence formats (this whole post) both matter right now; the kernel that decides cost is different per model, per workload, per chip generation, and the rate of new architectures shipping is faster than NVIDIA or AMD can write vendor kernels for each combination.

The same four models on three different chips: and why the answer changes.

The picture changes again when the chip changes. The same Kimi/K2-style MoE shape decoding 100K context behaves very differently on B200 vs H200 vs MI355X. Quick read across:

KIMI/K2-STYLE DECODE @ 100K CONTEXT 1× B200             1× H200             1× MI355X
                                     (192 GB / 8 TB/s)   (141 GB / 4.8 TB/s) (288 GB / 8 TB/s)
─────────────────────────────────    ─────────────       ─────────────       ──────────────
Active weights (32B BF16)            64 GB               64 GB               64 GB
KV cache for 100K (MLA)              ~103 GB             ~103 GB             ~103 GB
TOTAL HBM needed                     167 GB              167 GB              167 GB
Fits in 1 chip?                      Yes (25 GB free)    NO - must shard     Yes (121 GB free)
                                                          across 2 H200s
Bottleneck                           BW-bound on the     BW-bound + NVLink   BW-bound on
                                     decode kernel       cross-chip KV       MI355X decode
                                                          fetch                kernel
Vendor reference kernel              cuBLASLt / DeepGEMM cuBLASLt / DeepGEMM AITER + ROCm-CK
exists?                              Yes (FlashMLA-      Yes (older)         Partial - Gluon
                                     shaped on B200)                          + tile-lang for
                                                                              MFMA scale path
R3/R4 work that pays off most        tcgen05.mma decode  wgmma decode kernel  v_mfma_scale
                                     kernel for MLA      for MLA              decode kernel for
                                     (tile shapes B200-   (tile shapes Hopper- MLA (CDNA4 layout)
                                     specific)            specific)

QWEN3-NEXT-80B DECODE @ 100K        1× B200             1× H200             1× MI355X
─────────────────────────────────    ─────────────       ─────────────       ──────────────
Active weights (3B BF16)             6 GB                6 GB                6 GB
KV cache (linear-attn, shrinking)   ~9 GB                ~9 GB               ~9 GB
TOTAL                                15 GB               15 GB               15 GB
Concurrent sessions fit              ~10                 ~7                  ~15
Bottleneck                           Linear-attn scan    Linear-attn scan    Linear-attn scan
                                     kernel              kernel              kernel
Vendor reference kernel              NO - open kernel    NO                  NO
exists?                              territory; Triton
                                     impl is the open
                                     baseline
R3/R4 work that pays off most        Custom Gated         Custom Gated        Custom Gated
                                     DeltaNet kernel      DeltaNet kernel     DeltaNet kernel
                                     in ThunderKittens    in ThunderKittens   in HipKittens or
                                     or CuTeDSL on B200   on H200             tile-lang on AMD

DEEPSEEK V4-PRO DECODE @ 1M         8× B200 NVL72      8× H200             8× MI355X
─────────────────────────────────    ─────────────       ─────────────       ──────────────
Active weights (49B FP4)             ~25 GB              N/A (no FP4 on      ~25 GB
                                                          H200 widely
                                                          shipped)
KV cache for 1M tokens (CSA+HCA)    ~75 GB               (V3.2 fallback,      ~75 GB
                                                          much higher KV)
Vendor reference for CSA+HCA?       NO - brand new       N/A                 NO
                                     shape; DeepEP +
                                     FlashInfer racing
                                     to ship
R3/R4 work that pays off most        Custom CSA+HCA       Run V3.2 instead,   Custom CSA+HCA
                                     attention kernel,    where MLA has a     kernel on CDNA4
                                     and the FP4-MoE      vendor reference   matrix cores
                                     expert GEMM
                                     (NVFP4 path)

The point this whole subsection drives at. "What's the right kernel optimization to do" is a six-variable question: model architecture (MLA / GQA / DeltaNet / CSA-HCA) × parameter count × expert sparsity × chip generation (B200 / H200 / MI355X) × workload shape (chat / coding / agent / RAG) × precision (BF16 / FP8 / FP4 / NVFP4 / INT4-QAT). Six variables, each with 3–10 values, gives roughly 5,000–10,000 combinations. NVIDIA and AMD cannot write hand-tuned vendor kernels for every cell in that table. Neither can any single research lab. That's the gap automated kernel generation is filling, and that's why the verification-and-evidence layer underneath it has to be solid: because nobody will accept "the agent said it's faster" as proof when the search space is 10,000 cells. This is the planning layer, not the measured-result layer. The measured layer is where a Kimi K2.5-style receipt names the exact workload, exact config, exact hardware, exact result row, money math, energy proxy, and caveat.

Touchdown's R3/R4 work targets exactly these cells. Not "GEMM on B200" generically; the specific MoE expert GEMM for FP4 weights on a 49B-active model with a hybrid sparse-attention frontend, or the CDNA4 MFMA-scale path for an MLA decode kernel at 100K context, or the linear-attention scan kernel for a Gated DeltaNet layer on Hopper-generation tile shapes. Each one is a different program. Each one needs an honest baseline (different per chip), an honest verification harness (different per architecture), and a replay-able evidence packet so the next team doesn't have to re-do the measurement from scratch. That is the open job. And once the job is structured this way, it is exactly the kind of work that gets automated.

The recipe analogy, one more time, this time for the integration argument.

Step back. The kernel landscape in 2026 looks chaotic because there are roughly forty named libraries and DSLs and compilers and inference engines (the full map is in §08.45). The R-axis collapses all forty of them into four positions, and the kitchen analogy collapses the four positions into one mental model anyone can hold. Every team in §08 is some flavor of cook working at one or more of the four levels. The doubleAI WarpSpeed team is a Michelin-grade restaurant working at R2 / R3 / R4 simultaneously, with the last-mile RL doing the sous-vide-and-pyrometer R4 work. MIT HAN Lab's kernel-design-agents is a brigade-style kitchen working at R2 / R3 with a tightly orchestrated mise-en-place (the Humanize harness). Meta KernelEvolve is running a multi-cuisine kitchen across NVIDIA, AMD, and MTIA, spanning R1 through R4. Touchdown is doing kernel research too, but the open-source lane we are leaning into is the testing, measurement, and evidence standard: the people who taste, time, weigh, and record what every dish did, so the whole kitchen can learn from data instead of intuition alone.

The point of the analogy isn't that one level is better than another. The point is that every level has its own ceiling, its own engineering cost, and its own degree of automation possibility: and getting clarity on which level you're at is the precondition for any honest claim about cost-per-token reduction.

Why this is solvable: kernels are a verifiable domain.

The reason any of this matters beyond GPU kernels is that automated kernel optimization is the strictest verifiable domain in software right now. Verifiable in the strict sense: every kernel either compiles or doesn't, either produces a numerically-correct output or doesn't, and either runs faster than a vendor baseline or doesn't. Those are three binary gates plus one continuous wall-clock measurement, on real silicon, with no human in the loop. That's the rarest property in machine learning today.

Most domains an LLM operates in (writing, summarization, customer service, sales emails, creative work) have no automatic verification. The reward signal is human preference, which is slow, expensive, biased, and saturates fast. RLHF has carried the field a long way but it has real ceilings. Verifiable domains don't have that ceiling because the reward signal is mechanical. A kernel optimizer can run 100,000 candidates per day on real B200s, score each one against a vendor baseline at wall-clock truth, and learn what actually works without any human ever looking at the candidates. That's why §08 exists. That's why 2026 is the year kernel optimization started getting genuinely automated.

The strongest published map of verifiability we know is in §03 (the five-level chart from "kernels on real silicon" at the strict end to "subjective evaluation" at the abstract end). Kernel optimization sits at level 1. Math and code with test suites are at level 3. Tool-use trajectories are at level 4. Subjective evaluation is at level 5. The whole field is moving up that map. The same techniques that work for kernel optimization at level 1 (gate-spec correctness, multi-objective reward shaping, anti-reward-hacking harness design, drift detection, evidence-packet replay) generalize one level up at a time. Today: kernels. Next: full inference-engine optimization. Then: serving-stack optimization. Then: scheduler-and-router optimization. Then: agent-trajectory optimization. Each level reuses the verification machinery from the level below.

The point this whole post is built around. Automated kernel optimization in 2026 is the foundation move for automated optimization of any inherently-verifiable domain. The reason is the same reason ImageNet was the foundation move for vision in 2012: once a domain has a measurable reward signal, a public benchmark, a hardened harness, and an evidence format, the optimization problem becomes a search problem, and search problems compound. The hardware is what makes the reward signal honest. The harness is what makes the reward signal hard to cheat. The evidence layer is what makes the result re-verifiable by anyone with the same silicon. All three together turn kernel optimization into a domain where automation is no longer the question: it's the inevitability.

And the wider read. Once a small number of teams crack automated optimization on a verifiable domain like kernels, the verification machinery they build: gate specs, harness design, drift telemetry, anti-reward-hacking discipline, evidence-format standardization: becomes the toolkit for tackling unverifiable domains too. The honest path from "LLMs are good at code" to "LLMs are reliably useful outside their training distribution" almost certainly runs through this work. Verifiable domains teach the field how to design reward signals that don't collapse under search pressure. That skill is what eventually lets RL work on less-verifiable problems: research, writing, judgment, taste, planning under uncertainty: without the rewards being instantly gamed. Kernel optimization is the laboratory where the discipline gets built. Everything above the laboratory benefits from how carefully the discipline is enforced inside it.

That's why the rest of this post is about the harness, not the generator. Every serious team in §08 has a generator that can produce candidates. What separates the ones whose results will hold up six months from now from the ones that won't is the harness. And the harness (gate spec, evidence format, replay command, hardware-state snapshot, named failure catalog, drift detector) is exactly what the verifiable-domain thesis says is the piece that holds it all up. Before we look at those teams, it helps to look at the silicon they are actually targeting.

The §07 figure compared three chips at the SM/CU level. Before the §08 cohort, zoom out, because the most expensive misreads in this space come from treating "a GPU" as one shape. A B200 is not a GB200, a GB200 is not an NVL72 rack, and an MI355X is none of those things. Four visualizations below, each at a different scale, mapping how the modern inference rack is actually put together: chip package, superchip module, rack-coherent fabric, and the CPU that sits next to it. Companion to the prose figures in §18 (NVIDIA Vera Rubin) and the AMD-side §08.556 walkthrough.

This is the business meaning of the CPU coming back: agentic workloads spend real wall-clock time outside the GPU. That CPU-side work affects latency, utilization, and cost per task.

Hardware topology is also why the Kimi K2.5 GB200 example later is not just a bigger-GPU story. The measured difference comes from the serving path using the rack fabric differently: expert parallelism, disaggregated prefill/decode, NVLink domain size, and communication placement. A CEO sees capacity. A CFO sees rack-equivalent headroom. A CTO sees architecture risk. A kernel engineer sees which local win still has to survive the fabric.

topology audit:
  hardware target       -> NeuronLink ring | NVSwitch | Infinity Fabric | TPU fabric | future ASIC
  model parallelism     -> TP width, EP width, sequence/ring attention, activation sharding
  communication path    -> all-reduce, all-gather, all-to-all, reduce-scatter, KV transfer
  software proof        -> compiler schedule, kernel path, runtime placement, p95/p99 receipt

B200-style question:
  does expert traffic spill outside the fast local island?

GB200 NVL72-style question:
  can the rack fabric keep wider expert parallelism useful?

audit question:
  did the workload get faster because math improved,
  or because topology stopped communication from dominating?

The total-cost story underneath: what SemiAnalysis just published (May 23, 2026).

The diagram only matters if it explains the bill. SemiAnalysis posted the receipts that make that mapping concrete the same week this post went out ("The Coding Assistant Breakdown: More Tokens Please"), and the numbers underneath sharpen every previous figure in this section.

Read this against FIG 07.5-D above. The Grace → Vera shift NVIDIA shipped is not architectural decoration; it's silicon shaped for the workload SemiAnalysis just measured. If 42% of agentic-coding wall-clock is CPU work, and Fairwater pulls 16% of its total power into CPU + storage for a GPU cluster, then the rack-coherent CPU pool sitting next to the GPUs is no longer a host. It's a co-equal compute tier with its own cost line on every successful task. The §02.55 / §02.6 / §02.65 framing earlier in this post is exactly the unit-economics version of that hardware shift, and SemiAnalysis' numbers are the operator-side receipts that the gap is already showing up on real bills.

Vera Rubin · step-by-step. What a coding-agent task actually does on the next-gen NVIDIA rack.

Concrete walkthrough, because the abstract picture only lands when you trace one task end-to-end. Use the same user-facing workload from the top of the post: Hermes/OpenClaw asks Claude Code to build a mobile app screen for AI skincare progress. The agent needs the repo, the Expo routes, the design system, TypeScript errors, file diffs, and screenshot feedback. Under the hood, assume a single coding-agent task with roughly 100K tokens of repo/app context, a 500-token patch response, and one tool-call/test cycle. Hardware: Vera Rubin NVL72 (72 Rubin GPUs + 36 Vera CPUs, NVLink 6 fabric, ~120 kW liquid-cooled rack).

The important split is context stability. Stable context is the system prompt, Claude Code rules, Hermes/OpenClaw tool schema, product requirements, and design-system rules. Semi-stable context is the file tree, package.json, app routes, React Native / Expo components, design tokens, API types, and state-management files. Volatile context is the latest user message, TypeScript error, Metro or simulator error, latest diff, latest test result, or screenshot feedback. The rack-level question is whether the stable and semi-stable state stays reusable while the volatile state changes.

VERA RUBIN NVL72 · ONE AGENTIC CODING TASK · STEP-BY-STEP

[t=0.00s]   USER TASK ARRIVES at the inference endpoint.
            Router picks a (Vera CPU, Rubin GPU) pair from the rack pool
            based on prefix-cache locality and current load.

[t=0.01s]   VERA CPU (88 Olympus cores) wakes a dedicated agent sandbox
            - one of 22,500+ concurrent CPU environments per rack.
            Loads system prompt + tool schemas + MCP defs from local memory.
            Reads 100K tokens of repo context from CPU LPDDR5X.

[t=0.03s]   VERA → RUBIN GPU HANDOFF over NVLink-C2C.
            Coherent address space: zero-copy of prompt + KV cache pointers.

[t=0.03s]   PREFILL on Rubin GPU.
            100K tokens · 49B active params (assume Kimi/K2-style MoE)
            → ~131 PFLOPs of FP8 work at ~3500 TFLOPS sustained
            → ~38 ms prefill wall time.
            KV cache written: ~103 GB into HBM4.

[t=0.07s]   DECODE on Rubin GPU.
            500 tokens generated, bandwidth-bound at HBM4.
            ~10 ms per token batch at p50 interactivity.
            → ~500 ms decode wall time.

[t=0.57s]   TOOL CALL ISSUED. Model emits {"tool": "run_tests", ...}
            GPU returns control to Vera CPU via NVLink-C2C.

[t=0.58s]   VERA CPU executes the tool. (HERE IS THE 42%.)
            Spawns sandboxed pytest run on local cores.
            Bash, file I/O, test harness, lint, package install.
            → 2.1 seconds of CPU wall time.
            GPU sits warm but idle on this sandbox - other sandboxes
            on the same NVL72 keep the GPU pool fed.

[t=2.68s]   TOOL RESULT RETURNED. Failed test output (1.2K tokens).
            Vera CPU re-enters the agent context with the failure
            appended to the prompt. New prefill ahead.

[t=2.69s]   PREFILL #2 on Rubin GPU.
            1.2K-token incremental prefill (prefix cache hit on the
            original 100K) → ~5 ms.

[t=2.70s]   DECODE #2. 800 tokens of patch + reasoning → ~800 ms.

[t=3.50s]   SUCCESSFUL TASK. Patch returned to user.
            Vera CPU writes evidence: trace, KV-cache footprint,
            tool-call result, per-step latency, success/failure flag.
            Sandbox parked, ready for next task in the queue.

────────────────────────────────────────────────────────────────────
TOTAL: 3.50 s · 1.34 s GPU work · 2.16 s CPU work · 1 retry cycle
WALL-CLOCK SPLIT: ~38% GPU · ~62% CPU (matches the 42/58 distribution
                  after amortizing across the queue of concurrent tasks
                  the same Vera+Rubin pool is serving)

Read the timing carefully. The GPU was busy for 1.34 of the 3.50 seconds. The CPU was busy for 2.16. For this one task, the CPU was the bigger consumer of wall-clock time. The reason Vera CPU exists as a first-class processor in the NVL72 rack is that the CPU isn't waiting on the GPU anymore; the GPU is increasingly waiting on the CPU, and at rack scale you want the CPU pool sized so that doesn't happen. The Fairwater 1:6 ratio above is the silicon-floor allocation that prevents this exact GPU-idle failure mode at scale.

AMD MI355X · step-by-step. Same task, different topology, no NVLink-coherent CPU.

Same workload, same model class, AMD-side. 8-OAM HGX-class platform (8× MI355X GPUs over Infinity Fabric XGMI, x86 host CPU over PCIe). The architectural difference is the CPU↔GPU path.

AMD MI355X 8-OAM · ONE AGENTIC CODING TASK · STEP-BY-STEP

[t=0.00s]   USER TASK ARRIVES at the inference endpoint.
            x86 host CPU (typically AMD EPYC) wakes a Linux sandbox.

[t=0.01s]   x86 HOST loads system prompt + tool schemas + MCP defs.
            Reads 100K tokens of repo context from DRAM.

[t=0.04s]   HOST → MI355X HANDOFF over PCIe 5.0 x16.
            ~32 GB/s per link (vs NVLink-C2C's 900 GB/s).
            Higher per-token transfer cost; matters most for
            cold-start prefill or cross-node KV moves.

[t=0.04s]   PREFILL on MI355X GPU.
            100K tokens · ~49B active params, MFMA-scale FP8 path
            on CDNA4 matrix cores → ~131 PFLOPs at ~3500 TFLOPS
            → ~38 ms prefill wall time (parity with Rubin estimate).
            KV cache: ~103 GB into HBM3e.

[t=0.08s]   DECODE on MI355X.
            500 tokens, bandwidth-bound at 8 TB/s HBM3e.
            ~10 ms per token batch → ~500 ms decode wall time.

[t=0.58s]   TOOL CALL ISSUED. GPU returns control to x86 host
            over PCIe - the boundary the NVIDIA side avoids.

[t=0.61s]   x86 HOST executes the tool. (THE 42% WORK.)
            Spawns sandboxed pytest run on host EPYC cores.
            → 2.1 seconds of CPU wall time.
            GPU sits warm but idle on this sandbox; same caveat as
            NVIDIA - other sandboxes keep the GPU pool fed.

[t=2.71s]   TOOL RESULT RETURNED to x86 host. Failed test output.
            Re-enter context, prepare prefill #2.

[t=2.74s]   PREFILL #2 + DECODE #2 → ~5 ms + ~800 ms.

[t=3.55s]   SUCCESSFUL TASK. Patch returned to user.
            x86 host writes trace + evidence.

────────────────────────────────────────────────────────────────────
TOTAL: 3.55 s · 1.34 s GPU work · 2.16 s CPU work · ~50 ms PCIe overhead
WALL-CLOCK SPLIT: same 38/62 GPU/CPU
RACK-COHERENCY: no - bounded at 8 GPUs per server

The comparison that matters. For this one task on a single server, the AMD MI355X path is within ~1.5% of the Vera Rubin path on total wall-clock: and the GPU portion is parity. The differences live elsewhere: (1) the CPU↔GPU boundary crosses PCIe instead of NVLink-C2C, which costs ~50 ms per round-trip and starts to matter on tool-call-heavy workloads; (2) AMD has no rack-coherent NVLink-equivalent at this rev, so workloads above 8 GPUs require model sharding across a network rather than treating 72 GPUs as one device; (3) the CPU pool is whatever x86 you put in the chassis, sized to host duty rather than purpose-built for agent environments. For workloads that fit in 8 GPUs and aren't crossing the CPU boundary every few seconds, MI355X is fully competitive. For workloads that span 72 GPUs in one coherent device or hit the CPU↔GPU boundary thousands of times per task, Vera Rubin's NVLink-coherent CPU+GPU pool is a real structural advantage NVIDIA gets to charge for.

Compute element zoom: one SM vs one CU, full internals.

One more level of detail because the kernel-writing work in §08.5 / §08.555 / §08.556 only makes sense if you've seen what a single SM and a single CU actually contain. NVIDIA calls it an SM (Streaming Multiprocessor); AMD calls it a CU (Compute Unit). Different names, similar role, very different internals: and the difference is what makes a kernel written for one structurally not portable to the other at the depth tier.

AMD MI355X 8-OAM platform: different shape at the server level.

AMD's verified wins on real workloads: receipts, not impressions.

The hardware diagrams above show the architecture. The spec table below shows the silicon. This block shows the actual benchmark results where AMD MI355X beats NVIDIA B200 in 2026, and where the honest caveats sit: because the only way to make the catch-up story credible is to name the workloads where the catch-up has already happened.

The pattern across all four sources. Where AMD MI355X wins today: single-node large MoE workloads where 288 GB removes cross-GPU shuffling (GPT-OSS 120B, DeepSeek R1, Llama 3.1 405B); FP8 disagg prefill at parity-to-better; high-concurrency throughput on perf-per-TCO; real-time interactive latency on Llama-class dense models. Where NVIDIA wins today: multi-node FP4 disaggregated serving (the composability gap), rack-coherent NVL72 workloads, software-ecosystem breadth (CUDA's 20-year compound), and workloads with heavy CPU↔GPU boundary crossing (the Vera Rubin advantage from FIG 07.5-D). The honest read: AMD's hardware is competitive-to-leading on the workloads where its strengths fit; NVIDIA's hardware extends further at the largest serving scales and on FP4-disagg. No vendor wins universally; the workload decides. Which is exactly the kind of question the §02.6 cost-per-successful-task framing exists to answer rather than guess.

The caveat has to stay attached to the receipt. MoRI does not mean AMD wins every workload. NVIDIA still has the broader CUDA ecosystem, the stronger rack-coherent NVLink/NVL72 story, broader production maturity, and deeper FP4 serving coverage. The real conclusion is workload-specific: AMD can win or become cost-competitive when the workload fits its HBM capacity, ROCm/AITER/MoRI/SGLang stack, and communication pattern. The buyer should not ask "which GPU is best?" The buyer should ask which full stack wins my workload at my latency, quality, energy, and margin target?

Spec table: one read, all four GPUs side by side.

One sentence to read this table by. AMD MI355X matches B300 on HBM3e capacity at the package level and beats B200 there outright, but NVIDIA's rack-scale NVLink coherency packs 9× the GPU count into one programmable domain: a structural advantage for the largest serving deployments that no AMD platform addresses yet. Both vendors get to ~8 TB/s memory bandwidth and ~20–30 PFLOPS FP4 at the chip; the difference compounds at rack scale, and that's where the unit-economics arguments earlier in this post matter most.

Spec table primary sources. AMD MI355X: amd.com Instinct MI355X product page (288 GB HBM3E, 8 TB/s, 256 CUs, 16,384 stream processors, 1024 matrix cores, 2.4 GHz peak, 10.1 PFLOPs dense MXFP4, 1400W TBP, 7× Infinity Fabric links @ 153 GB/s, launched 6/12/2025) + official MI355X product brochure PDF. AMD MI355X 8-OAM platform: amd.com MI355X Platform page (2.3 TB HBM3E platform total, 80.5 PFLOPs MXFP4 platform). NVIDIA Vera CPU: NVIDIA developer blog: Vera CPU for AI Factories (Mar 16, 2026) (88 Olympus cores, 176 SMT threads, 1.2 TB/s LPDDR5X, 3.4 TB/s SCF bisection, 1.5 TB max capacity, Armv9.2) + NVIDIA Vera CPU product page. Vera Rubin NVL72: NVIDIA Vera Rubin platform page (72 Rubin GPUs + 36 Vera CPUs, 6th-gen NVLink). 22,500-sandbox claim: Vera CPU rack hosts up to 256 Vera CPUs delivering "more than 22,500 concurrent reinforcement learning or agent sandbox environments" per the NVIDIA developer blog cited above. GB200 NVL72 specs: nvidia.com GB200 NVL72 page (72 B200 + 36 Grace, 13.5 TB total HBM3e = 72 × 192 GB, 1.4 exaFLOPS FP4 per rack). GB300 NVL72 (288 GB HBM3e per GPU): NVIDIA Blackwell Ultra announcement (GTC March 2025). GB200 NVL72 $3.3M cost figure: SemiAnalysis GB200 Hardware Architecture & BoM (Dylan Patel + Doug O'Laughlin, Jul 2024, original $10K subscriber report).

The precision path is now part of the topology. B200's FP4 story is NVFP4 and Blackwell tensor cores. MI355X's FP4 story is MXFP4/OCP-style microscaling and CDNA4 MFMA-scale. Same "FP4" label, different physical contract. For a buyer, this means FP4 is not a procurement checkbox. It only saves money if the model format, kernel path, serving engine, and hardware instruction actually line up.

The thing to take from this section. Modern inference racks are stacked at five distinguishable scales: chip package (FIG 07.5-A), SM/CU internals (FIG 07.5-E), superchip module (FIG 07.5-B), rack-coherent fabric (FIG 07.5-C / 07.5-F), and the CPU sitting next to it (FIG 07.5-D): and the right scale to think at depends on the question. Cost per token at small batch? Chip-level. Cost per useful task at production scale? Rack-coherent + CPU environment + the SemiAnalysis numbers above. The §02.55 three-generation framing earlier in this post is the abstract version of what these six figures and one spec table show concretely. Now the people building on top of that silicon start to make more sense.

The topology section shows where the CPU sits physically. The next section explains what the CPU actually does.

The CPU did not suddenly become better than the GPU at AI math. That is not the point.

The GPU still owns the dense math path: prefill, attention, expert GEMMs, tensor cores, matrix cores, FP8, FP4, NVFP4, MXFP4, and the hot kernels underneath vLLM, SGLang, TensorRT-LLM, Dynamo, MAX, and every serious serving engine.

The CPU is back because the workload changed. A chatbot is mostly model serving. An agent is model serving plus an operating-system loop: files, tools, bash, tests, linters, sandboxes, package installs, network calls, databases, queues, retries, memory allocation, tokenization, prompt assembly, MCP servers, observability, and routing. That is not tensor-core work. That is CPU work.

The same point now shows up inside vLLM itself. The vLLM CPU path is not a "CPU replaces H100" claim. It is a real enterprise inference lane for smaller models, moderate concurrency, flexible-latency reasoning, private/local workloads, long-context jobs that do not need instant response, embedding/ranking paths, and environments where Xeon servers already exist. The hard part is not only model support. It is NUMA locality, CPU pinning, OpenMP binding, Kubernetes CPU isolation, and KV-cache memory sizing. A serious vLLM-on-CPU receipt should record VLLM_CPU_KVCACHE_SPACE, VLLM_CPU_OMP_THREADS_BIND, CPU_VISIBLE_MEMORY_NODES, physical cores per rank, TP/PP/DP layout, AMX/AVX-512/AVX2 support, TTFT, TPOT, throughput, p95/p99, and cost per request. Otherwise the CPU number is mostly vibes.

The backend work is getting more serious too. The vLLM CPU FP8 attention path for AMX/AVX-512 is the useful signal: CPU serving is moving toward the same memory-pressure playbook as GPU serving, just with AMX tiles, AVX-512 vector paths, oneDNN/DNNL, NUMA placement, OpenMP binding, and cache-space sizing instead of tensor cores and HBM alone. That is why the right framing is not "CPU versus GPU." The right framing is which backend fits this workload path.

GPU = dense model math
CPU = control flow, tools, sandboxes, orchestration, and environment execution
LPU = deterministic low-latency token path
DPU / NIC / storage = movement, security, and I/O
Router = placement decision
Evidence layer = truth about what actually happened

This is why Gen 3 inference optimization cannot stop at the GPU. The expensive thing is no longer one forward pass. The expensive thing is the full task path.

GPU-visible work:
  prefill 8192 tokens -> decode 1024 tokens -> expert GEMMs -> attention kernels

CPU/control work:
  launch server -> route request -> tokenize/parse -> coordinate workers
  collect logs -> collect power/telemetry where available -> write result artifact

agent version:
  add repo files, grep/ripgrep, tests, tools, package installs, retries, summaries

First principles: what a CPU actually is.

A CPU is a general-purpose processor. It is built for messy work: branches, if/else logic, operating-system calls, networking, filesystems, databases, queues, scheduling, serialization, encryption, tool execution, service processes, containers, virtual machines, and subprocesses.

A CPU core is "smart" compared with a GPU lane. It has big caches, branch prediction, out-of-order execution, speculative execution, complex control logic, and strong single-thread performance. That is why CPUs are good at irregular work. A GPU is the opposite tradeoff: many simpler parallel lanes designed to run the same operation across lots of data. That is why GPUs dominate AI math.

CPU:
  few powerful cores
  great at messy control flow

GPU:
  many simpler parallel lanes
  great at massive repeated math

LPU:
  compiler-scheduled inference path
  great when token generation fits deterministic execution

So when people say "the CPU is back," do not read that as "CPU replaces GPU." Read it as: the task expanded outside the model.

What a normal data-center CPU does today.

A normal server CPU, usually AMD EPYC or Intel Xeon, handles the work around the accelerator. It runs the operating system, owns the process scheduler, handles API ingress and networking, launches GPU kernels, feeds data into GPU memory, runs tokenization and prompt assembly, manages inference-server processes, handles logging, metrics, auth, billing, and queues, and coordinates storage, NVMe, memory, and network traffic.

In classic deep learning training, that CPU path mattered, but the main bottleneck was still keeping the GPU fed with dense math. In agentic inference, the CPU path becomes more visible because the model keeps leaving the GPU.

Coding-agent loop:

GPU: decode tool call
CPU: parse JSON
CPU: open repo files
CPU: run grep / ripgrep
CPU: edit file
CPU: run pytest
CPU: capture error
CPU: compress error text
GPU: prefill updated context
GPU: decode patch
CPU: apply patch
CPU: run tests again
GPU: explain result

That is why CPU latency can turn into GPU waste. If the GPU is waiting for a CPU sandbox or tool result, your expensive accelerator is not doing useful model math for that task.

What the CPU has to do for agents.

For agentic AI, the CPU metric is not just "how many cores." The better question is: how much environment work can this CPU sustain per GPU without creating idle GPU time, queueing delay, cache churn, or p99 latency spikes?

Before, CPU sizing was mostly "enough host CPU to keep the GPU busy." Now, for agents, CPU sizing becomes "enough environment CPU to keep the task path moving."

Normal x86 CPU vs Grace vs Vera.

Normal x86 CPU, AMD EPYC or Intel Xeon. This is the standard data-center CPU path. It is flexible, mature, and supported by the whole software world: Linux, containers, VMs, databases, queues, enterprise software, and normal cloud workloads. For AMD MI300X / MI355X systems, the host CPU is usually EPYC. That makes sense: AMD's strength is a standard x86 server ecosystem plus very high-HBM GPUs. The tradeoff is that the CPU and GPU are still more separated than NVIDIA's tightly integrated Grace / Vera + GPU superchip path.

NVIDIA Grace CPU. Grace is NVIDIA's Arm data-center CPU designed to sit close to NVIDIA GPUs. Grace is the CPU for the GPU-serving era: feed the GPU, manage host memory, coordinate data movement, support coherent CPU-GPU memory paths, and reduce the boundary tax between CPU and GPU. Good shorthand: Grace = CPU as accelerator host.

NVIDIA Vera CPU. Vera is different. NVIDIA describes Vera as purpose-built for reinforcement learning and agentic AI, powering code, tools, and data workflows beyond the model. The same NVIDIA page lists 88 custom Olympus cores, Armv9.2 compatibility, FP8 support, LPDDR5X memory bandwidth, and control-heavy software environments at scale. Vera's job is not only to feed the GPU. Vera's job is to run the agent environment around the GPU. Good shorthand: Vera = CPU as agent environment processor.

The important difference: x86 is general. Grace is GPU-adjacent. Vera is agent-environment-native.

Why Vera matters for this argument.

I want to phrase this carefully. We are not saying we discovered the CPU coming back. NVIDIA clearly already sees it. Vera is literally a silicon answer to this workload shift: more code, more tools, more data workflows, more RL environments, more control-heavy software around the model. SemiAnalysis' CPU landscape piece by Gerald Wong and Dylan Patel says the same thing from the market and architecture side: RAG, agents, tool use, databases, and RL loops are increasing the need for general-purpose CPU compute next to accelerators.

So the point is not "Touchdown is smarter than NVIDIA." Obviously not. The point is that the hardware roadmap is confirming the same thing we keep seeing in workloads. The GPU can be extremely fast and the task can still be slow, because the task is waiting on a sandbox, a test runner, a tokenizer, a repo search, a queue, a cache lookup, or a router decision.

A faster decode kernel does not fix a slow test runner. A better FP4 path does not fix bad prompt assembly. A bigger HBM pool does not fix CPU queueing. A good KV cache does not help if the router sends the next turn to the wrong worker. That is the observation we care about: the bottleneck can move outside the GPU without leaving the inference path.

CPU bottlenecks in real agent workloads.

The CPU bottleneck is dangerous because it does not always show up as CPU utilization. It can show up as GPU idle time, cache misses, retry storms, or p99 latency.

USER REQUEST
  ↓
CPU: API ingress → auth → routing → queueing → session lookup → prefix-cache lookup
  ↓
CPU: prompt assembly → tool schemas → repo context → tokenizer
  ↓
GPU: prefill → attention → KV allocation → expert routing
  ↓
GPU: decode plan / tool call
  ↓
CPU: parse tool call → execute bash → edit files → run tests → run lints → capture logs
  ↓
CPU: compress tool result → update task state → decide retry
  ↓
GPU: incremental prefill → decode patch
  ↓
CPU: apply result → write trace → emit evidence packet
  ↓
SUCCESSFUL TASK

Where Groq / LPU / LPX fits.

Groq's LPU belongs in this section, but it should not become the main character. The LPU is not a CPU. The LPU is not a general GPU. Groq describes the LPU as an inference-focused architecture with hundreds of MB of SRAM as primary weight storage, a compiler-controlled architecture, static scheduling, and deterministic execution for predictable performance at scale. (Groq LPU architecture, Groq LPU explainer.)

Cerebras is the other useful visual here because it is so physically different. Cerebras' WSE-3 is a wafer-scale engine with 4 trillion transistors, 900,000 AI cores, and 44 GB of on-chip SRAM. The point is not that a wafer replaces a GPU either. The point is that memory placement is becoming a first-class architecture choice: GPU systems lean on HBM, Groq leans into SRAM and deterministic scheduling, and Cerebras pushes compute and SRAM across an entire wafer.

The practical framing is: CPU = general-purpose environment execution. GPU = large dense model math. LPU / LPX = deterministic low-latency inference path. Cerebras WSE = wafer-scale compute with distributed on-chip SRAM. Groq and Cerebras are not proof that GPUs disappear. They are proof that inference is splitting into specialized lanes.

FIGURE 07.6-B · SPECIALIZED LANES

real hardware, different memory contracts

Cerebras WSE-3

Wafer-scale compute with SRAM distributed across the chip, shown next to a B200-sized package for scale.

NVIDIA Groq 3 LPU

SRAM-heavy, deterministic token-generation lane for LPX, meant to complement Rubin GPUs rather than replace them.

FIGURE 07.6-C · SRAM VS HBM

real source comparison from Hot Chips 2024

The takeaway most token-only benchmarks miss. The slow seam is CPU ↔ GPU, not GPU ↔ HBM. An agentic coding task does many small CPU ↔ GPU round-trips (read file, build prompt, prefill, decode, read result, run test, retry, build next prompt). Each crossing pays the PCIe-or-C2C tax. If 42% of agentic wall-clock is CPU work and 58% is GPU work, the rack-coherent CPU pool sitting next to the GPUs is no longer just a host. It is a co-equal compute tier, and the benchmark that ignores it is mis-measuring the task.

The future inference rack is not one chip, one processor type, or one memory tier. It is a placement system.

Why this changes the benchmark.

Current benchmarks often measure GPU tokens. CPU-heavy agents need task-path benchmarks. The benchmark has to include GPU prefill time, GPU decode time, CPU tool-loop time, sandbox startup, test and lint runtime, tokenization time, prompt assembly time, KV-cache hit or miss, prefix-cache stability, CPU queueing delay, GPU idle time caused by CPU waits, retry count, task success, energy per task, and cost per successful task.

A benchmark that only measures tokens/sec misses the CPU. A benchmark that only measures CPU utilization misses the GPU. A benchmark that only measures cost per token misses the task. The right benchmark records the full path: CPU time, GPU time, cache movement, retries, latency, success, and energy.

This is where we want to contribute. Not by declaring the answer, and not by pretending the stack is simple. We need a standardized, open-source way to measure the CPU/GPU task path so different teams can compare evidence instead of arguing from dashboards. If Vera, EPYC, Xeon, Grace, Rubin GPUs, MI355X, LPUs, serving engines, KV-cache systems, and agent sandboxes are all part of the path, then the measurement format has to follow the whole path too.

The CPU coming back does not mean the GPU era is over. It means inference stopped being only model serving. The GPU still owns the model math. The CPU owns the agent environment. The LPU may own parts of the deterministic low-latency token path. The router decides placement. The evidence layer proves what happened.

That is the Gen 3 machine. Not CPU versus GPU. Not GPU versus LPU. Not NVIDIA versus AMD versus Groq as a religion.

The real question is the only one that matters: what is the smallest, cheapest, fastest, most reliable compute path that completes the task correctly?

That is why Touchdown Labs keeps using cost per successful task as the unit. The task is where CPU, GPU, cache, routing, tools, kernels, precision, and hardware placement all collapse into one number. And that is why we keep coming back to open evidence: if the field is going to optimize this honestly, we need a shared way to measure what happened, replay it, and compare it across hardware without hiding the bottleneck.

The CPU does more than run tools, sandboxes, queues, and observability. It also prepares the text the model sees. In a simple chatbot, that can feel like background work. In a search or RAG product with reranking fanout, it becomes part of the user-visible path.

Perplexity's pplx-unigram writeup is the recent production proof. The workload is reranking: one user query retrieves hundreds of candidate documents, then a smaller RoBERTa/XLM-R-style encoder scores query-document pairs before final answer generation. The encoder can be fast and the query can still wait. Every candidate still has to be normalized, tokenized, batched, and handed to the model. At high fanout, CPU tokenization becomes part of p95/p99 latency and CPU fleet cost.

This is why the megawatt frame cannot stop at GPU utilization. A data center can have expensive accelerators waiting while CPU preprocessing, tokenizer allocation, page walks, or batch construction hold the request. At site scale, those host-side waits reduce successful tasks per megawatt the same way a slow kernel does. The power envelope does not care whether the wasted path was CUDA, Python, Rust, or Unicode normalization.

That distinction is the whole point. Perplexity is not saying CPU work replaces GPU work. They are pointing at a production path where the model forward pass is short enough that everything around it becomes visible: Unicode normalization, Unigram tokenization, Viterbi DP, scratch allocation, batch construction, and handoff into the reranker. The user still experiences one answer. The system may have paid for hundreds of query-document preprocessing jobs first.

The algorithm is not the novel part. Unigram tokenization is a Viterbi dynamic program over byte positions. Vocabulary tokens are edges. Byte positions are graph layers. The encoder picks the highest-scoring segmentation path. The hot loop is the vocabulary trie walk from each byte position.

The implementation is where the cost leaked. Perplexity's article walks through the reference-shaped path: a reverse token_to_ids hash map, a byte-level trie with hash-map children per node, and an inner loop where each prefix match can materialize a string before looking up the token ID. Their measurements show the shape of the leak: thousands of allocations at 512-1K tokens and almost 300K allocations at 16K tokens in the reference path. That is not model intelligence. That is CPU boundary tax.

REFERENCE-SHAPED HOT PATH

for each byte position:
  walk trie
  find matched bytes
  allocate token string
  hash lookup token id
  fetch score
  update Viterbi DP table

cost leaks:
  string allocation
  hash lookup
  pointer chasing
  scattered cache lines
  TLB/page-walk pressure
  fresh per-request buffers

Source-shaped code hook. The exact upstream code is in Perplexity's article and repo; this shortened version shows the important contract without pasting the whole implementation:

// Reference-shaped Unigram trie loop, shortened from the public writeup.
// The expensive part is not Viterbi itself. It is allocation + lookup + pointer chase.

for starts_at in byte_positions(sentence) {
  for matched_bytes in trie.common_prefix_search(sentence[starts_at..]) {
    let token = String::from_utf8(matched_bytes)?;   // per-match materialization
    let id = token_to_ids.get(&token)?;              // side hash-map lookup
    let score = vocab[*id].score;                    // score found after lookup
    best.update(starts_at + matched_bytes.len(), score, *id);
  }
}

Perplexity's fix was state layout. They store token ID and score directly in trie nodes, reuse caller-owned scratch buffers, replace hash-map child nodes with a double-array trie, pack hot fields into cache-line-sized nodes, and back the larger trie with huge pages. Same broad algorithm. Different memory contract.

OPTIMIZED STATE PATH

tokenizer state:
  double-array trie
  token_id + score in trie node
  cache-line-packed hot fields
  huge-page-backed large trie

request state:
  caller-owned scratch buffers
  reused Viterbi / backtrace memory
  zero steady-state heap allocations

receipt:
  same tokenizer config
  exact token parity
  lower encode latency
  lower CPU utilization
  reranker latency moves

The two mechanical changes are small, but they are the whole mechanism. First, the token ID and score move into the trie path, so the encoder does not allocate matched bytes just to rediscover what token it found. Second, the trie stops acting like a heap of small maps and becomes a flat lookup structure. That is the CPU version of making KV pages or GPU tiles layout-aware: represent the state so the hardware can walk it cheaply.

// Optimized shape, shortened from the public writeup.
// Same Viterbi DP. Different state layout.

for pos in 0..n {
  if scratch.prep[pos].is_utf8_continuation() { continue; }
  let s0 = scratch.best_score[pos];
  if s0 == NEG_INF { continue; }

  let mut node = trie.root();
  for end in pos..n {
    let byte = scratch.prep[end];
    node = trie.next(node, byte)?;          // double-array / flat-node transition

    if trie.is_terminal(node) {
      let id = trie.token_id(node);         // already stored on trie node
      let score = trie.score(node);         // no side lookup
      scratch.update(end + 1, s0 + score, pos, id);
    }
  }
}

Perplexity reports roughly 5x p50 latency improvement versus Hugging Face tokenizers, about 2x versus SentencePiece C++, about 1.5x versus IREE C, zero steady-state heap allocations, 5-6x lower production CPU utilization in their inference stack, and double-digit milliseconds shaved from reranker latency. These are Perplexity-reported numbers, not Touchdown measurements.

This is exactly the same lesson as KV cache, only smaller and easier to see. A tokenizer trie is state. A Viterbi table is transient request state. A side hash map is an avoidable lookup boundary. A heap-scattered trie is bad locality. A 50 MB trie on 4 KB pages creates TLB pressure. Once the state is laid out correctly, the same task gets cheaper.

So the rule is not "optimize tokenizers." The rule is: profile the task path. If the model is small, the fanout is high, or the GPU forward pass is already fast, CPU preprocessing may be the bottleneck. The GPU runs the model. The CPU prepares the world the model sees. If that preparation path is slow, the task is slow.

Quantization sounds like a file-size trick. It is not. It is compression with a contract.

Quantization is not just smaller numbers. It is deciding what information the inference system must preserve, where that information lives, and which layer pays to recover it. That is the spine for this section. The dtype is only the visible label. The real system is the representation, the scales, the kernel, the dequant path, the cache behavior, the serving engine, and the task eval that says whether the compressed path still works.

This is the part that feels weird at first: how can four bits preserve quality from sixteen bits? The answer is that model tensors are not random. They have structure. Local ranges. Outliers. Important channels. Low-dimensional signal. Layers that tolerate error and layers that do not. Quantization works when the compressed representation preserves the computation that matters, not every raw number.

So accuracy preservation does not mean "the numbers are identical." It means the next-token distribution stays close enough, the attention ranking is preserved, MoE routers usually pick the same experts, logits and norms do not drift too far in sensitive layers, and the final task still succeeds. If a coding agent gets cheaper but fails more tasks, the precision path did not win. It just moved the cost somewhere harder to see.

This is why the current research is not just "compress harder." The hard part is figuring out which information is actually worth protecting. AWQ is the earlier example on the weight side: do not quantize every weight equally, protect the salient channels that dominate downstream activations. NVFP4 and MXFP4 push that idea into hardware-native math paths, where the scale layout and block size become part of the kernel contract. TurboQuant and SpectralQuant push the same idea into the KV-cache side, where the thing to preserve is not a weight channel but the attention geometry the next token depends on.

So there are really two problems. The compression problem is "how many bytes did we save?" The accuracy-preservation problem is "did we preserve the computation the model actually uses?" The first one is easy to report. The second one is what matters. It is also why a method can look great on reconstruction error and still hurt task quality, or look weird mathematically and still work because it preserves the right ranking, routing, or inner product.

The practical question is not "how low can the dtype go?" The real question is: for this model, this workload, this sequence length, this serving engine, this chip, and this quality target, which precision path gives the best cost per successful task?

This is why the public Kimi K2.5 spine uses the exact words NVFP4, Dynamo vLLM, GB200, 8k/1k, TP/EP layout. NVFP4 is not magic by itself. It only matters because the serving engine, kernels, scale layout, dequant path, expert routing, and hardware fabric all line up enough for the benchmark row to show useful capacity. The precision label is not the proof. The workload replay is the proof.

Interview-reported quantization note: FP4 redundant zero. Makora's SemiAnalysis interview adds a useful hardware-specific example for this audit. The idea is simple: FP4 spends two code points on positive and negative zero, so the redundant zero can be remapped to a special value that better fits LLM tensor distributions. The public RaZeR paper and code show the same family of idea as redundant-zero remapping with fused dequantization experiments. The systems caveat is the important part: on NVIDIA paths, compensating for the special value can add overhead if it becomes a second pass; on AMD-style paths, the interview argues that FP6/MXFP6 hardware support can make this more attractive. Treat that as source-reported research and product direction unless the exact paper, code, model, GPU, and serving trace are linked. It is not a customer savings claim.

Here is the step-by-step version.

1. First, you shrink the numbers. FP16/BF16 use 16 bits. FP8 uses 8. FP4 uses 4. That can cut memory traffic and make more of the model fit closer to the compute.
2. Then you have to recover the range. Four bits cannot represent enough values by itself, so NVFP4 and MXFP4 add scale metadata. The scale tells the kernel how to turn tiny stored values back into useful math.
3. Then the model fights back. Tensors have outliers, sensitive layers, routers, MLP projections, attention paths, norms, logits, and gradients. They do not all tolerate the same error.
4. Then the runtime has to actually hit the fast path. A checkpoint label that says FP4 is not enough. The serving engine has to use native kernels, fused dequant, the right layout, and the right hardware path. Otherwise the system may unpack FP4 into a slower path and lose the point.
5. Then the workload has to prove it. The only proof that matters is the replay: same task, same quality bar, lower p95/p99 latency, fewer GPU-seconds, fewer retries, lower energy, and lower cost per successful task.

So why not just use FP4 everywhere? Because there are three different jobs hiding under the same word. Post-training quantization takes a trained model and compresses it for serving. Quantization-aware distillation or fine-tuning teaches the compressed model to recover behavior from a higher-precision teacher. Pretraining in FP4 asks the model to learn from scratch while the forward pass, backward pass, gradients, optimizer path, and distributed communication all live near the edge of numerical stability. Those are not the same problem.

For inference, FP4 mostly has to preserve the function the model already learned. That is hard, but bounded. For pretraining, FP4 has to preserve the learning process itself. A small rounding bias in a gradient is not just one bad token. It changes the next update. Then the next update starts from a slightly different model. Over trillions of tokens, that can become divergence.

That is why the best FP4 training work looks less like "turn on a dtype" and more like a stability recipe. NVIDIA's NVFP4 pretraining report trained a 12B hybrid Mamba-Transformer on 10T tokens and matched an FP8 baseline closely, but only with a full recipe: selective high-precision layers, Random Hadamard transforms to smooth outliers, 2D block scaling so weights have consistent rowwise and columnwise representations, and stochastic rounding for gradients. The paper's own ablations say each part matters. Transformer Engine says it plainly too: NVFP4 is a more complex recipe than earlier low-precision modes. That matters. If the vendor docs call it a recipe, a buyer should not treat it like a checkbox.

The MXFP4 side tells the same story from a different angle. Pretraining Large Language Models with MXFP4 on Native FP4 Hardware studies Llama 3.1 8B on AMD MI355X and finds that weight-gradient quantization is the primary driver of convergence degradation. Forward propagation and activation gradients are easier. Wgrad is the dangerous path. That is a useful systems fact because it says where the recipe has to spend protection. In that work, deterministic Hadamard rotations restore stability better than the stochastic interventions they tested. In other words: even inside "FP4 training," the answer changes by format, hardware, model, and gradient path.

This is also why a serving team has to be careful with post-training FP4. Diagnosing FP4 Inference isolates NVFP4 and MXFP4 sensitivity across Qwen2.5 0.5B, 7B, and 14B and finds that sensitivity is not uniform. MLP up- and down-projection layers dominate, attention projections are much less sensitive, and sensitive blocks are not always only at the end. Another microscaling PTQ study finds that MXFP8 is generally near-lossless, while MXFP4 W4A4 is still risky; it also finds that scaling-factor error is a real source of damage and that methods built for INT4 can fail when moved naively to MXFP4. That is the direct answer to the "why not FP4 all the time?" question: because the model does not fail uniformly.

For an operator, this becomes a very simple audit. Do not ask only, "does the model support NVFP4 or MXFP4?" Ask where the FP4 path is used, what stays BF16 or FP8, which kernels run, where dequant happens, whether the engine falls back, and what the replay says on your real workload. FP4 is a win only when the whole path stays correct and fast.

How you verify the low-bit path did not break the product. You need three proofs at the same time. Not one benchmark. Not one vibe check. Not one screenshot of lower memory. A quantized model that looks cheaper in isolation but fails the eval suite is not cheaper. A quantized model that improves tokens per second but makes p99 worse is not done. A quantized model that stores FP4 weights but dequants through a slow path is a demo, not an operating win.

1. Quality replay. Run the same real tasks through the BF16/FP8 baseline and the FP4 candidate. Use the eval that matches the product: lm-evaluation-harness or HELM for broad capability checks, OpenAI Evals for workflow evals, RAGAS for retrieval quality and groundedness, SWE-bench or EvalPlus for code, and LongBench v2 or needle-style tests for long context. The point is not to collect trophies. The point is to catch the exact thing your product cannot afford to regress.
2. Latency replay. Run the same input lengths, output lengths, concurrency, streaming mode, cache state, and routing policy. Measure TTFT, ITL / TPOT, end-to-end request latency, throughput, and p95/p99 tails. NVIDIA GenAI-Perf, LLMPerf, SGLang bench_serving, and vLLM bench serve all exist because "tokens/sec" alone is not enough. If users wait longer for the first token, if streaming gets choppy, or if the tail falls apart under real concurrency, the precision path is not production-ready.
3. Runtime replay. Confirm what actually ran. Capture kernel names, precision path, dequant location, HBM use, KV-cache behavior, fallback kernels, GPU power/telemetry when available, retry rate, and task success. GenAI-Perf's analyze path can include p99 latency metrics and GPU telemetry; SGLang and vLLM can emit detailed serving records. This is the layer that catches the fake win: FP4 on disk, FP16 in the hot loop, bad cache behavior, or a fallback path that only shows up at load.

The latency piece matters because throughput can lie. A system can produce more tokens in aggregate while more users miss the latency target. That is why the serving literature has started using SLO attainment and goodput: how many requests finish while meeting the user-facing latency constraint. That is a much better question for quantization than "did tokens/sec go up?"

For a CEO or CFO, this is the plain version: quality is part of cost. If quantization saves GPU memory but increases failed tasks, support tickets, human review, retries, or escalations, the unit economics got worse. For a CTO, the question is whether the change can pass a rollback-safe gate: same workload, same success bar, lower p95/p99 latency or lower cost at the same latency. For an engineer, the question is whether the profiler and trace prove the low-bit path actually executed.

Make it concrete with a Kimi/K2-style MoE example. Earlier in §07.5 we use a Kimi-style architecture shape as the big open-weights MoE example: roughly 1T total parameters, 32B activated per token, 384 experts, MLA attention, 256K context. The activated number matters for compute. The total number still matters for residency, placement, expert paging, and how much memory pressure the serving system has to manage. Raw weights alone already show why quantization changes the system.

This table is deliberately rough: decimal TB, raw parameters only, no KV cache, no activations, no optimizer state, no routing buffers, no tensor-parallel padding, no fragmentation, no scale metadata, and no engine overhead. Real FP4 is not exactly 0.5 bytes per parameter either: NVFP4 adds a scale byte per 16 values plus a tensor scale, and MXFP4 adds a scale byte per 32 values. The point is not the exact deployment plan. The point is the shape: quantization first buys residency and placement, then the kernel and engine decide whether that residency becomes lower cost per successful task.

The simple ladder looks like this:

FP16 / BF16
  safe 16-bit baseline

FP8
  first serious low-precision training / serving acceleration layer

NVFP4 / MXFP4
  4-bit float era
  scale metadata + block layout + hardware kernel path become part of the system

TurboQuant / SpectralQuant
  KV-cache and state compression
  attention behavior matters more than raw tensor reconstruction

FP16 and BF16 are the safe 16-bit baseline. FP16 is the older GPU workhorse: fast, familiar, widely supported, but with less exponent range. BF16 keeps the 16-bit footprint while preserving an FP32-like exponent range, which is why it became the forgiving default for a lot of training and inference paths. NVIDIA's Transformer Engine docs treat BF16/FP16 as the simpler mixed-precision baseline, and then describe FP8, MXFP8, and NVFP4 as lower-precision recipes that require scaling factors to represent tensor ranges correctly. That is the key distinction: below 16-bit, the format is not just a dtype. It is a recipe.

FP8 is the first serious low-precision acceleration layer. On Hopper, NVIDIA pushed FP8 through Transformer Engine and TensorRT-LLM. The usual FP8 formats are E4M3 and E5M2, also standardized in the OCP OFP8 spec. The win is obvious when the workload is GEMM-heavy and the model tolerates the scaling path: less memory, higher throughput, and more model residency on the same hardware. But FP8 is not magic. Bad calibration, unsupported operators, format conversions, or fallback kernels can erase the benefit. That is why performance claims have to name the engine and hardware, not only the dtype.

Block scaling is the trick that makes the 4-bit float era practical. FP4 values are tiny. They only work because the values are stored with scale metadata that gives a local block a usable dynamic range. Smaller blocks usually preserve accuracy better because fewer values share one scale, but they add metadata overhead and more layout pressure. Larger blocks are cheaper, but one outlier can stretch the scale and waste the available precision for everything else in the block. That scale choice is no longer a math footnote. It becomes a kernel and hardware contract.

NVFP4 is the NVIDIA Blackwell-specific FP4 path. NVIDIA's NVFP4 docs describe 4-bit E2M1 values with an FP8 E4M3 block scale shared across 16 consecutive elements, plus a tensor-level FP32 scale. That scale metadata is not decoration; it is part of what the Blackwell FP4 tensor-core kernel must load, align, and use. The upside is a strong NVIDIA-native path when the runtime and kernel are actually using it. The limit is portability: NVFP4 is not automatically the same contract another vendor consumes.

MXFP4 is the OCP microscaling direction. OCP's MXFP4 takes a different standard path: FP4 E2M1 values in 32-value scaling blocks with E8M0 scale data. AMD's MI355X story points directly at that direction: AMD lists OCP-FP8, MXFP6, MXFP4/FP4 in the MI355X precision set, and AMD Quark documents MI355X support for float8 and mxfp4 @ mxfp4 compute. On AMD, the interesting part is that MXFP4 can line up with Quark, Gluon, and CDNA4 MFMA-scale paths instead of living as a model-file label that the kernel has to unpack slowly.

TurboQuant and SpectralQuant are a different branch. NVFP4 and MXFP4 are mostly the math path: weights, activations, GEMMs, expert routing, tensor cores, matrix cores. TurboQuant and SpectralQuant are the state path: KV cache, long-context memory, and attention behavior. The common principle is the same: preserve the information the downstream computation actually uses. But the thing being protected is different. For weights, it may be important channels or expert GEMMs. For KV cache, it is the attention geometry the next token will read from. §16 goes deep on that side so we do not have to repeat it here.

MoRI quantized all-to-all: quantization is not only weights and KV cache.

One more quantization surface matters for MoE serving: communication. In the AMD + SGLang + MoRI result, quantization is used not only for model weights or KV cache. It is used inside expert-parallel all-to-all. Tokens are dispatched to routed experts and then combined back. MoRI uses FP4 dispatch and FP8 combine to reduce communication volume while preserving task quality. That moves quantization from "model compression" into distributed-systems compression.

That matters because MoE inference can be communication-bound even when the GEMM kernels are strong. If every token has to move across GPUs or nodes to reach its experts, the expensive thing is not only math. It is bytes in flight. Reducing those bytes by quantizing dispatch/combine changes the serving economics directly.

communication quantization receipt:
  dispatch_dtype
  combine_dtype
  scale_layout
  routing_quality
  expert_output_error
  bandwidth_saved
  p95_p99_impact
  task_quality_impact

FP4 can fail in very ordinary ways. Bad calibration can move the distribution out from under the scales. Outliers can stretch a block scale until useful values collapse. Sensitive layers can be quantized too aggressively. The dequant path can be separate instead of fused. The serving engine can silently fallback. The kernel may not consume the compressed layout directly. Format conversion can eat the win before useful math begins. This is why "the checkpoint is FP4" is not an answer. The question is what path actually ran.

Now put that in a real workload. Take a long-context MoE coding-agent workload on B200/GB200 or MI355X. The precision path changes how much HBM the weights consume, how much KV pressure the runtime can absorb, whether expert GEMMs hit the right tensor-core or matrix-core instruction, whether the dequant path is fused or separated, whether the router can keep workers fed, whether prefix-cache reuse survives the layout, and whether the serving engine actually uses the intended kernel instead of silently falling back. The same model filename can produce very different cost per successful task depending on that chain.

The business read is pretty direct. For a CEO, the precision path can mean more useful AI per GPU, per rack, or per megawatt if the path works. For a CFO, lower memory footprint and higher throughput can mean fewer GPU-hours, fewer racks, less power, and lower cost per successful task. For a CTO, it means validating format, engine support, kernel path, fallback behavior, accuracy, and p95/p99 latency together. For an engineer, it means inspecting dtype, block size, scale layout, dequant location, kernel backend, profiler traces, and task evals instead of trusting the filename.

This is why evidence matters. A quantized model is not faster because the filename says FP4. It is faster only if the runtime, kernel, scale layout, dequant path, memory movement, and accuracy eval all line up on real hardware. If any one of those is wrong, the lower-bit model can become a more complicated way to be slower, less accurate, or harder to debug.

The lowest-bit model does not automatically win. The model that wins is the one whose precision path, kernel path, cache path, and workload path all line up.

Our take: quantization is one of the clearest examples of the whole post. You cannot optimize the model separately from the engine. You cannot optimize the engine separately from the kernel. You cannot optimize the kernel separately from the scale layout and hardware instruction. And you cannot trust the result without a replayable evidence packet that says what actually ran.

This section maps who is attacking which failure mode. The point is not that every team does the same thing. Different teams keep running into the same loop: generate, measure, verify, replay.

Teams whose work we admire, often on layers right next to ours. We've learned a lot from them, so think of this as part thank-you, part reading list. The common thread, in our reading, is evidence and harness. Mostly, we think these are serious people doing real work.

The list is broader than kernel generation on purpose. Serious groups are attacking adjacent layers of the same problem: compilers, runtimes, serving engines, KV cache, distributed communication, search, benchmarks, and hardware portability. That is the point. The full-stack CPU+GPU inference problem is too wide for one layer to explain by itself.

I'm putting the research labs first because they shape the language before the vendors productize it. The Berkeley systems line is the distributed-systems, serving, search, and communication line. MIT CSAIL is the LM-programming and externalized-state line. GEPA starts on the Berkeley Sky / cross-lab side, then bridges into the DSPy world through the optimizer integration. After that, the companies and open-source teams are easier to read because you can see which layer of the same problem each one is attacking.

I respect what these teams are doing. Seeing full-stack optimization compound is honestly incredible, and we want to congratulate every single founder and researcher here. Everything is moving to heterogeneous compute. With so many different chips, rewriting code every time is impossible. Portability has to be native.

There won't be one winner. Teams will own different layers of specific workloads. The demand for compute is massive, and the group working on this is still tiny.

I want to shout out every single one of them: Anne and Chris at Standard Kernel for their R-axis rubric, Chris and the team at Modular getting Mojo and MAX compiler portability out, Michael at Spectral building the SCALE cross-compiler, the Zyphra team making AWS Neuron topology and Domino-style overlap legible, Steven and Emilio at Wafer keeping the leaderboard transparent with KernelArena, the doubleAI team proving kernel correctness with WarpSpeed, Edward and the Cursor team running multi-agent planners on B200 clusters, the LMSYS / SGLang / AMD MoRI team making MI355X economics visible end to end instead of hiding behind a chip headline, and Ryan, Wei-Chen, and Di at Eigen AI co-optimizing from model down to Blackwell. We are a small group trying to make compute and intelligence cheap for everyone. We will continue.

When things move this fast, concrete proof is the only thing that keeps everyone honest. We need open traces of what these systems actually did. That is what doubleAI did with PAC-verification to prove kernel correctness from invariants instead of trusting cuGraph baselines. It is what Cursor did by publishing all 235 solutions to GitHub, and what Wafer is doing by exposing every pinned task and frozen harness on KernelArena. Having raw compiler logs, metrics, and hardware baselines sitting in the open is how we keep the ecosystem honest, making speedup claims believable across different chips. Even for teams only focusing on NVIDIA configs right now, these lessons are what allow the next wave of compute to happen. What we learn today will spread across whatever hardware and software the future holds. We are a single systems and infra community: our job is to productize groundbreaking architectures so everyday people can affordably use them without money or energy constraints. The next section zooms out from the people to the map of tools they are all reaching for.

Before walking through the four code artifacts in §08.5, pause and lay out the whole map. The kernel-and-compiler ecosystem in 2026 has roughly forty named pieces of software in active use: CUDA, ROCm, Triton, CUTLASS, PyTorch, JAX, TensorRT-LLM, vLLM, SGLang, FlashAttention, FlashMLA, TileLang, ThunderKittens, MAX, SCALE, AITER, cuDNN, cuBLAS, NCCL, RCCL, MIOpen, FlashInfer, Liger, Unsloth, Triton-AMD, Pallas, XLA, MLIR, LLVM, NVCC, HIPCC, ROCm-LLVM, cuTile, Gluon, HipKittens, Dynamo, TGI, LMDeploy, MAX Engine, and a dozen more. That looks like chaos until you see the shape. The shape is six layers stacked on top of three vendors, and every name above slots cleanly into one cell of that grid. Worth walking through, because once the grid is in your head the rest of this post stops being a pile of names and starts being a map.

For an infrastructure buyer, this map is not trivia. It is the dependency chain behind revenue per megawatt. A gigawatt campus does not buy one abstract unit called "AI compute." It buys racks whose useful output depends on framework graphs, compiler lowering, engine kernels, collectives, tokenizer paths, KV layout, and hardware-specific fallbacks. If one layer silently falls back or one vendor path lacks coverage, the site can have the same nameplate power and produce fewer successful tasks.

The clean definition: the AI kernel stack is the chain that turns model code into hardware work, then proves whether the work actually ran on the intended silicon path. For production inference, that chain is behind latency, GPU-hours, cache reuse, power, and cost per successful task. A model can be "on B200" or "on H100" and still miss the native path if the compiler, kernel library, runtime engine, or layout contract falls back.

First: why this landscape exists at all.

The whole kernel ecosystem traces back to one decision NVIDIA made in 2006 and has compounded for 20 years. Worth knowing because every name in §08.5 is downstream of it.

YEAR    EVENT                                                                  WHY IT MATTERS
─────   ──────────────────────────────────────────────────────────────────    ───────────────────────────────────
2006    NVIDIA announces CUDA with the GeForce 8800 / Tesla G80 era.          The programming model becomes
                                                                              public enough for the field to
                                                                              see GPUs as compute targets.
2007    CUDA Toolkit 1.0 ships. C-with-extensions for GPUs.                   GPUs become programmable, not
                                                                              just rasterizers.
2007    cuBLAS, cuFFT released - vendor BLAS / FFT primitives.                Vendor libraries become the
                                                                              de-facto baseline for "fast."
2009    OpenCL 1.0 - cross-vendor open standard.                              The first portability attempt.
                                                                              Loses on toolchain quality.
2012    AlexNet wins ImageNet on two GTX 580s. Deep learning + GPU is born.   Demand explodes; cuDNN follows.
2014    cuDNN 1.0 - deep-learning-specific primitives (conv, pool, RNN).      Frameworks (Caffe, Theano) all
                                                                              wire to cuDNN. Lock-in begins.
2016    PyTorch 0.1 - Soumith Chintala et al. at Meta.                        Eager-mode Python, autograd.
                                                                              Researchers move off TensorFlow.
2017    Volta (V100) - first tensor cores (HMMA). NVCC adds <code>mma.sync</code>.        The hardware MMA instruction is
                                                                              now first-class. Kernels split into
                                                                              "uses tensor cores" vs "doesn't".
2018    Triton begins as Philippe Tillet's project; 2019 paper;               First Pythonic tile DSL that
        OpenAI publicly releases Triton 1.0 in 2021.                          actually generated competitive PTX.
2019    AMD ROCm 2.x + HIP - "CUDA-like" portability layer.                   AMD's first credible response.
                                                                              Still source-port (HIPIFY), not
                                                                              binary compatibility.
2020    Ampere (A100, sm_80) - TF32, async copies, 3rd-gen tensor cores.      The chip everyone in this post
                                                                              still benchmarks on.
2022    Hopper (H100, sm_90) - wgmma, TMA, distributed shared memory,         The chip that made warp
        thread-block clusters.                                                specialization the default
                                                                              kernel structure.
        PyTorch 2.0 announced - torch.compile, Dynamo, Inductor, AOTAutograd. Compilation is suddenly the
                                                                              default expectation for PyTorch.
2022    FlashAttention 1 (Tri Dao + Hazy Research) - IO-aware attention.      Hand-written CUDA beats cuDNN by
                                                                              3× on attention; spawns the
                                                                              "kernel research" subfield.
2023    FlashAttention 2 ships; vLLM and SGLang spin up.                      Production inference engines start
                                                                              shipping bespoke kernels by default.
2024    Triton's PyTorch-Inductor backend goes GA - Inductor emits Triton.    PyTorch's "fast path" is now
                                                                              autogenerated Triton kernels.
        FlashAttention 3 (Hopper-specific, async wgmma + TMA).                The first hand-tuned-by-PhDs
                                                                              kernel that fully exploits Hopper.
2025    Blackwell (B200, sm_100) - tcgen05.mma, Tensor Memory (TMEM),         A new instruction (tcgen05) and a
        2-SM CTA groups, decoupled TMA descriptor format.                     new memory tier (TMEM). Hopper
                                                                              kernels do NOT port over for free.
        FlashMLA (DeepSeek), DeepGEMM, ThunderMLA (Together / Hazy),          The DeepSeek / Hazy / Together
        Together Inference Engine coding-agent benchmarks.                    triangle now reaches the engine:
                                                                              kernel work wired into a real
                                                                              serving benchmark, not only a
                                                                              standalone microbenchmark.
        Modular MAX 25.4 - first cross-vendor production stack                Apache-2.0 kernels that target
        (NVIDIA + AMD + Apple).                                               three vendors from one source.
2026    SOL-ExecBench, WarpSpeed, K-Search, kernel-design-agents, KernelEvolve. The agentic kernel-generation
                                                                              cohort this whole post is about.

The single sentence that ties it together. Every kernel library, every DSL, every compiler in 2026 is some answer to one of two questions opened by that timeline: (a) "how do I get my code to feed the tensor cores efficiently" (the post-Volta question), and (b) "how do I do that without rewriting the kernel for every new chip" (the post-Ampere question). The whole landscape sorts by which question each piece of software is answering.

The grid: six layers stacked on three vendors.

The easiest way to hold the ecosystem is a grid: vertical axis is abstraction (high → low), horizontal axis is vendor (NVIDIA, AMD, cross-vendor). Every name in this post fits into exactly one cell.

                          ┌────────────────────┬───────────────────┬─────────────────────────────┐
                          │ NVIDIA-ONLY        │ AMD-ONLY          │ CROSS-VENDOR / PORTABLE     │
┌─────────────────────────┼────────────────────┼───────────────────┼─────────────────────────────┤
│ L0  PYTHON FRAMEWORKS   │ -                  │ -                 │ PyTorch, JAX, TensorFlow,   │
│     (what users write)  │                    │                   │ Keras                       │
├─────────────────────────┼────────────────────┼───────────────────┼─────────────────────────────┤
│ L1  COMPILER-FROM-EAGER │ TorchScript        │ -                 │ torch.compile (Dynamo →     │
│     (auto-fuse, codegen)│                    │                   │ AOT Autograd → Inductor),   │
│                         │                    │                   │ XLA (JAX/TF), TVM           │
├─────────────────────────┼────────────────────┼───────────────────┼─────────────────────────────┤
│ L2  HIGH-LEVEL TILE DSL │ cuTile             │ -                 │ Triton (OpenAI),            │
│     (write tiles in py) │                    │                   │ TileLang (MS), Pallas (JAX),│
│                         │                    │                   │ Halide                      │
├─────────────────────────┼────────────────────┼───────────────────┼─────────────────────────────┤
│ L3  LOW-LEVEL TILE DSL  │ CuTeDSL Python,    │ Gluon,            │ Mojo (Modular, MLIR-based)  │
│     (cute tiles, R3)    │ ThunderKittens     │ HipKittens        │                             │
├─────────────────────────┼────────────────────┼───────────────────┼─────────────────────────────┤
│ L4  C++ TEMPLATE KERNEL │ CUTLASS 3.x        │ Composable Kernel │ -                           │
│     (named MMA atoms)   │ (CuTe in C++)      │ (CK), AITER       │                             │
├─────────────────────────┼────────────────────┼───────────────────┼─────────────────────────────┤
│ L5  RAW KERNEL SOURCE   │ CUDA C++,          │ HIP C++           │ SCALE (compiles CUDA→AMD)   │
│     (hand-coded)        │ inline PTX         │ + inline AMDGPU   │                             │
├─────────────────────────┼────────────────────┼───────────────────┼─────────────────────────────┤
│ L6  ASSEMBLY / MACHINE  │ PTX, SASS          │ AMDGPU ISA,       │ -                           │
│     (what the chip runs)│                    │ assembled GCN     │                             │
└─────────────────────────┴────────────────────┴───────────────────┴─────────────────────────────┘

There is one more layer that does not fit cleanly in the table: runtime. Standalone kernels are not the whole business. A serving engine owns routing, batching, KV-cache layout, prefix reuse, speculative decoding, prefill/decode split, fallback, and traffic-shaped kernel choice. That is why a fast FlashAttention, FlashInfer, DeepGEMM, or ThunderKittens kernel is only half the answer. The runtime still has to select it for the right request shape and preserve the receipt.

How to read the grid. A PyTorch user lives at L0. The moment they write torch.compile(model), they're at L1: Dynamo traces the Python, AOT Autograd splits forward/backward, Inductor codegen lowers to L2 Triton. If they write a custom Triton kernel by hand, they're at L2. If they pick up ThunderKittens or CuTeDSL Python, they're at L3. If they crack open CUTLASS templates in C++, they're at L4. If they hand-write a .cu file with inline PTX, they're at L5. Almost nobody writes L6 by hand in 2026; it's read with cuobjdump and nvdisasm for profiling, audit, and source/SASS correlation. The interesting thing about the grid is the AMD column. AMD has hardware competitive with Hopper-class NVIDIA silicon, but the AMD column at L2 is empty until very recently (Triton-AMD landed in 2024 but coverage is still partial), the L4 box has Composable Kernel (CK) and AITER but with a fraction of CUTLASS's operator coverage, and there's no AMD-native answer at L1. Cross-vendor portability (the third column) is where TileLang, Mojo, SCALE, and torch.compile-with-AMD-backend all live: and where most of the open research is happening in 2026.

The GPU MODE PTX/SASS review is useful because it treats L6 as an evidence layer, not a lifestyle. The point is not that a product team should become assembly people. The point is that when a claim depends on Hopper WGMMA, Blackwell TCGEN05, async copy, TMA, mbarrier discipline, vectorized global loads, or a no-spill path, someone has to check whether the emitted code actually contains that path. For executives, this is infrastructure diligence. It prevents the team from buying more GPUs, changing vendors, or rewriting a kernel because of a guess. For investors, it is one of the reasons a serious inference optimization company is more than a cost dashboard: the defensible work is tracing the claim through the layers until the machine-level receipt matches the business result.

PyTorch and torch.compile: the path most engineers actually walk.

This deserves its own pass, because ~80% of the people who'll touch a GPU kernel in 2026 will never write CUDA: they'll write PyTorch, call torch.compile, and let the compiler stack pick everything below. Understanding what torch.compile actually does is the difference between treating it as magic and treating it as a tool with knobs.

PyTorch started as a pure eager-mode framework: every Python line dispatches to a CUDA kernel (cuBLAS for matmul, cuDNN for conv, a per-op CUDA kernel for everything else). That's fine for research velocity and terrible for performance: every op is a separate kernel launch, every intermediate is materialized in HBM, the compiler sees nothing. PyTorch 2.0 (December 2022) introduced torch.compile, which built four pieces of new infrastructure on top of eager mode to fix that:

USER CODE
  │
  │  @torch.compile decorator (or torch.compile(fn))
  ▼
┌─────────────────────────────────────────────────────────────────────────┐
│ TORCHDYNAMO  - Python frame evaluation hook. Traces actual Python       │
│                bytecode at runtime, captures FX graph of tensor ops.    │
│                Falls back to eager on unsupported patterns ("graph      │
│                breaks"). Handles arbitrary Python, not just TorchScript │
│                subsets. This is the new piece that made the others      │
│                viable on real research code.                            │
└─────────────────────────────────────────────────────────────────────────┘
  │
  ▼
┌─────────────────────────────────────────────────────────────────────────┐
│ AOT AUTOGRAD  - Decomposes the FX graph into ATen prim ops, then        │
│                 jointly traces forward + backward into a single graph   │
│                 the compiler can fuse across the autograd boundary.     │
└─────────────────────────────────────────────────────────────────────────┘
  │
  ▼
┌─────────────────────────────────────────────────────────────────────────┐
│ TORCHINDUCTOR - The default compiler backend. Takes the prim-decomposed │
│                 graph and lowers to one of:                             │
│                   • Triton kernels  (NVIDIA + AMD GPU, default)         │
│                   • C++ + OpenMP     (CPU)                              │
│                   • cuBLAS / cuDNN   (when calling vendor libs wins)    │
│                 Fuses elementwise + reduction ops aggressively;         │
│                 keeps matmul / conv as library calls by default.        │
└─────────────────────────────────────────────────────────────────────────┘
  │
  ▼
┌─────────────────────────────────────────────────────────────────────────┐
│ TRITON         - OpenAI's Pythonic tile DSL. Inductor emits Triton      │
│                  source, Triton lowers through LLVM to PTX (NVIDIA) or  │
│                  AMDGPU ISA (via Triton-AMD on ROCm). The vast majority │
│                  of "compiled PyTorch" runtime today is Triton-emitted. │
└─────────────────────────────────────────────────────────────────────────┘
  │
  ▼
PTX → SASS  (NVIDIA)        or        AMDGPU ISA → assembled GCN  (AMD)

CFO read. "torch.compile" is the lever that turns the same PyTorch model into 1.5–3× faster code on the same hardware, for free, with one line of code, by handing the graph to a compiler that emits autotuned Triton instead of running ops one at a time. Engineer read. Inductor's default codegen is Triton at L2 of the grid above, which is why "writing a custom Triton kernel" is the next thing PyTorch users reach for when Inductor doesn't fuse what they need: they're operating at the same level Inductor itself targets. Kernel engineer read. The frontier moves to "what does Triton not compile well, and what shape of L3 / L4 / L5 code beats Triton on this specific op family", which is exactly the cohort question §08 is about.

For the JAX side of the world, the same story runs through XLA → HLO → MLIR → LLVM → PTX, with Pallas as the JAX-native L2 tile DSL (analogous to Triton). Google's TPU stack uses the same XLA front-end with a different lowering target. PyTorch's L1 layer is Inductor + Triton; JAX's L1 layer is XLA + Pallas; both are the same idea wearing different clothes.

The kernel libraries: sorted by who maintains them and what job they do.

One more piece of the map, because every named library in this post fits in one of five buckets. Worth knowing because the bucket tells you why the library exists, who maintains it, and how stable the interface is.

BUCKET                  EXAMPLES                                          WHY IT EXISTS / WHO MAINTAINS
─────────────────────   ────────────────────────────────────────────────  ──────────────────────────────────
1. VENDOR REFERENCE     NVIDIA: cuBLAS, cuDNN, cuFFT, NCCL,               Maintained by the chip vendor.
   (closed or partly-       cuSPARSE, TensorRT, TensorRT-LLM,             Closed-source historically;
    closed; the bar       cuTLASS (open Apache 2.0), Dynamo, NIM          CUTLASS is the open exception.
    everyone benchmarks   AMD:    rocBLAS, rocFFT, RCCL, MIOpen,          The "if you beat this, you have
    against)                 AITER, hipBLASLt, Composable Kernel (CK)     a real kernel" baseline.
                          INTEL:  oneDNN, oneMKL                          What <code>cuBLAS</code> gives you on H100
                                                                          is the bar K-Search, WarpSpeed,
                                                                          and KernelEvolve are trying
                                                                          to beat in §08.

2. RESEARCH KERNELS     FlashAttention 1/2/3 (Tri Dao + Hazy)             Built by university labs and
   (open, hand-written  FlashMLA (DeepSeek)                               independent researchers. Usually
    by PhDs, beat the   ThunderMLA (Together / Hazy)                      one author or small team. Set the
    vendor baseline by  DeepGEMM (DeepSeek)                               state-of-the-art on a narrow op
    knowing the chip)   ThunderKittens / HipKittens (Hazy)                family for ~6-18 months, then
                          xFormers (Meta)                                   vendor catches up or the next
                          FlashInfer (UW, paged-attention inference)       generation rewrites the rules.
                          Liger Kernels (LinkedIn, training-focused)       The "FlashAttention beats cuDNN
                          Unsloth (fine-tuning, MIT)                       by 3×" pattern, repeated.
                          DeepEP (DeepSeek, MoE expert parallel)
                          DeepSeek's open-week kernel drops

3. INFERENCE-ENGINE     vLLM (UC Berkeley)                                Built by serving-engine teams.
   KERNELS              SGLang (UCB / LMSys)                              Each engine maintains its own
   (engine-internal,    TensorRT-LLM (NVIDIA)                             custom kernels for the bottlenecks
    not standalone      TGI (HuggingFace)                                 it actually hits - paged attention,
    libraries)          LMDeploy (OpenMMLab)                              MoE routing, speculative decoding
                          Dynamo (NVIDIA)                                   draft heads, etc. Usually Triton or
                          MAX Engine (Modular)                              CUTLASS underneath. Often the
                          Together Inference Engine                         interesting kernel work lives here
                          (ThunderMLA + EAGLE + coding-agent workload)      and never ships as a separate lib.

4. KERNEL DSLs &        Triton (OpenAI, L2)                               Built by compiler researchers and
   COMPILERS            TileLang (Microsoft, L2)                          DSL groups. Each takes a position
   (write once, lower   cuTile (NVIDIA, L2)                               on (a) what abstraction is right,
    to one or more      Pallas (Google / JAX, L2)                         (b) which vendors to target, (c)
    targets)            Gluon (AMD, L3)                                   how aggressive to be on autotuning.
                          CuTeDSL Python (NVIDIA, L3)                       The §08.5 four artifacts are all
                          ThunderKittens (Hazy, L3 CUDA-embedded)           in this bucket.
                          HipKittens (AMD port of TK, L3)
                          Mojo (Modular, L3 cross-vendor via MLIR)
                          SCALE (Spectral, L5 source-fixed retarget)
                          Halide (MIT, image processing roots)
                          TVM (Apache, the OG of this category)

5. TRAINING-STACK       Megatron-LM (NVIDIA)                              Custom training-stack kernels
   KERNELS              DeepSpeed (Microsoft)                             that ship inside training
   (inside the          FSDP / FSDP2 (Meta)                               frameworks rather than as
    training framework, ColossalAI                                        standalone libraries. Heavy on
    not standalone)     PyTorch's distributed kernels                     collectives (NCCL / RCCL),
                          slime (TML)                                       gradient accumulation, optimizer
                          fairseq's kernels (Myle Ott, Meta legacy)         fusion, mixed precision.

The bucket pattern. The vendor libraries (bucket 1) set the bar. The research kernels (bucket 2) beat the bar on narrow op families for 6–18 months until the vendor catches up or the next chip generation resets the game. The inference engines (bucket 3) consume both vendor and research kernels and add their own where the bottleneck is engine-specific. The DSLs and compilers (bucket 4) are the productivity layer that turns kernel-writing from "PhD in CUDA" to "competent Python engineer with hardware intuition." The training-stack kernels (bucket 5) live inside training frameworks and rarely ship as standalone libraries. The cohort in §08 is overwhelmingly working in bucket 4 (the DSLs they emit code into) and bucket 2 (the research-kernel territory their generated code is competing with).

Three readings: one for each audience.

For the CEO / CFO. The kernel stack is a layer cake: your application code at the top, the silicon at the bottom, six layers of compilers and libraries in between. You don't get to skip layers; you only get to choose how many of them your team owns. Buying NVIDIA's full stack (PyTorch on H100 with cuDNN + TensorRT-LLM + Dynamo) gets you 70–80% of the achievable performance with zero kernel headcount. Custom kernel work (research kernels, custom Triton, hand-written CUDA) recovers the last 20–30%, but it requires PhD-grade talent that costs $500K–$2M/year fully loaded and ships measured wins on specific bottlenecks rather than across-the-board lift. The real question is not "should we do kernel work": it's "which 3–5 kernels are 80% of our inference bill, and is it cheaper to optimize them ourselves or to wait 6 months for the open-source / vendor stack to absorb the win." The cohort in §08 is making that wait shorter every quarter.

For the engineer who works in PyTorch. Start at L0, write the model in eager PyTorch. Profile with torch.profiler: find the 3 ops that take 70% of the runtime. Apply torch.compile first; that takes you to L1 (Inductor + Triton) and usually buys 1.5–3× for free. If a specific op is still the bottleneck, check FlashAttention / FlashInfer / Liger / Unsloth first: there's probably already a research kernel for it. Only after exhausting those does it pay to drop to L2 (write your own Triton kernel) or L3 (ThunderKittens / TileLang). The single most cost-effective skill in 2026 for a PyTorch engineer is not "write CUDA": it's "know which of the 40 libraries above to import for the op that's actually slow."

For the kernel engineer. The grid above tells you where the open jobs are. The cross-vendor portable column at L1 and L2 is where the next 18 months of high-leverage work lives. Triton-AMD coverage is still partial. Pallas-on-AMD doesn't exist. TileLang's AMD path is the most mature open option but still narrower than its NVIDIA path. CUTLASS has no AMD equivalent with the same template depth (CK is the closest but a different model). The Blackwell-specific kernels (tcgen05.mma, TMEM, 2-SM CTA groups) are 6 months into the cycle and the open kernel libraries are still catching up; the cohort in §08 is mostly working on B200, and the open source is downstream of them. The other thing the grid tells you: the verification layer underneath all of this is the part that hasn't shipped yet at all. That's the layer the rest of this post is about.

One sentence that holds the whole map. Every kernel library is some answer to "feed the tensor cores on this specific chip"; every compiler and DSL is some answer to "do that without rewriting the kernel for every new chip"; every inference engine is some answer to "schedule those kernels under real traffic without leaving throughput on the floor"; and the agentic cohort in §08 is some answer to "have a machine do all of the above so a small team can keep up with the hardware cadence." Hold those four sentences and the rest of this post stops being a pile of names.

Tile DSLs: the programming model the kernel cohort converges on.

The kernel % gains across the §08 cohort don't come from rewriting CUDA by hand five times across five chips. They come from tile-shaped DSLs compiled down to vendor-specific code by a portable lowering stack. That is the programming-model bridge into the code examples. The three that matter:

• TileLang (Microsoft, originally TileIR): TVM-based Pythonic DSL for GEMM, Dequant GEMM, FlashAttention, LinearAttention. Tested on NVIDIA H100 (Auto TMA/WGMMA), A100, V100, RTX 4090, RTX 3090, RTX A6000 + AMD MI250 (Auto MatrixCore) and MI300X (Async Copy). The headline data point: in April 2025, tile-lang shipped a FlashMLA implementation for AMD MI300X that hit parity with AITER's hand-optimized assembly kernels: same performance, written in Python instead of inline assembly. In March 2025, an 80-line Python MLA decoder hit H100 parity with FlashMLA. That's the multiplier: write once in TileLang, get vendor-grade performance on both NVIDIA and AMD.
• cuTile (NVIDIA, Tile-IR-based): NVIDIA's own tile programming model, currently 2K stars, requires NVIDIA Driver r580+. Blackwell and Ampere/Ada supported today, Hopper coming. pip install cuda-tile. cuTile is NVIDIA's bet that the tile-shaped programming model is the right level of abstraction for the next decade of kernel work.
• AITER (AMD): AMD's high-performance kernel library on ROCm. The reference AMD-vendor implementations the cohort benchmarks against. The bar to beat on AMD CDNA hardware. The tile-lang FlashMLA result above matters precisely because matching AITER's hand-optimized assembly in a portable DSL is the moment the tile abstraction starts paying for itself on the open-ISA side.

The cohort has a lot of names. The code is where the claims become concrete. The portability and observability arguments the rest of the post leans on are doing real work and the code is publicly inspectable. Four short artifacts, one each from Spectral SCALE, Microsoft TileLang, Stanford's ThunderKittens, and Modular's Structured Mojo Kernels. They matter for one reason: each one is what the compile and verify verbs of the kernel-RL loop actually look like, on the candidates an agent is most likely to write.

Why this exists now and not in 2020. The AMDGPU LLVM backend went compute-grade around 2023–2024 (CDNA3 / gfx942) and CDNA4 / gfx950 landed in mid-2025. ROCm 7.x made AMD's CUDA-X equivalents (rocBLAS, AITER) drop-in-replaceable. MI300X and MI355X gave the cross-vendor problem a real buyer. And LLM kernel generation made it urgent. The unlock was not hardware alone. It was the decision to augment CUDA and Mojo rather than replace either: keep the surface developers know, swap the lowering target underneath. That door opened in 2025.

The investor version is simple: cross-vendor code paths decide whether power is trapped behind one supplier, one compiler, or one missing kernel. If a team can replay the same workload on B200, GB200, MI355X, and future ASICs with evidence attached, procurement changes from a brand bet into a measured capacity decision: which path gives the most successful tasks per dollar and per megawatt?

The point of these four artifacts is not "look at four random projects." They are four different answers to the same portability/control problem. Each chooses what stays fixed and what changes underneath.

Artifact 1: Spectral SCALE: the same CUDA file, two completely different machines.

What it is. SCALE is a CUDA-compatibility compiler from Spectral. It takes a CUDA source file (including inline PTX assembly) and compiles it to native code for non-NVIDIA accelerators. Drop-in replacement for nvcc, invoked the same way, just with a different target environment sourced first.

How it works. SCALE parses the CUDA source (including inline asm blocks), lowers it through LLVM IR, and emits native machine code for the target architecture: sm_89 for NVIDIA or gfx950 for AMD MI355X, depending on which environment you activated. Same input file, different lowering target, zero source edits.

Why it's different from the other three. SCALE keeps the source fixed (CUDA, the language agents already know best) and varies only the backend. It's the only artifact below that accepts inline PTX and still ships to AMD. TileLang and Mojo make you rewrite into their own DSL. ThunderKittens stays on NVIDIA only. SCALE is the option for teams that have a CUDA codebase already and need it to run somewhere else without porting it.

old model:   neutral spec → vendor implementation → portable code
SCALE model: CUDA corpus → compiler/runtime implementation → multiple hardware targets
agent model: LLM proposes → compiler diagnoses → runtime executes → harness verifies

The smallest code receipt is normal CUDA. SCALE's basic example is vector add. The important line is not exotic. It is the ordinary CUDA work hierarchy: local thread id plus block id times block size. That is the Part 3 point in code: CUDA describes parallel work, not just a specific NVIDIA chip.

// From SCALE's basic example.
__global__ void basicSum(const int * a, const int * b, size_t n, int * out) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) out[idx] = a[idx] + b[idx];
}

// Same CUDA launch shape; SCALE changes the target underneath.
basicSum<<<N / 256 + 1, 256>>>(devA, devB, N, devOut);

Source: SCALE basic example. The example allocates device memory with cudaMalloc, copies with cudaMemcpy, launches the CUDA kernel, synchronizes, and validates the output.

The stronger receipt is library compatibility. Real AI software is not just one .cu file. It is CUDA runtime APIs, CUDA-X libraries, build systems, link flags, and runtime behavior. SCALE's BLAS example uses the cuBLAS API shape and forwards it to the relevant ROCm path on AMD.

// From SCALE's BLAS example.
CHECK_CUBLAS(cublasDdot(handle, A.size(), devA, strideA, devB, strideB, &result));

// CMake keeps CUDA-shaped linking.
target_link_libraries(example_blas PRIVATE cublas redscale)

Source: SCALE BLAS example. This is why portability is a runtime/library problem, not only a kernel syntax problem.

Real CUDA from Spectral's PTX example. Inline PTX inside a CUDA kernel, which is the kind of code that's supposed to be unportable by definition:

__device__ inline uint32_t ptx_add(uint32_t x, uint32_t y) {
    asm("add.u32 %0, %0, %1;" : "+r"(x) : "r"(y));
    return x;
}

__global__ void kernelAdd(const uint32_t* a, const uint32_t* b,
                          size_t n, uint32_t* out) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) out[idx] = ptx_add(a[idx], b[idx]);
}

Source: docs.scale-lang.com/unstable/examples/ptx/ (verified verbatim from SCALE's ptx.cu example, May 2026). Compiler details: Inline PTX support.

One .cu file. Two compile paths, same nvcc invocation:

# NVIDIA target - PTX → SASS for, e.g., sm_89:
source /opt/scale/bin/scaleenv sm_89
nvcc ptx.cu -o ptx_nvidia

# AMD target - native AMDGPU machine code for gfx950 (MI355X):
source /opt/scale/bin/scaleenv gfx950
nvcc ptx.cu -o ptx_amd          # same nvcc command, same file

Source: docs.scale-lang.com/stable/manual/tutorials/how-to-use/: SCALE docs explicitly show source /opt/scale/bin/scaleenv gfx1201 (AMD) and source /opt/scale/bin/scaleenv sm_89 (NVIDIA) as the canonical activation pattern; gfx950 is the MI355X-specific target ID per SCALE's supported-AMD-architectures list.

Zero source changes. No HIPIFY, no second codebase. Even the inline PTX block goes through: SCALE parses the asm, lowers it to LLVM IR, and emits native AMDGPU machine code. That's what "drop-in nvcc replacement" means in practice. The same compiler is what lets the AMD-side benchmarking from the cohort section above produce real, comparable numbers across NVIDIA and AMD targets.

source file
SCALE compiler version
target architecture: sm_* or gfx*
CUDA runtime APIs used
inline PTX blocks used
library calls resolved
NVIDIA baseline
AMD baseline
generated binary evidence
correctness tolerance
runtime profile

Artifact 2: TileLang: one Python kernel, NVIDIA and AMD by a single flag.

What it is. TileLang is a TVM-based Pythonic DSL for high-performance GPU kernels, originally from Microsoft. You write the kernel in Python using tile-shaped primitives; the TileLang compiler emits vendor-specific CUDA for NVIDIA targets or HIP for AMD targets at JIT time, depending on a single target flag.

How it works. The Python source describes the kernel in tile-level operations: T.alloc_shared for shared-memory tiles, T.alloc_fragment for register tiles, T.gemm for the matrix-multiply primitive, T.Pipelined for software pipelining stages. The compiler maps those primitives to vendor-specific instruction streams (CUTLASS-style CUDA with wgmma + TMA on H100, HIP with mfma + async copy on MI300X) without the developer touching the lowering. Single .py file, JIT-compiled per target.

Why it's different from the other three. TileLang varies both the source surface (Python, not CUDA) and the lowering target through one description. It's the most aggressive of the four on portability; the developer never sees vendor-specific code. Where SCALE keeps the source fixed and changes the backend, TileLang changes both. Where ThunderKittens is NVIDIA-only and Mojo requires Mojo, TileLang lets a Python-fluent kernel author target both vendors without learning either's instruction set.

Real code from tile-ai/tilelang on GitHub: a single GEMM-plus-ReLU kernel that retargets to cuda or hip at JIT time:

import tilelang
import tilelang.language as T

@tilelang.jit                 # target = "cuda" | "hip" | "cpu"
def matmul_relu(A, B,
                block_M: int = 64, block_N: int = 64, block_K: int = 64,
                dtype = T.float16, accum_dtype = T.float32):
    M, N, K = T.const('M, N, K')
    A: T.Tensor[[M, K], dtype]
    B: T.Tensor[[K, N], dtype]
    C = T.empty([M, N], dtype)

    with T.Kernel(T.ceildiv(N, block_N),
                  T.ceildiv(M, block_M),
                  threads = 128) as (bx, by):
        A_shared = T.alloc_shared((block_M, block_K), dtype)
        B_shared = T.alloc_shared((block_K, block_N), dtype)
        C_local  = T.alloc_fragment((block_M, block_N), accum_dtype)
        T.clear(C_local)

        for ko in T.Pipelined(T.ceildiv(K, block_K), num_stages = 3):
            T.copy(A[by * block_M, ko * block_K], A_shared)
            T.copy(B[ko * block_K, bx * block_N], B_shared)
            T.gemm(A_shared, B_shared, C_local)   # cute on NVIDIA, hip on AMD

        for i, j in T.Parallel(block_M, block_N):
            C_local[i, j] = T.max(C_local[i, j], 0)          # fused ReLU
        T.copy(C_local, C[by * block_M, bx * block_N])
    return C

Source: github.com/tile-ai/tilelang README.md: verified verbatim against the matmul_relu "GEMM Example with Annotations" section (the cute/hip on Nvidia/AMD GPUs comment is the project's own wording on the T.gemm primitive). Multi-target dispatch detail: PR #858 made the @tilelang.jit(target="cuda"|"hip"|"cpu") selector explicit in the docs (Sep 2025).

The decisive line is the comment in TileLang's own source: # Currently we dispatch to the cute/hip on Nvidia/AMD GPUs. The Python is identical for both targets. The TileLang compiler emits a CUTLASS-style CUDA path for an H100 with wgmma + TMA and a HIP path for an MI300X with mfma + async copy from the same .py file. One kernel definition, two vendors, zero source changes. That's the kernel-RL compile verb made cross-vendor at the source level: and it means the harness has to be ready for the same source to produce two different binaries that need to be measured honestly on two different chips.

TileLang source
target = cuda / hip / cpu / other
arch target
compiler version
generated CUDA/HIP/CuTe path
chosen tile sizes
pipeline stages
selected MMA primitive
vendor reference baseline
correctness tolerance
Nsight or rocprof trace
PTX/SASS or AMD ISA evidence where available

Artifact 3: ThunderKittens: a tile-based R3 DSL embedded directly in CUDA.

What it is. ThunderKittens is a tile-based DSL from Stanford's Hazy Research lab: Tri Dao and collaborators. Unlike TileLang or Mojo, it isn't a separate language compiled to CUDA. It's a C++ header library you #include directly into a regular .cu file. The kernel is still CUDA; ThunderKittens just gives you first-class tile types and operations sized to what NVIDIA tensor cores actually consume.

How it works. The library exposes named tile types (st_bf (shared-memory bfloat16 tile), rt_fl / rt_bf (register float / bfloat tiles)) sized to tensor-core MMA fragments. It handles the things kernel writers usually hand-tune: producer-consumer staging through shared memory, TMA loads, swizzling for bank-conflict avoidance, warp-group tensor-core dispatch. You write the kernel against the tile API; ThunderKittens generates the right MMA instructions and the right async copies underneath.

Why it's different from the other three. NVIDIA-only by design. ThunderKittens is the opposite portability bet from SCALE / TileLang / Mojo: stay on NVIDIA, but compress the gap between R3 (DSL-level) and R4 (instruction-level) so an agent can get vendor-grade performance without escaping CUDA. It's the narrowest, deepest, most NVIDIA-specific of the four, and the one a kernel-writing agent is most likely to author successfully because the abstraction stays close to the hardware. The reason ThunderKittens matters in this lineup is that it represents the "vertically deep on one vendor" answer, against three "horizontally portable across vendors" answers.