Skip to content

Commit c9da5c7

Browse files
turinevgenyturinevgeny
authored
[SYCL][ESIMD] Move BFN function from experimental namespace (#10606)
1 parent 9e3852b commit c9da5c7

File tree

4 files changed

+127
-86
lines changed

4 files changed

+127
-86
lines changed

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

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

1087+
/// @addtogroup sycl_esimd_logical
1088+
/// @{
1089+
1090+
/// This enum is used to encode all possible logical operations performed
1091+
/// on the 3 input operands. It is used as a template argument of the bfn()
1092+
/// function.
1093+
/// Example: d = bfn<~bfn_t::x & ~bfn_t::y & ~bfn_t::z>(s0, s1, s2);
1094+
enum class bfn_t : uint8_t { x = 0xAA, y = 0xCC, z = 0xF0 };
1095+
1096+
static constexpr bfn_t operator~(bfn_t x) {
1097+
uint8_t val = static_cast<uint8_t>(x);
1098+
uint8_t res = ~val;
1099+
return static_cast<bfn_t>(res);
1100+
}
1101+
1102+
static constexpr bfn_t operator|(bfn_t x, bfn_t y) {
1103+
uint8_t arg0 = static_cast<uint8_t>(x);
1104+
uint8_t arg1 = static_cast<uint8_t>(y);
1105+
uint8_t res = arg0 | arg1;
1106+
return static_cast<bfn_t>(res);
1107+
}
1108+
1109+
static constexpr bfn_t operator&(bfn_t x, bfn_t y) {
1110+
uint8_t arg0 = static_cast<uint8_t>(x);
1111+
uint8_t arg1 = static_cast<uint8_t>(y);
1112+
uint8_t res = arg0 & arg1;
1113+
return static_cast<bfn_t>(res);
1114+
}
1115+
1116+
static constexpr bfn_t operator^(bfn_t x, bfn_t y) {
1117+
uint8_t arg0 = static_cast<uint8_t>(x);
1118+
uint8_t arg1 = static_cast<uint8_t>(y);
1119+
uint8_t res = arg0 ^ arg1;
1120+
return static_cast<bfn_t>(res);
1121+
}
1122+
1123+
/// Performs binary function computation with three vector operands.
1124+
/// @tparam FuncControl boolean function control expressed with bfn_t
1125+
/// enum values.
1126+
/// @tparam T type of the input vector element.
1127+
/// @tparam N size of the input vector.
1128+
/// @param s0 First boolean function argument.
1129+
/// @param s1 Second boolean function argument.
1130+
/// @param s2 Third boolean function argument.
1131+
template <bfn_t FuncControl, typename T, int N>
1132+
__ESIMD_API std::enable_if_t<std::is_integral_v<T>, __ESIMD_NS::simd<T, N>>
1133+
bfn(__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
1134+
__ESIMD_NS::simd<T, N> src2) {
1135+
if constexpr ((sizeof(T) == 8) || ((sizeof(T) == 1) && (N % 4 == 0)) ||
1136+
((sizeof(T) == 2) && (N % 2 == 0))) {
1137+
// Bitcast Nx8-byte vectors to 2xN vectors of 4-byte integers.
1138+
// Bitcast Nx1-byte vectors to N/4 vectors of 4-byte integers.
1139+
// Bitcast Nx2-byte vectors to N/2 vectors of 4-byte integers.
1140+
auto Result = __ESIMD_NS::bfn<FuncControl>(
1141+
src0.template bit_cast_view<int32_t>().read(),
1142+
src1.template bit_cast_view<int32_t>().read(),
1143+
src2.template bit_cast_view<int32_t>().read());
1144+
return Result.template bit_cast_view<T>();
1145+
} else if constexpr (sizeof(T) == 2 || sizeof(T) == 4) {
1146+
constexpr uint8_t FC = static_cast<uint8_t>(FuncControl);
1147+
return __esimd_bfn<FC, T, N>(src0.data(), src1.data(), src2.data());
1148+
} else if constexpr (N % 2 == 0) {
1149+
// Bitcast Nx1-byte vectors (N is even) to N/2 vectors of 2-byte integers.
1150+
auto Result = __ESIMD_NS::bfn<FuncControl>(
1151+
src0.template bit_cast_view<int16_t>().read(),
1152+
src1.template bit_cast_view<int16_t>().read(),
1153+
src2.template bit_cast_view<int16_t>().read());
1154+
return Result.template bit_cast_view<T>();
1155+
} else {
1156+
// Odd number of 1-byte elements.
1157+
__ESIMD_NS::simd<T, N + 1> Src0, Src1, Src2;
1158+
Src0.template select<N, 1>() = src0;
1159+
Src1.template select<N, 1>() = src1;
1160+
Src2.template select<N, 1>() = src2;
1161+
auto Result = __ESIMD_NS::bfn<FuncControl>(Src0, Src1, Src2);
1162+
return Result.template select<N, 1>();
1163+
}
1164+
}
1165+
1166+
/// Performs binary function computation with three scalar operands.
1167+
/// @tparam FuncControl boolean function control expressed with bfn_t enum
1168+
/// values.
1169+
/// @tparam T type of the input vector element.
1170+
/// @param s0 First boolean function argument.
1171+
/// @param s1 Second boolean function argument.
1172+
/// @param s2 Third boolean function argument.
1173+
template <bfn_t FuncControl, typename T>
1174+
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<
1175+
__ESIMD_DNS::is_esimd_scalar<T>::value && std::is_integral_v<T>, T>
1176+
bfn(T src0, T src1, T src2) {
1177+
__ESIMD_NS::simd<T, 1> Src0 = src0;
1178+
__ESIMD_NS::simd<T, 1> Src1 = src1;
1179+
__ESIMD_NS::simd<T, 1> Src2 = src2;
1180+
__ESIMD_NS::simd<T, 1> Result =
1181+
esimd::bfn<FuncControl, T, 1>(Src0, Src1, Src2);
1182+
return Result[0];
1183+
}
1184+
1185+
/// @} sycl_esimd_logical
1186+
10871187
/// Performs add with carry of 2 unsigned 32-bit vectors.
10881188
/// @tparam N size of the vectors
10891189
/// @param carry vector that is going to hold resulting carry flag

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

Lines changed: 13 additions & 72 deletions
Original file line numberDiff line numberDiff line change
@@ -1667,45 +1667,15 @@ __ESIMD_NS::simd<T, N> dp4(__ESIMD_NS::simd<T, N> v1,
16671667

16681668
/// @} sycl_esimd_math
16691669

1670-
/// @defgroup sycl_esimd_systolic_array_api Systolic Array APIs.
1671-
/// APIs below are used to implement dot product accumulate systolic functions
1672-
/// @ingroup sycl_esimd
1673-
16741670
/// @addtogroup sycl_esimd_logical
16751671
/// @{
16761672

16771673
/// This enum is used to encode all possible logical operations performed
16781674
/// on the 3 input operands. It is used as a template argument of the bfn()
16791675
/// function.
16801676
/// Example: d = bfn<~bfn_t::x & ~bfn_t::y & ~bfn_t::z>(s0, s1, s2);
1681-
enum class bfn_t : uint8_t { x = 0xAA, y = 0xCC, z = 0xF0 };
1682-
1683-
static constexpr bfn_t operator~(bfn_t x) {
1684-
uint8_t val = static_cast<uint8_t>(x);
1685-
uint8_t res = ~val;
1686-
return static_cast<bfn_t>(res);
1687-
}
1688-
1689-
static constexpr bfn_t operator|(bfn_t x, bfn_t y) {
1690-
uint8_t arg0 = static_cast<uint8_t>(x);
1691-
uint8_t arg1 = static_cast<uint8_t>(y);
1692-
uint8_t res = arg0 | arg1;
1693-
return static_cast<bfn_t>(res);
1694-
}
1695-
1696-
static constexpr bfn_t operator&(bfn_t x, bfn_t y) {
1697-
uint8_t arg0 = static_cast<uint8_t>(x);
1698-
uint8_t arg1 = static_cast<uint8_t>(y);
1699-
uint8_t res = arg0 & arg1;
1700-
return static_cast<bfn_t>(res);
1701-
}
1702-
1703-
static constexpr bfn_t operator^(bfn_t x, bfn_t y) {
1704-
uint8_t arg0 = static_cast<uint8_t>(x);
1705-
uint8_t arg1 = static_cast<uint8_t>(y);
1706-
uint8_t res = arg0 ^ arg1;
1707-
return static_cast<bfn_t>(res);
1708-
}
1677+
using bfn_t __SYCL_DEPRECATED("Please use sycl::ext::intel::esimd::bfn_t") =
1678+
__ESIMD_NS::bfn_t;
17091679

17101680
/// Performs binary function computation with three vector operands.
17111681
/// @tparam FuncControl boolean function control expressed with bfn_t
@@ -1716,38 +1686,12 @@ static constexpr bfn_t operator^(bfn_t x, bfn_t y) {
17161686
/// @param s1 Second boolean function argument.
17171687
/// @param s2 Third boolean function argument.
17181688
template <bfn_t FuncControl, typename T, int N>
1719-
__ESIMD_API std::enable_if_t<std::is_integral_v<T>, __ESIMD_NS::simd<T, N>>
1720-
bfn(__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
1689+
__SYCL_DEPRECATED(
1690+
"Please use sycl::ext::intel::esimd::bfn<FuncControl>(src0, src1, src2);")
1691+
__ESIMD_API std::enable_if_t<std::is_integral_v<T>, __ESIMD_NS::simd<T, N>> bfn(
1692+
__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
17211693
__ESIMD_NS::simd<T, N> src2) {
1722-
if constexpr ((sizeof(T) == 8) || ((sizeof(T) == 1) && (N % 4 == 0)) ||
1723-
((sizeof(T) == 2) && (N % 2 == 0))) {
1724-
// Bitcast Nx8-byte vectors to 2xN vectors of 4-byte integers.
1725-
// Bitcast Nx1-byte vectors to N/4 vectors of 4-byte integers.
1726-
// Bitcast Nx2-byte vectors to N/2 vectors of 4-byte integers.
1727-
auto Result = __ESIMD_ENS::bfn<FuncControl>(
1728-
src0.template bit_cast_view<int32_t>().read(),
1729-
src1.template bit_cast_view<int32_t>().read(),
1730-
src2.template bit_cast_view<int32_t>().read());
1731-
return Result.template bit_cast_view<T>();
1732-
} else if constexpr (sizeof(T) == 2 || sizeof(T) == 4) {
1733-
constexpr uint8_t FC = static_cast<uint8_t>(FuncControl);
1734-
return __esimd_bfn<FC, T, N>(src0.data(), src1.data(), src2.data());
1735-
} else if constexpr (N % 2 == 0) {
1736-
// Bitcast Nx1-byte vectors (N is even) to N/2 vectors of 2-byte integers.
1737-
auto Result = __ESIMD_ENS::bfn<FuncControl>(
1738-
src0.template bit_cast_view<int16_t>().read(),
1739-
src1.template bit_cast_view<int16_t>().read(),
1740-
src2.template bit_cast_view<int16_t>().read());
1741-
return Result.template bit_cast_view<T>();
1742-
} else {
1743-
// Odd number of 1-byte elements.
1744-
__ESIMD_NS::simd<T, N + 1> Src0, Src1, Src2;
1745-
Src0.template select<N, 1>() = src0;
1746-
Src1.template select<N, 1>() = src1;
1747-
Src2.template select<N, 1>() = src2;
1748-
auto Result = __ESIMD_ENS::bfn<FuncControl>(Src0, Src1, Src2);
1749-
return Result.template select<N, 1>();
1750-
}
1694+
return __ESIMD_NS::bfn<FuncControl>(src0, src1, src2);
17511695
}
17521696

17531697
/// Performs binary function computation with three scalar operands.
@@ -1758,15 +1702,12 @@ bfn(__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
17581702
/// @param s1 Second boolean function argument.
17591703
/// @param s2 Third boolean function argument.
17601704
template <bfn_t FuncControl, typename T>
1761-
__ESIMD_API std::enable_if_t<
1762-
__ESIMD_DNS::is_esimd_scalar<T>::value && std::is_integral_v<T>, T>
1763-
bfn(T src0, T src1, T src2) {
1764-
__ESIMD_NS::simd<T, 1> Src0 = src0;
1765-
__ESIMD_NS::simd<T, 1> Src1 = src1;
1766-
__ESIMD_NS::simd<T, 1> Src2 = src2;
1767-
__ESIMD_NS::simd<T, 1> Result =
1768-
esimd::bfn<FuncControl, T, 1>(Src0, Src1, Src2);
1769-
return Result[0];
1705+
__SYCL_DEPRECATED(
1706+
"Please use sycl::ext::intel::esimd::bfn<FuncControl>(src0, src1, src2);")
1707+
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T>::value &&
1708+
std::is_integral_v<T>,
1709+
T> bfn(T src0, T src1, T src2) {
1710+
return __ESIMD_NS::bfn<FuncControl>(src0, src1, src2);
17701711
}
17711712

17721713
/// @} sycl_esimd_logical

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

Lines changed: 12 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -38,11 +38,11 @@ template <class T> struct InitOps {
3838

3939
// --- Test boolean control functions.
4040

41-
using bfn_t = experimental::esimd::bfn_t;
41+
using bfn_t = esimd::bfn_t;
4242

43-
constexpr experimental::esimd::bfn_t F1 = bfn_t::x | bfn_t::y | bfn_t::z;
44-
constexpr experimental::esimd::bfn_t F2 = bfn_t::x & bfn_t::y & bfn_t::z;
45-
constexpr experimental::esimd::bfn_t F3 = ~bfn_t::x | bfn_t::y ^ bfn_t::z;
43+
constexpr esimd::bfn_t F1 = bfn_t::x | bfn_t::y | bfn_t::z;
44+
constexpr esimd::bfn_t F2 = bfn_t::x & bfn_t::y & bfn_t::z;
45+
constexpr esimd::bfn_t F3 = ~bfn_t::x | bfn_t::y ^ bfn_t::z;
4646

4747
// --- Template functions calculating given boolean operation on host and device
4848

@@ -51,7 +51,7 @@ enum ArgKind {
5151
AllSca,
5252
};
5353

54-
template <class T, experimental::esimd::bfn_t Op> struct HostFunc;
54+
template <class T, esimd::bfn_t Op> struct HostFunc;
5555

5656
#define DEFINE_HOST_OP(FUNC_CTRL) \
5757
template <class T> struct HostFunc<T, FUNC_CTRL> { \
@@ -76,20 +76,19 @@ DEFINE_HOST_OP(F3);
7676

7777
// --- Specializations per each boolean operation.
7878

79-
template <class T, int N, experimental::esimd::bfn_t Op, int Args = AllVec>
80-
struct ESIMDf;
79+
template <class T, int N, esimd::bfn_t Op, int Args = AllVec> struct ESIMDf;
8180

8281
#define DEFINE_ESIMD_DEVICE_OP(FUNC_CTRL) \
8382
template <class T, int N> struct ESIMDf<T, N, FUNC_CTRL, AllVec> { \
8483
esimd::simd<T, N> \
8584
operator()(esimd::simd<T, N> X0, esimd::simd<T, N> X1, \
8685
esimd::simd<T, N> X2) const SYCL_ESIMD_FUNCTION { \
87-
return experimental::esimd::bfn<FUNC_CTRL, T, N>(X0, X1, X2); \
86+
return esimd::bfn<FUNC_CTRL, T, N>(X0, X1, X2); \
8887
} \
8988
}; \
9089
template <class T, int N> struct ESIMDf<T, N, FUNC_CTRL, AllSca> { \
9190
esimd::simd<T, N> operator()(T X0, T X1, T X2) const SYCL_ESIMD_FUNCTION { \
92-
return experimental::esimd::bfn<FUNC_CTRL, T, N>(X0, X1, X2); \
91+
return esimd::bfn<FUNC_CTRL, T, N>(X0, X1, X2); \
9392
} \
9493
};
9594

@@ -99,8 +98,8 @@ DEFINE_ESIMD_DEVICE_OP(F3);
9998

10099
// --- Generic kernel calculating a binary function operation on array elements.
101100

102-
template <class T, int N, experimental::esimd::bfn_t Op,
103-
template <class, int, experimental::esimd::bfn_t, int> class Kernel>
101+
template <class T, int N, esimd::bfn_t Op,
102+
template <class, int, esimd::bfn_t, int> class Kernel>
104103
struct DeviceFunc {
105104
const T *In0, *In1, *In2;
106105
T *Out;
@@ -136,8 +135,8 @@ struct DeviceFunc {
136135

137136
// --- Generic test function for boolean function.
138137

139-
template <class T, int N, experimental::esimd::bfn_t Op, int Range,
140-
template <class, int, experimental::esimd::bfn_t, int> class Kernel,
138+
template <class T, int N, esimd::bfn_t Op, int Range,
139+
template <class, int, esimd::bfn_t, int> class Kernel,
141140
typename InitF = InitOps<T>>
142141
bool test(queue &Q, const std::string &Name, InitF Init = InitOps<T>{}) {
143142
constexpr size_t Size = Range * N;

‎sycl/test/esimd/math_impl.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -59,7 +59,8 @@ esimd_math_emu(simd<float, 16> x) {
5959
SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd<int, 16>
6060
esimd_bfn(simd<int, 16> x, simd<int, 16> y, simd<int, 16> z) {
6161
simd<int, 16> v =
62-
experimental::esimd::bfn<~bfn_t::x & ~bfn_t::y & ~bfn_t::z>(x, y, z);
62+
esimd::bfn<~esimd::bfn_t::x & ~esimd::bfn_t::y & ~esimd::bfn_t::z>(x, y,
63+
z);
6364
//CHECK: call spir_func noundef <16 x i32> @_Z11__esimd_bfn
6465
return v;
6566
}

0 commit comments

Comments
 (0)