昇腾算子模板库开发教程

作者：陆璐课题组，瑾丞

目 录

    1. 开发环境准备与基础概念
    1. 昇腾算子模板库（CATLASS）架构解析
    1. 基于模板库开发GEMM算子全流程
    1. GEMV与量化算子开发
    1. Epilogue实现
    1. 性能优化核心技巧
    1. 算子调试与测试
    1. 典型算子开发案例解析
    1. 开发者常见问题与解决方案
    1. 编译与部署流程
    1. 实际开发工具与文档参考
    1. 开发者学习路径建议

(这是昇腾知识体系的配套预览材料，转载随意，如反馈bug请移步原文：链接)

前言

为了更高效地开发算子，昇腾平台近期提供了CATLASS算子模板库，本教程是模板库编程的基础入门介绍，从最基础的环境配置，到如何让昇腾芯片发挥最大性能，再到实际调试技巧，十二个章节层层递进。特别适合那些刚接触昇腾开发、但又不熟悉芯片编程的人。比如您不需要一开始就理解什么是"bank冲突"，但需要知道如何正确设置工具链。建议初学者先看第1-3章建立基础，再结合第11章和文末推荐的文档，逐步深入复杂场景。

第1章：开发环境准备与基础概念

本文面向初学者，旨在用最短篇幅讲清昇腾算子开发的入门要点。我们将从环境配置到编程模型，逐步拆解算子开发的前置知识。昇腾算子的开发工具包叫做CANN（Compute Architecture for Neural Networks），整体包含驱动、开发工具链、模板库和性能分析工具等。

这一章还会介绍昇腾的三级内存结构：全局内存（Global Memory）、局部内存（Local Memory）和统一缓冲区（UB）。它们就像AI芯片的“物流系统”，决定了数据如何搬运、存储和计算。比如，Global Memory存放原始数据，Local Memory用于中间结果缓存，UB则帮助Vector和Scalar计算。

最后，本章会强调核函数参数的传递规则：核函数的参数必须通过结构化封装绑定，不能直接修改指针。这是昇腾算子开发的关键规范，违反它会导致计算结果错乱或程序崩溃。

1.1 昇腾开发环境搭建

1.1.1 CANN软件安装

昇腾算子开发需安装CANN软件包（Compute Architecture for Neural Networks），其包含：

  • 驱动固件
  • 开发工具链
  • 算子模板库
  • 性能分析工具（msprof op）

安装步骤（以Ubuntu为例）：

# 安装依赖
apt-get install -y gcc make net-tools cmake python3 python3-dev python3-pip

# 下载CANN安装包（需替换为实际版本号）
chmod +x Ascend-cann-toolkit_XXX_linux-x86_64.run
./Ascend-cann-toolkit_XXX_Linux-x86_64.run --install

# 配置环境变量
source /usr/local/Ascend/ascend-toolkit/set_env.sh

Tips：若需在非昇腾设备上开发（如PC），需先安装驱动固件（Atlas系列NPU驱动）并配置仿真环境。具体可参考《CANN软件安装指南》的"附录B：非昇腾设备安装"章节。

1.1.2 CPU仿真调试 vs NPU实际模式

场景 CPU仿真调试 NPU实际模式
硬件依赖 无需安装昇腾驱动 需安装昇腾驱动固件
内存管理 使用Host侧内存模拟Device行为 实际使用NPU的Local Memory
编译器行为 生成CPU可执行代码 生成NPU可执行代码
调试工具 gdb npu-smi + msprof op
典型配置 source /usr/local/Ascend/set_env.sh aclrtSetDevice() + aclrtCreateStream()

注意：仿真调试的代码结构与NPU实际模式完全一致，仅通过GetBlockIdx()获取的核ID不同（仿真模式返回0）。开发者可先用CPU仿真调试代码逻辑，再部署到NPU。

1.2 Ascend C编程模型基础

1.2.1 SPMD模型与多核并行

昇腾算子开发采用SPMD编程模型（Single Program, Multiple Data）：

  • SPMD：所有核执行相同代码，但处理不同数据切片
  • 核函数：用__global__ __aicore__定义，在Device侧运行
  • 核调用：通过kernel_name<<<blockdim, nullptr, stream>>>;语法执行，blockDim指定核数量（GetBlockIdx()获取当前核ID）
// 核函数定义示例
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

// 核调用示例
add_custom<<<blockDim, nullptr, stream>>>(x, y, z);

1.2.2 内存访问层级

昇腾架构中内存访问遵循三级存储模型，开发者需明确不同存储的用途和访问方式：

内存类型 缩写 访问方式 典型用途
Global Memory GM GlobalTensor 存储原始输入/输出数据
Local Memory L1/L0 LocalTensor + DataCopy 中间计算结果暂存（L1为缓存，L0为计算单元）
Unified Buffer UB TQue + TBuf Vector/Scalar计算的统一缓冲区

关键接口说明：

  • GlobalTensor：SetGlobalBuffer()设置全局内存地址
  • LocalTensor：AllocTensor()/FreeTensor()分配/释放局部内存
  • DataCopy：DataCopyParams控制搬运策略（如srcStride、dstStride）

示例：通过SetGlobalBuffer为GlobalTensor绑定地址

GlobalTensor<float> xGM;
xGM.SetGlobalBuffer((__gm__ float*)x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);

1.2.3 核函数参数传递

核函数参数需通过结构化封装传递，开发者需区分：

  • 输入参数：只读，不可修改指针本身
  • 输出参数：可写，需通过SetGlobalBuffer初始化

禁止操作示例：

// 错误：直接修改核函数参数
query = tmpQueryPtr; // 会导致地址错乱
tilingData = tmpTilingDataPtr; // 会引发核函数执行错误

正确实践：使用SetGlobalBuffer()绑定参数地址

// 正确：仅读取参数，输出参数可写
inputQueryGMTensor.SetGlobalBuffer(query); // 输入参数仅读
outputAttentionGMTensor.SetGlobalBuffer(attention); // 输出参数可写

1.3 开发工具链与调试

1.3.1 必备工具

工具 用途 典型命令
npu-smi 查询硬件型号与状态 npu-smi info
msprof op 性能分析与瓶颈定位 msprof op --output=&quot;./out&quot; --ai-core=on
msopst 生成ST测试用例 ./msopst run -i case.json
gdb CPU仿真调试 gdb -ex run

1.3.2 测试流程

  1. Host侧精度验证：通过aclrtMemcpy将结果从Device拷回Host，与预期值对比
  2. Device侧性能分析：使用msprof op采集aic_mte2_time等指标，结合trace.json分析流水线冲突
  3. ST测试用例生成：编写AddCustom_case.json定义输入输出，调用msopst自动生成测试用例

1.4 本章小结

核心概念 关键操作 注意事项
SPMD模型 通过GetBlockIdx()实现多核并行 核函数参数不可修改，仅可读写绑定内存
内存三级结构 DataCopy控制数据搬运 避免bank冲突，512B对齐提升性能
核函数调用语法 kernel_name<<<blockdim, nullptr, stream>>>;调用 blockDim建议设置为实际核数的倍数
性能调试工具 msprof op采集性能数据 需结合trace.json分析流水线利用率

Next Step：掌握本章后，开发者可进入第2章《昇腾算子模板库（CATLASS）架构解析》，了解如何基于模板库快速开发GEMM算子。

第2章 昇腾算子模板库（CATLASS）架构解析

打个比方，假设把开发算子比喻为“造车”，那么CATLASS模板库就是现成的“汽车零件库”。这一章将重点介绍，昇腾算子模板库的五层架构如何让开发者像搭积木一样快速组装算子。

模板库的最底层是Basic层，它封装了最基本的API，比如Mmad（矩阵乘加）、DataCopy（数据搬运）。往上是Tile层，负责最小粒度的数据搬运和计算；再往上是Block层，它像汽车的发动机，定义单核执行的计算逻辑；Kernel层是更高级的模块化组合，能适配多核并行；最上层的Device层则是开发者直接调用的接口，屏蔽了硬件细节。

比如，GEMM（通用矩阵乘法）算子的三层嵌套循环（外层遍历矩阵块，内层执行Cube计算）在模板库中被拆分为Block层的主循环和Tile层的数据搬运。通过模板化设计，开发者只需调整参数，就能让算子适配不同矩阵规模。

这一章还会展示一个典型的算子组装流程：从Device层封装到Kernel层调用，再到Block层的主循环和Epilogue（后处理）。您会发现，模板库的设计让算子开发变得模块化，比如您可以复用已有的TileCopy或BlockMmad组件，而无需重新编写底层逻辑。

2.1 模板库架构介绍

昇腾算子模板库（CATLASS）采用五层分层架构，从底层硬件抽象到顶层算子调用，各层分工明确，通过模板化设计实现高性能代码复用。架构层级关系如表1所示：

层级名称 功能描述 关键组件示例
Device层 提供统一的Host侧调用接口，封装算子的启动、参数传递和多核并行调度逻辑。 MatmulUniversalAdapter
Kernel层 算子逻辑的核心实现层，通过组合Block和Tile模块构建完整的计算流程。 BasicMatmul、GroupedMatmul、QuantMatmul等
Block层 定义单核计算的基本单元，包含主循环计算和后处理操作。 BlockMmad、BlockEpilogue
Tile层 实现数据分片搬运和计算，通过模板参数控制不同层次的数据通路的粒度。 TileCopy、TileMmad
Basic层 提供Ascend C的基础API，如Mmad、DataCopy等，是所有模板库的基础。 AscendC::Mmad、AscendC::DataCopy

2.1.1 GEMM三层嵌套循环与模板库映射

GEMM（General Matrix Multiply）的三层嵌套循环结构与CATLASS的分层设计高度契合。以下是典型GEMM的伪代码结构：

for (int block_m = 0; block_m < MatmulM; block_m += BlockTileM) {
    for (int block_n = 0; block_n < MatmulN; block_n += BlockTileN) {
        for (int k_tile = 0; k_tile < MatmulK; k_tile++) {
            for (int tile_mma_m = 0; tile_mma_m < m; tile_mma_m++) {
                for (int tile_mma_n = 0; tile_mma_n < n; tile_mma_n++) {
                    for (int tile_mma_k = 0; tile_mma_k < k; tile_mma_k++) {
                        mmad.call(c, a, b);
                    }
                }
            }
        }
    }
}

CATLASS的层级映射关系如下：

  • 前两层（block_m、block_n）：对应Kernel层的多核并行性，每个AICore负责一个分块。
  • 第三层（k_tile）：对应Block层的矩阵乘加主循环（BlockMmad）。
  • 内层（tile_mma_m, tile_mma_n, tile_mma_k）：对应Tile层的分片搬运和计算。

2.1.2 各层组件调用关系

CATLASS的算子组装流程分为三步：

  1. Block层：定义单核计算逻辑（如矩阵乘加主循环和结果处理）。
  2. Kernel层：将多个Block组合，并处理多核同步和数据排布。
  3. Device层：封装Kernel调用，提供Host侧统一接口。

以下是典型算子的组装代码示例：

using DispatchPolicy = matmul::MmadAtlasA2Pingpong<true>;
using L1TileShape = MatmulShape<128, 256, 256>;
using L0TileShape = MatmulShape<128, 256, 64>;
using AType = matmul::MatmulType<ElementA, LayoutA>;
using BType = matmul::MatmulType<ElementB, LayoutB>;
using CType = matmul::MatmulType<ElementC, LayoutC>;

using BlockMmad = matmul::block::BlockMmad<DispatchPolicy,
    L1TileShape,
    L0TileShape,
    AType,
    BType,
    CType>;

using BlockEpilogue = void;

using TileScheduler = matmul::block::MatmulIdentityBlockSwizzle<>;

using MatmulKernel = matmul::kernel::BasicMatmul<BlockMmad, BlockEpilogue, TileScheduler>;

using MatmulHandle = CATLASS::matmul::device::MatmulUniversalAdapter<MatmulKernel>;

2.2 Block层组件详解

Block层是CATLASS的核心模块，负责单核计算逻辑的实现。主要组件包括：

  • BlockMmad：矩阵乘加主循环的模板化接口，支持多数据类型和排布。
  • BlockEpilogue：结果矩阵的尾处理，比如逐元素操作（如加Bias），支持自定义扩展。

2.2.1 BlockMmad的实现

template <
    class DispatchPolicy,
    class L1TileShape,
    class L0TileShape,
    class AType,
    class BType,
    class CType,
    class BiasType = void,
    class TileCopy = Gemm::Tile::TileCopy<typename DispatchPolicy::ArchTag, AType, BType, CType, BiasType>,
    class TileMmad = Gemm::Tile::TileMmad<typename DispatchPolicy::ArchTag, AType, BType, BiasType>
>
struct BlockMmad {
    static_assert(DEPENDENT_FALSE<DispatchPolicy>, "BlockMmad is not implemented for this DispatchPolicy");
};

BlockMmad的模板参数定义了Block层级的计算策略和数据排布：

参数名称 说明 典型取值示例
DispatchPolicy 控制Block的具体分配策略（比如MmadAtlasA2Pingpong） Gemm::MmadAtlasA2Pingpong<true>;
L1TileShape L1 Buffer的分片大小（M, N, K方向） GemmShape<128, 256, 256>;
L0TileShape L0 Buffer的分片大小（M, N, K方向） GemmShape<128, 256, 64>;
AType/BType/CType/BiasType 矩阵A/B/C/Bias的数据类型和排布（如half、RowMajor） Gemm::GemmType<half, LayoutA>;
TileCopy 使用到的Tile级别的数据搬运接口 Gemm::Tile::TileCopy<typename DispatchPolicy::ArchTag, AType, BType, CType, BiasType>;
TileMmad 使用的TIle级别的Mmad接口 Gemm::Tile::TileMmad<typename DispatchPolicy::ArchTag, AType, BType, BiasType>;
  • 接口中规定了Block层的调度策略、L1和L0 Buffer上基本块大小、全局内存上A、B、C矩阵和Bias向量的数据类型和数据排布、不同访存层级间的块粒度数据拷贝以及L0上基本块粒度的矩阵乘累加运算。其中L1TileShape和L0TileShape统一了分块大小，各种Type参数统一了数据类型和数据排布。通过这些模板参数不同的分块大小、数据类型和数据排布都可以使用这个统一的接口。
  • 以上是基础的BlockMmad设计，用户可参考自定义实现复杂的主循环组件以快速实现不同的自定义算子。

2.2.2 BlockEpilogue的实现

BlockEpilogue用于在Block计算完成后执行后处理操作，比如逐元素操作（如加Bias、ReLU）。通过模板化设计，用户可灵活扩展该层功能：

template <
    class CType_,
    class XType_,
    class DType_,
    class TileElemWiseEpilogue_,
    class TileCopy_
>
class BlockEpilogue <
    EpilogueAtlasA2ElemWiseOneSource,
    CType_,
    XType_,
    DType_,
    TileElemWiseEpilogue_,
    TileCopy_
> {
        // 使用模板参数来具体实现计算逻辑
    };

2.2.2 Block Dispatch Policies的实现

Dispatch_policy影响着Block级具体的实现，针对具体的架构和硬件会有不同的实现，其中AtlasA2架构上的基础实现如下：

template <bool ASYNC_ = false> // ASYNC_表示是否需要进行同步
struct MmadAtlasA2Base {
    using ArchTag = arch::AtlasA2;
    static constexpr uint32_t ASYNC = ASYNC_;
};
using MmadAtlasA2 = MmadAtlasA2Base<false>;
using MmadAtlasA2Async = MmadAtlasA2Base<true>;

其余的分配策略以继承的方式在MmadAtlasA2Base的基础上完成，比如双缓冲实现：

template <bool ENABLE_UNIT_FLAG_ = false> //  参数用于表示是否启用Mmad运算与L0C结果拷贝到全局内存的细粒度并行。
struct MmadAtlasA2Pingpong : public MmadAtlasA2  {
    static constexpr uint32_t STAGES = 2;
    static constexpr bool ENABLE_UNIT_FLAG = ENABLE_UNIT_FLAG_;
};

其中STAGES 参数使用户可以方便的调整多buffer场景的buffer片数，ENABLE_UNIT_FLAG 参数用于表示是否启用Mmad运算与L0C结果拷贝到全局内存的细粒度并行。

2.3 Tile层内核设计

Tile层负责数据搬运的最小粒度（分片）和计算逻辑的细粒度实现。主要组件包括：

  • TileCopy：定义GM与ub/L1/l0间的搬运接口（如copy_gm_to_ub、copy_ub_to_gm）。
  • TileMmad：基于Ascend C基础API的矩阵乘加指令封装tile级别的计算逻辑（Mmad指令调用）。

2.3.1 TileCopy的关键参数

TileCopy通过模板参数和配置结构体控制数据搬运路径和格式转换：

template <class ArchTag, class AType, class BType, class CType, class BiasType>
struct TileCopy {
    using ElementA = typename AType::Element;
    using ElementB = typename BType::Element;
    using ElementAccumulator = typename matmul::helper::ElementAccumulatorSelector<ElementA, ElementB>::ElementAccumulator;

    using CopyGmToL1A = matmul::tile::CopyGmToL1<ArchTag, AType>;
    using CopyGmToL1B = matmul::tile::CopyGmToL1<ArchTag, BType>;
    using CopyL1ToL0A = matmul::tile::CopyL1ToL0A<ArchTag, ElementA>;
    using CopyL1ToL0B = matmul::tile::CopyL1ToL0B<ArchTag, ElementB>;
    using CopyL0CToGm = matmul::tile::CopyL0CToGm<ArchTag, ElementAccumulator, CType>;
};

关键参数说明如表2所示：

参数名称 说明 典型场景
ArchTag 硬件架构标签（如AtlasA2） 控制不同架构的具体实现策略
ElementA/B/C 定义矩阵A、B、C的元素类型（如half、float） 数据类型选择需与硬件指令兼容
CopyGmToL1A/CopyGmToL1B等 不同硬件之间适用于特定架构和数据类型的搬运接口 需与硬件架构的访存特性匹配

2.3.2 TileMmad的关键参数

TileMmad重载()运算符，负责通过AscendC::Mmad调用矩阵运算单元完成计算。模板参数包括A、B矩阵和Bias的数据类型和数据排布，核心代码如下：

template <class ArchTag_, class AType_, class BType_, class BiasType_> // 新增ArchTag表示架构标签，屏蔽不同架构的差异。
struct TileMmad {
    using ElementA = typename AType_::Element;
    using ElementB = typename BType_::Element;
    using ElementAccumulator = typename matmul::helper::ElementAccumulatorSelector<ElementA, ElementB>::ElementAccumulator;
    void operator()(AscendC::LocalTensor<ElementAccumulator> const &l0CTensor,
                    AscendC::LocalTensor<ElementA> const &l0ATensor,
                    AscendC::LocalTensor<ElementB> const &l0BTensor,
                    uint32_t m, uint32_t n, uint32_t k,
                    bool initC = true, uint8_t unitFlag = 0) {
        // 构造mmadParams
        ...
        AscendC::Mmad(l0CTensor,
                      l0ATensor,
                      l0BTensor,
                      mmadParams);
        ...
    }
};

