cccl
CCCL 2.6.1
| 撰稿人指南 | 开发容器 | 不和谐 | Godbolt | Github项目 | 文档 |
|---|
欢迎来到CUDA CORE COMPUTE库（CCCL），我们的任务是使Cuda更令人愉悦。
该存储库将三个必需的CUDA C ++库统计到一个方便的存储库中：
CCCL的目标是为CUDA C ++开发人员提供构建块，使编写安全有效的代码变得更加容易。将这些图书馆融合在一起，简化您的开发过程，并扩大您利用CUDA C ++力量的能力。有关决定统一这些项目的决定的更多信息，请参见此处的公告。
CUDA Core Compute库（CCCL）的概念从多年来独立开发的推力，CUB和LIBCUDACXX项目中有机地增长，具有类似的目标：为CUDA开发人员提供高质量，高性能和易于使用的C ++摘要。自然，这三个项目之间存在很多重叠，很明显，通过将它们统一到一个存储库中，社区将得到更好的服务。
推力是C ++并行算法库，它启发了向C ++标准库引入并行算法的引入。 Throust的高级接口极大地提高了程序员的生产率，同时通过可配置的后端启用GPU和多核心CPU之间的性能可移植性，这些后端允许使用多个并行编程框架（例如CUDA，TBB和OpenMP）。
幼崽是一个较低级别的CUDA特定库，旨在在所有GPU架构上进行光速并行算法。除了整个设备的算法外，它还还提供合作算法，例如宽范围的减少和宽范围的扫描，为CUDA内核开发人员提供构建块，以创建光速，自定义内核。
libcudacxx是CUDA C ++标准库。它提供了在主机和设备代码中使用的C ++标准库的实现。此外，它为CUDA特异性硬件功能提供了抽象，例如同步原语，缓存控制，原子等。
CCCL的主要目的是填补标准C ++库填充标准C ++的类似角色：为CUDA C ++开发人员提供通用，光速工具，从而使他们专注于解决重要的问题。统一这些项目是实现这一目标的第一步。
这是一个简单的示例，说明了推力，幼崽和libcudacxx的CCCL功能的使用。
它显示了如何使用推力/CUB/LIBCUDACXX实现简单的平行还原核。每个线程块使用cub::BlockReduce计算数组子集的总和。然后，使用libcudacxx的cuda::atomic_ref添加原子添加，将每个块的总和还原为单个值。
然后，它显示了如何使用推力的reduce算法进行相同的还原并比较结果。
试试它在Godbolt上！
# include < thrust/execution_policy.h >
# include < thrust/device_vector.h >
# include < cub/block/block_reduce.cuh >
# include < cuda/atomic >
# include < cuda/cmath >
# include < cuda/std/span >
# include < cstdio >
template < int block_size>
__global__ void reduce (cuda::std::span< int const > data, cuda::std::span< int > result) {
using BlockReduce = cub::BlockReduce< int , block_size>;
__shared__ typename BlockReduce::TempStorage temp_storage;
int const index = threadIdx. x + blockIdx. x * blockDim. x ;
int sum = 0 ;
if ( index < data. size ()) {
sum += data[ index ];
}
sum = BlockReduce (temp_storage). Sum (sum);
if (threadIdx. x == 0 ) {
cuda::atomic_ref< int , cuda::thread_scope_device> atomic_result (result. front ());
atomic_result. fetch_add (sum, cuda::memory_order_relaxed);
}
}
int main () {
// Allocate and initialize input data
int const N = 1000 ;
thrust::device_vector< int > data (N);
thrust::fill (data. begin (), data. end (), 1 );
// Allocate output data
thrust::device_vector< int > kernel_result ( 1 );
// Compute the sum reduction of `data` using a custom kernel
constexpr int block_size = 256 ;
int const num_blocks = cuda::ceil_div (N, block_size);
reduce<block_size><<<num_blocks, block_size>>>(cuda::std::span< int const >( thrust::raw_pointer_cast (data. data ()), data. size ()),
cuda::std::span< int >( thrust::raw_pointer_cast (kernel_result. data ()), 1 ));
auto const err = cudaDeviceSynchronize ();
if (err != cudaSuccess) {
std::cout << " Error: " << cudaGetErrorString (err) << std::endl;
return - 1 ;
}
int const custom_result = kernel_result[ 0 ];
// Compute the same sum reduction using Thrust
int const thrust_result = thrust::reduce (thrust::device, data. begin (), data. end (), 0 );
// Ensure the two solutions are identical
std::printf ( " Custom kernel sum: %d n " , custom_result);
std::printf ( " Thrust reduce sum: %d n " , thrust_result);
assert (kernel_result[ 0 ] == thrust_result);
return 0 ;
}CCCL中的所有内容都是仅标头的。因此，用户只需要关注他们如何获取标头文件以及如何将其整合到构建系统中。
使用CCCL开始使用CCCL的最简单方法是通过包括CCCL标头的CUDA工具包。使用nvcc编译时，它会自动将CCCL标头添加到包含路径中，因此您可以简单地#include代码中的任何CCCL标头，而无需其他配置。
如果使用另一个编译器进行编译，则需要更新构建系统的搜索路径，以指向CTK安装中的CCCL标头（例如， /usr/local/cuda/include ）。
# include < thrust/device_vector.h >
# include < cub/cub.cuh >
# include < cuda/std/atomic > 鼓励想要保持在CCCL开发前沿的用户使用GitHub的CCCL。支持使用较新版本的CCCL具有较旧版本的CUDA工具包，但不反过来。有关CCCL和CUDA工具包之间兼容性的完整信息，请参见我们的平台支持。
CCCL中的所有内容都是仅标头的，因此克隆并将其包括在一个简单的项目中与以下内容一样容易：
git clone https://github.com/NVIDIA/cccl.git
nvcc -Icccl/thrust -Icccl/libcudacxx/include -Icccl/cub main.cu -o main注意使用
-I而不是-isystem，以避免与CUDA工具包的nvcc隐式包含的CCCL标头发生冲突。所有CCCL标头都使用#pragma system_header来确保警告仍然会被沉默，就像使用-isystem一样，请参阅＃527，有关更多信息。
只能使用install cmake预设来配置仅生成安装规则的最小构建：
git clone https://github.com/NVIDIA/cccl.git
cd cccl
cmake --preset install -DCMAKE_INSTALL_PREFIX=/usr/local/
cd build/install
ninja install要在安装中包括实验库，请使用install-unstable预设和构建目录。
要仅安装实验库，请使用install-unstable-only预设并构建目录。
CCCL还通过conda-forge频道提供每个版本的Conda包：
conda config --add channels conda-forge
conda install cccl这将安装最新的CCCL到Conda环境的$CONDA_PREFIX/include/和$CONDA_PREFIX/lib/cmake/ Directories。 Cmake可以通过find_package(CCCL)发现它，并且可以在Conda环境中使用任何编译器。有关更多信息，请参阅此简介Conda-Forge。
如果要使用使用特定CUDA工具包的相同的CCCL版本，例如CUDA 12.4，则可以使用以下方式安装CCCL
conda config --add channels conda-forge
conda install cuda-cccl cuda-version=12.4 cuda-cccl MetaPackage安装了使用与cuda-version相对应的CUDA工具包的cccl版本。如果您希望在安装cuda-cccl后更新到最新的cccl ，请在更新cccl之前卸载cuda-cccl ：
conda uninstall cuda-cccl
conda install -c conda-forge cccl请注意，还有带有诸如
cuda-cccl_linux-64类的名称的Conda软件包。这些软件包包含CCCL版本作为CUDA工具包的一部分，但旨在由CUDA工具包进行内部使用。改为安装cccl或cuda-cccl，以与CONDA编译器兼容。有关更多信息，请参阅CCCL Conda-Forge食谱。
CCCL使用CMAKE用于所有构建和安装基础架构，包括测试以及在其他CMAKE项目中链接的目标。因此，CMAKE是将CCCL集成到另一个项目的推荐方法。
有关如何使用CMAKE软件包管理器执行此操作的完整示例，请参阅我们的基本示例项目。
其他构建系统应起作用，但仅测试CMAKE。欢迎简化将CCCL集成到其他构建系统中的贡献。
有兴趣促进使CCCL更好吗？查看我们的贡献指南，以全面概述设置开发环境，进行更改，运行测试并提交PR所需的一切。
目的：本节介绍用户可以期望CCCL能够成功地运行的地方。
通常，CCCL应在任何地方都可以使用CUDA工具包，但是，魔鬼在细节上。下面的各节描述了CUDA工具包，主机编译器和C ++方言的不同版本的支持和测试的详细信息。
概括：
鼓励CCCL用户利用最新的增强功能，并始终使用最新版本的CCCL来“居住”。为了获得无缝的体验，您可以独立于整个CUDA工具包升级CCCL。这是可能的，因为CCCL与当前和以前的主要版本系列中的每个次要CTK版本的最新补丁发行版保持向后兼容性。在某些例外情况下，CUDA工具包版本的最低支持的次要版本可能需要比其主要版本系列中最古老的版本更新。例如，由于CTK 11.0中存在不可避免的编译器问题，CCCL需要11.x系列中的最低支持版本的11.1版本。
当发布新的主要CTK时，我们放弃了对最古老的支持主要版本的支持。
| CCCL版本 | 支持CUDA工具包 |
|---|---|
| 2.x | 11.1-11.8，12.x（仅最新的补丁发布） |
| 3.x（未来） | 12.x，13.x（仅最新补丁发布） |
使用最新的CCCL的行为良好的代码应使用任何支持的CTK版本成功编译并成功运行。取决于新的CTK功能的新功能可能会出现例外，因此这些功能在CTK的较旧版本上不起作用。例如，直到CUDA 12.0之前，C ++ 20的支持才不会添加到nvcc中，因此依赖C ++ 20的CCCL功能将无法与CTK 11.x一起使用。
用户可以将较新版本的CCCL集成到较旧的CTK中，但反之亦然。这意味着CCCL的较旧版本与较新的CTK不兼容。换句话说， CCCL永远不会与CUDA工具包兼容。
下表总结了CTK和CCCL的兼容性：
| CTK版本 | 包括CCCL版本 | 所需的CCCL | 支持？ | 笔记 |
|---|---|---|---|---|
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR.MINOR+n | ✅ | 一些新功能可能不起作用 |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR+1.MINOR | ✅ | 可能的休息；一些新功能可能无法使用 |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR+2.MINOR | CCCL仅支持两个CTK主要版本 | |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR.MINOR-n | CCCL不兼容 | |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR-n.MINOR | CCCL不兼容 |
有关CCCL版本控制，API/ABI兼容性和破坏更改的更多信息，请参见下面的版本控制部分。
除非另有说明，否则CCCL支持与CUDA工具包相同的所有操作系统，这些操作系统已记录在这里：
除非另有说明，否则CCCL支持与最新的CUDA工具包相同的主机编译器，这些编译器已记录在此处：
当使用较旧的CUDA工具包时，我们还仅支持最新的CUDA工具包的主机编译器，但至少是任何受支持的旧CUDA工具包的最新主机编译器。
我们可能会保留其他编译器的支持，并将接受通过合理修复的社区的相应补丁。但是，我们不会花费大量时间来为较老的编译器分类或解决问题。
本着“您只支持您测试的内容”的精神，请参阅我们的CI概述，以获取有关我们测试的内容的更多信息。
除非另有说明，否则CCCL支持与CUDA工具包相同的所有GPU架构/计算功能，这些功能在此处进行了记录：https：//docs.nvidia.com/cuda/cuda-cuda-cuda-cuda-cuda-cuda-cuda-cuda-cuda-cuda-cuda-cuda-cuda-cuda-cuda-cuda-cuda-coda-colgramming-guide/index.html#compute-copute-capapute-compute-copute-capapiality
请注意，某些功能只能支持某些架构/计算功能。
CCCL的测试策略在测试尽可能多的配置与保持合理的CI时间之间取得了平衡。
对于CUDA工具包，对最古老和最新支持的版本进行了测试。例如，如果最新版本的CUDA工具包为12.3，则针对11.1和12.3进行测试。对于每个CUDA版本，构建均可完成所有受支持的主机编译器，并具有所有受支持的C ++方言。
测试策略和矩阵不断发展。 ci/matrix.yaml中定义的矩阵。YAML文件是真实的确定来源。有关CI管道的更多信息，请参见此处。
目的：本节介绍了CCCL的版本，API/ABI稳定性保证以及兼容指南，以最大程度地减少升级头痛。
概括
cub::或thrust::名称空间cuda::名称空间中符号的变化，可能会随时发生，但会通过将ABI版本嵌入到所有cuda::符号的内联命名空间中来反映。可以同时支持多个ABI版本。注意：将推力，CUB和LIBCUDACXX合并到此存储库之前，每个库都是根据语义版本独立版本的。从2.1版本开始，所有三个库都在单独的存储库中同步了其发行版本。向前迈进，CCCL将继续以单个语义版本发布，其中2.2.0是NVIDIA/CCCL存储库的第一个版本。
打破变化是更改已发布版本之间明确支持的功能，这将需要用户进行工作以升级到较新版本。
在极限上，任何变化都有可能将某人闯入某个地方。结果，并非所有可能的源打破变化都被认为是打破了对公共API的破坏，保证撞击主要语义版本。
以下各节描述了破坏CCCL API和ABI的变化的细节。
CCCL的公共API是有意暴露于提供图书馆实用程序的全部功能。
换句话说，CCCL的公共API不仅超越了功能签名，还包括（但不限于）：
此外，CCCL的公共API不包括以下任何一个：
_或__前缀为前缀的任何符号detail符号，包括detail::名称空间或宏detail/目录或子目录中包含的任何标头文件总的来说，目标是避免在公共API中破坏任何东西。仅当这些更改为用户提供更好的性能，更易于理解的API和/或更一致的API时才进行。
对公共API的任何破坏都将需要碰撞CCCL的主要版本编号。为了与CUDA MINIL版本的兼容性保持一致，API破坏变化和CCCL主要版本颠簸只会与CUDA工具包的新的主要版本发行相吻合。
没有任何警告的任何事情都可能随时改变公共API。
所有CCCL组件的公共API共享MAJOR.MINOR.PATCH的统一语义版本。
仅支持最近发布的版本。通常，功能和错误修复程序不会退回到先前发布的版本或分支。
查询该版本的首选方法是使用CCCL_[MAJOR/MINOR/PATCH_]VERSION如下所述。为了向后兼容，可以使用推力/CUB/LIBCUDACXXX版本定义，并且将始终与CCCL_VERSION一致。请注意，推力/幼崽使用MMMmmmpp方案，而cccl和libcudacxx则使用MMMmmmppp 。
| CCCL | libcudacxx | 推力 | 幼兽 | |
|---|---|---|---|---|
| 标题 | <cuda/version> | <cuda/std/version> | <thrust/version.h> | <cub/version.h> |
| 主要版本 | CCCL_MAJOR_VERSION | _LIBCUDACXX_CUDA_API_VERSION_MAJOR | THRUST_MAJOR_VERSION | CUB_MAJOR_VERSION |
| 次要版本 | CCCL_MINOR_VERSION | _LIBCUDACXX_CUDA_API_VERSION_MINOR | THRUST_MINOR_VERSION | CUB_MINOR_VERSION |
| 补丁/simbinor版本 | CCCL_PATCH_VERSION | _LIBCUDACXX_CUDA_API_VERSION_PATCH | THRUST_SUBMINOR_VERSION | CUB_SUBMINOR_VERSION |
| 串联版本 | CCCL_VERSION (MMMmmmppp) | _LIBCUDACXX_CUDA_API_VERSION (MMMmmmppp) | THRUST_VERSION (MMMmmmpp) | CUB_VERSION (MMMmmmpp) |
应用程序二进制接口（ABI）是一组规则：
图书馆的ABI包括但不限于：
ABI破裂的变化是任何会导致函数或键入公共API的ABI的变化。例如，将新数据成员添加到结构是ABI破坏变化，因为它会更改类型的大小。
在CCCL中，关于ABI的保证如下：
thrust::和cub::名称空间中的符号可能随时打破ABI而不会警告。thrust::和cub::符号的ABI包括用于编译的CUDA架构。因此thrust::如果cub::-x cu ）与C ++源（ -x cpp ）cuda::名称空间中的符号也可能随时破坏ABI。但是， cuda::符号嵌入了一个ABI版本号，每当发生ABI中断时，该编号会增加。可以同时支持多个ABI版本，因此用户可以选择恢复为先前的ABI版本。有关更多信息，请参见此处。谁应该关心阿比？
通常，CCCL用户只需要担心在构建或使用二进制文物（例如共享库）时，其API直接或间接包含CCCL提供的类型。
例如，考虑是否使用CCCL版本X构建libA.so ，其公共API包含一个函数：
void foo (cuda::std::optional< int >);如果使用CCCL版本Y编制了另一个库libB.so ，并使用libA.so中的foo ，那么如果版本X和Y之间存在ABI断裂，则可能会失败。与API断开更改不同，ABI断裂通常不需要代码更改，只需要重新编译所有使用相同的ABI版本即可。
要了解有关ABI的更多信息，为什么它很重要，看看ABI是什么，C ++应该怎么做？
如上所述，并非所有可能的源破坏变化都构成了破坏变化，这将需要增加CCCL的API主要版本编号。
鼓励用户遵守以下准则，以最大程度地降低意外造成的干扰风险，具体取决于不属于公共API的CCCL部分：
thrust:: cub:: ， nv:: ，或cuda::名称空间，例如，例如， cuda::std::iterator_traitsthrust:: cub:: ， cuda::或nv:: namespaces中获取任何API的地址。thrust:: ” cub:: ， cuda::或nv::名称空间中的任何API。_ ， __的前缀或名称中任何地方的detail ，包括detail:: namespace或acro#include声明该符号的标头文件。换句话说，不要依靠其他标题隐含包含的标头。本节的一部分灵感来自Abseil的兼容指南。
在对公共API，ABI或修改支持的平台和编译器进行任何破坏更改之前，我们将尽最大努力通知用户。
违反的折旧将以可以禁用的程序性警告的形式出现。
贬值期将取决于变更的影响，但通常至少持续2个次要版本。
即将推出！
有关CI管道的详细概述，请参见Ci-overview.md。
与CCCL使Cuda更令人愉悦的使命相关的项目：
您的项目使用CCCL吗？打开PR将您的项目添加到此列表中！
