C++/CUDA → PTX (NVIDIA's portable IR) → SASS (the actual machine code that runs on the SM). Walk through both ISAs — instruction families, the FFMA/HMMA/WGMMA tensor-core instructions, LDMATRIX/STMATRIX, BAR.SYNC, predication — and how to read SASS to understand what the compiler did to your kernel.
Two ISAs, one toolchain. PTX is the portable IR you write or that nvcc emits; SASS is the per-architecture machine code that actually executes on the streaming multiprocessor. Reading either one tells you what the compiler really did with your kernel.
Between the CUDA C++ you write and the silicon that runs it sit two intermediate forms. PTX is a portable text IR; SASS is the binary machine code per architecture. The boundary between them is ptxas, NVIDIA's optimising assembler.
One PTX file targets a logical architecture (sm_90). At runtime, if the host driver finds no embedded SASS for the present GPU, it JITs PTX on the fly. That's how a binary built today runs on a future GPU class.
Encoded for a specific compute capability. sm_75 SASS will not run on Hopper; sm_90 SASS will not run on Turing. Fat binaries embed several SASS blobs plus a PTX fallback.
cuobjdump --dump-elf-symbols my_app lists every embedded SASS blob and the PTX fallback. cuobjdump --list-text is shorter; cuobjdump --dump-sass --gpu-architecture=sm_90 my_app prints just the Hopper SASS.
PTX (Parallel Thread Execution) is NVIDIA's published, stable, ASCII-text intermediate representation. It looks like a typed assembly. It's the only ISA NVIDIA fully documents — the PTX ISA Reference is the official contract between compiler authors and the driver.
Every instruction names its operand type. .s32, .u64, .f16, .f16x2, .bf16, .b128. Mismatched types are a PTX-level error, not a runtime crash.
.reg (virtual registers), .local (per-thread stack), .shared (per-block scratchpad), .global (DRAM), .const, .param, .tex. Every load and store is tagged with its address space.
Each .reg result is a fresh symbolic name. ptxas does the real register allocation later. PTX itself has no notion of "R7" — that's a SASS thing.
add, mul, mad, fma, ld, st, cvt, setp, selp, bra, bar.sync, mma.sync, wgmma.mma_async.sync, cp.async. Modifiers (.lo, .hi, .sat, .rn) tune behaviour.
.version 8.3
.target sm_90
.address_size 64
.visible .entry block_reduce(
.param .u64 in_ptr,
.param .u64 out_ptr )
{
.reg .b64 %rd<4>;
.reg .b32 %r<8>;
.reg .f32 %f<6>;
.reg .pred %p<3>;
.shared .align 4 .b8 sdata[128]; // 32 floats
ld.param.u64 %rd1, [in_ptr];
cvta.to.global.u64 %rd1, %rd1;
mov.u32 %r1, %tid.x;
mul.wide.u32 %rd2, %r1, 4;
add.s64 %rd2, %rd1, %rd2;
ld.global.f32 %f1, [%rd2]; // my element
st.shared.f32 [sdata + %r1*4], %f1;
bar.sync 0; // __syncthreads()
// log-step reduction, stride 16, 8, 4, 2, 1
mov.u32 %r2, 16;
$LOOP:
setp.lt.u32 %p1, %r1, %r2;
@!%p1 bra $SKIP;
add.s32 %r3, %r1, %r2;
ld.shared.f32 %f2, [sdata + %r1*4];
ld.shared.f32 %f3, [sdata + %r3*4];
add.f32 %f2, %f2, %f3;
st.shared.f32 [sdata + %r1*4], %f2;
$SKIP:
bar.sync 0;
shr.u32 %r2, %r2, 1;
setp.ne.u32 %p2, %r2, 0;
@%p2 bra $LOOP;
setp.ne.u32 %p1, %r1, 0;
@%p1 bra $DONE;
ld.param.u64 %rd3, [out_ptr];
cvta.to.global.u64 %rd3, %rd3;
ld.shared.f32 %f4, [sdata];
st.global.f32 [%rd3], %f4;
$DONE:
ret;
}
NVIDIA refuses to commit SASS to a stable encoding — they want the freedom to change it every generation. PTX gives compiler vendors (LLVM, Triton, JAX, Mojo) a stable target while preserving NVIDIA's ability to redesign the SM. It's the same idea as JVM bytecode, only tuned for a SIMT machine.
SASS (Streaming ASSembler, sometimes Source ASSembly) is the per-architecture binary that runs on the SM. Each generation has its own encoding, its own instruction list, and its own quirks. NVIDIA does not publish a complete SASS reference. What we know comes from years of community reverse engineering.
| Arch | SM | Notable SASS additions |
|---|---|---|
| Volta | sm_70 | First HMMA (8.8.4 FP16); independent thread scheduling. |
| Turing | sm_75 | HMMA INT8/INT4; uniform datapath (URF, UR registers). |
| Ampere | sm_80 / sm_86 / sm_87 | HMMA m16n8k16 FP16/BF16, m16n8k8 TF32; LDGSTS (cp.async, async global→shared); LDMATRIX; sparsity 2:4. |
| Ada Lovelace | sm_89 | FP8 HMMA on tensor cores; otherwise same SASS family as Ampere. |
| Hopper | sm_90 | WGMMA (warp-group async MMA); STMATRIX; UTMALDG/UTMASTG (TMA). |
| Blackwell | sm_100 / sm_120 | New tcgen05.mma SM-level tensor-core family with FP4 (MX-FP4 & NVFP4) / FP6 / FP8; second-gen TMA. |
cuobjdump --dump-sass on a cubin or fat binary; nvdisasm kernel.cubin for richer output (control-flow graph with -cfg, scheduling info with -c).
OP[.modifier] dest, src1, src2 ; — left-to-right destination first. Most ops take 3 register sources; a few take immediates or constant-bank refs c[0x0][0x28].
@P0 ADD ... runs only on threads where P0 is true. @!P0 is the negation. @PT means "always" (PT is hard-wired true). Every SASS line has an implicit predicate.
Community projects: Maxas (Maxwell), NervanaSys, CuAssembler (Volta→Hopper). Used to extract instruction encodings, latency tables, and schedule-control bits.
// CUDA: __global__ void add(float* a, float* b, float* c) {
// int i = threadIdx.x; c[i] = a[i] + b[i];
// }
MOV R1, c[0x0][0x28] ; // stack base
S2R R0, SR_TID.X ; // threadIdx.x
ULDC.64 UR4, c[0x0][0x208] ; // pointer a (uniform)
IMAD.WIDE R2, R0, 0x4, UR4 ; // &a[i] (64-bit)
LDG.E R5, [R2.64] ; // a[i]
LDG.E R7, [R2.64+0x800000] ; // b[i] (relative)
FADD R5, R5, R7 ;
STG.E [R2.64+0x1000000], R5 ;
@PT EXIT ;
SASS is a 3-operand RISC-ish ISA with predication, fixed 32-thread warps and constant-memory operands. If you can read ARM or RISC-V assembly, you can read SASS — the unfamiliar bits are the addressing modifiers (.E = extended/64-bit address form on LDG/STG) and the constant banks (c[0x0][..]).
Underneath all the high-level matmul libraries (cuBLAS, cuDNN, CUTLASS, TRT-LLM) sit three families of tensor-core instructions. Knowing which one you're emitting tells you exactly what hardware path your kernel uses.
| Family | Generations | Shapes (M.N.K) | Operand types | Notes |
|---|---|---|---|---|
| HMMA | Volta → Ampere → Ada | 8.8.4, 16.8.8, 16.8.16 | FP16, BF16, TF32, FP8 (Ada) | Hardware MMA. Sync, warp-scoped. |
| WGMMA | Hopper (sm_90+) | m64.N.K, N up to 256; K=16 (FP16/BF16), K=8 (TF32), K=32 (FP8/INT8) | FP8 (E4M3/E5M2), BF16, TF32, FP16, INT8 | Warp-group async MMA. Operands in shared memory. |
| WMMA | C++/PTX abstraction | maps to HMMA | same as HMMA | Frontend wrapper; not a SASS opcode. |
| IMMA | Turing → Hopper | 8.8.16, 16.8.32 | INT8, INT4 (Turing/Ampere) | Integer tensor-core MMA. |
| BMMA | Turing → Ampere | 8.8.128 | 1-bit operands | Binary MMA; XNOR + popcount. |
| tcgen05.mma | Blackwell (sm_100 / sm_120) | SM-level tile shapes | FP4 (MX-FP4 + NVFP4), FP6, FP8, BF16, FP16 | New SM-level tensor-core ops; introduces FP4 microscaling (MX-FP4 32-element E8M0; NVFP4 16-element E4M3 + per-tensor FP32). |
// 64x128 = 64x16 * 16x128 in one async warp-group MMA
wgmma.mma_async.sync.aligned.m64n128k16.f32.bf16.bf16
{%f0, %f1, %f2, %f3, %f4, %f5, %f6, %f7,
%f8, %f9, %f10, %f11, %f12, %f13, %f14, %f15,
%f16, %f17, %f18, %f19, %f20, %f21, %f22, %f23,
%f24, %f25, %f26, %f27, %f28, %f29, %f30, %f31,
%f32, %f33, %f34, %f35, %f36, %f37, %f38, %f39,
%f40, %f41, %f42, %f43, %f44, %f45, %f46, %f47,
%f48, %f49, %f50, %f51, %f52, %f53, %f54, %f55,
%f56, %f57, %f58, %f59, %f60, %f61, %f62, %f63},
desc_a, desc_b, 1, 1, 1, 0, 0;
wgmma.mma_async — warp-group, asynchronous MMA.sync.aligned — all 128 threads of the warp-group must hit it.m64n128k16 — tile shape 64×128 = 64×16 × 16×128.f32.bf16.bf16 — accum / A-input / B-input typesdesc_a, desc_b — 64-bit shared-memory descriptors (not raw pointers)WGMMA requires its A and B operands to live in shared memory in a specific swizzled layout. Instead of pointers it takes 64-bit descriptors that encode the base address, the leading-dim stride and the swizzle mode. That's why every Hopper matmul kernel uses LDMATRIX or TMA to set up the tiles first.
HMMA on Ampere is synchronous: the warp issues, then waits for the result registers. WGMMA on Hopper is non-blocking — it runs while the warp keeps doing other work, and you reap the accumulators later with wgmma.commit_group + wgmma.wait_group. That's how Hopper kernels overlap MMA with TMA loads.
Tensor cores need their inputs in a very specific register layout: each thread holds two or four FP16 values from a non-obvious slice of the tile. Doing that by hand — element-by-element — cost as much as the matmul itself. ldmatrix fixes this.
ldmatrix.sync.aligned.m8n8.x4.shared.b16 loads four 8×8 FP16 tiles from shared memory into registers, distributing each tile across 32 threads in exactly the layout HMMA / WGMMA expects.
The reverse op — takes the tensor-core accumulator layout in registers and writes it back to shared memory in tile order. Used to drain WGMMA results before a global store.
CUTLASS, cuBLAS, cuDNN and TRT-LLM all build their hot path around ldmatrix + mma. Without it, hand-shuffling 16-bit half-words between threads dominates the kernel; with it, the loads vanish into the same warp-scheduler slot as the FFMAs they feed.
SIMT only takes you so far. As soon as threads share data through shared memory or async copies, you need explicit fences. PTX exposes a small zoo of them, each with very specific semantics.
| Instruction | Scope | What it does |
|---|---|---|
bar.sync 0 | Block | All threads hit this barrier; classic __syncthreads(). |
bar.arrive | Block / warp | Non-blocking arrival; producer warps signal readiness without stalling. |
wgmma.fence / wgmma.commit_group / wgmma.wait_group | Warp group (4 warps = 128 threads) | Hopper warpgroup MMA pipelining: fence registers, commit issued WGMMAs into a group, then wait until all but N groups complete. |
mbarrier.init | Block | Set up an mbarrier with a thread arrival count. |
mbarrier.arrive | Block | Drop a token into the mbarrier; non-blocking. |
mbarrier.try_wait | Block | Block until barrier phase completes (or timeout). |
fence.sc | cta / gpu / sys | Sequentially-consistent fence; orders all loads and stores. |
fence.acq_rel | cta / gpu / sys | Acquire-release ordering (cheaper than SC). |
cp.async.commit_group | Warp | Close out a group of in-flight cp.async loads. |
cp.async.wait_group N | Warp | Wait until at most N async groups remain in flight. |
Hopper's TMA + WGMMA design needs asynchronous producers and consumers. mbarrier primitives + wgmma.commit_group/wgmma.wait_group let one warp group load tiles via TMA into shared memory while another warp group runs WGMMA on the previous tile — the canonical ping-pong used by CUTLASS Hopper kernels.
A barrier is a meeting point: thread N waits for thread M. A fence is a memory-ordering primitive: it doesn't wait for anyone, it only constrains the visible order of memory ops. Most bugs come from confusing the two.
.shared .align 8 .b64 mbar;
// once, by thread 0:
mbarrier.init.shared.b64 [mbar], 128; // 128 arrivals
// producer warp group: load a tile via TMA
cp.async.bulk.tensor.shared::cluster.global
[smem_dst], [tma_desc, {x, y}], [mbar];
// consumer warp group: wait then run WGMMA
mbarrier.try_wait.parity.shared.b64 %p, [mbar], %phase;
@!%p bra retry;
wgmma.mma_async.sync.aligned.m64n128k16.f32.bf16.bf16 ... ;
wgmma.commit_group.sync.aligned;
wgmma.wait_group.sync.aligned 1;
GPUs hate branches: a divergent branch inside a 32-thread warp forces both paths to execute, with the wrong-side lanes masked off. NVIDIA's first answer is to avoid branches entirely — predicate every instruction.
PTX exposes virtual .pred registers (declared by the kernel, e.g. %p0..); after register allocation, SASS uses the 7 architectural predicate regs P0..P6 plus the hard-wired-true PT. Comparison ops (setp, ISETP) write predicates; almost every other op can be guarded by one.
@%p0 add.s32 %r1, %r2, %r3 — only threads where p0 = true execute. The unselected lanes still consume the warp slot; the result write is suppressed.
@%p0 bra target in PTX, @P0 BRA target in SASS. If all lanes agree (uniform branch), a single warp jumps. If lanes disagree (divergent), the SM walks each side serially with a mask.
Pre-Volta, after a divergent branch the warp would only re-converge at a compiler-marked sync point. Volta+ tracks each thread's PC; lanes can re-converge naturally any time their PCs match again. That changed how races are written safely.
__global__ void pick(float* y, float a, float b, int* x) {
int i = threadIdx.x;
if (x[i] > 0)
y[i] = a;
else
y[i] = b;
}
S2R R0, SR_TID.X ;
ULDC.64 UR4, c[0x0][0x218] ; // ptr x
IMAD.WIDE R2, R0, 0x4, UR4 ;
LDG.E R5, [R2.64] ; // x[i]
ISETP.GT.AND P0, PT, R5, RZ, PT ; // P0 = (x[i] > 0)
MOV R7, c[0x0][0x210] ; // b
@P0 MOV R7, c[0x0][0x20c] ; // a (overrides if P0)
ULDC.64 UR4, c[0x0][0x208] ; // ptr y
IMAD.WIDE R2, R0, 0x4, UR4 ;
STG.E [R2.64], R7 ;
@PT EXIT ;
The CUDA author wrote a branch; the SASS has none. The compiler used ISETP + a predicated MOV over the "else" value to keep the warp dense. This is the preferred shape for short, balanced conditionals — you only see real BRAs when the two sides do substantially different work.
Take the simplest possible kernel and walk through its SASS line by line. Compare two architectures to see how the same source compiles differently.
__global__ void vec_add(const float* a, const float* b, float* c) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
c[i] = a[i] + b[i];
}
MOV R1, c[0x0][0x28] ; // stack base
S2R R0, SR_CTAID.X ; // blockIdx.x
S2R R3, SR_TID.X ; // threadIdx.x
IMAD R0, R0, c[0x0][0x0], R3 ; // i = bid*bdim + tid
IMAD.WIDE R2, R0, 0x4, c[0x0][0x210] ; // &a[i]
LDG.E R5, [R2.64] ; // load a[i]
LDG.E R7, [R2.64+0x800000] ; // load b[i]
FADD R5, R5, R7 ;
STG.E [R2.64+0x1000000], R5 ; // store c[i]
@PT EXIT ;
S2R — "special to register": reads built-ins like SR_CTAID.X.IMAD.WIDE — one fused 64-bit multiply-add for index-to-pointer arithmetic.LDG.E — load global with extended (64-bit) addressing; the +0x800000 is a relative offset from the same base.FADD — FP32 add.STG.E — store global with extended (64-bit) addressing.@PT EXIT — unconditional kernel exit.Three pointers (a, b, c) are folded into a single base + offset because ptxas recognised they're laid out contiguously in the device-side argument struct. The 0x800000 and 0x1000000 are stride bytes between buffers in this build — they'll vary.
MOV R1, c[0x0][0x28] ;
S2R R0, SR_CTAID.X ;
S2R R2, SR_TID.X ;
IMAD R0, R0, c[0x0][0x0], R2 ;
SHF.L.U64 R3, R0, 0x2, RZ ; // i << 2
IADD3 R2.CC, R3, c[0x0][0x140], RZ ; // + base_a (low)
IMAD.X R3, RZ, RZ, c[0x0][0x144] ; // + base_a (high carry)
LDG.E.SYS R5, [R2] ; // load a[i]
IADD3 R2.CC, R3, c[0x0][0x148], RZ ;
IMAD.X R3, RZ, RZ, c[0x0][0x14c] ;
LDG.E.SYS R7, [R2] ; // load b[i]
FADD R5, R5, R7 ;
IADD3 R2.CC, R3, c[0x0][0x150], RZ ;
IMAD.X R3, RZ, RZ, c[0x0][0x154] ;
STG.E.SYS [R2], R5 ;
EXIT ;
Turing has no IMAD.WIDE: 64-bit address arithmetic takes a shift + a 32-bit add with carry (IADD3.CC) followed by a high-half IMAD.X. Hopper folds it all into one instruction. The kernel went from 14 SASS ops to 9 — not because the work shrank, but because the encoding got denser.
Watch one matmul tile travel through three layers. The C++ frontend is a thin wrapper; PTX exposes the MMA almost identically; SASS is one tensor-core op.
#include <mma.h>
using namespace nvcuda;
__global__ void tile_matmul(const half* A, const half* B, float* C) {
__shared__ half sA[16][16], sB[16][16];
// ... cooperative load of the 16x16 tile into smem ...
wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;
wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag;
wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;
wmma::fill_fragment(c_frag, 0.0f);
wmma::load_matrix_sync(a_frag, &sA[0][0], 16);
wmma::load_matrix_sync(b_frag, &sB[0][0], 16);
wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
wmma::store_matrix_sync(C, c_frag, 16, wmma::mem_row_major);
}
// f16 a/b fragments are 4 x .b32 (= 8 fp16 elements per thread, packed half2)
wmma.load.a.sync.aligned.m16n16k16.row.shared.f16
{%a0, %a1, %a2, %a3}, [smem_a], 16;
wmma.load.b.sync.aligned.m16n16k16.col.shared.f16
{%b0, %b1, %b2, %b3}, [smem_b], 16;
// f32 accumulator fragment is 8 x .f32 per thread
wmma.mma.sync.aligned.row.col.m16n16k16.f32.f16.f16.f32
{%c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7},
{%a0, %a1, %a2, %a3},
{%b0, %b1, %b2, %b3},
{%c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7};
wmma.store.d.sync.aligned.m16n16k16.row.global.f32
[out], {%c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7}, 16;
LDSM.16.MT88.4 R12, [R20+0x0] ; // ldmatrix → A frag
LDSM.16.M88.4 R20, [R20+0x800] ; // ldmatrix → B frag
HMMA.16816.F32 R8, R12, R20, R8 ; // 16x8x16 HMMA, FP32 accum
STS.128 [R28], R8 ; // drain to shared
STG.E.128 [R30.64], R8 ; // store global
From sm_70 onwards the PTX MMA family is a near-direct mapping to a single tensor-core SASS instruction (HMMA or WGMMA). NVIDIA exposes the hardware almost transparently — the heavy lifting is in the operand layout, not the op itself.
HMMA on Ampere implements m16n8k16 in a single hardware tick; the 16×16×16 logical tile is two such ops issued back-to-back. cuobjdump sometimes shows them collapsed; sometimes split. Hopper's WGMMA jumps straight to m64n128k16.
ptxas is the optimising assembler that lowers PTX to SASS. It does dead-code elimination, register allocation, instruction scheduling for ILP, peephole rewriting, address-arithmetic folding, and constant propagation. Two of its diagnostic outputs are essential reading.
Reports per-kernel: registers used, shared memory bytes, spill loads, spill stores, stack usage, constant memory. The register count is what decides occupancy.
Warn when the register allocator spills locals to .local (which is L1/L2-backed DRAM, not on-chip). A spill in the inner loop is usually catastrophic.
Full optimisation. -O0 exists, mostly useless. -O3 -ftz=true for speed; -prec-div=false -prec-sqrt=false for the fastest FP throughput.
Caps the per-thread register count. Forces spills if the kernel really needed more — usually used to push occupancy up at the cost of memory traffic. Measure both ways.
-Xptxas -dlcm=ca sets the default load-cache mode (ca = cache all, cg = cache global only, cs = streaming). Per-load override via PTX ld.global.ca.f32.
Emits source-line metadata so cuobjdump, Nsight Compute and nvdisasm can map SASS back to your CUDA file. Negligible runtime cost; turn it on always.
nvcc -O3 -arch=sm_90 \
-Xptxas=-v,-warn-spills \
-lineinfo \
-o kernel.cubin kernel.cu
// example output:
// ptxas info : Compiling entry function 'tile_matmul' for 'sm_90'
// ptxas info : Function properties for tile_matmul
// 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
// ptxas info : Used 96 registers, used 0 barriers, 16384 bytes smem,
// 456 bytes cmem[0]
"Used 96 registers" on Hopper means each thread holds 96 32-bit regs. The SM has 65536 registers, so a 256-thread block uses 256×96 = 24576, leaving room for 2 concurrent blocks per SM (limit is occupancy, not regs). "0 bytes spill" means the register allocator was happy. If you see non-zero spills in a hot kernel, raise -maxrregcount or rewrite the kernel to use less state.
Pick a SASS mnemonic; see its operand shape, latency class, throughput, predicate behaviour, the PTX op it usually comes from, and a one-line performance note. Knowledge inlined as a JS object — everything runs locally.
"Latency" is how many cycles between issue and result-available. "Throughput" is how often the warp scheduler can issue another op of this kind — usually 1 per cycle for ALU ops, much less for memory ops. The scheduler hides latency by issuing other warps' instructions in between.