Star1334
1451
代码
代码
Issues115
Pull Requests257
Actions
项目讨论
Wiki
分析
Star1334
1451
已合并
迁移TransData 950的tiling/kernel，调整op_api目录 #1169
马琦钧创建于 2月6日
迁移TransData 950的tiling/kernel，调整op_api目录 #1169
已合并
马琦钧创建于 2月6日
从maqijun/ops-mathmv_trans_data合入到cann/ops-mathmaster
共 32 个文件变更+1019-154
Mconversion/trans_data/CMakeLists.txt+5-7
@@ -9,10 +9,8 @@
9# See LICENSE in the root of the software repository for the full text of the License.9# See LICENSE in the root of the software repository for the full text of the License.
10# ----------------------------------------------------------------------------10# ----------------------------------------------------------------------------
11 11 
12-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)12+# 设置算子定义时支持的芯片类型
R
RRuiWang_2月25日

tests目录的一起调整

likedislike
马琦钧
马琦钧
2月25日 评论：
13-list(REMOVE_ITEM CURRENT_DIRS tests)13+set(SUPPORT_COMPUTE_UNIT "ascend950")
14-foreach(SUB_DIR ${CURRENT_DIRS})14+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
15- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")15+set(SUPPORT_TILING_DIR "arch35")
16- add_subdirectory(${SUB_DIR})16+add_all_modules_sources(OPTYPE trans_data ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
17- endif()
18-endforeach()
Rconversion/trans_data/op_host/op_api/aclnn_trans_convolution_weight.cpp→conversion/trans_data/op_api/aclnn_trans_convolution_weight.cpp+0-0
文件重命名但无更改。
Rconversion/trans_data/op_host/op_api/aclnn_trans_convolution_weight.h→conversion/trans_data/op_api/aclnn_trans_convolution_weight.h+0-0
文件重命名但无更改。
Rconversion/trans_data/op_host/op_api/aclnn_trans_matmul_weight.cpp→conversion/trans_data/op_api/aclnn_trans_matmul_weight.cpp+0-0
文件重命名但无更改。
Rconversion/trans_data/op_host/op_api/aclnn_trans_matmul_weight.h→conversion/trans_data/op_api/aclnn_trans_matmul_weight.h+0-0
文件重命名但无更改。
Rconversion/trans_data/op_host/op_api/transdata.cpp→conversion/trans_data/op_api/transdata.cpp+0-0
文件重命名但无更改。
Mconversion/trans_data/op_graph/trans_data_proto.h+1-0
@@ -82,6 +82,7 @@ namespace ge
82| NCHW ====> FRACTAL_Z_C04 | float16, bfloat16 | 16 | 1 |\n82| NCHW ====> FRACTAL_Z_C04 | float16, bfloat16 | 16 | 1 |\n
83| FRACTAL_Z_C04 ====> NCHW | float32 | 16 | 1 |\n83| FRACTAL_Z_C04 ====> NCHW | float32 | 16 | 1 |\n
84| ND ====> FRACTAL_NZ_C0_16 | float32, uint32, int32 | 16 | 1 |\n84| ND ====> FRACTAL_NZ_C0_16 | float32, uint32, int32 | 16 | 1 |\n
85+| ND ====> FRACTAL_NZ_C0_32 | float4_e2m1 | 32 | 1 |\n
85*\n86*\n
86*87*
87*/88*/
Aconversion/trans_data/op_host/arch35/transdata_tiling_arch35.cpp+267-0
@@ -0,0 +1,267 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file transdata_tiling_arch35.cpp
13+ * \brief calc tiling data for transdata AscendC kernel
14+ */
15+ 
16+#include "transdata_tiling_arch35.h"
17+ 
18+#include <algorithm>
19+ 
20+#include "log/log.h"
21+#include "op_host/tiling_util.h"
22+#include "util/math_util.h"
23+#include "util/platform_util.h"
24+ 
25+using namespace Ops::Math::OpTiling;
26+using namespace Ops::Base;
27+ 
28+namespace optiling {
29+namespace transdata_asc {
30+ 
31+ge::graphStatus TransDataTilingAscendC::GetHardwareInfo()
32+{
33+ auto compileInfo = reinterpret_cast<const TransDataCompileInfo*>(context_->GetCompileInfo());
34+ OP_CHECK_NULL_WITH_CONTEXT(context_, compileInfo);
35+ coreNum_ = static_cast<uint32_t>(compileInfo->coreNum);
36+ ubSize_ = compileInfo->ubSize;
37+ OP_CHECK_IF(
38+ (coreNum_ <= 0U || ubSize_ <= 0L),
39+ OP_LOGE(context_->GetNodeName(), "TransData GetHardwareInfo failed, core num: %u, ub size: %ld", coreNum_,
40+ ubSize_),
41+ return ge::GRAPH_FAILED);
42+ return ge::GRAPH_SUCCESS;
43+}
44+ 
45+void TransDataTilingAscendC::ReshapeInShape()
46+{
47+ auto dimCnt = inShape.GetDimNum();
48+ if (dimCnt > 1) {
49+ return;
50+ }
51+ auto bakDim = inShape.GetDim(0);
52+ inShape.SetDim(0, 1);
53+ inShape.AppendDim(1);
54+ inShape.AppendDim(bakDim);
55+}
56+ 
57+bool TransDataTilingAscendC::GetShapeInfo()
58+{
59+ auto xStorage = context_->GetInputShape(0);
60+ OP_CHECK_NULL_WITH_CONTEXT(context_, xStorage);
61+ inShape = EnsureNotScalar(xStorage->GetStorageShape());
62+ auto yStorage = context_->GetOutputShape(0);
63+ OP_CHECK_NULL_WITH_CONTEXT(context_, yStorage);
64+ outShape = EnsureNotScalar(yStorage->GetStorageShape());
65+ 
66+ OP_CHECK_IF(
67+ inShape.GetShapeSize() == 0 || outShape.GetShapeSize() == 0,
68+ OP_LOGE(context_->GetNodeName(), "The input or output shape is empty!"),
69+ return false);
70+ 
71+ ReshapeInShape();
72+ return true;
73+}
74+ 
75+bool TransDataTilingAscendC::GetTransFormatAndDType()
76+{
77+ auto srcTd = context_->GetInputDesc(0);
78+ auto dstTd = context_->GetOutputDesc(0);
79+ OP_CHECK_NULL_WITH_CONTEXT(context_, srcTd);
80+ OP_CHECK_NULL_WITH_CONTEXT(context_, dstTd);
81+ auto srcFormat = static_cast<ge::Format>(ge::GetPrimaryFormat(srcTd->GetStorageFormat()));
82+ dstFormat = static_cast<ge::Format>(ge::GetPrimaryFormat(dstTd->GetStorageFormat()));
83+ srcDtype_ = srcTd->GetDataType();
84+ dtypeSize = ge::GetSizeByDataType(srcDtype_);
85+ 
86+ std::vector<ge::Format> supportSrcFormat = {ge::FORMAT_ND, ge::FORMAT_NCL, ge::FORMAT_NCHW, ge::FORMAT_NHWC};
87+ std::vector<ge::Format> supportDstFormat = {ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ_C0_16,
88+ ge::FORMAT_FRACTAL_NZ_C0_32};
89+ OP_CHECK_IF(
90+ (std::find(supportSrcFormat.begin(), supportSrcFormat.end(), srcFormat) == supportSrcFormat.end() ||
91+ std::find(supportDstFormat.begin(), supportDstFormat.end(), dstFormat) == supportDstFormat.end()),
92+ OP_LOGE(context_->GetNodeName(), "The input or output format is invalid!"),
93+ return false);
94+ return true;
95+}
96+ 
97+bool TransDataTilingAscendC::CalcC0Size()
98+{
99+ auto dimCnt = outShape.GetDimNum();
100+ auto tmpC0 = outShape.GetDim(dimCnt - 1);
101+ 
102+ OP_CHECK_IF(
103+ dstFormat == ge::FORMAT_FRACTAL_NZ_C0_16 && tmpC0 != C0_16,
104+ OP_LOGE(context_->GetNodeName(), "The c0 should be 16 when dst format is FRACTAL_NZ_C0_16!"),
105+ return false);
106+ OP_CHECK_IF(
107+ dstFormat == ge::FORMAT_FRACTAL_NZ_C0_32 && tmpC0 != C0_32,
108+ OP_LOGE(context_->GetNodeName(), "The c0 should be 32 when dst format is FRACTAL_NZ_C0_32!"),
109+ return false);
110+ 
111+ int64_t expectC0 = tmpC0;
112+ if (dtypeSize == 1U) {
113+ expectC0 = C0_32;
114+ } else if (dtypeSize == nTwo) {
115+ expectC0 = C0_16;
116+ }
117+ OP_CHECK_IF(
118+ dstFormat == ge::FORMAT_FRACTAL_NZ && dtypeSize == nTwo * nTwo && C0_8 != expectC0 && C0_16 != expectC0,
119+ OP_LOGE(context_->GetNodeName(), "The c0 should be 8 or 16 when dst format is FRACTAL_NZ and dtype size is %zu!",
120+ dtypeSize),
121+ return false);
122+ OP_CHECK_IF(
123+ dstFormat == ge::FORMAT_FRACTAL_NZ && tmpC0 != expectC0,
124+ OP_LOGE(context_->GetNodeName(), "The c0 should be %ld when dst format is FRACTAL_NZ and dtype size is %zu!",
125+ expectC0, dtypeSize),
126+ return false);
127+ 
128+ c0_ = tmpC0;
129+ if (srcDtype_ == ge::DT_FLOAT4_E2M1) {
130+ // 当作u8处理，c0缩小一半
131+ c0_ = c0_ >> 1;
132+ }
133+ return true;
134+}
135+ 
136+void TransDataTilingAscendC::CalcHSize()
137+{
138+ int64_t res = 1;
139+ auto dimCnt = inShape.GetDimNum();
140+ if (dimCnt <= nTwo) {
141+ h_ = res;
142+ return;
143+ }
144+ 
145+ for (size_t i = 0; i < dimCnt - nTwo; i++) {
146+ res *= inShape.GetDim(i);
147+ }
148+ h_ = res;
149+}
150+ 
151+void TransDataTilingAscendC::CalcNCSize()
152+{
153+ auto dimCnt = inShape.GetDimNum();
154+ n_ = inShape.GetDim(dimCnt - nTwo);
155+ c_ = inShape.GetDim(dimCnt - 1);
156+ if (srcDtype_ == ge::DT_FLOAT4_E2M1) {
157+ // 当作u8处理，c缩小一半
158+ c_ = c_ >> 1;
159+ }
160+}
161+ 
162+void TransDataTilingAscendC::CalcTilingKey()
163+{
164+ int64_t ni = 16;
165+ int64_t shapeSize = h_ * CeilAlign(n_, ni) * CeilAlign(c_, c0_);
166+ tilingKey_ = (shapeSize > MAX_INT32_SIZE) ? TILING_MODE_SIMT_LARGE_SHAPE : TILING_MODE_SIMT;
167+}
168+ 
169+void TransDataTilingAscendC::CalcBlockAndThreadNum()
170+{
171+ bNum_ = coreNum_;
172+ tNum_ = (tilingKey_ == TILING_MODE_SIMT) ? tNum512 : tNum256;
173+}
174+ 
175+ge::graphStatus TransDataTilingAscendC::CalcTilingData()
176+{
177+ OP_CHECK_IF(!GetShapeInfo(), OP_LOGE(context_->GetNodeName(), "Failed to get shape info!"), return ge::GRAPH_FAILED);
178+ OP_CHECK_IF(
179+ !GetTransFormatAndDType(), OP_LOGE(context_->GetNodeName(), "Failed to get format and dtype info!"),
180+ return ge::GRAPH_FAILED);
181+ OP_CHECK_IF(!CalcC0Size(), OP_LOGE(context_->GetNodeName(), "Failed to get c0 size!"), return ge::GRAPH_FAILED);
182+ CalcHSize();
183+ CalcNCSize();
184+ CalcTilingKey();
185+ CalcBlockAndThreadNum();
186+ 
187+ return ge::GRAPH_SUCCESS;
188+}
189+ 
190+void TransDataTilingAscendC::WriteTilingData()
191+{
192+ context_->SetBlockDim(bNum_);
193+ context_->SetTilingKey(tilingKey_);
194+ context_->SetLocalMemorySize(ubSize_ - SIMT_RSV_SIZE);
195+ 
196+ tilingData_.set_c0(c0_);
197+ tilingData_.set_h(h_);
198+ tilingData_.set_n(n_);
199+ tilingData_.set_c(c_);
200+ tilingData_.set_tNum(tNum_);
201+ tilingData_.SaveToBuffer(context_->GetRawTilingData()->GetData(), context_->GetRawTilingData()->GetCapacity());
202+ context_->GetRawTilingData()->SetDataSize(tilingData_.GetDataSize());
203+}
204+ 
205+std::string TransDataTilingAscendC::PrintTilingData()
206+{
207+ std::string tilingStr;
208+ tilingStr += std::to_string(c0_) + ",";
209+ tilingStr += std::to_string(h_) + ",";
210+ tilingStr += std::to_string(n_) + ",";
211+ tilingStr += std::to_string(c_) + ",";
212+ tilingStr += std::to_string(tNum_);
213+ return tilingStr;
214+}
215+ 
216+ge::graphStatus TransDataTilingAscendC::DoTiling()
217+{
218+ OP_CHECK_IF(
219+ (CalcTilingData() != ge::GRAPH_SUCCESS),
220+ OP_LOGE(context_->GetNodeName(), "TransDataTilingAscendC failed to calc tiling data."),
221+ return ge::GRAPH_FAILED);
222+ OP_CHECK_IF(ubSize_ <= SIMT_RSV_SIZE, OP_LOGE(context_->GetNodeName(), "UB size too small for SIMT reserved size."),
223+ return ge::GRAPH_FAILED);
224+ WriteTilingData();
225+ size_t* currentWorkspace = context_->GetWorkspaceSizes(1);
226+ OP_CHECK_NULL_WITH_CONTEXT(context_, currentWorkspace);
227+ currentWorkspace[0] = kSyncWorkSpaceSize;
228+ OP_LOGI(context_->GetNodeName(), "TransData tiling data: %s", PrintTilingData().c_str());
229+ return ge::GRAPH_SUCCESS;
230+}
231+ 
232+} // namespace transdata_asc
233+ 
234+ge::graphStatus Tiling4TransDataAscendC(gert::TilingContext* context)
235+{
236+ transdata_asc::TransDataTilingAscendC tdTiling(context);
237+ OP_CHECK_IF(
238+ (tdTiling.GetHardwareInfo() != ge::GRAPH_SUCCESS),
239+ OP_LOGE(context->GetNodeName(), "TransDataTilingAscendC failed to get hardware info."),
240+ return ge::GRAPH_FAILED);
241+ return tdTiling.DoTiling();
242+}
243+ 
244+static ge::graphStatus TilingPrepare4TransData(gert::TilingParseContext* context)
245+{
246+ auto compileInfo = context->GetCompiledInfo<TransDataCompileInfo>();
247+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
248+ auto platformInfo = context->GetPlatformInfo();
249+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
250+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
251+ 
252+ compileInfo->coreNum = ascendcPlatform.GetCoreNumAiv();
253+ OP_CHECK_IF((compileInfo->coreNum <= 0), OP_LOGE(context->GetNodeName(), "The core num is invalid."),
254+ return ge::GRAPH_FAILED);
255+ 
256+ uint64_t ubSize = 0;
257+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
258+ compileInfo->ubSize = static_cast<int64_t>(ubSize);
259+ OP_CHECK_IF((compileInfo->ubSize <= 0), OP_LOGE(context->GetNodeName(), "Failed to get ub size."),
260+ return ge::GRAPH_FAILED);
261+ 
262+ return ge::GRAPH_SUCCESS;
263+}
264+ 
265+IMPL_OP_OPTILING(TransData).Tiling(Tiling4TransDataAscendC).TilingParse<TransDataCompileInfo>(TilingPrepare4TransData);
266+ 
267+} // namespace optiling
Aconversion/trans_data/op_host/arch35/transdata_tiling_arch35.h+98-0
@@ -0,0 +1,98 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file transdata_tiling_arch35.h
13+ * \brief transdata tiling for ascendC impl
14+ */
15+#ifndef OPS_MATH_CONVERSION_TRANSDATA_TILING_ARCH35_H_
16+#define OPS_MATH_CONVERSION_TRANSDATA_TILING_ARCH35_H_
17+ 
18+#include <cstdint>
19+#include <string>
20+ 
21+#include "register/op_impl_registry.h"
22+#include "register/tilingdata_base.h"
23+ 
24+namespace optiling {
25+ 
26+struct TransDataCompileInfo {
27+ int64_t coreNum;
28+ int64_t ubSize;
29+};
30+ 
31+BEGIN_TILING_DATA_DEF(TransDataASCTilingData)
32+TILING_DATA_FIELD_DEF(int64_t, c0);
33+TILING_DATA_FIELD_DEF(int64_t, h);
34+TILING_DATA_FIELD_DEF(int64_t, n);
35+TILING_DATA_FIELD_DEF(int64_t, c);
36+TILING_DATA_FIELD_DEF(int64_t, tNum); // thread number
37+END_TILING_DATA_DEF;
38+ 
39+REGISTER_TILING_DATA_CLASS(TransData, TransDataASCTilingData);
40+ 
41+ge::graphStatus Tiling4TransDataAscendC(gert::TilingContext* context);
42+ 
43+namespace transdata_asc {
44+constexpr int64_t TILING_MODE_SIMT = 21000;
45+constexpr int64_t TILING_MODE_SIMT_LARGE_SHAPE = 21001;
46+constexpr int64_t MAX_INT32_SIZE = 0x7fffffff;
47+constexpr size_t nTwo = 2;
48+constexpr size_t kSyncWorkSpaceSize = static_cast<size_t>(16) * 1024 * 1024;
49+constexpr int64_t tNum256 = 256;
50+constexpr int64_t tNum512 = 512;
51+constexpr int64_t C0_8 = 8;
52+constexpr int64_t C0_16 = 16;
53+constexpr int64_t C0_32 = 32;
54+constexpr int64_t SIMT_RSV_SIZE = 128 * 1024L;
55+ 
56+class TransDataTilingAscendC {
57+public:
58+ explicit TransDataTilingAscendC(gert::TilingContext* context) : context_(context){};
59+ ge::graphStatus DoTiling();
60+ ge::graphStatus GetHardwareInfo();
61+ 
62+private:
63+ ge::graphStatus CalcTilingData();
64+ bool GetTransFormatAndDType();
65+ bool GetShapeInfo();
66+ bool CalcC0Size();
67+ void CalcHSize();
68+ void CalcNCSize();
69+ void CalcBlockAndThreadNum();
70+ void ReshapeInShape();
71+ void CalcTilingKey();
72+ void WriteTilingData();
73+ std::string PrintTilingData();
74+ 
75+private:
76+ gert::TilingContext* context_ = nullptr;
77+ TransDataASCTilingData tilingData_;
78+ gert::Shape inShape;
79+ gert::Shape outShape;
80+ ge::Format dstFormat;
81+ size_t dtypeSize;
82+ ge::DataType srcDtype_;
83+ 
84+ uint32_t coreNum_{1};
85+ uint32_t bNum_;
86+ int64_t ubSize_;
87+ int64_t tilingKey_{TILING_MODE_SIMT};
88+ 
89+ int64_t c0_;
90+ int64_t h_;
91+ int64_t n_;
92+ int64_t c_;
93+ int64_t tNum_;
94+};
95+} // namespace transdata_asc
96+ 
97+} // namespace optiling
98+#endif // OPS_MATH_CONVERSION_TRANSDATA_TILING_ARCH35_H_
Aconversion/trans_data/op_host/config/ascend950/trans_data_binary.json+341-0
@@ -0,0 +1,341 @@
1+{
2+ "op_type": "TransData",
3+ "op_list": [
4+ {
5+ "bin_filename": "TransData_cb4c683189d13c2ea9554d47def0c52c",
6+ "inputs": [
7+ {
8+ "name": "src",
9+ "index": 0,
10+ "dtype": "int8",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "dtype_match_mode": "DtypeByte"
17+ }
18+ ],
19+ "outputs": [
20+ {
21+ "name": "dst",
22+ "index": 0,
23+ "dtype": "int8",
24+ "format": "FRACTAL_NZ",
25+ "paramType": "required",
26+ "shape": [
27+ -2
28+ ],
29+ "dtype_match_mode": "DtypeByte"
30+ }
31+ ],
32+ "attrs": [
33+ {
34+ "name": "src_format",
35+ "dtype": "string",
36+ "value": null
37+ },
38+ {
39+ "name": "dst_format",
40+ "dtype": "string",
41+ "value": null
42+ },
43+ {
44+ "name": "src_subformat",
45+ "dtype": "int",
46+ "value": null
47+ },
48+ {
49+ "name": "dst_subformat",
50+ "dtype": "int",
51+ "value": null
52+ },
53+ {
54+ "name": "groups",
55+ "dtype": "int",
56+ "value": null
57+ }
58+ ]
59+ },
60+ {
61+ "bin_filename": "TransData_cb5c683189d13c2ea9554d47def0c52c",
62+ "inputs": [
63+ {
64+ "name": "src",
65+ "index": 0,
66+ "dtype": "float16",
67+ "format": "ND",
68+ "paramType": "required",
69+ "shape": [
70+ -2
71+ ],
72+ "dtype_match_mode": "DtypeByte"
73+ }
74+ ],
75+ "outputs": [
76+ {
77+ "name": "dst",
78+ "index": 0,
79+ "dtype": "float16",
80+ "format": "FRACTAL_NZ",
81+ "paramType": "required",
82+ "shape": [
83+ -2
84+ ],
85+ "dtype_match_mode": "DtypeByte"
86+ }
87+ ],
88+ "attrs": [
89+ {
90+ "name": "src_format",
91+ "dtype": "string",
92+ "value": null
93+ },
94+ {
95+ "name": "dst_format",
96+ "dtype": "string",
97+ "value": null
98+ },
99+ {
100+ "name": "src_subformat",
101+ "dtype": "int",
102+ "value": null
103+ },
104+ {
105+ "name": "dst_subformat",
106+ "dtype": "int",
107+ "value": null
108+ },
109+ {
110+ "name": "groups",
111+ "dtype": "int",
112+ "value": null
113+ }
114+ ]
115+ },
116+ {
117+ "bin_filename": "TransData_1d1b8d1178d795fa2de4c04b3bf7394a",
118+ "inputs": [
119+ {
120+ "name": "src",
121+ "index": 0,
122+ "dtype": "float32",
123+ "format": "ND",
124+ "paramType": "required",
125+ "shape": [
126+ -2
127+ ],
128+ "dtype_match_mode": "DtypeByte"
129+ }
130+ ],
131+ "outputs": [
132+ {
133+ "name": "dst",
134+ "index": 0,
135+ "dtype": "float32",
136+ "format": "FRACTAL_NZ",
137+ "paramType": "required",
138+ "shape": [
139+ -2
140+ ],
141+ "dtype_match_mode": "DtypeByte"
142+ }
143+ ],
144+ "attrs": [
145+ {
146+ "name": "src_format",
147+ "dtype": "string",
148+ "value": null
149+ },
150+ {
151+ "name": "dst_format",
152+ "dtype": "string",
153+ "value": null
154+ },
155+ {
156+ "name": "src_subformat",
157+ "dtype": "int",
158+ "value": null
159+ },
160+ {
161+ "name": "dst_subformat",
162+ "dtype": "int",
163+ "value": null
164+ },
165+ {
166+ "name": "groups",
167+ "dtype": "int",
168+ "value": null
169+ }
170+ ]
171+ },
172+ {
173+ "bin_filename": "TransData_cb7c683189d13c2ea9554d47def0c52c",
174+ "inputs": [
175+ {
176+ "name": "src",
177+ "index": 0,
178+ "dtype": "float32",
179+ "format": "ND",
180+ "paramType": "required",
181+ "shape": [
182+ -2
183+ ],
184+ "dtype_match_mode": "DtypeByte"
185+ }
186+ ],
187+ "outputs": [
188+ {
189+ "name": "dst",
190+ "index": 0,
191+ "dtype": "float32",
192+ "format": "FRACTAL_NZ_C0_16",
193+ "paramType": "required",
194+ "shape": [
195+ -2
196+ ],
197+ "dtype_match_mode": "DtypeByte"
198+ }
199+ ],
200+ "attrs": [
201+ {
202+ "name": "src_format",
203+ "dtype": "string",
204+ "value": null
205+ },
206+ {
207+ "name": "dst_format",
208+ "dtype": "string",
209+ "value": null
210+ },
211+ {
212+ "name": "src_subformat",
213+ "dtype": "int",
214+ "value": null
215+ },
216+ {
217+ "name": "dst_subformat",
218+ "dtype": "int",
219+ "value": null
220+ },
221+ {
222+ "name": "groups",
223+ "dtype": "int",
224+ "value": null
225+ }
226+ ]
227+ },
228+ {
229+ "bin_filename": "TransData_cb8c254963d13c2ac9554d47def0c52c",
230+ "inputs": [
231+ {
232+ "name": "src",
233+ "index": 0,
234+ "dtype": "float32",
235+ "format": "ND",
236+ "paramType": "required",
237+ "shape": [
238+ -2
239+ ],
240+ "dtype_match_mode": "DtypeByte"
241+ }
242+ ],
243+ "outputs": [
244+ {
245+ "name": "dst",
246+ "index": 0,
247+ "dtype": "float32",
248+ "format": "FRACTAL_NZ_C0_32",
249+ "paramType": "required",
250+ "shape": [
251+ -2
252+ ],
253+ "dtype_match_mode": "DtypeByte"
254+ }
255+ ],
256+ "attrs": [
257+ {
258+ "name": "src_format",
259+ "dtype": "string",
260+ "value": null
261+ },
262+ {
263+ "name": "dst_format",
264+ "dtype": "string",
265+ "value": null
266+ },
267+ {
268+ "name": "src_subformat",
269+ "dtype": "int",
270+ "value": null
271+ },
272+ {
273+ "name": "dst_subformat",
274+ "dtype": "int",
275+ "value": null
276+ },
277+ {
278+ "name": "groups",
279+ "dtype": "int",
280+ "value": null
281+ }
282+ ]
283+ },
284+ {
285+ "bin_filename": "TransData_float4e2m1_ND_NZ_C0_32",
286+ "inputs": [
287+ {
288+ "name": "src",
289+ "index": 0,
290+ "dtype": "float4_e2m1",
291+ "format": "ND",
292+ "paramType": "required",
293+ "shape": [
294+ -2
295+ ],
296+ "dtype_match_mode": "DtypeByte"
297+ }
298+ ],
299+ "outputs": [
300+ {
301+ "name": "dst",
302+ "index": 0,
303+ "dtype": "float4_e2m1",
304+ "format": "FRACTAL_NZ_C0_32",
305+ "paramType": "required",
306+ "shape": [
307+ -2
308+ ],
309+ "dtype_match_mode": "DtypeByte"
310+ }
311+ ],
312+ "attrs": [
313+ {
314+ "name": "src_format",
315+ "dtype": "string",
316+ "value": null
317+ },
318+ {
319+ "name": "dst_format",
320+ "dtype": "string",
321+ "value": null
322+ },
323+ {
324+ "name": "src_subformat",
325+ "dtype": "int",
326+ "value": null
327+ },
328+ {
329+ "name": "dst_subformat",
330+ "dtype": "int",
331+ "value": null
332+ },
333+ {
334+ "name": "groups",
335+ "dtype": "int",
336+ "value": null
337+ }
338+ ]
339+ }
340+ ]
341+}
Aconversion/trans_data/op_host/config/ascend950/trans_data_simplified_key.ini+2-0
@@ -0,0 +1,2 @@
1+[TransData]
2+default=0
Aconversion/trans_data/op_host/trans_data_def.cpp+58-0
@@ -0,0 +1,58 @@
1+/**
2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file trans_data_def.cpp
13+ * \brief op store info of TransData
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops
18+{
19+class TransData : public OpDef
20+{
21+public:
22+ const std::vector<ge::DataType> dType = {ge::DT_INT8, ge::DT_UINT8, ge::DT_INT16, ge::DT_UINT16, ge::DT_FLOAT16,
23+ ge::DT_BF16, ge::DT_INT32, ge::DT_UINT32, ge::DT_FLOAT, ge::DT_FLOAT,
24+ ge::DT_FLOAT, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT4_E2M1};
25+ const std::vector<ge::Format> srcFormat = {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
26+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
27+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
28+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
29+ const std::vector<ge::Format> dstFormat = {
30+ ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ,
31+ ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ,
32+ ge::FORMAT_FRACTAL_NZ_C0_16, ge::FORMAT_FRACTAL_NZ_C0_16, ge::FORMAT_FRACTAL_NZ_C0_16, ge::FORMAT_FRACTAL_NZ,
33+ ge::FORMAT_FRACTAL_NZ_C0_32, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ_C0_32};
34+ 
35+ explicit TransData(const char* name) : OpDef(name)
36+ {
37+ this->Input("src").ParamType(REQUIRED).DataType(dType).Format(srcFormat).UnknownShapeFormat(srcFormat);
38+ this->Output("dst").ParamType(REQUIRED).DataType(dType).Format(dstFormat).UnknownShapeFormat(dstFormat);
39+ this->Attr("src_format").AttrType(REQUIRED).String();
40+ this->Attr("dst_format").AttrType(REQUIRED).String();
41+ this->Attr("src_subformat").AttrType(OPTIONAL).Int(0);
42+ this->Attr("dst_subformat").AttrType(OPTIONAL).Int(0);
43+ this->Attr("groups").AttrType(OPTIONAL).Int(1);
44+ 
45+ OpAICoreConfig aicore_config;
46+ aicore_config.DynamicCompileStaticFlag(true)
47+ .DynamicFormatFlag(false)
48+ .DynamicRankSupportFlag(true)
49+ .DynamicShapeSupportFlag(true)
50+ .NeedCheckSupportFlag(false)
51+ .ExtendCfgInfo("opFile.value", "trans_data_apt");
52+ 
53+ this->AICore().AddConfig("ascend950", aicore_config);
54+ }
55+};
56+ 
57+OP_ADD(TransData);
58+} // namespace ops
Aconversion/trans_data/op_kernel/arch35/trans_data.h+62-0
@@ -0,0 +1,62 @@
1+/**
2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file trans_data.h
13+ * \brief expose trans_data implementation
14+ */
15+ 
16+#ifndef TRANS_DATA_HEAD_FILE__H_
17+#define TRANS_DATA_HEAD_FILE__H_
18+ 
19+#include <type_traits>
20+ 
21+#include "trans_data_with_simt.h"
22+ 
23+using namespace AscendC;
24+using namespace TRSD;
25+ 
26+#define TILING_MODE_SIMT 21000
27+#define TILING_MODE_SIMT_LARGE_SHAPE 21001
28+ 
29+__aicore__ void inline trans_data_impl(GM_ADDR src, GM_ADDR dst, GM_ADDR workspace, GM_ADDR tiling)
30+{
31+ if (workspace == nullptr) {
32+ return;
33+ }
34+ SetSysWorkspace(workspace);
35+ 
36+ constexpr auto b8 = sizeof(uint8_t);
37+ constexpr auto b16 = sizeof(uint16_t);
38+ constexpr auto b32 = sizeof(uint32_t);
39+ constexpr auto b64 = sizeof(uint64_t);
40+ constexpr auto tSize = sizeof(DTYPE_SRC);
41+ using DTYPE_SRC_ = std::conditional_t<
42+ tSize != b32,
43+ std::conditional_t<
44+ tSize == b8, uint8_t,
45+ std::conditional_t<tSize == b16, uint16_t, std::conditional_t<tSize == b64, uint64_t, DTYPE_SRC>>>,
46+ DTYPE_SRC>;
47+ 
48+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
49+ GET_TILING_DATA(tilingData, tiling);
50+ 
51+ if (TILING_KEY_IS(TILING_MODE_SIMT)) {
52+ TransWithSIMT<DTYPE_SRC_> op;
53+ op.Init(src, dst, &tilingData);
54+ op.Process<uint32_t>();
55+ } else if (TILING_KEY_IS(TILING_MODE_SIMT_LARGE_SHAPE)) {
56+ TransWithSIMT<DTYPE_SRC_> op;
57+ op.Init(src, dst, &tilingData);
58+ op.Process<uint64_t>();
59+ }
60+}
61+ 
62+#endif // TRANS_DATA_HEAD_FILE__H_
Aconversion/trans_data/op_kernel/arch35/trans_data_with_simt.h+116-0
@@ -0,0 +1,116 @@
1+/**
2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file trans_data_with_simt.h
13+ * \brief kernel of trans_data with simt
14+ */
15+ 
16+#ifndef TRANS_DATA_WITH_SIMT_IMP_H_
17+#define TRANS_DATA_WITH_SIMT_IMP_H_
18+ 
19+#include "kernel_operator.h"
20+#include "op_kernel/math_util.h"
21+ 
22+namespace TRSD
23+{
24+using namespace AscendC;
25+ 
26+constexpr size_t THREAD_BOUND = 2048;
27+ 
28+template <typename T>
29+class TransWithSIMT
30+{
31+public:
32+ __aicore__ inline TransWithSIMT(){};
33+ __aicore__ inline void Init(GM_ADDR src, GM_ADDR dst, const TransDataASCTilingData* tilingDataPtr);
34+ template <typename U>
35+ __aicore__ inline void Process();
36+ 
37+private:
38+ GlobalTensor<T> inGM;
39+ GlobalTensor<T> outGM;
40+ const TransDataASCTilingData* tdPtr = nullptr;
41+};
42+ 
43+template <typename T>
44+__aicore__ inline void TransWithSIMT<T>::Init(GM_ADDR src, GM_ADDR dst, const TransDataASCTilingData* tilingDataPtr)
45+{
46+ inGM.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(src));
47+ outGM.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(dst));
48+ tdPtr = tilingDataPtr;
49+}
50+ 
51+template <typename T, typename U>
52+__simt_vf__ LAUNCH_BOUND(THREAD_BOUND / sizeof(U)) __aicore__
53+ void SIMTTrans(__gm__ T* dst, __gm__ T* src, uint64_t shapeSize, U c1, U padN, U c0, U oriN, U oriC, U mPNC, U sPNC,
54+ U mPNC0, U sPNC0, U mC1, U sC1, U mC0, U sC0, U mPN, U sPN)
55+{
56+ uint64_t tNum = uint64_t(Simt::GetThreadNum());
57+ uint64_t blockID = uint64_t(Simt::GetBlockIdx());
58+ uint64_t bNum = uint64_t(Simt::GetBlockNum());
59+ U hIdx = 0;
60+ U c1Idx = 0;
61+ U nIdx = 0;
62+ U cIdx = 0;
63+ auto oriNC = oriN * oriC;
64+ for (uint64_t idx = Simt::GetThreadIdx() + blockID * tNum; idx < shapeSize; idx += bNum * tNum) {
65+ U idxU = U(idx);
66+ hIdx = Simt::UintDiv(idxU, mPNC, sPNC);
67+ U c1Cnt = Simt::UintDiv(idxU, mPNC0, sPNC0);
68+ c1Idx = c1Cnt - Simt::UintDiv(c1Cnt, mC1, sC1) * c1;
69+ U nCnt = Simt::UintDiv(idxU, mC0, sC0);
70+ nIdx = nCnt - Simt::UintDiv(nCnt, mPN, sPN) * padN;
71+ cIdx = idxU - nCnt * c0 + c1Idx * c0;
72+ if (nIdx >= oriN || cIdx >= oriC) {
73+ dst[idx] = T(0);
74+ } else {
75+ dst[idx] = src[hIdx * oriNC + nIdx * oriC + cIdx];
76+ }
77+ }
78+}
79+ 
80+template <typename T>
81+template <typename U>
82+__aicore__ inline void TransWithSIMT<T>::Process()
83+{
84+ __gm__ T* srcAddr = (__gm__ T*)inGM.GetPhyAddr();
85+ __gm__ T* dstAddr = (__gm__ T*)outGM.GetPhyAddr();
86+ 
87+ auto c0 = U(tdPtr->c0);
88+ auto oriN = U(tdPtr->n);
89+ auto oriC = U(tdPtr->c);
90+ auto c1 = U(Ops::Base::CeilDiv(oriC, c0));
91+ auto NI = U(16);
92+ auto padN = U(Ops::Base::CeilAlign(oriN, NI));
93+ uint64_t shapeSize = uint64_t(tdPtr->h) * padN * c1 * c0;
94+ int32_t tNum = int32_t(tdPtr->tNum);
95+ U mPNC = 0;
96+ U sPNC = 0;
97+ U mPNC0 = 0;
98+ U sPNC0 = 0;
99+ U mC1 = 0;
100+ U sC1 = 0;
101+ U mC0 = 0;
102+ U sC0 = 0;
103+ U mPN = 0;
104+ U sPN = 0;
105+ GetUintDivMagicAndShift(mPNC, sPNC, c1 * padN * c0);
106+ GetUintDivMagicAndShift(mPNC0, sPNC0, padN * c0);
107+ GetUintDivMagicAndShift(mC1, sC1, c1);
108+ GetUintDivMagicAndShift(mC0, sC0, c0);
109+ GetUintDivMagicAndShift(mPN, sPN, padN);
110+ Simt::VF_CALL<SIMTTrans<T, U>>(Simt::Dim3(tNum), dstAddr, srcAddr, shapeSize, c1, padN, c0, oriN, oriC, mPNC, sPNC,
111+ mPNC0, sPNC0, mC1, sC1, mC0, sC0, mPN, sPN);
112+}
113+ 
114+} // namespace TRSD
115+ 
116+#endif // TRANS_DATA_WITH_SIMT_IMP_H_
Aconversion/trans_data/op_kernel/trans_data_apt.cpp+20-0
@@ -0,0 +1,20 @@
1+/**
2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file trans_data_apt.cpp
13+ * \brief kernel of op trans_data
14+ */
15+#include "arch35/trans_data.h"
16+ 
17+extern "C" __global__ __aicore__ void trans_data(GM_ADDR src, GM_ADDR dst, GM_ADDR workspace, GM_ADDR tiling)
18+{
19+ trans_data_impl(src, dst, workspace, tiling);
20+}
Dconversion/trans_data/tests/CMakeLists.txt+0-16
@@ -1,16 +0,0 @@
1-# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6-# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9-# See LICENSE in the root of the software repository for the full text of the License.
10-# ----------------------------------------------------------------------------
11-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12-foreach(SUB_DIR ${CURRENT_DIRS})
13- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14- add_subdirectory(${SUB_DIR})
15- endif()
16-endforeach()
Dconversion/trans_data/tests/ut/CMakeLists.txt+0-17
@@ -1,17 +0,0 @@
1-# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6-# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9-# See LICENSE in the root of the software repository for the full text of the License.
10-# ----------------------------------------------------------------------------
11- 
12-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13-foreach(SUB_DIR ${CURRENT_DIRS})
14- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15- add_subdirectory(${SUB_DIR})
16- endif()
17-endforeach()
Dconversion/trans_data/tests/ut/graph_plugin/CMakeLists.txt+0-10
@@ -1,10 +0,0 @@
1-# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6-# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9-# See LICENSE in the root of the software repository for the full text of the License.
10-# ----------------------------------------------------------------------------
Dconversion/trans_data/tests/ut/graph_plugin/fusion_pass/CMakeLists.txt+0-10
@@ -1,10 +0,0 @@
1-# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6-# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9-# See LICENSE in the root of the software repository for the full text of the License.
10-# ----------------------------------------------------------------------------
Rconversion/trans_data/tests/ut/op_host/op_api/test_aclnn_trans_convolution_weight_l2.cpp→conversion/trans_data/tests/ut/op_api/test_aclnn_trans_convolution_weight_l2.cpp+21-2
@@ -11,13 +11,14 @@
11#include <vector>11#include <vector>
12#include "gtest/gtest.h"12#include "gtest/gtest.h"
13 13 
14-#include "aclnn_trans_convolution_weight.h"14+#include "conversion/trans_data/op_api/aclnn_trans_convolution_weight.h"
15 15 
16#include "op_api_ut_common/op_api_ut.h"16#include "op_api_ut_common/op_api_ut.h"
17#include "op_api_ut_common/scalar_desc.h"17#include "op_api_ut_common/scalar_desc.h"
18#include "op_api_ut_common/tensor_desc.h"18#include "op_api_ut_common/tensor_desc.h"
19+#include "opdev/platform.h"
19 20 
20- 21+using namespace op;
21using namespace std;22using namespace std;
22 23 
23class l2_trans_convolution_weight_test : public testing::Test {24class l2_trans_convolution_weight_test : public testing::Test {
@@ -27,6 +28,7 @@ class l2_trans_convolution_weight_test : public testing::Test {
27};28};
28 29 
29TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_normal_input_FP16) {30TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_normal_input_FP16) {
31+ SetPlatformSocVersion(SocVersion::ASCEND310P);
30 // 使用**Desc描述host api输入输出32 // 使用**Desc描述host api输入输出
31 auto x1_desc = TensorDesc({16, 4, 5, 5}, ACL_FLOAT16, ACL_FORMAT_NCHW);33 auto x1_desc = TensorDesc({16, 4, 5, 5}, ACL_FLOAT16, ACL_FORMAT_NCHW);
32 auto y_desc = TensorDesc(x1_desc);34 auto y_desc = TensorDesc(x1_desc);
@@ -39,6 +41,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_normal_input_FP16) {
39}41}
40 42 
41TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_normal_input_FP32) {43TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_normal_input_FP32) {
44+ SetPlatformSocVersion(SocVersion::ASCEND310P);
42 // 使用**Desc描述host api输入输出45 // 使用**Desc描述host api输入输出
43 auto x1_desc = TensorDesc({16, 4, 5, 5}, ACL_FLOAT, ACL_FORMAT_NCHW);46 auto x1_desc = TensorDesc({16, 4, 5, 5}, ACL_FLOAT, ACL_FORMAT_NCHW);
44 auto y_desc = TensorDesc({16, 4, 5, 5}, ACL_FLOAT16, ACL_FORMAT_NCHW);47 auto y_desc = TensorDesc({16, 4, 5, 5}, ACL_FLOAT16, ACL_FORMAT_NCHW);
@@ -51,6 +54,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_normal_input_FP32) {
51}54}
52 55 
53TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_wrong_format) {56TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_wrong_format) {
57+ SetPlatformSocVersion(SocVersion::ASCEND310P);
54 // 使用**Desc描述host api输入输出58 // 使用**Desc描述host api输入输出
55 auto x1_desc = TensorDesc({16, 4, 5, 5}, ACL_FLOAT, ACL_FORMAT_NHWC);59 auto x1_desc = TensorDesc({16, 4, 5, 5}, ACL_FLOAT, ACL_FORMAT_NHWC);
56 auto y_desc = TensorDesc(x1_desc);60 auto y_desc = TensorDesc(x1_desc);
@@ -63,6 +67,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_wrong_format) {
63}67}
64 68 
65TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_wrong_dtype_input_int8) {69TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_wrong_dtype_input_int8) {
70+ SetPlatformSocVersion(SocVersion::ASCEND310P);
66 // 使用**Desc描述host api输入输出71 // 使用**Desc描述host api输入输出
67 auto x1_desc = TensorDesc({16, 4, 5, 5}, ACL_INT8, ACL_FORMAT_NCHW);72 auto x1_desc = TensorDesc({16, 4, 5, 5}, ACL_INT8, ACL_FORMAT_NCHW);
68 auto y_desc = TensorDesc(x1_desc);73 auto y_desc = TensorDesc(x1_desc);
@@ -75,6 +80,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_wrong_dtype_input_int8
75}80}
76 81 
77TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_wrong_format_out) {82TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_wrong_format_out) {
83+ SetPlatformSocVersion(SocVersion::ASCEND310P);
78 // 使用**Desc描述host api输入输出84 // 使用**Desc描述host api输入输出
79 auto x1_desc = TensorDesc({16, 4, 5, 5}, ACL_FLOAT, ACL_FORMAT_NCHW);85 auto x1_desc = TensorDesc({16, 4, 5, 5}, ACL_FLOAT, ACL_FORMAT_NCHW);
80 auto y_desc = TensorDesc({16, 4, 5, 5}, ACL_FLOAT, ACL_FORMAT_NCHW);86 auto y_desc = TensorDesc({16, 4, 5, 5}, ACL_FLOAT, ACL_FORMAT_NCHW);
@@ -87,6 +93,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_wrong_format_out) {
87}93}
88 94 
89TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_wrong_dtype_input_int8_out) {95TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_wrong_dtype_input_int8_out) {
96+ SetPlatformSocVersion(SocVersion::ASCEND310P);
90 // 使用**Desc描述host api输入输出97 // 使用**Desc描述host api输入输出
91 auto x1_desc = TensorDesc({16, 4, 5, 5}, ACL_INT8, ACL_FORMAT_NCHW);98 auto x1_desc = TensorDesc({16, 4, 5, 5}, ACL_INT8, ACL_FORMAT_NCHW);
92 auto y_desc = TensorDesc({16, 4, 5, 5}, ACL_INT8, ACL_FORMAT_NCHW);99 auto y_desc = TensorDesc({16, 4, 5, 5}, ACL_INT8, ACL_FORMAT_NCHW);
@@ -99,6 +106,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P3_test_wrong_dtype_input_int8
99}106}
100 107 
101TEST_F(l2_trans_convolution_weight_test, ascend310P3_dim_larger_than_4) {108TEST_F(l2_trans_convolution_weight_test, ascend310P3_dim_larger_than_4) {
109+ SetPlatformSocVersion(SocVersion::ASCEND310P);
102 // 使用**Desc描述host api输入输出110 // 使用**Desc描述host api输入输出
103 auto x1_desc = TensorDesc({16, 16, 32, 2, 2}, ACL_FLOAT16, ACL_FORMAT_NCHW);111 auto x1_desc = TensorDesc({16, 16, 32, 2, 2}, ACL_FLOAT16, ACL_FORMAT_NCHW);
104 auto y_desc = TensorDesc(x1_desc);112 auto y_desc = TensorDesc(x1_desc);
@@ -112,6 +120,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P3_dim_larger_than_4) {
112 120 
113 121 
114TEST_F(l2_trans_convolution_weight_test, ascend310P3_group_0) {122TEST_F(l2_trans_convolution_weight_test, ascend310P3_group_0) {
123+ SetPlatformSocVersion(SocVersion::ASCEND310P);
115 // 使用**Desc描述host api输入输出124 // 使用**Desc描述host api输入输出
116 auto x1_desc = TensorDesc({16, 16, 32, 2}, ACL_FLOAT16, ACL_FORMAT_NCHW);125 auto x1_desc = TensorDesc({16, 16, 32, 2}, ACL_FLOAT16, ACL_FORMAT_NCHW);
117 auto y_desc = TensorDesc(x1_desc);126 auto y_desc = TensorDesc(x1_desc);
@@ -124,6 +133,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P3_group_0) {
124}133}
125 134 
126TEST_F(l2_trans_convolution_weight_test, ascend310P3_transpose_true) {135TEST_F(l2_trans_convolution_weight_test, ascend310P3_transpose_true) {
136+ SetPlatformSocVersion(SocVersion::ASCEND310P);
127 // 使用**Desc描述host api输入输出137 // 使用**Desc描述host api输入输出
128 auto x1_desc = TensorDesc({16, 16, 32, 2}, ACL_FLOAT16, ACL_FORMAT_NCHW);138 auto x1_desc = TensorDesc({16, 16, 32, 2}, ACL_FLOAT16, ACL_FORMAT_NCHW);
129 auto y_desc = TensorDesc(x1_desc);139 auto y_desc = TensorDesc(x1_desc);
@@ -136,6 +146,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P3_transpose_true) {
136}146}
137 147 
138TEST_F(l2_trans_convolution_weight_test, ascend310P3_in_nullptr) {148TEST_F(l2_trans_convolution_weight_test, ascend310P3_in_nullptr) {
149+ SetPlatformSocVersion(SocVersion::ASCEND310P);
139 // 使用**Desc描述host api输入输出150 // 使用**Desc描述host api输入输出
140 auto x1_desc = TensorDesc({16, 16, 32, 2}, ACL_FLOAT16, ACL_FORMAT_NCHW);151 auto x1_desc = TensorDesc({16, 16, 32, 2}, ACL_FLOAT16, ACL_FORMAT_NCHW);
141 auto y_desc = TensorDesc(x1_desc);152 auto y_desc = TensorDesc(x1_desc);
@@ -148,6 +159,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P3_in_nullptr) {
148}159}
149 160 
150TEST_F(l2_trans_convolution_weight_test, ascend310P3_out_nullptr) {161TEST_F(l2_trans_convolution_weight_test, ascend310P3_out_nullptr) {
162+ SetPlatformSocVersion(SocVersion::ASCEND310P);
151 // 使用**Desc描述host api输入输出163 // 使用**Desc描述host api输入输出
152 auto x1_desc = TensorDesc({16, 16, 32, 2}, ACL_FLOAT16, ACL_FORMAT_NCHW);164 auto x1_desc = TensorDesc({16, 16, 32, 2}, ACL_FLOAT16, ACL_FORMAT_NCHW);
153 auto y_desc = TensorDesc(x1_desc);165 auto y_desc = TensorDesc(x1_desc);
@@ -160,6 +172,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P3_out_nullptr) {
160}172}
161 173 
162TEST_F(l2_trans_convolution_weight_test, ascend310P_normal_shape) {174TEST_F(l2_trans_convolution_weight_test, ascend310P_normal_shape) {
175+ SetPlatformSocVersion(SocVersion::ASCEND310P);
163 aclIntArray* tensorShape = nullptr;176 aclIntArray* tensorShape = nullptr;
164 vector<int64_t> tensorShapeVec = {2, 2, 32, 16};177 vector<int64_t> tensorShapeVec = {2, 2, 32, 16};
165 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());178 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());
@@ -169,12 +182,14 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P_normal_shape) {
169}182}
170 183 
171TEST_F(l2_trans_convolution_weight_test, ascend310P_shape_nullptr) {184TEST_F(l2_trans_convolution_weight_test, ascend310P_shape_nullptr) {
185+ SetPlatformSocVersion(SocVersion::ASCEND310P);
172 uint64_t weightSize = 0;186 uint64_t weightSize = 0;
173 aclnnStatus aclRet = aclnnCalculateConvolutionWeightSize(nullptr, false, 1, ACL_FLOAT16, &weightSize);187 aclnnStatus aclRet = aclnnCalculateConvolutionWeightSize(nullptr, false, 1, ACL_FLOAT16, &weightSize);
174 EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_NULLPTR);188 EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_NULLPTR);
175}189}
176 190 
177TEST_F(l2_trans_convolution_weight_test, ascend310P_size_nullptr) {191TEST_F(l2_trans_convolution_weight_test, ascend310P_size_nullptr) {
192+ SetPlatformSocVersion(SocVersion::ASCEND310P);
178 aclIntArray* tensorShape = nullptr;193 aclIntArray* tensorShape = nullptr;
179 vector<int64_t> tensorShapeVec = {2, 2, 32, 16};194 vector<int64_t> tensorShapeVec = {2, 2, 32, 16};
180 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());195 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());
@@ -183,6 +198,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P_size_nullptr) {
183}198}
184 199 
185TEST_F(l2_trans_convolution_weight_test, ascend310P_datatype_fp32) {200TEST_F(l2_trans_convolution_weight_test, ascend310P_datatype_fp32) {
201+ SetPlatformSocVersion(SocVersion::ASCEND310P);
186 aclIntArray* tensorShape = nullptr;202 aclIntArray* tensorShape = nullptr;
187 vector<int64_t> tensorShapeVec = {2, 2, 32, 16};203 vector<int64_t> tensorShapeVec = {2, 2, 32, 16};
188 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());204 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());
@@ -192,6 +208,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P_datatype_fp32) {
192}208}
193 209 
194TEST_F(l2_trans_convolution_weight_test, ascend310P_group_0) {210TEST_F(l2_trans_convolution_weight_test, ascend310P_group_0) {
211+ SetPlatformSocVersion(SocVersion::ASCEND310P);
195 aclIntArray* tensorShape = nullptr;212 aclIntArray* tensorShape = nullptr;
196 vector<int64_t> tensorShapeVec = {2, 2, 32, 16};213 vector<int64_t> tensorShapeVec = {2, 2, 32, 16};
197 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());214 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());
@@ -201,6 +218,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P_group_0) {
201}218}
202 219 
203TEST_F(l2_trans_convolution_weight_test, ascend310P_wrong_dim) {220TEST_F(l2_trans_convolution_weight_test, ascend310P_wrong_dim) {
221+ SetPlatformSocVersion(SocVersion::ASCEND310P);
204 aclIntArray* tensorShape = nullptr;222 aclIntArray* tensorShape = nullptr;
205 vector<int64_t> tensorShapeVec = {2, 32, 16};223 vector<int64_t> tensorShapeVec = {2, 32, 16};
206 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());224 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());
@@ -210,6 +228,7 @@ TEST_F(l2_trans_convolution_weight_test, ascend310P_wrong_dim) {
210}228}
211 229 
212TEST_F(l2_trans_convolution_weight_test, ascend310P_transpose_true) {230TEST_F(l2_trans_convolution_weight_test, ascend310P_transpose_true) {
231+ SetPlatformSocVersion(SocVersion::ASCEND310P);
213 aclIntArray* tensorShape = nullptr;232 aclIntArray* tensorShape = nullptr;
214 vector<int64_t> tensorShapeVec = {2, 2, 32, 16};233 vector<int64_t> tensorShapeVec = {2, 2, 32, 16};
215 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());234 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());
Rconversion/trans_data/tests/ut/op_host/op_api/test_aclnn_trans_matmul_weight_l2.cpp→conversion/trans_data/tests/ut/op_api/test_aclnn_trans_matmul_weight_l2.cpp+17-1
@@ -11,12 +11,14 @@
11#include <vector>11#include <vector>
12#include "gtest/gtest.h"12#include "gtest/gtest.h"
13 13 
14-#include "aclnn_trans_matmul_weight.h"14+#include "conversion/trans_data/op_api/aclnn_trans_matmul_weight.h"
15 15 
16#include "op_api_ut_common/op_api_ut.h"16#include "op_api_ut_common/op_api_ut.h"
17#include "op_api_ut_common/scalar_desc.h"17#include "op_api_ut_common/scalar_desc.h"
18#include "op_api_ut_common/tensor_desc.h"18#include "op_api_ut_common/tensor_desc.h"
19+#include "opdev/platform.h"
19 20 
21+using namespace op;
20using namespace std;22using namespace std;
21 23 
22class l2_trans_matmul_weight_test : public testing::Test {24class l2_trans_matmul_weight_test : public testing::Test {
@@ -33,6 +35,7 @@ protected:
33 35 
34TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_normal_dim2_input)36TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_normal_dim2_input)
35{37{
38+ SetPlatformSocVersion(SocVersion::ASCEND910B);
36 // 使用**Desc描述host api输入输出39 // 使用**Desc描述host api输入输出
37 auto x1_desc = TensorDesc({16, 32}, ACL_FLOAT16, ACL_FORMAT_ND);40 auto x1_desc = TensorDesc({16, 32}, ACL_FLOAT16, ACL_FORMAT_ND);
38 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());41 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());
@@ -45,6 +48,7 @@ TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_normal_dim2_input)
45 48 
46TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_normal_dim3_input)49TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_normal_dim3_input)
47{50{
51+ SetPlatformSocVersion(SocVersion::ASCEND910B);
48 // 使用**Desc描述host api输入输出52 // 使用**Desc描述host api输入输出
49 auto x1_desc = TensorDesc({16, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND);53 auto x1_desc = TensorDesc({16, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND);
50 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());54 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());
@@ -57,6 +61,7 @@ TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_normal_dim3_input)
57 61 
58TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_empty)62TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_empty)
59{63{
64+ SetPlatformSocVersion(SocVersion::ASCEND910B);
60 // 使用**Desc描述host api输入输出65 // 使用**Desc描述host api输入输出
61 auto x1_desc = TensorDesc({16, 0}, ACL_FLOAT16, ACL_FORMAT_ND);66 auto x1_desc = TensorDesc({16, 0}, ACL_FLOAT16, ACL_FORMAT_ND);
62 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());67 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());
@@ -69,6 +74,7 @@ TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_empty)
69 74 
70TEST_F(l2_trans_matmul_weight_test, ascend910B2_dim_larger_than_3)75TEST_F(l2_trans_matmul_weight_test, ascend910B2_dim_larger_than_3)
71{76{
77+ SetPlatformSocVersion(SocVersion::ASCEND910B);
72 // 使用**Desc描述host api输入输出78 // 使用**Desc描述host api输入输出
73 auto x1_desc = TensorDesc({16, 16, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND);79 auto x1_desc = TensorDesc({16, 16, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND);
74 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());80 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());
@@ -81,6 +87,7 @@ TEST_F(l2_trans_matmul_weight_test, ascend910B2_dim_larger_than_3)
81 87 
82TEST_F(l2_trans_matmul_weight_test, ascend910B2_invalid_format)88TEST_F(l2_trans_matmul_weight_test, ascend910B2_invalid_format)
83{89{
90+ SetPlatformSocVersion(SocVersion::ASCEND910B);
84 // 使用**Desc描述host api输入输出91 // 使用**Desc描述host api输入输出
85 auto x1_desc = TensorDesc({16, 32}, ACL_INT8, ACL_FORMAT_FRACTAL_NZ);92 auto x1_desc = TensorDesc({16, 32}, ACL_INT8, ACL_FORMAT_FRACTAL_NZ);
86 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());93 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());
@@ -93,6 +100,7 @@ TEST_F(l2_trans_matmul_weight_test, ascend910B2_invalid_format)
93 100 
94TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_nullptr)101TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_nullptr)
95{102{
103+ SetPlatformSocVersion(SocVersion::ASCEND910B);
96 uint64_t weightSize = 0;104 uint64_t weightSize = 0;
97 aclnnStatus aclRet = aclnnCalculateMatmulWeightSize(nullptr, &weightSize);105 aclnnStatus aclRet = aclnnCalculateMatmulWeightSize(nullptr, &weightSize);
98 EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_NULLPTR);106 EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_NULLPTR);
@@ -100,6 +108,7 @@ TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_nullptr)
100 108 
101TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_nullptr_2)109TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_nullptr_2)
102{110{
111+ SetPlatformSocVersion(SocVersion::ASCEND910B);
103 aclIntArray* tensorShape = nullptr;112 aclIntArray* tensorShape = nullptr;
104 vector<int64_t> tensorShapeVec = {32, 16};113 vector<int64_t> tensorShapeVec = {32, 16};
105 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());114 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());
@@ -109,6 +118,7 @@ TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_nullptr_2)
109 118 
110TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_invalid)119TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_invalid)
111{120{
121+ SetPlatformSocVersion(SocVersion::ASCEND910B);
112 aclIntArray* tensorShape = nullptr;122 aclIntArray* tensorShape = nullptr;
113 vector<int64_t> tensorShapeVec = {16, 16, 32, 16};123 vector<int64_t> tensorShapeVec = {16, 16, 32, 16};
114 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());124 tensorShape = aclCreateIntArray(tensorShapeVec.data(), tensorShapeVec.size());
@@ -119,6 +129,7 @@ TEST_F(l2_trans_matmul_weight_test, ascend910B2_test_invalid)
119 129 
120TEST_F(l2_trans_matmul_weight_test, ascend950_dim_larger_than_2)130TEST_F(l2_trans_matmul_weight_test, ascend950_dim_larger_than_2)
121{131{
132+ SetPlatformSocVersion(SocVersion::ASCEND950);
122 // 使用**Desc描述host api输入输出133 // 使用**Desc描述host api输入输出
123 auto x1_desc = TensorDesc({16, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND);134 auto x1_desc = TensorDesc({16, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND);
124 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());135 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());
@@ -131,6 +142,7 @@ TEST_F(l2_trans_matmul_weight_test, ascend950_dim_larger_than_2)
131 142 
132TEST_F(l2_trans_matmul_weight_test, ascend310P_test_normal_input_int8)143TEST_F(l2_trans_matmul_weight_test, ascend310P_test_normal_input_int8)
133{144{
145+ SetPlatformSocVersion(SocVersion::ASCEND310P);
134 // 使用**Desc描述host api输入输出146 // 使用**Desc描述host api输入输出
135 auto x1_desc = TensorDesc({16, 32}, ACL_INT8, ACL_FORMAT_ND);147 auto x1_desc = TensorDesc({16, 32}, ACL_INT8, ACL_FORMAT_ND);
136 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());148 auto ut = OP_API_UT(aclnnTransMatmulWeight, INPUT(x1_desc), OUTPUT());
@@ -143,6 +155,7 @@ TEST_F(l2_trans_matmul_weight_test, ascend310P_test_normal_input_int8)
143 155 
144TEST_F(l2_trans_matmul_weight_test, ascend310P_test_nullptr)156TEST_F(l2_trans_matmul_weight_test, ascend310P_test_nullptr)
145{157{
158+ SetPlatformSocVersion(SocVersion::ASCEND310P);
146 uint64_t weightSize = 0;159 uint64_t weightSize = 0;
147 aclDataType dataType = aclDataType::ACL_INT8;160 aclDataType dataType = aclDataType::ACL_INT8;
148 aclnnStatus aclRet = aclnnCalculateMatmulWeightSizeV2(nullptr, dataType, &weightSize);161 aclnnStatus aclRet = aclnnCalculateMatmulWeightSizeV2(nullptr, dataType, &weightSize);
@@ -151,6 +164,7 @@ TEST_F(l2_trans_matmul_weight_test, ascend310P_test_nullptr)
151 164 
152TEST_F(l2_trans_matmul_weight_test, ascend310P_test_nullptr_2)165TEST_F(l2_trans_matmul_weight_test, ascend310P_test_nullptr_2)
153{166{
167+ SetPlatformSocVersion(SocVersion::ASCEND310P);
154 aclIntArray* tensorShape = nullptr;168 aclIntArray* tensorShape = nullptr;
155 vector<int64_t> tensorShapeVec = {32, 16};169 vector<int64_t> tensorShapeVec = {32, 16};
156 aclDataType dataType = aclDataType::ACL_INT8;170 aclDataType dataType = aclDataType::ACL_INT8;
@@ -161,6 +175,7 @@ TEST_F(l2_trans_matmul_weight_test, ascend310P_test_nullptr_2)
161 175 
162TEST_F(l2_trans_matmul_weight_test, ascend310P_test_invalid_shape)176TEST_F(l2_trans_matmul_weight_test, ascend310P_test_invalid_shape)
163{177{
178+ SetPlatformSocVersion(SocVersion::ASCEND310P);
164 aclIntArray* tensorShape = nullptr;179 aclIntArray* tensorShape = nullptr;
165 uint64_t weightSize = 0;180 uint64_t weightSize = 0;
166 vector<int64_t> tensorShapeVec = {0, 16};181 vector<int64_t> tensorShapeVec = {0, 16};
@@ -172,6 +187,7 @@ TEST_F(l2_trans_matmul_weight_test, ascend310P_test_invalid_shape)
172 187 
173TEST_F(l2_trans_matmul_weight_test, ascend310P_test_invalid_shape_dim)188TEST_F(l2_trans_matmul_weight_test, ascend310P_test_invalid_shape_dim)
174{189{
190+ SetPlatformSocVersion(SocVersion::ASCEND310P);
175 aclIntArray* tensorShape = nullptr;191 aclIntArray* tensorShape = nullptr;
176 uint64_t weightSize = 0;192 uint64_t weightSize = 0;
177 vector<int64_t> tensorShapeVec = {2, 2, 2, 2, 2, 32, 16};193 vector<int64_t> tensorShapeVec = {2, 2, 2, 2, 2, 32, 16};
Dconversion/trans_data/tests/ut/op_host/CMakeLists.txt+0-22
@@ -1,22 +0,0 @@
1-# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6-# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9-# See LICENSE in the root of the software repository for the full text of the License.
10-# ----------------------------------------------------------------------------
11- 
12-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13-foreach(SUB_DIR ${CURRENT_DIRS})
14- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15- add_subdirectory(${SUB_DIR})
16- endif()
17-endforeach()
18- 
19-if(UT_TEST_ALL OR OP_HOST_UT)
20- # add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
21- # add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
22-endif()
Dconversion/trans_data/tests/ut/op_host/op_api/CMakeLists.txt+0-14
@@ -1,14 +0,0 @@
1-# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6-# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9-# See LICENSE in the root of the software repository for the full text of the License.
10-# ----------------------------------------------------------------------------
11- 
12-if(UT_TEST_ALL OR OP_API_UT)
13- # add_modules_ut_sources(UT_NAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14-endif()
Dconversion/trans_data/tests/ut/op_host/op_api/test_aclnn_calculate_convolution_weight_size_l2.cpp+0-10
@@ -1,10 +0,0 @@
1-/**
2- * This program is free software, you can redistribute it and/or modify it.
3- * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4- * This file is a part of the CANN Open Software.
5- * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6- * Please refer to the License for details. You may not use this file except in compliance with the License.
7- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8- * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
9- * the software repository for the full text of the License.
10- */
Dconversion/trans_data/tests/ut/op_host/op_api/test_aclnn_calculate_matmul_weight_size_l2.cpp+0-10
@@ -1,10 +0,0 @@
1-/**
2- * This program is free software, you can redistribute it and/or modify it.
3- * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4- * This file is a part of the CANN Open Software.
5- * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6- * Please refer to the License for details. You may not use this file except in compliance with the License.
7- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8- * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
9- * the software repository for the full text of the License.
10- */
Dconversion/trans_data/tests/ut/op_host/op_api/test_aclnn_calculate_matmul_weight_size_v2_l2.cpp+0-10
@@ -1,10 +0,0 @@
1-/**
2- * This program is free software, you can redistribute it and/or modify it.
3- * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4- * This file is a part of the CANN Open Software.
5- * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6- * Please refer to the License for details. You may not use this file except in compliance with the License.
7- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8- * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
9- * the software repository for the full text of the License.
10- */
Dconversion/trans_data/tests/ut/op_kernel/CMakeLists.txt+0-17
@@ -1,17 +0,0 @@
1-# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6-# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9-# See LICENSE in the root of the software repository for the full text of the License.
10-# ----------------------------------------------------------------------------
11- 
12-file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
13-foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
14- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15- add_subdirectory(${SUB_DIR})
16- endif()
17-endforeach()
Mexperimental/CMakeLists.txt+4-4
@@ -1,10 +1,10 @@
1# ---------------------------------------------------------------------------------------------------------1# ---------------------------------------------------------------------------------------------------------
2# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 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").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 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. 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 
@@ -21,7 +21,7 @@ set(COMMON_L0_SRCS
21 ${CMAKE_SOURCE_DIR}/conversion/strided_slice/op_api/strided_slice.cpp21 ${CMAKE_SOURCE_DIR}/conversion/strided_slice/op_api/strided_slice.cpp
22 ${CMAKE_SOURCE_DIR}/conversion/broadcast_to/op_api/broadcast_to.cpp22 ${CMAKE_SOURCE_DIR}/conversion/broadcast_to/op_api/broadcast_to.cpp
23 ${CMAKE_SOURCE_DIR}/conversion/tensor_move/op_api/tensor_move.cpp23 ${CMAKE_SOURCE_DIR}/conversion/tensor_move/op_api/tensor_move.cpp
24- ${CMAKE_SOURCE_DIR}/conversion/trans_data/op_host/op_api/transdata.cpp24+ ${CMAKE_SOURCE_DIR}/conversion/trans_data/op_api/transdata.cpp
25 ${CMAKE_SOURCE_DIR}/conversion/transpose/op_api/transpose.cpp25 ${CMAKE_SOURCE_DIR}/conversion/transpose/op_api/transpose.cpp
26 ${CMAKE_SOURCE_DIR}/conversion/pad/op_api/pad.cpp26 ${CMAKE_SOURCE_DIR}/conversion/pad/op_api/pad.cpp
27 ${CMAKE_SOURCE_DIR}/conversion/fill/op_api/fill.cpp27 ${CMAKE_SOURCE_DIR}/conversion/fill/op_api/fill.cpp
Mscripts/ci/ascend950/ops_math_operator_list.yaml+1-0
@@ -163,6 +163,7 @@ operator_group_5:
163 - strided_slice_assign_v2163 - strided_slice_assign_v2
164 - strided_slice_grad164 - strided_slice_grad
165 - tensor_move165 - tensor_move
166+ - trans_data
166 - transpose167 - transpose
167 - tril168 - tril
168 - triu169 - triu
Mscripts/kernel/binary_config/all_ops_impl_mode.ini+2-1
@@ -78,4 +78,5 @@ BatchMatMul=high_performance,high_precision,enable_float_32_execution,enable_hi_
78BatchMatMulV2=high_performance,high_precision,enable_float_32_execution,enable_hi_float_32_execution78BatchMatMulV2=high_performance,high_precision,enable_float_32_execution,enable_hi_float_32_execution
79FullyConnection=high_performance,high_precision,enable_float_32_execution,enable_hi_float_32_execution79FullyConnection=high_performance,high_precision,enable_float_32_execution,enable_hi_float_32_execution
80GEMM=high_performance,high_precision,enable_float_32_execution,enable_hi_float_32_execution80GEMM=high_performance,high_precision,enable_float_32_execution,enable_hi_float_32_execution
81-Lerp=high_performance,high_precision81+Lerp=high_performance,high_precision
82+TransData=optional
Mscripts/kernel/binary_config/ascendc_config.json+1-1
@@ -440,7 +440,7 @@
440 {"name":"BroadcastTo", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode":"", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},440 {"name":"BroadcastTo", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode":"", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
441 {"name": "DynamicPartition", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},441 {"name": "DynamicPartition", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
442 {"name":"Tile", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode":"", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},442 {"name":"Tile", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode":"", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
443- {"name":"TransData", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode":""},443+ {"name":"TransData", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
444 {"name":"GeluQuant", "compute_units": ["ascend310p", "ascend910b", "ascend910", "ascend950"], "auto_sync" : true},444 {"name":"GeluQuant", "compute_units": ["ascend310p", "ascend910b", "ascend910", "ascend950"], "auto_sync" : true},
445 {"name":"DequantSwigluQuant", "compute_units": ["ascend910b","ascend910_93", "ascend950"], "auto_sync" : true},445 {"name":"DequantSwigluQuant", "compute_units": ["ascend910b","ascend910_93", "ascend950"], "auto_sync" : true},
446 {"name":"QuantBatchMatmulV3", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false},446 {"name":"QuantBatchMatmulV3", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false},
Mscripts/util/parse_ini_to_json.py+3-2
@@ -42,8 +42,9 @@ FORMAT_LIST = ["NCHW", "NHWC", "ND", "NC1HWC0", "FRACTAL_Z", "NC1C0HWPAD", "NHWC
42 "FRACTAL_Z_C04", "CHWN", "FRACTAL_DECONV_SP_STRIDE8_TRANS", "HWCN", "NC1KHKWHWC0", "BN_WEIGHT",42 "FRACTAL_Z_C04", "CHWN", "FRACTAL_DECONV_SP_STRIDE8_TRANS", "HWCN", "NC1KHKWHWC0", "BN_WEIGHT",
43 "FILTER_HWCK", "HASHTABLE_LOOKUP_LOOKUPS", "HASHTABLE_LOOKUP_KEYS", "HASHTABLE_LOOKUP_VALUE",43 "FILTER_HWCK", "HASHTABLE_LOOKUP_LOOKUPS", "HASHTABLE_LOOKUP_KEYS", "HASHTABLE_LOOKUP_VALUE",
44 "HASHTABLE_LOOKUP_OUTPUT", "HASHTABLE_LOOKUP_HITS", "C1HWNCoC0", "MD", "NDHWC", "FRACTAL_ZZ",44 "HASHTABLE_LOOKUP_OUTPUT", "HASHTABLE_LOOKUP_HITS", "C1HWNCoC0", "MD", "NDHWC", "FRACTAL_ZZ",
45- "FRACTAL_NZ", "NCDHW", "DHWCN", "NDC1HWC0", "FRACTAL_Z_3D", "CN", "NC", "DHWNC",45+ "FRACTAL_NZ", "FRACTAL_NZ_C0_16", "FRACTAL_NZ_C0_32", "NCDHW", "DHWCN", "NDC1HWC0", "FRACTAL_Z_3D",
46- "FRACTAL_Z_3D_TRANSPOSE", "FRACTAL_ZN_LSTM", "FRACTAL_ZN_RNN", "FRACTAL_Z_G", "NULL"]46+ "CN", "NC", "DHWNC", "FRACTAL_Z_3D_TRANSPOSE", "FRACTAL_ZN_LSTM", "FRACTAL_ZN_RNN",
47+ "FRACTAL_Z_G", "NULL"]
47 48 
48 49 
49def parse_ini_files(ini_files):50def parse_ini_files(ini_files):