你还在对着文档硬啃Ascend C API吗还在为一个算子调好几天吗别慌昇腾CANN给咱们提供了一个新工具——CANNBot。配合Claude Code这个AI代码助手开发算子这件事突然变得简单了。你只需要告诉AI你想要什么它就能帮你搞定项目结构、写出代码、跑通测试。这篇文章我就带你从头到尾走一遍看看怎么用这两个工具开发一个真正的Ascend C算子。一、认识一下这两个工具首先我们来介绍一下什么是CANNBot以及什么是Claude Code。1.1、什么是Skills在说CANNBot之前有必要先聊聊Skills这个概念本身。AI代码助手虽然强大但它的知识是通用性的。当你让它处理专业领域的问题时比如开发昇腾算子它可能给出一个能跑的解法但未必是最符合领域最佳实践的方案。这是因为通用大模型缺乏对特定领域知识的深入理解——它不知道昇腾的编程范式不清楚Ascend C的API使用规范也不了解NPU硬件上的性能调优策略。Skills正是为了解决这一领域知识注入问题而设计的。它本质上是一套结构化的领域知识载体通过将专业知识、经验积累、最佳实践封装为可复用的模块使AI在特定领域能够提供专家级的指导。Skills的设计遵循知识与执行分离的原则——模块本身只负责提供知识储备而具体的执行逻辑由AI根据上下文自行判断。这种设计既保证了知识的专业性和权威性又赋予了AI足够的灵活性来适配不同的应用场景。在实际工程中Skills通常以Markdown文件的形式组织包含该领域的关键概念解释、API使用规范、常见问题解答、最佳实践指南等内容。当AI处理相关任务时会自动识别并加载对应的Skills作为本次对话的知识背景从而显著提升回答的专业度和准确性。1.2、CANNBot昇腾算子开发有其领域特殊性。开发者需要理解NPU的层次化存储架构和计算模型掌握GlobalMemory与LocalMemory之间的数据搬运策略熟悉tiling切分方法论以应对片上内存受限的场景同时还需要具备精度调试和性能分析的能力。这些知识分散在官方文档、培训材料、开源代码和开发者的个人经验中缺乏系统性的整合。CANNBot 是面向CANN开发的用于提升开发效率的系列智能体本仓库为其提供可复用的Skills模块目前已实现 Ascend C / PyPTO / Triton/ TileLang等算子开发全流程覆盖未来将拓展至 CANN 更多技术领域。仓库提了可复用的Skills模块托管于https://atomgit.com/cann/cannbot-skillsCANNBot的价值在于将昇腾开发的全生命周期知识进行系统性沉淀涵盖环境检查、开发工作流、API最佳实践、NPU架构知识、tiling设计、精度调试、运行时调试、测试开发、代码检视等方面。通过这些技能的协同配合AI能够像资深昇腾开发者一样为用户提供系统性的开发指导。1.3、Claude CodeClaude Code是Anthropic推出的AI代码助手采用对话式交互方式开发者可以通过自然语言与它深入交流它会根据上下文理解意图并提供精准的帮助。当你跟Claude Code描述需求时它会自动分析输入内容识别关键意图和技术需求然后从已加载的Skills中选择最相关的模块作为本次对话的知识背景接着结合Skills提供的知识做出专业回应并执行相应操作。在昇腾算子开发场景中配合CANNBot Skills使用开发者无需记忆繁杂的API参数只需要用自然语言描述功能需求AI就会自动加载相关知识、生成代码、验证结果。这种协作模式的核心优势在于开发者无需关心具体应该调用哪个Skill一切由AI助手自动完成。你只需要说我想要一个能加法的算子AI就会像一个经验丰富的昇腾开发者一样给你专业的回答。二、环境准备2.1、安装步骤Claude Code的安装非常简便可以访问官网下载对应平台的安装包或者使用包管理工具进行安装配合CC Switch就可以在顺畅使用了这里不再赘述。安装完成后通过终端或命令行启动即可使用。CANNBot的安装同样简单灵活官方提供了两种方式。方式一通过自然语言安装你可以直接告诉 Claude Code 帮你安装 Skill例如帮我安装几个 skills项目地址是https://atomgit.com/cann/cannbot-skills方式二手动安装下载 Skill 的安装包然后直接复制到 .claude/skills/ 目录下即可.claude/ └── skills/ └── your-skill/ ├── SKILL.md ├── scripts/ └── references/这里我们选择自动安装中间可能会碰到拉取镜像的问题直接选yes就好了Claude会帮你安装好的这里深度使用的话推荐选择全部安装体验的话可以暂时先选择核心开发skills。安装完成后可以验证一下环境是否正确配置。可以试用/ascendc-触发也可以用Claude Code执行环境检查功能来确认一切就绪例如说帮我检查昇腾开发环境是否配置正确或者问它目前已加载的CANNBot Skills有哪些。如果环境有问题它会列出具体的问题清单和修复建议如果环境正常就可以开始开发了。2.2、算子开发环境配置说完了 Claude Code 和 CANNBot 的安装咱们再补一件很多个人开发者最关心的事没有昇腾 NPU 硬件咋办答案是用 CANN 自带的 CPU 仿真cannsim来跑算子。对刚接触昇腾、想学 Ascend C 核函数写法的读者对手里没卡、想先跑通算子验证思路的读者对做模型适配、算子原型的读者CPU 仿真都是性价比最高的起点。而且仿真模式下对学习 Ascend C 算子的运行机理可能反而比真机更友好。CPU 仿真路线只需要装 CANN Toolkit不需要装驱动、固件和 Kernels。去昇腾社区下到对应版本的Ascend-cann-toolkit_ver_linux-x86_64.runaarch64 也支持后给执行权限、跑安装脚本指定一个你能写进去的目录即可。装好后把set_env.shsource 一下环境变量就位了。完整流程如下# 1) 下社区版 toolkitchmodx Ascend-cann-toolkit_ver_linux-x86_64.run ./Ascend-cann-toolkit_ver_linux-x86_64.run--install--install-path$HOME/Ascend# 2) 注入环境变量source$HOME/Ascend/ascend-toolkit/set_env.sh# 3) 验环境没有 npu-smi 是正常的whichcannsimecho$ASCEND_HOMEtoolkit 装好后还需要告诉算子工程这次走 CPU 仿真。最直接的方式也是把这件事甩给 Claude Code——告诉它我在 CPU 仿真环境下开发帮我把 run.sh 改成 sim 模式ascendc-run-helper 这类 skill 会自动识别 RUN_MODE 约定并改写脚本# 在 Claude Code 里说下面进入 CPU 仿真环境下开发帮我把 run.sh 改成 sim 模式。然后就可以正式进入实战了三、实战开发一个Add算子理论讲完了来点实际的。下面通过一个完整案例展示如何使用Claude Code CANN Bot开发一个Ascend C Add算子。3.1、提出需求假设你是刚接触昇腾平台的开发者向Claude Code提出这样的需求“帮我开发一个Ascend C的Add算子输入是两个half类型的张量输出是它们的和。”对就这么简单一句话剩余的工作AI全包了。3.2、AI响应Claude Code收到需求后会自动加载相关Skills然后开始工作。其会自动调用ascendc-env-check技能检测CANN版本、Ascend C工具链、编译环境等各项指标等等。在确认环境就绪后会立即调用核心技能 ascendc-kernel-develop-workflow此技能提供了算子开发的标准化“路线图”用于指导后续操作。3.3、实现核函数进入具体编码阶段Claude Code调用了 ascendc-api-best-practices 技能。此技能是Ascend C API的“使用指南”和“优化手册”提供了函数接口的正确调用方式、内存操作的最佳实践以及性能优化技巧。核函数是算子的核心直接决定算子能做什么。Claude Code会根据API最佳实践帮你实现核函数代码Add是一个典型的Element-wise二元算子使用Ascend C提供的GlobalTensor和LocalTensor来管理数据。完成后生成的目标项目结构如下其中ops/add/ ├── add.asc# 算子实现~300行├── CMakeLists.txt# 构建脚本├── gen_golden.py# Golden数据生成├── run.sh# 运行脚本├── README.md# 项目说明└── docs/ ├── design.md# 设计文档├── acceptance_report.md# 验收报告└── environment.json# 环境检查这样的结构符合昇腾算子项目的最佳实践便于后续的编译、测试和部署。他写出来的核函数核心逻辑大概是这个样子classKernelAdd{public:__aicore__KernelAdd(){}__aicore__ inlinevoidInit(GM_ADDRx1GmAddr,GM_ADDRx2GmAddr,GM_ADDRyGmAddr,AddTilingtiling){tilingDatatiling;// 输入输出 Global Tensor - 使用传入的地址x1Gm.SetGlobalBuffer(reinterpret_cast__gm__ half*(x1GmAddr),tiling.totalLength);x2Gm.SetGlobalBuffer(reinterpret_cast__gm__ half*(x2GmAddr),tiling.totalLength);yGm.SetGlobalBuffer(reinterpret_cast__gm__ half*(yGmAddr),tiling.totalLength);// 分配 Local Buffer带32字节对齐uint32_talignLength(tiling.blockSize31)/32*32;pipe.InitBuffer(x1Local,alignLength*sizeof(half));pipe.InitBuffer(x2Local,alignLength*sizeof(half));pipe.InitBuffer(yLocal,alignLength*sizeof(half));}__aicore__ inlinevoidCopyIn(){// 计算当前核的处理范围uint32_tblockIdxGetBlockIdx();if(blockIdxtilingData.usedCoreNum){return;// 越界检查}uint32_tstartPosblockIdx*tilingData.blockSize;uint32_tremaintilingData.totalLength-startPos;uint32_tprocessSize(remaintilingData.blockSize)?remain:tilingData.blockSize;// GM - UB 拷贝使用DataCopyPad处理任意长度DataCopyPad(x1Local,x1Gm[startPos],processSize);DataCopyPad(x2Local,x2Gm[startPos],processSize);}__aicore__ inlinevoidCompute(){uint32_tblockIdxGetBlockIdx();if(blockIdxtilingData.usedCoreNum){return;}uint32_tstartPosblockIdx*tilingData.blockSize;uint32_tremaintilingData.totalLength-startPos;uint32_tprocessSize(remaintilingData.blockSize)?remain:tilingData.blockSize;// 逐元素加法: y x1 x2Add(yLocal,x1Local,x2Local,processSize);}__aicore__ inlinevoidCopyOut(){uint32_tblockIdxGetBlockIdx();if(blockIdxtilingData.usedCoreNum){return;}uint32_tstartPosblockIdx*tilingData.blockSize;uint32_tremaintilingData.totalLength-startPos;uint32_tprocessSize(remaintilingData.blockSize)?remain:tilingData.blockSize;// UB - GM 拷贝DataCopyPad(yGm[startPos],yLocal,processSize);}__aicore__ inlinevoidProcess(){CopyIn();Compute();CopyOut();}private:// Global TensorGlobalTensorhalfx1Gm;GlobalTensorhalfx2Gm;GlobalTensorhalfyGm;// Local TensorLocalTensorhalfx1Local;LocalTensorhalfx2Local;LocalTensorhalfyLocal;// Tiling 参数私有通过public方法访问AddTilingtilingData;public:// 公开的 tilingData 成员供外部访问AddTilingGetTilingData(){returntilingData;}};// Kernel函数实现externC__global__ __aicore__voidadd_custom(KernelAdd*kernel){kernel-Process();}可以看到他写的还不错这里涉及几个关键概念●GlobalTensor vs LocalTensorGlobalTensor对应设备显存Global Memory容量大但访问慢LocalTensor对应片上高速缓存Unified Buffer容量小但访问快。数据需要先从Global Memory搬到Unified Buffer算完再搬回去。●为什么要分块由于芯片内存容量有限大尺寸张量的算子需要采用分块计算策略通过多核并行处理数据每个核处理一个分块。●DataCopyPad数据搬运操作使用DataCopyPad完成GlobalMemory和Unified Buffer之间的数据传输支持自动填充处理非对齐长度。●多核并行通过GetBlockIdx()获取当前核的索引计算该核负责的数据范围实现多核协同处理。3.4、编写测试代码代码写完后如果需要测试验证Claude Code也能会帮你生成完整的测试代码和用例并完成不同数据规模的测试覆盖以下场景●基础功能测试验证常规的逐元素加法是否正确●大张量测试验证大数据量场景下的正确性●边界值测试检验零值、最大half值、极小值、正负数组合等特殊情况的处理3.5、编译运行项目结构、代码、测试都准备好后接下来进行编译cdadd_operator/build cmake..makeadd_kernel# 编译算子maketest_add# 编译测试程序./build/test_add# 运行测试如果所有测试用例都通过说明算子功能正确可以使用了。四、总结整个开发流程大致如下用自然语言向AI描述算子需求AI自动加载Skills并创建项目结构AI生成代码框架和测试用例开发者确认结果并根据需要进行微调运行测试验证功能正确性根据AI建议进行性能优化这种AI辅助开发模式的核心优势在于●降低门槛开发者无需死记硬背大量API和架构知识●提升效率很多重复性工作由AI自动完成●保障质量昇腾沉淀的最佳实践内置在Skills中●加速排错遇到问题可以直接向AI描述现象获得诊断建议当然AI不是万能的。复杂的算法逻辑设计、性能瓶颈的深入分析、架构层面的决策判断这些仍需要开发者凭借专业经验来完成。但有了CANNBot Claude Code那些繁琐的体力活可以交给AI处理开发者可以把更多精力投入到真正需要创造力的工作中。