Star1303
1419
代码
代码
Issues110
Pull Requests234
Actions
项目讨论
Wiki
分析
Star1303
1419
已合并
Add StridedSlice operator implementation (issue #248) #395
liujun创建于 2025年12月12日
Add StridedSlice operator implementation (issue #248) #395
已合并
liujun创建于 2025年12月12日
从kbryantttt/ops-mathstridedslicefianl合入到cann/ops-mathmaster
共 12 个文件变更+1674-0
Aexperimental/conversion/strided_slice/CMakeLists.txt+12-0
@@ -0,0 +1,12 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify it.
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+add_all_modules_sources(OPTYPE strided_slice ACLNNTYPE aclnn)
Aexperimental/conversion/strided_slice/examples/test_aclnn_strided_slice.cpp+164-0
@@ -0,0 +1,164 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Liu Jun <@kbryantttt>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+#include <iostream>
21+#include <vector>
22+#include "acl/acl.h"
23+#include "aclnn_strided_slice.h"
24+#define CHECK_RET(cond, return_expr) \
25+ do { \
26+ if (!(cond)) { \
27+ return_expr; \
28+ } \
29+ } while (0)
30+ 
31+#define LOG_PRINT(message, ...) \
32+ do { \
33+ printf(message, ##__VA_ARGS__); \
34+ } while (0)
35+ 
36+int64_t GetShapeSize(const std::vector<int64_t>& shape)
37+{
38+ int64_t shapeSize = 1;
39+ for (auto i : shape) {
40+ shapeSize *= i;
41+ }
42+ return shapeSize;
43+}
44+void PrintOutResult(std::vector<int64_t>& shape, void** deviceAddr)
45+{
46+ auto size = GetShapeSize(shape);
47+ std::vector<int32_t> resultData(size, 0);
48+ auto ret = aclrtMemcpy(
49+ resultData.data(), resultData.size() * sizeof(resultData[0]), *deviceAddr, size * sizeof(resultData[0]),
50+ ACL_MEMCPY_DEVICE_TO_HOST);
51+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return);
52+ for (int64_t i = 0; i < 120; i++) {
53+ LOG_PRINT("mean result[%ld] is: %d\n", i, resultData[i]);
54+ }
55+}
56+int Init(int32_t deviceId, aclrtStream* stream)
57+{
58+ // 固定写法，初始化
59+ auto ret = aclInit(nullptr);
60+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
61+ ret = aclrtSetDevice(deviceId);
62+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
63+ ret = aclrtCreateStream(stream);
64+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
65+ return 0;
66+}
67+template <typename T>
68+int CreateAclTensor(
69+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
70+ aclTensor** tensor)
71+{
72+ auto size = GetShapeSize(shape) * sizeof(T);
73+ // 2. 申请device侧内存
74+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
75+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
76+ // 3. 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
77+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
78+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
79+ 
80+ // 计算连续tensor的strides
81+ std::vector<int64_t> strides(shape.size(), 1);
82+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
83+ strides[i] = shape[i + 1] * strides[i + 1];
84+ }
85+ 
86+ // 调用aclCreateTensor接口创建aclTensor
87+ *tensor = aclCreateTensor(
88+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
89+ *deviceAddr);
90+ return 0;
91+}
92+int main()
93+{
94+ // 1. 调用acl进行device/stream初始化
95+ int32_t deviceId = 0;
96+ aclrtStream stream;
97+ auto ret = Init(deviceId, &stream);
98+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
99+ 
100+ // 2. 构造输入与输出，需要根据API的接口自定义构造
101+ aclTensor* selfX = nullptr;
102+ void* selfXDeviceAddr = nullptr;
103+ std::vector<int64_t> selfXShape = {40};
104+ std::vector<int32_t> selfXHostData={
105+ 1,2,3,4,5,6,7,8,9,10,11,12,13,14,15,16,17,18,19,20,
106+ 21,22,23,24,25,26,27,28,29,30,31,32,33,34,35,36,37,38,39,40,
107+ };
108+ ret = CreateAclTensor(selfXHostData, selfXShape, &selfXDeviceAddr, aclDataType::ACL_INT32, &selfX);
109+ CHECK_RET(ret == ACL_SUCCESS, return ret);
110+ aclTensor* out = nullptr;
111+ void* outDeviceAddr = nullptr;
112+ std::vector<int64_t> outShape = {13};
113+ std::vector<int32_t> outHostData(13, 0);
114+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_INT32, &out);
115+ CHECK_RET(ret == ACL_SUCCESS, return ret);
116+ 
117+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
118+ uint64_t workspaceSize = 0;
119+ aclOpExecutor* executor;
120+ int64_t start1 =0;
121+ int64_t start2 = 2;
122+ int64_t end1 = 1;
123+ int64_t end2 = 39;
124+ int64_t stride1 = 1;
125+ int64_t stride2 = 3;
126+ // 4. 调用aclnnStridedSlice第一段接口
127+ ret = aclnnStridedSliceGetWorkspaceSize(selfX,start1,start2,end1,end2,stride1,stride2,out, &workspaceSize, &executor);
128+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnStridedSliceGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
129+ 
130+ // 根据第一段接口计算出的workspaceSize申请device内存
131+ void* workspaceAddr = nullptr;
132+ if (workspaceSize > static_cast<uint64_t>(0)) {
133+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
134+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
135+ }
136+ 
137+ // 5. 调用aclnnStridedSlice第二段接口
138+ ret = aclnnStridedSlice(workspaceAddr, workspaceSize, executor, stream);
139+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnStridedSlice failed. ERROR: %d\n", ret); return ret);
140+ 
141+ // 6. （固定写法）同步等待任务执行结束
142+ ret = aclrtSynchronizeStream(stream);
143+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
144+ 
145+ // 7. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
146+ PrintOutResult(outShape, &outDeviceAddr);
147+ 
148+ // 8. 释放aclTensor，需要根据具体API的接口定义修改
149+ aclDestroyTensor(selfX);
150+ aclDestroyTensor(out);
151+ 
152+ // 9. 释放device资源
153+ aclrtFree(selfXDeviceAddr);
154+ aclrtFree(outDeviceAddr);
155+ if (workspaceSize > static_cast<uint64_t>(0)) {
156+ aclrtFree(workspaceAddr);
157+ }
158+ aclrtDestroyStream(stream);
159+ aclrtResetDevice(deviceId);
160+ 
161+ // 10. acl去初始化
162+ aclFinalize();
163+ return 0;
164+}
Aexperimental/conversion/strided_slice/op_host/config/ascend910b/strided_slice_binary.json+737-0
@@ -0,0 +1,737 @@
1+{
2+ "op_type": "StridedSlice",
3+ "op_list": [
4+ {
5+ "bin_filename": "StridedSlice_a1532827238e1555db7b997c7bce2912",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "float32",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "format_match_mode": "FormatAgnostic"
17+ }
18+ ],
19+ "attrs": [
20+ {
21+ "name": "start1",
22+ "dtype": "uint32",
23+ "value": null
24+ },
25+ {
26+ "name": "start2",
27+ "dtype": "uint32",
28+ "value": null
29+ },
30+ {
31+ "name": "end1",
32+ "dtype": "uint32",
33+ "value": null
34+ },
35+ {
36+ "name": "end2",
37+ "dtype": "uint32",
38+ "value": null
39+ },
40+ {
41+ "name": "stride1",
42+ "dtype": "uint32",
43+ "value": null
44+ },
45+ {
46+ "name": "stride2",
47+ "dtype": "uint32",
48+ "value": null
49+ }
50+ ],
51+ "outputs": [
52+ {
53+ "name": "z",
54+ "index": 0,
55+ "dtype": "float32",
56+ "format": "ND",
57+ "paramType": "required",
58+ "shape": [
59+ -2
60+ ],
61+ "format_match_mode": "FormatAgnostic"
62+ }
63+ ]
64+ },
65+ {
66+ "bin_filename": "StridedSlice_a1532827238e1555db7b997c7bce2911",
67+ "inputs": [
68+ {
69+ "name": "x",
70+ "index": 0,
71+ "dtype": "float16",
72+ "format": "ND",
73+ "paramType": "required",
74+ "shape": [
75+ -2
76+ ],
77+ "format_match_mode": "FormatAgnostic"
78+ }
79+ ],
80+ "attrs": [
81+ {
82+ "name": "start1",
83+ "dtype": "uint32",
84+ "value": null
85+ },
86+ {
87+ "name": "start2",
88+ "dtype": "uint32",
89+ "value": null
90+ },
91+ {
92+ "name": "end1",
93+ "dtype": "uint32",
94+ "value": null
95+ },
96+ {
97+ "name": "end2",
98+ "dtype": "uint32",
99+ "value": null
100+ },
101+ {
102+ "name": "stride1",
103+ "dtype": "uint32",
104+ "value": null
105+ },
106+ {
107+ "name": "stride2",
108+ "dtype": "uint32",
109+ "value": null
110+ }
111+ ],
112+ "outputs": [
113+ {
114+ "name": "z",
115+ "index": 0,
116+ "dtype": "float16",
117+ "format": "ND",
118+ "paramType": "required",
119+ "shape": [
120+ -2
121+ ],
122+ "format_match_mode": "FormatAgnostic"
123+ }
124+ ]
125+ },
126+ {
127+ "bin_filename": "StridedSlice_a1532827238e1555db7b997c7bce2920",
128+ "inputs": [
129+ {
130+ "name": "x",
131+ "index": 0,
132+ "dtype": "int32",
133+ "format": "ND",
134+ "paramType": "required",
135+ "shape": [
136+ -2
137+ ],
138+ "format_match_mode": "FormatAgnostic"
139+ }
140+ ],
141+ "attrs": [
142+ {
143+ "name": "start1",
144+ "dtype": "uint32",
145+ "value": null
146+ },
147+ {
148+ "name": "start2",
149+ "dtype": "uint32",
150+ "value": null
151+ },
152+ {
153+ "name": "end1",
154+ "dtype": "uint32",
155+ "value": null
156+ },
157+ {
158+ "name": "end2",
159+ "dtype": "uint32",
160+ "value": null
161+ },
162+ {
163+ "name": "stride1",
164+ "dtype": "uint32",
165+ "value": null
166+ },
167+ {
168+ "name": "stride2",
169+ "dtype": "uint32",
170+ "value": null
171+ }
172+ ],
173+ "outputs": [
174+ {
175+ "name": "z",
176+ "index": 0,
177+ "dtype": "int32",
178+ "format": "ND",
179+ "paramType": "required",
180+ "shape": [
181+ -2
182+ ],
183+ "format_match_mode": "FormatAgnostic"
184+ }
185+ ]
186+ },
187+ {
188+ "bin_filename": "StridedSlice_a1532827238e1555db7b997c7bce2929",
189+ "inputs": [
190+ {
191+ "name": "x",
192+ "index": 0,
193+ "dtype": "int16",
194+ "format": "ND",
195+ "paramType": "required",
196+ "shape": [
197+ -2
198+ ],
199+ "format_match_mode": "FormatAgnostic"
200+ }
201+ ],
202+ "attrs": [
203+ {
204+ "name": "start1",
205+ "dtype": "uint32",
206+ "value": null
207+ },
208+ {
209+ "name": "start2",
210+ "dtype": "uint32",
211+ "value": null
212+ },
213+ {
214+ "name": "end1",
215+ "dtype": "uint32",
216+ "value": null
217+ },
218+ {
219+ "name": "end2",
220+ "dtype": "uint32",
221+ "value": null
222+ },
223+ {
224+ "name": "stride1",
225+ "dtype": "uint32",
226+ "value": null
227+ },
228+ {
229+ "name": "stride2",
230+ "dtype": "uint32",
231+ "value": null
232+ }
233+ ],
234+ "outputs": [
235+ {
236+ "name": "z",
237+ "index": 0,
238+ "dtype": "int16",
239+ "format": "ND",
240+ "paramType": "required",
241+ "shape": [
242+ -2
243+ ],
244+ "format_match_mode": "FormatAgnostic"
245+ }
246+ ]
247+ },
248+ {
249+ "bin_filename": "StridedSlice_a1532827238e1555db7b997c7bce2928",
250+ "inputs": [
251+ {
252+ "name": "x",
253+ "index": 0,
254+ "dtype": "int8",
255+ "format": "ND",
256+ "paramType": "required",
257+ "shape": [
258+ -2
259+ ],
260+ "format_match_mode": "FormatAgnostic"
261+ }
262+ ],
263+ "attrs": [
264+ {
265+ "name": "start1",
266+ "dtype": "uint32",
267+ "value": null
268+ },
269+ {
270+ "name": "start2",
271+ "dtype": "uint32",
272+ "value": null
273+ },
274+ {
275+ "name": "end1",
276+ "dtype": "uint32",
277+ "value": null
278+ },
279+ {
280+ "name": "end2",
281+ "dtype": "uint32",
282+ "value": null
283+ },
284+ {
285+ "name": "stride1",
286+ "dtype": "uint32",
287+ "value": null
288+ },
289+ {
290+ "name": "stride2",
291+ "dtype": "uint32",
292+ "value": null
293+ }
294+ ],
295+ "outputs": [
296+ {
297+ "name": "z",
298+ "index": 0,
299+ "dtype": "int8",
300+ "format": "ND",
301+ "paramType": "required",
302+ "shape": [
303+ -2
304+ ],
305+ "format_match_mode": "FormatAgnostic"
306+ }
307+ ]
308+ },
309+ {
310+ "bin_filename": "StridedSlice_a1532827238e1555db7b997c7bce2927",
311+ "inputs": [
312+ {
313+ "name": "x",
314+ "index": 0,
315+ "dtype": "bool",
316+ "format": "ND",
317+ "paramType": "required",
318+ "shape": [
319+ -2
320+ ],
321+ "format_match_mode": "FormatAgnostic"
322+ }
323+ ],
324+ "attrs": [
325+ {
326+ "name": "start1",
327+ "dtype": "uint32",
328+ "value": null
329+ },
330+ {
331+ "name": "start2",
332+ "dtype": "uint32",
333+ "value": null
334+ },
335+ {
336+ "name": "end1",
337+ "dtype": "uint32",
338+ "value": null
339+ },
340+ {
341+ "name": "end2",
342+ "dtype": "uint32",
343+ "value": null
344+ },
345+ {
346+ "name": "stride1",
347+ "dtype": "uint32",
348+ "value": null
349+ },
350+ {
351+ "name": "stride2",
352+ "dtype": "uint32",
353+ "value": null
354+ }
355+ ],
356+ "outputs": [
357+ {
358+ "name": "z",
359+ "index": 0,
360+ "dtype": "bool",
361+ "format": "ND",
362+ "paramType": "required",
363+ "shape": [
364+ -2
365+ ],
366+ "format_match_mode": "FormatAgnostic"
367+ }
368+ ]
369+ },
370+ {
371+ "bin_filename": "StridedSlice_a1532827238e1555db7b997c7bce2926",
372+ "inputs": [
373+ {
374+ "name": "x",
375+ "index": 0,
376+ "dtype": "uint8",
377+ "format": "ND",
378+ "paramType": "required",
379+ "shape": [
380+ -2
381+ ],
382+ "format_match_mode": "FormatAgnostic"
383+ }
384+ ],
385+ "attrs": [
386+ {
387+ "name": "start1",
388+ "dtype": "uint32",
389+ "value": null
390+ },
391+ {
392+ "name": "start2",
393+ "dtype": "uint32",
394+ "value": null
395+ },
396+ {
397+ "name": "end1",
398+ "dtype": "uint32",
399+ "value": null
400+ },
401+ {
402+ "name": "end2",
403+ "dtype": "uint32",
404+ "value": null
405+ },
406+ {
407+ "name": "stride1",
408+ "dtype": "uint32",
409+ "value": null
410+ },
411+ {
412+ "name": "stride2",
413+ "dtype": "uint32",
414+ "value": null
415+ }
416+ ],
417+ "outputs": [
418+ {
419+ "name": "z",
420+ "index": 0,
421+ "dtype": "uint8",
422+ "format": "ND",
423+ "paramType": "required",
424+ "shape": [
425+ -2
426+ ],
427+ "format_match_mode": "FormatAgnostic"
428+ }
429+ ]
430+ },
431+ {
432+ "bin_filename": "StridedSlice_a1532827238e1555db7b997c7bce2925",
433+ "inputs": [
434+ {
435+ "name": "x",
436+ "index": 0,
437+ "dtype": "uint16",
438+ "format": "ND",
439+ "paramType": "required",
440+ "shape": [
441+ -2
442+ ],
443+ "format_match_mode": "FormatAgnostic"
444+ }
445+ ],
446+ "attrs": [
447+ {
448+ "name": "start1",
449+ "dtype": "uint32",
450+ "value": null
451+ },
452+ {
453+ "name": "start2",
454+ "dtype": "uint32",
455+ "value": null
456+ },
457+ {
458+ "name": "end1",
459+ "dtype": "uint32",
460+ "value": null
461+ },
462+ {
463+ "name": "end2",
464+ "dtype": "uint32",
465+ "value": null
466+ },
467+ {
468+ "name": "stride1",
469+ "dtype": "uint32",
470+ "value": null
471+ },
472+ {
473+ "name": "stride2",
474+ "dtype": "uint32",
475+ "value": null
476+ }
477+ ],
478+ "outputs": [
479+ {
480+ "name": "z",
481+ "index": 0,
482+ "dtype": "uint16",
483+ "format": "ND",
484+ "paramType": "required",
485+ "shape": [
486+ -2
487+ ],
488+ "format_match_mode": "FormatAgnostic"
489+ }
490+ ]
491+ },
492+ {
493+ "bin_filename": "StridedSlice_a1532827238e1555db7b997c7bce2924",
494+ "inputs": [
495+ {
496+ "name": "x",
497+ "index": 0,
498+ "dtype": "uint32",
499+ "format": "ND",
500+ "paramType": "required",
501+ "shape": [
502+ -2
503+ ],
504+ "format_match_mode": "FormatAgnostic"
505+ }
506+ ],
507+ "attrs": [
508+ {
509+ "name": "start1",
510+ "dtype": "uint32",
511+ "value": null
512+ },
513+ {
514+ "name": "start2",
515+ "dtype": "uint32",
516+ "value": null
517+ },
518+ {
519+ "name": "end1",
520+ "dtype": "uint32",
521+ "value": null
522+ },
523+ {
524+ "name": "end2",
525+ "dtype": "uint32",
526+ "value": null
527+ },
528+ {
529+ "name": "stride1",
530+ "dtype": "uint32",
531+ "value": null
532+ },
533+ {
534+ "name": "stride2",
535+ "dtype": "uint32",
536+ "value": null
537+ }
538+ ],
539+ "outputs": [
540+ {
541+ "name": "z",
542+ "index": 0,
543+ "dtype": "uint32",
544+ "format": "ND",
545+ "paramType": "required",
546+ "shape": [
547+ -2
548+ ],
549+ "format_match_mode": "FormatAgnostic"
550+ }
551+ ]
552+ },
553+ {
554+ "bin_filename": "StridedSlice_a1532827238e1555db7b997c7bce2923",
555+ "inputs": [
556+ {
557+ "name": "x",
558+ "index": 0,
559+ "dtype": "bfloat16",
560+ "format": "ND",
561+ "paramType": "required",
562+ "shape": [
563+ -2
564+ ],
565+ "format_match_mode": "FormatAgnostic"
566+ }
567+ ],
568+ "attrs": [
569+ {
570+ "name": "start1",
571+ "dtype": "uint32",
572+ "value": null
573+ },
574+ {
575+ "name": "start2",
576+ "dtype": "uint32",
577+ "value": null
578+ },
579+ {
580+ "name": "end1",
581+ "dtype": "uint32",
582+ "value": null
583+ },
584+ {
585+ "name": "end2",
586+ "dtype": "uint32",
587+ "value": null
588+ },
589+ {
590+ "name": "stride1",
591+ "dtype": "uint32",
592+ "value": null
593+ },
594+ {
595+ "name": "stride2",
596+ "dtype": "uint32",
597+ "value": null
598+ }
599+ ],
600+ "outputs": [
601+ {
602+ "name": "z",
603+ "index": 0,
604+ "dtype": "bfloat16",
605+ "format": "ND",
606+ "paramType": "required",
607+ "shape": [
608+ -2
609+ ],
610+ "format_match_mode": "FormatAgnostic"
611+ }
612+ ]
613+ },
614+ {
615+ "bin_filename": "StridedSlice_a1532827238e1555db7b997c7bce2922",
616+ "inputs": [
617+ {
618+ "name": "x",
619+ "index": 0,
620+ "dtype": "uint64",
621+ "format": "ND",
622+ "paramType": "required",
623+ "shape": [
624+ -2
625+ ],
626+ "format_match_mode": "FormatAgnostic"
627+ }
628+ ],
629+ "attrs": [
630+ {
631+ "name": "start1",
632+ "dtype": "uint32",
633+ "value": null
634+ },
635+ {
636+ "name": "start2",
637+ "dtype": "uint32",
638+ "value": null
639+ },
640+ {
641+ "name": "end1",
642+ "dtype": "uint32",
643+ "value": null
644+ },
645+ {
646+ "name": "end2",
647+ "dtype": "uint32",
648+ "value": null
649+ },
650+ {
651+ "name": "stride1",
652+ "dtype": "uint32",
653+ "value": null
654+ },
655+ {
656+ "name": "stride2",
657+ "dtype": "uint32",
658+ "value": null
659+ }
660+ ],
661+ "outputs": [
662+ {
663+ "name": "z",
664+ "index": 0,
665+ "dtype": "uint64",
666+ "format": "ND",
667+ "paramType": "required",
668+ "shape": [
669+ -2
670+ ],
671+ "format_match_mode": "FormatAgnostic"
672+ }
673+ ]
674+ },
675+ {
676+ "bin_filename": "StridedSlice_a1532827238e1555db7b997c7bce2921",
677+ "inputs": [
678+ {
679+ "name": "x",
680+ "index": 0,
681+ "dtype": "int64",
682+ "format": "ND",
683+ "paramType": "required",
684+ "shape": [
685+ -2
686+ ],
687+ "format_match_mode": "FormatAgnostic"
688+ }
689+ ],
690+ "attrs": [
691+ {
692+ "name": "start1",
693+ "dtype": "uint32",
694+ "value": null
695+ },
696+ {
697+ "name": "start2",
698+ "dtype": "uint32",
699+ "value": null
700+ },
701+ {
702+ "name": "end1",
703+ "dtype": "uint32",
704+ "value": null
705+ },
706+ {
707+ "name": "end2",
708+ "dtype": "uint32",
709+ "value": null
710+ },
711+ {
712+ "name": "stride1",
713+ "dtype": "uint32",
714+ "value": null
715+ },
716+ {
717+ "name": "stride2",
718+ "dtype": "uint32",
719+ "value": null
720+ }
721+ ],
722+ "outputs": [
723+ {
724+ "name": "z",
725+ "index": 0,
726+ "dtype": "int64",
727+ "format": "ND",
728+ "paramType": "required",
729+ "shape": [
730+ -2
731+ ],
732+ "format_match_mode": "FormatAgnostic"
733+ }
734+ ]
735+ }
736+ ]
737+}
Aexperimental/conversion/strided_slice/op_host/strided_slice_def.cpp+104-0
@@ -0,0 +1,104 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Liu Jun <@kbryantttt>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+/*!
21+ * \file strided_slice.cpp
22+ * \brief
23+ */
24+#include "register/op_def_registry.h"
25+namespace ops {
26+class StridedSlice : public OpDef {
27+public:
28+ explicit StridedSlice(const char* name) : OpDef(name)
29+ {
30+ // 输入参数说明
31+ this->Input("x") // 输入x1定义
32+ .DataType({ ge::DT_FLOAT, ge::DT_INT32, ge::DT_UINT32, ge::DT_INT64,
33+ ge::DT_UINT64,ge::DT_FLOAT16, ge::DT_INT16, ge::DT_UINT16,
34+ ge::DT_BF16,ge::DT_INT8, ge::DT_UINT8, ge::DT_BOOL,
35+ ge::DT_FLOAT, ge::DT_INT32, ge::DT_UINT32, ge::DT_INT64,
36+ ge::DT_UINT64,ge::DT_FLOAT16, ge::DT_INT16, ge::DT_UINT16,
37+ ge::DT_BF16,ge::DT_INT8, ge::DT_UINT8, ge::DT_BOOL})
38+ .Format({ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
39+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
40+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
41+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
42+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
43+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND})
44+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
45+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
46+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
47+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
48+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
49+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND})
50+ .AutoContiguous(); // 内存自动连续化
51+ /* ...此处补充其他输入输出参数说明 */
52+ this->Attr("start1")
53+ .AttrType(OPTIONAL)
54+ .Int();
55+ this->Attr("start2")
56+ .AttrType(OPTIONAL)
57+ .Int();
58+ this->Attr("end1")
59+ .AttrType(OPTIONAL)
60+ .Int();
61+ this->Attr("end2")
62+ .AttrType(OPTIONAL)
63+ .Int();
64+ this->Attr("stride1")
65+ .AttrType(OPTIONAL)
66+ .Int();
67+ this->Attr("stride2")
R
RRuiWang_2025年12月18日

建议改为数组ListInt

likedislike
liujun
liujun
2025年12月19日 评论：
R
RuiWang_
2025年12月22日 评论：
liujun
liujun
2025年12月22日 评论：
68+ .AttrType(OPTIONAL)
69+ .Int();
70+ // 输出参数说明
71+ this->Output("z") // 输出y定义
72+ .ParamType(REQUIRED)
73+ .DataType({ ge::DT_FLOAT, ge::DT_INT32, ge::DT_UINT32, ge::DT_INT64,
74+ ge::DT_UINT64,ge::DT_FLOAT16, ge::DT_INT16, ge::DT_UINT16,
75+ ge::DT_BF16,ge::DT_INT8, ge::DT_UINT8, ge::DT_BOOL,
76+ ge::DT_FLOAT, ge::DT_INT32, ge::DT_UINT32, ge::DT_INT64,
77+ ge::DT_UINT64,ge::DT_FLOAT16, ge::DT_INT16, ge::DT_UINT16,
78+ ge::DT_BF16,ge::DT_INT8, ge::DT_UINT8, ge::DT_BOOL})
79+ .Format({ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
80+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
81+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
82+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
83+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
84+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND})
85+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
86+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
87+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
88+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
89+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND,
90+ ge::FORMAT_ND, ge::FORMAT_ND,ge::FORMAT_ND, ge::FORMAT_ND})
91+ .AutoContiguous();
92+ OpAICoreConfig aicoreConfig;
93+ aicoreConfig.DynamicCompileStaticFlag(true)
94+ .DynamicFormatFlag(false)
95+ .DynamicRankSupportFlag(true)
96+ .DynamicShapeSupportFlag(true)
97+ .NeedCheckSupportFlag(false)
98+ .PrecisionReduceFlag(true)
99+ .ExtendCfgInfo("opFile.value", "strided_slice");
100+ this->AICore().AddConfig("ascend910b", aicoreConfig);
101+ }
102+};
103+OP_ADD(StridedSlice); // 添加算子信息库
104+} // namespace ops
Aexperimental/conversion/strided_slice/op_host/strided_slice_infershape.cpp+83-0
@@ -0,0 +1,83 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Liu Jun <@kbryantttt>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+/*!
21+ * \file strided_slice_infer.cpp
22+ * \brief
23+ */
24+#include "register/op_impl_registry.h"
25+#include "log/log.h"
26+using namespace ge;
27+namespace ops {
28+static constexpr int64_t IDX_0 = 0;
29+static ge::graphStatus InferShapeStridedSlice(gert::InferShapeContext* context)
30+{
31+ const gert::Shape* xShape = context->GetInputShape(IDX_0);
32+ auto xShapeSize = xShape->GetDimNum();
33+ uint32_t start1 = 0;
34+ uint32_t start2 = 0;
35+ uint32_t end1 = 0;
36+ uint32_t end2 = 0;
37+ uint32_t stride1 = 0;
38+ uint32_t stride2 = 0;
39+ auto attrs = context->GetAttrs();
40+ if (attrs) {
41+ const int64_t* start1Ptr = attrs->GetInt(0);
42+ if (start1Ptr) {start1 = static_cast<uint32_t>(*start1Ptr);}
43+ if (xShapeSize == 1){
44+ start1 = static_cast<uint32_t>(0);
45+ }
46+ const int64_t* start2Ptr = attrs->GetInt(1);
47+ if (start2Ptr) {start2 = static_cast<uint32_t>(*start2Ptr);}
48+ const int64_t* end1Ptr = attrs->GetInt(2);
49+ if (end1Ptr) {end1 = static_cast<uint32_t>(*end1Ptr);}
50+ if (xShapeSize == 1){
51+ end1 = static_cast<uint32_t>(1);
52+ }
53+ const int64_t* end2Ptr = attrs->GetInt(3);
54+ if (end2Ptr) {end2 = static_cast<uint32_t>(*end2Ptr);}
55+ const int64_t* stride1Ptr = attrs->GetInt(4);
56+ if (stride1Ptr) {stride1 = static_cast<uint32_t>(*stride1Ptr);}
57+ if (xShapeSize == 1){
58+ stride1 = static_cast<uint32_t>(1);
59+ }
60+ const int64_t* stride2Ptr = attrs->GetInt(5);
61+ if (stride2Ptr) {stride2 = static_cast<uint32_t>(*stride2Ptr);}
62+ }
63+ OP_LOGD(context->GetNodeName(), "Begin to do InferShapeStridedSlice");
64+ uint32_t yRows;
65+ uint32_t yCols = (end2 - start2 + stride2 - 1) / stride2;
66+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
67+ gert::Shape* yShape = context->GetOutputShape(IDX_0);
68+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
69+ // 填充输出shape大小
70+
71+ yShape->SetDimNum(xShapeSize);
72+ if(xShapeSize == 1){
73+ yShape->SetDim(0, yCols);
74+ }else{
75+ yRows = (end1 - start1 + stride1 - 1) / stride1;
76+ yShape->SetDim(0, yRows);
77+ yShape->SetDim(1, yCols);
78+ }
79+ OP_LOGD(context->GetNodeName(), "End to do InferShapeStridedSlice");
80+ return GRAPH_SUCCESS;
81+}
82+IMPL_OP_INFERSHAPE(StridedSlice).InferShape(InferShapeStridedSlice);
83+} // namespace ops
Aexperimental/conversion/strided_slice/op_host/strided_slice_tiling.cpp+205-0
@@ -0,0 +1,205 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Liu Jun <@kbryantttt>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+/*!
21+ * \file strided_slice_tiling.cpp
22+ * \brief
23+ */
24+#include "log/log.h"
25+#include "util/math_util.h"
26+#include "tiling_base/tiling_util.h"
27+#include "tiling_base/tiling_templates_registry.h"
28+#include "../op_kernel/strided_slice_tiling_data.h"
29+#include "../op_kernel/strided_slice_tiling_key.h"
30+namespace optiling {
31+using namespace Ops::Math::OpTiling;
32+const uint32_t BLOCK_SIZE = 32;
33+const uint32_t BUFFER_NUM = 2;
34+const uint32_t WS_SYS_SIZE = 16U * 1024U * 1024U;
35+uint32_t type_Size = 0;
36+struct StridedSliceCompileInfo {};
37+// 获取平台信息如ubSize, coreNum
38+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
39+{
40+ // 获取ubsize coreNum
41+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
42+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
43+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
44+ coreNum = ascendcPlatform.GetCoreNumAiv();
45+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
46+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
47+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
48+ return ge::GRAPH_SUCCESS;
49+}
50+// 获取属性，shape信息
51+static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t& totalIdx, ge::DataType& dataType)
52+{
53+ // 获取输入shape信息
54+ auto inputX = context->GetInputShape(0);
55+ OP_CHECK_NULL_WITH_CONTEXT(context, inputX);
56+ totalIdx = inputX->GetStorageShape().GetShapeSize();
57+ // dtype校验
58+ const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT, ge::DT_INT32, ge::DT_UINT32, ge::DT_INT64, // int32对应12种数据类型
59+ ge::DT_UINT64,ge::DT_FLOAT16, ge::DT_INT16, ge::DT_UINT16,
60+ ge::DT_BF16,ge::DT_INT8, ge::DT_UINT8, ge::DT_BOOL};
61+ auto inputDesc = context->GetInputDesc(0);
62+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
63+ dataType = inputDesc->GetDataType();
64+ if (supportedDtype.count(dataType) == 0) {
65+ OP_LOGE(context, "invalid dtype");
66+ return ge::GRAPH_FAILED;
67+ }
68+ return ge::GRAPH_SUCCESS;
69+}
70+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
71+{
72+ auto ascendcPlatform = platform_ascendc:: PlatformAscendC(context->GetPlatformInfo());
73+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
74+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
75+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
76+ currentWorkspace[0] = WS_SYS_SIZE + sysWorkspaceSize;
77+ return ge::GRAPH_SUCCESS;
78+}
79+// tiling 分发入口
80+static ge::graphStatus StridedSliceTilingFunc(gert::TilingContext* context)
L
Lllimwang2025年12月19日

