贡献者指南	开发容器	Discord	Godbolt	GitHub 项目	文档

CUDA C++ 核心库（CCCL）

欢迎来到 CUDA C++ 核心库（CCCL），我们的使命是让 CUDA C++ 更加令人愉悦。

这个仓库将三个重要的 CUDA C++ 库统一到一个方便的仓库中：

CCCL 的目标是为 CUDA C++ 开发者提供构建模块，使编写安全高效的代码变得更加容易。将这些库整合在一起简化了您的开发过程，并拓宽了您利用 CUDA C++ 能力的范围。有关统一这些项目决定的更多信息，请参阅此处的公告。

概述

CUDA C++ 核心库（CCCL）的概念源于 Thrust、CUB 和 libcudacxx 项目，这些项目多年来独立开发，但都有相似的目标：为 CUDA 开发者提供高质量、高性能且易于使用的 C++ 抽象。自然地，这三个项目之间存在很多重叠，将它们统一到一个仓库中显然能更好地服务于社区。

Thrust 是 C++ 并行算法库，它启发了 C++ 标准库引入并行算法。Thrust 的高级接口大大提高了程序员的生产力，同时通过可配置的后端实现了 GPU 和多核 CPU 之间的性能可移植性，允许使用多种并行编程框架（如 CUDA、TBB 和 OpenMP）。
CUB 是一个较低级别的、专门用于 CUDA 的库，旨在为所有 GPU 架构提供极速并行算法。除了设备范围的算法外，它还提供了诸如块级归约和线程束级扫描等协作算法，为 CUDA 内核开发者提供了创建极速自定义内核的构建模块。
libcudacxx 是 CUDA C++ 标准库。它提供了一个可在主机和设备代码中使用的 C++ 标准库实现。此外，它还为 CUDA 特定的硬件功能提供了抽象，如同步原语、缓存控制、原子操作等。

CCCL 的主要目标是扮演类似于标准 C++ 库对标准 C++ 的角色：为 CUDA C++ 开发者提供通用、极速的工具，使他们能够专注于解决真正重要的问题。统一这些项目是实现这一目标的第一步。

示例

这是一个简单的示例，演示了如何使用来自 Thrust、CUB 和 libcudacxx 的 CCCL 功能。

它展示了如何使用 Thrust/CUB/libcudacxx 实现一个简单的并行归约内核。每个线程块使用 cub::BlockReduce 计算数组子集的总和。然后使用 libcudacxx 中的 cuda::atomic_ref 通过原子加法将每个块的总和归约为单个值。

接着，它展示了如何使用 Thrust 的 reduce 算法完成相同的归约，并比较结果。

在 Godbolt 上实时尝试！

#include <thrust/execution_policy.h>
#include <thrust/device_vector.h>
#include <cub/block/block_reduce.cuh>
#include <cuda/atomic>
#include <cstdio>

constexpr int block_size = 256;

__global__ void reduce(int const* data, int* result, int N) {
  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 < N) {
    sum += data[index];
  }
  sum = BlockReduce(temp_storage).Sum(sum);

  if (threadIdx.x == 0) {
    cuda::atomic_ref<int, cuda::thread_scope_device> atomic_result(*result);
    atomic_result.fetch_add(sum, cuda::memory_order_relaxed);
  }
}

int main() {

  // 分配并初始化输入数据
  int const N = 1000;
  thrust::device_vector<int> data(N);
  thrust::fill(data.begin(), data.end(), 1);

  // 分配输出数据
  thrust::device_vector<int> kernel_result(1);

  // 使用自定义内核计算 `data` 的总和归约
  int const num_blocks = (N + block_size - 1) / block_size;
  reduce<<<num_blocks, block_size>>>(thrust::raw_pointer_cast(data.data()),
                                     thrust::raw_pointer_cast(kernel_result.data()),
                                     N);

  auto const err = cudaDeviceSynchronize();
  if (err != cudaSuccess) {
    std::cout << "错误：" << cudaGetErrorString(err) << std::endl;
    return -1;
  }

  int const custom_result = kernel_result[0];

  // 使用 Thrust 计算相同的总和归约
  int const thrust_result = thrust::reduce(thrust::device, data.begin(), data.end(), 0);

  // 确保两种解决方案的结果相同
  std::printf("自定义内核总和：%d\n", custom_result);
  std::printf("Thrust reduce 总和：%d\n", thrust_result);
  assert(kernel_result[0] == thrust_result);
  return 0;
}

入门

用户

CCCL 中的所有内容都是仅头文件的。因此，用户只需关注如何获取头文件以及如何将它们纳入自己的构建系统。

CUDA 工具包

开始使用 CCCL 最简单的方法是通过 CUDA 工具包，它包含了 CCCL 头文件。当您使用 nvcc 编译时，它会自动将 CCCL 头文件添加到您的包含路径中，因此您可以在代码中直接 #include 任何 CCCL 头文件，无需额外配置。

如果使用其他编译器，您需要更新构建系统的包含搜索路径，以指向 CTK 安装中的 CCCL 头文件（例如，/usr/local/cuda/include）。

#include <thrust/device_vector.h>
#include <cub/cub.cuh>
#include <cuda/std/atomic>

GitHub

希望紧跟 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，以避免与 nvcc 从 CUDA Toolkit 中隐式包含的 CCCL 头文件发生冲突。所有 CCCL 头文件都使用 #pragma system_header 来确保警告仍会被静默处理，就像使用 -isystem 一样，详情请参见 https://github.com/NVIDIA/cccl/issues/527。

Conda

CCCL 还通过 conda-forge 频道提供每个发布版本的 conda 包：

conda config --add channels conda-forge
conda install cccl

这将把最新的 CCCL 安装到 conda 环境的 $CONDA_PREFIX/include/ 和 $CONDA_PREFIX/lib/cmake/ 目录中。它可以通过 CMake 的 find_package(CCCL) 被发现，并可以被 conda 环境中的任何编译器使用。更多信息，请参见这篇 conda-forge 介绍。

如果你想使用与特定 CUDA Toolkit 一起发布的相同 CCCL 版本，例如 CUDA 12.4，你可以这样安装 CCCL：

conda config --add channels conda-forge
conda install cuda-cccl cuda-version=12.4

cuda-cccl 元包会安装与 cuda-version 对应的 CUDA Toolkit 一起发布的 cccl 版本。如果你想在安装 cuda-cccl 后更新到最新的 cccl，请在更新 cccl 之前卸载 cuda-cccl：

conda uninstall cuda-cccl
conda install -c conda-forge cccl

注意还有一些名称类似 cuda-cccl_linux-64 的 conda 包。这些包包含作为 CUDA Toolkit 一部分发布的 CCCL 版本，但它们是为 CUDA Toolkit 的内部使用而设计的。请安装 cccl 或 cuda-cccl，以兼容 conda 编译器。更多信息，请参见 cccl conda-forge 配方。

CMake 集成

CCCL 使用 CMake 进行所有构建和安装基础设施，包括测试以及在其他 CMake 项目中链接的目标。因此，CMake 是将 CCCL 集成到其他项目中的推荐方式。

有关如何使用 CMake 包管理器实现这一点的完整示例，请参见我们的示例项目。

其他构建系统应该也能工作，但只有 CMake 经过测试。欢迎贡献简化 CCCL 集成到其他构建系统的方法。

贡献者

有兴趣为改进 CCCL 做出贡献吗？查看我们的贡献指南，了解设置开发环境、进行更改、运行测试和提交 PR 所需的全面概述。

平台支持

目标： 本节描述用户可以期望 CCCL 在哪里成功编译和运行。

一般来说，CCCL 应该在 CUDA Toolkit 支持的任何地方都能工作，但细节是关键。以下部分描述了对不同版本的 CUDA Toolkit、主机编译器和 C++ 方言的支持和测试详情。

CUDA Toolkit (CTK) 兼容性

总结：

最新版本的 CCCL 向后兼容当前和前一个 CTK 主要版本系列
CCCL 永远不会向前兼容任何版本的 CTK。始终使用与您的 CTK 相同或更新的版本
次要版本的 CCCL 升级不会破坏现有代码，但新功能可能不支持所有 CTK 版本

鼓励 CCCL 用户通过始终使用最新版本的 CCCL 来利用最新的增强功能并"活在前沿"。为了获得无缝体验，您可以独立于整个 CUDA Toolkit 升级 CCCL。这是可能的，因为 CCCL 保持向后兼容当前和前一个主要版本系列中每个次要 CTK 发布的最新补丁版本。在某些特殊情况下，CUDA Toolkit 发布的最低支持次要版本可能需要比其主要版本系列中最旧的版本更新。例如，由于 CTK 11.0 中存在不可避免的编译器问题，CCCL 要求 11.x 系列的最低支持版本为 11.1。

当发布新的主要 CTK 时，我们会放弃对最旧支持的主要版本的支持。

CCCL 版本	支持的 CUDA Toolkit 版本
2.x	11.1 - 11.8, 12.x（仅最新补丁版本）
3.x（未来）	12.x, 13.x（仅最新补丁版本）

使用最新 CCCL 的行为良好的代码应该能够在任何受支持的 CTK 版本上成功编译和运行。例外情况可能发生在依赖新 CTK 功能的新特性上，因此这些特性在较旧版本的 CTK 上不会工作。例如，C++20 支持直到 CUDA 12.0 才被添加到 nvcc，所以依赖 C++20 的 CCCL 特性不会在 CTK 11.x 上工作。

