100 DAYS OF GPU Course preview
Day 46Module 5
in-technical-review

Reading PTX and SASS, and which one runs

Someone compared two kernels, did not believe the register count, and dumped the intermediate code to check:

--ptxas-options=-v -arch=sm_20 -O2: kernel 1 = 21 registers

kernel 1: .reg .u32 %r<49>; .reg .u64 %rd<56>; .reg .f32 %f<81>; .reg .f64 %fd<94>; .reg .pred %p<8>;

"looking at the ptx white paper I assume that this means kernel 1 is allocating a total of 288 'virtual' registers"

https://forums.developer.nvidia.com/t/compiler-option-ptxas-options-v-gives-wrong-register-count/17728 (checked 2026-08-30)

21 against 288. Neither number is wrong. They are counts of two different programs: the one the compiler was given and the one it produced, and only the second one ever reaches a GPU.

This lesson shows how to distinguish those two outputs. You will can find an unrolled loop and a register spill in real machine code, with two commands, no profiler and no GPU.

Two compilation stages

nvcc manages several compilation steps. It sends your host code to the system compiler and compiles your device code separately into becomes PTX: "a low-level parallel thread execution virtual machine and instruction set architecture", whose programs "are translated at install time to the target hardware instruction set", and one of whose stated goals is to "Provide a stable ISA that spans multiple GPU generations" (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html , checked 2026-08-30).

A virtual machine has no register file, so PTX does not allocate one. It names values, not places. The front end emits it in static single assignment form, where every write gets a fresh name, which is why a .reg directive can declare hundreds of them for a kernel that ends up using twenty.

ptxas is the second compiler. It takes that PTX and a target architecture, then produces SASS, the instructions that the target GPU issues. At this stage, the compiler assigns real registers, schedules instructions, and moves values that do not fit to local memory.

ptxas has its own optimiser and its own level: ptxas's --opt-level N (-O) is "Specify optimization level" with "Default value: 3", reached from the command line through -Xptxas, "Specify options directly to ptxas, the PTX optimizing assembler" (https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html , checked 2026-08-30).

That matters if you have been passing -O3 and expecting it to tune your kernel. -O3 is an option for the host compiler. The device optimiser is a separate program with a separate flag, and it is on by default.

Both outputs ship in the binary. -arch=sm_75 is shorthand for -gencode arch=compute_75,code=sm_75, so the fatbinary carries sm_75 machine code and compute_75 PTX beside it, the second existing so a future card with no matching cubin can JIT one at run time (day 3 covers the flag, and gencode is the term). That is why one file can show you both.

Diagram: one source file, two compilers, two outputs in one binary. A left-to-right pipeline of four boxes with the tool that produces each arrow named above it, and a dashed line marking where the architecture is first known. Box 1, ptx_sass.cu: one for loop over 64 taps, written once. Caption "one source, no architecture named yet." Box 2, compute_75 PTX, produced by the nvcc front end: typed virtual registers, one fresh name per write. Caption "virtual registers, no allocation, one file for every card from Turing on." Box 3, sm_75 cubin, produced by ptxas: real registers, scheduled instructions, spills if they did not fit. Caption "one architecture, and the only code an SM ever issues." Box 4, the fatbinary: boxes 2 and 3 packed together, with cuobjdump -ptx and cuobjdump -sass drawn as two arrows back out of it. Alt text: "One CUDA file becomes two device outputs. The nvcc front end emits PTX with virtual registers and no card named; ptxas turns that into sm_75 machine code with real registers. Both ship in one binary."

Why PTX can be mistaken for machine code

PTX looks like assembly. It has one operation per line, opcodes, operands, and numbered registers. It is also easy to get: nvcc -ptx prints it, and it is what Compiler Explorer puts in the compiler pane unless you go and turn SASS on, which is old enough to have its own feature request (https://github.com/compiler-explorer/compiler-explorer/issues/1856 , checked 2026-08-30).

This can lead readers to treat PTX as machine code and use it to predict performance.

The register count can differ by an order of magnitude. "The register allocation in PTX is completely irrelevant to the final register consumption of the kernel ... A piece of PTX with hundreds of registers can compile into a kernel with only a few registers" (https://stackoverflow.com/questions/11483321/what-kind-of-variables-consume-registers-in-cuda , checked 2026-08-30).

PTX instruction counts also differ from SASS because ptxas reorders, folds and unrolls, so the number of lines in the PTX is not the number of instructions that issue.

PTX still shows what the front end did with your source. You can check which loads it moved, whether it honored #pragma unroll, and whether a multiply and an add became one fma.

Read PTX for the front end's output. Read SASS for the instructions that the GPU executes.

Three kernels, chosen for what they leave in the machine code

Full program in code/day46-ptx-sass/ptx_sass.cu. It is a 64-tap decaying fold, written three ways, and none of the three is there because it is fast.

Three decisions matter here.

The artifact comes out of the binary, not out of the program. cuobjdump "extracts information from CUDA binary files ... and presents them in human readable format", -ptx will "Dump PTX for all listed device functions" and -sass will "Dump CUDA assembly" (https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html , checked 2026-08-30). Neither needs the program to run, so this whole lesson works on a machine with no NVIDIA GPU in it.

Two of the kernels must agree element for element. The rolled version takes its trip count as an argument and is handed the same constant the unrolled version was compiled with, so the two are one program compiled twice. The program checks that and prints the count of elements where they differ.

One thing changes per kernel. The first two differ in a trip count and a pragma, the third in a launch bound. Nothing else changes.

The fold whose trip count nobody can see:

__global__ void decayRolled(const float* __restrict__ in,
                            float* __restrict__ out, size_t n, int taps) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (i < n) {
        float acc = 0.0f;
        for (int k = taps - 1; k >= 0; --k) {
            acc = acc * kDecay + in[i + static_cast<size_t>(k)];
        }
        out[i] = acc;
    }
}

The same fold with a fixed count and a pragma that asks for unrolling:

__global__ void decayUnrolled(const float* __restrict__ in,
                              float* __restrict__ out, size_t n) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (i < n) {
        float acc = 0.0f;
#pragma unroll
        for (int k = kTaps - 1; k >= 0; --k) {
            acc = acc * kDecay + in[i + static_cast<size_t>(k)];
        }
        out[i] = acc;
    }
}

