CANN ops-nn 中 LpLoss（L1Loss）算子深度解析：从 aclnnL1Loss 接口到 NPU 内核实现

CANN ops-nn 中 LpLoss（L1Loss）算子深度解析：从 aclnnL1Loss 接口到 NPU 内核实现

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

本文聚焦 CANN（Compute Architecture for Neural Networks）神经网络算子库 ops-nn 中的 LpLoss 算子，系统讲解其在 Ascend NPU 上的功能定义、参数语义、约束条件、aclnn 两段式接口调用方式，并结合仓库源码深入剖析其从 Host 侧参数校验、图融合到 Device 侧内核调度的完整实现链路。读完本文，你将掌握如何在 CANN 环境下手写代码调用aclnnL1Loss完成 L1 损失计算，并理解该算子在 ops-nn 仓库中的源码结构与底层工作原理。

概述：LpLoss 与 L1Loss 的关系

LpLoss 是神经网络训练中常用的损失函数族，其一般形式为 $l_p = \left(\sum|x-y|^p\right)^{1/p}$。在 CANN ops-nn 仓库中，LpLoss 算子仅支持p=1的场景，此时即为经典的 L1 损失（L1Loss），对外暴露的 aclnn 接口名称为aclnnL1Loss。这一点在 loss/lp_loss/README.md 与 loss/lp_loss/docs/aclnnL1Loss.md 的约束说明中均有明确表述，且与 PyTorch 框架中的torch.nn.L1Loss语义对齐——loss/lp_loss/op_graph/lp_loss_proto.h 的算子原型注释中明确标注"Compatible with the Pytorch operator LpLoss"。

从仓库目录结构看，该算子是一个完整的三层实现工程，包含：

• op_api：对外暴露的 aclnn 接口层（aclnn_l1_loss.cpp/aclnn_l1_loss.h）；
• op_host：算子原型注册（lp_loss_def.cpp）、shape 推导（lp_loss_infershape.cpp）与 tiling 计算（op_host/arch35/lp_loss_tiling_arch35.cpp）；
• op_kernel：Device 侧内核实现（op_kernel/lp_loss.cpp及op_kernel/arch35/下的 DAG 计算图描述）；
• op_graph：算子原型定义（lp_loss_proto.h）；
• tests：包含 op_api 与 op_host 的单元测试、ST（系统测试）用例以及 golden 数据生成脚本。

产品支持情况

LpLoss（aclnnL1Loss）算子在不同硬件产品上的支持情况如下表所示（来源于 loss/lp_loss/README.md）：

从源码角度看，这一支持情况与 loss/lp_loss/op_host/lp_loss_def.cpp 中注册的 AICore 配置（AddConfig("ascend950", aicoreConfig)）以及 loss/lp_loss/op_api/aclnn_l1_loss.cpp 中按 SoC 版本区分的数据类型支持列表（CheckSocVersionIsSupportBf16）相互印证。

功能说明与计算公式

算子功能

LpLoss 算子计算输入self与目标target中每个元素之间的平均绝对误差（Mean Absolute Error，MAE）。reduction属性指定应用到输出的缩减方式，支持'none'、'mean'、'sum'三种：

• 'none'：不应用缩减，逐元素输出|x - y|；
• 'mean'：输出总和除以输出中的元素数；
• 'sum'：输出被求和。

计算公式

当reduction为none时：

$$ \ell(x, y) = L = {l_1,\dots,l_N}^\top, \quad l_n = \left| x_n - y_n \right|, $$

其中 $x$ 是self，$y$ 是target，$N$ 是 batch 的大小。如果reduction不是'none'，那么：

$$ \ell(x, y) = \begin{cases} \operatorname{mean}(L), & \text{if reduction} = \text{'mean';}\ \operatorname{sum}(L), & \text{if reduction} = \text{'sum'.} \end{cases} $$

源码中的公式落地

在内核实现中，这一公式被拆解为标准的向量算子 DAG（有向无环计算图）。以 loss/lp_loss/op_kernel/arch35/lp_loss_dag.h 为例，三种 reduction 模式分别对应三个计算图模板：

