AscendC Matmul SplitM 多核切 M 直调样例深度解析：AIC:AIV=1:N 架构下的 M 轴切分并行

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend

导读

本文基于 asc-devkit 仓库中的matmul_splitm直调样例，系统讲解 AscendC 高级 Matmul 算子中SplitM（多核切 M）策略的实现原理与完整使用流程。该策略面向 Ascend 950PR&950DT 系列产品，解决"AIC 完成一次 Iterate 计算后，中间结果需按 M 轴切分给多个 AIV 并行做后续处理"的场景。读完本文，你将掌握 SplitM 策略的适用条件、样例的编译运行方法、内核代码中 M/N 双维度分核寻址逻辑，以及 Tiling 侧如何通过MultiCoreMatmulTiling生成切分参数，从而能够独立开展 SplitM 场景的算子开发与调试。

一、SplitM 策略解决的问题与适用场景

在传统 Matmul 多核并行方案中，常见的是把 M 轴、N 轴同时切分到不同核（SplitMN），或仅在 N 轴切分（SplitN）。而本样例演示的SplitM 策略面向一种特殊的硬件协作模式——AscendC 分离模式下的 AIC:AIV=1:N 架构：

• AIC（AI Core - Cube 侧）负责大算力的矩阵乘累加（Iterate）计算；
• 一次 Iterate 结束后，中间结果并不直接写回 GM，而是按 M 轴切分成多份；
• 切分后的多份结果分别交给N 个 AIV（AI Core - Vector 侧）做后续的向量类处理（本样例中即为搬出到 Global Memory）。

这种"一份结果、多个接收者"的搬运方式，在底层对应的是 Matmul 内部以McgShfMode::DUAL_DST_SPLIT_M描述的双目标（DUAL DST）数据分发模式。从仓库源码看，该模式在 impl/adv_api/detail/matmul/policy/matmul_policy.h 中由SplitMMatmulPolicy策略类引入：

/* SplitMMatmulPolicy is considered entirely experimental. We retain the freedom to make incompatible changes, but do not guarantee the stability. SplitMMatmulPolicy is only for internal usage, does not support extension or customized specialization! */ template <const auto& MM_CFG, typename IMPL, typename A_TYPE, typename B_TYPE, typename C_TYPE, typename BIAS_TYPE> struct SplitMMatmulPolicy : public MatmulPolicy<MM_CFG, IMPL, A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE> { public: using CopyCubeOut = AscendC::Impl::Detail::CopyCubeOut<IMPL, A_TYPE, B_TYPE, C_TYPE, MM_CFG, McgShfMode::DUAL_DST_SPLIT_M>; };

与默认MatmulPolicy相比，SplitMMatmulPolicy的核心差异在于将CopyCubeOut（L0C 结果搬出阶段）替换为McgShfMode::DUAL_DST_SPLIT_M模式，从而支持把一块 Cube 计算结果按 M 轴拆分输出到两个不同的目的地址。需要注意的是，源码注释明确声明该策略仅限内部使用、处于实验性质，不承诺 API 稳定性，不支持用户扩展或自定义特化，因此在实际工程中应将其视为特定硬件形态下的专用通道而非通用方案。

二、支持的产品与 CANN 版本

对应的 NPU 架构代号为dav-3510，这也是样例 CMake 配置中CMAKE_ASC_ARCHITECTURES的默认取值。

三、环境准备与编译运行

3.1 配置环境变量

请根据当前环境上 CANN 开发套件包的安装方式完成安装后，配置环境变量：

source ${install_path}/cann/set_env.sh

说明：${install_path}为 CANN 包安装目录，未指定安装目录时默认安装至/usr/local/Ascend下。

3.2 NPU 模式构建与执行

在样例根目录examples/01_simd_cpp_api/04_advanced_api/00_matmul/matmul_splitm/下依次执行：

mkdir -p build && cd build; # 创建并进入build目录 cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程，默认npu模式 python3 ../scripts/gen_data.py # 生成测试输入数据 ./demo # 执行编译生成的可执行程序，执行样例 python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确

执行成功后，精度校验脚本会打印：

test pass!

3.3 CPU 调试与 NPU 仿真模式

需要本地调试或环境受限时，可添加-DCMAKE_ASC_RUN_MODE参数切换到 CPU 调试或 NPU 仿真模式：

cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # CPU调试模式 cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # NPU仿真模式

注意：切换编译模式前需清理 cmake 缓存，可在 build 目录下执行rm CMakeCache.txt后重新 cmake。

3.4 编译选项说明

上述两个选项在 matmul_splitm/CMakeLists.txt 中声明为可缓存字符串变量，并通过find_package(ASC REQUIRED)引入 ASC 编译工具链，最终把matmul_splitm.asc编译为名为demo的可执行程序，链接tiling_api、register、platform等运行库，并携带--npu-arch=${CMAKE_ASC_ARCHITECTURES}与-DENABLE_CV_COMM_VIA_SSBUF=true编译选项。其中ENABLE_CV_COMM_VIA_SSBUF与 AIC、AIV 之间通过 SSBUF（共享片上缓冲区）通信的实现相关，与 SplitM 场景下的跨核数据分发配套使用。

四、样例结构与数据规格

样例目录结构如下：

matmul_splitm/ ├── CMakeLists.txt # 构建配置 ├── README.md / README_en.md ├── data_utils.h # 文件读写等辅助工具 ├── matmul_splitm.asc # 内核源码（含 Host 侧 main） └── scripts/ ├── gen_data.py # 生成输入与 golden 数据 └── verify_result.py # 输出精度校验

本样例计算的矩阵乘法为C = A × B + Bias，默认形状为：

constexpr uint32_t M = 127; constexpr uint32_t N = 127; constexpr uint32_t K = 63;

• A：M × K = 127 × 63，half（FP16），非转置；
• B：K × N = 63 × 127，half（FP16），非转置；
• Bias：1 × N = 127，float（FP32）；
• C：M × N = 127 × 127，float（FP32），输出位置为 VECIN。

选择 127 / 63 这类非对齐奇数尺寸，正是为了覆盖 M、N 方向上singleCoreM、singleCoreN无法整除时的尾块（tail）处理路径。

五、内核实现深度解析

内核主体位于 matmul_splitm.asc，整体分为四层：Host 侧main、Tiling 生成GenerateTiling、全局入口matmul_splitm_custom、以及核心计算类MatmulKernel。

5.1 Matmul 对象模板参数：如何启用 SplitM 策略

MatmulKernel类的核心是成员matmulObj，其模板参数完整声明了参与计算的矩阵属性与策略：

AscendC::Matmul< AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, AType, false, LayoutMode::NONE, true>, AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, BType, false, LayoutMode::NONE, true>, AscendC::MatmulType<AscendC::TPosition::VECIN, CubeFormat::ND_ALIGN, CType>, AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, BiasType>, CFG_NORM, AscendC::MatmulCallBackFunc<nullptr, nullptr, nullptr>, AscendC::Impl::Detail::SplitMMatmulPolicy> matmulObj;

逐项解读：

其中C 矩阵位置为VECIN与最后一个策略参数SplitMMatmulPolicy必须配套使用，二者共同构成"AIC 算完、结果留在向量侧、按 M 切分给多个 AIV"的完整数据通路。

5.2 初始化：GM 张量与分核偏移计算

Init中首先将 A/B/C/Bias 的GlobalTensor绑定到 GM 地址：

aGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ AType*>(a), tiling.M * tiling.Ka); bGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ BType*>(b), tiling.Kb * tiling.N); cGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ CType*>(c), tiling.M * tiling.N); biasGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ BiasType*>(bias), tiling.N);

随后按核号计算每个核负责的数据分片偏移：

// V0 V1 SameAB所以/2 CalcOffset(AscendC::GetBlockIdx() / 2, offsetA, offsetB, offsetC, offsetBias);

CalcOffset是 SplitM 分核逻辑的核心，实现了M 方向优先、N 方向其次的二维分核：

auto mSingleBlocks = (tiling.M + tiling.singleCoreM - 1) / tiling.singleCoreM; // split M into mSingleBlocks cores mCoreIndex = blockIdx % mSingleBlocks; nCoreIndex = blockIdx / mSingleBlocks; offsetA = mCoreIndex * tiling.Ka * tiling.singleCoreM; if (isTransA) { offsetA = mCoreIndex * tiling.singleCoreM; } offsetB = nCoreIndex * tiling.singleCoreN; if (isTransB) { offsetB = nCoreIndex * tiling.Kb * tiling.singleCoreN; } offsetC = mCoreIndex * tiling.N * tiling.singleCoreM + nCoreIndex * tiling.singleCoreN; offsetBias = nCoreIndex * tiling.singleCoreN;