2.4 Kernel层组装逻辑

Kernel层通过模块化设计block级别的主循环、后处理、计算顺序提高了算子开发效率及泛化性，用户可以单独实现特定的组件，通过模板参数传入，便可以快速开发对应的内核。Kernel层主要功能包括：

  • 对包含的不同的block逻辑进行组合，加入必要的同步逻辑，使用PipeBarrier、SetFlag/WaitFlag实现AICore与AIV核间的流水线同步。
  • 处理不同block与全局内存数据的对应关系（Swizzling操作）
  • 对输入数据进行block粒度切分。

2.4.1 Kernel的模块化实现

以BasicMatmul的内核实现为例：

#ifndef ACT_GEMM_KERNEL_MATMUL_HPP
#define ACT_GEMM_KERNEL_MATMUL_HPP

#include "act/act.hpp"
#include "act/arch/resource.hpp"
#include "act/coord.hpp"
#include "act/gemm_coord.hpp"
#include "act/matrix_coord.hpp"

namespace Act::Gemm::Kernel {

// Template for Matmul kernel. Compute C = A * B
template <
    class BlockMmad_, // block主循环
    class BlockEpilogue_, // block后处理
    class BlockScheduler_ // block计算顺序
>
class BasicMatmul {
public:
    using BlockMmad = BlockMmad_;
    using ArchTag = typename BlockMmad::ArchTag;
    using L1TileShape = typename BlockMmad::L1TileShape;
    using ElementA = typename BlockMmad::ElementA;
    using LayoutA = typename BlockMmad::LayoutA;
    using ElementB = typename BlockMmad::ElementB;
    using LayoutB = typename BlockMmad::LayoutB;
    using ElementC = typename BlockMmad::ElementC;
    using LayoutC = typename BlockMmad::LayoutC;
    using ElementAccumulator = typename BlockMmad::ElementAccumulator;

    using BlockScheduler = BlockScheduler_;

    /// Parameters structure
    struct Params {
        // Data members
        GemmCoord problemShape;
        GM_ADDR ptrA;
        LayoutA layoutA;
        GM_ADDR ptrB;
        LayoutB layoutB;
        GM_ADDR ptrC;
        LayoutC layoutC;

        // Methods
        ACT_DEVICE
        Params() {}

        ACT_DEVICE
        Params(GemmCoord const &problemShape_, GM_ADDR ptrA_, LayoutA layoutA_, GM_ADDR ptrB_,
               LayoutB layoutB_, GM_ADDR ptrC_, LayoutC layoutC_)
            : problemShape(problemShape_), ptrA(ptrA_), layoutA(layoutA_), ptrB(ptrB_), layoutB(layoutB_),
              ptrC(ptrC_), layoutC(layoutC_) {}
    };

    // Methods
    ACT_DEVICE
    BasicMatmul() {}

    template <int32_t CORE_TYPE = g_coreType>
    ACT_DEVICE
    void operator()(Params const &params);

    /// Executes one Matmul
    template <>
    ACT_DEVICE
    void operator()<AscendC::AIC>(Params const &params) {
        BlockScheduler matmulBlockScheduler(params.problemShape, MakeCoord(L1TileShape::M, L1TileShape::N));
        uint32_t coreLoops = matmulBlockScheduler.GetCoreLoops();

        Arch::Resource<ArchTag> resource;
        BlockMmad blockMmad(resource);

        // Represent the full gm
        AscendC::GlobalTensor<ElementA> gmA;
        gmA.SetGlobalBuffer((__gm__ ElementA *)params.ptrA);
        AscendC::GlobalTensor<ElementB> gmB;
        gmB.SetGlobalBuffer((__gm__ ElementB *)params.ptrB);
        AscendC::GlobalTensor<ElementC> gmC;
        gmC.SetGlobalBuffer((__gm__ ElementC *)params.ptrC);

        for (uint32_t loopIdx = AscendC::GetBlockIdx(); loopIdx < coreLoops; loopIdx += AscendC::GetBlockNum()) {
            // Compute block location
            GemmCoord blockCoord = matmulBlockScheduler.GetBlockCoord(loopIdx);
            GemmCoord actualBlockShape = matmulBlockScheduler.GetActualBlockShape(blockCoord);

            // Compute initial location in logical coordinates
            MatrixCoord offsetA{blockCoord.m() * L1TileShape::M, blockCoord.k() * L1TileShape::K};
            MatrixCoord offsetB{blockCoord.k() * L1TileShape::K, blockCoord.n() * L1TileShape::N};
            MatrixCoord offsetC{blockCoord.m() * L1TileShape::M, blockCoord.n() * L1TileShape::N};
            int64_t gmOffsetA = params.layoutA.GetOffset(offsetA);
            int64_t gmOffsetB = params.layoutB.GetOffset(offsetB);
            int64_t gmOffsetC = params.layoutC.GetOffset(offsetC);

            // Compute block-scoped matrix multiply-add
            blockMmad(gmA[gmOffsetA], params.layoutA,
                      gmB[gmOffsetB], params.layoutB,
                      gmC[gmOffsetC], params.layoutC,
                      actualBlockShape);
        }
    }

    template <>
    ACT_DEVICE
    void operator()<AscendC::AIV>(Params const &params) {}
};

} // namespace Act::Gemm::Kernel

#endif // ACT_GEMM_KERNEL_MATMUL_HPP

2.4.2 多核同步控制

多核场景下，开发者需注意同步控制对性能的影响。例如，使用SetFlag和WaitFlag确保数据搬运与计算的顺序性：

// 示例：多核同步控制
PipeBarrier<PIPE_ALL>();
SetFlag<HardEvent::MTE2>(eventId);
WaitFlag<HardEvent::Mte2>(eventId);

2.5 Device层调用流程

Device层是Host侧调用算子的统一入口，通过封装Kernel的执行逻辑，实现算子的高效调用。典型组件包括：

  • MatmulUniversalAdapter：封装GEMM算子的启动、参数传递和多核调度。

以下是Device层的典型调用代码：

// 组装BlockMmad、BlockEpilogue、TileScheduler和对用的MatmulKernel
using BlockMmad = matmul::block::BlockMmad<DispatchPolicy,
    L1TileShape,
    L0TileShape,
    AType,
    BType,
    CType>;

using BlockEpilogue = void;

using TileScheduler = matmul::block::MatmulIdentityBlockSwizzle<>;

using MatmulKernel = matmul::kernel::BasicMatmul<BlockMmad, BlockEpilogue, TileScheduler>;
// 组装device适配器
using MatmulHandle = matmul::device::MatmulUniversalAdapter<MatmulKernel>;
MatmulHandle matmulAdapter;
matmulAdapter(args, workspace, stream);

2.6 各层交互示意图

以下展示了典型GEMM算子的数据流和控制流：

GEMM算子分层交互示意图
[Host侧调用] -> [Device层] -> [Kernel层] -> [Block层] -> [Tile层] -> [Basic层]
  • 数据流：从Global Memory（GM）经L1、L0AB、L0C逐层搬运。
  • 控制流：通过模板参数配置各层计算策略，最终由Kernel层调用Mmad完成核心计算。

2.7 小结与建议

2.7.1 分层设计的优势

CATLASS的分层架构通过模板化设计显著提升了算子开发效率：

  • Device层：统一调用接口，减少开发者对硬件细节的依赖。
  • Kernel层：模块化设计使多核调度和同步控制更易实现。
  • Block & Tile层：模板参数化支持灵活适配不同数据类型和排布。

2.7.2 开发者实践建议

建议点 具体操作
模板参数选择 根据硬件规格（如L1大小）合理配置L1TileShape和L0TileShape。
分块粒度优化 优先使用512B对齐的分块策略，提升搬运效率。
同步控制 多核场景中务必插入PipeBarrier，避免因未同步导致的错误计算。
调试建议 通过msopst工具生成ST测试用例，验证精度和性能（参考第7章）。

[示例代码]
以下代码展示了从Device层到Basic层的完整调用关系：

// Device层适配器
using MatmulHandle = matmul::device::MatmulUniversalAdapter<MatmulKernel>;
MatmulHandle matmulAdapter;

//Kernel层核心
using MatmulKernel = matmul::kernel::BasicMatmul<BlockMmad, BlockEpilogue, TileScheduler>;

//Block层主循环
using Block::BlockMmad<DispatchPolicy, L1TileShape, L0TileShape, AType, BType, CType>;

//Tile层搬运
using TileCopy<ArchTag, AType, BType, CType>;

通过上述分层设计，开发者可以快速基于模板库实现GEMM、GEMV、量化算子等复杂逻辑，后续章节将深入具体算子的开发流程。

第3章 基于模板库开发GEMM算子全流程

这章聚焦于GEMM算子（通用矩阵乘法）的开发，这是深度学习模型中最核心的计算之一。开发GEMM的关键在于分块策略的设计——将矩阵分解为多个小块，每个块由不同的核心并行处理。想象您在做一道大菜，把食材分成小份，每个厨师（核心）负责一份，最后合并结果。

分块策略需要权衡多个因素：比如BlockTileM和BlockTileN的大小决定了每个核心处理的数据量，而L1TileShape和L0TileShape则决定了局部内存的分配。过大或过小的分块都会影响性能。比如，如果分块太大，局部内存可能装不下；如果分块太小，搬运数据的开销反而会抵消并行计算的收益。教程中提供的对比表显示，合理调整分块大小后，性能可提升数倍。

另一个核心是L2缓存的优化。昇腾AI芯片的L2缓存容量有限（如192MB），而GM的带宽（1.6TB/s）却很高。如果数据无法命中缓存，性能会大打折扣。教程建议通过512字节对齐的地址设计和Swizzling（数据分块排列策略）来最大化缓存利用率。比如，Swizzling能重新排列分块的计算顺序，让相同的数据被多个核心复用，减少重复搬运。

最后，教程提供了完整的开发步骤示例，包括核函数定义和Host侧调用。如果您刚开始接触，建议先用CPU仿真调试验证逻辑，再部署到真实硬件上。这章的实战部分（GroupGEMM开发）演示了如何处理多矩阵输入，通过动态调整每个矩阵的分块策略，实现多核并行计算。

3.1 GEMM分块策略设计

3.1.1 分块粒度选择

GEMM（通用矩阵乘法）的性能优化核心在于分块策略。根据矩阵规模（M, N, K）选择合适的分块粒度，将计算分解为多个可并行的子任务。昇腾NPU的分块策略通常涉及以下层级：

  • Block级（结果块分块的大小）：将矩阵按M和N方向划分为多个BlockTileM和BlockTileN大小的基块。每个基块由一个AICore独立计算。
  • Tile级（分块加载仅L1/L0的分块大小）：在Block内部进一步将矩阵划分为更小的Tile（L0TileM, L0TileN, L0TileK），适应片上缓存（L1/L0）的存储限制。

关键参数对比：

参数名称 含义 典型值（Atlas A2） 优化建议
BlockTileM Block在M方向的分块大小 128~512 根据L1缓存容量调整
BlockTileN Block在N方向的分块大小 256~1024 与BlockTileM保持比例
L1TileShape L1缓存中的Tile形状（M,N,K） (128,256,256) 优化为512B对齐
L0TileShape L0缓存中的Tile形状（m,n,k） (64,64,64) 根据Cube计算单元特性调整

3.1.2 L2Cache命中优化

昇腾AI处理器的L2Cache容量（如192MB）与GM带宽（1.6TB/s）存在差异。若数据无法命中L2Cache，会导致性能瓶颈。优化策略包括：

  • 512B对齐：GM->L1/L0的数据搬运时，确保地址对齐。对齐后带宽利用率提升30%以上。
  • Swizzling：调整分块计算顺序，最大化L2Cache命中率。例如在CATLASS模板库中Swizzling由由模板类MatmulIdentityBlockSwizzle控制，其中排列顺序由模板参数SwizzleDirection和SwizzleOffset决定。

示例代码（Swizzling实现）：

// Swizzling操作核心代码，根据taskIdx获取左右矩阵行列索引mIdx/nIdx
template <uint32_t SwizzleOffset = 1, uint32_t SwizzleDirection = 0>
struct GemmIdentityBlockSwizzle {
    /// Data members
    GemmCoord problemShape;
    MatrixCoord tileMN;
    MatrixCoord loopsMN;

    /// Methods
    ACT_DEVICE
    GemmIdentityBlockSwizzle() {}

    ACT_DEVICE
    GemmIdentityBlockSwizzle(GemmCoord const &problemShape_, MatrixCoord const &tileMN_)
        : problemShape(problemShape_), tileMN(tileMN_)
    {
        loopsMN = CeilDiv(MatrixCoord(problemShape.GetCoordMN()), tileMN);
    }

    ACT_DEVICE
    GemmIdentityBlockSwizzle(GemmCoord const &problemShape_, MatrixCoord const &tileMN_,
        MatrixCoord const &loopsMN_)
        : problemShape(problemShape_), tileMN(tileMN_), loopsMN(loopsMN_) {}

    ACT_DEVICE
    void Update(GemmCoord const &problemShape_, MatrixCoord const &tileMN_)
    {
        problemShape = problemShape_;
        tileMN = tileMN_;

        loopsMN = CeilDiv(MatrixCoord(problemShape.GetCoordMN()), tileMN);
    }

    ACT_DEVICE
    void Update(GemmCoord const &problemShape_, MatrixCoord const &tileMN_, MatrixCoord const &loopsMN_)
    {
        problemShape = problemShape_;
        tileMN = tileMN_;
        loopsMN = loopsMN_;
    }

    ACT_DEVICE
    uint32_t GetCoreLoops() const
    {
        return loopsMN.row() * loopsMN.column();
    }

    ACT_DEVICE
    uint32_t GetBatchIdx(uint32_t taskIdx)
    {
        return taskIdx / (GetCoreLoops());
    }

    ACT_DEVICE
    GemmCoord GetBlockCoord(uint32_t taskIdx)
    {
        uint32_t innerIdx = taskIdx % GetCoreLoops();
        if constexpr (SwizzleDirection == 0) { // Zn
            uint32_t tileBlockLoop = CeilDiv(loopsMN.row(), SwizzleOffset);
            uint32_t tileBlockIdx = innerIdx / (SwizzleOffset * loopsMN.column());
            uint32_t inTileBlockIdx = innerIdx % (SwizzleOffset * loopsMN.column());

            uint32_t nRow = SwizzleOffset;
            if (tileBlockIdx == tileBlockLoop - 1) {
                nRow = loopsMN.row() - SwizzleOffset * tileBlockIdx;
            }
            uint32_t mIdx = tileBlockIdx * SwizzleOffset + inTileBlockIdx % nRow;
            uint32_t nIdx = inTileBlockIdx / nRow;
            if (tileBlockIdx % 2 == 1) {
                nIdx = loopsMN.column() - nIdx - 1;
            }
            return GemmCoord{mIdx, nIdx, 0};
        } else if constexpr (SwizzleDirection == 1) { // Nz
            uint32_t tileBlockLoop = CeilDiv(loopsMN.column(), SwizzleOffset);
            uint32_t tileBlockIdx = innerIdx / (SwizzleOffset * loopsMN.row());
            uint32_t inTileBlockIdx = innerIdx % (SwizzleOffset * loopsMN.row());
            uint32_t nCol = SwizzleOffset;
            if (tileBlockIdx == tileBlockLoop - 1) {
                nCol = loopsMN.column() - SwizzleOffset * tileBlockIdx;
            }
            uint32_t mIdx = inTileBlockIdx / nCol;
            uint32_t nIdx = tileBlockIdx * SwizzleOffset + inTileBlockIdx % nCol;
            if (tileBlockIdx % 2 == 1) {
                mIdx = loopsMN.row() - mIdx - 1;
            }
            return GemmCoord{mIdx, nIdx, 0};
        }
    }

    ACT_DEVICE
    GemmCoord GetActualBlockShape(GemmCoord blockCoord)
    {
        uint32_t mActual = (blockCoord.m() == (loopsMN.row() - 1)) ?
            (problemShape.m() - blockCoord.m() * tileMN.row()) : tileMN.row();
        uint32_t nActual = (blockCoord.n() == (loopsMN.column() - 1)) ?
            (problemShape.n() - blockCoord.n() * tileMN.column()) : tileMN.column();
        uint32_t kActual = problemShape.k();
        return GemmCoord{mActual, nActual, kActual};
    }
};

CATLASS模板库中Swizzling方案的使用：

using BlockScheduler = typename Gemm::Block::GemmIdentityBlockSwizzle<3, 0>;

// 将Swizzling方案传入Kernel中
using MatmulKernel = Gemm::Kernel::BasicMatmul<BlockMmad, BlockEpilogue, BlockScheduler>;

3.2 GEMM数据搬运与计算流程

3.2.1 数据搬运层级

CATLASS模板库中的数据搬运由前面介绍的TileCopy组件完成，数据搬运遵循三层架构，确保数据在不同存储层级间的高效流转：

  • GM -> L1：使用CopyGmToL1接口，将大矩阵切片后搬运至L1缓存。
  • L1 -> L0：通过CopyL1ToL0A和CopyL1ToL0B接口，将数据进一步搬运到L0A/L0B。
  • L0 -> GM：计算结果通过CopyL0CToGm接口写回GM，同时支持随路格式转换（如zN到行优先ND）。

3.2.2 Cube计算单元调用

矩阵乘计算通过Mmad指令完成。其核心是定义M, N, K的维度，并配置计算参数（如是否初始化结果矩阵）。

示例代码（Mmad调用）：

template <class ArchTag_, class AType_, class BType_, class BiasType_>
struct TileMmad {
    void operator()(LocalTensor<ElementAccumulator> l0CTensor,
                    LocalTensor<ElementA> l0ATensor,
                    LocalTensor<ElementB> l0BTensor,
                    uint32_t m, uint32_t n, uint32_t k,
                    bool initC = true, uint8_t unitFlag = 0) {
        MmadParams mmadParams;
        mmadParams.m = m;
        mmadParams.n = n;
        mmadParams.k = k;
        AscendC::Mmad(l0CTensor, l0ATensor, l0BTensor, mmadParams);
    }
};

3.2.3 搬运与计算的协同

在分块策略下，需确保搬运和计算的流水线并行：

  • 搬运优先级：每次搬运需覆盖完整的Tile数据（如L0TileShape定义的尺寸）。
  • 计算依赖：Cube计算需等待数据从L1搬运到L0完成。

3.3 SplitK方案实现

3.3.1 K方向分块

当矩阵规模较小（如K << M/N）时，通过K_TILE将K方向分块，结合GetCoreNum()实现多核并行，避免核利用率不足。

性能收益对比（示例）：

场景 核利用率（单核） 核利用率（SplitK） 性能提升
K=256 20% 85% 4x
K=1024 55% 98% 1.8x

3.4 GEMM模板库快速开发

3.4.1 模板特化优势

昇腾CATLASS提供了多种高性能组件，并利用DispatchPolicy组件将复杂的内核实现解耦，DispatchPolicy可以与多个主循环与尾处理快速组合成不同实现策略的算子，以快速实现GEMM算子。以双缓冲策略为例：

