A24b_昇腾算子注册与模型迁移

本文介绍了昇腾NPU模型迁移的全流程与方法论，重点解析了迁移的价值、挑战及解决方案。主要内容包括：迁移价值：昇腾NPU提供更高算力密度和能效比，如船脸识别模型迁移后推理速度提升100倍以上。核心挑战：包括算子适配、动态Shape、精度差异和性能调优等问题，并给出了具体解决方案。四阶段方法：迁移分析：环境准备和兼容性核查代码适配：CUDA接口替换和自定义算子开发精度验证：功能验证和精度比

jieph01

1232人浏览 · 2025-05-24 10:00:48

jieph01 · 2025-05-24 10:00:48 发布

昇腾算子注册与模型迁移

作者：陆璐课题组，瑾丞

目录

昇腾NPU模型迁移全流程解析
未支持算子识别与适配方案
Ascend C算子开发核心实践
算子注册与部署全流程
精度调试与性能调优
实战案例：船脸识别模型迁移

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

前言

在人工智能领域，芯片是决定模型性能的关键因素。华为昇腾（Ascend NPU）作为一款专注于AI计算的高效能芯片，正在成为助力AI开发者的又一个强大工具。昇腾芯片通过其独特的的多核并行计算架构和异构计算能力，为AI模型的训练和推理提供了卓越的能效比。然而，将一个在其他平台上训练好的模型迁移到昇腾平台，需遵循一定的操作步骤。这涉及到模型中算子的适配与开发，特别是当模型中包含昇腾未支持的算子时，开发者需要了解如何自行开发这些算子。

本教程旨在为昇腾AI芯片的初学者提供一个清晰的指南，帮助理解模型迁移的全流程，以及如何处理未支持的算子。我们将通过六个章节，逐步引导您完成从模型迁移分析到算子注册与部署的全过程，帮助您掌握昇腾AI芯片的使用，并能够独立开发和优化算子。通过这个教程内容，我们希望让每一个对昇腾AI芯片感兴趣的开发者都能轻松入门，并在实际应用中取得优秀的性能表现。

1. 昇腾NPU模型迁移全流程解析

在这一章中，我们将探讨把AI模型从其他平台迁移到昇腾NPU的完整流程。模型迁移不仅是将代码从一个平台转移到另一个平台，更是对模型结构和算子的重新审视和调整。我们将会介绍迁移的四个阶段：迁移分析、代码适配、精度验证和性能优化，并详细描述每个阶段的具体的操作步骤和使用的工具。例如，在迁移分析阶段，我们使用迁移分析工具来检测模型中的算子支持度，而在代码适配阶段，我们将演示如何通过PyTorch NPU插件来替换CUDA接口。

通过这些步骤，开发者可以逐步解决精度、性能等问题，最终实现AI模型在昇腾NPU上的高效运行。这一章不仅为后续的算子开发奠定了基础，也为开发者提供了一个清晰的迁移路线图。

1.1 模型迁移价值与挑战

为什么需要迁移模型到昇腾NPU
昇腾NPU通过多核并行计算架构和异构计算特性，提供比传统CPU更高的算力密度，比部分GPU更优的能效比。对于AI模型二次训练、推理优化场景，迁移到昇腾可实现：

计算吞吐量提升：例如船脸识别模型从CPU迁移到NPU后，单步推理耗时从0.55秒降至0.04秒（加速比超100倍）；
全栈兼容性：CANN（Compute Architecture for Neural Networks）工具链支持PyTorch、ONNX等主流框架，降低迁移成本。

核心挑战

挑战类型	典型问题	解决方案
算子适配	GPU模型中含昇腾未支持算子（如`aten::flash_attention`）	通过算子注册流程实现自定义算子开发
动态Shape	模型输入输出形状不固定（如NLP任务）	使用`tiling`策略和动态Shape声明
精度差异	浮点计算精度与GPU存在微小差异	通过`TensorProbe`工具和混合精度配置验证
性能调优	初次迁移可能未充分利用NPU硬件特性	采用`msprof`工具分析瓶颈，结合内存优化策略

1.2 迁移四阶段方法论

阶段1：迁移分析

环境准备：安装CANN（含驱动固件）、PyTorch NPU插件（torch_npu）；
工具扫描：使用迁移分析工具检测模型中的算子支持度（如aten::max_pool2d是否被NPU支持）；
兼容性核查：确认模型依赖的三方库（如detectron2）是否适配昇腾。

阶段2：代码适配

自动迁移：通过import torch_npu替换CUDA接口（如model.cuda() → model.npu()）；
工具迁移：使用msopgen生成算子工程模板，减少手动开发量；
手工迁移：对不支持算子进行自定义开发（如Custom表达或Ascend C）。

阶段3：精度验证

功能验证：运行模型并检查迭代日志（如loss/ppl是否毛刺）；
精度比对：使用TensorProbe比对NPU与GPU的输出差异；
异常处理：CPU回退算子配置（修改npu_native_functions.yaml并重编译框架插件）。

阶段4：性能优化

瓶颈定位：通过msprof采集性能数据，分析op_summary_*.csv中的aic_mte2_time字段；
内存优化：减少GM（Global Memory）与UB（Unified Buffer）搬运次数（如512B对齐、UB缓冲区融合）；
并行策略：配置多核负载均衡（如SetBlockDim(MAX_AICORE_NUM)）。

迁移闭环流程图

迁移分析 → 代码适配 → 精度验证 → 性能优化 → 部署上线

1.3 迁移工具链全景图

核心工具及其协同关系

工具名称	功能	使用场景示例
CANN	算子开发编译环境	`msopgen`工程生成、`cmake`编译动态库
PyTorch NPU插件	框架级适配	`import torch_npu`替换CUDA接口
ONNX导出工具	模型结构转换	将`.pth`文件导出为ONNX格式进行分析
msprof	性能数据采集	`msprof --output="./out" --ai-core=on`
TensorProbe	精度调试	检测`aten::softmax`算子的输出差异
msopst	算子ST测试	验证`add_custom`算子的逻辑正确性

工具链协同流程

模型分析：通过msopgen生成迁移报告，明确需替换的算子清单；
算子开发：使用Ascend C实现自定义算子（如MaxPool2d）；
部署验证：msopst测试算子功能，msprof采集性能数据；
迭代优化：根据op_summary.csv中aic_mac_ratio等指标调整Tiling策略。

1.4 昇腾NPU模型迁移示例

案例：船脸识别模型迁移

原始模型：基于PyTorch的VOLOv5模型，使用GPU训练；
迁移目标：在昇腾NPU上实现推理加速。

迁移步骤

迁移分析
- 使用迁移分析工具扫描模型，发现aten::max_pool2d等算子需替换；
- 确认CANN 7.0支持该模型所需三方库。

代码适配

自动替换CUDA接口：

import torch
import torch_npu
from torch_npu.contrib import transfer_to_npu
model = model.npu()  # 替换CUDA的model.cuda()

手工替换未支持算子：

