Star1314
1432
代码
代码
Issues123
Pull Requests253
Actions
项目讨论
Wiki
分析
Star1314
1432
已合并
tan算子支持Ascend910b AscendC实现 #2132
wangweidong创建于 4月8日
tan算子支持Ascend910b AscendC实现 #2132
已合并
wangweidong创建于 4月8日
从gcw_YBIAEfqJ/ops-mathcann_0409合入到cann/ops-mathmaster
共 14 个文件变更+1206-0
Aexperimental/math/tan/CMakeLists.txt+22-0
@@ -0,0 +1,22 @@
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+ # NOTE: Portions of this code were AI-generated and have been
12+ # technically reviewed for functional accuracy and security
13+ 
14+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
15+if(NOT ENABLE_TEST)
16+ list(REMOVE_ITEM CURRENT_DIRS tests)
17+endif()
18+foreach(SUB_DIR ${CURRENT_DIRS})
19+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
20+ add_subdirectory(${SUB_DIR})
21+ endif()
22+endforeach()
Aexperimental/math/tan/README.md+242-0
@@ -0,0 +1,242 @@
1+# Tan 自定义算子
2+ 
3+## 功能说明
4+ 
5+Tan 算子计算输入张量的逐元素正切值：
6+ 
7+```
8+tan(x) = sin(x) / cos(x)
9+```
10+ 
11+**对标**：PyTorch `torch.tan`
12+ 
13+## 支持规格
14+ 
15+### 数据类型
16+ 
17+| 输入 x | 输出 out |
18+|--------|---------|
19+| float32 | float32 |
20+| float16 | float16 |
21+ 
22+注意：输出 out 的数据类型与输入 x 一致。
23+ 
24+### Shape
25+ 
26+- 支持任意维度（标量、1D ~ 8D 均已验证）
27+- 输出 shape 与输入 shape 相同
28+- 不支持广播（单输入算子）
29+- 支持空 tensor（元素数为 0）
30+- 支持动态 shape / 动态 rank
31+ 
32+### 目标芯片
33+ 
34+- Ascend 910B3（arch32 架构）
35+ 
36+## 使用方法
37+ 
38+### aclnn API 调用
39+ 
40+```cpp
41+#include "acl/acl.h"
42+#include "aclnn_tan.h"
43+ 
44+// 第一步：获取 workspace 大小和执行器
45+uint64_t workspaceSize = 0;
46+aclOpExecutor* executor = nullptr;
47+auto ret = aclnnTanGetWorkspaceSize(
48+ x_tensor, // 输入张量 x
49+ output_tensor, // 输出张量
50+ &workspaceSize, // 输出：workspace 大小
51+ &executor // 输出：执行器句柄
52+);
53+ 
54+// 第二步：分配 workspace 内存
55+void* workspace = nullptr;
56+if (workspaceSize > 0) {
57+ aclrtMalloc(&workspace, workspaceSize, ACL_MEM_MALLOC_NORMAL_ONLY);
58+}
59+ 
60+// 第三步：执行算子
61+ret = aclnnTan(workspace, workspaceSize, executor, stream);
62+ 
63+// 第四步：同步并释放资源
64+aclrtSynchronizeStream(stream);
65+if (workspace) aclrtFree(workspace);
66+```
67+ 
68+### 张量创建示例
69+ 
70+```cpp
71+// 创建 float32 张量，shape [2, 4]
72+std::vector<int64_t> shape = {2, 4};
73+std::vector<int64_t> strides = {4, 1};
74+aclTensor* tensor = aclCreateTensor(
75+ shape.data(), shape.size(),
76+ ACL_FLOAT, // 数据类型
77+ strides.data(), 0, // strides + offset
78+ ACL_FORMAT_ND, // 格式
79+ shape.data(), shape.size(),
80+ device_ptr // 设备内存指针
81+);
82+```
83+ 
84+### 完整示例参考
85+ 
86+- **aclnn 调用示例**：`examples/test_aclnn_tan.cpp`
87+- **GE IR 图模式示例**：`examples/test_geir_tan.cpp`
88+ 
89+运行示例：
90+ 
91+```bash
92+cd examples
93+bash run.sh # 运行 aclnn 调用示例（默认）
94+bash run.sh --graph # 运行图模式 (GE IR) 调用示例
95+```
96+ 
97+## 精度说明
98+ 
99+| 数据类型 | 精度标准 | 真实 NPU 表现 | 备注 |
100+|---------|---------|-------------|------|
101+| float32 | rtol=1e-4, atol=1e-6 | 100% 通过 | 全量 36 条用例 |
102+| float16 | rtol=1e-3, atol=1e-3 | 100% 通过 | 全量 20 条用例 |
103+ 
104+**说明**：float16 路径内部先 Cast 升至 float32 进行 Sin/Cos/Div 计算，再 Cast 回 float16，保证精度。
105+ 
106+## 构建方法
107+ 
108+### 前提条件
109+ 
110+- CANN Toolkit 已安装（路径：`/home/developer/Ascend/cann-9.0.0` 或其他版本）
111+- 已设置环境变量：`source /home/developer/Ascend/ascend-toolkit/set_env.sh`
112+ 
113+### 编译自定义算子包
114+ 
115+```bash
116+cd ops/tan
117+bash build.sh --soc=ascend910b --pkg
118+```
119+ 
120+编译成功后，算子包位于 `build/custom_opp_ubuntu_aarch64.run`。
121+ 
122+### 安装算子包
123+ 
124+```bash
125+bash build/custom_opp_ubuntu_aarch64.run
126+```
127+ 
128+算子安装到：`$ASCEND_HOME_PATH/opp/vendors/tan_custom/`
129+ 
130+## 测试方法
131+ 
132+### 运行 UT（单元测试）
133+ 
134+```bash
135+cd tests/ut
136+bash run.sh
137+```
138+ 
139+### 运行 ST（系统测试）
140+ 
141+```bash
142+cd tests/st
143+bash run.sh
144+```
145+ 
146+全量 56 条测试用例（36 条 float32 + 20 条 float16）。
147+ 
148+ST 测试支持两种模式：
149+- **Mock 模式**（CPU Golden）：无需 NPU，用于开发验证
150+- **真实 NPU 模式**：需要 NPU 设备，验证实际精度
151+ 
152+### 一键编译 + 测试
153+ 
154+```bash
155+bash build.sh --soc=ascend910b --pkg -a # 编译 + UT + ST
156+bash build.sh --soc=ascend910b --pkg -u # 编译 + 仅 UT
157+bash build.sh --soc=ascend910b --pkg -s # 编译 + 仅 ST
158+```
159+ 
160+## 目录结构
161+ 
162+```
163+ops/tan/
164+├── README.md # 本文档
165+├── CMakeLists.txt # 顶层构建脚本
166+├── build.sh # 编译脚本
167+├── op_host/ # Host 侧实现
168+│ ├── CMakeLists.txt
169+│ ├── tan_def.cpp # 算子定义（aclnn API 注册）
170+│ ├── tan_infershape.cpp # InferShape（输出 shape = 输入 shape）
171+│ └── arch32/
172+│ └── tan_tiling.cpp # Tiling 实现（多核切分 + UB 切分）
173+├── op_kernel/ # Device 侧实现
174+│ ├── CMakeLists.txt
175+│ ├── tan_arch32.cpp # Kernel 入口（模板分发）
176+│ └── arch32/
177+│ ├── tan.h # Kernel 实现（核心计算逻辑）
178+│ ├── tan_tiling_data.h # Tiling 数据结构
179+│ └── tan_tiling_key.h # Tiling Key 定义
180+├── op_graph/
181+│ └── tan_proto.h # GE IR 算子原型注册
182+├── examples/ # 调用示例
183+│ ├── CMakeLists.txt # aclnn 模式构建脚本
184+│ ├── CMakeLists_geir.txt # GE IR 模式构建脚本
185+│ ├── run.sh # 统一运行脚本（--eager / --graph）
186+│ ├── test_aclnn_tan.cpp # aclnn 调用示例
187+│ └── test_geir_tan.cpp # GE IR 图模式调用示例
188+├── tests/
189+│ ├── st/ # 系统测试
190+│ │ ├── test_aclnn_tan.cpp # ST 测试工程（56 条用例）
191+│ │ ├── CMakeLists.txt
192+│ │ └── run.sh # ST 运行脚本
193+│ └── ut/ # 单元测试
194+│ ├── run.sh # UT 运行脚本
195+│ ├── CMakeLists.txt
196+│ └── op_host/
197+│ ├── CMakeLists.txt
198+│ ├── test_op_host_main.cpp
199+│ ├── test_tan_infershape.cpp # InferShape UT
200+│ └── test_tan_tiling.cpp # Tiling UT
201+├── docs/
202+│ ├── aclnnTan.md # aclnn API 接口文档
203+│ ├── REQUIREMENT_ANALYSIS.md # 需求分析文档
204+│ ├── DETAILED_DESIGN.md # 详细设计文档
205+│ ├── TEST_DESIGN.md # 测试设计文档
206+│ ├── PRECISION_VERIFICATION_REPORT.md # 精度验收报告
207+│ └── DEVELOPMENT_LOG.md # 开发日志
208+└── build/ # 编译输出目录
209+ └── custom_opp_ubuntu_aarch64.run # 算子包
210+```
211+ 
212+## 实现说明
213+ 
214+### 核心计算步骤
215+ 
216+#### float32 路径
217+ 
218+```
219+sinVal = Sin(x) // 计算 sin(x)
220+cosVal = Cos(x) // 计算 cos(x)
221+y = Div(sinVal, cosVal) // tan(x) = sin(x) / cos(x)
222+```
223+ 
224+#### float16 路径（升精度计算）
225+ 
226+```
227+x_fp32 = Cast(x, CAST_NONE) // half -> float32
228+cosVal = Cos(x_fp32) // 先算 cos（避免输入被覆盖）
229+sinVal = Sin(x_fp32) // 再算 sin
230+result = Div(sinVal, cosVal) // tan = sin / cos
231+y = Cast(result, CAST_ROUND) // float32 -> half
232+```
233+ 
234+### Tiling 策略
235+ 
236+- **多核切分**：总元素数均匀分配到各 AI Core
237+ - `blockFactor = CeilDiv(totalNum, coreNum)`
238+ - `usedCoreNum = CeilDiv(totalNum, blockFactor)`
239+- **UB 切分**：每个 Core 内按 UB 容量分块处理
240+ - float32：`ubFactor = FloorAlign(ubCanUse / 4 / 6, ubBlockSize)`（6 块 buffer）
241+ - float16：`ubFactor = FloorAlign(ubCanUse / 4 / 4, ubBlockSize)`（4 块 buffer）
242+- **流水线**：双 buffer（BUFFER_NUM=2），CopyIn → Compute → CopyOut 三级流水
Aexperimental/math/tan/docs/aclnnTan.md+149-0
@@ -0,0 +1,149 @@
1+# aclnnTan
2+ 
3+## 支持的产品型号
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
8+ 
9+## 功能描述
10+ 
11+计算输入张量 `x` 的逐元素正切值，即 $out = \tan(x) = \frac{\sin(x)}{\cos(x)}$。
12+ 
13+- 不支持广播：`out` 的 shape 与 `x` 相同。
14+- 支持数据类型：float32、float16。
15+ 
16+## 函数原型
17+ 
18+```cpp
19+aclnnStatus aclnnTanGetWorkspaceSize(
20+ const aclTensor *x,
21+ const aclTensor *out,
22+ uint64_t *workspaceSize,
23+ aclOpExecutor **executor);
24+ 
25+aclnnStatus aclnnTan(
26+ void *workspace,
27+ uint64_t workspaceSize,
28+ aclOpExecutor *executor,
29+ aclrtStream stream);
30+```
31+ 
32+## aclnnTanGetWorkspaceSize
33+ 
34+### 参数说明
35+ 
36+| 参数名 | 输入/输出 | 描述 |
37+|-------|---------|------|
38+| x | 输入 | 数据类型：float32、float16。数据格式：ND。支持非连续 tensor。 |
39+| out | 输出 | 数据类型与 x 相同。数据格式：ND。shape 须与 x 相同。支持非连续 tensor。 |
40+| workspaceSize | 输出 | 算子执行所需 workspace 大小，单位为 Byte。由本函数返回，调用方须据此分配 workspace 内存。 |
41+| executor | 输出 | 算子执行器，包含算子计算流信息，由本函数返回后传入 aclnnTan 执行。 |
42+ 
43+### 返回值说明
44+ 
45+返回 `aclnnStatus` 错误码，详见 [aclnn 错误码](#错误码)。
46+ 
47+## aclnnTan
48+ 
49+### 参数说明
50+ 
51+| 参数名 | 输入/输出 | 描述 |
52+|-------|---------|------|
53+| workspace | 输入 | workspace 内存地址。若 workspaceSize 为 0，可传入 nullptr。 |
54+| workspaceSize | 输入 | workspace 大小，由 aclnnTanGetWorkspaceSize 返回。 |
55+| executor | 输入 | 算子执行器，由 aclnnTanGetWorkspaceSize 返回。 |
56+| stream | 输入 | ACL stream，用于异步调度算子执行。 |
57+ 
58+### 返回值说明
59+ 
60+返回 `aclnnStatus` 错误码，详见 [aclnn 错误码](#错误码)。
61+ 
62+## 错误码
63+ 
64+| 错误码 | 描述 |
65+|-------|------|
66+| ACLNN_SUCCESS（0） | 执行成功。 |
67+| ACLNN_ERR_PARAM_NULLPTR | 输入/输出 tensor 指针为空。 |
68+| ACLNN_ERR_PARAM_INVALID | 参数非法，包括：数据类型不支持、out shape 与 x 不一致等。 |
69+| ACLNN_ERR_INNER_CREATE_EXECUTOR | 内部创建算子执行器失败。 |
70+| ACLNN_ERR_INNER_NULLPTR | 内部 tensor 分配失败。 |
71+| ACLNN_ERR_INNER_INFERSHAPE_ERROR | 内部 InferShape 失败。 |
72+ 
73+## 约束说明
74+ 
75+- 输入 `x` 支持 float32 和 float16 数据类型。
76+- `out` 的数据类型须与 `x` 相同。
77+- `out` 的 shape 须与 `x` 相同（不支持广播）。
78+- 支持标量输入（内部转为 shape {1} 处理）。
79+- 支持空 tensor（元素数为 0），此时 workspaceSize 为 0，直接返回成功。
80+- workspace 须在调用 `aclnnTan` 之前分配，在 stream 中算子执行完成后方可释放。
81+- 当输入值接近 $\frac{\pi}{2} + k\pi$（$k$ 为整数）时，正切函数结果趋向无穷大，可能出现精度下降或溢出。
82+ 
83+## 调用示例
84+ 
85+以下示例展示了 Tan 算子的完整调用流程：
86+ 
87+```cpp
88+#include <cstdio>
89+#include <vector>
90+#include "acl/acl.h"
91+#include "aclnn_tan.h"
92+ 
93+int main() {
94+ // 1. 初始化 ACL 及设备
95+ aclInit(nullptr);
96+ aclrtSetDevice(0);
97+ aclrtStream stream;
98+ aclrtCreateStream(&stream);
99+ 
100+ // 2. 准备输入数据（fp32，shape=[2,4]）
101+ // x = [0.0, 0.5, 1.0, -1.0, 0.25, -0.5, 2.0, -2.0]
102+ // out = tan(x)
103+ int64_t shape[] = {2, 4};
104+ int64_t strides[] = {4, 1};
105+ float x_host[] = {0.0f, 0.5f, 1.0f, -1.0f, 0.25f, -0.5f, 2.0f, -2.0f};
106+ float out_host[8] = {0};
107+ 
108+ void *x_dev = nullptr, *out_dev = nullptr;
109+ size_t nbytes = 8 * sizeof(float);
110+ aclrtMalloc(&x_dev, nbytes, ACL_MEM_MALLOC_NORMAL_ONLY);
111+ aclrtMalloc(&out_dev, nbytes, ACL_MEM_MALLOC_NORMAL_ONLY);
112+ aclrtMemcpy(x_dev, nbytes, x_host, nbytes, ACL_MEMCPY_HOST_TO_DEVICE);
113+ 
114+ // 3. 创建 aclTensor
115+ aclTensor *x = aclCreateTensor(shape, 2, ACL_FLOAT, strides, 0,
116+ ACL_FORMAT_ND, shape, 2, x_dev);
117+ aclTensor *out = aclCreateTensor(shape, 2, ACL_FLOAT, strides, 0,
118+ ACL_FORMAT_ND, shape, 2, out_dev);
119+ 
120+ // 4. 查询 workspace 大小并分配
121+ uint64_t workspaceSize = 0;
122+ aclOpExecutor *executor = nullptr;
123+ aclnnTanGetWorkspaceSize(x, out, &workspaceSize, &executor);
124+ 
125+ void *workspace = nullptr;
126+ if (workspaceSize > 0)
127+ aclrtMalloc(&workspace, workspaceSize, ACL_MEM_MALLOC_NORMAL_ONLY);
128+ 
129+ // 5. 执行算子
130+ aclnnTan(workspace, workspaceSize, executor, stream);
131+ aclrtSynchronizeStream(stream);
132+ 
133+ // 6. 取回结果
134+ aclrtMemcpy(out_host, nbytes, out_dev, nbytes, ACL_MEMCPY_DEVICE_TO_HOST);
135+ printf("out = [%.4f, %.4f, %.4f, %.4f, %.4f, %.4f, %.4f, %.4f]\n",
136+ out_host[0], out_host[1], out_host[2], out_host[3],
137+ out_host[4], out_host[5], out_host[6], out_host[7]);
138+ // 期望: [0.0000, 0.5463, 1.5574, -1.5574, 0.2553, -0.5463, -2.1850, 2.1850]
139+ 
140+ // 7. 释放资源
141+ if (workspace) aclrtFree(workspace);
142+ aclrtFree(x_dev); aclrtFree(out_dev);
143+ aclDestroyTensor(x); aclDestroyTensor(out);
144+ aclrtDestroyStream(stream);
145+ aclrtResetDevice(0);
146+ aclFinalize();
147+ return 0;
148+}
149+```
Aexperimental/math/tan/examples/test_aclnn_tan.cpp+167-0
@@ -0,0 +1,167 @@
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+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+ /**
17+ * \file test_aclnn_tan.cpp
18+ * \brief Tan 算子 aclnn 调用示例（FP32）
19+ *
20+ * 计算: y = tan(x) = sin(x) / cos(x)
21+ */
22+ 
23+#include <iostream>
24+#include <vector>
25+#include <cmath>
26+#include "acl/acl.h"
27+#include "aclnn_tan.h"
28+ 
29+#define CHECK_RET(cond, return_expr) \
30+ do { \
31+ if (!(cond)) { \
32+ return_expr; \
33+ } \
34+ } while (0)
35+ 
36+#define LOG_PRINT(message, ...) \
37+ do { \
38+ printf(message, ##__VA_ARGS__); \
39+ } while (0)
40+ 
41+int64_t GetShapeSize(const std::vector<int64_t>& shape)
42+{
43+ int64_t shapeSize = 1;
44+ for (auto i : shape) {
45+ shapeSize *= i;
46+ }
47+ return shapeSize;
48+}
49+ 
50+int Init(int32_t deviceId, aclrtStream* stream)
51+{
52+ auto ret = aclInit(nullptr);
53+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
54+ ret = aclrtSetDevice(deviceId);
55+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
56+ ret = aclrtCreateStream(stream);
57+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
58+ return 0;
59+}
60+ 
61+template <typename T>
62+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
63+ aclDataType dataType, aclTensor** tensor)
64+{
65+ auto size = GetShapeSize(shape) * sizeof(T);
66+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
67+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
68+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
69+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
70+ 
71+ std::vector<int64_t> strides(shape.size(), 1);
72+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
73+ strides[i] = shape[i + 1] * strides[i + 1];
74+ }
75+ 
76+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0,
77+ aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), *deviceAddr);
78+ return 0;
79+}
80+ 
81+int main()
82+{
83+ // 1. 初始化
84+ int32_t deviceId = 0;
85+ aclrtStream stream;
86+ auto ret = Init(deviceId, &stream);
87+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
88+ 
89+ // 2. 构造输入与输出
90+ std::vector<int64_t> xShape = {2, 4};
91+ std::vector<int64_t> outShape = {2, 4};
92+ int64_t totalSize = GetShapeSize(outShape);
93+ 
94+ std::vector<float> xHostData = {0.0f, 0.5f, 1.0f, -1.0f, 0.25f, -0.5f, 2.0f, -2.0f};
95+ std::vector<float> outHostData(totalSize, 0.0f);
96+ 
97+ void* xDeviceAddr = nullptr;
98+ void* outDeviceAddr = nullptr;
99+ aclTensor* x = nullptr;
100+ aclTensor* out = nullptr;
101+ 
102+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_FLOAT, &x);
103+ CHECK_RET(ret == ACL_SUCCESS, return ret);
104+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
105+ CHECK_RET(ret == ACL_SUCCESS, return ret);
106+ 
107+ // 3. 调用 aclnnTan
108+ uint64_t workspaceSize = 0;
109+ aclOpExecutor* executor;
110+ ret = aclnnTanGetWorkspaceSize(x, out, &workspaceSize, &executor);
111+ CHECK_RET(ret == ACL_SUCCESS,
112+ LOG_PRINT("aclnnTanGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
113+ 
114+ void* workspaceAddr = nullptr;
115+ if (workspaceSize > 0) {
116+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
117+ CHECK_RET(ret == ACL_SUCCESS,
118+ LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
119+ }
120+ 
121+ ret = aclnnTan(workspaceAddr, workspaceSize, executor, stream);
122+ CHECK_RET(ret == ACL_SUCCESS,
123+ LOG_PRINT("aclnnTan failed. ERROR: %d\n", ret); return ret);
124+ 
125+ // 4. 同步
126+ ret = aclrtSynchronizeStream(stream);
127+ CHECK_RET(ret == ACL_SUCCESS,
128+ LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
129+ 
130+ // 5. 拷贝回 host 并比对
131+ std::vector<float> resultData(totalSize, 0.0f);
132+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(float),
133+ outDeviceAddr, totalSize * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);
134+ CHECK_RET(ret == ACL_SUCCESS,
135+ LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
136+ 
137+ int failCount = 0;
138+ const float atol = 1e-4f;
139+ const float rtol = 1e-4f;
140+ LOG_PRINT("=== aclnnTan result vs golden ===\n");
141+ for (int64_t i = 0; i < totalSize; i++) {
142+ float gold = std::tan(xHostData[i]);
143+ float diff = std::fabs(resultData[i] - gold);
144+ bool ok = diff <= (atol + rtol * std::fabs(gold));
145+ LOG_PRINT("[%ld] x=%8.4f out=%10.6f gold=%10.6f diff=%.2e %s\n",
146+ i, xHostData[i], resultData[i], gold, diff, ok ? "OK" : "FAIL");
147+ if (!ok) failCount++;
148+ }
149+ 
150+ // 6. 释放资源
151+ aclDestroyTensor(x);
152+ aclDestroyTensor(out);
153+ aclrtFree(xDeviceAddr);
154+ aclrtFree(outDeviceAddr);
155+ if (workspaceSize > 0) aclrtFree(workspaceAddr);
156+ aclrtDestroyStream(stream);
157+ aclrtResetDevice(deviceId);
158+ aclFinalize();
159+ 
160+ if (failCount == 0) {
161+ LOG_PRINT("=== PASS: all %ld elements match golden ===\n", totalSize);
162+ return 0;
163+ } else {
164+ LOG_PRINT("=== FAIL: %d / %ld elements mismatch ===\n", failCount, totalSize);
165+ return 1;
166+ }
167+}
Aexperimental/math/tan/op_graph/tan_proto.h+36-0
@@ -0,0 +1,36 @@
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+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+
17+/**
18+ * \file tan_proto.h
19+ * \brief Tan operator graph-mode (GE IR) proto registration
20+ */
21+#ifndef OPS_OP_PROTO_INC_TAN_H_
22+#define OPS_OP_PROTO_INC_TAN_H_
23+ 
24+#include "graph/operator_reg.h"
25+#include "graph/types.h"
26+ 
27+namespace ge {
28+ 
29+REG_OP(Tan)
30+ .INPUT(x, TensorType({DT_FLOAT, DT_FLOAT16}))
S
Ssu-yueming4月9日

数据类型还是不太对

likedislike
31+ .OUTPUT(y, TensorType({DT_FLOAT, DT_FLOAT16}))
32+ .OP_END_FACTORY_REG(Tan)
33+ 
34+} // namespace ge
35+ 
36+#endif // OPS_OP_PROTO_INC_TAN_H_
Aexperimental/math/tan/op_host/CMakeLists.txt+14-0
@@ -0,0 +1,14 @@
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+ # NOTE: Portions of this code were AI-generated and have been
12+ # technically reviewed for functional accuracy and security
13+ 
14+add_modules_sources(OPTYPE tan ACLNNTYPE aclnn)
Aexperimental/math/tan/op_host/tan_def.cpp+52-0
@@ -0,0 +1,52 @@
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+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+/**
17+ * \file tan_def.cpp
18+ * \brief Tan operator definition - declares inputs, outputs, and chip configuration
19+ */
20+#include "register/op_def_registry.h"
21+ 
22+namespace ops {
23+class Tan : public OpDef {
24+public:
25+ explicit Tan(const char* name) : OpDef(name)
26+ {
27+ this->Input("x")
28+ .ParamType(REQUIRED)
29+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
30+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
31+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
32+ .AutoContiguous();
33+ this->Output("y")
34+ .ParamType(REQUIRED)
35+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
36+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
37+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
38+ .AutoContiguous();
39+ 
40+ OpAICoreConfig aicoreConfig910B;
41+ aicoreConfig910B.DynamicCompileStaticFlag(true)
42+ .DynamicFormatFlag(false)
43+ .DynamicRankSupportFlag(true)
44+ .DynamicShapeSupportFlag(true)
45+ .NeedCheckSupportFlag(false)
46+ .PrecisionReduceFlag(true)
47+ .ExtendCfgInfo("opFile.value", "tan");
48+ this->AICore().AddConfig("ascend910b", aicoreConfig910B);
49+ }
50+};
51+OP_ADD(Tan);
52+} // namespace ops
Aexperimental/math/tan/op_host/tan_infershape.cpp+47-0
@@ -0,0 +1,47 @@
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+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+/**
17+ * \file tan_infershape.cpp
18+ * \brief Tan shape inference - output shape equals input shape
19+ */
20+ 
21+#include "register/op_impl_registry.h"
22+#include "exe_graph/runtime/infer_shape_context.h"
23+ 
24+using namespace ge;
25+ 
26+namespace ops {
27+ 
28+static ge::graphStatus InferShape4Tan(gert::InferShapeContext* context)
29+{
30+ const gert::Shape* input_shape = context->GetInputShape(0);
31+ if (input_shape == nullptr) {
32+ return ge::GRAPH_FAILED;
33+ }
34+ 
35+ gert::Shape* output_shape = context->GetOutputShape(0);
36+ if (output_shape == nullptr) {
37+ return ge::GRAPH_FAILED;
38+ }
39+ 
40+ *output_shape = *input_shape;
41+ 
42+ return ge::GRAPH_SUCCESS;
43+}
44+ 
45+IMPL_OP_INFERSHAPE(Tan).InferShape(InferShape4Tan);
46+ 
47+} // namespace ops
Aexperimental/math/tan/op_host/tan_tiling.cpp+157-0
@@ -0,0 +1,157 @@
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+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+/**
17+ * \file tan_tiling.cpp
18+ * \brief Tan Tiling implementation (arch32 - Ascend910B)
19+ *
20+ * Computes tiling parameters for Tan operator:
21+ * - Multi-core splitting: totalNum divided evenly across AI Cores
22+ * - UB splitting: each core processes data in chunks of ubFactor elements
23+ * - TilingKey selection: based on input dtype (float32 vs float16)
24+ */
25+ 
26+#include "register/op_def_registry.h"
27+#include "op_common/log/log.h"
28+#include "op_common/op_host/util/math_util.h"
29+#include "op_common/op_host/util/platform_util.h"
30+#include "../op_kernel/tan_tiling_data.h"
31+#include "../op_kernel/tan_tiling_key.h"
32+ 
33+namespace optiling {
34+ 
35+using Ops::Base::CeilDiv;
36+using Ops::Base::FloorDiv;
37+using Ops::Base::FloorAlign;
38+using Ops::Base::GetUbBlockSize;
39+ 
40+constexpr uint32_t WS_SYS_SIZE = 0U;
41+// float32: inputQueue(x2) + outputQueue(x2) + tmpBuf1(x1) + tmpBuf2(x1) = 6 float-sized buffers
42+// Total bytes = ubFactor * 6 * sizeof(float)
43+constexpr int64_t BUFFER_NUM_FP32 = 6;
44+// float16: inputQueue(x2,half) + outputQueue(x2,half) + tmpBuf1(x1,float) + tmpBuf2(x1,float)
45+// Total bytes = ubFactor * (2*2 + 2*2 + 4 + 4) = ubFactor * 16 = ubFactor * sizeof(float) * 4
46+constexpr int64_t BUFFER_NUM_FP16 = 4;
47+ 
48+static const gert::Shape g_vec_1_shape = {1};
49+ 
50+static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape)
51+{
52+ if (in_shape.GetDimNum() == 0) {
53+ return g_vec_1_shape;
54+ }
55+ return in_shape;
56+}
57+ 
58+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
59+{
60+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
61+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
62+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
63+ coreNum = ascendcPlatform.GetCoreNumAiv();
64+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
65+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
66+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
67+ return ge::GRAPH_SUCCESS;
68+}
69+ 
70+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
71+{
72+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
73+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
74+ currentWorkspace[0] = WS_SYS_SIZE;
75+ return ge::GRAPH_SUCCESS;
76+}
77+ 
78+static ge::graphStatus TanTilingFunc(gert::TilingContext* context)
79+{
80+ // 1. Get platform info
81+ uint64_t ubSize;
82+ int64_t coreNum;
83+ OP_CHECK_IF(
84+ GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
85+ OP_LOGE(context, "GetPlatformInfo error"),
86+ return ge::GRAPH_FAILED);
87+ 
88+ // 2. Get input shape and dtype
89+ auto inputShape = context->GetInputShape(0);
90+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
91+ auto storageShape = EnsureNotScalar(inputShape->GetStorageShape());
92+ int64_t totalNum = storageShape.GetShapeSize();
93+ 
94+ auto inputDesc = context->GetInputDesc(0);
95+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
96+ ge::DataType dtype = inputDesc->GetDataType();
97+ 
98+ // 3. Get workspace size
99+ OP_CHECK_IF(
100+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
101+ OP_LOGE(context, "GetWorkspaceSize error"),
102+ return ge::GRAPH_FAILED);
103+ 
104+ // 4. Set TilingData
105+ TanTilingData* tiling = context->GetTilingData<TanTilingData>();
106+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
107+ OP_CHECK_IF(
108+ memset_s(tiling, sizeof(TanTilingData), 0, sizeof(TanTilingData)) != EOK,
109+ OP_LOGE(context, "set tiling data error"),
110+ return ge::GRAPH_FAILED);
111+ 
112+ // Handle empty tensor
113+ if (totalNum == 0) {
114+ tiling->totalNum = 0;
115+ tiling->blockFactor = 0;
116+ tiling->ubFactor = 0;
117+ context->SetBlockDim(1);
118+ context->SetTilingKey(GET_TPL_TILING_KEY(TAN_TPL_SCH_MODE_0));
119+ return ge::GRAPH_SUCCESS;
120+ }
121+ 
122+ // 5. Multi-core splitting
123+ tiling->totalNum = totalNum;
124+ tiling->blockFactor = CeilDiv(totalNum, coreNum);
125+ int64_t usedCoreNum = CeilDiv(totalNum, tiling->blockFactor);
126+ 
127+ // 6. UB splitting and TilingKey selection
128+ int64_t ubCanUse = static_cast<int64_t>(ubSize);
129+ int64_t ubBlockSize = GetUbBlockSize(context);
130+ // Both paths compute internally in float32, so typeSize = 4
131+ constexpr int64_t typeSize = 4;
132+ 
133+ if (dtype == ge::DT_FLOAT) {
134+ tiling->ubFactor = FloorAlign(FloorDiv((ubCanUse / typeSize), BUFFER_NUM_FP32), ubBlockSize);
135+ context->SetTilingKey(GET_TPL_TILING_KEY(TAN_TPL_SCH_MODE_0));
136+ } else if (dtype == ge::DT_FLOAT16) {
137+ tiling->ubFactor = FloorAlign(FloorDiv((ubCanUse / typeSize), BUFFER_NUM_FP16), ubBlockSize);
138+ context->SetTilingKey(GET_TPL_TILING_KEY(TAN_TPL_SCH_MODE_1));
139+ } else {
140+ OP_LOGE(context, "Tan: unsupported dtype");
141+ return ge::GRAPH_FAILED;
142+ }
143+ 
144+ context->SetBlockDim(usedCoreNum);
145+ return ge::GRAPH_SUCCESS;
146+}
147+ 
148+static ge::graphStatus TilingParseForTan([[maybe_unused]] gert::TilingParseContext* context)
149+{
150+ return ge::GRAPH_SUCCESS;
151+}
152+ 
153+struct TanCompileInfo {};
154+ 
155+IMPL_OP_OPTILING(Tan).Tiling(TanTilingFunc).TilingParse<TanCompileInfo>(TilingParseForTan);
156+ 
157+} // namespace optiling
Aexperimental/math/tan/op_kernel/tan.cpp+44-0
@@ -0,0 +1,44 @@
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+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+/**
17+ * \file tan_arch32.cpp
18+ * \brief Tan kernel entry point (arch32 architecture - Ascend910B)
19+ */
20+ 
21+#include "tan.h"
22+ 
23+enum class TanTilingKey : uint32_t
24+{
25+ TILING_KEY_FLOAT32 = 0,
26+ TILING_KEY_FLOAT16 = 1,
27+};
28+ 
29+template <uint32_t schMode>
30+__global__ __aicore__ void tan(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
31+{
32+ REGISTER_TILING_DEFAULT(TanTilingData);
33+ GET_TILING_DATA_WITH_STRUCT(TanTilingData, tilingData, tiling);
34+ if constexpr (schMode == static_cast<uint32_t>(TanTilingKey::TILING_KEY_FLOAT32)) {
35+ NsTan::Tan<float> op;
36+ op.Init(x, y, &tilingData);
37+ op.Process();
38+ }
39+ if constexpr (schMode == static_cast<uint32_t>(TanTilingKey::TILING_KEY_FLOAT16)) {
40+ NsTan::Tan<half> op;
41+ op.Init(x, y, &tilingData);
42+ op.Process();
43+ }
44+}
Aexperimental/math/tan/op_kernel/tan.h+204-0
@@ -0,0 +1,204 @@
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+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+/**
17+ * \file tan.h
18+ * \brief Tan kernel class definition (arch32 - Ascend910B)
19+ *
20+ * Computes tan(x) = sin(x) / cos(x) for each element.
21+ * - float32 path: direct computation using Sin, Cos, Div
22+ * - float16 path: cast to float32 for computation, then cast back
23+ *
24+ * Uses double buffering (BUFFER_NUM=2) for pipeline parallelism:
25+ * - CopyIn: GM -> UB (input data transfer)
26+ * - Compute: UB vector computation (tan formula)
27+ * - CopyOut: UB -> GM (output data transfer)
28+ */
29+#ifndef TAN_H
30+#define TAN_H
31+ 
32+#include "kernel_operator.h"
33+#include "kernel_tiling/kernel_tiling.h"
34+#include "tan_tiling_data.h"
35+#include "tan_tiling_key.h"
36+ 
37+namespace NsTan {
38+ 
39+using namespace AscendC;
40+ 
41+constexpr int32_t BUFFER_NUM = 2;
42+ 
43+template <typename T>
44+class Tan {
45+public:
46+ __aicore__ inline Tan() {};
47+ 
48+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const TanTilingData* tilingData);
49+ __aicore__ inline void Process();
50+ 
51+private:
52+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum);
53+ __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum);
54+ __aicore__ inline void Compute(int64_t currentNum);
55+ 
56+private:
57+ TPipe pipe;
58+ TQue<QuePosition::VECIN, BUFFER_NUM> inputQueue;
59+ TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueue;
60+ TBuf<QuePosition::VECCALC> tmpBuf1; // stores sin(x) intermediate result
61+ TBuf<QuePosition::VECCALC> tmpBuf2; // stores cos(x) intermediate result
62+ 
63+ GlobalTensor<T> inputGM;
64+ GlobalTensor<T> outputGM;
65+ 
66+ int64_t blockLength_ = 0;
67+ int64_t ubLength_ = 0;
68+};
69+ 
70+// ============================================================================
71+// Init - Initialize GM pointers and allocate UB buffers
72+// ============================================================================
73+template <typename T>
74+__aicore__ inline void Tan<T>::Init(GM_ADDR x, GM_ADDR y, const TanTilingData* tilingData)
75+{
76+ int64_t remainderLength = tilingData->totalNum - tilingData->blockFactor * AscendC::GetBlockIdx();
77+ blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength;
78+ // Clamp to 0 for idle cores (when totalNum < coreNum, some cores have no work)
79+ if (blockLength_ < 0) {
80+ blockLength_ = 0;
81+ }
82+ ubLength_ = tilingData->ubFactor;
83+ 
84+ // Guard: empty tensor or idle core - skip buffer allocation
85+ if (blockLength_ <= 0 || ubLength_ <= 0) {
86+ return;
87+ }
88+ 
89+ inputGM.SetGlobalBuffer((__gm__ T*)x + tilingData->blockFactor * AscendC::GetBlockIdx(), blockLength_);
90+ outputGM.SetGlobalBuffer((__gm__ T*)y + tilingData->blockFactor * AscendC::GetBlockIdx(), blockLength_);
91+ 
92+ pipe.InitBuffer(inputQueue, BUFFER_NUM, ubLength_ * sizeof(T));
93+ pipe.InitBuffer(outputQueue, BUFFER_NUM, ubLength_ * sizeof(T));
94+ // Temporary buffers for sin and cos intermediate results
95+ // For float32: used as float buffers directly
96+ // For float16: used as float32 buffers for precision-promoted computation
97+ pipe.InitBuffer(tmpBuf1, ubLength_ * sizeof(float));
98+ pipe.InitBuffer(tmpBuf2, ubLength_ * sizeof(float));
99+}
100+ 
101+// ============================================================================
102+// CopyIn - Transfer data from GM to UB
103+// ============================================================================
104+template <typename T>
105+__aicore__ inline void Tan<T>::CopyIn(int64_t progress, int64_t currentNum)
106+{
107+ AscendC::LocalTensor<T> xLocal = inputQueue.AllocTensor<T>();
108+ AscendC::DataCopyParams copyParams;
109+ copyParams.blockCount = 1;
110+ copyParams.blockLen = currentNum * sizeof(T);
111+ copyParams.srcStride = 0;
112+ copyParams.dstStride = 0;
113+ AscendC::DataCopyPad(xLocal, inputGM[progress * ubLength_], copyParams, {false, 0, 0, 0});
114+ inputQueue.EnQue(xLocal);
115+}
116+ 
117+// ============================================================================
118+// CopyOut - Transfer data from UB to GM
119+// ============================================================================
120+template <typename T>
121+__aicore__ inline void Tan<T>::CopyOut(int64_t progress, int64_t currentNum)
122+{
123+ AscendC::LocalTensor<T> yLocal = outputQueue.DeQue<T>();
124+ AscendC::DataCopyParams copyParams;
125+ copyParams.blockCount = 1;
126+ copyParams.blockLen = currentNum * sizeof(T);
127+ copyParams.srcStride = 0;
128+ copyParams.dstStride = 0;
129+ AscendC::DataCopyPad(outputGM[progress * ubLength_], yLocal, copyParams);
130+ outputQueue.FreeTensor(yLocal);
131+}
132+ 
133+// ============================================================================
134+// Compute - float32 specialization: direct tan computation
135+// tan(x) = sin(x) / cos(x)
136+// ============================================================================
137+template <>
138+__aicore__ inline void Tan<float>::Compute(int64_t currentNum)
139+{
140+ AscendC::LocalTensor<float> xLocal = inputQueue.DeQue<float>();
141+ AscendC::LocalTensor<float> yLocal = outputQueue.AllocTensor<float>();
142+ AscendC::LocalTensor<float> sinVal = tmpBuf1.Get<float>();
143+ AscendC::LocalTensor<float> cosVal = tmpBuf2.Get<float>();
144+ 
145+ // Step 1: sin(x) -> sinVal (tmpBuf1)
146+ AscendC::Sin(sinVal, xLocal, currentNum);
147+ // Step 2: cos(x) -> cosVal (tmpBuf2)
148+ AscendC::Cos(cosVal, xLocal, currentNum);
149+ // Step 3: sin(x) / cos(x) -> yLocal
150+ AscendC::Div(yLocal, sinVal, cosVal, currentNum);
151+ 
152+ outputQueue.EnQue<float>(yLocal);
153+ inputQueue.FreeTensor(xLocal);
154+}
155+ 
156+// ============================================================================
157+// Compute - float16 specialization: cast to float32, compute, cast back
158+// Flow: Cast(half->float) -> Cos -> Sin -> Div -> Cast(float->half)
159+// Key constraint: Cos must be computed before Sin to avoid overwriting input
160+// ============================================================================
161+template <>
162+__aicore__ inline void Tan<half>::Compute(int64_t currentNum)
163+{
164+ AscendC::LocalTensor<half> xLocal = inputQueue.DeQue<half>();
165+ AscendC::LocalTensor<half> yLocal = outputQueue.AllocTensor<half>();
166+ AscendC::LocalTensor<float> sinVal = tmpBuf1.Get<float>();
167+ AscendC::LocalTensor<float> cosVal = tmpBuf2.Get<float>();
168+ 
169+ // Step 1: Cast half -> float (store in sinVal as temp for x_float)
170+ AscendC::Cast(sinVal, xLocal, AscendC::RoundMode::CAST_NONE, currentNum);
171+ // Step 2: cos(x) - must compute before sin overwrites sinVal
172+ AscendC::Cos(cosVal, sinVal, currentNum);
173+ // Step 3: sin(x) - overwrites sinVal (cosVal already saved)
174+ AscendC::Sin(sinVal, sinVal, currentNum);
175+ // Step 4: sin(x) / cos(x) -> sinVal (reuse)
176+ AscendC::Div(sinVal, sinVal, cosVal, currentNum);
177+ // Step 5: Cast float -> half
178+ AscendC::Cast(yLocal, sinVal, AscendC::RoundMode::CAST_ROUND, currentNum);
179+ 
180+ outputQueue.EnQue<half>(yLocal);
181+ inputQueue.FreeTensor(xLocal);
182+}
183+ 
184+// ============================================================================
185+// Process - Main loop: iterate over tiles
186+// ============================================================================
187+template <typename T>
188+__aicore__ inline void Tan<T>::Process()
189+{
190+ // Guard: empty tensor or idle core - nothing to process
191+ if (blockLength_ <= 0 || ubLength_ <= 0) {
192+ return;
193+ }
194+ int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
195+ for (int64_t i = 0; i < loopCount; i++) {
196+ int64_t currentNum = (i == (loopCount - 1)) ? (blockLength_ - ubLength_ * i) : ubLength_;
197+ CopyIn(i, currentNum);
198+ Compute(currentNum);
199+ CopyOut(i, currentNum);
200+ }
201+}
202+ 
203+} // namespace NsTan
204+#endif // TAN_H
Aexperimental/math/tan/op_kernel/tan_tiling_data.h+30-0
@@ -0,0 +1,30 @@
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+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+/**
17+ * \file tan_tiling_data.h
18+ * \brief Tan TilingData structure definition
19+ */
20+ 
21+#ifndef _TAN_TILING_DATA_H_
22+#define _TAN_TILING_DATA_H_
23+ 
24+struct TanTilingData {
25+ int64_t totalNum = 0; // Total number of elements
26+ int64_t blockFactor = 0; // Number of elements per AI Core
27+ int64_t ubFactor = 0; // Number of elements per UB loop iteration
28+};
29+ 
30+#endif
Aexperimental/math/tan/op_kernel/tan_tiling_key.h+42-0
@@ -0,0 +1,42 @@
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+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+/**
17+ * \file tan_tiling_key.h
18+ * \brief Tan TilingKey definition
19+ *
20+ * TilingKey mapping:
21+ * - TAN_TPL_SCH_MODE_0 (0): FLOAT32 type
22+ * - TAN_TPL_SCH_MODE_1 (1): FLOAT16 type
23+ */
24+ 
25+#ifndef __TAN_TILING_KEY_H__
26+#define __TAN_TILING_KEY_H__
27+ 
28+#include "ascendc/host_api/tiling/template_argument.h"
29+ 
30+#define TAN_TPL_SCH_MODE_0 0 // FLOAT32 type
31+#define TAN_TPL_SCH_MODE_1 1 // FLOAT16 type
32+ 
33+ASCENDC_TPL_ARGS_DECL(
34+ Tan,
35+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST,
36+ TAN_TPL_SCH_MODE_0, TAN_TPL_SCH_MODE_1));
37+ 
38+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
39+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST,
40+ TAN_TPL_SCH_MODE_0, TAN_TPL_SCH_MODE_1)));
41+ 
42+#endif
Aexperimental/math/tan/tests/.gitkeep+0-0
The file is empty