Star1399
1485
代码
代码
Issues127
Pull Requests316
Actions
项目讨论
Wiki
分析
Star1399
1485
已合并
个人-AscendC实现Sqrt算子贡献 #93
liuxiqiang创建于 2025年10月23日
个人-AscendC实现Sqrt算子贡献 #93
已合并
liuxiqiang创建于 2025年10月23日
从liuxiqiang/ops-math_10222master合入到cann/ops-mathmaster
共 30 个文件变更+1478-1
Aexperimental/math/sqrt/CMakeLists.txt+20-0
@@ -0,0 +1,20 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify.
S
Ssunday2025年10月28日
likedislike
L
liuxiqiang
2025年10月28日 评论：
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+if(NOT ENABLE_TEST)
14+ list(REMOVE_ITEM CURRENT_DIRS tests)
15+endif()
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()
Aexperimental/math/sqrt/examples/test_aclnn_sqrt.cpp+119-0
@@ -0,0 +1,119 @@
1+#include <iostream>
2+#include <vector>
3+#include "acl/acl.h"
4+#include "aclnn_sqrt.h"
5+ 
6+#define CHECK_RET(cond, return_expr) \
7+ do { \
8+ if (!(cond)) { \
9+ return_expr; \
10+ } \
11+ } while (0)
12+#define LOG_PRINT(message, ...) \
13+ do { \
14+ printf(message, ##__VA_ARGS__); \
15+ } while (0)
16+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
17+ int64_t shapeSize = 1;
18+ for (auto i : shape) {
19+ shapeSize *= i;
20+ }
21+ return shapeSize;
22+}
23+int Init(int32_t deviceId, aclrtStream* stream) {
24+ // 固定写法，资源初始化
25+ auto ret = aclInit(nullptr);
26+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
27+ ret = aclrtSetDevice(deviceId);
28+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
29+ ret = aclrtCreateStream(stream);
30+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
31+ return 0;
32+}
33+ 
34+template <typename T>
35+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
36+ aclDataType dataType, aclTensor** tensor) {
37+ auto size = GetShapeSize(shape) * sizeof(T);
38+ // 调用aclrtMalloc申请device侧内存
39+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
40+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
41+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
42+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
43+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
44+ // 计算连续tensor的strides
45+ std::vector<int64_t> strides(shape.size(), 1);
46+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
47+ strides[i] = shape[i + 1] * strides[i + 1];
48+ }
49+ // 调用aclCreateTensor接口创建aclTensor
50+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
51+ shape.data(), shape.size(), *deviceAddr);
52+ return 0;
53+}
54+int main() {
55+ // 1. （固定写法）device/stream初始化，参考acl API手册
56+ // 根据自己的实际device填写deviceId
57+ int32_t deviceId = 0;
58+ aclrtStream stream;
59+ auto ret = Init(deviceId, &stream);
60+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
61+ // 2. 构造输入与输出，需要根据API的接口自定义构造
62+ std::vector<int64_t> selfShape = {2, 2};
63+ std::vector<int64_t> outShape = {2, 2};
64+ void* selfDeviceAddr = nullptr;
65+ void* outDeviceAddr = nullptr;
66+ aclTensor* self = nullptr;
67+ aclTensor* out = nullptr;
68+ std::vector<float> selfHostData = {0, 1, 2, 3};
69+ std::vector<float> outHostData = {0, 0, 0, 0};
70+ // 创建self aclTensor
71+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
72+ CHECK_RET(ret == ACL_SUCCESS, return ret);
73+ // 创建out aclTensor
74+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
75+ CHECK_RET(ret == ACL_SUCCESS, return ret);
76+ // 3. 调用CANN算子库API，需要修改为具体的API名称
77+ // aclnnSqrt接口调用示例
78+ uint64_t workspaceSize = 0;
79+ aclOpExecutor* executor;
80+ // 调用aclnnSqrt第一段接口
81+ ret = aclnnSqrtGetWorkspaceSize(self, out, &workspaceSize, &executor);
82+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSqrtGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
83+ // 根据第一段接口计算出的workspaceSize申请device内存
84+ void* workspaceAddr = nullptr;
85+ if (workspaceSize > 0) {
86+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
87+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
88+ }
89+ // 调用aclnnSqrt第二段接口
90+ ret = aclnnSqrt(workspaceAddr, workspaceSize, executor, stream);
91+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSqrt failed. ERROR: %d\n", ret); return ret);
92+ // 4. （固定写法）同步等待任务执行结束
93+ ret = aclrtSynchronizeStream(stream);
94+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
95+ 
96+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
97+ auto size = GetShapeSize(outShape);
98+ std::vector<float> resultData(size, 0);
99+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
100+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
101+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
102+ for (int64_t i = 0; i < size; i++) {
103+ LOG_PRINT("aclnnSqrt result[%ld] is: %f\n", i, resultData[i]);
104+ }
105+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
106+ aclDestroyTensor(self);
107+ aclDestroyTensor(out);
108+ 
109+ // 7. 释放device资源，需要根据具体API的接口定义修改
110+ aclrtFree(selfDeviceAddr);
111+ aclrtFree(outDeviceAddr);
112+ if (workspaceSize > 0) {
113+ aclrtFree(workspaceAddr);
114+ }
115+ aclrtDestroyStream(stream);
116+ aclrtResetDevice(deviceId);
117+ aclFinalize();
118+ return 0;
119+}
Aexperimental/math/sqrt/op_graph/CMakeLists.txt+12-0
@@ -0,0 +1,12 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify.
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+add_graph_plugin_sources()
Aexperimental/math/sqrt/op_graph/fusion_pass/CMakeLists.txt+10-0
@@ -0,0 +1,10 @@
1+# This program is free software, you can redistribute it and/or modify.
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.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ============================================================================
10+ 
Aexperimental/math/sqrt/op_graph/sqrt_graph_infer.cpp+38-0
@@ -0,0 +1,38 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This file is a part of the CANN Open Software.
5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file sqrt_graph_infer.cpp
14+ * \brief sqrt operater graph infer resource
15+ */
16+#include "register/op_impl_registry.h"
17+#include "log/log.h"
18+ 
19+namespace ops {
20+using namespace ge;
21+ 
22+static constexpr int64_t IDX_0 = 0;
23+ 
24+static ge::graphStatus InferDataTypeSqrt(gert::InferDataTypeContext* context)
25+{
26+ OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeSqrt");
Z
Zzhangzijie2025年10月28日

建议在使用context指针前判断其是否为nullptr

likedislike
L
liuxiqiang
2025年10月28日 评论：
27+ 
28+ // 设置输出的dtype
29+ ge::DataType sizeDtype = context->GetInputDataType(IDX_0);
30+ context->SetOutputDataType(IDX_0, sizeDtype);
31+ 
32+ OP_LOGD(context->GetNodeName(), "End to do InferDataTypeSqrt");
33+ return GRAPH_SUCCESS;
34+}
35+ 
36+IMPL_OP(Sqrt).InferDataType(InferDataTypeSqrt);
37+ 
38+};
Aexperimental/math/sqrt/op_graph/sqrt_proto.h+42-0
@@ -0,0 +1,42 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This file is a part of the CANN Open Software.
5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file sqrt_proto.h
14+ * \brief
15+*/
16+ 
17+#ifndef OP_PROTO_SQRT_PROTO_H_
18+#define OP_PROTO_SQRT_PROTO_H_
19+ 
20+#include "graph/operator_reg.h"
21+#include "graph/operator.h"
22+namespace ge {
23+/**
24+* @brief Computes the sqrt of a tensor.
25+ 
26+*@par Inputs:
27+* @li x: A tensor of type float16, float32, bf16.
28+*@par Outputs:
29+* @li y: A tensor of type float16, float32, bf16.
30+ 
31+*@par Third-party framework compatibility
32+* Compatible with the Pytorch operator Sqrt.
33+*/
34+ 
35+REG_OP(Sqrt)
36+ .INPUT(x, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
37+ .OUTPUT(y, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
38+ .OP_END_FACTORY_REG(Sqrt);
39+ 
40+} // namespace ge
41+ 
42+#endif // OP_PROTO_SQRT_PROTO_H_
Aexperimental/math/sqrt/op_host/CMakeLists.txt+10-0
@@ -0,0 +1,10 @@
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,
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+add_modules_sources(OPTYPE sqrt ACLNNTYPE aclnn)
Aexperimental/math/sqrt/op_host/config/ascend310b/sqrt_binary.json+93-0
@@ -0,0 +1,93 @@
1+{
2+ "op_type": "Sqrt",
3+ "op_list": [
4+ {
5+ "bin_filename": "Sqrt_a1532827238e1555db7b997c7bce2927",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "float32",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ]
16+ }
17+ ],
18+ "outputs": [
19+ {
20+ "name": "y",
21+ "index": 0,
22+ "dtype": "float32",
23+ "format": "ND",
24+ "paramType": "required",
25+ "shape": [
26+ -2
27+ ]
28+ }
29+ ],
30+ "attrs": [
31+ ]
32+ },
33+
34+ {
35+ "bin_filename": "Sqrt_a1532827238e1555db7b997c7bce2926",
36+ "inputs": [
37+ {
38+ "name": "x",
39+ "index": 0,
40+ "dtype": "float16",
41+ "format": "ND",
42+ "paramType": "required",
43+ "shape": [
44+ -2
45+ ]
46+ }
47+ ],
48+ "outputs": [
49+ {
50+ "name": "y",
51+ "index": 0,
52+ "dtype": "float16",
53+ "format": "ND",
54+ "paramType": "required",
55+ "shape": [
56+ -2
57+ ]
58+ }
59+ ],
60+ "attrs": [
61+ ]
62+ },
63+ {
64+ "bin_filename": "Sqrt_a1532827238e1555db7b997c7bce2925",
65+ "inputs": [
66+ {
67+ "name": "x",
68+ "index": 0,
69+ "dtype": "bfloat16",
70+ "format": "ND",
71+ "paramType": "required",
72+ "shape": [
73+ -2
74+ ]
75+ }
76+ ],
77+ "outputs": [
78+ {
79+ "name": "y",
80+ "index": 0,
81+ "dtype": "bfloat16",
82+ "format": "ND",
83+ "paramType": "required",
84+ "shape": [
85+ -2
86+ ]
87+ }
88+ ],
89+ "attrs": [
90+ ]
91+ }
92+ ]
93+}
Aexperimental/math/sqrt/op_host/config/ascend310b/sqrt_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascendxx为可选mode，如果不同芯片有差异化要求时，需要配置；
6+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置：各个版本按default配置
8+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子，由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+[Sqrt]
13+default=0
Aexperimental/math/sqrt/op_host/config/ascend910b/sqrt_binary.json+93-0
@@ -0,0 +1,93 @@
1+{
2+ "op_type": "Sqrt",
3+ "op_list": [
4+ {
5+ "bin_filename": "Sqrt_a1532827238e1555db7b997c7bce2928",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "float32",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ]
16+ }
17+ ],
18+ "outputs": [
19+ {
20+ "name": "y",
21+ "index": 0,
22+ "dtype": "float32",
23+ "format": "ND",
24+ "paramType": "required",
25+ "shape": [
26+ -2
27+ ]
28+ }
29+ ],
30+ "attrs": [
31+ ]
32+ },
33+
34+ {
35+ "bin_filename": "Sqrt_a1532827238e1555db7b997c7bce2929",
36+ "inputs": [
37+ {
38+ "name": "x",
39+ "index": 0,
40+ "dtype": "float16",
41+ "format": "ND",
42+ "paramType": "required",
43+ "shape": [
44+ -2
45+ ]
46+ }
47+ ],
48+ "outputs": [
49+ {
50+ "name": "y",
51+ "index": 0,
52+ "dtype": "float16",
53+ "format": "ND",
54+ "paramType": "required",
55+ "shape": [
56+ -2
57+ ]
58+ }
59+ ],
60+ "attrs": [
61+ ]
62+ },
63+ {
64+ "bin_filename": "Sqrt_a1532827238e1555db7b997c7bce2930",
65+ "inputs": [
66+ {
67+ "name": "x",
68+ "index": 0,
69+ "dtype": "bfloat16",
70+ "format": "ND",
71+ "paramType": "required",
72+ "shape": [
73+ -2
74+ ]
75+ }
76+ ],
77+ "outputs": [
78+ {
79+ "name": "y",
80+ "index": 0,
81+ "dtype": "bfloat16",
82+ "format": "ND",
83+ "paramType": "required",
84+ "shape": [
85+ -2
86+ ]
87+ }
88+ ],
89+ "attrs": [
90+ ]
91+ }
92+ ]
93+}
Aexperimental/math/sqrt/op_host/config/ascend910b/sqrt_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascendxx为可选mode，如果不同芯片有差异化要求时，需要配置；
6+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置：各个版本按default配置
8+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子，由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+[Sqrt]
13+default=0
Aexperimental/math/sqrt/op_host/sqrt_def.cpp+47-0
@@ -0,0 +1,47 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This file is a part of the CANN Open Software.
5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file sqrt.cpp
14+ * \brief
15+*/
16+#include "register/op_def_registry.h"
17+ 
18+namespace ops {
19+class Sqrt : public OpDef {
20+public:
21+ explicit Sqrt(const char* name) : OpDef(name)
22+ {
23+ this->Input("x")
24+ .ParamType(REQUIRED)
25+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
26+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
27+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
28+ this->Output("y")
29+ .ParamType(REQUIRED)
30+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
31+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
32+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
C
Cchenxingyu182025年10月27日

请确实是否支撑图模式，支持的话需要配置动态shape

likedislike
L
liuxiqiang
2025年10月29日 评论：
33+ 
34+ OpAICoreConfig aicoreConfig;
35+ aicoreConfig.DynamicCompileStaticFlag(true)
36+ .DynamicFormatFlag(false)
37+ .DynamicRankSupportFlag(true)
38+ .DynamicShapeSupportFlag(true)
39+ .NeedCheckSupportFlag(false)
40+ .PrecisionReduceFlag(true)
41+ .ExtendCfgInfo("opFile.value", "sqrt_apt");
42+ this->AICore().AddConfig("ascend910b")
S
Ssunday2025年10月28日

因为这个算子已经有对应的TBE实现文件了，当前默认生成的ascendc的adapter文件名为sqrt.py会覆盖以前的tbe实现文件，因此这个需要让adapter换一个名字，建议为sqrt_apt 参考：https://gitcode.com/cann/ops-math/blob/master/examples/add_example/op_host/add_example_def.cpp 增加一个.ExtendCfgInfo("opFile.value", "sqrt_apt");的配置。 详细说明：https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/83RC1alpha003/API/ascendcopapi/atlasascendc_api_07_0997.html

likedislike
L
liuxiqiang
2025年10月29日 评论：
43+ .AddConfig("ascend310b");
44+ }
45+};
46+OP_ADD(Sqrt); // 添加算子信息库
47+} // namespace ops
Aexperimental/math/sqrt/op_host/sqrt_infershape.cpp+44-0
@@ -0,0 +1,44 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This file is a part of the CANN Open Software.
5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file sqrt_infershape.cpp
14+ * \brief
15+*/
16+#include "register/op_impl_registry.h"
17+#include "log/log.h"
18+ 
19+using namespace ge;
20+ 
21+namespace ops {
22+static constexpr int64_t IDX_0 = 0;
23+ 
24+static ge::graphStatus InferShapeSqrt(gert::InferShapeContext* context)
25+{
26+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
27+ OP_LOGD(context->GetNodeName(), "Begin to do InferShapeSqrt");
Z
Zzhangzijie2025年10月28日

建议在使用context指针前检查是否为nullptr

likedislike
L
liuxiqiang
2025年10月28日 评论：
28+ 
29+ // get input shapes
30+ const gert::Shape* xShape = context->GetInputShape(IDX_0);
31+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
32+ 
33+ // get output shapes
34+ gert::Shape* yShape = context->GetOutputShape(IDX_0);
35+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
36+ 
37+ // 填充输出shape大小
38+ *yShape = *xShape;
39+ OP_LOGD(context->GetNodeName(), "End to do InferShapeSqrt");
40+ return GRAPH_SUCCESS;
41+}
42+ 
43+IMPL_OP_INFERSHAPE(Sqrt).InferShape(InferShapeSqrt);
44+}
Aexperimental/math/sqrt/op_host/sqrt_tiling.cpp+173-0
@@ -0,0 +1,173 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This file is a part of the CANN Open Software.
5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file sqrt_tiling.cpp
14+ * \brief
15+ */
16+#include "log/log.h"
17+#include "util/math_util.h"
18+#include "tiling_base/tiling_util.h"
19+#include "tiling/platform/platform_ascendc.h"
20+#include "register/op_impl_registry.h"
21+#include "tiling_base/tiling_templates_registry.h"
22+#include "../op_kernel/sqrt_tiling_data.h"
23+#include "../op_kernel/sqrt_tiling_key.h"
24+ 
25+namespace optiling {
26+ 
27+using namespace Ops::Math::OpTiling;
28+ 
29+#define BLOCK_SIZE 32U
30+#define UB_DATA_NUM_FLOAT 4U // 对应DT_FLOAT类型的ub分块数量
31+#define UB_DATA_NUM_OTHER 6U // 对应其他数据类型的ub分块数量
32+constexpr uint32_t BUFFER_NUM = 2;
33+constexpr uint32_t WS_SYS_SIZE = 0;
34+struct SqrtCompileInfo {};
35+ 
36+static ge::graphStatus TilingParseForSqrt([[maybe_unused]] gert::TilingParseContext* context)
37+{
38+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
39+ return ge::GRAPH_SUCCESS;
40+}
41+ 
42+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
43+{
44+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
45+ // 获取ubsize coreNum
46+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
47+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
48+ coreNum = ascendcPlatform.GetCoreNum();
49+ OP_CHECK_IF(coreNum <= 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
50+ OP_CHECK_IF(ubSize <= 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
51+ return ge::GRAPH_SUCCESS;
52+}
53+ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
54+{
55+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
56+ size_t usrSize = 0;
57+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
58+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
59+ size_t* currentWorkspace = context->GetWorkspaceSizes(
60+ 1); // 通过框架获取workspace的指针，GetWorkspaceSizes入参为所需workspace的块数。当前限制使用一块。
61+ currentWorkspace[0] = usrSize + sysWorkspaceSize;
62+ return ge::GRAPH_SUCCESS;
63+}
64+ 
65+ge::graphStatus GetShapeAttrsInfo(
66+ gert::TilingContext* context, uint64_t ubSize, uint64_t& inputNum, uint64_t& inputBytes, uint64_t& tileBlockNum,
67+ uint64_t& tileDataNum, uint64_t& inputLengthAlgin32)
68+{
69+ OP_CHECK_IF(
70+ context == nullptr || context->GetInputShape(0) == nullptr, OP_LOGE(context, "context is nullptr"),
71+ return ge::GRAPH_FAILED);
72+ inputNum = context->GetInputShape(0)->GetStorageShape().GetShapeSize();
73+ uint32_t typeLength = 0;
74+ ge::TypeUtils::GetDataTypeLength(context->GetInputDesc(0)->GetDataType(), typeLength);
75+ uint64_t inputLength = inputNum * typeLength;
76+ if (inputNum == 0) {
77+ return ge::GRAPH_FAILED;
78+ }
79+ inputBytes = inputLength / inputNum;
80+ uint64_t ubDataNumber =
81+ (context->GetInputDesc(0)->GetDataType() == ge::DT_FLOAT) ? UB_DATA_NUM_FLOAT : UB_DATA_NUM_OTHER;
82+ tileBlockNum = (ubSize / BLOCK_SIZE) / ubDataNumber;
83+ if (inputBytes == 0) {
84+ return ge::GRAPH_FAILED;
85+ }
86+ tileDataNum = (tileBlockNum * BLOCK_SIZE) / inputBytes;
87+ inputLengthAlgin32 = (((inputLength + BLOCK_SIZE - 1) / BLOCK_SIZE) * BLOCK_SIZE);
88+ return ge::GRAPH_SUCCESS;
89+}
90+ 
91+ge::graphStatus CalculateCoreBlockNums(
92+ uint64_t inputLengthAlgin32, int64_t coreNum, uint64_t tileBlockNum, uint64_t inputBytes, uint64_t tileDataNum,
93+ uint64_t& smallCoreDataNum, uint64_t& bigCoreDataNum, uint64_t& smallTailDataNum, uint64_t& bigTailDataNum,
94+ uint64_t& finalSmallTileNum, uint64_t& finalBigTileNum, uint64_t& tailBlockNum)
95+{
96+ if (0 == BLOCK_SIZE || 0 == coreNum || 0 == tileBlockNum || 0 == inputBytes) {
97+ return ge::GRAPH_FAILED;
98+ }
99+ uint64_t everyCoreInputBlockNum = inputLengthAlgin32 / BLOCK_SIZE / coreNum;
100+ tailBlockNum = (inputLengthAlgin32 / BLOCK_SIZE) % coreNum;
101+ smallCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
102+ uint64_t smallTileNum = everyCoreInputBlockNum / tileBlockNum;
103+ finalSmallTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? smallTileNum : smallTileNum + 1;
104+ smallTailDataNum = smallCoreDataNum - (tileDataNum * smallTileNum);
105+ smallTailDataNum = smallTailDataNum == 0 ? tileDataNum : smallTailDataNum;
106+ 
107+ everyCoreInputBlockNum += 1;
108+ bigCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
109+ uint64_t bigTileNum = everyCoreInputBlockNum / tileBlockNum;
110+ finalBigTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? bigTileNum : bigTileNum + 1;
111+ bigTailDataNum = bigCoreDataNum - tileDataNum * bigTileNum;
112+ bigTailDataNum = bigTailDataNum == 0 ? tileDataNum : bigTailDataNum;
113+ 
114+ return ge::GRAPH_SUCCESS;
115+}
116+ 
117+static ge::graphStatus SqrtTilingFunc(gert::TilingContext* context)
118+{
119+ // SqrtTilingData tiling;
120+ SqrtTilingData* tiling = context->GetTilingData<SqrtTilingData>();
121+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
122+ OP_CHECK_IF(
123+ memset_s(tiling, sizeof(SqrtTilingData), 0, sizeof(SqrtTilingData)) != EOK,
124+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
125+ // 获取平台运行信息
126+ uint64_t ubSize;
127+ int64_t coreNum;
128+ ge::graphStatus ret = GetPlatformInfo(context, ubSize, coreNum);
129+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
130+ // 获取输入数据信息
131+ uint64_t inputNum, inputBytes, tileBlockNum, tileDataNum, inputLengthAlgin32;
132+ ret = GetShapeAttrsInfo(context, ubSize, inputNum, inputBytes, tileBlockNum, tileDataNum, inputLengthAlgin32);
133+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
134+ 
135+ // 计算coreNum
136+ if (tileDataNum >= inputNum) {
137+ coreNum = 1;
138+ } else {
139+ // There is at least 32B of data on each core, satisfying several settings for several cores. The maximum number
140+ // of audits is the actual number of audits
141+ coreNum = (static_cast<uint64_t>(coreNum) < inputLengthAlgin32 / BLOCK_SIZE) ? coreNum :
142+ inputLengthAlgin32 / BLOCK_SIZE;
143+ }
144+ // 计算每个core处理的数据块数
145+ uint64_t smallCoreDataNum, bigCoreDataNum, smallTailDataNum, bigTailDataNum;
146+ uint64_t finalSmallTileNum, finalBigTileNum, tailBlockNum;
147+ ret = CalculateCoreBlockNums(
148+ inputLengthAlgin32, coreNum, tileBlockNum, inputBytes, tileDataNum, smallCoreDataNum, bigCoreDataNum,
149+ smallTailDataNum, bigTailDataNum, finalSmallTileNum, finalBigTileNum, tailBlockNum);
150+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "CalculateCoreBlockNums error"), return ge::GRAPH_FAILED);
151+ // 设置tiling数据
152+ tiling->smallCoreDataNum = static_cast<uint64_t>(smallCoreDataNum);
153+ tiling->bigCoreDataNum = static_cast<uint64_t>(bigCoreDataNum);
154+ tiling->tileDataNum = static_cast<uint64_t>(tileDataNum);
155+ tiling->smallTailDataNum = static_cast<uint64_t>(smallTailDataNum);
156+ tiling->bigTailDataNum = static_cast<uint64_t>(bigTailDataNum);
157+ tiling->finalSmallTileNum = static_cast<uint64_t>(finalSmallTileNum);
158+ tiling->finalBigTileNum = static_cast<uint64_t>(finalBigTileNum);
159+ tiling->tailBlockNum = static_cast<uint64_t>(tailBlockNum);
R
RRuiWang_2025年11月3日

不需要加static_cast，另考虑是否uint32_t类型已足够存放，可适当减小tilingdata大小

likedislike
L
liuxiqiang
2025年11月3日 评论：
160+ // 计算workspace大小
161+ OP_CHECK_IF(
162+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
163+ return ge::GRAPH_FAILED);
164+ uint64_t tilingKey = 0;
165+ tilingKey = GET_TPL_TILING_KEY(0);
166+ context->SetTilingKey(tilingKey);
167+ context->SetBlockDim(coreNum);
168+ return ge::GRAPH_SUCCESS;
169+}
170+ 
171+// tiling注册入口.
172+IMPL_OP_OPTILING(Sqrt).Tiling(SqrtTilingFunc).TilingParse<SqrtCompileInfo>(TilingParseForSqrt);
173+} // namespace optiling
Aexperimental/math/sqrt/op_kernel/sqrt.cpp+32-0
@@ -0,0 +1,32 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This file is a part of the CANN Open Software.
5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file sqrt.cpp
14+ * \brief
15+*/
16+ 
17+#include "sqrt.h"
18+ 
19+enum class SqrtTilingKey : uint32_t
20+{
21+ TILING_KEY_EXAMPLE_FLOAT = 0,
22+ TILING_KEY_EXAMPLE_OTHER = 1,
23+};
24+template <uint32_t schMode>
25+__global__ __aicore__ void sqrt(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
26+{
27+ REGISTER_TILING_DEFAULT(SqrtTilingData);
28+ GET_TILING_DATA_WITH_STRUCT(SqrtTilingData, tilingData, tiling);
29+ MySqrt::KernelSqrt<DTYPE_X,DTYPE_Y> op;
30+ op.Init(x, y,tilingData.smallCoreDataNum,tilingData.bigCoreDataNum, tilingData.finalBigTileNum,tilingData.finalSmallTileNum, tilingData.tileDataNum,tilingData.smallTailDataNum, tilingData.bigTailDataNum,tilingData.tailBlockNum); // 算子kernel实例初始化
31+ op.Process();
32+}
Aexperimental/math/sqrt/op_kernel/sqrt.h+141-0
@@ -0,0 +1,141 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This file is a part of the CANN Open Software.
5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file sqrt.h
14+ * \brief
15+ */
16+#ifndef SQRT_H
17+#define SQRT_H
18+ 
19+#include "kernel_operator.h"
20+#include "kernel_tiling/kernel_tiling.h"
21+#include "sqrt_tiling_data.h"
22+#include "sqrt_tiling_key.h"
23+ 
24+namespace MySqrt {
25+ 
26+using namespace AscendC;
27+ 
28+constexpr int32_t BUFFER_NUM = 2;
29+ 
30+template <typename TYPE_X, typename TYPE_Y>
31+class KernelSqrt {
32+public:
33+ __aicore__ inline KernelSqrt(){};
34+ 
35+ __aicore__ inline void Init(
36+ GM_ADDR x, GM_ADDR y, uint64_t smallCoreDataNum, uint64_t bigCoreDataNum, uint64_t finalBigTileNum,
37+ uint64_t finalSmallTileNum, uint64_t tileDataNum, uint64_t smallTailDataNum, uint64_t bigTailDataNum,
38+ uint64_t tailBlockNuma);
39+ __aicore__ inline void Process();
40+ 
41+private:
42+ __aicore__ inline void CopyIn(int32_t progress);
43+ __aicore__ inline void CopyOut(int32_t progress);
44+ __aicore__ inline void Compute(int32_t progress);
45+ 
46+private:
47+ AscendC::TPipe pipe;
48+ AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX;
49+ AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueY;
50+ AscendC::TBuf<AscendC::QuePosition::VECCALC> tmp1;
51+ AscendC::GlobalTensor<TYPE_X> xGm;
52+ AscendC::GlobalTensor<TYPE_Y> yGm;
53+ uint64_t coreDataNum;
54+ uint64_t tileNum;
55+ uint64_t tileDataNum;
56+ uint64_t tailDataNum;
57+ uint64_t processDataNum;
58+};
59+ 
60+template <typename TYPE_X, typename TYPE_Y>
61+__aicore__ inline void KernelSqrt<TYPE_X, TYPE_Y>::Init(
62+ GM_ADDR x, GM_ADDR y, uint64_t smallCoreDataNum, uint64_t bigCoreDataNum, uint64_t finalBigTileNum,
63+ uint64_t finalSmallTileNum, uint64_t tileDataNum, uint64_t smallTailDataNum, uint64_t bigTailDataNum,
64+ uint64_t tailBlockNum)
65+{
66+ ASSERT(AscendC::GetBlockNum() != 0 && "block dim can not be zero!");
67+ uint64_t coreId = AscendC::GetBlockIdx();
68+ uint64_t globalBufferIndex = bigCoreDataNum * coreId;
69+ this->tileDataNum = tileDataNum;
70+ if (coreId < tailBlockNum) {
71+ this->coreDataNum = bigCoreDataNum;
72+ this->tileNum = finalBigTileNum;
73+ this->tailDataNum = bigTailDataNum;
74+ } else {
75+ this->coreDataNum = smallCoreDataNum;
76+ this->tileNum = finalSmallTileNum;
77+ this->tailDataNum = smallTailDataNum;
78+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (coreId - tailBlockNum);
79+ }
80+ xGm.SetGlobalBuffer((__gm__ TYPE_X*)x + globalBufferIndex, this->coreDataNum);
81+ yGm.SetGlobalBuffer((__gm__ TYPE_Y*)y + globalBufferIndex, this->coreDataNum);
82+ pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
83+ pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
84+ if constexpr (!std::is_same_v<TYPE_X, float32_t>) {
85+ pipe.InitBuffer(tmp1, this->tileDataNum * sizeof(float));
86+ }
87+}
88+ 
89+template <typename TYPE_X, typename TYPE_Y>
90+__aicore__ inline void KernelSqrt<TYPE_X, TYPE_Y>::CopyIn(int32_t progress)
91+{
92+ AscendC::LocalTensor<TYPE_X> xLocal = inQueueX.AllocTensor<TYPE_X>();
93+ AscendC::DataCopy(xLocal, xGm[progress * this->tileDataNum], this->processDataNum);
94+ inQueueX.EnQue(xLocal);
95+}
96+ 
97+template <typename TYPE_X, typename TYPE_Y>
98+__aicore__ inline void KernelSqrt<TYPE_X, TYPE_Y>::CopyOut(int32_t progress)
99+{
100+ AscendC::LocalTensor<TYPE_Y> yLocal = outQueueY.DeQue<TYPE_Y>();
101+ AscendC::DataCopy(yGm[progress * this->tileDataNum], yLocal, this->processDataNum);
102+ outQueueY.FreeTensor(yLocal);
103+}
104+ 
105+template <typename TYPE_X, typename TYPE_Y>
106+__aicore__ inline void KernelSqrt<TYPE_X, TYPE_Y>::Compute(int32_t progress)
107+{
108+ AscendC::LocalTensor<TYPE_X> xLocal = inQueueX.DeQue<TYPE_X>();
109+ AscendC::LocalTensor<TYPE_Y> yLocal = outQueueY.AllocTensor<TYPE_Y>();
110+ if constexpr (!std::is_same_v<TYPE_X, float32_t>) {
R
RRuiWang_2025年11月3日

建议使用AscendC接口IsSameType

likedislike
L
liuxiqiang
2025年11月3日 评论：
111+ AscendC::LocalTensor<float> p1 = tmp1.Get<float>();
112+ AscendC::Cast(p1, xLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
113+ AscendC::PipeBarrier<PIPE_V>();
114+ AscendC::Sqrt(p1, p1, this->processDataNum);
115+ AscendC::PipeBarrier<PIPE_V>();
116+ AscendC::Cast(yLocal, p1, AscendC::RoundMode::CAST_CEIL, this->processDataNum);
117+ } else {
118+ AscendC::Sqrt(yLocal, xLocal, this->processDataNum);
119+ }
120+ outQueueY.EnQue<TYPE_Y>(yLocal);
121+ inQueueX.FreeTensor(xLocal);
122+}
123+ 
124+template <typename TYPE_X, typename TYPE_Y>
125+__aicore__ inline void KernelSqrt<TYPE_X, TYPE_Y>::Process()
126+{
127+ int32_t loopCount = this->tileNum;
128+ this->processDataNum = this->tileDataNum;
129+ for (int32_t i = 0; i < loopCount - 1; i++) {
130+ CopyIn(i);
131+ Compute(i);
132+ CopyOut(i);
133+ }
134+ this->processDataNum = this->tailDataNum;
135+ CopyIn(loopCount - 1);
136+ Compute(loopCount - 1);
137+ CopyOut(loopCount - 1);
138+}
139+ 
140+} // namespace MySqrt
141+#endif // SQRT_H
Aexperimental/math/sqrt/op_kernel/sqrt_tiling_data.h+28-0
@@ -0,0 +1,28 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This file is a part of the CANN Open Software.
5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file sqrt_tiling_data.h
14+ * \brief tiling data struct
15+*/
R
RRuiWang_2025年11月3日

缺少ifndef，宏名尽量包含到算子类别目录，防止冲突

likedislike
L
liuxiqiang
2025年11月3日 评论：
16+#ifndef SQRT_TILING_DATA_H_
17+#define SQRT_TILING_DATA_H_
18+struct SqrtTilingData{
19+ uint64_t smallCoreDataNum;
20+ uint64_t bigCoreDataNum;
21+ uint64_t finalBigTileNum;
22+ uint64_t finalSmallTileNum;
23+ uint64_t tileDataNum;
24+ uint64_t smallTailDataNum;
25+ uint64_t bigTailDataNum;
26+ uint64_t tailBlockNum;
27+} ;
28+#endif
Aexperimental/math/sqrt/op_kernel/sqrt_tiling_key.h+29-0
@@ -0,0 +1,29 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This file is a part of the CANN Open Software.
5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file sqrt_tiling_key.h
14+ * \brief sqrt tiling key declare
15+*/
16+#include "ascendc/host_api/tiling/template_argument.h"
17+ 
18+#define ELEMENTWISE_TPL_SCH_MODE_0 0
19+#define ELEMENTWISE_TPL_SCH_MODE_1 1
20+ 
21+ASCENDC_TPL_ARGS_DECL(
22+ Sqrt,
23+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1));
24+ 
25+ASCENDC_TPL_SEL(
26+ ASCENDC_TPL_ARGS_SEL(
27+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST,ELEMENTWISE_TPL_SCH_MODE_0,ELEMENTWISE_TPL_SCH_MODE_1)
28+ ),
29+);
Aexperimental/math/sqrt/tests/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify it.
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
9+# the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+foreach(SUB_DIR ${CURRENT_DIRS})
14+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+endforeach()
Aexperimental/math/sqrt/tests/ut/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify it.
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
9+# the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+foreach(SUB_DIR ${CURRENT_DIRS})
14+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+endforeach()
Aexperimental/math/sqrt/tests/ut/op_host/CMakeLists.txt+22-0
@@ -0,0 +1,22 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify it.
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
9+# the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+if(UT_TEST_ALL OR OP_HOST_UT)
13+ add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+ add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
15+endif()
16+ 
17+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
18+foreach(SUB_DIR ${CURRENT_DIRS})
19+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
20+ add_subdirectory(${SUB_DIR})
21+ endif()
22+endforeach()
Aexperimental/math/sqrt/tests/ut/op_host/sqrt_tiling.h+30-0
@@ -0,0 +1,30 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This file is a part of the CANN Open Software.
5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file sqrt_tiling_def.h
14+ * \brief
15+ */
16+ 
17+#ifndef SQRT_TILING_DEF_H
18+#define SQRT_TILING_DEF_H
19+ 
20+#include "register/tilingdata_base.h"
21+ 
22+namespace optiling {
23+struct SqrtCompileInfo {
24+ int32_t totalCoreNum = 0;
25+ int64_t ubSize = 0;
26+ bool isRegbase = false;
27+};
28+} // namespace optiling
29+ 
30+#endif // SQRT_TILING_DEF_H
Aexperimental/math/sqrt/tests/ut/op_host/test_sqrt_infershape.cpp+60-0
@@ -0,0 +1,60 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This file is a part of the CANN Open Software.
5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+#include <gtest/gtest.h>
13+#include <iostream>
14+#include "infershape_context_faker.h"
15+#include "infershape_case_executor.h"
16+ 
17+class SqrtInfershape : public testing::Test {
18+protected:
19+ static void SetUpTestCase()
20+ {
21+ std::cout << "SqrtInfershape SetUp" << std::endl;
22+ }
23+ 
24+ static void TearDownTestCase()
25+ {
26+ std::cout << "SqrtInfershape TearDown" << std::endl;
27+ }
28+};
29+ 
30+TEST_F(SqrtInfershape, sqrt_infershape_test1)
31+{
32+ gert::InfershapeContextPara infershapeContextPara(
33+ "Sqrt",
34+ {
35+ {{{3, 4}, {3, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
36+ },
37+ {
38+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
39+ });
40+ std::vector<std::vector<int64_t>> expectOutputShape = {
41+ {3, 4},
42+ };
43+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
44+}
45+ 
46+TEST_F(SqrtInfershape, sqrt_infershape_test2)
47+{
48+ gert::InfershapeContextPara infershapeContextPara(
49+ "Sqrt",
50+ {
51+ {{{5, -1}, {5, -1}}, ge::DT_FLOAT, ge::FORMAT_ND},
52+ },
53+ {
54+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
55+ });
56+ std::vector<std::vector<int64_t>> expectOutputShape = {
57+ {5, -1},
58+ };
59+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
60+}
Aexperimental/math/sqrt/tests/ut/op_host/test_sqrt_tiling.cpp+88-0
@@ -0,0 +1,88 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This file is a part of the CANN Open Software.
5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+#include <iostream>
13+#include <gtest/gtest.h>
14+#include "sqrt_tiling.h"
15+#include "../../../op_kernel/sqrt_tiling_data.h"
16+#include "../../../op_kernel/sqrt_tiling_key.h"
17+#include "tiling_context_faker.h"
18+#include "tiling_case_executor.h"
19+ 
20+using namespace std;
21+using namespace optiling;
22+ 
23+class SqrtTiling : public testing::Test {
24+protected:
25+ static void SetUpTestCase()
26+ {
27+ cout << "SqrtTiling SetUp" << endl;
28+ }
29+ 
30+ static void TearDownTestCase()
31+ {
32+ cout << "SqrtTiling TearDown " << endl;
33+ }
34+};
35+ 
36+TEST_F(SqrtTiling, ascend9101_test_tiling_fp16_001)
37+{
38+ optiling::SqrtCompileInfo compileInfo = {64, 262144, true};
39+ gert::TilingContextPara tilingContextPara(
40+ "Sqrt",
41+ {
42+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND},
43+ },
44+ {
45+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND},
46+ },
47+ &compileInfo);
48+ uint64_t expectTilingKey = 0;
49+ string expectTilingData = "8192 8208 1 1 21840 8192 8208 0 ";
50+ std::vector<size_t> expectWorkspaces = {16777216};
51+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
52+}
53+ 
54+TEST_F(SqrtTiling, ascend9101_test_tiling_bf16_002)
55+{
56+ optiling::SqrtCompileInfo compileInfo = {64, 262144, true};
57+ gert::TilingContextPara tilingContextPara(
58+ "Sqrt",
59+ {
60+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_BF16, ge::FORMAT_ND},
61+ },
62+ {
63+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_BF16, ge::FORMAT_ND},
64+ },
65+ &compileInfo);
66+ uint64_t expectTilingKey = 0;
67+ string expectTilingData = "8192 8208 1 1 21840 8192 8208 0 ";
68+ std::vector<size_t> expectWorkspaces = {16777216};
69+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
70+}
71+ 
72+TEST_F(SqrtTiling, ascend9101_test_tiling_fp32_003)
73+{
74+ optiling::SqrtCompileInfo compileInfo = {64, 262144, true};
75+ gert::TilingContextPara tilingContextPara(
76+ "Sqrt",
77+ {
78+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT, ge::FORMAT_ND},
79+ },
80+ {
81+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT, ge::FORMAT_ND},
82+ },
83+ &compileInfo);
84+ uint64_t expectTilingKey = 0;
85+ string expectTilingData = "8192 8200 1 1 16384 8192 8200 0 ";
86+ std::vector<size_t> expectWorkspaces = {16777216};
87+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
88+}
Aexperimental/math/sqrt/tests/ut/op_kernel/CMakeLists.txt+29-0
@@ -0,0 +1,29 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify it.
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
9+# the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+if (UT_TEST_ALL OR OP_KERNEL_UT)
13+ # 需要将Tiling依赖的文件添加到CMakeLists.txt中
14+ # set(elewise_common_tiling_files
15+ # ${CANN_ROOT}/ops/built-in/op_tiling/runtime/elewise_tiling.cc
16+ # )
17+ # 算子自己的tiling文件路径
18+ set(sqrt_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/sqrt_tiling.cpp
20+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/sqrt_infershape.cpp
21+ # ${elewise_common_tiling_files}
22+ )
23+ # 使用AddOpTestCase
24+ # param1：算子名称，以kernel方式命名
25+ # param2：soc版本，多个以分号分隔，例如："ascend910_9599;AscendB1"
26+ # param3：自定义编译选项，一般填写测试的一种典型数据类型组合，不需要则传入空字符串，例如："-DDTYPE_X=float"，多个使用空格分隔，例如："-DDTYPE_X=float -DDTYPE_Y=float"
27+ # param4：该算子依赖的所有tiling源码文件
28+ # # AddOpTestCase(is_finite "ascend910_9599" "-DDTYPE_X=float" "${is_finite_tiling_files}")
29+endif()
Aexperimental/math/sqrt/tests/ut/op_kernel/sqrt_data/compare_data.py+57-0
@@ -0,0 +1,57 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# ----------------------------------------------------------------------------
4+# This program is free software, you can redistribute it and/or modify it.
5+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6+# This file is a part of the CANN Open Software.
7+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
8+# Please refer to the License for details. You may not use this file except in compliance with the License.
9+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
10+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
11+# the software repository for the full text of the License.
12+# ----------------------------------------------------------------------------
13+ 
14+import sys
15+import numpy as np
16+import glob
17+import os
18+ 
19+curr_dir = os.path.dirname(os.path.realpath(__file__))
20+ 
21+def compare_data(golden_file_lists, output_file_lists, d_type):
22+ if d_type == "float16":
23+ np_dtype = np.float16
24+ elif d_type == "float32":
25+ np_dtype = np.float32
26+ else:
27+ raise ValueError("d_type must be float16 or float32")
28+
29+ data_same = True
30+ for gold, out in zip(golden_file_lists, output_file_lists):
31+ tmp_out = np.fromfile(out, np_dtype)
32+ tmp_gold = np.fromfile(gold, np_dtype)
33+ diff_res = np.isclose(tmp_out, tmp_gold, 0, 0, True)
34+ diff_idx = np.where(diff_res != True)[0]
35+ if len(diff_idx) == 0:
36+ print("PASSED!")
37+ else:
38+ print("FAILED!")
39+ for idx in diff_idx[:5]:
40+ print(f"index: {idx}, output: {tmp_out[idx]}, golden: {tmp_gold[idx]}")
41+ data_same = False
42+ return data_same
43+ 
44+def get_file_lists(dtype):
45+ golden_file_lists = sorted(glob.glob(curr_dir + "/*golden*.bin"))
46+ output_file_lists = sorted(glob.glob(curr_dir + "/*output*.bin"))
47+ return golden_file_lists, output_file_lists
48+ 
49+def process(d_type):
50+ golden_file_lists, output_file_lists = get_file_lists(d_type)
51+ result = compare_data(golden_file_lists, output_file_lists, d_type)
52+ print("compare result:", result)
53+ return result
54+ 
55+if __name__ == '__main__':
56+ ret = process(sys.argv[1])
57+ exit(0 if ret else 1)
Aexperimental/math/sqrt/tests/ut/op_kernel/sqrt_data/gen_data.py+53-0
@@ -0,0 +1,53 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# ----------------------------------------------------------------------------
4+# This program is free software, you can redistribute it and/or modify it.
5+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6+# This file is a part of the CANN Open Software.
7+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
8+# Please refer to the License for details. You may not use this file except in compliance with the License.
9+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
10+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
11+# the software repository for the full text of the License.
12+# ----------------------------------------------------------------------------
13+ 
14+import sys
15+import os
16+import numpy as np
17+import re
18+import tensorflow as tf
19+ 
20+ 
21+def parse_str_to_shape_list(shape_str):
22+ shape_str = shape_str.strip('(').strip(')')
23+ shape_list = [int(x) for x in shape_str.split(",")]
24+ return np.array(shape_list)
25+ 
26+ 
27+def gen_data_and_golden(shape_str, d_type="float32"):
28+ d_type_dict = {
29+ "float32": np.float32,
30+ "float16": np.float16,
31+ "bfloat16": tf.bfloat16.as_numpy_dtype
32+ }
33+ np_type = d_type_dict[d_type]
34+ shape = parse_str_to_shape_list(shape_str)
35+ size = np.prod(shape)
36+ tmp_input = np.random.choice([0, 0.5, 1, 65504, np.nan, np.inf], size=size)
37+ tmp_input = tmp_input.reshape(shape).astype(np_type)
38+ tmp_golden = np.sqrt(tmp_input)
39+ 
40+ tmp_input.astype(np_type).tofile(f"{d_type}_input_t_sqrt.bin")
41+ tmp_golden.astype(np_type).tofile(f"{d_type}_golden_t_sqrt.bin")
42+ 
43+ 
44+if __name__ == "__main__":
45+ if len(sys.argv) != 3:
46+ print("Param num must be 3.")
47+ exit(1)
48+ # 清理bin文件
49+ os.system("rm -rf *.bin")
50+ gen_data_and_golden(sys.argv[1], sys.argv[2])
51+ 
52+ 
53+ 
Aexperimental/math/sqrt/tests/ut/op_kernel/test_sqrt.cpp+146-0
@@ -0,0 +1,146 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This file is a part of the CANN Open Software.
5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file test_sqrt.cpp
14+ * \brief
15+ */
16+ 
17+#include <array>
18+#include <vector>
19+#include <iostream>
20+#include <string>
21+#include <cstdint>
22+#include "gtest/gtest.h"
23+#include "tikicpulib.h"
24+#include "data_utils.h"
25+#include "tiling_case_executor.h"
26+#include "../op_host/sqrt_tiling.h"
27+ 
28+using namespace std;
29+ 
30+extern "C" __global__ __aicore__ void sqrt(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
31+ 
32+class SqrtTest : public testing::Test {
33+protected:
34+ static void SetUpTestCase()
35+ {
36+ std::cout << "sqrt_test SetUp" << std::endl;
37+ const string cmd = "cp -rf " + dataPath + " ./";
38+ system(cmd.c_str());
39+ system("chmod -R 755 ./sqrt_data/");
40+ }
41+ static void TearDownTestCase()
42+ {
43+ std::cout << "sqrt_test TearDown" << std::endl;
44+ }
45+ 
46+private:
47+ const static std::string rootPath;
48+ const static std::string dataPath;
49+};
50+ 
51+const std::string SqrtTest::rootPath = "../../../../";
52+const std::string SqrtTest::dataPath = rootPath + "math/sqrt/tests/ut/op_kernel/sqrt_data";
53+ 
54+template <typename T1, typename T2>
55+inline T1 CeilAlign(T1 a, T2 b)
56+{
57+ return (a + b - 1) / b * b;
58+}
59+ 
60+TEST_F(SqrtTest, test_case_float16_1)
61+{
62+ optiling::SqrtCompileInfo compileInfo = {64, 262144, false};
63+ gert::TilingContextPara tilingContextPara(
64+ "Sqrt",
65+ {
66+ {{{128, 64}, {128, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND},
67+ },
68+ {
69+ {{{128, 64}, {128, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND},
70+ },
71+ &compileInfo);
72+ TilingInfo tilingInfo;
73+ auto tilingRet = ExecuteTiling(tilingContextPara, tilingInfo);
74+ EXPECT_EQ(tilingRet, true);
75+ 
76+ system("cd ./sqrt_data/ && python3 gen_data.py '(128, 64)' 'float16'");
77+ uint32_t dataCount = 128 * 64;
78+ size_t inputByteSize = dataCount * sizeof(half);
79+ std::string fileName = "./sqrt_data/float16_input_t_sqrt.bin";
80+ ;
81+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));
82+ ReadFile(fileName, inputByteSize, x, inputByteSize);
83+ size_t outputByteSize = dataCount * sizeof(half);
84+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(CeilAlign(outputByteSize, 32));
85+ 
86+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(tilingInfo.workspaceSizes[0]);
87+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingInfo.tilingDataSize);
88+ std::memcpy(tiling, tilingInfo.tilingData.get(), tilingInfo.tilingDataSize);
89+ ICPU_SET_TILING_KEY(tilingInfo.tilingKey);
90+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
91+ ICPU_RUN_KF(sqrt, tilingInfo.blockNum, x, y, workspace, tiling);
92+ 
93+ fileName = "./sqrt_data/float16_output_t_sqrt.bin";
94+ WriteFile(fileName, y, outputByteSize);
95+ 
96+ AscendC::GmFree((void*)(x));
97+ AscendC::GmFree((void*)(y));
98+ AscendC::GmFree((void*)workspace);
99+ AscendC::GmFree((void*)tiling);
100+ 
101+ system("cd ./sqrt_data/ && python3 compare_data.py 'float16'");
102+}
103+ 
104+TEST_F(SqrtTest, test_case_float32_1)
105+{
106+ optiling::SqrtCompileInfo compileInfo = {64, 262144, false};
107+ gert::TilingContextPara tilingContextPara(
108+ "Sqrt",
109+ {
110+ {{{256, 33}, {256, 33}}, ge::DT_FLOAT, ge::FORMAT_ND},
111+ },
112+ {
113+ {{{256, 33}, {256, 33}}, ge::DT_FLOAT, ge::FORMAT_ND},
114+ },
115+ &compileInfo);
116+ TilingInfo tilingInfo;
117+ auto tilingRet = ExecuteTiling(tilingContextPara, tilingInfo);
118+ EXPECT_EQ(tilingRet, true);
119+ 
120+ system("cd ./sqrt_data/ && python3 gen_data.py '(256, 33)' 'float32'");
121+ uint32_t dataCount = 256 * 33;
122+ size_t inputByteSize = dataCount * sizeof(float);
123+ std::string fileName = "./sqrt_data/float32_input_t_sqrt.bin";
124+ ;
125+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));
126+ ReadFile(fileName, inputByteSize, x, inputByteSize);
127+ size_t outputByteSize = dataCount * sizeof(float);
128+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(CeilAlign(outputByteSize, 32));
129+ 
130+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(tilingInfo.workspaceSizes[0]);
131+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingInfo.tilingDataSize);
132+ std::memcpy(tiling, tilingInfo.tilingData.get(), tilingInfo.tilingDataSize);
133+ ICPU_SET_TILING_KEY(tilingInfo.tilingKey);
134+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
135+ ICPU_RUN_KF(sqrt, tilingInfo.blockNum, x, y, workspace, tiling);
136+ 
137+ fileName = "./sqrt_data/float32_output_t_sqrt.bin";
138+ WriteFile(fileName, y, outputByteSize);
139+ 
140+ AscendC::GmFree((void*)(x));
141+ AscendC::GmFree((void*)(y));
142+ AscendC::GmFree((void*)workspace);
143+ AscendC::GmFree((void*)tiling);
144+ 
145+ system("cd ./sqrt_data/ && python3 compare_data.py 'float32'");
146+}
Mscripts/kernel/binary_config/ascendc_config.json+2-1
@@ -33,5 +33,6 @@
33 {"name":"Segsum", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false},33 {"name":"Segsum", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false},
34 {"name":"Sinkhorn", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : true},34 {"name":"Sinkhorn", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : true},
35 {"name":"STFT", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false},35 {"name":"STFT", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false},
36- {"name":"TransformBiasRescaleQkv", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false}36+ {"name":"TransformBiasRescaleQkv", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false},
37+ {"name":"Sqrt", "compute_units": ["ascend910b", "ascend310b"], "auto_sync" : true, "impl_mode" : "high_performance"}
37]38]