• mSingleBlocks = ceil(M / singleCoreM)：M 方向切分的核数；
• mCoreIndex = blockIdx % mSingleBlocks：当前核在 M 方向上的序号；
• nCoreIndex = blockIdx / mSingleBlocks：当前核在 N 方向上的序号；
• A 的偏移按mCoreIndex沿 M 平移singleCoreM行；B 的偏移按nCoreIndex沿 N 平移singleCoreN列；
• C 的偏移是 M、N 两个方向偏移的叠加，这正是"先切 M、再切 N"的二维 Tile 布局。

关于GetBlockIdx() / 2的注释"V0 V1 SameAB所以/2"：在 AIC:AIV=1:N 的硬件分组中，同一 AIC 会带动若干 AIV（本场景下可见 2 个），它们共享同一份 A/B 输入（SameAB），因此用GetBlockIdx() / 2把"同一个 AIC 绑定的两个 AIV"归并到同一个分核索引上，避免同一份输入被重复切分。GetBlockIdx与GetSubBlockIdx的系统实现在 impl/basic_api/dav_3510/kernel_operator_sys_var_impl.h 中，其中GetBlockIdx在 AIV 分支会展开为get_block_idx() * get_subblockdim() + get_subblockid()，即把 AIC 块号与 AIV 子块号合并为全局核号——这正是 SplitM 场景区分"哪个 AIC、哪个 AIV"的基础。

缓冲区初始化同样体现了"半块"设计：

// 初始化按base/2大小初始化 halfBaseSize = tiling.baseM / 2 * tiling.baseN; pipe->InitBuffer(vecin, 1, halfBaseSize * sizeof(CType)); pipe->InitBuffer(vecout, 1, halfBaseSize * sizeof(CType));

因为一次 Iterate 产生的baseM × baseN结果块会被切分成上下两半分别交给两个 AIV，所以每个 AIV 只需申请baseM/2 × baseN大小的向量缓冲区，这正是vecin、vecout两个队列按"半块"初始化的原因。

5.3 Process：Iterate 主循环与双目标搬出

Process完成一次完整的计算流水：设置输入 → 处理尾块 → 双层循环 Iterate → 从 VECIN 取结果 → 按子块索引搬出到 GM：

matmulObj.SetTensorA(aGlobal, isTransA); matmulObj.SetTensorB(bGlobal, isTransB); if (tiling.isBias) { matmulObj.SetBias(biasGlobal); }

随后计算当前核实际的单核尺寸，并在 M 或 N 出现尾块时通知 Matmul 对象：

int tailM = tiling.M - mCoreIndex * tiling.singleCoreM; tailM = tailM < tiling.singleCoreM ? tailM : tiling.singleCoreM; int tailN = tiling.N - nCoreIndex * tiling.singleCoreN; tailN = tailN < tiling.singleCoreN ? tailN : tiling.singleCoreN; uint32_t curSingleCoreM = (tailM < tiling.singleCoreM) ? tailM : tiling.singleCoreM; uint32_t curSingleCoreN = (tailN < tiling.singleCoreN) ? tailN : tiling.singleCoreN; if (tailM < tiling.singleCoreM || tailN < tiling.singleCoreN) { matmulObj.SetTail(curSingleCoreM, curSingleCoreN); }

主循环按orderm（M 在外层、N 在内层）顺序遍历：

