一：入门篇 —— 从零搭建 Ascend C 开发环境并实现 VectorAdd 算子

# Ascend C 算子开发入门：手把手教你搭建环境并跑通第一个算子！

> 在昇腾 AI 生态中，**Ascend C** 是直接操控 AI Core 的高性能编程语言。本文将带你完成开发环境配置，并实现一个最基础的 `VectorAdd` 自定义算子，迈出昇腾算子开发的第一步！

## 一、为什么需要自定义算子？

虽然 CANN 提供了丰富的内置算子（如 ACL 库），但在以下场景仍需自定义：
- 新型激活函数（如 SwiGLU、GeGLU）
- 融合操作（如 RMSNorm + Scale）
- 特定模型中的稀疏或定制化计算

通过 Ascend C 编写 Kernel，可减少多次 Kernel Launch 与中间内存分配，显著提升推理性能。

## 二、开发环境准备（Ubuntu 20.04/22.04）

确保已安装 **CANN Toolkit ≥ 8.0**（推荐 8.1.RC1）：

```bash
# 1. 设置环境变量（关键！）
source /usr/local/Ascend/ascend-toolkit/set_env.sh

# 2. 验证编译器
which ascend-clang++
# 输出应类似：/usr/local/Ascend/ascend-toolkit/bin/ascend-clang++

若无昇腾设备，可使用 CPU 仿真模式 进行功能验证（需在 CMake 中开启 -DENABLE_CPU_RUN=ON）。

三、编写第一个 Ascend C 算子：VectorAdd

1. 目录结构

Text

编辑

1vector_add/
2├── kernel/
3│   └── vector_add_kernel.cpp
4├── host/
5│   └── vector_add_host.cpp
6├── CMakeLists.txt
7└── test/
8    └── test_vector_add.py

2. Kernel 侧实现（kernel/vector_add_kernel.cpp）

Cpp

编辑

1#include "kernel_operator.h"
2using namespace AscendC;
3
4extern "C" __global__ __aicore__ void vector_add(
5    GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLen) {
6    
7    constexpr int32_t TILE_SIZE = 8192; // 分块大小（需 ≤ UB 容量）
8    
9    auto ubX = AllocTensor<float>(TILE_SIZE);
10    auto ubY = AllocTensor<float>(TILE_SIZE);
11    auto ubZ = AllocTensor<float>(TILE_SIZE);
12
13    for (uint32_t i = 0; i < totalLen; i += TILE_SIZE) {
14        int32_t processLen = min(TILE_SIZE, static_cast<int32_t>(totalLen - i));
15        
16        // CopyIn: Global Memory → Unified Buffer
17        DataCopy(ubX, reinterpret_cast<float*>(x) + i, processLen);
18        DataCopy(ubY, reinterpret_cast<float*>(y) + i, processLen);
19        
20        // Compute: VecAdd
21        VecAdd(ubZ, ubX, ubY, processLen);
22        
23        // CopyOut: Unified Buffer → Global Memory
24        DataCopy(reinterpret_cast<float*>(z) + i, ubZ, processLen);
25    }
26
27    FreeTensor(ubX);
28    FreeTensor(ubY);
29    FreeTensor(ubZ);
30}

3. Host 侧调用（简化）

Cpp

编辑

1// 注册算子（实际需通过 OpReg 或 Python 绑定）
2void launch_vector_add(float* x, float* y, float* z, uint32_t len, aclrtStream stream) {
3    vector_add<<<1, 0, stream>>>(x, y, z, len);
4    aclrtSynchronizeStream(stream);
5}

四、编译与测试

使用官方构建脚本或 CMake 编译后，通过 Python 验证：

Python

编辑

1import numpy as np
2from custom_op import vector_add  # 假设已封装为 Python 接口
3
4x = np.random.rand(1024).astype(np.float32)
5y = np.random.rand(1024).astype(np.float32)
6z = vector_add(x, y)
7assert np.allclose(z, x + y, atol=1e-5), "结果不一致！"
8print("✅ VectorAdd 算子验证通过！")

五、常见问题

• Q：编译报错 “GM_ADDR not defined”？
  A：确保包含头文件 "kernel_operator.h" 并链接 Ascend C 运行时库。

• Q：如何调试 Kernel？
  A：使用 msopst 工具进行单算子测试，或启用 CPU 仿真模式。

六、小结

本文完成了 Ascend C 开发环境搭建与首个算子实现。下一步，我们将深入昇腾 AI Core 架构，理解高性能编程的核心思想。

🔗 下期预告：《Ascend C 架构篇：揭秘 AI Core 与三大编程范式》
👉 如果你觉得有帮助，欢迎点赞、收藏、关注!

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