add avg_pool_grad-ops-nn-AtomGit

Apooling/avg_pool_grad/CMakeLists.txt+14-0

@@ -0,0 +1,14 @@
	1	+# ----------------------------------------------------------------------------
	2	+# This program is free software, you can redistribute it and/or modify.
	3	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+# This file is a part of the CANN Open Software.
	5	+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
	6	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	W W wangyongguang3月30日补充PR描述
	9	+# ----------------------------------------------------------------------------
	10	+# 设置算子定义时支持的芯片类型
	11	+set(SUPPORT_COMPUTE_UNIT "ascend950")
	12	+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
	13	+set(SUPPORT_TILING_DIR "arch35")
	14	+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} OPTYPE avg_pool_grad ACLNNTYPE aclnn_exclude DEPENDENCIES avg_pool_v2_grad)

Apooling/avg_pool_grad/op_graph/avg_pool_grad_proto.h+65-0

@@ -0,0 +1,65 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#ifndef OP_PROTO_AVG_POOL_GRAD_H_
	12	+#define OP_PROTO_AVG_POOL_GRAD_H_
	13	+
	14	+#include "graph/operator_reg.h"
	15	+#include "graph/operator.h"
	16	+
	17	+namespace ge {
	18	+
	19	+/**
	20	+* @brief Computes avgpoolgrad function.
	21	+
	22	+* @par Inputs:
	23	+* @li orig_input_shape: An one-dim tensor of type int32, which describes the
	24	+* original input shape [N,C,H,W] or [N,H,W,C] of forward AvgPool.
	25	+* @li input_grad: An NHWC or NCHW tensor of type float16, float32, double or bfloat16. \n
	26	+
	27	+* @par Attributes:
	28	+* @li ksize: A required tuple or list of ints,
	29	+* specifying the size of the window for each dimension of the input tensor.
	30	+* For Ascend 950PR/Ascend 950DT AI Processor: "ksize" length is 1, 2 or 4, must be greater than 0. \n
	31	+* @li strides: A required tuple or list of ints,
	32	+* specifying the stride of the sliding window for each dimension of the input tensor.
	33	+* For Ascend 950PR/Ascend 950DT AI Processor: "strides" length is 1, 2 or 4, must be greater than 0. \n
	34	+* @li padding: An optional string, specifying the type of the padding algorithm to use,
	35	+* either "VALID", "SAME".
	36	+* With "SAME" means that the outputs will have the same spatial dimensions as its inputs.
	37	+* With "VALID" means no padding.
	38	+* @li data_format: An optional string. Defaults to "NHWC". \n
	39	+* For Ascend 950PR/Ascend 950DT AI Processor: support "NCHW" or "NHWC". \n
	S S sunday3月30日问题: 注释中声明 data_format 默认值为 "NCHW"，但当前注册代码默认值为 "NHWC"，存在文档与实现不一致。建议: 统一注释与实现，建议将注释默认值修正为 "NHWC"（或按规范同步调整实现）。
	40	+
	41	+* @par Outputs:
	42	+* @li out_grad: A mutable tensor with the same shape as "orig_input_shape" and the same type as "input_grad". \n
	43	+*\n
	44	+* input_grad_height = (out_grad_height + pads_top + pads_bottom - ksize_height)
	45	+* / strides_h + 1
	46	+*\n
	47	+* input_grad_width = (out_grad_width + pads_left + pads_right - ksize_width)
	48	+* / strides_w + 1
	49	+*\n
	50	+
	51	+* @par Third-party framework compatibility
	52	+* @li Compatible with the TensorFlow operator AvgPoolGrad.
	53	+*/
	54	+REG_OP(AvgPoolGrad)
	55	+ .INPUT(orig_input_shape, TensorType({DT_INT32}))
	56	+ .INPUT(input_grad, TensorType({DT_FLOAT16, DT_FLOAT32, DT_DOUBLE, DT_BF16}))
	57	+ .OUTPUT(out_grad, TensorType({DT_FLOAT16, DT_FLOAT32, DT_DOUBLE, DT_BF16}))
	58	+ .REQUIRED_ATTR(ksize, ListInt)
	59	+ .REQUIRED_ATTR(strides, ListInt)
	60	+ .REQUIRED_ATTR(padding, String)
	61	+ .ATTR(data_format, String, "NHWC")
	62	+ .OP_END_FACTORY_REG(AvgPoolGrad)
	63	+
	64	+} // namespace ge
	65	+#endif // OP_PROTO_AVG_POOL_GRAD_H_

Apooling/avg_pool_grad/op_host/arch35/avg_pool_grad_nchw_tiling.cpp+31-0

@@ -0,0 +1,31 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file avg_pool_grad_nchw_tiling.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+#include "avg_pool_grad_nchw_tiling.h"
	17	+
	18	+namespace optiling {
	19	+
	20	+//////////////////////////////// AvgPoolGradNCHWTiling /////////////////////////////////
	21	+ge::graphStatus AvgPoolGradNCHWTiling::GetPlatformInfo() {
	22	+ return GetAvgPoolGradPlatformInfo(context_, ubSize, coreNum);
	23	+}
	24	+
	25	+ge::graphStatus AvgPoolGradNCHWTiling::GetShapeAttrsInfo() {
	26	+ return GetAvgPoolGradShapeAttrsInfo(context_, inputData);
	27	+}
	28	+
	29	+REGISTER_TILING_TEMPLATE("AvgPoolGrad", AvgPoolGradNCHWTiling, 0);
	30	+
	31	+} // namespace optiling

Apooling/avg_pool_grad/op_host/arch35/avg_pool_grad_nchw_tiling.h+38-0

@@ -0,0 +1,38 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file avg_pool_grad_nchw_tiling.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#ifndef AVG_POOL_GRAD_NCHW_TILING_H_
	17	+#define AVG_POOL_GRAD_NCHW_TILING_H_
	18	+
	19	+#include "avg_pool_grad_tiling_base.h"
	20	+#include "../../../avg_pool_v2_grad/op_host/arch35/avg_pool_v2_grad_nchw_tiling.h"
	21	+
	22	+namespace optiling {
	23	+
	24	+class AvgPoolGradNCHWTiling : public AvgPoolV2GradCommonNCHWTiling {
	25	+public:
	26	+ explicit AvgPoolGradNCHWTiling(gert::TilingContext* context) : AvgPoolV2GradCommonNCHWTiling(context)
	27	+ {}
	28	+ ~AvgPoolGradNCHWTiling() override
	29	+ {}
	30	+
	31	+private:
	32	+ ge::graphStatus GetPlatformInfo() override;
	33	+ ge::graphStatus GetShapeAttrsInfo() override;
	34	+};
	35	+
	36	+} // namespace optiling
	37	+
	38	+#endif

Apooling/avg_pool_grad/op_host/arch35/avg_pool_grad_nhwc_tiling.cpp+31-0

@@ -0,0 +1,31 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file avg_pool_grad_nhwc_tiling.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+#include "avg_pool_grad_nhwc_tiling.h"
	17	+
	18	+namespace optiling
	19	+{
	20	+
	21	+ge::graphStatus AvgPoolGradNHWCTiling::GetPlatformInfo() {
	22	+ return GetAvgPoolGradPlatformInfo(context_, ubSize, coreNum);
	23	+}
	24	+
	25	+ge::graphStatus AvgPoolGradNHWCTiling::GetShapeAttrsInfo() {
	26	+ return GetAvgPoolGradShapeAttrsInfo(context_, inputData);
	27	+}
	28	+
	29	+REGISTER_OPS_TILING_TEMPLATE(AvgPoolGrad, AvgPoolGradNHWCTiling, 3);
	30	+
	31	+} // namespace optiling

Apooling/avg_pool_grad/op_host/arch35/avg_pool_grad_nhwc_tiling.h+39-0

@@ -0,0 +1,39 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file avg_pool_grad_nhwc_tiling.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#ifndef AVG_POOL_GRAD_NHWC_TILING_H_
	17	+#define AVG_POOL_GRAD_NHWC_TILING_H_
	18	+
	19	+#include "avg_pool_grad_tiling_base.h"
	20	+#include "../../../avg_pool_v2_grad/op_host/arch35/avg_pool_v2_grad_nhwc_tiling.h"
	21	+
	22	+namespace optiling
	23	+{
	24	+
	25	+class AvgPoolGradNHWCTiling : public AvgPoolV2GradCommonNHWCTiling {
	26	+public:
	27	+ explicit AvgPoolGradNHWCTiling(gert::TilingContext* context) : AvgPoolV2GradCommonNHWCTiling(context)
	28	+ {}
	29	+ ~AvgPoolGradNHWCTiling() override
	30	+ {}
	31	+
	32	+private:
	33	+ ge::graphStatus GetPlatformInfo() override;
	34	+ ge::graphStatus GetShapeAttrsInfo() override;
	35	+};
	36	+
	37	+} // namespace optiling
	38	+
	39	+#endif

Apooling/avg_pool_grad/op_host/arch35/avg_pool_grad_simt_tiling.cpp+36-0

@@ -0,0 +1,36 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file max_pool_with_argmax_v3_simt_tiling.cpp
	13	+ * \brief
	14	+ */
	15	+#include <cctype>
	16	+#include <algorithm>
	17	+#include "platform/platform_ascendc.h"
	18	+#include "atvoss/broadcast/broadcast_tiling.h"
	19	+#include "op_host/tiling_templates_registry.h"
	20	+#include "avg_pool_grad_simt_tiling.h"
	21	+
	22	+using namespace AscendC;
	23	+using namespace ge;
	24	+
	25	+namespace optiling {
	26	+
	27	+ge::graphStatus AvgPoolGradTilingSIMT::GetPlatformInfo() {
	28	+ return GetAvgPoolGradPlatformInfo(context_, ubSize, coreNum);
	29	+}
	30	+
	31	+ge::graphStatus AvgPoolGradTilingSIMT::GetShapeAttrsInfo() {
	32	+ return GetAvgPoolGradShapeAttrsInfo(context_, inputData);
	33	+}
	34	+
	35	+REGISTER_OPS_TILING_TEMPLATE(AvgPoolGrad, AvgPoolGradTilingSIMT, 100);
	36	+} // namespace optiling

Apooling/avg_pool_grad/op_host/arch35/avg_pool_grad_simt_tiling.h+37-0

