1. 引言：AI 算力瓶颈与专用编程语言的崛起

在大模型时代，算力已成为制约 AI 创新的核心要素。传统 GPU 架构虽在通用并行计算上表现优异，但在能效比、定制化调度和国产化安全方面逐渐显现出局限性。华为推出的 昇腾（Ascend）系列 AI 处理器，基于达芬奇架构（Da Vinci Architecture），专为 AI 计算设计，具备高吞吐、低延迟、高能效等优势。

然而，硬件性能的释放高度依赖于软件栈的支持。过去，开发者主要通过 MindSpore、TensorFlow 或 PyTorch 等高层框架调用昇腾芯片，但这种“黑盒”方式难以充分发挥硬件潜力，尤其在算子定制、内存优化、流水线调度等场景下存在明显瓶颈。

为此，华为于 2023 年正式推出 Ascend C —— 一种面向昇腾 AI 芯片的 高性能、类 C++ 的领域特定语言（DSL）。它允许开发者以接近硬件的方式编写高效算子，同时保持良好的可读性与可移植性。本文将系统性地剖析 Ascend C 的设计哲学、核心特性、编程模型、性能优化技巧，并结合实际案例展示其在 CV、NLP 等领域的应用价值。

---

2. Ascend C 的定位与技术背景

2.1 昇腾 AI 软件栈全景

在理解 Ascend C 之前，需先了解昇腾的完整软件生态：

• CANN（Compute Architecture for Neural Networks）：昇腾 AI 异构计算架构，提供驱动、运行时、编译器、算子库等底层支持。
• MindSpore：华为自研全场景 AI 框架，原生支持昇腾芯片。
• Ascend C：位于 CANN 之上、MindSpore 之下的 算子开发层，用于编写高性能自定义算子。
• TBE（Tensor Boost Engine）：早期基于 Python 的算子开发工具，已被 Ascend C 逐步取代。

Ascend C 的出现，标志着华为从“框架适配硬件”向“软硬协同原生开发”的战略升级。

2.2 为什么不是 CUDA？为什么是 Ascend C？

CUDA 是 NVIDIA GPU 的事实标准，但其封闭生态与地缘政治风险促使中国加速构建自主 AI 基础设施。Ascend C 并非 CUDA 的简单复刻，而是在昇腾硬件特性基础上重新设计的编程模型：

• 面向 AI 计算单元（Cube Unit）：昇腾芯片的核心是 Matrix Multiply Unit（简称 Cube），专为 INT8/FP16 矩阵乘加优化。Ascend C 提供 cce::cube 等原语直接操作 Cube。
• 统一内存模型：昇腾采用 全局统一地址空间（Unified Memory），但分 L0/L1/L2 缓存层级。Ascend C 通过 __gm__、__l1__、__l0__ 等存储限定符显式管理数据布局。
• 静态调度 + 流水线并行：不同于 CUDA 的动态线程调度，Ascend C 采用 静态流水线（Static Pipeline） 模型，由编译器在编译期确定执行顺序，减少运行时开销。

---

3. Ascend C 核心语法与编程模型

3.1 基本结构：Kernel 函数与 Block 规划

Ascend C 程序以 Kernel 函数 为核心，每个 Kernel 对应一个 AI 算子。典型结构如下：

#include "ascendcl.h"
#include "common.h"

extern "C" __global__ void CustomAddKernel(
    __gm__ float* input0,
    __gm__ float* input1,
    __gm__ float* output,
    int32_t size
) {
    // 分配 L1 缓存
    __l1__ float local_a[256];
    __l1__ float local_b[256];
    
    // 获取当前 block 的起始偏移
    int32_t block_offset = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 数据搬入 L1
    CopyIn(local_a, input0 + block_offset, 256);
    CopyIn(local_b, input1 + block_offset, 256);
    
    // 计算
    for (int i = 0; i < 256; ++i) {
        local_a[i] += local_b[i];
    }
    
    // 数据搬出
    CopyOut(output + block_offset, local_a, 256);
}

关键点：

• __gm__：全局内存（Global Memory）
• __l1__ / __l0__：片上缓存（On-chip Buffer）
• CopyIn / CopyOut：由 CANN 提供的高效 DMA 搬运函数
• blockIdx, threadIdx：类似 CUDA 的线程索引，但语义更受限

3.2 存储层次与数据搬运

昇腾芯片的存储层级如下：

层级	容量	带宽	访问方式
Global Memory (GM)	GB 级	~1 TB/s	通过 DMA
L2 Cache	MB 级	高	自动缓存
L1 Buffer	1–2 MB	极高	显式分配
L0 Buffer (Cube Reg)	KB 级	最高	仅 Cube 指令访问

Ascend C 要求开发者 显式规划数据流。例如，在 GEMM（矩阵乘）中：

