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

Backgroud（背景信息）

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

Origin（信息来源）

vector

Benefit / Necessity （价值/作用）

Design（设计方案）

算子定义与 Shape 推导

• op_graph/max_pool_v3_grad_proto.h：REG_OP(MaxPoolV3Grad)，复用 canndev 原型，3 输入（RealNumberType）/ 1 输出，7 个属性（ksize、strides 必选；padding_mode、pads、data_format、global_pooling、ceil_mode 可选）。通过 GEIR 图模式调用，无 aclnn 接口。
• op_host/max_pool_v3_grad_def.cpp：class MaxPoolV3Grad : public OpDef，输入输出 DataType 收窄为 {DT_FLOAT16, DT_FLOAT}、Format 声明 FORMAT_ND，全部 AutoContiguous()（GE 框架自动转连续，与前向 MaxPoolV3 一致）。属性声明顺序与 tiling/infershape 的 GetAttrPointer<T>(index) 索引严格一致（0~6）。AddConfig("ascend950", ...) 开启动态 shape/rank、PrecisionReduce。
• op_graph/max_pool_v3_grad_graph_infer.cpp：Dtype 推导 out_grad.dtype = orig_input.dtype。
• op_host/max_pool_v3_grad_infershape.cpp：IMPL_OP_INFERSHAPE，强制 4D 校验（GetDimNum()==4），逐项校验 dtype 一致性/支持 dtype/data_format/padding_mode/global_pooling(必须 false)/ceil_mode(SAME|VALID 下必须 false)/ksize·strides 长度及 N·C 维==1 与 H·W∈[1,255]·[1,63]/pads(CALCULATED 下 ≥0)/grad.shape==orig_output.shape，推导 out_grad.shape = orig_input.shape。与 tiling 侧形成双重校验。

Host 端 Tiling 实现

op_host/arch35/max_pool_v3_grad_tiling.cpp 中 MaxPoolV3GradTilingFunc：

1. GetPlatformInfo 获取 UB 大小与 AIV 核数。
2. 强制 orig_input 4D，校验 ksize/strides/pads 长度==4，ValidateAttrs 按 data_format 解析 H/W 维做范围与 N/C 维==1 校验（对齐 canndev check_param）。
3. 按 data_format 提取 N/C/H/W 与 kh/kw/sh/sw；ComputePoolParams 严格对齐 Golden _compute_pool_params 计算 VALID/SAME/CALCULATED（含 ceil_mode）三种模式下的 padTop/padLeft 与 Ho/Wo，并用 FloorDiv（对齐 Python // 向负无穷取整）覆盖空 tensor 场景，Ho/Wo 钳位 max(.,0)。
4. 校验 grad.shape==orig_output.shape（4D 逐维）。
5. 重叠判定 overlap = (sh<kh)||(sw<kw)，设 overlapMode。
6. 核切分：totalOutputPos=N*C*Ho*Wo，物理核均分→下限保护 1024→对齐 32→反推 needCoreNum；校验不超 INT32_MAX。
7. Workspace：重叠模式分配 totalInputElements*sizeof(float)（FP32 累加缓冲），加系统 workspace。
8. FillTilingData 写入 16 字段结构体；SetBlockDim/SetScheduleMode(1)（两条路径均需 SyncAll）/SetTilingKey（仅编码 overlapMode，dtype 由 DTYPE_ 宏实例化）/SetLocalMemorySize(ubSize - 128KB DCache)。

SIMT 设备端 Kernel

op_kernel/arch35/max_pool_v3_grad_simt.h 采用 512 线程（constexpr THREAD_NUM=512，__launch_bounds__ 同一常量）的统一 SIMT 架构，纯 GM 计算模式依赖 128KB DCache 加速随机访问，通过 if constexpr (OVERLAP_MODE) 编译期分发非重叠/重叠路径：

