CANN ops-nn HardtanhGrad 算子解析Hardtanh 反向传播原理、aclnn 两段式调用与源码实现【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn本文以 ops-nn 仓库中 HardtanhGrad 算子文档 为核心系统讲解激活函数 Hardtanh 反向算子的数学定义、参数约束、aclnn 两段式调用方式与完整 C 调用样例并结合activation/hardtanh_grad目录下的算子定义、Tiling 与 AIV 内核源码剖析其底层实现细节帮助读者在模型训练反向传播场景中正确接入并验证该算子。一、算子功能与反向计算公式HardtanhGrad 是激活函数 Hardtanh硬限幅激活的反向传播算子。Hardtanh 的前向计算将输入限制在线性区间 [min, max] 内区间之外的取值被“削平”为常数因此其梯度在区间外为 0区间内为 1。反向算子的计算分两步第一步根据前向输入self所在位置计算自梯度grad_self$$ grad_self_{i} \begin{cases} 0, if \ self_{i}max \ 0, if \ self_{i}min \ 1, otherwise \end{cases} $$第二步将上游回传的梯度grad_output逐元素相乘得到本层输入梯度res$$ res_{i} grad_output_{i} \times grad_self_{i} $$从语义上看该算子本质上是一个“梯度门控”只有在前向输入落在 (min, max) 线性区域内的位置才允许梯度通过区间外的位置梯度被直接置零。这也是它在 PyTorch 等框架中作为hardtanh反向算子的对应实现——算子定义文件中明确标注了“Compatible with the PyTorch operator HardtanhGrad”见 hardtanh_grad_proto.h。二、产品支持情况根据 README 与 aclnn 接口文档该算子的产品支持情况如下产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√Atlas A2 训练系列产品/Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品√Atlas 推理系列产品√Atlas 训练系列产品√注意数据类型支持存在差异aclnn 接口文档指出Atlas 推理系列/训练系列产品910/310p的数据类型仅支持 FLOAT16、FLOAT不含 BFLOAT16其余产品支持 BFLOAT16、FLOAT16、FLOAT 三种类型。三、参数说明与约束3.1 参数总表参数名输入/输出/属性描述数据类型数据格式result输入反向传播过程中上一步输出的梯度公式中的 grad_outputBFLOAT16、FLOAT16、FLOATNDgrad输入正向的输入数据公式中的 selfBFLOAT16、FLOAT16、FLOATNDmin_val属性线性范围的下限公式中的 minFLOAT-max_val属性线性范围的上限公式中的 maxFLOAT-out输出计算得到梯度公式中的 resBFLOAT16、FLOAT16、FLOATND说明README 中的参数名result、grad是算子 IR图模式视角的命名在 aclnn 接口中对应的入参分别名为gradOutputresult和selfgradmin_val/max_val则以aclScalar*形式传入。两套命名指向同一语义阅读代码时注意对应关系。3.2 形状与类型约束result、grad 和 out 的 shape 和数据类型必须一致README 约束说明维度rank范围 0-8不支持空 Tensor三个 Tensor 均支持非连续strides 非紧凑的 Tensor——接口文档中“非连续Tensor”一列均标记为 √这一能力的底层支撑见第六节的计算路径内部自动插入l0op::Contiguous数据类型为 BFLOAT16、FLOAT16、FLOAT 之一且三者必须相同。从源码看这些约束在两个层面得到双重校验算子定义层hardtanh_grad_def.cpp 中声明了两个输入result、grad与输出y数据类型统一为{ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT}格式为 ND属性min_val、max_val为必选 Float 属性默认值分别为 -1.0 和 1.0即 PyTorchF.hardtanh的默认区间。该定义同时注册在ascend950与ascend350两款芯片上并开启了动态 shape/rank 支持DynamicShapeSupportFlag、DynamicRankSupportFlag均为 true说明算子支持运行时才确定形状的张量。Tiling 层hardtanh_grad_tiling_arch35.cpp 中的CalcInputDtype、CalcOutputDtype会逐一检查三个张量类型是否落在 {FLOAT16, BF16, FLOAT} 内且互相一致CheckShape则强制输入 result 与 grad、输入 result 与输出 y 的 shape 完全相同违反时返回GRAPH_FAILED并打印具体原因与 aclnn 接口文档中ACLNN_ERR_PARAM_INVALID的报错场景一一对应。3.3 形状推导hardtanh_grad_infershape.cpp 仅有一行核心实现IMPL_OP_INFERSHAPE(HardtanhGrad).InferShape(Ops::Base::InferShape4Elewise);即复用逐元素elewise算子的通用形状推导规则输出 shape 等于输入 shape。这与“反向算子不改变张量形状”的语义吻合。四、aclnn 两段式接口CANN 的 aclnn 算子均采用两段式接口先调用aclnnHardtanhBackwardGetWorkspaceSize获取计算所需 workspace 大小并构建执行器aclOpExecutor内含算子计算流程再调用aclnnHardtanhBackward在指定 stream 上执行计算。函数原型如下摘自 aclnnHardtanhBackward.mdaclnnStatus aclnnHardtanhBackwardGetWorkspaceSize( const aclTensor* gradOutput, const aclTensor* self, const aclScalar* min, const aclScalar* max, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)aclnnStatus aclnnHardtanhBackward( void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)4.1 第一段接口参数参数名输入/输出描述数据类型数据格式维度(shape)非连续TensorgradOutput输入反向传播过程中上一步输出的梯度公式中的 grad_outputBFLOAT16、FLOAT16、FLOATND0-8√self输入正向的输入数据公式中的 selfBFLOAT16、FLOAT16、FLOATND0-8√min输入线性范围的下限BFLOAT16、FLOAT16、FLOAT---max输入线性范围的上限BFLOAT16、FLOAT16、FLOAT---out输出计算得到的梯度结果BFLOAT16、FLOAT16、FLOATND0-8√workspaceSize输出返回需要在 Device 侧申请的 workspace 大小----executor输出返回 op 执行器包含算子计算流程----三个 TensorgradOutput、self、out均不支持空 Tensor且 shape 必须一致、数据类型必须一致。4.2 计算路径接口头文件 aclnn_hardtanh_backward.h 的注释给出了算子内部的计算 DAGgradOutput ── l0op::Contiguous ──┐ ├── l0op::HardtanhGrad ── l0op::ViewCopy ── out self ─────── l0op::Contiguous ──┘ ^ min ────────────────────────────────────────┤ max ────────────────────────────────────────┘也就是说用户传入的非连续 Tensor 会先被Contiguous规整化真正的核心计算由 Level0 算子l0op::HardtanhGrad完成其声明见 hardtanh_grad.h最后通过ViewCopy按输出 Tensor 的 strides 写回。这条路径解释了为什么接口层可以“支持非连续 Tensor”——连续性由框架内部自动补齐用户无需自行调用 contiguous。4.3 错误码第一段接口完成入参校验典型报错场景返回值错误码描述ACLNN_ERR_PARAM_NULLPTR161001传入的 gradOutput、self、out、min、max 存在空指针ACLNN_ERR_PARAM_INVALID161002gradOutput、self、out 的数据类型不在支持范围内ACLNN_ERR_PARAM_INVALID161002gradOutput 和 self 数据类型不同ACLNN_ERR_PARAM_INVALID161002gradOutput 和 self 的 shape 不同ACLNN_ERR_PARAM_INVALID161002gradOutput 或 self 的维度大于 8此外接口文档明确aclnnHardtanhBackward 为确定性实现默认确定性计算相同输入多次运行结果一致。五、完整调用样例仓库提供了可直接编译运行的完整样例 test_aclnn_hardtanh_grad.cpp覆盖“初始化 → 构造 Tensor → 两段式调用 → 同步取回 → 资源释放”全流程。样例使用 shape 为{4, 2}的 FLOAT 数据gradOutputHostData {0, 1, 2, 3, 4, 5, 6, 7}上游梯度selfHostData {1, 1, 1, 2, 1, 2, 3, 3}前向输入minValue 1.2fmaxValue 2.4f线性区间。按公式逐元素验证self 落在 (1.2, 2.4) 内的元素只有2共两处因此输出中仅对应位置的grad_output即 4 和 5保留其余为 0。#include iostream #include vector #include acl/acl.h #include aclnnop/aclnn_hardtanh_backward.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 shape_size 1; for (auto i : shape) { shape_size * i; } return shape_size; } 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; } // 将 host 数据上传到 device 并封装为 aclTensor含连续 strides 计算 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数据格式为 ND *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 填写 aclrtStream stream; auto ret Init(deviceId, stream); CHECK_RET(ret 0, LOG_PRINT(Init acl failed. ERROR: %d\n, ret); return ret); // 2. 构造输入与输出 std::vectorint64_t gradOutputShape {4, 2}; std::vectorint64_t selfShape {4, 2}; std::vectorint64_t outShape {4, 2}; void* gradOutputDeviceAddr nullptr; void* selfDeviceAddr nullptr; void* outDeviceAddr nullptr; aclTensor* gradOutput nullptr; aclTensor* self nullptr; aclTensor* out nullptr; aclScalar* minVal nullptr; aclScalar* maxVal nullptr; std::vectorfloat gradOutputHostData {0, 1, 2, 3, 4, 5, 6, 7}; std::vectorfloat selfHostData {1, 1, 1, 2, 1, 2, 3, 3}; std::vectorfloat outHostData {0, 0, 0, 0, 0, 0, 0, 0}; float minValue 1.2f; float maxValue 2.4f; ret CreateAclTensor(gradOutputHostData, gradOutputShape, gradOutputDeviceAddr, aclDataType::ACL_FLOAT, gradOutput); CHECK_RET(ret ACL_SUCCESS, return ret); ret CreateAclTensor(selfHostData, selfShape, selfDeviceAddr, aclDataType::ACL_FLOAT, self); CHECK_RET(ret ACL_SUCCESS, return ret); // min/max 通过 aclScalar 表达 minVal aclCreateScalar(minValue, aclDataType::ACL_FLOAT); CHECK_RET(minVal ! nullptr, return ret); maxVal aclCreateScalar(maxValue, aclDataType::ACL_FLOAT); CHECK_RET(maxVal ! nullptr, return ret); ret CreateAclTensor(outHostData, outShape, outDeviceAddr, aclDataType::ACL_FLOAT, out); CHECK_RET(ret ACL_SUCCESS, return ret); // 3. 调用 CANN 算子库 API两段式 uint64_t workspaceSize 0; aclOpExecutor* executor; // 第一段接口获取 workspace 大小并构建执行器 ret aclnnHardtanhBackwardGetWorkspaceSize(gradOutput, self, minVal, maxVal, out, workspaceSize, executor); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnHardtanhBackwardGetWorkspaceSize failed. ERROR: %d\n, ret); return ret); // 根据 workspaceSize 申请 device 内存可能为 0 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 aclnnHardtanhBackward(workspaceAddr, workspaceSize, executor, stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnHardtanhBackward 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(float), 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 和 aclScalar aclDestroyTensor(gradOutput); aclDestroyTensor(self); aclDestroyScalar(minVal); aclDestroyScalar(maxVal); aclDestroyTensor(out); // 7. 释放 device 资源 aclrtFree(gradOutputDeviceAddr); aclrtFree(selfDeviceAddr); aclrtFree(outDeviceAddr); if (workspaceSize 0) { aclrtFree(workspaceAddr); } aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0; }具体编译与执行过程可参考仓库内的编译与运行样例通用说明。集成到自己工程时需要修改的核心只有三处输入 shape/数据、minValue/maxValue取值以及 deviceId。六、源码实现剖析activation/hardtanh_grad目录采用 ops-nn 标准分层结构以下沿“定义 → Tiling → Kernel”链路剖析关键实现。6.1 算子注册与属性默认值hardtanh_grad_def.cpp 中的OpAICoreConfig配置值得注意OpAICoreConfig aicoreConfig; aicoreConfig.DynamicCompileStaticFlag(true) .DynamicFormatFlag(false) .DynamicRankSupportFlag(true) .DynamicShapeSupportFlag(true) .NeedCheckSupportFlag(false) .PrecisionReduceFlag(true) .ExtendCfgInfo(opFile.value, hardtanh_grad_apt); this-AICore().AddConfig(ascend950, aicoreConfig); this-AICore().AddConfig(ascend350, aicoreConfig);opFile.value指向hardtanh_grad_apt即内核源码 hardtanh_grad_apt.cpp 编译产物PrecisionReduceFlag(true)表明该算子允许精度收敛这与反向算子常以 FP32 中间精度计算、最后落回 BF16/FP16 输出的做法一致见下一节 Cast 环节属性min_val默认 -1.0、max_val默认 1.0与 PyTorch 默认区间一致。6.2 Tiling类型/形状校验与 workspacehardtanh_grad_tiling_arch35.cpp 的RunTiling按顺序执行CalcInputDtype()/CalcOutputDtype()校验三个张量类型合法且一致CheckShape()校验 shape 两两相同依据输出类型分发到HardtanhGradDaghalf/HardtanhGradDagbfloat16_t/HardtanhGradDagfloat通过ElewiseBaseTiling::DoTiling完成逐元素算子的通用 Tilingblock 数、每 AICore 处理的数据块划分等从运行时属性中读取min_val第 0 个属性与max_val第 1 个属性写入HardtanhGradTilingData随 Tiling 数据下传给内核SetTilingData()固定申请 16 MBASCEND_WORKSPACE 16777216workspace并设置tilingKey与blockDim。内核声明处有KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY)说明该算子完全在 AIV向量核上执行不涉及 AICube——这与“纯逐元素比较选择”的计算形态匹配。模板参数见 hardtanh_grad_struct.hschMode取 0/1 两种调度模式dType固定为HardtanhGrad_TPL由 Tiling 阶段选择具体调度模式后通过 tilingKey 路由到对应的内核实例。6.3 内核 DAG比较、NaN 分支与选择核心计算逻辑在 hardtanh_grad_dag.h 中以声明式 DAG 描述模板参数U为张量存储类型half/bfloat16_t/float中间计算统一提升到T floattemplate typename U, typename T float struct HardtanhGradDag { // copy in cast to float using OpCopyIn0 BindVec::CopyInU, Placeholder::In0U; using OpCopyIn1 BindVec::CopyInU, Placeholder::In1U; using OpCopyIn0Cast BindVec::CastT, U, CAST_MODE_NONE, OpCopyIn0; // compare using OpCompareNaN BindVec::Compareuint8_t, T, COMPARE_MODE_NE, OpCopyIn0Cast, OpCopyIn0Cast; using OpCompareLowerThanMax BindVec::Compareuint8_t, T, COMPARE_MODE_LT, OpCopyIn0Cast, Placeholder::VarT, 1; using OpCompareGreaterThanMin BindVec::Compareuint8_t, T, COMPARE_MODE_GT, OpCopyIn0Cast, Placeholder::VarT, 0; // mask using OpMaskAnd BindVec::Anduint8_t, OpCompareLowerThanMax, OpCompareGreaterThanMin; using OpMaskOr BindVec::Oruint8_t, OpMaskAnd, OpCompareNaN; // select using OpSelect BindVec::Selectuint8_t, U, SELECT_MODE_VS, OpMaskOr, OpCopyIn1, ConstValueZero; using OpCopyOut BindVec::CopyOutU, Placeholder::Out0U, OpSelect; using Outputs ElemsOpCopyOut; using MemCfg MemOptCfgMemLevel::LEVEL_2; using OpDag DAGSchOutputs, void, MemCfg; };结合内核入口 hardtanh_grad_apt.cpp 中SetVarfloat, 0(minVal)、SetVarfloat, 1(maxVal)可以把执行流水归纳为CopyIn Cast将两路输入拷入并转为 float保证比较在 FP32 精度下进行低精度类型下比较 1.2、2.4 这类小数边界FP32 中间精度可避免边界误判三路比较x min_valGT 模式、x max_valLT 模式以及x ! xNE 模式自比较即 NaN 检测Mask 合成(x max_val) (x min_val) || isnan(x)Selectmask 为真的位置输出grad第 1 路输入否则输出常量 0.0CopyOut以原类型写回输出。文件头注释给出了其意图的等价表达torch.where((result min_val) | (result max_val), 0.0, grad)。有两点实现细节值得注意NaN 处理x ! x仅在 x 为 NaN 时为真因此 NaN 位置的 mask 被置 1梯度grad会直接穿透输出。这符合 PyTorch 等框架中 NaN 需要继续向后传播以便调试与损失计算的惯例边界语义从源码看比较模式码使用的是严格比较COMPARE_MODE_GT/COMPARE_MODE_LT即x min_val或x max_val恰好落在边界上时 mask 为假、输出为 0而 README 公式按“otherwise 取 1”解读时边界点梯度应为 1。二者在边界取值上存在细微差异实际集成时若前向 Hardtanh 的输出作为 self 回传边界点恰好可能出现建议以 tests 目录 中的 golden 数据golden.py 与 kernel 测试核对边界用例的实际行为为准。6.4 模板与调度hardtanh_grad_apt.cpp 中template uint64_t schMode, uint64_t dType __global__ __aicore__ void hardtanh_grad(GM_ADDR result, GM_ADDR grad, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) { ... TPipe pipe; if constexpr (dType static_castuint64_t(HardtanhGrad_TPL)) { ElementwiseSchschMode, HardtanhGradOp::HardtanhGradDagDTYPE_RESULT::OpDag sch((tilingData.baseTiling), pipe); sch.template SetVarfloat, 0(static_castfloat(tilingData.minVal)); sch.template SetVarfloat, 1(static_castfloat(tilingData.maxVal)); sch.Init(result, grad, y); sch.Process(); } return; }ElementwiseSch是仓库内通用的逐元素算子调度器负责按 Tiling 划分的 block 逐块执行上述 DAGmin/max通过SetVar注入 DAG 中的Placeholder::Var占位符作为标量常量参与向量化比较。if constexpr保证不同 schMode 实例之间无运行时开销。七、图模式调用除 aclnn 调用外该算子支持通过算子 IR 构图调用。hardtanh_grad_proto.h 中的注册声明为REG_OP(HardtanhGrad) .INPUT(result, TensorType({DT_BF16, DT_FLOAT16, DT_FLOAT})) .INPUT(grad, TensorType({DT_BF16, DT_FLOAT16, DT_FLOAT})) .OUTPUT(y, TensorType({DT_BF16, DT_FLOAT16, DT_FLOAT})) .ATTR(min_val, Float, -1.0) .ATTR(max_val, Float, 1.0) .OP_END_FACTORY_REG(HardtanhGrad)即在图构建中以HardtanhGrad为名创建节点两个输入 一个输出 min_val/max_val两个可省略默认 -1.0/1.0的浮点属性即可与 aclnn 接口的语义完全对应。八、测试体系与验证该算子配套了完整的 UT/ST 测试可用于集成后的自验证aclnn 接口 UTtest_aclnn_hardtanh_backward.cpp验证接口行为与参数校验形状推导 UTtest_Hardtanh_grad_infershape.cpp覆盖 elewise 形状推导Tiling UTtest_hardtanh_grad_tiling.cpp验证类型/形状校验与 Tiling 数据生成Kernel UTtest_hardtanh_grad_apt.cpp配合 gen_data.py 生成输入数据并与 golden 对比ST 测试数据ttk_kernel_hardtanh_grad_st.csv 与 ttk_aclnn_hardtanh_backward_st.csv分别覆盖内核级与 aclnn 接口的系统级测试用例矩阵。集成时建议的验证顺序是先用第五节样例跑通最小用例再对照 README 公式手工核算小样本如样例中 self2 的位置保留梯度、其余置零最后参考 ST CSV 中的用例组合覆盖 BF16/FP16、多维 shape 与非连续 Tensor 场景。九、小结HardtanhGrad 是 ops-nn 中典型的逐元素反向算子数学上只是“区间掩码 × 上游梯度”工程上却完整覆盖了 CANN 算子开发的标准链路——op_def 声明含类型/属性/动态 shape 能力、elewise 形状推导、AIV 内核上的声明式 DAGCast→Compare→Mask→Select、以及 aclnn 两段式接口与非连续 Tensor 的自动规整化。对其实现细节FP32 中间精度比较、NaN 透传、边界语义的把握是正确集成并排查该算子精度问题的关键依据。【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考