-
Notifications
You must be signed in to change notification settings - Fork 731
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
[SYCL] Fix bitselect builtin for integer types #12598
Changes from all commits
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,59 @@ | ||
// RUN: %{build} -o %t.out | ||
// RUN: %{run} %t.out | ||
// RUN: %if preview-breaking-changes-supported %{ %{build} -fpreview-breaking-changes -o %t_preview.out %} | ||
// RUN: %if preview-breaking-changes-supported %{ %{run} %t_preview.out%} | ||
|
||
#include <sycl/sycl.hpp> | ||
|
||
template <typename... Ts, typename FuncTy> void TestTypes(FuncTy F) { | ||
(F(Ts{}), ...); | ||
} | ||
|
||
int main() { | ||
sycl::queue q; | ||
|
||
auto Test = [&](auto F, auto Expected, auto... Args) { | ||
#if defined(__GNUC__) || defined(__clang__) | ||
std::cout << __PRETTY_FUNCTION__ << std::endl; | ||
#endif | ||
std::tuple ArgsTuple{Args...}; | ||
auto Result = std::apply(F, ArgsTuple); | ||
static_assert(std::is_same_v<decltype(Expected), decltype(Result)>); | ||
assert(Expected == Result); | ||
|
||
sycl::buffer<bool, 1> ResultBuf{1}; | ||
q.submit([&](sycl::handler &cgh) { | ||
sycl::accessor Result{ResultBuf, cgh}; | ||
cgh.single_task([=]() { | ||
auto R = std::apply(F, ArgsTuple); | ||
static_assert(std::is_same_v<decltype(Expected), decltype(R)>); | ||
Result[0] = Expected == R; | ||
}); | ||
}); | ||
assert(sycl::host_accessor{ResultBuf}[0]); | ||
}; | ||
|
||
auto TestBitSelect = [&](auto type_val) { | ||
using T = decltype(type_val); | ||
auto BitSelect = [](auto... xs) { return sycl::bitselect(xs...); }; | ||
|
||
static_assert(std::is_integral_v<T>, | ||
"Only integer test is implemented here!"); | ||
Test(BitSelect, T{0b0110}, T{0b1100}, T{0b0011}, T{0b1010}); | ||
}; | ||
|
||
TestTypes<signed char, unsigned char, char, long, long long, unsigned long, | ||
unsigned long long>(TestBitSelect); | ||
|
||
auto TestSelect = [&](auto type_val) { | ||
using T = decltype(type_val); | ||
auto Select = [](auto... xs) { return sycl::select(xs...); }; | ||
|
||
Test(Select, T{0}, T{1}, T{0}, true); | ||
Test(Select, T{1}, T{1}, T{0}, false); | ||
}; | ||
|
||
TestTypes<signed char, unsigned char, char>(TestSelect); | ||
|
||
return 0; | ||
} |
Original file line number | Diff line number | Diff line change |
---|---|---|
|
@@ -88,6 +88,70 @@ int main() { | |
TEST2(sycl::any, int, EXPECTED(int32_t, 0), 3, va7); | ||
TEST(sycl::bitselect, float, EXPECTED(float, 1.0, 1.0), 2, va8, va9, va10); | ||
TEST(sycl::select, float, EXPECTED(float, 1.0, 2.0, 8.0), 3, va5, va6, c1); | ||
{ | ||
// Extra tests for select/bitselect due to special handling required for | ||
// integer return types. | ||
|
||
auto Test = [&](auto F, auto Expected, auto... Args) { | ||
std::tuple ArgsTuple{Args...}; | ||
auto Result = std::apply(F, ArgsTuple); | ||
static_assert(std::is_same_v<decltype(Expected), decltype(Result)>); | ||
|
||
// Note: operator==(vec, vec) return vec. | ||
auto Equal = [](auto x, auto y) { | ||
for (size_t i = 0; i < x.size(); ++i) | ||
if (x[i] != y[i]) | ||
return false; | ||
|
||
return true; | ||
}; | ||
|
||
assert(Equal(Result, Expected)); | ||
|
||
sycl::buffer<bool, 1> ResultBuf{1}; | ||
deviceQueue.submit([&](sycl::handler &cgh) { | ||
sycl::accessor Result{ResultBuf, cgh}; | ||
cgh.single_task([=]() { | ||
auto R = std::apply(F, ArgsTuple); | ||
static_assert(std::is_same_v<decltype(Expected), decltype(R)>); | ||
Result[0] = Equal(R, Expected); | ||
}); | ||
}); | ||
assert(sycl::host_accessor{ResultBuf}[0]); | ||
}; | ||
|
||
// Note that only int8_t/uint8_t are supported by the bitselect/select | ||
// builtins and not all three char data types. Also, use positive numbers | ||
// for the values below so that we could use the same for both | ||
// signed/unsigned tests. | ||
sycl::vec<uint8_t, 2> a{0b1100, 0b0011}; | ||
sycl::vec<uint8_t, 2> b{0b0011, 0b1100}; | ||
sycl::vec<uint8_t, 2> c{0b1010, 0b1010}; | ||
sycl::vec<uint8_t, 2> r{0b0110, 0b1001}; | ||
|
||
auto BitSelect = [](auto... xs) { return sycl::bitselect(xs...); }; | ||
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. I'm wondering do we have negative test cases for sycl::bitselect (with vector type inputs) to check if all the arguments to sycl::bitselect have same element types and the same vector length? There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. To be honest, I'm relying on CTS to catch every possible obscure corner case. Otherwise, the full testing for all the builtins would be taking way too much time. There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Makes sense. |
||
Test(BitSelect, r, a, b, c); | ||
[&](auto... xs) { | ||
Test(BitSelect, xs.template as<sycl::vec<int8_t, 2>>()...); | ||
}(r, a, b, c); | ||
|
||
auto Select = [](auto... xs) { return sycl::select(xs...); }; | ||
sycl::vec<uint8_t, 2> c2{0x7F, 0xFF}; | ||
sycl::vec<uint8_t, 2> r2{a[0], b[1]}; | ||
|
||
Test(Select, r2, a, b, c2); | ||
[&](auto... xs) { | ||
Test(Select, xs.template as<sycl::vec<int8_t, 2>>()..., c2); | ||
}(r2, a, b); | ||
|
||
// Assume that MSB of a signed data type is the leftmost bit (signbit). | ||
auto c3 = c2.template as<sycl::vec<int8_t, 2>>(); | ||
|
||
Test(Select, r2, a, b, c3); | ||
[&](auto... xs) { | ||
Test(Select, xs.template as<sycl::vec<int8_t, 2>>()..., c3); | ||
}(r2, a, b); | ||
} | ||
|
||
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES | ||
TEST(sycl::isequal, int32_t, EXPECTED(int32_t, 1, 1), 2, va11.swizzle<0, 1>(), | ||
|
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
nit: Perhaps you could add a comment here to state the objective(s) of this test case.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
It's not different from existing
marray_relational.cpp
andvec_relational.cpp
and it never caused troubles to understand what they are doing... Are you asking about preview specifically? If yes, then it's a common idiom and we don't expect to document it in every file that uses it.There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Nope, I was talking about providing a broader overview of this test case. To my inexperienced eyes, it wasn't readily clear.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I think it matches how other builtins tests are set up, so I'm going to ignore it in scope of this PR.