Assignment 4（Part 1）：Ascend C 编程接口 —— 向量加与矩阵转置

📖 因篇幅过长，本报告拆成两篇。本篇是 Part 1；Part 2（融合的 Convolution + MaxPool 层）见 asst4_part2.md。 两篇共用的环境配置（CANNLab 环境、CANN 环境变量、编译命令）在本篇 §0，Part 2 不再重复。

Assignment 4 要求为机器学习加速器实现并优化 kernel：Part 1（30 分） 用向量加熟悉编程模型与数据搬运，并自己写一个矩阵转置 kernel；Part 2（70 分） 实现融合的 convolution + maxpool 层，中间结果不能落到片外。

课程原始版本面向 AWS Trainium2（NKI）。本报告改为在 华为昇腾 上完成：CANNLab 的一个 NPU（A3 = Ascend 910C，Ascend910_9362）。两个平台的逐项对照、以及每一处"为什么必须这么改"写在配套文档里：

• 作业 README 的昇腾对照版：asst4-ascend/README.md
• Trainium2 ↔ 昇腾 的完整适配说明：asst4-ascend/ascend_adaptation.md

0. 环境配置

0.1 方案选择：CANNLab（替代 AWS Trainium2 容量块）

选 CANNLab 的原因很直接：不需要付费容量块、不需要申请白名单，代价是要把 NKI 代码用 Ascend C 重写一遍（这正是配套文档在做的事）。

两点使用约束：

• 本作业是单算子调优，必须选单卡/单 NPU 规格。CANNLab 按"卡时"计费，选 8 卡规格会按卡的倍数扣，100 卡时可用的墙钟时间直接缩到 1/8。
• 用完立刻关机（平台提供闲置自动关机开关）。

0.2 环境确认

登录 CANNLab 后第一件事是确认自己拿到的是哪一代卡 —— 后面 --soc-version、以及在 Tiling 里取的容量值都依赖它。

npu-smi info

A3卡输出：

developer@8de626ef107b43269a92e55bdfcaea3c:/mnt/workspace$ npu-smi info
+------------------------------------------------------------------------------------------------+
| npu-smi 25.5.5                   Version: 25.5.5                                               |
+---------------------------+---------------+----------------------------------------------------+
| NPU   Name                | Health        | Power(W)    Temp(C)           Hugepages-Usage(page)|
| Chip  Phy-ID              | Bus-Id        | AICore(%)   Memory-Usage(MB)  HBM-Usage(MB)        |
+===========================+===============+====================================================+
| 4     Ascend910           | OK            | 163.0       43                0    / 0             |
| 0     8                   | 0000:0B:00.0  | 0           0    / 0          2908 / 65536         |
+------------------------------------------------------------------------------------------------+
| 4     Ascend910           | OK            | -           42                0    / 0             |
| 1     9                   | 0000:0A:00.0  | 0           0    / 0          2870 / 65536         |
+===========================+===============+====================================================+
+---------------------------+---------------+----------------------------------------------------+
| NPU     Chip              | Process id    | Process name             | Process memory(MB)      |
+===========================+===============+====================================================+
| No running processes found in NPU 4                                                            |
+===========================+===============+====================================================+

A2 卡输出：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ npu-smi info
+------------------------------------------------------------------------------------------------+
| npu-smi 25.5.1                   Version: 25.5.1                                               |
+---------------------------+---------------+----------------------------------------------------+
| NPU   Name                | Health        | Power(W)    Temp(C)           Hugepages-Usage(page)|
| Chip                      | Bus-Id        | AICore(%)   Memory-Usage(MB)  HBM-Usage(MB)        |
+===========================+===============+====================================================+
| 4     910B4               | OK            | 86.4        34                0    / 0             |
| 0                         | 0000:81:00.0  | 0           0    / 0          2886 / 32768         |
+===========================+===============+====================================================+
+---------------------------+---------------+----------------------------------------------------+
| NPU     Chip              | Process id    | Process name             | Process memory(MB)      |
+===========================+===============+====================================================+
| No running processes found in NPU 4                                                            |
+===========================+===============+====================================================+

echo $ASCEND_HOME_PATH
cat $ASCEND_HOME_PATH/version.info

A3卡输出：

developer@8de626ef107b43269a92e55bdfcaea3c:/mnt/workspace$ echo $ASCEND_HOME_PATH
/home/developer/Ascend/cann-9.0.0
developer@8de626ef107b43269a92e55bdfcaea3c:/mnt/workspace$ cat $ASCEND_HOME_PATH/version.info
cat: /home/developer/Ascend/cann-9.0.0/version.info: No such file or directory

说明本机 $ASCEND_HOME_PATH 是 /home/developer/Ascend/cann-9.0.0（用户目录下的 CANN 9.0.0），版本文件不在根目录，与文档里默认的 $ASCEND_HOME_PATH/version.info 不一样。

所以换成先看目录、再按文件名找：

ls $ASCEND_HOME_PATH
find $ASCEND_HOME_PATH -maxdepth 4 -type f -name "version*"

输出：

developer@8de626ef107b43269a92e55bdfcaea3c:/mnt/workspace$ ls $ASCEND_HOME_PATH
aarch64-linux  asc  cann_uninstall.sh  compat    conf    fwkacllib  lib64  ops_uninstall.sh  python   set_env.sh  test-ops  tools
arm64-linux    bin  combo_script       compiler  devlib  include    opp    pkg_inc           runtime  share       toolkit   var

developer@8de626ef107b43269a92e55bdfcaea3c:/mnt/workspace$ find $ASCEND_HOME_PATH -maxdepth 4 -type f -name "version*"
/home/developer/Ascend/cann-9.0.0/opp/version.info
/home/developer/Ascend/cann-9.0.0/python/site-packages/te/version.py
/home/developer/Ascend/cann-9.0.0/tools/msmemscope/version.txt
/home/developer/Ascend/cann-9.0.0/share/info/tbe-tik/version.info
/home/developer/Ascend/cann-9.0.0/share/info/ops_nn/version.info
/home/developer/Ascend/cann-9.0.0/share/info/ge-compiler/version.info
/home/developer/Ascend/cann-9.0.0/share/info/msserviceprofiler/version.info
/home/developer/Ascend/cann-9.0.0/share/info/pyACL/version.info
/home/developer/Ascend/cann-9.0.0/share/info/opbase/version.info
/home/developer/Ascend/cann-9.0.0/share/info/ops_legacy/version.info
/home/developer/Ascend/cann-9.0.0/share/info/ops_cv/version.info
/home/developer/Ascend/cann-9.0.0/share/info/hccl/version.info
/home/developer/Ascend/cann-9.0.0/share/info/hixl/version.info
/home/developer/Ascend/cann-9.0.0/share/info/pto_isa/version.info
/home/developer/Ascend/cann-9.0.0/share/info/mindstudio-operator-tools/version.info
/home/developer/Ascend/cann-9.0.0/share/info/oam_tools/version.info
/home/developer/Ascend/cann-9.0.0/share/info/ops_transformer/version.info
/home/developer/Ascend/cann-9.0.0/share/info/ge-executor/version.info
/home/developer/Ascend/cann-9.0.0/share/info/mindstudio-opprof/version.info
/home/developer/Ascend/cann-9.0.0/share/info/runtime/version.info
/home/developer/Ascend/cann-9.0.0/share/info/dvpp/version.info
/home/developer/Ascend/cann-9.0.0/share/info/mstx/version.info
/home/developer/Ascend/cann-9.0.0/share/info/ms_fmk_transplt/version.info
/home/developer/Ascend/cann-9.0.0/share/info/mspti/version.info
/home/developer/Ascend/cann-9.0.0/share/info/operator_cmp/version.info
/home/developer/Ascend/cann-9.0.0/share/info/aoe/version.info
/home/developer/Ascend/cann-9.0.0/share/info/ops_math/version.info
/home/developer/Ascend/cann-9.0.0/share/info/simulator/version.info
/home/developer/Ascend/cann-9.0.0/share/info/graph_autofusion/version.info
/home/developer/Ascend/cann-9.0.0/share/info/ops_ras/version.info
/home/developer/Ascend/cann-9.0.0/share/info/dflow-executor/version.info
/home/developer/Ascend/cann-9.0.0/share/info/msprof/version.info
/home/developer/Ascend/cann-9.0.0/share/info/asc-devkit/version.info
/home/developer/Ascend/cann-9.0.0/share/info/acl_extend/version.info
/home/developer/Ascend/cann-9.0.0/share/info/metadef/version.info
/home/developer/Ascend/cann-9.0.0/share/info/mindstudio-sanitizer/version.info
/home/developer/Ascend/cann-9.0.0/share/info/msmemscope/version.info
/home/developer/Ascend/cann-9.0.0/share/info/asc-tools/version.info
/home/developer/Ascend/cann-9.0.0/share/info/ncs/version.info
/home/developer/Ascend/cann-9.0.0/share/info/mindstudio-debugger/version.info
/home/developer/Ascend/cann-9.0.0/share/info/hcomm/version.info
/home/developer/Ascend/cann-9.0.0/share/info/test-ops/version.info
/home/developer/Ascend/cann-9.0.0/share/info/bisheng-compiler/version.info
/home/developer/Ascend/cann-9.0.0/share/info/ascendnpu-ir/version.info

然后读版本号。上面 44 个命中里，对作业最有用的是编译器与算子库这两个：

cat $ASCEND_HOME_PATH/compiler/version.info
cat $ASCEND_HOME_PATH/opp/version.info
head -20 $ASCEND_HOME_PATH/aarch64-linux/include/version
# 顺手记下：0.2 的 ls 里有 aarch64-linux/，片上容量的配置文件就在它下面
ls $ASCEND_HOME_PATH/aarch64-linux/data/platform_config/ | head

输出：

developer@8de626ef107b43269a92e55bdfcaea3c:/mnt/workspace$ cat $ASCEND_HOME_PATH/compiler/version.info
Version=9.0.0
version_dir=cann
required_package_runtime_version="9.0"
required_package_ge-executor_version="9.0"
required_package_metadef_version="9.0"
required_package_opbase_version="9.0"
required_package_bisheng-compiler_version="9.0"
required_package_tbe-tik_version="9.0"
timestamp=20260428_134817545

developer@8de626ef107b43269a92e55bdfcaea3c:/mnt/workspace$ cat $ASCEND_HOME_PATH/opp/version.info
Version=9.0.0
version_dir=cann
required_package_runtime_version=">=8.5"
required_package_ge-executor_version=">=8.5"
required_package_metadef_version=">=8.5"
required_package_asc-devkit_version=">=8.5"
required_package_bisheng-compiler_version=">=8.5"
timestamp=20260428_134817545

developer@8de626ef107b43269a92e55bdfcaea3c:/mnt/workspace$ head -20 $ASCEND_HOME_PATH/aarch64-linux/include/version
head: error reading '/home/developer/Ascend/cann-9.0.0/aarch64-linux/include/version': Is a directory

developer@8de626ef107b43269a92e55bdfcaea3c:/mnt/workspace$ ls $ASCEND_HOME_PATH/aarch64-linux/data/platform_config/ | head
Ascend031.ini
Ascend035A.ini
Ascend035B.ini
Ascend035.ini
Ascend310B1.ini
Ascend310B2.ini
Ascend310B3.ini
Ascend310B4.ini
Ascend310P1.ini
Ascend310P3.ini

