CANN cann-samples 入门实战基于 Ascend C 与内核调用符的 HelloWorld 样例全解析【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples导读本文围绕 cann-samples 仓库中 hello_world 样例展开它位于Samples/0_Introduction/01_simd_cpp_api/00_quickstart/目录下是 C 语法风格simd_cpp_api编程接口的起点样例。该样例演示了如何用 Ascend C 编写核函数、通过内核调用符将其直接下发到 NPU 侧运行验证并借助printf在 AI Core 上打印输出。读完本文你将掌握 Ascend C 核函数直调的最小完整工程结构源码 CMake 构建 运行/调试三种模式理解dav-2201/dav-3510架构与产品、CANN 版本的对应关系并能独立把该样例在 Atlas A2/A3 系列或 Ascend 950 系列产品上编译运行起来。1. 样例定位Ascend C 核函数直调的第一步在 CANN 的编程体系中除了通过算子库如 aclnn间接执行算子外还有一种更底层的开发方式——核函数直调Kernel Direct Call开发者直接用 Ascend C 编写设备侧核函数再在 Host 侧通过类似 CUDA 的语法将核函数下发到 NPU 执行。这种方式适合算子原型验证、性能调优和教学演示是理解 NPU 上编译—加载—执行—同步全流程的最短路径。00_quickstart/hello_world正是这条路径上的Hello World核函数内部仅做一件事printf(Hello World!!!\n)通过8, 0, stream把核函数派发到 8 个 AIVAI VectorCore 上并行执行最终在 Host 侧同步流并销毁资源完成一次完整的 NPU 侧运行验证。该样例属于 01_simd_cpp_apiSIMD 向量编程接口、C 语法风格下的快速入门模块与其并列的还有 SIMT 线程编程风格版本 03_simt_api/00_quickstart/hello_world_simt读者可对照学习两种核函数并行模型的启动差异。2. 支持的产品与 CANN 软件版本运行本样例前需确认当前 NPU 产品与 CANN 版本满足要求。README 给出的支持矩阵如下产品CANN 软件版本Ascend 950PR/Ascend 950DT CANN 9.1.0Atlas A3 训练系列产品/Atlas A3 推理系列产品 CANN 9.0.0Atlas A2 训练系列产品/Atlas A2 推理系列产品 CANN 9.0.0该矩阵与编译选项中CMAKE_ASC_ARCHITECTURES的两个取值一一对应详见第 5 节dav-2201面向 Atlas A2/A3 训练与推理系列产品dav-3510面向 Ascend 950PR/Ascend 950DT。本仓库构建体系中顶层 CMake 也会通过 sample_common.cmake 中的cann_sample_check_arch宏按NPU_ARCH对样例做架构匹配性检查不支持的样例会被跳过Skip sample ...: NPU_ARCH... is not supported。3. 样例目录结构├── hello_world │ ├── CMakeLists.txt // 编译工程文件 │ ├── hello_world.asc // Ascend C样例实现 调用样例 │ └── README.md // 样例说明文档三个文件职责清晰.asc是 Ascend C 源码既包含设备侧核函数也包含 Host 侧调用入口mainCMakeLists.txt负责把.asc编译链接为可执行文件demoREADME.md为使用说明。整个仓库中所有算子类样例均遵循这一CMake .asc README的组织惯例。4. 源码逐段解析从核函数定义到 NPU 执行完整的样例实现位于 hello_world.asc核心代码只有约 20 行却覆盖了 Ascend C 直调的全部关键环节。4.1 头文件与核函数定义#include utils/debug/asc_printf.h #include acl/acl.h __global__ __vector__ void hello_world() { printf(Hello World!!!\n); }acl/acl.hCANN 的 ACLAscend Computing Language运行时头文件提供设备管理、流管理等 Host 侧接口utils/debug/asc_printf.hAscend C 在设备侧AI Core提供printf能力所需的调试头文件核函数中的printf输出会经运行时回传到 Host 侧终端__global__声明这是一个可被 Host 侧调用的设备核函数__vector__将该核函数标注为SIMD向量类型核函数运行在 AIVAI VectorCore 上这是 simd_cpp_api 编程风格的核心标识。与之对比hello_world_simt 样例 中的核函数只有__global__修饰并在函数体内使用threadIdx.x、blockIdx.x等线程索引体现 SIMT 标量核函数按线程并行的工作方式。4.2 Host 侧调用ACL 运行时五步曲int main(int argc, char const* argv[]) { aclrtSetDevice(0); // 运行管理资源申请。 aclrtStream stream nullptr; aclrtCreateStream(stream); hello_world8, 0, stream(); aclrtSynchronizeStream(stream); aclrtDestroyStream(stream); aclrtResetDevice(0); // 销毁运行资源。 return 0; }Host 侧流程可归纳为标准的五步申请设备资源aclrtSetDevice(0)绑定设备 0创建流aclrtCreateStream(stream)创建执行流核函数与同步操作均挂在流上下发核函数hello_world8, 0, stream()其中8表示启动 8 个 Block对应 8 个 AIV Core0表示不使用动态 UB统一缓冲区空间stream为执行流——这是与 CUDA 启动配置最相似的语法糖流同步aclrtSynchronizeStream(stream)阻塞等待流上所有任务含核函数及 printf 输出回传完成释放资源aclrtDestroyStream销毁流aclrtResetDevice复位设备。说明本样例未显式调用aclInit/aclFinalizeACL 运行时会在首次使用接口时自动完成初始化而 hello_world_simt 样例 中则显式调用了aclInit(nullptr)与aclFinalize()两种写法在当前版本下均可运行显式初始化更利于理解 ACL 生命周期。4.3 编译工程文件CMakeLists.txt 内容如下cmake_minimum_required(VERSION 3.16) set(CMAKE_ASC_ARCHITECTURES dav-2201 CACHE STRING NPU architecture: dav-2201, dav-3510) find_package(ASC REQUIRED) project(kernel_samples LANGUAGES ASC) add_executable(demo hello_world.asc ) target_compile_options(demo PRIVATE $$COMPILE_LANGUAGE:ASC:--npu-arch${CMAKE_ASC_ARCHITECTURES} )要点find_package(ASC REQUIRED)用于定位 CANN 开发套件包仓库通过 cmake/ascend.cmake 完成查找逻辑优先读取环境变量ASCEND_HOME_PATH未设置时按 root/非 root 用户分别探测/usr/local/Ascend/ascend-toolkit/latest、/usr/local/Ascend/latest或$HOME/Ascend/...等默认路径并自动将 Bisheng 编译器路径配置为 C/C 编译器与链接器project(kernel_samples LANGUAGES ASC)声明工程使用 ASC 语言Ascend C编译.asc源文件--npu-arch${CMAKE_ASC_ARCHITECTURES}把架构参数透传给编译器决定生成针对哪种 NPU 架构的二进制。5. 编译与运行三种模式的完整操作5.1 前置准备配置环境变量请根据当前环境上 CANN 开发套件包的安装方式完成环境配置CANN 开发套件安装与准备可参考仓库中 asc-devkit 子模块的快速入门文档见 third_party/asc-devkitsource ${install_path}/cann/set_env.sh说明${install_path}为 CANN 包安装目录未指定安装目录时默认安装至/usr/local/Ascend下。若未设置ASCEND_HOME_PATHCMake 也会按第 4.3 节所述规则自动探测套件路径。5.2 NPU 运行模式默认在本样例目录下执行如下命令mkdir -p build cd build; cmake -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # 编译工程默认npu模式 ./demo5.3 CPU 调试与 NPU 仿真模式使用CPU 调试或NPU 仿真模式时添加-DCMAKE_ASC_RUN_MODEcpu或-DCMAKE_ASC_RUN_MODEsim参数即可cmake -DCMAKE_ASC_RUN_MODEcpu -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # cpu调试模式 cmake -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # NPU仿真模式CPU 调试模式在无 NPU 的纯 CPU 环境上编译运行核函数便于单步调试逻辑相关进阶内容可参考 asc-devkit 中的06_cpu_debug样例NPU 仿真模式通过仿真器在 Host 侧模拟 NPU 行为适合无硬件时的功能验证。注意切换编译模式前需清理 cmake 缓存可在 build 目录下执行rm CMakeCache.txt后重新 cmake。5.4 编译选项说明选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU调试、NPU仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201 对应 Atlas A2 训练系列产品/Atlas A2 推理系列产品和 Atlas A3 训练系列产品/Atlas A3 推理系列产品dav-3510 对应 Ascend 950PR/Ascend 950DTCMAKE_ASC_ARCHITECTURES的默认值dav-2201已在 CMakeLists.txt 中以CACHE STRING形式预置在 Ascend 950 系列产品上运行时务必改为dav-3510同时确保 CANN 版本不低于 9.1.0参见第 2 节支持矩阵。6. 执行结果解读8 个 AIV Core 的并行验证NPU 模式下./demo的成功输出如下[AIV Block 0/8] Hello World!!! [AIV Block 1/8] Hello World!!! [AIV Block 2/8] Hello World!!! [AIV Block 3/8] Hello World!!! [AIV Block 4/8] Hello World!!! [AIV Block 5/8] Hello World!!! [AIV Block 6/8] Hello World!!! [AIV Block 7/8] Hello World!!!输出共 8 行每行对应一个 AIVAI VectorCore 的打印结果。[AIV Block X/8]表示当前为第 X 个核共 8 个核说明核函数已在 8 个 AIV Core 上成功执行。这段输出同时验证了三点核函数成功编译并加载到 NPU未出现编译或运行时错误8, ...的 Block 数生效8 个 Block 分别被调度到 8 个 AIV Core 上每个 Core 独立执行了printf设备侧printf输出回传链路正常asc_printf.h提供的打印能力经 ACL 流同步后正确显示在 Host 终端。7. 延伸从 HelloWorld 走向真实算子开发跑通本样例后可以沿着Samples/0_Introduction/01_simd_cpp_api/的目录梯度继续深入01_add/add第一个带数据搬运与向量加法的 SIMD 算子引入GlobalTensor、LocalTensor与DataCopy等概念02_matrix/matmul_*矩阵乘Matmul的 advanced/basic/tensor 三类 API 用法03_fusion_operation/matmul_leakyrelu_*Matmul 与激活算子融合04_reg_compute/add与gelu寄存器级RegBase标量计算样例。这些样例均复用本文所讲的同一套工程骨架与编译运行方式.ascCMakeLists.txt 三种 RUN_MODE区别仅在于核函数内部的张量操作与流水编排。可以说吃透 hello_world 的工程流程就掌握了整个 cann-samples 仓库中算子样例编译—运行—验证的通法。8. 常见问题与排障提示找不到 CANN 开发套件CMake 报Ascend toolkit not found ... Please set ASCEND_HOME_PATH时说明既未设置环境变量也未在默认路径找到套件。请先执行source ${install_path}/cann/set_env.sh并确认ASCEND_HOME_PATH已指向套件latest目录详见 cmake/ascend.cmake 的探测逻辑。架构不匹配在 Ascend 950 产品上使用默认的dav-2201编译可能导致运行失败请按产品型号选择dav-3510并确认 CANN 版本满足支持矩阵。切换模式后行为异常cpu/sim/npu之间切换前务必清理 build 目录中的CMakeCache.txt避免残留旧配置。无输出或输出不完整确认核函数前已aclrtCreateStream、启动后已aclrtSynchronizeStream设备侧printf结果依赖流同步回传同时确认已包含utils/debug/asc_printf.h头文件。总结hello_world样例虽小却是 Ascend C 核函数直调流程的完整缩影__global__ __vector__定义设备核函数、8, 0, stream下发、ACL 五步资源管理、CMake 中ASC语言与--npu-arch架构编译、三种运行模式切换以及基于 8 个 AIV Core 输出验证执行成功。掌握它就掌握了 cann-samples 仓库中全部算子样例的通用工程方法论为后续向量运算、矩阵乘、算子融合等高性能实战样例打下坚实基础。【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考