引言：AI 算力之争与自研芯片的崛起

近年来，随着人工智能（尤其是大模型）的飞速发展，对底层算力的需求呈指数级增长。传统 GPU 架构虽仍占据主流，但其通用性与能效比已逐渐难以满足特定场景下的极致性能需求。在此背景下，专用 AI 芯片（如华为昇腾 Ascend 系列）应运而生，成为国产算力自主可控的关键突破口。

华为昇腾（Ascend）系列 AI 处理器基于达芬奇架构，具备高吞吐、低延迟、高能效等优势，广泛应用于训练与推理场景。然而，要充分发挥其硬件潜力，离不开高效的软件栈支持。其中，Ascend C 作为华为推出的新一代面向昇腾 AI 芯片的高性能算子开发语言，正成为连接算法与硬件的核心桥梁。

本文将系统介绍 Ascend C 的设计哲学、核心特性、编程模型，并通过一个典型算子（如 GELU 激活函数）的开发流程，帮助读者掌握其使用方法，理解其如何实现“软硬协同、极致性能”。

一、什么是 Ascend C？

Ascend C 是华为为昇腾 AI 处理器（如 Ascend 910B、310P 等）量身打造的高性能 C++ 扩展语言。它并非一门全新的编程语言，而是基于标准 C++17/20，通过引入一系列内置函数（Intrinsics）、内存管理原语和并行编程模型，使开发者能够直接操作昇腾芯片的计算单元（如 AI Core 中的 Cube 单元、Vector 单元）和片上内存（如 Unified Buffer, UB），从而编写出高度优化的自定义算子（Custom Operator）。

核心目标：

• 极致性能：接近硬件理论峰值的计算效率。
• 开发友好：相比传统汇编或底层驱动接口，提供更高抽象层次。
• 可移植性：同一份 Ascend C 代码可在不同型号的昇腾芯片上运行（需适配）。
• 生态融合：无缝集成到 MindSpore、CANN（Compute Architecture for Neural Networks）等华为 AI 软件栈中。

二、Ascend C 的核心特性

1. 基于 Tile 的编程模型（Tiling Programming Model）

昇腾芯片采用“分块计算”思想。Ascend C 要求开发者将输入数据划分为多个 Tile（块），每个 Tile 的大小需匹配硬件资源（如 UB 容量）。这种模型强制开发者显式管理数据搬运与计算重叠，是实现高吞吐的关键。

// 示例：定义 Tile 大小
constexpr int32_t BLOCK_SIZE = 256;

2. 内存层级显式管理

昇腾芯片具有多级内存结构：

• Global Memory（GM）：片外 DDR，容量大但带宽有限。
• Unified Buffer（UB）：片上高速缓存，容量有限（通常几百 KB），但带宽极高。
• Local L1/L0 Cache：更靠近计算单元的缓存。

Ascend C 提供 CopyIn、CopyOut、DataCopy 等 API，要求开发者显式控制数据在 GM 与 UB 之间的搬运，避免隐式拷贝带来的性能损失。

3. 高性能计算原语（Intrinsics）

Ascend C 封装了大量针对昇腾硬件的向量化与矩阵运算指令，例如：

• vadd、vmul：向量加法/乘法（Vector 单元）
• mmad：矩阵乘加（Cube 单元，支持 FP16/BF16/INT8）
• vexp、vlog：超越函数加速

这些 Intrinsics 直接映射到底层硬件指令，避免了通用编译器优化不足的问题。

4. 多核并行与流水线调度

昇腾芯片包含多个 AI Core，每个 Core 可独立执行任务。Ascend C 支持通过 blockIdx、blockDim 实现多核并行；同时，通过 Double Buffering 和 Pipeline Scheduling 技术，隐藏数据搬运延迟，实现“计算”与“通信”重叠。

三、Ascend C 编程模型详解

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

1. Kernel 函数声明

使用 extern "C" 和特定属性声明，标识为昇腾内核函数。

extern "C" __global__ __aicore__ void gelu_custom_kernel(
    half* input, half* output, uint32_t total_size)
{
    // Kernel 逻辑
}

2. 数据搬运（Data Movement）

从 GM 加载数据到 UB，计算完成后写回 GM。

// 定义 UB 缓冲区
__ubuf__ half input_ub[BLOCK_SIZE];
__ubuf__ half output_ub[BLOCK_SIZE];

// 拷贝输入
DataCopy(input_ub, input + block_offset, BLOCK_SIZE);

// 计算...
// 拷贝输出
DataCopy(output + block_offset, output_ub, BLOCK_SIZE);

3. 计算逻辑（Compute Logic）

使用 Intrinsics 实现核心算法。

// GELU 近似：0.5 * x * (1 + tanh(sqrt(2/π) * (x + 0.044715 * x^3)))
for (int i = 0; i < BLOCK_SIZE; ++i) {
    half x = input_ub[i];
    half x3 = vmul(x, vmul(x, x)); // x^3
    half inner = vadd(x, vmul(0.044715_h, x3));
    half scaled = vmul(sqrt_2_over_pi, inner);
    half tanh_val = vtanh(scaled);
    half result = vmul(0.5_h, vmul(x, vadd(1.0_h, tanh_val)));
    output_ub[i] = result;
}

注：实际开发中需考虑向量化（如一次处理 16 个 half 元素）以提升效率。

4. 启动配置（Launch Configuration）

在 Host 端（CPU）调用时，需指定 Grid/Block 尺寸。

dim3 grid((total_size + BLOCK_SIZE - 1) / BLOCK_SIZE);
dim3 block(BLOCK_SIZE);
gelu_custom_kernel<<<grid, block>>>(input, output, total_size);

四、实战：使用 Ascend C 实现 GELU 算子

GELU（Gaussian Error Linear Unit）是 Transformer 等大模型中常用的激活函数。其数学表达为：

GELU(x)=x⋅Φ(x)≈0.5x(1+tanh(π2​​(x+0.044715x3)))

步骤 1：环境准备

• 安装 CANN Toolkit（>=7.0）
• 配置 Ascend C 编译器（aarch64-linux-gnu-g++ + Ascend C 插件）
• 创建项目目录结构

步骤 2：编写 Kernel 代码（gelu_kernel.cpp）

#include "kernel_operator.h"

using namespace AscendC;

constexpr int32_t BLOCK_SIZE = 256;
constexpr float SQRT_2_OVER_PI = 0.7978845608028654f;

class GELUKernel {
public:
    __aicore__ inline void Init(GM_ADDR input, GM_ADDR output, uint32_t totalSize) {
        this->input_gm.SetGlobalBuffer((__gm__ half*)input, totalSize);
        this->output_gm.SetGlobalBuffer((__gm__ half*)output, totalSize);
        this->tileNum = (totalSize + BLOCK_SIZE - 1) / BLOCK_SIZE;
    }

    __aicore__ inline void Process() {
        for (uint32_t i = 0; i < tileNum; ++i) {
            Pipe pipe;
            pipe.InitBuffer(input_ub, 1, BLOCK_SIZE * sizeof(half));
            pipe.InitBuffer(output_ub, 1, BLOCK_SIZE * sizeof(half));

            // Load
            DataCopy(input_ub[0], input_gm[i * BLOCK_SIZE], BLOCK_SIZE);
            pipe.Drain();

            // Compute
            for (int j = 0; j < BLOCK_SIZE; ++j) {
                half x = input_ub[0][j];
                half x3 = vmul(x, vmul(x, x));
                half inner = vadd(x, vmul(0.044715_h, x3));
                half scaled = vmul(SQRT_2_OVER_PI, inner);
                half tanh_val = vtanh(scaled);
                output_ub[0][j] = vmul(0.5_h, vmul(x, vadd(1.0_h, tanh_val)));
            }

            // Store
            DataCopy(output_gm[i * BLOCK_SIZE], output_ub[0], BLOCK_SIZE);
            pipe.Drain();
        }
    }

private:
    GlobalTensor<half> input_gm;
    GlobalTensor<half> output_gm;
    TPipe pipe;
    TBuf<half> input_ub[1];
    TBuf<half> output_ub[1];
    uint32_t tileNum;
};

extern "C" __global__ __aicore__ void gelu_custom_kernel(
    half* input, half* output, uint32_t totalSize) {
    GELUKernel kernel;
    kernel.Init(input, output, totalSize);
    kernel.Process();
}

步骤 3：编译与集成

使用 aoe 或 atc 工具链编译为 .o 文件，再链接到 MindSpore 自定义算子中。

步骤 4：性能验证

在 Ascend 910B 上测试，相比 PyTorch 默认实现，性能提升可达 2–3 倍，且内存占用更低。

五、Ascend C 的优势与挑战

优势：

• 极致性能：直接操控硬件，避免框架开销。
• 确定性：无动态调度，性能可预测。
• 生态支持：与 MindSpore 深度集成，支持自动微分、图优化。

挑战：

• 学习曲线陡峭：需理解昇腾硬件架构。
• 调试困难：缺乏成熟的 IDE 调试支持。
• 可移植性限制：代码强依赖昇腾平台。

结语

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