昇腾AI CANN训练营〔开源基础系列〕：从语法基础到算子实操，一站式构建 Ascend C 全链路编程能力

Ascend C 算子是什么

Ascend C 算子是基于 CANN 推出的支持 C/C++ 标准规范的编程语言 Ascend C 所开发的算子，编写的算子程序经编译器编译和运行时调度可在昇腾硬件上运行，助力开发者高效实现自定义创新算法；使用它开发自定义算子具有遵循 C/C++ 编程规范、自动并行调度获得最优执行性能、结构化核函数简化算子开发逻辑、CPU/NPU 孪生调试提升算子调试效率等优势。

Ascend C 编程模型

1、Ascend C 的核函数是算子在设备侧 AI Core 的执行入口，也是连接 CPU 与 NPU 的桥梁，编写核函数定义设备端计算逻辑，编译后可在昇腾硬件并行执行，是自定义算子的核心载体；采用 SPMD 模型，仅需写一份核函数代码，设备会自动分发到多 AI Core，各核心通过内置变量 block_idx 区分身份，独立处理数据分片，实现一份代码、多核心并行处理不同数据

2、核函数开发规则（需添加特定限定符，明确运行载体和类型）

• global：标识为可被主机端调用的设备函数
• aicore：明确在昇腾 AI Core 上执行（区别于 CUDA 的核函数）

__global__ __aicore__ void 核函数名(参数列表);  // 返回值必须为void

3、参数与变量规则

入参类型：仅支持指针（需用 gm 标识指向全局内存，如gm float*）或 C/C++ 内置类型（如int32_t）

#define GM_ADDR __gm__ uint8_t* __restrict__  // 定义全局内存指针宏
__global__ __aicore__ void my_kernel(GM_ADDR input, GM_ADDR output, int32_t size);  // 简化入参

4、调用流程核函数如何被主机端触发（主机端通过特定语法调用核函数，需指定并行配置）

• blockDim：指定参与执行的 AI Core 数量
• l2ctrl：保留参数，暂设为 nullptr
• stream：任务队列（ aclrtStream 类型），用于管理设备端任务的并行、串行执行

核函数名<<<blockDim, l2ctrl, stream>>>(参数列表);

5、实践示例HelloWorld 核函数全流程

• 核函数实现（设备侧逻辑）

#include "kernel_operator.h"
using namespace AscendC;

// 定义核函数：每个AI Core打印一次Hello World
extern "C" __global__ __aicore__ void hello_kernel() {
    // 利用block_idx区分不同核心，打印时带上核心编号
    AscendC::printf("Hello World from AI Core %d!\n", block_idx);
}

// 封装调用逻辑（供主机端调用）
void run_hello(uint32_t core_num, aclrtStream stream) {
    // 调用核函数：core_num个核心执行，使用指定stream
    hello_kernel<<<core_num, nullptr, stream>>>();
}

• 主机端调用（CPU 侧控制流程）

需配合 AscendCL（昇腾计算库）完成初始化、资源管理等步骤

#include "acl/acl.h"
extern void run_hello(uint32_t core_num, aclrtStream stream);

int main() {
    aclrtStream stream = nullptr;
    int device_id = 0;

    // 1. 初始化AscendCL
    aclInit(nullptr);
    aclrtSetDevice(device_id);  // 绑定设备
    aclrtCreateStream(&stream);  // 创建任务流

    // 2. 调用核函数：用4个AI Core并行执行
    run_hello(4, stream);
    aclrtSynchronizeStream(stream);  // 等待设备端任务完成

    // 3. 释放资源
    aclrtDestroyStream(stream);
    aclrtResetDevice(device_id);
    aclFinalize();

    return 0;
}

Ascend C 硬件架构抽象与编程范式

Ascend C 基于硬件抽象架构，采用流水线式编程范式，将算子拆分为多个流水任务，通过 Queue 实现任务通信同步、Pipe 统一管理内存资源；其中 Vector 编程范式明确 CopyIn、Compute、CopyOut 三阶段流程，结合逻辑存储位置 TPosition 和张量对象 Global/LocalTensor 管理数据，并通过 Pipe、TBuf 分别完成常规与临时变量的内存分配回收，以此支撑算子在昇腾硬件上的高效开发执行

• 基于硬件抽象架构屏蔽底层差异，采用流水线式编程范式，将算子拆分为流水任务并通过 Queue 实现任务通信同步，依托 Pipe 统一管理设备侧资源
• 核心的 Vector 编程范式明确 CopyIn（数据搬入）、Compute（向量计算）、CopyOut（数据搬出）三阶段流程，结合 TPosition及GlobalTensor/LocalTensor实现精细化数据管理
• 通过 Pipe 分配回收常规内存、TBuf 管理临时变量内存，保障内存高效利用，支撑算子在昇腾硬件上的高效开发与执行