@@ -0,0 +1,37 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file avg_pool_grad_simt_tiling.h
	13	+ * \brief simt imply foravg_pool_grad
	14	+ */
	15	+
	16	+#ifndef CANN_AVG_POOL_GRAD_SIMT_TILING_H
	17	+#define CANN_AVG_POOL_GRAD_SIMT_TILING_H
	18	+
	19	+#include "avg_pool_grad_tiling_base.h"
	20	+#include "../../../avg_pool_v2_grad/op_host/arch35/avg_pool_v2_grad_simt_tiling.h"
	21	+
	22	+namespace optiling {
	23	+
	24	+class AvgPoolGradTilingSIMT : public AvgPoolV2GradTilingSIMT {
	25	+public:
	26	+ explicit AvgPoolGradTilingSIMT(gert::TilingContext* context) : AvgPoolV2GradTilingSIMT(context)
	27	+ {}
	28	+ ~AvgPoolGradTilingSIMT() override
	29	+ {}
	30	+
	31	+protected:
	32	+ ge::graphStatus GetPlatformInfo() override;
	33	+ ge::graphStatus GetShapeAttrsInfo() override;
	34	+};
	35	+
	36	+} // namespace optiling
	37	+#endif // CANN_AVG_POOL_GRAD_SIMT_TILING_H

Apooling/avg_pool_grad/op_host/arch35/avg_pool_grad_tiling.cpp+46-0

@@ -0,0 +1,46 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file avg_pool_grad_tiling.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+#include "op_host/tiling_templates_registry.h"
	17	+#include "avg_pool_grad_tiling_base.h"
	18	+#include "error_util.h"
	19	+
	20	+namespace optiling
	21	+{
	22	+using Ops::NN::Optiling::TilingRegistry;
	23	+ge::graphStatus Tiling4AvgPoolGrad(gert::TilingContext* context)
	24	+{
	25	+ return TilingRegistry::GetInstance().DoTilingImpl(context);
	26	+}
	27	+
	28	+ge::graphStatus TilingPrepare4AvgPoolGrad(gert::TilingParseContext* context) {
	29	+ OP_LOGD("AvgPoolGrad", "TilingPrepare4AvgPoolGrad");
	30	+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
	31	+ OP_TILING_CHECK(platformInfoPtr == nullptr, CUBE_INNER_ERR_REPORT(context, "platformInfoPtr info is null"),
	32	+ return ge::GRAPH_FAILED);
	33	+ auto compileInfoPtr = context->GetCompiledInfo<AvgPoolGradCompileInfo>();
	34	+ OP_TILING_CHECK(compileInfoPtr == nullptr, CUBE_INNER_ERR_REPORT(context, "compileInfoPtr is null"),
	35	+ return ge::GRAPH_FAILED);
	36	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	37	+ compileInfoPtr->coreNum = ascendcPlatform.GetCoreNum();
	38	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfoPtr->ubSize);
	39	+ return ge::GRAPH_SUCCESS;
	40	+}
	41	+
	42	+IMPL_OP_OPTILING(AvgPoolGrad)
	43	+ .InputsDataDependency({0})
	44	+ .Tiling(Tiling4AvgPoolGrad)
	45	+ .TilingParse<AvgPoolGradCompileInfo>(TilingPrepare4AvgPoolGrad);
	46	+} // namespace optiling

Apooling/avg_pool_grad/op_host/arch35/avg_pool_grad_tiling_base.cpp+463-0

