THE CUDA C++ DEVELOPER'S TOOLBOX

Bryce Adelstein Lelbach, NVIDIA

目录

• 引言：CPU 与 GPU 架构差异

• 在 C++ 中进行 GPU 编程

  • CUDA C++ 概述
  • NVIDIA 加速库生态系统

• 从 C++ 到 CUDA C++ 的演进

  • 编程范例：热量传递模拟
  • 代码演进与编译流程
  • 执行空间说明符與主机-设备执行模型

• 使用 Thrust 库实现并行化

  • Thrust：C++并行算法库
  • 使用Thrust迭代器实现高级并行模式
  • Thrust 算法优化示例：计算最大差值

• libcu++: CUDA C++ 基础库

• 高级应用：热方程求解器

  • 模板计算 (Stencil Computation
  • 使用 Thrust 实现与 mdspan 重构

• 异步执行与性能优化

  • 同步与异步执行模型
  • CUDA 流 (Stream) 与事件
  • Thrust 异步接口与阻塞行为

• 底层库 CUB 简介

  • CUB - CUDA C++ 算法创作工具包
  • Thrust 与 CUB API 对比：Reduce 操作
  • 使用 CUB 优化模板计算

• 内存管理

  • thrust::universal_vector 的内存模型
  • Thrust 容器对比与数据传输
  • 优化后的执行流程

• 总结与生态系统

  • CUDA 软件生态系统
  • GTC 2025 CUDA 开发者会议

引言：CPU 与 GPU 架构差异

讲座首先对比了中央处理器（CPU）和图形处理器（GPU）的核心架构及其内存系统特性，以阐明两者在设计哲学和适用场景上的根本不同。

• 核心架构：CPU 由少量为串行处理优化的强大核心构成，而 GPU 则集成了大量为并行处理设计的核心。

• 内存系统性能：
  • CPU：其内存系统具有较低的访问延迟（约100纳秒）和适中的带宽（约100 GB/s），适合执行需要快速响应和复杂逻辑控制的延迟敏感型任务。
  • GPU：其内存系统虽然访问延迟较高（约500纳秒），但提供了极高的内存带宽（约1000 GB/s）。这一特性使其非常适合处理数据密集型、可大规模并行化的计算任务，通过高吞吐量来弥补延迟的不足。

在 C++ 中进行 GPU 编程

为了利用 GPU 的强大并行计算能力，开发者可以在 C++ 环境中通过以下两种主要方式进行编程：

1. CUDA C++
2. 调用加速库

CUDA C++ 概述

CUDA C++ 是标准 C++（Standard C++）的一个扩展，旨在支持编写能够同时在 CPU（主机，Host）和 GPU（设备，Device）上异构执行的程序。

• 主机代码 (Host Code)：在 CPU 上运行，支持所有标准 C++ 的特性。
• 设备代码 (Device Code)：在 GPU 上运行，支持几乎所有标准 C++ 的特性。

NVIDIA 加速库生态系统

除了直接使用 CUDA C++ 编写底层代码，NVIDIA 还提供了一个丰富且功能强大的加速库生态系统。这些库针对特定领域进行了深度优化，使开发者能够轻松地在其应用程序中集成 GPU 加速功能。该生态系统包括但不限于：

• cuBLAS: 基础线性代数子程序库。
• cuDNN: 深度神经网络库。
• cuFFT: 快速傅里叶变换库。
• cuSPARSE: 稀疏矩阵运算库。
• RAPIDS: 面向数据科学和机器学习的库套件。
• Thrust: 基于 C++ 标准库的并行算法库。
• NCCL: NVIDIA 集体通信库，用于多 GPU/多节点通信。
• CUTLASS: 用于实现高性能矩阵乘法的 CUDA C++ 模板库。

从 C++ 到 CUDA C++ 的演进

编程范例：热量传递模拟

为了具体说明编程概念，讲座以一个简单的热量传递模拟为例。

1. 问题设定

模拟三个不同初始温度的杯子在特定环境温度下的冷却过程。
- 初始温度分别为：42°C, 24°C, 50°C。

2. 物理模型

温度的更新遵循牛顿冷却定律的简化模型。下一时刻的温度由当前温度、环境温度和传热系数共同决定。

• 更新公式: 下一时刻温度 = 当前温度 + 传热系数 * (环境温度 - 当前温度)
• 参数:
  • 环境温度 (Ambient Temperature) = 20°C
  • 传热系数 (Heat Transfer Coefficient, k) = 0.5

根据此模型，经过一个时间步长后，三个杯子的温度将分别从 42°C, 24°C, 50°C 更新为 31°C, 22°C, 35°C。

3. 标准 C++ 实现

该模拟过程可以使用标准 C++ 代码实现。核心是定义一个更新操作，并在一个循环中对每个杯子的温度应用该操作。

#include <vector>
#include <numeric>   // For std::views::iota
#include <algorithm> // For std::transform
#include <print>     // For std::println (C++23)

int main() {
    int steps = 3;
    float k = 0.5f;
    float ambient_temp = 20.0f;
    std::vector<float> cups{42.0f, 24.0f, 50.0f};

    // 定义单步温度更新操作 (lambda 函数)
    auto op = [=](float t) {
      float diff = ambient_temp - t;
      return t + k * diff;
    };

    // 模拟多个时间步长
    for (int step : std::views::iota(0, steps)) {
      std::println("Step {}: {}", step, cups);
      // 对容器中的每个元素应用更新操作
      std::transform(cups.begin(), cups.end(),
                     cups.begin(), op);
    }
    std::println("Final: {}", cups);
}

代码中的 std::transform 算法是关键，它遍历输入范围（cups 容器），对每个元素应用 op 函数，并将结果写回目标位置。这是一个典型的可以被并行化的数据并行操作。

代码演进与编译流程

1. 标准 C++ 编译与执行

上述标准 C++ 代码通过 GCC (g++) 等传统 C++ 编译器进行编译。编译器将高级的 C++ 表达式（如 t + k * diff）翻译成 CPU 可以直接执行的底层机器指令（例如 ARM 架构下的 vmla.f32 指令）。最终生成一个在主机 CPU 上运行的可执行文件。

2. CUDA C++ 编译流程

在CUDA C++中，源代码通过NVIDIA C++编译器（NVCC）进行编译。NVCC能够区分主机（Host）代码和设备（Device）代码。它将这两种代码分离开来，分别编译成适用于CPU和GPU的可执行部分。

• 主机代码 (Host code): 由CPU执行。
• 设备代码 (Device code): 由GPU执行。

编译过程如下图所示，nvcc main.cpp -o a.out 命令会将一份CUDA C++源代码（例如 main.cpp）编译成两个部分：一部分是CPU可执行的指令（如vmla.f32），另一部分是GPU可执行的指令（如fma.rm.f32）。

执行空间说明符與主机-设备执行模型

执行空间说明符 (Execution Space Specifiers)

为了让GPU能够执行计算，我们需要告诉编译器哪些代码可以在设备上运行。这通过在函数或lambda表达式前添加执行空间说明符来实现。

• void a(); 或 __host__ void b();

  • 默认情况下，或显式使用 __host__ 关键字声明的函数，仅可由CPU执行。这包括所有未被annotated的函数。

• __device__ void c();

  • 使用 __device__ 声明的函数，仅可由GPU执行。

• __host__ __device__ void d();

  • 同时使用 __host__ 和 __device__ 声明的函数，是一个通用函数，既可由CPU执行，也可由GPU执行。编译器会为其生成两个版本的代码。

在热量模拟的例子中，我们为lambda表达式 op 添加了 __host__ __device__ 说明符，表示它可以同时被主机和设备编译和调用。

auto op = [=] __host__ __device__ (float t) {
    float diff = ambient_temp - t;
    return t + k * diff;
};

主机-设备执行模型

CUDA的编程模型明确区分了主机（CPU）和设备（GPU）的执行流程。

1. 执行始于主机：程序的 main 函数在CPU上启动。
   

2. 显式启动设备任务：主机代码通过特定的调用（例如内核启动语法 <<<...>>> 或调用并行库函数）来显式地在GPU上启动计算任务。
   

3. 主机与设备间的转换是显式的：代码执行从CPU到GPU的切换必须由程序员明确指定。
   

4. 设备上调用的函数停留在设备上：一旦执行流程转移到GPU，在设备上调用的 __device__ 函数会继续在GPU上执行，直到该设备任务完成。
   

使用 Thrust 库实现并行化

标准库算法（如std::transform）通常在主机上串行执行。为了在GPU上并行执行这些操作，我们可以使用CUDA生态系统中的并行算法库，例如 Thrust。

只需将 std::transform 替换为 thrust::transform，并提供一个执行策略（如thrust::cuda::par，代表在CUDA设备上并行执行），即可将计算任务调度到GPU上。

同时，std::vector 的内存分配在主机端，GPU通常无法直接访问。为了解决数据可访问性问题，Thrust 提供了 thrust::universal_vector。它使用统一内存（Unified Memory），使得数据对于主机和设备都是可见和可访问的。

// 完整的 CUDA C++ 并行化代码
int steps = 3;
float k = 0.5;
float ambient_temp = 20;
thrust::universal_vector<float> cups{42, 24, 50}; // 修改：使用 universal_vector

auto op = [=] __host__ __device__ (float t) { // 修改：添加执行空间说明符
    float diff = ambient_temp - t;
    return t + k * diff;
};

for (int step : std::views::iota(0, steps))
{
    std::print("{} {}
", step, cups);
    thrust::transform(thrust::cuda::par, cups.begin(), cups.end(), // 修改：使用 thrust::transform
                      cups.begin(), op);
}

Thrust：C++并行算法库

Thrust是一个基于C++标准库（STL）的CUDA C++并行算法库。它提供了丰富的高性能并行算法接口，使开发者可以轻松编写GPU加速代码。

• 官方网站: https://nvidia.github.io/cccl/thrust

Thrust库主要包含以下几个核心组件：

1. 算法（Algorithms）

Thrust提供了一系列与C++标准库兼容的算法，以及一些为并行计算设计的扩展算法。

• 标准算法（Standard Algorithms）：这些算法是并行版本的C++标准库算法，例如：thrust::transform_reduce, thrust::inclusive_scan, thrust::sort, thrust::copy。
• 扩展算法（Extended Algorithms）：这些是为并行编程模式设计的附加算法，例如：thrust::reduce_by_key, thrust::sort_by_key, thrust::tabulate, thrust::gather。

2. 容器（Containers）

Thrust提供了类似于STL容器的并行数据结构，用于管理主机（CPU）和设备（GPU）内存。

• thrust::device_vector：在设备内存中分配和管理的动态数组。
• thrust::host_vector：在主机内存中分配和管理的动态数组，其内存是“可分页锁定（page-locked）”的，可以实现与设备之间的高效数据传输。
• thrust::universal_vector：在统一内存（Universal Memory）中分配的动态数组，可以从主机和设备代码中直接访问。
• thrust::allocate_unique：用于分配和管理唯一所有权内存的智能指针。

thrust::universal_vector与std::vector相比具有显著优势。std::vector的内容通常只能在主机代码中访问，并且其构造和赋值操作是串行执行的。相比之下，thrust::universal_vector的内容可以在主机和设备代码中无缝访问，并且其构造和赋值操作是并行执行的，从而提高了性能。

3. 迭代器（Iterators）

迭代器是Thrust库的核心抽象，它将算法与容器解耦。除了常规的指针式迭代器，Thrust还提供了一系列特殊的迭代器适配器，用于创建复杂的数据序列而无需显式地在内存中存储它们。

• thrust::counting_iterator：生成一个递增的整数序列。
• thrust::transform_iterator：在访问序列元素时，动态地对另一个迭代器指向的元素应用一个函数。
• thrust::zip_iterator：将多个输入迭代器合并成一个元组（tuple）迭代器，使得算法可以同时处理多个数据流。

使用Thrust迭代器实现高级并行模式

Thrust的迭代器适配器是实现高效、可组合的并行代码的关键。

虚拟数据序列

通过使用特殊的迭代器，可以在不占用额外内存的情况下生成输入序列。

• thrust::make_counting_iterator(0) 会生成一个从0开始的整数序列 0, 1, 2, ...，常用于并行循环中获取元素的索引。
  

• thrust::make_constant_iterator(42) 会生成一个所有元素都为 42 的序列 42, 42, 42, ...，可用于初始化或作为算法的常量输入。
  

算法融合（Algorithm Fusion）

一个常见的并行模式是transform后紧跟一个reduce操作。传统实现方式需要一个临时向量来存储transform的中间结果，这会消耗宝贵的设备内存，并且在两个内核调用之间引入了不必要的同步点和内存流量。

通过使用thrust::transform_iterator，可以将transform和reduce操作融合成一个单一的内核调用。transform_iterator在reduce算法访问数据时，即时（on-the-fly）地应用转换函数，从而避免了临时内存的分配和中间数据的读写。

// 低效的实现方式：需要临时存储 tmp
thrust::universal_vector X(N), tmp(N);
thrust::transform(thrust::cuda::par, X.begin(), X.end(), tmp.begin(), f);
auto r = thrust::reduce(thrust::cuda::par, tmp.begin(), tmp.end(), T{}, g);

// 高效的融合实现：使用 transform_iterator 避免临时存储
thrust::universal_vector X(N);
auto tmp_iterator = thrust::make_transform_iterator(X.begin(), f);
auto r = thrust::reduce(thrust::cuda::par, tmp_iterator, tmp_iterator + N, T{}, g);

这个版本的代码将两个操作逻辑上融合在一起，提高了执行效率和内存利用率。

处理多数据流

thrust::zip_iterator能够将多个数据序列“压缩”在一起，使算法可以对每个位置上来自不同序列的元素进行操作。例如，可以将向量X和Y的元素配对成(X[i], Y[i])的元组流。

zip_iterator也非常适用于实现模板（stencil）操作，例如计算相邻元素的差分。通过将一个迭代器与其自身偏移一个位置后的迭代器进行压缩，可以方便地访问 X[i] 和 X[i+1]。

Thrust 算法优化示例：计算最大差值

本节通过一个计算两个向量A和B对应元素之差的最大值的例子，展示了如何使用 Thrust 库逐步优化代码，以减少内存占用和提高性能。

初始状态
首先，我们有两个 thrust::universal_vector<int> 类型的向量 A 和 B。

步骤 1: transform + reduce (使用临时存储)
第一种实现方法是分两步进行，需要分配一个额外的向量 diffs 来存储中间结果，这会增加内存开销。

步骤 2: 使用迭代器适配器 (Iterator Adaptors)
为了避免创建临时向量，我们可以使用 thrust::make_zip_iterator 和 thrust::make_transform_iterator。这种方法通过"迭代器融合"（iterator fusion）避免了中间存储，从而减少了内存占用并可能提高性能。

步骤 3: 使用 transform_reduce (算法融合)
最高效的方法是使用 thrust::transform_reduce 算法。这个算法将转换（transformation）和归约（reduction）两个操作合并成一个单一的内核调用。这种"算法融合"（algorithm fusion）是最高效的方式，代码更简洁，并且性能通常是最好的。

libcu++: CUDA C++ 基础库

libcu++ 是 CUDA C++ 的基础库，提供了可在主机（Host）和设备（Device）代码中使用的标准 C++ 功能。

• 官方文档: https://nvidia.github.io/cccl/libcudacxx

libcu++ 扩展了标准的 C++ 库，使其能够在异构计算环境中使用。其命名空间和头文件结构如下：

1. 主机编译器标准库 (Host Compiler's Standard Library)

   • 头文件: #include <...>
   • 命名空间: std::
   • 描述: 标准 C++ 库，仅限于在 __host__ 代码中使用。

2. libcu++ - 标准 C++ 子集

   • 头文件: #include <cuda/std/...>
   • 命名空间: cuda::std::
   • 描述: 提供了标准 C++ 的一个子集，可在 __host__ 和 __device__ 代码中安全使用。

3. libcu++ - CUDA C++ 扩展

   • 头文件: #include <cuda/...>
   • 命名空间: cuda::
   • 描述: 提供了可在 __host__ 和/或 __device__ 代码中使用的现代 C++ API。

4. libcu++ - 实验性功能

   • 头文件: #include <cuda/experimental/...>
   • 命名空间: cuda::experimental:: (或 cudax::)
   • 描述: 包含尚在开发中的 Beta 功能。

高级应用：热方程求解器

模板计算 (Stencil Computation)

该图展示了一个典型的模板（Stencil）计算模式。模板计算是科学计算和图像处理中常见的并行计算模式。新网格中的每个点都依赖于旧网格中的一个局部区域。这种数据依赖模式非常适合在 GPU 上进行并行化。

使用 Thrust 实现与 mdspan 重构

本节通过一个热方程求解器的例子，展示了如何使用 Thrust 库编写并行计算代码，并引入了 C++ 标准库中的 mdspan 来优化和简化多维数据处理。

1. 原始实现：手动索引计算

最初的实现使用 thrust::for_each_n 并行处理每个网格点。在 Lambda 函数内部，需要手动将一维索引 xy 转换回二维坐标 x 和 y，并使用 x*ny + y 这样的方式来访问数据。

2. mdspan 简介与代码重构

mdspan 是一个非拥有的（non-owning）多维数据句柄，它提供了一个多维数组的视图，但本身不管理内存。使用 cuda::std::mdspan 重构后的代码，有以下优点：

• 函数签名更清晰：直接使用 mdspan 类型，明确表示参数是二维视图。
• 索引更直观：使用 U(x, y) 替代了 U[x*ny + y]，代码更易读。
• 维度获取更安全：使用 U.extent(0) 和 U.extent(1) 获取维度大小。

3. mdspan 的高级用法：submdspan

mdspan 还支持创建子视图（subspan），这对于处理边界条件等场景非常有用。例如，在初始化函数 initialize_oven 中，可以使用 cuda::std::submdspan 从原始 mdspan 中创建代表顶部、中部和底部区域的视图，而无需复制数据。

void initialize_oven(auto policy, cuda::std::mdspan<float, cuda::std::dims<2>> U) {
    auto nx = U.extent(0);
    auto top = cuda::std::submdspan(U, 0, cuda::std::full_extent);
    auto mid = cuda::std::submdspan(U, std::tuple{1, nx-2}, cuda::std::full_extent);
    auto bot = cuda::std::submdspan(U, nx-1, cuda::std::full_extent);

    thrust::fill_n(policy, top.data(), top.size(), 90.0);
    thrust::fill_n(policy, mid.data(), mid.size(), 15.0);
    thrust::fill_n(policy, bot.data(), bot.size(), 90.0);
}

这种方法使得对多维数据特定区域的操作变得非常简洁和高效。

异步执行与性能优化

同步与异步执行模型

• 同步执行 (thrust::cuda::par): 每次启动一个CUDA核函数，主处理器（Host）都会被阻塞，直到该核函数执行完毕。这种模式简化了编程逻辑，但限制了CPU与GPU的并行能力。

• 异步执行 (thrust::cuda::par_nosync): Host启动核函数后会立即返回，不会等待其完成。这允许CPU和GPU并行工作。然而，当需要从Host访问GPU数据时，必须使用 cudaDeviceSynchronize() 或更细粒度的同步机制来确保相关的GPU计算已经完成。

CUDA 流 (Stream) 与事件

为了提升性能并实现更精细的控制，可以采用 CUDA 流（Stream）。一个流是设备上按顺序执行的一系列操作。

• 流的创建与使用: 可以通过 cudax::stream 对象创建流，并使用 policy.on(stream) 将 Thrust 算法提交到特定流。不同流中的任务可以并发执行。
• 流内同步: 调用流对象的 wait() 方法会使Host线程阻塞，直到该流中此前提交的所有任务全部完成。
• 流间同步: 使用事件（Event）可以在不同的 CUDA 流之间创建依赖关系，甚至可以跨多个 GPU。一个流可以记录事件，而另一个流可以等待该事件发生后再继续执行。

通过将全局同步 cudaDeviceSynchronize() 替换为针对特定流的 stream.wait()，可以实现更高效的异步执行，避免不必要的全局阻塞。

// 同步版本
for (auto write_step : std::views::iota(0, write_steps)) {
    cudaDeviceSynchronize(); // 全局同步点
    save_to_file(U);
    // ...
}

// 异步版本
cudax::stream stream;
auto policy(thrust::cuda::par_nosync.on(stream.get()));
// ...
for (auto write_step : std::views::iota(0, write_steps)) {
    stream.wait(); // 仅等待此流完成
    save_to_file(U);
    // ...
}

Thrust 异步接口与阻塞行为

在使用 par_nosync 异步策略时，并非所有 Thrust 算法都是非阻塞的。
- 非阻塞: 像 for_each_n 或 transform 这样不返回值或返回可预先计算值的算法，是真正的异步操作。
- 阻塞: 某些 Thrust 算法会阻塞主机线程，因为主机需要等待GPU的计算结果：
- 返回一个依赖于计算结果的值（如 reduce、copy_if、find_if）。
- 需要为计算分配临时存储空间（如 inclusive_scan）。

底层库 CUB 简介

CUB - CUDA C++ 算法创作工具包

CUB 是一个用于创作高性能 CUDA C++ 算法的工具包，为编写 CUDA C++ 内核提供了可组合的构建块。

• 官方网站: https://nvidia.github.io/cccl/cub
• CUB 提供了四个层次的并行原语：设备级、块级、Warp级和线程级。

Thrust 与 CUB API 对比：Reduce 操作

Thrust 提供更高级、更简洁的接口，而 CUB 提供更底层、更灵活的接口。以 reduce 为例，Thrust 一行代码即可完成，但它是一个阻塞操作。

使用 CUB 实现相同功能需要更明确的两步：
1. 查询临时存储大小：首次调用 cub::DeviceReduce::Reduce，传入 nullptr 来获取执行所需的临时存储大小。
2. 分配内存并执行归约：分配所需内存后，再次调用 cub::DeviceReduce::Reduce 并传入内存指针以执行计算。

这种模式将资源分配与算法执行解耦，允许开发者对内存进行复用或其他优化，更好地支持异步操作。

使用 CUB 优化模板计算

对于模板计算，可以使用 cub::DeviceFor::ForEachInExtents。这是一个专为多维数据结构设计的 CUB API。它会自动处理从一维线程索引到多维坐标的映射，消除了在 Lambda 内部手动计算坐标的需要，使代码更简洁、可读，并可能带来性能优势。

内存管理

thrust::universal_vector 的内存模型

thrust::universal_vector 利用统一内存（Unified Memory），为异构计算提供了统一的内存视图。
- 逻辑视图: 用户看到的是一个普通的向量容器。
- 物理视图与自动迁移: 在底层，它在主机和设备内存中维护数据副本。当数据在一端被修改并在另一端被访问时，系统会自动、按需（通常以页为单位）地将数据迁移到访问端，这个过程对用户是透明的。

Thrust 容器对比与数据传输

• universal_vector: 可在主机和设备访问，数据传输是隐式的。
• host_vector: 仅可在主机访问。
• device_vector: 仅可在设备访问。
• host_vector 和 device_vector 之间的数据移动需要通过 thrust::copy 进行显式传输。

隐式传输虽然方便，但可能会导致意外的同步和性能瓶颈。为了精确控制数据传输，可以采用显式传输的策略：使用一个 thrust::host_vector 作为主机端缓冲区，并通过 thrust::copy_n 异步地将数据从 device_vector 或 universal_vector 显式复制到该缓冲区。

优化后的执行流程

通过使用显式数据复制，可以构建一个更高效的主机与设备之间的工作流程。设备端可以异步地将计算结果复制到主机端缓冲区，而主机则在确保复制完成后再进行文件保存等操作。这避免了在关键计算循环中由隐式内存访问引起的阻塞。

总结与生态系统

CUDA 软件生态系统

NVIDIA CUDA 提供了一个由众多库、API 和工具组成的庞大生态系统，以支持各种领域的加速计算，包括线性代数、深度学习、数据分析、信号处理等。

GTC 2025 CUDA 开发者会议

GTC 大会提供了大量与 CUDA C++、CUDA Python、性能优化和多 GPU 编程相关的开发者会议和议题。

相关链接:
- GitHub: http://github.com/NVIDIA/accelerated-computing-hub
- GTC 开发者页面: http://nvidia.com/gtc/sessions/cuda-developer