已合并
math算子目录与代码中误用的arch32统一改为arch22 #4873
clinglai0517创建于 8月20日
math算子目录与代码中误用的arch32统一改为arch22 #4873
已合并
clinglai0517创建于 8月20日
共 40 个文件变更+138-138
@@ -71,7 +71,7 @@ math@ops-math:
71 - ops/ops-math/math/atanh/op_api/71 - ops/ops-math/math/atanh/op_api/
72 - ops/ops-math/math/axpy/op_api/72 - ops/ops-math/math/axpy/op_api/
73 - ops/ops-math/math/axpy_v2/op_api/73 - ops/ops-math/math/axpy_v2/op_api/
74- - ops/ops-math/math/bessel_i1e/op_host/arch32/74+ - ops/ops-math/math/bessel_i1e/op_host/arch22/
75 - ops/ops-math/math/bias_add_grad/op_kernel/75 - ops/ops-math/math/bias_add_grad/op_kernel/
76 - ops/ops-math/math/bincount/op_api/76 - ops/ops-math/math/bincount/op_api/
77 - ops/ops-math/math/bitwise_and/op_api/77 - ops/ops-math/math/bitwise_and/op_api/
@@ -125,7 +125,7 @@ math@ops-math:
125 - ops/ops-math/math/invert/op_api/125 - ops/ops-math/math/invert/op_api/
126 - ops/ops-math/math/is_close/op_api/126 - ops/ops-math/math/is_close/op_api/
127 - ops/ops-math/math/is_finite/op_api/127 - ops/ops-math/math/is_finite/op_api/
128- - ops/ops-math/math/is_finite/op_host/arch32/128+ - ops/ops-math/math/is_finite/op_host/arch22/
129 - ops/ops-math/math/is_finite/op_kernel/129 - ops/ops-math/math/is_finite/op_kernel/
130 - ops/ops-math/math/is_inf/op_host/130 - ops/ops-math/math/is_inf/op_host/
131 - ops/ops-math/math/is_inf/op_kernel/131 - ops/ops-math/math/is_inf/op_kernel/
@@ -10,5 +10,5 @@
10#Generated By CANNBot10#Generated By CANNBot
11 11 
12set(SUPPORT_COMPUTE_UNIT "ascend910b" "ascend910_93" "ascend950")12set(SUPPORT_COMPUTE_UNIT "ascend910b" "ascend910_93" "ascend950")
13-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch35")13+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch35")
14add_all_modules_sources(OPTYPE bessel_i1e ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)14add_all_modules_sources(OPTYPE bessel_i1e ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
Rmath/bessel_i1e/op_host/arch32/bessel_i1e_tiling_arch32.cpp→math/bessel_i1e/op_host/arch22/bessel_i1e_tiling_arch22.cpp+0-0
文件重命名但无更改。
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE bias_add ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE bias_add ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35")
15-add_all_modules_sources(OPTYPE cummin ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE cummin ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE cumsum ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE cumsum ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35")
15-add_all_modules_sources(OPTYPE diag_part ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE diag_part ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -9,7 +9,7 @@
9# -----------------------------------------------------------------------------------------------------------9# -----------------------------------------------------------------------------------------------------------
10 10 
11set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")11set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
12-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")12+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
13 13 
14add_all_modules_sources(OPTYPE div ACLNNTYPE aclnn_exclude DEPENDENCIES truncate_div14add_all_modules_sources(OPTYPE div ACLNNTYPE aclnn_exclude DEPENDENCIES truncate_div
15 COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}15 COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
@@ -9,7 +9,7 @@
9# -----------------------------------------------------------------------------------------------------------9# -----------------------------------------------------------------------------------------------------------
10 10 
11set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")11set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")
12-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35")12+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35")
13 13 
14add_all_modules_sources(OPTYPE div_no_nan ACLNNTYPE aclnn_exclude14add_all_modules_sources(OPTYPE div_no_nan ACLNNTYPE aclnn_exclude
15 COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}15 COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35")
15-add_all_modules_sources(OPTYPE fused_mul_add_n ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE fused_mul_add_n ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -14,5 +14,5 @@
14# 设置算子定义时支持的芯片类型14# 设置算子定义时支持的芯片类型
15set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")15set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")
16# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译16# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
17-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35")17+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35")
18-add_all_modules_sources(OPTYPE ger ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)18+add_all_modules_sources(OPTYPE ger ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE greater ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE greater ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15add_all_modules_sources(OPTYPE greater_equal ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE HOSTCPU TRUE)15add_all_modules_sources(OPTYPE greater_equal ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE HOSTCPU TRUE)
Rmath/histogram_v2/tests/ut/op_host/arch32/test_histogram_v2_tiling.cpp→math/histogram_v2/tests/ut/op_host/arch22/test_histogram_v2_tiling.cpp+3-3
@@ -8,7 +8,7 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11-// arch32 (non-RegBase: ascend310p / ascend910_93 / ascend910b) tiling UT. Each case fakes a11+// arch22 (non-RegBase: ascend310p / ascend910_93 / ascend910b) tiling UT. Each case fakes a
12// non-RegBase SoC ("ascend910b") via the histogram_v2_ut::RunTilingWithSoc helper, so12// non-RegBase SoC ("ascend910b") via the histogram_v2_ut::RunTilingWithSoc helper, so
13// IsRegbaseSocVersion() is false and the HistogramV2MembaseTiling template is selected regardless13// IsRegbaseSocVersion() is false and the HistogramV2MembaseTiling template is selected regardless
14// of the build SoC. This is compiled unconditionally alongside the arch35 (SIMT) UT, so a single14// of the build SoC. This is compiled unconditionally alongside the arch35 (SIMT) UT, so a single
@@ -33,10 +33,10 @@ using TD = gert::TilingContextPara::TensorDescription;
33 33 
34class HistogramV2MembaseTilingTest : public testing::Test {34class HistogramV2MembaseTilingTest : public testing::Test {
35protected:35protected:
36- static void SetUpTestCase() { std::cout << "HistogramV2MembaseTilingTest (arch32/MemBase) SetUp" << std::endl; }36+ static void SetUpTestCase() { std::cout << "HistogramV2MembaseTilingTest (arch22/MemBase) SetUp" << std::endl; }
37 static void TearDownTestCase()37 static void TearDownTestCase()
38 {38 {
39- std::cout << "HistogramV2MembaseTilingTest (arch32/MemBase) TearDown" << std::endl;39+ std::cout << "HistogramV2MembaseTilingTest (arch22/MemBase) TearDown" << std::endl;
40 }40 }
41};41};
42 42 
@@ -11,7 +11,7 @@
11// arch35 (RegBase / ascend950) tiling UT. Each case fakes a RegBase SoC ("ascend950") via the11// arch35 (RegBase / ascend950) tiling UT. Each case fakes a RegBase SoC ("ascend950") via the
12// histogram_v2_ut::RunTilingWithSoc helper, so IsRegbaseSocVersion() is true and the12// histogram_v2_ut::RunTilingWithSoc helper, so IsRegbaseSocVersion() is true and the
13// HistogramV2SimtTiling template is selected regardless of the build SoC. This is compiled13// HistogramV2SimtTiling template is selected regardless of the build SoC. This is compiled
14-// unconditionally alongside the arch32 (MemBase) UT, so a single build (any SoC) covers both14+// unconditionally alongside the arch22 (MemBase) UT, so a single build (any SoC) covers both
15// the SIMT and MemBase tiling templates without needing --soc=ascend950.15// the SIMT and MemBase tiling templates without needing --soc=ascend950.
16 16 
17#include <iostream>17#include <iostream>
@@ -15,7 +15,7 @@
15// The shared tiling_case_executor bakes the SoC from the compile-time BUILD_SOC_VERSION, so it can15// The shared tiling_case_executor bakes the SoC from the compile-time BUILD_SOC_VERSION, so it can
16// only exercise one template per build. This helper builds the tiling context the same way but sets16// only exercise one template per build. This helper builds the tiling context the same way but sets
17// the platform SoC from a per-case argument, so a single build (any BUILD_SOC_VERSION) can cover17// the platform SoC from a per-case argument, so a single build (any BUILD_SOC_VERSION) can cover
18-// both the SIMT (arch35) and MemBase (arch32) tiling sources. It is a local copy of the shared18+// both the SIMT (arch35) and MemBase (arch22) tiling sources. It is a local copy of the shared
19// DO_TILING flow — kept entirely inside histogram_v2 so no common test file is modified.19// DO_TILING flow — kept entirely inside histogram_v2 so no common test file is modified.
20 20 
21#ifndef HISTOGRAM_V2_TESTS_UT_OP_HOST_HISTOGRAM_V2_TILING_UT_UTIL_H21#ifndef HISTOGRAM_V2_TESTS_UT_OP_HOST_HISTOGRAM_V2_TILING_UT_UTIL_H
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "kirinx90" "kirin9030")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "kirinx90" "kirin9030")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch32" "arch32")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch22" "arch22")
15-add_all_modules_sources(OPTYPE is_finite ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR})15+add_all_modules_sources(OPTYPE is_finite ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR})
Rmath/is_finite/op_host/arch32/is_finite_tiling_arch32.cpp→math/is_finite/op_host/arch22/is_finite_tiling_arch22.cpp+21-20
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file is_finite_tiling_arch32.cpp12+ * \file is_finite_tiling_arch22.cpp
13 * \brief13 * \brief
14 */14 */
15#include "log/log.h"15#include "log/log.h"
@@ -17,7 +17,7 @@
17#include "register/op_impl_registry.h"17#include "register/op_impl_registry.h"
18#include "util/math_util.h"18#include "util/math_util.h"
19#include "is_finite_tiling_regbase.h"19#include "is_finite_tiling_regbase.h"
20-#include "is_finite_tiling_arch32.h"20+#include "is_finite_tiling_arch22.h"
21#include "op_host/tiling_base_util.h"21#include "op_host/tiling_base_util.h"
22 22 
23using namespace ge;23using namespace ge;
@@ -27,9 +27,9 @@ namespace optiling {
27 27 
28constexpr uint64_t WORK_SPACE_SIZE = 32 * 1024 * 1024;28constexpr uint64_t WORK_SPACE_SIZE = 32 * 1024 * 1024;
29 29 
30-static ge::graphStatus TilingPrepare4IsFiniteArch32(gert::TilingParseContext* context)30+static ge::graphStatus TilingPrepare4IsFiniteArch22(gert::TilingParseContext* context)
31{31{
32- auto compileInfo = context->GetCompiledInfo<IsFiniteCompileInfoArch32>();32+ auto compileInfo = context->GetCompiledInfo<IsFiniteCompileInfoArch22>();
33 OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);33 OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
34 auto platformInfo = context->GetPlatformInfo();34 auto platformInfo = context->GetPlatformInfo();
35 auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);35 auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
@@ -38,21 +38,20 @@ static ge::graphStatus TilingPrepare4IsFiniteArch32(gert::TilingParseContext* co
38 uint64_t ubSizePlatForm;38 uint64_t ubSizePlatForm;
39 ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatForm);39 ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatForm);
40 compileInfo->ubSize = static_cast<int64_t>(ubSizePlatForm);40 compileInfo->ubSize = static_cast<int64_t>(ubSizePlatForm);
41- OP_CHECK_IF(41+ OP_CHECK_IF((compileInfo->totalCoreNum <= 0 || compileInfo->ubSize <= 0),
42- (compileInfo->totalCoreNum <= 0 || compileInfo->ubSize <= 0),42+ OP_LOGE(context, "IsFinite GetHardwareInfo Failed, vectorCoreNum:%d, ubSize:%ld.",
43- OP_LOGE(43+ compileInfo->totalCoreNum, compileInfo->ubSize),
44- context, "IsFinite GetHardwareInfo Failed, vectorCoreNum:%d, ubSize:%ld.", compileInfo->totalCoreNum,44+ return ge::GRAPH_FAILED);
45- compileInfo->ubSize),
46- return ge::GRAPH_FAILED);
47 OP_LOGD(context, "Get totalCoreNum:%d, ubSize:%ld", compileInfo->totalCoreNum, compileInfo->ubSize);45 OP_LOGD(context, "Get totalCoreNum:%d, ubSize:%ld", compileInfo->totalCoreNum, compileInfo->ubSize);
48 return ge::GRAPH_SUCCESS;46 return ge::GRAPH_SUCCESS;
49}47}
50 48 
51-static ge::graphStatus IsFiniteTilingArch32(gert::TilingContext* tilingContext)49+static ge::graphStatus IsFiniteTilingArch22(gert::TilingContext* tilingContext)
52{50{
53- OP_CHECK_IF(tilingContext == nullptr, OP_LOGE("IsFiniteTiling", "tiling context is nullptr"), return ge::GRAPH_FAILED);51+ OP_CHECK_IF(tilingContext == nullptr, OP_LOGE("IsFiniteTiling", "tiling context is nullptr"),
54- OP_LOGD(tilingContext, "Entering IsFiniteTilingArch32");52+ return ge::GRAPH_FAILED);
55- auto compileInfo = reinterpret_cast<const IsFiniteCompileInfoArch32*>(tilingContext->GetCompileInfo());53+ OP_LOGD(tilingContext, "Entering IsFiniteTilingArch22");
54+ auto compileInfo = reinterpret_cast<const IsFiniteCompileInfoArch22*>(tilingContext->GetCompileInfo());
56 OP_CHECK_NULL_WITH_CONTEXT(tilingContext, compileInfo);55 OP_CHECK_NULL_WITH_CONTEXT(tilingContext, compileInfo);
57 if (compileInfo->isRegbase) {56 if (compileInfo->isRegbase) {
58 OP_LOGD(tilingContext, "Entering IsFiniteRegbaseTiling");57 OP_LOGD(tilingContext, "Entering IsFiniteRegbaseTiling");
@@ -66,21 +65,21 @@ static ge::graphStatus IsFiniteTilingArch32(gert::TilingContext* tilingContext)
66 OP_CHECK_IF(shape == nullptr, OP_LOGE(tilingContext, "InputShape == nullptr"), return ge::GRAPH_FAILED);65 OP_CHECK_IF(shape == nullptr, OP_LOGE(tilingContext, "InputShape == nullptr"), return ge::GRAPH_FAILED);
67 66 
68 IsFiniteNs::IsFiniteTilingData* tilingData = tilingContext->GetTilingData<IsFiniteNs::IsFiniteTilingData>();67 IsFiniteNs::IsFiniteTilingData* tilingData = tilingContext->GetTilingData<IsFiniteNs::IsFiniteTilingData>();
69- IsFiniteNs::IsFiniteTiling::IsFiniteCommonTiling<gert::Shape>(shape->GetStorageShape(), *tilingData, 68+ IsFiniteNs::IsFiniteTiling::IsFiniteCommonTiling<gert::Shape>(shape->GetStorageShape(), *tilingData,
70 compileInfo->totalCoreNum, compileInfo->ubSize);69 compileInfo->totalCoreNum, compileInfo->ubSize);
71 70 
72 uint32_t D_T_X = IS_FINITE_TPL_FP32, D_T_Y = IS_FINITE_TPL_BOOL;71 uint32_t D_T_X = IS_FINITE_TPL_FP32, D_T_Y = IS_FINITE_TPL_BOOL;
73 ge::DataType dtype_x = tilingContext->GetInputDesc(0)->GetDataType();72 ge::DataType dtype_x = tilingContext->GetInputDesc(0)->GetDataType();
74 if (dtype_x == ge::DataType::DT_FLOAT) {73 if (dtype_x == ge::DataType::DT_FLOAT) {
75 D_T_X = IS_FINITE_TPL_FP32;74 D_T_X = IS_FINITE_TPL_FP32;
76- } else if(dtype_x == ge::DataType::DT_FLOAT16) {75+ } else if (dtype_x == ge::DataType::DT_FLOAT16) {
77 D_T_X = IS_FINITE_TPL_FP16;76 D_T_X = IS_FINITE_TPL_FP16;
78- } else if(dtype_x == ge::DataType::DT_BF16) {77+ } else if (dtype_x == ge::DataType::DT_BF16) {
79 D_T_X = IS_FINITE_TPL_BF16;78 D_T_X = IS_FINITE_TPL_BF16;
80 }79 }
81 80 
82 tilingContext->SetBlockDim(tilingData->needCoreNum);81 tilingContext->SetBlockDim(tilingData->needCoreNum);
83- const uint64_t tilingKey = GET_TPL_TILING_KEY(D_T_X, D_T_Y); //模板参数配置82+ const uint64_t tilingKey = GET_TPL_TILING_KEY(D_T_X, D_T_Y); // 模板参数配置
84 tilingContext->SetTilingKey(tilingKey);83 tilingContext->SetTilingKey(tilingKey);
85 size_t* workspaces = tilingContext->GetWorkspaceSizes(1);84 size_t* workspaces = tilingContext->GetWorkspaceSizes(1);
86 if (workspaces != nullptr) {85 if (workspaces != nullptr) {
@@ -90,5 +89,7 @@ static ge::graphStatus IsFiniteTilingArch32(gert::TilingContext* tilingContext)
90 return ge::GRAPH_SUCCESS;89 return ge::GRAPH_SUCCESS;
91}90}
92 91 
93-IMPL_OP_OPTILING(IsFinite).Tiling(IsFiniteTilingArch32).TilingParse<IsFiniteCompileInfoArch32>(TilingPrepare4IsFiniteArch32);92+IMPL_OP_OPTILING(IsFinite)
94-} // namespace optiling93+ .Tiling(IsFiniteTilingArch22)
94+ .TilingParse<IsFiniteCompileInfoArch22>(TilingPrepare4IsFiniteArch22);
95+} // namespace optiling
Rmath/is_finite/op_host/arch32/is_finite_tiling_arch32.h→math/is_finite/op_host/arch22/is_finite_tiling_arch22.h+5-5
@@ -9,12 +9,12 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file is_finite_tiling_arch32.h12+ * \file is_finite_tiling_arch22.h
13 * \brief13 * \brief
14 */14 */
15 15 
16-#ifndef IS_FINITE_TILING_ARCH32_H16+#ifndef IS_FINITE_TILING_ARCH22_H
17-#define IS_FINITE_TILING_ARCH32_H17+#define IS_FINITE_TILING_ARCH22_H
18 18 
19#include "../../op_kernel/is_finite_struct.h"19#include "../../op_kernel/is_finite_struct.h"
20#include "torch_extension/tiling_utils_math.h"20#include "torch_extension/tiling_utils_math.h"
@@ -66,10 +66,10 @@ public:
66} // namespace IsFiniteNs66} // namespace IsFiniteNs
67 67 
68namespace optiling {68namespace optiling {
69-struct IsFiniteCompileInfoArch32 {69+struct IsFiniteCompileInfoArch22 {
70 int32_t totalCoreNum = 0;70 int32_t totalCoreNum = 0;
71 int64_t ubSize = 0;71 int64_t ubSize = 0;
72 bool isRegbase = false;72 bool isRegbase = false;
73};73};
74} // namespace optiling74} // namespace optiling
75-#endif // IS_FINITE_TILING_ARCH32_H75+#endif // IS_FINITE_TILING_ARCH22_H
Rmath/is_finite/op_host/arch32/is_finite_tiling_regbase.cpp→math/is_finite/op_host/arch22/is_finite_tiling_regbase.cpp+21-30
@@ -30,9 +30,8 @@ ge::graphStatus IsFiniteRegbaseTiling::CalcInputDtype()
30 this->inputDtype = inputDesc->GetDataType();30 this->inputDtype = inputDesc->GetDataType();
31 OP_CHECK_IF(31 OP_CHECK_IF(
32 this->inputDtype != ge::DT_BF16 && this->inputDtype != ge::DT_FLOAT16 && this->inputDtype != ge::DT_FLOAT,32 this->inputDtype != ge::DT_BF16 && this->inputDtype != ge::DT_FLOAT16 && this->inputDtype != ge::DT_FLOAT,
33- OP_LOGE(33+ OP_LOGE(tilingContext, "Input x dtype not supported, only support [DT_FLOAT, DT_FLOAT16, DT_BF16], got %s",
34- tilingContext, "Input x dtype not supported, only support [DT_FLOAT, DT_FLOAT16, DT_BF16], got %s",34+ ge::TypeUtils::DataTypeToSerialString(this->inputDtype).c_str()),
35- ge::TypeUtils::DataTypeToSerialString(this->inputDtype).c_str()),
36 return ge::GRAPH_FAILED);35 return ge::GRAPH_FAILED);
37 return ge::GRAPH_SUCCESS;36 return ge::GRAPH_SUCCESS;
38}37}
@@ -42,12 +41,10 @@ ge::graphStatus IsFiniteRegbaseTiling::CalcOutputDtype()
42 auto outputDesc = tilingContext->GetOutputDesc(0);41 auto outputDesc = tilingContext->GetOutputDesc(0);
43 OP_CHECK_NULL_WITH_CONTEXT(tilingContext, outputDesc);42 OP_CHECK_NULL_WITH_CONTEXT(tilingContext, outputDesc);
44 this->outputDtype = outputDesc->GetDataType();43 this->outputDtype = outputDesc->GetDataType();
45- OP_CHECK_IF(44+ OP_CHECK_IF(this->outputDtype != ge::DT_BOOL,
46- this->outputDtype != ge::DT_BOOL,45+ OP_LOGE(tilingContext, "Output y dtype not supported, only support [DT_BOOL], got %s",
47- OP_LOGE(46+ ge::TypeUtils::DataTypeToSerialString(this->outputDtype).c_str()),
48- tilingContext, "Output y dtype not supported, only support [DT_BOOL], got %s",47+ return ge::GRAPH_FAILED);
49- ge::TypeUtils::DataTypeToSerialString(this->outputDtype).c_str()),
50- return ge::GRAPH_FAILED);
51 return ge::GRAPH_SUCCESS;48 return ge::GRAPH_SUCCESS;
52}49}
53 50 
@@ -61,9 +58,8 @@ ge::graphStatus IsFiniteRegbaseTiling::CheckShape()
61 OP_CHECK_NULL_WITH_CONTEXT(tilingContext, outputStorageShape);58 OP_CHECK_NULL_WITH_CONTEXT(tilingContext, outputStorageShape);
62 const gert::Shape& outputYShape = Ops::Base::EnsureNotScalar(outputStorageShape->GetStorageShape());59 const gert::Shape& outputYShape = Ops::Base::EnsureNotScalar(outputStorageShape->GetStorageShape());
63 60 
64- OP_CHECK_IF(61+ OP_CHECK_IF(inputXShape != outputYShape, OP_LOGE(tilingContext, "Input x and output y shape not the same"),
65- inputXShape != outputYShape, OP_LOGE(tilingContext, "Input x and output y shape not the same"),62+ return ge::GRAPH_FAILED);
66- return ge::GRAPH_FAILED);
67 return ge::GRAPH_SUCCESS;63 return ge::GRAPH_SUCCESS;
68}64}
69 65 
@@ -85,19 +81,16 @@ ge::graphStatus IsFiniteRegbaseTiling::SetTilingData()
85ge::graphStatus IsFiniteRegbaseTiling::RunTiling()81ge::graphStatus IsFiniteRegbaseTiling::RunTiling()
86{82{
87 ElewiseBaseTiling elewiseBaseTiling(tilingContext);83 ElewiseBaseTiling elewiseBaseTiling(tilingContext);
88- OP_CHECK_IF(84+ OP_CHECK_IF(CalcInputDtype() == ge::GRAPH_FAILED, OP_LOGE(tilingContext, "Get input dtype failed"),
89- CalcInputDtype() == ge::GRAPH_FAILED, OP_LOGE(tilingContext, "Get input dtype failed"),85+ return ge::GRAPH_FAILED);
90- return ge::GRAPH_FAILED);86+ OP_CHECK_IF(CalcOutputDtype() == ge::GRAPH_FAILED, OP_LOGE(tilingContext, "Get output dtype failed"),
91- OP_CHECK_IF(87+ return ge::GRAPH_FAILED);
92- CalcOutputDtype() == ge::GRAPH_FAILED, OP_LOGE(tilingContext, "Get output dtype failed"),88+ OP_CHECK_IF(CheckShape() == ge::GRAPH_FAILED, OP_LOGE(tilingContext, "Check shape failed"),
93- return ge::GRAPH_FAILED);89+ return ge::GRAPH_FAILED);
94- OP_CHECK_IF(
95- CheckShape() == ge::GRAPH_FAILED, OP_LOGE(tilingContext, "Check shape failed"), return ge::GRAPH_FAILED);
96 90 
97 tiling_ = tilingContext->GetTilingData<EleBaseTilingDataV2>();91 tiling_ = tilingContext->GetTilingData<EleBaseTilingDataV2>();
98- OP_CHECK_IF(92+ OP_CHECK_IF((tiling_ == nullptr), OP_LOGE(tilingContext, "Get EleBaseTilingDataV2 from context failed"),
99- (tiling_ == nullptr), OP_LOGE(tilingContext, "Get EleBaseTilingDataV2 from context failed"),93+ return ge::GRAPH_FAILED);
100- return ge::GRAPH_FAILED);
101 94 
102 ge::graphStatus baseTilingResult = ge::GRAPH_FAILED;95 ge::graphStatus baseTilingResult = ge::GRAPH_FAILED;
103 if (this->inputDtype == ge::DT_FLOAT16) {96 if (this->inputDtype == ge::DT_FLOAT16) {
@@ -110,15 +103,13 @@ ge::graphStatus IsFiniteRegbaseTiling::RunTiling()
110 dType = TPL_FP32;103 dType = TPL_FP32;
111 baseTilingResult = elewiseBaseTiling.DoTiling<IsFiniteDag<float>::OpDag>(*tiling_, ASCEND_API_BUFFER);104 baseTilingResult = elewiseBaseTiling.DoTiling<IsFiniteDag<float>::OpDag>(*tiling_, ASCEND_API_BUFFER);
112 } else {105 } else {
113- OP_LOGE(106+ OP_LOGE(tilingContext, "Input x dtype not supported, only support [DT_FLOAT, DT_FLOAT16, DT_BF16], got %s",
114- tilingContext, "Input x dtype not supported, only support [DT_FLOAT, DT_FLOAT16, DT_BF16], got %s",107+ ge::TypeUtils::DataTypeToSerialString(this->inputDtype).c_str());
115- ge::TypeUtils::DataTypeToSerialString(this->inputDtype).c_str());
116 return ge::GRAPH_FAILED;108 return ge::GRAPH_FAILED;
117 }109 }
118- OP_CHECK_IF(110+ OP_CHECK_IF(baseTilingResult == ge::GRAPH_FAILED, OP_LOGE(tilingContext, "ElewiseBaseTiling failed"),
119- baseTilingResult == ge::GRAPH_FAILED, OP_LOGE(tilingContext, "ElewiseBaseTiling failed"),111+ return ge::GRAPH_FAILED);
120- return ge::GRAPH_FAILED);
121 112 
122 return SetTilingData();113 return SetTilingData();
123}114}
124-} // namespace optiling115+} // namespace optiling
Rmath/is_finite/op_host/arch32/is_finite_tiling_regbase.h→math/is_finite/op_host/arch22/is_finite_tiling_regbase.h+2-3
@@ -21,8 +21,7 @@
21namespace optiling {21namespace optiling {
22using namespace Ops::Base;22using namespace Ops::Base;
23 23 
24-class IsFiniteRegbaseTiling24+class IsFiniteRegbaseTiling {
25-{
26public:25public:
27 explicit IsFiniteRegbaseTiling(gert::TilingContext* context) : tilingContext(context) {};26 explicit IsFiniteRegbaseTiling(gert::TilingContext* context) : tilingContext(context) {};
28 ge::graphStatus RunTiling();27 ge::graphStatus RunTiling();
@@ -41,4 +40,4 @@ private:
41 ge::DataType outputDtype;40 ge::DataType outputDtype;
42};41};
43} // namespace optiling42} // namespace optiling
44-#endif // OPS_BUILD_IN_OP_TILING_RUNTIME_IS_FINITE_TILING_REGBASE_H43+#endif // OPS_BUILD_IN_OP_TILING_RUNTIME_IS_FINITE_TILING_REGBASE_H
@@ -15,7 +15,7 @@ if (UT_TEST_ALL OR OP_KERNEL_UT)
15 # )15 # )
16 # 算子自己的tiling文件路径16 # 算子自己的tiling文件路径
17 set(is_finite_tiling_files17 set(is_finite_tiling_files
18- ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch32/is_finite_tiling_arch32.cpp18+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch22/is_finite_tiling_arch22.cpp
19 ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/is_finite_tiling_arch35.cpp19 ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/is_finite_tiling_arch35.cpp
20 )20 )
21 # list(APPEND is_finite_tiling_files ${elewise_common_tiling_files})21 # list(APPEND is_finite_tiling_files ${elewise_common_tiling_files})
@@ -22,7 +22,7 @@
22#include "tikicpulib.h"22#include "tikicpulib.h"
23#include "data_utils.h"23#include "data_utils.h"
24#include "tiling_case_executor.h"24#include "tiling_case_executor.h"
25-#include "../../../op_host/arch32/is_finite_tiling_arch32.h"25+#include "../../../op_host/arch22/is_finite_tiling_arch22.h"
26 26 
27using namespace std;27using namespace std;
28 28 
@@ -30,15 +30,14 @@ extern "C" __global__ __aicore__ void is_finite(GM_ADDR x, GM_ADDR y, GM_ADDR wo
30 30 
31class IsFiniteTest : public testing::Test {31class IsFiniteTest : public testing::Test {
32protected:32protected:
33- static void SetUpTestCase() {33+ static void SetUpTestCase()
34+ {
34 std::cout << "is_finite_test SetUp" << std::endl;35 std::cout << "is_finite_test SetUp" << std::endl;
35 const string cmd = "cp -rf " + dataPath + " ./";36 const string cmd = "cp -rf " + dataPath + " ./";
36 system(cmd.c_str());37 system(cmd.c_str());
37 system("chmod -R 755 ./is_finite_data/");38 system("chmod -R 755 ./is_finite_data/");
38 }39 }
39- static void TearDownTestCase() {40+ static void TearDownTestCase() { std::cout << "is_finite_test TearDown" << std::endl; }
40- std::cout << "is_finite_test TearDown" << std::endl;
41- }
42 41 
43private:42private:
44 const static std::string rootPath;43 const static std::string rootPath;
@@ -49,15 +48,21 @@ const std::string IsFiniteTest::rootPath = "../../../../";
49const std::string IsFiniteTest::dataPath = rootPath + "math/is_finite/tests/ut/op_kernel/is_finite_data";48const std::string IsFiniteTest::dataPath = rootPath + "math/is_finite/tests/ut/op_kernel/is_finite_data";
50 49 
51template <typename T1, typename T2>50template <typename T1, typename T2>
52-inline T1 CeilAlign(T1 a, T2 b) {51+inline T1 CeilAlign(T1 a, T2 b)
52+{
53 return (a + b - 1) / b * b;53 return (a + b - 1) / b * b;
54}54}
55 55 
56-TEST_F(IsFiniteTest, test_case_float16_1) {56+TEST_F(IsFiniteTest, test_case_float16_1)
57- optiling::IsFiniteCompileInfoArch32 compileInfo = {64, 262144, false};57+{
58+ optiling::IsFiniteCompileInfoArch22 compileInfo = {64, 262144, false};
58 gert::TilingContextPara tilingContextPara("IsFinite",59 gert::TilingContextPara tilingContextPara("IsFinite",
59- {{{{128, 64}, {128, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND},},60+ {
60- {{{{128, 64}, {128, 64}}, ge::DT_BOOL, ge::FORMAT_ND},},61+ {{{128, 64}, {128, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND},
62+ },
63+ {
64+ {{{128, 64}, {128, 64}}, ge::DT_BOOL, ge::FORMAT_ND},
65+ },
61 &compileInfo);66 &compileInfo);
62 TilingInfo tilingInfo;67 TilingInfo tilingInfo;
63 auto tilingRet = ExecuteTiling(tilingContextPara, tilingInfo);68 auto tilingRet = ExecuteTiling(tilingContextPara, tilingInfo);
@@ -66,7 +71,8 @@ TEST_F(IsFiniteTest, test_case_float16_1) {
66 system("cd ./is_finite_data/ && python3 gen_data.py '(128, 64)' 'float16'");71 system("cd ./is_finite_data/ && python3 gen_data.py '(128, 64)' 'float16'");
67 uint32_t dataCount = 128 * 64;72 uint32_t dataCount = 128 * 64;
68 size_t inputByteSize = dataCount * sizeof(half);73 size_t inputByteSize = dataCount * sizeof(half);
69- std::string fileName = "./is_finite_data/float16_input_t_is_finite.bin";;74+ std::string fileName = "./is_finite_data/float16_input_t_is_finite.bin";
75+ ;
70 uint8_t* x = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));76 uint8_t* x = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));
71 ReadFile(fileName, inputByteSize, x, inputByteSize);77 ReadFile(fileName, inputByteSize, x, inputByteSize);
72 size_t outputByteSize = dataCount * sizeof(bool);78 size_t outputByteSize = dataCount * sizeof(bool);
@@ -90,11 +96,16 @@ TEST_F(IsFiniteTest, test_case_float16_1) {
90 system("cd ./is_finite_data/ && python3 compare_data.py 'float16'");96 system("cd ./is_finite_data/ && python3 compare_data.py 'float16'");
91}97}
92 98 
93-TEST_F(IsFiniteTest, test_case_float16_2) {99+TEST_F(IsFiniteTest, test_case_float16_2)
94- optiling::IsFiniteCompileInfoArch32 compileInfo = {64, 262144, false};100+{
101+ optiling::IsFiniteCompileInfoArch22 compileInfo = {64, 262144, false};
95 gert::TilingContextPara tilingContextPara("IsFinite",102 gert::TilingContextPara tilingContextPara("IsFinite",
96- {{{{256, 33}, {256, 33}}, ge::DT_FLOAT, ge::FORMAT_ND},},103+ {
97- {{{{256, 33}, {256, 33}}, ge::DT_BOOL, ge::FORMAT_ND},},104+ {{{256, 33}, {256, 33}}, ge::DT_FLOAT, ge::FORMAT_ND},
105+ },
106+ {
107+ {{{256, 33}, {256, 33}}, ge::DT_BOOL, ge::FORMAT_ND},
108+ },
98 &compileInfo);109 &compileInfo);
99 TilingInfo tilingInfo;110 TilingInfo tilingInfo;
100 auto tilingRet = ExecuteTiling(tilingContextPara, tilingInfo);111 auto tilingRet = ExecuteTiling(tilingContextPara, tilingInfo);
@@ -103,7 +114,8 @@ TEST_F(IsFiniteTest, test_case_float16_2) {
103 system("cd ./is_finite_data/ && python3 gen_data.py '(256, 33)' 'float32'");114 system("cd ./is_finite_data/ && python3 gen_data.py '(256, 33)' 'float32'");
104 uint32_t dataCount = 256 * 33;115 uint32_t dataCount = 256 * 33;
105 size_t inputByteSize = dataCount * sizeof(float);116 size_t inputByteSize = dataCount * sizeof(float);
106- std::string fileName = "./is_finite_data/float32_input_t_is_finite.bin";;117+ std::string fileName = "./is_finite_data/float32_input_t_is_finite.bin";
118+ ;
107 uint8_t* x = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));119 uint8_t* x = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));
108 ReadFile(fileName, inputByteSize, x, inputByteSize);120 ReadFile(fileName, inputByteSize, x, inputByteSize);
109 size_t outputByteSize = dataCount * sizeof(bool);121 size_t outputByteSize = dataCount * sizeof(bool);
@@ -144,7 +156,6 @@ TEST_F(IsFiniteTest, test_case_float16_2) {
144// ReadFile(fileName, inputByteSize, x, inputByteSize);156// ReadFile(fileName, inputByteSize, x, inputByteSize);
145// IsFiniteTilingData* tilingDatafromBin = reinterpret_cast<IsFiniteTilingData*>(tiling);157// IsFiniteTilingData* tilingDatafromBin = reinterpret_cast<IsFiniteTilingData*>(tiling);
146 158 
147- 
148// tilingDatafromBin->totalDataCount = dataCount;159// tilingDatafromBin->totalDataCount = dataCount;
149// tilingDatafromBin->usableUbSize = 32 * 1024;160// tilingDatafromBin->usableUbSize = 32 * 1024;
150// tilingDatafromBin->needCoreNum = 1;161// tilingDatafromBin->needCoreNum = 1;
@@ -156,7 +167,6 @@ TEST_F(IsFiniteTest, test_case_float16_2) {
156// AscendC::SetKernelMode(KernelMode::AIV_MODE);167// AscendC::SetKernelMode(KernelMode::AIV_MODE);
157// ICPU_RUN_KF(is_finite, numBlocks, x, y, workspace, tiling);168// ICPU_RUN_KF(is_finite, numBlocks, x, y, workspace, tiling);
158 169 
159- 
160// fileName = "./is_finite_data/bfloat16_output_t_is_finite.bin";170// fileName = "./is_finite_data/bfloat16_output_t_is_finite.bin";
161// WriteFile(fileName, y, outputByteSize);171// WriteFile(fileName, y, outputByteSize);
162 172 
@@ -167,4 +177,3 @@ TEST_F(IsFiniteTest, test_case_float16_2) {
167 177 
168// system("cd ./is_finite_data/ && python3 compare_data.py 'bfloat16'");178// system("cd ./is_finite_data/ && python3 compare_data.py 'bfloat16'");
169// }179// }
170- 
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35")
15-add_all_modules_sources(OPTYPE log1p ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE log1p ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35")
15-add_all_modules_sources(OPTYPE masked_scale ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE masked_scale ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE maximum ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE maximum ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE minimum ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE minimum ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE mul ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE mul ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE muls ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE DEPENDENCIES mul)15+add_all_modules_sources(OPTYPE muls ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE DEPENDENCIES mul)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE neg ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE neg ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE real_div ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE real_div ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE reduce_max ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE reduce_max ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE reduce_mean ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE reduce_mean ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE reduce_min ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE reduce_min ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE reduce_prod ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE reduce_prod ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35")
15-add_all_modules_sources(OPTYPE reduce_std_v2 ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE DEPENDENCIES reduce_var)15+add_all_modules_sources(OPTYPE reduce_std_v2 ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE DEPENDENCIES reduce_var)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35" "arch35")
15-add_all_modules_sources(OPTYPE reduce_sum ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE reduce_sum ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35")
15-add_all_modules_sources(OPTYPE reduce_var ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE reduce_var ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -11,5 +11,5 @@
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")12set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")
13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35")14+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35")
15-add_all_modules_sources(OPTYPE sort_with_index ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE DEPENDENCIES sort)15+add_all_modules_sources(OPTYPE sort_with_index ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE DEPENDENCIES sort)
@@ -9,7 +9,7 @@
9# -----------------------------------------------------------------------------------------------------------9# -----------------------------------------------------------------------------------------------------------
10 10 
11set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")11set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")
12-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35")12+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22" "arch35")
13 13 
14add_all_modules_sources(OPTYPE trunc ACLNNTYPE aclnn_exclude14add_all_modules_sources(OPTYPE trunc ACLNNTYPE aclnn_exclude
15 COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}15 COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}