[SYCL] Add half non-assign math operators #6061

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

Jump to bottom

Merged

pvchupin merged 9 commits into intel:sycl from AidanBeltonS:fix_half_operators

May 9, 2022

+297 −9

Contributor

AidanBeltonS commented Apr 26, 2022 •

edited

Loading

This PR adds missing half operations +-*/.

There are existing operations for the legacy class host_half_impl, however these were not extended to the half class. These operations were being performed as floating point operations via the implicit floating conversion. This results in the output being a float not half type.

The template is limited to arithmetic types to prevent ambiguous templating.
A minor change to built-in __fract is needed to ensure fmin is not ambiguous.

llvm-test-suite PR: intel/llvm-test-suite#1012
Fixes: #6028

AidanBeltonS requested a review from a team as a code owner

April 26, 2022 09:36

AidanBeltonS requested a review from v-klochkov

April 26, 2022 09:36

AidanBeltonS mentioned this pull request

sycl::half arithmetic operator promotes results to float #6028

Closed

v-klochkov reviewed

View reviewed changes

Contributor

v-klochkov left a comment

Thank you for doing this fix. I have some questions though.
Also, please add a LIT test to this PR.

@rolandschulz - would you please take a look at this fix?

sycl/include/CL/sycl/half_type.hpp Outdated Show resolved Hide resolved

