有个朋友想给昇腾写自定义算子，问我应该从哪里入手。
我告诉他：“那你得先学 Ascend C。”
他第一反应是：“算子开发语言？不是用 C++ 吗？”

这可能是很多人的误区。昇腾的算子确实可以用 C++ 写（通过 TBE 或 ACL），但 CANN 专门推出了一套 Ascend C 的编程范式。它比纯 C++ 更适合算子场景——向量化、内存管理、流水线，这些在 Ascend C 里都有原生支持，能帮你避开无数坑。

Ascend C 是昇腾 CANN 推出的专用算子程序开发语言，位于 CANN 五层架构的第一层——昇腾计算语言层。它是通往昇腾硬件极限性能的“金钥匙”。

---

一、Ascend C 是什么？核心定位

• 官方定义：面向昇腾 AI 处理器的高性能算子开发语言，基于 C/C++ 标准，提供了多层级 API。
• 仓库地址：https://atomgit.com/cann/asc-devkit
• 核心价值：把复杂的硬件细节（如数据搬运、双缓冲、流水线）封装成语言级支持，让开发者专注于算法逻辑。

为什么不用纯 C++？

深度学习算子的核心计算有两个特点：

1. 向量化 (SIMD/SIMT)：同一份逻辑要并行跑在成千上万个数据上。
2. 数据局部性：必须精心管理数据在 GM（显存）、L1（片上缓存）、寄存器之间的搬运。

用纯 C++ 写，你需要手动处理所有底层细节，代码极易出错且难以优化。Ascend C 把这些封装好了，你只需要告诉它“做什么”，它负责“怎么做”。

---

二、Ascend C 核心概念解析

理解 Ascend C，只需掌握四个核心概念。

1. 双计算单元：AI Core vs AI Vector

昇腾芯片有两个主要计算单元，Ascend C 同时支持：

单元	全称	用途	特点
AI Core	矩阵计算单元	大矩阵乘 (MatMul)、卷积 (Conv)	吞吐极高，适合“大活”
AI Vector	矢量计算单元	逐元素操作 (Add, Relu)、规约 (Reduce)	灵活性强，适合“小活”

示例：使用 AI Vector 进行加法

#include "kernel_operator.h"

class VecAddKernel {
public:
    __aicore__ inline void Process(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
        // 每次搬 256 个 float 到 L1 缓存
        LocalTensor<float> xLocal = x_gm.GetCubeToL1Ptr<float>(256);
        LocalTensor<float> yLocal = y_gm.GetCubeToL1Ptr<float>(256);
        LocalTensor<float> zLocal = z_gm.GetCubeToL1Ptr<float>(256);
        
        // 调用 AI Vector 指令：256 路并行加法
        Add(zLocal, xLocal, yLocal, 256);
        
        // 结果写回 GM
        z_gm.GetL1ToCubePtr<float>(zLocal, 256);
    }
};

2. 编程模型：TPipe → TQueues → Tensors

Ascend C 的编程模型是分层的，类似工厂流水线：

TPipe (总管道)
  └─ TQueues (任务队列)
        └─ Tensors (数据张量)
              └─ 计算逻辑

这种结构天然支持流水线并发和资源调度。

标准框架示例：

#include "kernel_operator.h"

class MyKernel {
public:
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
        // 初始化全局指针
        this->x_gm.SetGlobalBuffer((__gm__ float*)x);
        this->y_gm.SetGlobalBuffer((__gm__ float*)y);
        this->z_gm.SetGlobalBuffer((__gm__ float*)z);
    }
    
    __aicore__ inline void Process() {
        // 主体计算逻辑
    }
};

// 入口函数：extern "C" 确保链接正确
extern "C" __global__ __open_acl_op__ void my_kernel(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
    MyKernel op;
    op.Init(x, y, z);
    op.Process();
}

3. 内存层级与数据搬运

昇腾的内存层级决定了性能上限。Ascend C 的核心工作就是管理这三层数据的搬运：

GM (Global Memory)  ← 设备显存 (HBM)，容量大，速度慢
  ↑  (DMA 搬运)
L1 (Level 1 Cache)  ← 片上缓存 (SRAM)，容量小，速度极快
  ↑  (寄存器传输)
寄存器               ← 计算单元内部，速度最快

关键原则：减少 GM 访问次数，最大化 L1 利用率。

