---
tags: [CANN]
doc_type: knowledge-note
返回目录: "[[计算机系/昇腾与CANN/大模型与算子专题学习导航|大模型与算子专题学习导航]]"
上一篇: "[[计算机系/昇腾与CANN/Ascend C矩阵算子与数据布局|Ascend C矩阵算子与数据布局]]"
下一篇: "[[计算机系/昇腾与CANN/Ascend C接入PyTorch单算子|Ascend C接入PyTorch单算子]]"
前置知识: "[[计算机系/昇腾与CANN/Ascend C矩阵算子与数据布局|Ascend C矩阵算子与数据布局]]"
---

# Ascend C融合算子与核间协同

算子融合把多个逻辑运算合并进一次执行，例如 Matmul 后接 LeakyReLU。一般意义上的融合可以完全发生在 Vector 上；AIC/AIV 混合编程则专门讨论矩阵核和矢量核如何协作。

本篇建立在 [矩阵算子与数据布局](Ascend%20C矩阵算子与数据布局.md) 和 [Memory 矢量开发](Ascend%20C%20Memory矢量算子开发.md) 之上，只展开核间分工、通信和同步。

## 1. 一个核函数，不等于一个硬件执行单元

```cpp
__global__ __mix__(1, 2) void mixed_kernel(/* 参数 */)
{
    if ASCEND_IS_AIC {
        // 矩阵计算路径
    }
    if ASCEND_IS_AIV {
        // 矢量计算路径
    }
}
```

`__mix__(1,2)` 表达一个 AIC 配合两个 AIV 的组合。A2/A3 的 AIC 与 AIV 各自执行相应二进制，拥有独立的指令发射与执行过程。源码里 AIC 分支写在前面，并不保证 AIV 等到它结束后才开始。

公共代码会出现在两类执行路径中，因此 L1/L0 的分配放在 AIC 路径，UB 的分配放在 AIV 路径。不能在公共区域无条件执行仅适用于某一侧的内存或计算操作。

![混合计算与跨核协作](.assets/AscendC进阶/fusion-8.png)

## 2. 给每个 AIV 分配不同的数据

| 变量 | 含义 |
| --- | --- |
| `block_idx` | 当前逻辑 block 的编号 |
| `block_num` | 启动配置的 block 数 |
| `sub_block_idx` | 组内 AIV 的编号，1:2 配置中为 0 或 1 |
| `sub_block_num` | 组内 AIV 数量，1:2 配置中为 2 |

`idx` 是编号，`num` 是数量，不能颠倒。假设一个 AIC 产生 `baseM × baseN` 的结果，两个 AIV 沿 M 分成两半：

$$
\text{elementsPerAIV}=\frac{\text{baseM}}2\,\text{baseN},
$$

$$
\text{localOffset}=\text{sub\_block\_idx}\times\text{elementsPerAIV}.
$$

还需叠加 block 与 tile 在整个输出中的偏移。这个公式要求 baseM 可被 2 整除；尾块应按有效行数重新分配，不能让两个 AIV 重复处理同一范围。

## 3. GM 中转与同步各解决什么问题

A2/A3 上 AIC 与 AIV 的这类通信通过 GM 中转，不能直接假设 AIV 能读取 AIC 的 L0C。

```text
AIC：GM → L1 → L0A/L0B → Mmad → L0C → Fixpipe → GM
                                                       ↓
AIV：                                  GM → UB → 激活 → GM
```

GM 保存数据，核间通知表示“某个阶段已经完成”。通知本身既不携带矩阵，也不替代数据搬运。

对一个结果 tile，正确顺序为：

1. AIC 完成计算，并把中间结果写入 GM。
2. 在结果写回完成的依赖点发送通知。
3. 对应 AIV 等待通知后，把自己的分片搬入 UB。
4. AIV 计算激活，写回最终输出。
5. 若 GM 中间缓冲会复用，生产者必须等相关消费者读完，再覆盖该缓冲。

最后一步是反向保护：只有“写完通知读”而没有“读完允许下一次写”，循环时仍可能出现竞争。基本原理见 [流水线双向同步](流水线双向同步原理.md)，但这里参与者是不同核，不能用核内 Flag 冒充核间通知。

## 4. 选择合适的同步范围

![核间同步范围](.assets/AscendC进阶/fusion-11.png)

| 模式 | 同步范围 | C API 家族 |
| --- | --- | --- |
| Mode 0 | 全体参与的计算核 | `asc_sync_inter_arrive/wait` |
| Mode 1 | 同组两个 AIV | `asc_sync_subblock_arrive/wait` |
| Mode 2 | 同组 AIC 与全部 AIV | `asc_sync_block_arrive/wait` |
| Mode 4 | AIC 与指定 AIV | `asc_sync_intra_arrive/wait`，A2/A3 不支持 |

这些模式不能只按编号互换。调用时要确认参与者、流水依赖、flag 标识和通知/等待次数匹配。某个核提前退出、尾块多执行一次 wait，都会破坏协议。

基础 API 的 AIC 结果通知片段可以是：

```cpp
// 放在当前 tile 的 Fixpipe 写回操作之后
AscendC::CrossCoreSetFlag<0x2, PIPE_FIX>(0);
```

它把通知关联到 FIX 流水的完成状态。接收方必须采用对应硬件/API 规定的等待方式；设置和等待接口的模板参数不一定对称，不能从发送端形式机械推导接收端签名。

全核同步比组内同步覆盖范围大，也更容易让快核等待慢核。如果依赖只存在于同组 AIC/AIV，通常不需要所有组一起停下。

## 5. Matmul + LeakyReLU 的手工实现

计算为：

$$
D=AB,\qquad Y_{ij}=\begin{cases}
D_{ij},&D_{ij}\ge0,\\
\alpha D_{ij},&D_{ij}<0.
\end{cases}
$$

例如取 $\alpha=0.001$。AIC 负责 D，两个 AIV 分别负责 D 的半块激活。实现时需要显式管理：矩阵布局转换、Mmad 累加、Fixpipe 输出格式、GM 中间区、AIV 搬入和激活、结果回写、双方同步。

如果先把 float 累加结果转换成 half 再激活，与在 float 中激活后最后转换成 half，可能产生不同舍入结果。中间精度属于算子语义的一部分，应与参考结果一致。

融合的收益可能来自减少 Host 启动、缩短中间结果生命周期和改善流水协作，但在 A2/A3 上不能据此声称完全消除了 GM 访问。额外通信、同步或 AIC/AIV 负载失衡也可能抵消收益。

## 6. 高阶 Matmul 与自动 C/V 分离

![KFC Client/Server 交互](.assets/AscendC进阶/fusion-17.png)

混合配置中的 Matmul 高阶 API 允许从 AIV 视角请求矩阵计算。KFC Client 向 GM 中的通信区域提交请求，AIC 上的 KFC Server 读取请求并执行矩阵计算，再交付结果。

这封装了核间控制过程，但不改变 AIC/AIV 是独立执行单元的事实，也不自动消除物理通信成本。纯 Cube 的 Matmul 配置不应一概解释成这个 Client/Server 模式。

调用流程为：

```text
定义 Matmul 的类型、位置和格式
→ 注册 Matmul 对象与 tiling、系统 workspace
→ 设置输入和原始 Shape
→ Iterate 逐块计算
→ GetTensorC 取出当前结果
→ 队列移交给 Vector 激活
→ 搬出并释放缓冲
→ End
```

激活部分的关键接口：

```cpp
AscendC::LeakyRelu(
    reluOutLocal, mmOutLocal, static_cast<float>(0.001),
    tiling.baseM * tiling.baseN);
```

这里假定 `mmOutLocal` 已经由 Matmul 完成写入，输入输出缓冲已分配，元素数是当前有效计算范围。实际尾块不能无条件沿用完整 tile 大小。

`REGIST_MATMUL_OBJ` 使用系统 workspace；手工中间结果也可能另需 workspace。两者用途不同，Host 要保证大小、地址与异步执行生命周期有效。

选择手工路径是为了控制精细数据流；选择高阶路径是为了复用已经封装的矩阵计算和协作逻辑。两者都要验证数值正确性、尾块和实际性能。

完整样例：[手工融合](https://gitcode.com/cann/asc-devkit/tree/master/examples/01_simd_cpp_api/00_introduction/03_fusion_operation/matmul_leakyrelu_basic_api)、[高阶融合](https://gitcode.com/cann/asc-devkit/tree/master/examples/01_simd_cpp_api/00_introduction/03_fusion_operation/matmul_leakyrelu_advanced_api)。

## 学习随记

把多个融合算子结合成一个算子