Skip to content

[SYCL] Add is_native_function_object trait #2412

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

Closed
wants to merge 4 commits into from
Closed

[SYCL] Add is_native_function_object trait #2412

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
8bdc40b
[SYCL] Add is_native_function_object trait
Pennycook Sep 2, 2020
8102815
[SYCL] Use std::integral_constant
Pennycook Sep 2, 2020
666bcbf
[SYCL] Add support for const via std::decay
Pennycook Sep 2, 2020
2037dbb
[SYCL] Fix C++11 include
Pennycook Sep 2, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
22 changes: 22 additions & 0 deletions sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -190,6 +190,28 @@ using EnableIfIsNonNativeOp = cl::sycl::detail::enable_if_t<
!cl::sycl::detail::is_native_op<T, BinaryOperation>::value,
T>;

template <typename T>
struct is_native_function_object_impl : std::false_type {};

template <template <typename> class BinaryOperation, typename T>
struct is_native_function_object_impl<BinaryOperation<T>>
: std::integral_constant<
bool, (sycl::detail::is_scalar_arithmetic<T>::value ||
sycl::detail::is_vector_arithmetic<T>::value ||
std::is_same<T, void>::value) &&
Copy link
Contributor

Choose a reason for hiding this comment

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

In case of T is void (when real type will be auto-deduced on binary_op call) - how an identity will be auto-detected?
In general - it is not possible...

Copy link
Contributor

Choose a reason for hiding this comment

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

I agree that's the interesting question

Copy link
Contributor

Choose a reason for hiding this comment

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

Above John wrote:

A note to reviewers: handling void was interesting here, and highlighted an issue that we might need to consider further. We need is_native_function_object_v to return true for something like sycl::plus<>, but that means the test is not actually sufficient to identify whether a group algorithm is supported -- the caller needs to check both the function type and the data type. Maybe this is already implied by the specification.

Are you referring to the same or something different? If something different what identity do you mean? This is the reason I think we shouldn't test T at all and only test BinaryOperation.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

An identity value isn't always required by the group algorithms. For example, reduce and exclusive_scan all take an init parameter for the cases where the identity value is unknown.

We should introduce a separate type trait for whether a functor has a known identity, and it shouldn't be tied to whether the functor is declared in the standard. @v-klochkov's already has something like this in the reduction implementation; we need to make it official, and provide a mechanism for users to declare the identity values for custom functors.

Copy link
Contributor

Choose a reason for hiding this comment

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

An identity value isn't always required by the group algorithms. For example, reduce and exclusive_scan all take an init parameter for the cases where the identity value is unknown.

John, it seems that we may incorrect result in case of considering an init parameter instead of identity... Or you want to
say the group algorithms don't combine the partial "sums"? ('m not properly aware about the group algorithms. Please, correct me if I mistake.)

Copy link
Contributor Author

Choose a reason for hiding this comment

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

John, it seems that we may incorrect result in case of considering an init parameter instead of identity... Or you want to
say the group algorithms don't combine the partial "sums"? ('m not properly aware about the group algorithms. Please, correct me if I mistake.)

The kernel-wide reduction needs an identity value for optimization purposes, because it's possible for a work-item executing a kernel to contribute nothing to the reduction. In such cases, the fact that the work-item didn't contribute anything needs to be tracked somehow (e.g. by storing a std::optional or an extra bool), or the work-item's contribution can be initialized to a value that doesn't impact the result (i.e. the identity value).

With the most common forms of reduce and exclusive_scan, all work-items contribute a value to the operation, and so that initialization issue goes away. I think the only form that strictly requires a known identity is sycl::exclusive_scan(Group, T, BinaryOperation), because it's not obvious what value to return on the first work-item otherwise.

sycl::detail::is_native_op<T, BinaryOperation<T>>::value> {
};
Comment on lines +198 to +203
Copy link
Contributor

Choose a reason for hiding this comment

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

I mentioned it as a minor in the other thread. I don't think there is any value in checking anything about T. Nor does the spec say anything about it.

Suggested change
: std::integral_constant<
bool, (sycl::detail::is_scalar_arithmetic<T>::value ||
sycl::detail::is_vector_arithmetic<T>::value ||
std::is_same<T, void>::value) &&
sycl::detail::is_native_op<T, BinaryOperation<T>>::value> {
};
: sycl::detail::is_native_op<T, BinaryOperation<T>> {};

Copy link
Contributor Author

Choose a reason for hiding this comment

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

This is much harder than I anticipated. You're right that the specification doesn't say anything about checking T, but maybe it should.

You'd almost convinced me (and I started implementing your suggestion), but then I realized that if SYCL itself doesn't provide a trait that checks the combination of T and BinaryOperation it's much harder for the user to do it themselves.

Consider the below:

is_native_function_object_v<BinaryOperation> && /* T is a SYCL fundamental type */

With your proposed patch, this will return true if BinaryOperation is plus<T2> but the inputs are T. So we'd only be providing part of the logic ("is this function object one that SYCL knows something about?"). Users would have to write individual checks for each algorithm, right?


template <typename T>
struct is_native_function_object
: is_native_function_object_impl<typename std::decay<T>::type> {};

#if __cplusplus >= 201703L
template <typename T>
Copy link
Contributor

Choose a reason for hiding this comment

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

I think we may emulate inline variables even before c++17. The question is should we? If the answer is "Yes, we should" you may ask me how to do that. I suggested that way for one of the compiler features in the past.

To be clear, I am ok to leave it as is. Just wanted to point out that if it's important to have that API prior c++17 there is the way

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I'm leaning towards leaving it as it is. The _v traits in the standard library (e.g. std::is_execution_policy_v) are only enabled for C++17 onwards, and I don't expect people to mix C++17-style traits with earlier standards.

My opinion would be different if there was a precedent for emulating standard SYCL _v traits in our implementation. But I don't think there are any other standard traits implemented yet!

Copy link
Contributor

Choose a reason for hiding this comment

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

Variable templates were added in C++14. It is true that they weren't added to the STL until C++17. But I see no reason why that should impact us. Therefore I think this should be enabled for C++14 (but not C++11).

inline constexpr bool is_native_function_object_v =
is_native_function_object<T>::value;
#endif

template <typename Group> bool all_of(Group, bool pred) {
static_assert(sycl::detail::is_generic_group<Group>::value,
"Group algorithms only support the sycl::group and "
Expand Down
22 changes: 22 additions & 0 deletions sycl/test/group-algorithm/traits.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,22 @@
// RUN: %clangxx -std=c++17 -fsyntax-only -Xclang -verify %s -I %sycl_include -Xclang -verify-ignore-unexpected=note,warning
// expected-no-diagnostics
#include <CL/sycl.hpp>

using namespace sycl::ONEAPI;

struct CustomType {};

template <typename T> struct CustomFunctor {};

int main() {
static_assert(is_native_function_object_v<plus<>>);
static_assert(is_native_function_object_v<multiplies<>>);
static_assert(is_native_function_object_v<bit_or<>>);
static_assert(is_native_function_object_v<bit_xor<>>);
static_assert(is_native_function_object_v<bit_and<>>);
static_assert(is_native_function_object_v<const plus<>>);
static_assert(is_native_function_object_v<plus<float>>);
static_assert(is_native_function_object_v<plus<sycl::vec<float, 4>>>);
static_assert(!is_native_function_object_v<plus<CustomType>>);
static_assert(!is_native_function_object_v<CustomFunctor<float>>);
}