CANN ops-nn LpNormV3 算子全解析：Lp 范数计算与归一化在 NPU 上的实现及 aclnn 调用实践

CANN ops-nn LpNormV3 算子全解析：Lp 范数计算与归一化在 NPU 上的实现及 aclnn 调用实践

【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库，实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn

LpNormV3 是 CANN ops-nn 开源算子库（experimental/norm/lp_norm_v3）中提供的 Lp 范数计算与归一化算子：它按指定维度计算输入张量的 p 范数，并用该范数对输入逐元素归一化输出。本文以该模块的 README.md 为主线，结合算子定义、InferShape、Tiling、AscendC Kernel 与 aclnn 调用样例等源码，系统讲解其数学原理、参数语义、NPU 多核实现机制，并给出可复现的 FP32/FP16 调用与精度验证实践，帮助开发者在 Atlas 训练/推理产品上快速完成 LpNormV3 的集成与验证。

一、算子概述与产品支持情况

1.1 功能定位

LpNormV3 同时完成两项任务：

• Lp 范数计算：对输入张量按指定维度（axis）计算 p 范数；
• 归一化：将输入张量的每个元素除以对应位置的范数，输出与输入同形状的归一化结果。

该能力在深度学习中常用于特征归一化（如 L2 归一化）、权重/梯度正则化以及数值稳定性处理等场景。按 README 的描述，算子支持全局、按行、按列三种计算模式，这一"三种模式"的划分与源码中三个 Tiling Key 一一对应（详见下文）。

1.2 产品支持情况

README 明确列出的产品支持范围如下：

产品	是否支持
Atlas A2 训练系列产品 / Atlas 800I A2 推理产品	√

该信息与算子定义文件 lp_norm_v3_def.cpp 中this->AICore().AddConfig("ascend910b")的配置相互印证——算子面向 910B 系列 AICore 构建。需要说明的是，当前仓库代码仅覆盖上表所列产品，其他产品型号的支持情况以 CANN 官方发布为准。

二、数学原理与计算公式

设输入张量为 (x)，范数阶数为 (p)，数值稳定项为 (\epsilon)，范数的计算范围由axis参数决定：

Lp 范数计算：

[ \text{norm} = \left( \sum |x_i|^p + \epsilon \right)^{1/p} ]

归一化结果：

[ y_i = \frac{x_i}{\text{norm}} ]

2.1 源码对公式的实现印证

Kernel 中严格按上述两步流水实现，两阶段处理流程与公式一一对应：

1. Reduce 阶段（求 (\sum |x_i|^p)）：在 lp_norm_v3.h 的Reduce中，先对元素取绝对值（Abs），再做 p 次幂（Power），最后按轴语义累加进局部和localSum；
2. 开方阶段：在SumAndSyncAll中，核 0 对全局和加上 epsilon 后，利用Ln → Muls(1/p) → Exp的等价变换计算 ((sum+\epsilon)^{1/p})（见 lp_norm_v3.h），避免直接调用高开销的开方/幂函数；
3. Normalization 阶段：逐元素用Div（或标量除法）除以对应范数槽位，得到归一化输出（见 lp_norm_v3.h）。

2.2 无穷范数（L∞）的特殊处理

当p取正无穷或负无穷时，公式退化为取绝对值后的最大值/最小值，不再适用"求和开方"路径。Kernel 通过判断p的比特位（0x7F800000/0xFF800000）识别无穷情形（见 lp_norm_v3.h），在Process中分流到InfProcess/ReduceInf：用逐元素 max/min 替代求和，并以原子 max/min 完成跨核归约（见 lp_norm_v3.h）。调用样例中也预留了p = std::numeric_limits<float>::infinity()的写法（见 test_aclnn_lp_norm_v3.cpp）。

2.3 数值稳定性

README 将 (\epsilon) 定位为"数值稳定补偿项"，用于避免范数过小时除法溢出。从源码看，当前实现中 epsilon 是 Kernel 内硬编码常量1e-6（见 lp_norm_v3.h），并非运行时属性——算子定义仅注册了p与axis两个属性（见下文参数说明），这一点在对接时需与 README 参数表的口径区分开。

三、参数说明

README 给出的参数表如下：

参数名	输入/输出/属性	描述	数据类型	数据格式
x	输入	待进行 LpNormV3 计算的输入张量	FLOAT、FLOAT16	ND
p	属性	Lp 范数的阶数	FLOAT	-
axis	属性	范数计算维度：0（全局），1（列），2（行）	INT	-
epsilon	属性	数值稳定补偿项	FLOAT	-
y	输出	归一化后的输出张量	FLOAT、FLOAT16	ND

3.1 参数在源码中的真实语义（重要）

结合算子定义与 Tiling/Kernel 实现，各参数在仓库代码中的实际约定如下，与 README 表格存在两处需要开发者注意的差异：

