Reading native GPU instructions

Your first native listing

Start with this GPU kernel: out[i] = in[i] + 1. In the listing, find the load that produces the value, the add that consumes it, and the store that writes the result. For in[0] = 10, following those three instructions should take you from 10 to 11.

Download source and native listings (.zip)

The kit includes editable PTX and CUDA sources plus seven kernel listings captured with CUDA 13.3 for an RTX 5060 Laptop GPU (sm_120). In x64 Developer PowerShell, extract the kit and enter the directory containing increment.ptx. Build and run the same binary you disassemble.

ptxas -arch=sm_120 -v increment.ptx -o increment.cubin
cuobjdump --dump-sass increment.cubin
nvcc -std=c++17 run.cpp -lcuda -o run.exe
.\run.exe increment.cubin increment

You can read the captured .sass files and do the browser exercises on any computer. Rebuilding needs the CUDA Toolkit and a supported host compiler. Running a cubin requires compatible hardware and a CUDA driver.

NVIDIA CUDA Binary Utilities

Machine instructions and target GPUs

SASS is the native instruction set of a particular NVIDIA GPU architecture. CUDA C++ or handwritten PTX provides the input. The assembler optimizes it, allocates hardware registers, and writes machine code into a cubin. Disassembly renders the machine code as readable mnemonics and operands.

A fat binary can hold several native images and PTX for driver JIT. When investigating a binary from another machine, record the GPU, code image, entry point, toolkit and driver. The line .target sm_120 in this lab identifies the captured native image. Compiling the same PTX for another architecture can change the opcodes and register numbers.

/*00e0*/    IADD R7, R2, 0x1 ;

The hexadecimal comment is the instruction’s byte offset in this listing. The sample adds the value in R2 and immediate 1, then writes R7. The next instruction in this capture starts at 0xf0. Use the addresses from the selected binary when correlating profiler results.

NVIDIA instruction set reference

Inspecting the executed image

Run cuobjdump --dump-sass on an executable or cubin. Use nvdisasm on an extracted or directly assembled cubin when you need additional control-flow or register-liveness views. Keep the cubin with your measurement. If the application loads PTX text directly, the driver JIT supplies the actual machine code.

nvcc -arch=sm_120 -lineinfo -cubin increment.cu -o cpp-increment.cubin
cuobjdump --dump-sass cpp-increment.cubin
nvdisasm -g cpp-increment.cubin

Inspect the complete function even when a short excerpt answers the question. For example, a native address calculation might reuse a register that earlier held a pointer. Find the most recent defining instruction along the path you are tracing. Source correlation helps locate the operation. Several source expressions can contribute to one instruction.

NVCC compilation workflow

Registers arithmetic and predicates

Ordinary registers such as R7 hold per-lane data. UR4 names a uniform register. P0 is a predicate, PT is the always-true predicate, and RZ supplies zero as an input. A destination and a source can name the same register. Read the old source value before recording the new destination value.

/*0050*/    IMAD R7, R7, UR4, R0 ;
/*0060*/    ISETP.GE.U32.AND P0, PT, R7, UR5, PT ;
/*0070*/    @P0 EXIT ;

In increment, the multiply-add computes blockWidth * blockId + threadId. Suppose the width is 64, the block is 1 and the thread is 6. The computed index is 70. If the count in UR5 is 70, the unsigned greater-than-or-equal comparison sets P0 and that lane exits. Repeat with thread 5. It continues with index 69.

The .U32 modifier matters when high bits are set. Separately, FFMA combines an FP32 multiply and add with a single rounding. Replacing it with independently rounded operations can change results in cancellation cases. Decide the required numerical behavior before changing compiler flags.

Floating point and FMA on NVIDIA GPUs

Address construction and memory traffic

/*00b0*/    IMAD.WIDE.U32 R2, R7, 0x4, R2 ;
/*00c0*/    LDG.E R2, desc[UR4][R2.64] ;
/*00d0*/    IMAD.WIDE.U32 R4, R7, 0x4, R4 ;
/*00e0*/    IADD R7, R2, 0x1 ;
/*00f0*/    STG.E desc[UR4][R4.64], R7 ;

At 0xb0 the old register pair starting at R2 is the input base. The instruction adds four times the index to that base. The load takes the resulting address and writes a scalar word back to R2. By 0xe0, R2 denotes the loaded data. The destination address uses the pair starting at R4. For i = 69, each address is 276 bytes past its respective base.

To reason about traffic, list every participating lane’s address and access width. A full warp loading adjacent four-byte words from a 32-byte-aligned base touches four 32-byte sectors. A stride of eight words touches 32 sectors. The useful bytes are identical. Cache behavior and memory-system measurements determine the actual DRAM traffic.

NVIDIA Best Practices Guide on coalescing

