你还在对着文档硬啃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.sh source 一下,环境变量就位了。完整流程如下:

# 1) 下社区版 toolkit
chmod +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 是正常的
which cannsim
echo $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    # 环境检查

在这里插入图片描述

这样的结构符合昇腾算子项目的最佳实践,便于后续的编译、测试和部署。他写出来的核函数核心逻辑大概是这个样子:

 
class KernelAdd {
public:
    __aicore__ KernelAdd() {}
    __aicore__ inline void Init(GM_ADDR x1GmAddr, GM_ADDR x2GmAddr, GM_ADDR yGmAddr,
                                AddTiling tiling) {
        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_t alignLength = (tiling.blockSize + 31) / 32 * 32;
        pipe.InitBuffer(x1Local, alignLength * sizeof(half));
        pipe.InitBuffer(x2Local, alignLength * sizeof(half));
        pipe.InitBuffer(yLocal, alignLength * sizeof(half));
    }
 
    __aicore__ inline void CopyIn() {
        // 计算当前核的处理范围
        uint32_t blockIdx = GetBlockIdx();
        if (blockIdx >= tilingData.usedCoreNum) {
            return;  // 越界检查
        }
 
        uint32_t startPos = blockIdx * tilingData.blockSize;
        uint32_t remain = tilingData.totalLength - startPos;
        uint32_t processSize = (remain < tilingData.blockSize) ? remain : tilingData.blockSize;
 
        // GM -> UB 拷贝(使用DataCopyPad处理任意长度)
        DataCopyPad(x1Local, x1Gm[startPos], processSize);
        DataCopyPad(x2Local, x2Gm[startPos], processSize);
    }
 
    __aicore__ inline void Compute() {
        uint32_t blockIdx = GetBlockIdx();
        if (blockIdx >= tilingData.usedCoreNum) {
            return;
        }
 
        uint32_t startPos = blockIdx * tilingData.blockSize;
        uint32_t remain = tilingData.totalLength - startPos;
        uint32_t processSize = (remain < tilingData.blockSize) ? remain : tilingData.blockSize;
 
        // 逐元素加法: y = x1 + x2
        Add(yLocal, x1Local, x2Local, processSize);
    }
 
    __aicore__ inline void CopyOut() {
        uint32_t blockIdx = GetBlockIdx();
        if (blockIdx >= tilingData.usedCoreNum) {
            return;
        }
 
        uint32_t startPos = blockIdx * tilingData.blockSize;
        uint32_t remain = tilingData.totalLength - startPos;
        uint32_t processSize = (remain < tilingData.blockSize) ? remain : tilingData.blockSize;
 
        // UB -> GM 拷贝
        DataCopyPad(yGm[startPos], yLocal, processSize);
    }
 
    __aicore__ inline void Process() {
        CopyIn();
        Compute();
        CopyOut();
    }
 
private:
    // Global Tensor
    GlobalTensor<half> x1Gm;
    GlobalTensor<half> x2Gm;
    GlobalTensor<half> yGm;
 
    // Local Tensor
    LocalTensor<half> x1Local;
    LocalTensor<half> x2Local;
    LocalTensor<half> yLocal;
 
    // Tiling 参数(私有,通过public方法访问)
    AddTiling tilingData;
 
public:
    // 公开的 tilingData 成员供外部访问
    AddTiling& GetTilingData() { return tilingData; }
};
 
// Kernel函数实现
extern "C" __global__ __aicore__ void add_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、编译运行

项目结构、代码、测试都准备好后,接下来进行编译:

cd add_operator/build
cmake ..
make add_kernel   # 编译算子
make test_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处理,开发者可以把更多精力投入到真正需要创造力的工作中。

Logo

CANN开发者社区旨在汇聚广大开发者,围绕CANN架构重构、算子开发、部署应用优化等核心方向,展开深度交流与思想碰撞,携手共同促进CANN开放生态突破!

更多推荐