CANN Ascend C 静态 Tensor 编程实现 Matmul 基础 API 矩阵乘mmad_custom 多核样例全解析【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples导读本文基于 matmul_basic_api 样例系统讲解如何在 CANN 中以 Ascend C 基础 API 和静态 Tensor 编程范式编写一个运行在 Cube矩阵计算单元上的多核矩阵乘法核函数。读者将掌握LocalMemAllocator静态内存分配、GM→L1→L0A/L0B→L0C 的完整数据通路、Nd2NzParams/LoadData2DParams(V2)/MmadParams/FixpipeParamsV220五类核心结构体参数语义以及 MTE2/MTE1/M/FIX 流水间的SetFlag/WaitFlag同步方法最终能在 NPU、CPU 调试与仿真三种模式下完成编译、运行和精度验证。概述静态 Tensor 编程范式下的 Cube 矩阵乘本样例基于静态 Tensor 编程范式实现多核矩阵乘Matmul计算。与动态 Tensor 编程需要在 Kernel 侧手动管理内存地址不同静态编程通过LocalMemAllocator按申请顺序分配片上On-Chip内存避免开发者手工维护LocalTensor地址偏移代码结构更接近硬件数据流本身。矩阵乘的数学定义为$$ C A \times B $$其中 A 矩阵形状为[M, K]B 矩阵形状为[K, N]输出 C 矩阵形状为[M, N]。对输出矩阵 C 中每一个元素C[m, n]都累加 A 矩阵第m行与 B 矩阵第n列在 K 轴上的乘积。三个方向的定义贯穿全文M 方向矩阵 C 的行方向输出行N 方向矩阵 C 的列方向输出列K 方向矩阵乘法的内维即累加维度。支持的产品与 CANN 软件版本产品CANN 软件版本Ascend 950PR/Ascend 950DT CANN 9.1.0Atlas A3 训练系列产品/Atlas A3 推理系列产品 CANN 9.0.0Atlas A2 训练系列产品/Atlas A2 推理系列产品 CANN 9.0.0不同产品在 L0A 上的排布格式与LoadData参数结构体存在差异详见API 参数说明一节这正是本样例在代码中通过__NPU_ARCH__宏分别处理dav-2201与dav-3510两套分支的原因。样例目录结构matmul_basic_api │ ├── scripts │ │ ├── gen_data.py // 输入数据和真值数据生成脚本 │ │ └── verify_result.py // 真值对比脚本 │ ├── CMakeLists.txt // 编译工程文件 │ ├── data_utils.h // 数据读入写出函数 │ ├── matmul_basic_api.asc // Ascend C 样例实现 调用样例 │ └── README.md // 样例说明文档中文仓库中的实际路径为 Samples/0_Introduction/01_simd_cpp_api/02_matrix/matmul_basic_api 目录本目录同级的 matmul_advanced_apiMatmul 高阶 API与 matmul_tensor_apiTensor API构成了三种由底层到高层的矩阵乘实现对照本样例是最贴近硬件细节的基础版本。样例规格与多核分片策略本样例参数取M 256, N 256, K 64输入输出均为half类型、ND格式启动2 个核完成计算。任务只沿M 轴切分每个核负责输出矩阵 C 在 M 轴方向的 128 行、N 轴方向的全部 256 列第 0 个核计算 C 矩阵的第0~127行第 1 个核计算 C 矩阵的第128~255行。输入输出规格如下样例类型(OpType)Matmul样例输入nameshapedata typeformatA[M, K]halfNDB[K, N]halfND样例输出C[M, N]halfND核函数名mmad_custom从源码 matmul_basic_api.asc 可以看到这些参数在 Host 侧以constexpr形式声明M 256、K 64、N 256、singleCoreM 128单核计算量、以及矩阵乘执行单元的 Tile 参数baseM 128、baseK 64、baseN 256。其中baseM × baseN × baseK是单核一次Mmad计算的矩阵块大小——在本样例中恰好等于单核计算量因此每个核只需执行一次完整的搬运—计算—回写流程。Kernel 侧实现从 GM 到 L0C 的完整数据流mmad_custom是一个__global__ __cube__核函数表示该函数运行在 AI Core 的 Cube矩阵计算单元上。样例采用静态 Tensor 编程方式通过LocalMemAllocator创建LocalTensor。CUBE_BLOCK 16表示 half 数据类型的矩阵分形Fractal为16 × 16后续所有搬运均以16 × 16分形为单位进行。1. 创建 GlobalTensor 并计算核偏移AscendC::GlobalTensorhalf aGM, bGM, cGM; uint32_t mIterIdx AscendC::GetBlockIdx() % (M / singleCoreM); aGM.SetGlobalBuffer((__gm__ half*)a mIterIdx * singleCoreM * K); bGM.SetGlobalBuffer((__gm__ half*)b); cGM.SetGlobalBuffer((__gm__ half*)c mIterIdx * singleCoreM * N);通过AscendC::GetBlockIdx()获取当前核号并计算mIterIdx。由于只沿 M 轴切分任务GM 地址偏移规则为aGM偏移mIterIdx * singleCoreM * K使当前核读取自己负责的 A 矩阵行块bGM不偏移因为每个核都需要读取完整的 B 矩阵参与 K 轴累加cGM偏移mIterIdx * singleCoreM * N使当前核把结果写回 C 矩阵中自己负责的行块。2. LocalMemAllocator 静态分配片上内存AscendC::LocalMemAllocatorAscendC::Hardware::L1 l1Allocator; AscendC::LocalMemAllocatorAscendC::Hardware::L0A l0aAllocator; AscendC::LocalMemAllocatorAscendC::Hardware::L0B l0bAllocator; AscendC::LocalMemAllocatorAscendC::Hardware::L0C l0cAllocator; AscendC::LocalTensorhalf a1Local l1Allocator.AllocAscendC::TPosition::A1, half(baseM * baseK); AscendC::LocalTensorhalf b1Local l1Allocator.AllocAscendC::TPosition::B1, half(baseK * baseN); AscendC::LocalTensorhalf a2Local l0aAllocator.AllocAscendC::TPosition::A2, half(baseM * baseK); AscendC::LocalTensorhalf b2Local l0bAllocator.AllocAscendC::TPosition::B2, half(baseK * baseN); AscendC::LocalTensorfloat cLocal l0cAllocator.AllocAscendC::TPosition::CO1, float(baseM * baseN);各LocalTensor的角色如下LocalTensor所在存储作用a1LocalL1A 矩阵临时存储b1LocalL1B 矩阵临时存储与a1Local共用同一 L1 allocator按申请顺序分配避免手动维护 L1 地址偏移a2LocalL0AA 矩阵存储供Mmad读取b2LocalL0BB 矩阵存储供Mmad读取cLocalL0C矩阵乘结果临时存储float 累加精度注意cLocal的元素类型是floatCube 计算内部以更高精度累加half输入经计算后在 L0C 中保存为float结果最后由Fixpipe转回half写回 GM。3. GM→L1DataCopy 与 ND→Nz 格式转换AscendC::DataCopy(a1Local, aGM, AscendC::Nd2NzParams{1, baseM, baseK, 0, K, baseM, 1, 0}); AscendC::DataCopy(b1Local, bGM, AscendC::Nd2NzParams{1, baseK, baseN, 0, N, baseK, 1, 0});DataCopy属于MTE2流水。此处使用Nd2NzParams参数在搬运过程中直接把 GM 中ND格式的 A、B 矩阵转换为 Cube 计算需要的Nz格式存入 L1。4. MTE2→MTE1 流水同步AscendC::SetFlagAscendC::HardEvent::MTE2_MTE1(EVENT_ID0); AscendC::WaitFlagAscendC::HardEvent::MTE2_MTE1(EVENT_ID0);SetFlag用于生产者在完成当前任务后写入同步事件WaitFlag用于消费者等待该事件HardEvent模板参数描述同步方向EVENT_ID区分同类同步事件。DataCopy由 MTE2 流水完成后续LoadData由MTE1流水读取 L1 数据因此必须设置MTE2_MTE1事件避免LoadData读取到尚未搬运完成的数据。5. L1→L0A/L0BLoadData 搬运按架构分支LoadData属于MTE1流水将 A 矩阵从 L1 搬运到 L0A、B 矩阵从 L1 搬运到 L0B。L0A/L0B 是 Cube 矩阵计算单元直接读取的输入缓存。由于产品架构差异源码通过__NPU_ARCH__宏分两套实现dav-2201Atlas A2/A3 系列分支——L0A 上排布格式为Zz使用LoadData2DParams结构体以循环逐分形搬运// A矩阵 L1-L0ANz-Zz循环 baseM/CUBE_BLOCK 次 for (int i 0; i baseM / CUBE_BLOCK; i) { AscendC::LoadData( a2Local[i * baseK * CUBE_BLOCK], a1Local[i * 512 / sizeof(half)], AscendC::LoadData2DParams{0, baseK / CUBE_BLOCK, baseM / CUBE_BLOCK, 0, 0, false, 0}); } // B矩阵 L1-L0BNz-ZnifTransposetrue 完成转置搬运 for (int i 0; i baseK / CUBE_BLOCK; i) { AscendC::LoadData( b2Local[i * baseN * CUBE_BLOCK], b1Local[i * 512 / sizeof(half)], AscendC::LoadData2DParams{0, baseN / CUBE_BLOCK, baseK / CUBE_BLOCK, 0, 0, true, 0}); }dav-3510Ascend 950PR/950DT分支——L0A 上排布格式为Nz使用LoadData2DParamsV2结构体一次调用即可完成整个矩阵的搬运// A矩阵 L1-L0ANz-Nz一次完成 AscendC::LoadData(a2Local, a1Local, AscendC::LoadData2DParamsV2{ 0, 0, baseM / CUBE_BLOCK, baseK / CUBE_BLOCK, baseM / CUBE_BLOCK, baseM / CUBE_BLOCK, false, 0}); // B矩阵 L1-L0BNz-Zn一次完成转置搬运 AscendC::LoadData(b2Local, b1Local, AscendC::LoadData2DParamsV2{ 0, 0, baseK / CUBE_BLOCK, baseN / CUBE_BLOCK, baseK / CUBE_BLOCK, baseN / CUBE_BLOCK, true, 0});B 矩阵ifTransposetrue的本质Cube 计算要求 A 按行主序M×K、B 按列主序K×N排列B 在 L0B 上需要以 Zn 排布呈现因此在搬运时对每个分形做转置。6. MTE1→M 流水同步AscendC::SetFlagAscendC::HardEvent::MTE1_M(EVENT_ID0); AscendC::WaitFlagAscendC::HardEvent::MTE1_M(EVENT_ID0);LoadData属于 MTE1 流水后续Mmad属于M矩阵计算流水M 流水必须等待 MTE1 完成避免Mmad读取到尚未搬运完成的 L0A/L0B 数据。7. Mmad 执行矩阵乘AscendC::Mmad(cLocal, a2Local, b2Local, AscendC::MmadParams{baseM, baseN, baseK, 0, false, true});Mmad属于M流水从 L0A/L0B 读取输入、执行矩阵乘并累加到 L0C。此处baseM 128、baseN 256、baseK 64对应单核一次计算的矩阵块大小。8. M→FIX 流水同步与 Fixpipe 回写AscendC::SetFlagAscendC::HardEvent::M_FIX(EVENT_ID0); AscendC::WaitFlagAscendC::HardEvent::M_FIX(EVENT_ID0); AscendC::Fixpipe( cGM, cLocal, AscendC::FixpipeParamsV220{baseN, baseM, baseM, N, false, QuantMode_t::F322F16, 0, 1, 0, 0, 0});Fixpipe属于FIX流水将 L0C 中的float累加结果通过量化模式F322F16float→half转换为half并写回 GM 中 C 矩阵的输出位置。FIX 流水必须等待 M 流水完成避免读取到尚未计算完成的 L0C 结果。9. 收尾同步AscendC::PipeBarrierPIPE_ALL();最后调用PipeBarrierPIPE_ALL()确保当前核内所有相关流水任务全部完成为核函数退出与后续 Host 侧数据回读提供保障。从源码结构看整个 Kernel 侧形成了一条清晰的流水链MTE2DataCopy→ MTE1LoadData→ MMmad→ FIXFixpipe每两级之间用一对SetFlag/WaitFlag完成生产者—消费者同步这是 Ascend C Cube 编程中最核心的流水编排范式。API 参数说明五个核心结构体以下结构体均以花括号{}方式传参各字段含义如下字段顺序与 API 文档保持一致实际 struct 声明中部分字段顺序可能不同AscendC::Nd2NzParams —— DataCopy 接口使用描述 ND→Nz 格式转换参数struct Nd2NzParams { int32_t ndNum; // 传输ND矩阵的数目[0, 4095] uint16_t nValue; // ND矩阵的行数[0, 16384] int32_t dValue; // ND矩阵的列数[0, 65535] int32_t srcNdMatrixStride; // 相邻ND矩阵起始地址偏移单位元素[0, 65535] int32_t srcDValue; // 同一ND矩阵相邻行偏移单位元素[1, 65535] uint16_t dstNzC0Stride; // 目的Nz中同源行转换后多行相邻偏移单位C0_SIZE(32B)[1, 16384] uint16_t dstNzNStride; // 目的Nz中Z型矩阵相邻行偏移单位C0_SIZE(32B)[1, 16384] int32_t dstNzMatrixStride; // 目的Nz中相邻Nz矩阵起始地址偏移单位元素[1, 65535] };例如搬运 A 矩阵时{1, baseM, baseK, 0, K, baseM, 1, 0}将 baseM×baseK 的 ND 数据转为 Nz 格式。对照源码nValue baseM行数、dValue baseK列数、srcDValue KGM 中 A 矩阵相邻行的元素间隔即 A 的原始行宽、dstNzC0Stride baseM目的端分形行偏移。B 矩阵同理为{1, baseK, baseN, 0, N, baseK, 1, 0}。AscendC::LoadData2DParams —— LoadData 接口使用A2/A3 系列描述 Atlas A2 训练系列产品/Atlas A2 推理系列产品、Atlas A3 训练系列产品/Atlas A3 推理系列产品中 A 矩阵 L1→L0A 和 B 矩阵 L1→L0B 的数据搬运参数struct LoadData2DParams { int32_t startIndex; // 分形矩阵ID0为第1个单位512B[0, 65535] int32_t repeatTimes; // 迭代次数每个迭代处理512B[1, 255] int32_t srcStride; // 相邻迭代源分形起始地址间隔单位512B[0, 65535] int32_t sid; // 预留配置为0 int32_t dstGap; // 目的端相邻迭代分形间隔单位512B[0, 65535] bool ifTranspose; // 是否转置每个分形默认false bool addrMode; // 地址更新方式false递增true递减默认false };例如在 A2/A3 系列产品中L0A 上排布格式为 Zz搬运 A 矩阵时用{0, baseK / CUBE_BLOCK, baseM / CUBE_BLOCK, 0, 0, false, 0}即repeatTimes 4次迭代、每次处理一个16×16half 分形搬运 B 矩阵时ifTransposetrue完成 Nz→Zn 的转置搬运。注意 half 类型一个16×16分形恰好为 512B与字段单位吻合。AscendC::LoadData2DParamsV2 —— LoadData 接口使用Ascend 950PR/950DT描述 Ascend 950PR/Ascend 950DT 产品中 A 矩阵 L1→L0A 和 B 矩阵 L1→L0B 的数据搬运参数struct LoadData2DParamsV2 { uint32_t mStartPosition; // M方向起始位置单位512B uint32_t kStartPosition; // K方向起始位置单位512B uint16_t mStep; // M方向搬运分形数 uint16_t kStep; // K方向搬运分形数 int32_t srcStride; // 源端相邻K方向分形间隔单位512B uint16_t dstStride; // 目的端相邻K方向分形间隔单位512B bool ifTranspose; // 是否转置每个分形默认false uint8_t sid; // 预留配置为0 };Ascend 950PR/950DT 产品中 L0A 排布格式为 Nz搬运 A 矩阵时使用{0, 0, baseM / CUBE_BLOCK, baseK / CUBE_BLOCK, baseM / CUBE_BLOCK, baseM / CUBE_BLOCK, false, 0}一次完成 A 矩阵 Nz→Nz 搬运搬运 B 矩阵时使用{0, 0, baseK / CUBE_BLOCK, baseN / CUBE_BLOCK, baseK / CUBE_BLOCK, baseN / CUBE_BLOCK, true, 0}一次完成 B 矩阵 Nz→Zn 搬运。AscendC::MmadParams —— Mmad 接口使用描述矩阵乘参数struct MmadParams { uint16_t m; // 左矩阵HeightM维[0, 4095] uint16_t n; // 右矩阵WidthN维[0, 4095] uint16_t k; // 左矩阵Width/右矩阵HeightK维[0, 4095] uint16_t unitFlag; // Mmad与Fixpipe细粒度并行控制默认0 bool cmatrixSource; // C矩阵初始值来源falseCO1trueC2默认false bool cmatrixInitVal; // C矩阵初始值是否为0默认true };例如{baseM, baseN, baseK, 0, false, true}计算 baseM×baseN 输出块并在 K 方向累加 baseK 长度。其中cmatrixInitVal true表示 C 矩阵初始值为 0本样例为单次累加不需要叠加外部偏置或残差cmatrixSource false表示从 CO1 读取与前面cLocal分配在TPosition::CO1保持一致。AscendC::FixpipeParamsV220 —— Fixpipe 接口使用描述 L0C→GM 的数据搬运和精度转换参数struct FixpipeParamsV220 { int32_t nSize; // 源Nz矩阵N方向大小[1, 4095] uint16_t mSize; // 源Nz矩阵M方向大小Nz2ND时[1, 8192] uint16_t srcStride; // 源Nz相邻Z排布起始偏移单位C0_SIZE[0, 65535] int32_t dstStride; // Nz2ND时目的ND矩阵每行元素数单位element bool reluEn; // 是否使能ReLU QuantMode_t quantPre; // 量化模式F322F16表示float→half uint64_t deqScalar; // scalar量化参数单个scale值 int32_t ndNum; // 源Nz矩阵数目[1, 65535] int32_t srcNdStride; // 不同Nz矩阵起始地址间隔单位16×C0_SIZE[1, 512] int32_t dstNdStride; // 目的相邻ND矩阵偏移单位element[1, 65535] int32_t unitFlag; // Mmad与Fixpipe并行控制 };例如{baseN, baseM, baseM, N, false, F322F16, 0, 1, 0, 0, 0}将 L0C 中的 baseM×baseN float32 结果转为 half 并写回 GM。字段与源码对照nSize baseNN 方向 256、mSize baseMM 方向 128、srcStride baseM、dstStride NGM 中 C 矩阵每行元素数即回写行宽、quantPre F322F16float→half 精度转换、ndNum 1单个 Nz 矩阵。Host 侧调用实现Host 侧在 matmul_basic_api.asc 的main函数中完成完整的 ACL 生命周期管理声明矩阵规格常量M/K/N、单核计算量singleCoreM与 Tile 参数baseM/baseK/baseNaclInit→aclrtSetDevice→aclrtCreateStream完成设备初始化为 A、B、C 分配 Host 与 Device 内存通过ReadFile读取./input/x1_gm.bin、./input/x2_gm.bin再以aclrtMemcpyH2D拷入设备使用内核调用符启动核函数mmad_customM, K, N, singleCoreM, baseM, baseK, baseNnumBlocks, 0, stream(aDevice, bDevice, cDevice);调用时模板参数传入矩阵规格M、K、N、单核计算量singleCoreM和基础 Tile 大小baseM、baseK、baseN运行时参数传入 Device 侧 A、B、C 矩阵地址numBlocks 2即启动 2 个核aclrtSynchronizeStream等待执行完成aclrtMemcpyD2H回拷结果WriteFile写出./output/output.bin依次释放内存、销毁流、复位设备、aclFinalize收尾。编译与运行在本样例根目录下执行如下步骤编译并执行样例。1. 配置环境变量根据当前环境上 CANN 开发套件包的安装方式配置环境变量source ${install_path}/cann/set_env.sh说明${install_path}为 CANN 包安装目录未指定安装目录时默认安装至/usr/local/Ascend下。2. 样例执行在本样例目录下执行如下命令mkdir -p build cd build; # 创建并进入build目录 cmake -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # 编译工程默认npu模式 python3 ../scripts/gen_data.py # 生成测试输入数据 ./demo # 执行编译生成的可执行程序运行样例 python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确确认算法逻辑正确使用 CPU 调试或 NPU 仿真模式时添加-DCMAKE_ASC_RUN_MODEcpu或-DCMAKE_ASC_RUN_MODEsim参数即可cmake -DCMAKE_ASC_RUN_MODEcpu -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # CPU调试模式 cmake -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # NPU仿真模式注意切换编译模式前需清理 cmake 缓存可在 build 目录下执行rm CMakeCache.txt后重新 cmake。从 CMakeLists.txt 可见工程通过find_package(ASC REQUIRED)引入 CANN 的 ASC 语言编译支持并将--npu-arch${CMAKE_ASC_ARCHITECTURES}传入编译选项最终把单个 matmul_basic_api.asc 编译为可执行文件demo。3. 编译选项说明选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201 对应 Atlas A2 训练系列产品/Atlas A2 推理系列产品和 Atlas A3 训练系列产品/Atlas A3 推理系列产品dav-3510 对应 Ascend 950PR/Ascend 950DT4. 执行结果执行结果如下说明精度对比成功test pass!数据生成与精度验证机制gen_data.py与verify_result.py构成了本样例的自动化验证闭环scripts/gen_data.py按m256, n256, k64随机生成[-10, 10]区间的float16输入x1_gm、x2_gm并用float32 精度执行np.matmul得到 golden 结果先提升到 float32 计算再转回 float16模拟 Cube 累加精度分别落盘到./input/与./output/golden.binscripts/verify_result.py读取output.bin与golden.bin以相对容差1e-6、绝对容差1e-9逐元素比对统计误差率并与1e-4的容忍阈值比较误差率不超阈值即输出test pass!否则打印首个错误元素的位置、期望值、实际值与相对偏差并退出非零。功能调试printf 与 DumpTensorprintfprintf接口提供 CPU 域或 NPU 域调试场景下的格式化输出功能。在算子 Kernel 侧实现代码中需要输出日志信息的地方调用AscendC::printf(matmul blockIdx%d\n, AscendC::GetBlockIdx());[!CAUTION]注意 printfPRINTF接口打印功能会对算子实际运行的性能带来一定影响通常在调测阶段使用。开发者可以按需通过设置ASCENDC_DUMP0的方式关闭打印功能。DumpTensorDumpTensor接口用于 Dump 指定 Tensor 的内容同时支持打印自定义附加信息仅支持uint32_t数据类型的信息比如打印当前行号。在 Kernel 侧需要打印 Tensor 数据的地方调用AscendC::DumpTensor(cLocal, baseM * baseN);上述示例打印 L0C 中 baseM×baseN 个元素即当前核的完整计算结果。[!CAUTION]注意 DumpTensor 接口打印功能会对算子实际运行的性能带来一定影响通常在调测阶段使用。开发者可以按需通过设置ASCENDC_DUMP0来关闭打印功能。性能调试msOpProf 单算子性能分析msOpProf 工具是单算子性能分析工具包含msopprof和msopprof simulator两种使用方式。该工具协助用户定位算子内存、算子代码以及算子指令的异常实现全方位的算子调优当前支持基于不同运行模式上板或仿真和不同文件形式可执行文件或算子二进制.o文件进行性能数据的采集和自动解析。上板性能采集通过上板性能采集可以直接测定算子在昇腾 AI 处理器上的运行时间适合在板环境中快速定位算子性能问题。基于可执行文件demo通过msopprof执行算子调优msopprof ./demo性能数据说明命令完成后会在默认目录下生成以OPPROF_{timestamp}_XXX命名的文件夹性能数据文件夹结构示例如下├──dump # 原始的性能数据用户无需关注 ├──ArithmeticUtilization.csv # cube/vector指令cycle占比 ├──L2Cache.csv # L2 Cache命中率影响MTE2建议合理规划数据搬运逻辑增加命中率 ├──Memory.csv # UBL1和主存储器读写带宽速率 ├──MemoryL0.csv # L0AL0B和L0C读写带宽速率 ├──MemoryUB.csv # Vector和Scalar到UB的读写带宽速率 ├──OpBasicInfo.csv # 算子基础信息 ├──PipeUtilization.csv # 采集计算单元和搬运单元耗时和占比 ├──ResourceConflictRatio.csv # UB上的bank group、bank conflict和资源冲突率在所有指令中的占比 └──visualize_data.bin # MindStudio Insight呈现文件对于本样例的 Cube 矩阵乘场景重点应关注ArithmeticUtilization.csvCube 指令的 cycle 占比直接反映Mmad计算密度PipeUtilization.csvMTE2/MTE1/M/FIX 各流水耗时占比用于判断流水是否均衡、是否存在同步等待MemoryL0.csvL0A/L0B/L0C 的读写带宽验证LoadData与Fixpipe的搬运效率L2Cache.csv命中率影响 MTE2 的 GM→L1 搬运合理规划数据搬运逻辑如 B 矩阵的复用可提升命中率。总结本样例以最基础的方式完整展示了 Ascend C Cube 编程的核心链路mmad_custom核函数沿 M 轴切分任务到 2 个核通过LocalMemAllocator静态分配 L1/L0A/L0B/L0C 片上内存依次执行DataCopyMTE2ND→Nz、LoadDataMTE1L1→L0、MmadM矩阵乘、FixpipeFIXfloat→half 回写并在每级流水切换处用SetFlag/WaitFlag建立生产者—消费者同步。五个参数结构体Nd2NzParams、LoadData2DParams、LoadData2DParamsV2、MmadParams、FixpipeParamsV220的字段语义与取值规则是后续深入 Matmul 高阶 APImatmul_advanced_api乃至各类性能调优样例如 matmul_story的重要基础。【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考