变更摘要

此 PR 为 BNInferGrad 算子新增 Ascend950 平台的 AscendC 实现，补充了从算子定义、形状推导、Tiling 到 Kernel 计算的完整链路。算子计算逻辑为 x_backprop = grads * scale / sqrt(batch_variance + epsilon)，其中 scale 和 batch_variance 沿通道维广播。Kernel 采用 Broadcast 范式，通过 NDDMA 随路广播和 TBuf 管理实现多核 tile 切分，支持 fp16/fp32/bf16 三种数据类型，并覆盖 rank≤4 和 rank=5 两种 TilingKey 场景。

主要改动

• 算子定义与注册：新增 bn_infer_grad_def.cpp、bn_infer_grad_proto.h，定义 BnInferGrad 算子的 3 输入（grads/scale/batch_variance）、1 输出（x_backprop）和一个 epsilon 属性，grads 与 x_backprop 支持 DT_FLOAT16/DT_FLOAT/DT_BF16，scale 与 batch_variance 固定为 DT_FLOAT

• 形状与数据类型推导：新增 bn_infer_grad_infershape.cpp，实现 InferShape4BnInferGrad 和 InferDataType4BnInferGrad，使输出 x_backprop 的 shape 和 dtype 完全跟随输入 grads

• Tiling 框架与 UB 切分策略：新增 bn_infer_grad_tiling_data.h/tiling_key.h/tiling_capacity.h/tiling_arch35.h/tiling_arch35.cpp，实现 FindSplitAxis（C 维禁止作为切分轴但计入 inner 累乘）、PadAndSqueeze（广播维归一化）、MultiCoreSplit（多核 tile 分配）及容量计算，TilingKey 按 RANK_4/RANK_8 区分，与数据类型无关

• AscendC Kernel 实现：新增 bn_infer_grad.h 中的 BnInferGradKernel<T, RANK> 模板类，FP32 路径直接运算，FP16/BF16 路径先 Cast 到 FP32 再计算最后 Cast 回原类型；通过 CopyInBrc 借助 loopSrcStride=0 实现 NDDMA 随路广播，CopyOutOne 采用 DataCopyPad 搬出结果

• Kernel 入口与 TilingKey 分发：新增 bn_infer_grad.cpp，作为 Ascend950 唯一 __global__ 入口，根据 RANK 模板参数选择 TilingData4 或 TilingData8，实例化对应的 BnInferGradKernel

代码审查

经过对全部 16 个变更文件的逐文件详细审查，包括对以下维度的重点排查：

• Ascend C 整数溢出与类型收窄：所有乘法链均在 int64_t 域完成，NPU 64GB 显存约束下无溢出风险；FindSplitAxis、MultiCoreSplit、stride 计算等关键路径均已保护。
• 内存与 GM 偏移：GM 偏移全部使用 int64_t；workspace 固定 16MB 系统预留，无动态溢出风险。
• 并发与流水线同步：SetFlag/WaitFlag（MTE2↔V↔MTE3↔MTE2）逐级成对，无遗漏；空 tensor 路径通过 start>=end 早退安全。
• 算子接口 / Tiling：TilingKey 仅含 RANK，无历史字段顺序问题；epsilon 从 GetAttrs() 正确获取；shape/dtype 通过 GetInputShape/GetInputDesc 取得。
• Pipeline 槽位分配：FP32 路径用 UB0-UB4，FP16/BF16 路径同样分配 5 个 TBuf，无槽位冲突。
• 边界值：a_o、per_buf_elems、count 等关键变量在非空 tensor 路径上均正值且无除零风险。
• 测试文件：FloatToHalfBits / FloatToBf16Bits 转换正确，GE API 返回值均已校验，无资源泄漏。

未发现正确性、安全性、可靠性或破坏性变更方面的问题。

关闭摘要：

整体风险判定：低风险。代码结构清晰，遵循 AscendC Broadcast 范式与参考算子实现，流水线同步正确，tiling 与 kernel 数据流对齐，无发现性缺陷。

逐文件审查确认：

• docs/zh/op_list.md — 无问题
• norm/bn_infer_grad/CMakeLists.txt — 无问题
• norm/bn_infer_grad/README.md — 无问题
• norm/bn_infer_grad/examples/arch35/test_geir_bn_infer_grad.cpp — 无问题
• norm/bn_infer_grad/examples/test_geir_bn_infer_grad.cpp — 无问题
• norm/bn_infer_grad/op_graph/bn_infer_grad_proto.h — 无问题
• norm/bn_infer_grad/op_graph/fusion_pass/.geikeep — 无问题（空文件）
• norm/bn_infer_grad/op_host/arch35/bn_infer_grad_tiling_arch35.cpp — 无问题
• norm/bn_infer_grad/op_host/arch35/bn_infer_grad_tiling_arch35.h — 无问题
• norm/bn_infer_grad/op_host/arch35/bn_infer_grad_tiling_capacity.h — 无问题
• norm/bn_infer_grad/op_host/bn_infer_grad_def.cpp — 无问题
• norm/bn_infer_grad/op_host/bn_infer_grad_infershape.cpp — 无问题
• norm/bn_infer_grad/op_kernel/arch35/bn_infer_grad.h — 无问题
• norm/bn_infer_grad/op_kernel/arch35/bn_infer_grad_tiling_data.h — 无问题
• norm/bn_infer_grad/op_kernel/arch35/bn_infer_grad_tiling_key.h — 无问题
• norm/bn_infer_grad/op_kernel/bn_infer_grad.cpp — 无问题

⚠️ 已识别出整体风险，但无法提取行内评论，请参考整体评估。

Pull Request 已合并或已关闭。

If you want to solve this problem, you can click here to do it in the FAQs.