1. 背景：为何需要 Ascend C？

在大模型时代，AI 算力需求呈指数级增长。通用深度学习框架（如 PyTorch、TensorFlow）虽提供了丰富的高层 API，但在面对以下场景时往往力不从心：

• 框架未支持的新型算子（如稀疏注意力、定制量化）
• 性能瓶颈出现在非标准融合操作
• 需要极致压榨硬件吞吐以降低推理延迟或训练成本

此时，自定义算子成为突破性能天花板的关键路径。而针对昇腾 NPU，Ascend C 正是华为官方推荐的底层开发工具。

📌 关键定位：Ascend C 并非通用编程语言，而是一种 面向昇腾 AI Core 架构的领域特定语言（DSL），基于 C++ 语法扩展，深度融合了昇腾芯片的计算单元（如 Vector Core、Cube Unit）和存储层次。

2. 昇腾 NPU 架构简析：理解 Ascend C 的“舞台”

要写好 Ascend C，必须先理解其运行的硬件环境。昇腾 910B 等主流芯片采用 达芬奇架构（Da Vinci Architecture），核心特点包括：

2.1 三级存储体系

• Global Memory (GM)：片外 HBM/DDR，容量大（数十 GB），但带宽有限、延迟高。
• Unified Buffer (UB)：片上 SRAM，容量小（通常 1–2 MB），但带宽极高（TB/s 级），是数据搬运与计算的核心中转站。
• L0/L1 Cache：紧邻计算单元的寄存器级缓存，用于 Cube 矩阵乘等操作。

2.2 异构计算单元

• AI Core：包含多个 Vector Core（处理向量化操作）和 Cube Unit（专用于 INT8/FP16 矩阵乘）。
• Scalar Core：负责控制流、地址计算等标量任务。
• DMA Engine：高效搬运数据，支持 GM ↔ UB 之间的高带宽传输。

💡 设计哲学：昇腾芯片的性能瓶颈不在计算，而在 访存带宽与延迟。因此，Ascend C 的核心目标是 最大化数据复用、隐藏通信延迟、饱和计算单元。

3. Ascend C 核心编程模型

Ascend C 通过一套声明式 API 和编译器指令，将程序员意图映射到硬件行为。其核心抽象包括：

3.1 Tensor 抽象

• GlobalTensor<T>：指向 GM 的张量，仅用于数据输入/输出。
• LocalTensor<T>：分配在 UB 中的张量，用于中间计算。
• 所有计算操作均在 LocalTensor 上进行。

3.2 流水线执行（Pipeline Execution）

Ascend C 程序被划分为多个 Stage，典型三阶段模型如下：

Stage	操作	硬件资源
CopyIn	GM → UB 数据加载	DMA Engine
Compute	UB 上执行向量化/矩阵运算	Vector/Cube Core
CopyOut	UB → GM 结果写回	DMA Engine

通过 双缓冲（Double Buffering） 技术，Stage i 的 Compute 可与 Stage i+1 的 CopyIn 并行执行，实现 计算掩盖通信。

3.3 内存管理与生命周期

• TPipe 对象用于管理 UB 缓冲区。
• AllocTensor 在 UB 中分配连续内存块。
• 编译器自动插入同步指令，确保数据依赖正确。

4. 实战：从零实现一个高性能 Add 算子

我们以最简单的逐元素加法为例，展示完整的 Ascend C 开发流程。

4.1 环境准备

• 安装 CANN ≥ 7.0
• 配置 Ascend C SDK
• 确保 NPU 驱动正常（npu-smi info）

项目结构：

custom_add/
├── kernel/
│   └── add_kernel.cpp
├── python/
│   └── add_op.py
├── build.sh
└── test_add.py

4.2 Kernel 实现（add_kernel.cpp）

#include "kernel/inc/tikicp.h"
using namespace AscendC;

const int32_t BLOCK_SIZE = 256; // 根据 UB 容量调整

