asc-devkit 融合算子开发实战：基于 Ascend C 的 Matmul 与 LeakyRelu 融合计算全解析

原创2026-10-14 11:16:3795 阅读

文章标签：人工智能深度学习算子库CANNAscend

asc-devkit 融合算子开发实战：基于 Ascend C 的 Matmul 与 LeakyRelu 融合计算全解析

导读

本文以 asc-devkit（CANN 昇腾 AI 处理器算子开发语言仓库）中 examples/01_simd_cpp_api/00_introduction/03_fusion_operation 目录下的两个官方示例为核心，系统讲解在 Ascend C SIMD C++ API 编程模型下，如何将 Cube（矩阵计算）单元与 Vector（向量计算）单元融合起来，实现 Matmul 矩阵乘法与 LeakyRelu 激活函数的协同计算。读完本文，你将掌握两类编程范式的完整实现路径：基于 Matmul 高阶 API 的融合实现（自动封装数据搬运与 Cube 调度），以及基于静态 Tensor 编程模式 + 底层 ISASI 指令的手动融合实现（显式控制内存搬移、格式转换与流水同步），并了解 MIX 混合核编程的核间协同原理、编译运行流程与性能优化方向。

一、示例总览：两种实现一条主线

融合算子（Fusion Operator）是昇腾算子开发中的高频需求：矩阵乘法结果往往需要继续经过激活、归一化等逐元素运算才能交付下游使用。若按两个独立算子串行执行，矩阵结果需要多次往返全局内存（GM），引入大量额外搬运开销。本目录的两个示例将 Matmul（Cube 单元） 与 LeakyRelu（Vector 单元） 的计算融合进一个核函数，演示了 Cube 与 Vector 单元的协同计算编程模式。

目录名称	功能描述	支持产品
matmul_leakyrelu_advanced_api	基于 Matmul 高阶 API 实现 Matmul 矩阵乘法与 LeakyRelu 激活函数的融合计算	Ascend 950PR&950DT 产品 Atlas A3 产品 Atlas A2 产品
matmul_leakyrelu_basic_api	基于基础 API（静态 Tensor 编程模式）实现 Matmul 与 LeakyRelu 的融合计算	Ascend 950PR&950DT 产品 Atlas A3 产品 Atlas A2 产品

两个示例的核心差异在于抽象层次：

advanced_api 走 Matmul 高阶 API 路线：矩阵计算期间的数据移动、Cube 计算调度、基础流水同步全部由 matmul::Matmul 对象封装，开发者只需完成矩阵形状配置、tiling 生成、输入 Tensor 绑定、Vector 侧融合计算与结果写回。
basic_api 走静态 Tensor 编程模式：矩阵形状与内存布局在编译期确定，开发者手工编排 GM→L1→L0A/L0B→L0C→GM 的完整数据通路，显式调用 DataCopy、LoadData、Mmad、Fixpipe 等底层 ISASI 接口，并对 Cube 与 Vector 之间的核间同步负责。

无论哪种路线，两者的融合计算目标都是同一个公式（advanced 版本还多一步 Bias 加法）：

# advanced_api 版本：C = A * B + Bias；C = C > 0 ? C : C * 0.001
# basic_api 版本：C = A * B；C = C >= 0 ? C : C * 0.001

1.1 支持的软硬件环境

产品	CANN 软件版本
Ascend 950PR&950DT 产品	>= CANN 9.1.0
Atlas A3 产品	>= CANN 9.0.0
Atlas A2 产品	>= CANN 9.0.0

1.2 目录结构

两个示例的工程结构完全一致（以下以 basic_api 为例，advanced_api 仅文件名不同）：

matmul_leakyrelu_basic_api
├── scripts
│   ├── gen_data.py                 // 输入数据与 golden 数据生成脚本
│   └── verify_result.py            // golden 值比对脚本
├── CMakeLists.txt                  // 编译工程文件
├── data_utils.h                    // 数据读写函数
├── matmul_leakyrelu_basic_api.asc  // Ascend C 样例实现与调用示例
└── README.md                       // 样例文档

.asc 文件是 Ascend C 的算子源文件，同时包含 Device 侧核函数实现与 Host 侧调用代码（main 函数），使用 CMake 中的 ASC 语言支持编译，见 cmake/asc/asc_modules 下的 CMakeDetermineASCCompiler.cmake 等工具链定义。

二、进阶路线一：Matmul 高阶 API 融合实现（advanced_api）

2.1 样例规格与输入输出定义

该示例以 M = 512, N = 128, K = 128 为固定规格。代码中将矩阵 A 的 K 轴记为 Ka、矩阵 B 的 K 轴记为 Kb，本例中 Ka = Kb = K = 128。输入 A、B 为 half 类型，Bias 与输出 C 为 float 类型，格式均为 ND：

项	名称	Shape	数据类型	格式
算子类型（OpType）	MatmulLeakyRelu	-	-	-
输入	A	[M, K]	half	ND
输入	B	[K, N]	half	ND
输入	Bias	[N]	float	ND
输出	C	[M, N]	float	ND
核函数名	matmul_leakyrelu_custom	-	-	-

核函数以 __mix__(1, 2) 声明，Host 侧以 numBlocks = 1 启动 1 个 AI Core 组合（1 个 AIC + 2 个 AIV）。当前 tiling 只沿 N 轴分核切分为 2 份，M 轴与 K 轴均不分核，每个 AIV 逻辑核负责 M * singleCoreN = 512 * 64 的输出区域。

2.2 Host 侧 tiling 生成流程

tiling 生成位于 matmul_leakyrelu_advanced_api.asc 的 GenerateTiling 函数中，核心步骤如下：