• LpLossOp（reduction=none）：CopyIn → Cast → Sub（相减）→ Abs（取绝对值）→ Cast → CopyOut，逐元素输出|x-y|；
• LpLossSumDag（reduction=sum）：在Sub → Abs之后追加ReduceSumOp完成全量求和；
• LpLossMeanDag（reduction=mean）：在求和之后追加Muls（乘以均值倒数meanVar），其中meanVar由 tiling 阶段计算并注入。

而 loss/lp_loss/op_kernel/lp_loss.cpp 中的内核入口函数lp_loss则根据Reduction模板参数在编译期选择ElementwiseSch（逐元素调度）或ReduceSch（归约调度）完成计算，mean模式下还通过op.Process(static_cast<DTYPE_PREDICT>(NAN))处理空输入时输出 NAN 的语义。

参数说明

self（计算输入）

• 类型：aclTensor*，Device 侧的 aclTensor；
• 数据类型：与target满足数据类型推导（promote）规则（参见互推导关系）；
• shape：支持 0-8 维，且需要与target满足 broadcast 规则；
• 其他：支持非连续的 Tensor，数据格式 支持 ND；
• 各产品数据类型支持：Atlas A2 训练/推理系列、Ascend 950PR/950DT、Atlas A3 训练/推理系列支持 BFLOAT16、FLOAT16、FLOAT32、INT64。

target（计算输入）

• 类型：aclTensor*，Device 侧的 aclTensor；
• 数据类型：与self满足数据类型推导规则；
• shape：支持 0-8 维，需要与self满足 broadcast 规则；
• 其他：支持非连续 Tensor，数据格式支持 ND；
• 各产品数据类型支持：同上，支持 BFLOAT16、FLOAT16、FLOAT32、INT64。

reduction（属性）

• 类型：int64_t，Host 侧的整型属性；
• 取值：0('none') | 1('mean') | 2('sum')，其中'none'表示不缩减，'mean'表示输出总和除以元素数，'sum'表示输出求和。

out（计算输出）

• 类型：aclTensor*，Device 侧的 aclTensor；
• 数据类型：需要是self与target推导之后可转换的数据类型（参见互转换关系）；
• shape 规则：当reduction为 0 时，out的 shape 与self和targetbroadcast 后的 shape 一致；当reduction不为 0 时，out为 0 维 tensor；
• 其他：支持非连续 Tensor，数据格式支持 ND；
• 各产品数据类型支持：Atlas A2 训练/推理系列、Ascend 950PR/950DT、Atlas A3 训练/推理系列支持 BFLOAT16、FLOAT16、FLOAT32、INT64、COMPLEX64、COMPLEX128。

参数校验的源码实现

loss/lp_loss/op_api/aclnn_l1_loss.cpp 中的CheckParams函数完整实现了上述参数的逐项校验，包括：

1. 空指针检查（CheckNotNull3Tensor）；
2. reduction 取值范围检查（CheckReduction，仅接受 0/1/2）；
3. 数据类型推导与支持范围检查（CheckDtypeValid）；
4. shape 与 broadcast 规则检查（CheckShape，维度上限 8）。

其中CheckDtypeValid中还包含两条与 CUDA 行为保持一致的特殊校验：

• 当reduction='none'时，若self不是浮点类型，则target也不能是浮点类型；
• 当reduction='mean'时，self和target至少有一个是浮点类型（因为求均值需要除法）。

另外，CheckFormat会对FORMAT_FRACTAL_NZ格式给出精度告警日志，提示该格式可能导致精度问题。

约束说明

• 确定性计算：aclnnL1Loss默认确定性实现，同一输入多次执行结果一致；
• p 值限制：LpLoss 中p为计算 loss 的参数，当前只支持p=1，对外接口名称为aclnnL1Loss。这一定义同时体现在算子原型 loss/lp_loss/op_graph/lp_loss_proto.h 的REQUIRED_ATTR(p, Int)与 loss/lp_loss/op_host/lp_loss_def.cpp 的Attr("p").AttrType(REQUIRED).Int(1)中；
• 空输入语义：当self或target为空 tensor 时（IsEmpty()），loss/lp_loss/op_api/aclnn_l1_loss.cpp 的L1LossEmptyTensorCompute会按语义直接处理：reduction=none 返回空 tensor，reduction=mean 填充 NAN，reduction=sum 填充 0，不再进入内核计算。

