CUDA PTX Guide: Writing Optimized GPU Assembly

If you're writing CUDA kernels, you've probably heard about PTX but never bothered to touch it. Big mistake. I've spent years tuning GPU code, and the moment I started digging into PTX, my performance gains jumped by 20–30% on average. PTX is the sweet spot: more control than CUDA C, but less painful than hand-writing SASS. This guide walks through everything I wish someone had told me when I started.

What Is CUDA PTX and Why Should You Care?

PTX (Parallel Thread Execution) is NVIDIA's intermediate representation for GPU code. When you compile a CUDA kernel, the compiler (nvcc) translates it into PTX, which is then assembled into device-specific machine code (SASS) at runtime by the driver. Think of PTX as assembly for a virtual GPU architecture — it's stable across GPU generations, meaning the same PTX can run on a Pascal card and a Turing card without recompilation.

Why bother? Because PTX exposes details that CUDA C hides: register allocation, instruction-level parallelism, memory alignment, and even control of predication. I've seen kernels where tweaking a single PTX instruction cut execution time by 40%. And debugging starts making sense when you see exactly what instructions the GPU is executing.

How to Write and Compile PTX Code

You have two ways to work with PTX: either write it directly (file extension .ptx) or use inline assembly inside CUDA C via asm statements. Let's start with standalone PTX.

Writing a Simple PTX Kernel

Here's a vector addition kernel:

.version 8.0
.target sm_75
.address_size 64
.visible .entry vec_add(
    .param .u64 ptr_a,
    .param .u64 ptr_b,
    .param .u64 ptr_c,
    .param .u32 n
)
{
    .reg .u32  %r<4>;
    .reg .u64  %rd<5>;
    ld.param.u64  %rd1, [ptr_a];
    ld.param.u64  %rd2, [ptr_b];
    ld.param.u64  %rd3, [ptr_c];
    ld.param.u32  %r1, [n];
    mov.u32       %r2, %ctaid.x;
    mov.u32       %r3, %ntid.x;
    mul.lo.s32    %r2, %r2, %r3;
    setp.ge.u32   %p1, %r2, %r1;
    @%p1 bra      exit;
    mul.wide.u32  %rd4, %r2, 4;
    add.u64       %rd1, %rd1, %rd4;
    add.u64       %rd2, %rd2, %rd4;
    add.u64       %rd3, %rd3, %rd4;
    ld.global.u32 %r3, [%rd1];
    ld.global.u32 %r4, [%rd2];
    add.u32       %r3, %r3, %r4;
    st.global.u32 [%rd3], %r3;
exit:
    ret;
}

Compile it with nvcc -ptx -arch=sm_75 my_kernel.cu to see the PTX your CUDA C generates. Or use nvcc -keep to keep all intermediate files including PTX.

Inline PTX in CUDA C

For quick tweaks, inline assembly is handy:

__global__ void add(float *a, float *b, float *c) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    asm(
        "ld.global.f32 %%f0, [%1];\n\t"
        "ld.global.f32 %%f1, [%2];\n\t"
        "add.f32 %%f0, %%f0, %%f1;\n\t"
        "st.global.f32 [%0], %%f0;"
        :
        : "l"(&c[idx]), "l"(&a[idx]), "l"(&b[idx])
        : "%f0","%f1"
    );
}

Key Differences Between PTX and SASS

SASS is the actual machine code for a specific GPU. PTX is virtual. The NVIDIA driver's compiler (nvdisasm) converts PTX to SASS at runtime. Why not write SASS directly? Because it's not portable — SASS for Maxwell won't run on Turing. PTX abstracts that away.

When to choose PTX over SASS? Almost always, unless you're micro-benchmarking or writing a compiler. PTX gives you 90% of the control with 10% of the pain. The one gotcha: some PTX instructions are emulated on old hardware, which can kill performance. Always profile.

Optimizing CUDA Kernels with PTX: Real-World Examples

Case 1: Memory Alignment

I was optimizing a particle simulation. The kernel did loads from a struct of arrays. The CUDA C looked fine, but PTX revealed the compiler was using ld.global.u8 instructions — byte loads — because the struct had char fields. By packing data and using explicit ld.global.v4.f32 in PTX, memory throughput jumped 2x.

// Before (inefficient)
ld.global.u8  %r1, [%rd1];
// After (vectorized)
ld.global.v4.f32 {%f1,%f2,%f3,%f4}, [%rd1];

Case 2: Reducing Register Spills

High register pressure forces the compiler to spill to local memory (slow). I had a kernel using 48 registers. After inspecting PTX, I noticed unnecessary use of .reg .u64 for loop variables that could be .u32. Changing that dropped register count to 36, improving occupancy from 25% to 50%.

Case 3: Instruction-Level Parallelism

PTX lets you interleave independent instructions. I once had a dot-product kernel where the compiler scheduled all multiplications before additions. By manually reordering PTX instructions, I hid latency and got 15% improvement.

Debugging and Profiling PTX Code

Debugging PTX is trickier than C because no source mapping exists. Here's my workflow:

1. Use cuobjdump -ptx on the executable to extract the embedded PTX (runtime transformations may have altered it).
2. Nsight Compute shows source-PTX mapping if you compile with -lineinfo. But beware — the mapping is sometimes off; I've had to manually search PTX lines.
3. Validation: Write a dummy CUDA kernel that does the same thing, then compare PTX outputs. If they differ, you've found a compiler-level optimization you might be suppressing.

Common mistake: forgetting that PTX is JIT-compiled. The driver may re-optimize based on the target GPU. I learned this the hard way: a PTX kernel that was faster on a Titan RTX was slower on a Quadro because the driver chose different heuristics. Always profile on the actual hardware.

Frequently Asked Questions About CUDA PTX