Star869
1749
代码
代码
Issues214
Pull Requests313
Actions
项目讨论
Wiki
分析
Star869
1749
已合并
gather_v2、scatter_elements、max_pool_v3编译时间优化 #1490
xiaodong666创建于 2月4日
gather_v2、scatter_elements、max_pool_v3编译时间优化 #1490
已合并
xiaodong666创建于 2月4日
从xiaodong666/ops-nnbianyi_youhua合入到cann/ops-nnmaster
共 14 个文件变更+901-632
Mindex/gather_v2/op_kernel/arch35/gather_v2.h+1-0
@@ -20,6 +20,7 @@
20#endif20#endif
21 21 
22#include "kernel_operator.h"22#include "kernel_operator.h"
23+#include "basic_api/kernel_vec_intf.h"
23 24 
24#ifdef __DAV_FPGA__25#ifdef __DAV_FPGA__
25constexpr uint32_t THREAD_NUM_LAUNCH_BOUND = 512;26constexpr uint32_t THREAD_NUM_LAUNCH_BOUND = 512;
Mindex/gather_v2/op_kernel/arch35/gather_v2_after_gather_full_load.h+1-0
@@ -16,6 +16,7 @@
16#define GATHER_V2_AFTER_GATHER_FULL_LOAD16#define GATHER_V2_AFTER_GATHER_FULL_LOAD
17 17 
18#include "kernel_operator.h"18#include "kernel_operator.h"
19+#include "basic_api/kernel_vec_intf.h"
19#include "op_kernel/platform_util.h"20#include "op_kernel/platform_util.h"
20 21 
21#ifdef __DAV_FPGA__22#ifdef __DAV_FPGA__
Mindex/gather_v2/op_kernel/arch35/gather_v2_empty.h+1-1
@@ -15,7 +15,7 @@
15#ifndef GATHER_V2_EMPTY15#ifndef GATHER_V2_EMPTY
16#define GATHER_V2_EMPTY16#define GATHER_V2_EMPTY
17 17 
18-#include "kernel_operator.h"18+#include "basic_api/kernel_vec_intf.h"
19#include "op_kernel/platform_util.h"19#include "op_kernel/platform_util.h"
20 20 
21namespace gatherv2 {21namespace gatherv2 {
Mindex/gather_v2/op_kernel/arch35/gather_v2_ga_all_load.h+2-1
@@ -19,7 +19,8 @@
19#define K_MAX_SHAPE_DIM 019#define K_MAX_SHAPE_DIM 0
20#endif20#endif
21 21 
22-#include "kernel_operator.h"22+#include "basic_api/kernel_vec_intf.h"
23+#include "micro_api/kernel_micro_intf.h"
23#include "op_kernel/platform_util.h"24#include "op_kernel/platform_util.h"
24 25 
25namespace gatherv2 {26namespace gatherv2 {
Mindex/gather_v2/op_kernel/arch35/gather_v2_simd.h+1-1
@@ -19,7 +19,7 @@
19#define K_MAX_SHAPE_DIM 019#define K_MAX_SHAPE_DIM 0
20#endif20#endif
21 21 
22-#include "kernel_operator.h"22+#include "basic_api/kernel_vec_intf.h"
23#include "op_kernel/platform_util.h"23#include "op_kernel/platform_util.h"
24 24 
25namespace gatherv2 {25namespace gatherv2 {
Mindex/gather_v2/op_kernel/arch35/gather_v2_simd_last_gather.h+2-1
@@ -19,7 +19,8 @@
19#define K_MAX_SHAPE_DIM 019#define K_MAX_SHAPE_DIM 0
20#endif20#endif
21 21 
22-#include "kernel_operator.h"22+#include "basic_api/kernel_vec_intf.h"
23+#include "micro_api/kernel_micro_intf.h"
23#include "op_kernel/platform_util.h"24#include "op_kernel/platform_util.h"
24 25 
25namespace gatherv2 {26namespace gatherv2 {
Mindex/gather_v2/op_kernel/arch35/gather_v2_simd_two_dim.h+1-1
@@ -19,7 +19,7 @@
19 #define K_MAX_SHAPE_DIM 019 #define K_MAX_SHAPE_DIM 0
20 #endif20 #endif
21 21
22-#include "kernel_operator.h"22+#include "basic_api/kernel_vec_intf.h"
23#include "op_kernel/platform_util.h"23#include "op_kernel/platform_util.h"
24 24
25 namespace gatherv2 {25 namespace gatherv2 {
Mindex/gather_v2/op_kernel/arch35/gather_v2_simt_two_dim.h+1-0
@@ -20,6 +20,7 @@
20#endif20#endif
21 21 
22#include "kernel_operator.h"22#include "kernel_operator.h"
23+#include "basic_api/kernel_vec_intf.h"
23 24 
24#ifdef __DAV_FPGA__25#ifdef __DAV_FPGA__
25constexpr uint32_t THREAD_NUM_LAUNCH_BOUND_TWO_DIM = 512;26constexpr uint32_t THREAD_NUM_LAUNCH_BOUND_TWO_DIM = 512;
Mindex/gather_v2/op_kernel/gather_v2_apt.cpp+200-169
@@ -58,173 +58,204 @@ using namespace gatherv2;
58#define SIMT_TWO_DIM_B64_INDEX_SIZE_64_TILING_KEY 2000000108UL58#define SIMT_TWO_DIM_B64_INDEX_SIZE_64_TILING_KEY 2000000108UL
59 59 
60extern "C" __global__ __aicore__ void gather_v2(GM_ADDR x, GM_ADDR indices, GM_ADDR axis, GM_ADDR y,60extern "C" __global__ __aicore__ void gather_v2(GM_ADDR x, GM_ADDR indices, GM_ADDR axis, GM_ADDR y,
61- GM_ADDR workspace, GM_ADDR tiling) {61+ GM_ADDR workspace, GM_ADDR tiling)
62- TPipe pipe;62+{
63- KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); 63+ TPipe pipe;
64- if (TILING_KEY_IS(TILING_KEY_X_B8_INDEX_SIZE_64)) {64+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
65- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);65+ 
66- Gatherv2<int8_t, DTYPE_INDICES, uint64_t> gatherv2Op;66+ TILING_KEY_IS(SIMD_TILING_KEY);
67- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);67+ TILING_KEY_IS(SIMD_TILING_KEY_TWO_DIM);
68- gatherv2Op.Process();68+ TILING_KEY_IS(SIMD_LAST_GATHER_B8_TILING_KEY);
69- } else if (TILING_KEY_IS(TILING_KEY_X_B16_INDEX_SIZE_64)) {69+ TILING_KEY_IS(SIMD_LAST_GATHER_B16_TILING_KEY);
70- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);70+ TILING_KEY_IS(SIMD_LAST_GATHER_B32_TILING_KEY);
71- Gatherv2<int16_t, DTYPE_INDICES, uint64_t> gatherv2Op;71+ TILING_KEY_IS(SIMD_LAST_GATHER_B64_TILING_KEY);
72- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);72+ TILING_KEY_IS(SIMD_LAST_GATHER_B8_SUPPORT_NEG_INDICE_TILING_KEY);
73- gatherv2Op.Process();73+ TILING_KEY_IS(SIMD_LAST_GATHER_B16_SUPPORT_NEG_INDICE_TILING_KEY);
74- } else if (TILING_KEY_IS(TILING_KEY_X_B32_INDEX_SIZE_64)) {74+ TILING_KEY_IS(SIMD_LAST_GATHER_B32_SUPPORT_NEG_INDICE_TILING_KEY);
75- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);75+ TILING_KEY_IS(SIMD_LAST_GATHER_B64_SUPPORT_NEG_INDICE_TILING_KEY);
76- Gatherv2<int32_t, DTYPE_INDICES, uint64_t> gatherv2Op;76+ TILING_KEY_IS(SIMD_GA_ALL_LOAD_TILING_KEY);
77- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);77+ TILING_KEY_IS(SIMD_GA_ALL_LOAD_SUPPORT_NEG_INDICE_TILING_KEY);
78- gatherv2Op.Process();78+ TILING_KEY_IS(SIMT_TWO_DIM_B8_INDEX_SIZE_32_TILING_KEY);
79- } else if (TILING_KEY_IS(TILING_KEY_X_B64_INDEX_SIZE_64)) {79+ TILING_KEY_IS(SIMT_TWO_DIM_B16_INDEX_SIZE_32_TILING_KEY);
80- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);80+ TILING_KEY_IS(SIMT_TWO_DIM_B32_INDEX_SIZE_32_TILING_KEY);
81- Gatherv2<int64_t, DTYPE_INDICES, uint64_t> gatherv2Op;81+ TILING_KEY_IS(SIMT_TWO_DIM_B64_INDEX_SIZE_32_TILING_KEY);
82- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);82+ TILING_KEY_IS(SIMT_TWO_DIM_B8_INDEX_SIZE_64_TILING_KEY);
83- gatherv2Op.Process();83+ TILING_KEY_IS(SIMT_TWO_DIM_B16_INDEX_SIZE_64_TILING_KEY);
84- } else if (TILING_KEY_IS(TILING_KEY_X_B8_INDEX_SIZE_32)) {84+ TILING_KEY_IS(SIMT_TWO_DIM_B32_INDEX_SIZE_64_TILING_KEY);
85- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);85+ TILING_KEY_IS(SIMT_TWO_DIM_B64_INDEX_SIZE_64_TILING_KEY);
86- Gatherv2<int8_t, DTYPE_INDICES, uint32_t> gatherv2Op;86+ TILING_KEY_IS(TILING_KEY_X_B64_INDEX_SIZE_64);
87- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);87+ TILING_KEY_IS(TILING_KEY_X_B32_INDEX_SIZE_64);
88- gatherv2Op.Process();88+ TILING_KEY_IS(TILING_KEY_X_B16_INDEX_SIZE_64);
89- } else if (TILING_KEY_IS(TILING_KEY_X_B16_INDEX_SIZE_32)) {89+ TILING_KEY_IS(TILING_KEY_X_B8_INDEX_SIZE_64);
90- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);90+ TILING_KEY_IS(TILING_KEY_X_B64_INDEX_SIZE_32);
91- Gatherv2<int16_t, DTYPE_INDICES, uint32_t> gatherv2Op;91+ TILING_KEY_IS(TILING_KEY_X_B32_INDEX_SIZE_32);
92- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);92+ TILING_KEY_IS(TILING_KEY_X_B16_INDEX_SIZE_32);
93- gatherv2Op.Process();93+ TILING_KEY_IS(TILING_KEY_X_B8_INDEX_SIZE_32);
94- } else if (TILING_KEY_IS(TILING_KEY_X_B32_INDEX_SIZE_32)) {94+ TILING_KEY_IS(TILING_KEY_X_B64_INDEX_SIZE_2D);
95- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);95+ TILING_KEY_IS(TILING_KEY_X_B32_INDEX_SIZE_2D);
96- Gatherv2<int32_t, DTYPE_INDICES, uint32_t> gatherv2Op;96+ TILING_KEY_IS(TILING_KEY_X_B16_INDEX_SIZE_2D);
97- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);97+ TILING_KEY_IS(TILING_KEY_X_B8_INDEX_SIZE_2D);
98- gatherv2Op.Process();98+ TILING_KEY_IS(TILING_KEY_EMPTY);
99- } else if (TILING_KEY_IS(TILING_KEY_X_B64_INDEX_SIZE_32)) {99+ 
100- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);100+ #if TILING_KEY_VAR == SIMD_TILING_KEY
101- Gatherv2<int64_t, DTYPE_INDICES, uint32_t> gatherv2Op;101+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);
102- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);102+ Gatherv2Simd<DTYPE_INDICES> gatherv2Op(&pipe);
103- gatherv2Op.Process();103+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
104- } else if (TILING_KEY_IS(SIMD_TILING_KEY)) {104+ gatherv2Op.Process();
105- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);105+ #elif TILING_KEY_VAR == SIMD_TILING_KEY_TWO_DIM
106- Gatherv2Simd<DTYPE_INDICES> gatherv2Op(&pipe);106+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimdTwoDim, tilingDataIn, tiling);
107- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);107+ Gatherv2SimdTwoDim<DTYPE_INDICES> gatherv2Op(&pipe);
108- gatherv2Op.Process();108+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
109- } else if (TILING_KEY_IS(SIMD_TILING_KEY_TWO_DIM)) {109+ gatherv2Op.Process();
110- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimdTwoDim, tilingDataIn, tiling);110+ #elif TILING_KEY_VAR == SIMD_LAST_GATHER_B8_TILING_KEY
111- Gatherv2SimdTwoDim<DTYPE_INDICES> gatherv2Op(&pipe);111+ GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);
112- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);112+ Gatherv2SimdLastGather<int8_t, DTYPE_INDICES, false> gatherv2Op(&pipe);
113- gatherv2Op.Process();113+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
114- } else if (TILING_KEY_IS(SIMD_LAST_GATHER_B8_TILING_KEY)) {114+ gatherv2Op.Process();
115- GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);115+ #elif TILING_KEY_VAR == SIMD_LAST_GATHER_B32_TILING_KEY
116- Gatherv2SimdLastGather<int8_t, DTYPE_INDICES, false> gatherv2Op(&pipe);116+ GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);
117- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);117+ Gatherv2SimdLastGather<int32_t, DTYPE_INDICES, false> gatherv2Op(&pipe);
118- gatherv2Op.Process();118+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
119- } else if (TILING_KEY_IS(SIMD_LAST_GATHER_B16_TILING_KEY)) {119+ gatherv2Op.Process();
120- GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);120+ #elif TILING_KEY_VAR == SIMD_LAST_GATHER_B64_TILING_KEY
121- Gatherv2SimdLastGather<int16_t, DTYPE_INDICES, false> gatherv2Op(&pipe);121+ GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);
122- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);122+ Gatherv2SimdLastGather<int64_t, DTYPE_INDICES, false> gatherv2Op(&pipe);
123- gatherv2Op.Process();123+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
124- } else if (TILING_KEY_IS(SIMD_LAST_GATHER_B32_TILING_KEY)) {124+ gatherv2Op.Process();
125- GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);125+ #elif TILING_KEY_VAR == SIMD_LAST_GATHER_B8_SUPPORT_NEG_INDICE_TILING_KEY
126- Gatherv2SimdLastGather<int32_t, DTYPE_INDICES, false> gatherv2Op(&pipe);126+ GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);
127- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);127+ Gatherv2SimdLastGather<int8_t, DTYPE_INDICES, true> gatherv2Op(&pipe);
128- gatherv2Op.Process();128+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
129- } else if (TILING_KEY_IS(SIMD_LAST_GATHER_B64_TILING_KEY)) {129+ gatherv2Op.Process();
130- GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);130+ #elif TILING_KEY_VAR == SIMD_LAST_GATHER_B16_SUPPORT_NEG_INDICE_TILING_KEY
131- Gatherv2SimdLastGather<int64_t, DTYPE_INDICES, false> gatherv2Op(&pipe);131+ GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);
132- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);132+ Gatherv2SimdLastGather<int16_t, DTYPE_INDICES, true> gatherv2Op(&pipe);
133- gatherv2Op.Process();133+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
134- } else if (TILING_KEY_IS(SIMD_LAST_GATHER_B8_SUPPORT_NEG_INDICE_TILING_KEY)) {134+ gatherv2Op.Process();
135- GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);135+ #elif TILING_KEY_VAR == SIMD_LAST_GATHER_B32_SUPPORT_NEG_INDICE_TILING_KEY
136- Gatherv2SimdLastGather<int8_t, DTYPE_INDICES, true> gatherv2Op(&pipe);136+ GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);
137- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);137+ Gatherv2SimdLastGather<int32_t, DTYPE_INDICES, true> gatherv2Op(&pipe);
138- gatherv2Op.Process();138+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
139- } else if (TILING_KEY_IS(SIMD_LAST_GATHER_B16_SUPPORT_NEG_INDICE_TILING_KEY)) {139+ gatherv2Op.Process();
140- GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);140+ #elif TILING_KEY_VAR == SIMD_LAST_GATHER_B64_SUPPORT_NEG_INDICE_TILING_KEY
141- Gatherv2SimdLastGather<int16_t, DTYPE_INDICES, true> gatherv2Op(&pipe);141+ GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);
142- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);142+ Gatherv2SimdLastGather<int64_t, DTYPE_INDICES, true> gatherv2Op(&pipe);
143- gatherv2Op.Process();143+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
144- } else if (TILING_KEY_IS(SIMD_LAST_GATHER_B32_SUPPORT_NEG_INDICE_TILING_KEY)) {144+ gatherv2Op.Process();
145- GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);145+ #elif TILING_KEY_VAR == SIMD_GA_ALL_LOAD_TILING_KEY
146- Gatherv2SimdLastGather<int32_t, DTYPE_INDICES, true> gatherv2Op(&pipe);146+ GET_TILING_DATA_WITH_STRUCT(GatherV2GaAllLoadTilingData, tilingDataIn, tiling);
147- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);147+ Gatherv2GaAllLoad<DTYPE_INDICES, false> gatherv2Op(&pipe);
148- gatherv2Op.Process();148+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
149- } else if (TILING_KEY_IS(SIMD_LAST_GATHER_B64_SUPPORT_NEG_INDICE_TILING_KEY)) {149+ gatherv2Op.Process();
150- GET_TILING_DATA_WITH_STRUCT(GatherV2LastTilingData, tilingDataIn, tiling);150+ #elif TILING_KEY_VAR == SIMD_GA_ALL_LOAD_SUPPORT_NEG_INDICE_TILING_KEY
151- Gatherv2SimdLastGather<int64_t, DTYPE_INDICES, true> gatherv2Op(&pipe);151+ GET_TILING_DATA_WITH_STRUCT(GatherV2GaAllLoadTilingData, tilingDataIn, tiling);
152- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);152+ Gatherv2GaAllLoad<DTYPE_INDICES, true> gatherv2Op(&pipe);
153- gatherv2Op.Process();153+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
154- } else if (TILING_KEY_IS(SIMD_GA_ALL_LOAD_TILING_KEY)) {154+ gatherv2Op.Process();
155- GET_TILING_DATA_WITH_STRUCT(GatherV2GaAllLoadTilingData, tilingDataIn, tiling);155+ #elif TILING_KEY_VAR == SIMT_TWO_DIM_B8_INDEX_SIZE_64_TILING_KEY
156- Gatherv2GaAllLoad<DTYPE_INDICES, false> gatherv2Op(&pipe);156+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);
157- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);157+ Gatherv2SimtTwoDim<int8_t, DTYPE_INDICES, uint64_t> gatherv2Op;
158- gatherv2Op.Process();158+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
159- } else if (TILING_KEY_IS(SIMD_GA_ALL_LOAD_SUPPORT_NEG_INDICE_TILING_KEY)) {159+ gatherv2Op.Process();
160- GET_TILING_DATA_WITH_STRUCT(GatherV2GaAllLoadTilingData, tilingDataIn, tiling);160+ #elif TILING_KEY_VAR == SIMT_TWO_DIM_B16_INDEX_SIZE_64_TILING_KEY
161- Gatherv2GaAllLoad<DTYPE_INDICES, true> gatherv2Op(&pipe);161+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);
162- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);162+ Gatherv2SimtTwoDim<int16_t, DTYPE_INDICES, uint64_t> gatherv2Op;
163- gatherv2Op.Process();163+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
164- } else if (TILING_KEY_IS(TILING_KEY_X_B8_INDEX_SIZE_2D)) {164+ gatherv2Op.Process();
165- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);165+ #elif TILING_KEY_VAR == SIMT_TWO_DIM_B32_INDEX_SIZE_64_TILING_KEY
166- Gatherv2FullLoad<int8_t, DTYPE_INDICES, uint32_t> gatherv2Op(&pipe);166+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);
167- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);167+ Gatherv2SimtTwoDim<int32_t, DTYPE_INDICES, uint64_t> gatherv2Op;
168- gatherv2Op.Process();168+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
169- } else if (TILING_KEY_IS(TILING_KEY_X_B16_INDEX_SIZE_2D)) {169+ gatherv2Op.Process();
170- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);170+ #elif TILING_KEY_VAR == SIMT_TWO_DIM_B64_INDEX_SIZE_64_TILING_KEY
171- Gatherv2FullLoad<int16_t, DTYPE_INDICES, uint32_t> gatherv2Op(&pipe);171+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);
172- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);172+ Gatherv2SimtTwoDim<int64_t, DTYPE_INDICES, uint64_t> gatherv2Op;
173- gatherv2Op.Process();173+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
174- } else if (TILING_KEY_IS(TILING_KEY_X_B32_INDEX_SIZE_2D)) {174+ gatherv2Op.Process();
175- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);175+ #elif TILING_KEY_VAR == SIMT_TWO_DIM_B8_INDEX_SIZE_32_TILING_KEY
176- Gatherv2FullLoad<int32_t, DTYPE_INDICES, uint32_t> gatherv2Op(&pipe);176+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);
177- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);177+ Gatherv2SimtTwoDim<int8_t, DTYPE_INDICES, uint32_t> gatherv2Op;
178- gatherv2Op.Process();178+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
179- } else if (TILING_KEY_IS(TILING_KEY_X_B64_INDEX_SIZE_2D)) {179+ gatherv2Op.Process();
180- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);180+ #elif TILING_KEY_VAR == SIMT_TWO_DIM_B16_INDEX_SIZE_32_TILING_KEY
181- Gatherv2FullLoad<int64_t, DTYPE_INDICES, uint32_t> gatherv2Op(&pipe);181+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);
182- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);182+ Gatherv2SimtTwoDim<int16_t, DTYPE_INDICES, uint32_t> gatherv2Op;
183- gatherv2Op.Process();183+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
184- } else if (TILING_KEY_IS(SIMT_TWO_DIM_B8_INDEX_SIZE_64_TILING_KEY)) {184+ gatherv2Op.Process();
185- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);185+ #elif TILING_KEY_VAR == SIMT_TWO_DIM_B32_INDEX_SIZE_32_TILING_KEY
186- Gatherv2SimtTwoDim<int8_t, DTYPE_INDICES, uint64_t> gatherv2Op;186+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);
187- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);187+ Gatherv2SimtTwoDim<int32_t, DTYPE_INDICES, uint32_t> gatherv2Op;
188- gatherv2Op.Process();188+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
189- } else if (TILING_KEY_IS(SIMT_TWO_DIM_B16_INDEX_SIZE_64_TILING_KEY)) {189+ gatherv2Op.Process();
190- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);190+ #elif TILING_KEY_VAR == SIMT_TWO_DIM_B64_INDEX_SIZE_32_TILING_KEY
191- Gatherv2SimtTwoDim<int16_t, DTYPE_INDICES, uint64_t> gatherv2Op;191+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);
192- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);192+ Gatherv2SimtTwoDim<int64_t, DTYPE_INDICES, uint32_t> gatherv2Op;
193- gatherv2Op.Process();193+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
194- } else if (TILING_KEY_IS(SIMT_TWO_DIM_B32_INDEX_SIZE_64_TILING_KEY)) {194+ gatherv2Op.Process();
195- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);195+ #elif TILING_KEY_VAR == TILING_KEY_X_B8_INDEX_SIZE_64
196- Gatherv2SimtTwoDim<int32_t, DTYPE_INDICES, uint64_t> gatherv2Op;196+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);
197- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);197+ Gatherv2<int8_t, DTYPE_INDICES, uint64_t> gatherv2Op;
198- gatherv2Op.Process();198+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
199- } else if (TILING_KEY_IS(SIMT_TWO_DIM_B64_INDEX_SIZE_64_TILING_KEY)) {199+ gatherv2Op.Process();
200- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);200+ #elif TILING_KEY_VAR == TILING_KEY_X_B16_INDEX_SIZE_64
201- Gatherv2SimtTwoDim<int64_t, DTYPE_INDICES, uint64_t> gatherv2Op;201+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);
202- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);202+ Gatherv2<int16_t, DTYPE_INDICES, uint64_t> gatherv2Op;
203- gatherv2Op.Process();203+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
204- } else if (TILING_KEY_IS(SIMT_TWO_DIM_B8_INDEX_SIZE_32_TILING_KEY)) {204+ gatherv2Op.Process();
205- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);205+ #elif TILING_KEY_VAR == TILING_KEY_X_B32_INDEX_SIZE_64
206- Gatherv2SimtTwoDim<int8_t, DTYPE_INDICES, uint32_t> gatherv2Op;206+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);
207- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);207+ Gatherv2<int32_t, DTYPE_INDICES, uint64_t> gatherv2Op;
208- gatherv2Op.Process();208+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
209- } else if (TILING_KEY_IS(SIMT_TWO_DIM_B16_INDEX_SIZE_32_TILING_KEY)) {209+ gatherv2Op.Process();
210- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);210+ #elif TILING_KEY_VAR == TILING_KEY_X_B64_INDEX_SIZE_64
211- Gatherv2SimtTwoDim<int16_t, DTYPE_INDICES, uint32_t> gatherv2Op;211+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);
212- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);212+ Gatherv2<int64_t, DTYPE_INDICES, uint64_t> gatherv2Op;
213- gatherv2Op.Process();213+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
214- } else if (TILING_KEY_IS(SIMT_TWO_DIM_B32_INDEX_SIZE_32_TILING_KEY)) {214+ gatherv2Op.Process();
215- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);215+ #elif TILING_KEY_VAR == TILING_KEY_X_B8_INDEX_SIZE_32
216- Gatherv2SimtTwoDim<int32_t, DTYPE_INDICES, uint32_t> gatherv2Op;216+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);
217- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);217+ Gatherv2<int8_t, DTYPE_INDICES, uint32_t> gatherv2Op;
218- gatherv2Op.Process();218+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
219- } else if (TILING_KEY_IS(SIMT_TWO_DIM_B64_INDEX_SIZE_32_TILING_KEY)) {219+ gatherv2Op.Process();
220- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataSimtTwoDim, tilingDataIn, tiling);220+ #elif TILING_KEY_VAR == TILING_KEY_X_B16_INDEX_SIZE_32
221- Gatherv2SimtTwoDim<int64_t, DTYPE_INDICES, uint32_t> gatherv2Op;221+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);
222- gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);222+ Gatherv2<int16_t, DTYPE_INDICES, uint32_t> gatherv2Op;
223- gatherv2Op.Process();223+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
224- } else if (TILING_KEY_IS(TILING_KEY_EMPTY)) {224+ gatherv2Op.Process();
225- GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataEmptyInput, tilingDataIn, tiling);225+ #elif TILING_KEY_VAR == TILING_KEY_X_B32_INDEX_SIZE_32
226- Gatherv2Empty<int8_t> gatherv2Op;226+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);
227- gatherv2Op.Init(y, &tilingDataIn);227+ Gatherv2<int32_t, DTYPE_INDICES, uint32_t> gatherv2Op;
228- gatherv2Op.Process();228+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
229- }229+ gatherv2Op.Process();
230+ #elif TILING_KEY_VAR == TILING_KEY_X_B64_INDEX_SIZE_32
231+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);
232+ Gatherv2<int64_t, DTYPE_INDICES, uint32_t> gatherv2Op;
233+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
234+ gatherv2Op.Process();
235+ #elif TILING_KEY_VAR == TILING_KEY_X_B8_INDEX_SIZE_2D
236+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);
237+ Gatherv2FullLoad<int8_t, DTYPE_INDICES, uint32_t> gatherv2Op(&pipe);
238+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
239+ gatherv2Op.Process();
240+ #elif TILING_KEY_VAR == TILING_KEY_X_B16_INDEX_SIZE_2D
241+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);
242+ Gatherv2FullLoad<int16_t, DTYPE_INDICES, uint32_t> gatherv2Op(&pipe);
243+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
244+ gatherv2Op.Process();
245+ #elif TILING_KEY_VAR == TILING_KEY_X_B32_INDEX_SIZE_2D
246+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);
247+ Gatherv2FullLoad<int32_t, DTYPE_INDICES, uint32_t> gatherv2Op(&pipe);
248+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
249+ gatherv2Op.Process();
250+ #elif TILING_KEY_VAR == TILING_KEY_X_B64_INDEX_SIZE_2D
251+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingData, tilingDataIn, tiling);
252+ Gatherv2FullLoad<int64_t, DTYPE_INDICES, uint32_t> gatherv2Op(&pipe);
253+ gatherv2Op.Init(x, indices, axis, y, &tilingDataIn);
254+ gatherv2Op.Process();
255+ #elif TILING_KEY_VAR ==TILING_KEY_EMPTY
256+ GET_TILING_DATA_WITH_STRUCT(GatherV2TilingDataEmptyInput, tilingDataIn, tiling);
257+ Gatherv2Empty<int8_t> gatherv2Op;
258+ gatherv2Op.Init(y, &tilingDataIn);
259+ gatherv2Op.Process();
260+ #endif
230}261}
Mindex/scatter_elements/op_kernel/arch35/scatter_elements.h+2-0
@@ -16,6 +16,8 @@
16#define ASCENDC_SCATTER_ELEMENTS_H_16#define ASCENDC_SCATTER_ELEMENTS_H_
17 17 
18#include "kernel_operator.h"18#include "kernel_operator.h"
19+#include "basic_api/kernel_basic_intf.h"
20+#include "micro_api/kernel_micro_intf.h"
19 21 
20#include "../inc/platform.h"22#include "../inc/platform.h"
21#include "../inc/kernel_utils.h"23#include "../inc/kernel_utils.h"
Mindex/scatter_elements/op_kernel/arch35/scatter_elements_deterministic.h+2-0
@@ -16,6 +16,8 @@
16#define ASCENDC_SCATTER_ELEMENTS_DETERMINISTIC_H_16#define ASCENDC_SCATTER_ELEMENTS_DETERMINISTIC_H_
17 17 
18#include "kernel_operator.h"18#include "kernel_operator.h"
19+#include "basic_api/kernel_basic_intf.h"
20+#include "adv_api/sort/sort.h"
19 21 
20#include "../inc/platform.h"22#include "../inc/platform.h"
21#include "../inc/kernel_utils.h"23#include "../inc/kernel_utils.h"
Mindex/scatter_elements/op_kernel/scatter_elements_apt.cpp+662-449
@@ -127,509 +127,722 @@ extern "C" __global__ __aicore__ void scatter_elements(GM_ADDR x, GM_ADDR indice
127 GET_TILING_DATA(tilingData, tiling);127 GET_TILING_DATA(tilingData, tiling);
128 TPipe pipe;128 TPipe pipe;
129 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);129 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);
130- if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_NONE_B8)) {130+ 
131- if (tilingData.isDeterministic) {131+ #if (ORIG_DTYPE_DATA == DT_INT8 || ORIG_DTYPE_DATA == DT_UINT8 || ORIG_DTYPE_DATA == DT_BOOL) && ORIG_DTYPE_INDICES == DT_INT32
132- KernelScatterElementsDeterm<int8_t, int32_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);132+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_NONE_B8);
133- op.Init(x, indices, updates, y);133+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_NONE_B8);
134- op.Process();134+ 
135- } else {135+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_NONE_B8
136- KernelScatterElements<int8_t, int32_t, uint32_t, uint32_t, REDU_NONE> op(pipe);136+ if (tilingData.isDeterministic) {
137+ KernelScatterElementsDeterm<int8_t, int32_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);
138+ op.Init(x, indices, updates, y);
139+ op.Process();
140+ } else {
141+ KernelScatterElements<int8_t, int32_t, uint32_t, uint32_t, REDU_NONE> op(pipe);
142+ op.Init(x, indices, updates, y, userWS, &tilingData);
143+ op.Process();
144+ }
145+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_NONE_B8
146+ if (tilingData.isDeterministic) {
147+ KernelScatterElementsDeterm<int8_t, int32_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
148+ op.Init(x, indices, updates, y);
149+ op.Process();
150+ } else {
151+ KernelScatterElements<int8_t, int32_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
152+ op.Init(x, indices, updates, y, userWS, &tilingData);
153+ op.Process();
154+ }
155+ #endif
156+ #endif
157+ 
158+ #if (ORIG_DTYPE_DATA == DT_INT8 || ORIG_DTYPE_DATA == DT_UINT8 || ORIG_DTYPE_DATA == DT_BOOL) && ORIG_DTYPE_INDICES == DT_INT64
159+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_NONE_B8);
160+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_NONE_B8);
161+ 
162+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_NONE_B8
163+ if (tilingData.isDeterministic) {
164+ KernelScatterElementsDeterm<int8_t, int64_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);
165+ op.Init(x, indices, updates, y);
166+ op.Process();
167+ } else {
168+ KernelScatterElements<int8_t, int64_t, uint32_t, uint32_t, REDU_NONE> op(pipe);
169+ op.Init(x, indices, updates, y, userWS, &tilingData);
170+ op.Process();
171+ }
172+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_NONE_B8
173+ if (tilingData.isDeterministic) {
174+ KernelScatterElementsDeterm<int8_t, int64_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
175+ op.Init(x, indices, updates, y);
176+ op.Process();
177+ } else {
178+ KernelScatterElements<int8_t, int64_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
179+ op.Init(x, indices, updates, y, userWS, &tilingData);
180+ op.Process();
181+ }
182+ #endif
183+ #endif
184+ 
185+ #if (ORIG_DTYPE_DATA == DT_INT16 || ORIG_DTYPE_DATA == DT_FLOAT16 || ORIG_DTYPE_DATA == DT_BF16) && ORIG_DTYPE_INDICES == DT_INT32
186+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_NONE_B16);
187+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_NONE_B16);
188+ 
189+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_NONE_B16
190+ if (tilingData.isDeterministic) {
191+ KernelScatterElementsDeterm<int16_t, int32_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);
192+ op.Init(x, indices, updates, y);
193+ op.Process();
194+ } else {
195+ KernelScatterElements<int16_t, int32_t, uint32_t, uint32_t, REDU_NONE> op(pipe);
196+ op.Init(x, indices, updates, y, userWS, &tilingData);
197+ op.Process();
198+ }
199+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_NONE_B16
200+ if (tilingData.isDeterministic) {
201+ KernelScatterElementsDeterm<int16_t, int32_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
202+ op.Init(x, indices, updates, y);
203+ op.Process();
204+ } else {
205+ KernelScatterElements<int16_t, int32_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
206+ op.Init(x, indices, updates, y, userWS, &tilingData);
207+ op.Process();
208+ }
209+ #endif
210+ #endif
211+ 
212+ #if (ORIG_DTYPE_DATA == DT_INT16 || ORIG_DTYPE_DATA == DT_FLOAT16 || ORIG_DTYPE_DATA == DT_BF16) && ORIG_DTYPE_INDICES == DT_INT64
213+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_NONE_B16);
214+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_NONE_B16);
215+ 
216+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_NONE_B16
217+ if (tilingData.isDeterministic) {
218+ KernelScatterElementsDeterm<int16_t, int64_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);
219+ op.Init(x, indices, updates, y);
220+ op.Process();
221+ } else {
222+ KernelScatterElements<int16_t, int64_t, uint32_t, uint32_t, REDU_NONE> op(pipe);
223+ op.Init(x, indices, updates, y, userWS, &tilingData);
224+ op.Process();
225+ }
226+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_NONE_B16
227+ if (tilingData.isDeterministic) {
228+ KernelScatterElementsDeterm<int16_t, int64_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
229+ op.Init(x, indices, updates, y);
230+ op.Process();
231+ } else {
232+ KernelScatterElements<int16_t, int64_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
233+ op.Init(x, indices, updates, y, userWS, &tilingData);
234+ op.Process();
235+ }
236+ #endif
237+ #endif
238+ 
239+ #if (ORIG_DTYPE_DATA == DT_INT32 || ORIG_DTYPE_DATA == DT_FLOAT) && ORIG_DTYPE_INDICES == DT_INT32
240+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_NONE_B32);
241+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_NONE_B32);
242+ 
243+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_NONE_B32
244+ if (tilingData.isDeterministic) {
245+ KernelScatterElementsDeterm<int32_t, int32_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);
246+ op.Init(x, indices, updates, y);
247+ op.Process();
248+ } else {
249+ KernelScatterElements<int32_t, int32_t, uint32_t, uint32_t, REDU_NONE> op(pipe);
250+ op.Init(x, indices, updates, y, userWS, &tilingData);
251+ op.Process();
252+ }
253+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_NONE_B32
254+ if (tilingData.isDeterministic) {
255+ KernelScatterElementsDeterm<int32_t, int32_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
256+ op.Init(x, indices, updates, y);
257+ op.Process();
258+ } else {
259+ KernelScatterElements<int32_t, int32_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
260+ op.Init(x, indices, updates, y, userWS, &tilingData);
261+ op.Process();
262+ }
263+ #endif
264+ #endif
265+ 
266+ #if (ORIG_DTYPE_DATA == DT_INT32 || ORIG_DTYPE_DATA == DT_FLOAT) && ORIG_DTYPE_INDICES == DT_INT64
267+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_NONE_B32);
268+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_NONE_B32);
269+ 
270+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_NONE_B32
271+ if (tilingData.isDeterministic) {
272+ KernelScatterElementsDeterm<int32_t, int64_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);
273+ op.Init(x, indices, updates, y);
274+ op.Process();
275+ } else {
276+ KernelScatterElements<int32_t, int64_t, uint32_t, uint32_t, REDU_NONE> op(pipe);
277+ op.Init(x, indices, updates, y, userWS, &tilingData);
278+ op.Process();
279+ }
280+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_NONE_B32
281+ if (tilingData.isDeterministic) {
282+ KernelScatterElementsDeterm<int32_t, int64_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
283+ op.Init(x, indices, updates, y);
284+ op.Process();
285+ } else {
286+ KernelScatterElements<int32_t, int64_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
287+ op.Init(x, indices, updates, y, userWS, &tilingData);
288+ op.Process();
289+ }
290+ #endif
291+ #endif
292+ 
293+ #if (ORIG_DTYPE_DATA == DT_INT64 || ORIG_DTYPE_DATA == DT_DOUBLE) && ORIG_DTYPE_INDICES == DT_INT32
294+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_NONE_B64);
295+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_NONE_B64);
296+ 
297+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_NONE_B64
298+ if (tilingData.isDeterministic) {
299+ KernelScatterElementsDeterm<int64_t, int32_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);
300+ op.Init(x, indices, updates, y);
301+ op.Process();
302+ } else {
303+ KernelScatterElements<int64_t, int32_t, uint32_t, uint32_t, REDU_NONE> op(pipe);
304+ op.Init(x, indices, updates, y, userWS, &tilingData);
305+ op.Process();
306+ }
307+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_NONE_B64
308+ if (tilingData.isDeterministic) {
309+ KernelScatterElementsDeterm<int64_t, int32_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
310+ op.Init(x, indices, updates, y);
311+ op.Process();
312+ } else {
313+ KernelScatterElements<int64_t, int32_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
314+ op.Init(x, indices, updates, y, userWS, &tilingData);
315+ op.Process();
316+ }
317+ #endif
318+ #endif
319+ 
320+ #if (ORIG_DTYPE_DATA == DT_INT64 || ORIG_DTYPE_DATA == DT_DOUBLE) && ORIG_DTYPE_INDICES == DT_INT64
321+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_NONE_B64);
322+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_NONE_B64);
323+ 
324+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_NONE_B64
325+ if (tilingData.isDeterministic) {
326+ KernelScatterElementsDeterm<int64_t, int64_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);
327+ op.Init(x, indices, updates, y);
328+ op.Process();
329+ } else {
330+ KernelScatterElements<int64_t, int64_t, uint32_t, uint32_t, REDU_NONE> op(pipe);
331+ op.Init(x, indices, updates, y, userWS, &tilingData);
332+ op.Process();
333+ }
334+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_NONE_B64
335+ if (tilingData.isDeterministic) {
336+ KernelScatterElementsDeterm<int64_t, int64_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
337+ op.Init(x, indices, updates, y);
338+ op.Process();
339+ } else {
340+ KernelScatterElements<int64_t, int64_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
341+ op.Init(x, indices, updates, y, userWS, &tilingData);
342+ op.Process();
343+ }
344+ #endif
345+ #endif
346+ 
347+ #if ORIG_DTYPE_DATA == DT_FLOAT && ORIG_DTYPE_INDICES == DT_INT32
348+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_FP32);
349+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_FP32);
350+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_FP32);
351+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_FP32);
352+ 
353+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_ADD_FP32
354+ if (tilingData.isDeterministic) {
355+ KernelScatterElementsDeterm<float, int32_t, uint32_t, REDU_ADD> op(&tilingData, &pipe);
356+ op.Init(x, indices, updates, y);
357+ op.Process();
358+ } else {
359+ KernelScatterElements<float, int32_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
360+ op.Init(x, indices, updates, y, userWS, &tilingData);
361+ op.Process();
362+ }
363+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_MUL_FP32
364+ KernelScatterElements<float, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
137 op.Init(x, indices, updates, y, userWS, &tilingData);365 op.Init(x, indices, updates, y, userWS, &tilingData);
138 op.Process();366 op.Process();
139- }367+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_ADD_FP32
140- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_NONE_B16)) {368+ if (tilingData.isDeterministic) {
141- if (tilingData.isDeterministic) {369+ KernelScatterElementsDeterm<float, int32_t, uint64_t, REDU_ADD> op(&tilingData, &pipe);
142- KernelScatterElementsDeterm<int16_t, int32_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);370+ op.Init(x, indices, updates, y);
143- op.Init(x, indices, updates, y);371+ op.Process();
144- op.Process();372+ } else {
145- } else {373+ KernelScatterElements<float, int32_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
146- KernelScatterElements<int16_t, int32_t, uint32_t, uint32_t, REDU_NONE> op(pipe);374+ op.Init(x, indices, updates, y, userWS, &tilingData);
375+ op.Process();
376+ }
377+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_MUL_FP32
378+ KernelScatterElements<float, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
147 op.Init(x, indices, updates, y, userWS, &tilingData);379 op.Init(x, indices, updates, y, userWS, &tilingData);
148 op.Process();380 op.Process();
149- }381+ #endif
150- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_NONE_B32)) {382+ #endif
151- if (tilingData.isDeterministic) {383+ 
152- KernelScatterElementsDeterm<int32_t, int32_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);384+ #if ORIG_DTYPE_DATA == DT_FLOAT && ORIG_DTYPE_INDICES == DT_INT64
153- op.Init(x, indices, updates, y);385+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_FP32);
154- op.Process();386+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_FP32);
155- } else {387+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_FP32);
156- KernelScatterElements<int32_t, int32_t, uint32_t, uint32_t, REDU_NONE> op(pipe);388+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_FP32);
389+ 
390+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_ADD_FP32
391+ if (tilingData.isDeterministic) {
392+ KernelScatterElementsDeterm<float, int64_t, uint32_t, REDU_ADD> op(&tilingData, &pipe);
393+ op.Init(x, indices, updates, y);
394+ op.Process();
395+ } else {
396+ KernelScatterElements<float, int64_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
397+ op.Init(x, indices, updates, y, userWS, &tilingData);
398+ op.Process();
399+ }
400+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_MUL_FP32
401+ KernelScatterElements<float, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
157 op.Init(x, indices, updates, y, userWS, &tilingData);402 op.Init(x, indices, updates, y, userWS, &tilingData);
158 op.Process();403 op.Process();
159- }404+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_ADD_FP32
160- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_NONE_B64)) {405+ if (tilingData.isDeterministic) {
161- if (tilingData.isDeterministic) {406+ KernelScatterElementsDeterm<float, int64_t, uint64_t, REDU_ADD> op(&tilingData, &pipe);
162- KernelScatterElementsDeterm<int64_t, int32_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);407+ op.Init(x, indices, updates, y);
163- op.Init(x, indices, updates, y);408+ op.Process();
164- op.Process();409+ } else {
165- } else {410+ KernelScatterElements<float, int64_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
166- KernelScatterElements<int64_t, int32_t, uint32_t, uint32_t, REDU_NONE> op(pipe);411+ op.Init(x, indices, updates, y, userWS, &tilingData);
412+ op.Process();
413+ }
414+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_MUL_FP32
415+ KernelScatterElements<float, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
167 op.Init(x, indices, updates, y, userWS, &tilingData);416 op.Init(x, indices, updates, y, userWS, &tilingData);
168 op.Process();417 op.Process();
169- }418+ #endif
170- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_FP32)) {419+ #endif
171- if (tilingData.isDeterministic) {420+ 
172- KernelScatterElementsDeterm<float, int32_t, uint32_t, REDU_ADD> op(&tilingData, &pipe);421+ #if ORIG_DTYPE_DATA == DT_FLOAT16 && ORIG_DTYPE_INDICES == DT_INT32
173- op.Init(x, indices, updates, y);422+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_FP16);
174- op.Process();423+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_FP16);
175- } else {424+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_FP16);
176- KernelScatterElements<float, int32_t, uint32_t, uint32_t, REDU_ADD> op(pipe);425+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_FP16);
426+ 
427+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_ADD_FP16
428+ if (tilingData.isDeterministic) {
429+ KernelScatterElementsDeterm<half, int32_t, uint32_t, REDU_ADD> op(&tilingData, &pipe);
430+ op.Init(x, indices, updates, y);
431+ op.Process();
432+ } else {
433+ KernelScatterElements<half, int32_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
434+ op.Init(x, indices, updates, y, userWS, &tilingData);
435+ op.Process();
436+ }
437+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_MUL_FP16
438+ KernelScatterElements<half, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
177 op.Init(x, indices, updates, y, userWS, &tilingData);439 op.Init(x, indices, updates, y, userWS, &tilingData);
178 op.Process();440 op.Process();
179- }441+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_ADD_FP16
180- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_FP16)) {442+ if (tilingData.isDeterministic) {
181- if (tilingData.isDeterministic) {443+ KernelScatterElementsDeterm<half, int32_t, uint64_t, REDU_ADD> op(&tilingData, &pipe);
182- KernelScatterElementsDeterm<half, int32_t, uint32_t, REDU_ADD> op(&tilingData, &pipe);444+ op.Init(x, indices, updates, y);
183- op.Init(x, indices, updates, y);445+ op.Process();
184- op.Process();446+ } else {
185- } else {447+ KernelScatterElements<half, int32_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
186- KernelScatterElements<half, int32_t, uint32_t, uint32_t, REDU_ADD> op(pipe);448+ op.Init(x, indices, updates, y, userWS, &tilingData);
449+ op.Process();
450+ }
451+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_MUL_FP16
452+ KernelScatterElements<half, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
187 op.Init(x, indices, updates, y, userWS, &tilingData);453 op.Init(x, indices, updates, y, userWS, &tilingData);
188 op.Process();454 op.Process();
189- }455+ #endif
190- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_INT8)) {456+ #endif
191- KernelScatterElements<int8_t, int32_t, uint32_t, int32_t, REDU_ADD> op(pipe);457+ 
192- op.Init(x, indices, updates, y, userWS, &tilingData);458+ #if ORIG_DTYPE_DATA == DT_FLOAT16 && ORIG_DTYPE_INDICES == DT_INT64
193- op.Process();459+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_FP16);
194- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_INT32)) {460+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_FP16);
195- KernelScatterElements<int32_t, int32_t, uint32_t, uint32_t, REDU_ADD> op(pipe);461+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_FP16);
196- op.Init(x, indices, updates, y, userWS, &tilingData);462+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_FP16);
197- op.Process();463+ 
198- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_UINT8)) {464+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_ADD_FP16
199- KernelScatterElements<uint8_t, int32_t, uint32_t, uint32_t, REDU_ADD> op(pipe);465+ if (tilingData.isDeterministic) {
200- op.Init(x, indices, updates, y, userWS, &tilingData);466+ KernelScatterElementsDeterm<half, int64_t, uint32_t, REDU_ADD> op(&tilingData, &pipe);
201- op.Process();467+ op.Init(x, indices, updates, y);
202- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_INT16)) {468+ op.Process();
203- KernelScatterElements<int16_t, int32_t, uint32_t, int32_t, REDU_ADD> op(pipe);469+ } else {
204- op.Init(x, indices, updates, y, userWS, &tilingData);470+ KernelScatterElements<half, int64_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
205- op.Process();471+ op.Init(x, indices, updates, y, userWS, &tilingData);
206- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_BF16)) {472+ op.Process();
207- if (tilingData.isDeterministic) {473+ }
208- KernelScatterElementsDeterm<bfloat16_t, int32_t, uint32_t, REDU_ADD> op(&tilingData, &pipe);474+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_MUL_FP16
209- op.Init(x, indices, updates, y);475+ KernelScatterElements<half, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
210- op.Process();
211- } else {
212- KernelScatterElements<bfloat16_t, int32_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
213 op.Init(x, indices, updates, y, userWS, &tilingData);476 op.Init(x, indices, updates, y, userWS, &tilingData);
214 op.Process();477 op.Process();
215- }478+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_ADD_FP16
216- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_INT64)) {479+ if (tilingData.isDeterministic) {
217- KernelScatterElements<int64_t, int32_t, uint32_t, uint32_t, REDU_ADD> op(pipe);480+ KernelScatterElementsDeterm<half, int64_t, uint64_t, REDU_ADD> op(&tilingData, &pipe);
218- op.Init(x, indices, updates, y, userWS, &tilingData);481+ op.Init(x, indices, updates, y);
219- op.Process();482+ op.Process();
220- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_BOOL)) {483+ } else {
221- KernelScatterElements<uint8_t, int32_t, uint32_t, half, REDU_ADD> op(pipe);484+ KernelScatterElements<half, int64_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
222- op.Init(x, indices, updates, y, userWS, &tilingData);485+ op.Init(x, indices, updates, y, userWS, &tilingData);
223- op.Process();486+ op.Process();
224- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_FP32)) {487+ }
225- KernelScatterElements<float, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);488+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_MUL_FP16
226- op.Init(x, indices, updates, y, userWS, &tilingData);489+ KernelScatterElements<half, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
227- op.Process();
228- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_FP16)) {
229- KernelScatterElements<half, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
230- op.Init(x, indices, updates, y, userWS, &tilingData);
231- op.Process();
232- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_INT8)) {
233- KernelScatterElements<int8_t, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
234- op.Init(x, indices, updates, y, userWS, &tilingData);
235- op.Process();
236- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_INT32)) {
237- KernelScatterElements<int32_t, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
238- op.Init(x, indices, updates, y, userWS, &tilingData);
239- op.Process();
240- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_UINT8)) {
241- KernelScatterElements<uint8_t, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
242- op.Init(x, indices, updates, y, userWS, &tilingData);
243- op.Process();
244- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_INT16)) {
245- KernelScatterElements<int16_t, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
246- op.Init(x, indices, updates, y, userWS, &tilingData);
247- op.Process();
248- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_BF16)) {
249- KernelScatterElements<bfloat16_t, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
250- op.Init(x, indices, updates, y, userWS, &tilingData);
251- op.Process();
252- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_INT64)) {
253- KernelScatterElements<int64_t, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
254- op.Init(x, indices, updates, y, userWS, &tilingData);
255- op.Process();
256- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_NONE_B8)) {
257- if (tilingData.isDeterministic) {
258- KernelScatterElementsDeterm<int8_t, int64_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);
259- op.Init(x, indices, updates, y);
260- op.Process();
261- } else {
262- KernelScatterElements<int8_t, int64_t, uint32_t, uint32_t, REDU_NONE> op(pipe);
263 op.Init(x, indices, updates, y, userWS, &tilingData);490 op.Init(x, indices, updates, y, userWS, &tilingData);
264 op.Process();491 op.Process();
265- }492+ #endif
266- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_NONE_B16)) {493+ #endif
267- if (tilingData.isDeterministic) {494+ 
268- KernelScatterElementsDeterm<int16_t, int64_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);495+ #if ORIG_DTYPE_DATA == DT_BF16 && ORIG_DTYPE_INDICES == DT_INT32
269- op.Init(x, indices, updates, y);496+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_BF16);
270- op.Process();497+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_BF16);
271- } else {498+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_BF16);
272- KernelScatterElements<int16_t, int64_t, uint32_t, uint32_t, REDU_NONE> op(pipe);499+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_BF16);
500+ 
501+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_ADD_BF16
502+ if (tilingData.isDeterministic) {
503+ KernelScatterElementsDeterm<bfloat16_t, int32_t, uint32_t, REDU_ADD> op(&tilingData, &pipe);
504+ op.Init(x, indices, updates, y);
505+ op.Process();
506+ } else {
507+ KernelScatterElements<bfloat16_t, int32_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
508+ op.Init(x, indices, updates, y, userWS, &tilingData);
509+ op.Process();
510+ }
511+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_MUL_BF16
512+ KernelScatterElements<bfloat16_t, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
273 op.Init(x, indices, updates, y, userWS, &tilingData);513 op.Init(x, indices, updates, y, userWS, &tilingData);
274 op.Process();514 op.Process();
275- }515+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_ADD_BF16
276- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_NONE_B32)) {516+ if (tilingData.isDeterministic) {
277- if (tilingData.isDeterministic) {517+ KernelScatterElementsDeterm<bfloat16_t, int32_t, uint64_t, REDU_ADD> op(&tilingData, &pipe);
278- KernelScatterElementsDeterm<int32_t, int64_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);518+ op.Init(x, indices, updates, y);
279- op.Init(x, indices, updates, y);519+ op.Process();
280- op.Process();520+ } else {
281- } else {521+ KernelScatterElements<bfloat16_t, int32_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
282- KernelScatterElements<int32_t, int64_t, uint32_t, uint32_t, REDU_NONE> op(pipe);522+ op.Init(x, indices, updates, y, userWS, &tilingData);
523+ op.Process();
524+ }
525+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_MUL_BF16
526+ KernelScatterElements<bfloat16_t, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
283 op.Init(x, indices, updates, y, userWS, &tilingData);527 op.Init(x, indices, updates, y, userWS, &tilingData);
284 op.Process();528 op.Process();
285- }529+ #endif
286- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_NONE_B64)) {530+ #endif
287- if (tilingData.isDeterministic) {531+ 
288- KernelScatterElementsDeterm<int64_t, int64_t, uint32_t, REDU_NONE> op(&tilingData, &pipe);532+ #if ORIG_DTYPE_DATA == DT_BF16 && ORIG_DTYPE_INDICES == DT_INT64
289- op.Init(x, indices, updates, y);533+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_BF16);
290- op.Process();534+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_BF16);
291- } else {535+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_BF16);
292- KernelScatterElements<int64_t, int64_t, uint32_t, uint32_t, REDU_NONE> op(pipe);536+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_BF16);
537+ 
538+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_ADD_BF16
539+ if (tilingData.isDeterministic) {
540+ KernelScatterElementsDeterm<bfloat16_t, int64_t, uint32_t, REDU_ADD> op(&tilingData, &pipe);
541+ op.Init(x, indices, updates, y);
542+ op.Process();
543+ } else {
544+ KernelScatterElements<bfloat16_t, int64_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
545+ op.Init(x, indices, updates, y, userWS, &tilingData);
546+ op.Process();
547+ }
548+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_MUL_BF16
549+ KernelScatterElements<bfloat16_t, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
293 op.Init(x, indices, updates, y, userWS, &tilingData);550 op.Init(x, indices, updates, y, userWS, &tilingData);
294 op.Process();551 op.Process();
295- }552+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_ADD_BF16
296- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_FP32)) {553+ if (tilingData.isDeterministic) {
297- if (tilingData.isDeterministic) {554+ KernelScatterElementsDeterm<bfloat16_t, int64_t, uint64_t, REDU_ADD> op(&tilingData, &pipe);
298- KernelScatterElementsDeterm<float, int64_t, uint32_t, REDU_ADD> op(&tilingData, &pipe);555+ op.Init(x, indices, updates, y);
299- op.Init(x, indices, updates, y);556+ op.Process();
300- op.Process();557+ } else {
301- } else {558+ KernelScatterElements<bfloat16_t, int64_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
302- KernelScatterElements<float, int64_t, uint32_t, uint32_t, REDU_ADD> op(pipe);559+ op.Init(x, indices, updates, y, userWS, &tilingData);
560+ op.Process();
561+ }
562+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_MUL_BF16
563+ KernelScatterElements<bfloat16_t, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
303 op.Init(x, indices, updates, y, userWS, &tilingData);564 op.Init(x, indices, updates, y, userWS, &tilingData);
304 op.Process();565 op.Process();
305- }566+ #endif
306- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_FP16)) {567+ #endif
307- if (tilingData.isDeterministic) {568+ 
308- KernelScatterElementsDeterm<half, int64_t, uint32_t, REDU_ADD> op(&tilingData, &pipe);569+ #if ORIG_DTYPE_DATA == DT_INT8 && ORIG_DTYPE_INDICES == DT_INT32
309- op.Init(x, indices, updates, y);570+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_INT8);
310- op.Process();571+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_INT8);
311- } else {572+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_INT8);
312- KernelScatterElements<half, int64_t, uint32_t, uint32_t, REDU_ADD> op(pipe);573+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_INT8);
574+ 
575+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_ADD_INT8
576+ KernelScatterElements<int8_t, int32_t, uint32_t, int32_t, REDU_ADD> op(pipe);
313 op.Init(x, indices, updates, y, userWS, &tilingData);577 op.Init(x, indices, updates, y, userWS, &tilingData);
314 op.Process();578 op.Process();
315- }579+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_MUL_INT8
316- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_INT8)) {580+ KernelScatterElements<int8_t, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
317- KernelScatterElements<int8_t, int64_t, uint32_t, int32_t, REDU_ADD> op(pipe);
318- op.Init(x, indices, updates, y, userWS, &tilingData);
319- op.Process();
320- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_INT32)) {
321- KernelScatterElements<int32_t, int64_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
322- op.Init(x, indices, updates, y, userWS, &tilingData);
323- op.Process();
324- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_UINT8)) {
325- KernelScatterElements<uint8_t, int64_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
326- op.Init(x, indices, updates, y, userWS, &tilingData);
327- op.Process();
328- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_INT16)) {
329- KernelScatterElements<int16_t, int64_t, uint32_t, int32_t, REDU_ADD> op(pipe);
330- op.Init(x, indices, updates, y, userWS, &tilingData);
331- op.Process();
332- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_BF16)) {
333- if (tilingData.isDeterministic) {
334- KernelScatterElementsDeterm<bfloat16_t, int64_t, uint32_t, REDU_ADD> op(&tilingData, &pipe);
335- op.Init(x, indices, updates, y);
336- op.Process();
337- } else {
338- KernelScatterElements<bfloat16_t, int64_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
339 op.Init(x, indices, updates, y, userWS, &tilingData);581 op.Init(x, indices, updates, y, userWS, &tilingData);
340 op.Process();582 op.Process();
341- }583+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_ADD_INT8
342- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_INT64)) {584+ KernelScatterElements<int8_t, int32_t, uint64_t, int32_t, REDU_ADD> op(pipe);
343- KernelScatterElements<int64_t, int64_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
344- op.Init(x, indices, updates, y, userWS, &tilingData);
345- op.Process();
346- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_BOOL)) {
347- KernelScatterElements<uint8_t, int64_t, uint32_t, half, REDU_ADD> op(pipe);
348- op.Init(x, indices, updates, y, userWS, &tilingData);
349- op.Process();
350- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_FP32)) {
351- KernelScatterElements<float, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
352- op.Init(x, indices, updates, y, userWS, &tilingData);
353- op.Process();
354- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_FP16)) {
355- KernelScatterElements<half, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
356- op.Init(x, indices, updates, y, userWS, &tilingData);
357- op.Process();
358- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_INT8)) {
359- KernelScatterElements<int8_t, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
360- op.Init(x, indices, updates, y, userWS, &tilingData);
361- op.Process();
362- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_INT32)) {
363- KernelScatterElements<int32_t, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
364- op.Init(x, indices, updates, y, userWS, &tilingData);
365- op.Process();
366- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_UINT8)) {
367- KernelScatterElements<uint8_t, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
368- op.Init(x, indices, updates, y, userWS, &tilingData);
369- op.Process();
370- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_INT16)) {
371- KernelScatterElements<int16_t, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
372- op.Init(x, indices, updates, y, userWS, &tilingData);
373- op.Process();
374- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_BF16)) {
375- KernelScatterElements<bfloat16_t, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
376- op.Init(x, indices, updates, y, userWS, &tilingData);
377- op.Process();
378- } else if (TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_INT64)) {
379- KernelScatterElements<int64_t, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
380- op.Init(x, indices, updates, y, userWS, &tilingData);
381- op.Process();
382- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_NONE_B8)) {
383- if (tilingData.isDeterministic) {
384- KernelScatterElementsDeterm<int8_t, int32_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
385- op.Init(x, indices, updates, y);
386- op.Process();
387- } else {
388- KernelScatterElements<int8_t, int32_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
389 op.Init(x, indices, updates, y, userWS, &tilingData);585 op.Init(x, indices, updates, y, userWS, &tilingData);
390 op.Process();586 op.Process();
391- }587+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_MUL_INT8
392- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_NONE_B16)) {588+ KernelScatterElements<int8_t, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
393- if (tilingData.isDeterministic) {
394- KernelScatterElementsDeterm<int16_t, int32_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
395- op.Init(x, indices, updates, y);
396- op.Process();
397- } else {
398- KernelScatterElements<int16_t, int32_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
399 op.Init(x, indices, updates, y, userWS, &tilingData);589 op.Init(x, indices, updates, y, userWS, &tilingData);
400 op.Process();590 op.Process();
401- }591+ #endif
402- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_NONE_B32)) {592+ #endif
403- if (tilingData.isDeterministic) {593+ 
404- KernelScatterElementsDeterm<int32_t, int32_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);594+ #if ORIG_DTYPE_DATA == DT_INT8 && ORIG_DTYPE_INDICES == DT_INT64
405- op.Init(x, indices, updates, y);595+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_INT8);
406- op.Process();596+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_INT8);
407- } else {597+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_INT8);
408- KernelScatterElements<int32_t, int32_t, uint64_t, uint32_t, REDU_NONE> op(pipe);598+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_INT8);
599+ 
600+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_ADD_INT8
601+ KernelScatterElements<int8_t, int64_t, uint32_t, int32_t, REDU_ADD> op(pipe);
409 op.Init(x, indices, updates, y, userWS, &tilingData);602 op.Init(x, indices, updates, y, userWS, &tilingData);
410 op.Process();603 op.Process();
411- }604+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_MUL_INT8
412- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_NONE_B64)) {605+ KernelScatterElements<int8_t, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
413- if (tilingData.isDeterministic) {
414- KernelScatterElementsDeterm<int64_t, int32_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
415- op.Init(x, indices, updates, y);
416- op.Process();
417- } else {
418- KernelScatterElements<int64_t, int32_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
419 op.Init(x, indices, updates, y, userWS, &tilingData);606 op.Init(x, indices, updates, y, userWS, &tilingData);
420 op.Process();607 op.Process();
421- }608+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_ADD_INT8
422- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_FP32)) {609+ KernelScatterElements<int8_t, int64_t, uint64_t, int32_t, REDU_ADD> op(pipe);
423- if (tilingData.isDeterministic) {
424- KernelScatterElementsDeterm<float, int32_t, uint64_t, REDU_ADD> op(&tilingData, &pipe);
425- op.Init(x, indices, updates, y);
426- op.Process();
427- } else {
428- KernelScatterElements<float, int32_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
429 op.Init(x, indices, updates, y, userWS, &tilingData);610 op.Init(x, indices, updates, y, userWS, &tilingData);
430 op.Process();611 op.Process();
431- }612+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_MUL_INT8
432- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_FP16)) {613+ KernelScatterElements<int8_t, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
433- if (tilingData.isDeterministic) {
434- KernelScatterElementsDeterm<half, int32_t, uint64_t, REDU_ADD> op(&tilingData, &pipe);
435- op.Init(x, indices, updates, y);
436- op.Process();
437- } else {
438- KernelScatterElements<half, int32_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
439 op.Init(x, indices, updates, y, userWS, &tilingData);614 op.Init(x, indices, updates, y, userWS, &tilingData);
440 op.Process();615 op.Process();
441- }616+ #endif
442- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_INT8)) {617+ #endif
443- KernelScatterElements<int8_t, int32_t, uint64_t, int32_t, REDU_ADD> op(pipe);618+ 
444- op.Init(x, indices, updates, y, userWS, &tilingData);619+ #if ORIG_DTYPE_DATA == DT_UINT8 && ORIG_DTYPE_INDICES == DT_INT32
445- op.Process();620+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_UINT8);
446- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_INT32)) {621+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_UINT8);
447- KernelScatterElements<int32_t, int32_t, uint64_t, uint32_t, REDU_ADD> op(pipe);622+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_UINT8);
448- op.Init(x, indices, updates, y, userWS, &tilingData);623+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_UINT8);
449- op.Process();624+ 
450- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_UINT8)) {625+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_ADD_UINT8
451- KernelScatterElements<uint8_t, int32_t, uint64_t, uint32_t, REDU_ADD> op(pipe);626+ KernelScatterElements<uint8_t, int32_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
452- op.Init(x, indices, updates, y, userWS, &tilingData);
453- op.Process();
454- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_INT16)) {
455- KernelScatterElements<int16_t, int32_t, uint64_t, int32_t, REDU_ADD> op(pipe);
456- op.Init(x, indices, updates, y, userWS, &tilingData);
457- op.Process();
458- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_BF16)) {
459- if (tilingData.isDeterministic) {
460- KernelScatterElementsDeterm<bfloat16_t, int32_t, uint64_t, REDU_ADD> op(&tilingData, &pipe);
461- op.Init(x, indices, updates, y);
462- op.Process();
463- } else {
464- KernelScatterElements<bfloat16_t, int32_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
465 op.Init(x, indices, updates, y, userWS, &tilingData);627 op.Init(x, indices, updates, y, userWS, &tilingData);
466 op.Process();628 op.Process();
467- }629+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_MUL_UINT8
468- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_INT64)) {630+ KernelScatterElements<uint8_t, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
469- KernelScatterElements<int64_t, int32_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
470- op.Init(x, indices, updates, y, userWS, &tilingData);
471- op.Process();
472- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_BOOL)) {
473- KernelScatterElements<uint8_t, int32_t, uint64_t, half, REDU_ADD> op(pipe);
474- op.Init(x, indices, updates, y, userWS, &tilingData);
475- op.Process();
476- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_FP32)) {
477- KernelScatterElements<float, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
478- op.Init(x, indices, updates, y, userWS, &tilingData);
479- op.Process();
480- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_FP16)) {
481- KernelScatterElements<half, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
482- op.Init(x, indices, updates, y, userWS, &tilingData);
483- op.Process();
484- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_INT8)) {
485- KernelScatterElements<int8_t, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
486- op.Init(x, indices, updates, y, userWS, &tilingData);
487- op.Process();
488- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_INT32)) {
489- KernelScatterElements<int32_t, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
490- op.Init(x, indices, updates, y, userWS, &tilingData);
491- op.Process();
492- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_UINT8)) {
493- KernelScatterElements<uint8_t, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
494- op.Init(x, indices, updates, y, userWS, &tilingData);
495- op.Process();
496- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_INT16)) {
497- KernelScatterElements<int16_t, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
498- op.Init(x, indices, updates, y, userWS, &tilingData);
499- op.Process();
500- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_BF16)) {
501- KernelScatterElements<bfloat16_t, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
502- op.Init(x, indices, updates, y, userWS, &tilingData);
503- op.Process();
504- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_INT64)) {
505- KernelScatterElements<int64_t, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
506- op.Init(x, indices, updates, y, userWS, &tilingData);
507- op.Process();
508- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_NONE_B8)) {
509- if (tilingData.isDeterministic) {
510- KernelScatterElementsDeterm<int8_t, int64_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
511- op.Init(x, indices, updates, y);
512- op.Process();
513- } else {
514- KernelScatterElements<int8_t, int64_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
515 op.Init(x, indices, updates, y, userWS, &tilingData);631 op.Init(x, indices, updates, y, userWS, &tilingData);
516 op.Process();632 op.Process();
517- }633+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_ADD_UINT8
518- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_NONE_B16)) {634+ KernelScatterElements<uint8_t, int32_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
519- if (tilingData.isDeterministic) {
520- KernelScatterElementsDeterm<int16_t, int64_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
521- op.Init(x, indices, updates, y);
522- op.Process();
523- } else {
524- KernelScatterElements<int16_t, int64_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
525 op.Init(x, indices, updates, y, userWS, &tilingData);635 op.Init(x, indices, updates, y, userWS, &tilingData);
526 op.Process();636 op.Process();
527- }637+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_MUL_UINT8
528- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_NONE_B32)) {638+ KernelScatterElements<uint8_t, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
529- if (tilingData.isDeterministic) {
530- KernelScatterElementsDeterm<int32_t, int64_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);
531- op.Init(x, indices, updates, y);
532- op.Process();
533- } else {
534- KernelScatterElements<int32_t, int64_t, uint64_t, uint32_t, REDU_NONE> op(pipe);
535 op.Init(x, indices, updates, y, userWS, &tilingData);639 op.Init(x, indices, updates, y, userWS, &tilingData);
536 op.Process();640 op.Process();
537- }641+ #endif
538- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_NONE_B64)) {642+ #endif
539- if (tilingData.isDeterministic) {643+ 
540- KernelScatterElementsDeterm<int64_t, int64_t, uint64_t, REDU_NONE> op(&tilingData, &pipe);644+ #if ORIG_DTYPE_DATA == DT_UINT8 && ORIG_DTYPE_INDICES == DT_INT64
541- op.Init(x, indices, updates, y);645+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_UINT8);
542- op.Process();646+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_UINT8);
543- } else {647+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_UINT8);
544- KernelScatterElements<int64_t, int64_t, uint64_t, uint32_t, REDU_NONE> op(pipe);648+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_UINT8);
649+ 
650+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_ADD_UINT8
651+ KernelScatterElements<uint8_t, int64_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
545 op.Init(x, indices, updates, y, userWS, &tilingData);652 op.Init(x, indices, updates, y, userWS, &tilingData);
546 op.Process();653 op.Process();
547- }654+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_MUL_UINT8
548- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_FP32)) {655+ KernelScatterElements<uint8_t, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
549- if (tilingData.isDeterministic) {
550- KernelScatterElementsDeterm<float, int64_t, uint64_t, REDU_ADD> op(&tilingData, &pipe);
551- op.Init(x, indices, updates, y);
552- op.Process();
553- } else {
554- KernelScatterElements<float, int64_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
555 op.Init(x, indices, updates, y, userWS, &tilingData);656 op.Init(x, indices, updates, y, userWS, &tilingData);
556 op.Process();657 op.Process();
557- }658+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_ADD_UINT8
558- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_FP16)) {659+ KernelScatterElements<uint8_t, int64_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
559- if (tilingData.isDeterministic) {
560- KernelScatterElementsDeterm<half, int64_t, uint64_t, REDU_ADD> op(&tilingData, &pipe);
561- op.Init(x, indices, updates, y);
562- op.Process();
563- } else {
564- KernelScatterElements<half, int64_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
565 op.Init(x, indices, updates, y, userWS, &tilingData);660 op.Init(x, indices, updates, y, userWS, &tilingData);
566 op.Process();661 op.Process();
567- }662+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_MUL_UINT8
568- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_INT8)) {663+ KernelScatterElements<uint8_t, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
569- KernelScatterElements<int8_t, int64_t, uint64_t, int32_t, REDU_ADD> op(pipe);
570- op.Init(x, indices, updates, y, userWS, &tilingData);
571- op.Process();
572- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_INT32)) {
573- KernelScatterElements<int32_t, int64_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
574- op.Init(x, indices, updates, y, userWS, &tilingData);
575- op.Process();
576- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_UINT8)) {
577- KernelScatterElements<uint8_t, int64_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
578- op.Init(x, indices, updates, y, userWS, &tilingData);
579- op.Process();
580- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_INT16)) {
581- KernelScatterElements<int16_t, int64_t, uint64_t, int32_t, REDU_ADD> op(pipe);
582- op.Init(x, indices, updates, y, userWS, &tilingData);
583- op.Process();
584- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_BF16)) {
585- if (tilingData.isDeterministic) {
586- KernelScatterElementsDeterm<bfloat16_t, int64_t, uint64_t, REDU_ADD> op(&tilingData, &pipe);
587- op.Init(x, indices, updates, y);
588- op.Process();
589- } else {
590- KernelScatterElements<bfloat16_t, int64_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
591 op.Init(x, indices, updates, y, userWS, &tilingData);664 op.Init(x, indices, updates, y, userWS, &tilingData);
592 op.Process();665 op.Process();
593- }666+ #endif
594- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_INT64)) {667+ #endif
595- KernelScatterElements<int64_t, int64_t, uint64_t, uint32_t, REDU_ADD> op(pipe);668+ 
596- op.Init(x, indices, updates, y, userWS, &tilingData);669+ #if ORIG_DTYPE_DATA == DT_INT32 && ORIG_DTYPE_INDICES == DT_INT32
597- op.Process();670+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_INT32);
598- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_BOOL)) {671+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_INT32);
599- KernelScatterElements<uint8_t, int64_t, uint64_t, half, REDU_ADD> op(pipe);672+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_INT32);
600- op.Init(x, indices, updates, y, userWS, &tilingData);673+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_INT32);
601- op.Process();674+ 
602- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_FP32)) {675+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_ADD_INT32
603- KernelScatterElements<float, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);676+ KernelScatterElements<int32_t, int32_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
604- op.Init(x, indices, updates, y, userWS, &tilingData);677+ op.Init(x, indices, updates, y, userWS, &tilingData);
605- op.Process();678+ op.Process();
606- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_FP16)) {679+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_MUL_INT32
607- KernelScatterElements<half, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);680+ KernelScatterElements<int32_t, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
608- op.Init(x, indices, updates, y, userWS, &tilingData);681+ op.Init(x, indices, updates, y, userWS, &tilingData);
609- op.Process();682+ op.Process();
610- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_INT8)) {683+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_ADD_INT32
611- KernelScatterElements<int8_t, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);684+ KernelScatterElements<int32_t, int32_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
612- op.Init(x, indices, updates, y, userWS, &tilingData);685+ op.Init(x, indices, updates, y, userWS, &tilingData);
613- op.Process();686+ op.Process();
614- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_INT32)) {687+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_MUL_INT32
615- KernelScatterElements<int32_t, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);688+ KernelScatterElements<int32_t, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
616- op.Init(x, indices, updates, y, userWS, &tilingData);689+ op.Init(x, indices, updates, y, userWS, &tilingData);
617- op.Process();690+ op.Process();
618- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_UINT8)) {691+ #endif
619- KernelScatterElements<uint8_t, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);692+ #endif
620- op.Init(x, indices, updates, y, userWS, &tilingData);693+ 
621- op.Process();694+ #if ORIG_DTYPE_DATA == DT_INT32 && ORIG_DTYPE_INDICES == DT_INT64
622- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_INT16)) {695+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_INT32);
623- KernelScatterElements<int16_t, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);696+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_INT32);
624- op.Init(x, indices, updates, y, userWS, &tilingData);697+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_INT32);
625- op.Process();698+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_INT32);
626- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_BF16)) {699+ 
627- KernelScatterElements<bfloat16_t, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);700+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_ADD_INT32
628- op.Init(x, indices, updates, y, userWS, &tilingData);701+ KernelScatterElements<int32_t, int64_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
629- op.Process();702+ op.Init(x, indices, updates, y, userWS, &tilingData);
630- } else if (TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_INT64)) {703+ op.Process();
631- KernelScatterElements<int64_t, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);704+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_MUL_INT32
632- op.Init(x, indices, updates, y, userWS, &tilingData);705+ KernelScatterElements<int32_t, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
633- op.Process();706+ op.Init(x, indices, updates, y, userWS, &tilingData);
634- }707+ op.Process();
708+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_ADD_INT32
709+ KernelScatterElements<int32_t, int64_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
710+ op.Init(x, indices, updates, y, userWS, &tilingData);
711+ op.Process();
712+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_MUL_INT32
713+ KernelScatterElements<int32_t, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
714+ op.Init(x, indices, updates, y, userWS, &tilingData);
715+ op.Process();
716+ #endif
717+ #endif
718+ 
719+ #if ORIG_DTYPE_DATA == DT_INT16 && ORIG_DTYPE_INDICES == DT_INT32
720+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_INT16);
721+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_INT16);
722+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_INT16);
723+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_INT16);
724+ 
725+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_ADD_INT16
726+ KernelScatterElements<int16_t, int32_t, uint32_t, int32_t, REDU_ADD> op(pipe);
727+ op.Init(x, indices, updates, y, userWS, &tilingData);
728+ op.Process();
729+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_MUL_INT16
730+ KernelScatterElements<int16_t, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
731+ op.Init(x, indices, updates, y, userWS, &tilingData);
732+ op.Process();
733+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_ADD_INT16
734+ KernelScatterElements<int16_t, int32_t, uint64_t, int32_t, REDU_ADD> op(pipe);
735+ op.Init(x, indices, updates, y, userWS, &tilingData);
736+ op.Process();
737+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_MUL_INT16
738+ KernelScatterElements<int16_t, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
739+ op.Init(x, indices, updates, y, userWS, &tilingData);
740+ op.Process();
741+ #endif
742+ #endif
743+ 
744+ #if ORIG_DTYPE_DATA == DT_INT16 && ORIG_DTYPE_INDICES == DT_INT64
745+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_INT16);
746+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_INT16);
747+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_INT16);
748+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_INT16);
749+ 
750+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_ADD_INT16
751+ KernelScatterElements<int16_t, int64_t, uint32_t, int32_t, REDU_ADD> op(pipe);
752+ op.Init(x, indices, updates, y, userWS, &tilingData);
753+ op.Process();
754+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_MUL_INT16
755+ KernelScatterElements<int16_t, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
756+ op.Init(x, indices, updates, y, userWS, &tilingData);
757+ op.Process();
758+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_ADD_INT16
759+ KernelScatterElements<int16_t, int64_t, uint64_t, int32_t, REDU_ADD> op(pipe);
760+ op.Init(x, indices, updates, y, userWS, &tilingData);
761+ op.Process();
762+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_MUL_INT16
763+ KernelScatterElements<int16_t, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
764+ op.Init(x, indices, updates, y, userWS, &tilingData);
765+ op.Process();
766+ #endif
767+ #endif
768+ 
769+ #if ORIG_DTYPE_DATA == DT_INT64 && ORIG_DTYPE_INDICES == DT_INT32
770+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_INT64);
771+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_MUL_INT64);
772+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_INT64);
773+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_MUL_INT64);
774+ 
775+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_ADD_INT64
776+ KernelScatterElements<int64_t, int32_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
777+ op.Init(x, indices, updates, y, userWS, &tilingData);
778+ op.Process();
779+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_MUL_INT64
780+ KernelScatterElements<int64_t, int32_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
781+ op.Init(x, indices, updates, y, userWS, &tilingData);
782+ op.Process();
783+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_ADD_INT64
784+ KernelScatterElements<int64_t, int32_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
785+ op.Init(x, indices, updates, y, userWS, &tilingData);
786+ op.Process();
787+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_MUL_INT64
788+ KernelScatterElements<int64_t, int32_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
789+ op.Init(x, indices, updates, y, userWS, &tilingData);
790+ op.Process();
791+ #endif
792+ #endif
793+ 
794+ #if ORIG_DTYPE_DATA == DT_INT64 && ORIG_DTYPE_INDICES == DT_INT64
795+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_INT64);
796+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_MUL_INT64);
797+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_INT64);
798+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_MUL_INT64);
799+ 
800+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_ADD_INT64
801+ KernelScatterElements<int64_t, int64_t, uint32_t, uint32_t, REDU_ADD> op(pipe);
802+ op.Init(x, indices, updates, y, userWS, &tilingData);
803+ op.Process();
804+ #elif TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_MUL_INT64
805+ KernelScatterElements<int64_t, int64_t, uint32_t, uint32_t, REDU_MUL> op(pipe);
806+ op.Init(x, indices, updates, y, userWS, &tilingData);
807+ op.Process();
808+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_ADD_INT64
809+ KernelScatterElements<int64_t, int64_t, uint64_t, uint32_t, REDU_ADD> op(pipe);
810+ op.Init(x, indices, updates, y, userWS, &tilingData);
811+ op.Process();
812+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_MUL_INT64
813+ KernelScatterElements<int64_t, int64_t, uint64_t, uint32_t, REDU_MUL> op(pipe);
814+ op.Init(x, indices, updates, y, userWS, &tilingData);
815+ op.Process();
816+ #endif
817+ #endif
818+ 
819+ #if ORIG_DTYPE_DATA == DT_BOOL && ORIG_DTYPE_INDICES == DT_INT32
820+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX32_REDU_ADD_BOOL);
821+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX32_REDU_ADD_BOOL);
822+ 
823+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX32_REDU_ADD_BOOL
824+ KernelScatterElements<uint8_t, int32_t, uint32_t, half, REDU_ADD> op(pipe);
825+ op.Init(x, indices, updates, y, userWS, &tilingData);
826+ op.Process();
827+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX32_REDU_ADD_BOOL
828+ KernelScatterElements<uint8_t, int32_t, uint64_t, half, REDU_ADD> op(pipe);
829+ op.Init(x, indices, updates, y, userWS, &tilingData);
830+ op.Process();
831+ #endif
832+ #endif
833+
834+ #if ORIG_DTYPE_DATA == DT_BOOL && ORIG_DTYPE_INDICES == DT_INT64
835+ TILING_KEY_IS(SCAC_ELE_UINT32_IDX64_REDU_ADD_BOOL);
836+ TILING_KEY_IS(SCAC_ELE_UINT64_IDX64_REDU_ADD_BOOL);
837+ 
838+ #if TILING_KEY_VAR == SCAC_ELE_UINT32_IDX64_REDU_ADD_BOOL
839+ KernelScatterElements<uint8_t, int64_t, uint32_t, half, REDU_ADD> op(pipe);
840+ op.Init(x, indices, updates, y, userWS, &tilingData);
841+ op.Process();
842+ #elif TILING_KEY_VAR == SCAC_ELE_UINT64_IDX64_REDU_ADD_BOOL
843+ KernelScatterElements<uint8_t, int64_t, uint64_t, half, REDU_ADD> op(pipe);
844+ op.Init(x, indices, updates, y, userWS, &tilingData);
845+ op.Process();
846+ #endif
847+ #endif
635}848}
Mpooling/max_pool_v3/op_kernel/arch35/max_pool_v3_common.h+2-1
@@ -15,7 +15,8 @@
15#ifndef MAX_POOL_V3_COMMON_H_15#ifndef MAX_POOL_V3_COMMON_H_
16#define MAX_POOL_V3_COMMON_H_16#define MAX_POOL_V3_COMMON_H_
17 17 
18-#include "kernel_operator.h"18+#include "kernel_vec_intf.h"
19+#include "kernel_micro_intf.h"
19#include "op_kernel/platform_util.h"20#include "op_kernel/platform_util.h"
20#include "op_kernel/math_util.h"21#include "op_kernel/math_util.h"
21 22 
Mpooling/max_pool_v3/op_kernel/max_pool_v3_apt.cpp+23-8
@@ -13,7 +13,6 @@
13 */13 */
14 14 
15#include <cstdint>15#include <cstdint>
16-#include "kernel_operator.h"
17#include "arch35/max_pool_v3_big_kernel.h"16#include "arch35/max_pool_v3_big_kernel.h"
18#include "arch35/max_pool_v3_small_kernel.h"17#include "arch35/max_pool_v3_small_kernel.h"
19#include "arch35/max_pool_v3_small_kernel_pad.h"18#include "arch35/max_pool_v3_small_kernel_pad.h"
@@ -32,35 +31,51 @@ extern "C" __global__ __aicore__ void max_pool_v3(GM_ADDR x, GM_ADDR y, GM_ADDR
32{31{
33 AscendC::TPipe pipeBase;32 AscendC::TPipe pipeBase;
34 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);33 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
35- if (TILING_KEY_IS(BIG_KERNEL_FORMAT_NCHW)) {34+ TILING_KEY_IS(BIG_KERNEL_FORMAT_NCHW);
35+ TILING_KEY_IS(SMALL_KERNEL_NO_PADDING_FORMAT_NCHW);
36+ TILING_KEY_IS(SMALL_KERNEL_PADDING_FORMAT_NCHW);
37+ TILING_KEY_IS(BIG_KERNEL_FORMAT_NHWC);
38+ TILING_KEY_IS(SMALL_KERNEL_NO_PADDING_FORMAT_NHWC);
39+ TILING_KEY_IS(SMALL_KERNEL_PADDING_FORMAT_NHWC);
40+ #if TILING_KEY_X == BIG_KERNEL_FORMAT_NCHW
36 GET_TILING_DATA_WITH_STRUCT(MaxPoolV3BigKernelTilingData, tilingData, tiling);41 GET_TILING_DATA_WITH_STRUCT(MaxPoolV3BigKernelTilingData, tilingData, tiling);
37 MaxPoolV3::MaxPoolV3BigKernel<DTYPE_X> op(&pipeBase, &tilingData);42 MaxPoolV3::MaxPoolV3BigKernel<DTYPE_X> op(&pipeBase, &tilingData);
38 op.Init(x, y);43 op.Init(x, y);
39 op.Process();44 op.Process();
40- } else if (TILING_KEY_IS(SMALL_KERNEL_NO_PADDING_FORMAT_NCHW)) {45+ #endif
46+ 
47+ #if TILING_KEY_X == SMALL_KERNEL_NO_PADDING_FORMAT_NCHW
41 GET_TILING_DATA_WITH_STRUCT(MaxPoolV3SmallKernelTilingData, tilingData, tiling);48 GET_TILING_DATA_WITH_STRUCT(MaxPoolV3SmallKernelTilingData, tilingData, tiling);
42 MaxPoolV3::MaxPoolV3SmallKernel<DTYPE_X> op(&pipeBase, &tilingData);49 MaxPoolV3::MaxPoolV3SmallKernel<DTYPE_X> op(&pipeBase, &tilingData);
43 op.Init(x, y);50 op.Init(x, y);
44 op.Process();51 op.Process();
45- } else if (TILING_KEY_IS(SMALL_KERNEL_PADDING_FORMAT_NCHW)) {52+ #endif
53+ 
54+ #if TILING_KEY_X == SMALL_KERNEL_PADDING_FORMAT_NCHW
46 GET_TILING_DATA_WITH_STRUCT(MaxPoolV3SmallKernelTilingData, tilingData, tiling);55 GET_TILING_DATA_WITH_STRUCT(MaxPoolV3SmallKernelTilingData, tilingData, tiling);
47 MaxPoolV3::MaxPoolV3SmallPadKernel<DTYPE_X> op(&pipeBase, &tilingData);56 MaxPoolV3::MaxPoolV3SmallPadKernel<DTYPE_X> op(&pipeBase, &tilingData);
48 op.Init(x, y);57 op.Init(x, y);
49 op.Process();58 op.Process();
50- } else if (TILING_KEY_IS(BIG_KERNEL_FORMAT_NHWC)) {59+ #endif
60+ 
61+ #if TILING_KEY_X == BIG_KERNEL_FORMAT_NHWC
51 GET_TILING_DATA_WITH_STRUCT(MaxPoolV3NHWCBigKernelTilingData, tilingData, tiling);62 GET_TILING_DATA_WITH_STRUCT(MaxPoolV3NHWCBigKernelTilingData, tilingData, tiling);
52 MaxPoolV3::MaxPoolV3NHWCBigKernel<DTYPE_X> op(&pipeBase, &tilingData);63 MaxPoolV3::MaxPoolV3NHWCBigKernel<DTYPE_X> op(&pipeBase, &tilingData);
53 op.Init(x, y);64 op.Init(x, y);
54 op.Process();65 op.Process();
55- } else if (TILING_KEY_IS(SMALL_KERNEL_NO_PADDING_FORMAT_NHWC)) {66+ #endif
67+ 
68+ #if TILING_KEY_X == SMALL_KERNEL_NO_PADDING_FORMAT_NHWC
56 GET_TILING_DATA_WITH_STRUCT(MaxPoolV3NHWCSmallKernelTilingData, tilingData, tiling);69 GET_TILING_DATA_WITH_STRUCT(MaxPoolV3NHWCSmallKernelTilingData, tilingData, tiling);
57 MaxPoolV3::MaxPoolV3NHWCSmallKernel<DTYPE_X> op(&pipeBase, &tilingData);70 MaxPoolV3::MaxPoolV3NHWCSmallKernel<DTYPE_X> op(&pipeBase, &tilingData);
58 op.Init(x, y);71 op.Init(x, y);
59 op.Process();72 op.Process();
60- } else if (TILING_KEY_IS(SMALL_KERNEL_PADDING_FORMAT_NHWC)) {73+ #endif
74+ 
75+ #if TILING_KEY_X == SMALL_KERNEL_PADDING_FORMAT_NHWC
61 GET_TILING_DATA_WITH_STRUCT(MaxPoolV3NHWCSmallKernelTilingData, tilingData, tiling);76 GET_TILING_DATA_WITH_STRUCT(MaxPoolV3NHWCSmallKernelTilingData, tilingData, tiling);
62 MaxPoolV3::MaxPoolV3NHWCSmallPadKernel<DTYPE_X> op(&pipeBase, &tilingData);77 MaxPoolV3::MaxPoolV3NHWCSmallPadKernel<DTYPE_X> op(&pipeBase, &tilingData);
63 op.Init(x, y);78 op.Init(x, y);
64 op.Process();79 op.Process();
65- }80+ #endif
66}81}