@@ -0,0 +1,463 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file avg_pool_grad_tiling_base.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+#include <cstdint>
	17	+#include "op_host/tiling_templates_registry.h"
	18	+#include "log/log.h"
	19	+#include "error_util.h"
	20	+#include "platform/platform_info.h"
	21	+#include "avg_pool_grad_tiling_base.h"
	22	+
	23	+using namespace AscendC;
	24	+using namespace ge;
	25	+
	26	+namespace optiling
	27	+{
	28	+static const int32_t KERNEL_POS = 0;
	29	+static const int32_t STRIDE_POS = 1;
	30	+static const int32_t PADDING_POS = 2;
	31	+static const int32_t FORMAT_POS = 3;
	32	+
	33	+static const int32_t AVG_POOL_GRAD_DIM_ZERO = 0;
	34	+static const int32_t AVG_POOL_GRAD_DIM_ONE = 1;
	35	+static const int32_t AVG_POOL_GRAD_DIM_TWO = 2;
	36	+static const int32_t AVG_POOL_GRAD_DIM_THREE = 3;
	37	+static const int32_t INDEX_GRAD = 1;
	38	+
	39	+static const int32_t ONE = 1;
	40	+static const int32_t TWO = 2;
	41	+constexpr size_t ORIG_INPUT_SHAPE_INDEX = 0;
	42	+
	43	+static bool IsInvalidType(const DataType dtype)
	44	+{
	45	+ const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16};
	46	+ bool dtypeInValid = (supportedDtype.count(dtype) == 0);
	47	+ return dtypeInValid;
	48	+}
	49	+
	50	+static bool IsInvalidPaddingMode(std::string padMode)
	51	+{
	52	+ const std::set<std::string> supportedPadModeList = {"SAME", "VALID"};
	53	+ bool padModeInValid = (supportedPadModeList.count(padMode) == 0);
	54	+ return padModeInValid;
	55	+}
	56	+
	57	+static inline bool IsGreaterThanInt32Max(const AvgPoolV2GradInputInfo& inputData)
	58	+{
	59	+ int64_t totalSize = inputData.batches * inputData.channels * inputData.inputShape[H_DIM] * inputData.inputShape[W_DIM];
	60	+ return totalSize > static_cast<int64_t>(INT32_MAX);
	61	+}
	62	+
	63	+static ge::graphStatus GetPadInfo(gert::TilingContext* context,
	64	+ AvgPoolV2GradInputInfo& inputData, const AvgPoolGradCommon& commInfo)
	65	+{
	66	+
	67	+ if (commInfo.padModeStr == "VALID") {
	68	+ inputData.pad = {0, 0, 0, 0}; // top, bottom, left, right
	69	+ } else if (commInfo.padModeStr == "SAME") {
	70	+ int64_t hPadNeed = std::max(int64_t{0}, (inputData.gradShape[H_DIM] - 1) * inputData.stride[H_DIM] +
	71	+ inputData.kernelSize[H_DIM] - inputData.inputShape[H_DIM]);
	72	+ int64_t topPad = hPadNeed / TWO;
	73	+ int64_t bottomPad = hPadNeed - topPad;
	74	+
	75	+ int64_t wPadNeed = std::max(int64_t{0}, (inputData.gradShape[W_DIM] - 1) * inputData.stride[W_DIM] +
	76	+ inputData.kernelSize[W_DIM] - inputData.inputShape[W_DIM]);
	77	+ int64_t leftPad = wPadNeed / TWO;
	78	+ int64_t rightPad = wPadNeed - leftPad;
	79	+
	80	+ inputData.pad = {topPad, bottomPad, leftPad, rightPad};
	81	+ } else {
	82	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "AvgPoolGrad: not support padmode %s", commInfo.padModeStr.c_str());
	83	+ return ge::GRAPH_FAILED;
	84	+ }
	85	+ return ge::GRAPH_SUCCESS;
	86	+}
	87	+
	88	+static ge::graphStatus GetStrideInfo(gert::TilingContext* context, const gert::RuntimeAttrs* runtimeAttrs,
	89	+ AvgPoolV2GradInputInfo& inputData, const AvgPoolGradCommon& commInfo)
	90	+{
	91	+ auto stride = runtimeAttrs->GetListInt(STRIDE_POS);
	92	+ OPS_CHECK_NULL_WITH_CONTEXT(context, stride);
	93	+ auto strideDim = stride->GetSize();
	94	+ OP_TILING_CHECK(strideDim != ONE_DIMS && strideDim != HW_DIMS && strideDim != NCHW_DIMS,
	95	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "AvgPoolGrad: stride must have %d, %d, or %d elements ",
	96	+ ONE_DIMS, HW_DIMS, NCHW_DIMS),
	97	+ return ge::GRAPH_FAILED);
	98	+
	99	+ int64_t hStride = ONE;
	100	+ int64_t wStride = ONE;
	101	+ if (strideDim == ONE_DIMS) {
	102	+ hStride = stride->GetData()[AVG_POOL_GRAD_DIM_ZERO];
	103	+ wStride = stride->GetData()[AVG_POOL_GRAD_DIM_ZERO];
	104	+ } else if (strideDim == HW_DIMS) {
	105	+ hStride = stride->GetData()[AVG_POOL_GRAD_DIM_ZERO];
	106	+ wStride = stride->GetData()[AVG_POOL_GRAD_DIM_ONE];
	107	+ } else if (strideDim == NCHW_DIMS) {
	108	+ hStride = stride->GetData()[commInfo.hDim];
	109	+ wStride = stride->GetData()[commInfo.wDim];
	110	+ }
	111	+ inputData.stride = {hStride, wStride};
	112	+ OP_TILING_CHECK(hStride <= 0 \|\| wStride <= 0,
	113	+ VECTOR_INNER_ERR_REPORT_TILIING(context->GetNodeName(),
	114	+ "AvgPoolGrad: The stride of the H and W dimensions should be greater than 0, not support [%ld, %ld]",
	115	+ hStride, wStride),
	116	+ return ge::GRAPH_FAILED);
	117	+ return ge::GRAPH_SUCCESS;
	118	+}
	119	+
	120	+static ge::graphStatus GetKernelKsizeInfo(gert::TilingContext* context, const gert::RuntimeAttrs* runtimeAttrs,
	121	+ AvgPoolV2GradInputInfo& inputData, const AvgPoolGradCommon& commInfo)
	122	+{
	123	+ auto kernelSize = runtimeAttrs->GetListInt(KERNEL_POS);
	124	+ OPS_CHECK_NULL_WITH_CONTEXT(context, kernelSize);
	125	+ auto kSizeDim = kernelSize->GetSize();
	126	+ OP_TILING_CHECK(
	127	+ kSizeDim != ONE_DIMS && kSizeDim != HW_DIMS && kSizeDim != NCHW_DIMS,
	128	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "AvgPoolGrad: kernel_size must have %d, %d, or %d elements ",
	129	+ ONE_DIMS, HW_DIMS, NCHW_DIMS),
	130	+ return ge::GRAPH_FAILED);
	131	+ int64_t hKernelSize = 1;
	132	+ int64_t wKernelSize = 1;
	133	+ if (kSizeDim == ONE_DIMS) {
	134	+ hKernelSize = kernelSize->GetData()[AVG_POOL_GRAD_DIM_ZERO];
	135	+ wKernelSize = kernelSize->GetData()[AVG_POOL_GRAD_DIM_ZERO];
	136	+ } else if (kSizeDim == HW_DIMS) {
	137	+ hKernelSize = kernelSize->GetData()[AVG_POOL_GRAD_DIM_ZERO];
	138	+ wKernelSize = kernelSize->GetData()[AVG_POOL_GRAD_DIM_ONE];
	139	+ } else if (kSizeDim == NCHW_DIMS) {
	140	+ hKernelSize = kernelSize->GetData()[commInfo.hDim];
	141	+ wKernelSize = kernelSize->GetData()[commInfo.wDim];
	142	+ }
	143	+ inputData.kernelSize = {hKernelSize, wKernelSize};
	144	+
	145	+ OP_TILING_CHECK(hKernelSize <= 0 \|\| wKernelSize <= 0,
	146	+ VECTOR_INNER_ERR_REPORT_TILIING(context->GetNodeName(),
	147	+ "AvgPoolGrad: The ksize of the H and W dimensions should be greater than 0, not support [%ld, %ld]",
	148	+ hKernelSize, wKernelSize),
	149	+ return ge::GRAPH_FAILED);
	150	+ return ge::GRAPH_SUCCESS;
	151	+}
	152	+
	153	+static ge::graphStatus CheckShape(gert::TilingContext* context, gert::Shape& gradShape, gert::Shape& outputShape)
	154	+{
	155	+ OP_TILING_CHECK(
	156	+ gradShape.GetDimNum() != NCHW_DIMS && gradShape.GetDimNum() != CHW_DIMS,
	157	+ VECTOR_INNER_ERR_REPORT_TILIING(context->GetNodeName(), "AvgPoolGrad: input shape dim = %zu, should be equal 3 or 4",
	158	+ gradShape.GetDimNum()),
	159	+ return ge::GRAPH_FAILED);
	160	+ OP_TILING_CHECK(
	161	+ outputShape.GetDimNum() != NCHW_DIMS && outputShape.GetDimNum() != CHW_DIMS,
	162	+ VECTOR_INNER_ERR_REPORT_TILIING(context->GetNodeName(), "AvgPoolGrad: output shape dim = %zu, should be equal 3 or 4",
	163	+ outputShape.GetDimNum()),
	164	+ return ge::GRAPH_FAILED);
	165	+ if (gradShape.GetShapeSize() == 0 && outputShape.GetShapeSize() == 0) {
	166	+ return ge::GRAPH_SUCCESS;
	167	+ }
	168	+ OP_TILING_CHECK(gradShape.GetShapeSize() <= 0,
	169	+ VECTOR_INNER_ERR_REPORT_TILIING(context->GetNodeName(),
	170	+ "AvgPoolGrad: input shape size %ld less than zero failed",
	171	+ gradShape.GetShapeSize()),
	172	+ return ge::GRAPH_FAILED);
	173	+
	174	+ OP_TILING_CHECK(outputShape.GetShapeSize() <= 0,
	175	+ VECTOR_INNER_ERR_REPORT_TILIING(context->GetNodeName(),
	176	+ "AvgPoolGrad: output shape size %ld less than zero failed",
	177	+ outputShape.GetShapeSize()),
	178	+ return ge::GRAPH_FAILED);
	179	+ return ge::GRAPH_SUCCESS;
	180	+}
	181	+
	182	+ge::graphStatus GetFormat(gert::TilingContext* context, const gert::RuntimeAttrs* runtimeAttrs, AvgPoolV2GradInputInfo& inputData)
	183	+{
	184	+ std::string inputFormatStr("NHWC");
	185	+ const char* inputFormat = runtimeAttrs->GetAttrPointer<char>(FORMAT_POS);
	186	+ if (inputFormat != nullptr) {
	187	+ inputFormatStr = inputFormat;
	188	+ }
	189	+ if (inputFormatStr == "NCHW") {
	190	+ inputData.inputFormat = ge::Format::FORMAT_NCHW;
	191	+ } else if (inputFormatStr == "NHWC") {
	192	+ inputData.inputFormat = ge::Format::FORMAT_NHWC;
	193	+ } else {
	194	+ VECTOR_INNER_ERR_REPORT_TILIING(context,
	195	+ "AvgPoolGrad: only support NCHW、NHWC, not support format %s",
	196	+ inputFormatStr.c_str());
	197	+ return ge::GRAPH_FAILED;
	198	+ }
	199	+ return ge::GRAPH_SUCCESS;
	200	+}
	201	+
	202	+ge::graphStatus CalculateShapeInfo(gert::TilingContext* context, AvgPoolV2GradInputInfo& inputData, AvgPoolGradCommon& commInfo, const int32_t* shapeValue)
	203	+{
	204	+ auto inputShape0 = context->GetInputShape(0);
	205	+ auto shapeDim = inputShape0->GetStorageShape().GetDim(0);
	206	+ if (inputData.inputFormat == ge::Format::FORMAT_NCHW) {
	207	+ if (shapeDim == CHW_DIMS) {
	208	+ commInfo.cDim = AVG_POOL_GRAD_DIM_ZERO;
	209	+ commInfo.hDim = AVG_POOL_GRAD_DIM_ONE;
	210	+ commInfo.wDim = AVG_POOL_GRAD_DIM_TWO;
	211	+ inputData.batches = shapeValue[commInfo.cDim];
	212	+ } else {
	213	+ commInfo.nDim = AVG_POOL_GRAD_DIM_ZERO;
	214	+ commInfo.cDim = AVG_POOL_GRAD_DIM_ONE;
	215	+ commInfo.hDim = AVG_POOL_GRAD_DIM_TWO;
	216	+ commInfo.wDim = AVG_POOL_GRAD_DIM_THREE;
	217	+ inputData.batches = shapeValue[commInfo.nDim] * shapeValue[commInfo.cDim];
	218	+ }
	219	+ inputData.channels = ONE;
	220	+ } else if (inputData.inputFormat == ge::Format::FORMAT_NHWC) {
	221	+ if (shapeDim == CHW_DIMS) {
	222	+ commInfo.cDim = AVG_POOL_GRAD_DIM_TWO;
	223	+ commInfo.hDim = AVG_POOL_GRAD_DIM_ZERO;
	224	+ commInfo.wDim = AVG_POOL_GRAD_DIM_ONE;
	225	+ inputData.batches = ONE;
	226	+ } else {
	227	+ commInfo.nDim = AVG_POOL_GRAD_DIM_ZERO;
	228	+ commInfo.cDim = AVG_POOL_GRAD_DIM_THREE;
	229	+ commInfo.hDim = AVG_POOL_GRAD_DIM_ONE;
	230	+ commInfo.wDim = AVG_POOL_GRAD_DIM_TWO;
	231	+ inputData.batches = shapeValue[commInfo.nDim];
	232	+ }
	233	+ inputData.channels = shapeValue[commInfo.cDim];
	234	+ } else {
	235	+ VECTOR_INNER_ERR_REPORT_TILIING(context->GetNodeName(),
	236	+ "AvgPoolGrad: only support NCHW and NHWC, not support format.");
	237	+ return ge::GRAPH_FAILED;
	238	+ }
	239	+ inputData.inputShape = {shapeValue[commInfo.hDim], shapeValue[commInfo.wDim]};
	240	+ return ge::GRAPH_SUCCESS;
	241	+}
	242	+
	243	+ge::graphStatus CheckDimConsistency(gert::TilingContext* context, const int32_t* shapeValue, const AvgPoolGradCommon& commInfo)
	244	+{
	245	+ auto inputShape0 = context->GetInputShape(0);
	246	+ auto shapeDim = inputShape0->GetStorageShape().GetDim(0);
	247	+ auto outX = context->GetOutputShape(0);
	248	+ auto outShape = EnsureNotScalar(outX->GetStorageShape());
	249	+ if (shapeDim == NCHW_DIMS) {
	250	+ OP_TILING_CHECK(
	251	+ shapeValue[commInfo.nDim] != outShape.GetDim(commInfo.nDim),
	252	+ VECTOR_INNER_ERR_REPORT_TILIING(context->GetNodeName(),
	253	+ "AvgPoolGrad: input n-dim shape value is %d, but output n-dim shape value is %ld, should be same ", shapeValue[commInfo.nDim], outShape.GetDim(commInfo.nDim)),
	254	+ return ge::GRAPH_FAILED);
	255	+ }
	256	+ OP_TILING_CHECK(
	257	+ shapeValue[commInfo.cDim] != outShape.GetDim(commInfo.cDim),
	258	+ VECTOR_INNER_ERR_REPORT_TILIING(context->GetNodeName(),
	259	+ "AvgPoolGrad: input c-dim shape value is %d, but output c-dim shape value is %ld, should be same ", shapeValue[commInfo.cDim], outShape.GetDim(commInfo.cDim)),
	260	+ return ge::GRAPH_FAILED);
	261	+ OP_TILING_CHECK(
	262	+ shapeValue[commInfo.hDim] != outShape.GetDim(commInfo.hDim),
	263	+ VECTOR_INNER_ERR_REPORT_TILIING(context->GetNodeName(),
	264	+ "AvgPoolGrad: input h-dim shape value is %d, but output h-dim shape value is %ld, should be same ", shapeValue[commInfo.hDim], outShape.GetDim(commInfo.hDim)),
	265	+ return ge::GRAPH_FAILED);
	266	+ OP_TILING_CHECK(
	267	+ shapeValue[commInfo.wDim] != outShape.GetDim(commInfo.wDim),
	268	+ VECTOR_INNER_ERR_REPORT_TILIING(context->GetNodeName(),
	269	+ "AvgPoolGrad: input w-dim shape value is %d, but output w-dim shape value is %ld, should be same ", shapeValue[commInfo.wDim], outShape.GetDim(commInfo.wDim)),
	270	+ return ge::GRAPH_FAILED);
	271	+ return ge::GRAPH_SUCCESS;
	272	+}
	273	+
	274	+static ge::graphStatus GetShapeAndDtype(gert::TilingContext* context, const gert::RuntimeAttrs* runtimeAttrs,
	275	+ AvgPoolV2GradInputInfo& inputData, AvgPoolGradCommon& commInfo)
	276	+{
	277	+ // 输入值依赖input
	278	+ auto inputShape0 = context->GetInputShape(0);
	279	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape0);
	280	+ auto shapeDim = inputShape0->GetStorageShape().GetDim(0);
	281	+ OP_TILING_CHECK(
	282	+ shapeDim != NCHW_DIMS && shapeDim != CHW_DIMS,
	283	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "inputShapeDim must be 3 or 4, shapeDim: %ld", shapeDim),
	284	+ return ge::GRAPH_FAILED);
	285	+ // input_grad
	286	+ auto inputShape1 = context->GetInputShape(INDEX_GRAD);
	287	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape1);
	288	+ auto gradShape = EnsureNotScalar(inputShape1->GetStorageShape());
	289	+ // output_grad
	290	+ auto outX = context->GetOutputShape(0);
	291	+ OPS_CHECK_NULL_WITH_CONTEXT(context, outX);
	292	+ auto outShape = EnsureNotScalar(outX->GetStorageShape());
	293	+
	294	+ auto inputDesc = context->GetInputDesc(INDEX_GRAD);
	295	+ OPS_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
	296	+
	297	+ auto dtype = inputDesc->GetDataType();
	298	+ if (IsInvalidType(dtype)) {
	299	+ VECTOR_INNER_ERR_REPORT_TILIING(context->GetNodeName(), "AvgPoolGrad: invalid dtype");
	300	+ return ge::GRAPH_FAILED;
	301	+ }
	302	+ inputData.dtypeSize = ge::GetSizeByDataType(dtype);
	303	+ OP_TILING_CHECK(
	304	+ inputData.dtypeSize <= 0,
	305	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "inputData.dtypeSize must be greater than 0, dtypeSize: %ld", inputData.dtypeSize),
	306	+ return ge::GRAPH_FAILED);
	307	+ // 校验是否是3/4维
	308	+ OP_TILING_CHECK(CheckShape(context, gradShape, outShape) != ge::GRAPH_SUCCESS,
	309	+ VECTOR_INNER_ERR_REPORT_TILIING(context->GetNodeName(), "AvgPoolGrad: check shape failed"),
	310	+ return ge::GRAPH_FAILED);
	311	+ // 值依赖转换
	312	+ const gert::Tensor* shapeTensor = context->GetInputTensor(ORIG_INPUT_SHAPE_INDEX);
	313	+ OP_CHECK_NULL_WITH_CONTEXT(context, shapeTensor);
	314	+ const int32_t* shapeValue = shapeTensor->GetData<int32_t>();
	315	+ if (shapeValue == nullptr) {
	316	+ return ge::GRAPH_FAILED;
	317	+ }
	318	+
	319	+ ge::graphStatus ret = GetFormat(context, runtimeAttrs, inputData);
	320	+ OP_TILING_CHECK(ret != ge::GRAPH_SUCCESS,
	321	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "AvgPoolGrad: get format failed"), return ret);
	322	+
	323	+ ret = CalculateShapeInfo(context, inputData, commInfo, shapeValue);
	324	+ OP_TILING_CHECK(ret != ge::GRAPH_SUCCESS,
	325	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "AvgPoolGrad: calculate shape info failed"), return ret);
	326	+
	327	+ inputData.gradShape = {gradShape.GetDim(commInfo.hDim), gradShape.GetDim(commInfo.wDim)};
	328	+ ret = CheckDimConsistency(context, shapeValue, commInfo);
	329	+ OP_TILING_CHECK(ret != ge::GRAPH_SUCCESS,
	330	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "AvgPoolGrad: check dim consistency failed"), return ret);
	331	+
	332	+ inputData.outShape = {outShape.GetDim(commInfo.hDim), outShape.GetDim(commInfo.wDim)};
	333	+ return ge::GRAPH_SUCCESS;
	334	+}
	335	+
	336	+static ge::graphStatus GetAttrsInfo(gert::TilingContext* context, const gert::RuntimeAttrs* runtimeAttrs,
	337	+ AvgPoolV2GradInputInfo& inputData, AvgPoolGradCommon& commInfo)
	338	+{
	339	+ const char* padMode = runtimeAttrs->GetAttrPointer<char>(PADDING_POS);
	340	+ OPS_CHECK_NULL_WITH_CONTEXT(context, padMode);
	341	+ commInfo.padModeStr = padMode;
	342	+ OP_TILING_CHECK(
	343	+ IsInvalidPaddingMode(commInfo.padModeStr),
	344	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "AvgPoolGrad: not support padmode %s", commInfo.padModeStr.c_str()),
	345	+ return ge::GRAPH_FAILED);
	346	+
	347	+ // tensorflow 默认值对应 exclusive = true, 故countIncludePad为false , divisorOverride = 0, globalPooling = false, ceil_mode对AvgPoolV2Grad无影响
	348	+ inputData.countIncludePad = false;
	349	+ inputData.divisorOverride = 0;
	350	+ inputData.globalPooling = false;
	351	+
	352	+ return ge::GRAPH_SUCCESS;
	353	+}
	354	+
	355	+static ge::graphStatus CheckGradShapeForValid(gert::TilingContext* context, AvgPoolV2GradInputInfo& inputData)
	356	+{
	357	+ int64_t expectedH = (inputData.inputShape[H_DIM] - inputData.kernelSize[H_DIM] + inputData.stride[H_DIM]) /
	358	+ inputData.stride[H_DIM];
	359	+ int64_t expectedW = (inputData.inputShape[W_DIM] - inputData.kernelSize[W_DIM] + inputData.stride[W_DIM]) /
	360	+ inputData.stride[W_DIM];
	361	+ if (inputData.gradShape[H_DIM] != expectedH \|\| inputData.gradShape[W_DIM] != expectedW) {
	362	+ VECTOR_INNER_ERR_REPORT_TILIING(context,
	363	+ "AvgPoolGrad: when padmode is VALID, the gradshape in h-dim and w-dim should be [%ld] [%ld], but got [%ld] [%ld]",
	364	+ expectedH, expectedW, inputData.gradShape[H_DIM],
	365	+ inputData.gradShape[W_DIM]);
	366	+ return ge::GRAPH_FAILED;
	367	+ }
	368	+ return ge::GRAPH_SUCCESS;
	369	+}
	370	+
	371	+static ge::graphStatus CheckGradShapeForSame(gert::TilingContext* context, AvgPoolV2GradInputInfo& inputData)
	372	+{
	373	+ int64_t expectedH = (inputData.inputShape[H_DIM] + inputData.stride[H_DIM] - 1) / inputData.stride[H_DIM];
	374	+ int64_t expectedW = (inputData.inputShape[W_DIM] + inputData.stride[W_DIM] - 1) / inputData.stride[W_DIM];
	375	+ if (inputData.gradShape[H_DIM] != expectedH \|\| inputData.gradShape[W_DIM] != expectedW) {
	376	+ VECTOR_INNER_ERR_REPORT_TILIING(context,
	377	+ "AvgPoolGrad: when padmode is SAME, the gradshape in h-dim and w-dim should be [%ld] [%ld], but got [%ld] [%ld]",
	378	+ expectedH, expectedW, inputData.gradShape[H_DIM],
	379	+ inputData.gradShape[W_DIM]);
	380	+ return ge::GRAPH_FAILED;
	381	+ }
	382	+ return ge::GRAPH_SUCCESS;
	383	+}
	384	+
	385	+static ge::graphStatus CheckGradShape(gert::TilingContext* context, AvgPoolV2GradInputInfo& inputData,
	386	+ const AvgPoolGradCommon& commInfo)
	387	+{
	388	+ if (commInfo.padModeStr == "VALID") {
	389	+ return CheckGradShapeForValid(context, inputData);
	390	+ } else if (commInfo.padModeStr == "SAME") {
	391	+ return CheckGradShapeForSame(context, inputData);
	392	+ }
	393	+ VECTOR_INNER_ERR_REPORT_TILIING(context,
	394	+ "AvgPoolGrad: unsupported pad mode [%s], only VALID and SAME are supported",
	395	+ commInfo.padModeStr.c_str());
	396	+ return ge::GRAPH_FAILED;
	397	+}
	398	+
	399	+ge::graphStatus GetAvgPoolGradPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, uint64_t& coreNum)
	400	+{
	401	+ auto platformPtr = context->GetPlatformInfo();
	402	+ if (platformPtr == nullptr) {
	403	+ auto compileInfoPtr = reinterpret_cast<const AvgPoolGradCompileInfo*>(context->GetCompileInfo());
	404	+ OP_TILING_CHECK(
	405	+ compileInfoPtr == nullptr, CUBE_INNER_ERR_REPORT(context, "compile info is null"),
	406	+ return ge::GRAPH_FAILED);
	407	+ coreNum = compileInfoPtr->coreNum;
	408	+ ubSize = compileInfoPtr->ubSize;
	409	+ } else {
	410	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformPtr);
	411	+ coreNum = ascendcPlatform.GetCoreNumAiv();
	412	+
	413	+ uint64_t ubSizePlatform;
	414	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatform);
	415	+ ubSize = static_cast<int64_t>(ubSizePlatform);
	416	+ }
	417	+
	418	+ OP_TILING_CHECK(
	419	+ coreNum == 0, CUBE_INNER_ERR_REPORT(context, "coreNum is 0"), return ge::GRAPH_FAILED);
	420	+ return ge::GRAPH_SUCCESS;
	421	+}
	422	+
	423	+ge::graphStatus GetAvgPoolGradShapeAttrsInfo(gert::TilingContext *context, AvgPoolV2GradInputInfo& inputData)
	424	+{
	425	+ auto runtimeAttrs = context->GetAttrs();
	426	+ AvgPoolGradCommon commInfo;
	427	+ OPS_CHECK_NULL_WITH_CONTEXT(context, runtimeAttrs);
	428	+
	429	+ OP_TILING_CHECK(GetAttrsInfo(context, runtimeAttrs, inputData, commInfo) != ge::GRAPH_SUCCESS,
	430	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "GetAttrsInfo fail."),
	431	+ return ge::GRAPH_FAILED);
	432	+
	433	+ OP_TILING_CHECK(GetShapeAndDtype(context, runtimeAttrs, inputData, commInfo) != ge::GRAPH_SUCCESS,
	434	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "GetShapeAndDtype fail."),
	435	+ return ge::GRAPH_FAILED);
	436	+
	437	+ OP_TILING_CHECK(GetKernelKsizeInfo(context, runtimeAttrs, inputData, commInfo) != ge::GRAPH_SUCCESS,
	438	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "GetKernelKsizeInfo fail."),
	439	+ return ge::GRAPH_FAILED);
	440	+
	441	+ OP_TILING_CHECK(GetStrideInfo(context, runtimeAttrs, inputData, commInfo) != ge::GRAPH_SUCCESS,
	442	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "GetStrideInfo fail."),
	443	+ return ge::GRAPH_FAILED);
	444	+
	445	+ OP_TILING_CHECK(GetPadInfo(context, inputData, commInfo) != ge::GRAPH_SUCCESS,
	446	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "GetPadInfo fail."),
	447	+ return ge::GRAPH_FAILED);
	448	+
	449	+ OP_TILING_CHECK(CheckGradShape(context, inputData, commInfo) != ge::GRAPH_SUCCESS,
	450	+ VECTOR_INNER_ERR_REPORT_TILIING(context, "CheckGradShape fail."),
	451	+ return ge::GRAPH_FAILED);
	452	+
	453	+ if (IsGreaterThanInt32Max(inputData)) {
	454	+ inputData.isInt32Meet = 0;
	455	+ } else {
	456	+ inputData.isInt32Meet = ONE;
	457	+ }
	458	+ inputData.hasDivisor = 0;
	459	+
	460	+ return ge::GRAPH_SUCCESS;
	461	+}
	462	+
	463	+} // namespace optiling

