Star1303
1424
代码
代码
Issues113
Pull Requests238
Actions
项目讨论
Wiki
分析
Star1303
1424
已合并
RsqrtGrad算子贡献 #1910
松柏创建于 3月26日
RsqrtGrad算子贡献 #1910
已合并
松柏创建于 3月26日
从ilovescrapy/rsqrtgradmaster合入到cann/ops-mathmaster
共 17 个文件变更+1242-0
Aexperimental/math/rsqrt_grad/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+add_all_modules_sources(OPTYPE rsqrt_grad ACLNNTYPE aclnn)
Aexperimental/math/rsqrt_grad/README.md+73-0
@@ -0,0 +1,73 @@
1+# RsqrtGrad
2+ 
3+## 产品支持情况
4+ 
5+ 
6+| 产品 | 是否支持 |
7+| :------------------------------------------------------- | :------: |
8+| <term>Ascend 950PR/Ascend 950DT</term> | × |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10+| <term>Atlas 200I/500 A2 推理产品</term> | × |
11+| <term>Atlas 推理系列产品</term> | × |
12+| <term>Atlas 训练系列产品</term> | × |
13+ 
14+## 功能说明
15+ 
16+- 算子功能：完成梯度平方根计算。
17+- 计算公式：
18+ 
19+$$
20+z = y * y * y * dy * (-0.5)
21+$$
22+ 
23+## 参数说明
24+ 
25+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
26+ <col style="width: 100px">
27+ <col style="width: 150px">
28+ <col style="width: 280px">
29+ <col style="width: 330px">
30+ <col style="width: 120px">
31+ </colgroup>
32+ <thead>
33+ <tr>
34+ <th>参数名</th>
35+ <th>输入/输出/属性</th>
36+ <th>描述</th>
37+ <th>数据类型</th>
38+ <th>数据格式</th>
39+ </tr></thead>
40+ <tbody>
41+ <tr>
42+ <td>y</td>
43+ <td>输入</td>
44+ <td>待进行rsqrt_grad计算的入参，公式中的y。</td>
45+ <td>DT_FLOAT,DT_FLOAT_16,DT_BF16</td>
46+ <td>ND</td>
47+ </tr>
48+ <tr>
49+ <td>dy</td>
50+ <td>输入</td>
51+ <td>待进行rsqrt_grad计算的入参，公式中的dy。</td>
52+ <td>DT_FLOAT,DT_FLOAT_16,DT_BF16</td>
53+ <td>ND</td>
54+ </tr>
55+ <tr>
56+ <td>z</td>
57+ <td>输出</td>
58+ <td>待进行rsqrt_grad计算的出参，公式中的z。</td>
59+ <td>DT_FLOAT,DT_FLOAT_16,DT_BF16</td>
60+ <td>ND</td>
61+ </tr>
62+ </tbody></table>
63+ 
64+## 约束说明
65+ 
66+无
67+ 
68+## 调用说明
69+ 
70+ 
71+| 调用方式 | 调用样例 | 说明 |
72+| ---------- | --------------------------------------------------------- | ------------------------------------------------------------------------------- |
73+| aclnn调用 | [test_rsqrt_grad](./examples/test_aclnn_rsqrt_grad.cpp) | 使用自动生成的aclnn接口调用RsqrtGrad算子。 |
Aexperimental/math/rsqrt_grad/examples/test_aclnn_rsqrt_grad.cpp+144-0
@@ -0,0 +1,144 @@
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 <vector>
13+#include "acl/acl.h"
14+ 
15+#define CHECK_RET(cond, return_expr) \
16+ do { \
17+ if (!(cond)) { \
18+ return_expr; \
19+ } \
20+ } while (0)
21+ 
22+#define LOG_PRINT(message, ...) \
F
Ffulltower3月28日

函数之间加个空行

