大规模并行哈希图设计的关键考虑事项

本教程将深入介绍GPU加速哈希图（即哈希表）的工作原理、实现方式与性能优势，帮助你掌握在CUDA C++中高效使用哈希数据结构的具体方法。

哈希图基础知识

哈希映射是一种 关联容器，它以 键值对 的形式存储数据，其中每个键（key）都映射到对应的值（value）。通过查找键，即可快速检索出相应的值。例如，用哈希图实现电话簿时，将姓名作为键，电话号码作为值。

哈希映射与其他关联容器的核心区别在于：其插入或检索等操作的平均时间复杂度为 常数时间。C++ 标准模板库中的 std::map 通常采用二叉搜索树实现（而非哈希表），而 std::unordered_map 则更接近本文讨论的哈希表类型。在本文中，“哈希表”与“哈希图”视为同义词，可互换使用。

单值与多值哈希表比较

讨论哈希表时，一个重要的区分在于 是否允许重复键：

• 单值哈希表（如 std::unordered_map）要求键唯一。
• 多值哈希表（如 std::unordered_multimap）允许重复键存在。

用电话簿类比：一个人可以拥有多个电话号码。例如，一条记录为 (k=Alice, v=408-555-0148)，另一条记录为 (k=Alice, v=408-555-3847)，这就属于重复键的情形。

存储与检索机制

从概念上讲，哈希映射由一组 桶数组 构成，每个桶可以包含一个或多个键值对。插入新键值对时，执行以下步骤：

1. 将键输入哈希函数，生成哈希值。
2. 使用该哈希值选择一个桶。
3. 若桶可用，则存储该键值对。

例如，插入 (Alice, 408-555-0148) 时：hash(Alice)=4，选择桶 4 存储。检索时同样通过 hash(Alice) 定位桶 4，取出对应的值。

哈希冲突

如果桶的数量等于可能的键数量，则每个键能唯一映射到一个桶。但这通常不现实，因为键的数量未知，或为每个键预留存储会消耗过多空间。例如，电话簿不可能为宇宙中所有可能的名字保留条目。

因此，哈希函数通常难以做到完美，会导致 哈希冲突：两个不同的键映射到相同的哈希值（图1）。一个好的哈希函数会尽量降低冲突概率，但在大多数情况下，冲突仍不可避免。

图 1。两个不同的键（Alice 和 Bob）具有相同的哈希值，导致桶 4 处发生哈希冲突

开放寻址与线性探测

解决哈希冲突有多种策略，本文重点介绍 开放寻址（open addressing） 配合 线性探测（linear probing） 的方法。

开放寻址哈希表使用内存中连续的桶数组。线性探测策略如下：

• 若位置 i 的桶已被占用，则检查下一个相邻位置 i+1。
• 如果 i+1 也被占用，则继续检查 i+2，依此类推。
• 到达最后一个桶后，绕回数组开头。

这种探测方案对每个键是确定性的（图2）。

图 2：开放寻址通过探测方案，将冲突条目存储在不同位置，探测序列以确定性顺序遍历备选桶

小提示：线性探测的缓存效率较高，因为它访问的是连续内存位置。但如果负载因子（已填充桶与总桶的比率）过高，会因额外内存读取而导致性能下降。

例如，检索键 “Bob” 时：从位置 hash(Bob)=4 开始，沿探测序列查找，直到在位置 6 找到目标桶。若探测过程中遇到空桶，则说明该键不在映射中。

常见问题

Q：开放寻址与链地址法（分离链接）相比有哪些优缺点？
A：开放寻址能节省指针开销，缓存友好，但负载因子高时性能退化明显；链地址法处理冲突简单，但需额外内存存储链表指针，且随机访问模式对缓存不友好。

随机内存访问：GPU的优势

精心设计的哈希函数会最大化不同键分散到不同内存位置的概率，导致哈希表操作的内存访问模式 本质上具有随机性。理解随机内存访问的性能至关重要。

表1对比了现代 CPU 和 GPU 的理论峰值带宽与随机 64 位读取的实际带宽（通过 GUPS 基准测试测量）。

表1：带宽计算方式为访问大小 × 访问次数 ÷ 时间

随机内存访问大约比理论峰值带宽慢 10 倍，因为内存子系统针对顺序访问进行了优化。但 NVIDIA GPU 的随机访问吞吐量比现代 CPU 高出一个数量级。这表明：

• 性能最好的 CPU 哈希表，其性能可能比性能最好的 GPU 哈希表慢一个数量级。
• GPU 的随机访问能力使其非常适合哈希表操作。

小提示：若想在自己的系统上运行 GUPS GPU 基准测试，可参考 NVIDIA developer blog code samples GitHub 仓库。CPU 代码可在 ParRes/Kernels GitHub 仓库中找到。

GPU 哈希图实现

本节讨论在 GPU 上实现哈希表的具体方法，并与 CPU 实现进行对比。目标并非替代 std::unordered_map，而是为大规模并行、高吞吐量场景设计哈希表。

示例假设以下简化条件：

