What is CUB in CUDA?

Thrust's engine, exposed at block, warp and device scope so you can drop a tuned primitive inside your own kernel.

The scopes are the whole idea, and most people only ever find the first. Device scope (cub::DeviceReduce, cub::DeviceScan, cub::DeviceRadixSort) is a whole algorithm like Thrust's, but it takes raw pointers, uses a scratch buffer you own, and returns a cudaError_t instead of a value, so nothing synchronizes behind your back. Block scope (cub::BlockReduce, cub::BlockScan) and warp scope (cub::WarpScan) are different in kind: they go inside a kernel you still write. A reduction in the middle of a loop you already own cannot be a DeviceReduce call, because device-scope entry points are host functions that launch kernels; inside a kernel the answer is BlockReduce, and day 24's whole shared-memory tree becomes one line of it.

Device scope has a calling convention that trips everyone once. Every entry point is called twice: first with a null temp-storage pointer, which writes the byte count it needs and does no work, then with a real buffer. Miss the second call and your program reports success while the output keeps whatever was in it. The sizes are worth respecting too: on day 39's 64 MiB input, CUB asked for 3,583 bytes to reduce, 70,655 to scan, and 70,029,567 to sort. A radix sort wanting scratch about as large as its input is the kind of fact that changes a design.

CUB is fast because it is not generic where it matters: it carries a tuning policy per architecture, selected by the target you passed to nvcc, so the code that runs on a T4 is not the code that runs on an H100. One caveat from its own header: "Results are not deterministic for pseudo-associative operators (e.g., addition of floating-point types)" (https://github.com/NVIDIA/cccl/blob/v2.5.0/cub/cub/device/device_scan.cuh , checked 2026-08-30).

Measured

Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), CUB 2.5.0, built with nvcc -O3 -arch=sm_75. Captured 2026-08-30 on the project's verification node. Day 39 ran each algorithm the hand-written way and CUB's way on the same 16,776,605-element input in one process, checking every result:

algorithm	hand	CUB	vs hand
reduce	192.2 GB/s	DeviceReduce::Sum 264.4 GB/s	1.38
exclusive scan	100.8 GB/s	DeviceScan::ExclusiveSum 235.1 GB/s	2.33
radix sort	92.095 ms	DeviceRadixSort::SortKeys 4.484 ms	20.54

The sort gap is the big one and most of it is algorithmic rather than tuning: the hand sort deliberately moves one bit per pass, 32 passes over the keys, while CUB's sorts several bits at once and pays for a wider histogram per pass instead. A faster primitive cannot fix an algorithm doing the wrong amount of work, which is day 39's exercise in one sentence. The same run exercised block scope for correctness: a kernel using cub::BlockReduce plus a cuda::atomic_ref counted 8,388,303 odd keys on the device, and the host agreed exactly.

Code

From code/day39-cccl/cccl.cu, block scope inside a kernel the course wrote itself:

__global__ void countOddKeys(const unsigned int* __restrict__ keys,
                             int* __restrict__ total, size_t n) {
    using BlockReduce = cub::BlockReduce<int, kThreadsPerBlock>;
    __shared__ typename BlockReduce::TempStorage temp;

    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    const int mine = (i < n) ? static_cast<int>(keys[i] & 1u) : 0;
    const int blockSum = BlockReduce(temp).Sum(mine);

    if (threadIdx.x == 0) {
        cuda::atomic_ref<int, cuda::thread_scope_device> counter(*total);
        counter.fetch_add(blockSum, cuda::memory_order_relaxed);
    }
}

Related terms

• Thrust
• CCCL
• libcu++
• warp shuffle
• shared memory
• atomic contention

Where you meet this

• Day 39, when to stop hand-writing, the lesson that owns this term and produced the numbers above.
• Day 25, parallel reduction, the hand-written ladder BlockReduce replaces with one line.
• Day 31 and day 35, the scan and sort in the table.
• Day 26, atomics, where the block-then-atomic pattern in the code above was first built by hand.
• Colab setup, enough card to run the comparison yourself.

Sources

• CUB documentation, all three scopes: https://nvidia.github.io/cccl/unstable/cub/ (checked 2026-08-30)
• device_scan.cuh at v2.5.0, for the determinism note quoted above: https://github.com/NVIDIA/cccl/blob/v2.5.0/cub/cub/device/device_scan.cuh (checked 2026-08-30)

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. The numbers came off the verification node on 2026-08-30, and this entry stays a draft until a named author and a different named reviewer sign it.