template <bool ENABLE_UNIT_FLAG_ = false>
struct MmadAtlasA2Pingpong : public MmadAtlasA2  {
    static constexpr uint32_t STAGES = 2;
    static constexpr bool ENABLE_UNIT_FLAG = ENABLE_UNIT_FLAG_;
};

通过设置STAGES=2和UNIT_FLAG，优化流水线并行。在实现了双缓冲的DispatchPolicy后，只要实现该策略下的主循环和 后处理计算逻辑，可以快速在主循环或后处理中应用该策略来实现特定计算过程：

using DispatchPolicy = Gemm::MmadAtlasA2Pingpong<true>;
...
using BlockMmad = Gemm::Block::BlockMmad<DispatchPolicy, L1TileShape, L0TileShape, AType, BType, CType>;

模板参数说明：

参数名称 含义 典型值
STAGES 双缓冲缓存区数量 2
UNIT_FLAG 启用Mmad运算与L0C结果拷贝到全局内存的细粒度并行 true

3.4.2 封装算子逻辑

通过MatmulUniversalAdapter封装Kernel逻辑，实现Device层的调用接口。

示例代码（模板组装）：

using BlockMmad = matmul::block::BlockMmad<DispatchPolicy,
    L1TileShape,
    L0TileShape,
    AType,
    BType,
    CType>;

using BlockEpilogue = void;

using TileScheduler = matmul::block::MatmulIdentityBlockSwizzle<>;

using MatmulKernel = matmul::kernel::BasicMatmul<BlockMmad, BlockEpilogue, TileScheduler>;

using MatmulHandle = CATLASS::matmul::device::MatmulUniversalAdapter<MatmulKernel>;

3.5 小结：GEMM开发关键点

3.5.1 开发流程总结

步骤 操作说明 工具/API
分块设计 定义BlockTileM和BlockTileN MatmulShape模板
数据搬运 GM->L1->L0->GM TileCopy组件
Cube计算 调用Mmad指令 TileMmad模板
SplitK优化 K方向分块提升核利用率 GetCoreNum()函数
模板封装 使用MatmulUniversalAdapter CATLASS::MatmulUniversalAdapter

3.5.2 常见问题与解决方案

问题1： 数据搬运导致L2Cache miss率高
解决方案： 优化分块大小，确保512B对齐。例如：

Nd2NzParams dataCopyA1Params;
dataCopyA1Params.srcDValue = k;
dataCopyA1Params.dstNzC0Stride = m;

问题2： Cube计算耗时过长
解决方案： 调整Mmad参数，启用STAGES=2减少流水等待。

问题3： 多核同步异常（卡死）
解决方案：

  1. 检查BlockDim是否超过GetCoreNum()，确保BlockDim <= CoreNum。
  2. 检查同步语句是否使用正确、数量是否成对。

3.6 高性能GEMM开发建议

优化方向 实施建议 工具/API
数据搬运优化 一次性搬运大块数据，减少循环次数 DataCopy接口
内存对齐 GM地址512B对齐 DataCopyParams结构体
流水线并行 使用double buffer减少等待 TQue与TPipe
核间同步优化 谨慎使用PipeBarrier，避免依赖 SetFlag/WaitFlag
格式转换 利用Fixpipe随路转换 FixpipeParamsV220

3.7 附录：完整开发步骤

3.7.1 核函数定义与调用

定义核函数时需遵循__global__ __aicore__规则，并通过<<<blockDim, nullptr, stream>>>启动。

示例代码（核函数定义）：

extern "C" __global__ __aicore__ void MatmulKernel(__gm__ half* A, __gm__ half* B, __gm__ float* C, __gm__ uint8_t* workspace, __gm__ uint8_t* tilingData) {
    // 初始化Tensor地址
    A->SetGlobalBuffer((__gm__ half*)A + BlockTileM * GetBlockIdx(), BlockTileM);
    B->SetGlobalBuffer((__gm__ half*)B + BlockTileN * GetBlockIdx(), BlockTileN);
    C->SetGlobalBuffer((__gm__ float*)C + BlockTileM * GetBlockIdx(), BlockTileM);

    // 搬运与计算
    TileCopyParams copyParams;
    mmadParams.m = BlockTileM;
    mmadParams.n = BlockTileN;
    mmadParams.k = BlockTileK;
}

3.7.2 Host侧调用

Host侧需分配设备内存、设置分块参数，并调用MatmulUniversalAdapter执行算子或者直接执行Kernel。

示例代码（Host侧调用）：

// 通过device适配器调用
MatmulUniversalAdapter matmulAdapter;
matmulAdapter(args, workspace, stream);

// 通过Kernel直接调用
using MatmulKernel = Gemm::Kernel::[Kernel_name]<BlockMmad, BlockEpilogue, BlockScheduler>;

typename MatmulKernel::Params params{problemShape, gmA, layoutA, gmB, layoutB, gmC, layoutC};

MatmulKernel matmul;
matmul(params);

3.8 实战：GroupGEMM开发

GroupGEMM用于处理多个不同规模的矩阵乘法。其核心是将多个矩阵的分块统一分配到不同的核心上并行计算，需要处理每个核心上的分块逻辑。每一个分块的计算可以调用BlockMmad来完成，主要的工作在于计算每个分块需要加载数据的索引。

示例代码（GroupGEMM计算逻辑）：

GemmCoord problemShapeList[MAX_TENSOR_COUNT];
LayoutA layoutAList[MAX_TENSOR_COUNT];
LayoutB layoutBList[MAX_TENSOR_COUNT];
LayoutC layoutCList[MAX_TENSOR_COUNT];

// Get matmul information from parameters
detail::UnpackListParam(problemShapeList, params.ptrProblemShape, params.problemCount);
detail::UnpackListParam(layoutAList, params.ptrLayoutA, params.problemCount);
detail::UnpackListParam(layoutBList, params.ptrLayoutB, params.problemCount);
detail::UnpackListParam(layoutCList, params.ptrLayoutC, params.problemCount);

BlockScheduler matmulBlockScheduler;
Arch::Resource<ArchTag> resource;
BlockMmad blockMmad(resource);

// Represent the full gm
AscendC::GlobalTensor<ElementA> gmA;
gmA.SetGlobalBuffer((__gm__ ElementA *)params.ptrA);
AscendC::GlobalTensor<ElementB> gmB;
gmB.SetGlobalBuffer((__gm__ ElementB *)params.ptrB);
AscendC::GlobalTensor<ElementC> gmC;
gmC.SetGlobalBuffer((__gm__ ElementC *)params.ptrC);

uint32_t coreIdx = AscendC::GetBlockIdx();
uint32_t coreNum = AscendC::GetBlockNum();
int64_t inGroupOffsetA = 0;
int64_t inGroupOffsetB = 0;
int64_t inGroupOffsetC = 0;

uint32_t startCoreIdx = 0;
for (uint32_t groupIdx = 0; groupIdx < params.problemCount; ++groupIdx) {
    GemmCoord problemShape = problemShapeList[groupIdx];
    LayoutA layoutA = layoutAList[groupIdx];
    LayoutB layoutB = layoutBList[groupIdx];
    LayoutC layoutC = layoutCList[groupIdx];

    matmulBlockScheduler.Update(problemShape, MakeCoord(L1TileShape::M, L1TileShape::N));
    uint32_t coreLoops = matmulBlockScheduler.GetCoreLoops();

    // Determine the starting loopIdx of the current core under the current groupIdx
    uint32_t startLoopIdx;
    if (coreIdx < startCoreIdx) {
        startLoopIdx = coreIdx + coreNum - startCoreIdx;
    } else {
        startLoopIdx = coreIdx - startCoreIdx;
    }
    // Loop through the matmul of each groupIdx
    for (uint32_t loopIdx = startLoopIdx; loopIdx < coreLoops; loopIdx += coreNum) {
        // Compute block location
        GemmCoord blockCoord = matmulBlockScheduler.GetBlockCoord(loopIdx);
        GemmCoord actualBlockShape = matmulBlockScheduler.GetActualBlockShape(blockCoord);

        // Compute initial location in logical coordinates
        MatrixCoord offsetA{blockCoord.m() * L1TileShape::M, blockCoord.k() * L1TileShape::K};
        MatrixCoord offsetB{blockCoord.k() * L1TileShape::K, blockCoord.n() * L1TileShape::N};
        MatrixCoord offsetC{blockCoord.m() * L1TileShape::M, blockCoord.n() * L1TileShape::N};
        int64_t gmOffsetA = layoutA.GetOffset(offsetA);
        int64_t gmOffsetB = layoutB.GetOffset(offsetB);
        int64_t gmOffsetC = layoutC.GetOffset(offsetC);

        // Compute block-scoped matrix multiply-add
        blockMmad(
            gmA[inGroupOffsetA + gmOffsetA], layoutA,
            gmB[inGroupOffsetB + gmOffsetB], layoutB,
            gmC[inGroupOffsetC + gmOffsetC], layoutC,
            actualBlockShape);
    }

    inGroupOffsetA += problemShape.m() * problemShape.k();
    inGroupOffsetB += problemShape.k() * problemShape.n();
    inGroupOffsetC += problemShape.m() * problemShape.n();

    startCoreIdx = (startCoreIdx + coreLoops) % coreNum;
}

3.9 性能验证与调试

最后，通过msprof op工具采集性能数据，分析aic_mte2_time和aic_mac_ratio等指标。若发现瓶颈，可进一步优化分块策略或调整流水线并行度。

第4章 GEMV与量化算子开发

GEMV（通用矩阵向量乘法）和量化算子是GEMM的延伸，但它们的实现方式更灵活。GEMV常用于处理向量计算，比如NLP模型中的注意力机制。它的分块策略和GEMM不同，通常采用长方形分块，并结合Repeat参数减少指令发射次数。

量化算子则更复杂一些。它的核心是将高精度浮点数转换为低精度整数（比如FP32转INT8），从而减少内存占用和计算量。比如，在量化算子的后处理（Epilogue）中，您需要将整数结果乘以缩放因子（Scale），还原为浮点数。

本章会通过代码示例展示如何用模板库中的BlockMmad和BlockEpilogue组合出GEMV和量化算子。比如，SplitK方案在GEMV中同样适用，但需要结合AtomicAdd（原子加法）来合并多核结果。

本章还通过对比不同Repeat参数的优化效果，说明如何选择最佳分块策略。比如Repeat=4时，搬运次数减少到原来的1/4，但计算延迟也同步降低。这种权衡需要结合具体场景，比如在模型推理中，更倾向于牺牲少量精度换取更快的速度。

4.1 GEMV分块与Repeat优化

GEMV（General Matrix Vector multiplication）是一种特殊的矩阵乘向量操作，其核心在于如何高效地将矩阵与向量相乘。与GEMM的方形分块不同，GEMV通常采用长方形分块，以适应向量的计算需求。

4.1.1 分块策略

GEMV的分块策略主要分为两种：

  • 方形分块：适用于向量维度较小的场景，分块大小为 M 0 × K 0 M_0 \times K_0 M0​×K0​。
  • 长方形分块：适用于向量维度较大的场景，分块大小为 M 0 × K 0 M_0 \times K_0 M0​×K0​，但通过Repeat参数减少指令发射次数。
Repeat参数的作用

Repeat参数可以减少指令的发射次数，从而降低计算延迟。例如，在处理一个大小为 M 0 × K 0 M_0 \times K_0 M0​×K0​ 的矩阵块时，Repeat参数可以使得每个指令能够处理多个数据块，从而减少搬运与计算的总次数。

4.1.2 Repeat参数的优化

通过合理设置Repeat参数，可以显著提升GEMV的性能。例如，使用Add指令时，Repeat参数可以设置为2，从而减少搬运次数。

表格：Repeat参数优化效果对比
Repeat参数 搬运次数 计算延迟 性能提升
1 16 128us 20%
2 8 64us 40%
4 4 32us 60%

4.2 GEMV的SplitK方案

GEMV的SplitK方案类似于GEMM的SplitK优化，通过在K方向上进行分块，进一步增加多核并行性，提升计算效率。

4.2.1 SplitK的原理

在GEMV中，SplitK通过将向量的K维度分块，使得每个AI Core能够并行处理不同的K分块。每个AI Core在计算时，只需要处理一个K分块，从而减少计算延迟。

示例代码
// SplitK方案示例
void GEMVSplitK()
{
    LocalTensor<float> src0Local = inQueueSrc0.AllocTensor<float>();
    LocalTensor<float> src1Local = inQueueSrc1.AllocTensor<float>();
    LocalTensor<float> dstLocal = outQueueDst.AllocTensor<float>();

    // 分块计算
    for (int k = 0; k < K; k += K_TILE)
    {
        // 读取A矩阵的分块
        DataCopy(src0Local, aGM[k * K_TILE], K_TILE);

        // 读取B矩阵的分块
        DataCopy(src1Local, bGM[k * K_TILE], K_TILE);

        // 矩阵乘
        Mmad(dstLocal, src0Local, src1Local, m, n, K_TILE);

        // 将结果写回GM
        DataCopy(cGM[k * K_TILE], dstLocal, m * n);
    }

    // 原子加操作
    AtomicAdd(cGM, dstLocal, m * n);
}

4.2.2 SplitK的实现步骤

  1. 数据分块：将向量的K维度分块为多个小块，每个小块由不同的AI Core并行处理。
  2. 数据搬运：使用DataCopy将每个K分块搬运到Local Memory。
  3. 矩阵乘计算：调用Mmad指令进行矩阵乘操作。
  4. 结果合并：使用AtomicAdd将多个AI Core的计算结果合并到全局内存中。

4.3 GEMV的计算流程

GEMV的计算流程分为以下几个步骤：

  1. 数据搬运：将矩阵A和B的分块搬运到Local Memory。
  2. 矩阵乘操作：使用Mmad指令进行矩阵乘操作。
  3. 向量累加：通过AtomicAdd指令进行向量累加。
  4. 结果输出：将结果从Local Memory搬运到Global Memory。
示例代码
// GEMV计算流程
void GEMVCompute()
{
    LocalTensor<float> aLocal = inQueueA.AllocTensor<float>();
    LocalTensor<float> bLocal = inQueueB.AllocTensor<float>();
    LocalTensor<float> cLocal = outQueueC.AllocTensor<float>();

    // 搬运A矩阵分块
    DataCopy(aLocal, aGM[0], K_TILE);

    // 搬运B矩阵分块
    DataCopy(bLocal, bGM[0], K_TILE);

    // 矩阵乘
    Mmad(cLocal, aLocal, bLocal, m, n, K_TILE);

    // 原子加操作
    AtomicAdd(cGM, cLocal, m * n);

    // 搬运结果回GM
    DataCopy(cGM, cLocal, m * n);
}

4.3.1 AIV核的向量累加

在AIV核上，可以通过AtomicAdd指令实现向量累加。这种方式适用于需要在多核上进行结果累加的场景。

示例代码
// AIV核的向量累加
void AIVGEMV()
{
    LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
    LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
    LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();

    // 向量累加
    Add(zLocal, xLocal, yLocal, m * n);

    // 原子加操作
    AtomicAdd(zGM, zLocal, m * n);

    // 搬运结果
    DataCopy(zGM, zLocal, m * n);
}

4.4 量化算子设计

当前模型由于参数量庞大，大多数都需要使用量化技术来加快推理速度的同时保证精度不会下降太多。量化可以降低模型的存储空间、加速推理计算和降低功耗。

4.4.1 量化与反量化

量化是将浮点数转换为整数，反量化则是将整数还原为浮点数。常见的量化算法包括对称量化和非对称量化。量化算子中多使用对称量化方式，公式如下：

对称量化公式

X i n t = r o u n d ( X f l o a t S c a l e ) X_{int}=round(\frac{X_{float}}{Scale}) Xint​=round(ScaleXfloat​​),其中 S c a l e Scale Scale表示缩放因子

反量化公式

X f l o a t = X i n t ∗ S c a l e X_{float}=X_{int}*Scale Xfloat​=Xint​∗Scale

4.5 量化算子模板库实现

量化算子的实现可以基于模板库中的组件快速完成。通过合理使用模板库中的BlockMmad和BlockEpilogue，可以实现高效的量化算子。对应的量化参数封装在对应的结构体中，通常的量化类型参数包括每个向量或者通道的缩放因子Scale，还有适用于transformer每个token的PerTokenScale在kernel层组装普通的block mmad和block epilogue，epilogue中传入获取到的Scale或PerTokenScale参数进行运算。

核心代码
// 参数结构体中传入对应的量化参数
struct Params {
        // Data members
        GemmCoord problemShape;
        __gm__ ElementA *ptrA;
        LayoutA layoutA;
        __gm__ ElementB *ptrB;
        LayoutB layoutB;
        __gm__ ElementScale *ptrScale;
        LayoutScale layoutScale;
        __gm__ ElementPerTokenScale *ptrPerTokenScale;
        LayoutPerTokenScale layoutPerTokenScale;
        __gm__ ElementD *ptrD;
        LayoutD layoutD;
        GM_ADDR ptrWorkspace;

        // Methods
        ACT_DEVICE
        Params() {}

        ACT_DEVICE
        Params(
            GemmCoord problemShape_,
            GM_ADDR ptrA_, LayoutA layoutA_,
            GM_ADDR ptrB_, LayoutB layoutB_,
            GM_ADDR ptrScale_, LayoutScale layoutScale_,
            GM_ADDR ptrPerTokenScale_, LayoutPerTokenScale layoutPerTokenScale_,
            GM_ADDR ptrD_, LayoutD layoutD_,
            GM_ADDR ptrWorkspace_
        ) : problemShape(problemShape_),
            ptrA(reinterpret_cast<__gm__ ElementA *>(ptrA_)), layoutA(layoutA_),
            ptrB(reinterpret_cast<__gm__ ElementB *>(ptrB_)), layoutB(layoutB_),
            ptrScale(reinterpret_cast<__gm__ ElementScale *>(ptrScale_)), layoutScale(layoutScale_),
            ptrPerTokenScale(reinterpret_cast<__gm__ ElementPerTokenScale *>(ptrPerTokenScale_)),
            layoutPerTokenScale(layoutPerTokenScale_),
            ptrD(reinterpret_cast<__gm__ ElementD *>(ptrD_)), layoutD(layoutD_),
            ptrWorkspace(ptrWorkspace_) {}
    };