• axis取值：算子定义中axis为可选 INT 属性，默认值-1（见 lp_norm_v3_def.cpp）。Tiling 阶段依据axis生成三个 Tiling Key（见 lp_norm_v3_tiling.cpp 与 lp_norm_v3_tiling_key.h）：

  axis 实际取值	Tiling Key	语义（按源码）	workspace 槽位数
  -1	LP_NORM_AXIS_NONE	全局规约（整个张量一个范数）	1
  0	LP_NORM_AXIS_0	按行计算（每行一个范数）	rows
  1	LP_NORM_AXIS_1	按列计算（每列一个范数）	cols

  即源码实际使用-1 / 0 / 1三值，与 README 表中"0（全局），1（列），2（行）"的文字描述并不一致；调用样例 test_aclnn_lp_norm_v3.cpp 以DEFAULT_AXIS = 0表示按行归一化，同样印证了源码语义。开发者在写图或调用 aclnn 接口时应以源码为准：全局传-1，按行传0，按列传1。

• epsilon属性：README 参数表将其列为属性，但算子定义只注册了p与axis（见 lp_norm_v3_def.cpp），aclnn 接口签名aclnnLpNormV3GetWorkspaceSize(input, p, axis, output, ...)也没有 epsilon 入参。因此当前版本中 epsilon 为 Kernel 内部常量（1e-6），未向用户开放配置。

• p的默认值：p为可选 FLOAT 属性，默认2.0f（即默认做 L2 范数归一化，见 lp_norm_v3_def.cpp），调用样例同样使用DEFAULT_P = 2.0f。

• 输出形状：y与输入x同形状。InferShape 实现直接将输入各维度拷贝给输出（见 lp_norm_v3_infershape.cpp），这也与"逐元素归一化、形状不变"的语义一致。

• 输入维度限制：从 Tiling 源码看，rows = shape.dim(0)、cols = shape.dim(1)（见 lp_norm_v3_tiling.cpp），可以推断当前实现主要面向二维张量（按行/按列模式），全局模式同样通过二维形状的rows*cols展开处理；调用样例也明确限定"Only support 2D input"（见 test_aclnn_lp_norm_v3.cpp）。

四、源码结构总览

该算子模块按 CANN 算子标准布局组织，各目录职责清晰：

experimental/norm/lp_norm_v3/ ├── README.md # 算子说明文档（本文主体） ├── CMakeLists.txt # 模块构建入口，遍历子目录 ├── examples/ │ └── test_aclnn_lp_norm_v3.cpp # aclnn 接口调用样例（FP32/FP16） ├── op_host/ │ ├── CMakeLists.txt │ ├── lp_norm_v3_def.cpp # 算子原型定义（输入/输出/属性/平台） │ ├── lp_norm_v3_infershape.cpp # InferShape：输出形状推导 │ └── lp_norm_v3_tiling.cpp # Tiling：分块、多核调度、workspace 计算 ├── op_kernel/ │ ├── lp_norm_v3.cpp # Kernel 入口（模板实例化） │ ├── lp_norm_v3.h # Kernel 核心实现（Reduce/Normalize） │ ├── lp_norm_v3_tiling_data.h # Tiling 数据结构 │ └── lp_norm_v3_tiling_key.h # Tiling Key（三种轴模式） └── tests/ └── ut/ # 单测目录（当前为空占位）

模块 CMakeLists.txt 通过file(GLOB ...)枚举子目录并add_subdirectory，且仅在ENABLE_TEST或BENCHMARK开启时才包含tests目录，遵循仓库统一的算子模块组织方式。

五、算子定义与 InferShape：数据契约

5.1 算子原型定义

lp_norm_v3_def.cpp 通过OpDef注册算子元信息，构成数据契约：

• 输入x：REQUIRED（必选），支持DT_FLOAT、DT_FLOAT16，格式为FORMAT_ND，并对未确定 shape 的场景声明了UnknownShapeFormat，同时调用AutoContiguous()保证输入内存连续化；
• 输出y：同样必选，数据类型与格式与输入对齐；
• 属性p：OPTIONAL，FLOAT，默认2.0f；
• 属性axis：OPTIONAL，INT，默认-1；
• 平台配置：AICore().AddConfig("ascend910b")，与 README 产品支持表呼应。

OP_ADD(LpNormV3)将该算子注册进算子信息库，供框架侧构图与下发使用。

5.2 InferShape

lp_norm_v3_infershape.cpp 中InferShapeLpNormV3的逻辑非常简洁：读取输入x的形状，将xShapeSize与各维尺寸原样写入输出y。由于归一化不改变张量形状，这一推导是准确且高效的。

六、Tiling 设计：多核分块、workspace 与调度模式