调用说明：aclnn 两段式接口

LpLoss（aclnnL1Loss）属于标准的 CANN aclnn 两段式接口（参见两段式接口说明）：必须先调用aclnnL1LossGetWorkspaceSize获取计算所需 workspace 大小以及包含算子计算流程的执行器，再调用aclnnL1Loss执行计算。

调用方式汇总如下（完整样例见 loss/lp_loss/examples/test_aclnn_l1_loss.cpp）：

第一段接口：aclnnL1LossGetWorkspaceSize

aclnnStatus aclnnL1LossGetWorkspaceSize( const aclTensor* self, const aclTensor* target, int64_t reduction, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)

第一段接口完成入参校验，并返回 workspace 大小与执行器。参数细节如下：

该接口返回aclnnStatus状态码（具体参见 aclnn 返回码说明），并在以下场景报错：

第二段接口：aclnnL1Loss

aclnnStatus aclnnL1Loss( void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)

接口头文件与声明

接口声明位于 loss/lp_loss/op_api/aclnn_l1_loss.h，属于aclnn_ops_train域。样例代码通过#include "aclnnop/aclnn_l1_loss.h"引入该接口。仓库中同时保留了通用的lp_loss低阶算子封装（loss/lp_loss/op_api/lp_loss.cpp 与lp_loss.h），供框架层组合调用。

完整调用示例与执行流程

以下示例改编自仓库样例 loss/lp_loss/examples/test_aclnn_l1_loss.cpp（完整的可编译版本请直接参考该文件），展示了从环境初始化到结果回拷的完整调用流程。示例中self = {0, 1, 2, 3}，target = {1, 1, 1, 1}，reduction = 1（mean），因此预期输出为(|0-1| + |1-1| + |2-1| + |3-1|) / 4 = 1.0。

#include <iostream> #include <vector> #include "acl/acl.h" #include "aclnnop/aclnn_l1_loss.h" #define CHECK_RET(cond, return_expr) \ do { \ if (!(cond)) { \ return_expr; \ } \ } while (0) #define LOG_PRINT(message, ...) \ do { \ printf(message, ##__VA_ARGS__); \ } while (0) int64_t GetShapeSize(const std::vector<int64_t>& shape) { int64_t shapeSize = 1; for (auto i : shape) { shapeSize *= i; } return shapeSize; } int Init(int32_t deviceId, aclrtStream* stream) { // 固定写法，资源初始化 auto ret = aclInit(nullptr); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); ret = aclrtSetDevice(deviceId); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); ret = aclrtCreateStream(stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); return 0; } template <typename T> int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType, aclTensor** tensor) { auto size = GetShapeSize(shape) * sizeof(T); // 调用aclrtMalloc申请device侧内存 auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret); // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上 ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret); // 计算连续tensor的strides std::vector<int64_t> strides(shape.size(), 1); for (int64_t i = shape.size() - 2; i >= 0; i--) { strides[i] = shape[i + 1] * strides[i + 1]; } // 调用aclCreateTensor接口创建aclTensor *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), *deviceAddr); return 0; } int main() { // 1.（固定写法）device/stream初始化，参考acl API手册 int32_t deviceId = 0; // 根据自己的实际device填写deviceId aclrtStream stream; auto ret = Init(deviceId, &stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret); // 2. 构造输入与输出 std::vector<int64_t> selfShape = {2, 2}; std::vector<int64_t> targetShape = {2, 2}; std::vector<int64_t> outShape = {}; // reduction=1(mean) 时 out 为0维 void* selfDeviceAddr = nullptr; void* targetDeviceAddr = nullptr; void* outDeviceAddr = nullptr; aclTensor* self = nullptr; aclTensor* target = nullptr; aclTensor* out = nullptr; std::vector<float> selfHostData = {0, 1, 2, 3}; std::vector<float> targetHostData = {1, 1, 1, 1}; std::vector<float> outHostData = {0}; ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self); CHECK_RET(ret == ACL_SUCCESS, return ret); ret = CreateAclTensor(targetHostData, targetShape, &targetDeviceAddr, aclDataType::ACL_FLOAT, &target); CHECK_RET(ret == ACL_SUCCESS, return ret); ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out); CHECK_RET(ret == ACL_SUCCESS, return ret); int64_t reduction = 1; // 0('none') | 1('mean') | 2('sum') // 3. 调用CANN算子库API（两段式） uint64_t workspaceSize = 0; aclOpExecutor* executor; // 调用第一段接口，完成参数校验并获取workspace大小与执行器 ret = aclnnL1LossGetWorkspaceSize(self, target, reduction, out, &workspaceSize, &executor); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnL1LossGetWorkspaceSize failed. ERROR: %d\n", ret); return ret); // 根据第一段接口计算出的workspaceSize申请device内存 void* workspaceAddr = nullptr; if (workspaceSize > 0) { ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret); } // 调用第二段接口执行计算 ret = aclnnL1Loss(workspaceAddr, workspaceSize, executor, stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnL1Loss failed. ERROR: %d\n", ret); return ret); // 4.（固定写法）同步等待任务执行结束 ret = aclrtSynchronizeStream(stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧 auto size = GetShapeSize(outShape); std::vector<float> resultData(size, 0); ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret); for (int64_t i = 0; i < size; i++) { LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]); } // 6. 释放aclTensor aclDestroyTensor(self); aclDestroyTensor(target); aclDestroyTensor(out); // 7. 释放device资源 aclrtFree(selfDeviceAddr); aclrtFree(targetDeviceAddr); aclrtFree(outDeviceAddr); if (workspaceSize > 0) { aclrtFree(workspaceAddr); } aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0; }

具体编译与执行过程请参考仓库的编译与运行样例指南。仓库中还提供了另一个调用示例 loss/lp_loss/examples/test_aclnn_lp_loss.cpp，以及单元测试 loss/lp_loss/tests/ut/op_api/test_aclnn_l1_loss.cpp 供参考。

源码实现深度解析

Host 侧：算子原型与 shape 推导

算子的计算图原型定义在 loss/lp_loss/op_graph/lp_loss_proto.h，通过REG_OP(LpLoss)声明两个输入predict、label（支持 DT_FLOAT16、DT_FLOAT、DT_BF16），一个输出y，一个必选整型属性p和一个默认为"mean"的字符串属性reduction。

算子原型注册位于 loss/lp_loss/op_host/lp_loss_def.cpp，其中还声明了 AICore 运行配置：支持动态编译（DynamicCompileStaticFlag(true)）、动态 rank（DynamicRankSupportFlag(true)）、动态 shape（DynamicShapeSupportFlag(true)），且关闭精度缩减（PrecisionReduceFlag(false)）。

Shape 推导逻辑位于 loss/lp_loss/op_host/lp_loss_infershape.cpp 的InferShape4LpLoss：当reduction == "none"时，输出 shape 与输入 shape 一致；否则输出为 0 维标量。这与 README 中 out 的 shape 约束完全对应。

aclnn 接口层的组合逻辑

loss/lp_loss/op_api/aclnn_l1_loss.cpp 的aclnnL1LossGetWorkspaceSize是整个算子执行流程编排的核心，其逻辑依次为：

1. 参数校验：调用CheckParams完成空指针、reduction、dtype、shape 的逐项校验；
2. 空 tensor 快速路径：若self或target为空，直接按语义填充（见前述约束说明）；
3. 数据类型对齐：通过op::PromoteType计算推导类型，并调用l0op::Cast将self、target统一 cast 到推导类型；
4. 连续性处理：通过l0op::Contiguous将输入转换为连续 tensor，以兼容非连续输入；
5. Broadcast：当self与targetshape 不一致时，调用l0op::BroadcastTo将两者对齐到同一 shape；
6. 核心计算分发：若推导类型为 INT64，则走小算子拼接路径GetL1LossFromInt64（Sub → Abs → ReduceSumOp，因为 LpLoss 的 AiCore 内核不支持 INT64）；否则调用l0op::LpLoss直接走 AiCore 内核；
7. 输出处理：将计算结果Cast到 out 的 dtype，再通过l0op::ViewCopy写入 out（兼容非连续 out）；
8. 返回 workspace 大小：通过executor->GetWorkspaceSize()汇总整个计算流程所需的 workspace 并返回。

第二段接口aclnnL1Loss则调用CommonOpExecutorRun完成实际的异步执行。

LpLoss 低阶算子与 AiCore 分发

loss/lp_loss/op_api/lp_loss.cpp 实现了l0op::LpLoss低阶算子封装：IsAiCoreSupport判断输入 dtype 是否在 AiCore 支持列表（FLOAT、FLOAT16、BF16）且两者 dtype 相同、reduction 合法；随后调用ADD_TO_LAUNCHER_LIST_AICORE将算子加入 AiCore 启动队列，并根据 reduction 模式分配输出 tensor（none 模式为 broadcast 后 shape，其余为 0 维）。

Device 侧：内核与计算图

loss/lp_loss/op_kernel/lp_loss.cpp 中的__global__ __aicore__ void lp_loss是 Device 侧内核入口，通过模板参数Reduction与Dtype在编译期实例化三种模式：

• Reduction == 0（none）：ElementwiseSch逐元素调度，输出完整 shape；
• Reduction == 1（sum）：ReduceSch+LpLossSumDag，全量归约求和；
• Reduction == 2（mean）：ReduceSch+LpLossMeanDag，求和后乘均值倒数meanVar。

Tiling 参数由 loss/lp_loss/op_host/arch35/lp_loss_tiling_arch35.cpp 在 Host 侧计算，通过REGISTER_TILING_DEFAULT(LpLossTilingData)与GET_TILING_DATA_WITH_STRUCT传递到内核；tiling 数据结构定义于 loss/lp_loss/op_kernel/arch35/lp_loss_tiling_struct.h，内核按键定义于 loss/lp_loss/op_kernel/arch35/lp_loss_tiling_key.h。

测试覆盖

仓库为该算子提供了完善的测试体系：

• 单元测试：loss/lp_loss/tests/ut/op_api/test_aclnn_l1_loss.cpp 覆盖 aclnn 接口层，loss/lp_loss/tests/ut/op_host/test_lp_loss_infershape.cpp 覆盖 shape 推导，loss/lp_loss/tests/ut/op_host/arch35/test_lp_loss_tiling.cpp 覆盖 tiling 计算；
• 系统测试（ST）：loss/lp_loss/tests/st/aclnnL1Loss/ 下包含atk_aclnnL1Loss.json（ATK 测试配置）与executor_aclnnL1Loss.py（Python 执行脚本），另有 loss/lp_loss/tests/st/arch35/ttk_kernel_lp_loss_david_st.csv 内核测试用例；
• golden 数据生成：loss/lp_loss/tests/assets/golden.py 用于生成期望结果，供测试比对。

使用建议与注意事项

1. reduction 与 out shape 的匹配：reduction=0（none）时务必为 out 分配与 broadcast 后一致的 shape；其余模式 out 必须为 0 维 tensor，否则第一段接口将返回ACLNN_ERR_PARAM_INVALID（161002）。
2. INT64 输入：当输入推导类型为 INT64 时，算子通过小算子拼接（Sub + Abs + ReduceSum）实现，仅支持 none 与 sum 两种模式；此时 out 的数据类型也应相应选择 INT64 或可转换类型。
3. 非连续 Tensor：接口层通过Contiguous与ViewCopy自动兼容非连续输入与输出，无需用户在调用前手动做连续性处理。
4. 空输入行为：若self或target为空 tensor，mean 模式输出 NAN、sum 模式输出 0、none 模式输出空 tensor，这是与 PyTorch 行为对齐的语义，使用时应知晓。
5. 确定性：aclnnL1Loss为确定性实现，可放心用于对可复现性有要求的训练场景。
6. p 值限制：当前仅支持p=1，如需计算 L2 等其他范数损失，需要等待后续算子演进或使用其他算子组合实现。

小结

LpLoss（L1Loss）算子是 CANN ops-nn 仓库中一个典型且完整的 aclnn 算子实现范例。本文从产品支持矩阵、数学定义、参数语义、约束条件、两段式接口调用到 Host/Device 双层源码实现，完整梳理了该算子的使用方式与内部原理。无论是直接调用acclnnL1Loss完成 L1 损失计算，还是以此为模板理解 CANN 算子的三层工程结构（op_api / op_host / op_kernel），本文提供的代码示例与源码路径都能成为可靠的参考起点。

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