__aicore__ inline void Process() {
    // 循环分块处理
    for (int i = 0; i < totalSize; i += 128) {
        // GM -> L1：搬入数据块
        LocalTensor<float> localX = x_gm.GetCubeToL1Ptr<float>(128);
        
        // 在 L1 上进行高速计算
        // ...
        
        // L1 -> GM：写回结果
        z_gm.GetL1ToCubePtr<float>(localZ, 128);
    }
}

4. 双缓冲技术 (Double Buffering)

为了掩盖数据搬运的延迟，Ascend C 支持双缓冲技术：一边搬运下一块数据，一边计算当前块数据。这是实现计算与通信重叠的关键。

双缓冲示例：

__aicore__ inline void Process() {
    // 定义两个缓冲区，交替使用
    LocalTensor<float> bufA[2], bufB[2], bufZ[2];
    
    // 预加载第一块
    bufA[0] = x_gm.GetCubeToL1Ptr<float>(BLOCK_SIZE);
    bufB[0] = y_gm.GetCubeToL1Ptr<float>(BLOCK_SIZE);
    
    for (int i = 0; i < totalBlocks; i++) {
        int cur = i % 2;      // 当前块索引
        int nxt = (i + 1) % 2; // 下一块索引
        
        // 当前块：计算
        Add(bufZ[cur], bufA[cur], bufB[cur], BLOCK_SIZE);
        z_gm.GetL1ToCubePtr<float>(bufZ[cur], BLOCK_SIZE);
        
        // 下一块：预加载 (与计算重叠)
        if (i + 1 < totalBlocks) {
            bufA[nxt] = x_gm.GetCubeToL1Ptr<float>(BLOCK_SIZE);
            bufB[nxt] = y_gm.GetCubeToL1Ptr<float>(BLOCK_SIZE);
        }
    }
}

效果：计算时间 + 搬运时间 ≈ max(计算时间，搬运时间)。如果两者相当，效率提升近 100%。

---

三、开发环境搭建与项目结构

1. 安装 CANN 工具链

首先下载并安装 Ascend Development Kit (Ascend Toolkit)：

# 1. 下载 CANN 工具链 (从昇腾官网或联系华为获取)
# 2. 安装
./ascend-toolkit-*.run --full

# 3. 配置环境变量
source /usr/local/Ascend/ascend-toolkit/set_env.sh

# 4. 验证安装
ascendc -v
# 输出版本信息即成功

2. 标准项目结构

一个规范的 Ascend C 项目应包含：

my_operator/
├── CMakeLists.txt       # 编译脚本
├── kernel_meta.json     # 算子元信息（名称、输入输出类型等）
├── src/
│   ├── kernel_operator.h    # 头文件
│   └── kernel_operator.cpp   # 核心实现
└── test/
    ├── ut_test.cpp       # 单元测试 (Host 端验证)
    └── run_ut.sh         # 测试脚本

3. 编译与部署

# 创建构建目录
mkdir build && cd build

# 配置编译选项
cmake .. -DCMAKE_BUILD_TYPE=Release

# 编译
make

# 生成 .om 文件 (离线模型)
# 通常由 CANN 工具链自动处理，或通过 omc 命令转换

---

四、实战案例：手写一个 Fused MatMul+Bias+Relu

这是一个典型的 Transformer 前馈网络层，我们尝试用 Ascend C 将其融合为一个算子。

1. 定义元信息 (kernel_meta.json)

{
  "op_name": "FusedMatMulBiasRelu",
  "version": 1,
  "input": [
    {"name": "input_x", "type": "float16", "shape": ["N", "K"]},
    {"name": "weight_w", "type": "float16", "shape": ["K", "M"]},
    {"name": "bias_b", "type": "float16", "shape": ["M"]}
  ],
  "output": [
    {"name": "output_y", "type": "float16", "shape": ["N", "M"]}
  ]
}

2. 实现核心逻辑 (kernel_operator.cpp)

#include "kernel_operator.h"

class FusedMatMulBiasRelu {
private:
    LocalTensor<float16_t> x_local;
    LocalTensor<float16_t> w_local;
    LocalTensor<float16_t> b_local;
    LocalTensor<float16_t> y_local;
    
    // 分块大小
    static constexpr int TILE_N = 256;
    static constexpr int TILE_K = 256;
    static constexpr int TILE_M = 256;

public:
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR w, GM_ADDR b, GM_ADDR y) {
        x_gm.SetGlobalBuffer((__gm__ float16_t*)x);
        w_gm.SetGlobalBuffer((__gm__ float16_t*)w);
        b_gm.SetGlobalBuffer((__gm__ float16_t*)b);
        y_gm.SetGlobalBuffer((__gm__ float16_t*)y);
    }

    __aicore__ inline void Process() {
        // 1. 加载 Bias 到 L1 (复用)
        b_local = b_gm.GetCubeToL1Ptr<float16_t>(TILE_M);
        
        // 2. 分块计算
        for (int n = 0; n < N; n += TILE_N) {
            for (int m = 0; m < M; m += TILE_M) {
                // 预加载 X 和 W 块
                x_local = x_gm.GetCubeToL1Ptr<float16_t>(TILE_N * TILE_K);
                w_local = w_gm.GetCubeToL1Ptr<float16_t>(TILE_K * TILE_M);
                
                // 3. 执行 MatMul (AI Core)
                // 这里调用内置的 MatMul 指令
                MatMul(y_local, x_local, w_local, TILE_N, TILE_K, TILE_M);
                
                // 4. 加 Bias (AI Vector)
                Add(y_local, y_local, b_local, TILE_M);
                
                // 5. 激活函数 Relu (AI Vector)
                Relu(y_local, y_local, TILE_M);
                
                // 6. 写回 GM
                y_gm.GetL1ToCubePtr<float16_t>(y_local, TILE_N * TILE_M);
            }
        }
    }
};

extern "C" __global__ __open_acl_op__ void fused_matmul_bias_relu(
    GM_ADDR x, GM_ADDR w, GM_ADDR b, GM_ADDR y) {
    FusedMatMulBiasRelu op;
    op.Init(x, w, b, y);
    op.Process();
}

3. 编写单元测试 (ut_test.cpp)

#include "gtest/gtest.h"
#include "kernel_operator.h"

TEST(FusedMatMulTest, BasicCheck) {
    // 准备 Host 端数据
    std::vector<float16_t> h_x(N*K);
    std::vector<float16_t> h_w(K*M);
    std::vector<float16_t> h_b(M);
    std::vector<float16_t> h_y_ref(N*M);
    
    // ... 填充数据 ...
    
    // 调用 Ascend C 算子 (需配合 ACL 接口)
    // aclrtSetDevice(...)
    // my_kernel(...)
    
    // 对比结果
    // EXPECT_NEAR(h_y_host[i], h_y_ref[i], epsilon);
}

---

五、常见问题排查 (FAQ)

Q1: 编译报错 undefined reference to 'GetCubeToL1Ptr'

• 原因：头文件未包含或版本不匹配。
• 解决：确保 #include "kernel_operator.h" 且 CANN 版本与工具链一致。检查 CMakeLists.txt 是否正确链接了 libascendc.a。

Q2: 运行时 OOM (Out of Memory)

• 原因：L1 缓存分配过大，超过了硬件限制（通常 1MB-2MB）。
• 解决：减小 BLOCK_SIZE 或 TILE 参数。检查是否有多处重复分配。

Q3: 性能不如预期

• 原因：数据搬运频繁，没有利用双缓冲；或者计算单元空闲等待数据。
• 解决：
  1. 启用 双缓冲 技术。
  2. 使用 Profiling 工具分析 L1 带宽利用率。
  3. 调整 TILE 大小以匹配 AI Core 的最佳粒度。

Q4: 浮点精度不一致

• 原因：FP16 累加时溢出或下溢。
• 解决：在累加过程中使用 FP32 中间变量，最后转回 FP16。

---

六、进阶：如何成为社区贡献者？

如果你已经掌握了基础，并希望为昇腾生态做贡献：

1. 阅读源码：深入研读 asc-devkit 中的示例代码，特别是 matmul 和 softmax 的实现。
2. 提交算子：
   • 按照 kernel_meta.json 规范定义你的算子。
   • 编写完整的单元测试（覆盖边界情况）。
   • 提交 Pull Request 到 AtomGit 仓库。
3. 参与讨论：加入昇腾开发者社区，解答他人关于 Ascend C 的问题。

---

七、总结

Ascend C 是连接开发者与昇腾硬件的桥梁。它不仅仅是另一种 C++ 方言，更是一套针对 AI 计算特性优化的编程范式。

• 对于初学者：它是学习昇腾硬件原理的最佳切入点。
• 对于资深开发者：它是突破性能瓶颈、释放硬件潜力的终极武器。
• 对于社区：每一个优秀的 Ascend C 算子，都是对昇腾生态的宝贵贡献。

下一步行动：

1. 安装 CANN 工具链。
2. 克隆 asc-devkit 仓库。
3. 尝试复现上面的 FusedMatMul 案例。
4. 动手写一个你自己的自定义算子！

代码改变世界，而 Ascend C 让它在昇腾上跑得更快。