Skip to content

[SYCL][COMPAT] Add math extend_vcompare[2/4] to SYCLCompat #14079

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.

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 42 commits into from
Jun 14, 2024
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
42 commits
Select commit Hold shift + click to select a range
2f908c4
[SYCL][COMPAT] add math extend_v*2 to SYCLCompat
OuadiElfarouki May 29, 2024
f5b7e11
Merge branch 'sycl' into math_extend_v2
OuadiElfarouki May 30, 2024
3fa6edd
minor layout changes
OuadiElfarouki May 30, 2024
033164e
Update sycl/test-e2e/syclcompat/math/math_extend_v.cpp
OuadiElfarouki May 31, 2024
ca1f991
Fixed some gaps in test cases and other typo fixes
OuadiElfarouki May 31, 2024
a9d33b3
Explicit casting to extended uint type
OuadiElfarouki Jun 2, 2024
cd86f19
Minor fixes & reformatting
OuadiElfarouki Jun 2, 2024
d65ab3b
Extended some tests to cover mixed types
OuadiElfarouki Jun 3, 2024
065d920
fix to zero_or_signed_extend
OuadiElfarouki Jun 5, 2024
29ae741
Added extend_v2 APIs doc to readme
OuadiElfarouki Jun 5, 2024
5e577bf
Apply suggestions from code review
OuadiElfarouki Jun 6, 2024
7018d4f
fixed formatting
OuadiElfarouki Jun 6, 2024
18c0163
Merge branch 'sycl' into math_extend_v2
OuadiElfarouki Jun 6, 2024
d9d295c
Merge branch 'sycl' into math_extend_v2
OuadiElfarouki Jun 6, 2024
cfa08ae
[SYCL][COMPAT] add math extend_v*4 to SYCLCompat
OuadiElfarouki Jun 4, 2024
1e04a11
Added missing test
OuadiElfarouki Jun 6, 2024
b3953c1
[SYCL][COMPAT] add math extend_vcompare[2/4] to SYCLCompat
OuadiElfarouki Jun 4, 2024
a007050
Apply suggestions from code review
OuadiElfarouki Jun 6, 2024
f1de443
Added extend_v4 APIs doc
OuadiElfarouki Jun 6, 2024
b2c1c9c
Merge branch 'math_extend_v4' into math_extend_compare_v
OuadiElfarouki Jun 6, 2024
5fa44ba
Added vcompare[2/4] APIs doc
OuadiElfarouki Jun 6, 2024
bcd3698
replaced enable_if with static assertion for extend_vbinary2
OuadiElfarouki Jun 6, 2024
00c7f16
Merge branch 'math_extend_v2' into math_extend_v4
OuadiElfarouki Jun 6, 2024
00e97d1
replaced enable_if with static assertion for extend_vbinary4
OuadiElfarouki Jun 6, 2024
35192c9
Merge branch 'math_extend_v4' into math_extend_compare_v
OuadiElfarouki Jun 6, 2024
8594248
Merge branch 'sycl' into math_extend_v4
OuadiElfarouki Jun 7, 2024
3a84333
Merge branch 'math_extend_v4' into math_extend_compare_v
OuadiElfarouki Jun 7, 2024
6f0e5a0
typo fix
OuadiElfarouki Jun 10, 2024
55d5926
typo fixes
OuadiElfarouki Jun 10, 2024
3fd38fd
Merge branch 'sycl' into math_extend_compare_v
OuadiElfarouki Jun 10, 2024
d78dfc8
Merge branch 'sycl' into math_extend_v4
OuadiElfarouki Jun 10, 2024
2105e49
Update sycl/include/syclcompat/math.hpp
OuadiElfarouki Jun 11, 2024
92be153
Addressed PR comments
OuadiElfarouki Jun 11, 2024
b09a1ff
Merge branch 'math_extend_v4' into math_extend_compare_v
OuadiElfarouki Jun 11, 2024
274363d
tidy values concat using sycl::vec
OuadiElfarouki Jun 11, 2024
d6a1b45
Merge branch 'sycl' into math_extend_compare_v
joeatodd Jun 12, 2024
5e6ad05
Simplified int/uint static type check
OuadiElfarouki Jun 12, 2024
9a538ad
Merge branch 'sycl' into math_extend_v4
OuadiElfarouki Jun 12, 2024
4a79acf
Merge branch 'math_extend_v4' into math_extend_compare_v
OuadiElfarouki Jun 12, 2024
28d9a18
Merge branch 'sycl' into math_extend_compare_v
OuadiElfarouki Jun 12, 2024
edb5744
Break math_extend_v tests per vector size case
OuadiElfarouki Jun 13, 2024
3b60e6a
Merge branch 'sycl' into math_extend_compare_v
OuadiElfarouki Jun 13, 2024
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
67 changes: 67 additions & 0 deletions sycl/doc/syclcompat/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -2416,6 +2416,73 @@ template <typename RetT, typename AT, typename BT>
inline constexpr RetT extend_vavrg4_sat(AT a, BT b, RetT c);
```

Vectorized comparison APIs also provided in the math header behave similarly
and support a `std` comparison operator parameter which can be `greater`,
`less`, `greater_equal`, `less_equal`, `equal_to` or `not_equal_to`. These APIs
cover both the 2-elements *(16-bits each)* and 4-elements *(8-bits each)*
variants, as well as an additional `_add` variant that computes the sum of the
2/4 output elements.

```cpp
/// Extend \p a and \p b to 33 bit and vectorized compare input values using
/// specified comparison \p cmp .
///
/// \tparam [in] AT The type of the first value, can only be 32 bit integer
/// \tparam [in] BT The type of the second value, can only be 32 bit integer
/// \tparam [in] BinaryOperation The type of the compare operation
/// \param [in] a The first value
/// \param [in] b The second value
/// \param [in] cmp The comparsion operator
/// \returns The comparison result of the two extended values.
template <typename AT, typename BT, typename BinaryOperation>
inline constexpr unsigned extend_vcompare2(AT a, BT b, BinaryOperation cmp);

