CANN Samples HelloWorld 算子核函数直调指南使用内核调用符在 NPU 侧验证 Ascend C 算子【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples导读本篇指南围绕 CANN 开源样例仓库 cann-samples 中Samples/0_Introduction/02_simd_c_api/00_quickstart目录下的 HelloWorld 快速入门样例展开讲解基于 Ascend C 的算子核函数Kernel Function直调方法开发者如何在单个.asc源文件中同时编写__global__ __vector__核函数与 Host 侧main函数并通过内核调用符将核函数直接下发到 NPU 上运行验证。读完本篇你将掌握 Ascend C 核函数直调的最小可运行范式、ACL Runtime 初始化与资源生命周期管理、Ascend C 内置printf的片上打印用法以及基于 CMake Bisheng 编译器完成 NPU 侧编译、执行与结果校验的完整流程可作为后续 Add、MatMul 等算子样例开发的第一步。样例定位C 语言风格接口的入门第一站在 Samples/0_Introduction/02_simd_c_api 的样例体系中本样例承担的是C 语言风格接口C_API入门的角色。该目录下的样例按难度递进组织样例目录功能定位00_quickstart基于 Ascend C 的 HelloWorld 算子核函数直调方法从 NPU 侧运行核验证算子核函数展示核函数从调用到执行的整体流程01_add基于 Ascend C 的 Add 自定义 Vector 算子核函数直调实现两个输入张量的逐元素相加支持 main 函数和 kernel 函数在同一个 cpp 文件中实现04_reg_base_add_compute基于 Ascend C 的 Add 算子核函数直调方法RegBase 场景展示片上存储和寄存器层级的向量计算流程与更偏重概念讲解的 npu_execution纯文档、无代码拆解算子从 PyTorch 调用到芯片执行的完整链路以及 vector_add_c_api最简向量加法实现相比本样例聚焦于核函数直调这一最基础的执行模型不经过算子工程化编译不涉及张量搬入搬出只验证一个核函数能否被正确编译、下发并在 NPU 上执行是打通编译、运行、调试全链路的最小闭环。核心概念核函数直调与内核调用符什么是核函数直调算子核函数直调Kernel Direct Call是 Ascend C 提供的一种轻量级算子开发与验证方式开发者直接用__global__修饰符声明核函数在 Host 侧通过内核调用符完成调用无需像传统算子开发那样经历算子工程创建、算子原型定义OpType、算子信息库注册等繁琐步骤。它尤其适合快速验证算子算法逻辑的正确性进行算子性能的初步摸底与调优迭代作为学习 Ascend C 编程模型的入门路径。内核调用符的参数语义在 hello_world.asc 中核函数的定义与调用分别如下__global__ __vector__ void hello_world() { AscendC::printf([Block (%lu/%lu)]: Hello World!!!\n, AscendC::GetBlockIdx(), AscendC::GetBlockNum()); } // Host 侧 main 函数中调用 constexpr uint32_t numBlocks 8; hello_worldnumBlocks, 0, stream();内核调用符携带三个参数参数位置本样例取值含义第一个参数numBlocks 8本次调用启动的 AI Core 核数Block 数量即核函数将被多少个计算核并行执行第二个参数0核函数运行时的本地内存大小参数本样例未使用传 0第三个参数stream核函数执行所在的 ACL Runtime 流Stream核函数异步排队到该流上执行三个参数之间用逗号分隔。这里需要注意一个易混淆点该调用符与 CUDA 的grid, block网格/线程块模型在语义上并不相同Ascend C 直调场景下第一个参数直接对应 NPU 侧的 BlockAI Core 实例数量。源文件逐行解读核函数与 Host 主函数整个样例仅由一个源文件 hello_world.asc 构成完整展示了核函数 Host 主函数同文件的直调范式结构如下├── hello_world_npu │ ├── CMakeLists.txt // 编译工程文件 │ └── hello_world.asc // Ascend C算子实现 调用样例头文件引入#include kernel_operator.h #include acl/acl.hkernel_operator.h是 Ascend C 编程的核心头文件提供核函数声明、AscendC::printf、AscendC::GetBlockIdx()/AscendC::GetBlockNum()等内核侧基础设施acl/acl.h是 ACLAscend Computing LanguageRuntime 的 Host 侧接口头文件提供aclInit、aclrtSetDevice、aclrtCreateStream、aclrtSynchronizeStream等设备管理与流管理接口。核函数定义__global__ __vector__ void hello_world() { AscendC::printf([Block (%lu/%lu)]: Hello World!!!\n, AscendC::GetBlockIdx(), AscendC::GetBlockNum()); }__global__核函数修饰符声明该函数为可被 Host 侧调用的设备端核函数__vector__声明该核函数运行在 Vector Core向量计算核上对应 Vector 算力场景AscendC::GetBlockIdx()与AscendC::GetBlockNum()分别返回当前核函数执行实例的 Block 编号与本次调用启动的 Block 总数用于在并行执行时区分不同核的上下文AscendC::printfAscend C 提供的核内打印接口支持格式化字符串可将输出从 NPU 侧打印到 Host 侧终端是核函数调试最直接的观测手段。Host 主函数与 ACL Runtime 生命周期int32_t main(int argc, char const* argv[]) { aclInit(nullptr); int32_t deviceId 0; aclrtSetDevice(deviceId); aclrtStream stream nullptr; aclrtCreateStream(stream); constexpr uint32_t numBlocks 8; hello_worldnumBlocks, 0, stream(); aclrtSynchronizeStream(stream); aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0; }Host 侧代码完整覆盖了 ACL Runtime 的标准生命周期这一段是任何后续直调样例Add、MatMul 等都会复用的骨架阶段接口作用初始化aclInit(nullptr)初始化 ACL Runtime 运行环境nullptr表示使用默认配置设备绑定aclrtSetDevice(deviceId)将当前进程绑定到指定 device本样例为设备 0创建流aclrtCreateStream(stream)创建执行流核函数异步任务将挂载到该流下发任务hello_worldnumBlocks, 0, stream()将核函数异步下发到流的执行队列同步等待aclrtSynchronizeStream(stream)阻塞等待流上所有任务执行完毕保证在销毁资源前核函数已真正跑完销毁流aclrtDestroyStream(stream)释放流资源解绑设备aclrtResetDevice(deviceId)释放当前进程对设备的占用去初始化aclFinalize()销毁 ACL Runtime 环境值得特别强调的是aclrtSynchronizeStream的必要性由于调用是异步的若缺少同步步骤就直接退出或销毁流核函数的打印输出可能尚未完成导致观察不到执行结果。编译工程解读ASC 语言与 npu-arch 架构选择样例的编译由 CMakeLists.txt 驱动核心内容如下cmake_minimum_required(VERSION 3.16) find_package(ASC REQUIRED) project(kernel_samples LANGUAGES ASC CXX) add_executable(demo hello_world.asc ) target_compile_options(demo PRIVATE $$COMPILE_LANGUAGE:ASC:--npu-archdav-2201 # $$COMPILE_LANGUAGE:ASC:--npu-archdav-3510 )其中关键点find_package(ASC REQUIRED)由仓库根目录 cmake/ascend.cmake 提供支持该文件负责定位 CANN 开发套件优先读取环境变量ASCEND_HOME_PATH其次按 root 用户/usr/local/Ascend/ascend-toolkit/latest、非 root 用户$HOME/Ascend/ascend-toolkit/latest等默认路径探测并将CMAKE_C_COMPILER、CMAKE_CXX_COMPILER、CMAKE_LINKER统一指向 CANN 自带的 Bisheng 编译器${ASCEND_DIR}/${SYSTEM_PREFIX}/ccec_compiler/bin/bisheng从而让 CMake 具备编译.asc源文件的能力project(kernel_samples LANGUAGES ASC CXX)显式声明工程同时使用ASCAscend C 语言与CXX两种语言这是.asc文件能够参与编译的前提--npu-archdav-2201通过target_compile_options指定目标 NPU 硬件架构这是必须与实际部署环境匹配的关键参数。从仓库中 01_add 各样例的 CMakeLists.txt 可以看到部分样例使用${CMAKE_ASC_ARCHITECTURES}变量自动传入架构而本样例采用注释切换的方式默认dav-2201若部署在 Ascend 950 系列上则取消注释改用dav-3510。编译前请务必确认当前环境的芯片型号选择对应的npu-arch参数否则编译产物无法在目标 NPU 上运行。编译与运行完整操作步骤第一步配置 CANN 环境变量根据当前环境上 CANN 开发套件包的安装方式选择对应的环境变量配置命令# 默认路径root 用户安装 CANN 软件包 source /usr/local/Ascend/cann/set_env.sh # 默认路径非 root 用户安装 CANN 软件包 source $HOME/Ascend/cann/set_env.sh # 指定路径 install_path 安装 CANN 软件包 source ${install_path}/cann/set_env.shset_env.sh会完成ASCEND_HOME_PATH等关键环境变量的注入这也是 cmake/ascend.cmake 中定位工具链所依赖的输入。第二步编译与执行在样例根目录Samples/0_Introduction/02_simd_c_api/00_quickstart/hello_world_npu下执行mkdir -p build cd build; # 创建并进入 build 目录 cmake ..; make -j; # 编译工程 ./demo # 执行样例编译成功后生成可执行文件demo运行后输出如下说明核函数已在 NPU 侧成功执行[Block (0/8)]: Hello World!!! [Block (1/8)]: Hello World!!! [Block (2/8)]: Hello World!!! [Block (3/8)]: Hello World!!! [Block (4/8)]: Hello World!!! [Block (5/8)]: Hello World!!! [Block (6/8)]: Hello World!!! [Block (7/8)]: Hello World!!!结果解读验证了什么8 行输出与numBlocks 8一一对应每行开头的Block (i/8)正是核函数内部AscendC::GetBlockIdx()与AscendC::GetBlockNum()的实时取值。这意味着核函数被成功编译为可执行目标编译链路正确8 个 Block 均被调度到 NPU 侧并行执行调度与下发链路正确核内printf的输出成功回传到 Host 终端片上打印与数据回传链路正确由于输出完整无截断说明aclrtSynchronizeStream的同步等待逻辑生效。这四点共同构成了核函数从调用到执行的整体流程的完整验证。支持的产品与适用前提根据 hello_world_npu 样例说明本样例支持的产品包括Ascend 950PR / Ascend 950DTAtlas A3 训练系列产品 / Atlas A3 推理系列产品Atlas A2 训练系列产品 / Atlas A2 推理系列产品实际使用时需满足以下前提目标机器已安装与硬件匹配的 CANN 开发套件建议参考 CANN 配套的安装指引完成环境准备编译时--npu-arch参数与硬件架构匹配本样例默认dav-2201Ascend 950 平台需切换为dav-3510具备可用的 NPU 设备本样例固定使用 device 0。进阶方向从 HelloWorld 走向真实算子完成本样例后建议沿 02_simd_c_api 的目录顺序继续深入01_add/c_api_sync_add在直调骨架中加入真实的张量输入输出学习如何为核函数分配、搬运和回收 Device 内存实现两个输入张量的逐元素相加01_add/c_api_async_add在同步版本基础上引入异步执行与更精细的流管理理解 Host/Device 异步协作模型04_reg_base_add_compute/c_api_simd_add将计算下沉到寄存器层级RegBase 场景展示片上存储和寄存器层级的向量计算流程。这些样例都复用本篇讲解的直调与 ACL Runtime 生命周期骨架因此熟练掌握 HelloWorld 样例相当于拿到了整个 C_API 算子开发系列的第一把钥匙。小结本文以 00_quickstart 目录 的 HelloWorld 样例为核心完整梳理了 Ascend C 核函数直调的最小工程范式__global__ __vector__核函数定义、内核调用符的语义、核内printf与 Block 上下文接口的使用、ACL Runtime 的初始化—下发—同步—销毁全生命周期以及 CMake Bisheng 的编译配置与npu-arch架构选择。核函数直调是 Ascend C 算子开发的起点能力理解并跑通这一最小闭环后即可在此基础上叠加张量操作、数据搬运与性能调优向更复杂的 Vector/Cube 算子开发稳步推进。【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考