CUDA Register Mapping: From PTX to SASS

Introduction

Register allocation is one of the critical aspects of GPU programming. On CPUs, the hardware’s “out-of-order” execution engine hides inefficiencies through register renaming, dynamically managing hundreds of physical registers behind 16 visible ones(General Purpose Registers). GPUs work differently: the per-thread register footprint chosen by the compiler becomes a real occupancy resource, with no CPU-style dynamic register renaming safety net.

In this post, I’ll explain how CUDA registers flow from PTX (a virtual, portable ISA) → SASS (the actual hardware ISA), and why this two-layer mapping fundamentally changes how you should think about register usage. Understanding register mapping is essential for grasping CUDA occupancy, latency hiding, and performance optimization - themes we’ll explore in depth in future posts on PTX-to-SASS compilation.

CPU Registers as a Baseline

To understand GPU registers, it helps to first see how CPUs handle them, then contrast the two.

Architectural vs. Physical Registers

Modern CPUs expose a small set of architectural registers to software:

But under the hood, a modern CPU core has far more physical registers-typically 200–224 integer registers per core (e.g., Zen 4/5, Intel Alder Lake). These hidden registers enable out-of-order execution to break false dependencies and extract instruction-level parallelism (ILP).

Register Renaming: Why Hardware Needs More Than It Shows

Consider this code sequence:

ADD R1, R2, R3    ; R1 = R2 + R3
SUB R1, R4, R5    ; R1 = R4 - R5

In program order, the second instruction must wait for the first to complete (true dependency). But CPU cores don’t execute in strict program order-they use register renaming to decouple architectural names (R1) from physical storage:

After renaming (internal CPU logic):
ADD P1, P2, P3    ; Physical P1 ← P2 + P3  (R1 → P1)
SUB P7, P4, P5    ; Physical P7 ← P4 - P5  (R1 → P7)

Now the CPU scheduler sees that SUB only depends on P4 and P5-not on P1-so both can execute in parallel on independent execution units. The ROB (Reorder Buffer) retires them in order later, updating architectural R1 correctly.

Key insight: Renaming is dynamic (happens at runtime) and transparent to software. Assembly code never mentions P1, P2, etc.; it always uses the 16 visible names (RAX–R15).

How Partial Writes Complicate Things

x86-64 registers have an aliasing structure:

RAX (64-bit full)
├── EAX (32-bit lower half)
    ├── AX (16-bit lower quarter)
        ├── AH (bits 8–15)
        └── AL (bits 0–7)

If you write only the lower 16 bits (AX), the CPU must internally merge the new value with the untouched upper bits:

MOV AX, 0x1234        ; Write only lower 16 bits
ADD EAX, 1            ; Now read/write lower 32 bits (depends on old EAX!)

Internally, the CPU performs:

new_EAX = (old_EAX & 0xFFFF0000) | new_AX

This creates an artificial dependency: the ADD instruction must wait for the MOV to complete, even though they operate on different bit ranges. This breaks out-of-order execution and reduces ILP-a performance penalty sometimes called the partial register stall.

Modern CPUs mitigate this with wider renaming support and partial-register tracking, but it remains a pitfall: always prefer full-width operations when possible.

GPU Registers - A Fundamentally Different Model

Key Differences

GPU register allocation is simpler but more rigid:

GPU Register Hierarchies: Two ISA Levels

Unlike CPUs (which expose one ISA layer), CUDA has two:

1. PTX (Parallel Thread Execution): A virtual, portable ISA with virtual registers. Compilers target PTX.
2. SASS (Streaming Assembly): The actual hardware ISA with finite numbered registers chosen by the backend for the target GPU.

PTX: The Virtual Layer

When you write CUDA C++ and compile it with nvcc, the compiler generates PTX as an intermediate representation:

// PTX (virtual ISA)
ld.global.f32   %f1, [%rd1]    ; Load a[i] from global memory
ld.global.f32   %f2, [%rd2]    ; Load b[i] from global memory
add.f32         %f3, %f1, %f2  ; Add them
st.global.f32   [%rd3], %f3    ; Store result to global memory

