Star1295
1409
代码
代码
Issues104
Pull Requests209
Actions
项目讨论
Wiki
分析
Star1295
1409
已合并
feat: 新增truncate_div A5 实现 #2798
yefeicoding创建于 5月18日
feat: 新增truncate_div A5 实现 #2798
已合并
yefeicoding创建于 5月18日
从yefeicoding/ops-mathtruncate_div合入到cann/ops-mathmaster
共 21 个文件变更+2343-45
Mdocs/zh/op_list.md+10-0
@@ -586,6 +586,16 @@
586 <td>AI Core</td>586 <td>AI Core</td>
587 <td>张量除法计算。</td>587 <td>张量除法计算。</td>
588 </tr>588 </tr>
589+ <tr>
590+ <td>math</td>
591+ <td><a href="../../math/truncate_div/README.md">truncate_div</a></td>
592+ <td>√</td>
593+ <td>√</td>
594+ <td>√</td>
595+ <td>√</td>
596+ <td>AI Core</td>
597+ <td>完成截断除法计算，结果向零取整。</td>
598+ </tr>
589 <tr>599 <tr>
590 <td>math</td>600 <td>math</td>
591 <td><a href="../../math/div_no_nan/README.md">div_no_nan</a></td>601 <td><a href="../../math/div_no_nan/README.md">div_no_nan</a></td>
Mmath/div/CMakeLists.txt+1-1
@@ -11,7 +11,7 @@
11set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")11set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950" "mc62")
12set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")12set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35" "arch35")
13 13 
14-add_all_modules_sources(OPTYPE div ACLNNTYPE aclnn_exclude14+add_all_modules_sources(OPTYPE div ACLNNTYPE aclnn_exclude DEPENDENCIES truncate_div
15 COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}15 COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
16 TILING_DIR ${SUPPORT_TILING_DIR}16 TILING_DIR ${SUPPORT_TILING_DIR}
17 DISABLE_IN_OPP TRUE)17 DISABLE_IN_OPP TRUE)
Mmath/div/op_api/aclnn_div.cpp+126-44
@@ -15,6 +15,7 @@
15#include "math/real_div/op_api/realdiv.h"15#include "math/real_div/op_api/realdiv.h"
16#include "math/trunc/op_api/trunc.h"16#include "math/trunc/op_api/trunc.h"
17#include "math/muls/op_api/muls.h"17#include "math/muls/op_api/muls.h"
18+#include "math/truncate_div/op_api/truncate_div.h"
18#include "op_api/op_api_def.h"19#include "op_api/op_api_def.h"
19#include "op_api/aclnn_check.h"20#include "op_api/aclnn_check.h"
20#include "aclnn_kernels/common/op_error_check.h"21#include "aclnn_kernels/common/op_error_check.h"
@@ -33,11 +34,12 @@ using namespace op;
33extern "C" {34extern "C" {
34#endif35#endif
35 36 
36-op::DataType PromoteIntegerInputsToFloat(const op::DataType input) {37+op::DataType PromoteIntegerInputsToFloat(const op::DataType input)
37- if (IsIntegralType(input)) {38+{
38- return op::DataType::DT_FLOAT;39+ if (IsIntegralType(input)) {
39- }40+ return op::DataType::DT_FLOAT;
40- return input;41+ }
42+ return input;
41}43}
42 44 
43static op::DataType InnerTypeToComplexType(const op::DataType input)45static op::DataType InnerTypeToComplexType(const op::DataType input)
@@ -114,13 +116,31 @@ static const int MODE_REAL_DIV = 0;
114static const int MODE_TRUNC_DIV = 1;116static const int MODE_TRUNC_DIV = 1;
115static const int MODE_FLOOR_DIV = 2;117static const int MODE_FLOOR_DIV = 2;
116 118 
119+// TruncateDiv 支持的类型组合映射表（不需要类型提升）
120+static const std::initializer_list<std::pair<op::DataType, op::DataType>> TRUNC_DTYPE_MAPPING = {
121+ {op::DataType::DT_BF16, op::DataType::DT_BF16}, {op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT16},
122+ {op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT}, {op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16},
123+ {op::DataType::DT_FLOAT, op::DataType::DT_FLOAT}, {op::DataType::DT_FLOAT, op::DataType::DT_INT32},
124+ {op::DataType::DT_INT32, op::DataType::DT_INT32}, {op::DataType::DT_INT32, op::DataType::DT_FLOAT},
125+ {op::DataType::DT_UINT8, op::DataType::DT_UINT8}, {op::DataType::DT_INT8, op::DataType::DT_INT8},
126+ {op::DataType::DT_INT64, op::DataType::DT_INT64}, {op::DataType::DT_INT16, op::DataType::DT_INT16}};
127+ 
128+static bool isInTruncDtypeMapping(const op::DataType selfDtype, const op::DataType otherDtype)
129+{
130+ for (const auto& pair : TRUNC_DTYPE_MAPPING) {
131+ if (pair.first == selfDtype && pair.second == otherDtype) {
132+ return true;
133+ }
134+ }
135+ return false;
136+}
137+ 
117static const std::initializer_list<std::pair<op::DataType, op::DataType>> AllowedMixDtypePairs = {138static const std::initializer_list<std::pair<op::DataType, op::DataType>> AllowedMixDtypePairs = {
118- {op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT16}, {op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT}, 139+ {op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT16}, {op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT},
119- {op::DataType::DT_FLOAT16, op::DataType::DT_BF16}, {op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16}, 140+ {op::DataType::DT_FLOAT16, op::DataType::DT_BF16}, {op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16},
120- {op::DataType::DT_FLOAT, op::DataType::DT_FLOAT}, {op::DataType::DT_FLOAT, op::DataType::DT_BF16}, 141+ {op::DataType::DT_FLOAT, op::DataType::DT_FLOAT}, {op::DataType::DT_FLOAT, op::DataType::DT_BF16},
121- {op::DataType::DT_BF16, op::DataType::DT_FLOAT16}, {op::DataType::DT_BF16, op::DataType::DT_FLOAT}, 142+ {op::DataType::DT_BF16, op::DataType::DT_FLOAT16}, {op::DataType::DT_BF16, op::DataType::DT_FLOAT},
122- {op::DataType::DT_BF16, op::DataType::DT_BF16}143+ {op::DataType::DT_BF16, op::DataType::DT_BF16}};
123-};
124 144 
125static const std::initializer_list<DataType>& GetDtypeSupportList()145static const std::initializer_list<DataType>& GetDtypeSupportList()
126{146{
@@ -164,13 +184,14 @@ static inline aclnnStatus CheckDivModComplexDtype(const op::DataType promoteType
164static inline op::DataType InferDivModeDtype(184static inline op::DataType InferDivModeDtype(
165 const op::DataType selfDtype, const op::DataType otherDtype, const int mode)185 const op::DataType selfDtype, const op::DataType otherDtype, const int mode)
166{186{
187+ auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();
167 auto promoteType = op::PromoteType(selfDtype, otherDtype);188 auto promoteType = op::PromoteType(selfDtype, otherDtype);
168 // 下沉PTA入口操作将入参类型转化成FLOAT进行后续处理189 // 下沉PTA入口操作将入参类型转化成FLOAT进行后续处理
169 if (mode == MODE_REAL_DIV && promoteType != op::DataType::DT_INT32 && promoteType != op::DataType::DT_BOOL) {190 if (mode == MODE_REAL_DIV && promoteType != op::DataType::DT_INT32 && promoteType != op::DataType::DT_BOOL) {
170 // IterateBase 配置特殊处理191 // IterateBase 配置特殊处理
171 promoteType = PromoteIntegerInputsToFloat(promoteType);192 promoteType = PromoteIntegerInputsToFloat(promoteType);
172 }193 }
173- if (mode == MODE_TRUNC_DIV && promoteType == DataType::DT_DOUBLE) {194+ if (mode == MODE_TRUNC_DIV && promoteType == DataType::DT_DOUBLE && !IsRegBase(npuArch)) {
174 promoteType = DataType::DT_FLOAT;195 promoteType = DataType::DT_FLOAT;
175 }196 }
176 return promoteType;197 return promoteType;
@@ -276,9 +297,8 @@ static bool CheckPromoteType(const aclTensor* self, const aclTensor* other, cons
276{297{
277 // 检查self和other能否做数据类型推导298 // 检查self和other能否做数据类型推导
278 auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();299 auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();
279- auto promoteType = (IsRegBase(npuArch)) ?300+ auto promoteType = (IsRegBase(npuArch)) ? InferDivModeDtype(self->GetDataType(), other->GetDataType(), mode) :
280- InferDivModeDtype(self->GetDataType(), other->GetDataType(), mode) :301+ op::PromoteType(self->GetDataType(), other->GetDataType());
281- op::PromoteType(self->GetDataType(), other->GetDataType());
282 if (promoteType == DataType::DT_UNDEFINED) {302 if (promoteType == DataType::DT_UNDEFINED) {
283 OP_LOGE(303 OP_LOGE(
284 ACLNN_ERR_PARAM_INVALID, "Self dtype %s and other dtype %s can not promote dtype.",304 ACLNN_ERR_PARAM_INVALID, "Self dtype %s and other dtype %s can not promote dtype.",
@@ -287,9 +307,8 @@ static bool CheckPromoteType(const aclTensor* self, const aclTensor* other, cons
287 }307 }
288 308 
289 // 检查推导后的数据类型能否转换为输出的数据类型309 // 检查推导后的数据类型能否转换为输出的数据类型
290- bool outDtypeToFloat = IsRegBase()310+ bool outDtypeToFloat = IsRegBase() && mode == MODE_REAL_DIV &&
291- && mode == MODE_REAL_DIV311+ (promoteType == op::DataType::DT_INT32 || promoteType == op::DataType::DT_BOOL);
292- && (promoteType == op::DataType::DT_INT32 || promoteType == op::DataType::DT_BOOL);
293 auto computeDtype = outDtypeToFloat ? op::DataType::DT_FLOAT : promoteType;312 auto computeDtype = outDtypeToFloat ? op::DataType::DT_FLOAT : promoteType;
294 OP_CHECK_RESULT_DTYPE_CAST_FAILED(computeDtype, y->GetDataType(), return false);313 OP_CHECK_RESULT_DTYPE_CAST_FAILED(computeDtype, y->GetDataType(), return false);
295 return true;314 return true;
@@ -390,11 +409,11 @@ inline static bool isDivsMixDtypeSupport(const aclTensor* self, const aclScalar*
390 (self->GetDataType() == DataType::DT_BF16 && other->GetDataType() == DataType::DT_FLOAT16);409 (self->GetDataType() == DataType::DT_BF16 && other->GetDataType() == DataType::DT_FLOAT16);
391}410}
392 411 
393-inline static bool checkMixDtypeConditions(DataType selfDtype, DataType otherDtype){412+inline static bool checkMixDtypeConditions(DataType selfDtype, DataType otherDtype)
413+{
394 return std::find(414 return std::find(
395 AllowedMixDtypePairs.begin(), AllowedMixDtypePairs.end(),415 AllowedMixDtypePairs.begin(), AllowedMixDtypePairs.end(),
396- std::pair<op::DataType, op::DataType>(selfDtype, otherDtype)) !=416+ std::pair<op::DataType, op::DataType>(selfDtype, otherDtype)) != AllowedMixDtypePairs.end();
397- AllowedMixDtypePairs.end();
398}417}
399 418 
400inline static bool isMixDtypeScalarSupport(const aclTensor* self, const aclScalar* other)419inline static bool isMixDtypeScalarSupport(const aclTensor* self, const aclScalar* other)
@@ -416,8 +435,8 @@ inline static bool isMixDtypeTensorSupport(const aclTensor* self, const aclTenso
416}435}
417 436 
418static aclnnStatus HandleMixDataTypeDiv(437static aclnnStatus HandleMixDataTypeDiv(
419- const aclTensor* self, const aclTensor* other, aclOpExecutor* executor, const aclTensor** divOpOut438+ const aclTensor* self, const aclTensor* other, aclOpExecutor* executor, const aclTensor** divOpOut)
420-) {439+{
421 // 固定写法，将输入self转换成连续的tensor440 // 固定写法，将输入self转换成连续的tensor
422 auto selfContiguous = l0op::Contiguous(self, executor);441 auto selfContiguous = l0op::Contiguous(self, executor);
423 CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);442 CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
@@ -433,13 +452,13 @@ static aclnnStatus HandleMixDataTypeDiv(
433}452}
434 453 
435static aclnnStatus HandleNotMixDataTypeDiv(454static aclnnStatus HandleNotMixDataTypeDiv(
436- const aclTensor* self, const aclTensor* other, aclOpExecutor* executor, const aclTensor** divOpOut455+ const aclTensor* self, const aclTensor* other, aclOpExecutor* executor, const aclTensor** divOpOut)
437-) {456+{
438 // RealDiv算子需要对self和other两个输入做隐式数据类型转换，根据具体算子语义按需调用457 // RealDiv算子需要对self和other两个输入做隐式数据类型转换，根据具体算子语义按需调用
439 auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();458 auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();
440 auto promoteType = (!IsRegBase(npuArch)) ?459 auto promoteType = (!IsRegBase(npuArch)) ?
441- CompatibleInferDivDtype(self->GetDataType(), other->GetDataType()) :460+ CompatibleInferDivDtype(self->GetDataType(), other->GetDataType()) :
442- InferDivModeDtype(self->GetDataType(), other->GetDataType(), MODE_REAL_DIV);461+ InferDivModeDtype(self->GetDataType(), other->GetDataType(), MODE_REAL_DIV);
443 462 
444 // 处理self输入463 // 处理self输入
445 const aclTensor* selfProcessed = nullptr;464 const aclTensor* selfProcessed = nullptr;
@@ -653,13 +672,17 @@ aclnnStatus aclnnDivsGetWorkspaceSize(
653 CompatibleInferDivsDtype(self->GetDataType(), other->GetDataType()) :672 CompatibleInferDivsDtype(self->GetDataType(), other->GetDataType()) :
654 InferDivsModeDtype(self->GetDataType(), other->GetDataType(), MODE_REAL_DIV);673 InferDivsModeDtype(self->GetDataType(), other->GetDataType(), MODE_REAL_DIV);
655 promoteType = (IsFloatingType(self->GetDataType()) || IsComplexType(self->GetDataType())) ?674 promoteType = (IsFloatingType(self->GetDataType()) || IsComplexType(self->GetDataType())) ?
656- self->GetDataType() : op::DataType::DT_FLOAT;675+ self->GetDataType() :
676+ op::DataType::DT_FLOAT;
657 promoteType = (self->GetDataType() == op::DataType::DT_BOOL && other->GetDataType() == op::DataType::DT_BOOL) ?677 promoteType = (self->GetDataType() == op::DataType::DT_BOOL && other->GetDataType() == op::DataType::DT_BOOL) ?
658- self->GetDataType() : promoteType;678+ self->GetDataType() :
659- promoteType = (IsComplexType(other->GetDataType())) ? op::PromoteType(promoteType, other->GetDataType()) : promoteType;679+ promoteType;
680+ promoteType =
681+ (IsComplexType(other->GetDataType())) ? op::PromoteType(promoteType, other->GetDataType()) : promoteType;
660 if (IsRegBase(npuArch)) {682 if (IsRegBase(npuArch)) {
661- promoteType = op::PromoteType(self->GetDataType(), other->GetDataType()) == op::DataType::DT_INT32683+ promoteType = op::PromoteType(self->GetDataType(), other->GetDataType()) == op::DataType::DT_INT32 ?
662- ? op::DataType::DT_INT32 : promoteType;684+ op::DataType::DT_INT32 :
685+ promoteType;
663 }686 }
664 687 
665 bool canUseMuls = CanUseMuls(self, other);688 bool canUseMuls = CanUseMuls(self, other);
@@ -736,7 +759,36 @@ aclnnStatus aclnnDivModGetWorkspaceSize(
736 759 
737 bool isMixDataType = isMixDtypeTensorSupport(self, other);760 bool isMixDataType = isMixDtypeTensorSupport(self, other);
738 const aclTensor* divOpOut = nullptr;761 const aclTensor* divOpOut = nullptr;
739- if (isMixDataType) {762+ auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();
763+ 
764+ // TruncateDiv 特殊处理：IsRegBase && 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())) {
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());
774+ } else {
775+ op::DataType promoteType;
776+ promoteType = InferDivModeDtype(self->GetDataType(), other->GetDataType(), mode);
777+ bool needToFloat = (promoteType == op::DataType::DT_BOOL);
778+ promoteType = needToFloat ? op::DataType::DT_FLOAT : promoteType;
779+ OP_LOGI(
780+ "aclnnDivMod", "TruncateDiv cast path: selfDtype=%s -> %s, otherDtype=%s -> %s, promoteType=%s",
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());
785+ CHECK_RET(selfCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
786+ otherCasted = l0op::Cast(otherContiguous, promoteType, uniqueExecutor.get());
787+ CHECK_RET(otherCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
788+ divOpOut = l0op::TruncateDiv(selfCasted, otherCasted, uniqueExecutor.get());
789+ }
790+ CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
791+ } else if (isMixDataType) {
740 if (mode == MODE_FLOOR_DIV) {792 if (mode == MODE_FLOOR_DIV) {
741 divOpOut = l0op::FloorDiv(selfCasted, otherCasted, false, uniqueExecutor.get());793 divOpOut = l0op::FloorDiv(selfCasted, otherCasted, false, uniqueExecutor.get());
742 } else {794 } else {
@@ -752,7 +804,6 @@ aclnnStatus aclnnDivModGetWorkspaceSize(
752 op::DataType promoteType;804 op::DataType promoteType;
753 bool needToInt32 = false;805 bool needToInt32 = false;
754 op::DataType oriType = out->GetDataType();806 op::DataType oriType = out->GetDataType();
755- auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();
756 if (!IsRegBase(npuArch)) {807 if (!IsRegBase(npuArch)) {
757 auto promoteRet = CompatibleInferDivModeDtype(self->GetDataType(), other->GetDataType(), mode, promoteType);808 auto promoteRet = CompatibleInferDivModeDtype(self->GetDataType(), other->GetDataType(), mode, promoteType);
758 CHECK_RET(promoteRet == ACLNN_SUCCESS, promoteRet);809 CHECK_RET(promoteRet == ACLNN_SUCCESS, promoteRet);
@@ -765,9 +816,9 @@ aclnnStatus aclnnDivModGetWorkspaceSize(
765 promoteType = needToFloat ? op::DataType::DT_FLOAT : promoteType;816 promoteType = needToFloat ? op::DataType::DT_FLOAT : promoteType;
766 // aicore is not supported, aicpu has problems when div 0817 // aicore is not supported, aicpu has problems when div 0
767 needToInt32 = (promoteType == op::DataType::DT_INT16 && mode == MODE_FLOOR_DIV) ||818 needToInt32 = (promoteType == op::DataType::DT_INT16 && mode == MODE_FLOOR_DIV) ||
768- ((promoteType == op::DataType::DT_INT8 || promoteType == op::DataType::DT_UINT8 ||819+ ((promoteType == op::DataType::DT_INT8 || promoteType == op::DataType::DT_UINT8 ||
769 promoteType == op::DataType::DT_INT16) &&820 promoteType == op::DataType::DT_INT16) &&
770- mode == MODE_TRUNC_DIV);821+ mode == MODE_TRUNC_DIV);
771 oriType = promoteType;822 oriType = promoteType;
772 promoteType = needToInt32 ? op::DataType::DT_INT32 : promoteType;823 promoteType = needToInt32 ? op::DataType::DT_INT32 : promoteType;
773 }824 }
@@ -792,7 +843,7 @@ aclnnStatus aclnnDivModGetWorkspaceSize(
792 CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);843 CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
793 }844 }
794 }845 }
795- 846+ 
796 auto castOut = l0op::Cast(divOpOut, out->GetDataType(), uniqueExecutor.get());847 auto castOut = l0op::Cast(divOpOut, out->GetDataType(), uniqueExecutor.get());
797 CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);848 CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
798 849 
@@ -836,7 +887,38 @@ aclnnStatus aclnnDivModsGetWorkspaceSize(
836 auto selfCasted = selfContiguous;887 auto selfCasted = selfContiguous;
837 bool isMixDataType = isMixDtypeScalarSupport(self, other);888 bool isMixDataType = isMixDtypeScalarSupport(self, other);
838 const aclTensor* divOpOut = nullptr;889 const aclTensor* divOpOut = nullptr;
839- if (isMixDataType) {890+ auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();
891+ 
892+ // TruncateDiv 特殊处理：IsRegBase && mode=MODE_TRUNC_DIV && 类型组合在映射表中，不做类型提升
893+ 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())) {
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());
902+ CHECK_RET(otherConvert != nullptr, ACLNN_ERR_INNER_NULLPTR);
903+ divOpOut = l0op::TruncateDiv(selfContiguous, otherConvert, uniqueExecutor.get());
904+ } else {
905+ op::DataType promoteType;
906+ promoteType = InferDivModeDtype(self->GetDataType(), other->GetDataType(), mode);
907+ bool needToFloat = (promoteType == op::DataType::DT_BOOL);
908+ promoteType = needToFloat ? op::DataType::DT_FLOAT : promoteType;
909+ OP_LOGI(
910+ "aclnnDivMods", "TruncateDiv cast path: selfDtype=%s -> %s, otherDtype=%s -> %s, promoteType=%s",
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());
915+ CHECK_RET(selfCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
916+ auto otherCasted = uniqueExecutor.get()->ConvertToTensor(other, promoteType);
917+ CHECK_RET(otherCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
918+ divOpOut = l0op::TruncateDiv(selfCasted, otherCasted, uniqueExecutor.get());
919+ }
920+ CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
921+ } else if (isMixDataType) {
840 auto otherConvert = uniqueExecutor.get()->ConvertToTensor(other, other->GetDataType());922 auto otherConvert = uniqueExecutor.get()->ConvertToTensor(other, other->GetDataType());
841 CHECK_RET(otherConvert != nullptr, ACLNN_ERR_INNER_NULLPTR);923 CHECK_RET(otherConvert != nullptr, ACLNN_ERR_INNER_NULLPTR);
842 if (mode == MODE_FLOOR_DIV) {924 if (mode == MODE_FLOOR_DIV) {
@@ -853,9 +935,9 @@ aclnnStatus aclnnDivModsGetWorkspaceSize(
853 op::DataType promoteType;935 op::DataType promoteType;
854 bool needToInt32 = false;936 bool needToInt32 = false;
855 op::DataType oriType = out->GetDataType();937 op::DataType oriType = out->GetDataType();
856- auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();
857 if (!IsRegBase(npuArch)) {938 if (!IsRegBase(npuArch)) {
858- auto promoteRet = CompatibleInferDivsModeDtype(self->GetDataType(), other->GetDataType(), mode, promoteType);939+ auto promoteRet =
940+ CompatibleInferDivsModeDtype(self->GetDataType(), other->GetDataType(), mode, promoteType);
859 CHECK_RET(promoteRet == ACLNN_SUCCESS, promoteRet);941 CHECK_RET(promoteRet == ACLNN_SUCCESS, promoteRet);
860 } else {942 } else {
861 promoteType = InferDivsModeDtype(self->GetDataType(), other->GetDataType(), mode);943 promoteType = InferDivsModeDtype(self->GetDataType(), other->GetDataType(), mode);
@@ -866,9 +948,9 @@ aclnnStatus aclnnDivModsGetWorkspaceSize(
866 promoteType = needToFloat ? op::DataType::DT_FLOAT : promoteType;948 promoteType = needToFloat ? op::DataType::DT_FLOAT : promoteType;
867 // aicore is not supported, aicpu has problems when div 0949 // aicore is not supported, aicpu has problems when div 0
868 needToInt32 = (promoteType == op::DataType::DT_INT16 && mode == MODE_FLOOR_DIV) ||950 needToInt32 = (promoteType == op::DataType::DT_INT16 && mode == MODE_FLOOR_DIV) ||
869- ((promoteType == op::DataType::DT_INT8 || promoteType == op::DataType::DT_UINT8 ||951+ ((promoteType == op::DataType::DT_INT8 || promoteType == op::DataType::DT_UINT8 ||
870 promoteType == op::DataType::DT_INT16) &&952 promoteType == op::DataType::DT_INT16) &&
871- mode == MODE_TRUNC_DIV);953+ mode == MODE_TRUNC_DIV);
872 oriType = promoteType;954 oriType = promoteType;
873 promoteType = needToInt32 ? op::DataType::DT_INT32 : promoteType;955 promoteType = needToInt32 ? op::DataType::DT_INT32 : promoteType;
874 }956 }
@@ -893,7 +975,7 @@ aclnnStatus aclnnDivModsGetWorkspaceSize(
893 CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);975 CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
894 }976 }
895 }977 }
896- 978+ 
897 auto castOut = l0op::Cast(divOpOut, out->GetDataType(), uniqueExecutor.get());979 auto castOut = l0op::Cast(divOpOut, out->GetDataType(), uniqueExecutor.get());
898 CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);980 CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
899 981 
@@ -1045,7 +1127,7 @@ aclnnStatus aclnnInplaceDivModGetWorkspaceSize(
1045 divOpOut = l0op::InplaceTrunc(divOpOut, uniqueExecutor.get());1127 divOpOut = l0op::InplaceTrunc(divOpOut, uniqueExecutor.get());
1046 }1128 }
1047 CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);1129 CHECK_RET(divOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
1048- 1130+ 
1049 auto castOut = l0op::Cast(divOpOut, out->GetDataType(), uniqueExecutor.get());1131 auto castOut = l0op::Cast(divOpOut, out->GetDataType(), uniqueExecutor.get());
1050 CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);1132 CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
1051 1133 
Amath/truncate_div/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
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+set(SUPPORT_COMPUTE_UNIT "ascend950")
12+set(SUPPORT_TILING_DIR "arch35")
13+ 
14+add_all_modules_sources(OPTYPE truncate_div ACLNNTYPE aclnn_exclude
15+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
16+ TILING_DIR ${SUPPORT_TILING_DIR}
17+ DISABLE_IN_OPP TRUE)
Amath/truncate_div/README.md+80-0
@@ -0,0 +1,80 @@
1+# TruncateDiv
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10+| <term>Atlas 200I/500 A2 推理产品</term> | √ |
11+| <term>Atlas 推理系列产品</term> | × |
12+| <term>Atlas 训练系列产品</term> | √ |
13+ 
14+ 
15+## 功能说明
16+ 
17+- 接口功能：完成截断除法计算，结果向零取整
18+ 
19+- 计算公式：
20+ 
21+ $$
22+ out_i = trunc(\frac{x1_i}{x2_i})
23+ $$
24+ 
25+ 其中 `trunc` 表示向零取整（截断取整）。
26+ 
27+- 例外说明：无
28+ 
29+## 参数说明
30+ 
31+<table style="undefined;table-layout: fixed; width: 1005px"><colgroup>
32+ <col style="width: 170px">
33+ <col style="width: 170px">
34+ <col style="width: 352px">
35+ <col style="width: 213px">
36+ <col style="width: 100px">
37+ </colgroup>
38+ <thead>
39+ <tr>
40+ <th>参数名</th>
41+ <th>输入/输出/属性</th>
42+ <th>描述</th>
43+ <th>数据类型</th>
44+ <th>数据格式</th>
45+ </tr></thead>
46+ <tbody>
47+ <tr>
48+ <td>x1</td>
49+ <td>输入</td>
50+ <td>公式中的被除数。</td>
51+ <td>FLOAT16、FLOAT、INT32、UINT8、INT8、INT64、INT16</td>
52+ <td>ND</td>
53+ </tr>
54+ <tr>
55+ <td>x2</td>
56+ <td>输入</td>
57+ <td>公式中的除数。</td>
58+ <td>FLOAT16、FLOAT、INT32、UINT8、INT8、INT64、INT16</td>
59+ <td>ND</td>
60+ </tr>
61+ <tr>
62+ <td>y</td>
63+ <td>输出</td>
64+ <td>公式中的out，截断除法结果。</td>
65+ <td>FLOAT16、FLOAT、INT32、UINT8、INT8、INT64、INT16</td>
66+ <td>ND</td>
67+ </tr>
68+ </tbody></table>
69+ 
70+## 约束说明
71+ 
72+- 输入x1和x2需满足broadcast关系
73+- 当除数为0时，结果为未定义行为
74+ 
75+ 
76+## 调用说明
77+ 
78+| 调用方式 | 样例代码 | 说明 |
79+| ---------------- | --------------------------- | --------------------------------------------------- |
80+| 图模式调用 | [test_geir_truncate_div](./examples/test_geir_truncate_div.cpp) | 通过[算子IR](./op_graph/truncate_div_proto.h)构图方式调用TruncateDiv算子。 |
Amath/truncate_div/examples/test_geir_truncate_div.cpp+286-0
@@ -0,0 +1,286 @@
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 can 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 test_geir_truncate_div.cpp
13+ * \brief TruncateDiv算子GEIR测试示例
14+ */
15+ 
16+#include <iostream>
17+#include <fstream>
18+#include <string.h>
19+#include <stdint.h>
20+#include <vector>
21+#include <string>
22+#include <map>
23+#include "assert.h"
24+ 
25+#include "graph.h"
26+#include "types.h"
27+#include "tensor.h"
28+#include "ge_error_codes.h"
29+#include "ge_api_types.h"
30+#include "ge_api.h"
31+#include "array_ops.h"
32+#include "ge_ir_build.h"
33+ 
34+#include "experiment_ops.h"
35+#include "nn_other.h"
36+#include "../op_graph/truncate_div_proto.h"
37+ 
38+#define FAILED -1
39+#define SUCCESS 0
40+ 
41+using namespace ge;
42+using std::map;
43+using std::string;
44+using std::vector;
45+ 
46+#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape, value) \
47+ do { \
48+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
49+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
50+ TensorDesc placeholder##intputIndex##_desc = \
51+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
52+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
53+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
54+ Tensor tensor_placeholder##intputIndex; \
55+ ret = GenOnesDataFloat32( \
56+ placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, placeholder##intputIndex##_desc, \
57+ value); \
58+ if (ret != SUCCESS) { \
59+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
60+ return FAILED; \
61+ } \
62+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
63+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
64+ input.push_back(tensor_placeholder##intputIndex); \
65+ graph.AddOp(placeholder##intputIndex); \
66+ truncateDiv1.set_input_##intputName(placeholder##intputIndex); \
67+ inputs.push_back(placeholder##intputIndex); \
68+ } while (0)
69+ 
70+#define LOG_PRINT(message, ...) \
71+ do { \
72+ printf(message, ##__VA_ARGS__); \
73+ } while (0)
74+ 
75+string GetTime()
76+{
77+ time_t timep;
78+ time(&timep);
79+ char tmp[64];
80+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
81+ return tmp;
82+}
83+ 
84+uint32_t GetDataTypeSize(DataType dt)
85+{
86+ uint32_t dilation = 1;
87+ uint32_t oneByte = 1;
88+ uint32_t twoByte = 2;
89+ uint32_t fourByte = 4;
90+ uint32_t eightByte = 8;
91+ 
92+ if (dt == ge::DT_FLOAT) {
93+ dilation = fourByte;
94+ } else if (dt == ge::DT_FLOAT16) {
95+ dilation = twoByte;
96+ } else if (dt == ge::DT_BF16) {
97+ dilation = twoByte;
98+ } else if (dt == ge::DT_INT16) {
99+ dilation = twoByte;
100+ } else if (dt == ge::DT_UINT16) {
101+ dilation = twoByte;
102+ } else if (dt == ge::DT_INT32) {
103+ dilation = fourByte;
104+ } else if (dt == ge::DT_UINT32) {
105+ dilation = fourByte;
106+ } else if (dt == ge::DT_INT64) {
107+ dilation = eightByte;
108+ } else if (dt == ge::DT_UINT64) {
109+ dilation = eightByte;
110+ } else if (dt == ge::DT_INT8) {
111+ dilation = oneByte;
112+ }
113+ return dilation;
114+}
115+ 
116+int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, float value)
117+{
118+ input_tensor_desc.SetRealDimCnt(shapes.size());
119+ size_t size = 1;
120+ for (uint32_t i = 0; i < shapes.size(); i++) {
121+ size *= shapes[i];
122+ }
123+ uint32_t byteSizeFloat32 = 4;
124+ uint32_t data_len = size * byteSizeFloat32;
125+ float* pData = new (std::nothrow) float[size];
126+ 
127+ for (size_t i = 0; i < size; ++i) {
128+ *(pData + i) = value;
129+ }
130+ input_tensor = Tensor(input_tensor_desc, (uint8_t*)pData, data_len);
131+ return SUCCESS;
132+}
133+ 
134+int32_t GenOnesData(
135+ vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, DataType data_type, int value)
136+{
137+ input_tensor_desc.SetRealDimCnt(shapes.size());
138+ size_t size = 1;
139+ for (uint32_t i = 0; i < shapes.size(); i++) {
140+ size *= shapes[i];
141+ }
142+ uint32_t data_len = size * GetDataTypeSize(data_type);
143+ int32_t* pData = new (std::nothrow) int32_t[data_len];
144+ for (size_t i = 0; i < size; ++i) {
145+ *(pData + i) = value;
146+ }
147+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t*>(pData), data_len);
148+ return SUCCESS;
149+}
150+ 
151+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData)
152+{
153+ FILE* fp = fopen(bin_file.c_str(), "w");
154+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
155+ fclose(fp);
156+ return SUCCESS;
157+}
158+ 
159+int CreateOppInGraph(
160+ DataType inDtype, std::vector<ge::Tensor>& input, std::vector<Operator>& inputs, std::vector<Operator>& outputs,
161+ Graph& graph)
162+{
163+ Status ret = SUCCESS;
164+ auto truncateDiv1 = op::TruncateDiv("truncateDiv1");
165+ std::vector<int64_t> xShape = {4, 2};
166+ ADD_INPUT(1, x1, inDtype, xShape, 10.0f);
167+ ADD_INPUT(2, x2, inDtype, xShape, 3.0f);
168+ 
169+ outputs.push_back(truncateDiv1);
170+ return SUCCESS;
171+}
172+ 
173+void SaveInputOutput(std::vector<ge::Tensor>& input, std::vector<ge::Tensor>& output)
174+{
175+ int input_num = input.size();
176+ for (int i = 0; i < input_num; i++) {
177+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
178+ string input_file = "./tc_ge_irrun_test_truncate_div_npu_input_" + std::to_string(i) + ".bin";
179+ uint8_t* input_data_i = input[i].GetData();
180+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
181+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
182+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
183+ WriteDataToFile((const char*)input_file.c_str(), data_size, input_data_i);
184+ }
185+ 
186+ int output_num = output.size();
187+ for (int i = 0; i < output_num; i++) {
188+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
189+ string output_file = "./tc_ge_irrun_test_truncate_div_npu_output_" + std::to_string(i) + ".bin";
190+ uint8_t* output_data_i = output[i].GetData();
191+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
192+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
193+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
194+ WriteDataToFile((const char*)output_file.c_str(), data_size, output_data_i);
195+ int32_t* resultData = (int32_t*)output_data_i;
196+ for (int64_t j = 0; j < output_shape; j++) {
197+ LOG_PRINT("result[%ld] is: %d\n", j, resultData[j]);
198+ }
199+ }
200+}
201+ 
202+int main(int argc, char* argv[])
203+{
204+ const char* graph_name = "tc_ge_irrun_test_truncate_div";
205+ Graph graph(graph_name);
206+ std::vector<ge::Tensor> input;
207+ 
208+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
209+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
210+ Status ret = ge::GEInitialize(global_options);
211+ if (ret != SUCCESS) {
212+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
213+ return FAILED;
214+ }
215+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
216+ 
217+ std::vector<Operator> inputs{};
218+ std::vector<Operator> outputs{};
219+ 
220+ std::cout << argv[1] << std::endl;
221+ 
222+ DataType inDtype = DT_FLOAT;
223+ 
224+ std::cout << inDtype << std::endl;
225+ 
226+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
227+ if (ret != SUCCESS) {
228+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
229+ return FAILED;
230+ }
231+ 
232+ if (!inputs.empty() && !outputs.empty()) {
233+ graph.SetInputs(inputs).SetOutputs(outputs);
234+ }
235+ 
236+ std::map<AscendString, AscendString> build_options = {
237+ 
238+ };
239+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
240+ ge::Session* session = new Session(build_options);
241+ 
242+ if (session == nullptr) {
243+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
244+ return FAILED;
245+ }
246+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
247+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
248+ 
249+ std::map<AscendString, AscendString> graph_options = {
250+ 
251+ };
252+ uint32_t graph_id = 0;
253+ ret = session->AddGraph(graph_id, graph, graph_options);
254+ 
255+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
256+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
257+ std::string file_path = "./dump_truncate_div";
258+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
259+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
260+ std::vector<ge::Tensor> output;
261+ ret = session->RunGraph(graph_id, input, output);
262+ if (ret != SUCCESS) {
263+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
264+ delete session;
265+ GEFinalize();
266+ return FAILED;
267+ }
268+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
269+ 
270+ SaveInputOutput(input, output);
271+ 
272+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
273+ std::string error_str(error_msg.GetString());
274+ std::cout << "Error message: " << error_str << std::endl;
275+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
276+ std::string warning_str(warning_msg.GetString());
277+ std::cout << "Warning message: " << warning_str << std::endl;
278+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
279+ ret = ge::GEFinalize();
280+ if (ret != SUCCESS) {
281+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
282+ return FAILED;
283+ }
284+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
285+ return SUCCESS;
286+}
Amath/truncate_div/op_api/truncate_div.cpp+103-0
@@ -0,0 +1,103 @@
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+#include "truncate_div.h"
12+#include "opdev/make_op_executor.h"
13+#include "opdev/op_dfx.h"
14+#include "opdev/op_executor.h"
15+#include "opdev/op_log.h"
16+#include "opdev/shape_utils.h"
17+#include "opdev/aicpu/aicpu_task.h"
18+#include "op_api/aclnn_check.h"
19+ 
20+#include <unordered_map>
21+#include <unordered_set>
22+ 
23+using namespace op;
24+ 
25+namespace l0op {
26+ 
27+OP_TYPE_REGISTER(TruncateDiv);
28+ 
29+static const std::unordered_map<op::DataType, std::unordered_set<op::DataType>> DTYPE_MAPPING = {
30+ {op::DataType::DT_FLOAT, {op::DataType::DT_FLOAT, op::DataType::DT_INT32, op::DataType::DT_FLOAT16}},
31+ {op::DataType::DT_FLOAT16, {op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT}},
32+ {op::DataType::DT_BF16, {op::DataType::DT_BF16}},
33+ {op::DataType::DT_INT8, {op::DataType::DT_INT8}},
34+ {op::DataType::DT_INT16, {op::DataType::DT_INT16}},
35+ {op::DataType::DT_INT32, {op::DataType::DT_INT32, op::DataType::DT_FLOAT}},
36+ {op::DataType::DT_INT64, {op::DataType::DT_INT64}},
37+ {op::DataType::DT_UINT8, {op::DataType::DT_UINT8}}};
38+ 
39+static inline bool IsAiCoreSupport(const aclTensor* self, const aclTensor* other)
40+{
41+ auto selfDtype = self->GetDataType();
42+ auto otherDtype = other->GetDataType();
43+ 
44+ auto it = DTYPE_MAPPING.find(selfDtype);
45+ if (it == DTYPE_MAPPING.end()) {
46+ return false;
47+ }
48+ 
49+ return it->second.find(otherDtype) != it->second.end();
50+}
51+ 
52+static const aclTensor* TruncateDivAiCore(
53+ const aclTensor* self, const aclTensor* other, aclTensor* out, aclOpExecutor* executor)
54+{
55+ L0_DFX(TruncateDivAiCore, self, other, out);
56+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(TruncateDiv, OP_INPUT(self, other), OP_OUTPUT(out));
57+ OP_CHECK(
58+ ret == ACL_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "TruncateDivAiCore ADD_TO_LAUNCHER_LIST_AICORE failed."),
59+ return nullptr);
60+ return out;
61+}
62+ 
63+static const aclTensor* TruncateDivAiCpu(
64+ const aclTensor* self, const aclTensor* other, aclTensor* out, aclOpExecutor* executor)
65+{
66+ L0_DFX(TruncateDivAiCpu, self, other, out);
67+ static internal::AicpuTaskSpace space("TruncateDiv", ge::DEPEND_IN_SHAPE, true);
68+ auto ret = ADD_TO_LAUNCHER_LIST_AICPU(TruncateDiv, OP_ATTR_NAMES(), OP_INPUT(self, other), OP_OUTPUT(out));
69+ OP_CHECK(
70+ ret == ACL_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "TruncateDivAiCpu ADD_TO_LAUNCHER_LIST_AICPU failed."),
71+ return nullptr);
72+ return out;
73+}
74+ 
75+const aclTensor* TruncateDiv(const aclTensor* self, const aclTensor* other, aclOpExecutor* executor)
76+{
77+ L0_DFX(TruncateDiv, self, other);
78+ op::Shape broadcastShape;
79+ if (!BroadcastInferShape(self->GetViewShape(), other->GetViewShape(), broadcastShape)) {
80+ OP_LOGE(
81+ ACLNN_ERR_PARAM_INVALID, "Broadcast %s and %s failed.", op::ToString(self->GetViewShape()).GetString(),
82+ op::ToString(other->GetViewShape()).GetString());
83+ return nullptr;
84+ }
85+ 
86+ op::DataType outDataType;
87+ 
88+ if (self->GetDataType() == other->GetDataType()) {
89+ outDataType = self->GetDataType();
90+ } else {
91+ outDataType = op::DataType::DT_FLOAT;
92+ }
93+ 
94+ auto divOut = executor->AllocTensor(broadcastShape, outDataType);
95+ 
96+ if (IsAiCoreSupport(self, other)) {
97+ return TruncateDivAiCore(self, other, divOut, executor);
98+ } else {
99+ return TruncateDivAiCpu(self, other, divOut, executor);
100+ }
101+}
102+ 
103+} // namespace l0op
Amath/truncate_div/op_api/truncate_div.h+20-0
@@ -0,0 +1,20 @@
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+#ifndef PTA_NPU_OP_API_INC_LEVEL0_OP_TRUNCATE_DIV_OP_H_
12+#define PTA_NPU_OP_API_INC_LEVEL0_OP_TRUNCATE_DIV_OP_H_
13+ 
14+#include "opdev/op_executor.h"
15+ 
16+namespace l0op {
17+const aclTensor* TruncateDiv(const aclTensor* self, const aclTensor* other, aclOpExecutor* executor);
18+} // namespace l0op
19+ 
20+#endif // PTA_NPU_OP_API_INC_LEVEL0_OP_TRUNCATE_DIV_OP_H_
Amath/truncate_div/op_graph/truncate_div_proto.h+56-0
@@ -0,0 +1,56 @@
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 truncate_div_proto.h
13+ * \brief truncate_div operator prototype
14+ */
15+#ifndef OPS_OP_PROTO_INC_TRUNCATE_DIV_H_
16+#define OPS_OP_PROTO_INC_TRUNCATE_DIV_H_
17+ 
18+#include "graph/operator_reg.h"
19+#include "graph/types.h"
20+ 
21+namespace ge {
22+ 
23+/**
24+ * @brief Returns trunc(x1/x2) element-wise. Support broadcasting operations.
25+ * truncation is towards zero.
26+ 
27+ * @par Inputs:
28+ * Two inputs, including:
29+ * @li x1: A ND Tensor. Must be one of the following types:
30+ * bfloat16, float16, float32, double, int8, uint8, uint16, int16, int32, int64, complex64, complex128. The format can be ND.
31+ * @li x2: A ND Tensor. Has the same dtype and format as input "x1". \n
32+ 
33+ * @par Outputs:
34+ * y: A ND Tensor. Has the same dtype and format as input "x1". Has the same shape as the broadcast shape of x1 and x2. \n
35+ 
36+ * @par Third-party framework compatibility
37+ * Compatible with the PyTorch operator TruncateDiv.
38+ */
39+REG_OP(TruncateDiv)
40+ .INPUT(
41+ x1, TensorType(
42+ {DT_FLOAT, DT_FLOAT16, DT_BF16, DT_INT8, DT_UINT8, DT_INT32, DT_DOUBLE, DT_UINT16, DT_INT16, DT_INT64,
43+ DT_COMPLEX64, DT_COMPLEX128}))
44+ .INPUT(
45+ x2, TensorType(
46+ {DT_FLOAT, DT_FLOAT16, DT_BF16, DT_INT8, DT_UINT8, DT_INT32, DT_DOUBLE, DT_UINT16, DT_INT16, DT_INT64,
47+ DT_COMPLEX64, DT_COMPLEX128}))
48+ .OUTPUT(
49+ y, TensorType(
50+ {DT_FLOAT, DT_FLOAT16, DT_BF16, DT_INT8, DT_UINT8, DT_INT32, DT_DOUBLE, DT_UINT16, DT_INT16, DT_INT64,
51+ DT_COMPLEX64, DT_COMPLEX128}))
52+ .OP_END_FACTORY_REG(TruncateDiv)
53+ 
54+} // namespace ge
55+ 
56+#endif // OPS_OP_PROTO_INC_TRUNCATE_DIV_H_
Amath/truncate_div/op_host/arch35/truncate_div_tiling_arch35.cpp+300-0
@@ -0,0 +1,300 @@
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 truncate_div_tiling_arch35.cpp
13+ * \brief truncate_div_tiling source file
14+ */
15+ 
16+#include <graph/utils/type_utils.h>
17+#include "register/op_impl_registry.h"
18+#include "op_host/math_tiling_templates_registry.h"
19+#include "atvoss/broadcast/broadcast_tiling.h"
20+#include "math/truncate_div/op_kernel/arch35/truncate_div_dag.h"
21+#include "math/truncate_div/op_kernel/arch35/truncate_div_struct.h"
22+#include "truncate_div_tiling_arch35.h"
23+#include "op_host/util/fp16.h"
24+#include "util/bfloat16.h"
25+ 
26+using namespace Ops::Base;
27+using namespace ge;
28+ 
29+namespace optiling {
30+ 
31+constexpr static uint64_t TRUNCATE_DIV_COMMON_TILING_PRIORITY = 0;
32+constexpr static uint32_t INPUT_IDX_X1 = 0;
33+constexpr static uint32_t INPUT_IDX_X2 = 1;
34+constexpr static int64_t DCACHE_SIZE = 32 * 1024;
35+ 
36+ge::graphStatus TruncateDivTiling::GetShapeAttrsInfo()
37+{
38+ return ge::GRAPH_SUCCESS;
39+}
40+ 
41+bool TruncateDivTiling::IsCapable()
42+{
43+ return true;
44+}
45+ 
46+float TruncateDivTiling::GetReciprocal(float data)
47+{
48+ float invScalarVal = 0.0f;
49+ float scalarValue = data;
50+ if (scalarValue == 0.0f) {
51+ float reciprocalZero = 1.0f / scalarValue;
52+ if (reciprocalZero < 0) {
53+ invScalarVal = -INFINITY;
54+ } else {
55+ invScalarVal = INFINITY;
56+ }
57+ } else if (scalarValue == INFINITY) {
58+ invScalarVal = 0.0f;
59+ } else if (scalarValue == -INFINITY) {
60+ invScalarVal = -0.0f;
61+ } else {
62+ invScalarVal = 1.0f / scalarValue;
63+ invScalarVal = invScalarVal * (2.0f - scalarValue * invScalarVal);
64+ }
65+ return invScalarVal;
66+}
67+ 
68+template <typename T>
69+ge::graphStatus TruncateDivTiling::GetConstData(uint32_t inputIdx, T& data)
70+{
71+ auto tensor = context_->GetInputTensor(inputIdx);
72+ OP_CHECK_NULL_WITH_CONTEXT(context_, tensor);
73+ const T* value = tensor->GetData<T>();
74+ if (value == nullptr) {
75+ OP_LOGE(context_->GetNodeName(), "const tensor is null.");
76+ return ge::GRAPH_FAILED;
77+ }
78+ data = value[0];
79+ OP_LOGI(context_->GetNodeName(), "scalarData %f", data);
80+ return ge::GRAPH_SUCCESS;
81+}
82+ 
83+ge::graphStatus TruncateDivTiling::DoOpTiling()
84+{
85+ // 1. 获取输入描述
86+ auto x1Desc = context_->GetInputDesc(INPUT_IDX_X1);
87+ OP_CHECK_NULL_WITH_CONTEXT(context_, x1Desc);
88+ auto x2Desc = context_->GetInputDesc(INPUT_IDX_X2);
89+ OP_CHECK_NULL_WITH_CONTEXT(context_, x2Desc);
90+ 
91+ ge::DataType x1DType = x1Desc->GetDataType();
92+ ge::DataType x2DType = x2Desc->GetDataType();
93+ 
94+ // 2. 获取形状
95+ auto x1StorageShape = context_->GetInputShape(INPUT_IDX_X1);
96+ OP_CHECK_NULL_WITH_CONTEXT(context_, x1StorageShape);
97+ auto x2StorageShape = context_->GetInputShape(INPUT_IDX_X2);
98+ OP_CHECK_NULL_WITH_CONTEXT(context_, x2StorageShape);
99+ 
100+ auto x2Shape = x2StorageShape->GetStorageShape();
101+ 
102+// bool isScalar = x2Shape.IsScalar() || (x2Shape.GetDimNum() == 1 && x2Shape.GetDim(0) == 1);
103+ bool isScalar = x2Shape.IsScalar();
104+ bool canUseMul = isScalar && (x2DType == ge::DT_FLOAT || x2DType == ge::DT_FLOAT16 || x2DType == ge::DT_BF16);
105+ 
106+ OP_LOGI(context_->GetNodeName(), "canUseMul %d", canUseMul);
107+ if (canUseMul) {
zhanw_coding
zhanw_codingzhanw_coding5月29日

函数过大，建议按是否走倒数模板拆分为两个子函数

likedislike
108+ bool success = false;
109+ float scalarValue = 0.0f;
110+ switch (x2DType) {
111+ case ge::DT_FLOAT: {
112+ success = (GetConstData<float>(INPUT_IDX_X2, scalarValue) == ge::GRAPH_SUCCESS);
113+ if (success) {
114+ reciprocal_ = GetReciprocal(scalarValue);
115+ }
116+ break;
117+ }
118+ case ge::DT_FLOAT16: {
119+ uint16_t tmpValue = 0;
120+ success = (GetConstData<uint16_t>(INPUT_IDX_X2, tmpValue) == ge::GRAPH_SUCCESS);
121+ if (success) {
122+ scalarValue = float(*(reinterpret_cast<const fp16_t*>(&tmpValue)));
123+ reciprocal_ = GetReciprocal(scalarValue);
124+ }
125+ break;
126+ }
127+ case ge::DT_BF16: {
128+ uint16_t tmpValue = 0;
129+ success = (GetConstData<uint16_t>(INPUT_IDX_X2, tmpValue) == ge::GRAPH_SUCCESS);
130+ if (success) {
131+ scalarValue = float(*(reinterpret_cast<const bfloat16*>(&tmpValue)));
132+ reciprocal_ = GetReciprocal(scalarValue);
zhanw_coding
zhanw_codingzhanw_coding5月29日

取倒数代码重复出现，可以提取到 144行 之后

likedislike
133+ }
134+ break;
135+ }
136+ default:
137+ OP_LOGE(
138+ context_->GetNodeName(), "Unsupported scalar type for reciprocal: %s",
139+ ge::TypeUtils::DataTypeToSerialString(x2DType).c_str());
140+ return ge::GRAPH_FAILED;
141+ }
142+ if (!success) {
143+ return ge::GRAPH_FAILED;
144+ }
145+ OP_LOGI(context_->GetNodeName(), "scalar value = %f, reciprocal value = %f", scalarValue, reciprocal_);
146+ }
147+ 
148+ 
149+ ge::graphStatus ret = ge::GRAPH_SUCCESS;
150+ uint32_t schMode = 0;
151+ int64_t maxLiveNodeCnt = 0;
152+ int64_t extraBuf = DCACHE_SIZE;
153+ 
154+ // 5. 定义模板 lambda：封装重复的 tiling 执行逻辑
155+ auto execTiling = [&, this]<typename OpDag>(bool isScalarBranch = false) {
zhanw_coding
zhanw_codingzhanw_coding5月29日

直接定义一个成员函数即可，没必要定义成 lambda

likedislike
156+ BroadcastBaseTiling<OpDag> brcTiling(context_);
157+ if (isScalarBranch) {
158+ brcTiling.SetScalar(reciprocal_);
159+ }
160+ 
161+ ret = brcTiling.DoTiling();
162+ schMode = brcTiling.GetSchMode();
163+ 
164+ tilingKey_ = GET_TPL_TILING_KEY(schMode, canUseMul);
165+ };
166+ 
167+ // 6. 根据数据类型组合调用 lambda
168+ if (x1DType == ge::DT_FLOAT16 && x2DType == ge::DT_FLOAT) {
169+ if (canUseMul) {
170+ execTiling.template operator()<TruncateDivOp::TruncateDivFloatWithCastScalar<half, float, float>::OpDag>(
171+ true);
172+ } else {
173+ execTiling.template operator()<TruncateDivOp::TruncateDivFloatWithCast<half, float, float>::OpDag>(false);
174+ }
175+ } else if (x1DType == ge::DT_FLOAT16 || x1DType == ge::DT_BF16) {
176+ if (x2DType == x1DType) {
177+ if (canUseMul) {
178+ execTiling.template operator()<TruncateDivOp::TruncateDivFloat16Scalar<half, float>::OpDag>(true);
179+ } else {
180+ execTiling.template operator()<TruncateDivOp::TruncateDivFloat16<half, float>::OpDag>(false);
181+ }
182+ }
zhanw_coding
zhanw_codingzhanw_coding5月29日

else 呢？如果不支持这种场景就要报错，如果不可能有else，就去除该判断，并添加注释说明

likedislike
183+ } else if (x1DType == ge::DT_FLOAT) {
184+ if (x2DType == ge::DT_FLOAT) {
185+ if (canUseMul) {
186+ execTiling.template operator()<TruncateDivOp::TruncateDivFloatScalar<float>::OpDag>(true);
187+ } else {
188+ execTiling.template operator()<TruncateDivOp::TruncateDivFloat<float>::OpDag>(false);
189+ }
190+ } else if (x2DType == ge::DT_INT32) {
191+ execTiling.template operator()<TruncateDivOp::TruncateDivFloatToLowBit<float, int32_t, float>::OpDag>(
192+ false);
193+ } else if (x2DType == ge::DT_FLOAT16) {
194+ if (canUseMul) {
195+ execTiling.template
196+ operator()<TruncateDivOp::TruncateDivFloatWithCastScalar<float, half, float>::OpDag>(true);
197+ } else {
198+ execTiling.template operator()<TruncateDivOp::TruncateDivFloatToLowBit<float, half, float>::OpDag>(false);
199+ }
200+ }
201+ } else if (x1DType == ge::DT_INT8 || x1DType == ge::DT_UINT8) {
202+ // 这些分支不支持 scalar 优化（canUseMul 应该为 false）
203+ if (x1DType == ge::DT_INT8) {
204+ execTiling.template operator()<TruncateDivOp::TruncateDivIntS8<int8_t, half>::OpDag>(false);
205+ } else {
206+ execTiling.template operator()<TruncateDivOp::TruncateDivIntU8<uint8_t, uint16_t>::OpDag>(false);
207+ }
208+ } else if (x1DType == ge::DT_INT16) {
209+ execTiling.template operator()<TruncateDivOp::TruncateDivInt<int16_t>::OpDag>(false);
210+ } else if (x1DType == ge::DT_INT32 && x2DType == ge::DT_INT32) {
211+ execTiling.template operator()<TruncateDivOp::TruncateDivInt<int32_t>::OpDag>(false);
212+ } else if (x1DType == ge::DT_INT64) {
213+ BroadcastBaseTiling<TruncateDivOp::TruncateDivInt64<int64_t>::OpDag> brcTiling(context_);
214+ ret = brcTiling.DoTiling(extraBuf, maxLiveNodeCnt);
215+ schMode = brcTiling.GetSchMode();
216+ tilingKey_ = GET_TPL_TILING_KEY(schMode, canUseMul);
217+ } else if (x1DType == ge::DT_INT32 && x2DType == ge::DT_FLOAT) {
218+ execTiling.template operator()<TruncateDivOp::TruncateDivIntToFloat<int32_t, float, float>::OpDag>(false);
219+ } else {
220+ OP_LOGE(
221+ context_->GetNodeName(), "Unsupported dtype combination: self=%s, other=%s",
222+ ge::TypeUtils::DataTypeToSerialString(x1DType).c_str(),
223+ ge::TypeUtils::DataTypeToSerialString(x2DType).c_str());
224+ return ge::GRAPH_FAILED;
225+ }
226+ 
227+ // 7. 完成后打印结果
228+ OP_LOGI(
229+ context_, "TruncateDiv tiling completed: tilingKey_=%lu, schMode=%u, canUseMul=%d, reciprocal %f", tilingKey_, schMode,
230+ static_cast<int>(canUseMul), reciprocal_);
231+ 
232+ return ret;
233+}
234+ 
235+ge::graphStatus TruncateDivTiling::DoLibApiTiling()
236+{
237+ return ge::GRAPH_SUCCESS;
238+}
239+ 
240+uint64_t TruncateDivTiling::GetTilingKey() const
241+{
242+ return tilingKey_;
243+}
244+ 
245+ge::graphStatus TruncateDivTiling::GetWorkspaceSize()
246+{
247+ return ge::GRAPH_SUCCESS;
248+}
249+ 
250+ge::graphStatus TruncateDivTiling::PostTiling()
251+{
252+ context_->SetLocalMemorySize(static_cast<uint32_t>(ubSize_ - DCACHE_SIZE));
253+ return ge::GRAPH_SUCCESS;
254+}
255+ 
256+ge::graphStatus TruncateDivTiling::GetPlatformInfo()
257+{
258+ auto platformInfo = context_->GetPlatformInfo();
259+ if (platformInfo == nullptr) {
260+ auto compileInfoPtr = reinterpret_cast<const BroadcastCompileInfo*>(context_->GetCompileInfo());
261+ OP_CHECK_IF(compileInfoPtr == nullptr, OP_LOGE(context_, "compile info is null"), return ge::GRAPH_FAILED);
262+ ubSize_ = compileInfoPtr->ubSize;
263+ OP_LOGD(context_->GetNodeName(), "Get ubSize form compileInfo is: %ld", ubSize_);
264+ } else {
265+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
266+ uint64_t ubSizePlatform;
267+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatform);
268+ ubSize_ = static_cast<int64_t>(ubSizePlatform);
269+ OP_LOGD(context_->GetNodeName(), "Get ubSize form ascendcPlatform is: %ld", ubSize_);
270+ }
271+ return ge::GRAPH_SUCCESS;
272+}
273+ 
274+ge::graphStatus TilingForTruncateDiv(gert::TilingContext* context)
zhanw_coding
zhanw_codingzhanw_coding5月29日

非公开方法，使用 static 声明

likedislike
275+{
276+ OP_LOGD("TruncateDivTiling", "Enter TilingForTruncateDiv");
277+ if (context == nullptr) {
278+ OP_LOGE("TruncateDivTiling", "Tiling context is nullptr");
279+ return ge::GRAPH_FAILED;
280+ }
zhanw_coding
zhanw_codingzhanw_coding5月29日

直接用 OP_CHECK_NULL_WITH_CONTEXT(context, context);

likedislike
281+ 
282+ auto compileInfo = reinterpret_cast<const BroadcastCompileInfo*>(context->GetCompileInfo());
283+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
284+ 
285+ OP_LOGD(context, "Enter ascendc TruncateDivTiling");
286+ return Ops::Math::OpTiling::TilingRegistry::GetInstance().DoTilingImpl(context);
287+}
288+ 
289+ge::graphStatus TilingPrepareForTruncateDiv([[maybe_unused]] gert::TilingParseContext* context)
zhanw_coding
zhanw_codingzhanw_coding5月29日

非公开方法，使用 static 声明

likedislike
290+{
291+ return ge::GRAPH_SUCCESS;
292+}
293+ 
294+IMPL_OP_OPTILING(TruncateDiv)
295+ .Tiling(TilingForTruncateDiv)
296+ .TilingInputsDataDependency({INPUT_IDX_X2})
297+ .TilingParse<BroadcastCompileInfo>(TilingPrepareForTruncateDiv);
298+ 
299+REGISTER_OPS_TILING_TEMPLATE(TruncateDiv, TruncateDivTiling, TRUNCATE_DIV_COMMON_TILING_PRIORITY);
300+} // namespace optiling
Amath/truncate_div/op_host/arch35/truncate_div_tiling_arch35.h+56-0
@@ -0,0 +1,56 @@
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 truncate_div_tiling_arch35.h
13+ * \brief truncate_div_tiling head file
14+ */
15+ 
16+#ifndef OPS_BUILD_IN_OP_TILING_RUNTIME_TRUNCATE_DIV_TILING_H
17+#define OPS_BUILD_IN_OP_TILING_RUNTIME_TRUNCATE_DIV_TILING_H
18+ 
19+#include "log/log.h"
20+#include "platform/platform_info.h"
21+#include "register/op_def_registry.h"
22+#include "register/op_impl_registry.h"
23+#include "tiling/tiling_api.h"
24+#include "op_host/tiling_base_class.h"
25+ 
26+namespace optiling {
27+ 
28+class TruncateDivTiling : public Ops::Base::TilingBaseClass {
29+public:
30+ explicit TruncateDivTiling(gert::TilingContext* context) : Ops::Base::TilingBaseClass(context)
31+ {}
32+ 
33+protected:
34+ bool IsCapable() override;
35+ ge::graphStatus DoLibApiTiling() override;
36+ ge::graphStatus DoOpTiling() override;
37+ ge::graphStatus PostTiling() override;
38+ uint64_t GetTilingKey() const override;
39+ ge::graphStatus GetShapeAttrsInfo() override;
40+ ge::graphStatus GetPlatformInfo() override;
41+ ge::graphStatus GetWorkspaceSize() override;
42+ 
43+private:
44+ uint64_t tilingKey_ = 0;
45+ float reciprocal_ = 0.0f;
46+ int64_t ubSize_ = 0;
47+ 
48+ template <typename T>
49+ ge::graphStatus GetConstData(uint32_t inputIdx, T& data);
50+ 
51+ float GetReciprocal(float data);
52+};
53+ 
54+} // namespace optiling
55+ 
56+#endif // OPS_BUILD_IN_OP_TILING_RUNTIME_TRUNCATE_DIV_TILING_H
Amath/truncate_div/op_host/config/ascend950/truncate_div_binary.json+377-0
@@ -0,0 +1,377 @@
1+{
2+ "op_type": "TruncateDiv",
3+ "op_list": [
4+ {
5+ "bin_filename": "TruncateDiv_BF16",
6+ "inputs": [
7+ {
8+ "name": "x1",
9+ "index": 0,
10+ "dtype": "bfloat16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [-2]
zhanw_coding
zhanw_codingzhanw_coding5月29日

应该加上"format_match_mode": "FormatAgnostic"，二进制匹配与格式无关

likedislike
14+ },
15+ {
16+ "name": "x2",
17+ "index": 1,
18+ "dtype": "bfloat16",
19+ "format": "ND",
20+ "paramType": "required",
21+ "shape": [-2]
22+ }
23+ ],
24+ "outputs": [
25+ {
26+ "name": "y",
27+ "index": 0,
28+ "dtype": "bfloat16",
29+ "format": "ND",
30+ "paramType": "required",
31+ "shape": [-2]
32+ }
33+ ]
34+ },
35+ {
36+ "bin_filename": "TruncateDiv_FLOAT16",
37+ "inputs": [
38+ {
39+ "name": "x1",
40+ "index": 0,
41+ "dtype": "float16",
42+ "format": "ND",
43+ "paramType": "required",
44+ "shape": [-2]
45+ },
46+ {
47+ "name": "x2",
48+ "index": 1,
49+ "dtype": "float16",
50+ "format": "ND",
51+ "paramType": "required",
52+ "shape": [-2]
53+ }
54+ ],
55+ "outputs": [
56+ {
57+ "name": "y",
58+ "index": 0,
59+ "dtype": "float16",
60+ "format": "ND",
61+ "paramType": "required",
62+ "shape": [-2]
63+ }
64+ ]
65+ },
66+ {
67+ "bin_filename": "TruncateDiv_FLOAT16_FLOAT",
68+ "inputs": [
69+ {
70+ "name": "x1",
71+ "index": 0,
72+ "dtype": "float16",
73+ "format": "ND",
74+ "paramType": "required",
75+ "shape": [-2]
76+ },
77+ {
78+ "name": "x2",
79+ "index": 1,
80+ "dtype": "float32",
81+ "format": "ND",
82+ "paramType": "required",
83+ "shape": [-2]
84+ }
85+ ],
86+ "outputs": [
87+ {
88+ "name": "y",
89+ "index": 0,
90+ "dtype": "float32",
91+ "format": "ND",
92+ "paramType": "required",
93+ "shape": [-2]
94+ }
95+ ]
96+ },
97+ {
98+ "bin_filename": "TruncateDiv_FLOAT_FLOAT16",
99+ "inputs": [
100+ {
101+ "name": "x1",
102+ "index": 0,
103+ "dtype": "float32",
104+ "format": "ND",
105+ "paramType": "required",
106+ "shape": [-2]
107+ },
108+ {
109+ "name": "x2",
110+ "index": 1,
111+ "dtype": "float16",
112+ "format": "ND",
113+ "paramType": "required",
114+ "shape": [-2]
115+ }
116+ ],
117+ "outputs": [
118+ {
119+ "name": "y",
120+ "index": 0,
121+ "dtype": "float32",
122+ "format": "ND",
123+ "paramType": "required",
124+ "shape": [-2]
125+ }
126+ ]
127+ },
128+ {
129+ "bin_filename": "TruncateDiv_FLOAT",
130+ "inputs": [
131+ {
132+ "name": "x1",
133+ "index": 0,
134+ "dtype": "float32",
135+ "format": "ND",
136+ "paramType": "required",
137+ "shape": [-2]
138+ },
139+ {
140+ "name": "x2",
141+ "index": 1,
142+ "dtype": "float32",
143+ "format": "ND",
144+ "paramType": "required",
145+ "shape": [-2]
146+ }
147+ ],
148+ "outputs": [
149+ {
150+ "name": "y",
151+ "index": 0,
152+ "dtype": "float32",
153+ "format": "ND",
154+ "paramType": "required",
155+ "shape": [-2]
156+ }
157+ ]
158+ },
159+ {
160+ "bin_filename": "TruncateDiv_FLOAT_INT32",
161+ "inputs": [
162+ {
163+ "name": "x1",
164+ "index": 0,
165+ "dtype": "float32",
166+ "format": "ND",
167+ "paramType": "required",
168+ "shape": [-2]
169+ },
170+ {
171+ "name": "x2",
172+ "index": 1,
173+ "dtype": "int32",
174+ "format": "ND",
175+ "paramType": "required",
176+ "shape": [-2]
177+ }
178+ ],
179+ "outputs": [
180+ {
181+ "name": "y",
182+ "index": 0,
183+ "dtype": "float32",
184+ "format": "ND",
185+ "paramType": "required",
186+ "shape": [-2]
187+ }
188+ ]
189+ },
190+ {
191+ "bin_filename": "TruncateDiv_INT32",
192+ "inputs": [
193+ {
194+ "name": "x1",
195+ "index": 0,
196+ "dtype": "int32",
197+ "format": "ND",
198+ "paramType": "required",
199+ "shape": [-2]
200+ },
201+ {
202+ "name": "x2",
203+ "index": 1,
204+ "dtype": "int32",
205+ "format": "ND",
206+ "paramType": "required",
207+ "shape": [-2]
208+ }
209+ ],
210+ "outputs": [
211+ {
212+ "name": "y",
213+ "index": 0,
214+ "dtype": "int32",
215+ "format": "ND",
216+ "paramType": "required",
217+ "shape": [-2]
218+ }
219+ ]
220+ },
221+ {
222+ "bin_filename": "TruncateDiv_INT32_FLOAT",
223+ "inputs": [
224+ {
225+ "name": "x1",
226+ "index": 0,
227+ "dtype": "int32",
228+ "format": "ND",
229+ "paramType": "required",
230+ "shape": [-2]
231+ },
232+ {
233+ "name": "x2",
234+ "index": 1,
235+ "dtype": "float32",
236+ "format": "ND",
237+ "paramType": "required",
238+ "shape": [-2]
239+ }
240+ ],
241+ "outputs": [
242+ {
243+ "name": "y",
244+ "index": 0,
245+ "dtype": "float32",
246+ "format": "ND",
247+ "paramType": "required",
248+ "shape": [-2]
249+ }
250+ ]
251+ },
252+ {
253+ "bin_filename": "TruncateDiv_UINT8",
254+ "inputs": [
255+ {
256+ "name": "x1",
257+ "index": 0,
258+ "dtype": "uint8",
259+ "format": "ND",
260+ "paramType": "required",
261+ "shape": [-2]
262+ },
263+ {
264+ "name": "x2",
265+ "index": 1,
266+ "dtype": "uint8",
267+ "format": "ND",
268+ "paramType": "required",
269+ "shape": [-2]
270+ }
271+ ],
272+ "outputs": [
273+ {
274+ "name": "y",
275+ "index": 0,
276+ "dtype": "uint8",
277+ "format": "ND",
278+ "paramType": "required",
279+ "shape": [-2]
280+ }
281+ ]
282+ },
283+ {
284+ "bin_filename": "TruncateDiv_INT8",
285+ "inputs": [
286+ {
287+ "name": "x1",
288+ "index": 0,
289+ "dtype": "int8",
290+ "format": "ND",
291+ "paramType": "required",
292+ "shape": [-2]
293+ },
294+ {
295+ "name": "x2",
296+ "index": 1,
297+ "dtype": "int8",
298+ "format": "ND",
299+ "paramType": "required",
300+ "shape": [-2]
301+ }
302+ ],
303+ "outputs": [
304+ {
305+ "name": "y",
306+ "index": 0,
307+ "dtype": "int8",
308+ "format": "ND",
309+ "paramType": "required",
310+ "shape": [-2]
311+ }
312+ ]
313+ },
314+ {
315+ "bin_filename": "TruncateDiv_INT64",
316+ "inputs": [
317+ {
318+ "name": "x1",
319+ "index": 0,
320+ "dtype": "int64",
321+ "format": "ND",
322+ "paramType": "required",
323+ "shape": [-2]
324+ },
325+ {
326+ "name": "x2",
327+ "index": 1,
328+ "dtype": "int64",
329+ "format": "ND",
330+ "paramType": "required",
331+ "shape": [-2]
332+ }
333+ ],
334+ "outputs": [
335+ {
336+ "name": "y",
337+ "index": 0,
338+ "dtype": "int64",
339+ "format": "ND",
340+ "paramType": "required",
341+ "shape": [-2]
342+ }
343+ ]
344+ },
345+ {
346+ "bin_filename": "TruncateDiv_INT16",
347+ "inputs": [
348+ {
349+ "name": "x1",
350+ "index": 0,
351+ "dtype": "int16",
352+ "format": "ND",
353+ "paramType": "required",
354+ "shape": [-2]
355+ },
356+ {
357+ "name": "x2",
358+ "index": 1,
359+ "dtype": "int16",
360+ "format": "ND",
361+ "paramType": "required",
362+ "shape": [-2]
363+ }
364+ ],
365+ "outputs": [
366+ {
367+ "name": "y",
368+ "index": 0,
369+ "dtype": "int16",
370+ "format": "ND",
371+ "paramType": "required",
372+ "shape": [-2]
373+ }
374+ ]
375+ }
376+ ]
377+}
Amath/truncate_div/op_host/config/ascend950/truncate_div_simplified_key.ini+3-0
@@ -0,0 +1,3 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值
2+[TruncateDiv]
3+default=0
Amath/truncate_div/op_host/truncate_div_def.cpp+68-0
@@ -0,0 +1,68 @@
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 truncate_div_def.cpp
13+ * \brief truncate_div operator definition
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class TruncateDiv : public OpDef {
19+public:
20+ explicit TruncateDiv(const char* name) : OpDef(name)
21+ {
22+ this->Input("x1")
23+ .ParamType(REQUIRED)
24+ .DataType(
25+ {ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_INT32,
26+ ge::DT_INT32, ge::DT_UINT8, ge::DT_INT8, ge::DT_INT64, ge::DT_INT16})
27+ .Format(
28+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
29+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
30+ .UnknownShapeFormat(
31+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
32+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
33+ this->Input("x2")
34+ .ParamType(REQUIRED)
35+ .DataType(
36+ {ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT32,
37+ ge::DT_FLOAT, ge::DT_UINT8, ge::DT_INT8, ge::DT_INT64, ge::DT_INT16})
38+ .Format(
39+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
40+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
41+ .UnknownShapeFormat(
42+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
43+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}).ValueDepend(OPTIONAL);
44+ this->Output("y")
45+ .ParamType(REQUIRED)
46+ .DataType(
47+ {ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_INT32,
48+ ge::DT_FLOAT, ge::DT_UINT8, ge::DT_INT8, ge::DT_INT64, ge::DT_INT16})
49+ .Format(
50+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
51+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
52+ .UnknownShapeFormat(
53+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
54+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
55+ OpAICoreConfig aicoreConfig;
56+ aicoreConfig.DynamicCompileStaticFlag(true)
57+ .DynamicFormatFlag(false)
58+ .DynamicRankSupportFlag(true)
59+ .DynamicShapeSupportFlag(true)
60+ .NeedCheckSupportFlag(false)
61+ .PrecisionReduceFlag(true)
62+ .ExtendCfgInfo("opFile.value", "truncate_div_apt");
63+ this->AICore().AddConfig("ascend950", aicoreConfig);
64+ }
65+};
66+ 
67+OP_ADD(TruncateDiv);
68+} // namespace ops
Amath/truncate_div/op_host/truncate_div_infershape.cpp+23-0
@@ -0,0 +1,23 @@
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 truncate_div_infershape.cpp
13+ * \brief truncate_div infershape
14+ */
15+#include "op_host/infershape_broadcast_util.h"
16+#include "register/op_impl_registry.h"
17+ 
18+using namespace ge;
19+namespace ops {
20+ 
21+IMPL_OP_INFERSHAPE(TruncateDiv).InferShape(Ops::Base::InferShape4Broadcast);
22+ 
23+} // namespace ops
Amath/truncate_div/op_kernel/arch35/truncate_div_dag.h+377-0
@@ -0,0 +1,377 @@
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 truncate_div_dag.h
13+ * \brief truncate_div dag
14+ */
15+ 
16+#ifndef TRUNCATE_DIV_DAG_H
17+#define TRUNCATE_DIV_DAG_H
18+ 
19+#include "atvoss/util/dag.h"
20+#include "atvoss/util/vec.h"
21+#include "atvoss/util/placeholder.h"
22+#include "kernel_tiling/kernel_tiling.h"
23+#include "op_kernel/math_util.h"
24+#ifdef __CCE_AICORE__
25+#include "op_kernel/platform_util.h"
26+#include "simt_api/asc_simt.h"
27+#endif
28+ 
29+namespace TruncateDivOp {
30+using namespace Ops::Base;
31+constexpr int TRUNCATE_DIV_CAST_MODE_NONE = 0;
32+constexpr 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;
36+constexpr int8_t SAT_POS = 60;
37+constexpr int64_t INT64_MAX_VALUE = 9223372036854775807;
38+constexpr int64_t INT32_MAX_VALUE = 2147483647;
39+const uint32_t UINT32_SIGN = 0x80000000;
40+const uint16_t UINT16_SIGN = 0x8000;
41+const int16_t TRUNCATE_DIV_B16_SIGN = -32768;
zhanw_coding
zhanw_codingzhanw_coding5月29日

无用且重复的常量，删除

likedislike
42+ 
43+namespace TruncDag1 {
44+template <class T>
45+struct TruncCustom : public Vec::ElemwiseUnaryOP<T, T> {
46+ __aicore__ inline TruncCustom(LocalTensor<T>& dst, LocalTensor<T>& src, uint32_t count)
47+ {
48+#ifdef __CCE_AICORE__
49+ uint32_t dtypeSize = sizeof(T);
50+ uint32_t vl = VECTOR_REG_WIDTH / dtypeSize;
51+ uint16_t loopNum = CeilDivision(count, vl);
52+ uint32_t vlSize = vl;
53+ __ubuf__ T* srcAddr = (__ubuf__ T*)src.GetPhyAddr();
54+ __ubuf__ T* dstAddr = (__ubuf__ T*)dst.GetPhyAddr();
55+ 
56+ MicroAPI::RegTensor<T, MicroAPI::RegTraitNumOne> vregInput;
57+ MicroAPI::RegTensor<T, MicroAPI::RegTraitNumOne> vregOutput;
58+ MicroAPI::MaskReg mask;
59+ if constexpr (std::is_same_v<T, float>) {
60+ MicroAPI::RegTensor<uint32_t, MicroAPI::RegTraitNumOne> vregOutInt;
61+ __VEC_SCOPE__
62+ {
63+ for (uint16_t loopIdx = 0; loopIdx < loopNum; loopIdx++) {
64+ mask = MicroAPI::UpdateMask<T, MicroAPI::RegTraitNumOne>(count);
65+ // OpCopyIn
66+ MicroAPI::DataCopy(vregInput, (__ubuf__ T*)(srcAddr + loopIdx * vlSize));
67+ 
68+ MicroAPI::Truncate<T, RoundMode::CAST_TRUNC, MicroAPI::MaskMergeMode::ZEROING>(
69+ vregOutput, vregInput, mask);
70+ MicroAPI::Duplicate(vregOutInt, UINT32_SIGN, mask);
71+ MicroAPI::And(vregOutInt, vregOutInt, (MicroAPI::RegTensor<uint32_t>&)vregInput, mask);
72+ MicroAPI::Or(vregOutInt, vregOutInt, (MicroAPI::RegTensor<uint32_t>&)vregOutput, mask);
73+ 
74+ // OpCopyOut
75+ MicroAPI::DataCopy(
76+ (__ubuf__ T*)(dstAddr + loopIdx * vlSize), (MicroAPI::RegTensor<T>&)vregOutInt, mask);
77+ }
78+ }
79+ } else {
80+ MicroAPI::RegTensor<uint16_t, MicroAPI::RegTraitNumOne> vregOutInt;
81+ __VEC_SCOPE__
82+ {
83+ for (uint16_t loopIdx = 0; loopIdx < loopNum; loopIdx++) {
84+ mask = MicroAPI::UpdateMask<T, MicroAPI::RegTraitNumOne>(count);
85+ // OpCopyIn
86+ MicroAPI::DataCopy(vregInput, (__ubuf__ T*)(srcAddr + loopIdx * vlSize));
87+ 
88+ MicroAPI::Truncate<T, RoundMode::CAST_TRUNC, MicroAPI::MaskMergeMode::ZEROING>(
89+ vregOutput, vregInput, mask);
90+ MicroAPI::Duplicate(vregOutInt, UINT16_SIGN, mask);
91+ MicroAPI::And(vregOutInt, vregOutInt, (MicroAPI::RegTensor<uint16_t>&)vregInput, mask);
92+ MicroAPI::Or(vregOutInt, vregOutInt, (MicroAPI::RegTensor<uint16_t>&)vregOutput, mask);
93+ 
94+ // OpCopyOut
95+ MicroAPI::DataCopy(
96+ (__ubuf__ T*)(dstAddr + loopIdx * vlSize), (MicroAPI::RegTensor<T>&)vregOutInt, mask);
97+ }
98+ }
99+ }
100+#endif
101+ }
102+};
103+} // namespace TruncDag1
104+ 
105+template <class R, class T, int roundMode>
106+struct CastOverFlow : public Vec::ElemwiseUnaryOP<R, T> {
107+ __aicore__ inline CastOverFlow(LocalTensor<R>& dst, LocalTensor<T>& src, const uint32_t& count)
108+ {
109+#ifdef __CCE_AICORE__
110+ SetCtrlSpr<SAT_POS, SAT_POS>(0);
111+ constexpr static MicroAPI::CastTrait castTrait3 = {
112+ MicroAPI::RegLayout::ZERO, MicroAPI::SatMode::NO_SAT, MicroAPI::MaskMergeMode::ZEROING,
113+ RoundMode::CAST_TRUNC};
114+ __VEC_SCOPE__
115+ {
116+ MicroAPI::RegTensor<T> vreg0;
117+ MicroAPI::RegTensor<R> vreg1;
118+ MicroAPI::MaskReg preg0;
119+ uint32_t size = count;
120+ uint16_t vfLoopNum = (size + (VECTOR_REG_WIDTH / sizeof(T)) - 1) / (VECTOR_REG_WIDTH / sizeof(T));
121+ __local_mem__ T* bufferIn0Addr = (__local_mem__ T*)src.GetPhyAddr();
122+ __local_mem__ R* bufferOut0Addr = (__local_mem__ R*)dst.GetPhyAddr();
123+ for (uint16_t i = 0; i < vfLoopNum; i++) {
124+ preg0 = MicroAPI::UpdateMask<T>(size);
125+ MicroAPI::DataCopy<T, MicroAPI::LoadDist::DIST_NORM>(
126+ vreg0, bufferIn0Addr + i * (VECTOR_REG_WIDTH / sizeof(T)));
127+ MicroAPI::Cast<R, T, castTrait3>(vreg1, vreg0, preg0);
128+ MicroAPI::DataCopy<R, MicroAPI::StoreDist::DIST_PACK_B16>(
129+ bufferOut0Addr + i * (VECTOR_REG_WIDTH / sizeof(T)), vreg1, preg0);
130+ }
131+ }
132+ SetCtrlSpr<SAT_POS, SAT_POS>(1);
133+#endif
134+ }
135+};
136+ 
137+template <class T>
138+struct TruncIntPostCompute : public Vec::ElemwiseTernaryOP<T, T, T, T> {
139+ __aicore__ inline TruncIntPostCompute(
140+ const LocalTensor<T>& dst, const LocalTensor<T>& input1, const LocalTensor<T>& input2,
141+ const LocalTensor<T>& div, const uint32_t& count)
142+ {
143+#ifdef __CCE_AICORE__
144+ constexpr uint32_t VECTOR_LENGTH = GetVRegSize();
145+ constexpr uint32_t VL_T = VECTOR_LENGTH / sizeof(T);
146+ __local_mem__ T* input1Addr = (__local_mem__ T*)input1.GetPhyAddr();
147+ __local_mem__ T* input2Addr = (__local_mem__ T*)input2.GetPhyAddr();
148+ __local_mem__ T* divAddr = (__local_mem__ T*)div.GetPhyAddr();
149+ __local_mem__ T* dstAddr = (__local_mem__ T*)dst.GetPhyAddr();
150+ uint16_t loopTimes = CeilDiv(count, VL_T);
151+ 
152+ __VEC_SCOPE__
153+ {
154+ MicroAPI::RegTensor<T> zeroValue;
155+ MicroAPI::RegTensor<T> defaultValue;
156+ MicroAPI::RegTensor<T> input1Value;
157+ MicroAPI::RegTensor<T> input2Value;
158+ MicroAPI::RegTensor<T> divValue;
159+ MicroAPI::RegTensor<T> resValue;
160+ MicroAPI::MaskReg preg;
161+ MicroAPI::MaskReg cmpValue;
162+ uint32_t sregMask = count;
163+ 
164+ MicroAPI::Duplicate(zeroValue, T(0));
165+ MicroAPI::Duplicate(defaultValue, T(-1));
166+ 
167+ for (uint16_t j = 0; j < loopTimes; j++) {
168+ preg = MicroAPI::UpdateMask<T>(sregMask);
169+ MicroAPI::DataCopy<T, MicroAPI::LoadDist::DIST_NORM>(input2Value, input2Addr + VL_T * j);
170+ MicroAPI::DataCopy<T, MicroAPI::LoadDist::DIST_NORM>(divValue, divAddr + VL_T * j);
171+ MicroAPI::DataCopy<T, MicroAPI::LoadDist::DIST_NORM>(input1Value, input1Addr + VL_T * j);
172+ MicroAPI::Compare<T, CMPMODE::NE>(cmpValue, input2Value, zeroValue, preg);
173+ MicroAPI::Select(resValue, divValue, defaultValue, cmpValue);
174+ MicroAPI::DataCopy<T, MicroAPI::StoreDist::DIST_NORM>(dstAddr + VL_T * j, resValue, preg);
175+ }
176+ }
177+#endif
178+ }
179+};
180+ 
181+#ifdef __CCE_AICORE__
182+template <typename T>
183+__simt_vf__ __aicore__
184+ LAUNCH_BOUND(1024) inline void TruncDivInt_1(__ubuf__ T* dst, __ubuf__ T* src1, __ubuf__ T* src2, int count)
zhanw_coding
zhanw_codingzhanw_coding5月29日

命名太随意了，建议改为 TruncDivIntSimt

likedislike
185+{
186+ for (uint32_t index = static_cast<uint32_t>(threadIdx.x); index < count;
187+ index += static_cast<uint32_t>(blockDim.x)) {
188+ bool pos_div_zero = ((src1[index] >= 0) && (src1[index] < INT64_MAX_VALUE) && (src2[index] == 0));
189+ bool div_zero = (src2[index] == 0);
190+ if (pos_div_zero) {
191+ dst[index] = INT32_MAX_VALUE;
192+ } else if (div_zero) {
193+ dst[index] = -1;
194+ } else {
195+ dst[index] = src1[index] / src2[index];
196+ }
197+ }
198+}
199+#endif
200+ 
201+template <class T>
202+struct TruncDivInt64 : public Vec::ElemwiseBinaryOP<T, T, T> {
203+ __aicore__ inline TruncDivInt64(LocalTensor<T>& dst, LocalTensor<T>& src1, LocalTensor<T>& src2, int count)
204+ {
205+#ifdef __CCE_AICORE__
206+ __ubuf__ T* dst_1 = (__ubuf__ T*)dst.GetPhyAddr();
207+ __ubuf__ T* src1_1 = (__ubuf__ T*)src1.GetPhyAddr();
208+ __ubuf__ T* src2_1 = (__ubuf__ T*)src2.GetPhyAddr();
209+ asc_vf_call<TruncDivInt_1<T>>(dim3(1024), dst_1, src1_1, src2_1, count);
210+#endif
211+ }
212+};
213+ 
214+template <typename T1, typename T2, typename PromoteT>
215+struct TruncateDivFloatWithCast {
216+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
217+ using InputX2 = Bind<Vec::CopyInBrc<T2>, Placeholder::In1<T2>>;
218+ using CastX1 = Bind<Vec::Cast<PromoteT, T1, TRUNCATE_DIV_CAST_MODE_NONE>, InputX1>;
219+ using CastX2 = Bind<Vec::Cast<PromoteT, T2, TRUNCATE_DIV_CAST_MODE_NONE>, InputX2>;
220+ using DivValue = Bind<Vec::DivHighPrecision<PromoteT>, CastX1, CastX2>;
221+ using TruncateValue = Bind<TruncDag1::TruncCustom<PromoteT>, DivValue>;
222+ using CastOut = Bind<Vec::Cast<PromoteT, PromoteT, TRUNCATE_DIV_CAST_MODE_RINT>, TruncateValue>;
223+ using OpCopyOut = Bind<Vec::CopyOut<PromoteT>, Placeholder::Out0<PromoteT>, CastOut>;
224+ using Outputs = Elems<OpCopyOut>;
225+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
226+ using OpDag = DAGSch<Outputs, void, MemCfg>;
227+};
228+ 
229+// half and bfloat16
230+template <typename T1, typename PromoteT>
231+struct TruncateDivFloat16 {
232+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
233+ using InputX2 = Bind<Vec::CopyInBrc<T1>, Placeholder::In1<T1>>;
234+ using CastX1 = Bind<Vec::Cast<PromoteT, T1, TRUNCATE_DIV_CAST_MODE_NONE>, InputX1>;
235+ using CastX2 = Bind<Vec::Cast<PromoteT, T1, TRUNCATE_DIV_CAST_MODE_NONE>, InputX2>;
236+ using DivValue = Bind<Vec::DivHighPrecision<PromoteT>, CastX1, CastX2>;
237+ using TruncateValue = Bind<TruncDag1::TruncCustom<PromoteT>, DivValue>;
238+ using CastOut = Bind<Vec::Cast<T1, PromoteT, TRUNCATE_DIV_CAST_MODE_RINT>, TruncateValue>;
239+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, CastOut>;
240+ using Outputs = Elems<OpCopyOut>;
241+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
242+ using OpDag = DAGSch<Outputs, void, MemCfg>;
243+};
244+ 
245+template <typename T1>
246+struct TruncateDivFloat {
247+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
248+ using InputX2 = Bind<Vec::CopyInBrc<T1>, Placeholder::In1<T1>>;
249+ using DivValue = Bind<Vec::DivHighPrecision<T1>, InputX1, InputX2>;
250+ using TruncateValue = Bind<TruncDag1::TruncCustom<T1>, DivValue>;
251+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, TruncateValue>;
252+ using Outputs = Elems<OpCopyOut>;
253+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
254+ using OpDag = DAGSch<Outputs, void, MemCfg>;
255+};
256+ 
257+template <typename T1, typename T2, typename PromoteT>
258+struct TruncateDivFloatWithCastScalar {
259+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
260+ using CastX1 = Bind<Vec::Cast<PromoteT, T1, TRUNCATE_DIV_CAST_MODE_NONE>, InputX1>;
261+ using DivValue = Bind<Vec::Muls<PromoteT>, CastX1, Placeholder::Var<float, 0>>;
262+ using TruncateValue = Bind<TruncDag1::TruncCustom<PromoteT>, DivValue>;
263+ using CastOut = Bind<Vec::Cast<PromoteT, PromoteT, TRUNCATE_DIV_CAST_MODE_RINT>, TruncateValue>;
264+ using OpCopyOut = Bind<Vec::CopyOut<PromoteT>, Placeholder::Out0<PromoteT>, CastOut>;
265+ using Outputs = Elems<OpCopyOut>;
266+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
267+ using OpDag = DAGSch<Outputs, void, MemCfg>;
268+};
269+ 
270+template <typename T1, typename PromoteT>
271+struct TruncateDivFloat16Scalar {
272+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
273+ using CastX1 = Bind<Vec::Cast<PromoteT, T1, TRUNCATE_DIV_CAST_MODE_NONE>, InputX1>;
274+ using DivValue = Bind<Vec::Muls<PromoteT>, CastX1, Placeholder::Var<float, 0>>;
275+ using TruncateValue = Bind<TruncDag1::TruncCustom<PromoteT>, DivValue>;
276+ using CastOut = Bind<Vec::Cast<T1, PromoteT, TRUNCATE_DIV_CAST_MODE_RINT>, TruncateValue>;
277+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, CastOut>;
278+ using Outputs = Elems<OpCopyOut>;
279+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
280+ using OpDag = DAGSch<Outputs, void, MemCfg>;
281+};
282+ 
283+template <typename T1>
284+struct TruncateDivFloatScalar {
285+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
286+ using DivValue = Bind<Vec::Muls<T1>, InputX1, Placeholder::Var<float, 0>>;
287+ using TruncateValue = Bind<TruncDag1::TruncCustom<T1>, DivValue>;
288+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, TruncateValue>;
289+ using Outputs = Elems<OpCopyOut>;
290+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
291+ using OpDag = DAGSch<Outputs, void, MemCfg>;
292+};
293+ 
294+template <typename T1, typename T2>
295+struct TruncateDivIntS8 {
296+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
297+ using InputX2 = Bind<Vec::CopyInBrc<T1>, Placeholder::In1<T1>>;
298+ 
299+ using CastX1 = Bind<Vec::Cast<T2, T1, TRUNCATE_DIV_CAST_MODE_NONE>, InputX1>;
300+ using CastX2 = Bind<Vec::Cast<T2, T1, TRUNCATE_DIV_CAST_MODE_NONE>, InputX2>;
301+ using DivValue = Bind<Vec::Div<T2>, CastX1, CastX2>;
302+ using OpIntResult = Bind<CastOverFlow<T1, T2, TRUNCATE_DIV_CAST_MODE_RINT>, DivValue>;
303+ using ComputeValue = Bind<TruncIntPostCompute<T1>, InputX1, InputX2, OpIntResult>;
304+ 
305+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, ComputeValue>;
306+ using Outputs = Elems<OpCopyOut>;
307+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
308+ using OpDag = DAGSch<Outputs, void, MemCfg>;
309+};
310+ 
311+template <typename T1, typename T2>
312+struct TruncateDivIntU8 {
313+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
314+ using InputX2 = Bind<Vec::CopyInBrc<T1>, Placeholder::In1<T1>>;
315+ using CastX1 = Bind<Vec::Cast<T2, T1, TRUNCATE_DIV_CAST_MODE_NONE>, InputX1>;
316+ using CastX2 = Bind<Vec::Cast<T2, T1, TRUNCATE_DIV_CAST_MODE_NONE>, InputX2>;
317+ using DivValue = Bind<Vec::Div<T2>, CastX1, CastX2>;
318+ using CastOut = Bind<Vec::Cast<T1, T2, TRUNCATE_DIV_CAST_MODE_NONE>, DivValue>;
319+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, CastOut>;
320+ using Outputs = Elems<OpCopyOut>;
321+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
322+ using OpDag = DAGSch<Outputs, void, MemCfg>;
323+};
324+ 
325+template <typename T1>
326+struct TruncateDivInt {
327+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
328+ using InputX2 = Bind<Vec::CopyInBrc<T1>, Placeholder::In1<T1>>;
329+ using DivValue = Bind<Vec::Div<T1>, InputX1, InputX2>;
330+ using ComputeValue = Bind<TruncIntPostCompute<T1>, InputX1, InputX2, DivValue>;
331+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, ComputeValue>;
332+ using Outputs = Elems<OpCopyOut>;
333+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
334+ using OpDag = DAGSch<Outputs, void, MemCfg>;
335+};
336+ 
337+template <typename T1>
338+struct TruncateDivInt64 {
339+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
340+ using InputX2 = Bind<Vec::CopyInBrc<T1>, Placeholder::In1<T1>>;
341+ using DivValue = Bind<TruncDivInt64<T1>, InputX1, InputX2>;
342+ using OpCopyOut = Bind<Vec::CopyOut<T1>, Placeholder::Out0<T1>, DivValue>;
343+ using Outputs = Elems<OpCopyOut>;
344+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
345+ using OpDag = DAGSch<Outputs, void, MemCfg>;
346+};
347+ 
348+template <typename T1, typename T2, typename PromoteT>
349+struct TruncateDivIntToFloat {
350+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
351+ using InputX2 = Bind<Vec::CopyInBrc<T2>, Placeholder::In1<T2>>;
352+ using CastX1 = Bind<Vec::Cast<PromoteT, T1, TRUNCATE_DIV_CAST_MODE_NONE>, InputX1>;
353+ using CastX2 = Bind<Vec::Cast<PromoteT, T2, TRUNCATE_DIV_CAST_MODE_NONE>, InputX2>;
354+ using DivValue = Bind<Vec::DivHighPrecision<PromoteT>, CastX1, CastX2>;
355+ using TruncateValue = Bind<TruncDag1::TruncCustom<PromoteT>, DivValue>;
356+ using OpCopyOut = Bind<Vec::CopyOut<PromoteT>, Placeholder::Out0<PromoteT>, TruncateValue>;
357+ using Outputs = Elems<OpCopyOut>;
358+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
359+ using OpDag = DAGSch<Outputs, void, MemCfg>;
360+};
361+ 
362+template <typename T1, typename T2, typename PromoteT>
363+struct TruncateDivFloatToLowBit {
364+ using InputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
365+ using InputX2 = Bind<Vec::CopyInBrc<T2>, Placeholder::In1<T2>>;
366+ using CastX2 = Bind<Vec::Cast<PromoteT, T2, TRUNCATE_DIV_CAST_MODE_NONE>, InputX2>;
367+ using DivValue = Bind<Vec::DivHighPrecision<PromoteT>, InputX1, CastX2>;
368+ using TruncateValue = Bind<TruncDag1::TruncCustom<PromoteT>, DivValue>;
369+ using OpCopyOut = Bind<Vec::CopyOut<PromoteT>, Placeholder::Out0<PromoteT>, TruncateValue>;
370+ using Outputs = Elems<OpCopyOut>;
371+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
372+ using OpDag = DAGSch<Outputs, void, MemCfg>;
373+};
374+ 
375+} // namespace TruncateDivOp
376+ 
377+#endif // TRUNCATE_DIV_DAG_H
Amath/truncate_div/op_kernel/arch35/truncate_div_struct.h+27-0
@@ -0,0 +1,27 @@
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 truncate_div_struct.h
13+* \brief truncate_div struct
14+*/
15+#ifndef TRUNCATE_DIV_STRUCT_H_
16+#define TRUNCATE_DIV_STRUCT_H_
17+ 
18+#include "atvoss/broadcast/broadcast_base_struct.h"
19+ 
20+ASCENDC_TPL_ARGS_DECL(
21+ TruncateDiv, BRC_TEMP_SCH_MODE_KEY_DECL(schMode), ASCENDC_TPL_BOOL_DECL(canUseMul, 0, 1));
22+ 
23+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
24+ BRC_TEMP_SCH_MODE_KEY_SEL(schMode)),
25+ ASCENDC_TPL_BOOL_SEL(canUseMul, 0, 1)
26+);
27+#endif // TRUNCATE_DIV_STRUCT_H_
Amath/truncate_div/op_kernel/truncate_div_apt.cpp+94-0
@@ -0,0 +1,94 @@
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 truncate_div_apt.cpp
13+ * \brief truncate_div kernel
14+ */
15+ 
16+#include "kernel_operator.h"
17+#include "arch35/truncate_div_dag.h"
18+#include "arch35/truncate_div_struct.h"
19+#include "atvoss/broadcast/broadcast_sch.h"
20+ 
21+using namespace AscendC;
22+using namespace Ops::Base;
23+ 
24+template <uint64_t schMode, bool canUseMul>
25+__global__ __aicore__ void truncate_div(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
26+{
27+ if constexpr (std::is_same<DTYPE_X1, half>::value && std::is_same<DTYPE_X2, float>::value) {
28+ if constexpr (canUseMul) {
29+ using OpDag = TruncateDivOp::TruncateDivFloatWithCastScalar<half, float, float>::OpDag;
30+ BroadcastSch<schMode, OpDag> sch(tiling);
31+ sch.Process(x1, y);
32+ } else {
33+ using OpDag = TruncateDivOp::TruncateDivFloatWithCast<half, float, float>::OpDag;
34+ BroadcastSch<schMode, OpDag> sch(tiling);
35+ sch.Process(x1, x2, y);
36+ }
37+ } else if constexpr (std::is_same<DTYPE_X1, half>::value || std::is_same<DTYPE_X1, bfloat16_t>::value) {
38+ if constexpr (std::is_same<DTYPE_X2, DTYPE_X1>::value) {
39+ if constexpr (canUseMul) {
40+ using OpDag = TruncateDivOp::TruncateDivFloat16Scalar<DTYPE_X1, float>::OpDag;
41+ BroadcastSch<schMode, OpDag> sch(tiling);
42+ sch.Process(x1, y);
43+ } else {
44+ using OpDag = TruncateDivOp::TruncateDivFloat16<DTYPE_X1, float>::OpDag;
45+ BroadcastSch<schMode, OpDag> sch(tiling);
46+ sch.Process(x1, x2, y);
47+ }
48+ }
49+ } else if constexpr (std::is_same<DTYPE_X1, float>::value) {
50+ if constexpr (std::is_same<DTYPE_X2, float>::value) {
51+ if constexpr (canUseMul) {
52+ using OpDag = TruncateDivOp::TruncateDivFloatScalar<float>::OpDag;
53+ BroadcastSch<schMode, OpDag> sch(tiling);
54+ sch.Process(x1, y);
55+ } else {
56+ using OpDag = TruncateDivOp::TruncateDivFloat<float>::OpDag;
57+ BroadcastSch<schMode, OpDag> sch(tiling);
58+ sch.Process(x1, x2, y);
59+ }
60+ } else if constexpr (std::is_same<DTYPE_X2, int32_t>::value) {
61+ using OpDag = TruncateDivOp::TruncateDivFloatToLowBit<float, int32_t, float>::OpDag;
62+ BroadcastSch<schMode, OpDag> sch(tiling);
63+ sch.Process(x1, x2, y);
64+ } else if constexpr (std::is_same<DTYPE_X2, half>::value) {
65+ using OpDag = TruncateDivOp::TruncateDivFloatToLowBit<float, half, float>::OpDag;
66+ BroadcastSch<schMode, OpDag> sch(tiling);
67+ sch.Process(x1, x2, y);
68+ }
69+ } else if constexpr (std::is_same<DTYPE_X1, int8_t>::value || std::is_same<DTYPE_X1, uint8_t>::value) {
70+ if constexpr (std::is_same<DTYPE_X1, int8_t>::value) {
71+ using OpDag = TruncateDivOp::TruncateDivIntS8<int8_t, half>::OpDag;
72+ BroadcastSch<schMode, OpDag> sch(tiling);
73+ sch.Process(x1, x2, y);
74+ } else {
75+ using OpDag = TruncateDivOp::TruncateDivIntU8<uint8_t, uint16_t>::OpDag;
76+ BroadcastSch<schMode, OpDag> sch(tiling);
77+ sch.Process(x1, x2, y);
78+ }
79+ } else if constexpr (
80+ std::is_same<DTYPE_X2, DTYPE_X1>::value &&
81+ (std::is_same<DTYPE_X1, int16_t>::value || std::is_same<DTYPE_X1, int32_t>::value)) {
82+ using OpDag = TruncateDivOp::TruncateDivInt<DTYPE_X1>::OpDag;
83+ BroadcastSch<schMode, OpDag> sch(tiling);
84+ sch.Process(x1, x2, y);
85+ } else if constexpr (std::is_same<DTYPE_X1, int64_t>::value) {
86+ using OpDag = TruncateDivOp::TruncateDivInt64<DTYPE_X1>::OpDag;
87+ BroadcastSch<schMode, OpDag> sch(tiling);
88+ sch.Process(x1, x2, y);
89+ } else if constexpr (std::is_same<DTYPE_X1, int32_t>::value && std::is_same<DTYPE_X2, float>::value) {
90+ using OpDag = TruncateDivOp::TruncateDivIntToFloat<int32_t, float, float>::OpDag;
91+ BroadcastSch<schMode, OpDag> sch(tiling);
92+ sch.Process(x1, x2, y);
93+ }
94+}
Amath/truncate_div/tests/ut/op_host/arch35/test_truncate_div_tiling_arch35.cpp+182-0
@@ -0,0 +1,182 @@
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 can 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 test_truncate_div_tiling_arch35.cpp
13+ * \brief TruncateDiv tiling test
14+ */
15+ 
16+#include "math/truncate_div/op_host/arch35/truncate_div_tiling_arch35.h"
17+#include <iostream>
18+#include <gtest/gtest.h>
19+#include "tiling_context_faker.h"
20+#include "tiling_case_executor.h"
21+#include "atvoss/broadcast/broadcast_tiling.h"
22+ 
23+using namespace std;
24+using namespace ge;
25+using namespace Ops::Base;
26+ 
27+class TruncateDivTilingTest : public testing::Test {
28+protected:
29+ static void SetUpTestCase()
30+ {
31+ std::cout << "TruncateDivTilingTest SetUp" << std::endl;
32+ }
33+ 
34+ static void TearDownTestCase()
35+ {
36+ std::cout << "TruncateDivTilingTest TearDown" << std::endl;
37+ }
38+};
39+ 
40+TEST_F(TruncateDivTilingTest, truncate_div_fp16_fp_scalar)
41+{
42+ BroadcastCompileInfo compileInfo;
43+ compileInfo.coreNum = 64;
44+ compileInfo.ubSize = 245760;
45+ std::vector<int32_t> x2 = {4};
46+ gert::TilingContextPara tilingContextPara(
47+ "TruncateDiv",
48+ {
49+ {{{8, 128}, {8, 128}}, ge::DT_FLOAT16, ge::FORMAT_ND},
50+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND, true, x2.data()},
51+ },
52+ {
53+ {{{8, 128}, {8, 128}}, ge::DT_INT32, ge::FORMAT_ND},
54+ },
55+ &compileInfo);
56+ uint64_t expectTilingKey = 0b1'00000000'00000111;
57+ std::vector<size_t> expectWorkspaces = {16777216};
58+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
59+}
60+ 
61+TEST_F(TruncateDivTilingTest, truncate_div_fp_fp1)
62+{
63+ BroadcastCompileInfo compileInfo;
64+ compileInfo.coreNum = 64;
65+ compileInfo.ubSize = 245760;
66+ gert::TilingContextPara tilingContextPara(
67+ "TruncateDiv",
68+ {
69+ {{{5, 5, 64, 128}, {5, 5, 64, 128}}, ge::DT_FLOAT, ge::FORMAT_ND},
70+ {{{5, 5, 64, 128}, {5, 5, 64, 128}}, ge::DT_FLOAT, ge::FORMAT_ND},
71+ },
72+ {
73+ {{{5, 5, 64, 128}, {5, 5, 64, 128}}, ge::DT_FLOAT, ge::FORMAT_ND},
74+ },
75+ &compileInfo);
76+ uint64_t expectTilingKey = 0b0'00000000'00001000;
77+ std::vector<size_t> expectWorkspaces = {16777216};
78+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
79+}
80+ 
81+TEST_F(TruncateDivTilingTest, truncate_div_bf16_1)
82+{
83+ BroadcastCompileInfo compileInfo;
84+ compileInfo.coreNum = 64;
85+ compileInfo.ubSize = 245760;
86+ gert::TilingContextPara tilingContextPara(
87+ "TruncateDiv",
88+ {
89+ {{{17772, 1, 2, 1, 2, 1, 2, 1}, {17772, 1, 2, 1, 2, 1, 2, 1}}, ge::DT_BF16, ge::FORMAT_ND},
90+ {{{2, 2, 2, 2, 2, 2, 2}, {2, 2, 2, 2, 2, 2, 2}}, ge::DT_BF16, ge::FORMAT_ND},
91+ },
92+ {
93+ {{{17772, 2, 2, 2, 2, 2, 2, 2}, {17772, 2, 2, 2, 2, 2, 2, 2}}, ge::DT_BF16, ge::FORMAT_ND},
94+ },
95+ &compileInfo);
96+ uint64_t expectTilingKey = 0b0'00000000'00000001;
97+ std::vector<size_t> expectWorkspaces = {16777216};
98+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
99+}
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+ 
121+TEST_F(TruncateDivTilingTest, truncate_div_f32_scalar_3)
122+{
123+ BroadcastCompileInfo compileInfo;
124+ compileInfo.coreNum = 64;
125+ compileInfo.ubSize = 245760;
126+ std::vector<int32_t> x2 = {4};
127+ gert::TilingContextPara tilingContextPara(
128+ "TruncateDiv",
129+ {
130+ {{{17772, 1, 2, 1, 2, 1, 2, 1}, {17772, 1, 2, 1, 2, 1, 2, 1}}, ge::DT_FLOAT, ge::FORMAT_ND},
131+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND, true, x2.data()},
132+ },
133+ {
134+ {{{17772, 2, 2, 2, 2, 2, 2, 2}, {17772, 2, 2, 2, 2, 2, 2, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
135+ },
136+ &compileInfo);
137+ uint64_t expectTilingKey = 0b1'00000000'00000001;
138+ std::vector<size_t> expectWorkspaces = {16777216};
139+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
140+}
141+ 
142+TEST_F(TruncateDivTilingTest, truncate_div_f16_scalar_4)
143+{
144+ BroadcastCompileInfo compileInfo;
145+ compileInfo.coreNum = 64;
146+ compileInfo.ubSize = 245760;
147+ std::vector<int32_t> x2 = {4};
148+ gert::TilingContextPara tilingContextPara(
149+ "TruncateDiv",
150+ {
151+ {{{17772, 1, 2, 1, 2, 1, 2, 1}, {17772, 1, 2, 1, 2, 1, 2, 1}}, ge::DT_FLOAT16, ge::FORMAT_ND},
152+ {{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND, true, x2.data()},
153+ },
154+ {
155+ {{{17772, 2, 2, 2, 2, 2, 2, 2}, {17772, 2, 2, 2, 2, 2, 2, 2}}, ge::DT_FLOAT16, ge::FORMAT_ND},
156+ },
157+ &compileInfo);
158+ uint64_t expectTilingKey = 0b1'00000000'00000001;
159+ std::vector<size_t> expectWorkspaces = {16777216};
160+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
161+}
162+ 
163+TEST_F(TruncateDivTilingTest, truncate_div_bf16_scalar_5)
164+{
165+ BroadcastCompileInfo compileInfo;
166+ compileInfo.coreNum = 64;
167+ compileInfo.ubSize = 245760;
168+ std::vector<int32_t> x2 = {4};
169+ gert::TilingContextPara tilingContextPara(
170+ "TruncateDiv",
171+ {
172+ {{{17772, 1, 2, 1, 2, 1, 2, 1}, {17772, 1, 2, 1, 2, 1, 2, 1}}, ge::DT_BF16, ge::FORMAT_ND},
173+ {{{}, {}}, ge::DT_BF16, ge::FORMAT_ND, true, x2.data()},
174+ },
175+ {
176+ {{{17772, 2, 2, 2, 2, 2, 2, 2}, {17772, 2, 2, 2, 2, 2, 2, 2}}, ge::DT_BF16, ge::FORMAT_ND},
177+ },
178+ &compileInfo);
179+ uint64_t expectTilingKey = 0b1'00000000'00000001;
180+ std::vector<size_t> expectWorkspaces = {16777216};
181+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
182+}
Amath/truncate_div/tests/ut/op_host/test_truncate_div_infershape.cpp+136-0
@@ -0,0 +1,136 @@
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 can 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 test_truncate_div_infershape.cpp
13+ * \brief TruncateDiv infershape test
14+ */
15+ 
16+#include <iostream>
17+#include <gtest/gtest.h>
18+#include "infershape_context_faker.h"
19+#include "infershape_case_executor.h"
20+ 
21+using namespace ge;
22+ 
23+class TruncateDivInfershape : public testing::Test {
24+protected:
25+ static void SetUpTestCase()
26+ {
27+ std::cout << "TruncateDivInfershape SetUp" << std::endl;
28+ }
29+ 
30+ static void TearDownTestCase()
31+ {
32+ std::cout << "TruncateDivInfershape TearDown" << std::endl;
33+ }
34+};
35+ 
36+TEST_F(TruncateDivInfershape, truncate_div_infershape_test_0)
37+{
38+ gert::InfershapeContextPara infershapeContextPara(
39+ "TruncateDiv",
40+ {
41+ {{{2, 2, 1}, {2, 2, 1}}, ge::DT_FLOAT16, ge::FORMAT_ND},
42+ {{{2, 2, 3}, {2, 2, 3}}, ge::DT_FLOAT16, ge::FORMAT_ND},
43+ },
44+ {
45+ {{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND},
46+ });
47+ std::vector<std::vector<int64_t>> expectOutputShape = {
48+ {2, 2, 3},
49+ };
50+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
51+}
52+ 
53+TEST_F(TruncateDivInfershape, truncate_div_infershape_test_1)
54+{
55+ gert::InfershapeContextPara infershapeContextPara(
56+ "TruncateDiv",
57+ {
58+ {{{3, 4, 5, 6, -1}, {3, 4, 5, 6, -1}}, ge::DT_INT32, ge::FORMAT_ND},
59+ {{{3, 4, 5, 6, 1}, {3, 4, 5, 6, 1}}, ge::DT_INT32, ge::FORMAT_ND},
60+ },
61+ {
62+ {{{}, {}}, ge::DT_INT32, ge::FORMAT_ND},
63+ });
64+ std::vector<std::vector<int64_t>> expectOutputShape = {
65+ {3, 4, 5, 6, -1},
66+ };
67+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
68+}
69+ 
70+TEST_F(TruncateDivInfershape, truncate_div_infershape_test_2)
71+{
72+ gert::InfershapeContextPara infershapeContextPara(
73+ "TruncateDiv",
74+ {
75+ {{{4, 2}, {4, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
76+ {{{4, 2}, {4, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
77+ },
78+ {
79+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
80+ });
81+ std::vector<std::vector<int64_t>> expectOutputShape = {
82+ {4, 2},
83+ };
84+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
85+}
86+ 
87+TEST_F(TruncateDivInfershape, truncate_div_infershape_test_3)
88+{
89+ gert::InfershapeContextPara infershapeContextPara(
90+ "TruncateDiv",
91+ {
92+ {{{8, 16, 32}, {8, 16, 32}}, ge::DT_INT8, ge::FORMAT_ND},
93+ {{{8, 16, 32}, {8, 16, 32}}, ge::DT_INT8, ge::FORMAT_ND},
94+ },
95+ {
96+ {{{}, {}}, ge::DT_INT8, ge::FORMAT_ND},
97+ });
98+ std::vector<std::vector<int64_t>> expectOutputShape = {
99+ {8, 16, 32},
100+ };
101+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
102+}
103+ 
104+TEST_F(TruncateDivInfershape, truncate_div_infershape_test_4)
105+{
106+ gert::InfershapeContextPara infershapeContextPara(
107+ "TruncateDiv",
108+ {
109+ {{{1, 128, 256}, {1, 128, 256}}, ge::DT_UINT8, ge::FORMAT_ND},
110+ {{{1, 128, 256}, {1, 128, 256}}, ge::DT_UINT8, ge::FORMAT_ND},
111+ },
112+ {
113+ {{{}, {}}, ge::DT_UINT8, ge::FORMAT_ND},
114+ });
115+ std::vector<std::vector<int64_t>> expectOutputShape = {
116+ {1, 128, 256},
117+ };
118+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
119+}
120+ 
121+TEST_F(TruncateDivInfershape, truncate_div_infershape_test_5)
122+{
123+ gert::InfershapeContextPara infershapeContextPara(
124+ "TruncateDiv",
125+ {
126+ {{{1024}, {1024}}, ge::DT_INT64, ge::FORMAT_ND},
127+ {{{1024}, {1024}}, ge::DT_INT64, ge::FORMAT_ND},
128+ },
129+ {
130+ {{{}, {}}, ge::DT_INT64, ge::FORMAT_ND},
131+ });
132+ std::vector<std::vector<int64_t>> expectOutputShape = {
133+ {1024},
134+ };
135+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
136+}
Mscripts/kernel/binary_config/ascendc_config.json+1-0
@@ -144,6 +144,7 @@
144 {"name":"Mul", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend950": ["--cce-simd-vf-fusion=true", "-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},144 {"name":"Mul", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend950": ["--cce-simd-vf-fusion=true", "-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
145 {"name":"Muls", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},145 {"name":"Muls", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
146 {"name":"Div", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},146 {"name":"Div", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
147+ {"name":"TruncateDiv", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
147 {"name":"FloorDiv", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},148 {"name":"FloorDiv", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
148 {"name":"RealDiv", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},149 {"name":"RealDiv", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
149 {"name":"LogicalOr", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},150 {"name":"LogicalOr", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},