Skip to content

[SYCL][ESIMD] Move BFN function from experimental namespace #10606

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 5 commits into from
Jul 31, 2023
Merged

[SYCL][ESIMD] Move BFN function from experimental namespace #10606

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
3ba12e1
[SYCL][ESIMD] Move BFN function from experimental namespace
turinevgeny Jul 28, 2023
c225450
Employ 'using' to not duplicate operators for bfn
turinevgeny Jul 28, 2023
47c2592
Merge branch 'sycl' into MoveBFNFromExperimental
turinevgeny Jul 28, 2023
fa7160c
Remove unneeed static casts
turinevgeny Jul 28, 2023
fc13b9c
Merge branch 'sycl' into MoveBFNFromExperimental
turinevgeny Jul 29, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
100 changes: 100 additions & 0 deletions sycl/include/sycl/ext/intel/esimd/math.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1084,6 +1084,106 @@ ESIMD_INLINE ESIMD_NODEBUG T0 reduce(simd<T1, SZ> v, BinaryOperation op) {
}
}

/// @addtogroup sycl_esimd_logical
/// @{

/// This enum is used to encode all possible logical operations performed
/// on the 3 input operands. It is used as a template argument of the bfn()
/// function.
/// Example: d = bfn<~bfn_t::x & ~bfn_t::y & ~bfn_t::z>(s0, s1, s2);
enum class bfn_t : uint8_t { x = 0xAA, y = 0xCC, z = 0xF0 };

static constexpr bfn_t operator~(bfn_t x) {
uint8_t val = static_cast<uint8_t>(x);
uint8_t res = ~val;
return static_cast<bfn_t>(res);
}

static constexpr bfn_t operator|(bfn_t x, bfn_t y) {
uint8_t arg0 = static_cast<uint8_t>(x);
uint8_t arg1 = static_cast<uint8_t>(y);
uint8_t res = arg0 | arg1;
return static_cast<bfn_t>(res);
}

static constexpr bfn_t operator&(bfn_t x, bfn_t y) {
uint8_t arg0 = static_cast<uint8_t>(x);
uint8_t arg1 = static_cast<uint8_t>(y);
uint8_t res = arg0 & arg1;
return static_cast<bfn_t>(res);
}

static constexpr bfn_t operator^(bfn_t x, bfn_t y) {
uint8_t arg0 = static_cast<uint8_t>(x);
uint8_t arg1 = static_cast<uint8_t>(y);
uint8_t res = arg0 ^ arg1;
return static_cast<bfn_t>(res);
}

/// Performs binary function computation with three vector operands.
/// @tparam FuncControl boolean function control expressed with bfn_t
/// enum values.
/// @tparam T type of the input vector element.
/// @tparam N size of the input vector.
/// @param s0 First boolean function argument.
/// @param s1 Second boolean function argument.
/// @param s2 Third boolean function argument.
template <bfn_t FuncControl, typename T, int N>
__ESIMD_API std::enable_if_t<std::is_integral_v<T>, __ESIMD_NS::simd<T, N>>
bfn(__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
__ESIMD_NS::simd<T, N> src2) {
if constexpr ((sizeof(T) == 8) || ((sizeof(T) == 1) && (N % 4 == 0)) ||
((sizeof(T) == 2) && (N % 2 == 0))) {
// Bitcast Nx8-byte vectors to 2xN vectors of 4-byte integers.
// Bitcast Nx1-byte vectors to N/4 vectors of 4-byte integers.
// Bitcast Nx2-byte vectors to N/2 vectors of 4-byte integers.
auto Result = __ESIMD_NS::bfn<FuncControl>(
src0.template bit_cast_view<int32_t>().read(),
src1.template bit_cast_view<int32_t>().read(),
src2.template bit_cast_view<int32_t>().read());
return Result.template bit_cast_view<T>();
} else if constexpr (sizeof(T) == 2 || sizeof(T) == 4) {
constexpr uint8_t FC = static_cast<uint8_t>(FuncControl);
return __esimd_bfn<FC, T, N>(src0.data(), src1.data(), src2.data());
} else if constexpr (N % 2 == 0) {
// Bitcast Nx1-byte vectors (N is even) to N/2 vectors of 2-byte integers.
auto Result = __ESIMD_NS::bfn<FuncControl>(
src0.template bit_cast_view<int16_t>().read(),
src1.template bit_cast_view<int16_t>().read(),
src2.template bit_cast_view<int16_t>().read());
return Result.template bit_cast_view<T>();
} else {
// Odd number of 1-byte elements.
__ESIMD_NS::simd<T, N + 1> Src0, Src1, Src2;
Src0.template select<N, 1>() = src0;
Src1.template select<N, 1>() = src1;
Src2.template select<N, 1>() = src2;
auto Result = __ESIMD_NS::bfn<FuncControl>(Src0, Src1, Src2);
return Result.template select<N, 1>();
}
}

/// Performs binary function computation with three scalar operands.
/// @tparam FuncControl boolean function control expressed with bfn_t enum
/// values.
/// @tparam T type of the input vector element.
/// @param s0 First boolean function argument.
/// @param s1 Second boolean function argument.
/// @param s2 Third boolean function argument.
template <bfn_t FuncControl, typename T>
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<
__ESIMD_DNS::is_esimd_scalar<T>::value && std::is_integral_v<T>, T>
bfn(T src0, T src1, T src2) {
__ESIMD_NS::simd<T, 1> Src0 = src0;
__ESIMD_NS::simd<T, 1> Src1 = src1;
__ESIMD_NS::simd<T, 1> Src2 = src2;
__ESIMD_NS::simd<T, 1> Result =
esimd::bfn<FuncControl, T, 1>(Src0, Src1, Src2);
return Result[0];
}

/// @} sycl_esimd_logical

/// Performs add with carry of 2 unsigned 32-bit vectors.
/// @tparam N size of the vectors
/// @param carry vector that is going to hold resulting carry flag
Expand Down
85 changes: 13 additions & 72 deletions sycl/include/sycl/ext/intel/experimental/esimd/math.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1667,45 +1667,15 @@ __ESIMD_NS::simd<T, N> dp4(__ESIMD_NS::simd<T, N> v1,

/// @} sycl_esimd_math

/// @defgroup sycl_esimd_systolic_array_api Systolic Array APIs.
/// APIs below are used to implement dot product accumulate systolic functions
/// @ingroup sycl_esimd

/// @addtogroup sycl_esimd_logical
/// @{

/// This enum is used to encode all possible logical operations performed
/// on the 3 input operands. It is used as a template argument of the bfn()
/// function.
/// Example: d = bfn<~bfn_t::x & ~bfn_t::y & ~bfn_t::z>(s0, s1, s2);
enum class bfn_t : uint8_t { x = 0xAA, y = 0xCC, z = 0xF0 };

static constexpr bfn_t operator~(bfn_t x) {
uint8_t val = static_cast<uint8_t>(x);
uint8_t res = ~val;
return static_cast<bfn_t>(res);
}

static constexpr bfn_t operator|(bfn_t x, bfn_t y) {
uint8_t arg0 = static_cast<uint8_t>(x);
uint8_t arg1 = static_cast<uint8_t>(y);
uint8_t res = arg0 | arg1;
return static_cast<bfn_t>(res);
}

static constexpr bfn_t operator&(bfn_t x, bfn_t y) {
uint8_t arg0 = static_cast<uint8_t>(x);
uint8_t arg1 = static_cast<uint8_t>(y);
uint8_t res = arg0 & arg1;
return static_cast<bfn_t>(res);
}

static constexpr bfn_t operator^(bfn_t x, bfn_t y) {
uint8_t arg0 = static_cast<uint8_t>(x);
uint8_t arg1 = static_cast<uint8_t>(y);
uint8_t res = arg0 ^ arg1;
return static_cast<bfn_t>(res);
}
using bfn_t __SYCL_DEPRECATED("Please use sycl::ext::intel::esimd::bfn_t") =
__ESIMD_NS::bfn_t;

/// Performs binary function computation with three vector operands.
/// @tparam FuncControl boolean function control expressed with bfn_t
Expand All @@ -1716,38 +1686,12 @@ static constexpr bfn_t operator^(bfn_t x, bfn_t y) {
/// @param s1 Second boolean function argument.
/// @param s2 Third boolean function argument.
template <bfn_t FuncControl, typename T, int N>
__ESIMD_API std::enable_if_t<std::is_integral_v<T>, __ESIMD_NS::simd<T, N>>
bfn(__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
__SYCL_DEPRECATED(
"Please use sycl::ext::intel::esimd::bfn<FuncControl>(src0, src1, src2);")
__ESIMD_API std::enable_if_t<std::is_integral_v<T>, __ESIMD_NS::simd<T, N>> bfn(
__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
__ESIMD_NS::simd<T, N> src2) {
if constexpr ((sizeof(T) == 8) || ((sizeof(T) == 1) && (N % 4 == 0)) ||
((sizeof(T) == 2) && (N % 2 == 0))) {
// Bitcast Nx8-byte vectors to 2xN vectors of 4-byte integers.
// Bitcast Nx1-byte vectors to N/4 vectors of 4-byte integers.
// Bitcast Nx2-byte vectors to N/2 vectors of 4-byte integers.
auto Result = __ESIMD_ENS::bfn<FuncControl>(
src0.template bit_cast_view<int32_t>().read(),
src1.template bit_cast_view<int32_t>().read(),
src2.template bit_cast_view<int32_t>().read());
return Result.template bit_cast_view<T>();
} else if constexpr (sizeof(T) == 2 || sizeof(T) == 4) {
constexpr uint8_t FC = static_cast<uint8_t>(FuncControl);
return __esimd_bfn<FC, T, N>(src0.data(), src1.data(), src2.data());
} else if constexpr (N % 2 == 0) {
// Bitcast Nx1-byte vectors (N is even) to N/2 vectors of 2-byte integers.
auto Result = __ESIMD_ENS::bfn<FuncControl>(
src0.template bit_cast_view<int16_t>().read(),
src1.template bit_cast_view<int16_t>().read(),
src2.template bit_cast_view<int16_t>().read());
return Result.template bit_cast_view<T>();
} else {
// Odd number of 1-byte elements.
__ESIMD_NS::simd<T, N + 1> Src0, Src1, Src2;
Src0.template select<N, 1>() = src0;
Src1.template select<N, 1>() = src1;
Src2.template select<N, 1>() = src2;
auto Result = __ESIMD_ENS::bfn<FuncControl>(Src0, Src1, Src2);
return Result.template select<N, 1>();
}
return __ESIMD_NS::bfn<FuncControl>(src0, src1, src2);
}

/// Performs binary function computation with three scalar operands.
Expand All @@ -1758,15 +1702,12 @@ bfn(__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
/// @param s1 Second boolean function argument.
/// @param s2 Third boolean function argument.
template <bfn_t FuncControl, typename T>
__ESIMD_API std::enable_if_t<
__ESIMD_DNS::is_esimd_scalar<T>::value && std::is_integral_v<T>, T>
bfn(T src0, T src1, T src2) {
__ESIMD_NS::simd<T, 1> Src0 = src0;
__ESIMD_NS::simd<T, 1> Src1 = src1;
__ESIMD_NS::simd<T, 1> Src2 = src2;
__ESIMD_NS::simd<T, 1> Result =
esimd::bfn<FuncControl, T, 1>(Src0, Src1, Src2);
return Result[0];
__SYCL_DEPRECATED(
"Please use sycl::ext::intel::esimd::bfn<FuncControl>(src0, src1, src2);")
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T>::value &&
std::is_integral_v<T>,
T> bfn(T src0, T src1, T src2) {
return __ESIMD_NS::bfn<FuncControl>(src0, src1, src2);
}

/// @} sycl_esimd_logical
Expand Down
25 changes: 12 additions & 13 deletions sycl/test-e2e/ESIMD/bfn.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,11 +38,11 @@ template <class T> struct InitOps {

// --- Test boolean control functions.

using bfn_t = experimental::esimd::bfn_t;
using bfn_t = esimd::bfn_t;

constexpr experimental::esimd::bfn_t F1 = bfn_t::x | bfn_t::y | bfn_t::z;
constexpr experimental::esimd::bfn_t F2 = bfn_t::x & bfn_t::y & bfn_t::z;
constexpr experimental::esimd::bfn_t F3 = ~bfn_t::x | bfn_t::y ^ bfn_t::z;
constexpr esimd::bfn_t F1 = bfn_t::x | bfn_t::y | bfn_t::z;
constexpr esimd::bfn_t F2 = bfn_t::x & bfn_t::y & bfn_t::z;
constexpr esimd::bfn_t F3 = ~bfn_t::x | bfn_t::y ^ bfn_t::z;

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

Expand All @@ -51,7 +51,7 @@ enum ArgKind {
AllSca,
};

template <class T, experimental::esimd::bfn_t Op> struct HostFunc;
template <class T, esimd::bfn_t Op> struct HostFunc;

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

// --- Specializations per each boolean operation.

template <class T, int N, experimental::esimd::bfn_t Op, int Args = AllVec>
struct ESIMDf;
template <class T, int N, esimd::bfn_t Op, int Args = AllVec> struct ESIMDf;

#define DEFINE_ESIMD_DEVICE_OP(FUNC_CTRL) \
template <class T, int N> struct ESIMDf<T, N, FUNC_CTRL, AllVec> { \
esimd::simd<T, N> \
operator()(esimd::simd<T, N> X0, esimd::simd<T, N> X1, \
esimd::simd<T, N> X2) const SYCL_ESIMD_FUNCTION { \
return experimental::esimd::bfn<FUNC_CTRL, T, N>(X0, X1, X2); \
return esimd::bfn<FUNC_CTRL, T, N>(X0, X1, X2); \
} \
}; \
template <class T, int N> struct ESIMDf<T, N, FUNC_CTRL, AllSca> { \
esimd::simd<T, N> operator()(T X0, T X1, T X2) const SYCL_ESIMD_FUNCTION { \
return experimental::esimd::bfn<FUNC_CTRL, T, N>(X0, X1, X2); \
return esimd::bfn<FUNC_CTRL, T, N>(X0, X1, X2); \
} \
};

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

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

template <class T, int N, experimental::esimd::bfn_t Op,
template <class, int, experimental::esimd::bfn_t, int> class Kernel>
template <class T, int N, esimd::bfn_t Op,
template <class, int, esimd::bfn_t, int> class Kernel>
struct DeviceFunc {
const T *In0, *In1, *In2;
T *Out;
Expand Down Expand Up @@ -136,8 +135,8 @@ struct DeviceFunc {

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

template <class T, int N, experimental::esimd::bfn_t Op, int Range,
template <class, int, experimental::esimd::bfn_t, int> class Kernel,
template <class T, int N, esimd::bfn_t Op, int Range,
template <class, int, esimd::bfn_t, int> class Kernel,
typename InitF = InitOps<T>>
bool test(queue &Q, const std::string &Name, InitF Init = InitOps<T>{}) {
constexpr size_t Size = Range * N;
Expand Down
3 changes: 2 additions & 1 deletion sycl/test/esimd/math_impl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,7 +59,8 @@ esimd_math_emu(simd<float, 16> x) {
SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd<int, 16>
esimd_bfn(simd<int, 16> x, simd<int, 16> y, simd<int, 16> z) {
simd<int, 16> v =
experimental::esimd::bfn<~bfn_t::x & ~bfn_t::y & ~bfn_t::z>(x, y, z);
esimd::bfn<~esimd::bfn_t::x & ~esimd::bfn_t::y & ~esimd::bfn_t::z>(x, y,
z);
//CHECK: call spir_func noundef <16 x i32> @_Z11__esimd_bfn
return v;
}