// 使用Ascend C实现自定义MaxPool2d
extern "C" __global__ __aicore__ void max_pool2d_kernel(...) {
   // Tiling逻辑和UB数据搬运
   pipe.InitBuffer(ub_tensor, 1, input_size * sizeof(float));
   // Vector计算逻辑
   AscendC::MaxPool(...);
}

精度验证
- 使用TensorProbe比对GPU与NPU的loss曲线，确认无显著毛刺；
- 检查op_summary.csv中aic_mte2_ratio是否接近理论值。
性能优化
- 内存优化：调整DataCopy的blockLen为32KB以上，减少搬运次数；
- 并行策略：配置SetBlockDim(8)，利用8个AI Core并行计算；
- 最终效果：NPU推理耗时从0.55秒降至0.04秒，线性度提升至85%。

迁移效果对比

指标	CPU版本	NPU版本（未优化）	NPU优化后
单步耗时	0.55s	0.35s	0.04s
线性度	-	60%	85%

1.5 附录：关键参数速查表

TilingData结构优化建议

（1）减少冗余字段，示例代码：

BEGIN_TILING_DATA_DEF(TilingData);
TILING_DATA_FIELD_DEF(uint8_t, formerNum);
TILING_DATA_FIELD_DEF(uint32_t, formerLength);
END_TILING_DATA_DEF;

（2）字节对齐：pipe.InitBuffer(ub_tensor, 1, 32KB)（强制512B对齐）

性能优化优先级

计算密集型：优先优化Cube矩阵计算；
内存访问：减少GM→UB搬运次数；
同步控制：避免PipeBarrier在关键路径插入过多同步点。

小结

模型迁移至昇腾NPU的核心在于算子适配和硬件特性利用。通过四阶段闭环流程，开发者可逐步解决精度、性能问题，最终实现AI模型在昇腾上的高效运行。后续章节将详细拆解每个阶段的实现细节，包括Ascend C的内存管理、Tiling策略设计等。

2. 未支持算子识别与适配方案

在模型迁移过程中，识别模型中未被昇腾NPU支持的算子是一个关键步骤。本章重点介绍了算子支持度分析方法，包括如何使用迁移分析工具来生成未支持算子清单。我们还探讨了动态shape检测和三方库兼容性核查的方法，并提供了了选择算子适配策略的建议。

在昇腾NPU平台上，开发者可以选择自动迁移、工具迁移或手工迁移的方式来适配未支持算子。每个方法都有其适用场景和操作步骤，例如，自动迁移适用于简单模型，而工具迁移适合复杂模型，手工迁移则适用于高定制需求。本章通过具体案例，如船脸识别模型迁移，展示了如何识别和适配未支持算子，并提供了实用的开发建议。

2.1 算子支持度分析方法

2.1.1 迁移分析工具使用指南

操作步骤：

环境准备：确保已安装CANN工具包（含PyTorch NPU插件），并配置环境变量ASCEND_CANN_PATH。

执行分析命令：

torch::npu::analyze_model --input_script=your_model_script.py \
                           --output_dir=./analysis_results \
                           --device_count=8

查看分析结果：
- 未支持算子清单：工具会输出unsupported_ops.csv文件，记录所有不被NPU支持的算子及其调用路径。
- 动态shape检测：工具会标记是否存在动态输入，输出文件dynamic_shape_info.txt。
- 三方库兼容性核查：若模型依赖PyTorch以外的库（如torchvision），工具会检查兼容性并输出third_party_check.log。

关键字段说明：

字段	含义	用途
`OpName`	算子名称	识别需要替代的算子
`DynamicShape`	是否为动态shape	决定是否需要特殊处理
`ThirdParty`	依赖的三方库	核对是否需回退或替代

2.1.2 动态shape检测

操作流程：

模型脚本标注：在PyTorch模型中使用torch.nn.AdaptiveAvgPool2d等动态shape算子时，需在代码中添加dynamic_shape=True标注。
运行分析工具：工具会检测模型输入是否动态，并输出dynamic_shape_info.txt，标记各层输入shape变化范围。
适配策略：
- 静态shape模型：无需特殊处理。
- 动态shape模型：需在算子实现中声明动态shape支持，并使用TilingData的shape_range字段。

代码示例：