// 核心实现逻辑，矩阵乘计算过后通过Epilogue来完成反量化过程
// GEMM过程
void operator()<AscendC::AIC>(Params const &params)
    {
        BlockScheduler blockScheduler;
        blockScheduler.Update(params.problemShape, MakeCoord(L1TileShape::M, L1TileShape::N));
        uint32_t coreLoops = blockScheduler.GetCoreLoops();

        BlockMmad blockMmad(resource);

        // Represent the full gm
        AscendC::GlobalTensor<ElementA> gmA;
        gmA.SetGlobalBuffer(params.ptrA);
        AscendC::GlobalTensor<ElementB> gmB;
        gmB.SetGlobalBuffer(params.ptrB);
        AscendC::GlobalTensor<ElementC> gmC;
        gmC.SetGlobalBuffer(reinterpret_cast<__gm__ ElementC *>(params.ptrWorkspace));
        layout::RowMajor layoutC(params.problemShape.m(), params.problemShape.n());

        uint32_t coreIdx = AscendC::GetBlockIdx();
        uint32_t coreNum = AscendC::GetBlockNum();

        AicFinishSync aicFinishSync{this};

        for (uint32_t loopIdx = coreIdx; loopIdx < coreLoops; loopIdx += coreNum) {
            // Compute block location
            GemmCoord blockCoord = blockScheduler.GetBlockCoord(loopIdx);
            GemmCoord actualBlockShape = blockScheduler.GetActualBlockShape(blockCoord);

            // Compute initial location in logical coordinates
            MatrixCoord offsetA{blockCoord.m() * L1TileShape::M, blockCoord.k() * L1TileShape::K};
            MatrixCoord offsetB{blockCoord.k() * L1TileShape::K, blockCoord.n() * L1TileShape::N};
            MatrixCoord offsetC{blockCoord.m() * L1TileShape::M, blockCoord.n() * L1TileShape::N};
            int64_t gmOffsetA = params.layoutA.GetOffset(offsetA);
            int64_t gmOffsetB = params.layoutB.GetOffset(offsetB);
            int64_t gmOffsetC = layoutC.GetOffset(offsetC);

            // Compute block-scoped matrix multiply-add
            if constexpr (BlockMmad::DispatchPolicy::ASYNC) {
                blockMmad(
                    gmA[gmOffsetA], params.layoutA,
                    gmB[gmOffsetB], params.layoutB,
                    gmC[gmOffsetC], layoutC,
                    actualBlockShape, MakeCallback(&aicFinishSync)
                );
            } else {
                blockMmad(
                    gmA[gmOffsetA], params.layoutA,
                    gmB[gmOffsetB], params.layoutB,
                    gmC[gmOffsetC], layoutC,
                    actualBlockShape
                );
                aicFinishSync();
            }
        }

        if constexpr (BlockMmad::DispatchPolicy::ASYNC) {
            blockMmad.SynchronizeBlock();
        }
    }

// 反量化过程
template <>
    ACT_DEVICE
    void operator()<AscendC::AIV>(Params const &params)
    {
        BlockScheduler blockScheduler;
        BlockEpilogue blockEpilogue(resource);

        uint32_t coreIdx = AscendC::GetBlockIdx() / AscendC::GetSubBlockNum();
        uint32_t coreNum = AscendC::GetBlockNum();
        uint32_t subCoreIndex = AscendC::GetSubBlockIdx();

        AscendC::GlobalTensor<ElementC> gmC;
        gmC.SetGlobalBuffer(reinterpret_cast<__gm__ ElementC *>(params.ptrWorkspace));

        AivWaitSync aicFinishSync{this};

        LayoutC layoutC = LayoutC(params.problemShape.m(), params.problemShape.n());
        LayoutScale layoutScale = params.layoutScale;
        LayoutPerTokenScale layoutPerTokenScale =
            params.layoutPerTokenScale.GetTileLayout(params.problemShape.template GetCoordByAxis<0>());
        LayoutD layoutD = params.layoutD.GetTileLayout(params.problemShape.GetCoordMN());

        EpilogueParams epilogueParams{
            params.ptrScale, layoutScale,
            params.ptrPerTokenScale, layoutPerTokenScale,
            params.ptrD, layoutD
        };

        blockScheduler.Update(params.problemShape, L1TileShape::ToCoordMN());
        blockEpilogue.UpdateParams(epilogueParams);
        uint32_t coreLoops = blockScheduler.GetCoreLoops();

        GemmCoord blockShapeMNK = L1TileShape::ToCoord();
        for (uint32_t loopIdx = coreIdx; loopIdx < coreLoops; loopIdx += coreNum) {
            GemmCoord blockCoordMNK = blockScheduler.GetBlockCoord(loopIdx);
            GemmCoord actualBlockShapeMNK = blockScheduler.GetActualBlockShape(blockCoordMNK);

            auto gmBlockC = gmC[layoutC.GetOffset(blockCoordMNK.GetCoordMN() * blockShapeMNK.GetCoordMN())];
            auto layoutBlockC = layoutC.GetTileLayout(actualBlockShapeMNK.GetCoordMN());

            blockEpilogue(
                    blockShapeMNK, blockCoordMNK,
                    actualBlockShapeMNK, gmBlockC,
                    layoutBlockC, MakeCallback(&aicFinishSync)
                );
        }
    }

4.5.3 量化算子的性能收益

量化算子的性能收益主要体现在以下几点：

  • 减少计算复杂度：通过使用低精度的整数运算，降低计算复杂度。
  • 减少内存占用：量化后的数据占用更少的内存空间，提升内存利用率。

通过以上章节的详细讲解，开发者可以快速掌握GEMV与量化算子的开发步骤，并结合模板库中的组件进行高效的的算子实现。希望这些内容能够为您的开发提供有价值的的指导和支持。

第5章 Epilogue实现

Epilogue是算子的“收尾动作”，比如在矩阵乘结果后添加偏置、激活函数或归约操作。这一章教会您如何把这些后处理模块化——您可以把Epilogue看作是计算流水线的“最后工序”，它与核心计算逻辑解耦，但又能灵活扩展。比如，教程中的Add操作示例展示了如何将Bias向量逐元素相加到矩阵结果上，而Cast操作则负责数据类型的转换（如从浮点转为整数）。

Epilogue的实现需要结合Vector指令（向量指令）和Scalar计算。教程中的性能数据表明，使用Counter模式（自动计数）能减少指令发射的Scalar开销。比如，处理256个元素时，Counter模式的耗时比手动计算模式降低35%。此外，Epilogue还支持多算子融合——比如将多个后处理步骤合并到同一个核函数中，减少中间结果的搬运次数。

本章的高级内容涉及GroupGEMM的开发，即同时处理多个矩阵乘任务。通过统一的分块策略和内存管理，开发者可以高效地将多个任务分配到不同核心并行计算。这里的关键是确保每个核心的数据分块逻辑独立，同时协调多核间的同步。

5.1 Epilogue功能与实现

Epilogue是昇腾算子模板库中用于处理矩阵乘计算后逐元素操作的关键模块。它与Block层级的矩阵乘加主循环（如BlockMmad）解耦，允许开发者灵活组合不同后处理逻辑（如加Bias、反量化、激活函数等），为算子融合和性能优化提供基础。

Epilogue的核心特性

特性 描述
独立性 与核心矩阵乘加逻辑分离，避免重复代码
模板化 通过模板参数（如数据类型、计算策略）自定义后处理功能
灵活性 支持逐元素操作（Element-wise）、广播（Broadcast）、规约（Reduce）等场景
融合优化 与Vector指令结合时，可减少核间同步开销

标准Epilogue实现示例

namespace CATLASS::epilogue::block {

template <
    class CType_,
    class XType_,
    class DType_,
    class TileElemWiseEpilogue_,
    class TileCopy_
>
class BlockEpilogue <
    EpilogueAtlasA2ElemWiseOneSource,
    CType_,
    XType_,
    DType_,
    TileElemWiseEpilogue_,
    TileCopy_
> {
    // 实现逐元素操作逻辑
};
};

常见Epilogue场景对比

场景 功能描述
Add操作 矩阵乘结果与Bias向量逐元素相加
Cast操作 反量化（如FP16→INT8）
ReLU激活 按照阈值进行非线性变换
ReduceAdd 矩阵行内元素求和，输出一维向量

Add核心代码

using BlockEpilogue = CATLASS::epilogue::block::BlockEpilogue<...>;
BlockEpilogue epilogue;
epilogue.operator()(localC, localBias, m, n);

Cast核心代码

Cast(dstLocal, srcLocal, RoundMode::CAST_NONE, cSize);

激活函数核心代码

AscendC::ReLU(dstLocal, srcLocal, cSize);

Reduce核心代码

ReduceAdd(dstTensor, srcTensor, m, n, k);

5.2 常见Epilogue实现方案

5.2.1 LinearCombination和LinearCombinationGeneric

  • LinearCombination: 通过模板参数来指定分块的大小，默认分块大小与Gemm分块大小一致。首先需要从gm中读取source矩阵分块通过AIV进行乘β操作，然后从AIC和AIV同步GM临时空间读取AIC计算后的分块矩阵进行乘α操作，最后进行相加。大致流程如下：
    1. 将源矩阵C读入到UB，并执行mul指令。
    2. 将矩阵乘结果块从临时空间读入到UB，进行mul指令。
    3. 将前面运算得到的结果进行Add操作
  • LinearCombinationGeneric，与LinearCombination类似，AIV对LineraCombation得到的结果调用激活函数，再存回GM

5.2.2 EpilogueWithBroadCast

需要将偏置向量进行广播到与矩阵一致的维度，然后求两者之和。大体流程如下：

  • 首先将bias向量读入到UB，维度为(m,1)
  • AIV对bias向量进行broadcast，维度变为(m, n)；同时此时将源矩阵读入到UB，维度为(m, n)
  • AIV进行Add操作，并将结果存回GM

5.3 多算子GroupGEMM开发

GroupGEMm是昇腾模板库中实现多矩阵并行计算的高阶接口。通过统一的分块策略（Tiling）和内存管理（如aclrtMalloc），开发者可以高效地将多个Matmul任务分配到AICore上并行执行。

GroupGEMM开发流程

  1. 定义分块策略

    • 通过TilingData结构体定义每个矩阵的分块维度（BlockTileM, BlockTileN）
    • 支持动态shape的自动推导（infer_shape）

  2. 组装Kernel逻辑

    • 使用CATLASS::matmul::kernel::BasicMatmul封装多个BlockMmad和BlockEpilogue组件

    • 示例代码：

      using BlockMmad = matmul::block::BlockMmad<DispatchPolicy, L1TileShape, L0TileShape, AType, BType, CType>;
using BlockEpilogue = matmul::block::BlockEpilogue<...>;
using MatmulKernel = matmul::kernel::BasicMatmul<BlockMmad, BlockEpilogue, TileScheduler>;

  3. Host侧适配

    • 通过MatmulUniversalAdapter接口传递算子参数

    • 示例代码：

      CATLASS::MatmulUniversalAdapter<MatmulKernel> matmulAdapter;
matmulAdapter(args, workspace, stream);

多算子同步与内存管理

  • 多核同步：PipeBarrier&lt;PIPE_ALL&gt;()确保核间计算顺序

  • 内存排布优化：通过TilingData字段控制数据在GM-L1-L0C层级的搬运策略

  • 典型数据流：

    GM -> L1 -> L0C -> GM

  • 关键API：

    LocalTensor<fp16> c1Local = outQueueCO1.AllocTensor<fp16>();
PipeBarrier<PIPE_M>();
Mmad(c1Local, a2Local, b2Local, mmadParams);

5.4 常见问题排查

问题1：Epilogue导致流水冲突

  • 现象：aiv_vec_time占比过高
  • 解决：检查TQue的BUFFER_num配置，确保BUFFER_NUM=2启用Double Buffer

问题2：多算子分块不均

  • 现象：aic_mte2_ratio低于80%
  • 解决：调整TilingData的last_problem_size前缀和计算方式

通过Epilogue模块化设计和GroupGEMm的并行化能力，开发者可以快速实现算子融合，减少流水线空闲时间。建议在模板库开发初期优先使用MatmulUniversalAdapter等封装接口，逐步深入TileCopy和BlockSwizzle定制化逻辑。

第6章 性能优化核心技巧

开发算子的目标不仅是“能跑”，更是“跑得快”。这一章会聚焦于如何让昇腾算子的性能接近理论极限。

性能优化的第一步是减少数据搬运次数。比如，一次性搬运大块数据比多次小搬运更高效。此外，数据搬运的地址对齐（如512B对齐）能显著提升带宽利用率。

第二步是内存管理。昇腾的局部内存（L1/L0）采用bank结构，多个核同时访问同一bank会导致冲突。通过调整数据排布（如间隔8字节）或使用Double Buffer，可以规避冲突。

第三步是流水线优化。比如，在Vector计算中启用Counter模式，能减少Scalar指令的开销，从而提升速度。

最后，本章会通过对比Normal模式和Counter模式的耗时，说明不同策略的性能差异。比如，Counter模式的Scalar指令耗时可能比Normal模式降低17%。

6.1 数据搬运优化

数据搬运直接影响算子性能，尤其是在多核并行或大矩阵计算中。优化策略的核心在于减少搬运次数、提升单次搬运效率。

优先级对比

优化策略 优先级 描述 适用场景
一次性搬运大块数据 高 单次搬运16KB以上能更高效利用带宽 所有搬运场景
512B地址对齐 高 保证GM与Local Memory搬运时地址对齐 GM->UB、GM->L1等搬运
使用高效搬运API 高 通过srcStride、dstStride等参数实现固定间隔搬运 非连续数据搬运
避免冗余搬运 中 通过合并搬运逻辑减少重复调用 多步骤计算中

示例：GM->UB高效搬运

DataCopyParams copyParams;
copyParams.blockCount = imgHeight;
copyParams.blockLen = copyWidth / 8;
copyParams.srcStride = (imgWidth - copyWidth) / 8;
copyParams.dstStride = 0;
DataCopy(tensorIn, tensorGM, copyParams);

优化要点：

  • blockLen和blockCount定义搬运粒度，单位为DataBlock（32Byte）。
  • srcStride和dstStride控制搬运间隔，单位为DataBlock。
  • 一次性搬运比多次小量搬运节省更多时间。

6.2 内存管理与bank冲突规避

昇腾NPU的Local Memory（如UB、L1）采用bank结构，多核并行时bank冲突会导致性能下降。合理分配内存空间，可显著提升运算效率。

6.2.1 bank冲突分类

冲突类型 描述
读写冲突 读写同时访问同一bank
写写冲突 多个写操作同时访问同一bank
读读冲突 多个读操作同时访问同一bank

6.2.2 bank冲突规避策略

策略 优化效果
8字节对齐 避免bank内地址错位
避免主块步长为16 防止写冲突集中
合并bank group使用 提升bank利用率

示例：优化bank group使用

LocalTensor<float> src0Local = inQueueSrc0.AllocTensor<float>();
LocalTensor<float> src1Local = inQueueSrc1.AllocTensor<float>();
LocalTensor<float> dstLocal = outQueueDst.AllocTensor<float>();

AscendC::Add(dstLocal, src0Local, src1Local, mask, 1, addParams);

优化要点：

  • src0Local和src1Local的地址间隔为8，避免落入同一bank group。
  • dstLocal的地址间隔为32，确保bank group利用率最大化。

6.3 流水线优化与AtomicAdd

流水线优化的核心在于并行执行搬运、计算和同步操作。通过合理利用TQue队列和TPipe同步模块，可以降低流水线等待时间。

6.3.1 double buffer机制

pipe.InitBuffer(inQueueX, 2, sizeX);

优化要点：

  • BUFFER_NUM=2启用double buffer，使搬运和计算并行。
  • 搬运和计算交替进行，减少流水等待时间。
  • 适用于Vector计算与GM交互的场景。

6.3.2 AtomicAdd优化

// 启用AtomicAdd，减少搬运次数
mm.IterateAll(gm_c, 1);

优化要点：

  • 通过IterateAll接口的enAtomic参数，直接将Cube计算结果累加至GM。
  • 避免中间结果多次搬运，适用于向量加法等场景。
  • 在Atlas A2系列设备上性能收益显著。

6.4 Vector指令与Counter模式

Vector指令的性能优化主要集中在减少Scalar计算量和提升指令并行能力。Counter模式可显著简化代码逻辑，避免手动处理尾块。

6.4.1 Normal模式与Counter模式对比

模式 描述 优势 示例代码
Normal模式 需要手动处理主块与尾块 灵活处理非对齐数据 
AscendC::SetMaskNorm();
AscendC::SetVectorMask<float>(0, totalLength);
AscendC::Add(zLocal, xLocal, yLocal, totalLength);

| Counter模式 | 自动计算总元素个数 | 简化代码，减少Scalar计算 |

AscendC::SetMaskCount();
AscendC::SetVectorMask<float>(0, totalLength);
AscendC::Add(zLocal, xLocal, yLocal, totalLength);

性能数据：

  • Normal模式：aiv_scalar_time ≈ 281us
  • Counter模式：aiv_scalar_time ≈ 236us

优化要点：

  • 对于连续数据计算，推荐使用Counter模式，减少Scalar开销。
  • 若数据非对齐，需结合dataBlockStride和repeatStride进行调整。

6.5 归约指令优化

归约指令（Reduce）在处理大矩阵时，容易成为性能瓶颈。合理选择WholeReduce或BlockReduce，结合分块策略，可优化性能。

6.5.1 归约策略选择

归约类型 适用场景 优化建议
WholeReduce 单条指令归约整个数据 适用于小数据量，如累加256元素
BlockReduce 按块归约 适用于大矩阵计算，减少流水延迟

示例：WholeReduce优化

AscendC::WholeReduceSum<float, false>(zLocal, xLocal, totalLength);

优化要点：

  • WholeReduceSum对Vector指令的流水利用率比BlockReduceSum更高。
  • 对于256元素的float数据，使用WholeReduceSum+BlockReduceSum组合，总耗时从13us降至8.44us。

6.6 Tiling策略优化

Tiling决定了算子的执行次数和数据切分方式。优化Tiling策略可提升L2Cache命中率，减少GM访问。

6.6.1 L2Cache切分策略

constexpr int32_t TILE_NUM = 2;
constexpr int32_t TOTAL_LENGTH = 384 * 1024 * 1024 / sizeof(half);

优化要点：

  • 当输入数据量超过L2Cache容量时，启用Tiling切分。
  • TILE_NUM=2将数据分两次处理，确保每次数据可完全命中L2Cache。

6.6.2 核间负载均衡

反例：

// 核间负载不均，部分核空闲
for (int32_t i = 0; i < round * 2; ++i) {
  CopyIn(i);
  Compute(i);
  CopyOut(i);
}

正例：

// 通过调整`TILE_NUM`和`USE_CORE_NUM`实现负载均衡
for (int32_t i = 0; i < round; ++i) {
  CopyIn(i);
  Compute(i);
  CopyOut(i);
}

优化建议：

  • 计算量应均匀分配，避免部分核空闲。
  • 小数据量场景中，可合并多个分块处理，提升利用率。

6.7 内存共享优化

昇腾NPU的Local Memory资源有限，合理复用内存可减少冗余搬运，提升算子性能。

6.7.1 临时Buffer共享

示例：减少UB空间占用

TBuf<QuePosition::VECCALC> tmpSharedBuf;
pipe.InitBuffer(tmpSharedBuf, bufferSize);

LocalTensor<uint8_t> tmpSharedTensor = tmpSharedBuf.Get<uint8_t>(softmaxBufSize);
LocalTensor<float> tmpSumTensor = tmpSharedBuf.Get<float>(sumBufSize);

SoftMax<float, true, true>(dstTensor, expSumTensor, dstMaxTensor, srcTensor, tmpSharedTensor, tiling);
Add<float>(tmpSumTensor, src0Tensor, src1Tensor, count);

