CCCL实战教程：从基础到高级的10个并行算法示例-优快云博客

CCCL实战教程：从基础到高级的10个并行算法示例

【免费下载链接】cccl CUDA Core Compute Libraries 项目地址: https://gitcode.com/gh_mirrors/cc/cccl

CCCL（CUDA Core Compute Libraries）是NVIDIA提供的一套强大的并行计算库集合，包含Thrust、CUB和libcu++等组件，帮助开发者轻松实现高效的GPU加速算法。本教程将通过10个实用示例，带你从基础到高级掌握CCCL的核心功能，解锁GPU并行计算的强大能力。

一、环境准备与项目搭建

1.1 快速安装CCCL

首先通过Git克隆项目仓库：

git clone https://gitcode.com/gh_mirrors/cc/cccl
cd cccl

1.2 配置CMake构建环境

CCCL使用CMake进行项目管理，提供了丰富的预设配置。在VS Code中打开项目后，通过CMake工具选择适合的配置预设：

选择目标构建类型后，CMake会自动生成对应的构建目标列表：

二、基础并行算法示例

2.1 向量加法：GPU并行计算入门

向量加法是并行计算的"Hello World"，下面是使用CUDAX实现的高效向量加法示例：

// 核心代码片段 [examples/cudax/vector_add/vector_add.cu]
__global__ void vectorAdd(cudax::span<const float> A, cudax::span<const float> B, cudax::span<float> C) {
  int i = blockDim.x * blockIdx.x + threadIdx.x;
  if (i < A.size()) {
    C[i] = A[i] + B[i];
  }
}

// 主机端代码
cudax::stream stream(cuda::devices[0]);
cudax::vector<float> A(numElements), B(numElements), C(numElements);
// ... 初始化数据 ...
constexpr int threadsPerBlock = 256;
auto config = cuda::distribute<threadsPerBlock>(numElements);
cudax::launch(stream, config, vectorAdd, in(A), in(B), out(C));
stream.sync();

这个示例展示了CUDAX的现代C++接口，通过cudax::vector管理设备内存，使用cudax::launch简化内核启动流程，让并行编程更加直观。

2.2 数组求和：Thrust库的简洁实现

Thrust库提供了类似STL的接口，让并行算法实现变得异常简单。下面是使用Thrust进行数组求和的示例：

// 核心代码片段 [examples/thrust_flexible_device_system/example.cpp]
constexpr std::size_t N = 1000;
thrust::device_vector<int> data(N, 1);
const auto result = thrust::reduce(data.cbegin(), data.cend());
std::cout << "Sum: " << result << '\n';

仅需3行代码，Thrust就完成了从数据分配到并行求和的全部工作。Thrust支持多种后端系统，可自动检测并使用CUDA加速：

// 自动检测设备系统 [examples/thrust_flexible_device_system/example.cpp]
std::cout << "Detected device system: ";
#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
  std::cout << "CUDA" << '\n';
#elif THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_TBB
  std::cout << "TBB" << '\n';
#elif THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_OMP
  std::cout << "OMP" << '\n';
#elif THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CPP
  std::cout << "CPP" << '\n';
#endif

2.3 并行归约：CUB库的高性能实现

CUB库提供了高度优化的并行原语，下面是使用CUB的BlockReduce实现块级并行求和的示例：

// 核心代码片段 [examples/basic/example.cu]
__global__ void sumKernel(int const* data, int* result, std::size_t N) {
  using BlockReduce = cub::BlockReduce<int, block_size>;
  __shared__ typename BlockReduce::TempStorage temp_storage;
  
  int 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);
  }
}

这个示例结合了CUB的块归约和libcu++的原子操作，展示了CCCL组件间的无缝协作。

三、中级并行算法示例

3.1 并行排序：Thrust库的高效实现

Thrust提供了多种排序算法，下面是使用Thrust进行并行排序的示例：

#include <thrust/sort.h>
#include <thrust/device_vector.h>

int main() {
  thrust::device_vector<int> vec(10000);
  thrust::generate(vec.begin(), vec.end(), rand);
  thrust::sort(vec.begin(), vec.end());
  return 0;
}

Thrust的排序算法针对GPU架构进行了深度优化，性能远超CPU实现。

3.2 流压缩：筛选满足条件的元素

流压缩（Stream Compaction）是从数组中筛选出满足特定条件元素的操作，在稀疏数据处理中非常有用：

#include <thrust/device_vector.h>
#include <thrust/remove.h>

int main() {
  thrust::device_vector<int> data(1000);
  thrust::generate(data.begin(), data.end(), [](){ return rand() % 10; });
  
  // 保留大于5的元素
  auto new_end = thrust::remove_if(data.begin(), data.end(), 
    [] __device__ (int x) { return x <= 5; });
  
  data.resize(new_end - data.begin());
  return 0;
}

3.3 前缀和：并行扫描算法

前缀和（Prefix Sum），也称为扫描（Scan），是许多并行算法的基础：

#include <thrust/device_vector.h>
#include <thrust/scan.h>

int main() {
  thrust::device_vector<int> data{1, 2, 3, 4, 5};
  thrust::device_vector<int> result(data.size());
  
  // 计算包含性前缀和
  thrust::inclusive_scan(data.begin(), data.end(), result.begin());
  // result: [1, 3, 6, 10, 15]
  
  return 0;
}