Skip to content

Commit 9e3852b

Browse files
fineg74fineg74
authored
[ESIMD] Move addc() and subb() functions out of experimental namespace (#10589)
1 parent 5dec21d commit 9e3852b

File tree

4 files changed

+170
-54
lines changed

4 files changed

+170
-54
lines changed

‎sycl/include/sycl/ext/intel/esimd/math.hpp

Lines changed: 126 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1084,6 +1084,132 @@ ESIMD_INLINE ESIMD_NODEBUG T0 reduce(simd<T1, SZ> v, BinaryOperation op) {
10841084
}
10851085
}
10861086

1087+
/// Performs add with carry of 2 unsigned 32-bit vectors.
1088+
/// @tparam N size of the vectors
1089+
/// @param carry vector that is going to hold resulting carry flag
1090+
/// @param src0 first term
1091+
/// @param src1 second term
1092+
/// @return sum of 2 terms, carry flag is returned through \c carry parameter
1093+
template <int N>
1094+
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1095+
addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
1096+
__ESIMD_NS::simd<uint32_t, N> src1) {
1097+
std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
1098+
__ESIMD_DNS::vector_type_t<uint32_t, N>>
1099+
Result = __esimd_addc<uint32_t, N>(src0.data(), src1.data());
1100+
1101+
carry = Result.first;
1102+
return Result.second;
1103+
}
1104+
1105+
/// Performs add with carry of a unsigned 32-bit vector and scalar.
1106+
/// @tparam N size of the vectors
1107+
/// @param carry vector that is going to hold resulting carry flag
1108+
/// @param src0 first term
1109+
/// @param src1 second term
1110+
/// @return sum of 2 terms, carry flag is returned through \c carry parameter
1111+
template <int N>
1112+
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1113+
addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
1114+
uint32_t src1) {
1115+
__ESIMD_NS::simd<uint32_t, N> Src1V = src1;
1116+
return addc(carry, src0, Src1V);
1117+
}
1118+
1119+
/// Performs add with carry of a unsigned 32-bit scalar and vector.
1120+
/// @tparam N size of the vectors
1121+
/// @param carry vector that is going to hold resulting carry flag
1122+
/// @param src0 first term
1123+
/// @param src1 second term
1124+
/// @return sum of 2 terms, carry flag is returned through \c carry parameter
1125+
template <int N>
1126+
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1127+
addc(__ESIMD_NS::simd<uint32_t, N> &carry, uint32_t src0,
1128+
__ESIMD_NS::simd<uint32_t, N> src1) {
1129+
__ESIMD_NS::simd<uint32_t, N> Src0V = src0;
1130+
return addc(carry, Src0V, src1);
1131+
}
1132+
1133+
/// Performs add with carry of a unsigned 32-bit scalars.
1134+
/// @tparam N size of the vectors
1135+
/// @param carry scalar that is going to hold resulting carry flag
1136+
/// @param src0 first term
1137+
/// @param src1 second term
1138+
/// @return sum of 2 terms, carry flag is returned through \c carry parameter
1139+
__ESIMD_API uint32_t addc(uint32_t &carry, uint32_t src0, uint32_t src1) {
1140+
__ESIMD_NS::simd<uint32_t, 1> CarryV = carry;
1141+
__ESIMD_NS::simd<uint32_t, 1> Src0V = src0;
1142+
__ESIMD_NS::simd<uint32_t, 1> Src1V = src1;
1143+
__ESIMD_NS::simd<uint32_t, 1> Res = addc(CarryV, Src0V, Src1V);
1144+
carry = CarryV[0];
1145+
return Res[0];
1146+
}
1147+
1148+
/// Performs substraction with borrow of 2 unsigned 32-bit vectors.
1149+
/// @tparam N size of the vectors
1150+
/// @param borrow vector that is going to hold resulting borrow flag
1151+
/// @param src0 first term
1152+
/// @param src1 second term
1153+
/// @return difference of 2 terms, borrow flag is returned through \c borrow
1154+
/// parameter
1155+
template <int N>
1156+
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1157+
subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
1158+
__ESIMD_NS::simd<uint32_t, N> src1) {
1159+
std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
1160+
__ESIMD_DNS::vector_type_t<uint32_t, N>>
1161+
Result = __esimd_subb<uint32_t, N>(src0.data(), src1.data());
1162+
1163+
borrow = Result.first;
1164+
return Result.second;
1165+
}
1166+
1167+
/// Performs substraction with borrow of unsigned 32-bit vector and scalar.
1168+
/// @tparam N size of the vectors
1169+
/// @param borrow vector that is going to hold resulting borrow flag
1170+
/// @param src0 first term
1171+
/// @param src1 second term
1172+
/// @return difference of 2 terms, borrow flag is returned through \c borrow
1173+
/// parameter
1174+
template <int N>
1175+
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1176+
subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
1177+
uint32_t src1) {
1178+
__ESIMD_NS::simd<uint32_t, N> Src1V = src1;
1179+
return subb(borrow, src0, Src1V);
1180+
}
1181+
1182+
/// Performs substraction with borrow of unsigned 32-bit scalar and vector.
1183+
/// @tparam N size of the vectors
1184+
/// @param borrow vector that is going to hold resulting borrow flag
1185+
/// @param src0 first term
1186+
/// @param src1 second term
1187+
/// @return difference of 2 terms, borrow flag is returned through \c borrow
1188+
/// parameter
1189+
template <int N>
1190+
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1191+
subb(__ESIMD_NS::simd<uint32_t, N> &borrow, uint32_t src0,
1192+
__ESIMD_NS::simd<uint32_t, N> src1) {
1193+
__ESIMD_NS::simd<uint32_t, N> Src0V = src0;
1194+
return subb(borrow, Src0V, src1);
1195+
}
1196+
1197+
/// Performs substraction with borrow of 2 unsigned 32-bit scalars.
1198+
/// @tparam N size of the vectors
1199+
/// @param borrow scalar that is going to hold resulting borrow flag
1200+
/// @param src0 first term
1201+
/// @param src1 second term
1202+
/// @return difference of 2 terms, borrow flag is returned through \c borrow
1203+
/// parameter
1204+
__ESIMD_API uint32_t subb(uint32_t &borrow, uint32_t src0, uint32_t src1) {
1205+
__ESIMD_NS::simd<uint32_t, 1> BorrowV = borrow;
1206+
__ESIMD_NS::simd<uint32_t, 1> Src0V = src0;
1207+
__ESIMD_NS::simd<uint32_t, 1> Src1V = src1;
1208+
__ESIMD_NS::simd<uint32_t, 1> Res = subb(BorrowV, Src0V, Src1V);
1209+
borrow = BorrowV[0];
1210+
return Res[0];
1211+
}
1212+
10871213
/// @} sycl_esimd_math
10881214

