已合并
修复MoeInitRoutingQuantV2第三个输出结果精度问题 #8354
wuxiyuan创建于 7月7日
修复MoeInitRoutingQuantV2第三个输出结果精度问题 #8354
已合并
共 1 个文件变更+7-18
| @@ -38,9 +38,8 @@ private: | |||
| 38 | 38 | ||
| 39 | private: | 39 | private: |
| 40 | TPipe* pipe; | 40 | TPipe* pipe; |
| 41 | - AscendC::TBuf<AscendC::TPosition::VECCALC> expertIdCunsumTBuf_; | ||
| 42 | - | ||
| 43 | GlobalTensor<int32_t> expandedRowIdxGmGT_; | 41 | GlobalTensor<int32_t> expandedRowIdxGmGT_; |
| 42 | + GlobalTensor<int32_t> expandedExpertIdxGmGT_; | ||
| 44 | __gm__ int32_t* expandedRowIdxGm_; | 43 | __gm__ int32_t* expandedRowIdxGm_; |
| 45 | __gm__ int32_t* expandedExpertIdxGm_; | 44 | __gm__ int32_t* expandedExpertIdxGm_; |
| 46 | __gm__ int32_t* expertTokensBeforeCapacityGm_; | 45 | __gm__ int32_t* expertTokensBeforeCapacityGm_; |
| @@ -48,9 +47,6 @@ private: | |||
| 48 | GlobalTensor<int32_t> expertTokensCountOrCumsumGmOut; | 47 | GlobalTensor<int32_t> expertTokensCountOrCumsumGmOut; |
| 49 | __gm__ int32_t* expertFirstIndexGm_; | 48 | __gm__ int32_t* expertFirstIndexGm_; |
| 50 | 49 | ||
| 51 | - LocalTensor<int32_t> lastExpertIdCunsum_; | ||
| 52 | - __ubuf__ int32_t* lastExpertIdCunsumAddr_; | ||
| 53 | - | ||
| 54 | const InnerMoeV2GatherOutComputeTilingData* srcToDstTilingData; | 50 | const InnerMoeV2GatherOutComputeTilingData* srcToDstTilingData; |
| 55 | 51 | ||
| 56 | int64_t coreNum_; | 52 | int64_t coreNum_; |
| @@ -142,9 +138,7 @@ __aicore__ inline void MoeV2ExpertTokenOutSimt::Init( | |||
| 142 | 138 | ||
| 143 | expertFirstIndexGm_ = (__gm__ int32_t*)workspace + Align(this->totalLength_, sizeof(int32_t)) * 2; | 139 | expertFirstIndexGm_ = (__gm__ int32_t*)workspace + Align(this->totalLength_, sizeof(int32_t)) * 2; |
| 144 | expandedExpertIdxGm_ = (__gm__ int32_t*)workspace; | 140 | expandedExpertIdxGm_ = (__gm__ int32_t*)workspace; |
| 145 | - pipe->InitBuffer(expertIdCunsumTBuf_, BLOCK_BYTES); | 141 | + expandedExpertIdxGmGT_.SetGlobalBuffer((__gm__ int32_t*)workspace, Align(this->totalLength_, sizeof(int32_t))); |
| 146 | - lastExpertIdCunsum_ = expertIdCunsumTBuf_.Get<int32_t>(); | ||
| 147 | - lastExpertIdCunsumAddr_ = (__ubuf__ int32_t*)lastExpertIdCunsum_.GetPhyAddr(); | ||
| 148 | } | 142 | } |
| 149 | 143 | ||
| 150 | __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void ExpertFirstIndexComputeSimt( | 144 | __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void ExpertFirstIndexComputeSimt( |
| @@ -191,17 +185,13 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void TokensComputeSimt( | |||
| 191 | 185 | ||
| 192 | __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void CumsumComputeSimt( | 186 | __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void CumsumComputeSimt( |
| 193 | int64_t coreRows, int64_t startIndex, int64_t totalLength, __gm__ int32_t* expandedExpertIdxGm, | 187 | int64_t coreRows, int64_t startIndex, int64_t totalLength, __gm__ int32_t* expandedExpertIdxGm, |
| 194 | - __gm__ int32_t* expertTokensCountOrCumsumGm, __ubuf__ int32_t* lastExpertIdCunsumAddr) | 188 | + __gm__ int32_t* expertTokensCountOrCumsumGm) |
| 195 | { | 189 | { |
| 196 | for (int32_t index = static_cast<int32_t>(threadIdx.x); index < static_cast<int32_t>(coreRows); | 190 | for (int32_t index = static_cast<int32_t>(threadIdx.x); index < static_cast<int32_t>(coreRows); |
| 197 | index += static_cast<int32_t>(blockDim.x)) { | 191 | index += static_cast<int32_t>(blockDim.x)) { |
| 198 | int64_t curIndex = startIndex + index; | 192 | int64_t curIndex = startIndex + index; |
| 199 | int32_t curExpertId = expandedExpertIdxGm[curIndex]; | 193 | int32_t curExpertId = expandedExpertIdxGm[curIndex]; |
| 200 | if (curIndex == totalLength - 1) { // 最后一个数 | 194 | if (curIndex == totalLength - 1) { // 最后一个数 |
| 201 | - if (threadIdx.y == 0) { | ||
| 202 | - lastExpertIdCunsumAddr[0] = curExpertId; | ||
| 203 | - lastExpertIdCunsumAddr[1] = curIndex + 1; | ||
| 204 | - } | ||
| 205 | return; | 195 | return; |
| 206 | } | 196 | } |
| 207 | int32_t nextExpertId = expandedExpertIdxGm[curIndex + 1]; | 197 | int32_t nextExpertId = expandedExpertIdxGm[curIndex + 1]; |
| @@ -245,13 +235,12 @@ __aicore__ inline void MoeV2ExpertTokenOutSimt::Process() | |||
| 245 | THREAD_NUM / this->threadNum_ > MAX_THREAD_DIM_Y_NUM ? MAX_THREAD_DIM_Y_NUM : THREAD_NUM / this->threadNum_; | 235 | THREAD_NUM / this->threadNum_ > MAX_THREAD_DIM_Y_NUM ? MAX_THREAD_DIM_Y_NUM : THREAD_NUM / this->threadNum_; |
| 246 | asc_vf_call<CumsumComputeSimt>( | 236 | asc_vf_call<CumsumComputeSimt>( |
| 247 | dim3{static_cast<uint32_t>(this->threadNum_), static_cast<uint32_t>(threadDimY), 1}, this->coreRows_, | 237 | dim3{static_cast<uint32_t>(this->threadNum_), static_cast<uint32_t>(threadDimY), 1}, this->coreRows_, |
| 248 | - this->startIndex_, this->totalLength_, expandedExpertIdxGm_, expertTokensCountOrCumsumGm_, | 238 | + this->startIndex_, this->totalLength_, expandedExpertIdxGm_, expertTokensCountOrCumsumGm_); |
| 249 | - lastExpertIdCunsumAddr_); | ||
| 250 | } | 239 | } |
| 251 | 240 | ||
| 252 | if (this->expertCumsum_ && this->blockIdx_ == this->needCoreNum_ - 1) { | 241 | if (this->expertCumsum_ && this->blockIdx_ == this->needCoreNum_ - 1) { |
| 253 | - int32_t expertId = lastExpertIdCunsumAddr_[0]; | 242 | + int32_t expertId = expandedExpertIdxGmGT_.GetValue(this->totalLength_ - 1); |
| 254 | - int32_t cunsumNum = lastExpertIdCunsumAddr_[1]; | 243 | + int32_t cunsumNum = static_cast<int32_t>(this->totalLength_); |
| 255 | int32_t count = expertNum_ - expertId; | 244 | int32_t count = expertNum_ - expertId; |
| 256 | if (count > 0) { | 245 | if (count > 0) { |
| 257 | InitOutput(expertTokensCountOrCumsumGmOut[expertId], static_cast<uint32_t>(count), cunsumNum); | 246 | InitOutput(expertTokensCountOrCumsumGmOut[expertId], static_cast<uint32_t>(count), cunsumNum); |
| @@ -259,4 +248,4 @@ __aicore__ inline void MoeV2ExpertTokenOutSimt::Process() | |||
| 259 | } | 248 | } |
| 260 | } | 249 | } |
| 261 | } // namespace MoeInitRoutingQuantV2 | 250 | } // namespace MoeInitRoutingQuantV2 |
| 262 | -#endif // MOE_V2_QUANT_EXPERT_TOKEN_OUT_SIMT_H | 251 | +#endif // MOE_V2_QUANT_EXPERT_TOKEN_OUT_SIMT_H |