PTX uses virtual registers (%f1, %f2, etc.) as a compiler-facing abstraction. This is one reason PTX is portable, but portability still depends on the PTX version, target architecture, feature use, and driver/toolchain support.

SASS: The Hardware Layer

When ptxas (the PTX assembler) compiles PTX to SASS for a specific GPU, it performs real register allocation:

// Representative SASS
LDG.E R2, [R4]        ; Load a[i] into physical R2
LDG.E R4, [R6]        ; Load b[i] into physical R4
FADD R2, R2, R4       ; R2 = a[i] + b[i]
STG.E [R8], R2        ; Store result from R2

Now registers are real, finite resources. The compiler decides which virtual PTX registers map to the numbered SASS registers used by each thread. On many modern NVIDIA GPUs, the CUDA resource limit is up to 255 32-bit registers(R0-R254) per thread, but the practical number is usually much lower because high register usage reduces how many warps can fit on an SM.

Relative Slot Assignment

Here’s a useful mental model: when ptxas emits SASS code with, say, R0, R2, R4, those names are best read as per-thread register slots, not absolute chip-wide register addresses.

Consider an SM with 64K registers running 2048 threads, where each thread gets 32 registers:

SM Register File: 65,536 32-bit registers (64K)
├── Thread 0:    R0–R31  (conceptual slots 0–31)
├── Thread 1:    R0–R31  (conceptual slots 32–63)
├── Thread 2:    R0–R31  (conceptual slots 64–95)
├── ...
└── Thread 2047: R0–R31  (conceptual slots 65,504–65,535)

Each thread conceptually sees its own R0–R31. A simple way to picture the storage is:

Physical address = base_register_for_thread_N + 2

Treat that formula as an intuition, not as an official NVIDIA-documented addressing rule. NVIDIA documents the visible resource limits, while microarchitecture papers and benchmarks often infer lower-level register-file behavior experimentally. For performance work, the important consequence is still solid: the compiler reports a per-thread register count, and the SM has a finite register-file budget shared by all resident warps.

GPU Register Properties: Simplicity by Design

Uniform 32-bit Size

CUDA general-purpose registers are allocated as 32-bit registers. No aliasing:

// SASS
add.s32 r1, r2, r3    ; Add two 32-bit integers

This always:

1. Reads full 32-bit r2
2. Reads full 32-bit r3
3. Computes result
4. Writes full 32-bit r1

No merge logic. No partial-register penalties. No dependency ambiguity.

Lower-Precision Operations

For operations narrower than 32 bits, the full register is still used, but only the relevant bits participate:

add.s16 r1, r2, r3    ; Add two 16-bit integers (stored in 32-bit registers)

The upper 16 bits are typically undefined or zero-padded, depending on context.

64-bit Values

64-bit operations use pairs of registers:

add.s64 r1, r2, r3, r4    ; r1:r2 = r3:r4 (64-bit add)
; r1 = lower 32 bits
; r2 = upper 32 bits

Data Packing

For sub-32-bit types (e.g., int8), multiple values can pack into one 32-bit register:

dp4a r0, r1, r2, r3
; Dot product of 4×int8 vectors
; r0 = r3 + (r1[0]*r2[0] + r1[1]*r2[1] + r1[2]*r2[2] + r1[3]*r2[3])
; where r1 = [int8_0, int8_1, int8_2, int8_3]

This packing is powerful for efficiency but requires careful kernel design.

Register Pressure and Occupancy - The GPU Tradeoff

This is where GPU register allocation becomes a nuanced optimization problem.

The Core Tradeoff: Registers vs. Occupancy

When you write a kernel, the compiler (nvcc + ptxas) determines:

“This kernel needs N registers per thread.”

For this blog series, I am using my current GPU RTX 2060/Turing sm_75 limit as the main example:

• 65,536 32-bit registers per SM
• 2,048 maximum resident threads per SM
• 64 maximum resident warps per SM
• 255 maximum registers per thread

Say N = 32 registers/thread:

Register-limited threads = 65,536 / 32 = 2,048 threads
Register-limited warps   = 65,536 / (32 threads × 32 registers) = 64 warps

But if another kernel needs N = 64 registers/thread:

Register-limited threads = 65,536 / 64 = 1,024 threads
Register-limited warps   = 65,536 / (32 threads × 64 registers) = 32 warps

