Star1290
2055
代码
代码
Issues650
Pull Requests585
Actions
项目讨论
Wiki
分析
Star1290
2055
已合并
[CANNBot] BNInferGrad算子支持Ascend950 Ascendc实现 #7960
yulianjie创建于 7月25日
[CANNBot] BNInferGrad算子支持Ascend950 Ascendc实现 #7960
已合并
yulianjie创建于 7月25日
从lianjieyu/ops-nnbn_infer_grad合入到cann/ops-nnmaster
共 25 个文件变更+2861-0
Mdocs/zh/op_list.md+10-0
@@ -3067,6 +3067,16 @@
3067 <td>AI Core</td>3067 <td>AI Core</td>
3068 <td>对一个批次的数据做批量归一化处理，正则化之后生成的数据的统计结果为0均值、1标准差。</td>3068 <td>对一个批次的数据做批量归一化处理，正则化之后生成的数据的统计结果为0均值、1标准差。</td>
3069 </tr>3069 </tr>
3070+ <tr>
3071+ <td>norm</td>
3072+ <td><a href="../../norm/bn_infer_grad/README.md">bn_infer_grad</a></td>
3073+ <td>✓</td>
3074+ <td>✓</td>
3075+ <td>✗</td>
3076+ <td>✓</td>
3077+ <td>AI Core</td>
3078+ <td>计算Batch Normalization推理模式下输入的反向梯度。</td>
3079+ </tr>
3070 <tr>3080 <tr>
3071 <td>norm</td>3081 <td>norm</td>
3072 <td><a href="../../norm/bn_training_reduce/README.md">bn_training_reduce</a></td>3082 <td><a href="../../norm/bn_training_reduce/README.md">bn_training_reduce</a></td>
Anorm/bn_infer_grad/CMakeLists.txt+24-0
@@ -0,0 +1,24 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+# Generated By CANNBot
11+ 
12+set(SUPPORT_COMPUTE_UNIT "ascend950")
13+set(SUPPORT_TILING_DIR "arch35")
14+ 
15+add_modules_sources(
16+ HOSTNAME ${OP_HOST_NAME}
17+ MODE PRIVATE
18+ DIR ${CMAKE_CURRENT_SOURCE_DIR}
19+ OPTYPE bn_infer_grad
20+ ACLNNTYPE aclnn_exclude
21+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
22+ TILING_DIR ${SUPPORT_TILING_DIR}
23+ DISABLE_IN_OPP TRUE
24+)
Anorm/bn_infer_grad/README.md+97-0
@@ -0,0 +1,97 @@
1+# BNInferGrad
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------- | :------:|
7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</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+- 算子功能：计算Batch Normalization推理模式下输入的反向梯度。
17+- 计算公式：
18+ 
19+ $$
20+ x\_backprop = grads \times \frac{scale}{\sqrt{batch\_variance + epsilon}}
21+ $$
22+ 
23+## 参数说明
24+ 
25+<table style="table-layout: fixed; width: 1576px"><colgroup>
26+<col style="width: 170px">
27+<col style="width: 170px">
28+<col style="width: 520px">
29+<col style="width: 300px">
30+<col style="width: 170px">
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>grads</td>
43+ <td>输入</td>
44+ <td>上游梯度，支持2D至5D张量。</td>
45+ <td>FLOAT16、FLOAT、BFLOAT16</td>
46+ <td>ND</td>
47+ </tr>
48+ <tr>
49+ <td>scale</td>
50+ <td>输入</td>
51+ <td>缩放参数，必须为一维张量，长度等于grads的C维。</td>
52+ <td>FLOAT</td>
53+ <td>ND</td>
54+ </tr>
55+ <tr>
56+ <td>batch_variance</td>
57+ <td>输入</td>
58+ <td>推理阶段使用的方差，必须为一维张量，长度等于grads的C维。</td>
59+ <td>FLOAT</td>
60+ <td>ND</td>
61+ </tr>
62+ <tr>
63+ <td>epsilon</td>
64+ <td>可选属性</td>
65+ <td>防止除零的小常数，默认值为0.0001。</td>
66+ <td>FLOAT</td>
67+ <td>-</td>
68+ </tr>
69+ <tr>
70+ <td>x_backprop</td>
71+ <td>输出</td>
72+ <td>输入梯度，shape和数据类型与grads一致。</td>
73+ <td>FLOAT16、FLOAT、BFLOAT16</td>
74+ <td>ND</td>
75+ </tr>
76+</tbody></table>
77+ 
78+## 约束说明
79+ 
80+- grads的第二维必须等于scale和batch_variance的长度。
81+- scale和batch_variance仅支持FLOAT。
82+ 
83+## 调用说明
84+ 
85+<table><thead>
86+ <tr>
87+ <th>调用方式</th>
88+ <th>调用样例</th>
89+ <th>说明</th>
90+ </tr></thead>
91+<tbody>
92+ <tr>
93+ <td>图模式调用</td>
94+ <td><a href="./examples/test_geir_bn_infer_grad.cpp">test_geir_bn_infer_grad</a></td>
95+ <td>通过<a href="./op_graph/bn_infer_grad_proto.h">算子IR</a>构图并调用BNInferGrad算子。</td>
96+ </tr>
97+</tbody></table>
Anorm/bn_infer_grad/examples/arch35/test_geir_bn_infer_grad.cpp+274-0
@@ -0,0 +1,274 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file test_geir_bn_infer_grad.cpp
15+ * \brief BNInferGrad GEIR graph-mode example with numerical validation.
16+ */
17+ 
18+#include <cmath>
19+#include <cstdint>
20+#include <cstdio>
21+#include <cstring>
22+#include <ctime>
23+#include <iostream>
24+#include <map>
25+#include <string>
26+#include <vector>
27+ 
28+#include "array_ops.h"
29+#include "ge_api.h"
30+#include "ge_api_types.h"
31+#include "ge_error_codes.h"
32+#include "ge_ir_build.h"
33+#include "graph.h"
34+#include "tensor.h"
35+#include "types.h"
36+ 
37+#include "../../op_graph/bn_infer_grad_proto.h"
38+ 
39+#define FAILED -1
40+#define SUCCESS 0
41+ 
42+using namespace ge;
43+using std::map;
44+using std::string;
45+using std::vector;
46+ 
47+namespace {
48+constexpr float kEpsilon = 1.0F;
49+constexpr int64_t kChannels = 3;
50+constexpr int64_t kInnerSize = 4;
51+constexpr DataType kGradsDtype = DT_BF16;
52+constexpr float kTolerance = 2.0e-2F;
53+ 
54+vector<float> gGrads = {
55+ 0.25F, 0.50F, 0.75F, 1.00F, 1.25F, 1.50F, 1.75F, 2.00F, 2.25F, 2.50F, 2.75F, 3.00F,
56+ 3.25F, 3.50F, 3.75F, 4.00F, 4.25F, 4.50F, 4.75F, 5.00F, 5.25F, 5.50F, 5.75F, 6.00F,
57+};
58+vector<float> gScale = {1.0F, 2.0F, 3.0F};
59+vector<float> gBatchVariance = {3.0F, 8.0F, 15.0F};
60+vector<uint16_t> gGradsStorage;
61+ 
62+string GetTime()
63+{
64+ time_t timeValue;
65+ time(&timeValue);
66+ char buffer[64];
67+ strftime(buffer, sizeof(buffer), "%Y-%m-%d %H:%M:%S,000", localtime(&timeValue));
68+ return buffer;
69+}
70+ 
71+uint16_t FloatToBf16Bits(float value)
72+{
73+ uint32_t bits;
74+ std::memcpy(&bits, &value, sizeof(bits));
75+ if (((bits >> 23) & 0xFFU) == 0xFFU && (bits & 0x7FFFFFU) != 0U) {
76+ return static_cast<uint16_t>((bits >> 16) | 0x0040U);
77+ }
78+ const uint32_t rounded = bits + 0x7FFFU + ((bits >> 16) & 1U);
79+ return static_cast<uint16_t>(rounded >> 16);
80+}
81+ 
82+float Bf16BitsToFloat(uint16_t value)
83+{
84+ const uint32_t bits = static_cast<uint32_t>(value) << 16;
85+ float result;
86+ std::memcpy(&result, &bits, sizeof(result));
87+ return result;
88+}
89+ 
90+void InitializeGradsStorage()
91+{
92+ gGradsStorage.clear();
93+ gGradsStorage.reserve(gGrads.size());
94+ for (float value : gGrads) {
95+ gGradsStorage.push_back(FloatToBf16Bits(value));
96+ }
97+}
98+ 
99+Tensor MakeFloatTensor(const vector<int64_t>& shape, vector<float>& data, TensorDesc& desc)
100+{
101+ desc = TensorDesc(Shape(shape), FORMAT_ND, DT_FLOAT);
102+ desc.SetPlacement(kPlacementHost);
103+ desc.SetFormat(FORMAT_ND);
104+ desc.SetRealDimCnt(shape.size());
105+ return Tensor(desc, reinterpret_cast<uint8_t*>(data.data()), data.size() * sizeof(float));
106+}
107+ 
108+Tensor MakeGradsTensor(const vector<int64_t>& shape, TensorDesc& desc)
109+{
110+ desc = TensorDesc(Shape(shape), FORMAT_ND, kGradsDtype);
111+ desc.SetPlacement(kPlacementHost);
112+ desc.SetFormat(FORMAT_ND);
113+ desc.SetRealDimCnt(shape.size());
114+ return Tensor(desc, reinterpret_cast<uint8_t*>(gGradsStorage.data()), gGradsStorage.size() * sizeof(uint16_t));
115+}
116+ 
117+int CreateGraph(vector<Tensor>& input, vector<Operator>& inputs, vector<Operator>& outputs, Graph& graph,
118+ const vector<int64_t>& scaleShape, const vector<int64_t>& varianceShape)
119+{
120+ auto bnInferGrad = op::BNInferGrad("bnInferGrad");
121+ const vector<int64_t> gradsShape = {2, kChannels, 2, 2};
122+ 
123+ TensorDesc gradsDesc;
124+ Tensor gradsTensor = MakeGradsTensor(gradsShape, gradsDesc);
125+ auto grads = op::Data("grads").set_attr_index(0);
126+ grads.update_input_desc_x(gradsDesc);
127+ grads.update_output_desc_y(gradsDesc);
128+ graph.AddOp(grads);
129+ bnInferGrad.set_input_grads(grads);
130+ input.push_back(gradsTensor);
131+ inputs.push_back(grads);
132+ 
133+ TensorDesc scaleDesc;
134+ Tensor scaleTensor = MakeFloatTensor(scaleShape, gScale, scaleDesc);
135+ auto scale = op::Data("scale").set_attr_index(1);
136+ scale.update_input_desc_x(scaleDesc);
137+ scale.update_output_desc_y(scaleDesc);
138+ graph.AddOp(scale);
139+ bnInferGrad.set_input_scale(scale);
140+ input.push_back(scaleTensor);
141+ inputs.push_back(scale);
142+ 
143+ TensorDesc varianceDesc;
144+ Tensor varianceTensor = MakeFloatTensor(varianceShape, gBatchVariance, varianceDesc);
145+ auto batchVariance = op::Data("batchVariance").set_attr_index(2);
146+ batchVariance.update_input_desc_x(varianceDesc);
147+ batchVariance.update_output_desc_y(varianceDesc);
148+ graph.AddOp(batchVariance);
149+ bnInferGrad.set_input_batch_variance(batchVariance);
150+ input.push_back(varianceTensor);
151+ inputs.push_back(batchVariance);
152+ 
153+ TensorDesc outputDesc(Shape(gradsShape), FORMAT_ND, kGradsDtype);
154+ bnInferGrad.update_output_desc_x_backprop(outputDesc);
155+ bnInferGrad.set_attr_epsilon(kEpsilon);
156+ outputs.push_back(bnInferGrad);
157+ return SUCCESS;
158+}
159+ 
160+bool ValidateOutput(const vector<Tensor>& output)
161+{
162+ if (output.size() != 1) {
163+ std::cerr << "Expected one output, got " << output.size() << std::endl;
164+ return false;
165+ }
166+ const Tensor& tensor = output[0];
167+ const int64_t count = tensor.GetTensorDesc().GetShape().GetShapeSize();
168+ if (count != static_cast<int64_t>(gGrads.size())) {
169+ std::cerr << "Unexpected output element count: " << count << std::endl;
170+ return false;
171+ }
172+ 
173+ const uint16_t* result = reinterpret_cast<const uint16_t*>(tensor.GetData());
174+ bool valid = true;
175+ for (int64_t index = 0; index < count; ++index) {
176+ const int64_t channel = (index / kInnerSize) % kChannels;
177+ const float expected = gGrads[index] * gScale[channel] / std::sqrt(gBatchVariance[channel] + kEpsilon);
178+ const float actual = Bf16BitsToFloat(result[index]);
179+ const float difference = std::fabs(actual - expected);
180+ if (difference > kTolerance) {
181+ std::cerr << "Mismatch at " << index << ": actual=" << actual << ", expected=" << expected
182+ << ", diff=" << difference << std::endl;
183+ valid = false;
184+ }
185+ }
186+ return valid;
187+}
188+ 
189+bool ValidateInvalidVarianceRankRejected()
190+{
191+ Graph graph("bn_infer_grad_invalid_variance_rank");
192+ vector<Tensor> input;
193+ vector<Operator> inputs;
194+ vector<Operator> outputs;
195+ if (CreateGraph(input, inputs, outputs, graph, {kChannels}, {kChannels, 1}) != SUCCESS) {
196+ return false;
197+ }
198+ graph.SetInputs(inputs).SetOutputs(outputs);
199+ 
200+ map<AscendString, AscendString> sessionOptions;
201+ Session session(sessionOptions);
202+ const uint32_t graphId = 1;
203+ map<AscendString, AscendString> graphOptions;
204+ Status status = session.AddGraph(graphId, graph, graphOptions);
205+ if (status != SUCCESS) {
206+ return true;
207+ }
208+ vector<Tensor> output;
209+ return session.RunGraph(graphId, input, output) != SUCCESS;
210+}
211+} // namespace
212+ 
213+int main()
214+{
215+ InitializeGradsStorage();
216+ Graph graph("bn_infer_grad_geir_example");
217+ vector<Tensor> input;
218+ vector<Operator> inputs;
219+ vector<Operator> outputs;
220+ 
221+ map<AscendString, AscendString> globalOptions = {
222+ {"ge.exec.deviceId", "0"},
223+ {"ge.graphRunMode", "1"},
224+ };
225+ Status status = GEInitialize(globalOptions);
226+ if (status != SUCCESS) {
227+ std::cerr << GetTime() << " GEInitialize failed: " << status << std::endl;
228+ return FAILED;
229+ }
230+ 
231+ if (CreateGraph(input, inputs, outputs, graph, {kChannels}, {kChannels}) != SUCCESS) {
232+ GEFinalize();
233+ return FAILED;
234+ }
235+ graph.SetInputs(inputs).SetOutputs(outputs);
236+ 
237+ map<AscendString, AscendString> sessionOptions;
238+ Session session(sessionOptions);
239+ const uint32_t graphId = 0;
240+ map<AscendString, AscendString> graphOptions;
241+ status = session.AddGraph(graphId, graph, graphOptions);
242+ if (status != SUCCESS) {
243+ std::cerr << GetTime() << " AddGraph failed: " << status << std::endl;
244+ GEFinalize();
245+ return FAILED;
246+ }
247+ 
248+ vector<Tensor> output;
249+ status = session.RunGraph(graphId, input, output);
250+ if (status != SUCCESS) {
251+ std::cerr << GetTime() << " RunGraph failed: " << status << std::endl;
252+ GEFinalize();
253+ return FAILED;
254+ }
255+ 
256+ const bool valid = ValidateOutput(output);
257+ const bool invalidRankRejected = ValidateInvalidVarianceRankRejected();
258+ status = GEFinalize();
259+ if (status != SUCCESS) {
260+ std::cerr << GetTime() << " GEFinalize failed: " << status << std::endl;
261+ return FAILED;
262+ }
263+ if (!valid) {
264+ std::cerr << "BNInferGrad BF16 GEIR numerical validation failed." << std::endl;
265+ return FAILED;
266+ }
267+ if (!invalidRankRejected) {
268+ std::cerr << "BNInferGrad invalid batch_variance rank was not rejected." << std::endl;
269+ return FAILED;
270+ }
271+ std::cout << "BNInferGrad BF16 GEIR numerical validation passed; invalid batch_variance rank rejected."
272+ << std::endl;
273+ return SUCCESS;
274+}
Anorm/bn_infer_grad/examples/test_geir_bn_infer_grad.cpp+309-0
@@ -0,0 +1,309 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file test_geir_bn_infer_grad.cpp
15+ * \brief BNInferGrad GEIR graph-mode example with numerical validation.
16+ */
17+ 
18+#include <cmath>
19+#include <cstdint>
20+#include <cstdio>
21+#include <cstring>
22+#include <ctime>
23+#include <iostream>
24+#include <map>
25+#include <string>
26+#include <vector>
27+ 
28+#include "array_ops.h"
29+#include "ge_api.h"
30+#include "ge_api_types.h"
31+#include "ge_error_codes.h"
32+#include "ge_ir_build.h"
33+#include "graph.h"
34+#include "tensor.h"
35+#include "types.h"
36+ 
37+#include "../op_graph/bn_infer_grad_proto.h"
38+ 
39+#define FAILED -1
40+#define SUCCESS 0
41+ 
42+using namespace ge;
43+using std::map;
44+using std::string;
45+using std::vector;
46+ 
47+namespace {
48+constexpr float kEpsilon = 1.0F;
49+constexpr int64_t kBatchSize = 128;
50+constexpr int64_t kChannels = 3;
51+constexpr int64_t kHeight = 32;
52+constexpr int64_t kWidth = 32;
53+constexpr int64_t kInnerSize = kHeight * kWidth;
54+constexpr DataType kGradsDtype = DT_FLOAT16;
55+constexpr float kTolerance = 5.0e-3F;
56+ 
57+vector<float> gGrads;
58+vector<float> gScale = {0.75F, -1.25F, 2.5F};
59+vector<float> gBatchVariance = {0.5F, 2.0F, 7.0F};
60+vector<uint16_t> gGradsStorage;
61+ 
62+string GetTime()
63+{
64+ time_t timeValue;
65+ time(&timeValue);
66+ char buffer[64];
67+ strftime(buffer, sizeof(buffer), "%Y-%m-%d %H:%M:%S,000", localtime(&timeValue));
68+ return buffer;
69+}
70+ 
71+uint16_t FloatToHalfBits(float value)
72+{
73+ uint32_t bits;
74+ std::memcpy(&bits, &value, sizeof(bits));
75+ const uint32_t sign = (bits >> 31) & 1U;
76+ const uint32_t exponent = (bits >> 23) & 0xFFU;
77+ const uint32_t mantissa = bits & 0x7FFFFFU;
78+ if (exponent == 0xFFU) {
79+ return static_cast<uint16_t>((sign << 15) | 0x7C00U | (mantissa == 0 ? 0U : 0x0200U));
80+ }
81+ const int32_t halfExponent = static_cast<int32_t>(exponent) - 127 + 15;
82+ if (halfExponent <= 0) {
83+ return static_cast<uint16_t>(sign << 15);
84+ }
85+ if (halfExponent >= 31) {
86+ return static_cast<uint16_t>((sign << 15) | 0x7C00U);
87+ }
88+ uint16_t result = static_cast<uint16_t>((sign << 15) | (static_cast<uint32_t>(halfExponent) << 10) |
89+ (mantissa >> 13));
90+ const uint32_t remainder = mantissa & 0x1FFFU;
91+ if (remainder > 0x1000U || (remainder == 0x1000U && (result & 1U) != 0U)) {
92+ ++result;
93+ }
94+ return result;
95+}
96+ 
97+float HalfBitsToFloat(uint16_t value)
98+{
99+ const uint32_t sign = (value >> 15) & 1U;
100+ const uint32_t exponent = (value >> 10) & 0x1FU;
101+ const uint32_t mantissa = value & 0x03FFU;
102+ uint32_t bits;
103+ if (exponent == 0) {
104+ bits = sign << 31;
105+ } else if (exponent == 0x1FU) {
106+ bits = (sign << 31) | 0x7F800000U | (mantissa << 13);
107+ } else {
108+ bits = (sign << 31) | ((exponent - 15U + 127U) << 23) | (mantissa << 13);
109+ }
110+ float result;
111+ std::memcpy(&result, &bits, sizeof(result));
112+ return result;
113+}
114+ 
115+void InitializeGradsStorage()
116+{
117+ const int64_t elementCount = kBatchSize * kChannels * kInnerSize;
118+ gGrads.clear();
119+ gGrads.reserve(elementCount);
120+ for (int64_t index = 0; index < elementCount; ++index) {
121+ gGrads.push_back(0.25F + static_cast<float>(index % 31) * 0.03125F);
122+ }
123+ 
124+ gGradsStorage.clear();
125+ gGradsStorage.reserve(gGrads.size());
126+ for (float value : gGrads) {
127+ gGradsStorage.push_back(FloatToHalfBits(value));
128+ }
129+}
130+ 
131+Tensor MakeFloatTensor(const vector<int64_t>& shape, vector<float>& data, TensorDesc& desc)
132+{
133+ desc = TensorDesc(Shape(shape), FORMAT_ND, DT_FLOAT);
134+ desc.SetPlacement(kPlacementHost);
135+ desc.SetFormat(FORMAT_ND);
136+ desc.SetRealDimCnt(shape.size());
137+ return Tensor(desc, reinterpret_cast<uint8_t*>(data.data()), data.size() * sizeof(float));
138+}
139+ 
140+Tensor MakeGradsTensor(const vector<int64_t>& shape, TensorDesc& desc)
141+{
142+ desc = TensorDesc(Shape(shape), FORMAT_ND, kGradsDtype);
143+ desc.SetPlacement(kPlacementHost);
144+ desc.SetFormat(FORMAT_ND);
145+ desc.SetRealDimCnt(shape.size());
146+ return Tensor(desc, reinterpret_cast<uint8_t*>(gGradsStorage.data()), gGradsStorage.size() * sizeof(uint16_t));
147+}
148+ 
149+int CreateGraph(vector<Tensor>& input, vector<Operator>& inputs, vector<Operator>& outputs, Graph& graph,
150+ const vector<int64_t>& scaleShape, const vector<int64_t>& varianceShape)
151+{
152+ auto bnInferGrad = op::BNInferGrad("bnInferGrad");
153+ // This shape exceeds one per-core UB tile on Ascend950. Each participating
154+ // core rereads the same channel-varying scale/variance broadcast inputs.
155+ const vector<int64_t> gradsShape = {kBatchSize, kChannels, kHeight, kWidth};
156+ 
157+ TensorDesc gradsDesc;
158+ Tensor gradsTensor = MakeGradsTensor(gradsShape, gradsDesc);
159+ auto grads = op::Data("grads").set_attr_index(0);
160+ grads.update_input_desc_x(gradsDesc);
161+ grads.update_output_desc_y(gradsDesc);
162+ graph.AddOp(grads);
163+ bnInferGrad.set_input_grads(grads);
164+ input.push_back(gradsTensor);
165+ inputs.push_back(grads);
166+ 
167+ TensorDesc scaleDesc;
168+ Tensor scaleTensor = MakeFloatTensor(scaleShape, gScale, scaleDesc);
169+ auto scale = op::Data("scale").set_attr_index(1);
170+ scale.update_input_desc_x(scaleDesc);
171+ scale.update_output_desc_y(scaleDesc);
172+ graph.AddOp(scale);
173+ bnInferGrad.set_input_scale(scale);
174+ input.push_back(scaleTensor);
175+ inputs.push_back(scale);
176+ 
177+ TensorDesc varianceDesc;
178+ Tensor varianceTensor = MakeFloatTensor(varianceShape, gBatchVariance, varianceDesc);
179+ auto batchVariance = op::Data("batchVariance").set_attr_index(2);
180+ batchVariance.update_input_desc_x(varianceDesc);
181+ batchVariance.update_output_desc_y(varianceDesc);
182+ graph.AddOp(batchVariance);
183+ bnInferGrad.set_input_batch_variance(batchVariance);
184+ input.push_back(varianceTensor);
185+ inputs.push_back(batchVariance);
186+ 
187+ TensorDesc outputDesc(Shape(gradsShape), FORMAT_ND, kGradsDtype);
188+ bnInferGrad.update_output_desc_x_backprop(outputDesc);
189+ bnInferGrad.set_attr_epsilon(kEpsilon);
190+ outputs.push_back(bnInferGrad);
191+ return SUCCESS;
192+}
193+ 
194+bool ValidateOutput(const vector<Tensor>& output)
195+{
196+ if (output.size() != 1) {
197+ std::cerr << "Expected one output, got " << output.size() << std::endl;
198+ return false;
199+ }
200+ const Tensor& tensor = output[0];
201+ const int64_t count = tensor.GetTensorDesc().GetShape().GetShapeSize();
202+ if (count != static_cast<int64_t>(gGrads.size())) {
203+ std::cerr << "Unexpected output element count: " << count << std::endl;
204+ return false;
205+ }
206+ 
207+ const uint16_t* result = reinterpret_cast<const uint16_t*>(tensor.GetData());
208+ bool valid = true;
209+ for (int64_t index = 0; index < count; ++index) {
210+ const int64_t channel = (index / kInnerSize) % kChannels;
211+ const float expected = gGrads[index] * gScale[channel] / std::sqrt(gBatchVariance[channel] + kEpsilon);
212+ const float actual = HalfBitsToFloat(result[index]);
213+ const float difference = std::fabs(actual - expected);
214+ if (difference > kTolerance) {
215+ std::cerr << "Mismatch at " << index << ": actual=" << actual << ", expected=" << expected
216+ << ", diff=" << difference << std::endl;
217+ valid = false;
218+ }
219+ }
220+ return valid;
221+}
222+ 
223+bool ValidateInvalidScaleRankRejected()
224+{
225+ Graph graph("bn_infer_grad_invalid_scale_rank");
226+ vector<Tensor> input;
227+ vector<Operator> inputs;
228+ vector<Operator> outputs;
229+ if (CreateGraph(input, inputs, outputs, graph, {1, kChannels}, {kChannels}) != SUCCESS) {
230+ return false;
231+ }
232+ graph.SetInputs(inputs).SetOutputs(outputs);
233+ 
234+ map<AscendString, AscendString> sessionOptions;
235+ Session session(sessionOptions);
236+ const uint32_t graphId = 1;
237+ map<AscendString, AscendString> graphOptions;
238+ Status status = session.AddGraph(graphId, graph, graphOptions);
239+ if (status != SUCCESS) {
240+ return true;
241+ }
242+ vector<Tensor> output;
243+ return session.RunGraph(graphId, input, output) != SUCCESS;
244+}
245+} // namespace
246+ 
247+int main()
248+{
249+ InitializeGradsStorage();
250+ Graph graph("bn_infer_grad_geir_example");
251+ vector<Tensor> input;
252+ vector<Operator> inputs;
253+ vector<Operator> outputs;
254+ 
255+ map<AscendString, AscendString> globalOptions = {
256+ {"ge.exec.deviceId", "0"},
257+ {"ge.graphRunMode", "1"},
258+ };
259+ Status status = GEInitialize(globalOptions);
260+ if (status != SUCCESS) {
261+ std::cerr << GetTime() << " GEInitialize failed: " << status << std::endl;
262+ return FAILED;
263+ }
264+ 
265+ if (CreateGraph(input, inputs, outputs, graph, {kChannels}, {kChannels}) != SUCCESS) {
266+ GEFinalize();
267+ return FAILED;
268+ }
269+ graph.SetInputs(inputs).SetOutputs(outputs);
270+ 
271+ map<AscendString, AscendString> sessionOptions;
272+ Session session(sessionOptions);
273+ const uint32_t graphId = 0;
274+ map<AscendString, AscendString> graphOptions;
275+ status = session.AddGraph(graphId, graph, graphOptions);
276+ if (status != SUCCESS) {
277+ std::cerr << GetTime() << " AddGraph failed: " << status << std::endl;
278+ GEFinalize();
279+ return FAILED;
280+ }
281+ 
282+ vector<Tensor> output;
283+ status = session.RunGraph(graphId, input, output);
284+ if (status != SUCCESS) {
285+ std::cerr << GetTime() << " RunGraph failed: " << status << std::endl;
286+ GEFinalize();
287+ return FAILED;
288+ }
289+ 
290+ const bool valid = ValidateOutput(output);
291+ const bool invalidRankRejected = ValidateInvalidScaleRankRejected();
292+ status = GEFinalize();
293+ if (status != SUCCESS) {
294+ std::cerr << GetTime() << " GEFinalize failed: " << status << std::endl;
295+ return FAILED;
296+ }
297+ if (!valid) {
298+ std::cerr << "BNInferGrad FP16 GEIR numerical validation failed." << std::endl;
299+ return FAILED;
300+ }
301+ if (!invalidRankRejected) {
302+ std::cerr << "BNInferGrad invalid scale rank was not rejected." << std::endl;
303+ return FAILED;
304+ }
305+ std::cout << "BNInferGrad FP16 multi-core broadcast GEIR numerical validation passed; "
306+ "invalid scale rank rejected."
307+ << std::endl;
308+ return SUCCESS;
309+}
Anorm/bn_infer_grad/op_graph/bn_infer_grad_graph_infer.cpp+28-0
@@ -0,0 +1,28 @@
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+/* Generated By CANNBot */
12+ 
13+#include "op_common/log/log.h"
14+#include "register/op_impl_registry.h"
15+ 
16+using namespace ge;
17+namespace ops {
18+static graphStatus InferDataType4SameAsInput(gert::InferDataTypeContext* context)
19+{
20+ OP_LOGD(context->GetNodeName(), "InferDataType4SameAsInput enter");
21+ auto input_x_dtype = context->GetInputDataType(0);
22+ context->SetOutputDataType(0, input_x_dtype);
23+ OP_LOGD(context->GetNodeName(), "InferDataType4SameAsInput end");
24+ return GRAPH_SUCCESS;
25+}
26+ 
27+IMPL_OP(BNInferGrad).InferDataType(InferDataType4SameAsInput);
28+} // namespace ops
Anorm/bn_infer_grad/op_graph/bn_infer_grad_proto.h+60-0
@@ -0,0 +1,60 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+#ifndef OPS_BN_INFER_GRAD_PROTO_H_
14+#define OPS_BN_INFER_GRAD_PROTO_H_
15+ 
16+#include "graph/operator_reg.h"
17+ 
18+namespace ge {
19+ 
20+/**
21+ * @brief Performs the backpropagation of BatchNorm for inference .
22+ 
23+ 
24+ * @par Inputs:
25+ * Three inputs, including:
26+ * @li grads: A tensor of type float16 or float32 or bfloat16. Indicates the gradient of the BathNorm output parameter
27+ "y".
28+ * Shape support 4D and 5D. Format support NHWC, NCHW or NC1HWC0.
29+ * @li scale: A 1D tensor of type float32. Shape must be C channel.
30+ * Specifies the scaling factor. Has the same format as "grads".
31+ * @li batch_variance: A 1D tensor of type float32. Calculated variance, that is, the value of BatchNorm output
32+ parameter "batch_variance".
33+ * Shape must be C channel. Has the same format as "grads". \n
34+ 
35+ 
36+ * @par Attributes:
37+ * epsilon: An optional float32. Defaults to "0.0001". A small float number
38+ * added to the variance of "batch_variance". \n
39+ 
40+ 
41+ * @par Outputs:
42+ * x_backprop: A tensor of type float16 or float32 or bfloat16.
43+ * Indicates the gradient of BathNorm input data "x".
44+ * Has the same type, shape and format as "grads". \n
45+ 
46+ 
47+ * @attention Constraints:
48+ * The preceding layer of this operator must be operator BatchNorm.
49+ */
50+REG_OP(BNInferGrad)
51+ .INPUT(grads, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
52+ .INPUT(scale, TensorType({DT_FLOAT}))
53+ .INPUT(batch_variance, TensorType({DT_FLOAT}))
54+ .OUTPUT(x_backprop, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
55+ .ATTR(epsilon, Float, 0.0001)
56+ .OP_END_FACTORY_REG(BNInferGrad)
57+ 
58+} // namespace ge
59+ 
60+#endif // OPS_BN_INFER_GRAD_PROTO_H_
Anorm/bn_infer_grad/op_host/arch35/bn_infer_grad_tiling_arch35.cpp+537-0
@@ -0,0 +1,537 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file bn_infer_grad_tiling.cpp
15+ * \brief BNInferGrad Tiling 实现（arch35 架构，Broadcast 范式）
16+ *
17+ * 结构对齐参考算子 adam_apply_one_assign：单次归一化 → 映射 rank → 模板填充。
18+ * BNInferGrad 特有处理：
19+ * - scale / batch_variance 为 [C]，沿 grads 的 axis=1（C 维）广播；
20+ * Tiling 侧将其 reshape 为与 grads rank 对齐的 [1, C, 1, ...]，
21+ * 交给通用 PadAndSqueeze / PrecomputeInputStrides（size-1 轴 stride=0）自动生成广播 stride。
22+ * - epsilon 为 OpDef Attr，读入 TilingData.epsilon_value。
23+ */
24+#include "bn_infer_grad_tiling_arch35.h"
25+#include "bn_infer_grad_tiling_capacity.h"
26+#include "../../op_kernel/arch35/bn_infer_grad_tiling_key.h"
27+#include <algorithm>
28+#include <sstream>
29+#include <graph/utils/type_utils.h>
30+#include "register/op_impl_registry.h"
31+#include "op_common/log/log.h"
32+#include "tiling/platform/platform_ascendc.h"
33+ 
34+using namespace ge;
35+ 
36+// ============================================================
37+// Tiling 函数模块 — namespace bn_infer_grad（无状态）
38+// ============================================================
39+namespace bn_infer_grad {
40+ 
41+bool CheckBroadcastShape(const std::vector<std::vector<int64_t>>& padded_in,
42+ const std::vector<std::vector<int64_t>>& padded_out, int64_t max_rank)
43+{
44+ for (int64_t d = 0; d < max_rank; d++) {
45+ int64_t ref = -1;
46+ for (size_t i = 0; i < padded_in.size(); i++) {
47+ if (padded_in[i][d] != 1) {
48+ if (ref == -1)
49+ ref = padded_in[i][d];
50+ else if (padded_in[i][d] != ref)
51+ return false;
52+ }
53+ }
54+ for (size_t i = 0; i < padded_out.size(); i++) {
55+ if (padded_out[i][d] != 1) {
56+ if (ref == -1)
57+ ref = padded_out[i][d];
58+ else if (padded_out[i][d] != ref)
59+ return false;
60+ }
61+ }
62+ }
63+ return true;
64+}
65+ 
66+bool PadAndSqueeze(const std::vector<std::vector<int64_t>>& input_shapes,
67+ const std::vector<std::vector<int64_t>>& output_shapes, std::vector<int64_t>& maximum_bro_shape,
68+ std::vector<std::vector<int64_t>>& normal_input_shapes,
69+ std::vector<std::vector<int64_t>>& normal_output_shapes)
70+{
71+ int64_t num_inputs = (int64_t)input_shapes.size();
72+ int64_t num_outputs = (int64_t)output_shapes.size();
73+ int64_t max_rank = 0;
74+ for (auto& s : input_shapes)
75+ max_rank = std::max(max_rank, (int64_t)s.size());
76+ for (auto& s : output_shapes)
77+ max_rank = std::max(max_rank, (int64_t)s.size());
78+ 
79+ // 补 1 — 低 rank 在 shape 最前面补 1
80+ auto pad = [&](const std::vector<int64_t>& s) {
81+ std::vector<int64_t> p;
82+ p.assign(max_rank - (int64_t)s.size(), 1);
83+ p.insert(p.end(), s.begin(), s.end());
84+ return p;
85+ };
86+ std::vector<std::vector<int64_t>> padded_in(num_inputs), padded_out(num_outputs);
87+ for (int64_t i = 0; i < num_inputs; i++)
88+ padded_in[i] = pad(input_shapes[i]);
89+ for (int64_t i = 0; i < num_outputs; i++)
90+ padded_out[i] = pad(output_shapes[i]);
91+ 
92+ maximum_bro_shape.clear();
93+ normal_input_shapes.assign(num_inputs, std::vector<int64_t>());
94+ normal_output_shapes.assign(num_outputs, std::vector<int64_t>());
95+ for (int64_t d = 0; d < max_rank; d++) {
96+ bool all_one = true;
97+ int64_t max_dim = 0;
98+ for (int64_t i = 0; i < num_inputs; i++) {
99+ if (padded_in[i][d] != 1)
100+ all_one = false;
Anorm/bn_infer_grad/op_host/arch35/bn_infer_grad_tiling_arch35.h+81-0
@@ -0,0 +1,81 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file bn_infer_grad_tiling_arch35.h
15+ * \brief BNInferGrad Tiling — arch35 头文件
16+ *
17+ * 无状态 Tiling 函数模块（与 CANN 框架零耦合），对齐 Broadcast 范式参考算子 adam_apply_one_assign。
18+ */
19+#pragma once
20+#include <cstdint>
21+#include <vector>
22+#include <exe_graph/runtime/tiling_context.h>
23+#include "../../op_kernel/arch35/bn_infer_grad_tiling_data.h"
24+ 
25+namespace bn_infer_grad {
26+ 
27+// numpy 广播兼容性校验（补 1 后逐维核对）
28+bool CheckBroadcastShape(const std::vector<std::vector<int64_t>>& padded_in,
29+ const std::vector<std::vector<int64_t>>& padded_out, int64_t max_rank);
30+ 
31+// 去 size-1 维 + 前补 1，生成 max_bro_shape 及各输入/输出的 normal shape
32+bool PadAndSqueeze(const std::vector<std::vector<int64_t>>& input_shapes,
33+ const std::vector<std::vector<int64_t>>& output_shapes, std::vector<int64_t>& maximum_bro_shape,
34+ std::vector<std::vector<int64_t>>& normal_input_shapes,
35+ std::vector<std::vector<int64_t>>& normal_output_shapes);
36+ 
37+// UB 切分轴选择（per_buf = (UB/P)&~31）
38+// c_axis: 归一化 max_bro_shape 中 C 维下标；禁止选 C 为切分轴，但 C 大小仍计入 inner 累乘。C=1 时传 -1。
39+bool FindSplitAxis(const std::vector<int64_t>& max_bro_shape, int64_t dtype_size, int64_t ub_per_core,
40+ int64_t phys_nodes, int64_t c_axis, SplitResult& out);
41+ 
42+// 多核切分
43+bool MultiCoreSplit(const std::vector<int64_t>& max_bro_shape, const SplitResult& ub_split, int64_t max_cores,
44+ MultiCoreResult& out);
45+ 
46+// GM 物理 stride 预计算（size-1 轴 = 0，为广播轴）
47+bool PrecomputeInputStrides(const std::vector<int64_t>& s, std::vector<int64_t>& strides);
48+bool PrecomputeOutputStrides(const std::vector<int64_t>& s, std::vector<int64_t>& strides);
49+ 
50+} // namespace bn_infer_grad
51+ 
52+namespace optiling {
53+ 
54+struct BNInferGradCompileInfo {
55+ uint64_t coreNum;
56+ uint64_t ubSize;
57+};
58+ 
59+class BNInferGradTiling {
60+public:
61+ explicit BNInferGradTiling(gert::TilingContext* ctx);
62+ ge::graphStatus RunTiling();
63+ 
64+private:
65+ ge::graphStatus GetShapeInfo();
66+ template <int64_t R>
67+ ge::graphStatus DoTilingAndSet();
68+ 
69+ gert::TilingContext* ctx_;
70+ std::vector<std::vector<int64_t>> raw_input_shapes_;
71+ std::vector<std::vector<int64_t>> raw_output_shapes_;
72+ std::vector<int64_t> max_bro_shape_;
73+ std::vector<std::vector<int64_t>> normal_input_shapes_;
74+ std::vector<std::vector<int64_t>> normal_output_shapes_;
75+ int64_t dtype_size_ = 0;
76+ int64_t rank_ = 0;
77+ float epsilon_ = 0.0001f;
78+ bool is_empty_ = false; // grads 任意维度为 0 → 空 tensor，不执行 Kernel（DESIGN.md §5.4.1）
79+};
80+ 
81+} // namespace optiling
Anorm/bn_infer_grad/op_host/arch35/bn_infer_grad_tiling_capacity.h+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+/* Generated By CANNBot */
12+ 
13+#ifndef BN_INFER_GRAD_TILING_CAPACITY_H_
14+#define BN_INFER_GRAD_TILING_CAPACITY_H_
15+ 
16+#include <cstdint>
17+ 
18+namespace bn_infer_grad {
19+ 
20+struct TilingCapacity {
21+ int64_t perBufBytes;
22+ int64_t perBufElements;
23+};
24+ 
25+inline bool GetTilingCapacity(int64_t ubPerCore, int64_t physicalNodes, int64_t dtypeSize, TilingCapacity& capacity)
26+{
27+ capacity = {0, 0};
28+ if (ubPerCore <= 0 || physicalNodes <= 0 || dtypeSize <= 0) {
29+ return false;
30+ }
31+ 
32+ capacity.perBufBytes = (ubPerCore / physicalNodes) & ~31LL;
33+ if (capacity.perBufBytes <= 0) {
34+ return false;
35+ }
36+ 
37+ capacity.perBufElements = capacity.perBufBytes / dtypeSize;
38+ return capacity.perBufElements > 0;
39+}
40+ 
41+} // namespace bn_infer_grad
42+ 
43+#endif // BN_INFER_GRAD_TILING_CAPACITY_H_
Anorm/bn_infer_grad/op_host/bn_infer_grad_def.cpp+72-0
@@ -0,0 +1,72 @@
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+/* Generated By CANNBot - Iter2 */
12+ 
13+/*!
14+ * \file bn_infer_grad_def.cpp
15+ * \brief BNInferGrad 算子定义，声明输入输出、属性和算子配置
16+ *
17+ * 3 输入（grads, scale, batch_variance）+ 1 属性（epsilon, float）→ 1 输出（x_backprop）
18+ *
19+ * dtype 策略：OpDef 声明并由Kernel直接支持fp16/fp32/bf16。
20+ *
21+ * dtype 约束（spec.yaml）:
22+ * - grads / x_backprop: float16/float32/bfloat16
23+ * - scale / batch_variance: 固定 float32
24+ * - epsilon: float 标量，默认 0.0001
25+ */
26+#include "register/op_def_registry.h"
27+ 
28+namespace ops {
29+class BNInferGrad : public OpDef {
30+public:
31+ explicit BNInferGrad(const char* name) : OpDef(name)
32+ {
33+ this->Input("grads") // 输入0：上游梯度 [N,C,...]
34+ .ParamType(REQUIRED)
35+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16}) // 支持 float16/float32/bfloat16
36+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) // 每个 dtype 对应一个 Format
AlfengYuan
AlfengYuanAlfengYuan7月27日

只支持ND格式吗，和A2/A3上是一致的吗？NHWC，NC1HWC0， NCHW等特定格式确认不需要支持是吗？

likedislike
Y
yulianjie
7月27日 评论：
37+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
38+ .AutoContiguous();
39+ this->Input("scale") // 输入1：BN 缩放参数 [C]
40+ .ParamType(REQUIRED)
41+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})
42+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
43+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
44+ .AutoContiguous();
45+ this->Input("batch_variance") // 输入2：运行时方差 [C]
46+ .ParamType(REQUIRED)
47+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})
48+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
49+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
50+ .AutoContiguous();
51+ this->Output("x_backprop") // 输出0：输入 x 的梯度 [N,C,...]
52+ .ParamType(REQUIRED)
53+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16}) // 与 grads 同 dtype
54+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) // 每个 dtype 对应一个 Format
55+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
56+ .AutoContiguous();
57+ 
58+ this->Attr("epsilon").AttrType(OPTIONAL).Float(0.0001f); // 防除零常数，默认 1e-4
59+ 
60+ OpAICoreConfig aiCoreConfig;
61+ aiCoreConfig.DynamicCompileStaticFlag(true)
62+ .DynamicFormatFlag(true)
63+ .DynamicRankSupportFlag(true)
64+ .DynamicShapeSupportFlag(true)
65+ .NeedCheckSupportFlag(false)
66+ .PrecisionReduceFlag(true)
67+ .ExtendCfgInfo("opFile.value", "bn_infer_grad");
68+ this->AICore().AddConfig("ascend950", aiCoreConfig);
69+ }
70+};
71+OP_ADD(BNInferGrad);
72+} // namespace ops
Anorm/bn_infer_grad/op_host/bn_infer_grad_infershape.cpp+19-0
@@ -0,0 +1,19 @@
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+/* Generated By CANNBot */
12+ 
13+#include "infershape_elewise_util.h"
14+#include "register/op_impl_registry.h"
15+ 
16+using namespace ge;
17+namespace ops {
18+IMPL_OP_INFERSHAPE(BNInferGrad).InferShape(Ops::Base::InferShape4Elewise);
19+} // namespace ops
Anorm/bn_infer_grad/op_kernel/arch35/bn_infer_grad.h+342-0
@@ -0,0 +1,342 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file bn_infer_grad.h
15+ * \brief BNInferGrad Kernel 实现（arch35 / RegBase / Broadcast 范式）
16+ *
17+ * 迭代二：多 TilingKey（RANK4+RANK8）+ 全 dtype（float16/float32/bfloat16）
Y
Yyulianjie7月28日

多余的注释

likedislike
18+ *
19+ * 计算链（逐 tile，NDDMA 随路广播后所有 buffer 均为满 tile 大小）：
20+ * FP32 路径：
21+ * ① var_eps = batch_variance + epsilon (Adds)
22+ * ② var_sqrt = sqrt(var_eps) (Sqrt)
23+ * ③ rsqrt = 1.0 / var_sqrt (Duplicate(1.0) + Div，vdiv 而非 rsqrt)
24+ * ④ scale_rsqrt= scale * rsqrt (Mul)
25+ * ⑤ x_backprop = grads * scale_rsqrt (Mul)
26+ *
27+ * FP16/BF16 路径：
28+ * ① Cast grads → grads_fp32 (Cast, CAST_NONE)
29+ * ② var_eps = batch_variance + epsilon (Adds)
30+ * ③ var_sqrt = sqrt(var_eps) (Sqrt)
31+ * ④ rsqrt = 1.0 / var_sqrt (Duplicate(1.0) + Div)
32+ * ⑤ scale_rsqrt= scale * rsqrt (Mul)
33+ * ⑥ x_backprop_fp32 = grads_fp32 * scale_rsqrt (Mul)
34+ * ⑦ Cast x_backprop_fp32 → x_backprop (Cast, CAST_ROUND)
35+ *
36+ * 内存模型对齐 Broadcast 范式参考算子 adam_apply_one_assign：
37+ * TBuf + 分 dtype NDDMA DataCopy + 多维坐标系统 (GetCoreRange / FlatToEffectiveCoord / CalcInputOffset)。
38+ * scale / batch_variance [C] 通过 loopSrcStride=0 沿非 C 轴随路广播到满 tile。
39+ */
40+ 
41+#pragma once
42+#include "kernel_operator.h"
43+#include "bn_infer_grad_tiling_data.h"
44+#include "bn_infer_grad_tiling_key.h"
45+ 
46+// ============================================================
47+// Kernel 侧辅助函数 (int64_t* 版本, 无 std::vector) — 对齐 adam 参考实现
48+// ============================================================
49+ 
50+__aicore__ inline void GetCoreRange(int64_t core_id, int64_t tiles_main, int64_t cores_tail, int64_t& start,
51+ int64_t& end)
52+{
53+ if (core_id < cores_tail) {
54+ start = core_id * (tiles_main + 1);
55+ end = start + tiles_main + 1;
56+ } else {
57+ start = cores_tail * (tiles_main + 1) + (core_id - cores_tail) * tiles_main;
58+ end = start + tiles_main;
59+ }
60+}
61+ 
62+__aicore__ inline int64_t GetUBSplitRange(int64_t a_o_off, int64_t a_o, int64_t a_i, int64_t a_i_tail)
63+{
64+ return (a_o_off == a_o - 1) ? a_i_tail : a_i;
65+}
66+ 
67+__aicore__ inline bool FlatToEffectiveCoord(int64_t flat, const int64_t* max_bro_shape, int64_t rank,
68+ int64_t split_axis, int64_t a_i, int64_t a_o, int64_t* eff_coord)
69+{
70+ for (int64_t d = 0; d < rank; d++)
71+ eff_coord[d] = 0;
72+ int64_t a_o_off = flat % a_o;
73+ int64_t outer = flat / a_o;
74+ for (int64_t d = split_axis - 1; d >= 0; d--) {
75+ eff_coord[d] = outer % max_bro_shape[d];
76+ outer /= max_bro_shape[d];
77+ }
78+ eff_coord[split_axis] = a_o_off * a_i;
79+ return true;
80+}
81+ 
82+__aicore__ inline int64_t CalcInputOffset(const int64_t* eff_coord, const int64_t* strides, int64_t rank)
83+{
84+ int64_t offset = 0;
85+ for (int64_t d = 0; d < rank; d++)
86+ offset += eff_coord[d] * strides[d];
87+ return offset; // 元素个数，gmIn_[] 的 index
88+}
89+ 
90+__aicore__ inline int64_t CalcOutputTransferCount(const int64_t* normal_shape, int64_t rank, int64_t split_axis,
91+ int64_t a_i_seg)
92+{
93+ int64_t split_elems = (normal_shape[split_axis] == 1) ? 1 : a_i_seg;
94+ int64_t inner_elems = 1;
95+ for (int64_t d = split_axis + 1; d < rank; d++)
96+ inner_elems *= normal_shape[d];
97+ return split_elems * inner_elems;
98+}
99+ 
100+template <typename T, int64_t RANK>
101+class BNInferGradKernel {
102+ static constexpr int64_t kMaxNdDmaDims = 5;
103+ static constexpr int64_t ND = (RANK <= kMaxNdDmaDims) ? RANK : kMaxNdDmaDims;
104+ 
105+ AscendC::TPipe pipe_; // 流水线管理器
106+ const BNInferGradTilingData<RANK>* td_; // TilingData 指针
107+ AscendC::GlobalTensor<T> gmGrads_; // GM 输入 grads，dtype 跟随 T
108+ AscendC::GlobalTensor<float> gmScale_; // GM 输入 scale，固定 FP32
109+ AscendC::GlobalTensor<float> gmVariance_; // GM 输入 batch_variance，固定 FP32
110+ AscendC::GlobalTensor<T> gmOut_; // GM 输出 x_backprop
111+ // ⛔ Broadcast 算子一律用 TBuf 管理 UB，禁止 TQue / 裸指针
112+ AscendC::TBuf<AscendC::TPosition::VECCALC> buf_[kPhysNodes]; // UB buffer (P 槽位)
113+ AscendC::MultiCopyParams<T, ND> nddmaParamsGrads_;
114+ AscendC::MultiCopyParams<float, ND> nddmaParamsScale_;
115+ AscendC::MultiCopyParams<float, ND> nddmaParamsVariance_;
116+ int64_t nddmaOuterItersGrads_;
117+ int64_t nddmaOuterItersScale_;
118+ int64_t nddmaOuterItersVariance_;
119+ int64_t nddma_dims_; // 实际 NDDMA 维数
120+ 
121+ // GM 输入槽位语义
122+ static constexpr int IN_GRADS = 0;
123+ static constexpr int IN_SCALE = 1;
124+ static constexpr int IN_VARIANCE = 2;
125+ static constexpr int OUT_XBP = 0;
126+ 
127+public:
128+ __aicore__ inline BNInferGradKernel() {}
129+ 
130+ // GM 绑定 + TBuf 分配 + 分 dtype NDDMA 参数预计算，实现参照仓内混合 dtype Broadcast 算子
131+ __aicore__ inline void Init(GM_ADDR inputs[kMaxInputSlots], GM_ADDR outputs[kMaxOutputSlots],
132+ const BNInferGradTilingData<RANK>* td)
133+ {
134+ td_ = td;
135+ gmGrads_.SetGlobalBuffer((__gm__ T*)inputs[IN_GRADS]);
136+ gmScale_.SetGlobalBuffer((__gm__ float*)inputs[IN_SCALE]);
137+ gmVariance_.SetGlobalBuffer((__gm__ float*)inputs[IN_VARIANCE]);
138+ gmOut_.SetGlobalBuffer((__gm__ T*)outputs[OUT_XBP]);
139+ for (int i = 0; i < kPhysNodes; i++)
140+ pipe_.InitBuffer(buf_[i], td_->per_buf_bytes);
141+ 
142+ const int64_t* dstShape = td_->max_bro_shape;
143+ int64_t k = td_->split.axis;
144+ nddma_dims_ = (RANK - k <= ND) ? (RANK - k) : ND;
145+ SetupNddmaParams(nddmaParamsGrads_, IN_GRADS, nddmaOuterItersGrads_, k, dstShape);
146+ SetupNddmaParams(nddmaParamsScale_, IN_SCALE, nddmaOuterItersScale_, k, dstShape);
147+ SetupNddmaParams(nddmaParamsVariance_, IN_VARIANCE, nddmaOuterItersVariance_, k, dstShape);
148+ }
149+ 
150+ __aicore__ inline void Process();
151+ 
152+private:
153+ template <typename U>
154+ __aicore__ inline void SetupNddmaParams(AscendC::MultiCopyParams<U, ND>& params, int inputIdx, int64_t& outerIters,
155+ int64_t k, const int64_t* dstShape)
156+ {
157+ int64_t inner = 1;
158+ int64_t nd = 0;
159+ for (int64_t d = RANK - 1; d >= k && nd < ND; d--) {
160+ params.loopInfo.loopSize[nd] = (d == k) ? 0 : dstShape[d];
161+ params.loopInfo.loopSrcStride[nd] = td_->input_strides[inputIdx][d];
162+ params.loopInfo.loopDstStride[nd] = inner;
163+ params.loopInfo.loopLpSize[nd] = 0;
164+ params.loopInfo.loopRpSize[nd] = 0;
165+ inner *= (d == k) ? td_->split.a_i : dstShape[d];
166+ nd++;
167+ }
168+ for (; nd < ND; nd++) {
169+ params.loopInfo.loopSize[nd] = 1;
170+ params.loopInfo.loopSrcStride[nd] = 0;
171+ params.loopInfo.loopDstStride[nd] = inner;
172+ params.loopInfo.loopLpSize[nd] = 0;
173+ params.loopInfo.loopRpSize[nd] = 0;
174+ }
175+ // outer loop 只覆盖 flat 层和 NDDMA 之间的 gap: d = k .. RANK-nddma_dims-1
176+ outerIters = 1;
177+ for (int64_t d = k; d < RANK - nddma_dims_; d++)
178+ outerIters *= (d == k) ? td_->split.a_i : dstShape[d];
179+ }
180+ 
181+ // NDDMA 多维搬入（随路广播）：参照 adam_apply_one_assign_kernel.h:359-395
182+ template <typename U>
183+ __aicore__ inline void CopyInBrc(const int64_t* coord, int inputIdx, int slot, int64_t a_i_seg,
184+ AscendC::GlobalTensor<U>& gmInput,
185+ const AscendC::MultiCopyParams<U, ND>& baseParams, int64_t outerIters)
186+ {
187+ int64_t k = td_->split.axis;
188+ int64_t off = CalcInputOffset(coord, td_->input_strides[inputIdx], RANK);
189+ const int64_t* dstShape = td_->max_bro_shape;
190+ 
191+ auto params = baseParams;
192+ int64_t k_nd = RANK - 1 - k;
193+ int64_t inner = 1;
194+ for (int64_t nd = 0; nd < ND; nd++) {
195+ if (nd == k_nd)
196+ params.loopInfo.loopSize[nd] = a_i_seg;
197+ params.loopInfo.loopDstStride[nd] = inner;
198+ inner *= params.loopInfo.loopSize[nd];
199+ }
200+ 
201+ static constexpr AscendC::NdDmaConfig cfg = {false, AscendC::NdDmaConfig::unsetPad,
202+ AscendC::NdDmaConfig::unsetPad, false};
203+ 
204+ if constexpr (RANK <= kMaxNdDmaDims) {
205+ AscendC::DataCopy<U, ND, cfg>(buf_[slot].template Get<U>(), gmInput[off], params);
206+ } else {
207+ AscendC::LocalTensor<U> buf = buf_[slot].template Get<U>();
208+ int64_t elem_base = off;
209+ for (int64_t oi = 0; oi < outerIters; oi++) {
210+ int64_t elem_adj = 0, tmp = oi;
211+ for (int64_t d = RANK - nddma_dims_ - 1; d >= k; d--) {
212+ int64_t sz = (d == k) ? a_i_seg : dstShape[d];
213+ elem_adj += (tmp % sz) * td_->input_strides[inputIdx][d];
214+ tmp /= sz;
215+ }
216+ AscendC::DataCopy<U, ND, cfg>(buf[oi * inner], gmInput[elem_base + elem_adj], params);
217+ }
218+ }
219+ }
220+ 
221+ // UB→GM 搬出，用 DataCopyPad（不要求 32B 对齐）：参照 adam_apply_one_assign_kernel.h:397-408
222+ __aicore__ inline void CopyOutOne(const int64_t* coord, int outputIdx, int slot, int64_t a_i_seg)
223+ {
224+ int64_t off = CalcInputOffset(coord, td_->output_strides[outputIdx], RANK);
225+ int64_t cnt = CalcOutputTransferCount(td_->output_shapes[outputIdx], RANK, td_->split.axis, a_i_seg);
226+ AscendC::DataCopyExtParams extParams;
227+ extParams.blockCount = 1;
228+ extParams.blockLen = cnt * sizeof(T);
229+ extParams.srcStride = 0;
230+ extParams.dstStride = 0;
231+ AscendC::DataCopyPad(gmOut_[off], buf_[slot].Get<T>(), extParams);
232+ }
233+};
234+ 
235+// ============================================================
236+// Process — 多 dtype 计算链（Broadcast 范式，逐 tile）
237+// ============================================================
238+template <typename T, int64_t RANK>
239+__aicore__ inline void BNInferGradKernel<T, RANK>::Process()
240+{
241+ int32_t evMTE2toV = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE2_V));
242+ int32_t evVtoMTE3 = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::V_MTE3));
243+ int32_t evMTE3toMTE2 = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE3_MTE2));
244+ 
245+ int64_t start, end;
246+ GetCoreRange(AscendC::GetBlockIdx(), td_->multicore.tiles_main, td_->multicore.cores_tail, start, end);
247+ 
248+ // 空核范围早退（DESIGN §5.4.1 空 tensor 路径 / 多核数 > tile 数时的空闲核）：
249+ // 本核无 tile 时直接返回，不做任何 DataCopy / 计算，x_backprop 保持空。
250+ if (start >= end) {
251+ return;
252+ }
253+ 
254+ // NDDMA DataCache is per core. Invalidate it once before this core's first
255+ // multi-dimensional DataCopy; broadcast inputs make multiple cores read
256+ // the same scale/variance GM addresses.
257+ AscendC::NdDmaDci();
258+ 
259+ // 物理槽位分配（P=5，所有 dtype 路径）
260+ constexpr int UB0 = 0, UB1 = 1, UB2 = 2, UB3 = 3, UB4 = 4;
261+ 
262+ const float epsilon = td_->epsilon_value;
263+ 
264+ int64_t inner_count = 1;
265+ for (int64_t d = td_->split.axis + 1; d < RANK; d++)
266+ inner_count *= td_->max_bro_shape[d];
267+ 
268+ int64_t coord[8] = {};
269+ for (int64_t flat = start; flat < end; flat++) {
270+ int64_t a_i_seg = GetUBSplitRange(flat % td_->split.a_o, td_->split.a_o, td_->split.a_i, td_->split.a_i_tail);
271+ int64_t count = a_i_seg * inner_count;
272+ FlatToEffectiveCoord(flat, td_->max_bro_shape, RANK, td_->split.axis, td_->split.a_i, td_->split.a_o, coord);
273+ 
274+ // 上轮 CopyOut(MTE3) 结束 → 本轮 CopyIn(MTE2) 可以开始
275+ if (flat != start)
276+ AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(evMTE3toMTE2);
277+ 
278+ if constexpr (std::is_same_v<T, float>) {
279+ // ============================================================
280+ // FP32 路径（直接计算，无 Cast）
281+ // ============================================================
282+ // CopyIn: grads→UB0, variance→UB1, scale→UB2 (NDDMA 随路广播 scale/variance)
283+ CopyInBrc(coord, IN_GRADS, UB0, a_i_seg, gmGrads_, nddmaParamsGrads_, nddmaOuterItersGrads_);
284+ CopyInBrc(coord, IN_VARIANCE, UB1, a_i_seg, gmVariance_, nddmaParamsVariance_, nddmaOuterItersVariance_);
285+ CopyInBrc(coord, IN_SCALE, UB2, a_i_seg, gmScale_, nddmaParamsScale_, nddmaOuterItersScale_);
286+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
287+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
288+ 
289+ // ① var_eps = batch_variance + epsilon
290+ AscendC::Adds(buf_[UB3].Get<float>(), buf_[UB1].Get<float>(), epsilon, count);
291+ // ② var_sqrt = sqrt(var_eps)
292+ AscendC::Sqrt(buf_[UB3].Get<float>(), buf_[UB3].Get<float>(), count);
293+ // ③ rsqrt = 1.0 / var_sqrt
294+ AscendC::Duplicate(buf_[UB1].Get<float>(), 1.0f, count);
295+ AscendC::Div(buf_[UB4].Get<float>(), buf_[UB1].Get<float>(), buf_[UB3].Get<float>(), count);
296+ // ④ scale_rsqrt = scale * rsqrt
297+ AscendC::Mul(buf_[UB3].Get<float>(), buf_[UB2].Get<float>(), buf_[UB4].Get<float>(), count);
298+ // ⑤ x_backprop = grads * scale_rsqrt
299+ AscendC::Mul(buf_[UB4].Get<float>(), buf_[UB0].Get<float>(), buf_[UB3].Get<float>(), count);
300+ 
301+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
302+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
303+ CopyOutOne(coord, OUT_XBP, UB4, a_i_seg);
304+ 
305+ } else if constexpr (std::is_same_v<T, half> || std::is_same_v<T, bfloat16_t>) {
306+ // ============================================================
307+ // FP16/BF16 路径（Cast → FP32 计算 → Cast）
308+ // ============================================================
309+ // CopyIn: grads(fp16/bf16)→UB0, variance(fp32)→UB2, scale(fp32)→UB3
310+ CopyInBrc(coord, IN_GRADS, UB0, a_i_seg, gmGrads_, nddmaParamsGrads_, nddmaOuterItersGrads_);
311+ CopyInBrc(coord, IN_VARIANCE, UB2, a_i_seg, gmVariance_, nddmaParamsVariance_, nddmaOuterItersVariance_);
312+ CopyInBrc(coord, IN_SCALE, UB3, a_i_seg, gmScale_, nddmaParamsScale_, nddmaOuterItersScale_);
313+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
314+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
315+ 
316+ // ① Cast grads(fp16/bf16) → grads_fp32(UB1)
317+ AscendC::Cast(buf_[UB1].Get<float>(), buf_[UB0].Get<T>(), AscendC::RoundMode::CAST_NONE, count);
318+ 
319+ // ② var_eps = batch_variance + epsilon (fp32)
320+ AscendC::Adds(buf_[UB4].Get<float>(), buf_[UB2].Get<float>(), epsilon, count);
321+ // ③ var_sqrt = sqrt(var_eps)
322+ AscendC::Sqrt(buf_[UB4].Get<float>(), buf_[UB4].Get<float>(), count);
323+ // ④ rsqrt = 1.0 / var_sqrt (复用 UB2 为 ones)
324+ AscendC::Duplicate(buf_[UB2].Get<float>(), 1.0f, count);
325+ AscendC::Div(buf_[UB0].Get<float>(), buf_[UB2].Get<float>(), buf_[UB4].Get<float>(), count);
326+ // ⑤ scale_rsqrt = scale * rsqrt (复用 UB4)
327+ AscendC::Mul(buf_[UB4].Get<float>(), buf_[UB3].Get<float>(), buf_[UB0].Get<float>(), count);
328+ // ⑥ x_backprop_fp32 = grads_fp32 * scale_rsqrt (复用 UB2)
329+ AscendC::Mul(buf_[UB2].Get<float>(), buf_[UB1].Get<float>(), buf_[UB4].Get<float>(), count);
330+ 
331+ // ⑦ Cast x_backprop_fp32(UB2) → x_backprop(fp16/bf16, UB0)
332+ AscendC::Cast(buf_[UB0].Get<T>(), buf_[UB2].Get<float>(), AscendC::RoundMode::CAST_ROUND, count);
333+ 
334+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
335+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
336+ CopyOutOne(coord, OUT_XBP, UB0, a_i_seg);
337+ }
338+ 
339+ if (flat != end - 1)
340+ AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(evMTE3toMTE2);
341+ }
342+}
Anorm/bn_infer_grad/op_kernel/arch35/bn_infer_grad_tiling_data.h+66-0
@@ -0,0 +1,66 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot - Iter1 */
12+ 
13+/*!
14+ * \file bn_infer_grad_tiling_data.h
15+ * \brief BNInferGrad TilingData 结构体（arch35 架构，按 rank 模板化）
16+ *
17+ * 迭代一：TilingKey_RANK4 + float32 路径
18+ *
19+ * 结构对齐 Broadcast 范式参考算子 adam_apply_one_assign：
20+ * - split/multicore 由 Host 侧 FindSplitAxis / MultiCoreSplit 填充
21+ * - max_bro_shape / input_shapes / input_strides / output_* 为 NDDMA 坐标系
22+ * - epsilon_value 为算子特有标量（来自 OpDef Attr）
23+ */
24+ 
25+#ifndef BN_INFER_GRAD_TILING_DATA_H_
26+#define BN_INFER_GRAD_TILING_DATA_H_
27+ 
28+#include <cstdint>
29+ 
30+// === 算子特定常量 ===
31+// grads, scale, batch_variance 三个 GM 输入（epsilon 为 Attr 标量，不占 GM slot）
32+constexpr int64_t kMaxInputSlots = 3; // 本算子最大输入张量数
33+constexpr int64_t kMaxOutputSlots = 1; // 本算子最大输出张量数
34+constexpr int64_t kPhysNodes = 5; // 物理存活节点 P（= TBuf 槽位数，float32 路径）
35+ 
36+struct SplitResult {
37+ int64_t axis; // UB 切分轴
38+ int64_t a_i; // 内轴 tile 大小（元素数）
Y
Yyulianjie7月29日

变量需要使用小驼峰命名

likedislike
39+ int64_t a_o; // 外轴 tile 数
40+ int64_t a_i_tail; // 末块大小（元素数）
41+};
42+ 
43+struct MultiCoreResult {
44+ int64_t num_cores; // 参与计算的核数
45+ int64_t total_tiles; // tile 总数
46+ int64_t tiles_main; // 每核主 tile 数
47+ int64_t cores_tail; // 多处理一个 tile 的核数
48+};
49+ 
50+template <int64_t kRank>
51+struct BNInferGradTilingData {
52+ SplitResult split; // UB 切分结果（来源：FindSplitAxis）
53+ MultiCoreResult multicore; // 多核切分结果（来源：MultiCoreSplit）
54+ int64_t rank; // 实际有效 rank（去 1 补 1 后）
55+ int64_t per_buf_bytes; // 单 buffer 字节数 = (UB/P) & ~31
56+ int64_t max_bro_shape[kRank]; // 广播后各维大小（坐标系）
57+ int64_t num_inputs; // 输入张量数（3）
58+ int64_t num_outputs; // 输出张量数（1）
59+ int64_t input_shapes[kMaxInputSlots][kRank]; // 各输入补 1 后的 shape
60+ int64_t input_strides[kMaxInputSlots][kRank]; // 各输入的 GM stride（broadcast 轴 = 0）
61+ int64_t output_shapes[kMaxOutputSlots][kRank]; // 各输出补 1 后的 shape
62+ int64_t output_strides[kMaxOutputSlots][kRank]; // 各输出的 GM stride
63+ float epsilon_value; // epsilon 标量值（来自 OpDef Attr）
64+};
65+ 
66+#endif // BN_INFER_GRAD_TILING_DATA_H_
Anorm/bn_infer_grad/op_kernel/arch35/bn_infer_grad_tiling_key.h+36-0
@@ -0,0 +1,36 @@
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+/* Generated By CANNBot - Iter1 */
12+ 
13+/*!
14+ * \file bn_infer_grad_tiling_key.h
15+ * \brief BNInferGrad TilingKey 模板参数定义（arch35 架构）
16+ *
17+ * Broadcast 范式：TilingKey 只与 RANK 有关，与数据类型无关。
18+ * - RANK_4: rank <= 4（NC / NCH / NCHW），常态场景
19+ * - RANK_8: rank == 5（NCDHW），保底扩展场景
20+ */
21+ 
22+#ifndef BN_INFER_GRAD_TILING_KEY_H_
23+#define BN_INFER_GRAD_TILING_KEY_H_
24+ 
25+#include "ascendc/host_api/tiling/template_argument.h"
26+ 
27+#define BN_INFER_GRAD_RANK_4 4
28+#define BN_INFER_GRAD_RANK_8 8
29+ 
30+ASCENDC_TPL_ARGS_DECL(BNInferGrad,
31+ ASCENDC_TPL_UINT_DECL(RANK, 8, ASCENDC_TPL_UI_LIST, BN_INFER_GRAD_RANK_4, BN_INFER_GRAD_RANK_8));
32+ 
33+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(RANK, ASCENDC_TPL_UI_LIST, BN_INFER_GRAD_RANK_4)),
34+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(RANK, ASCENDC_TPL_UI_LIST, BN_INFER_GRAD_RANK_8)));
35+ 
36+#endif // BN_INFER_GRAD_TILING_KEY_H_
Anorm/bn_infer_grad/op_kernel/bn_infer_grad.cpp+54-0
@@ -0,0 +1,54 @@
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+/* Generated By CANNBot - Iter1 */
12+ 
13+/*!
14+ * \file bn_infer_grad.cpp
15+ * \brief BNInferGrad 算子 Kernel 入口（arch35 架构，RegBase / Broadcast 范式）
16+ *
17+ * 一芯片一算子一入口：arch35 唯一 __global__ 入口。
18+ *
19+ * Key 传递: TilingKey(RANK) → bn_infer_grad_tiling_key.h(ASCENDC_TPL RANK)
20+ * → bn_infer_grad.cpp(template<int RANK>) → TilingData<kRank> + Kernel<DTYPE, RANK>
21+ *
22+ * DTYPE 来自第一个输入 grads（spec dtype_policy=same_as_first_input），
23+ * 构建系统按输入名注入编译宏 DTYPE_GRADS。
24+ * epsilon 为 OpDef Attr，经 Tiling 写入 TilingData.epsilon_value，不占 GM 入参。
25+ *
26+ * 核函数参数顺序契约：输入(grads, scale, batch_variance) → 输出(x_backprop) → workspace → tiling。
27+ */
28+#include "arch35/bn_infer_grad.h"
29+ 
30+using TilingData4 = BNInferGradTilingData<4>; // RANK<=4: 数组维度 [4]
31+using TilingData8 = BNInferGradTilingData<8>; // RANK==5: 数组维度 [8]
32+ 
33+template <int RANK>
34+__global__ __aicore__ void bn_infer_grad(GM_ADDR grads, GM_ADDR scale, GM_ADDR batch_variance, GM_ADDR x_backprop,
35+ GM_ADDR workspace, GM_ADDR tiling)
36+{
37+ GM_ADDR ins[kMaxInputSlots] = {grads, scale, batch_variance};
38+ GM_ADDR outs[kMaxOutputSlots] = {x_backprop};
39+ 
40+ REGISTER_NONE_TILING;
41+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
42+ 
43+ if constexpr (RANK == BN_INFER_GRAD_RANK_4) {
44+ GET_TILING_DATA_WITH_STRUCT(TilingData4, td, tiling);
45+ BNInferGradKernel<DTYPE_GRADS, 4> kernel;
46+ kernel.Init(ins, outs, &td);
47+ kernel.Process();
48+ } else {
49+ GET_TILING_DATA_WITH_STRUCT(TilingData8, td, tiling);
50+ BNInferGradKernel<DTYPE_GRADS, 8> kernel;
51+ kernel.Init(ins, outs, &td);
52+ kernel.Process();
53+ }
54+}
Anorm/bn_infer_grad/tests/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# -----------------------------------------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")
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()
Anorm/bn_infer_grad/tests/assets/golden.py+64-0
@@ -0,0 +1,64 @@
1+#!/usr/bin/env python3
2+# -*- coding: UTF-8 -*-
3+# -----------------------------------------------------------------------------------------------------------
4+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
5+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6+# CANN Open Software License Agreement Version 2.0 (the "License").
7+# Please refer to the License for details. You may not use this file except in compliance with the License.
8+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10+# See LICENSE in the root of the software repository for the full text of the License.
11+# -----------------------------------------------------------------------------------------------------------
12+ 
13+import numpy as np
14+ 
15+ 
16+__golden__ = {"kernel": {"bn_infer_grad": "bn_infer_grad_golden"}}
17+ 
18+ 
19+def bn_infer_grad_golden(grads, scale, batch_variance, *, epsilon=0.0001, **kwargs):
20+ """
21+ Compute BNInferGrad through the PyTorch batch-norm inference backward API.
22+ 
23+ In inference mode, native_batch_norm_backward calculates:
24+ grad_input = grad_out * weight / sqrt(running_var + epsilon)
25+ 
26+ BNInferGrad does not expose input or running_mean because they do not
27+ affect grad_input in this mode, so zero placeholders are supplied.
28+ """
29+ import torch
30+ 
31+ def to_torch(array):
32+ array = np.asarray(array)
33+ if array.dtype.name == "bfloat16":
34+ return torch.from_numpy(array.astype(np.float32).copy()).to(torch.bfloat16)
35+ return torch.from_numpy(array.copy())
36+ 
37+ grads_array = np.asarray(grads)
38+ grads_torch = to_torch(grads_array)
39+ scale_torch = to_torch(scale)
40+ variance_torch = to_torch(batch_variance)
41+ input_torch = torch.zeros_like(grads_torch)
42+ running_mean_torch = torch.zeros_like(scale_torch)
43+ 
44+ grad_input, _, _ = torch.ops.aten.native_batch_norm_backward(
45+ grads_torch,
46+ input_torch,
47+ scale_torch,
48+ running_mean_torch,
49+ variance_torch,
50+ None,
51+ None,
52+ False,
53+ float(epsilon),
54+ [True, False, False],
55+ )
56+ if grad_input.dtype == torch.bfloat16:
57+ return (
58+ grad_input.detach()
59+ .to(torch.float32)
60+ .cpu()
61+ .numpy()
62+ .astype(grads_array.dtype)
63+ )
64+ return grad_input.detach().cpu().numpy()
Anorm/bn_infer_grad/tests/ut/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# -----------------------------------------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")
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()
Anorm/bn_infer_grad/tests/ut/op_host/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
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+# BNInferGrad 的 CMakeLists（norm/bn_infer_grad/CMakeLists.txt）声明 ACLNNTYPE aclnn_exclude，
12+# 不产出 aclnn L0/L2 接口，故此处不注册 op_api UT，仅覆盖 tiling + infershape。
13+file(GLOB CURRENT_DIR RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
14+if(UT_TEST_ALL OR OP_HOST_UT)
15+ add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
16+ add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
17+endif()
Anorm/bn_infer_grad/tests/ut/op_host/arch35/test_bn_infer_grad_tiling.cpp+283-0
@@ -0,0 +1,283 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file test_bn_infer_grad_tiling.cpp
15+ * \brief BNInferGrad Tiling UT（主线仓级范式）
16+ *
17+ * 挂载：nn_op_tiling_ut 聚合体（add_modules_ut_sources HOSTNAME ${OP_TILING_MODULE_NAME}）。
18+ * 框架：仓级 gert::TilingContextPara + ExecuteTiling（tests/ut/common/tiling_case_executor.*）。
19+ *
20+ * 覆盖矩阵：两 TilingKey（RANK_4 / RANK_8）× 三 dtype（fp16/fp32/bf16）
21+ * + scale/batch_variance 沿 C 广播
22+ * + 边界：空 tensor（某维 0）、单核小 shape、多核大 shape
23+ * + 负例：grads rank<2 / rank>5、scale/variance 非 rank1、C 长度不匹配。
24+ *
25+ * TilingKey 由 ASCENDC_TPL_SEL_PARAM(RANK) 生成，数值编码不硬编码；本 UT 通过
26+ * "同组一致、跨组相异" 的关系断言验证两档 RANK Key 正确分流：
27+ * - grads rank 2~4（有效 rank<=4）→ RANK_4
28+ * - grads rank 5（有效 rank 5，映射 R=8）→ RANK_8
29+ *
30+ * 注：BNInferGrad tiling 直接从 ctx->GetCompileInfo() 读取 coreNum/ubSize，
31+ * UT 不跑 TilingParse，故构造 BNInferGradCompileInfo 通过 CompileInfo() 传入。
32+ */
33+ 
34+#include <iostream>
35+#include <vector>
36+#include <gtest/gtest.h>
37+#include "tiling_context_faker.h"
38+#include "tiling_case_executor.h"
39+// CompileInfo 结构与 optiling 类声明（相对交付源 op_host/arch35）
40+#include "../../../../op_host/arch35/bn_infer_grad_tiling_arch35.h"
41+ 
42+namespace BNInferGradTilingUT {
43+using namespace std;
44+using namespace ge;
45+ 
46+static const std::string OP_NAME = "BNInferGrad";
47+ 
48+// DAV_3510：UB 248KB，核数取 32（值不影响 TilingKey 分流，仅影响切分粒度）
49+static constexpr uint64_t kUbSize = 248UL * 1024UL;
50+static constexpr uint64_t kCoreNum = 32UL;
51+static constexpr uint64_t kTilingDataSize = 4096UL;
52+static constexpr float kEpsilon = 0.0001f;
53+ 
54+// 由运行期 vector 构造 gert::StorageShape（rank=0 时为空 shape）
55+static gert::StorageShape MakeStorageShape(const std::vector<int64_t>& dims)
56+{
57+ gert::StorageShape ss;
58+ for (auto d : dims) {
59+ ss.MutableOriginShape().AppendDim(d);
60+ ss.MutableStorageShape().AppendDim(d);
61+ }
62+ return ss;
63+}
64+ 
65+// grads [N,C,...]（dtype 可变）；scale/batch_variance 固定 FP32、shape=[C]。
66+// scaleLen/varLen 默认取 grads 的 C（axis=1）；负例可显式改写以触发校验。
67+static bool RunOneCase(const std::vector<int64_t>& gradsShape, ge::DataType gradsDtype, ge::graphStatus expectStatus,
68+ TilingInfo& info, int64_t scaleLenOverride = -1, int64_t varLenOverride = -1,
69+ int64_t scaleRankOverride = 1, int64_t varRankOverride = 1)
70+{
71+ int64_t channels = (gradsShape.size() >= 2) ? gradsShape[1] : 1;
72+ int64_t scaleLen = (scaleLenOverride >= 0) ? scaleLenOverride : channels;
73+ int64_t varLen = (varLenOverride >= 0) ? varLenOverride : channels;
74+ 
75+ // scale / batch_variance 的 shape：默认 rank1 [C]；负例可构造 rank2 触发 rank 校验。
76+ std::vector<int64_t> scaleDims = (scaleRankOverride == 1) ? std::vector<int64_t>{scaleLen} :
77+ std::vector<int64_t>{scaleLen, 1};
78+ std::vector<int64_t> varDims = (varRankOverride == 1) ? std::vector<int64_t>{varLen} :
79+ std::vector<int64_t>{varLen, 1};
80+ 
81+ gert::StorageShape gShape = MakeStorageShape(gradsShape);
82+ gert::StorageShape sShape = MakeStorageShape(scaleDims);
83+ gert::StorageShape vShape = MakeStorageShape(varDims);
84+ 
85+ // 3 输入：grads(dtype), scale(fp32), batch_variance(fp32)
86+ std::vector<gert::TilingContextPara::TensorDescription> inputs = {
87+ {gShape, gradsDtype, ge::FORMAT_ND}, // grads
88+ {sShape, ge::DT_FLOAT, ge::FORMAT_ND}, // scale
89+ {vShape, ge::DT_FLOAT, ge::FORMAT_ND}, // batch_variance
90+ };
91+ // 1 输出：x_backprop（与 grads 同 shape/同 dtype）
92+ std::vector<gert::TilingContextPara::TensorDescription> outputs = {
93+ {gShape, gradsDtype, ge::FORMAT_ND},
94+ };
95+ // 属性：epsilon(0)
96+ std::vector<gert::TilingContextPara::OpAttr> attrs = {
97+ {"epsilon", Ops::NN::AnyValue::CreateFrom<float>(kEpsilon)},
98+ };
99+ 
100+ // GetShapeInfo/DoTilingAndSet 直接读取该 compileInfo（UT 不跑 TilingParse）
101+ optiling::BNInferGradCompileInfo compileInfo{kCoreNum, kUbSize};
102+ 
103+ gert::TilingContextPara para(OP_NAME, inputs, outputs, attrs, &compileInfo, kCoreNum, kUbSize, kTilingDataSize);
104+ 
105+ bool ok = ExecuteTiling(para, info);
106+ if (expectStatus == ge::GRAPH_SUCCESS) {
107+ EXPECT_TRUE(ok) << "tiling should succeed for the given case";
108+ } else {
109+ EXPECT_FALSE(ok) << "tiling should fail for the given case";
110+ }
111+ return ok;
112+}
113+ 
114+class BNInferGradTilingTest : public testing::Test {
115+protected:
116+ static void SetUpTestCase() { std::cout << "BNInferGradTilingTest SetUp." << std::endl; }
117+ static void TearDownTestCase() { std::cout << "BNInferGradTilingTest TearDown." << std::endl; }
118+};
119+ 
120+// ---------------------------------------------------------------------------
121+// RANK_4 组：grads rank 2~4，三 dtype
122+// ---------------------------------------------------------------------------
123+TEST_F(BNInferGradTilingTest, Rank4_Fp32_2D)
124+{
125+ TilingInfo info;
126+ ASSERT_TRUE(RunOneCase({4, 8}, ge::DT_FLOAT, ge::GRAPH_SUCCESS, info));
127+ EXPECT_GE(info.blockNum, 1u);
128+ EXPECT_NE(info.tilingKey, -1);
129+}
130+ 
131+TEST_F(BNInferGradTilingTest, Rank4_Fp16_3D)
132+{
133+ TilingInfo info;
134+ ASSERT_TRUE(RunOneCase({2, 16, 5}, ge::DT_FLOAT16, ge::GRAPH_SUCCESS, info));
135+ EXPECT_GE(info.blockNum, 1u);
136+}
137+ 
138+TEST_F(BNInferGradTilingTest, Rank4_Bf16_4D)
139+{
140+ TilingInfo info;
141+ ASSERT_TRUE(RunOneCase({4, 32, 14, 14}, ge::DT_BF16, ge::GRAPH_SUCCESS, info));
142+ EXPECT_GE(info.blockNum, 1u);
143+}
144+ 
145+// ---------------------------------------------------------------------------
146+// RANK_8 组：grads rank 5（有效 rank 5 → 映射 R=8）
147+// ---------------------------------------------------------------------------
148+TEST_F(BNInferGradTilingTest, Rank8_Fp32_5D)
149+{
150+ TilingInfo info;
151+ ASSERT_TRUE(RunOneCase({2, 8, 3, 5, 7}, ge::DT_FLOAT, ge::GRAPH_SUCCESS, info));
152+ EXPECT_GE(info.blockNum, 1u);
153+}
154+ 
155+TEST_F(BNInferGradTilingTest, Rank8_Fp16_5D)
156+{
157+ TilingInfo info;
158+ ASSERT_TRUE(RunOneCase({2, 16, 2, 4, 5}, ge::DT_FLOAT16, ge::GRAPH_SUCCESS, info));
159+ EXPECT_GE(info.blockNum, 1u);
160+}
161+ 
162+TEST_F(BNInferGradTilingTest, Rank8_Bf16_5D)
163+{
164+ TilingInfo info;
165+ ASSERT_TRUE(RunOneCase({3, 8, 2, 3, 4}, ge::DT_BF16, ge::GRAPH_SUCCESS, info));
166+ EXPECT_GE(info.blockNum, 1u);
167+}
168+ 
169+// ---------------------------------------------------------------------------
170+// TilingKey 分流关系：同组一致、跨组相异
171+// ---------------------------------------------------------------------------
172+TEST_F(BNInferGradTilingTest, TilingKey_RankGroupDispatch)
173+{
174+ TilingInfo r4a, r4b, r8a, r8b;
175+ ASSERT_TRUE(RunOneCase({4, 8}, ge::DT_FLOAT, ge::GRAPH_SUCCESS, r4a)); // rank2 -> RANK_4
176+ ASSERT_TRUE(RunOneCase({4, 32, 14, 14}, ge::DT_BF16, ge::GRAPH_SUCCESS, r4b)); // rank4 -> RANK_4
177+ ASSERT_TRUE(RunOneCase({2, 8, 3, 5, 7}, ge::DT_FLOAT, ge::GRAPH_SUCCESS, r8a)); // rank5 -> RANK_8
178+ ASSERT_TRUE(RunOneCase({2, 16, 2, 4, 5}, ge::DT_FLOAT16, ge::GRAPH_SUCCESS, r8b)); // rank5 -> RANK_8
179+ 
180+ EXPECT_EQ(r4a.tilingKey, r4b.tilingKey) << "grads rank<=4 用例应命中同一 RANK_4 Key";
181+ EXPECT_EQ(r8a.tilingKey, r8b.tilingKey) << "grads rank5 用例应命中同一 RANK_8 Key";
182+ EXPECT_NE(r4a.tilingKey, r8a.tilingKey) << "RANK_4 与 RANK_8 应为不同 TilingKey";
183+}
184+ 
185+// ---------------------------------------------------------------------------
186+// 广播：scale/batch_variance [C] 沿非 C 轴广播到 grads
187+// ---------------------------------------------------------------------------
188+TEST_F(BNInferGradTilingTest, Broadcast_ChannelScaleVariance)
189+{
190+ TilingInfo info;
191+ // grads [8,4,6,6]，C=4；scale/variance=[4] 广播。tiling 内部 reshape 为 [1,4,1,1]。
192+ ASSERT_TRUE(RunOneCase({8, 4, 6, 6}, ge::DT_FLOAT, ge::GRAPH_SUCCESS, info));
193+ EXPECT_GE(info.blockNum, 1u);
194+}
195+ 
196+// ---------------------------------------------------------------------------
197+// 边界：空 tensor（某维为 0）→ is_empty，total_tiles=0，block_dim 保底 1，SUCCESS
198+// ---------------------------------------------------------------------------
199+TEST_F(BNInferGradTilingTest, Boundary_EmptyTensor_N0)
200+{
201+ TilingInfo info;
202+ ASSERT_TRUE(RunOneCase({0, 8}, ge::DT_FLOAT, ge::GRAPH_SUCCESS, info));
203+ EXPECT_GE(info.blockNum, 1u); // launch 要求 block_dim >= 1
204+}
205+ 
206+TEST_F(BNInferGradTilingTest, Boundary_EmptyTensor_Inner0)
207+{
208+ TilingInfo info;
209+ ASSERT_TRUE(RunOneCase({2, 8, 0, 5}, ge::DT_FLOAT16, ge::GRAPH_SUCCESS, info));
210+ EXPECT_GE(info.blockNum, 1u);
211+}
212+ 
213+// ---------------------------------------------------------------------------
214+// 大 shape：多核切分（total_tiles > 核数）
215+// ---------------------------------------------------------------------------
216+TEST_F(BNInferGradTilingTest, LargeShape_MultiCore)
217+{
218+ TilingInfo info;
219+ ASSERT_TRUE(RunOneCase({256, 64, 28, 28}, ge::DT_FLOAT, ge::GRAPH_SUCCESS, info));
220+ EXPECT_GE(info.blockNum, 1u);
221+}
222+ 
223+// ---------------------------------------------------------------------------
224+// 负例：grads rank<2 应被拒绝（GetShapeInfo grads_rank<2 校验）
225+// ---------------------------------------------------------------------------
226+TEST_F(BNInferGradTilingTest, Reject_GradsRankTooSmall)
227+{
228+ TilingInfo info;
229+ ASSERT_FALSE(RunOneCase({16}, ge::DT_FLOAT, ge::GRAPH_FAILED, info,
230+ /*scaleLen*/ 16, /*varLen*/ 16));
231+}
232+ 
233+// ---------------------------------------------------------------------------
234+// 负例：grads rank>5 应被拒绝（GetShapeInfo grads_rank>5 校验）
235+// ---------------------------------------------------------------------------
236+TEST_F(BNInferGradTilingTest, Reject_GradsRankTooLarge)
237+{
238+ TilingInfo info;
239+ ASSERT_FALSE(RunOneCase({2, 4, 2, 3, 2, 2}, ge::DT_FLOAT, ge::GRAPH_FAILED, info));
240+}
241+ 
242+// ---------------------------------------------------------------------------
243+// 负例：scale 非 rank1（rank2）应被拒绝
244+// ---------------------------------------------------------------------------
245+TEST_F(BNInferGradTilingTest, Reject_ScaleNotRank1)
246+{
247+ TilingInfo info;
248+ ASSERT_FALSE(RunOneCase({4, 8}, ge::DT_FLOAT, ge::GRAPH_FAILED, info,
249+ /*scaleLen*/ 8, /*varLen*/ 8, /*scaleRank*/ 2, /*varRank*/ 1));
250+}
251+ 
252+// ---------------------------------------------------------------------------
253+// 负例：batch_variance 非 rank1（rank2）应被拒绝
254+// ---------------------------------------------------------------------------
255+TEST_F(BNInferGradTilingTest, Reject_VarianceNotRank1)
256+{
257+ TilingInfo info;
258+ ASSERT_FALSE(RunOneCase({4, 8}, ge::DT_FLOAT, ge::GRAPH_FAILED, info,
259+ /*scaleLen*/ 8, /*varLen*/ 8, /*scaleRank*/ 1, /*varRank*/ 2));
260+}
261+ 
262+// ---------------------------------------------------------------------------
263+// 负例：scale 长度 != grads C 应被拒绝
264+// ---------------------------------------------------------------------------
265+TEST_F(BNInferGradTilingTest, Reject_ScaleChannelMismatch)
266+{
267+ TilingInfo info;
268+ // grads C=8，但 scale 长度给 7 → 不匹配
269+ ASSERT_FALSE(RunOneCase({4, 8}, ge::DT_FLOAT, ge::GRAPH_FAILED, info,
270+ /*scaleLen*/ 7, /*varLen*/ 8));
271+}
272+ 
273+// ---------------------------------------------------------------------------
274+// 负例：batch_variance 长度 != grads C 应被拒绝
275+// ---------------------------------------------------------------------------
276+TEST_F(BNInferGradTilingTest, Reject_VarianceChannelMismatch)
277+{
278+ TilingInfo info;
279+ ASSERT_FALSE(RunOneCase({4, 8}, ge::DT_FLOAT, ge::GRAPH_FAILED, info,
280+ /*scaleLen*/ 8, /*varLen*/ 5));
281+}
282+ 
283+} // namespace BNInferGradTilingUT
Anorm/bn_infer_grad/tests/ut/op_host/test_bn_infer_grad_infershape.cpp+75-0
@@ -0,0 +1,75 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file test_bn_infer_grad_infershape.cpp
15+ * \brief BNInferGrad InferShape UT
16+ *
17+ * 契约（bn_infer_grad_infershape.cpp）：
18+ * - x_backprop.shape = grads.shape（scale/batch_variance 沿 C 广播，不影响输出 shape）
19+ */
20+#include <gtest/gtest.h>
21+#include "infershape_test_util.h"
22+#include "ut_op_common.h"
23+#include "../../../op_graph/bn_infer_grad_proto.h"
24+ 
25+class BNInferGradInferShapeTest : public testing::Test {
26+protected:
27+ static void SetUpTestCase() { std::cout << "BNInferGrad Proto Test SetUp" << std::endl; }
28+ static void TearDownTestCase() { std::cout << "BNInferGrad Proto Test TearDown" << std::endl; }
29+};
30+ 
31+// ============================================================
32+// InferShape: x_backprop.shape 完整拷贝 grads.shape
33+// ============================================================
34+ 
35+// 4D grads（NCHW 典型），scale/variance 为 [C]
36+TEST_F(BNInferGradInferShapeTest, BNInferGrad_infershape_4d)
37+{
38+ ge::op::BNInferGrad op;
39+ op.UpdateInputDesc("grads", create_desc({4, 32, 14, 14}, ge::DT_FLOAT));
40+ op.UpdateInputDesc("scale", create_desc({32}, ge::DT_FLOAT));
41+ op.UpdateInputDesc("batch_variance", create_desc({32}, ge::DT_FLOAT));
42+ 
43+ EXPECT_EQ(InferShapeTest(op), ge::GRAPH_SUCCESS);
44+ 
45+ std::vector<int64_t> expected = {4, 32, 14, 14};
46+ EXPECT_EQ(op.GetOutputDesc(0).GetShape().GetDims(), expected);
47+}
48+ 
49+// 2D grads
50+TEST_F(BNInferGradInferShapeTest, BNInferGrad_infershape_2d)
51+{
52+ ge::op::BNInferGrad op;
53+ op.UpdateInputDesc("grads", create_desc({8, 16}, ge::DT_FLOAT16));
54+ op.UpdateInputDesc("scale", create_desc({16}, ge::DT_FLOAT));
55+ op.UpdateInputDesc("batch_variance", create_desc({16}, ge::DT_FLOAT));
56+ 
57+ EXPECT_EQ(InferShapeTest(op), ge::GRAPH_SUCCESS);
58+ 
59+ std::vector<int64_t> expected = {8, 16};
60+ EXPECT_EQ(op.GetOutputDesc(0).GetShape().GetDims(), expected);
61+}
62+ 
63+// 5D grads（rank 上界）
64+TEST_F(BNInferGradInferShapeTest, BNInferGrad_infershape_5d)
65+{
66+ ge::op::BNInferGrad op;
67+ op.UpdateInputDesc("grads", create_desc({2, 8, 3, 5, 7}, ge::DT_BF16));
68+ op.UpdateInputDesc("scale", create_desc({8}, ge::DT_FLOAT));
69+ op.UpdateInputDesc("batch_variance", create_desc({8}, ge::DT_FLOAT));
70+ 
71+ EXPECT_EQ(InferShapeTest(op), ge::GRAPH_SUCCESS);
72+ 
73+ std::vector<int64_t> expected = {2, 8, 3, 5, 7};
74+ EXPECT_EQ(op.GetOutputDesc(0).GetShape().GetDims(), expected);
75+}
Anorm/bn_infer_grad/tests/ut/op_kernel/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
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+# param1: 算子名（kernel 命名）
12+# param2: soc 版本（arch35 / RegBase → ascend950pr_9599）
13+# param3: 编译宏。grads 的 dtype 由构建系统按输入名注入 DTYPE_GRADS；kernel 入口按此实例化。
14+# 默认组合取 float（FP32 直算路径），FP16/BF16 路径在测试内部另行覆盖。
15+if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)
16+ AddOpTestCase(bn_infer_grad "ascend950pr_9599" "-DDTYPE_GRADS=float")
17+endif()
Anorm/bn_infer_grad/tests/ut/op_kernel/bn_infer_grad_tiling_def.h+35-0
@@ -0,0 +1,35 @@
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+ * Generated By CANNBot
11+ */
12+ 
13+// BNInferGrad 的 Kernel UT tiling-definition shim。
14+//
15+// op_kernel 入口用 REGISTER_NONE_TILING，构建系统的 tiling-header 生成器不识别该宏
16+// （它只处理 REGISTER_TILING_DEFAULT / BroadcastSch），因此不会自动生成
17+// GET_TILING_DATA_WITH_STRUCT。当此 <op>_tiling_def.h 存在时，ut.cmake（cmake/ut.cmake:574-575）
18+// 会用 `-include` 强制注入本头文件并跳过自动生成，故在此定义该宏。
19+//
20+// BNInferGradTilingData<R> 为 POD（全 int64_t/float + 定长数组），
21+// 从原始 tiling GM buffer 按字节拷贝即等价于生成器行为
22+// （省去仅内部使用、UT 不需要的 REGISTER_TILINGDATA_SIZE 记账）。
23+ 
24+#ifndef BN_INFER_GRAD_TILING_DEF_H_
25+#define BN_INFER_GRAD_TILING_DEF_H_
26+ 
27+#include <cstring>
28+ 
29+#ifndef GET_TILING_DATA_WITH_STRUCT
30+#define GET_TILING_DATA_WITH_STRUCT(tiling_struct, tiling_data, tiling_arg) \
31+ tiling_struct tiling_data; \
32+ std::memcpy(reinterpret_cast<void*>(&tiling_data), reinterpret_cast<const void*>(tiling_arg), sizeof(tiling_struct))
33+#endif
34+ 
35+#endif // BN_INFER_GRAD_TILING_DEF_H_
Anorm/bn_infer_grad/tests/ut/op_kernel/test_bn_infer_grad.cpp+284-0
@@ -0,0 +1,284 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file test_bn_infer_grad.cpp
15+ * \brief BNInferGrad Kernel UT（arch35 / RegBase / Broadcast，CPU 孪生仿真）
16+ *
17+ * 覆盖：FP32 直算路径（DTYPE_GRADS=float，见 op_kernel/CMakeLists.txt），RANK_4（2D）与 RANK_8（5D）两档。
18+ *
19+ * 入口 bn_infer_grad 为 `template<int RANK>`（非 extern "C"），DTYPE 由构建宏 DTYPE_GRADS 注入，
20+ * 故沿用仓内 elu_grad/leaky_relu_grad 的模板入口范式：直接 #include kernel .cpp，
21+ * 用 lambda 显式实例化 bn_infer_grad<RANK> 后交给 ICPU_RUN_KF。
22+ *
23+ * 计算契约：x_backprop[i] = grads[i] * scale[c] / sqrt(batch_variance[c] + epsilon)，
24+ * c 为元素 i 的通道号（NCHW，C=axis1）。golden 独立于 kernel 实现，
25+ * 直接按公式在 host 侧计算后逐元素比对（双千精度，abs<=1e-3）。
26+ *
27+ * TilingData 手工填充：仅取"整块单 tile + 单核"配置（split.axis=dim0、a_o=1、num_cores=1），
28+ * 由与 host tiling 一致的规则（PadAndSqueeze/FindSplitAxis/PrecomputeInputStrides）推导，
29+ * 见 op_host/arch35/bn_infer_grad_tiling_arch35.cpp。要求 dim0>=2 且 C(dim1)>=2、各维>=2（避免 squeeze）。
30+ */
31+#include <array>
32+#include <vector>
33+#include <iostream>
34+#include <string>
35+#include <cstdint>
36+#include <cmath>
37+#include "gtest/gtest.h"
38+#include "tikicpulib.h"
39+#include "data_utils.h"
40+ 
41+#include "../../../op_kernel/bn_infer_grad.cpp"
42+#include <cstdint>
43+ 
44+using namespace std;
45+ 
46+class bn_infer_grad_test : public testing::Test {
47+protected:
48+ static void SetUpTestCase() { cout << "bn_infer_grad_test SetUp\n" << endl; }
49+ static void TearDownTestCase() { cout << "bn_infer_grad_test TearDown\n" << endl; }
50+};
51+ 
52+// 整块单 tile + 单核的 TilingData 填充（split.axis=dim0, a_o=1, num_cores=1）。
53+// eff 为有效（去 size-1 后）shape，要求 eff[0]>=2、eff[1](=C)>=2、其余>=2。
54+template <int64_t R>
55+static void FillSingleTileTiling(BNInferGradTilingData<R>* td, const std::vector<int64_t>& eff, float eps)
56+{
57+ int64_t rank = static_cast<int64_t>(eff.size());
58+ int64_t delta = R - rank;
59+ 
60+ td->split.axis = 0 + delta; // dim0 作切分轴，右移 delta
61+ td->split.a_i = eff[0];
62+ td->split.a_o = 1;
63+ td->split.a_i_tail = eff[0];
64+ 
65+ td->multicore.num_cores = 1;
66+ td->multicore.total_tiles = 1;
67+ td->multicore.tiles_main = 1;
68+ td->multicore.cores_tail = 0;
69+ 
70+ td->rank = rank;
71+ td->per_buf_bytes = 32768; // 32KB/slot，5 槽 <UB；单 tile 元素数远小于此
72+ td->num_inputs = 3;
73+ td->num_outputs = 1;
74+ 
75+ for (int64_t d = 0; d < delta; d++)
76+ td->max_bro_shape[d] = 1;
77+ for (int64_t d = 0; d < rank; d++)
78+ td->max_bro_shape[d + delta] = eff[d];
79+ 
80+ // grads(slot0) / output: 完整 shape + 行主序 stride（size-1 轴 stride=0）
81+ auto fillFull = [&](int64_t* shape, int64_t* strides) {
82+ for (int64_t d = 0; d < delta; d++) {
83+ shape[d] = 1;
84+ strides[d] = 0;
85+ }
86+ for (int64_t d = 0; d < rank; d++)
87+ shape[d + delta] = eff[d];
88+ int64_t prod = 1;
89+ for (int64_t d = rank - 1; d >= 0; d--) {
90+ strides[d + delta] = (eff[d] == 1) ? 0 : prod;
91+ prod *= eff[d];
92+ }
93+ };
94+ fillFull(td->input_shapes[0], td->input_strides[0]);
95+ fillFull(td->output_shapes[0], td->output_strides[0]);
96+ 
97+ // scale(slot1) / variance(slot2): reshape 为 [1,C,1,...]，仅 C 轴 stride=1，其余广播(=0)
98+ for (int i = 1; i <= 2; i++) {
99+ for (int64_t d = 0; d < R; d++) {
100+ td->input_shapes[i][d] = 1;
101+ td->input_strides[i][d] = 0;
102+ }
103+ td->input_shapes[i][1 + delta] = eff[1]; // C 落在 axis1（右移后 1+delta）
104+ td->input_strides[i][1 + delta] = 1;
105+ }
106+ 
107+ td->epsilon_value = eps;
108+}
109+ 
110+// 独立 golden + 逐元素比对：x_backprop[i] = grads[i]*scale[c]/sqrt(var[c]+eps)
111+static void CheckGolden(const std::vector<int64_t>& eff, const float* grads, const float* scale, const float* var,
112+ const float* out, float eps)
113+{
114+ int64_t total = 1;
115+ for (auto d : eff)
116+ total *= d;
117+ int64_t C = eff[1];
118+ int64_t innerAfterC = 1;
119+ for (size_t d = 2; d < eff.size(); d++)
120+ innerAfterC *= eff[d];
121+ 
122+ for (int64_t i = 0; i < total; i++) {
123+ int64_t c = (i / innerAfterC) % C;
124+ float golden = grads[i] * scale[c] / std::sqrt(var[c] + eps);
125+ EXPECT_NEAR(out[i], golden, 1e-3f) << "mismatch at flat index " << i << " (c=" << c << ")";
126+ }
127+}
128+ 
129+// RANK_4 路径（grads rank<=4），FP32，[N,C] 二维整块。
130+TEST_F(bn_infer_grad_test, fp32_rank4_2d_golden)
131+{
132+ std::vector<int64_t> eff = {2, 3}; // N=2, C=3
133+ const float eps = 0.0001f;
134+ int64_t total = eff[0] * eff[1];
135+ int64_t C = eff[1];
136+ 
137+ size_t gradsBytes = total * sizeof(float);
138+ size_t chBytes = C * sizeof(float);
139+ size_t outBytes = total * sizeof(float);
140+ size_t tilingBytes = sizeof(BNInferGradTilingData<4>);
141+ 
142+ uint8_t* grads = (uint8_t*)AscendC::GmAlloc(gradsBytes);
143+ uint8_t* scale = (uint8_t*)AscendC::GmAlloc(chBytes);
144+ uint8_t* variance = (uint8_t*)AscendC::GmAlloc(chBytes);
145+ uint8_t* xbp = (uint8_t*)AscendC::GmAlloc(outBytes);
146+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(16 * 1024 * 1024);
147+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingBytes);
148+ 
149+ float* g = reinterpret_cast<float*>(grads);
150+ float* s = reinterpret_cast<float*>(scale);
151+ float* v = reinterpret_cast<float*>(variance);
152+ for (int64_t i = 0; i < total; i++)
153+ g[i] = -1.5f + 0.3f * i;
154+ float sVals[3] = {0.5f, 1.25f, 2.0f};
155+ float vVals[3] = {0.1f, 1.0f, 4.0f};
156+ for (int64_t c = 0; c < C; c++) {
157+ s[c] = sVals[c];
158+ v[c] = vVals[c];
159+ }
160+ 
161+ auto* td = reinterpret_cast<BNInferGradTilingData<4>*>(tiling);
162+ FillSingleTileTiling<4>(td, eff, eps);
163+ 
164+ auto KernelBNInferGrad4 = [](GM_ADDR grads, GM_ADDR scale, GM_ADDR variance, GM_ADDR xbp, GM_ADDR ws,
165+ GM_ADDR tiling) {
166+ bn_infer_grad<BN_INFER_GRAD_RANK_4>(grads, scale, variance, xbp, ws, tiling);
167+ };
168+ ICPU_SET_TILING_KEY(BN_INFER_GRAD_RANK_4);
169+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
170+ ICPU_RUN_KF(KernelBNInferGrad4, td->multicore.num_cores, grads, scale, variance, xbp, workspace, tiling);
171+ 
172+ CheckGolden(eff, g, s, v, reinterpret_cast<float*>(xbp), eps);
173+ 
174+ AscendC::GmFree(grads);
175+ AscendC::GmFree(scale);
176+ AscendC::GmFree(variance);
177+ AscendC::GmFree(xbp);
178+ AscendC::GmFree(workspace);
179+ AscendC::GmFree(tiling);
180+}
181+ 
182+// RANK_4 路径，稍大 shape（多通道），仍单 tile。
183+TEST_F(bn_infer_grad_test, fp32_rank4_2d_larger_golden)
184+{
185+ std::vector<int64_t> eff = {8, 16}; // N=8, C=16
186+ const float eps = 0.0001f;
187+ int64_t total = eff[0] * eff[1];
188+ int64_t C = eff[1];
189+ 
190+ size_t gradsBytes = total * sizeof(float);
191+ size_t chBytes = C * sizeof(float);
192+ size_t tilingBytes = sizeof(BNInferGradTilingData<4>);
193+ 
194+ uint8_t* grads = (uint8_t*)AscendC::GmAlloc(gradsBytes);
195+ uint8_t* scale = (uint8_t*)AscendC::GmAlloc(chBytes);
196+ uint8_t* variance = (uint8_t*)AscendC::GmAlloc(chBytes);
197+ uint8_t* xbp = (uint8_t*)AscendC::GmAlloc(gradsBytes);
198+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(16 * 1024 * 1024);
199+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingBytes);
200+ 
201+ float* g = reinterpret_cast<float*>(grads);
202+ float* s = reinterpret_cast<float*>(scale);
203+ float* v = reinterpret_cast<float*>(variance);
204+ for (int64_t i = 0; i < total; i++)
205+ g[i] = 0.01f * ((i % 37) - 18);
206+ for (int64_t c = 0; c < C; c++) {
207+ s[c] = 0.25f + 0.1f * c;
208+ v[c] = 0.2f + 0.15f * c;
209+ }
210+ 
211+ auto* td = reinterpret_cast<BNInferGradTilingData<4>*>(tiling);
212+ FillSingleTileTiling<4>(td, eff, eps);
213+ 
214+ auto KernelBNInferGrad4 = [](GM_ADDR grads, GM_ADDR scale, GM_ADDR variance, GM_ADDR xbp, GM_ADDR ws,
215+ GM_ADDR tiling) {
216+ bn_infer_grad<BN_INFER_GRAD_RANK_4>(grads, scale, variance, xbp, ws, tiling);
217+ };
218+ ICPU_SET_TILING_KEY(BN_INFER_GRAD_RANK_4);
219+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
220+ ICPU_RUN_KF(KernelBNInferGrad4, td->multicore.num_cores, grads, scale, variance, xbp, workspace, tiling);
221+ 
222+ CheckGolden(eff, g, s, v, reinterpret_cast<float*>(xbp), eps);
223+ 
224+ AscendC::GmFree(grads);
225+ AscendC::GmFree(scale);
226+ AscendC::GmFree(variance);
227+ AscendC::GmFree(xbp);
228+ AscendC::GmFree(workspace);
229+ AscendC::GmFree(tiling);
230+}
231+ 
232+// RANK_8 路径（grads rank5），FP32，[N,C,D,H,W] 整块单 tile。
233+TEST_F(bn_infer_grad_test, fp32_rank8_5d_golden)
234+{
235+ std::vector<int64_t> eff = {2, 3, 2, 2, 2}; // N=2, C=3, D=H=W=2
236+ const float eps = 0.0001f;
237+ int64_t total = 1;
238+ for (auto d : eff)
239+ total *= d;
240+ int64_t C = eff[1];
241+ 
242+ size_t gradsBytes = total * sizeof(float);
243+ size_t chBytes = C * sizeof(float);
244+ size_t tilingBytes = sizeof(BNInferGradTilingData<8>);
245+ 
246+ uint8_t* grads = (uint8_t*)AscendC::GmAlloc(gradsBytes);
247+ uint8_t* scale = (uint8_t*)AscendC::GmAlloc(chBytes);
248+ uint8_t* variance = (uint8_t*)AscendC::GmAlloc(chBytes);
249+ uint8_t* xbp = (uint8_t*)AscendC::GmAlloc(gradsBytes);
250+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(16 * 1024 * 1024);
251+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingBytes);
252+ 
253+ float* g = reinterpret_cast<float*>(grads);
254+ float* s = reinterpret_cast<float*>(scale);
255+ float* v = reinterpret_cast<float*>(variance);
256+ for (int64_t i = 0; i < total; i++)
257+ g[i] = -1.0f + 0.13f * i;
258+ float sVals[3] = {0.4f, 1.1f, 1.8f};
259+ float vVals[3] = {0.3f, 1.5f, 3.2f};
260+ for (int64_t c = 0; c < C; c++) {
261+ s[c] = sVals[c];
262+ v[c] = vVals[c];
263+ }
264+ 
265+ auto* td = reinterpret_cast<BNInferGradTilingData<8>*>(tiling);
266+ FillSingleTileTiling<8>(td, eff, eps);
267+ 
268+ auto KernelBNInferGrad8 = [](GM_ADDR grads, GM_ADDR scale, GM_ADDR variance, GM_ADDR xbp, GM_ADDR ws,
269+ GM_ADDR tiling) {
270+ bn_infer_grad<BN_INFER_GRAD_RANK_8>(grads, scale, variance, xbp, ws, tiling);
271+ };
272+ ICPU_SET_TILING_KEY(BN_INFER_GRAD_RANK_8);
273+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
274+ ICPU_RUN_KF(KernelBNInferGrad8, td->multicore.num_cores, grads, scale, variance, xbp, workspace, tiling);
275+ 
276+ CheckGolden(eff, g, s, v, reinterpret_cast<float*>(xbp), eps);
277+ 
278+ AscendC::GmFree(grads);
279+ AscendC::GmFree(scale);
280+ AscendC::GmFree(variance);
281+ AscendC::GmFree(xbp);
282+ AscendC::GmFree(workspace);
283+ AscendC::GmFree(tiling);
284+}