企业官网建设流程全解析

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend

本文以 asc-devkit 仓库中 SIMT 实现 InsertHashTable 算子案例 为主线，完整剖析一个基于 SIMT（Single Instruction, Multiple Threads）编程模型、由大规模线程并发访问 Global Memory 的哈希表批量插入算子。读者通过本文可以掌握：哈希表与开放寻址法的基本工作原理、SIMT 多线程写冲突与核间数据同步两大难题的成因、如何用asc_atomic_cas()原子比较赋值实现无锁抢占、如何用asc_threadfence()内存栅栏保证跨核数据可见性，以及如何阅读、编译、运行并验证仓库中完整的insert_hash_table样例。

案例介绍

本案例对应仓库中的 insert_hash_table 样例，目标是实现一个InsertHashTable 算子：将 8×1024 个键值对批量插入容量为 16×1024 的哈希表，其中 key 为int64_t类型，value 为长度 32 的floatTensor。与传统的 SIMD 向量化算子不同，本算子采用 SIMT 编程模型，由数千个线程并发访问 Global Memory，是理解昇腾 AI 处理器上"大规模线程 + 离散访存 + 无锁同步"编程范式的典型范例。

从样例 README 的算子规格表（见 insert_hash_table/README.md）可以看到：

该样例支持Ascend 950PR & 950DT 系列产品，要求 CANN 软件版本不低于 9.1.0，编译目标架构为dav-3510。实现上采用"每个 warp 处理一个键值对、warp 内线程 0 负责寻址、线程 0-31 协作写入 value"的并行策略。

背景知识：哈希表与开放寻址法

哈希表的基本结构

哈希表（Hash Table，也称散列表）是一种高效的数据结构，它通过哈希函数将"键"（Key）映射到固定大小数组（桶数组）的特定位置，从而实现快速的查找、插入和删除操作。三个核心概念如下：

• 哈希函数：接收一个键作为输入，计算出一个整数（哈希值），该整数通过某种运算（通常是取模）被映射为数组的索引；
• 桶数组：一个数组，每个元素（称为"桶"，Bucket）用于存储键值对（Key-Value Pair）；
• 插入流程：先用哈希函数计算 key 的哈希值，再将哈希值映射为桶数组索引 index（通常是对容量取余），最后将键值对存入bucket[index]位置。

![] (docs/zh/guide/figures/hash_insert.png "向哈希表中插入元素")

哈希冲突与开放寻址法

由于哈希表容量（桶数组大小）远小于哈希函数的输出范围，不同的键可能被映射到同一个索引，即产生哈希冲突。常见的解决方法之一是开放寻址法（Open Addressing）：当冲突发生时，算法按线性探查（Linear Probing）顺序寻找下一个"空"的桶来存储数据。

以图 2 为例：根据 key 计算出的哈希值得到的初始位置索引为 3，当与桶bucket[3]中的已有数据发生冲突时，向后逐个探测，直到找到可用的桶；若遍历完整个桶数组仍未找到可用桶，则插入失败。

![] (docs/zh/guide/figures/open_addr.png "开放寻址法")

为什么哈希表适合 SIMT 而不是 SIMD

由于每个键值对的存储位置由哈希值决定，这些值通常是随机且离散的，加上哈希冲突的存在，实际写入数据前需要进行多个条件判断，因此哈希表不适合基于 SIMD 编程模型实现——SIMD 要求数据并行路径一致，分支和离散访存会严重损害效率。相比之下，SIMT 编程中每个线程都能独立处理分支判断并支持离散访存，因此在实现哈希表方面更具优势。

算法分析：SIMT 高并发带来的两大难题

SIMT 能通过大量线程的并发执行高效处理大批数据，但同时也带来两个必须解决的问题：多线程写冲突与核间数据同步。

多线程写冲突与原子操作

在多线程操作同一块内存时，资源冲突在所难免。当两个线程插入的 key 产生哈希冲突时，多个线程会尝试向桶数组的同一位置写入数据，因此需要一种机制确保仅有一个线程能获得桶的写入权限。在 CPU 上通常使用互斥锁，但 SIMT 架构中无法使用锁，只能通过原子操作实现无锁编程。

例如图 3 中，线程 0 和线程 256 处理的 key 发生哈希冲突，计算出的位置索引均为 3。此时应确保仅有一个线程能获得bucket[3]的写入权限，另一个线程则通过开放寻址法寻找下一个可用桶。

![] (docs/zh/guide/figures/mt_hash.png "多线程哈希冲突")

程序实现时，在Bucket结构体中添加标记位flag标记当前桶的写入权限，线程通过原子指令asc_atomic_cas修改 flag 来获取桶的写权限：

