Star1669
1618
代码
代码
Issues412
Pull Requests457
Actions
项目讨论
Wiki
分析
Star1669
1618
已合并
【CANN开源开放社区任务】【社区任务】AscendC实现Polar算子贡献 #2923
bububu创建于 5月23日
【CANN开源开放社区任务】【社区任务】AscendC实现Polar算子贡献 #2923
已合并
bububu创建于 5月23日
从LiJianhao2/ops-mathmaster合入到cann/ops-mathmaster
共 39 个文件变更+4650-0
Aexperimental/math/polar/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
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+add_all_modules_sources(OPTYPE polar ACLNNTYPE aclnn_exclude)
Aexperimental/math/polar/README.md+84-0
@@ -0,0 +1,84 @@
1+# Polar
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
8+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
9+| <term>Ascend 950PR/Ascend 950DT</term> | × |
10+| <term>Atlas 200I/500 A2 推理产品</term> | × |
11+| <term>Atlas 推理系列产品</term> | × |
12+| <term>Atlas 训练系列产品</term> | × |
13+ 
14+## 功能说明
15+ 
16+- 算子功能：由模长（abs）与幅角（angle）构造极坐标复数张量，对应 PyTorch 接口 `torch.polar(abs, angle)`。
17+ 
18+- 计算公式：
19+ 
20+ $$
21+ out_i = input_i \times (\cos(angle_i) + i \cdot \sin(angle_i))
22+ $$
23+ 
24+## 参数说明
25+ 
26+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
27+ <col style="width: 100px">
28+ <col style="width: 150px">
29+ <col style="width: 280px">
30+ <col style="width: 330px">
31+ <col style="width: 120px">
32+ </colgroup>
33+ <thead>
34+ <tr>
35+ <th>参数名</th>
36+ <th>输入/输出/属性</th>
37+ <th>描述</th>
38+ <th>数据类型</th>
39+ <th>数据格式</th>
40+ </tr></thead>
41+ <tbody>
42+ <tr>
43+ <td>input</td>
44+ <td>输入</td>
45+ <td>极坐标模长分量，公式中的input。</td>
46+ <td>FLOAT</td>
47+ <td>ND</td>
48+ </tr>
49+ <tr>
50+ <td>angle</td>
51+ <td>输入</td>
52+ <td>极坐标幅角（弧度），公式中的angle。</td>
53+ <td>FLOAT</td>
54+ <td>ND</td>
55+ </tr>
56+ <tr>
57+ <td>out</td>
58+ <td>输出</td>
59+ <td>计算结果，公式中的out；shape为input与angle广播后的shape。</td>
60+ <td>COMPLEX64</td>
61+ <td>ND</td>
62+ </tr>
63+ </tbody></table>
64+ 
65+## 约束说明
66+ 
67+- input与angle的数据类型必须一致（FLOAT）。
68+- input与angle的shape满足NumPy广播关系，out的shape为两者广播后的shape。
69+- input、angle的维度数不超过8维。
70+ 
71+## 调用说明
72+ 
73+| 调用方式 | 调用样例 | 说明 |
74+|--------------|------------------------------------------------------------------------|--------------------------------------------------------------|
75+| aclnn调用 | [test_aclnn_polar](./examples/test_aclnn_polar.cpp) | 通过[aclnnPolar](./docs/aclnnPolar.md)接口方式调用Polar算子。 |
76+| 图模式调用 | | |
77+ 
78+本目录包含Polar算子的aclnn接口及AscendC实现；关于该算子的设计与自测，请参考[设计文档](./docs/design.md)与[自测报告](./tests/自测报告.md)。
F
Ffulltower6月30日

README.md 缺少 ## 贡献说明 章节。根据 CANN 文档交付规范，experimental 目录下新算子必须在 README 中包含贡献说明章节，格式参考：

## 贡献说明
| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
| ---- | ---- | ---- | ---- | ---- |
| <姓名> | <单位/组织> | <算子名> | <YYYY/MM/DD> | <算子名>算子适配开源仓 |
likedislike
79+ 
80+## 贡献说明
81+ 
82+| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
83+| ---- | ---- | ---- | ---- | ---- |
84+| 开源社区贡献者 | 开源社区 | Polar | 2026/06/30 | Polar算子适配开源仓 |
Aexperimental/math/polar/docs/aclnnPolar.md+379-0
@@ -0,0 +1,379 @@
1+# aclnnPolar
F
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
8+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
9+| <term>Ascend 950PR/Ascend 950DT</term> | × |
10+| <term>Atlas 200I/500 A2 推理产品</term> | × |
11+| <term>Atlas 推理系列产品</term> | × |
12+| <term>Atlas 训练系列产品</term> | × |
13+ 
14+## 功能说明
15+ 
16+- 接口功能：由模长（abs）与幅角（angle）构造极坐标复数张量，对应 PyTorch 接口 `torch.polar(abs, angle)`。
17+ 
18+- 计算公式：
19+ 
20+ $$
21+ out_i = input_i \times (\cos(angle_i) + i \cdot \sin(angle_i))
22+ $$
23+ 
24+## 函数原型
25+ 
26+每个算子分为[两段式接口](../../../../docs/zh/context/两段式接口.md)，必须先调用 “aclnnPolarGetWorkspaceSize” 接口获取入参并根据计算流程计算所需workspace大小，再调用 “aclnnPolar” 接口执行计算。
27+ 
28+```Cpp
29+aclnnStatus aclnnPolarGetWorkspaceSize(
30+ const aclTensor* input,
31+ const aclTensor* angle,
32+ aclTensor* out,
33+ uint64_t* workspaceSize,
34+ aclOpExecutor** executor)
35+```
36+ 
37+```Cpp
38+aclnnStatus aclnnPolar(
39+ void* workspace,
40+ uint64_t workspaceSize,
41+ aclOpExecutor* executor,
42+ const aclrtStream stream)
43+```
44+ 
45+## aclnnPolarGetWorkspaceSize
46+ 
47+- **参数说明：**
48+ 
49+ <table style="undefined;table-layout: fixed; width: 1547px"><colgroup>
50+ <col style="width: 150px">
51+ <col style="width: 121px">
52+ <col style="width: 206px">
53+ <col style="width: 456px">
54+ <col style="width: 211px">
55+ <col style="width: 122px">
56+ <col style="width: 135px">
57+ <col style="width: 146px">
58+ </colgroup>
59+ <thead>
60+ <tr>
61+ <th>参数名</th>
62+ <th>输入/输出</th>
63+ <th>描述</th>
64+ <th>使用说明</th>
65+ <th>数据类型</th>
66+ <th>数据格式</th>
67+ <th>维度(shape)</th>
68+ <th>非连续Tensor</th>
69+ </tr></thead>
70+ <tbody>
71+ <tr>
72+ <td>input</td>
73+ <td>输入</td>
74+ <td>公式中的input（极坐标模长分量）。</td>
75+ <td>
76+ <ul>
77+ <li>数据类型需与angle一致。</li>
78+ <li>input与angle的shape满足<a href="../../../../docs/zh/context/broadcast关系.md" target="_blank">broadcast关系</a>。</li>
79+ </ul>
80+ </td>
81+ <td>FLOAT</td>
82+ <td>ND</td>
83+ <td>0-8</td>
84+ <td>√</td>
85+ </tr>
86+ <tr>
87+ <td>angle</td>
88+ <td>输入</td>
89+ <td>公式中的angle（极坐标幅角，弧度）。</td>
90+ <td>
91+ <ul>
92+ <li>数据类型需与input一致。</li>
93+ <li>input与angle的shape满足<a href="../../../../docs/zh/context/broadcast关系.md" target="_blank">broadcast关系</a>。</li>
94+ </ul>
95+ </td>
96+ <td>FLOAT</td>
97+ <td>ND</td>
98+ <td>0-8</td>
99+ <td>√</td>
100+ </tr>
101+ <tr>
102+ <td>out</td>
103+ <td>输出</td>
104+ <td>公式中的out。</td>
105+ <td>out的shape与input、angle做broadcast后的shape一致；数据类型恒为COMPLEX64。</td>
106+ <td>COMPLEX64</td>
107+ <td>ND</td>
108+ <td>0-8</td>
109+ <td>√</td>
110+ </tr>
111+ <tr>
112+ <td>workspaceSize</td>
113+ <td>输出</td>
114+ <td>返回需要在Device侧申请的workspace大小。</td>
115+ <td>-</td>
116+ <td>-</td>
117+ <td>-</td>
118+ <td>-</td>
119+ <td>-</td>
120+ </tr>
121+ <tr>
122+ <td>executor</td>
123+ <td>输出</td>
124+ <td>返回op执行器，包含了算子计算流程。</td>
125+ <td>-</td>
126+ <td>-</td>
127+ <td>-</td>
128+ <td>-</td>
129+ <td>-</td>
130+ </tr>
131+ </tbody>
132+ </table>
133+ 
134+- **返回值：**
135+ 
136+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
137+ 
138+ 第一段接口完成入参校验，出现如下场景时报错：
139+ 
140+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
141+ <col style="width: 267px">
142+ <col style="width: 124px">
143+ <col style="width: 775px">
144+ </colgroup>
145+ <thead>
146+ <tr>
147+ <th>返回值</th>
148+ <th>错误码</th>
149+ <th>描述</th>
150+ </tr></thead>
151+ <tbody>
152+ <tr>
153+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
154+ <td>161001</td>
155+ <td>传入的input、angle、out是空指针。</td>
156+ </tr>
157+ <tr>
158+ <td rowspan="5">ACLNN_ERR_PARAM_INVALID</td>
159+ <td rowspan="5">161002</td>
160+ <td>input、angle的数据类型或数据格式不在支持的范围之内。</td>
161+ </tr>
162+ <tr>
163+ <td>input与angle的数据类型不一致。</td>
164+ </tr>
165+ <tr>
166+ <td>input或angle的shape超过8维。</td>
167+ </tr>
168+ <tr>
169+ <td>input与angle的shape不满足broadcast推导关系。</td>
170+ </tr>
171+ <tr>
172+ <td>out的shape与input、angle做broadcast后的shape不一致，或out的数据类型不是COMPLEX64。</td>
173+ </tr>
174+ </tbody>
175+ </table>
176+ 
177+## aclnnPolar
178+ 
179+- **参数说明：**
180+ 
181+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
182+ <col style="width: 173px">
183+ <col style="width: 133px">
184+ <col style="width: 860px">
185+ </colgroup>
186+ <thead>
187+ <tr>
188+ <th>参数名</th>
189+ <th>输入/输出</th>
190+ <th>描述</th>
191+ </tr></thead>
192+ <tbody>
193+ <tr>
194+ <td>workspace</td>
195+ <td>输入</td>
196+ <td>在Device侧申请的workspace内存地址。</td>
197+ </tr>
198+ <tr>
199+ <td>workspaceSize</td>
200+ <td>输入</td>
201+ <td>在Device侧申请的workspace大小，由第一段接口aclnnPolarGetWorkspaceSize获取。</td>
202+ </tr>
203+ <tr>
204+ <td>executor</td>
205+ <td>输入</td>
206+ <td>op执行器，包含了算子计算流程。</td>
207+ </tr>
208+ <tr>
209+ <td>stream</td>
210+ <td>输入</td>
211+ <td>指定执行任务的Stream。</td>
212+ </tr>
213+ </tbody>
214+ </table>
215+ 
216+- **返回值：**
217+ 
218+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
219+ 
220+## 约束说明
221+ 
222+- input与angle的数据类型必须一致（均为FLOAT）。
223+- out的数据类型恒为COMPLEX64（与input无关）。
224+- input、angle的维度数不超过8维，数据格式仅支持ND。
225+- 确定性计算：aclnnPolar默认确定性实现。
226+ 
227+## 调用示例
228+ 
229+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../../docs/zh/context/编译与运行样例.md)。
230+ 
231+```Cpp
232+#include <iostream>
233+#include <vector>
234+#include "acl/acl.h"
235+#include "aclnnop/aclnn_polar.h"
236+ 
237+#define CHECK_RET(cond, return_expr) \
238+ do { \
239+ if (!(cond)) { \
240+ return_expr; \
241+ } \
242+ } while (0)
243+ 
244+#define LOG_PRINT(message, ...) \
245+ do { \
246+ printf(message, ##__VA_ARGS__); \
247+ } while (0)
248+ 
249+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
250+ int64_t shapeSize = 1;
251+ for (auto i : shape) {
252+ shapeSize *= i;
253+ }
254+ return shapeSize;
255+}
256+ 
257+int Init(int32_t deviceId, aclrtStream* stream) {
258+ // 固定写法，资源初始化
259+ auto ret = aclInit(nullptr);
260+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
261+ ret = aclrtSetDevice(deviceId);
262+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
263+ ret = aclrtCreateStream(stream);
264+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
265+ return 0;
266+}
267+ 
268+template <typename T>
269+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
270+ aclDataType dataType, aclTensor** tensor) {
271+ auto size = GetShapeSize(shape) * sizeof(T);
272+ // 调用aclrtMalloc申请device侧内存
273+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
274+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
275+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
276+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
277+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
278+ 
279+ // 计算连续tensor的strides
280+ std::vector<int64_t> strides(shape.size(), 1);
281+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
282+ strides[i] = shape[i + 1] * strides[i + 1];
283+ }
284+ 
285+ // 调用aclCreateTensor接口创建aclTensor
286+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
287+ shape.data(), shape.size(), *deviceAddr);
288+ return 0;
289+}
290+ 
291+int main() {
292+ // 1. （固定写法）device/stream初始化，参考acl API手册
293+ int32_t deviceId = 0;
294+ aclrtStream stream;
295+ auto ret = Init(deviceId, &stream);
296+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
297+ 
298+ // 2. 构造输入与输出，需要根据API的接口自定义构造
299+ std::vector<int64_t> inputShape = {4, 2};
300+ std::vector<int64_t> angleShape = {4, 2};
301+ std::vector<int64_t> outShape = {4, 2};
302+ void* inputDeviceAddr = nullptr;
303+ void* angleDeviceAddr = nullptr;
304+ void* outDeviceAddr = nullptr;
305+ aclTensor* input = nullptr;
306+ aclTensor* angle = nullptr;
307+ aclTensor* out = nullptr;
308+ std::vector<float> inputHostData = {1, 1, 1, 1, 2, 2, 2, 2};
309+ std::vector<float> angleHostData = {0, 1.5707964f, 3.1415927f, -1.5707964f, 0, 1.5707964f, 3.1415927f, -1.5707964f};
310+ 
311+ // 创建input/angle aclTensor（FLOAT）
312+ ret = CreateAclTensor(inputHostData, inputShape, &inputDeviceAddr, aclDataType::ACL_FLOAT, &input);
313+ CHECK_RET(ret == ACL_SUCCESS, return ret);
314+ ret = CreateAclTensor(angleHostData, angleShape, &angleDeviceAddr, aclDataType::ACL_FLOAT, &angle);
315+ CHECK_RET(ret == ACL_SUCCESS, return ret);
316+ 
317+ // 创建out aclTensor（COMPLEX64：每元素 = real/imag 两个 float）
318+ int64_t outNum = GetShapeSize(outShape);
319+ std::vector<float> outHostData(outNum * 2, 0);
320+ size_t outByte = static_cast<size_t>(outNum) * 2 * sizeof(float);
321+ ret = aclrtMalloc(&outDeviceAddr, outByte, ACL_MEM_MALLOC_HUGE_FIRST);
322+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
323+ ret = aclrtMemcpy(outDeviceAddr, outByte, outHostData.data(), outByte, ACL_MEMCPY_HOST_TO_DEVICE);
324+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
325+ std::vector<int64_t> outStrides(outShape.size(), 1);
326+ for (int64_t i = outShape.size() - 2; i >= 0; i--) {
327+ outStrides[i] = outShape[i + 1] * outStrides[i + 1];
328+ }
329+ out = aclCreateTensor(outShape.data(), outShape.size(), aclDataType::ACL_COMPLEX64, outStrides.data(), 0,
330+ aclFormat::ACL_FORMAT_ND, outShape.data(), outShape.size(), outDeviceAddr);
331+ 
332+ // 3. 调用CANN算子库API
333+ uint64_t workspaceSize = 0;
334+ aclOpExecutor* executor;
335+ // 调用aclnnPolar第一段接口
336+ ret = aclnnPolarGetWorkspaceSize(input, angle, out, &workspaceSize, &executor);
337+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnPolarGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
338+ // 根据第一段接口计算出的workspaceSize申请device内存
339+ void* workspaceAddr = nullptr;
340+ if (workspaceSize > 0) {
341+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
342+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
343+ }
344+ // 调用aclnnPolar第二段接口
345+ ret = aclnnPolar(workspaceAddr, workspaceSize, executor, stream);
346+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnPolar failed. ERROR: %d\n", ret); return ret);
347+ 
348+ // 4. （固定写法）同步等待任务执行结束
349+ ret = aclrtSynchronizeStream(stream);
350+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
351+ 
352+ // 5. 获取输出的值（complex64：每元素 real/imag 两个 float）
353+ std::vector<float> resultData(outNum * 2, 0);
354+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, outByte,
355+ ACL_MEMCPY_DEVICE_TO_HOST);
356+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
357+ for (int64_t i = 0; i < outNum; i++) {
358+ LOG_PRINT("out[%ld] = %f + %fi\n", i, resultData[2 * i], resultData[2 * i + 1]);
359+ }
360+ 
361+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
362+ aclDestroyTensor(input);
363+ aclDestroyTensor(angle);
364+ aclDestroyTensor(out);
365+ 
366+ // 7. 释放device资源
367+ aclrtFree(inputDeviceAddr);
368+ aclrtFree(angleDeviceAddr);
369+ aclrtFree(outDeviceAddr);
370+ if (workspaceSize > 0) {
371+ aclrtFree(workspaceAddr);
372+ }
373+ aclrtDestroyStream(stream);
374+ aclrtResetDevice(deviceId);
375+ aclFinalize();
376+ 
377+ return 0;
378+}
379+```
Aexperimental/math/polar/docs/design.md+239-0
@@ -0,0 +1,239 @@
1+# 需求背景（required）
2+ 
3+## 需求来源
4+ 
5+昇腾 CANN 训练营第二季社区任务（序号 **04-5 Polar**）：基于 Ascend C 实现 `aclnnPolar`，对齐开源仓 l0 拼接参考实现并**新增广播支持**，验收通过后贡献至 `cann/ops-math`。
6+ 
7+## 背景介绍
8+ 
9+`torch.polar(abs, angle)` 由极坐标构造复数：`out = input·(cos(angle)+i·sin(angle))`，即 `out.real=input·cos(angle)`、`out.imag=input·sin(angle)`。
10+ 
11+### Polar 算子实现优化
12+ 
13+> 本算子**无 TBE 历史实现**（任务书："暂无 tbe 实现，和 `aclnn_polar.cpp` 对齐即可"），以**开源仓 l0 拼接参考实现**为对齐基准与性能基线。
14+ 
15+参考源码：`cann/ops-math` 仓 `math/complex/op_host/op_api/aclnn_polar.cpp`（https://gitcode.com/cann/ops-math ），依赖 l0 API `sin.h/cos.h/mul.h/complex.h`（`l0op::Complex` 为已注册设备算子）。系统 `libopapi.so` 已导出 `aclnnPolar`/`aclnnPolarGetWorkspaceSize`。
16+ 
17+### Polar 参考（l0）实现现状分析
18+ 
19+| 参数 | 含义 | 支持 dtype | 约束 | 形状 |
20+| --- | --- | --- | --- | --- |
21+| input | 模长 abs | float32 | 与 angle 同 dtype | ≤8维 |
22+| angle | 角度(弧度) | float32 | 与 input 同 dtype | ≤8维，可与 input 广播 |
23+| out | 复数结果 | complex64 | 恒 complex64 | broadcast(input,angle) |
24+ 
25+**实现逻辑**（与源码一致）：`aclnnPolarGetWorkspaceSize` 为纯 host 端 l0 拼接 —— `CheckParams`(空指针/dtype/广播 shape) → `Contiguous`×2 → `Sin(angle)`、`Cos(angle)` → `Mul(·,input)`×2 得虚/实部 → `Complex` → `ViewCopy`；约束 `DTYPE={DT_FLOAT}`、`OUTPUT={DT_COMPLEX64}`、dtype 须匹配、`MAX_DIM=8`、广播经 `OP_CHECK_BROADCAST_AND_INFER_SHAPE`。
26+ 
27+**核心问题**：拆成 `Contiguous→Sin→Cos→Mul→Mul→Complex` 约 6 个独立设备算子，每个有 kernel 启动 dispatch 开销，中间结果必经 GM↔UB 搬运（≥10 次），小 shape 下启动延迟远超计算、带宽浪费。Ascend C 单 kernel 融合可解：CopyIn 一次搬入、Compute 在 UB 链式完成 `Cos→Sin→Mul→交织`、中间结果驻留 UB、CopyOut 一次写出 —— **1 次启动、2 次 GM↔UB 搬运**，广播在单 kernel CopyIn 内完成。
28+ 
29+**参考（l0）实现流程图**（与源码逻辑一致，重点关注）：
30+ 
31+```mermaid
32+flowchart TD
33+ A["input, angle, out"] --> B{"CheckParams: null/dtype/bcast"}
34+ B -->|Fail| Z["error code"]
35+ B -->|Pass| C{"empty?"}
36+ C -->|Yes| Y["ws=0, return"]
37+ C -->|No| D["Contiguous x2"]
38+ D --> F["Sin(angle)"]
39+ D --> G["Cos(angle)"]
40+ F --> H["Mul -> imag"]
41+ G --> I["Mul -> real"]
42+ H --> J["Complex -> output"]
43+ I --> J
44+ J --> K["ViewCopy -> out"]
45+```
46+ 
47+### 算子功能规格与原型
48+ 
49+| 名称 | 类别 | dtype | format | 介绍 |
50+|------|------|-------|--------|------|
51+| input | 输入 | float32 | ND | 极坐标模长 abs |
52+| angle | 输入 | float32 | ND | 极坐标幅角(弧度)，dtype 须与 input 一致 |
53+| out | 输出 | complex64 | ND | 复数 tensor，real/imag 交替存储 |
54+ 
55+数学公式 `out=input·cos(angle)+i·input·sin(angle)`；无属性(attr)。支持广播（input 与 angle 维度可不一致，numpy 规则；**本任务新增功能点**）。
56+ 
57+### 相关约束
58+ 
59+Atlas A2 训练系列产品(ascend910b) / Atlas A3 训练系列产品(ascend910_93) 双平台支持；input 与 angle dtype 必须一致(float32)；≤8 维；angle 为弧度无范围限制(周期函数)；不支持 fp16/bf16/fp64/complex128（任务范围外，参考亦不支持）。
60+ 
61+---
62+ 
63+# 需求分析（required）
64+ 
65+## 外部组件依赖
66+ 
67+不涉及外部组件依赖。
68+ 
69+## 内部适配模块
70+ 
71+适配 Aclnn 两段式接口直调：`aclnnPolarGetWorkspaceSize(input, angle, out, *workspaceSize, **executor)` 与 `aclnnPolar(workspace, workspaceSize, executor, stream)`。
72+ 
73+## 需求描述
74+ 
75+使用 Ascend C 将 `Cos+Sin+Mul+复数构造` 融合为单 kernel，减少 GM↔UB 搬运与 kernel 启动开销；**新增 NumPy 广播语义**（input 与 angle 维度可不一致）。
76+ 
77+## 需求拆解
78+ 
79+1. 实现 `Cos+Sin+Mul×2+复数交织` 全融合 kernel，消除小算子拼接瓶颈
80+2. 广播迁移到 Kernel 侧：host 算广播 shape 与 stride，kernel 矢量化索引映射加载，统一在单 kernel CopyIn 完成
81+3. fp32(input/angle)→complex64(out)，与参考一致，≤8 维
82+4. 性能 ≥ 参考 l0 的 95%（"所有核参与场景"为核心验收）；精度满足 AscendOpTest 默认阈值
83+ 
84+## AscendC 算子原型
85+ 
86+| 名称 | 类别 | dtype | format | shape | 介绍 |
87+| --- | --- | --- | --- | --- | --- |
88+| input | 输入 | float32 | ND | ≤8维 | 模长 abs |
89+| angle | 输入 | float32 | ND | ≤8维，可广播 | 角度(弧度) |
90+| out | 输出 | complex64 | ND | broadcast(input,angle) | 复数结果 |
91+ 
92+无 attr；input/angle 须同 dtype(float32)；out 恒 complex64。
93+ 
94+## AscendC 算子相关约束（与参考相比缺失功能）
95+ 
96+无功能缺失：参考支持范围（fp32→complex64 + numpy 广播 + ≤8 维）本实现全部对齐，并原生实现广播（任务新增点）。fp64→complex128 任务书未要求、参考亦不支持，按范围不适配。
97+ 
98+---
99+ 
100+# 详细设计（required）
101+ 
102+## 算子分析
103+ 
104+**数学公式**：`out.real=input·cos(angle)`，`out.imag=input·sin(angle)`。
105+ 
106+**算子特性**：逐元素、各元素独立（天然多核并行）；确定性计算（无 Reduce/Scatter）；每输出元素 2 次三角函数 + 2 次乘法、读 2 输入写 1 复数；输出元素间无依赖（无需跨核同步，单核流水线满排）。
107+ 
108+**支持数据类型**：input、angle = float32（须同 dtype）；out = complex64。910B vector 不原生支持 complex64，经 `ReinterpretCast` 视作 2N 个交织 fp32。
109+ 
110+**支持形状**：ND，≤8 维；input 与 angle 满足 numpy 广播，`out.shape=broadcast(input,angle)`。
111+ 
112+## 算子实现
113+ 
114+总体思路：将参考约 6 个串接 l0 设备算子**融合为单 Ascend C 算子**，一趟 UB 内完成 `Cos/Sin+Mul×2+复数交织`，消除中间 HBM 往返与多次 dispatch（性能提升来源）。算子分 Host 侧（参数校验、广播 shape/stride 计算、Tiling 下发）与 Kernel 侧（Init + Process 三段式 CopyIn→Compute→CopyOut）。
115+ 
116+### 3.2.1 host 侧设计
117+ 
118+**InferShape/InferDataType**：InferShape 按 numpy 广播推导 out shape（右对齐各轴取 max，对齐参考 `OP_CHECK_BROADCAST_AND_INFER_SHAPE`）；InferDataType 将 out 恒置 `DT_COMPLEX64`。
119+ 
120+**广播 Shape 计算**：右对齐、缺失轴首部补 1；逐维相等或其一为 1 否则失败；每维取 max。例：`[3,1,4]`+`[2,4]`→`[3,2,4]`。host 对 input/angle 分别算其在各输出轴的元素 stride（广播轴=0），并判定全形快路径标志 `操作数 numel==out numel ⟺ 未广播`。
121+ 
122+**1. 分核策略**：优先满核。取平台 AIV 核数；**小 case 少核**（每核 ≥`MIN_PER_CORE=2048` 元素，降 launch/同步开销）；**强制偶数核**（vector core 两两绑定，>2 时下取偶）。能均分则无大小核；不能均分则余块分到前几核（前 `rem` 核各 `per+1`，其余 `per`）。
123+ 
124+**2. 数据分块与内存优化**：充分利用 UB（910B=192KB）。`BUFFER_NUM=1`，尾块同路径（DataCopyPad 任意字节无碎片）。tileLen 计算：固定占用 ≈52×tileLen 字节，分布于 8 个 UB buffer —— `qAbs/qAng` 各 `tileLen·4B`(VECIN)、`qOut/bufPacked/bufOff` 各 `2·tileLen·4B`(复数/实虚 packed/交织偏移表)、`bufBC` `5·tileLen·4B`(矢量化 unravel scratch)、`bufInFull/bufAnFull` 广播操作数 numel×4B（全形时仅 32B 占位）。取 **tileLen=2048**（固定 ≈104KB ≤192KB，余量留广播操作数整块载入）。tile 循环 `while(done<coreLen){n=min(coreLen-done,tileLen);…done+=n;}`。
125+> 注：TPipe 单 kernel InitBuffer 张量上限 8（BUFFER_NUM=2 每个计 2）；本算子 BUFFER_NUM=1、8 buffer 恰在上限（超限直接 run failed 无报错）。
126+ 
127+**3. tilingkey 策略**：**单 kernel 全路径**，全形/广播按 tiling 的 `inFull/anFull` **运行期判断**，**不设多 tilingkey**（统一 0）—— 避免多 funcEntry 编译/维护成本，轻量分支无性能损失。仅 fp32→complex64，OpDef 已限 dtype，无 bf16 兜底。
128+ 
129+**4. TilingData 字段**：`totalLen, tileLen, bigCoreNum/coreNum(偶数), bigCoreLen/smallCoreLen, outRank, inFull/anFull(全形快路径标志), inNumel/anNumel(广播整块载入用), outDims[8], inStr[8]/anStr[8](各输出轴元素 stride，广播轴=0)`（`outDims/inStr/anStr` 经 `TILING_DATA_FIELD_DEF_ARR`）。
130+ 
131+### 3.2.2 kernel 侧设计
132+ 
133+单 kernel，Init + Process（CopyIn/Compute/CopyOut），全程 fp32。
134+ 
135+- **Init**：按 blockIdx 算本核区间；`BuildOffsetTable()` 纯 int32 矢量构造交织静态偏移表 `off[j]=4·(j>>1)+(j&1)·4T`（一次性、全 tile 复用、无逐元素标量）；广播操作数经 `LoadFull` 整块载入 UB。
136+- **CopyIn(FetchOperand)**：全形操作数→连续 `DataCopyPad` 快路径；广播操作数→**矢量化 unravel**（对 tile 输出线性 idx，逐输出轴 `Muls(1/D)+Floor` 求商 + 两次 `CompareScalar+Select` 精确 ±1 修正算坐标，累加偏移 → `Cast` int32 → `Gather` 从 UB 整块操作数取数，全矢量无逐元素标量）。
137+- **Compute**：`Cos(angle)`→packed 实部区、`Sin(angle)`→packed 虚部区 → 原地 `Mul ×abs` → `Gather` 用静态偏移表交织为 `(r,i,r,i,…)`。
138+- **CopyOut**：`DataCopyPad` 写 complex64 GM（视作 2N fp32 大块对齐）。
139+ 
140+**AscendC 实现流程图**（重点关注）：
141+ 
142+```mermaid
143+flowchart TD
144+ S["polar(): GET_TILING_DATA; Init; Process"] --> I1["Init: read TilingData + this-core range"]
145+ I1 --> I2["Init: InitBuffer x8 + BuildOffsetTable (once)"]
146+ I2 --> I3{"input/angle need broadcast?"}
147+ I3 -->|Yes| I4["LoadFull operand into UB"]
148+ I3 -->|No| L0["tile loop"]
149+ I4 --> L0
150+ L0 --> CI["CopyIn: FetchOperand abs/angle"]
151+ CI --> B1{"operand == output shape?"}
152+ B1 -->|Yes| F1["DataCopyPad contiguous (fast path)"]
153+ B1 -->|No| U1["vectorized unravel: Muls/Floor + CompareScalar/Select correction -> Cast int32 -> Gather"]
154+ F1 --> CM["Compute: Cos->real, Sin->imag; Mul x abs; Gather interleave (r,i,r,i)"]
155+ U1 --> CM
156+ CM --> CO["CopyOut: DataCopyPad complex64 GM"]
157+ CO --> CK{"all tiles done?"}
158+ CK -->|No| L0
159+ CK -->|Yes| E["End"]
160+```
161+ 
162+**复数交织方案选型**（910B/dav_c220 输出需将分离 real/imag 交织为 `[r0,i0,r1,i1,…]`）：
163+ 
164+| 方案 | 结果 | 原因 |
165+|------|------|------|
166+| `Copy<float,true>` mask | ❌ | mask 同时作用 src/dst，无法错位 |
167+| `Transpose` NCHW2NHWC | ✅但慢 | 受 16×16 分形约束，占绝大部分耗时 |
168+| `vintlv`/`Interleave` | ❌ | dav_c310+ 独有，c220 不支持 |
169+| strided `DataCopyPad` | ❌ | VECOUT 32B 粒度，逐 4B 写出 MTE 灾难 |
170+| **Gather+静态偏移表** | **✅选用** | 纯向量指令，偏移表 Init 一次构造全 tile 复用 |
171+ 
172+**与参考流程图的差异点及原因**：
173+ 
174+| # | 参考 l0 | 本 Ascend C | 原因 |
175+|---|---|---|---|
176+| 1 | 约 6 独立设备算子串接，各 HBM 往返+dispatch | 单 kernel 融合 | 省 ≥5 次中间 HBM+多 dispatch（核心场景快 1.78×） |
177+| 2 | `l0op::Complex` 构造复数 | Gather+静态偏移表交织 | 910B vector 不支持 complex64；规避 strided/Transpose 约束 |
178+| 3 | 广播由 l0 Mul 内部处理 | host 算 stride + kernel 矢量化 unravel | 原生融合需自实现广播；全矢量避免标量 |
179+| 4 | Contiguous+ViewCopy 处理非连续 | 框架传连续 GM 无需 | aclnn/测试输入连续，省拷贝 |
180+ 
181+**融合优势**：kernel 启动 约 6 次→1 次；GM↔UB 搬运 ≥10 次→2 次；中间结果 GM→UB（片上）；angle 一次读入 cos+sin 同 tile 计算（参考无法复用）。
182+ 
183+## 支持硬件
184+ 
185+| 芯片版本 | 勾选 |
186+| --- | --- |
187+| Atlas A2 训练系列产品 / Atlas A2 推理系列产品 (ascend910b) | √ |
188+| Atlas A3 训练系列产品 / Atlas A3 推理系列产品 (ascend910_93) | √ |
189+ 
190+## 使能方式
191+ 
192+| 上层框架 | 勾选 |
193+| --- | :---: |
194+| TF 训练/推理 | |
195+| Pytorch 训练/推理 | √（pybind 调 aclnn 测试） |
196+| ATC 推理 | |
197+| **Aclnn 直调** | **√** |
198+| OPAT 调优 | |
199+| SGAT 子图切分 | |
200+ 
201+主要适配 **ACLNN** 调用框架。
202+ 
203+## 算子约束限制
204+ 
205+- dtype：input/angle 仅 float32（须同 dtype）；out 仅 complex64。不支持 fp16/bf16/fp64/complex128
206+- ≤8 维；input/angle 须满足 numpy 广播
207+- 广播操作数 numel 需可整体载入 UB（验收 shape 量级满足；超大广播待分块优化）
208+- 无 attr
209+ 
210+# 特性交叉分析
211+ 
212+逐元素计算，输出仅依赖对应位置输入，不涉及跨元素聚合(Reduce/Scatter)、数据搬移(Reshape/Transpose 语义)、量化(Quant)，与现有特性无冲突。广播 ×（rank/shape/角度范围/正负 abs/32B 对齐性）已交叉验证：8D 满秩同形、5D 中间轴广播、双向多轴广播、标量 input×高维 angle、负 abs、大角度归约、非 32B 对齐 inner+广播、大向量×标量（用例 case9–16），均严格判据 PASS，无异常。
213+ 
214+---
215+ 
216+# 可维可测分析
217+ 
218+## 精度标准/性能标准
219+ 
220+| 验收标准 | 描述 | 来源 |
221+| --- | --- | --- |
222+| 精度 | 满足 AscendOpTest 默认阈值（不低于参考）：fp32→complex64 实/虚部各 max_abs=1e-4、rate=1e-4 | AscendOpTest（HIT1920/AscendOpTest） |
223+| 性能 | 所有核参与场景 ≥ 参考 l0 的 95%；小 shape 提供仿真分析 | 《算子任务书》 |
224+ 
225+**精度细化与实测**：complex64 在 AscendOpTest `accuracy_config` 无内置默认 → 用例 JSON 显式配 `err_threshold=[1e-4,1e-4]`（fp32 分量默认）；`compare_complex` 为实/虚部各自纯绝对误差判定。**官方 AscendOpTest 实跑 6 用例（同 shape 小/16M + 广播低→高/标量/双向 + 高维非对齐）全 PASS**（含 16M [4096,4096]）。全程 fp32 计算不引入 Cast 中间精度损失，complex64 经 ReinterpretCast 视作交织 fp32 分量精度等价 fp32。
226+ 
227+**性能实测**（vs 系统 l0 参考基线，官方 AscendOpTest msprof / 每调用设备时）：
228+ 
229+| 场景 | l0 基线 | 本算子 | 结论 |
230+|---|---:|---:|---|
231+| 小 [2,6,10] | 14.25 µs | 12.52 µs | 更快 ✓ |
232+| **16M [4096,4096]（所有核）** | **1665 µs** | **937.98 µs** | **快 1.78×，核心验收达标** ✓✓ |
233+| 广播 [4,1,8]×[4,5,8] | 15.17 µs | 14.92 µs | 更快 ✓ |
234+ 
235+全场景优于 l0 参考，远超 ≥95% 要求。
236+ 
237+## 兼容性分析
238+ 
239+新算子（昇腾仓原无 Ascend C 原生 Polar，亦无 TBE），不涉及历史兼容性。aclnn 接口遵循 CANN 标准规范，提供与 PyTorch `torch.polar` 等效功能，支持模型迁移。
Aexperimental/math/polar/docs/性能优化_inner_broadcast.md+104-0
@@ -0,0 +1,104 @@
1+# Polar inner-broadcast 性能优化（方案 A）
2+ 
3+> 2026-05-21。本轮针对 acceptance 50-case 中**唯一性能不达 baseline 的两个 case**（Test_036 / Test_044）
4+> 做 angle inner-broadcast 专用优化，从输基线 0.8x 反超到 1.7-3.5x，并使所有 angle 周期广播 case 普遍 ~2x 提速。
5+ 
6+## 1. 问题定位
7+ 
8+测试口径：pybind + EXEC_NPU_CMD(aclnnPolar)，msprof op_summary Task Duration（每调用设备时，50 次取中位）。
9+baseline = `torch.polar`（NPU 内置，Sin/Cos/Mul/Complex 基础算子拼接）。
10+ 
11+优化前全 50 case 仅这两个输基线：
12+ 
13+| Case | input × angle | out 元素 | 优化前 ours | baseline | vs base |
14+|---|---|---:|---:|---:|---:|
15+| Test_036 | [256,256,64] × **[64]** | 4M | 168.8 µs | 147.2 µs | **0.87x** ❌ |
16+| Test_044 | [128,128,512] × **[512]** | 8M | 331.4 µs | 276.6 µs | **0.83x** ❌ |
17+ 
18+**根因**：这两个是 "angle 为 input trailing-dim 的周期重复" 广播。原实现在 op_api 层用 `BroadcastTo`
19+把 angle 物理展开到 out shape（同 shape kernel 要求），于是：
20+- Test_036：先写 16 MB（angle 展开）→ kernel 再读 input 16 MB + 读展开后 angle 16 MB + 写 out 32 MB = **80 MB HBM**
21+- baseline 用 view（stride=0）读 angle，省掉展开的一次 HBM 写，故反而更快。
22+ 
23+## 2. 方案 A：inner-broadcast 周期 pre-compute（kernel 内运行时分支，无 tilingKey）
24+ 
25+**核心观察**：angle 周期长度 K（= angle.numel），每个 tile 取整数个 period →
26+**cos/sin tile pattern 对所有 tile 完全相同**。
27+ 
28+实现（三处资格判定严格一致：op_api / op_host tiling / pybind）：
29+```
30+资格 = inN==outN && anN<inN && inN%anN==0 && anN<=2048 && anN%8==0
31+```
32+ 
33+- **op_api (`aclnn_polar.cpp`)**：资格命中时 angle **不 BroadcastTo**，保留原 shape [K] 传 kernel
34+- **op_host (`polar_tiling.cpp`)**：识别 `bcastMode=1`，按 **K-block 分核**（每核负载 K 整数倍 → tile 起点 period 对齐）
35+- **op_kernel (`polar.h`)**：`Process()` 内 `if(bcastMode==1)` 分支（纯 data 字段，无需多 kernel 注册）
36+ - `PrecomputeBcastCosSin()`：读 angle[0:K] → Cos/Sin 得 cosT/sinT[K] → 循环 m=T/K 次 broadcast 填满
37+ 常驻 buffer（cosTile[0:T] / sinTile[T:2T]），**只做一次**
38+ - `ProcessTileBcast()`：main loop **只读 input** → `Mul(cosTile,input)`→real / `Mul(sinTile,input)`→imag
39+ → Gather 交织 → 写 out。无 angle 搬运、无每 tile Sin/Cos 重算。
40+ 
41+**HBM/算力收益（Test_036, 4M）**：80 MB → **48 MB**（省 angle 展开 16 MB + angle 读 16 MB→256 B）；
42+Sin/Cos 从"每 tile 重算"降到"一次 precompute"。
43+ 
44+> ⚠️ 实现陷阱：`DataCopy(srcStride=0)` **不是** "重复同段 broadcast"——它连续读 src，会把 cosT 之后的
45+> sinT 读进第二 period（症状 `npu.real==golden.imag`）。正解是把 broadcast 提到 Precompute 一次性循环填满。
46+ 
47+## 3. 优化结果
48+ 
49+### 关注的两个 case：从输基线反超
50+ 
51+| Case | K | 优化前 | **优化后** | kernel 提速 | baseline | vs base 前→后 |
52+|---|---|---:|---:|---:|---:|---:|
53+| Test_036 | 64 | 168.8 | **42.1** | **4.0x** | 147.2 | 0.87x → **3.50x** |
54+| Test_044 | 512 | 331.4 | **161.9** | 2.0x | 276.6 | 0.83x → **1.71x** |
55+ 
56+### 其它 angle 周期广播 case（顺带普遍 ~2x）
57+ 
58+| Case | out 元素 | K | 优化前 | 优化后 | baseline | vs base |
59+|---|---:|---:|---:|---:|---:|---:|
60+| Test_035 | 262K | 64 | 17.6 | 15.7 | 54.7 | 3.47x |
61+| Test_037 | 16M | 64 | 654 | 321 | 765 | 2.38x |
62+| Test_038 | 67M | 64 | 2595 | 1261 | 3340 | 2.65x |
63+| Test_039 | 134M | 32 | 5181 | 2531 | 6771 | 2.67x |
64+| Test_040 | 268M | 16 | 10354 | 5144 | 13633 | 2.65x |
65+| Test_041 | 537M | 8 | 20701 | 10287 | 27361 | 2.66x |
66+| Test_047 | 134M | 8 | 5181 | 2581 | 6784 | 2.63x |
67+| Test_048 | 33M | 32 | 1301 | 641 | 1639 | 2.56x |
68+ 
69+未启用优化（K=1/2/4 不满足 32B 对齐，走 same-shape 兜底）的 Test_042/045/046：因规模大（268-537M）
70+baseline 同为 HBM bound，same-shape 已 1.31-1.47x 胜基线，无需再优化。
71+ 
72+### 全 50-case 结论：无一输基线
73+ 
74+| 类别 | case | 加速比 |
75+|---|---|---|
76+| 小同 shape (1-12) | same-shape | 2.5-3.4x |
77+| 大同 shape (13-29) | same-shape | 2.5-3.1x |
78+| 小广播 (30-34, K<8) | same-shape 兜底 | 3.2-5.3x |
79+| 大广播 (35-50) | bcast 资格 ~2x kernel 提速；非资格规模大也胜 | 1.3-3.5x |
80+ 
81+## 4. 官方 AscendOpTest 验证（ops-math 工程，customize_math vendor）
82+ 
83+```
84+python run_test.py -i .../acceptance/op.json -c <Test_036,Test_044> \
85+ --op-type custom --op-path $ASCEND_OPP_PATH/vendors/customize_math/op_api --build
86+```
87+结果 `result.csv`：
88+ 
89+| case | compare_result |
90+|---|---|
91+| Test_036 | **pass** |
92+| Test_044 | **pass** |
93+ 
94+精度判据 err_threshold=[1e-3, 1e-3]，官方工具确认数值正确。
95+ 
96+## 5. 改动文件（ops-math）
97+ 
98+- `op_kernel/polar_tiling_data.h`：+ `inN/anN/bcastMode`
99+- `op_host/polar_tiling.cpp`：inner-bcast 检测 + K-block 分核
100+- `op_kernel/polar.h`：`Process()` 分支 + `PrecomputeBcastCosSin()` + `ProcessTileBcast()`
101+- `op_api/aclnn_polar.cpp`：资格命中时 angle 不 BroadcastTo
102+- `op_api/polar.cpp`：注释更正
103+ 
104+同 shape 与非 inner-bcast 广播路径（中间轴 / input 侧广播，如 Test_043/049/050）行为不变，由 same-shape kernel + op_api BroadcastTo 兜底。
Aexperimental/math/polar/examples/test_aclnn_polar.cpp+207-0
@@ -0,0 +1,207 @@
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+ * Polar 算子独立 aclnn 调用测试（精度 + 性能）
13+ *
14+ * 参考官方范式：cann/ops-math math/sin/examples/test_aclnn_sin.cpp（两段式 aclnn 标准调用）
15+ * 适配 Polar：2×float32 输入(input, angle) + complex64 输出(out=broadcast(input,angle))
16+ *
17+ * 编译（在已安装自定义 Polar 包的 NPU 环境，路径按实际 CANN/vendor 调整）：
18+ * source /home/ma-user/Ascend/cann-8.5.0/set_env.sh
19+ * OPP=$ASCEND_HOME_PATH/opp/vendors/customize
20+ * g++ -O2 -std=c++17 test_aclnn_polar.cpp -o test_aclnn_polar \
21+ * -I$ASCEND_HOME_PATH/include -I$OPP/op_api/include \
22+ * -L$ASCEND_HOME_PATH/lib64 -L$OPP/op_api/lib \
23+ * -lascendcl -lnnopbase -lcust_opapi
24+ * 运行：
25+ * export LD_LIBRARY_PATH=$OPP/op_api/lib:$LD_LIBRARY_PATH
26+ * ./test_aclnn_polar # 精度 + 自带计时
27+ * msprof --application="./test_aclnn_polar" # 设备侧性能采集
28+ */
29+#include <iostream>
30+#include <vector>
31+#include <cmath>
32+#include <chrono>
33+#include "acl/acl.h"
34+#include "aclnn_polar.h" // 自定义算子生成的 aclnn 头（custom vendor op_api/include）
35+ 
36+#define CHECK_RET(cond, expr) \
37+ do { \
38+ if (!(cond)) { \
39+ expr; \
40+ } \
41+ } while (0)
42+#define LOG(msg, ...) \
43+ do { \
44+ printf(msg, ##__VA_ARGS__); \
45+ } while (0)
46+ 
47+struct C64 {
48+ float re;
49+ float im;
50+}; // 与 complex64 内存布局一致（8 字节）
51+ 
52+static int64_t ShapeSize(const std::vector<int64_t>& s)
53+{
54+ int64_t n = 1;
55+ for (auto v : s)
56+ n *= v;
57+ return n;
58+}
59+ 
60+static int Init(int32_t deviceId, aclrtStream* stream)
61+{
62+ auto ret = aclInit(nullptr);
63+ CHECK_RET(ret == ACL_SUCCESS, LOG("aclInit failed %d\n", ret); return ret);
64+ ret = aclrtSetDevice(deviceId);
65+ CHECK_RET(ret == ACL_SUCCESS, LOG("aclrtSetDevice failed %d\n", ret); return ret);
66+ ret = aclrtCreateStream(stream);
67+ CHECK_RET(ret == ACL_SUCCESS, LOG("aclrtCreateStream failed %d\n", ret); return ret);
68+ return 0;
69+}
70+ 
71+template <typename T>
72+static int CreateAclTensor(const std::vector<T>& host, const std::vector<int64_t>& shape, void** devAddr,
73+ aclDataType dt, aclTensor** tensor)
74+{
75+ auto bytes = ShapeSize(shape) * sizeof(T);
76+ auto ret = aclrtMalloc(devAddr, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
77+ CHECK_RET(ret == ACL_SUCCESS, LOG("aclrtMalloc failed %d\n", ret); return ret);
78+ ret = aclrtMemcpy(*devAddr, bytes, host.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE);
79+ CHECK_RET(ret == ACL_SUCCESS, LOG("aclrtMemcpy H2D failed %d\n", ret); return ret);
80+ std::vector<int64_t> strides(shape.size(), 1);
81+ for (int64_t i = (int64_t)shape.size() - 2; i >= 0; --i)
82+ strides[i] = shape[i + 1] * strides[i + 1];
83+ *tensor = aclCreateTensor(shape.data(), shape.size(), dt, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(),
84+ shape.size(), *devAddr);
85+ return 0;
86+}
87+ 
88+// numpy 右对齐广播 out shape
89+static std::vector<int64_t> BroadcastShape(const std::vector<int64_t>& a, const std::vector<int64_t>& b)
90+{
91+ size_t r = std::max(a.size(), b.size());
92+ std::vector<int64_t> o(r, 1);
93+ for (size_t i = 0; i < r; ++i) {
94+ int64_t da = (i + a.size() >= r) ? a[i + a.size() - r] : 1;
95+ int64_t db = (i + b.size() >= r) ? b[i + b.size() - r] : 1;
96+ o[i] = std::max(da, db);
97+ }
98+ return o;
99+}
100+ 
101+// 一组用例：构造 input/angle，调用 aclnnPolar（warmup + 计时重复），校验 + 报告
102+static int RunCase(const char* name, aclrtStream stream, const std::vector<int64_t>& inShape,
103+ const std::vector<int64_t>& anShape, int repeat = 100)
104+{
105+ auto outShape = BroadcastShape(inShape, anShape);
106+ int64_t inN = ShapeSize(inShape), anN = ShapeSize(anShape), outN = ShapeSize(outShape);
107+ 
108+ std::vector<float> inH(inN), anH(anN);
109+ for (int64_t i = 0; i < inN; ++i)
110+ inH[i] = 0.5f + (float)(i % 17); // abs ≥ 0 区间
111+ for (int64_t i = 0; i < anN; ++i)
112+ anH[i] = -3.1415926f + (float)(i % 11); // 角度
113+ std::vector<C64> outH(outN, {0.f, 0.f});
114+ 
115+ void *inD = nullptr, *anD = nullptr, *outD = nullptr;
116+ aclTensor *in = nullptr, *an = nullptr, *out = nullptr;
117+ CHECK_RET(CreateAclTensor(inH, inShape, &inD, ACL_FLOAT, &in) == 0, return -1);
118+ CHECK_RET(CreateAclTensor(anH, anShape, &anD, ACL_FLOAT, &an) == 0, return -1);
119+ CHECK_RET(CreateAclTensor(outH, outShape, &outD, ACL_COMPLEX64, &out) == 0, return -1);
120+ 
121+ uint64_t wsSize = 0;
122+ aclOpExecutor* exec = nullptr;
123+ auto ret = aclnnPolarGetWorkspaceSize(in, an, out, &wsSize, &exec);
124+ CHECK_RET(ret == ACL_SUCCESS, LOG("[%s] GetWorkspaceSize failed %d\n", name, ret); return ret);
125+ void* wsAddr = nullptr;
126+ if (wsSize > 0) {
127+ ret = aclrtMalloc(&wsAddr, wsSize, ACL_MEM_MALLOC_HUGE_FIRST);
128+ CHECK_RET(ret == ACL_SUCCESS, LOG("[%s] ws malloc failed %d\n", name, ret); return ret);
129+ }
130+ 
131+ // warmup
132+ for (int i = 0; i < 10; ++i) {
133+ ret = aclnnPolar(wsAddr, wsSize, exec, stream);
134+ CHECK_RET(ret == ACL_SUCCESS, LOG("[%s] aclnnPolar failed %d\n", name, ret); return ret);
135+ }
136+ aclrtSynchronizeStream(stream);
137+ 
138+ // 计时（壁钟，msprof 另采设备时）
139+ auto t0 = std::chrono::high_resolution_clock::now();
140+ for (int i = 0; i < repeat; ++i) {
141+ ret = aclnnPolar(wsAddr, wsSize, exec, stream);
142+ CHECK_RET(ret == ACL_SUCCESS, LOG("[%s] aclnnPolar failed %d\n", name, ret); return ret);
143+ }
144+ aclrtSynchronizeStream(stream);
145+ auto t1 = std::chrono::high_resolution_clock::now();
146+ double us = std::chrono::duration<double, std::micro>(t1 - t0).count() / repeat;
147+ 
148+ // 读回 + CPU 参考校验（polar = abs*(cos+ i sin)，按广播规则）
149+ aclrtMemcpy(outH.data(), outN * sizeof(C64), outD, outN * sizeof(C64), ACL_MEMCPY_DEVICE_TO_HOST);
150+ int64_t err = 0;
151+ for (int64_t g = 0; g < outN; ++g) {
152+ // 由 g 反解广播下的 input/angle 元素（行主序，size==1 维不进位）
153+ int64_t inIdx = 0, anIdx = 0, accIn = 1, accAn = 1, rem = g;
154+ for (int64_t d = (int64_t)outShape.size() - 1; d >= 0; --d) {
155+ int64_t c = rem % outShape[d];
156+ rem /= outShape[d];
157+ int64_t di = (d + (int64_t)inShape.size() >= (int64_t)outShape.size()) ?
158+ inShape[d + inShape.size() - outShape.size()] :
159+ 1;
160+ int64_t da = (d + (int64_t)anShape.size() >= (int64_t)outShape.size()) ?
161+ anShape[d + anShape.size() - outShape.size()] :
162+ 1;
163+ if (di != 1)
164+ inIdx += (c % di) * accIn;
165+ if (da != 1)
166+ anIdx += (c % da) * accAn;
167+ accIn *= di;
168+ accAn *= da;
169+ }
170+ float a = inH[inIdx % inN], th = anH[anIdx % anN];
171+ float er = a * std::cos(th), ei = a * std::sin(th);
172+ if (std::fabs(er - outH[g].re) > 1e-4f || std::fabs(ei - outH[g].im) > 1e-4f)
173+ ++err;
174+ }
175+ LOG("[%-22s] in=%lldel an=%lldel out=%lldel avg=%.2f us/call err=%lld/%lld %s\n", name, (long long)inN,
176+ (long long)anN, (long long)outN, us, (long long)err, (long long)outN, (err <= outN * 1e-4 ? "PASS" : "FAIL"));
177+ 
178+ aclDestroyTensor(in);
179+ aclDestroyTensor(an);
180+ aclDestroyTensor(out);
181+ aclrtFree(inD);
182+ aclrtFree(anD);
183+ aclrtFree(outD);
184+ if (wsSize > 0)
185+ aclrtFree(wsAddr);
186+ return 0;
187+}
188+ 
189+int main()
190+{
191+ int32_t deviceId = 0;
192+ aclrtStream stream;
193+ auto ret = Init(deviceId, &stream);
194+ CHECK_RET(ret == ACL_SUCCESS, LOG("Init acl failed %d\n", ret); return ret);
195+ 
196+ // 多组用例：同 shape（小/大-性能）+ 广播（多种模式）
197+ RunCase("same-small [2,6,10]", stream, {2, 6, 10}, {2, 6, 10});
198+ RunCase("same-large [4096,4096]", stream, {4096, 4096}, {4096, 4096}, 50);
199+ RunCase("bcast [4,1,8]x[4,5,8]", stream, {4, 1, 8}, {4, 5, 8});
200+ RunCase("bcast scalar [1]x[3,4,5]", stream, {1}, {3, 4, 5});
201+ RunCase("bcast [8,1]x[1,7]", stream, {8, 1}, {1, 7});
202+ 
203+ aclrtDestroyStream(stream);
204+ aclrtResetDevice(deviceId);
205+ aclFinalize();
206+ return 0;
207+}
Aexperimental/math/polar/examples/test_aclnn_polar_ref.cpp+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+ * Polar 基线对比 REF 变体：调用系统 l0 参考 aclnnPolar（同测试/同计时）
13+ *
14+ * 参考官方范式：cann/ops-math math/sin/examples/test_aclnn_sin.cpp（两段式 aclnn 标准调用）
15+ * 适配 Polar：2×float32 输入(input, angle) + complex64 输出(out=broadcast(input,angle))
16+ *
17+ * 编译（在已安装自定义 Polar 包的 NPU 环境，路径按实际 CANN/vendor 调整）：
18+ * source /home/ma-user/Ascend/cann-8.5.0/set_env.sh
19+ * OPP=$ASCEND_HOME_PATH/opp/vendors/customize
20+ * g++ -O2 -std=c++17 test_aclnn_polar.cpp -o test_aclnn_polar \
21+ * -I$ASCEND_HOME_PATH/include -I$OPP/op_api/include \
22+ * -L$ASCEND_HOME_PATH/lib64 -L$OPP/op_api/lib \
23+ * -lascendcl -lnnopbase -lcust_opapi
24+ * 运行：
25+ * export LD_LIBRARY_PATH=$OPP/op_api/lib:$LD_LIBRARY_PATH
26+ * ./test_aclnn_polar # 精度 + 自带计时
27+ * msprof --application="./test_aclnn_polar" # 设备侧性能采集
28+ */
29+#include <iostream>
30+#include <vector>
31+#include <cmath>
32+#include <chrono>
33+#include "acl/acl.h"
34+#include "aclnn/aclnn_base.h"
35+// [REF 变体] 不含自定义头，前向声明 → 链接系统 libopapi.so 的 l0 参考 aclnnPolar
36+extern "C" {
37+aclnnStatus aclnnPolarGetWorkspaceSize(const aclTensor* input, const aclTensor* angle, aclTensor* out,
38+ uint64_t* workspaceSize, aclOpExecutor** executor);
39+aclnnStatus aclnnPolar(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
40+}
41+ 
42+#define CHECK_RET(cond, expr) \
43+ do { \
44+ if (!(cond)) { \
45+ expr; \
46+ } \
47+ } while (0)
48+#define LOG(msg, ...) \
49+ do { \
50+ printf(msg, ##__VA_ARGS__); \
51+ } while (0)
52+ 
53+struct C64 {
54+ float re;
55+ float im;
56+}; // 与 complex64 内存布局一致（8 字节）
57+ 
58+static int64_t ShapeSize(const std::vector<int64_t>& s)
59+{
60+ int64_t n = 1;
61+ for (auto v : s)
62+ n *= v;
63+ return n;
64+}
65+ 
66+static int Init(int32_t deviceId, aclrtStream* stream)
67+{
68+ auto ret = aclInit(nullptr);
69+ CHECK_RET(ret == ACL_SUCCESS, LOG("aclInit failed %d\n", ret); return ret);
70+ ret = aclrtSetDevice(deviceId);
71+ CHECK_RET(ret == ACL_SUCCESS, LOG("aclrtSetDevice failed %d\n", ret); return ret);
72+ ret = aclrtCreateStream(stream);
73+ CHECK_RET(ret == ACL_SUCCESS, LOG("aclrtCreateStream failed %d\n", ret); return ret);
74+ return 0;
75+}
76+ 
77+template <typename T>
78+static int CreateAclTensor(const std::vector<T>& host, const std::vector<int64_t>& shape, void** devAddr,
79+ aclDataType dt, aclTensor** tensor)
80+{
81+ auto bytes = ShapeSize(shape) * sizeof(T);
82+ auto ret = aclrtMalloc(devAddr, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
83+ CHECK_RET(ret == ACL_SUCCESS, LOG("aclrtMalloc failed %d\n", ret); return ret);
84+ ret = aclrtMemcpy(*devAddr, bytes, host.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE);
85+ CHECK_RET(ret == ACL_SUCCESS, LOG("aclrtMemcpy H2D failed %d\n", ret); return ret);
86+ std::vector<int64_t> strides(shape.size(), 1);
87+ for (int64_t i = (int64_t)shape.size() - 2; i >= 0; --i)
88+ strides[i] = shape[i + 1] * strides[i + 1];
89+ *tensor = aclCreateTensor(shape.data(), shape.size(), dt, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(),
90+ shape.size(), *devAddr);
91+ return 0;
92+}
93+ 
94+// numpy 右对齐广播 out shape
95+static std::vector<int64_t> BroadcastShape(const std::vector<int64_t>& a, const std::vector<int64_t>& b)
96+{
97+ size_t r = std::max(a.size(), b.size());
98+ std::vector<int64_t> o(r, 1);
99+ for (size_t i = 0; i < r; ++i) {
100+ int64_t da = (i + a.size() >= r) ? a[i + a.size() - r] : 1;
101+ int64_t db = (i + b.size() >= r) ? b[i + b.size() - r] : 1;
102+ o[i] = std::max(da, db);
103+ }
104+ return o;
105+}
106+ 
107+// 一组用例：构造 input/angle，调用 aclnnPolar（warmup + 计时重复），校验 + 报告
108+static int RunCase(const char* name, aclrtStream stream, const std::vector<int64_t>& inShape,
109+ const std::vector<int64_t>& anShape, int repeat = 100)
110+{
111+ auto outShape = BroadcastShape(inShape, anShape);
112+ int64_t inN = ShapeSize(inShape), anN = ShapeSize(anShape), outN = ShapeSize(outShape);
113+ 
114+ std::vector<float> inH(inN), anH(anN);
115+ for (int64_t i = 0; i < inN; ++i)
116+ inH[i] = 0.5f + (float)(i % 17); // abs ≥ 0 区间
117+ for (int64_t i = 0; i < anN; ++i)
118+ anH[i] = -3.1415926f + (float)(i % 11); // 角度
119+ std::vector<C64> outH(outN, {0.f, 0.f});
120+ 
121+ void *inD = nullptr, *anD = nullptr, *outD = nullptr;
122+ aclTensor *in = nullptr, *an = nullptr, *out = nullptr;
123+ CHECK_RET(CreateAclTensor(inH, inShape, &inD, ACL_FLOAT, &in) == 0, return -1);
124+ CHECK_RET(CreateAclTensor(anH, anShape, &anD, ACL_FLOAT, &an) == 0, return -1);
125+ CHECK_RET(CreateAclTensor(outH, outShape, &outD, ACL_COMPLEX64, &out) == 0, return -1);
126+ 
127+ uint64_t wsSize = 0;
128+ aclOpExecutor* exec = nullptr;
129+ auto ret = aclnnPolarGetWorkspaceSize(in, an, out, &wsSize, &exec);
130+ CHECK_RET(ret == ACL_SUCCESS, LOG("[%s] GetWorkspaceSize failed %d\n", name, ret); return ret);
131+ void* wsAddr = nullptr;
132+ if (wsSize > 0) {
133+ ret = aclrtMalloc(&wsAddr, wsSize, ACL_MEM_MALLOC_HUGE_FIRST);
134+ CHECK_RET(ret == ACL_SUCCESS, LOG("[%s] ws malloc failed %d\n", name, ret); return ret);
135+ }
136+ 
137+ // warmup
138+ for (int i = 0; i < 10; ++i) {
139+ ret = aclnnPolar(wsAddr, wsSize, exec, stream);
140+ CHECK_RET(ret == ACL_SUCCESS, LOG("[%s] aclnnPolar failed %d\n", name, ret); return ret);
141+ }
142+ aclrtSynchronizeStream(stream);
143+ 
144+ // 计时（壁钟，msprof 另采设备时）
145+ auto t0 = std::chrono::high_resolution_clock::now();
146+ for (int i = 0; i < repeat; ++i) {
147+ ret = aclnnPolar(wsAddr, wsSize, exec, stream);
148+ CHECK_RET(ret == ACL_SUCCESS, LOG("[%s] aclnnPolar failed %d\n", name, ret); return ret);
149+ }
150+ aclrtSynchronizeStream(stream);
151+ auto t1 = std::chrono::high_resolution_clock::now();
152+ double us = std::chrono::duration<double, std::micro>(t1 - t0).count() / repeat;
153+ 
154+ // 读回 + CPU 参考校验（polar = abs*(cos+ i sin)，按广播规则）
155+ aclrtMemcpy(outH.data(), outN * sizeof(C64), outD, outN * sizeof(C64), ACL_MEMCPY_DEVICE_TO_HOST);
156+ int64_t err = 0;
157+ for (int64_t g = 0; g < outN; ++g) {
158+ // 由 g 反解广播下的 input/angle 元素（行主序，size==1 维不进位）
159+ int64_t inIdx = 0, anIdx = 0, accIn = 1, accAn = 1, rem = g;
160+ for (int64_t d = (int64_t)outShape.size() - 1; d >= 0; --d) {
161+ int64_t c = rem % outShape[d];
162+ rem /= outShape[d];
163+ int64_t di = (d + (int64_t)inShape.size() >= (int64_t)outShape.size()) ?
164+ inShape[d + inShape.size() - outShape.size()] :
165+ 1;
166+ int64_t da = (d + (int64_t)anShape.size() >= (int64_t)outShape.size()) ?
167+ anShape[d + anShape.size() - outShape.size()] :
168+ 1;
169+ if (di != 1)
170+ inIdx += (c % di) * accIn;
171+ if (da != 1)
172+ anIdx += (c % da) * accAn;
173+ accIn *= di;
174+ accAn *= da;
175+ }
176+ float a = inH[inIdx % inN], th = anH[anIdx % anN];
177+ float er = a * std::cos(th), ei = a * std::sin(th);
178+ if (std::fabs(er - outH[g].re) > 1e-4f || std::fabs(ei - outH[g].im) > 1e-4f)
179+ ++err;
180+ }
181+ LOG("[%-22s] in=%lldel an=%lldel out=%lldel avg=%.2f us/call err=%lld/%lld %s\n", name, (long long)inN,
182+ (long long)anN, (long long)outN, us, (long long)err, (long long)outN, (err <= outN * 1e-4 ? "PASS" : "FAIL"));
183+ 
184+ aclDestroyTensor(in);
185+ aclDestroyTensor(an);
186+ aclDestroyTensor(out);
187+ aclrtFree(inD);
188+ aclrtFree(anD);
189+ aclrtFree(outD);
190+ if (wsSize > 0)
191+ aclrtFree(wsAddr);
192+ return 0;
193+}
194+ 
195+int main()
196+{
197+ int32_t deviceId = 0;
198+ aclrtStream stream;
199+ auto ret = Init(deviceId, &stream);
200+ CHECK_RET(ret == ACL_SUCCESS, LOG("Init acl failed %d\n", ret); return ret);
201+ 
202+ // 多组用例：同 shape（小/大-性能）+ 广播（多种模式）
203+ RunCase("same-small [2,6,10]", stream, {2, 6, 10}, {2, 6, 10});
204+ RunCase("same-large [4096,4096]", stream, {4096, 4096}, {4096, 4096}, 50);
205+ RunCase("bcast [4,1,8]x[4,5,8]", stream, {4, 1, 8}, {4, 5, 8});
206+ RunCase("bcast scalar [1]x[3,4,5]", stream, {1}, {3, 4, 5});
207+ RunCase("bcast [8,1]x[1,7]", stream, {8, 1}, {1, 7});
208+ 
209+ aclrtDestroyStream(stream);
210+ aclrtResetDevice(deviceId);
211+ aclFinalize();
212+ return 0;
213+}
Aexperimental/math/polar/op_api/aclnn_polar.cpp+189-0
@@ -0,0 +1,189 @@
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+ * \file aclnn_polar.cpp
13+ * \brief Polar L2 API：Contiguous + numpy 广播 (BroadcastTo) + L0 Polar + ViewCopy。
14+ * kernel 假设同 shape elementwise；所有广播/非连续处理在本层完成 ——
15+ * 因此任意 shape / 任意广播组合（≤8 维）都能正确处理，不会因 outNumel 大或
16+ * 中间轴广播触发 kernel 内的 fp32 unravel 失精或 UB 容量越界。
17+ */
18+ 
19+#include "aclnn_polar.h"
20+#include "polar.h"
21+ 
22+#include "aclnn_kernels/contiguous.h"
23+#include "aclnn_kernels/common/op_error_check.h"
24+#include "conversion/broadcast_to/op_api/broadcast_to.h" // l0op::BroadcastTo
25+#include "opdev/op_log.h"
26+#include "opdev/op_dfx.h"
27+#include "opdev/common_types.h"
28+#include "opdev/data_type_utils.h"
29+#include "opdev/make_op_executor.h"
30+#include "opdev/platform.h"
31+#include "opdev/shape_utils.h"
32+#include "op_api/aclnn_check.h"
33+#include "op_api/op_api_def.h"
34+ 
35+using namespace op;
36+ 
37+#ifdef __cplusplus
38+extern "C" {
39+#endif
40+ 
41+#define ACLNN_MAX_SHAPE_RANK 8
42+ 
43+static const std::initializer_list<op::DataType> INPUT_DTYPE_SUPPORT_LIST = {DataType::DT_FLOAT};
44+ 
45+static const std::initializer_list<op::DataType> OUTPUT_DTYPE_SUPPORT_LIST = {DataType::DT_COMPLEX64};
46+ 
47+static bool CheckNotNull(const aclTensor* input, const aclTensor* angle, const aclTensor* out)
48+{
49+ OP_CHECK_NULL(input, return false);
50+ OP_CHECK_NULL(angle, return false);
51+ OP_CHECK_NULL(out, return false);
52+ return true;
53+}
54+ 
55+static bool CheckDtypeValid(const aclTensor* input, const aclTensor* angle, const aclTensor* out)
56+{
57+ if (!CheckType(input->GetDataType(), INPUT_DTYPE_SUPPORT_LIST)) {
58+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "input dtype %s not in support list (expect FLOAT).",
59+ op::ToString(input->GetDataType()).GetString());
60+ return false;
61+ }
62+ if (input->GetDataType() != angle->GetDataType()) {
63+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "input dtype %s != angle dtype %s.",
64+ op::ToString(input->GetDataType()).GetString(), op::ToString(angle->GetDataType()).GetString());
65+ return false;
66+ }
67+ if (!CheckType(out->GetDataType(), OUTPUT_DTYPE_SUPPORT_LIST)) {
68+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "out dtype %s not in support list (expect COMPLEX64).",
69+ op::ToString(out->GetDataType()).GetString());
70+ return false;
71+ }
72+ return true;
73+}
74+ 
75+static bool CheckBroadcastShape(const aclTensor* input, const aclTensor* angle, const aclTensor* out)
76+{
77+ OP_CHECK_MAX_DIM(input, ACLNN_MAX_SHAPE_RANK, return false);
78+ OP_CHECK_MAX_DIM(angle, ACLNN_MAX_SHAPE_RANK, return false);
79+ OP_CHECK_MAX_DIM(out, ACLNN_MAX_SHAPE_RANK, return false);
80+ 
81+ OP_CHECK_BROADCAST(input, angle, return false);
82+ Shape broadcastShape;
83+ BroadcastInferShape(input->GetViewShape(), angle->GetViewShape(), broadcastShape);
84+ 
85+ if (broadcastShape != out->GetViewShape()) {
86+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "broadcast(input, angle) shape %s != out shape %s.",
87+ op::ToString(broadcastShape).GetString(), op::ToString(out->GetViewShape()).GetString());
88+ return false;
89+ }
90+ return true;
91+}
92+ 
93+static aclnnStatus CheckParams(const aclTensor* input, const aclTensor* angle, const aclTensor* out)
94+{
95+ CHECK_RET(CheckNotNull(input, angle, out), ACLNN_ERR_PARAM_NULLPTR);
96+ CHECK_RET(CheckDtypeValid(input, angle, out), ACLNN_ERR_PARAM_INVALID);
97+ CHECK_RET(CheckBroadcastShape(input, angle, out), ACLNN_ERR_PARAM_INVALID);
98+ return ACLNN_SUCCESS;
99+}
100+ 
101+static aclIntArray* GetShapeAsIntArray(const aclTensor* tensor, aclOpExecutor* executor)
102+{
103+ int64_t dimNum = static_cast<int64_t>(tensor->GetViewShape().GetDimNum());
104+ if (dimNum == 0) {
105+ int64_t shape[1] = {1};
106+ return executor->AllocIntArray(shape, 1);
107+ }
108+ std::vector<int64_t> shape(dimNum);
109+ for (int64_t i = 0; i < dimNum; i++) {
110+ shape[i] = tensor->GetViewShape()[i];
111+ }
112+ return executor->AllocIntArray(shape.data(), dimNum);
113+}
114+ 
115+aclnnStatus aclnnPolarGetWorkspaceSize(const aclTensor* input, const aclTensor* angle, aclTensor* out,
116+ uint64_t* workspaceSize, aclOpExecutor** executor)
117+{
118+ L2_DFX_PHASE_1(aclnnPolar, DFX_IN(input, angle), DFX_OUT(out));
119+ 
120+ auto uniqueExecutor = CREATE_EXECUTOR();
121+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
122+ 
123+ auto ret = CheckParams(input, angle, out);
124+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
125+ 
126+ if (input->IsEmpty() || angle->IsEmpty()) {
127+ *workspaceSize = 0;
128+ uniqueExecutor.ReleaseTo(executor);
129+ return ACLNN_SUCCESS;
130+ }
131+ 
132+ // 1. make contiguous（解决 view/transpose 等非连续输入）
133+ auto inputContiguous = l0op::Contiguous(input, uniqueExecutor.get());
134+ CHECK_RET(inputContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
135+ auto angleContiguous = l0op::Contiguous(angle, uniqueExecutor.get());
136+ CHECK_RET(angleContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
137+ 
138+ // 2. broadcast 到 output shape。
139+ // input 永远 BroadcastTo out（kernel 要求 input/out 同 shape）。
140+ // angle：满足 inner-broadcast 资格时**保留原 shape [K]**（kernel bcastMode=1 周期复用，省一次 BroadcastTo
141+ // 的 HBM 写 + 每 tile angle 搬运 + Sin/Cos 重算）；否则 BroadcastTo 兜底。
142+ // 资格判定与 op_host/polar_tiling.cpp 严格一致：inN==outN, anN<inN, inN%anN==0, anN<=2048, anN%8==0。
143+ auto outShape = GetShapeAsIntArray(out, uniqueExecutor.get());
144+ CHECK_RET(outShape != nullptr, ACLNN_ERR_INNER_NULLPTR);
145+ 
146+ if (inputContiguous->GetViewShape() != out->GetViewShape()) {
147+ inputContiguous = l0op::BroadcastTo(inputContiguous, outShape, uniqueExecutor.get());
148+ CHECK_RET(inputContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
149+ }
150+ 
151+ constexpr int64_t POLAR_TILE_LEN = 2048;
152+ auto numelOf = [](const op::Shape& s) -> int64_t {
153+ int64_t n = 1;
154+ for (size_t i = 0; i < s.GetDimNum(); ++i)
155+ n *= s[i];
156+ return n;
157+ };
158+ int64_t outN = numelOf(out->GetViewShape());
159+ int64_t inNn = numelOf(inputContiguous->GetViewShape());
160+ int64_t anNn = numelOf(angleContiguous->GetViewShape());
161+ bool innerBcast = (inNn == outN) && (anNn > 0) && (anNn < inNn) && (inNn % anNn == 0) && (anNn <= POLAR_TILE_LEN) &&
162+ (anNn % 8 == 0);
163+ if (!innerBcast && angleContiguous->GetViewShape() != out->GetViewShape()) {
164+ angleContiguous = l0op::BroadcastTo(angleContiguous, outShape, uniqueExecutor.get());
165+ CHECK_RET(angleContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
166+ }
167+ 
168+ // 3. 调 Polar kernel（同 shape elementwise）
169+ auto opResult = l0op::Polar(inputContiguous, angleContiguous, uniqueExecutor.get());
170+ CHECK_RET(opResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
171+ 
172+ // 4. 把 kernel 输出 copy 到调用者提供的 out
173+ auto viewCopyResult = l0op::ViewCopy(opResult, out, uniqueExecutor.get());
174+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
175+ 
176+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
177+ uniqueExecutor.ReleaseTo(executor);
178+ return ACLNN_SUCCESS;
179+}
180+ 
181+aclnnStatus aclnnPolar(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
182+{
183+ L2_DFX_PHASE_2(aclnnPolar);
184+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
185+}
186+ 
187+#ifdef __cplusplus
188+}
189+#endif
Aexperimental/math/polar/op_api/aclnn_polar.h+61-0
@@ -0,0 +1,61 @@
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+#ifndef OP_API_INC_LEVEL2_ACLNN_POLAR_H_
12+#define OP_API_INC_LEVEL2_ACLNN_POLAR_H_
13+ 
14+#include "aclnn/aclnn_base.h"
15+// 注意：aclnn_util.h 含 ACLNN_API、CHECK_RET 等内部宏，但**只在 .cpp 内引用**；
16+// 这里不在公共头暴露 SDK 内部头，避免下游 build_aclnn.sh (CMakeLists 不含 SDK 内部 include path) 找不到。
17+// ACLNN_API 在 SDK 内是 visibility=default 宏，下游若未定义这里给个安全 fallback。
18+#ifndef ACLNN_API
19+#define ACLNN_API __attribute__((visibility("default")))
20+#endif
21+ 
22+#ifdef __cplusplus
23+extern "C" {
24+#endif
25+ 
26+/**
27+ * @brief aclnnPolar 第一段接口：参数校验 + Contiguous + BroadcastTo + 调 L0 kernel + ViewCopy。
28+ * @domain aclnn_math
29+ *
30+ * 公式： out = input * (cos(angle) + i*sin(angle))
31+ *
32+ * 计算图：
33+ * ```mermaid
34+ * graph LR
35+ * A[(input fp32)] --> B([l0op::Contiguous]) --> C([l0op::BroadcastTo])
36+ * D[(angle fp32)] --> E([l0op::Contiguous]) --> F([l0op::BroadcastTo])
37+ * C --> G([l0op::Polar])
38+ * F --> G
39+ * G --> H([l0op::ViewCopy]) --> I[(out complex64)]
40+ * ```
41+ *
42+ * @param [in] input 模长分量，fp32
43+ * @param [in] angle 幅角（弧度），fp32；与 input 满足 numpy 广播
44+ * @param [out] out 复数张量，complex64；shape 等于 input/angle 广播后的 shape
45+ * @param [out] workspaceSize device 侧需要的 workspace 字节数
46+ * @param [out] executor op 执行器（第二段接口使用）
47+ * @return aclnnStatus
48+ */
49+ACLNN_API aclnnStatus aclnnPolarGetWorkspaceSize(const aclTensor* input, const aclTensor* angle, aclTensor* out,
50+ uint64_t* workspaceSize, aclOpExecutor** executor);
51+ 
52+/**
53+ * @brief aclnnPolar 第二段接口：执行计算
54+ */
55+ACLNN_API aclnnStatus aclnnPolar(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
56+ 
57+#ifdef __cplusplus
58+}
59+#endif
60+ 
61+#endif // OP_API_INC_LEVEL2_ACLNN_POLAR_H_
Aexperimental/math/polar/op_api/polar.cpp+50-0
@@ -0,0 +1,50 @@
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+ * \file polar.cpp
13+ * \brief Polar L0 API：分配 complex64 输出张量并调度 AICore kernel。
14+ */
15+ 
16+#include "polar.h"
17+#include "opdev/op_log.h"
18+#include "opdev/op_dfx.h"
19+#include "opdev/shape_utils.h"
20+#include "opdev/make_op_executor.h"
21+#include "op_api/aclnn_check.h"
22+ 
23+using namespace op;
24+ 
25+namespace l0op {
26+ 
27+OP_TYPE_REGISTER(Polar);
28+ 
29+static const aclTensor* PolarAiCore(const aclTensor* input, const aclTensor* angle, const aclTensor* out,
30+ aclOpExecutor* executor)
31+{
32+ L0_DFX(PolarAiCore, input, angle, out);
33+ 
34+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(Polar, OP_INPUT(input, angle), OP_OUTPUT(out));
35+ OP_CHECK(ret == ACLNN_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "PolarAiCore add to launcher failed."),
36+ return nullptr);
37+ return out;
38+}
39+ 
40+const aclTensor* Polar(const aclTensor* input, const aclTensor* angle, aclOpExecutor* executor)
41+{
42+ // input 经 op_api 层 BroadcastTo 到 out shape；angle 可能保留原 [K]（inner-bcast，kernel 周期复用）。
43+ // out shape = input.shape（= 广播后输出形状），dtype = complex64。
44+ Shape outShape = input->GetViewShape();
45+ const aclTensor* out = executor->AllocTensor(outShape, op::DataType::DT_COMPLEX64);
46+ OP_CHECK(out != nullptr, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "Polar alloc output tensor failed."), return nullptr);
47+ return PolarAiCore(input, angle, out, executor);
48+}
49+ 
50+} // namespace l0op
Aexperimental/math/polar/op_api/polar.h+26-0
@@ -0,0 +1,26 @@
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+#ifndef OP_API_INC_LEVEL0_POLAR_H_
12+#define OP_API_INC_LEVEL0_POLAR_H_
13+ 
14+#include "opdev/op_executor.h"
15+ 
16+namespace l0op {
17+ 
18+/**
19+ * Polar L0：input/angle 已通过 op_api 层 Contiguous + BroadcastTo 对齐到同 shape；
20+ * 此处分配 complex64 输出张量并调度 AICore kernel。
21+ */
22+const aclTensor* Polar(const aclTensor* input, const aclTensor* angle, aclOpExecutor* executor);
23+ 
24+} // namespace l0op
25+ 
26+#endif // OP_API_INC_LEVEL0_POLAR_H_
Aexperimental/math/polar/op_host/polar_def.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+ * \file polar_def.cpp
13+ * \brief Polar 算子原型定义：out = input * (cos(angle) + i*sin(angle))
14+ * input/angle 为 float32（可 numpy 广播），out 恒为 complex64。
15+ * 支持 Atlas A2 训练系列产品（ascend910b）/ Atlas A3 训练系列产品（ascend910_93）。
16+ */
17+#include "register/op_def_registry.h"
18+ 
19+namespace ops {
20+class Polar : public OpDef {
21+public:
22+ explicit Polar(const char* name) : OpDef(name)
23+ {
24+ this->Input("input")
25+ .ParamType(REQUIRED)
26+ .DataType({ge::DT_FLOAT})
27+ .Format({ge::FORMAT_ND})
28+ .UnknownShapeFormat({ge::FORMAT_ND})
29+ .AutoContiguous();
30+ this->Input("angle")
31+ .ParamType(REQUIRED)
32+ .DataType({ge::DT_FLOAT})
33+ .Format({ge::FORMAT_ND})
34+ .UnknownShapeFormat({ge::FORMAT_ND})
35+ .AutoContiguous();
36+ this->Output("out")
37+ .ParamType(REQUIRED)
38+ .DataType({ge::DT_COMPLEX64})
39+ .Format({ge::FORMAT_ND})
40+ .UnknownShapeFormat({ge::FORMAT_ND})
41+ .AutoContiguous();
42+ 
43+ this->AICore().AddConfig("ascend910b"); // Atlas A2 训练/推理系列产品
44+ this->AICore().AddConfig("ascend910_93"); // Atlas A3 训练/推理系列产品
45+ }
46+};
47+OP_ADD(Polar);
48+} // namespace ops
Aexperimental/math/polar/op_host/polar_infershape.cpp+54-0
@@ -0,0 +1,54 @@
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+ * \file polar_infershape.cpp
13+ * \brief Polar InferShape / InferDataType。
14+ * InferShape：numpy 广播规则推导 out shape（右对齐各轴取 max），
15+ * 对齐参考 cann/ops-math math/complex/op_host/op_api/aclnn_polar.cpp 的
16+ * OP_CHECK_BROADCAST_AND_INFER_SHAPE。
17+ * InferDataType：out 恒置 complex64（与 input dtype 无关）。
18+ */
19+#include "register/op_impl_registry.h"
20+#include "op_common/log/log.h"
21+ 
22+using namespace ge;
23+ 
24+namespace ops {
25+ 
26+static ge::graphStatus InferShape4Polar(gert::InferShapeContext* context)
27+{
28+ const gert::Shape* in = context->GetInputShape(0);
29+ OP_CHECK_NULL_WITH_CONTEXT(context, in);
30+ const gert::Shape* an = context->GetInputShape(1);
31+ OP_CHECK_NULL_WITH_CONTEXT(context, an);
32+ gert::Shape* out = context->GetOutputShape(0);
33+ OP_CHECK_NULL_WITH_CONTEXT(context, out);
34+ size_t inRank = in->GetDimNum();
35+ size_t anRank = an->GetDimNum();
36+ size_t outRank = inRank > anRank ? inRank : anRank;
37+ out->SetDimNum(outRank);
38+ for (size_t i = 0; i < outRank; i++) {
39+ int64_t di = (i + inRank >= outRank) ? in->GetDim(i + inRank - outRank) : 1;
40+ int64_t da = (i + anRank >= outRank) ? an->GetDim(i + anRank - outRank) : 1;
41+ out->SetDim(i, di > da ? di : da);
42+ }
43+ return GRAPH_SUCCESS;
44+}
45+ 
46+// 输出恒为 complex64（对齐参考 OUTPUT_DTYPE = {DT_COMPLEX64}）
47+static ge::graphStatus InferDataType4Polar(gert::InferDataTypeContext* context)
48+{
49+ context->SetOutputDataType(0, ge::DT_COMPLEX64);
50+ return ge::GRAPH_SUCCESS;
51+}
52+ 
53+IMPL_OP_INFERSHAPE(Polar).InferShape(InferShape4Polar).InferDataType(InferDataType4Polar);
54+} // namespace ops
Aexperimental/math/polar/op_host/polar_tiling.cpp+153-0
@@ -0,0 +1,153 @@
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+ * \file polar_tiling.cpp
13+ * \brief Polar Tiling —— 同 shape elementwise：仅做 totalLen 计算 + 分核（小 case 少核、强制偶数核）。
14+ * 广播由 op_api 层 BroadcastTo + Contiguous 完成；kernel 收到的 input/angle 已是同 shape。
15+ */
16+#include "register/op_def_registry.h"
17+#include "op_common/log/log.h"
18+#include "op_common/op_host/util/platform_util.h"
19+#include "lib/math/sin_tiling.h" // GetSinMaxMinTmpSize
20+#include "lib/math/cos_tiling.h" // GetCosMaxMinTmpSize
21+#include "../op_kernel/polar_tiling_data.h"
22+#include "../op_kernel/polar_tiling_key.h"
23+ 
24+namespace optiling {
25+ 
26+static constexpr uint32_t POLAR_TILE_LEN = 2048; // 每 tile fp32 元素数
27+static constexpr uint32_t MIN_PER_CORE = 2048u; // 每核最少元素数（小 case 收敛到少核，降 launch 开销）
28+ 
29+static uint64_t Numel(const gert::Shape& s)
30+{
31+ uint64_t n = 1;
32+ uint32_t r = static_cast<uint32_t>(s.GetDimNum());
33+ for (uint32_t i = 0; i < r; i++) {
34+ n *= (uint64_t)s.GetDim(i);
35+ }
36+ return n == 0 ? 1 : n;
37+}
38+ 
39+// angle inner-broadcast 判定（与 op_api/aclnn_polar.cpp 的判定保持一致）：
40+// inN==totalLen, anN<inN, inN%anN==0, anN<=tileLen, anN%8==0（32B 对齐）
41+static bool IsAngleInnerBcast(uint64_t inN, uint64_t anN, uint64_t totalLen)
42+{
43+ return inN == totalLen && anN > 0 && anN < inN && (inN % anN) == 0 && anN <= POLAR_TILE_LEN && (anN % 8u) == 0u;
44+}
45+ 
46+// 核数收敛：小 case 少核（want 上限）+ 强制偶数（>2 时下取偶）
47+static uint32_t ClampWantEven(uint32_t coreNum, uint32_t want)
48+{
49+ if (want < 1u)
50+ want = 1u;
51+ if (coreNum > want)
52+ coreNum = want;
53+ if (coreNum > 2u && (coreNum & 1u))
54+ coreNum -= 1u;
55+ return coreNum;
56+}
57+ 
58+// bcastMode=1 K-block 分核：每核负载是 K 整数倍 → coreStart % K == 0，tile 起点 period 对齐
59+static void SplitCoresBcast(PolarTilingData* tiling, uint64_t totalLen, uint64_t anN, uint32_t& coreNum)
60+{
61+ uint32_t K = (uint32_t)anN;
62+ uint64_t totalBlocks = totalLen / K;
63+ coreNum = ClampWantEven(coreNum, (uint32_t)totalBlocks);
64+ if (coreNum == 0)
65+ coreNum = 1;
66+ uint64_t perBlk = totalBlocks / coreNum;
67+ uint32_t remBlk = (uint32_t)(totalBlocks % coreNum);
68+ tiling->bigCoreNum = remBlk;
69+ tiling->bigCoreLen = (uint32_t)((perBlk + 1) * (uint64_t)K);
70+ tiling->smallCoreLen = (uint32_t)(perBlk * (uint64_t)K);
71+}
72+ 
73+// 同 shape elementwise 分核：核数 ① 平台 AIV ② 小 case 少核 ③ 强制偶数，且不超过 totalLen
74+static void SplitCoresElementwise(PolarTilingData* tiling, uint64_t totalLen, uint32_t& coreNum)
75+{
76+ coreNum = ClampWantEven(coreNum, (uint32_t)(totalLen / MIN_PER_CORE));
77+ if ((uint64_t)coreNum > totalLen)
78+ coreNum = totalLen > 0 ? (uint32_t)totalLen : 1u;
79+ if (coreNum == 0)
80+ coreNum = 1;
81+ uint64_t per = totalLen / coreNum;
82+ uint32_t rem = (uint32_t)(totalLen % coreNum);
83+ tiling->bigCoreNum = rem;
84+ tiling->bigCoreLen = (uint32_t)(per + 1);
85+ tiling->smallCoreLen = (uint32_t)per;
86+}
87+ 
88+// Sin/Cos 显式 sharedTmpBuffer 大小：按 tileLen 元素一次性处理（fp32 noreuse，需 3×N×typeSize 字节）。
89+// ★ 取 maxValue（充足）而非 minValue：minValue 仅"勉强够用"，Sin/Cos 会按小 buffer 分块多趟计算，
90+// 大 same-shape case 性能暴跌（实测 16M same-shape: cosMin 5800μs → cosMax ~650μs，~9x）。
91+// UB 余量充足（24KB << 192KB），用 maxValue 让 Sin/Cos 一次算完整 tile；再 clamp 到 UB 上限防越界。
92+static uint32_t CalcSinCosTmpSize(platform_ascendc::PlatformAscendC& plat)
93+{
94+ ge::Shape tileShape({(int64_t)POLAR_TILE_LEN});
95+ uint32_t cosMax = 0, cosMin = 0;
96+ AscendC::GetCosMaxMinTmpSize(tileShape, 4u, /*isReuseSource=*/false, cosMax, cosMin);
97+ uint64_t ubSize = 0;
98+ plat.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
99+ uint32_t tmpSize = cosMax;
100+ if ((uint64_t)tmpSize > ubSize)
101+ tmpSize = (uint32_t)ubSize;
102+ return tmpSize;
103+}
104+ 
105+static ge::graphStatus PolarTilingFunc(gert::TilingContext* context)
106+{
107+ PolarTilingData* tiling = context->GetTilingData<PolarTilingData>();
108+ OP_CHECK_IF(tiling == nullptr, OP_LOGE(context, "tiling is null"), return ge::GRAPH_FAILED);
109+ 
110+ // input/out 同 shape（aclnn 层把 input BroadcastTo out）；angle inner-bcast 时保留原 shape [K]
111+ auto inShapePtr = context->GetInputShape(0);
112+ OP_CHECK_NULL_WITH_CONTEXT(context, inShapePtr);
113+ auto anShapePtr = context->GetInputShape(1);
114+ OP_CHECK_NULL_WITH_CONTEXT(context, anShapePtr);
115+ uint64_t inN = Numel(inShapePtr->GetStorageShape());
116+ uint64_t anN = Numel(anShapePtr->GetStorageShape());
117+ uint64_t totalLen = inN > anN ? inN : anN;
118+ uint32_t bcastMode = IsAngleInnerBcast(inN, anN, totalLen) ? 1u : 0u;
119+ 
120+ auto plat = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
121+ uint32_t coreNum = plat.GetCoreNumAiv();
122+ if (coreNum == 0)
123+ coreNum = 8;
124+ 
125+ if (bcastMode == 1) {
126+ SplitCoresBcast(tiling, totalLen, anN, coreNum);
127+ } else {
128+ SplitCoresElementwise(tiling, totalLen, coreNum);
129+ }
130+ 
131+ tiling->totalLen = (uint32_t)totalLen;
132+ tiling->tileLen = POLAR_TILE_LEN;
133+ tiling->coreNum = coreNum;
134+ tiling->inN = (uint32_t)inN;
135+ tiling->anN = (uint32_t)anN;
136+ tiling->bcastMode = bcastMode;
137+ tiling->tmpBufferSize = CalcSinCosTmpSize(plat);
138+ 
139+ context->SetBlockDim(coreNum);
140+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
141+ currentWorkspace[0] = 0;
142+ return ge::GRAPH_SUCCESS;
143+}
144+ 
145+struct PolarCompileInfo {};
146+ 
147+static ge::graphStatus TilingParseForPolar([[maybe_unused]] gert::TilingParseContext* context)
148+{
149+ return ge::GRAPH_SUCCESS;
150+}
151+ 
152+IMPL_OP_OPTILING(Polar).Tiling(PolarTilingFunc).TilingParse<PolarCompileInfo>(TilingParseForPolar);
153+} // namespace optiling
Aexperimental/math/polar/op_kernel/polar.cpp+29-0
@@ -0,0 +1,29 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file polar.cpp
13+ * \brief Polar kernel 入口（单 kernel 全路径；全形/广播按 TilingData 运行期判断）
14+ */
15+#include "polar.h"
16+ 
17+using namespace NsPolar;
18+ 
19+template <uint32_t schMode>
20+__global__ __aicore__ void polar(GM_ADDR input, GM_ADDR angle, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling)
21+{
22+ REGISTER_TILING_DEFAULT(PolarTilingData);
23+ GET_TILING_DATA_WITH_STRUCT(PolarTilingData, tiling_data, tiling);
24+ 
25+ AscendC::TPipe pipe;
26+ KernelPolar op;
27+ op.Init(input, angle, out, tiling_data, &pipe);
28+ op.Process();
29+}
Aexperimental/math/polar/op_kernel/polar.h+233-0
@@ -0,0 +1,233 @@
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+ * \file polar.h
13+ * \brief Polar kernel —— 同 shape input/angle 逐元素：out = input*(cos(angle)+i*sin(angle))。
14+ * 广播由 op_api 层 BroadcastTo + Contiguous 完成后再调本 kernel；
15+ * kernel 假设 input.shape == angle.shape == out.shape；
16+ * out 为 complex64（real/imag 交织 fp32，用 Gather + 静态 offset 表组装）。
17+ * 平台：Atlas A2 (ascend910b) / A3 (ascend910_93)；vector 不支持 complex64，
18+ * complex64 当 2*fp32 交织处理。
19+ *
20+ * 资源布局（BUFFER_NUM=2 double buffer）：
21+ * qIn VECIN : 2 buffer × (2×T fp32) = 2 × 16KB —— [abs(0:n) | ang(0:n)] 合并
22+ * qOut VECOUT : 2 buffer × (2×T fp32) = 2 × 16KB —— complex64 交织
23+ * bufPacked VECCALC: 1 buffer × (2×T fp32) = 16KB —— [cos | sin] 中间结果
24+ * bufOff VECCALC: 1 buffer × (2×T u32 ) = 16KB —— Gather 静态 offset 表
25+ * bufSinCosTmp VECCALC: tilingData.tmpBufferSize —— Sin/Cos 显式 sharedTmpBuffer
26+ * （由 host GetSinMaxMinTmpSize 算）
27+ * TPipe Tensor 数：qIn(2) + qOut(2) + bufPacked(1) + bufOff(1) + bufSinCosTmp(1) = 7 (≤8 上限)。
28+ * UB 用量：32 + 32 + 16 + 16 + 24 = 120 KB / 192 KB。
29+ *
30+ * 合并 qAbs+qAng → qIn 的关键：每个 qIn buffer 容纳 2T = 4096 fp32；
31+ * CopyIn 阶段 DataCopyPad input → qIn[0:n]，DataCopyPad angle → qIn[T:T+n]；
32+ * Compute 阶段 Sin/Cos 取 qIn[T:T+n] 作为 source，Mul 取 qIn[0:n] 作为 abs。
33+ */
34+#ifndef POLAR_H_
35+#define POLAR_H_
36+ 
37+#include "kernel_operator.h"
38+#include "kernel_tiling/kernel_tiling.h"
39+#include "lib/math/sin.h"
40+#include "lib/math/cos.h"
41+#include "polar_tiling_data.h"
42+#include "polar_tiling_key.h"
43+ 
44+namespace NsPolar {
45+ 
46+constexpr int32_t BUFFER_NUM = 2;
47+ 
48+class KernelPolar {
49+public:
50+ __aicore__ inline KernelPolar() {}
51+ 
52+ __aicore__ inline void Init(GM_ADDR input, GM_ADDR angle, GM_ADDR out, const PolarTilingData& t,
53+ AscendC::TPipe* pipe)
54+ {
55+ tileLen_ = t.tileLen;
56+ bcastMode_ = t.bcastMode;
57+ K_ = t.anN;
58+ 
59+ uint32_t bid = AscendC::GetBlockIdx();
60+ if (bid < t.bigCoreNum) {
61+ coreLen_ = t.bigCoreLen;
62+ coreStart_ = bid * t.bigCoreLen;
63+ } else {
64+ coreLen_ = t.smallCoreLen;
65+ coreStart_ = t.bigCoreNum * t.bigCoreLen + (bid - t.bigCoreNum) * t.smallCoreLen;
66+ }
67+ 
68+ // input/out 永远是 outN 长度（aclnn 层已把 input BroadcastTo out）
69+ inGm_.SetGlobalBuffer((__gm__ float*)input + coreStart_, coreLen_);
70+ outGmF_.SetGlobalBuffer((__gm__ float*)out + 2u * coreStart_, 2u * coreLen_);
71+ // angle：bcastMode=1 取整段 [0:K]（不偏移，周期复用）；same-shape 按 coreStart 偏移
72+ if (bcastMode_ == 1u) {
73+ anGm_.SetGlobalBuffer((__gm__ float*)angle, K_);
74+ } else {
75+ anGm_.SetGlobalBuffer((__gm__ float*)angle + coreStart_, coreLen_);
76+ }
77+ 
78+ // qIn: 每 buffer 容 2T fp32（[abs|ang]），BUFFER_NUM=2 double buffer
79+ pipe->InitBuffer(qIn_, BUFFER_NUM, 2u * tileLen_ * sizeof(float));
80+ pipe->InitBuffer(qOut_, BUFFER_NUM, 2u * tileLen_ * sizeof(float));
81+ pipe->InitBuffer(bufPacked_, 2u * tileLen_ * sizeof(float)); // [cos | sin]
82+ pipe->InitBuffer(bufOff_, 2u * tileLen_ * sizeof(uint32_t)); // 交织 offset 表
83+ pipe->InitBuffer(bufSinCosTmp_, t.tmpBufferSize); // Sin/Cos sharedTmpBuffer / bcast 常驻 cos/sin tile
84+ 
85+ BuildOffsetTable();
86+ }
87+ 
88+ __aicore__ inline void Process()
89+ {
90+ if (bcastMode_ == 1u) {
91+ // angle inner-broadcast：周期 K，每 tile 整数个 period → cos/sin tile pattern 恒定，只 precompute 一次
92+ PrecomputeBcastCosSin();
93+ uint32_t done = 0;
94+ while (done < coreLen_) {
95+ uint32_t rem = coreLen_ - done;
96+ uint32_t n = rem < tileLen_ ? rem : tileLen_;
97+ if (n >= K_)
98+ n = (n / K_) * K_; // n 取 K 整数倍，保证 period 对齐
99+ ProcessTileBcast(done, n);
100+ done += n;
101+ }
102+ } else {
103+ uint32_t done = 0;
104+ while (done < coreLen_) {
105+ uint32_t n = (coreLen_ - done) < tileLen_ ? (coreLen_ - done) : tileLen_;
106+ ProcessTile(done, n);
107+ done += n;
108+ }
109+ }
110+ }
111+ 
112+private:
113+ // 交织 offset 表：off[j] = 4·(j>>1) + (j&1)·4T （元素 j 选自 cos 部分还是 sin 部分的字节地址）
114+ // = 4T·j + (4 - 8T)·(j>>1)，纯 int32 矢量构造，O(1) 一次 Init 后整 kernel 周期复用
115+ __aicore__ inline void BuildOffsetTable()
116+ {
117+ const int32_t M = (int32_t)(2u * tileLen_);
118+ const int32_t T = (int32_t)tileLen_;
119+ AscendC::LocalTensor<int32_t> A = bufOff_.Get<uint32_t>().ReinterpretCast<int32_t>();
120+ AscendC::LocalTensor<int32_t> B = bufPacked_.Get<float>().ReinterpretCast<int32_t>();
121+ AscendC::CreateVecIndex(A, (int32_t)0, M);
122+ AscendC::ShiftRight(B, A, (int32_t)1, M);
123+ AscendC::Muls(A, A, (int32_t)(4 * T), M);
124+ AscendC::Muls(B, B, (int32_t)(4 - 8 * T), M);
125+ AscendC::Add(A, A, B, M);
126+ }
127+ 
128+ // 交织 + 写出：Gather 静态 offset 表把 [real|imag] 重排为 [r0,i0,r1,i1,...]，complex64 视作 2*fp32、
129+ // DataCopyPad 写 2n 个 fp32。same-shape 与 inner-bcast 两路共用（packed 均为 real[0:n]/imag[T:T+n]）。
130+ __aicore__ inline void InterleaveAndStore(AscendC::LocalTensor<float>& packed, uint32_t off, uint32_t n)
131+ {
132+ AscendC::LocalTensor<uint32_t> off32 = bufOff_.Get<uint32_t>();
133+ AscendC::LocalTensor<float> outL = qOut_.AllocTensor<float>();
134+ AscendC::Gather(outL, packed, off32, (uint32_t)0, 2u * n);
135+ qOut_.EnQue(outL);
136+ outL = qOut_.DeQue<float>();
137+ AscendC::DataCopyExtParams cpo{1, 2u * n * (uint32_t)sizeof(float), 0, 0, 0};
138+ AscendC::DataCopyPad(outGmF_[2u * off], outL, cpo);
139+ qOut_.FreeTensor(outL);
140+ }
141+ 
142+ __aicore__ inline void ProcessTile(uint32_t off, uint32_t n)
143+ {
144+ // CopyIn —— input → qIn[0:n], angle → qIn[T:T+n]
145+ AscendC::LocalTensor<float> inL = qIn_.AllocTensor<float>();
146+ AscendC::DataCopyExtParams cpi{1, n * (uint32_t)sizeof(float), 0, 0, 0};
147+ AscendC::DataCopyPadExtParams<float> pad{false, 0, 0, 0};
148+ AscendC::DataCopyPad(inL, inGm_[off], cpi, pad);
149+ AscendC::DataCopyPad(inL[tileLen_], anGm_[off], cpi, pad);
150+ qIn_.EnQue(inL);
151+ 
152+ inL = qIn_.DeQue<float>();
153+ AscendC::LocalTensor<float> absSeg = inL; // abs at offset 0
154+ AscendC::LocalTensor<float> angSeg = inL[tileLen_]; // ang at offset T
155+ 
156+ // Compute —— Cos/Sin 用显式 sharedTmpBuffer（4-arg 重载），避免 PopStackBuffer 占 TPipe stack
157+ AscendC::LocalTensor<float> packed = bufPacked_.Get<float>();
158+ AscendC::LocalTensor<uint8_t> sharedTmp = bufSinCosTmp_.Get<uint8_t>();
159+ AscendC::Cos<float, false>(packed, angSeg, sharedTmp, n); // packed[0:n] = cos(ang)
160+ AscendC::Sin<float, false>(packed[tileLen_], angSeg, sharedTmp, n); // packed[T:T+n] = sin(ang)
161+ // 原地 *abs 得 real/imag
162+ AscendC::Mul(packed, packed, absSeg, n); // real = abs * cos
163+ AscendC::Mul(packed[tileLen_], packed[tileLen_], absSeg, n); // imag = abs * sin
164+ qIn_.FreeTensor(inL);
165+ 
166+ InterleaveAndStore(packed, off, n);
167+ }
168+ 
169+ // ===== bcastMode=1 inner-broadcast 路径 =====
170+ // angle 周期 K，每 tile = 整数个 period → cos/sin tile pattern 对所有 tile 相同。
171+ // Init 阶段 pre Cos/Sin angle[0:K] 并 broadcast 填满 bufSinCosTmp_ 的 cosTile[0:T]/sinTile[T:2T] 常驻；
172+ // main loop 只读 input → Mul×2 → Gather 交织，省掉每 tile 的 angle 搬运与 Sin/Cos 重算。
173+ __aicore__ inline void PrecomputeBcastCosSin()
174+ {
175+ // 1) angle[0:K] → qIn buffer
176+ AscendC::LocalTensor<float> aL = qIn_.AllocTensor<float>();
177+ AscendC::DataCopyExtParams cpi{1, K_ * (uint32_t)sizeof(float), 0, 0, 0};
178+ AscendC::DataCopyPadExtParams<float> pad{false, 0, 0, 0};
179+ AscendC::DataCopyPad(aL, anGm_, cpi, pad);
180+ qIn_.EnQue(aL);
181+ aL = qIn_.DeQue<float>();
182+ 
183+ // 2) Cos/Sin 写 bufPacked_[0:K]/[K:2K]，scratch 用 bufPacked_ 尾部 [2K:2T]（N=K 所需 tmp 远小于该区）
184+ AscendC::LocalTensor<float> packed = bufPacked_.Get<float>();
185+ AscendC::LocalTensor<uint8_t> scratch = bufPacked_.Get<uint8_t>()[2u * K_ * sizeof(float)];
186+ AscendC::Cos<float, false>(packed[0], aL, scratch, K_);
187+ AscendC::Sin<float, false>(packed[K_], aL, scratch, K_);
188+ qIn_.FreeTensor(aL);
189+ 
190+ // 3) broadcast cosT[0:K]→cosTile[0:T]、sinT[0:K]→sinTile[0:T]，常驻 bufSinCosTmp_
191+ // （一次性 m=T/K 次 K-段复制，不在 hot loop）
192+ AscendC::LocalTensor<float> tile = bufSinCosTmp_.Get<uint8_t>().ReinterpretCast<float>();
193+ const uint32_t mFull = tileLen_ / K_;
194+ for (uint32_t i = 0; i < mFull; ++i) {
195+ AscendC::DataCopy(tile[i * K_], packed[0], K_); // cosTile period i
196+ AscendC::DataCopy(tile[tileLen_ + i * K_], packed[K_], K_); // sinTile period i
197+ }
198+ }
199+ 
200+ // bcastMode=1 单 tile：input only HBM read；cos/sin tile 从 bufSinCosTmp_ 常驻区取（前 n 个即正确 period）。
201+ // off 必 K 对齐、n 是 K 整数倍（host K-block 分核 + main loop 保证）。
202+ __aicore__ inline void ProcessTileBcast(uint32_t off, uint32_t n)
203+ {
204+ // CopyIn input[off:off+n] only
205+ AscendC::LocalTensor<float> inL = qIn_.AllocTensor<float>();
206+ AscendC::DataCopyExtParams cpi{1, n * (uint32_t)sizeof(float), 0, 0, 0};
207+ AscendC::DataCopyPadExtParams<float> pad{false, 0, 0, 0};
208+ AscendC::DataCopyPad(inL, inGm_[off], cpi, pad);
209+ qIn_.EnQue(inL);
210+ inL = qIn_.DeQue<float>();
211+ 
212+ // Mul：real = cosTile * input → packed[0:n]；imag = sinTile * input → packed[T:T+n]
213+ AscendC::LocalTensor<float> packed = bufPacked_.Get<float>();
214+ AscendC::LocalTensor<float> tile = bufSinCosTmp_.Get<uint8_t>().ReinterpretCast<float>();
215+ AscendC::LocalTensor<float> cosTile = tile; // [0:T]
216+ AscendC::LocalTensor<float> sinTile = tile[tileLen_]; // [T:2T]
217+ AscendC::Mul(packed[0], cosTile, inL, n);
218+ AscendC::Mul(packed[tileLen_], sinTile, inL, n);
219+ qIn_.FreeTensor(inL);
220+ 
221+ InterleaveAndStore(packed, off, n);
222+ }
223+ 
224+ AscendC::GlobalTensor<float> inGm_, anGm_, outGmF_;
225+ AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> qIn_; // 合并 abs + ang，buffer 大小 2T
226+ AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> qOut_;
227+ AscendC::TBuf<AscendC::QuePosition::VECCALC> bufPacked_, bufOff_, bufSinCosTmp_;
228+ uint32_t tileLen_{0}, coreLen_{0}, coreStart_{0};
229+ uint32_t bcastMode_{0}, K_{0};
230+};
231+ 
232+} // namespace NsPolar
233+#endif // POLAR_H_
Aexperimental/math/polar/op_kernel/polar_tiling_data.h+34-0
@@ -0,0 +1,34 @@
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+ * \file polar_tiling_data.h
13+ * \brief Polar tiling data struct（host 填写 / kernel 读取共用同一 struct）。
14+ * 两条路径：bcastMode=0 同 shape elementwise；bcastMode=1 angle inner-broadcast
15+ * （angle 周期 K << input，trailing-dim 重复）—— 详见 op_kernel/polar.h。
16+ * 非 inner-bcast 的广播（中间轴/input 侧）仍由 aclnn L2 层 BroadcastTo + Contiguous 完成。
17+ */
18+#ifndef POLAR_TILING_DATA_H_
19+#define POLAR_TILING_DATA_H_
20+ 
21+struct PolarTilingData {
22+ uint32_t totalLen; // 输出元素总数
23+ uint32_t tileLen; // 每 tile 元素数
24+ uint32_t bigCoreNum; // 大核数量（前 bigCoreNum 个核各多处理 1 元素）
25+ uint32_t bigCoreLen; // 大核分块长度
26+ uint32_t smallCoreLen; // 小核分块长度
27+ uint32_t coreNum; // 实际启用核数（偶数）
28+ uint32_t tmpBufferSize; // Sin/Cos 显式 sharedTmpBuffer 字节数（由 Get{Sin,Cos}MaxMinTmpSize 算）
29+ uint32_t inN; // input.numel（bcastMode=1 时 = totalLen）
30+ uint32_t anN; // angle.numel（bcastMode=1 时 = K，周期长度）
31+ uint32_t bcastMode; // 0: same-shape；1: angle inner-broadcast
32+};
33+ 
34+#endif // POLAR_TILING_DATA_H_
Aexperimental/math/polar/op_kernel/polar_tiling_key.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+ * \file polar_tiling_key.h
13+ * \brief polar tiling key declare（单 kernel 全路径，全形/广播按 TilingData 运行期判断，
14+ * 不设多 tilingkey；此处保留 ops-math 模板参数样板以接入构建框架）
15+ */
16+#ifndef POLAR_TILING_KEY_H_
17+#define POLAR_TILING_KEY_H_
18+ 
19+#include "ascendc/host_api/tiling/template_argument.h"
20+ 
21+#define ELEMENTWISE_TPL_SCH_MODE_0 0
22+#define ELEMENTWISE_TPL_SCH_MODE_1 1
23+ 
24+ASCENDC_TPL_ARGS_DECL(Polar, ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0,
25+ ELEMENTWISE_TPL_SCH_MODE_1));
26+ 
27+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0,
28+ ELEMENTWISE_TPL_SCH_MODE_1)), );
29+ 
30+#endif // POLAR_TILING_KEY_H_
Aexperimental/math/polar/tests/Polar.json+40-0
@@ -0,0 +1,40 @@
1+[
F
Ffulltower6月30日

tests下面是补充UT测试代码，参考其他样例

likedislike
2+ {
3+ "op": "Polar",
4+ "language": "cpp",
5+ "input_desc": [
6+ {
7+ "name": "input",
8+ "param_type": "required",
9+ "format": [
10+ "ND"
11+ ],
12+ "type": [
13+ "float"
14+ ]
15+ },
16+ {
17+ "name": "angle",
18+ "param_type": "required",
19+ "format": [
20+ "ND"
21+ ],
22+ "type": [
23+ "float"
24+ ]
25+ }
26+ ],
27+ "output_desc": [
28+ {
29+ "name": "out",
30+ "param_type": "required",
31+ "format": [
32+ "ND"
33+ ],
34+ "type": [
35+ "complex64"
36+ ]
37+ }
38+ ]
39+ }
40+]
Aexperimental/math/polar/tests/aot/README.md+42-0
@@ -0,0 +1,42 @@
1+# AscendOpTest 官方精度/性能验收（Polar）
2+ 
3+任务要求"精度满足 AscendOpTest 工具默认阈值"，须用**官方工具**跑（非我们 pybind 复刻）。本目录为离线准备好的资产。
4+ 
5+## 资产
6+ 
7+| 文件 | 作用 |
8+|---|---|
9+| `../../../S8/Polar.json` | 算子描述文件（IR/原型，复用，`-i` 传入） |
10+| `polar_cases.json` | 测试用例（6 例：同 shape 小/16M + 广播低→高/标量/双向 + 高维非对齐），`-c` 传入 |
11+| `polar_golden.py` | `expect_func` CPU 基准（numpy 广播 polar），被 case json 引用 |
12+ 
13+**关键点**：complex64 在 AscendOpTest `accuracy_config` **无内置默认** → 每个 `output_desc` 已显式写 `"err_threshold":[0.0001,0.0001]`（fp32 分量默认；[绝对偏差,错误率]），否则 run 时 KeyError。
14+ 
15+## 在 NPU notebook 上运行
16+ 
17+```bash
18+# 0) 前置
19+pip install ml_dtypes
20+# 1) 部署自定义 Polar 算子（构建+安装 .run，同 S8/Polar 流程）
21+cd ~/work/S8/Polar && bash build.sh && ./build_out/*.run --quiet
22+# 2) 关键：使部署算子生效（README 红字强调）
23+source $ASCEND_HOME_PATH/opp/vendors/customize/bin/set_env.bash
24+export LD_LIBRARY_PATH=$ASCEND_HOME_PATH/opp/vendors/customize/op_api/lib:$LD_LIBRARY_PATH
25+# 3) 取官方工具
26+git clone https://gitcode.com/HIT1920/AscendOpTest.git && cd AscendOpTest
27+# 4) 跑精度（--build 首次/改用例后必加；--op-type custom 自定义算子）
28+python run_test.py \
29+ -i /home/ma-user/work/S8/Polar.json \
30+ -c /home/ma-user/work/case_910b/Polar/aot/polar_cases.json \
31+ --op-type custom --build
32+# 5) 跑性能（msprof application 形式）
33+python run_test.py -i .../Polar.json -c .../polar_cases.json --op-type custom --msprof
34+```
35+ 
36+结果在 `result.csv`（每用例 pass/fail）。`expect_func` 路径已按 `/home/ma-user/work/case_910b/Polar/aot/polar_golden.py:polar` 写死——若上传路径不同需同步改 `polar_cases.json`。
37+ 
38+## 注意
39+ 
40+- input/angle/out 的 name 与顺序须与 `Polar.json`、OpDef 一致（input, angle, out）——已对齐
41+- `Test_same_16M` 用例专门用于让官方工具**独立复核**之前 standalone 测出的 16M 非确定性问题（见项目记忆 ⚠️ 未决风险）
42+- 角度取 `[-3.14,3.14]`（单周期，干净查正确性）；abs 取 `[-10,10]`（含负值，验证负 abs 翻号合法）
Aexperimental/math/polar/tests/aot/polar_cases.json+86-0
@@ -0,0 +1,86 @@
1+[
2+ {
3+ "case_name": "Test_same_small",
4+ "op_name": "Polar",
5+ "case_path": "",
6+ "expect_func": "/home/ma-user/work/case_910b/Polar/aot/polar_golden.py:polar",
7+ "input_desc": [
8+ { "name": "input", "format": "ND", "data_type": "float", "param_type": "required", "shape": [2, 6, 10], "data_path": "", "value_range": [-10, 10] },
9+ { "name": "angle", "format": "ND", "data_type": "float", "param_type": "required", "shape": [2, 6, 10], "data_path": "", "value_range": [-3.14, 3.14] }
10+ ],
11+ "output_desc": [
12+ { "name": "out", "format": "ND", "data_type": "complex64", "param_type": "required", "shape": [2, 6, 10], "data_path": "", "golden_path": "", "err_threshold": [0.0001, 0.0001] }
13+ ],
14+ "attr_desc": []
15+ },
16+ {
17+ "case_name": "Test_same_16M",
18+ "op_name": "Polar",
19+ "case_path": "",
20+ "expect_func": "/home/ma-user/work/case_910b/Polar/aot/polar_golden.py:polar",
21+ "input_desc": [
22+ { "name": "input", "format": "ND", "data_type": "float", "param_type": "required", "shape": [4096, 4096], "data_path": "", "value_range": [-10, 10] },
23+ { "name": "angle", "format": "ND", "data_type": "float", "param_type": "required", "shape": [4096, 4096], "data_path": "", "value_range": [-3.14, 3.14] }
24+ ],
25+ "output_desc": [
26+ { "name": "out", "format": "ND", "data_type": "complex64", "param_type": "required", "shape": [4096, 4096], "data_path": "", "golden_path": "", "err_threshold": [0.0001, 0.0001] }
27+ ],
28+ "attr_desc": []
29+ },
30+ {
31+ "case_name": "Test_bcast_lowhigh",
32+ "op_name": "Polar",
33+ "case_path": "",
34+ "expect_func": "/home/ma-user/work/case_910b/Polar/aot/polar_golden.py:polar",
35+ "input_desc": [
36+ { "name": "input", "format": "ND", "data_type": "float", "param_type": "required", "shape": [4, 1, 8], "data_path": "", "value_range": [-10, 10] },
37+ { "name": "angle", "format": "ND", "data_type": "float", "param_type": "required", "shape": [4, 5, 8], "data_path": "", "value_range": [-3.14, 3.14] }
38+ ],
39+ "output_desc": [
40+ { "name": "out", "format": "ND", "data_type": "complex64", "param_type": "required", "shape": [4, 5, 8], "data_path": "", "golden_path": "", "err_threshold": [0.0001, 0.0001] }
41+ ],
42+ "attr_desc": []
43+ },
44+ {
45+ "case_name": "Test_bcast_scalar",
46+ "op_name": "Polar",
47+ "case_path": "",
48+ "expect_func": "/home/ma-user/work/case_910b/Polar/aot/polar_golden.py:polar",
49+ "input_desc": [
50+ { "name": "input", "format": "ND", "data_type": "float", "param_type": "required", "shape": [1], "data_path": "", "value_range": [-10, 10] },
51+ { "name": "angle", "format": "ND", "data_type": "float", "param_type": "required", "shape": [3, 4, 5], "data_path": "", "value_range": [-3.14, 3.14] }
52+ ],
53+ "output_desc": [
54+ { "name": "out", "format": "ND", "data_type": "complex64", "param_type": "required", "shape": [3, 4, 5], "data_path": "", "golden_path": "", "err_threshold": [0.0001, 0.0001] }
55+ ],
56+ "attr_desc": []
57+ },
58+ {
59+ "case_name": "Test_bcast_2way",
60+ "op_name": "Polar",
61+ "case_path": "",
62+ "expect_func": "/home/ma-user/work/case_910b/Polar/aot/polar_golden.py:polar",
63+ "input_desc": [
64+ { "name": "input", "format": "ND", "data_type": "float", "param_type": "required", "shape": [8, 1], "data_path": "", "value_range": [-10, 10] },
65+ { "name": "angle", "format": "ND", "data_type": "float", "param_type": "required", "shape": [1, 7], "data_path": "", "value_range": [-3.14, 3.14] }
66+ ],
67+ "output_desc": [
68+ { "name": "out", "format": "ND", "data_type": "complex64", "param_type": "required", "shape": [8, 7], "data_path": "", "golden_path": "", "err_threshold": [0.0001, 0.0001] }
69+ ],
70+ "attr_desc": []
71+ },
72+ {
73+ "case_name": "Test_highdim_unalign",
74+ "op_name": "Polar",
75+ "case_path": "",
76+ "expect_func": "/home/ma-user/work/case_910b/Polar/aot/polar_golden.py:polar",
77+ "input_desc": [
78+ { "name": "input", "format": "ND", "data_type": "float", "param_type": "required", "shape": [3, 5, 17, 269], "data_path": "", "value_range": [-10, 10] },
79+ { "name": "angle", "format": "ND", "data_type": "float", "param_type": "required", "shape": [3, 5, 17, 269], "data_path": "", "value_range": [-3.14, 3.14] }
80+ ],
81+ "output_desc": [
82+ { "name": "out", "format": "ND", "data_type": "complex64", "param_type": "required", "shape": [3, 5, 17, 269], "data_path": "", "golden_path": "", "err_threshold": [0.0001, 0.0001] }
83+ ],
84+ "attr_desc": []
85+ }
86+]
Aexperimental/math/polar/tests/aot/polar_golden.py+30-0
@@ -0,0 +1,30 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# ----------------------------------------------------------------------------
4+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
5+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6+# CANN Open Software License Agreement Version 2.0 (the "License").
7+# Please refer to the License for details. You may not use this file except in compliance with the License.
8+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10+# See LICENSE in the root of the software repository for the full text of the License.
11+# ----------------------------------------------------------------------------
12+"""
13+AscendOpTest expect_func（CPU 基准）—— Polar 算子
14+被 polar_cases.json 的 "expect_func":"<path>/polar_golden.py:polar" 引用。
15+ 
16+约束（AscendOpTest README）：
17+- 参数名/顺序须与算子描述文件(Polar.json)、case 文件 input_desc 一致：input, angle
18+- 返回 list，元素为 numpy.ndarray；输出 dtype 须为 complex64
19+- polar(abs, angle) = abs * (cos(angle) + i*sin(angle))，input 与 angle 按 numpy 广播
20+"""
21+import numpy as np
22+ 
23+ 
24+def polar(input, angle):
25+ a = np.asarray(input).astype(np.float32)
26+ th = np.asarray(angle).astype(np.float32)
27+ # numpy 广播（input.dim 与 angle.dim 可不一致，与算子 InferShape 对齐）
28+ a, th = np.broadcast_arrays(a, th)
29+ out = (a * (np.cos(th) + 1j * np.sin(th))).astype(np.complex64)
30+ return [np.ascontiguousarray(out)]
Aexperimental/math/polar/tests/pybind/README.md+58-0
@@ -0,0 +1,58 @@
1+# Polar 测试框架（S8 同款）
2+ 
3+`out = input · (cos(angle) + i·sin(angle))`，input(abs)/angle 为 fp32 可广播，out 为 complex64。
4+对齐基准：`cann/ops-math` 的 `math/complex/op_host/op_api/aclnn_polar.cpp`（l0 拼接版）。
5+ 
6+## 文件构成（与 S8 case_910b/<Op>/ 完全一致）
7+ 
8+| 文件 | 来源 | 说明 |
9+|---|---|---|
10+| `common/pytorch_npu_helper.hpp` | **逐字复用 S8** | `EXEC_NPU_CMD` 通用胶水（算子无关，动态 dlopen `aclnnPolar`） |
11+| `setup.py` / `get_time.py` | **逐字复用 S8** | pybind wheel 构建 / msprof `op_summary*.csv` 取 [20:40] 均值 |
12+| `extension/custom_op.cpp` | Polar 适配 | `EXEC_NPU_CMD(aclnnPolar, input, angle, result)`，result=complex64、shape=broadcast(input,angle)，50 轮供 msprof 取平均 |
13+| `test_op.py` | Polar 适配 | 8 个 case + 复数 verify（`view_as_real` 拆 real/imag 当 fp32，rtol/atol=1e-4） |
14+| `run.sh` | Polar 适配 | 重建 wheel → msprof 计时 → get_time → 基线校验；LD 指向 `custom_math` vendor |
15+ 
16+## 前置（NPU 环境，改 kernel 后重做）
17+ 
18+```bash
19+# 在 ops-math 仓编译并安装自定义 Polar 算子
20+bash build.sh --pkg --soc=ascend910b --ops=polar -j16 # A3: --soc=ascend910_93
21+./build_out/cann-ops-math-*linux*.run
22+```
23+装到 `${ASCEND_HOME_PATH}/opp/vendors/custom_math/`。
24+ 
25+## 用法
26+ 
27+```bash
28+cd case_910b/Polar
29+bash run.sh 2 # 跑 case2（广播主战场）
30+```
31+输出 `xxx verify result pass!` + `time_base = ... time_use = ...`。
32+ 
33+## case 覆盖
34+ 
35+| case | input shape | angle shape | 考点 |
36+|---|---|---|---|
37+| 1 | [2,6,10] | [2,6,10] | 同 shape 基础 |
38+| 2 | [4,1,8] | [4,5,8] | **广播：低维→高维（新增功能主战场）** |
39+| 3 | [1] | [3,4,5] | 广播：标量 input |
40+| 4 | [8,1] | [1,7] | 广播：双向 → [8,7] |
41+| 5 | [4096,4096] | [4096,4096] | 大 shape（性能） |
42+| 6 | [3,5,17,269] | 同 | 高维 + inner 非 32B 对齐（1076B/行） |
43+| 7 | [1] | [1] | 1 元素边界 |
44+| 8 | [64,1024] | [64,1024] | 大角度（Sin/Cos 范围归约压力） |
45+ 
46+新增 case 直接往 `test_op.py` 的 `case_data` 加键即可（与 S8 习惯一致）。
47+ 
48+## 与 S8 的差异（仅构建层）
49+ 
50+- S8：每算子自带 `S8/<Op>/build.sh`，`auto_submit.sh` 直接 build。
51+- Polar：构建走 **ops-math 仓 `build.sh --pkg --ops=polar`**（仓级），与本测试框架解耦。
52+- `auto_submit.sh -o Polar` 仍可复用其 upload/run 阶段；build 阶段需替换为 ops-math 命令（待算子工程落地后再接线）。
53+ 
54+## 注意
55+ 
56+- 精度阈值现为 fp32 rtol/atol=1e-4；任务要求"AscendOpTest 默认阈值"，待 AscendOpTest 工具实际阈值确认后于 `test_op.py:verify_result` 回填。
57+- 性能基线 `run.sh:time_base` 现为哨兵值（仅正确性门禁）；硬件实测出 l0 参考实现耗时后回填真实基线（任务要求 ≥ 基线 95%）。
58+- 本框架只能在 NPU 环境运行（依赖 `torch_npu`），本地无法自检——与 S8 相同。
Aexperimental/math/polar/tests/pybind/common/pytorch_npu_helper.hpp+575-0
@@ -0,0 +1,575 @@
1+/******************************************************************************
2+ * Copyright (c) 2022 Huawei Technologies Co., Ltd
3+ * All rights reserved.
4+ *
5+ * Licensed under the BSD 3-Clause License (the "License");
6+ * you may not use this file except in compliance with the License.
7+ * You may obtain a copy of the License at
8+ *
9+ * https://opensource.org/licenses/BSD-3-Clause
10+ *
11+ * Unless required by applicable law or agreed to in writing, software
12+ * distributed under the License is distributed on an "AS IS" BASIS,
13+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
14+ * See the License for the specific language governing permissions and
15+ * limitations under the License.
16+ ******************************************************************************/
17+ 
18+#ifndef PYTORCH_NPU_HELPER_HPP_
19+#define PYTORCH_NPU_HELPER_HPP_
20+ 
21+#include <ATen/Tensor.h>
22+#include <acl/acl_base.h>
23+#include <acl/acl_rt.h>
24+#include <c10/util/Exception.h>
25+#include <torch/extension.h>
26+#include <dlfcn.h>
27+#include <torch_npu/csrc/framework/utils/CalcuOpUtil.h>
28+#include <torch_npu/csrc/framework/utils/OpAdapter.h>
29+ 
30+#include <functional>
31+#include <type_traits>
32+#include <vector>
33+ 
34+#include "torch_npu/csrc/aten/NPUNativeFunctions.h"
35+#include "torch_npu/csrc/core/npu/NPUStream.h"
36+#include "torch_npu/csrc/framework/OpCommand.h"
37+#include "torch_npu/csrc/framework/interface/EnvVariables.h"
38+#include "torch_npu/csrc/framework/utils/CalcuOpUtil.h"
39+#include "torch_npu/csrc/framework/utils/OpPreparation.h"
40+ 
41+#define NPU_NAME_SPACE at_npu::native
42+ 
43+#define __FILENAME__ (strrchr("/" __FILE__, '/') + 1)
44+ 
45+typedef struct aclOpExecutor aclOpExecutor;
46+typedef struct aclTensor aclTensor;
47+typedef struct aclScalar aclScalar;
48+typedef struct aclIntArray aclIntArray;
49+typedef struct aclFloatArray aclFloatArray;
50+typedef struct aclBoolArray aclBoolArray;
51+typedef struct aclTensorList aclTensorList;
52+ 
53+typedef aclTensor* (*_aclCreateTensor)(const int64_t* view_dims, uint64_t view_dims_num, aclDataType data_type,
54+ const int64_t* stride, int64_t offset, aclFormat format,
55+ const int64_t* storage_dims, uint64_t storage_dims_num, void* tensor_data);
56+typedef aclScalar* (*_aclCreateScalar)(void* value, aclDataType data_type);
57+typedef aclIntArray* (*_aclCreateIntArray)(const int64_t* value, uint64_t size);
58+typedef aclFloatArray* (*_aclCreateFloatArray)(const float* value, uint64_t size);
59+typedef aclBoolArray* (*_aclCreateBoolArray)(const bool* value, uint64_t size);
60+typedef aclTensorList* (*_aclCreateTensorList)(const aclTensor* const* value, uint64_t size);
61+ 
62+typedef int (*_aclDestroyTensor)(const aclTensor* tensor);
63+typedef int (*_aclDestroyScalar)(const aclScalar* scalar);
64+typedef int (*_aclDestroyIntArray)(const aclIntArray* array);
65+typedef int (*_aclDestroyFloatArray)(const aclFloatArray* array);
66+typedef int (*_aclDestroyBoolArray)(const aclBoolArray* array);
67+typedef int (*_aclDestroyTensorList)(const aclTensorList* array);
68+ 
69+constexpr int kHashBufSize = 8192;
70+constexpr int kHashBufMaxSize = kHashBufSize + 1024;
71+extern thread_local char g_hashBuf[kHashBufSize];
72+extern thread_local int g_hashOffset;
73+ 
74+#define AT_ALL_SCALAR_TYPE_AND_ACL_DATATYPE_PAIR(_) \
75+ _(at::ScalarType::Byte, ACL_UINT8) \
76+ _(at::ScalarType::Char, ACL_INT8) \
77+ _(at::ScalarType::Short, ACL_INT16) \
78+ _(at::ScalarType::Int, ACL_INT32) \
79+ _(at::ScalarType::Long, ACL_INT64) \
80+ _(at::ScalarType::Half, ACL_FLOAT16) \
81+ _(at::ScalarType::Float, ACL_FLOAT) \
82+ _(at::ScalarType::Double, ACL_DOUBLE) \
83+ _(at::ScalarType::ComplexHalf, ACL_DT_UNDEFINED) \
84+ _(at::ScalarType::ComplexFloat, ACL_COMPLEX64) \
85+ _(at::ScalarType::ComplexDouble, ACL_COMPLEX128) \
86+ _(at::ScalarType::Bool, ACL_BOOL) \
87+ _(at::ScalarType::QInt8, ACL_DT_UNDEFINED) \
88+ _(at::ScalarType::QUInt8, ACL_DT_UNDEFINED) \
89+ _(at::ScalarType::QInt32, ACL_DT_UNDEFINED) \
90+ _(at::ScalarType::BFloat16, ACL_BF16) \
91+ _(at::ScalarType::QUInt4x2, ACL_DT_UNDEFINED) \
92+ _(at::ScalarType::QUInt2x4, ACL_DT_UNDEFINED) \
93+ _(at::ScalarType::Undefined, ACL_DT_UNDEFINED) \
94+ _(at::ScalarType::NumOptions, ACL_DT_UNDEFINED)
95+ 
96+constexpr aclDataType kATenScalarTypeToAclDataTypeTable[static_cast<int64_t>(at::ScalarType::NumOptions) + 1] = {
97+#define DEFINE_ENUM(_1, n) n,
98+ AT_ALL_SCALAR_TYPE_AND_ACL_DATATYPE_PAIR(DEFINE_ENUM)
99+#undef DEFINE_ENUM
100+};