Read the PTX of a Vector Add

Lesson

Every run on this site keeps the PTX and the SASS of your kernel. Run the starter once, open the run's code comparison view, and read the PTX from the top. The kernel begins with .visible .entry addSequential( followed by four .param lines: three .u64 pointers and one .u32 for N. Then come the .reg declarations, which name virtual registers by type: %r for 32-bit integers, %rd for 64-bit values such as addresses, %f for floats and %p for predicates. The numbers inside the angle brackets are only how many names the compiler used; the physical count ptxas settled on is the registers-per-thread value in the run's launch details.

The body is short enough to read line by line. ld.param.u64 fetches each pointer argument and ld.param.u32 fetches N. The index blockIdx.x * blockDim.x + threadIdx.x is three mov.u32 from the special registers %ctaid.x, %ntid.x and %tid.x and one mad.lo.s32. The bounds check if (i < N) is a setp.ge.s32 that writes a predicate, followed by @%p1 bra to the label at the end of the kernel: when the predicate is true the body is skipped. Inside the body each pointer passes through cvta.to.global.u64, the index is widened and scaled by mul.wide.s32 with the constant 4, the byte size of a float, and added to the base. Two ld.global.f32, one add.f32 and one st.global.f32 do the work. The label and ret close the kernel. The .loc lines in between name the source line each group of instructions came from; the comparison view uses them to tie listing to source.

Now change the kernel so each thread owns four consecutive elements, base = 4 * i through base + 3. Reinterpret the three pointers as float4 and load, add and store whole float4 values. In the new listing the multiplier in mul.wide.s32 becomes 16, each pair of scalar loads becomes one ld.global.v4.f32 with four destination registers listed in braces, the four add.f32 stay, and the store becomes one st.global.v4.f32. Two rules make the rewrite correct for every N. A 16-byte load needs a 16-byte aligned address, which holds here because cudaMalloc returns 256-byte aligned buffers and base is a multiple of four; a misaligned vector access is a fault, not a slow path. And the test sizes include 1000 and 4097, so when fewer than four elements remain the kernel must finish them with scalar instructions.

The launch changes with the kernel. The starter's grid expression ceilDiv(N, blockX) creates one thread per element. With four elements per thread, ceilDiv(N, blockX * 4) creates the right number; with the old expression the surplus threads exit at the bounds check and the answer is still right, but they are still scheduled.

What to expect in the numbers: the kernel moves the same bytes and touches the same sectors, so the SassMetrics pass reports the same sector count, and GPU time barely moves because a vector add is bound by memory bandwidth, not by instruction issue. What changes is the instruction count, a quarter as many loads and stores, and the SASS, where LDG.E.128 and STG.E.128 replace LDG.E and STG.E. This lab is judged on correctness and on those two vector instructions, which the verdict reads from your kernel's captured PTX once the run's listing has arrived; find them yourself in the listing too.

What you will learn

• Match each line of a kernel's PTX to the C++ statement that produced it: parameters, index arithmetic, the bounds predicate, loads, arithmetic and stores.
• Tell virtual registers (%r, %rd, %f, %p) from the physical register count ptxas reports.
• Produce a 128-bit vector load and store from CUDA C++ and recognize ld.global.v4.f32 and st.global.v4.f32 in the listing.
• Handle the tail and the alignment rule that make a vectorized kernel correct for every N.

Starter code

    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }