[SYCL] Add has_known_identity/known_identity #2528

Pennycook · 2020-09-23T22:26:30Z

These traits have been requested by developers who would like to
determine at compile-time whether the reduction() interface supports
their use-case, or whether they need to fall back to some other
implementation.

Exposes existing functionality of reduction implementation as part of
the public API for reductions.

Signed-off-by: John Pennycook john.pennycook@intel.com

These traits have been requested by developers who would like to determine at compile-time whether the reduction() interface supports their use-case, or whether they need to fall back to some other implementation. Signed-off-by: John Pennycook <john.pennycook@intel.com>

Exposes existing functionality of reduction implementation as part of the public API for reductions. Signed-off-by: John Pennycook <john.pennycook@intel.com>

sycl/doc/extensions/Reduction/Reduction.md

Signed-off-by: John Pennycook <john.pennycook@intel.com>

Pennycook · 2020-09-24T16:13:45Z

@Fznamznon, @AlexeySachkov: Could you please take a look at the build failures here and help identify a path forwards?

It looks to me like reduction_nd_ext_half.cpp is failing now because the calculation of negative infinity isn't constexpr for half, but I'm not sure how to fix it. Is there an alternative way to write the below that is constexpr for half, or do we need to add some new constexpr functions to the host implementation of half?

  static constexpr AccumulatorT value =
      std::numeric_limits<AccumulatorT>::has_infinity
          ? static_cast<AccumulatorT>(
                -std::numeric_limits<AccumulatorT>::infinity())
          : std::numeric_limits<AccumulatorT>::lowest();

I think unary - triggers a conversion to float which isn't constexpr. I tried making operator float() constexpr, but there's no host definition of this function that I could see.

v-klochkov

Looks good to me, I found only 1 typo.

sycl/doc/extensions/Reduction/Reduction.md

v-klochkov · 2020-09-24T20:01:11Z

sycl/include/CL/sycl/ONEAPI/reduction.hpp

