CCCL - NVIDIA CUDA 核心计算库，统一的 CUDA C++ 高性能编程基石CCCL (CUDA Core

CUDA Core Compute Libraries (CCCL)

欢迎使用 CUDA Core Compute Libraries (CCCL)。我们的使命是让 CUDA 编程更加愉悦和高效。该仓库将三个必不可少的 CUDA C++ 库统一到一个便捷的仓库中：

Thrust：C++ 并行算法库，灵感来源于 C++ 标准库并行算法。
CUB：C++ 块级并行原语库，为构建高性能内核提供基础。
libcudacxx：CUDA C++ 标准库，提供 CUDA 设备代码中的标准库支持。

统一这些库的目标是为 CUDA C++ 开发者提供构建模块，使他们能够更轻松地编写安全且高效的代码。将这三个项目合并可以简化您的开发流程，并拓宽您利用 CUDA C++ 强大功能的能力。

功能特性

统一的高质量库：在一个仓库中同时访问 Thrust 的高级算法、CUB 的块级原语以及 libcudacxx 的 CUDA 标准库支持。
高性能并行抽象：提供经过优化的并行算法，如排序、归约、扫描等，能够充分利用现代 GPU 的吞吐量。
模块化与可组合性：CUB 的块级原语允许开发者构建复杂的自定义内核，同时保持高性能。
现代化 C++ 支持：支持 C++14/17/20 标准，并与 CUDA C++ 特性深度集成。
持续集成与测试：项目通过广泛的矩阵测试（涵盖不同 CUDA 版本、编译器、GPU 架构和操作系统）来确保代码质量和兼容性。
开发容器支持：提供基于 Dev Containers 的一键式开发环境，简化了开发环境的设置，无论是在本地还是 CI 中。

安装指南

前提条件

支持 CUDA 的 GPU 和 NVIDIA 驱动程序
NVIDIA CUDA 工具包 (版本 11.0 或更高)
支持 C++14/17 的编译器 (GCC 7+, Clang 8+, MSVC 2019+)
CMake 3.18 或更高版本

使用 CMake 安装

CCCL 推荐使用 CMake 的 find_package 机制进行集成。

# 克隆仓库
git clone https://github.com/nvidia/cccl.git
cd cccl

# 创建一个构建目录
mkdir build && cd build

# 配置 CMake。将 <install_path> 替换为目标安装路径。
cmake .. -DCMAKE_INSTALL_PREFIX=<install_path>

# 编译并安装
cmake --build . --target install

安装完成后，您可以在您的 CMake 项目中通过以下方式找到 CCCL：

find_package(CCCL REQUIRED)

# 然后链接所需的库
target_link_libraries(your_target PRIVATE CCCL::thrust CCCL::cub)

使用 CMake 预设

CCCL 提供了 CMake 预设来简化构建过程。您可以使用以下命令列出所有可用的预设：

cmake --list-presets

要配置并使用特定的预设（例如 cub-cpp17），请运行：

cmake --preset cub-cpp17
cmake --build --preset cub-cpp17

使用说明

Thrust 示例：向量加法

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <iostream>

int main() {
    // 在主机上分配向量
    thrust::host_vector<int> h_a(5, 1);
    thrust::host_vector<int> h_b(5, 2);
    thrust::host_vector<int> h_c(5);

    // 将数据传输到设备
    thrust::device_vector<int> d_a = h_a;
    thrust::device_vector<int> d_b = h_b;
    thrust::device_vector<int> d_c = h_c;

    // 在设备上执行向量加法
    thrust::transform(d_a.begin(), d_a.end(), d_b.begin(), d_c.begin(),
                      thrust::plus<int>());

    // 将结果拷贝回主机
    thrust::copy(d_c.begin(), d_c.end(), h_c.begin());

    // 输出结果
    for (int i = 0; i < 5; i++) {
        std::cout << "c[" << i << "] = " << h_c[i] << std::endl;
    }
    return 0;
}

CUB 示例：设备级归约

#include <cub/cub.cuh>
#include <iostream>

int main() {
    const int N = 1024;
    int *d_in, *d_out;

    // 分配设备内存
    cudaMalloc(&d_in, N * sizeof(int));
    cudaMalloc(&d_out, sizeof(int));

    // 初始化输入数据 (此处省略初始化代码)

    // 确定临时存储大小
    void *d_temp_storage = nullptr;
    size_t temp_storage_bytes = 0;
    cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_in, d_out, N);

    // 分配临时存储
    cudaMalloc(&d_temp_storage, temp_storage_bytes);

    // 运行归约
    cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_in, d_out, N);

    // 将结果拷贝回主机
    int h_out;
    cudaMemcpy(&h_out, d_out, sizeof(int), cudaMemcpyDeviceToHost);
    std::cout << "Sum is: " << h_out << std::endl;

    cudaFree(d_temp_storage);
    cudaFree(d_in);
    cudaFree(d_out);
    return 0;
}

开发容器 (Dev Container) 快速启动

对于希望快速搭建开发环境的用户，CCCL 支持 Dev Containers。

前提条件：安装 VSCode、Remote - Containers 扩展和 Docker。
克隆仓库：git clone https://github.com/nvidia/cccl.git
在 VSCode 中打开：打开克隆后的目录。
重新在容器中打开：按 F1，选择 Rebuild and Reopen in Container。

VSCode 会自动构建并启动一个包含所有必要依赖（特定版本的 CUDA 工具包、编译器）的容器。

核心代码

以下是项目中的一些关键核心代码模块，展示了其底层机制和高级封装。

1. 迭代器状态机与 JIT 生成

为了支持 JIT (Just-In-Time) 编译，CCCL 的 C 并行库部分定义了迭代器的状态和操作。下面的代码展示了如何为状态迭代器生成 C++ 类型定义。