extern "C" __global__ __aicore__ void CustomAddKernel(
    uint32_t totalElements,
    GlobalTensor<float> x,
    GlobalTensor<float> y,
    GlobalTensor<float> output) 
{
    TPipe pipe;
    // 初始化双缓冲区（2 个 buffer，每个 BLOCK_SIZE * sizeof(float)）
    pipe.InitBuffer(pipe, 2, BLOCK_SIZE * sizeof(float));

    LocalTensor<float> xLocal = pipe.AllocTensor<float>(BLOCK_SIZE);
    LocalTensor<float> yLocal = pipe.AllocTensor<float>(BLOCK_SIZE);
    LocalTensor<float> outLocal = pipe.AllocTensor<float>(BLOCK_SIZE);

    uint32_t loopCount = (totalElements + BLOCK_SIZE - 1) / BLOCK_SIZE;

    for (uint32_t i = 0; i < loopCount; ++i) {
        // CopyIn: 从 GM 加载数据到 UB
        DataCopy(xLocal, x[i * BLOCK_SIZE], BLOCK_SIZE);
        DataCopy(yLocal, y[i * BLOCK_SIZE], BLOCK_SIZE);

        // Compute: 向量化加法
        Add(outLocal, xLocal, yLocal, BLOCK_SIZE);

        // CopyOut: 写回 GM
        DataCopy(output[i * BLOCK_SIZE], outLocal, BLOCK_SIZE);
    }
}

4.3 关键点说明

• __global__ __aicore__：标记该函数为可在 AI Core 上执行的核函数。
• DataCopy：由编译器映射为高效 DMA 指令，自动处理地址对齐。
• Add：调用 Vector Core 的 SIMD 加法指令，吞吐达 1024 FP32 ops/cycle。

4.4 编译与注册（Python 层）

# add_op.py
from mindspore import ops
from mindspore.ops import Custom

def custom_add(x, y):
    op = Custom(
        "./custom_add.so",          # 编译生成的 .so 文件
        out_shape=lambda a, b: a.shape,
        out_dtype=lambda a, b: a.dtype,
        func_type="aot"             # Ahead-of-Time 编译模式
    )
    return op(x, y)

使用 build.sh 调用 atc 或 aoe 工具链完成编译。

5. 性能优化进阶：从可用到极致

初始版本的 Add 算子可能仅达到理论带宽的 30%。如何提升？三大优化方向：

5.1 双缓冲流水线

将上述单缓冲改为 ping-pong 双缓冲，使 CopyIn 与 Compute 重叠：

// 分配两组缓冲区
LocalTensor<float> xPing = pipe.AllocTensor<float>(BLOCK_SIZE);
LocalTensor<float> xPong = pipe.AllocTensor<float>(BLOCK_SIZE);
// ... 类似定义 yPing/Pong, outPing/Pong

for (int i = 0; i < loopCount; ++i) {
    if (i % 2 == 0) {
        DataCopy(xPing, x[i*BLOCK_SIZE], BLOCK_SIZE);
        Add(outPing, xPing, yPing, BLOCK_SIZE);
        DataCopy(output[i*BLOCK_SIZE], outPing, BLOCK_SIZE);
    } else {
        DataCopy(xPong, x[i*BLOCK_SIZE], BLOCK_SIZE);
        Add(outPong, xPong, yPong, BLOCK_SIZE);
        DataCopy(output[i*BLOCK_SIZE], outPong, BLOCK_SIZE);
    }
}

✅ 效果：内存带宽利用率提升至 80%+。

5.2 数据类型优化

若精度允许，使用 half（FP16）可使带宽需求减半，吞吐翻倍：

GlobalTensor<half> x, y, output;
LocalTensor<half> xLocal, yLocal, outLocal;

5.3 Block Size 调优

BLOCK_SIZE 需根据 UB 容量和数据类型计算：

Max_BLOCK = UB_Size / (sizeof(T) * num_tensors)

例如 UB=1MB，FP16，3 个张量 → Max_BLOCK ≈ 170K，但实际受对齐限制，通常取 256~1024。

6. 调试与性能分析

• msadvisor：分析 Kernel 是否存在流水线 stall、UB 利用不足。
• Profiler：查看 GM 带宽、计算单元利用率。
• 边界处理：务必处理 totalElements % BLOCK_SIZE != 0 的情况，避免越界。

7. 小结

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

报名链接:https://www.hiascend.com/developer/activities/cann20252