Apooling/avg_pool_grad/op_host/arch35/avg_pool_grad_tiling_base.h+53-0

@@ -0,0 +1,53 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file avg_pool_grad_tiling_base.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#ifndef OP_IMPL_AVG_POOL_GRAD_TILING_BASE_H_
	17	+#define OP_IMPL_AVG_POOL_GRAD_TILING_BASE_H_
	18	+
	19	+#include <array>
	20	+
	21	+#include "register/op_def_registry.h"
	22	+#include "tiling/tiling_api.h"
	23	+#include "op_host/tiling_base.h"
	24	+#include "util/math_util.h"
	25	+#include "op_common/op_host/util/platform_util.h"
	26	+#include "../../../avg_pool_v2_grad/op_host/arch35/avg_pool_v2_grad_tiling_common.h"
	27	+
	28	+namespace optiling
	29	+{
	30	+struct AvgPoolGradCompileInfo {
	31	+ uint64_t coreNum;
	32	+ uint64_t ubSize;
	33	+};
	34	+
	35	+struct AvgPoolGradCommon {
	36	+ int64_t nDim;
	37	+ int64_t cDim;
	38	+ int64_t hDim;
	39	+ int64_t wDim;
	40	+ std::string padModeStr;
	41	+};
	42	+
	43	+ge::graphStatus Tiling4AvgPoolGrad(gert::TilingContext* context);
	44	+
	45	+ge::graphStatus TilingPrepare4AvgPoolGrad(gert::TilingParseContext* context);
	46	+
	47	+ge::graphStatus GetAvgPoolGradPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, uint64_t& coreNum);
	48	+
	49	+ge::graphStatus GetAvgPoolGradShapeAttrsInfo(gert::TilingContext* context, AvgPoolV2GradInputInfo& inputData);
	50	+
	51	+} // namespace optiling
	52	+
	53	+#endif