# 在模型定义中添加动态shape标注
self.dynamic_pool = torch.nn.AdaptiveAvgPool2d((7, 7))
self.dynamic_pool.to(&#x27;npu&#x27;)  # 升级到NPU并启用动态shape

2.1.3 三方库兼容性核查

检查方式：

自动检查：迁移分析工具会扫描代码中的三方库调用，并输出third_party_check.log。
人工核查：若工具未覆盖的三方库（如apex），需手动检查是否提供NPU适配版本。

常见问题场景：

三方库	问题描述	解决方案
`apex`	混合精度训练模块未适配NPU	替换为PyTorch内置的`torch.cuda.amp`，并修改为`torch.npu.amp`
`detectron2`	自定义ROIAlign算子未适配	使用`torch_npu.contrib.transfer_to_npu`工具自动迁移，或手写替代算子

2.2 算子替代策略选择

2.2.1 自动迁移（推荐简单模型）

适用场景：

模型结构简单，算子调用路径清晰。
无特殊算子需求（如自定义归约）。
三方库兼容性无问题。

操作步骤：

代码修改：在模型脚本头部添加必要插件：

import torch
import torch_npu
from torch_npu.contrib import transfer_to_npu

替换CUDA接口：
- 将model.cuda()替换为model.npu()。
- 将loss = loss.cuda()替换为loss = loss.npu()。

执行迁移：

python your_train_script.py --backend=npu

验证成功率：观察训练日志是否正常打印迭代信息。

注意事项：

环境变量配置：迁移前需设置ASCEND_DEVICE_ID=0，指定NPU设备。
回退机制：若出现算子不支持报错，需启用CPU回退（见2.2.3节）。

2.2.2 工具迁移（推荐复杂模型）

适用场景：

模型包含复杂结构（如多尺度ROIAlign）。
三方库部分适配，需人工干预。

操作步骤：

环境准备：

pip install ascend-cann-toolkit pandas libcst

执行迁移命令：

python -m torch_npu.migrate --input_script=your_model_script.py \
                            --output_script=your_model_npu.py \
                            --config_file=migrate_config.yaml

配置文件示例（migrate_config.yaml）：

third_party:
 - torchvision
 - apex
dynamic_shapes: True

验证迁移结果：
- 精度验证：运行迁移后的脚本，检查loss是否收敛。
- 性能验证：使用msprof采集性能数据，对比CPU版本。

2.2.3 手工迁移（推荐高定制需求）

适用场景：

模型包含自定义算子（如FP量化参数处理）。
需要精确控制算子行为（如多核负载均衡）。

操作步骤：

导入NPU相关库：
```
import torch_npu
```

指定NPU设备：

model = model.npu()  # 手工迁移模型

替换CUDA接口：

# 原CUDA接口
data = data.cuda()
# 替换为NPU接口
data = data.npu()

自定义算子开发：若模型中存在未支持算子（如FP量化参数），需手写算子（见2.3节）。

2.3 自定义算子开发路径

2.3.1 基于Custom表达的快速开发 vs Ascend C的高性能开发

对比表格：

特性	Custom表达	Ascend C
开发难度	低	高
性能表现	一般	优秀
适用场景	快速验证、小规模模型	高性能需求、大规模模型
开发语言	Python	C++
编译流程	无需编译，直接调用	需要编译为动态库

推荐选型建议：

Custom表达：适合快速验证或小规模模型。
Ascend C：适合高性能需求（如矩阵乘累加、双缓冲机制）。

2.3.2 Custom表达自定义算子开发

开发流程：

定义算子接口：

from torch_npu import ops
class CustomAdd(torch.autograd.Function):
   @staticmethod
   def forward(ctx, x, y):
       # 使用Custom表达注册算子
       return ops.custom_add(x, y)

注册算子：

# 使用PyTorch NPU插件注册算子
python register_custom_op.py

验证算子：
- 精度验证：使用TensorProbe工具对比CPU和NPU输出。
- 性能验证：使用msprof采集算子耗时，对比基线。

代码示例：

# 在模型中使用自定义算子
class MyModel(torch.nn.Module):
    def forward(self, x, y):
        return CustomAdd.apply(x, y)

2.3.3 Ascend C高性能算子开发

开发流程：

工程创建：

msopgen gen -i ./add_custom.json -c ai_core-ascend910b

编写核函数（以Add算子为例）：

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

编译部署：

./build.sh
./custom_opp_ubuntu_x86_64.run --install-path=/usr/local/Ascend/vendors/

调试验证：

ST测试用例：

./msopst run -i ./add_custom_case.json -soc ascend910b -out ./st_results

CPU仿真验证：

./msopgen --simulate-cpu ./add_custom.cpp

关键参数说明：

参数	含义	用途
`blockDim`	多核并行配置	指定调用核函数的核数
`TillingData`	Tiling结构	定义算子切分策略
`workspace`	临时内存	用于Cube和Fixpipe计算

2.3.4 CPU回退算子配置规范

配置方法：

修改配置文件（op_def.yaml）：

ops:
 - name: "CustomAdd"
   fallback: "cpu"
   priority: 1

重新编译模型：

python custom_opgen --config op_def.yaml

验证CPU回退：
- 运行模型时，观察未支持算子是否被回退到CPU执行。
- 使用TensorProbe确保回退算子精度与CPU版本一致。

代码示例：

# 在迁移脚本中指定回退
if not_npu.is:
    model = model.n_add()  # 使用CPU定义的CPU回退算子
else:
    model = model.npu()

2.4.5 算子替代策略选择

策略选择图：

选择建议：

自动迁移：适用于简单模型，无需修改代码。
工具迁移：适用于复杂模型，需人工检查三方库。
手工迁移：适用于高定制需求，需编写算子逻辑。

附：算子替代优先级策略

优先级	算子类型	适配策略	说明
高	简单算子（如Add）	自动迁移	无需修改代码
中	动态结构算子（如ROIAlign）	工具迁移	需要人工配置
低	**复杂算子（如FP量化）	手工迁移	需要编写Ascend C代码

通过上述步骤，开发者可以系统地识别未支持算子，并选择合适的的适配策略。后续章节将深入Ascend C算子开发的细节。

3. Ascend C算子开发核心实践

本章深入解析了Ascend C算子开发的核心实践，包括算子结构设计原则、数据流与内存优化，以及并行计算编程范式。Ascend C是一种类似于CUDA-C的编程工具，用于在昇腾NPU上开发算子。我们讨论了如何使用Ascend C开发算子，包括如何设计Host侧Tiling函数和Device侧Kernel函数，以及如何通过TilingData结构优化算子性能。

通过本章的学习，开发者可以了解如何在昇腾NPU上高效开发算子，以及如何通过优化内存使用和并行计算来提升算子性能。我们还提供了一些实用的开发流程总结和开发者速查表，帮助开发者快速定位和解决问题。

3.1 算子开发环境搭建

3.1.1 CANN工具链安装

Ascend C开发依赖CANN（Compute Architecture for Ascend）工具链，安装步骤如下：

安装依赖

Debian系列（Ubuntu、UOS）：

apt-get install -y gcc make net-tools cmake python3 python3-dev python3-pip

openEuler系列（CentOS、openEuler）：

yum install -y gcc make net-tools cmake python3 python3-devel 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

验证安装
- 检查CANN安装目录是否存在/usr/local/Ascend/ascend-toolkit/，并确认set_env.sh已生效。

3.1.2 仿真与真实设备调试

仿真环境：无需真实NPU设备即可验证逻辑，适用于早期开发阶段。
真实设备：需通过npu-smi info获取芯片型号（如Chip Name为Ascendxxx），并配置环境变量ASCEND_INSTALL_PATH。
调试工具：
- CPU仿真：使用ICPU_RUN_KF宏调用核函数。
- NPU调试：通过aclrtCreateStream和aclrtSynchronizeStream管理异步任务。

3.2 算子结构设计原则

3.2.1 Host侧Tiling函数设计

Tiling函数负责将全局内存（GM）中的数据切分为小块，供Device侧核函数处理。其设计需遵循以下原则：

输入输出约束：根据算子输入输出的shape动态计算切分策略。
多核并行支持：通过GetBlockNum获取可用核数，确保BlockDim不超过硬件限制（如MAX_AICORE_NUM）。
内存对齐：GM地址需512B对齐（如Atlas A2设备），否则带宽效率下降30%以上。

代码示例：

static ge::graphStatus TilingFunc(gert::TilingContext *context) {
    TilingData tiling;
    uint32_t totalLength = context->GetInputTensor(0)->GetShapeSize();
    context->SetBlockDim(BLOCK_DIM);
    tiling.formerNum = (totalLength / ALIGN_NUM) % BLOCK_DIM;
    tiling.tailNum = BLOCK_DIM - tiling.formerNum;
    context->SaveTilingData(&tiling);
    return GRAPH_SUCCESS;
}

3.2.2 Device侧Kernel函数设计

Kernel函数在Device侧执行计算逻辑，需遵循以下规范：

核函数定义：使用__global__ __aicore__限定符标识为设备侧函数，输入输出参数使用GM_ADDR宏定义。
内存管理：通过TPipe和TQue分配队列内存，确保TilingData结构大小合理（减少冗余字段，如使用uint8_t代替uint64_t）。
同步控制：默认使用auto sync编译选项，开发者需手动插入PipeBarrier或SetFlag/WaitFlag处理多核同步问题。

代码示例：

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

3.3 数据流与内存优化

3.3.1 GM地址对齐与性能差异

GM地址对齐直接影响数据搬运效率。以下对比Atlas A2设备上不同对齐方式的性能差异：

对齐方式	数据搬运量 (KB)	实际带宽利用率 (%)	耗时 (us)
512B对齐	16K	95%	13.2
32B对齐	16K	70%	18.9

优化建议：

使用GMEM_ALIGNMENT宏检查地址对齐，确保输入输出地址为512B的倍数。
通过SetShapeInfo减少不必要的shape信息传递，降低栈内存占用。

3.3.2 UB缓冲区融合策略

通过UB（Unified Buffer）缓冲区减少GM与UB之间的冗余搬运：

连续计算：将多个小算子（如Exp+Abs）融合为一个UB内计算，避免中间结果写回GM。
共享缓冲区：复用TBuf对象，减少TilingData结构中的冗余参数。

反例与正例对比：

// 反例：多次搬运导致性能损失
DataCopy(aLocal, aGM, size);
Exp(aLocal);
DataCopy(bLocal, aLocal, size);
Abs(bLocal);

// 正例：UB融合减少搬运次数
TBuf<QuePosition::VECCALC> sharedBuf;
pipe.InitBuffer(sharedBuf, size);
LocalTensor<float> temp = sharedBuf.Get<float>();
Exp(temp);
Abs(temp);

3.3.3 FP Buffer量化参数优化

对于量化算子（如FP16→INT8），利用FP Buffer（Fixpipe Buffer）减少数据搬运：

量化参数存储：将量化参数搬运至FP Buffer，通过Fixpipe接口直接完成量化计算。
性能提升：减少CO1→GM→UB的搬运步骤，提升内存带宽利用率。

代码示例：

// 量化参数搬运至FP buffer
LocalTensor<int8_t> dstLocal = outQueueDst.AllocTensor<int8_t>();
FixpipeParams fpParams;
fpParams.quantMode = QuantMode::FP16_TO_INT8;
Fixpipe(dstLocal, dstLocal, fpParams);

3.4 并行计算编程范式

3.4.1 多核负载均衡策略

多核并行时需确保每个核的计算量均衡，避免尾核（Tail Core）拖慢整体性能：

动态切分：根据BlockDim和数据总长度totalLength计算每个核的tileLength。
尾块优化：若尾块长度较小（如< 64B），将其分配到主核处理，尾核仅处理剩余数据。

代码示例：

constexpr uint32_t totalLength = 8 * 2048;
constexpr uint32_t blockLength = totalLength / BlockDim;
LocalTensor<float> xLocal = inQueueX.AllocTensor<float>(blockLength);

3.4.2 双缓冲机制实现

双缓冲（Double Buffering）通过并行搬运和计算隐藏数据搬运时间：

内存分配：使用TQueBind绑定输入输出队列，减少冗余拷贝。
循环次数：Process函数中需确保loopCount >= 2以启用双缓冲。

性能对比：

未启用双缓冲：Vector利用率33%（每次计算等待搬运完成）。
启用双缓冲：Vector利用率75%（搬运与计算并行）。

代码示例：

// 双缓冲配置
TQueBind<QuePosition::VECIN, QuePosition::VECOUT, 2> bindQue;
bindQue.Init(xGm, yGm, zGm);
for (uint32_t i = 0; i < loopCount; i++) {
    CopyIn(i);
    Compute(i);
    CopyOut(i);
}

3.4.3 异步计算与同步控制

在MIX场景（Cube+Vector混合计算）中，通过异步接口减少核间通信开销：

异步调用：使用Iterate<false>减少Cube与Vector核间的同步信号。
同步点控制：仅在必要时插入PipeBarrier，避免频繁同步。

代码示例：

// 异步调用Matmul
while (mm.Iterate<false>()) {
    LocalTensor<float> cLocal = outQueueCO1.AllocTensor<float>();
    Mmad(cLocal, aLocal, bLocal, mmParams);
    outQueueCO1.EnQue(cLocal);
}

3.4 Cube-Vector融合算子开发

3.4.1 L1/L0C内存规划

在分离架构中，Cube计算结果需暂存至L0C Buffer，随后通过FP Buffer直接量化搬运：

Cube输出累加：使用Mmad的cmatrixInitVal参数实现多块结果的累加。
FP Buffer配置：通过FixpipeParams设置量化模式，减少中间搬运。

代码示例：

// Cube输出累加
MmadParams mmadParams;
mmadParams.cmatrixInitVal = true; // 启用累加
Mmad(cLocal, aLocal, bLocal, clocal, mmadParams);

3.4.2 双缓冲区（CO1 buffer）复用

在Cube计算中，若尾核数据量较小，可将其复用至主核处理：

主尾块分离：主核处理完整数据块，尾核仅处理余数部分。
内存复用：通过TQueBind绑定主尾块内存，减少队列管理开销。

性能数据：

复用前：搬运次数2n（n次Vector计算）。
复用后：搬运次数2（仅首尾块搬运）。

3.5 开发者速查表

3.5.1 TilingData结构优化对比

字段	冗余字段	优化后字段	内存节省 (B)
formerNum	uint64_t	uint8_t	64 → 8
tailNum	uint64_t	uint8_t	64 → 8
formLength	uint64_t	uint32_t	64 → 4
总计	-	-	128 → 28

3.5.2 内存优化关键API

场景	推荐API	说明
小矩阵长驻L1	`LoadData2dTransposeParams`	减少大矩阵分块次数
量化参数搬运	`Fixpipe`	通过FP buffer减少搬运次数
冗余搬运避免	`TQueBind`	绑定输入输出队列，省略搬运

3.6 典型开发流程总结

环境准备：安装CANN并配置环境变量（如source /usr/local/Ascend/ascend-toolkit/set_env.sh）。
算子结构设计：分Host侧Tiling和Device侧Kernel，确保TilingData结构精简。
内存优化：实现512B对齐、UB缓冲区融合、FP量化参数搬运。
并行编程：通过双缓冲、异步计算提升Vector/Cube利用率。

开发者提示：

调试优先：使用仿真环境（如msprof工具）验证逻辑，再迁移至真实设备。
性能调优：优先优化计算密集型算子，再处理内存访问瓶颈。

通过以上实践，开发者可高效开发Ascend C算子，并在昇腾NPU上实现模型迁移。

4. 算子注册与部署全流程

本章将详细介绍算子注册和部署的完整流程。算子注册是将自定义算子加入到昇腾NPU算子库的过程，而部署则是将算子安装到指定路径，以便在昇腾NPU上运行。我们解释了如何使用msopgen工具生成算子工程，如何使用CMake进行编译配置，以及如何进行算子调试和验证。

算子注册与部署是算子开发的关键环节，通过本章的学习，开发者可以掌握如何将自定义算子注册并部署到昇腾NPU上，从而完成模型迁移的最后一步。我们还提供了TilingData结构优化的示例，以及算子部署路径的规范，确保开发者能够顺利进行算子开发和部署。

4.1 算子原型定义规范

算子原型定义是开发自定义算子的第一步，也是框架识别算子的基础。通过OpDef类定义算子原型，需要明确输入输出参数的约束条件和动态shape支持声明。

4.1.1 算子原型定义

使用OpDef类定义算子原型，通常在AddCustom类中完成。每个输入输出参数需声明名称、格式、数据类型等信息。

4.1.2 参数约束

必填参数：使用ParamType(REQUIRED)声明，如Input("x")。
支持的数据类型：如DataType({ge::DT_FLOAT16, ge::DT_FLOAT})。
支持的数据格式：如Format({ge::FORMAT_ND})。
动态shape支持：在输入输出定义中添加dynamic_shape属性。

4.1.3 示例代码

namespace ops {
    class AddCustom : public OpDef {
    public:
        explicit AddCustom(const char* name) : OpDef(name) {
            // 输入参数定义
            this->Input("x")
                .ParamType(REQUIRED)
                .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
                .Format({ge::FORMAT_ND});
            this->Input("y")
                .ParamType(REQUIRED)
                .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
                .Format({ge::FORMAT_ND});

            // 输出参数定义
            this->Output("z")
                .ParamType(REQUIRED)
                .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
                .Format({ge::FORMAT_ND});
        }

        // 其他方法定义
    };
}

4.1.4 动态shape声明

在算子原型定义中，动态shape的支持是通过dynamic_shape属性声明的。这使得算子能够适应不同输入形状，提高灵活性。

4.2 Tiling策略实现

Tiling策略是算子在Host侧进行数据切分的关键，它决定了如何在Device侧进行计算。

4.2.1 TilingData结构定义

TilingData结构需要定义切分所需的字段，如totalLength、tileNum等。结构设计需考虑字段的大小和对齐。

4.2.2 示例代码

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);
    TILING_DATA_FIELD_DEF(uint32_t, alignNum);