列表里有一处值得注意：cat $ASCEND_HOME_PATH/compiler/version.info 能读到内容，但它不在上面这份 -type f 列表里（列表里只有 opp/version.info 和一串 share/info/*/version.info）。多半是因为 compiler 是个符号链接，而 find 默认不跟随符号链接（要跟随得加 -L）。可用 ls -ld $ASCEND_HOME_PATH/compiler 验证（可选）。

版本号前两条已经给全：compiler/version.info 与 opp/version.info 都是 Version=9.0.0，与目录名 cann-9.0.0 一致。

另外 platform_config/ 那份列表是按字母序排的，head 只截到 Ascend310P3（031 / 035 / 310 系列）；910B、910_93xx（A3）、950 等都在后面，要用的时候直接 grep 910 或指定文件名读。

A2 卡输出：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ echo $ASCEND_HOME_PATH
cat $ASCEND_HOME_PATH/version.info
/usr/local/Ascend/cann-8.5.0
cat: /usr/local/Ascend/cann-8.5.0/version.info: No such file or directory

0.3 CANN 环境变量

CANN 的环境脚本有两个位置（根目录的 set_env.sh 与 bin/setenv.bash），先确认再 source。本机两个都在，等价可用：

ls $ASCEND_HOME_PATH/set_env.sh $ASCEND_HOME_PATH/bin/setenv.bash 2>/dev/null
source $ASCEND_HOME_PATH/set_env.sh
which bisheng msprof
# NPU ID 用 npu-smi info 里看到的那个（本机是 4，换个会话可能变，不要写死别的）
npu-smi info -t board -i 4 -c 0

输出：

developer@8de626ef107b43269a92e55bdfcaea3c:/mnt/workspace$ ls $ASCEND_HOME_PATH/set_env.sh $ASCEND_HOME_PATH/bin/setenv.bash 2>/dev/null
/home/developer/Ascend/cann-9.0.0/bin/setenv.bash  /home/developer/Ascend/cann-9.0.0/set_env.sh

developer@8de626ef107b43269a92e55bdfcaea3c:/mnt/workspace$ source $ASCEND_HOME_PATH/set_env.sh

developer@8de626ef107b43269a92e55bdfcaea3c:/mnt/workspace$ which bisheng msprof
/home/developer/Ascend/cann-9.0.0/bin/bisheng
/home/developer/Ascend/cann-9.0.0/bin/msprof

developer@8de626ef107b43269a92e55bdfcaea3c:/mnt/workspace$ npu-smi info -t board -i 4 -c 0
        NPU ID                         : 4
        NPU Name                       : 9362
        Chip ID                        : 0
        Chip Name                      : Ascend910
        Chip Version                   : V1
        Board ID                       : 0xb3
        PCB ID                         : NA
        BOM ID                         : 1
        VDie ID                        : 6F43F664 00C10159 7302F8D3 BE0B0B56 144301E3
        NDie ID                        : 00000000 00000000 00000000 00000000 00000000
        Chip Position ID               : 0
        PCIe Bus Info                  : 0000:0B:00.0
        Firmware Version               : NA

→ 由此得到 --soc-version=Ascend910_9362（Chip Name + _ + NPU Name）。后面 0.5 的构建和 Step 2b 的 msprof 都用这个值；换个会话机型可能变，所以这个值要每次现查，不要沿用。

从以上输出里要读出这几个值，后面每一步都会用到：

判机型只能靠 npu-smi。platform_config/ 目录里有 60 多份不同 SoC 的配置（含 910A/910B/A3/950/310/Kirin），它是编译器的查找表，不是当前机器的型号证据。

0.4 代码位置

Part 1 的代码在 asst4-ascend/part1/（对应 Trainium2 版的 asst4-trainium2/part1/）：

0.5 构建

--soc-version 用 0.2 里确定的那个值。

cd asst4-ascend/part1
bash run.sh --soc-version Ascend910_9362 --build-only

输出：

developer@8de626ef107b43269a92e55bdfcaea3c:.../workspace/asst4-ascend/part1$ bash run.sh --soc-version Ascend910_9362 --build-only
-- [npu_lib] ASCENDC_CMAKE_DIR = /home/developer/Ascend/cann-9.0.0/compiler/tikcpp/ascendc_kernel_cmake
-- [npu_lib] SOC_VERSION      = Ascend910_9362
-- [npu_lib] KERNEL_FILES     = /mnt/workspace/asst4-ascend/part1/kernels.cpp
...（省略 3 行 CMake / ExternalProject 样板输出）...
[  2%] Performing build step for 'ascendc_kernels_npu_precompile'
Consolidate compiler generated dependencies of target precompile_obj
[100%] Building CXX object CMakeFiles/precompile_obj.dir/mnt/workspace/asst4-ascend/part1/kernels.cpp.o
[100%] Built target precompile_obj
[100%] Built target check_src_template
[  4%] No install step for 'ascendc_kernels_npu_precompile'
[  6%] Completed 'ascendc_kernels_npu_precompile'
[ 18%] Built target ascendc_kernels_npu_precompile
[ 20%] Performing build step for 'ascendc_kernels_npu_preprocess'
...（省略 3 行 CMake / ExternalProject 样板输出）...
Consolidate compiler generated dependencies of target aic_obj
Consolidate compiler generated dependencies of target preprocess_obj
Consolidate compiler generated dependencies of target aiv_obj
[ 33%] Building CXX object CMakeFiles/aic_obj.dir/mnt/workspace/asst4-ascend/part1/kernels.cpp.o
[ 66%] Building CXX object CMakeFiles/preprocess_obj.dir/mnt/workspace/asst4-ascend/part1/kernels.cpp.o
[100%] Building CXX object CMakeFiles/aiv_obj.dir/mnt/workspace/asst4-ascend/part1/kernels.cpp.o
[100%] Built target preprocess_obj
[100%] Built target aic_obj
[100%] Built target aiv_obj
/home/developer/Ascend/cann-9.0.0/aarch64-linux/ccec_compiler/bin/ld.lld -m aicorelinux -Ttext=0 /mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels_npu_preprocess-prefix/src/ascendc_kernels_npu_preprocess-build/CMakeFiles/aic_obj.dir/mnt/workspace/asst4-ascend/part1/kernels.cpp.o -o /mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels_npu_preprocess-prefix/src/ascendc_kernels_npu_preprocess-build/CMakeFiles/aic_obj.dir/mnt/workspace/asst4-ascend/part1/kernels.cpp.o/home/developer/Ascend/cann-9.0.0/aarch64-linux/ccec_compiler/bin/ld.lld -m aicorelinux -Ttext=0 /mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels_npu_preprocess-prefix/src/ascendc_kernels_npu_preprocess-build/CMakeFiles/aiv_obj.dir/mnt/workspace/asst4-ascend/part1/kernels.cpp.o -o /mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels_npu_preprocess-prefix/src/ascendc_kernels_npu_preprocess-build/CMakeFiles/aiv_obj.dir/mnt/workspace/asst4-ascend/part1/kernels.cpp.o[100%] Built target merge_aiv_obj_text
[100%] Built target merge_aic_obj_text
[WARNING]: Multiple kernel functions are detected. It is recommended to define only one kernel function per file.
[100%] Built target _host_cpp
[ 22%] No install step for 'ascendc_kernels_npu_preprocess'
[ 25%] Completed 'ascendc_kernels_npu_preprocess'
[ 36%] Built target ascendc_kernels_npu_preprocess
...（省略 12 行 CMake / ExternalProject 样板输出）...
[ 65%] Performing configure step for 'ascendc_kernels_npu_aiv_device'
[ 68%] Performing configure step for 'ascendc_kernels_npu_aic_device'
[ 70%] Performing configure step for 'ascendc_kernels_npu_host'
...（省略 41 行 CMake / ExternalProject 样板输出）...
CMake Warning:
  Manually-specified variables were not used by the project:

    ASCEND_PYTHON_EXECUTABLE


CMake Warning:
  Manually-specified variables were not used by the project:

    BUILD_MODE


...（省略 4 行 CMake / ExternalProject 样板输出）...
[ 72%] Performing build step for 'ascendc_kernels_npu_aic_device'
[ 75%] Performing build step for 'ascendc_kernels_npu_host'
[ 77%] Performing build step for 'ascendc_kernels_npu_aiv_device'
[ 79%] No install step for 'ascendc_kernels_npu_aic_device'
[100%] Building CXX object CMakeFiles/host_bisheng_obj.dir/mnt/workspace/asst4-ascend/part1/kernels.cpp.o
[100%] Building CXX object CMakeFiles/device_aiv_obj.dir/mnt/workspace/asst4-ascend/part1/build-npu/auto_gen/ascendc_kernels_npu/auto_gen_kernels.cpp.o
[ 81%] Completed 'ascendc_kernels_npu_aic_device'
[ 81%] Built target ascendc_kernels_npu_aic_device
[100%] Built target host_bisheng_obj
[ 84%] Performing install step for 'ascendc_kernels_npu_host'
Consolidate compiler generated dependencies of target host_bisheng_obj
[100%] Building CXX object CMakeFiles/host_bisheng_obj.dir/mnt/workspace/asst4-ascend/part1/kernels.cpp.o
[100%] Built target host_bisheng_obj
Install the project...
-- Install configuration: "Release"
-- Installing: /mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels_npu_host_dir/./objects-Release/host_bisheng_obj/mnt/workspace/asst4-ascend/part1/kernels.cpp.o
[ 86%] Completed 'ascendc_kernels_npu_host'
[ 86%] Built target ascendc_kernels_npu_host
[100%] Built target device_aiv_obj
/home/developer/Ascend/cann-9.0.0/aarch64-linux/ccec_compiler/bin/ld.lld  -m aicorelinux -r  -Ttext=0 /mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels_npu_aiv_device-prefix/src/ascendc_kernels_npu_aiv_device-build/CMakeFiles/device_aiv_obj.dir/mnt/workspace/asst4-ascend/part1/build-npu/auto_gen/ascendc_kernels_npu/auto_gen_kernels.cpp.o -static -o /mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels_npu_aiv_device_dir/device_aiv.o
[100%] Built target merge_aiv_device_obj
[ 88%] No install step for 'ascendc_kernels_npu_aiv_device'
[ 90%] Completed 'ascendc_kernels_npu_aiv_device'
[ 90%] Built target ascendc_kernels_npu_aiv_device
/home/developer/Ascend/cann-9.0.0/aarch64-linux/ccec_compiler/bin/ld.lld -x -m aicorelinux   -Ttext=0 /mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels_npu_aiv_device_dir/device_aiv.o -static -o /mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels_npu_merge_obj_dir/device_aiv.o
[ 90%] Built target ascendc_kernels_npu_merge_obj
[ 93%] Building CXX object CMakeFiles/ascendc_kernels_npu_host_stub_obj.dir/auto_gen/ascendc_kernels_npu/host_stub.cpp.o
[ 93%] Built target ascendc_kernels_npu_host_stub_obj
[ 95%] Linking CXX shared library lib/libascendc_kernels_npu.so
/home/developer/Ascend/cann-9.0.0/bin/ascendc_pack_kernel /mnt/workspace/asst4-ascend/part1/build-npu/CMakeFiles/ascendc_kernels_npu_host_stub_obj.dir/auto_gen/ascendc_kernels_npu/host_stub.cpp.o /mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels_npu_merge_obj_dir/device_aiv.o 1 /mnt/workspace/asst4-ascend/part1/build-npu/CMakeFiles/ascendc_kernels_npu_host_stub_obj.dir/auto_gen/ascendc_kernels_npu/host_stub.cpp.o
recompile: /usr/bin/c++ -fPIC -O3 -DNDEBUG -Wl,-z,relro -Wl,-z,now -Wl,-z,noexecstack -s -shared -Wl,-soname,libascendc_kernels_npu.so -o lib/libascendc_kernels_npu.so CMakeFiles/ascendc_kernels_npu_host_stub_obj.dir/auto_gen/ascendc_kernels_npu/host_stub.cpp.o /mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels_npu_host_dir/objects-Release/host_bisheng_obj/mnt/workspace/asst4-ascend/part1/kernels.cpp.o    -L/home/developer/Ascend/cann-9.0.0/lib64  -L/home/developer/Ascend/cann-9.0.0/tools/simulator/Ascend910_9362/lib  /home/developer/Ascend/cann-9.0.0/lib64/libascendc_runtime.a -lascend_dump -lc_sec
[ 95%] Built target ascendc_kernels_npu
[ 97%] Building CXX object CMakeFiles/ascendc_kernels.dir/main.cpp.o
[100%] Linking CXX executable ascendc_kernels
[100%] Built target ascendc_kernels
[INFO ] 构建完成：/mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels

一处 WARNING，说明一下：

[WARNING]: Multiple kernel functions are detected. It is recommended to define only one kernel function per file.

因为 kernels.cpp 里定义了 3 个 __global__ 核函数（vector_add_single、vector_add_stream、matrix_transpose），而 CANN 的构建系统建议一个文件只放一个核函数。

这只是建议：紧接着的 [100%] Built target ascendc_kernels 表明构建照常成功。之所以这么组织，是为了让三个入口共用同一份类模板实现（见 Part 1 的代码说明）；等到 Part 2 核函数变多、或需要走 ACLRT_LAUNCH_KERNEL 自动生成启动桩时，再按核函数拆文件更稳妥。

运行期的库搜索路径（上板前必踩的一步）：kernel 被编成 build-npu/lib/libascendc_kernels_npu.so，可执行文件链接了它，但运行时不一定会去那个目录找 —— 直接跑会得到

build-npu/ascendc_kernels: error while loading shared libraries: libascendc_kernels_npu.so:
 cannot open shared object file: No such file or directory
 [ERROR] kernel 执行失败（exit 127）

原因是 NPU 侧的 ascendc_library() 不把输出位置告诉 CMake，所以 CMake 不会像 CPU 侧那样自动给可执行文件加 RPATH。官方样例的 run.sh 是在运行前 export LD_LIBRARY_PATH=$(pwd)/out/lib:$(pwd)/out/lib64:${_ASCEND_INSTALL_PATH}/lib64:$LD_LIBRARY_PATH。本仓库两处都补上了：run_benchmark.py 运行子进程时自己设（run_env()），CMakeLists.txt 里也设了 BUILD_RPATH。

另外三类 CMake Warning 也不是本工程的问题：

Manually-specified variables were not used by the project:
  DYNAMIC_MODE / INCLUDE_DIR        ← 出现在 ascendc_kernels_npu_precompile 子工程的 configure 之后
  BUILD_MODE                        ← 出现在 ascendc_kernels_npu_aic_device 子工程的 configure 之后
  ASCEND_PYTHON_EXECUTABLE          ← 出现在 ascendc_kernels_npu_host 子工程的 configure 之后

CANN 的 ascendc.cmake 会用 ExternalProject 拉起 5 个子工程（precompile / preprocess / aic_device / aiv_device / host），转发给每个子工程的变量是一整串，而子工程只用其中一部分 —— CMake 就替"用不到的那几个"各发一次提示。

判断依据：我们自己的那次 configure（打印 [npu_lib] ASCENDC_CMAKE_DIR = ... 那一段）后面没有这类 warning，它们全都紧跟在子工程的 Generating done 之后。属于"传多了"的无害提示，不需要修（要修就得改 CANN 安装目录下的 cmake 文件，得不偿失）。

0.6 CPU 域孪生调试（省卡时的关键路径）

Ascend C 支持把同一份 kernels.cpp 编到 CPU 上执行（ICPU_RUN_KF），用来做正确性对比。编译器错误、对齐问题、精度问题都应该在这一步解决，不要上板调试。

cd asst4-ascend/part1
python3 run_benchmark.py --kernel tiled -n 25600 --run-mode cpu

# 如果需要重新build的话
python3 run_benchmark.py --kernel tiled -n 25600 --run-mode cpu --build

输出：

developer@8de626ef107b43269a92e55bdfcaea3c:.../workspace/asst4-ascend/part1$ python3 run_benchmark.py --kernel tiled -n 25600 --run-mode cpu --build
[INFO ] 构建中（soc-version=Ascend910_9362）...
-- [cpu_lib] tikicpulib cmake dir = /home/developer/Ascend/cann-9.0.0/tools/tikicpulib/lib/cmake
-- Configuring done
-- Generating done
-- Build files have been written to: /mnt/workspace/asst4-ascend/part1/build-cpu
Consolidate compiler generated dependencies of target ascendc_kernels_cpu
[ 50%] Built target ascendc_kernels_cpu
Consolidate compiler generated dependencies of target ascendc_kernels
[ 75%] Building CXX object CMakeFiles/ascendc_kernels.dir/main.cpp.o
[100%] Linking CXX executable ascendc_kernels
[100%] Built target ascendc_kernels
[INFO ] 构建完成：/mnt/workspace/asst4-ascend/part1/build-cpu/ascendc_kernels
[INFO ] CPU 域执行：/mnt/workspace/asst4-ascend/part1/build-cpu/ascendc_kernels --kernel single --n 25600 --m 25600 --tile 25600 --block-dim 1 --out /mnt/workspace/asst4-ascend/part1/out/tiled.bin
[TmSim]: Run in serial mode.
[SUCCESS][CORE_0][pid 6064] exit success!
[INFO ] CPU 域执行完成（不产生耗时数据），输出已写入 /mnt/workspace/asst4-ascend/part1/out/tiled.bin

kernel      : tiled  (CPU 域孪生调试，无耗时数据)
tile        : 25600 个元素
[ OK ] 结果与 numpy 参考一致

结果：CPU 域对比通过 —— vector_add_single（tile = 25600、单缓冲、blockDim = 1）的输出与 numpy 参考逐元素一致。这一步验证的是数值逻辑（是否真的算了 z = x + y、有没有对齐/尾块问题），耗时数据要等上板。

1. Part 1: Learning the Ascend C Kernel Interface with Vector Add and Matrix Transpose（30 分）

Part 1 要回答三个问题：分块（Step 1）、搬运粒度与流水（Step 2a/2b）、矩阵转置（Step 3）。本报告逐一记录命令与实测数据。

被测代码与参数含义（对应 Trainium2 版的 ROW_CHUNK / FREE_DIM）：

单位约定：本文的内存容量一律用二进制单位（1 KiB = 1024 B、1 MiB = 1024 KiB），硬件文档里的 "192 KB" 指的就是 192 KiB = 196608 B。而 run_benchmark.py 打印的 MB 是十进制（moved_bytes / 1e6），带宽 GB/s 也是十进制。两者不要混——153600 B 既等于 150 KiB，也等于 0.154 MB（十进制），是同一个数。

一个必须先记住的平台差异：--tile 的最小值是 16，不是 1。因为昇腾的 DataCopy 最小搬运粒度是 32 B = 1 个 DataBlock = 16 个 half，所以 Trainium2 版里 ROW_CHUNK = 1 这种"一次搬一个元素"的写法在昇腾上连合法都做不到。两边的"最小分块"因此在数值上就差了 16 倍。

1.1 Step 1: Chunking Vectors to Parallelize（6 分）

对应原版 Step 1 的三问，在昇腾上改问：

1. 最小合法分块（tile = 16）跑 25600 个元素，耗时多少？
2. 逐步放大 tile 直到接近 UB 能装下的上限，耗时降到多少？相比第 1 问快了多少？为什么？
3. 继续放大到超过 UB 容量，会看到什么报错？为什么？

1.1.1 最小合法分块：tile = 16

cd asst4-ascend/part1
python3 run_benchmark.py --kernel naive -n 25600 --verify

输出：

developer@8de626ef107b43269a92e55bdfcaea3c:.../workspace/asst4-ascend/part1$ python3 run_benchmark.py --kernel naive -n 25600 --verify

kernel      : naive  (可执行文件里的入口: single)
n           : 25600
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 1 个 AI Core
latency     : median 404.687 us   min 400.827 us   avg 404.772 us
bandwidth   : 0.38 GB/s  (按访存量 0.154 MB 估算)
[ OK ] 结果与 numpy 参考一致

如果修改了代码，先做一次带重新配置的构建（RPATH 改动需要重新 cmake），然后上板：

python3 run_benchmark.py --kernel naive -n 25600 --verify --build

如果还想亲眼确认一下诊断，先跑这两条：

ls build-npu/lib/                                       # 应该能看到 libascendc_kernels_npu.so
ldd build-npu/ascendc_kernels | grep ascendc_kernels     # 修复前应显示 "not found"

1.1.2 放大 tile（单缓冲）

--tile 必须是 16 的整数倍、且能整除 -n。25600 = 2^10 × 5^2，满足条件的取值有 21 个，下面挑几档有代表性的逐档跑：

cd asst4-ascend/part1 

for T in 16 128 1024 6400 12800 25600; do echo "===== tile=$T"; python3 run_benchmark.py --kernel tiled -n 25600 --tile $T; done

输出：

developer@8de626ef107b43269a92e55bdfcaea3c:.../workspace/asst4-ascend/part1$ for T in 16 128 1024 6400 12800 25600; do echo "===== tile=$T"; python3 run_benchmark.py --kernel tiled -n 25600 --tile $T; done
===== tile=16

kernel      : tiled  (可执行文件里的入口: single)
n           : 25600
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 1 个 AI Core
latency     : median 403.687 us   min 400.818 us   avg 404.640 us
bandwidth   : 0.38 GB/s  (按访存量 0.154 MB 估算)
===== tile=128

kernel      : tiled  (可执行文件里的入口: single)
n           : 25600
tile        : 128 个元素  (单次搬运 256 B)
block-dim   : 1 个 AI Core
latency     : median 70.492 us   min 68.421 us   avg 70.479 us
bandwidth   : 2.18 GB/s  (按访存量 0.154 MB 估算)
===== tile=1024

kernel      : tiled  (可执行文件里的入口: single)
n           : 25600
tile        : 1024 个元素  (单次搬运 2048 B)
block-dim   : 1 个 AI Core
latency     : median 27.620 us   min 23.771 us   avg 27.738 us
bandwidth   : 5.56 GB/s  (按访存量 0.154 MB 估算)
===== tile=6400

kernel      : tiled  (可执行文件里的入口: single)
n           : 25600
tile        : 6400 个元素  (单次搬运 12800 B)
block-dim   : 1 个 AI Core
latency     : median 25.630 us   min 16.401 us   avg 21.790 us
bandwidth   : 5.99 GB/s  (按访存量 0.154 MB 估算)
===== tile=12800

kernel      : tiled  (可执行文件里的入口: single)
n           : 25600
tile        : 12800 个元素  (单次搬运 25600 B)
block-dim   : 1 个 AI Core
latency     : median 24.710 us   min 14.071 us   avg 22.090 us
bandwidth   : 6.22 GB/s  (按访存量 0.154 MB 估算)
===== tile=25600

kernel      : tiled  (可执行文件里的入口: single)
n           : 25600
tile        : 25600 个元素  (单次搬运 51200 B)
block-dim   : 1 个 AI Core
latency     : median 22.731 us   min 16.670 us   avg 24.314 us
bandwidth   : 6.76 GB/s  (按访存量 0.154 MB 估算)

汇总表（从上面的输出里抄 median_us）：

1.1.3 超过 UB 容量

单缓冲下三个张量（x/y/z）各占 tile × 2 B，UB 容量 192 KiB。先取 tile = 32768 —— 三个张量合计 3 × 32768 × 2 = 196608 B，正好把 192 KiB 占满，先看它能不能过：

python3 run_benchmark.py --kernel tiled -n 65536 --tile 32768

输出：

developer@8de626ef107b43269a92e55bdfcaea3c:.../workspace/asst4-ascend/part1$ python3 run_benchmark.py --kernel tiled -n 65536 --tile 32768

kernel      : tiled  (可执行文件里的入口: single)
n           : 65536
tile        : 32768 个元素  (单次搬运 65536 B)
block-dim   : 1 个 AI Core
latency     : median 32.601 us   min 22.121 us   avg 29.107 us
bandwidth   : 12.06 GB/s  (按访存量 0.393 MB 估算)

它没有挂 —— 三个张量正好把 192 KiB 占满，装得下。那就再取一个肯定超的（384 KiB，是 UB 的 2 倍）：

python3 run_benchmark.py --kernel tiled -n 131072 --tile 65536

输出：

developer@8de626ef107b43269a92e55bdfcaea3c:.../workspace/asst4-ascend/part1$ python3 run_benchmark.py --kernel tiled -n 131072 --tile 65536
/mnt/workspace/asst4-ascend/part1/main.cpp:252 aclError:507035
[ERROR] kernel 执行失败（exit 1）

main.cpp:252 那一行是 CHECK_ACL(aclrtSynchronizeStream(stream))，也就是说：kernel 启动指令下发成功了，是执行时在设备侧出错，等 stream 时把错误抛了上来。507035 在昇腾官方案例库里对应的现象是 "Synchronize stream failed"（设备侧执行失败，plog 里还能看到 aic error code=0x8000000000000000）。放到本次的上下文里，原因很直接：3 × 65536 × 2 B = 393216 B = 384 KiB，是 UB 容量的 2 倍。

UB 容量的两个口径、GetCoreMemSize() 的语义、以及 ReserveLocalMemory() 是否强制 kernel，这一组平台侧的调查整理在 asst4-ascend/README.md 的 Memory Hierarchy 一节。

1.1.4 结果分析

（1）加速比从哪来：摊薄的是"固定开销"，不是"数据量"

三次实验搬运的数据量完全相同（x + y + z = 3 × 51200 B = 153600 B = 150 KiB），唯一变量是一次搬多少元素。tile 从 16 涨到 25600（1600 倍），耗时只降到 1/17.8——因为瓶颈从来不是数据量，而是每次迭代的固定成本被支付了多少遍：

同一份工作量，指令条数差了 1600 倍。这就是 tile=16 → 128 那一档跳得最猛（5.73×）的原因：搬运的建立开销与指令发射开销被摊薄的次数一下子多了 8 倍。

（2）22.7 μs 不是 kernel 的净耗时，是"启动 + 同步"的地板

0.154 MB 的搬运量，哪怕一颗 AI Core 的 MTE 带宽只有 100 GB/s，也只需约 1.5 μs；而实测中位 22.7 μs、最小 14.1 μs。差一个数量级，说明测到的主要不是 kernel 本身：main.cpp 每轮是 launch() → aclrtSynchronizeStream() → 读时间，把 kernel 下发、调度、同步的往返全算了进去。

两处旁证：

• min 与 median 差距很大：tile = 6400 时 min 16.401 / median 25.630（差 56%），tile = 12800 时 min 14.071 / median 24.710（差 76%）。min 剔掉了宿主侧抖动，更接近 kernel 本体。
• 按 min 排序，最优点不在最大 tile：min 口径的加速比是 12800（28.5×）> 6400（24.4×）≈ 25600（24.0×）> 1024（16.9×）。也就是说 tile 超过 12800 之后，kernel 本体反而略微变慢。

这个反转有明确的物理原因：tile = 25600 时内核只迭代 1 次，MTE2（搬入）/ Vector（计算）/ MTE3（搬出）没有任何跨迭代的重叠机会，三段流水完全串行；而 tile = 12800 / 6400 有 2 / 4 次迭代，后一轮的搬入可以和前一轮的计算重叠起来。

另外 tile = 25600 时三个张量占 3 × 51200 B = 150 KiB，UB（192 KiB）已用掉 78%，也留不出双缓冲的空间。——这正是 Step 2a（流水）与 Step 2b（剖析）要展开的问题。

顺带印证：1.1.3 里 tile = 32768 能过，恰好是 3 × 64 KB = 192 KB，把 UB 正好占满。

（3）表里的带宽数字怎么读

0.38 GB/s → 6.76 GB/s 是按 153600 B ÷ median 算的有效带宽（GB/s 按十进制 10⁹ 计），因为分母含固定启动开销，它不代表这颗芯片或这个 kernel 的访存上限。真实的上限要用 msprof 看 MTE 带宽利用率（Step 2b）。

（4）一个被本步骤固定住的变量：只用了 1 个 AI Core

全程 --block-dim 1，即 20 个 AI Core 里只用 1 个。在Trainium2中，Step 1 的标题是 "Parallelize Across 128 Compute Lanes"，讲的是核内 128 条 lane 的并行；

而昇腾（高维切分 API）：

Vector 计算单元每个迭代会从 UB 中取出 8 个 datablock（每个 datablock 数据块内部地址连续，长度 32 Byte）

mask参数: 当前数据类型单次迭代时能处理的元素个数最大值为：256 / sizeof(数据类型)。当操作数占 16 位时，mask ∈ [1, 128]

昇腾核内没有"128 条 lane"这个固定说法：如上面引文所示，宽度是按字节给的（每次迭代 256 B，half 合 128 个元素、float 合 64 个），要处理更多元素靠 repeatTime 迭代。

昇腾另外多出一层核间并行——这是 NKI 单核语义里没有的旋钮。本步骤先固定 block-dim = 1，是为了与 Trainium2 版的单核语义保持可比；核间这一维留到后面（或作为加分项）再考虑。

本节结论：

1. 最小合法分块（tile = 16）跑 25600 个元素：median 403.687 μs（与 --kernel naive 的 404.687 μs 一致，两者本来就是同一个 kernel + 同一个 tile）。
2. 放大到 UB 装得下的最大值：median 22.731 μs，快 17.76×。快的原因是固定开销被摊薄（指令数从 6400 条降到 4 条），不是因为搬得更快。
3. 超过 UB 容量：kernel 会在设备侧执行失败（main.cpp:252 aclError:507035），而不是 host 侧的参数错误 —— aclrtSynchronizeStream 报 "Synchronize stream failed"。

1.2 Step 2a: Improved Data Streaming（4 分）

Trainium版 Step 2a 的做法是把一维向量 reshape 成 (128, FREE_DIM)，让一次 dma_copy 搬更多数据，从而摊薄 DMA 传输的建立开销。昇腾的 DataCopy 本身就是"一次搬 tileLength 个元素"，没有分区维/自由维这套二维语义，于是这一步在昇腾上要拆成两个独立的旋钮分别测量：

一个必须先算清的约束：双缓冲要占两份 UB，6 × tile × 2 B ≤ 196608 B，所以双缓冲下 tile ≤ 16384——正好是单缓冲上限（32768）的一半。

1.2.1 同 tile 下：单缓冲 vs 双缓冲

cd ~/asst4-ascend/part1 

for T in 1280 3200 6400 12800; do echo "== tile=$T 单缓冲"; python3 run_benchmark.py --kernel tiled -n 25600 --tile $T --verify; echo "== tile=$T 双缓冲"; python3 run_benchmark.py --kernel stream -n 25600 --tile $T --verify; done

输出：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ for T in 1280 3200 6400 12800; do echo "== tile=$T 单缓冲"; python3 run_benchmark.py --kernel tiled -n 25600 --tile $T --verify; echo "== tile=$T 双缓冲"; python3 run_benchmark.py --kernel stream -n 25600 --tile $T --verify; done
== tile=1280 单缓冲

kernel      : tiled  (可执行文件里的入口: single)
n           : 25600
tile        : 1280 个元素  (单次搬运 2560 B)
block-dim   : 1 个 AI Core
latency     : median 27.950 us   min 21.840 us   avg 28.856 us
bandwidth   : 5.50 GB/s  (按访存量 153600 B = 0.146 MiB 估算)
[ OK ] 结果与 numpy 参考一致
== tile=1280 双缓冲
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 1280 元素 × 2 B × 2 缓冲 = 15360 B  → 在 GetCoreMemSize 报的可用量之内
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。

kernel      : stream  (可执行文件里的入口: stream)
n           : 25600
tile        : 1280 个元素  (单次搬运 2560 B)
block-dim   : 1 个 AI Core
latency     : median 22.481 us   min 19.500 us   avg 23.222 us
bandwidth   : 6.83 GB/s  (按访存量 153600 B = 0.146 MiB 估算)
[ OK ] 结果与 numpy 参考一致
== tile=3200 单缓冲
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 3200 元素 × 2 B × 1 缓冲 = 19200 B  → 在 GetCoreMemSize 报的可用量之内
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。

kernel      : tiled  (可执行文件里的入口: single)
n           : 25600
tile        : 3200 个元素  (单次搬运 6400 B)
block-dim   : 1 个 AI Core
latency     : median 20.930 us   min 19.070 us   avg 23.313 us
bandwidth   : 7.34 GB/s  (按访存量 153600 B = 0.146 MiB 估算)
[ OK ] 结果与 numpy 参考一致
== tile=3200 双缓冲
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 3200 元素 × 2 B × 2 缓冲 = 38400 B  → 在 GetCoreMemSize 报的可用量之内
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。

kernel      : stream  (可执行文件里的入口: stream)
n           : 25600
tile        : 3200 个元素  (单次搬运 6400 B)
block-dim   : 1 个 AI Core
latency     : median 27.870 us   min 16.040 us   avg 24.734 us
bandwidth   : 5.51 GB/s  (按访存量 153600 B = 0.146 MiB 估算)
[ OK ] 结果与 numpy 参考一致
== tile=6400 单缓冲
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 6400 元素 × 2 B × 1 缓冲 = 38400 B  → 在 GetCoreMemSize 报的可用量之内
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。

kernel      : tiled  (可执行文件里的入口: single)
n           : 25600
tile        : 6400 个元素  (单次搬运 12800 B)
block-dim   : 1 个 AI Core
latency     : median 20.050 us   min 19.210 us   avg 22.898 us
bandwidth   : 7.66 GB/s  (按访存量 153600 B = 0.146 MiB 估算)
[ OK ] 结果与 numpy 参考一致
== tile=6400 双缓冲
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 6400 元素 × 2 B × 2 缓冲 = 76800 B  → 在 GetCoreMemSize 报的可用量之内
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。

kernel      : stream  (可执行文件里的入口: stream)
n           : 25600
tile        : 6400 个元素  (单次搬运 12800 B)
block-dim   : 1 个 AI Core
latency     : median 26.380 us   min 19.451 us   avg 25.060 us
bandwidth   : 5.82 GB/s  (按访存量 153600 B = 0.146 MiB 估算)
[ OK ] 结果与 numpy 参考一致
== tile=12800 单缓冲
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 12800 元素 × 2 B × 1 缓冲 = 76800 B  → 在 GetCoreMemSize 报的可用量之内
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。

kernel      : tiled  (可执行文件里的入口: single)
n           : 25600
tile        : 12800 个元素  (单次搬运 25600 B)
block-dim   : 1 个 AI Core
latency     : median 20.020 us   min 19.790 us   avg 22.166 us
bandwidth   : 7.67 GB/s  (按访存量 153600 B = 0.146 MiB 估算)
[ OK ] 结果与 numpy 参考一致
== tile=12800 双缓冲
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 12800 元素 × 2 B × 2 缓冲 = 153600 B  → 在 GetCoreMemSize 报的可用量之内
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。

kernel      : stream  (可执行文件里的入口: stream)
n           : 25600
tile        : 12800 个元素  (单次搬运 25600 B)
block-dim   : 1 个 AI Core
latency     : median 21.761 us   min 19.520 us   avg 23.508 us
bandwidth   : 7.06 GB/s  (按访存量 153600 B = 0.146 MiB 估算)
[ OK ] 结果与 numpy 参考一致

输出汇总：

1.2.2 双缓冲把上限砍半（验证容量约束）

# 应当通过档
python3 run_benchmark.py --kernel stream -n 32768 --tile 16384 --verify
# 越界档
python3 run_benchmark.py --kernel stream -n 32800 --tile 16400 --verify

输出：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel stream -n 32768 --tile 16384 --verify
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 16384 元素 × 2 B × 2 缓冲 = 196608 B  → 已超过 GetCoreMemSize 报的可用量
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。

kernel      : stream  (可执行文件里的入口: stream)
n           : 32768
tile        : 16384 个元素  (单次搬运 32768 B)
block-dim   : 1 个 AI Core
latency     : median 20.930 us   min 14.510 us   avg 23.051 us
bandwidth   : 9.39 GB/s  (按访存量 196608 B = 0.188 MiB 估算)
[ OK ] 结果与 numpy 参考一致

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel stream -n 32800 --tile 16400 --verify
[WARN ] 三个队列共需 196800 B（3 张量 × 16400 元素 × 2 B × 2 缓冲）> UB 上限 196608 B
         预计会在设备侧失败，报 aclError:507035（Synchronize stream failed）。
         真实上限可用 --ub-size 覆盖；910B/910C 实测为 196608 B。
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 16400 元素 × 2 B × 2 缓冲 = 196800 B  → 已超过 GetCoreMemSize 报的可用量
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。
/opt/atomgit/asst4-ascend/part1/main.cpp:329 aclError:507035
[ERROR] kernel 执行失败（exit 1）

注：越界档必须满足 tile 是 16 的整数倍 且能整除 n。上一版写的 --tile 16392 不是 16 的倍数（16×1024.5），被 harness 的校验拦下——这正是这条规则存在的意义。修正为 16400 = 16 × 1025，配 n = 32800 = 2 × 16400。

6 × 16384 × 2 = 196608 B（正好占满、双缓冲合法上限）通过；6 × 16400 × 2 = 196800 B（超 192 B）失败。

1.2.3 结果分析

1. 一句话结论：median 口径看不出双缓冲的收益，min 口径才看得出来 —— 这是老问题的延续（§1.1.4 的（2））：单次 launch 的"启动 + 同步"地板是 ~19–20 μs，而 25600 个元素的 kernel 本体只有几 μs，host 抖动把中位数淹了。

1. 为什么受益只出现在小 tile：双缓冲的收益 = "MTE2 提前搬下一份、和当前份的计算重叠"。tile 越小 → 迭代次数越多（1280 → 20 次、3200 → 8 次、6400 → 4 次、12800 → 2 次）→ 每次搬入的建立/等待开销占比越大 → 被双缓冲藏掉的越多。等 tile 大到只剩 2–4 次迭代，kernel 本体只剩 1–2 轮、可重叠的量很小，双方都贴着 ~19–20 μs 的地板。

2. min 口径压下去的 3 μs 是什么：tile = 3200 时 min 从 19.07 → 16.04，正好是"每轮都等 DataCopy(i) 搬完才开始计算"被省掉的时间——单缓冲时 Vector 等 MTE2、双缓冲时 MTE2 已在上一轮提前动手。

3. 反思：宿主端墙钟测量回答不了"流水有没有重叠"——median 被启动地板淹没、min 只露出冰山一角。要真正看到 MTE2 / Vector / MTE3 谁在等谁，必须上 msprof 的 timeline（§1.3）。

4. 对应Trainium2版的问题：Trainium2版 Step 2a 问"为什么沿 free 维流式搬更多反而快"（摊薄 DMA 建立开销）；但是昇腾版这里变成"为什么双缓冲在墙钟中位数里看不出快"——同样的手段（增大搬运粒度 / 双缓冲）在昇腾被 ~20 μs 的启动基线掩盖。

5. 1.2.2 的两个边界结果都符合预期：6 × 16384 × 2 = 196608 B 正好占满 → [ OK ]（min 14.510 μs，本组最优——两轮迭代、双缓冲在无 host 抖动时把流水叠满）；6 × 16400 × 2 = 196800 B → 启动前被 [WARN] 预判拦下、随后按预期在设备侧以 aclError:507035 失败。

   顺带第三次印证了"GetCoreMemSize 少报 256 B、但 TPipe 按物理容量分配"（16384 档的 [UB] 账目显示 196608 > 196352 仍通过）。

1.3 Step 2b: Learning to Use msProf（5 分）

Trainium上用 neuron-profile 的 GUI 看 timeline；昇腾侧的对应工具是 msprof op 与 MindStudio Insight（不需要 InfluxDB、不需要 SSH 端口转发，产出就是本地目录）。

cd ~/asst4-ascend/part1
python3 run_benchmark.py --kernel stream -n 25600 --tile 6400 --profile_name s2b_stream_t6400
python3 run_benchmark.py --kernel stream -n 25600 --tile 512  --profile_name s2b_stream_t512
python3 run_benchmark.py --kernel tiled  -n 25600 --tile 6400 --profile_name s2b_tiled_t6400

产出落在 part1/profile/ 下（每组一个目录）。打开方式：

• MindStudio Insight：导入 profile/s2b_stream_t6400*/ 目录看 timeline；
• 或者直接用 msprof 自带的统计输出看各项耗时。

指标口径对照（原版的 neuron-profile 指标 → 昇腾 msprof 指标）见 asst4-ascend/README.md 的 Step 2b 一节，摘录关键的几项：

输出：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel stream -n 25600 --tile 6400 --profile_name s2b_stream_t6400
[INFO ] msprof 采集（第 1 种调用形式）：msprof op --output=/opt/atomgit/asst4-ascend/part1/profile --application=/opt/atomgit/asst4-ascend/part1/profile/s2b_stream_t6400.sh
2026-10-09 19:00:54 [WARN]  The option "--application" will be deprecated in future version. It is recommended to specify <app> [arguments] instead.
2026-10-09 19:00:54 [INFO]  Op profiling analysis start.
2026-10-09 19:00:54 [INFO]  Running device task: DeviceProf
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 6400 元素 × 2 B × 2 缓冲 = 76800 B  → 在 GetCoreMemSize 报的可用量之内
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。
[INFO] <ProfData> Start profiling on device 0, kernel: vector_add_stream_2
[INFO] <WarmUp> Warm Up enabled. times:5
[INFO ] kernel=stream n=25600 tile=6400 blockDim=1
[INFO ] min=628.863 us  median=632.843 us  avg=634.555 us  (0.24 GB/s)
2026-10-09 19:01:00 [INFO]  Profiling running finished. All task success.
2026-10-09 19:01:00 [INFO]  Start analyze kernel on device: device0, kernel: vector_add_stream_2, name is: 0
2026-10-09 19:01:00 [INFO]  Save profiling data with metric [PipeUtilization] success.
2026-10-09 19:01:00 [INFO]  Save profiling data with metric [ArithmeticUtilization] success.
2026-10-09 19:01:00 [INFO]  Save profiling data with metric [L2Cache] success.
2026-10-09 19:01:00 [INFO]  Save profiling data with metric [Memory] success.
2026-10-09 19:01:00 [INFO]  Save profiling data with metric [ResourceConflictRatio] success.
2026-10-09 19:01:00 [INFO]  Save profiling data with metric [MemoryL0] success.
2026-10-09 19:01:00 [INFO]  Save profiling data with metric [MemoryUB] success.
2026-10-09 19:01:02 [INFO]  Performance Summary Report:

        1) MTE2 bandwidth utilization lower than 80% when active.
        2) MTE3 bandwidth utilization lower than 80% when active.
        3) aivector compute usage lower than 20%.

