Star927
1847
代码
代码
Issues234
Pull Requests356
Actions
项目讨论
Wiki
分析
Star927
1847
已合并
matmul v3 增加vector nz2nd模板 #5356
zhengyuhao3创建于 5月28日
matmul v3 增加vector nz2nd模板 #5356
已合并
zhengyuhao3创建于 5月28日
从zhengyuhao3/ops-nnpr_4809合入到cann/ops-nnmaster
共 10 个文件变更+826-419
Mmatmul/common/op_host/op_api/matmul_util.cpp+1-4
@@ -119,10 +119,7 @@ static inline bool CheckMMV3NzNzNdSupport(MmOpInfo& mmOpInfo)
119 ALIGN_UNIT_MAP.find(mmOpInfo.support_info.self_dtype) == ALIGN_UNIT_MAP.end()) {119 ALIGN_UNIT_MAP.find(mmOpInfo.support_info.self_dtype) == ALIGN_UNIT_MAP.end()) {
120 return false;120 return false;
121 }121 }
122- auto it = ALIGN_UNIT_MAP.find(mmOpInfo.support_info.self_dtype);122+ return op::GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_2201;
123- uint64_t alignUnit = it->second;
124- return (op::GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_2201) &&
125- (mmOpInfo.shapeInfo.nDim % alignUnit == 0);
126}123}
127 124 
128static bool CheckDtypeValid(125static bool CheckDtypeValid(
Mmatmul/mat_mul_v3/op_host/op_tiling/matmul_v3_base_tiling.cpp+31-2
@@ -631,6 +631,22 @@ void MatmulV3BaseTiling::CalL1Tiling()
631 runInfo_.singleCoreN = runInfo_.baseN;631 runInfo_.singleCoreN = runInfo_.baseN;
632 return;632 return;
633}633}
634+bool MatmulV3BaseTiling::CheckNzNzNdSupport() {
635+ bool isDAV2201 = (compileInfo_.npuArch == NpuArch::DAV_2201);
636+ if (isDAV2201 &&
637+ (args_.aFormat == ge::FORMAT_FRACTAL_NZ) &&
638+ (args_.bFormat == ge::FORMAT_FRACTAL_NZ) &&
639+ (args_.outFormat == ge::FORMAT_ND) &&
640+ (args_.nValue % BASIC_ALIGN_256 != 0)) {
641+ return true;
642+ }
643+ return false;
644+}
645+void MatmulV3BaseTiling::DoNzNzNdVectorTiling() {
646+ if (CheckNzNzNdSupport()) {
647+ tilingEnable_.tilingEnableFixOpti = TilingEnableFixOpti::VEC_NZ2ND_UNALIGNOUT;
648+ }
649+}
634 650 
635void MatmulV3BaseTiling::DoBasicTiling()651void MatmulV3BaseTiling::DoBasicTiling()
636{652{
@@ -673,6 +689,7 @@ void MatmulV3BaseTiling::DoBasicTiling()
673 DoSelectTiling();689 DoSelectTiling();
674 // add nd2nz tiling here690 // add nd2nz tiling here
675 DoNd2NzVectorTiling();691 DoNd2NzVectorTiling();
692+ DoNzNzNdVectorTiling();
676 if (args_.hasBias) {693 if (args_.hasBias) {
677 runInfo_.baseN = std::min(256UL, runInfo_.baseN); // 有bias时 baseN 小于256694 runInfo_.baseN = std::min(256UL, runInfo_.baseN); // 有bias时 baseN 小于256
678 if (tilingEnable_.tilingEnableSplitCore == TilingEnableSplitCore::BASE &&695 if (tilingEnable_.tilingEnableSplitCore == TilingEnableSplitCore::BASE &&
@@ -2761,6 +2778,18 @@ uint64_t MatmulV3BaseTiling::GetDeterministicSplitKWorkspaceSize(uint64_t aligne
2761 RPC_WORKSIZE * MB_SIZE;2778 RPC_WORKSIZE * MB_SIZE;
2762}2779}
2763 2780 
2781+void MatmulV3BaseTiling::GetVecNz2ndUnAlignedOutWorkspaceSize()
2782+{
2783+ if(CheckNzNzNdSupport())
2784+ {
2785+ workspaceSize_ += tilingData_.matmulTiling.baseN * tilingData_.matmulTiling.baseM *
2786+ tilingData_.matmulTiling.usedCoreNum * NUMBER_TWO * cDtypeSize_;
2787+ }else{
2788+ workspaceSize_ += ops::CeilAlign(args_.nValue, N_ALIGNED) * tilingData_.matmulTiling.baseM *
2789+ tilingData_.matmulTiling.usedCoreNum * NUMBER_TWO * cDtypeSize_;
2790+ }
2791+}
2792+ 
2764ge::graphStatus MatmulV3BaseTiling::GetWorkspaceSize()2793ge::graphStatus MatmulV3BaseTiling::GetWorkspaceSize()
2765{2794{
2766 uint64_t align256Byte = 256 / aDtypeSize_; // 256B 对齐shape2795 uint64_t align256Byte = 256 / aDtypeSize_; // 256B 对齐shape
@@ -2785,9 +2814,9 @@ ge::graphStatus MatmulV3BaseTiling::GetWorkspaceSize()
2785 tilingData_.matmulTiling.usedCoreNum * NUMBER_TWO * cDtypeSize_;2814 tilingData_.matmulTiling.usedCoreNum * NUMBER_TWO * cDtypeSize_;
2786 }2815 }
2787 if (tilingEnable_.tilingEnableFixOpti == TilingEnableFixOpti::VEC_NZ2ND_UNALIGNOUT) {2816 if (tilingEnable_.tilingEnableFixOpti == TilingEnableFixOpti::VEC_NZ2ND_UNALIGNOUT) {
2788- workspaceSize_ += ops::CeilAlign(args_.nValue, N_ALIGNED) * tilingData_.matmulTiling.baseM *2817+ GetVecNz2ndUnAlignedOutWorkspaceSize();
2789- tilingData_.matmulTiling.usedCoreNum * NUMBER_TWO * cDtypeSize_;
2790 }2818 }
2819+
2791 if (!compileInfo_.supportL0c2out) {2820 if (!compileInfo_.supportL0c2out) {
2792 return ge::GRAPH_SUCCESS;2821 return ge::GRAPH_SUCCESS;
2793 }2822 }
Mmatmul/mat_mul_v3/op_host/op_tiling/matmul_v3_base_tiling.h+3-0
@@ -76,6 +76,7 @@ protected:
76 ge::graphStatus SelectNZTiling();76 ge::graphStatus SelectNZTiling();
77 void DoTilingKey();77 void DoTilingKey();
78 uint64_t GetDeterministicSplitKWorkspaceSize(uint64_t alignedM, uint64_t alignedN);78 uint64_t GetDeterministicSplitKWorkspaceSize(uint64_t alignedM, uint64_t alignedN);
79+ void GetVecNz2ndUnAlignedOutWorkspaceSize();
79 void DoBasicTiling();80 void DoBasicTiling();
80 void FormulaicBaseBlockTiling();81 void FormulaicBaseBlockTiling();
81 void FormulaicTilingNoTrans();82 void FormulaicTilingNoTrans();
@@ -142,6 +143,8 @@ protected:
142 const std::vector<uint64_t> supportNd2nzList) const;143 const std::vector<uint64_t> supportNd2nzList) const;
143 bool NeedNd2NzVnchw(uint64_t outerSize, uint64_t innerSize, bool supportNd2NzOnTheWay,144 bool NeedNd2NzVnchw(uint64_t outerSize, uint64_t innerSize, bool supportNd2NzOnTheWay,
144 uint64_t dtypeSize, ge::Format matFormat) const;145 uint64_t dtypeSize, ge::Format matFormat) const;
146+ bool CheckNzNzNdSupport();
147+ void DoNzNzNdVectorTiling();
145private:148private:
146 MatmulTilingData tilingDataSelf_{};149 MatmulTilingData tilingDataSelf_{};
147protected:150protected:
Amatmul/mat_mul_v3/op_kernel/mat_mul_base_fixpipe_kernel.h+200-0
@@ -0,0 +1,200 @@
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 mat_mul_base_fixpipe_kernel.h
13+ * \brief Intermediate base class for fixpipe-based matmul kernels with AIC+AIV cooperation.
14+ * Extracts common code from MatmulBaseVectorNz2NdKernel and MatmulBaseUnalignedNKernel.
15+ */
16+ #ifndef __OP_KERNEL_MATMUL_V3_BASE_FIXPIPE_KERNEL_H__
17+ #define __OP_KERNEL_MATMUL_V3_BASE_FIXPIPE_KERNEL_H__
18+
19+ #include "mat_mul_base_kernel.h"
20+ #include "mat_mul_nz2nd.h"
21+
22+ using namespace AscendC;
23+ using namespace matmul;
24+
25+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE = MatmulBaseBlock,
26+ const MatmulConfig &MM_CFG = MM_CFG_NO_PRELOAD, class MM_CB = MatmulCallBackFunc<nullptr, nullptr, nullptr>>
27+ class MatmulBaseFixPipeKernel : public MatmulBaseKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB> {
28+ public:
29+ using C_T = typename C_TYPE::T;
30+ __aicore__ inline MatmulBaseFixPipeKernel() {}
31+
32+ __aicore__ inline void AicProcess(GlobalTensor<C_T>& cTensor, uint8_t enAtomic, bool aicNeedWaitAiv,
33+ uint8_t pingPongId);
34+ __aicore__ inline void UpdateOffsetC(bool isNd);
35+
36+ // AicProcess 的两段公共实现，拆分点即派生类可注入差异（如切换 orgM）的位置：
37+ // - AicLoadAndIterate: 设置 single shape、A/B/bias，执行 Iterate，并按需等待 AIV；
38+ // - AicStoreC: 将结果写出（GetTensorC）并发起 AIC->AIV 同步。
39+ __aicore__ inline void AicLoadAndIterate(bool aicNeedWaitAiv, uint8_t pingPongId);
40+ __aicore__ inline void AicStoreC(GlobalTensor<C_T>& cTensor, uint8_t enAtomic, uint8_t pingPongId);
41+
42+ // 计算 AIV 侧 NZ2ND 转换所需的公共块参数，返回 false 表示本核无需处理。
43+ __aicore__ inline bool PrepareNz2NdConversion(uint8_t pingPongId, LocalTensor<C_T>& tensorNZ,
44+ LocalTensor<C_T>& tensorND, uint64_t& ubProcessMNum, uint64_t& srcGmOffset, uint64_t& dstGmOffset);
45+
46+ // AIV 侧 NZ2ND 的公共拷入与转换流程。
47+ __aicore__ inline void Nz2NdCopyAndConvert(GlobalTensor<C_T> cNzGlobal, LocalTensor<C_T> tensorNZ,
48+ LocalTensor<C_T> tensorND, uint64_t ubProcessMNum, uint64_t srcGmOffset, uint8_t pingPongId,
49+ uint64_t gatherNBlocks);
50+
51+ // 公共的两级 tiling 主调度循环（AIC+AIV 混合）。DERIVED 为派生 kernel 类型，
52+ // 循环内通过 static_cast<DERIVED*> 直接调用派生类的 AicProcess 与 AivStep，
53+ template <class DERIVED>
54+ __aicore__ inline void RunFixpipeTileLoop(uint64_t index, uint8_t enAtomic);
55+
56+ protected:
57+ GlobalTensor<C_T> tempCGlobal_;
58+ GatherMaskParams params_;
59+ uint64_t baseSize_ = 0UL;
60+ uint64_t alignedN_ = 0UL;
61+ uint64_t c0Size_ = BLOCK_SIZE;
62+ };
63+
64+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig &MM_CFG,
65+ class MM_CB>
66+ __aicore__ inline void MatmulBaseFixPipeKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB>::AicLoadAndIterate(
67+ bool aicNeedWaitAiv, uint8_t pingPongId)
68+ {
69+ this->mm_.SetSingleShape(this->block_.params_.singleCoreM, this->block_.params_.singleCoreN,
70+ this->block_.matmulTilingData_->matmulTiling.singleCoreK);
71+ this->mm_.SetTensorA(this->aGlobal_[this->block_.offset_.offsetA], A_TYPE::isTrans);
72+ this->mm_.SetTensorB(this->bGlobal_[this->block_.offset_.offsetB], B_TYPE::isTrans);
73+ if (this->block_.matmulTilingData_->matmulTiling.isBias) {
74+ this->mm_.SetBias(this->biasGlobal_[this->block_.offset_.offsetBias]);
75+ }
76+ this->mm_.Iterate();
77+ if (aicNeedWaitAiv) {
78+ CrossCoreWaitFlag(AIC_SYNC_AIV_FLAG + pingPongId);
79+ }
80+ }
81+
82+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig &MM_CFG,
83+ class MM_CB>
84+ __aicore__ inline void MatmulBaseFixPipeKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB>::AicStoreC(
85+ GlobalTensor<C_T>& cTensor, uint8_t enAtomic, uint8_t pingPongId)
86+ {
87+ this->mm_.GetTensorC(cTensor, enAtomic);
88+ #if defined(__CCE_AICORE__) && __CCE_AICORE__ == 220
89+ CrossCoreSetFlag<0x2, PIPE_FIX>(AIV_SYNC_AIC_FLAG + pingPongId);
90+ #endif
91+ }
92+
93+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig &MM_CFG,
94+ class MM_CB>
95+ __aicore__ inline void MatmulBaseFixPipeKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB>::AicProcess(
96+ GlobalTensor<typename C_TYPE::T>& cTensor, uint8_t enAtomic, bool aicNeedWaitAiv, uint8_t pingPongId)
97+ {
98+ if ASCEND_IS_AIC {
99+ this->AicLoadAndIterate(aicNeedWaitAiv, pingPongId);
100+ this->AicStoreC(cTensor, enAtomic, pingPongId);
101+ }
102+ }
103+
104+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig &MM_CFG,
105+ class MM_CB>
106+ __aicore__ inline void MatmulBaseFixPipeKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB>::UpdateOffsetC(
107+ bool isNd)
108+ {
109+ Nz2NdUpdateOffsetC(this->block_.params_, this->block_.offset_, this->block_.matmulTilingData_, isNd);
110+ }
111+
112+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig &MM_CFG,
113+ class MM_CB>
114+ __aicore__ inline bool MatmulBaseFixPipeKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG,
115+ MM_CB>::PrepareNz2NdConversion(uint8_t pingPongId, LocalTensor<C_T>& tensorNZ, LocalTensor<C_T>& tensorND,
116+ uint64_t& ubProcessMNum, uint64_t& srcGmOffset, uint64_t& dstGmOffset)
117+ {
118+ uint32_t vBlockIndex = GetBlockIdx();
119+ if (vBlockIndex >= (this->block_.matmulTilingData_->matmulTiling.usedCoreNum * NUM_TWO)) {
120+ return false;
121+ }
122+ CrossCoreWaitFlag(AIV_SYNC_AIC_FLAG + pingPongId);
123+ LocalTensor<C_T> ubTensor = this->ubBuf_.template Get<C_T>();
124+ uint64_t cDtypeSize = sizeof(C_T);
125+ ubProcessMNum = min(MMV3CeilAlign(this->block_.params_.singleCoreM / NUM_TWO, this->c0Size_),
126+ static_cast<uint64_t>(this->block_.params_.singleCoreM));
127+ int64_t subIdx = GetSubBlockIdx();
128+ srcGmOffset = 0UL;
129+ dstGmOffset = 0UL;
130+ if (subIdx == 1) {
131+ srcGmOffset = ubProcessMNum * ALIGNED_H;
132+ dstGmOffset = ubProcessMNum * this->block_.matmulTilingData_->matmulTiling.N;
133+ ubProcessMNum = this->block_.params_.singleCoreM - ubProcessMNum;
134+ }
135+ uint64_t ndOffset = (TOTAL_UB_SIZE >> 2) / cDtypeSize;
136+ uint64_t pingpongOffset = (TOTAL_UB_SIZE >> 1) / cDtypeSize;
137+ tensorNZ = ubTensor[pingPongId * pingpongOffset];
138+ tensorND = ubTensor[pingPongId * pingpongOffset + ndOffset];
139+ return true;
140+ }
141+
142+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig &MM_CFG,
143+ class MM_CB>
144+ __aicore__ inline void MatmulBaseFixPipeKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG,
145+ MM_CB>::Nz2NdCopyAndConvert(GlobalTensor<C_T> cNzGlobal, LocalTensor<C_T> tensorNZ, LocalTensor<C_T> tensorND,
146+ uint64_t ubProcessMNum, uint64_t srcGmOffset, uint8_t pingPongId, uint64_t gatherNBlocks)
147+ {
148+ Nz2NdCopyInWithNz<C_T>(tensorNZ, cNzGlobal[srcGmOffset], ubProcessMNum, pingPongId,
149+ this->alignedN_, this->block_.matmulTilingData_->matmulTiling.baseM);
150+ CrossCoreSetFlag<0x2, PIPE_MTE2>(AIC_SYNC_AIV_FLAG + pingPongId);
151+ Nz2NdMulsAndGatherMask<C_T>(tensorND, tensorNZ, ubProcessMNum, pingPongId,
152+ this->alignedN_, this->c0Size_, gatherNBlocks, this->params_);
153+ this->UpdateOffsetC(true);
154+ }
155+
156+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig &MM_CFG,
157+ class MM_CB>
158+ template <class DERIVED>
159+ __aicore__ inline void MatmulBaseFixPipeKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG,
160+ MM_CB>::RunFixpipeTileLoop(uint64_t index, uint8_t enAtomic)
161+ {
162+ int8_t pingPongId = 0;
163+ bool aicNeedWaitAiv = false;
164+ bool reverse = true;
165+ GlobalTensor<C_T> tempCGlobal = this->tempCGlobal_;
166+ for (uint64_t mTileIndex = 0; mTileIndex < this->block_.params_.mTileCntL2; mTileIndex++) {
167+ reverse = !reverse;
168+ for (uint64_t nTileIndexTemp = 0; nTileIndexTemp < this->block_.params_.nTileCntL2; nTileIndexTemp++) {
169+ uint64_t nTileIndex = reverse ? (this->block_.params_.nTileCntL2 - nTileIndexTemp - 1) : nTileIndexTemp;
170+ this->block_.UpdateBlockCnt(mTileIndex, nTileIndex);
171+ this->block_.InitBlockIndex(index);
172+ SetFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG));
173+ SetFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + 1));
174+ for (uint64_t j = 0; j < this->block_.params_.realRound; j++) {
175+ tempCGlobal = this->tempCGlobal_[this->baseSize_ * (GetCurrentBlockIdx() * 2 + pingPongId)];
176+ if (this->block_.params_.rowOrder == 0) {
177+ this->block_.UpdateBasicIndex(j);
178+ }
179+ if (this->block_.params_.index < this->block_.params_.totalTileCnt) {
180+ this->block_.UpdateBlockParams(mTileIndex, nTileIndex);
181+ this->block_.template CalcGMOffset<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE>(mTileIndex, nTileIndex);
182+ static_cast<DERIVED*>(this)->AicProcess(tempCGlobal, enAtomic, aicNeedWaitAiv, pingPongId);
183+ static_cast<DERIVED*>(this)->AivStep(tempCGlobal, pingPongId);
184+ aicNeedWaitAiv = aicNeedWaitAiv || bool(pingPongId);
185+ pingPongId = (pingPongId + 1) & 1;
186+ }
187+ this->block_.UpdateBlockIndex();
188+ }
189+ WaitFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG));
190+ WaitFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + 1));
191+ }
192+ }
193+ if (this->block_.params_.isHf32) {
194+ this->mm_.SetHF32(false, 0);
195+ }
196+ PipeBarrier<PIPE_ALL>();
197+ }
198+
199+ #endif
200+
Amatmul/mat_mul_v3/op_kernel/mat_mul_base_kernel_vec_nz2nd.h+173-0
@@ -0,0 +1,173 @@
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 mat_mul_base_vector_nz2nd_kernel.h
13+ * \brief Matmul kernel with AIV-based NZ2ND conversion for half/bfloat16 output
14+ */
15+ #ifndef __OP_KERNEL_MATMUL_V3_BASE_VECTOR_NZ2ND_KERNEL_H__
16+ #define __OP_KERNEL_MATMUL_V3_BASE_VECTOR_NZ2ND_KERNEL_H__
17+
18+ #include "mat_mul_base_fixpipe_kernel.h"
19+
20+ using namespace AscendC;
21+ using namespace matmul;
22+
23+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE = MatmulBaseBlock,
24+ const MatmulConfig &MM_CFG = MM_CFG_NO_PRELOAD, class MM_CB = MatmulCallBackFunc<nullptr, nullptr, nullptr>>
25+ class MatmulBaseVectorNz2NdKernel
26+ : public MatmulBaseFixPipeKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB> {
27+ public:
28+ using C_T = typename C_TYPE::T;
29+ __aicore__ inline MatmulBaseVectorNz2NdKernel() {}
30+
31+ __aicore__ inline void Init(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR biasGM, GM_ADDR offsetWGM,
32+ GM_ADDR workspaceGM, const void *tilingData, TPipe *pipe);
33+
34+ __aicore__ inline void AivNz2NdProcess(GlobalTensor<C_T> cNzGlobal, uint8_t pingPongId);
35+
36+ // 读 A 与写 C 的 orgM 取值不同，需在 Iterate 与 GetTensorC 之间切换 orgM。
37+ __aicore__ inline void AicProcess(GlobalTensor<C_T>& cTensor, uint8_t enAtomic, bool aicNeedWaitAiv,
38+ uint8_t pingPongId);
39+
40+ // 供基类 RunFixpipeTileLoop 回调的每块 AIV 处理步骤。
41+ __aicore__ inline void AivStep(GlobalTensor<C_T>& cTensor, uint8_t pingPongId)
42+ {
43+ AivNz2NdProcess(cTensor, pingPongId);
44+ }
45+
46+ __aicore__ inline void UpdateGlobalTensor(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR biasGM, GM_ADDR offsetWGM,
47+ GM_ADDR workspaceGM);
48+
49+ __aicore__ inline void Process(uint64_t index = 0, uint8_t enAtomic = 0);
50+ };
51+
52+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig &MM_CFG,
53+ class MM_CB>
54+ __aicore__ inline void MatmulBaseVectorNz2NdKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB>::Init(
55+ GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR biasGM, GM_ADDR offsetWGM,
56+ GM_ADDR workspaceGM, const void *tilingData, TPipe *pipe)
57+ {
58+ GetSizeC0<C_T>(this->c0Size_);
59+ this->block_.template Init<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE>(tilingData);
60+ this->pipe_ = pipe;
61+ this->InitInputs(aGM, bGM, cGM, biasGM);
62+
63+ this->pipe_->InitBuffer(this->ubBuf_, TOTAL_UB_SIZE);
64+
65+ int64_t originShapeM = this->block_.matmulTilingData_->matmulTiling.baseM;
66+
67+ this->baseSize_ = this->block_.matmulTilingData_->matmulTiling.baseN *
68+ this->block_.matmulTilingData_->matmulTiling.baseM;
69+
70+ this->tempCGlobal_.SetGlobalBuffer(reinterpret_cast<__gm__ C_T*>(workspaceGM),
71+ this->baseSize_ * NUM_TWO * this->block_.matmulTilingData_->matmulTiling.usedCoreNum);
72+ this->mm_.SetSubBlockIdx(0);
73+ this->mm_.Init(&this->block_.matmulTilingData_->matmulTiling, pipe);
74+ this->mm_.SetUserDefInfo(reinterpret_cast<uint64_t>(tilingData));
75+
76+ this->mm_.SetOrgShape(this->block_.params_.alignedOriM, this->block_.params_.alignedOriN,
77+ this->block_.params_.alignedKaSize, this->block_.params_.alignedKbSize,
78+ AlignUp(this->block_.matmulTilingData_->matmulTiling.N, BLOCK_SIZE));
79+
80+ if (this->block_.params_.isHf32) {
81+ this->mm_.SetHF32(true, 1);
82+ } else {
83+ this->mm_.SetHF32(false, 0);
84+ }
85+ }
86+
87+
88+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig &MM_CFG,
89+ class MM_CB>
90+ __aicore__ inline void MatmulBaseVectorNz2NdKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB>::AicProcess(
91+ GlobalTensor<C_T>& cTensor, uint8_t enAtomic, bool aicNeedWaitAiv, uint8_t pingPongId)
92+ {
93+ if ASCEND_IS_AIC {
94+ // 读 A 阶段：A 为整块 FRACTAL_NZ，沿 K 方向跨步使用 orgM，需用 alignedOriM
95+ this->mm_.SetOrgShape(this->block_.params_.alignedOriM, this->block_.params_.alignedOriN,
96+ this->block_.params_.alignedKaSize, this->block_.params_.alignedKbSize,
97+ AlignUp(this->block_.matmulTilingData_->matmulTiling.N, BLOCK_SIZE));
98+ this->AicLoadAndIterate(aicNeedWaitAiv, pingPongId);
99+ // 写 C 阶段：C 写入每个 base 块独立的 NZ workspace，N 分形跨步使用 orgM，需用 baseM
100+ this->mm_.SetOrgShape(this->block_.matmulTilingData_->matmulTiling.baseM, this->block_.params_.alignedOriN,
101+ this->block_.params_.alignedKaSize, this->block_.params_.alignedKbSize,
102+ AlignUp(this->block_.params_.singleCoreN, BLOCK_SIZE));
103+ this->AicStoreC(cTensor, enAtomic, pingPongId);
104+ }
105+ }
106+
107+
108+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig &MM_CFG,
109+ class MM_CB>
110+ __aicore__ inline void MatmulBaseVectorNz2NdKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB>::AivNz2NdProcess(
111+ GlobalTensor<C_T> cNzGlobal, uint8_t pingPongId)
112+ {
113+ if ASCEND_IS_AIV {
114+ uint64_t nBlocks = this->block_.params_.singleCoreN;
115+ this->alignedN_ = AlignUp(nBlocks, BLOCK_SIZE);
116+
117+ LocalTensor<C_T> tensorNZ;
118+ LocalTensor<C_T> tensorND;
119+ uint64_t ubProcessMNum = 0UL;
120+ uint64_t srcGmOffset = 0UL;
121+ uint64_t dstGmOffset = 0UL;
122+ if (!this->PrepareNz2NdConversion(pingPongId, tensorNZ, tensorND, ubProcessMNum, srcGmOffset, dstGmOffset)) {
123+ return;
124+ }
125+
126+ if (ubProcessMNum == 0UL) {
127+ //补充同步信号
128+ WaitFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + pingPongId));
129+ CrossCoreSetFlag<0x2, PIPE_MTE2>(AIC_SYNC_AIV_FLAG + pingPongId);
130+ SetFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + pingPongId));
131+ return;
132+ }
133+
134+ WaitFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + pingPongId));
135+ if (std::is_same_v<C_T, bfloat16_t> || std::is_same_v<C_T, half>) {
136+ this->Nz2NdCopyAndConvert(cNzGlobal, tensorNZ, tensorND, ubProcessMNum, srcGmOffset, pingPongId,
137+ this->alignedN_);
138+
139+ DataCopyExtParams dataCopyExtParams;
140+ dataCopyExtParams.blockCount = ubProcessMNum;
141+ dataCopyExtParams.blockLen = static_cast<uint32_t>(nBlocks * sizeof(C_T));
142+ dataCopyExtParams.srcStride = 0;
143+ dataCopyExtParams.dstStride = static_cast<uint32_t>((this->block_.matmulTilingData_->matmulTiling.N - nBlocks) * sizeof(C_T));
144+ dataCopyExtParams.rsv = 0;
145+ DataCopyPad(this->cGlobal_[this->block_.offset_.offsetC + dstGmOffset], tensorND, dataCopyExtParams);
146+ } else {
147+ CrossCoreSetFlag<0x2, PIPE_MTE2>(AIC_SYNC_AIV_FLAG + pingPongId);
148+ }
149+ SetFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + pingPongId));
150+ }
151+ }
152+
153+
154+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig &MM_CFG,
155+ class MM_CB>
156+ __aicore__ inline void MatmulBaseVectorNz2NdKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG,
157+ MM_CB>::UpdateGlobalTensor(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR biasGM, GM_ADDR offsetWGM,
158+ GM_ADDR workspaceGM)
159+ {
160+ this->InitInputs(aGM, bGM, cGM, biasGM);
161+ }
162+
163+
164+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig &MM_CFG,
165+ class MM_CB>
166+ __aicore__ inline void MatmulBaseVectorNz2NdKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB>::Process(
167+ uint64_t index, uint8_t enAtomic)
168+ {
169+ this->template RunFixpipeTileLoop<
170+ MatmulBaseVectorNz2NdKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB>>(index, enAtomic);
171+ }
172+ #endif // __OP_KERNEL_MATMUL_V3_BASE_VECTOR_NZ2ND_KERNEL_H__
173+
Amatmul/mat_mul_v3/op_kernel/mat_mul_nz2nd.h+28-0
@@ -0,0 +1,28 @@
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 mat_mul_nz2nd.h
13+ * \brief NZ2ND process utilities for AIV-based NZ to ND format conversion.
14+ * Provides low-level utility functions that can be reused by different kernel templates.
15+ * Refer to mat_mul_nd2nz.h for the ND2NZ counterpart pattern.
16+ */
17+#ifndef __OP_KERNEL_MATMUL_V3_NZ2ND_H__
18+#define __OP_KERNEL_MATMUL_V3_NZ2ND_H__
19+ 
20+#include "mat_mul_nz2nd_util.h"
21+#include "mat_mul_base_block.h"
22+ 
23+using namespace AscendC;
24+using namespace matmul;
25+ 
26+const uint8_t AIV_DB_SYNC_FLAG = 0x2;
27+ 
28+#endif
Amatmul/mat_mul_v3/op_kernel/mat_mul_nz2nd_util.h+122-0
@@ -0,0 +1,122 @@
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 mat_mul_nz2nd_util.h
13+ * \brief NZ2ND utility functions for vector-based NZ to ND format conversion
14+ */
15+#ifndef __OP_KERNEL_MATMUL_V3_NZ2ND_UTIL_H__
16+#define __OP_KERNEL_MATMUL_V3_NZ2ND_UTIL_H__
17+ 
18+#include "mat_mul_v3_common.h"
19+#include "mat_mul_base_block.h"
20+ 
21+using namespace AscendC;
22+using namespace matmul;
23+ 
24+__aicore__ inline void Nz2NdUpdateOffsetC(BaseBlockArgs& params, BlockOffset& offset,
25+ const MatmulTilingData* tilingData, bool isNd)
26+{
27+ uint64_t mCntIndex = params.index / params.nCntUse;
28+ uint64_t nCntIndex = params.index % params.nCntUse;
29+ if (isNd) {
30+ offset.offsetC =
31+ (nCntIndex * tilingData->matmulTiling.singleCoreN +
32+ mCntIndex * tilingData->matmulTiling.singleCoreM *
33+ tilingData->matmulTiling.N +
34+ (params.mTileAddrOffset * tilingData->matmulTiling.N +
35+ params.nTileAddrOffset));
36+ } else {
37+ offset.offsetC =
38+ (nCntIndex * tilingData->matmulTiling.singleCoreN *
39+ tilingData->matmulTiling.M +
40+ mCntIndex * tilingData->matmulTiling.singleCoreM * BLOCK_SIZE +
41+ (params.mTileAddrOffset * BLOCK_SIZE +
42+ params.nTileAddrOffset * tilingData->matmulTiling.M));
43+ }
44+}
45+ 
46+template <class C_T>
47+__aicore__ inline void Nz2NdCopyInWithNz(LocalTensor<C_T> tensorNZ, GlobalTensor<C_T> cNzGlobal,
48+ uint64_t ubProcessM, uint8_t pingPongId, uint64_t alignedN, uint64_t baseM)
49+{
50+ size_t NfractualNum = alignedN / BLOCK_SIZE;
51+ DataCopyParams copyParams;
52+ copyParams.blockCount = NfractualNum;
53+ if constexpr (std::is_same_v<C_T, float>) {
54+ copyParams.blockLen = ubProcessM * NUM_TWO;
55+ copyParams.srcStride = (baseM - ubProcessM) * NUM_TWO;
56+ } else {
57+ copyParams.blockLen = ubProcessM;
58+ copyParams.srcStride = baseM - ubProcessM;
59+ }
60+ copyParams.dstStride = 0;
61+ DataCopy(tensorNZ, cNzGlobal, copyParams);
62+ SetFlag<HardEvent::MTE2_V>(static_cast<event_t>(pingPongId));
63+ WaitFlag<HardEvent::MTE2_V>(static_cast<event_t>(pingPongId));
64+}
65+ 
66+template <class C_T>
67+__aicore__ inline void Nz2NdMulsAndGatherMask(LocalTensor<C_T> tensorND, LocalTensor<C_T> tensorNZ,
68+ uint64_t ubProcessM, uint8_t pingPongId, uint64_t alignedN, uint64_t c0Size,
69+ uint64_t nBlocks, GatherMaskParams& params)
70+{
71+ size_t NfractalNum = alignedN / BLOCK_SIZE;
72+ uint8_t repeatTimes = MMV3DivCeil(ubProcessM, 8);
73+ uint64_t mask[2] = {UINT64_MAX, UINT64_MAX};
74+ UnaryRepeatParams mulsRepeatParams;
75+ 
76+ if constexpr (std::is_same_v<C_T, float>) {
77+ mulsRepeatParams.srcBlkStride = 2;
78+ mulsRepeatParams.dstBlkStride = alignedN / c0Size;
79+ mulsRepeatParams.dstRepStride = alignedN;
80+ mulsRepeatParams.srcRepStride = 16;
81+ for (size_t inLoop = 0; inLoop < NfractalNum; inLoop++) {
82+ Muls(tensorND[inLoop * BLOCK_SIZE], tensorNZ[inLoop * ubProcessM * BLOCK_SIZE], static_cast<C_T>(1.0), mask,
83+ repeatTimes, mulsRepeatParams);
84+ Muls(tensorND[inLoop * BLOCK_SIZE + c0Size], tensorNZ[inLoop * ubProcessM * BLOCK_SIZE + c0Size],
85+ static_cast<C_T>(1.0), mask, repeatTimes, mulsRepeatParams);
86+ }
87+ } else if constexpr (std::is_same_v<C_T, half>) {
88+ mulsRepeatParams.srcBlkStride = 1;
89+ mulsRepeatParams.dstBlkStride = alignedN / c0Size;
90+ mulsRepeatParams.dstRepStride = alignedN / NUM_TWO;
91+ mulsRepeatParams.srcRepStride = 8;
92+ for (size_t inLoop = 0; inLoop < NfractalNum; inLoop++) {
93+ Muls(tensorND[inLoop * BLOCK_SIZE], tensorNZ[inLoop * ubProcessM * BLOCK_SIZE], static_cast<half>(1.0), mask,
94+ repeatTimes, mulsRepeatParams);
95+ }
96+ } else if constexpr (std::is_same_v<C_T, bfloat16_t>) {
97+ mulsRepeatParams.srcBlkStride = 1;
98+ mulsRepeatParams.dstBlkStride = alignedN / c0Size;
99+ mulsRepeatParams.dstRepStride = alignedN / NUM_TWO;
100+ mulsRepeatParams.srcRepStride = 8;
101+ LocalTensor<half> tensorNDInterpretCast = tensorND.template ReinterpretCast<half>();
102+ LocalTensor<half> tensorNZInterpretCast = tensorNZ.template ReinterpretCast<half>();
103+ for (size_t inLoop = 0; inLoop < NfractalNum; inLoop++) {
104+ Muls(tensorNDInterpretCast[inLoop * BLOCK_SIZE], tensorNZInterpretCast[inLoop * ubProcessM * BLOCK_SIZE],
105+ static_cast<half>(1.0), mask, repeatTimes, mulsRepeatParams);
106+ }
107+ }
108+ 
109+ PipeBarrier<PIPE_V>();
110+ 
111+ uint64_t rsvdCnt = 0;
112+ params.src0BlockStride = 1;
113+ params.src0RepeatStride = alignedN / c0Size;
114+ params.src1RepeatStride = 0;
115+ params.repeatTimes = ubProcessM;
116+ GatherMask(tensorND, tensorND, 7, true, nBlocks, params, rsvdCnt);
117+ 
118+ SetFlag<HardEvent::V_MTE3>(static_cast<event_t>(pingPongId));
119+ WaitFlag<HardEvent::V_MTE3>(static_cast<event_t>(pingPongId));
120+}
121+ 
122+#endif
Mmatmul/mat_mul_v3/op_kernel/mat_mul_optimized_fixpipe_algorithm.h+250-412
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -10,415 +10,253 @@
10 10 
11/*!11/*!
12 * \file mat_mul_optimized_fixpipe_algorithm.h12 * \file mat_mul_optimized_fixpipe_algorithm.h
13- * \brief13+ * \brief Optimized fixpipe algorithm kernels with AIC+AIV cooperation.
14+ * MatmulBaseUnalignedNKernel inherits from MatmulBaseFixPipeKernel for shared AIC processing
15+ * and NZ2ND utilities. Supports both aligned output (AivProcess) and NZ2ND unaligned output
16+ * (AivNz2NdProcess) via FIXPIPE_OPT template parameter.
14 */17 */
15-#ifndef __OP_KERNEL_MATMUL_V3_OPTIMIZED_FIXPIPE_ALGORITHM_H__18+ #ifndef __OP_KERNEL_MATMUL_V3_OPTIMIZED_FIXPIPE_ALGORITHM_H__
16-#define __OP_KERNEL_MATMUL_V3_OPTIMIZED_FIXPIPE_ALGORITHM_H__19+ #define __OP_KERNEL_MATMUL_V3_OPTIMIZED_FIXPIPE_ALGORITHM_H__
17- 20+
18-#include "mat_mul_base_block.h"21+ #include "mat_mul_base_fixpipe_kernel.h"
19-#include "mat_mul_base_kernel.h"22+ #include "mat_mul_l1_full_load.h"
20-#include "mat_mul_l1_full_load.h"23+
21- 24+ using namespace AscendC;
22-using namespace AscendC;25+ using namespace matmul;
23-using namespace matmul;26+ using namespace MatmulV3;
24-using namespace MatmulV3;27+
25- 28+ #if defined(__CCE_KT_TEST__)
26-#if defined(__CCE_KT_TEST__)29+ using namespace std;
27-using namespace std;30+ #endif
28-#endif31+
29-// 512 byte32+ const uint32_t MM_ALIGN_SIZE = 512;
30-const uint32_t MM_ALIGN_SIZE = 512;33+
31-const uint8_t AIV_DB_SYNC_FLAG = 0x2;34+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE = MatmulBaseBlock,
32- 35+ const MatmulConfig& MM_CFG = MM_CFG_NO_PRELOAD, class MM_CB = MatmulCallBackFunc<nullptr, nullptr, nullptr>,
33-template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE = MatmulBaseBlock,36+ FIXPIPE_OPT_SELECT FIXPIPE_OPT = FIXPIPE_OPT_SELECT::BASE_ENABLE_ALIGNOUT>
34- const MatmulConfig& MM_CFG = MM_CFG_NO_PRELOAD, class MM_CB = MatmulCallBackFunc<nullptr, nullptr, nullptr>,37+ class MatmulBaseUnalignedNKernel : public MatmulBaseFixPipeKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE,
35- FIXPIPE_OPT_SELECT FIXPIPE_OPT = FIXPIPE_OPT_SELECT::BASE_ENABLE_ALIGNOUT>38+ BLOCK_TYPE, MM_CFG, MM_CB> {
36-class MatmulBaseUnalignedNKernel : public MatmulBaseKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE,39+ public:
37- BLOCK_TYPE, MM_CFG, MM_CB> {40+ using C_T = typename C_TYPE::T;
38-public:41+ __aicore__ inline MatmulBaseUnalignedNKernel() {}
39- using C_T = typename C_TYPE::T;42+ __aicore__ inline void Init(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR biasGM, GM_ADDR offsetWGM,
40- __aicore__ inline MatmulBaseUnalignedNKernel() {}43+ GM_ADDR workspaceGM, const void* tilingData, TPipe* pipe);
41- __aicore__ inline void Init(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR biasGM, GM_ADDR offsetWGM,44+ __aicore__ inline void AivProcess(GlobalTensor<C_T>& cTensor, uint8_t pingPongId);
42- GM_ADDR workspaceGM, const void* tilingData, TPipe* pipe);45+ __aicore__ inline void Process(uint64_t index = 0UL, uint8_t enAtomic = 0UL);
43- __aicore__ inline void AicProcess(GlobalTensor<C_T>& cTensor, uint8_t enAtomic, bool aicNeedWaitAiv,46+ __aicore__ inline void AivNz2NdProcess(GlobalTensor<C_T> cNzGlobal, uint8_t pingPongId);
44- uint8_t pingPongId);47+
45- __aicore__ inline void AivProcess(GlobalTensor<C_T>& cTensor, uint8_t pingPongId);48+ // 供基类 RunFixpipeTileLoop 回调的每块 AIV 处理步骤：按 FIXPIPE_OPT 选择 NZ2ND 或对齐输出路径。
46- __aicore__ inline void Process(uint64_t index = 0UL, uint8_t enAtomic = 0UL);49+ __aicore__ inline void AivStep(GlobalTensor<C_T>& cTensor, uint8_t pingPongId)
47- __aicore__ inline void AivNz2NdProcess(GlobalTensor<C_T> cNzGlobal, uint8_t pingPongId);50+ {
48- __aicore__ inline void CopyInWithNz(LocalTensor<C_T> tensorNZ, GlobalTensor<C_T> cNzGlobal, uint64_t ubProcessM,51+ if constexpr (FIXPIPE_OPT == FIXPIPE_OPT_SELECT::VEC_NZ2ND_UNALIGNOUT) {
49- uint8_t pingPongId);52+ AivNz2NdProcess(cTensor, pingPongId);
50- __aicore__ inline void Nz2NdAndGatherMask(LocalTensor<C_T> tensorND, LocalTensor<C_T> tensorNZ, uint64_t ubProcessM,53+ } else {
51- uint8_t pingPongId);54+ AivProcess(cTensor, pingPongId);
52- __aicore__ inline void UpdateOffsetC(bool isNd);55+ }
53- 56+ }
54-protected:57+ };
55- GlobalTensor<C_T> tempCGlobal_;58+
56- GatherMaskParams params_;59+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,
57- uint64_t baseSize_ = 0UL;60+ class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>
58- uint64_t alignedN_ = 0UL;61+ __aicore__ inline void
59- uint64_t c0Size_ = BLOCK_SIZE;62+ MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::Init(
60-};63+ GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR biasGM, GM_ADDR offsetWGM, GM_ADDR workspaceGM,
61- 64+ const void* tilingData, TPipe* pipe)
62-template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,65+ {
63- class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>66+ GetSizeC0<C_T>(this->c0Size_);
64-__aicore__ inline void67+ uint64_t cDtypeSize = sizeof(C_T);
65-MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::Init(68+ this->block_.template Init<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE>(tilingData);
66- GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR biasGM, GM_ADDR offsetWGM, GM_ADDR workspaceGM,69+ this->InitInputs(aGM, bGM, cGM, biasGM);
67- const void* tilingData, TPipe* pipe)70+
68-{71+ this->pipe_ = pipe;
69- GetSizeC0<C_T>(c0Size_);72+ this->pipe_->InitBuffer(this->ubBuf_, TOTAL_UB_SIZE);
70- uint64_t cDtypeSize = sizeof(C_T);73+
71- this->block_.template Init<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE>(tilingData);74+ int64_t originShapeM = 0;
72- this->InitInputs(aGM, bGM, cGM, biasGM);75+ if constexpr (FIXPIPE_OPT == FIXPIPE_OPT_SELECT::VEC_NZ2ND_UNALIGNOUT) {
73- 76+ this->alignedN_ = AlignUp(this->block_.matmulTilingData_->matmulTiling.N, BLOCK_SIZE);
74- this->pipe_ = pipe;77+ originShapeM = this->block_.matmulTilingData_->matmulTiling.baseM;
75- this->pipe_->InitBuffer(this->ubBuf_, TOTAL_UB_SIZE);78+ } else {
76- 79+ this->alignedN_ = AlignUp(this->block_.matmulTilingData_->matmulTiling.N, MM_ALIGN_SIZE / cDtypeSize);
77- int64_t originShapeM = 0;80+ originShapeM = this->block_.matmulTilingData_->matmulTiling.M;
78- if constexpr (FIXPIPE_OPT == FIXPIPE_OPT_SELECT::VEC_NZ2ND_UNALIGNOUT) {81+ }
79- alignedN_ = AlignUp(this->block_.matmulTilingData_->matmulTiling.N, BLOCK_SIZE);82+ this->baseSize_ = this->alignedN_ * this->block_.matmulTilingData_->matmulTiling.baseM;
80- // for MatMul highlevel API copyout NZ format consecutively83+ this->params_.src0BlockStride = 1;
81- originShapeM = this->block_.matmulTilingData_->matmulTiling.baseM;84+ this->params_.src0RepeatStride = this->alignedN_ / this->c0Size_;
82- } else {85+ this->params_.src1RepeatStride = 0;
83- alignedN_ = AlignUp(this->block_.matmulTilingData_->matmulTiling.N, MM_ALIGN_SIZE / cDtypeSize);86+
84- originShapeM = this->block_.matmulTilingData_->matmulTiling.M;87+ this->tempCGlobal_.SetGlobalBuffer(reinterpret_cast<__gm__ C_T*>(workspaceGM),
85- }88+ this->baseSize_ * NUM_TWO * this->block_.matmulTilingData_->matmulTiling.usedCoreNum);
86- baseSize_ = alignedN_ * this->block_.matmulTilingData_->matmulTiling.baseM;89+ this->mm_.SetSubBlockIdx(0);
87- params_.src0BlockStride = 1;90+ this->mm_.Init(&this->block_.matmulTilingData_->matmulTiling, pipe);
88- params_.src0RepeatStride = alignedN_ / c0Size_;91+ this->mm_.SetUserDefInfo(reinterpret_cast<uint64_t>(tilingData));
89- params_.src1RepeatStride = 0;92+ this->mm_.SetOrgShape(originShapeM, this->block_.params_.alignedOriN,
90- 93+ this->block_.matmulTilingData_->matmulTiling.singleCoreK, this->block_.params_.alignedKbSize,
91- tempCGlobal_.SetGlobalBuffer(reinterpret_cast<__gm__ C_T*>(workspaceGM),94+ this->alignedN_);
92- baseSize_ * NUM_TWO * this->block_.matmulTilingData_->matmulTiling.usedCoreNum);95+ if (this->block_.params_.isHf32) {
93- this->mm_.SetSubBlockIdx(0);96+ this->mm_.SetHF32(true, 1);
94- this->mm_.Init(&this->block_.matmulTilingData_->matmulTiling, pipe);97+ } else {
95- this->mm_.SetUserDefInfo(reinterpret_cast<uint64_t>(tilingData));98+ this->mm_.SetHF32(false, 0);
96- this->mm_.SetOrgShape(originShapeM, this->block_.params_.alignedOriN,99+ }
97- this->block_.matmulTilingData_->matmulTiling.singleCoreK, this->block_.params_.alignedKbSize,100+ }
98- this->alignedN_);101+
99- if (this->block_.params_.isHf32) {102+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,
100- this->mm_.SetHF32(true, 1);103+ class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>
101- } else {104+ __aicore__ inline void
102- this->mm_.SetHF32(false, 0);105+ MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::AivProcess(
103- }106+ GlobalTensor<C_T>& cTensor, uint8_t pingPongId)
104-}107+ {
105- 108+ if ASCEND_IS_AIV {
106-template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,109+ uint32_t vBlockIndex = GetBlockIdx();
107- class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>110+ if (vBlockIndex >= (this->block_.matmulTilingData_->matmulTiling.usedCoreNum * NUM_TWO)) {
108-__aicore__ inline void111+ return;
109-MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::AivProcess(112+ }
110- GlobalTensor<C_T>& cTensor, uint8_t pingPongId)113+ CrossCoreWaitFlag(AIV_SYNC_AIC_FLAG + pingPongId);
111-{114+
112- if ASCEND_IS_AIV {115+ uint64_t cDtypeSize = sizeof(C_T);
113- uint32_t vBlockIndex = GetBlockIdx();116+ LocalTensor<C_T> ubTensor = this->ubBuf_.template Get<C_T>();
114- if (vBlockIndex >= (this->block_.matmulTilingData_->matmulTiling.usedCoreNum * NUM_TWO)) {117+ uint64_t vecM = min(MMV3CeilAlign(this->block_.params_.singleCoreM / NUM_TWO, this->c0Size_),
115- return;118+ static_cast<uint64_t>(this->block_.params_.singleCoreM));
116- }119+ uint64_t subIdx = GetSubBlockIdx();
117- CrossCoreWaitFlag(0x4 + pingPongId);120+ uint64_t srcOffset = 0UL;
118- 121+ uint64_t dstOffset = 0UL;
119- uint64_t cDtypeSize = sizeof(C_T);122+ if (subIdx == 1) {
120- LocalTensor<C_T> ubTensor = this->ubBuf_.template Get<C_T>();123+ srcOffset = this->alignedN_ * vecM;
121- // aic : aiv is 1 : 2, singlecore cal half of baseM.124+ dstOffset = vecM * this->block_.matmulTilingData_->matmulTiling.N;
122- uint64_t vecM = min(MMV3CeilAlign(this->block_.params_.singleCoreM / NUM_TWO, c0Size_),125+ vecM = this->block_.params_.singleCoreM - vecM;
123- static_cast<uint64_t>(this->block_.params_.singleCoreM));126+ }
124- uint64_t subIdx = GetSubBlockIdx();127+ if (vecM == 0UL) {
125- uint64_t srcOffset = 0UL;128+ return;
126- uint64_t dstOffset = 0UL;129+ }
127- if (subIdx == 1) {130+ uint64_t ubOffset = (pingPongId * TOTAL_UB_SIZE >> 1) / cDtypeSize;
128- srcOffset = alignedN_ * vecM;131+ DataCopy<C_T>(ubTensor[ubOffset], cTensor[srcOffset], vecM * this->alignedN_);
129- dstOffset = vecM * this->block_.matmulTilingData_->matmulTiling.N;132+ CrossCoreSetFlag<0x2, PIPE_MTE2>(AIC_SYNC_AIV_FLAG + pingPongId);
130- vecM = this->block_.params_.singleCoreM - vecM;133+
131- }134+ SetFlag<HardEvent::MTE2_V>(static_cast<event_t>(pingPongId));
132- if (vecM == 0UL) {135+ WaitFlag<HardEvent::MTE2_V>(static_cast<event_t>(pingPongId));
133- return;136+ this->params_.repeatTimes = vecM;
134- }137+ uint64_t rsvdCnt = 0UL;
135- uint64_t ubOffset = (pingPongId * TOTAL_UB_SIZE >> 1) / cDtypeSize;138+ GatherMask(ubTensor[ubOffset], ubTensor[ubOffset], 7, true, this->block_.matmulTilingData_->matmulTiling.N,
136- DataCopy<C_T>(ubTensor[ubOffset], cTensor[srcOffset], vecM * alignedN_);139+ this->params_, rsvdCnt);
137- CrossCoreSetFlag<0x2, PIPE_MTE2>(0x6 + pingPongId);140+ SetFlag<HardEvent::V_MTE3>(static_cast<event_t>(pingPongId));
138- 141+ WaitFlag<HardEvent::V_MTE3>(static_cast<event_t>(pingPongId));
139- SetFlag<HardEvent::MTE2_V>(static_cast<event_t>(pingPongId));142+ DataCopy<C_T>(this->cGlobal_[this->block_.offset_.offsetC + dstOffset], ubTensor[ubOffset],
140- WaitFlag<HardEvent::MTE2_V>(static_cast<event_t>(pingPongId));143+ AlignUp(vecM * this->block_.matmulTilingData_->matmulTiling.N, this->c0Size_));
141- params_.repeatTimes = vecM;144+ SetFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(pingPongId));
142- uint64_t rsvdCnt = 0UL;145+ WaitFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(pingPongId));
143- // src1Pattern is 7; mask is this->block_.matmulTilingData_->matmulTiling.N;146+ }
144- GatherMask(ubTensor[ubOffset], ubTensor[ubOffset], 7, true, this->block_.matmulTilingData_->matmulTiling.N,147+ }
145- params_, rsvdCnt);148+
146- SetFlag<HardEvent::V_MTE3>(static_cast<event_t>(pingPongId));149+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,
147- WaitFlag<HardEvent::V_MTE3>(static_cast<event_t>(pingPongId));150+ class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>
148- DataCopy<C_T>(this->cGlobal_[this->block_.offset_.offsetC + dstOffset], ubTensor[ubOffset],151+ __aicore__ inline void
149- AlignUp(vecM * this->block_.matmulTilingData_->matmulTiling.N, c0Size_));152+ MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::AivNz2NdProcess(
150- SetFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(pingPongId));153+ GlobalTensor<C_T> cNzGlobal, uint8_t pingPongId)
151- WaitFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(pingPongId));154+ {
152- }155+ if ASCEND_IS_AIV {
153-}156+ LocalTensor<C_T> tensorNZ;
154- 157+ LocalTensor<C_T> tensorND;
155-template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,158+ uint64_t ubProcessMNum = 0UL;
156- class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>159+ uint64_t srcGmOffset = 0UL;
157-__aicore__ inline void160+ uint64_t dstGmOffset = 0UL;
158-MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::AicProcess(161+ if (!this->PrepareNz2NdConversion(pingPongId, tensorNZ, tensorND, ubProcessMNum, srcGmOffset, dstGmOffset)) {
159- GlobalTensor<typename C_TYPE::T>& cTensor, uint8_t enAtomic, bool aicNeedWaitAiv, uint8_t pingPongId)162+ return;
160-{163+ }
161- if ASCEND_IS_AIC {164+ if (ubProcessMNum == 0UL) {
162- this->mm_.SetSingleShape(this->block_.params_.singleCoreM, this->block_.params_.singleCoreN,165+ return;
163- this->block_.matmulTilingData_->matmulTiling.singleCoreK);166+ }
164- this->mm_.SetTensorA(this->aGlobal_[this->block_.offset_.offsetA], A_TYPE::isTrans);167+
165- this->mm_.SetTensorB(this->bGlobal_[this->block_.offset_.offsetB], B_TYPE::isTrans);168+ WaitFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + pingPongId));
166- if (this->block_.matmulTilingData_->matmulTiling.isBias) {169+ this->Nz2NdCopyAndConvert(cNzGlobal, tensorNZ, tensorND, ubProcessMNum, srcGmOffset, pingPongId,
167- this->mm_.SetBias(this->biasGlobal_[this->block_.offset_.offsetBias]);170+ this->block_.matmulTilingData_->matmulTiling.N);
168- }171+
169- this->mm_.Iterate();172+ DataCopy(this->cGlobal_[this->block_.offset_.offsetC + dstGmOffset], tensorND,
170- if (aicNeedWaitAiv) {173+ AlignUp(ubProcessMNum * this->block_.matmulTilingData_->matmulTiling.N, this->c0Size_));
171- CrossCoreWaitFlag(0x6 + pingPongId);174+ SetFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + pingPongId));
172- }175+ }
173- this->mm_.GetTensorC(cTensor, enAtomic);176+ }
174-#if defined(__CCE_AICORE__) && __CCE_AICORE__ == 220177+
175- CrossCoreSetFlag<0x2, PIPE_FIX>(0x4 + pingPongId);178+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,
176-#endif179+ class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>
177- }180+ __aicore__ inline void
178-}181+ MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::Process(
179- 182+ uint64_t index, uint8_t enAtomic)
180-template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,183+ {
181- class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>184+ ctx.isFirst = true;
182-__aicore__ inline void185+ ctx.inputDtypeSize = sizeof(typename A_TYPE::T);
183-MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::CopyInWithNz(186+ this->template RunFixpipeTileLoop<
184- LocalTensor<C_T> tensorNZ, GlobalTensor<C_T> cNzGlobal, uint64_t ubProcessM, uint8_t pingPongId)187+ MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>>(
185-{188+ index, enAtomic);
186- // AIV cNzGlobal start address has updated189+ }
187- size_t NfractualNum = alignedN_ / BLOCK_SIZE;190+
188- DataCopyParams copyParams;191+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE = MatmulBaseBlock,
189- copyParams.blockCount = NfractualNum;192+ const MatmulConfig& MM_CFG = MM_CFG_NO_PRELOAD, class MM_CB = MatmulCallBackFunc<nullptr, nullptr, nullptr>,
190- copyParams.blockLen = ubProcessM * NUM_TWO;193+ FIXPIPE_OPT_SELECT FIXPIPE_OPT = FIXPIPE_OPT_SELECT::BASE_ENABLE_ALIGNOUT>
191- copyParams.srcStride = (this->block_.matmulTilingData_->matmulTiling.baseM - ubProcessM) * NUM_TWO;194+ class MatmulBaseAToNZWithBL1FixpipeKernel
192- copyParams.dstStride = 0;195+ : public MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT> {
193- DataCopy(tensorNZ, cNzGlobal, copyParams);196+ struct BaseUnAlignedNKernelParams {
194- SetFlag<HardEvent::MTE2_V>(static_cast<event_t>(pingPongId));197+ GM_ADDR aGMNZ;
195- WaitFlag<HardEvent::MTE2_V>(static_cast<event_t>(pingPongId));198+ GM_ADDR workspaceGMNZ;
196-}199+ uint64_t baseAN;
197- 200+ uint64_t baseAD;
198-template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,201+ };
199- class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>202+
200-__aicore__ inline void203+ public:
201-MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::UpdateOffsetC(204+ __aicore__ inline MatmulBaseAToNZWithBL1FixpipeKernel() {}
202- bool isNd)205+
203-{206+ __aicore__ inline void Init(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR biasGM, GM_ADDR offsetWGM,
204- uint64_t mCntIndex = this->block_.params_.index / this->block_.params_.nCntUse;207+ GM_ADDR workspaceGM, const MatmulTilingData* tilingData, TPipe* pipe);
205- uint64_t nCntIndex = this->block_.params_.index % this->block_.params_.nCntUse;208+
206- if (isNd) {209+ __aicore__ inline void Process(uint64_t index = 0, uint8_t enAtomic = 0);
207- this->block_.offset_.offsetC =210+
208- (nCntIndex * this->block_.matmulTilingData_->matmulTiling.singleCoreN +211+ protected:
209- mCntIndex * this->block_.matmulTilingData_->matmulTiling.singleCoreM *212+ using C_T = typename C_TYPE::T;
210- this->block_.matmulTilingData_->matmulTiling.N +213+ BaseUnAlignedNKernelParams fixpipeInnerParams_;
211- (this->block_.params_.mTileAddrOffset * this->block_.matmulTilingData_->matmulTiling.N +214+ };
212- this->block_.params_.nTileAddrOffset));215+
213- } else {216+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,
214- this->block_.offset_.offsetC =217+ class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>
215- (nCntIndex * this->block_.matmulTilingData_->matmulTiling.singleCoreN *218+ __aicore__ inline void
216- this->block_.matmulTilingData_->matmulTiling.M +219+ MatmulBaseAToNZWithBL1FixpipeKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::Init(
217- mCntIndex * this->block_.matmulTilingData_->matmulTiling.singleCoreM * BLOCK_SIZE +220+ GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR biasGM, GM_ADDR offsetWGM, GM_ADDR workspaceGM,
218- (this->block_.params_.mTileAddrOffset * BLOCK_SIZE +221+ const MatmulTilingData* matmulTilingData, TPipe* pipe)
219- this->block_.params_.nTileAddrOffset * this->block_.matmulTilingData_->matmulTiling.M));222+ {
220- }223+ GetSizeC0<C_T>(this->c0Size_);
221-}224+ uint64_t cDtypeSize = sizeof(C_T);
222- 225+ uint64_t alignedN = AlignUp(matmulTilingData->matmulTiling.N, MM_ALIGN_SIZE / cDtypeSize);
223-template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,226+ uint64_t baseSize = alignedN * matmulTilingData->matmulTiling.baseM;
224- class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>227+
225-__aicore__ inline void228+ MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::Init(
226-MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB,229+ workspaceGM + baseSize * NUM_TWO * matmulTilingData->matmulTiling.usedCoreNum * cDtypeSize, bGM, cGM, biasGM,
227- FIXPIPE_OPT>::Nz2NdAndGatherMask(LocalTensor<C_T> tensorND, LocalTensor<C_T> tensorNZ,230+ offsetWGM, workspaceGM, matmulTilingData, pipe);
228- uint64_t ubProcessM, uint8_t pingPongId)231+ this->fixpipeInnerParams_.baseAN = this->block_.matmulTilingData_->baseAN;
229-{232+ this->fixpipeInnerParams_.baseAD = this->block_.matmulTilingData_->baseAD;
230- size_t NfractalNum = alignedN_ / BLOCK_SIZE;233+ this->fixpipeInnerParams_.aGMNZ = aGM;
231- uint8_t repeatTimes = MMV3DivCeil(ubProcessM, 8); // ub calc 8 row every repeat234+ this->fixpipeInnerParams_.workspaceGMNZ =
232- uint64_t mask[2] = {UINT64_MAX, UINT64_MAX};235+ workspaceGM + this->baseSize_ * NUM_TWO * this->block_.matmulTilingData_->matmulTiling.usedCoreNum * cDtypeSize;
233- UnaryRepeatParams mulsRepeatParams;236+ this->mm_.SetOrgShape(this->block_.params_.alignedOriM, this->block_.matmulTilingData_->matmulTiling.N,
234- mulsRepeatParams.srcBlkStride = 2;237+ this->block_.params_.alignedKaSize, this->block_.matmulTilingData_->matmulTiling.Kb,
235- mulsRepeatParams.dstBlkStride = alignedN_ / c0Size_;238+ this->alignedN_);
236- mulsRepeatParams.dstRepStride = alignedN_;239+ }
237- mulsRepeatParams.srcRepStride = 16;240+
238- // float calc 2 cols per loop241+ template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,
239- if constexpr (std::is_same_v<typename A_TYPE::T, float>) {242+ class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>
240- for (size_t inLoop = 0; inLoop < NfractalNum; inLoop++) {243+ __aicore__ inline void
241- Muls(tensorND[inLoop * BLOCK_SIZE], tensorNZ[inLoop * ubProcessM * BLOCK_SIZE], static_cast<C_T>(1.0), mask,244+ MatmulBaseAToNZWithBL1FixpipeKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::Process(
242- repeatTimes, mulsRepeatParams);245+ uint64_t index, uint8_t enAtomic)
243- Muls(tensorND[inLoop * BLOCK_SIZE + c0Size_], tensorNZ[inLoop * ubProcessM * BLOCK_SIZE + c0Size_],246+ {
244- static_cast<C_T>(1.0), mask, repeatTimes, mulsRepeatParams);247+ if ASCEND_IS_AIV {
245- }248+ MatrixAtoNZV2<typename A_TYPE::T>(this->fixpipeInnerParams_.workspaceGMNZ, this->fixpipeInnerParams_.aGMNZ,
246- }249+ this->block_.matmulTilingData_->matmulTiling, A_TYPE::isTrans, this->ubBuf_,
247- PipeBarrier<PIPE_V>();250+ this->fixpipeInnerParams_.baseAN, this->fixpipeInnerParams_.baseAD);
248- // GatherMask251+ SyncAll();
249- uint64_t rsvdCnt = 0;252+ CrossCoreSetFlag<0x2, PIPE_MTE3>(CV_FLAG);
250- params_.repeatTimes = ubProcessM;253+ }
251- GatherMask(tensorND, tensorND, 7, true, this->block_.matmulTilingData_->matmulTiling.N, params_, rsvdCnt);254+ if ASCEND_IS_AIC {
252- SetFlag<HardEvent::V_MTE3>(static_cast<event_t>(pingPongId));255+ CrossCoreWaitFlag(CV_FLAG);
253- WaitFlag<HardEvent::V_MTE3>(static_cast<event_t>(pingPongId));256+ }
254-}257+ MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::Process(
255- 258+ index, enAtomic);
256-template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,259+ }
257- class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>260+
258-__aicore__ inline void261+ #endif
259-MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::AivNz2NdProcess(262+
260- GlobalTensor<C_T> cNzGlobal, uint8_t pingPongId)
261-{
262- if ASCEND_IS_AIV {
263- uint32_t vBlockIndex = GetBlockIdx();
264- if (vBlockIndex >= (this->block_.matmulTilingData_->matmulTiling.usedCoreNum * NUM_TWO)) {
265- return;
266- }
267- CrossCoreWaitFlag(0x4 + pingPongId);
268- LocalTensor<C_T> ubTensor = this->ubBuf_.template Get<C_T>();
269- uint64_t cDtypeSize = sizeof(C_T);
270- uint64_t ubProcessMNum = min(MMV3CeilAlign(this->block_.params_.singleCoreM / NUM_TWO, c0Size_),
271- static_cast<uint64_t>(this->block_.params_.singleCoreM));
272- int64_t subIdx = GetSubBlockIdx();
273- uint64_t srcGmOffset = 0UL;
274- uint64_t dstGmOffset = 0UL;
275- 
276- if (subIdx == 1) {
277- srcGmOffset = ubProcessMNum * ALIGNED_H; // for aiv 1 offset
278- dstGmOffset = ubProcessMNum * this->block_.matmulTilingData_->matmulTiling.N;
279- ubProcessMNum = this->block_.params_.singleCoreM - ubProcessMNum;
280- }
281- if (ubProcessMNum == 0UL) {
282- return;
283- }
284- uint64_t ndOffset = (TOTAL_UB_SIZE >> 2) / cDtypeSize;
285- uint64_t pingpongOffset = (TOTAL_UB_SIZE >> 1) / cDtypeSize;
286- 
287- size_t NfractalNum = alignedN_ / ALIGNED_H;
288- 
289- LocalTensor<C_T> tensorNZ = ubTensor[pingPongId * pingpongOffset];
290- LocalTensor<C_T> tensorND = ubTensor[pingPongId * pingpongOffset + ndOffset];
291- WaitFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + pingPongId));
292- CopyInWithNz(tensorNZ, cNzGlobal[srcGmOffset], ubProcessMNum, pingPongId);
293- CrossCoreSetFlag<0x2, PIPE_MTE2>(0x6 + pingPongId); // 通知AIC MTE2可进行搬运
294- Nz2NdAndGatherMask(tensorND, tensorNZ, ubProcessMNum, pingPongId);
295- this->UpdateOffsetC(true);
296- 
297- DataCopy(this->cGlobal_[this->block_.offset_.offsetC + dstGmOffset], tensorND,
298- AlignUp(ubProcessMNum * this->block_.matmulTilingData_->matmulTiling.N, c0Size_));
299- SetFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + pingPongId));
300- }
301-}
302- 
303-template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,
304- class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>
305-__aicore__ inline void
306-MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::Process(
307- uint64_t index, uint8_t enAtomic)
308-{
309- bool reverse = true;
310- int8_t pingPongId = 0;
311- bool aicNeedWaitAiv = false;
312- ctx.isFirst = true;
313- ctx.inputDtypeSize = sizeof(typename A_TYPE::T);
314- GlobalTensor<C_T> tempCGlobal = tempCGlobal_;
315- for (uint64_t mTileIndex = 0; mTileIndex < this->block_.params_.mTileCntL2; mTileIndex++) {
316- reverse = !reverse;
317- for (uint64_t nTileIndexTemp = 0; nTileIndexTemp < this->block_.params_.nTileCntL2; nTileIndexTemp++) {
318- uint64_t nTileIndex = reverse ? (this->block_.params_.nTileCntL2 - nTileIndexTemp - 1) : nTileIndexTemp;
319- this->block_.UpdateBlockCnt(mTileIndex, nTileIndex);
320- this->block_.InitBlockIndex(index);
321- SetFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG));
322- SetFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + 1));
323- for (uint64_t j = 0; j < this->block_.params_.realRound; j++) {
324- tempCGlobal = tempCGlobal_[baseSize_ * (GetCurrentBlockIdx() * 2 + pingPongId)];
325- if (this->block_.params_.rowOrder == 0) {
326- this->block_.UpdateBasicIndex(j); // 使能错位分核更新Index
327- }
328- if (this->block_.params_.index < this->block_.params_.totalTileCnt) {
329- this->block_.UpdateBlockParams(mTileIndex, nTileIndex);
330- this->block_.template CalcGMOffset<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE>(mTileIndex, nTileIndex);
331- AicProcess(tempCGlobal, enAtomic, aicNeedWaitAiv, pingPongId);
332- if constexpr (FIXPIPE_OPT == FIXPIPE_OPT_SELECT::VEC_NZ2ND_UNALIGNOUT) {
333- AivNz2NdProcess(tempCGlobal, pingPongId);
334- } else {
335- AivProcess(tempCGlobal, pingPongId);
336- }
337- aicNeedWaitAiv = aicNeedWaitAiv || bool(pingPongId);
338- pingPongId = (pingPongId + 1) & 1;
339- }
340- this->block_.UpdateBlockIndex();
341- }
342- WaitFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG));
343- WaitFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + 1));
344- }
345- }
346- if (this->block_.params_.isHf32) {
347- this->mm_.SetHF32(false, 0);
348- }
349- PipeBarrier<PIPE_ALL>();
350- return;
351-}
352- 
353-// Current Kernel support only nd2nzA. No need to do nd2nz for B.
354-template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE = MatmulBaseBlock,
355- const MatmulConfig& MM_CFG = MM_CFG_NO_PRELOAD, class MM_CB = MatmulCallBackFunc<nullptr, nullptr, nullptr>,
356- FIXPIPE_OPT_SELECT FIXPIPE_OPT = FIXPIPE_OPT_SELECT::BASE_ENABLE_ALIGNOUT>
357-class MatmulBaseAToNZWithBL1FixpipeKernel
358- : public MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT> {
359- struct BaseUnAlignedNKernelParams {
360- GM_ADDR aGMNZ;
361- GM_ADDR workspaceGMNZ;
362- uint64_t baseAN;
363- uint64_t baseAD;
364- };
365- 
366-public:
367- __aicore__ inline MatmulBaseAToNZWithBL1FixpipeKernel() {}
368- 
369- __aicore__ inline void Init(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR biasGM, GM_ADDR offsetWGM,
370- GM_ADDR workspaceGM, const MatmulTilingData* tilingData, TPipe* pipe);
371- 
372- __aicore__ inline void Process(uint64_t index = 0, uint8_t enAtomic = 0);
373- 
374-protected:
375- using C_T = typename C_TYPE::T;
376- BaseUnAlignedNKernelParams fixpipeInnerParams_;
377-};
378- 
379-template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,
380- class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>
381-__aicore__ inline void
382-MatmulBaseAToNZWithBL1FixpipeKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::Init(
383- GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR biasGM, GM_ADDR offsetWGM, GM_ADDR workspaceGM,
384- const MatmulTilingData* matmulTilingData, TPipe* pipe)
385-{
386- GetSizeC0<C_T>(this->c0Size_);
387- uint64_t cDtypeSize = sizeof(C_T);
388- uint64_t alignedN = AlignUp(matmulTilingData->matmulTiling.N, MM_ALIGN_SIZE / cDtypeSize);
389- uint64_t baseSize = alignedN * matmulTilingData->matmulTiling.baseM;
390- 
391- MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB>::Init(
392- workspaceGM + baseSize * NUM_TWO * matmulTilingData->matmulTiling.usedCoreNum * cDtypeSize, bGM, cGM, biasGM,
393- offsetWGM, workspaceGM, matmulTilingData, pipe);
394- this->fixpipeInnerParams_.baseAN = this->block_.matmulTilingData_->baseAN;
395- this->fixpipeInnerParams_.baseAD = this->block_.matmulTilingData_->baseAD;
396- this->fixpipeInnerParams_.aGMNZ = aGM;
397- this->fixpipeInnerParams_.workspaceGMNZ =
398- workspaceGM + this->baseSize_ * NUM_TWO * this->block_.matmulTilingData_->matmulTiling.usedCoreNum * cDtypeSize;
399- this->mm_.SetOrgShape(this->block_.params_.alignedOriM, this->block_.matmulTilingData_->matmulTiling.N,
400- this->block_.params_.alignedKaSize, this->block_.matmulTilingData_->matmulTiling.Kb,
401- this->alignedN_);
402-}
403- 
404-template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE, class BLOCK_TYPE, const MatmulConfig& MM_CFG,
405- class MM_CB, FIXPIPE_OPT_SELECT FIXPIPE_OPT>
406-__aicore__ inline void
407-MatmulBaseAToNZWithBL1FixpipeKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::Process(
408- uint64_t index, uint8_t enAtomic)
409-{
410- if ASCEND_IS_AIV {
411- MatrixAtoNZV2<typename A_TYPE::T>(this->fixpipeInnerParams_.workspaceGMNZ, this->fixpipeInnerParams_.aGMNZ,
412- this->block_.matmulTilingData_->matmulTiling, A_TYPE::isTrans, this->ubBuf_,
413- this->fixpipeInnerParams_.baseAN, this->fixpipeInnerParams_.baseAD);
414- SyncAll();
415- CrossCoreSetFlag<0x2, PIPE_MTE3>(CV_FLAG);
416- }
417- if ASCEND_IS_AIC {
418- CrossCoreWaitFlag(CV_FLAG);
419- }
420- MatmulBaseUnalignedNKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK_TYPE, MM_CFG, MM_CB, FIXPIPE_OPT>::Process(
421- index, enAtomic);
422-}
423- 
424-#endif
Mmatmul/mat_mul_v3/op_kernel/mat_mul_v3.cpp+10-1
@@ -24,10 +24,13 @@
24#include "mat_mul_unaligned_deterministic_splitk_kernel.h"24#include "mat_mul_unaligned_deterministic_splitk_kernel.h"
25#include "mat_mul_unaligned_sc_splitk_kernel.h"25#include "mat_mul_unaligned_sc_splitk_kernel.h"
26#include "mat_mul_unaligned_sc_splitk_kernel_gm_to_l1.h"26#include "mat_mul_unaligned_sc_splitk_kernel_gm_to_l1.h"
27-#include "mat_mul_optimized_fixpipe_algorithm.h"
28#include "mat_mul_l1_full_load.h"27#include "mat_mul_l1_full_load.h"
29#include "mat_mul_v3_tiling_key.h"28#include "mat_mul_v3_tiling_key.h"
30 29 
30+#if defined(__CCE_AICORE__) && __CCE_AICORE__ >= 220
31+#include "mat_mul_base_kernel_vec_nz2nd.h"
32+#include "mat_mul_optimized_fixpipe_algorithm.h"
33+#endif
31 34 
32#if defined(__CCE_AICORE__) && __CCE_AICORE__ == 220 || (defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3003 || __NPU_ARCH__ == 3113))35#if defined(__CCE_AICORE__) && __CCE_AICORE__ == 220 || (defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3003 || __NPU_ARCH__ == 3113))
33#include "mat_mul_multi_core_splitk_kernel.h"36#include "mat_mul_multi_core_splitk_kernel.h"
@@ -206,6 +209,12 @@ __global__ __aicore__ void mat_mul_v3(
206 MMV3_IMPL_CLASS( 209 MMV3_IMPL_CLASS(
207 MatmulBaseKernel, format_x1, MatmulBaseBlock, MM_CFG_NO_PRELOAD 210 MatmulBaseKernel, format_x1, MatmulBaseBlock, MM_CFG_NO_PRELOAD
208 );211 );
212+ } else if constexpr (
213+ LOADMODE == MAT_MUL_V3_BASE_FULLLOAD && SPLITCOREMODE == MAT_MUL_V3_BASE_SPLIT_K &&
214+ FIXOPTI == MAT_MUL_V3_VEC_NZ2ND_UNALIGNOUT && MIXND2NZ == MAT_MUL_V3_MIXND2NZ_FALSE) {
215+ MMV3_IMPL_C_CLASS(
216+ MatmulBaseVectorNz2NdKernel, format_x1, CubeFormat::NZ, MatmulBaseBlock, MM_CFG_NO_PRELOAD
217+ );
209 }218 }
210#else219#else
211 if constexpr (220 if constexpr (
Mmatmul/mat_mul_v3/op_kernel/mat_mul_v3_tiling_key.h+8-0
@@ -98,6 +98,14 @@ ASCENDC_TPL_SEL(
98 ASCENDC_TPL_UINT_SEL(MIXND2NZ, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_MIXND2NZ_FALSE), 98 ASCENDC_TPL_UINT_SEL(MIXND2NZ, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_MIXND2NZ_FALSE),
99 ASCENDC_TPL_UINT_SEL(SPECIALOPT, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_K_SHIFT, MAT_MUL_V3_K_NOT_SHIFT), 99 ASCENDC_TPL_UINT_SEL(SPECIALOPT, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_K_SHIFT, MAT_MUL_V3_K_NOT_SHIFT),
100 ASCENDC_TPL_UINT_SEL(FP32ADDMM, ASCENDC_TPL_UI_LIST, FP32_ADDMM_DISABLE)),100 ASCENDC_TPL_UINT_SEL(FP32ADDMM, ASCENDC_TPL_UI_LIST, FP32_ADDMM_DISABLE)),
101+ ASCENDC_TPL_ARGS_SEL(
102+ ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
103+ ASCENDC_TPL_UINT_SEL(LOADMODE, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_BASE_FULLLOAD),
104+ ASCENDC_TPL_UINT_SEL(SPLITCOREMODE, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_BASE_SPLIT_K),
105+ ASCENDC_TPL_UINT_SEL(FIXOPTI, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_VEC_NZ2ND_UNALIGNOUT),
106+ ASCENDC_TPL_UINT_SEL(MIXND2NZ, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_MIXND2NZ_FALSE),
107+ ASCENDC_TPL_UINT_SEL(SPECIALOPT, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_K_NOT_SHIFT),
108+ ASCENDC_TPL_UINT_SEL(FP32ADDMM, ASCENDC_TPL_UI_LIST, FP32_ADDMM_DISABLE)),
101);109);
102#else110#else
103ASCENDC_TPL_SEL(111ASCENDC_TPL_SEL(