uint16_t nIter = (curSingleCoreN + tiling.baseN - 1) / tiling.baseN; uint16_t mIter = (curSingleCoreM + tiling.baseM - 1) / tiling.baseM; for (int j = 0; j < nIter; j++) { for (int i = 0; i < mIter; i++) { matmulObj.Iterate(); AscendC::LocalTensor<CType> vecinTensor = vecin.AllocTensor<CType>(); matmulObj.GetTensorC(vecinTensor, 0, true); // 从 VECIN 取回本次 Iterate 的 C 结果 vecin.EnQue(vecinTensor); AscendC::LocalTensor<CType> vecinLocal = vecin.DeQue<CType>(); AscendC::LocalTensor<CType> vecoutTensor = vecout.AllocTensor<CType>(); DataCopy(vecoutTensor, vecinLocal, halfBaseSize); // 为方便直接拷贝整个大小 vecout.EnQue(vecoutTensor); AscendC::LocalTensor<CType> vecoutLocal = vecout.DeQue<CType>(); vecin.FreeTensor(vecinLocal); ... } }

GetTensorC是 Matmul 高级 API 中"Iterate 后取出 C 矩阵分片"的标准接口（其声明位于 include/adv_api/matmul/matmul.h），此处第三个参数true对应"等待 Iterate 完成后再取数"的同步语义。由于本场景 C 位置为 VECIN，取回的结果直接进入向量侧队列vecin，再经DataCopy转入vecout统一准备搬出。

搬出阶段使用DataCopyPad按块（block）组织搬运，块计数按baseM / 2切分，从而把一整块baseM × baseN的结果拆成上下两半：

// 搬出主块参数 uint16_t blockCountV0 = tiling.baseM / 2; uint16_t blockCountV1 = blockCountV0; uint32_t blockLen = tiling.baseN * sizeof(CType); uint32_t srcStride = 0; uint32_t dstStride = (tiling.N - tiling.baseN) * sizeof(CType); if ((j == nIter - 1) && (i != mIter - 1)) { // n方向最后一块 blockLen = tailBaseN * sizeof(CType); dstStride = (tiling.N - tailBaseN) * sizeof(CType); } else if ((j != nIter - 1) && (i == mIter - 1)) { // m方向最后一块 blockCountV0 = (tailBaseM + 2 - 1) / 2; blockCountV1 = (tailBaseM % 2 == 0) ? blockCountV0 : (blockCountV0 - 1); } else if ((j == nIter - 1) && (i == mIter - 1)) { // 双方向最后一块 blockCountV0 = (tailBaseM + 2 - 1) / 2; blockCountV1 = (tailBaseM % 2 == 0) ? blockCountV0 : (blockCountV0 - 1); blockLen = tailBaseN * sizeof(CType); dstStride = (tiling.N - tailBaseN) * sizeof(CType); } uint64_t cGlobalOffsetV0 = j * tiling.baseN + i * tiling.baseM * tiling.N; uint64_t cGlobalOffsetV1 = cGlobalOffsetV0 + blockCountV0 * tiling.N; if (AscendC::GetSubBlockIdx() == 0) { // v0 DataCopyPad(cGlobal[cGlobalOffsetV0], vecoutLocal, {blockCountV0, blockLen, srcStride, dstStride, 0}); } else { // v1 if (blockCountV1 > 0) { DataCopyPad(cGlobal[cGlobalOffsetV1], vecoutLocal, {blockCountV1, blockLen, srcStride, dstStride, 0}); } }

这里的分工非常清晰：

• GetSubBlockIdx() == 0（V0）：把结果块的上半部分（前blockCountV0行）写回cGlobalOffsetV0；
• GetSubBlockIdx() == 1（V1）：把结果块的下半部分写回cGlobalOffsetV1 = cGlobalOffsetV0 + blockCountV0 * tiling.N，即紧接上半部分之后的行；
• 三种尾块分支分别覆盖"仅 N 方向越界"、"仅 M 方向越界"、"M、N 双方向越界"，通过调整blockLen（N 方向剩余长度）与blockCount（M 方向剩余行数）保证127这类非对齐尺寸也能精确落位。

5.4 内核入口与 Matmul 对象注册

全局入口matmul_splitm_custom首先从 GM 拷贝 Tiling 数据，然后实例化具体类型的 Matmul 内核并注册：

__global__ __aicore__ void matmul_splitm_custom( __gm__ uint8_t* a, __gm__ uint8_t* b, __gm__ uint8_t* bias, __gm__ uint8_t* c, __kfc_workspace__ __gm__ uint8_t* workspace, __gm__ uint8_t* tilingGm) { TCubeTiling tiling; CopyTiling(&tiling, tilingGm); MatmulKernel<half, half, float, float> matmulKernel; AscendC::TPipe pipe; REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), matmulKernel.matmulObj, &tiling); // init matmul kernel, isTransA=false, isTransB=false matmulKernel.Init(a, b, bias, c, tiling, false, false, &pipe); // matmul kernel process matmulKernel.Process(); }

REGIST_MATMUL_OBJ负责把 Matmul 对象与TPipe、系统工作空间（__kfc_workspace__）绑定，使 Matmul 内部能够据此分配 L1/L0C/UB 资源并驱动计算流水。CopyTiling把 GM 中以 32 位字存储的TCubeTiling逐字拷贝到片上结构体，完成 Tiling 参数的传递。

六、Tiling 侧：MultiCoreMatmulTiling 生成切分参数

Host 侧的GenerateTiling使用matmul_tiling::MultiCoreMatmulTiling生成内核所需的全部切分参数。该 Tiling API 是 Matmul 多核场景的官方接口，其类定义与接口说明见 include/adv_api/matmul/bmm_tiling.h，底层实现位于 impl/adv_api/tiling/matmul/bmm_tiling.cpp。

样例中的完整调用序列：

matmul_tiling::MultiCoreMatmulTiling tilingApi(*ascendcPlatform); tilingApi.SetDim(ascendcPlatform->GetCoreNumAic()); // 用 AIC 核数作为 blockDim tilingApi.SetAType(leftPosition, leftFormat, leftDtype, isTransA); tilingApi.SetBType(rightPosition, rightFormat, rightDtype, isTransB); tilingApi.SetBiasType(biasPosition, biasFormat, biasDtype); tilingApi.SetCType(resultPosition, resultFormat, resultDtype); tilingApi.SetOrgShape(M, N, K); tilingApi.SetShape(M, N, K); tilingApi.EnableBias(isBias); tilingApi.SetBufferSpace(-1, -1, -1); int64_t res = tilingApi.GetTiling(tilingData);

参数语义（结合 Tiling 源码与接口注释）：

得到的TCubeTiling通过SaveToBuffer序列化到 GM 缓冲区，其中usedCoreNum字段被用作内核启动时的numBlocks：

uint32_t numBlocks = reinterpret_cast<TCubeTiling*>(tilingBuf)->usedCoreNum; matmul_splitm_custom<<<numBlocks, 0, stream>>>( aDevice, bDevice, biasDevice, cDevice, workspaceDevice, tilingDevice);

也就是说，核数完全由 Tiling 阶段算出的usedCoreNum决定，它等于 M、N 两个方向切分核数的乘积，与 5.2 节CalcOffset中的分核公式一一对应。

从 Tiling 源码结构还可以推断：多核切分时每个核的singleCoreM/singleCoreN/singleCoreK由 Tiling 算法在DoMultiCoreSplitMNTiling（见 impl/adv_api/tiling/matmul/matmul_tiling_algorithm.cpp）中计算并写入TCubeTiling，内核侧直接读取使用。若用户希望手工指定单核形状或对齐粒度，MultiCoreMatmulTiling还提供了SetSingleShape、SetSingleRange、SetAlignSplit等接口，本文样例采用默认自动推导，未显式调用。

七、Host 侧执行流程与 ACL 内存管理

main函数完整展示了 SplitM 样例的 Host 驱动流程，可作为此类直调样例的标准模板：

1. 获取平台信息：通过PlatformAscendCManager::GetInstance获取硬件平台实例；CPU 调试模式下按__NPU_ARCH__选择平台名（如Ascend950PR_9589对应 dav-3510）。
2. 计算文件与工作空间大小：cFileSize = M * N * sizeof(float)；workspaceSize由用户空间（本例为 0）与GetLibApiWorkSpaceSize()得到的系统空间组成。
3. 生成 Tiling：GenerateTiling填充tilingBuf，并读出usedCoreNum作为启动核数。
4. 初始化 ACL 运行时：aclInit → aclrtSetDevice(0) → aclrtCreateStream。
5. 数据搬运：为 A/B/Bias 分别执行aclrtMallocHost（Host 内存）、aclrtMalloc（Device 内存，ACL_MEM_MALLOC_HUGE_FIRST）、从./input/*.bin读入并aclrtMemcpy到 Device；C 与 Tiling、workspace 同样按需分配。
6. 启动内核：以numBlocks为块数、0 为 AIV 数（使用默认分离模式）通过<<< >>>语法启动，随后aclrtSynchronizeStream等待完成。
7. 回读结果：aclrtMemcpy把 C 结果拷回 Host，WriteFile写入./output/output.bin。
8. 资源释放：依次aclrtFree/aclrtFreeHost/aclrtDestroyStream/aclrtResetDevice/aclFinalize。

八、数据生成与精度校验机制

8.1 测试数据生成

scripts/gen_data.py 生成与内核完全一致的 127×63 输入：

x1_gm = np.random.uniform(-1, 1, [m, k]).astype(np.float16) x2_gm = np.random.uniform(-1, 1, [k, n]).astype(np.float16) bias_gm = np.random.uniform(-1, 1, [1, n]).astype(np.float32) golden = np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32)).astype(np.float32) + bias_gm.astype(np.float32)

值得注意的细节是：golden 计算时先把 FP16 的 A/B 提升为 FP32 再做np.matmul，再与 FP32 的 Bias 相加。这一方面保证参考结果的高精度，另一方面又会在结果中体现 FP16 输入量化带来的微小误差——因此校验脚本必须使用容差比较而非严格相等，这也正是 8.2 节相对/绝对误差容限存在的原因。

8.2 结果校验

scripts/verify_result.py 以np.isclose逐元素比较输出与 golden，关键容差参数为：

RELATIVE_TOL = 1e-3 # 相对误差 ABSOLUTE_TOL = 1e-3 # 绝对误差 ERROR_TOL = 1e-3 # 允许的全局错误率上限

校验逻辑为：计算所有不满足isclose(output, golden, rtol=1e-3, atol=1e-3)的元素索引，打印最多前 100 个差异点的期望值/实际值/相对偏差，并统计error_ratio = 差异元素数 / 总元素数；当error_ratio <= 1e-3时判定通过并输出test pass!，否则抛出[ERROR] result error并返回非零退出码。因此本文第三章节中看到的test pass!实际意味着：SplitM 多核切分、V0/V1 半块搬出、尾块处理以及 Bias 累加全部正确，且误差率在千分之一以内。

九、总结与进阶指引

matmul_splitm样例完整演示了 SplitM 策略在 AscendC 分离模式下的落地方式，核心要点可归纳为：

1. 策略入口：Matmul 对象最后一个模板参数传AscendC::Impl::Detail::SplitMMatmulPolicy，C 矩阵位置设为VECIN，二者缺一不可；
2. 硬件假设：面向 AIC:AIV=1:N 架构（本样例按 1:2 设计），同一 AIC 的两个 AIV 共享 A/B（SameAB），因此分核索引使用GetBlockIdx() / 2归并，搬出时用GetSubBlockIdx()区分 V0/V1；
3. 分核模型：Tiling 侧SetDim(GetCoreNumAic())+MultiCoreMatmulTiling自动计算singleCoreM/singleCoreN，内核侧CalcOffset按blockIdx % mSingleBlocks与blockIdx / mSingleBlocks解出 M/N 双维核号并计算各矩阵偏移；
4. 半块流水：缓冲区按baseM/2 × baseN初始化，Iterate → GetTensorC → DataCopy → DataCopyPad构成主循环，尾块通过SetTail与三支分支精细处理非对齐尺寸；
5. 验证闭环：gen_data.py生成随机 FP16 输入与 FP32 golden，verify_result.py以 1e-3 容差比对，输出test pass!确认正确性。

该样例与仓库中同目录的其他 Matmul 高级样例（如 matmul 基础直调、Bias 融合、SplitK 等，参见 04_advanced_api/00_matmul）相互对照，可帮助开发者理解不同切分策略的差异。若需继续深入，可重点阅读以下仓库源码：

• SplitM 策略定义：impl/adv_api/detail/matmul/policy/matmul_policy.h
• 双目标切 M 搬出实现：impl/adv_api/detail/matmul/stage/copy_cube_out/copy_cube_out_fixpipe.h
• 多核 Tiling API 声明与实现：include/adv_api/matmul/bmm_tiling.h、impl/adv_api/tiling/matmul/bmm_tiling.cpp
• 核号系统接口：impl/basic_api/dav_3510/kernel_operator_sys_var_impl.h

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend