Star1358
1455
代码
代码
Issues116
Pull Requests260
Actions
项目讨论
Wiki
分析
Star1358
1455
已合并
feat: 新增acos_grad_v2算子A2(Ascend910B)适配实现 #4371
镜花水月1tachi创建于 7月31日
feat: 新增acos_grad_v2算子A2(Ascend910B)适配实现 #4371
已合并
镜花水月1tachi创建于 7月31日
从CDragon/ops-mathmaster合入到cann/ops-mathmaster
共 26 个文件变更+2164-0
Aexperimental/math/acos_grad_v2/CMakeLists.txt+25-0
@@ -0,0 +1,25 @@
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+# AcosGradV2: AcosGrad 算子的 A2 (Atlas A2 / Ascend910B, DAV_2201) 适配版本
12+# 对应 math/acos_grad README: z = -dy / sqrt(1 - y^2)
13+ 
14+# 设置算子定义时支持的芯片类型：A2 训练/推理系列产品 (ascend910b)
15+set(SUPPORT_COMPUTE_UNIT "ascend910b")
16+# 设置每种芯片类型对应的 tiling 文件目录：A2 -> arch32
17+set(SUPPORT_TILING_DIR "arch32")
18+add_all_modules_sources(OPTYPE acos_grad_v2 ACLNNTYPE aclnn_exclude
19+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
20+ 
21+# 显式注册 kernel 入口（arch32 / Ascend910B）
22+add_kernel_sources(
23+ KERNEL_SRC arch32/acos_grad_v2.cpp
24+ COMPUTE_UNITS ascend910b
25+)
Aexperimental/math/acos_grad_v2/README.md+75-0
@@ -0,0 +1,75 @@
1+# AcosGradV2
2+ 
3+AcosGrad 算子的 A2（Atlas A2 训练/推理系列产品，Ascend910B / DAV_2201）适配版本，基于
4+`math/acos_grad` 的功能定义实现，采用 aclnn（registry-invoke）工程结构，便于在 A2 设备上
5+编译、部署与精度验证。
6+ 
7+## 产品支持情况
8+ 
9+| 产品 | 是否支持 |
10+| :------------------------------------------------------- | :------: |
11+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
12+ 
13+> 本版本（v2）面向 A2（ascend910b）适配；原 `math/acos_grad` 面向 Ascend950（arch35）。
14+ 
15+## 功能说明
16+ 
17+- 算子功能：计算 Acos（反余弦）算子的反向梯度。
18+- 算子公式：
19+ 
20+ $$
21+ z_i = -1 \cdot dy_i \cdot \dfrac{1}{\sqrt{1 - y_i^2}}
22+ $$
23+ 
24+ 其中：
25+ - $y_i$ 为前向 Acos 算子的输入张量；
26+ - $dy_i$ 为上游传入的梯度；
27+ - $z_i$ 为对原始输入张量的梯度，等于上游梯度乘以 $-1/\sqrt{1 - y_i^2}$。
28+ 
29+- 超出定义域（$|y_i| > 1$）时：$1 - y_i^2 < 0$，平方根结果为 NaN，结果同样为 NaN。
30+ 
31+## 参数说明
32+ 
33+| 参数名 | 输入/输出/属性 | 描述 | 数据类型 | 数据格式 |
34+| :----: | :------------: | :------------------------------------------- | :----------------------- | :------: |
35+| y | 输入 | 前向 Acos 算子的输入张量。值域期望落在 [-1, 1]。 | FLOAT16, FLOAT32, BFLOAT16 | ND |
36+| dy | 输入 | 上游传入的梯度张量，shape 与 dtype 与 y 一致。 | FLOAT16, FLOAT32, BFLOAT16 | ND |
37+| z | 输出 | 对原始输入张量的梯度，shape 与 dtype 与 y 一致。 | FLOAT16, FLOAT32, BFLOAT16 | ND |
38+ 
39+## 约束说明
40+ 
41+- y 与 dy 的 shape 必须完全一致。
42+- y 与 dy 的 dtype 必须完全一致。
43+- 仅支持 ND 格式。
44+ 
45+## 性能说明
46+ 
47+CANN 当前未提供同名官方 `aclnnAcosGrad` 单算子，故以 **torch_npu 的 `torch.acos` 反向路径**
48+（用户实际调用 `autograd.grad` 时 NPU 上执行的实现）作为标杆。该路径在 NPU 上被拆分为多算子链：
49+`Acos（前向重算）+ Mul×2 + Neg×2 + Rsqrt + Adds`。本算子将其融合为单个 kernel。
50+ 
51+加速比 = 标杆反向链 device 核时之和 / 本算子单 kernel device 核时（均由 profiler 采集，
52+shape = [1024, 1024]，10 轮平均）。
53+ 
54+| 数据类型 | 标杆反向链 (us) | 本算子 (us) | 加速比 |
55+| :------: | :-------------: | :---------: | :----: |
56+| FP32 | 26.4 | 6.97 | 3.8x |
57+| FP16 | 22.6 | 7.90 | 2.9x |
58+| BF16 | 26.1 | 8.09 | 3.2x |
59+ 
60+- 三种数据类型平均加速比约 **3.3x**。
61+- 加速主要来自融合：将反向链的多次基础算子访存与启动开销合并为单次，减少 HBM 来回搬运。
62+ 
63+## A2 适配要点
64+ 
65+- 算子定义 `op_host/acos_grad_v2_def.cpp` 仅注册 `ascend910b`（AICore 配置）。
66+- `CMakeLists.txt` 中 `COMPUTE_UNIT=ascend910b`、`TILING_DIR=arch32`，并通过
67+ `add_kernel_sources` 显式注册 `arch32/acos_grad_v2.cpp` 入口。
68+- Kernel 使用标准 Ascend C 高阶向量 API（Cast/Mul/Muls/Adds/Sqrt/Div 等），
69+ FP16/BF16 先 Cast 到 FP32 计算再 Cast 回原类型，FP32 直接计算；这些 API 在 A2 上原生支持。
70+ 
71+## 调用说明
72+ 
73+| 调用方式 | 调用样例 | 说明 |
74+| ---------- | ------------------------------------------------- | ------------------------------------------------- |
75+| aclnn 调用 | [test_aclnn_acos_grad_v2](./examples/test_aclnn_acos_grad_v2.cpp) | 两段式 aclnn 调用并在 NPU 上验证精度。 |
Aexperimental/math/acos_grad_v2/examples/test_aclnn_acos_grad_v2.cpp+354-0
@@ -0,0 +1,354 @@
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 test_aclnn_acos_grad_v2.cpp
13+ * @brief aclnnAcosGradV2 调用示例与精度验证 (A2 / Ascend910B)
14+ *
15+ * 验证公式: z = -dy / sqrt(1 - y^2)
16+ * 覆盖 FP32 / FP16 / BF16 三种数据类型，y 取值落在 (-1, 1)。
17+ */
18+ 
19+#include <iostream>
20+#include <vector>
21+#include <cmath>
22+#include <cstring>
23+#include "acl/acl.h"
24+#include "../op_api/aclnn_acos_grad_v2.h"
25+ 
26+#define CHECK_RET(cond, return_expr) \
27+ do { \
28+ if (!(cond)) { \
29+ return_expr; \
30+ } \
31+ } while (0)
32+ 
33+#define LOG_PRINT(message, ...) \
34+ do { \
35+ printf(message, ##__VA_ARGS__); \
36+ } while (0)
37+ 
38+enum class TestDtype { FP32, FP16, BF16 };
39+ 
40+int64_t GetShapeSize(const std::vector<int64_t>& shape)
41+{
42+ int64_t shapeSize = 1;
43+ for (auto i : shape) {
44+ shapeSize *= i;
45+ }
46+ return shapeSize;
47+}
48+ 
49+int Init(int32_t deviceId, aclrtStream* stream)
50+{
51+ auto ret = aclInit(nullptr);
52+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
53+ ret = aclrtSetDevice(deviceId);
54+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
55+ ret = aclrtCreateStream(stream);
56+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
57+ return 0;
58+}
59+ 
60+// ---- dtype 编解码 ----
61+static void EncodeFp32(float v, uint8_t* out) { std::memcpy(out, &v, sizeof(float)); }
62+ 
63+static void EncodeFp16(float v, uint8_t* out)
64+{
65+ uint32_t x = 0;
66+ std::memcpy(&x, &v, sizeof(uint32_t));
67+ uint16_t sign = (x >> 31) & 0x1;
68+ int32_t exp = ((x >> 23) & 0xff) - 127 + 15;
69+ uint32_t mantissa = x & 0x7fffff;
70+ uint16_t h = 0;
71+ if (exp <= 0) {
72+ h = sign << 15;
73+ } else if (exp >= 31) {
74+ h = (sign << 15) | (0x1f << 10);
75+ } else {
76+ h = (sign << 15) | (exp << 10) | (mantissa >> 13);
77+ }
78+ std::memcpy(out, &h, sizeof(uint16_t));
79+}
80+ 
81+static float DecodeFp16(const uint8_t* in)
82+{
83+ uint16_t h = 0;
84+ std::memcpy(&h, in, sizeof(uint16_t));
85+ uint32_t sign = (h >> 15) & 0x1;
86+ uint32_t exp = (h >> 10) & 0x1f;
87+ uint32_t mantissa = h & 0x3ff;
88+ if (exp == 0) {
89+ if (mantissa == 0) {
90+ return sign ? -0.0f : 0.0f;
91+ }
92+ float val = mantissa / 1024.0f / 1024.0f;
93+ return sign ? -val : val;
94+ }
95+ if (exp == 31) {
96+ if (mantissa == 0) {
97+ return sign ? -INFINITY : INFINITY;
98+ }
99+ return NAN;
100+ }
101+ float val = (1.0f + mantissa / 1024.0f) * std::pow(2.0f, (int)exp - 15);
102+ return sign ? -val : val;
103+}
104+ 
105+// bfloat16: 截断 float32 低 16 位，round-to-nearest-even
106+static void EncodeBf16(float v, uint8_t* out)
107+{
108+ uint32_t x = 0;
109+ std::memcpy(&x, &v, sizeof(uint32_t));
110+ uint32_t lsb = (x >> 16) & 0x1;
111+ uint32_t rounding_bias = 0x7FFFU + lsb;
112+ uint16_t bf = static_cast<uint16_t>((x + rounding_bias) >> 16);
113+ std::memcpy(out, &bf, sizeof(uint16_t));
114+}
115+ 
116+static float DecodeBf16(const uint8_t* in)
117+{
118+ uint16_t bf = 0;
119+ std::memcpy(&bf, in, sizeof(uint16_t));
120+ uint32_t x = static_cast<uint32_t>(bf) << 16;
121+ float v = 0;
122+ std::memcpy(&v, &x, sizeof(float));
123+ return v;
124+}
125+ 
126+static void Encode(float v, TestDtype dt, uint8_t* out)
127+{
128+ switch (dt) {
129+ case TestDtype::FP32:
130+ EncodeFp32(v, out);
131+ break;
132+ case TestDtype::FP16:
133+ EncodeFp16(v, out);
134+ break;
135+ case TestDtype::BF16:
136+ EncodeBf16(v, out);
137+ break;
138+ }
139+}
140+ 
141+static float Decode(const uint8_t* in, TestDtype dt)
142+{
143+ switch (dt) {
144+ case TestDtype::FP32: {
145+ float v = 0;
146+ std::memcpy(&v, in, sizeof(float));
147+ return v;
148+ }
149+ case TestDtype::FP16:
150+ return DecodeFp16(in);
151+ case TestDtype::BF16:
152+ return DecodeBf16(in);
153+ }
154+ return 0.0f;
155+}
156+ 
157+static size_t DtypeSize(TestDtype dt) { return (dt == TestDtype::FP32) ? sizeof(float) : sizeof(uint16_t); }
158+ 
159+static aclDataType ToAclDtype(TestDtype dt)
160+{
161+ switch (dt) {
162+ case TestDtype::FP32:
163+ return aclDataType::ACL_FLOAT;
164+ case TestDtype::FP16:
165+ return aclDataType::ACL_FLOAT16;
166+ case TestDtype::BF16:
167+ return aclDataType::ACL_BF16;
168+ }
169+ return aclDataType::ACL_FLOAT;
170+}
171+ 
172+static const char* DtypeName(TestDtype dt)
173+{
174+ switch (dt) {
175+ case TestDtype::FP32:
176+ return "FP32";
177+ case TestDtype::FP16:
178+ return "FP16";
179+ case TestDtype::BF16:
180+ return "BF16";
181+ }
182+ return "?";
183+}
184+ 
185+// 创建一个 aclTensor：hostData 为已按 dtype 编码的字节流
186+int CreateAclTensor(const std::vector<uint8_t>& hostBytes, const std::vector<int64_t>& shape, void** deviceAddr,
187+ aclDataType dataType, aclTensor** tensor)
188+{
189+ auto size = hostBytes.size();
190+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
191+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
192+ ret = aclrtMemcpy(*deviceAddr, size, hostBytes.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
193+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
194+ 
195+ std::vector<int64_t> strides(shape.size(), 1);
196+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
197+ strides[i] = shape[i + 1] * strides[i + 1];
198+ }
199+ 
200+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
201+ shape.data(), shape.size(), *deviceAddr);
202+ return 0;
203+}
204+ 
205+// 运行单个 dtype 的精度验证，返回通过元素数
206+int RunOneDtype(TestDtype dt, aclrtStream stream, const std::vector<int64_t>& shape)
207+{
208+ int64_t totalNum = GetShapeSize(shape);
209+ size_t dsize = DtypeSize(dt);
210+ 
211+ std::vector<uint8_t> yBytes(totalNum * dsize, 0);
212+ std::vector<uint8_t> dyBytes(totalNum * dsize, 0);
213+ std::vector<float> yFloat(totalNum);
214+ std::vector<float> dyFloat(totalNum);
215+ 
216+ for (int64_t i = 0; i < totalNum; i++) {
217+ // y 取值 (-0.98, 0.98)，避开定义域边界
218+ float y_val = -0.98f + (1.96f * i) / totalNum;
219+ // dy 取值 (0.5, 1.5)
220+ float dy_val = 0.5f + (float)i / totalNum;
221+ 
222+ yFloat[i] = y_val;
223+ dyFloat[i] = dy_val;
224+ 
225+ Encode(y_val, dt, yBytes.data() + i * dsize);
226+ Encode(dy_val, dt, dyBytes.data() + i * dsize);
227+ }
228+ 
229+ aclTensor* yTensor = nullptr;
230+ void* yDeviceAddr = nullptr;
231+ aclTensor* dyTensor = nullptr;
232+ void* dyDeviceAddr = nullptr;
233+ aclTensor* zTensor = nullptr;
234+ void* zDeviceAddr = nullptr;
235+ void* workspaceAddr = nullptr;
236+ uint64_t workspaceSize = 0;
237+ // 提前声明所有带初始化的变量，避免 goto 跨越其初始化（gcc 严格模式）
238+ std::vector<uint8_t> zBytes(totalNum * dsize, 0);
239+ std::vector<uint8_t> resultBytes(totalNum * dsize, 0);
240+ aclOpExecutor* executor = nullptr;
241+ int passCount = 0;
242+ float relTol = (dt == TestDtype::FP32) ? 1e-5f : (dt == TestDtype::FP16 ? 1e-3f : 1e-2f);
243+ float absTol = (dt == TestDtype::FP32) ? 1e-5f : (dt == TestDtype::FP16 ? 1e-3f : 1e-2f);
244+ int retVal = -1; // 默认失败；成功路径末尾置为 passCount
245+ 
246+ auto ret = CreateAclTensor(yBytes, shape, &yDeviceAddr, ToAclDtype(dt), &yTensor);
247+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("CreateAclTensor y failed. ERROR: %d\n", ret); goto cleanup);
248+ 
249+ ret = CreateAclTensor(dyBytes, shape, &dyDeviceAddr, ToAclDtype(dt), &dyTensor);
250+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("CreateAclTensor dy failed. ERROR: %d\n", ret); goto cleanup);
251+ 
252+ ret = CreateAclTensor(zBytes, shape, &zDeviceAddr, ToAclDtype(dt), &zTensor);
253+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("CreateAclTensor z failed. ERROR: %d\n", ret); goto cleanup);
254+ 
255+ ret = aclnnAcosGradV2GetWorkspaceSize(yTensor, dyTensor, zTensor, &workspaceSize, &executor);
256+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAcosGradV2GetWorkspaceSize failed. ERROR: %d\n", ret); goto cleanup);
257+ 
258+ if (workspaceSize > 0) {
259+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
260+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); goto cleanup);
261+ }
262+ 
263+ ret = aclnnAcosGradV2(workspaceAddr, workspaceSize, executor, stream);
264+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAcosGradV2 failed. ERROR: %d\n", ret); goto cleanup);
265+ 
266+ ret = aclrtSynchronizeStream(stream);
267+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); goto cleanup);
268+ 
269+ ret = aclrtMemcpy(resultBytes.data(), resultBytes.size(), zDeviceAddr, resultBytes.size(),
270+ ACL_MEMCPY_DEVICE_TO_HOST);
271+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result failed. ERROR: %d\n", ret); goto cleanup);
272+ 
273+ // 精度比对：以“量化后的 y 和 dy”重算 golden，避免输入量化带入的偏差
274+ // 算子内部 BF16/FP16 -> FP32 计算 -> 舍回原类型，误差主要来自最后一步舍入(<=0.5 ULP)
275+ // passCount / relTol / absTol 已在函数开头声明
276+ for (int64_t i = 0; i < totalNum; i++) {
277+ float result = Decode(resultBytes.data() + i * dsize, dt);
278+ float yQuant = Decode(yBytes.data() + i * dsize, dt);
279+ float dyQuant = Decode(dyBytes.data() + i * dsize, dt);
280+ double one_minus_y2 = 1.0 - static_cast<double>(yQuant) * static_cast<double>(yQuant);
281+ float expQuant = (one_minus_y2 > 0.0) ? dyQuant * static_cast<float>(-1.0 / std::sqrt(one_minus_y2)) : NAN;
282+ float diff = std::fabs(result - expQuant);
283+ float tol = absTol + relTol * std::fabs(expQuant);
284+ if (diff <= tol) {
285+ passCount++;
286+ } else {
287+ static int printed = 0;
288+ if (printed < 8) {
289+ LOG_PRINT("FAIL[%s][%ld]: y=%.5f, dy=%.5f, expected=%.6f, result=%.6f, diff=%.6f\n", DtypeName(dt), i,
290+ yFloat[i], dyFloat[i], expQuant, result, diff);
291+ printed++;
292+ }
293+ }
294+ }
295+ 
296+ LOG_PRINT("[%s] 总元素数: %ld, 通过: %d %s\n", DtypeName(dt), totalNum, passCount,
297+ (passCount == totalNum) ? "(PASS)" : "(FAIL)");
298+ retVal = passCount;
299+ 
300+cleanup:
301+ if (yTensor != nullptr) {
302+ aclDestroyTensor(yTensor);
303+ }
304+ if (dyTensor != nullptr) {
305+ aclDestroyTensor(dyTensor);
306+ }
307+ if (zTensor != nullptr) {
308+ aclDestroyTensor(zTensor);
309+ }
310+ if (yDeviceAddr != nullptr) {
311+ aclrtFree(yDeviceAddr);
312+ }
313+ if (dyDeviceAddr != nullptr) {
314+ aclrtFree(dyDeviceAddr);
315+ }
316+ if (zDeviceAddr != nullptr) {
317+ aclrtFree(zDeviceAddr);
318+ }
319+ if (workspaceAddr != nullptr) {
320+ aclrtFree(workspaceAddr);
321+ }
322+ 
323+ return retVal;
324+}
atomgit-bot
atomgit-botatomgit-bot7月31日

🟡 Medium Priority

变更函数 RunOneDtype（第 199-302 行）在函数末尾（第 291-299 行）统一释放 yTensor/dyTensor/zTensor 及对应 deviceAddr 和 workspaceAddr。但函数中多条错误返回路径（第 225/230/236/242/247/251/254/259 行的 CHECK_RET 失败 return -1 / 第 241-242 行的 return -1）直接返回 -1，跳过了清理代码。

具体泄漏路径：

  • 第 230 行：yTensor 已创建，dyTensor 创建失败 → yDeviceAddr 和 yTensor 泄漏
  • 第 236 行：yTensor/dyTensor 已创建，zTensor 创建失败 → yDeviceAddr/dyDeviceAddr 和 yTensor/dyTensor 泄漏
  • 第 241-242 行：GetWorkspaceSize 失败 → 三个 tensor 和 deviceAddr 均泄漏
  • 第 247 行：workspace 分配失败 → 同上 + workspaceAddr（虽为 nullptr 但逻辑上资源仍在）
  • 第 251 行：aclnnAcosGradV2 执行失败
  • 第 254 行：同步失败
  • 第 259 行：结果拷贝失败

以上路径均导致已分配的 device 内存和 aclTensor 句柄泄漏。

建议：将资源清理逻辑抽取到函数末尾的公共标签（如 goto cleanup），所有错误返回路径跳转到清理标签统一释放 yTensor/dyTensor/zTensor 和对应的 deviceAddr，以及 workspaceAddr。或者使用 RAII 包装器自动管理资源生命周期。