END_TILING_DATA_DEF;

4.2.3 内存切分算法

内存切分算法需要根据输入输出数据的大小和硬件特性进行设计。常见的切分算法包括：

均匀切分：将数据均匀分配到每个核上。
非均匀切分：根据核数和数据大小进行非均匀分配，减少冗余字段。

4.3 算子编译部署

算子开发完成后，需要进行编译和部署，以便在昇腾AI处理器上运行。

4.3.1 msopgen工程生成

使用msopgen工具生成算子工程，根据原型定义文件生成开发模板。

${INSTALL_DIR}/python/site-packages/bin/msopgen gen -i $HOME/sample/add_custom.json -c ai_core-<soc_version>

4.3.2 CMake编译配置

在算子工程目录下，使用CMake进行编译配置。关键配置包括：

算子实现文件：add_custom.cpp
Tiling函数：TilingFunc
Shape推导函数：InferShape

// Tiling函数示例
static ge::graphStatus TilingFunc(gert::TilingContext *context) {
    TilingData tiling;
    uint32_t totalLength = context->GetInputTensor(0)->GetShapeSize();
    context->SetBlockDim(BLOCK_DIM);
    tiling.set_totalLength(totalLength);
    tiling.set_tileNum(TILE_NUM);
    return ge::GRAPH_SUCCESS;
}

