---

引言：为什么需要 Ascend C？

随着人工智能技术的迅猛发展，模型规模日益庞大，对底层硬件计算能力提出了前所未有的挑战。华为昇腾（Ascend）系列 AI 处理器正是在这一背景下应运而生，其以高能效比、高吞吐量和强大的 AI 专用计算单元（如 AI Core、Vector Core）著称。然而，要充分发挥昇腾芯片的性能潜力，仅靠高层框架（如 MindSpore、TensorFlow）是远远不够的——开发者往往需要深入到底层，编写高度优化的自定义算子。

为此，华为推出了 Ascend C —— 一种面向昇腾 AI 处理器的高性能 C++ 扩展编程语言。它并非传统意义上的“新语言”，而是基于 C++17 标准，结合昇腾硬件架构特性，提供了一套丰富的内置函数（Intrinsic）、内存管理机制和并行编程模型，使开发者能够以接近硬件的方式编写高效、可移植的 AI 算子。

本文将系统性地介绍 Ascend C 的核心概念、编程模型、内存管理、数据搬运、向量化计算等关键技术，并通过多个完整代码示例（包括 Vector Add、Matrix Multiply、Custom ReLU 等），带领读者从零开始掌握 Ascend C 开发流程。全文约 6500 字，适合具备 C++ 基础和一定 AI 背景的开发者阅读。

---

一、Ascend C 概述与开发环境搭建

1.1 什么是 Ascend C？

Ascend C 是华为为昇腾 AI 处理器（如 Ascend 910B）量身定制的高性能编程接口。它运行在 Device 端（即 AI Core 上），用于实现用户自定义的算子逻辑。其主要特点包括：

• 贴近硬件：直接操作昇腾芯片的计算单元（Cube Unit、Vector Unit）和片上内存（Unified Buffer, UB）。
• 自动流水线调度：通过 Pipe 机制实现计算与数据搬运的重叠，最大化硬件利用率。
• 强类型安全：基于 C++ 模板和类型系统，避免低级错误。
• 与 CCE（Compute Capability Engine）深度集成：编译后生成高效的二进制指令（.o 文件），供 Host 端调用。

注意：Ascend C 并非运行在 CPU 上，而是编译后部署到昇腾 NPU 的 Device 端执行。

1.2 开发环境准备

要进行 Ascend C 开发，需安装以下组件：

• CANN（Compute Architecture for Neural Networks）：华为昇腾 AI 软件栈，包含驱动、固件、编译器（aicpu-cce）、调试工具等。
• Ascend C SDK：通常随 CANN 一起安装，位于 /usr/local/Ascend/ascend-toolkit/latest/ascend-c/。
• 支持的 IDE：推荐使用 VS Code + Ascend 插件，或直接使用命令行编译。
• 硬件：昇腾 910/310 系列 AI 加速卡（或 Atlas 服务器）。

验证环境：

npu-smi info  # 查看 NPU 状态

Ascend C 项目结构示例：

my_custom_op/
├── src/
│   └── kernel/
│       ├── add_custom.cpp    # Ascend C 算子实现
│       └── add_custom.h
├── CMakeLists.txt
└── build.sh

---

二、Ascend C 核心编程模型

Ascend C 的编程模型围绕 “分块-流水-并行” 展开，核心组件包括：

• Global Memory（GM）：设备全局内存，容量大但带宽有限。
• Unified Buffer（UB）：片上高速缓存，容量小（通常 2MB/核）但带宽极高。
• Core Pipeline（Pipe）：用于协调数据搬运（Load/Store）与计算（Compute）的流水线。
• AI Core 内部单元：
  • Scalar Core：处理标量和控制流。
  • Vector Core：处理向量运算（如加法、乘法、激活函数）。
  • Cube Core：专用于矩阵乘（GEMM），支持 INT8/FP16 混合精度。

开发者需将输入数据从 GM 搬入 UB，再在 UB 上进行计算，最后将结果写回 GM。

2.1 基本程序结构

一个典型的 Ascend C 算子由以下部分组成：

#include "acl/acl_base.h"
#include "ascendc.h"
#include "common.h"

using namespace AscendC;

// 定义算子入口函数
extern "C" __global__ __aicore__ void add_custom(
    GlobalTensor<float> input1,
    GlobalTensor<float> input2,
    GlobalTensor<float> output,
    uint32_t totalLength) {

    // 1. 初始化 Pipe
    Pipe pipe;
    pipe.InitBuffer();

    // 2. 创建 LocalTensor（UB 中的张量）
    constexpr int32_t BUFFER_NUM = 2;
    LocalTensor<float> localInput1 = pipe.AllocTensor<float>(BUFFER_NUM);
    LocalTensor<float> localInput2 = pipe.AllocTensor<float>(BUFFER_NUM);
    LocalTensor<float> localOutput = pipe.AllocTensor<float>(BUFFER_NUM);

    // 3. 分块处理
    const int32_t TILE = 256; // 每次处理 256 个 float
    for (int32_t i = 0; i < totalLength; i += TILE) {
        int32_t processLen = (totalLength - i) > TILE ? TILE : (totalLength - i);

        // Load 数据到 UB
        DataCopy(localInput1, input1[i], processLen);
        DataCopy(localInput2, input2[i], processLen);

        // 同步：确保数据加载完成
        pipe.DrainAll();

        // 计算：向量加法
        Add(localOutput, localInput1, localInput2, processLen);

        // Store 结果回 GM
        DataCopy(output[i], localOutput, processLen);

        // 同步：确保写回完成
        pipe.DrainAll();
    }

    // 4. 释放资源
    pipe.FreeTensor(localInput1);
    pipe.FreeTensor(localInput2);
    pipe.FreeTensor(localOutput);
}

关键点说明：

• __global__ __aicore__：标识该函数将在 AI Core 上执行。
• GlobalTensor：指向 GM 的张量。
• LocalTensor：指向 UB 的张量。
• Pipe：管理 UB 缓冲区和流水线。
• DataCopy：封装了 DMA 搬运操作。
• Add：Vector Core 提供的向量加法 intrinsic。

---

三、内存管理与数据搬运

3.1 Unified Buffer（UB）管理

UB 是性能关键。Ascend C 通过 Pipe 对象管理 UB 的分配与回收。每个 Pipe 默认提供两个缓冲区（Double Buffering），用于隐藏数据搬运延迟。

Pipe pipe;
pipe.InitBuffer(); // 初始化默认大小的 UB

// 或指定大小（单位：字节）
pipe.InitBuffer(1024 * 1024); // 1MB

分配 LocalTensor 时，需指定缓冲区数量（通常为 2）：

LocalTensor<float> buf = pipe.AllocTensor<float>(2); // 双缓冲

3.2 DataCopy 机制

DataCopy 是 Ascend C 中最常用的数据搬运函数，其原型为：

void DataCopy(DstTensor dst, SrcTensor src, int32_t count);

• 支持 GM ↔ UB、UB ↔ UB 之间的拷贝。
• 自动处理地址对齐和突发传输（Burst Transfer）。
• 非阻塞调用，需配合 pipe.Drain() 或 pipe.DrainAll() 同步。

示例：分块加载大张量

const int32_t TILE_SIZE = 512; // 512 floats = 2KB
for (int32_t offset = 0; offset < totalSize; offset += TILE_SIZE) {
    int32_t len = min(TILE_SIZE, totalSize - offset);
    DataCopy(localBuf[0], globalInput[offset], len);
    pipe.Drain(); // 等待当前 copy 完成
    // ... 计算 ...
}

⚠️ 注意：DataCopy 的 count 参数单位是 元素个数，不是字节数！

---

四、向量化计算与 Intrinsic 函数

Ascend C 提供了丰富的 Intrinsic 函数，直接映射到 Vector Core 指令。常见操作包括：

操作	函数	说明
向量加	Add(dst, a, b, len)	dst = a + b
向量乘	Mul(dst, a, b, len)	dst = a * b
标量加	Adds(dst, src, scalar, len)	dst = src + scalar
ReLU	Active(dst, src, len, RELU)	ReLU 激活
最大值	ReduceMax(dst, src, len)	向量规约求最大值