• ProcessOneOutputPos：用 int64_t outPos 索引，分解 nC/hoWo/ho/wo/n/c，ComputeOutAddr/ComputeInAddr 按 inputFormat 区分 NCHW（outAddr=outPos，inAddr=nCHW+hiW+wi）与 NHWC（C 交错）；维护 matched 布尔标志按 (ki,kj) 行优先扫描，首匹配位置经 RouteGrad 路由梯度，越界位置直接 continue（等价 pad_value=-65500 方案且更高效）。== 运算符天然对齐 IEEE 754（NaN 不传播、Inf 传播、+0==-0）。
• RouteGrad：非重叠直接 outGrad[inAddr]=grad[outAddr]；重叠将 grad 提升为 FP32 后 asc_atomic_add(&wsGrad[inAddr], gradVal)。
• MaxPoolV3GradSimt：Grid-Stride 循环逐输出位置处理。
• InitPass（末尾 __builtin_cce_dcci 刷新 VF DCache）/ ConvertPass（FP32→输出 dtype，half 走 __float2half）。

op_kernel/arch35/max_pool_v3_grad.cpp 为入口：REGISTER_TILING_DEFAULT + GET_TILING_DATA_WITH_STRUCT，通过 AscendC::GetUserWorkspace(workspace) 获取用户区指针（跳过 ascend950 16MB 系统保留区，避免 overlap 路径踩踏），按 schMode 分发 Process<DTYPE_ORIG_INPUT, 0/1>。

双路径调度：非重叠 = InitPass(out_grad)→SyncAll→直接写；重叠 = InitPass(ws FP32)→SyncAll→atomicAdd ws→SyncAll→ConvertPass(ws→out_grad)。

Tiling 数据结构与编译配置

• op_kernel/arch35/max_pool_v3_grad_tiling_data.h：MaxPoolV3GradTilingData 共 16 字段（totalOutputPos、totalInputElements 为 int64_t；needCoreNum/HoWo/Wo/C/HW/W/inputFormat/kh/kw/sh/sw/padTop/padLeft 为 int64_t, overlapMode为int32_t），不含 threadNum（编译期常量）。
• op_kernel/arch35/max_pool_v3_grad_tiling_key.h：ASCENDC_TPL_ARGS_DECL/ASCENDC_TPL_SEL 声明 MAX_POOL_V3_GRAD_TPL_MODE_NON_OVERLAP(0) / MAX_POOL_V3_GRAD_TPL_MODE_OVERLAP(1) 两种场景模式，AIV_ONLY，dtype 不入 TilingKey。
• op_host/config/ascend950/max_pool_v3_grad_binary.json：dtype/format 与二进制映射配置。

关键设计决策与精度

• first-wins：与 CANN TBE / TensorFlow 一致（PyTorch all-wins 不采用）；matched 标志替代 TBE 的 mask_or/mask_not。
• 双路径：非重叠（stride≥kernel）直接写、确定；重叠（stride<kernel）atomicAdd、不保证累加确定性（与 canndev 对齐）。
• FP32 累加：重叠路径 FP16 不能直接 atomicAdd，经 FP32 workspace 中转（dtype 隔离 + GetUserWorkspace 指针正确性 + __builtin_cce_dcci DCache 一致性三重必要）。
• 精度标准：fp16 最大相对误差 <1e-3，fp32 <1e-5。

Golden 参考实现与单元测试

• tests/assets/golden.py：以 golden_max_pool_v3_grad_direct（纯 numpy，first-wins，完整支持 VALID/SAME/CALCULATED+ceil_mode，中间提升 float32 累加，Ho/Wo 钳位 0）为默认验收标杆；另提供 golden_max_pool_v3_grad_tf（tf.raw_ops.MaxPoolGrad，仅 SAME/VALID+ceil_mode=False 交叉验证）；TTK 入口 max_pool_v3_grad_golden。
• tests/ut/op_host/test_max_pool_v3_grad_infershape.cpp：覆盖 dtype 一致性、4D 校验、ksize/strides 约束、grad shape 校验等场景。
• tests/ut/op_host/arch35/test_max_pool_v3_grad_tiling.cpp：覆盖 NCHW/NHWC、VALID/SAME/CALCULATED、ceil_mode、重叠/非重叠、空 tensor、无效 dtype 等场景。