Ascend 910B2 / CANN 8.5.0：Kernel 心智模型与 Microbench v0.1

状态：v0.1 / 学习与基线阶段
目标环境：CANN 8.5.0 + Ascend 910B2。
当前若使用 910B4-1 训练机，可以提前验证编程模型、数据通路和 tiling 方法，但不要把 B4-1 上的绝对性能直接当成 910B2 的最终结论。

这份笔记的目标不是从零介绍 Ascend C，而是给已经熟悉 CUDA kernel 优化的人建立一套最短的迁移路径：先测清 MTE / UB / Vector，再进入 L1 / L0A / L0B / L0C / Cube。

1. 先建立两张数据通路图

Vector / AIV

Vector 核最重要的数据通路可以先记成：

GM
 ↓ DataCopy / MTE
UB
 ↓ Vector Unit
UB
 ↓ DataCopy / MTE
GM

这里不要强行寻找 threadIdx.x 或 warp lane 的一一对应关系。对性能分析更有用的单位是：

GM 与 UB 之间搬了多少字节；
UB tile 多大；
Vector 指令一次处理多少数据；
CopyIn / Compute / CopyOut 是否形成流水。

CANN 8.5.0 的 Vector 高维切分模型中，一个 repeat 会处理 8 个 datablock，每个 datablock 为 32 B，因此单 repeat 对应 256 B 数据。对 16-bit 数据即最多 128 个元素，对 32-bit 数据即最多 64 个元素。repeatTime 最大为 255。

普通 Vector 访问也不局限于完全连续地址：可以通过 datablock stride、repeat stride 和 mask 描述规则的 strided SIMD 访问；但它仍然不是任意地址的通用 gather/scatter 模型。

Cube / AIC

Cube 的核心心智模型是：

                 ┌── L0A ──┐
GM → L1 ─────────┤          │
                 │  Cube    ├──→ L0C → GM
GM → L1 ─────────┤          │
                 └── L0B ──┘
                       ↑
                     Mmad

对 CUDA Tensor Core 开发者，可以暂时做如下“编程心智模型”类比（不是物理结构一一等价）：

CUDA Tensor Core	Ascend Cube
shared-memory tile	L1 tile
`ldmatrix`/operand load	`LoadData` 等搬运/格式准备
A fragment	L0A
B fragment	L0B
accumulator fragment	L0C
`mma.sync`	`Mmad`

Cube 的性能重点通常不只是 Mmad 本身，而是：

GM → L1 怎么切 tile；
L1 → L0A/L0B 怎么准备格式；
M/N/K 基本 tile 如何选择；
MTE 与 Cube 是否 overlap；
L0C 如何写回或交给后续计算。

2. ZZ / ZN / NZ：Cube 的分形物理布局

CANN 8.5.0 对 Cube 相关本地存储给出的推荐布局是：

L0A : FRACTAL_ZZ
L0B : FRACTAL_ZN
L0C : FRACTAL_NZ
L1  : 推荐 FRACTAL_NZ
UB  : 无固定数据格式要求

同时对应的最小访问对齐要求为：

存储	对齐
UB	32 B
L1	32 B
L0A	512 B
L0B	512 B
L0C	64 B

Z/N 可以理解成分形层级上的 row-major / column-major 排布方向：

ZZ：fractal 之间 row-major，fractal 内部也 row-major；
ZN：fractal 之间 row-major，fractal 内部 column-major；
NZ：fractal 之间 column-major，fractal 内部 row-major。

对 FP16，基本 fractal 是 16×16，正好 512 B。以 C = A @ B 为例，Cube 不是去对任意地址执行矩阵乘，而是先把 A/B 准备成硬件适合读取的 L0A/L0B fractal，再执行 Mmad。

因此更准确的理解是：

复杂寻址和 layout 变换尽量在“喂 Cube”阶段解决；Cube 本体尽量只吃规整 fractal。

这和 CUDA 中把 global-memory 数据先排成 ldmatrix / Tensor Core 友好的 shared-memory layout，再执行 MMA 很相似。

3. 第一阶段 Microbench：先把比赛会用到的 primitive 测出来

第一阶段暂时不追 Cube 峰值，先完成下面 6 个 benchmark：

Microbench	主要测量对象	直接对应的算子类型
`copy`	GM↔UB / MTE	`cat`、纯搬运
`vector`	Add / Mul / Compare	`square`、`gt`
`reduce`	Vector reduction	`sum`
`stride-transpose`	stride / layout 重排	`permute`
`indexed-add`	irregular access / atomic	`index_add`
`pipeline`	MTE + Vector overlap	几乎所有 Vector kernel

最终希望形成一张属于目标机器自己的能力表，而不是只记厂商峰值数字。

4. Microbench 1：GM ↔ UB Copy

最小数据通路：

GM src
  ↓ DataCopy
 UB
  ↓ DataCopy
GM dst

建议 sweep：

problem size:
1 KiB, 4 KiB, 16 KiB, 64 KiB,
256 KiB, 1 MiB, 4 MiB, 16 MiB,
64 MiB, 256 MiB

同时 sweep:
- tile size
- block/core 数量
- single buffer / double buffer

对 N 个、每个 s 字节的元素，若 kernel 完成一次 GM read 和一次 GM write，则有效带宽定义为：

BW_eff = 2 * N * s / time

这条曲线会成为后续 cat / square / gt / permute 的最重要 bandwidth baseline。

小尺寸主要受 launch / pipeline latency 影响；尺寸足够大后，曲线应该逐渐进入带宽平台区。

5. Microbench 2：Vector Add / Mul / Compare

Square / Mul

GM x
 ↓
UB
 ↓ Mul(x, x)
UB
 ↓
GM y

若乘法按 1 FLOP 计：

traffic = 2 * N * s
FLOPs   = N
AI      = 1 / (2s)

FP16 下：

AI = 0.25 FLOP/B

因此纯 square 通常更应该首先拿 effective bandwidth 去分析，而不是追求一个看起来很高的 Vector FLOPS 数字。

Add

z = x + y 的 GM traffic 约为：

read x + read y + write z = 3 * N * s

FP16 下理论算术强度仅：

1 / 6 FLOP/B

Compare / gt

gt 更适合记录：

GElem/s
useful GB/s

而不是 FLOPS。输出 dtype / mask 的实际表示需要按最终 kernel 实现统计真实流量。

建议额外测 Vector repeat 边界附近的尺寸，例如 FP16：

127, 128, 129
255*128
255*128 + 1

用来观察尾块、mask 和多指令切分的开销。

6. Microbench 3：ReduceSum

最小模型：

GM
 ↓
UB tile
 ↓
Vector Reduce
 ↓
partial sum
 ↓
最终结果
 ↓
GM

Reduction 输出很小，所以不要用“输入+输出总字节”掩盖问题，优先定义：

BW_reduce = input_bytes / time

即每秒真正能吃掉多少输入数据。

需要 sweep：

输入长度 N；
每个 core 负责的数据量；
core 数量；
tile size；
partial reduction 的组织方式。

目标是找到 N/core 的甜点，并判断大尺寸时 reduction 能接近多少 GM read bandwidth。

7. Microbench 4：Stride 与 Transpose

Stride

测规则 strided access：

blockStride = 1, 2, 4, 8, 16, ...

输出：

stride → useful GB/s

这样可以直接看到 burst 破碎、有效带宽下降的速度。

Transpose

再写一个以小块 transpose 为 primitive 的大矩阵重排测试：

GM layout A
 ↓
UB tile
 ↓ transpose/reorder
UB tile
 ↓
GM layout B

建议至少覆盖：

1024²
4096²
8192²

这一组 microbench 直接对应 permute：它通常不是“算力问题”，而是 layout + 数据搬运问题。

8. Microbench 5：Indexed Add

目标语义：

j = index[i]
dst[j] += src[i]

不要只测一种 random index。至少准备四类模式：

A. contiguous
   index = 0,1,2,3,...

B. permutation
   地址随机，但不重复

C. random
   随机地址，允许不同程度的局部性

D. hotspot
   大量更新集中到少量 dst

这样可以把性能问题拆开：

A → baseline
B → 重排 / locality penalty
C → irregular memory penalty
D → atomic contention penalty

CANN 8.5.0 提供 Gather 类能力，同时也可以通过原子写回能力实现更新。这个 benchmark 的意义不是只得到一个数字，而是找出 index_add 到底被地址不规则性、GM 流量还是原子冲突中的哪一个限制。

推荐同时记录：

GUpdate/s
input GB/s
output/update GB/s
collision rate

9. Microbench 6：Double Buffer Pipeline

用最简单的 square 写两个版本。

Sequential

CopyIn tile0
Compute tile0
CopyOut tile0

CopyIn tile1
Compute tile1
CopyOut tile1

时间模型近似：

T_tile = T_load + T_compute + T_store

Pipelined

time →

MTE in : LOAD0   LOAD1   LOAD2   LOAD3
Vector :      MUL0    MUL1    MUL2    MUL3
MTE out:          STORE0  STORE1  STORE2

理想 steady state 下：

T_tile ≈ max(T_load, T_compute, T_store)

这一步要真正理解 TPipe / TQue 的意义：不仅是“写起来方便”，而是在表达 producer-consumer、内存复用和流水同步。

10. 比赛算子的 primitive 映射

当前这类 PyTorch 算子可以先按底层瓶颈分成：

torch.cat
    → MTE / copy

torch.square
    → copy + Vector Mul + pipeline

torch.gt
    → copy + Vector Compare + pipeline

torch.sum
    → Vector reduction

torch.permute
    → stride + transpose / layout movement

torch.index_add
    → irregular access + update/atomic

因此第一阶段 6 个 microbench 完成后，比赛 kernel 大部分工作应该变成：

把已经测清楚的 primitive 按算子语义组合起来，再针对 shape 做 tiling。

而不是到比赛时才重新学习 Ascend C 的基础数据通路。

11. 第二阶段再进入 Cube Microbench

等 Vector/MTE 基线稳定后，再增加：

GM → L1
L1 → L0A/L0B
Mmad
L0C → GM

并分别测：

GM→L1 带宽；
L1→L0A / L1→L0B 搬运；
单个 Mmad 基本吞吐；
K-loop accumulation；
Cube double buffer；
完整 GEMM 的 MTE/Cube overlap。

这个阶段再开始研究 baseM/baseN/baseK、fractal layout 和 L0 容量的共同约束。

12. 版本与实验纪律

S5 的 Ascend C 环境搭建手册可以作为工程结构、安装和官方 AddCustom 示例的参考，但手册使用的是 CANN 8.1 RC 系列，不能拿来证明 CANN 8.5.0 的具体性能或接口行为。
当前目标是 CANN 8.5.0 + 910B2；在 910B4-1 上得到的结果要保留芯片型号、频率、CANN/driver 版本等 metadata。
每个 benchmark 都应该保留：shape、dtype、tile、block/core 数、buffer depth、运行时间、有效带宽/吞吐和 profiler 证据。
优化前先建立最简单、可验证正确性的 baseline，再逐项引入 stride、double buffer、更多 core 和特殊 API。

13. 当前 Milestone

第一轮只完成：

[1] Copy
 ↓
[2] Add / Mul / Compare
 ↓
[3] ReduceSum
 ↓
[4] Stride + Transpose
 ↓
[5] Gather / Indexed Add / Atomic
 ↓
[6] Double Buffer Pipeline

完成后应得到一张机器实测能力表：

GM copy               xxx GB/s
Vector Add             xxx GElem/s
Vector Mul             xxx GElem/s
Vector Compare         xxx GElem/s
ReduceSum              xxx GB/s input
Transpose              xxx GB/s useful
Stride=1/2/4/...        xxx GB/s
Indexed update         xxx GUpdate/s
Atomic hotspot         xxx GUpdate/s
Pipelined square       xxx GB/s

这张表会成为后续比赛算子优化和 Roofline 分析的第一层基线。

参考

CANN 8.5.0 版本配套关系：https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/850/releasenote/releasenote_0000.html
CANN 8.5.0 Vector 高维切分 API：https://www.hiascend.com/document/detail/en/canncommercial/850/opdevg/Ascendcopdevg/atlas_ascendc_10_0022.html
CANN 8.5.0 架构规格（L0A/L0B/L0C 推荐格式与对齐）：https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/850alpha001/opdevg/Ascendcopdevg/atlas_ascendc_10_0011.html
Ascend C 数据排布格式（ZZ / ZN / NZ）：https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/82RC1/opdevg/Ascendcopdevg/atlas_ascendc_10_0099.html