优化要点：

  • 通过tmpSharedBuf复用内存，减少UB空间占用。
  • 避免多次分配临时Buffer，节省时间开销。

6.8 总结

性能优化是一个系统工程，需结合数据搬运、内存管理、流水线调度和指令模式进行综合分析。建议开发者遵循以下步骤：

  1. 通过msprof op工具获取性能数据，识别瓶颈。
  2. 优先优化大块搬运和地址对齐。
  3. 使用Counter模式简化Vector计算。
  4. 调整Tiling策略，提升L2Cache利用率。
  5. 最后通过bank冲突分析优化内存排布。

通过不断迭代以上步骤，您将逐步逼近昇腾算子的理论性能上限。

第7章 算子调试与测试

算子开发完成后，必须经过严格的调试和测试，否则可能会出现精度错误或性能瓶颈。这一章会教您如何验证算子的正确性，并通过工具分析性能问题。

首先，Host侧的精度验证是关键。您需要将昇腾设备上的结果拷贝回Host，用numpy对比误差。比如，若FP16的误差超过千分之一，就需要检查搬运逻辑或计算顺序。

其次，性能分析工具（msprof op）能生成流水线图，显示Cube和Vector的执行情况。比如，如果发现aic_mte2_time异常高，可能是搬运未对齐导致的带宽浪费。

最后，ST测试用例生成工具（msopst）能帮您自动化测试算子。比如，定义不同规模的矩阵输入，生成结构化测试用例，再运行对比结果。

本章会强调：在昇腾上调试算子时，先用CPU仿真环境验证逻辑，再部署到真实设备。这能大幅减少因地址错位导致的崩溃风险。

7.1 Host侧功能验证

7.1.1 同步控制与多核调试

关键操作：多核同步时，需确保BlockDim不超过实际核心数（GetCoreNum()）。否则框架会插入异常同步导致卡死。

正例：

// 正确示例
FlashAttentionScoreApiTiling(tilingData);
FlashAttentionScoreGetTensorSize(tilingData);
context->SetBlockDim(MAX_AICORE_NUM);

反例：

// 错误示例
FlashAttentionScoreApiTiling(tilingData);
FlashAttentionScoreGetTensorSize(tilingData);
context->SetBlockDim(MAX_AICORE_NUM + 1);

注意事项：

  • 多核同步：若算子涉及多核间数据依赖（如累加），需显式插入SetFlag/WaitFlag控制流水线同步。
  • 核间通信：确保多核间数据分块逻辑合理，避免因地址偏移错误导致覆盖或未初始化数据。

7.1.2 地址偏移与数据分块

核心步骤：

  1. 分块逻辑：Host侧通过TilingFunc计算分块参数（如formerNum、tailNum）。
  2. 地址偏移：每个核根据GetBlockIdx()获取独立的Global Memory地址，避免跨核数据冲突。

示例：

// 多核地址设置
xGm.SetGlobalBuffer((__gm__ float*)x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);

正例与反例对比：

场景 反例 正例
分块冗余 TilingDataUnalign包含冗余字段 TilingDataUnalign字段精简
地址对齐 TilingDataUnalign字段未对齐 TilingDataUnalign字段8字节对齐

7.1.3 浮点精度验证

关键原则：

  • 精度标准：浮点运算需容忍误差（如half类型误差不超过千分之一）。
  • 计算顺序：避免因并行计算导致的精度漂移（如(a + b) + c与a + (b + c)结果差异）。

验证代码：

ACL_CHECK(aclrtMemcpy(h_C, C_size, d_C, C_size, ACL_MEMCPY_DEVICE_TO_HOST));

int errorCount = 0;
for (int i = 0; i < length_m * length_n; i++) {
    float diff_res = (float)(* ((__fp16 *)h_C + i)) - (float)length_k;
    if (fabs(diff_res) > (float)length_k * 0.001) {
        errorCount++;
    }
}

if (errorCount < length_m * length_n * 0.001) {
    std::cout << "[Compare success]" << std::endl;
} else {
    std::cout << "[Compare failed]" << std::endl;
}

注意事项：

  • 若使用Cast转换低精度数据（如bfloat16到float），需确保转换逻辑与目标硬件一致。
  • 复合指令：如Axp（乘加合一）可能导致与分步计算的精度差异，需预先测试。

7.2 性能数据采集与分析

7.2.1 上板Profiling

工具：msprof op采集硬件级性能数据（如aic_mte2_time）。

采集命令：

msprof op --output="./out" --ai-core=on --aic-metrics="PipeUtilization" add_custom_npu

数据解析：

  • 瓶颈识别：
    • MTE2/MTE3带宽利用率：若低于95%，需优化数据搬运对齐（如512B对齐）。
    • Cube利用率：若aic_mac_ratio低于80%，需检查分块粒度（如BlockTileM、BlockTileN是否合理）。

对比示例：

算子类型 理论耗时 实际耗时 性能差距 优化方向
GEMM 111.8 us 37.28 us 30% 优化L1缓存命中率
Vector Add 335.5 us 350 us 4.4% 优化Double Buffer

7.2.2 仿真流水图分析

工具：msprof op生成trace.json，通过Chrome浏览器分析流水线冲突。

分析步骤：

  1. 加载文件：

    chrome://tracing/ # 拖入trace.json

  2. 关键指标：

    • 流水线利用率：观察Cube、Vector、MTE的并行性，若存在断流需检查数据依赖。
    • 头开销：Scalar计算前的同步或搬运延迟（如aic_mte2_time与aiv_vec_time对比）。

典型问题：

  • bank冲突：多核访问同一bank组导致流水线阻塞（需通过bank_group_id分析）。
  • 搬运等待：GM->L1未对齐导致带宽未充分利用（需调整DataCopyParams）。

7.3 ST测试用例生成

7.3.1 测试用例定义

工具：msopst生成结构化测试用例（如AddCustom_case.json）。

定义示例：

[
  {
    "case_name": "Test_AddCustom_001",
    "op": "AddCustom",
    "input_desc": [
      {
        "format": ["ND"],
        "type": ["float16"],
        "shape": [8, 2048],
        "data_distribute": ["uniform"],
        "value_range": [[0.1, 1.0]],
        "name": "x"
      }
    ],
    "output_desc": [
      {
        "format": ["ND"],
        "type": ["float16"],
        "shape": [8, 2048],
        "name": "z"
      }
    ]
  }
]

7.3.2 测试执行与验证

步骤：

  1. 配置环境变量：

    export DDK_PATH=${INSTALL_DIR}
export NPU_HOST_LIB=${INSTALL_DIR}/{arch-os}/devlib

  2. 运行测试：

    cd $HOME/Ascend/ascend-toolkit/latest/python/site-packages/bin
./msopst run -i $HOME/AddCustom_st/AddCustom_case.json -soc <soc_version> -out $HOME/

  3. 结果验证：

    // 测试报告生成路径
std::string reportPath = "xxxx/AddCustom_st/20230828202015/st_report.json";

精度验证：

  • half类型误差：numpy对比误差需小于0.1%（如numpy.allclose(h_C, npy, rtol=1e-3)）。
  • 结果一致性：若Host与Device结果差异过大，需检查SetGlobalBuffer地址是否正确。

7.4 调试工具链对比

工具 功能 适用场景 性能收益
msprof op 采集硬件流水线数据 性能瓶颈分析 优化搬运/计算流水利用率
msopst 生成ST测试用例 精度验证与回归测试 降低测试开销
gdb CPU仿真调试 初期功能开发 快速定位逻辑错误

7.5 常见问题与解决方案

7.5.1 参数修改陷阱

反例：

query = tmpQueryPtr; // 会覆盖输入参数，导致结果错误

正例：

inputQueryGMTensor.SetGlobalBuffer(query);

7.5.2 TilingData优化建议

精简字段：

  • 冗余字段：blockDim可通过SetBlockDim传递，无需写入结构体。
  • 对齐排布：formerNum、tailNum建议使用uint8_t并8字节对齐。

示例结构体：

BEGIN_TILING_DATA_DEF(TilingDataUnalign)
TILING_DATA_FIELD_DEF(uint8_t, formerNum);
TILING_DATA_FIELD_DEF(uint8_t, tailNum);
TILING_DATA_FIELD_DEF(uint32_t, formerLength);
TILING_DATA_END

7.5.3 核间资源冲突

bank冲突规避：

  • 规则：确保LocalTensor访问不同bank组（如src0与src1错开32B）。
  • 工具：通过AscendC::GetBankId检测冲突，调整DataCopyParams步长。

反例：

LocalTensor<float> src0Local = inQueueSrc0.AllocTensor<float>();
LocalTensor<float> src1Local = inQueueSrc1.AllocTensor<float>();
// src0Local与src1Local可能落在同一bank组

正例：

LocalTensor<float> src0Local = inQueueSrc0.AllocTensor<float>();
LocalTensor<float> src1Local = inQueueSrc1.AllocTensor<float>();
// 通过增加src1Local起始地址32B，规避bank冲突

7.6 总结

  • 功能验证：优先确保Host侧Tiling逻辑正确，结合numpy进行精度对比。
  • 性能调优：通过msprof op采集瓶颈，优化数据搬运对齐、分块粒度、bank冲突。
  • ST测试：定义结构化用例（如AddCustom_case.json），自动化验证结果一致性。

提示：调试时先使用CPU仿真，再迁移至NPU，逐步逼近真实场景。

第8章 典型算子开发案例解析

这一章通过三个典型算子（Matmul、QuantMatmul、GroupMatmul）演示如何利用CATLASS模板库快速开发高性能算子。Matmul是最基础的矩阵乘算子，开发者只需选择合适的的模板参数（如分块大小、数据类型）即可生成代码。教程中的代码示例显示，通过不同的分块策略（如使用双缓冲）就能自动适配不同硬件场景。

QuantMatmul的开发则在Matmul基础上增加了反量化步骤。教程中的代码通过模板参数控制量化类型（对称量化或非对称量化），并在Epilogue中插入Cast指令将整数结果转回浮点数。这种设计让量化算子的实现与普通算子几乎完全兼容，只需在模板库中添加少量逻辑。

GroupMatmul用于处理多个不同规模的矩阵乘法。教程中的实现通过动态计算每个矩阵的分块数，并将任务分配到不同核心。比如，当处理10个矩阵时，每个核心独立计算自己分配的分块，最后合并结果。这种并行化能力是昇腾模板库的核心优势，开发者无需手动管理多核调度，而是依赖模板参数自动分配。

8.1 Matmul算子实现

场景背景

矩阵乘是深度学习中的核心算子。昇腾模板库通过各层级组件的协同合作，可快速组装支持多矩阵输入的Matmul

算子，适用于多任务并行计算（如NLP模型的Attention模块）。

实现步骤

  1. 利用CATLASS组件中的Block组件组装对应的Kernel

    #ifndef ACT_GEMM_KERNEL_MATMUL_HPP
#define ACT_GEMM_KERNEL_MATMUL_HPP

#include "act/act.hpp"
#include "act/arch/resource.hpp"
#include "act/coord.hpp"
#include "act/gemm_coord.hpp"
#include "act/matrix_coord.hpp"

namespace Act::Gemm::Kernel {

// Template for Matmul kernel. Compute C = A * B
template <
    class BlockMmad_,
    class BlockEpilogue_,
    class BlockScheduler_
>
class BasicMatmul {
public:
    using BlockMmad = BlockMmad_;
    using ArchTag = typename BlockMmad::ArchTag;
    using L1TileShape = typename BlockMmad::L1TileShape;
    using ElementA = typename BlockMmad::ElementA;
    using LayoutA = typename BlockMmad::LayoutA;
    using ElementB = typename BlockMmad::ElementB;
    using LayoutB = typename BlockMmad::LayoutB;
    using ElementC = typename BlockMmad::ElementC;
    using LayoutC = typename BlockMmad::LayoutC;
    using ElementAccumulator = typename BlockMmad::ElementAccumulator;

    using BlockScheduler = BlockScheduler_;

    /// Parameters structure
    struct Params {
        // Data members
        GemmCoord problemShape;
        GM_ADDR ptrA;
        LayoutA layoutA;
        GM_ADDR ptrB;
        LayoutB layoutB;
        GM_ADDR ptrC;
        LayoutC layoutC;

        // Methods
        ACT_DEVICE
        Params() {}

        ACT_DEVICE
        Params(GemmCoord const &problemShape_, GM_ADDR ptrA_, LayoutA layoutA_, GM_ADDR ptrB_,
               LayoutB layoutB_, GM_ADDR ptrC_, LayoutC layoutC_)
            : problemShape(problemShape_), ptrA(ptrA_), layoutA(layoutA_), ptrB(ptrB_), layoutB(layoutB_),
              ptrC(ptrC_), layoutC(layoutC_) {}
    };

    // Methods
    ACT_DEVICE
    BasicMatmul() {}

    template <int32_t CORE_TYPE = g_coreType>
    ACT_DEVICE
    void operator()(Params const &params);

    /// Executes one Matmul
    template <>
    ACT_DEVICE
    void operator()<AscendC::AIC>(Params const &params) {
        BlockScheduler matmulBlockScheduler(params.problemShape, MakeCoord(L1TileShape::M, L1TileShape::N));
        uint32_t coreLoops = matmulBlockScheduler.GetCoreLoops();

        Arch::Resource<ArchTag> resource;
        BlockMmad blockMmad(resource);

        // Represent the full gm
        AscendC::GlobalTensor<ElementA> gmA;
        gmA.SetGlobalBuffer((__gm__ ElementA *)params.ptrA);
        AscendC::GlobalTensor<ElementB> gmB;
        gmB.SetGlobalBuffer((__gm__ ElementB *)params.ptrB);
        AscendC::GlobalTensor<ElementC> gmC;
        gmC.SetGlobalBuffer((__gm__ ElementC *)params.ptrC);

        for (uint32_t loopIdx = AscendC::GetBlockIdx(); loopIdx < coreLoops; loopIdx += AscendC::GetBlockNum()) {
            // Compute block location
            GemmCoord blockCoord = matmulBlockScheduler.GetBlockCoord(loopIdx);
            GemmCoord actualBlockShape = matmulBlockScheduler.GetActualBlockShape(blockCoord);

            // Compute initial location in logical coordinates
            MatrixCoord offsetA{blockCoord.m() * L1TileShape::M, blockCoord.k() * L1TileShape::K};
            MatrixCoord offsetB{blockCoord.k() * L1TileShape::K, blockCoord.n() * L1TileShape::N};
            MatrixCoord offsetC{blockCoord.m() * L1TileShape::M, blockCoord.n() * L1TileShape::N};
            int64_t gmOffsetA = params.layoutA.GetOffset(offsetA);
            int64_t gmOffsetB = params.layoutB.GetOffset(offsetB);
            int64_t gmOffsetC = params.layoutC.GetOffset(offsetC);

            // Compute block-scoped matrix multiply-add
            blockMmad(gmA[gmOffsetA], params.layoutA,
                      gmB[gmOffsetB], params.layoutB,
                      gmC[gmOffsetC], params.layoutC,
                      actualBlockShape);
        }
    }

    template <>
    ACT_DEVICE
    void operator()<AscendC::AIV>(Params const &params) {}
};

} // namespace Act::Gemm::Kernel

#endif // ACT_GEMM_KERNEL_MATMUL_HPP

    blockMmad具有多个实现，Kernel中需要对特定i的blockMmad进行适配，模板库中的blockMmad都存在gemm/block下，这里使用的是双缓冲的block_mmad。

  2. 通过block级别组件组装完Kernel后，便可以在主机侧进行调用，需要声明对应的block和Kernel组件

    template <
    class LayoutA,
    class LayoutB,
    class LayoutC
>
ACT_GLOBAL
void BasicMatmul(
    GemmCoord problemShape,
    GM_ADDR gmA, LayoutA layoutA,
    GM_ADDR gmB, LayoutB layoutB,
    GM_ADDR gmC, LayoutC layoutC
)
{
    using ArchTag = Arch::AtlasA2;
    using DispatchPolicy = Gemm::MmadAtlasA2Pingpong<true>;
    using L1TileShape = GemmShape<128, 256, 256>;
    using L0TileShape = GemmShape<128, 256, 64>;

    using AType = Gemm::GemmType<half, LayoutA>;
    using BType = Gemm::GemmType<half, LayoutB>;
    using CType = Gemm::GemmType<half, LayoutC>;

    using BlockMmad = Gemm::Block::BlockMmad<DispatchPolicy, L1TileShape, L0TileShape, AType, BType, CType>;
    using BlockEpilogue = void;

    if (problemShape.m() > problemShape.n()) {
        // Swizzle offset is 3 and direction is 0.
        using BlockScheduler = typename Gemm::Block::GemmIdentityBlockSwizzle<3, 0>;

        // kernel level
        using MatmulKernel = Gemm::Kernel::BasicMatmul<BlockMmad, BlockEpilogue, BlockScheduler>;

        typename MatmulKernel::Params params{problemShape, gmA, layoutA, gmB, layoutB, gmC, layoutC};

        // call a kernel
        MatmulKernel matmul;
        matmul(params);
    } else {
        // Swizzle offset is 3 and direction is 1.
        using BlockScheduler = typename Gemm::Block::GemmIdentityBlockSwizzle<3, 1>;

        // kernel level
        using MatmulKernel = Gemm::Kernel::BasicMatmul<BlockMmad, BlockEpilogue, BlockScheduler>;

        typename MatmulKernel::Params params{problemShape, gmA, layoutA, gmB, layoutB, gmC, layoutC};

        // call a kernel
        MatmulKernel matmul;
        matmul(params);
    }
}

性能优化点

  • SplitK方案：当K维度较大时，启用GetCoreNum()划分多核计算，提升利用率。
  • Double Buffer：在TileCopy中设置BUFFER_NUM=2，减少搬运等待时间。
  • 512B对齐：在Fixpipe接口中通过SetFixpipeNz2ndFlag保证地址对齐。

这些优化点对应的组件都已经在CATLASS模板库中实现，详细可参考开源源码。

8.2 QuantMatmul开发全流程

场景背景

量化算子（QuantMatmul）用于将FP32计算结果压缩为INT8，减少存储和计算开销，加快计算速度。

实现步骤

  1. 组装Kernel，与Matmul算子不同之处在于需要传入反量化参数Scale和PerTokenScale，其中反量化过程需要借助blockEpilogue组件完成，CATLASS模板库提供了多个Epilogue组件，存在epilogue目录下。总体实现是在8.1的基础Matmul中添加了AIV上的反量化实现，其中反量化过程为：

    template <>

    ACT_DEVICE
