- 人工智能
- 深度学习
- 算子库
- CANN
- Ascend
【免费下载链接】asc-devkit
本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。
导读
本文围绕 CANN asc-devkit 仓库中examples/02_simd_c_api/00_introduction/01_add下的三个 Add 算子样例展开，系统讲解如何使用昇腾 SIMD C_API（c_api/asc_simd.h）在单个.asc文件中同时实现 kernel 函数与 main 函数，并通过<<<>>>直调方式在主机侧拉起核函数。读完本文，你将掌握 C_API 的asc_copy_gm2ub/asc_copy_ub2gm数据搬运、asc_add向量计算、asc_sync全量同步以及asc_sync_notify/asc_sync_wait事件级流水同步三种编程范式，并能够独立完成编译、运行与精度校验。
样例总体介绍
仓库在 examples/02_simd_c_api/00_introduction/01_add 目录下提供了一个"基于 Ascend C 的 Add 算子<<<>>>直调方法"样例集，其核心特点是main 函数与 kernel 函数写在同一个 cpp（.asc）文件中，无需拆分 host/device 工程即可完成算子实现、调用与验证。样例集共包含三个子目录，分别演示三种不同的接口组合方式：
| 目录名称 | 功能描述 |
|---|---|
| c_api_async_add | 采用 C_API 接口实现 Add 算子，基于异步数据搬运与计算接口 |
| c_api_delicacy_async_add | 采用 C_API 接口实现 Add 算子，基于异步搬运、计算接口，并手动添加同步指令控制流水线依赖 |
| c_api_sync_add | 采用 C_API 接口实现 Add 算子，基于同步数据搬运与计算接口 |
三者计算逻辑完全一致，差异集中在核内流水线的同步策略上，非常适合作为理解昇腾 SIMD C_API 编程模型的入门阶梯。
算子规格与支持产品
三个样例的算子规格完全一致，如下表所示：
| 项目 | 内容 |
|---|---|
| 算子类型（OpType） | Add |
| 输入 x | shape2048*8（同步样例）/8*2048（两个异步样例），数据类型 float，格式 ND |
| 输入 y | 与 x 同 shape、同类型、同格式 |
| 输出 z | 与输入同 shape、同类型、同格式 |
| 核函数名称 | add_custom |
其中 shape 的书写顺序（2048*8与8*2048）不影响实际计算语义，因为核内均按"总长度 / block 数"进行分块；c_api_async_add中TOTAL_LENGTH = 8 * 2048 = 16384，与另外两个样例的元素总数完全一致。
算子功能即数学表达式z = x + y：将两个输入张量逐元素相加并返回结果。
支持产品范围：
- Atlas A3 训练系列产品 / Atlas A3 推理系列产品
- Atlas A2 训练系列产品 / Atlas A2 推理系列产品
构建配置方面，三个样例的 CMakeLists.txt 均要求 CMake 版本不低于 3.16，通过find_package(ASC REQUIRED)引入昇腾 C 语言扩展编译器，并将--npu-arch=dav-2201作为默认 NPU 架构编译选项；c_api_async_add的 CMakeLists 额外提供了CMAKE_ASC_ARCHITECTURES缓存变量，便于按实际部署的 NPU 硬件架构调整编译目标。
核内实现流程：三步式 Add 计算
C_API 编程模型中，设备侧数据不能直接被向量计算指令访问，输入必须先搬入片上存储（Local Memory / UB），计算完成后再搬回外部存储（Global Memory）。三个样例的核内实现均遵循以下三步：
- 第一步（搬入）：将输入
x、y从 Global Memory 搬运到 Local Memory，分别存入xLocal、yLocal； - 第二步（计算）：对
xLocal、yLocal执行逐元素加法，结果存入zLocal； - 第三步（搬出）：将
zLocal中的输出数据搬运回 Global Memory 的输出z。
在__vector__ __global__修饰的核函数入口处，首先调用asc_init()完成片上环境初始化，随后按block_idx计算当前核（AI Core 上运行的 block）负责的数据分片，再依次执行搬运—计算—搬运。
三种 C_API 实现范式源码解析
范式一：c_api_sync_add —— 同步搬运 + 显式asc_sync
同步样例的核函数源码位于 c_api_add.asc，核心代码段如下：
__vector__ __global__ __aicore__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z) { asc_init(); __ubuf__ float xLocal[TILE_LENGTH]; __ubuf__ float yLocal[TILE_LENGTH]; __ubuf__ float zLocal[TILE_LENGTH]; uint32_t blockLength = TILE_LENGTH * NUM_BLOCKS / block_num; asc_copy_gm2ub(xLocal, (x + block_idx * blockLength), blockLength * sizeof(float)); asc_sync(); asc_copy_gm2ub(yLocal, (y + block_idx * blockLength), blockLength * sizeof(float)); asc_sync(); asc_add(zLocal, xLocal, yLocal, blockLength); asc_sync(); asc_copy_ub2gm((z + block_idx * blockLength), zLocal, blockLength * sizeof(float)); asc_sync(); }要点解读：
__gm__表示 Global Memory 空间指针，__ubuf__表示片上 Unified Buffer（UB）空间数组；asc_copy_gm2ub(dst, src, size)以字节数为单位将数据从全局搬入 UB，asc_copy_ub2gm反向搬出；- 每次搬运/计算之后都紧跟
asc_sync()全量同步。asc_sync在 include/c_api/sync/sync.h 中声明，作用是等待流水线上所有已发出的任务执行完毕，保证后续指令能安全读取前序指令的产出数据； asc_add(zLocal, xLocal, yLocal, blockLength)为向量加法接口，声明于 include/c_api/vector_compute/compute/vector_arith.h，第四个参数count为参与计算的元素个数；blockLength = TILE_LENGTH * NUM_BLOCKS / block_num实现了数据在多核间的自适应切分：当实际启动的 block 数小于 8 时，每个 block 分到的数据量自动增大，保证总数据不遗漏。
这种"每步一同步"的写法最简单直观，代码可读性强、不易出错，但同步开销较大，流水线无法重叠执行。
范式二：c_api_async_add —— 异步搬运 + 批间同步
异步样例源码位于 c_api_add.asc，其核函数省去了__aicore__修饰，并把同步点从"每条指令之后"收敛为"阶段之间"：
__vector__ __global__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z) { asc_init(); constexpr uint32_t block_length = TOTAL_LENGTH / NUM_BLOCKS; __gm__ float* x_gm = x + block_idx * block_length; __gm__ float* y_gm = y + block_idx * block_length; __gm__ float* z_gm = z + block_idx * block_length; __ubuf__ float x_local[block_length]; __ubuf__ float y_local[block_length]; __ubuf__ float z_local[block_length]; asc_copy_gm2ub(x_local, x_gm, block_length * sizeof(float)); asc_copy_gm2ub(y_local, y_gm, block_length * sizeof(float)); asc_sync(); asc_add(z_local, x_local, y_local, block_length); asc_sync(); asc_copy_ub2gm(z_gm, z_local, block_length * sizeof(float)); asc_sync(); }要点解读：
- 两次
asc_copy_gm2ub（x 与 y 的搬入）之间不再插入同步。由于二者写入的是不同的 UB 缓冲区（x_local与y_local），彼此无数据依赖，可异步并发下发，由同一流水级顺序执行即可保证正确性； - 同步点仅保留三处：全部搬入完成之后、向量计算完成之后、搬出完成之后，分别对应"搬运→计算""计算→搬出"两个数据依赖边界；
- 这种写法比范式一少了一半同步指令，让数据搬运与后续指令的发射可以更早进行，属于"异步接口 + 粗粒度同步"的折中方案。
范式三：c_api_delicacy_async_add —— 异步接口 + 事件级手动同步
精细异步样例源码位于 c_api_add.asc，它不再依赖asc_sync()全量同步，而是把数据进一步切成C_API_TILE_NUM = 8个 tile，在循环中通过asc_sync_notify/asc_sync_wait精确控制MTE2（搬入）、Vector（计算）、MTE3（搬出）三条流水线之间的事件依赖，实现流水线级并行：
constexpr uint32_t C_API_ONE_BLOCK_SIZE = 32; constexpr uint32_t C_API_ONE_REPEAT_BYTE_SIZE = 256; constexpr uint32_t C_API_TOTAL_LENGTH = 16384; constexpr uint32_t C_API_TILE_NUM = 8; constexpr uint32_t C_API_TILE_LENGTH = 256; __vector__ __global__ __aicore__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z) { asc_init(); uint32_t blockLength = C_API_TOTAL_LENGTH / block_num; uint32_t tileLength = blockLength / C_API_TILE_NUM; __gm__ float* xGm = x + block_idx * blockLength; __gm__ float* yGm = y + block_idx * blockLength; __gm__ float* zGm = z + block_idx * blockLength; __ubuf__ float xLocal[C_API_TILE_LENGTH]; __ubuf__ float yLocal[C_API_TILE_LENGTH]; __ubuf__ float zLocal[C_API_TILE_LENGTH]; uint16_t burst_len = tileLength; for (uint32_t i = 0; i < C_API_TILE_NUM; i++) { if (i != 0) { asc_sync_wait(PIPE_V, PIPE_MTE2, EVENT_ID0); } burst_len = tileLength * sizeof(float) / C_API_ONE_BLOCK_SIZE; asc_copy_gm2ub(xLocal, xGm + i * tileLength, 1, burst_len, 0, 0); asc_copy_gm2ub(yLocal, yGm + i * tileLength, 1, burst_len, 0, 0); asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); if (i != 0) { asc_sync_wait(PIPE_MTE3, PIPE_V, EVENT_ID0); } asc_add(zLocal, xLocal, yLocal, tileLength * sizeof(float) / C_API_ONE_REPEAT_BYTE_SIZE, 1, 1, 1, 8, 8, 8); if (i != (C_API_TILE_NUM - 1)) { asc_sync_notify(PIPE_V, PIPE_MTE2, EVENT_ID0); } asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); asc_copy_ub2gm(zGm + i * tileLength, zLocal, 1, burst_len, 0, 0); if (i != (C_API_TILE_NUM - 1)) { asc_sync_notify(PIPE_MTE3, PIPE_V, EVENT_ID0); } } }要点解读：
asc_sync_notify(pipe, tpipe, id)与asc_sync_wait(pipe, tpipe, id)是事件级同步原语（声明于 include/c_api/sync/sync.h），其中pipe/tpipe为流水线类型（PIPE_V向量计算、PIPE_MTE2数据搬入、PIPE_MTE3数据搬出），id为事件号（样例统一复用EVENT_ID0）。语义为：notify方完成该流水级任务后发送事件，wait方在对应流水级等待该事件；- 循环体内形成了三阶段流水：第
i次迭代搬入第i个 tile → 计算第i个 tile → 搬出第i个 tile；同时通过asc_sync_notify(PIPE_V, PIPE_MTE2, ...)告知 Vector 计算已完成、可以搬入下一个 tile，实现前一个 tile 的计算与后一个 tile 的搬入重叠； asc_copy_gm2ub使用了扩展签名：asc_copy_gm2ub(dst, src, 1, burst_len, 0, 0)，其中burst_len = tileLength * sizeof(float) / 32表示按 32 字节 block 为单位的搬运突发长度，这是针对连续内存块的高效搬运写法；asc_add使用了扩展签名：asc_add(dst, src0, src1, repeat, 1, 1, 1, 8, 8, 8)，repeat = tileLength * sizeof(float) / 256表示重复执行次数（256 字节/次），后续参数为迭代间隔与重复间隔等硬件调度参数。结合 include/c_api/vector_compute/compute/vector_arith.h 中asc_add的多个重载可见，C_API 同时提供"按元素数"与"按 repeat 次数"两种粒度的计算接口；- 该样例的主机侧输入使用常量初始化（
x全为1.2f、y全为2.3f），golden 直接取1.2 + 2.3 = 3.5f，便于精确比对。
这是三种写法中性能上限最高的方案，代价是同步逻辑复杂、依赖关系需要开发者手工维护。
主机侧调用链：ACL Runtime 直调<<<>>>
三个样例的调用侧实现高度一致，均封装在kernel_add函数中，展示了一条完整的"host 侧数据准备 → 核函数直调 → 结果回拷 → 资源释放"链路，以同步样例为例：
aclInit(nullptr); aclrtSetDevice(deviceId); aclrtCreateStream(&stream); aclrtMallocHost((void**)(&zHost), totalByteSize); aclrtMalloc((void**)&xDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc((void**)&yDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc((void**)&zDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMemcpy((uint8_t*)xDevice, totalByteSize, xHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE); aclrtMemcpy((uint8_t*)yDevice, totalByteSize, yHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE); add_custom<<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice); aclrtSynchronizeStream(stream); aclrtMemcpy(zHost, totalByteSize, (uint8_t*)zDevice, totalByteSize, ACL_MEMCPY_DEVICE_TO_HOST);关键点：
- 依次执行
aclInit（初始化 ACL）、aclrtSetDevice（指定 device）、aclrtCreateStream（创建执行流）； - 设备侧内存通过
aclrtMalloc以ACL_MEM_MALLOC_HUGE_FIRST策略申请，主机侧通过aclrtMallocHost申请； - 核函数调用采用 CUDA 风格的
<<<numBlocks, 0, stream>>>直调语法：第一个参数为 grid 维 block 数（样例取NUM_BLOCKS = 8，与核内block_num对应），第二个参数为局部 block 维大小（取 0），第三个参数为 ACL stream； c_api_sync_add与c_api_delicacy_async_add在调用后使用aclrtSynchronizeStream(stream)等待流内任务完成，c_api_async_add使用aclrtSynchronizeDevice()等待设备整体同步——这是异步样例与另外两个样例在主机侧的细微差异；- 结果通过
aclrtMemcpy以ACL_MEMCPY_DEVICE_TO_HOST模式回拷到主机，随后按逆序释放设备内存、销毁 stream、aclrtResetDevice并aclFinalize收尾。
编译与运行
在任意一个样例目录的根目录下，按以下步骤即可完成编译与运行。
1. 配置环境变量：根据当前环境上 CANN 开发套件包的安装方式，选择对应的命令：
# 默认安装路径，root 用户安装 source /usr/local/Ascend/cann/set_env.sh # 默认安装路径，非 root 用户安装 source $HOME/Ascend/cann/set_env.sh # 自定义安装路径 install_path source ${install_path}/cann/set_env.sh2. 编译并执行：
mkdir -p build && cd build # 创建并进入 build 目录 cmake ..; make -j # 编译工程 ./c_api_add_example # 执行样例编译过程中，find_package(ASC REQUIRED)会通过 CANN 提供的 ASC 语言扩展编译工具链将.asc文件中的__vector__ __global__核函数与主机代码编译为可执行文件；--npu-arch=dav-2201指定目标 NPU 架构，实际部署时应按硬件修改该参数（详见各样例的 CMakeLists.txt）。
3. 预期输出：执行成功后打印精度比对结果，表示算子计算结果与 golden 完全一致：
[Success] Case accuracy is verification passed.精度校验逻辑
三个样例均内置了VerifyResult/verify_result函数完成端到端校验，流程为：
- 打印输出张量与 golden 张量的前 20 个元素（多于 20 个时以
...截断），便于人工观察； - 通过
std::equal逐元素比对设备计算结果与主机 golden； - 一致则打印
[Success] Case accuracy is verification passed.并返回 0，否则打印[Failed] Case accuracy is verification failed!并返回 1。
golden 的构造方式：同步与异步样例在主机侧初始化x[i] = i * 0.1f; y[i] = i * 0.1f;，golden 取x[i] + y[i]；精细异步样例则用常量初始化并直接以valueX + valueY构造 golden。整个 main 函数采用"构造数据 → 调用kernel_add上板计算 → 构造 golden → 精度比对 → 返回退出码"的结构，可直接作为回归用例复用。
三种范式的选型建议
| 实现范式 | 同步策略 | 代码复杂度 | 流水线重叠 | 适用场景 |
|---|---|---|---|---|
| c_api_sync_add | 每条指令后asc_sync() | 低 | 无 | 入门学习、验证算法正确性 |
| c_api_async_add | 阶段间asc_sync() | 中 | 搬运指令间部分重叠 | 常规算子开发，兼顾可读性与性能 |
| c_api_delicacy_async_add | asc_sync_notify/asc_sync_wait事件同步 | 高 | 搬入—计算—搬出三流水线深度重叠 | 追求极致性能、UB 复用受限的高阶场景 |
建议初学者先对照同步样例理解"搬运—计算—搬出"三步模型，再依次进阶到异步与精细异步写法；当算子存在多组独立数据分块时，事件级同步能带来最明显的性能收益，但务必核对每个notify/wait的流水线方向与事件号，避免引入数据竞争。更多 C_API 接口与算子开发规范可参考 docs/en/quick_start.md 及 include/c_api/asc_simd.h 接口声明。
- 人工智能
- 深度学习
- 算子库
- CANN
- Ascend
【免费下载链接】asc-devkit
本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。
相关推荐
CANN 昇腾 C_API 直调 Add 算子实战：同步、异步与精细流水线同步三种实现方式解析
CANN 昇腾 C_API 直调 Add 算子实战：同步、异步与精细流水线同步三种实现方式解析 导读 本文基于 CANN cann samples 仓库 Sam
示例工程CANNCANN cann-samples 实战：基于 Ascend C C_API 的 Add 算子核函数直调（异步、同步、精细同步三场景详解）
CANN cann samples 实战：基于 Ascend C C_API 的 Add 算子核函数直调（异步、同步、精细同步三场景详解） 导读 本文围绕 ca
示例工程CANN精雕异步：cann-samples 中基于 SIMD C_API 与显式流水同步的 Add 算子实现详解
精雕异步：cann samples 中基于 SIMD C_API 与显式流水同步的 Add 算子实现详解 本文以 cann samples 仓库中 c_api_
示例工程CANN
创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