Apooling/avg_pool_grad/op_host/avg_pool_grad_def.cpp+70-0

@@ -0,0 +1,70 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+* \file avg_pool_grad_def.cpp
	13	+* \brief
	14	+*/
	15	+
	16	+#include "register/op_def_registry.h"
	17	+
	18	+namespace ops {
	19	+
	20	+class AvgPoolGrad : public OpDef {
	21	+public:
	22	+ const std::vector<ge::DataType> AvgPoolGradXDataType = {ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16};
	23	+ const std::vector<ge::Format> AvgPoolGradXFormat = {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
	24	+ explicit AvgPoolGrad(const char* name) : OpDef(name) {
	25	+ this->Input("orig_input_shape")
	26	+ .ParamType(REQUIRED)
	27	+ .ValueDepend(OPTIONAL)
	28	+ .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32})
	29	+ .Format({AvgPoolGradXFormat})
	30	+ .UnknownShapeFormat({AvgPoolGradXFormat})
	31	+ .AutoContiguous();
	32	+ this->Input("input_grad")
	33	+ .ParamType(REQUIRED)
	34	+ .DataType(AvgPoolGradXDataType)
	35	+ .Format(AvgPoolGradXFormat)
	36	+ .UnknownShapeFormat(AvgPoolGradXFormat)
	37	+ .AutoContiguous();
	38	+ this->Output("out_grad")
	39	+ .ParamType(REQUIRED)
	40	+ .DataType(AvgPoolGradXDataType)
	41	+ .Format(AvgPoolGradXFormat)
	42	+ .UnknownShapeFormat(AvgPoolGradXFormat)
	43	+ .AutoContiguous();
	44	+ this->Attr("ksize")
	45	+ .AttrType(REQUIRED)
	46	+ .ListInt();
	47	+ this->Attr("strides")
	48	+ .AttrType(REQUIRED)
	49	+ .ListInt();
	50	+ this->Attr("padding")
	51	+ .AttrType(REQUIRED)
	52	+ .String();
	53	+ this->Attr("data_format")
	54	+ .AttrType(OPTIONAL)
	55	+ .String("NHWC");
	56	+
	57	+ OpAICoreConfig aiCoreConfig;
	58	+ aiCoreConfig.DynamicCompileStaticFlag(true)
	59	+ .DynamicFormatFlag(false)
	60	+ .DynamicRankSupportFlag(true)
	61	+ .DynamicShapeSupportFlag(true)
	62	+ .NeedCheckSupportFlag(false)
	63	+ .PrecisionReduceFlag(true)
	64	+ .ExtendCfgInfo("opFile.value", "avg_pool_grad_apt");
	65	+ this->AICore().AddConfig("ascend950", aiCoreConfig);
	66	+ }
	67	+};
	68	+
	69	+OP_ADD(AvgPoolGrad);
	70	+} // namespace ops

Apooling/avg_pool_grad/op_host/avg_pool_grad_infershape.cpp+198-0

@@ -0,0 +1,198 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file avg_pool_grad_infershape.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+#include <string>
	17	+#include "error_util.h"
	18	+#include "exe_graph/runtime/infer_shape_context.h"
	19	+#include "graph/utils/type_utils.h"
	20	+#include "log/log.h"
	21	+#include "register/op_impl_registry.h"
	22	+#include "util/shape_util.h"
	23	+
	24	+using namespace ge;
	25	+
	26	+namespace ops {
	27	+
	28	+static constexpr size_t IDX_ORIGIN_INPUT = 0;
	Y Y YuningYu3月23日变量清理
	29	+static constexpr size_t IDX_GRAD_INPUT = 1;
	30	+static constexpr size_t IDX_OUTPUT = 0;
	31	+static constexpr size_t IDX_ZERO = 0;
	32	+static constexpr size_t IDX_ONE = 1;
	33	+static constexpr size_t IDX_THREE = 3;
	34	+static constexpr size_t CHW_DIMS = 3;
	35	+static constexpr size_t NCHW_DIMS = 4;
	36	+static constexpr size_t ATTR_LIST_SHAPE_SIZE = 4;
	37	+static constexpr size_t ATTR_KERNEL_POS = 0;
	38	+static constexpr size_t ATTR_STRIDE_POS = 1;
	39	+static constexpr size_t ATTR_PADDING_POS = 2;
	40	+static constexpr size_t ATTR_FORMAT_POS = 3;
	41	+static constexpr size_t ONE = 1;
	42	+static constexpr int64_t UNKNOWN_DIM_VALUE = -1LL;
	43	+
	44	+inline bool IsConstTensor(const gert::Tensor* inputTensor) {
	45	+ if (inputTensor != nullptr) {
	46	+ if (inputTensor->GetAddr() == nullptr) {
	47	+ return inputTensor->GetShapeSize() == 0;
	48	+ }
	49	+ return true;
	50	+ }
	51	+ return false;
	52	+}
	53	+
	54	+inline ge::graphStatus SetAllUnknownDim(const int64_t rank, gert::Shape* output_shape)
	55	+{
	56	+ OP_CHECK_IF(
	57	+ output_shape == nullptr, OP_LOGD("SetAllUnknownDim", "the output_shape is nullptr, return unsuccess"),
	58	+ return ge::GRAPH_FAILED);
	59	+ output_shape->SetDimNum(rank);
	60	+ for (int64_t i = 0; i < rank; ++i) {
	61	+ output_shape->SetDim(i, UNKNOWN_DIM_VALUE);
	62	+ }
	63	+ OP_LOGD("SetAllUnknownDim", "set all dim = -1, output = %s", Ops::Base::ToString(*output_shape).c_str());
	64	+ return ge::GRAPH_SUCCESS;
	65	+}
	66	+
	67	+ge::graphStatus CheckKernelAndStrides(gert::InferShapeContext* context, const std::string& dataFormatStr)
	68	+{
	69	+ auto attrs = context->GetAttrs();
	70	+ auto ksize = attrs->GetAttrPointer<gert::ContinuousVector>(ATTR_KERNEL_POS);
	71	+ OP_CHECK_NULL_WITH_CONTEXT(context, ksize);
	72	+ OP_CHECK_IF(
	73	+ ksize->GetSize() != ATTR_LIST_SHAPE_SIZE,
	74	+ OP_LOGE(context->GetNodeName(), "Length of ksize %lu must be 4!", ksize->GetSize()), return GRAPH_FAILED);
	75	+ auto ksize_data = reinterpret_cast<const int64_t*>(ksize->GetData());
	76	+
	77	+ if (dataFormatStr == "NCHW") {
	78	+ OP_CHECK_IF(ksize_data[IDX_ZERO] != ONE,
	79	+ OP_LOGE(context->GetNodeName(), "Pooling ksize[0] %ld must be 1.", ksize_data[IDX_ZERO]),
	80	+ return GRAPH_FAILED);
	81	+ OP_CHECK_IF(ksize_data[IDX_ONE] != ONE,
	82	+ OP_LOGE(context->GetNodeName(), "Pooling ksize[1] %ld must be 1.", ksize_data[IDX_ONE]),
	83	+ return GRAPH_FAILED);
	84	+ } else if (dataFormatStr == "NHWC") {
	85	+ OP_CHECK_IF(ksize_data[IDX_ZERO] != ONE,
	86	+ OP_LOGE(context->GetNodeName(), "Pooling ksize[0] %ld must be 1.", ksize_data[IDX_ZERO]),
	87	+ return GRAPH_FAILED);
	88	+ OP_CHECK_IF(ksize_data[IDX_THREE] != ONE,
	89	+ OP_LOGE(context->GetNodeName(), "Pooling ksize[3] %ld must be 1.", ksize_data[IDX_THREE]),
	90	+ return GRAPH_FAILED);
	91	+ }
	92	+
	93	+ auto strides = attrs->GetAttrPointer<gert::ContinuousVector>(ATTR_STRIDE_POS);
	94	+ OP_CHECK_NULL_WITH_CONTEXT(context, strides);
	95	+ OP_CHECK_IF(
	96	+ strides->GetSize() != ATTR_LIST_SHAPE_SIZE,
	97	+ OP_LOGE(context->GetNodeName(), "Length of strides %lu must be 4!", strides->GetSize()), return GRAPH_FAILED);
	98	+ auto strides_data = reinterpret_cast<const int64_t*>(strides->GetData());
	99	+
	100	+ if (dataFormatStr == "NCHW") {
	101	+ OP_CHECK_IF(strides_data[IDX_ZERO] != ONE,
	102	+ OP_LOGE(context->GetNodeName(), "Pooling stride size[0] %ld must be 1.", strides_data[IDX_ZERO]),
	103	+ return GRAPH_FAILED);
	104	+ OP_CHECK_IF(strides_data[IDX_ONE] != ONE,
	105	+ OP_LOGE(context->GetNodeName(), "Pooling stride size[1] %ld must be 1.", strides_data[IDX_ONE]),
	106	+ return GRAPH_FAILED);
	107	+ } else if (dataFormatStr == "NHWC") {
	108	+ OP_CHECK_IF(strides_data[IDX_ZERO] != ONE,
	109	+ OP_LOGE(context->GetNodeName(), "Pooling stride size[0] %ld must be 1.", strides_data[IDX_ZERO]),
	110	+ return GRAPH_FAILED);
	111	+ OP_CHECK_IF(strides_data[IDX_THREE] != ONE,
	112	+ OP_LOGE(context->GetNodeName(), "Pooling stride size[3] %ld must be 1.", strides_data[IDX_THREE]),
	113	+ return GRAPH_FAILED);
	114	+ }
	115	+ return ge::GRAPH_SUCCESS;
	116	+}
	117	+
	118	+ge::graphStatus InferShape4AvgPoolGrad(gert::InferShapeContext* context)
	119	+{
	120	+ if (context == nullptr) {
	121	+ return GRAPH_FAILED;
	122	+ }
	123	+ OP_LOGD(context->GetNodeName(), "runtime2.0 AvgPoolGrad infershape running");
	124	+
	125	+ auto gradDesc = context->GetInputDesc(IDX_GRAD_INPUT);
	126	+ OP_CHECK_NULL_WITH_CONTEXT(context, gradDesc);
	127	+ auto gradOriFormat = gradDesc->GetOriginFormat();
	128	+
	129	+ OP_CHECK_IF(
	130	+ gradOriFormat != FORMAT_ND && gradOriFormat != FORMAT_NCHW && gradOriFormat != FORMAT_NHWC,
	131	+ OP_LOGE(context->GetNodeName(), "format only supports ND, NCHW, NHWC"), return GRAPH_FAILED);
	132	+
	133	+ auto attrs = context->GetAttrs();
	134	+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
	135	+
	136	+ const char* dataFormatPtr = attrs->GetAttrPointer<char>(ATTR_FORMAT_POS);
	137	+ OP_LOGE_IF(dataFormatPtr == nullptr, GRAPH_FAILED, context->GetNodeName(), "Get dataFormat failed.");
	138	+ std::string dataFormatStr(dataFormatPtr);
	139	+
	140	+ auto padding = attrs->GetAttrPointer<char>(ATTR_PADDING_POS);
	141	+ OP_CHECK_NULL_WITH_CONTEXT(context, padding);
	142	+ OP_CHECK_IF(
	S S sunday3月30日问题: padding 校验条件使用了 `!strcmp(padding, "SAME") && !strcmp(padding, "VALID")`，该条件恒为 false，非法 padding 不会被拦截。建议: 改为“既不是 SAME 也不是 VALID”时返回失败，例如 `strcmp(...) != 0 && strcmp(...) != 0`。
	143	+ strcmp(padding, "SAME") != 0 && strcmp(padding, "VALID") != 0,
	144	+ OP_LOGE(context->GetNodeName(),"attr padding(%s) only support SAME、 VALID", padding), return GRAPH_FAILED);
	145	+
	146	+ ge::graphStatus checkStatus = CheckKernelAndStrides(context, dataFormatStr);
	147	+ OP_CHECK_IF(checkStatus != ge::GRAPH_SUCCESS, OP_LOGD(context->GetNodeName(), "CheckKernelAndStrides failed"),
	148	+ return checkStatus);
	149	+
	150	+ const gert::Tensor* inputShape0 = context->GetInputTensor(IDX_ORIGIN_INPUT);
	151	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape0);
	152	+ size_t inputDimNum = static_cast<size_t>(inputShape0->GetOriginShape().GetShapeSize());
	153	+ const int32_t* shapeValue = inputShape0->GetData<int32_t>();
	154	+ OP_CHECK_IF(
	155	+ inputDimNum != CHW_DIMS && inputDimNum != NCHW_DIMS,
	156	+ OP_LOGE(context->GetNodeName(), "input dim num should be 3 or 4, but get %zu.", inputDimNum),
	157	+ return GRAPH_FAILED);
	158	+ const gert::Shape* inputShape1 = context->GetInputShape(IDX_ORIGIN_INPUT);
	159	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape1);
	160	+
	161	+ gert::Shape* OutShape = context->GetOutputShape(IDX_OUTPUT);
	162	+ OP_CHECK_NULL_WITH_CONTEXT(context, OutShape);
	163	+
	164	+ if (Ops::Base::IsUnknownShape(*inputShape1) \|\| !IsConstTensor(inputShape0)) {
	165	+ SetAllUnknownDim(inputDimNum, OutShape);
	166	+ }
	167	+
	168	+ if (Ops::Base::IsUnknownRank(*inputShape1)) {
	169	+ Ops::Base::SetUnknownRank(*OutShape);
	170	+ return ge::GRAPH_SUCCESS;
	171	+ }
	172	+
	173	+ OutShape->SetDimNum(inputDimNum);
	174	+
	175	+ for (size_t idx = 0; idx < inputDimNum; ++idx) {
	176	+ OutShape->SetDim(idx, shapeValue[idx]);
	177	+ }
	178	+
	179	+ OP_LOGD(context->GetNodeName(), "runtime2.0 end AvgPoolGrad infershape");
	180	+ return ge::GRAPH_SUCCESS;
	181	+}
	182	+
	183	+static ge::graphStatus InferDataType4AvgPoolGrad(gert::InferDataTypeContext* context)
	184	+{
	185	+ if (context == nullptr) {
	186	+ return GRAPH_FAILED;
	187	+ }
	188	+
	189	+ const ge::DataType xDtype = context->GetInputDataType(1);
	190	+ context->SetOutputDataType(0, xDtype);
	191	+ return GRAPH_SUCCESS;
	192	+}
	193	+
	194	+IMPL_OP_INFERSHAPE(AvgPoolGrad)
	195	+ .InputsDataDependency({0})
	196	+ .InferShape(InferShape4AvgPoolGrad)
	197	+ .InferDataType(InferDataType4AvgPoolGrad);
	198	+} // namespace ops

