[ESIMD] Add infra to support half, bfloat, etc, support sycl::half. #5123

kbobrovs · 2021-12-11T00:45:22Z

implement infrastructure for non-standard element type support
to support a new type, the following must be implemented:
- element type traits for the new type
- scalar and vector conversions to/from selected std C++ type
- std operations (or default with promotion to std C++ type can be used)
esimd::simd<sycl::half, N> can now be used on host and device
- most of the operations, except extended math are supported for half

Complementary E2E tests update - intel/llvm-test-suite#640

sndmitriev

Is there any document which describes the concept of supporting non-standard types which you have implemented here? If not you should probably create one. That will greatly simplify things for maintainers of this code in future:)

sndmitriev · 2021-12-23T06:35:19Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/type_format.hpp

+namespace sycl {
+namespace ext {
+namespace intel {
+namespace experimental {
+namespace esimd {
+namespace detail {


nit: these namespaces can be collapsed.

sndmitriev · 2021-12-23T06:37:26Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/types.hpp

+// template <typename To, typename From, int N>
+// vector_type_t<To, N> convert(vector_type_t<From, N> Val) {
+//   return convert<vector_type_t<To, N>>(Val);
+// }


Looks like this commented code can be removed.

sndmitriev · 2021-12-23T06:39:21Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/elem_type_traits.hpp

+namespace sycl {
+namespace ext {
+namespace intel {
+namespace experimental {
+namespace esimd {
+namespace detail {


nit: these namespaces can be collapsed

sndmitriev · 2021-12-23T06:44:37Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/elem_type_traits.hpp

+namespace esimd {
+namespace detail {
+
+enum class BinOp {


Looks like you forgot to add corresponding infra for unary ops.

right, will add

sycl/include/sycl/ext/intel/experimental/esimd/simd.hpp

sndmitriev · 2021-12-23T06:54:54Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/elem_type_traits.hpp

+#ifdef __SYCL_DEVICE_ONLY__
+using half = _Float16;
+#else
+using half = uint16_t;
+#endif // __SYCL_DEVICE_ONLY__


Looks like these definitions are not used and can be removed.

sndmitriev · 2021-12-23T06:56:04Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/elem_type_traits.hpp

+ESIMD_INLINE DstWrapperTy convert_scalar(SrcWrapperTy Val) {
+  if constexpr (std::is_same_v<SrcWrapperTy, DstWrapperTy>) {
+    return Val;
+  } else if constexpr (!detail::is_wrapper_elem_type_v<SrcWrapperTy> &&


detail:: can be removed here.

sndmitriev · 2021-12-23T06:56:31Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/elem_type_traits.hpp

+  } else if constexpr (!detail::is_wrapper_elem_type_v<SrcWrapperTy> &&
+                       !detail::is_wrapper_elem_type_v<DstWrapperTy>) {


detail:: can be removed.

sndmitriev · 2021-12-23T06:59:07Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/operators.hpp

-
+template <class T> using __st = element_storage_t<T>;
+
+#if 0


Looks like this code under #if 0 is not needed anymore and thus need to be removed.

right, thanks for the catch

sndmitriev · 2021-12-23T07:29:16Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/elem_type_traits.hpp

+template <class T, class SFINAE> struct element_type_traits {
+  // The raw element type of the underlying clang vector used as a
+  // storage.
+  using StorageT = invalid_storage_element_type;


Can probably rename StorageT to shorter RawT as you have used word 'raw' in many other places to describe storage type.

kbobrovs · 2021-12-24T04:30:25Z

Is there any document which describes the concept of supporting non-standard types which you have implemented here? If not you should probably create one. That will greatly simplify things for maintainers of this code in future:)

ok, will do

kbobrovs · 2021-12-28T05:17:31Z

@sndmitriev, all comments are fixed - please review

kbobrovs · 2021-12-28T05:18:22Z

@AlexeySachkov, please review sycl/include/CL/sycl/half_type.hpp

kbobrovs · 2021-12-28T05:29:08Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/type_format.hpp

+// Meta-functions to compute compile-time element type of a simd_view resulting
+// from format operations.
+//===----------------------------------------------------------------------===//
+


Content of this file was moved from types.hpp, mostly verbatim. The changes are noted below.

kbobrovs · 2021-12-28T05:31:00Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/type_format.hpp

+          template <typename, int> class SimdT>
+struct compute_format_type<SimdT<Ty, N>, EltTy>
+    : compute_format_type_impl<Ty, N, EltTy> {};
+


specialization for <simd_obj_impl<Ty, N, SimdT>, EltTy> was removed - not supposed to be used as argument

kbobrovs · 2021-12-28T05:31:22Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/type_format.hpp

+          template <typename, int> class SimdT>
+struct compute_format_type_2d<SimdT<Ty, N>, EltTy, Height, Width>
+    : compute_format_type_2d_impl<Ty, N, EltTy, Height, Width> {};
+


specialization for <simd_obj_impl<Ty, N, SimdT>, EltTy> was removed - not supposed to be used as argument

kbobrovs · 2021-12-28T05:32:39Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/types.hpp

-using simd_mask_storage_t = vector_type_t<simd_mask_elem_type, N>;
+// @{
+// Checks if given type T derives from simd_obj_impl or is equal to it.
+template <typename T>


moved around the file due to factoring out some parts of types.hpp

kbobrovs · 2021-12-28T05:33:51Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/types.hpp

+
+// must match simd_mask<N>::element_type
+template <int N>
+using simd_mask_storage_t = vector_type_t<simd_mask_elem_type, N>;


No functional changes in the code above (moved here from different part and renamed template parameter)

kbobrovs · 2021-12-28T05:34:42Z

sycl/include/sycl/ext/intel/experimental/esimd/detail/types.hpp

-struct computation_type<
-    T1, T2, std::enable_if_t<is_vectorizable_v<T1> && is_vectorizable_v<T2>>> {
-  using type = decltype(std::declval<T1>() + std::declval<T2>());
+// Determine element type of simd_obj_impl's Derived type w/o having to have


314-331 is new code

kbobrovs · 2021-12-28T07:31:05Z

Unrelated failures from SYCL / Default Linux / OCL x64 Test Suite (pull_request):
SYCL :: XPTI/buffer/sub_buffer.cpp

terminate called after throwing an instance of 'cl::sycl::invalid_object_error'
what(): Specified offset of the sub-buffer being constructed is not a multiple of the memory base address alignment -30 (CL_INVALID_VALUE)

Unrelated failures from SYCL / Default Linux / HIP AMD GPU Test Suite (pull_request):
SYCL :: Printf/char.cpp
SYCL :: Printf/float.cpp
SYCL :: Printf/int.cpp
SYCL :: Printf/long.cpp
SYCL :: Printf/mixed-address-space.cpp
SYCL :: Printf/percent-symbol.cpp

all like this:

/__w/llvm/llvm/toolchain/bin/../include/sycl/ext/oneapi/experimental/builtins.hpp:71:36: error: SYCL kernel cannot call a variadic function
return ::printf(__format, args...);

sndmitriev

Looks good.

bader · 2021-12-28T08:01:54Z

Unrelated failures from SYCL / Default Linux / OCL x64 Test Suite (pull_request): SYCL :: XPTI/buffer/sub_buffer.cpp

terminate called after throwing an instance of 'cl::sycl::invalid_object_error'
what(): Specified offset of the sub-buffer being constructed is not a multiple of the memory base address alignment -30 (CL_INVALID_VALUE)

Unrelated failures from SYCL / Default Linux / HIP AMD GPU Test Suite (pull_request): SYCL :: Printf/char.cpp SYCL :: Printf/float.cpp SYCL :: Printf/int.cpp SYCL :: Printf/long.cpp SYCL :: Printf/mixed-address-space.cpp SYCL :: Printf/percent-symbol.cpp

all like this:

/__w/llvm/llvm/toolchain/bin/../include/sycl/ext/oneapi/experimental/builtins.hpp:71:36: error: SYCL kernel cannot call a variadic function
return ::printf(__format, args...);

printf issues are fixed by intel/llvm-test-suite#685

I think SYCL / Default Linux / OCL x64 Test Suite (pull_request): SYCL :: XPTI/buffer/sub_buffer.cpp fails because your branch doesn't have 8f9d0d2. Please, merge sycl branch to your private branch.
@vladimirlaz, am I right?

kbobrovs · 2021-12-28T08:03:12Z

printf issues are fixed by intel/llvm-test-suite#685

what should I do then - ignore the failures?

I think SYCL / Default Linux / OCL x64 Test Suite (pull_request): SYCL :: XPTI/buffer/sub_buffer.cpp fails because your branch doesn't have 8f9d0d2. Please, merge sycl branch to your private branch.

ok, thanks

- implement infrastructure for non-standard element type support to support a new type, the following must be implemented: * element type traits for the new type * scalar and vector conversions to/from selected std C++ type * std operations (or default with promotion to std C++ type can be used) - esimd::simd<sycl::half, N> can now be used on host and device * most of the operations, except extended math are supported for half Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>

- fix vector conversion not to use CommonT - should result in fewer conversions and more in accordance to C++ - computation_type_t will now return sycl::half instead of _Float16 - Fix bug in simd(const SimdT &RHS) constructor - atomics are checked based on user type, not raw - EnclosingCppT for sycl::half is now _Float16 in device compiler - Few review comments addressed. Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>

Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>

- Add description of the non-std type support infra - Support unary ops in the infra - rename storage_type to raw_type - remove user_element_type - use element_type and raw_element_type instead of element_type and user_element_type Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>

Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>

bader · 2021-12-28T08:05:51Z

what should I do then - ignore the failures?

When you merge with the sycl branch, tests will be ran again and these tests should be skipped.

AlexeySachkov

half_type.hpp changes LGTM

kbobrovs · 2021-12-28T08:18:56Z

@dm-vodopyanov, @s-kanaev, @againull, @v-klochkov, @smaslov-intel - could you folks please review the only non-ESIMD change in the patch - sycl/include/CL/sycl/half_type.hpp - on behalf of llvm-reviewers-runtime?

kbobrovs · 2021-12-28T08:30:42Z

had to rebase and force-push to fix test failures (no code changes)

kbobrovs force-pushed the half branch from bb95580 to 41b2a74 Compare December 20, 2021 07:15

kbobrovs mentioned this pull request Dec 20, 2021

[ESIMD] Extend tests to cover sycl::half. intel/llvm-test-suite#640

Merged

kbobrovs changed the title ~~[WIP][ESIMD] Support sycl::half in device code.~~ [ESIMD] Add infra for half, bfloat, etc. support, support sycl::half. Dec 20, 2021

kbobrovs changed the title ~~[ESIMD] Add infra for half, bfloat, etc. support, support sycl::half.~~ [ESIMD] Add infra to support half, bfloat, etc, support sycl::half. Dec 20, 2021

kbobrovs force-pushed the half branch from 41b2a74 to 18b4b3b Compare December 23, 2021 03:58

kbobrovs marked this pull request as ready for review December 23, 2021 03:59

kbobrovs requested review from AlexeySachkov, kychendev and v-klochkov as code owners December 23, 2021 03:59

kbobrovs requested a review from sndmitriev December 23, 2021 03:59

sndmitriev reviewed Dec 23, 2021

View reviewed changes

kbobrovs force-pushed the half branch from c8ba99c to 6d1d7ea Compare December 27, 2021 07:42

kbobrovs requested a review from a team as a code owner December 27, 2021 23:48

kbobrovs requested a review from sndmitriev December 28, 2021 05:17

kbobrovs commented Dec 28, 2021

View reviewed changes

sndmitriev previously approved these changes Dec 28, 2021

View reviewed changes

kbobrovs added 3 commits December 28, 2021 00:05

Address review comments.

c2e1c75

Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>

kbobrovs added 6 commits December 28, 2021 00:05

clang format

c07e6ac

Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>

mark some APIs as ESIMD_INLINE to avoid symbol duplication linker error.

bd7ea4d

Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>

Fixed bug in 1-elem simd_view element type definition.

d189cdf

Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>

clang format

7f45e13

Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>

Fix test failures

0c9fe3e

Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>

AlexeySachkov previously approved these changes Dec 28, 2021

View reviewed changes

kbobrovs dismissed stale reviews from AlexeySachkov and sndmitriev via 0c9fe3e December 28, 2021 08:23

kbobrovs force-pushed the half branch from a0c7b65 to 0c9fe3e Compare December 28, 2021 08:23

sndmitriev approved these changes Dec 28, 2021

View reviewed changes

againull approved these changes Dec 28, 2021

View reviewed changes

againull merged commit f34ba2c into intel:sycl Dec 28, 2021

kbobrovs deleted the half branch December 28, 2021 20:44

		} else if constexpr (!detail::is_wrapper_elem_type_v<SrcWrapperTy> &&
		!detail::is_wrapper_elem_type_v<DstWrapperTy>) {

[ESIMD] Add infra to support half, bfloat, etc, support sycl::half. #5123

[ESIMD] Add infra to support half, bfloat, etc, support sycl::half. #5123

Conversation

kbobrovs commented Dec 11, 2021 • edited

sndmitriev left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

kbobrovs commented Dec 24, 2021

kbobrovs commented Dec 28, 2021

kbobrovs commented Dec 28, 2021

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

kbobrovs commented Dec 28, 2021

sndmitriev left a comment

Choose a reason for hiding this comment

bader commented Dec 28, 2021

kbobrovs commented Dec 28, 2021

bader commented Dec 28, 2021

AlexeySachkov left a comment

Choose a reason for hiding this comment

kbobrovs commented Dec 28, 2021

kbobrovs commented Dec 28, 2021

kbobrovs commented Dec 11, 2021 •

edited