v3.5.0

CCCL 3.5 Release

The CCCL team is excited to announce the 3.5 release of the CUDA Core Compute Libraries (CCCL). Highlights include public tuning APIs for CUB, reproducible floating-point scans, batched Top-K selection, and reductions that accept problem sizes stored on the GPU.

Public CUB Tuning APIs

CCCL 3.5 exposes public tuning policies for CUB device-wide algorithms. Applications can customize parameters such as threads per block, items per thread, vectorization, and algorithm selection by passing a policy selector through cuda::execution::tune(...) in an execution environment.

Policy selectors choose tuning parameters for a GPU compute capability and return an algorithm-specific policy such as cub::ReducePolicy or cub::ScanPolicy. This provides a supported way to specialize CUB for a workload through the public device-wide APIs.

For example, this example policy below uses 256 threads per block and selects the number of items per thread based on compute capability:

struct ReduceTuning {
  __host__ __device__ constexpr cub::ReducePolicy
  operator()(cuda::compute_capability cc) const {
    auto pass = cub::ReducePassPolicy{
      .threads_per_block = 256,
      .items_per_thread = cc >= cuda::compute_capability{10, 0} ? 20 : 16,
      .vec_size = 4,
      .reduce_algorithm = cub::BLOCK_REDUCE_WARP_REDUCTIONS,
      .load_modifier = cub::LOAD_LDG
    };
    return {.multi_tile = pass, .single_tile = pass};
  }
};

auto env = cuda::std::execution::env{
  cuda::stream_ref{stream}, cuda::execution::tune(ReduceTuning{})
};
auto status = cub::DeviceReduce::Sum(input.begin(), output.begin(), input.size(), env);

See the CUB tuning documentation for an example and the CUB environment documentation for more information on how to use environments. Each CUB device-wide algorithm also has a documentation example for how to customize its tuning, see for example how to tune cub::DeviceReduce.

Reproducible Floating Point Scans

cub::DeviceScan now supports run-to-run reproducibility for floating-point summation. Requesting cuda::execution::determinism::run_to_run selects a fixed reduction order, producing repeatable results on the same GPU with the same input and build configuration.

This applies to inclusive and exclusive sums and scans using cuda::std::plus.

See the cub::DeviceScan documentation.

auto env = cuda::std::execution::env{
  cuda::stream_ref{stream},
  cuda::execution::require(cuda::execution::determinism::run_to_run)
};
auto status = cub::DeviceScan::InclusiveSum(input.begin(), output.begin(), input.size(), env);

Batched Top-K Selection

CCCL 3.5 adds cub::DeviceBatchedTopK, which selects the smallest or largest K items independently from many segments. MinKeys, MaxKeys, MinPairs, and MaxPairs support fixed or variable segment sizes and a K value that can vary by segment.

In CCCL 3.5, each segment is processed in one thread block, with a 48 KiB shared-memory limit. With the default policies, the maximum is 8,192 float keys or 4,096 double keys per segment; limits for key-value pairs depend on both types.

Callers must provide a compile-time upper bound on segment size (for example, cuda::args::bounds<1, 8192>() for float keys) and explicitly request non-deterministic, unordered output using determinism::not_guaranteed, tie_break::unspecified, and output_ordering::unsorted through cuda::execution::require(...).

Select the two largest keys per four-element segment from existing device buffers:

constexpr int segment_size = 4, k = 2;
auto segments_in = cuda::make_strided_iterator(
  cuda::make_counting_iterator(input.begin()), segment_size);
auto segments_out = cuda::make_strided_iterator(
  cuda::make_counting_iterator(output.begin()), k);

auto env = cuda::std::execution::env{
  cuda::stream_ref{stream},
  cuda::execution::require(
    cuda::execution::determinism::not_guaranteed,
    cuda::execution::tie_break::unspecified,
    cuda::execution::output_ordering::unsorted)
};
auto status = cub::DeviceBatchedTopK::MaxKeys(
  segments_in, segments_out,
  cuda::args::constant<segment_size>{}, cuda::args::constant<k>{}, num_segments, env);

