CUDA C++ 核心计算库 (CCCL): 加速 CUDA C++ 开发的利器

CCCL 简介

CUDA C++ 核心计算库 (CUDA C++ Core Libraries, CCCL) 是 NVIDIA 为 CUDA C++ 开发者提供的一套综合性工具库,其使命是让 CUDA C++ 编程变得更加愉悦和高效。CCCL 将三个重要的 CUDA C++ 库整合到一个统一的仓库中:

• Thrust: C++ 并行算法库
• CUB: CUDA 优化的并行原语库
• libcudacxx: CUDA C++ 标准库实现

通过将这些库集成到一起,CCCL 为 CUDA C++ 开发者提供了一套全面的工具,使他们能够更轻松地编写安全高效的代码。

CCCL 的核心组件

Thrust

Thrust 是一个受 C++ 标准模板库 (STL) 启发的高级 C++ 并行算法库。它为开发者提供了一组易用的并行算法接口,大大提高了编程效率。Thrust 支持多种后端,包括 CUDA、TBB 和 OpenMP,实现了 GPU 和多核 CPU 之间的性能可移植性。

主要特性:

• 提供类似 STL 的高级接口
• 支持 GPU 和 CPU 后端
• 包含丰富的并行算法,如排序、归约等

CUB

CUB 是一个低级的、面向 CUDA 的库,专为在所有 GPU 架构上实现极致性能而设计。除了设备级算法外,CUB 还提供了块级和线程束级的协作算法,如块级归约和线程束级扫描等,为 CUDA 内核开发者提供了构建高性能自定义内核的基础构件。

主要特性:

• 提供设备级、块级和线程束级算法
• 针对不同 GPU 架构优化
• 支持自定义数据类型和操作符

libcudacxx

libcudacxx 是 CUDA C++ 标准库的实现。它提供了可在主机和设备代码中使用的 C++ 标准库实现。此外,它还为 CUDA 特有的硬件特性提供了抽象,如同步原语、缓存控制和原子操作等。

主要特性:

• 提供 C++ 标准库在 CUDA 环境下的实现
• 支持主机和设备代码
• 提供 CUDA 特有功能的高级抽象

CCCL 的优势

1. 统一性: CCCL 将多个重要的 CUDA C++ 库整合到一个仓库中,简化了依赖管理和版本控制。

2. 性能: 通过整合优化的算法和原语,CCCL 能够帮助开发者编写高性能的 CUDA 代码。

3. 可移植性: 特别是 Thrust 库,支持多种后端,使得代码可以在 GPU 和 CPU 之间轻松移植。

4. 易用性: CCCL 提供了高级抽象和易用的接口,大大降低了 CUDA 编程的学习曲线。

5. 标准化: libcudacxx 提供了与标准 C++ 兼容的接口,使得开发者可以在 CUDA 环境中使用熟悉的 C++ 标准库功能。

使用 CCCL

安装

CCCL 的安装非常简单,有以下几种方式:

1. CUDA Toolkit: CCCL 已包含在 CUDA Toolkit 中。使用 nvcc 编译时,CCCL 的头文件会自动被添加到包含路径中。

2. GitHub: 对于想要使用最新版本的开发者,可以直接从 GitHub 克隆 CCCL 仓库:

   git clone https://github.com/NVIDIA/cccl.git
   nvcc -Icccl/thrust -Icccl/libcudacxx/include -Icccl/cub main.cu -o main

3. Conda: CCCL 也提供了 conda 包:

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

简单示例

以下是一个使用 CCCL 功能的简单示例,展示了如何使用 Thrust、CUB 和 libcudacxx 实现并行归约:

#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 << "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;
}

这个例子展示了如何使用 CUB 的块级归约、libcudacxx 的原子操作以及 Thrust 的设备向量和归约算法。它同时实现了自定义内核和 Thrust 算法两种方式的并行归约,并比较了结果。

CCCL 版本和兼容性

CCCL 采用语义化版本控制,格式为 MAJOR.MINOR.PATCH。CCCL 的版本与 CUDA Toolkit 版本有一定的对应关系:

• CCCL 2.x 支持 CUDA Toolkit 11.1 - 11.8 和 12.x (仅最新补丁版本)
• 未来的 CCCL 3.x 将支持 CUDA Toolkit 12.x 和 13.x (仅最新补丁版本)

CCCL 保持向后兼容性,但不保证向前兼容性。这意味着用户可以使用较新版本的 CCCL 配合较旧版本的 CUDA Toolkit,但反之则不行。

CCCL 的应用场景

CCCL 在众多领域和项目中得到了广泛应用,以下是一些使用 CCCL 的知名项目:

1. 机器学习框架:

   • PyTorch: 使用 CCCL 进行张量运算和神经网络计算
   • TensorFlow: 利用 CCCL 进行端到端机器学习

2. 数据分析:

   • cuDF: 使用 CCCL 进行 ETL 数据分析
   • cuML: 利用 CCCL 实现机器学习算法和原语

3. 科学计算:

   • CuPy: 基于 CCCL 实现 GPU 版本的 NumPy 和 SciPy
   • HOOMD-blue: 使用 CCCL 进行分子动力学模拟

4. 图形处理:

   • cuGraph: 利用 CCCL 实现图分析算法
   • Visualization Toolkit (VTK): 使用 CCCL 进行渲染和可视化

5. 高性能计算:

   • AmgX: 使用 CCCL 实现多重网格线性求解器
   • QUDA: 利用 CCCL 进行格点量子色动力学计算

这些应用充分展示了 CCCL 在各个领域的广泛应用,从机器学习到科学计算,再到图形处理和高性能计算,CCCL 都发挥着重要作用。

结语

CUDA C++ 核心计算库 (CCCL) 为 CUDA C++ 开发者提供了一套强大而全面的工具集。通过整合 Thrust、CUB 和 libcudacxx,CCCL 简化了 CUDA 编程,提高了开发效率,并帮助开发者充分发挥 GPU 的计算能力。无论是在机器学习、数据分析、科学计算还是高性能计算领域,CCCL 都是一个不可或缺的工具。

随着 GPU 计算在各个领域的不断普及,CCCL 的重要性也将日益凸显。对于希望在 CUDA 编程中提高效率和性能的开发者来说,深入学习和使用 CCCL 无疑是一个明智的选择。NVIDIA 也在持续改进和扩展 CCCL 的功能,相信在未来,CCCL 将为 CUDA C++ 开发带来更多可能性。