[SYCL] Add is_native_function_object trait

[Pennycook]

These traits were mentioned in the GroupAlgorithms extension and the SYCL 2020
provisional specification, but weren't implemented anywhere.

Signed-off-by: John Pennycook [email protected]

[Pennycook]

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.

[rarutyun]

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

[Pennycook]

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!

[rolandschulz]

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).

[rarutyun]

Overall question. Does it work with std::plus and other similar Callables?

[Pennycook]

Overall question. Does it work with std::plus and other similar Callables?

For the time being and talking only about our implementation -- yes, because we currently define sycl::plus as an alias of std::plus. SYCL 2020 provisional doesn't explicitly say what should happen for the std:: functors, and the specification likely needs some clarification there.

[rolandschulz]

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).

[rolandschulz]

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>> {};

[Pennycook]

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?

[MikeDvorskiy]

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...

[rarutyun]

I agree that's the interesting question

[rolandschulz]

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.

[Pennycook]

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.

[MikeDvorskiy]

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.)

[Pennycook]

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.

[Pennycook]

Closing in favor of #2462.