Star1317
1435
代码
代码
Issues122
Pull Requests252
Actions
项目讨论
Wiki
分析
Star1317
1435
已合并
支持Div算子在下一代实现 #654
yuhwang创建于 1月4日
支持Div算子在下一代实现 #654
已合并
yuhwang创建于 1月4日
从yuhwang/ops-mathmaster合入到cann/ops-mathmaster
共 36 个文件变更+2652-357
Mmath/div/CMakeLists.txt+13-16
@@ -1,20 +1,17 @@
1-# ----------------------------------------------------------------------------1+# -----------------------------------------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").4+# CANN Open Software License Agreement Version 2.0 (the "License").
6# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
10-# ----------------------------------------------------------------------------9+# -----------------------------------------------------------------------------------------------------------
11 10 
12-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)11+set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend910_95")
13-if(NOT ENABLE_TEST AND NOT BENCHMARK)12+set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35")
14- list(REMOVE_ITEM CURRENT_DIRS tests)13+ 
15-endif()14+add_all_modules_sources(OPTYPE div ACLNNTYPE aclnn_exclude
16-foreach(SUB_DIR ${CURRENT_DIRS})15+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
17- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")16+ TILING_DIR ${SUPPORT_TILING_DIR}
18- add_subdirectory(${SUB_DIR})17+ DISABLE_IN_OPP TRUE)
19- endif()
20-endforeach()
Mmath/div/README.md+69-1
@@ -1,3 +1,71 @@
1# Div1# Div
2 2 
3-本目录仅包含Div算子对应的aclnn接口；如您想要贡献该算子的AscendC实现，请参考[贡献流程](../../CONTRIBUTING.md)。3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8+ 
9+## 功能说明
10+ 
11+- 接口功能：完成除法计算
12+ 
13+- 计算公式：
14+ 
15+ $$
16+ out_i = \frac{input_i}{other_i}
17+ $$
18+ 
19+- 例外说明：当涉及Complex运算时，且分母的模为0时，仍然采用通用的复数计算公式，不做特殊处理，此处表现和CPU/GPU存在差异。
20+ 
21+## 参数说明
22+ 
23+<table style="undefined;table-layout: fixed; width: 1005px"><colgroup>
24+ <col style="width: 170px">
25+ <col style="width: 170px">
26+ <col style="width: 352px">
27+ <col style="width: 213px">
28+ <col style="width: 100px">
29+ </colgroup>
30+ <thead>
31+ <tr>
32+ <th>参数名</th>
33+ <th>输入/输出/属性</th>
34+ <th>描述</th>
35+ <th>数据类型</th>
36+ <th>数据格式</th>
37+ </tr></thead>
38+ <tbody>
39+ <tr>
40+ <td>x1</td>
41+ <td>输入</td>
42+ <td>公式中的input。</td>
43+ <td>BFLOAT16、FLOAT16、FLOAT、INT32、UINT8、INT8、COMPLEX32、COMPLEX64</td>
44+ <td>ND</td>
45+ </tr>
46+ <tr>
47+ <td>x2</td>
48+ <td>输入</td>
49+ <td>公式中的other。</td>
50+ <td>BFLOAT16、FLOAT16、FLOAT、INT32、UINT8、INT8、COMPLEX32、COMPLEX64</td>
51+ <td>-</td>
52+ </tr>
53+ <tr>
54+ <td>y</td>
55+ <td>输出</td>
56+ <td>公式中的out。</td>
57+ <td>BFLOAT16、FLOAT16、FLOAT、INT32、UINT8、INT8、COMPLEX32、COMPLEX64</td>
58+ <td>ND</td>
59+ </tr>
60+ </tbody></table>
61+ 
62+## 约束说明
63+ 
64+无
65+ 
66+ 
67+## 调用说明
68+ 
69+| 调用方式 | 样例代码 | 说明 |
70+| ---------------- | --------------------------- | --------------------------------------------------- |
71+| aclnn接口 | [test_aclnn_div](examples/test_aclnn_div.cpp) | 通过[AclnnDiv](docs/aclnnDiv&aclnnInplaceDiv.md)接口方式调用Div算子。 |
Mmath/div/docs/aclnnDiv&aclnnInplaceDiv.md+13-3
@@ -4,8 +4,9 @@
4 4 
5| 产品 | 是否支持 |5| 产品 | 是否支持 |
6| :----------------------------------------------------------- | :------: |6| :----------------------------------------------------------- | :------: |
7-| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8-| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
9 10 
10## 功能说明11## 功能说明
11 12 
@@ -138,6 +139,10 @@ aclnnStatus aclnnInplaceDiv(
138 </tbody>139 </tbody>
139 </table>140 </table>
140 141 
142+ - <term>Ascend 950PR/Ascend 950DT</term>：
143+ - self与other推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT。
144+ - out不支持INT32、INT64、INT16、INT8、UINT8、BOOL数据类型。
145+ 
141- **返回值：**146- **返回值：**
142 147 
143 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。148 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
@@ -293,6 +298,11 @@ aclnnStatus aclnnInplaceDiv(
293 </tbody>298 </tbody>
294 </table>299 </table>
295 300 
301+ - <term>Ascend 950PR/Ascend 950DT</term>：
302+ - selfRef与other推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT。
303+ - selfRef不支持INT32、INT64、INT16、INT8、UINT8、BOOL。
304+ 
305+ 
296- **返回值：**306- **返回值：**
297 307 
298 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。308 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
@@ -556,4 +566,4 @@ int main() {
556 566 
557 return 0;567 return 0;
558}568}
559-```569+```
Mmath/div/docs/aclnnDivMod&aclnnInplaceDivMod.md+9-3
@@ -4,8 +4,9 @@
4 4 
5| 产品 | 是否支持 |5| 产品 | 是否支持 |
6| :----------------------------------------------------------- | :------: |6| :----------------------------------------------------------- | :------: |
7-| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8-| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
9 10 
10## 功能说明11## 功能说明
11 12 
@@ -34,12 +35,15 @@ $$
34 35
35 - self(aclTensor*, 计算输入)：表示被除数，公式中的输入`self`，Device侧的aclTensor，数据类型与other的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)），shape需要与other满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。36 - self(aclTensor*, 计算输入)：表示被除数，公式中的输入`self`，Device侧的aclTensor，数据类型与other的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)），shape需要与other满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
36 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16。37 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16。
38+ - <term>Ascend 950PR/Ascend 950DT</term>：数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16。当mode为0且self与other推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT。
37 - other(aclTensor*, 计算输入)：表示除数，公式中的输入`other`，Device侧的aclTensor，数据类型与self的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)），shape需要与self满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。39 - other(aclTensor*, 计算输入)：表示除数，公式中的输入`other`，Device侧的aclTensor，数据类型与self的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)），shape需要与self满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
38 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16。40 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16。
41+ - <term>Ascend 950PR/Ascend 950DT</term>：数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16。当mode为0且self与other推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT。
39 - mode(int, 计算输入)：表示对商的舍入模式的选择，Host侧的整型值，数据类型支持int整型，公式中的输入`mode`，枚举值如下：<br>0-对应None：默认不执行舍入。<br>1-对应trunc：将除法的小数部分舍入为零。<br>2-对应floor：向下舍入除法的结果。42 - mode(int, 计算输入)：表示对商的舍入模式的选择，Host侧的整型值，数据类型支持int整型，公式中的输入`mode`，枚举值如下：<br>0-对应None：默认不执行舍入。<br>1-对应trunc：将除法的小数部分舍入为零。<br>2-对应floor：向下舍入除法的结果。
40 43 
41 - out(aclTensor\*, 计算输出)：表示商，公式中的`out`，Device侧的aclTensor，数据类型需要是self与other推导之后可转换的数据类型，shape需要是self与other broadcast之后的shape。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。44 - out(aclTensor\*, 计算输出)：表示商，公式中的`out`，Device侧的aclTensor，数据类型需要是self与other推导之后可转换的数据类型，shape需要是self与other broadcast之后的shape。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
42 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16。45 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16。
46+ - <term>Ascend 950PR/Ascend 950DT</term>：mode为0时，数据类型支持FLOAT、FLOAT16、DOUBLE、BFLOAT16、COMPLEX128、COMPLEX64。mode为1或2时，数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、BFLOAT16、COMPLEX128、COMPLEX64。
43 - workspaceSize(uint64_t\*, 出参)：返回需要在Device侧申请的workspace大小。47 - workspaceSize(uint64_t\*, 出参)：返回需要在Device侧申请的workspace大小。
44 48 
45 - executor(aclOpExecutor\**, 出参)：返回op执行器，包含了算子计算流程。49 - executor(aclOpExecutor\**, 出参)：返回op执行器，包含了算子计算流程。
@@ -79,8 +83,10 @@ $$
79 83
80 - selfRef(aclTensor*, 计算输入|计算输出)：表示被除数和商，公式中的输入`input`和`out`，Device侧的aclTensor，数据类型与other的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)），数据类型需要是selfRef与other推导之后可转换的数据类型（参见[互转换关系](../../../docs/zh/context/互转换关系.md)），shape需要与other满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。84 - selfRef(aclTensor*, 计算输入|计算输出)：表示被除数和商，公式中的输入`input`和`out`，Device侧的aclTensor，数据类型与other的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)），数据类型需要是selfRef与other推导之后可转换的数据类型（参见[互转换关系](../../../docs/zh/context/互转换关系.md)），shape需要与other满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
81 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16。85 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16。
86+ - <term>Ascend 950PR/Ascend 950DT</term>：mode为0时，数据类型支持FLOAT、FLOAT16、DOUBLE、BFLOAT16，selfRef与other推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT。mode为1或2时，数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、BFLOAT16。
82 - other(aclTensor\*, 计算输入)：表示除数，公式中的输入`other`，Device侧的aclTensor，数据类型与selfRef的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)），shape需要与selfRef满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。87 - other(aclTensor\*, 计算输入)：表示除数，公式中的输入`other`，Device侧的aclTensor，数据类型与selfRef的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)），shape需要与selfRef满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
83 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、BFLOAT16。88 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、BFLOAT16。
89+ - <term>Ascend 950PR/Ascend 950DT器</term>：数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、BFLOAT16。当mode为0且selfRef与other推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT。
84 - mode(int, 计算输入)：表示对商的舍入模式的选择，Host侧的整型值，数据类型支持int整型，公式中的输入`mode`，枚举值如下：<br>0-对应None：默认不执行舍入。<br>1-对应trunc：将除法的小数部分舍入为零。<br>2-对应floor：向下舍入除法的结果。90 - mode(int, 计算输入)：表示对商的舍入模式的选择，Host侧的整型值，数据类型支持int整型，公式中的输入`mode`，枚举值如下：<br>0-对应None：默认不执行舍入。<br>1-对应trunc：将除法的小数部分舍入为零。<br>2-对应floor：向下舍入除法的结果。
85 91 
86 - workspaceSize(uint64_t\*，出参)：返回需要在Device侧申请的workspace大小。92 - workspaceSize(uint64_t\*，出参)：返回需要在Device侧申请的workspace大小。
@@ -296,4 +302,4 @@ int main() {
296 302 
297 return 0;303 return 0;
298}304}
299-```305+```
Mmath/div/docs/aclnnDivMods&aclnnInplaceDivMods.md+9-3
@@ -4,8 +4,9 @@
4 4 
5| 产品 | 是否支持 |5| 产品 | 是否支持 |
6| :----------------------------------------------------------- | :------: |6| :----------------------------------------------------------- | :------: |
7-| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8-| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
9 10 
10## 功能说明11## 功能说明
11 12 
@@ -33,11 +34,14 @@ $$
33 34
34 * self(aclTensor*, 计算输入)：表示被除数，公式中的input，Device侧的aclTensor。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。35 * self(aclTensor*, 计算输入)：表示被除数，公式中的input，Device侧的aclTensor。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
35 * <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16，且数据类型与other的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)）。36 * <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16，且数据类型与other的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)）。
37+ * <term>Ascend 950PR/Ascend 950DT</term>：FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16，且数据类型与other的数据类型需满足[TensorScalar互推导关系](../../../docs/zh/context/TensorScalar互推导关系.md)，推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT。
36 * other(aclScalar*, 计算输入)：表示除数，公式中的输入`other`，Device侧的aclScalar。38 * other(aclScalar*, 计算输入)：表示除数，公式中的输入`other`，Device侧的aclScalar。
37 * <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16，且数据类型与self的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)）。39 * <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16，且数据类型与self的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)）。
40+ * <term>Ascend 950PR/Ascend 950DT</term>：FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16，且数据类型与self的数据类型需满足[TensorScalar互推导关系](../../../docs/zh/context/TensorScalar互推导关系.md)，推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT。
38 * mode(int, 计算输入)：表示对商的舍入模式的选择，公式中的输入`mode`，Host侧的整型值，数据类型支持int整型，枚举值如下：<br>0-对应None：默认不执行舍入。<br>1-对应trunc：将除法的小数部分舍入为零。<br>2-对应floor：向下舍入除法的结果。 41 * mode(int, 计算输入)：表示对商的舍入模式的选择，公式中的输入`mode`，Host侧的整型值，数据类型支持int整型，枚举值如下：<br>0-对应None：默认不执行舍入。<br>1-对应trunc：将除法的小数部分舍入为零。<br>2-对应floor：向下舍入除法的结果。
39 * out(aclTensor\*, 计算输出)：表示商，公式中的`out`，Device侧的aclTensor，且数据类型需要是self与other推导之后可转换的数据类型，shape与self相同。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。42 * out(aclTensor\*, 计算输出)：表示商，公式中的`out`，Device侧的aclTensor，且数据类型需要是self与other推导之后可转换的数据类型，shape与self相同。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
40 * <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT1643 * <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16
44+ * <term>Ascend 950PR/Ascend 950DT</term>：mode为0时，支持FLOAT、FLOAT16、DOUBLE、BFLOAT16、COMPLEX128、COMPLEX64。mode为1或2时，支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、BFLOAT16、COMPLEX128、COMPLEX64。
41 * workspaceSize(uint64_t\*, 出参)：返回需要在Device侧申请的workspace大小。45 * workspaceSize(uint64_t\*, 出参)：返回需要在Device侧申请的workspace大小。
42 * executor(aclOpExecutor\**, 出参)：返回op执行器，包含了算子计算流程。46 * executor(aclOpExecutor\**, 出参)：返回op执行器，包含了算子计算流程。
43 47
@@ -75,8 +79,10 @@ $$
75 79
76 * selfRef(aclTensor\*, 计算输入|计算输出)：表示被除数和商，公式中的输入`input`和`out`，Device侧的aclTensor，数据类型需要是selfRef与other推导之后可转换的数据类型（参见[互转换关系](../../../docs/zh/context/互转换关系.md)）。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。80 * selfRef(aclTensor\*, 计算输入|计算输出)：表示被除数和商，公式中的输入`input`和`out`，Device侧的aclTensor，数据类型需要是selfRef与other推导之后可转换的数据类型（参见[互转换关系](../../../docs/zh/context/互转换关系.md)）。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
77 * <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16，且数据类型与other的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)）。81 * <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16，且数据类型与other的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)）。
82+ * <term>Ascend 950PR/Ascend 950DT</term>：mode为0时，支持FLOAT、FLOAT16、DOUBLE、BFLOAT16。mode为1或2时，支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、BFLOAT16。数据类型与other的数据类型需满足[TensorScalar互推导关系](../../../docs/zh/context/TensorScalar互推导关系.md)，当mode为0时，selfRef与other推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT。
78 * other(aclScalar*, 计算输入)：公式中的输入`other`，Device侧的aclScalar。83 * other(aclScalar*, 计算输入)：公式中的输入`other`，Device侧的aclScalar。
79 * <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16，且数据类型与selfRef的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)）。84 * <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16，且数据类型与selfRef的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)）。
85+ * <term>Ascend 950PR/Ascend 950DT</term>：FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、BFLOAT16，且数据类型与selfRef的数据类型需满足[TensorScalar互推导关系](../../../docs/zh/context/TensorScalar互推导关系.md)，当mode为0时，selfRef与other推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT。
80 * mode(int, 计算输入)：表示对商的舍入模式的选择，公式中的输入`mode`，Host侧的整型值，数据类型支持int整型，枚举值如下：<br>0-对应None：默认不执行舍入。<br>1-对应trunc：将除法的小数部分舍入为零。<br>2-对应floor：向下舍入除法的结果。86 * mode(int, 计算输入)：表示对商的舍入模式的选择，公式中的输入`mode`，Host侧的整型值，数据类型支持int整型，枚举值如下：<br>0-对应None：默认不执行舍入。<br>1-对应trunc：将除法的小数部分舍入为零。<br>2-对应floor：向下舍入除法的结果。
81 * workspaceSize(uint64_t\*, 出参)：返回需要在Device侧申请的workspace大小。87 * workspaceSize(uint64_t\*, 出参)：返回需要在Device侧申请的workspace大小。
82 * executor(aclOpExecutor\**, 出参)：返回op执行器，包含了算子计算流程。88 * executor(aclOpExecutor\**, 出参)：返回op执行器，包含了算子计算流程。
@@ -283,4 +289,4 @@ int main() {
283 aclFinalize();289 aclFinalize();
284 return 0;290 return 0;
285}291}
286-```292+```
Mmath/div/docs/aclnnDivs&aclnnInplaceDivs.md+10-3
@@ -4,8 +4,9 @@
4 4 
5| 产品 | 是否支持 |5| 产品 | 是否支持 |
6| :----------------------------------------------------------- | :------: |6| :----------------------------------------------------------- | :------: |
7-| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8-| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
9 10 
10## 功能说明11## 功能说明
11 12 
@@ -139,6 +140,9 @@ aclnnStatus aclnnInplaceDivs(
139 140 
140 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：141 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：
141 - self数据类型与other的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)）。142 - self数据类型与other的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)）。
143+ - <term>Ascend 950PR/Ascend 950DT</term>：
144+ - out数据类型不支持INT32、INT64、INT16、INT8、UINT8、BOOL。
145+ - self数据类型与other的数据类型需满足[TensorScalar互推导关系](../../../docs/zh/context/TensorScalar互推导关系.md)，推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT。
142- **返回值：**146- **返回值：**
143 147 
144 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。148 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
@@ -293,6 +297,9 @@ aclnnStatus aclnnInplaceDivs(
293 297 
294 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：298 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：
295 - self数据类型与other的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)）。299 - self数据类型与other的数据类型需满足数据类型推导规则（参见[互推导关系](../../../docs/zh/context/互推导关系.md)）。
300+ - <term>Ascend 950PR/Ascend 950DT</term>：
301+ - out数据类型不支持INT32、INT64、INT16、INT8、UINT8、BOOL。
302+ - self数据类型与other的数据类型需满足[TensorScalar互推导关系](../../../docs/zh/context/TensorScalar互推导关系.md)，推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT。
296- **返回值：**303- **返回值：**
297 304 
298 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。305 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
@@ -551,4 +558,4 @@ int main() {
551 558 
552 return 0;559 return 0;
553}560}
554-```561+```
Mmath/div/examples/test_aclnn_div.cpp+155-147
@@ -1,12 +1,12 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3+ * 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").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.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,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.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.8+ * See LICENSE in the root of the software repository for the full text of the License.
9- */9+ */
10 10 
11#include <iostream>11#include <iostream>
12#include <vector>12#include <vector>
@@ -14,167 +14,175 @@
14#include "aclnnop/aclnn_div.h"14#include "aclnnop/aclnn_div.h"
15 15 
16#define CHECK_RET(cond, return_expr) \16#define CHECK_RET(cond, return_expr) \
17- do { \17+ do { \
18- if (!(cond)) { \18+ if (!(cond)) { \
19- return_expr; \19+ return_expr; \
20- } \20+ } \
21- } while (0)21+ } while (0)
22 22 
23-#define LOG_PRINT(message, ...) \23+#define LOG_PRINT(message, ...) \
24- do { \24+ do { \
25- printf(message, ##__VA_ARGS__); \25+ printf(message, ##__VA_ARGS__); \
26- } while (0)26+ } while (0)
27 27 
28-int64_t GetShapeSize(const std::vector<int64_t>& shape) {28+int64_t GetShapeSize(const std::vector<int64_t>& shape)
29- int64_t shapeSize = 1;29+{
30- for (auto i : shape) {30+ int64_t shapeSize = 1;
31- shapeSize *= i;31+ for (auto i : shape) {
32- }32+ shapeSize *= i;
33- return shapeSize;33+ }
34+ return shapeSize;
34}35}
35 36 
36-int Init(int32_t deviceId, aclrtStream* stream) {37+int Init(int32_t deviceId, aclrtStream* stream)
37- // 固定写法，资源初始化38+{
38- auto ret = aclInit(nullptr);39+ // 固定写法，资源初始化
39- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);40+ auto ret = aclInit(nullptr);
40- ret = aclrtSetDevice(deviceId);41+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
41- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);42+ ret = aclrtSetDevice(deviceId);
42- ret = aclrtCreateStream(stream);43+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
43- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);44+ ret = aclrtCreateStream(stream);
44- return 0;45+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
46+ return 0;
45}47}
46 48 
47template <typename T>49template <typename T>
48-int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,50+int CreateAclTensor(
49- aclDataType dataType, aclTensor** tensor) {51+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
50- auto size = GetShapeSize(shape) * sizeof(T);52+ aclTensor** tensor)
51- // 调用aclrtMalloc申请device侧内存53+{
52- auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);54+ auto size = GetShapeSize(shape) * sizeof(T);
53- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);55+ // 调用aclrtMalloc申请device侧内存
54- // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上56+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
55- ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);57+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
56- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);58+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
59+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
60+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
57 61 
58- // 计算连续tensor的strides62+ // 计算连续tensor的strides
59- std::vector<int64_t> strides(shape.size(), 1);63+ std::vector<int64_t> strides(shape.size(), 1);
60- for (int64_t i = shape.size() - 2; i >= 0; i--) {64+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
61- strides[i] = shape[i + 1] * strides[i + 1];65+ strides[i] = shape[i + 1] * strides[i + 1];
62- }66+ }
63 67 
64- // 调用aclCreateTensor接口创建aclTensor68+ // 调用aclCreateTensor接口创建aclTensor
65- *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,69+ *tensor = aclCreateTensor(
66- shape.data(), shape.size(), *deviceAddr);70+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
67- return 0;71+ *deviceAddr);
72+ return 0;
68}73}
69 74 
70-int main() {75+int main()
71- // 1. （固定写法）device/stream初始化，参考acl API手册76+{
72- // 根据自己的实际device填写deviceId77+ // 1. （固定写法）device/stream初始化，参考acl API手册
73- int32_t deviceId = 0;78+ // 根据自己的实际device填写deviceId
74- aclrtStream stream;79+ int32_t deviceId = 0;
75- auto ret = Init(deviceId, &stream);80+ aclrtStream stream;
76- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);81+ auto ret = Init(deviceId, &stream);
82+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
77 83 
78- // 2. 构造输入与输出，需要根据API的接口自定义构造84+ // 2. 构造输入与输出，需要根据API的接口自定义构造
79- std::vector<int64_t> selfShape = {4, 2};85+ std::vector<int64_t> selfShape = {4, 2};
80- std::vector<int64_t> otherShape = {4, 2};86+ std::vector<int64_t> otherShape = {4, 2};
81- std::vector<int64_t> outShape = {4, 2};87+ std::vector<int64_t> outShape = {4, 2};
82- void* selfDeviceAddr = nullptr;88+ void* selfDeviceAddr = nullptr;
83- void* otherDeviceAddr = nullptr;89+ void* otherDeviceAddr = nullptr;
84- void* outDeviceAddr = nullptr;90+ void* outDeviceAddr = nullptr;
85- aclTensor* self = nullptr;91+ aclTensor* self = nullptr;
86- aclTensor* other = nullptr;92+ aclTensor* other = nullptr;
87- aclTensor* out = nullptr;93+ aclTensor* out = nullptr;
88- std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7};94+ std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7};
89- std::vector<float> otherHostData = {1, 1, 1, 2, 2, 2, 3, 3};95+ std::vector<float> otherHostData = {1, 1, 1, 2, 2, 2, 3, 3};
90- std::vector<float> outHostData(8, 0);96+ std::vector<float> outHostData(8, 0);
91- // 创建self aclTensor97+ // 创建self aclTensor
92- ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);98+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
93- CHECK_RET(ret == ACL_SUCCESS, return ret);99+ CHECK_RET(ret == ACL_SUCCESS, return ret);
94- // 创建other aclTensor100+ // 创建other aclTensor
95- ret = CreateAclTensor(otherHostData, otherShape, &otherDeviceAddr, aclDataType::ACL_FLOAT, &other);101+ ret = CreateAclTensor(otherHostData, otherShape, &otherDeviceAddr, aclDataType::ACL_FLOAT, &other);
96- CHECK_RET(ret == ACL_SUCCESS, return ret);102+ CHECK_RET(ret == ACL_SUCCESS, return ret);
97- // 创建out aclTensor103+ // 创建out aclTensor
98- ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);104+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
99- CHECK_RET(ret == ACL_SUCCESS, return ret);105+ CHECK_RET(ret == ACL_SUCCESS, return ret);
100 106 
101- uint64_t workspaceSize = 0;107+ uint64_t workspaceSize = 0;
102- aclOpExecutor* executor;108+ aclOpExecutor* executor;
103 109 
104- // aclnnDiv接口调用示例110+ // aclnnDiv接口调用示例
105- LOG_PRINT("test aclnnDiv\n");111+ LOG_PRINT("test aclnnDiv\n");
106 112 
107- // 3. 调用CANN算子库API，需要修改为具体的Api名称113+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
108- // 调用aclnnDiv第一段接口114+ // 调用aclnnDiv第一段接口
109- ret = aclnnDivGetWorkspaceSize(self, other, out, &workspaceSize, &executor);115+ ret = aclnnDivGetWorkspaceSize(self, other, out, &workspaceSize, &executor);
110- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnDivGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);116+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnDivGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
111- // 根据第一段接口计算出的workspaceSize申请device内存117+ // 根据第一段接口计算出的workspaceSize申请device内存
112- void* workspaceAddr = nullptr;118+ void* workspaceAddr = nullptr;
113- if (workspaceSize > 0) {119+ if (workspaceSize > 0) {
114- ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);120+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
115- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);121+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
116- }122+ }
117- // 调用aclnnDiv第二段接口123+ // 调用aclnnDiv第二段接口
118- ret = aclnnDiv(workspaceAddr, workspaceSize, executor, stream);124+ ret = aclnnDiv(workspaceAddr, workspaceSize, executor, stream);
119- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnDiv failed. ERROR: %d\n", ret); return ret);125+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnDiv failed. ERROR: %d\n", ret); return ret);
120 126 
121- // 4. （固定写法）同步等待任务执行结束127+ // 4. （固定写法）同步等待任务执行结束
122- ret = aclrtSynchronizeStream(stream);128+ ret = aclrtSynchronizeStream(stream);
123- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);129+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
124 130 
125- // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改131+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
126- auto size = GetShapeSize(outShape);132+ auto size = GetShapeSize(outShape);
127- std::vector<float> resultData(size, 0);133+ std::vector<float> resultData(size, 0);
128- ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,134+ ret = aclrtMemcpy(
129- size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);135+ resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(resultData[0]),
130- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);136+ ACL_MEMCPY_DEVICE_TO_HOST);
131- for (int64_t i = 0; i < size; i++) {137+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
132- LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);138+ for (int64_t i = 0; i < size; i++) {
133- }139+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
140+ }
134 141 
135- // aclnnInplaceDiv接口调用示例142+ // aclnnInplaceDiv接口调用示例
136- LOG_PRINT("\ntest aclnnInplaceDiv\n");143+ LOG_PRINT("\ntest aclnnInplaceDiv\n");
137 144 
138- // 3. 调用CANN算子库API，需要修改为具体的Api名称145+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
139- // 调用aclnnInplaceDiv第一段接口146+ // 调用aclnnInplaceDiv第一段接口
140- ret = aclnnInplaceDivGetWorkspaceSize(self, other, &workspaceSize, &executor);147+ ret = aclnnInplaceDivGetWorkspaceSize(self, other, &workspaceSize, &executor);
141- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnInplaceDivGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);148+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnInplaceDivGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
142- // 根据第一段接口计算出的workspaceSize申请device内存149+ // 根据第一段接口计算出的workspaceSize申请device内存
143- if (workspaceSize > 0) {150+ if (workspaceSize > 0) {
144- ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);151+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
145- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);152+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
146- }153+ }
147- // 调用aclnnInplaceDiv第二段接口154+ // 调用aclnnInplaceDiv第二段接口
148- ret = aclnnInplaceDiv(workspaceAddr, workspaceSize, executor, stream);155+ ret = aclnnInplaceDiv(workspaceAddr, workspaceSize, executor, stream);
149- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnInplaceDiv failed. ERROR: %d\n", ret); return ret);156+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnInplaceDiv failed. ERROR: %d\n", ret); return ret);
150 157 
151- // 4. （固定写法）同步等待任务执行结束158+ // 4. （固定写法）同步等待任务执行结束
152- ret = aclrtSynchronizeStream(stream);159+ ret = aclrtSynchronizeStream(stream);
153- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);160+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
154 161 
155- // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改162+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
156- ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), selfDeviceAddr,163+ ret = aclrtMemcpy(
157- size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);164+ resultData.data(), resultData.size() * sizeof(resultData[0]), selfDeviceAddr, size * sizeof(resultData[0]),
158- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);165+ ACL_MEMCPY_DEVICE_TO_HOST);
159- for (int64_t i = 0; i < size; i++) {166+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
160- LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);167+ for (int64_t i = 0; i < size; i++) {
161- }168+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
169+ }
162 170 
163- // 6. 释放aclTensor171+ // 6. 释放aclTensor
164- aclDestroyTensor(self);172+ aclDestroyTensor(self);
165- aclDestroyTensor(other);173+ aclDestroyTensor(other);
166- aclDestroyTensor(out);174+ aclDestroyTensor(out);
167 175 
168- // 7. 释放Device资源176+ // 7. 释放Device资源
169- aclrtFree(selfDeviceAddr);177+ aclrtFree(selfDeviceAddr);
170- aclrtFree(otherDeviceAddr);178+ aclrtFree(otherDeviceAddr);
171- aclrtFree(outDeviceAddr);179+ aclrtFree(outDeviceAddr);
172- if (workspaceSize > 0) {180+ if (workspaceSize > 0) {
173- aclrtFree(workspaceAddr);181+ aclrtFree(workspaceAddr);
174- }182+ }
175- aclrtDestroyStream(stream);183+ aclrtDestroyStream(stream);
176- aclrtResetDevice(deviceId);184+ aclrtResetDevice(deviceId);
177- aclFinalize();185+ aclFinalize();
178 186 
179- return 0;187+ return 0;
180}188}
Amath/div/op_api/aclnn_div.cpp+863-0
@@ -0,0 +1,863 @@
1+/**
2+ * Copyright (c) 2025 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 "aclnn_div.h"
12+#include "aclnn_kernels/cast.h"
13+#include "aclnn_kernels/contiguous.h"
14+#include "math/floor_div/op_api/floordiv.h"
15+#include "math/real_div/op_api/realdiv.h"
16+#include "math/trunc/op_host/op_api/trunc.h"
17+#include "math/muls/op_api/muls.h"
18+#include "common/op_api_def.h"
19+#include "aclnn_kernels/common/op_error_check.h"
20+#include "opdev/common_types.h"
21+#include "opdev/data_type_utils.h"
22+#include "opdev/format_utils.h"
23+#include "opdev/op_dfx.h"
24+#include "opdev/op_executor.h"
25+#include "opdev/op_log.h"
26+#include "opdev/shape_utils.h"
27+#include "opdev/tensor_view_utils.h"
28+#include "opdev/platform.h"
29+ 
30+using namespace op;
31+#ifdef __cplusplus
32+extern "C" {
33+#endif
34+ 
35+op::DataType PromoteIntegerInputsToFloat(const op::DataType input) {
36+ if (IsIntegralType(input)) {
37+ return op::DataType::DT_FLOAT;
38+ }
39+ return input;
40+}
41+ 
42+static op::DataType InnerTypeToComplexType(const op::DataType input)
43+{
44+ switch (input) {
45+ case op::DataType::DT_BF16:
46+ // BFloat16 has range equivalent to Float,
47+ // so we map it to ComplexFloat.
48+ return op::DataType::DT_COMPLEX64;
49+ case op::DataType::DT_FLOAT16:
50+ return op::DataType::DT_COMPLEX32;
51+ case op::DataType::DT_FLOAT:
52+ return op::DataType::DT_COMPLEX64;
53+ case op::DataType::DT_DOUBLE:
54+ return op::DataType::DT_COMPLEX128;
55+ case op::DataType::DT_COMPLEX32:
56+ return op::DataType::DT_COMPLEX32;
57+ case op::DataType::DT_COMPLEX64:
58+ return op::DataType::DT_COMPLEX64;
59+ case op::DataType::DT_COMPLEX128:
60+ return op::DataType::DT_COMPLEX128;
61+ default:
62+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Unknown Complex ScalarType for [%s]", ToString(input).GetString());
63+ return op::DataType::DT_UNDEFINED;
64+ }
65+}
66+ 
67+static op::DataType CombineCategoriesWithComplex(const op::DataType higher, const op::DataType lower)
68+{
69+ if (IsComplexType(higher)) {
70+ return higher;
71+ } else if (IsComplexType(lower)) {
72+ // preserve value type of higher if it is floating type.
73+ if (IsFloatingType(higher)) {
74+ return InnerTypeToComplexType(higher);
75+ }
76+ // in case of integral input
77+ // lower complex takes precedence.
78+ return lower;
79+ } else if (IsFloatingType(higher)) {
80+ return higher;
81+ }
82+ if (higher == op::DataType::DT_BOOL || IsFloatingType(lower)) {
83+ return op::PromoteType(higher, lower);
84+ }
85+ if (higher != op::DataType::DT_UNDEFINED) {
86+ return higher;
87+ }
88+ return lower;
89+}
90+ 
91+static op::DataType GetScalarDefaultDtype(const op::DataType input)
92+{
93+ if (IsComplexType(input)) {
94+ return op::DataType::DT_COMPLEX64;
95+ } else if (IsFloatingType(input)) {
96+ return op::DataType::DT_FLOAT;
97+ }
98+ return input;
99+}
100+ 
101+// 根据API定义，需要列出所能支持的所有dtype
102+static const std::initializer_list<op::DataType> ASCEND910_DTYPE_SUPPORT_LIST = {
103+ op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT, op::DataType::DT_INT64, op::DataType::DT_INT32,
104+ op::DataType::DT_INT16, op::DataType::DT_INT8, op::DataType::DT_UINT8, op::DataType::DT_DOUBLE,
105+ op::DataType::DT_BOOL, op::DataType::DT_COMPLEX64, op::DataType::DT_COMPLEX128};
106+ 
107+static const std::initializer_list<op::DataType> ASCEND910B_DTYPE_SUPPORT_LIST = {
108+ op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT, op::DataType::DT_INT64, op::DataType::DT_INT32,
109+ op::DataType::DT_INT16, op::DataType::DT_INT8, op::DataType::DT_UINT8, op::DataType::DT_DOUBLE,
110+ op::DataType::DT_BOOL, op::DataType::DT_BF16, op::DataType::DT_COMPLEX64, op::DataType::DT_COMPLEX128};
111+ 
112+static const int MODE_REAL_DIV = 0;
113+static const int MODE_TRUNC_DIV = 1;
114+static const int MODE_FLOOR_DIV = 2;
115+ 
116+static const std::initializer_list<DataType>& GetDtypeSupportList()
117+{
118+ auto socVersion = GetCurrentPlatformInfo().GetSocVersion();
119+ if (socVersion == SocVersion::ASCEND910B || socVersion == SocVersion::ASCEND910_93 ||
120+ socVersion == SocVersion::ASCEND910_95) {
121+ return ASCEND910B_DTYPE_SUPPORT_LIST;
122+ } else {
123+ return ASCEND910_DTYPE_SUPPORT_LIST;
124+ }
125+}
126+ 
127+static bool CheckNotNull(const aclTensor* self, const aclTensor* other, const aclTensor* out)
CANN-robot
CANN-robotCANN-robot1月4日

指针与引用安全: CheckNotNull函数中空指针检查顺序不合理。当前检查顺序是out->other->self，但通常应该按照参数传递顺序或重要性顺序检查。更重要的是，当self或other为空时，函数返回false，但调用者可能无法区分是哪个参数为空，因为错误信息不明确。

问题类型: 指针与引用安全 文件路径: math/div/op_api/aclnn_div.cpp 行号: 127 问题代码:

static bool CheckNotNull(const aclTensor* self, const aclTensor* other, const aclTensor* out)
{
    OP_CHECK_NULL(out, return false);
    OP_CHECK_NULL(other, return false);
    OP_CHECK_NULL(self, return false);
    return true;
}

修改建议:

建议统一检查顺序为self->other->out，与参数顺序一致。同时考虑在OP_CHECK_NULL宏中添加更详细的错误信息，或者使用单独的日志记录哪个参数为空。

此评论由代码审查工具自动生成

likedislike
128+{
129+ OP_CHECK_NULL(out, return false);
130+ OP_CHECK_NULL(other, return false);
131+ OP_CHECK_NULL(self, return false);
132+ return true;
133+}
134+ 
135+static inline op::DataType CompatibleInferDivDtype(const op::DataType selfDtype, const op::DataType otherDtype)
136+{
137+ // RealDiv算子需要对self和other两个输入做隐式数据类型转换，根据具体算子语义按需调用
138+ auto promoteType = op::PromoteType(selfDtype, otherDtype);
139+ // 下沉PTA入口操作将入参类型转化成FLOAT进行后续处理
140+ promoteType = (IsFloatingType(promoteType) || IsComplexType(promoteType) || promoteType == op::DataType::DT_BOOL) ?
141+ promoteType :
142+ op::DataType::DT_FLOAT;
143+ return promoteType;
144+}
145+ 
146+static inline op::DataType InferDivModeDtype(
147+ const op::DataType selfDtype, const op::DataType otherDtype, const int mode)
148+{
149+ auto promoteType = op::PromoteType(selfDtype, otherDtype);
150+ // 下沉PTA入口操作将入参类型转化成FLOAT进行后续处理
151+ if (mode == MODE_REAL_DIV && promoteType != op::DataType::DT_INT32 && promoteType != op::DataType::DT_BOOL) {
152+ // IterateBase 配置特殊处理
153+ promoteType = PromoteIntegerInputsToFloat(promoteType);
154+ }
155+ if (mode == MODE_TRUNC_DIV && promoteType == DataType::DT_DOUBLE) {
156+ promoteType = DataType::DT_FLOAT;
157+ }
158+ return promoteType;
159+}
160+ 
161+static inline op::DataType CompatibleInferDivsDtype(const op::DataType selfDtype, const op::DataType otherDtype)
162+{
163+ auto promoteType = (IsFloatingType(selfDtype) || IsComplexType(selfDtype)) ? selfDtype : op::DataType::DT_FLOAT;
164+ promoteType = (selfDtype == op::DataType::DT_BOOL && otherDtype == op::DataType::DT_BOOL) ? selfDtype : promoteType;
165+ promoteType = (IsComplexType(otherDtype)) ? op::PromoteType(promoteType, otherDtype) : promoteType;
166+ return promoteType;
167+}
168+ 
169+static aclnnStatus CompatibleInferDivModeDtype(
170+ const op::DataType selfDtype, const op::DataType otherDtype, const int mode, op::DataType& promoteType)
171+{
172+ promoteType = op::PromoteType(selfDtype, otherDtype);
173+ if (mode == MODE_TRUNC_DIV || mode == MODE_FLOOR_DIV) {
174+ CHECK_RET(selfDtype != op::DataType::DT_COMPLEX128, ACLNN_ERR_PARAM_INVALID);
175+ CHECK_RET(selfDtype != op::DataType::DT_COMPLEX64, ACLNN_ERR_PARAM_INVALID);
176+ }
177+ // 根据mode分三种场景调用算子计算
178+ if (mode == MODE_FLOOR_DIV) {
179+ promoteType = (promoteType == op::DataType::DT_BOOL) ? op::DataType::DT_FLOAT : promoteType;
180+ } else {
181+ promoteType = ((promoteType != op::DataType::DT_FLOAT) && (promoteType != op::DataType::DT_FLOAT16) &&
182+ (promoteType != op::DataType::DT_COMPLEX64) && (promoteType != op::DataType::DT_COMPLEX128) &&
183+ (promoteType != op::DataType::DT_BF16) && (promoteType != op::DataType::DT_BOOL)) ?
184+ op::DataType::DT_FLOAT :
185+ promoteType;
186+ }
187+ return ACLNN_SUCCESS;
188+}
189+ 
190+static aclnnStatus CompatibleInferDivsModeDtype(
191+ const op::DataType selfDtype, const op::DataType otherDtype, const int mode, op::DataType& promoteType)
192+{
193+ promoteType = op::PromoteType(selfDtype, otherDtype);
194+ if ((mode == MODE_TRUNC_DIV || mode == MODE_FLOOR_DIV) &&
195+ (promoteType == op::DataType::DT_COMPLEX128 || promoteType == op::DataType::DT_COMPLEX64)) {
196+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "promoteType do not support DT_COMPLEX128 or DT_COMPLEX64.");
197+ return ACLNN_ERR_PARAM_INVALID;
198+ }
199+ // 根据mode分三种场景调用算子计算
200+ if (mode == MODE_FLOOR_DIV) {
201+ promoteType = (promoteType == op::DataType::DT_BOOL) ? op::DataType::DT_FLOAT : promoteType;
202+ } else {
203+ promoteType = ((selfDtype != op::DataType::DT_FLOAT) && (selfDtype != op::DataType::DT_FLOAT16) &&
204+ (selfDtype != op::DataType::DT_BF16) && (promoteType != op::DataType::DT_BOOL)) ?
205+ op::DataType::DT_FLOAT :
206+ selfDtype;
207+ promoteType = (IsComplexType(selfDtype) || IsComplexType(otherDtype)) ? op::PromoteType(selfDtype, otherDtype) :
208+ promoteType;
209+ }
210+ return ACLNN_SUCCESS;
211+}
212+ 
213+static inline op::DataType InferDivsModeDtype(
214+ const op::DataType selfDtype, const op::DataType otherDtype, const int mode)
215+{
216+ auto scalarDefaultDtype = GetScalarDefaultDtype(otherDtype);
217+ auto promoteType = CombineCategoriesWithComplex(selfDtype, scalarDefaultDtype);
218+ if (mode == MODE_REAL_DIV && promoteType != op::DataType::DT_INT32 && promoteType != op::DataType::DT_BOOL) {
219+ // IterateBase 配置特殊处理
220+ promoteType = PromoteIntegerInputsToFloat(promoteType);
221+ }
222+ 
223+ if (mode == MODE_TRUNC_DIV && promoteType == DataType::DT_DOUBLE) {
224+ promoteType = DataType::DT_FLOAT;
225+ }
226+ 
227+ if (promoteType == DataType::DT_COMPLEX32) {
228+ promoteType = DataType::DT_COMPLEX64;
229+ }
230+ 
231+ return promoteType;
232+}
233+ 
234+static bool CheckDtypeValid(const aclTensor* self, const aclTensor* other)
235+{
236+ auto supportList = GetDtypeSupportList();
237+ // 检查other的数据类型是否在div算子的支持列表内
238+ OP_CHECK_DTYPE_NOT_SUPPORT(other, supportList, return false);
239+ 
240+ // 检查self的数据类型是否在div算子的支持列表内
241+ OP_CHECK_DTYPE_NOT_SUPPORT(self, supportList, return false);
242+ return true;
243+}
244+ 
245+static bool CheckDtypeValidScalar(const aclTensor* self, const aclScalar* other)
246+{
247+ auto supportList = GetDtypeSupportList();
248+ // 检查other的数据类型是否在div算子的支持列表内
249+ OP_CHECK_DTYPE_NOT_SUPPORT(other, supportList, return false);
250+ 
251+ // 检查self的数据类型是否在div算子的支持列表内
252+ OP_CHECK_DTYPE_NOT_SUPPORT(self, supportList, return false);
253+ return true;
254+}
255+ 
256+static bool CheckPromoteType(const aclTensor* self, const aclTensor* other, const aclTensor* y, const int mode)
257+{
258+ // 检查self和other能否做数据类型推导
259+ auto promoteType = (GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_95) ?
260+ InferDivModeDtype(self->GetDataType(), other->GetDataType(), mode) :
261+ op::PromoteType(self->GetDataType(), other->GetDataType());
262+ if (promoteType == DataType::DT_UNDEFINED) {
263+ OP_LOGE(
264+ ACLNN_ERR_PARAM_INVALID, "Self dtype %s and other dtype %s can not promote dtype.",
265+ op::ToString(self->GetDataType()).GetString(), op::ToString(other->GetDataType()).GetString());
266+ return false;
267+ }
268+ 
269+ // 检查推导后的数据类型能否转换为输出的数据类型
270+ OP_CHECK_RESULT_DTYPE_CAST_FAILED(promoteType, y->GetDataType(), return false);
271+ return true;
272+}
273+ 
274+static bool CheckShape(const aclTensor* self, const aclTensor* other, const aclTensor* y)
275+{
276+ // 输入维度不超过8维
277+ OP_CHECK_MAX_DIM(self, MAX_SUPPORT_DIMS_NUMS, return false);
278+ OP_CHECK_MAX_DIM(other, MAX_SUPPORT_DIMS_NUMS, return false);
279+ 
280+ // self和other需满足broadcast关系
281+ op::Shape broadcastShape;
282+ OP_CHECK_BROADCAST_AND_INFER_SHAPE(self, other, broadcastShape, return false);
283+ OP_CHECK_SHAPE_NOT_EQUAL_WITH_EXPECTED_SIZE(y, broadcastShape, return false);
284+ return true;
285+}
286+ 
287+static bool CheckMode(int mode)
288+{
289+ if (mode > MODE_FLOOR_DIV || mode < MODE_REAL_DIV) {
290+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "mode should be between 0 and 2, but current is %d", mode);
291+ return false;
292+ }
293+ return true;
294+}
295+ 
296+static bool CheckFormat(const aclTensor* self, const aclTensor* other, const aclTensor* out)
297+{
298+ // 格式不能是私有格式
299+ // 校验self格式
300+ if (IsPrivateFormat(self->GetStorageFormat())) {
301+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Format only support ND、NCHW、NHWC、HWCN、NDHWC、NCDHW.");
302+ return false;
303+ }
304+ // 校验other格式
305+ if (IsPrivateFormat(other->GetStorageFormat())) {
306+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Format only support ND、NCHW、NHWC、HWCN、NDHWC、NCDHW.");
307+ return false;
308+ }
309+ // 校验out格式
310+ if (IsPrivateFormat(out->GetStorageFormat())) {
311+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Format only support ND、NCHW、NHWC、HWCN、NDHWC、NCDHW.");
312+ return false;
313+ }
314+ 
315+ return true;
316+}
317+ 
318+static bool CheckFormatScalar(const aclTensor* self, const aclTensor* out)
319+{
320+ // 格式不能是私有格式
321+ // 校验out格式
322+ if (IsPrivateFormat(out->GetStorageFormat())) {
323+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Format only support ND、NCHW、NHWC、HWCN、NDHWC、NCDHW.");
324+ return false;
325+ }
326+ // 校验self格式
327+ if (IsPrivateFormat(self->GetStorageFormat())) {
328+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Format only support ND、NCHW、NHWC、HWCN、NDHWC、NCDHW.");
329+ return false;
330+ }
331+ return true;
332+}
333+ 
334+static aclnnStatus CheckParams(const aclTensor* self, const aclTensor* other, const aclTensor* y, const int mode)
335+{
336+ // 1. 检查参数是否为空指针
337+ CHECK_RET(CheckNotNull(self, other, y), ACLNN_ERR_PARAM_NULLPTR);
338+ 
339+ // 2. 检查双输入是否能broadcast
340+ CHECK_RET(CheckShape(self, other, y), ACLNN_ERR_PARAM_INVALID);
341+ 
342+ // 3. 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验
343+ CHECK_RET(CheckDtypeValid(self, other), ACLNN_ERR_PARAM_INVALID);
344+ 
345+ // 4. 检查self和other能否做数据类型推导以及推导的数据类型能否转换为输出数据类型
346+ CHECK_RET(CheckPromoteType(self, other, y, mode), ACLNN_ERR_PARAM_INVALID);
347+ 
348+ // 5. 检查数据格式是否支持
349+ CHECK_RET(CheckFormat(self, other, y), ACLNN_ERR_PARAM_INVALID);
350+ 
351+ return ACLNN_SUCCESS;
352+}
353+ 
354+inline static bool isDivsMixDtypeSupport(const aclTensor* self, const aclScalar* other)
355+{
356+ auto socVersion = GetCurrentPlatformInfo().GetSocVersion();
357+ if (socVersion != SocVersion::ASCEND910B && socVersion != SocVersion::ASCEND910_93) {
358+ return false;
359+ }
360+ return (self->GetDataType() == DataType::DT_FLOAT16 && other->GetDataType() == DataType::DT_FLOAT) ||
361+ (self->GetDataType() == DataType::DT_FLOAT && other->GetDataType() == DataType::DT_FLOAT16) ||
362+ (self->GetDataType() == DataType::DT_BF16 && other->GetDataType() == DataType::DT_FLOAT) ||
363+ (self->GetDataType() == DataType::DT_FLOAT && other->GetDataType() == DataType::DT_BF16) ||
364+ (self->GetDataType() == DataType::DT_BF16 && other->GetDataType() == DataType::DT_DOUBLE);
365+}
366+ 
367+aclnnStatus aclnnDivGetWorkspaceSize(
368+ const aclTensor* self, const aclTensor* other, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)
369+{
370+ L2_DFX_PHASE_1(aclnnDiv, DFX_IN(self, other), DFX_OUT(out));
371+ // 固定写法，创建OpExecutor
372+ auto uniqueExecutor = CREATE_EXECUTOR();
373+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
374+ 
375+ // 固定写法，参数检查
376+ auto ret = CheckParams(self, other, out, MODE_REAL_DIV);
377+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
378+ 
379+ // div算子的空tensor在kernel中支持，对标竞品根据算子实际情况补充
380+ if (self->IsEmpty() || other->IsEmpty()) {
381+ // 根据实际支持情况补充
382+ *workspaceSize = 0;
383+ uniqueExecutor.ReleaseTo(executor);
384+ return ACLNN_SUCCESS;
385+ }
386+ 
387+ // RealDiv算子需要对self和other两个输入做隐式数据类型转换，根据具体算子语义按需调用
388+ auto promoteType = (GetCurrentPlatformInfo().GetSocVersion() != SocVersion::ASCEND910_95) ?
389+ CompatibleInferDivDtype(self->GetDataType(), other->GetDataType()) :
390+ InferDivModeDtype(self->GetDataType(), other->GetDataType(), MODE_REAL_DIV);
391+ 
392+ bool isSupportNonContiguous = GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_95;
393+ 
394+ // 处理self输入
395+ const aclTensor* selfProcessed = nullptr;
396+ if (self->GetDataType() == promoteType && isSupportNonContiguous) {
397+ selfProcessed = uniqueExecutor.get()->CreateView(
398+ self, self->GetViewShape(), self->GetStorageShape(), self->GetViewStrides(), self->GetViewOffset());
399+ } else {
400+ // 固定写法，将输入self转换成连续的tensor
401+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
402+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
403+ 
404+ // 将输入self的数据类型转换成隐式数据类型，根据具体算子语义按需调用
405+ selfProcessed = l0op::Cast(selfContiguous, promoteType, uniqueExecutor.get());
406+ }
407+ CHECK_RET(selfProcessed != nullptr, ACLNN_ERR_INNER_NULLPTR);
408+ 
409+ // 处理other输入
410+ const aclTensor* otherProcessed = nullptr;
411+ if (other->GetDataType() == promoteType && isSupportNonContiguous) {
412+ otherProcessed = uniqueExecutor.get()->CreateView(
413+ other, other->GetViewShape(), other->GetStorageShape(), other->GetViewStrides(), other->GetViewOffset());
414+ } else {
415+ // 固定写法，将输入other转换成连续的tensor
416+ auto otherContiguous = l0op::Contiguous(other, uniqueExecutor.get());
417+ CHECK_RET(otherContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
418+ 
419+ // 将输入other的数据类型转换成隐式数据类型，根据具体算子语义按需调用
420+ otherProcessed = l0op::Cast(otherContiguous, promoteType, uniqueExecutor.get());
421+ }
422+ CHECK_RET(otherProcessed != nullptr, ACLNN_ERR_INNER_NULLPTR);
423+ 
424+ // 调用l0算子RealDiv进行计算
425+ auto divOpOut = l0op::RealDiv(selfProcessed, otherProcessed, uniqueExecutor.get());
426+ CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
427+ 
428+ // 固定写法，将计算结果转换成输出out的数据类型
429+ auto castOut = l0op::Cast(divOpOut, out->GetDataType(), uniqueExecutor.get());
430+ CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
431+ 
432+ // 固定写法，将计算结果拷贝到输出out上，out可能是非连续的tensor
433+ auto viewCopyResult = l0op::ViewCopy(castOut, out, uniqueExecutor.get());
434+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
435+ 
436+ // 固定写法，获取计算过程中需要使用的workspace大小
437+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
438+ uniqueExecutor.ReleaseTo(executor);
439+ return ACLNN_SUCCESS;
440+}
441+ 
442+aclnnStatus aclnnDiv(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
443+{
444+ L2_DFX_PHASE_2(aclnnDiv);
445+ // 固定写法，调用框架能力，完成计算
446+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
447+}
448+ 
449+static bool CheckNotNullScalar(const aclTensor* self, const aclScalar* other, const aclTensor* out)
450+{
451+ OP_CHECK_NULL(self, return false);
452+ OP_CHECK_NULL(other, return false);
453+ OP_CHECK_NULL(out, return false);
454+ return true;
455+}
456+ 
457+static bool CheckPromoteTypeScalar(const aclTensor* self, const aclScalar* other, const aclTensor* y, const int mode)
458+{
459+ if (GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_95) {
460+ // 检查self和other能否做数据类型推导
461+ auto promoteType = InferDivsModeDtype(self->GetDataType(), other->GetDataType(), mode);
462+ if (promoteType == DataType::DT_UNDEFINED) {
463+ OP_LOGE(
464+ ACLNN_ERR_PARAM_INVALID, "Self dtype %s and other dtype %s can not promote dtype.",
465+ op::ToString(self->GetDataType()).GetString(), op::ToString(other->GetDataType()).GetString());
466+ return false;
467+ }
468+ // 检查推导后的数据类型能否转换为输出的数据类型
469+ OP_CHECK_RESULT_DTYPE_CAST_FAILED(promoteType, y->GetDataType(), return false);
470+ return true;
471+ }
472+ // 检查self的数据类型能否转换为输出的数据类型
473+ OP_CHECK_RESULT_DTYPE_CAST_FAILED(self->GetDataType(), y->GetDataType(), return false);
474+ return true;
475+}
476+ 
477+static bool CheckShapeScalar(const aclTensor* self, const aclTensor* y)
478+{
479+ OP_CHECK_MAX_DIM(self, MAX_SUPPORT_DIMS_NUMS, return false);
480+ 
481+ if (self->GetViewShape() != y->GetViewShape()) {
482+ OP_LOGE(
483+ ACLNN_ERR_PARAM_INVALID, "Shape of out should be %s, but current is %s.",
484+ op::ToString(self->GetViewShape()).GetString(), op::ToString(y->GetViewShape()).GetString());
485+ return false;
486+ }
487+ return true;
488+}
489+ 
490+static aclnnStatus CheckParamsScalar(const aclTensor* self, const aclScalar* other, const aclTensor* y, const int mode)
491+{
492+ CHECK_RET(CheckNotNullScalar(self, other, y), ACLNN_ERR_PARAM_NULLPTR);
493+ 
494+ CHECK_RET(CheckFormatScalar(self, y), ACLNN_ERR_PARAM_INVALID);
495+ 
496+ CHECK_RET(CheckDtypeValidScalar(self, other), ACLNN_ERR_PARAM_INVALID);
497+ 
498+ CHECK_RET(CheckShapeScalar(self, y), ACLNN_ERR_PARAM_INVALID);
499+ 
500+ CHECK_RET(CheckPromoteTypeScalar(self, other, y, mode), ACLNN_ERR_PARAM_INVALID);
501+ 
502+ return ACLNN_SUCCESS;
503+}
504+ 
505+static bool CanUseMuls(const aclTensor* self, const aclScalar* other)
506+{
507+ if (GetCurrentPlatformInfo().GetSocVersion() != SocVersion::ASCEND910_95) {
508+ return false;
509+ }
510+ if (self->GetDataType() != op::DataType::DT_FLOAT16 && self->GetDataType() != op::DataType::DT_BF16 &&
511+ self->GetDataType() != op::DataType::DT_FLOAT) {
512+ return false;
513+ }
514+ if (other->GetDataType() != op::DataType::DT_FLOAT16 && other->GetDataType() != op::DataType::DT_BF16 &&
515+ other->GetDataType() != op::DataType::DT_FLOAT) {
516+ return false;
517+ }
518+ 
519+ return true;
520+}
521+ 
522+aclnnStatus aclnnDivsGetWorkspaceSize(
523+ const aclTensor* self, const aclScalar* other, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)
524+{
525+ L2_DFX_PHASE_1(aclnnDivs, DFX_IN(self, other), DFX_OUT(out));
526+ 
527+ auto uniqueExecutor = CREATE_EXECUTOR();
528+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
529+ 
530+ // 调用适配aclScalar参数检查
531+ auto ret = CheckParamsScalar(self, other, out, MODE_REAL_DIV);
532+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
533+ 
534+ if (self->IsEmpty()) {
535+ *workspaceSize = 0;
536+ uniqueExecutor.ReleaseTo(executor);
537+ return ACLNN_SUCCESS;
538+ }
539+ 
540+ bool isSupportNonContiguous = GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_95;
541+ 
542+ // 判断输入是否符合kernel支持的混合输入类型
543+ bool isMixDataType = isDivsMixDtypeSupport(self, other);
544+ const aclTensor* divOpOut = nullptr;
545+ if (isMixDataType) {
546+ // aclScalar转aclTensor
547+ auto promoteType =
548+ other->GetDataType() == op::DataType::DT_DOUBLE ? op::DataType::DT_FLOAT : other->GetDataType();
549+ auto otherConvert = uniqueExecutor.get()->ConvertToTensor(other, promoteType);
550+ CHECK_RET(otherConvert != nullptr, ACLNN_ERR_INNER_NULLPTR);
551+ auto selfProcessed = isSupportNonContiguous ? uniqueExecutor.get()->CreateView(
552+ self, self->GetViewShape(), self->GetStorageShape(),
553+ self->GetViewStrides(), self->GetViewOffset()) :
554+ l0op::Contiguous(self, uniqueExecutor.get());
555+ CHECK_RET(selfProcessed != nullptr, ACLNN_ERR_INNER_NULLPTR);
556+ divOpOut = l0op::RealDiv(selfProcessed, otherConvert, uniqueExecutor.get());
557+ CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
558+ } else {
559+ auto promoteType = (GetCurrentPlatformInfo().GetSocVersion() != SocVersion::ASCEND910_95) ?
CANN-robot
CANN-robotCANN-robot1月4日

代码结构与可维护性: aclnnDivsGetWorkspaceSize函数中类型推导逻辑重复且混乱。第560-570行重新计算promoteType，但之前已经计算过（第559-561行），且逻辑复杂，包含多个条件判断，难以理解和维护。

问题类型: 代码结构与可维护性 文件路径: math/div/op_api/aclnn_div.cpp 行号: 559 问题代码:

        auto promoteType = (GetCurrentPlatformInfo().GetSocVersion() != SocVersion::ASCEND910_95) ?
                               CompatibleInferDivsDtype(self->GetDataType(), other->GetDataType()) :
                               InferDivsModeDtype(self->GetDataType(), other->GetDataType(), MODE_REAL_DIV);
        promoteType = (IsFloatingType(self->GetDataType()) || IsComplexType(self->GetDataType())) ?
                          self->GetDataType() : op::DataType::DT_FLOAT;
        promoteType = (self->GetDataType() == op::DataType::DT_BOOL && other->GetDataType() == op::DataType::DT_BOOL) ?
                          self->GetDataType() : promoteType;
        promoteType = (IsComplexType(other->GetDataType())) ? op::PromoteType(promoteType, other->GetDataType()) : promoteType;
        if (GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_95) {
            promoteType = op::PromoteType(self->GetDataType(), other->GetDataType()) == op::DataType::DT_INT32
                          ? op::DataType::DT_INT32 : promoteType;
        }

修改建议:

建议重构类型推导逻辑，将复杂的条件判断提取到专门的函数中，避免重复计算和逻辑混乱。确保类型推导逻辑集中在一处，便于维护和测试。

此评论由代码审查工具自动生成

likedislike
560+ CompatibleInferDivsDtype(self->GetDataType(), other->GetDataType()) :
561+ InferDivsModeDtype(self->GetDataType(), other->GetDataType(), MODE_REAL_DIV);
562+ promoteType = (IsFloatingType(self->GetDataType()) || IsComplexType(self->GetDataType())) ?
CANN-robot
CANN-robotCANN-robot1月4日

代码冗余: 在aclnnDivsGetWorkspaceSize函数中，promoteType的计算逻辑存在冗余。第560-566行重新计算了promoteType，但第562-566行的计算与前面的逻辑重复。

问题类型: 代码冗余 文件路径: math/div/op_api/aclnn_div.cpp 行号: 562 问题代码:

        promoteType = (IsFloatingType(self->GetDataType()) || IsComplexType(self->GetDataType())) ?
                          self->GetDataType() : op::DataType::DT_FLOAT;
        promoteType = (self->GetDataType() == op::DataType::DT_BOOL && other->GetDataType() == op::DataType::DT_BOOL) ?
                          self->GetDataType() : promoteType;
        promoteType = (IsComplexType(other->GetDataType())) ? op::PromoteType(promoteType, other->GetDataType()) : promoteType;

修改建议:

1. 删除冗余的promoteType计算，直接使用第560行计算的结果
2. 或者重构代码，将数据类型推导逻辑提取到单独的函数中，避免重复

此评论由代码审查工具自动生成

likedislike
563+ self->GetDataType() : op::DataType::DT_FLOAT;
564+ promoteType = (self->GetDataType() == op::DataType::DT_BOOL && other->GetDataType() == op::DataType::DT_BOOL) ?
565+ self->GetDataType() : promoteType;
566+ promoteType = (IsComplexType(other->GetDataType())) ? op::PromoteType(promoteType, other->GetDataType()) : promoteType;
567+ if (GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_95) {
568+ promoteType = op::PromoteType(self->GetDataType(), other->GetDataType()) == op::DataType::DT_INT32
569+ ? op::DataType::DT_INT32 : promoteType;
570+ }
571+ 
572+ bool canUseMuls = CanUseMuls(self, other);
573+ if (self->GetDataType() == promoteType && isSupportNonContiguous && !canUseMuls) {
574+ // aclScalar转aclTensor
575+ auto otherConvert = uniqueExecutor.get()->ConvertToTensor(other, promoteType);
576+ CHECK_RET(otherConvert != nullptr, ACLNN_ERR_INNER_NULLPTR);
577+ auto selfWithStride = uniqueExecutor.get()->CreateView(
578+ self, self->GetViewShape(), self->GetStorageShape(), self->GetViewStrides(), self->GetViewOffset());
579+ divOpOut = l0op::RealDiv(selfWithStride, otherConvert, uniqueExecutor.get());
580+ CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
581+ } else {
582+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
583+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
584+ auto selfCasted = l0op::Cast(selfContiguous, promoteType, uniqueExecutor.get());
585+ CHECK_RET(selfCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
586+ if (canUseMuls) {
587+ float invB = static_cast<float>(1.0f) / (other->ToFloat());
588+ aclScalar* invBPtr = uniqueExecutor.get()->AllocScalar(invB);
589+ divOpOut = l0op::Muls(selfCasted, invBPtr->ToFloat(), uniqueExecutor.get());
590+ } else {
591+ // aclScalar转aclTensor
592+ auto otherConvert = uniqueExecutor.get()->ConvertToTensor(other, promoteType);
593+ CHECK_RET(otherConvert != nullptr, ACLNN_ERR_INNER_NULLPTR);
594+ divOpOut = l0op::RealDiv(selfCasted, otherConvert, uniqueExecutor.get());
595+ }
596+ CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
597+ }
598+ }
599+ auto castOut = l0op::Cast(divOpOut, out->GetDataType(), uniqueExecutor.get());
600+ CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
601+ 
602+ auto viewCopyResult = l0op::ViewCopy(castOut, out, uniqueExecutor.get());
603+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
604+ 
605+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
606+ uniqueExecutor.ReleaseTo(executor);
607+ return ACLNN_SUCCESS;
608+}
609+ 
610+aclnnStatus aclnnDivs(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
611+{
612+ L2_DFX_PHASE_2(aclnnDivs);
613+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
614+}
615+ 
616+aclnnStatus aclnnDivModGetWorkspaceSize(
617+ const aclTensor* self, const aclTensor* other, int mode, aclTensor* out, uint64_t* workspaceSize,
618+ aclOpExecutor** executor)
619+{
620+ L2_DFX_PHASE_1(aclnnDivMod, DFX_IN(self, other, mode), DFX_OUT(out));
621+ 
622+ auto uniqueExecutor = CREATE_EXECUTOR();
623+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
624+ 
625+ auto ret = CheckParams(self, other, out, mode);
626+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
627+ CHECK_RET(CheckMode(mode), ACLNN_ERR_PARAM_INVALID);
628+ 
629+ if (self->IsEmpty() || other->IsEmpty()) {
630+ *workspaceSize = 0;
631+ uniqueExecutor.ReleaseTo(executor);
632+ return ACLNN_SUCCESS;
633+ }
634+ 
635+ auto otherContiguous = l0op::Contiguous(other, uniqueExecutor.get());
636+ CHECK_RET(otherContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
637+ 
638+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
639+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
640+ 
641+ auto selfCasted = selfContiguous;
642+ auto otherCasted = otherContiguous;
643+ op::DataType promoteType;
644+ bool needToInt32 = false;
645+ op::DataType oriType = out->GetDataType();
646+ if (GetCurrentPlatformInfo().GetSocVersion() != SocVersion::ASCEND910_95) {
647+ auto promoteRet = CompatibleInferDivModeDtype(self->GetDataType(), other->GetDataType(), mode, promoteType);
648+ CHECK_RET(promoteRet == ACLNN_SUCCESS, promoteRet);
649+ } else {
650+ promoteType = InferDivModeDtype(self->GetDataType(), other->GetDataType(), mode);
651+ // customization
652+ bool needToFloat = (promoteType == op::DataType::DT_BOOL && mode == MODE_FLOOR_DIV);
653+ promoteType = needToFloat ? op::DataType::DT_FLOAT : promoteType;
654+ // aicore is not supported, aicpu has problems when div 0
655+ needToInt32 = (promoteType == op::DataType::DT_INT16 && mode == MODE_FLOOR_DIV) ||
656+ ((promoteType == op::DataType::DT_INT8 || promoteType == op::DataType::DT_UINT8 ||
657+ promoteType == op::DataType::DT_INT16) &&
658+ mode == MODE_TRUNC_DIV);
659+ oriType = promoteType;
660+ promoteType = needToInt32 ? op::DataType::DT_INT32 : promoteType;
661+ }
662+ selfCasted = l0op::Cast(selfContiguous, promoteType, uniqueExecutor.get());
663+ CHECK_RET(selfCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
664+ otherCasted = l0op::Cast(otherContiguous, promoteType, uniqueExecutor.get());
665+ CHECK_RET(otherCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
666+ const aclTensor* divOpOut = nullptr;
667+ // 根据mode分三种场景调用算子计算
668+ if (mode == MODE_FLOOR_DIV) {
669+ divOpOut = l0op::FloorDiv(selfCasted, otherCasted, uniqueExecutor.get());
670+ } else {
671+ divOpOut = l0op::RealDiv(selfCasted, otherCasted, uniqueExecutor.get());
672+ CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
673+ if (mode == MODE_TRUNC_DIV && divOpOut->GetDataType() != op::DataType::DT_INT64 &&
674+ divOpOut->GetDataType() != op::DataType::DT_INT16) {
675+ divOpOut = l0op::Trunc(divOpOut, uniqueExecutor.get());
676+ }
677+ }
678+ CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
679+ if (needToInt32) {
680+ divOpOut = l0op::Cast(divOpOut, oriType, uniqueExecutor.get());
681+ CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
682+ }
683+ auto castOut = l0op::Cast(divOpOut, out->GetDataType(), uniqueExecutor.get());
684+ CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
685+ 
686+ auto viewCopyResult = l0op::ViewCopy(castOut, out, uniqueExecutor.get());
687+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
688+ 
689+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
690+ uniqueExecutor.ReleaseTo(executor);
691+ return ACLNN_SUCCESS;
692+}
693+ 
694+aclnnStatus aclnnDivMod(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
695+{
696+ L2_DFX_PHASE_2(aclnnDivMod);
697+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
698+}
699+ 
700+aclnnStatus aclnnDivModsGetWorkspaceSize(
701+ const aclTensor* self, const aclScalar* other, int mode, aclTensor* out, uint64_t* workspaceSize,
702+ aclOpExecutor** executor)
703+{
704+ L2_DFX_PHASE_1(aclnnDivMods, DFX_IN(self, other, mode), DFX_OUT(out));
705+ 
706+ auto uniqueExecutor = CREATE_EXECUTOR();
707+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
708+ 
709+ // 调用适配aclScalar参数检查
710+ auto ret = CheckParamsScalar(self, other, out, mode);
711+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
712+ CHECK_RET(CheckMode(mode), ACLNN_ERR_PARAM_INVALID);
713+ 
714+ if (self->IsEmpty()) {
715+ *workspaceSize = 0;
716+ uniqueExecutor.ReleaseTo(executor);
717+ return ACLNN_SUCCESS;
718+ }
719+ 
720+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
721+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
722+ 
723+ auto selfCasted = selfContiguous;
724+ op::DataType promoteType;
725+ bool needToInt32 = false;
726+ op::DataType oriType = out->GetDataType();
727+ if (GetCurrentPlatformInfo().GetSocVersion() != SocVersion::ASCEND910_95) {
728+ auto promoteRet = CompatibleInferDivsModeDtype(self->GetDataType(), other->GetDataType(), mode, promoteType);
729+ CHECK_RET(promoteRet == ACLNN_SUCCESS, promoteRet);
730+ } else {
731+ promoteType = InferDivsModeDtype(self->GetDataType(), other->GetDataType(), mode);
732+ // customization
733+ bool needToFloat = (promoteType == op::DataType::DT_BOOL && mode == MODE_FLOOR_DIV);
734+ promoteType = needToFloat ? op::DataType::DT_FLOAT : promoteType;
735+ // aicore is not supported, aicpu has problems when div 0
736+ needToInt32 = (promoteType == op::DataType::DT_INT16 && mode == MODE_FLOOR_DIV) ||
737+ ((promoteType == op::DataType::DT_INT8 || promoteType == op::DataType::DT_UINT8 ||
738+ promoteType == op::DataType::DT_INT16) &&
739+ mode == MODE_TRUNC_DIV);
740+ oriType = promoteType;
741+ promoteType = needToInt32 ? op::DataType::DT_INT32 : promoteType;
742+ }
743+ selfCasted = l0op::Cast(selfContiguous, promoteType, uniqueExecutor.get());
744+ CHECK_RET(selfCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
745+ auto otherCasted = uniqueExecutor.get()->ConvertToTensor(other, promoteType);
746+ CHECK_RET(otherCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
747+ const aclTensor* divOpOut = nullptr;
748+ // 根据mode分三种场景调用算子计算
749+ if (mode == MODE_FLOOR_DIV) {
750+ divOpOut = l0op::FloorDiv(selfCasted, otherCasted, uniqueExecutor.get());
751+ } else {
752+ divOpOut = l0op::RealDiv(selfCasted, otherCasted, uniqueExecutor.get());
753+ CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
754+ if (mode == MODE_TRUNC_DIV && divOpOut->GetDataType() != op::DataType::DT_INT64 &&
755+ divOpOut->GetDataType() != op::DataType::DT_INT16) {
756+ divOpOut = l0op::Trunc(divOpOut, uniqueExecutor.get());
757+ }
758+ }
759+ CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
760+ if (needToInt32) {
761+ divOpOut = l0op::Cast(divOpOut, oriType, uniqueExecutor.get());
762+ CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
763+ }
764+ auto castOut = l0op::Cast(divOpOut, out->GetDataType(), uniqueExecutor.get());
765+ CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
766+ 
767+ auto viewCopyResult = l0op::ViewCopy(castOut, out, uniqueExecutor.get());
768+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
769+ 
770+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
771+ uniqueExecutor.ReleaseTo(executor);
772+ return ACLNN_SUCCESS;
773+}
774+ 
775+aclnnStatus aclnnDivMods(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
776+{
777+ L2_DFX_PHASE_2(aclnnDivMods);
778+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
779+}
780+ 
781+static inline aclnnStatus CheckInplace(const aclTensor* selfRef, const aclTensor* other)
782+{
783+ OP_CHECK(
784+ selfRef != nullptr, OP_LOGE(ACLNN_ERR_PARAM_NULLPTR, "Expected selfRef not to be null."),
785+ return ACLNN_ERR_PARAM_NULLPTR);
786+ OP_CHECK(
787+ other != nullptr, OP_LOGE(ACLNN_ERR_PARAM_NULLPTR, "Expected other not to be null."),
788+ return ACLNN_ERR_PARAM_NULLPTR);
789+ op::Shape broadcastShape;
790+ OP_CHECK(
791+ BroadcastInferShape(selfRef->GetViewShape(), other->GetViewShape(), broadcastShape),
792+ OP_LOGE(
793+ ACLNN_ERR_PARAM_INVALID, "Shape of selfRef and other can't broadcast, got %s, %s.",
794+ op::ToString(selfRef->GetViewShape()).GetString(), op::ToString(other->GetViewShape()).GetString()),
795+ return ACLNN_ERR_PARAM_INVALID);
796+ OP_CHECK(
CANN-robot
CANN-robotCANN-robot1月4日

错误处理与异常安全: CheckInplace函数中错误信息不一致。第799行错误日志中使用了ACLNN_ERR_PARAM_NULLPTR，但实际错误是形状不匹配，应该使用ACLNN_ERR_PARAM_INVALID。错误码与错误描述不匹配。

问题类型: 错误处理与异常安全 文件路径: math/div/op_api/aclnn_div.cpp 行号: 796 问题代码:

    OP_CHECK(
        selfRef->GetViewShape() == broadcastShape,
        OP_LOGE(
            ACLNN_ERR_PARAM_NULLPTR, "Expected shape of selfRef should be %s, but got %s.",
            op::ToString(broadcastShape).GetString(), op::ToString(selfRef->GetViewShape()).GetString()),
        return ACLNN_ERR_PARAM_INVALID);

修改建议:

将第800行的ACLNN_ERR_PARAM_NULLPTR改为ACLNN_ERR_PARAM_INVALID，以正确反映错误类型。错误码应该与错误描述一致，形状不匹配属于参数无效错误，而不是空指针错误。

此评论由代码审查工具自动生成

likedislike
797+ selfRef->GetViewShape() == broadcastShape,
798+ OP_LOGE(
799+ ACLNN_ERR_PARAM_NULLPTR, "Expected shape of selfRef should be %s, but got %s.",
800+ op::ToString(broadcastShape).GetString(), op::ToString(selfRef->GetViewShape()).GetString()),
801+ return ACLNN_ERR_PARAM_INVALID);
802+ return ACLNN_SUCCESS;
803+}
804+ 
805+aclnnStatus aclnnInplaceDivGetWorkspaceSize(
806+ aclTensor* selfRef, const aclTensor* other, uint64_t* workspaceSize, aclOpExecutor** executor)
807+{
808+ auto ret = CheckInplace(selfRef, other);
809+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
810+ auto out = const_cast<aclTensor*>(selfRef);
811+ return aclnnDivGetWorkspaceSize(selfRef, other, out, workspaceSize, executor);
812+}
813+ 
814+aclnnStatus aclnnInplaceDiv(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
815+{
816+ L2_DFX_PHASE_2(aclnnInplaceDiv);
817+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
818+}
819+ 
820+aclnnStatus aclnnInplaceDivsGetWorkspaceSize(
821+ aclTensor* selfRef, const aclScalar* other, uint64_t* workspaceSize, aclOpExecutor** executor)
822+{
823+ auto out = const_cast<aclTensor*>(selfRef);
824+ return aclnnDivsGetWorkspaceSize(selfRef, other, out, workspaceSize, executor);
825+}
826+ 
827+aclnnStatus aclnnInplaceDivs(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
828+{
829+ L2_DFX_PHASE_2(aclnnInplaceDivs);
830+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
831+}
832+ 
833+aclnnStatus aclnnInplaceDivModGetWorkspaceSize(
834+ aclTensor* selfRef, const aclTensor* other, int mode, uint64_t* workspaceSize, aclOpExecutor** executor)
835+{
836+ auto ret = CheckInplace(selfRef, other);
837+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
838+ auto out = const_cast<aclTensor*>(selfRef);
839+ return aclnnDivModGetWorkspaceSize(selfRef, other, mode, out, workspaceSize, executor);
840+}
841+ 
842+aclnnStatus aclnnInplaceDivMod(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
843+{
844+ L2_DFX_PHASE_2(aclnnInplaceDivMod);
845+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
846+}
847+ 
848+aclnnStatus aclnnInplaceDivModsGetWorkspaceSize(
849+ aclTensor* selfRef, const aclScalar* other, int mode, uint64_t* workspaceSize, aclOpExecutor** executor)
850+{
851+ auto out = const_cast<aclTensor*>(selfRef);
852+ return aclnnDivModsGetWorkspaceSize(selfRef, other, mode, out, workspaceSize, executor);
853+}
854+ 
855+aclnnStatus aclnnInplaceDivMods(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
856+{
857+ L2_DFX_PHASE_2(aclnnInplaceDivMods);
858+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
859+}
860+ 
861+#ifdef __cplusplus
862+}
863+#endif
Amath/div/op_api/aclnn_div.h+128-0
@@ -0,0 +1,128 @@
1+/**
2+ * Copyright (c) 2025 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_DIV_H_
12+#define OP_API_INC_DIV_H_
13+ 
14+#include "aclnn/aclnn_base.h"
15+#include "aclnn_util.h"
16+ 
17+#ifdef __cplusplus
18+extern "C" {
19+#endif
20+ 
21+/**
22+ * @brief aclnnDiv的第一段接口，根据具体的计算流程，计算workspace大小。
23+ * @domain aclnn_math
24+ *
25+ * @param [in] self: npu
26+ * device侧的aclTensor，数据类型支持整型，浮点类型，且数据类型需要与other构成互相推导关系，shape需要与other满足broadcast关系。
27+ * 支持非连续的Tensor，数据格式支持ND，且数据格式需要与other一致。
28+ * @param [in] other: npu
29+ * device侧的aclTensor，数据类型支持整型，浮点类型，且数据类型需要与self构成互相推导关系，shape需要与self满足broadcast关系。
30+ * 支持非连续的Tensor，数据格式支持ND，且数据格式需要与self一致。
31+ * @param [in]
32+ * mode：余数处理方式的整型常量，枚举值如下：0-默认不执行舍入。1-将除法的小数部分舍入为零。2-向下舍入除法的结果。
33+ * @param [in] out: npu
34+ * device侧的aclTensor，数据类型支持整型，浮点类型，且数据类型需要是self与other推导之后可转换的数据类型，shape需要是self与other
35+ * broadcast之后的shape，数据格式支持ND，且数据格式需要与self一致。
36+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
37+ * @param [out] executor: 返回op执行器，包含算子计算流程。
38+ * @return aclnnStatus: 返回状态码。
39+ */
40+ACLNN_API aclnnStatus aclnnDivGetWorkspaceSize(
41+ const aclTensor* self, const aclTensor* other, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);
42+ 
43+/**
44+ * @brief aclnnDivs的第一段接口，根据具体的计算流程，计算workspace大小。
45+ * @domain aclnn_math
46+ */
47+ACLNN_API aclnnStatus aclnnDivsGetWorkspaceSize(
48+ const aclTensor* self, const aclScalar* other, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);
49+ 
50+/**
51+ * @brief aclnnDivMod的第一段接口，根据具体的计算流程，计算workspace大小。
52+ * @domain aclnn_math
53+ */
54+ACLNN_API aclnnStatus aclnnDivModGetWorkspaceSize(
55+ const aclTensor* self, const aclTensor* other, int mode, aclTensor* out, uint64_t* workspaceSize,
56+ aclOpExecutor** executor);
57+ 
58+/**
59+ * @brief aclnnDivMods的第一段接口，根据具体的计算流程，计算workspace大小。
60+ * @domain aclnn_math
61+ */
62+ACLNN_API aclnnStatus aclnnDivModsGetWorkspaceSize(
63+ const aclTensor* self, const aclScalar* other, int mode, aclTensor* out, uint64_t* workspaceSize,
64+ aclOpExecutor** executor);
65+ 
66+/**
67+ * @brief aclnnInplaceDiv的第一段接口，根据具体的计算流程，计算workspace大小。
68+ * @domain aclnn_math
69+ */
70+ACLNN_API aclnnStatus aclnnInplaceDivGetWorkspaceSize(
71+ aclTensor* selfRef, const aclTensor* other, uint64_t* workspaceSize, aclOpExecutor** executor);
72+ 
73+/**
74+ * @brief aclnnInplaceDivs的第一段接口，根据具体的计算流程，计算workspace大小。
75+ * @domain aclnn_math
76+ */
77+ACLNN_API aclnnStatus aclnnInplaceDivsGetWorkspaceSize(
78+ aclTensor* selfRef, const aclScalar* other, uint64_t* workspaceSize, aclOpExecutor** executor);
79+ 
80+/**
81+ * @brief aclnnInplaceDivMod的第一段接口，根据具体的计算流程，计算workspace大小。
82+ * @domain aclnn_math
83+ */
84+ACLNN_API aclnnStatus aclnnInplaceDivModGetWorkspaceSize(
85+ aclTensor* selfRef, const aclTensor* other, int mode, uint64_t* workspaceSize, aclOpExecutor** executor);
86+ 
87+/**
88+ * @brief aclnnInplaceDivMods的第一段接口，根据具体的计算流程，计算workspace大小。
89+ * @domain aclnn_math
90+ */
91+ACLNN_API aclnnStatus aclnnInplaceDivModsGetWorkspaceSize(
92+ aclTensor* selfRef, const aclScalar* other, int mode, uint64_t* workspaceSize, aclOpExecutor** executor);
93+ 
94+/**
95+ * @brief aclnnDiv的第二段接口，用于执行计算。
96+ *
97+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
98+ * @param [in] workspace_size: 在npu device侧申请的workspace大小，由第一段接口aclnnDivGetWorkspaceSize获取。
99+ * @param [in] executor: op执行器，包含了算子计算流程。
100+ * @param [in] stream: acl stream流。
101+ * @return aclnnStatus: 返回状态码。
102+ */
103+ACLNN_API aclnnStatus aclnnDiv(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
CANN-robot
CANN-robotCANN-robot1月4日

参数命名不一致: 函数aclnnDiv的第二个参数在文档中命名为'workspace_size'（带下划线），但在函数声明中为'workspaceSize'（驼峰命名）。这种命名不一致会影响代码的可读性和维护性，可能给开发者带来困惑。

问题类型: 参数命名不一致 文件路径: math/div/op_api/aclnn_div.h 行号: 103 问题代码:

ACLNN_API aclnnStatus aclnnDiv(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);

修改建议:

统一参数命名风格。建议将函数声明中的'workspaceSize'改为'workspace_size'，以与文档描述保持一致。同时，检查所有第二段接口函数（aclnnDivs、aclnnDivMod等）的参数命名，确保整个文件中的命名风格统一。

此评论由代码审查工具自动生成

likedislike
104+ 
105+ACLNN_API aclnnStatus aclnnDivs(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
106+ 
107+ACLNN_API aclnnStatus aclnnDivMod(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
108+ 
109+ACLNN_API aclnnStatus
110+aclnnDivMods(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
111+ 
112+ACLNN_API aclnnStatus
113+aclnnInplaceDiv(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
114+ 
115+ACLNN_API aclnnStatus
116+aclnnInplaceDivs(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
117+ 
118+ACLNN_API aclnnStatus
119+aclnnInplaceDivMod(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
120+ 
121+ACLNN_API aclnnStatus
122+aclnnInplaceDivMods(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
123+ 
124+#ifdef __cplusplus
125+}
126+#endif
127+ 
128+#endif // OP_API_INC_DIV_H_
Amath/div/op_api/div.cpp+109-0
@@ -0,0 +1,109 @@
1+/**
2+ * Copyright (c) 2025 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 "div.h"
12+#include "opdev/aicpu/aicpu_task.h"
13+#include "opdev/make_op_executor.h"
14+#include "opdev/op_def.h"
15+#include "opdev/op_dfx.h"
16+#include "opdev/op_executor.h"
17+#include "opdev/op_log.h"
18+#include "opdev/platform.h"
19+#include "opdev/shape_utils.h"
20+ 
21+using namespace op;
22+ 
23+namespace l0op {
24+ 
25+OP_TYPE_REGISTER(Div);
26+ 
27+static const std::initializer_list<op::DataType> ASCEND910_AICORE_DTYPE_SUPPORT_LIST = {
28+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_INT32, op::DataType::DT_INT8,
29+ op::DataType::DT_UINT8};
30+ 
31+static const std::initializer_list<op::DataType> ASCEND910B_AICORE_DTYPE_SUPPORT_LIST = {
32+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_INT32,
33+ op::DataType::DT_INT8, op::DataType::DT_UINT8, op::DataType::DT_BF16};
34+ 
35+static inline const std::initializer_list<op::DataType>& GetAiCoreDtypeSupportListBySocVersion()
36+{
37+ auto socVersion = GetCurrentPlatformInfo().GetSocVersion();
38+ switch (socVersion) {
39+ case SocVersion::ASCEND910B:
40+ case SocVersion::ASCEND910_95:
41+ case SocVersion::ASCEND910_93: {
42+ return ASCEND910B_AICORE_DTYPE_SUPPORT_LIST;
43+ }
44+ case SocVersion::ASCEND910: {
45+ return ASCEND910_AICORE_DTYPE_SUPPORT_LIST;
46+ }
47+ default: {
48+ return ASCEND910_AICORE_DTYPE_SUPPORT_LIST;
49+ }
50+ }
51+}
52+ 
53+// 根据芯片类型、dtype判断算子是否支持走aicore
54+static bool IsAiCoreSupport(const aclTensor* self)
55+{
56+ // Div只需要判断dtype
57+ return CheckType(self->GetDataType(), GetAiCoreDtypeSupportListBySocVersion());
58+}
59+ 
60+// AICORE算子kernel
61+static const aclTensor* DivAiCore(
62+ const aclTensor* self, const aclTensor* other, aclTensor* divOut, aclOpExecutor* executor)
63+{
64+ L0_DFX(DivAiCore, self, other, divOut);
65+ // 使用框架宏ADD_TO_LAUNCHER_LIST_AICORE，将AiCore Div算子加入任务队列
66+ // Div是算子的OpType，self、other是算子的输入，divOut是算子的输出
67+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(Div, OP_INPUT(self, other), OP_OUTPUT(divOut));
68+ OP_CHECK(
69+ ret == ACL_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "DivAiCore ADD_TO_LAUNCHER_LIST_AICORE failed."),
70+ return nullptr);
71+ return divOut;
72+}
73+ 
74+// AICPU算子kernel
75+static const aclTensor* DivAiCpu(
76+ const aclTensor* self, const aclTensor* other, aclTensor* divOut, aclOpExecutor* executor)
77+{
78+ // 使用框架宏ADD_TO_LAUNCHER_LIST_AICPU，将AiCpu Div算子加入任务队列
79+ // Div是算子的OpType，self、other是算子的输入，divOut是算子的输出
80+ L0_DFX(DivAiCpu, self, other);
81+ 
82+ static internal::AicpuTaskSpace space("Div");
83+ auto ret = ADD_TO_LAUNCHER_LIST_AICPU(Div, OP_ATTR_NAMES(), OP_INPUT(self, other), OP_OUTPUT(divOut));
84+ OP_CHECK(
85+ ret == ACL_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "DivAiCpu ADD_TO_LAUNCHER_LIST_AICPU failed."),
86+ return nullptr);
87+ return divOut;
88+}
89+ 
90+const aclTensor* Div(const aclTensor* self, const aclTensor* other, aclOpExecutor* executor)
91+{
92+ op::Shape broadcastShape;
93+ if (!BroadcastInferShape(self->GetViewShape(), other->GetViewShape(), broadcastShape)) {
94+ OP_LOGE(
95+ ACLNN_ERR_PARAM_INVALID, "Broadcast %s and %s failed.", op::ToString(self->GetViewShape()).GetString(),
96+ op::ToString(other->GetViewShape()).GetString());
97+ return nullptr;
98+ }
99+ 
100+ auto divOut = executor->AllocTensor(broadcastShape, self->GetDataType());
101+ if (IsAiCoreSupport(self)) {
102+ return DivAiCore(self, other, divOut, executor);
103+ } else {
104+ return DivAiCpu(self, other, divOut, executor);
105+ }
106+ return divOut;
107+}
108+ 
109+} // namespace l0op
Amath/div/op_api/div.h+21-0
@@ -0,0 +1,21 @@
1+/**
2+ * Copyright (c) 2025 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 PTA_NPU_OP_API_INC_LEVEL0_OP_DIV_OP_H_
12+#define PTA_NPU_OP_API_INC_LEVEL0_OP_DIV_OP_H_
13+ 
14+#include "opdev/op_executor.h"
15+ 
16+namespace l0op {
17+const aclTensor* Div(const aclTensor* self, const aclTensor* other, aclOpExecutor* executor);
18+ 
19+}
20+ 
21+#endif // PTA_NPU_OP_API_INC_LEVEL0_OP_ADD_OP_H_
Amath/div/op_graph/div_proto.h+56-0
@@ -0,0 +1,56 @@
1+/**
2+ * Copyright (c) 2025 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 div_proto.h
13+ * \brief
14+ */
15+#ifndef OPS_OP_PROTO_INC_DIV_H_
16+#define OPS_OP_PROTO_INC_DIV_H_
17+ 
18+#include "graph/operator_reg.h"
19+#include "graph/types.h"
20+ 
21+namespace ge {
22+ 
23+/**
24+* @brief Returns x1/x2 element-wise. Support broadcasting operations.
25+ 
26+* @par Inputs:
27+* Two inputs, including:
28+* @li x1: A ND Tensor. Must be one of the following types:
29+* float16, float32, int32, int8, uint8, float64, int64, uint16, int16,
30+* complex32, complex64, complex128, bfloat16, the format can be [NCHW,NHWC,ND].
31+* @li x2: A ND Tensor. Has the same dtype and format as input "x1". \n
32+ 
33+* @par Outputs:
34+* y: A ND Tensor. Has the same dtype and format as input "x1". \n
35+ 
36+* @par Third-party framework compatibility
37+* Compatible with the TensorFlow operator Div.
38+*/
39+REG_OP(Div)
40+ .INPUT(
CANN-robot
CANN-robotCANN-robot1月4日

代码结构与可维护性: 数据类型列表在INPUT(x1)、INPUT(x2)和OUTPUT(y)中重复了三次，存在代码重复问题。如果未来需要添加或删除支持的数据类型，需要修改三个地方，容易出错且维护困难。

问题类型: 代码结构与可维护性 文件路径: math/div/op_graph/div_proto.h 行号: 40 问题代码:

.INPUT(
        x1, TensorType(
                {DT_FLOAT, DT_FLOAT16, DT_INT8, DT_UINT8, DT_INT32, DT_DOUBLE, DT_INT64, DT_UINT16, DT_INT16,
                 DT_COMPLEX64, DT_COMPLEX128, DT_BF16, DT_COMPLEX32}))
    .INPUT(
        x2, TensorType(
                {DT_FLOAT, DT_FLOAT16, DT_INT8, DT_UINT8, DT_INT32, DT_DOUBLE, DT_INT64, DT_UINT16, DT_INT16,
                 DT_COMPLEX64, DT_COMPLEX128, DT_BF16, DT_COMPLEX32}))
    .OUTPUT(
        y, TensorType(
               {DT_FLOAT, DT_FLOAT16, DT_INT8, DT_UINT8, DT_INT32, DT_DOUBLE, DT_INT64, DT_UINT16, DT_INT16,
                DT_COMPLEX64, DT_COMPLEX128, DT_BF16, DT_COMPLEX32}))

修改建议:

建议将支持的数据类型列表定义为宏或常量，然后在三个地方引用该定义。例如：
#define DIV_SUPPORTED_TYPES {DT_FLOAT, DT_FLOAT16, DT_INT8, DT_UINT8, DT_INT32, DT_DOUBLE, DT_INT64, DT_UINT16, DT_INT16, DT_COMPLEX64, DT_COMPLEX128, DT_BF16, DT_COMPLEX32}
然后使用：.INPUT(x1, TensorType(DIV_SUPPORTED_TYPES))

此评论由代码审查工具自动生成

likedislike
41+ x1, TensorType(
42+ {DT_FLOAT, DT_FLOAT16, DT_INT8, DT_UINT8, DT_INT32, DT_DOUBLE, DT_INT64, DT_UINT16, DT_INT16,
43+ DT_COMPLEX64, DT_COMPLEX128, DT_BF16, DT_COMPLEX32}))
44+ .INPUT(
45+ x2, TensorType(
46+ {DT_FLOAT, DT_FLOAT16, DT_INT8, DT_UINT8, DT_INT32, DT_DOUBLE, DT_INT64, DT_UINT16, DT_INT16,
47+ DT_COMPLEX64, DT_COMPLEX128, DT_BF16, DT_COMPLEX32}))
48+ .OUTPUT(
49+ y, TensorType(
50+ {DT_FLOAT, DT_FLOAT16, DT_INT8, DT_UINT8, DT_INT32, DT_DOUBLE, DT_INT64, DT_UINT16, DT_INT16,
51+ DT_COMPLEX64, DT_COMPLEX128, DT_BF16, DT_COMPLEX32}))
52+ .OP_END_FACTORY_REG(Div)
53+ 
54+} // namespace ge
55+ 
56+#endif // OPS_OP_PROTO_INC_DIV_H_
Amath/div/op_graph/fusion_pass/.gitkeep+0-0
The file is empty
Dmath/div/op_host/CMakeLists.txt+0-12
@@ -1,12 +0,0 @@
1-# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6-# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9-# See LICENSE in the root of the software repository for the full text of the License.
10-# ----------------------------------------------------------------------------
11- 
12-add_modules_sources(OPTYPE div ACLNNTYPE aclnn_exclude)
Amath/div/op_host/arch35/div_tiling_arch35.cpp+143-0
@@ -0,0 +1,143 @@
1+/**
2+ * Copyright (c) 2025 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 div_tiling_arch35.cc
13+ * \brief div_tiling source file
14+ */
15+ 
16+#include <graph/utils/type_utils.h>
17+#include "register/op_impl_registry.h"
18+#include "tiling_base/tiling_templates_registry.h"
19+#include "atvoss/broadcast/broadcast_tiling.h"
20+#include "math/div/op_kernel/arch35/div_dag.h"
21+#include "math/div/op_kernel/arch35/div_struct.h"
22+#include "div_tiling_arch35.h"
23+ 
24+using namespace Ops::Base;
25+using namespace ge;
26+ 
27+namespace optiling {
28+ 
29+constexpr static uint64_t DIV_COMMON_TILING_PRIORITY = 0;
30+ 
31+ge::graphStatus DivTiling::GetShapeAttrsInfo()
32+{
33+ return ge::GRAPH_SUCCESS;
34+}
35+ 
36+bool DivTiling::IsCapable()
37+{
38+ return true;
39+}
40+ 
41+ge::graphStatus DivTiling::DoOpTiling()
42+{
43+ auto input0Desc = context_->GetInputDesc(0);
44+ OP_CHECK_NULL_WITH_CONTEXT(context_, input0Desc);
45+ ge::DataType input0DType = input0Desc->GetDataType();
46+ auto input1Desc = context_->GetInputDesc(1);
47+ OP_CHECK_NULL_WITH_CONTEXT(context_, input1Desc);
48+ ge::DataType input1DType = input1Desc->GetDataType();
49+ if (input0DType != input1DType) {
50+ OP_LOGE(
51+ context_->GetNodeName(), "dtype of input0[%s] != dtype of input1[%s].",
52+ ge::TypeUtils::DataTypeToSerialString(input0DType).c_str(),
53+ ge::TypeUtils::DataTypeToSerialString(input1DType).c_str());
54+ return ge::GRAPH_FAILED;
55+ }
56+ ge::graphStatus ret = ge::GRAPH_SUCCESS;
57+ if (input0DType == ge::DT_INT32) {
58+ BroadcastBaseTiling<DivOp::DivIntegerS32<int32_t>::OpDag> brcBaseTiling(context_);
59+ ret = brcBaseTiling.DoTiling();
60+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());
61+ } else if (input0DType == ge::DT_FLOAT16 || input0DType == ge::DT_BF16) {
62+ BroadcastBaseTiling<DivOp::DivFloatWithCast<half, float>::OpDag> brcBaseTiling(context_);
63+ ret = brcBaseTiling.DoTiling();
64+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());
65+ } else if (input0DType == ge::DT_FLOAT) {
66+ BroadcastBaseTiling<DivOp::DivFloatWithoutCast<float>::OpDag> brcBaseTiling(context_);
67+ ret = brcBaseTiling.DoTiling();
68+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());
69+ } else if (input0DType == ge::DT_UINT8) {
70+ BroadcastBaseTiling<DivOp::DivIntegerU8<uint8_t, uint16_t>::OpDag> brcBaseTiling(context_);
71+ ret = brcBaseTiling.DoTiling();
72+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());
73+ } else if (input0DType == ge::DT_INT8) {
74+ BroadcastBaseTiling<DivOp::DivIntegerS8<int8_t, int16_t>::OpDag> brcBaseTiling(context_);
75+ ret = brcBaseTiling.DoTiling();
76+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());
77+ } else if (input0DType == ge::DT_COMPLEX32) {
78+ BroadcastBaseTiling<DivOp::DivComplexWithoutCast<int32_t>::OpDag> brcBaseTiling(context_);
79+ ret = brcBaseTiling.DoTiling();
80+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());
81+ } else if (input0DType == ge::DT_COMPLEX64) {
82+ BroadcastBaseTiling<DivOp::DivComplexWithoutCast<int64_t>::OpDag> brcBaseTiling(context_);
83+ ret = brcBaseTiling.DoTiling();
84+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());
85+ } else {
86+ OP_LOGE(
87+ context_->GetNodeName(),
88+ "input dtype is only support fp16, bf16, fp32, int32, uint8, int8, complex32, complex64, but got %s!",
89+ ge::TypeUtils::DataTypeToSerialString(input0DType).c_str());
90+ return ge::GRAPH_FAILED;
91+ }
92+ return ret;
93+}
94+ 
95+ge::graphStatus DivTiling::DoLibApiTiling()
96+{
97+ return ge::GRAPH_SUCCESS;
98+}
99+ 
100+uint64_t DivTiling::GetTilingKey() const
101+{
102+ return tilingKey;
103+}
104+ 
105+ge::graphStatus DivTiling::GetWorkspaceSize()
106+{
107+ return ge::GRAPH_SUCCESS;
108+}
109+ 
110+ge::graphStatus DivTiling::PostTiling()
111+{
112+ return ge::GRAPH_SUCCESS;
113+}
114+ 
115+ge::graphStatus DivTiling::GetPlatformInfo()
116+{
117+ return ge::GRAPH_SUCCESS;
118+}
119+ 
120+ge::graphStatus TilingForDiv(gert::TilingContext* context)
121+{
122+ OP_LOGD("DivTiling", "Enter TilingForDiv");
123+ if (context == nullptr) {
124+ OP_LOGE("DivTiling", "Tiling context is nullptr");
125+ return ge::GRAPH_FAILED;
126+ }
127+ 
128+ auto compileInfo = reinterpret_cast<const BroadcastCompileInfo*>(context->GetCompileInfo());
129+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
130+ 
131+ OP_LOGD(context, "Enter ascendc DivTiling");
132+ return Ops::Math::OpTiling::TilingRegistry::GetInstance().DoTilingImpl(context);
133+}
134+ 
135+ge::graphStatus TilingPrepareForDiv([[maybe_unused]] gert::TilingParseContext* context)
136+{
137+ return ge::GRAPH_SUCCESS;
138+}
139+ 
140+IMPL_OP_OPTILING(Div).Tiling(TilingForDiv).TilingParse<BroadcastCompileInfo>(TilingPrepareForDiv);
141+ 
142+REGISTER_OPS_TILING_TEMPLATE(Div, DivTiling, DIV_COMMON_TILING_PRIORITY);
143+} // namespace optiling
Amath/div/op_host/arch35/div_tiling_arch35.h+49-0
@@ -0,0 +1,49 @@
1+/**
2+ * Copyright (c) 2025 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 div_tiling_arch35.h
13+ * \brief div_tiling head file
14+ */
15+ 
16+#ifndef OPS_BUILD_IN_OP_TILING_RUNTIME_DIV_TILING_H
17+#define OPS_BUILD_IN_OP_TILING_RUNTIME_DIV_TILING_H
18+ 
19+#include "log/log.h"
20+#include "platform/platform_info.h"
21+#include "register/op_def_registry.h"
22+#include "register/op_impl_registry.h"
23+#include "tiling/tiling_api.h"
24+#include "tiling_base/tiling_base.h"
25+ 
26+namespace optiling {
27+ 
28+class DivTiling : public Ops::Math::OpTiling::TilingBaseClass {
29+public:
30+ explicit DivTiling(gert::TilingContext* context) : Ops::Math::OpTiling::TilingBaseClass(context)
31+ {}
32+ 
33+protected:
34+ bool IsCapable() override;
35+ ge::graphStatus DoLibApiTiling() override;
36+ ge::graphStatus DoOpTiling() override;
37+ ge::graphStatus PostTiling() override;
38+ uint64_t GetTilingKey() const override;
39+ ge::graphStatus GetShapeAttrsInfo() override;
40+ ge::graphStatus GetPlatformInfo() override;
41+ ge::graphStatus GetWorkspaceSize() override;
42+ 
43+private:
44+ uint64_t tilingKey = 0;
45+};
46+ 
47+} // namespace optiling
48+ 
49+#endif // OPS_BUILD_IN_OP_TILING_RUNTIME_DIV_TILING_H
Amath/div/op_host/config/ascend910_95/div_binary.json+301-0
@@ -0,0 +1,301 @@
1+{
2+ "op_type": "Div",
3+ "op_list": [
4+ {
5+ "bin_filename": "Div_FLOAT16",
6+ "inputs": [
7+ {
8+ "name": "x1",
9+ "index": 0,
10+ "dtype": "float16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ]
16+ },
17+ {
18+ "name": "x2",
19+ "index": 1,
20+ "dtype": "float16",
21+ "format": "ND",
22+ "paramType": "required",
23+ "shape": [
24+ -2
25+ ]
26+ }
27+ ],
28+ "outputs": [
29+ {
30+ "name": "y",
31+ "index": 0,
32+ "dtype": "float16",
33+ "format": "ND",
34+ "paramType": "required",
35+ "shape": [
36+ -2
37+ ]
38+ }
39+ ]
40+ },
41+ {
42+ "bin_filename": "Div_FLOAT32",
43+ "inputs": [
44+ {
45+ "name": "x1",
46+ "index": 0,
47+ "dtype": "float32",
48+ "format": "ND",
49+ "paramType": "required",
50+ "shape": [
51+ -2
52+ ]
53+ },
54+ {
55+ "name": "x2",
56+ "index": 1,
57+ "dtype": "float32",
58+ "format": "ND",
59+ "paramType": "required",
60+ "shape": [
61+ -2
62+ ]
63+ }
64+ ],
65+ "outputs": [
66+ {
67+ "name": "y",
68+ "index": 0,
69+ "dtype": "float32",
70+ "format": "ND",
71+ "paramType": "required",
72+ "shape": [
73+ -2
74+ ]
75+ }
76+ ]
77+ },
78+ {
79+ "bin_filename": "Div_INT32",
80+ "inputs": [
81+ {
82+ "name": "x1",
83+ "index": 0,
84+ "dtype": "int32",
85+ "format": "ND",
86+ "paramType": "required",
87+ "shape": [
88+ -2
89+ ]
90+ },
91+ {
92+ "name": "x2",
93+ "index": 1,
94+ "dtype": "int32",
95+ "format": "ND",
96+ "paramType": "required",
97+ "shape": [
98+ -2
99+ ]
100+ }
101+ ],
102+ "outputs": [
103+ {
104+ "name": "y",
105+ "index": 0,
106+ "dtype": "int32",
107+ "format": "ND",
108+ "paramType": "required",
109+ "shape": [
110+ -2
111+ ]
112+ }
113+ ]
114+ },
115+ {
116+ "bin_filename": "Div_INT8",
117+ "inputs": [
118+ {
119+ "name": "x1",
120+ "index": 0,
121+ "dtype": "int8",
122+ "format": "ND",
123+ "paramType": "required",
124+ "shape": [
125+ -2
126+ ]
127+ },
128+ {
129+ "name": "x2",
130+ "index": 1,
131+ "dtype": "int8",
132+ "format": "ND",
133+ "paramType": "required",
134+ "shape": [
135+ -2
136+ ]
137+ }
138+ ],
139+ "outputs": [
140+ {
141+ "name": "y",
142+ "index": 0,
143+ "dtype": "int8",
144+ "format": "ND",
145+ "paramType": "required",
146+ "shape": [
147+ -2
148+ ]
149+ }
150+ ]
151+ },
152+ {
153+ "bin_filename": "Div_UINT8",
154+ "inputs": [
155+ {
156+ "name": "x1",
157+ "index": 0,
158+ "dtype": "uint8",
159+ "format": "ND",
160+ "paramType": "required",
161+ "shape": [
162+ -2
163+ ]
164+ },
165+ {
166+ "name": "x2",
167+ "index": 1,
168+ "dtype": "uint8",
169+ "format": "ND",
170+ "paramType": "required",
171+ "shape": [
172+ -2
173+ ]
174+ }
175+ ],
176+ "outputs": [
177+ {
178+ "name": "y",
179+ "index": 0,
180+ "dtype": "uint8",
181+ "format": "ND",
182+ "paramType": "required",
183+ "shape": [
184+ -2
185+ ]
186+ }
187+ ]
188+ },
189+ {
190+ "bin_filename": "Div_BFLOAT16",
191+ "inputs": [
192+ {
193+ "name": "x1",
194+ "index": 0,
195+ "dtype": "bfloat16",
196+ "format": "ND",
197+ "paramType": "required",
198+ "shape": [
199+ -2
200+ ]
201+ },
202+ {
203+ "name": "x2",
204+ "index": 1,
205+ "dtype": "bfloat16",
206+ "format": "ND",
207+ "paramType": "required",
208+ "shape": [
209+ -2
210+ ]
211+ }
212+ ],
213+ "outputs": [
214+ {
215+ "name": "y",
216+ "index": 0,
217+ "dtype": "bfloat16",
218+ "format": "ND",
219+ "paramType": "required",
220+ "shape": [
221+ -2
222+ ]
223+ }
224+ ]
225+ },
226+ {
227+ "bin_filename": "Div_COMPLEX32",
228+ "inputs": [
229+ {
230+ "name": "x1",
231+ "index": 0,
232+ "dtype": "complex32",
233+ "format": "ND",
234+ "paramType": "required",
235+ "shape": [
236+ -2
237+ ]
238+ },
239+ {
240+ "name": "x2",
241+ "index": 1,
242+ "dtype": "complex32",
243+ "format": "ND",
244+ "paramType": "required",
245+ "shape": [
246+ -2
247+ ]
248+ }
249+ ],
250+ "outputs": [
251+ {
252+ "name": "y",
253+ "index": 0,
254+ "dtype": "complex32",
255+ "format": "ND",
256+ "paramType": "required",
257+ "shape": [
258+ -2
259+ ]
260+ }
261+ ]
262+ },
263+ {
264+ "bin_filename": "Div_COMPLEX64",
265+ "inputs": [
266+ {
267+ "name": "x1",
268+ "index": 0,
269+ "dtype": "complex64",
270+ "format": "ND",
271+ "paramType": "required",
272+ "shape": [
273+ -2
274+ ]
275+ },
276+ {
277+ "name": "x2",
278+ "index": 1,
279+ "dtype": "complex64",
280+ "format": "ND",
281+ "paramType": "required",
282+ "shape": [
283+ -2
284+ ]
285+ }
286+ ],
287+ "outputs": [
288+ {
289+ "name": "y",
290+ "index": 0,
291+ "dtype": "complex64",
292+ "format": "ND",
293+ "paramType": "required",
294+ "shape": [
295+ -2
296+ ]
297+ }
298+ ]
299+ }
300+ ]
301+}
Amath/div/op_host/config/ascend910_95/div_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascnedxx为可选mode，如果不同芯片有差异化要求时，需要配置；
6+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置：各个版本按default配置
8+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子，由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+[Div]
13+default=0
Amath/div/op_host/div_def.cpp+69-0
@@ -0,0 +1,69 @@
1+/**
2+ * Copyright (c) 2025 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 div_def.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class Div : public OpDef {
19+public:
20+ explicit Div(const char* name) : OpDef(name)
21+ {
22+ this->Input("x1")
23+ .ParamType(REQUIRED)
24+ .DataType(
25+ {ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_UINT8, ge::DT_INT8, ge::DT_COMPLEX32,
26+ ge::DT_COMPLEX64})
27+ .Format(
28+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
29+ ge::FORMAT_ND, ge::FORMAT_ND})
30+ .UnknownShapeFormat(
31+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
32+ ge::FORMAT_ND, ge::FORMAT_ND});
33+ this->Input("x2")
34+ .ParamType(REQUIRED)
35+ .DataType(
36+ {ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_UINT8, ge::DT_INT8, ge::DT_COMPLEX32,
37+ ge::DT_COMPLEX64})
38+ .Format(
39+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
40+ ge::FORMAT_ND, ge::FORMAT_ND})
41+ .UnknownShapeFormat(
42+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
43+ ge::FORMAT_ND, ge::FORMAT_ND});
44+ this->Output("y")
45+ .ParamType(REQUIRED)
46+ .DataType(
47+ {ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_UINT8, ge::DT_INT8, ge::DT_COMPLEX32,
48+ ge::DT_COMPLEX64})
49+ .Format(
50+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
51+ ge::FORMAT_ND, ge::FORMAT_ND})
52+ .UnknownShapeFormat(
53+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
54+ ge::FORMAT_ND, ge::FORMAT_ND});
55+ OpAICoreConfig aicoreConfig;
56+ aicoreConfig.DynamicCompileStaticFlag(true)
57+ .DynamicFormatFlag(false)
58+ .DynamicRankSupportFlag(true)
59+ .DynamicShapeSupportFlag(true)
60+ .NeedCheckSupportFlag(false)
61+ .PrecisionReduceFlag(true)
62+ .ExtendCfgInfo("opFile.value", "div_apt");
63+ this->AICore().AddConfig("ascend910_95", aicoreConfig);
64+ this->AICore().AddConfig("mc62cm12a", aicoreConfig);
65+ }
66+};
67+ 
68+OP_ADD(Div);
69+} // namespace ops
Amath/div/op_host/div_infershape.cpp+23-0
@@ -0,0 +1,23 @@
1+/**
2+ * Copyright (c) 2025 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 div_infershape.cpp
13+ * \brief
14+ */
15+#include "op_host/infershape_broadcast_util.h"
16+#include "register/op_impl_registry.h"
17+ 
18+using namespace ge;
19+namespace ops {
20+ 
21+IMPL_OP_INFERSHAPE(Div).InferShape(Ops::Base::InferShape4Broadcast);
22+ 
23+} // namespace ops
Mmath/div/op_host/op_api/aclnn_div.cpp+30-24
@@ -1,20 +1,20 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3+ * 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").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.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,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.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.8+ * See LICENSE in the root of the software repository for the full text of the License.
9- */9+ */
10 10 
11#include "aclnn_div.h"11#include "aclnn_div.h"
12#include "aclnn_kernels/cast.h"12#include "aclnn_kernels/cast.h"
13#include "aclnn_kernels/contiguous.h"13#include "aclnn_kernels/contiguous.h"
14-#include "math/floor_div/op_api/floordiv.h"14+#include "math/floor_div/op_host/op_api/floordiv.h"
15-#include "math/real_div/op_host/op_api/realdiv.h"15+#include "../../../real_div/op_api/realdiv.h"
16#include "math/trunc/op_host/op_api/trunc.h"16#include "math/trunc/op_host/op_api/trunc.h"
17-#include "math/muls/op_api/muls.h"17+#include "math/muls/op_host/op_api/muls.h"
18#include "common/op_api_def.h"18#include "common/op_api_def.h"
19#include "aclnn_kernels/common/op_error_check.h"19#include "aclnn_kernels/common/op_error_check.h"
20#include "opdev/common_types.h"20#include "opdev/common_types.h"
@@ -32,8 +32,8 @@ using namespace op;
32extern "C" {32extern "C" {
33#endif33#endif
34 34 
35-static op::DataType PromoteIntegerAndBoolInputsToFloat(const op::DataType input) {35+static op::DataType PromoteIntegerInputsToFloat(const op::DataType input) {
36- if (IsIntegralType(input) || (input == op::DataType::DT_BOOL)) {36+ if (IsIntegralType(input)) {
37 return op::DataType::DT_FLOAT;37 return op::DataType::DT_FLOAT;
38 }38 }
39 return input;39 return input;
@@ -141,9 +141,9 @@ static inline op::DataType InferDivModeDtype(const op::DataType selfDtype, const
141 const int mode) {141 const int mode) {
142 auto promoteType = op::PromoteType(selfDtype, otherDtype);142 auto promoteType = op::PromoteType(selfDtype, otherDtype);
143 // 下沉PTA入口操作将入参类型转化成FLOAT进行后续处理143 // 下沉PTA入口操作将入参类型转化成FLOAT进行后续处理
144- if (mode == MODE_REAL_DIV) {144+ if (mode == MODE_REAL_DIV && promoteType != op::DataType::DT_INT32 && promoteType != op::DataType::DT_BOOL) {
145 // IterateBase 配置特殊处理145 // IterateBase 配置特殊处理
146- promoteType = PromoteIntegerAndBoolInputsToFloat(promoteType);146+ promoteType = PromoteIntegerInputsToFloat(promoteType);
147 }147 }
148 if (mode == MODE_TRUNC_DIV && promoteType == DataType::DT_DOUBLE) {148 if (mode == MODE_TRUNC_DIV && promoteType == DataType::DT_DOUBLE) {
149 promoteType = DataType::DT_FLOAT;149 promoteType = DataType::DT_FLOAT;
@@ -205,10 +205,9 @@ static inline op::DataType InferDivsModeDtype(const op::DataType selfDtype, cons
205 const int mode) {205 const int mode) {
206 auto scalarDefaultDtype = GetScalarDefaultDtype(otherDtype);206 auto scalarDefaultDtype = GetScalarDefaultDtype(otherDtype);
207 auto promoteType = CombineCategoriesWithComplex(selfDtype, scalarDefaultDtype);207 auto promoteType = CombineCategoriesWithComplex(selfDtype, scalarDefaultDtype);
208- 208+ if (mode == MODE_REAL_DIV && promoteType != op::DataType::DT_INT32 && promoteType != op::DataType::DT_BOOL) {
209- if (mode == MODE_REAL_DIV) {
210 // IterateBase 配置特殊处理209 // IterateBase 配置特殊处理
211- promoteType = PromoteIntegerAndBoolInputsToFloat(promoteType);210+ promoteType = PromoteIntegerInputsToFloat(promoteType);
212 }211 }
213 212 
214 if (mode == MODE_TRUNC_DIV && promoteType == DataType::DT_DOUBLE) {213 if (mode == MODE_TRUNC_DIV && promoteType == DataType::DT_DOUBLE) {
@@ -369,11 +368,9 @@ aclnnStatus aclnnDivGetWorkspaceSize(const aclTensor* self, const aclTensor* oth
369 CompatibleInferDivDtype(self->GetDataType(), other->GetDataType()) :368 CompatibleInferDivDtype(self->GetDataType(), other->GetDataType()) :
370 InferDivModeDtype(self->GetDataType(), other->GetDataType(), MODE_REAL_DIV);369 InferDivModeDtype(self->GetDataType(), other->GetDataType(), MODE_REAL_DIV);
371 370 
372- bool isSupportNonContiguous = GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_95;
373- 
374 // 处理self输入371 // 处理self输入
375 const aclTensor* selfProcessed = nullptr;372 const aclTensor* selfProcessed = nullptr;
376- if (self->GetDataType() == promoteType && isSupportNonContiguous) {373+ if (self->GetDataType() == promoteType && l0op::IsRealDivSupportNonContiguous(self)) {
377 selfProcessed = uniqueExecutor.get()->CreateView(374 selfProcessed = uniqueExecutor.get()->CreateView(
378 self, self->GetViewShape(), self->GetStorageShape(), self->GetViewStrides(), self->GetViewOffset());375 self, self->GetViewShape(), self->GetStorageShape(), self->GetViewStrides(), self->GetViewOffset());
379 } else {376 } else {
@@ -388,7 +385,7 @@ aclnnStatus aclnnDivGetWorkspaceSize(const aclTensor* self, const aclTensor* oth
388 385 
389 // 处理other输入386 // 处理other输入
390 const aclTensor* otherProcessed = nullptr;387 const aclTensor* otherProcessed = nullptr;
391- if (other->GetDataType() == promoteType && isSupportNonContiguous) {388+ if (other->GetDataType() == promoteType && l0op::IsRealDivSupportNonContiguous(other)) {
392 otherProcessed = uniqueExecutor.get()->CreateView(389 otherProcessed = uniqueExecutor.get()->CreateView(
393 other, other->GetViewShape(), other->GetStorageShape(), other->GetViewStrides(), other->GetViewOffset());390 other, other->GetViewShape(), other->GetStorageShape(), other->GetViewStrides(), other->GetViewOffset());
394 } else {391 } else {
@@ -485,7 +482,10 @@ static bool CanUseMuls(const aclTensor* self, const aclScalar* other) {
485 return false;482 return false;
486 }483 }
487 if (other->GetDataType() != op::DataType::DT_FLOAT16 && other->GetDataType() != op::DataType::DT_BF16 &&484 if (other->GetDataType() != op::DataType::DT_FLOAT16 && other->GetDataType() != op::DataType::DT_BF16 &&
488- other->GetDataType() != op::DataType::DT_FLOAT) {485+ other->GetDataType() != op::DataType::DT_FLOAT && other->GetDataType() != op::DataType::DT_DOUBLE) {
486+ return false;
487+ }
488+ if (!op::IsContiguous(self) && other->GetDataType() == op::DataType::DT_DOUBLE) {
489 return false;489 return false;
490 }490 }
491 491 
@@ -538,8 +538,14 @@ aclnnStatus aclnnDivsGetWorkspaceSize(const aclTensor* self, const aclScalar* ot
538 promoteType = (IsComplexType(other->GetDataType()))538 promoteType = (IsComplexType(other->GetDataType()))
539 ? op::PromoteType(promoteType, other->GetDataType())539 ? op::PromoteType(promoteType, other->GetDataType())
540 : promoteType;540 : promoteType;
541+ if (GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_95) {
542+ promoteType = op::PromoteType(self->GetDataType(), other->GetDataType()) == op::DataType::DT_INT32
543+ ? op::DataType::DT_INT32
544+ : promoteType;
545+ }
546+ 
541 bool canUseMuls = CanUseMuls(self, other);547 bool canUseMuls = CanUseMuls(self, other);
542- if (self->GetDataType() == promoteType && isSupportNonContiguous && !canUseMuls) {548+ if (self->GetDataType() == promoteType && l0op::IsRealDivSupportNonContiguous(self) && !canUseMuls) {
543 // aclScalar转aclTensor549 // aclScalar转aclTensor
544 auto otherConvert = uniqueExecutor.get()->ConvertToTensor(other, promoteType);550 auto otherConvert = uniqueExecutor.get()->ConvertToTensor(other, promoteType);
545 CHECK_RET(otherConvert != nullptr, ACLNN_ERR_INNER_NULLPTR);551 CHECK_RET(otherConvert != nullptr, ACLNN_ERR_INNER_NULLPTR);
Mmath/div/op_host/op_api/aclnn_div.h+37-35
@@ -1,12 +1,12 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3+ * 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").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.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,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.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.8+ * See LICENSE in the root of the software repository for the full text of the License.
9- */9+ */
10 10 
11#ifndef OP_API_INC_DIV_H_11#ifndef OP_API_INC_DIV_H_
12#define OP_API_INC_DIV_H_12#define OP_API_INC_DIV_H_
@@ -37,57 +37,59 @@ extern "C" {
37 * @param [out] executor: 返回op执行器，包含算子计算流程。37 * @param [out] executor: 返回op执行器，包含算子计算流程。
38 * @return aclnnStatus: 返回状态码。38 * @return aclnnStatus: 返回状态码。
39 */39 */
40-ACLNN_API aclnnStatus aclnnDivGetWorkspaceSize(const aclTensor* self, const aclTensor* other, aclTensor* out,40+ACLNN_API aclnnStatus aclnnDivGetWorkspaceSize(
41- uint64_t* workspaceSize, aclOpExecutor** executor);41+ const aclTensor* self, const aclTensor* other, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);
42 42 
43/**43/**
44 * @brief aclnnDivs的第一段接口，根据具体的计算流程，计算workspace大小。44 * @brief aclnnDivs的第一段接口，根据具体的计算流程，计算workspace大小。
45 * @domain aclnn_math45 * @domain aclnn_math
46 */46 */
47-ACLNN_API aclnnStatus aclnnDivsGetWorkspaceSize(const aclTensor* self, const aclScalar* other, aclTensor* out,47+ACLNN_API aclnnStatus aclnnDivsGetWorkspaceSize(
48- uint64_t* workspaceSize, aclOpExecutor** executor);48+ const aclTensor* self, const aclScalar* other, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);
49 49 
50/**50/**
51 * @brief aclnnDivMod的第一段接口，根据具体的计算流程，计算workspace大小。51 * @brief aclnnDivMod的第一段接口，根据具体的计算流程，计算workspace大小。
52 * @domain aclnn_math52 * @domain aclnn_math
53 */53 */
54-ACLNN_API aclnnStatus aclnnDivModGetWorkspaceSize(const aclTensor* self, const aclTensor* other, int mode,54+ACLNN_API aclnnStatus aclnnDivModGetWorkspaceSize(
55- aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);55+ const aclTensor* self, const aclTensor* other, int mode, aclTensor* out, uint64_t* workspaceSize,
56+ aclOpExecutor** executor);
56 57 
57/**58/**
58 * @brief aclnnDivMods的第一段接口，根据具体的计算流程，计算workspace大小。59 * @brief aclnnDivMods的第一段接口，根据具体的计算流程，计算workspace大小。
59 * @domain aclnn_math60 * @domain aclnn_math
60 */61 */
61-ACLNN_API aclnnStatus aclnnDivModsGetWorkspaceSize(const aclTensor* self, const aclScalar* other, int mode,62+ACLNN_API aclnnStatus aclnnDivModsGetWorkspaceSize(
62- aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);63+ const aclTensor* self, const aclScalar* other, int mode, aclTensor* out, uint64_t* workspaceSize,
64+ aclOpExecutor** executor);
63 65 
64/**66/**
65 * @brief aclnnInplaceDiv的第一段接口，根据具体的计算流程，计算workspace大小。67 * @brief aclnnInplaceDiv的第一段接口，根据具体的计算流程，计算workspace大小。
66 * @domain aclnn_math68 * @domain aclnn_math
67 */69 */
68-ACLNN_API aclnnStatus aclnnInplaceDivGetWorkspaceSize(aclTensor* selfRef, const aclTensor* other,70+ACLNN_API aclnnStatus aclnnInplaceDivGetWorkspaceSize(
69- uint64_t* workspaceSize, aclOpExecutor** executor);71+ aclTensor* selfRef, const aclTensor* other, uint64_t* workspaceSize, aclOpExecutor** executor);
70 72 
71/**73/**
72 * @brief aclnnInplaceDivs的第一段接口，根据具体的计算流程，计算workspace大小。74 * @brief aclnnInplaceDivs的第一段接口，根据具体的计算流程，计算workspace大小。
73 * @domain aclnn_math75 * @domain aclnn_math
74 */76 */
75-ACLNN_API aclnnStatus aclnnInplaceDivsGetWorkspaceSize(aclTensor* selfRef, const aclScalar* other,77+ACLNN_API aclnnStatus aclnnInplaceDivsGetWorkspaceSize(
76- uint64_t* workspaceSize, aclOpExecutor** executor);78+ aclTensor* selfRef, const aclScalar* other, uint64_t* workspaceSize, aclOpExecutor** executor);
77 79 
78/**80/**
79 * @brief aclnnInplaceDivMod的第一段接口，根据具体的计算流程，计算workspace大小。81 * @brief aclnnInplaceDivMod的第一段接口，根据具体的计算流程，计算workspace大小。
80 * @domain aclnn_math82 * @domain aclnn_math
81 */83 */
82-ACLNN_API aclnnStatus aclnnInplaceDivModGetWorkspaceSize(aclTensor* selfRef, const aclTensor* other, int mode,84+ACLNN_API aclnnStatus aclnnInplaceDivModGetWorkspaceSize(
83- uint64_t* workspaceSize, aclOpExecutor** executor);85+ aclTensor* selfRef, const aclTensor* other, int mode, uint64_t* workspaceSize, aclOpExecutor** executor);
84 86 
85/**87/**
86 * @brief aclnnInplaceDivMods的第一段接口，根据具体的计算流程，计算workspace大小。88 * @brief aclnnInplaceDivMods的第一段接口，根据具体的计算流程，计算workspace大小。
87 * @domain aclnn_math89 * @domain aclnn_math
88 */90 */
89-ACLNN_API aclnnStatus aclnnInplaceDivModsGetWorkspaceSize(aclTensor* selfRef, const aclScalar* other, int mode,91+ACLNN_API aclnnStatus aclnnInplaceDivModsGetWorkspaceSize(
90- uint64_t* workspaceSize, aclOpExecutor** executor);92+ aclTensor* selfRef, const aclScalar* other, int mode, uint64_t* workspaceSize, aclOpExecutor** executor);
91 93 
92/**94/**
93 * @brief aclnnDiv的第二段接口，用于执行计算。95 * @brief aclnnDiv的第二段接口，用于执行计算。
@@ -104,23 +106,23 @@ ACLNN_API aclnnStatus aclnnDivs(void* workspace, uint64_t workspaceSize, aclOpEx
104 106 
105ACLNN_API aclnnStatus aclnnDivMod(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);107ACLNN_API aclnnStatus aclnnDivMod(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
106 108 
107-ACLNN_API aclnnStatus aclnnDivMods(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,109+ACLNN_API aclnnStatus
108- aclrtStream stream);110+aclnnDivMods(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
109 111 
110-ACLNN_API aclnnStatus aclnnInplaceDiv(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,112+ACLNN_API aclnnStatus
111- aclrtStream stream);113+aclnnInplaceDiv(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
112 114 
113-ACLNN_API aclnnStatus aclnnInplaceDivs(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,115+ACLNN_API aclnnStatus
114- aclrtStream stream);116+aclnnInplaceDivs(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
115 117 
116-ACLNN_API aclnnStatus aclnnInplaceDivMod(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,118+ACLNN_API aclnnStatus
117- aclrtStream stream);119+aclnnInplaceDivMod(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
118 120 
119-ACLNN_API aclnnStatus aclnnInplaceDivMods(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,121+ACLNN_API aclnnStatus
120- aclrtStream stream);122+aclnnInplaceDivMods(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
121 123 
122#ifdef __cplusplus124#ifdef __cplusplus
123}125}
124#endif126#endif
125 127 
126-#endif // OP_API_INC_DIV_H_128+#endif // OP_API_INC_DIV_H_
Mmath/div/op_host/op_api/div.cpp+8-8
@@ -1,12 +1,12 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3+ * 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").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.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,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.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.8+ * See LICENSE in the root of the software repository for the full text of the License.
9- */9+ */
10 10 
11#include "div.h"11#include "div.h"
12#include "opdev/aicpu/aicpu_task.h"12#include "opdev/aicpu/aicpu_task.h"
Mmath/div/op_host/op_api/div.h+10-10
@@ -1,12 +1,12 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3+ * 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").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.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,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.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.8+ * See LICENSE in the root of the software repository for the full text of the License.
9- */9+ */
10 10 
11#ifndef PTA_NPU_OP_API_INC_LEVEL0_OP_DIV_OP_H_11#ifndef PTA_NPU_OP_API_INC_LEVEL0_OP_DIV_OP_H_
12#define PTA_NPU_OP_API_INC_LEVEL0_OP_DIV_OP_H_12#define PTA_NPU_OP_API_INC_LEVEL0_OP_DIV_OP_H_
@@ -14,8 +14,8 @@
14#include "opdev/op_executor.h"14#include "opdev/op_executor.h"
15 15 
16namespace l0op {16namespace l0op {
17-const aclTensor *Div(const aclTensor *self, const aclTensor *other, aclOpExecutor *executor);17+const aclTensor* Div(const aclTensor* self, const aclTensor* other, aclOpExecutor* executor);
18 18 
19}19}
20 20 
21-#endif // PTA_NPU_OP_API_INC_LEVEL0_OP_ADD_OP_H_21+#endif // PTA_NPU_OP_API_INC_LEVEL0_OP_ADD_OP_H_
Amath/div/op_kernel/arch35/div_dag.h+174-0
@@ -0,0 +1,174 @@
1+/**
2+ * Copyright (c) 2025 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 div_dag.h
13+ * \brief div dag
14+ */
15+ 
16+#ifndef DIV_DAG_H
17+#define DIV_DAG_H
18+ 
19+#include "atvoss/util/dag.h"
20+#include "atvoss/util/vec.h"
21+#include "atvoss/util/placeholder.h"
22+ 
23+namespace DivOp {
24+using namespace Ops::Base;
25+constexpr int DIV_CAST_MODE_NONE = 0;
26+constexpr int DIV_CAST_MODE_RINT = 1;
27+constexpr int DIV_CMP_NE_MODE = 5;
28+constexpr int DIV_SEL_TENSOR_TENSOR_MODE = 2;
29+constexpr int8_t SAT_POS = 60;
30+constexpr int16_t DIV_B16_SIGN = -32768;
31+constexpr int32_t DIV_B32_SIGN = -2147483648;
32+ 
33+template <class R, class T, int roundMode>
34+struct CastOverFlow : public Vec::ElemwiseUnaryOP<R, T> {
35+ __aicore__ inline CastOverFlow(LocalTensor<R>& dst, LocalTensor<T>& src, const uint32_t& count)
36+ {
37+#ifdef __CCE_AICORE__
38+ SetCtrlSpr<SAT_POS, SAT_POS>(0);
39+ constexpr static MicroAPI::CastTrait castTrait3 = {
40+ MicroAPI::RegLayout::ZERO, MicroAPI::SatMode::NO_SAT, MicroAPI::MaskMergeMode::ZEROING,
41+ RoundMode::CAST_RINT};
42+ __VEC_SCOPE__
43+ {
44+ MicroAPI::RegTensor<T> vreg0;
45+ MicroAPI::RegTensor<R> vreg1;
46+ MicroAPI::MaskReg preg0;
47+ // sizeof(T) must > sizeof(R)
48+ uint32_t size = count;
49+ uint16_t vfLoopNum = (size + (VECTOR_REG_WIDTH / sizeof(T)) - 1) / (VECTOR_REG_WIDTH / sizeof(T));
50+ __local_mem__ T* bufferIn0Addr = (__local_mem__ T*)src.GetPhyAddr();
51+ __local_mem__ R* bufferOut0Addr = (__local_mem__ R*)dst.GetPhyAddr();
52+ for (uint16_t i = 0; i < vfLoopNum; i++) {
53+ preg0 = MicroAPI::UpdateMask<T>(size);
54+ MicroAPI::DataCopy<T, MicroAPI::LoadDist::DIST_NORM>(
55+ vreg0, bufferIn0Addr + i * (VECTOR_REG_WIDTH / sizeof(T)));
56+ MicroAPI::Cast<R, T, castTrait3>(vreg1, vreg0, preg0);
57+ MicroAPI::DataCopy<R, MicroAPI::StoreDist::DIST_PACK_B16>(
58+ bufferOut0Addr + i * (VECTOR_REG_WIDTH / sizeof(T)), vreg1, preg0);
59+ }
60+ }
61+ SetCtrlSpr<SAT_POS, SAT_POS>(1);
62+#endif
63+ }
64+};
65+ 
66+template <typename T1>
67+struct DivComplexWithoutCast {
68+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
69+ using InputX2 = Bind<Vec::CopyInBrc<T1>, Placeholder::In1<T1>>;
70+ using DivRes = Bind<Vec::Div<T1>, InputX1, InputX2>;
71+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, DivRes>;
72+ using Outputs = Elems<OpCopyOut>;
73+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
74+ using OpDag = DAGSch<Outputs, void, MemCfg>;
75+};
76+ 
77+template <typename T1>
78+struct DivFloatWithoutCast {
79+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
80+ using InputX2 = Bind<Vec::CopyInBrc<T1>, Placeholder::In1<T1>>;
81+ using DivRes = Bind<Vec::DivHighPrecision<T1>, InputX1, InputX2>;
82+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, DivRes>;
83+ using Outputs = Elems<OpCopyOut>;
84+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
85+ using OpDag = DAGSch<Outputs, void, MemCfg>;
86+};
87+ 
88+template <typename T1, typename T2>
89+struct DivFloatWithCast {
90+ // half and bfloat16
91+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
92+ using InputX2 = Bind<Vec::CopyInBrc<T1>, Placeholder::In1<T1>>;
93+ using CastX1 = Bind<Vec::Cast<T2, T1, DIV_CAST_MODE_NONE>, InputX1>;
94+ using CastX2 = Bind<Vec::Cast<T2, T1, DIV_CAST_MODE_NONE>, InputX2>;
95+ using DivRes = Bind<Vec::DivHighPrecision<T2>, CastX1, CastX2>;
96+ using CastOut = Bind<Vec::Cast<T1, T2, DIV_CAST_MODE_RINT>, DivRes>;
97+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, CastOut>;
98+ using Outputs = Elems<OpCopyOut>;
99+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
100+ using OpDag = DAGSch<Outputs, void, MemCfg>;
101+};
102+ 
103+template <typename T1, typename T2>
104+struct DivIntegerS8 {
105+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
106+ using InputX2 = Bind<Vec::CopyInBrc<T1>, Placeholder::In1<T1>>;
107+ using CastX1 = Bind<Vec::Cast<T2, T1, DIV_CAST_MODE_NONE>, InputX1>;
108+ using CastX2 = Bind<Vec::Cast<T2, T1, DIV_CAST_MODE_NONE>, InputX2>;
109+ using DivValue = Bind<Vec::Div<T2>, CastX1, CastX2>;
110+ using MulValue = Bind<Vec::Mul<T2>, CastX2, DivValue>;
111+ using SubValue = Bind<Vec::Sub<T2>, CastX1, MulValue>;
112+ using ConstZero = MAKE_CONST(T2, 0);
113+ using DupZero = Bind<Vec::Duplicate<T2>, ConstZero>;
114+ using RemMask = Bind<Vec::Compare<uint8_t, T2, DIV_CMP_NE_MODE>, SubValue, DupZero>;
115+ using ConstFlag = MAKE_CONST(T2, DIV_B16_SIGN);
116+ using DupFlag = Bind<Vec::Duplicate<T2>, ConstFlag>;
117+ using AndX1 = Bind<Vec::And<T2>, CastX1, DupFlag>;
118+ using AndX2 = Bind<Vec::And<T2>, CastX2, DupFlag>;
119+ using SignMask = Bind<Vec::Compare<uint8_t, T2, DIV_CMP_NE_MODE>, AndX1, AndX2>;
120+ using ResMask = Bind<Vec::And<uint8_t>, RemMask, SignMask>;
121+ using ConstOne = MAKE_CONST(T2, 1);
122+ using DupOne = Bind<Vec::Duplicate<T2>, ConstOne>;
123+ using SubValue1 = Bind<Vec::Sub<T2>, DivValue, DupOne>;
124+ using SelectRes = Bind<Vec::Select<uint8_t, T2, DIV_SEL_TENSOR_TENSOR_MODE>, ResMask, SubValue1, DivValue>;
125+ using CastOutHalf = Bind<Vec::Cast<half, T2, DIV_CAST_MODE_RINT>, SelectRes>;
126+ using CastOutInteger = Bind<CastOverFlow<T1, half, DIV_CAST_MODE_RINT>, CastOutHalf>;
127+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, CastOutInteger>;
128+ using Outputs = Elems<OpCopyOut>;
129+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
130+ using OpDag = DAGSch<Outputs, void, MemCfg>;
131+};
132+ 
133+template <typename T1, typename T2>
134+struct DivIntegerU8 {
135+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
136+ using InputX2 = Bind<Vec::CopyInBrc<T1>, Placeholder::In1<T1>>;
137+ using CastX1 = Bind<Vec::Cast<T2, T1, DIV_CAST_MODE_NONE>, InputX1>;
138+ using CastX2 = Bind<Vec::Cast<T2, T1, DIV_CAST_MODE_NONE>, InputX2>;
139+ using DivValue = Bind<Vec::Div<T2>, CastX1, CastX2>;
140+ using CastOut = Bind<Vec::Cast<T1, T2, DIV_CAST_MODE_NONE>, DivValue>;
141+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, CastOut>;
142+ using Outputs = Elems<OpCopyOut>;
143+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
144+ using OpDag = DAGSch<Outputs, void, MemCfg>;
145+};
146+ 
147+template <typename T1>
148+struct DivIntegerS32 {
149+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
150+ using InputX2 = Bind<Vec::CopyInBrc<T1>, Placeholder::In1<T1>>;
151+ using DivValue = Bind<Vec::Div<T1>, InputX1, InputX2>;
152+ using MulValue = Bind<Vec::Mul<T1>, InputX2, DivValue>;
153+ using SubValue = Bind<Vec::Sub<T1>, InputX1, MulValue>;
154+ using ConstZero = MAKE_CONST(T1, 0);
155+ using DupZero = Bind<Vec::Duplicate<T1>, ConstZero>;
156+ using RemMask = Bind<Vec::Compare<uint8_t, T1, DIV_CMP_NE_MODE>, SubValue, DupZero>;
157+ using ConstFlag = MAKE_CONST(T1, DIV_B32_SIGN);
158+ using DupFlag = Bind<Vec::Duplicate<T1>, ConstFlag>;
159+ using AndX1 = Bind<Vec::And<T1>, InputX1, DupFlag>;
160+ using AndX2 = Bind<Vec::And<T1>, InputX2, DupFlag>;
161+ using SignMask = Bind<Vec::Compare<uint8_t, T1, DIV_CMP_NE_MODE>, AndX1, AndX2>;
162+ using ResMask = Bind<Vec::And<uint8_t>, RemMask, SignMask>;
163+ using ConstOne = MAKE_CONST(T1, 1);
164+ using DupOne = Bind<Vec::Duplicate<T1>, ConstOne>;
165+ using SubValue1 = Bind<Vec::Sub<T1>, DivValue, DupOne>;
166+ using SelectRes = Bind<Vec::Select<uint8_t, T1, DIV_SEL_TENSOR_TENSOR_MODE>, ResMask, SubValue1, DivValue>;
167+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, SelectRes>;
168+ using Outputs = Elems<OpCopyOut>;
169+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
170+ using OpDag = DAGSch<Outputs, void, MemCfg>;
171+};
172+} // namespace DivOp
173+ 
174+#endif // DIV_DAG_H
Amath/div/op_kernel/arch35/div_struct.h+24-0
@@ -0,0 +1,24 @@
1+/**
2+ * Copyright (c) 2025 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 div_struct.h
13+ * \brief div_struct
14+ */
15+#ifndef DIV_STRUCT_H_
16+#define DIV_STRUCT_H_
17+ 
18+#include "atvoss/broadcast/broadcast_base_struct.h"
19+ 
20+// 算子自定义的tiling key字段
21+ASCENDC_TPL_ARGS_DECL(Div, BRC_TEMP_SCH_MODE_KEY_DECL(schMode));
22+ 
23+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(BRC_TEMP_SCH_MODE_KEY_SEL(schMode)));
24+#endif // DIV_STRUCT_H_
Amath/div/op_kernel/div_apt.cpp+60-0
@@ -0,0 +1,60 @@
1+/**
2+ * Copyright (c) 2025 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 div_apt.cpp
13+ * \brief div kernel
14+ */
15+ 
16+#include "kernel_operator.h"
17+#include "arch35/div_dag.h"
18+#include "arch35/div_struct.h"
19+#include "atvoss/broadcast/broadcast_sch.h"
20+ 
21+using namespace AscendC;
22+using namespace Ops::Base;
23+ 
24+template <uint64_t schMode>
25+__global__ __aicore__ void div(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
26+{
27+ if constexpr (std::is_same<DTYPE_X1, complex32>::value) {
28+ using OpDag = DivOp::DivComplexWithoutCast<complex32>::OpDag;
29+ BroadcastSch<schMode, OpDag> sch(tiling);
30+ sch.Process(x1, x2, y);
31+ } else if constexpr (std::is_same<DTYPE_X1, complex64>::value) {
32+ using OpDag = DivOp::DivComplexWithoutCast<complex64>::OpDag;
33+ BroadcastSch<schMode, OpDag> sch(tiling);
34+ sch.Process(x1, x2, y);
35+ } else if constexpr (std::is_same<DTYPE_X1, int8_t>::value) {
36+ using OpDag = DivOp::DivIntegerS8<int8_t, int16_t>::OpDag;
37+ BroadcastSch<schMode, OpDag> sch(tiling);
38+ sch.Process(x1, x2, y);
39+ } else if constexpr (std::is_same<DTYPE_X1, uint8_t>::value) {
40+ using OpDag = DivOp::DivIntegerU8<uint8_t, uint16_t>::OpDag;
41+ BroadcastSch<schMode, OpDag> sch(tiling);
42+ sch.Process(x1, x2, y);
43+ } else if constexpr (std::is_same<DTYPE_X1, int32_t>::value) {
44+ using OpDag = DivOp::DivIntegerS32<int32_t>::OpDag;
45+ BroadcastSch<schMode, OpDag> sch(tiling);
46+ sch.Process(x1, x2, y);
47+ } else if constexpr (std::is_same<DTYPE_X1, bfloat16_t>::value) {
48+ using OpDag = DivOp::DivFloatWithCast<bfloat16_t, float>::OpDag;
49+ BroadcastSch<schMode, OpDag> sch(tiling);
50+ sch.Process(x1, x2, y);
51+ } else if constexpr (std::is_same<DTYPE_X1, half>::value) {
52+ using OpDag = DivOp::DivFloatWithCast<half, float>::OpDag;
53+ BroadcastSch<schMode, OpDag> sch(tiling);
54+ sch.Process(x1, x2, y);
55+ } else {
56+ using OpDag = DivOp::DivFloatWithoutCast<DTYPE_X1>::OpDag;
57+ BroadcastSch<schMode, OpDag> sch(tiling);
58+ sch.Process(x1, x2, y);
59+ }
60+}
Mmath/div/tests/CMakeLists.txt+4-5
@@ -1,11 +1,10 @@
1# ----------------------------------------------------------------------------1# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").4+# CANN Open Software License Agreement Version 2.0 (the "License").
6# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
10# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
11 10 
Mmath/div/tests/ut/CMakeLists.txt+4-5
@@ -1,11 +1,10 @@
1# ----------------------------------------------------------------------------1# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").4+# CANN Open Software License Agreement Version 2.0 (the "License").
6# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
10# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
11 10 
Rmath/div/tests/ut/op_host/op_api/CMakeLists.txt→math/div/tests/ut/op_api/CMakeLists.txt+4-5
@@ -1,10 +1,9 @@
1# ----------------------------------------------------------------------------1# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").4+# CANN Open Software License Agreement Version 2.0 (the "License").
6# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
10# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
Rmath/div/tests/ut/op_host/op_api/test_aclnn_div.cpp→math/div/tests/ut/op_api/test_aclnn_div.cpp+36-28
@@ -1,18 +1,18 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3+ * 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").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.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,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.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.8+ * See LICENSE in the root of the software repository for the full text of the License.
9- */9+ */
10 10 
11#include <array>11#include <array>
12#include <vector>12#include <vector>
13#include "gtest/gtest.h"13#include "gtest/gtest.h"
14 14 
15-#include "aclnn_div.h"15+#include "math/div/op_api/aclnn_div.h"
16 16 
17#include "op_api_ut_common/inner/types.h"17#include "op_api_ut_common/inner/types.h"
18#include "op_api_ut_common/op_api_ut.h"18#include "op_api_ut_common/op_api_ut.h"
@@ -25,9 +25,15 @@ using namespace std;
25 25 
26class l2_div_test : public testing::Test {26class l2_div_test : public testing::Test {
27protected:27protected:
28- static void SetUpTestCase() { cout << "div_test SetUp" << endl; }28+ static void SetUpTestCase()
29+ {
30+ cout << "div_test SetUp" << endl;
31+ }
29 32 
30- static void TearDownTestCase() { cout << "div_test TearDown" << endl; }33+ static void TearDownTestCase()
34+ {
35+ cout << "div_test TearDown" << endl;
36+ }
31};37};
32 38 
33// 测试aicore:FLOAT,FLOAT32类型支持39// 测试aicore:FLOAT,FLOAT32类型支持
@@ -98,8 +104,8 @@ TEST_F(l2_div_test, case_bool_dtype_support)
98// 测试支持数据格式104// 测试支持数据格式
99TEST_F(l2_div_test, case_dtype_all_format)105TEST_F(l2_div_test, case_dtype_all_format)
100{106{
101- vector<aclFormat> format_list{ACL_FORMAT_ND, ACL_FORMAT_NCHW, ACL_FORMAT_NHWC, ACL_FORMAT_HWCN,107+ vector<aclFormat> format_list{ACL_FORMAT_ND, ACL_FORMAT_NCHW, ACL_FORMAT_NHWC,
102- ACL_FORMAT_NDHWC, ACL_FORMAT_NCDHW};108+ ACL_FORMAT_HWCN, ACL_FORMAT_NDHWC, ACL_FORMAT_NCDHW};
103 for (auto format : format_list) {109 for (auto format : format_list) {
104 auto self_tensor_desc = TensorDesc({3, 5}, ACL_FLOAT, format).ValueRange(10, 100);110 auto self_tensor_desc = TensorDesc({3, 5}, ACL_FLOAT, format).ValueRange(10, 100);
105 auto other_tensor_desc = TensorDesc({3, 5}, ACL_FLOAT, format).ValueRange(10, 100);111 auto other_tensor_desc = TensorDesc({3, 5}, ACL_FLOAT, format).ValueRange(10, 100);
@@ -323,26 +329,28 @@ TEST_F(l2_div_test, case_other_scalar_CheckDtypeValid)
323}329}
324 330 
325// 测试other为scalar超过8维的tensor331// 测试other为scalar超过8维的tensor
326-TEST_F(l2_div_test, case_scalar_shape_dim_9) {332+TEST_F(l2_div_test, case_scalar_shape_dim_9)
327- auto self_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);333+{
328- auto other_tensor_desc = ScalarDesc(2.0f);334+ auto self_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);
329- auto out_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);335+ auto other_tensor_desc = ScalarDesc(2.0f);
336+ auto out_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);
330 337 
331- auto ut = OP_API_UT(aclnnDivs, INPUT(self_tensor_desc, other_tensor_desc), OUTPUT(out_tensor_desc));338+ auto ut = OP_API_UT(aclnnDivs, INPUT(self_tensor_desc, other_tensor_desc), OUTPUT(out_tensor_desc));
332- uint64_t workspace_size = 0;339+ uint64_t workspace_size = 0;
333- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);340+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
334- EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);341+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
335}342}
336 343 
337// 测试超过8维的tensor344// 测试超过8维的tensor
338-TEST_F(l2_div_test, case_shape_dim_9) {345+TEST_F(l2_div_test, case_shape_dim_9)
339- auto self_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);346+{
340- auto out_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);347+ auto self_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);
348+ auto out_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);
341 349 
342- auto ut = OP_API_UT(aclnnDiv, INPUT(self_tensor_desc, self_tensor_desc), OUTPUT(out_tensor_desc));350+ auto ut = OP_API_UT(aclnnDiv, INPUT(self_tensor_desc, self_tensor_desc), OUTPUT(out_tensor_desc));
343- uint64_t workspace_size = 0;351+ uint64_t workspace_size = 0;
344- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);352+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
345- EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);353+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
346}354}
347 355 
348// 测试div_:other为scalar立即数输入356// 测试div_:other为scalar立即数输入
Rmath/div/tests/ut/op_host/op_api/test_aclnn_divMod.cpp→math/div/tests/ut/op_api/test_aclnn_divMod.cpp+50-40
@@ -1,39 +1,44 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3+ * 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").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.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,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.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.8+ * See LICENSE in the root of the software repository for the full text of the License.
9- */9+ */
10 10 
11#include <array>11#include <array>
12#include <vector>12#include <vector>
13#include "gtest/gtest.h"13#include "gtest/gtest.h"
14 14 
15-#include "aclnn_div.h"15+#include "math/div/op_api/aclnn_div.h"
16 16 
17#include "op_api_ut_common/inner/types.h"17#include "op_api_ut_common/inner/types.h"
18#include "op_api_ut_common/op_api_ut.h"18#include "op_api_ut_common/op_api_ut.h"
19#include "op_api_ut_common/scalar_desc.h"19#include "op_api_ut_common/scalar_desc.h"
20#include "op_api_ut_common/tensor_desc.h"20#include "op_api_ut_common/tensor_desc.h"
21 21 
22- 
23using namespace std;22using namespace std;
24 23 
25class l2_div_mod_test : public testing::Test {24class l2_div_mod_test : public testing::Test {
26protected:25protected:
27- static void SetUpTestCase() { cout << "div_test SetUp" << endl; }26+ static void SetUpTestCase()
27+ {
28+ cout << "div_test SetUp" << endl;
29+ }
28 30 
29- static void TearDownTestCase() { cout << "div_test TearDown" << endl; }31+ static void TearDownTestCase()
32+ {
33+ cout << "div_test TearDown" << endl;
34+ }
30};35};
31 36 
32// 测试所有数据格式支持37// 测试所有数据格式支持
33TEST_F(l2_div_mod_test, case_dtype_all_format)38TEST_F(l2_div_mod_test, case_dtype_all_format)
34{39{
35- vector<aclFormat> format_list{ACL_FORMAT_ND, ACL_FORMAT_NCHW, ACL_FORMAT_NHWC, ACL_FORMAT_HWCN,40+ vector<aclFormat> format_list{ACL_FORMAT_ND, ACL_FORMAT_NCHW, ACL_FORMAT_NHWC,
36- ACL_FORMAT_NDHWC, ACL_FORMAT_NCDHW};41+ ACL_FORMAT_HWCN, ACL_FORMAT_NDHWC, ACL_FORMAT_NCDHW};
37 for (auto format : format_list) {42 for (auto format : format_list) {
38 auto self_tensor_desc = TensorDesc({3, 5}, ACL_FLOAT, format).ValueRange(10, 100);43 auto self_tensor_desc = TensorDesc({3, 5}, ACL_FLOAT, format).ValueRange(10, 100);
39 auto other_tensor_desc = TensorDesc({3, 5}, ACL_FLOAT, format).ValueRange(10, 100);44 auto other_tensor_desc = TensorDesc({3, 5}, ACL_FLOAT, format).ValueRange(10, 100);
@@ -156,28 +161,30 @@ TEST_F(l2_div_mod_test, case_other_scalar_trunc_support)
156}161}
157 162 
158// 测试other为scalar超过8维的tensor163// 测试other为scalar超过8维的tensor
159-TEST_F(l2_div_mod_test, case_scalar_shape_dim_9) {164+TEST_F(l2_div_mod_test, case_scalar_shape_dim_9)
160- auto self_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);165+{
161- auto other_tensor_desc = ScalarDesc(2.0f);166+ auto self_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);
162- auto out_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);167+ auto other_tensor_desc = ScalarDesc(2.0f);
163- int mode = 2;168+ auto out_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);
169+ int mode = 2;
164 170 
165- auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));171+ auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));
166- uint64_t workspace_size = 0;172+ uint64_t workspace_size = 0;
167- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);173+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
168- EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);174+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
169}175}
170 176 
171// 测试超过8维的tensor177// 测试超过8维的tensor
172-TEST_F(l2_div_mod_test, case_shape_dim_9) {178+TEST_F(l2_div_mod_test, case_shape_dim_9)
173- auto self_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);179+{
174- auto out_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);180+ auto self_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);
175- int mode = 2;181+ auto out_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);
182+ int mode = 2;
176 183 
177- auto ut = OP_API_UT(aclnnDivMod, INPUT(self_tensor_desc, self_tensor_desc, mode), OUTPUT(out_tensor_desc));184+ auto ut = OP_API_UT(aclnnDivMod, INPUT(self_tensor_desc, self_tensor_desc, mode), OUTPUT(out_tensor_desc));
178- uint64_t workspace_size = 0;185+ uint64_t workspace_size = 0;
179- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);186+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
180- EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);187+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
181}188}
182 189 
183// 测试div_:other为scalar立即数输入190// 测试div_:other为scalar立即数输入
@@ -206,24 +213,26 @@ TEST_F(l2_div_mod_test, case_inplace_other_support)
206 EXPECT_EQ(aclRet, ACL_SUCCESS);213 EXPECT_EQ(aclRet, ACL_SUCCESS);
207}214}
208 215 
209-TEST_F(l2_div_mod_test, case_complex_mod_1_invalid) {216+TEST_F(l2_div_mod_test, case_complex_mod_1_invalid)
217+{
210 auto self_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).ValueRange(1, 2);218 auto self_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).ValueRange(1, 2);
211 auto other_tensor_desc = ScalarDesc(2.0f);219 auto other_tensor_desc = ScalarDesc(2.0f);
212 auto out_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).Precision(0.001, 0.001);220 auto out_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).Precision(0.001, 0.001);
213 int mode = 1;221 int mode = 1;
214- 222+ 
215 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));223 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));
216 uint64_t workspace_size = 0;224 uint64_t workspace_size = 0;
217 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);225 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
218 // EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);226 // EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
219}227}
220 228 
221-TEST_F(l2_div_mod_test, case_complex_mod_2_invalid) {229+TEST_F(l2_div_mod_test, case_complex_mod_2_invalid)
230+{
222 auto self_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).ValueRange(1, 2);231 auto self_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).ValueRange(1, 2);
223 auto other_tensor_desc = ScalarDesc(2.0f);232 auto other_tensor_desc = ScalarDesc(2.0f);
224 auto out_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).Precision(0.001, 0.001);233 auto out_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).Precision(0.001, 0.001);
225 int mode = 2;234 int mode = 2;
226- 235+ 
227 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));236 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));
228 uint64_t workspace_size = 0;237 uint64_t workspace_size = 0;
229 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);238 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
@@ -231,12 +240,13 @@ TEST_F(l2_div_mod_test, case_complex_mod_2_invalid) {
231}240}
232 241 
233// 测试mode异常值拦截242// 测试mode异常值拦截
234-TEST_F(l2_div_mod_test, case_mode_invalid) {243+TEST_F(l2_div_mod_test, case_mode_invalid)
244+{
235 auto self_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);245 auto self_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);
236 auto other_tensor_desc = ScalarDesc(2.0f);246 auto other_tensor_desc = ScalarDesc(2.0f);
237 auto out_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);247 auto out_tensor_desc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_NCHW);
238 int mode = 3;248 int mode = 3;
239- 249+ 
240 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));250 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));
241 uint64_t workspace_size = 0;251 uint64_t workspace_size = 0;
242 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);252 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
@@ -253,7 +263,7 @@ TEST_F(l2_div_mod_test, Ascend910_9589_case_real_div_dtype_modeNone)
253 auto other_tensor_desc = ScalarDesc(2);263 auto other_tensor_desc = ScalarDesc(2);
254 auto out_tensor_desc = TensorDesc({4, 5}, dtype, ACL_FORMAT_ND).Precision(0.001, 0.001);264 auto out_tensor_desc = TensorDesc({4, 5}, dtype, ACL_FORMAT_ND).Precision(0.001, 0.001);
255 int mode = 0;265 int mode = 0;
256- 266+ 
257 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));267 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));
258 uint64_t workspace_size = 0;268 uint64_t workspace_size = 0;
259 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);269 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
@@ -269,7 +279,7 @@ TEST_F(l2_div_mod_test, Ascend910_9589_case_divmods_trunc_div)
269 auto other_tensor_desc = ScalarDesc(2);279 auto other_tensor_desc = ScalarDesc(2);
270 auto out_tensor_desc = TensorDesc({4, 5}, dtype, ACL_FORMAT_ND).Precision(0.001, 0.001);280 auto out_tensor_desc = TensorDesc({4, 5}, dtype, ACL_FORMAT_ND).Precision(0.001, 0.001);
271 int mode = 1;281 int mode = 1;
272- 282+ 
273 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));283 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));
274 uint64_t workspace_size = 0;284 uint64_t workspace_size = 0;
275 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);285 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
@@ -285,7 +295,7 @@ TEST_F(l2_div_mod_test, Ascend910_9589_case_divmods_floor_div)
285 auto other_tensor_desc = ScalarDesc(2);295 auto other_tensor_desc = ScalarDesc(2);
286 auto out_tensor_desc = TensorDesc({4, 5}, dtype, ACL_FORMAT_ND).Precision(0.001, 0.001);296 auto out_tensor_desc = TensorDesc({4, 5}, dtype, ACL_FORMAT_ND).Precision(0.001, 0.001);
287 int mode = 2;297 int mode = 2;
288- 298+ 
289 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));299 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));
290 uint64_t workspace_size = 0;300 uint64_t workspace_size = 0;
291 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);301 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
Mmath/div/tests/ut/op_host/CMakeLists.txt+4-5
@@ -1,11 +1,10 @@
1# ----------------------------------------------------------------------------1# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").4+# CANN Open Software License Agreement Version 2.0 (the "License").
6# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
10# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
11 10 
Amath/div/tests/ut/op_host/arch35/test_div_tiling_arch35.cpp+81-0
@@ -0,0 +1,81 @@
1+ 
2+/**
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+ * CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file test_div_tiling_arch35.cpp
14+ * \brief
15+ */
16+ 
17+#include "math/div/op_host/arch35/div_tiling_arch35.h"
18+#include <iostream>
19+#include <gtest/gtest.h>
20+#include "tiling_context_faker.h"
21+#include "tiling_case_executor.h"
22+ 
23+using namespace std;
24+using namespace ge;
25+ 
26+struct DivCompileInfo{
27+ uint64_t coreNum;
28+ uint64_t ubSize;
29+};
30+ 
31+ 
32+class DivTiling : public testing::Test {
33+protected:
34+ static void SetUpTestCase()
35+ {
36+ std::cout << "DivTiling SetUp" << std::endl;
37+ }
38+ 
39+ static void TearDownTestCase()
40+ {
41+ std::cout << "DivTiling TearDown" << std::endl;
42+ }
43+};
44+ 
45+TEST_F(DivTiling, div_test_0)
46+{
47+ DivCompileInfo compileInfo = {64, 245760};
48+ gert::TilingContextPara tilingContextPara(
49+ "Div",
50+ {
51+ {{{5, 5, 64, 128}, {5, 5, 64, 128}}, ge::DT_BF16, ge::FORMAT_ND},
52+ {{{5, 5, 64, 128}, {5, 5, 64, 128}}, ge::DT_BF16, ge::FORMAT_ND},
53+ },
54+ {
55+ {{{5, 5, 64, 128}, {5, 5, 64, 128}}, ge::DT_BF16, ge::FORMAT_ND},
56+ },
57+ &compileInfo);
58+ uint64_t expectTilingKey = 8;
59+ string expectTilingData = "204800 137438956672 ";
60+ std::vector<size_t> expectWorkspaces = {16777216};
61+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
62+}
63+ 
64+TEST_F(DivTiling, div_test_1)
65+{
66+ DivCompileInfo compileInfo = {64, 245760};
67+ gert::TilingContextPara tilingContextPara(
68+ "Div",
69+ {
70+ {{{5, 5, 64, 128}, {5, 5, 64, 128}}, ge::DT_FLOAT16, ge::FORMAT_ND},
71+ {{{5, 5, 64, 128}, {5, 5, 64, 128}}, ge::DT_FLOAT16, ge::FORMAT_ND},
72+ },
73+ {
74+ {{{5, 5, 64, 128}, {5, 5, 64, 128}}, ge::DT_FLOAT16, ge::FORMAT_ND},
75+ },
76+ &compileInfo);
77+ uint64_t expectTilingKey = 8;
78+ string expectTilingData = "204800 137438956672 ";
79+ std::vector<size_t> expectWorkspaces = {16777216};
80+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
81+}
Amath/div/tests/ut/op_host/test_div_infershape.cpp+69-0
@@ -0,0 +1,69 @@
1+/**
2+ * Copyright (c) 2025 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_div_infershape.cpp
13+ * \brief
14+ */
15+ 
16+#include <iostream>
17+#include <gtest/gtest.h>
18+#include "infershape_context_faker.h"
19+#include "infershape_case_executor.h"
20+ 
21+using namespace ge;
22+ 
23+// ----------------DivInfershape-------------------
24+class DivInfershape : public testing::Test {
25+protected:
26+ static void SetUpTestCase()
27+ {
28+ std::cout << "DivInfershape SetUp" << std::endl;
29+ }
30+ 
31+ static void TearDownTestCase()
32+ {
33+ std::cout << "DivInfershape TearDown" << std::endl;
34+ }
35+};
36+ 
37+TEST_F(DivInfershape, div_infershape_test_0)
38+{
39+ gert::InfershapeContextPara infershapeContextPara(
40+ "Div",
41+ {
42+ {{{2, 2, 1}, {2, 2, 1}}, ge::DT_FLOAT16, ge::FORMAT_ND},
43+ {{{2, 2, 3}, {2, 2, 3}}, ge::DT_FLOAT16, ge::FORMAT_ND},
44+ },
45+ {
46+ {{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND},
47+ });
48+ std::vector<std::vector<int64_t>> expectOutputShape = {
49+ {2, 2, 3},
50+ };
51+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
52+}
53+ 
54+TEST_F(DivInfershape, div_infershape_test_1)
55+{
56+ gert::InfershapeContextPara infershapeContextPara(
57+ "Div",
58+ {
59+ {{{3, 4, 5, 6, -1}, {3, 4, 5, 6, -1}}, ge::DT_INT32, ge::FORMAT_ND},
60+ {{{3, 4, 5, 6, 1}, {3, 4, 5, 6, 1}}, ge::DT_INT32, ge::FORMAT_ND},
61+ },
62+ {
63+ {{{}, {}}, ge::DT_INT32, ge::FORMAT_ND},
64+ });
65+ std::vector<std::vector<int64_t>> expectOutputShape = {
66+ {3, 4, 5, 6, -1},
67+ };
68+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
69+}
Mmath/sub/docs/aclnnRsubs.md+4-4
@@ -19,7 +19,7 @@
19 19 
20## 函数原型20## 函数原型
21 21 
22-每个算子分为[两段式接口](common/两段式接口.md)，必须先调用“aclnnRsubsGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnRsubs”接口执行计算。22+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnRsubsGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnRsubs”接口执行计算。
23 23 
24- `aclnnStatus aclnnRsubsGetWorkspaceSize(const aclTensor *self, const aclScalar *other, const aclScalar *alpha, aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor)`24- `aclnnStatus aclnnRsubsGetWorkspaceSize(const aclTensor *self, const aclScalar *other, const aclScalar *alpha, aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor)`
25- `aclnnStatus aclnnRsubs(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)`25- `aclnnStatus aclnnRsubs(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)`
@@ -30,10 +30,10 @@
30 30 
31 - self(aclTensor*, 计算输入)：公式中的输入`self`，shape维度不高于8维。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。31 - self(aclTensor*, 计算输入)：公式中的输入`self`，shape维度不高于8维。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
32 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持INT8、UINT8、INT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16，且与other满足[互推导关系](../../../docs/zh/context/互推导关系.md)，如果self是FLOAT16、FLOAT32、BFLOAT16、DOUBLE其中一种，self和other推导后的数据类型为self的数据类型。32 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持INT8、UINT8、INT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16，且与other满足[互推导关系](../../../docs/zh/context/互推导关系.md)，如果self是FLOAT16、FLOAT32、BFLOAT16、DOUBLE其中一种，self和other推导后的数据类型为self的数据类型。
33- - <term>Ascend 950PR/Ascend 950DT</term>：数据类型支持INT8、UINT8、INT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16，且与other满足[TensorScalar互推导关系](common/TensorScalar互推导关系.md)。33+ - <term>Ascend 950PR/Ascend 950DT</term>：数据类型支持INT8、UINT8、INT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16，且与other满足[TensorScalar互推导关系](../../../docs/zh/context/TensorScalar互推导关系.md)。
34 - other(aclScalar*, 计算输入)：公式中的输入`other`。34 - other(aclScalar*, 计算输入)：公式中的输入`other`。
35 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持INT8、UINT8、INT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16，且与other满足[互推导关系](../../../docs/zh/context/互推导关系.md)。如果self的数据类型不在FLOAT16、FLOAT32、BFLOAT16、DOUBLE中，other的数据类型是FLOAT16、FLOAT32、BFLOAT16、DOUBLE其中一种，self和other[互推导关系](../../../docs/zh/context/互推导关系.md)；如果self和other的数据类型都不在FLOAT16、FLOAT32、BFLOAT16、DOUBLE中，推导后的数据类型为self的数据类型。35 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持INT8、UINT8、INT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16，且与other满足[互推导关系](../../../docs/zh/context/互推导关系.md)。如果self的数据类型不在FLOAT16、FLOAT32、BFLOAT16、DOUBLE中，other的数据类型是FLOAT16、FLOAT32、BFLOAT16、DOUBLE其中一种，self和other[互推导关系](../../../docs/zh/context/互推导关系.md)；如果self和other的数据类型都不在FLOAT16、FLOAT32、BFLOAT16、DOUBLE中，推导后的数据类型为self的数据类型。
36- - <term>Ascend 950PR/Ascend 950DT</term>：数据类型支持INT8、UINT8、INT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16，且与other满足[TensorScalar互推导关系](common/TensorScalar互推导关系.md)。36+ - <term>Ascend 950PR/Ascend 950DT</term>：数据类型支持INT8、UINT8、INT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16，且与other满足[TensorScalar互推导关系](../../../docs/zh/context/TensorScalar互推导关系.md)。
37 - alpha(aclScalar*, 计算输入)：公式中的`alpha`，数据类型需要可转换成self与other推导后的数据类型。37 - alpha(aclScalar*, 计算输入)：公式中的`alpha`，数据类型需要可转换成self与other推导后的数据类型。
38 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持INT8、UINT8、INT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16。38 - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持INT8、UINT8、INT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16。
39 - <term>Ascend 950PR/Ascend 950DT</term>：数据类型支持INT8、UINT8、INT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16。39 - <term>Ascend 950PR/Ascend 950DT</term>：数据类型支持INT8、UINT8、INT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16。
@@ -45,7 +45,7 @@
45 45 
46- **返回值：**46- **返回值：**
47 47 
48- aclnnStatus：返回状态码，具体参见[aclnn返回码](common/aclnn返回码.md)。48+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
49 49 
50 ```50 ```
51 第一段接口完成入参校验，出现以下场景时报错：51 第一段接口完成入参校验，出现以下场景时报错：