Star906
1825
代码
代码
Issues226
Pull Requests303
Actions
项目讨论
Wiki
分析
Star906
1825
已合并
simt算子风格改为c风格 #4243
chenfeng创建于 4月25日
simt算子风格改为c风格 #4243
已合并
chenfeng创建于 4月25日
从chenfeng61/ops-nnsimt_zg合入到cann/ops-nnmaster
共 11 个文件变更+221-208
Mconv/deformable_offsets/op_kernel/arch35/deformable_offsets.h+4-3
@@ -15,6 +15,7 @@
15#ifndef DEFORMABLE_OFFSET_H15#ifndef DEFORMABLE_OFFSET_H
16#define DEFORMABLE_OFFSET_H16#define DEFORMABLE_OFFSET_H
17#include "kernel_operator.h"17#include "kernel_operator.h"
18+#include "simt_api/asc_simt.h"
18namespace DeformableOffsets {19namespace DeformableOffsets {
19using namespace AscendC;20using namespace AscendC;
20const int32_t WIDTH_OFFSET_INDEX = 0;21const int32_t WIDTH_OFFSET_INDEX = 0;
@@ -127,7 +128,7 @@ __simt_vf__ LAUNCH_BOUND(VF_MAX_THREAD_NUM) __aicore__ void ComputeDeformableOff
127 T1 widthOffset = WIDTH_OFFSET_INDEX * offsetKernelElementStride;128 T1 widthOffset = WIDTH_OFFSET_INDEX * offsetKernelElementStride;
128 T1 weightOffset = POINT_WEIGHT_OFFSET_INDEX * offsetKernelElementStride;129 T1 weightOffset = POINT_WEIGHT_OFFSET_INDEX * offsetKernelElementStride;
129 130 
130- for (T1 index = blockId_ * VF_MAX_THREAD_NUM + Simt::GetThreadIdx(); index < numKernels;131+ for (T1 index = blockId_ * VF_MAX_THREAD_NUM + threadIdx.x; index < numKernels;
131 index += (blockNumber * VF_MAX_THREAD_NUM)) {132 index += (blockNumber * VF_MAX_THREAD_NUM)) {
132 // output info (N H K_h W K_w, groups, groupC)133 // output info (N H K_h W K_w, groups, groupC)
133 T1 batchNum, heightCol, widthCol, channelIndex, groupsIndex;134 T1 batchNum, heightCol, widthCol, channelIndex, groupsIndex;
@@ -185,8 +186,8 @@ __aicore__ inline void DeformableOffset<T, T1, T2>::Process()
185 GetUintDivMagicAndShift(mH_, shiftH_, static_cast<T2>(tiling_->imgOutWidth * tiling_->imgChannel));186 GetUintDivMagicAndShift(mH_, shiftH_, static_cast<T2>(tiling_->imgOutWidth * tiling_->imgChannel));
186 GetUintDivMagicAndShift(mW_, shiftW_, static_cast<T2>(tiling_->imgChannel));187 GetUintDivMagicAndShift(mW_, shiftW_, static_cast<T2>(tiling_->imgChannel));
187 GetUintDivMagicAndShift(mC_, shiftC_, static_cast<T2>(tiling_->imgChannel / tiling_->deformableGroups));188 GetUintDivMagicAndShift(mC_, shiftC_, static_cast<T2>(tiling_->imgChannel / tiling_->deformableGroups));
188- Simt::VF_CALL<ComputeDeformableOffset<T, T1, T2>>(189+ asc_vf_call<ComputeDeformableOffset<T, T1, T2>>(
189- Simt::Dim3{VF_MAX_THREAD_NUM, 1, 1}, (__gm__ T*)(inputImgGm_.GetPhyAddr()),190+ dim3{VF_MAX_THREAD_NUM, 1, 1}, (__gm__ T*)(inputImgGm_.GetPhyAddr()),
190 (__gm__ T*)(offsetsGm_.GetPhyAddr()), (__gm__ T*)(yGm_.GetPhyAddr()), tiling_->blockNum, tiling_->numKernels,191 (__gm__ T*)(offsetsGm_.GetPhyAddr()), (__gm__ T*)(yGm_.GetPhyAddr()), tiling_->blockNum, tiling_->numKernels,
191 tiling_->imgOutWidth, tiling_->imgChannel, tiling_->imgHeight, tiling_->imgWidth, tiling_->strideHeight,192 tiling_->imgOutWidth, tiling_->imgChannel, tiling_->imgHeight, tiling_->imgWidth, tiling_->strideHeight,
192 tiling_->strideWidth, tiling_->dilationHeight, tiling_->dilationWidth, tiling_->padsHeight, tiling_->padsWidth,193 tiling_->strideWidth, tiling_->dilationHeight, tiling_->dilationWidth, tiling_->padsHeight, tiling_->padsWidth,
Mconv/deformable_offsets_grad/op_kernel/arch35/deformable_offsets_grad.h+16-12
@@ -15,6 +15,10 @@
15#ifndef DEFORMABLE_OFFSET_H15#ifndef DEFORMABLE_OFFSET_H
16#define DEFORMABLE_OFFSET_H16#define DEFORMABLE_OFFSET_H
17#include "kernel_operator.h"17#include "kernel_operator.h"
18+#include "simt_api/asc_simt.h"
19+#include "simt_api/device_atomic_functions.h"
20+#include "simt_api/asc_fp16.h"
21+#include "simt_api/asc_bf16.h"
18using namespace AscendC;22using namespace AscendC;
19namespace DeformableOffsetsGrad {23namespace DeformableOffsetsGrad {
20const uint32_t MAX_THREAD_NUM = 512;24const uint32_t MAX_THREAD_NUM = 512;
@@ -110,7 +114,7 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void BilinearIn
110 T1 posOffset = lowH * inW_ * inChannel_ + lowW * inChannel_;114 T1 posOffset = lowH * inW_ * inChannel_ + lowW * inChannel_;
111 for (T1 c = 0; c < channelPerGroup; c++) {115 for (T1 c = 0; c < channelPerGroup; c++) {
112 T gradXValue1 = static_cast<T>(static_cast<float>(gradGmAddr[gradPrevIdx + c]) * gradXW1);116 T gradXValue1 = static_cast<T>(static_cast<float>(gradGmAddr[gradPrevIdx + c]) * gradXW1);
113- Simt::AtomicAdd(yGradXGmAddr + (yGradXIdx + posOffset + c), gradXValue1);117+ asc_atomic_add(yGradXGmAddr + (yGradXIdx + posOffset + c), gradXValue1);
114 v1 = static_cast<float>(inputXGmAddr[inputXIdx + posOffset + c]);118 v1 = static_cast<float>(inputXGmAddr[inputXIdx + posOffset + c]);
115 ComputeForGetFloorValueDeformableGrad<T, T1>(119 ComputeForGetFloorValueDeformableGrad<T, T1>(
116 gradGmAddr, c, gradPrevIdx, mask, gradOffsetsValueH, -v1, hw, gradOffsetsValueW, -v1, hh,120 gradGmAddr, c, gradPrevIdx, mask, gradOffsetsValueH, -v1, hw, gradOffsetsValueW, -v1, hh,
@@ -122,7 +126,7 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void BilinearIn
122 T1 posOffset = lowH * inW_ * inChannel_ + HighW * inChannel_;126 T1 posOffset = lowH * inW_ * inChannel_ + HighW * inChannel_;
123 for (T1 c = 0; c < channelPerGroup; c++) {127 for (T1 c = 0; c < channelPerGroup; c++) {
124 T gradXValue2 = static_cast<T>(static_cast<float>(gradGmAddr[gradPrevIdx + c]) * gradXW2);128 T gradXValue2 = static_cast<T>(static_cast<float>(gradGmAddr[gradPrevIdx + c]) * gradXW2);
125- Simt::AtomicAdd(yGradXGmAddr + (yGradXIdx + posOffset + c), gradXValue2);129+ asc_atomic_add(yGradXGmAddr + (yGradXIdx + posOffset + c), gradXValue2);
126 v2 = static_cast<float>(inputXGmAddr[inputXIdx + posOffset + c]);130 v2 = static_cast<float>(inputXGmAddr[inputXIdx + posOffset + c]);
127 ComputeForGetFloorValueDeformableGrad<T, T1>(131 ComputeForGetFloorValueDeformableGrad<T, T1>(
128 gradGmAddr, c, gradPrevIdx, mask, gradOffsetsValueH, -v2, lw, gradOffsetsValueW, v2, hh,132 gradGmAddr, c, gradPrevIdx, mask, gradOffsetsValueH, -v2, lw, gradOffsetsValueW, v2, hh,
@@ -134,7 +138,7 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void BilinearIn
134 T1 posOffset = HighH * inW_ * inChannel_ + lowW * inChannel_;138 T1 posOffset = HighH * inW_ * inChannel_ + lowW * inChannel_;
135 for (T1 c = 0; c < channelPerGroup; c++) {139 for (T1 c = 0; c < channelPerGroup; c++) {
136 T gradXValue3 = static_cast<T>(static_cast<float>(gradGmAddr[gradPrevIdx + c]) * gradXW3);140 T gradXValue3 = static_cast<T>(static_cast<float>(gradGmAddr[gradPrevIdx + c]) * gradXW3);
137- Simt::AtomicAdd(yGradXGmAddr + (yGradXIdx + posOffset + c), gradXValue3);141+ asc_atomic_add(yGradXGmAddr + (yGradXIdx + posOffset + c), gradXValue3);
138 v3 = static_cast<float>(inputXGmAddr[inputXIdx + posOffset + c]);142 v3 = static_cast<float>(inputXGmAddr[inputXIdx + posOffset + c]);
139 ComputeForGetFloorValueDeformableGrad<T, T1>(143 ComputeForGetFloorValueDeformableGrad<T, T1>(
140 gradGmAddr, c, gradPrevIdx, mask, gradOffsetsValueH, v3, hw, gradOffsetsValueW, -v3, lh,144 gradGmAddr, c, gradPrevIdx, mask, gradOffsetsValueH, v3, hw, gradOffsetsValueW, -v3, lh,
@@ -146,7 +150,7 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void BilinearIn
146 T1 posOffset = HighH * inW_ * inChannel_ + HighW * inChannel_;150 T1 posOffset = HighH * inW_ * inChannel_ + HighW * inChannel_;
147 for (T1 c = 0; c < channelPerGroup; c++) {151 for (T1 c = 0; c < channelPerGroup; c++) {
148 T gradXValue4 = static_cast<T>(static_cast<float>(gradGmAddr[gradPrevIdx + c]) * gradXW4);152 T gradXValue4 = static_cast<T>(static_cast<float>(gradGmAddr[gradPrevIdx + c]) * gradXW4);
149- Simt::AtomicAdd(yGradXGmAddr + (yGradXIdx + posOffset + c), gradXValue4);153+ asc_atomic_add(yGradXGmAddr + (yGradXIdx + posOffset + c), gradXValue4);
150 v4 = static_cast<float>(inputXGmAddr[inputXIdx + posOffset + c]);154 v4 = static_cast<float>(inputXGmAddr[inputXIdx + posOffset + c]);
151 ComputeForGetFloorValueDeformableGrad<T, T1>(155 ComputeForGetFloorValueDeformableGrad<T, T1>(
152 gradGmAddr, c, gradPrevIdx, mask, gradOffsetsValueH, v4, lw, gradOffsetsValueW, v4, lh,156 gradGmAddr, c, gradPrevIdx, mask, gradOffsetsValueH, v4, lw, gradOffsetsValueW, v4, lh,
@@ -205,8 +209,8 @@ __simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void ComputeSetValueGradX(
205 __gm__ T* gradGmAddr, __gm__ T* inputXGmAddr, __gm__ T* offsetsGmAddr, __gm__ T* yGradXGmAddr,209 __gm__ T* gradGmAddr, __gm__ T* inputXGmAddr, __gm__ T* offsetsGmAddr, __gm__ T* yGradXGmAddr,
206 __gm__ T* yGradOffsetsGmAddr, T1 blockClearProcessNum, T1 blockClearStartOffset)210 __gm__ T* yGradOffsetsGmAddr, T1 blockClearProcessNum, T1 blockClearStartOffset)
207{211{
208- for (T1 idx = static_cast<T1>(Simt::GetThreadIdx()); idx < blockClearProcessNum;212+ for (T1 idx = static_cast<T1>(threadIdx.x); idx < blockClearProcessNum;
209- idx += static_cast<T1>(Simt::GetThreadNum())) {213+ idx += static_cast<T1>(blockDim.x)) {
210 T1 curIdx = idx + blockClearStartOffset;214 T1 curIdx = idx + blockClearStartOffset;
211 yGradXGmAddr[curIdx] = static_cast<T>(0);215 yGradXGmAddr[curIdx] = static_cast<T>(0);
212 }216 }
@@ -221,8 +225,8 @@ __simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void ComputeDeformableOffset
221 T1 dilationW_, T1 padsH_, T1 padsW_, T1 batchSize_)225 T1 dilationW_, T1 padsH_, T1 padsW_, T1 batchSize_)
222{226{
223 T1 channelPerGroup = inChannel_ / deformableGroup_;227 T1 channelPerGroup = inChannel_ / deformableGroup_;
224- for (T1 idx = static_cast<T1>(Simt::GetThreadIdx()); idx < blockProcessNum;228+ for (T1 idx = static_cast<T1>(threadIdx.x); idx < blockProcessNum;
225- idx += static_cast<T1>(Simt::GetThreadNum())) {229+ idx += static_cast<T1>(blockDim.x)) {
226 T1 curIdx = idx + blockStartOffset;230 T1 curIdx = idx + blockStartOffset;
227 231 
228 // split multi-core and multi-thread according to the N * Ho * Wo * DeformableGroup * Kh *Kw232 // split multi-core and multi-thread according to the N * Ho * Wo * DeformableGroup * Kh *Kw
@@ -271,15 +275,15 @@ __aicore__ inline void DeformableOffsetGrad<T, T1>::Process()
271 }275 }
272 276 
273 if (blockId_ < tilingData_->clearGradXCoreNum) {277 if (blockId_ < tilingData_->clearGradXCoreNum) {
274- Simt::VF_CALL<ComputeSetValueGradX<T, T1>>(278+ asc_vf_call<ComputeSetValueGradX<T, T1>>(
275- Simt::Dim3{MAX_THREAD_NUM, 1, 1}, (__gm__ T*)gradGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),279+ dim3{MAX_THREAD_NUM, 1, 1}, (__gm__ T*)gradGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
276 (__gm__ T*)offsetsGm_.GetPhyAddr(), (__gm__ T*)yGradXGm_.GetPhyAddr(),280 (__gm__ T*)offsetsGm_.GetPhyAddr(), (__gm__ T*)yGradXGm_.GetPhyAddr(),
277 (__gm__ T*)yGradOffsetsGm_.GetPhyAddr(), blockClearProcessNum, blockClearStartOffset);281 (__gm__ T*)yGradOffsetsGm_.GetPhyAddr(), blockClearProcessNum, blockClearStartOffset);
278 }282 }
279 SyncAll();283 SyncAll();
280 if (blockId_ < tilingData_->realCoreNum) {284 if (blockId_ < tilingData_->realCoreNum) {
281- Simt::VF_CALL<ComputeDeformableOffsetGrad<T, T1>>(285+ asc_vf_call<ComputeDeformableOffsetGrad<T, T1>>(
282- Simt::Dim3{MAX_THREAD_NUM, 1, 1}, (__gm__ T*)gradGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),286+ dim3{MAX_THREAD_NUM, 1, 1}, (__gm__ T*)gradGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
283 (__gm__ T*)offsetsGm_.GetPhyAddr(), (__gm__ T*)yGradXGm_.GetPhyAddr(),287 (__gm__ T*)offsetsGm_.GetPhyAddr(), (__gm__ T*)yGradXGm_.GetPhyAddr(),
284 (__gm__ T*)yGradOffsetsGm_.GetPhyAddr(), blockProcessNum, blockStartOffset, tilingData_->dimKHeight,288 (__gm__ T*)yGradOffsetsGm_.GetPhyAddr(), blockProcessNum, blockStartOffset, tilingData_->dimKHeight,
285 tilingData_->dimKWidth, tilingData_->deformableGroups, offsetsChannel_, tilingData_->imgOutHeight,289 tilingData_->dimKWidth, tilingData_->deformableGroups, offsetsChannel_, tilingData_->imgOutHeight,
Mloss/ctc_loss_v2/op_kernel/arch35/ctc_loss_v2_base.h+1-0
@@ -19,6 +19,7 @@
19#include "kernel_operator.h"19#include "kernel_operator.h"
20#include "kernel_tiling/kernel_tiling.h"20#include "kernel_tiling/kernel_tiling.h"
21#include "ctc_loss_v2_tiling_key.h"21#include "ctc_loss_v2_tiling_key.h"
22+#include "simt_api/asc_simt.h"
22 23 
23#define INFINITY (__builtin_inff())24#define INFINITY (__builtin_inff())
24#define BSALIGNSIZE 3225#define BSALIGNSIZE 32
Mloss/ctc_loss_v2/op_kernel/arch35/ctc_loss_v2_fp16.h+16-16
@@ -67,12 +67,12 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void CalcLogAl
67 int32_t tgBatchStride, int32_t blank, int32_t tgTargetStride, __gm__ T *logProbsGm, __gm__ DataType *targetsGm, 67 int32_t tgBatchStride, int32_t blank, int32_t tgTargetStride, __gm__ T *logProbsGm, __gm__ DataType *targetsGm,
68 __gm__ DataType *inputLengthsGm, __gm__ DataType *targetLengthsGm, __gm__ T *negLogLikelihoodGm, __gm__ T *logAlphaGm, __gm__ float *tmpDataGm)68 __gm__ DataType *inputLengthsGm, __gm__ DataType *targetLengthsGm, __gm__ T *negLogLikelihoodGm, __gm__ T *logAlphaGm, __gm__ float *tmpDataGm)
69{69{
70- int32_t threadIdy = Simt::GetThreadIdx<1>();70+ int32_t threadIdy = threadIdx.y;
71- int32_t threadIdx = Simt::GetThreadIdx<0>();71+ int32_t thread_idx = threadIdx.x;
72- int32_t blockDimx = Simt::GetThreadNum<0>();72+ int32_t blockDimx = blockDim.x;
73- int32_t blockDimy = Simt::GetThreadNum<1>();73+ int32_t blockDimy = blockDim.y;
74- int32_t blkIdx = Simt::GetBlockIdx();74+ int32_t blkIdx = blockIdx.x;
75- int32_t blockNum = Simt::GetBlockNum();75+ int32_t blockNum = gridDim.x;
76 for (int32_t b = threadIdy + blkIdx * blockDimy; b < batchSize; b += blockNum * blockDimy) {76 for (int32_t b = threadIdy + blkIdx * blockDimy; b < batchSize; b += blockNum * blockDimy) {
77 ThreadType inputLength = inputLengthsGm[b];77 ThreadType inputLength = inputLengthsGm[b];
78 ThreadType targetLength = targetLengthsGm[b];78 ThreadType targetLength = targetLengthsGm[b];
@@ -81,7 +81,7 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void CalcLogAl
81 ThreadType tgBatchOffset = ProcessTgBatchOffsetsFp16<DataType, ThreadType>(targetLengthsGm, targetsDim, tgBatchStride, b);81 ThreadType tgBatchOffset = ProcessTgBatchOffsetsFp16<DataType, ThreadType>(targetLengthsGm, targetsDim, tgBatchStride, b);
82 82 
83 if (inputLength == 0) {83 if (inputLength == 0) {
84- if (Simt::GetThreadIdx<0>() == 0) {84+ if (threadIdx.x == 0) {
85 float log_likelihood = targetLength == 0 ? 0 : neginf;85 float log_likelihood = targetLength == 0 ? 0 : neginf;
86 negLogLikelihoodGm[b] = -static_cast<T>(log_likelihood);86 negLogLikelihoodGm[b] = -static_cast<T>(log_likelihood);
87 }87 }
@@ -89,7 +89,7 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void CalcLogAl
89 }89 }
90 90 
91 for (int32_t block_s = 0; block_s < laInputStride; block_s += blockDimx) {91 for (int32_t block_s = 0; block_s < laInputStride; block_s += blockDimx) {
92- int32_t s = threadIdx + block_s;92+ int32_t s = thread_idx + block_s;
93 float la = 0;93 float la = 0;
94 switch (s) {94 switch (s) {
95 case 0:95 case 0:
@@ -109,7 +109,7 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void CalcLogAl
109 }109 }
110 }110 }
111 for (int32_t block_s = 0; block_s < laInputStride; block_s += blockDimx) {111 for (int32_t block_s = 0; block_s < laInputStride; block_s += blockDimx) {
112- int32_t s = threadIdx + block_s;112+ int32_t s = thread_idx + block_s;
113 int32_t currentChar;113 int32_t currentChar;
114 bool haveThree;114 bool haveThree;
115 if (s < 2 * targetLength + 1 && targetLength > 0) {115 if (s < 2 * targetLength + 1 && targetLength > 0) {
@@ -121,7 +121,7 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void CalcLogAl
121 haveThree = false;121 haveThree = false;
122 }122 }
123 for (int32_t t = 1; t < maxInputLength; t++) {123 for (int32_t t = 1; t < maxInputLength; t++) {
124- Simt::ThreadBarrier();124+ asc_syncthreads();
125 if ((t < inputLength) && (s < 2 * targetLength + 1)) {125 if ((t < inputLength) && (s < 2 * targetLength + 1)) {
126 float la1 = tmpDataGm[laBatchOffset + laInputStride * (t - 1) + s];126 float la1 = tmpDataGm[laBatchOffset + laInputStride * (t - 1) + s];
127 float lamax = la1;127 float lamax = la1;
@@ -147,24 +147,24 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void CalcLogAl
147 lamax = 0;147 lamax = 0;
148 }148 }
149 tmpDataGm[laBatchOffset + laInputStride * t + s] =149 tmpDataGm[laBatchOffset + laInputStride * t + s] =
150- Simt::Log1p(Simt::Exp(la1 - lamax) + Simt::Exp(la2 - lamax) + Simt::Exp(la3 - lamax) - 1) +150+ log1pf(expf(la1 - lamax) + expf(la2 - lamax) + expf(la3 - lamax) - 1) +
151 lamax + static_cast<float>(logProbsGm[lpBatchOffset + t * lpInputStride + currentChar]);151 lamax + static_cast<float>(logProbsGm[lpBatchOffset + t * lpInputStride + currentChar]);
152 logAlphaGm[laBatchOffset + laInputStride * t + s] =152 logAlphaGm[laBatchOffset + laInputStride * t + s] =
153 static_cast<T>(tmpDataGm[laBatchOffset + laInputStride * t + s]);153 static_cast<T>(tmpDataGm[laBatchOffset + laInputStride * t + s]);
154 }154 }
155 }155 }
156 }156 }
157- Simt::ThreadBarrier();157+ asc_syncthreads();
158 158 
159 // compute the loss159 // compute the loss
160- if (threadIdx == 0) {160+ if (thread_idx == 0) {
161 float l1 = tmpDataGm[laBatchOffset + laInputStride * (inputLength - 1) + (targetLength * 2)];161 float l1 = tmpDataGm[laBatchOffset + laInputStride * (inputLength - 1) + (targetLength * 2)];
162 float l2 = targetLength > 0 ?162 float l2 = targetLength > 0 ?
163 tmpDataGm[laBatchOffset + laInputStride * (inputLength - 1) + (targetLength * 2 - 1)] :163 tmpDataGm[laBatchOffset + laInputStride * (inputLength - 1) + (targetLength * 2 - 1)] :
164 neginf;164 neginf;
165 float m = ((l1 > l2) ? l1 : l2);165 float m = ((l1 > l2) ? l1 : l2);
166 m = ((m == neginf) ? 0 : m);166 m = ((m == neginf) ? 0 : m);
167- float log_likelihood = Simt::Log1p(Simt::Exp(l1 - m) + Simt::Exp(l2 - m) - 1) + m;167+ float log_likelihood = log1pf(expf(l1 - m) + expf(l2 - m) - 1) + m;
168 negLogLikelihoodGm[b] = -static_cast<T>(log_likelihood);168 negLogLikelihoodGm[b] = -static_cast<T>(log_likelihood);
169 }169 }
170 }170 }
@@ -209,7 +209,7 @@ template <typename T, typename DataType, typename ThreadType> __aicore__ inline
209 int32_t tgBatchStride = this->tdPtr->tgBatchStride;209 int32_t tgBatchStride = this->tdPtr->tgBatchStride;
210 int32_t gridY = this->tdPtr->gridY;210 int32_t gridY = this->tdPtr->gridY;
211 if constexpr (sizeof(ThreadType) == sizeof(int32_t)) {211 if constexpr (sizeof(ThreadType) == sizeof(int32_t)) {
212- Simt::VF_CALL<SimtComputeFp16Int32<T, DataType, ThreadType>>(Simt::Dim3(blockDimX, blockDimY), batchSize, laInputStride, laBatchStride,212+ asc_vf_call<SimtComputeFp16Int32<T, DataType, ThreadType>>(dim3(blockDimX, blockDimY), batchSize, laInputStride, laBatchStride,
213 lpBatchStride, maxInputLength, lpInputStride, targetsDim, tgBatchStride, blank, tgTargetStride,213 lpBatchStride, maxInputLength, lpInputStride, targetsDim, tgBatchStride, blank, tgTargetStride,
214 (__gm__ T *)(this->logProbsDataGm.GetPhyAddr()), (__gm__ DataType *)(this->targetsDataGm.GetPhyAddr()),214 (__gm__ T *)(this->logProbsDataGm.GetPhyAddr()), (__gm__ DataType *)(this->targetsDataGm.GetPhyAddr()),
215 (__gm__ DataType *)(this->inputLengthsGm.GetPhyAddr()), (__gm__ DataType *)(this->targetLengthsGm.GetPhyAddr()),215 (__gm__ DataType *)(this->inputLengthsGm.GetPhyAddr()), (__gm__ DataType *)(this->targetLengthsGm.GetPhyAddr()),
@@ -217,7 +217,7 @@ template <typename T, typename DataType, typename ThreadType> __aicore__ inline
217 (__gm__ float *)(tmpFloatData.GetPhyAddr()));217 (__gm__ float *)(tmpFloatData.GetPhyAddr()));
218 }218 }
219 if constexpr (sizeof(ThreadType) == sizeof(int64_t)) {219 if constexpr (sizeof(ThreadType) == sizeof(int64_t)) {
220- Simt::VF_CALL<SimtComputeFp16<T, DataType, ThreadType>>(Simt::Dim3(blockDimX, blockDimY), batchSize, laInputStride, laBatchStride,220+ asc_vf_call<SimtComputeFp16<T, DataType, ThreadType>>(dim3(blockDimX, blockDimY), batchSize, laInputStride, laBatchStride,
221 lpBatchStride, maxInputLength, lpInputStride, targetsDim, tgBatchStride, blank, tgTargetStride,221 lpBatchStride, maxInputLength, lpInputStride, targetsDim, tgBatchStride, blank, tgTargetStride,
222 (__gm__ T *)(this->logProbsDataGm.GetPhyAddr()), (__gm__ DataType *)(this->targetsDataGm.GetPhyAddr()),222 (__gm__ T *)(this->logProbsDataGm.GetPhyAddr()), (__gm__ DataType *)(this->targetsDataGm.GetPhyAddr()),
223 (__gm__ DataType *)(this->inputLengthsGm.GetPhyAddr()), (__gm__ DataType *)(this->targetLengthsGm.GetPhyAddr()),223 (__gm__ DataType *)(this->inputLengthsGm.GetPhyAddr()), (__gm__ DataType *)(this->targetLengthsGm.GetPhyAddr()),
Mloss/ctc_loss_v2/op_kernel/arch35/ctc_loss_v2_fp32.h+18-18
@@ -89,12 +89,12 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void CalcLogAlp
89 int32_t tgTargetStride, __gm__ T *logProbsGm, __gm__ DataType *targetsGm, __gm__ DataType *inputLengthsGm,89 int32_t tgTargetStride, __gm__ T *logProbsGm, __gm__ DataType *targetsGm, __gm__ DataType *inputLengthsGm,
90 __gm__ DataType *targetLengthsGm, __gm__ T *negLogLikelihoodGm, __gm__ T *logAlphaGm, __ubuf__ DataType *tensor)90 __gm__ DataType *targetLengthsGm, __gm__ T *negLogLikelihoodGm, __gm__ T *logAlphaGm, __ubuf__ DataType *tensor)
91{91{
92- int32_t threadIdy = Simt::GetThreadIdx<1>();92+ int32_t threadIdy = threadIdx.y;
93- int32_t threadIdx = Simt::GetThreadIdx<0>();93+ int32_t thread_idx = threadIdx.x;
94- int32_t blockDimx = Simt::GetThreadNum<0>();94+ int32_t blockDimx = blockDim.x;
95- int32_t blockDimy = Simt::GetThreadNum<1>();95+ int32_t blockDimy = blockDim.y;
96- int32_t blkIdx = Simt::GetBlockIdx();96+ int32_t blkIdx = blockIdx.x;
97- int32_t blockNum = Simt::GetBlockNum();97+ int32_t blockNum = gridDim.x;
98 for (int32_t b = threadIdy + blkIdx * blockDimy; b < batchSize; b += blockNum * blockDimy) {98 for (int32_t b = threadIdy + blkIdx * blockDimy; b < batchSize; b += blockNum * blockDimy) {
99 ThreadType inputLength = inputLengthsGm[b];99 ThreadType inputLength = inputLengthsGm[b];
100 ThreadType targetLength = targetLengthsGm[b];100 ThreadType targetLength = targetLengthsGm[b];
@@ -103,14 +103,14 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void CalcLogAlp
103 ThreadType tgBatchOffset = ProcessTgBatchOffsets<DataType>(tensor, targetsDim, tgBatchStride, b);103 ThreadType tgBatchOffset = ProcessTgBatchOffsets<DataType>(tensor, targetsDim, tgBatchStride, b);
104 104 
105 if (inputLength == 0) {105 if (inputLength == 0) {
106- if (threadIdx == 0) {106+ if (thread_idx == 0) {
107 float log_likelihood = targetLength == 0 ? 0 : neginf;107 float log_likelihood = targetLength == 0 ? 0 : neginf;
108 negLogLikelihoodGm[b] = -log_likelihood;108 negLogLikelihoodGm[b] = -log_likelihood;
109 }109 }
110 continue;110 continue;
111 }111 }
112 for (int32_t block_s = 0; block_s < laInputStride; block_s += blockDimx) {112 for (int32_t block_s = 0; block_s < laInputStride; block_s += blockDimx) {
113- int32_t s = threadIdx + block_s;113+ int32_t s = thread_idx + block_s;
114 float la = 0;114 float la = 0;
115 if(s == 0) {115 if(s == 0) {
116 la = logProbsGm[lpBatchOffset + blank];116 la = logProbsGm[lpBatchOffset + blank];
@@ -127,7 +127,7 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void CalcLogAlp
127 }127 }
128 128 
129 for (int32_t block_s = 0; block_s < laInputStride; block_s += blockDimx) {129 for (int32_t block_s = 0; block_s < laInputStride; block_s += blockDimx) {
130- int32_t s = threadIdx + block_s;130+ int32_t s = thread_idx + block_s;
131 int32_t currentChar;131 int32_t currentChar;
132 bool haveThree;132 bool haveThree;
133 if (s < 2 * targetLength + 1 && targetLength > 0) {133 if (s < 2 * targetLength + 1 && targetLength > 0) {
@@ -139,7 +139,7 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void CalcLogAlp
139 haveThree = false;139 haveThree = false;
140 }140 }
141 for (int32_t t = 1; t < maxInputLength; t++) {141 for (int32_t t = 1; t < maxInputLength; t++) {
142- Simt::ThreadBarrier();142+ asc_syncthreads();
143 if ((t < inputLength) && (s < 2 * targetLength + 1)) {143 if ((t < inputLength) && (s < 2 * targetLength + 1)) {
144 float x = logProbsGm[lpBatchOffset + t * lpInputStride + currentChar];144 float x = logProbsGm[lpBatchOffset + t * lpInputStride + currentChar];
145 float la1 = logAlphaGm[laBatchOffset + laInputStride * (t - 1) + s];145 float la1 = logAlphaGm[laBatchOffset + laInputStride * (t - 1) + s];
@@ -147,27 +147,27 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void CalcLogAlp
147 float la2 = 0;147 float la2 = 0;
148 float la3 = 0;148 float la3 = 0;
149 la2 = (s > 0) ? logAlphaGm[laBatchOffset + laInputStride * (t - 1) + (s - 1)] : neginf;149 la2 = (s > 0) ? logAlphaGm[laBatchOffset + laInputStride * (t - 1) + (s - 1)] : neginf;
150- lamax = (s > 0) ? (Simt::Max(la2, lamax)) : lamax;150+ lamax = (s > 0) ? (fmaxf(la2, lamax)) : lamax;
151 la3 = (haveThree == true) ? logAlphaGm[laBatchOffset + laInputStride * (t - 1) + (s - 2)] : neginf;151 la3 = (haveThree == true) ? logAlphaGm[laBatchOffset + laInputStride * (t - 1) + (s - 2)] : neginf;
152- lamax = (haveThree == true) ? (Simt::Max(la3, lamax)) : lamax;152+ lamax = (haveThree == true) ? (fmaxf(la3, lamax)) : lamax;
153 lamax = (lamax == neginf) ? 0 : lamax;153 lamax = (lamax == neginf) ? 0 : lamax;
154 logAlphaGm[laBatchOffset + laInputStride * t + s] =154 logAlphaGm[laBatchOffset + laInputStride * t + s] =
155- Simt::Log1p(Simt::Exp(la1 - lamax) + Simt::Exp(la2 - lamax) + Simt::Exp(la3 - lamax) - 1) +155+ log1pf(expf(la1 - lamax) + expf(la2 - lamax) + expf(la3 - lamax) - 1) +
156 lamax + x;156 lamax + x;
157 }157 }
158 }158 }
159 }159 }
160- Simt::ThreadBarrier();160+ asc_syncthreads();
161 161 
162 // compute the loss162 // compute the loss
163- if (threadIdx == 0) {163+ if (thread_idx == 0) {
164 float l1 = logAlphaGm[laBatchOffset + laInputStride * (inputLength - 1) + (targetLength * 2)];164 float l1 = logAlphaGm[laBatchOffset + laInputStride * (inputLength - 1) + (targetLength * 2)];
165 float l2 = targetLength > 0 ?165 float l2 = targetLength > 0 ?
166 logAlphaGm[laBatchOffset + laInputStride * (inputLength - 1) + (targetLength * 2 - 1)] :166 logAlphaGm[laBatchOffset + laInputStride * (inputLength - 1) + (targetLength * 2 - 1)] :
167 neginf;167 neginf;
168 float m = ((l1 > l2) ? l1 : l2);168 float m = ((l1 > l2) ? l1 : l2);
169 m = ((m == neginf) ? 0 : m);169 m = ((m == neginf) ? 0 : m);
170- float log_likelihood = Simt::Log1p(Simt::Exp(l1 - m) + Simt::Exp(l2 - m) - 1) + m;170+ float log_likelihood = log1pf(expf(l1 - m) + expf(l2 - m) - 1) + m;
171 negLogLikelihoodGm[b] = -log_likelihood;171 negLogLikelihoodGm[b] = -log_likelihood;
172 }172 }
173 }173 }
@@ -213,7 +213,7 @@ template <typename T, typename DataType, typename ThreadType> __aicore__ inline
213 int32_t gridY = this->tdPtr->gridY;213 int32_t gridY = this->tdPtr->gridY;
214 214 
215 if constexpr (sizeof(ThreadType) == sizeof(int32_t)) {215 if constexpr (sizeof(ThreadType) == sizeof(int32_t)) {
216- Simt::VF_CALL<SimtComputeINT32<T, DataType, ThreadType>>(Simt::Dim3(blockDimX, blockDimY), batchSize, laInputStride, laBatchStride,216+ asc_vf_call<SimtComputeINT32<T, DataType, ThreadType>>(dim3(blockDimX, blockDimY), batchSize, laInputStride, laBatchStride,
217 lpBatchStride, maxInputLength, lpInputStride, targetsDim, tgBatchStride, blank, tgTargetStride,217 lpBatchStride, maxInputLength, lpInputStride, targetsDim, tgBatchStride, blank, tgTargetStride,
218 (__gm__ T *)(this->logProbsDataGm.GetPhyAddr()), (__gm__ DataType *)(this->targetsDataGm.GetPhyAddr()),218 (__gm__ T *)(this->logProbsDataGm.GetPhyAddr()), (__gm__ DataType *)(this->targetsDataGm.GetPhyAddr()),
219 (__gm__ DataType *)(this->inputLengthsGm.GetPhyAddr()), (__gm__ DataType *)(this->targetLengthsGm.GetPhyAddr()),219 (__gm__ DataType *)(this->inputLengthsGm.GetPhyAddr()), (__gm__ DataType *)(this->targetLengthsGm.GetPhyAddr()),
@@ -221,7 +221,7 @@ template <typename T, typename DataType, typename ThreadType> __aicore__ inline
221 (__ubuf__ DataType *)(targetOffsetTensor.GetPhyAddr()));221 (__ubuf__ DataType *)(targetOffsetTensor.GetPhyAddr()));
222 }222 }
223 if constexpr (sizeof(ThreadType) == sizeof(int64_t)) {223 if constexpr (sizeof(ThreadType) == sizeof(int64_t)) {
224- Simt::VF_CALL<SimtCompute<T, DataType, ThreadType>>(Simt::Dim3(blockDimX, blockDimY), batchSize, laInputStride, laBatchStride,224+ asc_vf_call<SimtCompute<T, DataType, ThreadType>>(dim3(blockDimX, blockDimY), batchSize, laInputStride, laBatchStride,
225 lpBatchStride, maxInputLength, lpInputStride, targetsDim, tgBatchStride, blank, tgTargetStride,225 lpBatchStride, maxInputLength, lpInputStride, targetsDim, tgBatchStride, blank, tgTargetStride,
226 (__gm__ T *)(this->logProbsDataGm.GetPhyAddr()), (__gm__ DataType *)(this->targetsDataGm.GetPhyAddr()),226 (__gm__ T *)(this->logProbsDataGm.GetPhyAddr()), (__gm__ DataType *)(this->targetsDataGm.GetPhyAddr()),
227 (__gm__ DataType *)(this->inputLengthsGm.GetPhyAddr()), (__gm__ DataType *)(this->targetLengthsGm.GetPhyAddr()),227 (__gm__ DataType *)(this->inputLengthsGm.GetPhyAddr()), (__gm__ DataType *)(this->targetLengthsGm.GetPhyAddr()),
Mloss/ctc_loss_v2_grad/op_kernel/arch35/ctc_loss_v2_grad.h+25-25
@@ -168,11 +168,11 @@ __simt_vf__ LAUNCH_BOUND(THREAD_NUM) __aicore__ void CalGradCompute(
168 __gm__ float* logBetaGm, __gm__ float* tempGradGm, ThreadType maxInputLength, ThreadType batchSize,168 __gm__ float* logBetaGm, __gm__ float* tempGradGm, ThreadType maxInputLength, ThreadType batchSize,
169 ThreadType symbolSet, ThreadType zeroInfinity)169 ThreadType symbolSet, ThreadType zeroInfinity)
170{170{
171- ThreadType threadIdx = AscendC::Simt::GetThreadIdx<0>();171+ ThreadType thread_idx = threadIdx.x;
172- ThreadType blockDimX = AscendC::Simt::GetThreadNum<0>();172+ ThreadType blockDimX = blockDim.x;
173 ThreadType gradBatchOffset = batchSize * symbolSet;173 ThreadType gradBatchOffset = batchSize * symbolSet;
174 ThreadType length = maxInputLength * batchSize * symbolSet;174 ThreadType length = maxInputLength * batchSize * symbolSet;
175- for (ThreadType index = threadIdx + block_idx * blockDimX; index < length; index += block_num * blockDimX) {175+ for (ThreadType index = thread_idx + block_idx * blockDimX; index < length; index += block_num * blockDimX) {
176 ThreadType t = index / gradBatchOffset;176 ThreadType t = index / gradBatchOffset;
177 ThreadType offset = index % gradBatchOffset;177 ThreadType offset = index % gradBatchOffset;
178 ThreadType b = offset / symbolSet;178 ThreadType b = offset / symbolSet;
@@ -201,10 +201,10 @@ __simt_vf__ LAUNCH_BOUND(THREAD_NUM) __aicore__ void UpdateLcabCompute(
201 ThreadType batchSize, ThreadType symbolSet, ThreadType zeroInfinity, ThreadType blank, ThreadType logAlphaT,201 ThreadType batchSize, ThreadType symbolSet, ThreadType zeroInfinity, ThreadType blank, ThreadType logAlphaT,
202 ThreadType alphaLength, ThreadType targetsDimNum, ThreadType sDimRange)202 ThreadType alphaLength, ThreadType targetsDimNum, ThreadType sDimRange)
203{203{
204- ThreadType threadIdx = AscendC::Simt::GetThreadIdx<0>();204+ ThreadType thread_idx = threadIdx.x;
205- ThreadType blockDimX = AscendC::Simt::GetThreadNum<0>();205+ ThreadType blockDimX = blockDim.x;
206 ThreadType length = maxInputLength * batchSize;206 ThreadType length = maxInputLength * batchSize;
207- for (ThreadType index = threadIdx + block_idx * blockDimX; index < length; index += block_num * blockDimX) {207+ for (ThreadType index = thread_idx + block_idx * blockDimX; index < length; index += block_num * blockDimX) {
208 ThreadType b = index / maxInputLength;208 ThreadType b = index / maxInputLength;
209 ThreadType t = index % maxInputLength;209 ThreadType t = index % maxInputLength;
210 if ((t >= maxInputLength) || (b >= batchSize)) {210 if ((t >= maxInputLength) || (b >= batchSize)) {
@@ -271,10 +271,10 @@ __simt_vf__ LAUNCH_BOUND(THREAD_NUM) __aicore__ void LogBetaCompute(
271 ThreadType targetsDimNum, ThreadType sDimRange)271 ThreadType targetsDimNum, ThreadType sDimRange)
272{272{
273 constexpr float neginf = -INFINITY;273 constexpr float neginf = -INFINITY;
274- ThreadType threadIdy = AscendC::Simt::GetThreadIdx<1>();274+ ThreadType threadIdy = threadIdx.y;
275- ThreadType threadIdx = AscendC::Simt::GetThreadIdx<0>();275+ ThreadType thread_idx = threadIdx.x;
276- ThreadType blockDimx = AscendC::Simt::GetThreadNum<0>();276+ ThreadType blockDimx = blockDim.x;
277- ThreadType blockDimy = AscendC::Simt::GetThreadNum<1>();277+ ThreadType blockDimy = blockDim.y;
278 for (ThreadType index = threadIdy + block_idx * blockDimy; index < batchSize; index += block_num * blockDimy) {278 for (ThreadType index = threadIdy + block_idx * blockDimy; index < batchSize; index += block_num * blockDimy) {
279 ThreadType b = index;279 ThreadType b = index;
280 ThreadType inputLength = inputLengthsGm[b];280 ThreadType inputLength = inputLengthsGm[b];
@@ -288,7 +288,7 @@ __simt_vf__ LAUNCH_BOUND(THREAD_NUM) __aicore__ void LogBetaCompute(
288 ProcessTgBatchOffsets<T, DataType, ThreadType>(b, targetLengthsGm, targetsDimNum, sDimRange);288 ProcessTgBatchOffsets<T, DataType, ThreadType>(b, targetLengthsGm, targetsDimNum, sDimRange);
289 for (ThreadType block_s = alphaLength - 1 - ((alphaLength - 1) % blockDimx); block_s >= 0;289 for (ThreadType block_s = alphaLength - 1 - ((alphaLength - 1) % blockDimx); block_s >= 0;
290 block_s -= blockDimx) {290 block_s -= blockDimx) {
291- ThreadType s = threadIdx + block_s;291+ ThreadType s = thread_idx + block_s;
292 float lb;292 float lb;
293 if (s == 2 * targetLength) {293 if (s == 2 * targetLength) {
294 lb = logProbsGm[logProbsBatchOffset + (inputLength - 1) * batchSize * symbolSet + blank];294 lb = logProbsGm[logProbsBatchOffset + (inputLength - 1) * batchSize * symbolSet + blank];
@@ -305,7 +305,7 @@ __simt_vf__ LAUNCH_BOUND(THREAD_NUM) __aicore__ void LogBetaCompute(
305 }305 }
306 for (ThreadType block_s = alphaLength - 1 - ((alphaLength - 1) % blockDimx); block_s >= 0;306 for (ThreadType block_s = alphaLength - 1 - ((alphaLength - 1) % blockDimx); block_s >= 0;
307 block_s -= blockDimx) {307 block_s -= blockDimx) {
308- ThreadType s = threadIdx + block_s;308+ ThreadType s = thread_idx + block_s;
309 ThreadType currentTargetPrime;309 ThreadType currentTargetPrime;
310 bool haveThree;310 bool haveThree;
311 if (s < 2 * targetLength + 1 && targetLength > 0) {311 if (s < 2 * targetLength + 1 && targetLength > 0) {
@@ -320,7 +320,7 @@ __simt_vf__ LAUNCH_BOUND(THREAD_NUM) __aicore__ void LogBetaCompute(
320 haveThree = false;320 haveThree = false;
321 }321 }
322 for (ThreadType t = maxInputLength - 2; t >= 0; t--) {322 for (ThreadType t = maxInputLength - 2; t >= 0; t--) {
323- Simt::ThreadBarrier();323+ asc_syncthreads();
324 if ((t < inputLength - 1) && (s < 2 * targetLength + 1)) {324 if ((t < inputLength - 1) && (s < 2 * targetLength + 1)) {
325 float lb1 = logBetaGm[logBetaBatchOffset + (t + 1) * alphaLength + s];325 float lb1 = logBetaGm[logBetaBatchOffset + (t + 1) * alphaLength + s];
326 float lbmax = lb1;326 float lbmax = lb1;
@@ -360,8 +360,8 @@ __aicore__ inline void CTCLossV2Grad<T, DataType, ThreadType>::Process()
360 ThreadType blockDimX = tilingData_->blockDimX;360 ThreadType blockDimX = tilingData_->blockDimX;
361 ThreadType blockDimY = tilingData_->blockDimY;361 ThreadType blockDimY = tilingData_->blockDimY;
362 if constexpr (sizeof(ThreadType) == INT_SIZE_32) {362 if constexpr (sizeof(ThreadType) == INT_SIZE_32) {
363- Simt::VF_CALL<LogBetaCompute<T, DataType, ThreadType, THREAD_NUM_1024>>(363+ asc_vf_call<LogBetaCompute<T, DataType, ThreadType, THREAD_NUM_1024>>(
364- Simt::Dim3(blockDimX, blockDimY), (__gm__ T*)(gradOutGm.GetPhyAddr()), (__gm__ T*)(logProbsGm.GetPhyAddr()),364+ dim3(blockDimX, blockDimY), (__gm__ T*)(gradOutGm.GetPhyAddr()), (__gm__ T*)(logProbsGm.GetPhyAddr()),
365 (__gm__ DataType*)(targetsGm.GetPhyAddr()), (__gm__ DataType*)(inputLengthsGm.GetPhyAddr()),365 (__gm__ DataType*)(targetsGm.GetPhyAddr()), (__gm__ DataType*)(inputLengthsGm.GetPhyAddr()),
366 (__gm__ DataType*)(targetLengthsGm.GetPhyAddr()), (__gm__ T*)(negLogLikelihoodGm.GetPhyAddr()),366 (__gm__ DataType*)(targetLengthsGm.GetPhyAddr()), (__gm__ T*)(negLogLikelihoodGm.GetPhyAddr()),
367 (__gm__ T*)(logAlphaGm.GetPhyAddr()), (__gm__ T*)(gradGm.GetPhyAddr()),367 (__gm__ T*)(logAlphaGm.GetPhyAddr()), (__gm__ T*)(gradGm.GetPhyAddr()),
@@ -371,8 +371,8 @@ __aicore__ inline void CTCLossV2Grad<T, DataType, ThreadType>::Process()
371 tilingData_->sDimRange);371 tilingData_->sDimRange);
372 372 
373 SyncAll();373 SyncAll();
374- Simt::VF_CALL<UpdateLcabCompute<T, DataType, ThreadType, THREAD_NUM_1024>>(374+ asc_vf_call<UpdateLcabCompute<T, DataType, ThreadType, THREAD_NUM_1024>>(
375- Simt::Dim3(tilingData_->updateLcabThreadNum, 1), (__gm__ T*)(gradOutGm.GetPhyAddr()),375+ dim3(tilingData_->updateLcabThreadNum, 1), (__gm__ T*)(gradOutGm.GetPhyAddr()),
376 (__gm__ T*)(logProbsGm.GetPhyAddr()), (__gm__ DataType*)(targetsGm.GetPhyAddr()),376 (__gm__ T*)(logProbsGm.GetPhyAddr()), (__gm__ DataType*)(targetsGm.GetPhyAddr()),
377 (__gm__ DataType*)(inputLengthsGm.GetPhyAddr()), (__gm__ DataType*)(targetLengthsGm.GetPhyAddr()),377 (__gm__ DataType*)(inputLengthsGm.GetPhyAddr()), (__gm__ DataType*)(targetLengthsGm.GetPhyAddr()),
378 (__gm__ T*)(negLogLikelihoodGm.GetPhyAddr()), (__gm__ T*)(logAlphaGm.GetPhyAddr()),378 (__gm__ T*)(negLogLikelihoodGm.GetPhyAddr()), (__gm__ T*)(logAlphaGm.GetPhyAddr()),
@@ -381,8 +381,8 @@ __aicore__ inline void CTCLossV2Grad<T, DataType, ThreadType>::Process()
381 tilingData_->symbolSet, tilingData_->zeroInfinity, tilingData_->BLANK, tilingData_->logAlphaT,381 tilingData_->symbolSet, tilingData_->zeroInfinity, tilingData_->BLANK, tilingData_->logAlphaT,
382 tilingData_->alphaLength, tilingData_->targetsDimNum, tilingData_->sDimRange);382 tilingData_->alphaLength, tilingData_->targetsDimNum, tilingData_->sDimRange);
383 SyncAll();383 SyncAll();
384- Simt::VF_CALL<CalGradCompute<T, DataType, ThreadType, THREAD_NUM_1024>>(384+ asc_vf_call<CalGradCompute<T, DataType, ThreadType, THREAD_NUM_1024>>(
385- Simt::Dim3(tilingData_->calGradThreadNum, 1), (__gm__ T*)(gradOutGm.GetPhyAddr()),385+ dim3(tilingData_->calGradThreadNum, 1), (__gm__ T*)(gradOutGm.GetPhyAddr()),
386 (__gm__ T*)(logProbsGm.GetPhyAddr()), (__gm__ DataType*)(targetsGm.GetPhyAddr()),386 (__gm__ T*)(logProbsGm.GetPhyAddr()), (__gm__ DataType*)(targetsGm.GetPhyAddr()),
387 (__gm__ DataType*)(inputLengthsGm.GetPhyAddr()), (__gm__ DataType*)(targetLengthsGm.GetPhyAddr()),387 (__gm__ DataType*)(inputLengthsGm.GetPhyAddr()), (__gm__ DataType*)(targetLengthsGm.GetPhyAddr()),
388 (__gm__ T*)(negLogLikelihoodGm.GetPhyAddr()), (__gm__ T*)(logAlphaGm.GetPhyAddr()),388 (__gm__ T*)(negLogLikelihoodGm.GetPhyAddr()), (__gm__ T*)(logAlphaGm.GetPhyAddr()),
@@ -392,8 +392,8 @@ __aicore__ inline void CTCLossV2Grad<T, DataType, ThreadType>::Process()
392 }392 }
393 393 
394 if constexpr (sizeof(ThreadType) == INT_SIZE_64) {394 if constexpr (sizeof(ThreadType) == INT_SIZE_64) {
395- Simt::VF_CALL<LogBetaCompute<T, DataType, ThreadType, THREAD_NUM_512>>(395+ asc_vf_call<LogBetaCompute<T, DataType, ThreadType, THREAD_NUM_512>>(
396- Simt::Dim3(blockDimX, blockDimY), (__gm__ T*)(gradOutGm.GetPhyAddr()), (__gm__ T*)(logProbsGm.GetPhyAddr()),396+ dim3(blockDimX, blockDimY), (__gm__ T*)(gradOutGm.GetPhyAddr()), (__gm__ T*)(logProbsGm.GetPhyAddr()),
397 (__gm__ DataType*)(targetsGm.GetPhyAddr()), (__gm__ DataType*)(inputLengthsGm.GetPhyAddr()),397 (__gm__ DataType*)(targetsGm.GetPhyAddr()), (__gm__ DataType*)(inputLengthsGm.GetPhyAddr()),
398 (__gm__ DataType*)(targetLengthsGm.GetPhyAddr()), (__gm__ T*)(negLogLikelihoodGm.GetPhyAddr()),398 (__gm__ DataType*)(targetLengthsGm.GetPhyAddr()), (__gm__ T*)(negLogLikelihoodGm.GetPhyAddr()),
399 (__gm__ T*)(logAlphaGm.GetPhyAddr()), (__gm__ T*)(gradGm.GetPhyAddr()),399 (__gm__ T*)(logAlphaGm.GetPhyAddr()), (__gm__ T*)(gradGm.GetPhyAddr()),
@@ -403,8 +403,8 @@ __aicore__ inline void CTCLossV2Grad<T, DataType, ThreadType>::Process()
403 tilingData_->sDimRange);403 tilingData_->sDimRange);
404 404 
405 SyncAll();405 SyncAll();
406- Simt::VF_CALL<UpdateLcabCompute<T, DataType, ThreadType, THREAD_NUM_512>>(406+ asc_vf_call<UpdateLcabCompute<T, DataType, ThreadType, THREAD_NUM_512>>(
407- Simt::Dim3(tilingData_->updateLcabThreadNum, 1), (__gm__ T*)(gradOutGm.GetPhyAddr()),407+ dim3(tilingData_->updateLcabThreadNum, 1), (__gm__ T*)(gradOutGm.GetPhyAddr()),
408 (__gm__ T*)(logProbsGm.GetPhyAddr()), (__gm__ DataType*)(targetsGm.GetPhyAddr()),408 (__gm__ T*)(logProbsGm.GetPhyAddr()), (__gm__ DataType*)(targetsGm.GetPhyAddr()),
409 (__gm__ DataType*)(inputLengthsGm.GetPhyAddr()), (__gm__ DataType*)(targetLengthsGm.GetPhyAddr()),409 (__gm__ DataType*)(inputLengthsGm.GetPhyAddr()), (__gm__ DataType*)(targetLengthsGm.GetPhyAddr()),
410 (__gm__ T*)(negLogLikelihoodGm.GetPhyAddr()), (__gm__ T*)(logAlphaGm.GetPhyAddr()),410 (__gm__ T*)(negLogLikelihoodGm.GetPhyAddr()), (__gm__ T*)(logAlphaGm.GetPhyAddr()),
@@ -413,8 +413,8 @@ __aicore__ inline void CTCLossV2Grad<T, DataType, ThreadType>::Process()
413 tilingData_->symbolSet, tilingData_->zeroInfinity, tilingData_->BLANK, tilingData_->logAlphaT,413 tilingData_->symbolSet, tilingData_->zeroInfinity, tilingData_->BLANK, tilingData_->logAlphaT,
414 tilingData_->alphaLength, tilingData_->targetsDimNum, tilingData_->sDimRange);414 tilingData_->alphaLength, tilingData_->targetsDimNum, tilingData_->sDimRange);
415 SyncAll();415 SyncAll();
416- Simt::VF_CALL<CalGradCompute<T, DataType, ThreadType, THREAD_NUM_512>>(416+ asc_vf_call<CalGradCompute<T, DataType, ThreadType, THREAD_NUM_512>>(
417- Simt::Dim3(tilingData_->calGradThreadNum, 1), (__gm__ T*)(gradOutGm.GetPhyAddr()),417+ dim3(tilingData_->calGradThreadNum, 1), (__gm__ T*)(gradOutGm.GetPhyAddr()),
418 (__gm__ T*)(logProbsGm.GetPhyAddr()), (__gm__ DataType*)(targetsGm.GetPhyAddr()),418 (__gm__ T*)(logProbsGm.GetPhyAddr()), (__gm__ DataType*)(targetsGm.GetPhyAddr()),
419 (__gm__ DataType*)(inputLengthsGm.GetPhyAddr()), (__gm__ DataType*)(targetLengthsGm.GetPhyAddr()),419 (__gm__ DataType*)(inputLengthsGm.GetPhyAddr()), (__gm__ DataType*)(targetLengthsGm.GetPhyAddr()),
420 (__gm__ T*)(negLogLikelihoodGm.GetPhyAddr()), (__gm__ T*)(logAlphaGm.GetPhyAddr()),420 (__gm__ T*)(negLogLikelihoodGm.GetPhyAddr()), (__gm__ T*)(logAlphaGm.GetPhyAddr()),
Mloss/nll_loss/op_kernel/arch35/nll_loss_mix.h+37-36
@@ -20,6 +20,7 @@
20#include <cstdint>20#include <cstdint>
21#include "kernel_operator.h"21#include "kernel_operator.h"
22#include "../../inc/platform.h"22#include "../../inc/platform.h"
23+#include "simt_api/asc_simt.h"
23 24 
24using namespace AscendC;25using namespace AscendC;
25 26 
@@ -37,7 +38,7 @@ constexpr uint32_t FOUR_NUMBER = 4;
37__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtZeroPadTo256B(38__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtZeroPadTo256B(
38 const uint64_t start, const uint64_t padSize, __ubuf__ float* outBuffer)39 const uint64_t start, const uint64_t padSize, __ubuf__ float* outBuffer)
39{40{
40- for (uint32_t ubIdx = Simt::GetThreadIdx(); ubIdx < padSize; ubIdx += Simt::GetThreadNum()) {41+ for (uint32_t ubIdx = threadIdx.x; ubIdx < padSize; ubIdx += blockDim.x) {
41 outBuffer[start + ubIdx] = static_cast<float>(0);42 outBuffer[start + ubIdx] = static_cast<float>(0);
42 }43 }
43}44}
@@ -45,7 +46,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtZeroPadTo256B(
45__simt_vf__ __aicore__ LAUNCH_BOUND(MIN_THREAD) inline void GetReduceValueGm(46__simt_vf__ __aicore__ LAUNCH_BOUND(MIN_THREAD) inline void GetReduceValueGm(
46 int64_t blockId_, __ubuf__ float* srcBuffer, __gm__ volatile float* dstBuffer)47 int64_t blockId_, __ubuf__ float* srcBuffer, __gm__ volatile float* dstBuffer)
47{48{
48- if (Simt::GetThreadIdx() == 0) {49+ if (threadIdx.x == 0) {
49 dstBuffer[blockId_] = srcBuffer[0];50 dstBuffer[blockId_] = srcBuffer[0];
50 }51 }
51}52}
@@ -53,7 +54,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(MIN_THREAD) inline void GetReduceValueGm(
53__simt_vf__ __aicore__ LAUNCH_BOUND(MIN_THREAD) inline void GetReduceValue(54__simt_vf__ __aicore__ LAUNCH_BOUND(MIN_THREAD) inline void GetReduceValue(
54 uint32_t level1Idx, __ubuf__ float* srcBuffer, __ubuf__ float* dstBuffer)55 uint32_t level1Idx, __ubuf__ float* srcBuffer, __ubuf__ float* dstBuffer)
55{56{
56- if (Simt::GetThreadIdx() == 0) {57+ if (threadIdx.x == 0) {
57 dstBuffer[level1Idx] = srcBuffer[0];58 dstBuffer[level1Idx] = srcBuffer[0];
58 }59 }
59}60}
@@ -63,7 +64,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(MIN_THREAD) inline void GetRes(
63 uint32_t reduction_, __gm__ T* yGm_, __gm__ T* totalWeight_, __ubuf__ float* midOutResUB_,64 uint32_t reduction_, __gm__ T* yGm_, __gm__ T* totalWeight_, __ubuf__ float* midOutResUB_,
64 __ubuf__ float* midWeightResUB_)65 __ubuf__ float* midWeightResUB_)
65{66{
66- if (Simt::GetThreadIdx() == 0) {67+ if (threadIdx.x == 0) {
67 if (reduction_ == 1) {68 if (reduction_ == 1) {
68 yGm_[0] = static_cast<T>(midOutResUB_[0] / midWeightResUB_[0]);69 yGm_[0] = static_cast<T>(midOutResUB_[0] / midWeightResUB_[0]);
69 } else {70 } else {
@@ -78,7 +79,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtMeanModeCompute2
78 const uint32_t maxSize, const uint64_t offset, int64_t ignoreIndex_, uint32_t xDimC_, uint32_t isWeightPresent_,79 const uint32_t maxSize, const uint64_t offset, int64_t ignoreIndex_, uint32_t xDimC_, uint32_t isWeightPresent_,
79 __ubuf__ float* outBuffer, __ubuf__ float* weightbBuffer, __gm__ U* targetGm_, __gm__ T* xGm_, __gm__ T* weightGm_)80 __ubuf__ float* outBuffer, __ubuf__ float* weightbBuffer, __gm__ U* targetGm_, __gm__ T* xGm_, __gm__ T* weightGm_)
80{81{
81- for (uint32_t ubIdx = Simt::GetThreadIdx(); ubIdx < maxSize; ubIdx += Simt::GetThreadNum()) {82+ for (uint32_t ubIdx = threadIdx.x; ubIdx < maxSize; ubIdx += blockDim.x) {
82 uint64_t i = ubIdx + offset;83 uint64_t i = ubIdx + offset;
83 U targetIndex = targetGm_[i];84 U targetIndex = targetGm_[i];
84 if (targetIndex == ignoreIndex_) {85 if (targetIndex == ignoreIndex_) {
@@ -99,7 +100,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtMeanModeCompute4
99 int64_t xDimH_, int64_t xDimW_, uint32_t isWeightPresent_, __ubuf__ float* outBuffer, __ubuf__ float* weightbBuffer,100 int64_t xDimH_, int64_t xDimW_, uint32_t isWeightPresent_, __ubuf__ float* outBuffer, __ubuf__ float* weightbBuffer,
100 __gm__ U* targetGm_, __gm__ T* xGm_, __gm__ T* weightGm_, int64_t productOfCHW_, int64_t productOfHW_)101 __gm__ U* targetGm_, __gm__ T* xGm_, __gm__ T* weightGm_, int64_t productOfCHW_, int64_t productOfHW_)
101{102{
102- for (uint32_t ubIdx = Simt::GetThreadIdx(); ubIdx < maxSize; ubIdx += Simt::GetThreadNum()) {103+ for (uint32_t ubIdx = threadIdx.x; ubIdx < maxSize; ubIdx += blockDim.x) {
103 uint64_t i = ubIdx + offset;104 uint64_t i = ubIdx + offset;
104 U targetIndex = targetGm_[i];105 U targetIndex = targetGm_[i];
105 auto n = i / productOfHW_;106 auto n = i / productOfHW_;
@@ -298,18 +299,18 @@ public:
298 ReduceSum256(outBuffer, outBuffer, 0);299 ReduceSum256(outBuffer, outBuffer, 0);
299 PipeBarrier<PIPE_ALL>();300 PipeBarrier<PIPE_ALL>();
300 301 
301- AscendC::Simt::VF_CALL<GetReduceValueGm>(302+ asc_vf_call<GetReduceValueGm>(
302- AscendC::Simt::Dim3{static_cast<uint32_t>(MIN_THREAD)}, blockId_, (__ubuf__ float*)outBuffer.GetPhyAddr(),303+ dim3{static_cast<uint32_t>(MIN_THREAD)}, blockId_, (__ubuf__ float*)outBuffer.GetPhyAddr(),
303 (__gm__ volatile float*)workspace.GetPhyAddr());304 (__gm__ volatile float*)workspace.GetPhyAddr());
304 }305 }
305 306 
306 __aicore__ inline void MeanModeSubProcess(uint32_t& level1Idx, uint32_t& level2Idx, uint32_t& level3Idx)307 __aicore__ inline void MeanModeSubProcess(uint32_t& level1Idx, uint32_t& level2Idx, uint32_t& level3Idx)
307 {308 {
308- AscendC::Simt::VF_CALL<GetReduceValue>(309+ asc_vf_call<GetReduceValue>(
309- AscendC::Simt::Dim3{static_cast<uint32_t>(MIN_THREAD)}, level1Idx, (__ubuf__ float*)outUB_.GetPhyAddr(),310+ dim3{static_cast<uint32_t>(MIN_THREAD)}, level1Idx, (__ubuf__ float*)outUB_.GetPhyAddr(),
310 (__ubuf__ float*)level1OutUB_.GetPhyAddr());311 (__ubuf__ float*)level1OutUB_.GetPhyAddr());
311- AscendC::Simt::VF_CALL<GetReduceValue>(312+ asc_vf_call<GetReduceValue>(
312- AscendC::Simt::Dim3{static_cast<uint32_t>(MIN_THREAD)}, level1Idx, (__ubuf__ float*)weightUB_.GetPhyAddr(),313+ dim3{static_cast<uint32_t>(MIN_THREAD)}, level1Idx, (__ubuf__ float*)weightUB_.GetPhyAddr(),
313 (__ubuf__ float*)level1WeightUB_.GetPhyAddr());314 (__ubuf__ float*)level1WeightUB_.GetPhyAddr());
314 ++level1Idx;315 ++level1Idx;
315 PipeBarrier<PIPE_ALL>();316 PipeBarrier<PIPE_ALL>();
@@ -337,8 +338,8 @@ public:
337 338 
338 __aicore__ inline void SumModeSubProcess(uint32_t& level1Idx, uint32_t& level2Idx, uint32_t& level3Idx)339 __aicore__ inline void SumModeSubProcess(uint32_t& level1Idx, uint32_t& level2Idx, uint32_t& level3Idx)
339 {340 {
340- AscendC::Simt::VF_CALL<GetReduceValue>(341+ asc_vf_call<GetReduceValue>(
341- AscendC::Simt::Dim3{static_cast<uint32_t>(MIN_THREAD)}, level1Idx, (__ubuf__ float*)outUB_.GetPhyAddr(),342+ dim3{static_cast<uint32_t>(MIN_THREAD)}, level1Idx, (__ubuf__ float*)outUB_.GetPhyAddr(),
342 (__ubuf__ float*)level1OutUB_.GetPhyAddr());343 (__ubuf__ float*)level1OutUB_.GetPhyAddr());
343 ++level1Idx;344 ++level1Idx;
344 PipeBarrier<PIPE_ALL>();345 PipeBarrier<PIPE_ALL>();
@@ -385,17 +386,17 @@ public:
385 mainReduceLen = tailMainReduceSize_;386 mainReduceLen = tailMainReduceSize_;
386 tailLen = tailRemainSize_;387 tailLen = tailRemainSize_;
387 if (tailSize_ % MID_RES_128 != 0) {388 if (tailSize_ % MID_RES_128 != 0) {
388- AscendC::Simt::VF_CALL<SimtZeroPadTo256B>(389+ asc_vf_call<SimtZeroPadTo256B>(
389- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, padSize_,390+ dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, padSize_,
390 (__ubuf__ float*)outUB_.GetPhyAddr());391 (__ubuf__ float*)outUB_.GetPhyAddr());
391- AscendC::Simt::VF_CALL<SimtZeroPadTo256B>(392+ asc_vf_call<SimtZeroPadTo256B>(
392- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, padSize_,393+ dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, padSize_,
393 (__ubuf__ float*)weightUB_.GetPhyAddr());394 (__ubuf__ float*)weightUB_.GetPhyAddr());
394 }395 }
395 }396 }
396 397 
397- AscendC::Simt::VF_CALL<SimtMeanModeCompute2d<U, T>>(398+ asc_vf_call<SimtMeanModeCompute2d<U, T>>(
398- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, offset, ignoreIndex_, xDimC_,399+ dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, offset, ignoreIndex_, xDimC_,
399 isWeightPresent_, (__ubuf__ float*)outUB_.GetPhyAddr(), (__ubuf__ float*)weightUB_.GetPhyAddr(),400 isWeightPresent_, (__ubuf__ float*)outUB_.GetPhyAddr(), (__ubuf__ float*)weightUB_.GetPhyAddr(),
400 (__gm__ U*)targetGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ T*)weightGm_.GetPhyAddr());401 (__gm__ U*)targetGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ T*)weightGm_.GetPhyAddr());
401 402 
@@ -413,8 +414,8 @@ public:
413 simtMoveSize += mainReduceSize_;414 simtMoveSize += mainReduceSize_;
414 }415 }
415 416 
416- AscendC::Simt::VF_CALL<SimtMeanModeCompute2d<U, T>>(417+ asc_vf_call<SimtMeanModeCompute2d<U, T>>(
417- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, offset, ignoreIndex_, xDimC_,418+ dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, offset, ignoreIndex_, xDimC_,
418 isWeightPresent_, (__ubuf__ float*)outUB_.GetPhyAddr(), (__ubuf__ float*)weightUB_.GetPhyAddr(),419 isWeightPresent_, (__ubuf__ float*)outUB_.GetPhyAddr(), (__ubuf__ float*)weightUB_.GetPhyAddr(),
419 (__gm__ U*)targetGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ T*)weightGm_.GetPhyAddr());420 (__gm__ U*)targetGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ T*)weightGm_.GetPhyAddr());
420 421 
@@ -433,8 +434,8 @@ public:
433 simtMoveSize += mainReduceSize_;434 simtMoveSize += mainReduceSize_;
434 }435 }
435 436 
436- AscendC::Simt::VF_CALL<SimtMeanModeCompute2d<U, T>>(437+ asc_vf_call<SimtMeanModeCompute2d<U, T>>(
437- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, offset, ignoreIndex_, xDimC_,438+ dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, offset, ignoreIndex_, xDimC_,
438 isWeightPresent_, (__ubuf__ float*)outUB_.GetPhyAddr(), (__ubuf__ float*)weightUB_.GetPhyAddr(),439 isWeightPresent_, (__ubuf__ float*)outUB_.GetPhyAddr(), (__ubuf__ float*)weightUB_.GetPhyAddr(),
439 (__gm__ U*)targetGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ T*)weightGm_.GetPhyAddr());440 (__gm__ U*)targetGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ T*)weightGm_.GetPhyAddr());
440 441 
@@ -485,17 +486,17 @@ public:
485 mainReduceLen = tailMainReduceSize_;486 mainReduceLen = tailMainReduceSize_;
486 tailLen = tailRemainSize_;487 tailLen = tailRemainSize_;
487 if (tailSize_ % MID_RES_128 != 0) {488 if (tailSize_ % MID_RES_128 != 0) {
488- AscendC::Simt::VF_CALL<SimtZeroPadTo256B>(489+ asc_vf_call<SimtZeroPadTo256B>(
489- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, padSize_,490+ dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, padSize_,
490 (__ubuf__ float*)outUB_.GetPhyAddr());491 (__ubuf__ float*)outUB_.GetPhyAddr());
491- AscendC::Simt::VF_CALL<SimtZeroPadTo256B>(492+ asc_vf_call<SimtZeroPadTo256B>(
492- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, padSize_,493+ dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, padSize_,
493 (__ubuf__ float*)weightUB_.GetPhyAddr());494 (__ubuf__ float*)weightUB_.GetPhyAddr());
494 }495 }
495 }496 }
496 497 
497- AscendC::Simt::VF_CALL<SimtMeanModeCompute4d<U, T>>(498+ asc_vf_call<SimtMeanModeCompute4d<U, T>>(
498- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, offset, ignoreIndex_, xDimC_,499+ dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, offset, ignoreIndex_, xDimC_,
499 xDimN_, xDimH_, xDimW_, isWeightPresent_, (__ubuf__ float*)outUB_.GetPhyAddr(),500 xDimN_, xDimH_, xDimW_, isWeightPresent_, (__ubuf__ float*)outUB_.GetPhyAddr(),
500 (__ubuf__ float*)weightUB_.GetPhyAddr(), (__gm__ U*)targetGm_.GetPhyAddr(),501 (__ubuf__ float*)weightUB_.GetPhyAddr(), (__gm__ U*)targetGm_.GetPhyAddr(),
501 (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ T*)weightGm_.GetPhyAddr(), productOfCHW_, productOfHW_);502 (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ T*)weightGm_.GetPhyAddr(), productOfCHW_, productOfHW_);
@@ -514,8 +515,8 @@ public:
514 simtMoveSize += mainReduceSize_;515 simtMoveSize += mainReduceSize_;
515 }516 }
516 517 
517- AscendC::Simt::VF_CALL<SimtMeanModeCompute4d<U, T>>(518+ asc_vf_call<SimtMeanModeCompute4d<U, T>>(
518- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, offset, ignoreIndex_, xDimC_,519+ dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, offset, ignoreIndex_, xDimC_,
519 xDimN_, xDimH_, xDimW_, isWeightPresent_, (__ubuf__ float*)outUB_.GetPhyAddr(),520 xDimN_, xDimH_, xDimW_, isWeightPresent_, (__ubuf__ float*)outUB_.GetPhyAddr(),
520 (__ubuf__ float*)weightUB_.GetPhyAddr(), (__gm__ U*)targetGm_.GetPhyAddr(),521 (__ubuf__ float*)weightUB_.GetPhyAddr(), (__gm__ U*)targetGm_.GetPhyAddr(),
521 (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ T*)weightGm_.GetPhyAddr(), productOfCHW_, productOfHW_);522 (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ T*)weightGm_.GetPhyAddr(), productOfCHW_, productOfHW_);
@@ -535,8 +536,8 @@ public:
535 simtMoveSize += mainReduceSize_;536 simtMoveSize += mainReduceSize_;
536 }537 }
537 538 
538- AscendC::Simt::VF_CALL<SimtMeanModeCompute4d<U, T>>(539+ asc_vf_call<SimtMeanModeCompute4d<U, T>>(
539- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, offset, ignoreIndex_, xDimC_,540+ dim3{static_cast<uint32_t>(THREAD_DIM)}, simtMoveSize, offset, ignoreIndex_, xDimC_,
540 xDimN_, xDimH_, xDimW_, isWeightPresent_, (__ubuf__ float*)outUB_.GetPhyAddr(),541 xDimN_, xDimH_, xDimW_, isWeightPresent_, (__ubuf__ float*)outUB_.GetPhyAddr(),
541 (__ubuf__ float*)weightUB_.GetPhyAddr(), (__gm__ U*)targetGm_.GetPhyAddr(),542 (__ubuf__ float*)weightUB_.GetPhyAddr(), (__gm__ U*)targetGm_.GetPhyAddr(),
542 (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ T*)weightGm_.GetPhyAddr(), productOfCHW_, productOfHW_);543 (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ T*)weightGm_.GetPhyAddr(), productOfCHW_, productOfHW_);
@@ -587,8 +588,8 @@ public:
587 ReduceSumAmongCores(midOutResUB_, level2OutUB_, workspaceOutGm_);588 ReduceSumAmongCores(midOutResUB_, level2OutUB_, workspaceOutGm_);
588 ReduceSumAmongCores(midWeightResUB_, level3OutUB_, workspaceWeightGm_);589 ReduceSumAmongCores(midWeightResUB_, level3OutUB_, workspaceWeightGm_);
589 PipeBarrier<PIPE_ALL>();590 PipeBarrier<PIPE_ALL>();
590- AscendC::Simt::VF_CALL<GetRes<T>>(591+ asc_vf_call<GetRes<T>>(
591- AscendC::Simt::Dim3{static_cast<uint32_t>(MIN_THREAD)}, reduction_, (__gm__ T*)yGm_.GetPhyAddr(),592+ dim3{static_cast<uint32_t>(MIN_THREAD)}, reduction_, (__gm__ T*)yGm_.GetPhyAddr(),
592 (__gm__ T*)totalWeight_.GetPhyAddr(), (__ubuf__ float*)midOutResUB_.GetPhyAddr(),593 (__gm__ T*)totalWeight_.GetPhyAddr(), (__ubuf__ float*)midOutResUB_.GetPhyAddr(),
593 (__ubuf__ float*)midWeightResUB_.GetPhyAddr());594 (__ubuf__ float*)midWeightResUB_.GetPhyAddr());
594 }595 }
Mloss/nll_loss/op_kernel/arch35/nll_loss_simt.h+70-67
@@ -19,6 +19,9 @@
19#include <cmath>19#include <cmath>
20#include <cstdint>20#include <cstdint>
21#include "kernel_operator.h"21#include "kernel_operator.h"
22+#include "simt_api/asc_simt.h"
23+#include "simt_api/asc_bf16.h"
24+#include "simt_api/asc_fp16.h"
22 25 
23using namespace AscendC;26using namespace AscendC;
24 27 
@@ -50,16 +53,16 @@ __simt_callee__ __aicore__ inline void TargetCheck(const U idx, const uint32_t x
50 53 
51__simt_callee__ __aicore__ inline void SimtComputeBinaryReduction(__ubuf__ float* tmpOut_, __ubuf__ float* tmpWeight_)54__simt_callee__ __aicore__ inline void SimtComputeBinaryReduction(__ubuf__ float* tmpOut_, __ubuf__ float* tmpWeight_)
52{55{
53- uint32_t countBR = Simt::GetThreadNum();56+ uint32_t countBR = blockDim.x;
54 while (countBR > 1) {57 while (countBR > 1) {
55 uint32_t halfBR = (countBR + 1) / BINARY_HALF;58 uint32_t halfBR = (countBR + 1) / BINARY_HALF;
56- if (Simt::GetThreadIdx() < halfBR) {59+ if (threadIdx.x < halfBR) {
57- tmpOut_[Simt::GetThreadIdx()] = tmpOut_[Simt::GetThreadIdx() + halfBR] + tmpOut_[Simt::GetThreadIdx()];60+ tmpOut_[threadIdx.x] = tmpOut_[threadIdx.x + halfBR] + tmpOut_[threadIdx.x];
58 61 
59- tmpWeight_[Simt::GetThreadIdx()] =62+ tmpWeight_[threadIdx.x] =
60- tmpWeight_[Simt::GetThreadIdx() + halfBR] + tmpWeight_[Simt::GetThreadIdx()];63+ tmpWeight_[threadIdx.x + halfBR] + tmpWeight_[threadIdx.x];
61 }64 }
62- Simt::ThreadBarrier();65+ asc_syncthreads();
63 countBR = halfBR;66 countBR = halfBR;
64 }67 }
65}68}
@@ -100,8 +103,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeNLLLoss2d
100 int64_t blockId_, int64_t xDimN_, int64_t blockNums_, __gm__ U* targetGm_, int64_t ignoreIndex_, __gm__ T* yGm_,103 int64_t blockId_, int64_t xDimN_, int64_t blockNums_, __gm__ U* targetGm_, int64_t ignoreIndex_, __gm__ T* yGm_,
101 uint32_t isWeightPresent_, __gm__ T* weightGm_, __gm__ T* xGm_, uint32_t xDimC_)104 uint32_t isWeightPresent_, __gm__ T* weightGm_, __gm__ T* xGm_, uint32_t xDimC_)
102{105{
103- for (int64_t i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < xDimN_;106+ for (int64_t i = blockId_ * blockDim.x + threadIdx.x; i < xDimN_;
104- i = i + blockNums_ * Simt::GetThreadNum()) {107+ i = i + blockNums_ * blockDim.x) {
105 U targetIndex = targetGm_[i];108 U targetIndex = targetGm_[i];
106 if (targetIndex == ignoreIndex_) {109 if (targetIndex == ignoreIndex_) {
107 yGm_[i] = static_cast<T>(0);110 yGm_[i] = static_cast<T>(0);
@@ -119,10 +122,10 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeNLLLoss2d
119 uint32_t isWeightPresent_, __gm__ T* weightGm_, __ubuf__ float* tmpOut_, __gm__ T* xGm_, __ubuf__ float* tmpWeight_,122 uint32_t isWeightPresent_, __gm__ T* weightGm_, __ubuf__ float* tmpOut_, __gm__ T* xGm_, __ubuf__ float* tmpWeight_,
120 __gm__ volatile float* tmpSumGm_, __gm__ volatile float* tmpWeightGm_)123 __gm__ volatile float* tmpSumGm_, __gm__ volatile float* tmpWeightGm_)
121{124{
122- for (int64_t i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < xDimN_;125+ for (int64_t i = blockId_ * blockDim.x + threadIdx.x; i < xDimN_;
123- i = i + blockNums_ * Simt::GetThreadNum()) {126+ i = i + blockNums_ * blockDim.x) {
124 U targetIndex = targetGm_[i];127 U targetIndex = targetGm_[i];
125- int64_t ubIdx = (i % (blockNums_ * Simt::GetThreadNum())) % THREAD_DIM;128+ int64_t ubIdx = (i % (blockNums_ * blockDim.x)) % THREAD_DIM;
126 if (targetIndex == ignoreIndex_) {129 if (targetIndex == ignoreIndex_) {
127 continue;130 continue;
128 }131 }
@@ -132,10 +135,10 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeNLLLoss2d
132 tmpOut_[ubIdx] - static_cast<float>(curWeight) * static_cast<float>(xGm_[i * xDimC_ + targetIndex]);135 tmpOut_[ubIdx] - static_cast<float>(curWeight) * static_cast<float>(xGm_[i * xDimC_ + targetIndex]);
133 tmpWeight_[ubIdx] = static_cast<float>(curWeight) + tmpWeight_[ubIdx];136 tmpWeight_[ubIdx] = static_cast<float>(curWeight) + tmpWeight_[ubIdx];
134 }137 }
135- Simt::ThreadBarrier();138+ asc_syncthreads();
136 SimtComputeBinaryReduction(tmpOut_, tmpWeight_);139 SimtComputeBinaryReduction(tmpOut_, tmpWeight_);
137 140 
138- if (Simt::GetThreadIdx() == 0) {141+ if (threadIdx.x == 0) {
139 tmpSumGm_[blockId_] = tmpOut_[0];142 tmpSumGm_[blockId_] = tmpOut_[0];
140 tmpWeightGm_[blockId_] = tmpWeight_[0];143 tmpWeightGm_[blockId_] = tmpWeight_[0];
141 }144 }
@@ -147,10 +150,10 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeNLLLoss2d
147 __ubuf__ float* tmpOut_, __ubuf__ float* tmpWeight_, uint32_t isWeightPresent_, __gm__ T* weightGm_, __gm__ T* xGm_,150 __ubuf__ float* tmpOut_, __ubuf__ float* tmpWeight_, uint32_t isWeightPresent_, __gm__ T* weightGm_, __gm__ T* xGm_,
148 __gm__ volatile float* tmpSumGm_, uint32_t xDimC_, __gm__ volatile float* tmpWeightGm_)151 __gm__ volatile float* tmpSumGm_, uint32_t xDimC_, __gm__ volatile float* tmpWeightGm_)
149{152{
150- for (int64_t i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < xDimN_;153+ for (int64_t i = blockId_ * blockDim.x + threadIdx.x; i < xDimN_;
151- i = i + blockNums_ * Simt::GetThreadNum()) {154+ i = i + blockNums_ * blockDim.x) {
152 U targetIndex = targetGm_[i];155 U targetIndex = targetGm_[i];
153- int64_t ubIdx = (i % (blockNums_ * Simt::GetThreadNum())) % THREAD_DIM;156+ int64_t ubIdx = (i % (blockNums_ * blockDim.x)) % THREAD_DIM;
154 if (targetIndex == ignoreIndex_) {157 if (targetIndex == ignoreIndex_) {
155 continue;158 continue;
156 }159 }
@@ -160,10 +163,10 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeNLLLoss2d
160 tmpOut_[ubIdx] - static_cast<float>(curWeight) * static_cast<float>(xGm_[i * xDimC_ + targetIndex]);163 tmpOut_[ubIdx] - static_cast<float>(curWeight) * static_cast<float>(xGm_[i * xDimC_ + targetIndex]);
161 tmpWeight_[ubIdx] = static_cast<float>(curWeight) + tmpWeight_[ubIdx];164 tmpWeight_[ubIdx] = static_cast<float>(curWeight) + tmpWeight_[ubIdx];
162 }165 }
163- Simt::ThreadBarrier();166+ asc_syncthreads();
164 SimtComputeBinaryReduction(tmpOut_, tmpWeight_);167 SimtComputeBinaryReduction(tmpOut_, tmpWeight_);
165 168 
166- if (Simt::GetThreadIdx() == 0) {169+ if (threadIdx.x == 0) {
167 tmpSumGm_[blockId_] = tmpOut_[0];170 tmpSumGm_[blockId_] = tmpOut_[0];
168 tmpWeightGm_[blockId_] = tmpWeight_[0];171 tmpWeightGm_[blockId_] = tmpWeight_[0];
169 }172 }
@@ -175,7 +178,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeNLLLoss4d
175 __gm__ T* weightGm, __gm__ T* xGm, uint32_t xDimC, int64_t dealingNum, int64_t productOfCHW,178 __gm__ T* weightGm, __gm__ T* xGm, uint32_t xDimC, int64_t dealingNum, int64_t productOfCHW,
176 int64_t productOfHW)179 int64_t productOfHW)
177{180{
178- for (int64_t i = Simt::GetThreadIdx(); i < dealingNum; i = i + Simt::GetThreadNum()) {181+ for (int64_t i = threadIdx.x; i < dealingNum; i = i + blockDim.x) {
179 int64_t index = offset + i;182 int64_t index = offset + i;
180 U targetIndex = targetGm[index]; //blockid_*183 U targetIndex = targetGm[index]; //blockid_*
181 if (targetIndex == ignoreIndex) {184 if (targetIndex == ignoreIndex) {
@@ -198,10 +201,10 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeNLLLoss4d
198 __gm__ T* xGm_, __ubuf__ float* tmpWeight_, __gm__ volatile float* tmpSumGm_, __gm__ volatile float* tmpWeightGm_,201 __gm__ T* xGm_, __ubuf__ float* tmpWeight_, __gm__ volatile float* tmpSumGm_, __gm__ volatile float* tmpWeightGm_,
199 int64_t productOfNHW_, int64_t productOfCHW_, int64_t productOfHW_)202 int64_t productOfNHW_, int64_t productOfCHW_, int64_t productOfHW_)
200{203{
201- for (int64_t i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < productOfNHW_;204+ for (int64_t i = blockId_ * blockDim.x + threadIdx.x; i < productOfNHW_;
202- i = i + blockNums_ * Simt::GetThreadNum()) {205+ i = i + blockNums_ * blockDim.x) {
203 U targetIndex = targetGm_[i];206 U targetIndex = targetGm_[i];
204- int64_t ubIdx = (i % (blockNums_ * Simt::GetThreadNum())) % THREAD_DIM;207+ int64_t ubIdx = (i % (blockNums_ * blockDim.x)) % THREAD_DIM;
205 if (targetIndex == ignoreIndex_) {208 if (targetIndex == ignoreIndex_) {
206 continue;209 continue;
207 }210 }
@@ -214,10 +217,10 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeNLLLoss4d
214 static_cast<float>(xGm_[n * productOfCHW_ + targetIndex * productOfHW_ + rem]);217 static_cast<float>(xGm_[n * productOfCHW_ + targetIndex * productOfHW_ + rem]);
215 tmpWeight_[ubIdx] = static_cast<float>(curWeight) + tmpWeight_[ubIdx];218 tmpWeight_[ubIdx] = static_cast<float>(curWeight) + tmpWeight_[ubIdx];
216 }219 }
217- Simt::ThreadBarrier();220+ asc_syncthreads();
218 SimtComputeBinaryReduction(tmpOut_, tmpWeight_);221 SimtComputeBinaryReduction(tmpOut_, tmpWeight_);
219 222 
220- if (Simt::GetThreadIdx() == 0) {223+ if (threadIdx.x == 0) {
221 tmpSumGm_[blockId_] = tmpOut_[0];224 tmpSumGm_[blockId_] = tmpOut_[0];
222 tmpWeightGm_[blockId_] = tmpWeight_[0];225 tmpWeightGm_[blockId_] = tmpWeight_[0];
223 }226 }
@@ -230,10 +233,10 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeNLLLoss4d
230 __gm__ T* weightGm_, __gm__ T* xGm_, __gm__ volatile float* tmpSumGm_, uint32_t xDimC_,233 __gm__ T* weightGm_, __gm__ T* xGm_, __gm__ volatile float* tmpSumGm_, uint32_t xDimC_,
231 __gm__ volatile float* tmpWeightGm_, int64_t productOfNHW_, int64_t productOfCHW_, int64_t productOfHW_)234 __gm__ volatile float* tmpWeightGm_, int64_t productOfNHW_, int64_t productOfCHW_, int64_t productOfHW_)
232{235{
233- for (int64_t i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < productOfNHW_;236+ for (int64_t i = blockId_ * blockDim.x + threadIdx.x; i < productOfNHW_;
234- i = i + blockNums_ * Simt::GetThreadNum()) {237+ i = i + blockNums_ * blockDim.x) {
235 U targetIndex = targetGm_[i];238 U targetIndex = targetGm_[i];
236- int64_t ubIdx = (i % (blockNums_ * Simt::GetThreadNum())) % THREAD_DIM;239+ int64_t ubIdx = (i % (blockNums_ * blockDim.x)) % THREAD_DIM;
237 if (targetIndex == ignoreIndex_) {240 if (targetIndex == ignoreIndex_) {
238 continue;241 continue;
239 }242 }
@@ -246,10 +249,10 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeNLLLoss4d
246 static_cast<float>(xGm_[n * productOfCHW_ + targetIndex * productOfHW_ + rem]);249 static_cast<float>(xGm_[n * productOfCHW_ + targetIndex * productOfHW_ + rem]);
247 tmpWeight_[ubIdx] = static_cast<float>(curWeight) + tmpWeight_[ubIdx];250 tmpWeight_[ubIdx] = static_cast<float>(curWeight) + tmpWeight_[ubIdx];
248 }251 }
249- Simt::ThreadBarrier();252+ asc_syncthreads();
250 SimtComputeBinaryReduction(tmpOut_, tmpWeight_);253 SimtComputeBinaryReduction(tmpOut_, tmpWeight_);
251 254 
252- if (Simt::GetThreadIdx() == 0) {255+ if (threadIdx.x == 0) {
253 tmpSumGm_[blockId_] = tmpOut_[0];256 tmpSumGm_[blockId_] = tmpOut_[0];
254 tmpWeightGm_[blockId_] = tmpWeight_[0];257 tmpWeightGm_[blockId_] = tmpWeight_[0];
255 }258 }
@@ -261,27 +264,27 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void CalBinaryReductionIn
261 __gm__ T* yGm_, __gm__ T* totalWeight_)264 __gm__ T* yGm_, __gm__ T* totalWeight_)
262{265{
263 uint32_t valueLarger2Power = blockNums_ - calNumPower_;266 uint32_t valueLarger2Power = blockNums_ - calNumPower_;
264- if (Simt::GetThreadIdx() < valueLarger2Power) {267+ if (threadIdx.x < valueLarger2Power) {
265- tmpSumGm_[Simt::GetThreadIdx()] =268+ tmpSumGm_[threadIdx.x] =
266- tmpSumGm_[Simt::GetThreadIdx() + calNumPower_] + tmpSumGm_[Simt::GetThreadIdx()];269+ tmpSumGm_[threadIdx.x + calNumPower_] + tmpSumGm_[threadIdx.x];
267- tmpWeightGm_[Simt::GetThreadIdx()] =270+ tmpWeightGm_[threadIdx.x] =
268- tmpWeightGm_[Simt::GetThreadIdx() + calNumPower_] + tmpWeightGm_[Simt::GetThreadIdx()];271+ tmpWeightGm_[threadIdx.x + calNumPower_] + tmpWeightGm_[threadIdx.x];
269 }272 }
270- Simt::ThreadBarrier();273+ asc_syncthreads();
271 274 
272 int64_t countBR = calNumPower_;275 int64_t countBR = calNumPower_;
273 while (countBR > 1) {276 while (countBR > 1) {
274 int64_t halfBR = (countBR + 1) / BINARY_HALF;277 int64_t halfBR = (countBR + 1) / BINARY_HALF;
275- if (Simt::GetThreadIdx() < halfBR) {278+ if (threadIdx.x < halfBR) {
276- tmpSumGm_[Simt::GetThreadIdx()] =279+ tmpSumGm_[threadIdx.x] =
277- tmpSumGm_[Simt::GetThreadIdx() + halfBR] + tmpSumGm_[Simt::GetThreadIdx()];280+ tmpSumGm_[threadIdx.x + halfBR] + tmpSumGm_[threadIdx.x];
278- tmpWeightGm_[Simt::GetThreadIdx()] =281+ tmpWeightGm_[threadIdx.x] =
279- tmpWeightGm_[Simt::GetThreadIdx() + halfBR] + tmpWeightGm_[Simt::GetThreadIdx()];282+ tmpWeightGm_[threadIdx.x + halfBR] + tmpWeightGm_[threadIdx.x];
280 }283 }
281- Simt::ThreadBarrier();284+ asc_syncthreads();
282 countBR = halfBR;285 countBR = halfBR;
283 }286 }
284- if (Simt::GetThreadIdx() == 0) {287+ if (threadIdx.x == 0) {
285 if constexpr (is_same<bfloat16_t, T>::value) {288 if constexpr (is_same<bfloat16_t, T>::value) {
286 if constexpr (reduction == REDUCTION_MEAN) {289 if constexpr (reduction == REDUCTION_MEAN) {
287 yGm_[0] = Simt::Cast<bfloat16_t, float, RoundMode::CAST_EVEN>((tmpSumGm_[0] / tmpWeightGm_[0]));290 yGm_[0] = Simt::Cast<bfloat16_t, float, RoundMode::CAST_EVEN>((tmpSumGm_[0] / tmpWeightGm_[0]));
@@ -418,8 +421,8 @@ public:
418 for (int32_t i = 0; i < loopTimes_; i++) {421 for (int32_t i = 0; i < loopTimes_; i++) {
419 LocalTensor<T> outUb = outQueue_.AllocTensor<T>();422 LocalTensor<T> outUb = outQueue_.AllocTensor<T>();
420 int64_t offset = offsetGmY_ + i * dealingNumOnce_;423 int64_t offset = offsetGmY_ + i * dealingNumOnce_;
421- AscendC::Simt::VF_CALL<SimtComputeNLLLoss4dNone<U, T>>(424+ asc_vf_call<SimtComputeNLLLoss4dNone<U, T>>(
422- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)},425+ dim3{static_cast<uint32_t>(THREAD_DIM)},
423 (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, (__ubuf__ T*)outUb.GetPhyAddr(), offset, isWeightPresent_,426 (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, (__ubuf__ T*)outUb.GetPhyAddr(), offset, isWeightPresent_,
424 (__gm__ T*)weightGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), xDimC_, dealingNumOnce_, productOfCHW_, productOfHW_);427 (__gm__ T*)weightGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), xDimC_, dealingNumOnce_, productOfCHW_, productOfHW_);
425 outQueue_.EnQue(outUb);428 outQueue_.EnQue(outUb);
@@ -428,8 +431,8 @@ public:
428 if (dealingNumTail_ != 0) {431 if (dealingNumTail_ != 0) {
429 LocalTensor<T> outUb = outQueue_.AllocTensor<T>();432 LocalTensor<T> outUb = outQueue_.AllocTensor<T>();
430 int64_t offset = offsetGmY_ + loopTimes_ * dealingNumOnce_;433 int64_t offset = offsetGmY_ + loopTimes_ * dealingNumOnce_;
431- AscendC::Simt::VF_CALL<SimtComputeNLLLoss4dNone<U, T>>(434+ asc_vf_call<SimtComputeNLLLoss4dNone<U, T>>(
432- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)},435+ dim3{static_cast<uint32_t>(THREAD_DIM)},
433 (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, (__ubuf__ T*)outUb.GetPhyAddr(), offset, isWeightPresent_,436 (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, (__ubuf__ T*)outUb.GetPhyAddr(), offset, isWeightPresent_,
434 (__gm__ T*)weightGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), xDimC_, dealingNumTail_, productOfCHW_, productOfHW_);437 (__gm__ T*)weightGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), xDimC_, dealingNumTail_, productOfCHW_, productOfHW_);
435 outQueue_.EnQue(outUb);438 outQueue_.EnQue(outUb);
@@ -440,45 +443,45 @@ public:
440 __aicore__ inline void Process()443 __aicore__ inline void Process()
441 {444 {
442 if constexpr (xDims == 1) {445 if constexpr (xDims == 1) {
443- AscendC::Simt::VF_CALL<SimtComputeNLLLoss1d<U, T, reduction>>(446+ asc_vf_call<SimtComputeNLLLoss1d<U, T, reduction>>(
444- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_,447+ dim3{static_cast<uint32_t>(THREAD_DIM)}, (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_,
445 (__gm__ T*)yGm_.GetPhyAddr(), (__gm__ T*)totalWeight_.GetPhyAddr(), isWeightPresent_,448 (__gm__ T*)yGm_.GetPhyAddr(), (__gm__ T*)totalWeight_.GetPhyAddr(), isWeightPresent_,
446 (__gm__ T*)weightGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), xDimC_);449 (__gm__ T*)weightGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), xDimC_);
447 } else if constexpr (xDims == NUMBER_TWO) {450 } else if constexpr (xDims == NUMBER_TWO) {
448 if (reduction == 0) {451 if (reduction == 0) {
449- AscendC::Simt::VF_CALL<SimtComputeNLLLoss2dNone<U, T>>(452+ asc_vf_call<SimtComputeNLLLoss2dNone<U, T>>(
450- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, blockId_, xDimN_, blockNums_,453+ dim3{static_cast<uint32_t>(THREAD_DIM)}, blockId_, xDimN_, blockNums_,
451 (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, (__gm__ T*)yGm_.GetPhyAddr(), isWeightPresent_,454 (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, (__gm__ T*)yGm_.GetPhyAddr(), isWeightPresent_,
452 (__gm__ T*)weightGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), xDimC_);455 (__gm__ T*)weightGm_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(), xDimC_);
453 } else if constexpr (reduction == REDUCTION_SUM) {456 } else if constexpr (reduction == REDUCTION_SUM) {
454 Duplicate(tmpOut_, 0.0f, THREAD_DIM);457 Duplicate(tmpOut_, 0.0f, THREAD_DIM);
455 Duplicate(tmpWeight_, 0.0f, THREAD_DIM);458 Duplicate(tmpWeight_, 0.0f, THREAD_DIM);
456- AscendC::Simt::VF_CALL<SimtComputeNLLLoss2dSum<U, T>>(459+ asc_vf_call<SimtComputeNLLLoss2dSum<U, T>>(
457- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, blockId_, xDimN_, blockNums_,460+ dim3{static_cast<uint32_t>(THREAD_DIM)}, blockId_, xDimN_, blockNums_,
458 (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, xDimC_, isWeightPresent_,461 (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, xDimC_, isWeightPresent_,
459 (__gm__ T*)weightGm_.GetPhyAddr(), (__ubuf__ float*)tmpOut_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(),462 (__gm__ T*)weightGm_.GetPhyAddr(), (__ubuf__ float*)tmpOut_.GetPhyAddr(), (__gm__ T*)xGm_.GetPhyAddr(),
460 (__ubuf__ float*)tmpWeight_.GetPhyAddr(), (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(),463 (__ubuf__ float*)tmpWeight_.GetPhyAddr(), (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(),
461 (__gm__ volatile float*)tmpWeightGm_.GetPhyAddr());464 (__gm__ volatile float*)tmpWeightGm_.GetPhyAddr());
462 SyncAll();465 SyncAll();
463 if (blockId_ == 0) {466 if (blockId_ == 0) {
464- AscendC::Simt::VF_CALL<CalBinaryReductionInterBlock<U, T, reduction>>(467+ asc_vf_call<CalBinaryReductionInterBlock<U, T, reduction>>(
465- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, blockNums_, calNumPower_,468+ dim3{static_cast<uint32_t>(THREAD_DIM)}, blockNums_, calNumPower_,
466 (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(), (__gm__ volatile float*)tmpWeightGm_.GetPhyAddr(),469 (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(), (__gm__ volatile float*)tmpWeightGm_.GetPhyAddr(),
467 (__gm__ T*)yGm_.GetPhyAddr(), (__gm__ T*)totalWeight_.GetPhyAddr());470 (__gm__ T*)yGm_.GetPhyAddr(), (__gm__ T*)totalWeight_.GetPhyAddr());
468 }471 }
469 } else if constexpr (reduction == REDUCTION_MEAN) {472 } else if constexpr (reduction == REDUCTION_MEAN) {
470 Duplicate(tmpOut_, 0.0f, THREAD_DIM);473 Duplicate(tmpOut_, 0.0f, THREAD_DIM);
471 Duplicate(tmpWeight_, 0.0f, THREAD_DIM);474 Duplicate(tmpWeight_, 0.0f, THREAD_DIM);
472- AscendC::Simt::VF_CALL<SimtComputeNLLLoss2dMean<U, T>>(475+ asc_vf_call<SimtComputeNLLLoss2dMean<U, T>>(
473- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, blockId_, xDimN_, blockNums_,476+ dim3{static_cast<uint32_t>(THREAD_DIM)}, blockId_, xDimN_, blockNums_,
474 (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, (__ubuf__ float*)tmpOut_.GetPhyAddr(),477 (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, (__ubuf__ float*)tmpOut_.GetPhyAddr(),
475 (__ubuf__ float*)tmpWeight_.GetPhyAddr(), isWeightPresent_, (__gm__ T*)weightGm_.GetPhyAddr(),478 (__ubuf__ float*)tmpWeight_.GetPhyAddr(), isWeightPresent_, (__gm__ T*)weightGm_.GetPhyAddr(),
476 (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(), xDimC_,479 (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(), xDimC_,
477 (__gm__ volatile float*)tmpWeightGm_.GetPhyAddr());480 (__gm__ volatile float*)tmpWeightGm_.GetPhyAddr());
478 SyncAll();481 SyncAll();
479 if (blockId_ == 0) {482 if (blockId_ == 0) {
480- AscendC::Simt::VF_CALL<CalBinaryReductionInterBlock<U, T, reduction>>(483+ asc_vf_call<CalBinaryReductionInterBlock<U, T, reduction>>(
481- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, blockNums_, calNumPower_,484+ dim3{static_cast<uint32_t>(THREAD_DIM)}, blockNums_, calNumPower_,
482 (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(), (__gm__ volatile float*)tmpWeightGm_.GetPhyAddr(),485 (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(), (__gm__ volatile float*)tmpWeightGm_.GetPhyAddr(),
483 (__gm__ T*)yGm_.GetPhyAddr(), (__gm__ T*)totalWeight_.GetPhyAddr());486 (__gm__ T*)yGm_.GetPhyAddr(), (__gm__ T*)totalWeight_.GetPhyAddr());
484 }487 }
@@ -492,8 +495,8 @@ public:
492 495 
493 Duplicate(tmpOut_, 0.0f, THREAD_DIM);496 Duplicate(tmpOut_, 0.0f, THREAD_DIM);
494 Duplicate(tmpWeight_, 0.0f, THREAD_DIM);497 Duplicate(tmpWeight_, 0.0f, THREAD_DIM);
495- AscendC::Simt::VF_CALL<SimtComputeNLLLoss4dSum<U, T>>(498+ asc_vf_call<SimtComputeNLLLoss4dSum<U, T>>(
496- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, blockId_, xDimN_, xDimH_, xDimW_,499+ dim3{static_cast<uint32_t>(THREAD_DIM)}, blockId_, xDimN_, xDimH_, xDimW_,
497 blockNums_, (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, xDimC_, isWeightPresent_,500 blockNums_, (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, xDimC_, isWeightPresent_,
498 (__gm__ T*)weightGm_.GetPhyAddr(), (__ubuf__ float*)tmpOut_.GetPhyAddr(),501 (__gm__ T*)weightGm_.GetPhyAddr(), (__ubuf__ float*)tmpOut_.GetPhyAddr(),
499 (__gm__ T*)xGm_.GetPhyAddr(), (__ubuf__ float*)tmpWeight_.GetPhyAddr(),502 (__gm__ T*)xGm_.GetPhyAddr(), (__ubuf__ float*)tmpWeight_.GetPhyAddr(),
@@ -501,16 +504,16 @@ public:
501 productOfNHW_, productOfCHW_, productOfHW_);504 productOfNHW_, productOfCHW_, productOfHW_);
502 SyncAll();505 SyncAll();
503 if (blockId_ == 0) {506 if (blockId_ == 0) {
504- AscendC::Simt::VF_CALL<CalBinaryReductionInterBlock<U, T, reduction>>(507+ asc_vf_call<CalBinaryReductionInterBlock<U, T, reduction>>(
505- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, blockNums_, calNumPower_,508+ dim3{static_cast<uint32_t>(THREAD_DIM)}, blockNums_, calNumPower_,
506 (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(), (__gm__ volatile float*)tmpWeightGm_.GetPhyAddr(),509 (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(), (__gm__ volatile float*)tmpWeightGm_.GetPhyAddr(),
507 (__gm__ T*)yGm_.GetPhyAddr(), (__gm__ T*)totalWeight_.GetPhyAddr());510 (__gm__ T*)yGm_.GetPhyAddr(), (__gm__ T*)totalWeight_.GetPhyAddr());
508 }511 }
509 } else if constexpr (reduction == REDUCTION_MEAN) {512 } else if constexpr (reduction == REDUCTION_MEAN) {
510 Duplicate(tmpOut_, 0.0f, THREAD_DIM);513 Duplicate(tmpOut_, 0.0f, THREAD_DIM);
511 Duplicate(tmpWeight_, 0.0f, THREAD_DIM);514 Duplicate(tmpWeight_, 0.0f, THREAD_DIM);
512- AscendC::Simt::VF_CALL<SimtComputeNLLLoss4dMean<U, T>>(515+ asc_vf_call<SimtComputeNLLLoss4dMean<U, T>>(
513- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, blockId_, xDimN_, xDimH_, xDimW_, blockNums_,516+ dim3{static_cast<uint32_t>(THREAD_DIM)}, blockId_, xDimN_, xDimH_, xDimW_, blockNums_,
514 (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, (__ubuf__ float*)tmpOut_.GetPhyAddr(),517 (__gm__ U*)targetGm_.GetPhyAddr(), ignoreIndex_, (__ubuf__ float*)tmpOut_.GetPhyAddr(),
515 (__ubuf__ float*)tmpWeight_.GetPhyAddr(), isWeightPresent_, (__gm__ T*)weightGm_.GetPhyAddr(),518 (__ubuf__ float*)tmpWeight_.GetPhyAddr(), isWeightPresent_, (__gm__ T*)weightGm_.GetPhyAddr(),
516 (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(), xDimC_,519 (__gm__ T*)xGm_.GetPhyAddr(), (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(), xDimC_,
@@ -518,8 +521,8 @@ public:
518 521 
519 SyncAll();522 SyncAll();
520 if (blockId_ == 0) {523 if (blockId_ == 0) {
521- AscendC::Simt::VF_CALL<CalBinaryReductionInterBlock<U, T, reduction>>(524+ asc_vf_call<CalBinaryReductionInterBlock<U, T, reduction>>(
522- AscendC::Simt::Dim3{static_cast<uint32_t>(THREAD_DIM)}, blockNums_, calNumPower_,525+ dim3{static_cast<uint32_t>(THREAD_DIM)}, blockNums_, calNumPower_,
523 (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(), (__gm__ volatile float*)tmpWeightGm_.GetPhyAddr(),526 (__gm__ volatile float*)tmpSumGm_.GetPhyAddr(), (__gm__ volatile float*)tmpWeightGm_.GetPhyAddr(),
524 (__gm__ T*)yGm_.GetPhyAddr(), (__gm__ T*)totalWeight_.GetPhyAddr());527 (__gm__ T*)yGm_.GetPhyAddr(), (__gm__ T*)totalWeight_.GetPhyAddr());
525 }528 }
@@ -564,4 +567,4 @@ private:
564 int64_t dealingNumTail_ = 0;567 int64_t dealingNumTail_ = 0;
565 int64_t offsetGmY_ = 0;568 int64_t offsetGmY_ = 0;
566};569};
567-#endif570+#endif
Mloss/nll_loss_grad/op_kernel/arch35/nll_loss_grad.h+28-27
@@ -18,6 +18,7 @@
18 18 
19#include "kernel_operator.h"19#include "kernel_operator.h"
20#include "../../inc/platform.h"20#include "../../inc/platform.h"
21+#include "simt_api/asc_simt.h"
21 22 
22namespace NLLLossGrad23namespace NLLLossGrad
23{24{
@@ -124,8 +125,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeNoReduce2
124 uint32_t blockNums_)125 uint32_t blockNums_)
125{126{
126 int64_t localClassNum = classNum_;127 int64_t localClassNum = classNum_;
127- for (uint64_t i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < batchNum_;128+ for (uint64_t i = blockId_ * blockDim.x + threadIdx.x; i < batchNum_;
128- i = i + blockNums_ * Simt::GetThreadNum()) {129+ i = i + blockNums_ * blockDim.x) {
129 int64_t chosenClass = target_[i];130 int64_t chosenClass = target_[i];
130 if (chosenClass == ignoreIdx_) {131 if (chosenClass == ignoreIdx_) {
131 continue;132 continue;
@@ -143,8 +144,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeReduceMea
143{144{
144 int64_t localClassNum = classNum_;145 int64_t localClassNum = classNum_;
145 ASSERT(localTotalWeight != 0 && "Input total weight must not be 0!");146 ASSERT(localTotalWeight != 0 && "Input total weight must not be 0!");
146- for (uint64_t i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < batchNum_;147+ for (uint64_t i = blockId_ * blockDim.x + threadIdx.x; i < batchNum_;
147- i = i + blockNums_ * Simt::GetThreadNum()) {148+ i = i + blockNums_ * blockDim.x) {
148 int64_t chosenClass = target_[i];149 int64_t chosenClass = target_[i];
149 if (chosenClass == ignoreIdx_) {150 if (chosenClass == ignoreIdx_) {
150 continue;151 continue;
@@ -161,8 +162,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeReduceSum
161 __gm__ T* outputGm_, __gm__ T* yGrad_, __gm__ T* weight_, uint32_t blockId_, uint32_t blockNums_)162 __gm__ T* outputGm_, __gm__ T* yGrad_, __gm__ T* weight_, uint32_t blockId_, uint32_t blockNums_)
162{163{
163 int64_t localClassNum = classNum_;164 int64_t localClassNum = classNum_;
164- for (uint64_t i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < batchNum_;165+ for (uint64_t i = blockId_ * blockDim.x + threadIdx.x; i < batchNum_;
165- i = i + blockNums_ * Simt::GetThreadNum()) {166+ i = i + blockNums_ * blockDim.x) {
166 int64_t chosenClass = target_[i];167 int64_t chosenClass = target_[i];
167 if (chosenClass == ignoreIdx_) {168 if (chosenClass == ignoreIdx_) {
168 continue;169 continue;
@@ -182,8 +183,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeNoReduce4
182 uint32_t blockNums_, uint64_t productOfNHW_, uint64_t productOfHW_)183 uint32_t blockNums_, uint64_t productOfNHW_, uint64_t productOfHW_)
183{184{
184 int64_t localClassNum = classNum_;185 int64_t localClassNum = classNum_;
185- for (uint64_t i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < productOfNHW_;186+ for (uint64_t i = blockId_ * blockDim.x + threadIdx.x; i < productOfNHW_;
186- i = i + blockNums_ * Simt::GetThreadNum()) {187+ i = i + blockNums_ * blockDim.x) {
187 int64_t chosenClass = target_[i];188 int64_t chosenClass = target_[i];
188 if (chosenClass == ignoreIdx_) {189 if (chosenClass == ignoreIdx_) {
189 continue;190 continue;
@@ -204,8 +205,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeReduceMea
204{205{
205 int64_t localClassNum = classNum_;206 int64_t localClassNum = classNum_;
206 ASSERT(localTotalWeight != 0 && "Input total weight must not be 0!");207 ASSERT(localTotalWeight != 0 && "Input total weight must not be 0!");
207- for (uint64_t i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < productOfNHW_;208+ for (uint64_t i = blockId_ * blockDim.x + threadIdx.x; i < productOfNHW_;
208- i = i + blockNums_ * Simt::GetThreadNum()) {209+ i = i + blockNums_ * blockDim.x) {
209 int64_t chosenClass = target_[i];210 int64_t chosenClass = target_[i];
210 if (chosenClass == ignoreIdx_) {211 if (chosenClass == ignoreIdx_) {
211 continue;212 continue;
@@ -224,8 +225,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void SimtComputeReduceSum
224 __gm__ T* outputGm_, __gm__ T* yGrad_, __gm__ T* weight_, uint32_t blockId_, uint32_t blockNums_, uint64_t productOfNHW_, uint64_t productOfHW_)225 __gm__ T* outputGm_, __gm__ T* yGrad_, __gm__ T* weight_, uint32_t blockId_, uint32_t blockNums_, uint64_t productOfNHW_, uint64_t productOfHW_)
225{226{
226 int64_t localClassNum = classNum_;227 int64_t localClassNum = classNum_;
227- for (uint64_t i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < productOfNHW_;228+ for (uint64_t i = blockId_ * blockDim.x + threadIdx.x; i < productOfNHW_;
228- i = i + blockNums_ * Simt::GetThreadNum()) {229+ i = i + blockNums_ * blockDim.x) {
229 int64_t chosenClass = target_[i];230 int64_t chosenClass = target_[i];
230 if (chosenClass == ignoreIdx_) {231 if (chosenClass == ignoreIdx_) {
231 continue;232 continue;
@@ -244,7 +245,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void ComputeSetValue(
244 uint64_t classNum_, uint64_t batchNum_, uint32_t blockId_, uint32_t blockNums_, __ubuf__ float* singleGrad_,245 uint64_t classNum_, uint64_t batchNum_, uint32_t blockId_, uint32_t blockNums_, __ubuf__ float* singleGrad_,
245 __gm__ T* outputGm_, __gm__ T* yGrad_, __ubuf__ float* localTotalWeight_, __gm__ T* total_weight_)246 __gm__ T* outputGm_, __gm__ T* yGrad_, __ubuf__ float* localTotalWeight_, __gm__ T* total_weight_)
246{247{
247- if (Simt::GetThreadIdx() == 0) {248+ if (threadIdx.x == 0) {
248 localTotalWeight_[0] = static_cast<float>(total_weight_[0]);249 localTotalWeight_[0] = static_cast<float>(total_weight_[0]);
249 singleGrad_[0] = static_cast<float>(yGrad_[0]);250 singleGrad_[0] = static_cast<float>(yGrad_[0]);
250 }251 }
@@ -281,8 +282,8 @@ __aicore__ inline void KernelNLLLossGrad<T, F>::Process()
281 tmpOut_ = tmpOutBuf_.Get<float>();282 tmpOut_ = tmpOutBuf_.Get<float>();
282 283 
283 if (this->isComputeCore_) {284 if (this->isComputeCore_) {
284- AscendC::Simt::VF_CALL<ComputeSetValue<T, F>>(285+ asc_vf_call<ComputeSetValue<T, F>>(
285- AscendC::Simt::Dim3{usedThread_}, classNum_, batchNum_, blockId_, blockNums_,286+ dim3{usedThread_}, classNum_, batchNum_, blockId_, blockNums_,
286 (__ubuf__ float*)singleGrad_.GetPhyAddr(), (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)yGrad_.GetPhyAddr(),287 (__ubuf__ float*)singleGrad_.GetPhyAddr(), (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)yGrad_.GetPhyAddr(),
287 (__ubuf__ float*)localTotalWeight_.GetPhyAddr(), (__gm__ T*)total_weight_.GetPhyAddr());288 (__ubuf__ float*)localTotalWeight_.GetPhyAddr(), (__gm__ T*)total_weight_.GetPhyAddr());
288 }289 }
@@ -297,35 +298,35 @@ __aicore__ inline void KernelNLLLossGrad<T, F>::Process()
297 PipeBarrier<PIPE_ALL>();298 PipeBarrier<PIPE_ALL>();
298 if(xDims_ == 1 || xDims_ == NUMBER_TWO){299 if(xDims_ == 1 || xDims_ == NUMBER_TWO){
299 if (reductionMode_ == NONE_MODE) {300 if (reductionMode_ == NONE_MODE) {
300- AscendC::Simt::VF_CALL<SimtComputeNoReduce2d<T, F>>(301+ asc_vf_call<SimtComputeNoReduce2d<T, F>>(
301- AscendC::Simt::Dim3{usedThread_}, classNum_, batchNum_, (__gm__ F*)target_.GetPhyAddr(), ignoreIdx_,302+ dim3{usedThread_}, classNum_, batchNum_, (__gm__ F*)target_.GetPhyAddr(), ignoreIdx_,
302 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)yGrad_.GetPhyAddr(), (__gm__ T*)weight_.GetPhyAddr(),303 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)yGrad_.GetPhyAddr(), (__gm__ T*)weight_.GetPhyAddr(),
303 blockId_, blockNums_);304 blockId_, blockNums_);
304 } else if (reductionMode_ == MEAN_MODE) {305 } else if (reductionMode_ == MEAN_MODE) {
305- AscendC::Simt::VF_CALL<SimtComputeReduceMean2d<T, F>>(306+ asc_vf_call<SimtComputeReduceMean2d<T, F>>(
306- AscendC::Simt::Dim3{usedThread_}, classNum_, batchNum_, (__gm__ F*)target_.GetPhyAddr(), ignoreIdx_,307+ dim3{usedThread_}, classNum_, batchNum_, (__gm__ F*)target_.GetPhyAddr(), ignoreIdx_,
307 (__ubuf__ float*)tmpOut_.GetPhyAddr(), (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)yGrad_.GetPhyAddr(),308 (__ubuf__ float*)tmpOut_.GetPhyAddr(), (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)yGrad_.GetPhyAddr(),
308 (__gm__ T*)weight_.GetPhyAddr(), blockId_, blockNums_);309 (__gm__ T*)weight_.GetPhyAddr(), blockId_, blockNums_);
309 } else if (reductionMode_ == SUM_MODE) {310 } else if (reductionMode_ == SUM_MODE) {
310- AscendC::Simt::VF_CALL<SimtComputeReduceSum2d<T, F>>(311+ asc_vf_call<SimtComputeReduceSum2d<T, F>>(
311- AscendC::Simt::Dim3{usedThread_}, classNum_, batchNum_, (__gm__ F*)target_.GetPhyAddr(), ignoreIdx_,312+ dim3{usedThread_}, classNum_, batchNum_, (__gm__ F*)target_.GetPhyAddr(), ignoreIdx_,
312 (__ubuf__ float*)singleGrad_.GetPhyAddr(), (__gm__ T*)outputGm_.GetPhyAddr(),313 (__ubuf__ float*)singleGrad_.GetPhyAddr(), (__gm__ T*)outputGm_.GetPhyAddr(),
313 (__gm__ T*)yGrad_.GetPhyAddr(), (__gm__ T*)weight_.GetPhyAddr(), blockId_, blockNums_);314 (__gm__ T*)yGrad_.GetPhyAddr(), (__gm__ T*)weight_.GetPhyAddr(), blockId_, blockNums_);
314 }315 }
315 }else if(xDims_ == NUMBER_FOUR){316 }else if(xDims_ == NUMBER_FOUR){
316 if (reductionMode_ == NONE_MODE) {317 if (reductionMode_ == NONE_MODE) {
317- AscendC::Simt::VF_CALL<SimtComputeNoReduce4d<T, F>>(318+ asc_vf_call<SimtComputeNoReduce4d<T, F>>(
318- AscendC::Simt::Dim3{usedThread_}, classNum_, batchNum_, height_, width_, (__gm__ F*)target_.GetPhyAddr(), ignoreIdx_,319+ dim3{usedThread_}, classNum_, batchNum_, height_, width_, (__gm__ F*)target_.GetPhyAddr(), ignoreIdx_,
319 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)yGrad_.GetPhyAddr(), (__gm__ T*)weight_.GetPhyAddr(),320 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)yGrad_.GetPhyAddr(), (__gm__ T*)weight_.GetPhyAddr(),
320 blockId_, blockNums_, productOfNHW_, productOfHW_);321 blockId_, blockNums_, productOfNHW_, productOfHW_);
321 } else if (reductionMode_ == MEAN_MODE) {322 } else if (reductionMode_ == MEAN_MODE) {
322- AscendC::Simt::VF_CALL<SimtComputeReduceMean4d<T, F>>(323+ asc_vf_call<SimtComputeReduceMean4d<T, F>>(
323- AscendC::Simt::Dim3{usedThread_}, classNum_, batchNum_, height_, width_, (__gm__ F*)target_.GetPhyAddr(), ignoreIdx_,324+ dim3{usedThread_}, classNum_, batchNum_, height_, width_, (__gm__ F*)target_.GetPhyAddr(), ignoreIdx_,
324 (__ubuf__ float*)tmpOut_.GetPhyAddr(), (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)yGrad_.GetPhyAddr(),325 (__ubuf__ float*)tmpOut_.GetPhyAddr(), (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)yGrad_.GetPhyAddr(),
325 (__gm__ T*)weight_.GetPhyAddr(), blockId_, blockNums_, productOfNHW_, productOfHW_);326 (__gm__ T*)weight_.GetPhyAddr(), blockId_, blockNums_, productOfNHW_, productOfHW_);
326 } else if (reductionMode_ == SUM_MODE) {327 } else if (reductionMode_ == SUM_MODE) {
327- AscendC::Simt::VF_CALL<SimtComputeReduceSum4d<T, F>>(328+ asc_vf_call<SimtComputeReduceSum4d<T, F>>(
328- AscendC::Simt::Dim3{usedThread_}, classNum_, batchNum_, height_, width_, (__gm__ F*)target_.GetPhyAddr(), ignoreIdx_,329+ dim3{usedThread_}, classNum_, batchNum_, height_, width_, (__gm__ F*)target_.GetPhyAddr(), ignoreIdx_,
329 (__ubuf__ float*)singleGrad_.GetPhyAddr(), (__gm__ T*)outputGm_.GetPhyAddr(),330 (__ubuf__ float*)singleGrad_.GetPhyAddr(), (__gm__ T*)outputGm_.GetPhyAddr(),
330 (__gm__ T*)yGrad_.GetPhyAddr(), (__gm__ T*)weight_.GetPhyAddr(), blockId_, blockNums_, productOfNHW_, productOfHW_);331 (__gm__ T*)yGrad_.GetPhyAddr(), (__gm__ T*)weight_.GetPhyAddr(), blockId_, blockNums_, productOfNHW_, productOfHW_);
331 }332 }
Mloss/sparse_softmax_cross_entropy_with_logits/op_kernel/arch35/sparse_softmax_cross_entropy_with_logits_full_load.h+3-2
@@ -20,6 +20,7 @@
20#include "kernel_operator.h"20#include "kernel_operator.h"
21#include "../inc/platform.h"21#include "../inc/platform.h"
22#include "sparse_softmax_cross_entropy_with_logits_tiling_data.h"22#include "sparse_softmax_cross_entropy_with_logits_tiling_data.h"
23+#include "simt_api/asc_simt.h"
23 24 
24namespace SparseSoftmaxCrossEntropyWithLogits {25namespace SparseSoftmaxCrossEntropyWithLogits {
25using namespace AscendC;26using namespace AscendC;
@@ -34,7 +35,7 @@ using AscendC::MicroAPI::UpdateMask;
34template <typename T1, typename T2, uint64_t schId, uint64_t db>35template <typename T1, typename T2, uint64_t schId, uint64_t db>
35__simt_vf__ __aicore__ LAUNCH_BOUND(1024) inline void UbSimtComputeLoopRFullLoad(__ubuf__ float* temp1Addr, __ubuf__ float* temp2Addr, __ubuf__ T2* labelsAddr, __ubuf__ float* gatherAddr, const int64_t tileNum, const int64_t rAlign, const int64_t rMax)36__simt_vf__ __aicore__ LAUNCH_BOUND(1024) inline void UbSimtComputeLoopRFullLoad(__ubuf__ float* temp1Addr, __ubuf__ float* temp2Addr, __ubuf__ T2* labelsAddr, __ubuf__ float* gatherAddr, const int64_t tileNum, const int64_t rAlign, const int64_t rMax)
36{37{
37- for (int64_t index = static_cast<int64_t>(Simt::GetThreadIdx()); index < tileNum; index += static_cast<int64_t>(Simt::GetThreadNum<0>())) {38+ for (int64_t index = static_cast<int64_t>(threadIdx.x); index < tileNum; index += static_cast<int64_t>(blockDim.x)) {
38 ASSERT((0 <= labelsAddr[index] && labelsAddr[index] < rMax) && "lable is not in [0, C)");39 ASSERT((0 <= labelsAddr[index] && labelsAddr[index] < rMax) && "lable is not in [0, C)");
39 int64_t offset = index * rAlign + labelsAddr[index];40 int64_t offset = index * rAlign + labelsAddr[index];
40 temp1Addr[offset] -= 1.0f;41 temp1Addr[offset] -= 1.0f;
@@ -209,7 +210,7 @@ __aicore__ inline void SparseSoftmaxCrossEntropyWithLogitsFullLoad<T1, T2, schId
209 VfBackProp(tileNum, r_, rUbNumFactor_, sumBuf, temp1Buf, logBuf, labelsBuf, backPropBuf, temp2Buf, subBuf);210 VfBackProp(tileNum, r_, rUbNumFactor_, sumBuf, temp1Buf, logBuf, labelsBuf, backPropBuf, temp2Buf, subBuf);
210 int64_t simtParam1 = rUbNumFactor_;211 int64_t simtParam1 = rUbNumFactor_;
211 int64_t simtParam2 = r_;212 int64_t simtParam2 = r_;
212- Simt::VF_CALL<UbSimtComputeLoopRFullLoad<T1, T2, schId, db>>(Simt::Dim3{1024}, (__ubuf__ float*)temp1Buf.GetPhyAddr(), (__ubuf__ float*)temp2Buf.GetPhyAddr(), (__ubuf__ T2*)labelsBuf.GetPhyAddr(), (__ubuf__ float*)maxBuf.GetPhyAddr(), tileNum, simtParam1, simtParam2);213+ asc_vf_call<UbSimtComputeLoopRFullLoad<T1, T2, schId, db>>(dim3{1024}, (__ubuf__ float*)temp1Buf.GetPhyAddr(), (__ubuf__ float*)temp2Buf.GetPhyAddr(), (__ubuf__ T2*)labelsBuf.GetPhyAddr(), (__ubuf__ float*)maxBuf.GetPhyAddr(), tileNum, simtParam1, simtParam2);
213 if constexpr (sizeof(T1) == 2) {214 if constexpr (sizeof(T1) == 2) {
214 AscendC::Cast(backPropBuf, temp1Buf, AscendC::RoundMode::CAST_RINT, tileNum * rUbNumFactor_);215 AscendC::Cast(backPropBuf, temp1Buf, AscendC::RoundMode::CAST_RINT, tileNum * rUbNumFactor_);
215 } else {216 } else {
Mloss/sparse_softmax_cross_entropy_with_logits/op_kernel/arch35/sparse_softmax_cross_entropy_with_logits_split_r.h+3-2
@@ -20,6 +20,7 @@
20#include "kernel_operator.h"20#include "kernel_operator.h"
21#include "../inc/platform.h"21#include "../inc/platform.h"
22#include "sparse_softmax_cross_entropy_with_logits_tiling_data.h"22#include "sparse_softmax_cross_entropy_with_logits_tiling_data.h"
23+#include "simt_api/asc_simt.h"
23 24 
24namespace SparseSoftmaxCrossEntropyWithLogits {25namespace SparseSoftmaxCrossEntropyWithLogits {
25using namespace AscendC;26using namespace AscendC;
@@ -28,7 +29,7 @@ using namespace AscendC::MicroAPI;
28template <typename T1, typename T2, uint64_t schId, uint64_t db>29template <typename T1, typename T2, uint64_t schId, uint64_t db>
29__simt_vf__ __aicore__ LAUNCH_BOUND(1024) inline void UbSimtComputeLoopRSplit(__ubuf__ T1* backPropAddr, __ubuf__ float* temp2Addr, __ubuf__ T2* labelsAddr, __ubuf__ float* gatherAddr, const int32_t tileNum, const int32_t rOnceNum, const int32_t rOnceNumAlign, const int64_t rMainNum, const int32_t ci, const int64_t cMax)30__simt_vf__ __aicore__ LAUNCH_BOUND(1024) inline void UbSimtComputeLoopRSplit(__ubuf__ T1* backPropAddr, __ubuf__ float* temp2Addr, __ubuf__ T2* labelsAddr, __ubuf__ float* gatherAddr, const int32_t tileNum, const int32_t rOnceNum, const int32_t rOnceNumAlign, const int64_t rMainNum, const int32_t ci, const int64_t cMax)
30{31{
31- for (int64_t index = static_cast<int64_t>(Simt::GetThreadIdx()); index < tileNum; index += static_cast<int64_t>(Simt::GetThreadNum<0>())) {32+ for (int64_t index = static_cast<int64_t>(threadIdx.x); index < tileNum; index += static_cast<int64_t>(blockDim.x)) {
32 ASSERT((0 <= labelsAddr[index] && labelsAddr[index] < cMax) && "labels is not in [0, C)");33 ASSERT((0 <= labelsAddr[index] && labelsAddr[index] < cMax) && "labels is not in [0, C)");
33 if (labelsAddr[index] >= rMainNum * ci && labelsAddr[index] < (rMainNum * ci + rOnceNum)) {34 if (labelsAddr[index] >= rMainNum * ci && labelsAddr[index] < (rMainNum * ci + rOnceNum)) {
34 int64_t offset = index * rOnceNumAlign + labelsAddr[index] - rMainNum * ci;35 int64_t offset = index * rOnceNumAlign + labelsAddr[index] - rMainNum * ci;
@@ -692,7 +693,7 @@ __aicore__ inline void SparseSoftmaxCrossEntropyWithLogitsSplitR<T1, T2, schId,
692 LocalTensor<T1> featureUb = featuresQueue_.DeQue<T1>();693 LocalTensor<T1> featureUb = featuresQueue_.DeQue<T1>();
693 LocalTensor<T1> backPropUb = backPropQueue_.AllocTensor<T1>();694 LocalTensor<T1> backPropUb = backPropQueue_.AllocTensor<T1>();
694 ComputeBackpropAndLoss(featureUb, backPropUb, lossUbFp32, tailNum, onceR, alignR, ci);695 ComputeBackpropAndLoss(featureUb, backPropUb, lossUbFp32, tailNum, onceR, alignR, ci);
695- Simt::VF_CALL<UbSimtComputeLoopRSplit<T1, T2, schId, db>>(Simt::Dim3{1024}, (__ubuf__ T1 *)backPropUb.GetPhyAddr(), (__ubuf__ float *)lossUbFp32.GetPhyAddr(), (__ubuf__ T2 *)labelsUb.GetPhyAddr(), (__ubuf__ float *)onceMaxUb.GetPhyAddr(), tailNum, onceR, alignR, rFactor, ci, cMax);696+ asc_vf_call<UbSimtComputeLoopRSplit<T1, T2, schId, db>>(dim3{1024}, (__ubuf__ T1 *)backPropUb.GetPhyAddr(), (__ubuf__ float *)lossUbFp32.GetPhyAddr(), (__ubuf__ T2 *)labelsUb.GetPhyAddr(), (__ubuf__ float *)onceMaxUb.GetPhyAddr(), tailNum, onceR, alignR, rFactor, ci, cMax);
696 featuresQueue_.FreeTensor(featureUb);697 featuresQueue_.FreeTensor(featureUb);
697 backPropQueue_.EnQue<T1>(backPropUb);698 backPropQueue_.EnQue<T1>(backPropUb);
698 CopyOutBackProb(tailNum, onceR, offset);699 CopyOutBackProb(tailNum, onceR, offset);