Assignment 4（Part 2）：融合的 Convolution + MaxPool 层

📖 本篇是 Assignment 4 实验报告的 Part 2；Part 1（Ascend C 编程接口 —— 向量加与矩阵转置）见 asst4_part1.md。 两篇共用的环境配置在 Part 1 的 §0，本篇不再重复。

2. Part 2: Implementing a Fused Convolution - Max Pool Layer（70 分）

算子常见数据格式:

字母含义:

同一个字母在不同格式里含义可能不同——尤其 C0 在 5HD 与"分形块"里数值一样（都是 16）但指的不是同一件事。

分形（Fractal）是 Cube 的最小计算粒度：half 为 16×16（256 元素 / 512 B），int8 为 16×32 或 32×16——2201 与 3510 全系列都遵循，差异只在 2201 的 L0A 用 Zz，3510 的 L0A 用 Nz。

NPU 架构版本 2201（A2/A3）下 Cube 的分形格式要求（Ascend C 编程指南《Cube 矩阵计算编程》）：

举个例子：NZ = FRACTAL_NZ，是升腾 Cube 单元专用的矩阵分形格式——把大矩阵切成 16×16（或 16×8 / 16×32）的小块，然后按"块间 N 字、块内 z 字"重排内存。官方叫它 Nz（大 N 小 z）格式。

逻辑矩阵 (H × W)
   ┌─────┬─────┬─────┐
   │ F00 │ F01 │ F02 │     每个 F 是一个 H0×W0 的分形
   ├─────┼─────┼─────┤
   │ F10 │ F11 │ F12 │
   └─────┴─────┴─────┘

内存里的存放顺序（column major，块间）：
     F00 → F10      → F01 → F11      → F02 → F12 ...
   └─ 第0列从上到下 ─┘ └─ 第1列 ─┘
   轨迹像字母 N

每个分形内部（row major，块内）：
   逐行走，轨迹像字母 z

CNN 的常见格式

图片只有两种排法：

• NCHW（通道在第 2 位）：PyTorch 默认、cuDNN/TensorRT、CANN 的 data_format=NCHW；
• NHWC（通道在最后）：TensorFlow/Keras 默认、PyTorch 的 channels_last、CANN 的 data_format=NHWC。

卷积核有两种习惯：

2.0 conv2d_v2 的 im2col 实现调研

分析 一下 ops-nn/conv/conv2d_v2 使用im2col做卷积的实现方式。该kernel不支持 Atlas A2/A3系列。

以下面这个图片为例子：

2.0.1 七步流水

halo: 算 N 行输出需要 N+2 行输入（filter 高 3 时），多出来的那 2 行就叫 halo，列方向同理。

搬运指令 DataCopy 以 32 B（一个 DataBlock） 为单位工作：源地址要 32 B 对齐、每次连续搬运的长度要是 32 B 的整数倍。H5D格式最后一维16个half，恰好是 32B 整数。

2.0.2 排布转换：从 (1,256,224,224) 到 符合 cube 输入的样子

cube 读矩阵的最小单位是 16×16 的小方块（fp16；fp32 是 8）。而 (1,256,224,224) 这种排布里，同一个通道的一整张图是连续的一大片，不同通道之间隔着 224×224 那么远——要凑一个小方块得跳着读几十次。所以官方把"通道"这一维按 16 个一组拆开、挪到最后：

两点注意：

1. 这是"换排布"，不是"改形状"：元素总数一模一样，只是"谁挨着谁"变了；
2. 这一步可以不在主机上做：它是图里挂的一个设备算子。上游算子的输出本来就是这种排布时，图优化会把整段省掉——所以它有时存在、有时不存在。

2.0.3 三级分块：核间 → L1 → L0

第一级：核间切分（每个核算哪一片输出）。用上面的例子、20 个核：输出通道切成 2 块（每块 128 个）、输出行切成 10 块（每块约 23 行），2 × 10 = 20 个核。优先级是"重叠越小越优先"：批量 → 输出行 → 输出通道 → 输出列 → 输入通道。

每个核由此得到它要读的输入范围：全部 256 个输入通道 ×（23 + 2）行 ×（222 + 2）列 ≈ 2.87 MB（多出来的 2 行 2 列就是 3×3 窗口的 halo）。

第二级：L1 分块。2.87 MB 塞不进 512 KiB 的 L1，所以每个核是"一块一块搬、一块一块算"，切三刀：

账：图块 8 × 4 × 114 × 16 = 58,368 个元素 = 116,736 B ≈ 114 KiB；权重块 128 × 128 = 16,384 个 = 32,768 B = 32 KiB；合起来 146 KiB，开双缓冲 292 KiB，512 KiB 的 L1 放得下（剩下的留给别的）。

第三级：L0 分块（也就是第五步）。L1 里那块是"图"，cube 要的是"矩阵块"，得再切一次：

地址要求：

2.0.4 循环骨架

// 每个核：
for (输出行块 : 本核负责的行)                       // 由核间切分决定
  for (输出列块 : 本核负责的列)
    for (输出通道块 : 128 个一条) {
        结果缓存清零;
        for (输入通道块 : 128 个一组) {              // 收缩维，结果在结果缓存上累加
            [第四步] 显存 -> L1：搬这块图（含 halo）+ 对应的权重
            for (窗口行 = 0..2)
              for (窗口列 = 0..2) {
                  [第五步] L1 -> 图片输入缓冲：按窗口位置展开成矩阵
                  [第五步] L1 -> 权重输入缓冲：取该窗口位置那份权重
                  [第六步] 两块矩阵相乘，累加到结果缓存
              }
        }
        [第七步] 结果缓存 -> 显存（可顺手加 bias）
    }

2.0.5 三个可供参考的设计

1. 主机不搬数据：官方 host 侧只有"指针 + 分块参数"，Conv2DTilingData 里没有任何数据指针；搬运动作全发生在片上（第四步、第七步），形状转换与展开都搭在搬运路上。
2. L1 存"图"、L0 存"矩阵"：展开后的矩阵带 FH·FW 倍冗余（同一份数据在九个窗口位置各出现一次），这种冗余只在"L1 → L0"那一瞬间存在；L1 及以下（含显存）始终只存不冗余的图。我们把展开提前到了主机上，显存里因此多存了 3 倍的数据（按列位移的 3 份）——这是两条路最本质的差别。
3. 双缓冲是重叠的全部秘密：第四步与第六步交替进行，开两块缓冲后"搬下一块"和"算当前块"就能同时进行；tiling 里所有"每级切多大"的参数，本质都在调"一块多大、能开几块"。

2.1 Matmul 高阶 API 使用

计算语义与使用流程

Matmul 的计算公式：C = A × B + Bias：

• A 为左矩阵，形状 [M, K]；B 为右矩阵，形状 [K, N]；C 存放结果，形状 [M, N]；Bias 形状 [1, N]，对 C 的每一行广播相加。
• 轴定义：M 轴 = A 矩阵纵向；K 轴 = A 横向 = B 纵向；N 轴 = B 横向；尾轴 = 矩阵最后一个维度。

Kernel 侧官方五步（Ascend C 程序里就是按这个顺序写）：

1. 创建 Matmul 对象          Matmul<aType, bType, cType, biasType> mm;   ← 模板实例化，无运行期参数
2. 初始化 Tiling              mm.Init(cubeTiling, tpipe);   ← 按 Tiling 做资源划分
3. 设置 A / B / Bias          mm.SetTensorA(aGlobal); mm.SetTensorB(bGlobal); mm.SetBias(biasGlobal);
4. 完成矩阵乘                 mm.Iterate();  mm.GetTensorC(cLocal);   ← 每次算一块 baseM×baseN
5. 结束                       mm.End();

1. 模板参数 MatmulType（决定 A/B/C/Bias 的类型信息）

template <AscendC::TPosition POSITION, CubeFormat FORMAT, typename TYPE,
          bool ISTRANS = false, LayoutMode LAYOUT = LayoutMode::NONE,
          bool IBSHARE = false, TPosition SRCPOS = TPosition::GM>
struct MatmulType { ... };

更多信息参考官方手册 Matmul性能优化策略总览

1. 对象创建与初始化

输入矩阵A/B：MatmulType<TPosition::GM, CubeFormat::ND 或 NZ, T>；C 可以落在 UB（TPosition::VECIN）

1. Tiling 结构体（TCubeTiling）关键字段

Tiling 侧（Host）生成、Kernel 侧只读：

更多诸如 模板参数 MatmulConfig（可选，配置模板与关键数值），成员函数参数明细参考资料: Matmul使用说明。

2.2 正确性实现

2.2.1 输入数据格式处理

符号与形状速查

① 用户侧（契约，与 PyTorch 一致）

② 派生尺寸（全部由 ① 推出；stride = 1、无 padding）

输出尺寸为什么要对齐到 16：

1. matmul 的 N 直接取它 ： matmul 的 N 直接取整输出行， Cube 一次算一整行输出。而Cube 的最小计算粒度：half 为 16×16（256 元素 / 512 B）
2. 32 B / DataBlock 对齐 —— 16 个 fp16 = 32 B 。

多出来的列（OWp − W_out）：是卷积算出来的无效值 （来自右侧 padding 列 ）—— 所以 host 在最后把 pad 列裁掉 （输出从 poolOutWidth 裁到 W_pool ）。

内部格式（HBM 里的实际排布，由前置重排产出；计时循环内不再重排）

图片：X′ (B, H, FW, C_in, OWp)     ← 每个输入行存 FW 个列位移副本
权重：W′ (FH*FW, C_out, C_in)      ← 每个 (i,j) 切片内部 (C_out,C_in) 连续

• 图片的段 (b, r, j) 内是连续的 (C_in, OWp)：段[c][ow'] = X[b][c][r][ow' + j]（ow' + j ≥ W 的列补 0）；
• 权重的第 (i,j) 个切片内是连续的 (C_out, C_in)：W′[(i·K_w + j)][co][c] = W[co][c][i][j]。

为什么必须是这两种排布 :

1. 矩阵乘的 A/B 只能通过一个基址传入（SetTensorA/B 不带 stride），框架按 基址 + 行号×行长度 + 列号 取数。

   因此两块都必须是"连续 ND 块"——即行距等于行长度、内部不夹空洞的行主序子矩阵。

1. 两个方向位移的处理方式不同（这是本设计的关键）：

高阶API matmul 要求读取的地址在 HBM里连续。这里只 shift(input, (i, j)) 按道理要存 FW* FH 次副本，但是行位移是免费的，只有列位移不免费：

取 (i,j) 的段  →  (r = oh + i, j)

  (i,j)      段索引
  ─────────────────────────────────────
  i=0,j=0 →  r=oh,   j=0
  i=0,j=1 →  r=oh,   j=1
  i=0,j=2 →  r=oh,   j=2      ┐
  i=1,j=0 →  r=oh+1, j=0      │ 行方向只是换 r
  i=1,j=1 →  r=oh+1, j=1      │ → 换基址，零成本 ✓
  i=1,j=2 →  r=oh+1, j=2      │
  i=2,j=0 →  r=oh+2, j=0      │
  i=2,j=1 →  r=oh+2, j=1      │
  i=2,j=2 →  r=oh+2, j=2      ┘

于是 K_h·K_w 次 partial sum 对应的是 K_h 个行 × FW 份副本 = 9 个段，而存储只需 FW 份（不是 9 份）。

2.2.2 分块设计

约束核对

三个片上缓存各自给出一条约束（ε = 每元素字节数；L0C 存的是 fp32 累加值，所以按 4 B 算）：

一次 tile 的矩阵乘形状是：

A = (baseM, baseK)   ×   B = (baseK, OWp)   →   C = (baseM, OWp)
    权重切片               图片段                 结果（L0C 累加）

切分的标准（按什么分、为什么这么分）

由此得到四条量化准则：

① baseM  = 不超过 { C_out, L0A 上限, L0C 上限, UB 预算 } 的最大 2 的幂
② baseK  = 不超过 { C_in,  L0A 上限, L0B 上限 }           的最大 2 的幂
③ row_blocks：让 B × M_blocks × row_blocks 是 20 的整数倍（各核等量），且尽量小（R 尽量大）
④ 分配：工作项 = (b, m_block, row_block)，按 item = blockIdx; item += 20 轮转

为什么是这个优先级：切分轴的"重叠/通信代价"递增 —— C_out 零代价 → 输出行有 halo → C_in 需要累加（所以干脆不切核间）。先把零代价的轴切够，再动有代价的轴；剩下的收缩维放到核内循环里解决。

① OWp：只由输出宽度决定

W_out = W − K_w + 1          ← 卷积核宽度在这里起作用
OWp   = align16(W_out)       ← 矩阵乘的 N；对齐到 16 是为了分形/32 B

② baseM：由 C_out、L0A/L0C 上限、UB 预算三者共同夹出来（取不超过上限的最大 2 的幂）

baseM ≤ C_out                     ← 本核的输出通道块不能超过总通道数
baseM × baseK ≤ 64 KB / ε         ← L0A
baseM × OWp  ≤ 128 KB / 4         ← L0C（fp32 累加）
2 × baseM × OWp × ε + C_out × ε ≤ 192 KB   ← UB（结果 + 池化暂存 + bias）

注意最后一条里出现了 OWp —— 所以 K_w 会经由 W_out → OWp → UB 预算间接影响 baseM：输出行越宽，baseM 就越容易被 UB 卡住。

③ baseK：由 L0A/L0B 联立（收缩维，C_in > baseK 时切段在 L0C 累加）

baseK ≤ C_in
baseM × baseK ≤ 64 KB / ε         ← L0A
baseK × OWp  ≤ 64 KB / ε          ← L0B（图片段，大图时是紧约束）

两个用例算出来就是

用例的完整规模：

C_in 与 C_out 各自决定"要不要切"：

M 块数 = C_out / baseM      ← C_out 决定核间在 M 方向能切几块
K 段数 = C_in  / baseK      ← C_in 决定收缩维要分几段（在 L0C 累加）
每个 tile 的 matmul 次数 = K 段数 × K_h·K_w

读法：C_out 只影响"核间切几块 M"（每块内部的矩阵乘形状不变），C_in 只影响"收缩维切几段 K"（每段一次矩阵乘、在 L0C 累加）——两者都不改变 A/B 的单块形状，只改变块的数量。

其余两个量与 K_h 的关系（与 baseM/baseK/OWp 无关）：

一个 tile 覆盖 R 个池化行 = R × pool 个卷积输出行     ← 由 H_out 与核数分配决定
该 tile 需要的输入行数 = R × pool + K_h − 1           ← K_h 只在这里起作用（halo）

一次矩阵乘只算一行输出 ：

一次 matmul：  C = (baseM, OWp)
                    └── N = OWp = 一行输出

它需要的输入：  3 行（K_h = 3）
                靠 9 次 (i,j) 迭代取 9 个段 = 3 行 × 3 个列副本
                  i=0 → 段 (r = oh+0, j)   ┐
                  i=1 → 段 (r = oh+1, j)   ├─ 3 行输入
                  i=2 → 段 (r = oh+2, j)   ┘
                  j=0,1,2 → 每行的 3 个列副本

3 行是顺序读进来的再累加到L0C 上的。

2.2.3 算法描述

shift reduction 算法本体：把卷积写成一串矩阵乘的累加。

对某个输出行 oh 和窗口位移 (i, j)：