• 表的容量固定，不能添加超出初始容量的键值对。
• 需要设置一个哨兵值，用于表示空桶。
• 键和值类型的大小总和 ≤ 8 字节。
• 插入后不能删除键值对。

这些限制并非绝对，cuCollections 库中的更高级实现可以克服这些限制。

基础实现：开放寻址 + 原子操作

示例哈希表采用开放寻址，桶数组初始化时每个桶都携带空哨兵。冲突解决采用线性探测。

为了支持多线程并发更新，每个桶定义为 cuda::std::atomic>，通过原子比较交换（CAS）操作避免数据竞争。

简化版的插入函数如下：

__device__ bool insert(Key k, Value v) {
    // 根据键的哈希值获取初始探测位置
    auto i = hash(k) % capacity;
    while (true) {
      // 加载当前探测位置桶的内容
      auto [old_k, old_v] = buckets[i].load(memory_order_relaxed);
      // 如果桶为空，则尝试插入键值对
      if (old_k == empty_sentinel) {
        // 尝试原子性地将桶的当前内容替换为输入键值对
        bool success = buckets[i].compare_exchange_strong(
                        {old_k, old_v}, {k,v}, memory_order_relaxed);
        if (success) {
          // 存储成功
          return true;
        }
      } else if (old_k == k) {
        // 输入键已存在于映射中
        return false;
      }
      // 如果桶已被占用，则移动到下一个（线性）探测位置
      // 使用模运算，在超出容量时绕回开头
      i = ++i % capacity;
    }
}

查找操作类似：沿着键的探测序列检查，直到找到目标或遇到空桶为止。

协作组优化

最初，为每个输入元素分配一个线程看似合理，但存在两个问题：

• 相邻键的探测位置在内存中完全随机，一个 warp 中的 32 个线程可能访问 32 个不同位置（最坏情况）。
• 线性探测使每个线程可能访问多个相邻桶，但单个线程无法利用合并加载进行预取。

利用 CUDA 协作组（Cooperative Groups） 模型，可以将一个输入元素分配给 warp 内的一组连续线程（例如 4 个线程）。对于给定键：

1. 使用单个合并加载，预取多个相邻桶的窗口。
2. 组内通过高效的 ballot 和 shfl 内建函数，协同确定窗口内的候选桶。

图 3：键 “Bob” 的组协作探测步骤及中间步骤

以下代码展示了使用 warp 内 4 个连续线程协同插入一个键的简化版本：

enum class probing_state { SUCCESS, DUPLICATE, CONTINUE };

__device__ bool insert(cg::thread_block_tile<4> group, Key k, Value v) {
    // 根据键的哈希值，结合组内线程排名获取初始探测位置
    auto i = (hash(k) + group.thread_rank()) % capacity;
    auto state = probing_state::CONTINUE;
    while (true) {
      // 以合并加载方式，读取每个线程当前探测位置桶的内容
      auto [old_k, old_v] = buckets[i].load(memory_order_relaxed);
      // 如果输入键已存在于映射中
      if(group.any(old_k == k)) return false;
      // 每个线程检查其当前桶是否为空，即是否为候选插入桶
      auto const empty_mask = group.ballot(old_k == empty_sentinel);
      // 如果当前探测窗口内存在空桶
      if(empty_mask) {
        // 选出一个候选线程（此处为掩码中排名最低的线程）
        auto const candidate = __ffs(empty_mask) - 1;
        if(group.thread_rank() == candidate) {
          // 尝试原子性地将输入键值对交换到桶中
          bool const success = buckets[i].compare_exchange_strong(
                          {old_k, old_v}, {k, v}, memory_order_relaxed);
          if (success) {
            // 插入成功
            state = probing_state::SUCCESS;
          } else if (old_k == k) {
            // 否则，重新检查当前探测位置是否已被插入重复键
            state = probing_state::DUPLICATE;
          }
        }
        // 将插入结果从候选线程广播到所有其他线程
        auto const candidate_state = group.shfl(state, candidate);
        if(candidate_state == probing_state::SUCCESS) return true;
        if(candidate_state == probing_state::DUPLICATE) return false;
      } else {
        // 否则，移动到下一个（线性）探测窗口
        i = (i + group.size()) % capacity;
      }
    }
}

这是 cuCollections 中 cuco::static_map 实际实现的简化版本。

图4展示了在 NVIDIA A100 80GB GPU 上，不同组大小和表占用率下，非协作与协作探测的性能对比（吞吐量以 GB/s 计，越高越好）。

图 4：通过协作探测，吞吐量（GB/s）越高越好。红色虚线为峰值 GUPS 结果（吞吐量上限）

观察结论：

• 负载因子较低时，非协作方法接近最优。
• 负载因子增加后，冲突增多、探测序列变长，非协作方法的吞吐量急剧下降。
• 协作探测（组大小为 4）在高负载因子下，插入吞吐量高出 13%，查找吞吐量高出 40%。

长探测序列也出现在多值场景（高键乘数）中，协作探测同样有效。更多信息可参见 Parallel Hashing on Multi-GPU Nodes 和 WarpCore: A Library for Fast Hash Tables on GPUs。

常见问题

Q：协作组的分组大小如何选择？为什么是4？
A：分组大小需要权衡：分组太小无法充分利用合并加载；分组太大会导致组内同步开销增大。实测表明，4 个线程的 tile 在多数负载下能达到最佳平衡。

现有 CPU 与 GPU 哈希图性能对比

多年来，已出现多种 C++ 哈希图实现：

• CPU 顺序实现：libstdc++/libc++ std::unordered_map、absl::flat_hash_map。
• CPU 并发实现：TBB 的 tbb::concurrent_hash_map、Folly 的 folly::AtomicHashMap。
• GPU 实现：Kokkos 的 kokkos::UnorderedMap、cuCollections 的 cuco::static_map。

基准设置：插入 2^27（约1GB）个唯一 4B 键 / 4B 值对，然后查询同一组键以检索值。目标负载因子为 50%。性能以内存吞吐量（GB/s）衡量。

结果如图5所示：

图 5：流行 CPU 和 GPU 哈希图实现的性能比较

• cuco::static_map 在单个 NVIDIA H100-80GB-SXM 上实现了插入 87.5 GB/s，查找 134.6 GB/s 的吞吐量。
• 比最快的 CPU 单线程/多线程实现快 超过一个数量级。
• 优于其他 GPU 实现：插入速度比 Kokkos 快 3.8 倍，查找速度快 2.6 倍。

注意：CPU 测试数据驻留于 CPU 内存，GPU 测试数据驻留于 GPU 内存。若数据必须在 CPU 与 GPU 间移动，可通过显式复制或 CUDA Unified Memory 实现。结果表明，cuCollections 的吞吐量远超 PCIe Gen4/Gen5 的实际带宽，能完全饱和 CPU-GPU 链路。

小提示：NVIDIA Grace Hopper Superchip 利用 NVLink-C2C 互连，可进一步释放哈希表的全部吞吐量，而 CPU 哈希映射通常达不到 PCIe 带宽。

实际应用：多列关系联接示例

cuDF 是 GPU 加速的数据分析库，它利用 cuCollections 哈希表实现 哈希连接算法。图6展示了内部连接的构建与探测阶段。

图 6：RAPIDS cuDF 中内部连接的构建和探测阶段

过程如下：

1. 使用内置哈希函数，将任意类型的行哈希为值。
2. 左侧表填充 cuco::static_multimap，键为行的哈希值，值为行索引。例如，行 24 映射到桶 47，行 25 映射到桶 48。
3. 探测阶段：右侧表行 200 的哈希值为 47，与桶 47 匹配。然后通过行相等函数 row_equal(200, 24) 比较实际行内容。若不匹配，则继续探测；若匹配，则输出结果。

最终，行 25 与行 200 匹配（哈希值相同且行相等检查通过）。

连接操作基准涉及许多选项（大小、选择性等），可参考 How to Get the Most out of GPU Accelerated Database Operators 和 Effective, Scalable Multi-GPU Joins。

如何在代码中使用 GPU 哈希图

GPU 凭借高带宽存储器（HBM）架构，在处理许多小随机读取和原子更新时，比 CPU 快一个数量级，这直接转化为高效的哈希表插入和探测性能。在设计大规模并行哈希图时，关键考虑因素包括：

1. 采用开放寻址的平面内存布局。
2. 利用协作组优化探测过程。
3. 合理选择负载因子与组大小。

作为 cuCollections 库的一部分，你可以在 GitHub 上找到快速灵活的哈希图实现（如 cuco::static_map、cuco::static_multimap 等）。此外，RAPIDS cuDF 集成了 GPU 哈希图，助力数据科学工作负载实现惊人加速。更多信息请访问 GitHub 上的 rapidsai/cudf 和 使用 Dask 和 RAPIDS 为自然语言处理加速 TF-IDF。

cuCollections 还可用于：推荐系统、流压缩、图形算法、基因组学、稀疏线性代数等领域。

小提示：开始使用 cuCollections 时，建议先尝试简单的固定容量 static_map，再逐步过渡到支持动态扩容的 dynamic_map 版本。

常见问题

Q：cuCollections 是否支持删除操作？
A：基础实现不支持删除，但 cuCollections 提供了更高级的 dynamic_map 等容器，支持动态插入、删除和扩容。具体可查阅官方文档。

Q：如何将 cuCollections 集成到自己的 CUDA 项目中？
A：cuCollections 是仅头文件的库，只需将 include 目录添加到编译器路径，并包含相应头文件（如 #include ）。确保使用 C++17 及支持 CUDA 12+ 的编译器。

你是一名 AI 行业编辑，请围绕下面这条热点输出一份资讯解读：
热点：大规模并行哈希图设计的关键考虑事项要求：
1. 先用一句话解释这条热点在讲什么
2. 再总结它为什么重要
3. 说明会影响哪些 AI 产品或内容方向
4. 最后给出 3 个适合资讯站使用的标题