See the CCCL 3.5 cub::DeviceBatchedTopK API for argument annotations, supported types, and complete examples.

CUB Arguments Stored on the GPU

The new <cuda/argument> header provides cuda::args::constant, immediate, deferred, and deferred_sequence, together with argument bounds. These annotations describe values known at compile time, values supplied by the host, and values read from device-accessible memory in stream order.

cub::DeviceReduce::{Reduce, Sum, Min, Max, TransformReduce} now accept a problem size supplied through cuda::args::deferred. A preceding kernel can produce the element count without a round trip to the host. Captured reductions can consume a different count on each CUDA Graph replay without updating or recapturing the graph.

Selection writes its output count to a one-element device buffer. The reduction consumes that count directly on the same stream:

auto selected = cuda::make_device_buffer<float>(stream, device, input.size(), cuda::no_init);
auto count = cuda::make_device_buffer<int>(stream, device, 1, cuda::no_init);
auto sum = cuda::make_device_buffer<float>(stream, device, 1, cuda::no_init);
auto env = cuda::std::execution::env{cuda::stream_ref{stream}};

auto status = cub::DeviceSelect::If(
  input.begin(), selected.begin(), count.begin(), input.size(), predicate, env);
status = cub::DeviceReduce::Sum(
  selected.begin(), sum.begin(), cuda::args::deferred{count.begin()}, env);

cub::DeviceScan::InclusiveScan also accepts an initial value supplied through cuda::args::deferred.

See the DeviceReduce documentation, the InclusiveScan addition, and #8789: support for device-resident problem sizes.

Faster Searches for Sorted Query Values

cub::DeviceFind::LowerBoundSortedValues and UpperBoundSortedValues accelerate batched lower- and upper-bound searches when both the searched range and the query values are sorted under the same comparator. They use a merge-path traversal with O(N + M) total work, where N is the range length and M is the number of queries.

See the cub::DeviceFind documentation.

Additional Library Improvements

• Added cuda::std::ranges::zip_view and cuda::std::views::zip, with expanded tuple-like construction and assignment support in cuda::std::tuple and cuda::std::pair.
• Added bit_compress, bit_expand, bit_repeat, and bit_reverse, plus shl and shr, to <cuda/std/bit>.
• Added cuda::isclose for approximate comparisons, cuda::ceil_ilog10, and cuda::device::warp_match_any for matching trivially copyable values across warp lanes.
• Added atomic_ref::address() to both cuda::atomic_ref and cuda::std::atomic_ref, and cuda::std::is_virtual_base_of where supported by the compiler.
• Added initializer-list overloads for buffer factories, additional memory-pool attribute queries with CUDA 13.3, and updated PTX wrappers for CUDA 13.4.
• Expanded CUB's use of hardware warp-reduction instructions, including floating-point min/max on supported Blackwell targets; added Programmatic Dependent Launch support to DeviceRadixSort and vectorized BlockLoad/BlockStore operations for contiguous iterators.
• Added GDB and LLDB pretty-printers for cuda::buffer, cuda::std::array, and cuda::std::complex and cuda::complex.
• C++23 extended cuda::std::tuple interface for tuple like types
• Various compile time improvements for type traits
• Added cuda::std::expected::has_error()
• Public macros to detect the host architecture CCCL_HOST_ARCH

Notable Fixes