out[co, oh, :] += A(i,j) × B(oh, i, j)
                  (baseM × C_in)   (C_in × OWp)

  A(i, j)     = W′[(i·K_w + j)][cout0 : cout0+baseM][:]    ← 权重切片（连续 ND 块）
  B(oh, i, j) = X′[b][oh + i][j][:][:]                     ← 图片段（连续 ND 块）
  收缩维       = C_in（被求和消掉）

落到昇腾上：一个工作项 = (b, M 块, 池化行)，其内部是

for 本 tile 的每个卷积输出行 ohLocal:
    for i in 0 .. K_h-1:
      for j in 0 .. K_w-1:
          SetTensorA( W′ 的第 (i,j) 切片 + cout0 偏移 )      // A = (baseM, C_in)
          SetTensorB( X′ 的段 (b, ohLocal + i, j) )          // B = (C_in, OWp)
          while (Iterate(enPartialSum = (i != 0 || j != 0)))
              GetTensorC(部分和 → UB)
    GetTensorC(最终和) → UB → +bias → maxpool → 写回 GM

三个关键点

1. 9 次累加都发生在 L0C（部分和）

   第 (i,j)=(0,0) 次 enPartialSum = false（先清零这块 L0C），其后 8 次为 true（累加到同一块）。

   所以一个输出行的完整结果只在最后一次从 L0C 取出 —— 这也是"卷积结果不写进输出张量"的实现方式。C_in > baseK 时，K 方向的切段同样累加在 L0C，与 (i,j) 的累加叠加。

2. 两个位移的处理方式不同（见 2.2.2）

   行位移 i 靠段索引（ohLocal + i，换基址即可，零成本）；列位移 j 靠预物化的 FW 份副本。于是 K_h·K_w 次矩阵乘只需要 FW 份存储（3×3 时 3 份，而不是 9 份）。

3. maxpool + bias 的融合

   GetTensorC 把卷积行（baseM × OWp）落到 UB → 行方向 Max（pool=2 时取相邻两卷积行）→ 列方向 GatherMask（取偶/奇索引，紧凑输出）→ +bias → 写回。

   bias 放在 max 之后是正确的：max(a+b, c+b) == max(a, c) + b。pool = 1 时两次 max 退化成恒等，走同一份代码路径（所以 pool ∈ {1,2} 由同一实现保证一致）。

2.2.4 上板验证命令

# ── 0. 环境准备（每个新 shell 一次）──────────────────────────────
export ASCEND_INSTALL_PATH=/usr/local/Ascend/cann-8.5.0    # 以本机 CANN 实际路径为准
source $ASCEND_INSTALL_PATH/bin/setenv.bash

# ── 1. 编译 ─────────────────────────────────────────────────────
# 手动编译（-v 是机型：A2 用 npu-smi 的 Name → Ascend910B4；
#        A3 用 ChipName_NPUName → Ascend910_9362，见 README-ASCEND.md 判机型）
bash run.sh --run-mode npu --build-only -v Ascend910B4 2>&1 | tail -3;  
# 需要查看错误的的话：
bash run.sh --run-mode npu --build-only -v Ascend910B4 2>&1 | tee /tmp/build.log >/dev/null; echo "=== conv2d.asc 的报错 ==="; grep -nE "conv2d\.asc:[0-9]+:[0-9]+: error" /tmp/build.log | head -20; echo "=== 所有 error（带上下文）==="; grep -nE "error:" -A 5 /tmp/build.log | head -60
# 产物：build-npu/ascendc_kernels

# ── 2. 对比（正确性对比，numpy + torch 双基准）──────────────────
# 改了代码后 → --build 强制重新编译；没改代码重复跑 → 直接 --quick/不带参数。
# ★ test_harness 现在自带"源码比二进制旧就重编"的护栏（会打印 [build] 那一行）：
#   原先只有"二进制不存在"才编译（if args.build or not os.path.exists(binary)），
#   于是改完 conv2d.asc 直接跑测试会静默跑旧二进制 —— 曾连续三轮拿到逐字相同的旧结果。
python3 test_harness.py --build --quick   # 强制重编 + 小图对比（bias=0/1 各 fp16/fp32）
python3 test_harness.py --quick           # 重复跑（二进制不旧就跳过编译）：只跑小图对比

python3 test_harness.py                   # 全量：再加 224×224 大图 + bias（二进制不旧则不重编）
python3 test_harness.py --test_maxpool    # 含 maxpool（pool=2 → 15×7 / 111×111）
python3 test_harness.py --perf            # 性能：6 个用例的 median / GFLOP/s / GB/s
python3 test_harness.py --show-tiling     # 把 kernel 的 [TILING]/[TILING-B] 诊断打出来（tiling 取证）

# 手动单跑 + 落盘，再对比（可调 tile 切分验证行切片正确性）：
./build-npu/ascendc_kernels --dtype float32 --batch 4 --in-channels 128 --out-channels 128 \
    --height 32 --width 16 --filter 3 --pool 1 --block-dim 1 --warmup 1 --iters 1 --bias \
    --dump-dir /tmp/asst4p2
# → /tmp/asst4p2 下 x.bin / w.bin / bias.bin / out.bin

# 手动跑记得 LD_LIBRARY_PATH 前缀追加 $PWD/build-npu/lib
# 想验证核间切分无边界问题，把核数改小重跑（--block-dim 1 / 4 / 8），结果应与 20 核时逐元素一致
# 工作项 =(batch, M 块, 池化行) 三分量正交，写回区域互不重叠。

输出结果：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part2$ cd ~/asst4-ascend/part2 && python3 test_harness.py --quick
[info] soc_version = Ascend910B4
[info] 参考实现：numpy + torch 双基准（两者必须一致）

=== 正确性用例 ===
[case] bias=0 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）
[case] bias=1 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）

正确性总判定： [ OK ] 全部通过
atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part2$ cd ~/asst4-ascend/part2 && for P in 1 2; do for D in float32 float16; do echo "=== pool=$P $D ==="; rm -rf /tmp/pv && mkdir -p /tmp/pv; ./build-npu/ascendc_kernels --dtype $D --batch 4 --in-channels 128 --out-channels 128 --height 32 --width 16 --filter 3 --pool $P --dump-dir /tmp/pv >/dev/null; python3 check_conv.py --dump-dir /tmp/pv --dtype $D --batch 4 --in-channels 128 --out-channels 128 --height 32 --width 16 --filter 3 --pool $P --no-bias | tail -2; done; done
=== pool=1 float32 ===
判据：|Δ| > 0.0323753（tol=0.0001 × max|ref|=323.753）
  [ OK ] 215040 个元素全部在容差内
=== pool=1 float16 ===
判据：|Δ| > 6.47093（tol=0.02 × max|ref|=323.547）
  [ OK ] 215040 个元素全部在容差内
=== pool=2 float32 ===
判据：|Δ| > 0.0323753（tol=0.0001 × max|ref|=323.753）
  [ OK ] 53760 个元素全部在容差内
=== pool=2 float16 ===
判据：|Δ| > 6.47093（tol=0.02 × max|ref|=323.547）
  [ OK ] 53760 个元素全部在容差内

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part2$ python3 test_harness.py 
[info] soc_version = Ascend910B4
[info] 参考实现：numpy + torch 双基准（两者必须一致）

=== 正确性用例 ===
[case] bias=0 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）
[case] bias=0 maxpool=0 大图 224×224
  [ OK ] float32: 50466816 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 50466816 个元素全部匹配（rtol=atol=0.01）
[case] bias=1 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）
[case] bias=1 maxpool=0 大图 224×224
  [ OK ] float32: 50466816 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 50466816 个元素全部匹配（rtol=atol=0.01）

正确性总判定： [ OK ] 全部通过

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part2$ python3 test_harness.py --test_maxpool
[info] soc_version = Ascend910B4
[info] 参考实现：numpy + torch 双基准（两者必须一致）

=== 正确性用例 ===
[case] bias=0 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）
[case] bias=0 maxpool=0 大图 224×224
  [ OK ] float32: 50466816 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 50466816 个元素全部匹配（rtol=atol=0.01）
[case] bias=0 maxpool=1 小图
  [ OK ] float32: 53760 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 53760 个元素全部匹配（rtol=atol=0.01）
[case] bias=0 maxpool=1 大图 224×224
  [ OK ] float32: 12616704 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 12616704 个元素全部匹配（rtol=atol=0.01）
[case] bias=1 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）
[case] bias=1 maxpool=0 大图 224×224
  [ OK ] float32: 50466816 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 50466816 个元素全部匹配（rtol=atol=0.01）
[case] bias=1 maxpool=1 小图
  [ OK ] float32: 53760 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 53760 个元素全部匹配（rtol=atol=0.01）
[case] bias=1 maxpool=1 大图 224×224
  [ OK ] float32: 12616704 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 12616704 个元素全部匹配（rtol=atol=0.01）

正确性总判定： [ OK ] 全部通过

2.3 Optimization

2.3.1 先量：当前实测与 torch 对照

对比口径：推理一次的 API + 执行时间。

python3 torch_conv2d_maxpool_baseline.py --kernel（B=1, C_in=C_out=256, k=3，warmup=5 / iters=20，取 median）：

• 910B4 的 fp16 峰值约 295 TFLOP/s → torch 达 47% MFU，本实现只有 5.2%；
• 差距在小图上只有 3.5×，在大图上放大到 9× —— 说明瓶颈不是"算不动"，而是流水/搬运没重叠：小图的问题是单次规模太小、固定开销占比高，大图的问题是每行的 9 次串行搬运 + 权重反复重读。

2.3.2 用 msprof 查算子：三层口径

本实现是两个算子：

1. 前置 transData —— 用户格式 → X′/W′，AIV_ONLY、只在 Init() 里发一次；
2. 主算子 conv+maxpool —— AIC、每轮调用发一次。

查性能分三层：先看"两个算子各占多少"，再钻进主算子内部。

msprof 有四个坑（任意一个都会让采集失败或采错对象，part2/prof.py 都封好了）：

1. 必须带 op 子命令 —— 是 msprof op ...，不是裸 msprof；

2. 输出目录不能是 group/other 可写的 —— msprof 有权限安全检查，/tmp（权限 1777）会被直接拒绝；脚本固定用仓库内 profile/ 并 chmod 700；

3. 不接受 --soc-version —— 传了直接报错（那不是 op 子命令的参数）；

4. 一次只剖析一个算子，且默认抓先发的那个 —— 设备侧现在有两个算子：pretrans_*（W_mat 重排，在 Init 里发一次）和主算子 fused_conv2d_maxpool_*（计时循环里发 warmup+iters 次）。不加选择参数时，msprof 抓到的是先发的 pretrans_* —— 它的 PipeUtilization.csv 里 aic_* 全是 NA（纯向量算子）、aiv_scalar_ratio ≈ 0.99。

   要主算子就加 --kernel-name=<mangled 名字>（必须用 = 形式，空格形式在本版本会报 argument --kernel-name miss value）；名字从二进制里取：

   grep -ao '_Z[0-9A-Za-z_]*fused_conv2d_maxpool_fp16[0-9A-Za-z_]*' build-npu/ascendc_kernels | sort -u
   

   若名字没匹配上，msprof 会把 kernel 整体跳过、连数据都不出（part2/prof.py 对这个情形有专门诊断）。

   也可用 --launch-skip-before-match=N 跳过前 N 次 launch。

第 1 层 · 整体时间线：两个算子各占多少、中间有没有空隙

用 part2/prof.py 采集：

cd ~/asst4-ascend/part2

# ① 前置 transData（默认抓"先启动的那个算子"）→ 产出 profile/pre_op/
python3 prof.py --profile_name pre_op

# ② 主算子（跳过第 1 次 launch 的前置算子）→ 产出 profile/main_fp16/
python3 prof.py --profile_name main_fp16

# ③ 主算子 + 片上细节指标 → 产出 profile/main_l0/
python3 prof.py --profile_name main_l0 \
    --aic-metrics=MemoryL0,MemoryUB,ResourceConflictRatio,L2Cache

# ④ 主算子、连续采集 5 次 → 产出 profile/main_avg/
python3 prof.py --profile_name main_avg --count 5

# ⑤ 只看一眼（不采集，确认命令拼对）
python3 prof.py --profile_name t --no-profile

# 打包带走
tar -czf profile.tar.gz profile/

脚本等价于下面的原始命令（理解原理用；--warmup 1 --iters 5 是故意的：时间线里出现 1 个前置 + 6 个主算子，便于分辨）：

mkdir -p profile && chmod 700 profile        # ★ 见上方"坑 2"
msprof op --output=$PWD/profile \
    ./build-npu/ascendc_kernels --dtype float16 --batch 1 \
    --in-channels 256 --out-channels 256 --height 224 --width 224 --filter 3 --pool 2 \
    --warmup 1 --iters 5

两个算子的区别（kernel 名是 mangled 的，靠这两条判据认）：

这一层看三件事：

1. 前置算子的实际耗时（它是"每张图一次"的代价，报告里单独列）；
2. 主算子之间有没有空隙（有 → host 侧调度或同步开销）；
3. 两个算子是否首尾相接（这是我们的设计：前置一次、主算子多次）。

⚠️ 计时口径：host 打印的 [INFO ] median=... µs 只含主算子 —— 前置算子在 Init() 里、一次性、不进计时循环（符合题目"计时循环里 host 不搬数据"的要求）。所以对比 torch 时比的是主算子的 median，前置算子的开销单独报。

第 2 层 · 算子级：主算子内部是"等搬运"还是"等计算"

cd ~/asst4-ascend/part2

# 主算子（单次）→ profile/main_fp16/
python3 prof.py --profile_name main_fp16

# 主算子 + 片上细节 → profile/main_l0/
python3 prof.py --profile_name main_l0 \
    --aic-metrics=MemoryL0,MemoryUB,ResourceConflictRatio,L2Cache

# 主算子、连续 5 次（看抖动）→ profile/main_avg/
python3 prof.py --profile_name main_avg --count 5

# 或手动（等价；注意 chmod 700 与 --launch-skip-before-match）
mkdir -p profile && chmod 700 profile
msprof op --output=$PWD/profile ./build-npu/ascendc_kernels \
    --dtype float16 --batch 1 --in-channels 256 --out-channels 256 \
    --height 224 --width 224 --filter 3 --pool 2 --warmup 1 --iters 5

采集完脚本会打印 OpBasicInfo.csv / Memory.csv / PipeUtilization.csv 的头几行；也可以自己看：

cat profile/main_fp16/*/OpBasicInfo.csv | head -3    # 含 Current Freq
cat profile/main_fp16/*/Memory.csv | head -8         # 搬运量与带宽利用率

第 3 层 · 片上细节：定位到具体 pipe / bank / 带宽

cd ~/asst4-ascend/part2 && python3 prof.py --profile_name main_l0 \

# 手动加指标（等价，注意仍要 chmod 700 的 profile/ 目录）
msprof op --output=$PWD/profile \
    --aic-metrics=MemoryL0,MemoryUB,ResourceConflictRatio,L2Cache \
    ./build-npu/ascendc_kernels --dtype float16 --batch 1 --in-channels 256 \
    --out-channels 256 --height 224 --width 224 --filter 3 --pool 2 --warmup 1 --iters 5

看什么：

• L0A/L0B/L0C 的读写量是否与理论值相符 —— 大图 fp16 每个输出行理论上要读 A: 9 × 64×128 × 2B = 144 KB、B: 9 × 256×224 × 2B = 1 MB；实测远大于此就说明有重复搬运；
• ResourceConflictRatio 高 → L1 bank 冲突（2.3.4 的 7）；
• L2Cache 命中率 → 相邻输出行的 halo 复用是否生效（2.3.4 的 6）。

