拓冰网站定制 
CANN ops-nn 算子 NpuClearFloatStatus 深度解析AI Core 浮点溢出状态寄存器清除原理、约束与图模式调用指南【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn导读NpuClearFloatStatus 是 CANN ops-nn 算子库control/npu_clear_float_status中一个特殊的硬件状态管理算子其作用是在 NPU 上清除每个 AI Core 的浮点溢出状态寄存器并固定输出 8 个 float32 零值。它不参与任何数值运算而是用于训练/推理流程中的浮点异常状态复位与状态查询前置。本文以该算子的官方说明文档为主体结合仓库内算子定义、Shape/DataType 推导、Tiling 与 Kernel 实现及单元测试完整讲解其功能语义、输入输出约束、产品支持情况、底层实现原理并给出基于 GE 图模式的编译调用与验证方法。算子功能与数学语义功能说明根据 README 的官方定义NpuClearFloatStatus 算子的功能为清除 NPU 每个 AI Core 的浮点溢出状态寄存器输出固定为 8 个 float32 零值。该算子的“计算”表达式非常简单且固定$$ data zeros(8, \text{dtype}float32) $$也就是说无论输入数据是什么内容算子输出恒为一个长度为 8、全为 0 的 float32 一维张量。其真实工作负载在于硬件侧通过触发向量计算单元执行写操作来清除浮点溢出状态标志详见下文 Kernel 实现。从语义上讲本算子通常与 npu_get_float_status读取浮点状态配合使用先清除状态寄存器再执行需要监控的运算最后读取状态以判断是否发生浮点溢出从而在不打断计算流的前提下实现溢出检测。参数说明算子共有一个输入和一个输出均无属性Attribute参数。官方参数表如下参数名输入/输出/属性描述数据类型数据格式addr输入地址占位符shape 为 (8,)数据内容不参与计算FLOATNDdata输出固定输出 8 个 float32 零值shape 为 (8,)FLOATNDaddr输入仅作为算子输入接口的地址占位符存在其 shape 为 (8,)但数据内容完全不影响计算结果。在 npu_clear_float_status_def.cpp 中可以看到该输入被注册为必选REQUIRED、数据类型ge::DT_FLOAT、格式ge::FORMAT_ND并开启AutoContiguous()。data输出固定输出 8 个 float32 零值shape 恒为 (8,)同样为 ND 格式npu_clear_float_status_def.cpp。约束说明使用该算子时必须满足以下约束来源READMEaddr 数据类型必须为 float32输出 data 固定为 8 个 float32 零值与输入数据内容无关addr 仅作为算子输入接口占位符其数据内容不参与计算。在源码中约束 1 被双重校验一方面在 InferShape 阶段 检查addr的 dtype 必须为DT_FLOAT否则返回GRAPH_FAILED另一方面在 Tiling 阶段 同时对输入addr和输出data的 dtype 进行校验两者均必须是 float32。产品支持情况NpuClearFloatStatus 的产品支持矩阵如下来源README产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品×Atlas A2 训练系列产品 / Atlas A2 推理系列产品×Atlas 200I/500 A2 推理产品√Atlas 推理系列产品√Atlas 训练系列产品√从源码结构看该算子仅注册了ascend950一种 AICore 配置npu_clear_float_status_def.cpp并在op_host/arch35与op_kernel/arch35目录下提供了 arch35 架构的 Tiling 与 Kernel 实现这与“Ascend 950PR / Ascend 950DT 支持”的产品定位一致Atlas A3/A2 系列不支持该算子。图编译期行为Shape、Format 与 DataType 推导在 GEGraph Engine图编译阶段算子的输出形状、格式与数据类型由三段独立的注册逻辑共同决定常量定义公共常量位于 npu_clear_float_status_common.hnamespace NpuCfs { constexpr int64_t OUTPUT_DIM 8; // 输出固定 8 个元素 constexpr size_t EXPECTED_DIM_NUM 1; // 输出为 1 维 constexpr int32_t ADDR_IDX 0; // input addr 索引 constexpr int32_t DATA_IDX 0; // output data 索引 }InferShape输出恒为 (8,)InferShape 实现 将输出 shape 的维度数固定为 1、第 0 维固定为 8完全不依赖输入 addr 的 shape——这与 README 中“addr 仅作占位符”的语义保持一致。InferFormat输出恒为 NDInferFormat 实现 将输出的原始格式与存储格式均固定为ge::FORMAT_ND。InferDataType输出恒为 float32输出数据类型推导位于 npu_clear_float_status_graph_infer.cpp固定将输出 data 的 dtype 设置为ge::DT_FLOAT与输入 addr 的 dtype 无关。源码注释中特别说明addr 是地址占位符而非数据 Tensor若输出跟随输入 dtype输入非 float32 时会导致图编译阶段下游 dtype 推导异常。因此输出 dtype 必须显式固定。值得注意的架构细节是本算子的 InferShape/InferFormat 与 InferDataType 分别注册在两个文件中前者通过IMPL_OP_INFERSHAPE注册npu_clear_float_status_infershape.cpp后者通过IMPL_OP注册npu_clear_float_status_graph_infer.cpp。Tiling 阶段全核启动与系统 WorkspaceTiling切分/编排逻辑位于 npu_clear_float_status_tiling_arch35.cpp其执行流程如下输入校验校验 addr/data 的 dtype 均为 float32。平台信息获取通过PlatformAscendC获取 AI CoreAIV核数coreNum与 UB 内存大小ubSize并做非零防御性检查。Tiling 计算needCoreNum coreNum即全核启动所有 AI Core 都必须执行因为要清除的是每个 AI Core 的溢出状态寄存器。在 ComputeTiling 中还对coreNum是否超过INT32_MAX做了防御性检查避免 int64→int32 强转溢出。Workspace 申请本算子无需用户 workspace仅申请系统 workspace通过GetLibApiWorkSpaceSize获取见 GetWorkspaceSize。设置 BlockDimSetBlockDim(needCoreNum)保证所有 AI Core 被调度执行。UB 配置采用 DCACHE_SIZE 128KB默认值、STATIC_UB_ESTIMATE 0 的策略要求动态 UB 池ubSize − 128KB不小于 TBuf 所需容量VEC_DUP_SIZE * sizeof(half) 38400 × 2 75KB并通过SetLocalMemorySize设置本地内存ApplyUbConfig。TilingKey本算子单 dtype 单场景固定使用NPU_CLEAR_FLOAT_STATUS_SCH_MODE_DEFAULT值为 0定义见 npu_clear_float_status_tiling_key.h。Tiling 数据结构的全部内容仅为一个字段npu_clear_float_status_tiling_data.hstruct NPUClearFloatStatusTilingData { int32_t needCoreNum 0; // 需要启动的核数 物理核数 };Kernel 实现原理向量指令触发状态清除 SIMT 写零Kernel 入口npu_clear_float_status.cpp通过模板参数schMode实例化默认调度模式调用NsNpuClearFloatStatus::ProcessDTYPE_ADDR(addr, data, tilingData)。核心实现位于 npu_clear_float_status_simt.h分为两个关键步骤步骤一向量单元写操作清除溢出状态标志由于 ascend950 平台不支持set_overflow_status接口实现上采用“间接触发”的方式在标量作用域内申请TBufQuePosition::VECCALCfloat16[38400]约 75KB然后连续执行 6 次Duplicate(dataUbInput, value, VEC_DUP_SIZE)value 依次取 3~8。填充值本身无实际意义目的是触发向量计算单元执行从而清除浮点溢出状态标志对应代码注释见 npu_clear_float_status_simt.h。步骤二SIMT VF 写 8 个 float32 零值到输出通过asc_vf_callOpNPUClearFloatStatusSimtT(dim3(THREAD_NUM), OUTPUT_SIZE, outputGm)启动 128 线程的 SIMT 向量函数采用Grid-Stride 循环向输出 GM 写入totalElements8个零值for (uint32_t idx blockIdx.x * blockDim.x threadIdx.x; idx static_castuint32_t(totalElements); idx blockDim.x * gridDim.x) { output[idx] static_castT(0); }由于totalElements8远小于 blockDim × gridDim实际仅 core 0 的前 8 个线程执行写入其余线程在循环条件判断后立即退出。代码对totalElements 0做了防御性检查避免负数经static_castuint32_t转为巨大无符号数导致循环越界。SIMT VF 写 GM 后由框架自动保证 cache 一致性无需显式DataCacheCleanAndInvalid。从源码结构可以推断输入addr在 Kernel 中被显式忽略(void)addr;见 npu_clear_float_status_simt.h这再次印证了 README 中“addr 数据内容不参与计算”的语义。图模式调用与验证调用方式README 中明确本算子支持图模式GE 图模式调用样例为 test_geir_npu_clear_float_status.cpp完整的算子编译与验证流程参见 算子调用指南。快速体验基于项目 build.sh根据 算子调用指南可无需搭建调用工程直接基于项目脚本执行样例# 基于 ops-nn 包执行图模式样例graph 模式无需指定 pkg_mode 和 vendor_name bash build.sh --run_example npu_clear_float_status graph参数说明${op}算子名小写下划线形式此处为npu_clear_float_status${mode}调用方式graph表示图模式调用${soc_version}可选NPU 型号设置为ascend950时会额外运行arch35目录下的示例文件${simulator}可选仿真模式目前仅支持 eageraclnn 调用场景。GE 图模式调用要点从 test_geir_npu_clear_float_status.cpp 可以看到图模式调用的完整骨架其关键步骤如下初始化 GE通过ge::GEInitialize(global_options)初始化其中{ge.exec.deviceId: 0, ge.graphRunMode: 0, ge.exec.precision_mode: must_keep_origin_dtype}。创建算子实例auto add1 op::NPUClearFloatStatus(add1)算子在图中无属性。构图使用Data占位算子接入输入addrshape 固定 (8,)输出data的 shape 固定为{8}与输入 addr 的 shape 无关对应 CreateOppInGraph。创建 Session 并运行session-AddGraph(graphId, graph, graphOptions)添加图session-RunGraph(graphId, input, output)执行最后GEFinalize()释放资源。结果检查样例运行后通过PrintReport输出各 Case 的 Build / RunGraph / OutputExists 状态。样例中还演示了静态 shapeS 模式与动态 shapeD 模式graph_shape 用 -1 占位两种构图形制并覆盖FP32 fixed_8的 dtype/shape 组合矩阵可供开发者在集成时参考。测试与 Golden 验证仓库为该算子提供了两层验证UT 单元测试test_npu_clear_float_status_infershape.cpp验证 InferShape 函数在输入 addr shape 为 (8,) 时返回成功。test_npu_clear_float_status_tiling.cpp以UB_SIZE: 262144, CORE_NUM: 64的平台信息构造 TilingContext断言 TilingKey 为 0、needCoreNum 64、BlockDim 64并校验系统 workspace 大小。Golden 脚本golden.py 定义了 Kernel 与 GEIR 共用的 Golden 函数输出恒为tf.zeros([8], dtypetf.float32)并声明精度标准为binary_equal输出为比特级精确的零属非计算型算子同时提供 TensorFlow 第三方实现用于 GEIR 交叉校验由于输出恒为零结果确定且与输入数据无关。总结NpuClearFloatStatus 是 CANN ops-nn 中一类典型的“硬件状态管理型”算子数学上它只是输出 8 个 float32 零值但真正的作用在于通过向量计算单元触发执行来清除每个 AI Core 的浮点溢出状态寄存器从而为后续的浮点溢出检测提供干净的状态起点。其设计要点可归纳为固定语义输出恒为zeros(8, float32)输入 addr 仅为占位符全核调度Tiling 将 BlockDim 设为物理核数确保所有 AI Core 的状态寄存器都被清除明确的约束与校验addr/data 必须为 float32、ND 格式Shape/Format/DataType 推导均在编译期固定无需依赖输入内容产品限定仅支持 Ascend 950PR/950DT、Atlas 200I/500 A2 推理产品、Atlas 推理/训练系列产品不支持 Atlas A3/A2 系列。开发者可参考 test_geir_npu_clear_float_status.cpp 在业务图中集成该算子并结合 npu_get_float_status 实现“清除 → 计算 → 查询”的浮点溢出监控闭环。【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN ops-nn 算子 NpuClearFloatStatus 深度解析：AI Core 浮点溢出状态寄存器清除原理、约束与图模式调用指南
CANN ops-nn 算子 NpuClearFloatStatus 深度解析AI Core 浮点溢出状态寄存器清除原理、约束与图模式调用指南【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn导读NpuClearFloatStatus 是 CANN ops-nn 算子库control/npu_clear_float_status中一个特殊的硬件状态管理算子其作用是在 NPU 上清除每个 AI Core 的浮点溢出状态寄存器并固定输出 8 个 float32 零值。它不参与任何数值运算而是用于训练/推理流程中的浮点异常状态复位与状态查询前置。本文以该算子的官方说明文档为主体结合仓库内算子定义、Shape/DataType 推导、Tiling 与 Kernel 实现及单元测试完整讲解其功能语义、输入输出约束、产品支持情况、底层实现原理并给出基于 GE 图模式的编译调用与验证方法。算子功能与数学语义功能说明根据 README 的官方定义NpuClearFloatStatus 算子的功能为清除 NPU 每个 AI Core 的浮点溢出状态寄存器输出固定为 8 个 float32 零值。该算子的“计算”表达式非常简单且固定$$ data zeros(8, \text{dtype}float32) $$也就是说无论输入数据是什么内容算子输出恒为一个长度为 8、全为 0 的 float32 一维张量。其真实工作负载在于硬件侧通过触发向量计算单元执行写操作来清除浮点溢出状态标志详见下文 Kernel 实现。从语义上讲本算子通常与 npu_get_float_status读取浮点状态配合使用先清除状态寄存器再执行需要监控的运算最后读取状态以判断是否发生浮点溢出从而在不打断计算流的前提下实现溢出检测。参数说明算子共有一个输入和一个输出均无属性Attribute参数。官方参数表如下参数名输入/输出/属性描述数据类型数据格式addr输入地址占位符shape 为 (8,)数据内容不参与计算FLOATNDdata输出固定输出 8 个 float32 零值shape 为 (8,)FLOATNDaddr输入仅作为算子输入接口的地址占位符存在其 shape 为 (8,)但数据内容完全不影响计算结果。在 npu_clear_float_status_def.cpp 中可以看到该输入被注册为必选REQUIRED、数据类型ge::DT_FLOAT、格式ge::FORMAT_ND并开启AutoContiguous()。data输出固定输出 8 个 float32 零值shape 恒为 (8,)同样为 ND 格式npu_clear_float_status_def.cpp。约束说明使用该算子时必须满足以下约束来源READMEaddr 数据类型必须为 float32输出 data 固定为 8 个 float32 零值与输入数据内容无关addr 仅作为算子输入接口占位符其数据内容不参与计算。在源码中约束 1 被双重校验一方面在 InferShape 阶段 检查addr的 dtype 必须为DT_FLOAT否则返回GRAPH_FAILED另一方面在 Tiling 阶段 同时对输入addr和输出data的 dtype 进行校验两者均必须是 float32。产品支持情况NpuClearFloatStatus 的产品支持矩阵如下来源README产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品×Atlas A2 训练系列产品 / Atlas A2 推理系列产品×Atlas 200I/500 A2 推理产品√Atlas 推理系列产品√Atlas 训练系列产品√从源码结构看该算子仅注册了ascend950一种 AICore 配置npu_clear_float_status_def.cpp并在op_host/arch35与op_kernel/arch35目录下提供了 arch35 架构的 Tiling 与 Kernel 实现这与“Ascend 950PR / Ascend 950DT 支持”的产品定位一致Atlas A3/A2 系列不支持该算子。图编译期行为Shape、Format 与 DataType 推导在 GEGraph Engine图编译阶段算子的输出形状、格式与数据类型由三段独立的注册逻辑共同决定常量定义公共常量位于 npu_clear_float_status_common.hnamespace NpuCfs { constexpr int64_t OUTPUT_DIM 8; // 输出固定 8 个元素 constexpr size_t EXPECTED_DIM_NUM 1; // 输出为 1 维 constexpr int32_t ADDR_IDX 0; // input addr 索引 constexpr int32_t DATA_IDX 0; // output data 索引 }InferShape输出恒为 (8,)InferShape 实现 将输出 shape 的维度数固定为 1、第 0 维固定为 8完全不依赖输入 addr 的 shape——这与 README 中“addr 仅作占位符”的语义保持一致。InferFormat输出恒为 NDInferFormat 实现 将输出的原始格式与存储格式均固定为ge::FORMAT_ND。InferDataType输出恒为 float32输出数据类型推导位于 npu_clear_float_status_graph_infer.cpp固定将输出 data 的 dtype 设置为ge::DT_FLOAT与输入 addr 的 dtype 无关。源码注释中特别说明addr 是地址占位符而非数据 Tensor若输出跟随输入 dtype输入非 float32 时会导致图编译阶段下游 dtype 推导异常。因此输出 dtype 必须显式固定。值得注意的架构细节是本算子的 InferShape/InferFormat 与 InferDataType 分别注册在两个文件中前者通过IMPL_OP_INFERSHAPE注册npu_clear_float_status_infershape.cpp后者通过IMPL_OP注册npu_clear_float_status_graph_infer.cpp。Tiling 阶段全核启动与系统 WorkspaceTiling切分/编排逻辑位于 npu_clear_float_status_tiling_arch35.cpp其执行流程如下输入校验校验 addr/data 的 dtype 均为 float32。平台信息获取通过PlatformAscendC获取 AI CoreAIV核数coreNum与 UB 内存大小ubSize并做非零防御性检查。Tiling 计算needCoreNum coreNum即全核启动所有 AI Core 都必须执行因为要清除的是每个 AI Core 的溢出状态寄存器。在 ComputeTiling 中还对coreNum是否超过INT32_MAX做了防御性检查避免 int64→int32 强转溢出。Workspace 申请本算子无需用户 workspace仅申请系统 workspace通过GetLibApiWorkSpaceSize获取见 GetWorkspaceSize。设置 BlockDimSetBlockDim(needCoreNum)保证所有 AI Core 被调度执行。UB 配置采用 DCACHE_SIZE 128KB默认值、STATIC_UB_ESTIMATE 0 的策略要求动态 UB 池ubSize − 128KB不小于 TBuf 所需容量VEC_DUP_SIZE * sizeof(half) 38400 × 2 75KB并通过SetLocalMemorySize设置本地内存ApplyUbConfig。TilingKey本算子单 dtype 单场景固定使用NPU_CLEAR_FLOAT_STATUS_SCH_MODE_DEFAULT值为 0定义见 npu_clear_float_status_tiling_key.h。Tiling 数据结构的全部内容仅为一个字段npu_clear_float_status_tiling_data.hstruct NPUClearFloatStatusTilingData { int32_t needCoreNum 0; // 需要启动的核数 物理核数 };Kernel 实现原理向量指令触发状态清除 SIMT 写零Kernel 入口npu_clear_float_status.cpp通过模板参数schMode实例化默认调度模式调用NsNpuClearFloatStatus::ProcessDTYPE_ADDR(addr, data, tilingData)。核心实现位于 npu_clear_float_status_simt.h分为两个关键步骤步骤一向量单元写操作清除溢出状态标志由于 ascend950 平台不支持set_overflow_status接口实现上采用“间接触发”的方式在标量作用域内申请TBufQuePosition::VECCALCfloat16[38400]约 75KB然后连续执行 6 次Duplicate(dataUbInput, value, VEC_DUP_SIZE)value 依次取 3~8。填充值本身无实际意义目的是触发向量计算单元执行从而清除浮点溢出状态标志对应代码注释见 npu_clear_float_status_simt.h。步骤二SIMT VF 写 8 个 float32 零值到输出通过asc_vf_callOpNPUClearFloatStatusSimtT(dim3(THREAD_NUM), OUTPUT_SIZE, outputGm)启动 128 线程的 SIMT 向量函数采用Grid-Stride 循环向输出 GM 写入totalElements8个零值for (uint32_t idx blockIdx.x * blockDim.x threadIdx.x; idx static_castuint32_t(totalElements); idx blockDim.x * gridDim.x) { output[idx] static_castT(0); }由于totalElements8远小于 blockDim × gridDim实际仅 core 0 的前 8 个线程执行写入其余线程在循环条件判断后立即退出。代码对totalElements 0做了防御性检查避免负数经static_castuint32_t转为巨大无符号数导致循环越界。SIMT VF 写 GM 后由框架自动保证 cache 一致性无需显式DataCacheCleanAndInvalid。从源码结构可以推断输入addr在 Kernel 中被显式忽略(void)addr;见 npu_clear_float_status_simt.h这再次印证了 README 中“addr 数据内容不参与计算”的语义。图模式调用与验证调用方式README 中明确本算子支持图模式GE 图模式调用样例为 test_geir_npu_clear_float_status.cpp完整的算子编译与验证流程参见 算子调用指南。快速体验基于项目 build.sh根据 算子调用指南可无需搭建调用工程直接基于项目脚本执行样例# 基于 ops-nn 包执行图模式样例graph 模式无需指定 pkg_mode 和 vendor_name bash build.sh --run_example npu_clear_float_status graph参数说明${op}算子名小写下划线形式此处为npu_clear_float_status${mode}调用方式graph表示图模式调用${soc_version}可选NPU 型号设置为ascend950时会额外运行arch35目录下的示例文件${simulator}可选仿真模式目前仅支持 eageraclnn 调用场景。GE 图模式调用要点从 test_geir_npu_clear_float_status.cpp 可以看到图模式调用的完整骨架其关键步骤如下初始化 GE通过ge::GEInitialize(global_options)初始化其中{ge.exec.deviceId: 0, ge.graphRunMode: 0, ge.exec.precision_mode: must_keep_origin_dtype}。创建算子实例auto add1 op::NPUClearFloatStatus(add1)算子在图中无属性。构图使用Data占位算子接入输入addrshape 固定 (8,)输出data的 shape 固定为{8}与输入 addr 的 shape 无关对应 CreateOppInGraph。创建 Session 并运行session-AddGraph(graphId, graph, graphOptions)添加图session-RunGraph(graphId, input, output)执行最后GEFinalize()释放资源。结果检查样例运行后通过PrintReport输出各 Case 的 Build / RunGraph / OutputExists 状态。样例中还演示了静态 shapeS 模式与动态 shapeD 模式graph_shape 用 -1 占位两种构图形制并覆盖FP32 fixed_8的 dtype/shape 组合矩阵可供开发者在集成时参考。测试与 Golden 验证仓库为该算子提供了两层验证UT 单元测试test_npu_clear_float_status_infershape.cpp验证 InferShape 函数在输入 addr shape 为 (8,) 时返回成功。test_npu_clear_float_status_tiling.cpp以UB_SIZE: 262144, CORE_NUM: 64的平台信息构造 TilingContext断言 TilingKey 为 0、needCoreNum 64、BlockDim 64并校验系统 workspace 大小。Golden 脚本golden.py 定义了 Kernel 与 GEIR 共用的 Golden 函数输出恒为tf.zeros([8], dtypetf.float32)并声明精度标准为binary_equal输出为比特级精确的零属非计算型算子同时提供 TensorFlow 第三方实现用于 GEIR 交叉校验由于输出恒为零结果确定且与输入数据无关。总结NpuClearFloatStatus 是 CANN ops-nn 中一类典型的“硬件状态管理型”算子数学上它只是输出 8 个 float32 零值但真正的作用在于通过向量计算单元触发执行来清除每个 AI Core 的浮点溢出状态寄存器从而为后续的浮点溢出检测提供干净的状态起点。其设计要点可归纳为固定语义输出恒为zeros(8, float32)输入 addr 仅为占位符全核调度Tiling 将 BlockDim 设为物理核数确保所有 AI Core 的状态寄存器都被清除明确的约束与校验addr/data 必须为 float32、ND 格式Shape/Format/DataType 推导均在编译期固定无需依赖输入内容产品限定仅支持 Ascend 950PR/950DT、Atlas 200I/500 A2 推理产品、Atlas 推理/训练系列产品不支持 Atlas A3/A2 系列。开发者可参考 test_geir_npu_clear_float_status.cpp 在业务图中集成该算子并结合 npu_get_float_status 实现“清除 → 计算 → 查询”的浮点溢出监控闭环。【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考