How the NVIDIA Compiler Moat Actually Works
Inside the NVIDIA Compiler Moat: ptxas, SASS, and the 21 Bits Nobody Else Can Write
Intro
I spent, more or less, about a year writing an x86-64 assembler in C. Did Lexer, expression evaluator, instruction encoder, ELF writer, relocations, symbol resolution, eventually a macro preprocessor. The thing I keep coming back to from that year is how little information the assembler needed to know.
x86-64 encoding is somwthing baroque. ModRM and SIB are a whole mess, REX prefixes leak into everything, and the exact same mnemonic can have six encodings depending on operand width and register numbering.
But none of that is semantic. I emit the bytes and the machine figures out the rest in some ways. If I put two dependent instructions back to back, the processor detects the hazard, stalls, and my program suddently becomes slow. It is almost never wrong. The whole apparatus of scoreboarding, register renaming and out-of-order issue exists so that the assembler can behave like a dumb table lookup with a symbol table stapled on.
Then I started reading SASS, and that assumption stopped holding.
On modern NVIDIA hardware, if the compiler gets the scheduling metadata wrong, you do not get a slow kernel, you get directly a wrong answer. The dependency interlock for fixed-latency instructions is not in the hardware.
It is basically in a 21-bit field that ptxas writes into every instruction, and that field is undocumented, unstable across various architectures, and produced by exactly one program on earth.
Just that, and not the C++ dialect, and not cudaMalloc, is where I believe the compiler moat actually lives. What follows is a personal attempt to be precise about which layer is load-bearing and which layers people keep mistaking for it, because the public version of this argument is usually pitched at the wrong altitude.
Five moats wearing one name
“CUDA” is, in various conversations, a giant bag holding at least five separate things, with wildly different levels of defensibility. We have:
1. The source language. __global__, threadIdx, <<<grid, block>>>, a C++ dialect with a few extensions. This is the" “easiest” thing everyone points at and, by design, it is the least defensible layer. Clang has compiled CUDA for more than a decade. The NVPTX backend is just upstream LLVM. libNVVM is a documented public API and NVVM IR is a specified subset of LLVM IR. NVIDIA gave the frontend away years ago, precisely because giving away the frontend costs them nothing.
2. The virtual ISA and the JIT distribution channel. PTX plus the compiler that ships inside libcuda.so. This is a genuine moat, one that I’ve personally studied in the past few months, but it is a distribution moat, not a performance one. It is the reason a fatbinary built in 2019 still runs on Blackwell.
3. ptxas and SASS. This is interesting. We’re talking about a closed optimizing compiler that does register allocation, scheduling, and the encoding of hardware control state, pointing to an instruction set with no public specification and no official assembler.
4. The libraries and the layout algebra. cuBLAS, cuDNN, CUTLASS, CuTe, NCCL, TensorRT-LLM. Thousands of pre-tuned kernel variants plus a template metaprogramming layer that encodes the tensor core’s operand layout requirements. It’s partly open, but the real thing is that it is entirely co-designed with silicon that is not public.
5. The numerics contract and the install base. Now we’re serious. Reduction orders, TF32 defaults, FP8 and NVFP4 scaling recipes, atomics behaviour, the accumulated set of hyperparameter recipes that were tuned on top of all of it. I don’t know why yobody writes about this one, because it’s worth more than you can think.
Layer 1 is defintely gone. Layer 2 is a legal and logistical lock, not a technical one. Layers 3, 4 and 5 are “the moat”, and they defend against completely different attacks.
The vast majority “CUDA alternative” projects pick a layer, do good work, and then basically die on a different layer they did not budget for.
The frontend left the building a long time ago
nvcc is not a compiler. It is a driver script, and you can watch it work:
nvcc --dryrun -arch=sm_90 kernel.cu 2>&1 | grep -v '^#\$ *[A-Z_]*='
What you get back is a pipeline: cudafe++ splits host from device, cicc (the NVVM-based device compiler) turns device code into PTX, the ptxas compiler turns PTX into a cubin, fatbinary staples the cubins and the PTX together, nvlink handles device linking, and your host compiler of choice gently handles all the rest. Six things doing 6 separate tasks.
One of those, cicc, is an LLVM. NVIDIA has been explicit about this for the last few years: NVVM IR is just LLVM IR with a documented set of intrinsics and address space conventions, and libNVVM is a shipped, documented library you can use.
If you don’t want to use it, Clang’s own CUDA support will be happy to emit PTX for you, and it has been extreme quality, at least since Google put it there to build TensorFlow.
So the frontend nowadays is a commodity. That is why Triton, Mojo, tinygrad, Julia etc… can emit PTX, all without hurting NVIDIA’s position. Every one of those projects walks up to the same closed door at the end of the hallway.
I just want to be careful here, because “the frontend is commoditized” is often stated as if it were an indictment of NVIDIA’s openness, but it’s not. Ceding the frontend was correct and (probably) a deliberate choice.
Various frontends just increase the number of paths into PTX, and every path into PTX eventually terminates in a program NVIDIA controls.
The 21-bit field
Here is the part that changed how I think about this.
Kepler version removed hardware dependency checking for fixed-latency instructions and moved the responsibility into the compiler. Then we had Maxwell refining it.
Volta made it permanent by widening the instruction word from 64 bits to 128 bits and embedding the control payload directly in every instruction rather than in separate scheduling words interleaved every few instructions.
The reconstructed layout of that payload, for Volta and later, looks more or less like this:
First. four bits of operand reuse cache hints. Then, a six-bit mask of scoreboard barriers this instruction waits on. Now on, two three-bit indices naming the barriers it sets for variable-latency results. A yield bit that tells the warp scheduler whether to prefer this warp or switch away.
And finally we have four bits of static stall count: the number of cycles the scheduler must wait before issuing the next instruction from this warp.
That last field is the one that matters for real. For fixed-latency instructions, the hardware does not check read-after-write hazards at all. The microarchitecture research on this is mostly unambiguous: if you set the stall counter wrong and the program produces incorrect results, because there is no scoreboard behind you to catch it.
The register status table and the wires from the issue logic to it were purposefully deleted from the design, and the area and energy they would have cost were spent on something else. The compiler is the scoreboard.
On x86, the ISA is a contract about meaning and the microarchitecture is free to reorder underneath it as long as the contract holds. On NVIDIA hardware since Kepler, a meaningful piece of the microarchitecture’s correctness lives in bits the compiler writes.
ptxas is not a tool that targets the machine, because it is a component of the machine that happens to run on your workstation. Those are the obvious consequences.
You cannot open source that boundary without publishing the timing characteristics of every functional unit in every SM variant of every SKU, including the ones you have not announced. And if you publish it, you have frozen it, because now third-party code depends on it and any change basically breaks correctness, not just performance.
The closedness is not primarily a business decision. It is downstream of an architectural decision that was made for area and power reasons around 2012, and the strategic value came along for the ride.
I find that way interesting than the general story, even thoigh it’s more depressing, because it means the moat is not a policy anyone can be lobbied to reverse.
Occupancy is a compiler policy, not a hardware property
The other thing ptxas decides on your behalf, and the one you can actually measure without a disassembler, is how many registers each thread gets.
We alredy know that an H100 SM has a 256 KB register file: 65,536 registers of 32 bits, allocated to warps at a granularity of eight registers per thread. The SM can hold at most 64 resident warps. Those three numbers set the whole occupancy calculation, and the only free variable is the register count ptxas chose.
Take a kernel where ptxas -v reports 168 registers per thread. Each warp consumes 168 x 32 = 5,376 registers, so the SM can hold floor(65536 / 5376) = 12 warps, which is 12 of 64, or 18.75 percent occupancy. Now suppose you pass -maxrregcount=128.
The same warp now costs 4,096 registers, 16 warps fit, occupancy goes to roughly 25 percent, and ptxas pays for the difference by pushing the excess to local memory, which is to say to L1 and then to L2 and then, if you are unlucky and you have a giant working set, to HBM.
Neither number is right on its own. It depends on what the kernel spends its time waiting for. A kernel that waits on memory wants more warps. When one warp stalls on a load, the scheduler switches to another one, and if enough warps are resident, the waiting never shows up in the runtime. Occupancy is what keeps that kernel fed.
On the opposite side, a kernel that lives in the tensor cores usually wants the registers. It already hides its own waiting, by fetching the next tile while the current one is being multiplied. It has a pipeline, so it has nothing to gain from having other warps to switch to, and it will run happily at eighteen percent occupancy.
That tradeoffs are the single most consequential decision in GPU performance work, and they are made by a heuristic inside a binary that I believe is impossible to read, tune, or replace. -maxrregcount and __launch_bounds__ are the two knobs you get, and they are blunt.
I want to highlight this because it is the everyday version of the argument. You don’t need to care about control bits to be affected by the closed backend. If you have ever watched a two-line source change move register count from 168 to 176 and cost you a resident warp, you have already been on the wrong side of this wall.
There is no official SASS assembler. Instead, there’s a disassembler, nvdisasm, and cuobjdump --dump-sass, and the documentation for the instruction set amounts to a table of mnemonics with one-line descriptions. The opcode encodings, the control field layout, and the exact latency table are reverse engineered by a small group of people, generation by generation.
The lineage of that work is short enough to list: asfermi for Fermi, Scott Gray’s maxas for Maxwell, KeplerAs, TuringAs, CuAssembler for Pascal all the way through Ampere, openptxas under the gpuocelot umbrella for Maxwell and Pascal. Every one of these is a heroic, partial, architecture-locked effort that goes stale the moment a new chip make it to the public.
The payoff of that is real, which is the frustrating part in my opinion. Gray’s maxas SGEMM reached figures close to theoretical peak on GM204, ahead of the cuBLAS of that era, purely by scheduling better than ptxas did.
More recently there is paper work applying reinforcement learning to SASS schedules directly (CuAsmRL), reordering instructions inside cubins that ptxas already produced and finding wins. If ptxas were optimal, that entire research direction would return zeros.
I have personally enjoyed one single detail, that I’m gonna show you below:
/*0080*/ LDG.E.128 R8, desc[UR6][R2.64] ; /* [----:B----:R-:W2:-:S01] */
/*0090*/ LDG.E.128 R12, desc[UR6][R4.64] ; /* [----:B----:R-:W3:-:S01] */
/*00a0*/ IMAD.WIDE R2, R0, R7, c[0x0][0x168] ;/* [----:B----:R-:-:-:S02] */
/*00b0*/ HMMA.16816.F32 R20, R8, R12, R20 ; /* [--23:B----:R-:-:-:S04] */
^ ^
| waits on write barriers 2 and 3
| set by the two loads above
stall 4 cycles before the next issue
This is a reconstructed fragment in the shape that cuobjdump --dump-sass prints, not literal output from a specific compilation, but the bracketed control payload follows the documented conventions.
Please take a deep look at the two loads first. Each one sets a write barrier, 2 and 3 respectively, because a global load has variable latency and the compiler cannot know statically when the data lands.
Then the IMAD in between carries no barrier at all (only a stall count) because integer multiply-add is fixed latency and the compiler knows exactly how long it takes. The tensor core instruction waits on both barriers before it issues.
Now notice what is doing the work. The S02 on the IMAD is not an optimization hint. It is the compiler asserting the pipeline depth of the integer unit on this specific SM. Change the SM, keep the number, and the machine reads a register before the previous instruction has written it.
One last thing: NVIDIA Research’s own SASS instrumentation framework, SASSI, was distributed as a closed-source fork of ptxas binaries. Their own researchers, inside the company, shipped a patched binary rather than source. That suggests you something about how the artifact is treated internally.
PTX is a treaty, not an ISA
I often read posts about people describing PTX as “NVIDIA’s assembly language,” which is so wrong in a way that hides the actual mechanism. PTX is a virtual ISA with unbounded virtual registers, no notion of the physical register file, no scheduling, and no control bits.
It is much closer to LLVM IR than to x86 assembly. Writing PTX by hand does not get you anywhere near the metal; it gets you a slightly lower-level input to the same closed optimizer.
DeepSeek’s V3 work is the canonical example, and it is routinely misreported. They partitioned twenty of the H800’s 132 SMs to handle inter-node communication, used warp specialization, and wrote custom PTX to reduce L2 pressure and interference with the compute kernels.
That is obv. excellent engineering under export-control constraints, but is not an escape from CUDA. It is a deeper commitment to CUDA: PTX pins you to NVIDIA harder than CUDA C++ does, simply because CUDA C++ at least has clean-room reimplementations and PTX has one consumer.
What PTX actually buys NVIDIA is the forward compatibility story, and it is a genuinely excellent piece of platform engineering. Ship PTX inside your fatbinary, and when your binary lands on hardware it has never seen, the driver compiles it at load time. libcuda.so contains a compiler. Every NVIDIA GPU in the world is shipped with a copy of the last-mile toolchain baked into the driver.
The strategic consequences of that are larger than the technical ones:
Any stack that emits PTX, no matter how open, has a closed NVIDIA compiler in its runtime path. Triton is open source and ships PTX to
ptxas. tinygrad inPTX=1mode does the same. Open at the top, closed at the bottom.Compatibility is a one-way ratchet. Old PTX runs on new hardware. New PTX never runs on old hardware. Every generation, the set of instructions you need to hit peak throughput moves into a
.target sm_XXagate that the previous generation’s driver cannot parse.The world’s binaries contain PTX. That corpus is a compatibility obligation NVIDIA has taken on, and it is also an asset nobody else can serve.
The compatibility promise does not cover the instructions that matter
This is the part of the PTX story I have not seen written down anywhere, but I do believe it changes the conclusion.
PTX forward compatibility applies to base architecture targets. sm_90, sm_100, and so on. But the instructions that reach peak throughput do not live there. wgmma requires sm_90a. The tcgen05 family requires sm_100a or sm_103a.
That “a suffix” means architecture-specific, and code compiled for an a target is explicitly not forward compatible: it runs on that architecture and nowhere else, ever. NVIDIA later added f family targets to soften this within a family, which is an admission that the problem is real.
Put the two facts next to each other. The forward compatibility that everyone cites as CUDA’s great platform virtue covers the instruction set you use when you do not care about performance. The moment you write a kernel that actually saturates the tensor cores, you have opted out of it, and you are shipping architecture-locked binaries exactly like everyone else.
So the JIT compatibility story is not a technical moat at all; it is a convenience for the long tail, and the long tail is not where the money is. What it does buy NVIDIA is that the world’s software ships PTX, and PTX has one consumer.
And then there is the legal layer, which people mention and then move past, maybe too quickly. The CUDA EULA has, since 2021 online and since the 11.6 installed files, carried a restriction on reverse engineering, decompiling or disassembling the output generated using SDK elements for the purpose of translating those artifacts to target a non-NVIDIA platform.
Read it as written: it is a restriction on what you may do with the output of the compiler, not just the compiler. Whatever its enforceability, and I am not a lawyer, its practical effect is a chilling one on exactly the class of project that would otherwise attack layer 2.
The libraries are an admission, not a flex
The standard framing is that cuBLAS and cuDNN are evidence of NVIDIA’s software superiority. I want you to read them the opposite way, like I do.
If nvcc and ptxas could reliably generate peak code from readable source, cuBLAS would be a thin wrapper over a few templated kernels. Instead it is a multi-hundred-megabyte binary containing thousands of pre-compiled kernel variants, selected at runtime by heuristics, many of them tuned by hand at the SASS level by people who work at NVIDIA.
cuDNN is even worse. The size of those .so files is a direct measurement of the gap between what the compiler can do and what the hardware can do.
That gap is the real product. NVIDIA is not selling you a compiler that makes your code fast. It is selling you the output of a large, expensive, permanently-employed team of kernel engineers, wrapped in an API, in a form you cannot fork.
And the gap is widening, on purpose
The tensor core programming model has changed shape three times in six years, and each change moved way further away from anything a general compiler can target from ordinary source.
Ampere gave us mma, warp-level, operands in registers, with layout requirements you could hold in your head if you tried. Hopper introduced wgmma, warpgroup-scoped and asynchronous, plus TMA for bulk asynchronous copies driven by descriptors, plus mbarrier for the synchronization that async now requires.
Blackwell deprecated wgmma entirely and introduced the tcgen05 family: a dedicated on-chip Tensor Memory for accumulators, a single thread issuing the MMA on behalf of everyone, and cta_group::2 letting two CTAs on a TPC cooperate on one MMA with shared operands.
Read that again as a compiler person. The CuTe MMA atoms for Blackwell use a thread ID layout of one, where Hopper used 128. At peak performance, the SIMT model is a fiction.
What is actually happening is that one elected lane fires a descriptor at an accelerator, and the rest of the code is a software pipeline built out of barriers, async copies and a memory space whose allocation you manage explicitly.
This is why the ceiling moved from the compiler to the layout algebra. You will see some numbers, because the abstraction hides how physical this is.
Tensor Memory on Blackwell is 256 KB per SM, organised as 128 lanes of 512 columns of 32 bits. You allocate it explicitly, in columns, in power-of-two units with a minimum of 32.
A single warp can only reach the 32 lanes matching its position in the warpgroup, which is why the copy instruction has multicast modes to duplicate a tile across all four lane quadrants.
Block-scaled narrow formats add their own geometry on top: MXFP8 carries one E8M0 scale per block of 32 elements, NVFP4 carries an E4M3 scale per block of 16 plus a second-level FP32 tensor scale, and the scale factors themselves have to be staged into tensor memory in a specific layout before the MMA will consume them.
None of that is expressible as “the compiler will figure it out”. It is a data layout problem with hardware-imposed alignment constraints that propagate back into your choice of tile shape, which propagates back into your choice of pipeline depth, which propagates back into your register budget.
This is why the ceiling moved from the compiler to the layout algebra. You cannot express a competitive Blackwell GEMM by writing loops and hoping.
You express it as a set of layouts, and the algebra of those layouts is CuTe, and CuTe is co-designed with the silicon by the people who designed the silicon, eighteen months before you can buy it.
The single most clarifying data point I have seen on this: FlashAttention 4, the first version written in CuTe DSL rather than C++ templates, reportedly hits about 1613 TFLOP/s on B200 for bf16 head-dim 128 causal attention at 8K sequence length, roughly 71 percent of peak, about 1.3 times faster than cuDNN 9.13, and around 2.7 times faster than Triton on the same hardware.
Three conclusions fall out of one benchmark. NVIDIA’s own closed, hand-tuned attention library was not the ceiling.
The tile-level DSL, not the C++ template library, is now the productive authoring layer. And Triton, the great portable hope, is nowhere near peak on the newest tensor cores, because its abstractions do not yet expose what tcgen05 and TMA require.
What is actually in a cubin
I want to discuss and be clear about one thing worth opening up, because it is the part that felt most familiar to me coming from ELF.
A cubin is an ELF64 object. S ame header layout you already know, with a CUDA-specific machine type. Inside, you get a .text section per kernel holding the SASS, a .nv.constant0 section per kernel holding the kernel parameters and grid constants (this is why parameters arrive via c[0x0][...] in the disassembly rather than in registers), .nv.info sections carrying metadata like register counts and parameter layout, plus the usual relocation sections.
The fatbinary that nvcc embeds in your host object is itself a container of these, in a .nv_fatbin section, with a small descriptor segment telling the runtime what is inside.
You can walk all of it here:
cuobjdump -elf a.out # fatbin contents, section by section
nvdisasm -elf kernel.cubin # ELF structure plus disassembly
readelf -S kernel.cubin # it really is just ELF
The greatest absent, once you have looked, is the same one as before. Every part of this container is inspectable: the relocations and the metadata are readable, the instruction stream is disassemblable. What you can’t do is produce a valid one from modified SASS, because the only program that writes the .text section is the one you do not have.
Coming from x86, that asymmetry is the strangest part of the whole stack. I can hand-assemble an ELF object, link it, and run it. Here I can read everything and write nothing.
What is not being told
Most of the commentary stops at “CUDA is sticky.” Here are the parts I think are underweighted, roughly in order of how much they change the conclusion.
The compiler is correctness-critical, not just performance-critical
Again, sorry for that, but I have not seen this stated plainly anywhere outside the microarchitecture literature. Every discussion of open SASS toolchains treats the problem as “we would generate slower code”.
The actual problem is that a third-party SASS assembler that mis-models the latency of one functional unit on one SKU generates silently incorrect results, in a data-dependent way, on a subset of hardware.
That is a categorically harder engineering problem than performance parity, and it is the reason no serious commercial effort has attempted it. The risk profile is closer to writing a memory controller than to writing a backend.
The moat is a labor market
Strip away the branding and ask what actually cannot be replicated and you’ll find out that is not the compiler source.
It is the few hundred people worldwide who can look at a cuobjdump listing, reason about the control bits, and know from experience that a particular shared memory access pattern will hit a bank conflict on this generation but not the last.
A large fraction of them work at NVIDIA. Most of the rest work at four or five labs and two or three GPU startups. That reframing matters because it tells you what the actual counter-technology is. Many have attemped to address with it, but no, it’s not a new language.
It is just search: autotuners, superoptimizers, and increasingly LLM-driven kernel generation, all of which convert scarce expert labour into abundant compute. If the moat is a hiring problem, then the thing that erodes it is the thing that makes kernel expertise reproducible.
There is a telemetry loop nobody else has
Nsight, DCGM, the driver, customer escalations from every large training run on the planet. NVIDIA sees the performance pathologies of the entire industry’s kernels, continuously, at a scale no competitor approaches.
Then it fixes the top ones in the next library release and, where the fix has to be in hardware, in the next architecture. Every “CUDA maturity” argument is really an argument about the number of laps this loop has run.
The collectives are compiled kernels, and they eat your SMs
I listed the interconnect under layer 5 and then almost walked past it, which would have been a mistake, because NCCL is where the compiler moat and the network moat turn out to be exactly the same thing.
A collective is not a library call that hands work to a NIC, it’s a set of CUDA kernels. NCCL builds rings and trees over the discovered topology, splits each collective across some number of channels, and every channel is a resident thread block.
Which means bandwidth on the wire is bought with SMs that are then not available for compute. It also means the protocol matters: NCCL picks between a low latency mode that trades payload efficiency for skipping memory fences, a 128 byte variant that recovers most of the bandwidth but wants NVLink underneath, and a plain mode that gets full bandwidth at higher latency.
The choice is made by heuristics against measured topology, and the tuning knobs are environment variables. We see probably at least two consequences that are worth holding onto.
First, this is exactly what DeepSeek was doing when they carved twenty SMs out of 132 for communication. Their work was quite revolutionary because they were taking manual control of a tradeoff that NCCL normally makes for you, because on export-limited H800s with halved inter-GPU bandwidth, NCCL’s defaults were wrong for their shape. All of that, as stated before, working with CUDA, not around it or without it.
Second, in-network reduction changes the arithmetic entirely. When the switch itself can perform the reduction, the collective stops being a kernel that moves the same bytes several times and becomes a kernel that ships bytes once.
That capability lives in the NVSwitch and InfiniBand silicon, it is exposed through the same library, and there is no version of “port your kernels” that gets you it. AMD’s answer, RCCL, is a real port of the same design, and the vLLM tuning guidance for it involves pushing the channel count up sharply on MI300X, which tells you how much of the performance is in the topology-specific constants rather than the algorithm.
If you are keeping a scorecard of which moats are compiler-shaped, this one is a hybrid, and it is the least portable thing in the stack.
Port a model from cuBLAS to hipBLASLt and the arithmetic is not bit-identical for a few reasons. Split-k strategies differ, even reduction orders do not behave the same.
TF32 is on by default in one place and not another.
FP8 scaling granularity, per-tensor versus per-block versus the MX and NVFP4 formats, is a semantic choice baked into the kernels.
For inference this usually shows up as a small quality wobble that people can get along with. For training, it shows up as a recipe that diverges at step 40,000 for some random reasons nobody can bisect, on a piece of hardware where you have less tooling.
The cost of that risk almost never appears in the total-cost-of-ownership spreadsheets that compare dollars per FLOP, and it is a real reason large labs pay the NVIDIA premium.
Watch which artifact they do not publish
There is a pattern in NVIDIA’s open sourcing, and once you see it you can predict the next move. They publish the specification and the dialect, while keeping the lowering.
NVVM IR: specified, with a public library. cicc: closed. PTX: exhaustively documented, hundreds of pages, updated every release. ptxas: closed. CUTLASS: open, permissively licensed, genuinely excellent.
The SASS it lowers to: closed. And now CUDA Tile IR: the MLIR dialect, the Python bindings, the bytecode format and the conformance suite are on GitHub. The compiler that turns tile bytecode into cubin is not.
This is a coherent, repeated strategy. Publishing the interface grows the number of producers. Keeping the lowering means every producer terminates in your binary.
It is the same move IBM made with channel architecture and the same move ARM makes with the ISA versus the implementations, and NVIDIA runs it more cleanly than either.
AMD is the counterexample that breaks the simple story
Here is the fact that should discipline every “closed compiler is the moat” argument: AMD publishes its ISA documents. The AMDGPU backend is upstream LLVM. ROCm is open source, top to bottom, including the compiler and the runtime.
You can read the machine encodings for CDNA3 and CDNA4 in a PDF on AMD’s website But (unexpectedly) even this has not collapsed the moat.
So openness is not the mechanism. If it were, AMD would have won already a few years ago. The mechanism is the co-design loop, the library labour, the numerics contract, the install base and the interconnect, and NVIDIA’s closed compiler is a symptom of the same architectural strategy that produces the rest of it.
Anyone whose plan is “make an open CUDA” has already been run as an experiment, by a company with $25 billion of revenue, for at least a decade.
The moat is thinning at the top and thickening at the bottom
Two things are true at once and they point in opposite directions, which is again interesting because it sounds counterintuitive and understandable at the same time.
At the top, the default path from PyTorch to the GPU is now a portable tile language. torch.compile emits Triton, not CUDA C++. Liger, FlexAttention and a large fraction of production fused kernels are…. yes, again, Triton. That layer is vendor-neutral by construction, and it is where most new kernel code is being written.
At the bottom, the instructions you need to reach peak on the newest silicon are less reachable from a generic compiler than at any point since 2007. TMA descriptors, tensor memory allocation, CTA-pair MMA, block-scaled narrow formats with alignment constraints that constrain your tile shapes.
The gap between “compiles and runs” and “hits peak” is the widest it has ever been in 19 years.
Now the interesting question is not whether the moat holds. It is more which layer the industry ends up standardizing on, and NVIDIA has just made a very large move to answer that question in its own favour.
Tile IR: the moat gets rebuilt one level up
The most important compiler event of the last twelve months got almost no coverage outside compiler circles and a restricted group of GPU programmers.
CUDA 13.1 introduced CUDA Tile, a tile-based programming model with its own intermediate representation. Tile IR is an MLIR dialect. The dialect, the Python bindings, the bytecode serializer and a conformance test suite are open source at NVIDIA/cuda-tile.
The bytecode format is explicitly specified as stable, versioned, and forward and backward compatible: bytecode emitted by an older compiler is readable by a newer compiler or driver, and a driver accepts bytecode up to its supported version.
That paragraph should sound familiar, because it is the PTX contract, restated one abstraction level higher, in MLIR, with a published binary encoding. Tile IR is the new PTX.
The user-facing language is cuTile, a Python DSL, with a Julia binding already shipping at near parity on simple kernels. And critically, NVIDIA wrote a Triton backend that targets Tile IR, so Triton programs can be compiled through the new path instead of through PTX.
Meta has filed an RFC to add a cuTile backend to PyTorch Inductor, with the stated motivation that Tile IR gives bytecode portability, performance portability across GPU generations, and access to a tile-specific optimizing compiler, and that it is a prerequisite for pointing Helion at Tile IR as well.
Pro tip: read the sequence of moves rather than the individual announcements.
Triton became the default codegen target of PyTorch. That made the tile abstraction, the layer where new kernels get written, and that layer is portable across vendors. It was the most serious structural threat to the moat in fifteen years, because it made the frontend irrelevant and gave AMD and Intel a credible path to the same source.
NVIDIA’s response was not to fight Triton: it was to build a better tile IR, open source the dialect and the spec, write the Triton backend themselves, and offer everyone a faster path through their own infrastructure. If it works, the portable tile layer that was supposed to route around NVIDIA becomes a frontend for NVIDIA’s IR, and the optimizing compiler underneath, the one that turns tiles into cubin, is closed, just like ptxas.
It is the cleanest execution of the “publish the interface, keep the lowering” strategy I have seen so far, and it happened while everyone was arguing about whether ROCm had caught up.
I want to flag my uncertainty honestly: it is early. Tile IR shipped in December 2025, the Inductor RFC is in progress, and I have not seen independent third-party benchmarks of cuTile against hand-written CuTe on Blackwell at production shapes.
NVIDIA’s own claim is that cuTile is competitive with cuBLAS for large GEMMs and cuDNN for large attention. Competitive with cuDNN is a lower bar than it used to be, given FA4.
The alternatives, honestly
Every project below is real engineering by capable people. What I want to do is say which layer each one attacks and where it runs out of road, because most of them are described by their advocates as if they attack all five.
HIP and ROCm
Attacks: layer 1 and layer 4. A CUDA-shaped API, a source translation tool, and a library set with matching names.
ROCm in 2026 is a genuinely different product from ROCm in 2023, and people who formed their opinion during the Frontier era are working from stale data.
ROCm 7.2 ships tuned hipBLASLt GEMM kernels for FP8, BF16 and FP16 on MI350 and MI355X. AITER provides fused kernels that, inside vLLM, deliver large multiples over the legacy ROCm attention path. vLLM now carries seven attention backends on ROCm.
MI355X ships 288GB of HBM3E, and capacity per package is a real architectural advantage that removes tensor parallelism from model sizes where NVIDIA needs it.
Where it runs out of road: the API-clone strategy structurally guarantees you arrive second.
hipifyignores inline PTX, which means exactly the kernels that were worth hand-optimizing are the ones that do not port. Matrix core layouts (MFMA) differ enough frommma,wgmmaandtcgen05that a “ported” kernel is a rewrite at the layout level. And the performance story has historically been per-model tuning: a specific configuration of a specific model at a specific batch size is fast because someone at AMD made it fast, which is a different asset from a compiler that makes your unseen kernel fast.Honest read: on standard transformer inference, the gap is small enough that the decision is now economic rather than technical. On training, on novel architectures, and on anything requiring bespoke kernels, it is not close.
SCALE (Spectral Compute)
Attacks: layer 1 and layer 3 simultaneously, which makes it the most technically interesting project in this space.
SCALE is a clean-room, Clang and LLVM based drop-in replacement for nvcc that compiles nvcc-dialect CUDA source, including inline PTX, directly to AMD machine code. This is not translation, not even middleware: an actual second implementation of the CUDA language.
The company is small, London based, founded 2018, and funded its own compiler work off consulting. Published benchmarks claim roughly 5.94 times over HIP on AMD hardware, and as of May 2026 they are claiming speedups over nvcc itself on B300 with CUDA 13.
The detail I keep turning over is that they joined NVIDIA Inception in June 2026. A company whose product is CUDA portability, partnering with NVIDIA. Which makes sense if you believe NVIDIA’s interest is in CUDA remaining the standard more than in CUDA remaining exclusive, but it is a strange shape in my view.
Where it runs out of road: the CUDA-X library surface. There are hundreds of libraries above the language, and a compiler that perfectly compiles the language still leaves you needing cuDNN, cuTENSOR, cuDF, NCCL and the rest. Spectral is working on delegating those to ROCm equivalents, which means SCALE inherits ROCm’s library ceiling at exactly the point where it matters most.
ZLUDA
Attacks: layer 2. Binary and PTX level translation, no recompilation.
ZLUDA is the project people cite most and understand least. It gained ROCm 7 support in December 2025 and shipped v6 in June 2026.
That release also announced that commercial funding had ended again and the project is back to being one person’s weekend work, which is why v6’s headline features are 32-bit PhysX and Blender textures rather than anything relevant to inference.
Where it runs out of road: everywhere, and the reasons are structural rather than technical. Binary translation is a treadmill against a vendor who ships a new virtual ISA extension every generation and a whole new IR every few years. The EULA restriction points directly at this approach. It has now lost funding twice, from two different backers. I would treat the binary translation path as closed, not because the engineering is bad, but because the economics have been tested twice and failed twice.
Triton
Attacks: layer 1 and, at least partially, layer 4.
Triton is the most consequential thing on this list because of where it sits: it is what torch.compile emits, which means it is the default kernel authoring layer for most of the industry whether they know it or not. It runs on NVIDIA, AMD and Intel from the same source.
Where it runs out of road: two places. On NVIDIA it emits PTX, so it never escapes
ptxas; it relocates the frontend, it does not cross the wall. And on the newest hardware it is a long way from peak, because the abstractions do not expose TMEM, CTA-pair MMA and the descriptor-driven async machinery that Blackwell needs. FlashAttention 4 leaving Triton for CuTe DSL is the load-bearing evidence there.
And now NVIDIA is offering to compile it through Tile IR, which resolves the second problem in exchange for reintroducing the first at a higher altitude.
Helion, TileLang, Pallas, and the search-based approach
Attacks: the labour market, which as argued above is the real moat.
Helion is a Python DSL from Meta’s PyTorch compiler team that sits above Triton: you write PyTorch-shaped code with tile loops and it autotunes over hundreds of generated Triton implementations, spending roughly ten minutes of compute per kernel to do it.
The reported numbers are the interesting part. Around 1.05 times over torch.compile and 1.44 times over hand-written Triton on AMD on average, with individual kernels much higher, and on H100 roughly 1.2 to 1.85 times over Triton.
Beating hand-written Triton on average, by searching. That is the anti-moat technology, and it is not a language feature. It is a compute substituting for expertise, and it gets better every time GPUs get cheaper, which is a nice recursion.
The same logic explains AMD’s agentic kernel work (GEAK), LLM kernel generators, and the growing body of RL-over-schedules research including the SASS-level work mentioned earlier. If your competitor’s advantage is several hundred irreplaceable engineers, the correct response is not to hire several hundred engineers. It is to make the engineers less necessary.
Where it runs out of road: search needs a correct, expressive, fast-to-evaluate space to search over, and on Blackwell the space that contains peak performance is only reachable through the vendor’s layout algebra. You can search brilliantly within Triton’s expressible set and still be 2.7 times off.
tinygrad
Attacks: layers 1, 2 and 3, by refusing the premise.
tinygrad generates its own kernels, can render directly to PTX, and has experimental backends that talk to AMD and NVIDIA hardware without the vendor userspace runtime at all.
Whatever you think of the framework’s ambitions, the AMD path is an existence proof that the vendor compiler and userspace stack are not physically required.
Where it runs out of road: peak tensor core utilization on the newest silicon, and the entire ecosystem surface. It is a demonstration of what is possible for a small team, not a production alternative for a lab training frontier models. Its real contribution is epistemic: it proves the stack is thinner than the marketing implies.
Mojo and MAX
Attacks: layers 1 and 4, with an MLIR-native language designed for kernel authoring across vendors.
Modular’s own writing on this is refreshingly unromantic: they note that write-once-run-everywhere in GPU land has usually meant write-once-run-slowly-everywhere, that CUTLASS does not attempt portability beyond NVIDIA and is often locked within a generation, and that Triton’s performance degrades off NVIDIA.
Their bet is that a properly designed compile-time metaprogramming language can express hardware-specific structure without forking the source.
Where it runs out of road: it is a proprietary language asking developers to adopt it on faith, competing against a Python DSL layer that PyTorch already emits by default. The technology is pretty good. The distribution problem is brutal.
The vertical integrators: TPU, Trainium, and friends
Attacks: the entire column, by not participating in the compatibility game at all.
Google’s TPU plus XLA is the only stack that has demonstrably trained frontier models at scale outside NVIDIA, sustained, for years. It did not do that by cloning CUDA.
It did it by owning the chip, the interconnect, the compiler, one framework, and one enormous first-party customer whose workloads shape the hardware roadmap. Trainium plus NKI is the same play at an earlier stage.
The lesson is uncomfortable for everyone selling portability: the proven counter to vertical integration is vertical integration. Compatibility layers have a zero-for-many record. Owning the column has a one-for-one record.
Where it runs out of road: you cannot buy this. It requires a decade, a captive workload, and a willingness to eat several generations of worse hardware.
The ones I am not covering properly, and why
Three more categories exist and deserve at least an honest pointer rather than silence.
Huawei’s CANN with the Ascend NPUs is the most complete non-Western vertical stack I’ve ever seen, with its own kernel language, its own graph compiler, and a captive domestic customer base that gives it the one thing compatibility layers never have, which is a workload that shapes the roadmap.
I do not have enough independent benchmark data to say anything useful about where it sits on peak utilization, and I would rather say that than guess randomly. The same caveat applies to Moore Threads and Biren, both of which ship CUDA-adjacent programming models.
On the standards side, chipStar compiles HIP and CUDA to SPIR-V for OpenCL and Level Zero targets, and Intel’s SYCLomatic migrates CUDA source to SYCL with a reported majority of code converting automatically and a manual remainder. Both are layer 1 plays with the layer 4 problem intact.
And there is a fourth category I have deliberately excluded from the comparison table because it is not a competitor at all: NVIDIA’s own Python surface. cuda.core, Numba, Warp, nvmath, and now cuTile are a coordinated effort to make sure that when kernel authoring moves fully into Python, which it is doing, it moves into NVIDIA’s Python rather than someone else’s.
The cautionary tale: OpenCL
Worth remembering, because the failure mode repeats. OpenCL standardized the API and the language and left the libraries, the tuning, and the last-mile compiler quality to each vendor.
Portable source, unportable performance, no cuBLAS equivalent, no tensor core story, committee cadence against NVIDIA’s annual co-design cadence.
SYCL is a better designed successor with the same structural problem: it is a standard for the layer that was never the moat.
What replacing the kernel layer actually costs
Let me put a rough number on the library moat, since that is the part people assume is unbuyable.
Take LLM inference specifically, not the whole CUDA-X surface. The kernels that matter for a modern decoder-only model at serving time are: dense GEMM in several shapes, grouped or batched GEMM for MoE, prefill attention, decode attention with paged KV, an MLA variant if you serve DeepSeek-shaped models, quantize and dequantize paths for FP8 and FP4, RMSNorm, RoPE, fused SwiGLU, sampling, KV cache gather and copy, speculative verification, and the collectives.
Call it 30 to 50 kernels that appear in profiles, of which eight or so carry most of the time.
Assume a fully loaded peak-quality kernel engineer at somewhere between $350,000 and $600,000, call it $450,000 to be some kind of “average”. Assume a top-tier attention or GEMM kernel for a new architecture takes three to nine engineer-months including autotuning, numerics validation and the long tail of shape coverage, and the secondary kernels take one to two.
Eight critical kernels at six months each: roughly $1.8M
Forty secondary kernels at 1.5 months each: roughly $2.2M
Test infrastructure, numerics validation, CI across shapes and dtypes: call it 50 percent overhead
That lands somewhere around $6M to $15M per hardware generation to build a competitive inference kernel layer from scratch. For a company selling accelerators, that is a rounding error. It is roughly one percent of a single generation’s mask set and tape-out cost.
Which explains something that otherwise looks strange: AMD is actually close on inference and not close on training. The inference kernel surface is small enough to buy.
The training surface, plus the full CUDA-X library set, plus the numerics validation across a decade of recipes, is somewhere between ten and fifty times that, and the co-design loop that makes the kernels good on day one instead of month nine is not purchasable at any price.
Treat these figures as a stated model with stated assumptions, not as a measurement. The point is the ratio, not the absolute value: the part everyone talks about is the cheap part.
Where I might be wrong
Four ways this argument could be badly calibrated, in descending order of how much they would cost me.
The backend gap may be small. My case for layer 3 rests partly on the maxas era, which was Maxwell, and partly on recent reinforcement learning work over SASS schedules that finds wins by reordering ptxas output. If the residual there is three to five percent on typical kernels rather than the double digits the Maxwell results implied, then ptxas is close enough to optimal that the backend is not the moat at all, the libraries are, and I have spent a lot of words on the wrong floor.
I have not seen a rigorous modern measurement of that residual across a representative kernel set, and I would very much like to.
The labour market framing may be a category error. I argue the scarce resource is people who can read control bits, and therefore that search and autotuning erode it.
But the advantage might be organisational rather than headcount: eighteen months of pre-silicon access, a hardware team down the hall, and the ability to change the chip when the kernel is awkward. If that is the real mechanism, then no amount of autotuning compute closes it, because the loop being run is design, not optimization.
Tile IR may be commoditization rather than capture. I read the open dialect, the published bytecode format and the conformance suite as a way of owning the layer Triton was about to own. The more generous reading is that an MLIR dialect with a stable binary encoding and a test suite is precisely the artifact another vendor needs in order to target the same source, and NVIDIA has just handed it over. Both readings are consistent with the facts as of today. Which one is right will be visible in whether anyone outside NVIDIA ships a Tile IR backend.
And the boring one. I do not have a B200. Everything in this piece at the Blackwell layer comes from documentation, from other people’s benchmarks, and from reading disassembly conventions rather than fresh disassembly. The control field layout is reverse engineered by researchers, not published.
Where I have marked things Tier B or C, that is what the marking means, and I would treat the Hopper and Blackwell specifics as directionally right rather than precisely right until someone with the hardware checks them.
Five dated calls
Predictions with dates, so this piece can be scored rather than admired.
1. No third party ships a production-quality SASS assembler for Hopper or Blackwell before 2030. Production quality meaning it passes a public correctness suite across a non-trivial kernel set on more than one SKU. Falsified by any project that does.
2. NVIDIA never open sources the Tile IR to cubin lowering. The dialect, the bytecode, the spec and the conformance suite stay open. The optimizing backend does not. Falsified trivially and publicly if wrong.
3. Before the end of 2027, an automated kernel generator beats a vendor library on a headline GEMM or attention shape on non-NVIDIA hardware, with reproducible numbers. The search-versus-labour thesis stands or falls on this one.
4. AMD reaches within fifteen percent of NVIDIA on tokens per second per dollar for standard transformer inference on a like-for-like generation by end of 2027, and does not close the equivalent gap on training in the same window. The asymmetry is the prediction, not the number.
5. At least one more CUDA compatibility project loses its funding before the end of 2027. ZLUDA has now done it twice. The economics of chasing a virtual ISA that moves every eighteen months have been tested and they do not work, and I expect the next attempt to discover this independently.
What would actually work (for real)
I am sorry to say it, but I do not think the moat gets breached. I think it gets routed around, in specific places, for specific workloads, and here is where I would put effort if that were my job.
Own the tile layer, and make it a spec with teeth. Not a language, not a framework: a specification with a versioned binary encoding and a conformance test suite that vendors must pass. NVIDIA just told you this layer is the strategic one by shipping exactly that artifact. The mistake would be to let Tile IR become the only one.
Make numerics a specification too. This is the gap nobody is filling. A conformance suite that pins reduction order tolerances, accumulation types, scaling granularity for narrow formats, and determinism guarantees, with reference outputs. Right now “matches cuBLAS” is folklore transmitted through issue threads. Turning it into a testable contract would remove a real and unpriced switching cost.
Attack at inference, not at parity. The kernel surface is small, the workloads are known, the customers are cost-sensitive, and memory capacity per package is a lever NVIDIA does not always win. Trying to be CUDA-complete is the losing strategy; being excellent at forty kernels is the winning one.
Spend compute instead of headcount. Every dollar into autotuning, superoptimization and learned scheduling attacks the actual scarce resource. Helion beating hand-written Triton on average is the proof of concept. This is the only lever in the list that gets cheaper over time.
And keep the compiler argument in proportion. The binding constraint on who gets to serve AI workloads in 2026 is HBM supply, advanced packaging capacity, rack-scale networking and power. A perfect compiler on a chip you cannot buy, in a rack you cannot cool, connected by a fabric that stalls at 72 GPUs, wins nothing. The compiler moat is real and it is roughly the fourth most important moat NVIDIA has.
Check it yourself
Nothing above requires trusting me. The relevant artifacts are all inspectable with the toolkit you already have installed.
See the driver script decompose:
nvcc --dryrun -arch=sm_90 kernel.cu
Every stage, in order, with the real command lines. cicc, ptxas, fatbinary, nvlink.
Watch the virtual ISA:
nvcc -arch=sm_90 -ptx kernel.cu -o kernel.ptx
Note the unbounded %r virtual registers and the total absence of scheduling information.
Watch the last mile do the work:
ptxas -arch=sm_90 -v kernel.ptx -o kernel.cubin
The -v output tells you registers used, spill stores, spill loads, shared memory. That register count is a policy decision ptxas made on your behalf, and it sets your occupancy.
Look at the control bits:
nvdisasm -c -g kernel.cubin
cuobjdump --dump-sass kernel.cubin
The bracketed fields next to each instruction are the scheduling payload. Compare the same kernel compiled with and without -maxrregcount and watch the stall counts and barrier usage change.
Confirm the frontend is a commodity:
clang++ -x cuda --cuda-gpu-arch=sm_90 --cuda-device-only -S kernel.cu -o clang.ptx
Diff that against nvcc‘s PTX. They differ. Then push both through ptxas and compare the SASS. That comparison is the whole argument of this piece in two commands: two independent frontends, one backend, and the backend is where the differences stop mattering.
Find the boundary: try to turn a modified SASS listing back into a cubin using only NVIDIA tooling. There is no such command. That absence is the moat.
Confidence dossier
Tier A, directly verifiable from primary sources or first-party documentation
nvccis a driver invokingcicc,ptxas,fatbinaryandnvlink; observable via--dryrunNVVM IR is a documented LLVM IR subset with a public
libNVVMAPI; Clang’s NVPTX backend is upstreamPTX is a virtual ISA; the driver JIT-compiles it, providing forward compatibility
There is no NVIDIA-provided SASS assembler;
nvdisasmandcuobjdumpdisassemble onlyBlackwell deprecated
wgmmaand introducedtcgen05, Tensor Memory, single-thread MMA issue andcta_group::2CUDA Tile IR shipped in CUDA 13.1; the MLIR dialect, bytecode format and conformance suite are open source; the bytecode is specified as stable and versioned
A Triton to Tile IR backend exists, authored by NVIDIA
The CUDA EULA restricts translating SDK-generated output artifacts to non-NVIDIA platforms
wgmmarequires thesm_90aarchitecture-specific target andtcgen05requiressm_100aorsm_103a; code foratargets is not forward compatibleA
cubinis an ELF64 object with per-kernel.text,.nv.constant0and.nv.infosectionsBlackwell Tensor Memory is 256 KB per SM, 128 lanes by 512 columns of 32 bits, allocated in power-of-two column counts
MXFP8 uses one E8M0 scale per 32 elements; NVFP4 uses an E4M3 scale per 16 elements plus a tensor-level FP32 scale
Tier B, well-supported by multiple independent secondary sources or peer-reviewed measurement
The Volta-and-later control field layout: reuse, wait mask, read and write barrier indices, yield, stall count, packed into the instruction word
Fixed-latency instructions lack hardware RAW interlocks; incorrect stall counts produce incorrect results
Instruction encoding widened from 64 to 128 bits at Volta
H100 register file of 65,536 32-bit registers per SM, 64 maximum resident warps, eight-register allocation granularity, and the occupancy arithmetic that follows from them
FlashAttention 4 in CuTe DSL outperforming cuDNN 9.13 and Triton by the margins cited
ROCm 7.2 library and AITER performance improvements on MI300 through MI355X
Helion’s reported speedups over Triton and
torch.compileNCCL’s channel-per-thread-block structure, its protocol selection, and the fact that collectives consume SMs that are then unavailable for compute
In-network reduction in NVSwitch and InfiniBand switch silicon
SCALE’s benchmark claims against HIP and
nvcc
Tier C, single-source, historical, or reported rather than independently verified
maxasSGEMM efficiency figures on GM204 versus contemporaneous cuBLASThe annotated SASS fragment, which is reconstructed in the documented format rather than captured from a specific compilation
DeepSeek’s specific SM partitioning and custom PTX details, as described in the V3 paper and subsequent reporting
Spectral Compute’s NVIDIA Inception membership and its strategic reading
cuTile’s competitiveness with cuBLAS and cuDNN at large shapes, which is a first-party claim awaiting third-party benchmarks
Tier D, my model or my opinion, flagged as such
The $6M to $15M per-generation estimate for an inference kernel layer, and the 10x to 50x multiplier for the full surface
The claim that layer 3 is the load-bearing wall rather than layer 1
The reading of Tile IR as a strategic response to Triton’s position in PyTorch
The assertion that the moat is best modelled as a labour market
The ranking of the compiler moat as roughly fourth behind supply, packaging and networking
All five dated calls, which are forecasts and should be read as such
The four self-criticisms, which are my own estimate of where this argument is weakest
Bibliography
Primary documentation
NVIDIA, Parallel Thread Execution ISA. https://docs.nvidia.com/cuda/parallel-thread-execution/
NVIDIA, CUDA Binary Utilities (
cuobjdump,nvdisasm, SASS instruction set tables). https://docs.nvidia.com/cuda/cuda-binary-utilities/NVIDIA, NVVM IR Specification. https://docs.nvidia.com/cuda/nvvm-ir-spec/
NVIDIA, CUDA Compiler Driver NVCC. https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/
NVIDIA, License Agreement for NVIDIA Software Development Kits. https://docs.nvidia.com/cuda/eula/index.html
NVIDIA, CUDA Tile. https://developer.nvidia.com/cuda/tile
NVIDIA, Tile IR: Binary Format. https://docs.nvidia.com/cuda/tile-ir/latest/sections/bytecode.html
NVIDIA, cuTile Python documentation. https://docs.nvidia.com/cuda/cutile-python/
NVIDIA, CUTLASS documentation: tcgen05 MMA Programming Guide. https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/mma_docs/tcgen05_programming.html
NVIDIA, CUTLASS: Blackwell SM100 functionality. https://docs.nvidia.com/cutlass/4.2.1/media/docs/cpp/blackwell_functionality.html
AMD, ROCm documentation: vLLM inference optimization. https://rocm.docs.amd.com/en/latest/how-to/rocm-for-ai/inference-optimization/vllm-optimization.html
Microarchitecture and reverse engineering
Jia, Maggioni, Staiger, Scarpazza, Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking, arXiv:1804.06826. https://arxiv.org/pdf/1804.06826
Huerta, Abaie et al., Analyzing Modern NVIDIA GPU cores, arXiv:2503.20481. https://arxiv.org/pdf/2503.20481
Dissecting and Modeling the Architecture of Modern GPU Cores, MICRO 58, 2025. https://dl.acm.org/doi/10.1145/3725843.3756041
Luo et al., Benchmarking and Dissecting the Nvidia Hopper GPU Architecture, arXiv:2402.13499
CuAsmRL: Optimizing GPU SASS Schedules via Deep Reinforcement Learning, arXiv:2501.08071. https://arxiv.org/html/2501.08071v1
GPA: A GPU Performance Advisor Based on Instruction Sampling, arXiv:2009.04061
Finn, What happens when you run a CUDA kernel. https://fergusfinn.com/blog/what-happens-when-you-run-a-gpu-kernel/
Assembler and toolchain projects
Gray, maxas, Maxwell assembler and control code notes. https://github.com/NervanaSystems/maxas/wiki/Control-Codes
CuAssembler. https://github.com/cloudcores/CuAssembler
openptxas, gpuocelot. https://github.com/gpuocelot/openptxas
NVlabs, SASSI: Flexible GPGPU instrumentation, distributed as a closed-source
ptxasfork. https://github.com/NVlabs/SASSINVIDIA, cuda-tile. https://github.com/NVIDIA/cuda-tile
Kernel authoring layers
Colfax Research, CUTLASS Tutorial: Writing GEMM Kernels Using Tensor Memory for NVIDIA Blackwell GPUs. https://research.colfax-intl.com/cutlass-tutorial-writing-gemm-kernels-using-tensor-memory-for-nvidia-blackwell-gpus/
Colfax Research, CUTLASS Tutorial: Hardware-supported Block-scaling with NVIDIA Blackwell GPUs. https://research.colfax-intl.com/cutlass-tutorial-hardware-supported-block-scaling-with-nvidia-blackwell-gpus/
SemiAnalysis, Dissecting Nvidia Blackwell: Tensor Cores, PTX Instructions, SASS, Floorsweep, Yield. https://newsletter.semianalysis.com/p/dissecting-nvidia-blackwell-tensor
PyTorch, Helion: A High-Level DSL for Performant and Portable ML Kernels. https://pytorch.org/blog/helion/
pytorch/helion repository. https://github.com/pytorch/helion
PyTorch RFC, Inductor cuTile Backend, issue 175311. https://github.com/pytorch/pytorch/issues/175311
NVIDIA Developer Blog, Advancing GPU Programming with the CUDA Tile IR Backend for OpenAI Triton. https://developer.nvidia.com/blog/?p=112268
Red Hat Emerging Technologies, From hand-tuned to generated: a reproducible Triton GPU kernel benchmark across different vendors. https://next.redhat.com/2026/02/12/from-hand-tuned-to-generated-a-reproducible-triton-gpu-kernel-benchmark-across-different-vendors/
Modular, Structured Mojo Kernels Part 4: Portability and the Road Ahead. https://www.modular.com/blog/structured-mojo-kernels-part-4-portability-and-the-road-ahead
Mojo: MLIR-Based Performance-Portable HPC Science Kernels on GPUs for the Python Ecosystem, arXiv:2509.21039
Alternatives and portability
Spectral Compute, SCALE. https://scale-lang.com/ and https://github.com/spectral-compute/scale-docs
HPCwire, Spectral Compute Aims to Set CUDA Free. Will It Succeed?, July 2026. https://www.hpcwire.com/2026/07/09/spectral-compute-aims-to-set-cuda-free-will-it-succeed/
SCALE: Ahead-Of-Time Compilation of CUDA for AMD GPUs, Middleware 2024 demo track. https://dl.acm.org/doi/10.1145/3704440.3704782
ZLUDA, Update Q1 and Q2 2026: back to the roots. https://vosen.github.io/ZLUDA/blog/zluda-update-q1q2-2026/
Phoronix, ZLUDA For CUDA On Non-NVIDIA GPUs Enables AMD ROCm 7 Support. https://www.phoronix.com/news/ZLUDA-ROCm-7
Tom’s Hardware, Nvidia bans using translation layers for CUDA software, March 2024. https://www.tomshardware.com/pc-components/gpus/nvidia-bans-using-translation-layers-for-cuda-software-to-run-on-other-chips-new-restriction-apparently-targets-zluda-and-some-chinese-gpu-makers
vLLM Blog, Beyond Porting: How vLLM Orchestrates High-Performance Inference on AMD ROCm. https://vllm.ai/blog/2026-02-27-rocm-attention-backend
AMD, Speed is the Moat: Inference Performance on AMD GPUs. https://www.amd.com/en/developer/resources/technical-articles/2026/inference-performance-on-amd-gpus.html
tinygrad repository and release notes. https://github.com/tinygrad/tinygrad
Pall et al., GROMACS on AMD GPU-Based HPC Platforms: Using SYCL for Performance and Portability, arXiv:2405.01420
Model and systems work referenced
DeepSeek-AI, DeepSeek-V3 Technical Report, arXiv:2412.19437
Tom’s Hardware, DeepSeek’s AI breakthrough bypasses industry-standard CUDA for some functions. https://www.tomshardware.com/tech-industry/artificial-intelligence/deepseeks-ai-breakthrough-bypasses-industry-standard-cuda-uses-assembly-like-ptx-programming-instead
Dao et al., FlashAttention series, most recently FlashAttention 4 on Blackwell in CuTe DSL
Bauer et al., Singe: Leveraging Warp Specialization for High Performance on GPUs, PPoPP 2014
Corrections and disagreements are welcome and get published. If you have measured ptxas output against hand-scheduled SASS on Hopper or Blackwell, I would like to see the numbers.