likedislike
23+ do { \
24+ printf(message, ##__VA_ARGS__); \
25+ } while (0)
26+ 
27+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
28+ int64_t shapeSize = 1;
29+ for (auto i : shape) {
30+ shapeSize *= i;
31+ }
32+ return shapeSize;
33+}
34+ 
35+int Init(int32_t deviceId, aclrtStream* stream) {
36+ // 固定写法，资源初始化
37+ auto ret = aclInit(nullptr);
38+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
39+ ret = aclrtSetDevice(deviceId);
40+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
41+ ret = aclrtCreateStream(stream);
42+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
43+ return 0;
44+}
45+ 
46+template <typename T>
47+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
48+ aclDataType dataType, aclTensor** tensor) {
49+ auto size = GetShapeSize(shape) * sizeof(T);
50+ // 调用aclrtMalloc申请device侧内存
51+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
52+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
53+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
54+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
55+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
56+ // 计算连续tensor的strides
57+ std::vector<int64_t> strides(shape.size(), 1);
58+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
59+ strides[i] = shape[i + 1] * strides[i + 1];
60+ }
61+ // 调用aclCreateTensor接口创建aclTensor
62+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
63+ shape.data(), shape.size(), *deviceAddr);
64+ return 0;
65+}
66+ 
67+int main() {
68+ // 1. （固定写法）device/stream初始化，参考acl API手册
69+ // 根据自己的实际device填写deviceId
70+ int32_t deviceId = 0;
71+ aclrtStream stream;
72+ auto ret = Init(deviceId, &stream);
73+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
74+ // 2. 构造输入与输出，需要根据API的接口自定义构造
75+ std::vector<int64_t> selfShape = {2, 2};
76+ std::vector<int64_t> otherShape = {2, 2};
77+ std::vector<int64_t> outShape = {2, 2};
78+ 
79+ void* selfDeviceAddr = nullptr;
80+ void* otherDeviceAddr = nullptr;
81+ void* outDeviceAddr = nullptr;
82+ aclTensor* self = nullptr;
83+ aclTensor* other = nullptr;
84+ aclTensor* out = nullptr;
85+ std::vector<float> selfHostData = {0, 1, 2, 3};
86+ std::vector<float> otherHostData = {0, 1, 2, 3};
87+ std::vector<float> outHostData = {0, 0, 0, 0};
88+ // 创建self aclTensor
89+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
90+ CHECK_RET(ret == ACL_SUCCESS, return ret);
91+ 
92+ // 创建other aclTensor
93+ ret = CreateAclTensor(otherHostData, otherShape, &otherDeviceAddr, aclDataType::ACL_FLOAT, &other);
94+ CHECK_RET(ret == ACL_SUCCESS, return ret);
95+ 
96+ // 创建out aclTensor
97+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
98+ CHECK_RET(ret == ACL_SUCCESS, return ret);
99+ // 3. 调用CANN算子库API，需要修改为具体的API名称
100+ // aclnnRsqrtGrad接口调用示例
101+ uint64_t workspaceSize = 0;
102+ aclOpExecutor* executor;
103+ // 调用aclnnRsqrtGrad第一段接口
104+ ret = aclnnRsqrtGradGetWorkspaceSize(self, other, out, &workspaceSize, &executor);
105+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnRsqrtGradGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
106+ // 根据第一段接口计算出的workspaceSize申请device内存
107+ void* workspaceAddr = nullptr;
108+ if (workspaceSize > 0) {
109+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
110+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
111+ }
112+ // 调用aclnnRsqrtGrad第二段接口
113+ ret = aclnnRsqrtGrad(workspaceAddr, workspaceSize, executor, stream);
114+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnRsqrtGrad failed. ERROR: %d\n", ret); return ret);
115+ // 4. （固定写法）同步等待任务执行结束
116+ ret = aclrtSynchronizeStream(stream);
117+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
118+ 
119+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
120+ auto size = GetShapeSize(outShape);
121+ std::vector<float> resultData(size, 0);
122+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
123+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
124+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
125+ for (int64_t i = 0; i < size; i++) {
126+ LOG_PRINT("aclnnRsqrtGrad result[%ld] is: %f\n", i, resultData[i]);
127+ }
128+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
129+ aclDestroyTensor(self);
130+ aclDestroyTensor(other);
131+ aclDestroyTensor(out);
132+ 
133+ // 7. 释放device资源，需要根据具体API的接口定义修改
134+ aclrtFree(selfDeviceAddr);
135+ aclrtFree(otherDeviceAddr);
136+ aclrtFree(outDeviceAddr);
137+ if (workspaceSize > 0) {
138+ aclrtFree(workspaceAddr);
139+ }
140+ aclrtDestroyStream(stream);
141+ aclrtResetDevice(deviceId);
142+ aclFinalize();
143+ return 0;
144+}
Aexperimental/math/rsqrt_grad/op_host/rsqrt_grad_def.cpp+41-0
@@ -0,0 +1,41 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+/*!
12+ * \file rsqrt_grad_def.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class RsqrtGrad : public OpDef {
19+public:
20+ explicit RsqrtGrad(const char* name) : OpDef(name)
21+ {
22+ this->Input("y")
23+ .ParamType(REQUIRED)
24+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT, ge::DT_INT8, ge::DT_INT32})
25+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
26+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
27+ this->Input("dy")
28+ .ParamType(REQUIRED)
29+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT, ge::DT_INT8, ge::DT_INT32})
30+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
31+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
32+ this->Output("z")
33+ .ParamType(REQUIRED)
34+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT, ge::DT_INT8, ge::DT_INT32})
35+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
36+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
37+ this->AICore().AddConfig("ascend910b"); // 其他的soc版本补充部分配置项
38+ }
39+};
40+OP_ADD(RsqrtGrad); // 添加算子信息库
41+} // namespace ops
Aexperimental/math/rsqrt_grad/op_host/rsqrt_grad_infershape.cpp+43-0
@@ -0,0 +1,43 @@
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 rsqrt_grad_infershape.cpp
13+ * \brief
14+*/
15+#include "register/op_impl_registry.h"
16+#include "log/log.h"
17+ 
18+using namespace ge;
19+ 
20+namespace ops {
21+static constexpr int64_t IDX_0 = 0;
22+ 
23+static ge::graphStatus InferShapeRsqrtGrad(gert::InferShapeContext* context)
24+{
25+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
26+ OP_LOGD(context->GetNodeName(), "Begin to do InferShapeRsqrtGrad");
27+ 
28+ // get input shapes
29+ const gert::Shape* xShape = context->GetInputShape(IDX_0);
30+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
31+ 
32+ // get output shapes
33+ gert::Shape* yShape = context->GetOutputShape(IDX_0);
34+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
35+ 
36+ // 填充输出shape大小
37+ *yShape = *xShape;
38+ OP_LOGD(context->GetNodeName(), "End to do InferShapeRsqrtGrad");
39+ return GRAPH_SUCCESS;
40+}
41+ 
42+IMPL_OP_INFERSHAPE(RsqrtGrad).InferShape(InferShapeRsqrtGrad);
43+}
Aexperimental/math/rsqrt_grad/op_host/rsqrt_grad_tiling.cpp+228-0
@@ -0,0 +1,228 @@
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 rsqrt_grad_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include "log/log.h"
17+#include "util/math_util.h"
18+#include "util/platform_util.h"
19+#include "op_host/tiling_util.h"
20+#include "op_host/tiling_templates_registry.h"
21+#include "tiling/platform/platform_ascendc.h"
22+#include "register/op_impl_registry.h"
23+#include "../op_kernel/rsqrt_grad_tiling_data.h"
24+#include "../op_kernel/rsqrt_grad_tiling_key.h"
25+ 
26+namespace optiling {
27+ 
28+using namespace Ops::Math::OpTiling;
29+//SINGLE_BUFFER
30+#define UB_NUM_INT32_ONE 1U
31+#define UB_NUM_FLOAT_ONE 3U
32+#define UB_NUM_INT8_ONE 22U
33+#define UB_NUM_FLOAT16_BF16_ONE 7U
34+//DOUBLE_BUFFER
35+#define UB_NUM_INT32_TWO 2U
36+#define UB_NUM_FLOAT_TWO 6U
37+#define UB_NUM_INT8_TWO 24U
38+#define UB_NUM_FLOAT16_BF16_TWO 10U
39+constexpr uint32_t BUFFER_NUM = 2;
40+constexpr uint32_t WS_SYS_SIZE = 0;
41+ 
42+struct RsqrtGradCompileInfo {};
43+ 
44+static ge::graphStatus TilingParseForRsqrtGrad([[maybe_unused]] gert::TilingParseContext* context)
45+{
46+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
47+ return ge::GRAPH_SUCCESS;
48+}
49+ 
50+// 获取平台信息如ubSize, coreNum
51+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum, uint64_t& BLOCK_SIZE)
52+{
53+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
54+ // 获取ubsize coreNum
55+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
56+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
57+ coreNum = ascendcPlatform.GetCoreNumAiv();
58+ BLOCK_SIZE = Ops::Base::GetUbBlockSize(context);
59+ OP_CHECK_IF(BLOCK_SIZE == 0, OP_LOGE(context, "BLOCK_SIZE is 0"), return ge::GRAPH_FAILED);
60+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
61+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
62+ return ge::GRAPH_SUCCESS;
63+}
64+ 
65+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
66+{
67+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
68+ size_t usrSize = 0;
69+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
70+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
71+ size_t* currentWorkspace = context->GetWorkspaceSizes(
72+ 1); // 通过框架获取workspace的指针，GetWorkspaceSizes入参为所需workspace的块数。当前限制使用一块。
73+ currentWorkspace[0] = usrSize + sysWorkspaceSize;
74+ return ge::GRAPH_SUCCESS;
75+}
76+ 
77+static ge::graphStatus GetShapeAttrsInfo(
78+ gert::TilingContext* context, uint64_t BLOCK_SIZE, uint64_t ubSize, int64_t coreNum, uint64_t& bufferOpen, uint64_t& inputNum, uint64_t& inputBytes, uint64_t& tileBlockNum,
79+ uint64_t& tileDataNum, uint64_t& inputLengthAlgin)
80+{
81+ OP_CHECK_IF(
82+ context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
83+ OP_CHECK_IF(
84+ context->GetInputShape(0) == nullptr, OP_LOGE(context, "context->GetInputShape(0) is nullptr"),
85+ return ge::GRAPH_FAILED);
86+ inputNum = context->GetInputShape(0)->GetStorageShape().GetShapeSize();
87+ uint32_t typeLength = 0;
88+ ge::TypeUtils::GetDataTypeLength(context->GetInputDesc(0)->GetDataType(), typeLength);
89+ uint64_t inputLength = inputNum * typeLength;
90+ if (inputNum == 0 || BLOCK_SIZE == 0) {
91+ OP_LOGE(context, "inputNum or BLOCK_SIZE is 0");
92+ return ge::GRAPH_FAILED;
93+ }
94+ inputBytes = inputLength / inputNum;
95+ uint64_t ubDataNumber = 0U;
96+ inputLengthAlgin = (((inputLength + BLOCK_SIZE - 1) / BLOCK_SIZE) * BLOCK_SIZE);
97+ //buffer 是否开启，每个核上对齐后的数据量小于ub空间时不开启，默认开启
98+ bufferOpen = 1;
99+ if (UB_NUM_INT8_ONE == 0 || UB_NUM_FLOAT16_BF16_ONE == 0) {
100+ OP_LOGE(context, "UB_NUM_INT8_ONE or UB_NUM_FLOAT16_BF16_ONE is 0");
101+ return ge::GRAPH_FAILED;
102+ }
103+ if (UB_NUM_INT32_ONE == 0 || UB_NUM_FLOAT_ONE == 0) {
104+ OP_LOGE(context, "UB_NUM_INT32_ONE or UB_NUM_FLOAT_ONE is 0");
105+ return ge::GRAPH_FAILED;
106+ }
107+ if (context->GetInputDesc(0)->GetDataType() == ge::DT_INT8) {
108+ if (inputLengthAlgin < coreNum * (((ubSize / BLOCK_SIZE) * BLOCK_SIZE) / UB_NUM_INT8_ONE)) {
109+ bufferOpen = 0;
110+ ubDataNumber = UB_NUM_INT8_ONE;
111+ } else {
112+ ubDataNumber = UB_NUM_INT8_TWO;
113+ }
114+ } else if ((context->GetInputDesc(0)->GetDataType() == ge::DT_BF16) || (context->GetInputDesc(0)->GetDataType() == ge::DT_FLOAT16)) {
115+ if (inputLengthAlgin < coreNum * (((ubSize / BLOCK_SIZE) * BLOCK_SIZE) / UB_NUM_FLOAT16_BF16_ONE)) {
116+ bufferOpen = 0;
117+ ubDataNumber = UB_NUM_FLOAT16_BF16_ONE;
118+ } else {
119+ ubDataNumber = UB_NUM_FLOAT16_BF16_TWO;
120+ }
121+ } else if (context->GetInputDesc(0)->GetDataType() == ge::DT_INT32) {
122+ if (inputLengthAlgin < coreNum * (((ubSize / BLOCK_SIZE) * BLOCK_SIZE) / UB_NUM_INT32_ONE)) {
123+ bufferOpen = 0;
124+ ubDataNumber = UB_NUM_INT32_ONE;
125+ } else {
126+ ubDataNumber = UB_NUM_INT32_TWO;
127+ }
128+ } else {
129+ if (inputLengthAlgin < coreNum * (((ubSize / BLOCK_SIZE) * BLOCK_SIZE) / UB_NUM_FLOAT_ONE)) {
130+ bufferOpen = 0;
131+ ubDataNumber = UB_NUM_FLOAT_ONE;
132+ } else {
133+ ubDataNumber = UB_NUM_FLOAT_TWO;
134+ }
135+ }
136+ if (ubDataNumber == 0 || inputBytes == 0) {
137+ OP_LOGE(context, "ubDataNumber or inputBytes is 0");
138+ return ge::GRAPH_FAILED;
139+ }
140+ tileBlockNum = (ubSize / BLOCK_SIZE) / ubDataNumber;
141+ tileDataNum = (tileBlockNum * BLOCK_SIZE) / inputBytes;
142+ return ge::GRAPH_SUCCESS;
143+}
144+ 
145+static ge::graphStatus CalculateCoreBlockNums(
146+ gert::TilingContext* context, uint64_t BLOCK_SIZE, uint64_t inputLengthAlgin, int64_t coreNum, uint64_t tileBlockNum, uint64_t inputBytes,
147+ uint64_t tileDataNum, uint64_t& smallCoreDataNum, uint64_t& bigCoreDataNum, uint64_t& smallTailDataNum, uint64_t& bigTailDataNum,
148+ uint64_t& finalSmallTileNum, uint64_t& finalBigTileNum, uint64_t& tailBlockNum)
149+{
150+ if (0 == BLOCK_SIZE || 0 == coreNum || 0 == tileBlockNum || 0 == inputBytes) {
151+ OP_LOGE(context, "BLOCK_SIZE or coreNum or tileBlockNum or inputBytes is 0");
152+ return ge::GRAPH_FAILED;
153+ }
154+ uint64_t everyCoreInputBlockNum = inputLengthAlgin / BLOCK_SIZE / coreNum;
155+ tailBlockNum = (inputLengthAlgin / BLOCK_SIZE) % coreNum;
156+ smallCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
157+ uint64_t smallTileNum = everyCoreInputBlockNum / tileBlockNum;
158+ finalSmallTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? smallTileNum : smallTileNum + 1;
159+ smallTailDataNum = smallCoreDataNum - (tileDataNum * smallTileNum);
160+ smallTailDataNum = smallTailDataNum == 0 ? tileDataNum : smallTailDataNum;
161+ 
162+ everyCoreInputBlockNum += 1;
163+ bigCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
164+ uint64_t bigTileNum = everyCoreInputBlockNum / tileBlockNum;
165+ finalBigTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? bigTileNum : bigTileNum + 1;
166+ bigTailDataNum = bigCoreDataNum - tileDataNum * bigTileNum;
167+ bigTailDataNum = bigTailDataNum == 0 ? tileDataNum : bigTailDataNum;
168+ 
169+ return ge::GRAPH_SUCCESS;
170+}
171+ 
172+// tiling 分发入口
173+static ge::graphStatus RsqrtGradTilingFunc(gert::TilingContext* context)
174+{
175+ // 1、获取平台运行信息
176+ uint64_t ubSize, BLOCK_SIZE;
177+ int64_t coreNum;
178+ uint64_t bufferOpen;
179+ ge::graphStatus ret = GetPlatformInfo(context, ubSize, coreNum, BLOCK_SIZE);
180+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
181+ // 2、获取shape、属性信息
182+ uint64_t inputNum, inputBytes, tileBlockNum, tileDataNum, inputLengthAlgin;
183+ ret = GetShapeAttrsInfo(context, BLOCK_SIZE, ubSize, coreNum, bufferOpen, inputNum, inputBytes, tileBlockNum, tileDataNum, inputLengthAlgin);
184+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
185+ // 3、获取WorkspaceSize信息
186+ OP_CHECK_IF(
187+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
188+ return ge::GRAPH_FAILED);
189+ // 4、设置tiling信息
190+ RsqrtGradTilingData* tiling = context->GetTilingData<RsqrtGradTilingData>();
191+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
192+ OP_CHECK_IF(
193+ memset_s(tiling, sizeof(RsqrtGradTilingData), 0, sizeof(RsqrtGradTilingData)) != EOK,
194+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
195+ 
196+ if (tileDataNum >= inputNum) {
197+ coreNum = 1;
198+ } else {
199+ coreNum = (static_cast<uint64_t>(coreNum) < inputLengthAlgin / BLOCK_SIZE) ? coreNum : inputLengthAlgin / BLOCK_SIZE;
200+ }
201+ // 计算每个core处理的数据块数
202+ uint64_t smallCoreDataNum, bigCoreDataNum, smallTailDataNum, bigTailDataNum, finalSmallTileNum, finalBigTileNum, tailBlockNum;
203+ ret = CalculateCoreBlockNums(
204+ context, BLOCK_SIZE, inputLengthAlgin, coreNum, tileBlockNum, inputBytes, tileDataNum, smallCoreDataNum, bigCoreDataNum,
205+ smallTailDataNum, bigTailDataNum, finalSmallTileNum, finalBigTileNum, tailBlockNum);
206+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "CalculateCoreBlockNums error"), return ge::GRAPH_FAILED);
207+ // 设置tiling数据
208+ tiling->smallCoreDataNum = static_cast<uint64_t>(smallCoreDataNum);
209+ tiling->bigCoreDataNum = static_cast<uint64_t>(bigCoreDataNum);
210+ tiling->tileDataNum = static_cast<uint64_t>(tileDataNum);
211+ tiling->smallTailDataNum = static_cast<uint64_t>(smallTailDataNum);
212+ tiling->bigTailDataNum = static_cast<uint64_t>(bigTailDataNum);
213+ tiling->finalSmallTileNum = static_cast<uint64_t>(finalSmallTileNum);
214+ tiling->finalBigTileNum = static_cast<uint64_t>(finalBigTileNum);
215+ tiling->tailBlockNum = static_cast<uint64_t>(tailBlockNum);
216+ 
217+ tiling->bufferOpen = static_cast<uint64_t>(bufferOpen);
218+ 
219+ context->SetBlockDim(coreNum);
220+ uint64_t tilingKey = 0;
221+ tilingKey = GET_TPL_TILING_KEY(ELEMENTWISE_TPL_SCH_MODE_0);
222+ context->SetTilingKey(tilingKey);
223+ return ge::GRAPH_SUCCESS;
224+}
225+ 
226+// tiling注册入口.
227+IMPL_OP_OPTILING(RsqrtGrad).Tiling(RsqrtGradTilingFunc).TilingParse<RsqrtGradCompileInfo>(TilingParseForRsqrtGrad);
228+} // namespace optiling
Aexperimental/math/rsqrt_grad/op_kernel/rsqrt_grad.cpp+27-0
@@ -0,0 +1,27 @@
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 rsqrt_grad.cpp
13+ * \brief
14+ */
15+ 
16+#include "rsqrt_grad.h"
17+ 
18+template <uint32_t schMode>
19+__global__ __aicore__ void rsqrt_grad(GM_ADDR y, GM_ADDR dy, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
20+{
21+ REGISTER_TILING_DEFAULT(RsqrtGradTilingData);
22+ GET_TILING_DATA_WITH_STRUCT(RsqrtGradTilingData, tilingData, tiling);
23+ NsRsqrtGrad::KernelRsqrtGrad<DTYPE_Y> op; // 算子kernel实例获取
24+ op.Init(y, dy, z, tilingData.smallCoreDataNum, tilingData.bigCoreDataNum, tilingData.finalBigTileNum, tilingData.finalSmallTileNum, tilingData.tileDataNum,
25+ tilingData.smallTailDataNum, tilingData.bigTailDataNum, tilingData.tailBlockNum, tilingData.bufferOpen);
26+ op.Process();
27+}
Aexperimental/math/rsqrt_grad/op_kernel/rsqrt_grad.h+229-0
@@ -0,0 +1,229 @@
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 rsqrt_grad.h
13+ * \brief
14+ */
15+#ifndef RSQRT_GRAD_H_
16+#define RSQRT_GRAD_H_
17+ 
18+#include <math.h>
19+ 
20+#include "kernel_operator.h"
21+#include "kernel_tiling/kernel_tiling.h"
22+#include "rsqrt_grad_tiling_data.h"
23+#include "rsqrt_grad_tiling_key.h"
24+ 
25+#include "kernel_operator.h"
26+ 
27+namespace NsRsqrtGrad {
28+ 
29+using namespace AscendC;
30+ 
31+constexpr int32_t DOUBLE_BUFFER_NUM = 2;
32+constexpr int32_t SINGLE_BUFFER_NUM = 1;
33+ 
34+template <typename TYPE_Y>
35+class KernelRsqrtGrad {
36+public:
37+ __aicore__ inline KernelRsqrtGrad(){};
38+ 
39+ __aicore__ inline void Init(GM_ADDR y, GM_ADDR dy, GM_ADDR z, uint64_t smallCoreDataNum, uint64_t bigCoreDataNum, uint64_t finalBigTileNum,
40+ uint64_t finalSmallTileNum, uint64_t tileDataNum, uint64_t smallTailDataNum, uint64_t bigTailDataNum, uint64_t tailBlockNum, uint64_t bufferOpen);
41+ __aicore__ inline void Process();
42+ 
43+private:
44+ __aicore__ inline void CopyIn(int32_t progress);
45+ __aicore__ inline void CopyOut(int32_t progress);
46+ __aicore__ inline void Compute(int32_t progress);
47+ 
48+private:
49+ AscendC::TPipe pipe;
50+ AscendC::TQue<AscendC::TPosition::VECIN, DOUBLE_BUFFER_NUM> inQueueY, inQueueDY;
51+ AscendC::TQue<AscendC::TPosition::VECOUT, DOUBLE_BUFFER_NUM> outQueueZ;
52+ AscendC::TBuf<AscendC::TPosition::VECCALC> tmpQueue0, tmpQueue1, tmpQueue2, tmpQueue3, tmpQueue4;
53+ 
54+ AscendC::GlobalTensor<TYPE_Y> yGm, dyGm, zGm;
55+ uint64_t coreDataNum = 0;
56+ uint64_t tileNum = 0;
57+ uint64_t tileDataNum = 0;
58+ uint64_t tailDataNum = 0;
59+ uint64_t processDataNum = 0;
60+};
61+ 
62+template <typename TYPE_Y>
63+__aicore__ inline void KernelRsqrtGrad<TYPE_Y>::Init(GM_ADDR y, GM_ADDR dy, GM_ADDR z, uint64_t smallCoreDataNum, uint64_t bigCoreDataNum, uint64_t finalBigTileNum,
64+ uint64_t finalSmallTileNum, uint64_t tileDataNum, uint64_t smallTailDataNum, uint64_t bigTailDataNum, uint64_t tailBlockNum, uint64_t bufferOpen)
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+ uint64_t BUFFER_NUM = DOUBLE_BUFFER_NUM;
71+ if (bufferOpen == 0) {
72+ BUFFER_NUM = SINGLE_BUFFER_NUM;
73+ }
74+ if (coreId < tailBlockNum) {
75+ this->coreDataNum = bigCoreDataNum;
76+ this->tileNum = finalBigTileNum;
77+ this->tailDataNum = bigTailDataNum;
78+ } else {
79+ this->coreDataNum = smallCoreDataNum;
80+ this->tileNum = finalSmallTileNum;
81+ this->tailDataNum = smallTailDataNum;
82+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (coreId - tailBlockNum);
83+ }
84+ yGm.SetGlobalBuffer((__gm__ TYPE_Y *)y + globalBufferIndex, this->coreDataNum);
85+ dyGm.SetGlobalBuffer((__gm__ TYPE_Y *)dy + globalBufferIndex, this->coreDataNum);
86+ zGm.SetGlobalBuffer((__gm__ TYPE_Y *)z + globalBufferIndex, this->coreDataNum);
87+ if constexpr (!std::is_same_v<TYPE_Y, int32_t>) {
88+ pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
89+ pipe.InitBuffer(inQueueDY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
90+ }
91+ pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
92+ 
93+ if constexpr (std::is_same_v<TYPE_Y, int8_t>) {
94+ pipe.InitBuffer(tmpQueue0, this->tileDataNum * sizeof(half));
95+ pipe.InitBuffer(tmpQueue1, this->tileDataNum * sizeof(float));
96+ pipe.InitBuffer(tmpQueue2, this->tileDataNum * sizeof(float));
97+ pipe.InitBuffer(tmpQueue3, this->tileDataNum * sizeof(int32_t));
98+ pipe.InitBuffer(tmpQueue4, this->tileDataNum * sizeof(int32_t));
99+ } else if constexpr (std::is_same_v<TYPE_Y, half> || std::is_same_v<TYPE_Y, bfloat16_t>) {
100+ pipe.InitBuffer(tmpQueue0, this->tileDataNum * sizeof(float));
101+ pipe.InitBuffer(tmpQueue1, this->tileDataNum * sizeof(float));
102+ }
103+}
104+ 
105+template <typename TYPE_Y>
106+__aicore__ inline void KernelRsqrtGrad<TYPE_Y>::CopyIn(int32_t progress)
107+{
108+ AscendC::LocalTensor<TYPE_Y> yLocal = inQueueY.AllocTensor<TYPE_Y>();
109+ AscendC::LocalTensor<TYPE_Y> dyLocal = inQueueDY.AllocTensor<TYPE_Y>();
110+ AscendC::DataCopy(yLocal, yGm[progress * this->tileDataNum], this->processDataNum);
111+ AscendC::DataCopy(dyLocal, dyGm[progress * this->tileDataNum], this->processDataNum);
112+ inQueueY.EnQue(yLocal);
113+ inQueueDY.EnQue(dyLocal);
114+}
115+ 
116+template <typename TYPE_Y>
117+__aicore__ inline void KernelRsqrtGrad<TYPE_Y>::CopyOut(int32_t progress)
118+{
119+ AscendC::LocalTensor<TYPE_Y> zLocal = outQueueZ.DeQue<TYPE_Y>();
120+ AscendC::DataCopy(zGm[progress * this->tileDataNum], zLocal, this->processDataNum);
121+ outQueueZ.FreeTensor(zLocal);
122+}
123+ 
124+template <typename TYPE_Y>
125+__aicore__ inline void KernelRsqrtGrad<TYPE_Y>::Compute(int32_t progress)
126+{
127+ if constexpr (std::is_same_v<TYPE_Y, float>) {
128+ AscendC::LocalTensor<float> yLocal = inQueueY.DeQue<float>();
129+ AscendC::LocalTensor<float> dyLocal = inQueueDY.DeQue<float>();
130+ AscendC::LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();
131+ AscendC::Mul(zLocal, yLocal, yLocal, this->processDataNum);
132+ AscendC::Mul(zLocal, zLocal, yLocal, this->processDataNum);
133+ AscendC::Mul(zLocal, zLocal, dyLocal, this->processDataNum);
134+ AscendC::Muls(zLocal, zLocal, static_cast<TYPE_Y>(-0.5), this->processDataNum);
135+ outQueueZ.EnQue<float>(zLocal);
136+ inQueueY.FreeTensor(yLocal);
137+ inQueueDY.FreeTensor(dyLocal);
138+ } else if constexpr (std::is_same_v<TYPE_Y, int32_t>) {
139+ AscendC::LocalTensor<int32_t> zLocal = outQueueZ.AllocTensor<int32_t>();
140+ AscendC::Duplicate(zLocal, static_cast<TYPE_Y>(-0.5), this->processDataNum);
141+ outQueueZ.EnQue<int32_t>(zLocal);
142+ } else if constexpr (std::is_same_v<TYPE_Y, half> || std::is_same_v<TYPE_Y, bfloat16_t>) {
143+ AscendC::LocalTensor<TYPE_Y> yLocal = inQueueY.DeQue<TYPE_Y>();
144+ AscendC::LocalTensor<TYPE_Y> dyLocal = inQueueDY.DeQue<TYPE_Y>();
145+ AscendC::LocalTensor<TYPE_Y> zLocal = outQueueZ.AllocTensor<TYPE_Y>();
146+ AscendC::LocalTensor<float> tmp0Local = tmpQueue0.AllocTensor<float>();
147+ AscendC::LocalTensor<float> tmp1Local = tmpQueue1.AllocTensor<float>();
148+ AscendC::Cast(tmp0Local, yLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
149+ AscendC::Mul(tmp1Local, tmp0Local, tmp0Local, this->processDataNum);
150+ AscendC::Mul(tmp1Local, tmp1Local, tmp0Local, this->processDataNum);
151+ AscendC::Cast(tmp0Local, dyLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
152+ AscendC::Mul(tmp1Local, tmp1Local, tmp0Local, this->processDataNum);
153+ AscendC::Muls(tmp1Local, tmp1Local, static_cast<float>(-0.5), this->processDataNum);
154+ AscendC::Cast(zLocal, tmp1Local, AscendC::RoundMode::CAST_RINT, this->processDataNum);
155+ outQueueZ.EnQue<TYPE_Y>(zLocal);
156+ inQueueY.FreeTensor(yLocal);
157+ inQueueDY.FreeTensor(dyLocal);
158+ } else {
159+ //int8类型处理分支
160+ AscendC::LocalTensor<int8_t> yLocal = inQueueY.DeQue<int8_t>();
161+ AscendC::LocalTensor<int8_t> dyLocal = inQueueDY.DeQue<int8_t>();
162+ AscendC::LocalTensor<int8_t> zLocal = outQueueZ.AllocTensor<int8_t>();
163+ AscendC::LocalTensor<half> tmp0Local = tmpQueue0.AllocTensor<half>();
164+ AscendC::LocalTensor<float> tmp1Local = tmpQueue1.AllocTensor<float>();
165+ AscendC::LocalTensor<float> tmp2Local = tmpQueue2.AllocTensor<float>();
166+ AscendC::LocalTensor<int32_t> tmp3Local = tmpQueue3.AllocTensor<int32_t>();
167+ AscendC::LocalTensor<int32_t> tmp4Local = tmpQueue4.AllocTensor<int32_t>();
168+ AscendC::Cast(tmp0Local, yLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
169+ AscendC::Cast(tmp1Local, tmp0Local, AscendC::RoundMode::CAST_NONE, this->processDataNum);
170+ AscendC::Mul(tmp2Local, tmp1Local, tmp1Local, this->processDataNum);
171+ AscendC::Mul(tmp2Local, tmp2Local, tmp1Local, this->processDataNum);
172+ AscendC::Cast(tmp0Local, dyLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
173+ AscendC::Cast(tmp1Local, tmp0Local, AscendC::RoundMode::CAST_NONE, this->processDataNum);
174+ AscendC::Mul(tmp2Local, tmp2Local, tmp1Local, this->processDataNum);
175+ AscendC::Muls(tmp2Local, tmp2Local, static_cast<float>(-0.5), this->processDataNum);
176+ 
177+ AscendC::Cast(tmp3Local, tmp2Local, AscendC::RoundMode::CAST_TRUNC, this->processDataNum);
178+ AscendC::Duplicate(tmp4Local, static_cast<int32_t>(255), this->processDataNum);
179+ AscendC::And(tmp3Local, tmp3Local, tmp4Local, this->processDataNum);
180+ AscendC::Cast(tmp1Local, tmp3Local, AscendC::RoundMode::CAST_NONE, this->processDataNum);
181+ //uint8_int8_overflow_proc
182+ AscendC::Adds(tmp2Local, tmp1Local, static_cast<float>(128.0), this->processDataNum);
183+ //tensormodint tmpscalar=1/256=0.00390625
184+ AscendC::Muls(tmp2Local, tmp2Local, static_cast<float>(0.00390625), this->processDataNum);
185+ AscendC::Cast(tmp4Local, tmp2Local, AscendC::RoundMode::CAST_FLOOR, this->processDataNum);
186+ AscendC::Duplicate(tmp3Local, static_cast<int32_t>(128), this->processDataNum);
187+ AscendC::Mul(tmp4Local, tmp4Local, tmp3Local, this->processDataNum);
188+ AscendC::Cast(tmp2Local, tmp4Local, AscendC::RoundMode::CAST_NONE, this->processDataNum);
189+ AscendC::Sub(tmp2Local, tmp1Local, tmp2Local, this->processDataNum);
190+ //tensormodint end
191+ AscendC::Adds(tmp2Local, tmp2Local, static_cast<float>(-128.0), this->processDataNum);
192+ AscendC::Cast(tmp0Local, tmp2Local, AscendC::RoundMode::CAST_NONE, this->processDataNum);
193+ AscendC::Cast(zLocal, tmp0Local, AscendC::RoundMode::CAST_TRUNC, this->processDataNum);
194+ //uint8_int8_overflow_proc end
195+ outQueueZ.EnQue<int8_t>(zLocal);
196+ inQueueY.FreeTensor(yLocal);
197+ inQueueDY.FreeTensor(dyLocal);
198+ }
199+}
200+ 
201+template <typename TYPE_Y>
202+__aicore__ inline void KernelRsqrtGrad<TYPE_Y>::Process()
203+{
204+ int32_t loopCount = this->tileNum;
205+ if constexpr (!std::is_same_v<TYPE_Y, int32_t>) {
206+ this->processDataNum = this->tileDataNum;
207+ for (int32_t i = 0; i < loopCount - 1; i++) {
208+ CopyIn(i);
209+ Compute(i);
210+ CopyOut(i);
211+ }
212+ this->processDataNum = this->tailDataNum;
213+ CopyIn(loopCount - 1);
214+ Compute(loopCount - 1);
215+ CopyOut(loopCount - 1);
216+ } else {
217+ this->processDataNum = this->tileDataNum;
218+ for (int32_t i = 0; i < loopCount - 1; i++) {
219+ Compute(i);
220+ CopyOut(i);
221+ }
222+ this->processDataNum = this->tailDataNum;
223+ Compute(loopCount - 1);
224+ CopyOut(loopCount - 1);
225+ }
226+}
227+ 
228+} // namespace NsRsqrtGrad
229+#endif // RSQRT_GRAD_H
Aexperimental/math/rsqrt_grad/op_kernel/rsqrt_grad_tiling_data.h+30-0
@@ -0,0 +1,30 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file rsqrt_grad_tiling_data.h
13+ * \brief tiling data struct
14+ */
15+ 
16+#ifndef RSQRT_GRAD_TILING_DATA_H_
17+#define RSQRT_GRAD_TILING_DATA_H_
18+ 
19+struct RsqrtGradTilingData {
20+ uint64_t smallCoreDataNum;
21+ uint64_t bigCoreDataNum;
22+ uint64_t finalBigTileNum;
23+ uint64_t finalSmallTileNum;
24+ uint64_t tileDataNum;
25+ uint64_t smallTailDataNum;
26+ uint64_t bigTailDataNum;
27+ uint64_t tailBlockNum;
28+ uint64_t bufferOpen;
29+};
30+#endif
Aexperimental/math/rsqrt_grad/op_kernel/rsqrt_grad_tiling_key.h+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+/*!
12+ * \file rsqrt_grad_tiling_key.h
13+ * \brief rsqrt_grad tiling key declare
14+ */
15+ 
16+#ifndef __RSQRT_GRAD_TILING_KEY_H__
17+#define __RSQRT_GRAD_TILING_KEY_H__
18+ 
19+#include "ascendc/host_api/tiling/template_argument.h"
20+ 
21+#define ELEMENTWISE_TPL_SCH_MODE_0 0
22+#define ELEMENTWISE_TPL_SCH_MODE_1 1
23+ 
24+ASCENDC_TPL_ARGS_DECL(
25+ RsqrtGrad,
26+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1));
27+ 
28+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
29+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1)));
30+ 
31+#endif
Aexperimental/math/rsqrt_grad/tests/ut/op_host/rsqrt_grad_tiling.h+29-0
@@ -0,0 +1,29 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file rsqrt_grad_tiling.h
13+ * \brief
14+ */
15+ 
16+#ifndef RSQRT_GRAD_TILING_H
17+#define RSQRT_GRAD_TILING_H
18+ 
19+#include "register/tilingdata_base.h"
20+ 
21+namespace optiling {
22+struct RsqrtGradCompileInfo {
23+ int32_t totalCoreNum = 0;
24+ int64_t ubSize = 0;
25+ bool isRegbase = false;
26+};
27+} // namespace optiling
28+ 
29+#endif // RSQRT_GRAD_TILING_H
Aexperimental/math/rsqrt_grad/tests/ut/op_host/test_rsqrt_grad_tiling.cpp+52-0
@@ -0,0 +1,52 @@
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 "rsqrt_grad_tiling.h"
14+#include "../../../op_kernel/rsqrt_grad_tiling_data.h"
15+#include "../../../op_kernel/rsqrt_grad_tiling_key.h"
16+#include "tiling_context_faker.h"
17+#include "tiling_case_executor.h"
18+ 
19+using namespace std;
20+using namespace optiling;
21+ 
22+class RsqrtGradTiling : public testing::Test {
23+protected:
24+ static void SetUpTestCase()
25+ {
26+ cout << "RsqrtGradTiling SetUp" << endl;
27+ }
28+ 
29+ static void TearDownTestCase()
30+ {
31+ cout << "RsqrtGradTiling TearDown " << endl;
32+ }
33+};
34+ 
35+TEST_F(RsqrtGradTiling, ascend910_test_tiling_FLOAT_001)
36+{
37+ optiling::RsqrtGradCompileInfo compileInfo = {40, 196608, false};
38+ gert::TilingContextPara tilingContextPara(
39+ "RsqrtGrad",
40+ {
41+ {{{1024, 1024}, {1024, 1024}}, ge::DT_FLOAT, ge::FORMAT_ND},
42+ {{{1024, 1024}, {1024, 1024}}, ge::DT_FLOAT, ge::FORMAT_ND},
43+ },
44+ {
45+ {{{1024, 1024}, {1024, 1024}}, ge::DT_FLOAT, ge::FORMAT_ND},
46+ },
47+ &compileInfo);
48+ uint64_t expectTilingKey = 0;
49+ string expectTilingData = "16384 16392 1 1 21824 16384 16392 0 0 ";
50+ std::vector<size_t> expectWorkspaces = {16777216};
51+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
52+}
Aexperimental/math/rsqrt_grad/tests/ut/op_kernel/CMakeLists.txt+29-0
@@ -0,0 +1,29 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+ 
12+if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)
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(rsqrt_grad_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/rsqrt_grad_tiling.cpp
20+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/rsqrt_grad_infershape.cpp
21+ # ${elewise_common_tiling_files}
22+ )
23+ # 使用AddOpTestCase
24+ # param1：算子名称，以kernel方式命名
25+ # param2：soc版本，多个以分号分隔，例如："ascend950pr_9599;AscendB1"
26+ # param3：自定义编译选项，一般填写测试的一种典型数据类型组合，不需要则传入空字符串，例如："-DDTYPE_y=float"，多个使用空格分隔，例如："-DDTYPE_Y=float -DDTYPE_Y=float"
27+ # param4：该算子依赖的所有tiling源码文件
28+ AddOpTestCase(rsqrt_grad "ascend910b" "-DDTYPE_Y=float_t" "${rsqrt_grad_tiling_files}")
29+endif()
Aexperimental/math/rsqrt_grad/tests/ut/op_kernel/rsqrt_grad_data/compare_data.py+57-0
@@ -0,0 +1,57 @@
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+import sys
12+import numpy as np
13+import glob
14+import os
15+ 
16+curr_dir = os.path.dirname(os.path.realpath(__file__))
17+ 
18+def compare_data(golden_file_lists, output_file_lists, d_type):
19+ if d_type == "float16":
20+ np_dtype = np.float16
21+ rtol = 1e-3
22+ atol = 1e-3
23+ elif d_type == "float":
24+ np_dtype = np.float32
25+ rtol = 1e-4
26+ atol = 1e-4
27+ else:
28+ raise ValueError("d_type must be float16 or float")
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, rtol, atol, 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/rsqrt_grad/tests/ut/op_kernel/rsqrt_grad_data/gen_data.py+46-0
@@ -0,0 +1,46 @@
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+import sys
12+import os
13+import numpy as np
14+import re
15+ 
16+def parse_str_to_shape_list(shape_str):
17+ shape_str = shape_str.strip('(').strip(')')
18+ shape_list = [int(x) for x in shape_str.split(",")]
19+ return np.array(shape_list)
20+ 
21+ 
22+def gen_data_and_golden(shape_str, d_type="float"):
23+ d_type_dict = {
24+ "float16": np.float16,
25+ "float": np.float32
26+ }
27+ np_type = d_type_dict[d_type]
28+ shape = parse_str_to_shape_list(shape_str)
29+ size = np.prod(shape)
30+ y = np.random.uniform(-10, 0, shape).astype(np_type)
31+ dy = np.random.uniform(-10, 0, shape).astype(np_type)
32+ 
33+ golden = (y * y * y * dy * (-0.5)).astype(np_type)
34+ 
35+ y.astype(np_type).tofile(f"{d_type}_input_y_rsqrt_grad.bin")
36+ dy.astype(np_type).tofile(f"{d_type}_input_dy_rsqrt_grad.bin")
37+ golden.astype(np_type).tofile(f"{d_type}_golden_rsqrt_grad.bin")
38+ 
39+ 
40+if __name__ == "__main__":
41+ if len(sys.argv) != 3:
42+ print("Param num must be 3.")
43+ exit(1)
44+ # 清理bin文件
45+ os.system("rm -rf *.bin")
46+ gen_data_and_golden(sys.argv[1], sys.argv[2])
Aexperimental/math/rsqrt_grad/tests/ut/op_kernel/rsqrt_grad_tiling.h+59-0
@@ -0,0 +1,59 @@
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 rsqrt_grad_tiling.h
13+ * \brief
14+ */
15+ 
16+#ifndef _RSQRT_GRAD_TILING_H_
17+#define _RSQRT_GRAD_TILING_H_
18+ 
19+#include <cstdint>
20+ 
21+#include "../../../op_kernel/rsqrt_grad_tiling_data.h"
22+#include "kernel_tiling/kernel_tiling.h"
23+#include "../../../op_kernel/rsqrt_grad.h"
24+ 
25+using namespace NsRsqrtGrad;
26+ 
27+#define __aicore__
28+#ifdef __NPU_TILING__
29+inline[aicore] void InitTilingData(const __gm__ uint8_t *tiling, RsqrtGradTilingData *constData)
30+{
31+ const __gm__ uint32_t *src = (const __gm__ uint32_t *)tiling;
32+ uint32_t *dst = (uint32_t *)constData;
33+ for (size_t i = 0; i < sizeof(RsqrtGradTilingData) / 4; i++) {
34+ *(dst + i) = *(src + i);
35+ }
36+}
37+#else
38+inline void InitTilingData(uint8_t *tiling, RsqrtGradTilingData *constData)
39+{
40+ memcpy(constData, tiling, sizeof(RsqrtGradTilingData));
41+}
42+#endif // __NPU_TILING__
43+ 
44+#define CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
45+ __ubuf__ tilingStruct *tilingDataPointer = \
46+ reinterpret_cast<__ubuf__ tilingStruct *>((__ubuf__ uint8_t *)(tilingPointer));
47+ 
48+#define INIT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
49+ CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer);
50+ 
51+#define GET_TILING_DATA_WITH_STRUCT(tilingStruct, tilingData, tilingArg) \
52+ tilingStruct tilingData; \
53+ InitTilingData(tilingArg, &tilingData)
54+ 
55+#define GET_TILING_DATA(tilingData, tilingArg) \
56+ RsqrtGradTilingData tilingData; \
57+ InitTilingData(tilingArg, &tilingData)
58+ 
59+#endif // _RSQRT_GRAD_TILING_H_
Aexperimental/math/rsqrt_grad/tests/ut/op_kernel/test_rsqrt_grad.cpp+113-0
@@ -0,0 +1,113 @@
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 <array>
12+#include <vector>
13+#include "gtest/gtest.h"
14+ 
15+#include <array>
16+#include <vector>
17+#include <iostream>
18+#include <string>
19+#include <cstdint>
20+#include "gtest/gtest.h"
21+#include "tikicpulib.h"
22+#include "data_utils.h"
23+ 
24+#include "../../../op_kernel/rsqrt_grad.cpp"
25+ 
26+using namespace std;
27+ 
28+constexpr uint32_t smallCoreDataNum = 1024;
29+constexpr uint32_t bigCoreDataNum = 1040;
30+constexpr uint32_t tileDataNum = 2048;
31+constexpr uint32_t smallTailDataNum = 1024;
32+constexpr uint32_t bigTailDataNum = 1040;
33+ 
34+extern "C" __global__ __aicore__ void rsqrt_grad(GM_ADDR y, GM_ADDR dy, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling);
35+ 
36+class RsqrtGradTest : public testing::Test {
37+protected:
38+ static void SetUpTestCase()
39+ {
40+ std::cout << "rsqrt_grad_test SetUp" << std::endl;
41+ const string cmd = "cp -rf " + dataPath + " ./";
42+ system(cmd.c_str());
43+ system("chmod -R 755 ./rsqrt_grad_data/");
44+ }
45+ static void TearDownTestCase()
46+ {
47+ std::cout << "rsqrt_grad_test TearDown" << std::endl;
48+ }
49+ 
50+private:
51+ const static std::string rootPath;
52+ const static std::string dataPath;
53+};
54+ 
55+const std::string RsqrtGradTest::rootPath = "../../../../experimental/";
56+const std::string RsqrtGradTest::dataPath = rootPath + "math/rsqrt_grad/tests/ut/op_kernel/rsqrt_grad_data";
57+ 
58+template <typename T1, typename T2>
59+inline T1 CeilAlign(T1 a, T2 b)
60+{
61+ return (a + b - 1) / b * b;
62+}
63+ 
64+TEST_F(RsqrtGradTest, test_case_float_1)
65+{
66+ uint32_t blockDim = 1;
67+ system("cd ./rsqrt_grad_data/ && python3 gen_data.py '(1024)' 'float'");
68+ uint32_t dataCount = 1024;
69+ size_t inputByteSize = dataCount * sizeof(float_t);
70+ 
71+ std::string y_fileName = "./rsqrt_grad_data/float_input_y_rsqrt_grad.bin";
72+ std::string dy_fileName = "./rsqrt_grad_data/float_input_dy_rsqrt_grad.bin";
73+ 
74+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));
75+ uint8_t* dy = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));
76+ 
77+ ReadFile(y_fileName, inputByteSize, y, inputByteSize);
78+ ReadFile(dy_fileName, inputByteSize, dy, inputByteSize);
79+ 
80+ size_t outputByteSize = dataCount * sizeof(float_t);
81+ uint8_t* z = (uint8_t*)AscendC::GmAlloc(CeilAlign(outputByteSize, 32));
82+ 
83+ size_t workspaceSize = 32 * 1024 * 1024;
84+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceSize);
85+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(RsqrtGradTilingData));
86+ 
87+ RsqrtGradTilingData* tilingData = reinterpret_cast<RsqrtGradTilingData*>(tiling);
88+ 
89+ tilingData->smallCoreDataNum = smallCoreDataNum;
90+ tilingData->bigCoreDataNum = bigCoreDataNum;
91+ tilingData->tileDataNum = tileDataNum;
92+ tilingData->smallTailDataNum = smallTailDataNum;
93+ tilingData->bigTailDataNum = bigTailDataNum;
94+ tilingData->finalSmallTileNum = 1;
95+ tilingData->finalBigTileNum = 1;
96+ tilingData->tailBlockNum = 0;
97+ tilingData->bufferOpen = 1;
98+ 
99+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
100+ auto func = rsqrt_grad<ELEMENTWISE_TPL_SCH_MODE_0>;
101+ ICPU_RUN_KF(func, blockDim, y, dy, z, workspace, (uint8_t*)(tilingData));
102+ 
103+ std::string fileName = "./rsqrt_grad_data/float_output_t_rsqrt_grad.bin";
104+ WriteFile(fileName, z, outputByteSize);
105+ 
106+ AscendC::GmFree((void*)(y));
107+ AscendC::GmFree((void*)(dy));
108+ AscendC::GmFree((void*)(z));
109+ AscendC::GmFree((void*)workspace);
110+ AscendC::GmFree((void*)tiling);
111+ 
112+ system("cd ./rsqrt_grad_data/ && python3 compare_data.py 'float'");
113+}