引言

随着人工智能技术的迅猛发展，专用 AI 芯片成为提升计算效率的关键。华为昇腾（Ascend）系列 AI 处理器凭借其高能效比和强大的并行计算能力，在大模型训练与推理领域崭露头角。为了充分发挥昇腾芯片的硬件潜力，华为推出了 Ascend C —— 一种专为昇腾 NPU（神经网络处理单元）设计的高性能编程语言。

Ascend C 并非传统意义上的 C/C++ 扩展，而是一种融合了底层硬件调度、内存管理、流水线控制等特性的领域特定语言（DSL），旨在让开发者能够以接近硬件的方式编写高效算子（Operator）。本文将带你从零开始，深入理解 Ascend C 的编程模型，并通过一个完整的自定义算子开发流程，展示如何利用 Ascend C 构建高性能 AI 计算单元。

目标读者：熟悉 C/C++ 编程、了解基本 AI 概念、希望深入昇腾生态进行底层优化的开发者。

一、Ascend C 是什么？

1.1 背景与定位

Ascend C 是华为 MindSpore 生态中用于编写 自定义算子（Custom Operator）的核心工具。它运行于昇腾 AI 处理器（如 Ascend 910/310）之上，直接操作 NPU 的计算单元（Cube 单元、Vector 单元等）和片上内存（Unified Buffer, UB）。

与 CUDA 或 OpenCL 不同，Ascend C 的抽象层级更贴近硬件，强调 数据流驱动 和 流水线并行。开发者需显式管理：

  • 数据搬运（Global Memory ↔ Unified Buffer）
  • 计算调度（Cube/Vector 指令）
  • 同步机制（PipeLine 控制）

1.2 核心优势

  • 极致性能：绕过框架开销，直接调用硬件指令。
  • 细粒度控制：可精确控制内存布局、计算顺序、流水阶段。
  • 与 MindSpore 无缝集成：编译后可作为标准算子接入 MindSpore 图执行。

二、Ascend C 编程模型详解

2.1 内存层次结构

昇腾 NPU 采用三级内存架构：

层级 名称 特点
L0 Global Memory (GM) 容量大（GB 级），带宽低，延迟高
L1 Unified Buffer (UB) 片上高速缓存（MB 级），带宽高，延迟低
L2 L1 Cache / Register 寄存器级，自动管理

关键原则：所有计算必须在 UB 中进行，GM 仅用于 I/O。

2.2 计算单元

  • Cube Unit：专用于矩阵乘（GEMM），支持 INT8/FP16/BF16。
  • Vector Unit：用于向量运算（加、乘、激活函数等）。
  • Scalar Unit：控制流、地址计算等。

2.3 流水线模型（Pipeline）

Ascend C 采用 三阶段流水线：

  1. CopyIn：从 GM 搬运数据到 UB
  2. Compute：在 UB 上执行计算
  3. CopyOut：将结果写回 GM

通过重叠这三个阶段，可实现高吞吐。

三、开发环境准备

3.1 软件依赖

  • CANN（Compute Architecture for Neural Networks）5.1+
  • Ascend C 编译器（atc）
  • MindSpore 2.0+

3.2 目录结构

custom_op/
├── kernel/
│   └── add_custom.cpp      # Ascend C 算子实现
├── op/
│   └── add_custom.py       # Python 注册接口
└── build.sh                # 编译脚本

四、实战：编写一个自定义 Add 算子

我们将实现一个 AddCustom 算子，功能为 C = A + B，但使用 Ascend C 优化内存访问与计算。

4.1 算子定义（Python 层）

# op/add_custom.py
import mindspore as ms
from mindspore.ops import PrimitiveWithInfer

class AddCustom(PrimitiveWithInfer):
    def __init__(self):
        super().__init__("AddCustom")
        self.init_prim_io_names(inputs=['x', 'y'], outputs=['output'])

    def infer_shape(self, x_shape, y_shape):
        return x_shape

    def infer_dtype(self, x_dtype, y_dtype):
        return x_dtype

4.2 Ascend C 核心实现

// kernel/add_custom.cpp
#include "kernel_operator.h"

using namespace AscendC;

constexpr int32_t BUFFER_NUM = 2; // 双缓冲
constexpr int32_t BLOCK_SIZE = 256;
constexpr int32_t TILE_SIZE = 64;

class AddCustom {
public:
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalSize) {
        this->xGm.SetGlobalBuffer((__gm__ float*)x, totalSize);
        this->yGm.SetGlobalBuffer((__gm__ float*)y, totalSize);
        this->zGm.SetGlobalBuffer((__gm__ float*)z, totalSize);
        this->totalSize = totalSize;

        // 初始化 UB 缓冲区
        DataCopy(xUb, xGm, totalSize);
        DataCopy(yUb, yGm, totalSize);
    }

    __aicore__ inline void Process() {
        // 计算需要多少个 tile
        uint32_t loopCount = (totalSize + TILE_SIZE - 1) / TILE_SIZE;

        for (uint32_t i = 0; i < loopCount; i++) {
            uint32_t processSize = (i == loopCount - 1) ? (totalSize - i * TILE_SIZE) : TILE_SIZE;

            // 从 GM 拷贝 A 和 B 到 UB
            CopyIn(i, processSize);

            // 执行加法
            Compute(i, processSize);

            // 写回结果
            CopyOut(i, processSize);
        }
    }