likedislike
不准确？
325+ 
326+int main()
327+{
328+ int32_t deviceId = 0;
329+ aclrtStream stream;
330+ auto ret = Init(deviceId, &stream);
331+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
332+ 
333+ // 用一个略大的 shape 以触发多核切分 (8192 元素)
334+ std::vector<int64_t> shape = {32, 256};
335+ 
336+ LOG_PRINT("=== AcosGradV2 (A2) 精度验证, shape={32,256} ===\n");
337+ int allPass = 0;
338+ int totalDtypes = 0;
339+ for (auto dt : {TestDtype::FP32, TestDtype::FP16, TestDtype::BF16}) {
340+ totalDtypes++;
341+ int passed = RunOneDtype(dt, stream, shape);
342+ if (passed == GetShapeSize(shape)) {
343+ allPass++;
344+ }
345+ }
346+ 
347+ LOG_PRINT("=== 汇总: %d/%d dtype 全部通过 ===\n", allPass, totalDtypes);
348+ 
349+ aclrtDestroyStream(stream);
350+ aclrtResetDevice(deviceId);
351+ aclFinalize();
352+ 
353+ return (allPass == totalDtypes) ? 0 : 1;
354+}
Aexperimental/math/acos_grad_v2/op_api/aclnn_acos_grad_v2.cpp+170-0
@@ -0,0 +1,170 @@
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_acos_grad_v2.cpp
13+ * @brief ACLNN L2 API 实现 - AcosGradV2 算子 (A2 / Ascend910B)
14+ *
15+ * z = -dy / sqrt(1 - y^2)
16+ */
17+ 
18+#include "aclnn_acos_grad_v2.h"
19+#include "acos_grad_v2.h"
20+#include "aclnn_kernels/contiguous.h"
21+#include "aclnn_kernels/common/op_error_check.h"
22+#include "opdev/op_log.h"
23+#include "opdev/op_dfx.h"
24+#include "opdev/common_types.h"
25+#include "opdev/data_type_utils.h"
26+#include "opdev/make_op_executor.h"
27+#include "opdev/platform.h"
28+#include "op_api/aclnn_check.h"
29+ 
30+using namespace op;
31+ 
32+#define ACLNN_MAX_SHAPE_RANK 8
33+ 
34+static const std::initializer_list<op::DataType> ACOS_GRAD_V2_DTYPE_SUPPORT_LIST = {
35+ DataType::DT_FLOAT16, DataType::DT_FLOAT, DataType::DT_BF16};
36+ 
37+static bool IsDtypeSupported(DataType dtype)
38+{
39+ auto npuArch = GetCurrentPlatformInfo().GetCurNpuArch();
40+ // A2 (Atlas A2 训练/推理系列产品) -> DAV_2201
41+ if (npuArch == NpuArch::DAV_2201) {
42+ return CheckType(dtype, ACOS_GRAD_V2_DTYPE_SUPPORT_LIST);
43+ }
44+ return false;
45+}
46+ 
47+static bool HasEmptyTensor(const aclTensor* y, const aclTensor* dy) { return y->IsEmpty() || dy->IsEmpty(); }
48+ 
49+static bool CheckNotNull(const aclTensor* y, const aclTensor* dy, const aclTensor* z)
50+{
51+ OP_CHECK_NULL(y, return false);
52+ OP_CHECK_NULL(dy, return false);
53+ OP_CHECK_NULL(z, return false);
54+ return true;
55+}
56+ 
57+static bool CheckDtypeValid(const aclTensor* y, const aclTensor* dy, const aclTensor* z)
58+{
59+ OP_CHECK_DTYPE_NOT_MATCH(dy, y->GetDataType(), return false);
60+ OP_CHECK_DTYPE_NOT_MATCH(z, y->GetDataType(), return false);
61+ 
62+ if (!IsDtypeSupported(y->GetDataType())) {
63+ auto npuArch = GetCurrentPlatformInfo().GetCurNpuArch();
64+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
65+ "AcosGradV2: Dtype not supported: dtype=%d, npuArch=%d. "
66+ "Supported: FLOAT16, FLOAT32, BF16.",
67+ static_cast<int>(y->GetDataType()), static_cast<int>(npuArch));
68+ return false;
69+ }
70+ return true;
71+}
72+ 
73+static bool CheckFormat(const aclTensor* y, const aclTensor* dy, const aclTensor* z)
74+{
75+ auto fmtY = y->GetStorageFormat();
76+ auto fmtDy = dy->GetStorageFormat();
77+ auto fmtZ = z->GetStorageFormat();
78+ 
79+ if (IsPrivateFormat(fmtY) || IsPrivateFormat(fmtDy) || IsPrivateFormat(fmtZ)) {
80+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "AcosGradV2: Private format not supported: y=%d, dy=%d, z=%d",
81+ static_cast<int>(fmtY), static_cast<int>(fmtDy), static_cast<int>(fmtZ));
82+ return false;
83+ }
84+ return true;
85+}
86+ 
87+static bool CheckShape(const aclTensor* y, const aclTensor* dy, const aclTensor* z)
88+{
89+ OP_CHECK_MAX_DIM(y, ACLNN_MAX_SHAPE_RANK, return false);
90+ OP_CHECK_MAX_DIM(dy, ACLNN_MAX_SHAPE_RANK, return false);
91+ OP_CHECK_MAX_DIM(z, ACLNN_MAX_SHAPE_RANK, return false);
92+ 
93+ auto yShape = y->GetViewShape();
94+ auto dyShape = dy->GetViewShape();
95+ auto zShape = z->GetViewShape();
96+ 
97+ if (yShape != dyShape || yShape != zShape) {
98+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "AcosGradV2: Shape mismatch: y=%s, dy=%s, z=%s",
99+ op::ToString(yShape).GetString(), op::ToString(dyShape).GetString(), op::ToString(zShape).GetString());
100+ return false;
101+ }
102+ return true;
103+}
104+ 
105+static aclnnStatus CheckParams(const aclTensor* y, const aclTensor* dy, const aclTensor* z)
106+{
107+ if (!CheckNotNull(y, dy, z)) {
108+ OP_LOGE(ACLNN_ERR_PARAM_NULLPTR, "AcosGradV2: CheckNotNull failed");
109+ return ACLNN_ERR_PARAM_NULLPTR;
110+ }
111+ if (!CheckDtypeValid(y, dy, z)) {
112+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "AcosGradV2: CheckDtypeValid failed: y_dtype=%d, dy_dtype=%d, z_dtype=%d",
113+ static_cast<int>(y->GetDataType()), static_cast<int>(dy->GetDataType()),
114+ static_cast<int>(z->GetDataType()));
115+ return ACLNN_ERR_PARAM_INVALID;
116+ }
117+ if (!CheckFormat(y, dy, z)) {
118+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "AcosGradV2: CheckFormat failed");
119+ return ACLNN_ERR_PARAM_INVALID;
120+ }
121+ if (!CheckShape(y, dy, z)) {
122+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "AcosGradV2: CheckShape failed");
123+ return ACLNN_ERR_PARAM_INVALID;
124+ }
125+ return ACLNN_SUCCESS;
126+}
127+ 
128+extern "C" aclnnStatus aclnnAcosGradV2GetWorkspaceSize(const aclTensor* y, const aclTensor* dy, const aclTensor* z,
129+ uint64_t* workspaceSize, aclOpExecutor** executor)
130+{
131+ L2_DFX_PHASE_1(aclnnAcosGradV2, DFX_IN(y, dy), DFX_OUT(z));
132+ 
133+ OP_CHECK_NULL(workspaceSize, return ACLNN_ERR_PARAM_NULLPTR);
134+ OP_CHECK_NULL(executor, return ACLNN_ERR_PARAM_NULLPTR);
135+ 
136+ auto uniqueExecutor = CREATE_EXECUTOR();
137+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
138+ 
139+ auto ret = CheckParams(y, dy, z);
140+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
141+ 
142+ if (HasEmptyTensor(y, dy)) {
143+ *workspaceSize = 0;
144+ uniqueExecutor.ReleaseTo(executor);
145+ return ACLNN_SUCCESS;
146+ }
147+ 
148+ auto yContiguous = l0op::Contiguous(y, uniqueExecutor.get());
149+ CHECK_RET(yContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
150+ 
151+ auto dyContiguous = l0op::Contiguous(dy, uniqueExecutor.get());
152+ CHECK_RET(dyContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
153+ 
154+ const aclTensor* opResult = l0op::AcosGradV2(yContiguous, dyContiguous, uniqueExecutor.get());
155+ CHECK_RET(opResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
156+ 
157+ auto viewCopyResult = l0op::ViewCopy(opResult, z, uniqueExecutor.get());
158+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
159+ 
160+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
161+ uniqueExecutor.ReleaseTo(executor);
162+ return ACLNN_SUCCESS;
163+}
164+ 
165+extern "C" aclnnStatus aclnnAcosGradV2(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
166+ aclrtStream stream)
167+{
168+ L2_DFX_PHASE_2(aclnnAcosGradV2);
169+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
170+}
Aexperimental/math/acos_grad_v2/op_api/aclnn_acos_grad_v2.h+41-0
@@ -0,0 +1,41 @@
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_acos_grad_v2.h
13+ * @brief ACLNN L2 API 接口声明 - AcosGradV2 算子 (A2 / Ascend910B)
14+ *
15+ * z = -dy / sqrt(1 - y^2)
16+ */
17+ 
18+#ifndef ACLNN_ACOS_GRAD_V2_H_
19+#define ACLNN_ACOS_GRAD_V2_H_
20+ 
21+#include "aclnn/aclnn_base.h"
22+ 
23+#ifndef ACLNN_API
24+#define ACLNN_API __attribute__((visibility("default")))
25+#endif
26+ 
27+#ifdef __cplusplus
28+extern "C" {
29+#endif
30+ 
31+ACLNN_API aclnnStatus aclnnAcosGradV2GetWorkspaceSize(const aclTensor* y, const aclTensor* dy, const aclTensor* z,
32+ uint64_t* workspaceSize, aclOpExecutor** executor);
33+ 
34+ACLNN_API aclnnStatus aclnnAcosGradV2(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
35+ aclrtStream stream);
36+ 
37+#ifdef __cplusplus
38+}
39+#endif
40+ 
41+#endif // ACLNN_ACOS_GRAD_V2_H_
Aexperimental/math/acos_grad_v2/op_api/acos_grad_v2.cpp+89-0
@@ -0,0 +1,89 @@
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 acos_grad_v2.cpp
13+ * @brief ACLNN L0 API 实现 - AcosGradV2 算子 (A2 / Ascend910B)
14+ *
15+ * z = -dy / sqrt(1 - y^2)
16+ */
17+ 
18+#include "acos_grad_v2.h"
19+#include "opdev/op_log.h"
20+#include "opdev/op_dfx.h"
21+#include "opdev/shape_utils.h"
22+#include "opdev/make_op_executor.h"
23+#include "op_api/aclnn_check.h"
24+ 
25+using namespace op;
26+ 
27+namespace l0op {
28+ 
29+OP_TYPE_REGISTER(AcosGradV2);
30+ 
31+static const std::initializer_list<op::DataType> ACOS_GRAD_V2_DTYPE_SUPPORT_LIST = {
32+ DataType::DT_FLOAT16, DataType::DT_FLOAT, DataType::DT_BF16};
33+ 
34+static bool IsAiCoreSupport(const aclTensor* y, const aclTensor* dy)
35+{
36+ auto npuArch = GetCurrentPlatformInfo().GetCurNpuArch();
37+ // A2 (Atlas A2 训练/推理系列产品) -> DAV_2201
38+ if (npuArch == NpuArch::DAV_2201) {
39+ return CheckType(y->GetDataType(), ACOS_GRAD_V2_DTYPE_SUPPORT_LIST) &&
40+ CheckType(dy->GetDataType(), ACOS_GRAD_V2_DTYPE_SUPPORT_LIST);
41+ }
42+ return false;
43+}
44+ 
45+static bool AcosGradV2InferShape(const op::Shape& yShape, const op::Shape& dyShape, op::Shape& outShape)
46+{
47+ if (yShape != dyShape) {
48+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "AcosGradV2: Shape mismatch: y=%s, dy=%s", op::ToString(yShape).GetString(),
49+ op::ToString(dyShape).GetString());
50+ return false;
51+ }
52+ outShape = yShape;
53+ return true;
54+}
55+ 
56+static const aclTensor* AcosGradV2AiCore(const aclTensor* y, const aclTensor* dy, const aclTensor* z,
57+ aclOpExecutor* executor)
58+{
59+ L0_DFX(AcosGradV2AiCore, y, dy, z);
60+ 
61+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(AcosGradV2, OP_INPUT(y, dy), OP_OUTPUT(z));
62+ OP_CHECK(ret == ACLNN_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "AcosGradV2AiCore failed."), return nullptr);
63+ return z;
64+}
65+ 
66+const aclTensor* AcosGradV2(const aclTensor* y, const aclTensor* dy, aclOpExecutor* executor)
67+{
68+ Shape outShape;
69+ const aclTensor* out = nullptr;
70+ 
71+ if (!AcosGradV2InferShape(y->GetViewShape(), dy->GetViewShape(), outShape)) {
72+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "AcosGradV2: Infer shape failed.");
73+ return nullptr;
74+ }
75+ 
76+ if (!IsAiCoreSupport(y, dy)) {
77+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
78+ "AcosGradV2 not supported: dtype y=%d, dy=%d. Supported dtypes: FLOAT16, FLOAT32, BF16.",
79+ static_cast<int>(y->GetDataType()), static_cast<int>(dy->GetDataType()));
80+ return nullptr;
81+ }
82+ 
83+ out = executor->AllocTensor(outShape, y->GetDataType());
84+ OP_CHECK(out != nullptr, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "AcosGradV2: AllocTensor failed."), return nullptr);
85+ 
86+ return AcosGradV2AiCore(y, dy, out, executor);
atomgit-bot
atomgit-botatomgit-bot7月31日

🟡 Medium Priority

变更行：op_api/acos_grad_v2.cpp 第 85 行 out = executor->AllocTensor(outShape, y->GetDataType()); 的返回值未做空指针检查即传入 AcosGradV2AiCore（第 87 行），继而传入 ADD_TO_LAUNCHER_LIST_AICORE。

影响：若 AllocTensor 分配失败返回 nullptr，下游宏/函数可能对空指针解引用，触发崩溃或 AIC Error。

同类算子 mul_no_nan（mul_no_nan/op_api/mul_no_nan.cpp:87）和 div_v3（div_v3/op_api/div_v3.cpp:45）均在此处做了 OP_CHECK(out != nullptr, ...)，本文件遗漏了该检查。

建议：在 AllocTensor 之后增加 OP_CHECK 空指针检查，失败时打印错误日志并返回 nullptr

改动建议
86
+ out = executor->AllocTensor(outShape, y->GetDataType());
87
+ OP_CHECK(out != nullptr, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "AcosGradV2: AllocTensor failed."), return nullptr);
88
+
86
89
  return AcosGradV2AiCore(y, dy, out, executor);
应用建议
likedislike
不准确？
87+}
88+ 
89+} // namespace l0op
Aexperimental/math/acos_grad_v2/op_api/acos_grad_v2.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 acos_grad_v2.h
13+ * @brief ACLNN L0 API 接口声明 - AcosGradV2 算子 (A2 / Ascend910B)
14+ *
15+ * 对齐 math/acos_grad README: 输入 y（前向 Acos 输入）、dy（上游梯度），输出 z（对原始输入的梯度）。
16+ * 公式: z = -dy / sqrt(1 - y^2)
17+ */
18+ 
19+#ifndef OP_API_INC_LEVEL0_ACOS_GRAD_V2_H_
20+#define OP_API_INC_LEVEL0_ACOS_GRAD_V2_H_
21+ 
22+#include "opdev/op_executor.h"
23+ 
24+namespace l0op {
25+ 
26+const aclTensor* AcosGradV2(const aclTensor* y, const aclTensor* dy, aclOpExecutor* executor);
27+ 
28+} // namespace l0op
29+ 
30+#endif // OP_API_INC_LEVEL0_ACOS_GRAD_V2_H_
Aexperimental/math/acos_grad_v2/op_host/acos_grad_v2_def.cpp+66-0
@@ -0,0 +1,66 @@
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 acos_grad_v2_def.cpp
13+ * \brief AcosGradV2 算子定义（A2 / Ascend910B 适配）
14+ *
15+ * 对齐 math/acos_grad README:
16+ * INPUT(y) : 前向 Acos 的输入张量，值域期望 [-1, 1]
17+ * INPUT(dy) : 上游梯度，shape/dtype 与 y 一致
18+ * OUTPUT(z) : 对原始输入的梯度，shape/dtype 与 y 一致
19+ *
20+ * 公式: z = -dy / sqrt(1 - y^2)
21+ */
22+#include "register/op_def_registry.h"
23+ 
24+namespace ops {
25+class AcosGradV2 : public OpDef {
26+public:
27+ explicit AcosGradV2(const char* name) : OpDef(name)
28+ {
29+ const std::vector<ge::DataType> AcosGradV2DataType = {ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16};
30+ const std::vector<ge::Format> AcosGradV2Format = {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
31+ 
32+ this->Input("y")
33+ .ParamType(REQUIRED)
34+ .DataType(AcosGradV2DataType)
35+ .Format(AcosGradV2Format)
36+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
37+ .AutoContiguous();
38+ 
39+ this->Input("dy")
40+ .ParamType(REQUIRED)
41+ .DataType(AcosGradV2DataType)
42+ .Format(AcosGradV2Format)
43+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
44+ .AutoContiguous();
45+ 
46+ this->Output("z")
47+ .ParamType(REQUIRED)
48+ .DataType(AcosGradV2DataType)
49+ .Format(AcosGradV2Format)
50+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
51+ .AutoContiguous();
52+ 
53+ OpAICoreConfig aiCoreConfig;
54+ aiCoreConfig.DynamicCompileStaticFlag(true)
55+ .DynamicFormatFlag(false)
56+ .DynamicShapeSupportFlag(true)
57+ .PrecisionReduceFlag(true)
58+ .NeedCheckSupportFlag(false)
59+ .DynamicRankSupportFlag(true)
60+ .ExtendCfgInfo("opFile.value", "acos_grad_v2");
61+ // A2 适配：仅注册 ascend910b (Atlas A2 训练/推理系列产品, DAV_2201)
62+ this->AICore().AddConfig("ascend910b", aiCoreConfig);
63+ }
64+};
65+OP_ADD(AcosGradV2);
66+} // namespace ops
Aexperimental/math/acos_grad_v2/op_host/acos_grad_v2_infershape.cpp+48-0
@@ -0,0 +1,48 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file acos_grad_v2_infershape.cpp
13+ * \brief AcosGradV2 形状/类型推导（z 的 shape 与 dtype 均等于 y）
14+ */
15+ 
16+#include "register/op_impl_registry.h"
17+#include "exe_graph/runtime/infer_shape_context.h"
18+ 
19+using namespace ge;
20+ 
21+namespace ops {
22+ 
23+static ge::graphStatus InferShape4AcosGradV2(gert::InferShapeContext* context)
24+{
25+ const gert::Shape* yShape = context->GetInputShape(0);
26+ if (yShape == nullptr) {
27+ return ge::GRAPH_FAILED;
28+ }
29+ 
30+ const gert::Shape* dyShape = context->GetInputShape(1);
31+ if (dyShape == nullptr) {
32+ return ge::GRAPH_FAILED;
33+ }
34+ 
35+ gert::Shape* outputShape = context->GetOutputShape(0);
36+ if (outputShape == nullptr) {
37+ return ge::GRAPH_FAILED;
38+ }
39+ 
40+ // 输出 z 的 shape 与 y 一致
41+ *outputShape = *yShape;
42+ 
43+ return ge::GRAPH_SUCCESS;
44+}
45+ 
46+IMPL_OP_INFERSHAPE(AcosGradV2).InferShape(InferShape4AcosGradV2);
47+ 
48+} // namespace ops
Aexperimental/math/acos_grad_v2/op_host/arch32/acos_grad_v2_tiling.cpp+180-0
@@ -0,0 +1,180 @@
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 acos_grad_v2_tiling.cpp
13+ * \brief AcosGradV2 Tiling — 多核 + UB 两级切分（arch32 / Ascend910B）
14+ *
15+ * 切分策略：
16+ * 1) 多核（blockDim）：按 ELEM_ALIGN(512) 对齐切 totalLength → blockFormer × blockNum
17+ * 2) 核内 UB：按 dtype 的 bytesPerElem / alignFactor 从动态获取的 UB size 推出 ubFormer
18+ * 3) 核内 loop/tail 由 kernel 侧按 blockLength_ 自行推导，tiling 只下发 blockFormer/ubFormer
19+ */
20+ 
21+#include "register/op_def_registry.h"
22+#include "log/log.h"
23+#include "util/math_util.h"
24+#include "util/platform_util.h"
25+#include "../../op_kernel/arch32/acos_grad_v2_tiling_data.h"
26+#include "../../op_kernel/arch32/acos_grad_v2_tiling_key.h"
27+ 
28+namespace optiling {
29+ 
30+using Ops::Base::CeilDiv;
31+using Ops::Base::FloorDiv;
32+ 
33+constexpr uint32_t WS_SYS_SIZE = 0U;
34+constexpr uint32_t ELEM_ALIGN = 512U;
35+constexpr uint32_t FP32_BYTES_PER_ELEM = 32U;
36+constexpr uint32_t FP32_ALIGN = 64U;
37+constexpr uint32_t LOWPREC_BYTES_PER_ELEM = 28U;
38+constexpr uint32_t LOWPREC_ALIGN = 128U;
39+ 
40+static const gert::Shape g_scalar_to_vec1 = {1};
41+ 
42+// ---- 核内 UB 切分计算：按 dtype 选 bytesPerElem / alignFactor，推出 ubFormer ----
43+// ubSize 由平台信息动态获取（不同款型 UB size 未必相同），不再写死 184KB
44+static void CalcAcosGradV2UbTiling(uint64_t totalLength, uint32_t availCoreNum, uint64_t ubSize, ge::DataType dataType,
45+ AcosGradV2TilingData* tiling)
46+{
47+ // 多核切分：每核至少 ELEM_ALIGN 个元素，向上对齐
48+ uint32_t coreNum = static_cast<uint32_t>(
49+ CeilDiv(static_cast<int64_t>(totalLength), static_cast<int64_t>(ELEM_ALIGN)));
50+ coreNum = std::min(coreNum, availCoreNum);
51+ coreNum = std::max(coreNum, 1U);
52+ 
53+ uint32_t blockFormerRaw = static_cast<uint32_t>(
54+ CeilDiv(static_cast<int64_t>(totalLength), static_cast<int64_t>(coreNum)));
55+ uint32_t blockFormer = static_cast<uint32_t>(
56+ CeilDiv(static_cast<int64_t>(blockFormerRaw), static_cast<int64_t>(ELEM_ALIGN)) * ELEM_ALIGN);
57+ blockFormer = std::max(blockFormer, ELEM_ALIGN);
58+ 
59+ uint32_t blockNum = static_cast<uint32_t>(
60+ CeilDiv(static_cast<int64_t>(totalLength), static_cast<int64_t>(blockFormer)));
61+ blockNum = std::max(blockNum, 1U);
62+ 
63+ // UB 切分：按 dtype 确定每元素占用字节数和对齐因子
64+ uint32_t bytesPerElem = (dataType == ge::DT_FLOAT) ? FP32_BYTES_PER_ELEM : LOWPREC_BYTES_PER_ELEM;
65+ uint32_t alignFactor = (dataType == ge::DT_FLOAT) ? FP32_ALIGN : LOWPREC_ALIGN;
66+ 
67+ uint32_t ubFormerRaw = static_cast<uint32_t>(ubSize / bytesPerElem);
68+ uint32_t ubFormer = static_cast<uint32_t>(
69+ FloorDiv(static_cast<int64_t>(ubFormerRaw), static_cast<int64_t>(alignFactor)) * alignFactor);
70+ ubFormer = std::max(ubFormer, alignFactor);
71+ ubFormer = std::min(ubFormer, blockFormer);
72+ 
73+ // 核内 loop/tail 由 kernel 侧按 blockLength_ 自行推导，无需在此预计算
74+ tiling->totalLength = totalLength;
75+ tiling->blockFormer = blockFormer;
76+ tiling->blockNum = blockNum;
77+ tiling->ubFormer = ubFormer;
78+}
79+ 
80+// 本算子硬件能力：AIV 核数 + UB 字节数合并查询（区别于模板的逐项 out-param 写法）
81+struct AcosGradV2HwCap {
82+ uint32_t aivCoreNum = 0;
83+ uint64_t ubBytes = 0;
84+};
85+ 
86+static ge::graphStatus QueryAcosGradV2HwCap(gert::TilingContext* context, AcosGradV2HwCap& cap)
87+{
88+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
89+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
90+ platform_ascendc::PlatformAscendC plat(platformInfoPtr);
91+ 
92+ int64_t aiv = plat.GetCoreNumAiv();
93+ OP_CHECK_IF(aiv <= 0, OP_LOGE(context, "AcosGradV2: invalid AIV core count %ld", aiv), return ge::GRAPH_FAILED);
94+ cap.aivCoreNum = static_cast<uint32_t>(aiv);
95+ 
96+ plat.GetCoreMemSize(platform_ascendc::CoreMemType::UB, cap.ubBytes);
97+ OP_CHECK_IF(cap.ubBytes == 0, OP_LOGE(context, "AcosGradV2: UB size unavailable on this SoC"),
98+ return ge::GRAPH_FAILED);
99+ return ge::GRAPH_SUCCESS;
100+}
101+ 
102+// 把存储 shape 中的标量(0维)规整为 {1}，便于按一维长度统一切分
103+static gert::Shape AsVecIfScalar(const gert::Shape& storageShape)
104+{
105+ return (storageShape.GetDimNum() == 0) ? g_scalar_to_vec1 : storageShape;
106+}
107+ 
108+// ---- tiling 入口 ----
109+static ge::graphStatus AcosGradV2TilingFunc(gert::TilingContext* context)
110+{
111+ OP_LOGI(context->GetNodeName(), "Enter AcosGradV2TilingFunc");
112+ 
113+ // 1) shape：标量视作 {1}，并校验 y/dy/z 三者元素数一致
114+ auto inputY = context->GetInputShape(0);
115+ OP_CHECK_NULL_WITH_CONTEXT(context, inputY);
116+ auto inputDy = context->GetInputShape(1);
117+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDy);
118+ auto outputZ = context->GetOutputShape(0);
119+ OP_CHECK_NULL_WITH_CONTEXT(context, outputZ);
120+ 
121+ gert::Shape yShape = AsVecIfScalar(inputY->GetStorageShape());
122+ gert::Shape dyShape = AsVecIfScalar(inputDy->GetStorageShape());
123+ gert::Shape zShape = AsVecIfScalar(outputZ->GetStorageShape());
124+ int64_t yElemCnt = yShape.GetShapeSize();
125+ OP_CHECK_IF(yElemCnt != dyShape.GetShapeSize() || yElemCnt != zShape.GetShapeSize(),
126+ OP_LOGE(context, "AcosGradV2: shape size mismatch: y=%ld, dy=%ld, z=%ld", yElemCnt,
127+ dyShape.GetShapeSize(), zShape.GetShapeSize()),
128+ return ge::GRAPH_FAILED);
129+ uint64_t totalLength = static_cast<uint64_t>(yElemCnt);
130+ 
131+ // 2) dtype 仅支持 fp16/fp32/bf16
132+ auto inputDesc = context->GetInputDesc(0);
133+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
134+ ge::DataType dataType = inputDesc->GetDataType();
135+ OP_CHECK_IF(dataType != ge::DT_FLOAT16 && dataType != ge::DT_FLOAT && dataType != ge::DT_BF16,
136+ OP_LOGE(context, "AcosGradV2: unsupported dtype %d", static_cast<int>(dataType)),
137+ return ge::GRAPH_FAILED);
138+ 
139+ // 3) 取 tiling 缓冲并清零
140+ AcosGradV2TilingData* tiling = context->GetTilingData<AcosGradV2TilingData>();
141+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
142+ OP_CHECK_IF(memset_s(tiling, sizeof(AcosGradV2TilingData), 0, sizeof(AcosGradV2TilingData)) != EOK,
143+ OP_LOGE(context, "AcosGradV2: memset_s tiling data error"), return ge::GRAPH_FAILED);
144+ 
145+ // 4) 空 tensor 无需查询硬件，单核直接返回
146+ if (totalLength == 0UL) {
147+ context->SetBlockDim(1U);
148+ ASCENDC_TPL_SEL_PARAM(context, static_cast<uint32_t>(dataType));
149+ return ge::GRAPH_SUCCESS;
150+ }
151+ 
152+ // 5) 查询硬件能力(AIV 核数 + UB)，再做多核 + UB 两级切分
153+ AcosGradV2HwCap cap;
154+ OP_CHECK_IF(QueryAcosGradV2HwCap(context, cap) != ge::GRAPH_SUCCESS,
155+ OP_LOGE(context, "AcosGradV2: query hardware capability failed"), return ge::GRAPH_FAILED);
156+ CalcAcosGradV2UbTiling(totalLength, cap.aivCoreNum, cap.ubBytes, dataType, tiling);
157+ context->SetBlockDim(tiling->blockNum);
158+ 
159+ // 6) workspace
160+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
161+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
162+ currentWorkspace[0] = WS_SYS_SIZE;
163+ 
164+ OP_LOGI(context, "[AcosGradV2 Tiling] totalLength=%lu, blockFormer=%u, blockNum=%u, ubFormer=%u",
165+ tiling->totalLength, tiling->blockFormer, tiling->blockNum, tiling->ubFormer);
166+ 
167+ ASCENDC_TPL_SEL_PARAM(context, static_cast<uint32_t>(dataType));
168+ return ge::GRAPH_SUCCESS;
169+}
170+ 
171+static ge::graphStatus TilingParseForAcosGradV2([[maybe_unused]] gert::TilingParseContext* context)
172+{
173+ return ge::GRAPH_SUCCESS;
174+}
175+ 
176+struct AcosGradV2CompileInfo {};
177+ 
178+IMPL_OP_OPTILING(AcosGradV2).Tiling(AcosGradV2TilingFunc).TilingParse<AcosGradV2CompileInfo>(TilingParseForAcosGradV2);
179+ 
180+} // namespace optiling
Aexperimental/math/acos_grad_v2/op_kernel/arch32/acos_grad_v2.cpp+33-0
@@ -0,0 +1,33 @@
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 acos_grad_v2.cpp
13+ * \brief AcosGradV2 算子 Kernel 入口（arch32 / Ascend910B）
14+ *
15+ * Inputs (对齐 def.cpp / README):
16+ * y : 前向 Acos 的输入张量
17+ * dy : 上游梯度
18+ * z : 输出梯度
19+ *
20+ * 公式: z = -dy / sqrt(1 - y^2)
21+ */
22+ 
23+#include "acos_grad_v2.h"
24+ 
25+template <typename D_T>
26+__global__ __aicore__ void acos_grad_v2(GM_ADDR y, GM_ADDR dy, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
27+{
28+ REGISTER_TILING_DEFAULT(AcosGradV2TilingData);
29+ GET_TILING_DATA_WITH_STRUCT(AcosGradV2TilingData, tilingData, tiling);
30+ NsAcosGradV2::KernelAcosGradV2<D_T> op;
31+ op.Init(y, dy, z, &tilingData);
32+ op.Process();
33+}
Aexperimental/math/acos_grad_v2/op_kernel/arch32/acos_grad_v2.h+378-0
@@ -0,0 +1,378 @@
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 acos_grad_v2.h
13+ * \brief KernelAcosGradV2 类实现（arch32 / Ascend910B）
14+ *
15+ * 对齐 math/acos_grad README:
16+ * y : 前向 Acos 的输入张量
17+ * dy : 上游梯度
18+ * z : 对原始输入的梯度
19+ *
20+ * 公式: z = -dy / sqrt(1 - y^2)
21+ *
22+ * 超出定义域（|y| > 1）时：y^2 > 1 -> 1-y^2 < 0 -> sqrt(负数) = NaN
23+ *
24+ * FP16/BF16 路径：Cast -> float32 计算 -> Cast 回原始类型
25+ * FP32 路径：直接计算
26+ */
27+ 
28+#ifndef ACOS_GRAD_V2_H
29+#define ACOS_GRAD_V2_H
30+ 
31+#include "kernel_operator.h"
32+#include "kernel_tiling/kernel_tiling.h"
33+#include "acos_grad_v2_tiling_data.h"
34+#include "acos_grad_v2_tiling_key.h"
35+ 
36+namespace NsAcosGradV2 {
37+ 
38+using namespace AscendC;
39+ 
40+template <typename T>
41+class KernelAcosGradV2 {
42+ static constexpr int32_t BUFFER_NUM = 2;
43+ static constexpr int32_t TMP_BUFFER_NUM = 1;
44+ 
45+public:
46+ __aicore__ inline KernelAcosGradV2() {}
47+ 
48+ __aicore__ inline void Init(GM_ADDR y, GM_ADDR dy, GM_ADDR z, const AcosGradV2TilingData* tilingData);
49+ __aicore__ inline void Process();
50+ 
51+private:
52+ __aicore__ inline void ProcessBlock(uint32_t ubLoop, uint32_t ubTail);
53+ __aicore__ inline void CopyIn(int64_t progress, uint32_t currentNum);
54+ __aicore__ inline void Compute(uint32_t currentNum);
55+ __aicore__ inline void CopyOut(int64_t progress, uint32_t currentNum);
56+ 
57+private:
58+ TPipe pipe_;
59+ 
60+ TQue<TPosition::VECIN, BUFFER_NUM> inQueueY_;
61+ TQue<TPosition::VECIN, BUFFER_NUM> inQueueDy_;
62+ TQue<TPosition::VECOUT, BUFFER_NUM> outQueueZ_;
63+ 
64+ TBuf<TPosition::VECCALC> yF32Buf_;
65+ TBuf<TPosition::VECCALC> dyF32Buf_;
66+ TBuf<TPosition::VECCALC> tmpF32ABuf_;
67+ TBuf<TPosition::VECCALC> tmpF32BBuf_;
68+ 
69+ GlobalTensor<T> yGm_;
70+ GlobalTensor<T> dyGm_;
71+ GlobalTensor<T> zGm_;
72+ 
73+ uint32_t blockIdx_ = 0;
74+ uint32_t blockFormer_ = 0;
75+ uint32_t ubFormer_ = 0;
76+ int64_t blockOffset_ = 0;
77+ uint32_t blockLength_ = 0;
78+};
79+ 
80+template <typename T>
81+__aicore__ inline void KernelAcosGradV2<T>::Init(GM_ADDR y, GM_ADDR dy, GM_ADDR z,
82+ const AcosGradV2TilingData* tilingData)
83+{
84+ blockIdx_ = GetBlockIdx();
85+ blockFormer_ = tilingData->blockFormer;
86+ ubFormer_ = tilingData->ubFormer;
87+ 
88+ uint64_t totalLength = tilingData->totalLength;
89+ uint64_t start = static_cast<uint64_t>(blockIdx_) * blockFormer_;
90+ 
91+ if (start >= totalLength) {
92+ blockLength_ = 0;
93+ return;
94+ }
95+ 
96+ uint64_t remaining = totalLength - start;
97+ blockLength_ = static_cast<uint32_t>((remaining < blockFormer_) ? remaining : blockFormer_);
98+ blockOffset_ = static_cast<int64_t>(start);
99+ 
100+ yGm_.SetGlobalBuffer((__gm__ T*)y + blockOffset_, blockLength_);
101+ dyGm_.SetGlobalBuffer((__gm__ T*)dy + blockOffset_, blockLength_);
102+ zGm_.SetGlobalBuffer((__gm__ T*)z + blockOffset_, blockLength_);
103+ 
104+ pipe_.InitBuffer(inQueueY_, BUFFER_NUM, ubFormer_ * sizeof(T));
105+ pipe_.InitBuffer(inQueueDy_, BUFFER_NUM, ubFormer_ * sizeof(T));
106+ pipe_.InitBuffer(outQueueZ_, BUFFER_NUM, ubFormer_ * sizeof(T));
107+ 
108+ pipe_.InitBuffer(yF32Buf_, ubFormer_ * sizeof(float));
109+ pipe_.InitBuffer(dyF32Buf_, ubFormer_ * sizeof(float));
110+ pipe_.InitBuffer(tmpF32ABuf_, ubFormer_ * sizeof(float));
111+ pipe_.InitBuffer(tmpF32BBuf_, ubFormer_ * sizeof(float));
112+}
113+ 
114+template <typename T>
115+__aicore__ inline void KernelAcosGradV2<T>::CopyIn(int64_t progress, uint32_t currentNum)
116+{
117+ LocalTensor<T> yLocal = inQueueY_.template AllocTensor<T>();
118+ LocalTensor<T> dyLocal = inQueueDy_.template AllocTensor<T>();
119+ 
120+ AscendC::DataCopyExtParams copyParams{1, static_cast<uint32_t>(currentNum * sizeof(T)), 0, 0, 0};
121+ AscendC::DataCopyPadExtParams<T> padParams{false, 0, 0, T(0)};
122+ 
123+ DataCopyPad(yLocal, yGm_[progress * static_cast<int64_t>(ubFormer_)], copyParams, padParams);
124+ DataCopyPad(dyLocal, dyGm_[progress * static_cast<int64_t>(ubFormer_)], copyParams, padParams);
125+ 
126+ inQueueY_.EnQue(yLocal);
127+ inQueueDy_.EnQue(dyLocal);
128+}
129+ 
130+template <typename T>
131+__aicore__ inline void KernelAcosGradV2<T>::Compute(uint32_t currentNum)
132+{
133+ LocalTensor<T> yLocal = inQueueY_.template DeQue<T>();
134+ LocalTensor<T> dyLocal = inQueueDy_.template DeQue<T>();
135+ LocalTensor<T> zLocal = outQueueZ_.template AllocTensor<T>();
136+ 
137+ LocalTensor<float> yF32 = yF32Buf_.Get<float>();
138+ LocalTensor<float> dyF32 = dyF32Buf_.Get<float>();
139+ LocalTensor<float> tmpF32A = tmpF32ABuf_.Get<float>();
140+ LocalTensor<float> tmpF32B = tmpF32BBuf_.Get<float>();
141+ 
142+ // FP16/BF16 -> FP32
143+ Cast(yF32, yLocal, RoundMode::CAST_NONE, currentNum);
144+ PipeBarrier<PIPE_V>();
145+ Cast(dyF32, dyLocal, RoundMode::CAST_NONE, currentNum);
146+ PipeBarrier<PIPE_V>();
147+ 
148+ // tmpF32A = 1 - y^2
149+ Mul(tmpF32A, yF32, yF32, currentNum);
150+ PipeBarrier<PIPE_V>();
151+ 
152+ Muls(tmpF32A, tmpF32A, static_cast<float>(-1.0f), currentNum);
153+ PipeBarrier<PIPE_V>();
154+ 
155+ Adds(tmpF32B, tmpF32A, static_cast<float>(1.0f), currentNum);
156+ PipeBarrier<PIPE_V>();
157+ 
158+ // tmpF32B = sqrt(1 - y^2)
159+ Sqrt(tmpF32B, tmpF32B, currentNum);
160+ PipeBarrier<PIPE_V>();
161+ 
162+ // dyF32 = -dy / sqrt(1 - y^2) = z
163+ Muls(dyF32, dyF32, static_cast<float>(-1.0f), currentNum);
164+ PipeBarrier<PIPE_V>();
165+ 
166+ Div(dyF32, dyF32, tmpF32B, currentNum);
167+ PipeBarrier<PIPE_V>();
168+ 
169+ // FP32 -> 原始类型
170+ Cast(zLocal, dyF32, RoundMode::CAST_RINT, currentNum);
171+ PipeBarrier<PIPE_V>();
172+ 
173+ outQueueZ_.template EnQue<T>(zLocal);
174+ inQueueY_.FreeTensor(yLocal);
175+ inQueueDy_.FreeTensor(dyLocal);
176+}
177+ 
178+template <typename T>
179+__aicore__ inline void KernelAcosGradV2<T>::CopyOut(int64_t progress, uint32_t currentNum)
180+{
181+ LocalTensor<T> zLocal = outQueueZ_.template DeQue<T>();
182+ 
183+ int64_t gmOffset = progress * static_cast<int64_t>(ubFormer_);
184+ AscendC::DataCopyExtParams copyParams{1, static_cast<uint32_t>(currentNum * sizeof(T)), 0, 0, 0};
185+ DataCopyPad(zGm_[gmOffset], zLocal, copyParams);
186+ 
187+ outQueueZ_.FreeTensor(zLocal);
188+}
189+ 
190+template <typename T>
191+__aicore__ inline void KernelAcosGradV2<T>::ProcessBlock(uint32_t ubLoop, uint32_t ubTail)
192+{
193+ for (uint32_t i = 0; i < ubLoop; i++) {
194+ CopyIn(static_cast<int64_t>(i), ubFormer_);
195+ Compute(ubFormer_);
196+ CopyOut(static_cast<int64_t>(i), ubFormer_);
197+ }
198+ if (ubTail > 0) {
199+ CopyIn(static_cast<int64_t>(ubLoop), ubTail);
200+ Compute(ubTail);
201+ CopyOut(static_cast<int64_t>(ubLoop), ubTail);
202+ }
203+}
204+ 
205+template <typename T>
206+__aicore__ inline void KernelAcosGradV2<T>::Process()
207+{
208+ if (blockLength_ == 0) {
209+ return;
210+ }
211+ 
212+ uint32_t ubLoop = blockLength_ / ubFormer_;
213+ uint32_t ubTail = blockLength_ % ubFormer_;
214+ 
215+ ProcessBlock(ubLoop, ubTail);
216+}
217+ 
218+// FP32 特化：无需 Cast，直接计算
219+template <>
220+class KernelAcosGradV2<float> {
221+ static constexpr int32_t BUFFER_NUM = 2;
222+ static constexpr int32_t TMP_BUFFER_NUM = 1;
223+ 
224+public:
225+ __aicore__ inline KernelAcosGradV2() {}
226+ 
227+ __aicore__ inline void Init(GM_ADDR y, GM_ADDR dy, GM_ADDR z, const AcosGradV2TilingData* tilingData);
228+ __aicore__ inline void Process();
229+ 
230+private:
231+ __aicore__ inline void ProcessBlock(uint32_t ubLoop, uint32_t ubTail);
232+ __aicore__ inline void CopyIn(int64_t progress, uint32_t currentNum);
233+ __aicore__ inline void Compute(uint32_t currentNum);
234+ __aicore__ inline void CopyOut(int64_t progress, uint32_t currentNum);
235+ 
236+private:
237+ TPipe pipe_;
238+ 
239+ TQue<TPosition::VECIN, BUFFER_NUM> inQueueY_;
240+ TQue<TPosition::VECIN, BUFFER_NUM> inQueueDy_;
241+ TQue<TPosition::VECOUT, BUFFER_NUM> outQueueZ_;
242+ 
243+ TQue<TPosition::VECCALC, TMP_BUFFER_NUM> tmpQueue1_;
244+ TQue<TPosition::VECCALC, TMP_BUFFER_NUM> tmpQueue2_;
245+ 
246+ GlobalTensor<float> yGm_;
247+ GlobalTensor<float> dyGm_;
248+ GlobalTensor<float> zGm_;
249+ 
250+ uint32_t blockIdx_ = 0;
251+ uint32_t blockFormer_ = 0;
252+ uint32_t ubFormer_ = 0;
253+ int64_t blockOffset_ = 0;
254+ uint32_t blockLength_ = 0;
255+};
256+ 
257+__aicore__ inline void KernelAcosGradV2<float>::Init(GM_ADDR y, GM_ADDR dy, GM_ADDR z,
258+ const AcosGradV2TilingData* tilingData)
259+{
260+ blockIdx_ = GetBlockIdx();
261+ blockFormer_ = tilingData->blockFormer;
262+ ubFormer_ = tilingData->ubFormer;
263+ 
264+ uint64_t totalLength = tilingData->totalLength;
265+ uint64_t start = static_cast<uint64_t>(blockIdx_) * blockFormer_;
266+ 
267+ if (start >= totalLength) {
268+ blockLength_ = 0;
269+ return;
270+ }
271+ 
272+ uint64_t remaining = totalLength - start;
273+ blockLength_ = static_cast<uint32_t>((remaining < blockFormer_) ? remaining : blockFormer_);
274+ blockOffset_ = static_cast<int64_t>(start);
275+ 
276+ yGm_.SetGlobalBuffer((__gm__ float*)y + blockOffset_, blockLength_);
277+ dyGm_.SetGlobalBuffer((__gm__ float*)dy + blockOffset_, blockLength_);
278+ zGm_.SetGlobalBuffer((__gm__ float*)z + blockOffset_, blockLength_);
279+ 
280+ pipe_.InitBuffer(inQueueY_, BUFFER_NUM, ubFormer_ * sizeof(float));
281+ pipe_.InitBuffer(inQueueDy_, BUFFER_NUM, ubFormer_ * sizeof(float));
282+ pipe_.InitBuffer(outQueueZ_, BUFFER_NUM, ubFormer_ * sizeof(float));
283+ pipe_.InitBuffer(tmpQueue1_, TMP_BUFFER_NUM, ubFormer_ * sizeof(float));
284+ pipe_.InitBuffer(tmpQueue2_, TMP_BUFFER_NUM, ubFormer_ * sizeof(float));
285+}
286+ 
287+__aicore__ inline void KernelAcosGradV2<float>::CopyIn(int64_t progress, uint32_t currentNum)
288+{
289+ LocalTensor<float> yLocal = inQueueY_.template AllocTensor<float>();
290+ LocalTensor<float> dyLocal = inQueueDy_.template AllocTensor<float>();
291+ 
292+ AscendC::DataCopyExtParams copyParams{1, static_cast<uint32_t>(currentNum * sizeof(float)), 0, 0, 0};
293+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.0f};
294+ 
295+ DataCopyPad(yLocal, yGm_[progress * static_cast<int64_t>(ubFormer_)], copyParams, padParams);
296+ DataCopyPad(dyLocal, dyGm_[progress * static_cast<int64_t>(ubFormer_)], copyParams, padParams);
297+ 
298+ inQueueY_.EnQue(yLocal);
299+ inQueueDy_.EnQue(dyLocal);
300+}
301+ 
302+__aicore__ inline void KernelAcosGradV2<float>::Compute(uint32_t currentNum)
303+{
304+ LocalTensor<float> yLocal = inQueueY_.template DeQue<float>();
305+ LocalTensor<float> dyLocal = inQueueDy_.template DeQue<float>();
306+ LocalTensor<float> zLocal = outQueueZ_.template AllocTensor<float>();
307+ 
308+ LocalTensor<float> tmpA = tmpQueue1_.template AllocTensor<float>();
309+ LocalTensor<float> tmpB = tmpQueue2_.template AllocTensor<float>();
310+ 
311+ // tmpA = 1 - y^2
312+ Mul(tmpA, yLocal, yLocal, currentNum);
313+ PipeBarrier<PIPE_V>();
314+ 
315+ Muls(tmpA, tmpA, static_cast<float>(-1.0f), currentNum);
316+ PipeBarrier<PIPE_V>();
317+ 
318+ Adds(tmpB, tmpA, static_cast<float>(1.0f), currentNum);
319+ PipeBarrier<PIPE_V>();
320+ 
321+ // tmpB = sqrt(1 - y^2)
322+ Sqrt(tmpB, tmpB, currentNum);
323+ PipeBarrier<PIPE_V>();
324+ 
325+ // z = -dy / sqrt(1 - y^2)
326+ Muls(tmpA, dyLocal, static_cast<float>(-1.0f), currentNum);
327+ PipeBarrier<PIPE_V>();
328+ 
329+ Div(zLocal, tmpA, tmpB, currentNum);
Y
Yyue-ma8月3日

Div指令后需要有PipeBarrier<PIPE_V>();

likedislike
镜花水月1tachi
镜花水月1tachi
8月3日 评论：
330+ PipeBarrier<PIPE_V>();
331+ 
332+ outQueueZ_.template EnQue<float>(zLocal);
atomgit-bot
atomgit-botatomgit-bot7月31日

🟡 Medium Priority

变更文件 op_kernel/arch32/acos_grad_v2.h，KernelAcosGradV2<float>::Compute()（第 302-336 行）：

第 329 行 Div(zLocal, tmpA, tmpB, currentNum) 是最后一条向量计算指令，其结果 zLocal 紧接着在第 331 行通过 outQueueZ_.template EnQue<float>(zLocal) 进入输出队列。但此处缺少 PipeBarrier<PIPE_V>() 保证 Div 结果写入完成后再执行 EnQue。

对比通用模板 KernelAcosGradV2<T>::Compute()（第 166-173 行），Div 之后有 PipeBarrier<PIPE_V>() 才执行后续的 Cast 和 EnQue。设计文档（design.md §5）也明确要求"每条向量指令后 PipeBarrier<PIPE_V>()，保证读写依赖"。

缺少该 barrier 可能导致 Div 尚未完成时 EnQue 已读取 zLocal 的脏数据，造成输出结果错误。当前精度测试通过可能得益于硬件流水线的隐式同步，但依赖未定义行为存在风险。

建议：在 Div 之后、EnQue 之前增加 PipeBarrier<PIPE_V>()，与其他向量指令后的同步模式保持一致

改动建议
332
+ Div(zLocal, tmpA, tmpB, currentNum);
333
+ PipeBarrier<PIPE_V>();
334
+
332
335
  outQueueZ_.template EnQue<float>(zLocal);
应用建议
likedislike
不准确？
333+ inQueueY_.FreeTensor(yLocal);
334+ inQueueDy_.FreeTensor(dyLocal);
335+ tmpQueue1_.FreeTensor(tmpA);
336+ tmpQueue2_.FreeTensor(tmpB);
337+}
338+ 
339+__aicore__ inline void KernelAcosGradV2<float>::CopyOut(int64_t progress, uint32_t currentNum)
340+{
341+ LocalTensor<float> zLocal = outQueueZ_.template DeQue<float>();
342+ 
343+ int64_t gmOffset = progress * static_cast<int64_t>(ubFormer_);
344+ AscendC::DataCopyExtParams copyParams{1, static_cast<uint32_t>(currentNum * sizeof(float)), 0, 0, 0};
345+ DataCopyPad(zGm_[gmOffset], zLocal, copyParams);
346+ 
347+ outQueueZ_.FreeTensor(zLocal);
348+}
349+ 
350+__aicore__ inline void KernelAcosGradV2<float>::ProcessBlock(uint32_t ubLoop, uint32_t ubTail)
351+{
352+ for (uint32_t i = 0; i < ubLoop; i++) {
353+ CopyIn(static_cast<int64_t>(i), ubFormer_);
354+ Compute(ubFormer_);
355+ CopyOut(static_cast<int64_t>(i), ubFormer_);
356+ }
357+ if (ubTail > 0) {
358+ CopyIn(static_cast<int64_t>(ubLoop), ubTail);
359+ Compute(ubTail);
360+ CopyOut(static_cast<int64_t>(ubLoop), ubTail);
361+ }
362+}
363+ 
364+__aicore__ inline void KernelAcosGradV2<float>::Process()
365+{
366+ if (blockLength_ == 0) {
367+ return;
368+ }
369+ 
370+ uint32_t ubLoop = blockLength_ / ubFormer_;
371+ uint32_t ubTail = blockLength_ % ubFormer_;
372+ 
373+ ProcessBlock(ubLoop, ubTail);
374+}
375+ 
376+} // namespace NsAcosGradV2
377+ 
378+#endif // ACOS_GRAD_V2_H
Aexperimental/math/acos_grad_v2/op_kernel/arch32/acos_grad_v2_tiling_data.h+29-0
@@ -0,0 +1,29 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file acos_grad_v2_tiling_data.h
13+ * \brief AcosGradV2 Tiling Data 结构（arch32 / Ascend910B）
14+ */
15+ 
16+#ifndef ACOS_GRAD_V2_TILING_DATA_H
17+#define ACOS_GRAD_V2_TILING_DATA_H
18+ 
19+#include <cstdint>
20+ 
21+struct AcosGradV2TilingData {
22+ uint64_t totalLength;
23+ uint32_t blockFormer;
24+ uint32_t blockNum;
25+ uint32_t ubFormer;
26+ // 核内 loop/tail 由 kernel 侧按 blockLength_ 自行推导，无需在 tiling 预计算
27+};
28+ 
29+#endif // ACOS_GRAD_V2_TILING_DATA_H
Aexperimental/math/acos_grad_v2/op_kernel/arch32/acos_grad_v2_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 acos_grad_v2_tiling_key.h
13+ * \brief AcosGradV2 TilingKey 定义（arch32 / Ascend910B）
14+ *
15+ * 模板参数 D_T 由输入 0 (y) 的 dtype 决定：FP32 / FP16 / BF16。
16+ */
17+ 
18+#ifndef ACOS_GRAD_V2_TILING_KEY_H
19+#define ACOS_GRAD_V2_TILING_KEY_H
20+ 
21+#include "ascendc/host_api/tiling/template_argument.h"
22+ 
23+ASCENDC_TPL_ARGS_DECL(AcosGradV2,
24+ ASCENDC_TPL_DATATYPE_DECL(D_T, C_DT_FLOAT, C_DT_FLOAT16, C_DT_BF16, ASCENDC_TPL_INPUT(0)), );
25+ 
26+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T, C_DT_FLOAT)),
27+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T, C_DT_FLOAT16)),
28+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T, C_DT_BF16)), );
29+ 
30+#endif // ACOS_GRAD_V2_TILING_KEY_H
Aexperimental/math/acos_grad_v2/tests/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# 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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+foreach(SUB_DIR ${CURRENT_DIRS})
13+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14+ add_subdirectory(${SUB_DIR})
15+ endif()
16+endforeach()
Aexperimental/math/acos_grad_v2/tests/eval_precision.cpp+217-0
@@ -0,0 +1,217 @@
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 eval_precision.cpp
13+ * @brief AcosGradV2 精度评测（ascendc-evaluation 方法论：多用例 NPU vs CPU-golden，atol/rtol）
14+ *
15+ * golden = fp64: z = -dy / sqrt(1 - y^2) （|y|>1 → NaN，跳过）
16+ * 容差: FP32 1e-5 / FP16 1e-3 / BF16 1e-2
17+ * 用例: 3 dtype × 11 shape = 33 例（1D-4D、小/中/大、对齐/非对齐）
18+ */
19+#include <iostream>
20+#include <vector>
21+#include <cmath>
22+#include <cstring>
23+#include "acl/acl.h"
24+#include "../op_api/aclnn_acos_grad_v2.h"
25+ 
26+enum class DT { FP32, FP16, BF16 };
27+static int64_t SS(const std::vector<int64_t>& s)
28+{
29+ int64_t n = 1;
30+ for (auto d : s)
31+ n *= d;
32+ return n;
33+}
34+static void Init(int d, aclrtStream* st)
35+{
36+ aclInit(nullptr);
37+ aclrtSetDevice(d);
38+ aclrtCreateStream(st);
39+}
40+ 
41+static void EncF(float v, uint8_t* o) { std::memcpy(o, &v, 4); }
42+static void EncH(float v, uint8_t* o)
43+{
44+ uint32_t x;
45+ std::memcpy(&x, &v, 4);
46+ uint16_t sg = (x >> 31) & 1;
47+ int32_t e = ((x >> 23) & 0xff) - 127 + 15;
48+ uint32_t m = x & 0x7fffff;
49+ uint16_t h = e <= 0 ? (sg << 15) : (e >= 31 ? ((sg << 15) | (0x1f << 10)) : ((sg << 15) | (e << 10) | (m >> 13)));
50+ std::memcpy(o, &h, 2);
51+}
52+static float DecH(const uint8_t* i)
53+{
54+ uint16_t h;
55+ std::memcpy(&h, i, 2);
56+ uint32_t sg = (h >> 15) & 1, e = (h >> 10) & 0x1f, m = h & 0x3ff;
57+ if (e == 0)
58+ return m == 0 ? (sg ? -0.0f : 0.0f) : (sg ? -1 : 1) * (m / 1024.0f / 1024.0f);
59+ if (e == 31)
60+ return m ? NAN : (sg ? -INFINITY : INFINITY);
61+ float v = (1.0f + m / 1024.0f) * std::pow(2.0f, (int)e - 15);
62+ return sg ? -v : v;
63+}
64+static void EncB(float v, uint8_t* o)
65+{
66+ uint32_t x;
67+ std::memcpy(&x, &v, 4);
68+ uint16_t b = (uint16_t)((x + 0x7FFFU + ((x >> 16) & 1)) >> 16);
69+ std::memcpy(o, &b, 2);
70+}
71+static float DecB(const uint8_t* i)
72+{
73+ uint16_t b;
74+ std::memcpy(&b, i, 2);
75+ uint32_t x = (uint32_t)b << 16;
76+ float v;
77+ std::memcpy(&v, &x, 4);
78+ return v;
79+}
80+static void Enc(float v, DT dt, uint8_t* o)
81+{
82+ if (dt == DT::FP32)
83+ EncF(v, o);
84+ else if (dt == DT::FP16)
85+ EncH(v, o);
86+ else
87+ EncB(v, o);
88+}
89+static float Dec(const uint8_t* i, DT dt)
90+{
91+ if (dt == DT::FP32) {
92+ float v;
93+ std::memcpy(&v, i, 4);
94+ return v;
95+ }
96+ return dt == DT::FP16 ? DecH(i) : DecB(i);
97+}
98+static size_t Dsz(DT dt) { return dt == DT::FP32 ? 4 : 2; }
99+static aclDataType AD(DT dt) { return dt == DT::FP32 ? ACL_FLOAT : (dt == DT::FP16 ? ACL_FLOAT16 : ACL_BF16); }
100+static const char* DN(DT dt) { return dt == DT::FP32 ? "FP32" : (dt == DT::FP16 ? "FP16" : "BF16"); }
101+ 
102+int MkT(const std::vector<uint8_t>& b, const std::vector<int64_t>& sh, void** d, aclDataType t, aclTensor** tt)
103+{
104+ aclrtMalloc(d, b.size(), ACL_MEM_MALLOC_HUGE_FIRST);
105+ aclrtMemcpy(*d, b.size(), b.data(), b.size(), ACL_MEMCPY_HOST_TO_DEVICE);
106+ std::vector<int64_t> st(sh.size(), 1);
107+ for (int64_t i = sh.size() - 2; i >= 0; i--)
108+ st[i] = sh[i + 1] * st[i + 1];
109+ *tt = aclCreateTensor(sh.data(), sh.size(), t, st.data(), 0, ACL_FORMAT_ND, sh.data(), sh.size(), *d);
110+ return 0;
111+}
112+ 
113+struct R {
114+ DT dt;
115+ int total;
116+ int pass;
117+ bool ok;
118+ double mx;
119+};
120+static R Run(const std::vector<int64_t>& sh, DT dt, float ylo, float yhi, float atol, float rtol, aclrtStream st)
121+{
122+ int64_t N = SS(sh);
123+ size_t ds = Dsz(dt);
124+ std::vector<uint8_t> yb(N * ds), dyb(N * ds);
125+ for (int64_t i = 0; i < N; i++) {
126+ float yv = ylo + (yhi - ylo) * (float)i / (float)N, dv = -1.0f + 2.0f * (float)i / (float)N;
127+ Enc(yv, dt, yb.data() + i * ds);
128+ Enc(dv, dt, dyb.data() + i * ds);
129+ }
130+ aclTensor* yT = nullptr;
131+ void* yD = nullptr;
132+ MkT(yb, sh, &yD, AD(dt), &yT);
133+ aclTensor* dyT = nullptr;
134+ void* dyD = nullptr;
135+ MkT(dyb, sh, &dyD, AD(dt), &dyT);
136+ std::vector<uint8_t> zb(N * ds, 0);
137+ aclTensor* zT = nullptr;
138+ void* zD = nullptr;
139+ MkT(zb, sh, &zD, AD(dt), &zT);
140+ uint64_t ws = 0;
141+ aclOpExecutor* ex = nullptr;
142+ int ret = aclnnAcosGradV2GetWorkspaceSize(yT, dyT, zT, &ws, &ex);
143+ void* wa = nullptr;
144+ if (ws > 0)
145+ aclrtMalloc(&wa, ws, ACL_MEM_MALLOC_HUGE_FIRST);
146+ if (ret == 0)
147+ ret = aclnnAcosGradV2(wa, ws, ex, st);
148+ aclrtSynchronizeStream(st);
149+ std::vector<uint8_t> rb(N * ds, 0);
150+ aclrtMemcpy(rb.data(), rb.size(), zD, rb.size(), ACL_MEMCPY_DEVICE_TO_HOST);
151+ int pass = 0, dom = 0;
152+ double mx = 0;
153+ for (int64_t i = 0; i < N; i++) {
154+ float yq = Dec(yb.data() + i * ds, dt), dq = Dec(dyb.data() + i * ds, dt), rs = Dec(rb.data() + i * ds, dt);
155+ double omy = 1.0 - (double)yq * (double)yq;
156+ if (omy <= 0)
157+ continue;
158+ dom++;
159+ double g = (double)dq * (-1.0 / std::sqrt(omy));
160+ if (std::isnan(rs) || std::isinf(rs))
161+ continue;
162+ double err = std::fabs((double)rs - g), tol = atol + rtol * std::fabs(g),
163+ re = std::fabs(g) > 0 ? err / std::fabs(g) : 0;
164+ if (re > mx)
165+ mx = re;
166+ if (err <= tol)
167+ pass++;
168+ }
169+ R r;
170+ r.dt = dt;
171+ r.total = dom > 0 ? dom : (int)N;
172+ r.pass = pass;
173+ r.mx = mx;
174+ r.ok = (ret == 0) && (pass == r.total);
175+ aclDestroyTensor(yT);
176+ aclDestroyTensor(dyT);
177+ aclDestroyTensor(zT);
178+ aclrtFree(yD);
179+ aclrtFree(dyD);
180+ aclrtFree(zD);
181+ if (ws > 0)
182+ aclrtFree(wa);
183+ return r;
184+}
185+ 
186+int main()
187+{
188+ int dev = 0;
189+ aclrtStream st;
190+ Init(dev, &st);
191+ std::vector<std::vector<int64_t>> sh = {{8}, {255}, {1024}, {4096}, {17, 31}, {32, 256},
192+ {64, 1024}, {3, 5, 7}, {2, 4, 8, 16}, {1, 3, 1, 5}, {1000, 1000}};
193+ struct T {
194+ DT dt;
195+ float a, r;
196+ } ts[] = {{DT::FP32, 1e-5f, 1e-5f}, {DT::FP16, 1e-3f, 1e-3f}, {DT::BF16, 1e-2f, 1e-2f}};
197+ printf("dtype,elem,max_relerr,atol/rtol,status\n");
198+ int ap = 0, ac = 0;
199+ int pd[3] = {0, 0, 0}, cd[3] = {0, 0, 0};
200+ for (auto& t : ts)
201+ for (auto& s : sh) {
202+ auto r = Run(s, t.dt, -0.99f, 0.99f, t.a, t.r, st);
203+ printf("%s,%d,%.4e,%.0e,%s\n", DN(t.dt), r.total, r.mx, t.a, r.ok ? "PASS" : "FAIL");
204+ cd[(int)t.dt]++;
205+ if (r.ok) {
206+ pd[(int)t.dt]++;
207+ ap++;
208+ }
209+ ac++;
210+ }
211+ printf("\n=== 总例数 %d, 通过 %d ===\nFP32 %d/%d FP16 %d/%d BF16 %d/%d\n结论: %s\n", ac, ap, pd[0], cd[0], pd[1],
212+ cd[1], pd[2], cd[2], ap == ac ? "PRECISION PASS" : "PRECISION FAIL");
213+ aclrtDestroyStream(st);
214+ aclrtResetDevice(dev);
215+ aclFinalize();
216+ return ap == ac ? 0 : 1;
217+}
Aexperimental/math/acos_grad_v2/tests/st/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# 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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+foreach(SUB_DIR ${CURRENT_DIRS})
13+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14+ add_subdirectory(${SUB_DIR})
15+ endif()
16+endforeach()
Aexperimental/math/acos_grad_v2/tests/st/aclnnAcosGradV2/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+# Resource holder for ATK json and executor files.
Aexperimental/math/acos_grad_v2/tests/st/aclnnAcosGradV2/all_aclnn_acos_grad_v2.yaml+45-0
@@ -0,0 +1,45 @@
1+api: pytorch
2+api_type: ascend_acos_grad_v2
3+version: v1.0
4+name: torch.acos.backward
5+aclnn_name: AcosGradV2
6+generate: ascend_aclnn_acos_grad_v2
7+dtype_numbers: 100
8+standard:
9+ acc:
10+ single_bm:
11+ type: high_performance
12+ perf: not_key
13+inputs:
14+ - name: y
15+ type: tensor
16+ required: true
17+ dtypes:
18+ values: [ fp16, fp32, bf16 ]
19+ shapes:
20+ dim_numbers:
21+ values: [ 1, 2, 3, 4 ]
22+ dim_values:
23+ values: [ [ 1, 10 ], [10, 100], [100, 255], 1024, 2048, 4096 ]
24+ weights: [0.8, 0.05, 0.005, 0.002, 0.001, 0.001]
25+ ranges:
26+ valid:
27+ values: [ [-1, 1] ]
28+ invalid:
29+ values: [ [ '-inf' ], [ 'inf' ], [ 'nan' ] ]
30+ - name: dy
31+ type: tensor
32+ required: true
33+ dtypes:
34+ values: [ fp16, fp32, bf16 ]
35+ shapes:
36+ dim_numbers:
37+ values: [ 1, 2, 3, 4 ]
38+ dim_values:
39+ values: [ [ 1, 10 ], [10, 100], [100, 255], 1024, 2048, 4096 ]
40+ weights: [0.8, 0.05, 0.005, 0.002, 0.001, 0.001]
41+ ranges:
42+ valid:
43+ values: [ [-10, 10] ]
44+ invalid:
45+ values: [ [ '-inf' ], [ 'inf' ], [ 'nan' ] ]
Aexperimental/math/acos_grad_v2/tests/st/aclnnAcosGradV2/executor_aclnnAcosGradV2.py+37-0
@@ -0,0 +1,37 @@
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+ 
14+import torch
15+from atk.configs.dataset_config import InputDataset
16+from atk.tasks.api_execute import register
17+from atk.tasks.api_execute.aclnn_base_api import AclnnBaseApi
18+from atk.tasks.api_execute.base_api import BaseApi
19+ 
20+ 
21+def reference(input_data: InputDataset):
22+ """Golden: z = -dy / sqrt(1 - y^2)"""
23+ y = input_data.kwargs["y"]
24+ dy = input_data.kwargs["dy"]
25+ return -dy / torch.sqrt(1 - y * y)
26+ 
27+ 
28+@register("ascend_acos_grad_v2")
29+class TorchAcosGradV2(BaseApi):
30+ def __call__(self, input_data: InputDataset, with_output: bool = False):
31+ return reference(input_data)
32+ 
33+ 
34+@register("aclnn_acos_grad_v2")
35+class AclnnAcosGradV2(AclnnBaseApi):
36+ def __call__(self):
37+ super().__call__()
Aexperimental/math/acos_grad_v2/tests/st/aclnnAcosGradV2/generator_aclnn_acos_grad_v2.py+25-0
@@ -0,0 +1,25 @@
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+ 
14+from atk.case_generator.generator.base_generator import CaseGenerator
15+from atk.case_generator.generator.generate_types import GENERATOR_REGISTRY
16+from atk.configs.case_config import CaseConfig
17+ 
18+ 
19+@GENERATOR_REGISTRY.register("ascend_aclnn_acos_grad_v2")
20+class AcosGradV2Generator(CaseGenerator):
21+ def after_case_config(self, case_config: CaseConfig) -> CaseConfig:
22+ # z = -dy / sqrt(1 - y^2): y and dy must share the same shape and dtype
23+ case_config.inputs[1].dtype = case_config.inputs[0].dtype
24+ case_config.inputs[1].shape = case_config.inputs[0].shape
25+ return case_config
Aexperimental/math/acos_grad_v2/tests/ut/op_api/test_aclnn_acos_grad_v2.cpp+95-0
@@ -0,0 +1,95 @@
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+#include "gtest/gtest.h"
12+#include "../../../op_api/aclnn_acos_grad_v2.h"
13+#include "op_api_ut_common/op_api_ut.h"
14+#include "op_api_ut_common/tensor_desc.h"
15+ 
16+using namespace op;
17+ 
18+class L2AcosGradV2Test : public testing::Test {};
19+ 
20+// 正常用例：FP32 同 shape（精度标杆 atol/rtol = 1e-4，与 acos 保持一致）
21+TEST_F(L2AcosGradV2Test, fp32_same_shape)
22+{
23+ auto y = TensorDesc({10, 10}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
24+ auto dy = TensorDesc({10, 10}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-10, 10);
25+ auto z = TensorDesc({10, 10}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
26+ auto ut = OP_API_UT(aclnnAcosGradV2, INPUT(y, dy), OUTPUT(z));
27+ uint64_t workspaceSize = 0;
28+ EXPECT_EQ(ut.TestGetWorkspaceSize(&workspaceSize), ACL_SUCCESS);
29+}
30+ 
31+// 正常用例：FP16
32+TEST_F(L2AcosGradV2Test, fp16_same_shape)
33+{
34+ auto y = TensorDesc({2, 3, 5}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-1, 1);
35+ auto dy = TensorDesc({2, 3, 5}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-10, 10);
36+ auto z = TensorDesc({2, 3, 5}, ACL_FLOAT16, ACL_FORMAT_ND).Precision(0.001, 0.001);
37+ auto ut = OP_API_UT(aclnnAcosGradV2, INPUT(y, dy), OUTPUT(z));
38+ uint64_t workspaceSize = 0;
39+ EXPECT_EQ(ut.TestGetWorkspaceSize(&workspaceSize), ACL_SUCCESS);
40+}
41+ 
42+// 正常用例：BF16
43+TEST_F(L2AcosGradV2Test, bf16_same_shape)
44+{
45+ auto y = TensorDesc({4, 8, 16}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-1, 1);
46+ auto dy = TensorDesc({4, 8, 16}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-10, 10);
47+ auto z = TensorDesc({4, 8, 16}, ACL_BF16, ACL_FORMAT_ND).Precision(0.004, 0.004);
48+ auto ut = OP_API_UT(aclnnAcosGradV2, INPUT(y, dy), OUTPUT(z));
49+ uint64_t workspaceSize = 0;
50+ EXPECT_EQ(ut.TestGetWorkspaceSize(&workspaceSize), ACL_SUCCESS);
51+}
52+ 
53+// 异常用例：dtype 不一致（y 为 FP32，dy 为 FP16）
54+TEST_F(L2AcosGradV2Test, invalid_dtype_mismatch)
55+{
56+ auto y = TensorDesc({10, 10}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
57+ auto dy = TensorDesc({10, 10}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-10, 10);
58+ auto z = TensorDesc({10, 10}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
59+ auto ut = OP_API_UT(aclnnAcosGradV2, INPUT(y, dy), OUTPUT(z));
60+ uint64_t workspaceSize = 0;
61+ EXPECT_EQ(ut.TestGetWorkspaceSize(&workspaceSize), ACLNN_ERR_PARAM_INVALID);
62+}
63+ 
64+// 异常用例：shape 不一致
65+TEST_F(L2AcosGradV2Test, invalid_shape_mismatch)
66+{
67+ auto y = TensorDesc({10, 10}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
68+ auto dy = TensorDesc({20, 5}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-10, 10);
69+ auto z = TensorDesc({10, 10}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
70+ auto ut = OP_API_UT(aclnnAcosGradV2, INPUT(y, dy), OUTPUT(z));
71+ uint64_t workspaceSize = 0;
72+ EXPECT_EQ(ut.TestGetWorkspaceSize(&workspaceSize), ACLNN_ERR_PARAM_INVALID);
73+}
74+ 
75+// 异常用例：不支持 INT32
76+TEST_F(L2AcosGradV2Test, invalid_dtype_int32)
77+{
78+ auto y = TensorDesc({10, 10}, ACL_INT32, ACL_FORMAT_ND).ValueRange(-1, 1);
79+ auto dy = TensorDesc({10, 10}, ACL_INT32, ACL_FORMAT_ND).ValueRange(-10, 10);
80+ auto z = TensorDesc({10, 10}, ACL_INT32, ACL_FORMAT_ND).Precision(0, 0);
81+ auto ut = OP_API_UT(aclnnAcosGradV2, INPUT(y, dy), OUTPUT(z));
82+ uint64_t workspaceSize = 0;
83+ EXPECT_NE(ut.TestGetWorkspaceSize(&workspaceSize), ACL_SUCCESS);
84+}
85+ 
86+// 正常用例：大 shape 多核
87+TEST_F(L2AcosGradV2Test, fp32_large_shape)
88+{
89+ auto y = TensorDesc({1024, 1024}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
90+ auto dy = TensorDesc({1024, 1024}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-10, 10);
91+ auto z = TensorDesc({1024, 1024}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
92+ auto ut = OP_API_UT(aclnnAcosGradV2, INPUT(y, dy), OUTPUT(z));
93+ uint64_t workspaceSize = 0;
94+ EXPECT_EQ(ut.TestGetWorkspaceSize(&workspaceSize), ACL_SUCCESS);
95+}
Aexperimental/math/acos_grad_v2/tests/ut/op_host/CMakeLists.txt+14-0
@@ -0,0 +1,14 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# 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+if(UT_TEST_ALL OR OP_HOST_UT)
12+ add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
13+ add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+endif()
Aexperimental/math/acos_grad_v2/tests/ut/op_host/test_acos_grad_v2_infershape.cpp+60-0
@@ -0,0 +1,60 @@
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+#include <gtest/gtest.h>
12+#include <iostream>
13+#include <vector>
14+#include "infershape_context_faker.h"
15+#include "infershape_case_executor.h"
16+ 
17+// AcosGradV2: 输出 z 的 shape = 输入 y 的 shape
18+class AcosGradV2Infershape : public testing::Test {};
19+ 
20+TEST_F(AcosGradV2Infershape, fp32_same_shape)
21+{
22+ gert::InfershapeContextPara para("AcosGradV2",
23+ {
24+ {{{4, 3, 4}, {4, 3, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
25+ {{{4, 3, 4}, {4, 3, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
26+ },
27+ {
28+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
29+ });
30+ std::vector<std::vector<int64_t>> expectOutputShape = {{4, 3, 4}};
31+ ExecuteTestCase(para, ge::GRAPH_SUCCESS, expectOutputShape);
32+}
33+ 
34+TEST_F(AcosGradV2Infershape, fp16_one_dim)
35+{
36+ gert::InfershapeContextPara para("AcosGradV2",
37+ {
38+ {{{1024}, {1024}}, ge::DT_FLOAT16, ge::FORMAT_ND},
39+ {{{1024}, {1024}}, ge::DT_FLOAT16, ge::FORMAT_ND},
40+ },
41+ {
42+ {{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND},
43+ });
44+ std::vector<std::vector<int64_t>> expectOutputShape = {{1024}};
45+ ExecuteTestCase(para, ge::GRAPH_SUCCESS, expectOutputShape);
46+}
47+ 
48+TEST_F(AcosGradV2Infershape, bf16_two_dim)
49+{
50+ gert::InfershapeContextPara para("AcosGradV2",
51+ {
52+ {{{32, 256}, {32, 256}}, ge::DT_BF16, ge::FORMAT_ND},
53+ {{{32, 256}, {32, 256}}, ge::DT_BF16, ge::FORMAT_ND},
54+ },
55+ {
56+ {{{}, {}}, ge::DT_BF16, ge::FORMAT_ND},
57+ });
58+ std::vector<std::vector<int64_t>> expectOutputShape = {{32, 256}};
59+ ExecuteTestCase(para, ge::GRAPH_SUCCESS, expectOutputShape);
60+}
Aexperimental/math/acos_grad_v2/tests/ut/op_host/test_acos_grad_v2_tiling.cpp+79-0
@@ -0,0 +1,79 @@
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+#include <gtest/gtest.h>
12+#include "tiling_case_executor.h"
13+#include "../../../op_kernel/arch32/acos_grad_v2_tiling_data.h"
14+ 
15+namespace optiling {
16+struct AcosGradV2CompileInfo {
17+ int32_t totalCoreNum = 0;
18+ int64_t ubSize = 0;
19+ bool isRegbase = false;
20+};
21+} // namespace optiling
22+ 
23+namespace NsAcosGradV2 {
24+bool operator==(const AcosGradV2TilingData& a, const AcosGradV2TilingData& b)
25+{
26+ return a.totalLength == b.totalLength && a.blockFormer == b.blockFormer && a.blockNum == b.blockNum &&
27+ a.ubFormer == b.ubFormer;
28+}
29+} // namespace NsAcosGradV2
30+ 
31+namespace {
32+constexpr size_t WORKSPACE_SIZE = 0;
33+}
34+ 
35+class AcosGradV2Tiling : public testing::Test {};
36+ 
37+// FP32 8192 元素，20 核：tiling 应成功，totalLength=8192，blockNum≥1，blockFormer 按 512 对齐
38+TEST_F(AcosGradV2Tiling, fp32_8192)
39+{
40+ optiling::AcosGradV2CompileInfo compileInfo{20, 192 * 1024, false};
41+ gert::TilingContextPara para("AcosGradV2",
42+ {
43+ {{{8192}, {8192}}, ge::DT_FLOAT, ge::FORMAT_ND},
44+ {{{8192}, {8192}}, ge::DT_FLOAT, ge::FORMAT_ND},
45+ },
46+ {
47+ {{{8192}, {8192}}, ge::DT_FLOAT, ge::FORMAT_ND},
48+ },
49+ {}, &compileInfo);
50+ 
51+ TilingInfo tilingInfo;
52+ ASSERT_TRUE(ExecuteTiling(para, tilingInfo));
53+ ASSERT_EQ(tilingInfo.tilingDataSize, sizeof(AcosGradV2TilingData));
54+ auto* td = reinterpret_cast<AcosGradV2TilingData*>(tilingInfo.tilingData.get());
55+ EXPECT_EQ(td->totalLength, 8192ULL);
56+ EXPECT_GE(td->blockNum, 1U);
57+ EXPECT_GE(td->blockFormer, 512U); // ELEM_ALIGN
58+}
59+ 
60+// FP16 多维：tiling 应成功
61+TEST_F(AcosGradV2Tiling, fp16_multi_dim)
62+{
63+ optiling::AcosGradV2CompileInfo compileInfo{20, 192 * 1024, false};
64+ gert::TilingContextPara para("AcosGradV2",
65+ {
66+ {{{2, 3, 5}, {2, 3, 5}}, ge::DT_FLOAT16, ge::FORMAT_ND},
67+ {{{2, 3, 5}, {2, 3, 5}}, ge::DT_FLOAT16, ge::FORMAT_ND},
68+ },
69+ {
70+ {{{2, 3, 5}, {2, 3, 5}}, ge::DT_FLOAT16, ge::FORMAT_ND},
71+ },
72+ {}, &compileInfo);
73+ 
74+ TilingInfo tilingInfo;
75+ ASSERT_TRUE(ExecuteTiling(para, tilingInfo));
76+ ASSERT_EQ(tilingInfo.tilingDataSize, sizeof(AcosGradV2TilingData));
77+ auto* td = reinterpret_cast<AcosGradV2TilingData*>(tilingInfo.tilingData.get());
78+ EXPECT_EQ(td->totalLength, 30ULL); // 2*3*5
79+}