4.3.3 动态库安装路径

编译完成后，生成的动态库会安装到指定路径。路径配置如下：

默认路径：ASCEND_CUSTOM_OPP_PATH
自定义路径：使用--install-path=<path>指定安装目录。

./custom_opp_<target_os>_<target_arch>.run --install-path=$HOME/Ascend/vendors/custom

4.4 算子调试验证

算子开发完成后，需要进行调试验证以确保其正确性和性能。主要分为ST测试用例编写、CPU仿真验证和NPU真实环境验证。

4.4.1 ST测试用例编写

ST测试用例定义文件通常以.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"
            },
            {
                "format": ["ND"],
                "type": ["float16"],
                "shape": [8, 2048],
                "data_distribute": ["uniform"],
                "value_range": [[0.1, 1.0]],
                "name": "y"
            }
        ],
        "output_desc": [
            {
                "format": ["ND"],
                "type": ["float16"],
                "shape": [8, 2048],
                "name": "z"
            }
        ]
    }
]

4.4.2 CPU仿真验证

在CPU仿真环境中验证算子功能，确保其在Host侧的正确性。

// CPU模式调试示例
AscendC::SetKernelMode(KernelMode::AIV_MODE);
AscendC::AddCustomKernel::Process(...);

4.4.3 NPU真实环境验证

将算子部署到NPU设备上进行验证，使用性能分析工具进行性能优化。

// NPU模式运行示例
aclrtCreateStream(&stream);
aclrtMemcpyDeviceToHost(...);
aclrtSynchronizeStream(stream);

4.5 算子注册与部署流程图

4.5.1 算子注册流程

算子注册流程包括定义原型、生成Tiling函数、编译部署、注册到算子库。

步骤	描述
1	使用`OpDef`类定义算子原型
2	生成Tiling函数和Shape推导函数
3	编译生成动态库并安装到指定路径
4	注册到昇腾算子库

4.5.2 部署流程

算子部署流程包括编译、安装、注册和验证。

步骤	描述
1	编译算子工程，生成动态库
2	安装动态库到指定路径
3	注册算子到昇腾算子库
4	在NPU设备上验证算子功能和性能

4.6 TilingData结构优化

在昇腾NPU上进行算子开发时，TilingData结构的优化是提高算子性能的重要的一环。优化策略包括：

减少冗余字段：根据实际需求移除不必要的字段。
选择合适的变量类型：根据计算逻辑选择合适的变量类型。
合理排布字段：确保字段在内存中的合理排布，减少对齐开销。

4.6.1 反例

BEGIN_TILING_DATA_DEF(TilingDataUnalign)
    TILING_DATA_FIELD_def(uint64_t, blockDim);
    TILING_DATA_field_def(uint64_t, formerNum);
    TILING_DATA_field_def(uint64_t, tailNum);
    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;

4.6.2 正例

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);
    TIL_DATA_field_def(uint32_t, alignNum);
END_TILE_DATA_DEF;

4.6.3 字段排布优化

在昇腾NPU上，内存对齐要求通常是512B，合理排布字段可以减少对齐开销。

字段	类型	描述
formerNum	uint8_t	表示前次计算涉及的核数
tailNum	uint8_t	表示尾次计算涉及的核数
formerLength	uint32_t	表示前次计算的数据长度
tailLength	uint32_t	表示尾块计算数据长度
alignNum	uint32_t	表示对齐参数

4.7 算子部署路径规范

算子部署时，需要确保动态库安装到正确的路径，并更新环境变量以使算子生效。

4.7.1 默认路径

ASCEND_CUSTOM_OPP_PATH=/usr/local/Ascend/vendors/custom

4.7.2 自定义路径

export ASCEND_INSTALL_PATH=$HOME/Ascend
./custom_opp_ubuntu_x86_64.run --install-path=$HOME/Ascend/Ascend/vendors/custom

4.7.3 部署后验证

source $HOME/Ascend/vendors/custom/bin/set_env.bash

4.8 算子调试工具使用

4.8.1 TensorProbe工具

TensorProbe工具用于调试算子的精度问题，特别是在计算结果是否符合预期。