/// Extend Inputs to 33 bit, and vectorized compare input values using specified
/// comparison \p cmp , then add the result with \p c .
///
/// \tparam [in] AT The type of the first value, can only be 32 bit integer
/// \tparam [in] BT The type of the second value, can only be 32 bit integer
/// \tparam [in] BinaryOperation The type of the compare operation
/// \param [in] a The first value
/// \param [in] b The second value
/// \param [in] c The third value
/// \param [in] cmp The comparsion operator
/// \returns The comparison result of the two extended values, and add the
/// result with \p c .
template <typename AT, typename BT, typename BinaryOperation>
inline constexpr unsigned extend_vcompare2_add(AT a, BT b, unsigned c,
BinaryOperation cmp);

/// Extend \p a and \p b to 33 bit and vectorized compare input values using
/// specified comparison \p cmp .
///
/// \tparam [in] AT The type of the first value, can only be 32 bit integer
/// \tparam [in] BT The type of the second value, can only be 32 bit integer
/// \tparam [in] BinaryOperation The type of the compare operation
/// \param [in] a The first value
/// \param [in] b The second value
/// \param [in] cmp The comparsion operator
/// \returns The comparison result of the two extended values.
template <typename AT, typename BT, typename BinaryOperation>
inline constexpr unsigned extend_vcompare4(AT a, BT b, BinaryOperation cmp);