A two-pass window keeps all 64 taps live at once under a __launch_bounds__ budget that cannot hold them. Every value that does not fit becomes a spill: day 17 counted those bytes, and this is the lesson that shows you the instructions moving them.

__global__ __launch_bounds__(kThreadsPerBlock, kMinBlocksPerSm) void
decayStaged(const float* __restrict__ in, float* __restrict__ out, size_t n) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (i < n) {
        float window[kTaps];
        float carry = 0.0f;
#pragma unroll
        for (int k = 0; k < kTaps; ++k) {
            carry = carry * kDecay + in[i + static_cast<size_t>(k)];
            window[k] = carry;
        }
        float acc = 0.0f;
#pragma unroll
        for (int k = kTaps - 1; k >= 0; --k) {
            acc = acc * kDecay + window[k];
        }
        out[i] = acc;
    }
}

Note. decayStaged folds the running window rather than the raw input, so it computes a different number from the other two on purpose and has its own CPU reference. Reading the window backwards is what keeps every tap live between the two loops, and that is the whole source of the register pressure. Take that away and there is nothing to spill.

Results

Re-verified on a Tesla T4 with driver 580.173.02 and CUDA 13.0 V13.0.88 on 2026-09-02. Correctness, registers, local storage, O3 instruction shapes and the one-line PTX difference under -Xptxas -O0 reproduced. Runtime roughly doubled, but no lesson conclusion changed.

The original CUDA 12.6 results and artifacts remain the baseline below; the CUDA 13 PTX, SASS, cubin, O0 and resource-usage artifacts are listed in front matter.

Originally measured on a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -std=c++17 -O3 -arch=sm_75; disassembled with cuobjdump and nvdisasm from the same toolkit. Captured 2026-09-01 on the project's verification node. The transcript is in the page's evidence file, and the six disassembly artifacts named in the README sit beside the code, because the disassembly is what the exercise is done against.

The run table first:

kernel          regs   local B        ms
-------------- ----- --------- ---------
decayRolled       42         0     0.114
decayUnrolled     62         0     0.113
decayStaged       64        48     0.157

all three kernels match their CPU reference; decayRolled and decayUnrolled
differ at 0 of 1049187 elements, so the two are one program compiled twice

And the shape table, filled from the SASS:

Kernel Back edge FFMA STL, LDL regs local B
decayRolled 2 29 0, 0 42 0
decayUnrolled 0 63 0, 0 62 0
decayStaged 0 126 11, 11 64 48

Against the five predictions:

  1. Held. decayRolled keeps two backward BRAs, a loop and its remainder.
  2. Two parts did not hold. The fold has no back edge, but "at least one FFMA per tap" came back 63 for 64 taps: the first tap multiplies an accumulator that still holds zero, so ptxas emitted one FADD and folded the dead multiply. And decayRolled's 29 FFMAs for 64 taps is the partial unroll the prediction named as its own most likely cause: the compiler unrolled the loop body and kept a smaller loop around it, exactly the allowed transformation.
  3. Held. STL and LDL appear only in decayStaged, 11 of each.
  4. Held. 64 registers per thread under __launch_bounds__(256, 4), with 48 bytes of local memory beside it: the staging array did not fit and the spill is real, at 0.157 ms against 0.113 for the same arithmetic in registers.
  5. Held except for one comment line. The -Xptxas -O0 rebuild changes the PTX in exactly one place, the ptxasOptions = -O0 header comment nvcc writes into the file, and no instruction. The SASS diff is 428 KB. The flag reaches only the second compiler.

Run it yourself

Two commands, on any machine with a CUDA toolkit, GPU or not:

nvcc -std=c++17 -O3 -arch=sm_75 -o ptx_sass ptx_sass.cu
cuobjdump -sass ./ptx_sass | less

