CUB BlockReduce: Match Hand-Tuned Performance In Five Lines

You have already hand-built the block reduction: stage values in shared memory, fold them in log2(blockDim) barrier steps, then issue one atomic per block. Production CUDA code rarely writes that by hand. NVIDIA ships the same algorithm, tuned per architecture, as cub::BlockReduce in the CUB library (part of the CUDA toolkit). Replace the contended atomic starter with a BlockReduce call and match the hand-tuned kernel's throughput with a fraction of the code.

Languages: CUDA C++, Rust

Lesson

The reduction labs before this one exist so you understand what a block reduction costs and why it works. This lab teaches the production habit: reach for CUB (CUDA Core Compute Libraries) before hand-rolling a primitive. cub::BlockReduce<float, 256> is a class template specialized on the element type and block size. You give it a per-block scratch area (its TempStorage, declared in shared memory) and every thread passes its value to Sum(); the library performs the staged shared-memory folds - including architecture-specific choices like warp-shuffle final steps that are painful to tune by hand - and returns the block total to thread 0. Your kernel keeps only the parts that are genuinely yours: loading the value, and deciding what to do with the block total (here, one atomicAdd per block, exactly like the hand-rolled solution).

Two contracts matter. First, the block size in the template argument must match the launch (this lab pins Block X = 256). Second, BlockReduce's Sum() involves internal synchronization, so every thread of the block must call it - threads past the end of the array contribute 0.0f instead of returning early.

Equal SM throughput thresholds do not imply equal speed. Compare the handwritten and CUB versions on identical inputs and record their elapsed times.

What you will learn

Know that the CUDA toolkit ships tuned block/warp primitives (CUB) and when to prefer them over hand-rolled code.
Use cub::BlockReduce with a shared-memory TempStorage and a pinned block size.
Keep every thread on the collective-call path (contribute identity values instead of early-returning).
Confirm with the profiler that the library version matches the hand-tuned kernel's SM throughput.

Starter code

    // Correct but contended: one global atomic per element (the same starter
    // you fixed by hand two labs ago). This time, fix it with the library.
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        atomicAdd(out, X[i]);
    }