三层的分工与迭代节奏：

第 1 层 ── 确认"两个算子各占多少、算子间有无空隙"
   ↓
第 2 层 ── 定"搬运 vs 计算"（mte2_ratio vs mac_ratio）
   ↓
第 3 层 ── 定位到具体 pipe / bank / 带宽
   ↓
改一项优化 → 回到第 2 层做前后对比（看比值是否真的变了）

第一次采集结果：

• 左图（耗时）：一次完整调用 20.29 ms 里，前置 transData 占 16.85 ms（83%）、主算子 3.45 ms（17%）—— 所以优化顺序应该是先前置、后主算子（而不是盯着主算子那 9× 的差距不放）。
• 右图（各单元忙碌比例）——两个算子的瓶颈完全不同：

第一次采集就抓到主要问题（B=1, C_in=C_out=256, 224×224, k=3, pool=2, fp16）：

前置算子占了"一次完整调用"的 81%（16.8 / (16.8+4.0)）：

注意 ratio 的含义是"该单元活跃时间 / 算子总时间"，各单元之和不等于 1 —— 它回答的是"谁在忙"，不是时间切分。主算子里 cube 只有 13.2% 说明 cube 大部分时间在等（等 fixpipe 或等搬运）。

原因是我们用逐元素 GetValue/SetValue 做列位移，规模是 B·H·FW·C_in·OWp = 1×224×3×256×224 ≈ 3850 万次标量操作。

所以优化顺序应当是先前置、后主算子。把逐像素标量拷贝能换成批量向量拷贝 DataCopy。

而主算子这边，关键列与判据：

由此确定两项优化的先后与目标：

2.3.3 TransData 放在 CPU 上

问题：列位移 out[ow] = in[ow + j] 的本质是错位拷贝 —— 源地址偏移 1~2 个元素（fp16 下 2~4 B），而设备侧 DataCopy 要求两端 32 B 对齐、长度是 32 B 整数倍，做不到；

这正是 5HD / Load3D 存在的意义：硬件专为窗口寻址设计。

于是原实现退化成Device 逐元素 GetValue/SetValue（而不能成片拷贝），把 3850 万次标量操作全压在 scalar pipe 上。实测 aiv_scalar_ratio = 0.954，已到 ~1 op/cycle 的极限。

CPU 的 memcpy 不在乎对齐，可以成片拷贝，（且这一步在计时循环之外，）于是把它挪到 RepackImHost。

1. 先整块搬图片矩阵，再在缓存里错位

std::vector<uint8_t> plane(H * W * elemBytes);          // 例如 224²×2B = 100 KB
for (b) for (c) {
    std::memcpy(plane.data(), X[b][c], H*W*elemBytes);  // ★ 一次搬一整个通道平面（连续、大块）
    for (r) for (j) {                                   //   源在 NCHW 里本来就连续 ✓
        memcpy(dst_seg, plane + r*W + j, n*elemBytes);  //   后面的小拷贝全都命中缓存 ✓
        ...
    }
}

1. CPU 多线程 transData

关于实际可用核数，这里是虚拟机，受到配额影响：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part2$ echo "nproc                : $(nproc 2>/dev/null)"
echo "nproc --all          : $(nproc --all 2>/dev/null)"
echo "getconf ONLN         : $(getconf _NPROCESSORS_ONLN 2>/dev/null)"
echo "affinity(实际可用核数) : $(python3 -c 'import os;print(len(os.sched_getaffinity(0)))' 2>/dev/null)"
echo "lscpu CPU(s)         : $(lscpu 2>/dev/null | awk -F: '/^CPU\(s\)/{print $2}' | tr -d ' ')"
echo "cgroup v2 cpu.max    : $(cat /sys/fs/cgroup/cpu.max 2>/dev/null)"
echo "cgroup v1 quota/period: $(cat /sys/fs/cgroup/cpu/cpu.cfs_quota_us 2>/dev/null) / $(cat /sys/fs/cgroup/cpu/cpu.cfs_period_us 2>/dev/null)"
echo "cpuset.cpus          : $(cat /sys/fs/cgroup/cpuset/cpuset.cpus 2>/dev/null) $(cat /sys/fs/cgroup/cpuset/cpuset.effective_cpus 2>/dev/null)"
echo "loadavg              : $(cat /proc/loadavg 2>/dev/null)"
nproc                : 1
nproc --all          : 192
getconf ONLN         : 192
affinity(实际可用核数) : 16
lscpu CPU(s)         : 192
cgroup v2 cpu.max    : 
cgroup v1 quota/period: 1600000 / 100000
cpuset.cpus          : 1-4,29-40 1-4,29-40
loadavg              : 11.03 11.31 12.54 6/6608 11694

虚拟机上只能相信：

• affinity(实际可用核数)：当前进程真正能跑在几个 CPU 上（亲和性掩码 ）
• cgroup v2 cpu.max：格式 <quota> <period> → quota/period = 配额核数 ；若显示 max 100000 = 不限额

优化 1 实测：前置 transData 16.85 ms → 310.6 µs（54×）

同一用例（B=1, C_in=C_out=256, 224×224, k=3, pool=2, fp16）改动前后：

前置算子改后仍是 scalar-bound（0.978），因为设备侧只剩权重重排（TransWeight 同样是逐元素 GetValue/SetValue）——但规模只有 589824 个元素，310 µs 已可忽略（占 7%）；若要彻底消掉，可同样搬到 host。

阶段 2：权重也搬到 host —— 设备前置算子整个删除

阶段 1 之后，设备侧前置算子只剩权重重排（TransWeight，310 µs），而它同样是逐元素 GetValue/SetValue（aiv_scalar_ratio 仍为 0.978）。既然图片已经证明"host 做重排 + 一次上传"完全可行，权重也照做，于是设备侧也可以去除这个算子：

1. 新增 host 侧 RepackWHost（逐元素重排，照抄 TransWeight 的逻辑）：

// W (C_out,C_in,FH,FW) → W′ (FH,C_out,FW*C_in)，W′[i][co][j*C_in+c] = W[co][c][i][j]
const uint8_t *s = src + (((co*Ci + c)*K + i)*K + j) * elemBytes;  // 源：c 的步长是 K*K
uint8_t       *d = dst + (((i*Co + co)*K + j)*Ci + c) * elemBytes;
std::memcpy(d, s, elemBytes);

1. LaunchPreTrans() 变成纯 host 重排 + 两次上传，不再发射任何前置 kernel：

{   // X′
    std::vector<uint8_t> x5Host(x5dBytes_);
    RepackImHost(imHost_, x5Host.data(), ...);
    CHECK_ACL(aclrtMemcpy(x5dDev_, x5dBytes_, x5Host.data(), x5dBytes_, ACL_MEMCPY_HOST_TO_DEVICE));
}
{   // W′
    std::vector<uint8_t> wPackHost(wPackedBytes_);
    RepackWHost(wHost_, wPackHost.data(), a_.outChannels, a_.inChannels, a_.filter, elemBytes_);
    CHECK_ACL(aclrtMemcpy(wDev_, wPackedBytes_, wPackHost.data(), wPackedBytes_, ACL_MEMCPY_HOST_TO_DEVICE));
}

1. 删除设备前置算子：TransDataOp 类、pretrans_fp16/fp32 两个 kernel 入口、两个 _do 发射器（按花括号配对整块删除，之后 grep 确认这些符号的引用数均为 0）。文件从 1435 行降到 1303 行，全局只剩：

434: __global__ __aicore__ void fused_conv2d_maxpool_fp16(...)
450: __global__ __aicore__ void fused_conv2d_maxpool_fp32(...)

连带影响：现在只有一个算子，msprof 采集时不能再带 --skip 1（那会把主算子也跳过去）。prof.py 的 --skip 默认值应改回 0。

性能测试：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part2$ bash run.sh --run-mode npu --build-only -v Ascend910B4 2>&1 | tail -2 && python3 test_harness.py --quick && time ./build-npu/ascendc_kernels --dtype float16 --batch 1 --in-channels 256 --out-channels 256 --height 224 --width 224 --filter 3 --pool 2 --warmup 0 --iters 1 2>&1 | grep -E "^\[HOST\]|median"
[100%] Built target ascendc_kernels
[INFO ] 构建完成：/opt/atomgit/asst4-ascend/part2/build-npu/ascendc_kernels
[info] soc_version = Ascend910B4
[info] 参考实现：numpy + torch 双基准（两者必须一致）

=== 正确性用例 ===
[case] bias=0 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）
[case] bias=1 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）

正确性总判定： [ OK ] 全部通过
[HOST] host 侧重排线程数 = 8（usableCpus：亲和性/cgroup；hardware_concurrency 报 192）
[INFO ] min=19117.516 us  median=19117.516 us  avg=19117.516 us  (blockDim=20, 工作项=222)

real    0m3.261s
user    0m0.736s
sys     0m0.249s

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part2$  python3 torch_conv2d_maxpool_baseline.py --kernel
device  : Ascend910B4   (torch.npu.is_available() = True)
baseline: torch F.conv2d → F.max_pool2d（两次独立调用，**不融合**）  B=1 C_in=256 C_out=256 k=3
口径    : warmup=5 iters=20，每次前后 npu.synchronize()，取 median

dtype          H×W pool |  torch median  GFLOP/s | kernel median  GFLOP/s     加速比    端到端 torch       端到端 我们 |   torch 多搬
-------------------------------------------------------------------------------------------------------------------------
float16      32×16    1 |         153.4  3230.51 |         437.0  1133.75   0.35× |         0.9ms        2.26s |       0.4 MB
float32      32×16    1 |         178.7  2773.14 |         170.5  2905.52   1.05× |         1.2ms        2.25s |       0.9 MB
float16    224×224    1 |         429.5 135348.64 |        2451.3 23716.96   0.18× |         2.4ms        3.26s |      50.5 MB
float32    224×224    1 |         997.8 58265.06 |        5453.8 10660.08   0.18× |         4.4ms        2.25s |     100.9 MB
path string is NULLpath string is NULL.float16    224×224    2 |         480.7 120945.78 |        2347.9 24761.34   0.20× |         3.2ms        3.26s |      50.5 MB
float32    224×224    2 |        1192.4 48758.22 |        4676.8 12431.13   0.25× |         4.5ms        3.26s |     100.9 MB

注：'torch 多搬' = 不融合时多出的中间张量往返（conv 结果写一遍 + 读一遍）= 2 × B·C_out·H_out·W_out·ε；融合 kernel 的'最小访存'不含这一项。

那 2~3 s 的"端到端"是 ACL/显存的一次性固定开销，不是 host 重排。

最初我们以为"端到端慢"是 host 重排（逐通道 memcpy）导致的，于是做了两轮优化：

1. 逐平面整块搬（HBM 读从 B·H·K·C ≈ 17.2 万 次散读降为 B·C = 256 次大块拷贝）
2. 多线程（线程数取容器真实可用核数的一半；容器里 hardware_concurrency() 会虚报宿主核数 192，实测 sched_getaffinity/cgroup 配额才是 16 → 用 8）。

结果 real 从 3.267 s 只降到 3.261 s —— 几乎没有变化。这否定了原猜想，并指向真正的来源：

数据量差 60 倍，耗时只差 40% —— 若拷贝主导，两者应差几十倍。更有力的证据是：同一个用例、同一份二进制，两次运行分别得到 2.25 s 与 3.26 s（本机 loadavg ≈ 11，差值属环境噪声）。因此这 2~3 s 与重排无关，而是进程级的固定开销：

aclInit + aclrtSetDevice（加载驱动、建 device 上下文）
+ aclrtMalloc（X′ 77 MB、原图 25.7 MB、输出、workspace，HUGE_FIRST 需凑大页）
+ aclFinalize / aclrtResetDevice（拆上下文）

这几项任何昇腾程序都要付（torch_npu 在 import 和首次上设备时已经付过），且全部在计时循环之外，所以不污染 median：

• host 打印的 median = 一次 launch + aclrtSynchronizeStream，只含主算子 —— 这是与 torch 的 F.conv2d + F.max_pool2d 一一对应的口径；
• 而"整个进程墙钟"含上面那笔一次性初始化，不能与 torch 的进程内计时并列比较（否则会得出"慢 4000 倍"的错误印象）。

据此把基准脚本的那一列从"端到端 我们"改名为"进程墙钟"，并在表尾加了口径说明（ACL 初始化 / 77 MB 显存分配 / 收尾占 2~3 s 且与用例规模无关；对比还是使用 kernel median）。

为什么 CPU 侧反而快：重排本质是数据搬运，CPU 有 memcpy（不在乎 32 B 对齐）和多线程，而设备侧受制于"DataCopy 要求两端 32 B 对齐、长度为 32 B 整数倍"—— 列位移恰好破坏对齐，于是设备侧只能退化成逐元素标量读写，把标量流水打满（aiv_scalar_ratio ≈ 0.99），16.85 ms 几乎全是标量指令的开销。

2.3.4 把 FW 个列位移副本折进 K 维

问题：GetTensorC 在 while(Iterate) 里被调 9 次/输出行 —— 每算完一个 (i,j) 就把 L0C 的部分和搬一遍。profile 显示主算子 fixpipe 占 59.8% 而 cube 只有 13.2%（cube 大部分时间在等 fixpipe）。

关键：SetTensorA/B 只是告诉库"这块数据在 GM 哪里、逻辑形状多大" —— 它们不一次性进 L0 。库拿到之后会沿 K（收缩维） 自己切段、流式喂进 L0 。

改法：X′ (B,H,FW,C_in,OWp) 里同一行的 FW 个段本来就是连续的，于是可以把它们当成一个 (FW·C_in, OWp) 的矩阵读 —— 列位移从"9 次独立矩阵乘"变成"K 维里的 3 段"：

改前： 9 次矩阵乘   A=(baseM, C_in)      B=(C_in, OWp)         → fixpipe 9 次/行
改后： 3 次矩阵乘   A=(baseM, FW·C_in)   B=(FW·C_in, OWp)      → fixpipe 3 次/行
                   （每个 i 一次；K=768 由 Matmul 库内部按 baseK 切段、累加在 L0C）

改前：

for (i = 0; i < FH; ++i)          // 3 次
  for (j = 0; j < FW; ++j) {      // × 3 次  → 共 9 次
      SetTensorA(wGm[WSliceOff(i, j, cout0)]);   // A = (baseM, C_in)   = 128 × 256
      SetTensorB(imGm[imOff(i, j)], false);      // B = (C_in, OWp)     = 256 × 224
      while (Iterate(enPartialSum = (i || j))) GetTensorC(dst);   // 9 次取 C
  }

改后：

for (i = 0; i < FH; ++i) {                        // 只剩 3 次
      SetTensorA(wGm[WSliceOffI(i, cout0)]);     // A = (baseM, FW*C_in) = 128 × 768   ← 3 倍大
      SetTensorB(imGm[imOffI(i)], false);        // B = (FW*C_in, OWp)   = 768 × 224   ← 3 倍大
      while (Iterate(enPartialSum = (i != 0))) GetTensorC(dst);   // 3 次取 C
}

优化 2 实测

先通过数据可以确定一件事情：