2026-10-09 19:01:02 [INFO]  Operator Basic Information:

        Op Name: vector_add_stream_2
        Op Type: vector
        Task Duration(us): 7.340293
        Block Dim: 1
        Mix Block Dim: 
        Device Id: 0
        Pid: 5239
        Current Freq: 1650
        Rated Freq: 1650

2026-10-09 19:01:02 [INFO]  Profiling results saved in /opt/atomgit/asst4-ascend/part1/profile/OPPROF_20261009190054_AWSXYMXLTSZJJYRP
2026-10-09 19:01:02 [INFO]  Profiling data parse finished.
2026-10-09 19:01:02 [INFO]  Op profiling finish. Welcome to next use.
[INFO ] 本次剖析产出：/opt/atomgit/asst4-ascend/part1/profile/s2b_stream_t6400（msprof 默认命名为 OPPROF_<时间戳>_<随机串>，已重命名）
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 6400 元素 × 2 B × 2 缓冲 = 76800 B  → 在 GetCoreMemSize 报的可用量之内
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。

kernel      : stream  (可执行文件里的入口: stream)
n           : 25600
tile        : 6400 个元素  (单次搬运 12800 B)
block-dim   : 1 个 AI Core
latency     : median 26.750 us   min 24.841 us   avg 28.189 us
bandwidth   : 5.74 GB/s  (按访存量 153600 B = 0.146 MiB 估算)

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel stream -n 25600 --tile 512  --profile_name s2b_stream_t512
[INFO ] msprof 采集（第 1 种调用形式）：msprof op --output=/opt/atomgit/asst4-ascend/part1/profile --application=/opt/atomgit/asst4-ascend/part1/profile/s2b_stream_t512.sh
2026-10-09 19:01:08 [WARN]  The option "--application" will be deprecated in future version. It is recommended to specify <app> [arguments] instead.
2026-10-09 19:01:08 [INFO]  Op profiling analysis start.
2026-10-09 19:01:08 [INFO]  Running device task: DeviceProf
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 512 元素 × 2 B × 2 缓冲 = 6144 B  → 在 GetCoreMemSize 报的可用量之内
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。
[INFO] <ProfData> Start profiling on device 0, kernel: vector_add_stream_2
[INFO] <WarmUp> Warm Up enabled. times:5
[INFO ] kernel=stream n=25600 tile=512 blockDim=1
[INFO ] min=637.123 us  median=641.253 us  avg=643.298 us  (0.24 GB/s)
2026-10-09 19:01:14 [INFO]  Profiling running finished. All task success.
2026-10-09 19:01:14 [INFO]  Start analyze kernel on device: device0, kernel: vector_add_stream_2, name is: 0
2026-10-09 19:01:14 [INFO]  Save profiling data with metric [PipeUtilization] success.
2026-10-09 19:01:14 [INFO]  Save profiling data with metric [ArithmeticUtilization] success.
2026-10-09 19:01:14 [INFO]  Save profiling data with metric [L2Cache] success.
2026-10-09 19:01:14 [INFO]  Save profiling data with metric [Memory] success.
2026-10-09 19:01:14 [INFO]  Save profiling data with metric [ResourceConflictRatio] success.
2026-10-09 19:01:14 [INFO]  Save profiling data with metric [MemoryL0] success.
2026-10-09 19:01:14 [INFO]  Save profiling data with metric [MemoryUB] success.
2026-10-09 19:01:15 [INFO]  Performance Summary Report:

        1) MTE2 bandwidth utilization lower than 80% when active.
        2) MTE3 bandwidth utilization lower than 80% when active.
        3) aivector compute usage lower than 20%.

2026-10-09 19:01:15 [INFO]  Operator Basic Information:

        Op Name: vector_add_stream_2
        Op Type: vector
        Task Duration(us): 15.100604
        Block Dim: 1
        Mix Block Dim: 
        Device Id: 0
        Pid: 5313
        Current Freq: 1650
        Rated Freq: 1650

2026-10-09 19:01:15 [INFO]  Profiling results saved in /opt/atomgit/asst4-ascend/part1/profile/OPPROF_20261009190108_DGQJBUHSUQMXOYWM
2026-10-09 19:01:16 [INFO]  Profiling data parse finished.
2026-10-09 19:01:16 [INFO]  Op profiling finish. Welcome to next use.
[INFO ] 本次剖析产出：/opt/atomgit/asst4-ascend/part1/profile/s2b_stream_t512（msprof 默认命名为 OPPROF_<时间戳>_<随机串>，已重命名）
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 512 元素 × 2 B × 2 缓冲 = 6144 B  → 在 GetCoreMemSize 报的可用量之内
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。

kernel      : stream  (可执行文件里的入口: stream)
n           : 25600
tile        : 512 个元素  (单次搬运 1024 B)
block-dim   : 1 个 AI Core
latency     : median 33.730 us   min 24.310 us   avg 34.637 us
bandwidth   : 4.55 GB/s  (按访存量 153600 B = 0.146 MiB 估算)

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel tiled  -n 25600 --tile 6400 --profile_name s2b_tiled_t6400
[INFO ] msprof 采集（第 1 种调用形式）：msprof op --output=/opt/atomgit/asst4-ascend/part1/profile --application=/opt/atomgit/asst4-ascend/part1/profile/s2b_tiled_t6400.sh
2026-10-09 19:01:21 [WARN]  The option "--application" will be deprecated in future version. It is recommended to specify <app> [arguments] instead.
2026-10-09 19:01:21 [INFO]  Op profiling analysis start.
2026-10-09 19:01:21 [INFO]  Running device task: DeviceProf
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 6400 元素 × 2 B × 1 缓冲 = 38400 B  → 在 GetCoreMemSize 报的可用量之内
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。
[INFO] <ProfData> Start profiling on device 0, kernel: vector_add_single_1
[INFO] <WarmUp> Warm Up enabled. times:5
[INFO ] kernel=single n=25600 tile=6400 blockDim=1
[INFO ] min=627.293 us  median=631.763 us  avg=633.088 us  (0.24 GB/s)
2026-10-09 19:01:26 [INFO]  Profiling running finished. All task success.
2026-10-09 19:01:26 [INFO]  Start analyze kernel on device: device0, kernel: vector_add_single_1, name is: 0
2026-10-09 19:01:26 [INFO]  Save profiling data with metric [PipeUtilization] success.
2026-10-09 19:01:26 [INFO]  Save profiling data with metric [ArithmeticUtilization] success.
2026-10-09 19:01:26 [INFO]  Save profiling data with metric [L2Cache] success.
2026-10-09 19:01:26 [INFO]  Save profiling data with metric [Memory] success.
2026-10-09 19:01:26 [INFO]  Save profiling data with metric [ResourceConflictRatio] success.
2026-10-09 19:01:26 [INFO]  Save profiling data with metric [MemoryL0] success.
2026-10-09 19:01:26 [INFO]  Save profiling data with metric [MemoryUB] success.
2026-10-09 19:01:28 [INFO]  Performance Summary Report:

        1) MTE2 bandwidth utilization lower than 80% when active.
        2) MTE3 bandwidth utilization lower than 80% when active.
        3) aivector compute usage lower than 20%.

2026-10-09 19:01:28 [INFO]  Operator Basic Information:

        Op Name: vector_add_single_1
        Op Type: vector
        Task Duration(us): 8.740350
        Block Dim: 1
        Mix Block Dim: 
        Device Id: 0
        Pid: 5386
        Current Freq: 1650
        Rated Freq: 1650

2026-10-09 19:01:28 [INFO]  Profiling results saved in /opt/atomgit/asst4-ascend/part1/profile/OPPROF_20261009190121_OEONOOMWNBUAYEDI
2026-10-09 19:01:28 [INFO]  Profiling data parse finished.
2026-10-09 19:01:28 [INFO]  Op profiling finish. Welcome to next use.
[INFO ] 本次剖析产出：/opt/atomgit/asst4-ascend/part1/profile/s2b_tiled_t6400（msprof 默认命名为 OPPROF_<时间戳>_<随机串>，已重命名）
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] kernel 本次要分配：3 张量 × 6400 元素 × 2 B × 1 缓冲 = 38400 B  → 在 GetCoreMemSize 报的可用量之内
[UB] 未声明预留（--reserve none）。加 --reserve 8k 可做预留实验。

kernel      : tiled  (可执行文件里的入口: single)
n           : 25600
tile        : 6400 个元素  (单次搬运 12800 B)
block-dim   : 1 个 AI Core
latency     : median 24.950 us   min 18.010 us   avg 25.206 us
bandwidth   : 6.16 GB/s  (按访存量 153600 B = 0.146 MiB 估算)

msprof op 输出的各 CSV 字段速览

表头为本机 CANN 实测；字段语义依据：op_summary（算子详细信息） 里的 Task Duration 定义、l2_cache 命中率文档，以及 MTE/Vector/Scalar 引擎的基准语义。

OpBasicInfo.csv —— 算子基础信息

• Op Name：被采集的算子入口名（本作业：vector_add_stream_2 / matrix_transpose_0）
• Op Type：算子类别（本作业为 vector；矩阵类为 cube）
• Task Duration(us)：该算子一次执行的设备侧总耗时（官方文档：可按其排序找出高耗时算子）
• Block Dim：本次启动的块数（--block-dim）
• Mix Block Dim / Device Id / Pid：混合块维 / 设备号 / 进程号
• Current Freq / Rated Freq：采集时实际核频 / 额定核频（本机 1650/1650 —— 换算 [AICoreMemoryRates] 的 Byte/cycle 用）

PipeUtilization.csv —— 流水（引擎）占用

