Reverse Engineering Nvidia

NVIDIA's GPU Is Closed Hardware With a Compiler Inside It. Here's What Reverse Engineering Actually Found.

NVIDIA documents PTX and keeps SASS, the machine code your kernels actually run as, undocumented. Everything below that line was learned the hard way, by timing hand-written instructions on real silicon. This covers what that work found, who did it, and which results I'd trust.

The Compiler Is Part of the Hardware

The best single paper here is Analyzing Modern NVIDIA GPU Cores from UPC Barcelona. Its method is simple. Write short sequences of hand-written SASS, wrap them in reads of the GPU clock counter, and use the elapsed cycles to confirm or refute a hypothesis. The authors used CuAssembler for this, the latest in a line of third-party assemblers that began with MaxAS.

The central finding is that modern NVIDIA hardware doesn't check most hazards itself. It relies on the compiler to encode dependencies in control bits on every instruction. Three mechanisms do the work:

• Stall counter. For fixed-latency producers, the compiler sets a per-warp counter, and the hardware doesn't check RAW hazards for these instructions at all. A wrong value gives a wrong answer, not a stall.
• Dependence counters. For variable-latency work like memory loads, each warp has six counters (SB0–SB5), each counting up to 63, with a 6-bit mask per instruction saying which to wait on. One counter can be released at register read (for WAR hazards) and another at write-back, which matters because source operands are often read hundreds of cycles before the result lands.
• Reuse bit. The register file cache is software-managed through a per-operand reuse bit.

Two findings contradict what academic simulators assumed. There is no operand collector stage, since variable operand-fetch time would make fixed instruction latencies impossible. And the scheduler policy is "Compiler Guided Greedy Then Youngest", where the compiler's control bits steer a greedy-then-youngest warp selector.

The bank-conflict experiment is a good example of the method. Two back-to-back FFMAs took five cycles with odd source registers, six with one even, and seven with two even. That timing exposes a two-bank register file with one read port per bank.

The cost argument is where it gets interesting. Per the paper, the control-bit scheme needs 41 bits per warp, about 0.09% of an SM's register file, versus 5.32% for a scoreboard design with 63 consumers per entry. The authors' simulator model reached 13.98% mean error against an RTX A6000, 18.24 points better than Accel-Sim's baseline. They compare it to AMD, where a waitcnt instruction requires an explicit extra instruction rather than control bits.

Limits to keep in mind: the scheduler behavior is confirmed only for warps within one CTA, and the authors say they couldn't find a read policy matching every case. It's validated on Ampere and Turing, and the Hopper and Blackwell generations extend it with new tensor instructions.

Who Has Mapped the Rest

Mesa's NVK external hardware docs are the best index of public NVIDIA reverse engineering. They credit Kuter Dinel's nv_isa_solver with Hopper and Ada instruction encodings, point to nv-shader-tools for disassembly and fuzzing, and list envytools for Maxwell and earlier. They also describe tooling for dumping command sequences from the proprietary driver (envyhooks) and disassembling them (nv_push_dump).

That work feeds a real driver. Per the nouveau status page, NVK became a conformant Vulkan 1.4 implementation in December 2024, and was enabled for Blackwell and Kepler starting with Mesa 25.2.

The Driver Split Almost Nobody Mentions

The open-gpu-kernel-modules release changed what's reversible. NVIDIA started shipping open kernel modules with the R515 driver branch. Behind them sits a firmware boundary: the GPU System Processor is a RISC-V coprocessor, and NVIDIA's Resource Manager is split into a host CPU portion and a GSP portion. That has two consequences:

1. Nouveau on RTX 40-series and later requires the GSP firmware blobs.
2. There was no stable ABI for the GSP firmware at that point, so each driver version brought revised firmware. Treat the ABI point as dated, since it describes 2023.

The reverse-engineering target therefore moved. You can read the host side of the RM interface, but the firmware stays closed.

A Clean-Room CUDA Stack, With Footnotes

The most striking recent claim is a preprint on a clean-room toolchain for the GB10 (DGX Spark), with code in the CUDA++ repo. Its reported approach has three parts:

• A roughly 900-line driver client written against the open kernel module's public headers, reproducing the object, memory, channel and submission sequence libcuda performs over ioctl.
• Instruction encodings recovered from three sources: NVIDIA's compiler as an offline oracle on probe kernels, the open-source Mesa NAK compiler, and the chip itself, with arithmetic models fitted until output matched bit for bit.
• Launches assembled in the order Mesa's NVK uses, with the launch descriptor (QMD) built directly.

The headline result: a 35-billion-parameter mixture-of-experts model served with zero CUDA libraries mapped in the process, its output identical to a host reference. Read the caveats before repeating that. It's a single-author Zenodo preprint, not peer reviewed, and the author says it was built by one engineer directing AI coding agents. The paper states its own boundaries: NVIDIA's open kernel module remains in the path, and GSP firmware stays NVIDIA's because the chip verifies its signature. So this is a replacement for the user-space stack, not for the whole trust chain.

Blackwell: Measured Versus Marketed

Two microbenchmark papers cover Blackwell, and they differ in how much to trust them. The GB203 study reports that low-precision mma.sync PTX instructions translate to new SASS opcodes (QMMA and OMMA) on Blackwell, where Hopper consistently uses HMMA.

The B200 study describes the architectural shift in more detail. Per the paper, Blackwell replaces warp-synchronous MMA with tcgen05.mma, a single-thread instruction, with operands from shared memory or a new Tensor Memory (TMEM). TMEM is 256 KB per SM, laid out as 128 lanes by 512 columns of 32-bit cells, managed in software through the tcgen05 instruction family. Its PTX-to-SASS table maps FP4 to OMMA and FP6/FP8 to QMMA, and FP64 is not supported by tcgen05.mma.

I'd verify its numbers before quoting them. The paper says it couldn't share code under double-blind review, uses hedged phrasing such as "consistent with a spatial array design" for its microarchitecture inferences, and says dual-die behavior across the NV-HBI link remains future work. A reasonable way to use it: take the instruction mappings and TMEM structure as leads, and re-measure any latency or throughput figure on your own hardware.

What Reverse Engineering Can't Tell You

It shows structure, not intent. Microbenchmarks recover latencies, port counts and policies, and they say nothing about why NVIDIA chose them. They also age fast, since each generation changes tensor instructions while the control-bit scheme carries forward. As a cross-check, the UPC authors note NVIDIA's own diagrams show significant changes limited mostly to tensor cores, ray tracing units and minor features.

The practical rule for this space: trust a claim only if it traces to a timing experiment, an open driver header, or a disassembly you can reproduce.

Follow for more coverage on MCP, agentic AI, and AI infrastructure.