main_kfold（优化2 大图，20 AIC + 40 AIV 的均值）
  AIC: L0C_to_L1_datas = 0.000 KB       ← 根本没有 L0C→L1
       L0C_to_GM_datas = 7123.2 KB      ← ★ 非零：fixpipe 在往 GM 搬
  AIV: GM_to_UB_datas  = 3562.1 KB      ← ★ 非零：再从 GM 搬回 UB

  ★ 3562.1 × 2（每个 AIC 对应 2 个 AIV）= 7124.2  ≈ 7123.2   ← 完全对上 ✓✓
ia_small 同样：L0C_to_GM = 44.0，GM_to_UB = 22.25 × 2 = 44.5 ✓

A2 上 CType=VECIN 就是 L0C →(fixpipe)→ GM →(MTE2)→ UB —— 不是直接 L0C→UB 。而且 GM_to_UB 恰好是 L0C_to_GM 的一半（2 个 AIV 各读一份 ） —— 数字严丝合缝 。

• 左图：主算子 Task Duration 3445 → 2323 µs（1.48×，host 侧 median 3983 → 2846 µs）；
• 右图：各单元忙碌比例。最忙的一始终是 AIC scalar（深红） —— 改前 0.912（比 fixpipe 的 0.598 更高），改后 0.689。

把忙碌比例换算成绝对时间更清楚：

两个结论：

1. K 折叠确实按预期把 fixpipe 砍半 （2062 → 824 µs），而且额外把 AIC scalar 也砍半 —— 因为 Matmul 库内部的循环次数同样从 9 次降到 3 次；
2. MTE2 没变 —— 它现在与 AIC scalar 并列成为最大项（1473 vs 1600 µs）。

新瓶颈：四个单元各自的忙碌比例是 cube 0.22 / fixpipe 0.39 / MTE2 0.68 / scalar 0.76 —— 四个都没跑满，且之和远大于 1(≈ 2.2)，说明它们是串行等待而不是并行重叠，AIV 全程在岗却 96% 空闲（在等 AIC 的 GetTensorC）：

aic_time = 2320 µs（cube0）
aiv( vector1 ) aiv_time = 2319 µs，但 aiv_vec_ratio = 0.0157、aiv_scalar_ratio = 0.0378
   ↑ AIV 全程"在岗"却 96% 空闲 —— 它在等 AIC 的 GetTensorC（队列依赖）

即：流水没有重叠 —— AIC 算完一行 → AIV 池化 → AIC 再算下一行，两者互相等待。下一步应做乒乓缓冲（C 结果用两块 UB 交替：AIV 池化第 N 行时，AIC 已在算第 N+1 行）以及 L1 双缓冲 / MTE2 预取。

改完之后的perf：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part2$ for D in float16 float32; do echo "=== $D ==="; ./build-npu/ascendc_kernels --dtype $D --batch 1 --in-channels 256 --out-channels 256 --height 224 --width 224 --filter 3 --pool 2 --warmup 1 --iters 5 2>&1 | grep -E "median"; done
=== float16 ===
[INFO ] min=2318.902 us  median=2328.652 us  avg=2331.366 us  (blockDim=20, 工作项=222)
=== float32 ===
[INFO ] min=4457.652 us  median=4642.744 us  avg=4605.761 us  (blockDim=20, 工作项=444)

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part2$ python3 torch_conv2d_maxpool_baseline.py --kernel
device  : Ascend910B4   (torch.npu.is_available() = True)
baseline: torch F.conv2d → F.max_pool2d（两次独立调用，**不融合**）  B=1 C_in=256 C_out=256 k=3
口径    : warmup=5 iters=20，每次前后 npu.synchronize()，取 median

dtype          H×W pool |  torch median  GFLOP/s | kernel median  GFLOP/s     加速比    端到端 torch       端到端 我们 |   torch 多搬
-------------------------------------------------------------------------------------------------------------------------
float16      32×16    1 |         134.6  3680.35 |         437.0  1133.80   0.31× |         0.7ms        3.26s |       0.4 MB
float32      32×16    1 |         172.9  2865.57 |         170.4  2906.88   1.01× |         0.7ms        3.26s |       0.9 MB
float16    224×224    1 |         418.8 138829.49 |        2610.2 22273.08   0.16× |         2.4ms        3.26s |      50.5 MB
float32    224×224    1 |         983.9 59087.01 |        5561.6 10453.40   0.18× |         4.5ms        3.26s |     100.9 MB
float16    224×224    2 |         472.6 123013.57 |        2339.1 24854.72   0.20× |         2.7ms        3.26s |      50.5 MB
float32    224×224    2 |        1178.8 49317.38 |        4677.2 12429.96   0.25× |         4.2ms        2.25s |     100.9 MB

注：'torch 多搬' = 不融合时多出的中间张量往返（conv 结果写一遍 + 读一遍）= 2 × B·C_out·H_out·W_out·ε；融合 kernel 的'最小访存'不含这一项。

两列端到端的口径（都是"从 host 上的输入图片开始，到池化结果落在 NPU 上为止"）：

• 端到端 torch = x.npu() + F.conv2d + F.max_pool2d + npu.synchronize()，含 H2D 上传；
• 端到端 本实现 = 整个进程的墙钟（--warmup 0 --iters 1）：aclInit + host 侧重排 X′/W′ + 上传 + 一次执行。

两点读法：

1. 小图 fp32 已经追平 torch（1.01×） —— 说明分块/格式/融合的设计是对的；
2. 差距集中在大图（0.16~0.25×），是规模效应：大图里每行的 9 段输入、按 (i,j) 的 GetTensorC 次数、以及 AIC 标量与 MTE2 的绝对量都被放大。
3. 端到端被 host 侧重排主导 ✗ —— 本实现 2.25~3.26 s 几乎全是 RepackImHost/RepackWHost 的耗时（逐通道 memcpy，每次仅 448 B、共约 17 万次），在计时循环之外但属于"从输入图片开始"的真实代价；torch 侧同一口径只有 0.7~4.5 ms。这一项是纯 host 改动、低风险，预期可降到几十毫秒级（见 2.3.9 的优化候选）。

2.3.5 AIV和AIC重复计算

发现了一个冗余的计算：向量侧（池化 / bias / 写回）写在了 ConvRow 返回之后，而 ConvRow 内部才是 while (Iterate)；于是所有进入 OneTile 的核都会执行向量侧（1 AIC + 2 AIV，三份）

__aicore__ inline void OneTile(uint32_t b, uint32_t cout0, uint32_t ph){
...
   ConvRow(...)                  ← 库管这一段（AIC 算、AIV 收）✓ 分工正确
   Max / PairMaxCols / +bias / DataCopy(outGm)   ← 我们自己写的 ✗ 没有区分AIV和AIC的守卫
                                                     → 三个核都照着做了一遍 ✗
...                                                     
}

PairMaxCols / +bias / EnQue / DataCopy(outGm)在三个核上都会执行到。

用 ++dbgWrite_; 测试写次数。同时Process() 里增加测试代码：

// ---- MIX 分工诊断（临时）----
uint32_t dbgTile_ = 0, dbgConv_ = 0, dbgWrite_ = 0;
__aicore__ inline void DumpMixRole() const
{
    if (blockIdx != 0) return;            // 只看 0 号核（避免刷屏）
    if ASCEND_IS_AIC {
        AscendC::printf("[MIX] AIC block=%u blockNum=%u | OneTile=%u ConvRow=%u write=%u\n",
                        blockIdx, blockNum, dbgTile_, dbgConv_, dbgWrite_);
    } else {
        AscendC::printf("[MIX] AIV block=%u blockNum=%u | OneTile=%u ConvRow=%u write=%u\n",
                        blockIdx, blockNum, dbgTile_, dbgConv_, dbgWrite_);
    }
}

// 然后又测试了这个接口
__aicore__ inline void DumpMixRole() const
{
    if (blockIdx != 0) return;
    // 用 sub-block 号分辨 MIX 分组里的 1 个 AIC + 2 个 AIV
    //（ASCEND_IS_AIC 在本环境恒为假，三行会走同一分支，测不出分工）
    AscendC::printf("[MIX] sub=%u block=%u blockNum=%u | OneTile=%u ConvRow=%u write=%u\n",
                    (uint32_t)AscendC::GetSubBlockIdx(), blockIdx, blockNum,
                    dbgTile_, dbgConv_, dbgWrite_);
}

输出都是：

[MIX] sub=0 ... | OneTile=6 ConvRow=6 write=6     ← 同一个 blockIdx 下的 3 个核
[MIX] sub=0 ... | OneTile=6 ConvRow=6 write=6        都在跑完整流程 ✓
[MIX] sub=0 ... | OneTile=6 ConvRow=6 write=6

现在三个核（1 AIC + 2 AIV ）都在跑完整的"矩阵乘 + 池化 + bias + 写回" → 向量侧被算了 3 遍 （值相同所以结果对 ，但白做）。

能不能加守卫省掉这 2/3？ 直觉上可以 —— 昇腾的 MIX 内核（cube + vector 混在同一份代码里 ）有个特点：同一份 kernel 源码会在 AIC 和 AIV 上各执行一遍 。于是很自然想到：要让"cube 相关的代码只在 AIC 跑、向量相关的只在 AIV 跑"，自己加分支就行了 ：

if (g_coreType == AIC) {          // ← 这就是"AIC 守卫"
    ... 只让 cube 干的事（Matmul 的 Iterate 等）...
}
if (g_coreType == AIV) {          // ← "AIV 守卫"
    ... 只让向量单元干的事（池化、bias、写回）...
}

但是在 MIX 内核里，AIC/AIV 的分工由框架经 KFC 负责 —— 用户代码不需要、也不应该自己加核类型守卫 。参考官方案例：

// ① 核类型与比例：写在【核函数声明】上，是编译期的分类，不是运行期分支
//    " __mix__(1, 2) 表示核函数按 Cube:Vector = 1:2 的比例执行。"
__global__ __mix__(1, 2) void matmul_leakyrelu_custom(...)

// ② 固定 baseShape，让 Cube 单次 Iterate 产出一个完整的输出块
//    " 通过 SetFixSplit(128, 64, 128) 固定 baseShape，Cube 侧单次 Iterate 计算
//      baseM * baseN = 128 * 64 的输出块，每个逻辑核需要 4 次 Iterate 完成
//      自身负责的 512 * 64 输出区域。"
tilingApi.SetFixSplit(128, 64, 128);

// ③ ★ 向量处理写在 while (Iterate) 循环体内部，紧跟 GetTensorC
while (matmulObj.template Iterate<true>()) {
    LocalTensor<float> mmOutLocal = mmOutQueue.AllocTensor<float>();   // VECIN（UB）
    matmulObj.template GetTensorC<true>(mmOutLocal, false, true);      // C 直接落 UB
    ...
    AscendC::LeakyRelu(reluOutLocal, mmOutLocal, 0.001, tiling.baseM * tiling.baseN);
    ... 搬回 GM ...
}

三条结论：

1. __mix__(1, 2) 是官方的"编译器分类"机制 —— 核类型与 Cube:Vector 比例写在核函数声明上（编译期），而不是运行期 if ASCEND_IS_AIC 分支。本实现的 Block Dim = 20 / Mix Block Dim = 40 恰好也是 1:2，因此把两处 kernel 入口从 __global__ __aicore__ void 改成 __global__ __mix__(1, 2) void。
2. GetTensorC 落 UB 是官方支持的 —— TQue<TPosition::VECIN, 1> + GetTensorC(localTensor, false, true)，与我们的 CType = VECIN + GetTensorC(cLocal) 完全一致，说明这条路线选对了（C 不必绕 GM）；
3. 向量处理必须写在 while (Iterate) 循环体内 —— 这是分核的正解。官方注释写明"MIX 分离模式下 Matmul API 由 AIV 侧发起。AIV 执行 Iterate 时通知 AIC 完成 Cube 矩阵计算"，即 Iterate() 只在拥有该 C 块的核（AIV）上返回 true，AIC 上立即返回 false，因此循环体天然只在 AIV 上执行，不需要任何核类型守卫。

实测（CANN 8.5 / 910B4）：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part2$ bash run.sh --run-mode npu --build-only -v Ascend910B4 2>&1 | tail -2 && python3 test_harness.py --quick && ./build-npu/ascendc_kernels --dtype float16 --batch 4 --in-channels 128 --out-channels 128 --height 32 --width 16 --filter 3 --pool 1 --warmup 1 --iters 2 2>&1 | grep -E "median"
[100%] Built target ascendc_kernels
[INFO ] 构建完成：/opt/atomgit/asst4-ascend/part2/build-npu/ascendc_kernels
[info] soc_version = Ascend910B4
[info] 参考实现：numpy + torch 双基准（两者必须一致）

=== 正确性用例 ===
[case] bias=0 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）
[case] bias=1 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）

正确性总判定： [ OK ] 全部通过
[INFO ] min=495.176 us  median=508.416 us  avg=501.796 us  (blockDim=20, 工作项=120)

四个小图用例全部 [ OK ]；小图 median = 508.4 µs，与改前（512.2 µs）持平 —— 说明该属性在 8.5 上受支持且语义中性（它只是把编译器对核类型的推断改为显式声明，不改变计算结果，也不改变调度比例）。

2.3.6 Bias 和 maxpooling 与 矩阵乘法流水优化

思考目前的瓶颈：

aic_cube_ratio   0.20   ← cube 只忙 20%
aic_scalar_ratio 0.69   ← 最忙的是 AIC 的标量（库内部寻址/循环）
aic_mte2_ratio   0.63
aic_fixpipe_ratio 0.36
aiv_vec_ratio    0.014  ← AIV 只忙 1.4%
aiv_scalar_ratio 0.034  ← 3.4%

阶段 1 用 Iterate() 的返回值当"谁持有 GetTensorC"的判据，消灭三核重复

向量侧（池化 / bias / 写回）写在了 ConvRow 返回之后，而 ConvRow 内部才是 while (Iterate)；于是所有进入 OneTile 的核都会执行向量侧（1 AIC + 2 AIV，三份），而 ASCEND_IS_AIC / GetSubBlockIdx() 在本环境都不可用作运行期判据。

"MIX 分离模式下 Matmul API 由 AIV 侧发起；AIV 执行 Iterate 时通知 AIC 完成 Cube 计算"

—— Iterate() 只在持有该 C 块的核（AIV）上返回 true，在 AIC 上立即返回 false、循环体不执行。于是把"是否持有 C"这个事实从循环里带出来即可：

// ConvRow 增加出参 gotC：Iterate 成功过就置真
__aicore__ inline void ConvRow(const LocalTensor<T> &dst, uint32_t b, uint32_t cout0,
                               uint32_t oh, bool *gotC = nullptr)
{
    for (uint32_t i = 0; i < tiling.filterHeight; ++i) {
        mm.SetTensorA(wGm[WSliceOff(i, cout0)]);
        mm.SetTensorB(imGm[imOff(i)], false);
        while (mm.template Iterate<true>(i != 0)) {
            mm.template GetTensorC<true>(dst, false, true);   // C 落 UB（VECIN）
            if (gotC) *gotC = true;                            // ★ 只有持有 C 的核会走到这里
        }
    }
}

// OneTile：向量侧整体用 gotC 门控（行 max / 列 max / bias / EnQue / 写回）
bool gotC = false;
ConvRow(cLocal, b, cout0, ph * pool, &gotC);
for (uint32_t p = 1; p < pool; ++p) {
    ConvRow(tmp, b, cout0, ph * pool + p, &gotC);
    if (gotC) AscendC::Max(tmp, tmp, cLocal, tiling.ubTileElems);
}
if (gotC && pool > 1) { ... PairMaxCols ... }
if (gotC) { ... +bias ... EnQue/DeQue ... DataCopy(outGm, ...) ... }

