Star718
269
代码
代码
Issues66
Pull Requests162
Actions
项目讨论
Wiki
分析
Star718
269
已合并
【fix】: Asin/Acos/Remainder修改为Simt接口 #2045
gaoxin创建于 19 天前
【fix】: Asin/Acos/Remainder修改为Simt接口 #2045
已合并
gaoxin创建于 19 天前
从gcw_V3YyYBt1/graph-autofusionfix/asin-acos-simt合入到cann/graph-autofusiondevelop
共 10 个文件变更+133-43
Mautofuse/common/ascgraph_info_complete.cpp+8-1
@@ -80,7 +80,11 @@ void CompleteStoreApiInfo(af::AscNodePtr &node) {
80void CompleteElewiseApiInfo(af::AscNodePtr &node) {80void CompleteElewiseApiInfo(af::AscNodePtr &node) {
81 node->attr.api.type = af::ApiType::kAPITypeCompute;81 node->attr.api.type = af::ApiType::kAPITypeCompute;
82 node->attr.api.unit = af::ComputeUnit::kUnitVector;82 node->attr.api.unit = af::ComputeUnit::kUnitVector;
83- if (af::ops::IsOps<Expm1>(node) || af::ops::IsOps<Sin>(node) || af::ops::IsOps<Cos>(node)) {83+ if (af::ops::IsOps<Expm1>(node) || af::ops::IsOps<Sin>(node) || af::ops::IsOps<Cos>(node) ||
84+ af::ops::IsOps<Asin>(node) || af::ops::IsOps<Acos>(node)) {
85+ (void)::ascir::SetDcacheSize(node, kSimtDcacheSize);
86+ }
87+ if (af::ops::IsOps<Remainder>(node) && node->inputs[0].attr.dtype == ge::DT_INT32) {
84 (void)::ascir::SetDcacheSize(node, kSimtDcacheSize);88 (void)::ascir::SetDcacheSize(node, kSimtDcacheSize);
85 }89 }
86}90}
@@ -168,6 +172,8 @@ static const std::map<std::string, af::ComputeType> kOpTypeToComputeType = {
168 {IsFinite::Type, af::ComputeType::kComputeElewise},172 {IsFinite::Type, af::ComputeType::kComputeElewise},
169 {Sin::Type, af::ComputeType::kComputeElewise},173 {Sin::Type, af::ComputeType::kComputeElewise},
170 {Cos::Type, af::ComputeType::kComputeElewise},174 {Cos::Type, af::ComputeType::kComputeElewise},
175+ {Asin::Type, af::ComputeType::kComputeElewise},
176+ {Acos::Type, af::ComputeType::kComputeElewise},
171 {Ln::Type, af::ComputeType::kComputeElewise},177 {Ln::Type, af::ComputeType::kComputeElewise},
172 {Expm1::Type, af::ComputeType::kComputeElewise},178 {Expm1::Type, af::ComputeType::kComputeElewise},
173 {LogicalNot::Type, af::ComputeType::kComputeElewise},179 {LogicalNot::Type, af::ComputeType::kComputeElewise},
@@ -176,6 +182,7 @@ static const std::map<std::string, af::ComputeType> kOpTypeToComputeType = {
176 {Sub::Type, af::ComputeType::kComputeElewise},182 {Sub::Type, af::ComputeType::kComputeElewise},
177 {Mul::Type, af::ComputeType::kComputeElewise},183 {Mul::Type, af::ComputeType::kComputeElewise},
178 {Div::Type, af::ComputeType::kComputeElewise},184 {Div::Type, af::ComputeType::kComputeElewise},
185+ {Remainder::Type, af::ComputeType::kComputeElewise},
179 {TrueDiv::Type, af::ComputeType::kComputeElewise},186 {TrueDiv::Type, af::ComputeType::kComputeElewise},
180 {Minimum::Type, af::ComputeType::kComputeElewise},187 {Minimum::Type, af::ComputeType::kComputeElewise},
181 {Maximum::Type, af::ComputeType::kComputeElewise},188 {Maximum::Type, af::ComputeType::kComputeElewise},
Mautofuse/tests/v35/st/backend_e2e_v2/acos_bf16_test/acos_bf16_backend_generator.cpp+8-0
@@ -24,6 +24,10 @@
24#include "runtime_stub.h"24#include "runtime_stub.h"
25#include "common/platform_context.h"25#include "common/platform_context.h"
26 26 
27+namespace {
28+constexpr int64_t kAcosSimtDcacheSize = 40 * 1024;
29+}
30+ 
27class TestBackendAcosBf16E2e : public testing::Test {31class TestBackendAcosBf16E2e : public testing::Test {
28 protected:32 protected:
29 void SetUp() override {33 void SetUp() override {
@@ -68,6 +72,10 @@ TEST_F(TestBackendAcosBf16E2e, AcosBf16E2eCodegen) {
68 EXPECT_EQ(optimizer.Optimize(graph, fused_schedule_result), 0);72 EXPECT_EQ(optimizer.Optimize(graph, fused_schedule_result), 0);
69 codegen::CodegenResult result;73 codegen::CodegenResult result;
70 EXPECT_EQ(codegen.Generate(shape_info, fused_schedule_result, result), 0);74 EXPECT_EQ(codegen.Generate(shape_info, fused_schedule_result, result), 0);
75+ EXPECT_NE(result.kernel.find("AcosExtend"), std::string::npos);
76+ EXPECT_NE(result.kernel.find("AcosSimtCompute"), std::string::npos);
77+ EXPECT_NE(result.kernel.find("simt_api/cpp/kernel_simt_intf.h"), std::string::npos);
78+ EXPECT_NE(result.tiling.find(std::to_string(kAcosSimtDcacheSize)), std::string::npos);
71 kernel_file << tilig_stub << RemoveSubDirInclude(result.kernel);79 kernel_file << tilig_stub << RemoveSubDirInclude(result.kernel);
72 tiling_file << result.tiling;80 tiling_file << result.tiling;
73 tiling_data_file << result.tiling_data;81 tiling_data_file << result.tiling_data;
Mautofuse/tests/v35/st/backend_e2e_v2/asin_bf16_test/asin_bf16_backend_generate.cpp+8-0
@@ -13,6 +13,10 @@
13#include "runtime_stub.h"13#include "runtime_stub.h"
14#include "common/platform_context.h"14#include "common/platform_context.h"
15 15 
16+namespace {
17+constexpr int64_t kAsinSimtDcacheSize = 40 * 1024;
18+}
19+ 
16class TestBackendAsinBf16E2e : public testing::Test {20class TestBackendAsinBf16E2e : public testing::Test {
17 protected:21 protected:
18 void SetUp() override {22 void SetUp() override {
@@ -57,6 +61,10 @@ TEST_F(TestBackendAsinBf16E2e, AsinBf16E2eCodegen) {
57 EXPECT_EQ(optimizer.Optimize(graph, fused_schedule_result), 0);61 EXPECT_EQ(optimizer.Optimize(graph, fused_schedule_result), 0);
58 codegen::CodegenResult result;62 codegen::CodegenResult result;
59 EXPECT_EQ(codegen.Generate(shape_info, fused_schedule_result, result), 0);63 EXPECT_EQ(codegen.Generate(shape_info, fused_schedule_result, result), 0);
64+ EXPECT_NE(result.kernel.find("AsinExtend"), std::string::npos);
65+ EXPECT_NE(result.kernel.find("AsinSimtCompute"), std::string::npos);
66+ EXPECT_NE(result.kernel.find("simt_api/cpp/kernel_simt_intf.h"), std::string::npos);
67+ EXPECT_NE(result.tiling.find(std::to_string(kAsinSimtDcacheSize)), std::string::npos);
60 kernel_file << tilig_stub << RemoveSubDirInclude(result.kernel);68 kernel_file << tilig_stub << RemoveSubDirInclude(result.kernel);
61 tiling_file << result.tiling;69 tiling_file << result.tiling;
62 tiling_data_file << result.tiling_data;70 tiling_data_file << result.tiling_data;
Mautofuse/tests/v35/st/backend_e2e_v2/remainder_int32_store_test/remainder_int32_store_backend_generator.cpp+4-0
@@ -41,6 +41,7 @@ class TestBackendRemainderInt32StoreE2e : public testing::Test {
41 41 
42TEST_F(TestBackendRemainderInt32StoreE2e, RemainderInt32StoreE2eCodegen) {42TEST_F(TestBackendRemainderInt32StoreE2e, RemainderInt32StoreE2eCodegen) {
43 bool gen_success = true;43 bool gen_success = true;
44+ constexpr int64_t kRemainderSimtDcacheSize = 40 * 1024;
44 std::string remainder_tiling_stub = R"(45 std::string remainder_tiling_stub = R"(
45#define REGISTER_TILING_DEFAULT(tiling)46#define REGISTER_TILING_DEFAULT(tiling)
46#define GET_TILING_DATA(t, tiling) AutofuseTilingData t = *(AutofuseTilingData*)tiling;47#define GET_TILING_DATA(t, tiling) AutofuseTilingData t = *(AutofuseTilingData*)tiling;
@@ -66,7 +67,10 @@ TEST_F(TestBackendRemainderInt32StoreE2e, RemainderInt32StoreE2eCodegen) {
66 codegen::CodegenResult remainder_codegen_result;67 codegen::CodegenResult remainder_codegen_result;
67 EXPECT_EQ(remainder_codegen.Generate(remainder_shape_info, remainder_fused_result, remainder_codegen_result), 0);68 EXPECT_EQ(remainder_codegen.Generate(remainder_shape_info, remainder_fused_result, remainder_codegen_result), 0);
68 EXPECT_NE(remainder_codegen_result.kernel.find("RemainderExtend"), std::string::npos);69 EXPECT_NE(remainder_codegen_result.kernel.find("RemainderExtend"), std::string::npos);
70+ EXPECT_NE(remainder_codegen_result.kernel.find("RemainderIntSimtCompute"), std::string::npos);
71+ EXPECT_NE(remainder_codegen_result.kernel.find("simt_api/cpp/kernel_simt_intf.h"), std::string::npos);
69 EXPECT_NE(remainder_codegen_result.kernel.find("tmp_buf_"), std::string::npos);72 EXPECT_NE(remainder_codegen_result.kernel.find("tmp_buf_"), std::string::npos);
73+ EXPECT_NE(remainder_codegen_result.tiling.find(std::to_string(kRemainderSimtDcacheSize)), std::string::npos);
70 remainder_kernel_stream << remainder_tiling_stub << RemoveSubDirInclude(remainder_codegen_result.kernel);74 remainder_kernel_stream << remainder_tiling_stub << RemoveSubDirInclude(remainder_codegen_result.kernel);
71 remainder_tiling_stream << remainder_codegen_result.tiling;75 remainder_tiling_stream << remainder_codegen_result.tiling;
72 remainder_data_stream << remainder_codegen_result.tiling_data;76 remainder_data_stream << remainder_codegen_result.tiling_data;
Mautofuse/v35/ascendc/api_regbase/CMakeLists.txt+2-0
@@ -100,6 +100,8 @@ set(ascendc_api_regbase_extend_src
100 sinh.h100 sinh.h
101 asinh.h101 asinh.h
102 i1.h102 i1.h
103+ asin.h
104+ acos.h
103)105)
104 106 
105set(generated_regbase_files)107set(generated_regbase_files)
Aautofuse/v35/ascendc/api_regbase/acos.h+29-0
@@ -0,0 +1,29 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+#ifndef __ASCENDC_API_REGBASE_ACOS_H__
11+#define __ASCENDC_API_REGBASE_ACOS_H__
12+ 
13+constexpr uint32_t ACOS_THREAD_NUM = 1024;
14+ 
15+template <typename T>
16+__simt_vf__ __aicore__ LAUNCH_BOUND(ACOS_THREAD_NUM) inline void AcosSimtCompute(__ubuf__ T *x, __ubuf__ T *y,
17+ const int64_t total_num) {
18+ for (int64_t i = threadIdx.x; i < total_num; i += blockDim.x) {
19+ y[i] = Simt::Acos(x[i]);
20+ }
21+}
22+ 
23+template <typename T>
24+__aicore__ inline void AcosExtend(const LocalTensor<T> &dst, const LocalTensor<T> &src,
25+ const LocalTensor<uint8_t> &tmp_buf, const uint32_t calc_cnt) {
26+ AscendC::Simt::VF_CALL<AcosSimtCompute<T>>(AscendC::Simt::Dim3(ACOS_THREAD_NUM), (__ubuf__ T *)src.GetPhyAddr(),
27+ (__ubuf__ T *)dst.GetPhyAddr(), calc_cnt);
28+}
29+#endif // __ASCENDC_API_REGBASE_ACOS_H__
Aautofuse/v35/ascendc/api_regbase/asin.h+29-0
@@ -0,0 +1,29 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+#ifndef __ASCENDC_API_REGBASE_ASIN_H__
11+#define __ASCENDC_API_REGBASE_ASIN_H__
12+ 
13+constexpr uint32_t ASIN_THREAD_NUM = 1024;
14+ 
15+template <typename T>
16+__simt_vf__ __aicore__ LAUNCH_BOUND(ASIN_THREAD_NUM) inline void AsinSimtCompute(__ubuf__ T *x, __ubuf__ T *y,
17+ const int64_t total_num) {
18+ for (int64_t i = threadIdx.x; i < total_num; i += blockDim.x) {
19+ y[i] = Simt::Asin(x[i]);
20+ }
21+}
22+ 
23+template <typename T>
24+__aicore__ inline void AsinExtend(const LocalTensor<T> &dst, const LocalTensor<T> &src,
25+ const LocalTensor<uint8_t> &tmp_buf, const uint32_t calc_cnt) {
26+ AscendC::Simt::VF_CALL<AsinSimtCompute<T>>(AscendC::Simt::Dim3(ASIN_THREAD_NUM), (__ubuf__ T *)src.GetPhyAddr(),
27+ (__ubuf__ T *)dst.GetPhyAddr(), calc_cnt);
28+}
29+#endif // __ASCENDC_API_REGBASE_ASIN_H__
Mautofuse/v35/ascendc/api_regbase/remainder.h+18-40
@@ -12,6 +12,21 @@
12 12 
13// Remainder(x1, x2) = x1 - x2 * floor(x1/x2)13// Remainder(x1, x2) = x1 - x2 * floor(x1/x2)
14 14 
15+// int32 走 SIMT 实现（参考 ops-math FloorModInt_1）
16+constexpr uint32_t REMAINDER_THREAD_NUM = 1024;
17+ 
18+template <typename T>
19+__simt_vf__ __aicore__ LAUNCH_BOUND(REMAINDER_THREAD_NUM) inline void RemainderIntSimtCompute(__ubuf__ T *dst,
20+ __ubuf__ T *src1,
21+ __ubuf__ T *src2,
22+ const int64_t total_num) {
23+ for (int64_t i = threadIdx.x; i < total_num; i += blockDim.x) {
24+ const auto rem = src1[i] % src2[i];
25+ const bool signs_differ = (rem < 0) != (src2[i] < 0);
26+ dst[i] = (signs_differ && rem != 0) ? rem + src2[i] : rem;
27+ }
28+}
29+ 
15template <typename T>30template <typename T>
16__aicore__ inline void RemainderImplVF(__ubuf__ T *dst, __ubuf__ T *src1, __ubuf__ T *src2, uint32_t count,31__aicore__ inline void RemainderImplVF(__ubuf__ T *dst, __ubuf__ T *src1, __ubuf__ T *src2, uint32_t count,
17 uint16_t repeat_time) {32 uint16_t repeat_time) {
@@ -42,42 +57,6 @@ __aicore__ inline void RemainderImplVF(__ubuf__ T *dst, __ubuf__ T *src1, __ubuf
42 }57 }
43}58}
44 59 
45-template <typename T>
46-__aicore__ inline void RemainderIntImplVF(__ubuf__ T *dst, __ubuf__ T *src1, __ubuf__ T *src2, uint32_t count,
47- uint16_t repeatTime) {
48- constexpr uint32_t oneRepElm = static_cast<uint32_t>(AscendC::GetVecLen() / sizeof(T));
49- 
50- for (uint16_t i = 0; i < repeatTime; i++) {
51- AscendC::Reg::RegTensor<T> srcReg1;
52- AscendC::Reg::RegTensor<T> srcReg2;
53- AscendC::Reg::RegTensor<T> remReg;
54- AscendC::Reg::MaskReg mask;
55- AscendC::Reg::MaskReg signDiffMask;
56- AscendC::Reg::MaskReg src1Mask;
57- 
58- mask = AscendC::Reg::UpdateMask<T>(count);
59- 
60- AscendC::Reg::DataCopy(srcReg1, src1 + i * oneRepElm);
61- AscendC::Reg::DataCopy(srcReg2, src2 + i * oneRepElm);
62- 
63- // Calculate x1 mod x2
64- AscendC::Reg::Div(remReg, srcReg1, srcReg2, mask);
65- AscendC::Reg::Mul(remReg, srcReg2, remReg, mask);
66- AscendC::Reg::Sub(remReg, srcReg1, remReg, mask);
67- 
68- // sign diff mask
69- AscendC::Reg::Compares<T, CMPMODE::GT>(src1Mask, srcReg1, 0, mask);
70- AscendC::Reg::Compares<T, CMPMODE::GT>(signDiffMask, srcReg2, 0, mask);
71- AscendC::Reg::Xor(signDiffMask, src1Mask, signDiffMask, mask);
72- AscendC::Reg::Compares<T, CMPMODE::NE>(src1Mask, remReg, 0, mask);
73- AscendC::Reg::And(signDiffMask, signDiffMask, src1Mask, mask);
74- 
75- AscendC::Reg::Add<T, AscendC::Reg::MaskMergeMode::MERGING>(remReg, remReg, srcReg2, signDiffMask);
76- 
77- AscendC::Reg::DataCopy(dst + i * oneRepElm, remReg, mask);
78- }
79-}
80- 
81template <typename T>60template <typename T>
82__aicore__ inline void RemainderExtend(const AscendC::LocalTensor<T> &dst, const AscendC::LocalTensor<T> &src1,61__aicore__ inline void RemainderExtend(const AscendC::LocalTensor<T> &dst, const AscendC::LocalTensor<T> &src1,
83 const AscendC::LocalTensor<T> &src2, const uint32_t size) {62 const AscendC::LocalTensor<T> &src2, const uint32_t size) {
@@ -93,12 +72,11 @@ __aicore__ inline void RemainderExtend(const AscendC::LocalTensor<T> &dst, const
93 const uint32_t size) {72 const uint32_t size) {
94 static_assert(SupportType<T, int32_t, float>(),73 static_assert(SupportType<T, int32_t, float>(),
95 "RemainderExtend only support int32_t/float data type on current device!");74 "RemainderExtend only support int32_t/float data type on current device!");
96- constexpr uint32_t oneRepElm = static_cast<uint32_t>(AscendC::GetVecLen() / sizeof(T));
97- uint16_t repeatTime = static_cast<uint16_t>(AscendC::CeilDivision(size, oneRepElm));
98 75 
99 if constexpr (IsSameType<T, int32_t>::value) {76 if constexpr (IsSameType<T, int32_t>::value) {
100- VF_CALL<RemainderIntImplVF<T>>((__ubuf__ T *)dst.GetPhyAddr(), (__ubuf__ T *)src1.GetPhyAddr(),77+ AscendC::Simt::VF_CALL<RemainderIntSimtCompute<T>>(AscendC::Simt::Dim3(REMAINDER_THREAD_NUM),
101- (__ubuf__ T *)src2.GetPhyAddr(), size, repeatTime);78+ (__ubuf__ T *)dst.GetPhyAddr(), (__ubuf__ T *)src1.GetPhyAddr(),
79+ (__ubuf__ T *)src2.GetPhyAddr(), size);
102 } else {80 } else {
103 RemainderExtend(dst, src1, src2, size);81 RemainderExtend(dst, src1, src2, size);
104 }82 }
Mautofuse/v35/ascir/generator/v2_ascir_codegen_impl.h+19-2
@@ -3880,7 +3880,7 @@ class AcosAscIrCodegenImplV2 : public SimtFloatUnaryAscIrCodegenImplV2 {
3880 return "UnaryApiTmpCall";3880 return "UnaryApiTmpCall";
3881 }3881 }
3882 [[nodiscard]] std::string GetApiName() const override {3882 [[nodiscard]] std::string GetApiName() const override {
3883- return "Acos";3883+ return "AcosExtend";
3884 }3884 }
3885 [[nodiscard]] std::string GetSimtScalarApiName() const override {3885 [[nodiscard]] std::string GetSimtScalarApiName() const override {
3886 return "Acos";3886 return "Acos";
@@ -3890,12 +3890,18 @@ class AcosAscIrCodegenImplV2 : public SimtFloatUnaryAscIrCodegenImplV2 {
3890 const AscNode &node) {3890 const AscNode &node) {
3891 std::map<ge::DataType, ge::DataType> dtype_conversion_map = {3891 std::map<ge::DataType, ge::DataType> dtype_conversion_map = {
3892 {DT_BF16, DT_FLOAT},3892 {DT_BF16, DT_FLOAT},
3893+ {DT_FLOAT16, DT_FLOAT},
3893 };3894 };
3894 return GetConversionFromDtypeMap(node, dtype_conversion_map);3895 return GetConversionFromDtypeMap(node, dtype_conversion_map);
3895 }3896 }
3897+ [[nodiscard]] std::vector<std::string> LoadApiHeaderFiles([[maybe_unused]] bool is_dynamic) const override {
3898+ return {"acos_reg_base.h"};
3899+ }
3896 [[nodiscard]] std::vector<std::string> IncludeApiHeaderFiles() const override {3900 [[nodiscard]] std::vector<std::string> IncludeApiHeaderFiles() const override {
3897 return {3901 return {
3898 "adv_api/math/acos.h",3902 "adv_api/math/acos.h",
3903+ "simt_api/cpp/kernel_simt_intf.h",
3904+ "simt_api/math_functions.h",
3899 };3905 };
3900 }3906 }
3901 [[nodiscard]] bool IsNodeValid(const AscNode &node) const override {3907 [[nodiscard]] bool IsNodeValid(const AscNode &node) const override {
@@ -3987,7 +3993,7 @@ class AsinAscIrCodegenImplV2 : public SimtFloatUnaryAscIrCodegenImplV2 {
3987 return "UnaryApiTmpCall";3993 return "UnaryApiTmpCall";
3988 }3994 }
3989 [[nodiscard]] std::string GetApiName() const override {3995 [[nodiscard]] std::string GetApiName() const override {
3990- return "Asin";3996+ return "AsinExtend";
3991 }3997 }
3992 [[nodiscard]] std::string GetSimtScalarApiName() const override {3998 [[nodiscard]] std::string GetSimtScalarApiName() const override {
3993 return "Asin";3999 return "Asin";
@@ -3996,12 +4002,18 @@ class AsinAscIrCodegenImplV2 : public SimtFloatUnaryAscIrCodegenImplV2 {
3996 const AscNode &node) override {4002 const AscNode &node) override {
3997 std::map<ge::DataType, ge::DataType> dtype_conversion_map = {4003 std::map<ge::DataType, ge::DataType> dtype_conversion_map = {
3998 {DT_BF16, DT_FLOAT},4004 {DT_BF16, DT_FLOAT},
4005+ {DT_FLOAT16, DT_FLOAT},
3999 };4006 };
4000 return GetConversionFromDtypeMap(node, dtype_conversion_map);4007 return GetConversionFromDtypeMap(node, dtype_conversion_map);
4001 }4008 }
4009+ [[nodiscard]] std::vector<std::string> LoadApiHeaderFiles([[maybe_unused]] bool is_dynamic) const override {
4010+ return {"asin_reg_base.h"};
4011+ }
4002 [[nodiscard]] std::vector<std::string> IncludeApiHeaderFiles() const override {4012 [[nodiscard]] std::vector<std::string> IncludeApiHeaderFiles() const override {
4003 return {4013 return {
4004 "adv_api/math/asin.h",4014 "adv_api/math/asin.h",
4015+ "simt_api/cpp/kernel_simt_intf.h",
4016+ "simt_api/math_functions.h",
4005 };4017 };
4006 }4018 }
4007 [[nodiscard]] bool IsNodeValid(const AscNode &node) const override {4019 [[nodiscard]] bool IsNodeValid(const AscNode &node) const override {
@@ -5363,6 +5375,11 @@ class RemainderAscIrCodegenImplV2 : public AscIrCodegenV2 {
5363 [[nodiscard]] std::vector<std::string> LoadApiHeaderFiles([[maybe_unused]] bool is_dynamic) const override {5375 [[nodiscard]] std::vector<std::string> LoadApiHeaderFiles([[maybe_unused]] bool is_dynamic) const override {
5364 return {"remainder_reg_base.h"};5376 return {"remainder_reg_base.h"};
5365 }5377 }
5378+ [[nodiscard]] std::vector<std::string> IncludeApiHeaderFiles() const override {
5379+ return {
5380+ "simt_api/cpp/kernel_simt_intf.h",
5381+ };
5382+ }
5366 [[nodiscard]] std::pair<std::vector<ge::DataType>, std::vector<ge::DataType>> GetConversionDtype(5383 [[nodiscard]] std::pair<std::vector<ge::DataType>, std::vector<ge::DataType>> GetConversionDtype(
5367 const AscNode &node) override {5384 const AscNode &node) override {
5368 std::map<ge::DataType, ge::DataType> dtype_conversion_map = {{DT_BF16, DT_FLOAT}};5385 std::map<ge::DataType, ge::DataType> dtype_conversion_map = {{DT_BF16, DT_FLOAT}};
Mautofuse/v35/codegen/ascendc_reg_base_api_register.cpp+8-0
@@ -317,6 +317,12 @@ Register::Register() {
317 };317 };
318 const std::string kAscendcAtanRegBaseStr = {318 const std::string kAscendcAtanRegBaseStr = {
319#include "atan_reg_base.h"319#include "atan_reg_base.h"
320+ };
321+ const std::string kAscendcAsinRegBaseStr = {
322+#include "asin_reg_base.h"
323+ };
324+ const std::string kAscendcAcosRegBaseStr = {
325+#include "acos_reg_base.h"
320 };326 };
321 std::unordered_map<std::string, std::string> api_to_file{327 std::unordered_map<std::string, std::string> api_to_file{
322 {"cast_reg_base.h", kAscendcCastRegStr},328 {"cast_reg_base.h", kAscendcCastRegStr},
@@ -413,6 +419,8 @@ Register::Register() {
413 {"sin_reg_base.h", kAscendcSinRegBaseStr},419 {"sin_reg_base.h", kAscendcSinRegBaseStr},
414 {"cos_reg_base.h", kAscendcCosRegBaseStr},420 {"cos_reg_base.h", kAscendcCosRegBaseStr},
415 {"atan_reg_base.h", kAscendcAtanRegBaseStr},421 {"atan_reg_base.h", kAscendcAtanRegBaseStr},
422+ {"asin_reg_base.h", kAscendcAsinRegBaseStr},
423+ {"acos_reg_base.h", kAscendcAcosRegBaseStr},
416 };424 };
417 425 
418 AscendCApiRegistry::GetInstance().RegisterApi(api_to_file);426 AscendCApiRegistry::GetInstance().RegisterApi(api_to_file);