Ascend C API 体系：基础与高层 API 的分层赋能

Ascend C 的 API 体系分为基础 API和高层 API，分别支撑底层功能灵活组合与上层算法快速落地的开发需求

基础 API：底层能力的灵活拼接

1、计算类 API：分为标量（Scalar 单元）、向量（Vector 单元）、矩阵（Cube 单元）三类，适配不同粒度的计算场景

using namespace AscendC;
LocalTensor<half> scalarA(1), scalarB(1), scalarC(1);  // 标量（1元素）
LocalTensor<half> vectorA(128), vectorB(128), vectorC(128);  // 向量（128元素）
LocalTensor<half> matrixA(16, 16), matrixB(16, 16), matrixC(16, 16);  // 矩阵（16x16）
ScalarAdd(scalarA, scalarB, scalarC);  // 标量计算：单个元素加法
VectorAdd(vectorA, vectorB, vectorC);  // 向量计算：128元素并行加法
CubeMatmul(matrixA, matrixB, matrixC);  // 矩阵计算：16x16矩阵乘法

2、数据搬运 API：以DataCopy为核心，实现Global Memory与Local Memory间的数据迁移

// 从GlobalTensor搬入LocalTensor
LocalTensor<half> localIn = pipe.AllocTensor<half>(len);
DataCopy(localIn, globalIn);

// 计算后，从LocalTensor搬出至GlobalTensor
DataCopy(globalOut, localOut);
pipe.FreeTensor(localIn);
pipe.FreeTensor(localOut);

3、内存管理 API：通过AllocTensor/FreeTensor管理内存生命周期

TPipe pipe;
TQue<TPosition::VECIN, 2> que;
pipe.InitBuffer(que, 4, 1024);  // 初始化队列内存
LocalTensor<half> tensor = que.AllocTensor<half>();  // 分配张量
que.FreeTensor(tensor);  // 回收内存

4、任务同步 API：通过EnQue/DeQue实现任务间通信

TQue<TPosition::VECIN, 2> que;
LocalTensor<half> localTensor = ...;
que.EnQue(localTensor);  // 入队
LocalTensor<half> outTensor = que.DeQue<half>();  // 出队

高层 API：上层算法的高效封装

高层 API 封装了 Matmul、Softmax 等常用算法逻辑，借助 “对象化封装 + 流程化调用” 模式，将复杂算法逻辑转化为简洁的 API 调用，既减少了重复开发工作，又大幅提升了开发效率

// 1. 定义Matmul对象（支持数据类型、存储格式定制）
typedef MatmulType<TPosition::GM, CubeFormat::ND, half> AType;
typedef MatmulType<TPosition::GM, CubeFormat::ND, half> BType;
typedef MatmulType<TPosition::GM, CubeFormat::ND, float> CType;
Matmul<AType, BType, CType, TPosition::GM, CubeFormat::ND, float> mm;

// 2. 初始化Matmul
mm.Init(&tiling, &pipe, &blasType);

// 3. 绑定输入输出张量
mm.SetTensorA(globalA);  // 左矩阵A
mm.SetTensorB(globalB);  // 右矩阵B
mm.SetBias(globalBias);  // Bias（可选）

// 4. 执行矩阵乘（迭代或批量模式）
while (mm.Iterate()) {
    mm.GetTensorC(globalC);
}
// 批量执行：mm.IterateAll(globalC);

// 5. 结束矩阵乘
mm.End();

基于 Kernel 直调工程的算子开发

核函数定义

用 global__ __aicore 限定符定义设备侧入口，调用算子类的初始化与处理函数

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

算子类实现

通过CopyIn（数据从 Global 搬入 Local）、Compute（向量加法）、CopyOut（结果搬出至 Global）三个流水任务实现逻辑，结合Queue做任务同步、Pipe做内存管理

class KernelAdd {
public:
    __aicore__ inline KernelAdd() {}  // 设备侧构造函数，在AI Core上执行

    // 初始化：绑定全局内存张量，初始化队列（同步）和管道（资源管理）
    __aicore__ inline void Init(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z) {}

    // 主流程：调度CopyIn、Compute、CopyOut，循环处理多批次数据
    __aicore__ inline void Process() {}

private:
    // 数据搬入：全局内存→本地内存，通过队列同步
    __aicore__ inline void CopyIn(int32_t progress) {}

    // 计算：本地内存数据执行向量加法，结果入队
    __aicore__ inline void Compute(int32_t progress) {}