// 文件: c/parallel/include/kernels/iterators.h
// 功能: 为 JIT 编译生成输入迭代器的 C++ 类型定义
constexpr std::string_view format_template = R"XXX(
#define DIFF_T {0}
#define OP_ALIGNMENT {1}
#define OP_SIZE {2}
#define VALUE_T {3}
#define DEREF {4}
#define ADVANCE {5}

// Kernel Source
{6}

#undef DIFF_T
#undef OP_ALIGNMENT
#undef OP_SIZE
#undef VALUE_T
#undef DEREF
#undef ADVANCE
)XXX";

std::string make_kernel_input_iterator(
  std::string_view diff_t,
  size_t alignment,
  size_t size,
  std::string_view iterator_name,
  std::string_view value_t,
  std::string_view deref,
  std::string_view advance)
{
  const std::string iter_def = std::format(R"XXX(
extern "C" __device__ void DEREF(const void *self_ptr, VALUE_T* result);
extern "C" __device__ void ADVANCE(void *self_ptr, DIFF_T offset);
struct __align__(OP_ALIGNMENT) {0} {{
  // ... 迭代器接口实现 ...
  __device__ inline value_type operator*() const {{
    value_type result;
    DEREF(data, &result);
    return result;
  }}
  // ... 其他成员函数 ...
}})XXX", /* ... 参数 ... */);
  // ... 返回完整代码字符串
}

2. 运行时策略提取

CCCL 能够从 CUB 的编译时策略中提取信息，用于运行时决策。这通过 ptx-json 库实现，将设备端的策略信息序列化为 JSON。

// 文件: c/parallel/src/util/runtime_policy.cc
// 功能: 通过编译一个小的翻译单元并解析生成的 PTX 中的 JSON 数据来获取策略信息。
nlohmann::json
get_policy(std::string_view policy_wrapper_expr, std::string_view translation_unit, std::span<const char*> args)
{
  // ... 设置编译参数 ...
  fixed_args.push_back("-DCUB_ENABLE_POLICY_PTX_JSON");
  fixed_args.push_back("-std=c++20");
  fixed_args.push_back("-default-device");

  std::string fixed_source = std::format(
    "#include <cub/detail/ptx-json/json.h>\n"
     "{0}\n"
     "__global__ void ptx_json_emitting_kernel()\n"
     "{{\n"
     "  [[maybe_unused]] auto wrapped = {1};\n"
     "  ptx_json::id<ptx_json::string(\"{2}\")>() = wrapped.EncodedPolicy();\n"
     "}}\n",
    translation_unit,
    policy_wrapper_expr,
    tag_name);

  // 编译并链接以获得 PTX
  auto nvrtc_ptx = begin_linking_nvrtc_program(0, nullptr)
                     ->add_program(nvrtc_translation_unit{...})
                     ->compile_program(...)
                     ->get_program_ptx();

  // 解析 PTX 中的 JSON
  return cub::detail::ptx_json::parse(tag_name, nvrtc_ptx.ptx.get());
}

3. 基准测试分析工具

CCCL 包含一套强大的 Python 工具，用于分析和比较不同版本之间的基准测试结果，帮助识别性能回归。

# 文件: ci/bench_analysis.py
# 功能: 比较两个基准测试数据库文件，并标记性能变化
def compare():
    args = parse_args()
    reference_df = alg_dfs(args.reference)
    compare_df = alg_dfs(args.compare)
    for alg in sorted(reference_df.keys() & compare_df.keys()):
        # ... 合并两个数据库的数据 ...
        df = pd.merge(reference_df[alg], compare_df[alg], on=merge_columns, suffixes=("Ref", "Cmp"))
        df["Abs. Diff"] = df["MeanCmp"] - df["MeanRef"]
        df["Rel. Diff"] = (df["Abs. Diff"] / df["MeanRef"]) * 100
        df["Status"] = list(map(status, df["Rel. Diff"], df["NoiseRef"], df["NoiseCmp"]))
        # ... 打印表格形式的结果 ...
    # 输出统计摘要
    print("# Summary\n")
    print("- Total Matches: %d" % config_count)
    print("  - Pass    (diff <= min_noise): %d" % pass_count)
    print("  - Faster  (diff > min_noise):  %d" % faster_count)
    print("  - Slower  (diff > min_noise):  %d" % slower_count)

4. 编译缓存与构建加速

CI 脚本利用 sccache 来缓存编译结果，加速构建过程。以下脚本片段用于计算并输出 sccache 的命中率。

# 文件: ci/sccache_stats.sh
# 功能: 计算两次调用 sccache --show-stats 之间的缓存命中率
# 使用方法: source ci/sccache_stats.sh start  # 在构建前
#          source ci/sccache_stats.sh end    # 在构建后

case $mode in
  start)
    export SCCACHE_START_HITS=$(sccache --show-stats | awk '/^[ \t]*Cache hits[ \t]+[0-9]+/ {print $3}')
    export SCCACHE_START_MISSES=$(sccache --show-stats | awk '/^[ \t]*Cache misses[ \t]+[0-9]+/ {print $3}')
    ;;
  end)
    # ... 获取结束时的统计数据 ...
    hits=$((final_hits - SCCACHE_START_HITS))
    misses=$((final_misses - SCCACHE_START_MISSES))
    total=$((hits + misses))

    if (( total > 0 )); then
      hit_rate=$(awk -v hits="$hits" -v total="$total" 'BEGIN { printf "%.2f", (hits / total) * 100 }')
      echo "sccache hits: $hits | misses: $misses | hit rate: $hit_rate%"
    fi
    ;;
esac

svaqnhrkxBCO3lCiEoYbYnWVOPw59Va9feAuss5k76c=