4.8.2 性能分析工具

性能分析工具用于采集和分析算子的性能数据，识别性能瓶颈。

msprof --output="./out" --ai-core=on --aic-m=on

通过以上步骤，开发者可以将自定义算子，将其注册并部署到昇腾NPU上，完成从模型迁移到性能优化的全过程。希望本章能够清晰展示算子注册与部署的完整流程，为后续性能优化和调试提供坚实基础。

5. 精度调试与性能调优

在这一章中，我们将介绍如何进行算子的精度验证和性能调优。精度验证是确保算子计算结果与CPU/GPU版本一致的重要步骤，而性能调优则是为了最大化算子在昇腾NPU上的计算效率。我们介绍了使用TensorProbe工具进行精度验证，以及使用msprof工具进行性能分析。

通过本章的学习，开发者可以了解如何使用工具链进行瓶颈定位和优化，并掌握性能调优的优先级策略。我们还通过典型优化场景示例，展示了如何通过双缓冲机制、Vector指令模式切换和L2Cache切分策略来提升算子性能。实用的调试与优化技巧帮助开发者更好地理解和优化算子性能。

5.1 算子精度验证方法

精度验证是算子开发中不可忽视的环节。昇腾AI处理器遵循IEEE 754浮点标准，但并行计算中浮点运算的顺序可能引发精度差异。

5.1.1 TensorProbe工具的使用

TensorProbe是精度验证的核心工具，通过采集输入输出张量数据，对比昇腾与CPU/GPU的计算结果差异。

操作步骤：

启用TensorProbe：在训练脚本中设置环境变量TENSOR_PROBE_ENABLE=1，并指定输出路径。
采集数据：通过TensorProbe接口将关键算子的中间结果保存到指定路径。
对比分析：使用TensorProbe配套工具对采集结果进行差值计算，判断是否满足误差容忍范围。

# 示例：在PyTorch中启用TensorProbe（需结合昇腾适配版本）
import torch_npu
from torch_npu.contrib import TensorProbe

TensorProbe.enable(path="./probe_results")

示例输出：

算子名称	误差类型	误差范围	建议措施
MatMulV2	累积误差	< 1e-3	保留浮点精度计算
SoftMax	位移误差	< 1e-5	使用FP32中间计算

5.1.2 混合精度配置检查

混合精度计算是提升性能的关键，但若配置不当可能导致精度异常。

检查要点：

输入输出类型一致性：确保算子输入输出的数据类型与混合精度模式匹配。
量化参数合法性：检查量化参数是否在合理范围内（如RoundMode设置是否为CAST_FLOOR或CAST_CEIL）。
同步控制：在涉及多核并行的场景中，检查同步指令是否插入正确（如PipeBarrier的使用）。

代码片段（需结合TensorProbe验证结果调整）：

AscendC::SetMaskCount();
AscendC::SetVectorMask<DTYPE_X, AscendC::MaskMode::COUNTER>(ELE_SIZE);
AscendC::Add<DTYPE_X, false>(zLocal, xLocal, yLocal, ELE_SIZE);

5.1.3 浮点计算差异处理

浮点计算的顺序差异可能导致精度偏差。例如，A + B + C在昇腾上可能与CPU/GPU结果略有不同。

处理策略：

严格顺序计算：在涉及浮点计算的场景中，避免并行计算引发的顺序差异。
高精度中间计算：将部分计算提升到FP32精度后，再转换回目标精度（如FP16）。
容忍范围设置：在验证环节允许一定误差（如<1e-5），避免因微小差异导致误判。

代码片段（FP32中间计算示例）：

// 将FP16计算提升到FP32再降回FP16
LocalTensor<float> tmpLocal = inQueueTmp.AllocTensor<float>();
Cast(tmpLocal, srcLocal, RoundMode::CAST_FLOOR, size);
Add(dstLocal, tmpLocal, biasLocal, size);
Cast(dstLocal, tmpLocal, RoundMode::CAST_FLOOR, size);

5.2 性能分析工具链

性能调优需要依赖工具链精准定位瓶颈。昇腾生态提供了msprof、op_summary.csv等工具，帮助开发者快速分析算子执行情况。

5.2.1 msprof性能采集

msprof支持采集算子的流水线利用率、内存带宽使用情况等关键指标。

操作步骤：

启用性能采集：在训练脚本中设置--ai-core=on和--aic-metrics="PipeUtilization"。
执行训练：通过msprof采集性能数据，输出路径由--output指定。
解析结果：使用msprof配套工具解析op_summary_*.csv文件，识别性能瓶颈。

示例命令：

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

5.2.2 op_summary.csv解读

op_summary_*.csv文件记录了算子执行的详细流水线信息，如aic_mte2_ratio、aic_cube_ratio等字段。

关键字段说明：

字段名称	含义	建议
`aic_mte2_ratio`	MTE2流水利用率	> 90% 为理想状态
`aic_cube_ratio`	Cube流水利用率	> 80% 为理想状态
`aic_mte3_ratio`	MTE3流水利用率	> 85% 为理想状态

示例数据：

算子名称	Duration	Cube Ratio	MTE2 Ratio	MTE3 Ratio
MatMulV2	1350us	80%	95%	85%
BiasAdd	210us	92%	70%	88%

5.2.3 瓶颈定位方法

性能瓶颈通常出现在流水线利用率不足或数据搬运频繁的场景。

定位策略：

流水线利用率不足：检查aic_cube_ratio、aic_mte2_ratio是否低于预期。
数据搬运频繁：观察aic_mte2_time是否接近理论值（如搬运16KB数据需<100us）。
冗余同步：检查PipeBarrier插入是否合理（如PipeBarrier<PIPE_ALL>是否在多核计算后引入）。

代码片段（冗余同步示例）：

// 反例：未使用PipeBarrier导致流水线串行
Add(dstLocal, src0Local, src1Local, size);
// 正例：插入PipeBarrier隐藏流水线等待
PipeBarrier<PIPE_V>();
Add(dstLocal, src0Local, src1Local, size);

5.3 性能优化优先级策略

性能调优需分层进行，优先优化计算密集型部分，再解决内存和通信问题。

5.3.1 计算密集型优化

优化目标：最大化Cube和Vector计算单元的利用率。

关键措施：

Cube优化：确保矩阵乘法使用Mmad接口，且aic_cube_ratio接近100%。
Vector优化：使用Add、Exp等指令时，确保aic_vec_ratio不低于85%。

代码片段（Cube优化示例）：

// 使用Mmad接口进行矩阵乘
MmadParams mmadParams;
mmadParams.m = m;
mmadParams.n = n;
mmadParams.k = k;
Mmad(c1Local, a2Local, b2Local, mmadParams);

5.3.2 内存访问优化

优化目标：减少数据搬运次数，提升UB/FP Buffer利用率。

关键措施：

512B对齐：对GM地址进行512B对齐，提升搬运带宽利用率。
UB融合：将连续计算的中间结果暂存于UB中，避免多次搬运。
FP Buffer量化参数：将量化参数提前搬运到FP buffer，减少计算中搬运开销。