之后又做了一组"为乒乓流水铺路"的小改动，其中一项是把 C 块的 UB 队列深度从 1 调到 2。它在小图上完全看不出问题（508.4 → 503.0 µs），于是被当作无害改动保留了下来。但它造成了大图上一次静默的 UB 溢出：

大图 fp16 的 UB 账（ubTileElems = baseM × padOutWidth = 128 × 224 = 28672 元素 = 57 344 B）

症状极具迷惑性：

修复与防复发：

1. 深度恢复为 1，并把四个队列深度改为单一来源的命名常量，避免"改了一处、核算没跟"：

constexpr uint32_t kOutQDepth = 1, kTmpQDepth = 1, kPoolQDepth = 1, kBiasQDepth = 1;
pipe.InitBuffer(outQueue,  kOutQDepth,  tiling.ubTileElems * sizeof(T));
pipe.InitBuffer(tmpQueue,  kTmpQDepth,  tiling.ubTileElems * sizeof(T));
if (tiling.poolSize > 1) pipe.InitBuffer(poolQueue, kPoolQDepth, tiling.baseM * tiling.poolOutWidth * sizeof(T));
pipe.InitBuffer(biasQueue, kBiasQDepth, tiling.outChannels * sizeof(T));

1. 紧邻着加运行时核算与溢出告警：按真实的 (深度, 元素数, sizeof(T)) 求和，超限直接在设备日志报出：

AscendC::printf("[UB] 自用 = %llu B（outQ %u×%u + tmpQ %u×%u + poolQ %llu + bias %llu）"
                " / 上限 %llu B%s\n", ubUsed, ...,
                (ubUsed > kUbLimit) ? "   ✗✗ UB 溢出！会静默劣化，必须改深度或减小 tile" : "   ✓");

修复后实测：大图 fp16 2 404.7 µs、fp32 4 594.7 µs、小图 fp16 511.4 µs、--quick 4/4 ✓ —— 大图回到 34× 之前的水平。

阶段 2：沿B矩阵的列N切分， 让库把"算"与"取"流水起来

阶段 1 只去掉了重复，没有解决串行：profile 显示 cube 0.20 / fixpipe 0.36 / MTE2 0.63 / AIC scalar 0.69，忙碌比例之和 ≈ 1.9 > 1，说明各单元是互相等待而不是重叠 —— 根因是 baseN 等于整个 N（= OWp），一个 base 块，库内部没有可流水的东西。

官方样例（matmul_leakyrelu_advanced_api）用 SetFixSplit(128, 64, 128) 把每个核的输出切成 4 个 base 块，于是 Iterate() 每块返回一次，库可以"算第 N+1 块的同时，把第 N 块交给向量侧"。

现在（baseN = 224 → 1 个 base 块）：
   [cube 算整行 128×224]  →  [fixpipe 取回 UB]  →  [池化+bias+写回]  →  下一项
   ↑ cube 和 fixpipe/AIV 严格【交替】✗  ← profile 里各 pipe 忙碌和 ≈1.9 > 1 就是它

切 N 之后（baseN = 64 → 4 个 base 块）：
   [cube 块0] → [cube 块1 ‖ fixpipe 块0] → [cube 块2 ‖ fixpipe 块1] → ...
                        ↑ 重叠 ✓（库内部用 KFC 队列自己排）

对应到本实现：

host 一行：  tilingApi.SetFixSplit(t.baseM, /*baseN=*/64, -1);   // N=224 → 4 个 base 块
device：     把"取 C → 池化 → bias → 写回"改成【按列段】执行

实现方式：host SetFixSplit(baseM, 64, -1)；device 在 while (Iterate) 里为每个 base 块分配 segQueue（baseM × 64）承接 GetTensorC，再用 DataCopyParams{blockCount = baseM, blockLen = 64·ε/32, srcStride = 0, dstStride = (padOutWidth−64)·ε/32} 把该段跨步拷进整行缓冲的对应列范围 —— 这样下游的池化 / bias / 写回一行都不用改。

按列段执行要重算的东西：

1. 列方向池化必须在段内完成：baseN 取 64、pool = 2 时段内产生 32 个输出列 —— 必须保证 baseN 是 pool × 16 的整数倍，段边界与池化窗口对齐（64 ✓）；
2. 写回参数按段重算：blockLen = (baseN/pool) × sizeof(T) / 32 个 DataBlock，dstStride 用整行跨度 outPoolHeight × PN − 段宽（单位仍是 32 B）；
3. C 缓冲需要多份：N/baseN = 4 个 baseM × baseN 的 UB 缓冲（或队列深度 ≥ 4），因为库会连续吐出多个块；
4. 仍是"最后一次 i 才有效"：C 是 FH 次累加的部分和，只有 i == FH-1 的那轮 Iterate 拿到的是完整结果，池化/写回依然只能在那时做。

测试结果：

这次 UB 没有溢出（新增的 [UB] 核算打印 176 640 B / 196 352 B），正确性 4/4 通过，但性能明显变差：

结论：沿输出列方向，把 N 切成多个 base 块，在本形态下只增加了 Iterate/GetTensorC 的次数（每行 3 → 12 次）和一次额外的段内搬运，并没有换来 cube 与 fixpipe 的重叠。

原因：本实现的 C 需要落在 UB（供池化）、且有"每次 i 累加、只有最后一次有效"的结构，库在本形态下没有可重排的余地。

如果不切分baseN的话（还是保持128，不用减半为64），考虑过异步 Iterate 这条路，被官方接口的注释判了"不可用"

官方注释（matmul.h:226-231）：

/**
 * @brief Calculate a C matrix of size baseM * baseN
 * @tparam [in] sync: set to synchronous or asynchronous mode
 * @param  [in] enPartialSum: whether to accumulate the result of Iterate into CO1 data
 */
template <bool sync = true> __aicore__ inline bool Iterate(bool enPartialSum = false)

在头文件里找到了三条硬约束，matmul_client.h：

// :737
ASSERT(cacheWorkspaceAddr != 0);   // The cache address must be configured in asynchronous mode.

// :767-771
// Only support the mode that the IterateAll is asynchronous and GM output is continuous.
__aicore__ inline void WaitIterateAll()
{
    ASSERT(!isSyncGetC);           // Must be asynchronous mode
}

结论：高阶 API 的异步模式只服务"IterateAll + 落 GM 且连续"这一种形态。本实现"逐块 Iterate + C 落 UB + 每次 i 累加"的形态结构上不适用 —— 不是调参问题。

官方的 [cube 块0] → [cube 块1 ‖ fixpipe 块0] 用的低阶API做的：

KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIC_ONLY);          // 核类型：官方用这个宏（不是 __mix__，也不是运行期分支）
uint8_t l1Mutex  = AscendC::AllocMutexID();   // MTE2 <-> MTE1
uint8_t l0Mutex  = AscendC::AllocMutexID();   // MTE1 <-> M
uint8_t l0cMutex = AscendC::AllocMutexID();   // M    <-> FIX
...
AscendC::Mutex::Lock<PIPE_M>(l0cMutex);                   // 手工流水
for (iter0 = 0; iter0 < kL1TileNum; ++iter0)              // K 分两级：L1 一块 → L0 若干块
    AscendC::Te::Copy(copyGM2L1, tensorAL1, tensorAGmTile);   // Te API：声明式搬运

即：手工管 L1/L0 分级 + AllocMutexID/Mutex::Lock 互斥量 + Te::Copy/布局对象。

至此，"打破 cube 与 fixpipe 串行"的可用手段已经逐条验证完毕：

因此设备侧的优化在"高阶 API"这一层已经穷尽：能拿到的两笔（优化 1 的 54×、优化 2 的 1.40×）已经拿到，其余方向要么代价是重写（低阶手工流水），要么实测为负（上表四行）。当前保留版本：大图 fp16 2 404.7 µs、fp32 4 594.7 µs、小图 511.4 µs，正确性 4/4。

2.3.7 使用异步模式 IterateAll接口——im2col算法

2.3.8.1 fixpipe 不能直接把 L0C 搬进 UB

最开始以为 CType = VECIN 意味着 " "。实测数据不支持这个说法 —— 从 Memory.csv 的列可以直接看出来（AIC 20 个核 / AIV 40 个子核的均值）：

★ 关键的自洽性检验：每个 AIC 对应 2 个 AIV（Mix Block Dim 40 / Block Dim 20），因此

L0C_to_GM（每 AIC 一份）      = 7123.2 KB
GM_to_UB（每 AIV 一份）× 2    = 3562.1 × 2 = 7124.2 KB   ← 与上式吻合到 0.02%

两路读数按 2:1 严丝合缝，说明 C 的路径确实是：L0C --fixpipe--> GM --MTE2--> UB，两个 AIV 各从 GM 读一份。小图（ia_small）同样：L0C_to_GM = 44.0，GM_to_UB = 22.25 × 2 = 44.5 。

这条事实改变了之前几轮优化的成本判断：

1. "让 C 走 GM"不是新增代价，而是现状。路线 C（我们自己的 GM 双缓冲 + 异步 IterateAll）与当前实现的搬运量完全相同（约 7 MB / 全程，L0C_to_GM_bw_usage_rate 仅 1.54%），区别只在于谁拥有那块 GM 缓冲、以及是否同步等待。
2. 串行的根源因此更清楚了：不是内存通路（通路本来就一样），而是 GetTensorC 的同步语义 —— 库把 C 写进它自己的 GM 缓冲后，阻塞等 AIV 取走，取完才轮到下一次 Iterate。换句话说，库已经在"落 GM"，只是不把这步交给我们做流水。
3. 这解释了 MTE2 为何是最忙的 AIC 管道（0.634，绝对时间 1472.7 µs）：它既要搬 A/B（GM_to_L1 = 67209 KB），又承接了 C 的落盘（L0C_to_GM = 7123 KB）；read_main_memory_datas = 74313 KB ≈ 67209 + 7104 。

2.3.8.2 AIC 与AIV 工作流程

MIX 默认流程的描写是：「用户从 AIV 侧发起消息，通过消息通信框架中转消息后，在 AIC 侧执行 Matmul 计算」。代价：「这套消息处理机制会带来额外的 Scalar 性能开销」。

AIV: Iterate(N) ──→ AIC 算 N ──→ GetTensorC(N)【阻塞等取回】──→ 池化(N) ──→ Iterate(N+1)
                                                              ↑ 这段时间 AIC 在【空闲】✗

官方注释说"Matmul API 由 AIV 侧发起" —— 所以 AIC 是 AIV 的"下游" ：AIV 不开口，AIC 就不动 。而我们的代码里 Iterate（发起 ）和 GetTensorC（取回 ）紧挨着 → 于是"池化期间 AIC 空转" 。

每个核调用一次 IterateAll，会计算出 singleCoreM singleCoreN 大小的 C 矩阵”，它只算当前核分到的那一块。切分是 *host 侧 tiling 决定的（MultiCoreMatmulTiling + SetDim + GetTiling）：A 沿 M 轴切、B 沿 N 轴切，K 轴要显式 EnableMultiCoreSplitK 才切。

每调用一次 Iterate，会计算出一块 baseM baseN 的 C 矩阵，一次singleCoreM singleCoreN 大小的 矩阵乘切开后有 N 个 base 块，同步模式就是 N 条消息 + N 次等待。

参考 IterateAll 手册, 不切 K 时：

• A 矩阵沿 M 轴切分 → singleCoreM，单核处理 SingleCoreM × K
• B 矩阵沿 N 轴切分 → singleCoreN，单核处理 K × SingleCoreN
• C = SingleCoreM × SingleCoreN

2.3.8.3 一次错误的尝试

最开始尝试了在 X′ 口径上，只把列位移折进了 K （N = OWp = 224、K = FW·C_in = 768 ，行位移还在外层 i 循环里 ✓）"把切分权交给库" ✓：

SetDim(16);                              // 告诉库：可用 16 个 AIV
SetOrgShape(C_out, OWp, K);              // 原始形状
SetCType(TT::GM, CF::ND, dtype);         // C 落 GM
// ★ 去掉 SetFixSplit —— 让库自由切

kernel 侧改用 IterateAll、每核写回自己那片 M 。

库的回答：

[TILING-B] C_out=256 OWp=224 K=768 | usedCoreNum=14
           singleCoreM=256   singleCoreN=16   singleCoreK=768

→ 切的是 N（不是 M） ，而且切到分形下限 16 → 每核拿到 256 × 16 这个"巨大 M、极小 N"的形状

按 N 切 → 每个核拿到的 M 是【完整的 C_out = 256】，只有 N 是 16 列
       → 每核都要算全部 256 个输出通道
       → 就必须把【整个 A (256 × K)】读进自己的 L1
       → 14 个核 → A 被读 14 遍 ✗

→ 结果错 50 倍、耗时 82 ms 。

据此确定的路线 （IterateAll + 自管 GM 缓冲）。

2.3.8.4 IterateAll 实现 iml2col 算法

问题起点：用户指定的输入是 X (B, C_in, H, W)、权重 W (C_out, C_in, K_h, K_w)。把卷积化为矩阵乘时，$$\mathrm{out}[co, o_h, o_w] = \sum{c,i,j} W[c_o,\, c,\, i,\, j] \cdot X[c,\, o_h+i,\, o_w+j]$$ 。

shift reduction 实现只需 3 份：行位移是免费的

列位移 j：段内 [j, j+OWp) → 源地址偏移 j 个元素（fp16 下 2~4 B）✗ 不满足 DataCopy 的 32 B 对齐 ✗
         → 必须在 host 上物化 ✓（这正是"优化 1/2"里那 FW 份副本的由来）
行位移 i：取段 (r = oh+i) → 只是换一个基址 ✓ 零成本 ✓
         → 所以 X′ 里【根本没有 i 这一维】✓，它由"读第 oh+i 行"隐式实现

为什么 IterateAll 会把这份"免费"收回去

IterateAll 接口要求 B 必须是一整块连续的 (K, N)，而收缩维 K 轴按语义是 (i, j, c)。于是 FH·FW 组数据必须在内存里连续挨着。唯一办法就是使用im2col算法那样，把它们都复制到连续位置，这就是那 9 份（FH·FW）。

因此三条路互斥：

2.3.8.5 落地实测：im2col + IterateAll

设计要点（从输入输出看）

• 切分不交给库：SetDim(1) + SetSingleShape(M=C_out, N=N_s, K=FH·FW·C_in)，库只负责执行（实测 usedCoreNum=1）。

• 工作项 = (batch, 行条带)：把整图的 N（= H_out·W_out）沿空间维切成 nStrips 刀，每刀大小 N_s = rowsPerStrip × padOutWidth（大图 6×224 = 1344，共 37 刀）；

  总共工作数量items = batch × nStrips（batch = 图片张数，大图 batch=1 → items=37，小图 batch=4 → items=20）。

• 并行只切 AIV：items 按 AIV 序号轮转（步长 blockNum × GetSubBlockNum() = 40）。大图 37 项 ⇒ 37/40 个 AIV 各 1 项（3 个空转）；小图 20 项 ⇒ 20/40 各 1 项。

