What changes between CUDA, PTX, and SASS?

If a CUDA kernel adds one to an array element, where does that addition appear in the code the GPU executes? Each representation answers a different part of the question. CUDA C++ describes the update a thread should perform. PTX makes its loads, arithmetic, and stores explicit using virtual registers. SASS shows the native instructions selected for a particular GPU architecture.

One thread, one increment

The collection’s increment example has an input array, a separate output array, and a length. Each thread computes an index. A thread whose index falls inside the array executes out[i] = in[i] + 1u;. If element 7 contains 41, the thread responsible for index 7 writes 42 to output element 7. The source tells us which value must change and which threads may write.

The accompanying handwritten PTX describes the same update. It calculates the input and output byte addresses, then expresses the value’s journey in three instructions. The .u32 suffix identifies a 32-bit unsigned value. Square brackets around an address identify a memory operand.

ld.global.u32 %value, [%src];
add.u32       %value, %value, 1;
st.global.u32 [%dst], %value;

The load puts 41 into %value. The add reads that value and replaces it with 42. The store copies 42 to the address in %dst. A compiler can optimize this PTX for its target and assign the live values to machine registers. When an application distributes PTX, a compatible driver can also compile it at run time. NVIDIA PTX ISA · CUDA compilation

Follow the value through the native registers

Here are four consecutive instructions from the collection’s disassembly of that PTX, assembled with CUDA Toolkit 13.3 for sm_120. The hexadecimal numbers on the left are byte offsets within the function. Original source and disassembly (ZIP)

/*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;

The load delivers the input value to R2. Before the first use of that value, an independent instruction computes the output address using the element index in R7. Four bytes separate neighboring elements. Once the address calculation has consumed the index, the addition can reuse R7 for the result. The store reads the result from R7. This explains how a single PTX name, %value, corresponds to a loaded value in one native register and an incremented value in another.

A different target or compiler build can change both the register numbers and the instruction sequence. The dependencies still matter. The addition needs the input value, and the store needs both the updated value and the output address. NVIDIA CUDA Binary Utilities

Where measurement begins

The listing shows a global load and a global store. Whether those accesses limit execution time depends on the workload and hardware behavior, including cache hits and how neighboring lanes’ addresses group into memory transactions. Runtime measurements can distinguish a kernel that saturates memory bandwidth from one that spends time waiting for a dependent load. NVIDIA Nsight Compute Profiling Guide