---
tags: [昇腾, AscendC, 算子开发]
doc_type: knowledge-note
aliases: [Ascend算子基础]
abstract: "以 Hello World 和 Add 为主线，理解 Ascend C 的执行模型、内存搬运、分块和同步。"
finished: false
creation: 2026-10-15
返回目录: "[[计算机系/昇腾与CANN/大模型与算子专题学习导航|大模型与算子专题学习导航]]"
下一篇: "[[计算机系/昇腾与CANN/CMake与Ascend C编译入门|CMake与Ascend C编译入门]]"
前置知识: "[[计算机系/昇腾与CANN/异构计算基础与CANN架构整理|异构计算基础与CANN架构整理]]"
后续学习: ["[[计算机系/昇腾与CANN/CANNBot智能体与算子开发工作流|CANNBot智能体与算子开发工作流]]"]
---

# Ascend C 算子基础整理

前置阅读：[异构计算基础与 CANN 架构整理](异构计算基础与CANN架构整理.md)。算子定义、Host/Device 协作、SPMD 与 SIMD/SIMT 的区别，以及 CANN 架构集中在基础篇。本篇讲解 Ascend C 接口和核函数实现。

## 1. Ascend C 与核函数
![[CANN3.png|448]]

**Ascend C 是基于 C/C++ 扩展的昇腾 AI 处理器算子编程语言。** 它提供多层 API，支持编写设备端计算代码，并通过 `<<<>>>` 语法从 Host 侧直接启动核函数。

相关概念可以这样区分：

| 概念 | 含义 |
| --- | --- |
| 算子 | 要完成的计算功能，如 Add |
| Kernel／核函数 | 在 Device 上执行该计算任务的函数入口 |
| Ascend C | 表达算子计算和硬件操作的语言与接口体系 |
| 毕昇编译器 | 编译相关代码的工具 |
| Runtime | 管理运行资源、提交任务和等待完成的运行时能力 |

## 2. Add 的数据位置与搬运路径

Add 核函数还涉及另一层搬运：**设备全局内存与计算单元本地存储之间的搬运**。

| 存储位置 | 含义 | Add 示例中的表现 |
| --- | --- | --- |
| Host Memory | CPU 一侧的内存 | 输入准备与结果接收的位置，未在 Add 核函数片段中展示 |
| Global Memory，简称 GM |  Device Memory | `x`、`y`、`z` 指向的设备全局数据 |
| Unified Buffer，简称 UB | AI Core 的片上本地存储 | `x_local`、`y_local`、`z_local` |

完整的数据路径可以写成：

```text
Host Memory
    ↓ Host 将输入复制到设备
Device Global Memory（GM）
    ↓ Kernel 内执行 GM → UB
片上 UB：输入数据
    ↓ SIMD 计算
片上 UB：结果数据
    ↓ Kernel 内执行 UB → GM
Device Global Memory（GM）
    ↓ Host 将结果复制回主机
Host Memory
```

**CPU 与 NPU 之间的搬运、GM 与 UB 之间的搬运是两个不同层次。**`asc_copy_gm2ub` 处理后者，不能把它理解成从 CPU 内存直接拷贝到 NPU。

## 3. Add 的分片索引计算

每个 block 都运行 `add_custom`，通过 `block_idx` 计算自己的数据起点。这是 SPMD 在 Add 中的具体实现。

Add 示例采用均匀分片：

```cpp
constexpr uint32_t block_length = TOTAL_LENGTH / NUM_BLOCKS;
```

设总元素数为 $N$，block 数为 $B$，每个 block 处理 $L=N/B$ 个元素。编号为 $b$ 的 block，其起始偏移是：

$$
\mathrm{offset}_b = bL
$$

对应代码：

```cpp
__gm__ float* x_gm = x + block_idx * block_length;
```

这里偏移的单位是 **float 元素**，因为 C/C++ 指针加法按所指向的类型移动。

由这段代码可以直接看出两个适用条件：

- 简单整除方案需要考虑 `TOTAL_LENGTH` 能否被 `NUM_BLOCKS` 整除；这段代码没有尾部数据处理逻辑。
- 每个 block 的本地数组需要放进可用 UB；数据量更大时，不能无限扩大本地数组。

## 4. Ascend C 的 API 层次

![[CANN4.png]]

不同层次的 API 提供不同粒度的编程抽象。

| 层次或组件 | 作用 |
| --- | --- |
| SIMD C API | 基于指针进行编程，表达搬运、计算等操作 |
| C++ 基础 API／Tensor API | 基于 Tensor 组织数据与计算，Tensor API 结合 Layout 描述数据布局 |
| TPipe／TQue 框架 | 管理同步与内存 |
| 高阶 API | 封装单核公共算法 |
| 算子模板库 | 通过模板提供算子的完整实现 |
| PyAsc | Python 前端编程能力 |

**Add 示例用的是基于指针的 C API。** 因此代码中显式出现 `__gm__`、`__ubuf__`、搬运函数和 `asc_sync()`。不要在阅读这个示例时，把另一套 Tensor 或 TPipe/TQue 写法混入其中。

## 5. Hello World：理解核函数如何启动

### 5.1 核函数与启动代码

```cpp
__global__ __vector__ void hello_world()
{
    printf("Hello World!!!\n");
}

int main(int argc, char const* argv[])
{
    aclrtSetDevice(0);
    aclrtStream stream = nullptr;
    aclrtCreateStream(&stream);

    // Launch kernel <<<numBlocks, dynUBufSize, stream>>>
    hello_world<<<8, 0, stream>>>();

    aclrtSynchronizeStream(stream);
    aclrtDestroyStream(stream);
    aclrtResetDevice(0);
    return 0;
}
```

### 5.2 函数修饰符

| 修饰符 | 含义 |
| --- | --- |
| `__global__` | 标识算子核函数入口 |
| `__vector__` | 指定该算子仅在向量计算单元上执行 |

这里的 `__global__` 是函数修饰符，后面的 `__gm__` 是数据地址空间修饰符，两者用途不同。

### 5.3 Stream 与启动参数

**Stream（计算流）是有序任务队列。** 队列中的任务遵循入队顺序运行。

```cpp
hello_world<<<8, 0, stream>>>();
```

三个启动参数分别是：

| 参数 | 示例值 | 含义 |
| --- | --- | --- |
| `numBlocks` | `8` | 启动的 block 数量 |
| `dynUBufSize` | `0` | 动态 Unified Buffer 大小 |
| `stream` | `stream` | 提交任务使用的计算流 |

`8` 是此示例的启动参数，不能据此判断整颗芯片只有 8 个物理核。

### 5.4 为什么要等待

核函数启动是**异步的**。Host 提交任务之后，不能仅凭调用返回就认为 Device 已完成计算。

```cpp
aclrtSynchronizeStream(stream);
```

这一步让 Host 等待指定 stream 中的任务完成，然后再销毁 stream、重置设备。它针对指定的计算流，不应笼统理解成所有设备、所有流的全局同步。

本例用来观察核函数启动与设备端输出，还没有输入、输出张量，所以没有展示 Host 与 Device 之间的数据拷贝。

## 6. Add：完整理解 SIMD Kernel 的四个步骤

### 6.1 Add 核函数

Add 对两个输入逐元素相加：

$$
z_i = x_i + y_i
$$

`TOTAL_LENGTH` 表示总元素数，`NUM_BLOCKS` 表示分片数。

```cpp
__vector__ __global__ void add_custom(
    __gm__ float* x,
    __gm__ float* y,
    __gm__ float* z)
{
    asc_init();
    constexpr uint32_t block_length = TOTAL_LENGTH / NUM_BLOCKS;

    // Determine the data that each block needs to process
    __gm__ float* x_gm = x + block_idx * block_length;
    __gm__ float* y_gm = y + block_idx * block_length;
    __gm__ float* z_gm = z + block_idx * block_length;

    // Allocate on-chip UB memory
    __ubuf__ float x_local[block_length];
    __ubuf__ float y_local[block_length];
    __ubuf__ float z_local[block_length];

    // Copy input data from GM to UB
    asc_copy_gm2ub(x_local, x_gm, block_length * sizeof(float));
    asc_copy_gm2ub(y_local, y_gm, block_length * sizeof(float));
    asc_sync();

    // Perform element-wise addition
    asc_add(z_local, x_local, y_local, block_length);
    asc_sync();

    // Copy the result from UB to GM
    asc_copy_ub2gm(z_gm, z_local, block_length * sizeof(float));
    asc_sync();
}
```

### 6.2 第一步：分片与本地存储准备

核函数先调用 `asc_init()`，随后计算当前 block 的输入、输出起点。

```cpp
__gm__ float* x_gm = x + block_idx * block_length;
__gm__ float* y_gm = y + block_idx * block_length;
__gm__ float* z_gm = z + block_idx * block_length;
```

三个偏移相同，保证当前 block 处理的是同一组下标上的 $x_i$、$y_i$ 和 $z_i$。

```cpp
__ubuf__ float x_local[block_length];
__ubuf__ float y_local[block_length];
__ubuf__ float z_local[block_length];
```

这些数组位于 UB，分别存放两个输入分片和输出分片。从代码推算，三个数组占用的总字节数为：

$$
3L\times\operatorname{sizeof}(\texttt{float})
$$

这只是三个数组本身的占用，不能直接当作任意程序的全部 UB 需求。

### 6.3 第二步：数据搬入

```cpp
asc_copy_gm2ub(x_local, x_gm, block_length * sizeof(float));
asc_copy_gm2ub(y_local, y_gm, block_length * sizeof(float));
asc_sync();
```

输入先从 GM 搬到 UB，向量计算再使用 UB 中的数据。数据搬运与计算是异步的，需要同步保证依赖顺序。这里的第一次 `asc_sync()` 位于输入搬运之后、计算之前。

### 6.4 第三步：执行 SIMD 计算

```cpp
asc_add(z_local, x_local, y_local, block_length);
asc_sync();
```

`asc_add` 对本地分片进行逐元素相加，将结果写入 `z_local`。第二次同步保证计算完成后，后续输出搬运才能使用结果。

### 6.5 第四步：结果搬出

```cpp
asc_copy_ub2gm(z_gm, z_local, block_length * sizeof(float));
asc_sync();
```

结果由 UB 写回当前 block 对应的 GM 分片。第三次同步放在输出搬运之后。

**执行完这一阶段，结果位于设备 GM 中。**Host 若需要读取结果，还需要在相应任务完成后安排设备到主机的数据拷贝。

### 6.6 最容易混淆的长度单位

| 位置 | 表达式 | 单位 |
| --- | --- | --- |
| 指针偏移 | `block_idx * block_length` | float 元素数 |
| 搬运长度 | `block_length * sizeof(float)` | 字节数 |
| Add 计算长度 | `block_length` | 元素数 |

这也是阅读算子代码时应先检查的地方：相同的逻辑长度，传给不同接口时可能使用不同单位。

## 7. 两种同步分别解决什么问题

| 同步调用 | 所在位置 | 解决的问题 |
| --- | --- | --- |
| `aclrtSynchronizeStream(stream)` | Host 端 | 等待指定 stream 中的任务完成 |
| `asc_sync()` | Device 核函数内部 | 保证示例中的搬运、计算、写回按依赖顺序完成 |

Host 等待整个任务完成，不能替代核函数内部的依赖管理。核函数内部完成同步，也不代表结果已经拷贝回 Host。

## 8. 实践入口与验证范围

### 8.1 前置知识

前置知识包括 C/C++ 基本语法、Makefile 基本概念和 Git 常见操作。

学习这两个示例时，还应能读懂指针偏移、数组长度、`sizeof`、函数修饰符，以及 Host/Device 的执行边界。

### 8.2 完整样例

- [Hello World 完整源码](https://gitcode.com/cann/asc-devkit/blob/master/examples/01_simd_cpp_api/00_introduction/00_quickstart/hello_world/hello_world.asc)
- [Add 完整源码](https://gitcode.com/cann/asc-devkit/blob/master/examples/02_simd_c_api/00_introduction/01_add/c_api_async_add/c_api_add.asc)
- [SIMD C++ API 入门目录](https://gitcode.com/cann/asc-devkit/tree/master/examples/01_simd_cpp_api/00_introduction)
- [SIMD C API 入门目录](https://gitcode.com/cann/asc-devkit/tree/master/examples/02_simd_c_api/00_introduction)

### 8.3 阅读或运行时检查什么

以下检查点由示例逻辑归纳：

1. **Hello World**：确认核函数确实在设备端执行，理解提交与等待的区别。
2. **Add 数据分片**：确认启动 block 数与分片使用的 `NUM_BLOCKS` 一致，检查分片能否覆盖输入。
3. **长度与存储**：确认搬运字节数、计算元素数与本地数组长度相符，UB 能容纳分片。
4. **计算结果**：将返回 Host 的结果与逐元素相加的参考结果比较；浮点误差判定遵循完整样例的验证方式。
5. **同步位置**：分别检查 Host 等待与核函数内部依赖，不能只检查其中一层。

## 9. 课后问题与复习答案

### 9.1 为什么 Hello World 使用 `__vector__`？

`__vector__` 指定核函数在向量计算单元上执行，`__global__` 标识核函数入口。

### 9.2 Device Memory 与 UB 相同吗？

不同。示例用 Device Memory／GM 保存设备全局输入与输出，用片上 UB 保存当前计算使用的本地分片，并显式在两者之间搬运。

### 9.3 Add 核函数的四个基本步骤是什么？

**Tiling 分片、数据搬入、执行计算、数据搬出。**代码中还要为本地数据安排存储，并在有依赖的异步操作之间同步。

### 9.4 能否直接复制本笔记代码运行？

运行时需要补齐头文件、构建配置，以及 Add 的常量定义、Host 侧内存管理与调用代码，可参考完整样例工程。

## 10. 后续学习

后续可以进一步学习硬件抽象、内存层级、细粒度同步，以及算子端到端编译流程。

参考资料：

- [Ascend C API 实现与样例](https://gitcode.com/cann/asc-devkit)
- [Ascend C 编程指南](https://asc.gitcode.com)
- [CANN 算子领域样例](https://gitcode.com/cann/cann-samples/tree/master/Samples/)
- [CANN Learning Hub](https://gitcode.com/cann/cann-learning-hub/)