-
Notifications
You must be signed in to change notification settings - Fork 130
[SYCL] Add tests for native math extension #895
Conversation
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.
LGTM but someone else should review
intel/llvm#5747 changes are not covered by existing pre-commit runs, but new test passes, so it's not clear what is the value of this test if it passes w/o changes it supposed to validate. In other words, I expect the test to fail in pre-commit until intel/llvm#5747 is merged. |
Yes, the tests are performed just if the extension is supported, i.e. |
I wonder if we can make the status of the test more clear. Right now it report pass even if nothing was tested. |
I added an assert if the extension is not supported, it this case the test is reported as failed.
In that case, we need all extensions listed as features in SYCL/lit.cfg.py.
Yes, this is an open question. |
/verify with intel/llvm#5747 |
This patch extends the native math definitions in order to include builtins out of the current SYCL specification. In particular, this patch adds a ``tanh`` builtin for floats/halfs and a exp2 builtin for ``halfs`` which are mapped to instructions introduced for ``sm_75`` and above. Tests in intel/llvm-test-suite#895
native_exp2_tester<sycl::half4>(q); | ||
native_exp2_tester<sycl::half8>(q); | ||
native_exp2_tester<sycl::half16>(q); | ||
} |
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.
You are missing tests for marray
in all of these cases. According to the SYCL 2020 spec, the "genfloatf" and "genfloath" pseudo-types include marray
with all possible values of N. I'd suggest testing the same set as vec
(2, 3, 4, 8, 16) and also at least one weird one (e.g. 43).
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.
Unfortunately, I don't think this is possible. At the moment, all built-ins (not only the ones tested by this patch) don't accept marray
's as arguments.
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.
Yes, that does seem to be the case. I created an internal issue to track this missing feature.
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.
Hi @gmlueck : a lot of marray support has been added here: intel/llvm#6038
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.
Thanks for this! I added a note to our internal issue with a link to that PR.
q.submit([&](sycl::handler &cgh) { | ||
auto AccR = BufR.template get_access<sycl::access::mode::write>(cgh); | ||
cgh.single_task([=]() { | ||
AccR[0] = sycl::ext::oneapi::experimental::native::tanh(T(1.0f)); |
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.
This only tests the case when all vector elements have the same value. It would be good to test the case when each element has a different value. That will also give you better coverage for a variety of input values.
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 took an approach similar to the tests present in SYCL/DeviceLib/built-ins
and I generalized it for considering not only float2
but also all the possible vector sizes.
Considering different values is possible, but will complicate the code in a way that could potentially bring problems since the native functions don't have strict accuracy constraints, i.e., each value could have a different range of approximation depending on the device.
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.
Your assert_out_of_bound
checks already test an arbitrary precision range, so it seems like you could use the same technique for other values. I noticed that the range was pretty big, so presumably any reasonable device would compute a value within that range.
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.
Sure, but it is one single range. So, it will be easier to tweak in case of problems. If every value is different up to 16 ranges need to be defined.
If this is necessary, I can modify the code.
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.
Testing only the case when all elements of the vec
have the same value seems like a weak test for this feature. I think it should be expanded to test the case when each element has a different value.
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.
Done.
80cda53
to
1fbd48a
Compare
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.
Looks good, thanks!
Would it be possible to merge this PR? Thanks! |
CMPLRTST-16192: update test driver to support compile llvm_test_suite_esimd_embargo only
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.
Sorry for the late review, @pgorlani .
This patch adds tests for intel/llvm#5747
This patch adds tests for intel#5747
This patch adds tests for intel/llvm#5747