GPU上哈希表必须用CUDA C++重写，禁用std::unordered_map等CPU结构；需用__device__函数、POD数据、开放寻址法、原子操作及2的幂容量哈希，避免动态分配与分支。
不能直接在 GPU 上用 std::unordered_map 或手写 CPU 哈希表跑起来——它会编译失败或运行崩溃，因为 GPU 不支持动态内存分配、递归、异常、虚函数调用，也不认 std::vector 或 new。
GPU 上哈希表必须用 CUDA C++ 重写，且只用 device-side 可执行的代码
你写的每个函数（插入、查找、删除）都得标上 __device__，甚至多数还得加 __host__ __device__；所有数据结构必须是 POD（plain old data），不能含指针链表、std::string、std::unique_ptr。常见错误现象包括：
- CUDA 编译器报错：
error: calling a __host__ function("std::vector::push_back") from a __device__ function is not allowed - 运行时卡死或返回乱码：因 host 分配的内存没用
cudaMalloc拷贝到 device，或没同步cudaDeviceSynchronize() - 多个线程同时写同一 bucket 导致数据损坏：GPU 线程块内无原子锁机制，必须显式用
atomicCAS或atomicAdd
典型做法是把整个哈希表布局为连续 device 内存块：key[]、value[]、state[]（EMPTY/EXIST/DELETE），用数组索引代替指针跳转。
开放寻址法比拉链法更适合 GPU，尤其线性探测 + 墓碑清理
拉链法依赖动态链表节点分配，在 GPU 上几乎不可行（cudaMalloc 在 kernel 里调用极慢且不推荐）；而开放寻址法所有数据平铺在固定大小数组中，天然契合 GPU 的 SIMD 和 coalesced memory access 模式。但要注意：
- 线性探测易产生“聚集”，导致 cache miss 飙升——实测在 80% 负载下，平均探测长度可能 >5，严重拖慢吞吐
- 必须保留
DELETE状态位，否则删除后查找路径断裂，后续 key 永远查不到 - 负载因子建议压到 ≤0.7；超过 0.75 后性能断崖下跌，此时必须触发异步扩容（host 端 realloc + 数据重哈希）
- 探测步长别硬写
+1，可用(hash + i * i) % capacity（二次探测）缓解聚集，但会增加分支和计算开销
哈希函数必须无分支、无除法、全整数运算
GPU 的 warp 执行模型要求所有线程走相同指令路径，任何 if/else 或取模 % 都会导致 warp divergence。所以：
- 禁用
key % capacity：改用capacity为 2 的幂，用hash & (capacity - 1) - 字符串哈希不能遍历每个 char：改用常量展开（如最多 16 字节 unroll）、SIMD load（
__ldg）或预计算 hash 后传入 device - 推荐 MurmurHash3 的 CUDA 移植版，或更轻量的 xxHash3 的 device 版本；避免
std::hash<:string>—— 它内部有分支和 heap 访问 - 对 int/float key，可用
__funnelshift_r（Turing+ 架构）做快速混洗，比乘加更高效
并发插入必须靠原子操作 + 失败重试，不能靠锁
GPU 没有 mutex，也没有 sleep；多线程竞争同一 slot 时，必须用原子原语保障一致性：
- 插入前先
atomicCAS(&state[idx], EMPTY, EXIST)，成功才写 key/value；失败则按探测策略跳到下一个位置重试 - 删除用
atomicCAS(&state[idx], EXIST, DELETE)，确保状态切换原子性 - 查找不用原子，但需循环读
state[idx]直到遇到EMPTY或匹配 key——注意加__ldg()提升只读缓存命中率 - 避免无限重试：单次 kernel 中最大探测次数建议设为 32～64，超限则返回失败，由 host 侧重试或丢弃
真正难的不是单次操作，而是当数万个线程同时插入不同 key 时，大量线程卡在相同热点 bucket 上反复 CAS 失败——这需要结合分段哈希（partitioned hash table）或两级索引（coarse-grained lock per segment）来缓解，但会增加实现复杂度和内存占用。
C++免费学习笔记（深入）：立即使用
在学习笔记中，你将探索 C++ 的入门与实战技巧！
