Star86
179
代码
代码
Issues9
Pull Requests50
Actions
项目讨论
Wiki
分析
Star86
179
已合并
fix A2 bevpoolv3 bug with large channel #2075
gitdzy创建于 6月3日
fix A2 bevpoolv3 bug with large channel #2075
已合并
gitdzy创建于 6月3日
从gitdzy/DrivingSDK_a5bevpoolv3grad合入到Ascend/DrivingSDKmaster
共 5 个文件变更+300-170
Mkernels/bev_pool_v3/op_host/bev_pool_v3.cpp+1-75
@@ -1,5 +1,5 @@
1/*1/*
2- * Copyright (c) Huawei Technologies Co., Ltd. 2022-2023. All rights reserved.2+ * Copyright (c) Huawei Technologies Co., Ltd. 2026. All rights reserved.
3 */3 */
4#include <graph/types.h>4#include <graph/types.h>
5#include <log/log.h>5#include <log/log.h>
@@ -178,81 +178,7 @@ class BEVPoolV3 : public OpDef {
178 }178 }
179};179};
180 180 
181-/**
182- * @brief: BEVPoolGrad, the backward of bev_pool
183- * @par Inputs:
184- * grad_out: input grad, 5D tensor(b, d, h, w, c), dtype: float32, format:
185- * NDHWC, ND geom_feat: input coords, 2D tensor(n, 4), dtype: int32, format: ND
186- * interval_starts: starting position for pooled point, 1D tensor(n_interval),
187- * dtype: int32, format: ND interval_lengths: the number of points in each
188- * interval, 1D tensor(n_interval), dtype: int32, format: ND
189- * @par Outputs:
190- * grad_feat: output grad, 2D tensor(n, c), dtype: float32, format: ND
191- * @par Attributes:
192- **/
193-class BEVPoolV3Grad : public OpDef {
194- public:
195- explicit BEVPoolV3Grad(const char *name) : OpDef(name) {
196- this->Input("grad_out")
197- .ParamType(REQUIRED)
198- .DataType({ge::DT_FLOAT})
199- .Format({ge::FORMAT_ND})
200- .AutoContiguous()
201- .UnknownShapeFormat({ge::FORMAT_ND});
202- this->Input("depth")
203- .ParamType(OPTIONAL)
204- .DataType({ge::DT_FLOAT})
205- .Format({ge::FORMAT_ND})
206- .AutoContiguous()
207- .UnknownShapeFormat({ge::FORMAT_ND});
208- this->Input("feat")
209- .ParamType(REQUIRED)
210- .DataType({ge::DT_FLOAT})
211- .Format({ge::FORMAT_ND})
212- .AutoContiguous()
213- .UnknownShapeFormat({ge::FORMAT_ND});
214- this->Input("ranks_depth")
215- .ParamType(OPTIONAL)
216- .DataType({ge::DT_INT32})
217- .Format({ge::FORMAT_ND})
218- .AutoContiguous()
219- .UnknownShapeFormat({ge::FORMAT_ND});
220- this->Input("ranks_feat")
221- .ParamType(OPTIONAL)
222- .DataType({ge::DT_INT32})
223- .Format({ge::FORMAT_ND})
224- .AutoContiguous()
225- .UnknownShapeFormat({ge::FORMAT_ND});
226- this->Input("ranks_bev")
227- .ParamType(REQUIRED)
228- .DataType({ge::DT_INT32})
229- .Format({ge::FORMAT_ND})
230- .AutoContiguous()
231- .UnknownShapeFormat({ge::FORMAT_ND});
232- 
233- this->Attr("with_depth").Bool();
234- 
235- this->Output("grad_depth")
236- .ParamType(OPTIONAL)
237- .DataType({ge::DT_FLOAT})
238- .Format({ge::FORMAT_ND})
239- .UnknownShapeFormat({ge::FORMAT_ND});
240- this->Output("grad_feat")
241- .ParamType(REQUIRED)
242- .DataType({ge::DT_FLOAT})
243- .Format({ge::FORMAT_ND})
244- .UnknownShapeFormat({ge::FORMAT_ND});
245- 
246- this->AICore().SetTiling(optiling::TilingForBEVPoolV3<true>);
247- this->AICore().AddConfig("ascend910b");
248- this->AICore().AddConfig("ascend910_93");
249-#if __DRIVING_HOST_AICORE__ == 310
250- this->AICore().AddConfig("ascend950");
251-#endif
252- }
253-};
254IMPL_OP_INFERSHAPE(BEVPoolV3).InferShape(InferShapeForBEVPoolV3).InferDataType(InferDataTypeForBEVPoolV3);181IMPL_OP_INFERSHAPE(BEVPoolV3).InferShape(InferShapeForBEVPoolV3).InferDataType(InferDataTypeForBEVPoolV3);
255 182 
256OP_ADD(BEVPoolV3);183OP_ADD(BEVPoolV3);
257-OP_ADD(BEVPoolV3Grad);
258} // namespace ops184} // namespace ops
Akernels/bev_pool_v3/op_host/bev_pool_v3_grad.cpp+204-0
@@ -0,0 +1,204 @@
1+/*
2+ * Copyright (c) Huawei Technologies Co., Ltd. 2022-2023. All rights reserved.
3+ */
4+#include <graph/types.h>
5+#include <log/log.h>
6+#include <register/op_def.h>
7+ 
8+#include <cstdint>
9+ 
10+#include "bev_pool_v3_grad_tiling.h"
11+#include "ge/utils.h"
12+#include "common/op_host/common.h"
13+#include "register/op_def_registry.h"
14+#include "tiling/platform/platform_ascendc.h"
15+ 
16+namespace {
17+constexpr size_t INPUT_FEAT = 1;
18+constexpr size_t INPUT_FEAT_GRAD = 2;
19+constexpr size_t INPUT_RANKS_BEV = 4;
20+constexpr size_t INPUT_RANKS_BEV_GRAD = 5;
21+constexpr uint64_t RANK_NUM_PER_TASK = 1024;
22+constexpr int32_t ONE_BLK_SIZE = 8;
23+constexpr int32_t RESERVE_UB = 10 * 1024; // 10 KB
24+constexpr size_t ATTR_B_IDX = 1;
25+constexpr size_t ATTR_D_IDX = 2;
26+constexpr size_t ATTR_H_IDX = 3;
27+constexpr size_t ATTR_W_IDX = 4;
28+constexpr size_t ATTR_C_IDX = 5;
29+constexpr size_t DOUBLE_BUFFER = 2;
30+constexpr size_t BLOCK_BYTE_SIZE = 32;
31+constexpr size_t RANK_KIND = 3;
32+constexpr size_t PRESET_RANK_NUM = 512;
33+constexpr int32_t MAX_RANK_STEP = 32;
34+constexpr int32_t MAX_LOOP_RANK_NUM = 1024;
35+} // namespace
36+ 
37+namespace optiling {
38+template <bool is_grad> static ge::graphStatus TilingForBEVPoolGradV3(gert::TilingContext *context) {
39+ CHECK_NULLPTR(context);
40+ BEVPoolGradV3TilingData tiling;
41+ CHECK_NULLPTR(context->GetPlatformInfo());
42+ auto platform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
43+ uint64_t ubSize;
44+ platform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
45+ auto coreNum = platform.GetCoreNum();
46+ auto featShape = context->GetRequiredInputShape(is_grad ? INPUT_FEAT_GRAD : INPUT_FEAT);
47+ auto ranksBevShape = context->GetRequiredInputShape(is_grad ? INPUT_RANKS_BEV_GRAD : INPUT_RANKS_BEV);
48+ if (featShape == nullptr || ranksBevShape == nullptr) {
49+ return ge::GRAPH_FAILED;
50+ }
51+ auto attrsPtr = context->GetAttrs();
52+ CHECK_NULLPTR(attrsPtr);
53+ auto withDepthPtr = attrsPtr->GetBool(0);
54+ CHECK_NULLPTR(withDepthPtr);
55+ bool withDepth = *withDepthPtr;
56+ context->SetTilingKey(withDepth);
57+ 
58+ auto channel =
59+ featShape->GetOriginShape().GetDim(featShape->GetOriginShape().GetDimNum() - 1); // channel要求32B对齐
60+ uint64_t ranks = ranksBevShape->GetOriginShape().GetDim(0);
61+ 
62+ int32_t fp32ByteSize = sizeof(float); // 当前bevpool只支持fp32类型
63+ int32_t fp32AlignNum = BLOCK_BYTE_SIZE / fp32ByteSize; // 8
64+ 
65+ ubSize = ubSize - RESERVE_UB;
66+ int32_t channelUBSize =
67+ (DOUBLE_BUFFER * channel * 3 + DOUBLE_BUFFER * BLOCK_BYTE_SIZE * 3 + channel * 2 + BLOCK_BYTE_SIZE + 1) *
68+ fp32ByteSize;
69+ 
70+ int32_t rankStep = (ubSize - DOUBLE_BUFFER * RANK_KIND * PRESET_RANK_NUM * fp32ByteSize) / channelUBSize;
Z
Zzhang_qianqian466月3日

针对当前修复的bug，建议后续提供ut进行看护

likedislike
gitdzy
gitdzy
6月5日 评论：
gitdzy
gitdzy
6月5日 评论：
71+ rankStep = FloorAlign(rankStep, fp32AlignNum);
72+ rankStep = std::max(rankStep, fp32AlignNum);
73+ 
74+ int32_t eachLoopRankNum = (ubSize - channelUBSize * rankStep) / (DOUBLE_BUFFER * RANK_KIND * fp32ByteSize);
75+ eachLoopRankNum = FloorAlign(eachLoopRankNum, fp32AlignNum);
76+ eachLoopRankNum = std::max(eachLoopRankNum, fp32AlignNum);
77+ 
78+ int32_t eachCoreRankNum = DivCeil(ranks, static_cast<uint64_t>(coreNum));
79+ eachCoreRankNum = CeilAlign(eachCoreRankNum, fp32AlignNum);
80+ eachLoopRankNum = std::min(eachLoopRankNum, eachCoreRankNum); // 每个iterLoop处理的数量不能超过每个core处理的数量
81+ rankStep = std::min(rankStep, eachLoopRankNum); // 每个innerLoop处理的数量不能超过iterLoop
82+ 
83+ // 32 RankStep性能较优
84+ rankStep = std::min(rankStep, MAX_RANK_STEP);
85+ eachLoopRankNum = std::min(eachLoopRankNum, MAX_LOOP_RANK_NUM);
86+ 
87+ uint64_t avgRankNum = withDepth
88+ ? eachLoopRankNum
89+ : (ubSize - RESERVE_UB) / (sizeof(float) * (channel + 1) * 2) / ONE_BLK_SIZE * ONE_BLK_SIZE;
90+ avgRankNum = std::min(avgRankNum, ranks);
91+ if (avgRankNum == 0) {
92+ return ge::GRAPH_FAILED;
93+ }
94+ 
95+ auto totalTaskNum = (ranks + avgRankNum - 1) / avgRankNum;
96+ uint64_t usedCoreNum = std::min(static_cast<uint64_t>(coreNum), totalTaskNum);
97+ if (usedCoreNum == 0) {
98+ return ge::GRAPH_FAILED;
99+ }
100+ context->SetBlockDim(usedCoreNum);
101+ 
102+ auto avgTaskNum = totalTaskNum / usedCoreNum;
103+ auto tailTaskNum = totalTaskNum % usedCoreNum;
104+ auto tailRankNum = ranks - (totalTaskNum - 1) * avgRankNum;
105+ tiling.set_usedCoreNum(usedCoreNum);
106+ tiling.set_totalTaskNum(totalTaskNum);
107+ tiling.set_avgTaskNum(avgTaskNum);
108+ tiling.set_tailTaskNum(tailTaskNum);
109+ tiling.set_avgRankNum(avgRankNum); // avgRankNum表示每个iterLoop处理的rank数量
110+ tiling.set_tailRankNum(tailRankNum);
111+ tiling.set_channel(channel);
112+ tiling.set_rankStep(rankStep);
113+ MX_DRIVING_LOGI(
114+ "BEVPoolGradV3 tiling: usedCoreNum=%d, totalTaskNum=%d, avgTaskNum=%d, tailTaskNum=%d, avgRankNum=%d, "
115+ "tailRankNum=%d, channel=%d, rankStep=%d",
116+ usedCoreNum, totalTaskNum, avgTaskNum, tailTaskNum, avgRankNum, tailRankNum, channel, rankStep);
117+ 
118+ ADD_TILING_DATA(context, tiling);
119+ 
120+ uint32_t sysWorkspaceSize = platform.GetLibApiWorkSpaceSize();
121+ size_t *currentWorkspace = context->GetWorkspaceSizes(1);
122+ CHECK_NULLPTR(currentWorkspace);
123+ currentWorkspace[0] = sysWorkspaceSize;
124+ return ge::GRAPH_SUCCESS;
125+}
126+} // namespace optiling
127+ 
128+namespace ops {
129+/**
130+ * @brief: BEVPoolGrad, the backward of bev_pool
131+ * @par Inputs:
132+ * grad_out: input grad, 5D tensor(b, d, h, w, c), dtype: float32, format:
133+ * NDHWC, ND geom_feat: input coords, 2D tensor(n, 4), dtype: int32, format: ND
134+ * interval_starts: starting position for pooled point, 1D tensor(n_interval),
135+ * dtype: int32, format: ND interval_lengths: the number of points in each
136+ * interval, 1D tensor(n_interval), dtype: int32, format: ND
137+ * @par Outputs:
138+ * grad_feat: output grad, 2D tensor(n, c), dtype: float32, format: ND
139+ * @par Attributes:
140+ **/
141+class BEVPoolV3Grad : public OpDef {
142+ public:
143+ explicit BEVPoolV3Grad(const char *name) : OpDef(name) {
144+ this->Input("grad_out")
145+ .ParamType(REQUIRED)
146+ .DataType({ge::DT_FLOAT})
147+ .Format({ge::FORMAT_ND})
148+ .AutoContiguous()
149+ .UnknownShapeFormat({ge::FORMAT_ND});
150+ this->Input("depth")
151+ .ParamType(OPTIONAL)
152+ .DataType({ge::DT_FLOAT})
153+ .Format({ge::FORMAT_ND})
154+ .AutoContiguous()
155+ .UnknownShapeFormat({ge::FORMAT_ND});
156+ this->Input("feat")
157+ .ParamType(REQUIRED)
158+ .DataType({ge::DT_FLOAT})
159+ .Format({ge::FORMAT_ND})
160+ .AutoContiguous()
161+ .UnknownShapeFormat({ge::FORMAT_ND});
162+ this->Input("ranks_depth")
163+ .ParamType(OPTIONAL)
164+ .DataType({ge::DT_INT32})
165+ .Format({ge::FORMAT_ND})
166+ .AutoContiguous()
167+ .UnknownShapeFormat({ge::FORMAT_ND});
168+ this->Input("ranks_feat")
169+ .ParamType(OPTIONAL)
170+ .DataType({ge::DT_INT32})
171+ .Format({ge::FORMAT_ND})
172+ .AutoContiguous()
173+ .UnknownShapeFormat({ge::FORMAT_ND});
174+ this->Input("ranks_bev")
175+ .ParamType(REQUIRED)
176+ .DataType({ge::DT_INT32})
177+ .Format({ge::FORMAT_ND})
178+ .AutoContiguous()
179+ .UnknownShapeFormat({ge::FORMAT_ND});
180+ 
181+ this->Attr("with_depth").Bool();
182+ 
183+ this->Output("grad_depth")
184+ .ParamType(OPTIONAL)
185+ .DataType({ge::DT_FLOAT})
186+ .Format({ge::FORMAT_ND})
187+ .UnknownShapeFormat({ge::FORMAT_ND});
188+ this->Output("grad_feat")
189+ .ParamType(REQUIRED)
190+ .DataType({ge::DT_FLOAT})
191+ .Format({ge::FORMAT_ND})
192+ .UnknownShapeFormat({ge::FORMAT_ND});
193+ 
194+ this->AICore().SetTiling(optiling::TilingForBEVPoolGradV3<true>);
195+ this->AICore().AddConfig("ascend910b");
196+ this->AICore().AddConfig("ascend910_93");
197+#if __DRIVING_HOST_AICORE__ == 310
198+ this->AICore().AddConfig("ascend950");
199+#endif
200+ }
201+};
202+ 
203+OP_ADD(BEVPoolV3Grad);
204+} // namespace ops
Akernels/bev_pool_v3/op_host/bev_pool_v3_grad_tiling.h+22-0
@@ -0,0 +1,22 @@
1+/*
2+ * Copyright (c) Huawei Technologies Co., Ltd. 2026. All rights reserved.
3+ */
4+#ifndef BEV_POOL_V3_GRAD_TILING_H
5+#define BEV_POOL_V3_GRAD_TILING_H
6+#include "register/tilingdata_base.h"
7+ 
8+namespace optiling {
9+BEGIN_TILING_DATA_DEF(BEVPoolGradV3TilingData)
10+TILING_DATA_FIELD_DEF(uint64_t, usedCoreNum)
11+TILING_DATA_FIELD_DEF(uint64_t, avgTaskNum)
12+TILING_DATA_FIELD_DEF(uint64_t, tailTaskNum)
13+TILING_DATA_FIELD_DEF(uint64_t, totalTaskNum)
14+TILING_DATA_FIELD_DEF(uint64_t, avgRankNum)
15+TILING_DATA_FIELD_DEF(uint64_t, tailRankNum)
16+TILING_DATA_FIELD_DEF(uint64_t, channel)
17+TILING_DATA_FIELD_DEF(uint64_t, rankStep)
18+END_TILING_DATA_DEF
19+ 
20+REGISTER_TILING_DATA_CLASS(BEVPoolV3Grad, BEVPoolGradV3TilingData)
21+} // namespace optiling
22+#endif // BEV_POOL_V3_GRAD_TILING_H
Mkernels/bev_pool_v3/op_host/bev_pool_v3_tiling.h+1-2
@@ -1,5 +1,5 @@
1/*1/*
2- * Copyright (c) Huawei Technologies Co., Ltd. 2022-2023. All rights reserved.2+ * Copyright (c) Huawei Technologies Co., Ltd. 2026. All rights reserved.
3 */3 */
4#ifndef BEV_POOL_V3_TILING_H4#ifndef BEV_POOL_V3_TILING_H
5#define BEV_POOL_V3_TILING_H5#define BEV_POOL_V3_TILING_H
@@ -17,6 +17,5 @@ TILING_DATA_FIELD_DEF(uint64_t, channel)
17END_TILING_DATA_DEF17END_TILING_DATA_DEF
18 18 
19REGISTER_TILING_DATA_CLASS(BEVPoolV3, BEVPoolV3TilingData)19REGISTER_TILING_DATA_CLASS(BEVPoolV3, BEVPoolV3TilingData)
20-REGISTER_TILING_DATA_CLASS(BEVPoolV3Grad, BEVPoolV3TilingData)
21} // namespace optiling20} // namespace optiling
22#endif // BEV_POOL_V3_TILING_H21#endif // BEV_POOL_V3_TILING_H
Mkernels/bev_pool_v3/op_kernel/bev_pool_v3_grad.cpp+72-93
@@ -1,25 +1,21 @@
1#include "kernel_operator.h"1#include "kernel_operator.h"
2using namespace AscendC;2using namespace AscendC;
3 3 
4- 4+// 00000001 00000001 00000001 00000001
5-static constexpr uint32_t RANK_STEP = 16;
6- // 00000001 00000001 00000001 00000001
7static constexpr uint32_t PATTERN8_0 = 16843009;5static constexpr uint32_t PATTERN8_0 = 16843009;
8// feature\depth\bev6// feature\depth\bev
9static constexpr uint32_t RANK_KIND = 3;7static constexpr uint32_t RANK_KIND = 3;
10static constexpr uint32_t DOUBLE_BUFFER = 2;8static constexpr uint32_t DOUBLE_BUFFER = 2;
11 9 
12-template<bool with_depth>10+template <bool with_depth> class BEVPoolV3GradKernel {
13-class BEVPoolV3GradKernel {11+ public:
14-public:
15 __aicore__ inline BEVPoolV3GradKernel() = delete;12 __aicore__ inline BEVPoolV3GradKernel() = delete;
16 13 
17 __aicore__ inline ~BEVPoolV3GradKernel() = default;14 __aicore__ inline ~BEVPoolV3GradKernel() = default;
18 15 
19- __aicore__ inline BEVPoolV3GradKernel(TPipe* pipe, GM_ADDR gradOut, GM_ADDR depth, GM_ADDR feat, GM_ADDR ranksDepth,16+ __aicore__ inline BEVPoolV3GradKernel(TPipe *pipe, GM_ADDR gradOut, GM_ADDR depth, GM_ADDR feat, GM_ADDR ranksDepth,
20- GM_ADDR ranksFeat, GM_ADDR ranksBev, GM_ADDR gradDepth, GM_ADDR gradFeat, const BEVPoolV3TilingData& tiling)17+ GM_ADDR ranksFeat, GM_ADDR ranksBev, GM_ADDR gradDepth, GM_ADDR gradFeat, const BEVPoolGradV3TilingData &tiling)
21- : pipe_(pipe), blkIdx_(GetBlockIdx()), channel_(tiling.channel)18+ : pipe_(pipe), blkIdx_(GetBlockIdx()), channel_(tiling.channel) {
22- {
23 InitTask(tiling);19 InitTask(tiling);
24 InitOffset();20 InitOffset();
25 InitGM(gradOut, depth, feat, ranksDepth, ranksFeat, ranksBev, gradDepth, gradFeat);21 InitGM(gradOut, depth, feat, ranksDepth, ranksFeat, ranksBev, gradDepth, gradFeat);
@@ -29,9 +25,8 @@ public:
29 __aicore__ inline void Process();25 __aicore__ inline void Process();
30 __aicore__ inline void ProcessWithoutDepth();26 __aicore__ inline void ProcessWithoutDepth();
31 27 
32-private:28+ private:
33- __aicore__ inline void InitTask(const BEVPoolV3TilingData& tiling)29+ __aicore__ inline void InitTask(const BEVPoolGradV3TilingData &tiling) {
34- {
35 int32_t avgTaskNum = tiling.avgTaskNum;30 int32_t avgTaskNum = tiling.avgTaskNum;
36 int32_t tailTaskNum = tiling.tailTaskNum;31 int32_t tailTaskNum = tiling.tailTaskNum;
37 totalTaskNum_ = tiling.totalTaskNum;32 totalTaskNum_ = tiling.totalTaskNum;
@@ -44,10 +39,11 @@ private:
44 taskStartIdx_ = blkIdx_ * avgTaskNum + tailTaskNum;39 taskStartIdx_ = blkIdx_ * avgTaskNum + tailTaskNum;
45 taskEndIdx_ = taskStartIdx_ + avgTaskNum;40 taskEndIdx_ = taskStartIdx_ + avgTaskNum;
46 }41 }
42+ 
43+ RANK_STEP = tiling.rankStep;
47 }44 }
48 45 
49- __aicore__ inline void InitOffset()46+ __aicore__ inline void InitOffset() {
50- {
51 rankSize_ = AlignUp(avgRankNum_, B32_DATA_NUM_PER_BLOCK);47 rankSize_ = AlignUp(avgRankNum_, B32_DATA_NUM_PER_BLOCK);
52 rankBevOffset_ = 0;48 rankBevOffset_ = 0;
53 49 
@@ -58,7 +54,7 @@ private:
58 54 
59 inFeatOffset_ = B32_DATA_NUM_PER_BLOCK;55 inFeatOffset_ = B32_DATA_NUM_PER_BLOCK;
60 inBevOffset_ = inFeatOffset_ + channel_;56 inBevOffset_ = inFeatOffset_ + channel_;
61- 57+ 
62 srcS[0] = RANK_STEP;58 srcS[0] = RANK_STEP;
63 srcS[1] = 1;59 srcS[1] = 1;
64 dstS[0] = RANK_STEP;60 dstS[0] = RANK_STEP;
@@ -76,22 +72,20 @@ private:
76 }72 }
77 73 
78 __aicore__ inline void InitGM(GM_ADDR gradOut, GM_ADDR depth, GM_ADDR feat, GM_ADDR ranksDepth, GM_ADDR ranksFeat,74 __aicore__ inline void InitGM(GM_ADDR gradOut, GM_ADDR depth, GM_ADDR feat, GM_ADDR ranksDepth, GM_ADDR ranksFeat,
79- GM_ADDR ranksBev, GM_ADDR gradDepth, GM_ADDR gradFeat)75+ GM_ADDR ranksBev, GM_ADDR gradDepth, GM_ADDR gradFeat) {
80- {76+ gradOutGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float *>(gradOut));
81- gradOutGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(gradOut));77+ featGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float *>(feat));
82- featGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(feat));78+ ranksBevGm_.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t *>(ranksBev));
83- ranksBevGm_.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(ranksBev));79+ gradFeatGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float *>(gradFeat));
84- gradFeatGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(gradFeat));
85 if (with_depth) {80 if (with_depth) {
86- depthGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(depth));81+ depthGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float *>(depth));
87- ranksDepthGm_.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(ranksDepth));82+ ranksDepthGm_.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t *>(ranksDepth));
88- ranksFeatGm_.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(ranksFeat));83+ ranksFeatGm_.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t *>(ranksFeat));
89- gradDepthGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(gradDepth));84+ gradDepthGm_.SetGlobalBuffer(reinterpret_cast<__gm__ float *>(gradDepth));
90 }85 }
91 }86 }
92 87 
93- __aicore__ inline void InitBuffer()88+ __aicore__ inline void InitBuffer() {
94- {
95 if (with_depth) {89 if (with_depth) {
96 pipe_->InitBuffer(ranksBuf_, DOUBLE_BUFFER * rankBatchSize_ * sizeof(int32_t));90 pipe_->InitBuffer(ranksBuf_, DOUBLE_BUFFER * rankBatchSize_ * sizeof(int32_t));
97 pipe_->InitBuffer(featBuf_, DOUBLE_BUFFER * batchChannel_ * sizeof(float));91 pipe_->InitBuffer(featBuf_, DOUBLE_BUFFER * batchChannel_ * sizeof(float));
@@ -101,10 +95,11 @@ private:
101 95 
102 pipe_->InitBuffer(depthBuf_, batchChannel_ * sizeof(float));96 pipe_->InitBuffer(depthBuf_, batchChannel_ * sizeof(float));
103 pipe_->InitBuffer(depthGatherTmpBuf_, RANK_STEP * B32_DATA_NUM_PER_BLOCK * sizeof(float));97 pipe_->InitBuffer(depthGatherTmpBuf_, RANK_STEP * B32_DATA_NUM_PER_BLOCK * sizeof(float));
104- pipe_->InitBuffer(patternBuf_, RANK_STEP * B32_DATA_NUM_PER_BLOCK * sizeof(uint32_t));98+ pipe_->InitBuffer(patternBuf_, RANK_STEP * B32_DATA_NUM_PER_BLOCK * sizeof(uint32_t));
105- 99+ 
106- pipe_->InitBuffer(gradTempDepthBuf_, batchChannel_* sizeof(float));100+ pipe_->InitBuffer(gradTempDepthBuf_, batchChannel_ * sizeof(float));
107- pipe_->InitBuffer(gradDepthBuf_, (RANK_STEP + DOUBLE_BUFFER * RANK_STEP * B32_DATA_NUM_PER_BLOCK) * sizeof(float));101+ pipe_->InitBuffer(
102+ gradDepthBuf_, (RANK_STEP + DOUBLE_BUFFER * RANK_STEP * B32_DATA_NUM_PER_BLOCK) * sizeof(float));
108 103 
109 ranksBev_ = ranksBuf_.Get<int32_t>();104 ranksBev_ = ranksBuf_.Get<int32_t>();
110 ranksFeat_ = ranksBev_[rankSize_];105 ranksFeat_ = ranksBev_[rankSize_];
@@ -141,13 +136,14 @@ private:
141 __aicore__ inline void CopyInStage1(uint8_t ping, uint32_t step, uint64_t off, uint64_t depOff);136 __aicore__ inline void CopyInStage1(uint8_t ping, uint32_t step, uint64_t off, uint64_t depOff);
142 __aicore__ inline void ComputeStage1(uint8_t ping, uint32_t step, uint64_t featOff, uint64_t depOff);137 __aicore__ inline void ComputeStage1(uint8_t ping, uint32_t step, uint64_t featOff, uint64_t depOff);
143 __aicore__ inline void CopyOutStage1(uint8_t ping, uint32_t step, uint64_t off, uint64_t featOff);138 __aicore__ inline void CopyOutStage1(uint8_t ping, uint32_t step, uint64_t off, uint64_t featOff);
144- __aicore__ inline void ProcessDualRank(int32_t rankOffset, uint32_t step0, uint32_t step1, uint64_t featOff1, uint64_t depOff1, int32_t idx);139+ __aicore__ inline void ProcessDualRank(
140+ int32_t rankOffset, uint32_t step0, uint32_t step1, uint64_t featOff1, uint64_t depOff1, int32_t idx);
145 __aicore__ inline void ProcessSingle(uint8_t ping, uint64_t taskIdx, uint32_t rankNum);141 __aicore__ inline void ProcessSingle(uint8_t ping, uint64_t taskIdx, uint32_t rankNum);
146 142 
147 __aicore__ inline void ProcessSingleWithoutDepth(uint64_t taskIdx, uint32_t rankNum);143 __aicore__ inline void ProcessSingleWithoutDepth(uint64_t taskIdx, uint32_t rankNum);
148 144 
149-private:145+ private:
150- TPipe* pipe_;146+ TPipe *pipe_;
151 int32_t blkIdx_;147 int32_t blkIdx_;
152 GlobalTensor<float> gradOutGm_, depthGm_, featGm_, gradDepthGm_, gradFeatGm_;148 GlobalTensor<float> gradOutGm_, depthGm_, featGm_, gradDepthGm_, gradFeatGm_;
153 GlobalTensor<int32_t> ranksDepthGm_, ranksFeatGm_, ranksBevGm_;149 GlobalTensor<int32_t> ranksDepthGm_, ranksFeatGm_, ranksBevGm_;
@@ -164,30 +160,21 @@ private:
164 int32_t channel_, batchChannel_;160 int32_t channel_, batchChannel_;
165 uint32_t rankSize_, avgRankNum_, tailRankNum_, rankBatchSize_;161 uint32_t rankSize_, avgRankNum_, tailRankNum_, rankBatchSize_;
166 uint64_t rankDepthOffset_, rankFeatOffset_, rankBevOffset_, inFeatOffset_, inBevOffset_;162 uint64_t rankDepthOffset_, rankFeatOffset_, rankBevOffset_, inFeatOffset_, inBevOffset_;
167- 163+ 
164+ uint32_t RANK_STEP;
168 uint32_t srcS[2], dstS[2];165 uint32_t srcS[2], dstS[2];
169 uint32_t srcSDepth[2], dstSDepth[2];166 uint32_t srcSDepth[2], dstSDepth[2];
170 SumParams sumParams;167 SumParams sumParams;
171- DataCopyParams cpSingleParams_ {1, B32_BYTE_SIZE, 0, 0};168+ DataCopyParams cpSingleParams_{1, B32_BYTE_SIZE, 0, 0};
172 TEventID cpInEvtID_, cpOutEvtID_;169 TEventID cpInEvtID_, cpOutEvtID_;
173 170 
174- TBuf<TPosition::VECCALC> 171+ TBuf<TPosition::VECCALC> ranksBuf_, featBuf_, depthBuf_, depthGatherTmpBuf_, depthCopyTmpBuf_, patternBuf_, outBuf_,
175- ranksBuf_, 172+ gradOutBuf_, gradFeatBuf_, gradTempDepthBuf_, gradDepthBuf_;
176- featBuf_,
177- depthBuf_,
178- depthGatherTmpBuf_,
179- depthCopyTmpBuf_,
180- patternBuf_,
181- outBuf_,
182- gradOutBuf_,
183- gradFeatBuf_,
184- gradTempDepthBuf_,
185- gradDepthBuf_;
186};173};
187 174 
188-template<bool with_depth>175+template <bool with_depth>
189-__aicore__ inline void BEVPoolV3GradKernel<with_depth>::CopyInStage0(uint8_t ping, uint32_t step, uint64_t off, uint64_t featOff)176+__aicore__ inline void BEVPoolV3GradKernel<with_depth>::CopyInStage0(
190-{177+ uint8_t ping, uint32_t step, uint64_t off, uint64_t featOff) {
191 WaitFlag<HardEvent::V_MTE2>(ping + 4);178 WaitFlag<HardEvent::V_MTE2>(ping + 4);
192 for (int32_t j = 0; j < step; j++) {179 for (int32_t j = 0; j < step; j++) {
193 uint64_t rf = ranksFeat_.GetValue(off + j);180 uint64_t rf = ranksFeat_.GetValue(off + j);
@@ -198,9 +185,8 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::CopyInStage0(uint8_t pin
198 SetFlag<HardEvent::MTE2_V>(ping + 4);185 SetFlag<HardEvent::MTE2_V>(ping + 4);
199}186}
200 187 
201-template<bool with_depth>188+template <bool with_depth>
202-__aicore__ inline void BEVPoolV3GradKernel<with_depth>::ComputeStage0(uint8_t ping, uint32_t step, uint64_t featOff)189+__aicore__ inline void BEVPoolV3GradKernel<with_depth>::ComputeStage0(uint8_t ping, uint32_t step, uint64_t featOff) {
203-{
204 LocalTensor<float> broadTmp = gradDepthBroad_[ping * RANK_STEP * B32_DATA_NUM_PER_BLOCK];190 LocalTensor<float> broadTmp = gradDepthBroad_[ping * RANK_STEP * B32_DATA_NUM_PER_BLOCK];
205 WaitFlag<HardEvent::MTE2_V>(ping + 4);191 WaitFlag<HardEvent::MTE2_V>(ping + 4);
206 Mul(gradTempDepthLocal_, gradOutLocal_[featOff], featLocal_[featOff], channel_ * step);192 Mul(gradTempDepthLocal_, gradOutLocal_[featOff], featLocal_[featOff], channel_ * step);
@@ -211,9 +197,8 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::ComputeStage0(uint8_t pi
211 SetFlag<HardEvent::V_MTE3>(ping);197 SetFlag<HardEvent::V_MTE3>(ping);
212}198}
213 199 
214-template<bool with_depth>200+template <bool with_depth>
215-__aicore__ inline void BEVPoolV3GradKernel<with_depth>::CopyOutStage0(uint8_t ping, uint32_t step, uint64_t off)201+__aicore__ inline void BEVPoolV3GradKernel<with_depth>::CopyOutStage0(uint8_t ping, uint32_t step, uint64_t off) {
216-{
217 LocalTensor<float> broadTmp = gradDepthBroad_[ping * RANK_STEP * B32_DATA_NUM_PER_BLOCK];202 LocalTensor<float> broadTmp = gradDepthBroad_[ping * RANK_STEP * B32_DATA_NUM_PER_BLOCK];
218 WaitFlag<HardEvent::V_MTE3>(ping);203 WaitFlag<HardEvent::V_MTE3>(ping);
219 SetAtomicAdd<float>();204 SetAtomicAdd<float>();
@@ -225,10 +210,9 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::CopyOutStage0(uint8_t pi
225 SetFlag<HardEvent::MTE3_V>(ping);210 SetFlag<HardEvent::MTE3_V>(ping);
226}211}
227 212 
228- 213+template <bool with_depth>
229-template<bool with_depth>214+__aicore__ inline void BEVPoolV3GradKernel<with_depth>::CopyInStage1(
230-__aicore__ inline void BEVPoolV3GradKernel<with_depth>::CopyInStage1(uint8_t ping, uint32_t step, uint64_t off, uint64_t depOff)215+ uint8_t ping, uint32_t step, uint64_t off, uint64_t depOff) {
231-{
232 WaitFlag<HardEvent::V_MTE2>(ping + 6);216 WaitFlag<HardEvent::V_MTE2>(ping + 6);
233 for (int32_t j = 0; j < step; j++) {217 for (int32_t j = 0; j < step; j++) {
234 uint64_t rd = ranksDepth_.GetValue(off + j);218 uint64_t rd = ranksDepth_.GetValue(off + j);
@@ -237,13 +221,14 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::CopyInStage1(uint8_t pin
237 SetFlag<HardEvent::MTE2_V>(ping + 6);221 SetFlag<HardEvent::MTE2_V>(ping + 6);
238}222}
239 223 
240-template<bool with_depth>224+template <bool with_depth>
241-__aicore__ inline void BEVPoolV3GradKernel<with_depth>::ComputeStage1(uint8_t ping, uint32_t step, uint64_t featOff, uint64_t depOff)225+__aicore__ inline void BEVPoolV3GradKernel<with_depth>::ComputeStage1(
242-{226+ uint8_t ping, uint32_t step, uint64_t featOff, uint64_t depOff) {
243 WaitFlag<HardEvent::MTE2_V>(ping + 6);227 WaitFlag<HardEvent::MTE2_V>(ping + 6);
244 228 
245 uint64_t rsvdCnt = 0;229 uint64_t rsvdCnt = 0;
246- GatherMask(depthGather_, depthTmp_[depOff], patternLocal_, false, 0, { 1, RANK_STEP / 8, 8, 0 }, rsvdCnt);230+ GatherMask(depthGather_, depthTmp_[depOff], patternLocal_, false, 0,
231+ {1, static_cast<uint16_t>(RANK_STEP / 8), 8, 0}, rsvdCnt);
247 BroadCast<float, 2, 1>(depthLocal_, depthGather_, dstS, srcS);232 BroadCast<float, 2, 1>(depthLocal_, depthGather_, dstS, srcS);
248 233 
249 WaitFlag<HardEvent::MTE3_V>(ping + 2);234 WaitFlag<HardEvent::MTE3_V>(ping + 2);
@@ -254,9 +239,9 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::ComputeStage1(uint8_t pi
254 SetFlag<HardEvent::V_MTE2>(ping + 6);239 SetFlag<HardEvent::V_MTE2>(ping + 6);
255}240}
256 241 
257-template<bool with_depth>242+template <bool with_depth>
258-__aicore__ inline void BEVPoolV3GradKernel<with_depth>::CopyOutStage1(uint8_t ping, uint32_t step, uint64_t off, uint64_t featOff)243+__aicore__ inline void BEVPoolV3GradKernel<with_depth>::CopyOutStage1(
259-{244+ uint8_t ping, uint32_t step, uint64_t off, uint64_t featOff) {
260 WaitFlag<HardEvent::V_MTE3>(ping + 2);245 WaitFlag<HardEvent::V_MTE3>(ping + 2);
261 SetAtomicAdd<float>();246 SetAtomicAdd<float>();
262 for (int32_t j = 0; j < step; j++) {247 for (int32_t j = 0; j < step; j++) {
@@ -267,14 +252,14 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::CopyOutStage1(uint8_t pi
267 SetFlag<HardEvent::MTE3_V>(ping + 2);252 SetFlag<HardEvent::MTE3_V>(ping + 2);
268}253}
269 254 
270-template<bool with_depth>255+template <bool with_depth>
271-__aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessDualRank(int32_t rankOffset, uint32_t step0, uint32_t step1, uint64_t featOff1, uint64_t depOff1, int32_t idx)256+__aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessDualRank(
272-{257+ int32_t rankOffset, uint32_t step0, uint32_t step1, uint64_t featOff1, uint64_t depOff1, int32_t idx) {
273 uint64_t off0 = rankOffset + idx * RANK_STEP;258 uint64_t off0 = rankOffset + idx * RANK_STEP;
274 uint64_t off1 = off0 + RANK_STEP;259 uint64_t off1 = off0 + RANK_STEP;
275 CopyInStage0(0, step0, off0, 0);260 CopyInStage0(0, step0, off0, 0);
276 CopyInStage0(1, step1, off1, featOff1);261 CopyInStage0(1, step1, off1, featOff1);
277- 262+ 
278 ComputeStage0(0, step0, 0);263 ComputeStage0(0, step0, 0);
279 CopyInStage1(0, step0, off0, 0);264 CopyInStage1(0, step0, off0, 0);
280 CopyOutStage0(0, step0, off0);265 CopyOutStage0(0, step0, off0);
@@ -288,20 +273,20 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessDualRank(int32_t
288 CopyOutStage1(1, step1, off1, featOff1);273 CopyOutStage1(1, step1, off1, featOff1);
289}274}
290 275 
291-template<bool with_depth>276+template <bool with_depth>
292-__aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessSingle(uint8_t ping, uint64_t taskIdx, uint32_t actualRankNum)277+__aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessSingle(
293-{278+ uint8_t ping, uint64_t taskIdx, uint32_t actualRankNum) {
294 int32_t rankNum = AlignUp(actualRankNum, B32_DATA_NUM_PER_BLOCK);279 int32_t rankNum = AlignUp(actualRankNum, B32_DATA_NUM_PER_BLOCK);
295 int32_t rankOffset = ping * rankBatchSize_;280 int32_t rankOffset = ping * rankBatchSize_;
296 WaitFlag<HardEvent::V_MTE2>(ping);281 WaitFlag<HardEvent::V_MTE2>(ping);
297 DataCopy(ranksBev_[rankOffset], ranksBevGm_[taskIdx * avgRankNum_], rankNum);282 DataCopy(ranksBev_[rankOffset], ranksBevGm_[taskIdx * avgRankNum_], rankNum);
298 DataCopy(ranksFeat_[rankOffset], ranksFeatGm_[taskIdx * avgRankNum_], rankNum);283 DataCopy(ranksFeat_[rankOffset], ranksFeatGm_[taskIdx * avgRankNum_], rankNum);
299 SetFlag<HardEvent::MTE2_V>(ping);284 SetFlag<HardEvent::MTE2_V>(ping);
300- 285+ 
301 WaitFlag<HardEvent::MTE2_V>(ping);286 WaitFlag<HardEvent::MTE2_V>(ping);
302 // 2 * rankSize_ -> ranksBev_、ranksFeat_287 // 2 * rankSize_ -> ranksBev_、ranksFeat_
303 Muls(ranksBev_[rankOffset], ranksBev_[rankOffset], channel_, 2 * rankSize_);288 Muls(ranksBev_[rankOffset], ranksBev_[rankOffset], channel_, 2 * rankSize_);
304- 289+ 
305 WaitFlag<HardEvent::V_MTE2>(ping + 2);290 WaitFlag<HardEvent::V_MTE2>(ping + 2);
306 DataCopy(ranksDepth_[rankOffset], ranksDepthGm_[taskIdx * avgRankNum_], rankNum);291 DataCopy(ranksDepth_[rankOffset], ranksDepthGm_[taskIdx * avgRankNum_], rankNum);
307 292 
@@ -311,7 +296,7 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessSingle(uint8_t pi
311 uint32_t step1 = RANK_STEP;296 uint32_t step1 = RANK_STEP;
312 uint64_t featOff1 = batchChannel_;297 uint64_t featOff1 = batchChannel_;
313 uint64_t depOff1 = RANK_STEP * B32_DATA_NUM_PER_BLOCK;298 uint64_t depOff1 = RANK_STEP * B32_DATA_NUM_PER_BLOCK;
314- if ((cnt % 2) == 0) {299+ if ((cnt % 2) == 0) {
315 for (int32_t i = 0; i < cnt; i += 2) {300 for (int32_t i = 0; i < cnt; i += 2) {
316 if (unlikely(i == cnt - 2)) {301 if (unlikely(i == cnt - 2)) {
317 step1 = tail;302 step1 = tail;
@@ -324,7 +309,7 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessSingle(uint8_t pi
324 }309 }
325 {310 {
326 step0 = tail;311 step0 = tail;
327- int32_t i = cnt - 1;312+ int32_t i = cnt - 1;
328 uint64_t off0 = rankOffset + i * RANK_STEP;313 uint64_t off0 = rankOffset + i * RANK_STEP;
329 CopyInStage0(0, step0, off0, 0);314 CopyInStage0(0, step0, off0, 0);
330 ComputeStage0(0, step0, 0);315 ComputeStage0(0, step0, 0);
@@ -339,9 +324,7 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessSingle(uint8_t pi
339 SetFlag<HardEvent::V_MTE2>(ping + 2);324 SetFlag<HardEvent::V_MTE2>(ping + 2);
340}325}
341 326 
342-template<bool with_depth>327+template <bool with_depth> __aicore__ inline void BEVPoolV3GradKernel<with_depth>::Process() {
343-__aicore__ inline void BEVPoolV3GradKernel<with_depth>::Process()
344-{
345 uint8_t ping = 0;328 uint8_t ping = 0;
346 SetFlag<HardEvent::V_MTE2>(0);329 SetFlag<HardEvent::V_MTE2>(0);
347 SetFlag<HardEvent::V_MTE2>(1);330 SetFlag<HardEvent::V_MTE2>(1);
@@ -357,7 +340,6 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::Process()
357 SetFlag<HardEvent::MTE3_V>(2);340 SetFlag<HardEvent::MTE3_V>(2);
358 SetFlag<HardEvent::MTE3_V>(3);341 SetFlag<HardEvent::MTE3_V>(3);
359 342 
360- 
361 for (uint32_t i = taskStartIdx_; i < taskEndIdx_; ++i) {343 for (uint32_t i = taskStartIdx_; i < taskEndIdx_; ++i) {
362 uint32_t actualRankNum = avgRankNum_;344 uint32_t actualRankNum = avgRankNum_;
363 if (unlikely(i == totalTaskNum_ - 1)) {345 if (unlikely(i == totalTaskNum_ - 1)) {
@@ -382,9 +364,9 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::Process()
382 WaitFlag<HardEvent::MTE3_V>(3);364 WaitFlag<HardEvent::MTE3_V>(3);
383}365}
384 366 
385-template<bool with_depth>367+template <bool with_depth>
386-__aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessSingleWithoutDepth(uint64_t taskIdx, uint32_t actualRankNum)368+__aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessSingleWithoutDepth(
387-{369+ uint64_t taskIdx, uint32_t actualRankNum) {
388 int32_t rankNum = AlignUp(actualRankNum, B32_DATA_NUM_PER_BLOCK);370 int32_t rankNum = AlignUp(actualRankNum, B32_DATA_NUM_PER_BLOCK);
389 LocalTensor<int32_t> ranks = ranksQue_.AllocTensor<int32_t>();371 LocalTensor<int32_t> ranks = ranksQue_.AllocTensor<int32_t>();
390 LocalTensor<int32_t> rankBev = ranks[rankBevOffset_];372 LocalTensor<int32_t> rankBev = ranks[rankBevOffset_];
@@ -407,9 +389,7 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessSingleWithoutDept
407 ranksQue_.FreeTensor(ranks);389 ranksQue_.FreeTensor(ranks);
408}390}
409 391 
410-template<bool with_depth>392+template <bool with_depth> __aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessWithoutDepth() {
411-__aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessWithoutDepth()
412-{
413 for (uint32_t i = taskStartIdx_; i < taskEndIdx_; ++i) {393 for (uint32_t i = taskStartIdx_; i < taskEndIdx_; ++i) {
414 uint32_t actualRankNum = avgRankNum_;394 uint32_t actualRankNum = avgRankNum_;
415 if (unlikely(i == totalTaskNum_ - 1)) {395 if (unlikely(i == totalTaskNum_ - 1)) {
@@ -420,8 +400,7 @@ __aicore__ inline void BEVPoolV3GradKernel<with_depth>::ProcessWithoutDepth()
420}400}
421 401 
422extern "C" __global__ __aicore__ void bev_pool_v3_grad(GM_ADDR gradOut, GM_ADDR depth, GM_ADDR feat, GM_ADDR ranksDepth,402extern "C" __global__ __aicore__ void bev_pool_v3_grad(GM_ADDR gradOut, GM_ADDR depth, GM_ADDR feat, GM_ADDR ranksDepth,
423- GM_ADDR ranksFeat, GM_ADDR ranksBev, GM_ADDR gradDepth, GM_ADDR gradFeat, GM_ADDR workspace, GM_ADDR tiling)403+ GM_ADDR ranksFeat, GM_ADDR ranksBev, GM_ADDR gradDepth, GM_ADDR gradFeat, GM_ADDR workspace, GM_ADDR tiling) {
424-{
425 GET_TILING_DATA(bevPoolTiling, tiling);404 GET_TILING_DATA(bevPoolTiling, tiling);
426 TPipe pipe;405 TPipe pipe;
427 if (TILING_KEY_IS(0)) {406 if (TILING_KEY_IS(0)) {
@@ -433,4 +412,4 @@ extern "C" __global__ __aicore__ void bev_pool_v3_grad(GM_ADDR gradOut, GM_ADDR
433 &pipe, gradOut, depth, feat, ranksDepth, ranksFeat, ranksBev, gradDepth, gradFeat, bevPoolTiling);412 &pipe, gradOut, depth, feat, ranksDepth, ranksFeat, ranksBev, gradDepth, gradFeat, bevPoolTiling);
434 kernel.Process();413 kernel.Process();
435 }414 }
436-}415+}