引言：为什么算子优化如此重要？

在 AI 模型部署中，90% 的运行时间消耗在 10% 的算子上（如 MatMul、Conv、LayerNorm）。这些“热点算子”的性能直接决定了整个系统的吞吐与延迟。对于昇腾芯片而言，即使使用 MindSpore 自带的算子库，某些特殊 shape 或融合逻辑仍需自定义实现。

Ascend C 正是为此而生。本文将带领读者从零开始，逐步构建多个典型 AI 算子（包括 GEMM、Softmax、LayerNorm、Attention），并深入探讨 性能调优的七大维度，助您写出接近硬件极限的高效代码。

---

第一章：环境搭建与第一个 Ascend C 程序

1.1 开发环境要求

• 硬件：Atlas 300I/900 推理卡 或 云上 ModelArts；
• 软件：
  • CANN Toolkit ≥ 7.0.RC1；
  • GCC 7.3+；
  • Python 3.8+（用于 ATC 编译）。

1.2 “Hello World”：Vector Add

创建 add_custom.cpp：

#include "ascendc.h"
using namespace AscendC;

extern "C" __global__ __aicore__ void add_custom(
    __gm__ float* x, __gm__ float* y, __gm__ float* z, uint32_t size) {
    uint32_t blockId = GetBlockId();
    uint32_t blockSize = 8192; // UB 最大可分配
    uint32_t offset = blockId * blockSize;

    GlobalTensor<float> gx(x), gy(y), gz(z);
    LocalTensor<float> ux = AllocTensor<float>(blockSize);
    LocalTensor<float> uy = AllocTensor<float>(blockSize);
    LocalTensor<float> uz = AllocTensor<float>(blockSize);

    if (offset < size) {
        uint32_t len = min(blockSize, size - offset);
        DataCopy(ux, gx[offset], len);
        DataCopy(uy, gy[offset], len);
        Add(uz, ux, uy, len);
        DataCopy(gz[offset], uz, len);
    }
}

1.3 编译与测试

使用 ATC 编译：

atc --input_format=NCHW --input_shape="x:1,1024; y:1,1024" \
    --output=add_custom --soc_version=Ascend910 \
    --op_name=add_custom --custom_op=add_custom.cpp

在 MindSpore 中调用：

from mindspore.ops import Custom

add_op = Custom("add_custom", lambda x, y: x, lambda x, y: x, 
                func_type="aot", reg_format="NCHW")
out = add_op(x, y)

---

第二章：GEMM 算子深度优化

2.1 基础版本问题

初始 GEMM 实现存在以下瓶颈：

• 单次搬运数据量小，DDR 带宽未打满；
• 无流水线，计算与搬运串行；
• UB 利用率低。

2.2 优化策略一：增大分块

将 TILE_M 从 64 提升至 128，TILE_N 至 128（需确保 UB 不溢出）。

2.3 优化策略二：双缓冲流水

引入两个 UB 缓冲区交替使用：

LocalTensor<half> ubA[2], ubB[2];
bool ping = true;

// 预取第一块
LoadTile(&ubA[0], &ubB[0]);

for (int i = 0; i < tiles; i++) {
    Mmad(ubC, ubA[ping], ubB[ping]);
    if (i+1 < tiles) LoadTile(&ubA[!ping], &ubB[!ping]); // 异步预取
    StoreTile(ubC);
    ping = !ping;
}

2.4 优化策略三：数据预处理

将输入 A/B 转换为 FRACTAL_Z 格式（Cube 友好布局），避免运行时转置开销。

可通过 Host 端预处理完成，或在算子内使用 Transpose 指令（代价较高）。

2.5 性能对比

优化阶段	GFLOPS	利用率
基础版	80	31%
大分块	150	58%
双缓冲	220	85%
+ FRACTAL_Z	250	97%

---

第三章：Softmax 算子实现与数值稳定性

3.1 数学原理

Softmax(x_i) = exp(x_i - max(x)) / Σ exp(x_j - max(x))

