CANN ops-nn 中 LpLossL1Loss算子深度解析从 aclnnL1Loss 接口到 NPU 内核实现【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn本文聚焦 CANNCompute 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 算子仅支持p1的场景此时即为经典的 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.hop_host算子原型注册lp_loss_def.cpp、shape 推导lp_loss_infershape.cpp与 tiling 计算op_host/arch35/lp_loss_tiling_arch35.cppop_kernelDevice 侧内核实现op_kernel/lp_loss.cpp及op_kernel/arch35/下的 DAG 计算图描述op_graph算子原型定义lp_loss_proto.htests包含 op_api 与 op_host 的单元测试、ST系统测试用例以及 golden 数据生成脚本。产品支持情况LpLossaclnnL1Loss算子在不同硬件产品上的支持情况如下表所示来源于 loss/lp_loss/README.md产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品×Atlas 推理系列产品√Atlas 训练系列产品√从源码角度看这一支持情况与 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 ErrorMAE。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 模式分别对应三个计算图模板LpLossOpreductionnoneCopyIn → Cast → Sub相减→ Abs取绝对值→ Cast → CopyOut逐元素输出|x-y|LpLossSumDagreductionsum在Sub → Abs之后追加ReduceSumOp完成全量求和LpLossMeanDagreductionmean在求和之后追加Muls乘以均值倒数meanVar其中meanVar由 tiling 阶段计算并注入。而 loss/lp_loss/op_kernel/lp_loss.cpp 中的内核入口函数lp_loss则根据Reduction模板参数在编译期选择ElementwiseSch逐元素调度或ReduceSch归约调度完成计算mean模式下还通过op.Process(static_castDTYPE_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_tHost 侧的整型属性取值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函数完整实现了上述参数的逐项校验包括空指针检查CheckNotNull3Tensorreduction 取值范围检查CheckReduction仅接受 0/1/2数据类型推导与支持范围检查CheckDtypeValidshape 与 broadcast 规则检查CheckShape维度上限 8。其中CheckDtypeValid中还包含两条与 CUDA 行为保持一致的特殊校验当reductionnone时若self不是浮点类型则target也不能是浮点类型当reductionmean时self和target至少有一个是浮点类型因为求均值需要除法。另外CheckFormat会对FORMAT_FRACTAL_NZ格式给出精度告警日志提示该格式可能导致精度问题。约束说明确定性计算aclnnL1Loss默认确定性实现同一输入多次执行结果一致p 值限制LpLoss 中p为计算 loss 的参数当前只支持p1对外接口名称为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会按语义直接处理reductionnone 返回空 tensorreductionmean 填充 NANreductionsum 填充 0不再进入内核计算。调用说明aclnn 两段式接口LpLossaclnnL1Loss属于标准的 CANN aclnn 两段式接口参见两段式接口说明必须先调用aclnnL1LossGetWorkspaceSize获取计算所需 workspace 大小以及包含算子计算流程的执行器再调用aclnnL1Loss执行计算。调用方式汇总如下完整样例见 loss/lp_loss/examples/test_aclnn_l1_loss.cpp调用方式样例代码说明aclnn 接口test_aclnn_l1_loss.cpp通过aclnnL1Loss接口调用 LpLoss 算子详见 aclnnL1Loss 接口文档第一段接口aclnnL1LossGetWorkspaceSizeaclnnStatus aclnnL1LossGetWorkspaceSize( const aclTensor* self, const aclTensor* target, int64_t reduction, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)第一段接口完成入参校验并返回 workspace 大小与执行器。参数细节如下参数名输入/输出描述使用说明数据类型数据格式维度(shape)非连续TensorselfaclTensor*输入公式中的输入 selfshape 需与 target 满足 broadcast 关系dtype 满足数据类型推导规则FLOAT、FLOAT16、BFLOAT16、INT64ND1-8√targetaclTensor*输入真实的标签shape 需与 self 满足 broadcast 关系dtype 满足数据类型推导规则FLOAT、FLOAT16、BFLOAT16、INT64ND1-8√reductionint64_t输入指定应用到输出的缩减支持 0(none) | 1(mean) | 2(sum)INT64--√outaclTensor*输出输出 tensor存放 L1Loss 计算结果reduction0 时 shape 与 broadcast 后一致reduction≠0 时为 0 维 tensorFLOAT、FLOAT16、BFLOAT16、INT64-0-8√workspaceSizeuint64_t*输出返回需要在 Device 侧申请的 workspace 大小-----executoraclOpExecutor**输出返回 op 执行器包含算子计算流程-----该接口返回aclnnStatus状态码具体参见 aclnn 返回码说明并在以下场景报错返回值错误码描述ACLNN_ERR_PARAM_NULLPTR161001self、target 或 out 是空指针ACLNN_ERR_PARAM_INVALID161002self 和 target 的数据类型不满足推导规则或推导后 dtype 不在支持范围之内ACLNN_ERR_PARAM_INVALID161002推导后的类型无法 cast 为 out 的数据类型ACLNN_ERR_PARAM_INVALID161002self 或 target 的维度大于 8ACLNN_ERR_PARAM_INVALID161002self 和 target 的 shape 不满足 broadcast 规则ACLNN_ERR_PARAM_INVALID161002reduction 值不在 0~2 范围之内ACLNN_ERR_PARAM_INVALID161002reduction0 时broadcast 后的 shape 与 out 的 shape 不一致ACLNN_ERR_PARAM_INVALID161002reduction≠0 时out 的维度大于 0ACLNN_ERR_PARAM_INVALID161002reductionnone、self 非浮点时target 是浮点类型ACLNN_ERR_PARAM_INVALID161002reductionmean 时self 和 target 均非浮点类型第二段接口aclnnL1LossaclnnStatus aclnnL1Loss( void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)参数名输入/输出描述workspace输入在 Device 侧申请的 workspace 内存地址workspaceSize输入在 Device 侧申请的 workspace 大小由第一段接口获取executor输入op 执行器包含算子计算流程stream输入指定执行任务的 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 1mean因此预期输出为(|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::vectorint64_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::vectorT hostData, const std::vectorint64_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::vectorint64_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::vectorint64_t selfShape {2, 2}; std::vectorint64_t targetShape {2, 2}; std::vectorint64_t outShape {}; // reduction1(mean) 时 out 为0维 void* selfDeviceAddr nullptr; void* targetDeviceAddr nullptr; void* outDeviceAddr nullptr; aclTensor* self nullptr; aclTensor* target nullptr; aclTensor* out nullptr; std::vectorfloat selfHostData {0, 1, 2, 3}; std::vectorfloat targetHostData {1, 1, 1, 1}; std::vectorfloat 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::vectorfloat 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)、动态 rankDynamicRankSupportFlag(true)、动态 shapeDynamicShapeSupportFlag(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是整个算子执行流程编排的核心其逻辑依次为参数校验调用CheckParams完成空指针、reduction、dtype、shape 的逐项校验空 tensor 快速路径若self或target为空直接按语义填充见前述约束说明数据类型对齐通过op::PromoteType计算推导类型并调用l0op::Cast将self、target统一 cast 到推导类型连续性处理通过l0op::Contiguous将输入转换为连续 tensor以兼容非连续输入Broadcast当self与targetshape 不一致时调用l0op::BroadcastTo将两者对齐到同一 shape核心计算分发若推导类型为 INT64则走小算子拼接路径GetL1LossFromInt64Sub → Abs → ReduceSumOp因为 LpLoss 的 AiCore 内核不支持 INT64否则调用l0op::LpLoss直接走 AiCore 内核输出处理将计算结果Cast到 out 的 dtype再通过l0op::ViewCopy写入 out兼容非连续 out返回 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 模式分配输出 tensornone 模式为 broadcast 后 shape其余为 0 维。Device 侧内核与计算图loss/lp_loss/op_kernel/lp_loss.cpp 中的__global__ __aicore__ void lp_loss是 Device 侧内核入口通过模板参数Reduction与Dtype在编译期实例化三种模式Reduction 0noneElementwiseSch逐元素调度输出完整 shapeReduction 1sumReduceSchLpLossSumDag全量归约求和Reduction 2meanReduceSchLpLossMeanDag求和后乘均值倒数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 计算系统测试STloss/lp_loss/tests/st/aclnnL1Loss/ 下包含atk_aclnnL1Loss.jsonATK 测试配置与executor_aclnnL1Loss.pyPython 执行脚本另有 loss/lp_loss/tests/st/arch35/ttk_kernel_lp_loss_david_st.csv 内核测试用例golden 数据生成loss/lp_loss/tests/assets/golden.py 用于生成期望结果供测试比对。使用建议与注意事项reduction 与 out shape 的匹配reduction0none时务必为 out 分配与 broadcast 后一致的 shape其余模式 out 必须为 0 维 tensor否则第一段接口将返回ACLNN_ERR_PARAM_INVALID161002。INT64 输入当输入推导类型为 INT64 时算子通过小算子拼接Sub Abs ReduceSum实现仅支持 none 与 sum 两种模式此时 out 的数据类型也应相应选择 INT64 或可转换类型。非连续 Tensor接口层通过Contiguous与ViewCopy自动兼容非连续输入与输出无需用户在调用前手动做连续性处理。空输入行为若self或target为空 tensormean 模式输出 NAN、sum 模式输出 0、none 模式输出空 tensor这是与 PyTorch 行为对齐的语义使用时应知晓。确定性aclnnL1Loss为确定性实现可放心用于对可复现性有要求的训练场景。p 值限制当前仅支持p1如需计算 L2 等其他范数损失需要等待后续算子演进或使用其他算子组合实现。小结LpLossL1Loss算子是 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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考