• AIC 不参与切割：它由框架经 KFC 执行其配对的 2 个 AIV（1 AIC : 2 AIV）发来的 IterateAll 请求。

• 每核私有 scratch ⇒ 核间零同步：matmul的输出 C 的 scratch 按 blockIdx 私有，池化只碰自己那几行，不需要 SyncAll / 原子操作。

Tiling 画个表格：

由于 singleCoreK = 整个 K，所以「每一个工作项都要用完整的 A」。所以 37 个工作项、每个都要把 A 从头到尾走一遍 （A 净大小只有 1.18 MB，实际搬运 6.75 MB）。

1. 一个工作项的 A 范围singleCoreM × singleCoreK = 256 × 2304 —— 就是整块 W_mat（每个条带都得用它）
2. 库在 (M,K) 平面上自己找块：
   • 沿 M 切 ceil(256/baseM=128) = 2 块
   • 沿 K 切 ceil(2304/baseK=64) = 36 刀。
   • 每块 baseM×baseK = 128×64 经 GM→L1→MTE1→L0A
3. B 同理，沿 (K,N) 切每块 baseK×baseN 进 L0B，记作n。A 的每一块要和某一个 n 块配对，而 baseN < N_s ⇒ 库要绕 n 循环 ⇒ 每绕一圈，A 就要重头再来一次
4. Mmad 累加进 L0C，K 方向 36 刀全部累加完，L0C 里才是最终结果；然后 fixpipe 写出去。

四条约束（都属于"不满足就静默错"）

maxpool 的上下行边界：用"整除"消掉，不需要任何跨核通信

pool=2、rowsPerStrip = 6 ⇒ 每个条带的 6 个卷积行恰好配成 3 个池化行：

条带 0：卷积行 0..5    → 池化行 0,1,2     配对 (0,1) (2,3) (4,5)
条带 1：卷积行 6..11   → 池化行 3,4,5     配对 (6,7) (8,9) (10,11)
                                  ↑ 所有配对都落在条带**内部**，不跨 5|6 这条缝

• rowsPerStrip 被 poolSize 整除 ⇒ 池化窗口永不跨条带 ⇒ 每个核只需要自己那块 scratch ⇒ 这正是"核间零同步"的物理来源。
• 行数守恒：37 刀 × 3 行 = 111 = H_out / pool = 222 / 2 。
• 反面：若 rowsPerStrip 取奇数（如 3），窗口就会跨缝，相邻条带之间产生真实的数据依赖（要读邻居的行、或加同步）。
• 列方向同理：池化在同一行内成对（ow, ow+1），W_out = 222 为偶数 ⇒ 恰好 111 对；padOutWidth = 224 多出来的 2 列是纯 0，配成一对后在写回时由 host 裁掉。

性能测试

python3 test_harness.py --perf

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part2$ python3 test_harness.py --perf
[info] soc_version = Ascend910B4
[info] 参考实现：numpy + torch 双基准（两者必须一致）
[build] /opt/atomgit/asst4-ascend/part2/build-npu/ascendc_kernels 比 conv2d.asc 旧 → 自动重编（源码改动必须重新编译才生效）
[build] bash /opt/atomgit/asst4-ascend/part2/run.sh --run-mode npu --build-only --soc-version Ascend910B4

=== 正确性用例 ===
[case] bias=0 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）
[case] bias=0 maxpool=0 大图 224×224
  [ OK ] float32: 50466816 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 50466816 个元素全部匹配（rtol=atol=0.01）
[case] bias=1 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）
[case] bias=1 maxpool=0 大图 224×224
  [ OK ] float32: 50466816 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 50466816 个元素全部匹配（rtol=atol=0.01）

=== 性能用例 ===
注意：Trainium2 版给出的参考阈值（4596 / 1002 μs 等）是特定硬件的实测值，
      在昇腾上无意义。下面是本机的实测值，MFU 需要你自己按机型峰值算。

  float16  32× 16 pool=1  median=    111.0 us  4465.50 GFLOP/s   14.93 GB/s  baseM/N/K=128×16×128  blockDim=20
  float32  32× 16 pool=1  median=    167.6 us  2957.01 GFLOP/s   19.78 GB/s  baseM/N/K=64×16×64  blockDim=20
  float16  224×224 pool=1  median=   1076.9 us  53984.58 GFLOP/s   48.38 GB/s  baseM/N/K=128×224×128  blockDim=20
  float32  224×224 pool=1  median=   2166.2 us  26838.76 GFLOP/s   48.11 GB/s  baseM/N/K=64×224×64  blockDim=20
  float16  224×224 pool=2  median=   1000.9 us  58085.96 GFLOP/s   33.15 GB/s  baseM/N/K=128×224×128  blockDim=20
  float32  224×224 pool=2  median=   2123.1 us  27383.73 GFLOP/s   31.25 GB/s  baseM/N/K=64×224×64  blockDim=20

正确性总判定： [ OK ] 全部通过

1. 性能：大图 fp16 从 X′ 版的 2 424 µs 降到 1 076.9 µs（−55.6%）；同口径的含 maxpool 一组从 X′ 版的 2 339.1 µs 降到 1 000.9 µs（−57.2%）；fp32 为 2 166.2 / 2 123.1 µs。含 maxpool 与不含 maxpool 的两组基本持平 —— 池化不再是叠加在卷积之后的一段独立开销。
2. 异步双缓冲已试（§2.3.8.4，--async-iterate 2）：机制成立、功能正确（大图 50 466 816 元素全匹配），但大图零收益（±0.6%）—— 因为 items=37 / aivNum=40 ⇒ 每核恰好 1 个工作项，流水只有填充与排空、没有稳态。所以 1 000.9 µs 在当前工作项粒度下就是这条路的地板；要突破得先改粒度（rows-per-strip）或转向 MTE2 那 2.52× 重读。

一个"静默失败"：

UB 溢出是静默的。InitBuffer 分配失败只返回 false、不报错。本轮为补 MTE2→V 同步新增了一块 inQueue（UB 里的 tile 从 2 块变成 3 块），再叠加 pool>1 才分配的 poolQueue，大图（padOutWidth = 224）的 UB 预算变成：

fp16: 3×57 344 + 28 672（poolQueue）+ 512（bias） = 201 216 B  >  196 608 B（910B/910C 上限）
fp32: 3×57 344 + 28 672 + 1 024                    = 201 728 B  >  196 608 B

而 pool=1（不分配 poolQueue）只要 172 544 B —— 于是症状恰好是"只有大图 + maxpool 挂"，且错的地方实际值恒为 0（那片输出根本没被写），而不是 NaN 或垃圾值。

修法是不申请第三块缓冲：池化列归约的"偶列 / 奇列各取一半"可以落在目标缓冲自己的后半段（dst 的容量 = baseM × padOutWidth = 2 × baseM × poolOutWidth），省掉整块 28 672 B。

另外 host 侧的 g_ubNeed 原先按"2 块"估算而内核实际 3 块（现在两者对齐，超上限时 [TILING] UB 预算 打印真值并给 [WARN]）。

2.3.8.6 引擎画像：msprof 逐引擎实测

采集命令（两个算子各采一份。）

cd ~/asst4-ascend/part2
KM='_Z25fused_conv2d_maxpool_fp16PhS_S_S_S_S_N7AscendC6tiling11TCubeTilingE12Conv2dTiling'
KP='_ZN12_GLOBAL__N_113pretrans_fp16EPhS0_S0_S0_NS_15TransDataTilingE'
python3 prof.py --profile_name pretrans_w        --kernel-name="$KP"
python3 prof.py --profile_name im2col_fp16_p2    --kernel-name="$KM"
python3 prof.py --profile_name im2col_fp16_p1    --pool 1 --kernel-name="$KM"
python3 prof.py --profile_name im2col_fp16_p2_l0 \
    --aic-metrics=MemoryL0,MemoryUB,ResourceConflictRatio,L2Cache --kernel-name="$KM"

直接用Proile 脚本:

cd ~/asst4-ascend/part2

# ① 主口径：大图 fp16 pool=2 —— prof.py 的默认参数就是这套（batch=1, in/out=256, 224×224, k=3, warmup=1, iters=5）
python3 prof.py --profile_name im2col_fp16_p2

# ② 同图 pool=1（--perf 里 1076.9 µs 那条）
python3 prof.py --profile_name im2col_fp16_p1 --pool 1

# ③ 片上细节：L0/UB 带宽 + bank 冲突 + L2 命中（和 ① 同参，另存一份）
python3 prof.py --profile_name im2col_fp16_p2_l0 \
    --aic-metrics=MemoryL0,MemoryUB,ResourceConflictRatio,L2Cache

# ④ 打包下载
tar -C profile -czf im2col_prof.tgz im2col_fp16_p2 im2col_fp16_p1 im2col_fp16_p2_l0 pretrans_w
ls -lh /tmp/im2col_prof.tgz

一次完整前向 = 1.88 ms，前置算子占 17%

算子                耗时(ms)      占比   主要忙碌单元
pretrans_w         0.311     17%   scalar_ratio=0.993, mte2_ratio=0.006, mte3_ratio=0.003  (blockDim=40)
im2col_fp16_p2     1.567     83%   mte2_ratio=0.664, scalar_ratio=0.281, cube_ratio=0.224  (blockDim=20)

pretrans_* 是纯标量算子（aiv_scalar_ratio = 0.993、aiv_vec_ratio = 0.0005、AIC 侧全是 NA）：它把 1.18 MB 权重逐元素排成 W_mat，耗时 311 µs。

发现 A：kernel 把工作项重复执行了一遍半，而且负载严重倾斜

PipeUtilization.csv 的 aic_time 按块排开是这样的五级阶梯：

块 0–7   : 1560 µs ×8     块 8 : 1206 µs     块 9–17 : 890 µs ×9
块 18    : 496 µs          块 19 : 17 µs（空转）

只有一种模型能推出这个形状：循环步长与实际核号范围不匹配： GetBlockNum() 与 GetBlockIdx() 不在同一个计数空间！

// ✗ 修复前：AIV 侧序号空间到 39，步长却是 20
for (i = GetBlockIdx(); i < items; i += GetBlockNum()) { ... }

// ✓ 修复后：本类核的实际核数 = 配置块数 × 一个 AI Core 上本类核的数量
const uint32_t aivNum = blockNum * static_cast<uint32_t>(AscendC::GetSubBlockNum());  // 20 × 2 = 40
for (i = GetBlockIdx(); i < items; i += aivNum) { ... }

AIV 侧 GetBlockIdx() 给出 [0, 39]（MIX_AIC_1_2 下 AIV 是 AIC 的 2 倍），但 GetBlockNum() 给出的是 20（声明的 AIC blockDim）。于是 AIV k 与 AIV k+20 各算了一遍 item k（k ≤ 16），合计执行 54 次工作项而不是 37 次。

逐块推演与实测逐一对上：

单次执行的耗时约 390 µs（1560 ÷ 4），所以把 stride 修正后主算子的下界约 800 µs（每核 2 份而不是 4 份）。

为什么正确性一直没暴露：重复执行的两份输入相同、结果相同，写的是同一片输出；C 的 GM scratch 又是按 blockIdx 私有的，所以既不错也不冲突。

修复前后（python3 test_harness.py --perf，median）：

小图不变恰好是一次自洽性检验：小图 items = batch(4) × 条带(5) = 20 = blockNum，旧步长根本没有回绕；只有大图 items=37 > blockNum=20 才会让 AIV k 与 AIV k+20 撞车。修复只在大图生效 —— 与"错位只在 items > blockNum 时发生"的机制完全吻合，不是碰巧。

逐引擎对比（同用例：大图 fp16 pool=2）

为什么用绝对时间而不是"平均忙碌比例"：负载不均匀时，"对 20 个块求平均的 ratio"会被轻载块拉低，跨版本不可比。绝对时间（Σ aic_*_time ÷ 20）才是能跨版本对齐的口径；上图右栏的堆叠柱是平均 ratio，只用于看"谁高谁低"的相对形态。

三个版本在这张表里读作一条线：X′ 版 → fixpipe/scalar 两条大头被 IterateAll 砍掉（825→170、1601→441）但负载失衡（利用率 70.6%）→ 修 stride 后时间再降 42.7%、利用率回到 91.7%。而 MTE2 的占比一路涨到 0.852 —— 搬运已经成为唯一的天花板。

工况矩阵（均为修 stride 之后的数）

• 单位：cube/fixpipe/MTE1/MTE2/AIC scalar 是单块平均绝对时间（µs）（Σ aic_*_time ÷ 20）；AIV 忙碌和、核时利用率 无量纲；GM→L1 为 20 块合计。
• 小图的 AIV 忙碌和 0.73~0.93 看起来很高，其实是分母太小（aiv_time 只有 83~186 µs，绝大部分是固定开销），不能与大图的 0.13 直接比。

大图可读出的两件事：

1. MTE2 在六行里全是最大值（大图占 op 时间 0.85，小图占 0.07~0.09 的绝对量级但比值也最高）；
2. 核时利用率 大图 91.7~92.6%（已接近"37 项 / 20 个 AIC 块 = 2:1 分块"能给的 92.5% 上限），而小图只有 19~24% —— 小图 items=5（fp16）/6，20 个 AIC 块里绝大多数领不到活，那是工作项太粗造成的，与 stride 无关。

这一节的三条结论：

• 成果：IterateAll 确实把 §2.3.4 记录的那场 GetTensorC 风暴扑灭了 —— fixpipe 825→170 µs、AIC scalar 1601→441 µs，两者在图上是肉眼可见的两条带子塌下去。
• 修 stride 的收益：核时利用率 从 70.6% 回到 91.7%（见发现 A），主算子 1567 → 898.6 µs。值得注意的是这一刀只削掉了"重复"，没有改变"单次成本"：GM→L1 从 1002.7 MB 降到 687.1 MB，恰好是 ×37/54 = −31.5%，单次工作项仍是 18.6 MB（理想 7.37 MB 的 2.52×） —— 两笔账（重复执行 / L1 分块重读）互不污染，这正好说明发现 B 那个 2.5× 是独立存在的，没有被这次修复顺带解决掉。
• 没做到：AIV 忙碌和 几乎没动（0.113 → 0.120 → 0.129，仍 87% 空闲）。也就是说 "cube 在 AIV 池化期间空转"这个最初的动机，在同步 waitIterateAll = true 下并没有被消除 —— 池化仍然串在中间，只是被"一次 IterateAll 算完整条带"盖住了大半。异步双缓冲（§2.3.10）不能省。

发现 B：扣掉重复之后，MTE2 的真实重读仍有 2.5×

理论下限 = 37 项 × (A = W_mat 256×2304×2B = 1.18 MB + B = 条带 2304×1344×2B = 6.19 MB) = 272.7 MB
实测 GM→L1 总计                                                                    = 1002.7 MB  → 3.7×
扣除重复执行（54/37 = 1.46×）后                                                                → 2.5×
单次执行 = 1002.7 MB ÷ 54 次 = 18.6 MB，理想 7.37 MB                                           → 2.5×

即重复执行与 L1 分块的重读是两笔独立的账：前者是 bug，后者是 baseM=128 / baseN=256 / baseK=64 这套分块下的固有代价（A 按 baseK 切片的行式访问每行只有 128 B）。

2.3.8.7 异步双缓冲：机制成立，但当前粒度下零收益

