NVIDIA/cccl v3.1.0

v3.1.0

Repository: NVIDIA/cccl

Tag: v3.1.0

Published: 2025-10-14T22:04:54Z

Prerelease: no

Release notes:

Highlights

New options for deterministic reductions in cub::DeviceReduce

Due to non-associativity of floating point addition, cub::DeviceReduce historically only guaranteed bitwise identical results run-to-run on the same GPU.

Starting with CCCL 3.1, formalizes three different levels of determinism with difference performance trade-offs

• Not-guaranteed (new!) - new single-pass reduction using atomics
• Run-to-run (status quo) - existing two-pass implementation
• GPU-to-GPU (new!) - based on reproducible reduction in @maddyscientis GTC 2024 talk

// Pick your desired trade-off of performance and determinism
// auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed);
// auto env = cuda::execution::require(cuda::execution::determinism::run_to_run);
// auto env = cuda::execution::require(cuda::execution::determinism::gpu_to_gpu);
cub::DeviceReduce::Sum(..., env);

More convenient single-phase CUB APIs

Nearly every CUB algorithm requires temporary storage for intermediate scratch space to carry out the algorithm.

Historically, it was the users responsibility to query and allocate the necessary temporary storage through a two-phase call pattern that is cumbersome and error-prone if arguments aren’t passed the same between two invocations.

CCCL 3.1 adds new overloads of some CUB algorithms that accept a memory resource so you skip the temp-storage query/allocate/free pattern.

Before

// determine temporary storage size
cub::DeviceScan::ExclusiveSum(d_temp_storage,
temp_storage_bytes,
nullptr, ...);

// Allocate the required temporary storage
cudaMallocAsync(&d_temp_storage,
temp_storage_bytes, stream);

// run the actual scan
cub::DeviceScan::ExclusiveSum(d_temp_storage,
temp_storage_bytes,
d_input...);

// Free the temporary storage
cudaFreeAsync(temp_storage, stream);

After

// Pool mr uses cudaMallocAsync under the hood
cuda::device_memory_pool mr{cuda::devices[0]};

// Single call. Temp storage is handled by the pool.
cub::DeviceScan::ExclusiveSum(d_input,..., mr);

What's Changed

🚀 Thrust / CUB

• par_nosync now uses async allocations by default.
• New reduce_into algorithm (PR #4355).
• Added strided_iterator (PR #4014).
• thrust::device_vector now supports default-init and skip-init constructors (PR #4183).
• New overloads for cub::WarpReduce (PR #3884).
• Tuned cub::ThreadReduce (PR #3441).

libcu++

• Added host/device/managed mdspan and accessors.
• New pointer utilities: is_aligned, align_up, align_down, ptr_rebind.
• New math utilities: ceil_ilog2, power-of-two helpers, fast_mod_div.
• New PTX primitive cuda::ptx::elect.sync.
• New warp primitive cuda::device::warp_match_all.
• Added compile-time loop utility cuda::static_for.

📚 Libcudacxx

• Enable device assertions in CUDA debug mode nvcc -G by @fbusato in https://github.com/NVIDIA/cccl/pull/4444
• avoid EDG bug by moving diagnostic push & pop out of templates by @ericniebler in https://github.com/NVIDIA/cccl/pull/4416
• Add host/device/managed mdspan and accessors by @fbusato in https://github.com/NVIDIA/cccl/pull/3686
• Add cuda::ptx::elect.sync by @fbusato in https://github.com/NVIDIA/cccl/pull/4445
• Add pointer utilities cuda::is_aligned, cuda::align_up, cuda::align_down, cuda::ptr_rebind by @fbusato in https://github.com/NVIDIA/cccl/pull/5037
• Add cuda::ceil_ilog2 by @fbusato in https://github.com/NVIDIA/cccl/pull/4485
• Add cuda::is_power_of_two, cuda::next_power_of_two, cuda::prev_power_of_two by @fbusato in https://github.com/NVIDIA/cccl/pull/4627
• Add...