Star938
1856
代码
代码
Issues242
Pull Requests352
Actions
项目讨论
Wiki
分析
Star938
1856
已合并
950支持AdaptiveAvgPool3D big kernel模板，解决当kernel Size大于128时，走Simt模板性能较差 #1572
SimonZzz创建于 2月5日
950支持AdaptiveAvgPool3D big kernel模板，解决当kernel Size大于128时，走Simt模板性能较差 #1572
已合并
SimonZzz创建于 2月5日
从SimonZzz/ops-nnadaptiveavgpool3d合入到cann/ops-nnmaster
共 7 个文件变更+576-10
Mpooling/adaptive_avg_pool3d/op_kernel/adaptive_avg_pool3d_apt.cpp+6-0
@@ -15,6 +15,7 @@
15 15 
16#include "kernel_operator.h"16#include "kernel_operator.h"
17#include "kernel_tiling/kernel_tiling.h"17#include "kernel_tiling/kernel_tiling.h"
18+#include "../adaptive_pool3d_common/arch35/adaptive_avg_pool3d_big_kernel.h"
18#include "../adaptive_pool3d_common/arch35/adaptive_avg_pool3d_simt.h"19#include "../adaptive_pool3d_common/arch35/adaptive_avg_pool3d_simt.h"
19#include "../adaptive_pool3d_common/arch35/adaptive_avg_pool3d_parall_pool.h"20#include "../adaptive_pool3d_common/arch35/adaptive_avg_pool3d_parall_pool.h"
20#include "../adaptive_pool3d_common/arch35/adaptive_pool3d_tiling_struct.h"21#include "../adaptive_pool3d_common/arch35/adaptive_pool3d_tiling_struct.h"
@@ -47,6 +48,11 @@ __global__ __aicore__ void adaptive_avg_pool3d(
47 AdaptivePool3d::AdaptiveAvgPool3dParaPool<DTYPE_X, int64_t> op(tilingData, pipeBase);48 AdaptivePool3d::AdaptiveAvgPool3dParaPool<DTYPE_X, int64_t> op(tilingData, pipeBase);
48 op.Init(x, y);49 op.Init(x, y);
49 op.Process();50 op.Process();
51+ } else if constexpr (TEMPLATE_MODE == TPL_MODE_1 && DYTPE_MODE == TPL_DTYPE_0 && MULTI_MODE == TPL_MULTI_MODE_0) {
52+ GET_TILING_DATA_WITH_STRUCT(AdaptivePool3DTiling::AdaptivePool3dBigKernelTilingData, tilingData, tiling);
53+ AdaptivePool3d::AdaptiveAvgPool3dBigKernel<DTYPE_X> op(tilingData, pipeBase);
54+ op.Init(x, y);
55+ op.Process();
50 }56 }
51 return;57 return;
52}58}
Apooling/adaptive_pool3d_common/op_host/arch35/adaptive_avg_pool3d_big_kernel_tiling.cpp+152-0
@@ -0,0 +1,152 @@
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 adaptive_avg_pool3d_tiling_big_kernel.cpp
13+ * \brief
14+ */
15+#include <cctype>
16+#include <algorithm>
17+#include "log/log.h"
18+#include "util/math_util.h"
19+#include "error_util.h"
20+#include "op_host/tiling_base.h"
21+#include "op_host/tiling_templates_registry.h"
22+#include "adaptive_avg_pool3d_big_kernel_tiling.h"
23+#include "op_common/op_host/util/platform_util.h"
24+#include "platform/platform_ascendc.h"
25+#include "register/op_def_registry.h"
26+#include "platform/platform_info.h"
27+#include "register/op_impl_registry.h"
28+ 
29+namespace optiling {
30+ 
31+static constexpr int64_t ADAPTIVE_AVG_POOL3D_BIG_KERNEL_THERSHOLD = 128;
32+static constexpr int64_t UB_MAX_INDICES_USE_COUNT = 1024;
33+static constexpr int64_t BUFFER_NUM = 2;
34+static constexpr int64_t FLOAT16_BYTPES = 2;
35+static constexpr int64_t FLOAT32_BYTPES = 4;
36+static constexpr int64_t STORE_ADD_SIZE = 1024;
37+ 
38+ge::graphStatus AdaptiveAvgPool3dBigKernelTiling::CheckOutputDtypeInfo()
39+{
40+ auto opNodeName = context_->GetNodeName();
41+ OP_LOGD(opNodeName, "CheckOutputDtypeInfo begin.");
42+ 
43+ auto outputShape = context_->GetOutputShape(0);
44+ OP_CHECK_NULL_WITH_CONTEXT(context_, outputShape);
45+ auto outputDesc = context_->GetOutputDesc(0);
46+ OP_CHECK_NULL_WITH_CONTEXT(context_, outputDesc);
47+ auto outputDtype = outputDesc->GetDataType();
48+ OP_CHECK_IF((outputDtype != ge::DT_FLOAT && outputDtype != ge::DT_FLOAT16 && outputDtype != ge::DT_BF16),
49+ OP_LOGE(opNodeName, "output datatype only support only supports float, float16, bfloat16"), return ge::GRAPH_FAILED);
50+ return ge::GRAPH_SUCCESS;
51+}
52+ 
53+bool AdaptiveAvgPool3dBigKernelTiling::IsCapable()
54+{
55+ OP_TILING_CHECK(
56+ GetAndCheckDataFormat() != ge::GRAPH_SUCCESS,
57+ VECTOR_INNER_ERR_REPORT_TILIING(context_, "GetDataFormatAttrInfo fail."),
58+ return ge::GRAPH_FAILED);
59+ // 按照搬运对齐的大小全载UB, 判断是否走当前模板
60+ OP_LOGD(context_->GetNodeName(), "AdaptiveAvgPool3dBigKernelTiling IsCapable check.");
61+ uint64_t kernelDMax = CalKernelSizeOneDimMax(input_.dIn, input_.dOut);
62+ uint64_t kernelHMax = CalKernelSizeOneDimMax(input_.hIn, input_.hOut);
63+ uint64_t kernelWMax = CalKernelSizeOneDimMax(input_.wIn, input_.wOut);
64+ avgBigKernelInfo.kernelMaxDHW = kernelDMax * kernelHMax * kernelWMax;
65+ bool isCapable = (avgBigKernelInfo.kernelMaxDHW >= ADAPTIVE_AVG_POOL3D_BIG_KERNEL_THERSHOLD) && (input_.dataFormat == ge::Format::FORMAT_NCDHW);
66+ OP_LOGD(context_->GetNodeName(), "AdaptiveAvgPool3dBigKernelTiling IsCapable check: %s", isCapable ? "true" : "false");
67+ return isCapable;
68+}
69+ 
70+uint64_t AdaptiveAvgPool3dBigKernelTiling::GetTilingKey() const
71+{
72+ return GET_TPL_TILING_KEY(TPL_MODE_1, TPL_DTYPE_0, TPL_MULTI_MODE_0, TPL_DATA_FORMAT_MODE_1);
73+}
74+ 
75+void AdaptiveAvgPool3dBigKernelTiling::DoBlockTiling()
76+{
77+ avgBigKernelInfo.totalIdx = input_.nIn * input_.cIn * input_.dOut * input_.hOut * input_.wOut;
78+ avgBigKernelInfo.blockFactor = avgBigKernelInfo.totalIdx / input_.coreNum;
79+ avgBigKernelInfo.blockTail = avgBigKernelInfo.totalIdx % input_.coreNum;
80+ 
81+ if (avgBigKernelInfo.blockFactor == 0) {
82+ avgBigKernelInfo.coreNums = avgBigKernelInfo.totalIdx;
83+ } else {
84+ avgBigKernelInfo.coreNums = input_.coreNum;
85+ }
86+ 
87+ int64_t vRegSize = Ops::Base::GetVRegSize(context_);
88+ auto xDtypeSize = input_.xDtype == ge::DT_FLOAT ? FLOAT32_BYTPES : FLOAT16_BYTPES;
89+ int64_t ubAvailable = input_.ubSize - xDtypeSize * UB_MAX_INDICES_USE_COUNT - FLOAT32_BYTPES * STORE_ADD_SIZE;
90+ int64_t defaultMaxSize = Ops::Base::FloorAlign(ubAvailable / BUFFER_NUM, vRegSize);
91+ avgBigKernelInfo.maxCount = defaultMaxSize / FLOAT32_BYTPES;
92+ avgBigKernelInfo.batchCount = avgBigKernelInfo.maxCount / avgBigKernelInfo.kernelMaxDHW;
93+}
94+ 
95+void AdaptiveAvgPool3dBigKernelTiling::PrintTilingData() const
96+{
97+ std::ostringstream info;
98+ info << "nc: " << input_.nIn * input_.cIn;
99+ info << ", dInDim: " << input_.dIn;
100+ info << ", hInDim: " << input_.hIn;
101+ info << ", wInDim: " << input_.wIn;
102+ info << ", dOutDim: " << input_.dOut;
103+ info << ", hOutDim: " << input_.hOut;
104+ info << ", wOutDim: " << input_.wOut;
105+ info << ", coreNums: " << avgBigKernelInfo.coreNums;
106+ info << ", blockFactor: " << avgBigKernelInfo.blockFactor;
107+ info << ", blockTail: " << avgBigKernelInfo.blockTail;
108+ info << ", totalIdx: " << avgBigKernelInfo.totalIdx;
109+ info << ", maxCount: " << avgBigKernelInfo.maxCount;
110+ info << ", batchCount: " << avgBigKernelInfo.batchCount;
111+ info << std::endl;
112+ 
113+ OP_LOGI("AdaptiveAvgPool3dBigKernel", "%s", info.str().c_str());
114+}
115+void AdaptiveAvgPool3dBigKernelTiling::SetTilingData()
116+{
117+ AdaptivePool3DTiling::AdaptivePool3dBigKernelTilingData* tilingData = context_->GetTilingData<AdaptivePool3dBigKernelTilingData>();
118+ tilingData->nc = input_.nIn * input_.cIn;
119+ tilingData->dInDim = input_.dIn;
120+ tilingData->hInDim = input_.hIn;
121+ tilingData->wInDim = input_.wIn;
122+ tilingData->dOutDim = input_.dOut;
123+ tilingData->hOutDim = input_.hOut;
124+ tilingData->wOutDim = input_.wOut;
125+ tilingData->blockFactor = avgBigKernelInfo.blockFactor;
126+ tilingData->blockTail = avgBigKernelInfo.blockTail;
127+ tilingData->totalIdx = avgBigKernelInfo.totalIdx;
128+ tilingData->coreNums = avgBigKernelInfo.coreNums;
129+ tilingData->maxCount = avgBigKernelInfo.maxCount;
130+ tilingData->batchCount = avgBigKernelInfo.batchCount;
131+}
132+ 
133+ge::graphStatus AdaptiveAvgPool3dBigKernelTiling::DoOpTiling()
134+{
135+ OP_LOGD(context_->GetNodeName(), "AdaptiveAvgPool3dBigKernelTiling DoOpTiling start.");
136+ OP_CHECK_IF(CheckOutputDtypeInfo() != ge::GRAPH_SUCCESS,
137+ OP_LOGE(context_->GetNodeName(), "AdaptiveAvgPool3d indices dtype unexpected"), return ge::GRAPH_FAILED);
138+ DoBlockTiling();
139+ SetTilingData();
140+ PrintTilingData();
141+ return ge::GRAPH_SUCCESS;
142+}
143+ 
144+ge::graphStatus AdaptiveAvgPool3dBigKernelTiling::PostTiling()
145+{
146+ context_->SetBlockDim(avgBigKernelInfo.coreNums);
147+ return ge::GRAPH_SUCCESS;
148+}
149+ 
150+REGISTER_OPS_TILING_TEMPLATE(AdaptiveAvgPool3d, AdaptiveAvgPool3dBigKernelTiling, 1);
151+ 
152+}
Apooling/adaptive_pool3d_common/op_host/arch35/adaptive_avg_pool3d_big_kernel_tiling.h+60-0
@@ -0,0 +1,60 @@
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 adaptive_avg_pool3d_big_kernel_tiling.h
13+ * \brief
14+ */
15+ 
16+#ifndef AIR_CXX_RUNTIME_V2_OP_IMPL_ADAPTIVE_AVG_POOL3D_TILING_BIG_KERNEL_H_
17+#define AIR_CXX_RUNTIME_V2_OP_IMPL_ADAPTIVE_AVG_POOL3D_TILING_BIG_KERNEL_H_
18+ 
19+#include "register/op_def_registry.h"
20+#include "tiling/tiling_api.h"
21+#include "op_host/tiling_base.h"
22+#include "util/math_util.h"
23+#include "op_common/op_host/util/platform_util.h"
24+#include "adaptive_pool3d_tiling.h"
25+#include "../op_kernel/arch35/adaptive_pool3d_tiling_struct.h"
26+ 
27+using namespace std;
28+using namespace AdaptivePool3DTiling;
29+ 
30+namespace optiling {
31+ 
32+struct AdaptiveAvgPool3dBigKernelInfo {
33+ int64_t blockFactor {0};
34+ int64_t blockTail {0};
35+ int64_t totalIdx {0};
36+ int64_t coreNums {0};
37+ int64_t maxCount {0};
38+ int64_t kernelMaxDHW {0};
39+ int64_t batchCount {1};
40+};
41+ 
42+class AdaptiveAvgPool3dBigKernelTiling : public AdaptivePool3dBaseTiling {
43+public:
44+ explicit AdaptiveAvgPool3dBigKernelTiling(gert::TilingContext* context) : AdaptivePool3dBaseTiling(context) {}
45+ 
46+private:
47+ ge::graphStatus CheckOutputDtypeInfo();
48+ void DoBlockTiling();
49+ void SetTilingData();
50+ int64_t CalKernelSize(int64_t inSize, int64_t outSize);
51+ void PrintTilingData() const;
52+ bool IsCapable() override;
53+ ge::graphStatus DoOpTiling() override;
54+ ge::graphStatus PostTiling() override;
55+ uint64_t GetTilingKey() const override;
56+
57+ AdaptiveAvgPool3dBigKernelInfo avgBigKernelInfo;
58+};
59+} // namespace optiling
60+#endif
Apooling/adaptive_pool3d_common/op_kernel/arch35/adaptive_avg_pool3d_big_kernel.h+342-0
@@ -0,0 +1,342 @@
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 adaptive_avg_pool3d_big_kernel.h
13+ * \brief
14+ */
15+#ifndef ADAPTIVE_AVG_POOL3D_BIG_KERNEL_H_
16+#define ADAPTIVE_AVG_POOL3D_BIG_KERNEL_H_
17+ 
18+#include "adaptive_pool3d_big_kernel.h"
19+ 
20+namespace AdaptivePool3d{
21+using namespace AscendC;
22+constexpr int32_t STORE_ADD_BUFFER = 1024;
23+ 
24+template <typename T, typename U>
25+__aicore__ inline void StoreOneValue(const __local_mem__ void* dstAddr, MicroAPI::RegTensor<U>& srcReg,
26+ MicroAPI::MaskReg& maskReg, uint32_t offset)
27+{
28+ auto addr = (__local_mem__ T*)dstAddr + offset;
29+ if constexpr (IsSameType<T, half>::value) {
30+ MicroAPI::RegTensor<half> regfp16;
31+ MicroAPI::Cast<half, float, CASTB4TOB2>(regfp16, srcReg, maskReg);
32+ MicroAPI::DataCopy<half, MicroAPI::StoreDist::DIST_FIRST_ELEMENT_B16>(addr, regfp16, maskReg);
33+ } else if constexpr (IsSameType<T, bfloat16_t>::value) {
34+ MicroAPI::RegTensor<bfloat16_t> regBf16;
35+ MicroAPI::Cast<bfloat16_t, float, CASTB4TOB2>(regBf16, srcReg, maskReg);
36+ MicroAPI::DataCopy<bfloat16_t, MicroAPI::StoreDist::DIST_FIRST_ELEMENT_B16>(addr, regBf16, maskReg);
37+ } else if constexpr (sizeof(T) == DIGHT4) {
38+ MicroAPI::DataCopy<T, MicroAPI::StoreDist::DIST_FIRST_ELEMENT_B32>(addr, (MicroAPI::RegTensor<T>&)srcReg, maskReg);
39+ } else {
40+ MicroAPI::UnalignReg uReg;
41+ MicroAPI::DataCopyUnAlign(addr, srcReg, uReg, 1);
42+ MicroAPI::DataCopyUnAlignPost(addr, uReg, 0);
43+ }
44+}
45+ 
46+template <typename U>
47+__aicore__ inline void LoadOneValue(const __local_mem__ void* srcAddr, MicroAPI::RegTensor<U>& dstReg,
48+ MicroAPI::MaskReg& preg, uint32_t offset)
49+{
50+ auto addr = (__local_mem__ U*)srcAddr + offset;
51+ if constexpr (sizeof(U) == DIGHT4) {
52+ MicroAPI::DataCopy<U, MicroAPI::LoadDist::DIST_BRC_B32>(dstReg, addr);
53+ } else {
54+ MicroAPI::UnalignReg ureg;
55+ MicroAPI::DataCopyUnAlignPre(ureg, addr);
56+ MicroAPI::DataCopyUnAlign(dstReg, ureg, addr, 1);
57+ }
58+}
59+ 
60+template <typename T, typename U>
61+__aicore__ inline void LoadXLocalToReg(const __local_mem__ void* srcAddr, MicroAPI::RegTensor<U>& dstReg,
62+ MicroAPI::MaskReg& preg, MicroAPI::AddrReg& offset)
63+{
64+ if constexpr (IsSameType<T, half>::value) {
65+ MicroAPI::RegTensor<half> regfp16;
66+ MicroAPI::DataCopy<half, MicroAPI::LoadDist::DIST_UNPACK_B16>(regfp16, (__local_mem__ half*)srcAddr, offset);
67+ MicroAPI::Cast<float, half, CASTB2TOB4>(dstReg, regfp16, preg);
68+ } else if constexpr (IsSameType<T, bfloat16_t>::value) {
69+ MicroAPI::RegTensor<bfloat16_t> regBf16;
70+ MicroAPI::DataCopy<bfloat16_t, MicroAPI::LoadDist::DIST_UNPACK_B16>(regBf16, (__local_mem__ bfloat16_t*)srcAddr, offset);
71+ MicroAPI::Cast<float, bfloat16_t, CASTB2TOB4>(dstReg, regBf16, preg);
72+ } else {
73+ MicroAPI::DataCopy(dstReg, (__local_mem__ float*)srcAddr, offset);
74+ }
75+}
76+ 
77+template <typename U>
78+__aicore__ inline void UpdateSum(MicroAPI::RegTensor<U>& res, const __local_mem__ U* storeLocalAddr, int32_t offset)
79+{
80+ // get data from local mem
81+ MicroAPI::MaskReg pregOne = MicroAPI::CreateMask<U, MicroAPI::MaskPattern::VL1>();
82+ MicroAPI::RegTensor<U> lastRes;
83+ 
84+ // get last res from local mem
85+ LoadOneValue<U>(storeLocalAddr, lastRes, pregOne, offset);
86+ 
87+ //calc sum
88+ MicroAPI::Add(res, res, lastRes, pregOne);
89+ MicroAPI::LocalMemBar<MicroAPI::MemType::VEC_LOAD, MicroAPI::MemType::VEC_STORE>();
90+}
91+ 
92+template <typename T>
93+class AdaptiveAvgPool3dBigKernel : public AdaptivePool3dBigKernel<T>
94+{
95+public:
96+ __aicore__ inline AdaptiveAvgPool3dBigKernel(const AdaptivePool3DTiling::AdaptivePool3dBigKernelTilingData &tilingData, TPipe &pipe) :
97+ AdaptivePool3dBigKernel<T>(tilingData, pipe) {};
98+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y);
99+ __aicore__ inline void Process();
100+ 
101+private:
102+ __aicore__ inline void InitOutputBuffer();
103+ __aicore__ inline void BaseCompute(int64_t curIdx);
104+ __aicore__ inline void NoSplitProcess(int64_t curIdx);
105+ __aicore__ inline void SplitProcess(int64_t curIdx);
106+ __aicore__ inline void ComputeSplitD(int64_t curIdx);
107+ __aicore__ inline void ComputeSplitH(int64_t curIdx);
108+ __aicore__ inline void ComputeSplitW(int64_t curIdx);
109+ template <int32_t SPLIT_MODE, typename U>
110+ __aicore__ inline void ComputeSum(LocalTensor<T> xLocal, int64_t localCurIdx, int64_t dataCount);
111+ template <typename U>
112+ __aicore__ inline void ComputeAvg(LocalTensor<U> storeAddLocal, int64_t curIdx);
113+ 
114+protected:
115+ TBuf<QuePosition::VECCALC> storeAddUB_;
116+};
117+ 
118+template <typename T>
119+__aicore__ inline void AdaptiveAvgPool3dBigKernel<T>::InitOutputBuffer()
120+{
121+ event_t eventIdMTE3toV = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE3_V));
122+ SetFlag<HardEvent::MTE3_V>(eventIdMTE3toV);
123+ WaitFlag<HardEvent::MTE3_V>(eventIdMTE3toV);
124+ LocalTensor<T> avgOutLocal = this->outputUB_.template Get<T>();
125+ LocalTensor<float> avgStoreOutLocal = this->storeAddUB_.template Get<float>();
126+ __local_mem__ T* avgOutAddr = (__local_mem__ T*)avgOutLocal.GetPhyAddr();
127+ __local_mem__ float* avgStoreOutAddr = (__local_mem__ float*)avgStoreOutLocal.GetPhyAddr();
128+ 
129+ uint32_t maxOutCount = BATCH_COPYOUT_COUNT;
130+ uint32_t maxVfCount = platform::GetVRegSize() / sizeof(T);
131+ uint16_t repeatMaxTimes = ops::CeilDiv(static_cast<uint32_t>(maxOutCount), maxVfCount);
132+ 
133+ __VEC_SCOPE__
134+ {
135+ MicroAPI::RegTensor<T> avgOutReg;
136+ MicroAPI::RegTensor<float> avgStoreOutReg;
137+ MicroAPI::Duplicate(avgOutReg, static_cast<T>(0));
138+ MicroAPI::Duplicate(avgStoreOutReg, static_cast<float>(0));
139+ for (uint16_t i = 0; i < repeatMaxTimes; i++) {
140+ MicroAPI::MaskReg avgOutMask = MicroAPI::UpdateMask<T>(maxOutCount);
141+ MicroAPI::MaskReg avgStoreOutMask = MicroAPI::UpdateMask<float>(maxOutCount);
142+ MicroAPI::AddrReg offsetReg = MicroAPI::CreateAddrReg<T>(i, maxVfCount);
143+ MicroAPI::AddrReg offsetStoreReg = MicroAPI::CreateAddrReg<float>(i, maxVfCount);
144+ MicroAPI::DataCopy(avgOutAddr, avgOutReg, offsetReg, avgOutMask);
145+ MicroAPI::DataCopy(avgStoreOutAddr, avgStoreOutReg, offsetStoreReg, avgStoreOutMask);
146+ }
147+ }
148+}
149+ 
150+template<typename T>
151+template<typename U>
152+__aicore__ inline void AdaptiveAvgPool3dBigKernel<T>::ComputeAvg(LocalTensor<U> storeAddLocal, int64_t curIdx)
153+{
154+ LocalTensor<T> outputLocal = this->outputUB_.template Get<T>();
155+ __local_mem__ U* storeLocalAddr = (__local_mem__ U*)storeAddLocal.GetPhyAddr();
156+ __local_mem__ T* dstLocalAddr = (__local_mem__ T*)outputLocal.GetPhyAddr();
157+ U divNum = static_cast<U>(this->curkDHW_);
158+ 
159+ __VEC_SCOPE__
160+ {
161+ MicroAPI::MaskReg pregOne = MicroAPI::CreateMask<U, MicroAPI::MaskPattern::VL1>();
162+ MicroAPI::RegTensor<U> disiv;
163+ MicroAPI::RegTensor<U> lastRes;
164+ 
165+ MicroAPI::Duplicate(disiv, divNum);
166+ LoadOneValue<U>(storeLocalAddr, lastRes, pregOne, curIdx);
167+ MicroAPI::Div(lastRes, lastRes, disiv, pregOne);
168+ 
169+ StoreOneValue<T, U>(dstLocalAddr, lastRes, pregOne, curIdx);
170+ }
171+}
172+ 
173+template <typename T>
174+template <int32_t SPLIT_MODE, typename U>
175+__aicore__ inline void AdaptiveAvgPool3dBigKernel<T>::ComputeSum(
176+ LocalTensor<T> xLocal, int64_t localCurIdx, int64_t dataCount)
177+{
178+ LocalTensor<U> storeAddLocal = this->storeAddUB_.template Get<U>();
179+ __local_mem__ T* xLocalAddr = (__local_mem__ T*)xLocal.GetPhyAddr();
180+ __local_mem__ U* storeLocalAddr = (__local_mem__ U*)storeAddLocal.GetPhyAddr();
181+
182+ uint32_t repeatCount = platform::GetVRegSize() / sizeof(U); //一个vf需要的次数
183+ uint16_t repeatTimes = ops::CeilDiv(static_cast<uint32_t>(dataCount), repeatCount); //上取整，获取repeatCount的整数倍
184+ uint32_t dataCount_ = dataCount;
185+ 
186+ __VEC_SCOPE__
187+ {
188+ MicroAPI::RegTensor<U> vd0;
189+ MicroAPI::RegTensor<U> vd1;
190+ MicroAPI::RegTensor<U> res;
191+ MicroAPI::MaskReg sumMask = MicroAPI::CreateMask<U, MicroAPI::MaskPattern::VL1>();
192+ MicroAPI::Duplicate(res, static_cast<U>(0));
193+ for (uint16_t i = 0; i < repeatTimes; i++) {
194+ MicroAPI::MaskReg p0 = MicroAPI::UpdateMask<U>(dataCount_); //一次处理数量
195+ MicroAPI::AddrReg offset = MicroAPI::CreateAddrReg<T>(i, repeatCount); //搬运偏移
196+ LoadXLocalToReg<T, U>(xLocalAddr, vd0, p0, offset);
197+ MicroAPI::ReduceSum(vd1, vd0, p0);
198+ MicroAPI::Add(res, res, vd1, sumMask);
199+ }
200+ if constexpr (SPLIT_MODE != NO_SPLIT) {
201+ UpdateSum<U>(res, storeLocalAddr, localCurIdx);
202+ }
203+ StoreOneValue<U, U>(storeLocalAddr, res, sumMask, localCurIdx);
204+ }
205+}
206+ 
207+template <typename T>
208+__aicore__ inline void AdaptiveAvgPool3dBigKernel<T>::ComputeSplitD(int64_t curIdx)
209+{
210+ int64_t dFactor = this->tilingData_.maxCount / this->curkHW_;
211+ int64_t dLoops = ops::CeilDiv(this->curkD_, dFactor);
212+ int64_t dTail = this->curkD_ - (dLoops - DIGHT1) * dFactor;
213+ int64_t inputOffset = this->curInOffset_;
214+ for (int64_t dLoop = 0; dLoop < dLoops; dLoop++) {
215+ int32_t curDFactor = dLoop == (dLoops - 1) ? dTail : dFactor;
216+ AdaptivePool3dBigKernel<T>::CopyIn(inputOffset, this->curkW_, this->curkH_, curDFactor);
217+ LocalTensor<T> xLocal = this->inputQue_.template DeQue<T>();
218+ ComputeSum<SPLIT_D, float>(xLocal, curIdx, this->curkHW_ * curDFactor);
219+ inputOffset += curDFactor * this->inHW_;
220+ this->inputQue_.template FreeTensor<T>(xLocal);
221+ }
222+}
223+ 
224+template <typename T>
225+__aicore__ inline void AdaptiveAvgPool3dBigKernel<T>::ComputeSplitH(int64_t curIdx)
226+{
227+ int64_t hFactor = this->tilingData_.maxCount / this->curkW_;
228+ int64_t hLoops = ops::CeilDiv(this->curkH_, hFactor);
229+ int64_t hTail = this->curkH_ - (hLoops - DIGHT1) * hFactor;
230+ for (int64_t dLoop = 0; dLoop < this->curkD_; dLoop++) {
231+ int64_t inputOffset = this->curInOffset_ + dLoop * this->inHW_;
232+ for (int64_t hLoop = 0; hLoop < hLoops; hLoop++) {
233+ int64_t curHFactor = hLoop == (hLoops - 1) ? hTail : hFactor;
234+ AdaptivePool3dBigKernel<T>::CopyIn(inputOffset, this->curkW_, curHFactor, DIGHT1);
235+ LocalTensor<T> xLocal = this->inputQue_.template DeQue<T>();
236+ ComputeSum<SPLIT_H, float>(xLocal, curIdx, this->curkW_ * curHFactor);
237+ inputOffset += hFactor * this->tilingData_.wInDim;
238+ this->inputQue_.template FreeTensor<T>(xLocal);
239+ }
240+ }
241+}
242+ 
243+template <typename T>
244+__aicore__ inline void AdaptiveAvgPool3dBigKernel<T>::ComputeSplitW(int64_t curIdx)
245+{
246+ int64_t wFactor = this->tilingData_.maxCount;
247+ int64_t wLoops = ops::CeilDiv(this->curkW_, wFactor);
248+ int64_t wTail = this->curkW_ - (wLoops - DIGHT1) * wFactor;
249+ for (int64_t dLoop = 0; dLoop < this->curkD_; dLoop++) {
250+ int64_t dOffset = this->curInOffset_ + dLoop * this->inHW_;
251+ for (int64_t hLoop = 0; hLoop < this->curkH_; hLoop++) {
252+ int64_t inputOffset = dOffset + hLoop * this->tilingData_.wInDim;
253+ for (int64_t wLoop = 0; wLoop < wLoops; wLoop++) {
254+ int64_t curWFactor = wLoop == (wLoops - 1) ? wTail : wFactor;
255+ AdaptivePool3dBigKernel<T>::CopyIn(inputOffset, curWFactor, DIGHT1, DIGHT1);
256+ LocalTensor<T> xLocal = this->inputQue_.template DeQue<T>();
257+ ComputeSum<SPLIT_W, float>(xLocal, curIdx, curWFactor);
258+ inputOffset += curWFactor;
259+ this->inputQue_.template FreeTensor<T>(xLocal);
260+ }
261+ }
262+ }
263+}
264+ 
265+template <typename T>
266+__aicore__ inline void AdaptiveAvgPool3dBigKernel<T>::NoSplitProcess(int64_t curIdx)
267+{
268+ AdaptivePool3dBigKernel<T>::CopyIn(this->curInOffset_, this->curkW_, this->curkH_, this->curkD_);
269+ LocalTensor<T> xLocal = this->inputQue_.template DeQue<T>();
270+ ComputeSum<NO_SPLIT, float>(xLocal, curIdx, this->curkDHW_);
271+ this->inputQue_.template FreeTensor<T>(xLocal);
272+}
273+ 
274+template <typename T>
275+__aicore__ inline void AdaptiveAvgPool3dBigKernel<T>::SplitProcess(int64_t curIdx)
276+{
277+ if (this->curkHW_ <= this->tilingData_.maxCount) {
278+ ComputeSplitD(curIdx);
279+ } else if (this->curkW_ <= this->tilingData_.maxCount) {
280+ ComputeSplitH(curIdx);
281+ } else {
282+ ComputeSplitW(curIdx);
283+ }
284+}
285+ 
286+template <typename T>
287+__aicore__ inline void AdaptiveAvgPool3dBigKernel<T>::BaseCompute(int64_t curIdx)
288+{
289+ LocalTensor<float> storeAddLocal = this->storeAddUB_.template DeQue<float>();
290+ if (this->curkDHW_ <= this->tilingData_.maxCount) {
291+ NoSplitProcess(curIdx);
292+ } else {
293+ SplitProcess(curIdx);
294+ }
295+ ComputeAvg<float>(storeAddLocal, curIdx);
296+}
297+ 
298+template <typename T>
299+__aicore__ inline void AdaptiveAvgPool3dBigKernel<T>::Init(GM_ADDR x, GM_ADDR y)
300+{
301+ // AdaptivePool3dBigKernel init
302+ AdaptivePool3dBigKernel<T>::Init(x, y);
303+ this->pipe_.InitBuffer(storeAddUB_, STORE_ADD_BUFFER);
304+ // set half overflow
305+ if constexpr (IsSameType<T, half>::value) {
306+ SetCtrlSpr<HALF_OVERFLOW_MODE_CTRL, HALF_OVERFLOW_MODE_CTRL>(1);
307+ }
308+}
309+ 
310+template <typename T>
311+__aicore__ inline void AdaptiveAvgPool3dBigKernel<T>::Process()
312+{
313+ int64_t beginIdx = 0;
314+ int64_t endIdx = 0;
315+ if (GetBlockIdx() < this->tilingData_.blockTail) {
316+ beginIdx = GetBlockIdx() * (this->tilingData_.blockFactor + 1);
317+ endIdx = beginIdx + this->tilingData_.blockFactor + 1;
318+ } else {
319+ beginIdx = GetBlockIdx() * this->tilingData_.blockFactor + this->tilingData_.blockTail;
320+ endIdx = beginIdx + this->tilingData_.blockFactor;
321+ }
322+ 
323+ InitOutputBuffer();
324+ int64_t curLocalIdx = 0;
325+ int64_t outputOffset = beginIdx;
326+ for (int64_t outIdx = beginIdx; outIdx < endIdx; outIdx++) {
327+ AdaptivePool3dBigKernel<T>::CalcWindowSize(outIdx);
328+ BaseCompute(curLocalIdx);
329+ curLocalIdx++;
330+ if (curLocalIdx == BATCH_COPYOUT_COUNT) {
331+ AdaptivePool3dBigKernel<T>::CopyOut(curLocalIdx, outputOffset);
332+ InitOutputBuffer();
333+ outputOffset = outIdx + 1;
334+ curLocalIdx = 0;
335+ }
336+ }
337+ if (curLocalIdx != 0) {
338+ AdaptivePool3dBigKernel<T>::CopyOut(curLocalIdx, outputOffset);
339+ }
340+}
341+} // namespace AdaptivePool3d
342+#endif // ADAPTIVE_AVG_POOL3D_BIG_KERNEL_H
Mpooling/adaptive_pool3d_common/op_kernel/arch35/adaptive_max_pool3d_big_kernel.h+4-4
@@ -239,7 +239,7 @@ private:
239 __aicore__ inline void CopyOutIndices(int64_t copyCount, int64_t offset);239 __aicore__ inline void CopyOutIndices(int64_t copyCount, int64_t offset);
240 240 
241 // indices need241 // indices need
242- TBuf<> indicesUB_;242+ TBuf<QuePosition::VECCALC> indicesUB_;
243 GlobalTensor<TINDEX> indicesGm_;243 GlobalTensor<TINDEX> indicesGm_;
244};244};
245 245 
@@ -249,7 +249,7 @@ __aicore__ inline void AdaptiveMaxPool3dBigKernel<T, TINDEX>::InitOutputBuffer()
249 event_t eventIdMTE3toV = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE3_V));249 event_t eventIdMTE3toV = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE3_V));
250 SetFlag<HardEvent::MTE3_V>(eventIdMTE3toV);250 SetFlag<HardEvent::MTE3_V>(eventIdMTE3toV);
251 WaitFlag<HardEvent::MTE3_V>(eventIdMTE3toV);251 WaitFlag<HardEvent::MTE3_V>(eventIdMTE3toV);
252- LocalTensor<T> maxOutLocal = this->maxUB_.template Get<T>();252+ LocalTensor<T> maxOutLocal = this->outputUB_.template Get<T>();
253 __local_mem__ T* maxOutAddr = (__local_mem__ T*)maxOutLocal.GetPhyAddr();253 __local_mem__ T* maxOutAddr = (__local_mem__ T*)maxOutLocal.GetPhyAddr();
254 LocalTensor<TINDEX> idxOutLocal = indicesUB_.Get<TINDEX>();254 LocalTensor<TINDEX> idxOutLocal = indicesUB_.Get<TINDEX>();
255 __local_mem__ TINDEX* idxOutAddr = (__local_mem__ TINDEX*)idxOutLocal.GetPhyAddr();255 __local_mem__ TINDEX* idxOutAddr = (__local_mem__ TINDEX*)idxOutLocal.GetPhyAddr();
@@ -302,10 +302,10 @@ template <int32_t SPLIT_MODE, typename U, typename UINDEX>
302__aicore__ inline void AdaptiveMaxPool3dBigKernel<T, TINDEX>::ComputeMax(302__aicore__ inline void AdaptiveMaxPool3dBigKernel<T, TINDEX>::ComputeMax(
303 LocalTensor<T> xLocal, int64_t localCurIdx, int64_t dataCount, int64_t curOffset)303 LocalTensor<T> xLocal, int64_t localCurIdx, int64_t dataCount, int64_t curOffset)
304{304{
305- LocalTensor<T> maxLocal = this->maxUB_.template Get<T>();305+ LocalTensor<T> outputLocal = this->outputUB_.template Get<T>();
306 LocalTensor<TINDEX> indicesLocal = indicesUB_.Get<TINDEX>();306 LocalTensor<TINDEX> indicesLocal = indicesUB_.Get<TINDEX>();
307 __local_mem__ T* xLocalAddr = (__local_mem__ T*)xLocal.GetPhyAddr();307 __local_mem__ T* xLocalAddr = (__local_mem__ T*)xLocal.GetPhyAddr();
308- __local_mem__ T* dstLocalAddr = (__local_mem__ T*)maxLocal.GetPhyAddr();308+ __local_mem__ T* dstLocalAddr = (__local_mem__ T*)outputLocal.GetPhyAddr();
309 __local_mem__ TINDEX* indicesLocalAddr = (__local_mem__ TINDEX*)indicesLocal.GetPhyAddr();309 __local_mem__ TINDEX* indicesLocalAddr = (__local_mem__ TINDEX*)indicesLocal.GetPhyAddr();
310 310
311 T minValue = this->minValue_;311 T minValue = this->minValue_;
Mpooling/adaptive_pool3d_common/op_kernel/arch35/adaptive_pool3d_big_kernel.h+4-4
@@ -80,7 +80,7 @@ private:
80protected:80protected:
81 TPipe pipe_;81 TPipe pipe_;
82 TQue<QuePosition::VECIN, USE_BUFFER_NUM> inputQue_;82 TQue<QuePosition::VECIN, USE_BUFFER_NUM> inputQue_;
83- TBuf<> maxUB_;83+ TBuf<QuePosition::VECCALC> outputUB_;
84 GlobalTensor<T> xGm_, yGm_;84 GlobalTensor<T> xGm_, yGm_;
85 const AdaptivePool3DTiling::AdaptivePool3dBigKernelTilingData tilingData_;85 const AdaptivePool3DTiling::AdaptivePool3dBigKernelTilingData tilingData_;
86 86 
@@ -149,13 +149,13 @@ __aicore__ inline void AdaptivePool3dBigKernel<T>::CopyOut(int64_t copyCount, in
149 event_t eventIdVtoMTE3 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE3));149 event_t eventIdVtoMTE3 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE3));
150 SetFlag<HardEvent::V_MTE3>(eventIdVtoMTE3);150 SetFlag<HardEvent::V_MTE3>(eventIdVtoMTE3);
151 WaitFlag<HardEvent::V_MTE3>(eventIdVtoMTE3);151 WaitFlag<HardEvent::V_MTE3>(eventIdVtoMTE3);
152- LocalTensor<T> maxLocal = maxUB_.Get<T>();152+ LocalTensor<T> outputLocal = outputUB_.Get<T>();
153 DataCopyExtParams extParams;153 DataCopyExtParams extParams;
154 extParams.blockCount = DIGHT1;154 extParams.blockCount = DIGHT1;
155 extParams.blockLen = copyCount * sizeof(T);155 extParams.blockLen = copyCount * sizeof(T);
156 extParams.srcStride = 0;156 extParams.srcStride = 0;
157 extParams.dstStride = 0;157 extParams.dstStride = 0;
158- DataCopyPad(yGm_[offset], maxLocal, extParams);158+ DataCopyPad(yGm_[offset], outputLocal, extParams);
159}159}
160 160 
161template <typename T>161template <typename T>
@@ -171,7 +171,7 @@ __aicore__ inline void AdaptivePool3dBigKernel<T>::Init(GM_ADDR x, GM_ADDR y)
171 xGm_.SetGlobalBuffer((__gm__ T*)x);171 xGm_.SetGlobalBuffer((__gm__ T*)x);
172 yGm_.SetGlobalBuffer((__gm__ T*)y);172 yGm_.SetGlobalBuffer((__gm__ T*)y);
173 pipe_.InitBuffer(inputQue_, USE_BUFFER_NUM, tilingData_.maxCount * sizeof(T));173 pipe_.InitBuffer(inputQue_, USE_BUFFER_NUM, tilingData_.maxCount * sizeof(T));
174- pipe_.InitBuffer(maxUB_, BATCH_COPYOUT_COUNT * sizeof(T));174+ pipe_.InitBuffer(outputUB_, BATCH_COPYOUT_COUNT * sizeof(T));
175}175}
176 176 
177template <typename T>177template <typename T>
Mpooling/adaptive_pool3d_common/op_kernel/arch35/adaptive_pool3d_tiling_struct.h+8-2
@@ -64,9 +64,15 @@ ASCENDC_TPL_SEL(
64 ASCENDC_TPL_UINT_SEL(FORMAT_MODE, ASCENDC_TPL_UI_LIST, TPL_DATA_FORMAT_MODE_0),64 ASCENDC_TPL_UINT_SEL(FORMAT_MODE, ASCENDC_TPL_UI_LIST, TPL_DATA_FORMAT_MODE_0),
65 ASCENDC_TPL_TILING_STRUCT_SEL(AdaptivePool3DTiling::AdaptivePool3dParaKernelTilingData)65 ASCENDC_TPL_TILING_STRUCT_SEL(AdaptivePool3DTiling::AdaptivePool3dParaKernelTilingData)
66 ),66 ),
67+ ASCENDC_TPL_ARGS_SEL(
68+ ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_AIV_ONLY),
69+ ASCENDC_TPL_UINT_SEL(TEMPLATE_MODE, ASCENDC_TPL_UI_LIST, TPL_MODE_1),
70+ ASCENDC_TPL_UINT_SEL(DYTPE_MODE, ASCENDC_TPL_UI_LIST, TPL_DTYPE_0),
71+ ASCENDC_TPL_UINT_SEL(MULTI_MODE, ASCENDC_TPL_UI_LIST, TPL_MULTI_MODE_0),
72+ ASCENDC_TPL_UINT_SEL(FORMAT_MODE, ASCENDC_TPL_UI_LIST, TPL_DATA_FORMAT_MODE_1)
73+ ASCENDC_TPL_TILING_STRUCT_SEL(AdaptivePool3DTiling::AdaptivePool3dBigKernelTilingData)
74+ ),
67);75);
68- 
69- 
70class AdaptivePool3DSimtTilingData {76class AdaptivePool3DSimtTilingData {
71public:77public:
72 int64_t nDim = 0;78 int64_t nDim = 0;