How to find out which instructions a new GPU accepts?

At Latent One, we're deeply curious about the hardware the current AI models run on, the curiosity is obsessive enough to let us go all the way down to single assembly instructions. So we asked a plain question about the GB10, which instructions does its compiler accept, and does the answer change with the exact chip target you compile for?

Part 1#

How do you check which instructions a GPU compiler target accepts?#

We wanted to know which PTX instructions the compiler accepts for the GB10, so we wrote one tiny PTX file per instruction and asked ptxas to compile each one for two targets. That was 976 compiles in 8.9 seconds, and nothing ran on the GPU. The compiler accepted 406 of 481 instruction forms for sm_121 and 409 for sm_121a, and the three it only accepted for sm_121a were setmaxnreg.inc, setmaxnreg.dec and one ldmatrix form.

NVIDIA publishes long instruction lists, and a list says nothing about the exact chip target you compile for. So we wanted a quick way to ask the compiler, one instruction at a time, and to repeat the question for every target.

PTX is the portable text form CUDA compiles to, and ptxas is the tool that turns PTX into machine code for one chip. A target such as sm_121 names the chip, and the a suffix in sm_121a asks for features specific to that chip.

If ptxas accepts an instruction for a target, the compiler allows it there. If it rejects the instruction, you get an error message and the check stops.

The check#

Each probe is a PTX file with a .target line and a single instruction inside an empty kernel. The image shows the one for setmaxnreg.inc, followed by the command that compiles it and throws the output away.

A PTX probe for setmaxnreg.inc: .version 9.0, .target sm_121a and .address_size 64, then an empty kernel test_kernel holding setmaxnreg.inc.sync.aligned.u32 64 and ret, followed by the command ptxas -arch=sm_121a probe.ptx -o /dev/null.

Part 2#

SASSquatch, the open-source auditor we used, generates 488 of these probes and runs the same command for each target. We ran it with the ptxas from CUDA 13.0 on sm_121 and sm_121a, which is 976 compiles that finished in 8.9 seconds.

What it found?#

The compiler accepted 406 of 481 instruction forms for sm_121 and 409 for sm_121a. So the chip-specific target added exactly three. setmaxnreg.inc and setmaxnreg.dec let a warp group raise or lower its register budget while it runs, and the third is an ldmatrix form (m8n16.x4, b4x16_p64) that loads packed 4-bit values from shared memory.

Which instructions ptxas accepts on GB10: sm_121 accepted 406 of 481 instruction forms and sm_121a accepted 409 of 481. Only sm_121a accepted three: setmaxnreg.inc.sync.aligned.u32, setmaxnreg.dec.sync.aligned.u32 and an ldmatrix m8n16.x4 b4x16_p64 form.

Part 3#

Matrix multiply and block-scaled matrix multiply were accepted in full (8 of 8 each), and so were the TMA copies (7 of 7). The one group that failed completely was tensor memory, where both tcgen05 probes were rejected on sm_121.

Rejection is information in disguise!#

A rejection says what the compiler took, which is a different thing from what the chip can do. Some misses look like our probe wrote the instruction in a form ptxas does not take, for example fma.f16 with no rounding mode, which PTX requires for half precision. The accepted instructions were also never executed in this pass, so acceptance only says the compiler allows them.

We found out afterwards that these numbers are already public. The author of SASSquatch posted the same result on the NVIDIA developer forum in February, with 409 forms compiling and 72 failing, and public GitHub threads already say tcgen05 does not exist on sm_121. So this was a repeat of that check on our own chip with the ptxas from CUDA 13.0. A public report also shows an FP4 conversion that sm_121 rejects and sm_121a accepts, and it is not among our three, so those three may depend on the ptxas version.

A compile loop tells you which instruction forms ptxas accepts for a target, and it does that in seconds. We have not sorted the 75 rejections on sm_121 into forms the chip lacks and forms we wrote wrongly, and we have not run any of the accepted instructions on the GPU.

SASSquatch can also read the machine code ptxas builds and patch an opcode in a compiled kernel to run it on the live GPU. We think comparing what the hardware runs with what the compiler emits could point to compiler optimizations, and we have not tried it.