前言

在深度学习训练与推理的生产环境中，算子（Operator）是模型计算图的基本执行单元。每一层卷积、每一次矩阵乘法、每一个激活函数，本质上都是一个个算子在硬件上的具体实现。昇腾NPU作为华为面向AI场景自研的专用加速器，其软件栈核心之一便是CANN（Compute Architecture for Neural Networks）。CANN不仅负责底层硬件资源调度，更承载了完整的算子生态体系——而ops-nn，正是这个体系中提供神经网络高阶算子的开源仓库。

ops-nn的全称是"基于CANN技术的神经网络高阶算子库"，它以开源的方式向开发者公开了矩阵乘法类（matmul）、激活函数类（activation）、索引类（index）、损失函数类（loss）等核心算子的完整实现。开发者不仅可以直接使用这些预置算子，还能基于项目提供的工程模板和Ascend C编程框架，自主开发符合昇腾NPU架构规范的自定义算子，并在真实硬件上进行验证和部署。

对于希望在昇腾生态中深耕的工程师而言，了解ops-nn的架构设计与开发流程，是掌握NPU算子定制能力的关键一步。项目源码托管于AtomGit，遵循CANN软件版本配套机制，开发者需根据目标CANN版本选择对应的Git标签分支，以确保API兼容性和功能完整性。

ops-nn在昇腾软件栈中的位置

理解ops-nn，首先要理解CANN的整体架构。CANN是昇腾NPU的软件基础层，向上对接主流AI框架（如MindSpore、PyTorch、TensorFlow），向下直接管理AI Core、AI CPU等硬件计算单元。在CANN的算子体系内部，根据算子的实现方式和调用层次，可以划分为多个层级：

最底层是L0级接口，这是硬件直接识别的原生指令接口，性能最优但开发门槛最高；往上一层是aclnn接口（AI Compiler Neural Network），它对L0接口进行了封装，提供了面向神经网络算子的标准化C语言API，开发者无需直面硬件指令细节，通过aclnn即可完成算子的调用与编排；再往上是图模式（Graph Mode）接口，允许用户将多个算子组合为计算图进行优化和执行。

ops-nn中的算子同时提供了aclnn接口和图模式接口两套调用方式。以矩阵乘法类算子为例，quant_batch_matmul_v4、weight_quant_batch_matmul_v2等均支持通过aclnn接口直接调用，同时在op_graph目录下保留了图融合的相关实现。项目的设计哲学是：优先提供高性能的aclnn接口，同时兼顾图模式的可组合性。

从支持的硬件范围来看，ops-nn当前开源的算子支持Atlas A2、A3系列训练和推理产品，以及Ascend 950PR、Ascend 950DT等新一代芯片，并可通过CANN Simulator仿真工具在没有真实NPU硬件的环境中进行开发调试。这一特性对开发者的日常开发和CI/CD验证流程极为友好。

项目目录结构解析

ops-nn的目录结构体现了算子开发的标准工程化思路。理解这套结构，是掌握算子开发和贡献流程的基础。

项目根目录下，最核心的目录包括算子分类目录（如matmul、activation、loss、index等）、examples示例目录、docs文档目录以及scripts脚本目录。算子分类目录下，每个算子都有自己独立的子目录，子目录内遵循固定的组织规范。

单个算子目录的典型结构如下：

${op_name}/
├── examples/
│   ├── test_aclnn_${op_name}.cpp    # aclnn调用示例
│   └── test_geir_${op_name}.cpp      # 图模式调用示例
├── op_host/
│   ├── ${op_name}_def.cpp            # 算子信息库
│   ├── ${op_name}_infershape.cpp     # 输出形状推导
│   └── ${op_name}_tiling.cpp         # Tiling切分策略
├── op_kernel/
│   ├── ${op_name}.cpp                # Kernel入口
│   ├── ${op_name}.h                  # Kernel实现
│   ├── ${op_name}_tiling_key.h       # TilingKey定义
│   └── ${op_name}_tiling_data.h      # Tiling参数结构体
├── op_api/
│   ├── aclnn_${op_name}.cpp          # aclnn接口实现
│   └── ${op_name}.cpp                # L0接口实现
├── op_graph/
│   └── ${op_name}_proto.h            # 算子原型定义
└── tests/ut/                          # 单元测试