Tiling 在 lp_norm_v3_tiling.cpp 中完成，是连接 Host 侧与 Device 侧 Kernel 的桥梁，核心要点如下：

6.1 平台信息与 UB 容量

GetPlatformInfo通过PlatformAscendC获取片上UB（Unified Buffer）大小与可用核数（lp_norm_v3_tiling.cpp）。单核单次搬运的数据量tileDataNum依据UB_SIZE / BUFFER_NUM / blockSize / 10计算（lp_norm_v3_tiling.cpp），其中BUFFER_NUM = 2对应 Kernel 侧的双缓冲队列设计。

6.2 多核负载均衡

CalculateCoreBlockNums与LpNormV3TilingFunc共同决定核数及每个核处理的数据块：

• 若单 tile 即可容纳全部输入（tileDataNum >= inputNum），则退化为单核执行；
• 否则在"可用核数"与"对齐后数据块数"之间取较小值，保证每个核至少分到 32B 数据（lp_norm_v3_tiling.cpp）；
• 为应对数据量不能被核数整除的情况，Tiling 数据中同时输出smallCoreDataNum / bigCoreDataNum、smallTailDataNum / bigTailDataNum、finalSmallTileNum / finalBigTileNum等成对参数，让前tailBlockNum个核处理"大核"数据、其余核处理"小核"数据，实现负载均衡。这些字段的结构定义见 lp_norm_v3_tiling_data.h。

6.3 Workspace 分配策略

GetWorkspaceSize依据轴模式决定 workspace 大小（lp_norm_v3_tiling.cpp）：

• 全局模式：1 个范数槽位；
• 按行模式：rows个槽位；
• 按列模式：cols个槽位。

每个范数以 float 存储并做 64B 对齐（SLOT_STRIDE = 64 / sizeof(float) = 16，即每个槽位 16 个 float），注释明确说明：多核并发写全局缓存时，若多个核在同一个 64B 内同时操作会导致随机覆写，因此按"范数数量 × 16 float"分配（lp_norm_v3_tiling.cpp）。workspace 总大小为"用户槽位 + 系统 workspace（GetLibApiWorkSpaceSize）"。

6.4 调度模式与 Tiling Key

由于 LpNormV3 需要跨核归约（先求和、再广播范数），Tiling 通过context->SetScheduleMode(1)声明为核间同步算子（lp_norm_v3_tiling.cpp），并依据axis设置LP_NORM_AXIS_NONE / LP_NORM_AXIS_0 / LP_NORM_AXIS_1三个编译模板参数（Tiling Key，见 lp_norm_v3_tiling_key.h），从而在编译期展开不同的轴处理分支，避免运行时分支开销。

七、Kernel 实现原理：两阶段流水 + 原子归约

Kernel 入口 lp_norm_v3.cpp 以__global__ __aicore__模板函数实例化NsLpNormV3::LpNormV3<DTYPE_X, schMode>，核心实现在 lp_norm_v3.h。

7.1 整体流程

Process根据p是否为无穷选择NormalProcess或InfProcess，两者结构一致，均分两个阶段：

1. 规约阶段：循环CopyIn → Reduce（或ReduceInf），将输入分 tile 搬入 UB，计算|x|^p（或 max/min）累加进各核本地localSum；
2. 同步与广播阶段：SumAndSyncAll（或SumAndSyncAllInf）通过 workspace + 原子操作完成跨核归约，核 0 计算最终范数并写回 workspace，随后SyncAll同步所有核；
3. 归一化阶段：各核再次CopyIn → Normalization → CopyOut，按轴语义读取对应范数槽位，逐元素相除得到输出。

7.2 轴语义在 Kernel 中的映射

Reduce与Normalization中，通过全局线性索引反推行/列（lp_norm_v3.h）：

• 按行（AXIS_0）：rowIdx = globalIdx / cols，行内元素累加到同一行范数；
• 按列（AXIS_1）：colIdx = globalIdx % cols，列内元素累加到同一列范数；
• 全局（AXIS_NONE）：所有元素累加为单一标量。

归一化阶段反向使用同样的索引取范数并相除（lp_norm_v3.h）。

7.3 跨核归约的可靠性设计

• 原子加：各核通过SetAtomicAdd<float>()+DataCopy将本地和累加到 workspace 槽位，完成后SetAtomicNone()恢复（lp_norm_v3.h）；无穷范数路径则使用SetAtomicMax/SetAtomicMin；
• 缓存一致性：核 0 在写回范数后调用DataCacheCleanAndInvalid<...>(workGm[...])逐槽位刷缓存，配合SyncAll与PipeBarrier<PIPE_V>保证其他核读到的是最新范数值（lp_norm_v3.h）；
• 数据类型适配：FP16 输入在Reduce中先Cast到 float 再计算，规避半精度累加误差（lp_norm_v3.h）。

八、aclnn 调用实践：从样例代码到精度验证

