Star1339
1454
代码
代码
Issues119
Pull Requests262
Actions
项目讨论
Wiki
分析
Star1339
1454
已合并
添加 950 算子 reduce_log_sum #2628
sakuraqqz创建于 5月7日
添加 950 算子 reduce_log_sum #2628
已合并
sakuraqqz创建于 5月7日
从sakuraqqz/ops-mathreduceLogSum合入到cann/ops-mathmaster
共 27 个文件变更+1960-761
Mmath/reduce_log_sum/CMakeLists.txt+5-8
@@ -1,5 +1,5 @@
1# This program is free software, you can redistribute it and/or modify it.1# This program is free software, you can redistribute it and/or modify it.
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3# This file is a part of the CANN Open Software.3# This file is a part of the CANN Open Software.
4# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").4# Licensed under 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.
@@ -7,10 +7,7 @@
7# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of7# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
8# the software repository for the full text of the License.8# the software repository for the full text of the License.
9 9 
10-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)10+set(SUPPORT_COMPUTE_UNIT "ascend950" )
11-list(REMOVE_ITEM CURRENT_DIRS tests)11+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
12-foreach(SUB_DIR ${CURRENT_DIRS})12+set(SUPPORT_TILING_DIR "arch35" )
13- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")13+add_all_modules_sources(OPTYPE reduce_log_sum ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
14- add_subdirectory(${SUB_DIR})
15- endif()
16-endforeach()
Mmath/reduce_log_sum/docs/aclnnReduceLogSum.md+4-2
@@ -4,7 +4,7 @@
4 4 
5| 产品 | 是否支持 |5| 产品 | 是否支持 |
6| :----------------------------------------------------------- | :------: |6| :----------------------------------------------------------- | :------: |
7-| <term>Ascend 950PR/Ascend 950DT</term> | × |7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |8| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |9| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10| <term>Atlas 200I/500 A2 推理产品</term> | × |10| <term>Atlas 200I/500 A2 推理产品</term> | × |
@@ -69,7 +69,7 @@ aclnnStatus aclnnReduceLogSum(
69 <td>输入</td>69 <td>输入</td>
70 <td>表示参与计算的目标张量。</td>70 <td>表示参与计算的目标张量。</td>
71 <td>支持空Tensor。</td>71 <td>支持空Tensor。</td>
72- <td>FLOAT16、FLOAT32</td>72+ <td>FLOAT16、FLOAT32、BFLOAT16</td>
73 <td>ND</td>73 <td>ND</td>
74 <td>0-8</td>74 <td>0-8</td>
75 <td>√</td>75 <td>√</td>
@@ -136,6 +136,8 @@ aclnnStatus aclnnReduceLogSum(
136 </tr>136 </tr>
137 </tbody></table>137 </tbody></table>
138 138 
139+ - <term>Atlas 推理系列产品</term>、<term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：`data`、`reduce` 不支持BFLOAT16。
140+ 
139- **返回值：**141- **返回值：**
140 142 
141 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。143 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
Mmath/reduce_log_sum/examples/test_aclnn_reduce_log_sum.cpp+120-113
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions 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.
@@ -13,132 +13,139 @@
13#include "aclnnop/aclnn_reduce_log_sum.h"13#include "aclnnop/aclnn_reduce_log_sum.h"
14 14 
15#define CHECK_RET(cond, return_expr) \15#define CHECK_RET(cond, return_expr) \
16- do { \16+ do { \
17- if (!(cond)) { \17+ if (!(cond)) { \
18- return_expr; \18+ return_expr; \
19- } \19+ } \
20- } while (0)20+ } while (0)
21 21 
22-#define LOG_PRINT(message, ...) \22+#define LOG_PRINT(message, ...) \
23- do { \23+ do { \
24- printf(message, ##__VA_ARGS__); \24+ printf(message, ##__VA_ARGS__); \
25- } while (0)25+ } while (0)
26 26 
27-int64_t GetShapeSize(const std::vector<int64_t>& shape) {27+int64_t GetShapeSize(const std::vector<int64_t>& shape)
28- int64_t shapeSize = 1;28+{
29- for (auto i : shape) {29+ int64_t shapeSize = 1;
30- shapeSize *= i;30+ for (auto i : shape) {
31- }31+ shapeSize *= i;
32- return shapeSize;32+ }
33+ return shapeSize;
33}34}
34 35 
35-int Init(int32_t deviceId, aclrtStream* stream) {36+int Init(int32_t deviceId, aclrtStream* stream)
36- // 固定写法，资源初始化37+{
37- auto ret = aclInit(nullptr);38+ // 固定写法，资源初始化
38- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);39+ auto ret = aclInit(nullptr);
39- ret = aclrtSetDevice(deviceId);40+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
40- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);41+ ret = aclrtSetDevice(deviceId);
41- ret = aclrtCreateStream(stream);42+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
42- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);43+ ret = aclrtCreateStream(stream);
43- return 0;44+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
45+ return 0;
44}46}
45 47 
46template <typename T>48template <typename T>
47-int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,49+int CreateAclTensor(
48- aclDataType dataType, aclTensor** tensor) {50+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
49- auto size = GetShapeSize(shape) * sizeof(T);51+ aclTensor** tensor)
50- // 调用aclrtMalloc申请device侧内存52+{
51- auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);53+ auto size = GetShapeSize(shape) * sizeof(T);
52- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);54+ // 调用aclrtMalloc申请device侧内存
53- // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上55+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
54- ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);56+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
55- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);57+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
58+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
59+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
56 60 
57- // 计算连续tensor的strides61+ // 计算连续tensor的strides
58- std::vector<int64_t> strides(shape.size(), 1);62+ std::vector<int64_t> strides(shape.size(), 1);
59- for (int64_t i = shape.size() - 2; i >= 0; i--) {63+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
60- strides[i] = shape[i + 1] * strides[i + 1];64+ strides[i] = shape[i + 1] * strides[i + 1];
61- }65+ }
62 66 
63- // 调用aclCreateTensor接口创建aclTensor67+ // 调用aclCreateTensor接口创建aclTensor
64- *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,68+ *tensor = aclCreateTensor(
65- shape.data(), shape.size(), *deviceAddr);69+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
66- return 0;70+ *deviceAddr);
71+ return 0;
67}72}
68 73 
69-int main() {74+int main()
70- // 1. （固定写法）device/context/stream初始化，参考acl API手册75+{
71- // 根据自己的实际device填写deviceId76+ // 1. （固定写法）device/context/stream初始化，参考acl API手册
72- int32_t deviceId = 0;77+ // 根据自己的实际device填写deviceId
73- aclrtStream stream;78+ int32_t deviceId = 0;
74- auto ret = Init(deviceId, &stream);79+ aclrtStream stream;
75- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init failed. ERROR: %d\n", ret); return ret);80+ auto ret = Init(deviceId, &stream);
81+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init failed. ERROR: %d\n", ret); return ret);
76 82 
77- // 2. 构造输入与输出，需要根据API的接口自定义构造83+ // 2. 构造输入与输出，需要根据API的接口自定义构造
78- std::vector<int64_t> dataShape = {4, 2};84+ std::vector<int64_t> dataShape = {1};
79- std::vector<int64_t> outShape = {2};85+ std::vector<int64_t> outShape = {1};
80- void* dataDeviceAddr = nullptr;86+ void* dataDeviceAddr = nullptr;
81- void* outDeviceAddr = nullptr;87+ void* outDeviceAddr = nullptr;
82- aclTensor* data = nullptr;88+ aclTensor* data = nullptr;
83- aclIntArray* axes = nullptr;89+ aclIntArray* axes = nullptr;
84- aclTensor* out = nullptr;90+ aclTensor* out = nullptr;
85- std::vector<float> dataHostData = {0, 1, 2, 3, 4, 5, 6, 7};91+ std::vector<float> dataHostData = {9.7656};
86- std::vector<float> outHostData = {0, 0};92+ std::vector<float> outHostData = {0};
87- std::vector<int64_t> axesData = {0};93+ std::vector<int64_t> axesData = {0};
88- bool keepDims = false;94+ bool keepDims = true;
89- bool noopWithEmptyAxes = false;95+ bool noopWithEmptyAxes = false;
90- // 创建data aclTensor96+ // 创建data aclTensor
91- ret = CreateAclTensor(dataHostData, dataShape, &dataDeviceAddr, aclDataType::ACL_FLOAT, &data);97+ ret = CreateAclTensor(dataHostData, dataShape, &dataDeviceAddr, aclDataType::ACL_FLOAT, &data);
92- CHECK_RET(ret == ACL_SUCCESS, return ret);98+ CHECK_RET(ret == ACL_SUCCESS, return ret);
93- // 创建out aclTensor99+ // 创建out aclTensor
94- ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);100+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
95- CHECK_RET(ret == ACL_SUCCESS, return ret);101+ CHECK_RET(ret == ACL_SUCCESS, return ret);
96- // 创建axes aclIntArray102+ // 创建axes aclIntArray
97- axes = aclCreateIntArray(axesData.data(), 1);103+ axes = aclCreateIntArray(axesData.data(), 1);
98- CHECK_RET(axes != nullptr, return ret);104+ CHECK_RET(axes != nullptr, return ret);
99- // 3. 调用CANN算子库API，需要修改为具体的API名称105+ // 3. 调用CANN算子库API，需要修改为具体的API名称
100- uint64_t workspaceSize = 0;106+ uint64_t workspaceSize = 0;
101- aclOpExecutor* executor;107+ aclOpExecutor* executor;
102- // 调用aclnnReduceLogSum第一段接口108+ // 调用aclnnReduceLogSum第一段接口
103- ret = aclnnReduceLogSumGetWorkspaceSize(data, axes, keepDims, noopWithEmptyAxes, out, &workspaceSize, &executor);109+ ret = aclnnReduceLogSumGetWorkspaceSize(data, axes, keepDims, noopWithEmptyAxes, out, &workspaceSize, &executor);
104- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnReduceLogSumGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);110+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnReduceLogSumGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
105- // 根据第一段接口计算出的workspaceSize申请device内存111+ // 根据第一段接口计算出的workspaceSize申请device内存
106- void* workspaceAddr = nullptr;112+ void* workspaceAddr = nullptr;
107- if (workspaceSize > 0) {113+ if (workspaceSize > 0) {
108- ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);114+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
109- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);115+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
110- }116+ }
111- // 调用aclnnReduceLogSum第二段接口117+ // 调用aclnnReduceLogSum第二段接口
112- ret = aclnnReduceLogSum(workspaceAddr, workspaceSize, executor, stream);118+ ret = aclnnReduceLogSum(workspaceAddr, workspaceSize, executor, stream);
113- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnReduceLogSum failed. ERROR: %d\n", ret); return ret);119+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnReduceLogSum failed. ERROR: %d\n", ret); return ret);
114 120 
115- // 4. （固定写法）同步等待任务执行结束121+ // 4. （固定写法）同步等待任务执行结束
116- ret = aclrtSynchronizeStream(stream);122+ ret = aclrtSynchronizeStream(stream);
117- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);123+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
118 124 
119- // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改125+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
120- auto size = GetShapeSize(outShape);126+ auto size = GetShapeSize(outShape);
121- std::vector<float> resultData(size, 0);127+ std::vector<float> resultData(size, 0);
122- ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,128+ ret = aclrtMemcpy(
123- size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);129+ resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(resultData[0]),
124- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);130+ ACL_MEMCPY_DEVICE_TO_HOST);
125- for (int64_t i = 0; i < size; i++) {131+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
126- LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);132+ for (int64_t i = 0; i < size; i++) {
127- }133+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
134+ }
128 135 
129- // 6. 释放aclTensor和aclIntArray，需要根据具体API的接口定义修改136+ // 6. 释放aclTensor和aclIntArray，需要根据具体API的接口定义修改
130- aclDestroyTensor(data);137+ aclDestroyTensor(data);
131- aclDestroyIntArray(axes);138+ aclDestroyIntArray(axes);
132- aclDestroyTensor(out);139+ aclDestroyTensor(out);
133 140 
134- // 7. 释放device资源，需要根据具体API的接口定义修改141+ // 7. 释放device资源，需要根据具体API的接口定义修改
135- aclrtFree(dataDeviceAddr);142+ aclrtFree(dataDeviceAddr);
136- aclrtFree(outDeviceAddr);143+ aclrtFree(outDeviceAddr);
137- if (workspaceSize > 0) {144+ if (workspaceSize > 0) {
138- aclrtFree(workspaceAddr);145+ aclrtFree(workspaceAddr);
139- }146+ }
140- aclrtDestroyStream(stream);147+ aclrtDestroyStream(stream);
141- aclrtResetDevice(deviceId);148+ aclrtResetDevice(deviceId);
142- aclFinalize();149+ aclFinalize();
143- return 0;150+ return 0;
144}151}
Amath/reduce_log_sum/op_api/aclnn_reduce_log_sum.cpp+218-0
@@ -0,0 +1,218 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+#include <bitset>
11+#include "aclnn_reduce_log_sum.h"
12+#include "reduce_log_sum.h"
13+#include "aclnn_kernels/cast.h"
14+#include "aclnn_kernels/contiguous.h"
15+#include "aclnn_kernels/common/op_error_check.h"
16+#include "opdev/common_types.h"
17+#include "opdev/shape_utils.h"
18+#include "opdev/data_type_utils.h"
19+#include "opdev/format_utils.h"
20+#include "opdev/op_dfx.h"
21+#include "opdev/op_executor.h"
22+#include "opdev/op_log.h"
23+#include "opdev/tensor_view_utils.h"
24+#include "opdev/platform.h"
25+#include "conversion/fill/op_api/fill.h"
26+#include "op_api/op_api_def.h"
27+#include "op_api/aclnn_check.h"
28+ 
29+using namespace op;
30+#ifdef __cplusplus
31+extern "C" {
32+#endif
33+ 
34+constexpr size_t MAX_MASK_LEN = 64;
35+ 
36+// 根据API定义，需要列出所能支持的所有dtype
37+static const std::initializer_list<op::DataType> DTYPE_SUPPORT_LIST = {
38+ op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT};
39+ 
40+static bool CheckNotNull(const aclTensor* data, const aclIntArray* axes, const aclTensor* reduce)
41+{
42+ OP_CHECK_NULL(data, return false);
43+ OP_CHECK_NULL(axes, return false);
44+ OP_CHECK_NULL(reduce, return false);
45+ return true;
46+}
47+ 
48+static bool CheckDtypeValid(const aclTensor* data, const aclTensor* reduce)
49+{
50+ // 检查data和reduce的数据类型是否在支持列表内
51+ OP_CHECK_DTYPE_NOT_SUPPORT(data, DTYPE_SUPPORT_LIST, return false);
52+ OP_CHECK_DTYPE_NOT_SUPPORT(reduce, DTYPE_SUPPORT_LIST, return false);
53+ 
54+ return true;
55+}
56+ 
57+static bool CheckMaxDimension(const aclTensor* data)
58+{
59+ OP_CHECK_MAX_DIM(data, MAX_SUPPORT_DIMS_NUMS, return false);
60+ return true;
61+}
62+ 
63+static inline uint64_t GetPosDim(int64_t dim, int64_t dimNum)
64+{
65+ if (dimNum <= 0) {
66+ dimNum = 1;
67+ }
68+ return dim >= 0 ? dim : dim + dimNum;
69+}
70+ 
71+static bool CheckAxesValid(const aclTensor* data, const aclIntArray* axes)
72+{
73+ auto dataViewShape = data->GetViewShape();
74+ auto dataDimNum = static_cast<int64_t>(dataViewShape.GetDimNum());
75+ // data为标量时，axes range [-1, 0]
76+ if (dataDimNum <= 0) {
77+ dataDimNum = 1;
78+ }
79+ // axes为负时需要转正校验
80+ std::bitset<MAX_MASK_LEN> axesMask = std::bitset<MAX_MASK_LEN>();
81+ 
82+ for (size_t i = 0; i < axes->Size(); i++) {
83+ int64_t curDim = (*axes)[i];
84+ if (curDim >= dataDimNum || curDim < (-dataDimNum)) {
85+ OP_LOGE(
86+ ACLNN_ERR_PARAM_INVALID, "Provided axes %ld not in the range of input tensor size %ld.", curDim,
87+ dataDimNum);
88+ return false;
89+ }
90+ uint64_t index = GetPosDim(curDim, dataDimNum);
91+ if (axesMask[index]) {
92+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Axes %lu appears multiple times in the list of axes", index);
93+ }
94+ axesMask.set(index);
95+ }
96+ 
97+ return true;
98+}
99+ 
100+static aclnnStatus CheckParams(const aclTensor* data, const aclIntArray* axes, const aclTensor* reduce)
101+{
102+ // 1. 检查参数是否为空指针
103+ CHECK_RET(CheckNotNull(data, axes, reduce), ACLNN_ERR_PARAM_NULLPTR);
104+ 
105+ // 2. 检查data、reduce的数据类型是否合法
106+ CHECK_RET(CheckDtypeValid(data, reduce), ACLNN_ERR_PARAM_INVALID);
107+ 
108+ // 3. 检查最大维度是否超过8
109+ CHECK_RET(CheckMaxDimension(data), ACLNN_ERR_PARAM_INVALID);
110+ 
111+ // 4. 检查reduce的轴是否超出data维度范围
112+ CHECK_RET(CheckAxesValid(data, axes), ACLNN_ERR_PARAM_INVALID);
113+ 
114+ return ACLNN_SUCCESS;
115+}
116+ 
117+static aclnnStatus FillScalar(aclTensor* reduce, float val, aclOpExecutor* executor)
118+{
119+ FVector<int64_t> shape;
120+ size_t axesNum = reduce->GetViewShape().GetDimNum();
121+ 
122+ if (reduce->IsEmpty()) {
123+ return ACLNN_SUCCESS;
124+ }
125+ 
126+ for (size_t idx = 0; idx < axesNum; idx++) {
127+ int64_t tmpVal = reduce->GetViewShape().GetDim(idx);
128+ shape.push_back(tmpVal);
129+ }
130+ 
131+ auto axes = executor->ConvertToTensor(shape.data(), shape.size(), DataType::DT_INT64);
132+ auto shapeArray = executor->AllocIntArray(shape.data(), shape.size());
133+ 
134+ FVector<float> valVector = {val};
135+ auto valTensor = executor->ConvertToTensor(valVector.data(), valVector.size(), reduce->GetDataType());
136+ auto fillOut = l0op::Fill(axes, valTensor, shapeArray, executor);
137+ CHECK_RET(fillOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
138+ auto viewCopyResult = l0op::ViewCopy(fillOut, reduce, executor);
139+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
140+ return ACLNN_SUCCESS;
141+}
142+ 
143+aclnnStatus aclnnReduceLogSumGetWorkspaceSize(
144+ const aclTensor* data, const aclIntArray* axes, bool keepDims, bool noopWithEmptyAxes, aclTensor* reduce,
145+ uint64_t* workspaceSize, aclOpExecutor** executor)
146+{
147+ L2_DFX_PHASE_1(aclnnReduceLogSum, DFX_IN(data, axes, keepDims, noopWithEmptyAxes), DFX_OUT(reduce));
148+ 
149+ // 固定写法，创建OpExecutor
150+ auto uniqueExecutor = CREATE_EXECUTOR();
151+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
152+ 
153+ // 固定写法，参数检查
154+ auto ret = CheckParams(data, axes, reduce);
155+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
156+ 
157+ // 输入self为空tensor时，直接返回dtype类型的空tensor
158+ if (data->IsEmpty()) {
159+ ret = FillScalar(reduce, 0.0f, uniqueExecutor.get());
160+ if (ret == ACLNN_SUCCESS) {
161+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
162+ uniqueExecutor.ReleaseTo(executor);
163+ }
164+ return ret;
165+ }
166+ 
167+ op::Shape shape = data->GetViewShape();
168+ 
169+ //固定写法，将输入的data转换成连续的tensor
170+ auto dataContiguous = l0op::Contiguous(data, uniqueExecutor.get());
171+ const aclTensor* reduceOut = nullptr;
172+ 
173+ CHECK_RET(dataContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
174+ 
175+ if (axes->Size() == 0) {
176+ if (noopWithEmptyAxes == false) {
177+ size_t axesDum = shape.GetDimNum();
178+ std::vector<int64_t> appendDim(axesDum);
179+ for (size_t i = 0; i < axesDum; i++) {
180+ appendDim[i] = i;
181+ }
182+ axes = uniqueExecutor.get()->AllocIntArray(appendDim.data(), axesDum);
183+ reduceOut = l0op::ReduceLogSum(dataContiguous, axes, keepDims, uniqueExecutor.get());
184+ } else {
185+ //固定写法，将计算结果拷贝到输出reduce上，reduce可能是非连续的tensor
186+ auto viewCopyResult = l0op::ViewCopy(dataContiguous, reduce, uniqueExecutor.get());
187+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
188+ //固定写法，获取计算过程中需要使用的workspace大小
189+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
190+ uniqueExecutor.ReleaseTo(executor);
191+ return ACLNN_SUCCESS;
192+ }
193+ } else {
194+ reduceOut = l0op::ReduceLogSum(dataContiguous, axes, keepDims, uniqueExecutor.get());
195+ }
196+ 
197+ CHECK_RET(reduceOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
198+ CHECK_RET(CheckShapeAndScalarSame(reduceOut, reduce), ACLNN_ERR_PARAM_INVALID);
199+ 
200+ //固定写法，将计算结果拷贝到输出reduce上，reduce可能是非连续的tensor
201+ auto viewCopyResult = l0op::ViewCopy(reduceOut, reduce, uniqueExecutor.get());
202+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
203+ //固定写法，获取计算过程中需要使用的workspace大小
204+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
205+ uniqueExecutor.ReleaseTo(executor);
206+ return ACLNN_SUCCESS;
207+}
208+ 
209+aclnnStatus aclnnReduceLogSum(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
210+{
211+ L2_DFX_PHASE_2(aclnnReduceLogSum);
212+ // 固定写法，调用框架能力，完成计算
213+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
214+}
215+ 
216+#ifdef __cplusplus
217+}
218+#endif
Rmath/reduce_log_sum/op_host/op_api/aclnn_reduce_log_sum.h→math/reduce_log_sum/op_api/aclnn_reduce_log_sum.h+13-10
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions 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.
@@ -19,20 +19,23 @@ extern "C" {
19 * @brief aclnnReduceLogSum的第一段接口，根据具体的计算流程，计算workspace大小。19 * @brief aclnnReduceLogSum的第一段接口，根据具体的计算流程，计算workspace大小。
20 * @domain aclnn_math20 * @domain aclnn_math
21 * * 算子功能：使用输入边界的反射填充输入tensor。21 * * 算子功能：使用输入边界的反射填充输入tensor。
22- * @param [in] data: 表示参与计算的目标张量，维度小于8维，Device侧的aclTensor，支持[非连续的Tensor](../../../../docs/zh/context/非连续的Tensor.md)，22+ * @param [in] data:
23+ * 表示参与计算的目标张量，维度小于8维，Device侧的aclTensor，支持[非连续的Tensor](../../../../docs/zh/context/非连续的Tensor.md)，
23 * 数据类型支持FLOAT16、FLOAT32，[数据格式](../../../../docs/zh/context/数据格式.md)支持ND。24 * 数据类型支持FLOAT16、FLOAT32，[数据格式](../../../../docs/zh/context/数据格式.md)支持ND。
24 * @param [in] axes: 指定计算维度，Host侧的aclIntArray，数据类型支持INT64，取值范围为[-self.dim(), self.dim()-1]。25 * @param [in] axes: 指定计算维度，Host侧的aclIntArray，数据类型支持INT64，取值范围为[-self.dim(), self.dim()-1]。
25 * @param [in] keepDims: 指定是否在输出张量中保留输入张量的维度，Host侧的BOOL值。26 * @param [in] keepDims: 指定是否在输出张量中保留输入张量的维度，Host侧的BOOL值。
26- * @param [in] noopWithEmptyAxes: 指定axes为空时的行为：false即对所有轴进行计算；true即不进行计算，输出张量等于输入张量，Host侧的BOOL值。27+ * @param [in] noopWithEmptyAxes:
27- * @param [in] reduce: 表示计算后的结果，维度小于8维，Device侧的aclTensor，支持[非连续的Tensor](../../../../docs/zh/context/非连续的Tensor.md)，28+ * 指定axes为空时的行为：false即对所有轴进行计算；true即不进行计算，输出张量等于输入张量，Host侧的BOOL值。
29+ * @param [in] reduce:
30+ * 表示计算后的结果，维度小于8维，Device侧的aclTensor，支持[非连续的Tensor](../../../../docs/zh/context/非连续的Tensor.md)，
28 * 数据类型支持FLOAT16、FLOAT32，需与data一致，[数据格式](../../../../docs/zh/context/数据格式.md)支持ND。31 * 数据类型支持FLOAT16、FLOAT32，需与data一致，[数据格式](../../../../docs/zh/context/数据格式.md)支持ND。
29 * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。32 * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
30 * @param [out] executor: 返回op执行器，包含算子计算流程。33 * @param [out] executor: 返回op执行器，包含算子计算流程。
31 * @return aclnnStatus: 返回状态码。34 * @return aclnnStatus: 返回状态码。
32 */35 */
33-ACLNN_API aclnnStatus aclnnReduceLogSumGetWorkspaceSize(const aclTensor* data, const aclIntArray* axes, bool keepDims,36+ACLNN_API aclnnStatus aclnnReduceLogSumGetWorkspaceSize(
34- bool noopWithEmptyAxes, aclTensor* reduce, uint64_t* workspaceSize,37+ const aclTensor* data, const aclIntArray* axes, bool keepDims, bool noopWithEmptyAxes, aclTensor* reduce,
35- aclOpExecutor** executor);38+ uint64_t* workspaceSize, aclOpExecutor** executor);
36 39 
37/**40/**
38 * @brief aclnnReduceLogSum的第二段接口，用于执行计算。41 * @brief aclnnReduceLogSum的第二段接口，用于执行计算。
@@ -42,11 +45,11 @@ ACLNN_API aclnnStatus aclnnReduceLogSumGetWorkspaceSize(const aclTensor* data, c
42 * @param [in] executor: op执行器，包含了算子计算流程。45 * @param [in] executor: op执行器，包含了算子计算流程。
43 * @return aclnnStatus: 返回状态码。46 * @return aclnnStatus: 返回状态码。
44 */47 */
45-ACLNN_API aclnnStatus aclnnReduceLogSum(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,48+ACLNN_API aclnnStatus
46- aclrtStream stream);49+aclnnReduceLogSum(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
47 50 
48#ifdef __cplusplus51#ifdef __cplusplus
49}52}
50#endif53#endif
51 54 
52-#endif // OP_API_INC_REDUCE_LOG_SUM_H_55+#endif // OP_API_INC_REDUCE_LOG_SUM_H_
Rmath/reduce_log_sum/op_host/op_api/reduce_log_sum.cpp→math/reduce_log_sum/op_api/reduce_log_sum.cpp+36-36
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions 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.
@@ -16,7 +16,6 @@
16#include "opdev/make_op_executor.h"16#include "opdev/make_op_executor.h"
17#include "opdev/shape_utils.h"17#include "opdev/shape_utils.h"
18 18 
19- 
20using namespace op;19using namespace op;
21 20 
22namespace l0op {21namespace l0op {
@@ -25,46 +24,47 @@ OP_TYPE_REGISTER(ReduceLogSum);
25static const std::initializer_list<DataType> AICORE_DTYPE_SUPPORT_LIST = {DataType::DT_FLOAT, DataType::DT_FLOAT16};24static const std::initializer_list<DataType> AICORE_DTYPE_SUPPORT_LIST = {DataType::DT_FLOAT, DataType::DT_FLOAT16};
26 25 
27//根据芯片类型，dtype判断算子是否支持走aicore26//根据芯片类型，dtype判断算子是否支持走aicore
28-static inline bool IsAiCoreSupport(DataType inputDtype) {27+static inline bool IsAiCoreSupport(DataType inputDtype)
29- //只需要判断dtype28+{
30- return CheckType(inputDtype, AICORE_DTYPE_SUPPORT_LIST);29+ //只需要判断dtype
30+ return CheckType(inputDtype, AICORE_DTYPE_SUPPORT_LIST);
31}31}
32 32 
33+// AICORE算子kernel
34+static inline const aclTensor* ReduceLogSumAiCore(
35+ const aclTensor* data, const aclTensor* axes, aclTensor* reduce, bool keepDims, aclOpExecutor* executor)
36+{
37+ L0_DFX(ReduceLogSumAiCore, data, axes, reduce, keepDims);
38+ //使用框架宏ADD_TO_LAUNCHER_LIST_AICORE,将AiCore ReduceLogSum算子加入任务队列
39+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(ReduceLogSum, OP_INPUT(data, axes), OP_OUTPUT(reduce), OP_ATTR(keepDims));
40+ OP_CHECK(
41+ ret == ACLNN_SUCCESS,
42+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "ReduceLogSumAiCore ADD_TO_LAUNCHER_LIST_AICORE failed."), return nullptr);
33 43 
34-//AICORE算子kernel44+ return reduce;
35-static inline const aclTensor* ReduceLogSumAiCore(const aclTensor* data,
36- const aclTensor* axes,
37- aclTensor* reduce,
38- bool keepDims,
39- aclOpExecutor* executor) {
40- L0_DFX(ReduceLogSumAiCore, data, axes, reduce, keepDims);
41- //使用框架宏ADD_TO_LAUNCHER_LIST_AICORE,将AiCore ReduceLogSum算子加入任务队列
42- auto ret = ADD_TO_LAUNCHER_LIST_AICORE(ReduceLogSum, OP_INPUT(data,axes), OP_OUTPUT(reduce), OP_ATTR(keepDims));
43- OP_CHECK(ret == ACLNN_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "ReduceLogSumAiCore ADD_TO_LAUNCHER_LIST_AICORE failed."), return nullptr);
44- 
45- return reduce;
46}45}
47 46 
48-const aclTensor* ReduceLogSum(const aclTensor* data, const aclIntArray* axes, bool keepDims, aclOpExecutor* executor) {47+const aclTensor* ReduceLogSum(const aclTensor* data, const aclIntArray* axes, bool keepDims, aclOpExecutor* executor)
49- auto dims = executor->ConvertToTensor(axes, op::DataType::DT_INT64);48+{
50- auto output = executor->AllocTensor(data->GetDataType(), op::Format::FORMAT_ND, op::Format::FORMAT_ND);49+ auto dims = executor->ConvertToTensor(axes, op::DataType::DT_INT64);
51- INFER_SHAPE(ReduceLogSum, OP_INPUT(data, dims), OP_OUTPUT(output), OP_ATTR(keepDims));50+ auto output = executor->AllocTensor(data->GetDataType(), op::Format::FORMAT_ND, op::Format::FORMAT_ND);
51+ INFER_SHAPE(ReduceLogSum, OP_INPUT(data, dims), OP_OUTPUT(output), OP_ATTR(keepDims));
52 52 
53- op::Shape outShape = data->GetViewShape();53+ op::Shape outShape = data->GetViewShape();
54- auto count = axes->Size();54+ auto count = axes->Size();
55- size_t dimNum = outShape.GetDimNum();55+ size_t dimNum = outShape.GetDimNum();
56- if (keepDims) {56+ if (keepDims) {
57- for (uint64_t i = 0; i < count; i++) {57+ for (uint64_t i = 0; i < count; i++) {
58- int64_t dimIndex = static_cast<int64_t>((*axes)[i]);58+ int64_t dimIndex = static_cast<int64_t>((*axes)[i]);
59- int64_t dimNew = dimIndex >= 0 ? dimIndex : dimIndex + dimNum;59+ int64_t dimNew = dimIndex >= 0 ? dimIndex : dimIndex + dimNum;
60- outShape.SetDim(dimNew, 1);60+ outShape.SetDim(dimNew, 1);
61+ }
62+ output->SetViewShape(outShape);
61 }63 }
62- output->SetViewShape(outShape);
63- }
64 64 
65- if (IsAiCoreSupport(data->GetDataType())) {65+ if (IsAiCoreSupport(data->GetDataType())) {
66- return ReduceLogSumAiCore(data, dims, output, keepDims, executor);66+ return ReduceLogSumAiCore(data, dims, output, keepDims, executor);
67- }67+ }
68- return output;68+ return output;
69}69}
70-} // namespace l0op70+} // namespace l0op
Rmath/reduce_log_sum/op_host/op_api/reduce_log_sum.h→math/reduce_log_sum/op_api/reduce_log_sum.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions 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.
@@ -16,4 +16,4 @@ namespace l0op {
16const aclTensor* ReduceLogSum(const aclTensor* data, const aclIntArray* axes, bool keepDims, aclOpExecutor* executor);16const aclTensor* ReduceLogSum(const aclTensor* data, const aclIntArray* axes, bool keepDims, aclOpExecutor* executor);
17}17}
18 18 
19-#endif // OP_API_INC_LEVEL0_REDUCE_LOG_SUM_H19+#endif // OP_API_INC_LEVEL0_REDUCE_LOG_SUM_H
Mmath/reduce_log_sum/op_host/CMakeLists.txt+5-3
@@ -1,10 +1,12 @@
1+# ----------------------------------------------------------------------------
1# This program is free software, you can redistribute it and/or modify it.2# This program is free software, you can redistribute it and/or modify it.
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3# This file is a part of the CANN Open Software.4# This file is a part of the CANN Open Software.
4# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").5# Licensed under 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# 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, INCLUDING7# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
7-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-# the software repository for the full text of the License.9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
9 11 
10add_modules_sources()12add_modules_sources()
Amath/reduce_log_sum/op_host/arch35/reduce_log_sum_tiling_arch35.cpp+92-0
@@ -0,0 +1,92 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file reduce_log_sum_tiling_arch35.cpp
13+ * \brief tiling for reduce sum
14+ */
15+ 
16+#include <vector>
17+#include "log/log.h"
18+#include "register/op_impl_registry.h"
19+#include "atvoss/reduce/reduce_tiling.h"
20+#include "atvoss/reduce/reduce_tiling_data.h"
21+#include "op_host/tiling_base_util.h"
22+#include "math/reduce_log_sum/op_kernel/arch35/reduce_log_sum_dag.h"
23+#include "math/reduce_log_sum/op_kernel/arch35/reduce_log_sum_tiling_key.h"
24+ 
25+using namespace Ops::Base;
26+ 
27+namespace optiling {
28+// 支持的数据类型字节数
29+static constexpr int32_t SIZE8 = 8;
C
Cchenxingyu185月27日

建议添加注释说明其作用

likedislike
S
sakuraqqz
5月27日 评论：
30+static constexpr int32_t SIZE4 = 4;
31+static constexpr int32_t SIZE2 = 2;
32+static ge::graphStatus DoTiling(gert::TilingContext* context, ReduceOpInputParam& opInput, ReduceTilingKey& key)
33+{
34+ ge::graphStatus status = ge::GRAPH_FAILED;
35+ 
36+ if (ge::GetSizeByDataType(opInput.inputDtype) == SIZE8) {
37+ status = Tiling4ReduceOp<ReduceLogSum::ReduceLogSumDag<int64_t, int64_t>::OpDag>(context, opInput, key);
38+ } else if (ge::GetSizeByDataType(opInput.inputDtype) == SIZE4) {
39+ status = Tiling4ReduceOp<ReduceLogSum::ReduceLogSumDag<float, float>::OpDag>(context, opInput, key);
40+ } else if (ge::GetSizeByDataType(opInput.inputDtype) == SIZE2) {
41+ status = Tiling4ReduceOp<ReduceLogSum::ReduceLogSumDag<half, float>::OpDag>(context, opInput, key);
42+ }
43+ OP_CHECK_IF(
44+ (status == ge::GRAPH_FAILED),
45+ OP_LOGE(
46+ context->GetNodeName(), "ReduceOp Tiling failed, dtype shoude be in (bfloat16/float16/float/int32/int64)"),
47+ return ge::GRAPH_FAILED);
48+ return status;
49+}
50+ 
51+static ge::graphStatus Tiling4ReduceLogSum(gert::TilingContext* context)
52+{
53+ auto compileInfo = reinterpret_cast<const ReduceOpCompileInfo*>(context->GetCompileInfo());
54+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
55+ 
56+ ReduceOpInputParam opInput;
57+ OP_CHECK_IF(
58+ (ReduceOpTmpl::GetInputParam(context, opInput, 0, 1, 0) == ge::GRAPH_FAILED),
59+ OP_LOGE(context->GetNodeName(), "ReduceOp get x input param failed"), return ge::GRAPH_FAILED);
60+ 
61+ if (opInput.axes.empty()) {
62+ auto attrs = context->GetAttrs();
63+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
64+ const bool isNoopWithEmpty = *(attrs->GetAttrPointer<bool>(1));
65+ if (!isNoopWithEmpty) {
66+ opInput.axes.resize(opInput.shape.size());
67+ for (size_t i = 0; i < opInput.shape.size(); i++) {
68+ opInput.axes[i] = i;
69+ }
70+ }
71+ }
72+ ReduceTilingKey key;
73+ OP_CHECK_IF(
74+ (DoTiling(context, opInput, key) == ge::GRAPH_FAILED),
75+ OP_LOGE(context->GetNodeName(), "DoTiling Failed for ReduceLogSum"), return ge::GRAPH_FAILED);
76+ uint64_t tilingKey;
77+ GEN_REDUCE_TILING_KEY(tilingKey, key);
78+ OP_LOGI(
79+ context->GetNodeName(), "patternID:%u, loopARCount:%u, loopInnerARCount:%u, isContiguous:%d, Tiling Key is:%lu",
80+ key.patternID, key.loopARCount, key.loopInnerARCount, key.isContiguous ? 1 : 0, tilingKey);
81+ context->SetTilingKey(tilingKey);
82+ return ge::GRAPH_SUCCESS;
83+}
84+ 
85+static ge::graphStatus TilingPrepare4ReduceLogSum(gert::TilingParseContext* context)
86+{
87+ (void)context;
88+ return ge::GRAPH_SUCCESS;
89+}
90+ 
91+IMPL_OP_OPTILING(ReduceLogSum).Tiling(Tiling4ReduceLogSum).TilingParse<ReduceOpCompileInfo>(TilingPrepare4ReduceLogSum);
92+} // namespace optiling
Amath/reduce_log_sum/op_host/config/ascend950/reduce_log_sum_binary.json+287-0
@@ -0,0 +1,287 @@
1+{
2+ "op_type": "ReduceLogSum",
3+ "op_list": [
4+ {
5+ "bin_filename": "ReduceLogSum_float16_int32",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "float16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "format_match_mode": "FormatAgnostic"
17+ },
18+ {
19+ "name": "axes",
20+ "index": 1,
21+ "dtype": "int32",
22+ "format": "ND",
23+ "paramType": "required",
24+ "shape": [
25+ -2
26+ ],
27+ "format_match_mode": "FormatAgnostic"
28+ }
29+ ],
30+ "outputs": [
31+ {
32+ "name": "y",
33+ "index": 0,
34+ "dtype": "float16",
35+ "format": "ND",
36+ "paramType": "required",
37+ "shape": [
38+ -2
39+ ],
40+ "format_match_mode": "FormatAgnostic"
41+ }
42+ ],
43+ "attrs": [
44+ {
45+ "name": "keep_dims",
46+ "dtype": "bool",
47+ "value": null
48+ }
49+ ]
50+ },
51+ {
52+ "bin_filename": "ReduceLogSum_float16_int64",
53+ "inputs": [
54+ {
55+ "name": "x",
56+ "index": 0,
57+ "dtype": "float16",
58+ "format": "ND",
59+ "paramType": "required",
60+ "shape": [
61+ -2
62+ ],
63+ "format_match_mode": "FormatAgnostic"
64+ },
65+ {
66+ "name": "axes",
67+ "index": 1,
68+ "dtype": "int64",
69+ "format": "ND",
70+ "paramType": "required",
71+ "shape": [
72+ -2
73+ ],
74+ "format_match_mode": "FormatAgnostic"
75+ }
76+ ],
77+ "outputs": [
78+ {
79+ "name": "y",
80+ "index": 0,
81+ "dtype": "float16",
82+ "format": "ND",
83+ "paramType": "required",
84+ "shape": [
85+ -2
86+ ],
87+ "format_match_mode": "FormatAgnostic"
88+ }
89+ ],
90+ "attrs": [
91+ {
92+ "name": "keep_dims",
93+ "dtype": "bool",
94+ "value": null
95+ }
96+ ]
97+ },
98+ {
99+ "bin_filename": "ReduceLogSum_float32_int32",
100+ "inputs": [
101+ {
102+ "name": "x",
103+ "index": 0,
104+ "dtype": "float32",
105+ "format": "ND",
106+ "paramType": "required",
107+ "shape": [
108+ -2
109+ ],
110+ "format_match_mode": "FormatAgnostic"
111+ },
112+ {
113+ "name": "axes",
114+ "index": 1,
115+ "dtype": "int32",
116+ "format": "ND",
117+ "paramType": "required",
118+ "shape": [
119+ -2
120+ ],
121+ "format_match_mode": "FormatAgnostic"
122+ }
123+ ],
124+ "outputs": [
125+ {
126+ "name": "y",
127+ "index": 0,
128+ "dtype": "float32",
129+ "format": "ND",
130+ "paramType": "required",
131+ "shape": [
132+ -2
133+ ],
134+ "format_match_mode": "FormatAgnostic"
135+ }
136+ ],
137+ "attrs": [
138+ {
139+ "name": "keep_dims",
140+ "dtype": "bool",
141+ "value": null
142+ }
143+ ]
144+ },
145+ {
146+ "bin_filename": "ReduceLogSum_float32_int64",
147+ "inputs": [
148+ {
149+ "name": "x",
150+ "index": 0,
151+ "dtype": "float32",
152+ "format": "ND",
153+ "paramType": "required",
154+ "shape": [
155+ -2
156+ ],
157+ "format_match_mode": "FormatAgnostic"
158+ },
159+ {
160+ "name": "axes",
161+ "index": 1,
162+ "dtype": "int64",
163+ "format": "ND",
164+ "paramType": "required",
165+ "shape": [
166+ -2
167+ ],
168+ "format_match_mode": "FormatAgnostic"
169+ }
170+ ],
171+ "outputs": [
172+ {
173+ "name": "y",
174+ "index": 0,
175+ "dtype": "float32",
176+ "format": "ND",
177+ "paramType": "required",
178+ "shape": [
179+ -2
180+ ],
181+ "format_match_mode": "FormatAgnostic"
182+ }
183+ ],
184+ "attrs": [
185+ {
186+ "name": "keep_dims",
187+ "dtype": "bool",
188+ "value": null
189+ }
190+ ]
191+ },
192+ {
193+ "bin_filename": "ReduceLogSum_bfloat16_int32",
194+ "inputs": [
195+ {
196+ "name": "x",
197+ "index": 0,
198+ "dtype": "bfloat16",
199+ "format": "ND",
200+ "paramType": "required",
201+ "shape": [
202+ -2
203+ ],
204+ "format_match_mode": "FormatAgnostic"
205+ },
206+ {
207+ "name": "axes",
208+ "index": 1,
209+ "dtype": "int32",
210+ "format": "ND",
211+ "paramType": "required",
212+ "shape": [
213+ -2
214+ ],
215+ "format_match_mode": "FormatAgnostic"
216+ }
217+ ],
218+ "outputs": [
219+ {
220+ "name": "y",
221+ "index": 0,
222+ "dtype": "bfloat16",
223+ "format": "ND",
224+ "paramType": "required",
225+ "shape": [
226+ -2
227+ ],
228+ "format_match_mode": "FormatAgnostic"
229+ }
230+ ],
231+ "attrs": [
232+ {
233+ "name": "keep_dims",
234+ "dtype": "bool",
235+ "value": null
236+ }
237+ ]
238+ },
239+ {
240+ "bin_filename": "ReduceLogSum_bfloat16_int64",
241+ "inputs": [
242+ {
243+ "name": "x",
244+ "index": 0,
245+ "dtype": "bfloat16",
246+ "format": "ND",
247+ "paramType": "required",
248+ "shape": [
249+ -2
250+ ],
251+ "format_match_mode": "FormatAgnostic"
252+ },
253+ {
254+ "name": "axes",
255+ "index": 1,
256+ "dtype": "int64",
257+ "format": "ND",
258+ "paramType": "required",
259+ "shape": [
260+ -2
261+ ],
262+ "format_match_mode": "FormatAgnostic"
263+ }
264+ ],
265+ "outputs": [
266+ {
267+ "name": "y",
268+ "index": 0,
269+ "dtype": "bfloat16",
270+ "format": "ND",
271+ "paramType": "required",
272+ "shape": [
273+ -2
274+ ],
275+ "format_match_mode": "FormatAgnostic"
276+ }
277+ ],
278+ "attrs": [
279+ {
280+ "name": "keep_dims",
281+ "dtype": "bool",
282+ "value": null
283+ }
284+ ]
285+ }
286+ ]
287+}
Amath/reduce_log_sum/op_host/config/ascend950/reduce_log_sum_simplified_key.ini+2-0
@@ -0,0 +1,2 @@
1+[ReduceLogSum]
2+default=0
Dmath/reduce_log_sum/op_host/op_api/aclnn_reduce_log_sum.cpp+0-209
@@ -1,209 +0,0 @@
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-#include <bitset>
11-#include "aclnn_reduce_log_sum.h"
12-#include "reduce_log_sum.h"
13-#include "aclnn_kernels/cast.h"
14-#include "aclnn_kernels/contiguous.h"
15-#include "aclnn_kernels/common/op_error_check.h"
16-#include "opdev/common_types.h"
17-#include "opdev/shape_utils.h"
18-#include "opdev/data_type_utils.h"
19-#include "opdev/format_utils.h"
20-#include "opdev/op_dfx.h"
21-#include "opdev/op_executor.h"
22-#include "opdev/op_log.h"
23-#include "opdev/tensor_view_utils.h"
24-#include "opdev/platform.h"
25-#include "conversion/fill/op_api/fill.h"
26-#include "op_api/op_api_def.h"
27-#include "op_api/aclnn_check.h"
28- 
29-using namespace op;
30-#ifdef __cplusplus
31-extern "C" {
32-#endif
33- 
34-constexpr size_t MAX_MASK_LEN = 64;
35- 
36-// 根据API定义，需要列出所能支持的所有dtype
37-static const std::initializer_list<op::DataType> DTYPE_SUPPORT_LIST = {
38- op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT};
39- 
40-static bool CheckNotNull(const aclTensor *data, const aclIntArray *axes, const aclTensor *reduce) {
41- OP_CHECK_NULL(data, return false);
42- OP_CHECK_NULL(axes, return false);
43- OP_CHECK_NULL(reduce, return false);
44- return true;
45-}
46- 
47-static bool CheckDtypeValid(const aclTensor *data, const aclTensor *reduce) {
48- // 检查data和reduce的数据类型是否在支持列表内
49- OP_CHECK_DTYPE_NOT_SUPPORT(data, DTYPE_SUPPORT_LIST, return false);
50- OP_CHECK_DTYPE_NOT_SUPPORT(reduce, DTYPE_SUPPORT_LIST, return false);
51- 
52- return true;
53-}
54- 
55-static bool CheckMaxDimension(const aclTensor *data) {
56- OP_CHECK_MAX_DIM(data, MAX_SUPPORT_DIMS_NUMS, return false);
57- return true;
58-}
59- 
60-static inline uint64_t GetPosDim(int64_t dim, int64_t dimNum) {
61- if (dimNum <= 0) {
62- dimNum = 1;
63- }
64- return dim >= 0 ? dim : dim + dimNum;
65-}
66- 
67-static bool CheckAxesValid(const aclTensor* data, const aclIntArray* axes) {
68- auto dataViewShape = data->GetViewShape();
69- auto dataDimNum = static_cast<int64_t>(dataViewShape.GetDimNum());
70- //data为标量时，axes range [-1, 0]
71- if (dataDimNum <= 0) {
72- dataDimNum = 1;
73- }
74- // axes为负时需要转正校验
75- std::bitset<MAX_MASK_LEN> axesMask = std::bitset<MAX_MASK_LEN>();
76- 
77- for (size_t i = 0; i < axes->Size(); i++) {
78- int64_t curDim = (*axes)[i];
79- if (curDim >= dataDimNum || curDim < (-dataDimNum)) {
80- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Provided axes %ld not in the range of input tensor size %ld.", curDim,
81- dataDimNum);
82- return false;
83- }
84- uint64_t index = GetPosDim(curDim, dataDimNum);
85- if (axesMask[index]) {
86- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Axes %lu appears multiple times in the list of axes", index);
87- }
88- axesMask.set(index);
89- }
90- 
91- return true;
92-}
93- 
94-static aclnnStatus CheckParams(const aclTensor *data, const aclIntArray *axes, const aclTensor *reduce) {
95- // 1. 检查参数是否为空指针
96- CHECK_RET(CheckNotNull(data, axes, reduce), ACLNN_ERR_PARAM_NULLPTR);
97- 
98- // 2. 检查data、reduce的数据类型是否合法
99- CHECK_RET(CheckDtypeValid(data, reduce), ACLNN_ERR_PARAM_INVALID);
100- 
101- // 3. 检查最大维度是否超过8
102- CHECK_RET(CheckMaxDimension(data), ACLNN_ERR_PARAM_INVALID);
103- 
104- // 4. 检查reduce的轴是否超出data维度范围
105- CHECK_RET(CheckAxesValid(data, axes), ACLNN_ERR_PARAM_INVALID);
106- 
107- return ACLNN_SUCCESS;
108-}
109- 
110-static aclnnStatus FillScalar(aclTensor *reduce, float val, aclOpExecutor *executor)
111-{
112- FVector<int64_t> shape;
113- size_t axesNum = reduce->GetViewShape().GetDimNum();
114- 
115- if (reduce->IsEmpty()) {
116- return ACLNN_SUCCESS;
117- }
118- 
119- for (size_t idx = 0; idx < axesNum; idx++) {
120- int64_t tmpVal = reduce->GetViewShape().GetDim(idx);
121- shape.push_back(tmpVal);
122- }
123- 
124- auto axes = executor->ConvertToTensor(shape.data(), shape.size(), DataType::DT_INT64);
125- auto shapeArray = executor->AllocIntArray(shape.data(), shape.size());
126- 
127- FVector<float> valVector = {val};
128- auto valTensor = executor->ConvertToTensor(valVector.data(), valVector.size(), reduce->GetDataType());
129- auto fillOut = l0op::Fill(axes, valTensor, shapeArray, executor);
130- CHECK_RET(fillOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
131- auto viewCopyResult = l0op::ViewCopy(fillOut, reduce, executor);
132- CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
133- return ACLNN_SUCCESS;
134-}
135- 
136-aclnnStatus aclnnReduceLogSumGetWorkspaceSize(const aclTensor *data, const aclIntArray *axes, bool keepDims,
137- bool noopWithEmptyAxes, aclTensor *reduce, uint64_t *workspaceSize,
138- aclOpExecutor **executor) {
139- L2_DFX_PHASE_1(aclnnReduceLogSum, DFX_IN(data, axes, keepDims, noopWithEmptyAxes), DFX_OUT(reduce));
140- 
141- // 固定写法，创建OpExecutor
142- auto uniqueExecutor = CREATE_EXECUTOR();
143- CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
144- 
145- // 固定写法，参数检查
146- auto ret = CheckParams(data, axes, reduce);
147- CHECK_RET(ret == ACLNN_SUCCESS, ret);
148- 
149- // 输入self为空tensor时，直接返回dtype类型的空tensor
150- if (data->IsEmpty()) {
151- ret = FillScalar(reduce, 0.0f, uniqueExecutor.get());
152- if (ret == ACLNN_SUCCESS) {
153- *workspaceSize = uniqueExecutor->GetWorkspaceSize();
154- uniqueExecutor.ReleaseTo(executor);
155- }
156- return ret;
157- }
158- 
159- op::Shape shape = data->GetViewShape();
160- 
161- //固定写法，将输入的data转换成连续的tensor
162- auto dataContiguous = l0op::Contiguous(data, uniqueExecutor.get());
163- const aclTensor* reduceOut = nullptr;
164- 
165- CHECK_RET(dataContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
166- 
167- if (axes->Size() == 0) {
168- if (noopWithEmptyAxes == false) {
169- size_t axesDum = shape.GetDimNum();
170- std::vector<int64_t> appendDim(axesDum);
171- for (size_t i = 0; i < axesDum; i++) {
172- appendDim[i] = i;
173- }
174- axes = uniqueExecutor.get()->AllocIntArray(appendDim.data(), axesDum);
175- reduceOut = l0op::ReduceLogSum(dataContiguous, axes, keepDims, uniqueExecutor.get());
176- } else {
177- //固定写法，将计算结果拷贝到输出reduce上，reduce可能是非连续的tensor
178- auto viewCopyResult = l0op::ViewCopy(dataContiguous, reduce, uniqueExecutor.get());
179- CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
180- //固定写法，获取计算过程中需要使用的workspace大小
181- *workspaceSize = uniqueExecutor->GetWorkspaceSize();
182- uniqueExecutor.ReleaseTo(executor);
183- return ACLNN_SUCCESS;
184- }
185- } else {
186- reduceOut = l0op::ReduceLogSum(dataContiguous, axes, keepDims, uniqueExecutor.get());
187- }
188- 
189- CHECK_RET(reduceOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
190- CHECK_RET(CheckShapeAndScalarSame(reduceOut, reduce), ACLNN_ERR_PARAM_INVALID);
191- 
192- //固定写法，将计算结果拷贝到输出reduce上，reduce可能是非连续的tensor
193- auto viewCopyResult = l0op::ViewCopy(reduceOut, reduce, uniqueExecutor.get());
194- CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
195- //固定写法，获取计算过程中需要使用的workspace大小
196- *workspaceSize = uniqueExecutor->GetWorkspaceSize();
197- uniqueExecutor.ReleaseTo(executor);
198- return ACLNN_SUCCESS;
199-}
200- 
201-aclnnStatus aclnnReduceLogSum(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream) {
202- L2_DFX_PHASE_2(aclnnReduceLogSum);
203- // 固定写法，调用框架能力，完成计算
204- return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
205-}
206- 
207-#ifdef __cplusplus
208-}
209-#endif
Amath/reduce_log_sum/op_host/reduce_log_sum_def.cpp+48-0
@@ -0,0 +1,48 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+/*!
11+ * \file reduce_log_sum.cpp
12+ * \brief aicore info for ReduceLogSum op
13+ */
14+#include "register/op_def_registry.h"
15+ 
16+namespace ops {
17+static const std::vector<ge::DataType> dataType = {ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16,
18+ ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16};
19+ 
20+static const std::vector<ge::Format> format = {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
21+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
22+ 
23+static const std::vector<ge::DataType> axesDataType = {ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
24+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64};
25+ 
26+class ReduceLogSum : public OpDef {
27+public:
28+ explicit ReduceLogSum(const char* name) : OpDef(name)
29+ {
30+ this->Input("x").ParamType(REQUIRED).DataType(dataType).UnknownShapeFormat(format);
31+ 
32+ this->Input("axes").ParamType(REQUIRED).ValueDepend(OPTIONAL).DataType(axesDataType).UnknownShapeFormat(format);
33+ 
34+ this->Output("y").ParamType(REQUIRED).DataType(dataType).UnknownShapeFormat(format);
35+ 
36+ this->Attr("keep_dims").AttrType(OPTIONAL).Bool(false);
37+ 
38+ OpAICoreConfig aicoreConfig;
39+ aicoreConfig.DynamicCompileStaticFlag(true)
40+ .DynamicRankSupportFlag(true)
41+ .DynamicShapeSupportFlag(true)
42+ .ExtendCfgInfo("opFile.value", "reduce_log_sum_apt");
43+ this->AICore().AddConfig("ascend950", aicoreConfig);
44+ }
45+};
46+ 
47+OP_ADD(ReduceLogSum);
48+} // namespace ops
Amath/reduce_log_sum/op_host/reduce_log_sum_infershape.cpp+62-0
@@ -0,0 +1,62 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file reduce_infer.cc
13+ * \brief
14+ */
15+#include "log/log.h"
16+#include "register/op_impl_registry.h"
17+#include "op_host/util/shape_util.h"
18+#include "op_api/op_util.h"
19+#include "util/math_util.h"
20+#include "op_host/infershape_reduce_util.h"
21+ 
22+using namespace ge;
23+using namespace Ops::Base;
24+namespace ops {
25+static ge::graphStatus InferShape4ReduceCommon(gert::InferShapeContext* context)
26+{
27+ auto in_shape = context->GetInputShape(0);
28+ OP_CHECK_NULL_WITH_CONTEXT(context, in_shape);
29+ auto axes_tensor = context->GetInputTensor(1);
30+ OP_CHECK_NULL_WITH_CONTEXT(context, axes_tensor);
31+ auto out_shape = context->GetOutputShape(0);
32+ OP_CHECK_NULL_WITH_CONTEXT(context, out_shape);
33+ auto attrs = context->GetAttrs();
34+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
35+ 
36+ const bool* keep_dims = attrs->GetAttrPointer<bool>(0);
37+ OP_CHECK_NULL_WITH_CONTEXT(context, keep_dims);
38+ 
39+ auto axes_size = static_cast<int32_t>(axes_tensor->GetShapeSize());
40+ 
41+ OP_CHECK_IF(
42+ axes_size < 0, OP_LOGE(context->GetNodeName(), "axes num cannot be less than 0!"), return ge::GRAPH_FAILED);
43+ 
44+ if (axes_size == 0) {
45+ *out_shape = *in_shape;
46+ OP_LOGD(context->GetNodeName(), "axes is empty tensor, will ignore infer, set output shape = input shape");
47+ return ge::GRAPH_SUCCESS;
48+ }
49+ 
50+ auto dtype = axes_tensor->GetDataType();
51+ OP_CHECK_IF(
52+ dtype != ge::DT_INT32 && dtype != ge::DT_INT64,
53+ OP_LOGE(context->GetNodeName(), "axes datatype %s must in (int32, int64)", ToString(dtype).c_str()),
54+ return ge::GRAPH_FAILED);
55+ if (dtype == ge::DT_INT32) {
56+ return ReduceDims<int32_t>(in_shape, axes_tensor, axes_size, *keep_dims, out_shape);
57+ }
58+ return ReduceDims<int64_t>(in_shape, axes_tensor, axes_size, *keep_dims, out_shape);
59+}
60+ 
61+IMPL_OP_INFERSHAPE(ReduceLogSum).InferShape(InferShape4ReduceCommon).InputsDataDependency({1});
62+} // namespace ops
Amath/reduce_log_sum/op_kernel/arch35/reduce_log_sum_dag.h+49-0
@@ -0,0 +1,49 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file reduce_log_sum_dag.h
13+ * \brief reduce_log_sum_dag
14+ */
15+ 
16+#ifndef CANN_CUSTOM_OPS_REDUCE_LOG_SUM_EXP_DAG_H
17+#define CANN_CUSTOM_OPS_REDUCE_LOG_SUM_EXP_DAG_H
18+ 
19+#ifndef INFINITY
20+#define INFINITY (__builtin_inff())
21+#endif
22+ 
23+#include "atvoss/util/elems.h"
24+#include "atvoss/util/dag.h"
25+#include "atvoss/util/vec.h"
26+#include "atvoss/util/placeholder.h"
27+#include "atvoss/reduce/reduce_operator.h"
28+ 
29+namespace ReduceLogSum {
30+using namespace AscendC;
31+using namespace Ops::Base;
32+ 
33+template <typename T, typename PromteT = float>
34+struct ReduceLogSumDag {
35+ using OpCopyIn0 = Bind<Vec::CopyIn<T>, Placeholder::In0<T>>;
36+ using OpCopyIn0Cast = Bind<Vec::Cast<PromteT, T, 0>, OpCopyIn0>;
37+ 
38+ using ReduceOp0 = Bind<Vec::ReduceSumOp<PromteT>, OpCopyIn0Cast>;
39+ using OpLog = Bind<Vec::Log<PromteT>, ReduceOp0>;
40+ using OpLogCast = Bind<Vec::Cast<T, PromteT, 1>, OpLog>;
41+ 
42+ using OpCopyOut = Bind<Vec::CopyOut<T>, Placeholder::Out0<T>, OpLogCast>;
43+ using Outputs = Elems<OpCopyOut>;
44+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
45+ using OpDag = DAGSch<Outputs, void, MemCfg>;
46+};
47+} // namespace ReduceLogSum
48+ 
49+#endif // CANN_CUSTOM_OPS_REDUCE_LOG_SUM_EXP_DAG_H
Amath/reduce_log_sum/op_kernel/arch35/reduce_log_sum_tiling_key.h+24-0
@@ -0,0 +1,24 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * 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 reduce_log_sum_tiling_key.h
13+ * \brief reduce_log_sum tiling key
14+ */
15+ 
16+#ifndef _REDUCE_LOG_SUM_EXP_TILING_KEY_H_
17+#define _REDUCE_LOG_SUM_EXP_TILING_KEY_H_
18+ 
19+#include "atvoss/reduce/reduce_tiling_key_decl.h"
20+#include "atvoss/reduce/reduce_tiling_key_sel.h"
21+ 
22+ASCENDC_TPL_ARGS_DECL(ReduceLogSum, REDUCE_TPL_KEY_DECL());
23+ 
24+#endif
Amath/reduce_log_sum/op_kernel/reduce_log_sum_apt.cpp+45-0
@@ -0,0 +1,45 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file reduce_log_sum_apt.cpp
13+ * \brief reduce_log_sum
14+ */
15+ 
16+#include "atvoss/reduce/reduce_sch.h"
17+#include "arch35/reduce_log_sum_dag.h"
18+#include "arch35/reduce_log_sum_tiling_key.h"
19+ 
20+using namespace Ops::Base::ReduceOpTmpl;
21+using namespace AscendC;
22+ 
23+template <REDUCE_TPL_PARAM>
24+__global__ __aicore__ void reduce_log_sum(GM_ADDR x, GM_ADDR axes, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
25+{
26+ if (g_coreType == AIC) {
27+ return;
28+ }
29+ if (workspace == nullptr) {
30+ return;
31+ }
32+ SetSysWorkspace(workspace);
33+ GM_ADDR userWS = GetUserWorkspace(workspace);
34+ if (userWS == nullptr) {
35+ return;
36+ }
37+ REGISTER_TILING_DEFAULT(ReduceOpTilingData);
38+ GET_TILING_DATA_WITH_STRUCT(ReduceOpTilingData, tilingData, tiling);
39+ TPipe pipe;
40+ using PromoteType = __reduceType::GetPromoteType<DTYPE_X>::T;
41+ using Op = ReduceSch<REDUCE_TPL_VALUE, ReduceLogSum::ReduceLogSumDag<DTYPE_X, PromoteType>::OpDag>;
42+ Op op(&tilingData);
43+ op.Init(&pipe, x, y, userWS);
44+ op.Process();
45+}
Dmath/reduce_log_sum/tests/CMakeLists.txt+0-15
@@ -1,15 +0,0 @@
1-# This program is free software, you can redistribute it and/or modify it.
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This file is a part of the CANN Open Software.
4-# Licensed under 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, INCLUDING
7-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
8-# the software repository for the full text of the License.
9- 
10-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
11-foreach(SUB_DIR ${CURRENT_DIRS})
12- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
13- add_subdirectory(${SUB_DIR})
14- endif()
15-endforeach()
Dmath/reduce_log_sum/tests/ut/CMakeLists.txt+0-15
@@ -1,15 +0,0 @@
1-# This program is free software, you can redistribute it and/or modify it.
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This file is a part of the CANN Open Software.
4-# Licensed under 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, INCLUDING
7-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
8-# the software repository for the full text of the License.
9- 
10-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
11-foreach(SUB_DIR ${CURRENT_DIRS})
12- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
13- add_subdirectory(${SUB_DIR})
14- endif()
15-endforeach()
Amath/reduce_log_sum/tests/ut/op_api/test_aclnn_reduce_log_sum.cpp+318-0
@@ -0,0 +1,318 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include "math/reduce_log_sum/op_api/aclnn_reduce_log_sum.h"
12+#include <vector>
13+#include <array>
14+#include "gtest/gtest.h"
15+#include "op_api_ut_common/tensor_desc.h"
16+#include "op_api_ut_common/op_api_ut.h"
17+#include "acl/acl.h"
18+ 
19+using namespace std;
20+ 
21+class l2_reduce_log_sum_test : public testing::Test {
22+protected:
23+ static void SetUpTestCase()
24+ {
25+ cout << "reduce_log_sum_test SetUp" << endl;
26+ }
27+ 
28+ static void TearDownTestCase()
29+ {
30+ cout << "reduce_log_sum_test TearDown" << endl;
31+ }
32+};
33+ 
34+// 正常场景 - float16
35+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_float16)
36+{
37+ auto xDesc = TensorDesc({2, 4}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(1, 8);
38+ auto dim = IntArrayDesc(vector<int64_t>{0});
39+ bool keep_dim = true;
40+ bool noopWithEmptyAxes = false;
41+ auto outTensorDesc = TensorDesc({1, 4}, ACL_FLOAT16, ACL_FORMAT_ND).ValidCount(4);
42+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
43+ uint64_t workspaceSize = 0;
44+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
45+ EXPECT_EQ(aclRet, ACL_SUCCESS);
46+}
47+ 
48+// 正常场景 - float32
49+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_float32)
50+{
51+ auto xDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
52+ auto dim = IntArrayDesc(vector<int64_t>{0});
53+ bool keep_dim = true;
54+ bool noopWithEmptyAxes = false;
55+ auto outTensorDesc = TensorDesc({1, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValidCount(4);
56+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
57+ uint64_t workspaceSize = 0;
58+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
59+ EXPECT_EQ(aclRet, ACL_SUCCESS);
60+}
61+ 
62+// 正常场景 - dim为-1
63+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_dim_negative_one)
64+{
65+ auto xDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
66+ auto dim = IntArrayDesc(vector<int64_t>{-1});
67+ bool keep_dim = true;
68+ bool noopWithEmptyAxes = false;
69+ auto outTensorDesc = TensorDesc({2, 1}, ACL_FLOAT, ACL_FORMAT_ND).ValidCount(2);
70+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
71+ uint64_t workspaceSize = 0;
72+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
73+ EXPECT_EQ(aclRet, ACL_SUCCESS);
74+}
75+ 
76+// 正常场景 - keep_dim为false
77+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_keep_dim_false)
78+{
79+ auto xDesc = TensorDesc({2, 3, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
80+ auto dim = IntArrayDesc(vector<int64_t>{-1, 0});
81+ bool keep_dim = false;
82+ bool noopWithEmptyAxes = false;
83+ auto outTensorDesc = TensorDesc({3}, ACL_FLOAT, ACL_FORMAT_ND).ValidCount(3);
84+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
85+ uint64_t workspaceSize = 0;
86+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
87+ EXPECT_EQ(aclRet, ACL_SUCCESS);
88+}
89+ 
90+// 正常场景 - 多维度reduce
91+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_multi_dim_keep_dim_true)
92+{
93+ auto xDesc = TensorDesc({2, 3, 4, 5}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
94+ auto dim = IntArrayDesc(vector<int64_t>{1, 2});
95+ bool keep_dim = true;
96+ bool noopWithEmptyAxes = false;
97+ auto outTensorDesc = TensorDesc({2, 1, 1, 5}, ACL_FLOAT, ACL_FORMAT_ND).ValidCount(10);
98+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
99+ uint64_t workspaceSize = 0;
100+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
101+ EXPECT_EQ(aclRet, ACL_SUCCESS);
102+}
103+ 
104+// 正常场景 - 多维度reduce keep_dim为false
105+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_multi_dim_keep_dim_false)
106+{
107+ auto xDesc = TensorDesc({2, 3, 4, 5}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
108+ auto dim = IntArrayDesc(vector<int64_t>{1, 2});
109+ bool keep_dim = false;
110+ bool noopWithEmptyAxes = false;
111+ auto outTensorDesc = TensorDesc({2, 5}, ACL_FLOAT, ACL_FORMAT_ND).ValidCount(10);
112+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
113+ uint64_t workspaceSize = 0;
114+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
115+ EXPECT_EQ(aclRet, ACL_SUCCESS);
116+}
117+ 
118+// 正常场景 - dim为空 noopWithEmptyAxes为true
119+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_empty_dim_noop_true)
120+{
121+ auto xDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
122+ auto dim = IntArrayDesc(vector<int64_t>{});
123+ bool keep_dim = false;
124+ bool noopWithEmptyAxes = true;
125+ auto outTensorDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValidCount(8);
126+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
127+ uint64_t workspaceSize = 0;
128+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
129+ EXPECT_EQ(aclRet, ACL_SUCCESS);
130+}
131+ 
132+// 正常场景 - dim为空 noopWithEmptyAxes为false
133+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_empty_dim_noop_false)
134+{
135+ auto xDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
136+ auto dim = IntArrayDesc(vector<int64_t>{});
137+ bool keep_dim = false;
138+ bool noopWithEmptyAxes = false;
139+ auto outTensorDesc = TensorDesc({1}, ACL_FLOAT, ACL_FORMAT_ND).ValidCount(1);
140+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
141+ uint64_t workspaceSize = 0;
142+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
143+ EXPECT_EQ(aclRet, ACL_SUCCESS);
144+}
145+ 
146+// 正常场景 - 所有Format
147+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_all_format)
148+{
149+ vector<aclFormat> formats{ACL_FORMAT_ND, ACL_FORMAT_NCHW, ACL_FORMAT_NC, ACL_FORMAT_NCL};
150+ bool keep_dim = true;
151+ bool noopWithEmptyAxes = false;
152+ for (auto format : formats) {
153+ auto xDesc = TensorDesc({2, 2, 2, 3}, ACL_FLOAT, format).ValueRange(-50, 50);
154+ auto dim = IntArrayDesc(vector<int64_t>{0});
155+ auto outTensorDesc = TensorDesc({1, 2, 2, 3}, ACL_FLOAT, format).ValidCount(12);
156+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
157+ uint64_t workspaceSize = 0;
158+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
159+ EXPECT_EQ(aclRet, ACL_SUCCESS);
160+ }
161+}
162+ 
163+// 异常场景 - data为空指针
164+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_nullptr_self)
165+{
166+ auto dim = IntArrayDesc(vector<int64_t>{0});
167+ bool keep_dim = true;
168+ bool noopWithEmptyAxes = false;
169+ auto outDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);
170+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(nullptr, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outDesc));
171+ uint64_t workspaceSize = 0;
172+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
173+ EXPECT_NE(getWorkspaceResult, ACLNN_ERR_INNER_NULLPTR);
174+}
175+ 
176+// 异常场景 - axes为空指针
177+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_nullptr_dim)
178+{
179+ auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);
180+ bool keep_dim = true;
181+ bool noopWithEmptyAxes = false;
182+ auto outDesc = TensorDesc({1, 4}, ACL_FLOAT, ACL_FORMAT_ND);
183+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(selfDesc, nullptr, keep_dim, noopWithEmptyAxes), OUTPUT(outDesc));
184+ uint64_t workspaceSize = 0;
185+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
186+ EXPECT_NE(getWorkspaceResult, ACLNN_ERR_INNER_NULLPTR);
187+}
188+ 
189+// 异常场景 - reduce为空指针
190+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_nullptr_reduce)
191+{
192+ auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);
193+ auto dim = IntArrayDesc(vector<int64_t>{0});
194+ bool keep_dim = true;
195+ bool noopWithEmptyAxes = false;
196+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(selfDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(nullptr));
197+ uint64_t workspaceSize = 0;
198+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
199+ EXPECT_NE(getWorkspaceResult, ACLNN_ERR_INNER_NULLPTR);
200+}
201+ 
202+// 异常场景 - 不支持的self数据类型
203+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_self_dtype_not_support)
204+{
205+ vector<aclDataType> dtypes{ACL_INT64, ACL_INT8, ACL_BOOL, ACL_DOUBLE, ACL_INT32,
206+ ACL_UINT8, ACL_INT16, ACL_COMPLEX64, ACL_COMPLEX128};
207+ for (auto dtype : dtypes) {
208+ auto selfDesc = TensorDesc({2, 4}, dtype, ACL_FORMAT_ND);
209+ auto dim = IntArrayDesc(vector<int64_t>{0});
210+ bool keep_dim = true;
211+ bool noopWithEmptyAxes = false;
212+ auto outDesc = TensorDesc({1, 4}, dtype, ACL_FORMAT_ND);
213+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(selfDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outDesc));
214+ uint64_t workspaceSize = 0;
215+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
216+ EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);
217+ }
218+}
219+ 
220+// 异常场景 - dim超出范围
221+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_dim_out_of_range)
222+{
223+ auto xDesc = TensorDesc({2, 3, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
224+ auto dim = IntArrayDesc(vector<int64_t>{1, 3});
225+ bool keep_dim = true;
226+ bool noopWithEmptyAxes = false;
227+ auto outTensorDesc = TensorDesc({2, 1, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValidCount(8);
228+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
229+ uint64_t workspaceSize = 0;
230+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
231+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
232+}
233+ 
234+// 异常场景 - 维度超过8维
235+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_dim_exceeds_limit)
236+{
237+ auto xDesc = TensorDesc({2, 2, 1, 1, 1, 1, 1, 1, 1, 1}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
238+ auto dim = IntArrayDesc(vector<int64_t>{0});
239+ bool keep_dim = true;
240+ bool noopWithEmptyAxes = false;
241+ auto outTensorDesc = TensorDesc({1, 2, 1, 1, 1, 1, 1, 1, 1, 1}, ACL_FLOAT, ACL_FORMAT_ND).ValidCount(2);
242+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
243+ uint64_t workspaceSize = 0;
244+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
245+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
246+}
247+ 
248+// 异常场景 - reduce shape不匹配
249+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_reduce_shape_not_match)
250+{
251+ auto xDesc = TensorDesc({2, 3, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
252+ auto dim = IntArrayDesc(vector<int64_t>{0});
253+ bool keep_dim = true;
254+ bool noopWithEmptyAxes = false;
255+ auto outTensorDesc = TensorDesc({99, 99}, ACL_FLOAT, ACL_FORMAT_ND).ValidCount(10);
256+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
257+ uint64_t workspaceSize = 0;
258+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
259+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
260+}
261+ 
262+// 边界场景 - 空tensor
263+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_empty_tensor)
264+{
265+ auto xDesc = TensorDesc({2, 0}, ACL_FLOAT, ACL_FORMAT_ND);
266+ auto dim = IntArrayDesc(vector<int64_t>{0});
267+ bool keep_dim = true;
268+ bool noopWithEmptyAxes = false;
269+ auto outTensorDesc = TensorDesc({1, 0}, ACL_FLOAT, ACL_FORMAT_ND);
270+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
271+ uint64_t workspaceSize = 0;
272+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
273+ EXPECT_EQ(aclRet, ACL_SUCCESS);
274+}
275+ 
276+// 边界场景 - 0维tensor
277+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_zero_dim_tensor)
278+{
279+ auto xDesc = TensorDesc({}, ACL_FLOAT, ACL_FORMAT_ND);
280+ auto dim = IntArrayDesc(vector<int64_t>{0});
281+ bool keep_dim = true;
282+ bool noopWithEmptyAxes = false;
283+ auto outTensorDesc = TensorDesc({}, ACL_FLOAT, ACL_FORMAT_ND);
284+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
285+ uint64_t workspaceSize = 0;
286+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
287+ EXPECT_EQ(aclRet, ACL_SUCCESS);
288+}
289+ 
290+// 不连续场景 - self discontinues
291+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_self_discontinues)
292+{
293+ auto xDesc =
294+ TensorDesc({2, 2, 5, 4}, ACL_FLOAT, ACL_FORMAT_ND, {40, 20, 1, 5}, 0, {2, 2, 4, 5}).ValueRange(-50, 50);
295+ auto dim = IntArrayDesc(vector<int64_t>{0});
296+ bool keep_dim = true;
297+ bool noopWithEmptyAxes = false;
298+ auto outTensorDesc = TensorDesc({1, 2, 5, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValidCount(40);
299+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
300+ uint64_t workspaceSize = 0;
301+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
302+ EXPECT_EQ(aclRet, ACL_SUCCESS);
303+}
304+ 
305+// 不连续场景 - out discontinues
306+TEST_F(l2_reduce_log_sum_test, l2_reduce_log_sum_out_discontinues)
307+{
308+ auto xDesc = TensorDesc({2, 2, 5, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-50, 50);
309+ auto dim = IntArrayDesc(vector<int64_t>{0});
310+ bool keep_dim = true;
311+ bool noopWithEmptyAxes = false;
312+ auto outTensorDesc =
313+ TensorDesc({1, 2, 5, 4}, ACL_FLOAT, ACL_FORMAT_ND, {40, 20, 1, 5}, 0, {1, 2, 4, 5}).ValidCount(40);
314+ auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outTensorDesc));
315+ uint64_t workspaceSize = 0;
316+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
317+ EXPECT_EQ(aclRet, ACL_SUCCESS);
318+}
Dmath/reduce_log_sum/tests/ut/op_host/CMakeLists.txt+0-20
@@ -1,20 +0,0 @@
1-# This program is free software, you can redistribute it and/or modify it.
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This file is a part of the CANN Open Software.
4-# Licensed under 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, INCLUDING
7-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
8-# the software repository for the full text of the License.
9- 
10-if(UT_TEST_ALL OR OP_HOST_UT)
11- add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
12- add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
13-endif()
14- 
15-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
16-foreach(SUB_DIR ${CURRENT_DIRS})
17- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
18- add_subdirectory(${SUB_DIR})
19- endif()
20-endforeach()
Amath/reduce_log_sum/tests/ut/op_host/arch35/test_reduce_log_sum_tiling.cpp+31-0
@@ -0,0 +1,31 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <gtest/gtest.h>
13+#include "tiling_context_faker.h"
14+#include "tiling_case_executor.h"
15+#include "atvoss/reduce/reduce_tiling.h"
16+ 
17+using namespace std;
18+using namespace ge;
19+ 
20+class ReduceLogSumTiling : public testing::Test {
21+protected:
22+ static void SetUpTestCase()
23+ {
24+ std::cout << "ReduceLogSumTiling SetUp" << std::endl;
25+ }
26+ 
27+ static void TearDownTestCase()
28+ {
29+ std::cout << "ReduceLogSumTiling TearDown" << std::endl;
30+ }
31+};
Dmath/reduce_log_sum/tests/ut/op_host/op_api/CMakeLists.txt+0-12
@@ -1,12 +0,0 @@
1-# This program is free software, you can redistribute it and/or modify it.
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This file is a part of the CANN Open Software.
4-# Licensed under 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, INCLUDING
7-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
8-# the software repository for the full text of the License.
9- 
10-if(UT_TEST_ALL OR OP_API_UT)
11- add_modules_ut_sources(UT_NAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
12-endif()
Dmath/reduce_log_sum/tests/ut/op_host/op_api/test_reduce_log_sum.cpp+0-316
@@ -1,316 +0,0 @@
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 <float.h>
12-#include <vector>
13-#include <array>
14-#include "gtest/gtest.h"
15- 
16-#include "aclnn_reduce_log_sum.h"
17-#include "op_api_ut_common/tensor_desc.h"
18-#include "op_api_ut_common/op_api_ut.h"
19-#include "opdev/platform.h"
20-#include <cstdlib>
21-#include <ctime>
22- 
23-using namespace std;
24- 
25-class l2_reduce_log_sum_test : public testing::Test {
26-protected:
27- static void SetUpTestCase() {cout << "l2_reduce_log_sum_test SetUp" << endl;}
28- 
29- static void TearDownTestCase() { cout << "l2_reduce_log_sum_test TearDown" << endl; }
30-};
31- 
32-// data为空指针
33-TEST_F(l2_reduce_log_sum_test, case_1) {
34- auto outDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);
35- auto dim = IntArrayDesc(vector<int64_t>{0});
36- bool keep_dim = true;
37- bool noopWithEmpty = false;
38- 
39- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(nullptr, dim, keep_dim, noopWithEmpty), OUTPUT(outDesc));
40- 
41- uint64_t workspaceSize = 0;
42- aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
43- EXPECT_NE(getWorkspaceResult, ACLNN_ERR_INNER_NULLPTR);
44-}
45- 
46-// axes为空指针
47-TEST_F(l2_reduce_log_sum_test, case_2) {
48- auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);
49- auto outDesc = TensorDesc({1, 4}, ACL_FLOAT, ACL_FORMAT_ND);
50- bool keep_dim = true;
51- bool noopWithEmptyAxes = false;
52- 
53- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(selfDesc, nullptr, keep_dim, noopWithEmptyAxes), OUTPUT(outDesc));
54- 
55- uint64_t workspaceSize = 0;
56- aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
57- EXPECT_NE(getWorkspaceResult, ACLNN_ERR_INNER_NULLPTR);
58-}
59- 
60-// reduce为空指针
61-TEST_F(l2_reduce_log_sum_test, case_3) {
62- auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);
63- auto dim = IntArrayDesc(vector<int64_t>{0});
64- bool keep_dim = true;
65- bool noopWithEmptyAxes = false;
66- 
67- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(selfDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(nullptr));
68- 
69- uint64_t workspaceSize = 0;
70- aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
71- EXPECT_NE(getWorkspaceResult, ACLNN_ERR_INNER_NULLPTR);
72-}
73- 
74-// 数据为int64
75-TEST_F(l2_reduce_log_sum_test, case_4) {
76- auto selfDesc = TensorDesc({2, 4}, ACL_INT64, ACL_FORMAT_ND);
77- auto outDesc = TensorDesc({1, 4}, ACL_INT64, ACL_FORMAT_ND);
78- auto dim = IntArrayDesc(vector<int64_t>{0});
79- bool keep_dim = true;
80- bool noopWithEmptyAxes = false;
81- 
82- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(selfDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outDesc));
83- 
84- uint64_t workspaceSize = 0;
85- aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
86- EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);
87-}
88- 
89-// 数据为int8
90-TEST_F(l2_reduce_log_sum_test, case_5) {
91- auto selfDesc = TensorDesc({2, 4}, ACL_INT8, ACL_FORMAT_ND);
92- auto outDesc = TensorDesc({1, 4}, ACL_INT8, ACL_FORMAT_ND);
93- auto dim = IntArrayDesc(vector<int64_t>{0});
94- bool keep_dim = true;
95- bool noopWithEmptyAxes = false;
96- 
97- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(selfDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outDesc));
98- 
99- uint64_t workspaceSize = 0;
100- aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
101- EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);
102-}
103-// 数据为bool
104-TEST_F(l2_reduce_log_sum_test, case_6) {
105- auto selfDesc = TensorDesc({2, 4}, ACL_BOOL, ACL_FORMAT_ND);
106- auto outDesc = TensorDesc({1, 4}, ACL_BOOL, ACL_FORMAT_ND);
107- auto dim = IntArrayDesc(vector<int64_t>{0});
108- bool keep_dim = true;
109- bool noopWithEmptyAxes = false;
110- 
111- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(selfDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outDesc));
112- 
113- uint64_t workspaceSize = 0;
114- aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
115- EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);
116-}
117-// 数据为double
118-TEST_F(l2_reduce_log_sum_test, case_7) {
119- auto selfDesc = TensorDesc({2, 4}, ACL_DOUBLE, ACL_FORMAT_ND);
120- auto outDesc = TensorDesc({1, 4}, ACL_DOUBLE, ACL_FORMAT_ND);
121- auto dim = IntArrayDesc(vector<int64_t>{0});
122- bool keep_dim = true;
123- bool noopWithEmptyAxes = false;
124- 
125- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(selfDesc, dim, keep_dim, noopWithEmptyAxes), OUTPUT(outDesc));
126- 
127- uint64_t workspaceSize = 0;
128- aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
129- EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);
130-}
131- 
132-// float16
133-TEST_F(l2_reduce_log_sum_test, case_8) {
134- auto xDesc = TensorDesc({2, 4}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(1, 8);
135- auto dim = IntArrayDesc(vector<int64_t>{0});
136- bool keep_dim = true;
137- bool noopWithEmptyAxes = false;
138- aclDataType dType = ACL_FLOAT16;
139- const vector<int64_t>& outShape = {1, 4};
140- auto outTensorDesc = TensorDesc(outShape, dType, ACL_FORMAT_ND).ValidCount(4);
141- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes),
142- OUTPUT(outTensorDesc));
143- 
144- uint64_t workspaceSize = 0;
145- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
146- EXPECT_EQ(aclRet, ACL_SUCCESS);
147-}
148- 
149- 
150-// float32
151-TEST_F(l2_reduce_log_sum_test, case_9) {
152- auto xDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
153- auto dim = IntArrayDesc(vector<int64_t>{0});
154- bool keep_dim = true;
155- bool noopWithEmptyAxes = false;
156- aclDataType dType = ACL_FLOAT;
157- const vector<int64_t>& outShape = {1, 4};
158- auto outTensorDesc = TensorDesc(outShape, dType, ACL_FORMAT_ND).ValidCount(4);
159- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes),
160- OUTPUT(outTensorDesc));
161- 
162- uint64_t workspaceSize = 0;
163- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
164- EXPECT_EQ(aclRet, ACL_SUCCESS);
165-}
166-// float32 dim为-1
167-TEST_F(l2_reduce_log_sum_test, case_10) {
168- auto xDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
169- auto dim = IntArrayDesc(vector<int64_t>{-1});
170- bool keep_dim = true;
171- bool noopWithEmptyAxes = false;
172- aclDataType dType = ACL_FLOAT;
173- const vector<int64_t>& outShape = {2, 1};
174- auto outTensorDesc = TensorDesc(outShape, dType, ACL_FORMAT_ND).ValidCount(2);
175- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes),
176- OUTPUT(outTensorDesc));
177- 
178- uint64_t workspaceSize = 0;
179- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
180- EXPECT_EQ(aclRet, ACL_SUCCESS);
181-}
182- 
183-// float32 keep_dim为false
184-TEST_F(l2_reduce_log_sum_test, case_11) {
185- auto xDesc = TensorDesc({2, 3, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
186- auto dim = IntArrayDesc(vector<int64_t>{-1, 0});
187- bool keep_dim = false;
188- bool noopWithEmptyAxes = false;
189- aclDataType dType = ACL_FLOAT;
190- const vector<int64_t>& outShape = {3};
191- auto outTensorDesc = TensorDesc(outShape, dType, ACL_FORMAT_ND).ValidCount(3);
192- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes),
193- OUTPUT(outTensorDesc));
194- 
195- uint64_t workspaceSize = 0;
196- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
197- EXPECT_EQ(aclRet, ACL_SUCCESS);
198-}
199-// float32 dim为空，noopWithEmptyAxes为true
200-TEST_F(l2_reduce_log_sum_test, case_12) {
201- auto xDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
202- auto dim = IntArrayDesc(vector<int64_t>{});
203- bool keep_dim = false;
204- bool noopWithEmptyAxes = true;
205- aclDataType dType = ACL_FLOAT;
206- const vector<int64_t>& outShape = {2, 4};
207- auto outTensorDesc = TensorDesc(outShape, dType, ACL_FORMAT_ND).ValidCount(8);
208- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes),
209- OUTPUT(outTensorDesc));
210- 
211- uint64_t workspaceSize = 0;
212- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
213- EXPECT_EQ(aclRet, ACL_SUCCESS);
214-}
215- 
216-// float32 dim为空，noopWithEmptyAxes为false
217-TEST_F(l2_reduce_log_sum_test, case_13) {
218- auto xDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
219- auto dim = IntArrayDesc(vector<int64_t>{});
220- bool keep_dim = false;
221- bool noopWithEmptyAxes = false;
222- aclDataType dType = ACL_FLOAT;
223- const vector<int64_t>& outShape = {1};
224- auto outTensorDesc = TensorDesc(outShape, dType, ACL_FORMAT_ND).ValidCount(1);
225- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes),
226- OUTPUT(outTensorDesc));
227- 
228- uint64_t workspaceSize = 0;
229- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
230- EXPECT_EQ(aclRet, ACL_SUCCESS);
231-}
232- 
233-//dim重复
234-TEST_F(l2_reduce_log_sum_test, case_14) {
235- auto xDesc = TensorDesc({2, 3, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
236- auto dim = IntArrayDesc(vector<int64_t>{1, 1});
237- bool keep_dim = true;
238- bool noopWithEmptyAxes = false;
239- aclDataType dType = ACL_FLOAT;
240- const vector<int64_t>& outShape = {};
241- auto outTensorDesc = TensorDesc(outShape, dType, ACL_FORMAT_ND).ValidCount(1);
242- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes),
243- OUTPUT(outTensorDesc));
244- 
245- uint64_t workspaceSize = 0;
246- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
247- EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
248-}
249- 
250-//dim超出范围
251-TEST_F(l2_reduce_log_sum_test, case_15) {
252- auto xDesc = TensorDesc({2, 3, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
253- auto dim = IntArrayDesc(vector<int64_t>{1, 3});
254- bool keep_dim = true;
255- bool noopWithEmptyAxes = false;
256- aclDataType dType = ACL_FLOAT;
257- const vector<int64_t>& outShape = {};
258- auto outTensorDesc = TensorDesc(outShape, dType, ACL_FORMAT_ND).ValidCount(1);
259- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes),
260- OUTPUT(outTensorDesc));
261- 
262- uint64_t workspaceSize = 0;
263- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
264- EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
265-}
266- 
267-//维度超出范围
268-TEST_F(l2_reduce_log_sum_test, case_16) {
269- auto xDesc = TensorDesc({2, 2, 1, 1, 1, 1, 1, 1, 1, 1}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(1, 8);
270- auto dim = IntArrayDesc(vector<int64_t>{0});
271- bool keep_dim = true;
272- bool noopWithEmptyAxes = false;
273- aclDataType dType = ACL_FLOAT;
274- const vector<int64_t>& outShape = {1, 2, 1, 1, 1, 1, 1, 1, 1, 1};
275- auto outTensorDesc = TensorDesc(outShape, dType, ACL_FORMAT_ND).ValidCount(1);
276- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes),
277- OUTPUT(outTensorDesc));
278- 
279- uint64_t workspaceSize = 0;
280- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
281- EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
282-}
283- 
284-//self为空tensor
285-TEST_F(l2_reduce_log_sum_test, case_17) {
286- auto xDesc = TensorDesc({2, 0}, ACL_FLOAT, ACL_FORMAT_ND);
287- auto dim = IntArrayDesc(vector<int64_t>{0});
288- bool keep_dim = true;
289- bool noopWithEmptyAxes = false;
290- aclDataType dType = ACL_FLOAT;
291- const vector<int64_t>& outShape = {1, 0};
292- auto outTensorDesc = TensorDesc(outShape, dType, ACL_FORMAT_ND);
293- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes),
294- OUTPUT(outTensorDesc));
295- 
296- uint64_t workspaceSize = 0;
297- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
298- EXPECT_EQ(aclRet, ACL_SUCCESS);
299-}
300- 
301-// self为0维tensor
302-TEST_F(l2_reduce_log_sum_test, case_18) {
303- auto xDesc = TensorDesc({}, ACL_FLOAT, ACL_FORMAT_ND);
304- auto dim = IntArrayDesc(vector<int64_t>{0});
305- bool keep_dim = true;
306- bool noopWithEmptyAxes = false;
307- aclDataType dType = ACL_FLOAT;
308- const vector<int64_t>& outShape = {};
309- auto outTensorDesc = TensorDesc(outShape, dType, ACL_FORMAT_ND);
310- auto ut = OP_API_UT(aclnnReduceLogSum, INPUT(xDesc, dim, keep_dim, noopWithEmptyAxes),
311- OUTPUT(outTensorDesc));
312- 
313- uint64_t workspaceSize = 0;
314- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
315- EXPECT_EQ(aclRet, ACL_SUCCESS);
316-}
Amath/reduce_log_sum/tests/ut/op_host/test_reduce_log_sum_infershape.cpp+207-0
@@ -0,0 +1,207 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <gtest/gtest.h>
12+#include <iostream>
13+#include "infershape_context_faker.h"
14+#include "infershape_case_executor.h"
15+ 
16+class ReduceLogSumInferShape : public testing::Test {
17+protected:
18+ static void SetUpTestCase()
19+ {
20+ std::cout << "ReduceLogSumInferShape SetUp" << std::endl;
21+ }
22+ 
23+ static void TearDownTestCase()
24+ {
25+ std::cout << "ReduceLogSumInferShape TearDown" << std::endl;
26+ }
27+};
28+ 
29+// 单轴reduce, keep_dims=true
30+TEST_F(ReduceLogSumInferShape, reduce_log_sum_infershape_test_0)
31+{
32+ std::vector<int64_t> axesValue = {0};
33+ gert::InfershapeContextPara infershapeContextPara(
34+ "ReduceLogSum",
35+ {
36+ {{{2, 4}, {2, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
37+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, axesValue.data()},
38+ },
39+ {
40+ {{{1, 4}, {1, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
41+ },
42+ {gert::InfershapeContextPara::OpAttr("keep_dims", Ops::Math::AnyValue::CreateFrom<bool>(true))});
43+ std::vector<std::vector<int64_t>> expectOutputShape = {
44+ {1, 4},
45+ };
46+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
47+}
48+ 
49+// 多轴reduce, keep_dims=true
50+TEST_F(ReduceLogSumInferShape, reduce_log_sum_infershape_test_1)
51+{
52+ std::vector<int64_t> axesValue = {0, 1};
53+ gert::InfershapeContextPara infershapeContextPara(
54+ "ReduceLogSum",
55+ {
56+ {{{2, 3, 4}, {2, 3, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
57+ {{{2}, {2}}, ge::DT_INT64, ge::FORMAT_ND, true, axesValue.data()},
58+ },
59+ {
60+ {{{1, 1, 4}, {1, 1, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
61+ },
62+ {gert::InfershapeContextPara::OpAttr("keep_dims", Ops::Math::AnyValue::CreateFrom<bool>(true))});
63+ std::vector<std::vector<int64_t>> expectOutputShape = {
64+ {1, 1, 4},
65+ };
66+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
67+}
68+ 
69+// 负轴reduce, keep_dims=false
70+TEST_F(ReduceLogSumInferShape, reduce_log_sum_infershape_test_2)
71+{
72+ std::vector<int64_t> axesValue = {-1};
73+ gert::InfershapeContextPara infershapeContextPara(
74+ "ReduceLogSum",
75+ {
76+ {{{2, 3, 4}, {2, 3, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
77+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, axesValue.data()},
78+ },
79+ {
80+ {{{2, 3}, {2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
81+ },
82+ {gert::InfershapeContextPara::OpAttr("keep_dims", Ops::Math::AnyValue::CreateFrom<bool>(false))});
83+ std::vector<std::vector<int64_t>> expectOutputShape = {
84+ {2, 3},
85+ };
86+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
87+}
88+ 
89+// FLOAT16
90+TEST_F(ReduceLogSumInferShape, reduce_log_sum_infershape_test_3)
91+{
92+ std::vector<int64_t> axesValue = {0};
93+ gert::InfershapeContextPara infershapeContextPara(
94+ "ReduceLogSum",
95+ {
96+ {{{2, 4}, {2, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
97+ {{{1}, {1}}, ge::DT_INT32, ge::FORMAT_ND, true, axesValue.data()},
98+ },
99+ {
100+ {{{1, 4}, {1, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
101+ },
102+ {gert::InfershapeContextPara::OpAttr("keep_dims", Ops::Math::AnyValue::CreateFrom<bool>(true))});
103+ std::vector<std::vector<int64_t>> expectOutputShape = {
104+ {1, 4},
105+ };
106+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
107+}
108+ 
109+// 空axes (保持原shape)
110+TEST_F(ReduceLogSumInferShape, reduce_log_sum_infershape_test_4)
111+{
112+ std::vector<int64_t> axesValue = {};
113+ gert::InfershapeContextPara infershapeContextPara(
114+ "ReduceLogSum",
115+ {
116+ {{{2, 3, 4, 5}, {2, 3, 4, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
117+ {{{0}, {0}}, ge::DT_INT64, ge::FORMAT_ND, true, axesValue.data()},
118+ },
119+ {
120+ {{{2, 3, 4, 5}, {2, 3, 4, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
121+ },
122+ {gert::InfershapeContextPara::OpAttr("keep_dims", Ops::Math::AnyValue::CreateFrom<bool>(false))});
123+ std::vector<std::vector<int64_t>> expectOutputShape = {
124+ {2, 3, 4, 5},
125+ };
126+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
127+}
128+ 
129+// 动态shape, keep_dims=true
130+TEST_F(ReduceLogSumInferShape, reduce_log_sum_infershape_test_5)
131+{
132+ std::vector<int64_t> axesValue = {1};
133+ gert::InfershapeContextPara infershapeContextPara(
134+ "ReduceLogSum",
135+ {
136+ {{{-1, 2}, {-1, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
137+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, axesValue.data()},
138+ },
139+ {
140+ {{{-1, 1}, {-1, 1}}, ge::DT_FLOAT, ge::FORMAT_ND},
141+ },
142+ {gert::InfershapeContextPara::OpAttr("keep_dims", Ops::Math::AnyValue::CreateFrom<bool>(true))});
143+ std::vector<std::vector<int64_t>> expectOutputShape = {
144+ {-1, 1},
145+ };
146+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
147+}
148+ 
149+// BF16, 多轴, keep_dims=true
150+TEST_F(ReduceLogSumInferShape, reduce_log_sum_infershape_test_6)
151+{
152+ std::vector<int64_t> axesValue = {2, 4};
153+ gert::InfershapeContextPara infershapeContextPara(
154+ "ReduceLogSum",
155+ {
156+ {{{1, 2, 3, 4, 5}, {1, 2, 3, 4, 5}}, ge::DT_BF16, ge::FORMAT_ND},
157+ {{{2}, {2}}, ge::DT_INT64, ge::FORMAT_ND, true, axesValue.data()},
158+ },
159+ {
160+ {{{1, 2, 1, 4, 1}, {1, 2, 1, 4, 1}}, ge::DT_BF16, ge::FORMAT_ND},
161+ },
162+ {gert::InfershapeContextPara::OpAttr("keep_dims", Ops::Math::AnyValue::CreateFrom<bool>(true))});
163+ std::vector<std::vector<int64_t>> expectOutputShape = {
164+ {1, 2, 1, 4, 1},
165+ };
166+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
167+}
168+ 
169+// 多轴, keep_dims=false
170+TEST_F(ReduceLogSumInferShape, reduce_log_sum_infershape_test_7)
171+{
172+ std::vector<int64_t> axesValue = {2, 4};
173+ gert::InfershapeContextPara infershapeContextPara(
174+ "ReduceLogSum",
175+ {
176+ {{{1, 2, 3, 4, 5}, {1, 2, 3, 4, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
177+ {{{2}, {2}}, ge::DT_INT64, ge::FORMAT_ND, true, axesValue.data()},
178+ },
179+ {
180+ {{{1, 2, 4}, {1, 2, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
181+ },
182+ {gert::InfershapeContextPara::OpAttr("keep_dims", Ops::Math::AnyValue::CreateFrom<bool>(false))});
183+ std::vector<std::vector<int64_t>> expectOutputShape = {
184+ {1, 2, 4},
185+ };
186+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
187+}
188+ 
189+// 全动态shape(-1), keep_dims=true
190+TEST_F(ReduceLogSumInferShape, reduce_log_sum_infershape_test_8)
191+{
192+ std::vector<int64_t> axesValue = {0};
193+ gert::InfershapeContextPara infershapeContextPara(
194+ "ReduceLogSum",
195+ {
196+ {{{-1}, {-1}}, ge::DT_FLOAT, ge::FORMAT_ND},
197+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, axesValue.data()},
198+ },
199+ {
200+ {{{1}, {1}}, ge::DT_FLOAT, ge::FORMAT_ND},
201+ },
202+ {gert::InfershapeContextPara::OpAttr("keep_dims", Ops::Math::AnyValue::CreateFrom<bool>(true))});
203+ std::vector<std::vector<int64_t>> expectOutputShape = {
204+ {1},
205+ };
206+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
207+}
Amath/reduce_log_sum/tests/ut/op_kernel/arch35/test_reduce_log_sum.cpp+183-0
@@ -0,0 +1,183 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file test_reduce_log_sum.cpp
13+ * \brief reduce_log_sum opkernel unit test
14+ */
15+ 
16+#include <array>
17+#include <vector>
18+#include <string>
19+#include <cstdint>
20+#include "gtest/gtest.h"
21+#include "tikicpulib.h"
22+#include "reduce_log_sum_apt.cpp"
23+ 
24+using namespace std;
25+ 
26+class ReduceLogSumKernel : public testing::Test {
27+protected:
28+ static void SetUpTestCase()
29+ {
30+ cout << "ReduceLogSumKernel SetUp\n" << endl;
31+ }
32+ static void TearDownTestCase()
33+ {
34+ cout << "ReduceLogSumKernel TearDown\n" << endl;
35+ }
36+};
37+ 
38+// float32, 4x64 reduce axis=1 -> output 4
39+TEST_F(ReduceLogSumKernel, test_case_float32_4x64)
40+{
41+ uint64_t tilingKey = 0;
42+ uint32_t numBlocks = 4;
43+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
44+ 
45+ size_t xSize = 4 * 64 * sizeof(float);
46+ size_t ySize = 4 * sizeof(float);
47+ size_t workspaceFileSize = 16 * 1024 * 1024;
48+ 
49+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(xSize);
50+ uint8_t* axes = (uint8_t*)AscendC::GmAlloc(sizeof(int32_t));
51+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(ySize);
52+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceFileSize);
53+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(ReduceOpTilingData));
54+ 
55+ Ops::Base::ReduceOpTilingData* tilingData = reinterpret_cast<Ops::Base::ReduceOpTilingData*>(tiling);
56+ tilingData->factorACntPerCore = 1;
57+ tilingData->factorATotalCnt = 4;
58+ tilingData->ubFactorA = 1;
59+ tilingData->factorRCntPerCore = 1;
60+ tilingData->factorRTotalCnt = 1;
61+ tilingData->ubFactorR = 1;
62+ tilingData->groupR = 1;
63+ tilingData->outSize = 4;
64+ tilingData->basicBlock = 51200;
65+ tilingData->coreNum = 64;
66+ tilingData->meanVar = 0.015625f;
67+ tilingData->shape[0] = 4;
68+ tilingData->shape[1] = 64;
69+ tilingData->stride[0] = 64;
70+ tilingData->stride[1] = 1;
71+ tilingData->dstStride[0] = 1;
72+ tilingData->dstStride[1] = 1;
73+ 
74+ ICPU_SET_TILING_KEY(tilingKey);
75+ auto reduce_log_sum_func = [](GM_ADDR x, GM_ADDR axes, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
76+ reduce_log_sum<true, 0, 1, 1>(x, axes, y, workspace, tiling);
77+ };
78+ ICPU_RUN_KF(reduce_log_sum_func, numBlocks, x, axes, y, workspace, tiling);
79+ 
80+ AscendC::GmFree((void*)x);
81+ AscendC::GmFree((void*)axes);
82+ AscendC::GmFree((void*)y);
83+ AscendC::GmFree((void*)workspace);
84+ AscendC::GmFree((void*)tiling);
85+}
86+ 
87+// float16, 2x4 reduce axis=0 -> output 4 (keep_dims)
88+TEST_F(ReduceLogSumKernel, test_case_float16_2x4)
89+{
90+ uint64_t tilingKey = 0;
91+ uint32_t numBlocks = 2;
92+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
93+ 
94+ size_t xSize = 2 * 4 * sizeof(uint16_t);
95+ size_t ySize = 1 * 4 * sizeof(uint16_t);
96+ size_t workspaceFileSize = 16 * 1024 * 1024;
97+ 
98+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(xSize);
99+ uint8_t* axes = (uint8_t*)AscendC::GmAlloc(sizeof(int32_t));
100+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(ySize);
101+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceFileSize);
102+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(ReduceOpTilingData));
103+ 
104+ Ops::Base::ReduceOpTilingData* tilingData = reinterpret_cast<Ops::Base::ReduceOpTilingData*>(tiling);
105+ tilingData->factorACntPerCore = 1;
106+ tilingData->factorATotalCnt = 1;
107+ tilingData->ubFactorA = 1;
108+ tilingData->factorRCntPerCore = 1;
109+ tilingData->factorRTotalCnt = 1;
110+ tilingData->ubFactorR = 1;
111+ tilingData->groupR = 1;
112+ tilingData->outSize = 4;
113+ tilingData->basicBlock = 51200;
114+ tilingData->coreNum = 64;
115+ tilingData->meanVar = 0.015625f;
116+ tilingData->shape[0] = 2;
117+ tilingData->shape[1] = 4;
118+ tilingData->stride[0] = 4;
119+ tilingData->stride[1] = 1;
120+ tilingData->dstStride[0] = 1;
121+ tilingData->dstStride[1] = 1;
122+ 
123+ ICPU_SET_TILING_KEY(tilingKey);
124+ auto reduce_log_sum_func = [](GM_ADDR x, GM_ADDR axes, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
125+ reduce_log_sum<true, 0, 1, 1>(x, axes, y, workspace, tiling);
126+ };
127+ ICPU_RUN_KF(reduce_log_sum_func, numBlocks, x, axes, y, workspace, tiling);
128+ 
129+ AscendC::GmFree((void*)x);
130+ AscendC::GmFree((void*)axes);
131+ AscendC::GmFree((void*)y);
132+ AscendC::GmFree((void*)workspace);
133+ AscendC::GmFree((void*)tiling);
134+}
135+ 
136+// float32, 16x256 reduce axis=0 -> output 256
137+TEST_F(ReduceLogSumKernel, test_case_float32_16x256)
138+{
139+ uint64_t tilingKey = 0;
140+ uint32_t numBlocks = 16;
141+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
142+ 
143+ size_t xSize = 16 * 256 * sizeof(float);
144+ size_t ySize = 1 * 256 * sizeof(float);
145+ size_t workspaceFileSize = 16 * 1024 * 1024;
146+ 
147+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(xSize);
148+ uint8_t* axes = (uint8_t*)AscendC::GmAlloc(sizeof(int32_t));
149+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(ySize);
150+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceFileSize);
151+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(ReduceOpTilingData));
152+ 
153+ Ops::Base::ReduceOpTilingData* tilingData = reinterpret_cast<Ops::Base::ReduceOpTilingData*>(tiling);
154+ tilingData->factorACntPerCore = 1;
155+ tilingData->factorATotalCnt = 1;
156+ tilingData->ubFactorA = 1;
157+ tilingData->factorRCntPerCore = 1;
158+ tilingData->factorRTotalCnt = 1;
159+ tilingData->ubFactorR = 1;
160+ tilingData->groupR = 1;
161+ tilingData->outSize = 256;
162+ tilingData->basicBlock = 51200;
163+ tilingData->coreNum = 64;
164+ tilingData->meanVar = 0.015625f;
165+ tilingData->shape[0] = 16;
166+ tilingData->shape[1] = 256;
167+ tilingData->stride[0] = 256;
168+ tilingData->stride[1] = 1;
169+ tilingData->dstStride[0] = 1;
170+ tilingData->dstStride[1] = 1;
171+ 
172+ ICPU_SET_TILING_KEY(tilingKey);
173+ auto reduce_log_sum_func = [](GM_ADDR x, GM_ADDR axes, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
174+ reduce_log_sum<true, 0, 1, 1>(x, axes, y, workspace, tiling);
175+ };
176+ ICPU_RUN_KF(reduce_log_sum_func, numBlocks, x, axes, y, workspace, tiling);
177+ 
178+ AscendC::GmFree((void*)x);
179+ AscendC::GmFree((void*)axes);
180+ AscendC::GmFree((void*)y);
181+ AscendC::GmFree((void*)workspace);
182+ AscendC::GmFree((void*)tiling);
183+}
Ascripts/ci/ascend950/ops_math_operator_list.yaml+209-0
@@ -0,0 +1,209 @@
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+# 分组编译，当前支持5个分组，需将每个算子配置到对应的分组里，此配置文件仅仅在CI上编译出包有效
12+# 配置方式：math仓当前有三个子目录，分别为math、conversion、random，将对应目录内的算子按照字母表的顺序从前往后配置即可
13+operator_group_1:
14+ - signbit
15+ - abs
16+ - acos
17+ - add
18+ - addcdiv
19+ - addcmul
20+ - add_n
21+ - addr
22+ - adjacent_difference
23+ - arg_max_v2
24+ - arg_max_with_value
25+ - arg_min
26+ - arg_min_with_value
27+ - assign_add
28+ - assign_sub
29+ - asin
30+ - atan
31+ - axpy
32+ - axpy_v2
33+ - bias_add
34+ - bias_add_grad
35+ - bincount
36+ - bitwise_and
37+ - bitwise_or
38+ - bitwise_xor
39+ - cast
40+ - cdist
41+ - ceil
42+ - cholesky
43+ - cos
44+ - cosh
45+ - cumsum
46+ - cummin
47+ - diag
48+ - diag_flat
49+ - diag_v2
50+ - diag_part
51+ - div
52+ - div_no_nan
53+ - dot
54+ - dynamic_partition
55+ - dynamic_stitch
56+ - equal
57+ - erf
58+ - exp
59+ - expand
60+ - expm1
61+ - exp_segsum_grad
62+ - eye
63+ - floor
64+ - floor_div
65+ - floor_mod
66+ - fused_mul_add_n
67+ - gcd
68+ - greater
69+ - ger
70+ - greater_equal
71+ - grouped_bias_add_grad
72+ - histogram_v2
73+ - invert
74+ - kl_div_v2
75+ - pad
76+ - tile
77+ - is_close
78+ - is_finite
79+ - is_inf
80+ - is_nan
81+ - is_neg_inf
82+ - is_pos_inf
83+operator_group_2:
84+ - lerp
85+ - less
86+ - less_equal
87+ - lin_space
88+ - log1p
89+ - log
90+ - logical_and
91+ - logical_not
92+ - logical_or
93+ - masked_scale
94+ - maximum
95+ - minimum
96+ - mod
97+ - mul
98+ - muls
99+ - mul_addn
100+ - nan_to_num
101+ - neg
102+ - non_finite_check
103+ - not_equal
104+ - one_hot
105+ - ones_like
106+ - pow
107+ - pows
108+ - range
109+ - real_div
110+ - reciprocal
111+ - reduce_max
112+ - reduce_mean
113+ - reduce_min
114+ - reduce_prod
115+operator_group_3:
116+ - reduce_std_v2
117+ - reduce_sum
118+ - reduce_var
119+ - square_sum_v1
120+ - reduce_all
121+ - reduce_any
122+ - reduce_log_sum
123+ - reduce_log_sum_exp
124+operator_group_4:
125+ - rfft1_d
126+ - right_shift
127+ - round
128+ - rsqrt
129+ - rsqrt_grad
130+ - select
131+ - select_v2
132+ - sign
133+ - sin
134+ - sinkhorn
135+ - sort
136+ - sort_with_index
137+ - sqrt
138+ - sqrt_grad
139+ - square
140+ - squared_difference
141+ - stateless_randperm # 复用sort。要写在同一分组下，否则CI编译kernel报错。
142+ - sub
143+ - tanh
144+ - tanh_grad
145+ - tensor_equal
146+ - top_k_v2
147+operator_group_5:
148+ - transform_bias_rescale_qkv
149+ - trunc
150+ - assign
151+ - as_strided
152+ - batch_to_space_nd
153+ - broadcast_to
154+ - clip_by_value
155+ - clip_by_value_v2
156+ - concat
157+ - concat_d
158+ - concat_v2
159+ - confusion_transpose_d
160+ - depth_to_space
161+ - fill
162+ - fills
163+ - masked_fill
164+ - masked_select_v3
165+ - matrix_diag
166+ - matrix_set_diag
167+ - im2col
168+ - mem_set
169+ - mem_set_v2
170+ - mirror_pad
171+ - pack
172+ - pad_v2
173+ - pad_v3
174+ - pad_v3_grad
175+ - reflection_pad3d_grad
176+ - roll
177+ - slice
178+ - space_to_depth
179+ - split
180+ - split_v
181+ - strided_slice
182+ - strided_slice_v2
183+ - strided_slice_v3
184+ - strided_slice_assign_v2
185+ - strided_slice_grad
186+ - tensor_move
187+ - trans_data
188+ - transpose
189+ - tril
190+ - triu
191+ - unfold_grad
192+ - unpack
193+ - zeros_like
194+ - drop_out_do_mask
195+ - drop_out_do_mask_v3
196+ - drop_out_do_mask_v3_d
197+ - drop_out_v3
198+ - random_standard_normal_v2
199+ - random_uniform_v2
200+ - random_uniform_int_v2
201+ - stateless_bernoulli
202+ - stateless_drop_out_gen_mask
203+ - stateless_random_normal_v2
204+ - stateless_random_uniform_v2
205+ - truncated_normal_v2
206+ - view_copy
207+ - stateless_random_choice_with_mask
208+ - stateless_random_uniform_v3
209+ - stateless_random_normal_v3