inline int32_t asc_atomic_cas(int32_t *address, int32_t compare, int32_t val)

根据 asc_atomic_cas API 文档，该接口对指定地址上的数据执行原子比较赋值：若address上的值等于compare，则将address赋值为val，否则地址值不变。接口返回地址上的初始数据，并保证多个线程并发修改同一地址时最多只有一个线程修改成功（但不保证固定由哪个线程成功）。头文件为simt_api/device_atomic_functions.h。

在本案例中，flag的取值包括DEFAULT_FLAG（桶为空）、HOLD_FLAG（桶中有数据）和WRITE_FLAG（桶正在被写入）。使用asc_atomic_cas修改 flag 时：

• compare = DEFAULT_FLAG、val = WRITE_FLAG：尝试获取空桶的写权限；
• compare = HOLD_FLAG、val = WRITE_FLAG：尝试获取非空桶的写权限（用于更新已有键值对）。

根据返回值与 compare 是否相等即可判断是否成功获得写权限。样例源码中定义了三个标志位常量（见 insert_hash_table.asc）：

constexpr uint32_t DEFAULT_FLAG = 0; constexpr uint32_t HOLD_FLAG = 0x57277275; constexpr uint32_t WRITE_FLAG = 0x39555593;

值得注意的是，原子操作文档 还提示了性能约束：Global Memory 原子操作经过 L2 Cache 处理，L2 Cache 以 512B Cache Line 为缓存管理单位，每个 Cache Line 包含 4 个 128B 扇区（Sector），原子操作以 128B Sector 为处理粒度。目标地址集中在一个 Sector 内时处理效率较低；业务允许时，建议将原子操作的目标地址分散到更多 Sector 中。

核间数据同步与内存栅栏

当发生哈希冲突时，线程需要判断当前桶中存储的 key 是否与待插入的 key 相同，这要求当前线程能读取到其他线程写入的数据，并确保数据完整。SIMT 基于**弱内存模型（Weak Memory Model）**实现：线程的内存操作顺序对其他线程而言是不确定的，且 Data Cache 与 Global Memory 之间不保证一致性——一个核中的线程写入 Global Memory 的数据，另一个核可能无法立即读到。因此，当多个核操作同一块 Global Memory 时，需要机制保证核间数据同步。

弱内存模型允许硬件对内存访问指令进行重排序以追求更高性能。考虑如下代码，假设A和flag是共享的全局变量，初始值都为 0：

// core 0 A = 1; // store A flag = 1; // store flag // core 1 while (flag != 1); // load flag (spin until flag=1) int r1 = A; // load A

在弱内存模型下，core 0 的store A和store flag可能被重排序，导致其他核先看到flag = 1，而A = 1的更新尚未到达共享缓存，此时 core 1 可能读到r1 = 0，与程序顺序执行逻辑不符。常用解法是使用内存屏障强制保证屏障两侧的内存操作不被重排序，asc_threadfence()接口正是为此设计。

根据 asc_threadfence API 文档，该接口（头文件simt_api/device_sync_functions.h）用于保证所有线程对同一份全局、共享内存访问过程中写入操作的时序性和可见性，关键特征如下：

• 不阻塞线程，仅保证内存操作的可见性顺序；
• 作用范围是全局的：调用后，确保调用线程在asc_threadfence()之前的所有全局内存和共享内存写操作对其他线程可见。

典型适用场景包括：生产者-消费者同步（确保数据写入在标志位设置之前完成）、原子计数器更新前的数据落盘、链表/树等指针结构的数据一致性，以及多线程共享数据的可见性。

![] (docs/zh/guide/figures/mem_model.png "内存模型示意图")

回到本案例：程序在Bucket结构体中添加标记位state标识 key 的写入状态。写线程写入 key 后将 state 置 1，并在二者之间调用asc_threadfence()，确保当 state 被设置为 1 时 key 的写入已完成；读线程通过 while 循环轮询 state，直到 state 为 1 后再读取 key，判断当前 key 是否与桶中 key 一致：

// 写线程 bucket->key = key; asc_threadfence(); bucket->state = 1; // 读线程 while (*reinterpret_cast<volatile uint32_t*>(&bucket->state) != 1) {} if (key == *reinterpret_cast<volatile int64_t*>(&bucket->key)) ...

另一个关键点是volatile 关键字的使用。由于 Data Cache 与 Global Memory 之间不保证一致性，核 0 写入 Global Memory 的数据会更新到 L2 Cache，但核 1 读取时若 Data Cache 命中，会直接使用 Data Cache 中的旧值而非从 L2 Cache 获取新值。因此，需要读取其他核写入 Global Memory 的数据时，应使用volatile关键字，确保每次读取都从 Global Memory 获取而非 Data Cache。

