PTX Programming from Your First Kernel to Tensor Cores

Start with a kernel whose result you can predict on paper, out[i] = in[i] + 1. Follow one thread from its ID to the four bytes it reads and the four bytes it writes. Then change one instruction and run the same test. That exercise gives you a way to learn the rest of PTX, from exchanging a value with another lane to assembling a matrix product across a warp.

The labs assume you can read C++ and work with pointers. The commands use Windows PowerShell, CUDA Toolkit 13.3, and an RTX 5060 Laptop GPU with compute capability 12.0. The first three labs share a small C++ launcher. The last uses inline PTX inside a CUDA C++ kernel. Unzip the companion lab files and run the commands from the ptx-lab directory.

PTX programming labs ↓

PTX 101 with one thread

PTX stands for Parallel Thread Execution. It is NVIDIA’s virtual instruction set. A CUDA C++ compiler can emit PTX as text. The offline assembler (ptxas) or the CUDA driver’s just-in-time compiler turns PTX into instructions for a specific GPU. A .cubin holds compiled device code, and disassembling it shows the native instructions commonly called SASS. This gives you three views of the same computation. Use C++ to express the result, PTX to reason about each thread, and SASS to see the implementation selected for one target. NVIDIA compiler guide · CUDA Binary Utilities

Predict one output

A kernel is a function the host asks many GPU threads to execute. A launch groups threads into blocks and blocks into a grid. In a one-dimensional launch, each thread has a number inside its block. Multiplying the block number by its width and adding the thread number gives a unique array index for the sizes in these labs. CUDA programming model

Read the entire first kernel. extern "C" gives the entry point a predictable name for the launcher, and __global__ marks a function the host can launch on the GPU.

extern "C" __global__
void increment(unsigned* out, const unsigned* in, unsigned n) {
    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        out[i] = in[i] + 1u;
}

Use n = 70, two blocks of 64 threads, and in[i] = 10*i. All block and thread numbers start at zero.

Thread 5 in block 1 reads at in + 276 bytes and writes at out + 276 bytes. Write down both addresses before looking at the assembly. The value to add is 1. The distance between elements is 4 bytes. Changing one of those constants changes the answer in a different way.

Compile and run

Open x64 Developer PowerShell for Visual Studio with the CUDA Toolkit on your path. Check the installed compiler with nvcc --version and the device with the first command below. For this GPU the numeric suffix is 120. On another machine choose a target listed by nvcc --list-gpu-arch and nvcc --list-gpu-code that works with its GPU. The supplied handwritten modules declare an sm_80 baseline. NVCC target options

nvidia-smi --query-gpu=name,compute_cap --format=csv
nvcc -ptx -arch=compute_120 increment.cu -o generated.ptx
nvcc -std=c++17 run.cpp -lcuda -o run.exe
.\run.exe generated.ptx increment
.\run.exe increment.ptx increment

Both runs should print PASS for nine lengths from 0 to 129. The launcher fills the input, marks spare output slots with 0xdeadbeef, launches the kernel, and compares every slot with a CPU calculation. At length zero it deliberately launches one block so you exercise the bounds check. Start here before reading the host code.

Read the handwritten version

Open generated.ptx beside increment.ptx. The compiler chooses names such as %rd5. The handwritten version names the same jobs explicitly. Brackets mean a memory access, and most arithmetic instructions put the destination first. Read the code once for the index, once for the addresses, and once for the value.

.version 7.0
.target sm_80
.address_size 64

.visible .entry increment(
    .param .u64 out_ptr,
    .param .u64 in_ptr,
    .param .u32 count
)
{
    .reg .pred %past_end;
    .reg .u32 %block, %width, %thread, %i, %n, %value;
    .reg .u64 %in, %out, %offset, %src, %dst;

    mov.u32     %block, %ctaid.x;
    mov.u32     %width, %ntid.x;
    mov.u32     %thread, %tid.x;
    mad.lo.u32  %i, %block, %width, %thread;

    ld.param.u32 %n, [count];
    setp.ge.u32  %past_end, %i, %n;
    @%past_end bra DONE;

    ld.param.u64 %in, [in_ptr];
    ld.param.u64 %out, [out_ptr];
    mul.wide.u32 %offset, %i, 4;
    add.u64      %src, %in, %offset;
    add.u64      %dst, %out, %offset;
    ld.global.u32 %value, [%src];
    add.u32       %value, %value, 1;
    st.global.u32 [%dst], %value;
DONE:
    ret;
}