In a browser, Compiler Explorer maps each source line to its generated instructions. NVIDIA describes the view as one that "correlates each line of your source code with the corresponding generated instructions", with colour coding (https://developer.nvidia.com/blog/compiler-explorer-the-kernel-playground-for-cuda-developers/ , checked 2026-08-30).

Pick an NVCC compiler rather than NVRTC and select a released compiler version. Set the target architecture for the code you want to inspect, then turn on Compile to binary object to show SASS instead of PTX in the device pane.

The shipped artifact used nvcc133, targeted sm_75, and passed -std=c++17 -O3 -arch=sm_75. Use the compiler version and architecture that match the binary you need to inspect.

This lesson needs no profiler. Nsight Compute can add per-instruction counters to SASS, which show where warps stall. Those counters need performance-counter access.

Plain ncu reports ERR_NVGPUCTRPERM, in full: The user running <tool_name/application_name> does not have permission to access NVIDIA GPU Performance Counters or the Hardware Event System on the target device (https://developer.nvidia.com/nvidia-development-tools-solutions-err_nvgpuctrperm-permission-issue-performance-counters , checked 2026-08-30). On the measured node, sudo ncu works because RmProfilingAdminOnly: 1 restricts access. This proves the ordinary-user restriction on that node only.

Whether free Colab currently permits NCU counters remains open; the shipped disassembly needs no counters and is sufficient for the exercise.

Exercise

Two hunts and one question, all done on a disassembly you generate yourself or on the one shipped with the lesson.

  1. Find the loop. In cuobjdump -sass output, one of decayRolled and decayUnrolled still has a loop and one does not. Name the instruction that proves it and say how many times the fold's body appears in each.
  2. Find the spill. In decayStaged, name the two instructions that move a spilled value, say which direction each moves it, and say why the register count on its own would not have told you.
  3. Rebuild with -Xptxas -O0. Diff the PTX against the first build, then diff the SASS. What do the two diffs, taken together, prove about which program optimised your kernel?

Time: 25 to 40 minutes. Submit: two instruction names, two counts, and one sentence for question 3.

Check: the quiz in content/quizzes/day46.toml explains each option. The answers to the three tasks above are in the reveal, and each names the instruction, the documentation line or the diff that settles it.

Questions 1 and 2 are grep -c on your own dump, so a wrong answer is a number you can go back and recount rather than an opinion. Question 3 is a diff that is either empty or not, and the README gives the exact two builds.

Hint 1

A loop that survives compilation has to be able to get back to its own first instruction. Every line of the dump has its address on the left. What would going backwards look like from there?

Hint 2

For the spill, find two three-letter opcodes that both end in L; the prior letter names the memory. For question 3, use the meaning of -Xptxas to decide which compiler stage receives the flag.

Solution

1. The back edge. BRA is a relative branch. A loop uses it to jump to a lower address, with a comparison such as ISETP setting its predicate.

decayRolled has that pair and its body appears once. decayUnrolled has neither, and its body appears once per tap. The unrolled version has more instructions but removes a comparison and branch per tap.

2. STL and LDL. "Store to Local Memory" writes a value, and "Load within Local Memory Window" reads it. Local memory is device memory with a per-thread address.

The register count alone cannot identify a spill because a spilling kernel may use fewer registers than the same kernel without the limit. The local memory byte count identifies the spill.

3. The PTX diff is empty and the SASS diff is not. -Xptxas passes options to the second compiler only, so the first stage produced the same file both times. The flag changed only the SASS stage.

The front end produces PTX. ptxas then chooses and schedules the target instructions.

The GPU executes SASS, not source or PTX.

Pitfalls

You counted registers in the PTX. Those are virtual, in single static assignment form, and a new one appears every time a value is written. The number that binds occupancy comes from ptxas: -Xptxas -v at build time, or cuobjdump -res-usage out of a finished binary. Day 17 does that arithmetic.

nvdisasm refused your executable. It "extracts information from standalone cubin files"; cuobjdump is the one that reaches inside a host binary and finds the cubins embedded in it (https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html , checked 2026-08-30). Point nvdisasm at a .cubin from nvcc --cubin, or use cuobjdump on the executable.

Your SASS has no source lines in it. cuobjdump -sass gives instructions and addresses. For lines you need -lineinfo at build time and nvdisasm -g, which annotates "with source line information obtained from .debug_line section, if present". Compiler Explorer draws the same mapping in colour.

You compared SASS from two architectures. Opcodes, register counts and spill bytes are all per -arch, and a build with several -gencode targets disassembles into one section per target. Going below the toolkit's floor does not compare, it stops: nvcc fatal : Unsupported gpu architecture 'sm_60', three spaces before the colon, which matters if you paste it into a search box (error page).

Your #pragma unroll did nothing. It needs a trip count the compiler can see. A bound that arrives as a kernel argument cannot be unrolled by anyone, which is what decayRolled is for, and day 25 has the same limit with blockDim.x.

You read -Xptxas -O0 output and drew a performance conclusion. The default level is 3. At 0 you are looking at a kernel nobody will ever run. Use it to see what ptxas did, never to decide what to write.

Go deeper

Next

Day 47 uses the same disassembly and changes one flag. -use_fast_math and -fmad change the SASS instructions and the numerical result. Day 45 set a register budget; SASS shows the resulting instructions.