-
Notifications
You must be signed in to change notification settings - Fork 769
[SYCL] Add complex support to group algorithms #5394
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] Add complex support to group algorithms #5394
Conversation
…to sycl::plus binary op Signed-off-by: Chris Perkins <chris.perkins@intel.com>
Signed-off-by: Chris Perkins <chris.perkins@intel.com>
Signed-off-by: Chris Perkins <chris.perkins@intel.com>
Signed-off-by: Chris Perkins <chris.perkins@intel.com>
… code duplication. Removed duplicates with a small bit of SFINAE instead. Signed-off-by: Chris Perkins <chris.perkins@intel.com>
Signed-off-by: Chris Perkins <chris.perkins@intel.com>
/verify intel/llvm-test-suite#767 |
Signed-off-by: Chris Perkins <chris.perkins@intel.com>
// ---- identity_for_ga_op | ||
// the group algorithms support std::complex, limited to sycl::plus operation | ||
// get the correct identity for group algorithm operation. | ||
template <typename T, class BinaryOperation> | ||
constexpr detail::enable_if_t< | ||
(is_complex<T>::value && | ||
std::is_same<BinaryOperation, sycl::plus<T>>::value), | ||
T> | ||
identity_for_ga_op() { | ||
return {0, 0}; | ||
} |
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.
We should consider whether it makes sense to extend sycl::known_identity_v
to std::complex
types. I'd like to hear from @v-klochkov -- how hard would it be to do this?
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.
@Pennycook the reason I didn't extend known_identity was that I didn't want to accidentally introduce complex support to a bunch of places where it might not be intended (all the callers of known_identity), plus this particular identity of {0,0} only works because the only operation we are worried about is plus. But that wouldn't be correct for the more general known_identity.
That said, I guess it's worth asking.
Signed-off-by: Chris Perkins <chris.perkins@intel.com>
Signed-off-by: Chris Perkins <chris.perkins@intel.com>
/verify intel/llvm-test-suite#767 |
@@ -257,7 +334,7 @@ joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) { | |||
std::is_same<decltype(binary_op(init, *first)), float>::value), | |||
"Result type of binary_op must match reduction accumulation type."); | |||
#ifdef __SYCL_DEVICE_ONLY__ | |||
T partial = sycl::known_identity_v<BinaryOperation, T>; | |||
T partial = detail::identity_for_ga_op<T, BinaryOperation>(); |
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.
isn't it better to update sycl::known_identity to let it accept std::complex? Otherwise, there may be confusion when known_identity states it is unknown but functionality is supported.
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.
We are asking this question of @v-klochkov above. My feeling is that known_identity is used by the reductions and other parts of SYCL where std::complex is not supported. Adding complex there may not be wise. Furthermore, the identity for complex numbers is {0,0} only in the case of addition - which is the only binary operation the group algorithms support. But known_identity accepts all the binary operations - so having unsupported ones may change its API in ways that confuse it or needlessly increase its test surface.
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.
My feeling is that known_identity is used by the reductions and other parts of SYCL where std::complex is not supported. Adding complex there may not be wise
This is true, but if it's simple to add std::complex
support to has_known_identity
then it might be worth us revising the extension specification to say that we're also adding support for reduction
with std::complex
and plus<>
.
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 am ok to NOT extend known_identity with the complex type. but it is worth commenting with motivation on the definition of identity_for_ga_op.
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 was hoping v-klochkov would be able to comment on this, but he is not going to be available to comment any time soon. We should not wait on him.
I have expanded the comment on identity_for_ga_op
and left a //TODO that we should eliminate it once complex number are supported by the other reductions.
@vladimirlaz and @Pennycook would this be enough for now?
@Pennycook do you want to open a ticket for complex support for the reductions? (I think they are the primary other caller of known_identity right now).
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 agree that sycl::has_known_identity should return true for std::plusstd::complex.
It allows to enable std::complex+std::plus for reductions for free.
Signed-off-by: Chris Perkins <chris.perkins@intel.com>
Signed-off-by: Chris Perkins <chris.perkins@intel.com>
/verify with intel/llvm-test-suite#767 |
This PR proposes to augment sycl::has_known_identity to return true for std::complex and std::plus operator. This will have two benefits: 1. It enables support for complex numbers in sycl::reduction without the user having to explicitly pass identity. 2. sycl::known_identity can now be used to simplify the implementation of group algorithms (See PR #5394). Also, this PR addresses the Github issue #5477. Test Case PR: intel/llvm-test-suite#1609
Many of the group algorithms already support std::complex data types, but not all. There is a new extension that defines their expanded use in the reduce and scan group algorithms. ( here ). This PR implements that new extension.
Tests have been updated here: intel/llvm-test-suite#767