尧图网页设计 
人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载导读本文围绕 CANN asc-devkit 仓库中examples/02_simd_c_api/00_introduction/01_add下的三个 Add 算子样例展开系统讲解如何使用昇腾 SIMD C_APIc_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 编程模型的入门阶梯。算子规格与支持产品三个样例的算子规格完全一致如下表所示项目内容算子类型OpTypeAdd输入 xshape2048*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-archdav-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 BufferUB空间数组asc_copy_gm2ub(dst, src, size)以字节数为单位将数据从全局搬入 UBasc_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_gm2ubx 与 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.3fgolden 直接取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_customnumBlocks, 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 streamc_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-archdav-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_addasc_sync_notify/asc_sync_wait事件同步高搬入—计算—搬出三流水线深度重叠追求极致性能、UB 复用受限的高阶场景建议初学者先对照同步样例理解搬运—计算—搬出三步模型再依次进阶到异步与精细异步写法当算子存在多组独立数据分块时事件级同步能带来最明显的性能收益但务必核对每个notify/wait的流水线方向与事件号避免引入数据竞争。更多 C_API 接口与算子开发规范可参考 docs/en/quick_start.md 及 include/c_api/asc_simd.h 接口声明。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐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上一篇aws-amplify数据迁移策略从传统数据库到云原生存储的无缝过渡下一篇VSCode-GitLens访问错误AccessDeniedError与权限处理创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 昇腾 C_API 编写 Add 算子实战：同步、异步与精细异步三种调用范式解析（基于 CANN asc-devkit）
人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载导读本文围绕 CANN asc-devkit 仓库中examples/02_simd_c_api/00_introduction/01_add下的三个 Add 算子样例展开系统讲解如何使用昇腾 SIMD C_APIc_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 编程模型的入门阶梯。算子规格与支持产品三个样例的算子规格完全一致如下表所示项目内容算子类型OpTypeAdd输入 xshape2048*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-archdav-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 BufferUB空间数组asc_copy_gm2ub(dst, src, size)以字节数为单位将数据从全局搬入 UBasc_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_gm2ubx 与 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.3fgolden 直接取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_customnumBlocks, 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 streamc_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-archdav-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_addasc_sync_notify/asc_sync_wait事件同步高搬入—计算—搬出三流水线深度重叠追求极致性能、UB 复用受限的高阶场景建议初学者先对照同步样例理解搬运—计算—搬出三步模型再依次进阶到异步与精细异步写法当算子存在多组独立数据分块时事件级同步能带来最明显的性能收益但务必核对每个notify/wait的流水线方向与事件号避免引入数据竞争。更多 C_API 接口与算子开发规范可参考 docs/en/quick_start.md 及 include/c_api/asc_simd.h 接口声明。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐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上一篇aws-amplify数据迁移策略从传统数据库到云原生存储的无缝过渡下一篇VSCode-GitLens访问错误AccessDeniedError与权限处理创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考