sycl/include/CL/sycl/half_type.hpp Outdated

    
                __SYCL_CONSTEXPR_HALF friend half operator op(const half &lhs,               \

                                                              const half &rhs) {             \

                  half rtn = lhs;                                                            \

                  rtn op1 rhs;                                                               \

Contributor

v-klochkov Apr 26, 2022

Suggested change

      
                rtn op1 rhs;                                                               \
          
                rtn op_eq rhs;                                                             \

Contributor Author

AidanBeltonS Apr 27, 2022

Done

sycl/include/CL/sycl/half_type.hpp Outdated Show resolved Hide resolved

sycl/include/CL/sycl/half_type.hpp Outdated Show resolved Hide resolved

AidanBeltonS force-pushed the fix_half_operators branch from bd9353c to 8d9106c Compare

April 27, 2022 10:02

Contributor Author

AidanBeltonS commented Apr 27, 2022

Also, please add a LIT test to this PR.

I have added a test to check for different output types for various inputs.

dkhaldi reviewed

View reviewed changes

sycl/include/CL/sycl/half_type.hpp

    
                OP(+, +=)

                OP(-, -=)

                OP(*, *=)

                OP(/, /=)

Contributor

dkhaldi Apr 27, 2022 •

edited

Loading

Can you also add support for:
OP(bool, ==)
OP(bool, !=)
OP(bool, <)
OP(bool, >)
OP(bool, <=)
OP(bool, >=)
?
Like this, we will have a more complete list similar to what we have for bfloat16 support currently.

Contributor Author

AidanBeltonS Apr 27, 2022

I have added the logical operators and expanded the test

rolandschulz previously approved these changes

View reviewed changes

AidanBeltonS dismissed rolandschulz’s stale review via

aa1ccc1

April 27, 2022 16:09

rolandschulz previously approved these changes

View reviewed changes

v-klochkov previously approved these changes

View reviewed changes

dkhaldi previously approved these changes

View reviewed changes

Contributor

dkhaldi left a comment

Thank you for adding the missing operators

aidan.belton added 6 commits

April 29, 2022 13:39


          Add half non-assign math operators

910e66f


          Make half operations+-*/ constexpr


          Add test and update based on feedback

b000431


          Change op1 to op_eq

b3c7c40


          add logical half operators

e0401ae


          add unsigned types to half operators

fb7f65c

AidanBeltonS dismissed stale reviews from dkhaldi, v-klochkov, and rolandschulz via

fb7f65c

May 2, 2022 09:33

AidanBeltonS force-pushed the fix_half_operators branch from aa1ccc1 to fb7f65c Compare

May 2, 2022 09:33


          clang format

b05e58a

AidanBeltonS mentioned this pull request

[SYCL]Update tests for half operator intel/llvm-test-suite#1012

Merged

Contributor Author

AidanBeltonS commented May 2, 2022

/verify with intel/llvm-test-suite#1012

Contributor Author

AidanBeltonS commented May 2, 2022

Some llvm-test-suite tests were failing due to the half changes. I have updated the test-suite to resolve the issues. The changes relate to ambiguous namespaces with ESIMD and an increase in check tolerance now that 16-bit operations are being performed.

llvm-test-suite PR: intel/llvm-test-suite#1012

aelovikov-intel reviewed

View reviewed changes

sycl/test/type_traits/half_operator_types.cpp Outdated Show resolved Hide resolved

sycl/test/type_traits/half_operator_types.cpp Show resolved Hide resolved


          fix neg, div mistake. Add device-code testing

45d93e6

aelovikov-intel reviewed

View reviewed changes

sycl/test/type_traits/half_operator_types.cpp Outdated

Comment on lines 16 to 24

+                static_assert(std::is_same_v<decltype(T1(1) + sycl::half(1)), T_rtn>);
+                static_assert(std::is_same_v<decltype(T1(1) - sycl::half(1)), T_rtn>);
+                static_assert(std::is_same_v<decltype(T1(1) * sycl::half(1)), T_rtn>);
+                static_assert(std::is_same_v<decltype(T1(1) / sycl::half(1)), T_rtn>);
+                static_assert(std::is_same_v<decltype(sycl::half(1) + T1(1)), T_rtn>);
+                static_assert(std::is_same_v<decltype(sycl::half(1) - T1(1)), T_rtn>);
+                static_assert(std::is_same_v<decltype(sycl::half(1) * T1(1)), T_rtn>);
+                static_assert(std::is_same_v<decltype(sycl::half(1) / T1(1)), T_rtn>);

Contributor

aelovikov-intel May 3, 2022

Can we outline that into a helper called from both host/device code to avoid duplication?

Contributor Author

AidanBeltonS May 3, 2022

Yes, that would make the code much clearer. I have updated the test.

sycl/test/type_traits/half_operator_types.cpp Outdated

+              void check_half_math_operator_types(sycl::queue Queue) {
+                // Test on host
+                static_assert(std::is_same_v<decltype(T1(1) + sycl::half(1)), T_rtn>);

Contributor

aelovikov-intel May 3, 2022

nit: that could also be (I think)

decltype(std::declval<T1>() + std::declval<sycl::half>())

But it's probably not worth decreased signal/noise ratio to be that pedantic here.

Contributor Author

AidanBeltonS May 3, 2022

Done

sycl/test/type_traits/half_operator_types.cpp Outdated

Comment on lines 46 to 58

+                static_assert(std::is_same_v<decltype(T1(1) == sycl::half(1)), bool>);
+                static_assert(std::is_same_v<decltype(T1(1) != sycl::half(1)), bool>);
+                static_assert(std::is_same_v<decltype(T1(1) > sycl::half(1)), bool>);
+                static_assert(std::is_same_v<decltype(T1(1) < sycl::half(1)), bool>);
+                static_assert(std::is_same_v<decltype(T1(1) <= sycl::half(1)), bool>);
+                static_assert(std::is_same_v<decltype(T1(1) <= sycl::half(1)), bool>);
+                static_assert(std::is_same_v<decltype(sycl::half(1) == T1(1)), bool>);
+                static_assert(std::is_same_v<decltype(sycl::half(1) != T1(1)), bool>);
+                static_assert(std::is_same_v<decltype(sycl::half(1) > T1(1)), bool>);
+                static_assert(std::is_same_v<decltype(sycl::half(1) < T1(1)), bool>);
+                static_assert(std::is_same_v<decltype(sycl::half(1) <= T1(1)), bool>);
+                static_assert(std::is_same_v<decltype(sycl::half(1) <= T1(1)), bool>);

Contributor

aelovikov-intel May 3, 2022

Likewise about helper routine.

Contributor Author

AidanBeltonS May 3, 2022

Updated same as above.

sycl/test/type_traits/half_operator_types.cpp Outdated

+              }
+              template <typename T1>
+              void check_half_logical_operator_types(sycl::queue Queue) {

Contributor

aelovikov-intel May 3, 2022

I am very surprised copy-ctor isn't deleted... Regardless of what the standard says, can we turn this into a reference?

Contributor Author

AidanBeltonS May 3, 2022

Yes that makes sense, I have adjusted this to be pass by reference.


          update test based on feedback

b71bb8d

aelovikov-intel approved these changes

View reviewed changes

Contributor Author

AidanBeltonS commented May 3, 2022

/verify with intel/llvm-test-suite#1012

Contributor Author

AidanBeltonS commented May 9, 2022

@pvchupin would it be possible to find out what is failing in the llvm-test-suite when run with the intel/llvm-test-suite#1012 changes? Thank you.

Contributor

pvchupin commented May 9, 2022

@pvchupin would it be possible to find out what is failing in the llvm-test-suite when run with the intel/llvm-test-suite#1012 changes? Thank you.

It actually didn't run and failed earlier due to some infrastructure issue. Tagging @tfzhu.

Contributor

pvchupin commented May 9, 2022

@AidanBeltonS, I'm restarting it now.

Contributor

pvchupin commented May 9, 2022

@AidanBeltonS, restart passed, I'll merge both this and test changes.

pvchupin pushed a commit to intel/llvm-test-suite that referenced this pull request


          [SYCL] Update tests for half operator (#1012)

5dd32c5

This PR makes some small changes due to the introduction of some missing operators.
It changes the namespaces in ext_math.cpp as it was ambiguous as to if the sycl::abs() or sycl::ext::intel::esimd::abs() function should be used.

There is also an adjustment of the load store's check tolerance. The checks tolerance is slightly increased, as previously these operations would have been performed as 32-bit floats.

Depends on: intel/llvm#6061

pvchupin merged commit 059efbc into intel:sycl

Contributor

yanfeng3721 commented May 10, 2022

@AidanBeltonS , could you please check the following post commit error in Windows Build? Thanks.
https://github.com/intel/llvm/runs/6361728718?check_suite_focus=true

[2700/2961] Building CXX object tools\sycl\source\CMakeFiles\sycl_object.dir\detail\reduction.cpp.obj
Alignment of class vec is not in accordance with SYCL specification requirements, a limitation of the MSVC compiler(Error C2719).Requested alignment applied, limited at 64.

Contributor

pvchupin commented May 10, 2022

I think real issue is this one:

D:\github_work\llvm\llvm\src\sycl\source\detail\builtins_relational.cpp(230): error C2220: the following warning is treated as an error
D:\github_work\llvm\llvm\src\sycl\source\detail\builtins_relational.cpp(230): warning C4804: '-': unsafe use of type 'bool' in operation
D:\github_work\llvm\llvm\src\sycl\source\detail\builtins_relational.cpp(24): warning C4804: '-': unsafe use of type 'bool' in operation
D:\github_work\llvm\llvm\src\sycl\source\detail\builtins_relational.cpp(155): note: see reference to function template instantiation 'T cl::__host_std::`anonymous-namespace'::__vFOrdEqualcl::sycl::cl_half(T,T)' being compiled
with
[
T=cl::sycl::cl_half
]

Contributor

yanfeng3721 commented May 10, 2022

I think real issue is this one:

D:\github_work\llvm\llvm\src\sycl\source\detail\builtins_relational.cpp(230): error C2220: the following warning is treated as an error
D:\github_work\llvm\llvm\src\sycl\source\detail\builtins_relational.cpp(230): warning C4804: '-': unsafe use of type 'bool' in operation
D:\github_work\llvm\llvm\src\sycl\source\detail\builtins_relational.cpp(24): warning C4804: '-': unsafe use of type 'bool' in operation
D:\github_work\llvm\llvm\src\sycl\source\detail\builtins_relational.cpp(155): note: see reference to function template instantiation 'T cl::__host_std::`anonymous-namespace'::__vFOrdEqualcl::sycl::cl_half(T,T)' being compiled
with
[
T=cl::sycl::cl_half
]

Thanks for correction, I did not notice "Error C2719" exist before this commit.

AidanBeltonS mentioned this pull request

[SYCL] Add static cast for half builtins relational #6133

Merged

pvchupin pushed a commit that referenced this pull request


          [SYCL] Add static cast for half builtins relational (#6133)

This PR adds a static cast to builtins_relational functions which apply a negative operation to a boolean type.
This is only applied to functions which are specialised for cl_half. The cast prevents windows error message C4804.

This resolves a windows build issue introduced by #6061

Contributor

pvchupin commented May 11, 2022

@AidanBeltonS, it seems there was some collision of this change with recent pulldown. Can you please take a look at the latest post-commit issue?
https://github.com/intel/llvm/runs/6377899714?check_suite_focus=true

pvchupin mentioned this pull request

post-commit fail on Windows: use of overloaded operator is ambiguous #6142

Closed

aelovikov-intel pushed a commit to aelovikov-intel/llvm that referenced this pull request


          [SYCL] Update tests for half operator (intel/llvm-test-suite#1012)

3ad1cc6

This PR makes some small changes due to the introduction of some missing operators.
It changes the namespaces in ext_math.cpp as it was ambiguous as to if the sycl::abs() or sycl::ext::intel::esimd::abs() function should be used.

There is also an adjustment of the load store's check tolerance. The checks tolerance is slightly increased, as previously these operations would have been performed as 32-bit floats.

Depends on: intel#6061

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet