Star934
1856
代码
代码
Issues238
Pull Requests339
Actions
项目讨论
Wiki
分析
Star934
1856
已合并
SeluGrad算子对齐竞品实现，修复int类型结果溢出问题 #8771
yulianjie创建于 25 天前
SeluGrad算子对齐竞品实现，修复int类型结果溢出问题 #8771
已合并
yulianjie创建于 25 天前
从lianjieyu/ops-nnfix_selu_grad合入到cann/ops-nnmaster
共 12 个文件变更+334-611
Mactivation/selu_grad/examples/arch35/test_aclnn_selu_grad.cpp+5-5
@@ -118,16 +118,16 @@ int main()
118 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result failed. ERROR: %d\n", ret); return ret);118 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result failed. ERROR: %d\n", ret); return ret);
119 119 
120 LOG_PRINT("\n=== SeluGrad Results ===\n");120 LOG_PRINT("\n=== SeluGrad Results ===\n");
121- LOG_PRINT("outputs > 0: y = SCALE * gradients = %.6f * grad\n", SCALE);121+ LOG_PRINT("outputs >= 0: y = SCALE * gradients = %.6f * grad\n", SCALE);
122- LOG_PRINT("outputs <= 0: y = grad * (outputs + SCALE_ALPHA) = grad * (out + %.6f)\n\n", SCALE_ALPHA_PRODUCT);122+ LOG_PRINT("outputs < 0: y = grad * (outputs + SCALE_ALPHA) = grad * (out + %.6f)\n\n", SCALE_ALPHA_PRODUCT);
123 123 
124 int pass = 0, fail = 0;124 int pass = 0, fail = 0;
125 for (int64_t i = 0; i < size; i++) {125 for (int64_t i = 0; i < size; i++) {
126 float expected;126 float expected;
127- if (outHostData[i] > 0) {127+ if (outHostData[i] < 0) {
128- expected = SCALE * gradHostData[i];
129- } else {
130 expected = gradHostData[i] * (outHostData[i] + SCALE_ALPHA_PRODUCT);128 expected = gradHostData[i] * (outHostData[i] + SCALE_ALPHA_PRODUCT);
129+ } else {
130+ expected = SCALE * gradHostData[i];
131 }131 }
132 bool ok = std::fabs(resultData[i] - expected) < 0.01f;132 bool ok = std::fabs(resultData[i] - expected) < 0.01f;
133 if (ok)133 if (ok)
Mactivation/selu_grad/op_graph/selu_grad_proto.h+1-1
@@ -36,7 +36,7 @@ namespace ge {
36 * y: A Tensor. Has the same type, shape and format as "gradients".36 * y: A Tensor. Has the same type, shape and format as "gradients".
37 *37 *
38 * @par Third-party framework compatibility38 * @par Third-party framework compatibility
39- * @li Compatible with the Pytorch operator selu_backward.39+ * @li Compatible with the TensorFlow operator SeluGrad.
40 */40 */
41#ifndef OPS_PROTO_DEF_SELUGRAD41#ifndef OPS_PROTO_DEF_SELUGRAD
42#define OPS_PROTO_DEF_SELUGRAD42#define OPS_PROTO_DEF_SELUGRAD
Mactivation/selu_grad/op_host/arch35/selu_grad_tiling_arch35.cpp+17-197
@@ -14,12 +14,10 @@
14 * \file selu_grad_tiling.cpp14 * \file selu_grad_tiling.cpp
15 * \brief SeluGrad 算子 Tiling 实现（arch35 架构）15 * \brief SeluGrad 算子 Tiling 实现（arch35 架构）
16 *16 *
17- * TilingKey_0 (OneDim) + TilingKey_1 (Broadcast)17+ * 仅保留 TilingKey_0 (OneDim)
18 * 支持 float16, float32, bfloat1618 * 支持 float16, float32, bfloat16
19 *19 *
20- * TilingKey 判定：20+ * gradients、outputs 和 y 的 shape 必须完全一致。
21- * - 合轴后仅 1 维（shape 完全一致或标量广播）→ TilingKey_0 (OneDim)
22- * - 合轴后 > 1 维（需要多维广播）→ TilingKey_1 (Broadcast)
23 */21 */
24 22 
25#include "register/op_def_registry.h"23#include "register/op_def_registry.h"
@@ -39,22 +37,11 @@ using Ops::Base::GetUbBlockSize;
39 37 
40constexpr uint32_t WS_SYS_SIZE = 0U;38constexpr uint32_t WS_SYS_SIZE = 0U;
41constexpr size_t WORKSPACE_NUM = 1;39constexpr size_t WORKSPACE_NUM = 1;
42-constexpr int32_t MAX_RANK = 8;
43 40 
44// Buffer 数量常量41// Buffer 数量常量
45constexpr int64_t SELECT_UB_RESERVE = 8192; // Select 8K 预留42constexpr int64_t SELECT_UB_RESERVE = 8192; // Select 8K 预留
46constexpr int64_t DEFAULT_BYTES_PER_ELEM = 40; // 默认每元素 UB 字节数（兜底）43constexpr int64_t DEFAULT_BYTES_PER_ELEM = 40; // 默认每元素 UB 字节数（兜底）
47 44 
48-static const gert::Shape g_vec_1_shape = {1};
49- 
50-static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape)
51-{
52- if (in_shape.GetDimNum() == 0) {
53- return g_vec_1_shape;
54- }
55- return in_shape;
56-}
57- 
58// 获取平台信息45// 获取平台信息
59static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t* ubSize, int64_t* coreNum)46static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t* ubSize, int64_t* coreNum)
60{47{
@@ -68,9 +55,8 @@ static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t* u
68 return ge::GRAPH_SUCCESS;55 return ge::GRAPH_SUCCESS;
69}56}
70 57 
71-// 获取 shape、dtype 信息（含广播推导）58+// 获取并校验 shape、dtype 信息
72-static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t* totalElements, ge::DataType* dataType,59+static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t* totalElements, ge::DataType* dataType)
73- gert::Shape* gradShapeOut, gert::Shape* outShapeOut, gert::Shape* yShapeOut)
74{60{
75 auto inputGrad = context->GetInputShape(0);61 auto inputGrad = context->GetInputShape(0);
76 OP_CHECK_NULL_WITH_CONTEXT(context, inputGrad);62 OP_CHECK_NULL_WITH_CONTEXT(context, inputGrad);
@@ -84,9 +70,8 @@ static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t*
84 OP_CHECK_NULL_WITH_CONTEXT(context, outputY);70 OP_CHECK_NULL_WITH_CONTEXT(context, outputY);
85 auto yShape = outputY->GetStorageShape();71 auto yShape = outputY->GetStorageShape();
86 72 
87- *gradShapeOut = gradShape;73+ OP_CHECK_IF(gradShape != outShape || gradShape != yShape,
88- *outShapeOut = outShape;74+ OP_LOGE(context, "gradients, outputs and y must have the same shape"), return ge::GRAPH_FAILED);
89- *yShapeOut = yShape;
90 75 
91 // 计算输出总元素数76 // 计算输出总元素数
92 if (yShape.GetDimNum() == 0) {77 if (yShape.GetDimNum() == 0) {
@@ -116,75 +101,7 @@ static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
116 return ge::GRAPH_SUCCESS;101 return ge::GRAPH_SUCCESS;
117}102}
118 103 
119-// 判断是否需要广播（shape 不完全一致）104+// 每元素 UB 字节数（buffer 数量随 dtype 的 Cast 路径而定）
120-static bool NeedsBroadcast(const gert::Shape& gradShape, const gert::Shape& outShape)
121-{
122- if (gradShape.GetDimNum() != outShape.GetDimNum()) {
123- return true;
124- }
125- for (size_t i = 0; i < gradShape.GetDimNum(); i++) {
126- if (gradShape.GetDim(i) != outShape.GetDim(i)) {
127- return true;
128- }
129- }
130- return false;
131-}
132- 
133-// 计算广播 stride（广播轴 stride=0）
134-static void ComputeBroadcastStrides(const gert::Shape& yShape, const gert::Shape& inputShape, int64_t* strides)
135-{
136- int64_t yRank = static_cast<int64_t>(yShape.GetDimNum());
137- int64_t inputRank = static_cast<int64_t>(inputShape.GetDimNum());
138- 
139- // 先计算 input 自身的 stride（从右到左累积）
140- int64_t inputStrides[MAX_RANK] = {0};
141- if (inputRank > 0) {
142- inputStrides[inputRank - 1] = 1;
143- for (int64_t d = inputRank - 2; d >= 0; d--) {
144- inputStrides[d] = inputStrides[d + 1] * inputShape.GetDim(d + 1);
145- }
146- }
147- 
148- // 映射到 output 维度（右对齐）
149- for (int64_t d = 0; d < yRank; d++) {
150- int64_t inputDimIdx = d - (yRank - inputRank);
151- if (inputDimIdx < 0) {
152- // 该维度在 input 中不存在（补维），stride = 0
153- strides[d] = 0;
154- } else if (inputShape.GetDim(inputDimIdx) == 1 && yShape.GetDim(d) > 1) {
155- // 广播轴：input dim=1, output dim>1 → stride=0
156- strides[d] = 0;
157- } else {
158- // 非广播轴：使用 input 原始 stride
159- strides[d] = inputStrides[inputDimIdx];
160- }
161- }
162-}
163- 
164-// 计算连续内维大小（从最内层开始，连续非广播轴）
165-// 修复 B2/B3: 当任一输入在此维度有 stride=0（广播轴）时，内维到此结束
166-// 这确保了内维中的两个输入都是连续的，可以安全地用 DataCopyPad 搬入
167-static int64_t ComputeInnerSize(const gert::Shape& yShape, const int64_t* gradStrides, const int64_t* outStrides)
168-{
169- int64_t rank = static_cast<int64_t>(yShape.GetDimNum());
170- if (rank == 0) {
171- return 1;
172- }
173- 
174- // 从最内层开始，找到第一个在任一输入中是广播轴的维度
175- // 内维 = 从最内层到第一个"任一广播轴"之间的所有维度
176- int64_t innerSize = 1;
177- for (int64_t d = rank - 1; d >= 0; d--) {
178- // 如果该维度在任一输入中是广播轴（stride=0），则内维到此为止
179- if (gradStrides[d] == 0 || outStrides[d] == 0) {
180- break;
181- }
182- innerSize *= yShape.GetDim(d);
183- }
184- return innerSize;
185-}
186- 
187-// 每元素 UB 字节数（OneDim 与 Broadcast 共用；buffer 数量随 dtype 的 Cast 路径而定）
188static int64_t GetBytesPerElem(ge::DataType dataType)105static int64_t GetBytesPerElem(ge::DataType dataType)
189{106{
190 switch (dataType) {107 switch (dataType) {
@@ -223,104 +140,10 @@ static void ComputeOneDimTiling(SeluGradTilingData* tiling, int64_t totalElement
223 }140 }
224}141}
225 142 
226-// Broadcast 内维分块 + UB/多核切分（已知 innerSize/totalRows/bytesPerElem 后）143+// 按元素数设置使用核数
227-static void ComputeBroadcastUbSplit(SeluGradTilingData* tiling, int64_t bytesPerElem, int64_t availableUb,144+static void SetUsedCoreNum(gert::TilingContext* context, const SeluGradTilingData* tiling, int64_t totalElements)
228- int64_t coreNum, int64_t ubBlockSize)
229{145{
230- // 内维分块：如果 innerSize 太大无法放入 UB，则分块处理146+ int64_t usedCoreNum = CeilDiv(totalElements, tiling->blockFormer);
231- if (bytesPerElem <= 0) {
232- return;
233- }
234- int64_t maxInnerSize = availableUb / bytesPerElem;
235- if (maxInnerSize < 1) {
236- return;
237- }
238- 
239- int64_t innerChunkSize;
240- int32_t numInnerChunks;
241- if (tiling->innerSize <= maxInnerSize) {
242- innerChunkSize = tiling->innerSize;
243- numInnerChunks = 1;
244- } else {
245- innerChunkSize = FloorAlign(maxInnerSize, ubBlockSize);
246- if (innerChunkSize < 1) {
247- innerChunkSize = 1;
248- }
249- numInnerChunks = static_cast<int32_t>((tiling->innerSize + innerChunkSize - 1) / innerChunkSize);
250- }
251- 
252- tiling->innerChunkSize = innerChunkSize;
253- tiling->numInnerChunks = numInnerChunks;
254- 
255- // 每个 "work item" 处理一个 chunk（innerChunkSize 个元素）
256- int64_t totalSubRows = tiling->totalRows * numInnerChunks;
257- 
258- // ubFormer = 每个 UB 能容纳的 work item 数
259- int64_t bytesPerItem = innerChunkSize * bytesPerElem;
260- tiling->ubFormer = availableUb / bytesPerItem;
261- if (tiling->ubFormer < 1) {
262- tiling->ubFormer = 1;
263- }
264- 
265- // 多核切分（按 work item 数），blockFormer 不超过 ubFormer
266- tiling->blockFormer = CeilDiv(totalSubRows, coreNum);
267- if (tiling->blockFormer > tiling->ubFormer) {
268- tiling->blockFormer = tiling->ubFormer;
269- }
270- if (tiling->blockFormer < 1) {
271- tiling->blockFormer = 1;
272- }
273-}
274- 
275-// TilingKey_1 (Broadcast) 路径参数计算
276-static void ComputeBroadcastTiling(SeluGradTilingData* tiling, const gert::Shape& gradShape,
277- const gert::Shape& outShape, const gert::Shape& yShape, int64_t totalElements,
278- ge::DataType dataType, uint64_t ubSize, int64_t coreNum, int64_t ubBlockSize)
279-{
280- tiling->totalElements = totalElements;
281- tiling->needBroadcast = 1;
282- tiling->shapeLen = static_cast<int32_t>(yShape.GetDimNum());
283- 
284- // 填充 outputDims
285- for (int32_t d = 0; d < tiling->shapeLen; d++) {
286- tiling->outputDims[d] = yShape.GetDim(d);
287- }
288- 
289- // 计算广播 stride
290- ComputeBroadcastStrides(yShape, gradShape, tiling->gradStrides);
291- ComputeBroadcastStrides(yShape, outShape, tiling->outStrides);
292- 
293- // 计算连续内维大小
294- tiling->innerSize = ComputeInnerSize(yShape, tiling->gradStrides, tiling->outStrides);
295- tiling->totalRows = totalElements / tiling->innerSize;
296- 
297- // UB 切分：根据 dtype 每元素内存，支持内维分块
298- int64_t availableUb = static_cast<int64_t>(ubSize) - SELECT_UB_RESERVE;
299- int64_t bytesPerElem = GetBytesPerElem(dataType);
300- ComputeBroadcastUbSplit(tiling, bytesPerElem, availableUb, coreNum, ubBlockSize);
301-}
302- 
303-// 判定 schMode 并计算对应路径的 tiling 参数，返回 schMode；needBroadcast 经出参回传
304-static uint32_t DispatchTiling(SeluGradTilingData* tiling, const gert::Shape& gradShape, const gert::Shape& outShape,
305- const gert::Shape& yShape, int64_t totalElements, ge::DataType dataType, uint64_t ubSize,
306- int64_t coreNum, int64_t ubBlockSize, bool* needBroadcast)
307-{
308- *needBroadcast = NeedsBroadcast(EnsureNotScalar(gradShape), EnsureNotScalar(outShape));
309- if (*needBroadcast) {
310- ComputeBroadcastTiling(tiling, EnsureNotScalar(gradShape), EnsureNotScalar(outShape), yShape, totalElements,
311- dataType, ubSize, coreNum, ubBlockSize);
312- return static_cast<uint32_t>(SELU_GRAD_BROADCAST);
313- }
314- ComputeOneDimTiling(tiling, totalElements, dataType, ubSize, coreNum, ubBlockSize);
315- return static_cast<uint32_t>(SELU_GRAD_ONE_DIM);
316-}
317- 
318-// 按 work item 数（Broadcast: 行×内维块；OneDim: 元素）设置使用核数
319-static void SetUsedCoreNum(gert::TilingContext* context, const SeluGradTilingData* tiling, int64_t totalElements,
320- bool needBroadcast)
321-{
322- int64_t workItems = needBroadcast ? (tiling->totalRows * tiling->numInnerChunks) : totalElements;
323- int64_t usedCoreNum = CeilDiv(workItems, tiling->blockFormer);
324 if (usedCoreNum < 1) {147 if (usedCoreNum < 1) {
325 usedCoreNum = 1;148 usedCoreNum = 1;
326 }149 }
@@ -339,10 +162,8 @@ static ge::graphStatus SeluGradTilingFunc(gert::TilingContext* context)
339 // 2. 获取 shape、属性信息162 // 2. 获取 shape、属性信息
340 int64_t totalElements;163 int64_t totalElements;
341 ge::DataType dataType;164 ge::DataType dataType;
342- gert::Shape gradShape, outShape, yShape;165+ OP_CHECK_IF(GetShapeAttrsInfo(context, &totalElements, &dataType) != ge::GRAPH_SUCCESS,
343- OP_CHECK_IF(166+ OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
344- GetShapeAttrsInfo(context, &totalElements, &dataType, &gradShape, &outShape, &yShape) != ge::GRAPH_SUCCESS,
345- OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
346 167 
347 // 3. 获取 WorkspaceSize168 // 3. 获取 WorkspaceSize
348 OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),169 OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
@@ -364,18 +185,17 @@ static ge::graphStatus SeluGradTilingFunc(gert::TilingContext* context)
364 return ge::GRAPH_SUCCESS;185 return ge::GRAPH_SUCCESS;
365 }186 }
366 187 
367- // 5. 判定 TilingKey 并计算 tiling188+ // 5. 计算 OneDim tiling
368- bool needBroadcast = false;189+ ComputeOneDimTiling(tiling, totalElements, dataType, ubSize, coreNum, ubBlockSize);
369- uint32_t schMode = DispatchTiling(tiling, gradShape, outShape, yShape, totalElements, dataType, ubSize, coreNum,190+ uint32_t schMode = static_cast<uint32_t>(SELU_GRAD_ONE_DIM);
370- ubBlockSize, &needBroadcast);
371 191 
372 // 6. 校验 tiling 参数有效性192 // 6. 校验 tiling 参数有效性
373 if (tiling->blockFormer < 1 || tiling->ubFormer < 1) {193 if (tiling->blockFormer < 1 || tiling->ubFormer < 1) {
374 return ge::GRAPH_FAILED;194 return ge::GRAPH_FAILED;
375 }195 }
376 196 
377- // 7. 设置核数 + TilingKey（仅 schMode，dtype 由 def 驱动展开）197+ // 7. 设置核数 + TilingKey
378- SetUsedCoreNum(context, tiling, totalElements, needBroadcast);198+ SetUsedCoreNum(context, tiling, totalElements);
379 ASCENDC_TPL_SEL_PARAM(context, schMode);199 ASCENDC_TPL_SEL_PARAM(context, schMode);
380 200 
381 return ge::GRAPH_SUCCESS;201 return ge::GRAPH_SUCCESS;
Mactivation/selu_grad/op_host/selu_grad_infershape.cpp+23-44
@@ -13,48 +13,35 @@
13/*!13/*!
14 * \file selu_grad_infershape.cpp14 * \file selu_grad_infershape.cpp
15 * \brief SeluGrad 算子形状推导实现15 * \brief SeluGrad 算子形状推导实现
16- *
17- * 迭代二：实现 numpy broadcast 形状推导
18- * 输出 shape = broadcast(gradients.shape, outputs.shape)
19 */16 */
20 17 
21#include "register/op_impl_registry.h"18#include "register/op_impl_registry.h"
22#include "exe_graph/runtime/infer_shape_context.h"19#include "exe_graph/runtime/infer_shape_context.h"
23#include "op_common/log/log.h"20#include "op_common/log/log.h"
21+#include "util/shape_util.h"
24 22 
25using namespace ge;23using namespace ge;
26 24 
27namespace ops {25namespace ops {
28 26 
29-constexpr int64_t MAX_SUPPORTED_RANK = 8;27+constexpr size_t MAX_SUPPORTED_RANK = 8;
28+constexpr int64_t UNKNOWN_DIM_VALUE = -1;
30 29 
31-static bool NumpyBroadcastShape(const gert::Shape& gradShape, const gert::Shape& outShape, gert::Shape& yShape)30+static bool AreShapesCompatibleWithoutBroadcast(const gert::Shape& gradientsShape, const gert::Shape& outputsShape)
32{31{
33- int64_t gradRank = static_cast<int64_t>(gradShape.GetDimNum());32+ if (Ops::Base::IsUnknownRank(gradientsShape) || Ops::Base::IsUnknownRank(outputsShape)) {
34- int64_t outRank = static_cast<int64_t>(outShape.GetDimNum());33+ return true;
35- int64_t maxRank = (gradRank > outRank) ? gradRank : outRank;34+ }
36- 35+ if (gradientsShape.GetDimNum() != outputsShape.GetDimNum()) {
37- if (maxRank > MAX_SUPPORTED_RANK) {
38 return false;36 return false;
39 }37 }
40- 38+ for (size_t i = 0; i < gradientsShape.GetDimNum(); ++i) {
41- gert::Shape result;39+ const int64_t gradientsDim = gradientsShape.GetDim(i);
42- for (int64_t i = 0; i < maxRank; i++) {40+ const int64_t outputsDim = outputsShape.GetDim(i);
43- int64_t gradDim = (i < maxRank - gradRank) ? 1 : gradShape.GetDim(i - (maxRank - gradRank));41+ if (gradientsDim != UNKNOWN_DIM_VALUE && outputsDim != UNKNOWN_DIM_VALUE && gradientsDim != outputsDim) {
44- int64_t outDim = (i < maxRank - outRank) ? 1 : outShape.GetDim(i - (maxRank - outRank));42+ return false;
45- 
46- if (gradDim == outDim) {
47- result.AppendDim(gradDim);
48- } else if (gradDim == 1) {
49- result.AppendDim(outDim);
50- } else if (outDim == 1) {
51- result.AppendDim(gradDim);
52- } else {
53- return false; // 不可广播
54 }43 }
55 }44 }
56- 
57- yShape = result;
58 return true;45 return true;
59}46}
60 47 
@@ -69,25 +56,17 @@ static ge::graphStatus InferShape4SeluGrad(gert::InferShapeContext* context)
69 gert::Shape* yShape = context->GetOutputShape(0);56 gert::Shape* yShape = context->GetOutputShape(0);
70 OP_CHECK_NULL_WITH_CONTEXT(context, yShape);57 OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
71 58 
72- // 处理 rank=0 标量59+ OP_CHECK_IF(!Ops::Base::IsUnknownRank(*gradShape) && gradShape->GetDimNum() > MAX_SUPPORTED_RANK,
73- if (gradShape->GetDimNum() == 0 && outShape->GetDimNum() == 0) {60+ OP_LOGE(context, "The rank of gradients must not exceed %zu", MAX_SUPPORTED_RANK),
74- // 两个标量输入，输出也是标量
75- return ge::GRAPH_SUCCESS;
76- }
77- 
78- // 处理 rank=0 标量 + tensor 广播
79- if (gradShape->GetDimNum() == 0) {
80- *yShape = *outShape;
81- return ge::GRAPH_SUCCESS;
82- }
83- if (outShape->GetDimNum() == 0) {
84- *yShape = *gradShape;
85- return ge::GRAPH_SUCCESS;
86- }
87- 
88- // numpy broadcast 推导
89- OP_CHECK_IF(!NumpyBroadcastShape(*gradShape, *outShape, *yShape), OP_LOGE(context, "Shape broadcast failed"),
90 return ge::GRAPH_FAILED);61 return ge::GRAPH_FAILED);
62+ OP_CHECK_IF(!Ops::Base::IsUnknownRank(*outShape) && outShape->GetDimNum() > MAX_SUPPORTED_RANK,
63+ OP_LOGE(context, "The rank of outputs must not exceed %zu", MAX_SUPPORTED_RANK),
64+ return ge::GRAPH_FAILED);
65+ OP_CHECK_IF(!AreShapesCompatibleWithoutBroadcast(*gradShape, *outShape),
66+ OP_LOGE(context, "gradients and outputs must have the same shape without broadcasting"),
67+ return ge::GRAPH_FAILED);
68+ 
69+ *yShape = *gradShape;
91 70 
92 return ge::GRAPH_SUCCESS;71 return ge::GRAPH_SUCCESS;
93}72}
Mactivation/selu_grad/op_kernel/arch35/selu_grad.h+53-245
@@ -14,14 +14,12 @@
14 * \file selu_grad.h14 * \file selu_grad.h
15 * \brief SeluGrad 算子 Kernel 类定义（arch35 架构）15 * \brief SeluGrad 算子 Kernel 类定义（arch35 架构）
16 *16 *
17- * 全 dtype × 2 SCH_MODE = 12 TilingKey。17+ * 全 dtype × 1 SCH_MODE = 6 个编译实例。
18 *18 *
19- * 设计：计算逻辑封装为两个 "Kit"，搬运/切分逻辑封装为两个驱动类，二者组合：19+ * 设计：计算逻辑封装为两个 "Kit"，搬运/切分逻辑封装为连续 OneDim 驱动类：
20- * - SeluGradDirectKit<T> : float 直算（在 T==float 上原地计算）20+ * - SeluGradDirectKit<T> : float/half/bfloat16 原 dtype 直算
21- * - SeluGradTransitKit<T> : half/bfloat16/int32/int8/uint8 FP32 中转21+ * - SeluGradTransitKit<T> : int8/uint8 在 FP16 计算，int32 在 FP32 计算
22- * (half/bf16 单步 Cast；int32/int8/uint8 经 half 两步 Cast)
23 * - SeluGradOneDim<T, Kit> : 连续分块搬运 + 多核切分22 * - SeluGradOneDim<T, Kit> : 连续分块搬运 + 多核切分
24- * - SeluGradBroadcast<T, Kit> : 多维 stride 偏移 + 内维分块逐行搬运
25 *23 *
26 * 公式：24 * 公式：
27 * y = SCALE * gradients if outputs >= 025 * y = SCALE * gradients if outputs >= 0
@@ -44,13 +42,12 @@ using namespace AscendC;
44constexpr float SCALE_F = 1.0507009873554804934193349852946f;42constexpr float SCALE_F = 1.0507009873554804934193349852946f;
45constexpr float SCALE_ALPHA_PRODUCT_F = 1.7580993408473768599402175208123f;43constexpr float SCALE_ALPHA_PRODUCT_F = 1.7580993408473768599402175208123f;
46 44 
47-// 需要经 half 两步 Cast（int → half → float）的 dtype（DAV_3510 不支持 int↔float 直接 Cast）45+// 整数计算类型：int8/uint8 上浮到 FP16，int32 上浮到 FP32。
48template <typename T>46template <typename T>
49-constexpr bool kNeedsHalfTransit = std::is_same_v<T, int32_t> || std::is_same_v<T, int8_t> ||47+constexpr bool kIntegerTransit = std::is_same_v<T, int32_t> || std::is_same_v<T, int8_t> || std::is_same_v<T, uint8_t>;
50- std::is_same_v<T, uint8_t>;
51 48 
52// ============================================================================49// ============================================================================
53-// 公共自由函数：搬运参数 / 计算核 / Cast / 广播偏移50+// 公共自由函数：搬运参数 / 计算核 / Cast
54// ============================================================================51// ============================================================================
55 52 
56template <typename T>53template <typename T>
@@ -71,101 +68,34 @@ __aicore__ inline void SeluGradSelectFp32(const LocalTensor<float>& yFp32, const
71 const LocalTensor<float>& branchB, const LocalTensor<float>& tmp,68 const LocalTensor<float>& branchB, const LocalTensor<float>& tmp,
72 const LocalTensor<uint8_t>& selMask, int64_t n)69 const LocalTensor<uint8_t>& selMask, int64_t n)
73{70{
74- CompareScalar(selMask, outFp32, (float)0.0f, CMPMODE::LE, n);71+ CompareScalar(selMask, outFp32, (float)0.0f, CMPMODE::LT, n);
atomgit-bot
atomgit-botatomgit-bot25 天前

🟡 Medium Priority

改动建议
71
- CompareScalar(selMask, outFp32, (float)0.0f, CMPMODE::LT, n);
71
+ CompareScalar(selMask, outFp32, (float)0.0f, CMPMODE::LE, n);
应用建议
likedislike
不准确？
Y
yulianjie
25 天前 评论：
75 Muls(branchA, gradFp32, SCALE_F, n);72 Muls(branchA, gradFp32, SCALE_F, n);
76 Adds(tmp, outFp32, SCALE_ALPHA_PRODUCT_F, n);73 Adds(tmp, outFp32, SCALE_ALPHA_PRODUCT_F, n);
77 Mul(branchB, gradFp32, tmp, n);74 Mul(branchB, gradFp32, tmp, n);
78 Select(yFp32, selMask, branchB, branchA, SELMODE::VSEL_TENSOR_TENSOR_MODE, n);75 Select(yFp32, selMask, branchB, branchA, SELMODE::VSEL_TENSOR_TENSOR_MODE, n);
79}76}
80 77 
81-// 输入 Cast: T → float（整型经 half 两步，浮点单步）78+// TensorFlow 对 half/bfloat16 直接在输入 dtype 上计算，并以 outputs < 0 选择负分支。
82template <typename T>79template <typename T>
83-__aicore__ inline void SeluGradCastInToFp32(const LocalTensor<float>& gradFp32, const LocalTensor<float>& outFp32,80+__aicore__ inline void SeluGradSelectTfNative(const LocalTensor<T>& y, const LocalTensor<T>& grad,
84- const LocalTensor<T>& gradLocal, const LocalTensor<T>& outLocal,81+ const LocalTensor<T>& out, const LocalTensor<T>& branchA,
85- const LocalTensor<half>& gradHalf, const LocalTensor<half>& outHalf,82+ const LocalTensor<T>& branchB, const LocalTensor<T>& tmp,
86- int64_t n)83+ const LocalTensor<uint8_t>& selMask, int64_t n)
87{84{
88- if constexpr (kNeedsHalfTransit<T>) {85+ const T scale = static_cast<T>(SCALE_F);
89- Cast(gradHalf, gradLocal, RoundMode::CAST_NONE, n);86+ const T scaleAlphaProduct = static_cast<T>(SCALE_ALPHA_PRODUCT_F);
90- Cast(outHalf, outLocal, RoundMode::CAST_NONE, n);87+ CompareScalar(selMask, out, static_cast<T>(0.0f), CMPMODE::LT, n);
91- Cast(gradFp32, gradHalf, RoundMode::CAST_NONE, n);88+ Muls(branchA, grad, scale, n);
92- Cast(outFp32, outHalf, RoundMode::CAST_NONE, n);89+ Adds(tmp, out, scaleAlphaProduct, n);
93- } else {90+ Mul(branchB, grad, tmp, n);
94- Cast(gradFp32, gradLocal, RoundMode::CAST_NONE, n);91+ Select(y, selMask, branchB, branchA, SELMODE::VSEL_TENSOR_TENSOR_MODE, n);
95- Cast(outFp32, outLocal, RoundMode::CAST_NONE, n);
96- }
97-}
98- 
99-// 输出 Cast: float → T（整型经 half 两步，浮点单步），统一 CAST_RINT
100-template <typename T>
101-__aicore__ inline void SeluGradCastFp32ToOut(const LocalTensor<T>& yLocal, const LocalTensor<float>& yFp32,
102- const LocalTensor<half>& yHalf, int64_t n)
103-{
104- if constexpr (kNeedsHalfTransit<T>) {
105- Cast(yHalf, yFp32, RoundMode::CAST_RINT, n);
106- Cast(yLocal, yHalf, RoundMode::CAST_RINT, n);
107- } else {
108- Cast(yLocal, yFp32, RoundMode::CAST_RINT, n);
109- }
110-}
111- 
112-// 广播路径：flatIdx → multiIdx 分解
113-__aicore__ inline void FlatIdxToMultiIdx(int64_t flatIdx, int32_t shapeLen, const int64_t* dims, int64_t* multiIdx)
114-{
115- for (int32_t d = shapeLen - 1; d >= 0; d--) {
116- if (dims[d] > 0) {
117- multiIdx[d] = flatIdx % dims[d];
118- flatIdx = flatIdx / dims[d];
119- } else {
120- multiIdx[d] = 0;
121- }
122- }
123-}
124- 
125-__aicore__ inline int64_t ComputeStrideOffset(int32_t shapeLen, const int64_t* multiIdx, const int64_t* strides)
126-{
127- int64_t offset = 0;
128- for (int32_t d = 0; d < shapeLen; d++) {
129- offset += multiIdx[d] * strides[d];
130- }
131- return offset;
132-}
133- 
134-// 广播行的 GM 偏移与本块元素数（chunkOffset 已折入三个偏移）
135-struct RowGeom {
136- int64_t gradOffset;
137- int64_t outOffset;
138- int64_t yOffset;
139- int64_t count;
140-};
141- 
142-__aicore__ inline RowGeom ComputeRowGeom(int64_t rowIdx, int32_t chunkIdx, int64_t innerSize, int32_t shapeLen,
143- int64_t innerChunkSize, const int64_t* outputDims, const int64_t* gradStrides,
144- const int64_t* outStrides)
145-{
146- int64_t flatIdx = rowIdx * innerSize;
147- int64_t multiIdx[SELU_GRAD_MAX_DIM];
148- FlatIdxToMultiIdx(flatIdx, shapeLen, outputDims, multiIdx);
149- 
150- int64_t chunkOffset = (int64_t)chunkIdx * innerChunkSize;
151- int64_t count = innerChunkSize;
152- if (chunkOffset + count > innerSize) {
153- count = innerSize - chunkOffset;
154- }
155- 
156- RowGeom geom;
157- geom.gradOffset = ComputeStrideOffset(shapeLen, multiIdx, gradStrides) + chunkOffset;
158- geom.outOffset = ComputeStrideOffset(shapeLen, multiIdx, outStrides) + chunkOffset;
159- geom.yOffset = rowIdx * innerSize + chunkOffset;
160- geom.count = count;
161- return geom;
162}92}
163 93 
164// ============================================================================94// ============================================================================
165// 计算 Kit：封装计算所需缓冲与一次 (grad, out) -> y 的计算95// 计算 Kit：封装计算所需缓冲与一次 (grad, out) -> y 的计算
166// ============================================================================96// ============================================================================
167 97 
168-// 直算 Kit（half/float；实际仅 float 实例化）98+// 原 dtype 直算 Kit（float 保持原逻辑；half/bfloat16 对齐 TensorFlow）
169template <typename T>99template <typename T>
170struct SeluGradDirectKit {100struct SeluGradDirectKit {
171 TQue<QuePosition::VECCALC, 1> branchAQueue, branchBQueue, tmpQueue, selMaskQueue;101 TQue<QuePosition::VECCALC, 1> branchAQueue, branchBQueue, tmpQueue, selMaskQueue;
@@ -188,7 +118,11 @@ struct SeluGradDirectKit {
188 LocalTensor<T> tmp = tmpQueue.template AllocTensor<T>();118 LocalTensor<T> tmp = tmpQueue.template AllocTensor<T>();
189 LocalTensor<uint8_t> selMask = selMaskQueue.template AllocTensor<uint8_t>();119 LocalTensor<uint8_t> selMask = selMaskQueue.template AllocTensor<uint8_t>();
190 120 
191- SeluGradSelectFp32(yLocal, gradLocal, outLocal, branchA, branchB, tmp, selMask, n);121+ if constexpr (std::is_same_v<T, float>) {
122+ SeluGradSelectFp32(yLocal, gradLocal, outLocal, branchA, branchB, tmp, selMask, n);
123+ } else {
124+ SeluGradSelectTfNative<T>(yLocal, gradLocal, outLocal, branchA, branchB, tmp, selMask, n);
125+ }
192 126 
193 branchAQueue.FreeTensor(branchA);127 branchAQueue.FreeTensor(branchA);
194 branchBQueue.FreeTensor(branchB);128 branchBQueue.FreeTensor(branchB);
@@ -197,62 +131,49 @@ struct SeluGradDirectKit {
197 }131 }
198};132};
199 133 
200-// FP32 中转 Kit（half/bfloat16/int32/int8/uint8）134+// 整数中转 Kit：按对应浮点精度计算。
135+// int32 使用向零截断回铸；int8/uint8 使用普通转换回铸。
201template <typename T>136template <typename T>
202struct SeluGradTransitKit {137struct SeluGradTransitKit {
203- TBuf<TPosition::VECCALC> gradHalfBuf, outHalfBuf, yHalfBuf;138+ using ComputeT = std::conditional_t<std::is_same_v<T, int32_t>, float, half>;
204- TBuf<TPosition::VECCALC> gradFp32Buf, outFp32Buf, branchABuf, branchBBuf, tmpBuf, yFp32Buf, maskBuf;139+ 
140+ TBuf<TPosition::VECCALC> gradBuf, outBuf, branchABuf, branchBBuf, tmpBuf, yBuf, maskBuf;
205 141 
206 __aicore__ inline void InitBufs(TPipe& pipe, int64_t n)142 __aicore__ inline void InitBufs(TPipe& pipe, int64_t n)
207 {143 {
208- if constexpr (kNeedsHalfTransit<T>) {144+ pipe.InitBuffer(gradBuf, n * sizeof(ComputeT));
209- pipe.InitBuffer(gradHalfBuf, n * sizeof(half));145+ pipe.InitBuffer(outBuf, n * sizeof(ComputeT));
210- pipe.InitBuffer(outHalfBuf, n * sizeof(half));146+ pipe.InitBuffer(branchABuf, n * sizeof(ComputeT));
211- if constexpr (!std::is_same_v<T, int8_t>) {147+ pipe.InitBuffer(branchBBuf, n * sizeof(ComputeT));
212- pipe.InitBuffer(yHalfBuf, n * sizeof(half)); // int8 复用 outHalfBuf148+ pipe.InitBuffer(tmpBuf, n * sizeof(ComputeT));
213- }149+ pipe.InitBuffer(yBuf, n * sizeof(ComputeT));
214- }
215- pipe.InitBuffer(gradFp32Buf, n * sizeof(float));
216- pipe.InitBuffer(outFp32Buf, n * sizeof(float));
217- pipe.InitBuffer(branchABuf, n * sizeof(float));
218- pipe.InitBuffer(branchBBuf, n * sizeof(float));
219- pipe.InitBuffer(tmpBuf, n * sizeof(float));
220- pipe.InitBuffer(yFp32Buf, n * sizeof(float));
221 pipe.InitBuffer(maskBuf, (n / 8) + 32);150 pipe.InitBuffer(maskBuf, (n / 8) + 32);
222 }151 }
223 152 
224 __aicore__ inline void Compute(const LocalTensor<T>& gradLocal, const LocalTensor<T>& outLocal,153 __aicore__ inline void Compute(const LocalTensor<T>& gradLocal, const LocalTensor<T>& outLocal,
225 const LocalTensor<T>& yLocal, int64_t n)154 const LocalTensor<T>& yLocal, int64_t n)
226 {155 {
227- LocalTensor<float> gradFp32 = gradFp32Buf.template Get<float>();156+ LocalTensor<ComputeT> grad = gradBuf.template Get<ComputeT>();
228- LocalTensor<float> outFp32 = outFp32Buf.template Get<float>();157+ LocalTensor<ComputeT> out = outBuf.template Get<ComputeT>();
229- LocalTensor<float> branchA = branchABuf.template Get<float>();158+ LocalTensor<ComputeT> branchA = branchABuf.template Get<ComputeT>();
230- LocalTensor<float> branchB = branchBBuf.template Get<float>();159+ LocalTensor<ComputeT> branchB = branchBBuf.template Get<ComputeT>();
231- LocalTensor<float> tmp = tmpBuf.template Get<float>();160+ LocalTensor<ComputeT> tmp = tmpBuf.template Get<ComputeT>();
161+ LocalTensor<ComputeT> y = yBuf.template Get<ComputeT>();
232 LocalTensor<uint8_t> selMask = maskBuf.template Get<uint8_t>();162 LocalTensor<uint8_t> selMask = maskBuf.template Get<uint8_t>();
233- LocalTensor<float> yFp32 = yFp32Buf.template Get<float>();
234 163 
235- LocalTensor<half> gradHalf;164+ Cast(grad, gradLocal, RoundMode::CAST_NONE, n);
236- LocalTensor<half> outHalf;165+ Cast(out, outLocal, RoundMode::CAST_NONE, n);
237- LocalTensor<half> yHalf;166+ SeluGradSelectTfNative<ComputeT>(y, grad, out, branchA, branchB, tmp, selMask, n);
238- if constexpr (kNeedsHalfTransit<T>) {167+ if constexpr (std::is_same_v<T, int32_t>) {
239- gradHalf = gradHalfBuf.template Get<half>();168+ Cast(yLocal, y, RoundMode::CAST_TRUNC, n);
240- outHalf = outHalfBuf.template Get<half>();169+ } else {
241- if constexpr (std::is_same_v<T, int8_t>) {170+ Cast(yLocal, y, RoundMode::CAST_NONE, n);
242- yHalf = outHalfBuf.template Get<half>();
243- } else {
244- yHalf = yHalfBuf.template Get<half>();
245- }
246 }171 }
247- 
248- SeluGradCastInToFp32<T>(gradFp32, outFp32, gradLocal, outLocal, gradHalf, outHalf, n);
249- SeluGradSelectFp32(yFp32, gradFp32, outFp32, branchA, branchB, tmp, selMask, n);
250- SeluGradCastFp32ToOut<T>(yLocal, yFp32, yHalf, n);
251 }172 }
252};173};
253 174 
254// ============================================================================175// ============================================================================
255-// 驱动类：连续分块搬运（OneDim） / 多维广播逐行搬运（Broadcast）176+// 驱动类：连续分块搬运（OneDim）
256// ============================================================================177// ============================================================================
257 178 
258template <typename T, typename Kit>179template <typename T, typename Kit>
@@ -333,123 +254,10 @@ private:
333 Kit kit_;254 Kit kit_;
334};255};
335 256 
336-template <typename T, typename Kit>257+// dtype 调度别名：float/half/bfloat16 走 Direct，整数按对应计算精度走 Transit。
337-class SeluGradBroadcast {
338-public:
339- __aicore__ inline void Init(GM_ADDR gradients, GM_ADDR outputs, GM_ADDR y, const SeluGradTilingData* tilingData)
340- {
341- totalElements_ = tilingData->totalElements;
342- if (totalElements_ == 0) {
343- return;
344- }
345- innerSize_ = tilingData->innerSize;
346- totalRows_ = tilingData->totalRows;
347- shapeLen_ = tilingData->shapeLen;
348- blockRows_ = tilingData->blockFormer;
349- innerChunkSize_ = tilingData->innerChunkSize;
350- numInnerChunks_ = tilingData->numInnerChunks;
351- 
352- for (int32_t d = 0; d < shapeLen_; d++) {
353- outputDims_[d] = tilingData->outputDims[d];
354- gradStrides_[d] = tilingData->gradStrides[d];
355- outStrides_[d] = tilingData->outStrides[d];
356- }
357- 
358- int64_t totalItems = totalRows_ * numInnerChunks_;
359- startItem_ = (int64_t)AscendC::GetBlockIdx() * blockRows_;
360- endItem_ = startItem_ + blockRows_;
361- if (endItem_ > totalItems) {
362- endItem_ = totalItems;
363- }
364- 
365- gradBase_ = (__gm__ T*)gradients;
366- outBase_ = (__gm__ T*)outputs;
367- yBase_ = (__gm__ T*)y;
368- 
369- pipe.InitBuffer(gradQueue, 1, innerChunkSize_ * sizeof(T));
370- pipe.InitBuffer(outQueue, 1, innerChunkSize_ * sizeof(T));
371- pipe.InitBuffer(yQueue, 1, innerChunkSize_ * sizeof(T));
372- kit_.InitBufs(pipe, innerChunkSize_);
373- }
374- 
375- __aicore__ inline void Process()
376- {
377- if (totalElements_ == 0) {
378- return;
379- }
380- for (int64_t item = startItem_; item < endItem_; item++) {
381- ProcessRow(item / numInnerChunks_, static_cast<int32_t>(item % numInnerChunks_));
382- }
383- }
384- 
385-private:
386- __aicore__ inline void ProcessRow(int64_t rowIdx, int32_t chunkIdx)
387- {
388- RowGeom geom = ComputeRowGeom(rowIdx, chunkIdx, innerSize_, shapeLen_, innerChunkSize_, outputDims_,
389- gradStrides_, outStrides_);
390- gradGM.SetGlobalBuffer(gradBase_ + geom.gradOffset, geom.count);
391- outGM.SetGlobalBuffer(outBase_ + geom.outOffset, geom.count);
392- yGM.SetGlobalBuffer(yBase_ + geom.yOffset, geom.count);
393- 
394- // CopyIn
395- LocalTensor<T> gradLocal = gradQueue.template AllocTensor<T>();
396- LocalTensor<T> outLocal = outQueue.template AllocTensor<T>();
397- DataCopyParams copyParams = MakeCopyParams<T>(geom.count);
398- DataCopyPad(gradLocal, gradGM, copyParams, {false, 0, 0, 0});
399- DataCopyPad(outLocal, outGM, copyParams, {false, 0, 0, 0});
400- gradQueue.EnQue(gradLocal);
401- outQueue.EnQue(outLocal);
402- 
403- // Compute
404- gradLocal = gradQueue.template DeQue<T>();
405- outLocal = outQueue.template DeQue<T>();
406- LocalTensor<T> yLocal = yQueue.template AllocTensor<T>();
407- kit_.Compute(gradLocal, outLocal, yLocal, geom.count);
408- yQueue.template EnQue<T>(yLocal);
409- gradQueue.FreeTensor(gradLocal);
410- outQueue.FreeTensor(outLocal);
411- 
412- // CopyOut
413- yLocal = yQueue.template DeQue<T>();
414- DataCopyPad(yGM, yLocal, MakeCopyParams<T>(geom.count));
415- yQueue.FreeTensor(yLocal);
416- }
417- 
418- TPipe pipe;
419- TQue<QuePosition::VECIN, 1> gradQueue, outQueue;
420- TQue<QuePosition::VECOUT, 1> yQueue;
421- GlobalTensor<T> gradGM, outGM, yGM;
422- 
423- int64_t totalElements_ = 0;
424- int64_t innerSize_ = 0;
425- int64_t totalRows_ = 0;
426- int64_t blockRows_ = 0;
427- int32_t shapeLen_ = 0;
428- int64_t innerChunkSize_ = 0;
429- int32_t numInnerChunks_ = 0;
430- 
431- int64_t outputDims_[SELU_GRAD_MAX_DIM];
432- int64_t gradStrides_[SELU_GRAD_MAX_DIM];
433- int64_t outStrides_[SELU_GRAD_MAX_DIM];
434- 
435- int64_t startItem_ = 0;
436- int64_t endItem_ = 0;
437- 
438- __gm__ T* gradBase_;
439- __gm__ T* outBase_;
440- __gm__ T* yBase_;
441- 
442- Kit kit_;
443-};
444- 
445-// dtype 调度别名：float 走 Direct，其余走 Transit
446template <typename T>258template <typename T>
447using SeluGradOneDimOp = SeluGradOneDim<259using SeluGradOneDimOp = SeluGradOneDim<
448- T, std::conditional_t<std::is_same_v<T, float>, SeluGradDirectKit<T>, SeluGradTransitKit<T>>>;260+ T, std::conditional_t<kIntegerTransit<T>, SeluGradTransitKit<T>, SeluGradDirectKit<T>>>;
449- 
450-template <typename T>
451-using SeluGradBroadcastOp = SeluGradBroadcast<
452- T, std::conditional_t<std::is_same_v<T, float>, SeluGradDirectKit<T>, SeluGradTransitKit<T>>>;
453 261 
454} // namespace NsSeluGrad262} // namespace NsSeluGrad
455 263 
Mactivation/selu_grad/op_kernel/arch35/selu_grad_tiling_data.h+3-22
@@ -13,34 +13,15 @@
13/*!13/*!
14 * \file selu_grad_tiling_data.h14 * \file selu_grad_tiling_data.h
15 * \brief TilingData 结构体定义（arch35 架构）15 * \brief TilingData 结构体定义（arch35 架构）
16- *
17- * 迭代二：TilingKey_0 (OneDim) + TilingKey_1 (Broadcast)
18- * 支持 float16, float32, bfloat16 三种浮点 dtype
19 */16 */
20 17 
21#ifndef _SELU_GRAD_TILING_DATA_H_18#ifndef _SELU_GRAD_TILING_DATA_H_
22#define _SELU_GRAD_TILING_DATA_H_19#define _SELU_GRAD_TILING_DATA_H_
23 20 
24-constexpr int32_t SELU_GRAD_MAX_DIM = 8;
25- 
26struct SeluGradTilingData {21struct SeluGradTilingData {
27- // === 基础信息（TilingKey_0 和 TilingKey_1 共用） ===22+ int64_t totalElements = 0; // 总元素数
28- int64_t totalElements = 0; // 输出总元素数（broadcast 后）23+ int64_t blockFormer = 0; // 每核元素数
29- int64_t blockFormer = 0; // OneDim: 每核元素数; Broadcast: 每核行数24+ int64_t ubFormer = 0; // 每次 UB 处理的元素数
30- int64_t ubFormer = 0; // OneDim: 每次 UB 元素数; Broadcast: 每次 UB 行数
31- 
32- // === 广播信息（TilingKey_1 使用） ===
33- int32_t needBroadcast = 0; // 0=无需广播, 1=需要广播
34- int32_t shapeLen = 0; // 合轴后维度数
35- int64_t innerSize = 0; // 连续内维大小（最内层连续非广播轴及其右侧元素积）
36- int64_t totalRows = 0; // 外层行数 = totalElements / innerSize
37- int64_t innerChunkSize = 0; // 内维分块大小（每块元素数，<= innerSize）
38- int32_t numInnerChunks = 0; // 内维分块数 = CeilDiv(innerSize, innerChunkSize)
39- 
40- // === 多维广播参数（仅 TilingKey_1 使用） ===
41- int64_t outputDims[SELU_GRAD_MAX_DIM]; // 合轴后输出 shape
42- int64_t gradStrides[SELU_GRAD_MAX_DIM]; // gradients 合轴后 stride（0=广播轴）
43- int64_t outStrides[SELU_GRAD_MAX_DIM]; // outputs 合轴后 stride（0=广播轴）
44};25};
45 26 
46#endif // _SELU_GRAD_TILING_DATA_H_27#endif // _SELU_GRAD_TILING_DATA_H_
Mactivation/selu_grad/op_kernel/arch35/selu_grad_tiling_key.h+3-7
@@ -14,8 +14,7 @@
14 * \file selu_grad_tiling_key.h14 * \file selu_grad_tiling_key.h
15 * \brief TilingKey 模板参数定义（arch35 架构）15 * \brief TilingKey 模板参数定义（arch35 架构）
16 *16 *
17- * dtype 由 def 驱动展开，tiling_key 仅编码 SCH_MODE 维度。17+ * dtype 由 def 驱动展开，仅保留连续 OneDim 调度模式。
18- * SCH_MODE: OneDim(0), Broadcast(1)
19 */18 */
20 19 
21#ifndef __SELU_GRAD_TILING_KEY_H__20#ifndef __SELU_GRAD_TILING_KEY_H__
@@ -24,12 +23,9 @@
24#include "ascendc/host_api/tiling/template_argument.h"23#include "ascendc/host_api/tiling/template_argument.h"
25 24 
26#define SELU_GRAD_ONE_DIM 025#define SELU_GRAD_ONE_DIM 0
27-#define SELU_GRAD_BROADCAST 1
28 26 
29-ASCENDC_TPL_ARGS_DECL(SeluGrad,27+ASCENDC_TPL_ARGS_DECL(SeluGrad, ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, SELU_GRAD_ONE_DIM));
30- ASCENDC_TPL_UINT_DECL(schMode, 8, ASCENDC_TPL_UI_LIST, SELU_GRAD_ONE_DIM, SELU_GRAD_BROADCAST));
31 28 
32-ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, SELU_GRAD_ONE_DIM)),29+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, SELU_GRAD_ONE_DIM)));
33- ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, SELU_GRAD_BROADCAST)));
34 30 
35#endif // __SELU_GRAD_TILING_KEY_H__31#endif // __SELU_GRAD_TILING_KEY_H__
Mactivation/selu_grad/op_kernel/selu_grad.cpp+1-3
@@ -15,7 +15,7 @@
15 * \brief SeluGrad 算子 Kernel 入口（arch35 架构）15 * \brief SeluGrad 算子 Kernel 入口（arch35 架构）
16 *16 *
17 * dtype 由 def 驱动展开，通过 DTYPE_X 宏获取实际类型。17 * dtype 由 def 驱动展开，通过 DTYPE_X 宏获取实际类型。
18- * tiling_key 仅编码 SCH_MODE（OneDim / Broadcast）。18+ * tiling_key 仅保留 OneDim 调度模式。
19 */19 */
20 20 
21#include "arch35/selu_grad.h"21#include "arch35/selu_grad.h"
@@ -36,7 +36,5 @@ __global__ __aicore__ void selu_grad(GM_ADDR gradients, GM_ADDR outputs, GM_ADDR
36 36 
37 if constexpr (SCH_MODE == SELU_GRAD_ONE_DIM) {37 if constexpr (SCH_MODE == SELU_GRAD_ONE_DIM) {
38 RunSeluGrad<NsSeluGrad::SeluGradOneDimOp<DTYPE_GRADIENTS>>(gradients, outputs, y, &tilingData);38 RunSeluGrad<NsSeluGrad::SeluGradOneDimOp<DTYPE_GRADIENTS>>(gradients, outputs, y, &tilingData);
39- } else if constexpr (SCH_MODE == SELU_GRAD_BROADCAST) {
40- RunSeluGrad<NsSeluGrad::SeluGradBroadcastOp<DTYPE_GRADIENTS>>(gradients, outputs, y, &tilingData);
41 }39 }
42}40}
Mactivation/selu_grad/tests/ut/op_host/CMakeLists.txt+7-2
@@ -1,13 +1,18 @@
1# Copyright (c) 2025 Huawei Technologies Co., Ltd.1# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").3# CANN Open Software License Agreement Version 2.0 (the "License").
4# Please refer to the License for details. You may not use this file except in compliance with the License.4# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.6# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.7# See LICENSE in the root of the software repository for the full text of the License.
8#/8#/
9 9 
10message(STATUS "=== Debug: start ops.activation.selu_grad.tests.ut.CMakeLists.txt ")10message(STATUS "=== Debug: start ops.activation.selu_grad.tests.ut.CMakeLists.txt ")
11+if(UT_TEST_ALL OR OP_HOST_UT)
12+ add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
13+ add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+endif()
15+ 
11file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)16file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
12message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")17message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
13foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})18foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
Mactivation/selu_grad/tests/ut/op_host/arch35/test_selu_grad_tiling.cpp+108-82
@@ -8,112 +8,138 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11-#include <gtest/gtest.h>
12- 
13#include <map>11#include <map>
14#include <string>12#include <string>
15 13 
14+#include <gtest/gtest.h>
15+ 
16+#include "exe_graph/runtime/storage_format.h"
16#include "exe_graph/runtime/storage_shape.h"17#include "exe_graph/runtime/storage_shape.h"
17#include "kernel_run_context_facker.h"18#include "kernel_run_context_facker.h"
18#include "platform/platform_infos_def.h"19#include "platform/platform_infos_def.h"
20+#include "register/op_impl_registry.h"
19#include "test_cube_util.h"21#include "test_cube_util.h"
22+#include "ut_op_common.h"
23+#include "ut_op_util.h"
20 24 
21-namespace optiling {25+using namespace ge;
22-struct SeluGradCompileInfo {};26+using namespace ut_util;
23-} // namespace optiling
24 27 
25namespace {28namespace {
26-constexpr const char* OP_TYPE = "SeluGrad";
27-constexpr const char* COMPILE_INFO = R"({
28- "hardware_info": {
29- "BT_SIZE": 0,
30- "load3d_constraints": "1",
31- "Intrinsic_fix_pipe_l0c2out": false,
32- "Intrinsic_data_move_l12ub": true,
33- "Intrinsic_data_move_l0c2ub": true,
34- "Intrinsic_data_move_out2l1_nd2nz": false,
35- "UB_SIZE": 245760,
36- "L2_SIZE": 33554432,
37- "L1_SIZE": 524288,
38- "L0A_SIZE": 65536,
39- "L0B_SIZE": 65536,
40- "L0C_SIZE": 131072,
41- "CORE_NUM": 64
42- }
43-})";
44 29 
45-ge::graphStatus RunTiling(ge::DataType gradientsDtype, ge::DataType outputsDtype, ge::DataType yDtype)30+struct SeluGradTilingCompileInfo {};
31+ 
32+std::string GetCompileInfo()
46{33{
47- gert::StorageShape shape = {{16}, {16}};34+ return R"({
35+ "hardware_info": {
36+ "BT_SIZE": 0,
37+ "load3d_constraints": "1",
38+ "Intrinsic_fix_pipe_l0c2out": false,
39+ "Intrinsic_data_move_l12ub": true,
40+ "Intrinsic_data_move_l0c2ub": true,
41+ "Intrinsic_data_move_out2l1_nd2nz": false,
42+ "UB_SIZE": 196608,
43+ "L2_SIZE": 33554432,
44+ "L1_SIZE": 524288,
45+ "L0A_SIZE": 65536,
46+ "L0B_SIZE": 65536,
47+ "L0C_SIZE": 131072,
48+ "CORE_NUM": 48
49+ }
50+ })";
51+}
52+ 
53+void ExpectTilingStatus(const gert::StorageShape& gradientsShape, const gert::StorageShape& outputsShape,
54+ const gert::StorageShape& yShape, ge::graphStatus expectedStatus)
55+{
56+ std::string compileInfoString = GetCompileInfo();
48 std::map<std::string, std::string> socInfos;57 std::map<std::string, std::string> socInfos;
49 std::map<std::string, std::string> aicoreSpec;58 std::map<std::string, std::string> aicoreSpec;
50 std::map<std::string, std::string> intrinsics;59 std::map<std::string, std::string> intrinsics;
51- std::map<std::string, std::string> version = {{"Short_SoC_version", "Ascend950"}, {"NpuArch", "3510"}};60+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
52- GetPlatFormInfos(COMPILE_INFO, socInfos, aicoreSpec, intrinsics);
53 61 
54 fe::PlatFormInfos platformInfo;62 fe::PlatFormInfos platformInfo;
55 platformInfo.Init();63 platformInfo.Init();
56- optiling::SeluGradCompileInfo compileInfo;64+ SeluGradTilingCompileInfo compileInfo;
57- 
58- auto* impl = gert::OpImplRegistry::GetInstance().GetOpImpl(OP_TYPE);
59- if (impl == nullptr || impl->tiling_parse == nullptr || impl->tiling == nullptr) {
60- ADD_FAILURE() << "SeluGrad tiling callbacks are not registered";
61- return ge::GRAPH_FAILED;
62- }
63- 
64- auto parseContextHolder = gert::KernelRunContextFaker()
65- .KernelIONum(2, 1)
66- .Inputs({const_cast<char*>(COMPILE_INFO), reinterpret_cast<void*>(&platformInfo)})
67- .Outputs({&compileInfo})
68- .Build();
69- auto* parseContext = parseContextHolder.GetContext<gert::TilingParseContext>();
70- if (parseContext == nullptr || parseContext->GetPlatformInfo() == nullptr ||
71- !parseContext->GetPlatformInfo()->Init()) {
72- ADD_FAILURE() << "Failed to create the SeluGrad tiling parse context";
73- return ge::GRAPH_FAILED;
74- }
75- auto* parsePlatformInfo = parseContext->GetPlatformInfo();
76- parsePlatformInfo->SetPlatformRes("SoCInfo", socInfos);
77- parsePlatformInfo->SetPlatformRes("AICoreSpec", aicoreSpec);
78- parsePlatformInfo->SetCoreNumByCoreType("AICore");
79- parsePlatformInfo->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
80- parsePlatformInfo->SetPlatformRes("version", version);
81- if (impl->tiling_parse(parseContextHolder.GetContext<gert::KernelContext>()) != ge::GRAPH_SUCCESS) {
82- ADD_FAILURE() << "SeluGrad tiling parse failed";
83- return ge::GRAPH_FAILED;
84- }
85- 
86 auto tilingData = gert::TilingData::CreateCap(4096);65 auto tilingData = gert::TilingData::CreateCap(4096);
87- auto workspaceHolder = gert::ContinuousVector::Create<size_t>(4096);66+ auto workspaceSizeHolder = gert::ContinuousVector::Create<size_t>(1);
88- if (tilingData == nullptr || workspaceHolder == nullptr) {67+ auto* workspaceSizes = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHolder.get());
89- ADD_FAILURE() << "Failed to allocate SeluGrad tiling test buffers";68+ ASSERT_NE(tilingData, nullptr);
90- return ge::GRAPH_FAILED;69+ ASSERT_NE(workspaceSizes, nullptr);
91- }70+ 
92- auto* workspace = reinterpret_cast<gert::ContinuousVector*>(workspaceHolder.get());71+ auto* opImpl = gert::OpImplRegistry::GetInstance().GetOpImpl("SeluGrad");
93- auto contextHolder = gert::TilingContextFaker()72+ ASSERT_NE(opImpl, nullptr);
94- .SetOpType(OP_TYPE)73+ ASSERT_NE(opImpl->tiling, nullptr);
95- .NodeIoNum(2, 1)74+ 
96- .IrInstanceNum({1, 1})75+ auto holder = gert::TilingContextFaker()
97- .InputShapes({&shape, &shape})76+ .SetOpType("SeluGrad")
98- .OutputShapes({&shape})77+ .NodeIoNum(2, 1)
99- .CompileInfo(&compileInfo)78+ .IrInstanceNum({1, 1})
100- .PlatformInfo(reinterpret_cast<char*>(&platformInfo))79+ .InputShapes({const_cast<gert::StorageShape*>(&gradientsShape),
101- .NodeInputTd(0, gradientsDtype, ge::FORMAT_ND, ge::FORMAT_ND)80+ const_cast<gert::StorageShape*>(&outputsShape)})
102- .NodeInputTd(1, outputsDtype, ge::FORMAT_ND, ge::FORMAT_ND)81+ .OutputShapes({const_cast<gert::StorageShape*>(&yShape)})
103- .NodeOutputTd(0, yDtype, ge::FORMAT_ND, ge::FORMAT_ND)82+ .CompileInfo(&compileInfo)
104- .TilingData(tilingData.get())83+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
105- .Workspace(workspace)84+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
106- .Build();85+ .NodeInputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
107- return impl->tiling(contextHolder.GetContext<gert::TilingContext>());86+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
87+ .TilingData(tilingData.get())
88+ .Workspace(workspaceSizes)
89+ .Build();
90+ auto* context = holder.GetContext<gert::TilingContext>();
91+ ASSERT_NE(context, nullptr);
92+ context->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
93+ context->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
94+ context->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
95+ context->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
96+ 
97+ EXPECT_EQ(opImpl->tiling(context), expectedStatus);
108}98}
99+ 
109} // namespace100} // namespace
110 101 
111-TEST(SeluGradTilingTest, AcceptsMatchingSupportedDtypes)102+TEST(SeluGradTilingTest, SameShapeSucceeds)
112{103{
113- EXPECT_EQ(RunTiling(ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT), ge::GRAPH_SUCCESS);104+ const gert::StorageShape gradientsShape({2, 3, 4}, {2, 3, 4});
105+ const gert::StorageShape outputsShape({2, 3, 4}, {2, 3, 4});
106+ const gert::StorageShape yShape({2, 3, 4}, {2, 3, 4});
107+ 
108+ ExpectTilingStatus(gradientsShape, outputsShape, yShape, ge::GRAPH_SUCCESS);
114}109}
115 110 
116-TEST(SeluGradTilingTest, RejectsMismatchedSupportedInputDtypes)111+TEST(SeluGradTilingTest, SameEmptyShapeSucceeds)
117{112{
118- EXPECT_EQ(RunTiling(ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_FLOAT), ge::GRAPH_FAILED);113+ const gert::StorageShape gradientsShape({0, 3, 4}, {0, 3, 4});
114+ const gert::StorageShape outputsShape({0, 3, 4}, {0, 3, 4});
115+ const gert::StorageShape yShape({0, 3, 4}, {0, 3, 4});
116+ 
117+ ExpectTilingStatus(gradientsShape, outputsShape, yShape, ge::GRAPH_SUCCESS);
118+}
119+ 
120+TEST(SeluGradTilingTest, BroadcastableShapeFails)
121+{
122+ const gert::StorageShape gradientsShape({2, 3, 4}, {2, 3, 4});
123+ const gert::StorageShape outputsShape({1, 3, 1}, {1, 3, 1});
124+ const gert::StorageShape yShape({2, 3, 4}, {2, 3, 4});
125+ 
126+ ExpectTilingStatus(gradientsShape, outputsShape, yShape, ge::GRAPH_FAILED);
127+}
128+ 
129+TEST(SeluGradTilingTest, EmptyBroadcastableShapeFails)
130+{
131+ const gert::StorageShape gradientsShape({0, 3, 4}, {0, 3, 4});
132+ const gert::StorageShape outputsShape({1, 3, 1}, {1, 3, 1});
133+ const gert::StorageShape yShape({0, 3, 4}, {0, 3, 4});
134+ 
135+ ExpectTilingStatus(gradientsShape, outputsShape, yShape, ge::GRAPH_FAILED);
136+}
137+ 
138+TEST(SeluGradTilingTest, OutputShapeMismatchFails)
139+{
140+ const gert::StorageShape gradientsShape({2, 3, 4}, {2, 3, 4});
141+ const gert::StorageShape outputsShape({2, 3, 4}, {2, 3, 4});
142+ const gert::StorageShape yShape({2, 3, 5}, {2, 3, 5});
143+ 
144+ ExpectTilingStatus(gradientsShape, outputsShape, yShape, ge::GRAPH_FAILED);
119}145}
Mactivation/selu_grad/tests/ut/op_host/op_api/test_aclnn_selubackward.cpp+3-3
@@ -120,11 +120,11 @@ TEST_F(selu_backward_test, test_selubackward_format_3)
120 }120 }
121}121}
122 122 
123-TEST_F(selu_backward_test, test_selubackward_inconsistent_shape)123+TEST_F(selu_backward_test, test_selubackward_broadcast_shape_error)
124{124{
125 auto gradoutput = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND);125 auto gradoutput = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND);
126- auto result = TensorDesc({2, 16, 32, 18}, ACL_FLOAT, ACL_FORMAT_ND);126+ auto result = TensorDesc({1, 16, 1, 16}, ACL_FLOAT, ACL_FORMAT_ND);
127- auto gradinput = TensorDesc({2, 16, 32, 15}, ACL_FLOAT, ACL_FORMAT_ND);127+ auto gradinput = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND);
128 128 
129 auto ut = OP_API_UT(aclnnSeluBackward, INPUT(gradoutput, result), OUTPUT(gradinput));129 auto ut = OP_API_UT(aclnnSeluBackward, INPUT(gradoutput, result), OUTPUT(gradinput));
130 130 
Aactivation/selu_grad/tests/ut/op_host/test_selu_grad_infershape.cpp+110-0
@@ -0,0 +1,110 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <gtest/gtest.h>
12+ 
13+#include "exe_graph/runtime/storage_shape.h"
14+#include "kernel_run_context_facker.h"
15+#include "register/op_impl_registry.h"
16+ 
17+class SeluGradInferShapeTest : public testing::Test {};
18+ 
19+TEST_F(SeluGradInferShapeTest, SameShapeSucceeds)
20+{
21+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("SeluGrad")->infer_shape;
22+ gert::Shape gradientsShape = {2, 3, 4};
23+ gert::Shape outputsShape = {2, 3, 4};
24+ gert::Shape yShape = {};
25+ 
26+ auto holder = gert::InferShapeContextFaker()
27+ .NodeIoNum(2, 1)
28+ .IrInstanceNum({1, 1})
29+ .InputShapes({&gradientsShape, &outputsShape})
30+ .OutputShapes({&yShape})
31+ .Build();
32+ 
33+ auto context = holder.GetContext<gert::InferShapeContext>();
34+ ASSERT_EQ(inferShapeFunc(context), ge::GRAPH_SUCCESS);
35+ const gert::Shape* inferredYShape = context->GetOutputShape(0);
36+ ASSERT_NE(inferredYShape, nullptr);
37+ EXPECT_EQ(*inferredYShape, gradientsShape);
38+}
39+ 
40+TEST_F(SeluGradInferShapeTest, BroadcastableShapeFails)
41+{
42+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("SeluGrad")->infer_shape;
43+ gert::Shape gradientsShape = {2, 3, 4};
44+ gert::Shape outputsShape = {1, 3, 1};
45+ gert::Shape yShape = {};
46+ 
47+ auto holder = gert::InferShapeContextFaker()
48+ .NodeIoNum(2, 1)
49+ .IrInstanceNum({1, 1})
50+ .InputShapes({&gradientsShape, &outputsShape})
51+ .OutputShapes({&yShape})
52+ .Build();
53+ 
54+ EXPECT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_FAILED);
55+}
56+ 
57+TEST_F(SeluGradInferShapeTest, CompatibleDynamicDimensionSucceeds)
58+{
59+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("SeluGrad")->infer_shape;
60+ gert::Shape gradientsShape = {-1, 3, 4};
61+ gert::Shape outputsShape = {2, 3, 4};
62+ gert::Shape yShape = {};
63+ 
64+ auto holder = gert::InferShapeContextFaker()
65+ .NodeIoNum(2, 1)
66+ .IrInstanceNum({1, 1})
67+ .InputShapes({&gradientsShape, &outputsShape})
68+ .OutputShapes({&yShape})
69+ .Build();
70+ 
71+ auto context = holder.GetContext<gert::InferShapeContext>();
72+ ASSERT_EQ(inferShapeFunc(context), ge::GRAPH_SUCCESS);
73+ const gert::Shape* inferredYShape = context->GetOutputShape(0);
74+ ASSERT_NE(inferredYShape, nullptr);
75+ EXPECT_EQ(*inferredYShape, gradientsShape);
76+}
77+ 
78+TEST_F(SeluGradInferShapeTest, DynamicDimensionDoesNotHideKnownMismatch)
79+{
80+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("SeluGrad")->infer_shape;
81+ gert::Shape gradientsShape = {-1, 3, 4};
82+ gert::Shape outputsShape = {2, 5, 4};
83+ gert::Shape yShape = {};
84+ 
85+ auto holder = gert::InferShapeContextFaker()
86+ .NodeIoNum(2, 1)
87+ .IrInstanceNum({1, 1})
88+ .InputShapes({&gradientsShape, &outputsShape})
89+ .OutputShapes({&yShape})
90+ .Build();
91+ 
92+ EXPECT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_FAILED);
93+}
94+ 
95+TEST_F(SeluGradInferShapeTest, UnknownGradientsRankDoesNotBypassOutputsRankLimit)
96+{
97+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("SeluGrad")->infer_shape;
98+ gert::Shape gradientsShape = {-2};
99+ gert::Shape outputsShape = {1, 1, 1, 1, 1, 1, 1, 1, 1};
100+ gert::Shape yShape = {};
101+ 
102+ auto holder = gert::InferShapeContextFaker()
103+ .NodeIoNum(2, 1)
104+ .IrInstanceNum({1, 1})
105+ .InputShapes({&gradientsShape, &outputsShape})
106+ .OutputShapes({&yShape})
107+ .Build();
108+ 
109+ EXPECT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_FAILED);
110+}