Star0
0
代码
代码
Issues
Pull Requests
Actions
项目讨论
Wiki
分析
Star0
0
已合并
fix: def-driven dtype + InferDataType separation for Celu/DataFormatDimMap/ApplyProximalGradientDescent #1
babeiee创建于 8月6日
fix: def-driven dtype + InferDataType separation for Celu/DataFormatDimMap/ApplyProximalGradientDescent #1
已合并
babeiee创建于 8月6日
从fix/def-driven-dtype-and-infer-celu-dfdm-apgd合入到master
共 20 个文件变更+395-127
Aactivation/celu/examples/arch35/test_geir_celu_dynamic.cpp+240-0
@@ -0,0 +1,240 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under
4+* the terms of conditions of CANN Open Software License Agreement Version 2.0.
5+* Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+*/
7+ 
8+#include <iostream>
9+#include <vector>
10+#include <string>
11+#include <map>
12+#include <cmath>
13+#include <ctime>
14+ 
15+#include "graph.h"
16+#include "types.h"
17+#include "tensor.h"
18+#include "ge_error_codes.h"
19+#include "ge_api_types.h"
20+#include "ge_api.h"
21+#include "array_ops.h"
22+#include "nonlinear_fuc_ops.h"
23+#include "ge_ir_build.h"
24+ 
25+#define FAILED -1
26+#define SUCCESS 0
27+ 
28+using namespace ge;
29+using std::map;
30+using std::string;
31+using std::vector;
32+ 
33+static string GetTime()
34+{
35+ time_t timep;
36+ time(&timep);
37+ char tmp[64];
38+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
39+ return tmp;
40+}
41+ 
42+static Tensor MakeFloatTensor(const vector<int64_t>& shape, TensorDesc& desc, float fillValue)
43+{
44+ desc = TensorDesc(ge::Shape(shape), FORMAT_ND, DT_FLOAT);
45+ desc.SetPlacement(ge::kPlacementHost);
46+ desc.SetFormat(FORMAT_ND);
47+ desc.SetRealDimCnt(shape.size());
48+ size_t n = 1;
49+ for (auto d : shape) n *= static_cast<size_t>(d);
50+ uint32_t dataLen = n * sizeof(float);
51+ float* pData = new (std::nothrow) float[n];
52+ for (size_t i = 0; i < n; ++i) pData[i] = fillValue;
53+ Tensor tensor(desc, reinterpret_cast<uint8_t*>(pData), dataLen);
54+ delete[] pData;
55+ return tensor;
56+}
57+ 
58+static bool VerifyOutput(const ge::Tensor& tensor, const vector<int64_t>& expectedShape,
59+ float inputValue, float tolerance)
60+{
61+ auto outShape = tensor.GetTensorDesc().GetShape();
62+ if (outShape.GetDimNum() != expectedShape.size()) return false;
63+ for (size_t i = 0; i < expectedShape.size(); ++i) {
64+ if (outShape.GetDim(i) != expectedShape[i]) return false;
65+ }
66+ if (tensor.GetTensorDesc().GetDataType() != DT_FLOAT) return false;
67+ if (tensor.GetData() == nullptr) return false;
68+ const float* data = reinterpret_cast<const float*>(tensor.GetData());
69+ int64_t n = outShape.GetShapeSize();
70+ float alpha1 = 1.0f, alpha2 = 1.0f, alpha3 = 1.0f;
71+ float bound = 87.0f;
72+ float ratio = inputValue / alpha2;
73+ if (ratio > bound) ratio = bound;
74+ float expected;
75+ if (inputValue >= 0.0f) {
76+ expected = alpha3 * inputValue;
77+ } else {
78+ expected = alpha1 * (expf(ratio) - 1.0f);
79+ }
80+ for (int64_t i = 0; i < n; ++i) {
81+ if (std::abs(data[i] - expected) > tolerance) return false;
82+ }
83+ return true;
84+}
85+ 
86+static int RunScenario(Session* session, uint32_t graphId,
87+ const vector<vector<int64_t>>& shapes,
88+ const string& scenarioName)
89+{
90+ float inputValue = -0.5f;
91+ float tolerance = 1e-5f;
92+ int passed = 0;
93+ 
94+ printf("Scenario %s, declared shape [-1,-1]\n", scenarioName.c_str());
95+ for (const auto& s : shapes) {
96+ printf("Run concrete shape [%ld,%ld]\n", s[0], s[1]);
97+ TensorDesc desc;
98+ Tensor tensor = MakeFloatTensor(s, desc, inputValue);
99+ vector<ge::Tensor> feeds{tensor};
100+ vector<ge::Tensor> output;
101+ Status ret = session->RunGraph(graphId, feeds, output);
102+ if (ret != SUCCESS || output.empty()) {
103+ printf("Shape, dtype and values FAILED for [%ld,%ld]\n", s[0], s[1]);
104+ continue;
105+ }
106+ if (VerifyOutput(output[0], s, inputValue, tolerance)) {
107+ printf("Shape, dtype and values PASSED for [%ld,%ld]\n", s[0], s[1]);
108+ passed++;
109+ } else {
110+ printf("Shape, dtype and values FAILED for [%ld,%ld]\n", s[0], s[1]);
111+ }
112+ }
113+ printf("Scenario %s summary: %d/%d passed\n", scenarioName.c_str(), passed, static_cast<int>(shapes.size()));
114+ return passed == static_cast<int>(shapes.size()) ? SUCCESS : FAILED;
115+}
116+ 
117+static int RunRankScenario(Session* session, uint32_t graphId,
118+ const vector<vector<int64_t>>& shapes)
119+{
120+ float inputValue = 0.3f;
121+ float tolerance = 1e-5f;
122+ int passed = 0;
123+ 
124+ printf("Scenario unknown_rank_minus_2, declared shape [-2]\n");
125+ for (const auto& s : shapes) {
126+ printf("Run concrete shape [");
127+ for (size_t i = 0; i < s.size(); ++i) {
128+ printf("%ld%s", s[i], i + 1 < s.size() ? "," : "");
129+ }
130+ printf("]\n");
131+ TensorDesc desc;
132+ Tensor tensor = MakeFloatTensor(s, desc, inputValue);
133+ vector<ge::Tensor> feeds{tensor};
134+ vector<ge::Tensor> output;
135+ Status ret = session->RunGraph(graphId, feeds, output);
136+ if (ret != SUCCESS || output.empty()) {
137+ printf("Shape, dtype and values FAILED for [");
138+ for (size_t i = 0; i < s.size(); ++i) printf("%ld%s", s[i], i + 1 < s.size() ? "," : "");
139+ printf("]\n");
140+ continue;
141+ }
142+ if (VerifyOutput(output[0], s, inputValue, tolerance)) {
143+ printf("Shape, dtype and values PASSED for [");
144+ for (size_t i = 0; i < s.size(); ++i) printf("%ld%s", s[i], i + 1 < s.size() ? "," : "");
145+ printf("]\n");
146+ passed++;
147+ } else {
148+ printf("Shape, dtype and values FAILED for [");
149+ for (size_t i = 0; i < s.size(); ++i) printf("%ld%s", s[i], i + 1 < s.size() ? "," : "");
150+ printf("]\n");
151+ }
152+ }
153+ printf("Scenario unknown_rank_minus_2 summary: %d/%d passed\n", passed, static_cast<int>(shapes.size()));
154+ return passed == static_cast<int>(shapes.size()) ? SUCCESS : FAILED;
155+}
156+ 
157+static Graph BuildDynamicGraph(const string& name, const vector<int64_t>& dynShape)
158+{
159+ Graph graph(name.c_str());
160+ TensorDesc dynDesc(ge::Shape(dynShape), FORMAT_ND, DT_FLOAT);
161+ dynDesc.SetFormat(FORMAT_ND);
162+ dynDesc.SetRealDimCnt(dynShape.size());
163+ 
164+ auto data_op = op::Data("placeholder_x").set_attr_index(0);
165+ data_op.update_input_desc_x(dynDesc);
166+ data_op.update_output_desc_y(dynDesc);
167+ 
168+ auto celu_op = op::Celu("celu");
169+ celu_op.set_input_x(data_op);
170+ celu_op.set_attr_alpha1(1.0f);
171+ celu_op.set_attr_alpha2(1.0f);
172+ celu_op.set_attr_alpha3(1.0f);
173+ TensorDesc outDesc(ge::Shape(dynShape), FORMAT_ND, DT_FLOAT);
174+ celu_op.update_output_desc_y(outDesc);
175+ 
176+ graph.AddOp(data_op);
177+ graph.AddOp(celu_op);
178+ vector<Operator> inputs{data_op};
179+ vector<Operator> outputs{celu_op};
180+ graph.SetInputs(inputs).SetOutputs(outputs);
181+ return graph;
182+}
183+ 
184+int main(int argc, char* argv[])
185+{
186+ (void)argc;
187+ (void)argv;
188+ 
189+ printf("%s - INFO - [XIR]: Start to initialize ge\n", GetTime().c_str());
190+ std::map<AscendString, AscendString> global_options = {
191+ {"ge.exec.deviceId", "1"}, {"ge.graphRunMode", "1"}};
192+ Status ret = ge::GEInitialize(global_options);
193+ if (ret != SUCCESS) {
194+ printf("Initialize ge failed\n");
195+ return FAILED;
196+ }
197+ 
198+ ge::Session* session = new Session(std::map<AscendString, AscendString>{});
199+ if (session == nullptr) {
200+ printf("Create session failed\n");
201+ GEFinalize();
202+ return FAILED;
203+ }
204+ 
205+ Graph dimGraph = BuildDynamicGraph("celu_dyn_dim", {-1, -1});
206+ uint32_t graph_id1 = 0;
207+ ret = session->AddGraph(graph_id1, dimGraph, std::map<AscendString, AscendString>{});
208+ if (ret != SUCCESS) {
209+ printf("AddGraph failed\n");
210+ delete session;
211+ GEFinalize();
212+ return FAILED;
213+ }
214+ printf("AddGraph success (dynamic shape [-1,-1])\n");
215+ 
216+ int dimResult = RunScenario(session, graph_id1, {{4, 2}, {1, 8}, {3, 5}}, "unknown_dim_minus_1");
217+ 
218+ Graph rankGraph = BuildDynamicGraph("celu_dyn_rank", {-2});
219+ uint32_t graph_id2 = 1;
220+ ret = session->AddGraph(graph_id2, rankGraph, std::map<AscendString, AscendString>{});
221+ if (ret != SUCCESS) {
222+ printf("AddGraph2 failed\n");
223+ delete session;
224+ GEFinalize();
225+ return FAILED;
226+ }
227+ printf("AddGraph2 success (dynamic rank [-2])\n");
228+ 
229+ int rankResult = RunRankScenario(session, graph_id2, {{8}, {4, 2}, {2, 3, 4}});
230+ 
231+ delete session;
232+ GEFinalize();
233+ 
234+ if (dimResult == SUCCESS && rankResult == SUCCESS) {
235+ printf("Celu dynamic GEIR verification PASSED (-1 and -2)\n");
236+ return SUCCESS;
237+ }
238+ printf("Celu dynamic GEIR verification FAILED (-1 and -2)\n");
239+ return FAILED;
240+}
Rindex/data_format_dim_map/op_host/data_format_dim_map_infershape.cpp→activation/celu/op_graph/celu_graph_infer.cpp+17-14
@@ -3,33 +3,36 @@
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS FILE 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.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.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11/* Generated By CANNBot */11/* Generated By CANNBot */
12 12 
13+/*!
14+ * \file celu_graph_infer.cpp
15+ * \brief Celu operator graph infer resource
16+ */
13#include "register/op_impl_registry.h"17#include "register/op_impl_registry.h"
14-#include "exe_graph/runtime/infer_shape_context.h"18+#include "log/log.h"
15-#include "op_common/log/log.h"
16- 
17-using namespace ge;
18 19 
19namespace ops {20namespace ops {
21+using namespace ge;
20 22 
21-static ge::graphStatus InferShape4DataFormatDimMap(gert::InferShapeContext* context)23+static constexpr int64_t IDX_0 = 0;
24+ 
25+static ge::graphStatus InferDataTypeCelu(gert::InferDataTypeContext* context)
22{26{
23- const gert::Shape* inputShape = context->GetInputShape(0);27+ OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeCelu");
24- OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
25 28 
26- gert::Shape* outputShape = context->GetOutputShape(0);29+ ge::DataType inputDtype = context->GetInputDataType(IDX_0);
27- OP_CHECK_NULL_WITH_CONTEXT(context, outputShape);30+ context->SetOutputDataType(IDX_0, inputDtype);
28 31 
29- *outputShape = *inputShape;32+ OP_LOGD(context->GetNodeName(), "End to do InferDataTypeCelu");
30- return ge::GRAPH_SUCCESS;33+ return GRAPH_SUCCESS;
31}34}
32 35 
33-IMPL_OP_INFERSHAPE(DataFormatDimMap).InferShape(InferShape4DataFormatDimMap);36+IMPL_OP(Celu).InferDataType(InferDataTypeCelu);
34 37 
35-} // namespace ops38+}; // namespace ops
Mactivation/celu/op_host/arch35/celu_tiling_arch35.cpp+9-5
@@ -62,6 +62,10 @@ static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t&
62 auto inputX = context->GetInputShape(0);62 auto inputX = context->GetInputShape(0);
63 OP_CHECK_NULL_WITH_CONTEXT(context, inputX);63 OP_CHECK_NULL_WITH_CONTEXT(context, inputX);
64 auto inputShapeX = EnsureNotScalar(inputX->GetStorageShape());64 auto inputShapeX = EnsureNotScalar(inputX->GetStorageShape());
65+ OP_CHECK_IF(inputShapeX.GetDimNum() > 8,
66+ OP_LOGE(context, "Celu: input dimensions %zu exceeds platform limit 8",
67+ inputShapeX.GetDimNum()),
68+ return ge::GRAPH_FAILED);
65 auto outY = context->GetOutputShape(0);69 auto outY = context->GetOutputShape(0);
66 OP_CHECK_NULL_WITH_CONTEXT(context, outY);70 OP_CHECK_NULL_WITH_CONTEXT(context, outY);
67 auto outShapeY = EnsureNotScalar(outY->GetStorageShape());71 auto outShapeY = EnsureNotScalar(outY->GetStorageShape());
@@ -76,7 +80,7 @@ static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t&
76 OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);80 OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
77 dataType = inputDesc->GetDataType();81 dataType = inputDesc->GetDataType();
78 if (supportedDtype.count(dataType) == 0) {82 if (supportedDtype.count(dataType) == 0) {
79- OP_LOGE(context, "invalid dtype");83+ OP_LOGE(context, "unsupported dtype: %d", static_cast<int>(dataType));
80 return ge::GRAPH_FAILED;84 return ge::GRAPH_FAILED;
81 }85 }
82 return ge::GRAPH_SUCCESS;86 return ge::GRAPH_SUCCESS;
@@ -142,8 +146,8 @@ static ge::graphStatus CeluTilingFunc(gert::TilingContext* context)
142 tiling->blockFactor = 0;146 tiling->blockFactor = 0;
143 tiling->ubFactor = 0;147 tiling->ubFactor = 0;
144 context->SetBlockDim(1);148 context->SetBlockDim(1);
145- uint32_t dType = static_cast<uint32_t>(dataType);149+ uint32_t mode = 0;
146- ASCENDC_TPL_SEL_PARAM(context, dType);150+ ASCENDC_TPL_SEL_PARAM(context, mode);
147 return ge::GRAPH_SUCCESS;151 return ge::GRAPH_SUCCESS;
148 }152 }
149 153 
@@ -174,8 +178,8 @@ static ge::graphStatus CeluTilingFunc(gert::TilingContext* context)
174 178 
175 context->SetBlockDim(usedCoreNum);179 context->SetBlockDim(usedCoreNum);
176 180 
177- uint32_t dType = static_cast<uint32_t>(dataType);181+ uint32_t mode = 0;
178- ASCENDC_TPL_SEL_PARAM(context, dType);182+ ASCENDC_TPL_SEL_PARAM(context, mode);
179 return ge::GRAPH_SUCCESS;183 return ge::GRAPH_SUCCESS;
180}184}
181 185 
Mactivation/celu/op_host/celu_def.cpp+1-1
@@ -39,7 +39,7 @@ public:
39 .DynamicRankSupportFlag(true)39 .DynamicRankSupportFlag(true)
40 .DynamicShapeSupportFlag(true)40 .DynamicShapeSupportFlag(true)
41 .NeedCheckSupportFlag(false)41 .NeedCheckSupportFlag(false)
42- .PrecisionReduceFlag(true)42+ .PrecisionReduceFlag(false)
43 .ExtendCfgInfo("opFile.value", "celu");43 .ExtendCfgInfo("opFile.value", "celu");
44 this->AICore().AddConfig("ascend950", aiCoreConfig);44 this->AICore().AddConfig("ascend950", aiCoreConfig);
45 }45 }
Mactivation/celu/op_kernel/arch35/celu_tiling_key.h+3-4
@@ -3,7 +3,7 @@
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS FILE 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.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.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
@@ -15,9 +15,8 @@
15 15 
16#include "ascendc/host_api/tiling/template_argument.h"16#include "ascendc/host_api/tiling/template_argument.h"
17 17 
18-ASCENDC_TPL_ARGS_DECL(Celu, ASCENDC_TPL_DATATYPE_DECL(D_T, C_DT_FLOAT16, C_DT_FLOAT, ASCENDC_TPL_INPUT(0)));18+ASCENDC_TPL_ARGS_DECL(Celu, ASCENDC_TPL_UINT_DECL(MODE, 8, ASCENDC_TPL_UI_LIST, 0));
19 19 
20-ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T, C_DT_FLOAT16)),20+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(MODE, ASCENDC_TPL_UI_LIST, 0)), );
21- ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T, C_DT_FLOAT)), );
22 21 
23#endif // __CELU_TILING_KEY_H__22#endif // __CELU_TILING_KEY_H__
Mactivation/celu/op_kernel/celu.cpp+3-3
@@ -3,7 +3,7 @@
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS FILE 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.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.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
@@ -12,12 +12,12 @@
12 12 
13#include "arch35/celu.h"13#include "arch35/celu.h"
14 14 
15-template <typename D_T>15+template <int MODE>
16__global__ __aicore__ void celu(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)16__global__ __aicore__ void celu(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
17{17{
18 REGISTER_TILING_DEFAULT(CeluTilingData);18 REGISTER_TILING_DEFAULT(CeluTilingData);
19 GET_TILING_DATA_WITH_STRUCT(CeluTilingData, tilingData, tiling);19 GET_TILING_DATA_WITH_STRUCT(CeluTilingData, tilingData, tiling);
20- NsCelu::Celu<D_T> op;20+ NsCelu::Celu<DTYPE_X> op;
21 op.Init(x, y, &tilingData);21 op.Init(x, y, &tilingData);
22 op.Process();22 op.Process();
23}23}
Mexperimental/index/data_format_dim_map/op_host/data_format_dim_map_tiling.cpp+2-2
@@ -141,8 +141,8 @@ static ge::graphStatus DataFormatDimMapTilingFunc(gert::TilingContext* context)
141 }141 }
142 142 
143 context->SetBlockDim(CeilDiv(totalIdx, tiling->blockFactor));143 context->SetBlockDim(CeilDiv(totalIdx, tiling->blockFactor));
144- uint32_t dTypeX = static_cast<uint32_t>(dataType);144+ uint32_t mode = 0;
145- ASCENDC_TPL_SEL_PARAM(context, dTypeX);145+ ASCENDC_TPL_SEL_PARAM(context, mode);
146 return ge::GRAPH_SUCCESS;146 return ge::GRAPH_SUCCESS;
147}147}
148 148 
Mexperimental/index/data_format_dim_map/op_kernel/data_format_dim_map.cpp+3-3
@@ -3,7 +3,7 @@
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS FILE 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.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.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
@@ -15,12 +15,12 @@
15 15 
16#include "data_format_dim_map.h"16#include "data_format_dim_map.h"
17 17 
18-template <typename D_T_X>18+template <int MODE>
19__global__ __aicore__ void data_format_dim_map(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)19__global__ __aicore__ void data_format_dim_map(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
20{20{
21 REGISTER_TILING_DEFAULT(DataFormatDimMapTilingData);21 REGISTER_TILING_DEFAULT(DataFormatDimMapTilingData);
22 GET_TILING_DATA_WITH_STRUCT(DataFormatDimMapTilingData, tilingData, tiling);22 GET_TILING_DATA_WITH_STRUCT(DataFormatDimMapTilingData, tilingData, tiling);
23- NsDataFormatDimMap::DataFormatDimMap<D_T_X> op;23+ NsDataFormatDimMap::DataFormatDimMap<DTYPE_X> op;
24 op.Init(x, y, &tilingData);24 op.Init(x, y, &tilingData);
25 op.Process();25 op.Process();
26}26}
Mexperimental/index/data_format_dim_map/op_kernel/data_format_dim_map_tiling_key.h+3-4
@@ -3,7 +3,7 @@
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS FILE 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.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.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
@@ -18,9 +18,8 @@
18 18 
19#include "ascendc/host_api/tiling/template_argument.h"19#include "ascendc/host_api/tiling/template_argument.h"
20 20 
21-ASCENDC_TPL_ARGS_DECL(DataFormatDimMap, ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_INT32, C_DT_INT64, ASCENDC_TPL_INPUT(0)));21+ASCENDC_TPL_ARGS_DECL(DataFormatDimMap, ASCENDC_TPL_UINT_DECL(MODE, 8, ASCENDC_TPL_UI_LIST, 0));
22 22 
23-ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_INT32)),23+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(MODE, ASCENDC_TPL_UI_LIST, 0)), );
24- ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_INT64)), );
25 24 
26#endif25#endif
Mexperimental/optim/apply_proximal_gradient_descent/op_host/apply_proximal_gradient_descent_tiling.cpp+2-8
@@ -135,10 +135,8 @@ static ge::graphStatus ApplyProximalGradientDescentTilingFunc(gert::TilingContex
135 tiling->blockFactor = 0;135 tiling->blockFactor = 0;
136 tiling->ubFactor = 0;136 tiling->ubFactor = 0;
137 context->SetBlockDim(1);137 context->SetBlockDim(1);
138- // 仍需选择一个合法 TilingKey，缺省选 FP32+DB 即可（不进 kernel）
139- uint32_t dTypeX = static_cast<uint32_t>(dataType);
140 uint64_t useDoubleBuffer = 1;138 uint64_t useDoubleBuffer = 1;
141- ASCENDC_TPL_SEL_PARAM(context, dTypeX, useDoubleBuffer);139+ ASCENDC_TPL_SEL_PARAM(context, useDoubleBuffer);
142 return ge::GRAPH_SUCCESS;140 return ge::GRAPH_SUCCESS;
143 }141 }
144 142 
@@ -169,11 +167,7 @@ static ge::graphStatus ApplyProximalGradientDescentTilingFunc(gert::TilingContex
169 context->SetBlockDim(usedCoreNum);167 context->SetBlockDim(usedCoreNum);
170 168 
171 // TilingKey 参数：D_T_X + BUFFER_MODE（与 ASCENDC_TPL_ARGS_DECL 顺序一致）169 // TilingKey 参数：D_T_X + BUFFER_MODE（与 ASCENDC_TPL_ARGS_DECL 顺序一致）
172- // 编码：tilingKey = dtype | (buffer_mode << 8)170+ ASCENDC_TPL_SEL_PARAM(context, useDoubleBuffer);
173- // FP32+SB=0 FP16+SB=1 FP32+DB=256 FP16+DB=257
174- // 与 binary 索引（_0/_1/_256/_257）一一对齐
175- uint32_t dTypeX = static_cast<uint32_t>(dataType);
176- ASCENDC_TPL_SEL_PARAM(context, dTypeX, useDoubleBuffer);
177 return ge::GRAPH_SUCCESS;171 return ge::GRAPH_SUCCESS;
178}172}
179 173 
Mexperimental/optim/apply_proximal_gradient_descent/op_kernel/apply_proximal_gradient_descent.cpp+5-13
@@ -3,37 +3,29 @@
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS FILE 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.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.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11-/**
12- * NOTE: Portions of this code were AI-generated and have been
13- * technically reviewed for functional accuracy and security
14- */
15- 
16/*!11/*!
17 * \file apply_proximal_gradient_descent.cpp12 * \file apply_proximal_gradient_descent.cpp
18 * \brief ApplyProximalGradientDescent Kernel 入口13 * \brief ApplyProximalGradientDescent Kernel 入口
19 *14 *
20- * 模板参数：15+ * dtype 由 def 驱动：DTYPE_VAR=var dtype（T），在入口处实例化 op 类的 T 模板实参。
21- * - D_T_X: 数据类型（由 ASCENDC_TPL_DATATYPE_DECL 定义）16+ * 模板参数 BUFFER_MODE: 缓冲模式（0=单缓冲, 1=双缓冲）
22- * - BUFFER_MODE: 缓冲模式（0=单缓冲, 1=双缓冲）
23- *
24- * 参数顺序（固定）：输入 (var, alpha, l1, l2, delta) → 输出 (varOut) → workspace → tiling
25 */17 */
26 18 
27#include "apply_proximal_gradient_descent.h"19#include "apply_proximal_gradient_descent.h"
28 20 
29-template <typename D_T_X, int BUFFER_MODE>21+template <int BUFFER_MODE>
30__global__ __aicore__ void apply_proximal_gradient_descent(GM_ADDR var, GM_ADDR alpha, GM_ADDR l1, GM_ADDR l2,22__global__ __aicore__ void apply_proximal_gradient_descent(GM_ADDR var, GM_ADDR alpha, GM_ADDR l1, GM_ADDR l2,
31 GM_ADDR delta, GM_ADDR varOut, GM_ADDR workspace,23 GM_ADDR delta, GM_ADDR varOut, GM_ADDR workspace,
32 GM_ADDR tiling)24 GM_ADDR tiling)
33{25{
34 REGISTER_TILING_DEFAULT(ApplyProximalGradientDescentTilingData);26 REGISTER_TILING_DEFAULT(ApplyProximalGradientDescentTilingData);
35 GET_TILING_DATA_WITH_STRUCT(ApplyProximalGradientDescentTilingData, tilingData, tiling);27 GET_TILING_DATA_WITH_STRUCT(ApplyProximalGradientDescentTilingData, tilingData, tiling);
36- NsApplyProximalGradientDescent::ApplyProximalGradientDescent<D_T_X, BUFFER_MODE> op;28+ NsApplyProximalGradientDescent::ApplyProximalGradientDescent<DTYPE_VAR, BUFFER_MODE> op;
37 op.Init(var, alpha, l1, l2, delta, varOut, &tilingData);29 op.Init(var, alpha, l1, l2, delta, varOut, &tilingData);
38 op.Process();30 op.Process();
39}31}
Mexperimental/optim/apply_proximal_gradient_descent/op_kernel/apply_proximal_gradient_descent_tiling_key.h+4-14
@@ -3,23 +3,17 @@
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS FILE 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.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.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11-/**
12- * NOTE: Portions of this code were AI-generated and have been
13- * technically reviewed for functional accuracy and security
14- */
15- 
16/*!11/*!
17 * \file apply_proximal_gradient_descent_tiling_key.h12 * \file apply_proximal_gradient_descent_tiling_key.h
18 * \brief Tiling 模板参数定义13 * \brief Tiling 模板参数定义
19 *14 *
20- * 模板参数：15+ * dtype 由 def 驱动（构建系统按 def 的 DataType profile 生成 DTYPE_VAR 宏），
21- * - D_T_X: var 的数据类型（C_DT_FLOAT / C_DT_FLOAT16）16+ * tiling_key 不再编码 dtype 维度，只保留 BUFFER_MODE。
22- * - BUFFER_MODE: 缓冲模式（0=单缓冲, 1=双缓冲）
23 */17 */
24 18 
25#ifndef __APPLY_PROXIMAL_GRADIENT_DESCENT_TILING_KEY_H__19#ifndef __APPLY_PROXIMAL_GRADIENT_DESCENT_TILING_KEY_H__
@@ -28,12 +22,8 @@
28#include "ascendc/host_api/tiling/template_argument.h"22#include "ascendc/host_api/tiling/template_argument.h"
29 23 
30ASCENDC_TPL_ARGS_DECL(ApplyProximalGradientDescent,24ASCENDC_TPL_ARGS_DECL(ApplyProximalGradientDescent,
31- ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_FLOAT, C_DT_FLOAT16, ASCENDC_TPL_INPUT(0)),
32 ASCENDC_TPL_UINT_DECL(BUFFER_MODE, 8, ASCENDC_TPL_UI_LIST, 0, 1));25 ASCENDC_TPL_UINT_DECL(BUFFER_MODE, 8, ASCENDC_TPL_UI_LIST, 0, 1));
33 26 
34-ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT),27+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(BUFFER_MODE, ASCENDC_TPL_UI_LIST, 0, 1)), );
35- ASCENDC_TPL_UINT_SEL(BUFFER_MODE, ASCENDC_TPL_UI_LIST, 0, 1)),
36- ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT16),
37- ASCENDC_TPL_UINT_SEL(BUFFER_MODE, ASCENDC_TPL_UI_LIST, 0, 1)), );
38 28 
39#endif29#endif
Roptim/apply_proximal_gradient_descent/op_host/apply_proximal_gradient_descent_infershape.cpp→index/data_format_dim_map/op_graph/data_format_dim_map_graph_infer.cpp+17-22
@@ -3,7 +3,7 @@
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS FILE 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.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.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
@@ -11,33 +11,28 @@
11/* Generated By CANNBot */11/* Generated By CANNBot */
12 12 
13/*!13/*!
14- * \file apply_proximal_gradient_descent_infershape.cpp14+ * \file data_format_dim_map_graph_infer.cpp
15- * \brief ApplyProximalGradientDescent 算子形状推导15+ * \brief DataFormatDimMap operator graph infer resource
16- *
17- * 逻辑：varOut.shape = var.shape, varOut.dtype = var.dtype
18 */16 */
19- 
20#include "register/op_impl_registry.h"17#include "register/op_impl_registry.h"
21-#include "exe_graph/runtime/infer_shape_context.h"18+#include "log/log.h"
22- 
23-using namespace ge;
24 19 
25namespace ops {20namespace ops {
21+using namespace ge;
26 22 
27-static ge::graphStatus InferShape4ApplyProximalGradientDescent(gert::InferShapeContext* context)23+static constexpr int64_t IDX_0 = 0;
24+ 
25+static ge::graphStatus InferDataTypeDataFormatDimMap(gert::InferDataTypeContext* context)
28{26{
29- const gert::Shape* varShape = context->GetInputShape(0);27+ OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeDataFormatDimMap");
30- if (varShape == nullptr) {28+ 
31- return ge::GRAPH_FAILED;29+ ge::DataType inputDtype = context->GetInputDataType(IDX_0);
32- }30+ context->SetOutputDataType(IDX_0, inputDtype);
33- gert::Shape* outShape = context->GetOutputShape(0);31+ 
34- if (outShape == nullptr) {32+ OP_LOGD(context->GetNodeName(), "End to do InferDataTypeDataFormatDimMap");
35- return ge::GRAPH_FAILED;33+ return GRAPH_SUCCESS;
36- }
37- *outShape = *varShape;
38- return ge::GRAPH_SUCCESS;
39}34}
40 35 
41-IMPL_OP_INFERSHAPE(ApplyProximalGradientDescent).InferShape(InferShape4ApplyProximalGradientDescent);36+IMPL_OP(DataFormatDimMap).InferDataType(InferDataTypeDataFormatDimMap);
42 37 
43-} // namespace ops38+}; // namespace ops
Mindex/data_format_dim_map/op_host/arch35/data_format_dim_map_tiling_arch35.cpp+13-4
@@ -64,9 +64,18 @@ static ge::graphStatus GetInputInfo(gert::TilingContext* context, int64_t& total
64 auto inputShapePtr = context->GetInputShape(0);64 auto inputShapePtr = context->GetInputShape(0);
65 OP_CHECK_NULL_WITH_CONTEXT(context, inputShapePtr);65 OP_CHECK_NULL_WITH_CONTEXT(context, inputShapePtr);
66 totalIdx = EnsureNotScalar(inputShapePtr->GetStorageShape()).GetShapeSize();66 totalIdx = EnsureNotScalar(inputShapePtr->GetStorageShape()).GetShapeSize();
67+ OP_CHECK_IF(EnsureNotScalar(inputShapePtr->GetStorageShape()).GetDimNum() > 8,
68+ OP_LOGE(context, "DataFormatDimMap: input dimensions %zu exceeds platform limit 8",
69+ EnsureNotScalar(inputShapePtr->GetStorageShape()).GetDimNum()),
70+ return ge::GRAPH_FAILED);
67 auto inputDesc = context->GetInputDesc(0);71 auto inputDesc = context->GetInputDesc(0);
68 OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);72 OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
69 dataType = inputDesc->GetDataType();73 dataType = inputDesc->GetDataType();
74+ const std::set<ge::DataType> supportedDtype = {ge::DT_INT32, ge::DT_INT64};
75+ OP_CHECK_IF(supportedDtype.count(dataType) == 0,
76+ OP_LOGE(context, "DataFormatDimMap: unsupported dtype=%d (allowed: INT32/INT64)",
77+ static_cast<int>(dataType)),
78+ return ge::GRAPH_FAILED);
70 return ge::GRAPH_SUCCESS;79 return ge::GRAPH_SUCCESS;
71}80}
72 81 
@@ -124,8 +133,8 @@ static ge::graphStatus DataFormatDimMapTilingFunc(gert::TilingContext* context)
124 memset_s(tilingZero, sizeof(DataFormatDimMapTilingData), 0, sizeof(DataFormatDimMapTilingData)) != EOK,133 memset_s(tilingZero, sizeof(DataFormatDimMapTilingData), 0, sizeof(DataFormatDimMapTilingData)) != EOK,
125 OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);134 OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
126 context->SetBlockDim(1);135 context->SetBlockDim(1);
127- uint32_t dTypeXZero = static_cast<uint32_t>(dataType);136+ uint32_t modeZero = 0;
128- ASCENDC_TPL_SEL_PARAM(context, dTypeXZero);137+ ASCENDC_TPL_SEL_PARAM(context, modeZero);
129 return ge::GRAPH_SUCCESS;138 return ge::GRAPH_SUCCESS;
130 }139 }
131 140 
@@ -152,8 +161,8 @@ static ge::graphStatus DataFormatDimMapTilingFunc(gert::TilingContext* context)
152 }161 }
153 162 
154 context->SetBlockDim(CeilDiv(totalIdx, tiling->blockFactor));163 context->SetBlockDim(CeilDiv(totalIdx, tiling->blockFactor));
155- uint32_t dTypeX = static_cast<uint32_t>(dataType);164+ uint32_t mode = 0;
156- ASCENDC_TPL_SEL_PARAM(context, dTypeX);165+ ASCENDC_TPL_SEL_PARAM(context, mode);
157 return ge::GRAPH_SUCCESS;166 return ge::GRAPH_SUCCESS;
158}167}
159 168 
Mindex/data_format_dim_map/op_kernel/arch35/data_format_dim_map_tiling_key.h+7-5
@@ -3,21 +3,23 @@
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS FILE 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.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.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11-/* Generated By CANNBot */11+/**
12+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
12 15 
13#ifndef __DATA_FORMAT_DIM_MAP_TILING_KEY_H__16#ifndef __DATA_FORMAT_DIM_MAP_TILING_KEY_H__
14#define __DATA_FORMAT_DIM_MAP_TILING_KEY_H__17#define __DATA_FORMAT_DIM_MAP_TILING_KEY_H__
15 18 
16#include "ascendc/host_api/tiling/template_argument.h"19#include "ascendc/host_api/tiling/template_argument.h"
17 20 
18-ASCENDC_TPL_ARGS_DECL(DataFormatDimMap, ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_INT32, C_DT_INT64, ASCENDC_TPL_INPUT(0)));21+ASCENDC_TPL_ARGS_DECL(DataFormatDimMap, ASCENDC_TPL_UINT_DECL(MODE, 8, ASCENDC_TPL_UI_LIST, 0));
19 22 
20-ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_INT32)),23+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(MODE, ASCENDC_TPL_UI_LIST, 0)), );
21- ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_INT64)), );
22 24 
23#endif25#endif
Mindex/data_format_dim_map/op_kernel/data_format_dim_map_apt.cpp+2-2
@@ -12,12 +12,12 @@
12 12 
13#include "arch35/data_format_dim_map.h"13#include "arch35/data_format_dim_map.h"
14 14 
15-template <typename D_T_X>15+template <int MODE>
16__global__ __aicore__ void data_format_dim_map(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)16__global__ __aicore__ void data_format_dim_map(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
17{17{
18 REGISTER_TILING_DEFAULT(DataFormatDimMapTilingData);18 REGISTER_TILING_DEFAULT(DataFormatDimMapTilingData);
19 GET_TILING_DATA_WITH_STRUCT(DataFormatDimMapTilingData, tilingData, tiling);19 GET_TILING_DATA_WITH_STRUCT(DataFormatDimMapTilingData, tilingData, tiling);
20- NsDataFormatDimMap::DataFormatDimMap<D_T_X> op;20+ NsDataFormatDimMap::DataFormatDimMap<DTYPE_X> op;
21 op.Init(x, y, &tilingData);21 op.Init(x, y, &tilingData);
22 op.Process();22 op.Process();
23}23}
Aoptim/apply_proximal_gradient_descent/op_graph/apply_proximal_gradient_descent_graph_infer.cpp+38-0
@@ -0,0 +1,38 @@
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 FILE 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 apply_proximal_gradient_descent_graph_infer.cpp
15+ * \brief ApplyProximalGradientDescent operator graph infer resource
16+ */
17+#include "register/op_impl_registry.h"
18+#include "log/log.h"
19+ 
20+namespace ops {
21+using namespace ge;
22+ 
23+static constexpr int64_t IDX_0 = 0;
24+ 
25+static ge::graphStatus InferDataTypeApplyProximalGradientDescent(gert::InferDataTypeContext* context)
26+{
27+ OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeApplyProximalGradientDescent");
28+ 
29+ ge::DataType inputDtype = context->GetInputDataType(IDX_0);
30+ context->SetOutputDataType(IDX_0, inputDtype);
31+ 
32+ OP_LOGD(context->GetNodeName(), "End to do InferDataTypeApplyProximalGradientDescent");
33+ return GRAPH_SUCCESS;
34+}
35+ 
36+IMPL_OP(ApplyProximalGradientDescent).InferDataType(InferDataTypeApplyProximalGradientDescent);
37+ 
38+}; // namespace ops
Moptim/apply_proximal_gradient_descent/op_host/arch35/apply_proximal_gradient_descent_tiling_arch35.cpp+20-10
@@ -66,6 +66,10 @@ static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t&
66 auto varStorage = context->GetInputShape(0);66 auto varStorage = context->GetInputShape(0);
67 OP_CHECK_NULL_WITH_CONTEXT(context, varStorage);67 OP_CHECK_NULL_WITH_CONTEXT(context, varStorage);
68 auto varShape = EnsureNotScalar(varStorage->GetStorageShape());68 auto varShape = EnsureNotScalar(varStorage->GetStorageShape());
69+ OP_CHECK_IF(varShape.GetDimNum() > 8,
70+ OP_LOGE(context, "ApplyProximalGradientDescent: input dimensions %zu exceeds platform limit 8",
71+ varShape.GetDimNum()),
72+ return ge::GRAPH_FAILED);
69 73 
70 auto deltaStorage = context->GetInputShape(4);74 auto deltaStorage = context->GetInputShape(4);
71 OP_CHECK_NULL_WITH_CONTEXT(context, deltaStorage);75 OP_CHECK_NULL_WITH_CONTEXT(context, deltaStorage);
@@ -89,9 +93,22 @@ static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t&
89 OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);93 OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
90 dataType = inputDesc->GetDataType();94 dataType = inputDesc->GetDataType();
91 if (supportedDtype.count(dataType) == 0) {95 if (supportedDtype.count(dataType) == 0) {
92- OP_LOGE(context, "invalid dtype: %d", static_cast<int>(dataType));96+ OP_LOGE(context, "unsupported dtype: %d", static_cast<int>(dataType));
93 return ge::GRAPH_FAILED;97 return ge::GRAPH_FAILED;
94 }98 }
99+ for (int32_t i = 1; i <= 4; i++) {
100+ auto descI = context->GetInputDesc(static_cast<size_t>(i));
101+ OP_CHECK_NULL_WITH_CONTEXT(context, descI);
102+ auto dtI = descI->GetDataType();
103+ OP_CHECK_IF(supportedDtype.count(dtI) == 0,
104+ OP_LOGE(context, "ApplyProximalGradientDescent: input %d unsupported dtype=%d",
105+ i, static_cast<int>(dtI)),
106+ return ge::GRAPH_FAILED);
107+ OP_CHECK_IF(dtI != dataType,
108+ OP_LOGE(context, "ApplyProximalGradientDescent: dtype mismatch var=%d vs input%d=%d",
109+ static_cast<int>(dataType), i, static_cast<int>(dtI)),
110+ return ge::GRAPH_FAILED);
111+ }
95 return ge::GRAPH_SUCCESS;112 return ge::GRAPH_SUCCESS;
96}113}
97 114 
@@ -133,10 +150,8 @@ static ge::graphStatus ApplyProximalGradientDescentTilingFunc(gert::TilingContex
133 tiling->blockFactor = 0;150 tiling->blockFactor = 0;
134 tiling->ubFactor = 0;151 tiling->ubFactor = 0;
135 context->SetBlockDim(1);152 context->SetBlockDim(1);
136- // 仍需选择一个合法 TilingKey，缺省选 FP32+DB 即可（不进 kernel）
137- uint32_t dTypeX = static_cast<uint32_t>(dataType);
138 uint64_t useDoubleBuffer = 1;153 uint64_t useDoubleBuffer = 1;
139- ASCENDC_TPL_SEL_PARAM(context, dTypeX, useDoubleBuffer);154+ ASCENDC_TPL_SEL_PARAM(context, useDoubleBuffer);
140 return ge::GRAPH_SUCCESS;155 return ge::GRAPH_SUCCESS;
141 }156 }
142 157 
@@ -167,12 +182,7 @@ static ge::graphStatus ApplyProximalGradientDescentTilingFunc(gert::TilingContex
167 182 
168 context->SetBlockDim(usedCoreNum);183 context->SetBlockDim(usedCoreNum);
169 184 
170- // TilingKey 参数：D_T_X + BUFFER_MODE（与 ASCENDC_TPL_ARGS_DECL 顺序一致）185+ ASCENDC_TPL_SEL_PARAM(context, useDoubleBuffer);
171- // 编码：tilingKey = dtype | (buffer_mode << 8)
172- // FP32+SB=0 FP16+SB=1 FP32+DB=256 FP16+DB=257
173- // 与 binary 索引（_0/_1/_256/_257）一一对齐
174- uint32_t dTypeX = static_cast<uint32_t>(dataType);
175- ASCENDC_TPL_SEL_PARAM(context, dTypeX, useDoubleBuffer);
176 return ge::GRAPH_SUCCESS;186 return ge::GRAPH_SUCCESS;
177}187}
178 188 
Moptim/apply_proximal_gradient_descent/op_kernel/apply_proximal_gradient_descent.cpp+2-2
@@ -23,14 +23,14 @@ extern "C" __global__ __aicore__ void apply_proximal_gradient_descent(GM_ADDR va
23 op.Process();23 op.Process();
24}24}
25#else25#else
26-template <typename D_T_X, int BUFFER_MODE>26+template <int BUFFER_MODE>
27__global__ __aicore__ void apply_proximal_gradient_descent(GM_ADDR var, GM_ADDR alpha, GM_ADDR l1, GM_ADDR l2,27__global__ __aicore__ void apply_proximal_gradient_descent(GM_ADDR var, GM_ADDR alpha, GM_ADDR l1, GM_ADDR l2,
28 GM_ADDR delta, GM_ADDR varOut, GM_ADDR workspace,28 GM_ADDR delta, GM_ADDR varOut, GM_ADDR workspace,
29 GM_ADDR tiling)29 GM_ADDR tiling)
30{30{
31 REGISTER_TILING_DEFAULT(ApplyProximalGradientDescentTilingData);31 REGISTER_TILING_DEFAULT(ApplyProximalGradientDescentTilingData);
32 GET_TILING_DATA_WITH_STRUCT(ApplyProximalGradientDescentTilingData, tilingData, tiling);32 GET_TILING_DATA_WITH_STRUCT(ApplyProximalGradientDescentTilingData, tilingData, tiling);
33- NsApplyProximalGradientDescent::ApplyProximalGradientDescent<D_T_X, BUFFER_MODE> op;33+ NsApplyProximalGradientDescent::ApplyProximalGradientDescent<DTYPE_VAR, BUFFER_MODE> op;
34 op.Init(var, alpha, l1, l2, delta, varOut, &tilingData);34 op.Init(var, alpha, l1, l2, delta, varOut, &tilingData);
35 op.Process();35 op.Process();
36}36}
Moptim/apply_proximal_gradient_descent/op_kernel/arch35/apply_proximal_gradient_descent_tiling_key.h+4-11
@@ -3,20 +3,17 @@
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS FILE 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.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.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11-/* Generated By CANNBot */
12- 
13/*!11/*!
14 * \file apply_proximal_gradient_descent_tiling_key.h12 * \file apply_proximal_gradient_descent_tiling_key.h
15 * \brief Tiling 模板参数定义13 * \brief Tiling 模板参数定义
16 *14 *
17- * 模板参数：15+ * dtype 由 def 驱动（构建系统按 def 的 DataType profile 生成 DTYPE_VAR 宏），
18- * - D_T_X: var 的数据类型（C_DT_FLOAT / C_DT_FLOAT16）16+ * tiling_key 不再编码 dtype 维度，只保留 BUFFER_MODE。
19- * - BUFFER_MODE: 缓冲模式（0=单缓冲, 1=双缓冲）
20 */17 */
21 18 
22#ifndef __APPLY_PROXIMAL_GRADIENT_DESCENT_TILING_KEY_H__19#ifndef __APPLY_PROXIMAL_GRADIENT_DESCENT_TILING_KEY_H__
@@ -25,12 +22,8 @@
25#include "ascendc/host_api/tiling/template_argument.h"22#include "ascendc/host_api/tiling/template_argument.h"
26 23 
27ASCENDC_TPL_ARGS_DECL(ApplyProximalGradientDescent,24ASCENDC_TPL_ARGS_DECL(ApplyProximalGradientDescent,
28- ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_FLOAT, C_DT_FLOAT16, ASCENDC_TPL_INPUT(0)),
29 ASCENDC_TPL_UINT_DECL(BUFFER_MODE, 8, ASCENDC_TPL_UI_LIST, 0, 1));25 ASCENDC_TPL_UINT_DECL(BUFFER_MODE, 8, ASCENDC_TPL_UI_LIST, 0, 1));
30 26 
31-ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT),27+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(BUFFER_MODE, ASCENDC_TPL_UI_LIST, 0, 1)), );
32- ASCENDC_TPL_UINT_SEL(BUFFER_MODE, ASCENDC_TPL_UI_LIST, 0, 1)),
33- ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT16),
34- ASCENDC_TPL_UINT_SEL(BUFFER_MODE, ASCENDC_TPL_UI_LIST, 0, 1)), );
35 28 
36#endif29#endif