void operator()<AscendC::AIV>(Params const &params)
{
    BlockScheduler blockScheduler;
    BlockEpilogue blockEpilogue(resource);

    uint32_t coreIdx = AscendC::GetBlockIdx() / AscendC::GetSubBlockNum();
    uint32_t coreNum = AscendC::GetBlockNum();
    uint32_t subCoreIndex = AscendC::GetSubBlockIdx();

    AscendC::GlobalTensor<ElementC> gmC;
    gmC.SetGlobalBuffer(reinterpret_cast<__gm__ ElementC *>(params.ptrWorkspace));

    AivWaitSync aicFinishSync{this};

    LayoutC layoutC = LayoutC(params.problemShape.m(), params.problemShape.n());
    LayoutScale layoutScale = params.layoutScale;
    LayoutPerTokenScale layoutPerTokenScale =
        params.layoutPerTokenScale.GetTileLayout(params.problemShape.template GetCoordByAxis<0>());
    LayoutD layoutD = params.layoutD.GetTileLayout(params.problemShape.GetCoordMN());

    EpilogueParams epilogueParams{
        params.ptrScale, layoutScale,
        params.ptrPerTokenScale, layoutPerTokenScale,
        params.ptrD, layoutD
    };

    blockScheduler.Update(params.problemShape, L1TileShape::ToCoordMN());
    blockEpilogue.UpdateParams(epilogueParams);
    uint32_t coreLoops = blockScheduler.GetCoreLoops();

    GemmCoord blockShapeMNK = L1TileShape::ToCoord();
    for (uint32_t loopIdx = coreIdx; loopIdx < coreLoops; loopIdx += coreNum) {
        GemmCoord blockCoordMNK = blockScheduler.GetBlockCoord(loopIdx);
        GemmCoord actualBlockShapeMNK = blockScheduler.GetActualBlockShape(blockCoordMNK);

        auto gmBlockC = gmC[layoutC.GetOffset(blockCoordMNK.GetCoordMN() * blockShapeMNK.GetCoordMN())];
        auto layoutBlockC = layoutC.GetTileLayout(actualBlockShapeMNK.GetCoordMN());

        blockEpilogue(
                blockShapeMNK, blockCoordMNK,
                actualBlockShapeMNK, gmBlockC,
                layoutBlockC, MakeCallback(&aicFinishSync)
            );
    }
}

    重载括号操作符是一个模板函数，模板参数为核心类型，可以看到该实现是在AIV上进行的，由AIV进行反量化过程。

  2. 在主机侧需要声明对应的blockMmad、blockEpilogue组件来组成Kernel。

    void QuantMatmul(
    uint64_t fftsAddr,
    GemmCoord problemShape,
    GM_ADDR gmA, LayoutA layoutA,
    GM_ADDR gmB, LayoutB layoutB,
    GM_ADDR gmScale, layout::VectorLayout layoutScale,
    GM_ADDR gmPerTokenScale, layout::VectorLayout layoutPerTokenScale,
    GM_ADDR gmD, layout::RowMajor layoutD,
    GM_ADDR gmWorkspace
)
{
    AscendC::SetSyncBaseAddr(fftsAddr);
    using ArchTag = Arch::AtlasA2;
    constexpr uint32_t preloadStages = 1;
    constexpr uint32_t l1Stages = 2;
    constexpr uint32_t l0AStages = 2;
    constexpr uint32_t l0BStages = 2;
    constexpr uint32_t l0CStages = 1;
    constexpr bool enableUnitFlag = false;
    constexpr bool enableShuffleK = true;
    using DispatchPolicy = Gemm::MmadAtlasA2PreloadAsyncWithCallback<
        preloadStages,
        l1Stages, l0AStages, l0BStages, l0CStages,
        enableUnitFlag, enableShuffleK
    >;
    using L0TileShape = GemmShape<128, 256, 128>;

    using AType = Gemm::GemmType<int8_t, LayoutA>;
    using BType = Gemm::GemmType<int8_t, LayoutB>;
    using CType = Gemm::GemmType<int32_t, layout::RowMajor>;

    using BlockMmad = Gemm::Block::BlockMmad<DispatchPolicy, L1TileShape, L0TileShape, AType, BType, CType>;

    constexpr uint32_t ubStages = 2;
    using EpilogueDispatchPolicy = Epilogue::EpilogueAtlasA2PerTokenDequant<ubStages>;
    using ScaleType = Gemm::GemmType<bfloat16_t, layout::VectorLayout>;
    using PerTokenScaleType = Gemm::GemmType<bfloat16_t, layout::VectorLayout>;
    using DType = Gemm::GemmType<bfloat16_t, layout::RowMajor>;

    using RowBroadcastMulType = Gemm::GemmType<float, layout::RowMajor>;
    using BroadcastOneBlkType = Gemm::GemmType<float, layout::RowMajor>;
    using OneBlkColumnBroadcastMulType = Gemm::GemmType<float, layout::RowMajor>;

    using EpilogueTileShape = MatrixShape<32, 256>;
    using TileRowBroadcastMul = Epilogue::Tile::TileRowBroadcastMul<ArchTag, RowBroadcastMulType, EpilogueTileShape>;
    using TileBroadcastOneBlk = Epilogue::Tile::TileBroadcastOneBlk<ArchTag, BroadcastOneBlkType,
        EpilogueTileShape::ROW>;
    using TileOneBlkColumnBroadcastMul = Epilogue::Tile::TileOneBlkColumnBroadcastMul<ArchTag,
        OneBlkColumnBroadcastMulType, EpilogueTileShape>;
    using TileCopy = Epilogue::Tile::TileCopy<ArchTag, CType, ScaleType, PerTokenScaleType, DType>;
    using BlockScheduler = Epilogue::Tile::EpilogueHorizontalTileSwizzle;

    using BlockEpilogue = Epilogue::Block::BlockEpilogue<EpilogueDispatchPolicy, CType, ScaleType, PerTokenScaleType,
        DType, TileRowBroadcastMul, TileBroadcastOneBlk, TileOneBlkColumnBroadcastMul, TileCopy, BlockScheduler>;

    if (problemShape.m() > problemShape.n()) {
        using BlockScheduler = typename Gemm::Block::GemmIdentityBlockSwizzle<3, 0>;

        // kernel level
        using MatmulKernel = Gemm::Kernel::QuantMatmulMultiStageWorkspace<BlockMmad, BlockEpilogue,
            BlockScheduler, workspaceStages>;

        typename MatmulKernel::Params params{
            problemShape,
            gmA, layoutA,
            gmB, layoutB,
            gmScale, layoutScale,
            gmPerTokenScale, layoutPerTokenScale,
            gmD, layoutD,
            gmWorkspace
        };

        // call a kernel
        MatmulKernel matmul;
        matmul(params);
    } else {
        using BlockScheduler = typename Gemm::Block::GemmIdentityBlockSwizzle<3, 1>;

        // kernel level
        using MatmulKernel = Gemm::Kernel::QuantMatmulMultiStageWorkspace<BlockMmad, BlockEpilogue,
            BlockScheduler, workspaceStages>;

        typename MatmulKernel::Params params{
            problemShape,
            gmA, layoutA,
            gmB, layoutB,
            gmScale, layoutScale,
            gmPerTokenScale, layoutPerTokenScale,
            gmD, layoutD,
            gmWorkspace
        };

        // call a kernel
        MatmulKernel matmul;
        matmul(params);
    }
}

8.3 GroupMatmul算子开发

场景背景

GroupMatmul是一次完成多个不同规模的矩阵乘法，在实际场景应用广泛，CATLASS模板库为实现新的算子提供了遍历，可以基于已有的接口快速实现。

实现步骤

  1. GroupGEMM的参数结构体封装了矩阵数目、每个矩阵的分块数、每个矩阵的维度以及每个矩阵的排布，每个矩阵的对应参数变为指针传入，在计算时逐个读取，参数结构体如下：

    struct Params {

        // Data members
    uint32_t problemCount;
    GM_ADDR ptrProblemShape;
    GM_ADDR ptrA;
    GM_ADDR ptrLayoutA;
    GM_ADDR ptrB;
    GM_ADDR ptrLayoutB;
    GM_ADDR ptrC;
    GM_ADDR ptrLayoutC;

    // Methods
    ACT_DEVICE
    Params() {}

    ACT_DEVICE
    Params(
        uint32_t problemCount_, GM_ADDR ptrProblemShape_,
        GM_ADDR ptrA_, GM_ADDR ptrLayoutA_,
        GM_ADDR ptrB_, GM_ADDR ptrLayoutB_,
        GM_ADDR ptrC_, GM_ADDR ptrLayoutC_
    ) : problemCount(problemCount_), ptrProblemShape(ptrProblemShape_),
        ptrA(ptrA_), ptrLayoutA(ptrLayoutA_),
        ptrB(ptrB_), ptrLayoutB(ptrLayoutB_),
        ptrC(ptrC_), ptrLayoutC(ptrLayoutC_)
    {
    }
};

  2. Kernel层的具体计算可以复用blockMmad，只需要实现计算地址偏移的逻辑

    template <>
   ACT_DEVICE
   void operator()<AscendC::AIC>(Params const &params)
   {
       GemmCoord problemShapeList[MAX_TENSOR_COUNT];
       LayoutA layoutAList[MAX_TENSOR_COUNT];
       LayoutB layoutBList[MAX_TENSOR_COUNT];
       LayoutC layoutCList[MAX_TENSOR_COUNT];

       // Get matmul information from parameters
       detail::UnpackListParam(problemShapeList, params.ptrProblemShape, params.problemCount);
       detail::UnpackListParam(layoutAList, params.ptrLayoutA, params.problemCount);
       detail::UnpackListParam(layoutBList, params.ptrLayoutB, params.problemCount);
       detail::UnpackListParam(layoutCList, params.ptrLayoutC, params.problemCount);

       BlockScheduler matmulBlockScheduler;
       Arch::Resource<ArchTag> resource;
       BlockMmad blockMmad(resource);

       // Represent the full gm
       AscendC::GlobalTensor<ElementA> gmA;
       gmA.SetGlobalBuffer((__gm__ ElementA *)params.ptrA);
       AscendC::GlobalTensor<ElementB> gmB;
       gmB.SetGlobalBuffer((__gm__ ElementB *)params.ptrB);
       AscendC::GlobalTensor<ElementC> gmC;
       gmC.SetGlobalBuffer((__gm__ ElementC *)params.ptrC);

       uint32_t coreIdx = AscendC::GetBlockIdx();
       uint32_t coreNum = AscendC::GetBlockNum();
       int64_t inGroupOffsetA = 0;
       int64_t inGroupOffsetB = 0;
       int64_t inGroupOffsetC = 0;

       uint32_t startCoreIdx = 0;
       for (uint32_t groupIdx = 0; groupIdx < params.problemCount; ++groupIdx) {
           GemmCoord problemShape = problemShapeList[groupIdx];
           LayoutA layoutA = layoutAList[groupIdx];
           LayoutB layoutB = layoutBList[groupIdx];
           LayoutC layoutC = layoutCList[groupIdx];

           matmulBlockScheduler.Update(problemShape, MakeCoord(L1TileShape::M, L1TileShape::N));
           uint32_t coreLoops = matmulBlockScheduler.GetCoreLoops();

           // Determine the starting loopIdx of the current core under the current groupIdx
           uint32_t startLoopIdx;
           if (coreIdx < startCoreIdx) {
               startLoopIdx = coreIdx + coreNum - startCoreIdx;
           } else {
               startLoopIdx = coreIdx - startCoreIdx;
           }
           // Loop through the matmul of each groupIdx
           for (uint32_t loopIdx = startLoopIdx; loopIdx < coreLoops; loopIdx += coreNum) {
               // Compute block location
               GemmCoord blockCoord = matmulBlockScheduler.GetBlockCoord(loopIdx);
               GemmCoord actualBlockShape = matmulBlockScheduler.GetActualBlockShape(blockCoord);

               // Compute initial location in logical coordinates
               MatrixCoord offsetA{blockCoord.m() * L1TileShape::M, blockCoord.k() * L1TileShape::K};
               MatrixCoord offsetB{blockCoord.k() * L1TileShape::K, blockCoord.n() * L1TileShape::N};
               MatrixCoord offsetC{blockCoord.m() * L1TileShape::M, blockCoord.n() * L1TileShape::N};
               int64_t gmOffsetA = layoutA.GetOffset(offsetA);
               int64_t gmOffsetB = layoutB.GetOffset(offsetB);
               int64_t gmOffsetC = layoutC.GetOffset(offsetC);

               // Compute block-scoped matrix multiply-add
               blockMmad(
                   gmA[inGroupOffsetA + gmOffsetA], layoutA,
                   gmB[inGroupOffsetB + gmOffsetB], layoutB,
                   gmC[inGroupOffsetC + gmOffsetC], layoutC,
                   actualBlockShape);
           }

           inGroupOffsetA += problemShape.m() * problemShape.k();
           inGroupOffsetB += problemShape.k() * problemShape.n();
           inGroupOffsetC += problemShape.m() * problemShape.n();

           startCoreIdx = (startCoreIdx + coreLoops) % coreNum;
       }

       if constexpr (BlockMmad::DispatchPolicy::ASYNC) {
           blockMmad.SynchronizeBlock();
       }
   }

  3. 主机侧实例化对应的组件，并作为模板参数传入组件，进行内核调用

    using L1TileShape = GemmShape<128, 256, 256>;
   using L0TileShape = GemmShape<128, 256, 64>;

   using AType = Gemm::GemmType<half, LayoutA>;
   using BType = Gemm::GemmType<half, LayoutB>;
   using CType = Gemm::GemmType<half, LayoutC>;

   using BlockMmad = Gemm::Block::BlockMmad<DispatchPolicy, L1TileShape, L0TileShape, AType, BType, CType>;
   using BlockEpilogue = void;
   using BlockScheduler = typename Gemm::Block::GemmIdentityBlockSwizzle<3, 1>;

   // kernel level
   using MatmulKernel = Gemm::Kernel::GroupedMatmul<BlockMmad, BlockEpilogue, BlockScheduler>;

   typename MatmulKernel::Params params{
       problemCount, ptrProblemShape, ptrA, ptrLayoutA, ptrB, ptrLayoutB, ptrC, ptrLayoutC
   };

   // call a kernel
   MatmulKernel matmul;
   matmul(params);

8.4 开发者实践建议

案例选择

算子类型 推荐场景 注意事项
GEMM 多任务并行、L2Cache优化 需关注TileScheduler与SplitK的配合
QuantGEMM 模型压缩、低精度推理 量化参数需在FP_BUFFER中对齐
融合算子 非连续计算、高吞吐需求 优先使用TQueBind优化Vector计算

调试技巧

  • 精度验证：在Host侧使用aclrtMemcpy与numpy对比结果，误差阈值设为0.1%。
  • 性能分析：通过msprof op工具采集aic_mte2_time和aic_scalar_ratio，定位流水线瓶颈。
  • bank冲突规避：在TilingData中通过TILING_DATA_FIELD_DEF控制字段对齐，避免Vector读写冲突。

8.5 拓展案例：GEMV与多核同步

场景背景

GEMV（General Matrix-Vector Multiplication）通常需要结合AtomicAdd实现多核结果累加。例如，当输出向量y由多个AICore并行计算时，需确保写入同一位置的线程性。

实现步骤

步骤 操作内容 关键代码片段
1. 启用AtomicAdd 在Mmad中设置enAtomic=1 
MmadParams mmadParams;
mmadParams.enAtomic = 1;

| 2. 同步控制 | 使用SetFlag和WaitFlag管理多核依赖 |

SetFlag<HardEvent::MTE2>(eventId); 
WaitFlag<HardEvent::Mte2>(eventId);

| 3. 向量扩展 | 通过AscendC::Add实现多核结果累加 |

AscendC::Add(dstLocal, srcLocal, biasLocal, cSize);

性能优化点

  • Repeat参数：在GEMV中使用repeat减少Vector指令发射次数。
  • Bias暂存：将bias数据搬运至BT Buffer，避免重复从GM读取。

8.6 小结

本章通过三个典型算子开发案例，展示了昇腾模板库的分层设计优势：

  • GEMM：通过GroupedMatmul快速实现多任务并行。
  • QuantGEMM：FP_BUFFER和Cast接口降低量化成本。
  • 融合算子：Cube与Vector流水线合并，减少空闲时间。

开发者可结合msprof op工具的op_summary_*.csv和仿真trace.json，持续迭代优化分块策略与bank排布。下一章将深入性能调优工具的使用方法，帮助您量化改进效果。

第9章 开发者常见问题与解决方案

这章总结了算子开发中的最易犯的错误及解决方法。比如，核函数参数被修改的陷阱：核函数的输入输出地址由SetGlobalBuffer绑定，直接修改参数指针会导致数据错乱。教程中的反例展示了如何错误地重新赋值指针，而正例则通过绑定地址实现安全访问。

另一个常见问题是TillingData结构体的冗余和排布。教程建议用最小的数据类型（如uint8_t）描述字段，并确保结构体8字节对齐。比如，冗余的字段会增加内存拷贝开销，而对齐则能避免bank冲突。

教程还分析了Cube与Vector资源冲突的场景。比如，当两个向量同时访问同一个bank group时，会导致流水线等待。解决方案是通过增加地址偏移（如32字节）让它们落在不同bank。这种优化需要结合npu-smi查询硬件bank结构，并用msprof op分析冲突情况。

9.1 核函数参数修改陷阱

问题描述
在昇腾算子开发中，核函数的参数通常是全局内存指针（__gm__类型），直接修改这些指针会导致数据访问错误或性能下降。例如，反例中错误地将参数指针重新赋值，导致后续计算无法正确获取数据。

反例代码

// 核函数参数重新赋值导致错误
__aicore__ __global__ void FlashAttentionKernel(__gm__ uint8_t* query, __gm__ uint8_t* key, ..., __gm__ uint8_t* tilingData) {
    query = tmpQueryPtr;  // 错误：修改指针本身
    key = tmpKeyPtr;       // 错误：修改指针本身
    tilingData = tmpTilingDataPtr;  // 错误：修改指针本身
}

正例代码

// 正确使用SetGlobalBuffer设置数据地址
inputQueryGMTensor.SetGlobalBuffer(query);  // 仅读取query指向的数据
outputAttentionGMTensor.SetGlobalBuffer(attention);  // 仅写入attention指向数据

解决方案

  • 禁止修改参数指针：核函数的参数指针应视为只读，开发者不应尝试修改其指向。
  • 正确设置数据地址：使用SetGlobalBuffer方法绑定全局内存地址，确保数据读写正确。
  • 调试建议：若参数设置后出现数据错误，优先检查是否遗漏了SetGlobalBuffer调用。

9.2 TilingData冗余与排布

问题描述
TilingData结构体若设计不当，会导致内存浪费和性能下降。例如，字段类型选择过大或未对齐8字节，会增加内存拷贝时间，并可能引发bank冲突。

反例代码

// TilingData结构冗余且未对齐
BEGIN_TILING_DATA_DEF(TilingDataUnalign)
TILING_DATA_FIELD_DEF(uint64_t, blockDim);    // 实际只需uint8_t
TILING_DATA_FIELD_DEF(uint64_t, formerNum);   // 实际只需uint8_t
TILING_DATA_FIELD_DEF(uint64_t, tailNum);      // 实际只需uint8_t
TILING_DATA_FIELD_DEF(uint64_t, formerLength);
TILING_DATA_field_DEF(uint64_t, tailLength);
TILING_DATA_field_DEF(uint64_t, alignNum);
END_TILING_DATA_DEF;

正例代码

