Skip to content
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

Add operators for +,-,*,/ for half with other scalar types #969

Merged
merged 7 commits into from
Mar 30, 2023

Conversation

nilsfriess
Copy link
Collaborator

This also required removing the explicit keyword for the constructors and adding three new constructors for unsigned int, unsigned long and long.

@illuhad
Copy link
Collaborator

illuhad commented Mar 13, 2023

Can you comment more on the use case? What is the code you attempt to make work?

@nilsfriess
Copy link
Collaborator Author

Things like

sycl::half a(1);
float b = 1;
auto res = a * b;

currently do not compile because the multiplication is ambiguous (the compiler could first create a half from b and then use operator*(const half& a, const half& b) or convert a to float and use the built-in operator*(float, float)).
There are a few places in the CTS where scalars are multiplied with/divided by half and to compile the CTS it would probably suffice to just implement the operators for float. I just added the other ones for completeness.

@illuhad
Copy link
Collaborator

illuhad commented Mar 23, 2023

The SYCL WG agreed that half should indeed only have the float conversion operator: KhronosGroup/SYCL-CTS#559 (comment)

return fp16::builtin_add(a._data, b._data),
return fp16::builtin_add(a._data, b._data))
data = fp16::builtin_add(a._data, b._data),
data = fp16::builtin_add(a._data, b._data));
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Couldn't you simplify this along the lines of:

return detail::create_half(__hipsycl_backend_switch(
  fp16::builtin_add(a._data, b._data),
  ....
));

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Unfortunately not, because __hipsycl_backend_switch adds semicolons to its arguments, so this becomes something like

return detail::create_half(fp16::builtin_add(a._data, b._data););

which does not compile (this could of course be fixed by removing the semicolons in the definition of __hipsycl_backend_switch and adding them in the places where the macro is used). However, for the sscp backend, __hipsycl_backend_switch becomes an if statement, which again doesn't compile, and we can't just change that to a ternary operator because sometimes __hipsycl_backend_switch is used with return statements as its arguments and these are not allowed as expressions in the ternary operator.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ugh... even if SSCP might in principle work with ternary operator, it probably won't work with nvc++, where we need if target() branching.. So you are totally right :(

@illuhad
Copy link
Collaborator

illuhad commented Mar 28, 2023

Can you resolve merge conflicts?

This also required removing the `explicit` keyword for the
constructors and adding three new constructors for `unsigned int`,
`unsigned long` and `long`.
We can rely on implicit conversion here, because according to
https://en.cppreference.com/w/cpp/language/static_cast, the static
cast performs the same conversion as an implicit conversion (and
implicit conversions from {integral types, double} -> float are valid
according to https://en.cppreference.com/w/cpp/language/implicit_conversion)
Since we removed the half(fp16::half_storage) constructors, we must
explicitly construct the half that is returned. Unfortunately, the
macro `__hipsycl_backend_switch` contains an if statement for the host
pass in the generic backend, so we cannot just write
`auto data = __hipsycl_backend_switch(...)`.
@illuhad illuhad merged commit 373d311 into AdaptiveCpp:develop Mar 30, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

2 participants