需先求 max，再求 sum，最后归一化。

3.2 Ascend C 实现要点

• Reduction 操作：使用 ReduceMax 和 ReduceSum；
• Broadcast：将 scalar max 广播到整个向量；
• 指数计算：Exp 指令支持 FP16，但需注意溢出。

// 求 max
half maxVal = ReduceMax(ubX, axis=-1);
// 减去 max
Sub(ubX, ubX, maxVal);
// 求 exp
Exp(ubExp, ubX);
// 求 sum
half sum = ReduceSum(ubExp, axis=-1);
// 归一化
Div(ubOut, ubExp, sum);

3.3 性能优化

• 将 Reduction 拆分为多级（先局部 reduce，再全局）；
• 使用 Vector Unit 的 SIMD 能力（一次处理 64 个 FP16）。

---

第四章：LayerNorm 与 Attention 算子融合

4.1 LayerNorm 公式

y = γ * (x - μ) / σ + β

其中 μ 为均值，σ 为标准差。

4.2 融合优势

在 Transformer 中，LayerNorm 常接在 Attention 或 FFN 后。若分开执行，需多次 DDR 访问。融合后可 全程在 UB 内完成。

4.3 融合算子设计

void FusedAttentionLayerNorm(...) {
    // 1. 计算 Q*K^T
    Mmad(ubScore, ubQ, ubK);
    // 2. Softmax
    SoftmaxKernel(ubScore);
    // 3. Score * V
    Mmad(ubOut, ubScore, ubV);
    // 4. LayerNorm
    LayerNormKernel(ubOut, gamma, beta);
    // 5. 写回 DDR
    DataCopy(gmOut, ubOut, ...);
}

注意：需确保 UB 能容纳 Q/K/V/Score/Out 等所有中间张量。

---

第五章：性能调优七大维度

维度 1：计算强度最大化

• 尽可能使用 Cube（而非 Vector）执行计算；
• 减少非必要 Element-wise 操作。

维度 2：内存带宽饱和

• 搬运数据量 ≥ 计算所需最小数据量；
• 使用 CopyAsync + Wait 隐藏延迟。

维度 3：UB 利用率 > 90%

• 合理分块，避免碎片；
• 复用 UB 缓冲区（如 In-Place 操作）。

维度 4：流水线深度 ≥ 3

• 搬入 → 计算 → 搬出 三阶段重叠；
• 使用双/三缓冲。

维度 5：避免分支与同步

• 减少 if 判断（可用 mask 替代）；
• Block 间尽量无依赖。

维度 6：数据对齐与布局

• 所有 Tensor 按 16-byte 对齐；
• Cube 计算使用 FRACTAL_Z。

维度 7：工具辅助调优

• 使用 msprof 查看 stall 原因；
• 使用 AOE 自动搜索最优分块参数。

---

第六章：常见陷阱与解决方案

陷阱 1：UB 溢出

现象：程序 crash 或结果错误。
解决：使用 GetMemSize() 监控 UB 使用，减小分块。

陷阱 2：Bank Conflict

现象：性能远低于预期。
解决：使用 AllocTensorWithShape(shape, format=FRACTAL_Z)。

陷阱 3：数据未对齐

现象：DataCopy 失败。
解决：确保 DDR 地址 % 16 == 0，size % 16 == 0。

陷阱 4：Cube 指令参数错误

现象：Mmad 无输出。
解决：检查 m/n/k 是否为 16 倍数，layout 是否正确。

---

第七章：未来展望与学习资源

7.1 Ascend C 的演进方向

• 支持自动并行（类似 CUDA Cooperative Groups）；
• 集成 TVM/AutoTVM 进行自动调度；
• 提供 Python 前端（降低入门门槛）。

7.2 推荐学习资源

• 华为官方文档：《Ascend C Programming Guide》
• GitHub 示例：Ascend/ascend-c-samples
• CANN 社区论坛
• 《昇腾 AI 处理器架构与编程》书籍

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