+                           typename std::enable_if<IsMinimumIdentityOp<
+                               AccumulatorT, BinaryOperation>::value>::type> {


Minor comment, definitely not a request to change.
This could be a little bit shorter (but same number of lines though):

Suggested change

typename std::enable_if<IsMinimumIdentityOp<

AccumulatorT, BinaryOperation>::value>::type> {

enable_if_t<IsMinimumIdentityOp<

AccumulatorT, BinaryOperation>::value>> {

I tried this during development and couldn't get it to work. I chalked it up to me not understanding enough about how our implementation of enable_if_t differs from std::enable_if. Can you compile locally if you make this change? Maybe I'm doing something wrong.

I checked our enable_if_t implementation. It is the same as in C++ standard. So, there should not be any problem. If you need help we may have a call or IM

Signed-off-by: John Pennycook <john.pennycook@intel.com>

rarutyun · 2020-09-29T10:56:04Z

sycl/doc/extensions/Reduction/Reduction.md

+If an implementation can identify the identity value for a given combination of accumulator type `AccumulatorT` and function object type `BinaryOperation`, the value is defined as a member of the `known_identity` trait class:
+```c++
+template <typename BinaryOperation, typename AccumulatorT>
+struct known_identity {


I am not sure we should telling that known_identity is a struct on the spec level. It may be implemented as the using to something. I would also try to align the API with integral_constant. I understand that this is not possible for 100% but traits with values usually have some run-time API as well. The comment is applicable to other traits

All of the traits in C++ are defined as a struct, aren't they? Is there a good reason that known_identity should be different?

I'm not sure what you're imagining when aligning with integral_constant. In the general case, I don't think we can guarantee that the identity value is something that can be passed as a template parameter (e.g. it may be a floating-point type).

sycl/include/CL/sycl/ONEAPI/reduction.hpp

rarutyun · 2020-09-29T15:01:44Z

sycl/include/CL/sycl/ONEAPI/reduction.hpp

+                           typename std::enable_if<IsMaximumIdentityOp<
+                               AccumulatorT, BinaryOperation>::value>::type> {
+  static constexpr AccumulatorT value =
+      std::numeric_limits<AccumulatorT>::has_infinity


It seems like AccumulatorT is the the fundamental C++ type only (because it's used in std::numeric_limits).

Two questions basing on that:

What about custom types?

If only fundamental types should be supported, can we introduce the integral_constant semantics for known_identity API as I suggested for has_known_identity?

This is a short-term fix. Our reduction implementation currently only detects identity values automatically for fundamental C++ types (and half), so that's all we've decided to cover with the traits for now.

Eventually we want developers to be able to declare the identity value for their own types and function object types. The current thinking is that for combinations of fundamental types and known function object types (e.g. std::plus<>) the traits will continue to work as implemented here, but for anything else they'll check the function object for something like an identity member.

There are still some open questions about exactly what this should look like for transparent functors and what additional checks are needed (e.g. if we can determine an identity for T1 and T1 is convertible to T2, is it safe to assume that the identity for T1 can be used?). These are important questions, but I don't think the first implementation of the traits should be blocked by them.

rarutyun · 2020-09-29T15:13:42Z

sycl/include/CL/sycl/ONEAPI/reduction.hpp

+                           typename std::enable_if<IsMinimumIdentityOp<
+                               AccumulatorT, BinaryOperation>::value>::type> {


I checked our enable_if_t implementation. It is the same as in C++ standard. So, there should not be any problem. If you need help we may have a call or IM

Signed-off-by: John Pennycook <john.pennycook@intel.com>

rolandschulz · 2020-10-02T21:32:20Z

sycl/include/CL/sycl/half_type.hpp

  uint16_t Buf;
+  friend std::numeric_limits<class cl::sycl::detail::half_impl::half>;


I didn't realize all is fwd-declared and therefore

Suggested change

friend std::numeric_limits<class cl::sycl::detail::half_impl::half>;

friend std::numeric_limits<half_impl::half>;

or

Suggested change

friend std::numeric_limits<class cl::sycl::detail::half_impl::half>;

friend std::numeric_limits<sycl::half>;

works too. Do we prefer one of these shorter options?

I have no preference.

Pennycook · 2020-10-14T14:20:01Z

@rolandschulz: Are you still working on fixing the broken cases here?

bader · 2020-10-14T14:23:35Z

We also need resolve conflicts with #2624.

rolandschulz · 2021-01-12T02:06:19Z

sycl/include/CL/sycl/half_type.hpp

@@ -26,6 +26,11 @@
 #else
 #define __SYCL_CONSTEXPR_ON_DEVICE
 #endif
+#if __cplusplus >= 201402L


Why is there the sycl-include-gnu11.cpp test which checks it compiles with C++11? I thought we only support compiling as >=C++14.

I think the RT library is C++14, but the headers are C++11.

@bader , @romanovvlad - Can't we assume C++14 or newer in SYCL header files?

I don't have a simple "yes/no" answer, but I'm aware of a couple of factors restricting usage of C++ features:

SYCL specification: SYCL-1.2.1 relies on C++11 and SYCL-2020 relies on C++17 features. It seems to me that DPC++ is not going to support "strict SYCL-1.2.1" mode where C++14 (or later) features are not supported.

Runtime library ABI compatibility (see the note in our contribution guide). It's hard to say if using C++14 features in SYCL headers breaks runtime ABI, but it possible and should be prevented by regression tests.

We have C++ version documented at:
https://github.com/intel/llvm/blob/sycl/sycl/doc/GetStartedGuide.md#c-standard
but it doesn't answer whether C++11 mode will be functional or broken.

I think it would be good to know and catch the case when we really want to break C++11 compatibility and have some discussion on that. Maybe that's what sycl-include-gnu11.cpp test is doing?

The test only checks that the header can be included. But not that any template in the header can be instantiated in C++11 mode. The current version of this PR passes all tests . But trying to instantiate known_identity for half in C++11 will give a compiler error (not tested).

It shouldn't have any ABI issues to use any C++14 features in the header. The problematic release for ABI issues is C++11 (e.g. string). There weren't any such things in C++14.

C++14 is required now. The doc changes have been merged 10 days ago: #3053

Pennycook · 2021-01-12T15:04:12Z

Thanks for the fixes, @rolandschulz. Opening up to review now.

* Add note that SYCL 2020 implementation is in progress (see issue intel#2971) * Add note that DPC++ runtime and headers require C++14 (see PR intel#2528) Signed-off-by: Pavel V Chupin <pavel.v.chupin@intel.com>

* Add note that SYCL 2020 implementation is in progress (see issue #2971) * Add note that DPC++ runtime and headers require C++14 (see PR #2528) Signed-off-by: Pavel V Chupin <pavel.v.chupin@intel.com>

v-klochkov

Looks good to me. I suggest merging this patch as is now.

A bit later I'll fix the c++11 LIT test (the fix in that LIT test is included into #3123) and also will remove the macro _CPP14_CONSTEXPR .

@Pennycook please approve the fix if you agree with the changes and the idea to merge the fix in the current state.

AlexeySachkov

half_type.hpp changes look good to me, one minor code style comment

sycl/include/CL/sycl/half_type.hpp

Pennycook · 2021-02-03T16:08:03Z

Looks good to me. I suggest merging this patch as is now.

A bit later I'll fix the c++11 LIT test (the fix in that LIT test is included into #3123) and also will remove the macro _CPP14_CONSTEXPR .

@Pennycook please approve the fix if you agree with the changes and the idea to merge the fix in the current state.

Apparently GitHub doesn't let people approve their own PRs, but I agree with your proposal. Thanks for picking up the required LIT test and macro changes.

…same (#2528) Original commit: KhronosGroup/SPIRV-LLVM-Translator@4a6a1746c61f917

Pennycook added 2 commits September 23, 2020 17:41

[SYCL] Implement has_known_identity/known_identity

be643d0

Exposes existing functionality of reduction implementation as part of the public API for reductions. Signed-off-by: John Pennycook <john.pennycook@intel.com>

Pennycook added enhancement New feature or request spec extension All issues/PRs related to extensions specifications labels Sep 23, 2020

Pennycook requested review from rolandschulz, rarutyun and v-klochkov September 23, 2020 22:26

keryell reviewed Sep 24, 2020

View reviewed changes

sycl/doc/extensions/Reduction/Reduction.md Outdated Show resolved Hide resolved

rolandschulz reviewed Sep 24, 2020

View reviewed changes

sycl/doc/extensions/Reduction/Reduction.md Outdated Show resolved Hide resolved

sycl/doc/extensions/Reduction/Reduction.md Outdated Show resolved Hide resolved

[SYCL][Doc] Fix typo: T => AccumulatorT

c2ff9dc

Signed-off-by: John Pennycook <john.pennycook@intel.com>

v-klochkov reviewed Sep 24, 2020

View reviewed changes

Pennycook added 2 commits September 24, 2020 16:41

[SYCL][Doc] Fix typo again: T => AccumulatorT

aa2222a

Signed-off-by: John Pennycook <john.pennycook@intel.com>

[SYCL][Doc] Define has in terms of known

0136df2

Signed-off-by: John Pennycook <john.pennycook@intel.com>

rarutyun reviewed Sep 29, 2020

View reviewed changes

Pennycook and others added 2 commits September 30, 2020 10:24

[SYCL] Use std::integral_constant

6ca8965

Signed-off-by: John Pennycook <john.pennycook@intel.com>

Add half inf/-inf support

b791034

rolandschulz reviewed Oct 2, 2020

View reviewed changes

rolandschulz added 4 commits January 11, 2021 12:27

Merge remote-tracking branch 'origin/sycl' into HEAD

3717a98

Add 0/1 half known_idenitfy for host

d5b3766

Apply work-around for half only to half

bb3c4c2

Make compile with C++11

93d9499

rolandschulz reviewed Jan 12, 2021

View reviewed changes

Pennycook marked this pull request as ready for review January 12, 2021 15:03

Pennycook requested review from AlexeySachkov and a team as code owners January 12, 2021 15:03

pvchupin mentioned this pull request Jan 18, 2021

[SYCL][Doc] Update for get started guide #3053

Merged

v-klochkov approved these changes Feb 3, 2021

View reviewed changes

AlexeySachkov approved these changes Feb 3, 2021

View reviewed changes

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

Merge remote-tracking branch 'intel_llvm/sycl' into reduction-identity

857bfd8

v-klochkov merged commit 0c7bd24 into intel:sycl Feb 3, 2021

Pennycook deleted the reduction-identity branch February 3, 2021 19:06

v-klochkov mentioned this pull request Feb 8, 2021

[SYCL][NFC] Remove some work-arounds for C++11 #3168

Merged

jsji pushed a commit that referenced this pull request Jun 27, 2024

True Label and False Label of an OpBranchConditional must not be the …

061de13

…same (#2528) Original commit: KhronosGroup/SPIRV-LLVM-Translator@4a6a1746c61f917

		typename std::enable_if<IsMinimumIdentityOp<
		AccumulatorT, BinaryOperation>::value>::type> {

		uint16_t Buf;
		friend std::numeric_limits<class cl::sycl::detail::half_impl::half>;

	friend std::numeric_limits<class cl::sycl::detail::half_impl::half>;
	friend std::numeric_limits<half_impl::half>;

	friend std::numeric_limits<class cl::sycl::detail::half_impl::half>;
	friend std::numeric_limits<sycl::half>;

[SYCL] Add has_known_identity/known_identity #2528

[SYCL] Add has_known_identity/known_identity #2528

Uh oh!

Conversation

Pennycook commented Sep 23, 2020

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Pennycook commented Sep 24, 2020

Uh oh!

v-klochkov left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Pennycook commented Oct 14, 2020

Uh oh!

bader commented Oct 14, 2020

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Pennycook commented Jan 12, 2021

Uh oh!

v-klochkov left a comment

Choose a reason for hiding this comment

Uh oh!

AlexeySachkov left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Pennycook commented Feb 3, 2021

Uh oh!

Uh oh!