Star1295
1409
代码
代码
Issues104
Pull Requests209
Actions
项目讨论
Wiki
分析
Star1295
1409
已合并
[CANNBot]TanV3、SinhV2算子支持A2 AscendC实现 #1969
Hana77创建于 3月28日
[CANNBot]TanV3、SinhV2算子支持A2 AscendC实现 #1969
已合并
Hana77创建于 3月28日
从Hana77/ops-mathmaster合入到cann/ops-mathmaster
共 22 个文件变更+1527-0
Aexperimental/math/sinh_v2/CMakeLists.txt+24-0
@@ -0,0 +1,24 @@
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+ 
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,
8+# INCLUDING 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+#
13+# NOTE: Portions of this code were AI-generated and have been
14+# technically reviewed for functional accuracy and security
15+ 
16+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
17+if(NOT ENABLE_TEST)
18+ list(REMOVE_ITEM CURRENT_DIRS tests)
19+endif()
20+foreach(SUB_DIR ${CURRENT_DIRS})
21+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
22+ add_subdirectory(${SUB_DIR})
23+ endif()
24+endforeach()
Aexperimental/math/sinh_v2/examples/test_aclnn_sinh_v2.cpp+143-0
@@ -0,0 +1,143 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+#include <iostream>
18+#include <vector>
19+#include "acl/acl.h"
20+#include "aclnnop/aclnn_sinh_v2.h"
21+ 
22+#define CHECK_RET(cond, return_expr) \
23+ do { \
24+ if (!(cond)) { \
25+ return_expr; \
26+ } \
27+} while (0)
28+ 
29+#define LOG_PRINT(message, ...) \
30+ do { \
31+ printf(message, ##__VA_ARGS__); \
32+} while (0)
33+ 
34+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
35+ int64_t shape_size = 1;
36+ for (auto i : shape) {
37+ shape_size *= i;
38+ }
39+ return shape_size;
40+}
41+ 
42+int Init(int32_t deviceId, aclrtStream* stream) {
43+ // 固定写法，资源初始化
44+ auto ret = aclInit(nullptr);
45+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
46+ ret = aclrtSetDevice(deviceId);
47+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
48+ ret = aclrtCreateStream(stream);
49+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
50+ return 0;
51+}
52+ 
53+template <typename T>
54+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
55+ aclDataType dataType, aclTensor** tensor) {
56+ auto size = GetShapeSize(shape) * sizeof(T);
57+ // 调用aclrtMalloc申请Device侧内存
58+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
59+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
60+ // 调用aclrtMemcpy将Host侧数据拷贝到Device侧内存上
61+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
62+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
63+ // 计算连续tensor的strides
64+ std::vector<int64_t> strides(shape.size(), 1);
65+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
66+ strides[i] = shape[i + 1] * strides[i + 1];
67+ }
68+ // 调用aclCreateTensor接口创建aclTensor
69+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
70+ shape.data(), shape.size(), *deviceAddr);
71+ return 0;
72+}
73+ 
74+int main() {
75+ // 1. （固定写法）device/stream初始化, 参考acl对外接口列表
76+ // 根据自己的实际device填写deviceId
77+ int32_t deviceId = 0;
78+ aclrtStream stream;
79+ auto ret = Init(deviceId, &stream);
80+ // check根据自己的需要处理
81+ CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
82+ 
83+ // 2. 构造输入与输出，需要根据API的接口自定义构造
84+ std::vector<int64_t> xShape = {4, 2};
85+ std::vector<int64_t> outShape = {4, 2};
86+ void* xDeviceAddr = nullptr;
87+ void* outDeviceAddr = nullptr;
88+ aclTensor* x = nullptr;
89+ aclTensor* out = nullptr;
90+ std::vector<float> xHostData = {1, 2, 3, 4, 5, 6, 7, 8};
91+ std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0};
92+ // 创建x aclTensor
93+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_FLOAT, &x);
94+ CHECK_RET(ret == ACL_SUCCESS, return ret);
95+ // 创建out aclTensor
96+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
97+ CHECK_RET(ret == ACL_SUCCESS, return ret);
98+ 
99+ // 3. 调用CANN算子库API
100+ uint64_t workspaceSize = 0;
101+ aclOpExecutor* executor;
102+ // 调用aclnnSinhV2第一段接口
103+ ret = aclnnSinhV2GetWorkspaceSize(x, out, &workspaceSize, &executor);
104+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSinhV2GetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
105+ // 根据第一段接口计算出的workspaceSize申请device内存
106+ void* workspaceAddr = nullptr;
107+ if (workspaceSize > 0) {
108+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
109+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret;);
110+ }
111+ // 调用aclnnSinhV2第二段接口
112+ ret = aclnnSinhV2(workspaceAddr, workspaceSize, executor, stream);
113+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSinhV2 failed. ERROR: %d\n", ret); return ret);
114+ 
115+ // 4. （固定写法）同步等待任务执行结束
116+ ret = aclrtSynchronizeStream(stream);
117+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
118+ 
119+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
120+ auto size = GetShapeSize(outShape);
121+ std::vector<float> resultData(size, 0);
122+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(float),
123+ ACL_MEMCPY_DEVICE_TO_HOST);
124+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
125+ for (int64_t i = 0; i < size; i++) {
126+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
127+ }
128+ 
129+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
130+ aclDestroyTensor(x);
131+ aclDestroyTensor(out);
132+ 
133+ // 7. 释放device资源，需要根据具体API的接口定义修改
134+ aclrtFree(xDeviceAddr);
135+ aclrtFree(outDeviceAddr);
136+ if (workspaceSize > 0) {
137+ aclrtFree(workspaceAddr);
138+ }
139+ aclrtDestroyStream(stream);
140+ aclrtResetDevice(deviceId);
141+ aclFinalize();
142+ return 0;
143+}
Aexperimental/math/sinh_v2/op_host/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
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+
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,
8+# INCLUDING 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+#
13+# NOTE: Portions of this code were AI-generated and have been
14+# technically reviewed for functional accuracy and security
15+ 
16+add_modules_sources(OPTYPE sinh_v2 ACLNNTYPE aclnn)
Aexperimental/math/sinh_v2/op_host/sinh_v2_def.cpp+53-0
@@ -0,0 +1,53 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file sinh_v2_def.cpp
19+ * \brief SinhV2 operator definition - declares inputs, outputs, and chip configuration
20+ */
21+#include "register/op_def_registry.h"
22+ 
23+namespace ops {
24+class SinhV2 : public OpDef {
25+public:
26+ explicit SinhV2(const char* name) : OpDef(name)
27+ {
28+ this->Input("x")
29+ .ParamType(REQUIRED)
30+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
31+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
32+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
33+ .AutoContiguous();
34+ this->Output("y")
35+ .ParamType(REQUIRED)
36+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
37+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
38+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
39+ .AutoContiguous();
40+ 
41+ OpAICoreConfig aicoreConfig910B;
42+ aicoreConfig910B.DynamicCompileStaticFlag(true)
43+ .DynamicFormatFlag(false)
44+ .DynamicRankSupportFlag(true)
45+ .DynamicShapeSupportFlag(true)
46+ .NeedCheckSupportFlag(false)
47+ .PrecisionReduceFlag(true)
48+ .ExtendCfgInfo("opFile.value", "sinh_v2");
49+ this->AICore().AddConfig("ascend910b", aicoreConfig910B);
50+ }
51+};
52+OP_ADD(SinhV2);
53+} // namespace ops
Aexperimental/math/sinh_v2/op_host/sinh_v2_infershape.cpp+48-0
@@ -0,0 +1,48 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file sinh_v2_infershape.cpp
19+ * \brief SinhV2 shape inference - output shape equals input shape
20+ */
21+ 
22+#include "register/op_impl_registry.h"
23+#include "exe_graph/runtime/infer_shape_context.h"
24+ 
25+using namespace ge;
26+ 
27+namespace ops {
28+ 
29+static ge::graphStatus InferShape4SinhV2(gert::InferShapeContext* context)
30+{
31+ const gert::Shape* input_shape = context->GetInputShape(0);
32+ if (input_shape == nullptr) {
33+ return ge::GRAPH_FAILED;
34+ }
35+ 
36+ gert::Shape* output_shape = context->GetOutputShape(0);
37+ if (output_shape == nullptr) {
38+ return ge::GRAPH_FAILED;
39+ }
40+ 
41+ *output_shape = *input_shape;
42+ 
43+ return ge::GRAPH_SUCCESS;
44+}
45+ 
46+IMPL_OP_INFERSHAPE(SinhV2).InferShape(InferShape4SinhV2);
47+ 
48+} // namespace ops
Aexperimental/math/sinh_v2/op_host/sinh_v2_tiling.cpp+158-0
@@ -0,0 +1,158 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file sinh_v2_tiling.cpp
19+ * \brief SinhV2 Tiling implementation (arch32 - Ascend910B)
20+ *
21+ * Computes tiling parameters for SinhV2 operator:
22+ * - Multi-core splitting: totalNum divided evenly across AI Cores
23+ * - UB splitting: each core processes data in chunks of ubFactor elements
24+ * - TilingKey selection: based on input dtype (float32 vs float16)
25+ */
26+ 
27+#include "register/op_def_registry.h"
28+#include "op_common/log/log.h"
29+#include "op_common/op_host/util/math_util.h"
30+#include "op_common/op_host/util/platform_util.h"
31+#include "../op_kernel/sinh_v2_tiling_data.h"
32+#include "../op_kernel/sinh_v2_tiling_key.h"
33+ 
34+namespace optiling {
35+ 
36+using Ops::Base::CeilDiv;
37+using Ops::Base::FloorDiv;
38+using Ops::Base::FloorAlign;
39+using Ops::Base::GetUbBlockSize;
40+ 
41+constexpr uint32_t WS_SYS_SIZE = 0U;
42+// float32: inputQueue(x2) + outputQueue(x2) + tmpBuf1(x1) + tmpBuf2(x1) = 6 float-sized buffers
43+// Total bytes = ubFactor * 6 * sizeof(float)
44+constexpr int64_t BUFFER_NUM_FP32 = 6;
45+// float16: inputQueue(x2,half) + outputQueue(x2,half) + tmpBuf1(x1,float) + tmpBuf2(x1,float)
46+// Total bytes = ubFactor * (2*2 + 2*2 + 4 + 4) = ubFactor * 16 = ubFactor * sizeof(float) * 4
47+constexpr int64_t BUFFER_NUM_FP16 = 4;
48+ 
49+static const gert::Shape g_vec_1_shape = {1};
50+ 
51+static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape)
52+{
53+ if (in_shape.GetDimNum() == 0) {
54+ return g_vec_1_shape;
55+ }
56+ return in_shape;
57+}
58+ 
59+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
60+{
61+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
62+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
63+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
64+ coreNum = ascendcPlatform.GetCoreNumAiv();
65+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
66+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
67+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
68+ return ge::GRAPH_SUCCESS;
69+}
70+ 
71+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
72+{
73+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
74+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
75+ currentWorkspace[0] = WS_SYS_SIZE;
76+ return ge::GRAPH_SUCCESS;
77+}
78+ 
79+static ge::graphStatus SinhV2TilingFunc(gert::TilingContext* context)
80+{
81+ // 1. Get platform info
82+ uint64_t ubSize;
83+ int64_t coreNum;
84+ OP_CHECK_IF(
85+ GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
86+ OP_LOGE(context, "GetPlatformInfo error"),
87+ return ge::GRAPH_FAILED);
88+ 
89+ // 2. Get input shape and dtype
90+ auto inputShape = context->GetInputShape(0);
91+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
92+ auto storageShape = EnsureNotScalar(inputShape->GetStorageShape());
93+ int64_t totalNum = storageShape.GetShapeSize();
94+ 
95+ auto inputDesc = context->GetInputDesc(0);
96+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
97+ ge::DataType dtype = inputDesc->GetDataType();
98+ 
99+ // 3. Get workspace size
100+ OP_CHECK_IF(
101+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
102+ OP_LOGE(context, "GetWorkspaceSize error"),
103+ return ge::GRAPH_FAILED);
104+ 
105+ // 4. Set TilingData
106+ SinhV2TilingData* tiling = context->GetTilingData<SinhV2TilingData>();
107+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
108+ OP_CHECK_IF(
109+ memset_s(tiling, sizeof(SinhV2TilingData), 0, sizeof(SinhV2TilingData)) != EOK,
110+ OP_LOGE(context, "set tiling data error"),
111+ return ge::GRAPH_FAILED);
112+ 
113+ // Handle empty tensor
114+ if (totalNum == 0) {
115+ tiling->totalNum = 0;
116+ tiling->blockFactor = 0;
117+ tiling->ubFactor = 0;
118+ context->SetBlockDim(1);
119+ context->SetTilingKey(GET_TPL_TILING_KEY(SINHV2_TPL_SCH_MODE_0));
120+ return ge::GRAPH_SUCCESS;
121+ }
122+ 
123+ // 5. Multi-core splitting
124+ tiling->totalNum = totalNum;
125+ tiling->blockFactor = CeilDiv(totalNum, coreNum);
126+ int64_t usedCoreNum = CeilDiv(totalNum, tiling->blockFactor);
127+ 
128+ // 6. UB splitting and TilingKey selection
129+ int64_t ubCanUse = static_cast<int64_t>(ubSize);
130+ int64_t ubBlockSize = GetUbBlockSize(context);
131+ // Both paths compute internally in float32, so typeSize = 4
132+ constexpr int64_t typeSize = 4;
133+ 
134+ if (dtype == ge::DT_FLOAT) {
135+ tiling->ubFactor = FloorAlign(FloorDiv((ubCanUse / typeSize), BUFFER_NUM_FP32), ubBlockSize);
136+ context->SetTilingKey(GET_TPL_TILING_KEY(SINHV2_TPL_SCH_MODE_0));
137+ } else if (dtype == ge::DT_FLOAT16) {
138+ tiling->ubFactor = FloorAlign(FloorDiv((ubCanUse / typeSize), BUFFER_NUM_FP16), ubBlockSize);
139+ context->SetTilingKey(GET_TPL_TILING_KEY(SINHV2_TPL_SCH_MODE_1));
140+ } else {
141+ OP_LOGE(context, "SinhV2: unsupported dtype");
142+ return ge::GRAPH_FAILED;
143+ }
144+ 
145+ context->SetBlockDim(usedCoreNum);
146+ return ge::GRAPH_SUCCESS;
147+}
148+ 
149+static ge::graphStatus TilingParseForSinhV2([[maybe_unused]] gert::TilingParseContext* context)
150+{
151+ return ge::GRAPH_SUCCESS;
152+}
153+ 
154+struct SinhV2CompileInfo {};
155+ 
156+IMPL_OP_OPTILING(SinhV2).Tiling(SinhV2TilingFunc).TilingParse<SinhV2CompileInfo>(TilingParseForSinhV2);
157+ 
158+} // namespace optiling
Aexperimental/math/sinh_v2/op_kernel/sinh_v2.cpp+45-0
@@ -0,0 +1,45 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file sinh_v2_arch32.cpp
19+ * \brief SinhV2 kernel entry point (arch32 architecture - Ascend910B)
20+ */
21+ 
22+#include "sinh_v2.h"
23+ 
24+enum class SinhV2TilingKey : uint32_t
25+{
26+ TILING_KEY_FLOAT32 = 0,
27+ TILING_KEY_FLOAT16 = 1,
28+};
29+ 
30+template <uint32_t schMode>
31+__global__ __aicore__ void sinh_v2(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
32+{
33+ REGISTER_TILING_DEFAULT(SinhV2TilingData);
34+ GET_TILING_DATA_WITH_STRUCT(SinhV2TilingData, tilingData, tiling);
35+ if constexpr (schMode == static_cast<uint32_t>(SinhV2TilingKey::TILING_KEY_FLOAT32)) {
36+ NsSinhV2::SinhV2<float> op;
37+ op.Init(x, y, &tilingData);
38+ op.Process();
39+ }
40+ if constexpr (schMode == static_cast<uint32_t>(SinhV2TilingKey::TILING_KEY_FLOAT16)) {
41+ NsSinhV2::SinhV2<half> op;
42+ op.Init(x, y, &tilingData);
43+ op.Process();
44+ }
45+}
Aexperimental/math/sinh_v2/op_kernel/sinh_v2.h+213-0
@@ -0,0 +1,213 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file sinh_v2.h
19+ * \brief SinhV2 kernel class definition (arch32 - Ascend910B)
20+ *
21+ * Computes sinh(x) = (exp(x) - exp(-x)) / 2 for each element.
22+ * - float32 path: direct computation
23+ * - float16 path: cast to float32 for computation, then cast back
24+ *
25+ * Uses double buffering (BUFFER_NUM=2) for pipeline parallelism:
26+ * - CopyIn: GM -> UB (input data transfer)
27+ * - Compute: UB vector computation (sinh formula)
28+ * - CopyOut: UB -> GM (output data transfer)
29+ */
30+#ifndef SINH_V2_H
31+#define SINH_V2_H
32+ 
33+#include "kernel_operator.h"
34+#include "kernel_tiling/kernel_tiling.h"
35+#include "sinh_v2_tiling_data.h"
36+#include "sinh_v2_tiling_key.h"
37+ 
38+namespace NsSinhV2 {
39+ 
40+using namespace AscendC;
41+ 
42+constexpr int32_t BUFFER_NUM = 2;
43+ 
44+template <typename T>
45+class SinhV2 {
46+public:
47+ __aicore__ inline SinhV2() {};
48+ 
49+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const SinhV2TilingData* tilingData);
50+ __aicore__ inline void Process();
51+ 
52+private:
53+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum);
54+ __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum);
55+ __aicore__ inline void Compute(int64_t currentNum);
56+ 
57+private:
58+ TPipe pipe;
59+ TQue<QuePosition::VECIN, BUFFER_NUM> inputQueue;
60+ TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueue;
61+ TBuf<QuePosition::VECCALC> tmpBuf1;
62+ TBuf<QuePosition::VECCALC> tmpBuf2;
63+ 
64+ GlobalTensor<T> inputGM;
65+ GlobalTensor<T> outputGM;
66+ 
67+ int64_t blockLength_ = 0;
68+ int64_t ubLength_ = 0;
69+};
70+ 
71+// ============================================================================
72+// Init - Initialize GM pointers and allocate UB buffers
73+// ============================================================================
74+template <typename T>
75+__aicore__ inline void SinhV2<T>::Init(GM_ADDR x, GM_ADDR y, const SinhV2TilingData* tilingData)
76+{
77+ int64_t remainderLength = tilingData->totalNum - tilingData->blockFactor * AscendC::GetBlockIdx();
78+ blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength;
79+ // Clamp to 0 for idle cores (when totalNum < coreNum, some cores have no work)
80+ if (blockLength_ < 0) {
81+ blockLength_ = 0;
82+ }
83+ ubLength_ = tilingData->ubFactor;
84+ 
85+ // Guard: empty tensor or idle core - skip buffer allocation
86+ if (blockLength_ <= 0 || ubLength_ <= 0) {
87+ return;
88+ }
89+ 
90+ inputGM.SetGlobalBuffer((__gm__ T*)x + tilingData->blockFactor * AscendC::GetBlockIdx(), blockLength_);
91+ outputGM.SetGlobalBuffer((__gm__ T*)y + tilingData->blockFactor * AscendC::GetBlockIdx(), blockLength_);
92+ 
93+ pipe.InitBuffer(inputQueue, BUFFER_NUM, ubLength_ * sizeof(T));
94+ pipe.InitBuffer(outputQueue, BUFFER_NUM, ubLength_ * sizeof(T));
95+ // Temporary buffers for intermediate computation (exp results)
96+ // For float32: used as float buffers directly
97+ // For float16: used as float32 buffers for precision-promoted computation
98+ pipe.InitBuffer(tmpBuf1, ubLength_ * sizeof(float));
99+ pipe.InitBuffer(tmpBuf2, ubLength_ * sizeof(float));
100+}
101+ 
102+// ============================================================================
103+// CopyIn - Transfer data from GM to UB
104+// ============================================================================
105+template <typename T>
106+__aicore__ inline void SinhV2<T>::CopyIn(int64_t progress, int64_t currentNum)
107+{
108+ AscendC::LocalTensor<T> xLocal = inputQueue.AllocTensor<T>();
109+ AscendC::DataCopyParams copyParams;
110+ copyParams.blockCount = 1;
111+ copyParams.blockLen = currentNum * sizeof(T);
112+ copyParams.srcStride = 0;
113+ copyParams.dstStride = 0;
114+ AscendC::DataCopyPad(xLocal, inputGM[progress * ubLength_], copyParams, {false, 0, 0, 0});
115+ inputQueue.EnQue(xLocal);
116+}
117+ 
118+// ============================================================================
119+// CopyOut - Transfer data from UB to GM
120+// ============================================================================
121+template <typename T>
122+__aicore__ inline void SinhV2<T>::CopyOut(int64_t progress, int64_t currentNum)
123+{
124+ AscendC::LocalTensor<T> yLocal = outputQueue.DeQue<T>();
125+ AscendC::DataCopyParams copyParams;
126+ copyParams.blockCount = 1;
127+ copyParams.blockLen = currentNum * sizeof(T);
128+ copyParams.srcStride = 0;
129+ copyParams.dstStride = 0;
130+ AscendC::DataCopyPad(outputGM[progress * ubLength_], yLocal, copyParams);
131+ outputQueue.FreeTensor(yLocal);
132+}
133+ 
134+// ============================================================================
135+// Compute - float32 specialization: direct sinh computation
136+// sinh(x) = (exp(x) - exp(-x)) / 2
137+// ============================================================================
138+template <>
139+__aicore__ inline void SinhV2<float>::Compute(int64_t currentNum)
140+{
141+ AscendC::LocalTensor<float> xLocal = inputQueue.DeQue<float>();
142+ AscendC::LocalTensor<float> yLocal = outputQueue.AllocTensor<float>();
143+ AscendC::LocalTensor<float> expPos = tmpBuf1.Get<float>();
144+ AscendC::LocalTensor<float> expNeg = tmpBuf2.Get<float>();
145+ 
146+ // Step 1: exp(x)
147+ AscendC::Exp(expPos, xLocal, currentNum);
148+ // Step 2: -x
149+ AscendC::Muls(xLocal, xLocal, static_cast<float>(-1.0f), currentNum);
150+ // Step 3: exp(-x)
151+ AscendC::Exp(expNeg, xLocal, currentNum);
152+ // Step 4: exp(x) - exp(-x)
153+ AscendC::Sub(expPos, expPos, expNeg, currentNum);
154+ // Step 5: (exp(x) - exp(-x)) * 0.5
155+ AscendC::Muls(yLocal, expPos, static_cast<float>(0.5f), currentNum);
156+ 
157+ outputQueue.EnQue<float>(yLocal);
158+ inputQueue.FreeTensor(xLocal);
159+}
160+ 
161+// ============================================================================
162+// Compute - float16 specialization: cast to float32, compute, cast back
163+// ============================================================================
164+template <>
165+__aicore__ inline void SinhV2<half>::Compute(int64_t currentNum)
166+{
167+ AscendC::LocalTensor<half> xLocal = inputQueue.DeQue<half>();
168+ AscendC::LocalTensor<half> yLocal = outputQueue.AllocTensor<half>();
169+ AscendC::LocalTensor<float> xFloat = tmpBuf1.Get<float>();
170+ AscendC::LocalTensor<float> expNeg = tmpBuf2.Get<float>();
171+ 
172+ // Step 1: Cast half -> float for precision
173+ AscendC::Cast(xFloat, xLocal, AscendC::RoundMode::CAST_NONE, currentNum);
174+ // Step 2: exp(x) - store in xLocal's position (we reuse xFloat after saving exp result)
175+ // We need expPos, so let's compute exp(x) into expNeg first as temp, then swap
176+ // Actually: compute exp(x) into expNeg (as temp for expPos), then compute -x, then exp(-x)
177+ AscendC::Exp(expNeg, xFloat, currentNum); // expNeg temporarily holds exp(x)
178+ // Step 3: -x
179+ AscendC::Muls(xFloat, xFloat, static_cast<float>(-1.0f), currentNum);
180+ // Step 4: exp(-x) - store back into xFloat (reuse)
181+ AscendC::Exp(xFloat, xFloat, currentNum); // xFloat now holds exp(-x)
182+ // Step 5: exp(x) - exp(-x) : expNeg has exp(x), xFloat has exp(-x)
183+ AscendC::Sub(expNeg, expNeg, xFloat, currentNum);
184+ // Step 6: * 0.5
185+ AscendC::Muls(xFloat, expNeg, static_cast<float>(0.5f), currentNum);
186+ // Step 7: Cast float -> half
187+ AscendC::Cast(yLocal, xFloat, AscendC::RoundMode::CAST_ROUND, currentNum);
188+ 
189+ outputQueue.EnQue<half>(yLocal);
190+ inputQueue.FreeTensor(xLocal);
191+}
192+ 
193+// ============================================================================
194+// Process - Main loop: iterate over tiles
195+// ============================================================================
196+template <typename T>
197+__aicore__ inline void SinhV2<T>::Process()
198+{
199+ // Guard: empty tensor or idle core - nothing to process
200+ if (blockLength_ <= 0 || ubLength_ <= 0) {
201+ return;
202+ }
203+ int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
204+ for (int64_t i = 0; i < loopCount; i++) {
205+ int64_t currentNum = (i == (loopCount - 1)) ? (blockLength_ - ubLength_ * i) : ubLength_;
206+ CopyIn(i, currentNum);
207+ Compute(currentNum);
208+ CopyOut(i, currentNum);
209+ }
210+}
211+ 
212+} // namespace NsSinhV2
213+#endif // SINH_V2_H
Aexperimental/math/sinh_v2/op_kernel/sinh_v2_tiling_data.h+31-0
@@ -0,0 +1,31 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file sinh_v2_tiling_data.h
19+ * \brief SinhV2 TilingData structure definition
20+ */
21+ 
22+#ifndef _SINH_V2_TILING_DATA_H_
23+#define _SINH_V2_TILING_DATA_H_
24+ 
25+struct SinhV2TilingData {
26+ int64_t totalNum = 0; // Total number of elements
27+ int64_t blockFactor = 0; // Number of elements per AI Core
28+ int64_t ubFactor = 0; // Number of elements per UB loop iteration
29+};
30+ 
31+#endif
Aexperimental/math/sinh_v2/op_kernel/sinh_v2_tiling_key.h+43-0
@@ -0,0 +1,43 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file sinh_v2_tiling_key.h
19+ * \brief SinhV2 TilingKey definition
20+ *
21+ * TilingKey mapping:
22+ * - SINHV2_TPL_SCH_MODE_0 (0): FLOAT32 type
23+ * - SINHV2_TPL_SCH_MODE_1 (1): FLOAT16 type
24+ */
25+ 
26+#ifndef __SINH_V2_TILING_KEY_H__
27+#define __SINH_V2_TILING_KEY_H__
28+ 
29+#include "ascendc/host_api/tiling/template_argument.h"
30+ 
31+#define SINHV2_TPL_SCH_MODE_0 0 // FLOAT32 type
32+#define SINHV2_TPL_SCH_MODE_1 1 // FLOAT16 type
33+ 
34+ASCENDC_TPL_ARGS_DECL(
35+ SinhV2,
36+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST,
37+ SINHV2_TPL_SCH_MODE_0, SINHV2_TPL_SCH_MODE_1));
38+ 
39+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
40+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST,
41+ SINHV2_TPL_SCH_MODE_0, SINHV2_TPL_SCH_MODE_1)));
42+ 
43+#endif
Aexperimental/math/sinh_v2/tests/.gitkeep+0-0
The file is empty
Aexperimental/math/tan_v3/CMakeLists.txt+24-0
@@ -0,0 +1,24 @@
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+ 
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,
8+# INCLUDING 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+#
13+# NOTE: Portions of this code were AI-generated and have been
14+# technically reviewed for functional accuracy and security
15+ 
16+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
17+if(NOT ENABLE_TEST)
18+ list(REMOVE_ITEM CURRENT_DIRS tests)
19+endif()
20+foreach(SUB_DIR ${CURRENT_DIRS})
21+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
22+ add_subdirectory(${SUB_DIR})
23+ endif()
24+endforeach()
Aexperimental/math/tan_v3/examples/test_aclnn_tan_v3.cpp+130-0
@@ -0,0 +1,130 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+#include <iostream>
18+#include <vector>
19+#include "acl/acl.h"
20+#include "aclnnop/aclnn_tan_v3.h"
21+ 
22+#define CHECK_RET(cond, return_expr) \
23+ do { \
24+ if (!(cond)) { \
25+ return_expr; \
26+ } \
27+ } while (0)
28+ 
29+#define LOG_PRINT(message, ...) \
30+ do { \
31+ printf(message, ##__VA_ARGS__); \
32+ } while (0)
33+ 
34+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
35+ int64_t shape_size = 1;
36+ for (auto i : shape) {
37+ shape_size *= i;
38+ }
39+ return shape_size;
40+}
41+ 
42+int Init(int32_t deviceId, aclrtStream* stream) {
43+ auto ret = aclInit(nullptr);
44+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
45+ ret = aclrtSetDevice(deviceId);
46+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
47+ ret = aclrtCreateStream(stream);
48+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
49+ return 0;
50+}
51+ 
52+template <typename T>
53+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
54+ aclDataType dataType, aclTensor** tensor) {
55+ auto size = GetShapeSize(shape) * sizeof(T);
56+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
57+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
58+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
59+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
60+ 
61+ std::vector<int64_t> strides(shape.size(), 1);
62+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
63+ strides[i] = shape[i + 1] * strides[i + 1];
64+ }
65+ 
66+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
67+ shape.data(), shape.size(), *deviceAddr);
68+ return 0;
69+}
70+ 
71+int main() {
72+ int32_t deviceId = 0;
73+ aclrtStream stream;
74+ auto ret = Init(deviceId, &stream);
75+ CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
76+ 
77+ std::vector<int64_t> xShape = {4, 2};
78+ std::vector<int64_t> outShape = {4, 2};
79+ void* xDeviceAddr = nullptr;
80+ void* outDeviceAddr = nullptr;
81+ aclTensor* x = nullptr;
82+ aclTensor* out = nullptr;
83+ std::vector<float> xHostData = {1, 2, 3, 4, 5, 6, 7, 8};
84+ std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0};
85+ 
86+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_FLOAT, &x);
87+ CHECK_RET(ret == ACL_SUCCESS, return ret);
88+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
89+ CHECK_RET(ret == ACL_SUCCESS, return ret);
90+ 
91+ uint64_t workspaceSize = 0;
92+ aclOpExecutor* executor = nullptr;
93+ ret = aclnnTanV3GetWorkspaceSize(x, out, &workspaceSize, &executor);
94+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnTanV3GetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
95+ 
96+ void* workspaceAddr = nullptr;
97+ if (workspaceSize > 0) {
98+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
99+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret;);
100+ }
101+ 
102+ ret = aclnnTanV3(workspaceAddr, workspaceSize, executor, stream);
103+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnTanV3 failed. ERROR: %d\n", ret); return ret);
104+ 
105+ ret = aclrtSynchronizeStream(stream);
106+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
107+ 
108+ auto size = GetShapeSize(outShape);
109+ std::vector<float> resultData(size, 0);
110+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(float),
111+ ACL_MEMCPY_DEVICE_TO_HOST);
112+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
113+ 
114+ for (int64_t i = 0; i < size; i++) {
115+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
116+ }
117+ 
118+ aclDestroyTensor(x);
119+ aclDestroyTensor(out);
120+ 
121+ aclrtFree(xDeviceAddr);
122+ aclrtFree(outDeviceAddr);
123+ if (workspaceSize > 0) {
124+ aclrtFree(workspaceAddr);
125+ }
126+ aclrtDestroyStream(stream);
127+ aclrtResetDevice(deviceId);
128+ aclFinalize();
129+ return 0;
130+}
Aexperimental/math/tan_v3/op_host/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
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+ 
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,
8+# INCLUDING 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+#
13+# NOTE: Portions of this code were AI-generated and have been
14+# technically reviewed for functional accuracy and security
15+ 
16+add_modules_sources(OPTYPE tan_v3 ACLNNTYPE aclnn)
Aexperimental/math/tan_v3/op_host/tan_v3_def.cpp+53-0
@@ -0,0 +1,53 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file tan_v3_def.cpp
19+ * \brief TanV3 operator definition - declares inputs, outputs, and chip configuration
20+ */
21+#include "register/op_def_registry.h"
22+ 
23+namespace ops {
24+class TanV3 : public OpDef {
25+public:
26+ explicit TanV3(const char* name) : OpDef(name)
27+ {
28+ this->Input("x")
29+ .ParamType(REQUIRED)
30+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
CANN-robot
CANN-robotCANN-robot3月28日
代码健壮性与防御性编程: 输入和输出的数据类型（DT_FLOAT, DT_FLOAT16）和格式（FORMAT_ND）是通过初始化列表硬编码的。虽然这是算子定义阶段的常见做法，但缺乏对列表长度一致性的运行时检查或静态断言。如果未来修改时，DataType列表和Format列表的长度不匹配，可能导致未定义行为。
问题类型: 代码健壮性与防御性编程
文件路径: experimental/math/tan_v3/op_host/tan_v3_def.cpp
行号: 31
问题代码:
.DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
.Format({ge::FORMAT_ND, ge::FORMAT_ND})
修改建议:
1. 确保框架层（OpDef类）的Input/Output方法内部有对列表长度一致性的校验。2. 在本代码层面，可以考虑使用静态断言（如果C++版本支持）或在添加每个数据类型时显式关联其默认格式，以降低出错风险。
---
此评论由代码审查工具自动生成
likedislike
31+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
32+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
33+ .AutoContiguous();
34+ this->Output("y")
35+ .ParamType(REQUIRED)
36+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
37+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
38+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
39+ .AutoContiguous();
40+ 
41+ OpAICoreConfig aicoreConfig910B;
42+ aicoreConfig910B.DynamicCompileStaticFlag(true)
43+ .DynamicFormatFlag(false)
44+ .DynamicRankSupportFlag(true)
45+ .DynamicShapeSupportFlag(true)
46+ .NeedCheckSupportFlag(false)
47+ .PrecisionReduceFlag(true)
48+ .ExtendCfgInfo("opFile.value", "tan_v3");
49+ this->AICore().AddConfig("ascend910b", aicoreConfig910B);
50+ }
51+};
52+OP_ADD(TanV3);
53+} // namespace ops
Aexperimental/math/tan_v3/op_host/tan_v3_infershape.cpp+48-0
@@ -0,0 +1,48 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file tan_v3_infershape.cpp
19+ * \brief TanV3 shape inference - output shape equals input shape
20+ */
21+ 
22+#include "register/op_impl_registry.h"
23+#include "exe_graph/runtime/infer_shape_context.h"
24+ 
25+using namespace ge;
26+ 
27+namespace ops {
28+ 
29+static ge::graphStatus InferShape4TanV3(gert::InferShapeContext* context)
30+{
31+ const gert::Shape* input_shape = context->GetInputShape(0);
32+ if (input_shape == nullptr) {
33+ return ge::GRAPH_FAILED;
34+ }
35+ 
36+ gert::Shape* output_shape = context->GetOutputShape(0);
37+ if (output_shape == nullptr) {
38+ return ge::GRAPH_FAILED;
39+ }
40+ 
41+ *output_shape = *input_shape;
42+ 
43+ return ge::GRAPH_SUCCESS;
44+}
45+ 
46+IMPL_OP_INFERSHAPE(TanV3).InferShape(InferShape4TanV3);
47+ 
48+} // namespace ops
Aexperimental/math/tan_v3/op_host/tan_v3_tiling.cpp+158-0
@@ -0,0 +1,158 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file tan_v3_tiling.cpp
19+ * \brief Tanv3 Tiling implementation (arch32 - Ascend910B)
20+ *
21+ * Computes tiling parameters for Tanv3 operator:
22+ * - Multi-core splitting: totalNum divided evenly across AI Cores
23+ * - UB splitting: each core processes data in chunks of ubFactor elements
24+ * - TilingKey selection: based on input dtype (float32 vs float16)
25+ */
26+ 
27+#include "register/op_def_registry.h"
28+#include "op_common/log/log.h"
29+#include "op_common/op_host/util/math_util.h"
30+#include "op_common/op_host/util/platform_util.h"
31+#include "../op_kernel/tan_v3_tiling_data.h"
32+#include "../op_kernel/tan_v3_tiling_key.h"
33+ 
34+namespace optiling {
35+ 
36+using Ops::Base::CeilDiv;
37+using Ops::Base::FloorDiv;
38+using Ops::Base::FloorAlign;
39+using Ops::Base::GetUbBlockSize;
40+ 
41+constexpr uint32_t WS_SYS_SIZE = 0U;
42+// float32: inputQueue(x2) + outputQueue(x2) + tmpBuf1(x1) + tmpBuf2(x1) = 6 float-sized buffers
43+// Total bytes = ubFactor * 6 * sizeof(float)
44+constexpr int64_t BUFFER_NUM_FP32 = 6;
45+// float16: inputQueue(x2,half) + outputQueue(x2,half) + tmpBuf1(x1,float) + tmpBuf2(x1,float)
46+// Total bytes = ubFactor * (2*2 + 2*2 + 4 + 4) = ubFactor * 16 = ubFactor * sizeof(float) * 4
47+constexpr int64_t BUFFER_NUM_FP16 = 4;
48+ 
49+static const gert::Shape g_vec_1_shape = {1};
50+ 
51+static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape)
52+{
53+ if (in_shape.GetDimNum() == 0) {
54+ return g_vec_1_shape;
55+ }
56+ return in_shape;
57+}
58+ 
59+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
60+{
61+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
62+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
63+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
64+ coreNum = ascendcPlatform.GetCoreNumAiv();
65+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
66+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
67+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
68+ return ge::GRAPH_SUCCESS;
69+}
70+ 
71+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
72+{
73+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
74+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
75+ currentWorkspace[0] = WS_SYS_SIZE;
76+ return ge::GRAPH_SUCCESS;
77+}
78+ 
79+static ge::graphStatus TanV3TilingFunc(gert::TilingContext* context)
80+{
81+ // 1. Get platform info
82+ uint64_t ubSize;
83+ int64_t coreNum;
84+ OP_CHECK_IF(
85+ GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
86+ OP_LOGE(context, "GetPlatformInfo error"),
87+ return ge::GRAPH_FAILED);
88+ 
89+ // 2. Get input shape and dtype
90+ auto inputShape = context->GetInputShape(0);
91+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
92+ auto storageShape = EnsureNotScalar(inputShape->GetStorageShape());
93+ int64_t totalNum = storageShape.GetShapeSize();
94+ 
95+ auto inputDesc = context->GetInputDesc(0);
96+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
97+ ge::DataType dtype = inputDesc->GetDataType();
98+ 
99+ // 3. Get workspace size
100+ OP_CHECK_IF(
101+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
102+ OP_LOGE(context, "GetWorkspaceSize error"),
103+ return ge::GRAPH_FAILED);
104+ 
105+ // 4. Set TilingData
106+ TanV3TilingData* tiling = context->GetTilingData<TanV3TilingData>();
107+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
108+ OP_CHECK_IF(
109+ memset_s(tiling, sizeof(TanV3TilingData), 0, sizeof(TanV3TilingData)) != EOK,
110+ OP_LOGE(context, "set tiling data error"),
111+ return ge::GRAPH_FAILED);
112+ 
113+ // Handle empty tensor
114+ if (totalNum == 0) {
115+ tiling->totalNum = 0;
116+ tiling->blockFactor = 0;
117+ tiling->ubFactor = 0;
118+ context->SetBlockDim(1);
119+ context->SetTilingKey(GET_TPL_TILING_KEY(TANV3_TPL_SCH_MODE_0));
120+ return ge::GRAPH_SUCCESS;
121+ }
122+ 
123+ // 5. Multi-core splitting
124+ tiling->totalNum = totalNum;
125+ tiling->blockFactor = CeilDiv(totalNum, coreNum);
126+ int64_t usedCoreNum = CeilDiv(totalNum, tiling->blockFactor);
127+ 
128+ // 6. UB splitting and TilingKey selection
129+ int64_t ubCanUse = static_cast<int64_t>(ubSize);
130+ int64_t ubBlockSize = GetUbBlockSize(context);
131+ // Both paths compute internally in float32, so typeSize = 4
132+ constexpr int64_t typeSize = 4;
133+ 
134+ if (dtype == ge::DT_FLOAT) {
135+ tiling->ubFactor = FloorAlign(FloorDiv((ubCanUse / typeSize), BUFFER_NUM_FP32), ubBlockSize);
136+ context->SetTilingKey(GET_TPL_TILING_KEY(TANV3_TPL_SCH_MODE_0));
137+ } else if (dtype == ge::DT_FLOAT16) {
138+ tiling->ubFactor = FloorAlign(FloorDiv((ubCanUse / typeSize), BUFFER_NUM_FP16), ubBlockSize);
139+ context->SetTilingKey(GET_TPL_TILING_KEY(TANV3_TPL_SCH_MODE_1));
140+ } else {
141+ OP_LOGE(context, "TanV3: unsupported dtype");
142+ return ge::GRAPH_FAILED;
143+ }
144+ 
145+ context->SetBlockDim(usedCoreNum);
146+ return ge::GRAPH_SUCCESS;
147+}
148+ 
149+static ge::graphStatus TilingParseForTanV3([[maybe_unused]] gert::TilingParseContext* context)
150+{
151+ return ge::GRAPH_SUCCESS;
152+}
153+ 
154+struct TanV3CompileInfo {};
155+ 
156+IMPL_OP_OPTILING(TanV3).Tiling(TanV3TilingFunc).TilingParse<TanV3CompileInfo>(TilingParseForTanV3);
157+ 
158+} // namespace optiling
Aexperimental/math/tan_v3/op_kernel/tan_v3.cpp+45-0
@@ -0,0 +1,45 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file tan_v3_arch32.cpp
19+ * \brief TanV3 kernel entry point (arch32 architecture - Ascend910B)
20+ */
21+ 
22+#include "tan_v3.h"
23+ 
24+enum class TanV3TilingKey : uint32_t
25+{
26+ TILING_KEY_FLOAT32 = 0,
27+ TILING_KEY_FLOAT16 = 1,
28+};
29+ 
30+template <uint32_t schMode>
31+__global__ __aicore__ void tan_v3(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
32+{
33+ REGISTER_TILING_DEFAULT(TanV3TilingData);
34+ GET_TILING_DATA_WITH_STRUCT(TanV3TilingData, tilingData, tiling);
35+ if constexpr (schMode == static_cast<uint32_t>(TanV3TilingKey::TILING_KEY_FLOAT32)) {
36+ NsTanV3::TanV3<float> op;
37+ op.Init(x, y, &tilingData);
38+ op.Process();
39+ }
40+ if constexpr (schMode == static_cast<uint32_t>(TanV3TilingKey::TILING_KEY_FLOAT16)) {
41+ NsTanV3::TanV3<half> op;
42+ op.Init(x, y, &tilingData);
43+ op.Process();
44+ }
45+}
Aexperimental/math/tan_v3/op_kernel/tan_v3.h+205-0
@@ -0,0 +1,205 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file tan_v3.h
19+ * \brief TanV3 kernel class definition (arch32 - Ascend910B)
20+ *
21+ * Computes tan(x) = sin(x) / cos(x) for each element.
22+ * - float32 path: direct computation using Sin, Cos, Div
23+ * - float16 path: cast to float32 for computation, then cast back
24+ *
25+ * Uses double buffering (BUFFER_NUM=2) for pipeline parallelism:
26+ * - CopyIn: GM -> UB (input data transfer)
27+ * - Compute: UB vector computation (tan formula)
28+ * - CopyOut: UB -> GM (output data transfer)
29+ */
30+#ifndef TAN_V3_H
31+#define TAN_V3_H
32+ 
33+#include "kernel_operator.h"
34+#include "kernel_tiling/kernel_tiling.h"
35+#include "tan_v3_tiling_data.h"
36+#include "tan_v3_tiling_key.h"
37+ 
38+namespace NsTanV3 {
39+ 
40+using namespace AscendC;
41+ 
42+constexpr int32_t BUFFER_NUM = 2;
43+ 
44+template <typename T>
45+class TanV3 {
46+public:
47+ __aicore__ inline TanV3() {};
48+ 
49+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const TanV3TilingData* tilingData);
50+ __aicore__ inline void Process();
51+ 
52+private:
53+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum);
54+ __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum);
55+ __aicore__ inline void Compute(int64_t currentNum);
56+ 
57+private:
58+ TPipe pipe;
59+ TQue<QuePosition::VECIN, BUFFER_NUM> inputQueue;
60+ TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueue;
61+ TBuf<QuePosition::VECCALC> tmpBuf1; // stores sin(x) intermediate result
62+ TBuf<QuePosition::VECCALC> tmpBuf2; // stores cos(x) intermediate result
63+ 
64+ GlobalTensor<T> inputGM;
65+ GlobalTensor<T> outputGM;
66+ 
67+ int64_t blockLength_ = 0;
68+ int64_t ubLength_ = 0;
69+};
70+ 
71+// ============================================================================
72+// Init - Initialize GM pointers and allocate UB buffers
73+// ============================================================================
74+template <typename T>
75+__aicore__ inline void TanV3<T>::Init(GM_ADDR x, GM_ADDR y, const TanV3TilingData* tilingData)
76+{
77+ int64_t remainderLength = tilingData->totalNum - tilingData->blockFactor * AscendC::GetBlockIdx();
78+ blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength;
79+ // Clamp to 0 for idle cores (when totalNum < coreNum, some cores have no work)
80+ if (blockLength_ < 0) {
81+ blockLength_ = 0;
82+ }
83+ ubLength_ = tilingData->ubFactor;
84+ 
85+ // Guard: empty tensor or idle core - skip buffer allocation
86+ if (blockLength_ <= 0 || ubLength_ <= 0) {
87+ return;
88+ }
89+ 
90+ inputGM.SetGlobalBuffer((__gm__ T*)x + tilingData->blockFactor * AscendC::GetBlockIdx(), blockLength_);
91+ outputGM.SetGlobalBuffer((__gm__ T*)y + tilingData->blockFactor * AscendC::GetBlockIdx(), blockLength_);
92+ 
93+ pipe.InitBuffer(inputQueue, BUFFER_NUM, ubLength_ * sizeof(T));
94+ pipe.InitBuffer(outputQueue, BUFFER_NUM, ubLength_ * sizeof(T));
95+ // Temporary buffers for sin and cos intermediate results
96+ // For float32: used as float buffers directly
97+ // For float16: used as float32 buffers for precision-promoted computation
98+ pipe.InitBuffer(tmpBuf1, ubLength_ * sizeof(float));
99+ pipe.InitBuffer(tmpBuf2, ubLength_ * sizeof(float));
100+}
101+ 
102+// ============================================================================
103+// CopyIn - Transfer data from GM to UB
104+// ============================================================================
105+template <typename T>
106+__aicore__ inline void TanV3<T>::CopyIn(int64_t progress, int64_t currentNum)
107+{
108+ AscendC::LocalTensor<T> xLocal = inputQueue.AllocTensor<T>();
109+ AscendC::DataCopyParams copyParams;
110+ copyParams.blockCount = 1;
111+ copyParams.blockLen = currentNum * sizeof(T);
112+ copyParams.srcStride = 0;
113+ copyParams.dstStride = 0;
114+ AscendC::DataCopyPad(xLocal, inputGM[progress * ubLength_], copyParams, {false, 0, 0, 0});
115+ inputQueue.EnQue(xLocal);
116+}
117+ 
118+// ============================================================================
119+// CopyOut - Transfer data from UB to GM
120+// ============================================================================
121+template <typename T>
122+__aicore__ inline void TanV3<T>::CopyOut(int64_t progress, int64_t currentNum)
123+{
124+ AscendC::LocalTensor<T> yLocal = outputQueue.DeQue<T>();
125+ AscendC::DataCopyParams copyParams;
126+ copyParams.blockCount = 1;
127+ copyParams.blockLen = currentNum * sizeof(T);
128+ copyParams.srcStride = 0;
129+ copyParams.dstStride = 0;
130+ AscendC::DataCopyPad(outputGM[progress * ubLength_], yLocal, copyParams);
131+ outputQueue.FreeTensor(yLocal);
132+}
133+ 
134+// ============================================================================
135+// Compute - float32 specialization: direct tan computation
136+// tan(x) = sin(x) / cos(x)
137+// ============================================================================
138+template <>
139+__aicore__ inline void TanV3<float>::Compute(int64_t currentNum)
140+{
141+ AscendC::LocalTensor<float> xLocal = inputQueue.DeQue<float>();
142+ AscendC::LocalTensor<float> yLocal = outputQueue.AllocTensor<float>();
143+ AscendC::LocalTensor<float> sinVal = tmpBuf1.Get<float>();
144+ AscendC::LocalTensor<float> cosVal = tmpBuf2.Get<float>();
145+ 
146+ // Step 1: sin(x) -> sinVal (tmpBuf1)
147+ AscendC::Sin(sinVal, xLocal, currentNum);
148+ // Step 2: cos(x) -> cosVal (tmpBuf2)
149+ AscendC::Cos(cosVal, xLocal, currentNum);
150+ // Step 3: sin(x) / cos(x) -> yLocal
151+ AscendC::Div(yLocal, sinVal, cosVal, currentNum);
152+ 
153+ outputQueue.EnQue<float>(yLocal);
154+ inputQueue.FreeTensor(xLocal);
155+}
156+ 
157+// ============================================================================
158+// Compute - float16 specialization: cast to float32, compute, cast back
159+// Flow: Cast(half->float) -> Cos -> Sin -> Div -> Cast(float->half)
160+// Key constraint: Cos must be computed before Sin to avoid overwriting input
161+// ============================================================================
162+template <>
163+__aicore__ inline void TanV3<half>::Compute(int64_t currentNum)
164+{
165+ AscendC::LocalTensor<half> xLocal = inputQueue.DeQue<half>();
166+ AscendC::LocalTensor<half> yLocal = outputQueue.AllocTensor<half>();
167+ AscendC::LocalTensor<float> sinVal = tmpBuf1.Get<float>();
168+ AscendC::LocalTensor<float> cosVal = tmpBuf2.Get<float>();
169+ 
170+ // Step 1: Cast half -> float (store in sinVal as temp for x_float)
171+ AscendC::Cast(sinVal, xLocal, AscendC::RoundMode::CAST_NONE, currentNum);
172+ // Step 2: cos(x) - must compute before sin overwrites sinVal
173+ AscendC::Cos(cosVal, sinVal, currentNum);
174+ // Step 3: sin(x) - overwrites sinVal (cosVal already saved)
175+ AscendC::Sin(sinVal, sinVal, currentNum);
176+ // Step 4: sin(x) / cos(x) -> sinVal (reuse)
177+ AscendC::Div(sinVal, sinVal, cosVal, currentNum);
178+ // Step 5: Cast float -> half
179+ AscendC::Cast(yLocal, sinVal, AscendC::RoundMode::CAST_ROUND, currentNum);
180+ 
181+ outputQueue.EnQue<half>(yLocal);
182+ inputQueue.FreeTensor(xLocal);
183+}
184+ 
185+// ============================================================================
186+// Process - Main loop: iterate over tiles
187+// ============================================================================
188+template <typename T>
189+__aicore__ inline void TanV3<T>::Process()
190+{
191+ // Guard: empty tensor or idle core - nothing to process
192+ if (blockLength_ <= 0 || ubLength_ <= 0) {
193+ return;
194+ }
195+ int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
196+ for (int64_t i = 0; i < loopCount; i++) {
197+ int64_t currentNum = (i == (loopCount - 1)) ? (blockLength_ - ubLength_ * i) : ubLength_;
198+ CopyIn(i, currentNum);
199+ Compute(currentNum);
200+ CopyOut(i, currentNum);
201+ }
202+}
203+ 
204+} // namespace NsTanV3
205+#endif // TAN_V3_H
Aexperimental/math/tan_v3/op_kernel/tan_v3_tiling_data.h+31-0
@@ -0,0 +1,31 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file tan_v3_tiling_data.h
19+ * \brief TanV3 TilingData structure definition
20+ */
21+ 
22+#ifndef _TAN_V3_TILING_DATA_H_
23+#define _TAN_V3_TILING_DATA_H_
24+ 
25+struct TanV3TilingData {
26+ int64_t totalNum = 0; // Total number of elements
27+ int64_t blockFactor = 0; // Number of elements per AI Core
28+ int64_t ubFactor = 0; // Number of elements per UB loop iteration
29+};
30+ 
31+#endif
Aexperimental/math/tan_v3/op_kernel/tan_v3_tiling_key.h+43-0
@@ -0,0 +1,43 @@
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+
11+/**
12+*
13+* NOTE: Portions of this code were AI-generated and have been
14+* technically reviewed for functional accuracy and security
15+*/
16+ 
17+/**
18+ * \file tan_v3_tiling_key.h
19+ * \brief TanV3 TilingKey definition
20+ *
21+ * TilingKey mapping:
22+ * - TANV3_TPL_SCH_MODE_0 (0): FLOAT32 type
23+ * - TANV3_TPL_SCH_MODE_1 (1): FLOAT16 type
24+ */
25+ 
26+#ifndef __TAN_V3_TILING_KEY_H__
27+#define __TAN_V3_TILING_KEY_H__
28+ 
29+#include "ascendc/host_api/tiling/template_argument.h"
30+ 
31+#define TANV3_TPL_SCH_MODE_0 0 // FLOAT32 type
32+#define TANV3_TPL_SCH_MODE_1 1 // FLOAT16 type
33+ 
34+ASCENDC_TPL_ARGS_DECL(
35+ TanV3,
36+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST,
37+ TANV3_TPL_SCH_MODE_0, TANV3_TPL_SCH_MODE_1));
38+ 
39+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
40+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST,
41+ TANV3_TPL_SCH_MODE_0, TANV3_TPL_SCH_MODE_1)));
42+ 
43+#endif
Aexperimental/math/tan_v3/tests/.gitkeep+0-0
The file is empty