你还在对着文档硬啃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-skills

CANNBot的价值在于将昇腾开发的全生命周期知识进行系统性沉淀，涵盖环境检查、开发工作流、API最佳实践、NPU架构知识、tiling设计、精度调试、运行时调试、测试开发、代码检视等方面。通过这些技能的协同配合，AI能够像资深昇腾开发者一样为用户提供系统性的开发指导。

1.3、Claude Code

Claude 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.run（aarch64 也支持）后，给执行权限、跑安装脚本，指定一个你能写进去的目录即可。装好后把set_env.shsource 一下，环境变量就位了。完整流程如下：

# 1) 下社区版 toolkitchmod+x 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_HOME

toolkit 装好后，还需要告诉算子工程"这次走 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){tilingData=tiling;// 输入输出 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.blockSize+31)/32*32;pipe.InitBuffer(x1Local,alignLength*sizeof(half));pipe.InitBuffer(x2Local,alignLength*sizeof(half));pipe.InitBuffer(yLocal,alignLength*sizeof(half));}__aicore__ inlinevoidCopyIn(){// 计算当前核的处理范围uint32_tblockIdx=GetBlockIdx();if(blockIdx>=tilingData.usedCoreNum){return;// 越界检查}uint32_tstartPos=blockIdx*tilingData.blockSize;uint32_tremain=tilingData.totalLength-startPos;uint32_tprocessSize=(remain<tilingData.blockSize)?remain:tilingData.blockSize;// GM -> UB 拷贝（使用DataCopyPad处理任意长度）DataCopyPad(x1Local,x1Gm[startPos],processSize);DataCopyPad(x2Local,x2Gm[startPos],processSize);}__aicore__ inlinevoidCompute(){uint32_tblockIdx=GetBlockIdx();if(blockIdx>=tilingData.usedCoreNum){return;}uint32_tstartPos=blockIdx*tilingData.blockSize;uint32_tremain=tilingData.totalLength-startPos;uint32_tprocessSize=(remain<tilingData.blockSize)?remain:tilingData.blockSize;// 逐元素加法: y = x1 + x2Add(yLocal,x1Local,x2Local,processSize);}__aicore__ inlinevoidCopyOut(){uint32_tblockIdx=GetBlockIdx();if(blockIdx>=tilingData.usedCoreNum){return;}uint32_tstartPos=blockIdx*tilingData.blockSize;uint32_tremain=tilingData.totalLength-startPos;uint32_tprocessSize=(remain<tilingData.blockSize)?remain:tilingData.blockSize;// UB -> GM 拷贝DataCopyPad(yGm[startPos],yLocal,processSize);}__aicore__ inlinevoidProcess(){CopyIn();Compute();CopyOut();}private:// Global TensorGlobalTensor<half>x1Gm;GlobalTensor<half>x2Gm;GlobalTensor<half>yGm;// Local TensorLocalTensor<half>x1Local;LocalTensor<half>x2Local;LocalTensor<half>yLocal;// Tiling 参数（私有，通过public方法访问）AddTilingtilingData;public:// 公开的 tilingData 成员供外部访问AddTiling&GetTilingData(){returntilingData;}};// Kernel函数实现extern"C"__global__ __aicore__voidadd_custom(KernelAdd*kernel){kernel->Process();}

可以看到，他写的还不错，这里涉及几个关键概念：

●GlobalTensor vs LocalTensor：GlobalTensor对应设备显存（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# 运行测试

如果所有测试用例都通过，说明算子功能正确可以使用了。

四、总结

整个开发流程大致如下：

1. 用自然语言向AI描述算子需求

2. AI自动加载Skills并创建项目结构

3. AI生成代码框架和测试用例

4. 开发者确认结果并根据需要进行微调

5. 运行测试验证功能正确性

6. 根据AI建议进行性能优化

这种AI辅助开发模式的核心优势在于：

●降低门槛：开发者无需死记硬背大量API和架构知识

●提升效率：很多重复性工作由AI自动完成

●保障质量：昇腾沉淀的最佳实践内置在Skills中

●加速排错：遇到问题可以直接向AI描述现象获得诊断建议

当然，AI不是万能的。复杂的算法逻辑设计、性能瓶颈的深入分析、架构层面的决策判断，这些仍需要开发者凭借专业经验来完成。但有了CANNBot+ Claude Code，那些繁琐的"体力活"可以交给AI处理，开发者可以把更多精力投入到真正需要创造力的工作中。