创建 tiling 对象：实例化 matmul_tiling::MultiCoreMatmulTiling tilingApi(*ascendcPlatform)，用于生成多核 Matmul 所需的切分参数。
SetDim(2)：告知 tiling 侧配置 2 个 AIV 逻辑核参与 Matmul 计算启动。在 MIX 分离架构下，Matmul API 由 AIV 侧发起，因此 tiling 从 AIV 视角切分任务。
类型设置：SetAType/SetBType/SetCType/SetBiasType 分别设置 A、B、C、Bias 的数据源位置、数据格式与数据类型。其中 C 类型被设置为 VECIN，使得 Matmul 结果直接落到 Vector 侧可读取的位置，便于后续 LeakyRelu 直接消费中间结果。
形状设置：SetOrgShape(M, N, K) 设置原始完整矩阵形状；SetShape(M, N, K) 设置本次实际参与 Matmul 计算的大小（本例计算完整矩阵，两者一致）。
EnableBias(true)：使能 Bias 输入，让 Matmul 计算阶段完成 A * B + Bias。
SetFixSplit(128, 64, 128)：固定 Matmul 基础分块大小（baseM=128、baseN=64、baseK=128），保证 Kernel 侧 baseM/baseN 相关搬出逻辑与 tiling 结果一致。
SetBufferSpace(-1, -1, ubSize - usedUb)：设置 Matmul 可用的 L1/L0C/UB 空间。L1 与 L0C 传 -1 表示使用当前 AI 处理器对应缓冲的默认大小；UB 空间则先通过 ascendcPlatform->GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize) 查询可用 UB 大小，再减去 usedUb = baseM * baseN * sizeof(float)——这 baseM * baseN * 4 字节的 UB 缓存是为 LeakyRelu 预留的，避免 Matmul tiling 生成时把 Vector 融合计算所需的缓存全部占用。
GetTiling(tilingData)：生成最终 tiling 结果，返回值 -1 表示 tiling 计算失败，结果不可继续使用。

tiling 计算所依赖的平台信息由 platform_ascendc::PlatformAscendCManager::GetInstance(...) 获取，CPU 调试模式下会按 __NPU_ARCH__ 选择具体设备型号（如 Ascend910B1、Ascend950PR_9589），NPU 直调模式则使用默认实例。

2.3 Kernel 侧整体实现

核函数签名如下（摘自 matmul_leakyrelu_advanced_api.asc）：

__global__ __mix__(1, 2) void matmul_leakyrelu_custom(
    __gm__ uint8_t* a, __gm__ uint8_t* b, __gm__ uint8_t* bias, __gm__ uint8_t* c,
    __kfc_workspace__ __gm__ uint8_t* workspace, AscendC::tiling::TCubeTiling tiling)

实现要点逐条拆解：

参数约定：A、B、Bias、C 均以 __gm__ uint8_t* 字节地址传入，Kernel 内部将 A/B 固定为 half、C/Bias 固定为 float，再转换为对应的 GlobalTensor（GlobalTensor 只描述 GM 上的地址与元素个数，真实的数据搬运与 L1/L0 分块由 Matmul 高阶 API 结合 tiling 完成）。workspace 使用 __kfc_workspace__ 修饰符，表示 Host 侧分配并传入的系统 workspace，注册 Matmul 对象时通过 GetSysWorkSpacePtr() 传入。
Matmul 对象模板参数：matmul::Matmul<MatmulType<GM, ND, half>, MatmulType<GM, ND, half>, MatmulType<VECIN, ND, float>, MatmulType<GM, ND, float>> 依次描述 A、B、C、Bias 的位置、格式与数据类型——C 使用 VECIN 与 tiling 侧 SetCType 保持一致。
注册 Matmul 对象：REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), matmulObj, &tiling) 将 TPipe、系统 workspace、Matmul 对象与 TCubeTiling 一并交给高阶 API 内部使用。
核间偏移计算：用 GetBlockIdx() 取当前 AIV 侧逻辑任务号，nCoreOffset = blockIdx * tiling.singleCoreN 得到 N 方向分片起点。矩阵 A 在所有逻辑核上复用起始地址（每个逻辑块都需要完整 M、K 轴输入参与 K 轴累加），而矩阵 B、Bias、C 需要按 AIV 逻辑核号做 N 向偏移。
输入绑定：SetTensorA(aGlobal)、SetTensorB(bGlobal)、SetBias(biasGlobal) 绑定当前逻辑块需要读取的输入。
队列与缓存：TQue<TPosition::VECIN, 1> mmOutQueue 管理 Matmul 输出到 Vector 侧的中间结果缓存，TQue<TPosition::VECOUT, 1> reluOutQueue 管理 LeakyRelu 输出缓存；pipe.InitBuffer 为每个 baseM * baseN 输出块分配片上缓存。

2.4 Matmul 与 LeakyRelu 融合主循环

核心融合逻辑在 while (matmulObj.template Iterate<true>()) 循环中完成，每次 Iterate 产生一个 baseM * baseN = 128 * 64 输出块。由于 baseM=128、baseN=64 保持不变，每个 AIV 逻辑核需要 4 次 Iterate<true>() 才能完成自身 512 * 64 的输出区域。循环体内依次完成：

mmOutQueue.AllocTensor<float>() 从 Vector 输入队列申请缓存，承接 Matmul 输出块。
matmulObj.template GetTensorC<true>(mmOutLocal, false, true) 从 Matmul 对象取回当前块的 C 中间结果（第二参 false 表示不使能原子累加，第三参 true 表示按连续地址写入）。由于 C 类型为 VECIN，该中间结果可直接作为 Vector 侧 LeakyRelu 的输入。
AscendC::LeakyRelu(reluOutLocal, mmOutLocal, static_cast<float>(0.001), tiling.baseM * tiling.baseN) 对当前块执行逐元素激活（负值乘 0.001），结果写入独立输出缓存 reluOutLocal。
mmOutQueue.EnQue/DeQue 建立 Matmul 结果写入与后续 LeakyRelu 读取之间的数据依赖；reluOutQueue.EnQue/DeQue 建立 Vector 计算与搬出之间的数据依赖。
AscendC::DataCopy(cGlobal[cOffset], reluOutLocal, AscendC::DataCopyParams{...}) 将激活后的块写回 GM 输出矩阵对应位置。
matmulObj.End() 通知 Matmul 对象当前核计算结束，释放高阶 API 内部状态。

2.5 分核与地址偏移详解

0 号 AIV 逻辑核处理矩阵 C 第 0~63 列，1 号 AIV 逻辑核处理第 64~127 列。
矩阵 A [512, 128]：N 向分核时每个逻辑块都需要完整 M、K 轴输入，因此无分核偏移。
矩阵 B [128, 128]（代码中对应 [Kb, N]）：起始偏移为 blockIdx * tiling.singleCoreN，指向当前 N 向分片的起始列。
Bias [128]：对应 N 方向的 128 列，起始偏移同样是 blockIdx * tiling.singleCoreN。
矩阵 C [512, 128]：每行长度为 tiling.N = 128，逻辑块起始偏移为 blockIdx * tiling.singleCoreN。
每个逻辑块内部含 singleCoreM / baseM * singleCoreN / baseN = 4 * 1 = 4 个 baseM * baseN 输出块。由于 baseN 等于 singleCoreN，单核内部只沿 M 方向遍历 base 块；写回时按当前 Iterate 轮次计算行偏移 cOffset = computeRound * baseM * tiling.N（见 matmul_leakyrelu_advanced_api.asc）。
M 轴与 N 轴均能被对应分块整除，因此代码中不含尾块处理逻辑。

2.6 调用实现

Host 侧使用核函数调用运算符 <<<>>>（matmul_leakyrelu_custom<<<numBlocks, 0, stream>>>(...)）启动核函数，运行期参数依次为：Device 侧矩阵 A 地址、矩阵 B 地址、Bias 地址、矩阵 C 地址、workspace 地址、Host 侧生成的 tiling 数据。numBlocks = 1 表示启动 1 个 AI Core 组合，对应 2 个 AIV 侧逻辑任务，与 SetDim(2) 配置的 2 个 AIV 逻辑核精确匹配。

2.7 MIX 场景下的 SetDim 与启动核数配置

在 MIX 场景（同时包含矩阵计算与向量计算）的分离模式下（本示例支持的产品基本架构均为分离模式），Matmul API 从 AIV 侧发起：AIV 执行 Iterate 时通知 AIC 执行矩阵计算，AIC 完成后通知 AIV 取结果并继续 Vector 侧融合计算。因此：

tiling 侧从 AIV 视角切分任务：SetDim(2) 生成面向 2 个 AIV 逻辑核的 Matmul 分核参数；
__mix__(1, 2) 表示每个 AI Core 组合包含 1 个 AIC 与 2 个 AIV；
本示例为 Kernel 直调（direct call）场景，Host 侧 numBlocks = 1 启动 1 个 AI Core 组合，对应 2 个 AIV 逻辑任务，正好匹配 SetDim(2)。

三、进阶路线二：静态 Tensor 编程模式实现（basic_api）

3.1 样例规格与核心分布逻辑

该示例固定 M = 512, K = 128, N = 128，调用 2 个 Cube 核 + 4 个 Vector 核完成计算：

项	名称	Shape	数据类型	格式
算子类型（OpType）	Matmul+LeakyRelu Fusion	-	-	-
输入	A（左矩阵）	[512, 128]	half	ND
输入	B（右矩阵）	[128, 128]	half	ND
输出	C	[512, 128]	half	ND
核函数名	mmad_vec_custom	-	-	-

核心分布逻辑：输出矩阵 C 沿 M 方向（第一维，行方向）划分为 2 块，N 方向（第二维，列方向）不划分：

         K                N
     ┌─────────┐      ┌──────────┐
     │         │      │          │
   M │    A    │  × K │    B     │  =  C (M × N)
     │         │      │          │
     └─────────┘      └──────────┘

M 方向分块数：M / singleCoreM = 512 / 256 = 2
N 方向分块数：N / singleCoreN = 128 / 128 = 1
Cube 核总数：2 × 1 = 2
Vector 核总数：2 × 2 = 4（每个 Cube 核对应 2 个 Vector 核）

每个 Cube 核产生一个 baseM × baseN = 256 × 128 结果块，交由 2 个 Vector 核各自处理 baseM/2 × baseN = 128 × 128 的数据。

3.2 昇腾 AI Core 内存层级与数据通路

理解 basic_api 示例必须先建立昇腾 AI Core 的内存层级与数据通路认知（该示例文档对此有系统说明）：

GM（Global Memory）：芯片外部 DDR，存放输入/输出数据，容量最大但访问延迟最高。
L1：片上缓存（主要存放 Cube 侧的 A/B 矩阵数据），由 MTE2 引擎从 GM 搬入。
L0A/L0B：Cube 矩阵计算单元的输入缓存，由 MTE1 引擎从 L1 搬入。
L0C：Cube 矩阵计算单元的输出缓存，存放 Mmad 计算结果。
UB（Unified Buffer，对应 TPosition::VECCALC）：Vector 侧的片上缓存，存放 Vector 计算所需的输入/输出数据。

数据在搬运过程中需要格式转换（ND→Nz/Zn），因为 Cube 计算单元要求 Fractal（分形）布局而非 ND 布局。

同步机制：

核内同步：SetFlag/WaitFlag 用于同一核内不同流水引擎（MTE2、MTE1、M、V、FIX、MTE3）之间的依赖同步，例如 MTE2 搬移完成后通知 MTE1 才能开始搬运。
核间同步：CrossCoreSetFlag/CrossCoreWaitFlag 用于不同核之间的依赖同步，例如 Cube 核计算完成后通知 Vector 核才能开始读取数据。

