[SYCL] Added support of rounding modes for floating and integer types #1576

fadeeval · 2020-04-23T07:34:04Z

Implementing rounding models for cl::sycl::vec type for non host devices.

Signed-off-by: Aleksander Fadeev aleksander.fadeev@intel.com

bader · 2020-04-23T09:32:58Z

+@Naghasan to review SPIRV built-ins tablegen changes.

Fznamznon · 2020-04-23T14:03:33Z

clang/lib/Sema/SPIRVBuiltins.td

+    if !ne(OutType.ElementSize, InType.ElementSize) then {      
+      def : SPVBuiltin<"SConvert_R" # OutType.Name, [OutType, InType], Attr.Const>;      


Suggested change

if !ne(OutType.ElementSize, InType.ElementSize) then {

def : SPVBuiltin<"SConvert_R" # OutType.Name, [OutType, InType], Attr.Const>;

if !ne(OutType.ElementSize, InType.ElementSize) then {

def : SPVBuiltin<"SConvert_R" # OutType.Name, [OutType, InType], Attr.Const>;

Fznamznon · 2020-04-23T14:07:20Z

sycl/include/CL/sycl/types.hpp

@@ -7,7 +7,7 @@
 //===----------------------------------------------------------------------===//

 // Implements vec and __swizzled_vec__ classes.
-
+#include <typeinfo>


I don't see why you do need this include.

AlexeySachkov

Part about conversions is getting bigger and bigger. We should probably consider outlining it into a separate header file, just to reduce the size of types.hpp and make it easier to read

AlexeySachkov · 2020-04-23T17:31:45Z

sycl/include/CL/sycl/types.hpp

@@ -199,6 +199,34 @@ using is_int_to_int =
    std::integral_constant<bool, std::is_integral<T>::value &&
                                     std::is_integral<R>::value>;

+template <typename T, typename R>
+using is_sint_to_sint = std::integral_constant<
+    bool, std::is_integral<T>::value && !(std::is_unsigned<T>::value) &&


!istd::is_unsigned -> std::is_signed

std::is_signed include floating point types also, is not?

Yes, is_signed includes is_arithmetic, but you already have is_integral here - I guess this should be enough to reject floating-point types, right?

I will use "is_sigeninteger" in new version any way.

AlexeySachkov · 2020-04-23T17:37:27Z

sycl/include/CL/sycl/types.hpp

+template <typename T, typename R>
+using is_sint_to_sint = std::integral_constant<
+    bool, std::is_integral<T>::value && !(std::is_unsigned<T>::value) &&
+              std::is_integral<R>::value && !(std::is_unsigned<R>::value)>;


I think that you can use the following type trait that we have:

llvm/sycl/include/CL/sycl/detail/generic_type_traits.hpp

Line 163 in 04a360a

using is_sigeninteger = is_contained<T, gtl::scalar_signed_integer_list>;

AlexeySachkov · 2020-04-23T17:48:08Z

sycl/include/CL/sycl/types.hpp

+// unsigned to unsigned
+#define __SYCL_GENERATE_CONVERT_IMPL(DestType)                                 \
+  template <typename T, typename R, rounding_mode roundingMode>                \
+  detail::enable_if_t<                                                         \
+      !std::is_same<T, R>::value && is_uint_to_uint<T, R>::value &&            \
+          std::is_same<R, DestType>::value &&                                  \
+          std::is_same<cl::sycl::detail::ConvertToOpenCLType_t<T>, R>::value,  \
+      R>                                                                       \
+  convertImpl(T Value) {                                                       \
+    return static_cast<R>(Value);                                              \
+  }


You have very similar chunk of code for signed to signed. I suggest you to adjust enable_if a bit so it can be re-used for this unsigned to unsigned as well.

Something like

detail::enable_if_t< \ !std::is_same<T, R>::value && (is_sint_to_sint<T, R>::value || is_uint_to_uint<T, R>::value) && \ std::is_same<R, DestType>::value && \ std::is_same<cl::sycl::detail::ConvertToOpenCLType_t<T>, R>::value, \ R>

AlexeySachkov · 2020-04-23T17:51:15Z

sycl/include/CL/sycl/types.hpp

+  detail::enable_if_t<                                                         \
+      !std::is_same<T, R>::value && is_sint_to_sint<T, R>::value &&            \
+          std::is_same<R, DestType>::value &&                                  \
+          std::is_same<cl::sycl::detail::ConvertToOpenCLType_t<T>, R>::value,  \


This looks like a lot of work on templates for the compiler. Can we somehow simplify it?

Is it possible to convert SYCL types to OpenCL types somewhere else in the call stack (let's say in a function which calls convertImpl) and re-design convertImpl so it already operates on OpenCL types?

AlexeySachkov · 2020-04-23T18:06:15Z

sycl/include/CL/sycl/types.hpp

+  detail::enable_if_t<                                                         \
+      is_sint_to_float<T, R>::value && std::is_same<R, DestType>::value, R>    \
+  convertImpl(T Value) {                                                       \
+    using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t<T>;                \
+    OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value);  \
+    return __spirv_Convert##SPIRVOp##_R##DestType(OpValue);                    \
+  }


This std::is_same<R, DestType::value doesn't seem like a right thing to do here. Is it possible to avoid adding this check into each enable_if?

What about creating partial specializations of convertImpl?

template <typename T, typename R, rounding_mode roundingMode> R convertImp(T Value) { // the most generic one // static_cast as fallback? } template <typename T, rounding_mode roundingMode> half convertImpl<T, half, roundingMode>(half Value) { // actual implementation here } // Actual implementations can still be generated by preprocessor: #define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType) \ detail::enable_if_t< is_signed_arithmetic<T>::value, T> \ template <typename T, rounding_mode roundingMode> \ convertImpl<T, DestType, roundingMode>(DestType Value) { \ // actual implementation here \ } __SYCL_GENERATE_CONVERT_IMPL(SToF, half) __SYCL_GENERATE_CONVERT_IMPL(SToF, float) __SYCL_GENERATE_CONVERT_IMPL(SToF, double) #undef __SYCL_GENERATE_CONVERT_IMPL

How about just get name of a type as string and add it to a spirv func name?

I don't understand how to implement your idea

C++ doesn't support function template partial specialization.

AlexeySachkov · 2020-04-23T18:06:49Z

sycl/test/basic_tests/vec_convert.cpp

-  test<float, half, 8, rounding_mode::rte>(
-      float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
-      half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});
+  test<double, float, 8, rounding_mode::automatic>(


Suggested change

test<double, float, 8, rounding_mode::automatic>(

test<double, float, 8, rounding_mode::rte>(

AlexeySachkov · 2020-04-23T18:08:18Z

sycl/test/basic_tests/vec_convert.cpp

-  test<float, half, 8, rounding_mode::automatic>(
-      float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
-      half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});
+  test<double, float, 8, rounding_mode::automatic>(


Since you are adding separate tests for particular data types, I suggest you to just remove this test case from here - it is anyway duplicated in vec_convert_f_to_f.cpp

AlexeySachkov · 2020-04-23T18:11:12Z

sycl/test/basic_tests/vec_convert_f_to_f.cpp

+template <int N>
+struct helper;
+
+template <>
+struct helper<0> {
+  template <typename T, int NumElements>
+  static void compare(const vec<T, NumElements> &x,
+                      const vec<T, NumElements> &y) {
+    const T xs = x.template swizzle<0>();
+    const T ys = y.template swizzle<0>();
+    assert(xs == ys);
+  }
+};
+
+template <int N>
+struct helper {
+  template <typename T, int NumElements>
+  static void compare(const vec<T, NumElements> &x,
+                      const vec<T, NumElements> &y) {
+    const T xs = x.template swizzle<N>();
+    const T ys = y.template swizzle<N>();
+    helper<N - 1>::compare(x, y);
+    assert(xs == ys);
+  }
+};
+


Please outline duplicated code into a header file with helpers

There should I place the header?

AlexeySachkov · 2020-04-23T18:13:27Z

sycl/test/basic_tests/vec_convert.cpp

-      half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});
+  test<double, float, 8, rounding_mode::automatic>(
+      double8{+2.3, +2.5, +2.7, -2.3, -2.5, -2.7, 0., 0.},
+      float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});

  // rte
  test<int, int, 8, rounding_mode::rte>(


Since you are adding separate files for conversions between different types, I suggest you to leave in this file only float to int conversions - just to avoid testing the same combinations in different files

Maybe I should rename this file then?

Maybe I should rename this file then?

Makes sense

Naghasan · 2020-04-24T13:40:03Z

clang/lib/Sema/SPIRVBuiltins.td

@@ -745,7 +745,7 @@ foreach IType = [Char, Short, Int, Long] in {
  }
 }

-foreach InType = TLAll.List in {
+foreach InType = TLUnsignedInts.List in {


Why this restriction ? It is valid to convert a signed int to an unsigned one.

Yes, it is valid, but SatConvertSToU exists for this purpose.

UConvert/SConvert does not saturate the result unless decorated...

I will check it, but I have doubt, that UConvert gets non unsigned arguments.

When I run test that converts int to uint, it writes "call to '__spirv_UConvert_Ruint' is ambiguous", that means it make implicit conversion form signed to unsigned so that push int argument, but compiler have many options for that that's why it writes the error.

I think the document https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html#OpUConvert proves my gusses.

I think the document https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html#OpUConvert proves my gusses.

I hardly see what it proves, from the spec This is either a truncate or a zero extend. whether the destination type is bigger or smaller, this is not a saturating operation: Uconvert 0x0F00 to uchar yields 0x0, SatConvertStoU yields 0xFF.

that converts int to uint, it writes "call to '__spirv_UConvert_Ruint' is ambiguous"

This is expected as this operation does not exist in SPIR-V, so the overload does not exist. You are trying to do a 32 to 32 bits conversion (see https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html#OpUConvert: The component width cannot equal the component width in Result Type.). The operation you are looking for in this cast is simply a reinterpret_cast in that case.

Ok, I will remake

AlexeySachkov · 2020-04-27T08:42:41Z

sycl/include/CL/sycl/types.hpp

@@ -7,7 +7,7 @@
 //===----------------------------------------------------------------------===//

 // Implements vec and __swizzled_vec__ classes.
-#include <typeinfo>
+


Suggested change

AlexeySachkov · 2020-04-27T08:46:39Z

sycl/include/CL/sycl/types.hpp

-#define __SYCL_GENERATE_CONVERT_IMPL(DestType)                                 \
-  template <typename T, typename R, rounding_mode roundingMode>                \
+// convert signed and unsigned types with an equal size and diff names
+#define __SYCL_GENERATE_CONVERT_IMPL()                                 \


Why do you need this as macro? The main idea was to auto-generate some code by using pre-processor features. Now you have no arguments to this macro and call it only once - in that case having a macro just doesn't make sense

fadeeval · 2020-04-29T12:40:31Z

@turinevgeny, @erichkeane, @Fznamznon, @Naghasan, please, make review.

erichkeane

The CFE changes are fine, but RT reviewers are going to need to do approval for the rest.

turinevgeny

types.hpp looks fine.

sycl/include/CL/sycl/types.hpp

sycl/test/basic_tests/vec_convert_half.cpp

sycl/test/basic_tests/vec_convert_f_to_f.cpp

sycl/test/basic_tests/vec_convert_half.cpp

bader · 2020-05-01T16:23:03Z

@fadeeval, please, resolve merge conflict.

bader · 2020-05-10T10:26:08Z

sycl/test/basic_tests/vec_convert_f_to_f.cpp

+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+//==------------ vec_convert.cpp - SYCL vec class convert method test ------==//


vec_convert.cpp -> vec_convert_f_to_f.cpp

Please, fix the comment in a separate PR.

bader · 2020-05-10T10:26:40Z

sycl/test/basic_tests/vec_convert_f_to_i.cpp

+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+//==------------ vec_convert.cpp - SYCL vec class convert method test ------==//


vec_convert.cpp -> vec_convert_f_to_i.cpp

Please, fix the comment in a separate PR.

Naghasan · 2020-05-11T14:36:14Z

sycl/test/basic_tests/vec_convert_f_to_i.cpp

+  test<float, int, 8, rounding_mode::automatic>(
+      float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
+      int8{2, 2, 3, -2, -2, -3, 0, 0});


Suggested change

test<float, int, 8, rounding_mode::automatic>(

float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},

int8{2, 2, 3, -2, -2, -3, 0, 0});

test<float, int, 8, rounding_mode::automatic>(

float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},

int8{2, 2, 2, -2, -2, -2, 0, 0});

RTZ is the default for conversions to integer

I think so too, but in the sycl-1.2.1.pdf spec on the 226 page the automatic description is following: "Default rounding mode for the SYCL vec class element type. rtz (round toward zero) for integer types and rte (round to nearest even) for floating-point types", which sounds confusing. And the person who wrote conversions for HOST implemented so that RTE is automatic mode for converting from floating-point types. That is why I followed the idea that automatic mode is RTE. Should I change it to RTZ?

"Default rounding mode for the SYCL vec class element type. rtz (round toward zero) for integer types and rte (round to nearest even) for floating-point types", which sounds confusing

I agree, RTZ for conversions to integer types could be a better wording for this.

Should I change it to RTZ?

Well that's what the spec mandates. But according to what you are saying, changing this will go well out of scope of your patch. The bug is already there anyway and the patch consistent with it, so I would personally lean toward merging as is and do a PR to fully fix the automatic mode in one go after. But that's more a CO decision now.

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

fadeeval · 2020-05-13T08:12:51Z

@turinevgeny, @erichkeane, @Fznamznon, @Naghasan, @sergey-semenov, make review, please,.

turinevgeny

LGTM in general, but I'd like someone else to approve who knows more details.

fadeeval · 2020-05-14T09:46:03Z

@erichkeane, @Fznamznon, @Naghasan, @sergey-semenov, approve if no objections, please.

Fznamznon · 2020-05-14T09:58:51Z

I don't see any FE changes, please don't wait for approve from me.

Naghasan

Looks ok to me. I think RTZ issue for integers should be handled separately.

sergey-semenov

LGTM aside from the unresolved file header comments by @bader, but those can be addressed separately.

fadeeval requested review from erichkeane, Fznamznon, turinevgeny and a team as code owners April 23, 2020 07:34

fadeeval requested review from sergey-semenov, bader and AlexeySachkov April 23, 2020 07:34

fadeeval force-pushed the private/fadeeval/vec_convert_API_rounding_mode_consideration branch from 67329a5 to 1cd7097 Compare April 23, 2020 08:21

Fznamznon reviewed Apr 23, 2020

View reviewed changes

bader mentioned this pull request Apr 23, 2020

[SPIR-V] Correct/improve declaration of some SPIR-V builtins #1519

Merged

AlexeySachkov reviewed Apr 23, 2020

View reviewed changes

Naghasan requested changes Apr 24, 2020

View reviewed changes

AlexeySachkov reviewed Apr 27, 2020

View reviewed changes

fadeeval requested a review from Fznamznon April 29, 2020 10:56

erichkeane previously approved these changes Apr 29, 2020

View reviewed changes

fadeeval requested a review from romanovvlad April 29, 2020 12:58

turinevgeny previously approved these changes Apr 29, 2020

View reviewed changes

sergey-semenov reviewed Apr 29, 2020

View reviewed changes

sycl/include/CL/sycl/types.hpp Outdated Show resolved Hide resolved

sycl/test/basic_tests/vec_convert_half.cpp Outdated Show resolved Hide resolved

sycl/test/basic_tests/vec_convert_f_to_f.cpp Outdated Show resolved Hide resolved

fadeeval dismissed stale reviews from turinevgeny and erichkeane via da36f2a April 30, 2020 09:41

sergey-semenov reviewed Apr 30, 2020

View reviewed changes

sycl/test/basic_tests/vec_convert_half.cpp Outdated Show resolved Hide resolved

sergey-semenov previously approved these changes Apr 30, 2020

View reviewed changes

bader reviewed May 10, 2020

View reviewed changes

Naghasan requested changes May 11, 2020

View reviewed changes

fadeeval added 3 commits May 12, 2020 11:10

Unsupported on cuda buffer_dev_to_dev.cpp

8b4058d

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

Comments fix

b70cd8b

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

Fix

8d0b71e

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

fadeeval added 13 commits May 12, 2020 12:07

redoing templates

b8f770b

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

Tests reconstruction

aa76104

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

Fix convert to long long type

6d204e0

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

Convert to long long correction

62db4e5

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

Builtins fix

9038cc9

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

generic_type_traits modifying

39c7c19

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

tepes.hpp improving

58c128b

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

types.hpp improvment 2

1ee5c1a

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

Formatting

27b79ef

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

Formatting 2

dbad799

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

FIX

3a3fcbe

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

vec_convert_half.cpp fix

c8484e0

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

tepes.cpp fix

d2c7a22

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

fadeeval dismissed sergey-semenov’s stale review via d2c7a22 May 12, 2020 09:13

fadeeval force-pushed the private/fadeeval/vec_convert_API_rounding_mode_consideration branch from f834594 to d2c7a22 Compare May 12, 2020 09:13

FIX rebase

7bb9767

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

fadeeval requested review from turinevgeny and AlexeySachkov May 12, 2020 11:48

turinevgeny approved these changes May 13, 2020

View reviewed changes

bader requested a review from sergey-semenov May 14, 2020 10:28

Naghasan approved these changes May 14, 2020

View reviewed changes

sergey-semenov approved these changes May 14, 2020

View reviewed changes

bader merged commit 096d0a0 into intel:sycl May 14, 2020

fadeeval mentioned this pull request May 21, 2020

[SYCL] vec convert of long long types correction. #1734

Merged

fadeeval deleted the private/fadeeval/vec_convert_API_rounding_mode_consideration branch May 26, 2020 08:25

dnmokhov mentioned this pull request Jan 26, 2021

[SYCL] Fix long long support in vec::convert on Windows #3097

Merged

		if !ne(OutType.ElementSize, InType.ElementSize) then {
		def : SPVBuiltin<"SConvert_R" # OutType.Name, [OutType, InType], Attr.Const>;

	test<double, float, 8, rounding_mode::automatic>(
	test<double, float, 8, rounding_mode::rte>(

[SYCL] Added support of rounding modes for floating and integer types #1576

[SYCL] Added support of rounding modes for floating and integer types #1576

Uh oh!

Conversation

fadeeval commented Apr 23, 2020

Uh oh!

bader commented Apr 23, 2020

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

AlexeySachkov left a comment

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!

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!

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!

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!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

fadeeval commented Apr 29, 2020

Uh oh!

erichkeane left a comment

Choose a reason for hiding this comment

Uh oh!

turinevgeny left a comment

Choose a reason for hiding this comment

fadeeval May 12, 2020 •

edited

Loading