Star944
1860
代码
代码
Issues229
Pull Requests324
Actions
项目讨论
Wiki
分析
Star944
1860
已合并
segment_sum算子性能优化 #2127
z30075199创建于 2月26日
segment_sum算子性能优化 #2127
已合并
z30075199创建于 2月26日
从z30075199/ops-nn_zxmodify_ss_simd合入到cann/ops-nnmaster
共 5 个文件变更+51-43
Mindex/segment_sum/op_host/arch35/segment_sum_simd_tiling.cpp+9-6
@@ -25,6 +25,8 @@ static constexpr int64_t INNER_ADD_NUM = 128;
25static constexpr int64_t BASE_BLOCK_ALIGN = 512;25static constexpr int64_t BASE_BLOCK_ALIGN = 512;
26static constexpr int64_t SINGLE_CORE_THRESHOLD = 4 * 1024;26static constexpr int64_t SINGLE_CORE_THRESHOLD = 4 * 1024;
27static constexpr int64_t NUM_TWO = 2;27static constexpr int64_t NUM_TWO = 2;
28+static constexpr int64_t Y_BUFFER_NUM = 3;
29+static constexpr int64_t MIN_OUTTERS = 8;
CANN-robot
CANN-robotCANN-robot2月26日

命名规范: 常量名 'MIN_OUTTERS' 存在拼写错误，应为 'MIN_OUTERS'（'outer' 的正确拼写）。该常量用于表示最小外部维度数，拼写错误会影响代码可读性和维护性。

问题类型: 命名规范 文件路径: index/segment_sum/op_host/arch35/segment_sum_simd_tiling.cpp 行号: 29 问题代码:

static constexpr int64_t MIN_OUTTERS = 8;

修改建议:

将常量名修改为正确的拼写：static constexpr int64_t MIN_OUTERS = 8; 并同步更新所有使用该常量的地方（如第154行）。

此评论由代码审查工具自动生成