Apooling/avg_pool_grad/op_host/config/ascend950/avg_pool_grad_binary.json+191-0

@@ -0,0 +1,191 @@
	1	+{
	2	+ "op_type": "AvgPoolGrad",
	3	+ "op_list": [
	4	+ {
	5	+ "bin_filename": "AvgPoolGrad_aa39a01270e5cc309f704228d56100aa",
	6	+ "inputs": [
	7	+ {
	8	+ "name": "orig_input_shape",
	9	+ "index": 0,
	10	+ "dtype": "int32",
	11	+ "format": "ND",
	12	+ "paramType": "required",
	13	+ "shape": [
	14	+ -2
	15	+ ],
	16	+ "format_match_mode": "FormatAgnostic"
	17	+ },
	18	+ {
	19	+ "name": "input_grad",
	20	+ "index": 1,
	21	+ "dtype": "float16",
	22	+ "format": "ND",
	23	+ "paramType": "required",
	24	+ "shape": [
	25	+ -2
	26	+ ],
	27	+ "format_match_mode": "FormatAgnostic"
	28	+ }
	29	+ ],
	30	+ "outputs": [
	31	+ {
	32	+ "name": "output_grad",
	33	+ "index": 0,
	34	+ "dtype": "float16",
	35	+ "format": "ND",
	36	+ "paramType": "required",
	37	+ "shape": [
	38	+ -2
	39	+ ],
	40	+ "format_match_mode": "FormatAgnostic"
	41	+ }
	42	+ ],
	43	+ "attrs": [
	44	+ {
	45	+ "name": "ksize",
	46	+ "dtype": "list_int",
	47	+ "value": null
	48	+ },
	49	+ {
	50	+ "name": "strides",
	51	+ "dtype": "list_int",
	52	+ "value": null
	53	+ },
	54	+ {
	55	+ "name": "padding",
	56	+ "dtype": "string",
	57	+ "value": null
	58	+ },
	59	+ {
	60	+ "name": "data_format",
	61	+ "dtype": "string",
	62	+ "value": null
	63	+ }
	64	+ ]
	65	+ },
	66	+ {
	67	+ "bin_filename": "AvgPoolGrad_bb9c950763acb46236b38d30f1a100bb",
	68	+ "inputs": [
	69	+ {
	70	+ "name": "orig_input_shape",
	71	+ "index": 0,
	72	+ "dtype": "int32",
	73	+ "format": "ND",
	74	+ "paramType": "required",
	75	+ "shape": [
	76	+ -2
	77	+ ],
	78	+ "format_match_mode": "FormatAgnostic"
	79	+ },
	80	+ {
	81	+ "name": "input_grad",
	82	+ "index": 1,
	83	+ "dtype": "bfloat16",
	84	+ "format": "ND",
	85	+ "paramType": "required",
	86	+ "shape": [
	87	+ -2
	88	+ ],
	89	+ "format_match_mode": "FormatAgnostic"
	90	+ }
	91	+ ],
	92	+ "outputs": [
	93	+ {
	94	+ "name": "output_grad",
	95	+ "index": 0,
	96	+ "dtype": "bfloat16",
	97	+ "format": "ND",
	98	+ "paramType": "required",
	99	+ "shape": [
	100	+ -2
	101	+ ],
	102	+ "format_match_mode": "FormatAgnostic"
	103	+ }
	104	+ ],
	105	+ "attrs": [
	106	+ {
	107	+ "name": "ksize",
	108	+ "dtype": "list_int",
	109	+ "value": null
	110	+ },
	111	+ {
	112	+ "name": "strides",
	113	+ "dtype": "list_int",
	114	+ "value": null
	115	+ },
	116	+ {
	117	+ "name": "padding",
	118	+ "dtype": "string",
	119	+ "value": null
	120	+ },
	121	+ {
	122	+ "name": "data_format",
	123	+ "dtype": "string",
	124	+ "value": null
	125	+ }
	126	+ ]
	127	+ },
	128	+ {
	129	+ "bin_filename": "AvgPoolGrad_cc333e1910a01784d66f7e60fd0200cc",
	130	+ "inputs": [
	131	+ {
	132	+ "name": "orig_input_shape",
	133	+ "index": 0,
	134	+ "dtype": "int32",
	135	+ "format": "ND",
	136	+ "paramType": "required",
	137	+ "shape": [
	138	+ -2
	139	+ ],
	140	+ "format_match_mode": "FormatAgnostic"
	141	+ },
	142	+ {
	143	+ "name": "input_grad",
	144	+ "index": 1,
	145	+ "dtype": "float32",
	146	+ "format": "ND",
	147	+ "paramType": "required",
	148	+ "shape": [
	149	+ -2
	150	+ ],
	151	+ "format_match_mode": "FormatAgnostic"
	152	+ }
	153	+ ],
	154	+ "outputs": [
	155	+ {
	156	+ "name": "output_grad",
	157	+ "index": 0,
	158	+ "dtype": "float32",
	159	+ "format": "ND",
	160	+ "paramType": "required",
	161	+ "shape": [
	162	+ -2
	163	+ ],
	164	+ "format_match_mode": "FormatAgnostic"
	165	+ }
	166	+ ],
	167	+ "attrs": [
	168	+ {
	169	+ "name": "ksize",
	170	+ "dtype": "list_int",
	171	+ "value": null
	172	+ },
	173	+ {
	174	+ "name": "strides",
	175	+ "dtype": "list_int",
	176	+ "value": null
	177	+ },
	178	+ {
	179	+ "name": "padding",
	180	+ "dtype": "string",
	181	+ "value": null
	182	+ },
	183	+ {
	184	+ "name": "data_format",
	185	+ "dtype": "string",
	186	+ "value": null
	187	+ }
	188	+ ]
	189	+ }
	190	+ ]
	191	+ }