4.1 实战：自定义 ReLU 算子

虽然 Ascend 已内置 ReLU，但我们可以手动实现以理解流程：

extern "C" __global__ __aicore__ void custom_relu(
    GlobalTensor<float> input,
    GlobalTensor<float> output,
    uint32_t totalLength) {

    Pipe pipe;
    pipe.InitBuffer();

    LocalTensor<float> inBuf = pipe.AllocTensor<float>(2);
    LocalTensor<float> outBuf = pipe.AllocTensor<float>(2);

    const int32_t TILE = 256;
    for (uint32_t i = 0; i < totalLength; i += TILE) {
        uint32_t len = (totalLength - i) > TILE ? TILE : (totalLength - i);

        DataCopy(inBuf, input[i], len);
        pipe.DrainAll();

        // 方法1：使用内置 Active
        // Active(outBuf, inBuf, len, RELU);

        // 方法2：手动实现（max(x, 0)）
        LocalTensor<float> zero = pipe.AllocTensor<float>(1);
        Fill(zero, 0.0f, len); // 填充零向量
        Max(outBuf, inBuf, zero, len); // out = max(in, 0)
        pipe.FreeTensor(zero);

        DataCopy(output[i], outBuf, len);
        pipe.DrainAll();
    }

    pipe.FreeTensor(inBuf);
    pipe.FreeTensor(outBuf);
}

此处展示了两种实现方式：直接调用 Active，或用 Max + Fill 组合实现。后者虽冗余，但有助于理解底层机制。

---

五、矩阵乘（GEMM）与 Cube Core 编程

对于深度学习中最核心的 GEMM 操作，Ascend C 提供了 Matmul intrinsic，利用 Cube Core 实现高性能矩阵乘。

5.1 Matmul 接口

void Matmul(
    LocalTensor<DST_DTYPE> dst,
    LocalTensor<FACTOR_DTYPE> a,
    LocalTensor<FACTOR_DTYPE> b,
    LocalTensor<DST_DTYPE> bias, // 可选
    bool isBias = false,
    MM_MODE mode = MM_MODE::MATMUL);

• 输入要求：
  • A: [M, K]，B: [K, N]
  • 必须按 tiling 格式 存储（如 fractal Z 格式）
• 精度支持：FP16、INT8、BF16 等

5.2 实战：FP16 矩阵乘

extern "C" __global__ __aicore__ void matmul_fp16(
    GlobalTensor<half> inputA,
    GlobalTensor<half> inputB,
    GlobalTensor<half> outputC,
    uint32_t M, uint32_t N, uint32_t K) {

    Pipe pipe;
    pipe.InitBuffer();

    // 分块大小（需满足 Cube 要求）
    constexpr int32_t BLOCK_M = 16;
    constexpr int32_t BLOCK_N = 16;
    constexpr int32_t BLOCK_K = 16;

    // 分配 UB 缓冲区
    LocalTensor<half> aTile = pipe.AllocTensor<half>(2, BLOCK_M * BLOCK_K);
    LocalTensor<half> bTile = pipe.AllocTensor<half>(2, BLOCK_K * BLOCK_N);
    LocalTensor<half> cTile = pipe.AllocTensor<half>(2, BLOCK_M * BLOCK_N);

    for (int32_t m = 0; m < M; m += BLOCK_M) {
        for (int32_t n = 0; n < N; n += BLOCK_N) {
            // 初始化 C 为零
            Fill(cTile, static_cast<half>(0.0), BLOCK_M * BLOCK_N);

            for (int32_t k = 0; k < K; k += BLOCK_K) {
                // Load A[m:m+BLOCK_M, k:k+BLOCK_K]
                // Load B[k:k+BLOCK_K, n:n+BLOCK_N]
                // 注意：实际需处理地址偏移和格式转换

                DataCopy(aTile, inputA[m * K + k], BLOCK_M * BLOCK_K);
                DataCopy(bTile, inputB[k * N + n], BLOCK_K * BLOCK_N);
                pipe.DrainAll();

                // 执行矩阵乘：cTile += aTile * bTile
                Matmul(cTile, aTile, bTile, cTile, false, MM_MODE::MATMUL);
                pipe.DrainAll();
            }

            // Store C[m:m+BLOCK_M, n:n+BLOCK_N]
            DataCopy(outputC[m * N + n], cTile, BLOCK_M * BLOCK_N);
            pipe.DrainAll();
        }
    }

    pipe.FreeTensor(aTile);
    pipe.FreeTensor(bTile);
    pipe.FreeTensor(cTile);
}

⚠️ 注意：实际使用中，输入矩阵通常需预转为 fractal Z 格式，否则性能极差。华为提供了 FormatHelper 工具辅助转换。

---

六、性能优化技巧

6.1 双缓冲（Double Buffering）

通过双缓冲隐藏数据搬运延迟：

// Buffer 0 用于计算，Buffer 1 用于加载下一块数据
DataCopy(localBuf[1], nextInput, len); // 异步加载
Compute(localBuf[0]);                  // 同时计算
pipe.Switch();                         // 切换缓冲区

6.2 向量化对齐

确保数据地址按 32 字节对齐，避免 bank conflict：

// 使用 alignas(32) 或编译器指令
__attribute__((aligned(32))) float data[256];

6.3 减少 UB 访问冲突

避免多个 core 同时访问同一 UB bank。可通过数据分片或 padding 解决。

6.4 使用内置高性能算子

优先使用 Active、Reduce、Matmul 等内置函数，而非手动循环。

---

七、调试与性能分析

7.1 日志输出

Ascend C 支持 printf（仅限 Scalar Core）：

if (blockIdx.x == 0 && threadIdx.x == 0) {
    printf("Debug: offset=%d\n", offset);
}

7.2 Profiling 工具

使用 Ascend Profiler 或 msadvisor 分析：

• UB 利用率
• Pipe 流水线效率
• 计算单元饱和度

---

八、完整工程示例：Vector Add 算子集成到 MindSpore

8.1 编写 Ascend C 算子（add_custom.cpp）

（见第二节代码）

8.2 编写 Host 端注册代码（Python）

import mindspore as ms
from mindspore.ops import Custom

def add_custom(input1, input2):
    output = ms.numpy.zeros_like(input1)
    func = Custom(
        "./add_custom.so",
        out_shape=lambda x, y: x,
        out_dtype=lambda x, y: x,
        func_type="aot"
    )
    return func(input1, input2, output, input1.size)

8.3 编译脚本（build.sh）

#!/bin/bash
ATC --mode=op \
    --input_shape="input1:1024;input2:1024;output:1024" \
    --op_name=add_custom \
    --source_file=./src/kernel/add_custom.cpp \
    --out=./add_custom.so

---

九、总结与展望

Ascend C 作为昇腾生态的核心编程接口，为 AI 算子开发者提供了接近硬件的控制能力。通过合理利用 UB、Pipe 流水线、Vector/Cube Core intrinsic，开发者可实现远超框架默认算子的性能。

未来，随着 CANN 版本迭代，Ascend C 将进一步简化编程模型（如自动 tiling、自动格式转换），并支持更多数据类型（如 FP8、INT4）。对于追求极致性能的 AI 工程师而言，掌握 Ascend C 已成为一项关键技能。

建议学习路径：

1. 熟悉 C++17 和模板编程
2. 阅读《Ascend C Programming Guide》官方文档
3. 在 Atlas 服务器上实践 Vector Add / Matmul
4. 尝试优化 ResNet / BERT 中的瓶颈算子

---

参考文献

1. Huawei CANN Documentation – Ascend C Programming Guide (v7.0.RC1)
2. 《昇腾 AI 处理器架构与编程》 – 华为技术有限公司
3. MindSpore Custom Operator Development Guide
4. GitHub: Ascend Samples (https://github.com/Ascend/samples)

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

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

---