likedislike
28static constexpr uint64_t NUM_FOUR = 4;30static constexpr uint64_t NUM_FOUR = 4;
29static constexpr uint64_t MIN_INNER_SIZE = 256;31static constexpr uint64_t MIN_INNER_SIZE = 256;
30static constexpr size_t WS_SYS_SIZE = static_cast<size_t>(16 * 1024 * 1024);32static constexpr size_t WS_SYS_SIZE = static_cast<size_t>(16 * 1024 * 1024);
@@ -95,9 +97,9 @@ void SegmentSumSimdTiling::AutoTilingRowCol(int64_t& rowTileNum, int64_t& colTil
95 }97 }
96 98 
97 std::sort(allTiling.begin(), allTiling.end(), [](const std::vector<int64_t>& a, const std::vector<int64_t>& b) {99 std::sort(allTiling.begin(), allTiling.end(), [](const std::vector<int64_t>& a, const std::vector<int64_t>& b) {
98- constexpr int MIndex = 0;100+ constexpr int NIndex = 1;
99 constexpr int DeltaIndex = 3;101 constexpr int DeltaIndex = 3;
100- return std::make_pair(a[DeltaIndex], a[MIndex]) < std::make_pair(b[DeltaIndex], b[MIndex]);102+ return std::make_pair(a[NIndex], a[DeltaIndex]) < std::make_pair(b[NIndex], b[DeltaIndex]);
101 });103 });
102 104 
103 while (allTiling.size() > 1 && outerDim_ / allTiling[0][0] < std::min(NUM_FOUR, outerDim_)) {105 while (allTiling.size() > 1 && outerDim_ / allTiling[0][0] < std::min(NUM_FOUR, outerDim_)) {
@@ -149,14 +151,15 @@ void SegmentSumSimdTiling::DoSplitColUBTiling(int64_t availableUbsize)
149 int64_t innerSizeAlign = Ops::Base::CeilAlign(normalCoreInnerNum_ * valueTypeBytes_, ubBlockSize_);151 int64_t innerSizeAlign = Ops::Base::CeilAlign(normalCoreInnerNum_ * valueTypeBytes_, ubBlockSize_);
150 int64_t tmpColSize = BLOCK_TILING_THRES + idTypeBytes_;152 int64_t tmpColSize = BLOCK_TILING_THRES + idTypeBytes_;
151 int64_t tmpRowNum = availableUbsize / tmpColSize;153 int64_t tmpRowNum = availableUbsize / tmpColSize;
152- while (tmpRowNum - 1 > normalCoreOutterNum_) {154+ int64_t minOutter = std::min(MIN_OUTTERS, normalCoreOutterNum_);
155+ while (tmpRowNum - Y_BUFFER_NUM > minOutter) {
153 if (tmpColSize >= static_cast<int64_t>(innerSizeAlign + idTypeBytes_)) {156 if (tmpColSize >= static_cast<int64_t>(innerSizeAlign + idTypeBytes_)) {
154 break;157 break;
155 }158 }
156 tmpColSize += INNER_ADD_NUM;159 tmpColSize += INNER_ADD_NUM;
157 tmpRowNum = availableUbsize / tmpColSize;160 tmpRowNum = availableUbsize / tmpColSize;
158 }161 }
159- int64_t rowNumInUb = std::min(tmpRowNum - 1, normalCoreOutterNum_);162+ int64_t rowNumInUb = std::min(tmpRowNum - Y_BUFFER_NUM, normalCoreOutterNum_);
160 int64_t colSizeInUb = std::min(static_cast<int64_t>(tmpColSize - idTypeBytes_), innerSizeAlign);163 int64_t colSizeInUb = std::min(static_cast<int64_t>(tmpColSize - idTypeBytes_), innerSizeAlign);
161 164 
162 xBufferSize_ = rowNumInUb * colSizeInUb;165 xBufferSize_ = rowNumInUb * colSizeInUb;
@@ -199,7 +202,7 @@ void SegmentSumSimdTiling::DoUBTiling()
199 int64_t innerSizeAlign = Ops::Base::CeilAlign(normalCoreInnerNum_ * valueTypeBytes_, ubBlockSize_);202 int64_t innerSizeAlign = Ops::Base::CeilAlign(normalCoreInnerNum_ * valueTypeBytes_, ubBlockSize_);
200 int64_t tmpColSize = innerSizeAlign + idTypeBytes_;203 int64_t tmpColSize = innerSizeAlign + idTypeBytes_;
201 int64_t tmpRowNum = availableUbsize / tmpColSize;204 int64_t tmpRowNum = availableUbsize / tmpColSize;
202- int64_t rowNumInUb = std::min(tmpRowNum - 1, normalCoreOutterNum_);205+ int64_t rowNumInUb = std::min(tmpRowNum - Y_BUFFER_NUM, normalCoreOutterNum_);
203 206 
204 xBufferSize_ = rowNumInUb * innerSizeAlign;207 xBufferSize_ = rowNumInUb * innerSizeAlign;
205 segmentIdBufferSize_ = Ops::Base::CeilAlign(rowNumInUb * idTypeBytes_, ubBlockSize_);208 segmentIdBufferSize_ = Ops::Base::CeilAlign(rowNumInUb * idTypeBytes_, ubBlockSize_);
@@ -239,7 +242,7 @@ void SegmentSumSimdTiling::DoMultCoreAddTiling()
239 int64_t mulAddUbsize = ubSize_;242 int64_t mulAddUbsize = ubSize_;
240 multAddIdsBufferSize_ = Ops::Base::CeilAlign(NUM_TWO * blockNumInRow_ * idTypeBytes_, ubBlockSize_);243 multAddIdsBufferSize_ = Ops::Base::CeilAlign(NUM_TWO * blockNumInRow_ * idTypeBytes_, ubBlockSize_);
241 mulAddUbsize -= multAddIdsBufferSize_;244 mulAddUbsize -= multAddIdsBufferSize_;
242- mulAddUbsize /= blockNumInRow_ * NUM_TWO + 1;245+ mulAddUbsize /= blockNumInRow_ * NUM_TWO + Y_BUFFER_NUM;
243 int64_t availableInnerUb = Ops::Base::FloorAlign(mulAddUbsize, static_cast<int64_t>(ubBlockSize_));246 int64_t availableInnerUb = Ops::Base::FloorAlign(mulAddUbsize, static_cast<int64_t>(ubBlockSize_));
244 247 
245 int64_t innerNumInUb = availableInnerUb / valueTypeBytes_;248 int64_t innerNumInUb = availableInnerUb / valueTypeBytes_;
Mindex/segment_sum/op_kernel/arch35/segment_sum_simd.h+16-13
@@ -19,7 +19,8 @@
19namespace SegmentSum {19namespace SegmentSum {
20using namespace AscendC;20using namespace AscendC;
21 21 
22-constexpr uint32_t BUFFER_NUM = 1;22+constexpr uint32_t X_BUFFER_NUM = 1;
23+constexpr uint32_t TMP_BUFFER_NUM = 2;
23 24 
24 25 
25template <typename T1, typename T2>26template <typename T1, typename T2>
@@ -40,8 +41,9 @@ private:
40 GlobalTensor<T2> segmentIdsGm_;41 GlobalTensor<T2> segmentIdsGm_;
41 GlobalTensor<T1> yGm_;42 GlobalTensor<T1> yGm_;
42 43
43- TQue<QuePosition::VECIN, BUFFER_NUM> xQue_;44+ TQue<QuePosition::VECIN, X_BUFFER_NUM> xQue_;
44- TQue<QuePosition::VECIN, BUFFER_NUM> segmentIdsQue_;45+ TQue<QuePosition::VECIN, X_BUFFER_NUM> segmentIdsQue_;
46+ TQue<QuePosition::VECOUT, TMP_BUFFER_NUM> tmpQue_;
45 TBuf<QuePosition::VECCALC> yBuf_;47 TBuf<QuePosition::VECCALC> yBuf_;
46 48 
47 const SegmentSumSimdTilingData* tilingData_;49 const SegmentSumSimdTilingData* tilingData_;
@@ -101,8 +103,9 @@ __aicore__ inline void SegmentSumSimd<T1, T2>::Init(
101 segmentIdsGm_.SetGlobalBuffer((__gm__ T2*)segmentIds + rowGmOffset_);103 segmentIdsGm_.SetGlobalBuffer((__gm__ T2*)segmentIds + rowGmOffset_);
102 yGm_.SetGlobalBuffer((__gm__ T1*)y + colGmOffset_);104 yGm_.SetGlobalBuffer((__gm__ T1*)y + colGmOffset_);
103 105 
104- pipeIn.InitBuffer(xQue_, BUFFER_NUM, tilingData_->xBufferSize); // 需要满足double block对齐106+ pipeIn.InitBuffer(xQue_, X_BUFFER_NUM, tilingData_->xBufferSize); // 需要满足double block对齐
105- pipeIn.InitBuffer(segmentIdsQue_, BUFFER_NUM, tilingData_->segmentIdBufferSize);107+ pipeIn.InitBuffer(segmentIdsQue_, X_BUFFER_NUM, tilingData_->segmentIdBufferSize);
108+ pipeIn.InitBuffer(tmpQue_, TMP_BUFFER_NUM, tilingData_->yBufferSize);
106 pipeIn.InitBuffer(yBuf_, tilingData_->yBufferSize);109 pipeIn.InitBuffer(yBuf_, tilingData_->yBufferSize);
107 110 
108}111}
@@ -176,20 +179,20 @@ __aicore__ inline void SegmentSumSimd<T1, T2>::ComputeSumAndCopyOut(LocalTensor<
176 Copy(yLocal, xLocal[i * curLoopInnersAlign], curLoopInners);179 Copy(yLocal, xLocal[i * curLoopInnersAlign], curLoopInners);
177 preId_ = curId;180 preId_ = curId;
178 } else { // curId != preId_181 } else { // curId != preId_
179- event_t eventId = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE3));182+ LocalTensor<T1> tmpLocal = tmpQue_.AllocTensor<T1>();
180- SetFlag<HardEvent::V_MTE3>(eventId);183+ Copy(tmpLocal, yLocal, curLoopInners);
181- WaitFlag<HardEvent::V_MTE3>(eventId);184+ tmpQue_.EnQue(tmpLocal);
185+ LocalTensor<T1> outLocal = tmpQue_.DeQue<T1>();
186+ 
182 if (isFirstId_ && !isStartRowCore_) {187 if (isFirstId_ && !isStartRowCore_) {
183 SetAtomicAdd<T1>();188 SetAtomicAdd<T1>();
184- CopyOutY(yLocal, curLoopInners, preId_, colOffset);189+ CopyOutY(outLocal, curLoopInners, preId_, colOffset);
185 SetAtomicNone();190 SetAtomicNone();
186 isFirstId_ = false;191 isFirstId_ = false;
187 } else {192 } else {
188- CopyOutY(yLocal, curLoopInners, preId_, colOffset);193+ CopyOutY(outLocal, curLoopInners, preId_, colOffset);
189 }194 }
190- event_t eventId1 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE3_V));195+ tmpQue_.FreeTensor(outLocal);
191- SetFlag<HardEvent::MTE3_V>(eventId1);
192- WaitFlag<HardEvent::MTE3_V>(eventId1);
193 196 
194 preId_ = curId;197 preId_ = curId;
195 Copy(yLocal, xLocal[i * curLoopInnersAlign], curLoopInners);198 Copy(yLocal, xLocal[i * curLoopInnersAlign], curLoopInners);
Mindex/segment_sum/op_kernel/arch35/segment_sum_simd_determ.h+14-12
@@ -44,8 +44,9 @@ private:
44 GlobalTensor<T1> sumWorkspace_;44 GlobalTensor<T1> sumWorkspace_;
45 GlobalTensor<T2> segIdWorkspace_;45 GlobalTensor<T2> segIdWorkspace_;
46 46
47- TQue<QuePosition::VECIN, BUFFER_NUM> xQue_;47+ TQue<QuePosition::VECIN, X_BUFFER_NUM> xQue_;
48- TQue<QuePosition::VECIN, BUFFER_NUM> segmentIdsQue_;48+ TQue<QuePosition::VECIN, X_BUFFER_NUM> segmentIdsQue_;
49+ TQue<QuePosition::VECOUT, TMP_BUFFER_NUM> tmpQue_;
49 TBuf<QuePosition::VECCALC> yBuf_;50 TBuf<QuePosition::VECCALC> yBuf_;
50 TBuf<QuePosition::VECCALC> tmpBuf_;51 TBuf<QuePosition::VECCALC> tmpBuf_;
51 52 
@@ -111,8 +112,9 @@ __aicore__ inline void SegmentSumSimdDeterm<T1, T2>::Init(
111 sumWorkspace_.SetGlobalBuffer((__gm__ T1*)workspace + rowCoreIdx_ * DOUBLE * tilingData_->innerDim + colGmOffset_);112 sumWorkspace_.SetGlobalBuffer((__gm__ T1*)workspace + rowCoreIdx_ * DOUBLE * tilingData_->innerDim + colGmOffset_);
112 segIdWorkspace_.SetGlobalBuffer((__gm__ T2*)workspace + segIdAddrOffset + rowCoreIdx_ * DOUBLE);113 segIdWorkspace_.SetGlobalBuffer((__gm__ T2*)workspace + segIdAddrOffset + rowCoreIdx_ * DOUBLE);
113 114 
114- pipeIn.InitBuffer(xQue_, BUFFER_NUM, tilingData_->xBufferSize);115+ pipeIn.InitBuffer(xQue_, X_BUFFER_NUM, tilingData_->xBufferSize);
115- pipeIn.InitBuffer(segmentIdsQue_, BUFFER_NUM, tilingData_->segmentIdBufferSize);116+ pipeIn.InitBuffer(segmentIdsQue_, X_BUFFER_NUM, tilingData_->segmentIdBufferSize);
117+ pipeIn.InitBuffer(tmpQue_, TMP_BUFFER_NUM, tilingData_->yBufferSize);
116 pipeIn.InitBuffer(yBuf_, tilingData_->yBufferSize);118 pipeIn.InitBuffer(yBuf_, tilingData_->yBufferSize);
117 pipeIn.InitBuffer(tmpBuf_, platform::GetUbBlockSize()); // 放头尾id119 pipeIn.InitBuffer(tmpBuf_, platform::GetUbBlockSize()); // 放头尾id
118 120 
@@ -209,19 +211,19 @@ __aicore__ inline void SegmentSumSimdDeterm<T1, T2>::ComputeSumAndCopyOut(LocalT
209 Copy(yLocal, xLocal[i * curLoopInnersAlign], curLoopInners);211 Copy(yLocal, xLocal[i * curLoopInnersAlign], curLoopInners);
210 preId_ = curId;212 preId_ = curId;
211 } else { // curId != preId_213 } else { // curId != preId_
212- event_t eventId = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE3));214+ LocalTensor<T1> tmpLocal = tmpQue_.AllocTensor<T1>();
213- SetFlag<HardEvent::V_MTE3>(eventId);215+ Copy(tmpLocal, yLocal, curLoopInners);
214- WaitFlag<HardEvent::V_MTE3>(eventId);216+ tmpQue_.EnQue(tmpLocal);
217+ LocalTensor<T1> outLocal = tmpQue_.DeQue<T1>();
218+ 
215 if (isFirstId_ && !isStartRowCore_) {219 if (isFirstId_ && !isStartRowCore_) {
216- CopyOutSumWorkspace(yLocal, curLoopInners, colOffset, 0);220+ CopyOutSumWorkspace(outLocal, curLoopInners, colOffset, 0);
217 position0_ = preId_;221 position0_ = preId_;
218 isFirstId_ = false;222 isFirstId_ = false;
219 } else {223 } else {
220- CopyOutY(yLocal, curLoopInners, preId_, colOffset);224+ CopyOutY(outLocal, curLoopInners, preId_, colOffset);
221 }225 }
222- event_t eventId1 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE3_V));226+ tmpQue_.FreeTensor(outLocal);
223- SetFlag<HardEvent::MTE3_V>(eventId1);
224- WaitFlag<HardEvent::MTE3_V>(eventId1);
225 227 
226 preId_ = curId;228 preId_ = curId;
227 Copy(yLocal, xLocal[i * curLoopInnersAlign], curLoopInners);229 Copy(yLocal, xLocal[i * curLoopInnersAlign], curLoopInners);
Mindex/segment_sum/op_kernel/arch35/segment_sum_simd_mult_core_add.h+11-11
@@ -36,7 +36,8 @@ private:
36 GlobalTensor<T1> sumWorkspace_;36 GlobalTensor<T1> sumWorkspace_;
37 GlobalTensor<T2> segIdWorkspace_;37 GlobalTensor<T2> segIdWorkspace_;
38 38 
39- TQue<QuePosition::VECIN, BUFFER_NUM> xQue_;39+ TQue<QuePosition::VECIN, X_BUFFER_NUM> xQue_;
40+ TQue<QuePosition::VECOUT, TMP_BUFFER_NUM> tmpQue_;
40 TBuf<QuePosition::VECCALC> segmentIdsBuf_;41 TBuf<QuePosition::VECCALC> segmentIdsBuf_;
41 TBuf<QuePosition::VECCALC> yBuf_;42 TBuf<QuePosition::VECCALC> yBuf_;
42 43 
@@ -77,7 +78,8 @@ __aicore__ inline void SegmentSumMultiCoreAdd<T1, T2>::Init(GM_ADDR y, GM_ADDR w
77 sumWorkspace_.SetGlobalBuffer((__gm__ T1*)workspace + colGmOffset_);78 sumWorkspace_.SetGlobalBuffer((__gm__ T1*)workspace + colGmOffset_);
78 segIdWorkspace_.SetGlobalBuffer((__gm__ T2*)workspace + segIdAddrOffset);79 segIdWorkspace_.SetGlobalBuffer((__gm__ T2*)workspace + segIdAddrOffset);
79 80 
80- pipeIn.InitBuffer(xQue_, BUFFER_NUM, tilingData_->multAddXBufferSize);81+ pipeIn.InitBuffer(xQue_, X_BUFFER_NUM, tilingData_->multAddXBufferSize);
82+ pipeIn.InitBuffer(tmpQue_, TMP_BUFFER_NUM, tilingData_->multAddYBufferSize);
81 pipeIn.InitBuffer(segmentIdsBuf_, tilingData_->multAddIdsBufferSize);83 pipeIn.InitBuffer(segmentIdsBuf_, tilingData_->multAddIdsBufferSize);
82 pipeIn.InitBuffer(yBuf_, tilingData_->multAddYBufferSize);84 pipeIn.InitBuffer(yBuf_, tilingData_->multAddYBufferSize);
83}85}
@@ -154,17 +156,15 @@ __aicore__ inline void SegmentSumMultiCoreAdd<T1, T2>::ComputeAndCopyOut(LocalTe
154 } else if (curId == -1) {156 } else if (curId == -1) {
155 continue;157 continue;
156 } else { // curId != preId158 } else { // curId != preId
157- event_t eventId2 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE3));159+ LocalTensor<T1> tmpLocal = tmpQue_.AllocTensor<T1>();
158- SetFlag<HardEvent::V_MTE3>(eventId2);160+ Copy(tmpLocal, yLocal, curLoopInners);
159- WaitFlag<HardEvent::V_MTE3>(eventId2);161+ tmpQue_.EnQue(tmpLocal);
162+ LocalTensor<T1> outLocal = tmpQue_.DeQue<T1>();
163+ 
164+ CopyOutY(outLocal, curLoopInners, preId, colOffset);
165+ tmpQue_.FreeTensor(outLocal);
160 166 
161- CopyOutY(yLocal, curLoopInners, preId, colOffset);
162 preId = curId;167 preId = curId;
163- 
164- event_t eventId3 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE3_V));
165- SetFlag<HardEvent::MTE3_V>(eventId3);
166- WaitFlag<HardEvent::MTE3_V>(eventId3);
167- 
168 Copy(yLocal, xLocal[i * curLoopInnersAlign], curLoopInners);168 Copy(yLocal, xLocal[i * curLoopInnersAlign], curLoopInners);
169 }169 }
170 }170 }
Mindex/segment_sum/op_kernel/arch35/segment_sum_simt.h+1-1
@@ -208,7 +208,7 @@ __simt_vf__ __launch_bounds__(MAX_THREAD_NUM) inline void ComputeInWs(__gm__ TX*
208 if (segmentIdsWs[i] == -1) {208 if (segmentIdsWs[i] == -1) {
209 continue;209 continue;
210 }210 }
211- if (threadIdx.y > 0) {211+ if (i > 0) {
212 int32_t prevId = (segmentIdsWs[i - 1] == -1) ? (i - 2) : (i - 1);212 int32_t prevId = (segmentIdsWs[i - 1] == -1) ? (i - 2) : (i - 1);
213 if (segmentIdsWs[i] == segmentIdsWs[prevId]) { // 如果当前线程组对应id非0且跟前一行id相同，则跳过213 if (segmentIdsWs[i] == segmentIdsWs[prevId]) { // 如果当前线程组对应id非0且跟前一行id相同，则跳过
214 continue;214 continue;