• Fixed a race condition in DeviceTopK.
• Fixed out-of-bounds writes in DeviceHistogram for large inputs and in streaming reduce-by-key and run-length encoding for single-partition inputs.
• Fixed Thrust OpenMP scans processing unused block sums, and improved cuda::device_buffer iterator compatibility in Thrust and [DeviceAdjacentDifference](https://github.com...

v3.4.3

What's Changed

🔄 Other Changes

• Bump branch/3.4.x to 3.4.3. by @wmaxey in #10616
• [backport 3.4.x][cub] Always set smem limit to max for lookahead scan (#10570) by @davebayer in #10571
• [Backport branch/3.4.x] Fixes a race condition in DeviceTopK by @github-actions[bot] in #10682
• [Backport branch/3.4.x] Add inputs for controlling which repo and branch are checked out by @github-actions[bot] in #11008
• [Backport branch/3.4.x] [libcu++] Add cuda::ptx PTX 9.4 instructions by @fbusato in #11037

Full Changelog: v3.4.2...v3.4.3

CCCL Python Libraries (1.2.1)

CCCL Python Libraries (v1.2.1)

Previous release: v1.2.0.

These are the release notes for the cuda-cccl Python package version 1.2.1. This is a maintenance release with no new public Python APIs. The main changes are a self-contained Windows wheel, a workaround for a numba-cuda-mlir regression, clearer diagnostics for unsupported operations on structured data, and reduced Python-side overhead in cuda.compute calls.

Installation

Please refer to the install instructions here.

Packaging and dependencies

• Windows wheels now bundle msvcp140.dll (#11495)

  The Windows cuda-cccl wheel is repaired with delvewheel to include the MSVC C++ runtime DLL needed by cccl.c.parallel.dll. It no longer relies on whichever msvcp140.dll happens to be installed on the user's machine. The Windows cuda.compute test lanes now run without installing that DLL separately. Other packages, such as CuPy, may still have their own runtime requirements.

• numba-cuda-mlir 0.5.3 excluded (#11617)

  The cu12, cu13, sysctk12, and sysctk13 extras now require numba-cuda-mlir>=0.5.2,!=0.5.3,<0.6. Version 0.5.3 can fail when a stateful operator with the same state dtype and shape is compiled a second time in one process.

Bug fixes

• Clearer errors for unsupported built-in operations on structs (#11404)

  On the current V1 backend, using a built-in OpKind with a struct or other opaque storage type now raises a descriptive TypeError directing users to provide a custom operator, instead of failing later with a generic NVRTC compilation error. This does not change operations that use a built-in comparator on ordinary keys while carrying struct values.

Performance

• Less per-call overhead in cuda.compute (#11497)

  Shared iterator handling now caches whether an iterator represents a pointer and reuses its Pointer wrapper. Device-array pointer lookup also caches the appropriate accessor for each array type. These changes apply across algorithms without changing their public APIs or results. The PR measured approximately 9–20% less host-call overhead in representative reduce, scan, segmented-reduce, and histogram benchmarks; those measurements exclude GPU kernel execution.

Notes

• v1.2.1 includes the Python-package changes since v1.2.0; there were no intervening cuda-cccl releases.

CCCL Python Libraries (1.2.0)

CCCL Python Libraries (v1.2.0)

Previous release: v1.1.1.

These are the release notes for the cuda-cccl Python package version 1.2.0.

The headline changes in this release are: bfloat16 support in cuda.compute and free-threaded (3.14t) wheel support for the cuda-compute-minimal extra. Under the hood, cuda.compute's JIT backend has migrated from numba-cuda to numba-cuda-mlir, with no public API changes.

Installation

Please refer to the install instructions here.

Packaging / breaking changes

• cuda.coop._experimental removed from cuda-cccl (#10104)

  The experimental cooperative-primitives module (cuda.coop, including cuda.coop._experimental) has been removed from cuda-cccl wheels entirely — implementation, tests, benchmarks, examples, docs, and CI jobs. It depended on the classic numba-cuda JIT path, which is not compatible with free-threaded Python, and removing it clears the way for publishing free-threaded cuda.compute wheels. A significantly improved cuda.coop is in development, targeted for a future release. This does not affect cuda.compute, cuda.cccl, or any C++ cooperative APIs (CUDA Experimental, CUB).

• cuda.compute's JIT backend migrated from numba-cuda to numba-cuda-mlir (#9421)

  User-operator and gpu_struct compilation now goes through numba-cuda-mlir, the MLIR-based successor to numba-cuda. This is an internal backend swap with no change to cuda.compute's public API or performance (compile-time or runtime). Installing cuda-cccl[cu12]/[cu13] now pulls in numba-cuda-mlir instead of numba-cuda.

Features

• bfloat16 support in cuda.compute (#10933)

  Algorithms in cuda.compute now accept bfloat16 arrays (_nv_bfloat16 in CUDA C++). NumPy has no native bfloat16 dtype, so this relies on the ml_dtypes package, which is not a required dependency — install it yourself if you want to use bfloat16. Verified compatible with CuPy arrays and PyTorch tensors. Python callables cannot be JIT-compiled for bfloat16 (Numba has no support), so operators must be either well-known OpKind operations or pre-compiled device code.

  import numpy as np
  import ml_dtypes
  import cuda.compute as cc
  
  BFLOAT16 = np.dtype(ml_dtypes.bfloat16)
  h_in = np.array([1.0, 2.0, 3.0], dtype=BFLOAT16)
  # reduce, scan, sort, transform, histogram, etc. now accept bfloat16 arrays
  cc.reduce_into(d_in=h_in, d_out=..., op=cc.OpKind.PLUS, h_init=np.array([0], dtype=BFLOAT16))

• Free-threaded (Python 3.14t) wheel support (#9475, #11311, #10803, #10789, #11311)

  cuda.compute is now thread-safe and can run with the GIL disabled.

• reduce_into accepts h_init=None (#10953)

  reduce_into no longer requires an explicit identity/init value — pass h_init=None to reduce using the first element as the seed, mirroring cub::DeviceReduce's no-init overload. (Not supported together with Determinism.NOT_GUARANTEED.)

  cc.reduce_into(d_in=d_in, d_out=d_out, op=cc.OpKind.MAX, h_init=None)

Bug Fixes

• Fixed TransformIterator ignoring stateful-op state (#11213) — TransformIterator previously composed its state bytes from only the underlying iterator's state, dropping the state of a stateful op. It now combines both correctly.

Internal / CI

• Simplified the Python CMake configuration (#9883).
• Replaced device-array usage outside examples with a wrapper independent of CuPy/numba-cuda (#9653).
• Added CI coverage: thread-sanitizer job for c.parallel (#9986), pytest-run-parallel (#9886), benchmark smoke tests (#9885), tests against the system CTK instead of a pip-installed one (#10500), running Python test payloads in a minimal sibling container (#11018).
• Enabled precompiled headers for cuda.compute HostJIT builds (#10628); split the Clang/LLD HostJIT library out separately (#9583); serialized HostJIT builds and improved related tests (#9998).
• Fixed mypy cache_dir handling and re-enabled mypy on pre-commit.ci (#10876); applied mypy fixes to benchmark infra scripts (#11099).
• pre-commit.ci autoupdate (#10729); fixed broken links to the Python docs on github.io (#10585).
• Fixed an overflow bug in the segmented-reduce test suite (#11327).

Notes

• v1.2.0 includes all commits from v1.1.1 plus the above; there were no releases between v1.1.1 and v1.2.0.

v3.4.2

What's Changed

🔄 Other Changes

• Bump branch/3.4.x to 3.4.2. by @wmaxey in #10524
• Disable Tile compilation mode by @jrhemstad in #10522

Full Changelog: v3.4.1...v3.4.2

v3.4.1

What's Changed

🔄 Other Changes

• [Backport to 3.4] Align tile state for warpspeed scan by @bernhardmgruber in #9781
• [Backport branch/3.4.x] Do not vectorize large type reductions by @github-actions[bot] in #9775
• [Backport branch/3.4.x] Update __cccl_ptx_isa by @fbusato in #9797
• [Backport branch/3.4.x] Remove _CCCL_GRID_CONSTANT from DeviceSelectSweepKernel parameters by @github-actions[bot] in #9799
• Bump branch/3.4.x to 3.4.1. by @wmaxey in #9721
• Fix mypy CI error on 3.4 branch by @bernhardmgruber in #9849
• [Backport to 3.4] Use plus tunings for lookahead scan more widely (#9813) by @bernhardmgruber in #9848
• [Backport to 3.4] Fix scanning out-of-bounds items in OpenMP scan (#9759) by @bernhardmgruber in #9819
• [Backport to 3.4] Remove most _CCCL_GRID_CONSTANT annotations, including merge sort by @bernhardmgruber in #9864

Full Changelog: v3.4.0...v3.4.1

CCCL Python Libraries (1.1.1)

CCCL Python Libraries (v1.1.1)

Previous release: v1.1.0.

These are the release notes for the cuda-cccl Python package version 1.0.1, dated July 16th, 2026.

Bug Fixes / Packaging

• Fixed an issue with the new warpspeed scan on windows + sm_120

CCCL Python Libraries (v1.1.0)

CCCL Python Libraries (v1.1.0)

Previous release: v1.0.2.

These are the release notes for the cuda-cccl Python package version 1.1.0.

The headline of this release is support for serialization and ahead-of-time (AoT) compilation of cuda.compute algorithm objects. There are no breaking changes to the public API.

Installation

Please refer to the install instructions here.

Features

• Serialization of cuda.compute algorithm objects (#9644)

  New top-level cuda.compute.serialize() and cuda.compute.deserialize() functions let you turn a built algorithm object into bytes and reconstruct it later. These are low-level building blocks for ahead-of-time compilation, on-disk caches, and cross-node communication of algorithm objects.

  Compile once and write to disk:

  import cuda.compute as cc, cupy as cp
  
  d_in = cp.empty(1)
  d_out = cp.empty(1)
  op = lambda x: 2 * x
  transformer = cc.make_unary_transform(d_in=d_in, d_out=d_out, op=op)
  with open("transform.cclb", "wb") as f:
      f.write(cc.serialize(transformer))

  Load and run from a subsequent process (no rebuild):

  import cuda.compute as cc, cupy as cp
  
  with open("transform.cclb", "rb") as f:
      transformer = cc.deserialize(f.read())
  
  d_in = cp.asarray([1., 2, 3])
  d_out = cp.empty_like(d_in)
  transformer(d_in=d_in, d_out=d_out, op=lambda x: 2 * x, num_items=len(d_in))
  # d_out == [2., 4., 6.]

• Ahead-of-time (AoT) compilation for multiple compute capabilities, including GPU-less builds (#9732)

  You can now compile cuda.compute algorithms ahead of time for several compute capabilities at once, and on machines that have no GPU at all. This builds on separate compile and load steps in the underlying C layer (#8484), plus serialize()/deserialize() support in cccl.c (#9568). Documentation and examples for the serialization / AoT workflows were added as part of this work.

Bug Fixes / Performance

• Fixed build-cache misses for closures over Python scalars (#9680, closes #9626) — an op that closed over a plain Python int/float/bool was JIT-rebuilt on every call because the cache keyed those scalars by id(). They are now keyed by value.

• Relaxed the histogram build-cache key (#9596, closes #9594) — avoids unnecessary recompilations.

• Fixed the v2 (host-JIT) backend rejecting some well-known operators (#9649).

Packaging

• Pinned numba < 0.66 in the CUDA extras (cuda-cccl[cu12] / [cu13]) to work around a temporary incompatibility with numba.cuda.types.NPDatetime (#9692).

Internal / CI

• Added benchmarks to measure cuda.compute host-side overhead (#9432).
• Added a CI job for the minimal cuda-cccl extra and explicit no-numba tests, decoupling more of the test suite from numba.cuda in favor of CuPy / cuda.core (#9434).

Notes

• v1.1.0 also includes all fixes previously shipped in the 1.0.1 and 1.0.2 patch releases (e.g. NVRTC 13.3 warning fix #9171, reduce_into type-mismatch fix #9206, building against CCCL C v2 #9200, and numpy-scalar cache fix #9469).

v3.4.0

The CCCL team is excited to announce the 3.4 release of the CUDA Core Compute Libraries (CCCL). Highlights include 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.

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.

CCCL Python Libraries (1.0.2)

CCCL Python Libraries (v1.0.2)

Previous release: v1.0.1.

These are the release notes for the cuda-cccl Python package version 1.0.1, dated June 16th, 2026.

Bug Fixes / Packaging

• Fixed an issue with numpy scalars leading to unnecessary recompilations with stateful cuda.compute ops