Thanks for sending an requirement! Please fill in the following template to help quickly solve your problem.

Backgroud（背景信息）

长尾算子开发，NPUGetFloatStatus算子新增950 SIMT实现

Origin（信息来源）

vector

Benefit / Necessity （价值/作用）

NPUGetFloatStatus 与 NPUAllocFloatStatus 配套，用于混合精度训练（AMP）loss scaling 场景读取 NPU 浮点溢出状态。torch_npu 等三方框架迁移 Ascend 950 时依赖该控制流算子补齐算子清单，缺失会导致 AMP 溢出检测图构建失败。本仓交付 950 实现（暂不支持溢出探测的空实现），保证图编译与执行链路可用，语义与 canndev 对齐（无溢出时 addr 不被修改、data 恒为全零）；其余产品（Atlas 训练/推理系列、200I/500 A2）由 canndev 既有实现覆盖。

Design（设计方案）

算子定义与 Shape 推导

• op_graph/npu_get_float_status_proto.h：REG_OP(NPUGetFloatStatus)，复用 canndev 原型（npu_loss_scale_ops.h:122），1 输入 addr（TensorType{DT_FLOAT}）/ 1 输出 data（TensorType{DT_FLOAT}），无属性。通过 GEIR 图模式调用，无 aclnn 接口（CMake 声明 aclnn_exclude）。
• op_host/npu_get_float_status_def.cpp：class NPUGetFloatStatus : public OpDef，输入输出 DataType 收窄为 {ge::DT_FLOAT}、Format 声明 FORMAT_ND、全部 AutoContiguous()。AddConfig("ascend950", ...) 开启动态 shape/rank（DynamicRankSupportFlag/DynamicShapeSupportFlag），NeedCheckSupportFlag(false)，opFile.value=npu_get_float_status_apt。
• op_graph/npu_get_float_status_graph_infer.cpp：Dtype 推导 data.dtype = addr.dtype。
• op_host/npu_get_float_status_infershape.cpp：IMPL_OP_INFERSHAPE，输出 shape 固定 {8}，不依赖输入 dim 值，unknown rank/unknown shape 输入下推导天然正确，无需显式短路。

Host 端 Tiling 实现

op_host/arch35/npu_get_float_status_tiling.cpp 中 NPUGetFloatStatusTilingFunc：

• GetCoreNumAiv 获取 AIV 核数并校验 >0（内部输入，OP_LOGE）。
• 校验输入 addr shape size==8（OP_LOGE_FOR_INVALID_SHAPESIZE）与 dtype==DT_FLOAT（OP_LOGE_FOR_INVALID_DTYPE），外部输入统一 FOR 系列日志，与 inferShape 无输入校验形成分工（输出不依赖输入）。
• SetBlockDim(coreNum)；分配系统 workspace（GetLibApiWorkSpaceSize，950 为 16MB）；memset_s 清零 TilingData；SetTilingKey(DEFAULT=0)。整体对齐 canndev Tiling4NPUGetFloatStatus（SetBlockDim(core_num)），950 侧额外补齐输入校验。

SIMT 设备端 Kernel

• op_kernel/arch35/npu_get_float_status_simt.h：950 空实现。VF kernel NPUGetFloatStatusSimt 采用 32 线程（constexpr THREAD_NUM=32，launch_bounds 同一常量），Grid-Stride 循环（tid += blockDim.x，线程 0-7 有效）将 dataGm[0:7] 写 0.0。不调用 get_overflow_status()（溢出状态恒按 0 处理）、不修改输入 addr（无写回 side effect），全核幂等写入。get_overflow_status() 为 scalar-scope 编译器内建函数、不能在 simt_vf 内调用——后续支持溢出探测时应在 Process（scalar 作用域）读取并将结果作为标量参数传入 VF，addr 形参即为此预留。
• Process：scalar 作用域入口，asc_vf_call<NPUGetFloatStatusSimt>(dim3(THREAD_NUM), dataGm)，C 风格 API。
• op_kernel/npu_get_float_status_apt.cpp 为入口：REGISTER_TILING_DEFAULT + GET_TILING_DATA_WITH_STRUCT，TilingData 为空，直接调用 NsNPUGetFloatStatus::Process(addr, data)，schMode 恒为 DEFAULT(0) 无分支分发。

Tiling 数据结构与编译配置

• op_kernel/arch35/npu_get_float_status_tiling_data.h：NPUGetFloatStatusTilingData 为空结构体（输入输出 shape 固定 [8] 编译期已知，无数据面切分参数，BlockDim 经 SetBlockDim 直传）。
• op_kernel/arch35/npu_get_float_status_tiling_key.h：ASCENDC_TPL_ARGS_DECL/ASCENDC_TPL_SEL 声明唯一场景 NPU_GET_FLOAT_STATUS_TPL_DEFAULT(0)（schMode 1 bit），AIV_ONLY，dtype 不入 TilingKey。
• op_host/config/ascend950/npu_get_float_status_binary.json：dtype/format 与二进制映射配置（shape -2 动态、FormatAgnostic）。

关键设计决策与精度

• 空实现语义：canndev 原版（tik 实现）读取 0x40000 状态寄存器，溢出时将 addr 8 元素写 1.0；950 硬件暂不支持溢出探测，status 恒为 0 → addr 不被修改、data 恒全零，与"无溢出"行为一致，AMP loss scaling 在 950 上退化为不跳步（可接受降级，非错误结果）。
• 确定性：无原子操作、无跨核同步，多核写同一常量 0.0 幂等，天然确定性。
• 注释一致性：simt.h/apt.cpp/def.cpp 注释均明确空实现语义（不读寄存器、不写 addr、data 全零），避免误导后续维护。
• 精度标准：输出固定全零，float32 tolerance cross_check（固定全零下实际 bit-exact）。

Golden 参考实现与单元测试

• tests/assets/golden.py：spec 注册 NPUGetFloatStatusTestSpec；golden 返回固定 [8] float32 全零（与输入值和硬件状态无关）；CPU 环境无法读取 NPU 硬件溢出寄存器，addr side-effect 不在 golden 比对范围；ThirdPartyImpl（torch provider）独立返回全零供 GEIR 远端派发；tolerance float32 cross_check。
• tests/ut/op_host/test_npu_get_float_status_infershape.cpp：覆盖输出 shape 固定 {8} 推导（Ascend950 平台）。
• tests/ut/op_host/arch35/test_npu_get_float_status_tiling.cpp：覆盖 Ascend950 tiling 成功路径（tilingKey==0、16MB 系统 workspace、float32 dtype）。
• examples/test_geir_npu_get_float_status.cpp：GEIR 图模式 S/D（静态/动态 shape）双模式调用验证，float 输入按浮点值生成、Tensor 深拷贝构造无泄漏。