How CUDA Binaries Actually Work
A byte-level surgical dissection of the cubin and fatbin formats, and the second encoding nobody documents.
Intro
I spent a good part of last year writing an x86-64 assembler by hand, in C, for no more reason than wanting to know what an object file really is. The thing that surprised me was not the instruction encoding, but was the metadata.
A modern object file is essentially a negotiation between the compiler and the loader about few things the machine code itself cannot say properly: where the arguments are, how much stack to reserve, which symbols are entry points, what has to be patched before anything runs.
So when I started reading GPU binaries, the question I kept asking more and more was not “what does this instruction do.” It was “what is this file promising the driver.”
The answer turns out to be… well, a lot, and almost none of it is written down. NVIDIA just documents the container in one sentence and the disassembler in twelve pages. And that’s basically it.
The CUDA Binary Utilities manual says a cubin is “an ELF-formatted file which consists of CUDA executable code sections as well as other sections containing symbols, relocators, debug info, etc.” That is the specification, and the word “etc.” is where the whole platform lives.
I want you to imagine this piece as a medical dissection, if it’s possible to say so. Every single number, hex dump and section listing below, came out of a container on my own machine, and the exact commands are in the appendix so you can disagree with me precisely.
Two things made it possible: first thing is that you do not need a GPU to compile, link or disassemble CUDA binaries: ptxas, fatbinary, nvdisasm and cuobjdump are a bunch of ordinary x86 programs. The second is that the entire toolchain is on PyPI, so getting a byte-exact CUDA 13.3.73 install is actually just one pip install.
Methodology
Everything measured here is from CUDA 13.3.73 (built 9 June 2026) installed from the nvidia-cuda-nvcc, nvidia-cuda-cuobjdump and nvidia-cuda-nvdisasm wheels on x86-64 Linux, plus the cuBLAS 13.6.0.2 wheel for the library dissection.
I do not currently have a GPU, which means every claim here is about what the toolchain emits and not about what silicon does with it. There are a few parts where I’m reading structure out of bytes instead of documentation, dont’t worry: I’ll say so and I give it like a confidence tier at the end.
NVIDIA publishes no specification for the cubin section layout, the .nv.info attribute encoding, the fatbin entry header or the Mercury sections, and all the interpretations are mine, unless specified.
Seven processes to compile one file
I want you to introduce the format and the pipeline now, because the shape of the output is derived from it. Keep in mind that nvcc is not a compiler, it’s just a driver that orchestrates other programs, and it will show you its plan.
$ nvcc -arch=sm_90 -c -o k.o saxpy.cu --dryrun
gcc -D__CUDA_ARCH_LIST__=900 -E -x c++ -D__CUDACC__ ... # host preprocess
cudafe++ --c++17 --static-host-stub --device-hidden-visibility ... # split host/device
gcc -D__CUDA_ARCH__=900 -E -x c++ -DCUDA_DOUBLE_MATH_FUNCTIONS # device preprocess
cicc ... saxpy.cudafe1.gpu -o saxpy.ptx # C++ front end -> PTX
ptxas -arch=sm_90 -m64 saxpy.ptx -o saxpy.sm_90.cubin # PTX -> SASS
fatbinary -64 --cicc-cmdline=... --image3=... -o saxpy.fatbin.c # package
gcc -c -x c++ saxpy.cudafe1.cpp -o k.o # host compileWe have two preprocessor passes over the same source with “different macros”, then a source-to-source splitter, a front end that emits PTX, an assembler that (you know) emits machine code, a packager that turns the machine code into a C array, and a host compiler that swallows the result.
The artifacts in the middle are real files you can save with --keep, and each one is a place where the format is decided.
The only thing that matters is where the close part actually is. We alredy know that cicc is the NVVM-based front end and it produces PTX, which is a published virtual instruction set, with all docs that you can read. But for ptxas,which takes PTX and produces the thing this article is about, we don’t know nothing.
Everything that works below PTX is the vendor’s private business, and everything upstream has alredy been made public: Clang, Triton, Julia, Mojo and tinygrad all emit PTX, and NVIDIA itself ships libNVVM, so they can do it.
The interesting feature in the CUDA stack is just one process wide.
The container is an ELF file that lies about being one
Now, let’s compile a kernel to a standalone cubin and we will see that the file recognizes it immediately.
$ nvcc -arch=sm_90 -cubin -o saxpy.sm90.cubin saxpy.cu
$ file saxpy.sm90.cubin
saxpy.sm90.cubin: ELF 64-bit LSB executable, NVIDIA CUDA architecture, version 1,
statically linked, not strippedIt is a real ELF64, little endian, and readelf will walk its section table happily. Three fields in the 64-byte header carry the CUDA-specific part.
e_ident: 7f 45 4c 46 02 01 01 41 08 00 00 00 00 00 00 00
| |
| +-- EI_ABIVERSION = 8
+----- EI_OSABI = 0x41
e_type = 2 (ET_EXEC; ET_REL when compiled with -rdc=true)
e_machine = 190 (0xbe, EM_CUDA)
e_flags = 0x06005a04EM_CUDA = 190 is the one part of this that is available to everybody: it sits in LLVM’s BinaryFormat/ELF.h alongside every other machine type, because LLVM needs to read these files.
The OS ABI byte is where things get strange. That same header defines two CUDA values, ELFOSABI_CUDA = 51 and ELFOSABI_CUDA_V2 = 41, both written in decimal. Every cubin my 13.3 toolchain produces carries 65, which is 0x41. 51 is 0x33, so the first constant is the older value written in decimal, and the second…. looks like 0x41 transcribed as if it were decimal.
Obv. I don’t know where is the mistake, but if you are writing a tool that sniffs cubins, sniff for the byte 0x41 and don’t trust either constant. The ABI version byte is 8.
I know that e_flags is where the architecture lives, and the field names are recoverable from the disassembler: EF_CUDA_SM, EF_CUDA_PTX_SM, EF_CUDA_64BIT_ADDRESS and EF_CUDA_ACCELERATORS, alongside an EF_CUDA_SM10 through EF_CUDA_SM121 enumeration that covers fifteen years of hardware in one list.
The layout is easy peasy to recover by sweeping every target the compiler supports.
The same four-line kernel, one target per row. Bits 8 to 15 hold the compute capability as a plain integer. Byte 0 flips from 0x04 to 0x02 at the Blackwell boundary, and the file grows by 44% across that same boundary for identical source.
This is actually very good, but at least three things don’t match reality:
First, the virtual architecture is not in
e_flagsanymore, despite the name in the field. Compile the same kernel three ways, with the PTX architecture set tocompute_75,compute_80andcompute_90but the real target fixed atsm_90, and all three cubins carrye_flags = 0x06005a04.The virtual architecture actually moved into a note section, where
cuobjdumpreports it asCUDA Virtual SM: sm_75and so on. This used to be visible: CUDA 8-era disassembly printed a line reading.headerflags @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)", and currentnvdisasmprints a plain.target sm_90instead. If you have old tooling that reads the PTX architecture out of the flags word, you have seen zero for some years.
Second, the architecture-conditional suffixes do not show up here either.
sm_90andsm_90aproduce identicale_flags, and so do all three ofsm_100,sm_100aandsm_100f. For a kernel that uses no special features thesm_100andsm_100fcubins are the same program byte for byte, and the 210 bytes that do differ between the two files are all downstream of one string: the note section records-arch sm_100finstead of-arch sm_100, which is one byte longer and shifts everything after it.That is in line with what NVIDIA says about the feature, which is that
code=sm_100andcode=sm_100fare aliases producing the same cubin when no family-specific features are in play. The suffix is recorded, but the recording is somewhere else and is “triggered” only when the kernel actually uses something.
Third, CUDA 13’s supported target list is very short, and that surprised even me. Turing is now the floor.
ptxasin 13.3 accepts exactlysm_75, sm_80, sm_86, sm_87, sm_88, sm_89, sm_90, sm_90a, then the three-digit generation:sm_100, sm_103, sm_110, sm_120, sm_121each withaandfvariants, plus the matchinglto_*targets.Maxwell, Pascal and Volta were removed in CUDA 13.0, which NVIDIA’s release notes state directly: offline compilation and library support for those architectures are gone completely, and 12.x is the last line that can target them.
You have to know at least two numbers if you ever have to explain a build matrix. sm_110 is Jetson AGX Thor, renamed from sm_101 when CUDA 13.0 shipped, which means build scripts that hardcoded 101 broke on a rename rather than on a chip. NVIDIA states the rename in the nvCOMPDx release notes rather than anywhere prominent.
And sm_88, compute capability 8.8, is said by the community compatibility references to be the Nintendo Switch 2, which I cannot verify from the toolchain itself: ptxas will happily build for it and tells you nothing about what it is.
If that claim is right, the CUDA binary format has a target for a games console sitting in the same list as your H100.
Fifteen sections, but only one is code
Below I will show you the full section table of a cubin containing one trivial kernel, a SAXPY with four parameters, compiled for Hopper.
idx name type off size align
1 .shstrtab STRTAB 64 295 1
2 .strtab STRTAB 406 369 1
3 .symtab SYMTAB 776 240 8
4 .debug_frame PROGBITS 1016 104 1
5 .note.nv.tkinfo NOTE 1120 164 4
6 .note.nv.cuinfo NOTE 1284 32 4
7 .nv.info CUDA_INFO 1316 36 4
8 .nv.compat CUDA_COMPAT_INFO
1352 28 4
9 .nv.info._Z5saxpyifPKfPf CUDA_INFO 1380 132 4
10 .nv.callgraph CUDA_CALLGRAPH
1512 32 4
11 .rela.debug_frame RELA 1544 24 8
12 .text._Z5saxpyifPKfPf PROGBITS 1664 512 128
13 .nv.shared.reserved.0 NOBITS 2176 0 1
14 .nv.constant0._Z5saxpyifPKfPf PROGBITS 2176 552 4We have 512 bytes of machine code inside a 3,968 byte file. The italicised sections are the CUDA-specific ones, with section types in the 0x70000000 range that generic ELF tools print as “unknown”.
The code section is said to be “per kernel” and named after the mangled symbol, which is why a file with ten kernels has ten .text.* sections and ten of most other things too. Alignment is 128 bytes, one instruction cache line’s worth of paranoia.
Then, in a random order of how much they surprised me:
.note.nv.tkinfo records how the binary was built
This is a standard ELF note with owner string NVIDIA Corp, and cuobjdump will print it for you.
$ cuobjdump -elf saxpy.sm90.cubin | grep -A5 “CUDA Toolkit Information”
NVIDIA Corp 140 NVIDIA CUDA Toolkit Information
Note Version: 2
Tool Name: ptxas
Tool Version: Cuda compilation tools, release 13.3, V13.3.73
Tool Branch: Build cuda_13.3.r13.3/compiler.38244171_0
Tool Command Line Arguments: -arch sm_90 -m 64Every cubin has the exact compiler build that produced it and the flags it was invoked with. That's a provenance record, and it’s sitting in every shipped binary on every machine learning platform you have ever installed and used.
If you want to know which toolkit built the kernels inside somebody’s wheel, no need to ask them. There is also a -verbose-tkinfo mode in the assembler that emits, in its own words, “object name and command line arguments which contains all arguments having file format,” meaning paths. I havent seen it on by default, and I’d look before shipping a binary built with it.
.nv.constant0 is the launch frame, and it keeps growing
Kernel parameters don’t come in registers, but in constant bank 0, and the cubin declares a per-kernel section for it. On Hopper we said this section is 552 bytes for a kernel whose parameters occupy 24. The user arguments start at offset 0x210; everything below that is owned by the driver and the ABI.
You can watch the exact boundary move across generations. It is at 0x160 for Turing through Ada, 0x210 for Hopper, and 0x380 for Blackwell and later, with the section growing from 380 to 556 to 924 bytes for the same kernel.
This is the sort of constant that reverse-engineering projects hardcode and then rediscover every two years.
You can see the effect in the disassembly, which is the easiest way I know to make constant bank 0 sort of…. real. Here is the same kernel on Hopper and on Thor.
sm_90 sm_110
LDC R1, c[0x0][0x28] LDC R1, c[0x0][0x37c]
S2R R0, SR_TID.X S2R R0, SR_TID.X
S2UR UR4, SR_CTAID.X S2UR UR4, SR_CTAID.X
LDC R7, c[0x0][RZ] LDCU UR5, c[0x0][0x380]
IMAD R7, R7, UR4, R0 LDC R7, c[0x0][0x360]
ULDC UR4, c[0x0][0x210] IMAD R7, R7, UR4, R0
ISETP.GE.AND P0, PT, R7, UR4, PT ISETP.GE.AND P0, PT, R7, UR5, PT
@P0 EXIT @P0 EXITLeft: the first parameter is loaded from 0x210. Right: from 0x380, using LDCU, an instruction that loads a constant straight into a uniform register and that does not exist in the Hopper instruction set reference. Even the stack pointer moved, from 0x28 to 0x37c.
The symbol table has its own vocabulary
Kernels are ordinary global function symbols with two CUDA-specific decorations, but if you look carefully, ther's one symbol in a cubin that has nothing to do with your code.
index value size info other shndx name
0x3 0 0 0x3 0 0xc .text._Z5saxpyifPKfPf
0x4 0 0x4 0x21 0 0 .nv.reservedSmem.offset0
0x5 0 0 0x20 0xa0 0xd __nv_reservedSMEM_offset_0_alias
0x8 0 0x200 0x12 0x10 0xc _Z5saxpyifPKfPf
0x9 0 0 0x3 0 0xe .nv.constant0._Z5saxpyifPKfPfThe kernel symbol carries st_other = 0x10, which nvdisasm prints as STO_CUDA_ENTRY STV_DEFAULT: a visibility byte reused to mark launchable entry points, which is how the driver tells a kernel from a device function without parsing names.
The reserved shared memory symbols are the interesting thing here. Even a kernel that declares “no shared memory” unexpectedly gets a reservation record, and on Blackwell that reservation has a non-zero size: somehw you see 64 bytes of shared memory that belong to the platform, not to you, aliased under a name beginning with __nv_. Anyone computing occupancy from source is computing it from the wrong number.
The same source is not the same program
I decided to write a whole paragraph to the architecture sweep, because it shows how much the cubin is influenced by the specific chip, rather than by your code.
The four-line SAXPY compiles to 24 instruction slots on sm_86 and 48 on sm_87, Orin, and the difference is not optimization.
sm_86 opcode histogram: 8 NOP 2 S2R 2 MOV 2 LDG.E 2 IMAD.WIDE ...
sm_87 opcode histogram: 15 BMOV.32.CLEAR 12 NOP 2 S2R 2 LDG.E ...We see fifteen instructions clearing barrier state at kernel entry, on one embedded part, for a kernel that uses no barriers.
In other words, that piece of code inside the kernel isn’t there because your program desperately needs it. It’s here because NVIDIA is paying, through software, the price of a silicon limitation. And that proves that SASS is dependent not only on the PTX, but also on the hardware, the drivers and the workarounds embedded into the toolchain.
.nv.callgraph exists because device code can call things
Fixed 8-byte entries, and for a leaf kernel it is four rows of negative sentinels:
.nv.callgraph
<0,-1> <0,-2> <0,-3> <0,-4>The sentinels are the interesting part: a leaf kernel does not get an empty section, but four explicit “no callee of kind N” markers. This suggests that .nv.callgraph is not simply a list of callees, but a structured description of call relationships understood by the runtime.
That metadata matters because features such as indirect calls, dynamic parallelism, and separately compiled device functions require the driver to know the worst-case call depth before launch so it can size the per-thread stack correctly.
Even a kernel that calls nothing still participates in this contract, which is why absence is encoded explicitly rather than by omitting the section entirely.
.nv.info is where the kernel describes itself
This is the section that let you see a loadable cubin, rather than merely readable. The format is a simple tag-length-value stream: one byte of format, one byte of attribute, then either nothing, one byte, two bytes, or a 16-bit length followed by that many bytes.
And….thirty-six bytes of it, at file level, look like this:
04 2f 08 00 | 08 00 00 00 0a 00 00 00 EIATTR_REGCOUNT
04 11 08 00 | 08 00 00 00 00 00 00 00 EIATTR_FRAME_SIZE
04 12 08 00 | 08 00 00 00 00 00 00 00 EIATTR_MIN_STACK_SIZE
^ ^ ^
| | +-- value length (EIFMT_SVAL)
| +----- attribute code
+-------- format: 1 = none, 2 = one byte, 3 = two bytes, 4 = length-prefixedThe first value that you read is a symbol table index, the second is the payload. The symbol 8 in each column is the kernel, and it needs 10 registers, no stack frame, no minimum stack. You can check that decode against the tool, because cuobjdump -elf knows the names.
$ cuobjdump -elf saxpy.sm90.cubin
.nv.info
Attribute: EIATTR_REGCOUNT Value: function: _Z5saxpyifPKfPf(0x8) register count: 10
Attribute: EIATTR_FRAME_SIZE Value: function: _Z5saxpyifPKfPf(0x8) frame size: 0x0
Attribute: EIATTR_MIN_STACK_SIZE Value: function: _Z5saxpyifPKfPf(0x8) min stack size: 0x0The per-kernel section is way richer. Let’s see a four-parameter SAXPY:
.nv.info._Z5saxpyifPKfPf
EIATTR_LANGUAGE PTX
EIATTR_CUDA_API_VERSION 0x85 (133 = CUDA 13.3)
EIATTR_KPARAM_INFO Ordinal 0x3 Offset 0x10 Size 0x8 Space CBANK
EIATTR_KPARAM_INFO Ordinal 0x2 Offset 0x8 Size 0x8 Space CBANK
EIATTR_KPARAM_INFO Ordinal 0x1 Offset 0x4 Size 0x4 Space CBANK
EIATTR_KPARAM_INFO Ordinal 0x0 Offset 0x0 Size 0x4 Space CBANK
EIATTR_SPARSE_MMA_MASK 0x0
EIATTR_MAXREG_COUNT 0xff
EIATTR_MERCURY_ISA_VERSION 1.1
EIATTR_EXIT_INSTR_OFFSETS 0x70 0x120
EIATTR_CBANK_PARAM_SIZE 0x18 (24 bytes of parameters)
EIATTR_PARAM_CBANK 0x9 0x180210 (symbol 9, size 0x18 at 0x210)
EIATTR_SW_WAR 0x8
EIATTR_NVSAL_SW_WAR 0x1Kernel parameters are described entirely through size-and-offset metadata, with entries stored in something called “reverse ordinal order”.
This means the host runtime doesn’t have to understand the original programming language or the function signature. It can construct the kernel’s parameter buffer directly from the metadata, placing each argument at the correct location before launch. The compiled binary therefore exposes a language-independent ABI that the driver and runtime can consume.
EIATTR_EXIT_INSTR_OFFSETS contains the byte offset of every EXIT instruction generated for the kernel. At first glance this is considered to be like a minor detail, but it’s very useful for tooling.
A profiler, tracer, or instrumentation framework can find every kernel return point immediately, without performing a full disassembly or reconstructing the control-flow graph. If a tool have to inject timing code, coverage probes, or custom telemetry before kernel termination, these offsets provide the exact insertion points.
EIATTR_SW_WAR is a different kind of purpose here. It encodes software workarounds for known hardware errata as a bitmask. Modern GPUs sometimes require compiler- or driver-level mitigations to avoid wrong behavior in specific corner cases.
Rather than hard-coding these decisions somewhere else, the compiler do a record of the workarounds that are needed, in the binary metadata. The driver then inspects the flags at load time and enables the correct mitigation paths for the architecture in place.
All these attributes show you that CUDA metadata is not just a bunch of descriptions; it forms actively some parts of the contract between the compiler, runtime, driver, profiling tools, and the GPU itself, carrying the information needed to launch kernels, instrument execution, and safely navigate hardware-specific quirks.
Now a question arises naturally: how big is this vocabulary? The disassembler contains the enum names, so you can just ask it, and you’ll find out.
$ strings -a nvdisasm | grep -o “EIATTR_[A-Z0-9_]*” | sort -u | wc -l
113Now you see it: 113 attribute kinds. A few of them are a decent map of what the hardware has learned to do since 2010.
EIATTR_TCGEN05_1CTA_USEDandEIATTR_TCGEN05_2CTA_USEDflag use of Blackwell’s fifth-generation tensor cores, and the fact that there are separate flags for the one-CTA and two-CTA forms tells you the driver has to care which.EIATTR_CTA_PER_CLUSTER,EIATTR_MAX_CLUSTER_RANK,EIATTR_EXPLICIT_CLUSTERandEIATTR_BLOCKS_ARE_CLUSTERSare the Hopper cluster model.EIATTR_STACK_CANARY_TRAP_OFFSETSmeans device code has stack canaries now, and indeednvcchas a flag for them.EIATTR_COROUTINE_RESUME_ID_OFFSETSis there for something that has not been announced.And a small family of attributes are named after bug numbers:
EIATTR_SW1850030_WAR,EIATTR_SW2393858_WAR,EIATTR_SW2861232_WAR,EIATTR_WAR5829587_NEEDED.
There is also an internal ticket for each somewhere, and the workaround is load-bearing enough to have its own slot in the binary format.
The format is not a description of a program. It is a contract between a compiler and a driver that ship on different schedules and are written by the same company.
Twenty-one bits per instruction that the hardware does not check
Starting with Volta, SASS instructions are 128 bits, and a piece of every instruction word is not the instruction at all, because it’s the scheduling metadata that the assembler computes and the hardware obeys without verifying.
You will find plenty of documentation about the structure of the field in the microbenchmarking literature, rather than by the vendor: a paper by Jia and colleagues describes stall counts, a yield flag, separate read and write dependency barriers, a wait mask that is a bitmask because an instruction can wait on several barriers at once, and reuse flags.
The exact bit positions below are the ones I decoded, and I trust them because the result is coherent with the docs: stall count in bits 105 to 108, yield flag at 109, write barrier index at 110 to 112, read barrier index at 113 to 115, a six-bit wait mask at 116 to 121, and four reuse flags at 122 to 125. Bits 126 and 127 come out zero on every instruction of every kernel I looked at, on five architectures.
You can decode it yourself with twenty lines of Python, and I think everyone who works on inference should do it at least once, because the mechanism explains more about the CUDA moat than any benchmark. Here is a real Hopper GEMM prologue with the fields pulled out.
off stall y wr rd wait reuse instruction
0000 1 1 - - 000000 0000 LDC R1, c[0x0][0x28]
0010 1 1 0 - 000000 0000 S2R R9, SR_CTAID.Y
0020 1 1 - - 000000 0000 ULDC UR4, c[0x0][0x228]
0030 1 1 - - 000000 0000 ULDC.64 UR6, c[0x0][0x208]
0040 1 1 - - 000000 0000 MOV R0, UR4
0050 1 1 0 - 000000 0000 S2R R13, SR_TID.Y
0060 2 1 - - 000000 0000 HFMA2.MMA R16, -RZ, RZ, 0, 0
0070 1 1 - - 000000 0000 ISETP.GE.AND P0, PT, R0, 0x1, PT
0080 1 1 1 - 000000 0000 S2R R17, SR_CTAID.X
0090 1 1 1 - 000000 0000 S2R R7, SR_TID.X
00a0 2 0 - - 000001 0000 LEA R9, R9, R13, 0x5
00b0 8 0 - - 000010 0000 LEA R11, R17, R7, 0x5Two special register reads allocate barrier 0, and the LEA that consumes them waits on barrier 0. Two more allocate barrier 1, and the next LEA waits on barrier 1. The dependency is not discovered at run time, but written into the binary.
That is the whole argument, in twelve instructions. For variable-latency instructions the assembler allocates one of six barriers and makes consumers wait on a bitmask.
For fixed-latency instructions it does not use barriers at all, it just writes a number of cycles into the stall field and the scheduler holds off that long. There is no interlock checking this: if the number is wrong, you do not get a stall, you get wrong answers.
Across a full tiled GEMM kernel the distribution looks like this.
The linker can rewrite the control bits
The relocation namespace makes the point better than I can. Pull the type names out of nvlink and there are 119 of them, and most encode a bit position in the name.
$ strings -a nvlink | grep -o “R_CUDA_[A-Z0-9_]*” | sort -u | wc -l
119
R_CUDA_ABS32_23 patch 32 bits at bit offset 23
R_CUDA_ABS32_HI_32 high half of a 64-bit address, at bit 32
R_CUDA_CONST_FIELD19_20 19-bit constant bank field at bit 20
R_CUDA_ABS55_16_34 55-bit value split across bits 16 and 34
R_CUDA_INSTRUCTION128 a whole 128-bit instruction word
R_CUDA_PCREL_IMM24_23 24-bit program counter relative branch at bit 23
R_CUDA_YIELD_CLEAR_PRED4_87 clear 4 bits at bit 87
R_CUDA_YIELD_OPCODE9_0 rewrite a 9-bit opcode field at bit 0On a CPU, relocations patch whole bytes because operands are byte-aligned. In this case, the operands are bit fields inside a 128-bit word, so the relocation type has to name the field, and there is a separate type for the same width at each position it can occur.
The last two are the ones that really stopped me. A relocation that clears the yield predicate, and one that rewrites an opcode, mean the device linker’s contract includes editing scheduling and instruction selection after the assembler has finished.
The 21 control bits are not a compiler-internal detail, which is nested at the end of ptxas, because they are a key part of the object format, and a later stage is allowed to change them.
I have argued many times before that this field, not the language and not PTX, is the load-bearing part of the CUDA moat, and taking cubins apart has not changed my mind. It has sharpened one point though.
The reason the assembler is fully close is not that NVIDIA wants to hide optimizations and reverse-engineering, at least not the first point. What we also need to keep in mind is that the ptxas is correctness-critical infrastructure.
An external tool that emits SASS is not competing with a compiler on quality of code, it is taking over responsibility for a hazard-avoidance protocol with no runtime check and no error reporting, which is extremely dangerous.
NVIDIA’s own programming guide is blunt about the consequence: you read that binary compatibility is promised only for binaries created by NVIDIA tools: manual editing or generating binary code is not supported, and compatibility promises are invalidated if binaries are modified in any way.
Read that as a description of the support boundary rather than a threat, and the closure looks less like strategy and more like the only sentence a vendor can write once it has moved the interlock into software.
Blackwell cubins carry a second copy of every kernel
When I thought about writing this part, I didn’t even know what to expect. It was a complete surprise to me. Let’s do a concrete example: compile the same trivial kernel for Hopper and for Blackwell and count sections: 15 for sm_90, 22 for sm_100. The seven extra sections are not debug info.
idx name type flags size
12 .text._Z5saxpyifPKfPf PROGBITS 6 512
13 .nv.shared.reserved.0 NOBITS 3 64
14 .nv.constant0._Z5saxpyifPKfPf PROGBITS 42 920
15 .nv.capmerc.text._Z5saxpyifPKfPf 0x70000016 10000000 258
16 .nv.merc.debug_frame PROGBITS 10000000 112
17 .nv.merc.nv.info 0x70000083 10000000 36
18 .nv.merc.nv.info._Z5saxpyifPKfPf 0x70000083 10000040 144
19 .nv.merc.rela.debug_frame 0x70000082 10000040 24
20 .nv.merc.nv.shared.reserved.0 0x70000015 10000003 0
21 .nv.merc.symtab 0x70000085 10000000 216You are now staring at a complete parallel object: its own symbol table, its own info sections, its own relocations, its own shared memory reservation, and a “capmerc” text section. Section flag bit 0x10000000 marks the whole family.
The prefix is merc, and the name shows up in three other places. It is in the per-kernel info as EIATTR_MERCURY_ISA_VERSION 1.1. It is in a new section called .nv.compat, which exists on every target but says more on Blackwell, and it is all over the assembler binary.
$ cuobjdump -elf saxpy.sm100a.cubin | sed -n ‘/nv.compat/,/nv.info\._/p’
.nv.compat
EICOMPAT_ATTR_CUDA_ACCELERATOR_TARGET 0x1
EICOMPAT_ATTR_ISA_CLASS 0x1
EICOMPAT_ATTR_INST_TCGEN05_MMA 0x5
EICOMPAT_ATTR_MERCURY_ISA_MAJOR_MINOR_VERSION_V2 1.1
EICOMPAT_ATTR_MERCURY_ISA_MAJOR_MINOR_VERSION_V1 1.1
EICOMPAT_ATTR_INST_TENSORMAP_V1 0x0
EICOMPAT_ATTR_CAN_FASTPATH_FINALIZE 0x9 0x0There is the missing suffix. EICOMPAT_ATTR_CUDA_ACCELERATOR_TARGET is 1 for sm_100a and 0 for both sm_100 and sm_100f, and the accelerated build declares EICOMPAT_ATTR_INST_TCGEN05_MMA, a capability it was permitted to use even though this kernel does not use it.
To really find out what the block really tracks, you need to build a kernel that needs a family-specific instruction. One line of inline PTX will do: tcgen05.fence::before_thread_sync, part of Blackwell’s fifth-generation tensor core interface.
Compiling it for five targets gives the feature tiers as a straight experiment rather than as a diagram.
$ for a in sm_100 sm_100f sm_100a sm_103f sm_120a; do nvcc -arch=$a -cubin -o t.cubin tc.cu; done
sm_100 Instruction ‘tcgen05.fence’ not supported on .target ‘sm_100’
sm_100f OK (5136 bytes)
sm_100a OK (5136 bytes)
sm_103f OK (5136 bytes)
sm_120a Instruction ‘tcgen05.fence’ not supported on .target ‘sm_120a’The baseline target refuses the instruction, but the family target accepts it, and so does the family target for the other member of the same family.
And the fully accelerated target for consumer Blackwell refuses it, because sm_120 is a different family and doesn’t have the hardware, which is the concrete answer to the frequently asked question of why a 5090 is not a small B200.
Now look at what the successful builds wrote into the object.
That is the piece I was missing when I first wrote this section, infact I had to rebuild it. The compat block is not a record of your compiler flags; it is a graded statement about what the machine code inside requires, computed from the code, and the flags only set the ceiling on what the assembler was allowed to reach for.
A resolver reading ISA_CLASS = 1 knows this object is safe anywhere in the generation; reading 2, it knows to check the family. This is the minimum viable metadata for the compatibility promise NVIDIA made in CUDA 12.9, and it is sitting in a 32-byte section nobody documents.
What reads it becomes obvious once you look inside ptxas. The strings are unambiguous.
$ strings -a ptxas | grep -iE “merc|finaliz” | sort -u
[Finalizer] fastpath optimization applied for off-target %u -> %u finalization
Generate Capsule Mercury
Specify the type of target ELF binary kind. Default on sm100+ is capmerc
Self check for capsule mercury (capmerc)
Specify the opportunistic finalization level. 0=default, 1=no opportunistic
finalization, 2=intra family finalization only, or 3=intra and inter family
finalization
Turns off the fast-path finalization optimization (allows normal refinaization)
Specify the ‘sm_’ name of the target architecture. If not specified, default
behavior is on-target finalization
R_MERCURY_NONE R_MERCURY_G64 R_MERCURY_ABS64 R_MERCURY_ABS32 R_MERCURY_ABS16 ...Read together with the sections, this describes a two-stage back end. Mercury is an encoding of a kernel that is below PTX, but above final machine code.
A capsule wraps it with the metadata a finalizer needs: its own symbol table, register and barrier counts, shared memory usage, and its own relocation types. “Finalization” is the step that turns a capsule into executable SASS for a specific chip. On-target finalization is the ordinary case.
Off-target finalization is the same capsule being finalized for a different chip than the one it was compiled for, with a fast path when the assembler can prove the existing encoding is already valid. And the levels of “opportunistic finalization” go from none, to within a family, to across families.
That is what the f targets are made of. NVIDIA introduced family-conditional compilation in CUDA 12.9 and actually described it in terms of feature sets: an f binary may use the architecture-specific features that are common to a whole family, and will run on later members of that family.
The public framing is a compatibility promise. The mechanism, as far as I can see it in the binary, is that the cubin ships a re-finalizable representation next to the SASS, and the driver can produce correct machine code for a family member the compiler never saw.
Where I am reading, not knowing
Unfortunately, I havent observed finalization happen. At the moment, don’t own a Blackwell GPU, and none of this machinery has the ptxas command line: passing --binary-kind gets you Unknown option, so the option table I am quoting belongs to an internal or driver-side entry point.
What I am confident about is the presence and shape of the artifacts. An independent reverse-engineering effort on ptxas 13.0 reached the same reading and adds detail I could not verify, including a 328-byte capsule descriptor and a compilation-knob snapshot inside it. Treat the purpose as tier C, so a speculative claim.
There is more of it in the section type table than any single cubin shows. cuobjdump has to be able to print a name for every section type it might meet, so the names are in the binary, and the Mercury family is large.
$ strings -a cuobjdump | grep -oE “CUDA_[A-Z_]{3,}” | sort -u | grep -i merc
CUDA_CAPMERC
CUDA_MERCURY
CUDA_MERCURY_CONSTANT_DRIVER CUDA_MERCURY_CONSTANT_PARAMS
CUDA_MERCURY_CONSTANT_IMGHDR CUDA_MERCURY_CONSTANT_PIC
CUDA_MERCURY_CONSTANT_OPT CUDA_MERCURY_CONSTANT_TOOLS
CUDA_MERCURY_CONSTANT_USER
CUDA_MERCURY_RESOLVED_RELA
CUDA_MERCURY_SASS_MAPSeven distinct constant bank section types, split by who owns the data: driver, image header, optimizer, parameters, position independent code, tools, user. A resolved-relocation type. And a section type called CUDA_MERCURY_SASS_MAP, which is hard to read as anything other than a correspondence between Mercury entities and SASS.
A format that needs a map to SASS is not an annotation on SASS, because it’s a representation of the program in its own right, and the map exists so that debuggers, profilers and the finalizer can move between the two.
All of which makes the size behaviour the really awkward part, because it argues against the simplest reading of what I just described.
If the capsule were a second copy of the instruction stream, in whatever denser encoding, its size would have to follow the SASS. Longer kernel, longer capsule. That is not what happens. Across seven kernels the SASS spans 384 to 2,048 bytes, a factor of 5.3, while the capsules span only 130 to 350, a factor of 2.7.
Correlations are carried almost entirely by a single long kernel: the coefficient across all seven is 0.76, and dropping that one point takes it to 0.33. Now, if you rank the kernels by capsule size and you get an ordering with no obvious relationship to how much code they contain.
The cleanest way to see it is a controlled pair. Two of these kernels compile to exactly 384 bytes of SASS. One stores a float, the other multiplies and adds in double precision. Same code size, same instruction count, and their capsules are 130 and 194 bytes, a 49% difference. Meanwhile a kernel using sqrtf compiles to 896 bytes of SASS, more than twice the double-precision kernel, and carries a smaller capsule at 184.
I want to be careful here, because when I first wrote this section I claimed the ordering tracked the variety of operations a kernel performs, and then I checked. It does not. Correlating capsule size against the number of distinct opcodes in each kernel gives 0.03, which is nothing at all. I could not find anything that predicts capsule size, and with seven kernels I would not trust a pattern even if I had found one.
What the double-precision pair does suggest is that whatever the capsule enumerates is closer to a set of requirements than to a program, and double precision is a plausible thing to find in such a set. Double-precision throughput is one of the sharpest differences within a Blackwell family: the datacenter parts and the consumer parts are not the same machine in that respect.
A representation whose job is to let a driver produce correct code for a family member the compiler never saw would need to record exactly this sort of thing, and would not need to grow much when you add another hundred instructions of the same kind.
So the sizes push against “a second copy of every instruction” and toward “a description of what this code needs.” They cannot separate the two readings that remain. A compact whole-kernel encoding in a much denser format, and a fixup table over the existing SASS naming only the sites a re-finalizer would have to revisit, would both be small and would both fail to scale with instruction count.
The section type called CUDA_MERCURY_SASS_MAP is the one piece of evidence that leans toward the first, since a table of fixups would not obviously need a map. I cannot settle it from sizes alone, and I doubt anyone can without a driver to watch.
What it costs is easier to state. On the ten-kernel workload, Mercury sections add 4.2 KB to a 57.8 KB cubin, about 7%, and Blackwell cubins are 2.0x the size of Turing cubins for identical source.
There is one more hint about direction of travel, in the disassembler rather than the assembler. nvdisasm in 13.3 has an option called --no-vliw, described as “conventional mode; disassemble paired instructions in normal syntax, instead of VLIW syntax.”
Nothing I compiled for any current target produced paired output. An option to turn off VLIW syntax implies a target where VLIW syntax is the default.
I would not build a thesis on a command-line flag, but if a future architecture issues statically paired instructions, the tooling has already been taught the word.
The fatbin is a filesystem, and the driver is the resolver
A cubin runs on exactly one compute capability. Shipping software means shipping several, plus PTX for the ones that do not exist yet, and the container for that is the fatbin.
It starts with a 16-byte header whose magic is 0xBA55ED50, and it is followed by a run of entries, each with its own header and payload.
$ nvcc -gencode arch=compute_90,code=sm_90 \
-gencode arch=compute_100,code=[sm_100,compute_100] -fatbin -o saxpy.fatbin saxpy.cu
kind arch hdrsz payload compressed uncompressed flags
CUBIN 90 64 3968 0 0 0x11
CUBIN 100 112 5712 0 0 0x1000011
PTX 100 80 400 399 922 0x8011The entry header fields are recoverable by differential compilation: change one thing about the build, diff the bytes, name the field.
Reconstructed by construction, not from documentation. Every row was fixed by building the same kernel with one thing changed and diffing the header bytes. Fields I did not manage to move are omitted.
PTX is not the cheap option
The standard advice is to always ship PTX for the newest virtual architecture so future GPUs have something to JIT. The advice generally is right, and the reason people give for it is usually wrong.
PTX is not a compact representation. For the ten-kernel workload the PTX is 37.3 KB against a 35.9 KB Hopper cubin, and on Blackwell it is 48.0 KB against 57.8 KB. It is the same order of magnitude in both directions, and after compression the gap narrows further because PTX is text and compresses beautifully.
What PTX buys is coverage of chips that do not exist yet, at the cost of a JIT compile on first use, per process, unless the driver’s compute cache is warm.
What it does not buy is what most people think: forward compatibility does not extend to the instructions anyone actually cares about, because PTX built for an architecture-conditional target is not forward compatible at all.
If your kernel needs wgmma or tcgen05, it needs sm_90a or sm_100a, and the compatibility story ends there. The f targets exist precisely because that cliff was too steep, and the Mercury capsule is how they made it less steep.
That extended header is the good part. The sm_100 entry above has a 112-byte header instead of 64, and the extra 48 bytes are a length-prefixed copy of the .nv.compat capability block from inside the cubin.
header bytes at offset 0x40:
48 00 00 00 20 00 00 00 -- payload at 0x48, length 0x20
02 09 00 00 02 02 01 00 03 0d 01 01 03 07 01 01
02 03 00 00 04 0b 08 00 09 00 00 00 00 00 00 00
^^ the same TLV stream that .nv.compat holds inside the ELFThe capability claim is duplicated into the index so that whatever picks an entry never has to open the ELF. That is a loader optimization, and it is also a small confirmation of what the compat block is for: it is the thing a resolver matches against a device before committing to a payload.
Compression is off by default for the part that matters
CUDA has a --compress-mode flag with values none, speed, balance, size and default. Running all five over the same two-cubin-plus-PTX fatbin gives a result worth knowing.
The flag values behave like two different algorithms rather than one algorithm at three settings: speed sets flag bit 13 and gets 504 bytes out of 922, while balance and size set bit 15 and get 400.
On the larger ten-kernel workload across twelve targets, size mode takes 537 KB down to 75 KB, an 86% reduction, for no measurable build time cost at this scale.
If you ship CUDA binaries and have not tried it, that is the cheapest win in this article.
How it gets into your executable
The host side is one of the few places NVIDIA actually publishes the format, in fatbinary_section.h in the toolkit include directory.
#define FATBINC_MAGIC 0x466243B1
#define FATBINC_VERSION 1
typedef struct {
int magic;
int version;
const unsigned long long* data;
void *filename_or_fatbins;
} __fatBinC_Wrapper_t;
#define FATBIN_CONTROL_SECTION_NAME “.nvFatBinSegment”
#define FATBIN_DATA_SECTION_NAME “.nv_fatbin”Compile a .cu file to an object and you get exactly that: a 24-byte wrapper in .nvFatBinSegment holding the magic and a relocation pointing at the fatbin bytes in .nv_fatbin, plus a section called __nv_module_id, plus a constructor.
$ readelf -x .nvFatBinSegment saxpy.o
0x00000000 b1436246 01000000 00000000 00000000 .CbF............
$ readelf -rW saxpy.o | grep nvFatBinSegment
R_X86_64_64 .nv_fatbin + 0
$ cat saxpy.cudafe1.stub.c # generated by nvcc --keep
static void __nv_cudaEntityRegisterCallback(void **__T0) {
__cudaRegisterEntry(__T0, (void(*)(int,float,const float*,float*))saxpy,
_Z5saxpyifPKfPf, (-1)); }
static void __sti____cudaRegisterAll(void) __attribute__((__constructor__));
static void __sti____cudaRegisterAll(void) {
__cudaRegisterBinary(__nv_cudaEntityRegisterCallback); }
void __device_stub__Z5saxpyifPKfPf(int p0, float p1, const float *p2, float *p3){
__cudaLaunchPrologue(4);
__cudaSetupArgSimple(p0, 0UL); __cudaSetupArgSimple(p1, 4UL);
__cudaSetupArgSimple(p2, 8UL); __cudaSetupArgSimple(p3, 16UL);
__cudaLaunch((char *)saxpy); }So a CUDA program registers its device code from an ELF constructor, before main. The argument offsets in the stub, 0, 4, 8 and 16, are the same offsets the cubin declared in EIATTR_KPARAM_INFO.
Host and device agree on the layout because both were generated from the same front end, and the binary carries the agreement in two places.
Two variants are worth noting because they trip people up. With -rdc=true, the device code is relocatable, the cubin becomes ET_REL with .rela.text.* sections and an EIATTR_EXTERNS attribute, and it lands in a differently named section, __nv_relfatbin, until the device link step turns it into a normal .nv_fatbin. The cost shows up in the machine code, and it is not subtle:
$ nvcc -arch=sm_90 -rdc=true -c a.cu b.cu && nvcc -arch=sm_90 -dlink -o dl.o a.o b.o
$ cuobjdump -sass dl.o
Function : _Z5applyPfPKffi
/*0100*/ CALL.ABS.NOINC 0x0 ;
Function : _Z5scaleff
/*0000*/ FFMA R4, R4, R5, 1 ;
/*0010*/ RET.ABS.NODEC R20 0x0 ;A one-instruction device function became a real call, with a stack frame, a return address register and a separate .text section, where whole-program compilation would have inlined it into a single FFMA.
That is what -dlto exists to undo: with LTO the fatbin entry is kind 8 and the instruction selection is deferred to link time so the inlining can happen after all translation units are visible.
In my sandbox the LTO device link fails for want of a driver-side link library, so I can report the container shape (1,944 compressed bytes expanding to 2,548) but not the linked output.
What the driver does with all this
The loading path is the reason the format looks the way it does, and it is worth stating plainly because the layering is easy to lose. The constructor calls __cudaRegisterBinary with a pointer to the wrapper.
The runtime hands the fatbin to the driver, which reads the index and selects one payload for the device it is about to use: an exact-architecture cubin if there is one, otherwise a compatible one, otherwise PTX to be compiled on the spot. The selected payload becomes a module.
The module’s .nv.info tells the driver how many registers each kernel needs, how much shared memory to reserve, how big a stack to allocate given the call graph, and where in constant bank 0 to write the arguments. Only then can a launch happen, and a launch is essentially a memcpy of the argument block into the frame the cubin described, followed by a grid dispatch.
Everything in the binary that is not machine code exists to let that sequence happen without the driver consulting anything else. That is why the parameter layout is in the object rather than in a header, why the exit offsets are precomputed, why the capability claim is duplicated into the fatbin index, and why the reserved shared memory has a symbol. A loader that had to infer any of it would be slower and more fragile.
Two changes to this path are worth knowing. Since CUDA 12.0 there is a second module abstraction, the library and kernel API (cuLibraryLoadData, cuKernel), which decouples a loaded artifact from a specific context, and it exists largely because the old model made lazy loading awkward.
And since 12.2 on Linux and 12.3 everywhere, lazy loading is the default: modules load on first use of a symbol from them, and kernels within a module load on first cuModuleGetFunction.
The requirement is a runtime of 11.7 or later, statically linked into whatever built the library, which means an old dependency in your stack can quietly opt you out. NVIDIA’s own documentation puts the requirement plainly: libraries compiled against older runtimes load all modules eagerly.
What the format costs, in bytes and seconds
All of the above has a price, and it is paid by everyone who ships a wheel.
This is why PyTorch’s release engineering discussions read the way they do. When the CUDA 12.8 build matrix had to add Blackwell, the team explicitly rejected keeping the older architectures because of binary size. The trade is not subtle: adding a target adds a copy of every kernel in the library, and machine learning libraries have a great many kernels.
You can see the endpoint of that logic by taking apart what NVIDIA itself ships. I pulled the cuBLAS 13.6.0.2 wheel from PyPI and parsed the .nv_fatbin sections directly.
Two libraries, 6,860 fatbin entries, 178 MB on disk expanding to 1.8 GB of device code. And every single entry is compressed, which tells you that NVIDIA does not ship its own libraries with the default compression mode.
The right-hand panel is worth sitting with. In cuBLASLt, 101 MB is host code, and 78 MB is a section called .cask_resource.
One detail worth noting before anyone adds those numbers up: the library also declares 99 MB of .bss and .lbss, and those are NOBITS sections, which means they cost 99 MB of memory at run time and zero bytes on disk.
Counting only the sections that actually occupy file space gets you 491.4 MB of the 492.7 MB file, and the rest is headers and padding. I cannot tell you what CASK stands for, because NVIDIA has never said, and the name only surfaces publicly in cuDNN and TensorRT diagnostics.
What I can tell you is that the strings in that section include 2,429 mangled C++ names beginning N11cask, so it is a namespace, and that the rest of the section is a catalog. Pull strings out of it and you get entries like this:
cutlass3x_sm103_bstensorop_s256x256x96gemm_block_scaled_ue4m3xe2m1_ue4m3xe2m1_
f32_bf16_ue4m3xe2m1_256x256x768_0_tnn_align...A CUTLASS 3.x kernel for sm_103 doing a block-scaled GEMM with FP8 and FP4 operand types at a 256x256x96 tile.
Counting distinct kernel-shaped names across the resource and read-only sections gives 25,401, with the generation prefixes still visible in the histogram: 5,329 mentioning sm90, 5,015 cutlass3x, 1,956 ampere, 808 volta, 274 turing.
That number is the moat, expressed as an artifact count rather than an argument. Not one clever kernel. Twenty-five thousand parameterised ones, plus 78 MB of metadata describing which to pick, plus 101 MB of host code doing the picking. A competitor can match the code generator.
Matching the catalog and the selection heuristics is a different kind of work, and it is the kind that does not benefit much from being smart.
NVIDIA is visibly feeling the weight of it. cuDNN 9.5 shipped a build configuration called GRAPH_JIT_ONLY that keeps the runtime kernel generation engines and drops the precompiled ones specifically to cut binary size, which is the vendor arriving at the same trade its users have been making by hand: generate at run time or ship the archive.
It also explains why lazy loading became the default. If a process eagerly loaded all 5,309 cubins in cuBLASLt it would spend its startup budget on kernels it will never call.
Lazy loading arrived as opt-in in CUDA 11.7, became the Linux default in 12.2 and the universal default in 12.3, and defers module and kernel loading until first use.
In a world where a single library holds thousands of independently loadable objects, that stops being an optimization and becomes a requirement.
Debug information is not free either
For the same ten kernels on sm_90: 35.9 KB release, 75.6 KB with -lineinfo, 205.9 KB with -G. The last one is not just extra sections.
It changes code generation, taking .text from 16.0 KB to 57.5 KB, because full device debug turns off most optimization. -lineinfo is the option you want in production builds: it doubles the cubin and leaves the code alone.
Why any of this matters if you are not writing a compiler
Three reasons, in increasing order of how much I care about them.
The first is operational. Binary size and load time are real costs in inference deployments, container images and cold starts, and they are almost entirely determined by choices in the format: how many targets, PTX or not, which compression mode, lazy or eager. Those are four flags. Most teams have never touched them.
The second is diagnostic. A cubin tells you its register count, its shared memory usage, its stack frame, its parameter layout, the toolkit that built it and the flags it was built with, and you can read all of that out of a shipped wheel without running anything.
If you are trying to work out why somebody else’s kernel spills, or which architectures a vendor actually optimized for versus which they merely support, the binary is a better source than the changelog. The -res-usage flag on cuobjdump is the fastest occupancy audit I know.
Auditing somebody else’s wheel
cuobjdump -lelf libfoo.so # which architectures, how many cubins
cuobjdump -lptx libfoo.so # is there PTX for anything newer
cuobjdump -res-usage libfoo.so # registers, shared memory, stack, per kernel
cuobjdump -elf libfoo.so | grep -A5 “Toolkit Information” # which ptxas built it
readelf -SW libfoo.so | grep nv_fatbin # how much of the file is device codeFive commands, no GPU, no source. The last one is the one that surprises people: if .nv_fatbin is 85% of a library, you are shipping a kernel archive with a thin API on top, and any conversation about image size has to start there.
The third is strategic, and it is why I went looking in the first place.
There is a long line of work reverse-engineering these formats: decuda for the earliest chips, asfermi for Fermi, KeplerAs, Scott Gray’s maxas for Maxwell (the toolchain behind the fast convolution kernels of the deep learning era), turingas, CuAssembler for everything through Ampere, NVBit for dynamic instrumentation, and now work lifting SASS back into typed compiler IR.
Every one of these projects had to rediscover the container before it could touch the code, and every one of them targets an architecture at least two generations behind current silicon. That lag is the moat’s actual shape.
It is not that the format is unknowable, instead it’s that knowing it is a full-time job that resets on a yearly cadence, and the number of people doing it is small enough to name.
Mercury is the part of this that I think changes the picture, and not in the direction I expected. Publishing an interface and keeping the lowering is NVIDIA’s oldest move: PTX is public, ptxas is not.
What the capsule does is insert a second, private layer at exactly the point where a competitor would want to attach, and it does it for a reason customers asked for: one binary that keeps working on the next chip in the family. The compatibility benefit is real.
The side effect is that the interesting representation of a Blackwell kernel is no longer the SASS you can disassemble, and the format that carries it has no public name, no documented encoding, and its own relocation types.
Four things I would change, and one I would not
Having read a few hundred of these files, I have opinions about the format that have nothing to do with the moat.
Publish the metadata, keep the encoding. There is no competitive information in EIATTR_KPARAM_INFO. It is a calling convention. Documenting the .nv.info attribute codes and the fatbin entry header would cost NVIDIA nothing and would stop every profiler, virtualization layer, checkpointer and build tool from rediscovering the same twenty structs. The instruction encoding is a different argument and I understand the refusal. The container is not.
Make size the default compression mode. The current default compresses PTX and leaves the SASS raw, which is the wrong way round: SASS is the bulk and the part that never gets read by a human. NVIDIA already ships its own libraries with everything compressed. Every framework wheel on PyPI is paying for a default its vendor does not use.
Give the toolkit a way to strip PTX. There is nvprune for removing architectures, but the common case in a container build is “I know exactly which GPUs this image runs on, remove the JIT fallback.” Doing that today means unpacking and repacking fatbins with nvFatbin, which is a strange amount of work for what should be a link flag.
Version the format visibly. Between the ABI version byte, the OS ABI byte that changed value, the toolkit note, the API version attribute and the Mercury ISA version, there are five different notions of “which format is this” in one file, and no single field that a tool can check to know whether it will understand what follows. A cubin that a future tool cannot parse should say so in its first sixteen bytes.
What I would not change personally is the decision to make the binary self-describing rather than pushing the information into headers or a side-channel.
It is the reason NVBit can instrument a kernel it has never seen, the reason cuobjdump -res-usage works on a stranger’s wheel, and the reason this article was possible on a machine with no GPU in it. Verbose, redundant, self-describing binaries are a gift to everyone downstream, including the people trying to work out what you shipped.
Most of what is wrong with the CUDA binary format is that it is undocumented. Very little of it is that it is badly designed.
Where I might be wrong
The capsule might be much less than I think. If it turns out to be a small fixup table rather than a re-encodable program, then “second copy of every kernel” is the wrong phrase and the right one is “compatibility annotations.” The sublinear size scaling in the capsule plot above is evidence for the smaller reading; the existence of a section type called
CUDA_MERCURY_SASS_MAPis evidence for the larger one. I could not settle it.I have no Blackwell hardware. Nothing here was executed. A claim about what a driver does with a capsule is, from my position, a claim about what an artifact appears designed for. Everything about run time behaviour, including JIT cost and the actual effect of lazy loading, is cited rather than measured.
The OS ABI byte is a mess and I may be reading the mess wrong. LLVM lists 51 and 41; cubins carry 65. My reading is that the second constant is
0x41written as decimal, but it is equally possible NVIDIA changed the value again and LLVM has not caught up. I only checked 13.3 output.The fatbin flag bits are inferred from a handful of data points. I set them by construction using the compression modes and the LTO path; other bits in that word certainly mean things I did not exercise. Byte 0 of
e_flags, which flips from 4 to 2 at Blackwell, I could not explain at all.My library numbers are one version of one vendor library. cuBLAS is the extreme case for kernel count. Generalizing “27% device code” to other libraries would be wrong.
The feature-tier experiment used one instruction.
tcgen05.fencebehaves as described, and I have no reason to think the mechanism differs for other family-specific instructions, but one probe is one probe.
Predictions
By the end of 2027, at least one open-source project will publish a working parser for the Mercury capsule format, and it will be motivated by GPU virtualization or checkpointing rather than by performance.
NVIDIA will not document the capsule encoding, on the same terms it has never documented SASS encodings, through at least CUDA 15.
By the end of 2027,
-compress-mode=sizeor its equivalent will be the default in at least one major framework’s release build, driven by wheel size limits rather than by anyone reading the flag documentation.The next architecture after Blackwell will ship a SASS encoding with statically paired instructions, and
nvdisasm‘s VLIW syntax will be why we find out.Family-conditional targets will absorb the architecture-conditional ones in practice: by the end of 2028, the fraction of shipped cubins carrying
asuffixes will fall as libraries move tof, because one binary per family is a better deal than one per chip.
Dossier
Tier A, measured directly and reproducible from the appendix.
ELF header field values and the
e_flagsarchitecture encoding, including the absence of the virtual architecture from that word.Section inventories and sizes for all twelve targets.
The
.nv.infoand.nv.compatattribute contents as decoded bycuobjdump, including theISA_CLASStransition when a family-specific instruction is used.The rejection of
tcgen05.fenceonsm_100andsm_120aand its acceptance onsm_100f,sm_100aandsm_103f.The 113 attribute names and the 119 relocation names. Parameter base offsets and constant bank sizes.
Control field bit positions and their per-kernel distribution.
Fatbin entry field offsets, compression flag bits and all size and timing measurements.
The
CALL.ABS.NOINCcost of relocatable device code.The cuBLAS section and entry statistics, the 25,401 kernel names and the 2,429
caskmangled names.The presence, sizes and scaling of the Mercury sections.
Tier B, documented by NVIDIA, in shipped headers, or in third-party source.
The fatbin wrapper struct and section names (
fatbinary_section.h).The family and architecture-conditional target semantics, including the statement that
sm_100andsm_100fare aliases absent family features (Programming Guide, PTX ISA, the CUDA 12.9 blog post).Lazy loading history and defaults.
Dropped architecture support in CUDA 13.0 and the Thor renumbering.
EM_CUDA,ELFOSABI_CUDAandELFOSABI_CUDA_V2in LLVM.cuDNN’s
GRAPH_JIT_ONLYconfiguration.The structure of the control field, which is documented in the literature rather than by the vendor.
Tier C, inferred from artifacts.
Mercury is a re-finalizable intermediate representation and that the capsule exists to serve family-conditional compatibility.
That the duplicated compat block in the fatbin entry header is a loader fast path.
ptxas‘s hidden finalizer options correspond to a driver-side code path..cask_resourceis a kernel selection catalog.That the
ELFOSABI_CUDA_V2constant is a hexadecimal value written as decimal.
Tier D, speculation, flagged as such in the text.
--no-vliwimplies a future paired-issue ISA.EIATTR_COROUTINE_RESUME_ID_OFFSETSpoints at an unannounced feature.sm_88is the Nintendo Switch 2, which comes from community references and not from anything I could check.
Reproduce this
No GPU required. On x86-64 Linux:
pip install nvidia-cuda-nvcc==13.3.73 nvidia-cuda-cuobjdump==13.3.73 \
nvidia-cuda-nvdisasm==13.3.73 nvidia-cuda-crt==13.3.73 \
nvidia-cuda-runtime nvidia-cuda-cccl
export PATH=$(python3 -c “import sysconfig,os;print(os.path.join(sysconfig.get_paths()[’purelib’],’nvidia/cu13/bin’))”):$PATH
cat > saxpy.cu <<’EOF’
#include <cuda_runtime.h>
__global__ void saxpy(int n, float a, const float* x, float* y) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) y[i] = a * x[i] + y[i];
}
EOF
# the container and its metadata
nvcc -arch=sm_90 -cubin -o k90.cubin saxpy.cu
nvcc -arch=sm_100 -cubin -o k100.cubin saxpy.cu
cuobjdump -elf k90.cubin | less
diff <(cuobjdump -elf k90.cubin | sed -n ‘/^Sections:/,/^$/p’) \
<(cuobjdump -elf k100.cubin | sed -n ‘/^Sections:/,/^$/p’)
# the capability claim, and the accelerator bit
for a in sm_100 sm_100a sm_100f; do nvcc -arch=$a -cubin -o c.cubin saxpy.cu
echo “== $a”; cuobjdump -elf c.cubin | grep -A1 EICOMPAT | grep -E “Attr|Value”; done
# the attribute namespace and the finalizer options
strings -a $(command -v nvdisasm) | grep -o “EIATTR_[A-Z0-9_]*” | sort -u
strings -a $(command -v ptxas) | grep -iE “capmerc|finaliz|R_MERCURY”
# the feature tiers, as an experiment rather than a diagram
cat > tc.cu <<’EOF’
#include <cuda_runtime.h>
__global__ void k(unsigned* o){
asm volatile(”tcgen05.fence::before_thread_sync;” ::: “memory”);
o[0]=1;
}
EOF
for a in sm_100 sm_100f sm_100a sm_103f sm_120a; do printf “%-9s “ $a
nvcc -arch=$a -cubin -o t.cubin tc.cu 2>&1 | head -1; done
for a in sm_100f sm_100a; do nvcc -arch=$a -cubin -o t_$a.cubin tc.cu
echo “== $a”; cuobjdump -elf t_$a.cubin | grep -A2 EICOMPAT | grep -E “Attribute|Value”; done
# is the virtual architecture in e_flags? (no)
for v in 75 80 90; do nvcc -gencode arch=compute_$v,code=sm_90 -cubin -o v.cubin saxpy.cu
python3 -c “import struct;print(’compute_$v -> 0x%08x’%struct.unpack_from(’<I’,open(’v.cubin’,’rb’).read(),48)[0])”
cuobjdump -elf v.cubin | grep “Virtual SM”; done
# the field names nvdisasm knows
strings -a $(command -v nvdisasm) | grep -o “EF_CUDA_[A-Z0-9_]*” | sort -u
# scheduling control bits, 21 bits starting at 105 of each 128-bit word
nvdisasm -c -hex k90.cubin | head -40
python3 - <<’PY’
import struct
d=open(’k90.cubin’,’rb’).read()
o,=struct.unpack_from(’<Q’,d,0x28); es,n,si=struct.unpack_from(’<HHH’,d,0x3a)
sh=[struct.unpack_from(’<IIQQQQIIQQ’,d,o+i*es) for i in range(n)]
st=sh[si][4]; nm=lambda x: d[st+x:d.index(b’\0’,st+x)].decode()
for s in sh:
if not nm(s[0]).startswith(’.text.’): continue
b=d[s[4]:s[4]+s[5]]
for i in range(0,len(b),16):
lo,hi=struct.unpack_from(’<QQ’,b,i); w=(hi<<64)|lo
print(f”{i:04x} stall={(w>>105)&0xf} yield={(w>>109)&1} “
f”wr={(w>>110)&7} rd={(w>>113)&7} wait={(w>>116)&0x3f:06b} “
f”reuse={(w>>122)&0xf:04b}”)
PY
# the fatbin container and the cost of coverage
nvcc -gencode arch=compute_90,code=sm_90 \
-gencode arch=compute_100,code=[sm_100,compute_100] -fatbin -o f.fatbin saxpy.cu
for m in none speed balance size default; do
nvcc -compress-mode=$m -gencode arch=compute_90,code=sm_90 \
-gencode arch=compute_100,code=[sm_100,compute_100] \
-fatbin -o f_$m.fatbin saxpy.cu; ls -l f_$m.fatbin; done
# the host side
nvcc -arch=sm_90 -c -o k.o saxpy.cu
readelf -SW k.o | grep -E “nv_fatbin|nvFatBinSegment|module_id”
readelf -x .nvFatBinSegment k.o
nvcc -arch=sm_90 --keep -c -o /dev/null saxpy.cu && cat saxpy.cudafe1.stub.cThe fatbin entry parser used for the library dissection is thirty lines: walk from the 0xBA55ED50 header, then for each entry read the kind at offset 0, header size at 4, padded payload size at 8, compressed size at 0x10, architecture at 0x1c, flags at 0x28, and skip header_size + payload_size to the next one.