10891215
} // namespace ext::intel::esimd

‎sycl/include/sycl/ext/intel/experimental/esimd/math.hpp

Lines changed: 42 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -497,77 +497,69 @@ imul(T &rmd, T0 src0, T1 src1) {
497497
}
498498

499499
template <int N>
500-
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
501-
addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
502-
__ESIMD_NS::simd<uint32_t, N> src1) {
503-
std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
504-
__ESIMD_DNS::vector_type_t<uint32_t, N>>
505-
Result = __esimd_addc<uint32_t, N>(src0.data(), src1.data());
506-
507-
carry = Result.first;
508-
return Result.second;
500+
__SYCL_DEPRECATED(
501+
"Please use sycl::ext::intel::esimd::addc(carry, src0, src1);")
502+
__ESIMD_API __ESIMD_NS::simd<uint32_t, N> addc(
503+
__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
504+
__ESIMD_NS::simd<uint32_t, N> src1) {
505+
return __ESIMD_NS::addc(carry, src0, src1);
509506
}
510507

511508
template <int N>
512-
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
513-
addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
514-
uint32_t src1) {
515-
__ESIMD_NS::simd<uint32_t, N> Src1V = src1;
516-
return addc(carry, src0, Src1V);
509+
__SYCL_DEPRECATED(
510+
"Please use sycl::ext::intel::esimd::addc(carry, src0, src1);")
511+
__ESIMD_API __ESIMD_NS::simd<uint32_t, N> addc(
512+
__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
513+
uint32_t src1) {
514+
return __ESIMD_NS::addc(carry, src0, src1);
517515
}
518516

519517
template <int N>
520-
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
521-
addc(__ESIMD_NS::simd<uint32_t, N> &carry, uint32_t src0,
522-
__ESIMD_NS::simd<uint32_t, N> src1) {
523-
__ESIMD_NS::simd<uint32_t, N> Src0V = src0;
524-
return addc(carry, Src0V, src1);
518+
__SYCL_DEPRECATED(
519+
"Please use sycl::ext::intel::esimd::addc(carry, src0, src1);")
520+
__ESIMD_API __ESIMD_NS::simd<uint32_t, N> addc(
521+
__ESIMD_NS::simd<uint32_t, N> &carry, uint32_t src0,
522+
__ESIMD_NS::simd<uint32_t, N> src1) {
523+
return __ESIMD_NS::addc(carry, src0, src1);
525524
}
526525

526+
__SYCL_DEPRECATED(
527+
"Please use sycl::ext::intel::esimd::addc(carry, src0, src1);")
527528
__ESIMD_API uint32_t addc(uint32_t &carry, uint32_t src0, uint32_t src1) {
528-
__ESIMD_NS::simd<uint32_t, 1> CarryV = carry;
529-
__ESIMD_NS::simd<uint32_t, 1> Src0V = src0;
530-
__ESIMD_NS::simd<uint32_t, 1> Src1V = src1;
531-
__ESIMD_NS::simd<uint32_t, 1> Res = addc(CarryV, Src0V, Src1V);
532-
carry = CarryV[0];
533-
return Res[0];
529+
return __ESIMD_NS::addc(carry, src0, src1);
534530
}
535531

536532
template <int N>
537-
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
538-
subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
539-
__ESIMD_NS::simd<uint32_t, N> src1) {
540-
std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
541-
__ESIMD_DNS::vector_type_t<uint32_t, N>>
542-
Result = __esimd_subb<uint32_t, N>(src0.data(), src1.data());
543-
544-
borrow = Result.first;
545-
return Result.second;
533+
__SYCL_DEPRECATED(
534+
"Please use sycl::ext::intel::esimd::subb(borrow, src0, src1);")
535+
__ESIMD_API __ESIMD_NS::simd<uint32_t, N> subb(
536+
__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
537+
__ESIMD_NS::simd<uint32_t, N> src1) {
538+
return __ESIMD_NS::subb(borrow, src0, src1);
546539
}
547540

548541
template <int N>
549-
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
550-
subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
551-
uint32_t src1) {
552-
__ESIMD_NS::simd<uint32_t, N> Src1V = src1;
553-
return subb(borrow, src0, Src1V);
542+
__SYCL_DEPRECATED(
543+
"Please use sycl::ext::intel::esimd::subb(borrow, src0, src1);")
544+
__ESIMD_API __ESIMD_NS::simd<uint32_t, N> subb(
545+
__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
546+
uint32_t src1) {
547+
return __ESIMD_NS::subb(borrow, src0, src1);
554548
}
555549

556550
template <int N>
557-
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
558-
subb(__ESIMD_NS::simd<uint32_t, N> &borrow, uint32_t src0,
559-
__ESIMD_NS::simd<uint32_t, N> src1) {
560-
__ESIMD_NS::simd<uint32_t, N> Src0V = src0;
561-
return subb(borrow, Src0V, src1);
551+
__SYCL_DEPRECATED(
552+
"Please use sycl::ext::intel::esimd::subb(borrow, src0, src1);")
553+
__ESIMD_API __ESIMD_NS::simd<uint32_t, N> subb(
554+
__ESIMD_NS::simd<uint32_t, N> &borrow, uint32_t src0,
555+
__ESIMD_NS::simd<uint32_t, N> src1) {
556+
return __ESIMD_NS::subb(borrow, src0, src1);
562557
}
563558

559+
__SYCL_DEPRECATED(
560+
"Please use sycl::ext::intel::esimd::subb(borrow, src0, src1);")
564561
__ESIMD_API uint32_t subb(uint32_t &borrow, uint32_t src0, uint32_t src1) {
565-
__ESIMD_NS::simd<uint32_t, 1> BorrowV = borrow;
566-
__ESIMD_NS::simd<uint32_t, 1> Src0V = src0;
567-
__ESIMD_NS::simd<uint32_t, 1> Src1V = src1;
568-
__ESIMD_NS::simd<uint32_t, 1> Res = subb(BorrowV, Src0V, Src1V);
569-
borrow = BorrowV[0];
570-
return Res[0];
562+
return __ESIMD_NS::subb(borrow, src0, src1);
571563
}
572564

573565
/// Integral quotient (vector version)

‎sycl/test-e2e/ESIMD/addc.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,6 @@
2020

2121
using namespace sycl;
2222
using namespace sycl::ext::intel::esimd;
23-
namespace iesimd = sycl::ext::intel::experimental::esimd;
2423

2524
template <int N, bool AIsVector, bool BIsVector> bool test(sycl::queue Q) {
2625
static_assert(AIsVector || BIsVector || N == 1,
@@ -84,7 +83,7 @@ template <int N, bool AIsVector, bool BIsVector> bool test(sycl::queue Q) {
8483
using ResType = std::conditional_t<AIsVector || BIsVector,
8584
simd<uint32_t, N>, uint32_t>;
8685
ResType Carry = 0;
87-
ResType Res = iesimd::addc(Carry, A, B);
86+
ResType Res = addc(Carry, A, B);
8887

8988
if constexpr (AIsVector || BIsVector) {
9089
Carry.copy_to(CarryMatrixPtr + (ValuesToTrySize * AI + BI) * N);

‎sycl/test-e2e/ESIMD/subb.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,6 @@
2020

2121
using namespace sycl;
2222
using namespace sycl::ext::intel::esimd;
23-
namespace iesimd = sycl::ext::intel::experimental::esimd;
2423

2524
template <int N, bool AIsVector, bool BIsVector> bool test(sycl::queue Q) {
2625
static_assert(AIsVector || BIsVector || N == 1,
@@ -85,7 +84,7 @@ template <int N, bool AIsVector, bool BIsVector> bool test(sycl::queue Q) {
8584
using ResType = std::conditional_t<AIsVector || BIsVector,
8685
simd<uint32_t, N>, uint32_t>;
8786
ResType Borrow = 0;
88-
ResType Res = iesimd::subb(Borrow, A, B);
87+
ResType Res = subb(Borrow, A, B);
8988

9089
if constexpr (AIsVector || BIsVector) {
9190
Borrow.copy_to(BorrowMatrixPtr + (ValuesToTrySize * AI + BI) * N);

0 commit comments

Comments
 (0)