算子实现：从数据结构到核心逻辑

数据结构设计

根据上述分析，桶结构体设计如下（Bucket需要 8 字节对齐）：

struct Bucket { int64_t key; // 键 uint32_t state; // key值写入状态标记 uint32_t flag; // 原子操作标志位 float value[32]; // 值 };

样例 README 进一步说明了封装思路：哈希表插入涉及哈希表首地址、容量、桶大小、输入 key/value 等多项信息，若在函数间逐一传递会令接口冗长、耦合增加，因此样例用结构体分别封装哈希表上下文与输入数据。Bucket是哈希表实际存储键值对的目标内存结构；KeyValueData封装待插入键值对信息（输入 key、输入 value、value 维度、键值对数量，并提供load_key/load_value读取方法）；HashTableContext保存哈希表自身参数（首地址、容量、桶大小），提供get_bucket（获取桶地址）、query（查找并占用桶）、insert（写入 value）三个核心方法。

哈希函数：MurmurHash3

样例采用 MurmurHash3 算法计算 key 的哈希值（见 insert_hash_table.asc）：对int64_t类型的 key 按 32 位分块迭代混合，最终经三次雪崩式（avalanche）操作（hash ^= hash >> 16、hash *= 0x85ebca6b等）得到分布均匀的 32 位哈希值。哈希质量直接影响冲突概率与查找效率，MurmurHash3 以良好的随机化特性保证了键在桶数组中的离散分布。

核心逻辑：query 无锁寻址流程

算子核心逻辑的伪代码如下（摘自关联文档）：

uint32_t hash = murmur_hash_3(key); // 采用MurmurHash3算法计算哈希值 uint32_t index = hash & (capacity - 1); // 通过hash得到初始位置索引 Bucket* bucket = nullptr; for (uint32_t i = 0; i < capacity; i++) { bucket = &buckets[index]; int32_t old_flag = asc_atomic_cas(&bucket->flag, DEFAULT_FLAG, WRITE_FLAG); // 获取空桶的写权限 if (old_flag == DEFAULT_FLAG) { // 成功获取空桶写权限 bucket->key = key; asc_threadfence(); // 在写入key和state置1中间插入内存栅栏，保证执行顺序 bucket->state = 1; bucket->value = value; bucket->flag = HOLD_FLAG; } else { while (*reinterpret_cast<volatile uint32_t*>(&bucket->state) != 1) {} // 通过state标记位保证key值的有效性 if (key == *reinterpret_cast<volatile int64_t*>(&bucket->key)) { // 桶中存储的key与要插入的key相同 old_flag = asc_atomic_cas(&bucket->flag, HOLD_FLAG, WRITE_FLAG); // 获取已有数据的bucket的写权限 if (old_flag == HOLD_FLAG) { // 成功获取写权限 bucket->value = value; // 更新value bucket->flag = HOLD_FLAG; } } else { index = (index + 1) % capacity; // 线性探测下一个可用桶 } } }

对照样例源码HashTableContext::query（insert_hash_table.asc），实现细节完全一致，且补充了两个关键点：

1. 桶索引映射采用位运算hash & (capacity - 1)而非取余——这要求容量为 2 的幂（源码注释明确HASH_TABLE_CAPACITY = 16 * 1024; // must be 2^n），位运算在算子内比除法取余更高效；
2. 线性探测同样用(bucket_idx + 1) & (capacity - 1)实现环形回绕；若遍历完整个桶数组仍未找到可用桶，query返回nullptr表示插入失败。

query的返回值通过success出参告知调用者是否成功获得桶，其语义为"查找并占用桶"——成功时要么是新插入（flag 从 DEFAULT_FLAG 抢占为 WRITE_FLAG），要么是命中已有 key 并抢占更新权（flag 从 HOLD_FLAG 抢占为 WRITE_FLAG）。

并发写入与 warp 协作

真正的并发写入发生在insert与insert_one_kv中。样例采用每个 warp（32 线程）处理一个键值对的分工：

• 线程 0（lane_id == 0）负责加载 key、调用query寻找可用桶；
• 全部 32 个线程通过asc_shfl（warp 内洗牌指令）广播success与桶地址；
• 获取桶后，insert让线程j处理value[j]，通过j += warpSize的步长循环协作写入长度 32 的 value（正好一个 warp 一轮写满）；
• 写入完成后调用asc_threadfence()，最后由线程 0 将 flag 置回HOLD_FLAG，释放桶并对外发布"数据已就绪"状态。

__aicore__ inline void insert(Bucket* bucket, const KeyValueData* input, uint32_t kv_idx, int32_t lane_id) { for (uint32_t j = lane_id; j < input->value_dim; j += warpSize) { bucket->value[j] = input->load_value(kv_idx, j); } asc_threadfence(); if (lane_id == 0) { bucket->flag = HOLD_FLAG; } }

核函数insert_hash_table的线程编排为（见 insert_hash_table.asc）：全局线程号blockIdx.x * blockDim.x + threadIdx.x除以 warpSize 得到键值对索引，越界线程直接返回；__launch_bounds__(1024)声明每 block 线程数上限为 1024，即每 block 含 32 个 warp，可并行处理 32 个键值对。

多线程冲突处理流程详解

特别说明多线程发生哈希冲突时的完整处理流程。设线程 0 和线程 256 分别计算key[0]与key[8]的初始位置时发生哈希冲突，初始位置索引均为 3：

1. 抢占阶段：线程 0 和线程 256 通过asc_atomic_cas同时抢占bucket[3]的写权限，原子操作保证只有一个线程成功——线程 0 成功获得写权限，将数据写入bucket[3]；线程 256 竞争失败，向后线性探测可用桶；
2. 探测阶段：线程 256 查询bucket[4]，发现其中已有数据（且 key 与待插入 key 不同），继续向后寻找，最终发现bucket[5]可用；
3. 写入阶段：线程 256 通过asc_atomic_cas获取bucket[5]的写权限并写入数据。

整个过程无需任何锁，完全由原子比较赋值 + 线性探测 + 内存栅栏协同完成，体现了"无锁编程"在 SIMT 架构下的完整实践。

编译、运行与结果验证

编译与运行

样例的目录结构如下：

insert_hash_table ├── CMakeLists.txt // cmake编译文件 ├── insert_hash_table.asc // Ascend C算子实现 & 调用样例 └── README.md

在样例根目录下按以下步骤编译执行：

1. 配置环境变量：根据当前环境上 CANN 开发套件包的安装方式配置环境变量，CANN 包未指定安装目录时默认安装至/usr/local/Ascend下：

