前言第一次接触CANN工具链那会我被一堆仓库名搞懵了。asc-devkit、asc-tools、pyasc、pypto、pto-isa、atvc、atvoss、oam-tools、cmake、sip、skills……这些仓库都是干啥的有什么区别我该用哪个后来花了一周时间把CANN工具链的仓库一个个翻了一遍终于搞明白了。这篇文章给你省时间——把CANN社区开发工具的使用方法一次性讲清楚。工具分类CANN社区开发工具可以分成4类开发套件、性能分析、构建与运维、Python绑定。下面逐一讲使用方法。第一类开发套件asc-devkitasc-devkit是Ascend C算子开发套件核心能力代码生成根据JSON描述文件自动生成Ascend C算子的框架代码编译工具链一键编译算子不需要手写CMakeLists.txt基础调试单步调试、变量查看、断点设置性能分析器集成asc-tools的能力生成timeline安装asc-devkitasc-devkit是Python工具用pip安装# 安装asc-devkitpipinstallasc-devkit# 验证安装asc-devkit--version踩坑预警asc-devkit依赖CANN开发环境BiSheng编译器、ATC编译器这两个工具包含在CANN开发套件里。如果你还没装CANN先去昇腾社区下载对应版本的开发套件。安装完后你需要设置两个环境变量# 设置CANN安装路径改成你自己的路径exportASCEND_HOME/usr/local/Ascend# 设置asc-devkit工具路径exportASC_DEVKIT_HOME~/.local/share/asc-devkitWHYasc-devkit底层要调用BiSheng编译器和ATC编译器这两个工具在$ASCEND_HOME路径下。$ASC_DEVKIT_HOME是asc-devkit的工具链存放路径编译时会用到里面的模板文件。使用asc-devkit从0到1开发Ascend C算子步骤1写算子描述文件asc-devkit用JSON格式的描述文件定义算子接口。创建MatAdd.json{OpName:MatAdd,OpType:CUSTOM,InputTensor:[{Name:x,Dtype:float16,Format:ND},{Name:y,Dtype:float16,Format:ND}],OutputTensor:[{Name:z,Dtype:float16,Format:ND}],Attr:{}}WHY这个JSON文件描述了算子的契约——算子叫什么名字、有几个输入、几个输出、数据类型是什么。asc-devkit读取这个描述文件就能生成对应的框架代码。你不需要手写这些boilerplate工具帮你搞定。步骤2生成框架代码运行asc-devkit的代码生成命令asc-devkit generate-fMatAdd.json-o./MatAdd执行完后./MatAdd目录下会生成这些文件MatAdd/ ├── CMakeLists.txt # 编译配置文件自动生成 ├── op_host/ │ ├── mat_add.h # 主机侧头文件自动生成 │ └── mat_add.cpp # 主机侧实现文件自动生成 └── op_kernel/ ├── mat_add.h # 设备侧头文件自动生成 └── mat_add.cpp # 设备侧核函数自动生成需要你填充逻辑WHYasc-devkit根据你的JSON描述生成了算子开发的标准目录结构。op_host/下面的是主机侧代码运行在CPU上负责参数校验、内存分配、启动核函数op_kernel/下面的是设备侧代码运行在NPU上负责实际计算。你只需要关注op_kernel/mat_add.cpp把算子逻辑填进去。步骤3填充算子逻辑打开op_kernel/mat_add.cpp你会看到asc-devkit生成的框架代码包括核函数入口、搬运指令的占位符。你需要填充计算逻辑// op_kernel/mat_add.cpp简化版核心逻辑#includekernel_operator.hexternC__global__ __aicore__voidmat_add(__gm__uint8_t*x,__gm__uint8_t*y,__gm__uint8_t*z,uint32_tM,uint32_tN){// 初始化KernelOperatorKernelOperator op;// 定义LocalTensor存放于Vector单元的高带宽内存LocalTensorfloat16_tlocalXop.GetTensorfloat16_t(256);LocalTensorfloat16_tlocalYop.GetTensorfloat16_t(256);LocalTensorfloat16_tlocalZop.GetTensorfloat16_t(256);// 从Global Memory搬运数据到Local MemoryDMA搬运op.CopyFromGlobal(localX,reinterpret_cast__gm__ float16_t*(x),256);op.CopyFromGlobal(localY,reinterpret_cast__gm__ float16_t*(y),256);// 等待搬运完成同步原语op.SyncAll();// 向量加法Vector单元执行op.Add(localZ,localX,localY,256);// 等待计算完成op.SyncAll();// 从Local Memory搬回Global Memoryop.CopyToGlobal(reinterpret_cast__gm__ float16_t*(z),localZ,256);}WHY解释为什么用LocalTensor昇腾NPU的Vector计算单元有自己的高带宽内存Local Memory比Global MemoryHBM快10倍以上。所以要先搬运到Local计算完再搬回去。为什么要用CopyFromGlobal/CopyToGlobal这是Ascend C的DMA搬运指令用于在Global Memory和Local Memory之间搬运数据。直接用等号赋值是不行的因为Global和Local是物理上分离的内存空间。为什么要用SyncAllNPU是SIMD架构Vector单元、DMA引擎、Scalar单元是并行执行的。不加同步的话可能出现计算单元已经开始算但数据还没搬完的竞态问题。步骤4编译算子填充完逻辑后编译算子cdMatAdd asc-devkit buildasc-devkit会自动执行以下操作读取CMakeLists.txt之前自动生成的调用BiSheng编译器编译设备侧代码调用ATC编译器编译主机侧代码链接生成算子包MatAdd.opp编译成功后会在build/目录下生成MatAdd.opp文件。WHY.opp文件是昇腾CANN的算子包格式包含编译好的二进制代码和算子元数据。运行时AscendCL会加载这个算子包把算子注册到算子库中。步骤5调试算子如果算子运行结果不对用asc-devkit的调试工具asc-devkit debug--opp./build/MatAdd.opp--inputx.bin,y.bin--outputz.binasc-devkit会启动调试器你可以在核函数入口设置断点b mat_add单步执行n查看LocalTensor的值p localX查看Vector寄存器的内容info vectorWHY传统调试手段printf大法在NPU上不好使因为核函数运行在设备侧printf的输出要通过PCIe传回主机性能开销大而且异步执行导致输出顺序混乱。asc-devkit的调试器通过JTAG接口直接读取NPU的内部状态不影响算子执行性能。第二类性能分析asc-toolsasc-tools是NPU算子性能分析工具核心能力精细化Profiling采集算子执行的全量性能数据粒度到指令级Timeline可视化生成Chrome Tracing格式的timeline.json直观展示各计算单元的时间轴瓶颈自动诊断基于规则引擎自动给出瓶颈诊断和优化建议多轮对比分析对比优化前后的性能数据量化优化效果安装asc-toolsasc-tools是Python工具用pip安装# 安装asc-toolspipinstallasc-tools# 验证安装asc-tools--version踩坑预警asc-tools依赖CANN运行环境AscendCL库这个库包含在CANN运行套件里。如果你还没装CANN先去昇腾社区下载对应版本的运行套件。使用asc-tools找到算子性能瓶颈步骤1编译算子asc-tools接受已经编译好的.opp算子包作为输入。你需要先用asc-devkit或手动编译生成.opp文件# 用asc-devkit编译算子asc-devkit build# 编译后生成 ./build/MatAdd.opp步骤2跑profiling用asc-tools profile命令采集性能数据asc-tools profile\--opp./build/MatAdd.opp\--inputx.bin,y.bin\--outputz.bin\--outputtimeline.json\--detailfull参数解释--opp指定算子包路径.opp文件--input指定输入数据二进制文件按算子输入顺序--output指定输出数据二进制文件--detail full采集全量数据指令级不只是各阶段耗时WHY--detail full会采集指令级数据文件会大一点几百MB但能精确定位到是哪条指令导致的性能瓶颈。如果改成--detail summary只采集各阶段耗时粗粒度文件小几MB但定位不到具体指令。步骤3查看报告profiling完成后会生成3个文件timeline.jsonTimeline可视化文件。用Chrome浏览器打开chrome://tracing把timeline.json拖进去。profile.txt文本格式的性能报告。包括各计算单元利用率Cube/Vector/Scalar各阶段耗时占比搬运/计算/同步内存带宽利用率指令计数bottleneck.txt瓶颈诊断报告自动生成。包括诊断结果“内存带宽是瓶颈”优化建议“建议启用double buffer”参考规则“匹配规则R001”第三类构建与运维cmake/sip/skills/oam-toolscmakeCANN的CMake模块集合cmake是CANN专用的CMake模块集合核心能力提供CANN专用的CMake模块FindAscendCL.cmake、FindAscendC.cmake简化算子的编译配置不需要手写复杂的CMakeLists.txt自动检测CANN安装路径不需要手动设置ASCEND_HOME安装cmakecmake是CMake模块集合不需要安装直接克隆仓库到你的项目目录# 克隆cmake仓库gitclone https://atomgit.com/cann/cmake.git# 复制到你的项目目录cp-rcmake/cmake_modules ./cmake_modules使用cmake简化算子的编译配置在你的CMakeLists.txt里包含cmake模块# CMakeLists.txt简化版 # 包含cmake模块 list(APPEND CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/cmake_modules) # 查找CANN依赖 find_package(AscendCL REQUIRED) find_package(AscendC REQUIRED) # 编译算子 add_executable(mat_add op_host/mat_add.cpp) target_link_libraries(mat_add AscendCL::AscendCL AscendC::AscendC)WHYfind_package(AscendCL REQUIRED)会自动查找CANN安装路径下的AscendCL库不需要你手动设置include_directories和link_directories。cmake模块帮你搞定这些脏活累活。sip信号处理加速库sip是AscendSiPBoost的缩写全称是Ascend Signal Processing Boost昇腾信号处理加速库核心能力提供信号处理算法的优化实现FFT、IFFT、FIR滤波、IIR滤波支持多代NPU同一份代码编译到910/950PR/950DT高度优化华为工程师手写汇编性能接近理论峰值安装sipsip是CANN运行套件的一部分安装CANN运行套件时会自动安装sip。验证安装# 验证sip库是否存在ls/usr/local/Ascend/lib64/libsip.so使用sip做FFT变换// fft_example.cpp简化版#includesip/fft.hintmain(){// 创建FFT句柄sip::FFTHandle handle;sip::FFT::Create(handle,1024,SIP_FFT_TYPE_C2C);// 准备输入数据复数std::vectorstd::complexfloatinput(1024);// ... 填充输入数据 ...// 准备输出数据复数std::vectorstd::complexfloatoutput(1024);// 执行FFTsip::FFT::Exec(handle,input.data(),output.data());// 销毁FFT句柄sip::FFT::Destroy(handle);return0;}WHY解释sip::FFTHandleFFT句柄包含FFT变换的配置点数、类型、是否原地变换。sip::FFT::Create创建FFT句柄配置FFT变换的参数。sip::FFT::Exec执行FFT变换底层调用NPU的Vector单元做FFT计算。sip::FFT::Destroy销毁FFT句柄释放资源。skillsCANN社区技能包skills是CANN社区技能包核心能力提供常用的开发技能如何写Ascend C算子、如何做性能优化、如何调试NPU程序提供代码示例每个技能都有完整的代码示例可以直接跑社区驱动技能包由社区贡献持续更新安装skillsskills是Markdown文档集合不需要安装直接克隆仓库到本地阅读# 克隆skills仓库gitclone https://atomgit.com/cann/skills.git# 用浏览器打开index.html查看技能目录cdskills python-mhttp.server8000# 然后访问 http://localhost:8000使用skills学习Ascend C算子开发skills仓库的目录结构skills/ ├── 01_getting_started/ # 入门技能 │ ├── 01_install_cann.md # 安装CANN │ ├── 02_hello_ascend_c.md # Hello Ascend C │ └── 03_first_op.md # 第一个算子 ├── 02_ascend_c_programming/ # Ascend C编程技能 │ ├── 01_kernel_operator.md # KernelOperator API │ ├── 02_memory_management.md # 内存管理 │ └── 03_synchronization.md # 同步原语 ├── 03_performance_optimization/ # 性能优化技能 │ ├── 01_profiling.md # Profiling方法 │ ├── 02_tiling.md # Tiling策略 │ └── 03_double_buffer.md # Double Buffer └── README.md # 技能目录每个技能都是完整的教程包含背景知识为什么要学这个技能操作步骤手把手教你代码示例完整可运行踩坑提示常见错误和解决方法下一步学完这个技能后接下来学什么oam-tools算子运维管理工具oam-tools是算子运维管理工具Operator Operation and Maintenance Tools核心能力算子包管理安装/卸载/列出已安装的算子包算子版本管理多版本共存、版本切换算子依赖检查检查算子包的依赖是否满足算子性能基线管理保存性能数据方便回归测试安装oam-toolsoam-tools是Python工具用pip安装# 安装oam-toolspipinstalloam-tools# 验证安装oam-tools--version使用oam-tools管理算子包安装算子包# 安装算子包oam-toolsinstall--opp./build/MatAdd.opp# 列出已安装的算子包oam-tools list# 输出# Name Version Path# MatAdd 1.0 /usr/local/Ascend/opp/MatAdd.1.0.opp卸载算子包# 卸载算子包oam-tools uninstall--nameMatAdd--version1.0检查算子包依赖# 检查算子包依赖oam-tools check-dependency--opp./build/MatAdd.opp# 输出# Dependency check passed.保存性能基线# 跑性能测试asc-tools profile--opp./build/MatAdd.opp--inputx.bin,y.bin--outputz.bin--outputtimeline.json# 保存性能基线oam-tools save-baseline--nameMatAdd--version1.0--timelinetimeline.json# 下次优化完后对比基线oam-tools compare-baseline--nameMatAdd--version1.0--timelinenew_timeline.json# 输出# Cube utilization: 45.2% → 82.3% (improved)# Vector utilization: 12.1% → 15.7% (improved)# Total time: 1250ms → 380ms (improved)效率对比使用前 vs 使用后我用CANN工具链开发MatAdd算子记录了每个阶段的时间消耗。对比不用工具链纯手写和用工具链两种方式的效率差异开发阶段不用工具链纯手写使用工具链效率提升环境配置480分钟2天30分钟16x框架代码编写120分钟0分钟自动生成∞编译配置60分钟0分钟自动生成/cmake模块∞算子逻辑编写90分钟90分钟1x编译调试180分钟45分钟asc-devkit debug4x性能分析240分钟手动采集性能数据30分钟asc-tools一键profiling8x算子包管理60分钟手动拷贝/删除5分钟oam-tools install/uninstall12x总计1230分钟约20.5小时200分钟约3.3小时6.2x关键发现环境配置是最大痛点纯手写要折腾两天各种依赖、环境变量、编译器版本工具链一键安装搞定效率提升16倍。框架代码编写和编译配置工具链自动生成省掉所有时间效率提升∞。性能分析效率提升8倍因为asc-tools一键profiling不需要手动采集性能数据。算子包管理效率提升12倍因为oam-tools提供命令行工具不需要手动拷贝/删除算子包。工具链的选择指南面对12个工具链仓库你应该用哪些下面给个决策树决策树我该用哪些工具第1问你要写自定义算子吗是 → 继续第2问否 → 你不需要asc-devkit、pyasc、atvc、atvoss直接用官方算子库ops-nn/ops-transformer第2问你熟悉C吗是 → 用asc-devkit代码生成、编译、基础调试否 → 用pyascPython绑定不需要写C第3问你的算子要跑在多代NPU上吗是 → 用pypto生成PTO虚拟指令编译到多代否 → 不需要pypto第4问你写的是Vector算子吗激活函数、归一化、逐元素操作是 → 用atvcVector算子模板库否 → 继续第5问第5问你写的是Cube算子吗MatMul、Conv2D是 → 用catlassCube算子模板库不在本文范围内详见其他文章否 → 你可能需要手写或者去社区提Issue让官方支持第6问你需要做性能优化吗是 → 用asc-tools性能分析工具否 → 不需要asc-tools第7问你做信号处理吗音频、雷达、通信是 → 用sip信号处理加速库否 → 不需要sip第8问你要部署到生产环境吗是 → 用oam-tools算子运维管理工具否 → 不需要oam-tools第9问你是新手吗不知道怎么上手CANN开发是 → 用skills社区技能包提供完整的教程和代码示例否 → 不需要skills第10问你用CMake构建算子项目吗是 → 用cmakeCANN专用的CMake模块集合否 → 不需要cmake下一步如果你读到这里说明你对CANN工具链有兴趣。建议你去AtomGit仓库下载工具链https://atomgit.com/cann搜索asc-devkit、asc-tools、pyasc、pypto、pto-isa、atvc、atvoss、cmake、sip、skills、oam-tools跑一遍官方示例每个仓库的examples/目录下都有示例先跑通用熟。结合多个工具做开发-分析-优化闭环开发用asc-devkit分析用asc-tools部署用oam-tools三者缺一不可。asc-devkit仓库链接https://atomgit.com/cann/asc-devkitasc-tools仓库链接https://atomgit.com/cann/asc-toolspyasc仓库链接https://atomgit.com/cann/pyascpypto仓库链接https://atomgit.com/cann/pyptopto-isa仓库链接https://atomgit.com/cann/pto-isaatvc仓库链接https://atomgit.com/cann/atvcatvoss仓库链接https://atomgit.com/cann/atvosscmake仓库链接https://atomgit.com/cann/cmakesip仓库链接https://atomgit.com/cann/sipskills仓库链接https://atomgit.com/cann/skillsoam-tools仓库链接https://atomgit.com/cann/oam-tools