The three header lines answer different questions. .version 7.0 selects the PTX language version, .target sm_80 names the feature baseline used here, and .address_size 64 sets the module’s address size. With the local CUDA 13.3 compiler, generated.ptx instead starts with PTX 9.3 and sm_120 for the command above. Both files implement the same kernel. PTX ISA for CUDA 13.3

Follow %i. The three mov instructions read the block ID, block width, and thread ID. mad.lo.u32 multiplies and adds, keeping the low 32 bits. setp.ge.u32 produces a one-bit predicate. A true predicate on @%past_end bra DONE takes the branch. The lab’s indices fit in 32 bits. For larger launches, widen the block multiplication before it overflows.

Follow %src. ld.param.u64 fetches the pointer passed by the host. mul.wide.u32 forms a 64-bit byte offset from the unsigned index. ld.global.u32 then fetches the actual element from device memory. The argument slot and the array element are two different locations. The final three instructions on the data path load the value, increment it, and store it.

First edit. Make a copy named plus2.ptx. Change the final operand of add.u32 %value, %value, 1 to 2. Run .\run.exe plus2.ptx increment without recompiling the launcher. Predict the first failure. At n=0 the guard still passes. At n=1, i=0 you should get 2 where the original CPU test expects 1. Update the CPU expectation only after you have explained that difference.

From PTX to the instructions on the GPU

Assemble the original file and run the same CPU test against the binary. First check that the kernel gives the expected answer. Then look at the instructions the offline assembler emitted. CUDA Binary Utilities

ptxas -arch=sm_120 -v increment.ptx -o increment.cubin
.\run.exe increment.cubin increment
cuobjdump -sass increment.cubin

On the RTX 5060 Laptop build, ptxas reported 10 registers per thread and zero spill loads or stores. Look for IMAD.WIDE.U32 forming an address, LDG.E loading, IADD adding one, and STG.E storing. The assembler folded a multiply and an address add together. Its register allocation also reused registers after values died. That is why counting virtual registers or PTX lines gives a poor prediction of execution time.

The Driver API launcher loads a module with cuModuleLoad, looks up its .entry symbol with cuModuleGetFunction, and calls cuLaunchKernel. Its argument list is {&dst, &src, &n}. These are host addresses of variables holding the two device pointers and the count. Their order and widths match the three .param declarations. When changing a parameter, change both sides. A device-callable PTX subroutine uses .func and call. For mixed-language code, match the documented calling convention, parameter alignment, and symbol names. CUDA Driver API launch contract · PTX Writer’s Guide to Interoperability

For deployment, decide which real architectures receive native cubins and which virtual target supplies PTX for driver JIT. Check both the driver’s supported PTX version and the hardware feature required by each instruction. Architecture-specific targets with an a suffix have additional compatibility restrictions. A PTX file that assembled for one target can still fail when loaded on another. NVIDIA architecture compatibility guide

Arithmetic types and control flow

Type suffixes describe how an instruction interprets its operands. .u32 and .s32 specify unsigned and signed integers. .b32 holds a 32-bit pattern. .f32 participates in floating-point arithmetic. .pred holds a condition. A mov.b32 can preserve a bit pattern across compatible 32-bit registers. cvt.rn.f32.s32 instead converts a signed integer’s numerical value to floating point using round-to-nearest-even. Try this distinction with integer 1. Numeric conversion produces floating-point 1.0, whose encoding is 0x3f800000. PTX operands and instructions

Use labels with a predicated bra to build loops. A grid-stride loop starts at the same index as increment, does one item, adds gridDim.x * blockDim.x, and repeats while the index is in range. In PTX, %nctaid.x supplies the grid width in blocks. For production-sized arrays use sufficiently wide counters for the initial product, stride, and byte offset. Exercise. Extend the kernel to use a grid-stride loop. In the launcher, first set the output capacity to std::max(n, blocks * width) + 1 so the output still holds all 129 elements when you set blocks = 1. Use the new block count in the launch. Thread 0 must cover indices 0, 64, and 128. If 128 retains the guard, inspect your loop condition and step.

