Star906
1825
代码
代码
Issues226
Pull Requests303
Actions
项目讨论
Wiki
分析
Star906
1825
已合并
FM非全载场景增加小kernel模板 #8376
yangyang4536创建于 8月6日
FM非全载场景增加小kernel模板 #8376
已合并
yangyang4536创建于 8月6日
从yangyang4536/ops-nnmaster合入到cann/ops-nnmaster
共 7 个文件变更+413-11
Mconv/common/op_kernel/arch35/conv_tilingkey.h+6-5
@@ -75,16 +75,17 @@ namespace ConvKey {
75 75 
76#define CONV_NOT_SMALL_KERNEL 076#define CONV_NOT_SMALL_KERNEL 0
77#define CONV_SMALL_KERNEL 177#define CONV_SMALL_KERNEL 1
78+#define CONV_SMALL_KERNEL_FM_PARTLOAD 2
78 79 
79#define CONV_NORMAL_KERNEL 080#define CONV_NORMAL_KERNEL 0
80#define CONV_BIG_KERNEL 181#define CONV_BIG_KERNEL 1
81 82 
82#if defined(FORMAT_FILTER) && (FORMAT_FILTER == FORMAT_FRACTAL_Z || FORMAT_FILTER == FORMAT_FRACTAL_Z_C04)83#if defined(FORMAT_FILTER) && (FORMAT_FILTER == FORMAT_FRACTAL_Z || FORMAT_FILTER == FORMAT_FRACTAL_Z_C04)
83-#define CONV2D_SCALAR_OPT_SEL(...) \84+#define CONV2D_SCALAR_OPT_SEL(...) \
84- , ASCENDC_TPL_UINT_SEL(BatchOne, ASCENDC_TPL_UI_LIST, CONV_MULTI_BATCH, CONV_ONE_BATCH), \85+ , ASCENDC_TPL_UINT_SEL(BatchOne, ASCENDC_TPL_UI_LIST, CONV_MULTI_BATCH, CONV_ONE_BATCH), \
85- ASCENDC_TPL_UINT_SEL(NoPad, ASCENDC_TPL_UI_LIST, CONV_HAS_PAD, CONV_NO_PAD), \86+ ASCENDC_TPL_UINT_SEL(NoPad, ASCENDC_TPL_UI_LIST, CONV_HAS_PAD, CONV_NO_PAD), \
86- ASCENDC_TPL_UINT_SEL(SmallWeight, ASCENDC_TPL_UI_LIST, __VA_ARGS__) \87+ ASCENDC_TPL_UINT_SEL(SmallWeight, ASCENDC_TPL_UI_LIST, __VA_ARGS__) ASCENDC_TPL_UINT_SEL( \
87- ASCENDC_TPL_UINT_SEL(SmallKernel, ASCENDC_TPL_UI_LIST, CONV_NOT_SMALL_KERNEL, CONV_SMALL_KERNEL)88+ SmallKernel, ASCENDC_TPL_UI_LIST, CONV_NOT_SMALL_KERNEL, CONV_SMALL_KERNEL, CONV_SMALL_KERNEL_FM_PARTLOAD)
88#else89#else
89#define CONV2D_SCALAR_OPT_SEL(...) \90#define CONV2D_SCALAR_OPT_SEL(...) \
90 , ASCENDC_TPL_UINT_SEL(BatchOne, ASCENDC_TPL_UI_LIST, 0), ASCENDC_TPL_UINT_SEL(NoPad, ASCENDC_TPL_UI_LIST, 0), \91 , ASCENDC_TPL_UINT_SEL(BatchOne, ASCENDC_TPL_UI_LIST, 0), ASCENDC_TPL_UINT_SEL(NoPad, ASCENDC_TPL_UI_LIST, 0), \
Mconv/conv2d_v2/op_host/op_tiling/arch35/conv2d_v2_base_tiling_tilingkey.cpp+11-0
@@ -197,6 +197,17 @@ uint64_t Conv2dBaseTiling::GetSmallKernelVal()
197 al1Fullload = kAL1FullloadFlag && hoL1FullloadFlag && woL1FullloadFlag;197 al1Fullload = kAL1FullloadFlag && hoL1FullloadFlag && woL1FullloadFlag;
198 }198 }
199 199 
200+ // FmPartload: FM not fullload L1, Weight fullload L1, NZ format, FP16*FP16 or INT8*INT8.
201+ // IsSmallKernelBlocked() already guarantees singleCoreBatch==1, pad<=kernel, not FP16*INT8, nL0==nBL1.
202+ bool fmPartloadCond = !al1Fullload && bl1Fullload && groupOk;
203+ bool dtypeFmPartloadOk = (descInfo_.fMapDtype == ge::DataType::DT_FLOAT16 &&
204+ descInfo_.weightDtype == ge::DataType::DT_FLOAT16) ||
205+ (descInfo_.fMapDtype == ge::DataType::DT_INT8 &&
206+ descInfo_.weightDtype == ge::DataType::DT_INT8);
207+ if (fmPartloadCond && dtypeFmPartloadOk && IsWeightNZFormat(descInfo_.weightFormat)) {
208+ return CONV_SMALL_KERNEL_FM_PARTLOAD;
209+ }
210+ 
200 if (!(al1Fullload && bl1Fullload && groupOk)) {211 if (!(al1Fullload && bl1Fullload && groupOk)) {
201 return CONV_NOT_SMALL_KERNEL;212 return CONV_NOT_SMALL_KERNEL;
202 }213 }
Aconv/conv2d_v2/op_kernel/arch35/conv2d_small_kernel_fm_partload.h+371-0
@@ -0,0 +1,371 @@
1+/**
2+ * Copyright (c) Huawei Technologies Co., Ltd. 2026-2026. All rights reserved.
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 CONV2D_SMALL_KERNEL_FM_PARTLOAD_H
12+#define CONV2D_SMALL_KERNEL_FM_PARTLOAD_H
13+ 
14+#include "conv2d_small_kernel_parallelism.h"
15+ 
16+using namespace AscendC;
17+ 
18+constexpr uint32_t FMP_BLOCK_BYTES = 32; // DataCopy block size in bytes
19+// cinL1_ is derived from tiling-side decided kAL1: cinL1_ = kAL1 / kernelHxkernelW
20+ 
21+// Entry constraints for this template:
22+// 1. singleCoreBatch == 1
23+// 2. FM not fullload L1 (FM is chunked over Cin), Weight fullload L1 (one-shot)
24+// 3. N axis fullload L0 (singleCoreCo loaded to L0B at once, no N-axis split)
25+// 4. pad <= kernel (no pad larger than kernel supported)
26+// 5. dtype: FP16*FP16 and INT8*INT8 only (FP16*INT8 not supported)
27+// Requirements:
28+// - tiling supports M mode and HW mode
29+// - supports both conv2dv2 and extendconv2d (via ExtendParams)
30+// - format supports NHWC and NCHW for both input and output
31+// - FM chunked loading (shared with Conv2dSmallKernelParallelism)
32+// - Weight one-shot loading into L1 (this subclass's only difference)
33+ 
34+namespace {
35+static constexpr event_t FMP_EVT_WBS_DONE = static_cast<event_t>(0);
36+} // namespace
37+ 
38+// FmPartload inherits from Parallelism to reuse the FM-chunked helpers
39+// (CalcChunkFmap / CalcChunkFmapW / LoadFmapL1Chunk / SetupLoad3DForChunk /
40+// PrepareCinBlock / RunKL0Loop / ProcessCinBlocks / CopyOutResult and shared members).
41+// It only overrides Init and Process, and adds its own LoadWeightL1Full:
42+// - Weight is loaded ONCE into L1 (one-shot), not per Cin chunk.
43+// - kL0 is taken directly from tiling (no per-chunk kL0 recomputation).
44+template <typename FmapType, typename weightType, typename biasType, typename out0Type, typename out1Type = half,
45+ bool isNHWCin = false, bool isNHWCout = false, bool IsHwMode = false>
46+class Conv2dSmallKernelFmPartload : public Conv2dSmallKernelParallelism<FmapType, weightType, biasType, out0Type,
47+ out1Type, isNHWCin, isNHWCout, IsHwMode> {
48+public:
49+ using BaseT = Conv2dSmallKernelParallelism<FmapType, weightType, biasType, out0Type, out1Type, isNHWCin, isNHWCout,
50+ IsHwMode>;
51+ using L0cT = typename BaseT::L0cT;
52+ __aicore__ inline void Init(const Conv2DTilingData& tiling);
53+ __aicore__ inline void Process(GM_ADDR x, GM_ADDR filter, GM_ADDR bias, GM_ADDR y,
54+ const ExtendParams* extendParams);
55+ 
56+private:
57+ __aicore__ inline void LoadWeightL1Full(GM_ADDR filter);
58+ __aicore__ inline void ProcessHwMode(uint32_t kL0, uint32_t kL0Iters, uint32_t kernelHxW, uint32_t mmadN,
59+ uint64_t hwOut, GM_ADDR y, const ExtendParams* extendParams,
60+ LocalTensor<weightType>& bl1Full);
61+ __aicore__ inline void ProcessMMode(uint32_t kL0, uint32_t kL0Iters, uint32_t kernelHxW, uint32_t mmadN,
62+ uint64_t hwOut, GM_ADDR y, const ExtendParams* extendParams,
63+ LocalTensor<weightType>& bl1Full);
64+ __aicore__ inline void ProcessCinBlocks(LocalTensor<L0cT>& cl0, MmadParams& mp, LocalTensor<weightType>& bl1Full,
65+ uint32_t kL0, uint32_t kL0Iters, uint32_t kernelHxW, uint32_t curHi,
66+ uint32_t padTop, uint32_t padBottom, uint32_t hiLoadOff, uint32_t curWi,
67+ uint32_t wiLoadOff, uint32_t curM, uint32_t setupMOff, uint32_t setupWoOff,
68+ int32_t padLeft, int32_t padRight, bool loadWeight);
69+};
70+ 
71+template <typename FmapType, typename weightType, typename biasType, typename out0Type, typename out1Type,
72+ bool isNHWCin, bool isNHWCout, bool IsHwMode>
73+__aicore__ inline void Conv2dSmallKernelFmPartload<FmapType, weightType, biasType, out0Type, out1Type, isNHWCin,
74+ isNHWCout, IsHwMode>::Init(const Conv2DTilingData& tiling)
75+{
76+ this->InitCommon(tiling);
77+ if (!this->coreActive_) {
78+ return;
79+ }
80+ 
81+ if constexpr (IsHwMode) {
82+ // HW-mode: fmap W range from this core's Wo range (base InitCommon sets orgWin_ = win).
83+ uint32_t woStart = this->woIdxStart_;
84+ uint32_t woEnd = this->woIdxStart_ + this->actualWo_ - 1;
85+ uint32_t wiStart = woStart * this->tiling_->strideW;
86+ uint32_t wiEnd = woEnd * this->tiling_->strideW + this->tiling_->dilationW * (this->tiling_->kw - 1);
87+ uint32_t wiTotal = wiEnd - wiStart + 1;
88+ 
89+ this->padLeftL1_ = 0;
90+ this->padRightL1_ = 0;
91+ this->curWiLoadL1_ = wiTotal;
92+ this->wiLoadStart_ = wiStart;
93+ if (wiStart < this->tiling_->padLeft) {
94+ this->padLeftL1_ = this->tiling_->padLeft - wiStart;
95+ this->curWiLoadL1_ -= this->padLeftL1_;
96+ this->wiLoadStart_ = 0;
97+ } else {
98+ this->wiLoadStart_ = wiStart - this->tiling_->padLeft;
99+ }
100+ if (wiEnd >= static_cast<uint32_t>(this->tiling_->win) + this->tiling_->padLeft) {
101+ this->padRightL1_ = wiEnd - (static_cast<uint32_t>(this->tiling_->win) + this->tiling_->padLeft) + 1;
102+ this->curWiLoadL1_ -= this->padRightL1_;
103+ }
104+ this->orgWin_ = this->curWiLoadL1_;
105+ }
106+ 
107+ // Per-group cout: each core handles groupsPerCore groups.
108+ uint32_t groupCoutPar = this->tiling_->cout / this->tiling_->groups;
109+ uint32_t groupsPerCorePar = this->tiling_->groups / this->tiling_->groupDim;
110+ this->coutAligned_ = AlignB(groupsPerCorePar * groupCoutPar, GN0);
111+ 
112+ this->n1Total_ = this->coutAligned_ / GN0;
113+ 
114+ // Use tiling-side decided kAL1 for Cin L1 partitioning.
115+ // cinL1_ derived from kAL1: kAL1 = cinL1 * kernelHxkernelW.
116+ // hoL0_/woL0_ are already set by InitCommon from tiling_->hoL0/woL0.
117+ this->cinL1_ = AlignB(this->tiling_->kAL1 / this->tiling_->kernelHxkernelW, this->GK0);
118+ this->cinL1Blocks_ = CeilDiv(this->cinAligned_, this->cinL1_);
119+ 
120+ this->al1BufBytes_ = this->tiling_->aL1SpaceSize;
121+ this->al1ElemPerBuf_ = this->tiling_->aL1SpaceSize / sizeof(FmapType);
122+ 
123+ // L1 layout: [FM pingpong 2 bufs][Weight fullload][Bias][Scale0][ReluWeight0][Scale1][ReluWeight1]
124+ // Reuse SetupL1SplitLayout from parent (identical layout computation).
125+ this->SetupL1SplitLayout();
126+}
127+ 
128+template <typename FmapType, typename weightType, typename biasType, typename out0Type, typename out1Type,
129+ bool isNHWCin, bool isNHWCout, bool IsHwMode>
130+__aicore__ inline void Conv2dSmallKernelFmPartload<FmapType, weightType, biasType, out0Type, out1Type, isNHWCin,
131+ isNHWCout, IsHwMode>::LoadWeightL1Full(GM_ADDR filter)
132+{
133+ // Weight one-shot fullload into L1 (after FM pingpong buffers).
134+ // N axis fullload L0: singleCoreCo (per-core N partition) is fully loaded; no further N split.
135+ GlobalTensor<weightType> filterGm;
136+ filterGm.SetGlobalBuffer(reinterpret_cast<__gm__ weightType*>(filter),
137+ this->k1Total_ * this->n1Total_ * GN0 * this->GK0);
138+ 
139+ LocalTensor<weightType> bl1(TPosition::B1, this->bl1OffBytes_, this->bl1ElemCount_);
140+ 
141+ if (this->tiling_->nDim == 1) {
142+ // Whole weight fits in this core's N partition: direct copy.
143+ DataCopy(bl1, filterGm[0], this->bl1ElemCount_);
144+ } else {
145+ // N-axis partitioning across cores, but per-core N is fullload L0 (one-shot).
146+ uint32_t n1Start = this->nIdx_ * this->tiling_->singleCoreCo / GN0;
147+ uint32_t tileBytes = GN0 * this->GK0 * sizeof(weightType);
148+ uint32_t srcGmOff = n1Start * GN0 * this->GK0;
149+ uint16_t blkLen = static_cast<uint16_t>((this->n1PerCore_ * tileBytes) / FMP_BLOCK_BYTES);
150+ uint16_t srcGap = static_cast<uint16_t>(((this->n1Total_ - this->n1PerCore_) * tileBytes) / FMP_BLOCK_BYTES);
151+ DataCopyParams cp(static_cast<uint16_t>(this->k1Total_), blkLen, srcGap, 0);
152+ DataCopy(bl1, filterGm[srcGmOff], cp);
153+ }
154+}
155+ 
156+template <typename FmapType, typename weightType, typename biasType, typename out0Type, typename out1Type,
157+ bool isNHWCin, bool isNHWCout, bool IsHwMode>
158+__aicore__ inline void Conv2dSmallKernelFmPartload<FmapType, weightType, biasType, out0Type, out1Type, isNHWCin,
159+ isNHWCout, IsHwMode>::Process(GM_ADDR x, GM_ADDR filter,
160+ GM_ADDR bias, GM_ADDR y,
161+ const ExtendParams* extendParams)
162+{
163+ if (!this->coreActive_ || this->actualCo_ == 0) {
164+ return;
165+ }
166+ 
167+ // kL0 from tiling directly (N axis fullload L0: nl0 = nbl1, no per-chunk recomputation).
168+ uint32_t kL0 = this->tiling_->kL0;
169+ uint32_t kL0Iters = CeilDiv(this->kTotal_, kL0);
170+ uint32_t kernelHxW = this->tiling_->kh * this->tiling_->kw;
171+ 
172+ LocalTensor<weightType> bl1Full(TPosition::B1, this->bl1OffBytes_, this->bl1ElemCount_);
173+ uint32_t mmadN = AlignB(this->actualCo_, GN0);
174+ uint64_t hwOut = static_cast<uint64_t>(this->tiling_->hout) * this->tiling_->wout;
175+ 
176+ // Stage 1: Load Bias/Scale/ReluWeight into L1 (small, channel-wise).
177+ this->LoadBiasScaleL1(bias, extendParams);
178+ SetFlag<HardEvent::MTE2_MTE1>(FMP_EVT_WBS_DONE);
179+ WaitFlag<HardEvent::MTE2_MTE1>(FMP_EVT_WBS_DONE);
180+ 
181+ SetFlag<HardEvent::MTE2_FIX>(static_cast<event_t>(0));
182+ WaitFlag<HardEvent::MTE2_FIX>(static_cast<event_t>(0));
183+ 
184+ // Stage 2: Setup FM GM buffer (singleCoreBatch == 1 constraint).
185+ if constexpr (isNHWCin) {
186+ uint64_t batchFmapOff = static_cast<uint64_t>(this->batchIdx_) * this->tiling_->hin * this->tiling_->win *
187+ this->tiling_->cin;
188+ this->fmapGm_.SetGlobalBuffer(reinterpret_cast<__gm__ FmapType*>(x) + batchFmapOff);
189+ } else {
190+ uint64_t batchFmapOff = static_cast<uint64_t>(this->batchIdx_) * this->tiling_->cin * this->tiling_->hin *
191+ this->tiling_->win;
192+ this->fmapGm_.SetGlobalBuffer(reinterpret_cast<__gm__ FmapType*>(x) + batchFmapOff);
193+ }
194+ 
195+ // Stage 3: Weight fullload L1 (one-shot, no per-K splitting).
196+ // N axis fullload L0: per-core singleCoreCo loaded at once.
197+ LoadWeightL1Full(filter);
198+ SetFlag<HardEvent::MTE2_MTE1>(FMP_EVT_WBS_DONE);
199+ WaitFlag<HardEvent::MTE2_MTE1>(FMP_EVT_WBS_DONE);
200+ 
201+ if (this->tiling_->hasBias) {
202+ this->LoadBiasToBT();
203+ }
204+ 
205+ // Stage 4: Dispatch to HW-mode or M-mode compute loop.
206+ // FM is chunked over Cin (shared helper via ProcessCinBlocks); Weight is already
207+ // fully in L1, so loadWeight = false inside both modes (no per-chunk weight reload).
208+ if constexpr (IsHwMode) {
209+ ProcessHwMode(kL0, kL0Iters, kernelHxW, mmadN, hwOut, y, extendParams, bl1Full);
210+ } else {
211+ ProcessMMode(kL0, kL0Iters, kernelHxW, mmadN, hwOut, y, extendParams, bl1Full);
212+ }
213+}
214+ 
215+template <typename FmapType, typename weightType, typename biasType, typename out0Type, typename out1Type,
216+ bool isNHWCin, bool isNHWCout, bool IsHwMode>
217+__aicore__ inline void Conv2dSmallKernelFmPartload<FmapType, weightType, biasType, out0Type, out1Type, isNHWCin,
218+ isNHWCout, IsHwMode>::ProcessHwMode(uint32_t kL0, uint32_t kL0Iters,
219+ uint32_t kernelHxW,
220+ uint32_t mmadN, uint64_t hwOut,
221+ GM_ADDR y,
222+ const ExtendParams* extendParams,
223+ LocalTensor<weightType>& bl1Full)
224+{
225+ // HW-mode: nested Ho/Wo-chunk loop; each chunk accumulates over all cin blocks.
226+ // Weight already fully loaded into L1 (one-shot), so loadWeight = false.
227+ bool needRowSplit = (this->actualWo_ < static_cast<uint32_t>(this->tiling_->wout));
228+ for (uint32_t hoOff = 0; hoOff < this->actualHo_; hoOff += this->hoL0_) {
229+ uint32_t curHo = this->hoL0_;
230+ if (hoOff + curHo > this->actualHo_) {
231+ curHo = this->actualHo_ - hoOff;
232+ }
233+ for (uint32_t woOff = 0; woOff < this->actualWo_; woOff += this->woL0_) {
234+ uint32_t curWo = this->woL0_;
235+ if (woOff + curWo > this->actualWo_) {
236+ curWo = this->actualWo_ - woOff;
237+ }
238+ uint32_t curM = curHo * curWo;
239+ uint32_t curMAlign = AlignB(curM, GM0);
240+ 
241+ uint32_t curHi, padTop, padBottom, hiLoadOff;
242+ this->CalcChunkFmap(hoOff, curHo, curHi, padTop, padBottom, hiLoadOff);
243+ 
244+ uint32_t curWi;
245+ int32_t padLeft, padRight;
246+ uint32_t wiLoadOff;
247+ this->CalcChunkFmapW(woOff, curWo, curWi, padLeft, padRight, wiLoadOff);
248+ 
249+ LocalTensor<L0cT> cl0(TPosition::CO1, 0, this->L0C_ELEMS);
250+ MmadParams mp;
251+#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)
252+ if constexpr (AscendC::IsSameType<FmapType, half>::value) {
253+ mp.fixShiftVal = this->tiling_->fixedShiftValue;
254+ }
255+#endif
256+ mp.m = curMAlign;
257+ mp.n = mmadN;
258+ mp.cmatrixInitVal = !(this->tiling_->hasBias);
259+ mp.cmatrixSource = (this->tiling_->hasBias != 0);
260+ 
261+ this->ProcessCinBlocks(cl0, mp, bl1Full, kL0, kL0Iters, kernelHxW, curHi, padTop, padBottom, hiLoadOff,
262+ curWi, wiLoadOff, curM, hoOff, woOff, padLeft, padRight, false);
263+ 
264+ uint32_t outOff = (this->hoIdxStart_ + hoOff) * static_cast<uint32_t>(this->tiling_->wout) +
265+ this->woIdxStart_ + woOff;
266+ uint32_t fpMSize = needRowSplit ? curWo : curM;
267+ uint32_t fpDnNum = needRowSplit ? curHo : 1;
268+ uint32_t fpDstDnStride = needRowSplit ? static_cast<uint32_t>(this->tiling_->wout) :
269+ static_cast<uint32_t>(hwOut);
270+ this->CopyOutResult(cl0, y, extendParams, outOff, fpMSize, curMAlign, fpDnNum, fpDstDnStride);
271+ }
272+ }
273+}
274+ 
275+template <typename FmapType, typename weightType, typename biasType, typename out0Type, typename out1Type,
276+ bool isNHWCin, bool isNHWCout, bool IsHwMode>
277+__aicore__ inline void Conv2dSmallKernelFmPartload<FmapType, weightType, biasType, out0Type, out1Type, isNHWCin,
278+ isNHWCout, IsHwMode>::ProcessMMode(uint32_t kL0, uint32_t kL0Iters,
279+ uint32_t kernelHxW,
280+ uint32_t mmadN, uint64_t hwOut,
281+ GM_ADDR y,
282+ const ExtendParams* extendParams,
283+ LocalTensor<weightType>& bl1Full)
284+{
285+ // M-mode: M-loop -> CinL1 chunk loop -> KL0 inner loop -> Fixpipe.
286+ // Weight already fully loaded into L1 (one-shot), so loadWeight = false.
287+ // N axis fullload L0: actualCo_ used in full (no per-group reload).
288+ for (uint32_t mOff = 0; mOff < this->actualM_; mOff += this->hoL0_) {
289+ uint32_t curM = this->hoL0_;
290+ if (mOff + curM > this->actualM_) {
291+ curM = this->actualM_ - mOff;
292+ }
293+ 
294+ uint32_t curHi, padTop, padBottom, hiLoadOff;
295+ this->CalcChunkFmap(mOff, curM, curHi, padTop, padBottom, hiLoadOff);
296+ 
297+ uint32_t curMAlign = AlignB(curM, GM0);
298+ LocalTensor<L0cT> cl0(TPosition::CO1, 0, this->L0C_ELEMS);
299+ 
300+ MmadParams mp;
301+#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)
302+ if constexpr (AscendC::IsSameType<FmapType, half>::value) {
303+ mp.fixShiftVal = this->tiling_->fixedShiftValue;
304+ }
305+#endif
306+ mp.m = curMAlign;
307+ mp.n = mmadN; // N axis fullload L0: actualCo_ used in full.
308+ mp.cmatrixInitVal = !(this->tiling_->hasBias);
309+ mp.cmatrixSource = (this->tiling_->hasBias != 0);
310+ 
311+ // ProcessCinBlocks handles the CinL1 chunk loop + KL0 inner loop internally.
312+ // loadWeight = false: weight already fully loaded, skip per-chunk LoadWeightL1Block.
313+ printf("fm partload ProcessCinBlocks mOff=%d\n", mOff);
314+ this->ProcessCinBlocks(cl0, mp, bl1Full, kL0, kL0Iters, kernelHxW, curHi, padTop, padBottom, hiLoadOff,
315+ this->orgWin_, 0, curM, mOff, 0, 0, 0, false);
316+ 
317+ // Fixpipe out (supports NHWC and NCHW output formats).
318+ this->CopyOutResult(cl0, y, extendParams, this->mIdxStart_ + mOff, curM, curMAlign, 1,
319+ static_cast<uint32_t>(hwOut));
320+ }
321+}
322+ 
323+template <typename FmapType, typename weightType, typename biasType, typename out0Type, typename out1Type,
324+ bool isNHWCin, bool isNHWCout, bool IsHwMode>
325+__aicore__ inline void
326+Conv2dSmallKernelFmPartload<FmapType, weightType, biasType, out0Type, out1Type, isNHWCin, isNHWCout,
327+ IsHwMode>::ProcessCinBlocks(LocalTensor<L0cT>& cl0, MmadParams& mp,
328+ LocalTensor<weightType>& bl1Full, uint32_t kL0,
329+ uint32_t kL0Iters, uint32_t kernelHxW, uint32_t curHi,
330+ uint32_t padTop, uint32_t padBottom, uint32_t hiLoadOff,
331+ uint32_t curWi, uint32_t wiLoadOff, uint32_t curM,
332+ uint32_t setupMOff, uint32_t setupWoOff, int32_t padLeft,
333+ int32_t padRight, bool loadWeight)
334+{
335+ SetFlag<HardEvent::MTE1_MTE2>(EVT_FMAP_BUF0);
336+ SetFlag<HardEvent::MTE1_MTE2>(EVT_FMAP_BUF1);
337+ for (uint32_t kl1 = 0; kl1 < this->cinL1Blocks_; kl1++) {
338+ uint32_t cinOff;
339+ uint32_t curCin;
340+ uint32_t curCinOri;
341+ uint32_t kOff;
342+ uint32_t curKL1;
343+ uint32_t kl1Buf;
344+ event_t kl1Ev;
345+ this->PrepareCinBlock(kl1, kernelHxW, cinOff, curCin, curCinOri, kOff, curKL1, kl1Buf, kl1Ev);
346+ 
347+ WaitFlag<HardEvent::MTE1_MTE2>(kl1Ev);
348+ 
349+ this->LoadFmapL1Chunk(kl1Buf, curHi, hiLoadOff, padTop, padBottom, curWi, wiLoadOff, cinOff, curCinOri);
350+ if (loadWeight) {
351+ this->LoadWeightL1Block(kOff, curKL1);
352+ }
353+ 
354+ SetFlag<HardEvent::MTE2_MTE1>(kl1Ev);
355+ WaitFlag<HardEvent::MTE2_MTE1>(kl1Ev);
356+ 
357+ this->SetupLoad3DForChunk(curHi, setupMOff, curM, padTop, padBottom, setupWoOff, padLeft, padRight, curWi);
358+ 
359+ uint32_t al1ElemCount = curHi * curWi * curCin;
360+ uint32_t al1BufOff = kl1Buf * this->al1BufBytes_;
361+ LocalTensor<FmapType> al1(TPosition::A1, al1BufOff, al1ElemCount);
362+ 
363+ this->RunKL0Loop(al1, bl1Full, cl0, mp, kOff, curKL1, kl1, kL0, kL0Iters);
364+ 
365+ SetFlag<HardEvent::MTE1_MTE2>(kl1Ev);
366+ }
367+ 
368+ WaitFlag<HardEvent::MTE1_MTE2>(EVT_FMAP_BUF0);
369+ WaitFlag<HardEvent::MTE1_MTE2>(EVT_FMAP_BUF1);
370+}
371+#endif // CONV2D_SMALL_KERNEL_FM_PARTLOAD_H
Mconv/conv2d_v2/op_kernel/arch35/conv2d_small_kernel_parallelism.h+2-2
@@ -35,7 +35,7 @@ public:
35 __aicore__ inline void Process(GM_ADDR x, GM_ADDR filter, GM_ADDR bias, GM_ADDR y,35 __aicore__ inline void Process(GM_ADDR x, GM_ADDR filter, GM_ADDR bias, GM_ADDR y,
36 const ExtendParams* extendParams);36 const ExtendParams* extendParams);
37 37 
38-private:38+protected:
39 using Base = Conv2dSmallKernel<FmapType, weightType, biasType, out0Type, out1Type, isNHWCin, isNHWCout,39 using Base = Conv2dSmallKernel<FmapType, weightType, biasType, out0Type, out1Type, isNHWCin, isNHWCout,
40 ConvFormat::FRACTAL_Z, IsHwMode>;40 ConvFormat::FRACTAL_Z, IsHwMode>;
41 using Output0T = typename Base::Output0T;41 using Output0T = typename Base::Output0T;
@@ -228,7 +228,7 @@ __aicore__ inline void Conv2dSmallKernelParallelism<FmapType, weightType, biasTy
228 hiLoadOff = hiStart - this->tiling_->padTop;228 hiLoadOff = hiStart - this->tiling_->padTop;
229 }229 }
230 230 
231- uint32_t needHi = hoRelEnd * this->tiling_->strideH + this->tiling_->dilationH * (this->tiling_->kh - 1) + padTop +231+ uint32_t needHi = hoRelEnd * this->tiling_->strideH + this->tiling_->dilationH * (this->tiling_->kh - 1) - padTop +
232 1;232 1;
233 uint32_t maxGmRows = static_cast<uint32_t>(this->tiling_->hin) - hiLoadOff;233 uint32_t maxGmRows = static_cast<uint32_t>(this->tiling_->hin) - hiLoadOff;
234 curHi = (needHi < maxGmRows) ? needHi : maxGmRows;234 curHi = (needHi < maxGmRows) ? needHi : maxGmRows;
Mconv/conv2d_v2/op_kernel/arch35/conv2d_v2_tilingkey.h+2-2
@@ -56,8 +56,8 @@ ASCENDC_TPL_ARGS_DECL(
56 ASCENDC_TPL_UINT_DECL(NoPad, ASCENDC_TPL_1_BW, ASCENDC_TPL_UI_LIST, CONV_HAS_PAD, CONV_NO_PAD),56 ASCENDC_TPL_UINT_DECL(NoPad, ASCENDC_TPL_1_BW, ASCENDC_TPL_UI_LIST, CONV_HAS_PAD, CONV_NO_PAD),
57 ASCENDC_TPL_UINT_DECL(SmallWeight, ASCENDC_TPL_2_BW, ASCENDC_TPL_UI_LIST, CONV_NOT_SMALL_WEIGHT,57 ASCENDC_TPL_UINT_DECL(SmallWeight, ASCENDC_TPL_2_BW, ASCENDC_TPL_UI_LIST, CONV_NOT_SMALL_WEIGHT,
58 CONV_FULLLOAD_KL1_NL0, CONV_WEIGHT_SMALLER_THAN_BL0),58 CONV_FULLLOAD_KL1_NL0, CONV_WEIGHT_SMALLER_THAN_BL0),
59- ASCENDC_TPL_UINT_DECL(SmallKernel, ASCENDC_TPL_1_BW, ASCENDC_TPL_UI_LIST, CONV_NOT_SMALL_KERNEL,59+ ASCENDC_TPL_UINT_DECL(SmallKernel, ASCENDC_TPL_2_BW, ASCENDC_TPL_UI_LIST, CONV_NOT_SMALL_KERNEL, CONV_SMALL_KERNEL,
60- CONV_SMALL_KERNEL));60+ CONV_SMALL_KERNEL_FM_PARTLOAD));
61 61 
62#if defined(FORMAT_FILTER) && (FORMAT_FILTER == FORMAT_FRACTAL_Z || FORMAT_FILTER == FORMAT_FRACTAL_Z_C04)62#if defined(FORMAT_FILTER) && (FORMAT_FILTER == FORMAT_FRACTAL_Z || FORMAT_FILTER == FORMAT_FRACTAL_Z_C04)
63#define CONV2D_SMALL_KERNEL_NCHW_SEL()63#define CONV2D_SMALL_KERNEL_NCHW_SEL()
Mconv/conv2d_v2/op_kernel/conv2d_v2.cpp+9-1
@@ -20,6 +20,7 @@
20#include "arch35/conv2d_v2_tilingkey.h"20#include "arch35/conv2d_v2_tilingkey.h"
21#include "arch35/conv2d_small_kernel.h"21#include "arch35/conv2d_small_kernel.h"
22#include "arch35/conv2d_small_kernel_parallelism.h"22#include "arch35/conv2d_small_kernel_parallelism.h"
23+#include "arch35/conv2d_small_kernel_fm_partload.h"
23#include "arch35/conv2d_v2_depthwise_simplify.h"24#include "arch35/conv2d_v2_depthwise_simplify.h"
24 25 
25using namespace AscendC;26using namespace AscendC;
@@ -112,7 +113,14 @@ __global__ __aicore__ void conv2dv2(GM_ADDR x, GM_ADDR filter, GM_ADDR bias, GM_
112#endif113#endif
113 using scaleType = ConvType<TPosition::GM, scaleFormat, uint64_t>;114 using scaleType = ConvType<TPosition::GM, scaleFormat, uint64_t>;
114 115 
115- if constexpr (SmallKernel == 1) {116+ if constexpr (SmallKernel == CONV_SMALL_KERNEL_FM_PARTLOAD) {
117+ constexpr bool isNHWCin = (fmapFormat == ConvFormat::NHWC);
118+ constexpr bool isNHWCout = (outputFormat == ConvFormat::NHWC);
119+ constexpr bool isHw = (OutputOrder == static_cast<int8_t>(ConvOutputOrder::HW_MODE));
120+ Conv2dSmallKernelFmPartload<DTYPE_X, DTYPE_FILTER, biasType::T, DTYPE_Y, half, isNHWCin, isNHWCout, isHw> op;
121+ op.Init(tilingData);
122+ op.Process(x, filter, bias, y, nullptr);
123+ } else if constexpr (SmallKernel == CONV_SMALL_KERNEL) {
116 constexpr bool isNHWCin = (fmapFormat == ConvFormat::NHWC);124 constexpr bool isNHWCin = (fmapFormat == ConvFormat::NHWC);
117 constexpr bool isNHWCout = (outputFormat == ConvFormat::NHWC);125 constexpr bool isNHWCout = (outputFormat == ConvFormat::NHWC);
118 constexpr bool isHw = (OutputOrder == static_cast<int8_t>(ConvOutputOrder::HW_MODE));126 constexpr bool isHw = (OutputOrder == static_cast<int8_t>(ConvOutputOrder::HW_MODE));
Mconv/extend_conv2d/op_kernel/extend_conv2d.cpp+12-1
@@ -20,6 +20,7 @@
20#include "../conv2d_v2/arch35/conv2d_v2_tilingkey.h"20#include "../conv2d_v2/arch35/conv2d_v2_tilingkey.h"
21#include "../conv2d_v2/arch35/conv2d_small_kernel.h"21#include "../conv2d_v2/arch35/conv2d_small_kernel.h"
22#include "../conv2d_v2/arch35/conv2d_small_kernel_parallelism.h"22#include "../conv2d_v2/arch35/conv2d_small_kernel_parallelism.h"
23+#include "../conv2d_v2/arch35/conv2d_small_kernel_fm_partload.h"
23 24 
24using namespace AscendC;25using namespace AscendC;
25 26 
@@ -119,7 +120,17 @@ __global__ __aicore__ void extend_conv2d(GM_ADDR x, GM_ADDR filter, GM_ADDR bias
119#endif120#endif
120 121 
121 ExtendParams extendParams(scale0, relu_weight0, clip_value0, scale1, relu_weight1, clip_value1, y1);122 ExtendParams extendParams(scale0, relu_weight0, clip_value0, scale1, relu_weight1, clip_value1, y1);
122- if constexpr (SmallKernel == 1 && weightFormat == ConvFormat::FRACTAL_Z) {123+ 
124+ if constexpr (SmallKernel == CONV_SMALL_KERNEL_FM_PARTLOAD) {
125+ constexpr bool isNHWCin = (fmapFormat == ConvFormat::NHWC);
126+ constexpr bool isNHWCout = (outputFormat == ConvFormat::NHWC);
127+ constexpr bool isHw = (OutputOrder == static_cast<int8_t>(ConvOutputOrder::HW_MODE));
128+ Conv2dSmallKernelFmPartload<DTYPE_X, DTYPE_FILTER, biasType::T, DTYPE_Y0, output1Type, isNHWCin, isNHWCout,
129+ isHw>
130+ op;
131+ op.Init(tilingData);
132+ op.Process(x, filter, bias, y0, &extendParams);
133+ } else if constexpr (SmallKernel == CONV_SMALL_KERNEL && weightFormat == ConvFormat::FRACTAL_Z) {
123 constexpr bool isNHWCin = (fmapFormat == ConvFormat::NHWC);134 constexpr bool isNHWCin = (fmapFormat == ConvFormat::NHWC);
124 constexpr bool isNHWCout = (outputFormat == ConvFormat::NHWC);135 constexpr bool isNHWCout = (outputFormat == ConvFormat::NHWC);
125 constexpr bool isHw = (OutputOrder == static_cast<int8_t>(ConvOutputOrder::HW_MODE));136 constexpr bool isHw = (OutputOrder == static_cast<int8_t>(ConvOutputOrder::HW_MODE));