---
tags: [CANN]
doc_type: knowledge-note
返回目录: "[[计算机系/昇腾与CANN/大模型与算子专题学习导航|大模型与算子专题学习导航]]"
上一篇: "[[计算机系/昇腾与CANN/Ascend C Memory矢量算子开发|Ascend C Memory矢量算子开发]]"
下一篇: "[[计算机系/昇腾与CANN/Ascend C融合算子与核间协同|Ascend C融合算子与核间协同]]"
前置知识: "[[计算机系/昇腾与CANN/Ascend C SIMD硬件内存与同步|Ascend C SIMD硬件内存与同步]]"
后续学习: ["[[计算机系/昇腾与CANN/Ascend C融合算子与核间协同|Ascend C融合算子与核间协同]]"]
---

# Ascend C矩阵算子与数据布局

矩阵乘的数学接口可以写成函数 $C=AB$，但高效实现还要决定分块、内存格式、搬运顺序与累加精度。纯 AIC 路径使用 `__global__ __cube__`，通过 L1、L0A、L0B、L0C 和 Cube 完成计算。

基础内存概念见 [SIMD 硬件内存与同步](Ascend%20C%20SIMD硬件内存与同步.md)，Tensor 与队列生命周期见 [Memory 矢量开发](Ascend%20C%20Memory矢量算子开发.md)。

## 1. 从矩阵公式理解硬件工作

对于 $A\in\mathbb{R}^{M\times K}$、$B\in\mathbb{R}^{K\times N}$：

$$
C_{ij}=C^{\mathrm{init}}_{ij}+\sum_{k=0}^{K-1}A_{ik}B_{kj}.
$$

普通矩阵乘从零开始累加；带偏置时可以从广播后的 Bias 开始。K 是归约维，M、N 决定输出位置。

![矩阵乘加数学关系](.assets/AscendC进阶/matrix-10.png)

沿 M/N 分块可以让不同核计算互不重叠的 C；沿 K 分块则产生同一输出的部分和，必须继续累加。不能在每个 K 块开始时清零 C，否则最后只剩最后一块的结果。

## 2. 数据路径与流水分工

![矩阵数据流与分块](.assets/AscendC进阶/matrix-11.png)

| 阶段 | 内存路径 | 流水 | 基础 API |
| --- | --- | --- | --- |
| CopyIn | GM → L1 | MTE2 | `DataCopy`，可伴随 ND2Nz |
| Split | L1 → L0A/L0B | MTE1 | `LoadData` |
| Compute/Accumulate | L0A × L0B → L0C | M | `Mmad` |
| CopyOut | L0C → GM，部分路径可到 L1 | FIX | `Fixpipe` |

MTE2 搬完后 MTE1 才能读取；MTE1 搬完后 Cube 才能计算；Cube 产生有效结果后 FIX 才能搬出。循环覆盖 L1、L0A/B、L0C 时还要保护上一轮消费者。用普通源代码的先后顺序不能代替跨流水依赖。

矩阵搬运是有方向限制的硬件通路，不能把所有局部内存都当作能双向 `memcpy` 的普通数组。L0C 回到 L1 也不等于 L0C 可以直接搬到 L0A/B。

## 3. Shape、Stride 和分形格式

Shape 描述逻辑尺寸，Stride 描述相邻逻辑索引在内存中跨越的元素数。二维张量的地址为：

$$
\operatorname{offset}(i,j)=i\,s_0+j\,s_1.
$$

Shape 为 `(2,4)` 时，行主序 Stride 是 `(4,1)`，列主序是 `(1,2)`。相同 Shape 不保证相同物理排列。

分形格式先把大矩阵分成小矩阵，再分别规定块间与块内顺序：

| 字母 | 含义 |
| --- | --- |
| 大写 Z | 小矩阵之间按行主序排列 |
| 大写 N | 小矩阵之间按列主序排列 |
| 小写 z | 每个小矩阵内部按行主序排列 |
| 小写 n | 每个小矩阵内部按列主序排列 |

![ND 与分形格式](.assets/AscendC进阶/matrix-13.png)

因此 Zz 与 Zn 的区别在块内排列，Nz 与 Zz 的区别在块间排列。格式转换改变存储顺序，通常不改变逻辑矩阵的数学值；转置则改变逻辑索引含义，二者不能混为一谈。

A2/A3 的典型 Mmad 输入输出要求是 **A 为 Zz，B 为 Zn，C 为 Nz**。

![A、B、C 的布局要求](.assets/AscendC进阶/matrix-14.png)

GM 的 ND 数据可以在搬到 L1 时转换为 Nz，再通过 LoadData 组织成 L0 的目标格式。必须按最终消费接口检查格式，不能认为“L1 中是 Nz，所以 L0A 也必然是 Nz”。

![搬运中的布局变换](.assets/AscendC进阶/matrix-23.png)

## 4. 计算粒度、数据类型与尾块

典型小块的计算尺寸为：

$$
(16\times K_0)(K_0\times16)\rightarrow16\times16,
\qquad K_0=32/\operatorname{sizeof}(T).
$$

对 half，$K_0=16$；对 int8，$K_0=32$。这是理解分块的硬件粒度，不意味着一次 `Mmad` 只能计算一个小块；接口可根据 M/N/K 完成多个小块。

常见组合是 half/bfloat16 输入、float 累加，以及 int8 输入、int32 累加。float 输入的支持方式和精度模式需要与目标架构的接口配置匹配，不能仅按类型名称推断吞吐或精度。

![非对齐矩阵与 GEMV](.assets/AscendC进阶/matrix-26.png)

例如 half 输入的 M=17、K=58、N=30，不满足完整小块要求。物理缓冲需要按相关粒度分配；K 方向的无效输入必须正确填零，避免污染有效输出。最终只写回逻辑范围内的 C。不能因为输出会裁剪，就允许输入越界读取。

M=1 时可考虑受支持的 GEMV 模式，减少按完整 16 行矩阵计算造成的浪费。启用方式与 `disable_gemv` 参数含义要对应，不能只凭变量名设置布尔值。

## 5. 手工 Tensor 实现的关键点

```cpp
AscendC::LocalMemAllocator<AscendC::Hardware::L1> l1;
AscendC::LocalMemAllocator<AscendC::Hardware::L0A> l0a;
AscendC::LocalMemAllocator<AscendC::Hardware::L0B> l0b;
AscendC::LocalMemAllocator<AscendC::Hardware::L0C> l0c;
auto a1 = l1.Alloc<half, baseM * baseK>();
auto b1 = l1.Alloc<half, baseK * baseN>();
auto a2 = l0a.Alloc<half, baseM * baseK>();
auto b2 = l0b.Alloc<half, baseK * baseN>();
auto c = l0c.Alloc<float, baseM * baseN>();
```

这里的尺寸使用已适配格式和对齐的 tile。half 输入不意味着 L0C 也应分配成 half；累加类型往往是 float。

典型相关地址对齐为 L1 32B、L0A/L0B 512B、L0C 搬出源地址 64B。地址对齐、搬运长度粒度与矩阵形状对齐分别检查。

Mmad 参数片段：

```cpp
AscendC::MmadParams params;
params.m = curM;
params.n = curN;
params.k = baseK;
params.cmatrixInitVal = (kBlockIdx == 0);
AscendC::Mmad(c, a2, b2, params);
```

这段位于已经完成加载与同步的 K 循环内。`cmatrixInitVal` 只在第一块启用；最后通过 Fixpipe 写出结果，必要时进行 float→half 转换。选择具名字段比记忆聚合初始化的位置更不易出错。

缓冲预算要分别满足每一级容量。例如 L0C 至少容纳 `baseM*baseN*sizeof(float)`，双缓冲还要乘以对应份数。不能只检查三种内存占用之和。

## 6. 队列路径与 Matmul 高阶接口

![矩阵队列阶段和逻辑位置](.assets/AscendC进阶/matrix-41.png)

矩阵队列分为 CopyIn、Split、Compute、Aggregate、CopyOut。`A1/B1` 对应 L1，`A2/B2` 对应 L0A/B，`CO1` 对应 L0C。逻辑位置描述用途，不能跨架构把所有逻辑位置固定解释为同一种物理内存。

高阶 `Matmul` 把上述流程封装起来，但仍需要 Host 提供合理的 tiling 和类型配置。典型调用顺序：

```cpp
REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), mm, &tiling);
mm.SetOrgShape(tiling.M, tiling.N, tiling.Ka, tiling.Kb);
mm.SetTensorA(aGlobal[blockIdx * tiling.singleCoreM * tiling.Ka], false);
mm.SetTensorB(bGlobal[0], false);
mm.IterateAll(cGlobal[blockIdx * tiling.singleCoreM * tiling.N]);
mm.End();
```

该偏移方式假定仅沿 M 划核、每核覆盖完整 N，且数据为相应 ND 排列。A 按行偏移，B 被各核共享读取，C 按输出行偏移；沿 N 或 M/N 二维划分时需要重新计算，不能照搬。

Matmul 模板中的位置、格式、数据类型必须与 Host tiling 配置一致。纯 Cube 配置与 AIV 发起的混合配置也不同；后者的 KFC 协同见 [融合算子与核间协同](Ascend%20C融合算子与核间协同.md)。

## 7. 练习的规模如何转化为设计约束

对 M=K=N=8192、half 输入、float 输出：A、B 各占 128 MiB，C 占 256 MiB；总计算量约为 $2MNK\approx1.10\times10^{12}$ FLOPs。

这样的矩阵不可能整体放入片上缓存。设计时先选输出 tile 的 M/N，再选 K 分块；检查各级容量，安排多核分工，最后尝试数据复用和双缓冲。验证除了大矩阵，还应包括非整除形状与 M=1 等特殊情况。

实现入口：[矩阵基础与高阶 API 样例](https://gitcode.com/cann/asc-devkit/tree/master/examples/01_simd_cpp_api/00_introduction/02_matrix)。