// TilingData结构精简且对齐8字节
BEGIN_TILING_DATA_DEF(TilingDataUnalign)
TILING_DATA_FIELD_DEF(uint8_t, formerNum);     // 仅需8位
TILING_DATA_FIELD_DEF(uint8_t, tailNum);         // 仅需8位
TILING_DATA_FIELD_DEF(uint32_t, formerLength);      // 仅需32位
TILING_DATA_FIELD_DEF(uint32_t, tailLength);      // 仅需32位
TILING_DATA_FIELD_DEF(uint32_t, alignNum);       // 仅需32位
END_TILING_DATA_DEF;

解决方案

问题 措荐做法 原因
冗余字段 使用uint8_t、uint32_t等最小类型 减少内存占用
8字节对齐 按字段排布顺序调整结构体 避免编译器自动补齐增加额外内存
冗余搬运 通过TilingDataUnalign减少GM访问 降低GetTilingData拷贝耗时

调试建议

  • 使用msprof op工具分析aic_mte2_time，若搬运耗时显著高于理论值，检查TilingData排布。
  • 在Host侧通过aclrtMemcpy验证数据是否正确拷贝至Device内存。

9.3 Cube与Vector资源冲突

问题描述
Cube计算（如Mmad）与Vector计算（如Add）若共享同一bank group，会导致资源冲突，降低性能。例如，未合理划分bank group的代码可能引发读写冲突，导致流水线等待。

反例代码

// 未合理划分bank group的Vector计算
LocalTensor<float> src0Local = inQueueSrc0.AllocTensor<float>();
LocalTensor<float> biasLocal = inQueueBias.DeQue<float>();
LocalTensor<float> dstLocal = outQueueDst.AllocTensor<float>();

Add(dstLocal, src0Local, biasLocal, 32, m, addRepeatParams);  // 可能导致bank冲突

正例代码

// 通过地址偏移规避bank冲突
LocalTensor<float> src0Local = inQueueSrc0.AllocTensor<float>();
LocalTensor<float> biasLocal = inQueueBias.DeQue<float>();
LocalTensor<float> dstLocal = outQueueDst.AllocTensor<float>();

// biasLocal与src0Local错开32字节，避免bank group冲突
AscendC::Add(dstLocal, src0Local, biasLocal[(8 * 1024 + 32) / sizeof(float)], mask, (8 * 1024) / 256, params);

解决方案

场景 反例 正例 效果
双输入向量 两向量起始地址相同bank group 增加32字节偏移 降低bank冲突率
Cube与Vector同步 未使用PipeBarrier 插入PipeBarrier&lt;PIPE_ALL&gt;() 确保Cube先完成
内存共享 重复申请UB缓冲区 通过TBuf共享Buffer 减少搬运次数

性能优化建议

  1. bank冲突分析：使用npu-smi查询硬件bank group布局，通过msprof op的bank_group_id字段定位冲突。
  2. bank group排布：在TQue中通过QuePosition::A1、QuePosition::B2等逻辑位置控制bank group分配。
  3. 双输入对齐：确保两个输入向量的起始地址差至少为32字节（AscendC::GetBlockIdx()生成偏移）。

9.4 其他常见问题

9.4.1 L2Cache命中不足

问题：未启用L2Cache搬运导致GM访问延迟。
解决方案：在DataCopy中使用512B对齐，并通过SetFlag同步。
代码示例：

// L2Cache搬运优化
DataCopyParams params;
params.srcStride = 512;  // 512B对齐
params.dstStride = 512;
DataCopy(tensorIn, tensorGM, params);

9.4.2 Tiling数据切分不合理

问题：未根据矩阵规模调整分块大小，导致多核利用率低。
解决方案：动态计算BlockTileM、BlockTileN，确保每个分块大小适配L1/L0 Buffer。
代码示例：

// 动态分块策略
constexpr uint32_t K_TILE = 256;  // K方向分块
uint32_t coreNum = GetCoreNum();   // 获取可用核数
uint32_t blockM = MatmulM / coreNum;  // 按核数切分M方向
uint32_t blockN = MatmulN / coreNum;   // 按核数切分N方向

9.5 调试工具推荐

工具 用途 使用场景
msprof op 性能分析 采集aic_mte2_time、aic_scalar_ratio等指标
npu-smi 硬件信息查询 查看bank group数和L2Cache大小
aclrtMemcpy 数据校验 在Host侧比较GM与CPU内存结果

调试流程

  1. 初步验证：使用aclrtSynchronizeStream确保核函数执行完毕，再通过aclrtMemcpy获取结果。

  2. 性能瓶颈定位：

    msprof op --output="./out" --ai-core=on --aic-metrics="PipeUtilization" add_custom_npu

  3. 仿真流水图分析：

    • 生成trace.json文件，通过Chrome浏览器的chrome://tracing页面分析流水线效率。
    • 重点检查MTE2与Vector的利用率差异。

9.6 文档参考

文档名称 关键章节
《Ascend C API参考》 Mmad、DataCopy的参数说明
《昇腾AI处理器硬件架构指南》 bank group结构与cacheline对齐规则
《性能调优工具指南》 msprof op工具的详细使用

学习建议

  • 初学者：优先掌握TilingData排布和PipeBarrier同步机制。
  • 进阶者：深入TQueBind和TBuf的内存管理策略，结合AtomicAdd减少搬运。

总结
昇腾算子开发中，参数修改、内存排布和资源冲突是高频问题。通过遵循模板库规范，合理使用bank group、bank冲突规避和异步接口，可显著提升算子性能。若遇到异常，可结合msprof op和npu-smi工具快速定位瓶颈，参考上述反例与正例调整实现。

第10章 编译与部署流程

算子写好之后，如何编译到昇腾设备上？这一章主要拆解整个流程，走完算子从代码到硬件运行的“最后一公里”。

首先，您需要用CMakeLists配置编译选项，指定昇腾的硬件版本（如Ascend910）和编译模式（Shared库或Python扩展）。其次，通过自定义算子包的安装路径（ASCEND_CUSTOM_OPP_PATH）将算子部署到昇腾的算子库目录。

编译时需注意：soc_version必须与目标设备匹配。比如，训练用的Ascend910和推理用的Ascend310需要不同的配置。此外，启用K_MAX_SHAPE_DIM宏能减少内存结构体的占用，避免栈溢出。

部署完成后，用msprof op检查算子是否被正确加载。如果发现算子未注册，可能是环境变量未设置正确，或安装路径未包含在ASCEND_CUSTOM_OPP_PATH。

10.1 算子编译工具链

昇腾算子的编译流程依赖CMake与Ascend C工具链，核心步骤如下：

阶段 工具/参数 说明
1. 环境准备 source /usr/local/Ascend/ascend-toolkit/set_env.sh 同步环境变量，确保工具链路径正确。
2. 编译入口 ./build.sh 脚本调用CMakeLists.txt生成Makefile并执行编译。
3. CMake配置 CMakeLists.txt 需指定aicore与soc_version（如Ascend910B），确保编译适配目标硬件。
4. 编译模式 -DCMAKE_BUILD_TYPE=Shared 生成共享库（.so）；若需构建Python扩展，使用-DCMAKE_BUILD_TYPE=Python。
5. 编译输出 build_out/ 编译结果（如libadd_custom.so）存于该目录，后续部署时使用。

典型编译命令示例

# 生成共享库
./build.shmul.sh -t Shared -v Ascend910B -a x86_64
# 生成Python扩展
./buildmatmul.sh -t Python -v Ascend310P -a aarch64

CMakeLists关键配置

set(ASCEND_TOOLKIT_PATH "/usr/local/Ascend/ascend-toolkit/latest")
set(SOC_VERSION "Ascend910B")
set(ARCH_TYPE "x86_64")
include(${ASCEND_TOOLKIT_PATH}/cmake/ascend_c.cmake)

10.2 算子安装与环境变量

编译完成后，需将算子包部署至昇腾AI处理器的算子库路径。以下为关键步骤：

步骤1：安装自定义算子包

通过编译生成的.run脚本执行安装：

./custom_opp_run --install-path=/usr/local/Ascend/opp/custom
  • 关键参数说明：
    • --install-path：指定算子包部署路径，若不指定则默认为ASCEND_OPP_PATH。

步骤2：配置环境变量

在Host侧运行时，需设置以下环境变量以定位自定义算子：

环境变量 说明
ASCEND_CUSTOM_OPP_PATH 指定自定义算子库路径，需包含lib目录（如/usr/local/Ascend/opp/custom）。
ASCEND_INSTALL_PATH CANN安装路径（如/usr/local/Ascend/ascend-toolkit/latest）。 
export ASCEND_CUSTOM_OPP_PATH=/usr/local/Ascend/opp/custom
export ASCEND_INSTALL_PATH=/usr/local/Ascend/ascend-toolkit/latest

步骤3：验证部署

安装完成后，可通过msprof op工具检查算子是否成功注册：

msprof op --query-op add_custom

若输出包含算子信息，则表示部署成功。

10.3 编译与部署注意事项

1. soc_version适配

不同硬件平台需指定对应的soc_version，例如：

  • 训练场景：Ascend910B、Ascend920
  • 推理场景：Ascend310P、Atlas 800I A2

2. 编译模式选择

模式 适用场景 典型配置
Shared 高性能算子，与CANN库绑定 -t Shared
Python Python扩展，灵活调试 -t Python

3. 编译依赖管理

若算子依赖其他模块（如TilingData），需在CMakeLists.txt中显式声明依赖：

find_package(ascend_c REQUIRED)
ascend_c_add_operator(add_custom 
    SOURCES add_custom.cpp 
    TILING_SOURCE add_custom_tiling.h
    DEPENDS ascend_c::matmul)

4. 环境变量冲突规避

多算子部署时，需确保ASCEND_CUSTOM_OPP_PATH优先级高于系统默认路径：

# 修改环境变量顺序
export ASCEND_CUSTOM_OPP_PATH=/usr/local/Ascend/opp/custom:/usr/local/Ascend/opp/built-in

5. 调试与运行模式切换

编译时需区分CPU仿真与NPU实际运行：

  • CPU调试：无需硬件，仅需配置npu-smi仿真环境。
  • NPU运行：需实际设备，且soc_version需与硬件型号一致。

6. 性能优化编译选项

通过CMakeLists.txt传递编译宏，优化性能：

add_definitions(-DK_MAX_SHAPE_DIM=0)

此参数可减少ShapeInfo结构体的栈空间占用。

10.4 典型部署结构示例

部署后算子库目录结构如下：

/usr/local/Ascend/opp/custom/
├── framework
│   └── ai_core
│       └── lib
│           └── libadd_custom.so
├── op_proto
│   └── inc
│       └── add_custom.proto.h
└── config
    └── aicore_add_custom.json

运行时依赖检查

  • 共享库路径：/usr/local/Ascend/ascend-toolkit/latest/lib64
  • 算子配置文件：aicore_add_custom.json需包含op_type与op_version字段。

10.5 基于CATLASS模板库进行算子编译

CATLASS提供了对应的编译脚本，可以快速编译基于CATLASS模板库开发的算子，对应的scripts/build.sh编译脚本如下：

SCRIPT_PATH=$(dirname $(realpath $0))
CMAKE_SOURCE_PATH=$(realpath $SCRIPT_PATH/..)

CMAKE_BUILD_PATH=$CMAKE_SOURCE_PATH/build

OUTPUT_PATH=$CMAKE_SOURCE_PATH/output