用户可以将较新版本的 CCCL 集成到较旧的 CTK 中，但反之则不行。这意味着较旧版本的 CCCL 与较新的 CTK 不兼容。换句话说，CCCL 永远不会向前兼容 CUDA Toolkit。

下表总结了 CTK 和 CCCL 的兼容性：

CTK 版本	包含的 CCCL 版本	期望的 CCCL	是否支持？	注释
CTK `X.Y`	CCCL `MAJOR.MINOR`	CCCL `MAJOR.MINOR+n`	✅	一些新特性可能无法工作
CTK `X.Y`	CCCL `MAJOR.MINOR`	CCCL `MAJOR+1.MINOR`	✅	可能会有破坏性变更；一些新特性可能不可用
CTK `X.Y`	CCCL `MAJOR.MINOR`	CCCL `MAJOR+2.MINOR`	❌	CCCL 只支持两个 CTK 主要版本
CTK `X.Y`	CCCL `MAJOR.MINOR`	CCCL `MAJOR.MINOR-n`	❌	CCCL 不向前兼容
CTK `X.Y`	CCCL `MAJOR.MINOR`	CCCL `MAJOR-n.MINOR`	❌	CCCL 不向前兼容

有关 CCCL 版本控制、API/ABI 兼容性和破坏性变更的更多信息，请参见下面的版本控制部分。

操作系统

除非另有说明，CCCL 支持与 CUDA Toolkit 相同的所有操作系统，这些都记录在这里：

主机编译器

除非另有说明，CCCL 支持与 CUDA Toolkit 相同的所有主机编译器，这些都记录在这里：

Linux
Windows 秉承"你只支持你测试过的内容"的精神，有关我们具体测试内容的更多信息，请参阅我们的CI概述。

C++方言

C++11（在Thrust/CUB中已弃用，将在下一个主要版本中移除）
C++14（在Thrust/CUB中已弃用，将在下一个主要版本中移除）
C++17
C++20

GPU架构

除非另有说明，CCCL支持与CUDA工具包相同的所有GPU架构/计算能力，相关文档可在此处找到：https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capability

请注意，某些功能可能仅支持特定架构/计算能力。

测试策略

CCCL的测试策略在尽可能测试多种配置和保持合理的CI时间之间取得平衡。

对于CUDA工具包版本，我们会对最旧和最新的支持版本进行测试。例如，如果CUDA工具包的最新版本是12.3，则会对11.1和12.3进行测试。对于每个CUDA版本，我们会针对所有支持的主机编译器和所有支持的C++方言完成构建。

测试策略和矩阵不断演进。ci/matrix.yaml文件中定义的矩阵是最终的真实来源。有关我们CI管道的更多信息，请参见此处。

版本控制

目标： 本节描述CCCL的版本控制方式、API/ABI稳定性保证以及兼容性指南，以最大程度减少升级带来的麻烦。

摘要

CCCL的整个API在所有组件中共享一个通用的语义版本
仅支持最新发布的版本，不会将修复回溯到先前的版本
API破坏性变更和CCCL主版本号的增加仅会与CUDA工具包的新主版本发布同步
并非所有源代码破坏性变更都被视为公共API的破坏性变更，从而需要提升主版本号
不要依赖cub::或thrust::命名空间中实体的ABI稳定性
cuda::命名空间中符号的ABI破坏性变更可能随时发生，但会通过增加嵌入在所有cuda::符号的内联命名空间中的ABI版本来反映。可能同时支持多个ABI版本。

注意： 在将Thrust、CUB和libcudacxx合并到此存储库之前，每个库都按照语义版本控制独立进行版本控制。从2.1版本开始，这三个库在各自的存储库中同步了它们的发布版本。今后，CCCL将继续以单一的语义版本发布，2.2.0是nvidia/cccl存储库的第一个发布版本。

破坏性变更

破坏性变更是指已发布版本之间对明确支持的功能的更改，这些更改要求用户为升级到新版本而进行工作。

极端情况下，任何变更都有可能在某处破坏某些内容。因此，并非所有可能的源代码破坏性变更都被视为公共API的破坏性变更，从而需要提升主语义版本。

以下各节描述了CCCL的API和ABI破坏性变更的详细信息。

应用程序编程接口（API）

CCCL的公共API是为提供库实用性而有意暴露的全部功能。

换句话说，CCCL的公共API不仅仅包括函数签名，还包括（但不限于）：

预期直接包含在用户代码中的头文件的位置和名称
预期在用户代码中直接使用的命名空间
位于头文件中并预期直接在用户代码中使用的函数、类和变量的声明和/或定义
预期直接在用户代码中使用的函数、类和变量的语义