动机：实测 AIV 忙碌和只有 0.129（约 87% 空闲）—— 同步 IterateAll 下，AIV 必须等本次矩阵乘回来才能开始池化，而池化期间 AIC 无事可做。双缓冲的思路：给每个核两块 C scratch，交替使用，于是"项 k+1 的矩阵乘（AIC）"与"项 k 的池化（AIV）"可以重叠。

实现（开关 --async-iterate 2，默认路径不动）

同步路径下 AIV 的 IterateAll 是"发一次、等它算完"，池化只能排在它后面 ⇒ AIC 在 AIV 池化期间空转。双缓冲就是把"等"从发起点挪走，让第 k 项的矩阵乘与第 k−1 项的池化同时进行：

1. 矩阵乘的输出矩阵C 的 scratch 每核 2 块 ：

   cScratch  +  blockIdx × 2 × slotElems  +  slot × slotElems
      ↑              ↑                            ↑
    GM 整块首地址    跳过前面核（每核 2 块）      本核内部选第 0 块还是第 1 块
   （host 一次性 malloc）
   

   • slotElems = "一块 slot 放多少元素" = C_out × N_s （大图 256 × 1344 = 344,064 元素 = 688 KB fp16 ）；
   • blockIdx = 本核编号 （AIV 侧 0~39 ）→ 于是每个核独占 2 × slotElems 一块区域
   • slot = 序号 & 1 交替（序号按 2 取余）。
   • 于是正在被写的那块与正在被读的那块永远不同，不需要任何额外同步。

2. 调用 IterateAll不等待 —— IterateAll<false>(dst, 0, false, /*waitIterateAll=*/true)，发完立即返回。

   第 4 个参数 waitIterateAll必须给 true；给 false 却调 WaitIterateAll() 等于等一个从未被武装的信号 —— 死锁。

   waitIterateAll：仅在异步场景下使用:

   • true：需要通过 WaitIterateAll 接口等待 IterateAll 执行结束。

   • false：不需要通过 WaitIterateAll 接口等待，开发者自行处理等待的过程。