The same hardware, running the same kernel, but with double the registers → half the active threads.

Occupancy: Why It Matters

GPU latency hiding depends on having enough active warps to hide memory stalls. When a warp stalls (waiting for a load), the scheduler switches to another ready warp. If register pressure forces occupancy down, fewer warps are available, stalls aren’t hidden as well, and performance collapses.

This table is still simplified. Real occupancy also depends on block size, shared memory, maximum blocks per SM, launch bounds, and allocation granularity. The principle is the part to remember: more registers per thread can mean fewer resident warps, and fewer resident warps can expose latency.

Register Spilling on GPU

Unlike CPUs, where stack spills often hit in nearby CPU caches, CUDA register spills go to local memory. Local memory is private to each thread in the CUDA programming model, but it is backed by device memory and may be cached by the GPU memory hierarchy.

That means a spilled value can be much more expensive than a register value:

• a cached local-memory access may be tolerable,
• an uncached or poorly reused local-memory access can behave like normal global-memory traffic,
• and a hot loop with repeated spill loads/stores can become memory-latency bound.

Compilers try hard to avoid spilling, but sometimes it’s unavoidable. When it happens:

1. Local-memory traffic increases
2. Instruction count increases
3. Latency hiding can fail if there are not enough ready warps
4. Performance can drop sharply

Measuring Register Usage

Use nvcc flags to query register usage:

nvcc -o kernel.ptx --ptx my_kernel.cu
ptxas -v my_kernel.ptx   # Shows registers used

Or directly:

nvcc --resource-usage my_kernel.cu

Output example:

ptxas info    : 32 bytes gmem
ptxas info    : Compiling entry function '_Z11my_kernelPf' for 'sm_75'
ptxas info    : Function properties for _Z11my_kernelPf
    0 bytes stack frame, ...
    96 bytes spill stores (3 spills)
    96 bytes spill loads (3 spills)
    32 registers

The 32 registers tells you how many registers each thread uses. High numbers or spill counts are warning signs.

Practical Implications for Kernel Optimization

Write Register-Efficient Kernels

1. Minimize temporary variables: Each live variable occupies a register.

   // Bad: creates many temporaries
   float a = x + y;
   float b = a * z;
   float c = b - w;
   result = c / 2;
   
   // Better: reuse or inline
   result = ((x + y) * z - w) / 2;
   

2. Use -maxrregcount cautiously to force the compiler to reduce register usage:

   nvcc -maxrregcount=64 my_kernel.cu
   

   This can improve occupancy if register pressure is the limiting resource, but it can also create spills and make the kernel slower. Use it as an experiment, then compare register count, spills, achieved occupancy, and runtime.

3. Check PTX and SASS: Always inspect what the compiler generated:

   nvcc -ptx my_kernel.cu   # See the PTX layer
   ptxas -v my_kernel.ptx   # See the SASS mapping and resource usage
   

4. Profile occupancy: Use NVIDIA’s profiling tools (Nsight Compute, Nsight Systems) to see the actual occupancy your kernel achieves vs. theoretical max.

Interaction with Thread Block Size

Occupancy also depends on thread block size. For example:

• Block size 128 threads, 64 regs/thread = 8,192 registers per block
• Block size 256 threads, 64 regs/thread = 16,384 registers per block

Larger blocks use more registers. On GPUs with limited register files, smaller blocks + more blocks sometimes outperform one large block (even though ILP per block might be lower).

Key Takeaways

The GPU register model is simpler than CPUs but demands discipline: every additional register your kernel uses can reduce the number of threads that can run simultaneously when registers are the limiting resource. Understanding this tradeoff is fundamental to writing efficient CUDA code.

Read More

• NVIDIA CUDA C++ Programming Guide: Occupancy
• PTX ISA Documentation
• NVIDIA Nsight Compute Documentation
• Dissecting the NVIDIA Turing T4 GPU via Microbenchmarking
• Volkov, V. (2010). “Better Performance at Lower Occupancy”. GTC presentation on GPU register allocation and occupancy optimization.

GPU Microarchitecture Series Navigation

Previous: CUDA PTX: Learning to Read NVIDIA’s Virtual ISA