README 指明算子通过 aclnn 接口调用，对应样例为 test_aclnn_lp_norm_v3.cpp。该样例完整演示了 ACL 环境初始化 → 数据准备 → 张量创建 → workspace 申请 → 算子执行 → 结果回拷 → 精度验证的完整闭环，可直接作为集成参考。

8.1 环境初始化与张量创建

// ACL 初始化：aclInit -> aclrtSetDevice -> aclrtCreateStream int Init(int32_t deviceId, aclrtStream* stream); // 创建输入/输出张量：申请 Device 内存 + Host->Device 拷贝 + aclCreateTensor template <typename T> int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType, aclTensor** tensor) { auto ret = aclrtMalloc(deviceAddr, elemNum * sizeof(T), ACL_MEM_MALLOC_HUGE_FIRST); ret = aclrtMemcpy(*deviceAddr, memSize, hostData.data(), memSize, ACL_MEMCPY_HOST_TO_DEVICE); // 计算连续张量 strides，shape 为 ND 格式 *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, ACL_FORMAT_ND, shape.data(), shape.size(), *deviceAddr); }

样例以固定随机种子（12345）在 [1, 5] 区间生成输入，输入形状为{2, 3}，规避范数为 0 的退化场景，便于精度比对（test_aclnn_lp_norm_v3.cpp）。

8.2 算子执行四步曲

// 1. 获取 workspace 大小与 executor（注意接口入参为 p 与 axis） aclnnLpNormV3GetWorkspaceSize(inputTensor, DEFAULT_P /*2.0f*/, DEFAULT_AXIS /*0*/, outputTensor, &workspaceSize, &executor); // 2. 申请 workspace（可能为 0） if (workspaceSize > 0) { aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); } // 3. 执行算子 aclnnLpNormV3(workspaceAddr, workspaceSize, executor, stream); // 4. 流同步，确保算子完成 aclrtSynchronizeStream(stream);

执行完毕后通过aclrtMemcpy(..., ACL_MEMCPY_DEVICE_TO_HOST)将结果回拷 Host，FP32 直接以 float 打印，FP16 需经aclFloat16ToFloat转回 float 展示（test_aclnn_lp_norm_v3.cpp）。

8.3 Golden 计算与精度验证思路

样例的ComputeGoldenData在 Host 侧按公式手算范数与归一化结果作为基准（golden），支持有限 p 与无穷 p 两种分支（test_aclnn_lp_norm_v3.cpp）。VerifyResult逐元素比较算子输出与 golden：

• 统计Max Error / Avg Error；
• 阈值：FP32 为1e-5，FP16 为1e-3（test_aclnn_lp_norm_v3.cpp）；
• 超差时打印前 8 个失配元素的下标、输出值、golden 值与误差，便于定位。

主函数默认执行 FP32 用例；FP16 用例通过ENABLE_FP16_TEST开关启用（默认关闭，见 test_aclnn_lp_norm_v3.cpp）。FP16 路径的 golden 以 FP32 精度计算后与半精度输出比对，从而把误差来源收敛到半精度表示与累加精度上。

8.4 构建与运行

该模块通过仓库顶层 CMake 统一构建（experimental 目录下算子模块均以子目录方式挂接），模块 CMakeLists.txt 负责遍历并add_subdirectory各子目录；样例代码需在具备 CANN 工具链（含 acl/acl.h 与生成的 aclnn_lp_norm_v3.h 接口头文件）与对应 NPU 设备（README 所列 Atlas A2 系列）的环境下编译运行。测试目录tests/ut当前为占位状态，可在开启ENABLE_TEST后按仓库 tests/ut 的既有范式补充单测。

九、从 README 到源码的几点实践提示

1. axis 语义以源码为准：全局传-1、按行传0、按列传1，与 README 参数表的文字描述存在差异，构图或调用前务必确认；
2. epsilon 当前不可配置：README 列为属性，但实际为 Kernel 内部常量1e-6，若业务需要其他补偿值，需关注后续版本是否开放该属性；
3. 输入维度：当前实现按二维（rows/cols）组织归约，高维张量场景需先确认算子是否支持或自行 reshape；
4. 无穷范数支持：p取 ±inf 时走独立的 max/min 路径，样例中已预留该配置方式；
5. 产品边界：当前仅支持 Atlas A2 训练系列 / Atlas 800I A2 推理产品，对应ascend910bAICore 配置。

十、贡献信息

按 README 记录，LpNormV3 由个人开发者 shixiangyang 于 2025/11/8 适配贡献至开源仓（贡献方：个人开发者），代码版权归哈尔滨工业大学 AISS Group（OpenBOAT 项目）所有，并遵循 CANN Open Software License Agreement Version 2.0 开源许可（见仓库根目录 LICENSE）。

【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库，实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn