Star1329
1447
代码
代码
Issues117
Pull Requests256
Actions
项目讨论
Wiki
分析
Star1329
1447
已合并
修改A5算子truncatediv的接口文档，部分代码重构 #3096
yefeicoding创建于 6月2日
修改A5算子truncatediv的接口文档，部分代码重构 #3096
已合并
yefeicoding创建于 6月2日
从yefeicoding/ops-mathrefactor合入到cann/ops-mathmaster
共 17 个文件变更+392-240
Mmath/div/docs/aclnnDivMods&aclnnInplaceDivMods.md+7-7
@@ -96,7 +96,7 @@ aclnnStatus aclnnInplaceDivMods(
96 <td>self（aclTensor*）</td>96 <td>self（aclTensor*）</td>
97 <td>输入</td>97 <td>输入</td>
98 <td>表示被除数，公式中的input。</td>98 <td>表示被除数，公式中的input。</td>
99- <td>数据类型与other的数据类型需满足数据类型推导规则（参见<a href="../../../docs/zh/context/互推导关系.md" target="_blank">互推导关系</a>）。</td>99+ <td>数据类型与other的数据类型需满足数据类型推导规则（参见<a href="../../../docs/zh/context/TensorScalar互推导关系.md" target="_blank">TensorScalar互推导关系</a>）。</td>
100 <td>FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16</td>100 <td>FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16</td>
101 <td>ND</td>101 <td>ND</td>
102 <td>不超过8维</td>102 <td>不超过8维</td>
@@ -106,7 +106,7 @@ aclnnStatus aclnnInplaceDivMods(
106 <td>other（aclScalar*）</td>106 <td>other（aclScalar*）</td>
107 <td>输入</td>107 <td>输入</td>
108 <td>表示除数，公式中的输入other。</td>108 <td>表示除数，公式中的输入other。</td>
109- <td>数据类型与self的数据类型需满足数据类型推导规则（参见<a href="../../../docs/zh/context/互推导关系.md" target="_blank">互推导关系</a>）。</td>109+ <td>数据类型与self的数据类型需满足数据类型推导规则（参见<a href="../../../docs/zh/context/TensorScalar互推导关系.md" target="_blank">TensorScalar互推导关系</a>）。</td>
110 <td>FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16</td>110 <td>FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16</td>
111 <td>-</td>111 <td>-</td>
112 <td>-</td>112 <td>-</td>
@@ -157,9 +157,9 @@ aclnnStatus aclnnInplaceDivMods(
157 - <term>Atlas 训练系列产品</term>：不支持BFLOAT16数据类型。157 - <term>Atlas 训练系列产品</term>：不支持BFLOAT16数据类型。
158 158 
159 - <term>Ascend 950PR/Ascend 950DT</term>：159 - <term>Ascend 950PR/Ascend 950DT</term>：
160- - 数据类型需满足<a href="../../../docs/zh/context/TensorScalar互推导关系.md" target="_blank">TensorScalar互推导关系</a>，推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT。160+ - 数据类型需满足<a href="../../../docs/zh/context/TensorScalar互推导关系.md" target="_blank">TensorScalar互推导关系</a>
161- - out在mode为0时，支持FLOAT、FLOAT16、DOUBLE、BFLOAT16、COMPLEX128、COMPLEX64。161+ - mode为0时，out数据类型支持FLOAT、FLOAT16、DOUBLE、BFLOAT16、COMPLEX128、COMPLEX64，self与other推导之后的数据类型为整数类型或布尔类型时，推导之后的数据类型会转换为FLOAT；
162- - mode为1或2时，支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、BFLOAT16、COMPLEX128、COMPLEX64。162+ - mode为1或2时，out数据类型支持FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、BFLOAT16、COMPLEX128、COMPLEX64。
163 163 
164- **返回值：**164- **返回值：**
165 165 
@@ -280,7 +280,7 @@ aclnnStatus aclnnInplaceDivMods(
280 <td>selfRef（aclTensor*）</td>280 <td>selfRef（aclTensor*）</td>
281 <td>输入/输出</td>281 <td>输入/输出</td>
282 <td>表示被除数和商，公式中的输入input和out。</td>282 <td>表示被除数和商，公式中的输入input和out。</td>
283- <td>数据类型需要是selfRef与other推导之后可转换的数据类型（参见<a href="../../../docs/zh/context/互转换关系.md" target="_blank">互转换关系</a>）。</td>283+ <td>数据类型需要是selfRef与other推导之后可转换的数据类型（参见<a href="../../../docs/zh/context/TensorScalar互推导关系.md" target="_blank">TensorScalar互推导关系</a>）。</td>
284 <td>FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16</td>284 <td>FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16</td>
285 <td>ND</td>285 <td>ND</td>
286 <td>不超过8维</td>286 <td>不超过8维</td>
@@ -290,7 +290,7 @@ aclnnStatus aclnnInplaceDivMods(
290 <td>other（aclScalar*）</td>290 <td>other（aclScalar*）</td>
291 <td>输入</td>291 <td>输入</td>
292 <td>表示除数，公式中的输入other。</td>292 <td>表示除数，公式中的输入other。</td>
293- <td>数据类型与selfRef的数据类型需满足数据类型推导规则（参见<a href="../../../docs/zh/context/互推导关系.md" target="_blank">互推导关系</a>）。</td>293+ <td>数据类型与selfRef的数据类型需满足数据类型推导规则（参见<a href="../../../docs/zh/context/TensorScalar互推导关系.md" target="_blank">TensorScalar互推导关系</a>）。</td>
294 <td>FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16</td>294 <td>FLOAT、FLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL、COMPLEX128、COMPLEX64、BFLOAT16</td>
295 <td>-</td>295 <td>-</td>
296 <td>-</td>296 <td>-</td>
Mmath/div/op_api/aclnn_div.cpp+4-22
@@ -763,24 +763,15 @@ aclnnStatus aclnnDivModGetWorkspaceSize(
763 763 
764 // TruncateDiv 特殊处理：IsRegBase && mode=MODE_TRUNC_DIV && 类型组合在映射表中，不做类型提升764 // TruncateDiv 特殊处理：IsRegBase && mode=MODE_TRUNC_DIV && 类型组合在映射表中，不做类型提升
765 if (IsRegBase(npuArch) && mode == MODE_TRUNC_DIV) {765 if (IsRegBase(npuArch) && mode == MODE_TRUNC_DIV) {
766- OP_LOGI(
767- "aclnnDivMod", "Enter TruncateDiv branch, selfDtype=%s, otherDtype=%s",
768- op::ToString(self->GetDataType()).GetString(), op::ToString(other->GetDataType()).GetString());
769 if (isInTruncDtypeMapping(self->GetDataType(), other->GetDataType())) {766 if (isInTruncDtypeMapping(self->GetDataType(), other->GetDataType())) {
770- OP_LOGI(
771- "aclnnDivMod", "TruncateDiv direct path: no type promotion, selfDtype=%s, otherDtype=%s",
772- op::ToString(self->GetDataType()).GetString(), op::ToString(other->GetDataType()).GetString());
773 divOpOut = l0op::TruncateDiv(selfContiguous, otherContiguous, uniqueExecutor.get());767 divOpOut = l0op::TruncateDiv(selfContiguous, otherContiguous, uniqueExecutor.get());
774 } else {768 } else {
775 op::DataType promoteType;769 op::DataType promoteType;
776 promoteType = InferDivModeDtype(self->GetDataType(), other->GetDataType(), mode);770 promoteType = InferDivModeDtype(self->GetDataType(), other->GetDataType(), mode);
777 bool needToFloat = (promoteType == op::DataType::DT_BOOL);771 bool needToFloat = (promoteType == op::DataType::DT_BOOL);
778 promoteType = needToFloat ? op::DataType::DT_FLOAT : promoteType;772 promoteType = needToFloat ? op::DataType::DT_FLOAT : promoteType;
779- OP_LOGI(773+ auto complexRet = CheckDivModComplexDtype(promoteType, mode);
780- "aclnnDivMod", "TruncateDiv cast path: selfDtype=%s -> %s, otherDtype=%s -> %s, promoteType=%s",774+ CHECK_RET(complexRet == ACLNN_SUCCESS, complexRet);
781- op::ToString(self->GetDataType()).GetString(), op::ToString(promoteType).GetString(),
782- op::ToString(other->GetDataType()).GetString(), op::ToString(promoteType).GetString(),
783- op::ToString(promoteType).GetString());
784 selfCasted = l0op::Cast(selfContiguous, promoteType, uniqueExecutor.get());775 selfCasted = l0op::Cast(selfContiguous, promoteType, uniqueExecutor.get());
785 CHECK_RET(selfCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);776 CHECK_RET(selfCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
786 otherCasted = l0op::Cast(otherContiguous, promoteType, uniqueExecutor.get());777 otherCasted = l0op::Cast(otherContiguous, promoteType, uniqueExecutor.get());
@@ -891,13 +882,7 @@ aclnnStatus aclnnDivModsGetWorkspaceSize(
891 882 
892 // TruncateDiv 特殊处理：IsRegBase && mode=MODE_TRUNC_DIV && 类型组合在映射表中，不做类型提升883 // TruncateDiv 特殊处理：IsRegBase && mode=MODE_TRUNC_DIV && 类型组合在映射表中，不做类型提升
893 if (IsRegBase(npuArch) && mode == MODE_TRUNC_DIV) {884 if (IsRegBase(npuArch) && mode == MODE_TRUNC_DIV) {
894- OP_LOGI(
895- "aclnnDivMods", "Enter TruncateDiv branch, selfDtype=%s, otherDtype=%s",
896- op::ToString(self->GetDataType()).GetString(), op::ToString(other->GetDataType()).GetString());
897 if (isInTruncDtypeMapping(self->GetDataType(), other->GetDataType())) {885 if (isInTruncDtypeMapping(self->GetDataType(), other->GetDataType())) {
898- OP_LOGI(
899- "aclnnDivMods", "TruncateDiv direct path: no type promotion, selfDtype=%s, otherDtype=%s",
900- op::ToString(self->GetDataType()).GetString(), op::ToString(other->GetDataType()).GetString());
901 auto otherConvert = uniqueExecutor.get()->ConvertToTensor(other, other->GetDataType());886 auto otherConvert = uniqueExecutor.get()->ConvertToTensor(other, other->GetDataType());
902 CHECK_RET(otherConvert != nullptr, ACLNN_ERR_INNER_NULLPTR);887 CHECK_RET(otherConvert != nullptr, ACLNN_ERR_INNER_NULLPTR);
903 divOpOut = l0op::TruncateDiv(selfContiguous, otherConvert, uniqueExecutor.get());888 divOpOut = l0op::TruncateDiv(selfContiguous, otherConvert, uniqueExecutor.get());
@@ -906,11 +891,8 @@ aclnnStatus aclnnDivModsGetWorkspaceSize(
906 promoteType = InferDivModeDtype(self->GetDataType(), other->GetDataType(), mode);891 promoteType = InferDivModeDtype(self->GetDataType(), other->GetDataType(), mode);
907 bool needToFloat = (promoteType == op::DataType::DT_BOOL);892 bool needToFloat = (promoteType == op::DataType::DT_BOOL);
908 promoteType = needToFloat ? op::DataType::DT_FLOAT : promoteType;893 promoteType = needToFloat ? op::DataType::DT_FLOAT : promoteType;
909- OP_LOGI(894+ auto complexRet = CheckDivModComplexDtype(promoteType, mode);
910- "aclnnDivMods", "TruncateDiv cast path: selfDtype=%s -> %s, otherDtype=%s -> %s, promoteType=%s",895+ CHECK_RET(complexRet == ACLNN_SUCCESS, complexRet);
911- op::ToString(self->GetDataType()).GetString(), op::ToString(promoteType).GetString(),
912- op::ToString(other->GetDataType()).GetString(), op::ToString(promoteType).GetString(),
913- op::ToString(promoteType).GetString());
914 selfCasted = l0op::Cast(selfContiguous, promoteType, uniqueExecutor.get());896 selfCasted = l0op::Cast(selfContiguous, promoteType, uniqueExecutor.get());
915 CHECK_RET(selfCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);897 CHECK_RET(selfCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
916 auto otherCasted = uniqueExecutor.get()->ConvertToTensor(other, promoteType);898 auto otherCasted = uniqueExecutor.get()->ConvertToTensor(other, promoteType);
Mmath/div/tests/ut/op_api/test_aclnn_divMod.cpp+32-6
@@ -213,7 +213,7 @@ TEST_F(l2_div_mod_test, case_inplace_other_support)
213 EXPECT_EQ(aclRet, ACL_SUCCESS);213 EXPECT_EQ(aclRet, ACL_SUCCESS);
214}214}
215 215 
216-TEST_F(l2_div_mod_test, case_complex_mod_1_invalid)216+TEST_F(l2_div_mod_test, case_complex_mods_1_invalid)
217{217{
218 auto self_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).ValueRange(1, 2);218 auto self_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).ValueRange(1, 2);
219 auto other_tensor_desc = ScalarDesc(2.0f);219 auto other_tensor_desc = ScalarDesc(2.0f);
@@ -223,20 +223,46 @@ TEST_F(l2_div_mod_test, case_complex_mod_1_invalid)
223 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));223 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));
224 uint64_t workspace_size = 0;224 uint64_t workspace_size = 0;
225 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);225 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
226- // EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);226+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
227}227}
228 228 
229-TEST_F(l2_div_mod_test, case_complex_mod_2_invalid)229+TEST_F(l2_div_mod_test, case_complex_mods_2_invalid)
230{230{
231- auto self_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).ValueRange(1, 2);231+ auto self_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX128, ACL_FORMAT_ND).ValueRange(1, 2);
232 auto other_tensor_desc = ScalarDesc(2.0f);232 auto other_tensor_desc = ScalarDesc(2.0f);
233- auto out_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).Precision(0.001, 0.001);233+ auto out_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX128, ACL_FORMAT_ND).Precision(0.001, 0.001);
234 int mode = 2;234 int mode = 2;
235 235 
236 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));236 auto ut = OP_API_UT(aclnnDivMods, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));
237 uint64_t workspace_size = 0;237 uint64_t workspace_size = 0;
238 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);238 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
239- // EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);239+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
240+}
241+ 
242+TEST_F(l2_div_mod_test, case_complex_mod_1_invalid)
243+{
244+ auto self_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).ValueRange(1, 2);
245+ auto other_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).ValueRange(1, 2);
246+ auto out_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX64, ACL_FORMAT_ND).Precision(0.001, 0.001);
247+ int mode = 2;
248+ 
249+ auto ut = OP_API_UT(aclnnDivMod, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));
250+ uint64_t workspace_size = 0;
251+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
252+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
253+}
254+ 
255+TEST_F(l2_div_mod_test, case_complex_mod_2_invalid)
256+{
257+ auto self_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX128, ACL_FORMAT_ND).ValueRange(1, 2);
258+ auto other_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX128, ACL_FORMAT_ND).ValueRange(1, 2);
259+ auto out_tensor_desc = TensorDesc({2, 3}, ACL_COMPLEX128, ACL_FORMAT_ND).Precision(0.001, 0.001);
260+ int mode = 2;
261+ 
262+ auto ut = OP_API_UT(aclnnDivMod, INPUT(self_tensor_desc, other_tensor_desc, mode), OUTPUT(out_tensor_desc));
263+ uint64_t workspace_size = 0;
264+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
265+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
240}266}
241 267 
242// 测试mode异常值拦截268// 测试mode异常值拦截
Mmath/truncate_div/examples/test_geir_truncate_div.cpp+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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 can not use this file except in compliance with the License.5 * Please refer to the License for details. You can not use this file except in compliance with the License.
Mmath/truncate_div/op_api/truncate_div.cpp+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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.
Mmath/truncate_div/op_api/truncate_div.h+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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.
Mmath/truncate_div/op_graph/truncate_div_proto.h+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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.
Mmath/truncate_div/op_host/arch35/truncate_div_tiling_arch35.cpp+175-130
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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.
@@ -77,9 +77,168 @@ ge::graphStatus TruncateDivTiling::GetConstData(uint32_t inputIdx, T& data)
77 return ge::GRAPH_SUCCESS;77 return ge::GRAPH_SUCCESS;
78}78}
79 79 
80+template <typename OpDag>
81+ge::graphStatus TruncateDivTiling::ExecTiling(bool isScalarBranch)
82+{
83+ BroadcastBaseTiling<OpDag> brcTiling(context_);
84+ if (isScalarBranch) {
85+ brcTiling.SetScalar(reciprocal_);
86+ }
87+ 
88+ auto ret = brcTiling.DoTiling();
89+ schMode_ = brcTiling.GetSchMode();
90+ tilingKey_ = GET_TPL_TILING_KEY(schMode_, canUseMul_);
91+ 
92+ return ret;
93+}
94+ 
95+ge::graphStatus TruncateDivTiling::GetScalarReciprocal(ge::DataType x2DType)
96+{
97+ bool success = false;
98+ float scalarValue = 0.0f;
99+ 
100+ switch (x2DType) {
101+ case ge::DT_FLOAT: {
102+ success = (GetConstData<float>(INPUT_IDX_X2, scalarValue) == ge::GRAPH_SUCCESS);
103+ break;
104+ }
105+ case ge::DT_FLOAT16: {
106+ uint16_t tmpValue = 0;
107+ success = (GetConstData<uint16_t>(INPUT_IDX_X2, tmpValue) == ge::GRAPH_SUCCESS);
108+ if (success) {
109+ scalarValue = float(*(reinterpret_cast<const fp16_t*>(&tmpValue)));
110+ }
111+ break;
112+ }
113+ case ge::DT_BF16: {
114+ uint16_t tmpValue = 0;
115+ success = (GetConstData<uint16_t>(INPUT_IDX_X2, tmpValue) == ge::GRAPH_SUCCESS);
116+ if (success) {
117+ scalarValue = float(*(reinterpret_cast<const bfloat16*>(&tmpValue)));
118+ }
119+ break;
120+ }
121+ default: {
122+ OP_LOGE_FOR_INVALID_DTYPE_WITH_REASON(
123+ context_->GetNodeName(), "x2", ToString(x2DType).c_str(),
124+ "The dtype of x2 must be within the range DT_FLOAT, DT_FLOAT16 and DT_BF16");
125+ return ge::GRAPH_FAILED;
126+ }
127+ }
128+ 
129+ if (!success) {
130+ return ge::GRAPH_FAILED;
131+ }
132+ 
133+ reciprocal_ = GetReciprocal(scalarValue);
134+ OP_LOGI(context_->GetNodeName(), "scalar value = %f, reciprocal value = %f", scalarValue, reciprocal_);
135+ return ge::GRAPH_SUCCESS;
136+}
137+ 
138+ge::graphStatus TruncateDivTiling::HandleFloat16WithFloat()
139+{
140+ return canUseMul_ ? ExecTiling<TruncateDivOp::TruncateDivFloatWithCastScalar<half, float, float>::OpDag>(true) :
141+ ExecTiling<TruncateDivOp::TruncateDivFloatWithCast<half, float, float>::OpDag>(false);
142+}
143+ 
144+ge::graphStatus TruncateDivTiling::HandleFloat16OrBf16()
145+{
146+ return canUseMul_ ? ExecTiling<TruncateDivOp::TruncateDivFloat16Scalar<half, float>::OpDag>(true) :
147+ ExecTiling<TruncateDivOp::TruncateDivFloat16<half, float>::OpDag>(false);
148+}
149+ 
150+ge::graphStatus TruncateDivTiling::HandleFloat(ge::DataType x2DType)
151+{
152+ if (x2DType == ge::DT_FLOAT) {
153+ return canUseMul_ ? ExecTiling<TruncateDivOp::TruncateDivFloatScalar<float>::OpDag>(true) :
154+ ExecTiling<TruncateDivOp::TruncateDivFloat<float>::OpDag>(false);
155+ }
156+ if (x2DType == ge::DT_INT32) {
157+ return ExecTiling<TruncateDivOp::TruncateDivFloatToLowBit<float, int32_t, float>::OpDag>(false);
158+ }
159+ if (x2DType == ge::DT_FLOAT16) {
160+ return canUseMul_ ? ExecTiling<TruncateDivOp::TruncateDivFloatScalar<float>::OpDag>(true) :
161+ ExecTiling<TruncateDivOp::TruncateDivFloatToLowBit<float, half, float>::OpDag>(false);
162+ }
163+ return ge::GRAPH_PARAM_INVALID;
164+}
165+ 
166+ge::graphStatus TruncateDivTiling::HandleIntTypes(ge::DataType x1DType, ge::DataType x2DType)
167+{
168+ if (x1DType != x2DType) {
169+ return ge::GRAPH_PARAM_INVALID;
170+ }
171+ 
172+ if (x1DType == ge::DT_INT8) {
173+ return ExecTiling<TruncateDivOp::TruncateDivIntS8<int8_t, half>::OpDag>(false);
174+ }
175+ if (x1DType == ge::DT_UINT8) {
176+ return ExecTiling<TruncateDivOp::TruncateDivIntU8<uint8_t, uint16_t>::OpDag>(false);
177+ }
178+ if (x1DType == ge::DT_INT16) {
179+ return ExecTiling<TruncateDivOp::TruncateDivInt<int16_t>::OpDag>(false);
180+ }
181+ if (x1DType == ge::DT_INT32) {
182+ return ExecTiling<TruncateDivOp::TruncateDivInt<int32_t>::OpDag>(false);
183+ }
184+ 
185+ return ge::GRAPH_PARAM_INVALID;
186+}
187+ 
188+ge::graphStatus TruncateDivTiling::HandleInt64()
189+{
190+ int64_t maxLiveNodeCnt = 0;
191+ int64_t extraBuf = DCACHE_SIZE;
192+ BroadcastBaseTiling<TruncateDivOp::TruncateDivInt64<int64_t>::OpDag> brcTiling(context_);
193+ auto ret = brcTiling.DoTiling(extraBuf, maxLiveNodeCnt);
194+ schMode_ = brcTiling.GetSchMode();
195+ tilingKey_ = GET_TPL_TILING_KEY(schMode_, canUseMul_);
196+ return ret;
197+}
198+ 
199+ge::graphStatus TruncateDivTiling::SelectAndExecTiling(ge::DataType x1DType, ge::DataType x2DType)
200+{
201+ ge::graphStatus ret = ge::GRAPH_PARAM_INVALID;
202+ 
203+ if (x1DType == ge::DT_FLOAT16 && x2DType == ge::DT_FLOAT) {
204+ return HandleFloat16WithFloat();
205+ }
206+ 
207+ if ((x1DType == ge::DT_FLOAT16 || x1DType == ge::DT_BF16) && x2DType == x1DType) {
208+ return HandleFloat16OrBf16();
209+ }
210+ 
211+ if (x1DType == ge::DT_FLOAT) {
212+ ret = HandleFloat(x2DType);
213+ if (ret != ge::GRAPH_PARAM_INVALID) {
214+ return ret;
215+ }
216+ }
217+ 
218+ if (x1DType == ge::DT_INT32 && x2DType == ge::DT_FLOAT) {
219+ return ExecTiling<TruncateDivOp::TruncateDivIntToFloat<int32_t, float, float>::OpDag>(false);
220+ }
221+ 
222+ ret = HandleIntTypes(x1DType, x2DType);
223+ if (ret != ge::GRAPH_PARAM_INVALID) {
224+ return ret;
225+ }
226+ 
227+ if (x1DType == ge::DT_INT64) {
228+ return HandleInt64();
229+ }
230+ 
231+ std::string errorDtype = ToString(x1DType) + ", " + ToString(x2DType);
232+ std::string errorMsg =
233+ std::string("The dtypes of these parameters support only the following combinations: ") +
234+ "((DT_FLOAT16, DT_FLOAT), (DT_FLOAT16, DT_BF16), all DT_FLOAT, (DT_FLOAT, DT_INT32), (DT_FLOAT, DT_FLOAT16), " +
235+ "all DT_INT8, all DT_UINT8, all DT_INT16, all DT_INT32, all DT_INT64 and (DT_INT32, DT_FLOAT))";
236+ OP_LOGE_FOR_INVALID_DTYPES_WITH_REASON(context_->GetNodeName(), "x1, x2", errorDtype.c_str(), errorMsg.c_str());
237+ return ge::GRAPH_FAILED;
238+}
239+ 
80ge::graphStatus TruncateDivTiling::DoOpTiling()240ge::graphStatus TruncateDivTiling::DoOpTiling()
81{241{
82- // 1. 获取输入描述
83 auto x1Desc = context_->GetInputDesc(INPUT_IDX_X1);242 auto x1Desc = context_->GetInputDesc(INPUT_IDX_X1);
84 OP_CHECK_NULL_WITH_CONTEXT(context_, x1Desc);243 OP_CHECK_NULL_WITH_CONTEXT(context_, x1Desc);
85 auto x2Desc = context_->GetInputDesc(INPUT_IDX_X2);244 auto x2Desc = context_->GetInputDesc(INPUT_IDX_X2);
@@ -88,146 +247,33 @@ ge::graphStatus TruncateDivTiling::DoOpTiling()
88 ge::DataType x1DType = x1Desc->GetDataType();247 ge::DataType x1DType = x1Desc->GetDataType();
89 ge::DataType x2DType = x2Desc->GetDataType();248 ge::DataType x2DType = x2Desc->GetDataType();
90 249 
91- // 2. 获取形状
92 auto x1StorageShape = context_->GetInputShape(INPUT_IDX_X1);250 auto x1StorageShape = context_->GetInputShape(INPUT_IDX_X1);
93 OP_CHECK_NULL_WITH_CONTEXT(context_, x1StorageShape);251 OP_CHECK_NULL_WITH_CONTEXT(context_, x1StorageShape);
94 auto x2StorageShape = context_->GetInputShape(INPUT_IDX_X2);252 auto x2StorageShape = context_->GetInputShape(INPUT_IDX_X2);
95 OP_CHECK_NULL_WITH_CONTEXT(context_, x2StorageShape);253 OP_CHECK_NULL_WITH_CONTEXT(context_, x2StorageShape);
96 254 
97 auto x2Shape = x2StorageShape->GetStorageShape();255 auto x2Shape = x2StorageShape->GetStorageShape();
98- 
99 bool isScalar = x2Shape.IsScalar();256 bool isScalar = x2Shape.IsScalar();
100- bool canUseMul = isScalar && (x2DType == ge::DT_FLOAT || x2DType == ge::DT_FLOAT16 || x2DType == ge::DT_BF16);257+ canUseMul_ = isScalar && (x2DType == ge::DT_FLOAT || x2DType == ge::DT_FLOAT16 || x2DType == ge::DT_BF16);
101 258 
102- OP_LOGI(context_->GetNodeName(), "canUseMul %d", canUseMul);259+ OP_LOGI(context_->GetNodeName(), "canUseMul_ %d", canUseMul_);
103- if (canUseMul) {260+ if (canUseMul_) {
104- bool success = false;261+ auto ret = GetScalarReciprocal(x2DType);
105- float scalarValue = 0.0f;262+ if (ret != ge::GRAPH_SUCCESS) {
106- switch (x2DType) {263+ return ret;
107- case ge::DT_FLOAT: {
108- success = (GetConstData<float>(INPUT_IDX_X2, scalarValue) == ge::GRAPH_SUCCESS);
109- if (success) {
110- reciprocal_ = GetReciprocal(scalarValue);
111- }
112- break;
113- }
114- case ge::DT_FLOAT16: {
115- uint16_t tmpValue = 0;
116- success = (GetConstData<uint16_t>(INPUT_IDX_X2, tmpValue) == ge::GRAPH_SUCCESS);
117- if (success) {
118- scalarValue = float(*(reinterpret_cast<const fp16_t*>(&tmpValue)));
119- reciprocal_ = GetReciprocal(scalarValue);
120- }
121- break;
122- }
123- case ge::DT_BF16: {
124- uint16_t tmpValue = 0;
125- success = (GetConstData<uint16_t>(INPUT_IDX_X2, tmpValue) == ge::GRAPH_SUCCESS);
126- if (success) {
127- scalarValue = float(*(reinterpret_cast<const bfloat16*>(&tmpValue)));
128- reciprocal_ = GetReciprocal(scalarValue);
129- }
130- break;
131- }
132- default:
133- OP_LOGE_FOR_INVALID_DTYPE_WITH_REASON(
134- context_->GetNodeName(),
135- "x2", ToString(x2DType).c_str(),
136- "The dtype of x2 must be within the range DT_FLOAT, DT_FLOAT16 and DT_BF16");
137- return ge::GRAPH_FAILED;
138 }264 }
139- if (!success) {
140- return ge::GRAPH_FAILED;
141- }
142- OP_LOGI(context_->GetNodeName(), "scalar value = %f, reciprocal value = %f", scalarValue, reciprocal_);
143 }265 }
144 266 
145- 267+ auto ret = SelectAndExecTiling(x1DType, x2DType);
146- ge::graphStatus ret = ge::GRAPH_SUCCESS;268+ if (ret != ge::GRAPH_SUCCESS) {
147- uint32_t schMode = 0;269+ return ret;
148- int64_t maxLiveNodeCnt = 0;
149- int64_t extraBuf = DCACHE_SIZE;
150- 
151- // 5. 定义模板 lambda：封装重复的 tiling 执行逻辑
152- auto execTiling = [&, this]<typename OpDag>(bool isScalarBranch = false) {
153- BroadcastBaseTiling<OpDag> brcTiling(context_);
154- if (isScalarBranch) {
155- brcTiling.SetScalar(reciprocal_);
156- }
157- 
158- ret = brcTiling.DoTiling();
159- schMode = brcTiling.GetSchMode();
160- 
161- tilingKey_ = GET_TPL_TILING_KEY(schMode, canUseMul);
162- };
163- 
164- // 6. 根据数据类型组合调用 lambda
165- if (x1DType == ge::DT_FLOAT16 && x2DType == ge::DT_FLOAT) {
166- if (canUseMul) {
167- execTiling.template operator()<TruncateDivOp::TruncateDivFloatWithCastScalar<half, float, float>::OpDag>(
168- true);
169- } else {
170- execTiling.template operator()<TruncateDivOp::TruncateDivFloatWithCast<half, float, float>::OpDag>(false);
171- }
172- } else if (x1DType == ge::DT_FLOAT16 || x1DType == ge::DT_BF16) {
173- if (x2DType == x1DType) {
174- if (canUseMul) {
175- execTiling.template operator()<TruncateDivOp::TruncateDivFloat16Scalar<half, float>::OpDag>(true);
176- } else {
177- execTiling.template operator()<TruncateDivOp::TruncateDivFloat16<half, float>::OpDag>(false);
178- }
179- }
180- } else if (x1DType == ge::DT_FLOAT) {
181- if (x2DType == ge::DT_FLOAT) {
182- if (canUseMul) {
183- execTiling.template operator()<TruncateDivOp::TruncateDivFloatScalar<float>::OpDag>(true);
184- } else {
185- execTiling.template operator()<TruncateDivOp::TruncateDivFloat<float>::OpDag>(false);
186- }
187- } else if (x2DType == ge::DT_INT32) {
188- execTiling.template operator()<TruncateDivOp::TruncateDivFloatToLowBit<float, int32_t, float>::OpDag>(
189- false);
190- } else if (x2DType == ge::DT_FLOAT16) {
191- if (canUseMul) {
192- execTiling.template operator()<TruncateDivOp::TruncateDivFloatScalar<float>::OpDag>(true);
193- } else {
194- execTiling.template operator()<TruncateDivOp::TruncateDivFloatToLowBit<float, half, float>::OpDag>(false);
195- }
196- }
197- } else if (x1DType == ge::DT_INT8 || x1DType == ge::DT_UINT8) {
198- // 这些分支不支持 scalar 优化（canUseMul 应该为 false）
199- if (x1DType == ge::DT_INT8) {
200- execTiling.template operator()<TruncateDivOp::TruncateDivIntS8<int8_t, half>::OpDag>(false);
201- } else {
202- execTiling.template operator()<TruncateDivOp::TruncateDivIntU8<uint8_t, uint16_t>::OpDag>(false);
203- }
204- } else if (x1DType == ge::DT_INT16) {
205- execTiling.template operator()<TruncateDivOp::TruncateDivInt<int16_t>::OpDag>(false);
206- } else if (x1DType == ge::DT_INT32 && x2DType == ge::DT_INT32) {
207- execTiling.template operator()<TruncateDivOp::TruncateDivInt<int32_t>::OpDag>(false);
208- } else if (x1DType == ge::DT_INT64) {
209- BroadcastBaseTiling<TruncateDivOp::TruncateDivInt64<int64_t>::OpDag> brcTiling(context_);
210- ret = brcTiling.DoTiling(extraBuf, maxLiveNodeCnt);
211- schMode = brcTiling.GetSchMode();
212- tilingKey_ = GET_TPL_TILING_KEY(schMode, canUseMul);
213- } else if (x1DType == ge::DT_INT32 && x2DType == ge::DT_FLOAT) {
214- execTiling.template operator()<TruncateDivOp::TruncateDivIntToFloat<int32_t, float, float>::OpDag>(false);
215- } else {
216- std::string errorDtype = ToString(x1DType) + ", " + ToString(x2DType);
217- std::string errorMsg = std::string("The dtypes of these parameters support only the following combinations: ") +
218- "((DT_FLOAT16, DT_FLOAT), (DT_FLOAT16, DT_BF16), all DT_FLOAT, (DT_FLOAT, DT_INT32), (DT_FLOAT, DT_FLOAT16), " +
219- "all DT_INT8, all DT_UINT8, all DT_INT16, all DT_INT32, all DT_INT64 and (DT_INT32, DT_FLOAT))";
220- OP_LOGE_FOR_INVALID_DTYPES_WITH_REASON(
221- context_->GetNodeName(), "x1, x2", errorDtype.c_str(), errorMsg.c_str());
222- return ge::GRAPH_FAILED;
223 }270 }
224 271 
225- // 7. 完成后打印结果
226 OP_LOGI(272 OP_LOGI(
227- context_, "TruncateDiv tiling completed: tilingKey_=%lu, schMode=%u, canUseMul=%d, reciprocal %f", tilingKey_, schMode,273+ context_, "TruncateDiv tiling completed: tilingKey_=%lu, schMode_=%u, canUseMul_=%d, reciprocal %f", tilingKey_,
228- static_cast<int>(canUseMul), reciprocal_);274+ schMode_, static_cast<int>(canUseMul_), reciprocal_);
229 275 
230- return ret;276+ return ge::GRAPH_SUCCESS;
231}277}
232 278 
233ge::graphStatus TruncateDivTiling::DoLibApiTiling()279ge::graphStatus TruncateDivTiling::DoLibApiTiling()
@@ -269,9 +315,8 @@ ge::graphStatus TruncateDivTiling::GetPlatformInfo()
269 return ge::GRAPH_SUCCESS;315 return ge::GRAPH_SUCCESS;
270}316}
271 317 
272-ge::graphStatus TilingForTruncateDiv(gert::TilingContext* context)318+static ge::graphStatus TilingForTruncateDiv(gert::TilingContext* context)
273{319{
274- OP_LOGD("TruncateDivTiling", "Enter TilingForTruncateDiv");
275 OP_CHECK_NULL_WITH_CONTEXT(context, context);320 OP_CHECK_NULL_WITH_CONTEXT(context, context);
276 321 
277 auto compileInfo = reinterpret_cast<const BroadcastCompileInfo*>(context->GetCompileInfo());322 auto compileInfo = reinterpret_cast<const BroadcastCompileInfo*>(context->GetCompileInfo());
@@ -281,7 +326,7 @@ ge::graphStatus TilingForTruncateDiv(gert::TilingContext* context)
281 return Ops::Math::OpTiling::TilingRegistry::GetInstance().DoTilingImpl(context);326 return Ops::Math::OpTiling::TilingRegistry::GetInstance().DoTilingImpl(context);
282}327}
283 328 
284-ge::graphStatus TilingPrepareForTruncateDiv([[maybe_unused]] gert::TilingParseContext* context)329+static ge::graphStatus TilingPrepareForTruncateDiv([[maybe_unused]] gert::TilingParseContext* context)
285{330{
286 return ge::GRAPH_SUCCESS;331 return ge::GRAPH_SUCCESS;
287}332}
Mmath/truncate_div/op_host/arch35/truncate_div_tiling_arch35.h+16-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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.
@@ -44,11 +44,26 @@ private:
44 uint64_t tilingKey_ = 0;44 uint64_t tilingKey_ = 0;
45 float reciprocal_ = 0.0f;45 float reciprocal_ = 0.0f;
46 int64_t ubSize_ = 0;46 int64_t ubSize_ = 0;
47+ uint32_t schMode_ = 0;
48+ bool canUseMul_ = false;
47 49 
48 template <typename T>50 template <typename T>
49 ge::graphStatus GetConstData(uint32_t inputIdx, T& data);51 ge::graphStatus GetConstData(uint32_t inputIdx, T& data);
50 52 
51 float GetReciprocal(float data);53 float GetReciprocal(float data);
54+ 
55+ template <typename OpDag>
56+ ge::graphStatus ExecTiling(bool isScalarBranch = false);
57+ 
58+ ge::graphStatus GetScalarReciprocal(ge::DataType x2DType);
59+ 
60+ ge::graphStatus SelectAndExecTiling(ge::DataType x1DType, ge::DataType x2DType);
61+ 
62+ ge::graphStatus HandleFloat16WithFloat();
63+ ge::graphStatus HandleFloat16OrBf16();
64+ ge::graphStatus HandleFloat(ge::DataType x2DType);
65+ ge::graphStatus HandleIntTypes(ge::DataType x1DType, ge::DataType x2DType);
66+ ge::graphStatus HandleInt64();
52};67};
53 68 
54} // namespace optiling69} // namespace optiling
Mmath/truncate_div/op_host/config/ascend950/truncate_div_binary.json+144-36
@@ -10,7 +10,10 @@
10 "dtype": "bfloat16",10 "dtype": "bfloat16",
11 "format": "ND",11 "format": "ND",
12 "paramType": "required",12 "paramType": "required",
13- "shape": [-2]13+ "shape": [
14+ -2
15+ ],
16+ "format_match_mode": "FormatAgnostic"
14 },17 },
15 {18 {
16 "name": "x2",19 "name": "x2",
@@ -18,7 +21,10 @@
18 "dtype": "bfloat16",21 "dtype": "bfloat16",
19 "format": "ND",22 "format": "ND",
20 "paramType": "required",23 "paramType": "required",
21- "shape": [-2]24+ "shape": [
25+ -2
26+ ],
27+ "format_match_mode": "FormatAgnostic"
22 }28 }
23 ],29 ],
24 "outputs": [30 "outputs": [
@@ -28,7 +34,10 @@
28 "dtype": "bfloat16",34 "dtype": "bfloat16",
29 "format": "ND",35 "format": "ND",
30 "paramType": "required",36 "paramType": "required",
31- "shape": [-2]37+ "shape": [
38+ -2
39+ ],
40+ "format_match_mode": "FormatAgnostic"
32 }41 }
33 ]42 ]
34 },43 },
@@ -41,7 +50,10 @@
41 "dtype": "float16",50 "dtype": "float16",
42 "format": "ND",51 "format": "ND",
43 "paramType": "required",52 "paramType": "required",
44- "shape": [-2]53+ "shape": [
54+ -2
55+ ],
56+ "format_match_mode": "FormatAgnostic"
45 },57 },
46 {58 {
47 "name": "x2",59 "name": "x2",
@@ -49,7 +61,10 @@
49 "dtype": "float16",61 "dtype": "float16",
50 "format": "ND",62 "format": "ND",
51 "paramType": "required",63 "paramType": "required",
52- "shape": [-2]64+ "shape": [
65+ -2
66+ ],
67+ "format_match_mode": "FormatAgnostic"
53 }68 }
54 ],69 ],
55 "outputs": [70 "outputs": [
@@ -59,7 +74,10 @@
59 "dtype": "float16",74 "dtype": "float16",
60 "format": "ND",75 "format": "ND",
61 "paramType": "required",76 "paramType": "required",
62- "shape": [-2]77+ "shape": [
78+ -2
79+ ],
80+ "format_match_mode": "FormatAgnostic"
63 }81 }
64 ]82 ]
65 },83 },
@@ -72,7 +90,10 @@
72 "dtype": "float16",90 "dtype": "float16",
73 "format": "ND",91 "format": "ND",
74 "paramType": "required",92 "paramType": "required",
75- "shape": [-2]93+ "shape": [
94+ -2
95+ ],
96+ "format_match_mode": "FormatAgnostic"
76 },97 },
77 {98 {
78 "name": "x2",99 "name": "x2",
@@ -80,7 +101,10 @@
80 "dtype": "float32",101 "dtype": "float32",
81 "format": "ND",102 "format": "ND",
82 "paramType": "required",103 "paramType": "required",
83- "shape": [-2]104+ "shape": [
105+ -2
106+ ],
107+ "format_match_mode": "FormatAgnostic"
84 }108 }
85 ],109 ],
86 "outputs": [110 "outputs": [
@@ -90,7 +114,10 @@
90 "dtype": "float32",114 "dtype": "float32",
91 "format": "ND",115 "format": "ND",
92 "paramType": "required",116 "paramType": "required",
93- "shape": [-2]117+ "shape": [
118+ -2
119+ ],
120+ "format_match_mode": "FormatAgnostic"
94 }121 }
95 ]122 ]
96 },123 },
@@ -103,7 +130,10 @@
103 "dtype": "float32",130 "dtype": "float32",
104 "format": "ND",131 "format": "ND",
105 "paramType": "required",132 "paramType": "required",
106- "shape": [-2]133+ "shape": [
134+ -2
135+ ],
136+ "format_match_mode": "FormatAgnostic"
107 },137 },
108 {138 {
109 "name": "x2",139 "name": "x2",
@@ -111,7 +141,10 @@
111 "dtype": "float16",141 "dtype": "float16",
112 "format": "ND",142 "format": "ND",
113 "paramType": "required",143 "paramType": "required",
114- "shape": [-2]144+ "shape": [
145+ -2
146+ ],
147+ "format_match_mode": "FormatAgnostic"
115 }148 }
116 ],149 ],
117 "outputs": [150 "outputs": [
@@ -121,7 +154,10 @@
121 "dtype": "float32",154 "dtype": "float32",
122 "format": "ND",155 "format": "ND",
123 "paramType": "required",156 "paramType": "required",
124- "shape": [-2]157+ "shape": [
158+ -2
159+ ],
160+ "format_match_mode": "FormatAgnostic"
125 }161 }
126 ]162 ]
127 },163 },
@@ -134,7 +170,10 @@
134 "dtype": "float32",170 "dtype": "float32",
135 "format": "ND",171 "format": "ND",
136 "paramType": "required",172 "paramType": "required",
137- "shape": [-2]173+ "shape": [
174+ -2
175+ ],
176+ "format_match_mode": "FormatAgnostic"
138 },177 },
139 {178 {
140 "name": "x2",179 "name": "x2",
@@ -142,7 +181,10 @@
142 "dtype": "float32",181 "dtype": "float32",
143 "format": "ND",182 "format": "ND",
144 "paramType": "required",183 "paramType": "required",
145- "shape": [-2]184+ "shape": [
185+ -2
186+ ],
187+ "format_match_mode": "FormatAgnostic"
146 }188 }
147 ],189 ],
148 "outputs": [190 "outputs": [
@@ -152,7 +194,10 @@
152 "dtype": "float32",194 "dtype": "float32",
153 "format": "ND",195 "format": "ND",
154 "paramType": "required",196 "paramType": "required",
155- "shape": [-2]197+ "shape": [
198+ -2
199+ ],
200+ "format_match_mode": "FormatAgnostic"
156 }201 }
157 ]202 ]
158 },203 },
@@ -165,7 +210,10 @@
165 "dtype": "float32",210 "dtype": "float32",
166 "format": "ND",211 "format": "ND",
167 "paramType": "required",212 "paramType": "required",
168- "shape": [-2]213+ "shape": [
214+ -2
215+ ],
216+ "format_match_mode": "FormatAgnostic"
169 },217 },
170 {218 {
171 "name": "x2",219 "name": "x2",
@@ -173,7 +221,10 @@
173 "dtype": "int32",221 "dtype": "int32",
174 "format": "ND",222 "format": "ND",
175 "paramType": "required",223 "paramType": "required",
176- "shape": [-2]224+ "shape": [
225+ -2
226+ ],
227+ "format_match_mode": "FormatAgnostic"
177 }228 }
178 ],229 ],
179 "outputs": [230 "outputs": [
@@ -183,7 +234,10 @@
183 "dtype": "float32",234 "dtype": "float32",
184 "format": "ND",235 "format": "ND",
185 "paramType": "required",236 "paramType": "required",
186- "shape": [-2]237+ "shape": [
238+ -2
239+ ],
240+ "format_match_mode": "FormatAgnostic"
187 }241 }
188 ]242 ]
189 },243 },
@@ -196,7 +250,10 @@
196 "dtype": "int32",250 "dtype": "int32",
197 "format": "ND",251 "format": "ND",
198 "paramType": "required",252 "paramType": "required",
199- "shape": [-2]253+ "shape": [
254+ -2
255+ ],
256+ "format_match_mode": "FormatAgnostic"
200 },257 },
201 {258 {
202 "name": "x2",259 "name": "x2",
@@ -204,7 +261,10 @@
204 "dtype": "int32",261 "dtype": "int32",
205 "format": "ND",262 "format": "ND",
206 "paramType": "required",263 "paramType": "required",
207- "shape": [-2]264+ "shape": [
265+ -2
266+ ],
267+ "format_match_mode": "FormatAgnostic"
208 }268 }
209 ],269 ],
210 "outputs": [270 "outputs": [
@@ -214,7 +274,10 @@
214 "dtype": "int32",274 "dtype": "int32",
215 "format": "ND",275 "format": "ND",
216 "paramType": "required",276 "paramType": "required",
217- "shape": [-2]277+ "shape": [
278+ -2
279+ ],
280+ "format_match_mode": "FormatAgnostic"
218 }281 }
219 ]282 ]
220 },283 },
@@ -227,7 +290,10 @@
227 "dtype": "int32",290 "dtype": "int32",
228 "format": "ND",291 "format": "ND",
229 "paramType": "required",292 "paramType": "required",
230- "shape": [-2]293+ "shape": [
294+ -2
295+ ],
296+ "format_match_mode": "FormatAgnostic"
231 },297 },
232 {298 {
233 "name": "x2",299 "name": "x2",
@@ -235,7 +301,10 @@
235 "dtype": "float32",301 "dtype": "float32",
236 "format": "ND",302 "format": "ND",
237 "paramType": "required",303 "paramType": "required",
238- "shape": [-2]304+ "shape": [
305+ -2
306+ ],
307+ "format_match_mode": "FormatAgnostic"
239 }308 }
240 ],309 ],
241 "outputs": [310 "outputs": [
@@ -245,7 +314,10 @@
245 "dtype": "float32",314 "dtype": "float32",
246 "format": "ND",315 "format": "ND",
247 "paramType": "required",316 "paramType": "required",
248- "shape": [-2]317+ "shape": [
318+ -2
319+ ],
320+ "format_match_mode": "FormatAgnostic"
249 }321 }
250 ]322 ]
251 },323 },
@@ -258,7 +330,10 @@
258 "dtype": "uint8",330 "dtype": "uint8",
259 "format": "ND",331 "format": "ND",
260 "paramType": "required",332 "paramType": "required",
261- "shape": [-2]333+ "shape": [
334+ -2
335+ ],
336+ "format_match_mode": "FormatAgnostic"
262 },337 },
263 {338 {
264 "name": "x2",339 "name": "x2",
@@ -266,7 +341,10 @@
266 "dtype": "uint8",341 "dtype": "uint8",
267 "format": "ND",342 "format": "ND",
268 "paramType": "required",343 "paramType": "required",
269- "shape": [-2]344+ "shape": [
345+ -2
346+ ],
347+ "format_match_mode": "FormatAgnostic"
270 }348 }
271 ],349 ],
272 "outputs": [350 "outputs": [
@@ -276,7 +354,10 @@
276 "dtype": "uint8",354 "dtype": "uint8",
277 "format": "ND",355 "format": "ND",
278 "paramType": "required",356 "paramType": "required",
279- "shape": [-2]357+ "shape": [
358+ -2
359+ ],
360+ "format_match_mode": "FormatAgnostic"
280 }361 }
281 ]362 ]
282 },363 },
@@ -289,7 +370,10 @@
289 "dtype": "int8",370 "dtype": "int8",
290 "format": "ND",371 "format": "ND",
291 "paramType": "required",372 "paramType": "required",
292- "shape": [-2]373+ "shape": [
374+ -2
375+ ],
376+ "format_match_mode": "FormatAgnostic"
293 },377 },
294 {378 {
295 "name": "x2",379 "name": "x2",
@@ -297,7 +381,10 @@
297 "dtype": "int8",381 "dtype": "int8",
298 "format": "ND",382 "format": "ND",
299 "paramType": "required",383 "paramType": "required",
300- "shape": [-2]384+ "shape": [
385+ -2
386+ ],
387+ "format_match_mode": "FormatAgnostic"
301 }388 }
302 ],389 ],
303 "outputs": [390 "outputs": [
@@ -307,7 +394,10 @@
307 "dtype": "int8",394 "dtype": "int8",
308 "format": "ND",395 "format": "ND",
309 "paramType": "required",396 "paramType": "required",
310- "shape": [-2]397+ "shape": [
398+ -2
399+ ],
400+ "format_match_mode": "FormatAgnostic"
311 }401 }
312 ]402 ]
313 },403 },
@@ -320,7 +410,10 @@
320 "dtype": "int64",410 "dtype": "int64",
321 "format": "ND",411 "format": "ND",
322 "paramType": "required",412 "paramType": "required",
323- "shape": [-2]413+ "shape": [
414+ -2
415+ ],
416+ "format_match_mode": "FormatAgnostic"
324 },417 },
325 {418 {
326 "name": "x2",419 "name": "x2",
@@ -328,7 +421,10 @@
328 "dtype": "int64",421 "dtype": "int64",
329 "format": "ND",422 "format": "ND",
330 "paramType": "required",423 "paramType": "required",
331- "shape": [-2]424+ "shape": [
425+ -2
426+ ],
427+ "format_match_mode": "FormatAgnostic"
332 }428 }
333 ],429 ],
334 "outputs": [430 "outputs": [
@@ -338,7 +434,10 @@
338 "dtype": "int64",434 "dtype": "int64",
339 "format": "ND",435 "format": "ND",
340 "paramType": "required",436 "paramType": "required",
341- "shape": [-2]437+ "shape": [
438+ -2
439+ ],
440+ "format_match_mode": "FormatAgnostic"
342 }441 }
343 ]442 ]
344 },443 },
@@ -351,7 +450,10 @@
351 "dtype": "int16",450 "dtype": "int16",
352 "format": "ND",451 "format": "ND",
353 "paramType": "required",452 "paramType": "required",
354- "shape": [-2]453+ "shape": [
454+ -2
455+ ],
456+ "format_match_mode": "FormatAgnostic"
355 },457 },
356 {458 {
357 "name": "x2",459 "name": "x2",
@@ -359,7 +461,10 @@
359 "dtype": "int16",461 "dtype": "int16",
360 "format": "ND",462 "format": "ND",
361 "paramType": "required",463 "paramType": "required",
362- "shape": [-2]464+ "shape": [
465+ -2
466+ ],
467+ "format_match_mode": "FormatAgnostic"
363 }468 }
364 ],469 ],
365 "outputs": [470 "outputs": [
@@ -369,7 +474,10 @@
369 "dtype": "int16",474 "dtype": "int16",
370 "format": "ND",475 "format": "ND",
371 "paramType": "required",476 "paramType": "required",
372- "shape": [-2]477+ "shape": [
478+ -2
479+ ],
480+ "format_match_mode": "FormatAgnostic"
373 }481 }
374 ]482 ]
375 }483 }
Mmath/truncate_div/op_host/truncate_div_def.cpp+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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.
Mmath/truncate_div/op_host/truncate_div_infershape.cpp+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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.
Mmath/truncate_div/op_kernel/arch35/truncate_div_dag.h+3-7
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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.
@@ -30,15 +30,11 @@ namespace TruncateDivOp {
30using namespace Ops::Base;30using namespace Ops::Base;
31constexpr int TRUNCATE_DIV_CAST_MODE_NONE = 0;31constexpr int TRUNCATE_DIV_CAST_MODE_NONE = 0;
32constexpr int TRUNCATE_DIV_CAST_MODE_RINT = 1;32constexpr int TRUNCATE_DIV_CAST_MODE_RINT = 1;
33-constexpr int CAST_ROUND_MODE_TRUNC = 5;
34-constexpr int TRUNCATE_DIV_CMP_NE_MODE = 5;
35-constexpr int SEL_MODE_TENSOR_SCALAR = 1;
36constexpr int8_t SAT_POS = 60;33constexpr int8_t SAT_POS = 60;
37constexpr int64_t INT64_MAX_VALUE = 9223372036854775807;34constexpr int64_t INT64_MAX_VALUE = 9223372036854775807;
38constexpr int64_t INT32_MAX_VALUE = 2147483647;35constexpr int64_t INT32_MAX_VALUE = 2147483647;
39const uint32_t UINT32_SIGN = 0x80000000;36const uint32_t UINT32_SIGN = 0x80000000;
40const uint16_t UINT16_SIGN = 0x8000;37const uint16_t UINT16_SIGN = 0x8000;
41-const int16_t TRUNCATE_DIV_B16_SIGN = -32768;
42 38 
43namespace TruncDag1 {39namespace TruncDag1 {
44template <class T>40template <class T>
@@ -181,7 +177,7 @@ struct TruncIntPostCompute : public Vec::ElemwiseTernaryOP<T, T, T, T> {
181#ifdef __CCE_AICORE__177#ifdef __CCE_AICORE__
182template <typename T>178template <typename T>
183__simt_vf__ __aicore__179__simt_vf__ __aicore__
184- LAUNCH_BOUND(1024) inline void TruncDivInt_1(__ubuf__ T* dst, __ubuf__ T* src1, __ubuf__ T* src2, int count)180+ LAUNCH_BOUND(1024) inline void TruncDivInt_SIMT(__ubuf__ T* dst, __ubuf__ T* src1, __ubuf__ T* src2, int count)
185{181{
186 for (uint32_t index = static_cast<uint32_t>(threadIdx.x); index < count;182 for (uint32_t index = static_cast<uint32_t>(threadIdx.x); index < count;
187 index += static_cast<uint32_t>(blockDim.x)) {183 index += static_cast<uint32_t>(blockDim.x)) {
@@ -206,7 +202,7 @@ struct TruncDivInt64 : public Vec::ElemwiseBinaryOP<T, T, T> {
206 __ubuf__ T* dst_1 = (__ubuf__ T*)dst.GetPhyAddr();202 __ubuf__ T* dst_1 = (__ubuf__ T*)dst.GetPhyAddr();
207 __ubuf__ T* src1_1 = (__ubuf__ T*)src1.GetPhyAddr();203 __ubuf__ T* src1_1 = (__ubuf__ T*)src1.GetPhyAddr();
208 __ubuf__ T* src2_1 = (__ubuf__ T*)src2.GetPhyAddr();204 __ubuf__ T* src2_1 = (__ubuf__ T*)src2.GetPhyAddr();
209- asc_vf_call<TruncDivInt_1<T>>(dim3(1024), dst_1, src1_1, src2_1, count);205+ asc_vf_call<TruncDivInt_SIMT<T>>(dim3(1024), dst_1, src1_1, src2_1, count);
210#endif206#endif
211 }207 }
212};208};
Mmath/truncate_div/op_kernel/arch35/truncate_div_struct.h+1-1
@@ -1,5 +1,5 @@
1/**1/**
2-* Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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.
Mmath/truncate_div/op_kernel/truncate_div_apt.cpp+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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.
Mmath/truncate_div/tests/ut/op_host/arch35/test_truncate_div_tiling_arch35.cpp+2-22
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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 can not use this file except in compliance with the License.5 * Please refer to the License for details. You can not use this file except in compliance with the License.
@@ -98,26 +98,6 @@ TEST_F(TruncateDivTilingTest, truncate_div_bf16_1)
98 ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);98 ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
99}99}
100 100 
101-TEST_F(TruncateDivTilingTest, truncate_div_int8_scalar)
102-{
103- BroadcastCompileInfo compileInfo;
104- compileInfo.coreNum = 64;
105- compileInfo.ubSize = 245760;
106- gert::TilingContextPara tilingContextPara(
107- "TruncateDiv",
108- {
109- {{{17772, 1, 2, 1, 2, 1, 2, 1}, {17772, 1, 2, 1, 2, 1, 2, 1}}, ge::DT_INT8, ge::FORMAT_ND},
110- {{{}, {}}, ge::DT_INT16, ge::FORMAT_ND},
111- },
112- {
113- {{{17772, 2, 2, 2, 2, 2, 2, 2}, {17772, 2, 2, 2, 2, 2, 2, 2}}, ge::DT_FLOAT16, ge::FORMAT_ND},
114- },
115- &compileInfo);
116- uint64_t expectTilingKey = 0b0'00000000'00000001;
117- std::vector<size_t> expectWorkspaces = {16777216};
118- ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
119-}
120- 
121TEST_F(TruncateDivTilingTest, truncate_div_f32_scalar_3)101TEST_F(TruncateDivTilingTest, truncate_div_f32_scalar_3)
122{102{
123 BroadcastCompileInfo compileInfo;103 BroadcastCompileInfo compileInfo;
@@ -202,7 +182,7 @@ TEST_F(TruncateDivTilingTest, test_ez0020_scalar_dtype_int32)
202 uint64_t expectTilingKey = 0;182 uint64_t expectTilingKey = 0;
203 std::vector<size_t> expectWorkspaces = {0};183 std::vector<size_t> expectWorkspaces = {0};
204 184 
205- ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);185+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectWorkspaces);
206}186}
207 187 
208TEST_F(TruncateDivTilingTest, test_ez0021_dtype_combination_double)188TEST_F(TruncateDivTilingTest, test_ez0021_dtype_combination_double)
Mmath/truncate_div/tests/ut/op_host/test_truncate_div_infershape.cpp+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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 can not use this file except in compliance with the License.5 * Please refer to the License for details. You can not use this file except in compliance with the License.