op_host目录存放Host侧（CPU端）代码，负责算子的元信息定义、形状推导和Tiling策略计算。op_kernel目录存放Device侧（AI Core端）代码，即真正的计算核函数，使用Ascend C语言编写。op_api目录存放aclnn接口的封装实现。op_graph目录则用于图融合场景下的算子原型定义。

值得注意的是，项目中并非每个算子都包含上述全部子目录。README中明确说明，若某个算子缺少op_kernel目录，可能是因为它复用了其他算子的Kernel实现；若缺少op_api目录，则表明该算子暂不支持aclnn调用。这种灵活的设计允许算子之间存在复用关系，避免重复开发。

环境准备与源码获取

ops-nn的运行依赖昇腾CANN软件环境。官方推荐两种部署方式：CANNLab云开发环境和Docker容器环境。CANNLab直接提供最新商发版CANN包及配套的ops-nn源码，进入环境后即可开始开发，适合快速上手；Docker方式则提供了更好的环境隔离性，适合需要严格控制依赖版本的场景。

若选择手动搭建环境，需要先安装NPU驱动和CANN包，然后从AtomGit仓库下载与CANN版本配套的ops-nn源码分支。版本配套关系可参考release仓库中的说明文档。使用master分支存在版本不匹配风险，官方强烈建议开发者使用与目标CANN版本精确对应的标签分支。

源码下载的通用命令格式为：

git clone -b ${tag_version} https://atomgit.com/cann/ops-nn.git && cd ops-nn

以CANN 9.0.0版本为例，执行git clone -b 9.0.0 https://atomgit.com/cann/ops-nn.git即可获取配套的源码。若已在开发环境中预置了源码，可通过git branch查询当前版本，通过git checkout ${tag_version}切换到目标分支。

编译与快速运行

ops-nn提供了一套统一的编译脚本build.sh，极大简化了算子编译流程。项目支持两种编译模式：单算子编译和全量编译。

单算子编译适合日常开发迭代，只构建目标算子，编译速度快。以AddExample算子为例，编译命令如下：

bash build.sh --pkg --soc=${soc_version} --ops=add_example -j16

参数--soc指定目标NPU型号。Atlas A2系列取值为ascend910b，Atlas A3系列取值为ascend910_93，950系列取值为ascend950。若提示Self-extractable archive "cann-ops-nn-custom_linux-${arch}.run" successfully created，表明编译成功，run包位于项目根目录的build_out目录下。

编译成功后，安装自定义算子包只需执行run包：

./build_out/cann-ops-nn-*linux*.run

安装路径为${ASCEND_HOME_PATH}/opp/vendors。安装完成后，配置动态库路径：

export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/custom_nn/op_api/lib:${LD_LIBRARY_PATH}

验证算子是否正常工作，运行example样例：

bash build.sh --run_example add_example eager cust --vendor_name=custom

预期输出为算子的实际计算结果，证明算子已成功部署并正确执行。

/// WHY：单算子编译模式将开发闭环压缩到最小范围——修改Kernel代码后只需重新执行编译、安装、验证三步即可看到效果，大幅缩短了开发-验证-迭代的周期。-j16参数利用多核并行编译，对于含有多个算子的大项目能显著节省等待时间。

AI Core算子开发全流程

工程创建

ops-nn提供了build.sh脚本的--genop参数，用于自动创建算子目录骨架。进入项目根目录，执行：

bash build.sh --genop=${op_class}/${op_name}

其中${op_class}为算子分类（如nn、matmul、activation），${op_name}为算子名的小写下划线形式。若指定的分类目录尚不存在，需要在项目根CMakeLists.txt中添加add_subdirectory(${op_class})指令（experimental目录下的算子除外，可直接通过--experimental编译参数处理）。

工程创建成功后，脚本会自动生成完整的目录结构和占位文件。开发者只需在各文件中填充对应逻辑，即可完成一个基础算子的开发。

Tiling切分策略

Tiling是NPU算子开发中最核心的概念之一。由于AI Core内部的 Unified Buffer（UB）容量有限，无法一次性将整个张量数据全部加载到计算单元中进行处理。因此，需要将输入张量切分为多个小块（Tile），逐块加载、计算、写出。用于指导这种数据切分的算法策略，称为Tiling策略。