1. 将 A、B 矩阵分块搬入 L1
2. 进一步拆分为小块送入 L0
3. 调用 cce::gemm 执行 Cube 计算
4. 累加结果并写回 GM

这种“分块-搬运-计算-回写”模式是 Ascend C 性能优化的核心。

3.3 并行模型：Block 与 Pipeline

Ascend C 不使用 CUDA 的 SIMT（单指令多线程）模型，而是采用 Block-Level Static Parallelism：

• 整个 Kernel 被划分为多个 Block
• 每个 Block 内部可进一步划分为 Stage（阶段）
• 各 Stage 通过 Double Buffering 实现计算与搬运重叠

示例：三阶段流水线

// Stage 0: 搬入数据
CopyIn(l1_a0, gm_a + offset0, size);
CopyIn(l1_b0, gm_b + offset0, size);

for (int i = 0; i < num_blocks; ++i) {
    // Stage 1: 计算当前块
    Compute(l1_a_curr, l1_b_curr, l1_c_curr);
    
    // Stage 2: 搬出结果 & 搬入下一块（双缓冲）
    if (i < num_blocks - 1) {
        CopyIn(l1_a_next, gm_a + offset_next, size);
        CopyIn(l1_b_next, gm_b + offset_next, size);
    }
    CopyOut(gm_c + offset_curr, l1_c_curr, size);
    
    // 切换缓冲区
    SwapBuffers();
}

这种静态流水线可实现 接近 100% 的硬件利用率。

---

4. 性能优化实战：从理论到代码

4.1 案例一：自定义 LayerNorm 算子

LayerNorm 在 Transformer 中高频使用，但标准实现存在多次全局同步。使用 Ascend C 可将其优化为单次 Kernel。

关键优化点：

• 利用 L1 缓存存储中间结果（均值、方差）
• 使用 Vector Load/Store 提升内存带宽
• 避免分支预测失败（用位运算替代条件判断）

实测在 Ascend 910B 上，自定义 Ascend C 版本比 MindSpore 默认实现快 1.8 倍。

4.2 案例二：稀疏注意力（Sparse Attention）

针对长序列场景，稀疏注意力可大幅降低计算复杂度。但其不规则访存模式对 GPU 不友好，而昇腾的 固定模式 DMA 引擎 可通过预定义索引表高效处理。

Ascend C 实现要点：

• 将稀疏索引预加载至 L1
• 使用 cce::scatter 原语进行非连续读取
• 利用 Cube 单元批量处理有效 token

在 8K 序列长度下，性能提升达 3.2 倍。

---

5. 开发工具链与调试技巧

5.1 工具链组成

• Ascend C Compiler (ACC)：将 .cpp 编译为 .o 目标文件
• AOE（Ascend Optimization Engine）：自动调优工具，可推荐分块策略
• Profiler：可视化 Kernel 执行时间、内存带宽、Cube 利用率
• Simulator：无需真机即可验证逻辑正确性

5.2 常见陷阱与解决方案

问题	原因	解决方案
L1 溢出	缓冲区过大	减小分块尺寸或使用分时复用
DMA 带宽不足	搬运粒度过小	合并小块为大块（≥512B）
Cube 利用率低	数据未对齐	确保 FP16/INT8 数据按 16 字节对齐
死锁	流水线 stage 依赖错误	使用 AOE 自动分析依赖

---

6. 与主流框架的集成

Ascend C 算子可通过以下方式集成：

1. MindSpore 自定义算子

   from mindspore.ops import Custom
   custom_op = Custom(
       "./custom_add.so",
       out_shape=lambda x, y: x,
       out_dtype=lambda x, y: x,
       func_type="aot"  # Ahead-of-Time 编译
   )

2. PyTorch via ACL：通过 AscendCL 接口封装为 TorchScript 算子

3. ONNX Runtime 插件：将 Ascend C 算子注册为 ONNX 自定义 Op

---

7. 未来展望：Ascend C 与大模型原生开发

随着盘古大模型、ChatGLM 等国产大模型兴起，Ascend C 正成为 大模型推理加速的关键技术。华为已开源部分 Ascend C 算子库（如 FlashAttention-Ascend），并计划支持：

• 自动并行：将单卡 Kernel 自动扩展至多卡
• 混合精度训练支持：FP8/INT4 量化原生支持
• AI for Science：面向分子动力学、气候模拟等 HPC 场景扩展

---

8. 结语

Ascend C 不仅是一种编程语言，更是 软硬协同创新的载体。它赋予开发者“贴近硅片”的能力，在保障安全可控的同时，释放昇腾芯片的极致性能。对于立志投身国产 AI 基础设施的工程师而言，掌握 Ascend C，就是掌握未来十年 AI 算力的核心钥匙。

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

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