Warp communication and active lanes

/*0100*/    SHFL.BFLY PT, R9, R0, 0x1, 0x1f ;

In pair, all 32 physical lanes reach this exchange. Each lane receives its neighboring lane’s R0 into R9. The XOR distance is one, so lanes 6 and 7 exchange values. The immediate 0x1f expresses the lane bound for this operation. Recover the participation contract from the controlling flow and PTX source.

For n = 65, lane 0 in the final block wants data from lane 1. That lane is outside the logical array, so it enters the exchange with zero. Guard global accesses and preserve participation. Sketching the data value and participation separately helps you detect a faulty early exit.

SIMT kernels and warp behavior

Shared tiles and handoffs

/*0120*/    LDGSTS.E [R7], desc[UR6][R2.64] ;
// intervening instructions omitted
/*0170*/    LDGDEPBAR ;
/*0180*/    LEA R0, R0, UR4, 0x2 ;
/*0190*/    DEPBAR.LE SB0, 0x0 ;
/*01a0*/    BAR.SYNC.DEFER_BLOCKING 0x0 ;

The mirror kernel stages a 64-element tile in shared memory. A valid lane copies a global word to its slot. A tail lane initializes its slot to zero. After copy completion and the block handoff, each valid lane reads slot 63 - threadIdx.x. This puts the producer and consumer in different warps for many pairs.

In this capture, the native dependency instructions and block barrier implement different pieces of the PTX protocol. When adapting a pipeline, distinguish transfer completion, peer visibility, and permission to overwrite a tile. An early overwrite can race a slow consumer even when every incoming copy completed.

CUDA asynchronous data copies

Atomics ordering and scope

In the records sample, each lane adds one to a bin and uses the old value as a ticket. The listing contains ATOMG.E.ADD.STRONG.GPU. To understand what the program guarantees, pair the native instruction with its PTX semantic qualifiers, memory space and scope, and inspect how the returned value is consumed.

Suppose 32 lanes increment one initially empty bin. Tickets must be unique and together cover 0 through 31. Lane IDs do not determine ticket order. Contention changes timing. Counting independent events and publishing a separate payload require different reasoning.

For a producer writing a payload followed by a ready flag, check the producer’s release, the consumer’s acquire, and whether the acquire observes that release at a scope covering both parties. A single atomic-looking mnemonic in a disassembly does not establish that whole relationship.

PTX memory consistency model

Tensor Core instructions

/*02c0*/    HMMA.16816.F32 R4, R4, R12, RZ ;

In mma.cu, 32 lanes jointly multiply a 16-by-16 FP16 tile by a 16-by-8 tile and accumulate into FP32. The inline PTX instruction lowers to this native HMMA on the test GPU. A short printed operand list can designate groups of registers. Its interpretation follows the specific instruction form and fragment layout.

The reference sets A to the identity, assigns B a distinct integer at each coordinate, and initializes C to zero. The four results owned by lane 5 must be 11, 12, 75 and 76. A matrix of ones could hide a permutation of the columns. Verify the numerical result as well as the presence of the native instruction.

nvcc -arch=sm_120 mma.cu -o mma.exe
.\mma.exe
cuobjdump --dump-sass mma.exe | Select-String HMMA

PTX matrix instructions and fragment layouts

Profiling and production decisions

Begin with a failing test or measured bottleneck. Inspect the optimized machine code that was actually run. If a register cap shortens register lifetimes but introduces spills, look at the register report, local-memory traffic and elapsed kernel time. The number of printed instructions alone cannot tell you which kernel wins.

For bandwidth, count useful bytes at the algorithm level and identify the memory interface your hardware counters measure. Warm up both variants, record CUDA events in one stream, and synchronize the stop event. Report the working set, target GPU and variation. Use source and SASS correlation to find which instruction issues a costly memory request.

When chasing incorrect output, run small inputs at warp and block boundaries. Check invalid accesses with memcheck, shared-memory hazards with racecheck, and supported synchronization errors with synccheck. An in-bounds load from the wrong row still needs a CPU oracle to expose it.

Nsight Compute profiling guide

Compute Sanitizer

Implementing algorithms in the simulator

Open Algorithms to write a native-style instruction fragment. Each task defines the initial register and memory layout and the output the checker inspects. Trace the sample once by hand, implement the algorithm, and choose Run tests. A failure identifies a mismatching case or an instruction that needs correction. Editing and rerunning keeps the problem open.

The simulator documents its supported operations in the following reference section. Its bounded execution checks values on test inputs. Scheduling, timing and execution on a physical GPU require a native toolchain and hardware. To connect an exercise to a real program, express the same computation in PTX or CUDA, assemble it for your GPU, and compare the resulting disassembly and numerical outputs.

Download source and native listings (.zip)