一、昇腾好啊，昇腾得学。

昇腾训练营报名链接：
[https://www.hiascend.com/developer/activities/cann20252#cann-camp-2502-intro]
训练营简介：2025 年昇腾 CANN 训练营第二季，基于 CANN 开源开放全场景，推出 0 基础入门系列、码力全开特辑、开发者案例等专题课程，助力不同阶段开发者快速提升算子开发技能。获得 Ascend C 算子中级认证，即可领取精美证书，完成社区任务更有机会赢取华为手机，平板、开发板等大奖。

二、引言：从 “是什么” 到 “为什么学”—— 昇腾 CANN 与 AscendC 入门

先帮小白搞清楚两个核心概念：

• 昇腾 CANN：简单说就是连接 “开发者写的代码” 和 “昇腾 AI 芯片” 的 “翻译官”，它能把我们写的程序优化后，让芯片跑得更快、更高效。

• AscendC：专门给昇腾芯片 “量身定制” 的编程工具（基于大家熟悉的 C/C++），用它写的算子（可以理解为 “芯片要执行的具体计算任务”，比如加法、乘法），能完美适配芯片，发挥最大性能。

在人工智能时代，昇腾芯片是很多 AI 项目的 “算力心脏”，而 AscendC 就是让这颗心脏 “高效跳动” 的关键。本次 “2025CANN 训练营开源基础系列课程”，就是从零基础带大家掌握 AscendC 编程，不仅能学会写算子，还能搞懂芯片底层逻辑，不管是想入门 AI 开发，还是考认证、找工作，这都是重要的基础。

三、核心知识点解析：从基础到核心， AscendC 体系

3.1 先搞懂：AscendC 到底有啥好？—— 四大核心价值

很多小白会问：“我会 C/C++，为啥还要学 AscendC？” 答案就在它的四大优势里，专门解决 “普通编程适配昇腾芯片难” 的问题：

1. 上手快：完全沿用 C/C++ 惯用语法，比如写个变量、循环，和你平时写 C++ 几乎一样，不用重新学一套全新语法；

2. 性能优：自带 “自动并行调度” 功能。比如要处理 1000 个数据，普通代码可能要 “一个一个算”，AscendC 能让芯片 “多个单元一起算”，不用你手动控制，性能直接拉满；

3. 逻辑简：支持 “结构化核函数”。核函数是 AscendC 里最核心的代码（芯片实际执行的计算逻辑），结构化设计能把复杂的计算步骤拆成清晰的模块，比如 “拿数据→算数据→存结果”，小白也能理清逻辑；

4. 调试易：有 “CPU/NPU 孪生调试” 功能。没有昇腾芯片也能学！可以先用电脑的 CPU 模拟芯片运行，提前验证代码对不对，不用等拿到实体硬件才发现问题。

3.2 再明白：昇腾芯片 “内部构造”—— 硬件架构映射（Da Vinci 架构）

要写好 AscendC 代码，得先知道芯片 “怎么干活”。昇腾芯片用的 Da Vinci 架构，就像一个 “小型工厂”，里面有三个关键 “部门”：

硬件模块	作用（用工厂类比）	对应代码操作
AI CPU	“厂长”：统筹数据搬运、任务分配	控制数据从内存到计算单元的流动
AI Core	“车间”：核心计算单元，专门算 AI 任务	执行加法、乘法等算子计算
多层次存储系统	“仓库”：分等级存数据，快的近、慢的远	用不同内存（Global/Local/UB）存数据

这个 “工厂” 能高效运转，关键靠三个 “流水线” 协同：

• 异步指令流（ASync Instruction Flow）：“厂长” 发指令不用等 “车间” 做完，能同时发多个指令，不浪费时间；

• 同步信号流（Sync Signal Flow）：“车间” 和 “仓库” 之间有 “信号灯”，确保 “仓库” 给数据时 “车间” 刚好需要，不弄错顺序；

• 计算数据流（Compute Data Flow）：数据在 “仓库” 和 “车间” 之间流动顺畅，不会断供或堆积。

3.3 核心技能：写核函数 ——SPMD 模型与三阶段流程

核函数是 AscendC 的 “灵魂”，小白记住：写核函数必须用 SPMD 模型，而且流程是固定的，跟着做就不会错。

3.3.1 SPMD 模型：“一份代码，多人干活”

SPMD（Single Program Multiple Data）翻译过来是 “单程序多数据”，简单说就是：你写一份核函数代码，芯片会自动让多个 “计算小单元”（类似工厂里的 “工人”）同时执行，每个 “工人” 处理一部分数据。

举个例子：要处理 2048 个数据，用 8 个 “工人”（8 核并发），芯片会自动把数据分成 8 份，每个 “工人” 算 256 个，不用你手动拆分。

代码里怎么实现？关键靠两个 “标识”：

1. 核函数开头必须写 __global__ __aicore__ void：告诉芯片 “这是要在 AI Core 上跑的核函数”；

2. 用 blockIdx 获取 “工人编号”：比如 blockIdx.x 就是 “第几个工人”，通过编号确定每个 “工人” 处理哪部分数据。

3.3.2 核函数三阶段流程：“拿数据→算数据→存结果”

不管写什么算子，核函数都要走这三步，就像 “工人” 干活的固定流程：

1. copy-in 阶段：“工人” 从 “大仓库”（Global Memory）拿数据，放到 “自己身边的小仓库”（Local Memory/UB 缓冲区），方便快速用；

2. compute 阶段：“工人” 在 “车间”（AI Core）里算数据，比如做向量加法、矩阵乘法；

3. copy-out 阶段：“工人” 把算好的结果放回 “大仓库”（Global Memory），供后续使用。

这三步要顺畅，得靠 “队列（inQ/dQ）” 解决 “数据依赖”：比如 “算数据” 必须等 “拿数据” 完成，队列就像 “传送带”，确保前一步做完数据才到下一步，不混乱。

3.4 性能优化：内存与双缓冲 —— 让芯片 “不摸鱼”

小白写代码先保证 “能跑通”，再学优化。优化的核心是：不让 AI Core（“车间”）等数据，让它一直算。

3.4.1 内存管理：用 Pipe 和 tBuffer 管好 “仓库”

芯片里的内存分等级，速度从快到慢是：UB 缓冲区 > Local Memory > Global Memory。优化的关键是 “尽量用快内存”，而且要 “少申请、多复用”。

代码里怎么管内存？用两个工具：

1. Pipe 内存管理器：统一调度 Local Memory（“车间旁的小仓库”），比如你要存 100 个数据，Pipe 会帮你分配好空间，不用你手动算大小；

2. tBuffer 模板：动态分配内存，用 tb.get() 接口就能拿到内存空间，用完自动回收，不会浪费。

3.4.2 双缓冲机制（Double Buffer）：“边算边拿，不耽误”

双缓冲就像 “两个小推车”：当 “工人” 用 “推车 A” 里的数据计算时，“仓库” 同时把下一批数据装到 “推车 B” 里；等 “推车 A” 用完，直接用 “推车 B”，不用等 “仓库” 装货，AI Core 就一直有数据算，不空闲。

具体用的是 UB 缓冲区（最快的内存）的 A、B 两块区域，代码里会通过切换 “用 A 还是用 B” 实现，后面实操案例会详细讲。

四、进阶实战：两种开发模式 —— 小白选 “快速验证”，部署用 “标准”

很多小白怕 “一上来就写复杂代码”，AscendC 提供了两种开发模式，小白可以从简单的来：

4.1 快速验证模式：小白首选，免硬件、代码少

适合场景：刚写好算子逻辑，想快速验证 “对不对”，不用实体昇腾芯片（用 CPU 仿真）。

优点：

• 步骤少：只需要 “配置环境 + 写核函数”，不用管复杂的部署流程；

• 代码量少：比标准模式少，比如写个向量加法，几十行代码就够；

• 调试快：CPU 上就能跑，错了马上改，不用等芯片。

4.2 标准部署模式：项目落地用，完整流程

适合场景：算子逻辑验证通过后，要放到实际昇腾芯片上跑，用于项目部署。

步骤：比快速验证多 “数据预处理、结果后处理、硬件适配” 等步骤，但核心还是 “核函数三阶段流程”，只是多了一些 “和芯片交互” 的代码（比如申请设备内存、同步数据），后面实操案例会讲。

五、实操案例：向量加法算子 —— 从 0 到 1 跑通（小白版）

下面用 “向量加法”（两个数组的对应元素相加，比如 [1,2] + [3,4] = [4,6]）为例，用 “快速验证模式” 和 “标准部署模式” 分别实现，小白跟着做就能成功。

5.1 环境准备：先搭好 “工作台”

不管哪种模式，先装环境（以 Linux 系统为例）：

1. 装 CANN 8.0+ 和 Ascend Toolkit 8.0+（训练营会给安装包，跟着教程点下一步就行）；

2. 配置环境变量：打开终端，复制粘贴下面的命令，按回车（作用是告诉电脑 “去哪里找 AscendC 的工具”）：

source /usr/local/Ascend/ascend-toolkit/set\_env.sh

5.2 快速验证模式：3 步跑通向量加法（小白重点学）

步骤 1：写核函数（核心代码，只关注 “算” 的逻辑）

新建一个文件 add_kernel.cpp，复制下面的代码，每一行都有注释，小白能看懂：

#include "ascendc/ascendc.h"  // 引入AscendC工具库
#include <vector>
#include <iostream>

// 1. 核函数：向量加法，必须加__global__ __aicore__标识
__global__ __aicore__ void AddKernel_float16(
    const float16* input1,  // 输入向量1（float16是昇腾常用数据类型，占16位，比float快）
    const float16* input2,  // 输入向量2
    float16* output,        // 输出向量（结果）
    uint32_t dataSize       // 向量长度（有多少个元素）
) {
    // 2. SPMD模型：获取当前“工人编号”（threadId），确定处理哪部分数据
    uint32_t threadId = blockIdx.x * blockDim.x + threadIdx.x;
    if (threadId >= dataSize) return;  // 防止“工人”编号超过数据总数，越界报错

    // 3. 核函数三阶段：copy-in→compute→copy-out（这里简化，直接用快速验证的内存）
    // copy-in：因为是快速验证，数据已经在Local Memory，不用手动搬
    // compute：计算 input1[threadId] + input2[threadId]
    float16 result = input1[threadId] + input2[threadId];
    // copy-out：把结果存到输出
    output[threadId] = result;
}

// 4. 测试函数：用CPU仿真跑核函数（免硬件）
void TestAddKernel() {
    uint32_t dataSize = 4;  // 测试用小数据，4个元素，方便看结果
    std::vector<float16> input1 = {1.0f, 2.0f, 3.0f, 4.0f};  // 输入1
    std::vector<float16> input2 = {5.0f, 6.0f, 7.0f, 8.0f};  // 输入2
    std::vector<float16> output(dataSize, 0.0f);  // 输出，初始为0

    // 5. 配置“工人”数量：blockDim=2（每个“组”2个工人），gridDim=2（共2个组），总4个工人
    dim3 blockDim(2);
    dim3 gridDim((dataSize + blockDim.x - 1) / blockDim.x);  // 计算需要多少个组

    // 6. 调用核函数：CPU仿真运行
    AddKernel_float16<<<gridDim, blockDim>>>(input1.data(), input2.data(), output.data(), dataSize);

    // 7. 打印结果，看对不对
    std::cout << "输入1：1.0, 2.0, 3.0, 4.0" << std::endl;
    std::cout << "输入2：5.0, 6.0, 7.0, 8.0" << std::endl;
    std::cout << "输出结果：";
    for (auto val : output) {
        std::cout << (float)val << " ";  // 转成float方便看
    }
    std::cout << std::endl;
}

int main() {
    TestAddKernel();  // 运行测试
    return 0;
}

步骤 2：编译代码（把代码变成可执行程序）

新建一个 build.sh 文件，复制下面的代码（小白不用改，直接用）：

#!/bin/bash
# 告诉电脑用g++编译，链接AscendC的库
g++ -std=c++11 add_kernel.cpp -o add_test \
    -I/usr/local/Ascend/ascend-toolkit/include \
    -L/usr/local/Ascend/ascend-toolkit/lib64 \
    -lascendc -lpthread

然后在终端执行：

chmod +x build.sh  # 给脚本权限

./build.sh         # 执行编译

步骤 3：运行程序，看结果

终端输入 ./add_test，如果输出下面的内容，说明成功了：

输入1：1.0, 2.0, 3.0, 4.0
输入2：5.0, 6.0, 7.0, 8.0
输出结果：6 8 10 12 

5.3 标准部署模式：加硬件交互，项目能用

快速验证通过后，要放到昇腾芯片上跑，需要加 “和芯片交互” 的代码（比如申请设备内存、同步数据）。核心是在快速验证的基础上，加 “内存申请→数据传芯片→跑核函数→结果传回来→释放内存” 这几步。

新建 add_standard.cpp，核心代码如下（重点看和快速验证的区别）：

#include "acl/acl.h"  // 新增：和昇腾芯片交互的库
#include "ascendc/ascendc.h"
#include <vector>
#include <iostream>

// 核函数和快速验证一样，不变
__global__ __aicore__ void AddKernel_float16(
    const float16* input1,
    const float16* input2,
    float16* output,
    uint32_t dataSize
) {
    uint32_t threadId = blockIdx.x * blockDim.x + threadIdx.x;
    if (threadId >= dataSize) return;
    output[threadId] = input1[threadId] + input2[threadId];
}

// 标准部署函数：加芯片交互逻辑
aclError VectorAddStandard(
    const std::vector<float16>& input1,
    const std::vector<float16>& input2,
    std::vector<float16>& output
) {
    aclError ret = ACL_SUCCESS;
    uint32_t dataSize = input1.size();
    if (dataSize != input2.size()) return ACL_ERROR_INVALID_PARAM;  // 输入长度不一致报错

    // 1. 初始化ACL（和芯片建立连接）
    ret = aclInit(nullptr);
    if (ret != ACL_SUCCESS) {
        std::cerr << "和芯片连接失败，错误码：" << ret <</doubaocanvas>

6.1 学习资源

学习官网：https://www.hiascend.com/activities 学习仓库：https://gitcode.com/cann