人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载本篇技术指南以 asc-devkit 仓库中examples/02_simd_c_api/00_introduction/00_quickstart的 HelloWorld 样例为研究对象完整讲解如何在 NPU 侧通过内核调用符直接调用 Ascend C 算子核函数、在核函数内使用printf打印输出并借助 CANN 的 ACLAscendCL运行时接口完成从「调用 → 执行 → 同步 → 资源回收」的完整闭环。读完本文你将掌握核函数直调Kernel Direct Call的最小可运行工程结构、hello_world.asc双端Device/ Host代码的组织方式、CMake 构建配置要点以及编译运行与结果验证的完整操作步骤。样例概述一次最小化的 NPU 侧核函数验证HelloWorld 是理解 Ascend C 算子开发的最短路径样例。与需要Tiling、Context、CopyIn/CopyOut等完整算子框架的复杂样例不同本样例只做一件事在 NPU 侧直接调用一个核函数Kernel Function并在核函数内通过printf打印Hello World!!!。根据 README_en.md 的描述该样例用于演示基于 Ascend C 的 HelloWorld 算子的核函数直调方法从 NPU 侧运行核NPU 上执行验证算子核函数的正确性展示核函数从调用到执行的整体流程。它位于02_simd_c_api的00_introduction分类下属于 SIMD向量计算C API 系列的入门样例。整个02_simd_c_api目录定位是基于 Ascend C 的简单样例通过 Ascend C 编程语言演示自定义函数实现并给出对应的 C_API 实现见 02_simd_c_api/README_en.md。样例列表目录名称功能描述hello_world_npu通过使用内核调用符完成算子核函数在 NPU 侧运行验证的基础流程核函数内通过printf打印输出结果支持的产品与运行前提该样例在 NPU 真机侧运行官方声明支持以下产品见 hello_world_npu/README_en.mdAscend 950PR / Ascend 950DTAtlas A3 训练系列产品 / Atlas A3 推理系列产品Atlas A2 训练系列产品 / Atlas A2 推理系列产品运行前需要确认当前环境满足以下前提已安装与硬件匹配的 CANN 开发套件包包括 CANN toolkit 包必需与 CANN ops 包部分样例编译依赖建议安装。具体安装方式可参考 docs/en/quick_start.md 中的 Download and Install CANN Package 章节。NPU 驱动与固件可用可通过npu-smi info验证设备是否正常枚举。编译工具链齐全根据 docs/en/quick_start.md仓库源码编译与样例编译运行依赖python 3.9.0、gcc/g 7.3.0两者版本需一致、cmake 3.16.0、pkg-config 0.29.0。若使用 CANN 官方 Docker 镜像或云端开发环境这些依赖通常已内置。目录结构一个最小可运行工程的骨架hello_world_npu ├── CMakeLists.txt // 构建工程文件 └── hello_world.asc // Ascend C 算子实现 调用样例Device 侧核函数 Host 侧入口整个样例只有两个文件却完整覆盖了 Ascend C 核函数直调所需的全部要素核函数定义、Host 侧调用入口、CMake 构建配置。这对初学者而言是非常理想的最小工程模板。核心代码逐段解析hello_world.aschello_world.asc 是样例的灵魂文件代码量不足 40 行却串联起了 Ascend C 核函数编程的两大核心环节。下面分段讲解。1. 头文件引入#include kernel_operator.h #include acl/acl.hkernel_operator.h是 Ascend C 编程的核心头文件位于仓库 include/kernel_operator.h聚合了基本算子开发所需的类、宏与内置函数如printf、GetBlockIdx、GetBlockNum等。acl/acl.h提供 Host 侧使用的 AscendCLACL运行时 API 声明对应仓库 tools/include/acl/acl.h。2. Device 侧核函数定义__global__ __vector__ void hello_world() { AscendC::printf([Block (%lu/%lu)]: Hello World!!!\n, AscendC::GetBlockIdx(), AscendC::GetBlockNum()); }这段代码包含三个关键点__global__限定符声明这是一个可由 Host 侧发起调用的核函数Kernel Function编译器会为其生成可在 NPU 上执行的设备代码。__vector__限定符声明该核函数面向向量SIMD计算场景对应本仓库 SIMD C API 系列。类似的还有矩阵Cube计算等场景的限定符。AscendC::printf与内置函数核函数在 NPU 上并行执行printf会将格式化字符串与参数输出到 Host 侧控制台AscendC::GetBlockIdx()返回当前执行的 Block 索引AscendC::GetBlockNum()返回本次调用启动的 Block 总数格式串中分别以%lu承接。从源码实现看printf是 Ascend C 的设备侧打印能力在仓库中由 impl/utils/debug/asc_simd_printf_impl.h、impl/utils/debug/asc_aicore_printf_impl.h 等头文件提供实现printf_impl系列模板函数支持可变参数模板并最终将格式化输出经设备调试通道回传 Host。它本质上等价于一套面向 NPU 的格式化输出基础设施常用于算子调试阶段。3. Host 侧ACL 运行时初始化与核函数调用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 侧逻辑呈现了 AscendCL 运行时调用的标准生命周期aclInit(nullptr)初始化 ACL 运行时参数为nullptr即使用默认配置。aclrtSetDevice(deviceId)绑定 0 号 NPU 设备后续的流与内核执行都发生在这台设备上。aclrtCreateStream(stream)创建执行流Stream。AscendC 的核函数调用必须挂载在某个流上执行。核函数调用hello_worldnumBlocks, 0, stream()这是本样例的核心语法—— 内核调用符。三个参数的含义分别是第一个参数numBlocks 8启动的Block 数量即核函数在 NPU 上并行实例化的份数第二个参数0预留参数本样例未使用置 0第三个参数stream核函数提交到的ACL 流。aclrtSynchronizeStream(stream)阻塞等待流中所有任务即核函数执行完成保证后续回收资源时设备侧任务已结束同时确保printf输出已全部回传。aclrtDestroyStream(stream)/aclrtResetDevice(deviceId)/aclFinalize()依次销毁流、复位设备、终止 ACL 运行时完成资源回收。4. 执行流程全景将上述代码串起来核函数直调的整体流程为aclInit → aclrtSetDevice → aclrtCreateStream → hello_world8, 0, stream 提交 8 个 Block 到流 → aclrtSynchronizeStream 等待执行完成 → aclrtDestroyStream → aclrtResetDevice → aclFinalize每个 Block 独立执行hello_world核函数GetBlockIdx()输出各自的 Block 编号GetBlockNum()输出总数 8。CMake 构建配置解析CMakeLists.txtCMakeLists.txt 展示了如何用 CMake 驱动 Ascend C 编译器构建 NPU 可执行程序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)引入 CANN 为 Ascend C 语言提供的 CMake 包其查找逻辑与配置定义位于仓库 cmake/asc/asc_modules/FindASC.cmake编译器判定与信息收集则见 cmake/asc/asc_modules/CMakeDetermineASCCompiler.cmake。project(kernel_samples LANGUAGES ASC CXX)显式声明工程使用ASCAscend C 语言与CXX两种语言CMake 会据此选择 ASC 编译器链处理.asc源文件。add_executable(demo hello_world.asc)将.asc源文件编译链接为名为demo的可执行程序。target_compile_options(... --npu-archdav-2201)通过生成器表达式仅对 ASC 语言源文件传入--npu-arch选项指定目标 NPU 的架构代号arch。仓库注释明确提示需根据实际部署的 NPU 硬件架构选择对应的npu-arch参数。例如dav-2201对应 Atlas A2 系列dav-3510对应 Ascend 950 系列注释中给出了dav-3510的示例。这一参数与「支持的产品」列表Ascend 950PR/950DT、Atlas A2/A3 系列直接对应编译前务必按实际硬件核对。编译运行与结果验证第一步配置环境变量根据当前环境上 CANN 开发套件包的安装方式详见 docs/en/quick_start.md 的 Environment Variable Configuration 章节选择对应的环境变量配置命令默认路径root 用户安装 CANN 软件包source /usr/local/Ascend/cann/set_env.sh默认路径非 root 用户安装 CANN 软件包source $HOME/Ascend/cann/set_env.sh指定安装路径install_pathsource ${install_path}/cann/set_env.sh提示使用 CANN 官方 Docker 镜像或云端开发环境时环境变量通常已自动配置可跳过此步。第二步构建并执行在样例根目录hello_world_npu下执行mkdir -p build cd build; # 创建并进入 build 目录 cmake ..; make -j; # 配置并编译工程 ./demo # 执行样例cmake会根据CMakeLists.txt中的--npu-arch配置调用对应的 ASC 编译器链将hello_world.asc编译成同时包含 Device 侧核函数与 Host 侧入口的可执行文件demo./demo启动后Host 侧完成 ACL 初始化并把 8 个 Block 的核函数提交到 NPU 上执行。第三步验证执行结果执行成功时控制台输出如下共 8 个 Block编号 07[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!!!每一行由核函数内AscendC::printf打印[Block (当前Block编号/Block总数)]其中编号来自GetBlockIdx()、总数来自GetBlockNum()。看到 8 行输出即代表核函数被成功提交并在 NPU 侧执行8 个 Block 均正常并行运行设备侧printf的调试输出通道工作正常整条「调用 → 执行 → 同步 → 回收」链路畅通。这标志着一个最小化的 Ascend C 算子核函数直调流程验证完成。延伸从 HelloWorld 走向完整算子开发本样例虽小却是后续所有 Ascend C 样例的公共底座。在 02_simd_c_api 系列中00_introduction之后的目录在此基础上逐步叠加能力01_utilities引入printf、assert、DumpTensor等调试工具类 API支撑算子调试阶段02_features演示使用 Framework、C_API 等特性实现样例03_c_api系统介绍数据搬移Data Movement与计算Compute类 C API 的用法05_best_practices基于 Ascend C C API 的最佳实践样例如使用 SSBuffer 通信实现融合计算。若希望深入学习 C API 体系可以继续阅读 include/c_api/asc_simd.h——它是 SIMD 场景 C API 的聚合头文件按架构条件引入asc_cube.h、asc_memory_vector.h、asc_reg_vector.h等子模块帮助理解核函数内部可用的向量搬移与计算接口层级。结语HelloWorld 样例用最精简的代码量完整演示了 Ascend C 核函数直调的两条主线Device 侧如何用__global__ __vector__与printf定义可打印输出的核函数以及Host 侧如何用 ACL 运行时 API 与 调用符把核函数投递到 NPU 并回收资源。配合 CMakeLists.txt 中--npu-arch的架构选择、docs/en/quick_start.md 中的环境准备开发者可以在 10 分钟内跑通第一个 NPU 侧算子程序为后续数据搬移、向量计算等完整算子开发打下坚实基础。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐CANN HelloWorld NPU 直调样例实战基于 内核调用符的 Ascend C 核函数快速验证指南CANN HelloWorld NPU 直调样例实战基于 内核调用符的 Ascend C 核函数快速验证指南 本篇技术指南围绕 cann sample示例工程CANNCANN Samples HelloWorld 算子核函数直调指南使用 内核调用符在 NPU 侧验证 Ascend C 算子CANN Samples HelloWorld 算子核函数直调指南使用 内核调用符在 NPU 侧验证 Ascend C 算子 导读 本篇指南围绕 CA示例工程CANNCANN Ascend C 入门实战基于 内核调用符的 HelloWorld 算子 NPU 直调与 printf 验证CANN Ascend C 入门实战基于 内核调用符的 HelloWorld 算子 NPU 直调与 printf 验证 导读 本文以 cann sam示例工程CANN上一篇GameEngineFromScratch高性能计算ISPC并行编程实战指南下一篇gh_mirrors/server117/server REST API全解析开发必备手册创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考