Tiling策略的实现涉及三个文件：${op_name}_tiling.cpp（Host侧Tiling计算入口）、${op_name}_tiling_key.h（标识不同的Tiling实现路径）和${op_name}_tiling_data.h（存储切分参数的结构体）。Tiling与Kernel之间通过TilingData结构体传递信息，包括总数据量totalLength、切块数量tileNum、每个Tile处理的数据长度等。

在Tiling计算主入口函数中，首先获取平台信息（可用核数、UB大小），然后根据输入张量的shape和数据类型计算切分参数，最后将结果写入TilingData供Kernel侧读取：

static ge::graphStatus TilingFunc(gert::TilingContext* context) {
    // 获取平台信息
    uint64_t ubSize; int64_t coreNum;
    GetPlatformInfo(context, ubSize, coreNum);
    // 获取输入shape
    auto inputX = context->GetInputShape(0);
    auto inputDesc = context->GetInputDesc(0);
    dataType = inputDesc->GetDataType();
    // 计算Tiling参数
    ...
    // 设置TilingData
    auto tiling = context->GetTilingData<AddExampleTilingData>();
    tiling->totalLength = totalIdx;
    tiling->tileNum = TILE_NUM;
}
IMPL_OP_OPTILING(AddExample).Tiling(TilingFunc).TilingParse<CompileInfo>(TilingParse);

/// WHY：Tiling策略直接决定了算子在硬件上的并行度和内存利用效率。不同的shape和数据类型组合可能需要不同的切分策略，因此通过TilingKey可以在Kernel侧根据TilingData中的参数选择最优的计算路径。掌握Tiling是理解NPU算子性能优化的大门。

Kernel实现

Kernel是算子在AI Core上执行的核心代码，使用Ascend C语言编写。Ascend C是华为为昇腾硬件定制的C++扩展语言，其语法与标准C++接近，但提供了针对AI Core硬件特性的内置原语（如向量加法AscendC::Add、矩阵乘法AscendC::Matmul等）。

Kernel实现的完整流程如下：

// 1. 核函数定义
template <uint32_t schMode>
__global__ __aicore__ void add_example(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) {
    // 注册TilingData
    REGISTER_TILING_DEFAULT(AddExampleTilingData);
    GET_TILING_DATA_WITH_STRUCT(AddExampleTilingData, tilingData, tiling);
    // 根据TilingKey选择Kernel实现路径
    if constexpr (schMode == static_cast<uint32_t>(AddExampleTilingKey::TILING_KEY_EXAMPLE_FLOAT)) {
        NsAddExample::AddExample<float> op;
        op.Init(x, y, z, &tilingData);
        op.Process();
    }
}

// 2. Kernel类Process函数（数据流编排）
template <typename T>
__aicore__ inline void AddExample<T>::Process() {
    int32_t loopCount = tileNum_ * BUFFER_NUM;
    for (int32_t i = 0; i < loopCount; i++) {
        CopyIn(i);   // GM -> LM 数据搬入
        Compute(i);  // 计算
        CopyOut(i);  // LM -> GM 数据搬出
    }
}

/// WHY：Process函数中的CopyIn-Compute-CopyOut三步构成了一条流水线。通过开启Double Buffer（BUFFER_NUM=2），可以在当前Tile计算的同时异步搬入下一个Tile的数据，实现计算与通信的流水并行，从而充分隐藏数据搬移的延迟。理解这一流水线思想是编写高效NPU算子的核心。

aclnn接口适配

aclnn是CANN提供的一套标准化算子调用接口，定义了算子的输入输出描述、内存管理和执行调度。算子开发完成后，需要在scripts/kernel/binary_config/ascendc_config.json中注册算子的NPU型号和实现模式：

{
  "name": "AddExample",
  "compute_units": ["${soc_version}"],
  "auto_sync": true,
  "impl_mode": "high_performance"
}

注册完成后，重新编译即会自动生成对应的aclnn接口，开发者可以在应用程序中直接通过aclnn接口调用该算子，无需关心底层L0接口的细节。

调试与验证

ops-nn提供了完整的调试和验证体系，涵盖打印、性能采集和单元测试三个维度。

打印调试方面，Ascend C提供了AscendC::PRINTF接口（支持Scalar类型）和DumpTensor接口（支持打印张量内容），可在Kernel代码中直接调用，在不中断执行流程的情况下输出中间状态：

AscendC::PRINTF("Tiling blockLength is %llu\n", blockLength_);
DumpTensor(zLocal, 0, 128);