/// Extend Inputs to 33 bit, and vectorized compare input values using specified
/// comparison \p cmp , then add the result with \p c .
///
/// \tparam [in] AT The type of the first value, can only be 32 bit integer
/// \tparam [in] BT The type of the second value, can only be 32 bit integer
/// \tparam [in] BinaryOperation The type of the compare operation
/// \param [in] a The first value
/// \param [in] b The second value
/// \param [in] c The third value
/// \param [in] cmp The comparsion operator
/// \returns The comparison result of the two extended values, and add the
/// result with \p c .
template <typename AT, typename BT, typename BinaryOperation>
inline constexpr unsigned extend_vcompare4_add(AT a, BT b, unsigned c,
BinaryOperation cmp);
```

The math header file provides APIs for bit-field insertion (`bfi_safe`) and
bit-field extraction (`bfe_safe`). These are bounds-checked variants of
underlying `detail` APIs (`detail::bfi`, `detail::bfe`) which, in future
Expand Down
66 changes: 66 additions & 0 deletions sycl/include/syclcompat/math.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -1856,6 +1856,39 @@ inline constexpr RetT extend_vavrg2_sat(AT a, BT b, RetT c) {
return detail::extend_vbinary2<RetT, true, false>(a, b, c, detail::average());
}

/// Extend \p a and \p b to 33 bit and vectorized compare input values using
/// specified comparison \p cmp .
///
/// \tparam [in] AT The type of the first value, can only be 32 bit integer
/// \tparam [in] BT The type of the second value, can only be 32 bit integer
/// \tparam [in] BinaryOperation The type of the compare operation
/// \param [in] a The first value
/// \param [in] b The second value
/// \param [in] cmp The comparsion operator
/// \returns The comparison result of the two extended values.
template <typename AT, typename BT, typename BinaryOperation>
inline constexpr unsigned extend_vcompare2(AT a, BT b, BinaryOperation cmp) {
return detail::extend_vbinary2<unsigned, false, false>(a, b, 0, cmp);
}

/// Extend Inputs to 33 bit, and vectorized compare input values using specified
/// comparison \p cmp , then add the result with \p c .
///
/// \tparam [in] AT The type of the first value, can only be 32 bit integer
/// \tparam [in] BT The type of the second value, can only be 32 bit integer
/// \tparam [in] BinaryOperation The type of the compare operation
/// \param [in] a The first value
/// \param [in] b The second value
/// \param [in] c The third value
/// \param [in] cmp The comparsion operator
/// \returns The comparison result of the two extended values, and add the
/// result with \p c .
template <typename AT, typename BT, typename BinaryOperation>
inline constexpr unsigned extend_vcompare2_add(AT a, BT b, unsigned c,
BinaryOperation cmp) {
return detail::extend_vbinary2<unsigned, false, true>(a, b, c, cmp);
}

/// Compute vectorized addition of \p a and \p b, with each value treated as a
/// 4 elements vector type and extend each element to 9 bit.
/// \tparam [in] RetT The type of the return value, can only be 32 bit integer
Expand Down Expand Up @@ -2121,4 +2154,37 @@ inline constexpr RetT extend_vavrg4_sat(AT a, BT b, RetT c) {
return detail::extend_vbinary4<RetT, true, false>(a, b, c, detail::average());
}

/// Extend \p a and \p b to 33 bit and vectorized compare input values using
/// specified comparison \p cmp .
///
/// \tparam [in] AT The type of the first value, can only be 32 bit integer
/// \tparam [in] BT The type of the second value, can only be 32 bit integer
/// \tparam [in] BinaryOperation The type of the compare operation
/// \param [in] a The first value
/// \param [in] b The second value
/// \param [in] cmp The comparsion operator
/// \returns The comparison result of the two extended values.
template <typename AT, typename BT, typename BinaryOperation>
inline constexpr unsigned extend_vcompare4(AT a, BT b, BinaryOperation cmp) {
return detail::extend_vbinary4<unsigned, false, false>(a, b, 0, cmp);
}

/// Extend Inputs to 33 bit, and vectorized compare input values using specified
/// comparison \p cmp , then add the result with \p c .
///
/// \tparam [in] AT The type of the first value, can only be 32 bit integer
/// \tparam [in] BT The type of the second value, can only be 32 bit integer
/// \tparam [in] BinaryOperation The type of the compare operation
/// \param [in] a The first value
/// \param [in] b The second value
/// \param [in] c The third value
/// \param [in] cmp The comparsion operator
/// \returns The comparison result of the two extended values, and add the
/// result with \p c .
template <typename AT, typename BT, typename BinaryOperation>
inline constexpr unsigned extend_vcompare4_add(AT a, BT b, unsigned c,
BinaryOperation cmp) {
return detail::extend_vbinary4<unsigned, false, true>(a, b, c, cmp);
}

} // namespace syclcompat
Loading
Loading