此外，CCCL的公共API不包括以下任何内容：

任何以_或__为前缀的符号
任何名称包含detail的符号，包括detail::命名空间或宏
任何包含在detail/目录或其子目录中的头文件
公共API部分的任何头文件隐式包含的头文件

通常，我们的目标是避免破坏公共API中的任何内容。只有在为用户提供更好的性能、更易理解的API和/或更一致的API时，才会进行此类更改。

对公共API的任何破坏性变更都需要提升CCCL的主版本号。为了与CUDA次要版本兼容性保持一致，API破坏性变更和CCCL主版本号的提升只会与CUDA工具包的新主版本发布同时进行。

不属于公共API的任何内容可能随时更改，恕不另行通知。

API版本控制

CCCL所有组件的公共API共享一个统一的主版本.次版本.补丁版本语义版本。

仅支持最新发布的版本。通常，功能和错误修复不会回溯到先前发布的版本或分支。

查询版本的首选方法是使用下面描述的CCCL_[MAJOR/MINOR/PATCH_]VERSION。为了向后兼容，Thrust/CUB/libcudacxxx版本定义仍然可用，并且始终与CCCL_VERSION保持一致。请注意，Thrust/CUB使用MMMmmmpp方案，而CCCL和libcudacxx使用MMMmmmppp方案。

	CCCL	libcudacxx	Thrust	CUB
头文件	`<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`
补丁/子次版本	`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包括但不限于：

函数和类型的修饰名称
对象和类型的大小和对齐
对象二进制表示中字节的语义

ABI破坏性变更是指导致公共API中函数或类型的ABI发生变化的任何更改。例如，向结构体添加新的数据成员是一个ABI破坏性变更，因为它改变了类型的大小。

在CCCL中，关于ABI的保证如下：

thrust:: 和 cub:: 命名空间中的符号可能随时在没有警告的情况下破坏 ABI。
thrust:: 和 cub:: 符号的 ABI 包括用于编译的 CUDA 架构。因此，在以下情况下，thrust:: 或 cub:: 符号可能具有不同的 ABI：
- 使用不同的架构编译
- 作为 CUDA 源文件（-x cu）与 C++ 源文件（-x cpp）编译
cuda:: 命名空间中的符号也可能随时破坏 ABI。但是，cuda:: 符号嵌入了 ABI 版本号，每当发生 ABI 破坏时，该版本号就会递增。可以同时支持多个 ABI 版本，因此用户可以选择恢复到先前的 ABI 版本。更多信息请参见此处。

谁应该关心 ABI？

通常，CCCL 用户只需在构建或使用 API 直接或间接包含 CCCL 提供的类型的二进制工件（如共享库）时才需要担心 ABI 问题。

例如，假设 libA.so 是使用 CCCL 版本 X 构建的，其公共 API 包括如下函数：

void foo(cuda::std::optional<int>);

如果另一个库 libB.so 使用 CCCL 版本 Y 编译并使用 libA.so 中的 foo，那么如果版本 X 和 Y 之间发生了 ABI 破坏，这可能会失败。与 API 破坏性变更不同，ABI 破坏通常不需要代码更改，只需要重新编译所有内容以使用相同的 ABI 版本。

要了解更多关于 ABI 及其重要性的信息，请参阅《ABI 是什么，C++ 应该如何处理它？》。

兼容性指南

如上所述，并非所有可能的源代码破坏性变更都构成需要增加 CCCL API 主版本号的破坏性变更。

鼓励用户遵循以下指南，以最大程度地减少因意外依赖不属于公共 API 的 CCCL 部分而导致的中断风险：

不要向 thrust::、cub::、nv:: 或 cuda:: 命名空间添加任何声明，也不要专门化其中的任何模板，除非某个特定符号有例外说明，例如专门化 cuda::std::iterator_traits
- 理由：这会导致与同名的符号或专门化发生冲突。
不要获取 thrust::、cub::、cuda:: 或 nv:: 命名空间中任何 API 的地址。
- 理由：这会阻止为这些 API 添加重载。
不要前向声明 thrust::、cub::、cuda:: 或 nv:: 命名空间中的任何 API。
- 理由：这会阻止为这些 API 添加重载。
不要直接引用任何以 _、__ 开头或名称中包含 detail 的符号，包括 detail:: 命名空间或宏
- 理由：这些符号仅供内部使用，可能随时更改，恕不另行通知。
包含你所使用的内容。对于你使用的每个 CCCL 符号，直接 #include 声明该符号的头文件。换句话说，不要依赖于其他头文件隐式包含的头文件。
- 理由：内部包含可能随时更改。