Apooling/avg_pool_grad/op_host/config/ascend950/avg_pool_grad_simplified_key.ini+13-0

@@ -0,0 +1,13 @@
	1	+; 该文件主要影响 opc 工具编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
	2	+; [某算子]
	3	+; default=xx
	4	+; ascendxx=xx
	5	+; 其中，default为默认mode，ascendxx为可选mode，如果不同芯片有差异化要求时，需要配置；
	6	+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
	7	+; 2)如果仅有default配置：各个版本按default配置
	8	+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
	9	+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
	10	+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
	11	+; 6)是否是AscendC算子，由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
	12	+[AvgPoolGrad]
	13	+default=0

Apooling/avg_pool_grad/op_kernel/avg_pool_grad_apt.cpp+72-0

@@ -0,0 +1,72 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file avg_pool_v2_grad.cpp
	13	+ * \brief
	14	+ */
	15	+#include "../avg_pool_v2_grad/arch35/avg_pool_v2_grad_simt.h"
	16	+#include "../avg_pool_v2_grad/arch35/avg_pool_v2_grad_nhwc_kernel.h"
	17	+#include "../avg_pool_v2_grad/arch35/avg_pool_v2_grad_nchw_kernel.h"
	18	+#include "../avg_pool_v2_grad/arch35/avg_pool_v2_grad_tiling_data.h"
	19	+#include "../avg_pool_v2_grad/arch35/avg_pool_v2_grad_tiling_key.h"
	20	+using namespace AscendC;
	21	+using namespace AvgPoolV2Grad;
	22	+using namespace AvgPoolV2GradNHWCNameSpace;
	23	+using namespace AvgPoolV2GradNCHWNameSpace;
	24	+template <
	25	+ uint32_t schMode, uint32_t format, uint32_t isInt32Meet, uint32_t isPad, uint32_t isCheckRange,
	26	+ uint32_t countIncludePad, uint32_t hasDivsor>
	27	+__global__ __aicore__ void avg_pool_grad(
	28	+ GM_ADDR orig_input_shape, GM_ADDR input_grad, GM_ADDR out_grad, GM_ADDR workspace, GM_ADDR tiling)
	29	+{
	30	+ AscendC::TPipe pipe;
	31	+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
	32	+ REGISTER_TILING_DEFAULT(AvgPoolV2GradTilingData);
	33	+ if constexpr (schMode == TPL_SIMT_KERNEL) {
	34	+ REGISTER_TILING_FOR_TILINGKEY("schMode == TPL_SIMT_KERNEL", AvgPoolV2GradSimtTilingData);
	YL Y YuningYu3月23日放在struct内 L liuchuangdev3月23日注册移出到tiling声明的地方
	35	+ GET_TILING_DATA_WITH_STRUCT(AvgPoolV2GradSimtTilingData, tilingData, tiling);
	36	+ if constexpr (isInt32Meet == TPL_INT32) {
	37	+ AvgPoolV2GradSimtNamespace::AvgPoolV2GradSimt<DTYPE_INPUT_GRAD, int32_t, format, countIncludePad, hasDivsor> op(&pipe, &tilingData);
	38	+ op.Init(input_grad, out_grad);
	39	+ op.Process();
	40	+ } else {
	41	+ AvgPoolV2GradSimtNamespace::AvgPoolV2GradSimt<DTYPE_INPUT_GRAD, int64_t, format, countIncludePad, hasDivsor> op(&pipe, &tilingData);
	42	+ op.Init(input_grad, out_grad);
	43	+ op.Process();
	44	+ }
	45	+ } else if constexpr (schMode == TPL_NCHW_KERNEL) {
	46	+ REGISTER_TILING_FOR_TILINGKEY("schMode == TPL_NCHW_KERNEL", AvgPoolV2GradNCHWTilingData);
	47	+ if constexpr (isInt32Meet == 1) {
	48	+ GET_TILING_DATA_WITH_STRUCT(AvgPoolV2GradNCHWTilingData, tilingData, tiling);
	49	+ AvgPoolV2GradNCHWKernel<DTYPE_INPUT_GRAD, int32_t, hasDivsor, isCheckRange, countIncludePad> op(&pipe, &tilingData);
	50	+ op.Init(input_grad, out_grad);
	51	+ op.Process();
	52	+ } else {
	53	+ GET_TILING_DATA_WITH_STRUCT(AvgPoolV2GradNCHWTilingData, tilingData, tiling);
	54	+ AvgPoolV2GradNCHWKernel<DTYPE_INPUT_GRAD, int64_t, hasDivsor, isCheckRange, countIncludePad> op(&pipe, &tilingData);
	55	+ op.Init(input_grad, out_grad);
	56	+ op.Process();
	57	+ }
	58	+ } else if constexpr (schMode == TPL_NHWC_KERNEL) { //NHWC
	59	+ REGISTER_TILING_FOR_TILINGKEY("schMode == TPL_NHWC_KERNEL", AvgPoolV2GradNHWCTilingData);
	60	+ if constexpr (isInt32Meet == TPL_INT32){
	61	+ GET_TILING_DATA_WITH_STRUCT(AvgPoolV2GradNHWCTilingData, tilingData, tiling);
	62	+ AvgPoolV2GradNHWCNameSpace::AvgPoolV2GradKernelNHWC<DTYPE_INPUT_GRAD, int32_t, hasDivsor, isCheckRange, countIncludePad> op(&pipe, &tilingData);
	63	+ op.Init(input_grad, out_grad);
	64	+ op.Process();
	65	+ } else {
	66	+ GET_TILING_DATA_WITH_STRUCT(AvgPoolV2GradNHWCTilingData, tilingData, tiling);
	67	+ AvgPoolV2GradNHWCNameSpace::AvgPoolV2GradKernelNHWC<DTYPE_INPUT_GRAD, int64_t, hasDivsor, isCheckRange, countIncludePad> op(&pipe, &tilingData);
	68	+ op.Init(input_grad, out_grad);
	69	+ op.Process();
	70	+ }
	71	+ }
	72	+}

Apooling/avg_pool_grad/tests/CMakeLists.txt+18-0

@@ -0,0 +1,18 @@
	1	+# ----------------------------------------------------------------------------
	2	+# This program is free software, you can redistribute it and/or modify.
	3	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+# This file is a part of the CANN Open Software.
	5	+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
	6	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------
	10	+
	11	+message(STATUS "=== Debug: start ops.pooling.avg_pool_grad.tests.CMakeLists.txt ")
	12	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	13	+message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")
	14	+foreach(SUB_DIR ${CURRENT_DIRS})
	15	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	16	+ add_subdirectory(${SUB_DIR})
	17	+ endif()
	18	+endforeach()

Apooling/avg_pool_grad/tests/ut/CMakeLists.txt+18-0

@@ -0,0 +1,18 @@
	1	+# ----------------------------------------------------------------------------
	2	+# This program is free software, you can redistribute it and/or modify.
	3	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+# This file is a part of the CANN Open Software.
	5	+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
	6	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------
	10	+
	11	+# 每个目录下需要生成的可执行文件，具体参考：ops/built-in/test/CMakeLists.txt： 50~124
	12	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	13	+message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")
	14	+foreach(SUB_DIR ${CURRENT_DIRS})
	15	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	16	+ add_subdirectory(${SUB_DIR})
	17	+ endif()
	18	+endforeach()

Apooling/avg_pool_grad/tests/ut/op_host/CMakeLists.txt+15-0

@@ -0,0 +1,15 @@
	1	+# ----------------------------------------------------------------------------
	2	+# This program is free software, you can redistribute it and/or modify.
	3	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+# This file is a part of the CANN Open Software.
	5	+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
	6	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------
	10	+
	11	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	12	+if(UT_TEST_ALL OR OP_HOST_UT)
	13	+ add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
	14	+ add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
	15	+endif()

Apooling/avg_pool_grad/tests/ut/op_host/arch35/test_avg_pool_grad_tiling.cpp+211-0

