4 — Below the thread

Requirements: chapter 3. Experiments 1 to 3 need only the CUDA Toolkit, with no GPU. Experiment 4 needs a GPU.

The compiler

When you compile a .cu file with nvcc:

Host and device code are separated. Host code goes to your normal C++ compiler.
cudafe++ handles the split, and cicc generates and optimizes PTX, a virtual instruction set (HPC course).
ptxas compiles PTX into SASS, the machine code for one GPU architecture.

PTX has virtual registers, which ptxas maps to physical ones, spilling if there are too many (PTX ISA).

Targets look like sm_90 and sm_90a. sm_90 is the architecture, and sm_90a is a feature target with no backward or forward compatibility (CUDA guide). Normally a compiled binary is embedded in your program. The driver only JIT-compiles when PTX is loaded at runtime, and it caches the result.

How the compiler schedules

A CPU keeps much of its scheduling logic in silicon. A GPU moves part of that into the compiler. Each 128-bit SASS instruction reserves about 21 bits for scheduling, which ptxas fills in (TuringAs, maxas):

a stall count
a yield flag
a write barrier index, set by an instruction whose result arrives late
a read barrier index
a wait mask, listing the barriers to wait for

NVIDIA doesn’t publish this layout. It comes from reverse engineering, so treat exact bit positions as a working model. Older GPUs, from Kepler to Pascal, stored this in a separate control word shared by several instructions, which is why old guides read differently.

Programming the warp

Threads in a warp can copy register values from each other with no memory involved, using the __shfl_*_sync family:

__device__ inline float warpReduceSum(float v) {
    for (int off = 16; off > 0; off /= 2)
        v += __shfl_down_sync(0xffffffff, v, off);
    return v;  // lane 0 now holds the sum of all 32 threads
}

On the first pass, lane 0 adds lane 16’s value. On the second, it adds lane 8, which already holds a sum. Each pass halves the distance, so 32 values collapse in $\log_2 32 = 5$ steps.

Three details matter:

0xffffffff is a participant mask, saying which threads take part. It isn’t the active mask, which is a separate query, __activemask().
Reading an inactive thread is undefined, so this reduction is undefined in a partly filled warp.
These intrinsics don’t imply a memory barrier (CUDA guide).

Only the non-_sync forms are deprecated. __reduce_add_sync exists too, but it handles only integers, so it doesn’t replace this loop.

Moving data into shared memory

Kernels often copy tiles from global memory into fast shared memory. There are three generations of this:

Plain loads. Before Ampere, threads simply loaded data and stored it.
cp.async on Ampere (sm_80 and later), which copies 4, 8 or 16 bytes per instruction (PTX ISA).
The Tensor Memory Accelerator (TMA) on Hopper. NVIDIA calls it an asynchronous copy engine (Hopper Tuning Guide). One thread issues one cp.async.bulk.tensor instruction, and the hardware works out bounds and strides. Out-of-bounds values can be filled with zero or NaN.

If you’ve written drivers, think of programmed I/O versus DMA. It’s only an analogy, but the idea carries over: stop spending compute threads on moving bytes.

The tensor instructions differ by generation. Hopper uses wgmma (on sm_90a), and Blackwell uses tcgen05 (on sm_100a). CUTLASS warns that Blackwell kernels have “a substantially different warp-specialization structure” (CUTLASS). Warp specialization is called Spatial Partitioning in NVIDIA’s guide: some warps fetch data while others compute.

The memory model

CUDA memory is weakly ordered, and data races are undefined. Don’t use volatile to synchronize. Use cuda::atomic and the memory model in CCCL.

Hands-on

Experiment 1: from C++ to SASS

Needs: CUDA Toolkit, no GPU.

Create add.cu:

__global__ void add(float *A, float *B, float *C) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    C[i] = A[i] + B[i];
}

Compile it to PTX, then to an object file, and disassemble:

nvcc -ptx -arch=compute_90 add.cu -o add.ptx
nvcc -c -arch=sm_90 add.cu -o add.o
cuobjdump -sass add.o

Look for: ld.global.f32, add.f32 and st.global.f32 in the PTX, and LDG, FADD and STG in the SASS.
Compare with the _Z3addPfS_S_ example in NVIDIA’s Binary Utilities docs, which is sm_100. Yours will differ, so note how.

Experiment 2: see the real pipeline

Needs: CUDA Toolkit, no GPU.

List the steps without running them, then keep the intermediate files:

nvcc --dryrun -c add.cu
mkdir -p build && nvcc -c -keep -keep-dir build add.cu
ls build

Look for: cudafe++, cicc, ptxas and fatbinary, plus files such as .ptx and .cubin. Report what your toolkit actually prints.

Experiment 3: decode the scheduling bits

Needs: Python only.

Save this as decode.py. It uses the reverse-engineered layout from Step 2:

def decode(hi):
    return dict(stall=(hi >> 41) & 15, yld=(hi >> 45) & 1, wr=(hi >> 46) & 7,
                rd=(hi >> 49) & 7, wait=format((hi >> 52) & 63, "06b"),
                reuse=(hi >> 58) & 15)

print(decode(0x002ea2000c1e1900))   # LDG.E in NVIDIA's sm_100 sample
print(decode(0x004fca0000000000))   # FADD in NVIDIA's sm_100 sample

Run python3 decode.py.
Expected: the LDG.E should have wr=2, and the FADD wait mask should be 000100, which is barrier 2. So the add waits for the load. That’s the arithmetic result, worked out from NVIDIA’s published hex.
Repeat with instructions from your Experiment 1 output. For every wait bit, find the earlier instruction that set that barrier.

Experiment 4: a warp reduction

Needs: a GPU.

Create shfl.cu:

#include <cstdio>
#include <cuda_runtime.h>

__device__ inline float warpReduceSum(float v) {
    for (int off = 16; off > 0; off /= 2)
        v += __shfl_down_sync(0xffffffff, v, off);
    return v;
}

__global__ void reduce(float *out) {
    float s = warpReduceSum((float)(threadIdx.x + 1));
    if (threadIdx.x == 0) out[0] = s;
}

int main() {
    float *out, host;
    cudaMalloc(&out, sizeof(float));
    reduce<<<1, 32>>>(out);
    cudaMemcpy(&host, out, sizeof(float), cudaMemcpyDeviceToHost);
    printf("%f\n", host);
    cudaFree(out);
}

Build and run it, then count the shuffles:

nvcc -O3 shfl.cu -o shfl
./shfl
cuobjdump -sass shfl | grep SHFL

Expected: 1 + 2 + … + 32 = 528, so the program must print 528. Expect up to five SHFL instructions and no shared-memory loads in the kernel body.

TMA and wgmma experiments aren’t included. They need Hopper or Blackwell hardware, and the CUTLASS examples would be the place to start.