Skip to content

v3.1.0

Choose a tag to compare

View all tags
@github-actions github-actions released this 14 Oct 22:04
· 2868 commits to main since this release
v3.1.0
ecfd3ad
This commit was created on GitHub.com and signed with GitHub’s verified signature.
GPG key ID: B5690EEEBB952194
Verified
Learn about vigilant mode.

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);
image image

  Not-Guaranteed (new!) Run-to-run (status quo) GPU-to-GPU (new!)
Determinism Varies per run Varies per GPU Constant
Performance Best Better Good

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

libcu++

📚 Libcudacxx

📝 Documentation

  • Making extended API documentation slightly more uniform by @fbusato in #4965
  • Add memory space note to cuda::memory documentation by @fbusato in #5151
  • Better specify lane_mask::all_active() behavior by @fbusato in #5183

🔄 Other Changes

New Contributors

Full Changelog: v3.0.3...v3.1.0

Contributors

cryos, alliepiper, and 35 other contributors
Assets 6
Loading