@@ -0,0 +1,211 @@
	1	+ /**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#include <iostream>
	12	+#include <fstream>
	13	+#include <vector>
	14	+#include <gtest/gtest.h>
	15	+#include "log/log.h"
	16	+#include "kernel_run_context_facker.h"
	17	+#include "test_cube_util.h"
	18	+#include "exe_graph/runtime/storage_format.h"
	19	+#include "exe_graph/runtime/storage_shape.h"
	20	+#include "platform/platform_infos_def.h"
	21	+#include "ut_op_util.h"
	22	+#include "../../../../op_host/arch35/avg_pool_grad_tiling_base.h"
	23	+
	24	+using namespace ut_util;
	25	+using namespace std;
	26	+using namespace ge;
	27	+
	28	+class AvgPoolGradTiling : public testing::Test {
	29	+protected:
	30	+ static void SetUpTestCase()
	31	+ {
	32	+ std::cout << "AvgPoolGradTiling SetUp" << std::endl;
	33	+ }
	34	+
	35	+ static void TearDownTestCase()
	36	+ {
	37	+ std::cout << "AvgPoolGradTiling TearDown" << std::endl;
	38	+ }
	39	+};
	40	+
	41	+template <typename T>
	42	+void SetConstInput(
	43	+ size_t const_index, ge::DataType dtype, T* const_data, int64_t data_size,
	44	+ std::vector<std::pair<size_t, std::unique_ptr<uint8_t[]>>>& const_tensors)
	45	+{
	46	+ std::unique_ptr<uint8_t[]> input_tensor_holder =
	47	+ std::unique_ptr<uint8_t[]>(new uint8_t[sizeof(gert::Tensor) + sizeof(T) * data_size]);
	48	+ auto input_tensor = reinterpret_cast<gert::Tensor*>(input_tensor_holder.get());
	49	+ gert::Tensor tensor(
	50	+ {{data_size}, {data_size}}, {ge::FORMAT_ND, ge::FORMAT_ND, {}}, gert::kFollowing, dtype, nullptr);
	51	+ std::memcpy(input_tensor, &tensor, sizeof(gert::Tensor));
	52	+ auto tensor_data = reinterpret_cast<T*>(input_tensor + 1);
	53	+ for (int64_t i = 0; i < data_size; i++) {
	54	+ tensor_data[i] = const_data[i];
	55	+ }
	56	+ input_tensor->SetData(gert::TensorData{tensor_data});
	57	+ auto pair = std::make_pair(const_index, std::move(input_tensor_holder));
	58	+ const_tensors.push_back(std::move(pair));
	59	+}
	60	+
	61	+static void ExecuteTestCase(
	62	+ gert::StorageShape xShape, gert::StorageShape yShape, gert::StorageShape gradShape,
	63	+ std::vector<int64_t> ksize, std::vector<int64_t> strides, std::string padding,
	64	+ std::string data_format,
	65	+ ge::DataType dtype, ge::DataType dtypeIdx, uint64_t except_tilingkey,
	66	+ int32_t* shape_data)
	67	+{
	68	+ dlog_setlevel(0, 0, 0);
	69	+
	70	+ string compile_info_string = R"({
	71	+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
	72	+ "Intrinsic_fix_pipe_l0c2out": false,
	73	+ "Intrinsic_data_move_l12ub": true,
	74	+ "Intrinsic_data_move_l0c2ub": true,
	75	+ "Intrinsic_data_move_out2l1_nd2nz": false,
	76	+ "UB_SIZE": 245760, "L2_SIZE": 33554432, "L1_SIZE": 524288,
	77	+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
	78	+ "CORE_NUM": 64}
	79	+ })";
	80	+ map<string, string> soc_infos;
	81	+ map<string, string> aicore_spec;
	82	+ map<string, string> intrinsics;
	83	+ GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
	84	+ std::map<std::string, std::string> soc_version_infos = {{"Short_SoC_version", "ascend950"}};
	85	+ // platform info
	86	+ fe::PlatFormInfos platform_info;
	87	+ platform_info.Init();
	88	+ // compile info
	89	+ optiling::AvgPoolGradCompileInfo compile_info;
	90	+
	91	+ std::string op_type("AvgPoolGrad");
	92	+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str()), nullptr);
	93	+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
	94	+ auto tiling_parse_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling_parse;
	95	+
	96	+ std::vector<std::pair<size_t, std::unique_ptr<uint8_t[]>>> const_tensors;
	97	+ SetConstInput(0, DT_INT32, shape_data, 4, const_tensors);
	98	+
	99	+ // tilingParseFunc simulate
	100	+ auto kernel_holder =
	101	+ gert::KernelRunContextFaker()
	102	+ .KernelIONum(2, 1)
	103	+ .Inputs({const_cast<char>(compile_info_string.c_str()), reinterpret_cast<void>(&platform_info)})
	104	+ .Outputs(std::vector<void*>{&compile_info})
	105	+ .Build();
	106	+
	107	+ ASSERT_TRUE((kernel_holder.GetContext<gert::TilingParseContext>()) ->GetPlatformInfo()->Init());
	108	+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
	109	+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
	110	+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
	111	+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes(
	112	+ "AICoreintrinsicDtypeMap", intrinsics);
	113	+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes(
	114	+ "version", soc_version_infos);
	115	+ ASSERT_EQ(tiling_parse_func((kernel_holder.GetContext<gert::KernelContext>())), ge::GRAPH_SUCCESS);
	116	+
	117	+ // tilingFunc simulate
	118	+ auto param = gert::TilingData::CreateCap(4096);
	119	+ auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
	120	+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
	121	+ ASSERT_NE(param, nullptr);
	122	+ auto holder = gert::TilingContextFaker()
	123	+ .SetOpType(op_type)
	124	+ .NodeIoNum(2, 1)
	125	+ .IrInstanceNum({1, 1})
	126	+ .InputShapes({&xShape, &gradShape})
	127	+ .OutputShapes({&yShape})
	128	+ .CompileInfo(&compile_info)
	129	+ .PlatformInfo(reinterpret_cast<char*>(&platform_info))
	130	+ .NodeInputTd(0, dtypeIdx, ge::FORMAT_ND, ge::FORMAT_ND)
	131	+ .NodeInputTd(1, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
	132	+ .NodeOutputTd(0, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
	133	+ .NodeAttrs(
	134	+ {{"ksize", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>(ksize)},
	135	+ {"strides", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>(strides)},
	136	+ {"padding", Ops::NN::AnyValue::CreateFrom<std::string>(padding)},
	137	+ {"data_format", Ops::NN::AnyValue::CreateFrom<std::string>(data_format)}})
	138	+ .TilingData(param.get())
	139	+ .ConstInput(const_tensors)
	140	+ .Workspace(ws_size)
	141	+ .Build();
	142	+
	143	+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
	144	+ ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
	145	+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
	146	+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
	147	+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
	148	+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
	149	+
	150	+ // workspaces nullptr return failed
	151	+ EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
	152	+ auto tiling_key = tiling_context->GetTilingKey();
	153	+ ASSERT_EQ(tiling_key, except_tilingkey);
	154	+}
	155	+
	156	+TEST_F(AvgPoolGradTiling, AvgPoolGradTiling_Test_1)
	157	+{
	158	+ gert::StorageShape xShape = {{4}, {4}};
	159	+ gert::StorageShape gradShape = {{1, 1, 1, 1}, {1, 1, 1, 1}};
	160	+ gert::StorageShape yShape = {{1, 3, 3, 1}, {1, 3, 3, 1}};
	161	+ std::vector<int64_t> ksize = {3, 3};
	162	+ std::vector<int64_t> strides = {1, 1};
	163	+ std::string padding = "VALID";
	164	+ ge::DataType dtype = ge::DT_FLOAT;
	165	+ ge::DataType dtypeIdx = ge::DT_INT32;
	166	+ std::string data_format = "NHWC";
	167	+ uint64_t except_tilingkey = 274;
	168	+ int shape_data[4] = {1, 3, 3, 1};
	169	+
	170	+ ExecuteTestCase(
	171	+ xShape, yShape, gradShape, ksize, strides, padding,
	172	+ data_format, dtype, dtypeIdx, except_tilingkey, shape_data);
	173	+}
	174	+
	175	+TEST_F(AvgPoolGradTiling, AvgPoolGradTiling_Test_2)
	176	+{
	177	+ gert::StorageShape xShape = {{4}, {4}};
	178	+ gert::StorageShape gradShape = {{1, 1, 1, 1}, {1, 1, 1, 1}};
	179	+ gert::StorageShape yShape = {{1, 1, 3, 3}, {1, 1, 3, 3}};
	180	+ std::vector<int64_t> ksize = {3, 3};
	181	+ std::vector<int64_t> strides = {1, 1};
	182	+ std::string padding = "VALID";
	183	+ ge::DataType dtype = ge::DT_FLOAT;
	184	+ ge::DataType dtypeIdx = ge::DT_INT32;
	185	+ std::string data_format = "NCHW";
	186	+ uint64_t except_tilingkey = 258;
	187	+ int shape_data[4] = {1, 1, 3, 3};
	188	+
	189	+ ExecuteTestCase(
	190	+ xShape, yShape, gradShape, ksize, strides, padding,
	191	+ data_format, dtype, dtypeIdx, except_tilingkey, shape_data);
	192	+}
	193	+
	194	+TEST_F(AvgPoolGradTiling, AvgPoolGradTiling_Test_3)
	195	+{
	196	+ gert::StorageShape xShape = {{4}, {4}};
	197	+ gert::StorageShape gradShape = {{1, 3, 3, 1}, {1, 3, 3, 1}};
	198	+ gert::StorageShape yShape = {{1, 5, 5, 1}, {1, 5, 5, 1}};
	199	+ std::vector<int64_t> ksize = {2, 2};
	200	+ std::vector<int64_t> strides = {2, 2};
	201	+ std::string padding = "SAME";
	202	+ ge::DataType dtype = ge::DT_FLOAT;
	203	+ ge::DataType dtypeIdx = ge::DT_INT32;
	204	+ std::string data_format = "NHWC";
	205	+ uint64_t except_tilingkey = 1297;
	206	+ int shape_data[4] = {1, 5, 5, 1};
	207	+
	208	+ ExecuteTestCase(
	209	+ xShape, yShape, gradShape, ksize, strides, padding,
	210	+ data_format, dtype, dtypeIdx, except_tilingkey, shape_data);
	211	+}

Apooling/avg_pool_grad/tests/ut/op_kernel/avg_pool_grad_tiling_def.h+26-0

@@ -0,0 +1,26 @@
	1	+/**
	2	+ * This program is free software, you can redistribute it and/or modify.
	3	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+ * This file is a part of the CANN Open Software.
	5	+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
	6	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file max_pool3d_with_argmax_v2_tiling.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#ifndef _GE_AVG_POOL_GRAD_TILING_DEF_H_
	17	+#define _GE_AVG_POOL_GRAD_TILING_DEF_H_
	18	+#include <cstdint>
	19	+#include <cstring>
	20	+#include "kernel_tiling/kernel_tiling.h"
	21	+
	22	+#define DT_BF16 bfloat16_t
	23	+#define ORIG_DTYPE_START DT_BF16
	24	+#define __CCE_UT_TEST__
	25	+
	26	+#define DTYPE_INPUT_GRAD float

Mscripts/kernel/binary_config/ascendc_config.json+1-0

	@@ -9,6 +9,7 @@
9	{"name":"AddLora", "compute_units": ["ascend310p", "ascend910b"], "auto_sync":true},	9	{"name":"AddLora", "compute_units": ["ascend310p", "ascend910b"], "auto_sync":true},
10	{"name":"AddExample", "compute_units": ["ascend910b", "ascend910_93", "ascend950"], "auto_sync":true, "impl_mode" : ""},	10	{"name":"AddExample", "compute_units": ["ascend910b", "ascend910_93", "ascend950"], "auto_sync":true, "impl_mode" : ""},
11	{"name":"AvgPool3DGrad", "compute_units": ["ascend910b", "ascend910_93", "ascend950"], "auto_sync" : true},	11	{"name":"AvgPool3DGrad", "compute_units": ["ascend910b", "ascend910_93", "ascend950"], "auto_sync" : true},
		12	+ {"name":"AvgPoolGrad", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend950":["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
12	{"name":"AvgPoolV2Grad", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend950":["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},	13	{"name":"AvgPoolV2Grad", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend950":["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
13	{"name":"ForeachNonFiniteCheckAndUnscale", "compute_units": ["ascend310p", "ascend910b", "ascend910_93", "ascend950", "kirinx90", "kirin9030"], "auto_sync": false, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},	14	{"name":"ForeachNonFiniteCheckAndUnscale", "compute_units": ["ascend310p", "ascend910b", "ascend910_93", "ascend950", "kirinx90", "kirin9030"], "auto_sync": false, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
14	{"name":"AdamApplyOne", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},	15	{"name":"AdamApplyOne", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},