github NVIDIA/cccl v3.5.0
on GitHub

5 hours ago

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

Notable Fixes

Deprecations and Migration Notes

cub::ChainedPolicy, the CUB dispatch structs and the agent policy types listed below are deprecated. Custom tuning should use policy selectors through the execution environment of the corresponding public Device* API. Both single-call environment APIs and traditional two-phase temporary-storage APIs remain supported. See the tuning migration examples.

All type names in the table are in namespace cub.

Deprecated customization types Public tuning policy
DispatchAdjacentDifference, AgentAdjacentDifferencePolicy AdjacentDifferencePolicy
AgentBatchMemcpyPolicy BatchedCopyPolicy
DispatchHistogram, AgentHistogramPolicy HistogramPolicy
DispatchMergeSort, AgentMergeSortPolicy MergeSortPolicy
DispatchRadixSort, AgentRadixSortDownsweepPolicy, AgentRadixSortUpsweepPolicy, AgentRadixSortHistogramPolicy, AgentRadixSortExclusiveSumPolicy, AgentRadixSortOnesweepPolicy RadixSortPolicy
DispatchReduce, DispatchTransformReduce, AgentReducePolicy ReducePolicy
DispatchReduceByKey, AgentReduceByKeyPolicy ReduceByKeyPolicy
DeviceRleDispatch, AgentRlePolicy RleEncodePolicy or RleNonTrivialRunsPolicy, depending on the operation
DispatchScan, AgentScanPolicy ScanPolicy
DispatchScanByKey, AgentScanByKeyPolicy ScanByKeyPolicy
DispatchSegmentedRadixSort SegmentedRadixSortPolicy
DispatchSegmentedReduce, AgentWarpReducePolicy SegmentedReducePolicy
DispatchSegmentedSort, AgentSubWarpMergeSortPolicy SegmentedSortPolicy
DispatchSelectIf, AgentSelectIfPolicy SelectPolicy or PartitionPolicy, depending on the operation
DispatchThreeWayPartitionIf, AgentThreeWayPartitionPolicy ThreeWayPartitionPolicy
DispatchUniqueByKey, AgentUniqueByKeyPolicy UniqueByKeyPolicy

Full changelog from v3.4.0 to the 3.5 release branch

What's Changed

🚀 Thrust / CUB

📚 Libcudacxx

🔄 Other Changes

New Contributors

Full Changelog: v3.5.0.dev...v3.5.0

Check out latest releases or
releases around NVIDIA/cccl v3.5.0

Don't miss a new cccl release

NewReleases is sending notifications on new releases.

Get notifications