• block_id / sub_block_id：块号 / 子块号（多核时每个 block 一行）
• aiv_time(us) / aiv_total_cycles：该块 AIV 侧（向量核）总耗时 / 总周期；aiv_* 的所有占比都以它为分母
• aiv_vec_time / aiv_vec_ratio：Vector 计算占用 μs / 占比（转置只有 ~2.5%）
• aiv_scalar_time / aiv_scalar_ratio：标量（指令发射、地址/同步等）占用 / 占比（转置 96.6%→83.5%，全报告的主角）
• aiv_mte2_time /aiv_mte2_ratio / aiv_mte2_active_bw(GB/s)：MTE2（GM→UB 数据搬入）占用 / 占比 / 有效带宽
• aiv_mte3_time / aiv_mte3_ratio / aiv_mte3_active_bw(GB/s)：MTE3（UB→GM 数据搬出）同上
• aic_*：Cube（AIC）侧对应项（纯 Vector 算子为 NA）
• aiv_icache_miss_rate：指令缓存缺失率

ResourceConflictRatio.csv —— 资源/访存冲突占比

• aiv_vec_total_cflt_ratio：Vector 访存总冲突占比
• aiv_vec_bankgroup_cflt_ratio：bank group 冲突（同一迭代 8 块挤进同组）
• aiv_vec_bank_cflt_ratio：严格同 bank 冲突（Part 0 强调的那种）
• aiv_vec_resc_cflt_ratio / aiv_vec_mte_cflt_ratio：资源冲突 / 与 MTE 引擎的冲突
• aiv_vec_wait_ratio：Vector 等待占比（等数据/同步；≠ bank 冲突）
• aiv_mte2_wait_ratio /aiv_mte3_wait_ratio：搬入 / 搬出引擎的等待占比

Memory.csv —— 搬运与主存带宽

• aiv_gm_to_ub_bw / aiv_ub_to_gm_bw(GB/s)：MTE2 搬入 / MTE3 搬出的有效带宽
• aiv_main_mem_read_bw / aiv_main_mem_write_bw(GB/s)：对 HBM 的读写带宽
• aic_l1_read_bw / aic_l1_write_bw(GB/s)：Cube 侧 L1 带宽（纯 Vector 为 NA）
• aiv_mte2_instructions / aiv_mte3_instructions：搬运指令条数（本作业转置 = 1048576 / 1048576）
• read_main_memory_datas / write_main_memory_datas(KB)：实际读写主存总量（转置单核 ≈ 4× 实数据：32 B 粒度 × 128 B 行）
• GM_to_UB_bw_usage_rate / UB_to_GM_bw_usage_rate(%)：搬入 / 搬出带宽利用率（据此可反推 msprof 参考带宽 ≈ 195 GB/s，见 README）

MemoryUB.csv / MemoryL0.csv —— UB 与 L0 读写带宽

• aiv_ub_read_bw_vector / aiv_ub_write_bw_vector(GB/s)：Vector 读写 UB 的带宽（t512 实测 6.6 / 3.3 GB/s）
• aic_l0a/l0b/l0c_*（MemoryL0.csv）：Cube 侧 L0A/L0B/L0C 带宽（纯 Vector 为 NA）

ArithmeticUtilization.csv —— 计算利用率

• aiv_vec_ratio：Vector 计算占比（与 PipeUtilization 的 vec 一致）
• aiv_vec_fp16_ratio / *_int32_ratio 等：各精度指令占比
• aiv_vec_fops：Vector 执行的 fops（指令宽度口径：小 tile 时掩码浪费会被算进去，跨 tile 比较 AI 要小心——§1.3b 的 t16 即此例）
• aic_cube_*：Cube 侧（转置为 NA）

L2Cache.csv —— L2 命中率

• aiv_write_hit_rate /aiv_read_hit_rate /aiv_total_hit_rate(%)：写 / 读 / 总命中率（本作业 ~94–97%）
• aiv_r0/r1_read_cache_hit/_miss_allocate 等：命中与未命中（含上行/归并）计数

主要关心指标

上图把三组各自的四根竖排横条叠在一起看：t512 的"标量"条几乎顶到 Task（95%），而 t6400 的两组四面都短——重叠存在（条长之和远超 Task），但 t512 的重叠全都挤在指令发射与同步的等待里。这也是用 CSV 直接画的"timeline"（无需 MindStudio Insight）。

Insight Details 页的数字，CSV 里已经有了：L2 总命中率 1.73%（每个元素只读一遍、纯流式）；UB 读 / 写带宽 6.6 / 3.3 GB/s；Vector 计算占比 8.6%（fp16 4.2%）、总 fops 32128。Roofline 判定：既不是计算瓶颈（vec 占比 8.6%），也不是内存带宽瓶颈（UB/L2 利用率均为个位数 GB/s），而是指令 / 标量开销瓶颈（scalar 95%）——Insight Details 的热力图结论与此一致。

可以对比platform_config/<SoC>.ini里的数据：

atomgit@6ac4a72a17ae8e63171149c9:~$ grep -inE "bw|bandwidth|rate|pll" $ASCEND_HOME_PATH/aarch64-linux/data/platform_config/Ascend910B4.ini

[AICoreMemoryRates]
ddr_rate=32          ddr_read_rate=32    ddr_write_rate=32    // HBM（读/写同 32 B/cycle）
l2_rate=110          l2_read_rate=110    l2_write_rate=86     // L2
l1_to_l0_a_rate=512  l1_to_l0_b_rate=256 l1_to_ub_rate=128    // L1 → 各计算缓冲
l0_c_to_ub_rate=256  ub_to_l2_rate=64    ub_to_ddr_rate=64    // 回写路径
ub_to_l1_rate=128

换算成 GB/s：Byte/cycle × 实际时钟。核频 1650 MHz 取自 OpBasicInfo.csv 的 Current Freq / Rated Freq 两列（dump/op_basic_info.txt 同值，见 §1.3 三组输出）。以它计：

• HBM ≈ 32 × 1.65 ≈ 53 GB/s、L2 ≈ 110 × 1.65 ≈ 182 GB/s；
• 用 1.8 GHz 计 L2 ≈ 198 GB/s

对比Memory.csv 的 GM_to_UB_bw_usage_rate(%) / UB_to_GM_bw_usage_rate(%) 是 msprof 按某个参考带宽算的利用率。用"实测带宽 ÷ 利用率"可反推参考值：

GM→UB：6.59 GB/s ÷ 3.38% ≈ 195 GB/s
UB→GM：3.29 GB/s ÷ 1.87% ≈ 176 GB/s

~195 GB/s，同量级，基本可以认为 msprof 的参考就是 L2 通路这条 rate。

三问回答

（数据来自 PipeUtilization.csv，aiv_* 是 AIV 侧时间轴上的占用；ratio = 该项 / aiv_time）：

问1：大 tile 与小 tile 的搬运次数与总耗时

总字节不变（153600 B），小 tile 的搬运指令数多 12.5 倍，Task 耗时也大一倍——"次数多"与"慢"是同一个原因的两面。

问2：总结耗时与搬运次数的关系

搬运次数多的时候慢——主要不是搬得慢，是建立与同步开销（scalar）

t512 的 Vector 实际只有 8.6% 占用，标量单元却占了 95.2%——几百条 32 B 级小搬运，每条都要过一遍标量单元的指令发射与 WaitFlag 同步（ResourceConflictRatio.csv 里 aiv_vec_wait_ratio = 0.768，Vector 四分之三的时间在等）。大 tile 把指令数压到 1/12.5，标量占用降到 36%，MTE2/MTE3 有效带宽升到 22.3 / 64.7 GB/s。时间花在"每次搬运的建立+同步"，不在搬字节本身。

Trainium侧 "FREE_DIM = 1000 比 2000 快" 的对照 —— 不矛盾，是同一机制的两个区间：

Trainium FREE_DIM = 1000 比 2000 快的答案（Hint: pipelining）其实写在题面自己的引言里（asst4-trainium2/README.md 的 Step 2a）：

"Too large of a tile size often leads to inefficient pipelining between engines and high memory pressure in SBUF in cases of data reuse"

Trainium 侧并没有昇腾那样的显式"双缓冲"API（没有 InitBuffer(que, num=2) 这类 buffer 份数参数）—— 重叠来自引擎级并发 + 编译器流水：每个 NeuronCore 有 16 个 DMA 引擎可与计算引擎并行，NKI 的 dma_move 是异步的，编译器在循环迭代之间把"下一份的搬运"与"当前份的计算"重叠。而重叠的前提是 SBUF 能同时驻留"正在算的一份"和"正在搬的下一份"（隐式双缓冲 staging）。FREE_DIM=2000 时每个 (128, 2000) tile 占用过大，两份放不下 → memory pressure → 编译器无法重叠 → 搬运与计算串行；FREE_DIM=1000 的两份放得下 → 流水成立 → 尽管 DMA setup 次数更多（每次 setup 有固定开销），总时间反而更短。

昇腾这次的数据同时落到了两个区间：

所以结论统一为一句：搬运次数本身不是性能指标；决定快慢的是"每笔搬运的固定开销能否被流水（重叠）藏住"。Trainium 的 1000-vs-2000 落在"藏得住"一侧（SBUF 容得下两代数据、流水能建立）；昇腾的 512 落在"藏不住、还叠加建立开销"一侧；而昇腾 t1280 档的双缓冲对照，正是原版那个问题在昇腾上的直接体现。

三段重叠的证据（ratio 之和 > 1）

如果三段严格串行，ratio 之和最多 = 1；两组的和都 > 1，证明 MTE2 / MTE3 / Vector 确实在同一段时间内并行。双缓冲与单缓冲（stream vs tiled @ 6400）的差别也在此现形：7.34 vs 8.74 μs（-16%），MTE2 有效带宽 22.3 vs 18.3 GB/s——墙钟中位数看不到（§1.2.3），剖析数据里一目了然。

问3：UB 资源冲突占比（bank 冲突）

三组的 ResourceConflictRatio.csv（aiv_vec_* 为 Vector 访存冲突占比）：

结论：tileLength 与 LocalTensor 分配方式没有引入额外 bank 冲突——最严苛的"同 bank 冲突"（Part 0 说的"8 个 DataBlock 全落同一 bank group"那种）只有 0–0.18%，总冲突 ≤ 2%。原因：我们的 x / y / z 是连续分配的，DataCopy 用 count 形式（连续搬运），Add 用 blkStride = 1（迭代内 8 个 DataBlock 顺序落在交错的 bank 上）——恰好避开 Part 0 例子里"跳读（srcBlkStride = 16）+ 连续写"与"同组写"那两类坑。那 1–2% 的 bankgroup 级冲突来自三个张量同时活跃时偶发的读-读/写-写重叠，量级上可忽略。

顺带澄清：表中 70–77% 的 vec 等待 ≠ bank 冲突——那是 Vector 在等 MTE 搬完数据 / 同步（§1.3 三问 2 里"scalar 95%、vec 只占 8.6%"的另一面），ResourceConflictRatio 里它单独一列，和 bank 冲突是两回事。

1.3b Roofline：tileLength 扫描（补充实验）

用 msprof op 对 9 组 tile（tiled×7 + stream×2）逐一采集，每档提取 Task Duration、aiv_vec_fops、read/write_main_memory_datas，得到 Roofline 图（横轴 AI = FLOP/B，纵轴实现吞吐；内存斜率按 msprof 参考带宽 195 GB/s，计算天花板按 vector 128 lane × 核频（OpBasicInfo.csv 的 Current Freq = 1650 MHz）≈ 211 GFLOP/s）：

之前测算的memory bound = 195 GB/s（内存斜率）:

6.59 GB/s ÷ 3.38% ≈ 195 GB/s   ← msprof 内部参考口径（推断）

再估算一下算数强度：211 GFLOP/s（计算横线）

211 = 128 lane × 1.65 GHz

三个假设叠加：① vector 单元每次迭代 128 条 lane（fp16 的 mask ≤ 128）；② 每 lane 每周期 1 次运算（没算 FMA 双发、没算流水效率）；③ 频率取 OpBasicInfo.csv 的 1650 MHz。

没有任何一个是已测或已公布的数字——这是图上最不确定的一条线。

label   kernel tile      task_us     GFLOP/s     GB/s      AI(FLOP/B)
t16     tiled     16    382.160     1.072      1.617     0.663   ← 见注
t128    tiled    128     50.420      1.018      3.115     0.327
t1024   tiled   1024     14.040      2.060     11.177     0.184
t3200   tiled   3200      9.100      2.940     17.245     0.171
t6400   tiled   6400      7.980      3.288     19.665     0.167
t12800  tiled  12800      7.580      3.428     20.703     0.166
t25600  tiled  25600      7.880      3.281     19.931     0.165
s1280   stream 1280       8.940      3.164     17.553     0.180
s6400   stream 6400       6.680      3.928     23.473     0.167

四条结论：

1. AI 几乎不随 tile 变（t128 以上收敛到 0.165–0.18 FLOP/B），且全部远低于 ridge（211 ÷ 195 ≈ 1.08）→ 全员 memory-bound。AI 是 kernel 的结构属性（fp16 加法，按 profiler 的 fops 口径），tile 改不动它。
2. tile 增大 = 点沿 memory-bound 斜线上移：实现带宽 3.1 → 20.7 GB/s（t12800 最优），GFLOP/s 1.0 → 3.4；耗时 50.4 → 7.6 μs 的地板。这正是"摊薄固定开销"在 Roofline 上的形态：x 不动，y 沿斜线爬。
3. stream（双缓冲）同 tile 高一档：s6400 6.68 μs / 23.47 GB/s vs t6400 7.98 / 19.67——双缓冲把实现带宽再抬 ≈19%，又一次在剖析数据里确认 §1.2.3 的结论。
4. 绝对带宽仍远低于 195 GB/s 参考（最高 23.5 ≈ 12%）：Roofline 展示的不是"带宽打满"，而是"每次搬运的建立/同步开销（scalar）拖住了实现带宽"——单核、footprint 0.15 MB 远小于 L2 的流式小算子就是这种形态。

注：t16 一行的 bytes ≈ 4× 正常值、task 382 μs——profiler 对小 tile 的多次 launch 与"按指令宽度（128 lane）计 fops"叠加了口径偏差，仅作"小 tile 极端"参考，不计入主趋势。

图的两个峰值是 --bw-gbps 195 --peak-gflops 211 的可调参数（910B4 无官方逐级公开值）；峰值改动不影响点的相对位置与"全员 memory-bound"的结论。

1.4 Step 3: Matrix Transpose（15 分 = 10 编码 + 5 写作 (+ 1 EC)）

1.4.1 实现说明

Vector计算单元每个迭代会从UB中取出 8个DataBlock，每个DataBlock数据块内部地址连续，长度 32 Byte，以 half（2B）作为元素类型：一个 DataBlock 装 16 个元素，一次 repeat 处理 128 个。

kernels.cpp 里的 KernelTranspose 是一个基准实现，按 16×16 分块逐块做「搬入 → 转置 → 搬出」：

多核切分：按块列 bj 轮转分给各核（bj = blockIdx; bj < blocksN; bj += blockNum），这样不要求块列数能被核数整除。

1. 搬入：16 行各 32 B，用 连续搬运形式的 DataCopy 逐行搬（首地址 32 B 对齐、长度正好一个 DataBlock）；逐行 32 B 搬运的地址算术最直白、不必先算 gap 与对齐，所以基准版这么写；
2. 转置：AscendC::Transpose(dst, src) —— 该接口的粒度就是 16×16（m、n 都必须是 16 的整数倍，Host 侧已校验）；
3. 搬出：转置结果同样按 16 行各 32 B 写回 out 的对应位置。

分块内的索引映射:

块级上 (行块 bi, 列块 bj) → (行块 bj, 列块 bi)，块内 (r, c) → (c, r)——两次行列互换刚好合成一次转置，*n 与 *m 的差别正是"输入行宽"与"输出行宽"的互换。

设输入块的行、列起点分别为 bi*16、bj*16（bi = 行块索引、bj = 列块索引），块内坐标记为 (r, c) ∈ [0,16)²：

当前实现要求 输入矩阵大小 m、n 都是 16 的整数倍：

• Transpose(dst, src) 官方规定粒度就是 16×16（A2/A3 只支持 16×16 的 uint16_t/half/int16_t 转置）
• 输入行首地址 aGm[(bi*16 + r) * n + bj*16] 元素偏移必须 16 的倍数（= 32 B 对齐）。bj*16 ✓、(bi*16+r)*16 ✓，但 ×(n) 要能被 16 整除 ⇒ n 必须是 16 的倍数
• 输出行首地址 outGm[(bj*16 + c) * m + bi*16]同理 ⇒ m 必须是 16 的倍数。

如果要支持任意尺寸：

DataCopy 函数入参:

1. 连续搬运（最常用）

// GM -> UB
template <typename T> __aicore__ inline void
DataCopy(const LocalTensor<T>&  dst, const GlobalTensor<T>& src, const uint32_t count);

// UB -> GM
template <typename T> __aicore__ inline void
DataCopy(const GlobalTensor<T>& dst, const LocalTensor<T>&  src, const uint32_t count);

// Local -> Local
template <typename T> __aicore__ inline void
DataCopy(const LocalTensor<T>&  dst, const LocalTensor<T>&  src, const uint32_t count);

1. 非连续搬运（带 DataCopyParams）

template <typename T> __aicore__ inline void
DataCopy(const LocalTensor<T>&  dst, const GlobalTensor<T>& src, const DataCopyParams& repeatParams);

template <typename T> __aicore__ inline void
DataCopy(const GlobalTensor<T>& dst, const LocalTensor<T>&  src, const DataCopyParams& repeatParams);

template <typename T> __aicore__ inline void
DataCopy(const LocalTensor<T>&  dst, const LocalTensor<T>&  src, const DataCopyParams& repeatParams);

「一次搬多行」正是 EC 的优化方向（见 §1.4.4 第 1 条，可把每块 32 次调用降到 2 次）。

DataCopyParams 结构体:

struct DataCopyParams {
    uint16_t blockCount = 1;                    // DEFAULT_DATA_COPY_NBURST
    uint16_t blockLen   = 0;
    union { uint16_t srcGap = 0; uint16_t srcStride; };   // 同一内存，两个名字
    union { uint16_t dstGap = 0; uint16_t dstStride; };
};

示例场景：GM 上 [128, 256] 的 half 矩阵，取前 64 行、每行 128 个元素，搬到 UB 连续存放。

AscendC::DataCopyParams params;
params.blockCount = 64;                            // 搬 64 行 → 64 轮
params.blockLen   = 128 * sizeof(half) / 32;       // 每行 256B = 8 个 DB
params.srcGap     = (256 - 128) * sizeof(half)/32; // 源行末要跳过 128 个 half = 8 个 DB
params.dstGap     = 0;                             // UB 里连续存放，无缝
AscendC::DataCopy(ubLocal, srcGm[0], params);

src (GM, 每行 256 half = 16 DB):
行0: [取前8个DB | 跳过8个DB]
行1: [取前8个DB | 跳过8个DB]     ← blockCount=64 轮，每轮跨 8+8=16 个 DB
...

dst (UB):
[行0前128个][行1前128个]...  ← dstGap=0，紧密排列，共 64×8 = 512 个 DB = 16 KB

反过来，把 UB 的连续数据写回大矩阵的某个子块，就是把 srcGap=0（UB 连续）、dstGap=行间隔（GM 跳着写）。

注：CANN 8.5 和 CANN 9.0 的 DataCopyParams 是同一个结构体，字段没变。 看到的命名差异（srcStride/dstStride vs srcGap/dstGap）不是版本差异，是新旧命名并存——头文件里它们是一个 union 的两个别名：

// 源码：${INSTALL_DIR}/.../kernel_struct_data_copy.h
struct DataCopyParams {
    __aicore__ DataCopyParams() {}
    __aicore__ DataCopyParams(const uint16_t count, const uint16_t len,
                              const uint16_t srcStrideIn, const uint16_t dstStrideIn)
        : blockCount(count), blockLen(len), srcStride(srcStrideIn), dstStride(dstStrideIn) {}
    uint16_t blockCount = DEFAULT_DATA_COPY_NBURST;   // 默认 1
    uint16_t blockLen = 0;
    // srcStride and dstStride will be deprecated, use srcGap and dstGap instead.
    union { uint16_t srcGap = 0;  uint16_t srcStride; };
    union { uint16_t dstGap = 0;  uint16_t dstStride; };
};

所以 params.srcGap = 4 和 params.srcStride = 4 写入的是同一块内存，完全等价。

官方从 9.0 起在文档上统一推 srcGap/dstGap，并把 srcStride/dstStride 标记为 "will be deprecated"。

1. 增强搬运（8.5 起 / 9.0 保留并扩充）

template <typename T, typename U> __aicore__ inline void
DataCopy(const LocalTensor<T>& dst, const LocalTensor<U>& src,
         const DataCopyParams&         intriParams,
         const DataCopyEnhancedParams& enhancedParams);

地址对齐约束：

1.4.2 运行与耗时

目标：扫 --block-dim 找出"最快"档。转置是 Vector（AIV）+ MTE + 标量 通路（不碰 Cube）；本机物理单元为 Cube × 20 / Vector × 40（platform_config core_type_list=CubeCore,VectorCore，探针 AIC = 20 / AIV = 40）。--block-dim 是逻辑核（[1, 65535]），于是要实测两个问题：① 加速曲线在哪一档停止下降（20 = 按 Cube 计 / 40 = 按 Vector 计）？② 超过物理单元后是否回升？

cd ~/asst4-ascend/part1

# 0) 首次跑建议 --clean 一次（kernels.cpp 改过，避免旧 .o）
python3 run_benchmark.py --kernel transpose -n 1024 --verify --clean

# 1) 扫 block-dim：1 → 128，每档都 --verify（正确性把关）
for B in 1 2 4 8 16 20 24 32 40 48 64 96 128; do echo "===== block-dim=$B"; python3 run_benchmark.py --kernel transpose -n 1024 --block-dim $B --verify; done

# 2) 大矩阵（4096² = 64 MB 访存）挑重点档复核（太重，不必全扫）
for B in 1 4 8 20 40 64 128; do echo "===== block-dim=$B"; python3 run_benchmark.py --kernel transpose -n 4096 --block-dim $B --verify; done

用双NPU的A3的环境运行（每个NPU 40 个 AIV)：

developer@8de626ef107b43269a92e55bdfcaea3c:.../workspace/asst4-ascend/part1$ npu-smi info
+------------------------------------------------------------------------------------------------+
| npu-smi 25.5.5                   Version: 25.5.5                                               |
+---------------------------+---------------+----------------------------------------------------+
| NPU   Name                | Health        | Power(W)    Temp(C)           Hugepages-Usage(page)|
| Chip  Phy-ID              | Bus-Id        | AICore(%)   Memory-Usage(MB)  HBM-Usage(MB)        |
+===========================+===============+====================================================+
| 7     Ascend910           | OK            | 169.2       47                0    / 0             |
| 0     14                  | 0000:0A:00.0  | 0           0    / 0          3108 / 65536         |
+------------------------------------------------------------------------------------------------+
| 7     Ascend910           | OK            | -           46                0    / 0             |
| 1     15                  | 0000:0B:00.0  | 0           0    / 0          2870 / 65536         |
+===========================+===============+====================================================+
+---------------------------+---------------+----------------------------------------------------+
| NPU     Chip              | Process id    | Process name             | Process memory(MB)      |
+===========================+===============+====================================================+
| No running processes found in NPU 7                                                            |
+===========================+===============+====================================================+

输出：

developer@8de626ef107b43269a92e55bdfcaea3c:.../workspace/asst4-ascend/part1$ python3 run_benchmark.py --kernel transpose -n 1024 --verify --clean
[INFO ] --clean：删除 build-npu/ 做一次干净构建
[INFO ] 构建中（soc-version=Ascend910_9362）...
...
INFO ] 构建完成：/mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 1 个 AI Core
latency     : median 1347.092 us   min 1342.963 us   avg 1346.540 us
bandwidth   : 3.11 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致


developer@8de626ef107b43269a92e55bdfcaea3c:.../workspace/asst4-ascend/part1$ for B in 1 2 4 8 16 20 24 32 40 48 64 96 128; do echo "===== block-dim=$B"; python3 run_benchmark.py --kernel transpose -n 1024 --block-dim $B --verify; done
===== block-dim=1
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 1 个 AI Core
latency     : median 1346.834 us   min 1342.733 us   avg 1346.372 us
bandwidth   : 3.11 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=2
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 2 个 AI Core
latency     : median 683.997 us   min 680.037 us   avg 686.389 us
bandwidth   : 6.13 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=4
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 4 个 AI Core
latency     : median 354.535 us   min 351.734 us   avg 354.387 us
bandwidth   : 11.83 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=8
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 8 个 AI Core
latency     : median 185.677 us   min 184.847 us   avg 185.739 us
bandwidth   : 22.59 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=16
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 16 个 AI Core
latency     : median 105.414 us   min 104.134 us   avg 105.420 us
bandwidth   : 39.79 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=20
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 20 个 AI Core
latency     : median 103.434 us   min 101.994 us   avg 103.957 us
bandwidth   : 40.55 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=24
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 24 个 AI Core
latency     : median 84.333 us   min 83.523 us   avg 84.659 us
bandwidth   : 49.74 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=32
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 32 个 AI Core
latency     : median 65.563 us   min 64.432 us   avg 65.975 us
bandwidth   : 63.97 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=40
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 40 个 AI Core
latency     : median 65.172 us   min 63.702 us   avg 75.830 us
bandwidth   : 64.36 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=48
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 48 个 AI Core
latency     : median 84.983 us   min 83.934 us   avg 85.391 us
bandwidth   : 49.35 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=64
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 64 个 AI Core
latency     : median 66.432 us   min 65.573 us   avg 66.731 us
bandwidth   : 63.14 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=96
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 96 个 AI Core
latency     : median 67.303 us   min 66.263 us   avg 68.080 us
bandwidth   : 62.32 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=128
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 1024
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 128 个 AI Core
latency     : median 67.643 us   min 66.732 us   avg 69.999 us
bandwidth   : 62.01 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致

developer@8de626ef107b43269a92e55bdfcaea3c:.../workspace/asst4-ascend/part1$ for B in 1 4 8 20 40 64 128; do echo "===== block-dim=$B"; python3 run_benchmark.py --kernel transpose -n 4096 --block-dim $B --verify; done
===== block-dim=1
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 4096
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 1 个 AI Core
latency     : median 21154.948 us   min 21111.956 us   avg 21152.699 us
bandwidth   : 3.17 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=4
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 4096
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 4 个 AI Core
latency     : median 5307.902 us   min 5305.062 us   avg 5309.387 us
bandwidth   : 12.64 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=8
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 4096
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 8 个 AI Core
latency     : median 2666.387 us   min 2663.316 us   avg 2666.403 us
bandwidth   : 25.17 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=20
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 4096
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 20 个 AI Core
latency     : median 1097.654 us   min 1094.864 us   avg 1097.813 us
bandwidth   : 61.14 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=40
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 4096
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 40 个 AI Core
latency     : median 600.884 us   min 598.874 us   avg 601.871 us
bandwidth   : 111.68 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=64
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 4096
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 64 个 AI Core
latency     : median 684.777 us   min 682.977 us   avg 685.214 us
bandwidth   : 98.00 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致
===== block-dim=128
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 4096
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 128 个 AI Core
latency     : median 686.188 us   min 682.827 us   avg 686.029 us
bandwidth   : 97.80 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致

汇总（1024²；median_us）：

大矩阵（4096²；median_us）：

实测结论：最优 --block-dim = 40，膝盖点对应的是 Vector（AIV = 40）而不是 Cube（AIC = 20）。

1. 4096² 达到 35.2× 是决定性证据：如果并行执行单元只有 20 个（按 Cube 计），加速比不可能超过 20×。实测 35× → 该 kernel 实际并行在 40 个 Vector 单元上（转置不碰 Cube；core_type_list=CubeCore,VectorCore、探针 AIC = 20 / AIV = 40）——修正了"最多 20"的直觉。
2. 超过 40 后持平或略回：64 / 96 / 128 时 1024² 停在 ~66–68 μs、4096² 停在 ~685 μs（比 40 的 601 还差一点）——多出的 block 与其他 block 共享物理 Vector 单元，不再加速；块列数不够分时多出的 block 直接空转退出，也不添乱。
3. 负载均衡注释（1024² 的 48 档异常峰 85.0 μs）：1024² 只有 64 个块列，48 核时 64 = 48×1 + 16，16 个核要扛 2 列 → 不均衡；而 40 / 64 核（64 列 ≈ 1.6 / 1 列）回到 65–66。4096² = 256 列，40 核 6.4 列/核、轮转分配基本均衡 → 近线性（1→4→8→20→40 的加速比 1 / 3.99 / 7.94 / 19.3 / 35.2）。块列数远大于核数时，轮转切分没有均衡性问题。
4. 绝对值仍低：4096² 单核 21 ms → 40 核 0.6 ms（带宽 ~112 GB/s，仍远低于参考），因为每个 16×16 块做 32 次 32 B 级 DataCopy——基准实现的"碎搬迁"开销，EC 优化方向不变（带步长 DataCopy / TransDataTo5HD）。

1.4.3 剖析与"是 memory-bound 还是 compute-bound"

取同一规模（4096²）、单核 vs 最优 40 核各剖析一次——对比引擎占比的变化，才能讲清楚"并行买到了什么"：

cd ~/asst4-ascend/part1
python3 run_benchmark.py --kernel transpose -n 4096 --profile_name s3_tpose_4096_bd1
python3 run_benchmark.py --kernel transpose -n 4096 --block-dim 40 --profile_name s3_tpose_4096_bd40

如果需要Timeline的话，仿真流水图指令：

cd ~/asst4-ascend/part1

# 单核版（用 512² 看形态；要更小就用 --n 256）
msprof op simulator --soc-version=Ascend910_9362 --output=./sim_tpose_bd1 \
    ./build-npu/ascendc_kernels --kernel transpose -n 512 -m 512 --block-dim 1 --iters 1 --warmup 0 --reserve none

# 40 核版
msprof op simulator --soc-version=Ascend910_9362 --output=./sim_tpose_bd40 \
    ./build-npu/ascendc_kernels --kernel transpose -n 512 -m 512 --block-dim 40 --iters 1 --warmup 0 --reserve none

注意：上板 msprof op 不产生指令流水的 timeline——timeline（指令级 trace）只有仿真 msprof op simulator 才出 trace.json。所以"给这两个跑 timeline"要用仿真，而且 仿真 4096² 的 210 万条指令会跑到天荒地老，必须用小矩阵看形态

输出：

developer@8de626ef107b43269a92e55bdfcaea3c:.../workspace/asst4-ascend/part1$ python3 run_benchmark.py --kernel transpose -n 4096 --profile_name s3_tpose_4096_bd1
[INFO ] msprof 采集（第 1 种调用形式）：msprof op --output=/mnt/workspace/asst4-ascend/part1/profile /mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels --kernel transpose --n 4096 --m 4096 --tile 16 --block-dim 1 --warmup 5 --iters 20 --reserve none --out /mnt/workspace/asst4-ascend/part1/out/transpose.bin
2026-10-10 01:06:15 [INFO]  Op profiling analysis start.
2026-10-10 01:06:15 [INFO]  Running device task: DeviceProf
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）
[INFO] <ProfData> Start profiling on device 0, kernel: matrix_transpose_0
[INFO] <WarmUp> Warm Up enabled. times:5
[INFO ] kernel=transpose n=4096 tile=16 blockDim=1
[INFO ] min=21410.747 us  median=21444.918 us  avg=21448.480 us  (3.13 GB/s)
2026-10-10 01:06:20 [INFO]  Profiling running finished. All task success.
2026-10-10 01:06:20 [INFO]  Start analyze kernel on device: device0, kernel: matrix_transpose_0, name is: 0
2026-10-10 01:06:21 [INFO]  Save profiling data with metric [PipeUtilization] success.
2026-10-10 01:06:21 [INFO]  Save profiling data with metric [ArithmeticUtilization] success.
2026-10-10 01:06:21 [INFO]  Save profiling data with metric [L2Cache] success.
2026-10-10 01:06:21 [INFO]  Save profiling data with metric [Memory] success.
2026-10-10 01:06:21 [INFO]  Save profiling data with metric [ResourceConflictRatio] success.
2026-10-10 01:06:21 [INFO]  Save profiling data with metric [MemoryL0] success.
2026-10-10 01:06:21 [INFO]  Save profiling data with metric [MemoryUB] success.
2026-10-10 01:06:21 [INFO]  Performance Summary Report:

        1) MTE2 bandwidth utilization lower than 80% when active.
        2) MTE3 bandwidth utilization lower than 80% when active.
        3) aivector compute usage lower than 20%.

