CCCL 运行时：面向 CUDA 的现代 C++ 运行时 | NVIDIA 技术博客
CCCL 运行时：面向 CUDA 的现代 C++ 运行时 | NVIDIA 技术博客
CCCL 运行时：为 CUDA 打造的现代 C++ 运行时
NVIDIA CUDA 核心计算库（CCCL） 为 CUDA 开发者提供了用 C++ 和 Python 编写的、高效且令人愉悦的抽象。它包含以下特性：
- 并行算法 – 主机启动的算法，包括排序、扫描和归约，免去了为常见操作编写自定义内核的需要。
- 协同算法 – 设备端算法，例如块级或线程束级的归约或扫描，简化了自定义内核的开发。
- 符合语言习惯的 CUDA 抽象 – 针对 CUDA 特定操作（包括内存分配、资源管理和硬件特性）的基础抽象。
本文介绍 CCCL 中的一组新功能，它为 CUDA 编程模型的基本概念提供了现代化的 C++ 抽象，使 CUDA C++ 开发更安全、更便捷。
什么是 CCCL 运行时？
NVIDIA CCCL 运行时是一组新的、符合习惯的 C++ API，从 CUDA 13.2 开始可用，实现了核心 CUDA 功能：流管理、内存分配、内核启动等。
大家熟悉的 NVIDIA CUDA 运行时最初是作为 CUDA 驱动 API 之上的便利层开发的。新的 CCCL 运行时旨在成为一个具有相同目标的替代方案，但采用与现代 C++ 对齐的更新设计。下图展示了上述三个 CUDA API 层之间的关系：
CCCL 运行时是 CCCL 内部的一组头文件，例如 <cuda/stream>、<cuda/memory> 和 <cuda/launch>。它利用现代 C++ 特性，提供了比传统 CUDA 运行时 API（受限于 C 源代码兼容性约束）更便捷、更健壮的抽象。
我们还借此机会将过去 20 年 CUDA 演变过程中积累的经验融入到了 API 设计中。即使有这些变化，CCCL 运行时也提供了兼容性辅助工具，让开发者可以逐步采用它，而无需重写使用 CUDA 运行时 API 的现有代码。
随着 CUDA 程序变得越来越复杂，多个库共享设备、流和内存，对能够干净组合并使依赖关系显式的 API 的需求变得日益迫切。这正是 CCCL 运行时旨在填补的领域。
代码
下面是用新的 CCCL 运行时 API 实现的经典 vectorAdd 示例。如果你之前写过 CUDA，整体结构会似曾相识：请关注不同之处。不要试图一下子理解所有内容，本文的其余部分将逐步讲解这个示例，解释 CCCL 运行时背后的语义和设计选择。
#include <cuda/stream>
#include <cuda/memory>
#include <cuda/launch>
#include <cuda/grid>
#include <cuda/std/span>
#include <cuda/thread>
struct kernel {
template <typename Config>
__device__ void operator()(Config config,
cuda::std::span<const int> A,
cuda::std::span<const int> B,
cuda::std::span<int> C) {
auto tid = cuda::gpu_thread.rank(cuda::grid, config);
if (tid < A.size())
C[tid] = A[tid] + B[tid];
}
};
int main() {
// 1. 设备和流
cuda::device_ref device = cuda::devices[0];
cuda::stream stream{device};
// 2. 内存分配
auto pool = cuda::device_default_memory_pool(device);
int num_elements = 1000;
auto A = cuda::make_buffer(stream, pool, num_elements, 1);
auto B = cuda::make_buffer(stream, pool, num_elements, 2);
auto C = cuda::make_buffer(stream, pool, num_elements, cuda::no_init);
// 3. 内核启动
constexpr int threads_per_block = 256;
auto config = cuda::distribute(num_elements);
cuda::launch(stream, config, kernel{}, A, B, C);
// 让 CPU 线程等待 GPU 工作完成。
stream.sync();
return 0;
}
该示例可分为以下三个主要部分：
1）设备和流
考虑使用 CUDA 运行时 API 创建流，如下代码片段所示：
cudaStream_t stream;
cudaStreamCreate(&stream); // 与碰巧是“当前”的设备关联
注意，它创建了一个流，但该流与调用 cudaStreamCreate 时“当前”的设备相关联。仅凭这一调用，你无法知道该流与哪个设备关联。
对比之下，使用 CCCL 运行时 API 的创建方式如下所示：
cuda::device_ref device = cuda::devices[0];
cuda::stream stream{device};
上述代码片段展示了如何在特定设备上创建流。第一行说明了一个核心设计原则：CCCL 运行时使用专用类型而非原始标识符。设备是 device_ref 而非普通整数；流是一个对象而非不透明指针。整个 API 的强类型有助于在编译时捕获错误，而不是在运行时排查。
第二行说明了另一个原则：使依赖关系显式化。在 CCCL 运行时和 CUDA 运行时 API 中，流都与设备关联。区别在于关联的方式。这里，cuda::stream 构造函数将设备作为显式参数，而 CUDA 运行时 API 中，流与创建时处于活动状态的设备关联。
显式依赖关系支持局部推理。阅读函数时，无需追踪全局状态就能理解其功能。它们还提高了可组合性：当多个库一起使用时，无需在调用之间保存和恢复隐式状态以避免相互干扰。
一个相关的后果是，CCCL 运行时不暴露默认流。管理默认流的含义需要追踪当前设备，这正是我们要摒弃的隐式状态。虽然 CUDA 运行时 API 的默认流仍然可以包装到 CCCL 运行时类型中，但不推荐使用；任何涉及默认流的操作都应直接通过 CUDA 运行时 API 处理。由于 API 中没有默认流，因此“阻塞流”的概念不再适用，所有 CCCL 运行时流都创建为非阻塞流。
资源所有权：拥有类型和引用类型
借鉴 std::string 和 std::string_view 的例子，CCCL 运行时中的许多 CUDA 对象都有两种类型：拥有类型和带有 _ref 后缀的非拥有类型；cuda::stream 拥有底层的 cudaStream_t 句柄，并在其析构函数中销毁它。cuda::stream_ref 持有句柄但不管理其生命周期，并且是平凡可复制的。
_ref 类型对于与现有代码的可组合性至关重要。如果流句柄的生命周期在其他地方管理，cudaStream_t 可以隐式转换为 cuda::stream_ref，并且可以通过 .get() 获取原始句柄。要转移所有权，cuda::stream::from_native_handle 将原始句柄包装到拥有类型中，而 .release() 则放弃所有权。
void stream_type_example(cudaStream_t handle) {
cuda::stream_ref non_owning{handle};
assert(handle == non_owning.get());
cuda::stream owning = cuda::stream::from_native_handle(handle);
assert(handle == owning.get());
assert(handle == owning.release());
}
同样的模式适用于事件、内存池和其他 CUDA 对象：cuda::device_ref 没有对应的拥有类型，因为没有需要拥有的设备状态。
2）内存分配
auto pool = cuda::device_default_memory_pool(device);
auto A = cuda::make_buffer(stream, pool, num_elements, 1);
auto B = cuda::make_buffer(stream, pool, num_elements, 2);
auto C = cuda::make_buffer(stream, pool, num_elements, cuda::no_init);
下一部分展示了异步