此函数过大，建议按功能拆分开

likedislike
81+{
82+ // 1、获取平台运行信息
83+ uint64_t ubSize;
84+ int64_t coreNum;
85+ OP_CHECK_IF(
86+ GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"),
87+ return ge::GRAPH_FAILED);
88+ // 2、获取shape、属性信息
89+ int64_t totalIdx=0;
90+ ge::DataType dataType;
91+ OP_CHECK_IF(
92+ GetShapeAttrsInfo(context, totalIdx, dataType) != ge::GRAPH_SUCCESS,
93+ OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
94+ // handle empty input
95+ if (totalIdx <= 0) {
96+ StridedSliceTilingData* tiling = context->GetTilingData<StridedSliceTilingData>();
97+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
98+ memset_s(tiling, sizeof(StridedSliceTilingData), 0, sizeof(StridedSliceTilingData));
99+ context->SetBlockDim(1);
100+ context->SetTilingKey(GET_TPL_TILING_KEY(ELEMENTWISE_TPL_SCH_MODE_0));
101+ return ge::GRAPH_SUCCESS;
102+ }
103+ // 3、获取WorkspaceSize信息
104+ OP_CHECK_IF(
105+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
106+ return ge::GRAPH_FAILED);
107+ // 4、设置tiling信息
108+ StridedSliceTilingData* tiling = context->GetTilingData<StridedSliceTilingData>();
109+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
110+ OP_CHECK_IF(
111+ memset_s(tiling, sizeof(StridedSliceTilingData), 0, sizeof(StridedSliceTilingData)) != EOK,
112+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
113+ ge::TypeUtils::GetDataTypeLength(context->GetInputDesc(0)->GetDataType(), type_Size);
114+ if (type_Size == 0) {
115+ OP_LOGE(context, "type_Size is 0");
116+ return ge::GRAPH_FAILED;
117+ }
118+ uint32_t start1 = 0;
119+ uint32_t start2 = 0;
120+ uint32_t end1 = 0;
121+ uint32_t end2 = 0;
122+ uint32_t stride1 = 0;
123+ uint32_t stride2 = 0;
124+ const auto xShape = context->GetInputTensor(0)->GetOriginShape();
125+ int64_t dim = static_cast<int64_t>(xShape.GetDimNum());
126+ auto attrs = context->GetAttrs();
127+ if (attrs) {
128+ const int64_t* start1Ptr =attrs->GetInt(0);
129+ if (start1Ptr) {start1 = static_cast<uint32_t>(*start1Ptr);}
130+ if (dim == 1){
131+ start1 = static_cast<uint32_t>(0);
132+ }
133+ const int64_t* start2Ptr =attrs->GetInt(1);
134+ if (start2Ptr) {start2 = static_cast<uint32_t>(*start2Ptr);}
135+ const int64_t* end1Ptr =attrs->GetInt(2);
136+ if (end1Ptr) {end1 = static_cast<uint32_t>(*end1Ptr);}
137+ if (dim == 1){
138+ end1 = static_cast<uint32_t>(1);
139+ }
140+ const int64_t* end2Ptr =attrs->GetInt(3);
141+ if (end2Ptr) {end2 = static_cast<uint32_t>(*end2Ptr);}
142+ const int64_t* stride1Ptr =attrs->GetInt(4);
143+ if (stride1Ptr) {stride1 = static_cast<uint32_t>(*stride1Ptr);}
144+ if (dim == 1){
145+ stride1 = static_cast<uint32_t>(1);
146+ }
147+ const int64_t* stride2Ptr =attrs->GetInt(5);
148+ if (stride2Ptr) {stride2 = static_cast<uint32_t>(*stride2Ptr);}
149+ }
150+ uint32_t ubDataNumber = 9;
151+ auto inputx = context->GetInputShape(0);
152+ auto inputShapeX = inputx->GetStorageShape();
153+ uint32_t rows;
154+ if(dim == 1){
155+ rows=1;
156+ }else{
157+ rows = static_cast<uint32_t>(inputShapeX.GetDim(0));
158+ }
159+ uint32_t cols = static_cast<uint32_t>(inputShapeX.GetDim(1));
160+ uint32_t rowBytes = cols * type_Size;
161+ uint32_t rowBytesAligned = ((rowBytes + BLOCK_SIZE - 1) / BLOCK_SIZE) * BLOCK_SIZE;
162+ uint32_t blocksPerRow = rowBytesAligned / BLOCK_SIZE;
163+ coreNum = (coreNum > rows) ? rows : coreNum;
164+ coreNum = (coreNum >= 1) ? coreNum : 1;
165+ uint32_t baseRowsPerCore = rows / coreNum;
166+ uint32_t tailRows = rows % coreNum;
167+ uint32_t bigTotalRows = baseRowsPerCore + 1;
168+ bigTotalRows = std::min(bigTotalRows, rows);
169+ uint32_t totalUbBlocks = (ubSize / BLOCK_SIZE) / BUFFER_NUM / ubDataNumber;
170+ uint32_t tileRows = (blocksPerRow == 0) ? 1 : (totalUbBlocks / blocksPerRow);
171+ tileRows = (tileRows == 0) ? 1 : tileRows;
172+ uint32_t smallTileNum = baseRowsPerCore / tileRows;
173+ uint32_t finalSmallTileNum = (baseRowsPerCore % tileRows == 0) ? smallTileNum : smallTileNum + 1;
174+ uint32_t smallTailRows = (baseRowsPerCore % tileRows == 0) ? tileRows : (baseRowsPerCore % tileRows);
175+ uint32_t bigTileNum = bigTotalRows / tileRows;
176+ uint32_t finalBigTileNum = (bigTotalRows % tileRows == 0) ? bigTileNum : bigTileNum + 1;
177+ uint32_t bigTailRows = (bigTotalRows % tileRows == 0) ? tileRows : (bigTotalRows % tileRows);
178+ tiling->cols=cols;
179+ tiling->rows=rows;
180+ tiling->start1=start1;
181+ tiling->start2=start2;
182+ tiling->end1=end1;
183+ tiling->end2=end2;
184+ tiling->stride1=stride1;
185+ tiling->stride2=stride2;
186+ tiling->tileRows=tileRows;
187+ tiling->smallTailRows=smallTailRows;
188+ tiling->bigTailRows=bigTailRows;
189+ tiling->finalSmallTileNum=finalSmallTileNum;
190+ tiling->finalBigTileNum=finalBigTileNum;
191+ tiling->baseRowsPerCore=baseRowsPerCore;
192+ tiling->bigTotalRows=bigTotalRows;
193+ tiling->tailRows=tailRows;
194+ context->SetBlockDim(coreNum);
195+ return ge::GRAPH_SUCCESS;
196+}
197+ 
198+static ge::graphStatus TilingParseForStridedSlice([[maybe_unused]] gert::TilingParseContext* context)
199+{
200+ return ge::GRAPH_SUCCESS;
201+}
202+ 
203+// tiling注册入口.
204+IMPL_OP_OPTILING(StridedSlice).Tiling(StridedSliceTilingFunc).TilingParse<StridedSliceCompileInfo>(TilingParseForStridedSlice);
205+} // namespace optiling
Aexperimental/conversion/strided_slice/op_kernel/strided_slice.cpp+41-0
@@ -0,0 +1,41 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Liu Jun <@kbryantttt>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+/*!
21+ * \file strided_slice.cpp
22+ * \brief
23+ */
24+ 
25+#include "strided_slice.h"
26+ 
27+enum class StridedSliceTilingKey : uint32_t
28+{
29+ TILING_KEY_EXAMPLE_FLOAT = 0,
30+ TILING_KEY_EXAMPLE_INT32 = 1,
31+};
32+ 
33+template <uint32_t schMode>
34+__global__ __aicore__ void strided_slice(GM_ADDR x, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
35+{
36+ REGISTER_TILING_DEFAULT(StridedSliceTilingData);
37+ GET_TILING_DATA_WITH_STRUCT(StridedSliceTilingData, tilingData, tiling);
38+ NsStridedSlice::StridedSlice<DTYPE_X> op; // 算子kernel实例获取
39+ op.Init(x, z, &tilingData); // 算子kernel实例初始化
40+ op.Process(); // 算子kernel实例执行
41+}
Aexperimental/conversion/strided_slice/op_kernel/strided_slice.h+238-0
@@ -0,0 +1,238 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Liu Jun <@kbryantttt>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+/*!
21+ * \file strided_slice.h
22+ * \brief
23+ */
24+#ifndef __STRIDED_SLICE_H__
25+#define __STRIDED_SLICE_H__
26+ 
27+#include "kernel_operator.h"
28+#include "kernel_tiling/kernel_tiling.h"
29+#include "strided_slice_tiling_data.h"
30+#include "strided_slice_tiling_key.h"
31+ 
32+namespace NsStridedSlice {
33+ 
34+using namespace AscendC;
35+ 
36+constexpr int32_t BUFFER_NUM = 2;
37+constexpr uint32_t MAX_TILE_NUM = 64;
38+constexpr uint32_t MAX_TILE_OUT_ROWS = 64;
39+ 
40+template <typename T>
41+class StridedSlice {
42+public:
43+ __aicore__ inline StridedSlice(){};
44+ 
45+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, const StridedSliceTilingData* tilingData);
46+ __aicore__ inline void Process();
47+ 
48+private:
49+ __aicore__ inline void CopyIn(int32_t tile_idx);
50+ __aicore__ inline void CopyOut(int32_t tile_idx);
51+ __aicore__ inline uint32_t CalcCoreOutRows(uint32_t core_start, uint32_t core_end);
52+ __aicore__ inline uint32_t CalcPreCoreOutTotal(uint32_t current_core_id, uint32_t total_core_num);
53+ __aicore__ inline void InitTileArrays();
54+private:
55+ AscendC::TPipe pipe;
56+ AscendC::TQueBind<TPosition::VECIN, TPosition::VECOUT, BUFFER_NUM> inQueueX;
57+ AscendC::GlobalTensor<T> xGm;
58+ AscendC::GlobalTensor<T> zGm;
59+ uint32_t tile_start_rows[MAX_TILE_NUM];
60+ uint32_t tile_end_rows[MAX_TILE_NUM];
61+ uint32_t tile_out_rows[MAX_TILE_NUM];
62+ uint32_t tile_out_start_indices[MAX_TILE_NUM];
63+ uint32_t tile_in_row_offsets[MAX_TILE_NUM][MAX_TILE_OUT_ROWS];
64+ uint32_t cols;
65+ uint32_t rows;
66+ uint32_t tileRows;
67+ uint32_t baseRowsPerCore;
68+ uint32_t bigTotalRows;
69+ uint32_t tailRows;
70+ uint32_t tileTailRows;
71+ uint32_t tileNum;
72+ uint32_t tileDataNum;
73+ uint32_t start1;
74+ uint32_t end1;
75+ uint32_t stride1;
76+ uint32_t start2;
77+ uint32_t end2;
78+ uint32_t stride2;
79+ uint32_t outNum_cols;
80+ uint32_t core_start_row;
81+ uint32_t core_last_row;
82+};
83+ 
84+template <typename T>
85+__aicore__ inline void StridedSlice<T>::Init(GM_ADDR x, GM_ADDR z, const StridedSliceTilingData* tilingData)
86+{
87+ uint32_t coreId = AscendC::GetBlockIdx();
88+ uint32_t totalCoreNum = AscendC::GetBlockNum();
89+ this->cols = tilingData->cols;
90+ this->rows = tilingData->rows;
91+ this->tileRows = tilingData->tileRows;
92+ this->baseRowsPerCore = tilingData->baseRowsPerCore;
93+ this->bigTotalRows = tilingData->bigTotalRows;
94+ this->tailRows = tilingData->tailRows;
95+ this->start1 = tilingData->start1;
96+ this->end1 = tilingData->end1;
97+ this->stride1 = tilingData->stride1;
98+ this->start2 = tilingData->start2;
99+ this->end2 = tilingData->end2;
100+ this->stride2 = tilingData->stride2;
101+ this->outNum_cols = ((this->end2 - this->start2 + this->stride2 - 1) / this->stride2);
102+ uint32_t globalBufferIndex = bigTotalRows * this->cols * coreId;
103+ this->tileDataNum = tileRows * this->cols;
104+ uint32_t coreRows;
105+ if (coreId < tailRows) {
106+ coreRows = bigTotalRows;
107+ this->tileNum = tilingData->finalBigTileNum;
108+ this->core_start_row = coreId * bigTotalRows;
109+ this->tileTailRows = tilingData->bigTailRows;
110+ }
111+ else {
112+ coreRows = baseRowsPerCore;
113+ this->tileNum = tilingData->finalSmallTileNum;
114+ this->core_start_row = tailRows * bigTotalRows + (coreId - tailRows) * baseRowsPerCore;
115+ globalBufferIndex -= (bigTotalRows - baseRowsPerCore) * this->cols * (coreId - tailRows);
116+ this->tileTailRows = tilingData->smallTailRows;
117+ }
118+ this->core_last_row = this->core_start_row + coreRows - 1;
119+ uint32_t coreOutRows = CalcCoreOutRows(core_start_row, core_last_row);
120+ uint32_t core_out_count = coreOutRows * this->outNum_cols;
121+ uint32_t core_out_offset = CalcPreCoreOutTotal(coreId, totalCoreNum);
122+ InitTileArrays();
123+ 
124+ xGm.SetGlobalBuffer((__gm__ T*)x + globalBufferIndex, coreRows*this->cols);
125+ zGm.SetGlobalBuffer((__gm__ T*)z+core_out_offset,core_out_count);
126+ pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileDataNum*(32/sizeof(T)) * sizeof(T));
127+}
128+ 
129+template <typename T>
130+__aicore__ inline void StridedSlice<T>::CopyIn(int32_t tile_idx)
131+{
132+ AscendC::LocalTensor<T> xLocal = inQueueX.AllocTensor<T>();
133+ AscendC::DataCopyExtParams copyParams{static_cast<uint16_t>(this->outNum_cols), static_cast<uint32_t>(sizeof(T)), static_cast<uint32_t>((this->stride2-1)*sizeof(T)), 0,0};
134+ AscendC::DataCopyPadExtParams<T> padParams{true, 0, static_cast<uint8_t>((32/sizeof(T))-1), 0};
135+ for(int i=0;i<this->tileRows;i++){
136+ AscendC::DataCopyPad(xLocal[i*this->outNum_cols*(32/sizeof(T))], xGm[tile_idx*this->tileDataNum+i*this->cols+this->start2],copyParams,padParams);
137+ }
138+ inQueueX.EnQue(xLocal);
139+ AscendC::SyncAll();
L
Lllimwang2025年12月19日

这里的多核同步是必须的吗？如果想实现MTE2到MTE3的同步，使用TQueBind，EnQue DeQue即可。

likedislike
liujun
liujun
2025年12月22日 评论：
140+}
141+ 
142+template <typename T>
143+__aicore__ inline void StridedSlice<T>::CopyOut(int32_t tile_idx)
144+{
145+ uint32_t tile_out_row_cnt = tile_out_rows[tile_idx];
146+ AscendC::SyncAll();
147+ uint32_t tile_out_start = tile_out_start_indices[tile_idx];
148+ AscendC::LocalTensor<T> xLocal = inQueueX.DeQue<T>();
149+ if (tile_out_row_cnt == 0) {
150+ inQueueX.FreeTensor(xLocal);
151+ return;
152+ }
153+ AscendC::DataCopyExtParams copyParams{static_cast<uint16_t>(this->outNum_cols), static_cast<uint32_t>(sizeof(T)), 0, 0, 0};
154+ for (uint32_t out_row_idx = 0; out_row_idx < tile_out_row_cnt; out_row_idx++) {
155+ uint32_t local_in_offset = tile_in_row_offsets[tile_idx][out_row_idx] * this->outNum_cols*(32/sizeof(T));
156+ uint32_t core_global_out_row = tile_out_start + out_row_idx;
157+ uint32_t global_out_offset = core_global_out_row * this->outNum_cols;
158+ AscendC::DataCopyPad(zGm[global_out_offset], xLocal[local_in_offset], copyParams);
159+ }
160+ inQueueX.FreeTensor(xLocal);
161+}
162+ 
163+template <typename T>
164+__aicore__ inline uint32_t StridedSlice<T>::CalcCoreOutRows(uint32_t core_start, uint32_t core_end)
165+{
166+ const uint32_t slice_start = this->start1;
167+ const uint32_t slice_end = this->end1 - 1;
168+ if (core_start > slice_end || core_end < slice_start) {
169+ return 0;
170+ }
171+ uint32_t clip_start = AscendC::Std::max(core_start, slice_start);
172+ uint32_t clip_end = AscendC::Std::min(core_end, slice_end);
173+ uint32_t k_first = (clip_start - slice_start + this->stride1 - 1) / this->stride1;
174+ uint32_t k_last = (clip_end - slice_start) / this->stride1;
175+ return k_last - k_first + 1;
176+}
177+ 
178+template <typename T>
179+__aicore__ inline uint32_t StridedSlice<T>::CalcPreCoreOutTotal(uint32_t current_core_id, uint32_t total_core_num)
180+{
181+ uint32_t pre_total = 0;
182+ for (uint32_t pre_core_id = 0; pre_core_id < current_core_id; pre_core_id++) {
183+ uint32_t pre_core_rows = (pre_core_id < tailRows) ? bigTotalRows : baseRowsPerCore;
184+ uint32_t pre_core_start = (pre_core_id < tailRows) ?
185+ (pre_core_id * bigTotalRows) :
186+ (tailRows * bigTotalRows + (pre_core_id - tailRows) * baseRowsPerCore);
187+ uint32_t pre_core_end = pre_core_start + pre_core_rows - 1;
188+ uint32_t pre_core_out_rows = CalcCoreOutRows(pre_core_start, pre_core_end);
189+ pre_total += pre_core_out_rows * this->outNum_cols;
190+ }
191+ return pre_total;
192+}
193+ 
194+template <typename T>
195+__aicore__ inline void StridedSlice<T>::InitTileArrays(/*参数列表*/)
196+{
197+ uint32_t current_core_row = this->core_start_row;
198+ uint32_t core_global_out_idx = 0;
199+ for (uint32_t tile_idx = 0; tile_idx < this->tileNum; tile_idx++) {
200+ uint32_t tile_rows = (tile_idx == this->tileNum - 1) ? this->tileTailRows : this->tileRows;
201+ tile_start_rows[tile_idx] = current_core_row;
202+ tile_end_rows[tile_idx] = current_core_row + tile_rows - 1;
203+ tile_end_rows[tile_idx] = AscendC::Std::min(tile_end_rows[tile_idx], this->core_last_row);
204+ const uint32_t slice_start = this->start1;
205+ const uint32_t slice_end = this->end1 - 1;
206+ const uint32_t tile_clip_start = AscendC::Std::max(tile_start_rows[tile_idx], slice_start);
207+ const uint32_t tile_clip_end = AscendC::Std::min(tile_end_rows[tile_idx], slice_end);
208+ if (tile_clip_start > tile_clip_end || this->stride1 == 0) {
209+ tile_out_rows[tile_idx] = 0;
210+ tile_out_start_indices[tile_idx] = core_global_out_idx;
211+ current_core_row = tile_end_rows[tile_idx] + 1;
212+ continue;
213+ }
214+ const uint32_t offset_from_slice_start = tile_clip_start - slice_start;
215+ const uint32_t k_first = (offset_from_slice_start + this->stride1 - 1) / this->stride1;
216+ const uint32_t k_last = (tile_clip_end - slice_start) / this->stride1;
217+ tile_out_rows[tile_idx] = k_last - k_first + 1;
218+ tile_out_start_indices[tile_idx] = core_global_out_idx;
219+ for (uint32_t out_row_idx = 0; out_row_idx < tile_out_rows[tile_idx]; out_row_idx++) {
220+ uint32_t global_in_row = slice_start + (k_first + out_row_idx) * this->stride1;
221+ tile_in_row_offsets[tile_idx][out_row_idx] = global_in_row - tile_start_rows[tile_idx];
222+ }
223+ core_global_out_idx += tile_out_rows[tile_idx];
224+ current_core_row = tile_end_rows[tile_idx] + 1;
225+ }
226+}
227+template <typename T>
228+__aicore__ inline void StridedSlice<T>::Process()
229+{
230+ int32_t loopCount = this->tileNum;
231+ for (int32_t i = 0; i < loopCount; i++) {
232+ CopyIn(i);
233+ CopyOut(i);
234+ }
235+}
236+ 
237+} // namespace NsStridedSlice
238+#endif // STRIDED_SLICE_H
Aexperimental/conversion/strided_slice/op_kernel/strided_slice_tiling_data.h+46-0
@@ -0,0 +1,46 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Liu Jun <@kbryantttt>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+/*!
21+ * \file strided_slice_tiling_data.h
22+ * \brief tiling data struct
23+ */
24+ 
25+#ifndef __STRIDED_SLICE_TILLING_DATA_H__
26+#define __STRIDED_SLICE_TILLING_DATA_H__
27+ 
28+struct StridedSliceTilingData {
29+ uint32_t cols;
30+ uint32_t rows;
31+ uint32_t tileRows;
32+ uint32_t start1;
33+ uint32_t start2;
34+ uint32_t end1;
35+ uint32_t end2;
36+ uint32_t stride1;
37+ uint32_t stride2;
38+ uint32_t smallTailRows;
39+ uint32_t bigTailRows;
40+ uint32_t finalSmallTileNum;
41+ uint32_t finalBigTileNum;
42+ uint32_t baseRowsPerCore;
43+ uint32_t bigTotalRows;
44+ uint32_t tailRows;
45+};
46+#endif
Aexperimental/conversion/strided_slice/op_kernel/strided_slice_tiling_key.h+43-0
@@ -0,0 +1,43 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Liu Jun <@kbryantttt>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+/*!
21+ * \file strided_slice_tiling_key.h
22+ * \brief strided_slice tiling key declare
23+ */
24+ 
25+#ifndef __STRIDED_SLICE_TILING_KEY_H__
26+#define __STRIDED_SLICE_TILING_KEY_H__
27+ 
28+#include "ascendc/host_api/tiling/template_argument.h"
29+ 
30+/* Mode场景定义 */
31+#define ELEMENTWISE_TPL_SCH_MODE_0 0
32+#define ELEMENTWISE_TPL_SCH_MODE_1 1
33+ 
34+/* 模板参数 */
35+ASCENDC_TPL_ARGS_DECL(
36+ StridedSlice,
37+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1));
38+ 
39+/* 模板参数组合 */
40+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
41+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1)));
42+ 
43+#endif
Mscripts/kernel/binary_config/ascendc_config.json+1-0
@@ -40,6 +40,7 @@
40 {"name":"Transposev", "compute_units": ["ascend910b"], "auto_sync" : true, "impl_mode" : "high_performance"},40 {"name":"Transposev", "compute_units": ["ascend910b"], "auto_sync" : true, "impl_mode" : "high_performance"},
41 {"name":"SelectV2", "compute_units": ["ascend910b"], "auto_sync" : true, "impl_mode" : "high_performance"},41 {"name":"SelectV2", "compute_units": ["ascend910b"], "auto_sync" : true, "impl_mode" : "high_performance"},
42 {"name":"TanV2", "compute_units": ["ascend910b", "ascend310b"], "auto_sync" : true, "impl_mode" : "high_performance"},42 {"name":"TanV2", "compute_units": ["ascend910b", "ascend310b"], "auto_sync" : true, "impl_mode" : "high_performance"},
43+ {"name":"StridedSlice", "compute_units": ["ascend910b", "ascend310b"], "auto_sync" : true, "impl_mode" : "high_performance"},
43 {"name":"AxpyV3", "compute_units": ["ascend910b"], "auto_sync" : true, "impl_mode" : "high_performance"},44 {"name":"AxpyV3", "compute_units": ["ascend910b"], "auto_sync" : true, "impl_mode" : "high_performance"},
44 {"name":"Log1pV2", "compute_units": ["ascend910b"], "auto_sync" : true, "impl_mode" : "high_performance"},45 {"name":"Log1pV2", "compute_units": ["ascend910b"], "auto_sync" : true, "impl_mode" : "high_performance"},
45 {"name":"Arange", "compute_units": ["ascend910b", "ascend310b"], "auto_sync" : true, "impl_mode" : "high_performance"},46 {"name":"Arange", "compute_units": ["ascend910b", "ascend310b"], "auto_sync" : true, "impl_mode" : "high_performance"},