    // 数据搬出：本地内存→全局内存，通过队列出队获取数据
    __aicore__ inline void CopyOut(int32_t progress) {}

private:
    TPipe pipe;  // 管理设备侧内存与资源调度
    TQue<TPosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;  // 输入数据同步队列
    TQue<TPosition::VECCOUT, BUFFER_NUM> outQueueZ;  // 输出数据同步队列
    GlobalTensor<half> xGM, yGM, zGM;  // 绑定全局内存的张量对象
};

init实现

KernelAdd 类的 Init 方法主要做两件事：通过 block_idx 给当前 AI Core 划分全局内存数据，把 xGM、yGM、zGM 绑定到当前核心的处理区域，实现多核并行；用 pipe 初始化输入输出队列的双缓冲内存，按 TILE_LENGTH 分配内存块，队列深度设为 BUFFER_NUM，支撑流水线并行

// 常量定义：数据分片与双缓冲相关参数
constexpr int32_t TOTAL_LENGTH = 8 * 2048;
constexpr int32_t USE_CORE_NUM = 8;
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;
constexpr int32_t TILE_NUM = 8;
constexpr int32_t BUFFER_NUM = 2;
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM;

class KernelAdd {
    // （类中其他成员声明保持不变）
    __aicore__ inline void Init(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z) {
        // 多核并行：为当前核设置全局内存数据分片的起始地址
        xGM.SetGlobalBuffer((__gm__ half*)x + block_idx * BLOCK_LENGTH, BLOCK_LENGTH);
        yGM.SetGlobalBuffer((__gm__ half*)y + block_idx * BLOCK_LENGTH, BLOCK_LENGTH);
        zGM.SetGlobalBuffer((__gm__ half*)z + block_idx * BLOCK_LENGTH, BLOCK_LENGTH);

        // 双缓冲内存初始化：为输入/输出队列分配内存块（每块容纳TILE_LENGTH个half类型数据）
        pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half));
        pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half));
        pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(half));
    }
    // （类中其他成员声明保持不变）
};

Process() 实现

KernelAdd 类的 Process () 方法通过循环调度 CopyIn、Compute、CopyOut 三个阶段，结合双缓冲实现流水线并行

流程中，CopyIn 将全局内存数据搬入本地并通过队列同步，Compute 执行向量加法并将结果入队，CopyOut 将结果搬回全局内存，借助双缓冲和流水线机制让数据搬运与计算并行，提升昇腾 AI Core 的 Vector 单元利用率

__aicore__ inline void Process() {
    constexpr int32_t loopCount = TILE_NUM * BUFFER_NUM;
    for (int32_t i = 0; i < loopCount; ++i) {
        CopyIn(i);
        Compute(i);
        CopyOut(i);
    }
}

__aicore__ inline void CopyIn(int32_t progress) {
    LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
    DataCopy(xLocal, xGM[progress * TILE_LENGTH], TILE_LENGTH);
    inQueueX.EnQue(xLocal);
    // yLocal同理...
}

__aicore__ inline void Compute(int32_t progress) {
    LocalTensor<half> xLocal = inQueueX.DeQue<half>();
    LocalTensor<half> yLocal = inQueueY.DeQue<half>();
    LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();
    Add(zLocal, xLocal, yLocal, TILE_LENGTH);
    outQueueZ.EnQue(zLocal);
    inQueueX.FreeTensor(xLocal);
    // ...
}