3. 池化段拆出来单独可调 —— 把原来的 OneStrip（发起 `IterateAll+ 池化）拆成 StripMatmul（发起到指定 slot）与 OneStripFinish（从指定 slot 回读、做行/列 max + bias、写回），后者才能被安插在"发起之后"。

4. 时序：关键就是这三行的先后

for (本核领到的每个工作项 k) {
    mm.WaitIterateAll();                    // ① 等第 k−1 项的 cube 算完
    StripMatmul(k, slot = k & 1);           // ② 发起第 k 项，不等 → AIC 立刻开始算 cube k
    OneStripFinish(k-1, slot = (k-1) & 1);  // ③ AIV 同时池化第 k−1 项（与 ② 并行）
}
// 循环末尾补一次 WaitIterateAll() + OneStripFinish(最后一项)  // 收尾

② 的执行者是 AIC（cube），③ 的执行者是 AIV（vector：GM→UB 回读 + 行 max + 列 max + bias + 写回）—— 两者落在不同引擎上，所以 ② 一旦发出、③ 就能同时跑 。这正是"cube k 在算的同时，k−1 的输出正在被池化"。

这套流水里，WaitIterateAll() 是唯一的串行点 ：

正确性（两种异步形态都过）：

python3 test_harness.py --quick --test_maxpool --async-iterate 1   # 4/4 ✓
python3 test_harness.py --quick --test_maxpool --async-iterate 2   # 4/4 ✓
python3 test_harness.py --perf --async-iterate 2                   # 含大图 50 466 816 元素全匹配 ✓

性能 A/B（test_harness.py --perf，median，µs）：

大图四行全部落在噪声里（±0.6%），小图反而快 3~4%。

前提：这套流水要有收益，必须每核有 ≥2 个工作项 —— 收益 ≈ (每核项数 − 1) × 单次池化时间（首项的 cube 无物可叠、末项的池化无物可叠）。而本项目两种用例目前每核都只有 1 项（大图 items=37/40、小图 20/40），所以实测是零收益（见下表）—— 机制成立，但没有可重叠的项。

这一节的结论：异步双缓冲机制成立、功能正确，但在"一行条带一个工作项"的粒度下没有可重叠的东西。要让它生效，前提是先让每核有多个工作项（--rows-per-strip 调小 → 大图 111 条 / 每核 ~2.8 项），代价是 N_s 从 1344 降到 448、每项固定开销上升，两者对冲。留作开关备用；真正该动的仍是 MTE2 那 2.52× 重读（§2.3.8.3 发现 B）—— 它才是六个用例共同的、且已到 0.852 的天花板。

截止目前为止的 profile:

表 1 · fp32 · 大图 224×224 vs 小图 32×16（均为 pool=2 全融合）

表 2 · fp16 · 大图 224×224 vs 小图 32×16（均为 pool=2 全融合）

两个尺寸不同病。

大图 85~88% 的时间压在 MTE2 上，是搬运受限；

小图最忙的换成了 AIC 标量（是 MTE2 的 2.8~3.2 倍），是固定开销受限 —— 而它的 核时利用率 只有 29~30%，意味着七成时间是「开了工但没活干」。

同一份代码、同一个算子：大图要动的是 tile 形状（§2.3.8.3 发现 B 那 2.52× 重读），小图要动的是工作项粒度（5 项 ÷ 40 个 AIV）。

2.3.8 W 权重常驻 L1 Cache

动机：§2.3.8.3 发现 B 已经指明 —— 大图 85~88% 的时间压在 MTE2（GM→L1）上，而这笔搬运里 A（W_mat）占 36%，并且是被重复读 6 次读出来的（重读倍数 = N_s / baseN）。A 在整个 kernel 里是同一份数据，那能不能让它常驻 L1、根本不再走 GM 通路？

CANN 有两条路，机制完全不同：

UB 是 AIV 自己的存储，L1 属于 AI Core 里 Cube 侧。在 MIX 场景下 Matmul 是由 AIV 侧发起的 —— 「AIV 发送消息到 AIC 启动 Matmul 计算」，MIX 下 AIC 与它的 2 个 AIV 同属一个 AI Core； IBShare 是两个 AIV 把同一份数据放进 L1 的同一块位置、只搬一次。

官方参数说明：scmBlockGroup 为 false 时「VECTOR0 与 VECTOR1 映射到 CUBE 上的内存位置错开」，为 true 时「相同」。

TSCM 不是把搬运取消了，而是把搬运从「Matmul 库内部那条自动搬入路径」换成「自己写的一句 DataCopy」：

参考资料：TSCM 输入的矩阵乘 · Matmul 高阶 API 使能 IBShare 模板共享 B 矩阵数据 · TPosition · TSCM 简介

门槛核算（A = W_mat，M = C_out、K = FH·FW·C_in，fp16）：

结论：大图上 IBShare 与 TSCM 都开不了。 这和两者的参数细节无关，是「L1 装不下整块 A」的一票否决 —— 两条路撞的是同一堵墙。TSCM 还多一层：官方明说「用户管理整块 TSCM 内存空间」，所以它不是凭空多出 1152 KB，B 的 L1 缓冲要从同一块 512 KB 里划。

为什么不能"降 K 换常驻"：把 K 从 2304 拆小确实能让 A 装进 L1（K = 256 时 256×256×2 = 128 KB），但那等于回到「9 次小 K 矩阵乘累加」，C 要写 9 次（fixpipe ≈ 229 MB）—— 而 fixpipe 与 AIC scalar 正是 IterateAll 这条路线当初花大力气砍下去的两个大头（§2.3.4：fixpipe 825 → 170 µs、AIC scalar 1601 → 441 µs）。这是拿一个已解决的问题去换另一个，不划算。

所以这是AIC的L0A和L0B大小的限制，导致了需要频繁 GM→L1，在L0C上去求局部和。

2.5 总结

2.5.1 工程结构

重新整理了代码，Assigenment README中的算法和课件里im2col算法都实现：

part2/
├── conv2d_common.h        三边共享：Args / Conv2dTiling / TransDataTiling / LCG 造数 /
│                          ReadCoreMem / WriteFile / MakeCubeTiling{SR,Im2col} / RepackIm{SR,2col}Host
├── op_shiftreduce.asc     算子①：KernelFusedConv2dMaxpoolSR<T> + TransDataSR<T> + 4 个 __global__
├── op_im2col.asc          算子②：KernelFusedConv2dMaxpoolIm2col<T> + TransDataIm2col<T> + 4 个 __global__
└── main_conv2d.asc        唯一 host：抽象基类 + 两个子类 + RunFusedConv2dMaxpool + main

两个开关都在构建期生效（不是运行期参数）：

--impl/--transdata 都会触发重编（test_harness.py 因此改为每次运行都构建，不再用"二进制比源码旧"判断 —— 换开关时源码时间戳可能不变，靠时间戳会静默跑上一个算子的二进制）。

2.5.2 命令小结

cd ~/asst4-ascend/part2

# ① 构建（只编一个算子；三种组合）
bash run.sh --run-mode npu --build-only -v Ascend910B4 --impl im2col
bash run.sh --run-mode npu --build-only -v Ascend910B4 --impl sr
bash run.sh --run-mode npu --build-only -v Ascend910B4 --impl sr --transdata kernel

# ② 正确性：小图 4 用例（bias × maxpool）× fp16/fp32
python3 test_harness.py --quick --test_maxpool
python3 test_harness.py --quick --test_maxpool --impl sr
python3 test_harness.py --quick --test_maxpool --impl sr --transdata kernel

# ③ 大图正确性（50 466 816 元素）+ 性能（六用例）
python3 test_harness.py --perf --impl im2col
python3 test_harness.py --perf --impl sr
python3 test_harness.py --perf --impl sr --transdata kernel

# ③b 与 torch 基线的**端到端**对照（torch 侧：F.conv2d → F.max_pool2d 两次独立调用，不融合）
#     脚本会**先按同一套编译期开关重编**并打印 `[info] 本实现 = impl=…，transData=…`，
#     避免把"上一轮构建出来的另一个算子"当成本轮结果（那是编译期开关最容易踩的坑）
python3 torch_conv2d_maxpool_baseline.py --kernel --impl im2col
python3 torch_conv2d_maxpool_baseline.py --kernel --impl sr --transdata host
python3 torch_conv2d_maxpool_baseline.py --kernel --impl sr --transdata kernel

# profile 采集
# ⚠️ 用当前 build 的 mangled 名（从 dump/aicore_binary.o 里 nm 出来的；
#    不是旧版那份 _Z25fused_conv2d_maxpool_fp16… —— 代码拆成 op_im2col.asc 后改名了）
KM16='_Z32fused_conv2d_maxpool_im2col_fp16PhS_S_S_S_S_N7AscendC6tiling11TCubeTilingE12Conv2dTiling'
KM32='_Z32fused_conv2d_maxpool_im2col_fp32PhS_S_S_S_S_N7AscendC6tiling11TCubeTilingE12Conv2dTiling'
python3 prof.py --profile_name im2col_fp16_s2 --dtype float16 \
    --batch 1 --in-channels 128 --out-channels 128 --height 32 --width 16 --pool 2 \
    --kernel-name="$KM16"

python3 prof.py --profile_name im2col_fp32_s2 --dtype float32 \
    --batch 1 --in-channels 128 --out-channels 128 --height 32 --width 16 --pool 2 \
    --kernel-name="$KM32"

test_harness.py 用它自己的 run.sh 调 cmake（--impl → -DOP_IMPL、--transdata → -DTRANSDATA），并把 [info] 算子实现（编译期）：im2col|sr 打在输出第一行 —— 那行就是"这一轮编的是哪个算子"的直接证据。

为什么默认是 im2col：

im2col 用 9× 的输入显存换来 IterateAll 这条通路，实测大图快 2.4× 。

2.5.3 低阶 Load3D 路线：结论与搁置理由

低阶路线是本作业最早尝试、也最先被放弃的一条。它的思路本身是对的，但被 A2 上关键原语的可用性与参数语义卡住。当时的结论完整记录在 conv2d_load3d.asc.bak（975 行，文件头即结论清单），摘要如下。

为什么值得试：5HD 是硬件原生的卷积特征图布局 —— NC1HWC0 (N, C1, H, W, C0)，C1 = C/16、C0 = 16 正好 32 B，于是列位移 j 的地址偏移就是 j × 32 B，天然对齐，任何位移副本都不需要物化（0× 内存）；而 Load3D 正是硬件为"带窗口取数"设计的原语（窗口大小、步长、补零、偏移全是参数，元素级寻址，没有 32 B 限制）。

A2 上实测确认的原语可用性（依据 asc-devkit impl/basic_api/dav_c220/）：

✓ LoadData + LoadData3DParamsV2   (L1 → L0A/L0B)
✗ LoadData + LoadData3DParamsV1   (未实现：ReportNotSupport)

即 V1 在 A2 上根本没有实现，只能用 V2。

卡住的地方是参数语义，而不是算法：

⚠️ 待上机确认的两处（最终未能确认）：
  1. LoadData3DParamsV2 的 mExtension / kExtension 与目的矩阵 (k × n) 的对应；
  2. LoadData2DParams 的 repeatTimes / srcStride 语义。

这两处没有可对照的 A2 实例或文档，参数只能靠试，主 kernel 始终未跑出正确结果。

但这条路上确立下来的东西可以复用：

搁置的三个理由：

1. 关键原语的参数无参考（V1 未实现；V2 的 mExtension/kExtension 与 LoadData2DParams 的语义在 A2 上查不到对照），只能反复试；
2. 一切都要手工管：L1/L0 分块与双缓冲、SetFlag/WaitFlag 同步、(C1, C0) 分形排布 —— 成本等于自己实现一个 Matmul 库；
3. 收益不确定：即便跑通，省下的是 X′ 那 3× 物化（77 MB），而高阶 API 路线已经拿到 8× 的总提升。

补充一条容易被忽略的事实：即便有 5HD，也不能直接喂给高阶 Matmul —— 5HD 的块内是"通道"，而高阶 Matmul 对 B 要求块内是"位置"，需要块级转置，框架读不了（实测表现为确定性全 0）。这正是本实现转向自定义 ND 块 X′ 的直接原因，详见 §2.2.2。

结论：低阶 Load3D 路线的价值在于它揭示了"5HD 天然对齐、零物化"这一硬件事实，但它要求的环境与参数资料在 A2 上不具备，故搁置；本实现最终采用"高阶 Matmul + host 侧物化 X′（3×）"的路线。

2.6 排障总结：易错点与注意事项

2.6.1 排障实录：双缓冲流水的 host 缓冲复用竞态

现象：行切片 + 双缓冲流水落地后第一次跑，python3 test_harness.py --build --quick：

[case] bias=0 maxpool=0 小图
  [FAIL] float32: 最大绝对误差 32.7，首个不符位置 67235（实际 276.136，期望 308.833）
[case] bias=1 maxpool=0 小图
  [FAIL] float32: 最大绝对误差 322.5，首个不符位置 68223（实际 0.197121，期望 322.657）
正确性总判定： [FAIL] 存在不通过的用例

第 1 步：先把「布局 / 重排」排除掉（没有卡也能做）

• kernel 的输入是确定性 LCG 造的（seed 12345 / 67890 / 24680）。按同一套种子在 numpy 里重建落盘张量，直接复现出报错里的两个「期望」：308.83258（@67235）与 322.65738（@68223）⇒ 两次运行喂给 kernel 的 x/w 逐字节相同，两者之间只差 bias 那 128 个数。
• 把 host 的 RepackIm / RepackW 与 kernel 的 imOff / dstOff 地址公式原样搬到 numpy 重放整条链路：与参考实现 max|Δ| = 1.2e-4（fp32 尾数级）⇒ 「行位移 i 折进行索引 + tile 行切片」这次重排在数学上等价，布局可以排除。
• 反解失败位置：67235 = ((1*128+32)*30+2)*14+7、68223 = ((1*128+34)*30+13)*14+1；「实际 0.197121」恰好等于 bias[34]，而 322.5 正好是该输出行 conv 的最大值 ⇒ 整行 (b=1, oh=13) 的 conv 贡献全为 0。把整行置 0 重放：坏点 1792、max err 322.5、argmax 落在 (co=34, ow=1)，与报错逐位一致。
• 但 case1 的最大误差只有 32.7（≈ 一个 (i,j) 项的量级），没有任何一行整行为 0 ⇒ 同一份输入、同一条 kernel，两次运行的损伤形态不同 ⇒ 输出不是输入的确定性函数。

第 2 步：用 diag_part2.py 三分法二分（Q1 确定性 / Q2 损伤图 / Q3 退化流水切块）

修复前：

$ python3 diag_part2.py
=== Q1：同一参数连跑 3 次，out.bin 是否一致 ===
  run0: md5=afc3b9d52efcdf1f30a1fe1a325a6a74
  run1: md5=bf24bd2a1cf0e33cc453656395c64e5c
  run2: md5=bf24bd2a1cf0e33cc453656395c64e5c
  结论：**不确定**（每次不同）→ 设备侧同步/竞态（AIC↔AIV 或流水覆写）

=== Q2：损伤图 ===
[case] bias=0
  坏点 107440 / 215040（49.963%），最大绝对误差 208.9
  按 batch 分布（每 batch 53760 点）：[0, 53726, 53714, 0]
  '输出 == bias'（该点 conv 贡献为 0）的点数：0
[case] bias=1
  坏点 107436 / 215040（49.961%），最大绝对误差 32.7
  按 batch 分布（每 batch 53760 点）：[0, 53722, 53714, 0]
# （损伤行列表略：b=1/b=2 的每一行都是 ~1790/1792 坏，b=0/b=3 一行都没有）

=== Q3：改流水切块，看损伤是否消失 ===
[单行 tile（--tile-rows 30）]      坏点 0 / 215040（0.000%），最大绝对误差 0.0001221
[--tile-rows 7]                    坏点 162907 / 215040（75.757%），最大绝对误差 33.02
[单 batch tile（--tile-batch 4）]  坏点 0 / 215040（0.000%），最大绝对误差 0.0001221
[两者都退化]                       坏点 0 / 215040（0.000%），最大绝对误差 0.0001221

三条判据：

1. Q1 三次 md5 不一致 ⇒ 竞态，不是纯逻辑 bug；
2. Q2 坏点在 batch 维度成块（[0, 53726, 53714, 0]：b=0/b=3 全对、b=1/b=2 几乎全错），且最大误差落在"整行 conv 为 0"这一形态上；
3. Q3 一旦退化成 ≤4 个流水单元就全部正确（残差只剩 fp32 尾数 1.2e-4）⇒ 故障在 host 双缓冲流水，与 kernel 的计算逻辑无关（期间怀疑过的「9 次 GetTensorC + UB Add 累加」与「缺 MIX 任务类型声明」由此排除）。

第 3 步：根因 —— host 缓冲只有 2 份，而 CPU 不受 stream 顺序约束

• 把「unit 2/3 的 H2D 实际读到的是 unit 6/7（即 b=3）的打包数据」建成模型，逐行最大误差与 diag 观测的 24 行中 22 行逐位相同（18.8 24.6 32.7 32.7 25.8 … / 21.8 19.3 25.9 27.1 …）——不是巧合。
• 机制：imPackedHost_[p] 只有 2 份，而 RepackIm 是先 memset 全 0、再逐段填；H2D 是异步的，又被 aclrtStreamWaitEvent(copyStream_, calcEv_[p]) 挡在 copyStream 队列里可能尚未执行；与此同时 CPU 侧的三个调用全是异步入队，CPU 会一路 RepackIm 跑到后面第 4 个单元去。等 DMA 真正开跑，缓冲里已经是后面那个单元的数据——甚至只是 memset 出来的全 0（这正是"整行 conv 为 0"、输出恰好等于 bias 的来源）。
• 为什么首尾 batch 看起来没事：unit 0/1 的拷贝排在空队列上立即执行；unit 6/7 之后 CPU 已无后续 RepackIm。只有中间"被卡在队列里"的单元才会读到被提前覆盖的缓冲。
• 设备缓冲那半条保护是对的（aclrtStreamWaitEvent(copyStream_, calcEv_[p])），缺的是 host 这半条。

第 4 步：修复（conv2d.asc，两处）

// ① RepackIm 之前：host 阻塞，等 unit-2 那次 H2D 真正读完这块 host 缓冲
if (unitIdx >= 2) {
    aclrtSynchronizeEvent(copyEv_[p]);
}

// ② 每轮收尾：连 copyStream 一起同步（跨轮次复用 host 缓冲的窗口）
aclrtSynchronizeStream(calcStream_);
aclrtSynchronizeStream(copyStream_);

代价：CPU 不再能"领先 2 个单元以上"（本来就是 2 份缓冲的语义）。要恢复重叠度，把 host 缓冲扩成 4 份、每份配一个 event 即可，代价只是几 MB 主机内存。

第 5 步：修复后复跑

$ python3 diag_part2.py --skip-variants
=== Q1：同一参数连跑 3 次，out.bin 是否一致 ===
  run0: md5=a90e72214417ebd4d7adc3e577cf7623
  run1: md5=a90e72214417ebd4d7adc3e577cf7623
  run2: md5=a90e72214417ebd4d7adc3e577cf7623
  结论：**确定性**（每次完全相同）→ 查地址/布局（读到别人的数据/越界）
# （Q1 的结论行是脚本的机械提示语：确定性 + 下面 Q2 的 0 坏点 = 已修好）

=== Q2：损伤图 ===
[case] bias=0
  坏点 0 / 215040（0.000%），最大绝对误差 0.0001221
  按 batch 分布（每 batch 53760 点）：[0, 0, 0, 0]
  '输出 == bias'（该点 conv 贡献为 0）的点数：0
[case] bias=1
  坏点 0 / 215040（0.000%），最大绝对误差 0.0001221
  按 batch 分布（每 batch 53760 点）：[0, 0, 0, 0]
  '输出 == bias'（该点 conv 贡献为 0）的点数：0

$ python3 test_harness.py --quick
[info] soc_version = Ascend910B4
[info] 参考实现：numpy + torch 双基准（两者必须一致）

=== 正确性用例 ===
[case] bias=0 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）
[case] bias=1 maxpool=0 小图
  [ OK ] float32: 215040 个元素全部匹配（rtol=atol=0.01）
  [ OK ] float16: 215040 个元素全部匹配（rtol=atol=0.01）

正确性总判定： [ OK ] 全部通过

这一条的普适教训：异步拷贝有两端缓冲，两端要分别写保护——设备缓冲靠 stream 事件（aclrtStreamWaitEvent），host 缓冲只能靠 host 侧阻塞（aclrtSynchronizeEvent）。只做前者时，只要 CPU 比"被事件卡住的拷贝"跑得快，就会把还在队列里等着被读的 host 缓冲覆盖掉；症状是"数值看起来像另一个 tile / 一片 0"，且随流水切块方式而变——先看 --tile-rows / --tile-batch 退化后是否消失，就能一刀切开 host 流水与 kernel 逻辑。

2.6.2 排障实录：构建/启动 · 数据排布 · kernel 结构 · 运行调试

构建 / 启动层（症状：kernel 不执行、进度字恒 0、stream 挂）：

1. 构建流：CANN 8.5 用官方第二代 ASC 语言流（find_package(ASC) + .asc 单文件，kernel 与 host main 同文件），不要用第一代 ascendc_library(auto_gen)——它对 MIX/高阶 API kernel 的注册失败（RegisterAscendBinary mix ret 107000），任务贴了但设备端不执行。
2. 链接器：用 bisheng 链接（纯 ASC target 自动、混合 target 设 LINKER_LANGUAGE ASC）。它自带 CANN 库搜索路径并解析 .so 传递依赖；用 GNU ld 会一路缺 libregister / libascendcl / libunified_dlog / libc_sec。
3. tiling 传参：kernel 入口的 tiling 必须按值结构体传（AscendC::tiling::TCubeTiling tiling）。GM 缓冲 + 入口 DecodeTiling 的形式在 8.5 下任务静默不执行。
4. REGIST 位置与 workspace：REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), mm, &tiling) 必须在入口函数内、workspace 用 GetSysWorkSpacePtr()（__kfc_workspace__ 参数）。把自 malloc 的 workspace 值传进去会在 REGIST 死锁。
5. 不要用 isTransposeB=true：实测在 8.5 下 C 只会算出第 1 列（N 方向丢失，其余为垃圾）。B 直接给 (K,N) row-major 连续段、isTransposeB=false——官方样例（02/03）全部如此。"把转置交给框架"的小聪明在小 N 上翻车。

数据排布（症状：数值系统性偏差、无 pattern）：

1. B 段 = (Cin, OWp) 连续（host 预排布）：每个 (i,j) 一个段，段内通道主序、OWp 列连续；输入列越界（ow'+j ≥ W）补 0。这才是"零搬运 shift"的官方形态（配合 5 的 isTransposeB=false）。
2. RepackW 必须逐元素拷贝：权重是 4 维 (C_out, C_in, FH, FW)，c 的存储步长是 FH*FW=9——一行 memcpy 拷 C_in 个"连续"字节会把 c 拷成错序。所有多维重排都要按源步长逐元素写，不能靠连续 memcpy。症状：值量级全对（5–35% 随机偏差、co 相关），排查极难。
3. DataCopyParams 的单位是 32 B DataBlock：blockLen / srcStride / dstStride 都按「字节数 / 32」算，不是「元素数 / 16」。fp32 按 OWp/16 写 blockLen 会每行只搬半行（列 0..7 有值、其余 0）。

kernel 结构（症状：输出 ≈ 单次 (i,j)、或卡死）：

1. 跨 FH*FW 次 (i,j) 复用同一 Matmul 对象：GetTensorC 的"首块"判定用跨循环的 step（cLocal 是否已初始化），不要用 while 内的 subCnt——否则每个 (i,j) 都覆盖 cLocal、FH*FW 项从不累加。
2. while(Iterate) 每轮必须 GetTensorC：只 Iterate 不取、或多次 Iterate 后取，会挂起。GetTensorC(dst, 0, true) 的 enSequentialWrite=true 照官方。
3. 写回：输出用 padded 布局（matmul 的 N 直接取 OWp=align16(W_out)），一条带步长 DataCopy 落 GM、host 裁列——避免奇数列的 DataCopyPad 问题（8.5 的 DataCopyPad 只有 GM→UB 方向）。

运行 / 调试：

1. 环境变量：setenv.bash 的完整 LD_LIBRARY_PATH 不能覆盖（CANN 库散在 lib64、devlib、lib64/plugin/*、opp/... 等多处）；手动跑用「前缀追加」$PWD/build-npu/lib:$LD_LIBRARY_PATH；ASCEND_INSTALL_PATH 每开新 shell 要重设。
2. kernel 挂起时的调试手段：AscendC::printf 在挂起时不 flush（看不到）→ 用「GM 进度字 + host 不 sync 轮询」（kernel 每阶段写 magic 到 GM 尾部，host aclrtMemcpy 轮询读回）；隔离写回用「Duplicate 填 UB 全 1.5」实验；对账用「设备采样 vs host 按同一布局重现」——这次正是靠后者揪出 RepackW 的错序拷贝。

2.6.3 排障实录：三个"静默失败"与两个环境坑

场景：把前置算子的产出从 5HD (B,C1,H,W_pad,C0) 改成 X′ (B,H,FW,C_in,OWp)、主算子从 CubeFormat::NZ 改成 ND。改造只动了两个函数，但从"输出整片为 0"到"数值正确"花了 8 轮 —— 因为踩到的全是不报错的坑。

① MIX 声明让 AIV 把输出覆盖成 0

• 现象：20 核输出整片 0，单核（--block-dim 1）却正常；加 --bias 时输出恰好等于 bias。
• 根因：kernel 声明了 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2) → 函数体在 AIC 与 AIV 上都执行 → AIV 侧执行到 DataCopy(outGm, cLocal, ...) 时，cLocal（UB）里是未初始化的 0 → 把 AIC 算好的结果覆盖成 0。
• 修法：用高阶 Matmul 时不要声明 MIX —— C → UB 那段是 Matmul 库内部经 KFC 机制（REGIST_MATMUL_OBJ + __kfc_workspace__）完成的，不需要用户代码跑在 AIV 上。反之，自己发 LoadData/Mmad 的低阶写法必须拆成 AIC_ONLY / AIV_ONLY 两个 kernel。

② 前置算子改了产出格式，却忘了同步放大 UB 缓冲

• 现象：X′ 验证 100% 正确，但 W′ 是整整齐齐的全 0（不是乱码）。
• 根因：outQ 仍按旧 5HD 的 widthPad*C0 = 512 个元素分配，而新格式每个工作项要写 C_in*OWp = 2048 → 越界 4 倍 → 溢出踩坏了相邻的 wInQ/wOutQ 队列元数据 → 权重那次 DataCopy 到 GM 从未发生 → W′ 保持 aclrtMalloc 时的零。
• 为什么 X′ 反而正确：溢出只写到 out 之外，没有改 out 自身的内容。
• 判别特征："某个下游缓冲整齐全 0" —— 这是"没搬"而不是"算错"。

③ coPerCore 用了函数参数的副本（此时还是 0）

• 现象：W′ 仍然全 0，而这次缓冲大小已经修好。
• 定位：设备端 printf 打出 [W] coPerCore=0 → 权重循环一次都没跑。
• 根因：td_.coPerCore = (outChannels + a.blockDim - 1) / a.blockDim; 用的是函数参数 a 的副本，而真实核数赋给了成员 a_.blockDim；此时 a.blockDim == 0 → coPerCore = 0 → for (co = 0; co < 0; ...) 静默不执行。
• 修法：用真实核数（a_.blockDim，或前置算子实际发射的核数）。

三个坑的共同点：都是"静默失败" —— 不崩、不报错、数据就是没了。由此得到三条方法论：

1. 把每一个中间产物都 dump 出来逐元素验证 —— 这一轮就是 X′ 100% → W′ 全 0 → 设备打印 [W] 三步逼出根因；
2. 判别"算错"还是"没算"：加 --bias 再跑。bias 是在 matmul 之后加的：若输出恰好等于 bias，说明写回路径正常、C 恒为 0（问题在 matmul 侧）；若仍全 0，则写回路径坏了；
3. 多核 / 单核对比：--block-dim 1 正常而 20 核不正常 → 一定与多核有关（竞态或 AIV 覆盖）。

3. 参考资料

• Cube矩阵计算编程
• Matmul 使用说明（公式、轴定义、实现原理）
• Matmul 模板参数（MatmulType 全参数）
• MatmulConfig（模板选择与配置参数）
• Matmul Tiling 参数说明（baseM/stepM/singleCoreM 等的图形化解释）
• Batch Matmul 基础功能
• Matmul 成员函数
• Matmul 更多样例
• Conv2DV2