3.3 Cube 侧数据流与实现代码

Cube 侧数据流（数据以 A 矩阵为例，B 类似）：

GM(A:ND,half) -> L1(A:Nz,half) -> L0A(A:Nz,half) -
            │                  │                 │
          DataCopy            LoadData            │
          ND->Nz              Nz->Zz/Nz           │
                                                  │--->L0C(Nz,float) -> GM(C:ND,half)
                                                  │  │                  │
                                                  │ Mmad             Fixpipe
                                                  │ C=A×B       Nz->ND, float32->half
                                                  │
GM(B:ND,half) -> L1(B:Nz,half) -> L0B(B:Zn,half) -
            │                  │
          DataCopy            LoadData
          ND->Nz              Nz->Zn(transpose)

注意：L0A 的 Fractal 格式在不同产品上有差异——Ascend 950PR&950DT 产品的 L0A Fractal 格式为 Nz，Atlas A2/A3 系列产品为 Zz。这直接影响 LoadData 的参数写法，代码中用 __NPU_ARCH__ == 2201（对应 A2/A3）与 __NPU_ARCH__ == 3510（对应 950PR）分支处理。

核函数与类骨架（摘自 matmul_leakyrelu_basic_api.asc）：

class KernelMatmul {
public:
    __aicore__ inline void Init(__gm__ uint8_t* xMatrix, __gm__ uint8_t* yMatrix, __gm__ uint8_t* zMatrix)
    {
        // 设置 GlobalTensor 起始地址，指向 GM 中对应输入/输出矩阵
        aGM.SetGlobalBuffer((__gm__ half*)xMatrix);
        bGM.SetGlobalBuffer((__gm__ half*)yMatrix);
        cGM.SetGlobalBuffer((__gm__ half*)zMatrix);

        // Cube 核沿 M 方向分核，每个 Cube 核负责 singleCoreM * N 输出块
        if ASCEND_IS_AIC {
            aGM = aGM[AscendC::GetBlockIdx() * singleCoreM * K];  // 偏移到当前核负责的 A 矩阵行
            cGM = cGM[AscendC::GetBlockIdx() * singleCoreM * N];  // 偏移到当前核负责的 C 矩阵行
        }
        // Vector 核数是 Cube 核数的 2 倍，1 个 Cube 核的结果由 2 个 Vector 核处理
        // GetBlockIdx() / 2 将 Vector 核索引映射回对应 Cube 核输出块起始地址
        if ASCEND_IS_AIV {
            cGM = cGM[AscendC::GetBlockIdx() / 2 * singleCoreM * N];
        }
    }
    // ...
};

Cube 侧完整计算流程（Process() 中 if ASCEND_IS_AIC 分支）：

LocalTensor 创建：使用 LocalMemAllocator（片上内存分配器，按申请顺序自动分配，避免手动维护地址偏移）为各片上缓存创建 LocalTensor。L1 中 A/B 矩阵临时空间由同一个 L1 分配器按申请顺序分配。
GM → L1：调用 DataCopy（GM 到 L1，ND→Nz），搬运同时完成 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});
// 核内同步：MTE2 搬运完成后通知 MTE1 可开始搬运
AscendC::SetFlag<AscendC::HardEvent::MTE2_MTE1>(EVENT_ID0);
AscendC::WaitFlag<AscendC::HardEvent::MTE2_MTE1>(EVENT_ID0);

L1 → L0A/L0B：调用 LoadData（2D），B 矩阵需要转置（Nz→Zn）。A2/A3（2201）架构 L0A 为 Zz 布局，需逐块搬运；950PR（3510）架构 L0A 为 Nz 布局，可一次性完成：

#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 2201)
    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});
    }
    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});
    }
#elif defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510)
    AscendC::LoadData(a2Local, a1Local,
        AscendC::LoadData2DParamsV2{0, 0, baseM / CUBE_BLOCK, baseK / CUBE_BLOCK,
            baseM / CUBE_BLOCK, baseM / CUBE_BLOCK, false, 0});
    AscendC::LoadData(b2Local, b1Local,
        AscendC::LoadData2DParamsV2{0, 0, baseK / CUBE_BLOCK, baseN / CUBE_BLOCK,
            baseK / CUBE_BLOCK, baseN / CUBE_BLOCK, true, 0});
#endif
AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(EVENT_ID0);
AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(EVENT_ID0);

L0A/L0B → L0C：调用 Mmad（矩阵乘加指令），沿 K 方向累加所有数据块：

AscendC::Mmad(cLocal, a2Local, b2Local,
    AscendC::MmadParams{baseM, baseN, baseK, 0, false, true});
// 注意：cmatrixInitVal=true 表示首次计算时把 L0C 累加器清零

L0C → GM：调用 Fixpipe（L0C 到 GM），搬运同时完成 Nz→ND 格式转换与 float32→half 精度转换：

AscendC::SetFlag<AscendC::HardEvent::M_FIX>(EVENT_ID0);
AscendC::WaitFlag<AscendC::HardEvent::M_FIX>(EVENT_ID0);
AscendC::Fixpipe(cGM, cLocal,
    AscendC::FixpipeParamsV220{baseN, baseM, baseM, N, false, QuantMode_t::F322F16, 0, 1, 0, 0, 0});

核间同步：AscendC::CrossCoreSetFlag<0x2, PIPE_FIX>(0) 通知 Vector 核数据已就绪。

3.4 Vector 侧数据流与实现代码

Vector 侧数据流（1 个 Cube 核的结果由 2 个 Vector 核处理）：

Cube core:  Fixpipe -> GM(C:ND,half)
                    │
                    │ CrossCoreSetFlag / CrossCoreWaitFlag
                    ▼
Vector core: GM(C:ND,half) -> UB(VECCALC,half) -> UB(VECCALC,half) -> GM(C:ND,half)
                      │                   │                   │
                  DataCopyPad          LeakyRelu          DataCopyPad
            MTE2 transfer(baseM/2×baseN)  VEC compute       MTE3 write out

Process() 中 if ASCEND_IS_AIV 分支与私有方法 LeakyReluCompute 完成 Vector 侧工作：

核间同步：AscendC::CrossCoreWaitFlag<0x2>(0) 阻塞等待 Cube 核完成 Fixpipe 写回，确保读到完整的 Matmul 结果。
UB 分配：用 UB 分配器在 VECCALC 位置创建 LocalTensor，每个 Vector 核处理 baseM / 2 * baseN 个元素。
GM → UB：调用 DataCopyPad（GM 到 UB）搬运半块结果。gmOffset = GetBlockIdx() % 2 * (baseM / 2 * N) 用 GetBlockIdx() % 2 区分同一逻辑核内 2 个 AIV 各自处理的半块区域：

uint32_t blockLen = baseN * sizeof(half);
uint32_t srcStride = (N - baseN) * sizeof(half);  // GM 中相邻行之间的跳变步长
AscendC::DataCopyPad<half>(
    vecLocal, cGM[gmOffset],
    AscendC::DataCopyExtParams{static_cast<uint16_t>(baseM / 2), blockLen, srcStride, 0, 0},
    AscendC::DataCopyPadExtParams<half>{true, 0, 0, 0});
AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);

UB 计算：调用 LeakyRelu 执行激活，负值乘 0.001：

AscendC::LeakyRelu(vecLocal, vecLocal, (half)0.001, baseM / 2 * baseN);
AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);

UB → GM：调用 DataCopyPad（UB 到 GM）写回激活结果，完成融合计算。

3.5 核函数声明与调用

核函数使用 __mix__(1, 2) 实现 Cube 与 Vector 核的协同调用，参数 (1, 2) 表示每个逻辑核由 1 个 AIC（Cube 核）+ 2 个 AIV（Vector 核） 组成，Cube:Vector 核数比为 1:2：

// __global__：核函数声明修饰符
// __mix__(1,2)：混合核函数，1 个 AIC(Cube) + 2 个 AIV(Vector) 组成一个逻辑核
__global__ __mix__(1, 2) void mmad_vec_custom(__gm__ uint8_t* xMatrix, __gm__ uint8_t* yMatrix, __gm__ uint8_t* zMatrix)
{
    AscendC::InitSocState();  // 初始化硬件状态
    KernelMatmul<M, K, N, singleCoreM, singleCoreK, singleCoreN,
                 baseM, baseK, baseN> op;
    op.Init(xMatrix, yMatrix, zMatrix);   // GlobalTensor 初始化与分核偏移
    op.Process();        // AIC 侧执行 Matmul，AIV 侧执行 LeakyRelu
    AscendC::PipeBarrier<PIPE_ALL>();  // 等待所有流水阶段完成
}

Host 侧启动时 numBlocks = (M / singleCoreM) * (N / singleCoreN) = 2，模板实参依次为 M, K, N, singleCoreM, singleCoreK, singleCoreN, baseM, baseK, baseN（见 matmul_leakyrelu_basic_api.asc）。

3.6 约束条件

静态 Tensor 融合实现必须满足以下约束（源自示例文档与源码结构）：

baseM/baseK/baseN 满足 16 对齐；
M/N 能被 singleCoreM/singleCoreN 整除；
singleCoreM/singleCoreN 能被 baseM/baseN 整除，K 能被 baseK 整除，不支持非整除场景；
Vector 核数为 Cube 核数的 2 倍（与 __mix__(1, 2) 的 1:2 配置匹配）。

四、核心接口参数详解

两个示例都以 {} 花括号方式传参。以下结构体字段说明与源码参数一一对应，字段顺序与 API 文档一致。

4.1 Nd2NzParams（DataCopy：GM→L1 格式转换）

用于 DataCopy 接口描述 ND→Nz 格式转换参数（详见 DataCopy_GMToL1_ND2NZ）：

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]
};

例如 {1, baseM, baseK, 0, K, baseM, 1, 0} 将 baseM×baseK 的 ND 数据转换为 Nz 格式。

4.2 LoadData2DParams（LoadData：L1→L0A/L0B，A2/A3）

用于 Atlas A2 产品与 Atlas A3 产品（详见 LoadData_2D）：

struct LoadData2DParams {
    int32_t startIndex;   // 分形矩阵 ID（0 为第一个），单位：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}；搬运 B 矩阵需置 ifTranspose=true 完成 Nz→Zn 转置搬运。

4.3 LoadData2DParamsV2（LoadData：L1→L0A/L0B，Ascend 950PR&950DT）

用于 Ascend 950PR&950DT 产品（详见 LoadData_2D_V2）：

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
};

950PR&950DT 产品 L0A 布局为 Nz：A 矩阵使用 {0, 0, baseM / CUBE_BLOCK, baseK / CUBE_BLOCK, baseM / CUBE_BLOCK, baseM / CUBE_BLOCK, false, 0} 一步完成 Nz→Nz 搬运；B 矩阵使用 {0, 0, baseK / CUBE_BLOCK, baseN / CUBE_BLOCK, baseK / CUBE_BLOCK, baseN / CUBE_BLOCK, true, 0} 一步完成 Nz→Zn 转置搬运。

4.4 MmadParams（Mmad：矩阵乘加）

用于 Mmad 接口描述矩阵乘法参数（详见 Mmad）：

struct MmadParams {
    uint16_t m;               // 左矩阵高度（M 维），[0, 4095]
    uint16_t n;               // 右矩阵宽度（N 维），[0, 4095]
    uint16_t k;               // 左矩阵宽度/右矩阵高度（K 维），[0, 4095]
    uint16_t unitFlag;        // Mmad 与 Fixpipe 的细粒度并行控制，默认 0
    bool     cmatrixSource;   // C 矩阵初始值来源，false=CO1，true=C2，默认 false
    bool     cmatrixInitVal;  // C 矩阵初始值是否为 0，默认 true
};

例如 {baseM, baseN, baseK, 0, false, true} 计算 baseM×baseN 输出块，并沿 K 方向按 baseK 长度累加。

4.5 FixpipeParamsV220（Fixpipe：L0C→GM）

用于 Fixpipe 接口描述 L0C 到 GM 的搬运与精度转换参数（详见 Fixpipe_L0CToGM）：

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 矩阵每行元素个数，单位：元素
    bool        reluEn;       // 是否使能 ReLU
    QuantMode_t quantPre;     // 量化模式，F322F16 表示 float 转 half
    uint64_t    deqScalar;    // 标量量化参数，单一缩放值
    int32_t     ndNum;        // 源 Nz 矩阵数量，[1, 65535]
    int32_t     srcNdStride;  // 不同 Nz 矩阵起始地址间隔，单位：16×C0_SIZE，[1, 512]
    int32_t     dstNdStride;  // 相邻目的 ND 矩阵之间的偏移，单位：元素，[1, 65535]
    int32_t     unitFlag;     // Mmad 与 Fixpipe 的并行控制
};

例如 {baseN, baseM, baseM, N, false, F322F16, 0, 1, 0, 0, 0} 将 L0C 中 baseM×baseN 的 float32 结果转换为 half 并写回 GM。

4.6 DataCopyParams（DataCopy：UB→GM 搬出）

用于 DataCopy 接口描述 Local Memory 到 Global Memory 的搬运参数（advanced_api 搬出矩阵 C 分块时使用）：

struct DataCopyParams {
    uint16_t blockCount;  // 连续搬运的数据块数量，[1, 4095]
    uint16_t blockLen;    // 每个连续数据块的长度，单位：DataBlock（32B），[1, 65535]
    uint16_t srcGap;      // 源操作数相邻连续数据块之间的间隔，单位：DataBlock（32B）
    uint16_t dstGap;      // 目的操作数相邻连续数据块之间的间隔，单位：DataBlock（32B）
};

advanced_api 搬出矩阵 C 时使用 {tiling.baseM, tiling.baseN * sizeof(float) / AscendC::DEFAULT_C0_SIZE, 0, (tiling.N - tiling.baseN) * sizeof(float) / AscendC::DEFAULT_C0_SIZE}：blockCount 表示搬出 baseM 行，blockLen 表示每行对应 baseN 个矩阵 C 元素的数据量，dstGap 表示目的端相邻行之间的间隔（单位 DataBlock），用于把 baseM * baseN 块写回完整 M * N 矩阵的对应列区间。

4.7 DataCopyExtParams 与 DataCopyPadExtParams（DataCopyPad：GM↔UB）

用于 DataCopyPad 接口描述 GM 与 UB 之间的块搬运参数（详见 DataCopyPad_GMToUB）：

struct DataCopyExtParams {
    uint16_t blockCount;  // 连续搬运的数据块数量
    uint32_t blockLen;    // 每个数据块的长度，单位：Byte
    uint32_t srcStride;   // 源端相邻数据块之间的间隔，单位：Byte
    uint32_t dstStride;   // 目的端相邻数据块之间的间隔，单位：Byte
    uint32_t rsv;         // 预留字段，填 0
};

template <typename T>
struct DataCopyPadExtParams {
    bool isPad;       // 是否使能 padding
    uint8_t leftPad;  // 左侧 padding 长度
    uint8_t rightPad; // 右侧 padding 长度
    T paddingValue;   // padding 填充值
};

例如 GM→UB 时 {static_cast<uint16_t>(baseM / 2), blockLen, srcStride, 0, 0} 表示每个 Vector 核读取 baseM/2 行结果；{true, 0, 0, 0} 表示使能 padding 能力但左右均不额外填充。

五、编译运行与验证

5.1 配置环境变量

按照 CANN 开发套件安装方式配置当前环境：

source ${install_path}/cann/set_env.sh

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

5.2 构建与执行

在两个示例根目录下执行（以 basic_api 为例）：

mkdir -p build && cd build;                                               # 创建并进入 build 目录
cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j;                      # 编译工程（默认 npu 模式）
python3 ../scripts/gen_data.py                                            # 生成测试输入数据
./demo                                                                    # 执行编译产物运行样例
python3 ../scripts/verify_result.py output/output.bin output/golden.bin   # 校验输出正确性

gen_data.py 会生成 input/x1_gm.bin（矩阵 A）、input/x2_gm.bin（矩阵 B），advanced_api 还会生成 input/bias.bin，并分别按对应精度在 output/golden.bin 中写入 golden 参考值（advanced 版本 golden 为 float32 且包含 Bias 加法，basic 版本为 float16）。verify_result.py 使用 np.isclose 逐元素比对，输出错误比例并在比例不超过 ERROR_TOL = 1e-4 时打印 test pass!。

CPU 调试或 NPU 仿真模式需追加 -DCMAKE_ASC_RUN_MODE 参数：

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

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

5.3 编译选项说明

选项	可选值	说明
`CMAKE_ASC_RUN_MODE`	`npu`（默认）、`cpu`、`sim`	运行模式：NPU 运行、CPU 调试、NPU 仿真
`CMAKE_ASC_ARCHITECTURES`	`dav-2201`（默认）、`dav-3510`	NPU 架构：dav-2201 对应 Atlas A2 产品与 Atlas A3 产品，dav-3510 对应 Ascend 950PR&950DT 产品

注意：advanced_api 的 CMakeLists.txt 中，当 CMAKE_ASC_ARCHITECTURES=dav-3510 时构建工程会自动为 ASC 编译追加 -DENABLE_CV_COMM_VIA_SSBUF=true 宏以启用 SSBUF 通信，dav-2201 场景不会追加该编译宏。

5.4 执行结果

精度比对通过时输出：

test pass!

六、实现流程逐阶段分析（basic_api）

以下两张表按时间顺序拆解 Cube 核与 Vector 核的每一步操作及其目的。

6.1 Cube 核流程

阶段	数据流/行为	实现目的/原因
初始化	使用 LocalMemAllocator 为 L1、L0A/L0B、L0C 等片上缓存分配 LocalTensor	按申请顺序自动分配片上内存，避免手动维护地址偏移
GM → L1	DataCopy 将 A/B 矩阵从 GM 搬运到 L1，同时完成 ND→Nz 格式转换	Cube 计算单元要求 Nz 分形布局，搬运时转换可避免额外转换开销
L1 → L0A/L0B	LoadData 将 A 矩阵从 L1 搬运到 L0A（Nz→Zz/Nz），B 矩阵从 L1 搬运到 L0B（Nz→Zn 转置）	Mmad 指令要求 B 矩阵以 Zn（转置 Nz）格式输入，因此 B 需要转置
L0A/L0B → L0C	Mmad 执行矩阵乘加，沿 K 方向累加所有数据块	完成 A×B 矩阵乘法计算，K 向分块累加保证正确性
核内同步	SetFlag/WaitFlag 确保数据搬运完成后再启动下一步操作	避免 LoadData 读取未搬运完成的数据，避免 Mmad 读取未装载完成的数据
L0C → GM	Fixpipe 将 L0C 结果搬出到 GM，同时完成 Nz→ND 格式转换与 fp32→fp16 精度转换	输出结果需回到 GM 供 Vector 核读取；格式需转为 ND，精度需由 fp32 降为 fp16 匹配输出要求
核间同步	CrossCoreSetFlag 通知 Vector 核数据就绪	确保 Vector 核在 Fixpipe 完成前不会开始读取 GM 数据

6.2 Vector 核流程

阶段	数据流/行为	实现目的/原因
核间同步	CrossCoreWaitFlag 等待 Cube 核 Fixpipe 完成	阻塞至 Cube 核通知数据就绪，保证读到完整 Matmul 结果
初始化	使用 UB 分配器在 `VECCALC` 位置分配 LocalTensor	为 GM→UB 搬运与 Vector 计算分配 UB 缓冲
GM → UB	DataCopyPad 将 Matmul 结果从 GM 搬运到 UB，每个 Vector 核读取 baseM/2 行、每行 baseN 个元素	Vector 核从 GM 读取 Cube 核写回的 Matmul 结果，各处理半块数据
UB 计算	LeakyRelu 执行激活计算，负值乘以 0.001	对 Matmul 结果施加 LeakyRelu 激活函数，完成融合计算
UB → GM	DataCopyPad 将 LeakyRelu 结果写回 GM	输出最终计算结果到 GM 供后续使用

七、优化方向分析

两个示例文档结合源码指出了当前实现可优化的方向，可直接指导读者在真实算子工程中做性能迭代：

优化方向	当前实现的问题	预期优化收益
L1/L0 双缓冲流水	Cube 核串行执行 GM→L1→L0A/L0B→L0C→GM，搬运与计算无法重叠	Ping-Pong 双缓冲使 MTE2 搬运与 MTE1/Mmad 计算并行，显著提升吞吐
Fixpipe 与 Mmad 细粒度并行	当前 Mmad 与 Fixpipe 块级串行，Mmad 计算期间 Fixpipe 空闲，Fixpipe 搬出期间 Mmad 空闲	拆成更细粒度块，Mmad 与 Fixpipe 交替执行，减少流水气泡
UB 双缓冲	Vector 核串行执行 GM→UB→计算→GM	UB Ping-Pong 双缓冲使 MTE2 搬运与 VEC 计算并行
Vector 核 VF 融合	当前使用 MemBase API，中间结果需写回 UB	使用 RegBase + asc_vf_call 做 VF 融合，中间计算在寄存器内完成，减少 UB 读写次数
大块搬运	当前仅搬运单块 baseM×baseN 数据	增大 singleCoreM/singleCoreN，减少搬运次数，提升带宽利用率
GM 搬运优化	Cube 核结果需经 GM 中转给 Vector 核（L0C→GM→UB），两次搬运开销高	Ascend 950PR 支持 Fixpipe 直写 UB（L0C→UB），省去 GM 中转开销

八、功能调试与性能分析

8.1 printf 调试

printf 接口提供 CPU 域/NPU 域调试场景下的格式化输出能力。在算子 Kernel 侧实现中需要输出日志信息的位置调用：

AscendC::printf("matmul blockIdx=%d\n", AscendC::GetBlockIdx());

注意：printf 打印功能会对算子实际性能产生一定影响，通常用于调试阶段；开发者可按需设置 ASCENDC_DUMP=0 关闭打印功能。

8.2 DumpTensor 调试

对于基于算子工程开发的算子，可用 DumpTensor 接口 Dump 指定 LocalTensor 的内容，也支持打印自定义附加信息（仅支持 uint32_t 类型信息），例如打印当前行号。在 Kernel 侧需要打印 Tensor 数据的位置调用：

// Vector 核 LeakyRelu 结果
AscendC::LeakyRelu(zLocal, xLocal, ..., len);
AscendC::DumpTensor(zLocal, 1, 16);

注意：DumpTensor 打印功能同样影响算子实际性能，通常用于调试阶段，可按需设置 ASCENDC_DUMP=0 关闭。

8.3 msOpProf 单算子性能分析

msOpProf 是单算子性能分析工具，提供 msopprof 与 msopprof simulator 两种使用模式，帮助用户定位算子内存、代码、指令层面的异常，支持设备运行与仿真等不同运行模式、可执行文件与算子二进制 .o 文件等文件类型的性能数据采集与自动解析。

设备侧性能采集：直接在昇腾 AI 处理器上测量算子执行时间，适用于快速定位设备环境下的算子性能问题。对 demo 可执行文件执行：

msopprof ./demo

命令完成后默认目录会生成 OPPROF_{timestamp}_XXX 文件夹，性能数据目录结构如下：

├── dump                       # 原始性能数据，用户无需查看
├── ArithmeticUtilization.csv  # Cube/Vector 指令周期占比
├── L2Cache.csv                # L2 Cache 命中率；影响 MTE2，合理规划搬运逻辑提升命中率
├── Memory.csv                 # UB、L1、主存读写带宽率
├── MemoryL0.csv               # L0A、L0B、L0C 读写带宽率
├── MemoryUB.csv               # Vector 与 Scalar 到 UB 的读写带宽率
├── OpBasicInfo.csv            # 算子基本信息
├── PipeUtilization.csv        # 计算与搬运单元的耗时与占比
├── ResourceConflictRatio.csv  # 全部指令中 UB bank 组、bank 冲突、资源冲突占比
└── visualize_data.bin         # MindStudio Insight 展示文件

查看详细性能分析结果：

# 查看 Task Duration 等指标
cat ./OPPROF_*/PipeUtilization.csv

九、两种实现路线对比与选型建议

对比维度	Matmul 高阶 API（advanced_api）	静态 Tensor 编程（basic_api）
编程抽象	数据搬运、Cube 调度、基础流水同步由 `matmul::Matmul` 封装	手工编排 GM→L1→L0A/L0B→L0C→GM 全链路与全部流水同步
开发工作量	低，聚焦 tiling 配置、Tensor 绑定与 Vector 侧融合	高，需显式处理格式转换、架构差异（Zz/Nz）、核间同步
灵活性	受高阶 API 能力边界约束	可精细控制每一级缓存的搬运与调度
跨架构适配	tiling 侧统一配置，适配成本低	需按 `__NPU_ARCH__` 分支处理 LoadData 参数（2201/3510）
典型定位	快速实现融合算子、业务迭代期	极致性能调优、深入理解硬件流水时

选型建议：业务功能验证与快速交付优先选择 Matmul 高阶 API；当性能瓶颈集中在数据搬运与流水重叠时，可参照 basic_api 的手动编排方式做精细化控制，并结合第七节的优化方向（双缓冲、Fixpipe 直写 UB、VF 融合等）逐项迭代。

十、仓库证据索引

以下为本篇文章引用到的仓库核心文件，便于读者继续深入：

示例索引与实现：03_fusion_operation 目录、advanced_api 内核源码、basic_api 内核源码
构建配置：advanced_api CMakeLists.txt
数据脚本：advanced gen_data.py、basic gen_data.py、verify_result.py
高阶 API 实现：Matmul 相关头文件、tiling 实现
API 参考文档：LeakyRelu、Mmad、Fixpipe_L0CToGM、DataCopy_GMToL1_ND2NZ、LoadData_2D、LoadData_2D_V2、DataCopyPad_GMToUB、DataCopyPad_UBToGM、GlobalTensor、LocalTensor_intro、LocalMemAllocator、GetBlockIdx、SetFlag_WaitFlag_ISASI、CrossCoreSetFlag_ISASI、CrossCoreWaitFlag_ISASI、asc_vf_call、DumpTensor

登录后查看全文

人工智能深度学习算子库CANNAscend

asc-devkit

asc-devkit 融合算子开发实战：基于 Ascend C 的 Matmul 与 LeakyRelu 融合计算全解析

asc-devkit 融合算子开发实战：基于 Ascend C 的 Matmul 与 LeakyRelu 融合计算全解析

导读

一、示例总览：两种实现一条主线

1.1 支持的软硬件环境

1.2 目录结构

二、进阶路线一：Matmul 高阶 API 融合实现（advanced_api）

2.1 样例规格与输入输出定义

2.2 Host 侧 tiling 生成流程

2.3 Kernel 侧整体实现

2.4 Matmul 与 LeakyRelu 融合主循环

2.5 分核与地址偏移详解

2.6 调用实现

2.7 MIX 场景下的 SetDim 与启动核数配置

三、进阶路线二：静态 Tensor 编程模式实现（basic_api）

3.1 样例规格与核心分布逻辑

3.2 昇腾 AI Core 内存层级与数据通路

3.3 Cube 侧数据流与实现代码

3.4 Vector 侧数据流与实现代码

3.5 核函数声明与调用

3.6 约束条件

四、核心接口参数详解

4.1 Nd2NzParams（DataCopy：GM→L1 格式转换）

4.2 LoadData2DParams（LoadData：L1→L0A/L0B，A2/A3）

4.3 LoadData2DParamsV2（LoadData：L1→L0A/L0B，Ascend 950PR&950DT）

4.4 MmadParams（Mmad：矩阵乘加）

4.5 FixpipeParamsV220（Fixpipe：L0C→GM）

4.6 DataCopyParams（DataCopy：UB→GM 搬出）

4.7 DataCopyExtParams 与 DataCopyPadExtParams（DataCopyPad：GM↔UB）

五、编译运行与验证

5.1 配置环境变量

5.2 构建与执行

5.3 编译选项说明

5.4 执行结果

六、实现流程逐阶段分析（basic_api）

6.1 Cube 核流程

6.2 Vector 核流程

七、优化方向分析

八、功能调试与性能分析

8.1 printf 调试

8.2 DumpTensor 调试

8.3 msOpProf 单算子性能分析

九、两种实现路线对比与选型建议

十、仓库证据索引

相关推荐

热门文章