if [[ $# -eq 0 ]]; then
    echo "Usage: bash build.sh [--clean] [target]"
    exit 0
fi

TARGET=${!#}
echo "Target is: $TARGET"
CMAKE_BUILD_TYPE=Release

mkdir -p $CMAKE_BUILD_PATH

while [[ $# -gt 0 ]]; do
    case "$1" in
        --clean)
            rm -rf build
            rm -rf output
            ;;
        --debug)
            echo "Hint: only python extension support debug mode."
            CMAKE_BUILD_TYPE=Debug
            ;;
        --*)
            echo "Unknown option: $1"
            ;;
    esac
    shift
done

function build_shared_lib() {
    cd $CMAKE_SOURCE_PATH/examples/shared_lib
    rm -rf build
    cmake --no-warn-unused-cli -B build -DCMAKE_BUILD_TYPE=$CMAKE_BUILD_TYPE -DCMAKE_INSTALL_PREFIX=$OUTPUT_PATH/shared_lib -DACT_INCLUDE_DIR=$CMAKE_SOURCE_PATH/include
    cmake --build build -j
    cmake --install build
    cd $CMAKE_SOURCE_PATH
}

function build_torch_library() {
    cd $CMAKE_SOURCE_PATH/examples/python_extension
    rm -rf build
    cmake --no-warn-unused-cli -B build -DCMAKE_BUILD_TYPE=$CMAKE_BUILD_TYPE -DCMAKE_INSTALL_PREFIX=$OUTPUT_PATH/python_extension -DACT_INCLUDE_DIR=$CMAKE_SOURCE_PATH/include -DPython3_EXECUTABLE=$(which python3) -DBUILD_TORCH_LIB=True
    cmake --build build -j
    cmake --install build
    cd $CMAKE_SOURCE_PATH
}

function build_python_extension() {
    cd $CMAKE_SOURCE_PATH/examples/python_extension
    rm -rf build
    python3 setup.py bdist_wheel --dist-dir $OUTPUT_PATH/python_extension
    cd $CMAKE_SOURCE_PATH
}

if [[ "$TARGET" == "shared_lib" ]]; then
    build_shared_lib
elif [[  "$TARGET" == "lib_cmake" ]]; then
    cmake -DENABLE_LIB=ON -S $CMAKE_SOURCE_PATH -B $CMAKE_BUILD_PATH
    cmake --build $CMAKE_BUILD_PATH
elif [[ "$TARGET" == "python_extension" ]]; then
    build_python_extension
elif [[ "$TARGET" == "torch_library" ]]; then
    build_torch_library
else
    cmake --no-warn-unused-cli -S$CMAKE_SOURCE_PATH -B$CMAKE_BUILD_PATH
    cmake --build $CMAKE_BUILD_PATH --target $TARGET -j
fi

模板库支持三中方式编译，单算子、共享库和torch_library，接下来重点介绍单算子编译方式，其他两种可参考开源网站。

10.5.1 添加自定义算子的CMakelists

假设主机侧调用都在examples/目录下，首先在examples下的CMakeLists下添加对应算子子目录，假设当前在examples目录下实现了00_basic_matmul算子，下是一个示例：

目录结构为

-examples
|----00_basic_matmul
   |---basic_matmul.cpp
   |---CMakelists.txt
|----CMakelists.txt

examples/CMakelists.txt内容示例为：

set(ACT_EXAMPLES_COMMON_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/common)
set(BISHENG_COMPILER_OPTIONS
    -O2 -std=c++17 -xcce
    -mllvm -cce-aicore-stack-size=0x8000
    -mllvm -cce-aicore-function-stack-size=0x8000
    -mllvm -cce-aicore-record-overflow=true
    -mllvm -cce-aicore-addr-transform
    -mllvm -cce-aicore-dcci-insert-for-scalar=false
    -DL2_CACHE_HINT
    -I${ASCEND_HOME_PATH}/compiler/tikcpp
    -I${ASCEND_HOME_PATH}/compiler/tikcpp/tikcfw
    -I${ASCEND_HOME_PATH}/compiler/tikcpp/tikcfw/impl
    -I${ASCEND_HOME_PATH}/compiler/tikcpp/tikcfw/interface
    -I${ASCEND_HOME_PATH}/include
    -I${ASCEND_HOME_PATH}/include/experiment/runtime
    -I${ASCEND_HOME_PATH}/include/experiment/msprof
    -I${ACT_EXAMPLES_COMMON_SOURCE_DIR}
    -I${CMAKE_SOURCE_DIR}/include
    -L${ASCEND_HOME_PATH}/lib64
    -Wno-macro-redefined -Wno-ignored-attributes
    -lruntime -lstdc++ -lascendcl -lm -ltiling_api -lplatform -lc_sec -ldl -lnnopbase
)

if(DEFINED PROF)
    list(APPEND ${BISHENG_COMPILER_OPTIONS} -lprofapi)
endif()

file(GLOB_RECURSE ACT_INCLUDE_FILES ${CMAKE_SOURCE_DIR}/include/*.hpp)
file(GLOB_RECURSE ACT_EXAMPLES_COMMON_INCLUDE_FILES ${ACT_EXAMPLES_COMMON_SOURCE_DIR}/*.hpp)
add_custom_target(act_examples)

function(act_example_add_executable NAME)
    if (${NAME} STREQUAL "17_gemv_aiv")
        set(ARCH "dav-c220-vec")
    else()
        set(ARCH "dav-c220")
    endif()

    add_custom_command(
        OUTPUT ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}/${NAME}
        COMMAND ${CMAKE_BISHENG_COMPILER} --cce-aicore-arch=${ARCH} ${BISHENG_COMPILER_OPTIONS} ${ARGN} -o ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}/${NAME}
        DEPENDS ${ARGN} ${ACT_INCLUDE_FILES} ${ACT_EXAMPLES_COMMON_INCLUDE_FILES}
        WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
        COMMENT "Compiling executable kernel: ${NAME}"
    )

    add_custom_target(${NAME} ALL DEPENDS ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}/${NAME})
    add_dependencies(act_examples ${NAME})
endfunction()

if(ENABLE_LIB)
    add_subdirectory(lib_cmake)
else()

foreach(EXAMPLE
    // 算子子目录
    00_basic_matmul
)
    add_subdirectory(${EXAMPLE})
endforeach()

endif()

该文件定义了对应的编译命令和参与编译的子目录，同时需要在自定义算子目录下的CMakelists.txt添加可执行文件

act_example_add_executable(
    00_basic_matmul
    basic_matmul.cpp
)

10.5.2 进行单算子编译

进行单算子的编译，获得可执行文件

bash scripts/build.sh 00_basic_matmul

10.5.3 执行对应的算子

./00_basic_matmul [参数列表]

10.6 常见问题与解决方案

问题1：算子未加载

  • 现象：msprof op未识别自定义算子。
  • 解决方案：检查ASCEND_CUSTOM_OPP_PATH是否包含部署路径。

问题2：编译报错TilingData未定义

  • 原因：TilingData结构体未在Host侧正确定义或未注册。
  • 修复：确保TilingData已通过REGISTER_TILING_DATA宏注册。

问题3：多核调用失败

  • 现象：blockDim设置超过硬件核数。
  • 调试：检查GetCoreNum()输出，确保blockDim ≤ GetCoreNum()。

10.7 总结

昇腾算子的编译与部署需遵循以下核心原则：

  1. 工具链一致性：soc_version、arch与目标硬件匹配。
  2. 环境变量优先级：自定义算子路径需优先于内置算子路径。
  3. 编译优化：合理设置K_MAX_SHAPE_DIM、STAGES等宏，减少冗余资源占用。
  4. 多平台适配：区分训练/推理场景，选择GroupedMatmul或QuantGEMM模板库。

掌握上述流程后，开发者可快速将自定义算子集成到昇腾环境中，并通过msprof op等工具进行性能验证与调优。

第11章 实际开发工具与文档参考

这章汇总了算子开发所需的的工具和文档。首先是性能分析工具：msprof op能采集硬件运行时的详细指标（如搬运耗时、计算单元利用率），并生成trace.json文件用于Chrome浏览器查看流水线。其次是是ST测试工具（msopst），它根据定义的用例文件自动生成测试代码，并对比numpy结果验证精度。比如，用例文件描述了输入输出的shape和类型，测试工具会自动生成数据并运行算子。

文档方面，教程推荐了三类关键资料：Ascend C API手册（解释基础指令参数）、硬件架构指南（说明bank冲突规避规则）、性能调优工具手册（指导如何解析msprof数据）。此外，教程还提供了编译器选项（如soc_version）和环境变量配置（如ddk_path）的详细说明，帮助开发者避免因配置不当导致的编译错误。

工具链的使用需要循序渐进：初学者应从msprof op的流水图分析开始，逐步掌握bank冲突分析和性能瓶颈定位。比如，先用msprof op采集数据，再在Chrome浏览器中放大关键计算步骤，观察Cube和Vector的并行性。

11.1 推荐工具列表

工具名称 工具用途 使用场景示例
npu-smi 查询硬件型号与状态 npu-smi info 获取AI处理器型号（如Atlas A2）。
msprof op 性能分析与流水图仿真 msprof op --output=./out --ai-core=on --aic-metrics=PipeUtilization add_custom_npu 采集性能数据。
msopst ST测试用例生成 msopst run -i AddCustom_case.json 生成并运行测试用例。
aclrtMalloc 设备内存分配 aclrtMalloc((void**)&amp;d_C, C_size, ACL_MEM_MALLOC_HUGE_FIRST);
aclrtMemcpy 主机与设备侧数据拷贝 aclrtMemcpy(h_C, C_size, d_C, C_size, ACL_MEMCPY_DEVICE_TO_HOST);

11.1.1 性能分析工具链

msprof op 工作流

  1. 采集性能数据

    msprof op --output=./out --ai-core=on --aic-metrics=PipeUtilization add_custom_npu

  2. 查看流水图

    • 使用 Chrome 浏览器打开 trace.json，通过 w（放大）、s（缩小）查看指令级流水线。
性能数据文件解析
  • op_summary_*.csv
    包含算子执行的详细指标，如 aic_mte2_time（MTE2搬运耗时）、aic_mte2_ratio（MTE2带宽利用率）。
  • trace.json
    仿真流水图，展示各计算单元（Cube/Vector）的指令执行顺序与资源占用。

11.2 关键文档与资源

文档名称 核心内容 获取方式
《Ascend C API 参考》 基础指令参数说明（如 DataCopy 的 srcStride 配置）。 ${install_path}/docs/ascendc_api_reference.pdf
《昇腾AI处理器硬件架构指南》 Bank结构、cacheline对齐规则（如UB分配时需避免bank冲突）。 ${install_path}/docs/ascendc_hardware_glossary.pdf
《性能调优工具指南》 msprof op 工具使用详解（如如何解析 aic_mte2_ratio）。 ${install_path}/docs/performance_tuning_guide.pdf
《算子开发工具指南》 TilingData结构定义规范、资源冲突占比文件说明。 ${install_path}/docs/operator_development_tools.pdf

11.2.1 文档使用示例

msprof op 性能分析示例
// 帮助命令
msprof op --help // 获取该工具的说明书

以basic_matmul算子为例，可以执行

msprof op --application="./20_lab1_pingpong 1024 2048 512 1" --output="../../examples/20_lab1/"

其中重要参数包括：

  • --application，表示可执行程序的执行命令
  • --output，性能文件的存储目录

在输出目录下会存在类似OPPROF_***的目录下，存在各个性能指标，常见目录结构如下：

OPPROF_****
|--dump
|--ArithmeticUtilization.csv
|--L2Cache.csv
|--Memory.csv
|--OpBasicInfo.csv
|--PipeUtilization.csv
....

可以看到包括L2Cache，Memory等性能指标。

11.3 开发环境配置要点

11.3.1 CANN软件安装

  1. 安装依赖

    # Ubuntu环境示例
apt-get install -y gcc make cmake python3 python3-pip

  2. 安装CANN套件

    chmod +x Ascend-cann-toolkit_XXX_linux-x86_64.run
./Ascend-cann-toolkit_XXX_linux-x86_64.run --install
source /usr/local/Ascend/ascend-toolkit/set_env.sh

  3. 验证安装

    npu-smi info  # 查询芯片型号（如Atlas A2）

11.3.2 编译器选项

编译选项 用途说明 示例配置（CMakeLists.txt）
aicore 启用AI Core计算优化 set(CMAKE_CXX_FLAGS &quot;-DAICORE&quot;)
soc_version 指定目标AI处理器型号（如Ascend910） add_definitions(-DSOC_VERSION=AtlasA2)
K_MAX_SHAPE_DIM 限制ShapeInfo维度大小以减少栈空间占用 add_definitions(-DK_MAX_SHAPE_DIM=0)

11.4 常见问题排查工具

11.4.1 内存对齐验证

// 检查GM地址是否512B对齐
if (address % 512 != 0) {
    printf("Address not 512B aligned!\n");
    return ACL_ERROR;
}

11.4.2 Bank冲突分析

  • 工具命令

    msprof op --output=./out --bank-conflict=on add_custom_npu

  • 冲突类型

    冲突类型 描述 解决方案示例
    读写冲突 同一bank同时读写 修改 dstBlkStride 为8B步长。
    写写冲突 同一bank_group内并发写入 使用 TQueBind 减少UB数据搬运。
    读读冲突 同一bank_group内并发读取（双src场景） 为 src0 和 src1 分配不同bank_group。

11.5 开发流程模板化工具

11.5.1 msopst 测试用例生成

  1. 定义测试用例

    [
 {
   "case_name": "Test_AddCustom_001",
   "op": "AddCustom",
   "input_desc": [
     {
       "name": "x",
       "shape": [8, 2048],
       "data_type": "float16",
       "data_distribute": "uniform"
     }
   ],
   "output_desc": [
     {
       "name": "z",
       "shape": [8, 2048],
       "data_type": "float16"
     }
   ]
 }
]

  2. 执行测试

    ./msopst run -i ./AddCustom_st/AddCustom_case.json -soc Ascend910 -out ./reports

11.6 性能调优关键指标

11.6.1 核心性能字段解析

字段名 含义 优化方向
aic_mte2_time MTE2搬运耗时（单位us） 优化 srcStride 和 blockLen 提高带宽利用率。
aic_mac_ratio Cube单元利用率（0~100%） 增加 STAGES=2 使用pingpong buffer减少依赖。
aic_scalar_ratio Scalar指令执行占比 减少 if/else 逻辑，使用 IterateAll&lt;false&gt; 异步执行。

11.7 环境变量管理

11.7.1 安装路径配置

export DDK_PATH=/usr/local/Ascend/cann-toolkit
export NPU_HOST_LIB=$DDK_PATH/x86_64-linux/devlib

11.7.2 自定义算子路径注册

# 安装自定义算子包
./custom_opp_run --install-path=$DDK_PATH/opp/vendors/my_vendor

# 注册到环境变量（需在运行前执行）
source $DDK_PATH/opp/vendors/my_vendor/bin/set_env.sh

11.8 学习资源推荐

11.8.1 入门级学习路径

  1. 阅读《Ascend C算子开发指南》
    • 掌握 __global__ __aicore__ 函数定义与调用规则。
  2. 实践Add算子模板
    • 从 CopyIn、Compute、CopyOut 三级流水线开始。
  3. 使用msprof op分析瓶颈
    • 通过 aic_mte2_ratio 识别搬运带宽未达标场景。

11.8.2 高级调优技巧

  • FP16转bfloat16精度测试

    // 量化参数搬运优化
LocalTensor<uint64_t> deqLocal = inQueueDeq.AllocTensor<uint64_t>();
DataCopy(deqLocal, deqGM, QuantSize);

  • 双核同步控制

    PipeBarrier<PIPE_ALL>();
SetFlag<HardEvent::MTE2>(eventId);
WaitFlag<HardEvent::Mte2>(eventId);

11.9 开发者工具速查表

工具名称 功能描述 常用命令/参数
npu-smi 硬件状态查询 npu-smi info 查看芯ID（如910B/Atlas A2）。
msprof op 性能分析 --ai-core=on 启用AI Core性能数据采集，--bank-conflict=on 检查bank冲突。
msopst ST测试用例管理 -i 指定测试用例文件，-soc 指定AI处理器型号。
GmAlloc CPU仿真调试内存分配 AscendC::GmAlloc(size) 用于 AddCustom_case.json 数据生成。
build.sh 模板库算子编译脚本 bash build.sh gemm_grouped 生成GroupedMatmul算子库。

11.10 附录：性能优化工具对比

工具名称 适用阶段 优势 限制
msprof op 性能分析 支持仿真流水图、实时数据采集 依赖完整CANN环境配置。
TilingScheduler Host侧分块策略 通过 problemCount 实现多矩阵并行 需手动调优 TilingData 字段对齐。
msopst 功能验证 自动对比numpy精度误差 仅支持FP16/FP32等部分数据类型。

通过以上工具与文档的协同使用，开发者可快速完成算子开发、验证与性能调优。工具链的熟练掌握是高效开发的关键，建议从 msprof op 流水图分析开始，逐步深入TilingData对齐优化与bank冲突规避。

第12章 开发者学习路径建议

算子开发需要循序渐进的学习路径。

第一步：掌握Ascend C的基础语法，比如如何定义核函数、如何绑定内存地址。

第二步：熟悉模板库的分层架构，学会如何用BlockMmad、TileCopy等组件搭建算子。

第三步：进入性能优化阶段，用msprof op分析瓶颈，调整分块策略和bank排布。

第四步：通过ST测试用例验证算子的正确性，并逐步迁移到真实昇腾设备。

第五步：深入复杂场景，比如开发GroupedMatmul或QuantMatmul，结合SplitK和Double Buffer提升效率。

最后，推荐多读模板库的源码，了解各组件的实现逻辑。比如，模板库中的的MatmulUniversalAdapter封装了多核调度逻辑，是优化性能的关键。

12.1 从零到模板库使用：循序渐进掌握核心能力

第一步：掌握Ascend C与C++的差异

昇腾算子开发的核心语言是Ascend C，它基于C++扩展了硬件特化语法。初学者需重点关注以下差异：

C++特性 Ascend C特性 注意事项
LocalTensor/GlobalTensor 基础数据类型（如half、float） LocalTensor用于描述AI Core内部存储，GlobalTensor用于描述Global Memory地址。使用LocalTensor时需注意其访问权限（如VECIN、L0A等）
LocalTensor需通过AllocTensor/FreeTensor管理 普通C++变量直接赋值 LocalTensor必须通过TQue/TBuf接口操作，不可直接赋值
__aicore__函数类型限定符 void普通函数 __aicore__声明的函数必须在NPU上执行，__global__声明的核函数需通过&lt;&lt;&lt;blockDim, stream&gt;&gt;&gt;调用
GetBlockIdx()函数 std::thread::get_id() GetBlockIdx()是Ascend C提供的核心函数，用于多核并行中的数据切分（如x + BLOCK_LENGTH * GetBlockIdx()）

示例代码

__aicore__ inline void CopyIn()
{
    LocalTensor<half> a1Local = inQueueA1.AllocTensor<half>();
    LocalTensor<half> b1Local = inQueueB1.AllocTensor<half>();
    LocalTensor<uint64_t> deq1Local = inQueueDeq1.AllocTensor<uint64_t>();
    DataCopy(a1Local, aGM, dataCopyA1Params);
    DataCopy(b1Local, bGM, dataCopyB1Params);
    DataCopy(deq1Local, deqGM, QuantSize);
    inQueueA1.EnQue(a1Local);
    inQueueB1.EnQue(b1Local);
    inQueueDeq1.EnQue(deq1Local);
}

第二步：理解SPMD模型与核函数调用

昇腾算子开发采用SPMD（Single Program, Multiple Data）模型，每个AI Core独立处理数据分片。关键步骤如下：

  1. 多核并行：通过GetBlockIdx()获取当前核的索引，计算数据偏移地址。
  2. 核函数调用：使用&<<<blockDim, stream>>>;启动核函数，blockDim需与实际AI Core数量匹配。
  3. 数据切分：通过TilingData结构定义分块逻辑，确保每个核的数据规模合理（如TilingDataUnalign）。

第三步：参考模板库快速开发算子

模板库通过分层设计（如BlockMmad、TileCopy）封装了算子核心逻辑。开发者可直接使用模板参数化配置，减少重复代码。例如，MatmulKernel的组装逻辑：

using DispatchPolicy = matmul::MmadAtlasA2Pingpong<true>;
using L1TileShape = MatmulShape<128, 256, 256>;
using L0TileShape = MatmulShape<128, 256, 64>;
using AType = matmul::MatmulType<ElementA, LayoutA>;
using BType = matmul::MatmulType<ElementB, LayoutB>;
using CType = matmul::MatmulType<ElementC, LayoutC>;

using BlockMmad = matmul::block::BlockMmad<DispatchPolicy, L1TileShape, L0TileShape, AType, BType, CType>;
using BlockEpilogue = void;
using TileScheduler = matmul::block::MatmulIdentityBlockSwizzle<>;
using MatmulKernel = matmul::kernel::BasicMatmul<BlockMmad, BlockEpilogue, TileScheduler>;
using MatmulHandle = CATLASS::matmul::device::MatmulUniversalAdapter<MatmulKernel>;

12.2 深入性能优化：从瓶颈分析到高效实践

性能调优核心流程

性能优化是一个采集-分析-优化-验证的闭环流程。以下是关键工具和步骤：

阶段 工具 操作示例
性能数据采集 msprof op msprof op --output=&quot;./out&quot; --ai-core=on --aic-metrics=&quot;PipeUtilization&quot; add_custom_npu
瓶颈分析 op_summary_*.csv 查看aic_mte2_time、aic_scalar_ratio等字段，定位搬运或计算瓶颈
流水图分析 Chrome Tracing 使用trace.json文件在Chrome浏览器中分析流水线并行性（快捷键：w放大，s缩小）

实战建议

  • 搬运优化：使用DataCopy接口代替for循环搬运数据，减少搬运次数。例如：

    DataCopyParams copyParams;
copyParams.blockCount = imgHeight;
copyParams.blockLen = copyWidth / 8;
copyParams.srcStride = (imgWidth - copyWidth) / 8;
copyParams.dstStride = 0;
DataCopy(tensorIn, tensorGM, copyParams);  // 一次搬运32KB

  • 内存管理：通过K_MAX_SHAPE_DIM宏缩减栈空间，例如将ShapeInfo的维度设置为0：

    #define K_MAX_SHAPE_DIM 0
struct ShapeInfo {
    uint32_t shape[K_MAX_SHAPE_DIM];
    uint32_t originalShape[K_MAX_SHAPE_DIM];
};

高阶技巧：bank冲突规避与Double Buffer

  • bank冲突分析：Unified Buffer的bank group冲突需通过地址对齐规避。例如，定义TilingData时需保证字段对齐：

    BEGIN_TILING_DATA_DEF(TilingDataUnalign)
TILING_DATA_FIELD_DEF(uint8_t, formerNum);  // 字节对齐优化
TILING_DATA_FIELD_DEF(uint8_t, tailNum);
TILING_DATA_FIELD_DEF(uint32_t, formerLength);
TILING_DATA_FIELD_DEF(uint32_t, tailLength);
END_TILING_DATA_DEF;

    反例：若字段类型为uint64_t而未对齐，可能导致额外10字节的浪费。
    正例：通过字段排布优化，减少冗余字节（如formerNum和tailNum为uint8_t）。

  • Double Buffer：在TQue初始化时设置BUFFER_NUM=2，实现搬运与计算的并行：

    TQue<QuePosition::VECIN, 2> inQueueSrc0;  // 使能Double Buffer

附录：学习路径推荐工具与文档

工具方面，教程推荐熟练使用npu-smi（查询硬件状态）、msprof op（性能分析）、msopst（测试用例生成）、gdb（仿真调试）等工具。每个工具都有特定场景下的的使用建议：比如npu-smi用于确认硬件型号，msprof op则用于采集性能数据。

文档方面，教程推荐三类核心资料：API手册（如DataCopy的参数说明）、硬件架构指南（bank冲突规避规则）、工具使用指南（msprof op的详细操作）。开发者可在编译和调试时随时查阅这些文档，比如在遇到bank冲突时参考架构指南的内存排布建议。

工具列表

工具 用途 使用建议
npu-smi 查询硬件型号与状态 npu-smi info获取AI处理器型号
msprof op 性能分析与Profiling 采集aic_mte2_time耗时等关键指标
msopst ST测试用例生成 生成AddCustom_case.json
gdb CPU仿真调试 用于Host侧代码调试

文档推荐

  • 《Ascend C API参考》：基础指令参数说明（如Mmad、DataCopy）。
  • 《昇腾AI处理器硬件架构指南》：bank结构、cacheline对齐规则。
  • 《性能调优工具指南》：msprof op工具的详细使用（如op_summary_*.csv解析）。

总结
开发者应遵循**“模板化开发→性能分析→优化→融合算子设计”**的路径，逐步掌握昇腾算子开发的核心能力。建议初期通过模板库快速实现算子，后期深入性能调优与bank冲突规避，并结合仿真工具持续迭代优化。

# 人工智能 # 昇腾
Logo
鲲鹏昇腾开发者社区是面向全社会开放的“联接全球计算开发者，聚合华为+生态”的社区，内容涵盖鲲鹏、昇腾资源，帮助开发者快速获取所需的知识、经验、软件、工具、算力，支撑开发者易学、好用、成功，成为核心开发者。

鲲鹏昇腾开发者社区是面向全社会开放的“联接全球计算开发者，聚合华为+生态”的社区，内容涵盖鲲鹏、昇腾资源，帮助开发者快速获取所需的知识、经验、软件、工具、算力，支撑开发者易学、好用、成功，成为核心开发者。

更多推荐

昇腾平台MindSpore模型训练优化心得体会

MindSpore作为昇腾AI生态的核心深度学习框架，凭借自动微分、动静结合、端边云全场景部署等特性，成为昇腾平台上模型开发的首选工具。在实际模型训练过程中，开发者常面临训练速度慢、显存占用高、资源利用率低等问题。本文结合MindSpore框架特性与昇腾硬件优势，从数据预处理、网络结构优化、训练策略调整、显存优化四个核心维度，分享模型训练的优化思路与实战方法，助力开发者在昇腾平台上高效完成模型训练

avatar 鲲鹏昇腾开发者社区

昇腾AI创新大赛-昇思模型开发挑战赛（S1赛季）-MultiModal赛道铜奖方案

本文档详细记录了针对 Qwen2-VL 和 janus_pro 模型的关键性能优化点，并附带了相应的核心代码实现。

avatar 鲲鹏昇腾开发者社区
cover

[嵌入式AI从0开始到入土]21_基于昇腾310P RC模式的Pi0模型部署实践

avatar 鲲鹏昇腾开发者社区