贡献者指南 (Contributor Guide)	开发容器 (Dev Containers)	Discord	Godbolt	GitHub 项目 (GitHub Project)	文档 (Documentation)

CUDA 核心计算库（CUDA Core Compute Libraries, CCCL）

欢迎来到 CUDA 核心计算库（CCCL），我们的使命是让 CUDA 开发更加愉悦。

本仓库将三个关键的 CUDA C++ 库统一到一个便捷的代码仓库中：

• Thrust（原仓库）
• CUB（原仓库）
• libcudacxx（原仓库）

CCCL 的目标是为 CUDA C++ 开发者提供构建模块，使其更容易编写安全且高效的代码。
将这些库整合在一起可简化您的开发流程，并增强您利用 CUDA C++ 强大能力的灵活性。
有关统一这些项目的决策详情，请参阅此处公告。

概述

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

• Thrust 是一个 C++ 并行算法库，它启发了 C++ 标准库引入并行算法。Thrust 的高层接口极大地提升了程序员的生产力，同时通过可配置的后端（支持多种并行编程框架，如 CUDA、TBB 和 OpenMP）实现了在 GPU 与多核 CPU 之间的性能可移植性。

• CUB 是一个底层的、专为 CUDA 设计的库，旨在为所有 GPU 架构提供极致性能的并行算法。除了设备级（device-wide）算法外，它还提供了协作式算法（cooperative algorithms），例如块级（block-wide）归约和线程束级（warp-wide）扫描，为 CUDA 内核开发者提供构建极致性能自定义内核所需的构建模块。

• libcudacxx 是 CUDA C++ 标准库。它提供了可在主机（host）和设备（device）代码中使用的 C++ 标准库实现。此外，它还提供了针对 CUDA 特定硬件功能的抽象，例如同步原语（synchronization primitives）、缓存控制（cache control）、原子操作（atomics）等。

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 <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() {

  // 分配并初始化输入数据
  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` 的归约和
  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];

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

  // 确保两种方法的结果一致
  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 中的所有内容均为头文件（header-only）形式。
因此，用户只需关注如何获取头文件以及如何将其集成到自己的构建系统中即可。

CUDA Toolkit

使用 CCCL 最简单的方式是通过 CUDA Toolkit，其中已包含 CCCL 的头文件。
当你使用 nvcc 编译时，它会自动将 CCCL 头文件路径添加到包含路径中，因此你可以直接在代码中 #include 任意 CCCL 头文件，无需额外配置。

如果使用其他编译器进行编译，则需要更新构建系统的包含路径（include search path），使其指向 CUDA Toolkit 安装目录中的 CCCL 头文件（例如 /usr/local/cuda/include）。

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

GitHub

希望使用最新版 CCCL 的用户建议直接从 GitHub 获取。
注意：使用较新版本的 CCCL 配合较旧版本的 CUDA Toolkit 是支持的，但反过来则不支持。
有关 CCCL 与 CUDA Toolkit 兼容性的完整信息，请参阅 平台支持。

由于 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，确保即使使用 -I 也能像使用 -isystem 一样屏蔽警告。更多信息请参见 https://github.com/NVIDIA/cccl/issues/527。

安装

默认的 CMake 选项仅生成安装规则，因此熟悉的 cmake . && make install 工作流可直接使用：

git clone https://github.com/NVIDIA/cccl.git
cd cccl
cmake . -DCMAKE_INSTALL_PREFIX=/usr/local
make install

我们还提供了一个便捷脚本：

ci/install_cccl.sh /usr/local

使用预设（presets）进行高级安装

CMake 预设（presets）也已提供，并支持包含实验性库的选项：

cmake --preset install -DCMAKE_INSTALL_PREFIX=/usr/local
cmake --build --preset install --target install

使用 install-unstable 预设可包含实验性库，或使用 install-unstable-only 预设仅安装实验性库。

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（例如 CUDA 12.4）配套发布的 CCCL 版本，可以运行以下命令安装：

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

cuda-cccl 元包（metapackage）会安装与 cuda-version 对应的 CUDA Toolkit 中配套发布的 cccl 版本。
如果你在安装 cuda-cccl 后希望升级到最新的 cccl，请先卸载 cuda-cccl，再安装 cccl：

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

注意
还存在名为 cuda-cccl_linux-64 等的 conda 包。
这些包包含作为 CUDA Toolkit 一部分发布的 CCCL 版本，但专为 CUDA Toolkit 内部使用而设计。
为了与 conda 编译器兼容，请安装 cccl 或 cuda-cccl。
更多信息请参阅 cccl conda-forge recipe。

CMake 集成

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

关于如何通过 CMake Package Manager 实现集成的完整示例，请参阅 我们的基础示例项目。

其他构建系统理论上也可工作，但只有 CMake 经过测试。
欢迎提交改进 CCCL 在其他构建系统中集成体验的贡献。

贡献者

有兴趣为 CCCL 的改进做出贡献？请查阅我们的 贡献指南，其中全面介绍了设置开发环境、修改代码、运行测试和提交 PR 所需了解的一切内容。

平台支持

目标：本节描述用户可以在哪些平台上预期 CCCL 能成功编译和运行。

通常情况下，CCCL 应能在 CUDA Toolkit 支持的所有平台上正常工作，但细节决定成败。
以下各小节详细说明了不同版本的 CUDA Toolkit、主机编译器（host compilers）和 C++ 方言（C++ dialects）的支持与测试情况。

CUDA Toolkit (CTK) 兼容性

摘要：

• 最新版本的 CCCL（CUDA C++ Core Libraries）向后兼容当前及上一个 CTK 主版本系列
• CCCL 从不向前兼容任何版本的 CTK。请始终使用与您的 CTK 所附带版本相同或更新的 CCCL 版本
• CCCL 的次版本升级不会破坏现有代码，但新功能可能不支持所有 CTK 版本

我们鼓励 CCCL 用户充分利用最新改进，并通过始终使用最新版 CCCL 来 "live at head"（即始终使用主干最新代码）。 为了获得无缝体验，您可以独立于整个 CUDA Toolkit 升级 CCCL。 这是可行的，因为 CCCL 对当前和上一个主版本系列中每个 CTK 次版本的最新补丁版本都保持向后兼容。 在某些特殊情况下，所支持的 CUDA Toolkit 发布版本的最低次版本可能需要比其主版本系列中最旧的发布版本更新。

当新的 CTK 主版本发布时，我们将停止对最旧受支持主版本的支持。

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

使用最新 CCCL 编写的行为良好的代码应能在任何受支持的 CTK 版本上成功编译和运行。 对于依赖新 CTK 功能的新特性，可能会出现例外情况，因此这些特性在较旧版本的 CTK 上将无法工作。

用户可以将新版 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 相同的所有操作系统，具体文档如下：

• Linux
• Windows

主机编译器（Host Compilers）

除非另有说明，CCCL 支持与最新 CUDA Toolkit 相同的主机编译器，具体文档如下：

• Linux
• Windows

对于 Linux 上的 GCC，至少需要 7.x 版本。

使用较旧 CUDA Toolkit 时，我们也仅支持最新 CUDA Toolkit 的主机编译器， 但至少会支持任何受支持的旧版 CUDA Toolkit 中最新的主机编译器。

我们可能会保留对额外编译器的支持，并接受社区提供的合理修复补丁。 但我们不会投入大量时间对旧编译器的问题进行排查或修复。

本着“只支持经过测试的内容”的原则，有关我们确切测试内容的更多信息，请参阅我们的 CI 概览。

GPU 架构

CCCL 支持所有 当前 主版本 CUDA Toolkit (CTK) 所支持的 GPU 架构。

明确说明：尽管 CCCL 可以使用当前和上一个 CTK 主版本进行编译，但我们不会测试或验证仅在旧版 CTK 中支持的架构。

这些架构可能仍然可以工作——我们不会故意破坏它们——但它们不在我们的常规 CI 覆盖范围内。此外，新功能也无法保证在这些架构上正常工作。

我们欢迎社区贡献合理的修复方案，以帮助这些旧架构上的用户解决问题。

例如，CCCL 3.0 支持使用 CTK 12.x 和 13.x 进行编译，其中：

• CUDA Toolkit 13.x 支持 >=sm_75
• CUDA Toolkit 12.x 支持 >=sm_50

在此场景下，使用 CTK 12.x 编译 CCCL 3.0 并针对低于 sm_75 的架构可能是可行的，但这些配置不属于我们的常规测试范围。

C++ 方言（Dialects）

• C++17
• C++20

测试策略

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

对于 CUDA Toolkit 版本，测试会同时针对最旧和最新的受支持版本进行。 例如，如果 CUDA Toolkit 的最新版本是 12.6，则测试会针对 11.1 和 12.6 进行。 对于每个 CUDA 版本，会使用所有受支持的主机编译器和所有受支持的 C++ 方言完成构建。

测试策略和矩阵在不断演进。 ci/matrix.yaml 文件中定义的矩阵是权威信息来源。 有关我们 CI 流水线的更多信息，请参阅 此处。

版本管理（Versioning）

目标： 本节描述了 CCCL 的版本管理方式、API/ABI 稳定性保证以及兼容性指南，以尽量减少升级时的麻烦。

摘要

• CCCL 的整个 API 在所有组件中共享一个统一的语义化版本（semantic version）
• 仅最新发布的版本受支持，修复不会向后移植到之前的版本
• API 的破坏性变更（breaking changes）和 CCCL 主版本号（major version）的递增，仅会在 CUDA Toolkit 发布新主版本时同步进行
• 并非所有源代码层面的破坏性变更都被视为公共 API 的破坏性变更，从而触发主版本号的递增
• 不要依赖 cub:: 或 thrust:: 命名空间中实体的 ABI（Application Binary Interface，应用程序二进制接口）稳定性
• cuda:: 命名空间中的符号可能发生 ABI 破坏性变更，但会通过递增嵌入在 cuda:: 符号内联命名空间（inline namespace）中的 ABI 版本来体现。多个 ABI 版本可能被同时支持。

注意： 在将 Thrust、CUB 和 libcudacxx 合并到本仓库之前，每个库都根据语义化版本规范独立进行版本管理。
从 2.1 版本开始，这三个库在其各自的仓库中同步了发布版本号。
今后，CCCL 将继续使用单一的 语义化版本 进行发布，其中 2.2.0 是首个来自 nvidia/cccl 仓库的版本。

破坏性变更（Breaking Change）

破坏性变更是指在已发布版本之间，对明确支持的功能所做的更改，该更改要求用户必须进行额外工作才能升级到新版本。

从极限角度看，_任何_更改 都有可能在某处破坏某些用户的使用。
因此，并非所有可能导致源代码不兼容的更改都被视为公共 API 的破坏性变更，从而触发主版本号的递增。

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

应用程序编程接口（Application Programming Interface, API）

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

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

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

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

• 任何以下划线 _ 或双下划线 __ 开头的符号
• 名称中包含 detail 的任何符号，包括 detail:: 命名空间或宏
• 位于 detail/ 目录或其子目录中的任何头文件
• 公共 API 中任何头文件隐式包含的头文件

总体而言，我们的目标是避免对公共 API 做出破坏性更改。
只有当此类更改能为用户提供更好的性能、更易理解的 API 和/或更一致的 API 时，才会进行。

任何对公共 API 的破坏性变更都将导致 CCCL 主版本号的递增。
根据 CUDA 次版本兼容性（CUDA Minor Version Compatibility） 的原则，
API 破坏性变更和 CCCL 主版本号的递增，仅会在 CUDA Toolkit 发布新主版本时同步发生。

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

API 版本管理

CCCL 所有组件的公共 API 共享统一的语义化版本号 MAJOR.MINOR.PATCH。

仅最新发布的版本受支持。
原则上，新功能和错误修复不会向后移植到先前发布的版本或分支中。

推荐的版本查询方式是使用如下所述的 CCCL_[MAJOR/MINOR/PATCH_]VERSION 宏。
为了向后兼容，Thrust/CUB/libcudacxx 的版本定义仍然可用，并且始终与 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, Application Binary Interface）

应用程序二进制接口（ABI）是一组规则，用于规定：

• 库的组件在机器码中如何表示
• 这些组件如何在不同的翻译单元（translation units）之间交互

一个库的 ABI 包括但不限于以下内容：

• 函数和类型的名称修饰（mangled names）
• 对象和类型的大小与对齐方式（size and alignment）
• 对象二进制表示中字节的语义（semantics）

ABI 破坏性变更（ABI Breaking Change） 是指任何导致公共 API 中函数或类型的 ABI 发生变化的修改。
例如，向结构体（struct）中添加一个新的数据成员就是一种 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 主版本号的“破坏性变更”。

为尽量减少因意外依赖 CCCL 非公共 API 部分而导致的问题，建议用户遵循以下准则：

• 除非针对特定符号另有说明（例如特化 cuda::std::iterator_traits），否则不要向 thrust::、cub::、nv:: 或 cuda:: 命名空间中添加任何声明，也不要特化其中的任何模板
  • 理由：如果将来添加了同名的符号或特化，会导致冲突。
• 不要获取 thrust::、cub::、cuda:: 或 nv:: 命名空间中任何 API 的地址
  • 理由：这会阻碍为这些 API 添加重载。
• 不要对 thrust::、cub::、cuda:: 或 nv:: 命名空间中的任何 API 进行前向声明
  • 理由：这会阻碍为这些 API 添加重载。
• 不要直接引用任何以下划线 _、双下划线 __ 开头，或名称中包含 detail（包括 detail:: 命名空间或宏）的符号
  • 理由：这些符号仅供内部使用，可能随时无预警地更改。
• 使用什么就包含什么。对于你使用的每个 CCCL 符号，请直接 #include 声明该符号的头文件。换句话说，不要依赖其他头文件隐式包含的头文件。
  • 理由：内部包含关系可能随时更改。

本部分内容部分参考了 Abseil 的兼容性指南。

弃用策略

我们将尽最大努力在对公共 API、ABI 或支持的平台和编译器进行任何破坏性变更之前通知用户。

在适当的情况下，弃用将以可禁用的程序化警告形式出现。

弃用周期将根据变更的影响程度而定，但通常至少持续 2 个次要版本（minor version）发布。

与 CTK 版本的对应关系

CCCL 版本	CTK 版本
3.2	13.2
3.1	13.1
3.0	13.0
2.8	12.9
2.7	12.8
2.5	12.6
2.4	12.5
2.3	12.4

自我测试：https://cuda.godbolt.org/z/K818M4Y9f

12.4 之前的 CTK 将 Thrust、CUB 和 libcudacxx 作为独立库发布。

Thrust/CUB/libcudacxx 版本	CTK 版本
2.2	12.3
2.1	12.2
2.0/2.0/1.9	12.1
2.0/2.0/1.9	12.0

CI 流水线概览

有关 CI 流水线的详细概述，请参阅 ci-overview.md。

相关项目

以下项目与 CCCL 使命（让 CUDA 更加愉悦）相关：

• cuCollections - GPU 加速的数据结构，如哈希表
• NVBench - 专为 CUDA 应用定制的基准测试库
• stdexec - Senders 异步编程模型的参考实现

使用 CCCL 的项目

您的项目是否使用了 CCCL？欢迎提交 PR 将您的项目添加到此列表中！

• AmgX - 多重网格线性求解器库
• ColossalAI - 用于编写分布式深度学习模型的工具
• cuDF - 用于 ETL 数据分析的算法和文件读取器
• cuGraph - 图分析算法
• cuML - 机器学习算法和基础组件（primitives）
• cuOpt - 加速决策优化
• CuPy - GPU 上的 NumPy 与 SciPy
• cuSOLVER - 稠密与稀疏线性求解器
• CUSP - 稀疏矩阵运算、迭代方法和代数多重网格
• cuVS - 近似聚类与向量搜索
• GooFit - 最大似然拟合库
• HeavyDB - SQL 数据库引擎
• HOOMD - 蒙特卡洛与分子动力学模拟
• HugeCTR - GPU 加速的推荐系统框架
• Hydra - 高能物理数据分析
• Hypre - 多重网格线性求解器
• LightSeq - 序列处理与生成的训练和推理
• MatX - 使用表达式模板（expression templates）提供高效、类似 Python 语法的数值计算库
• PyTorch - 张量与神经网络计算
• Qiskit - 量子电路高性能模拟器
• QUDA - 格点量子色动力学（QCD）计算
• RAFT - 机器学习算法和基础组件（primitives）
• SGLang - 大语言模型（LLM）服务框架
• TensorFlow - 端到端机器学习平台
• TensorRT - 深度学习推理
• TensorRT-LLM - 优化的大语言模型（LLM）推理
• tsne-cuda - 随机邻域嵌入（Stochastic Neighborhood Embedding）库
• 可视化工具包（VTK, Visualization Toolkit） - 渲染与可视化库
• vLLM - 大语言模型（LLM）推理与服务
• XGBoost - 梯度提升机器学习算法

CCCL 快速上手指南

CUDA Core Compute Libraries（CCCL）是 NVIDIA 提供的统一 CUDA C++ 基础库，整合了 Thrust、CUB 和 libcudacxx 三大核心组件，帮助开发者高效编写高性能 CUDA 代码。

---

环境准备

• 操作系统：Linux（推荐 Ubuntu 20.04+）、Windows 或 macOS（仅支持部分功能）
• GPU：支持 CUDA 的 NVIDIA GPU（计算能力 ≥ 3.5）
• 依赖项：
  • CUDA Toolkit 11.8 或更高版本（推荐 12.x）
  • 支持 C++17 的主机编译器（如 GCC ≥ 9、Clang ≥ 12）
  • （可选）CMake ≥ 3.20（用于高级集成）

💡 国内用户建议通过 清华镜像站 安装 CUDA Toolkit 加速下载。

---

安装步骤

方法一：使用 CUDA Toolkit（推荐）

CCCL 已内置在 CUDA Toolkit 中，无需额外安装。只需确保已正确安装 CUDA：

nvcc --version

编译时直接包含头文件即可（nvcc 自动配置路径）。

---

方法二：从 GitHub 获取最新版

适用于需要最新特性的开发者：

git clone https://github.com/NVIDIA/cccl.git

编译示例程序：

nvcc -Icccl/thrust -Icccl/libcudacxx/include -Icccl/cub main.cu -o main

⚠️ 注意：使用 -I 而非 -isystem，避免与 CUDA Toolkit 内置头文件冲突。

---

方法三：通过 Conda 安装（推荐国内用户）

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

若需匹配特定 CUDA 版本（如 12.4）：

conda install cuda-cccl cuda-version=12.4

---

方法四：CMake 安装（适用于项目集成）

git clone https://github.com/NVIDIA/cccl.git
cd cccl
cmake . -DCMAKE_INSTALL_PREFIX=/usr/local
make install

或使用预设：

cmake --preset install -DCMAKE_INSTALL_PREFIX=/usr/local
cmake --build --preset install --target install

---

基本使用

以下是一个结合 Thrust、CUB 和 libcudacxx 实现并行归约的完整示例：

#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() {
  int const N = 1000;
  thrust::device_vector<int> data(N);
  thrust::fill(data.begin(), data.end(), 1);

  thrust::device_vector<int> kernel_result(1);
  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)
  );

  cudaDeviceSynchronize();
  int const custom_result = kernel_result[0];
  int const thrust_result = thrust::reduce(thrust::device, data.begin(), data.end(), 0);

  std::printf("Custom kernel sum: %d\n", custom_result);
  std::printf("Thrust reduce sum: %d\n", thrust_result);
  assert(custom_result == thrust_result);
  return 0;
}

编译并运行：

nvcc -Icccl/thrust -Icccl/libcudacxx/include -Icccl/cub example.cu -o example
./example

✅ 输出应为：

Custom kernel sum: 1000
Thrust reduce sum: 1000