The CCCL team is excited to announce the 3.4 release of the CUDA Core Compute Libraries (CCCL). Highlights include support for CUDA Tile C++, major cub::DeviceScan performance improvements, convenient single-call overloads for all CUB device-wide algorithms, a new batched warp reduce, and cuda::std parallel STL algorithms.
CUDA Tile C++ Support
CCCL 3.4 includes support for using a subset of CCCL headers in CUDA Tile kernels. This first release includes support for fundamental device-side features like <cuda/std/type_traits> and <cuda/std/numeric_limits>. Additional features and headers, including CUB device-side cooperative algorithms, will be added in future releases.
#include <cuda/std/type_traits>
#include <cuda/std/limits>
template<typename T>
__tile_global__ void tile_kernel(T* out) {
static_assert(cuda::std::is_integral_v<T>);
out[0] = cuda::std::numeric_limits<T>::max();
}
See the CCCL CUDA Tile support documentation for the full availability table and restrictions, and the CUDA Tile C++ documentation for more about CUDA Tile.
Faster cub::DeviceScan on Blackwell
CCCL 3.4 comes with significant performance improvements for cub::DeviceScan on the NVIDIA Blackwell GPU architecture thanks to an all new warp-specialized implementation that leverages the Tensor Memory Accelerator (TMA).
CUB Single-Call Environment APIs
CCCL 3.4 completes the rollout of single-call, environment-based overloads for CUB device-wide algorithms.
Environment overloads remove the need to manually split a CUB call into a temporary-storage query phase and an execution phase. Instead, temporary storage is obtained from a memory resource in the execution environment. Streams, memory resources, and execution requirements can be composed into a single cuda::std::execution::env argument.
auto device = cuda::devices[0];
auto stream = cuda::stream{device};
auto pool = cuda::device_default_memory_pool(device);
auto env = cuda::std::execution::env{
cuda::stream_ref{stream},
pool
};
cub::DeviceReduce::Sum(d_in, d_out, num_items, env);
The traditional two-phase APIs are not deprecated and remain available for users who want explicit temporary-storage control.
For background, see Streamlining CUB with a Single-Call API. The full CUB environment API is documented in CUB Device-Wide Primitives.
New cub::WarpReduceBatched Warp Primitive
CCCL 3.4 adds cub::WarpReduceBatched, a new CUB warp-wide collective for reducing multiple independent batches of values partitioned across a warp. This can be up to 5x faster than issuing multiple independent cub::WarpReduce operations.
Thrust Iterator Migration to cuda:: Iterators
CCCL 3.4 continues consolidating shared iterator utilities under the cuda:: namespace.
The following Thrust iterator APIs are now deprecated in favor of their shared CCCL equivalents:
- thrust::constant_iterator -> cuda::constant_iterator
- thrust::make_constant_iterator -> cuda::make_constant_iterator
- thrust::strided_iterator -> cuda::strided_iterator
- thrust::make_strided_iterator -> cuda::make_strided_iterator
See the CCCL Fancy Iterators documentation.
Memory Resource Alignment Deprecation
CCCL 3.4 deprecates default-alignment overloads on type-erased memory resource wrappers.
Users should now pass allocation alignment explicitly when using resource wrappers such as cuda::mr::resource_ref, cuda::mr::synchronous_resource_ref, cuda::mr::any_resource, and cuda::mr::any_synchronous_resource.
See the Memory Resources documentation and cuda::resource_ref.
cuda::std Parallel STL Algorithms
CCCL 3.4 introduces support for the C++ Standard Library parallel algorithms model in cuda::std.
These algorithms are selected with the cuda::execution::gpu execution policy and operate on device-accessible ranges, bringing familiar standard algorithm interfaces to CUDA C++ while using CCCL/CUB machinery underneath.
See the cuda::std Algorithms Library documentation and the cuda::std parallel algorithms tracking issue.
Deprecations and Migration Notes
- thrust::constant_iterator and thrust::make_constant_iterator are deprecated. Use cuda::constant_iterator and cuda::make_constant_iterator.
- thrust::strided_iterator and thrust::make_strided_iterator are deprecated. Use cuda::strided_iterator and cuda::make_strided_iterator.
- thrust::integer_sequence, thrust::index_sequence, and related helpers are deprecated in favor of cuda::std equivalents where available.
- Default-alignment overloads on type-erased memory resource wrappers are deprecated. Pass alignment explicitly.
- cuda::arch_id relational operators <, <=, >, and >= are deprecated. Compare cuda::compute_capability values instead.
- cub::TexObjInputIterator is deprecated.
- cub::AlignBytes<T> is deprecated. Use alignof(T) directly.
- cuda::std::rel_ops is deprecated in C++20 mode.