private:
    __aicore__ inline void CopyIn(uint32_t loopIndex, uint32_t processSize) {
        // 使用 Pipe 让拷贝异步
        auto pipe = GetPipe();
        pipe.CopyIn(xUb[loopIndex % BUFFER_NUM], xGm[loopIndex * TILE_SIZE], processSize);
        pipe.CopyIn(yUb[loopIndex % BUFFER_NUM], yGm[loopIndex * TILE_SIZE], processSize);
        pipe.WaitPipe();
    }

    __aicore__ inline void Compute(uint32_t loopIndex, uint32_t processSize) {
        auto dst = zUb[loopIndex % BUFFER_NUM];
        auto src0 = xUb[loopIndex % BUFFER_NUM];
        auto src1 = yUb[loopIndex % BUFFER_NUM];

        // Vector 加法
        VecAdd(dst, src0, src1, processSize);
    }

    __aicore__ inline void CopyOut(uint32_t loopIndex, uint32_t processSize) {
        auto pipe = GetPipe();
        pipe.CopyOut(zGm[loopIndex * TILE_SIZE], zUb[loopIndex % BUFFER_NUM], processSize);
        pipe.WaitPipe();
    }

private:
    GlobalTensor<float> xGm, yGm, zGm;
    TBuf<float> xUb[BUFFER_NUM], yUb[BUFFER_NUM], zUb[BUFFER_NUM];
    uint32_t totalSize = 0;
};

extern "C" __global__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalSize) {
    AscendC::SetSysCtrl();
    auto taskId = GetBlockIdx();
    if (taskId != 0) return;

    AddCustom op;
    op.Init(x, y, z, totalSize);
    op.Process();
}

4.3 关键代码解析

  • GlobalTensor：绑定 GM 地址。
  • TBuf：UB 中的临时缓冲区，支持双缓冲（BUFFER_NUM=2）。
  • VecAdd：Ascend C 内置向量加法指令，自动向量化。
  • Pipe：实现 CopyIn/Compute/CopyOut 的流水线重叠。

注意：实际开发中需考虑对齐（16-byte）、分块策略、多核调度等。

五、性能优化技巧

5.1 内存对齐

昇腾要求 UB 访问地址 32-byte 对齐。使用 ALIGN_SIZE 宏：

constexpr int32_t ALIGN_SIZE = 32;
int32_t alignedSize = ((processSize + ALIGN_SIZE - 1) / ALIGN_SIZE) * ALIGN_SIZE;

5.2 双缓冲（Double Buffering）

通过 BUFFER_NUM=2，当前计算使用 buffer 0，同时下一组数据拷贝到 buffer 1，隐藏 I/O 延迟。

5.3 向量化与分块

  • 尽量使用 Vec* 系列指令（VecAdd, VecMul 等）。
  • 分块大小（TILE_SIZE）应匹配 UB 容量（通常 64~256KB）。

六、编译与部署

6.1 编译脚本

# build.sh
atc --om=add_custom.om \
    --framework=5 \
    --model=add_custom.pb \
    --output=add_custom \
    --soc_version=Ascend910

6.2 在 MindSpore 中调用

from op.add_custom import AddCustom

add_op = AddCustom()
a = ms.Tensor([1.0, 2.0, 3.0], dtype=ms.float32)
b = ms.Tensor([4.0, 5.0, 6.0], dtype=ms.float32)
c = add_op(a, b)
print(c)  # [5.0, 7.0, 9.0]

七、进阶：实现 GELU 激活函数

GELU 是 Transformer 中常用激活函数：
GELU(x)=x⋅Φ(x)≈x⋅0.5⋅(1+tanh(2/π​(x+0.044715x3)))

Ascend C 实现要点：

  • 使用 VecMul, VecAdd, VecTanh 组合
  • 预计算常数（√(2/π) ≈ 0.79788456）
void GeluCompute(TBuf<float>& dst, TBuf<float>& src, uint32_t size) {
    TBuf<float> tmp1, tmp2, tmp3;
    constexpr float coef = 0.044715f;
    constexpr float sqrt_2_over_pi = 0.79788456f;

    VecMul(tmp1, src, src, size);           // x^2
    VecMul(tmp1, tmp1, src, size);          // x^3
    VecMuls(tmp2, tmp1, coef, size);        // 0.044715 * x^3
    VecAdd(tmp2, tmp2, src, size);          // x + ...
    VecMuls(tmp2, tmp2, sqrt_2_over_pi, size); // √(2/π) * ...
    VecTanh(tmp3, tmp2, size);              // tanh(...)
    VecAdds(tmp3, tmp3, 1.0f, size);        // 1 + tanh
    VecMuls(tmp3, tmp3, 0.5f, size);        // 0.5 * (1 + tanh)
    VecMul(dst, src, tmp3, size);           // x * ...
}

八、总结与展望

Ascend C 为昇腾开发者提供了接近硬件的编程能力，虽学习曲线陡峭，但性能收益显著。未来，随着 CANN 和 MindSpore 的演进，Ascend C 将支持更多高级特性（如自动分块、AI 编译优化），降低开发门槛。

建议：

  • 从简单算子入手（Add、Relu）
  • 善用 Profiling 工具（msadvisor）
  • 关注华为官方文档与样例库

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

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

# c语言 # 开发语言
Logo
CANN开发者社区

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

更多推荐

cover

CANNBot + DeepSeek-V4 实操：30 分钟生成可达理论性能极限的 MXFP8 Matmul + Add 融合算子

avatar CANN开发者社区
cover

5月14号直播丨​多模态生成技术优化实践第二期--并行和Cache篇

avatar CANN开发者社区
cover

多模态生成技术优化实践第一期

avatar CANN开发者社区