Adjacent threads execute in groups called warps. On the lab GPU each warp has 32 lanes. When neighboring lanes take different branches, only the lanes on a given path perform its work. A predicated store (@%valid st.global…) also lets you express short conditional work. selp selects between two available values. When one choice would read past the end of an array, guard the load itself. Selecting its result afterward leaves the invalid access in place. CUDA SIMT execution

For floating point, specify the numerical contract before changing the instructions. fma.rn.f32 d,a,b,c computes a*b+c with one final rounding. Separate rounded multiply and add can give a different result. Try a = 1 + 2^-12, b = 1 - 2^-12, and c = -1 in FP32. Their exact product is 1 - 2^-24, which is representable. Now change both exponents to -13. The product rounds to 1 when rounded separately, while the fused result retains -2^-26. Check the emitted PTX when using CUDA C++ because the compiler may contract operations. NVIDIA guide to floating point and FMA

Memory spaces and access patterns

Read every load or store by asking who owns the storage, who can see it, and how the address was formed. A register value belongs to a thread. Local memory also belongs to a thread, but arrays requiring addressable storage and spilled registers can use the device memory hierarchy. Shared memory belongs to a block in these labs. Global memory holds the input and output buffers. Kernel arguments arrive through parameter space. Constant memory is useful for read-only values, particularly when a warp reads the same address. PTX state spaces

Trace a warp. In increment, the first 32 threads access 32 adjacent four-byte words. With a 32-byte-aligned base, that covers four 32-byte sectors. Shift the base by one word and the same 128 useful bytes span five sectors. Use a stride of eight words and the lanes now touch 32 sectors. This predicts the addresses requested by the warp. Cache hits and reuse affect the resulting DRAM traffic. A wider instruction such as ld.global.v4.u32 also requires a valid, naturally aligned 16-byte access and a tail path for a short final group. CUDA Best Practices Guide on coalescing

Shared-memory exercise. For a tile of four-byte words with 32 banks, sketch tile[32][32] in row-major order. If lane l reads tile[l][0], its word index is 32*l, so every lane requests a different word in bank 0. With a pitch of 33 words, the bank becomes (33*l) mod 32 = l. This explains a common padding trick for the transpose of a tile. The case where all lanes read the same word can use broadcast. CUDA Best Practices Guide on shared memory

Address conversions. A pointer passed to CUDA inline assembly is generally a generic address. Establish which space it points into before choosing a space-specific instruction. For a pointer known to point to shared memory, __cvta_generic_to_shared can produce its shared-space address. In mirror.ptx the symbol tile is declared in shared space, so the code takes its shared address directly. Cache operators such as .ca and .cg belong to the memory access policy. For communication between threads, choose explicit synchronization. NVIDIA inline PTX guide · CUDA programming guide

Exchanging values within a warp

Run .\run.exe pair.ptx pair. This kernel swaps adjacent pairs. Output 0 receives input 1, output 1 receives input 0, and so on. In a warp, the partner’s lane number is lane XOR 1. For an odd-sized input, the last valid element receives zero. The essential part of the kernel is five instructions.

mov.u32 %value, 0;
@%valid ld.global.u32 %value, [%src];
shfl.sync.bfly.b32 %peer, %value, 1, 31, 0xffffffff;
@%valid st.global.u32 [%dst], %peer;
ret;

bfly chooses the partner with XOR. Here 1 flips the lowest lane bit, 31 sets the upper lane bound with a full-warp segment, and the final mask names all 32 lanes. The launcher uses two full warps per block. Lanes beyond the logical end still execute the shuffle with value zero. For n=65, lane 0 of the last block therefore receives zero from lane 1. An early return before the shuffle would remove that source. When adapting the code to a subgroup, make sure every source lane you read participates with the correct mask. PTX warp shuffle instruction

Exercise. Turn the exchange into a warp sum. Start inactive data lanes at zero, then for XOR distances 1, 2, 4, 8, and 16, add the partner’s current subtotal to your own. After five steps every lane has the sum of its warp. With inputs 0…31 each lane should have 496. To reduce a whole block, save one total per warp in shared memory, synchronize the block, then let one warp combine those totals.

Coordinating two warps with shared memory

Run .\run.exe mirror.ptx mirror. This kernel reverses each fixed group of 64 elements. It treats missing elements in a short final group as zero, then writes only the first n logical outputs. In the first full block, thread 0 reads the value that thread 63 staged. Those threads are in different warps. A block-wide handoff makes that exchange possible.

