Star869
1749
代码
代码
Issues214
Pull Requests313
Actions
项目讨论
Wiki
分析
Star869
1749
已合并
simt-c风格转换: gather_elements/index/index_put_with_sort_v2/unsorted_segment_sum #4744
renshuojian创建于 5月11日
simt-c风格转换: gather_elements/index/index_put_with_sort_v2/unsorted_segment_sum #4744
已合并
renshuojian创建于 5月11日
从rsj007/ops-nn_7417master合入到cann/ops-nnmaster
共 12 个文件变更+314-298
Mindex/gather_elements/op_kernel/arch35/gather_elements.h+23-22
@@ -16,6 +16,7 @@
16#define GATHER_ELEMENTS_H16#define GATHER_ELEMENTS_H
17 17 
18#include "kernel_operator.h"18#include "kernel_operator.h"
19+#include "simt_api/asc_simt.h"
19#include "gather_elements_common.h"20#include "gather_elements_common.h"
20 21 
21namespace GatherElements {22namespace GatherElements {
@@ -68,13 +69,13 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_2048) inline void GatherDim1Compu
68 __gm__ const GatherElementsTilingData* tiling)69 __gm__ const GatherElementsTilingData* tiling)
69{70{
70 GET_TILING_DATA_PTR_WITH_STRUCT(GatherElementsTilingData, tilingData, tiling);71 GET_TILING_DATA_PTR_WITH_STRUCT(GatherElementsTilingData, tilingData, tiling);
71- if (Simt::GetBlockIdx() >= tilingData->usedCore) {72+ if (blockIdx.x >= tilingData->usedCore) {
72 return;73 return;
73 }74 }
74- U coreOffset = static_cast<U>(Simt::GetBlockIdx() * tilingData->perCoreNum);75+ U coreOffset = static_cast<U>(blockIdx.x * tilingData->perCoreNum);
75- U batchNum = static_cast<U>((Simt::GetBlockIdx() == tilingData->usedCore - 1) ? tilingData->tailCoreNum :76+ U batchNum = static_cast<U>((blockIdx.x == tilingData->usedCore - 1) ? tilingData->tailCoreNum :
76 tilingData->perCoreNum) + coreOffset;77 tilingData->perCoreNum) + coreOffset;
77- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {78+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
78 yAddr[i] = xAddr[indexAddr[i]];79 yAddr[i] = xAddr[indexAddr[i]];
79 }80 }
80}81}
@@ -84,17 +85,17 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_2048) inline void GatherDim2Compu
84 __gm__ const GatherElementsTilingData* tiling)85 __gm__ const GatherElementsTilingData* tiling)
85{86{
86 GET_TILING_DATA_PTR_WITH_STRUCT(GatherElementsTilingData, tilingData, tiling);87 GET_TILING_DATA_PTR_WITH_STRUCT(GatherElementsTilingData, tilingData, tiling);
87- if (Simt::GetBlockIdx() >= tilingData->usedCore) {88+ if (blockIdx.x >= tilingData->usedCore) {
88 return;89 return;
89 }90 }
90- U coreOffset = static_cast<U>(Simt::GetBlockIdx() * tilingData->perCoreNum);91+ U coreOffset = static_cast<U>(blockIdx.x * tilingData->perCoreNum);
91- U batchNum = static_cast<U>((Simt::GetBlockIdx() == tilingData->usedCore - 1) ? tilingData->tailCoreNum :92+ U batchNum = static_cast<U>((blockIdx.x == tilingData->usedCore - 1) ? tilingData->tailCoreNum :
92 tilingData->perCoreNum) + coreOffset;93 tilingData->perCoreNum) + coreOffset;
93 U magic = tilingData->magic[MS_IDX6];94 U magic = tilingData->magic[MS_IDX6];
94 U shift = tilingData->shift[MS_IDX6];95 U shift = tilingData->shift[MS_IDX6];
95 U indexStride = tilingData->indexStrideArr[MS_IDX6];96 U indexStride = tilingData->indexStrideArr[MS_IDX6];
96 U xStride = tilingData->xStrideArr[MS_IDX6];97 U xStride = tilingData->xStrideArr[MS_IDX6];
97- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {98+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
98 U indexVal = indexAddr[i];99 U indexVal = indexAddr[i];
99 100 
100 U dim0Index = Simt::UintDiv(i, magic, shift);101 U dim0Index = Simt::UintDiv(i, magic, shift);
@@ -112,7 +113,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_2048) inline void GatherDim3Compu
112 __gm__ X_T* yGm, __gm__ X_T* xGm, U m, U shift, U m1, U shift1, U indexStride, U xStridex,113 __gm__ X_T* yGm, __gm__ X_T* xGm, U m, U shift, U m1, U shift1, U indexStride, U xStridex,
113 U indexStride1, U xStridex1, U batchNum, U coreOffset)114 U indexStride1, U xStridex1, U batchNum, U coreOffset)
114{115{
115- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {116+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
116 U indexVal = indexGm[i];117 U indexVal = indexGm[i];
117 118 
118 U dim0Index = Simt::UintDiv(i, m, shift);119 U dim0Index = Simt::UintDiv(i, m, shift);
@@ -134,7 +135,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_1024) inline void GatherDim4Compu
134 __gm__ X_T* yGm, __gm__ X_T* xGm, __local_mem__ U* mAndShiftAddr, __local_mem__ U* strideAddr,135 __gm__ X_T* yGm, __gm__ X_T* xGm, __local_mem__ U* mAndShiftAddr, __local_mem__ U* strideAddr,
135 U batchNum, U coreOffset)136 U batchNum, U coreOffset)
136{137{
137- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {138+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
138 U indexVal = indexGm[i];139 U indexVal = indexGm[i];
139 140 
140 U dim0Index = Simt::UintDiv(i, mAndShiftAddr[MS_IDX4], mAndShiftAddr[M_SHIFT_OFFSET + MS_IDX4]);141 U dim0Index = Simt::UintDiv(i, mAndShiftAddr[MS_IDX4], mAndShiftAddr[M_SHIFT_OFFSET + MS_IDX4]);
@@ -171,7 +172,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_1024) inline void GatherDim5Compu
171 __gm__ X_T* yGm, __gm__ X_T* xGm, __local_mem__ U* mAndShiftAddr, __local_mem__ U* strideAddr,172 __gm__ X_T* yGm, __gm__ X_T* xGm, __local_mem__ U* mAndShiftAddr, __local_mem__ U* strideAddr,
172 U batchNum, U coreOffset)173 U batchNum, U coreOffset)
173{174{
174- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {175+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
175 U indexVal = indexGm[i];176 U indexVal = indexGm[i];
176 177 
177 U dim0Index = Simt::UintDiv(i, mAndShiftAddr[MS_IDX3], mAndShiftAddr[M_SHIFT_OFFSET + MS_IDX3]);178 U dim0Index = Simt::UintDiv(i, mAndShiftAddr[MS_IDX3], mAndShiftAddr[M_SHIFT_OFFSET + MS_IDX3]);
@@ -215,7 +216,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_512) inline void GatherDim6Comput
215 __gm__ X_T* yGm, __gm__ X_T* xGm, __local_mem__ U* mAndShiftAddr, __local_mem__ U* strideAddr,216 __gm__ X_T* yGm, __gm__ X_T* xGm, __local_mem__ U* mAndShiftAddr, __local_mem__ U* strideAddr,
216 U batchNum, U coreOffset)217 U batchNum, U coreOffset)
217{218{
218- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {219+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
219 U indexVal = indexGm[i];220 U indexVal = indexGm[i];
220 221 
221 U dim0Index = Simt::UintDiv(i, mAndShiftAddr[MS_IDX2], mAndShiftAddr[M_SHIFT_OFFSET + MS_IDX2]);222 U dim0Index = Simt::UintDiv(i, mAndShiftAddr[MS_IDX2], mAndShiftAddr[M_SHIFT_OFFSET + MS_IDX2]);
@@ -271,7 +272,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_512) inline void GatherDim7Comput
271 __gm__ X_T* yGm, __gm__ X_T* xGm, __local_mem__ U* mAndShiftAddr, __local_mem__ U* strideAddr,272 __gm__ X_T* yGm, __gm__ X_T* xGm, __local_mem__ U* mAndShiftAddr, __local_mem__ U* strideAddr,
272 U batchNum, U coreOffset)273 U batchNum, U coreOffset)
273{274{
274- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {275+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
275 U indexVal = indexGm[i];276 U indexVal = indexGm[i];
276 277 
277 U dim0Index = Simt::UintDiv(i, mAndShiftAddr[MS_IDX1], mAndShiftAddr[M_SHIFT_OFFSET + MS_IDX1]);278 U dim0Index = Simt::UintDiv(i, mAndShiftAddr[MS_IDX1], mAndShiftAddr[M_SHIFT_OFFSET + MS_IDX1]);
@@ -336,7 +337,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_512) inline void GatherDim8Comput
336 __gm__ X_T* yGm, __gm__ X_T* xGm, __local_mem__ U* mAndShiftAddr, __local_mem__ U* strideAddr,337 __gm__ X_T* yGm, __gm__ X_T* xGm, __local_mem__ U* mAndShiftAddr, __local_mem__ U* strideAddr,
337 U batchNum, U coreOffset)338 U batchNum, U coreOffset)
338{339{
339- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {340+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
340 U indexVal = indexGm[i];341 U indexVal = indexGm[i];
341 342 
342 U dim0Index = Simt::UintDiv(i, mAndShiftAddr[MS_IDX0], mAndShiftAddr[M_SHIFT_OFFSET + MS_IDX0]);343 U dim0Index = Simt::UintDiv(i, mAndShiftAddr[MS_IDX0], mAndShiftAddr[M_SHIFT_OFFSET + MS_IDX0]);
@@ -417,10 +418,10 @@ __aicore__ inline void GatherElementsKernel<X_T, INDEX_T, COM_T, DIM_NUM, AXIS>:
417 __gm__ const GatherElementsTilingData* tilingData)418 __gm__ const GatherElementsTilingData* tilingData)
418{419{
419 if constexpr (DIM_NUM == DIM1) {420 if constexpr (DIM_NUM == DIM1) {
420- Simt::VF_CALL<GatherDim1Compute<X_T, INDEX_T, COM_T>>(Simt::Dim3(THREAD_DIM_2048),421+ asc_vf_call<GatherDim1Compute<X_T, INDEX_T, COM_T>>(dim3(THREAD_DIM_2048),
421 xAddr, indexAddr, yAddr, tilingData);422 xAddr, indexAddr, yAddr, tilingData);
422 } else if constexpr (DIM_NUM == DIM2) {423 } else if constexpr (DIM_NUM == DIM2) {
423- Simt::VF_CALL<GatherDim2Compute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_2048),424+ asc_vf_call<GatherDim2Compute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_2048),
424 xAddr, indexAddr, yAddr, tilingData);425 xAddr, indexAddr, yAddr, tilingData);
425 }426 }
426}427}
@@ -455,26 +456,26 @@ __aicore__ inline void GatherElementsKernel<X_T, INDEX_T, COM_T, DIM_NUM, AXIS>:
455 __local_mem__ COM_T* strideAddr = (__local_mem__ COM_T*)(strideLocal.GetPhyAddr());456 __local_mem__ COM_T* strideAddr = (__local_mem__ COM_T*)(strideLocal.GetPhyAddr());
456 457 
457 if constexpr (DIM_NUM == DIM3) {458 if constexpr (DIM_NUM == DIM3) {
458- Simt::VF_CALL<GatherDim3Compute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_2048),459+ asc_vf_call<GatherDim3Compute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_2048),
459 indexAddr, yAddr, xAddr, m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],460 indexAddr, yAddr, xAddr, m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],
460 static_cast<COM_T>(tilingData_->indexStrideArr[MS_IDX5]),461 static_cast<COM_T>(tilingData_->indexStrideArr[MS_IDX5]),
461 static_cast<COM_T>(tilingData_->xStrideArr[MS_IDX5]),462 static_cast<COM_T>(tilingData_->xStrideArr[MS_IDX5]),
462 static_cast<COM_T>(tilingData_->indexStrideArr[MS_IDX6]),463 static_cast<COM_T>(tilingData_->indexStrideArr[MS_IDX6]),
463 static_cast<COM_T>(tilingData_->xStrideArr[MS_IDX6]), batchNum, coreOffset);464 static_cast<COM_T>(tilingData_->xStrideArr[MS_IDX6]), batchNum, coreOffset);
464 } else if constexpr (DIM_NUM == DIM4) {465 } else if constexpr (DIM_NUM == DIM4) {
465- Simt::VF_CALL<GatherDim4Compute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_1024),466+ asc_vf_call<GatherDim4Compute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_1024),
466 indexAddr, yAddr, xAddr, mAndShiftAddr, strideAddr, batchNum, coreOffset);467 indexAddr, yAddr, xAddr, mAndShiftAddr, strideAddr, batchNum, coreOffset);
467 } else if constexpr (DIM_NUM == DIM5) {468 } else if constexpr (DIM_NUM == DIM5) {
468- Simt::VF_CALL<GatherDim5Compute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_1024),469+ asc_vf_call<GatherDim5Compute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_1024),
469 indexAddr, yAddr, xAddr, mAndShiftAddr, strideAddr, batchNum, coreOffset);470 indexAddr, yAddr, xAddr, mAndShiftAddr, strideAddr, batchNum, coreOffset);
470 } else if constexpr (DIM_NUM == DIM6) {471 } else if constexpr (DIM_NUM == DIM6) {
471- Simt::VF_CALL<GatherDim6Compute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_512),472+ asc_vf_call<GatherDim6Compute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_512),
472 indexAddr, yAddr, xAddr, mAndShiftAddr, strideAddr, batchNum, coreOffset);473 indexAddr, yAddr, xAddr, mAndShiftAddr, strideAddr, batchNum, coreOffset);
473 } else if constexpr (DIM_NUM == DIM7) {474 } else if constexpr (DIM_NUM == DIM7) {
474- Simt::VF_CALL<GatherDim7Compute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_512),475+ asc_vf_call<GatherDim7Compute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_512),
475 indexAddr, yAddr, xAddr, mAndShiftAddr, strideAddr, batchNum, coreOffset);476 indexAddr, yAddr, xAddr, mAndShiftAddr, strideAddr, batchNum, coreOffset);
476 } else {477 } else {
477- Simt::VF_CALL<GatherDim8Compute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_512),478+ asc_vf_call<GatherDim8Compute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_512),
478 indexAddr, yAddr, xAddr, mAndShiftAddr, strideAddr, batchNum, coreOffset);479 indexAddr, yAddr, xAddr, mAndShiftAddr, strideAddr, batchNum, coreOffset);
479 }480 }
480}481}
Mindex/gather_elements/op_kernel/arch35/gather_elements_no_contiguos.h+25-24
@@ -16,6 +16,7 @@
16#define GATHER_ELEMENTS_NO_CONTIGUOUS_H16#define GATHER_ELEMENTS_NO_CONTIGUOUS_H
17 17 
18#include "kernel_operator.h"18#include "kernel_operator.h"
19+#include "simt_api/asc_simt.h"
19#include "gather_elements_common.h"20#include "gather_elements_common.h"
20 21 
21namespace GatherElements {22namespace GatherElements {
@@ -60,7 +61,7 @@ template <typename X_T, typename INDEX_T, typename U>
60__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_2048) inline void GatherDim1NoContiguousCompute(__gm__ INDEX_T* indexGm, __gm__ X_T* yGm,61__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_2048) inline void GatherDim1NoContiguousCompute(__gm__ INDEX_T* indexGm, __gm__ X_T* yGm,
61 __gm__ X_T* xGm, U indexStride, U xStridex, U batchNum, U coreOffset)62 __gm__ X_T* xGm, U indexStride, U xStridex, U batchNum, U coreOffset)
62{63{
63- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {64+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
64 U indexVal = indexGm[i * indexStride];65 U indexVal = indexGm[i * indexStride];
65 yGm[i] = xGm[indexVal * xStridex];66 yGm[i] = xGm[indexVal * xStridex];
66 }67 }
@@ -71,7 +72,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_2048) inline void GatherDim2NoCon
71 __gm__ X_T* xGm, U m, U shift, U indexStride, U xStridex, U indexContiguousStride, U indexStride1, U xStridex1, U batchNum,72 __gm__ X_T* xGm, U m, U shift, U indexStride, U xStridex, U indexContiguousStride, U indexStride1, U xStridex1, U batchNum,
72 U coreOffset)73 U coreOffset)
73{74{
74- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {75+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
75 U dim0Index = Simt::UintDiv(i, m, shift);76 U dim0Index = Simt::UintDiv(i, m, shift);
76 U dim1Index = i - dim0Index * indexContiguousStride;77 U dim1Index = i - dim0Index * indexContiguousStride;
77 U indexVal = indexGm[dim0Index * indexStride + dim1Index * indexStride1];78 U indexVal = indexGm[dim0Index * indexStride + dim1Index * indexStride1];
@@ -89,7 +90,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_2048) inline void GatherDim3NoCon
89 __gm__ X_T* xGm, U m, U shift, U m1, U shift1, U indexStride, U xStridex, U indexContiguousStride, U indexStride1, U xStridex1,90 __gm__ X_T* xGm, U m, U shift, U m1, U shift1, U indexStride, U xStridex, U indexContiguousStride, U indexStride1, U xStridex1,
90 U indexContiguousStride1, U indexStride2, U xStridex2, U batchNum, U coreOffset)91 U indexContiguousStride1, U indexStride2, U xStridex2, U batchNum, U coreOffset)
91{92{
92- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {93+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
93 U dim0Index = Simt::UintDiv(i, m, shift);94 U dim0Index = Simt::UintDiv(i, m, shift);
94 U newIdx2 = i - dim0Index * indexContiguousStride;95 U newIdx2 = i - dim0Index * indexContiguousStride;
95 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);96 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);
@@ -110,7 +111,7 @@ template <typename X_T, typename INDEX_T, typename U, int32_t AXIS>
110__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_2048) inline void GatherDim3NoContiguousComputeB64(__gm__ INDEX_T* indexGm, __gm__ X_T* yGm,111__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_2048) inline void GatherDim3NoContiguousComputeB64(__gm__ INDEX_T* indexGm, __gm__ X_T* yGm,
111 __gm__ X_T* xGm, __ubuf__ U* stride, U m, U shift, U m1, U shift1, U batchNum, U coreOffset)112 __gm__ X_T* xGm, __ubuf__ U* stride, U m, U shift, U m1, U shift1, U batchNum, U coreOffset)
112{113{
113- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {114+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
114 U dim0Index = Simt::UintDiv(i, m, shift);115 U dim0Index = Simt::UintDiv(i, m, shift);
115 U newIdx2 = i - dim0Index * stride[0];116 U newIdx2 = i - dim0Index * stride[0];
116 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);117 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);
@@ -133,7 +134,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_1024) inline void GatherDim4NoCon
133 U indexStride1, U xStridex1, U indexContiguousStride1, U indexStride2, U xStridex2, U indexContiguousStride2, U indexStride3, U xStridex3,134 U indexStride1, U xStridex1, U indexContiguousStride1, U indexStride2, U xStridex2, U indexContiguousStride2, U indexStride3, U xStridex3,
134 U batchNum, U coreOffset)135 U batchNum, U coreOffset)
135{136{
136- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {137+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
137 U dim0Index = Simt::UintDiv(i, m, shift);138 U dim0Index = Simt::UintDiv(i, m, shift);
138 U newIdx2 = i - dim0Index * indexContiguousStride;139 U newIdx2 = i - dim0Index * indexContiguousStride;
139 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);140 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);
@@ -158,7 +159,7 @@ template <typename X_T, typename INDEX_T, typename U, int32_t AXIS>
158__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_1024) inline void GatherDim4NoContiguousComputeB64(__gm__ INDEX_T* indexGm, __gm__ X_T* yGm, 159__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_1024) inline void GatherDim4NoContiguousComputeB64(__gm__ INDEX_T* indexGm, __gm__ X_T* yGm,
159 __gm__ X_T* xGm, __ubuf__ U* stride, U m, U shift, U m1, U shift1, U m2, U shift2, U batchNum, U coreOffset)160 __gm__ X_T* xGm, __ubuf__ U* stride, U m, U shift, U m1, U shift1, U m2, U shift2, U batchNum, U coreOffset)
160{161{
161- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {162+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
162 U dim0Index = Simt::UintDiv(i, m, shift);163 U dim0Index = Simt::UintDiv(i, m, shift);
163 U newIdx2 = i - dim0Index * stride[0];164 U newIdx2 = i - dim0Index * stride[0];
164 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);165 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);
@@ -185,7 +186,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_1024) inline void GatherDim5NoCon
185 U indexStride1, U xStridex1, U indexContiguousStride1, U indexStride2, U xStridex2, U indexContiguousStride2, 186 U indexStride1, U xStridex1, U indexContiguousStride1, U indexStride2, U xStridex2, U indexContiguousStride2,
186 U indexStride3, U xStridex3, U indexContiguousStride3, U indexStride4, U xStridex4, U batchNum, U coreOffset)187 U indexStride3, U xStridex3, U indexContiguousStride3, U indexStride4, U xStridex4, U batchNum, U coreOffset)
187{188{
188- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {189+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
189 U dim0Index = Simt::UintDiv(i, m, shift);190 U dim0Index = Simt::UintDiv(i, m, shift);
190 U newIdx2 = i - dim0Index * indexContiguousStride;191 U newIdx2 = i - dim0Index * indexContiguousStride;
191 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);192 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);
@@ -218,7 +219,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_512) inline void GatherDim6NoCont
218 U indexStride3, U xStridex3, U indexContiguousStride3, U indexStride4, U xStridex4, U indexContiguousStride4,219 U indexStride3, U xStridex3, U indexContiguousStride3, U indexStride4, U xStridex4, U indexContiguousStride4,
219 U indexStride5, U xStridex5, U batchNum, U coreOffset)220 U indexStride5, U xStridex5, U batchNum, U coreOffset)
220{221{
221- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {222+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
222 U dim0Index = Simt::UintDiv(i, m, shift);223 U dim0Index = Simt::UintDiv(i, m, shift);
223 U newIdx2 = i - dim0Index * indexContiguousStride;224 U newIdx2 = i - dim0Index * indexContiguousStride;
224 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);225 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);
@@ -262,7 +263,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_512) inline void GatherDim7NoCont
262 U indexStride4, U xStridex4, U indexContiguousStride4, U indexStride5, U xStridex5, U indexContiguousStride5,263 U indexStride4, U xStridex4, U indexContiguousStride4, U indexStride5, U xStridex5, U indexContiguousStride5,
263 U indexStride6, U xStridex6, U batchNum, U coreOffset)264 U indexStride6, U xStridex6, U batchNum, U coreOffset)
264{265{
265- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {266+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
266 U dim0Index = Simt::UintDiv(i, m, shift);267 U dim0Index = Simt::UintDiv(i, m, shift);
267 U newIdx2 = i - dim0Index * indexContiguousStride;268 U newIdx2 = i - dim0Index * indexContiguousStride;
268 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);269 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);
@@ -311,7 +312,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_512) inline void GatherDim8NoCont
311 U indexStride4, U xStridex4, U indexContiguousStride4, U indexStride5, U xStridex5, U indexContiguousStride5,312 U indexStride4, U xStridex4, U indexContiguousStride4, U indexStride5, U xStridex5, U indexContiguousStride5,
312 U indexStride6, U xStridex6, U indexContiguousStride6, U indexStride7, U xStridex7, U batchNum, U coreOffset)313 U indexStride6, U xStridex6, U indexContiguousStride6, U indexStride7, U xStridex7, U batchNum, U coreOffset)
313{314{
314- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {315+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
315 U dim0Index = Simt::UintDiv(i, m, shift);316 U dim0Index = Simt::UintDiv(i, m, shift);
316 U newIdx2 = i - dim0Index * indexContiguousStride;317 U newIdx2 = i - dim0Index * indexContiguousStride;
317 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);318 U dim1Index = Simt::UintDiv(newIdx2, m1, shift1);
@@ -362,7 +363,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_2048) inline void GatherDim3Compu
362 __gm__ X_T* xGm, U m, U shift, U m1, U shift1, U indexStride, U xStridex, U indexContiguousStride, U yStride, U indexStride1,363 __gm__ X_T* xGm, U m, U shift, U m1, U shift1, U indexStride, U xStridex, U indexContiguousStride, U yStride, U indexStride1,
363 U xStridex1, U indexContiguousStride1, U yStride1, U indexStride2, U xStridex2, U batchNum, U coreOffset)364 U xStridex1, U indexContiguousStride1, U yStride1, U indexStride2, U xStridex2, U batchNum, U coreOffset)
364{365{
365- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {366+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
366 // [a, b, c] -> [b, a, c]367 // [a, b, c] -> [b, a, c]
367 U dim1Index = Simt::UintDiv(i, m, shift);368 U dim1Index = Simt::UintDiv(i, m, shift);
368 U newIdx2 = i - dim1Index * indexContiguousStride;369 U newIdx2 = i - dim1Index * indexContiguousStride;
@@ -384,7 +385,7 @@ template <typename X_T, typename INDEX_T, typename U, int32_t AXIS>
384__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_2048) inline void GatherDim3ComputeTransPoseB64(__gm__ INDEX_T* indexGm, __gm__ X_T* yGm,385__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_2048) inline void GatherDim3ComputeTransPoseB64(__gm__ INDEX_T* indexGm, __gm__ X_T* yGm,
385 __gm__ X_T* xGm, __ubuf__ U* stride, U m, U shift, U m1, U shift1, U batchNum, U coreOffset)386 __gm__ X_T* xGm, __ubuf__ U* stride, U m, U shift, U m1, U shift1, U batchNum, U coreOffset)
386{387{
387- for (U i = Simt::GetThreadIdx() + coreOffset; i < batchNum; i += Simt::GetThreadNum()) {388+ for (U i = threadIdx.x + coreOffset; i < batchNum; i += blockDim.x) {
388 // [a, b, c] -> [b, a, c]389 // [a, b, c] -> [b, a, c]
389 U dim1Index = Simt::UintDiv(i, m, shift);390 U dim1Index = Simt::UintDiv(i, m, shift);
390 U newIdx2 = i - dim1Index * stride[0];391 U newIdx2 = i - dim1Index * stride[0];
@@ -424,14 +425,14 @@ __aicore__ inline void GatherElementsKernelNoContiguous<X_T, INDEX_T, COM_T, DIM
424 __gm__ X_T* xAddr = (__gm__ X_T*)(xGm_.GetPhyAddr());425 __gm__ X_T* xAddr = (__gm__ X_T*)(xGm_.GetPhyAddr());
425 426 
426 if constexpr (DIM_NUM == DIM1) {427 if constexpr (DIM_NUM == DIM1) {
427- Simt::VF_CALL<GatherDim1NoContiguousCompute<X_T, INDEX_T, COM_T>>(Simt::Dim3(THREAD_DIM_2048),428+ asc_vf_call<GatherDim1NoContiguousCompute<X_T, INDEX_T, COM_T>>(dim3(THREAD_DIM_2048),
428 indexAddr, yAddr, xAddr, 429 indexAddr, yAddr, xAddr,
429 static_cast<COM_T>(tilingData_->indexStride[MS_IDX7]),430 static_cast<COM_T>(tilingData_->indexStride[MS_IDX7]),
430 static_cast<COM_T>(tilingData_->xStride[MS_IDX7]), 431 static_cast<COM_T>(tilingData_->xStride[MS_IDX7]),
431 batchNum, 432 batchNum,
432 coreOffset);433 coreOffset);
433 } else if constexpr (DIM_NUM == DIM2) {434 } else if constexpr (DIM_NUM == DIM2) {
434- Simt::VF_CALL<GatherDim2NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_2048),435+ asc_vf_call<GatherDim2NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_2048),
435 indexAddr, yAddr, xAddr, m_[MS_IDX6], shift_[MS_IDX6],436 indexAddr, yAddr, xAddr, m_[MS_IDX6], shift_[MS_IDX6],
436 static_cast<COM_T>(tilingData_->indexStride[MS_IDX6]),437 static_cast<COM_T>(tilingData_->indexStride[MS_IDX6]),
437 static_cast<COM_T>(tilingData_->xStride[MS_IDX6]),438 static_cast<COM_T>(tilingData_->xStride[MS_IDX6]),
@@ -442,7 +443,7 @@ __aicore__ inline void GatherElementsKernelNoContiguous<X_T, INDEX_T, COM_T, DIM
442 coreOffset); 443 coreOffset);
443 } else if constexpr (DIM_NUM == DIM3) {444 } else if constexpr (DIM_NUM == DIM3) {
444 if constexpr (sizeof(COM_T) == B32) {445 if constexpr (sizeof(COM_T) == B32) {
445- Simt::VF_CALL<GatherDim3NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_2048),446+ asc_vf_call<GatherDim3NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_2048),
446 indexAddr, yAddr, xAddr, m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],447 indexAddr, yAddr, xAddr, m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],
447 static_cast<COM_T>(tilingData_->indexStride[MS_IDX5]),448 static_cast<COM_T>(tilingData_->indexStride[MS_IDX5]),
448 static_cast<COM_T>(tilingData_->xStride[MS_IDX5]),449 static_cast<COM_T>(tilingData_->xStride[MS_IDX5]),
@@ -464,7 +465,7 @@ __aicore__ inline void GatherElementsKernelNoContiguous<X_T, INDEX_T, COM_T, DIM
464 }465 }
465 DataSyncBarrier<MemDsbT::UB>();466 DataSyncBarrier<MemDsbT::UB>();
466 __local_mem__ COM_T* strideAddr = (__local_mem__ COM_T*)(strideLocal.GetPhyAddr());467 __local_mem__ COM_T* strideAddr = (__local_mem__ COM_T*)(strideLocal.GetPhyAddr());
467- Simt::VF_CALL<GatherDim3NoContiguousComputeB64<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_2048),468+ asc_vf_call<GatherDim3NoContiguousComputeB64<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_2048),
468 indexAddr, yAddr, xAddr, strideAddr, 469 indexAddr, yAddr, xAddr, strideAddr,
469 m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],470 m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],
470 batchNum, 471 batchNum,
@@ -472,7 +473,7 @@ __aicore__ inline void GatherElementsKernelNoContiguous<X_T, INDEX_T, COM_T, DIM
472 }473 }
473 } else if constexpr (DIM_NUM == DIM4) {474 } else if constexpr (DIM_NUM == DIM4) {
474 if constexpr (sizeof(COM_T) == B32) {475 if constexpr (sizeof(COM_T) == B32) {
475- Simt::VF_CALL<GatherDim4NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_1024),476+ asc_vf_call<GatherDim4NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_1024),
476 indexAddr, yAddr, xAddr, m_[MS_IDX4], shift_[MS_IDX4],477 indexAddr, yAddr, xAddr, m_[MS_IDX4], shift_[MS_IDX4],
477 m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],478 m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],
478 static_cast<COM_T>(tilingData_->indexStride[MS_IDX4]),479 static_cast<COM_T>(tilingData_->indexStride[MS_IDX4]),
@@ -498,14 +499,14 @@ __aicore__ inline void GatherElementsKernelNoContiguous<X_T, INDEX_T, COM_T, DIM
498 }499 }
499 DataSyncBarrier<MemDsbT::UB>();500 DataSyncBarrier<MemDsbT::UB>();
500 __local_mem__ COM_T* strideAddr = (__local_mem__ COM_T*)(strideLocal.GetPhyAddr());501 __local_mem__ COM_T* strideAddr = (__local_mem__ COM_T*)(strideLocal.GetPhyAddr());
501- Simt::VF_CALL<GatherDim4NoContiguousComputeB64<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_1024),502+ asc_vf_call<GatherDim4NoContiguousComputeB64<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_1024),
502 indexAddr, yAddr, xAddr, strideAddr, m_[MS_IDX4], shift_[MS_IDX4],503 indexAddr, yAddr, xAddr, strideAddr, m_[MS_IDX4], shift_[MS_IDX4],
503 m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6], 504 m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],
504 batchNum, 505 batchNum,
505 coreOffset);506 coreOffset);
506 }507 }
507 } else if constexpr (DIM_NUM == DIM5) {508 } else if constexpr (DIM_NUM == DIM5) {
508- Simt::VF_CALL<GatherDim5NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_1024),509+ asc_vf_call<GatherDim5NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_1024),
509 indexAddr, yAddr, xAddr, m_[MS_IDX3], shift_[MS_IDX3], m_[MS_IDX4], shift_[MS_IDX4],510 indexAddr, yAddr, xAddr, m_[MS_IDX3], shift_[MS_IDX3], m_[MS_IDX4], shift_[MS_IDX4],
510 m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],511 m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],
511 static_cast<COM_T>(tilingData_->indexStride[MS_IDX3]),512 static_cast<COM_T>(tilingData_->indexStride[MS_IDX3]),
@@ -525,7 +526,7 @@ __aicore__ inline void GatherElementsKernelNoContiguous<X_T, INDEX_T, COM_T, DIM
525 batchNum, 526 batchNum,
526 coreOffset);527 coreOffset);
527 } else if constexpr (DIM_NUM == DIM6) {528 } else if constexpr (DIM_NUM == DIM6) {
528- Simt::VF_CALL<GatherDim6NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_512),529+ asc_vf_call<GatherDim6NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_512),
529 indexAddr, yAddr, xAddr, m_[MS_IDX2], shift_[MS_IDX2], m_[MS_IDX3], shift_[MS_IDX3], m_[MS_IDX4], shift_[MS_IDX4],530 indexAddr, yAddr, xAddr, m_[MS_IDX2], shift_[MS_IDX2], m_[MS_IDX3], shift_[MS_IDX3], m_[MS_IDX4], shift_[MS_IDX4],
530 m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],531 m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],
531 static_cast<COM_T>(tilingData_->indexStride[MS_IDX2]),532 static_cast<COM_T>(tilingData_->indexStride[MS_IDX2]),
@@ -548,7 +549,7 @@ __aicore__ inline void GatherElementsKernelNoContiguous<X_T, INDEX_T, COM_T, DIM
548 batchNum, 549 batchNum,
549 coreOffset);550 coreOffset);
550 } else if constexpr (DIM_NUM == DIM7) {551 } else if constexpr (DIM_NUM == DIM7) {
551- Simt::VF_CALL<GatherDim7NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_512),552+ asc_vf_call<GatherDim7NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_512),
552 indexAddr, yAddr, xAddr, m_[MS_IDX1], shift_[MS_IDX1], m_[MS_IDX2], shift_[MS_IDX2], m_[MS_IDX3], shift_[MS_IDX3], 553 indexAddr, yAddr, xAddr, m_[MS_IDX1], shift_[MS_IDX1], m_[MS_IDX2], shift_[MS_IDX2], m_[MS_IDX3], shift_[MS_IDX3],
553 m_[MS_IDX4], shift_[MS_IDX4], m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],554 m_[MS_IDX4], shift_[MS_IDX4], m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],
554 static_cast<COM_T>(tilingData_->indexStride[MS_IDX1]),555 static_cast<COM_T>(tilingData_->indexStride[MS_IDX1]),
@@ -574,7 +575,7 @@ __aicore__ inline void GatherElementsKernelNoContiguous<X_T, INDEX_T, COM_T, DIM
574 batchNum, 575 batchNum,
575 coreOffset);576 coreOffset);
576 } else if constexpr (DIM_NUM == DIM8) {577 } else if constexpr (DIM_NUM == DIM8) {
577- Simt::VF_CALL<GatherDim8NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_512),578+ asc_vf_call<GatherDim8NoContiguousCompute<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_512),
578 indexAddr, yAddr, xAddr, m_[MS_IDX0], shift_[MS_IDX0], m_[MS_IDX1], shift_[MS_IDX1], m_[MS_IDX2], shift_[MS_IDX2], 579 indexAddr, yAddr, xAddr, m_[MS_IDX0], shift_[MS_IDX0], m_[MS_IDX1], shift_[MS_IDX1], m_[MS_IDX2], shift_[MS_IDX2],
579 m_[MS_IDX3], shift_[MS_IDX3], m_[MS_IDX4], shift_[MS_IDX4], m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],580 m_[MS_IDX3], shift_[MS_IDX3], m_[MS_IDX4], shift_[MS_IDX4], m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],
580 static_cast<COM_T>(tilingData_->indexStride[MS_IDX0]),581 static_cast<COM_T>(tilingData_->indexStride[MS_IDX0]),
@@ -629,7 +630,7 @@ __aicore__ inline void GatherElementsKernelNoContiguous<X_T, INDEX_T, COM_T, DIM
629 630
630 if constexpr (DIM_NUM == DIM3) {631 if constexpr (DIM_NUM == DIM3) {
631 if constexpr (sizeof(COM_T) == B32) {632 if constexpr (sizeof(COM_T) == B32) {
632- Simt::VF_CALL<GatherDim3ComputeTransPose<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_2048),633+ asc_vf_call<GatherDim3ComputeTransPose<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_2048),
633 indexAddr, yAddr, xAddr, m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],634 indexAddr, yAddr, xAddr, m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],
634 static_cast<COM_T>(tilingData_->indexStride[MS_IDX5]),635 static_cast<COM_T>(tilingData_->indexStride[MS_IDX5]),
635 static_cast<COM_T>(tilingData_->xStride[MS_IDX5]),636 static_cast<COM_T>(tilingData_->xStride[MS_IDX5]),
@@ -654,7 +655,7 @@ __aicore__ inline void GatherElementsKernelNoContiguous<X_T, INDEX_T, COM_T, DIM
654 }655 }
655 DataSyncBarrier<MemDsbT::UB>(); 656 DataSyncBarrier<MemDsbT::UB>();
656 __local_mem__ COM_T* strideAddr = (__local_mem__ COM_T*)(strideLocal.GetPhyAddr()); 657 __local_mem__ COM_T* strideAddr = (__local_mem__ COM_T*)(strideLocal.GetPhyAddr());
657- Simt::VF_CALL<GatherDim3ComputeTransPoseB64<X_T, INDEX_T, COM_T, AXIS>>(Simt::Dim3(THREAD_DIM_2048),658+ asc_vf_call<GatherDim3ComputeTransPoseB64<X_T, INDEX_T, COM_T, AXIS>>(dim3(THREAD_DIM_2048),
658 indexAddr, yAddr, xAddr, strideAddr, m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],659 indexAddr, yAddr, xAddr, strideAddr, m_[MS_IDX5], shift_[MS_IDX5], m_[MS_IDX6], shift_[MS_IDX6],
659 batchNum, 660 batchNum,
660 coreOffset);661 coreOffset);
Mindex/index/op_kernel/arch35/index.h+91-87
@@ -19,6 +19,10 @@
19#include "kernel_operator.h"19#include "kernel_operator.h"
20#include "op_kernel/platform_util.h"20#include "op_kernel/platform_util.h"
21 21 
22+#include "simt_api/asc_simt.h"
23+#include "simt_api/device_atomic_functions.h"
24+#include "simt_api/asc_fp16.h"
25+#include "simt_api/asc_bf16.h"
22template <typename T>26template <typename T>
23struct calcParams {27struct calcParams {
24 uint64_t indexedDimStride_[8] = {0, 0, 0, 0, 0, 0, 0, 0};28 uint64_t indexedDimStride_[8] = {0, 0, 0, 0, 0, 0, 0, 0};
@@ -106,7 +110,7 @@ struct IndexPutAdd {
106 output[idx] = true;110 output[idx] = true;
107 }111 }
108 } else {112 } else {
109- Simt::AtomicAdd(output + idx, input[i]);113+ asc_atomic_add(output + idx, input[i]);
110 }114 }
111 }115 }
112};116};
@@ -148,8 +152,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(LAUNCH_BOUND_LIMIT) inline void SimtComputeC
148 localIndexedDimStride[i] = calcParamsPtr_->indexedDimStride_[i];152 localIndexedDimStride[i] = calcParamsPtr_->indexedDimStride_[i];
149 }153 }
150 154 
151- for (T2 i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < outputLength_;155+ for (T2 i = blockId_ * blockDim.x + threadIdx.x; i < outputLength_;
152- i = i + blockNums_ * Simt::GetThreadNum()) {156+ i = i + blockNums_ * blockDim.x) {
153 T2 inputIndex = 0;157 T2 inputIndex = 0;
154 T2 remLength = i;158 T2 remLength = i;
155 if (formerNonIndexStride_ != 1) {159 if (formerNonIndexStride_ != 1) {
@@ -219,8 +223,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(LAUNCH_BOUND_LIMIT) inline void SimtComputeN
219 localNonIndexedStride[i] = calcParamsPtr_->nonIndexedStride_[i];223 localNonIndexedStride[i] = calcParamsPtr_->nonIndexedStride_[i];
220 }224 }
221 225 
222- for (T2 i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < outputLength_;226+ for (T2 i = blockId_ * blockDim.x + threadIdx.x; i < outputLength_;
223- i = i + blockNums_ * Simt::GetThreadNum()) {227+ i = i + blockNums_ * blockDim.x) {
224 // quick div for i / innerLoopLength_228 // quick div for i / innerLoopLength_
225 T2 outLoopIdx = AscendC::Simt::UintDiv(i, m3_, shift3_);229 T2 outLoopIdx = AscendC::Simt::UintDiv(i, m3_, shift3_);
226 230 
@@ -417,202 +421,202 @@ __aicore__ inline void KernelIndex<T, F, P, T2>::Process()
417 if (indexContinue_) {421 if (indexContinue_) {
418 if (indexedDimNum_ == DIM_NUMS_ONE) {422 if (indexedDimNum_ == DIM_NUMS_ONE) {
419 if (isSmallIndexSize) {423 if (isSmallIndexSize) {
420- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_ONE, T2, THREAD_DIM_SMALL>>(424+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_ONE, T2, THREAD_DIM_SMALL>>(
421- AscendC::Simt::Dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,425+ dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,
422 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),426 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
423 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);427 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
424 } else {428 } else {
425- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_ONE, T2, THREAD_DIM>>(429+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_ONE, T2, THREAD_DIM>>(
426- AscendC::Simt::Dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,430+ dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,
427 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),431 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
428 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);432 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
429 }433 }
430 } else if (indexedDimNum_ == DIM_NUMS_TWO) {434 } else if (indexedDimNum_ == DIM_NUMS_TWO) {
431 if (isSmallIndexSize) {435 if (isSmallIndexSize) {
432- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_TWO, T2, THREAD_DIM_SMALL>>(436+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_TWO, T2, THREAD_DIM_SMALL>>(
433- AscendC::Simt::Dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,437+ dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,
434 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),438 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
435 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);439 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
436 } else {440 } else {
437- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_TWO, T2, THREAD_DIM>>(441+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_TWO, T2, THREAD_DIM>>(
438- AscendC::Simt::Dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,442+ dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,
439 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),443 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
440 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);444 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
441 }445 }
442 } else if (indexedDimNum_ == DIM_NUMS_THREE) {446 } else if (indexedDimNum_ == DIM_NUMS_THREE) {
443 if(isSmallIndexSize) {447 if(isSmallIndexSize) {
444- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_THREE, T2, THREAD_DIM_SMALL>>(448+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_THREE, T2, THREAD_DIM_SMALL>>(
445- AscendC::Simt::Dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,449+ dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,
446 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),450 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
447 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);451 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
448 } else {452 } else {
449- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_THREE, T2, THREAD_DIM>>(453+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_THREE, T2, THREAD_DIM>>(
450- AscendC::Simt::Dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,454+ dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,
451 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),455 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
452 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);456 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
453 }457 }
454 } else if (indexedDimNum_ == DIM_NUMS_FOUR) {458 } else if (indexedDimNum_ == DIM_NUMS_FOUR) {
455 if(isSmallIndexSize) {459 if(isSmallIndexSize) {
456- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_FOUR, T2, THREAD_DIM_SMALL>>(460+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_FOUR, T2, THREAD_DIM_SMALL>>(
457- AscendC::Simt::Dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,461+ dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,
458 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),462 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
459 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);463 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
460 } else {464 } else {
461- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_FOUR, T2, THREAD_DIM>>(465+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_FOUR, T2, THREAD_DIM>>(
462- AscendC::Simt::Dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,466+ dim3{threadDim_}, blockId_, outputLength_, blockNums_, secondThirdLoopLength_,
463 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),467 formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
464 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);468 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
465 }469 }
466 } else if (indexedDimNum_ == DIM_NUMS_FIVE) {470 } else if (indexedDimNum_ == DIM_NUMS_FIVE) {
467 if (isSmallIndexSize) {471 if (isSmallIndexSize) {
468- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_FIVE, T2, SMALL_THREAD_DIM>>(472+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_FIVE, T2, SMALL_THREAD_DIM>>(
469- AscendC::Simt::Dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,473+ dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,
470 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),474 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
471 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);475 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
472 } else {476 } else {
473- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_FIVE, T2, SMALL_THREAD_DIM_LAUNCH_BOUND>>(477+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_FIVE, T2, SMALL_THREAD_DIM_LAUNCH_BOUND>>(
474- AscendC::Simt::Dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,478+ dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,
475 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),479 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
476 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);480 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
477 }481 }
478 } else if (indexedDimNum_ == DIM_NUMS_SIX) {482 } else if (indexedDimNum_ == DIM_NUMS_SIX) {
479 if (isSmallIndexSize) {483 if (isSmallIndexSize) {
480- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_SIX, T2, SMALL_THREAD_DIM>>(484+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_SIX, T2, SMALL_THREAD_DIM>>(
481- AscendC::Simt::Dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,485+ dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,
482 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),486 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
483 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);487 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
484 } else {488 } else {
485- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_SIX, T2, SMALL_THREAD_DIM_LAUNCH_BOUND>>(489+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_SIX, T2, SMALL_THREAD_DIM_LAUNCH_BOUND>>(
486- AscendC::Simt::Dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,490+ dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,
487 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),491 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
488 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);492 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
489 }493 }
490 } else if (indexedDimNum_ == DIM_NUMS_SEVEN) {494 } else if (indexedDimNum_ == DIM_NUMS_SEVEN) {
491 if (isSmallIndexSize) {495 if (isSmallIndexSize) {
492- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_SEVEN, T2, SMALL_THREAD_DIM>>(496+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_SEVEN, T2, SMALL_THREAD_DIM>>(
493- AscendC::Simt::Dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,497+ dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,
494 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),498 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
495 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);499 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
496 } else {500 } else {
497- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_SEVEN, T2, SMALL_THREAD_DIM_LAUNCH_BOUND>>(501+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_SEVEN, T2, SMALL_THREAD_DIM_LAUNCH_BOUND>>(
498- AscendC::Simt::Dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,502+ dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,
499 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),503 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
500 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);504 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
501 }505 }
502 } else if (indexedDimNum_ == DIM_NUMS_EIGHT) {506 } else if (indexedDimNum_ == DIM_NUMS_EIGHT) {
503 if (isSmallIndexSize) {507 if (isSmallIndexSize) {
504- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_EIGHT, T2, SMALL_THREAD_DIM>>(508+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_EIGHT, T2, SMALL_THREAD_DIM>>(
505- AscendC::Simt::Dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,509+ dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,
506 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),510 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
507 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);511 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
508 } else {512 } else {
509- AscendC::Simt::VF_CALL<SimtComputeContinue<T, F, P, DIM_NUMS_EIGHT, T2, SMALL_THREAD_DIM_LAUNCH_BOUND>>(513+ asc_vf_call<SimtComputeContinue<T, F, P, DIM_NUMS_EIGHT, T2, SMALL_THREAD_DIM_LAUNCH_BOUND>>(
510- AscendC::Simt::Dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,514+ dim3{smallThreadDim_}, blockId_, outputLength_, blockNums_,
511 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),515 secondThirdLoopLength_, formerNonIndexStride_, thirdLoopLength_, (__gm__ T*)outputGm_.GetPhyAddr(),
512 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);516 (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
513 }517 }
514 }518 }
515 } else {519 } else {
516 if (indexedDimNum_ == DIM_NUMS_ONE && nonIndexedDimNum_ == DIM_NUMS_ONE) {520 if (indexedDimNum_ == DIM_NUMS_ONE && nonIndexedDimNum_ == DIM_NUMS_ONE) {
517- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_ONE, DIM_NUMS_ONE, T2, THREAD_DIM>>(521+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_ONE, DIM_NUMS_ONE, T2, THREAD_DIM>>(
518- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,522+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
519 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);523 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
520 } else if (indexedDimNum_ == DIM_NUMS_ONE && nonIndexedDimNum_ == DIM_NUMS_TWO) {524 } else if (indexedDimNum_ == DIM_NUMS_ONE && nonIndexedDimNum_ == DIM_NUMS_TWO) {
521- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_ONE, DIM_NUMS_TWO, T2, THREAD_DIM>>(525+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_ONE, DIM_NUMS_TWO, T2, THREAD_DIM>>(
522- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,526+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
523 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);527 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
524 } else if (indexedDimNum_ == DIM_NUMS_ONE && nonIndexedDimNum_ == DIM_NUMS_THREE) {528 } else if (indexedDimNum_ == DIM_NUMS_ONE && nonIndexedDimNum_ == DIM_NUMS_THREE) {
525- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_ONE, DIM_NUMS_THREE, T2, THREAD_DIM>>(529+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_ONE, DIM_NUMS_THREE, T2, THREAD_DIM>>(
526- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,530+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
527 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);531 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
528 } else if (indexedDimNum_ == DIM_NUMS_ONE && nonIndexedDimNum_ == DIM_NUMS_FOUR) {532 } else if (indexedDimNum_ == DIM_NUMS_ONE && nonIndexedDimNum_ == DIM_NUMS_FOUR) {
529- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_ONE, DIM_NUMS_FOUR, T2, THREAD_DIM>>(533+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_ONE, DIM_NUMS_FOUR, T2, THREAD_DIM>>(
530- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,534+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
531 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);535 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
532 } else if (indexedDimNum_ == DIM_NUMS_TWO && nonIndexedDimNum_ == DIM_NUMS_ONE) {536 } else if (indexedDimNum_ == DIM_NUMS_TWO && nonIndexedDimNum_ == DIM_NUMS_ONE) {
533- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_TWO, DIM_NUMS_ONE, T2, THREAD_DIM>>(537+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_TWO, DIM_NUMS_ONE, T2, THREAD_DIM>>(
534- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,538+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
535 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);539 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
536 } else if (indexedDimNum_ == DIM_NUMS_TWO && nonIndexedDimNum_ == DIM_NUMS_TWO) {540 } else if (indexedDimNum_ == DIM_NUMS_TWO && nonIndexedDimNum_ == DIM_NUMS_TWO) {
537- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_TWO, DIM_NUMS_TWO, T2, THREAD_DIM>>(541+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_TWO, DIM_NUMS_TWO, T2, THREAD_DIM>>(
538- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,542+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
539 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);543 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
540 } else if (indexedDimNum_ == DIM_NUMS_TWO && nonIndexedDimNum_ == DIM_NUMS_THREE) {544 } else if (indexedDimNum_ == DIM_NUMS_TWO && nonIndexedDimNum_ == DIM_NUMS_THREE) {
541- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_TWO, DIM_NUMS_THREE, T2, THREAD_DIM>>(545+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_TWO, DIM_NUMS_THREE, T2, THREAD_DIM>>(
542- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,546+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
543 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);547 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
544 } else if (indexedDimNum_ == DIM_NUMS_TWO && nonIndexedDimNum_ == DIM_NUMS_FOUR) {548 } else if (indexedDimNum_ == DIM_NUMS_TWO && nonIndexedDimNum_ == DIM_NUMS_FOUR) {
545- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_TWO, DIM_NUMS_FOUR, T2, THREAD_DIM>>(549+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_TWO, DIM_NUMS_FOUR, T2, THREAD_DIM>>(
546- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,550+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
547 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);551 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
548 } else if (indexedDimNum_ == DIM_NUMS_THREE && nonIndexedDimNum_ == DIM_NUMS_ONE) {552 } else if (indexedDimNum_ == DIM_NUMS_THREE && nonIndexedDimNum_ == DIM_NUMS_ONE) {
549- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_THREE, DIM_NUMS_ONE, T2, THREAD_DIM>>(553+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_THREE, DIM_NUMS_ONE, T2, THREAD_DIM>>(
550- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,554+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
551 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);555 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
552 } else if (indexedDimNum_ == DIM_NUMS_THREE && nonIndexedDimNum_ == DIM_NUMS_TWO) {556 } else if (indexedDimNum_ == DIM_NUMS_THREE && nonIndexedDimNum_ == DIM_NUMS_TWO) {
553- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_THREE, DIM_NUMS_TWO, T2, THREAD_DIM>>(557+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_THREE, DIM_NUMS_TWO, T2, THREAD_DIM>>(
554- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,558+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
555 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);559 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
556 } else if (indexedDimNum_ == DIM_NUMS_THREE && nonIndexedDimNum_ == DIM_NUMS_THREE) {560 } else if (indexedDimNum_ == DIM_NUMS_THREE && nonIndexedDimNum_ == DIM_NUMS_THREE) {
557- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_THREE, DIM_NUMS_THREE, T2, THREAD_DIM>>(561+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_THREE, DIM_NUMS_THREE, T2, THREAD_DIM>>(
558- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,562+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
559 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);563 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
560 } else if (indexedDimNum_ == DIM_NUMS_THREE && nonIndexedDimNum_ == DIM_NUMS_FOUR) {564 } else if (indexedDimNum_ == DIM_NUMS_THREE && nonIndexedDimNum_ == DIM_NUMS_FOUR) {
561- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_THREE, DIM_NUMS_FOUR, T2, THREAD_DIM>>(565+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_THREE, DIM_NUMS_FOUR, T2, THREAD_DIM>>(
562- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,566+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
563 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);567 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
564 } else if (indexedDimNum_ == DIM_NUMS_FOUR && nonIndexedDimNum_ == DIM_NUMS_ONE) {568 } else if (indexedDimNum_ == DIM_NUMS_FOUR && nonIndexedDimNum_ == DIM_NUMS_ONE) {
565- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_FOUR, DIM_NUMS_ONE, T2, THREAD_DIM>>(569+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_FOUR, DIM_NUMS_ONE, T2, THREAD_DIM>>(
566- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,570+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
567 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);571 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
568 } else if (indexedDimNum_ == DIM_NUMS_FOUR && nonIndexedDimNum_ == DIM_NUMS_TWO) {572 } else if (indexedDimNum_ == DIM_NUMS_FOUR && nonIndexedDimNum_ == DIM_NUMS_TWO) {
569- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_FOUR, DIM_NUMS_TWO, T2, THREAD_DIM>>(573+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_FOUR, DIM_NUMS_TWO, T2, THREAD_DIM>>(
570- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,574+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
571 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);575 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
572 } else if (indexedDimNum_ == DIM_NUMS_FOUR && nonIndexedDimNum_ == DIM_NUMS_THREE) {576 } else if (indexedDimNum_ == DIM_NUMS_FOUR && nonIndexedDimNum_ == DIM_NUMS_THREE) {
573- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_FOUR, DIM_NUMS_THREE, T2, THREAD_DIM>>(577+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_FOUR, DIM_NUMS_THREE, T2, THREAD_DIM>>(
574- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,578+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
575 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);579 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
576 } else if (indexedDimNum_ == DIM_NUMS_FOUR && nonIndexedDimNum_ == DIM_NUMS_FOUR) {580 } else if (indexedDimNum_ == DIM_NUMS_FOUR && nonIndexedDimNum_ == DIM_NUMS_FOUR) {
577- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_FOUR, DIM_NUMS_FOUR, T2, THREAD_DIM>>(581+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_FOUR, DIM_NUMS_FOUR, T2, THREAD_DIM>>(
578- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,582+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
579 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);583 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
580 } else if (indexedDimNum_ == DIM_NUMS_FIVE && nonIndexedDimNum_ == DIM_NUMS_TWO) {584 } else if (indexedDimNum_ == DIM_NUMS_FIVE && nonIndexedDimNum_ == DIM_NUMS_TWO) {
581 // add missing non-continuous scenarios585 // add missing non-continuous scenarios
582- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_FIVE, DIM_NUMS_TWO, T2, THREAD_DIM>>(586+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_FIVE, DIM_NUMS_TWO, T2, THREAD_DIM>>(
583- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,587+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
584 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);588 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
585 } else if (indexedDimNum_ == DIM_NUMS_FIVE && nonIndexedDimNum_ == DIM_NUMS_THREE) {589 } else if (indexedDimNum_ == DIM_NUMS_FIVE && nonIndexedDimNum_ == DIM_NUMS_THREE) {
586- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_FIVE, DIM_NUMS_THREE, T2, THREAD_DIM>>(590+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_FIVE, DIM_NUMS_THREE, T2, THREAD_DIM>>(
587- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,591+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
588 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);592 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
589 } else if (indexedDimNum_ == DIM_NUMS_THREE && nonIndexedDimNum_ == DIM_NUMS_FIVE) {593 } else if (indexedDimNum_ == DIM_NUMS_THREE && nonIndexedDimNum_ == DIM_NUMS_FIVE) {
590- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_THREE, DIM_NUMS_FIVE, T2, THREAD_DIM>>(594+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_THREE, DIM_NUMS_FIVE, T2, THREAD_DIM>>(
591- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,595+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
592 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);596 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
593 } else if (indexedDimNum_ == DIM_NUMS_TWO && nonIndexedDimNum_ == DIM_NUMS_FIVE) {597 } else if (indexedDimNum_ == DIM_NUMS_TWO && nonIndexedDimNum_ == DIM_NUMS_FIVE) {
594- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_TWO, DIM_NUMS_FIVE, T2, THREAD_DIM>>(598+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_TWO, DIM_NUMS_FIVE, T2, THREAD_DIM>>(
595- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,599+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
596 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);600 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
597 } else if (indexedDimNum_ == DIM_NUMS_TWO && nonIndexedDimNum_ == DIM_NUMS_SIX) {601 } else if (indexedDimNum_ == DIM_NUMS_TWO && nonIndexedDimNum_ == DIM_NUMS_SIX) {
598- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_TWO, DIM_NUMS_SIX, T2, THREAD_DIM>>(602+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_TWO, DIM_NUMS_SIX, T2, THREAD_DIM>>(
599- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,603+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
600 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);604 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
601 } else if (indexedDimNum_ == DIM_NUMS_SIX && nonIndexedDimNum_ == DIM_NUMS_TWO) {605 } else if (indexedDimNum_ == DIM_NUMS_SIX && nonIndexedDimNum_ == DIM_NUMS_TWO) {
602- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_SIX, DIM_NUMS_TWO, T2, THREAD_DIM>>(606+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_SIX, DIM_NUMS_TWO, T2, THREAD_DIM>>(
603- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,607+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
604 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);608 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
605 } else if (indexedDimNum_ == DIM_NUMS_SIX && nonIndexedDimNum_ == DIM_NUMS_ONE) {609 } else if (indexedDimNum_ == DIM_NUMS_SIX && nonIndexedDimNum_ == DIM_NUMS_ONE) {
606- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_SIX, DIM_NUMS_ONE, T2, THREAD_DIM>>(610+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_SIX, DIM_NUMS_ONE, T2, THREAD_DIM>>(
607- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,611+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
608 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);612 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
609 } else if (indexedDimNum_ == DIM_NUMS_SEVEN && nonIndexedDimNum_ == DIM_NUMS_ONE) {613 } else if (indexedDimNum_ == DIM_NUMS_SEVEN && nonIndexedDimNum_ == DIM_NUMS_ONE) {
610- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_SEVEN, DIM_NUMS_ONE, T2, THREAD_DIM>>(614+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_SEVEN, DIM_NUMS_ONE, T2, THREAD_DIM>>(
611- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,615+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
612 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);616 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
613 } else if (indexedDimNum_ == DIM_NUMS_FIVE && nonIndexedDimNum_ == DIM_NUMS_ONE) {617 } else if (indexedDimNum_ == DIM_NUMS_FIVE && nonIndexedDimNum_ == DIM_NUMS_ONE) {
614- AscendC::Simt::VF_CALL<SimtComputeNonContinue<T, F, P, DIM_NUMS_FIVE, DIM_NUMS_ONE, T2, THREAD_DIM>>(618+ asc_vf_call<SimtComputeNonContinue<T, F, P, DIM_NUMS_FIVE, DIM_NUMS_ONE, T2, THREAD_DIM>>(
615- AscendC::Simt::Dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,619+ dim3{THREAD_DIM}, blockId_, outputLength_, blockNums_, innerLoopLength_, m3_, shift3_,
616 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);620 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
617 }621 }
618 }622 }
Mindex/index/op_kernel/arch35/index_no_continuous.h+83-82
@@ -18,6 +18,7 @@
18#define ASCENDC_NOCONTINUOUS_INDEX_SIMT_H_18#define ASCENDC_NOCONTINUOUS_INDEX_SIMT_H_
19 19 
20#include "kernel_operator.h"20#include "kernel_operator.h"
21+#include "simt_api/asc_simt.h"
21#include "op_kernel/platform_util.h"22#include "op_kernel/platform_util.h"
22 23 
23 24 
@@ -69,8 +70,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_LAUNCH_BOUND) inline void SimtCom
69 localIndexList[i] = (__gm__ P*)calcParamsPtr_->indexList_[i];70 localIndexList[i] = (__gm__ P*)calcParamsPtr_->indexList_[i];
70 }71 }
71 72 
72- for (T2 i = blockId_ * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < outputLength_;73+ for (T2 i = blockId_ * blockDim.x + threadIdx.x; i < outputLength_;
73- i = i + blockNums_ * Simt::GetThreadNum()) {74+ i = i + blockNums_ * blockDim.x) {
74 T2 inputIndex = 0;75 T2 inputIndex = 0;
75 T2 length = i;76 T2 length = i;
76 T2 index = calcParamsPtr_-> isIndexPut_ == false ? i : 0;77 T2 index = calcParamsPtr_-> isIndexPut_ == false ? i : 0;
@@ -279,164 +280,164 @@ __aicore__ inline void KernelIndexNoContiguous<T, F, P, T2>::Process()
279{280{
280 DataSyncBarrier<MemDsbT::UB>();281 DataSyncBarrier<MemDsbT::UB>();
281 if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_ONE) {282 if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_ONE) {
282- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_ONE, T2>>(283+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_ONE, T2>>(
283- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,284+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
284 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_); 285 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
285 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_TWO) {286 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_TWO) {
286- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_TWO, T2>>(287+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_TWO, T2>>(
287- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,288+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
288 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_); 289 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
289 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_THREE) {290 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_THREE) {
290- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_THREE, T2>>(291+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_THREE, T2>>(
291- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,292+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
292 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);293 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
293 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {294 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
294- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_FOUR, T2>>(295+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_FOUR, T2>>(
295- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,296+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
296 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);297 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
297 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_ONE) {298 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_ONE) {
298- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_ONE, T2>>(299+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_ONE, T2>>(
299- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,300+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
300 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);301 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
301 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_TWO) {302 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_TWO) {
302- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_TWO, T2>>(303+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_TWO, T2>>(
303- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,304+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
304 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);305 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
305 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_THREE) {306 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_THREE) {
306- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_THREE, T2>>(307+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_THREE, T2>>(
307- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,308+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
308 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);309 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
309 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {310 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
310- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_FOUR, T2>>(311+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_FOUR, T2>>(
311- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,312+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
312 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);313 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
313 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_ONE) {314 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_ONE) {
314- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_ONE, T2>>(315+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_ONE, T2>>(
315- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,316+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
316 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);317 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
317 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_TWO) {318 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_TWO) {
318- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_TWO, T2>>(319+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_TWO, T2>>(
319- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,320+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
320 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);321 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
321 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_THREE) {322 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_THREE) {
322- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_THREE, T2>>(323+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_THREE, T2>>(
323- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,324+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
324 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);325 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
325 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {326 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
326- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_FOUR, T2>>(327+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_FOUR, T2>>(
327- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,328+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
328 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);329 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
329 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_ONE) {330 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_ONE) {
330- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_ONE, T2>>(331+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_ONE, T2>>(
331- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,332+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
332 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);333 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
333 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_TWO) {334 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_TWO) {
334- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_TWO, T2>>(335+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_TWO, T2>>(
335- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,336+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
336 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);337 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
337 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_THREE) {338 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_THREE) {
338- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_THREE, T2>>(339+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_THREE, T2>>(
339- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,340+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
340 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);341 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
341 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {342 } else if (indexedNum_ == NOCON_COUNT_NUMS_ONE && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
342- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_FOUR, T2>>(343+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_ONE, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_FOUR, T2>>(
343- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,344+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
344 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_); 345 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
345 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_TWO) {346 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_TWO) {
346- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_TWO, T2>>(347+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_TWO, T2>>(
347- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,348+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
348 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_); 349 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
349 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_THREE) {350 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_THREE) {
350- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_THREE, T2>>(351+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_THREE, T2>>(
351- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,352+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
352 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);353 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
353 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {354 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
354- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_FOUR, T2>>(355+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_FOUR, T2>>(
355- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,356+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
356 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);357 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
357 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_TWO) {358 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_TWO) {
358- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_TWO, T2>>(359+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_TWO, T2>>(
359- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,360+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
360 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);361 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
361 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_THREE) {362 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_THREE) {
362- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_THREE, T2>>(363+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_THREE, T2>>(
363- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,364+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
364 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);365 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
365 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {366 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
366- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_FOUR, T2>>(367+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_FOUR, T2>>(
367- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,368+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
368 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);369 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
369 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_TWO) {370 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_TWO) {
370- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_TWO, T2>>(371+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_TWO, T2>>(
371- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,372+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
372 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);373 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
373 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_THREE) {374 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_THREE) {
374- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_THREE, T2>>(375+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_THREE, T2>>(
375- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,376+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
376 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);377 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
377 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {378 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
378- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_FOUR, T2>>(379+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_FOUR, T2>>(
379- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,380+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
380 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);381 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
381 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_TWO) {382 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_TWO) {
382- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_TWO, T2>>(383+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_TWO, T2>>(
383- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,384+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
384 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);385 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
385 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_THREE) {386 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_THREE) {
386- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_THREE, T2>>(387+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_THREE, T2>>(
387- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,388+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
388 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);389 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
389 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {390 } else if (indexedNum_ == NOCON_COUNT_NUMS_TWO && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
390- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_FOUR, T2>>(391+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_TWO, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_FOUR, T2>>(
391- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,392+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
392 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);393 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
393 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_THREE) {394 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_THREE) {
394- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_THREE, T2>>(395+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_THREE, T2>>(
395- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,396+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
396 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);397 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
397 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {398 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
398- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_FOUR, T2>>(399+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_FOUR, T2>>(
399- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,400+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
400 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);401 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
401 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_THREE) {402 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_THREE) {
402- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_THREE, T2>>(403+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_THREE, T2>>(
403- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,404+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
404 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);405 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
405 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {406 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
406- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_FOUR, T2>>(407+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_FOUR, T2>>(
407- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,408+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
408 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);409 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
409 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_THREE) {410 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_THREE) {
410- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_THREE, T2>>(411+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_THREE, T2>>(
411- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,412+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
412 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);413 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
413 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {414 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
414- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_FOUR, T2>>(415+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_FOUR, T2>>(
415- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,416+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
416 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);417 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
417 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_THREE) {418 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_THREE) {
418- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_THREE, T2>>(419+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_THREE, T2>>(
419- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,420+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
420 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);421 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
421 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {422 } else if (indexedNum_ == NOCON_COUNT_NUMS_THREE && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
422- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_FOUR, T2>>(423+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_THREE, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_FOUR, T2>>(
423- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,424+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
424 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);425 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
425 } else if (indexedNum_ == NOCON_COUNT_NUMS_FOUR && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {426 } else if (indexedNum_ == NOCON_COUNT_NUMS_FOUR && indexedDimNum_ == NOCON_DIM_NUMS_ONE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
426- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_FOUR, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_FOUR, T2>>(427+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_FOUR, NOCON_DIM_NUMS_ONE, NOCON_DIM_NUMS_FOUR, T2>>(
427- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,428+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
428 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);429 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
429 } else if (indexedNum_ == NOCON_COUNT_NUMS_FOUR && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {430 } else if (indexedNum_ == NOCON_COUNT_NUMS_FOUR && indexedDimNum_ == NOCON_DIM_NUMS_TWO && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
430- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_FOUR, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_FOUR, T2>>(431+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_FOUR, NOCON_DIM_NUMS_TWO, NOCON_DIM_NUMS_FOUR, T2>>(
431- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,432+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
432 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);433 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
433 } else if (indexedNum_ == NOCON_COUNT_NUMS_FOUR && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {434 } else if (indexedNum_ == NOCON_COUNT_NUMS_FOUR && indexedDimNum_ == NOCON_DIM_NUMS_THREE && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
434- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_FOUR, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_FOUR, T2>>(435+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_FOUR, NOCON_DIM_NUMS_THREE, NOCON_DIM_NUMS_FOUR, T2>>(
435- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,436+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
436 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);437 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
437 } else if (indexedNum_ == NOCON_COUNT_NUMS_FOUR && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {438 } else if (indexedNum_ == NOCON_COUNT_NUMS_FOUR && indexedDimNum_ == NOCON_DIM_NUMS_FOUR && inputDimNum_ == NOCON_DIM_NUMS_FOUR) {
438- AscendC::Simt::VF_CALL<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_FOUR, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_FOUR, T2>>(439+ asc_vf_call<SimtComputeNoContinuous<T, F, P, NOCON_COUNT_NUMS_FOUR, NOCON_DIM_NUMS_FOUR, NOCON_DIM_NUMS_FOUR, T2>>(
439- AscendC::Simt::Dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,440+ dim3{NONCON_THREAD_DIM}, blockId_, outputLength_, blockNums_, firstIndexedDim_,
440 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);441 (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(), calcParamsPtr_);
441 }442 }
442}443}
Mindex/index/op_kernel/arch35/index_perf.h+5-4
@@ -17,6 +17,7 @@
17#define ASCENDC_INDEX_PERF_H_17#define ASCENDC_INDEX_PERF_H_
18 18 
19#include "kernel_operator.h"19#include "kernel_operator.h"
20+#include "simt_api/asc_simt.h"
20#include "op_kernel/platform_util.h"21#include "op_kernel/platform_util.h"
21 22 
22namespace Index {23namespace Index {
@@ -40,8 +41,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_LAUNCH) inline void SimtCompute(
40 int64_t inputShape0 = static_cast<int64_t>(tilingData->inputShape0);41 int64_t inputShape0 = static_cast<int64_t>(tilingData->inputShape0);
41 T2 inputShape1 = static_cast<T2>(tilingData->inputShape1);42 T2 inputShape1 = static_cast<T2>(tilingData->inputShape1);
42 43 
43- for (T2 i = Simt::GetBlockIdx() * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < outputLength;44+ for (T2 i = blockIdx.x * blockDim.x + threadIdx.x; i < outputLength;
44- i = i + Simt::GetBlockNum() * Simt::GetThreadNum()) {45+ i = i + gridDim.x * blockDim.x) {
45 int64_t curIdx0 = indexTensor0[i];46 int64_t curIdx0 = indexTensor0[i];
46 curIdx0 = curIdx0 < 0 ? (curIdx0 + inputShape0) : curIdx0;47 curIdx0 = curIdx0 < 0 ? (curIdx0 + inputShape0) : curIdx0;
47 int64_t curIdx1 = indexTensor1[i];48 int64_t curIdx1 = indexTensor1[i];
@@ -68,8 +69,8 @@ __aicore__ inline void Process(
68 // get dynamci input tensor69 // get dynamci input tensor
69 __gm__ P* indexTensor0 = GetInputTensorAddr<P>(indices, 0);70 __gm__ P* indexTensor0 = GetInputTensorAddr<P>(indices, 0);
70 __gm__ P* indexTensor1 = GetInputTensorAddr<P>(indices, 1);71 __gm__ P* indexTensor1 = GetInputTensorAddr<P>(indices, 1);
71- AscendC::Simt::VF_CALL<SimtCompute<T, P, T2>>(72+ asc_vf_call<SimtCompute<T, P, T2>>(
72- AscendC::Simt::Dim3{USED_THREAD}, output, inputX, indexTensor0, indexTensor1, tiling);73+ dim3{USED_THREAD}, output, inputX, indexTensor0, indexTensor1, tiling);
73}74}
74} // namespace Index75} // namespace Index
75 76 
Mindex/index/op_kernel/arch35/index_perf_no_continuous.h+31-30
@@ -17,6 +17,7 @@
17#define INDEX_PERF_NO_CONTINUOUS_H_17#define INDEX_PERF_NO_CONTINUOUS_H_
18 18 
19#include "kernel_operator.h"19#include "kernel_operator.h"
20+#include "simt_api/asc_simt.h"
20#include "op_kernel/platform_util.h"21#include "op_kernel/platform_util.h"
21#include "op_kernel/math_util.h"22#include "op_kernel/math_util.h"
22#include "kernel_operator_list_tensor_intf.h"23#include "kernel_operator_list_tensor_intf.h"
@@ -50,8 +51,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_LAUNCH_BOUND) inline void SimtDim
50 __gm__ T* outputGm_, __gm__ T* inputXGm_, __gm__ P* indexTensor0, __gm__ P* indexTensor1,51 __gm__ T* outputGm_, __gm__ T* inputXGm_, __gm__ P* indexTensor0, __gm__ P* indexTensor1,
51 __ubuf__ SimtParams *simtParamsPtr_, uint64_t outputLength)52 __ubuf__ SimtParams *simtParamsPtr_, uint64_t outputLength)
52{53{
53- for (T2 i = Simt::GetBlockIdx() * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < outputLength;54+ for (T2 i = blockIdx.x * blockDim.x + threadIdx.x; i < outputLength;
54- i = i + Simt::GetBlockNum() * Simt::GetThreadNum()) {55+ i = i + gridDim.x * blockDim.x) {
55 T2 indexDim[NOCON_IDX_MAX_DIM_NUM] = {0, 0, 0, 0};56 T2 indexDim[NOCON_IDX_MAX_DIM_NUM] = {0, 0, 0, 0};
56 T2 offset = i;57 T2 offset = i;
57 T2 inputIndex = 0;58 T2 inputIndex = 0;
@@ -92,8 +93,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_LAUNCH_BOUND) inline void SimtDim
92 __gm__ T* outputGm_, __gm__ T* inputXGm_, __gm__ P* indexTensor0, __gm__ P* indexTensor1,93 __gm__ T* outputGm_, __gm__ T* inputXGm_, __gm__ P* indexTensor0, __gm__ P* indexTensor1,
93 __gm__ P* indexTensor2, __ubuf__ SimtParams *simtParamsPtr_, uint64_t outputLength)94 __gm__ P* indexTensor2, __ubuf__ SimtParams *simtParamsPtr_, uint64_t outputLength)
94{95{
95- for (T2 i = Simt::GetBlockIdx() * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < outputLength;96+ for (T2 i = blockIdx.x * blockDim.x + threadIdx.x; i < outputLength;
96- i = i + Simt::GetBlockNum() * Simt::GetThreadNum()) {97+ i = i + gridDim.x * blockDim.x) {
97 T2 indexDim[NOCON_IDX_MAX_DIM_NUM] = {0, 0, 0, 0};98 T2 indexDim[NOCON_IDX_MAX_DIM_NUM] = {0, 0, 0, 0};
98 T2 offset = i;99 T2 offset = i;
99 T2 inputIndex = 0;100 T2 inputIndex = 0;
@@ -145,8 +146,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM_LAUNCH_BOUND) inline void SimtDim
145 __gm__ P* indexTensor2, __gm__ P* indexTensor3, __ubuf__ SimtParams *simtParamsPtr_,146 __gm__ P* indexTensor2, __gm__ P* indexTensor3, __ubuf__ SimtParams *simtParamsPtr_,
146 uint64_t outputLength)147 uint64_t outputLength)
147{148{
148- for (T2 i = Simt::GetBlockIdx() * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < outputLength;149+ for (T2 i = blockIdx.x * blockDim.x + threadIdx.x; i < outputLength;
149- i = i + Simt::GetBlockNum() * Simt::GetThreadNum()) {150+ i = i + gridDim.x * blockDim.x) {
150 T2 indexDim[NOCON_IDX_MAX_DIM_NUM] = {0, 0, 0, 0};151 T2 indexDim[NOCON_IDX_MAX_DIM_NUM] = {0, 0, 0, 0};
151 T2 offset = i;152 T2 offset = i;
152 T2 inputIndex = 0;153 T2 inputIndex = 0;
@@ -276,68 +277,68 @@ __aicore__ inline void KernelIndexNoContiguousPerf<T, P, T2>::Process()
276 __gm__ P* indexTensor1 = GetInputTensorAddr(indices_, LIST_INDEX_ONE);277 __gm__ P* indexTensor1 = GetInputTensorAddr(indices_, LIST_INDEX_ONE);
277 if (inputDimNum_ == INPUTDIMNUMTWO) {278 if (inputDimNum_ == INPUTDIMNUMTWO) {
278 if (indexedDimNum_ == INDEX_DIM_NUM_ONE) {279 if (indexedDimNum_ == INDEX_DIM_NUM_ONE) {
279- AscendC::Simt::VF_CALL<SimtDim2Compute<T, P, T2, INDEX_DIM_NUM_ONE>>(280+ asc_vf_call<SimtDim2Compute<T, P, T2, INDEX_DIM_NUM_ONE>>(
280- AscendC::Simt::Dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),281+ dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
281 indexTensor0, indexTensor1, simtParamsPtr_, outputLength_);282 indexTensor0, indexTensor1, simtParamsPtr_, outputLength_);
282 } else if (indexedDimNum_ == INDEX_DIM_NUM_TWO) {283 } else if (indexedDimNum_ == INDEX_DIM_NUM_TWO) {
283- AscendC::Simt::VF_CALL<SimtDim2Compute<T, P, T2, INDEX_DIM_NUM_TWO>>(284+ asc_vf_call<SimtDim2Compute<T, P, T2, INDEX_DIM_NUM_TWO>>(
284- AscendC::Simt::Dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),285+ dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
285 indexTensor0, indexTensor1, simtParamsPtr_, outputLength_);286 indexTensor0, indexTensor1, simtParamsPtr_, outputLength_);
286 } else if (indexedDimNum_ == INDEX_DIM_NUM_THREE) {287 } else if (indexedDimNum_ == INDEX_DIM_NUM_THREE) {
287- AscendC::Simt::VF_CALL<SimtDim2Compute<T, P, T2, INDEX_DIM_NUM_THREE>>(288+ asc_vf_call<SimtDim2Compute<T, P, T2, INDEX_DIM_NUM_THREE>>(
288- AscendC::Simt::Dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),289+ dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
289 indexTensor0, indexTensor1, simtParamsPtr_, outputLength_);290 indexTensor0, indexTensor1, simtParamsPtr_, outputLength_);
290 } else if (indexedDimNum_ == INDEX_DIM_NUM_FOUR) {291 } else if (indexedDimNum_ == INDEX_DIM_NUM_FOUR) {
291- AscendC::Simt::VF_CALL<SimtDim2Compute<T, P, T2, INDEX_DIM_NUM_FOUR>>(292+ asc_vf_call<SimtDim2Compute<T, P, T2, INDEX_DIM_NUM_FOUR>>(
292- AscendC::Simt::Dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),293+ dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
293 indexTensor0, indexTensor1, simtParamsPtr_, outputLength_);294 indexTensor0, indexTensor1, simtParamsPtr_, outputLength_);
294 }295 }
295 } else if (inputDimNum_ == INPUTDIMNUMTHREE) {296 } else if (inputDimNum_ == INPUTDIMNUMTHREE) {
296 if (indexedDimNum_ == INDEX_DIM_NUM_ONE) {297 if (indexedDimNum_ == INDEX_DIM_NUM_ONE) {
297 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);298 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);
298- AscendC::Simt::VF_CALL<SimtDim3Compute<T, P, T2, INDEX_DIM_NUM_ONE>>(299+ asc_vf_call<SimtDim3Compute<T, P, T2, INDEX_DIM_NUM_ONE>>(
299- AscendC::Simt::Dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),300+ dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
300 indexTensor0, indexTensor1, indexTensor2, simtParamsPtr_, outputLength_);301 indexTensor0, indexTensor1, indexTensor2, simtParamsPtr_, outputLength_);
301 } else if (indexedDimNum_ == INDEX_DIM_NUM_TWO) {302 } else if (indexedDimNum_ == INDEX_DIM_NUM_TWO) {
302 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);303 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);
303- AscendC::Simt::VF_CALL<SimtDim3Compute<T, P, T2, INDEX_DIM_NUM_TWO>>(304+ asc_vf_call<SimtDim3Compute<T, P, T2, INDEX_DIM_NUM_TWO>>(
304- AscendC::Simt::Dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),305+ dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
305 indexTensor0, indexTensor1, indexTensor2, simtParamsPtr_, outputLength_);306 indexTensor0, indexTensor1, indexTensor2, simtParamsPtr_, outputLength_);
306 } else if (indexedDimNum_ == INDEX_DIM_NUM_THREE) {307 } else if (indexedDimNum_ == INDEX_DIM_NUM_THREE) {
307 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);308 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);
308- AscendC::Simt::VF_CALL<SimtDim3Compute<T, P, T2, INDEX_DIM_NUM_THREE>>(309+ asc_vf_call<SimtDim3Compute<T, P, T2, INDEX_DIM_NUM_THREE>>(
309- AscendC::Simt::Dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),310+ dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
310 indexTensor0, indexTensor1, indexTensor2, simtParamsPtr_, outputLength_);311 indexTensor0, indexTensor1, indexTensor2, simtParamsPtr_, outputLength_);
311 } else if (indexedDimNum_ == INDEX_DIM_NUM_FOUR) {312 } else if (indexedDimNum_ == INDEX_DIM_NUM_FOUR) {
312 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);313 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);
313- AscendC::Simt::VF_CALL<SimtDim3Compute<T, P, T2, INDEX_DIM_NUM_FOUR>>(314+ asc_vf_call<SimtDim3Compute<T, P, T2, INDEX_DIM_NUM_FOUR>>(
314- AscendC::Simt::Dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),315+ dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
315 indexTensor0, indexTensor1, indexTensor2, simtParamsPtr_, outputLength_);316 indexTensor0, indexTensor1, indexTensor2, simtParamsPtr_, outputLength_);
316 }317 }
317 } else if (inputDimNum_ == INPUTDIMNUMFOUR) {318 } else if (inputDimNum_ == INPUTDIMNUMFOUR) {
318 if (indexedDimNum_ == INDEX_DIM_NUM_ONE) {319 if (indexedDimNum_ == INDEX_DIM_NUM_ONE) {
319 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);320 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);
320 __gm__ P* indexTensor3 = GetInputTensorAddr(indices_, LIST_INDEX_THREE);321 __gm__ P* indexTensor3 = GetInputTensorAddr(indices_, LIST_INDEX_THREE);
321- AscendC::Simt::VF_CALL<SimtDim4Compute<T, P, T2, INDEX_DIM_NUM_ONE>>(322+ asc_vf_call<SimtDim4Compute<T, P, T2, INDEX_DIM_NUM_ONE>>(
322- AscendC::Simt::Dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),323+ dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
323 indexTensor0, indexTensor1, indexTensor2, indexTensor3, simtParamsPtr_, outputLength_);324 indexTensor0, indexTensor1, indexTensor2, indexTensor3, simtParamsPtr_, outputLength_);
324 } else if (indexedDimNum_ == INDEX_DIM_NUM_TWO) {325 } else if (indexedDimNum_ == INDEX_DIM_NUM_TWO) {
325 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);326 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);
326 __gm__ P* indexTensor3 = GetInputTensorAddr(indices_, LIST_INDEX_THREE);327 __gm__ P* indexTensor3 = GetInputTensorAddr(indices_, LIST_INDEX_THREE);
327- AscendC::Simt::VF_CALL<SimtDim4Compute<T, P, T2, INDEX_DIM_NUM_TWO>>(328+ asc_vf_call<SimtDim4Compute<T, P, T2, INDEX_DIM_NUM_TWO>>(
328- AscendC::Simt::Dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),329+ dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
329 indexTensor0, indexTensor1, indexTensor2, indexTensor3, simtParamsPtr_, outputLength_);330 indexTensor0, indexTensor1, indexTensor2, indexTensor3, simtParamsPtr_, outputLength_);
330 } else if (indexedDimNum_ == INDEX_DIM_NUM_THREE) {331 } else if (indexedDimNum_ == INDEX_DIM_NUM_THREE) {
331 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);332 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);
332 __gm__ P* indexTensor3 = GetInputTensorAddr(indices_, LIST_INDEX_THREE);333 __gm__ P* indexTensor3 = GetInputTensorAddr(indices_, LIST_INDEX_THREE);
333- AscendC::Simt::VF_CALL<SimtDim4Compute<T, P, T2, INDEX_DIM_NUM_THREE>>(334+ asc_vf_call<SimtDim4Compute<T, P, T2, INDEX_DIM_NUM_THREE>>(
334- AscendC::Simt::Dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),335+ dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
335 indexTensor0, indexTensor1, indexTensor2, indexTensor3, simtParamsPtr_, outputLength_);336 indexTensor0, indexTensor1, indexTensor2, indexTensor3, simtParamsPtr_, outputLength_);
336 } else if (indexedDimNum_ == INDEX_DIM_NUM_FOUR) {337 } else if (indexedDimNum_ == INDEX_DIM_NUM_FOUR) {
337 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);338 __gm__ P* indexTensor2 = GetInputTensorAddr(indices_, LIST_INDEX_TWO);
338 __gm__ P* indexTensor3 = GetInputTensorAddr(indices_, LIST_INDEX_THREE);339 __gm__ P* indexTensor3 = GetInputTensorAddr(indices_, LIST_INDEX_THREE);
339- AscendC::Simt::VF_CALL<SimtDim4Compute<T, P, T2, INDEX_DIM_NUM_FOUR>>(340+ asc_vf_call<SimtDim4Compute<T, P, T2, INDEX_DIM_NUM_FOUR>>(
340- AscendC::Simt::Dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),341+ dim3{USED_THREAD}, (__gm__ T*)outputGm_.GetPhyAddr(), (__gm__ T*)inputXGm_.GetPhyAddr(),
341 indexTensor0, indexTensor1, indexTensor2, indexTensor3, simtParamsPtr_, outputLength_);342 indexTensor0, indexTensor1, indexTensor2, indexTensor3, simtParamsPtr_, outputLength_);
342 }343 }
343 }344 }
Mindex/index_put_with_sort_v2/op_kernel/arch35/index_put_with_sort_v2.h+11-10
@@ -20,6 +20,7 @@
20#include "kernel_tiling/kernel_tiling.h"20#include "kernel_tiling/kernel_tiling.h"
21#include "kernel_operator.h"21#include "kernel_operator.h"
22#include "op_kernel/platform_util.h"22#include "op_kernel/platform_util.h"
23+#include "simt_api/asc_simt.h"
23#include "index_put_with_sort_v2_struct.h"24#include "index_put_with_sort_v2_struct.h"
24 25 
25namespace AscendC26namespace AscendC
@@ -67,11 +68,11 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(LAUNCH_BOUND_LIMIT) inline void SimtIndexPut
67 // 动态计算外层循环起始位置和步长68 // 动态计算外层循环起始位置和步长
68 uint64_t outerStart, outerStep;69 uint64_t outerStart, outerStep;
69 if constexpr (INDEXED_BLOCK_MODE) {70 if constexpr (INDEXED_BLOCK_MODE) {
70- outerStart = Simt::GetBlockIdx() * Simt::GetThreadNum<0>() + Simt::GetThreadIdx<0>();71+ outerStart = blockIdx.x * blockDim.x + threadIdx.x;
71- outerStep = Simt::GetBlockNum() * Simt::GetThreadNum<0>();72+ outerStep = gridDim.x * blockDim.x;
72 } else {73 } else {
73- outerStart = Simt::GetThreadIdx<0>();74+ outerStart = threadIdx.x;
74- outerStep = Simt::GetThreadNum<0>();75+ outerStep = blockDim.x;
75 }76 }
76 #pragma unroll77 #pragma unroll
77 for (uint64_t idxedIdx = outerStart; idxedIdx < indexedDimSize; idxedIdx += outerStep) {78 for (uint64_t idxedIdx = outerStart; idxedIdx < indexedDimSize; idxedIdx += outerStep) {
@@ -96,11 +97,11 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(LAUNCH_BOUND_LIMIT) inline void SimtIndexPut
96 uint64_t idxedValueIdx = posIdx[curIdxedIdx] * idxedValueStride;97 uint64_t idxedValueIdx = posIdx[curIdxedIdx] * idxedValueStride;
97 uint64_t innerStart, innerStep;98 uint64_t innerStart, innerStep;
98 if constexpr (INDEXED_BLOCK_MODE) {99 if constexpr (INDEXED_BLOCK_MODE) {
99- innerStart = Simt::GetThreadIdx<1>();100+ innerStart = threadIdx.y;
100- innerStep = Simt::GetThreadNum<1>();101+ innerStep = blockDim.y;
101 } else {102 } else {
102- innerStart = Simt::GetBlockIdx() * Simt::GetThreadNum<1>() + Simt::GetThreadIdx<1>();103+ innerStart = blockIdx.x * blockDim.y + threadIdx.y;
103- innerStep = Simt::GetBlockNum() * Simt::GetThreadNum<1>();104+ innerStep = gridDim.x * blockDim.y;
104 }105 }
105 for (uint64_t k = innerStart; k < nonIndexedDimSize; k += innerStep) {106 for (uint64_t k = innerStart; k < nonIndexedDimSize; k += innerStep) {
106 uint64_t nonIdxedSelfIdx = 0;107 uint64_t nonIdxedSelfIdx = 0;
@@ -167,8 +168,8 @@ public:
167 calcParamsPtr->m_[i] = m, calcParamsPtr->shift_[i] = shift;168 calcParamsPtr->m_[i] = m, calcParamsPtr->shift_[i] = shift;
168 }169 }
169 constexpr uint16_t THREAD_NUM_LIMIT = ALL_INDEXED ? THREAD_NUM_FULL : THREAD_NUM_HALF;170 constexpr uint16_t THREAD_NUM_LIMIT = ALL_INDEXED ? THREAD_NUM_FULL : THREAD_NUM_HALF;
170- Simt::VF_CALL<SimtIndexPutV2<TX, TIDX, ACCUMULATE, ALL_INDEXED, INDEXED_BLOCK_MODE, THREAD_NUM_LIMIT>>(171+ asc_vf_call<SimtIndexPutV2<TX, TIDX, ACCUMULATE, ALL_INDEXED, INDEXED_BLOCK_MODE, THREAD_NUM_LIMIT>>(
171- Simt::Dim3{tiling_->indexedThreadNum, tiling_->nonIndexedThreadNum, 1}, output, sortIndices, posIdx,172+ dim3{tiling_->indexedThreadNum, tiling_->nonIndexedThreadNum, 1}, output, sortIndices, posIdx,
172 values, tilingUbAddr, calcParamsPtr);173 values, tilingUbAddr, calcParamsPtr);
173 }174 }
174 175 
Mindex/unsorted_segment_sum/op_kernel/arch35/unsorted_segment_add.h+2-2
@@ -134,8 +134,8 @@ public:
134 __local_mem__ TX* xUbLocalPtr = (__local_mem__ TX*)xUbLocal.GetPhyAddr();134 __local_mem__ TX* xUbLocalPtr = (__local_mem__ TX*)xUbLocal.GetPhyAddr();
135 __local_mem__ TX* midResPtr = (__local_mem__ TX*)midRes.GetPhyAddr();135 __local_mem__ TX* midResPtr = (__local_mem__ TX*)midRes.GetPhyAddr();
136 DataSyncBarrier<MemDsbT::UB>();136 DataSyncBarrier<MemDsbT::UB>();
137- AscendC::Simt::VF_CALL<SimtGatherValue<TX, Index>>(137+ asc_vf_call<SimtGatherValue<TX, Index>>(
138- AscendC::Simt::Dim3{static_cast<uint32_t>(innerDimSize), static_cast<uint32_t>(ROW_NUM)},138+ dim3{static_cast<uint32_t>(innerDimSize), static_cast<uint32_t>(ROW_NUM)},
139 midResPtr, xUbLocalPtr, (__ubuf__ Index*)indexUb.GetPhyAddr(), outputOuterDimSize, innerDimSize,139 midResPtr, xUbLocalPtr, (__ubuf__ Index*)indexUb.GetPhyAddr(), outputOuterDimSize, innerDimSize,
140 needIndexOneUb, oneRowOutNumAlign);140 needIndexOneUb, oneRowOutNumAlign);
141 inQueueX.FreeTensor(xUbLocal);141 inQueueX.FreeTensor(xUbLocal);
Mindex/unsorted_segment_sum/op_kernel/arch35/unsorted_segment_base.h+20-14
@@ -19,6 +19,12 @@
19#include "kernel_tiling/kernel_tiling.h"19#include "kernel_tiling/kernel_tiling.h"
20#include "op_kernel/platform_util.h"20#include "op_kernel/platform_util.h"
21#include "op_kernel/math_util.h"21#include "op_kernel/math_util.h"
22+#include "simt_api/asc_simt.h"
23+#include "simt_api/device_atomic_functions.h"
24+#include "simt_api/asc_fp16.h"
25+#include "simt_api/asc_bf16.h"
26+#include "simt_api/math_functions.h"
27+#include "simt_api/device_functions.h"
22 28 
23namespace UnsortedSegmentSum {29namespace UnsortedSegmentSum {
24using namespace AscendC;30using namespace AscendC;
@@ -77,8 +83,8 @@ template <typename TX>
77__simt_vf__ __aicore__ LAUNCH_BOUND(SIMT_THREAD_DIM_LAUNCH_BOUND) inline void ComputeSetValue(83__simt_vf__ __aicore__ LAUNCH_BOUND(SIMT_THREAD_DIM_LAUNCH_BOUND) inline void ComputeSetValue(
78 __gm__ TX* outputGm, const uint32_t blockNums, const uint32_t outputLength)84 __gm__ TX* outputGm, const uint32_t blockNums, const uint32_t outputLength)
79{85{
80- for (uint32_t outputIndex = block_idx * Simt::GetThreadNum() + Simt::GetThreadIdx(); outputIndex < outputLength;86+ for (uint32_t outputIndex = block_idx * blockDim.x + threadIdx.x; outputIndex < outputLength;
81- outputIndex = outputIndex + blockNums * Simt::GetThreadNum()) {87+ outputIndex = outputIndex + blockNums * blockDim.x) {
82 outputGm[outputIndex] = static_cast<TX>(0);88 outputGm[outputIndex] = static_cast<TX>(0);
83 }89 }
84}90}
@@ -88,15 +94,15 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(SORT_THREAD_DIM_LAUNCH_BOUND) inline void Si
88 __ubuf__ TX* midResPtr, __ubuf__ TX* xUbLocalPtr, __ubuf__ Index* indexUb, const uint32_t outputOuterDimSize,94 __ubuf__ TX* midResPtr, __ubuf__ TX* xUbLocalPtr, __ubuf__ Index* indexUb, const uint32_t outputOuterDimSize,
89 const uint32_t innerDimSize, const uint32_t needIndexOneUb, const uint32_t outputOffset)95 const uint32_t innerDimSize, const uint32_t needIndexOneUb, const uint32_t outputOffset)
90{96{
91- Index midBaseOffset = Simt::GetThreadIdx<1>() * outputOffset;97+ Index midBaseOffset = threadIdx.y * outputOffset;
92- Index offset = Simt::GetThreadIdx<1>();98+ Index offset = threadIdx.y;
93 for (; offset < needIndexOneUb; offset += ROW_NUM) {99 for (; offset < needIndexOneUb; offset += ROW_NUM) {
94 Index indexVal = indexUb[offset];100 Index indexVal = indexUb[offset];
95 if (indexVal >= 0 && indexVal < outputOuterDimSize) {101 if (indexVal >= 0 && indexVal < outputOuterDimSize) {
96 Index midResOffSet = indexVal * innerDimSize;102 Index midResOffSet = indexVal * innerDimSize;
97 Index xUbLocalOffSet = offset * innerDimSize;103 Index xUbLocalOffSet = offset * innerDimSize;
98- midResPtr[midBaseOffset + midResOffSet + Simt::GetThreadIdx<0>()] +=104+ midResPtr[midBaseOffset + midResOffSet + threadIdx.x] +=
99- xUbLocalPtr[xUbLocalOffSet + Simt::GetThreadIdx<0>()];105+ xUbLocalPtr[xUbLocalOffSet + threadIdx.x];
100 }106 }
101 }107 }
102}108}
@@ -106,8 +112,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(SIMT_THREAD_DIM_LAUNCH_BOUND) inline void Si
106 __gm__ TX* xGm, __gm__ Index* segmentIdsGm, __gm__ TX* outputGm, const uint32_t blockNums, const COM_T inputLength,112 __gm__ TX* xGm, __gm__ Index* segmentIdsGm, __gm__ TX* outputGm, const uint32_t blockNums, const COM_T inputLength,
107 const COM_T innerDimSize, const uint64_t outputOuterDimSize, const COM_T magic, const COM_T shift)113 const COM_T innerDimSize, const uint64_t outputOuterDimSize, const COM_T magic, const COM_T shift)
108{114{
109- for (COM_T inputIndex = block_idx * Simt::GetThreadNum() + Simt::GetThreadIdx(); inputIndex < inputLength;115+ for (COM_T inputIndex = block_idx * blockDim.x + threadIdx.x; inputIndex < inputLength;
110- inputIndex = inputIndex + blockNums * Simt::GetThreadNum()) {116+ inputIndex = inputIndex + blockNums * blockDim.x) {
111 COM_T inputSegmentIndex = Simt::UintDiv(inputIndex, magic, shift);117 COM_T inputSegmentIndex = Simt::UintDiv(inputIndex, magic, shift);
112 COM_T segmentOffset = inputIndex - inputSegmentIndex * innerDimSize;118 COM_T segmentOffset = inputIndex - inputSegmentIndex * innerDimSize;
113 const Index outputSegmentIndex = segmentIdsGm[inputSegmentIndex];119 const Index outputSegmentIndex = segmentIdsGm[inputSegmentIndex];
@@ -115,7 +121,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(SIMT_THREAD_DIM_LAUNCH_BOUND) inline void Si
115 continue;121 continue;
116 }122 }
117 const uint64_t outputIndex = outputSegmentIndex * innerDimSize + segmentOffset;123 const uint64_t outputIndex = outputSegmentIndex * innerDimSize + segmentOffset;
118- Simt::AtomicAdd(outputGm + outputIndex, xGm[inputIndex]);124+ asc_atomic_add(outputGm + outputIndex, xGm[inputIndex]);
119 }125 }
120}126}
121 127 
@@ -125,10 +131,10 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(SORT_THREAD_DIM_LAUNCH_BOUND) inline void Se
125 __ubuf__ uint32_t* cumSumAddr, __gm__ TX* outputAddr, int32_t uniqueIndexNum, uint32_t lastDim,131 __ubuf__ uint32_t* cumSumAddr, __gm__ TX* outputAddr, int32_t uniqueIndexNum, uint32_t lastDim,
126 uint32_t outputOuterDimSize)132 uint32_t outputOuterDimSize)
127{133{
128- int32_t blockIdx = Simt::GetThreadIdx<1>();134+ int32_t blockIdxVal = threadIdx.y;
129- int32_t blockNum = Simt::GetThreadNum<1>();135+ int32_t blockNum = blockDim.y;
130- int32_t innerOffset = Simt::GetThreadIdx<0>();136+ int32_t innerOffset = threadIdx.x;
131- for (int32_t i = blockIdx; i < uniqueIndexNum; i += blockNum) {137+ for (int32_t i = blockIdxVal; i < uniqueIndexNum; i += blockNum) {
132 if (sortedAddr[cumSumAddr[i]] < 0 || sortedAddr[cumSumAddr[i]] >= outputOuterDimSize) {138 if (sortedAddr[cumSumAddr[i]] < 0 || sortedAddr[cumSumAddr[i]] >= outputOuterDimSize) {
133 continue;139 continue;
134 }140 }
@@ -138,7 +144,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(SORT_THREAD_DIM_LAUNCH_BOUND) inline void Se
138 result += inputAddr[srcOffset];144 result += inputAddr[srcOffset];
139 }145 }
140 int64_t gmDstOffset = sortedAddr[cumSumAddr[i]] * lastDim + innerOffset;146 int64_t gmDstOffset = sortedAddr[cumSumAddr[i]] * lastDim + innerOffset;
141- Simt::AtomicAdd(outputAddr + gmDstOffset, result);147+ asc_atomic_add(outputAddr + gmDstOffset, result);
142 }148 }
143 return;149 return;
144}150}
Mindex/unsorted_segment_sum/op_kernel/arch35/unsorted_segment_sort_simt.h+2-2
@@ -157,8 +157,8 @@ __aicore__ inline void KernelUnsortedSegmentSortSimt<TX, Index, CAST_T, castType
157 __ubuf__ uint32_t* cumSumAddr = (__ubuf__ uint32_t*)cumSumLocal.GetPhyAddr();157 __ubuf__ uint32_t* cumSumAddr = (__ubuf__ uint32_t*)cumSumLocal.GetPhyAddr();
158 __gm__ TX* outputGm = (__gm__ TX*)outputGm_.GetPhyAddr();158 __gm__ TX* outputGm = (__gm__ TX*)outputGm_.GetPhyAddr();
159 159 
160- AscendC::Simt::VF_CALL<SegmentReduceSortSimt<TX, CAST_T>>(160+ asc_vf_call<SegmentReduceSortSimt<TX, CAST_T>>(
161- Simt::Dim3({static_cast<uint32_t>(tilingData_->innerDim), static_cast<uint32_t>(threadBlock)}), inputAddr,161+ dim3({static_cast<uint32_t>(tilingData_->innerDim), static_cast<uint32_t>(threadBlock)}), inputAddr,
162 sortedIndexAddr, sortedSengmentAddr, cumSumAddr, outputGm, uniqueIndexNum, tilingData_->innerDim,162 sortedIndexAddr, sortedSengmentAddr, cumSumAddr, outputGm, uniqueIndexNum, tilingData_->innerDim,
163 tilingData_->outputOuterDim);163 tilingData_->outputOuterDim);
164 164 
Mindex/unsorted_segment_sum/op_kernel/arch35/unsorted_segment_sum.h+2-2
@@ -58,8 +58,8 @@ public:
58 COM_T magic = 1;58 COM_T magic = 1;
59 COM_T shift = 1;59 COM_T shift = 1;
60 GetUintDivMagicAndShift(magic, shift, static_cast<COM_T>(innerDimSizeTmp));60 GetUintDivMagicAndShift(magic, shift, static_cast<COM_T>(innerDimSizeTmp));
61- AscendC::Simt::VF_CALL<SimtComputeSegment<T, Index, COM_T>>(61+ asc_vf_call<SimtComputeSegment<T, Index, COM_T>>(
62- Simt::Dim3(static_cast<uint32_t>(td_->maxThread)), input, segmentIds, output, blockNums, inputLength,62+ dim3(static_cast<uint32_t>(td_->maxThread)), input, segmentIds, output, blockNums, inputLength,
63 innerDimSizeTmp, outputOuterDimSizeTmp, magic, shift);63 innerDimSizeTmp, outputOuterDimSizeTmp, magic, shift);
64 }64 }
65 65 
Mindex/unsorted_segment_sum/op_kernel/arch35/uss_deterministic.h+19-19
@@ -91,7 +91,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(USED_THREAD_NUMS) inline void SimtAtomicComp
91 __gm__ uint64_t* workspaceNValue_, __local_mem__ T* xTensor, __local_mem__ U* sortedIdTensor,91 __gm__ uint64_t* workspaceNValue_, __local_mem__ T* xTensor, __local_mem__ U* sortedIdTensor,
92 __local_mem__ uint32_t* sortedIdIndexTensor)92 __local_mem__ uint32_t* sortedIdIndexTensor)
93{93{
94- for (int32_t j = Simt::GetThreadIdx(); j < innerDim_; j += Simt::GetThreadNum()) {94+ for (int32_t j = threadIdx.x; j < innerDim_; j += blockDim.x) {
95 // 找第一个有效的 startSegId95 // 找第一个有效的 startSegId
96 int64_t startSegId = -1;96 int64_t startSegId = -1;
97 for (uint32_t k = 0; k < curProcessRowsNum; k++) {97 for (uint32_t k = 0; k < curProcessRowsNum; k++) {
@@ -116,25 +116,25 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(USED_THREAD_NUMS) inline void SimtAtomicComp
116 if (curSegId < 0 || curSegId >= static_cast<int64_t>(segmentNum)) {116 if (curSegId < 0 || curSegId >= static_cast<int64_t>(segmentNum)) {
117 if (i == curProcessRowsNum - 1 && num > 0) {117 if (i == curProcessRowsNum - 1 && num > 0) {
118 if (j == 0) {118 if (j == 0) {
119- Simt::AtomicAdd(workspaceNValue_ + static_cast<uint64_t>(startSegId), num);119+ asc_atomic_add(workspaceNValue_ + static_cast<uint64_t>(startSegId), num);
120 }120 }
121- Simt::AtomicMax(workspaceMValue_ + static_cast<uint64_t>(startSegId) * innerDim_ + j, maxValue);121+ asc_atomic_max(workspaceMValue_ + static_cast<uint64_t>(startSegId) * innerDim_ + j, maxValue);
122 }122 }
123 continue;123 continue;
124 }124 }
125 125 
126 int32_t tmpXIndex = sortedIdIndexTensor[i] * innerDim_ + j;126 int32_t tmpXIndex = sortedIdIndexTensor[i] * innerDim_ + j;
127- float tmpValue = Simt::Abs(static_cast<float>(xTensor[tmpXIndex]));127+ float tmpValue = fabsf(static_cast<float>(xTensor[tmpXIndex]));
128 128 
129 if (curSegId == startSegId) {129 if (curSegId == startSegId) {
130 num++;130 num++;
131 maxValue = maxValue > tmpValue ? maxValue : tmpValue; 131 maxValue = maxValue > tmpValue ? maxValue : tmpValue;
132 } else {132 } else {
133 if (j == 0) {133 if (j == 0) {
134- Simt::AtomicAdd(workspaceNValue_ + static_cast<uint64_t>(startSegId), num);134+ asc_atomic_add(workspaceNValue_ + static_cast<uint64_t>(startSegId), num);
135 }135 }
136 136 
137- Simt::AtomicMax(workspaceMValue_ + static_cast<uint64_t>(startSegId) * innerDim_ + j, maxValue);137+ asc_atomic_max(workspaceMValue_ + static_cast<uint64_t>(startSegId) * innerDim_ + j, maxValue);
138 startSegId = curSegId;138 startSegId = curSegId;
139 num = 1;139 num = 1;
140 maxValue = tmpValue;140 maxValue = tmpValue;
@@ -142,9 +142,9 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(USED_THREAD_NUMS) inline void SimtAtomicComp
142 142 
143 if (i == curProcessRowsNum - 1) {143 if (i == curProcessRowsNum - 1) {
144 if (j == 0) {144 if (j == 0) {
145- Simt::AtomicAdd(workspaceNValue_ + static_cast<uint64_t>(startSegId), num);145+ asc_atomic_add(workspaceNValue_ + static_cast<uint64_t>(startSegId), num);
146 }146 }
147- Simt::AtomicMax(workspaceMValue_ + static_cast<uint64_t>(startSegId) * innerDim_ + j, maxValue);147+ asc_atomic_max(workspaceMValue_ + static_cast<uint64_t>(startSegId) * innerDim_ + j, maxValue);
148 break;148 break;
149 }149 }
150 }150 }
@@ -250,7 +250,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(USED_THREAD_NUMS) inline void QuantizeAndCop
250{250{
251 float scaling = static_cast<float>(1L << 30);251 float scaling = static_cast<float>(1L << 30);
252 252 
253- for (uint32_t i = Simt::GetThreadIdx(); i < innerDim; i += Simt::GetThreadNum()) {253+ for (uint32_t i = threadIdx.x; i < innerDim; i += blockDim.x) {
254 // 找第一个有效的 startSegId254 // 找第一个有效的 startSegId
255 int64_t startSegId = -1;255 int64_t startSegId = -1;
256 for (uint32_t k = 0; k < curProcessRowsNum; k++) {256 for (uint32_t k = 0; k < curProcessRowsNum; k++) {
@@ -280,7 +280,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(USED_THREAD_NUMS) inline void QuantizeAndCop
280 int32_t quantizedResInt =280 int32_t quantizedResInt =
281 AscendC::Simt::Cast<int32_t, float, RoundMode::CAST_RINT, AscendC::Simt::SatMode::SAT>(281 AscendC::Simt::Cast<int32_t, float, RoundMode::CAST_RINT, AscendC::Simt::SatMode::SAT>(
282 sumValue);282 sumValue);
283- Simt::AtomicAdd(workspaceOutput32 + tmpOffset, quantizedResInt);283+ asc_atomic_add(workspaceOutput32 + tmpOffset, quantizedResInt);
284 }284 }
285 continue;285 continue;
286 }286 }
@@ -297,7 +297,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(USED_THREAD_NUMS) inline void QuantizeAndCop
297 297 
298 sumValue = sumValue / (mValueFloat * nValueFloat) * scaling;298 sumValue = sumValue / (mValueFloat * nValueFloat) * scaling;
299 int32_t quantizedResInt = AscendC::Simt::Cast<int32_t, float, RoundMode::CAST_RINT, AscendC::Simt::SatMode::SAT>(sumValue);299 int32_t quantizedResInt = AscendC::Simt::Cast<int32_t, float, RoundMode::CAST_RINT, AscendC::Simt::SatMode::SAT>(sumValue);
300- Simt::AtomicAdd(workspaceOutput32 + tmpOffset, quantizedResInt);300+ asc_atomic_add(workspaceOutput32 + tmpOffset, quantizedResInt);
301 301 
302 startSegId = curSegId;302 startSegId = curSegId;
303 sumValue = tmpValue;303 sumValue = tmpValue;
@@ -309,7 +309,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(USED_THREAD_NUMS) inline void QuantizeAndCop
309 float nValueFloat = static_cast<float>(nValueWorkSpace[static_cast<uint64_t>(startSegId)]);309 float nValueFloat = static_cast<float>(nValueWorkSpace[static_cast<uint64_t>(startSegId)]);
310 sumValue = sumValue / (mValueFloat * nValueFloat) * scaling;310 sumValue = sumValue / (mValueFloat * nValueFloat) * scaling;
311 int32_t quantizedResInt = AscendC::Simt::Cast<int32_t, float, RoundMode::CAST_RINT, AscendC::Simt::SatMode::SAT>(sumValue);311 int32_t quantizedResInt = AscendC::Simt::Cast<int32_t, float, RoundMode::CAST_RINT, AscendC::Simt::SatMode::SAT>(sumValue);
312- Simt::AtomicAdd(workspaceOutput32 + tmpOffset, quantizedResInt);312+ asc_atomic_add(workspaceOutput32 + tmpOffset, quantizedResInt);
313 break;313 break;
314 }314 }
315 }315 }
@@ -329,8 +329,8 @@ __aicore__ inline void KernelUSSDeterministic<T, U>::FirstUbProcess(uint64_t cur
329 AscendC::Sort<U, false, sortConfig>(sortedIdTensor, sortedIdIndexTensor, segmentIdTensor, sharedTmpTensor,329 AscendC::Sort<U, false, sortConfig>(sortedIdTensor, sortedIdIndexTensor, segmentIdTensor, sharedTmpTensor,
330 static_cast<uint32_t>(curProcessRowsNum));330 static_cast<uint32_t>(curProcessRowsNum));
331 331 
332- AscendC::Simt::VF_CALL<SimtAtomicComputeNValueAndMValue<T, U>>(332+ asc_vf_call<SimtAtomicComputeNValueAndMValue<T, U>>(
333- Simt::Dim3(MAX_THREAD), curProcessRowsNum, innerDim_, segmentNum_,333+ dim3(MAX_THREAD), curProcessRowsNum, innerDim_, segmentNum_,
334 (__gm__ float*)(workspaceMValue_.GetPhyAddr()), (__gm__ uint64_t*)(workspaceNValue_.GetPhyAddr()),334 (__gm__ float*)(workspaceMValue_.GetPhyAddr()), (__gm__ uint64_t*)(workspaceNValue_.GetPhyAddr()),
335 (__local_mem__ T*)(xLocal.GetPhyAddr()), (__local_mem__ U*)(sortedIdTensor.GetPhyAddr()),335 (__local_mem__ T*)(xLocal.GetPhyAddr()), (__local_mem__ U*)(sortedIdTensor.GetPhyAddr()),
336 (__local_mem__ uint32_t*)(sortedIdIndexTensor.GetPhyAddr()));336 (__local_mem__ uint32_t*)(sortedIdIndexTensor.GetPhyAddr()));
@@ -352,8 +352,8 @@ __aicore__ inline void KernelUSSDeterministic<T, U>::SecondUbProcess(uint64_t cu
352 AscendC::Sort<U, false, sortConfig>(sortedIdTensor, sortedIdIndexTensor, segmentIdTensor, sharedTmpTensor,352 AscendC::Sort<U, false, sortConfig>(sortedIdTensor, sortedIdIndexTensor, segmentIdTensor, sharedTmpTensor,
353 static_cast<uint32_t>(curProcessRowsNum));353 static_cast<uint32_t>(curProcessRowsNum));
354 354 
355- AscendC::Simt::VF_CALL<QuantizeAndCopyOut<T, U>>(355+ asc_vf_call<QuantizeAndCopyOut<T, U>>(
356- Simt::Dim3(MAX_THREAD), (__gm__ float*)(workspaceMValue_.GetPhyAddr()),356+ dim3(MAX_THREAD), (__gm__ float*)(workspaceMValue_.GetPhyAddr()),
357 (__gm__ uint64_t*)(workspaceNValue_.GetPhyAddr()), (__gm__ int32_t*)(workspaceOutput_.GetPhyAddr()),357 (__gm__ uint64_t*)(workspaceNValue_.GetPhyAddr()), (__gm__ int32_t*)(workspaceOutput_.GetPhyAddr()),
358 (__local_mem__ T*)(xLocal.GetPhyAddr()), (__local_mem__ U*)(sortedIdTensor.GetPhyAddr()),358 (__local_mem__ T*)(xLocal.GetPhyAddr()), (__local_mem__ U*)(sortedIdTensor.GetPhyAddr()),
359 (__local_mem__ uint32_t*)(sortedIdIndexTensor.GetPhyAddr()), curProcessRowsNum, innerDim_, segmentNum_);359 (__local_mem__ uint32_t*)(sortedIdIndexTensor.GetPhyAddr()), curProcessRowsNum, innerDim_, segmentNum_);
@@ -369,8 +369,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(USED_THREAD_NUMS) inline void Dequantize(
369{369{
370 float scaling = static_cast<float>(1L << 30);370 float scaling = static_cast<float>(1L << 30);
371 371
372- for (uint32_t i = blockId * Simt::GetThreadNum() + Simt::GetThreadIdx(); i < segmentNum * innerDim; 372+ for (uint32_t i = blockId * blockDim.x + threadIdx.x; i < segmentNum * innerDim;
373- i += blockNum * Simt::GetThreadNum()) { 373+ i += blockNum * blockDim.x) {
374 uint32_t row = i / innerDim;374 uint32_t row = i / innerDim;
375 uint32_t col = i % innerDim;375 uint32_t col = i % innerDim;
376 376 
@@ -387,7 +387,7 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(USED_THREAD_NUMS) inline void Dequantize(
387template <typename T, typename U>387template <typename T, typename U>
388__aicore__ inline void KernelUSSDeterministic<T, U>::ThirdUbProcess()388__aicore__ inline void KernelUSSDeterministic<T, U>::ThirdUbProcess()
389{389{
390- Simt::VF_CALL<Dequantize<T,U>>(Simt::Dim3(MAX_THREAD), (__gm__ float*)(workspaceMValue_.GetPhyAddr()), 390+ asc_vf_call<Dequantize<T,U>>(dim3(MAX_THREAD), (__gm__ float*)(workspaceMValue_.GetPhyAddr()),
391 (__gm__ uint64_t*)(workspaceNValue_.GetPhyAddr()), (__gm__ int32_t*)(workspaceOutput_.GetPhyAddr()),391 (__gm__ uint64_t*)(workspaceNValue_.GetPhyAddr()), (__gm__ int32_t*)(workspaceOutput_.GetPhyAddr()),
392 (__gm__ T*)(yGm_.GetPhyAddr()), segmentNum_, innerDim_, blockId_, usedCoreNum_);392 (__gm__ T*)(yGm_.GetPhyAddr()), segmentNum_, innerDim_, blockId_, usedCoreNum_);
393}393}