代码片段（UB融合示例）：

// 反例：Exp后搬运到GM再搬运回UB
Exp(dstLocal, srcLocal, size);
DataCopy(gmOut, dstLocal, size);
DataCopy(srcLocal, gmOut, size);
// 正例：Exp结果直接用于Abs计算
Exp(dstLocal, srcLocal, size);
Abs(dstLocal, srcLocal, size);

5.3.3 通信效率优化

优化目标：减少核间通信开销，提升多核并行效率。

关键措施：

双缓冲机制：通过pipe.InitBuffer(..., 2)设置双缓冲，减少流水线等待。
核间负载均衡：确保BlockDim不超过实际核数（如MAX_AICORE_NUM）。

代码片段（双缓冲机制示例）：

// 启用双缓冲时，内存块数设为2
pipe.InitBuffer(inQueueA1, 2, aSize * sizeof(half));

5.4 典型优化场景示例

5.4.1 双缓冲配置优化

场景描述：算子执行时，Vector计算单元等待数据搬运时间占比过高。

优化方法：使用双缓冲机制（double buffer），将数据搬运与计算并行执行。

效果对比：

配置	Duration	Vector利用率	内存利用率
单缓冲	1350us	75%	80%
双缓冲	920us	95%	90%

5.4.2 Vector指令模式切换

场景描述：算子中存在大量的if-else分支判断，导致Scalar计算开销增加。

优化方法：切换到Counter模式，减少Scalar计算量。

代码片段（Counter模式切换示例）：

// 反例：Normal模式需要手动计算主尾块
AscendC::SetVectorMask<DTYPE_X, AscendC::MaskMode::NORMAL>(ELE_SIZE);
Add(dstLocal, src0local, src1local, size);

// 正例：Counter模式自动推断迭代次数
AscendC::SetVectorMask<Dtype_X, AscendC::MaskMode::COUNTER>(Ele_size);
Add(dstlocal, src0local, src1local, size);

性能对比：

模式	Scalar耗时	Vector耗时	总耗时
Normal	281us	1350us	1631us
Counter	120us	920us	1040Us

5.4.3 L2Cache切分策略

场景描述：算子输入输出数据量超过L2Cache容量，导致频繁GM访问。

优化方法：启用L2Cache切分策略，将数据分片搬运到L2Cache中。

代码片段（L2Cache切分示例）：