   source ${install_path}/cann/set_env.sh

2. 编译与执行：

   mkdir -p build && cd build; # 创建并进入build目录 cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..; make -j; # 编译工程 ./demo # 执行样例

3. NPU 仿真模式：添加-DCMAKE_ASC_RUN_MODE=sim参数即可：

   cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..; make -j;

   注意：切换编译模式前需清理 cmake 缓存，可在 build 目录下删除CMakeCache.txt后重新 cmake。

编译选项说明：

结果验证

样例的main函数流程为：用随机数生成器产生KEY_NUM个int64_tkey 与对应 value → 通过 ACL 运行时（aclrtMalloc/aclrtMemcpy/ 内核调用符<<<>>>启动核函数 /aclrtSynchronizeStream同步）在设备端完成批量插入 → 将哈希表内存拷回主机 → 调用check_result逐 key 校验。

check_result（见 insert_hash_table.asc）按与设备端一致的哈希与线性探测顺序查找每个 key，校验桶中state == 1、flag == HOLD_FLAG，并以epsilon = 1e-5f的容差比较 value 的每个分量。全部校验通过后输出：

[Success] find all key-value in hash table.

该输出表明：8×1024 个键值对全部成功插入容量为 16×1024 的哈希表，无 key 丢失、无 value 错写，多线程写冲突与核间同步处理均正确。

总结

本案例展示了基于 SIMT 编程模型向哈希表批量插入元素的完整实践，核心经验可归纳为两点：

1. 用原子指令解决多线程写冲突：SIMT 架构无法使用互斥锁，通过asc_atomic_cas()对桶的flag标志位执行原子比较赋值，保证同一时刻最多一个线程获得桶写权限，配合开放寻址法的线性探测完成冲突转移；
2. 用内存栅栏实现核间数据同步：弱内存模型下写序可能被重排，且 Data Cache 与 Global Memory 不保证一致，通过asc_threadfence()约束 key 写入与state置位的顺序，并用volatile强制从 Global Memory 读取其他核写入的数据，确保读线程轮询state后总能拿到完整、有效的 key。

将三者（原子操作 + 内存栅栏 + volatile 轮询）组合使用，即可在昇腾 AI 处理器上写出正确、高效、无锁的并发数据结构算子。这一套方法论不仅适用于哈希表，也可推广到计数、栈、队列、链表等任意需要多线程并发访问共享内存的场景。感兴趣的读者可继续阅读 asc_atomic_cas API 文档 与 asc_threadfence API 文档 了解接口的完整约束，或直接查看样例源码 insert_hash_table.asc 深入研读实现细节。

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend