---
tags: [AscendC, SIMD, 内存, 同步]
doc_type: knowledge-note
abstract: "在Add入门基础上理解AIC/AIV、片上内存、流水线同步、核类型、事件及仿真编译。"
creation: 2026-10-15
返回目录: "[[计算机系/昇腾与CANN/大模型与算子专题学习导航|大模型与算子专题学习导航]]"
上一篇: "[[计算机系/昇腾与CANN/CMake与Ascend C编译入门|CMake与Ascend C编译入门]]"
下一篇: "[[计算机系/昇腾与CANN/流水线双向同步原理|流水线双向同步原理]]"
后续学习: ["[[计算机系/昇腾与CANN/Ascend C矩阵算子与数据布局|Ascend C矩阵算子与数据布局]]", "[[计算机系/昇腾与CANN/Ascend C Memory矢量算子开发|Ascend C Memory矢量算子开发]]"]
---

# Ascend C SIMD 硬件内存与同步

前置：[算子基础](Ascend%20C算子基础整理.md) · 构建：[CMake 与 Ascend C](CMake与Ascend%20C编译入门.md)

本篇沿着“计算在哪执行、数据放在哪里、依赖怎样满足”的顺序深入 SIMD。SPMD/SIMD/SIMT 的基本定义见 [异构计算基础](异构计算基础与CANN架构整理.md)，这里不重复比较表。

## 1. 编程视角与硬件抽象

SIMD 适合规整、密集的向量或矩阵计算；SIMT 适合离散访问、复杂分支和稀疏逻辑。A2/A3 主线使用 SIMD，950PR/950DT 扩展 SIMT 及混合能力。

混合模式可在核函数内部以 Vector Function 为粒度切换：SIMD VF 调用向量指令，SIMT VF 抽象为线程块并按线程编程。它不是任意平台都支持的功能。

![AI Core硬件抽象](.assets/大模型与算子专题/AI-Core硬件抽象.png)

| 单元 | 职责 |
| --- | --- |
| Scalar | 地址计算、循环控制、发射计算搬运与同步指令 |
| Vector | 向量计算 |
| Cube | 矩阵计算 |
| Local Memory | 保存核内计算的数据和中间结果 |
| DMA / MTE | 执行不同存储位置之间的搬运 |

### A2/A3 的 AIC 与 AIV

![A2/A3硬件架构](.assets/大模型与算子专题/A2-A3硬件架构.png)

逻辑组合中 AIC 与 AIV 比例为 1:2，二者各有 Scalar 控制能力。AIV 的 Memory Vector 计算使用 UB，AIC 使用 L1、L0A、L0B 和 L0C 完成矩阵计算。

矩阵计算可表示为 $C\leftarrow AB+C$：L0A 保存 A，L0B 保存 B，L0C 保存累加结果。FP16 的典型 Cube 运算块为 $16\times16$，不能将这个块大小直接等同于任意矩阵乘的总耗时。

实际设备核数依具体型号变化，应读取设备能力，而不是把 `numBlocks` 写成所有产品通用的最大核数。

## 2. 核函数类型与索引

| 执行空间修饰符 | 使用单元 | 示例 |
| --- | --- | --- |
| `__vector__` | AIV | Softmax、逐元素运算 |
| `__cube__` | AIC | MatMul |
| `__mix__(1, 1)` / `__mix__(1, 2)` | AIC 与 AIV 配合 | 矩阵与向量阶段融合 |

核函数入口使用 `__global__`，返回类型为 `void`，参数可包含基础类型、结构体或指针等受支持类型。

```cpp
__global__ __vector__ void vector_kernel(__gm__ float* input);
```

启动形式：

```cpp
vector_kernel<<<numBlocks, dynUBufSize, stream>>>(input);
```

`block_idx` 标识当前逻辑 block。混合模式还通过 `sub_block_idx` 区分组合内的 Vector 子核，1:2 配比下子核索引为 0、1。它与全局 block 索引属于不同层次。

![核函数调用关系](.assets/大模型与算子专题/核函数调用.png)

图展示 Host 调用核函数、核函数再调用设备函数的关系；图中的 `__aicore__` 是另一套接口表达，不能直接替换当前示例的执行空间修饰符。

## 3. 内存层级与地址空间

![SIMD内存层级](.assets/大模型与算子专题/SIMD内存层级.png)

| 存储 | 地址空间修饰符 | 用途 | 对齐要求示例 |
| --- | --- | --- | --- |
| GM | `__gm__` | 全局输入、输出、临时空间 | 基础地址空间为字节寻址，具体搬运 API 可能要求更严格 |
| UB | `__ubuf__` | Vector 输入、输出和临时数据 | 32 B |
| L1 | `__cbuf__` | 矩阵输入缓存 | 32 B |
| L0A | `__ca__` | 左矩阵 A | 512 B |
| L0B | `__cb__` | 右矩阵 B | 512 B |
| L0C | `__cc__` | 矩阵累加结果 | 64 B |
| Fixpipe Buffer | `__fbuf__` | 量化参数暂存 | 64 B |
| Bias Table | `__biasbuf__` | 偏置缓冲 | 64 B |

![存储规格表](.assets/大模型与算子专题/存储规格原表.png)

A2/A3 的示例容量包括 UB 192K、L1 512K、L0A/L0B 各 64K，特殊缓冲示例为 6KB 和 4KB。**L0C 容量存在资料冲突：本表写 128K，前面的硬件概览写 256K。应以具体芯片和开发工具提供的能力值为准，不能将任一数字当作所有 A2/A3 的固定上限。**

### 对齐不只看数组起点

即使基地址对齐，增加元素偏移后也可能破坏对齐。对于需要 32B 对齐的 `float` 地址，偏移 8 个元素是 32B，偏移 1 个元素只是 4B。

对齐要求、搬运长度要求和尾部处理是三个相关但不同的问题，应分别检查。`__gm__` 可以字节寻址，不表示任何搬运 API 都支持任意字节数。

## 4. 片上内存怎样分配

![片上内存分配](.assets/大模型与算子专题/片上内存分配.png)

### 静态 UB 数组

```cpp
__ubuf__ half staticBuf[1024];
```

长度在编译时确定。多个输入、输出、临时数组共同占用预算，不能各自按完整 UB 容量申请。

### 动态 UB 空间

```cpp
extern __ubuf__ half dynBuf[];
```

启动参数 `dynUBufSize` 以字节为单位提供动态 UB 空间，它只针对 UB，不是所有片上存储的统一动态容量。

另一种方式是使用 TPipe/TQue 或 LocalMemoryAllocator 管理分配。动态数组与框架分配不应未经协调同时管理同一动态区域，以免重叠。

### 预留与可用预算

当前示例的编译器与 ASC 预留 UB 为 8K，用于寄存器溢出和临时空间。因此 192K UB 下的应用预算示例是：

$$
192K-8K=184K
$$

若三份 float 分片占用 $3L\times4$ 字节，需在可用预算内再考虑对齐、其他临时数据与流水线缓冲。容量与预留策略随产品和工具链变化。

`--cce-disable-asc-reserved-ubuf` 可以关闭相应预留，但可能让依赖临时空间的 API 编译失败，不宜作为默认省内存手段。

L1/L0A/L0B/L0C 不通过启动参数的动态 UB 数组申请；示例使用框架分配器管理。静态数组支持也需与所用工具链版本匹配。

## 5. 流水线与数据路径

| 流水线 | 任务 |
| --- | --- |
| `PIPE_S` | 标量计算和控制 |
| `PIPE_V` | 向量计算 |
| `PIPE_M` | 矩阵计算 |
| `PIPE_MTE1` | L1 到 L0A/L0B |
| `PIPE_MTE2` | GM 到 UB/L1 等受支持路径 |
| `PIPE_MTE3` | UB 到 GM，以及受支持的其他路径 |
| `PIPE_FIX` | L0C 结果搬出与转换 |

多个流水线可以异步推进，但前提是没有未满足的数据依赖。搬运 B 分片与计算 A 分片可以重叠，计算 A 分片却必须等待 A 的输入搬完。

搬运能力需要按芯片细分：通用流水线表中的 UB→L1、L0C→L1 等路径不应无条件套到 A2/A3。A2/A3 的 AIC/AIV 交换按这里的模型通过 GM 中转。

## 6. 同步要匹配依赖范围

![同步类型](.assets/大模型与算子专题/同步类型.png)

| 范围 | 需要解决的依赖 |
| --- | --- |
| 核内跨流水线 | 搬运结果就绪后才能计算 |
| 核内同流水线 | 有相关数据依赖的前后指令 |
| 内存访问顺序 | 写入与后续读取、覆盖之间的可见性和次序 |
| 核间 | 不同 AIC/AIV 的数据生产与消费 |
| Host 与设备流 | Host 使用结果前等待对应设备任务 |

同一 Stream 有序，不等于核函数内部的全部硬件流水线自动满足依赖。普通全核同步也不能任意代替特定生产者/消费者的同步协议。

### Notify/Wait 的双向依赖

A2/A3 使用 Notify/Wait 或 Set/WaitFlag 类机制。一个可复用缓冲区既有“数据已生产”的依赖，也有“缓冲区可以覆盖”的依赖：

```text
生产者 A：等待缓冲空闲 → 搬入数据 → 通知 B 数据就绪
消费者 B：等待数据就绪 → 使用数据 → 通知 A 缓冲可复用
```

省掉第二个方向，下一轮搬入可能覆盖尚未计算完的数据。循环首轮、末轮需要处理事件配对，错误的等待可能死锁。

950 增加的锁式机制不能直接迁移到只支持 Notify/Wait 的 A2/A3。这里的 lock/notify/wait 是机制描述，实际 API 名称与参数需要与具体编程层匹配。

## 7. 算子类型与接口选择

Memory Vector 计算直接面向 UB，Reg Vector 计算面向寄存器；后者在支持的 950 系列中使用 VF，相关 SIMD VF 由 `__simd_vf__` 修饰。

![SIMD接口层次](.assets/大模型与算子专题/SIMD接口层次.png)

| 接口层 | 编程关注点 |
| --- | --- |
| 指针 C API | 显式地址、搬运、计算和同步 |
| Tensor / Layout API | 数据布局和编译期抽象，Arch/Atom/Algorithm 分层 |
| TPipe/TQue 框架 | 按队列和 Tensor 生命周期组织内存与依赖 |

框架降低手工管理复杂度，但仍需要正确设计分块、缓冲大小和生产消费顺序。学习时可以用同一个 Add 对比不同接口，避免把 API 名称不同误解成数学计算不同。

## 8. Stream 与 Event

Stream 表达任务队列的顺序，Event 是插入流中的执行进度标记。Host 发出“记录事件”的调用，不等于事件已经在设备上完成。

常见用途包括等待特定阶段、建立不同流之间的依赖，以及统计设备时间。计时常结合 `aclrtRecordEvent` 和 `aclrtEventElapsedTime`，并在读取耗时前确认结束事件完成；具体事件创建标志也要支持计时。

典型生命周期是选择设备、创建流、分配内存、搬入输入、启动 kernel、等待、回拷结果、释放内存与流。具体 Host 示例见 [算子基础](Ascend%20C算子基础整理.md)，不要将其中的 `asc_sync()` 与 Runtime Event 混为同一层接口。

## 9. 编译、架构与仿真

`.asc` 由毕昇按 Ascend C 处理；`.cpp` 默认作为 C++，按 ASC 处理时需要相应语言选项，例如 `-xasc`。头文件可以供 Host、Device 或混合源码包含，并非独立可执行单元。

物理 NPU 示例：

```bash
bisheng main.cpp example.asc -o main --npu-arch=dav-2201
```

仿真配置可使用 CMake 变量：

```bash
cmake -S . -B build-sim -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201
cmake --build build-sim
```

前提是项目的 ASC 包支持这些变量，且环境具备对应仿真组件。`CMAKE_ASC_RUN_MODE=npu` 表示物理设备模式。不同模式建议使用不同构建目录，避免缓存混用。

可以在加载 ASC 包前设置默认架构：

```cmake
set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture")
find_package(ASC REQUIRED)
project(kernel_samples LANGUAGES ASC)
```

它与 [CMake 最小配置与编译选项](01%20CMake最小配置与编译选项.md)中的目标级架构选项都是设置架构的入口，不应给同一编译步骤传入相互矛盾的值。仿真用于功能与开发验证，性能数字不能直接代表真实设备。

## 10. 学习实践

用 C API 和基于 Tensor/框架的 API 分别组织形状 `[1024, 1024]` 的 Add：确认元素总量、核间分片、核内分块、UB 预算、对齐与同步，再与参考结果比较。

复习重点是解释：动态 UB 参数为何不是 L1 容量；32B 对齐为何会被偏移破坏；循环缓冲为什么需要双向同步；Event 记录与事件完成为什么不同。

参考：[Ascend C 编程指南](https://asc.gitcode.com)、[asc-devkit](https://gitcode.com/cann/asc-devkit)。