大家好，我是CANN训练营的新兵。在上一篇文章里，我们一起磕磕绊绊地把Ascend C开发环境给搭了起来。当终端里终于弹出“Build Success”时，我激动了没三秒，就陷入了新的迷茫：这个黑乎乎的命令行，怎么就能变出能在昇腾AI处理器上跑的算子呢？

 

别慌，这篇笔记，就带你踏出这“从0到1”的关键一步——亲手创建并运行你的第一个Ascend C算子：向量加法。这就像编程世界的“Hello World”，简单，却蕴含着整个生态最基本的逻辑。相信我，当你看到两个向量在AI Core上成功相加的那一刻，整个世界都会清晰起来。

 

这正是 [2025年昇腾CANN训练营第二季] “0基础入门系列”课程的精妙之处，它用一个最经典的案例，帮你打通任督二脉。好了，深呼吸，打开你的VSCode，我们的实战开始了！>> 如果你还没上车，快来：点击加入CANN训练营**

第一章：磨刀不误砍柴工——认识“向量加法”工程

在开始敲代码前，训练营的老师花了很大力气让我们理解一个核心概念：Ascend C是一种异构编程模型。啥意思？就是我们的程序会分两部分跑：

1. Host (主机)：就是通用的CPU。它负责准备工作，比如申请内存、装填数据，然后像个指挥官一样，告诉Device“开始干活！”。

2. Device (设备)：就是昇腾AI Core。它是专门干重活的，执行我们写的核心计算代码（核函数）。

    

我们的“向量加法”工程，就是让Host准备两个数组，然后命令Device把这两个数组对应位置相加，结果存到第三个数组里。

用VSCode打开训练营提供的样例工程，你会看到类似这样的结构：

vector_add/
├── src/
│   ├── vector_add_kernel.cpp    // **核心！核函数在这里**
│   └── main.cpp                 // Host侧代码，负责调度
├── scripts/
│   └── build.sh                 // 编译脚本
└── CMakeLists.txt

这个结构非常清晰，我们主要就要对付 src/ 下的两个文件。

第二章：庖丁解牛——读懂你的第一份核函数

怀着敬畏的心情，我打开了 vector_add_kernel.cpp。第一眼是有点懵的，但老师让我们一行行看。

// 1. 引入Ascend C的核心头文件
#include "kernel_operator.h"
​
// 2. 定义核函数 - 这是关键中的关键！
// 注意这里的 __global__ 和函数名 vector_add_custom
extern "C" __global__ __aicore__ void vector_add_custom(
    uint32_t totalLength,   // 数据总长度
    uint32_t tileNum,       // 任务块数量（先理解为1）
    uint8_t* x1,            // 输入向量1的指针
    uint8_t* x2,            // 输入向量2的指针
    uint8_t* y)             // 输出向量的指针
{
    // 3. 初始化内核，为后续操作做准备
    INIT_KERNEL(totalLength, tileNum);
​
    // 4. 计算当前核函数实例要处理的数据块
    // 比如有1000个数据，用100个核来处理，每个核就知道自己该处理哪10个
    int32_t blockIdx = GET_BLOCK_IDX();  // 我是第几个核？
    int32_t blockDim = GET_BLOCK_NUM();  // 总共有多少个核？
    int32_t remainNum = totalLength % blockDim;
    int32_t currentLength = totalLength / blockDim + (blockIdx < remainNum ? 1 : 0);
    int32_t currentOffset = ... // 计算偏移量（略复杂，先不深究）
​
    // 5. 申请本地内存 (Local Memory)
    // 要把数据从Global Memory搬到这里才能计算
    __gm__ uint8_t* globalX1 = x1;
    __gm__ uint8_t* globalX2 = x2;
    __gm__ uint8_t* globalY = y;
​
    constexpr int32_t BUFFER_SIZE = 1024; // 本地缓冲区大小
    uint8_t localX1[BUFFER_SIZE];
    uint8_t localX2[BUFFER_SIZE];
    uint8_t localY[BUFFER_SIZE];
​
    // 6. 数据搬运：把数据从Global Memory拷贝到Local Memory
    for (int32_t i = 0; i < currentLength; ++i) {
        localX1[i] = globalX1[currentOffset + i];
        localX2[i] = globalX2[currentOffset + i];
    }
​
    // 7. 【最核心的一步！】执行计算：向量加法
    for (int32_t i = 0; i < currentLength; ++i) {
        localY[i] = localX1[i] + localX2[i]; // 看，就是这里！
    }
​
    // 8. 结果回写：将计算结果从Local Memory拷贝回Global Memory
    for (int32_t i = 0; i < currentLength; ++i) {
        globalY[currentOffset + i] = localY[i];
    }
}

刚开始，我对 __gm__、Local Memory这些概念一头雾水。但老师打了个比方：Global Memory好比仓库，Local Memory好比车间的工作台。CPU把原料（数据）从仓库搬到工作台，AI Core在工作台上加工（计算），最后再把成品搬回仓库。

这样一想，整个流程瞬间就通了！

第三章：运筹帷幄——编写Host侧的调度官

核函数是干活的士兵，那Host侧（main.cpp）就是发号施令的指挥官。

#include <iostream>
#include <cstdlib>
#include "vector_add_kernel.h" // 包含核函数声明
​
int main() {
    constexpr uint32_t TOTAL_LENGTH = 8; // 我们处理8个数据
    constexpr uint32_t TILE_NUM = 1;     // 先只用一个任务块
​
    // 1. 指挥官(Host)在CPU上申请内存，并准备好数据
    uint8_t* hostX1 = new uint8_t[TOTAL_LENGTH]{1, 2, 3, 4, 5, 6, 7, 8};
    uint8_t* hostX2 = new uint8_t[TOTAL_LENGTH]{8, 7, 6, 5, 4, 3, 2, 1};
    uint8_t* hostY = new uint8_t[TOTAL_LENGTH]{0}; // 输出初始为0
​
    // 2. 指挥官在Device上（昇腾卡）也申请同样大小的内存
    // 这里用了rtMalloc等运行时API，是训练营下一课的重点，先知道是干这个用的
    uint8_t* deviceX1 = nullptr;
    uint8_t* deviceX2 = nullptr;
    uint8_t* deviceY = nullptr;
    rtMalloc(&deviceX1, TOTAL_LENGTH, ...);
    rtMalloc(&deviceX2, TOTAL_LENGTH, ...);
    rtMalloc(&deviceY, TOTAL_LENGTH, ...);
​
    // 3. 把数据从Host内存拷贝到Device内存
    rtMemcpy(deviceX1, hostX1, TOTAL_LENGTH, RT_MEMCPY_HOST_TO_DEVICE);
    rtMemcpy(deviceX2, hostX2, TOTAL_LENGTH, RT_MEMCPY_HOST_TO_DEVICE);
​
    // 4. 【最激动人心的命令！】启动核函数，让Device开始计算！
    vector_add_custom<<<TILE_NUM, nullptr>>>(TOTAL_LENGTH, TILE_NUM, deviceX1, deviceX2, deviceY);
​
    // 5. 等待计算完成...
    rtDeviceSynchronize();
​
    // 6. 把计算结果从Device内存拷回Host内存
    rtMemcpy(hostY, deviceY, TOTAL_LENGTH, RT_MEMCPY_DEVICE_TO_HOST);
​
    // 7. 验货！打印结果
    std::cout << "向量加法结果: ";
    for (int i = 0; i < TOTAL_LENGTH; ++i) {
        std::cout << static_cast<int>(hostY[i]) << " ";
    }
    std::cout << std::endl;
​
    // 8. 打扫战场，释放内存
    delete[] hostX1;
    ... // 其他释放操作
​
    return 0;
}

第四章：编译、运行与“翻车”现场

代码写完了（其实是读懂了），怀着虔诚的心，在终端里执行：

cd scripts
bash build.sh

当屏幕上再次出现 “Build success” 时，我的心跳加速了。然后运行生成的可执行文件。

理想情况：你应该看到 向量加法结果: 9 9 9 9 9 9 9 9。因为 1+8=9, 2+7=9...

而我的实际情况：第一次，我报错了！提示内存拷贝失败。我慌了一秒，但马上想起训练营老师说的“排查三步法”：

1. 检查环境：ascend-cann-version 正常，环境变量没问题。

2. 检查资源：运行 npu-smi info，确认昇腾卡状态是正常的。

3. 检查代码：最后发现，是我在模仿样例时，rtMalloc 的大小写错了！

    

修正后，再次运行。当屏幕上整齐地打印出八个“9”的时候，我对着屏幕傻笑了半天。

结语：从“Hello World”到星辰大海

这个简单的 vector_add，就像你学会的第一个单词“妈妈”。它本身很简单，但通过它，你理解了 核函数、Host/Device分工、数据搬运、本地内存 这一整套Ascend C的核心工作流程。

训练营的巧妙之处就在于此，它不急于灌输复杂的Tiling、流水线，而是用一个最直观的案例，让你先感受到“它跑起来了！”的成就感。这份成就感，是支撑我继续啃下后面更硬骨头的最大动力。

我知道，前面的路还很长，动态Shape、多核并行、性能优化…一个个关卡都在等着我。但至少现在，我已经成功发出了第一声啼哭。

下一期，我计划深入探索《Ascend C内存迷宫通关指南》，彻底搞懂GM, LM和Register那些事儿。希望你能继续关注我的学习笔记。一起加油！

通往算子开发高手之路，从训练营开始 >> 立即报名2025年CANN训练营第二季