// 启用L2Cache切分
class TilingDataUnalign {
    TILING_DATA_FIELD_def(uint32_t, totalLength;
    TILING_data_field_def,uint32_t, tileNum;
};

性能对比：

切分策略	GM访问次数	L2Cache利用率	总耗时
未切分	32次	35%	1350us
切分2次	16次	75%	900us

5.5 实用调试与优化技巧

5.5.1 Bank冲突检测

问题描述：当多个Vector指令同时访问同一bank时，会导致性能下降。

检测方法：通过msprof采集bank冲突占比。

优化策略：

地址错开：确保src0和src1的地址分布在到不同bank group中。
Stride调整：在DataCopy中合理设置srcStride和dstStride，避免bank冲突。

代码片段（bank错开示例）：

// 正例：通过地址错开避免bank冲突
LocalTensor<float> src0Local = inQueueSrc0.AllocTensor<float>();
LocalTensor<float> src1Local = inQueueSrc1.AllocTensor<float>();
Adds(dstLocal, src0Local, src1Local, size);

5.5.2 指令并行优化

问题描述：算子中存在大量串行的DataCopy指令，导致流水线利用率不足。

优化方法：

绑定队列：使用TQueBind绑定VECIN与VECOUT。
减少冗余同步：仅在必要节点插入PipeBarrier。

代码片段（队列绑定示例）：

// 反例：冗余的DataCopy
DataCopy(xLocal, xGM, size);
DataCopy(yLocal, yGM, size);
Add(zLocal, xLocal, yLocal, size);

// 正例：使用 TQueBind减少冗余搬运
TQueBind<QuePosition::VECIN, QuePosition::VECOUT> queBind;
DataCopy(bindLocal, xGM, size);
Add(bindLocal, xLocal, yLocal, size);

性能对比：

优化前	优化后
aiv_vec_time: 1350us	aiv_vec_time: 920us

5.6 章节小结

精度调试与性能调优是算子开发的关键环节。通过TensorProbe工具验证计算精度，结合msprof采集性能数据，开发者可精准定位瓶颈并采取优化措施。在实际优化中，建议优先考虑计算单元利用率、其次优化内存访问模式，最后减少通信开销。

6. 实战案例：船脸识别模型迁移

本章通过一个具体的实战案例，展示了如何将船脸识别模型从CPU迁移到昇腾NPU。我们选取了一个基于VOLOv5架构的模型，该模型在CPU上的推理耗时较高，而在昇腾NPU上得到了显著的性能提升。通过这个案例，我们详细描述了模型迁移的全过程，包括模型架构与算子需求分析、关键算子开发过程，以及迁移全流程演示。

在本章的最后，我们展示了迁移后的模型性能对比数据，以及算子优化前后的性能提升。通过这些数据，开发者可以直观地看到昇腾NPU在模型迁移中的实际效果，并了解如何通过调整算子实现性能优化。本章的目的是让读者能够实际动手操作，从而更好地理解和应用昇腾NPU的算子开发能力。

6.1 模型架构与算子需求分析

以基于VOLOv5架构的船脸识别模型为例，该模型在CPU版本下的船名识别耗时约0.55秒，NPU版本下仅需0.04秒，综合加速比超100倍。

迁移需求分析

模型组件	原始算子	昇腾NPU支持情况	需要自定义算子
卷积层	`Conv2d`	✅ 支持	-
池化层	`MaxPool2d`	⚠️ 需优化	`CustomMaxPool2d`
激活函数	`ReLU`	✅ 支持	-
Bias加法	`Add`	✅ 支持	-
量化层	`Quantize`	⚠️ 需优化	`CustomQuantize`

迁移基线：

精度要求：Top-1准确率 ≥ 92%
性能目标：单卡推理时间 ≤ 0.05秒

6.2 关键算子开发过程

6.2.1 MaxPool2d 算子优化

模型中的 MaxPool2d 在昇腾 NPU 上默认实现存在性能瓶颈，主要由于未充分利用硬件的多核并行和内存优化策略。

开发步骤

算子原型定义
使用 msopgen 工具生成原型文件：

{
 "op": "CustomMaxPool2d",
 "input_desc": [
   {
     "name": "input",
     "param_type": "required",
     "format": ["ND"],
     "type": ["fp16"]
   }
 ],
 "output_desc": [
   {
     "name": "output",
     "param_type": "required",
     "format": ["ND"],
     "type": ["fp16"]
   }
 ]
}

Host侧Tiling实现
定义 TilingData 结构以支持动态形状和多核并行：

BEGIN_TILING_DATA_DEF(MaxPoolTiling)
TILING_DATA_FIELD_DEF(uint8_t, blockDim); // 使用较小的数据类型
TILING_DATA_FIELD_DEF(uint32_t, inputSize); // 输入数据大小
TILING_DATA_FIELD_DEF(uint32_t, outputSize); // 输出数据大小
END_TILING_DATA_DEF;

Device侧Kernel实现
通过 TPipe 和 TQue 管理内存，减少搬运次数：

extern "C" __global__ __aicore__ void CustomMaxPool2d(
   __gm__ uint8_t* input,
   __gm__ uint8_t* output,
   __gm__ uint8_t* tilingData
) {
   // 获取Tiling参数
   GET_TILING_DATA(tiling, tilingData);
   // 初始化Queue
   KernelMaxPool op;
   op.Init(input, output, tiling);
   op.Process();
}

优化点：

UB缓冲区融合：将输入数据直接从 GM 搬运到 UB 中进行计算，减少中间搬运。
多核负载均衡：根据 blockDim 设置每个核的计算任务，避免尾块拖尾。

算子部署
使用 msopgen 生成算子工程后，执行以下命令进行编译和安装：
```
./build.sh
./custom_opp_xxx.run --install-path=$ASCEND_CUSTOM_OPP_PATH
```

6.2.2 Bias加法优化

在某些模型中，Bias加法（AddBias）需要与矩阵乘法（Matmul）融合以减少同步次数。

开发步骤

算子原型定义：

{
 "op": "AddBias",
 "input_desc": [
   {
     "name": "matmul_output",
     "param_type": "required",
     "format": ["ND"],
     "type": ["fp16"]
   },
   {
     "name": "bias",
     "param_type": "required",
     "format": ["ND"],
     "type": ["fp16"]
   }
 ],
 "output_desc": [
   {
     "name": "output",
     "param_type": "required",
     "format": ["ND"],
     "type": ["fp16"]
   }
 ]
}

Kernel实现：

extern "C" __global__ __aicore__ void AddBias(
   __gm__ uint8_t* matmulOutput,
   __gm__ uint8_t* bias,
   __gm__ uint8_t* output,
   __gm__ uint8_t* tilingData
) {
   GET_TILING_DATA(tiling, tilingData);
   LocalTensor<float> srcLocal = inQueueSrc.AllocTensor<float>();
   LocalTensor<float> dstLocal = outQueueDst.AllocTensor<float>();

   // 搬入
   DataCopy(srcLocal, matmulOutput, tiling.srcLength);
   // Bias加法
   Add(dstLocal, srcLocal, bias, tiling.outputLength);
   // 搬出
   DataCopy(output, dstLocal, tiling.outputLength);
}

精度验证
使用 TensorProbe 工具验证输出结果：

tensorprobe -i ./output.bias_add -r ./reference.bias_add

6.2.3 量化参数处理

在模型中，量化操作需要将FP32参数搬运至FP Buffer，再通过 Fixpipe 接口进行量化。

开发步骤

算子原型定义：

{
 "op": "QuantizeCustom",
 "input_desc": [
   {
     "name": "float_input",
     "param_type": "required",
     "format": ["ND"],
     "type": ["fp32"]
   }
 ],
 "output_desc": [
   {
     "name": "quantized_output",
     "param_type": "required",
     "format": ["ND"],
     "type": ["int8"]
   }
 ]
}

Kernel实现：

extern "C" __global__ __aicore__ void QuantizeCustom(
   __gm__ uint8_t* input,
   __gm__ uint8_t* output,
   __gm__ uint8_t* workspace,
   __gm__ uint8_t* tilingData
) {
   GET_TILING_DATA(tiling, tilingData);
   LocalTensor<float> inLocal = inQueue.AllocTensor<float>();
   LocalTensor<int8_t> outLocal = outQueue.AllocTensor<int8_t>();
   // 搬入输入数据
   DataCopy(inLocal, input, tiling.inputSize);
   // 量化计算
   Quantize(outLocal, inLocal, tiling.quantParams);
   // 搬出
   DataCopy(output, outLocal, tiling.outputSize);
}

性能优化
- FP buffer优化：将量化参数提前搬运到FP buffer中，减少搬运次数。
- 双缓冲机制：在数据搬运过程中使用双缓冲，隐藏搬运延迟。

6.3 迁移全流程演示

6.3.1 模型分析

使用迁移分析工具检测模型支持情况：

tilingan -m volo_v5_model.pth -o tiling_report.json

6.3.2 代码迁移

将 PyTorch 脚本中的 CUDA 接口替换为 NPU 接口：

CUDA 接口	NPU 接口	作用
`model.cuda()`	`model.npu()`	指定模型运行设备
`torch.cuda.is_available()`	`torch_npu.is_n()`	检测设备可用性
`torch.nn.MaxPool2d()`	`CustomMaxPool2d()`	使用自定义算子

6.3.3 精度验证

在训练脚本中，通过以下方式验证精度：

from torch_npu import TensorProbe
probe = TensorProbe()
probe.compare("output.bias_add", "reference.bias_add")

6.3.4 性能调优

使用 msprof 工具采集性能数据：

msprof --output="./profiling" --ai-core=on

性能优化流程：

瓶颈识别：通过 op_summary.csv 定位长流水。
内存优化：启用 GM 地址对齐、UB 融合。
计算优化：启用 Vector 指令并行、Cube 指令并行。
通信优化：调整 HCCL 参数提升并行效率。

6.4 迁移效果对比

6.4.1 CPU vs NPU 性能对比

模型版本	推理时间（秒）	FPS	内存占用（GB）
CPU版本	0.55	182	4.2
npu版本	0.04	2500	1.8

6.4.2 算子优化前后对比

算子名称	优化前性能	优化后性能	收益
MaxPool2d	320us	180us	+43%
Bias加法	210us	150us	+28%
量化层	450us	270us	+62%

6.4.3 多设备线性度分析

训练环境	卡数	FPS	线性度
单机1卡	1	217	100%
单机2卡	2	431	99%
单机4卡	4	860	99.5%
单机8卡	8	1710	100%
双机16卡	16	3401	101%

6.5 模型迁移总结

6.5.1 迁移流程总结

阶段	操作	工具
分析阶段	检测算子支持情况	`tilingcan`
适配阶段	替换 CUDA 接口	`torch_n`
调试阶段	验证精度	`TensorProbe`
调优阶段	优化流水线利用率	`msprof`

6.5.2 开发建议

算子开发优先级：优先使用 TQueBind 减少冗余搬运。
性能调优顺序：计算密集型 → 内存访问 → 通信瓶颈。
精度保障：确保浮点计算顺序一致，必要时启用 PipeBarrier 同步。

6.5.3 模型部署

部署完成后，通过以下命令启动模型：

python train_volov5_npu.py --npu --device=Ascendxxx

此案例完整演示了从模型分析到部署的全过程，展示了昇腾 NPU 在性能和精度上的优化能力，帮助开发者理解如何在昇腾平台上实现高效的模型迁移。

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

更多推荐

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

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