性能采集方面，通过msprof工具对算子可执行文件进行profiling，采集结果会自动解析并导出性能数据文件，包括各阶段的执行时间占比、内存带宽利用率等关键指标，帮助开发者定位性能瓶颈。

单元测试方面，项目在tests/ut目录下为op_host、op_kernel和op_api分别提供了独立的测试工程。UT验证无需NPU硬件环境，适合在开发迭代过程中进行快速功能回归验证。

算子分类与核心算子

ops-nn目前开放的算子类别涵盖了深度学习模型中最常用的计算类型。矩阵乘法类（matmul）是整个算子库中数量最多、变体最丰富的类别，包括常规的batch_matmul、量化矩阵乘法quant_batch_matmul_v4、权重量化版本weight_quant_batch_matmul_v2，以及稀疏4:2量化sparse4to2quant_matmul等。量化类算子支持fp8、mxfp8、hifp8、mxfp4等多种低比特数据类型，并支持per_tensor、per_channel、per_token、per_group、per_block等不同的量化粒度组合。

激活函数类（activation）提供了常用非线性激活算子的高性能实现。索引类（index）包括index_fill、masked_scatter、scatter、tf_scatter_add等数据索引和修改操作。损失函数类（loss）则包含了fused_cross_entropy_loss_with_max_sum等融合损失算子，可在单个算子中完成多个数学步骤的融合计算，减少中间结果的访存开销。

experimental目录的引入，允许开发者将自定义实验性算子存放在此处进行调试和验证。通过提交Pull Request，这些算子有机会被正式纳入开源算子库。这一机制为社区贡献者提供了低门槛的参与途径。

使用前vs使用后

对于AI框架开发者和深度学习模型工程师而言，理解ops-nn的价值可以从以下几个维度来对比：

对比维度	使用前（依赖框架预置算子）	使用后（基于ops-nn开发）
算子定制能力	受限于框架提供的通用算子集合，无法针对特定硬件特性调优	可以根据昇腾NPU硬件特性编写专用Kernel，充分利用硬件流水线和UB资源
性能优化空间	通用算子实现，难以针对特定shape或数据类型做极致优化	Tiling策略和Kernel实现完全可控，可针对目标场景的数据分布定制切分方案
量化与稀疏支持	依赖框架层面的量化通路，可能存在多次数据转换开销	低比特量化算子（如mxfp8/mxfp4）和稀疏算子（如sparse4to2quant_matmul）原生集成，减少转换链路
开发环境	仅能在已有算子基础上组合使用	通过CANN Simulator可在无硬件环境下完成算子开发、仿真和调试，降低开发门槛
调试手段	算子作为黑盒运行，内部状态不可见	DumpTensor和PRINTF直接输出Kernel内部数据流，配合msprof性能分析，调试粒度精确到Tile级别
算子复用与组合	新算子需求需要向框架侧提交，等待上游合入	开发者可独立维护自定义算子包，通过vendor_name参数实现与官方算子隔离部署

基于ops-nn进行算子开发后，开发者获得的是对昇腾NPU硬件执行全链路的透明可见性。Tiling切分决定数据如何在UB与GM之间流转，Kernel中的Compute函数直接控制AI Core的计算行为，aclnn接口负责Host-Device协同调度。这套全链路可控的能力，是深度学习框架在使用通用算子时无法获取的。

对于需要极致性能的场景（如大模型推理中的矩阵乘法瓶颈、量化模型的精度调优、稀疏模型的硬件加速），掌握ops-nn的算子开发能力意味着拥有了直接对话硬件的钥匙。而对于需要在昇腾NPU上验证新算法、新结构的学术研究人员，ops-nn提供的QuickStart模板和experimental目录提供了一个低摩擦的起点——从编译运行到自定义开发，整个流程可以在数小时内完成闭环。

结语

ops-nn作为CANN算子生态的开源窗口，正在持续扩展昇腾NPU的算子覆盖范围和应用深度。项目通过标准化的工程结构、完善的文档体系和模块化的开发流程，将原本高度专业化的NPU算子开发工作拆解为可操作、可验证的工程步骤。从AddExample的简单加法到matmul系列的高性能矩阵乘法，从激活函数到融合损失算子，每一类算子的实现都遵循统一的开发范式。

仓库地址：https://atomgit.com/cann/ops-nn