__aicore__ inline void CopyOut(int32_t progress) {
    LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
    DataCopy(zGM[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
    outQueueZ.FreeTensor(zLocal);
}

ddCustom算子的主机侧 main.cpp

• CPU 模式逻辑（用于算法逻辑快速验证，不依赖昇腾硬件，直接在 CPU 上执行计算）

#include "ascend/ascendc.h"
#include <fstream>

int main() {
    // 1. 计算内存大小（输入输出均为8×2048的half类型）
    size_t inputByteSize = 8 * 2048 * sizeof(uint16_t);
    size_t outputByteSize = 8 * 2048 * sizeof(uint16_t);

    // 2. 分配内存
    uint8_t* x = reinterpret_cast<uint8_t*>(ascend::GmAlloc(inputByteSize));
    uint8_t* y = reinterpret_cast<uint8_t*>(ascend::GmAlloc(inputByteSize));
    uint8_t* z = reinterpret_cast<uint8_t*>(ascend::GmAlloc(outputByteSize));

    // 3. 读取输入数据
    ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize);
    ReadFile("./input/input_y.bin", inputByteSize, y, inputByteSize);

    // 4. 执行CPU侧加法（模拟算子逻辑）
    CPU_RUN([&]() {
        half* xHalf = reinterpret_cast<half*>(x);
        half* yHalf = reinterpret_cast<half*>(y);
        half* zHalf = reinterpret_cast<half*>(z);
        for (int i = 0; i < 8 * 2048; ++i) {
            zHalf[i] = xHalf[i] + yHalf[i];
        }
    });

    // 5. 输出结果并释放内存
    WriteFile("./output/output_z.bin", z, outputByteSize);
    ascend::GmFree(x);
    ascend::GmFree(y);
    ascend::GmFree(z);

    return 0;
}

• NPU 模式主机侧逻辑（用于在昇腾 NPU 硬件上执行算子，充分利用硬件加速能力）

#include "acl/acl.h"
#include <fstream>

int main() {
    // 1. ACL初始化：设备、上下文、流
    aclInit(nullptr);
    int32_t deviceId = 0;
    aclrtSetDevice(deviceId);
    aclrtContext context;
    aclrtCreateContext(&context, deviceId);
    aclrtStream stream;
    aclrtCreateStream(&stream);

    // 2. 内存分配（主机侧+设备侧）
    size_t inputByteSize = 8 * 2048 * sizeof(uint16_t);
    size_t outputByteSize = 8 * 2048 * sizeof(uint16_t);
    uint8_t *xHost, *yHost, *zHost;
    uint8_t *xDevice, *yDevice, *zDevice;
    aclrtMallocHost(&xHost, inputByteSize);
    aclrtMallocHost(&yHost, inputByteSize);
    aclrtMallocHost(&zHost, outputByteSize);
    aclMalloc(&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclMalloc(&yDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclMalloc(&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);

    // 3. 读取主机数据并拷贝到NPU设备
    ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize);
    ReadFile("./input/input_y.bin", inputByteSize, yHost, inputByteSize);
    aclMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
    aclMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE);

    // 4. 调用NPU侧自定义加法算子
    add_custom_do(xDevice, yDevice, zDevice, inputByteSize, stream);
    aclrtSynchronizeStream(stream);  // 等待计算完成

    // 5. 结果从设备拷贝回主机并输出
    aclMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
    WriteFile("./output/output_z.bin", zHost, outputByteSize);

    // 6. 释放资源
    aclFree(xDevice);
    aclFree(yDevice);
    aclFree(zDevice);
    aclrtFreeHost(xHost);
    aclrtFreeHost(yHost);
    aclrtFreeHost(zHost);
    aclrtDestroyStream(stream);
    aclrtDestroyContext(context);
    aclrtResetDevice(deviceId);
    aclFinalize();

    return 0;
}

数据生成脚本gen_data.py

基于 NumPy 编写，用于生成 AddCustom 算子的输入数据和真值数据，为算子的功能验证提供标准化的输入与预期结果参考

实现逻辑：

• 生成两个形状为（8, 2048）的 float16 随机输入数据input_x 和 input_y（数值范围 1-100）
• 计算加法算子的真值数据 golden（即 input_x + input_y）
• 将输入数据和真值数据以二进制格式分别保存到指定路径，供后续算子调用、验证时使用

用途：

为 AddCustom 算子的开发、调试和测试提供一致的输入源与真值基准，确保算子计算结果的准确性

#!/usr/bin/python3
# -*- coding:utf-8 -*-
# Copyright 2022-2023 Huawei Technologies Co., Ltd
import numpy as np

def gen_golden_data_simple():
    input_x = np.random.uniform(1, 100, [8, 2048]).astype(np.float16)
    input_y = np.random.uniform(1, 100, [8, 2048]).astype(np.float16)
    golden = (input_x + input_y).astype(np.float16)

    input_x.tofile("./input/input_x.bin")
    input_y.tofile("./input/input_y.bin")
    golden.tofile("./output/golden.bin")

if __name__ == "__main__":
    gen_golden_data_simple()

简单工程执行

Ascend C 算子的简易工程执行方式：将编译和执行命令封装到run.sh脚本中，可通过不同参数实现 CPU 和 NPU 模式下的算子运行调试

• CPU 模式：执行 bash run.sh -v Ascend910x -r cpu ，通过 md5 校验 output_z.bin 与 golden.bin 一致性验证执行成功；架构依赖 Host APP、CPU 调用库、算子 kernel 程序和 AscendC 类库
• NPU 模式：执行 bash run.sh -v Ascend910x -r npu，通过 md5 校验验证编译、执行成功；架构依赖 Host APP、AscendCL API 库及设备侧算子 kernel 程序、AscendC 类库，实现主机 - 设备交互与功能验证

总结

2025年昇腾CANN训练营第二季，基于CANN开源开放全场景，推出0基础入门系列、码力全开特辑、开发者案例等专题课程，助力不同阶段开发者快速提升算子开发技能。获得Ascend C算子中级认证，即可领取精美证书，完成社区任务更有机会赢取华为手机，平板、开发板等大奖

👉昇腾训练营报名链接：
https://www.hiascend.com/developer/activities/cann20252#cann-camp-2502-intro