2026-10-10 01:06:21 [INFO]  Operator Basic Information:

        Op Name: matrix_transpose_0
        Op Type: vector
        Task Duration(us): 21545.347656
        Block Dim: 1
        Mix Block Dim: 
        Device Id: 0
        Pid: 10803
        Current Freq: 1650
        Rated Freq: 1650

2026-10-10 01:06:21 [INFO]  Profiling results saved in /mnt/workspace/asst4-ascend/part1/profile/OPPROF_20261010010615_QWUWKTPYVLAVCBRS
2026-10-10 01:06:21 [INFO]  Profiling data parse finished.
2026-10-10 01:06:21 [INFO]  Op profiling finish. Welcome to next use.
[INFO ] 本次剖析产出：/mnt/workspace/asst4-ascend/part1/profile/s3_tpose_4096_bd1（msprof 默认命名为 OPPROF_<时间戳>_<随机串>，已重命名）
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 4096
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 1 个 AI Core
latency     : median 21153.957 us   min 21111.005 us   avg 21144.814 us
bandwidth   : 3.17 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)

developer@8de626ef107b43269a92e55bdfcaea3c:.../workspace/asst4-ascend/part1$ python3 run_benchmark.py --kernel transpose -n 4096 --block-dim 40 --profile_name s3_tpose_4096_bd40
[INFO ] msprof 采集（第 1 种调用形式）：msprof op --output=/mnt/workspace/asst4-ascend/part1/profile /mnt/workspace/asst4-ascend/part1/build-npu/ascendc_kernels --kernel transpose --n 4096 --m 4096 --tile 16 --block-dim 40 --warmup 5 --iters 20 --reserve none --out /mnt/workspace/asst4-ascend/part1/out/transpose.bin
2026-10-10 01:06:52 [INFO]  Op profiling analysis start.
2026-10-10 01:06:52 [INFO]  Running device task: DeviceProf
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）
[INFO] <ProfData> Start profiling on device 0, kernel: matrix_transpose_0
[INFO] <WarmUp> Warm Up enabled. times:5
[INFO ] kernel=transpose n=4096 tile=16 blockDim=40
[INFO ] min=872.885 us  median=879.095 us  avg=881.531 us  (76.34 GB/s)
2026-10-10 01:06:57 [INFO]  Profiling running finished. All task success.
2026-10-10 01:06:57 [INFO]  Start analyze kernel on device: device0, kernel: matrix_transpose_0, name is: 0
2026-10-10 01:06:57 [INFO]  Save profiling data with metric [PipeUtilization] success.
2026-10-10 01:06:57 [INFO]  Save profiling data with metric [ArithmeticUtilization] success.
2026-10-10 01:06:57 [INFO]  Save profiling data with metric [L2Cache] success.
2026-10-10 01:06:57 [INFO]  Save profiling data with metric [Memory] success.
2026-10-10 01:06:57 [INFO]  Save profiling data with metric [ResourceConflictRatio] success.
2026-10-10 01:06:57 [INFO]  Save profiling data with metric [MemoryL0] success.
2026-10-10 01:06:57 [INFO]  Save profiling data with metric [MemoryUB] success.
2026-10-10 01:06:57 [INFO]  Performance Summary Report:

        1) MTE2 bandwidth utilization lower than 80% when active.
        2) MTE3 bandwidth utilization lower than 80% when active.
        3) aivector compute usage lower than 20%.

2026-10-10 01:06:57 [INFO]  Operator Basic Information:

        Op Name: matrix_transpose_0
        Op Type: vector
        Task Duration(us): 690.526184
        Block Dim: 40
        Mix Block Dim: 
        Device Id: 0
        Pid: 10981
        Current Freq: 1650
        Rated Freq: 1650

2026-10-10 01:06:57 [INFO]  Profiling results saved in /mnt/workspace/asst4-ascend/part1/profile/OPPROF_20261010010652_SQIUEXZAKLXMUFGU
2026-10-10 01:06:57 [INFO]  Profiling data parse finished.
2026-10-10 01:06:57 [INFO]  Op profiling finish. Welcome to next use.
[INFO ] 本次剖析产出：/mnt/workspace/asst4-ascend/part1/profile/s3_tpose_4096_bd40（msprof 默认命名为 OPPROF_<时间戳>_<随机串>，已重命名）
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 4096
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 40 个 AI Core
latency     : median 603.894 us   min 600.844 us   avg 603.666 us
bandwidth   : 111.13 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)

汇总（从各自目录的 OpBasicInfo.csv / PipeUtilization.csv 抄）：

搬运指令数计算：

块数 = (4096/16)² = 65536；每块逐行搬 16 行 → 单向 65536×16 = 1048576，搬入/搬出各一份；Transpose 每块 1 次 = 65536 次。

均有实测佐证：Memory.csv 的 aiv_mte2_instructions = 1048576、aiv_mte3_instructions = 1048576，反推块数 1048576÷16 = 65536，与 (4096/16)² 两条独立路径吻合

一眼就能看出：单核时"标量"条（20.8 ms）几乎顶满 Task（21.5 ms）；40 核后同样这四根条一起缩到亚毫秒，标量仍是四根里最长（572 μs / 690 μs）——多核把 210 万条指令的发射与同步摊进了 40 个 Vector 单元，但"每次搬运的固定开销"结构没变。

判断路径（先用推理、再用 msprof 验证）：

1. 不是 compute-bound：转置的算术强度 ≈ 每元素 0 次浮点（aiv_vec 应接近 0 —— 唯一"算"是搬移本身），跟 §1.3b Roofline 的结论一致（AI ≪ ridge）；
2. 是不是 bandwidth-bound 要看：单核版 aiv_mte2/mte3 占比与有效带宽（预期很低，如 ~3 GB/s）→ 不是被带宽压着，而是被搬运指令数 + 标量发射/同步压着（每 16×16 块 32 次 32 B 级 DataCopy，4096² 共 65536 块 × 32 = 2×1048576 ≈ 210 万条——profiler 实测 aiv_mte2_instructions = 1048576、aiv_mte3_instructions = 1048576）；
3. 40 核那一档的对比价值：并行把楼上那坨指令摊到 40 个 Vector 单元后，aiv_mte2 有效带宽应大幅上升（预期 ~110 GB/s 量级，与 wall 带宽一致）、标量占比相对下降——这组对比直接回答"多核买到了什么"：买的是搬运指令的并行执行，而不是更高单核带宽。

实测结论（msprof 证据就是上面两列数字）：

1. 不是 compute-bound：两侧 aiv_vec 都只有 ~2.5%（转置 ≈ 0 FLOP/元素）——msprof 的自动摘要也给出同样的判定（aivector compute usage lower than 20%）；
2. 单核不是带宽-bound：MTE2/3 有效带宽只有 9.6 / 10.4 GB/s、usage_rate 2.6–3.1%（lower than 80%）——真正压死时间的是 210 万条 32 B 级 DataCopy 的指令发射与同步（scalar 96.6%，aiv_time 几乎被标量占满）；
3. 40 核把指令流平分给 40 个 Vector 单元：21545 → 690 μs（31×，与 §1.4.2 的 35.2× 同量级）。比例变成 mte2 74.3% / scalar 83.5%——并行买的是"搬运指令的并行执行"，不是"单笔搬运变快"：每核有效带宽仍只有 7–12 GB/s，wall 带宽 111 GB/s 是 40 核叠加的结果；
4. L2 命中 ≈ 94–97%：32 B 粒度 × 128 B 行 → 同一行被 4 次读取复用（顺带解释了 read datas ≈ 4× 实数据）；复用率高反而说明"不是没搬完，而是每次搬得太碎"；
5. 推论（接 §1.4.4）：scalar（指令条数）才是瓶颈 → 优化方向 = 把每条 DataCopy 搬大（带步长一次搬 16 行、或 TransDataTo5HD），指令数降 64×，直接攻击 83–96% 这个标量占比；这比调 mask / 换算法更对症。

1.4.4 额外加分：优化方向

昇腾上没有原版那样的现成阈值（原文是 4096×4096 < 700 μs），所以要自建基线：

• 基线 1：torch.transpose(x, 0, 1).contiguous() 在 NPU 上的耗时（CANN 自带实现）；
• 基线 2：本节的朴素分块实现（1.4.2 的数据）。

达标建议：打败基线 1，或达到其 1.2 倍以内。可用的优化方向（按性价比排序）：

1. 换带步长的搬运：用 DataCopyParams 一次搬 16 行（把每块的 32 次调用降到 2 次）。搬入侧 {blockCount=16, blockLen=1, srcGap=n/16-1, dstGap=0}、搬出侧 {16, 1, 0, m/16-1}（blockLen 与 gap 单位都是 DataBlock = 32 B；gap 是「前块尾→后块头」的间隔，等于行距减块长）。写 srcStride/dstStride 与写 srcGap/dstGap 等价（同一 union），随 9.0 文档用新名即可；
2. 换 TransDataTo5HD：官方明说"Transpose 仅支持 16×16 大小的矩阵"，更大的块要用它；
3. 消除物理转置（最有"平台味道"）：如果转置只是后续算子的输入，直接在下一跳用 Matmul 的转置参数或换内存布局，比真搬一遍快得多——对应 NKI 里的 lhsT 约定；
4. 多核并行：--block-dim 调大（当前基准版已支持按块列切核）。

EC 基线：CANN 自带实现（torch.transpose(...).contiguous()）

原版的 EC 要求在 4096×4096 上打到 <700 μs；昇腾侧没有现成阈值，先用 CANN 自带的 torch 实现建基线（脚本 part1/torch_transpose_baseline.py，计时口径与本节一致：warmup 5 + iters 20，每次前后 npu.synchronize()，取 median/min/avg）：

cd ~/asst4-ascend/part1
python3 torch_transpose_baseline.py --n 4096 2>&1 | grep -v "path string is NULL"        # 默认 fp16，与我们的 kernel 对齐
python3 torch_transpose_baseline.py --n 1024 2>&1 | grep -v "path string is NULL"         # 小规模对照
# （可选）python3 torch_transpose_baseline.py --n 4096 --fp32   # 好奇 fp32 差多少

脚本同时会打印一个对照组：纯 x.transpose(0,1)（只切视图、不搬数据）的耗时——用来区分"搬运开销"与"算子调度/下发开销"。

可比性说明：torch 基线没有 --tile / --block-dim 这类旋钮（脚本只接受 --n/--iters/--warmup/--fp32；传 --tile 会直接报 unrecognized arguments）。它的 tiling 与用核数由 CANN 算子内部决定（默认就是用满设备）。所以我们这边的 --block-dim 40 才是与它同口径的配置；单核（--block-dim 1）的数字不能拿去和 torch 比。--tile 是我们自己的分块参数——比的是"我们手动调优" vs "CANN 内部调优"。

输出：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 torch_transpose_baseline.py --n 1024 2>&1 | grep -v "path string is NULL"
device       : Ascend910B4   (torch.npu.is_available() = True)
baseline: x.transpose(0,1).contiguous()   shape=(1024, 1024)  dtype=torch.float16  device=npu:0
result       : device=npu:0  contiguous=True  shape=(1024, 1024)   ← 证明数据搬运确实发生在 NPU 上
latency     : median     90.700 us   min    84.933 us   avg    99.280 us   (20 iters)
bandwidth   :   46.244 GB/s  (按访存量 4.2 MB 估算)
纯 transpose（仅视图，不搬数据）:    6.165 us —— 对照用：contiguous() 才是真实搬运开销

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 torch_transpose_baseline.py --n 4096 2>&1 | grep -v "path string is NULL" 
device       : Ascend910B4   (torch.npu.is_available() = True)
baseline: x.transpose(0,1).contiguous()   shape=(4096, 4096)  dtype=torch.float16  device=npu:0
result       : device=npu:0  contiguous=True  shape=(4096, 4096)   ← 证明数据搬运确实发生在 NPU 上
latency     : median    177.883 us   min   164.043 us   avg   178.136 us   (20 iters)
bandwidth   :  377.265 GB/s  (按访存量 67.1 MB 估算)
纯 transpose（仅视图，不搬数据）:    6.138 us —— 对照用：contiguous() 才是真实搬运开销

总结结果：

scalar 瓶颈解读

§1.4.3 结论是「scalar（标量单元）占 96.6% / 83.5%，是真正的瓶颈」。 aiv_scalar 是标量单元在 AIV 时间轴上的占用，涵盖：

但实测表明它不是"发射吞吐打满"，而是"每条搬运指令太贵"：

也就是说：不是"指令发得太慢"，而是"每条小搬运的固定代价太高"（地址/描述符构造 + 与 MTE 的同步握手）。所以优化杠杆只有一个方向——让每条搬运覆盖更多字节，而不是去抠发射周期。

TransDataTo5HD 不是搬运，是数据排布转换，特别的，也可以用于二维矩阵数据块的转置。参考 TransDataTo5HD 功能说明。

一次 repeat 内可处理 512 Byte（16 个 DataBlock）；指令内部循环 16 次，每次从 16 个指定 DataBlock 的对应位置取值，拼成一个新的 DataBlock 写到目的地址。

对于half，下面图例里每行就是一个DataBlock，正好16个元素。一次取16个DataBlock做转置：

TransDataTo5HDParams 五个字段:

struct TransDataTo5HDParams {
    bool     dstHighHalf   = false;
    bool     srcHighHalf   = false;
    uint8_t  repeatTimes   = 1;
    uint16_t dstRepStride  = 0;
    uint16_t srcRepStride  = 0;
};

• 当 repeatTimes > 1 时，第 1 次 repeat 的有效起始位置 = 地址序列起始位置，第 2 次起每次加上 dstRepStride/srcRepStride；
• 当 repeatTimes = 1 时，有效起始位置 = 序列起始位置 + repStride —— 所以想让起始位置就是序列起点，必须把 repStride 置 0。

TransDataTo5HD 是 UB 内的搬运（dstList/srcList 的 TPosition 只允许 VECIN/VECCALC/VECOUT），数据必须驻留 UB。官方约束：

• 单次 repeat 处理 512 B（16 个 DataBlock），repeatTimes ≤ 255 → 单次调用最多 127.5 KiB 输入；
• 地址列表固定 16 项（一维恒为 16）、起始地址 32 B 对齐、src/dst 需独立缓冲（未见可原地复用声明）。

于是占用 UB 天花板：单缓冲 2S ≤ 192 KiB → S ≤ 96 KiB；src/dst 各双缓冲 4S ≤ 192 KiB → S ≤ 48 KiB。

本作业按 16×16（512 B）或 16×64（2 KiB）分块，src+dst 仅 1–4 KiB，离 UB 上限三个数量级——这条路的真正瓶颈仍是搬运指令数。

优化后的预期路径（用于 EC 定标）：①+② 把单核指令数压到 ~1/64，标量不再是瓶颈 → 新的地板变成 MTE 搬运带宽（单核实测有效带宽 ~10 GB/s，64 MiB 往返 → 约 3–6 ms）；再乘上 ⑥ 的多核并行（40×），量级可到 ~100–200 μs。届时再用 msprof 复核：aiv_scalar_ratio 应显著下降、aiv_mte2/3_active_bw 上升——即"瓶颈从标量转移到带宽"。

反面提醒：不要指望通过编译选项/循环展开来救——aiv_icache_miss_rate ≈ 0 已排除取指瓶颈，抠发射周期解决不了"每条搬运 ~10 周期"的固定代价；只有"每条搬更多字节"才动得了它。

优化算法实现说明

1. 带步长 DataCopyParams：一个 tile 搬入/搬出各 1 条指令。

// 搬入：16 行 × K 个 DataBlock，行间距 (n/16 − K) 个 DataBlock
inParams  = {blockCount=16,        blockLen=K, srcStride=n/16-K, dstStride=0};
// 搬出：TILE_COLS 行 × 1 个 DataBlock（dst 在 UB 里连续）
outParams = {blockCount=TILE_COLS, blockLen=1, srcStride=0,     dstStride=m/16-1};

（注释里点明了 stride 的语义是"间隔"而非"间距"、以及 srcStride/dstStride 与 srcGap/dstGap 是同一 union。）

1. TransDataTo5HD 取代 Transpose

for (g = 0; g < K; ++g) {                       // 每个 16 列组一次
    for (i = 0; i < 16; ++i) {
        srcList[i] = src[i*TILE_COLS + g*16].GetPhyAddr();  // 第 i 行、第 g 组的起点
        dstList[i] = dst[(g*16 + i)*16].GetPhyAddr();       // 目的第 (g*16+i) 行
    }
    TransDataTo5HD<half>(dstList, srcList, params);          // 地址列表固定 16 项
}

关键点（注释里重点写了）：Transpose 只搬运连续排布的 16×16，而 TransDataTo5HD 能直接搬运"行地址带间距"的 16 个 DataBlock —— 这正是 ① 换来的大 tile 所必需的配套能力，否则还得先 gather，优化就白做了。

指令数计算：

4096²、K=4（tile = 16 行 × 64 列）：

UB 占用：K=4 时 src/dst 各 2 KiB，K=16 时各 8 KiB——离 192 KiB 三个数量级。

实测：优化版 KernelTransposeOpt 的结果（EC）

按上面的 ①+②+③ 写了新 kernel（原版一行未动，新类 KernelTransposeOpt<K> + 新入口 matrix_transpose_opt，--kernel transpose_opt --tile 64 即 K=4）：

cd ~/asst4-ascend/part1

# 先做正确性（小矩阵，快）
python3 run_benchmark.py --kernel transpose_opt -n 1024 --tile 64 --verify --clean

# 再对比性能（同规模、同核数）
python3 run_benchmark.py --kernel transpose     -n 4096 --block-dim 40 --verify
python3 run_benchmark.py --kernel transpose_opt -n 4096 --tile 64  --block-dim 40 --verify
python3 run_benchmark.py --kernel transpose_opt -n 4096 --tile 256 --block-dim 40 --verify   # K=16

# 剖析验证"瓶颈是否从 scalar 转移"
python3 run_benchmark.py --kernel transpose_opt -n 4096 --tile 64 --block-dim 40 --profile_name s3_opt_t64
# tar -czf profile.tar.gz profile/

输出结果：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel transpose_opt -n 1024 --tile 64 --verify --clean
......
[100%] Built target ascendc_kernels
[INFO ] 构建完成：/opt/atomgit/asst4-ascend/part1/build-npu/ascendc_kernels
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose_opt  (可执行文件里的入口: transpose_opt)
n           : 1024
tile        : 16 行 × 64 列  (每个 tile 1024 个元素 = 2048 B)
block-dim   : 1 个 AI Core
latency     : median 344.046 us   min 334.716 us   avg 345.417 us
bandwidth   : 12.19 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel transpose     -n 4096 --block-dim 40 --verify
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose  (可执行文件里的入口: transpose)
n           : 4096
tile        : 16 个元素  (单次搬运 32 B)
block-dim   : 40 个 AI Core
latency     : median 599.880 us   min 597.070 us   avg 599.989 us
bandwidth   : 111.87 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel transpose_opt -n 4096 --tile 64  --block-dim 40 --verify
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose_opt  (可执行文件里的入口: transpose_opt)
n           : 4096
tile        : 16 行 × 64 列  (每个 tile 1024 个元素 = 2048 B)
block-dim   : 40 个 AI Core
latency     : median 228.665 us   min 222.874 us   avg 227.304 us
bandwidth   : 293.48 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel transpose_opt -n 4096 --tile 256 --block-dim 40 --verify
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose_opt  (可执行文件里的入口: transpose_opt)
n           : 4096
tile        : 16 行 × 256 列  (每个 tile 4096 个元素 = 8192 B)
block-dim   : 40 个 AI Core
latency     : median 369.896 us   min 363.657 us   avg 368.428 us
bandwidth   : 181.43 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致

时间 / 带宽汇总（上面的输出结果里抄的 median 与带宽）：

三个可复现点：K=4 是最优列档（K=16 反而倒车到 181 GB/s）；优化 v1 在 4096² 稳定在 ~228–229 μs / 293–294 GB/s；小规模（1024²）仍被固定开销压制（只有 12.2 GB/s）。

msprof op 分析：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel transpose_opt -n 4096 --tile 64 --block-dim 40 --profile_name s3_opt_t64
[INFO ] msprof 采集（第 1 种调用形式）：msprof op --output=/opt/atomgit/asst4-ascend/part1/profile /opt/atomgit/asst4-ascend/part1/build-npu/ascendc_kernels --kernel transpose_opt --n 4096 --m 4096 --tile 64 --block-dim 40 --warmup 5 --iters 20 --reserve none --out /opt/atomgit/asst4-ascend/part1/out/transpose_opt.bin
......
[INFO ] 本次剖析产出：/opt/atomgit/asst4-ascend/part1/profile/s3_opt_t64（msprof 默认命名为 OPPROF_<时间戳>_<随机串>，已重命名）
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose_opt  (可执行文件里的入口: transpose_opt)
n           : 4096
tile        : 16 行 × 64 列  (每个 tile 1024 个元素 = 2048 B)
block-dim   : 40 个 AI Core
latency     : median 228.024 us   min 227.674 us   avg 229.002 us
bandwidth   : 294.31 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)

实测对比（数据源 report/asst4_profiles/profile-EC/s3_opt_t64；基准版为 report/asst4_profiles/profile-Step3）：

40 核后"标量"仍是四根里最长（572 μs / 690 μs）；优化 v1 里它缩到最短（99 μs / 218 μs），MTE2 顶上（179 μs）——正是表格里 scalar_ratio 46.8%、mte2_ratio 84.2% 的图形化。

结论：

1. scalar 不再是瓶颈：aiv_scalar_ratio 走完 96.6% → 83.5% → 46.8% 的三代迁移轨迹（基准 → 多核 → v1）；

2. 瓶颈转移到搬运：mte2_ratio 升到 84.2%、mte3_ratio 78.8%，Vector 计算仍只有 7.5%（非 compute-bound）；

3. 单缓冲串行的铁证：vec_wait = 96.9%、mte2_wait = 99.1%、mte3_wait = 99.5%——Vector 几乎永远在等 DMA。

   可以进一步优化，使用双缓冲要吃掉的部分（三次 wait 各 ~99% 说明流水里几乎没有任何重叠）；

4. L2 读命中 93.7% → 19.5% 是"变好"而不是"变差"：基准版的 93.7% 来自"32 B 粒度 × 128 B 行"导致一行被重读 4 次；v1 读取接近"每字节一次"，DRAM 侧不再浪费；

5. 列方向加大反而更差（实测）：--tile 256（K=16）只有 369.9 μs / 181 GB/s，远不如 K=4 的 228.7 μs / 293 GB/s——列方向的增益已经吃尽，这一路接下来的空间在"行方向（v2 的 128 行×256 B 行段）与搬运重叠"，而不是继续加大 K。

优化 v2：双缓冲 + 128×128 大 tile（实现说明）

v1 的实测把剩下的差距指到了两处，v2 就是针对它们：

指令数（K=8，tile = 128×128）：搬入 1 条 + 搬出 1 条 + 转置 8 条（每 16 行 band 一条，靠 repeatTimes=8 覆盖 128 列）= 10 条 / 16384 元素，比 v1 的 3 条 / 1024 元素再降 4.8×/每元素。

UB 预算：单份 buffer = 128×128×2 B = 32 KiB；两队列 × 双缓冲 = 4 份 = 128 KiB ≤ 192 KiB ✓（K=16 会到 256 KiB，超限——所以列方向最大 K=8）。

转置的地址推进：对第 b 个 16 行 band，repeatTimes=K、srcRepStride=1（每次前进 1 个 DataBlock = 下一个 16 列组）、dstRepStride=TILE_COLS（每次前进 16 行 × TILE_COLS 元素 = TILE_COLS 个 DataBlock）；16 个源地址 = 该 band 的 16 行行首，16 个目的地址 = 第 0 列组的 16 个目的行、偏移到"本 band 的 16 元素切片"处。

实现是新增的 KernelTransposeOpt2<K> + 新入口 matrix_transpose_opt2：

cd ~/asst4-ascend/part1

# 正确性优先（小矩阵）
python3 run_benchmark.py --kernel transpose_opt2 -n 1024 --tile 128 --verify --clean

# 与 v1、torch 基线对比（同规模、同核数）
python3 run_benchmark.py --kernel transpose_opt  -n 4096 --tile 64  --block-dim 40 --verify
python3 run_benchmark.py --kernel transpose_opt2 -n 4096 --tile 128 --block-dim 40 --verify
python3 run_benchmark.py --kernel transpose_opt2 -n 4096 --tile 64  --block-dim 40 --verify   # 对照：小列数

# 剖析（看双缓冲是否让 MTE2/MTE3 重叠、scalar 占比是否进一步下降）
python3 run_benchmark.py --kernel transpose_opt2 -n 4096 --tile 128 --block-dim 40 --profile_name s3_opt2_t128
# tar -czf profile.tar.gz profile/

约束：--tile 只支持 64/128，n % tile == 0，且 m 必须是 128 的整数倍（行方向固定 128 行一个 tile）。

输出：

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel transpose_opt2 -n 1024 --tile 128 --verify --clean
[INFO ] --clean：删除 build-npu/ 做一次干净构建
[INFO ] 构建中（soc-version=Ascend910B4）...
......
INFO ] 构建完成：/opt/atomgit/asst4-ascend/part1/build-npu/ascendc_kernels
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose_opt2  (可执行文件里的入口: transpose_opt2)
n           : 1024
tile        : 128 行 × 128 列  (双缓冲 2 份；每个 tile 16384 个元素 = 32768 B, UB 共 128 KiB)
block-dim   : 1 个 AI Core
latency     : median 70.011 us   min 67.631 us   avg 69.436 us
bandwidth   : 59.91 GB/s  (按访存量 4194304 B = 4.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel transpose_opt  -n 4096 --tile 64  --block-dim 40 --verify
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose_opt  (可执行文件里的入口: transpose_opt)
n           : 4096
tile        : 16 行 × 64 列  (每个 tile 1024 个元素 = 2048 B)
block-dim   : 40 个 AI Core
latency     : median 227.963 us   min 227.523 us   avg 228.108 us
bandwidth   : 294.38 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel transpose_opt2 -n 4096 --tile 128 --block-dim 40 --verify
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose_opt2  (可执行文件里的入口: transpose_opt2)
n           : 4096
tile        : 128 行 × 128 列  (双缓冲 2 份；每个 tile 16384 个元素 = 32768 B, UB 共 128 KiB)
block-dim   : 40 个 AI Core
latency     : median 54.061 us   min 51.460 us   avg 55.547 us
bandwidth   : 1241.35 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel transpose_opt2 -n 4096 --tile 64  --block-dim 40 --verify
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose_opt2  (可执行文件里的入口: transpose_opt2)
n           : 4096
tile        : 128 行 × 64 列  (双缓冲 2 份；每个 tile 8192 个元素 = 16384 B, UB 共 64 KiB)
block-dim   : 40 个 AI Core
latency     : median 75.401 us   min 73.631 us   avg 75.267 us
bandwidth   : 890.03 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)
[ OK ] 结果与 numpy 参考一致

atomgit@6ac4a72a17ae8e63171149c9:~/asst4-ascend/part1$ python3 run_benchmark.py --kernel transpose_opt2 -n 4096 --tile 128 --block-dim 40 --profile_name s3_opt2_t128
[INFO ] msprof 采集（第 1 种调用形式）：msprof op --output=/opt/atomgit/asst4-ascend/part1/profile /opt/atomgit/asst4-ascend/part1/build-npu/ascendc_kernels --kernel transpose_opt2 --n 4096 --m 4096 --tile 128 --block-dim 40 --warmup 5 --iters 20 --reserve none --out /opt/atomgit/asst4-ascend/part1/out/transpose_opt2.bin
2026-10-10 22:03:24 [INFO]  Op profiling analysis start.
2026-10-10 22:03:24 [INFO]  Running device task: DeviceProf
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）
[INFO] <ProfData> Start profiling on device 0, kernel: matrix_transpose_opt2_2
[INFO] <WarmUp> Warm Up enabled. times:5
[INFO ] kernel=transpose_opt2 n=4096 tile=128 blockDim=40
[INFO ] min=714.219 us  median=718.499 us  avg=720.149 us  (93.40 GB/s)
2026-10-10 22:03:28 [INFO]  Profiling running finished. All task success.
2026-10-10 22:03:28 [INFO]  Start analyze kernel on device: device0, kernel: matrix_transpose_opt2_2, name is: 0
2026-10-10 22:03:28 [INFO]  Save profiling data with metric [PipeUtilization] success.
2026-10-10 22:03:28 [INFO]  Save profiling data with metric [ArithmeticUtilization] success.
2026-10-10 22:03:28 [INFO]  Save profiling data with metric [L2Cache] success.
2026-10-10 22:03:28 [INFO]  Save profiling data with metric [Memory] success.
2026-10-10 22:03:28 [INFO]  Save profiling data with metric [ResourceConflictRatio] success.
2026-10-10 22:03:28 [INFO]  Save profiling data with metric [MemoryL0] success.
2026-10-10 22:03:28 [INFO]  Save profiling data with metric [MemoryUB] success.
2026-10-10 22:03:30 [INFO]  Performance Summary Report:

        1) MTE2 bandwidth utilization lower than 80% when active.
        2) MTE3 bandwidth utilization lower than 80% when active.
        3) aivector compute usage lower than 20%.

2026-10-10 22:03:30 [INFO]  Operator Basic Information:

        Op Name: matrix_transpose_opt2_2
        Op Type: vector
        Task Duration(us): 91.001823
        Block Dim: 40
        Mix Block Dim: 
        Device Id: 0
        Pid: 2237
        Current Freq: 1650
        Rated Freq: 1650

2026-10-10 22:03:30 [INFO]  Profiling results saved in /opt/atomgit/asst4-ascend/part1/profile/OPPROF_20261010220324_BBZQDJOXEJJIRCUI
2026-10-10 22:03:30 [INFO]  Profiling data parse finished.
2026-10-10 22:03:30 [INFO]  Op profiling finish. Welcome to next use.
[INFO ] 本次剖析产出：/opt/atomgit/asst4-ascend/part1/profile/s3_opt2_t128（msprof 默认命名为 OPPROF_<时间戳>_<随机串>，已重命名）
[UB] Tiling 侧账本：GetCoreMemSize(UB) = 196352 B，GetCoreMemSize(L1) = 524032 B
[UB] （转置 kernel 不用这个公式，本次不比对）

kernel      : transpose_opt2  (可执行文件里的入口: transpose_opt2)
n           : 4096
tile        : 128 行 × 128 列  (双缓冲 2 份；每个 tile 16384 个元素 = 32768 B, UB 共 128 KiB)
block-dim   : 40 个 AI Core
latency     : median 54.420 us   min 51.991 us   avg 56.430 us
bandwidth   : 1233.17 GB/s  (按访存量 67108864 B = 64.000 MiB 估算)

结果：

四行的演化一目了然：标量条 20.8 ms → 572 μs → 99 μs → 15 μs；v2 里 Vector（20 μs）第一次与 MTE 逼近（因为流水重叠让它和搬运并行在跑）。

v2 剖析（profile-EC/s3_opt2_t128）证实 "双缓冲生效"：

EC 判定：达标——v2 的 54.1 μs 远快于 torch 基线（162.4 μs，3.0×），也远低于原版 EC 阈值（4096² < 700 μs，12.9×），且两次 --verify 均 [OK]。