Inline PTX: Warp Shuffle Reduction

Lesson

Read the driver tab first: this kernel writes one partial sum per block into partial[blockIdx.x], and the host adds those values to judge you. The cross-block step is not your job.

The starter is correct and serial. Thread 0 adds the block's elements one at a time while the other 255 threads skip the whole body. Its PTX is a loop you can read: a label, ld.global.f32, add.f32, an add.s32 on the counter, setp.lt.s32 and a backward @%p bra. The cost that the gate measures hides in the load. Memory arrives in 32-byte sectors, and one lane reading one float still fetches a whole sector, so the starter issues one sector per element, 4,210,688 of them at the profile size. The gate asks for at most 1,500,000.

The fix is three instructions, and you will write each of them in PTX. First, every lane loads its own element: v = (idx < N) ? X[idx] : 0.0f, which the compiler turns into a predicated ld.global.f32. Thirty-two lanes then read 128 consecutive bytes, four sectors, and every byte is used.

Second, the fold. shfl.sync.down.b32 d, a, delta, clamp, mask reads the 32-bit value a from the lane delta positions above the caller and writes it to d. clamp is 31 for a full warp: a lane whose source would fall off the top reads its own value instead, which the fold never uses. mask is 0xffffffff, the promise that all 32 lanes execute the instruction together. Five of these with delta 16, 8, 4, 2 and 1, each followed by an add, leave the warp's total in lane 0. The instruction is typed .b32 because it moves bits, not floats, so send __float_as_int(v) through an "r" input, receive the result in an "=r" output, and add __int_as_float of it. That is exactly what __shfl_down_sync does inside; its PTX is this instruction with the same clamp and mask.

Third, the join. A 256-thread block has eight warps, each holding its total in lane 0, and registers are private to a warp, so lane 0 of each warp writes its value to a small shared array. Then bar.sync 0: every thread of the block arrives, and shared-memory writes made before it are visible after it. That instruction is what __syncthreads() compiles to; write it as asm volatile("bar.sync 0;" ::: "memory") so the compiler knows memory changed across it. Warp 0 reads the eight values, lanes beyond the warp count read zero, folds them with the same five shuffles, and lane 0 writes partial[blockIdx.x]. Because a shuffle names a fixed lane, keep the block size a multiple of 32.

In the PTX pane, find your shuffles between the inline asm markers; after the run, find them in the captured listing by their opcode, with the matching SHFL.DOWN instructions in the SASS and BAR.SYNC where you wrote bar.sync 0. The SassMetrics sector count is what passes the gate: a correct warp-parallel load measures about 540,672 sectors, roughly 7.8 times fewer than the starter. One aside for integers: compute capability 8.0 and newer has redux.sync.add.s32, which folds a warp of integers in one instruction; a float sum still needs the shuffle ladder on every GPU in the fleet.

What you will learn

• Write shfl.sync.down.b32 by hand with its delta, clamp and member mask, and explain what each operand does.
• Move a float through a 32-bit shuffle as its bit pattern with __float_as_int and __int_as_float.
• Write the bar.sync that __syncthreads compiles to, with the memory clobber that makes it safe.
• Explain why the sector count falls when every lane loads, and why the shuffles themselves do not change it.

Starter code

    // Serial and wasteful: only thread 0 works while the rest of the block idles.
    int blockStart = blockIdx.x * blockDim.x;
    if (threadIdx.x == 0) {
        float sum = 0.0f;
        for (int i = 0; i < blockDim.x; ++i) {
            int idx = blockStart + i;
            if (idx < N) sum += X[idx];
        }
        partial[blockIdx.x] = sum;
    }