The core of mirror.ptx is below. Each thread has a four-byte aligned slot in a 256-byte shared tile. A valid thread issues one asynchronous global-to-shared copy. A tail thread writes zero to its own slot.

@%valid  cp.async.ca.shared.global [%slot], [%src], 4;
@!%valid st.shared.u32 [%slot], 0;
cp.async.commit_group;
cp.async.wait_group 0;
bar.sync 0;

There are two waits to reason about. cp.async.wait_group 0 waits for the calling thread’s own committed copy groups. bar.sync 0 brings the block together before a peer reads the tile. After that handoff, each thread reads slot 63 - threadIdx.x. If you put this in a loop that reuses one tile, add another block handoff after the consumers finish and before a producer overwrites a slot. This cp.async form requires sm_80 or later. NVIDIA asynchronous data copies · PTX asynchronous copy instructions

Predict the tail. At n=70, the second block has valid threads 0 through 5. They read slots 63 through 58, all of which were filled with zero. The last six logical outputs must therefore be zero. To reverse only those six valid values, compute the current tile’s valid length and use validLength - 1 - threadIdx.x for valid consumers after the barrier. Change the CPU reference to match.

Build a real pipeline. First replace the copy with a normal load and shared store and keep the block handoff. Check the outputs. Next use two shared tiles. Stage tile 1 while consuming tile 0, wait before reading tile 1, and let all consumers finish before reusing tile 0. The teaching kernel waits immediately after issuing the copy, so its purpose is to establish correct ownership and completion. A speedup requires enough independent work or data movement to overlap. Larger tensor-memory transfers have their own completion protocols, such as transaction-counted mbarrier operations. Read the rules for the exact copy form when changing mechanisms. CUDA pipelines · CUDA asynchronous barriers

Atomic operations and memory ordering

Consider a histogram. Two threads load a bin containing 7, both add 1, and both write 8. The intended value was 9. An atomic add makes each read-modify-write on the bin indivisible for the participating threads at the chosen scope. atom gives the caller the previous value. red is useful for supported reductions when the old value is unnecessary. A heavily contended bin can still serialize updates. A common experiment is to build per-block histograms in shared memory and merge them into global bins. Compare both uniform input and an input that always selects bin zero. PTX memory consistency model

Publishing a payload uses a different property. Suppose a flag starts at zero, one producer writes a payload once, and one consumer checks whether it can read it. Here are the relevant fragments for two threads on the same GPU using global memory.

// Producer
st.global.u32 [%payload], 37;
st.release.gpu.global.u32 [%flag], 1;

// Consumer, in a different thread
ld.acquire.gpu.global.u32 %seen, [%flag];
setp.eq.u32 %ready, %seen, 1;
@%ready ld.global.u32 %result, [%payload];

If the acquire reads the value from the release, the consumer can use the preceding payload write. If it reads zero, this attempt consumes nothing. The example assumes aligned words, matching scopes, and no further writers. .cta covers a block, .gpu covers the device, and .sys broadens the scope for supported system participants and allocations. Choose a scope that includes both parties. A relaxed atomic counter suffices to count independent updates, but publishing other data requires an ordering protocol. An ordinary launch can schedule blocks in waves, so building a grid barrier by spinning on unscheduled blocks can deadlock. For two whole-grid phases, launch a second kernel in the same stream. CUDA memory model · CUDA programming model

Inline PTX and a Tensor Core tile

Inline PTX lets a CUDA C++ kernel own the control flow and memory management while you choose one PTX instruction. In this helper %0 is the output, %1 is the input, and r selects a 32-bit integer register. A literal PTX special register in the format string needs two percent signs, such as %%laneid.

__device__ unsigned countBits(unsigned x) {
    unsigned bits;
    asm("popc.b32 %0, %1;" : "=r"(bits) : "r"(x));
    return bits;
}

Use "=r" for write-only output and "+r" when the instruction also needs its old value. Match floating-point registers with "f" and 64-bit integers with "l". Enclose temporary PTX declarations in braces so repeated inlining keeps names local. For assembly with hidden memory effects, express the effects to the C++ compiler with the appropriate volatile and "memory" clobber. Inter-thread visibility still requires the device ordering operation appropriate to the algorithm. NVIDIA inline PTX assembly guide

Multiply one tile with a full warp

Build and run the fourth lab. It checks every position in a 16-by-8 matrix.

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

The operation is D = A × B + C. The shape m16n8k16 means A is 16×16, B is 16×8, and D is 16×8. The instruction uses FP16 input elements packed two per 32-bit register and FP32 accumulators. For this shape each of 32 lanes contributes four packed registers for A and two for B, and receives four FP32 outputs. The full warp executes the same mma.sync.aligned instruction in converged control flow. Treat the lane mapping as part of the instruction contract. PTX matrix instructions and fragment layouts

The lab makes A the identity, gives B a distinct value 8*row + col + 1 at each position, and sets C to zero. D must reproduce B, so a transposed or misplaced output is visible. Lane 5 has group = 5/4 = 1 and pair = 2*(5%4) = 2. It writes D at rows 1 and 9, columns 2 and 3. Predict 11, 12, 75, and 76. All 128 output values passed on the local GPU, and the binary contains HMMA.16816.F32. Exercise. Replace A with a matrix that reverses the row order. The first output row should now hold 121 through 128.

To turn this into a larger matrix multiply, tile the output by block, iterate over K tiles, stage inputs, and accumulate multiple MMA operations before storing. If you move fragments from shared memory with ldmatrix, derive lane addresses from the chosen load shape and transpose modifier. For newer instructions such as wgmma or tcgen05, check the exact variant’s target notes, operand layouts, issuer group, and completion rules. Support varies across GPU families. NVIDIA’s CUTLASS kernels show how these choices fit together. CUTLASS matrix multiplication · NVIDIA compute capabilities

Debugging and measuring a kernel

Start with the smallest input that separates two explanations. A failure at 1 often reveals a bad value or parameter. A failure at 65 after success at 64 points you toward a block boundary or invalid source lane. An overwritten output guard localizes a store past the logical end. The runner allocates its input separately and checks CUDA calls, including synchronization, so asynchronous errors surface near the launch.

compute-sanitizer --tool memcheck .\run.exe mirror.ptx mirror
compute-sanitizer --tool racecheck .\run.exe mirror.ptx mirror
compute-sanitizer --tool synccheck .\run.exe mirror.ptx mirror

Use memcheck for invalid or misaligned accesses, racecheck for shared-memory hazards, and synccheck for supported synchronization errors. Use initcheck when tracking uninitialized memory reads. Keep the CPU reference and boundary cases even when these tools are clean. A legal load from the wrong array element can still give the wrong answer. NVIDIA Compute Sanitizer

Performance experiment. Write a CUDA C++ version of the 64-element block mirror with the same zero-padding behavior. Compare it with the synchronous shared-memory PTX variant and the async variant on a large, identical input. Allocate once, warm up each version, and time repeated launches using CUDA events in the same stream. Synchronize the stop event. Report the median across batches, input size, block size, and GPU. Exclude allocations, transfers, and JIT when reporting kernel-only time. For N full elements, this operation requests 8*N logical bytes from global memory, so its effective bandwidth in decimal GB/s is 8*N / seconds / 1e9. For a bandwidth experiment, record your working set relative to L2 cache. CUDA Best Practices Guide on timing and bandwidth

Use ptxas -v for registers, shared memory, and spills. Compare SASS for differences in address calculation and load/store instructions. In Nsight Compute begin with Launch Statistics, Occupancy, and Memory Workload Analysis. More active warps can help hide latency, but forcing a lower register count can cause spills. Check the relevant traffic or stall counters and measure elapsed time again after each change. The tiny correctness cases in run.cpp are designed for finding indexing errors. Use the large-input harness for performance conclusions. Nsight Compute Profiling Guide

Capstone. Implement an out-of-place tiled transpose for a non-square matrix. Use a CPU transpose and a straightforward CUDA kernel as references. Verify 1×1, 31×33, 32×32, 33×31, and a large rectangle with distinct values. The output for input coordinate (r,c) must be stored at c*rows+r. Start with global loads and stores, introduce a shared tile with a uniform block barrier, then compare shared pitches of 32 and 33 words. Hand in one correct implementation, a lane-by-lane explanation of its addresses, a disassembly excerpt supporting your change, and a timing comparison. At that point you can take an unfamiliar instruction from the PTX manual, state its contract, test it, and decide whether it earns its place in a kernel.