[SYCL][ESIMD][EMU] Enable dpas with ESIMD_EMULATOR backend #6475

dongkyunahn-intel · 2022-07-25T21:18:26Z

+fix for 'saturate<>' for ESIMD_EMULATOR backend

- +fix for 'saturate<>' for ESIMD_EMULATOR backend

dongkyunahn-intel · 2022-07-25T21:19:56Z

Note : DPASW support will be enabled later as it requires updates in esimd_emulator backend interface as well as CM_EMU library for runtime support

The patch caused compilation fail at esimd/dpas.cpp. Need additional changes.

- First argument for '__esimd_dpas_inner' is changed from 'const __ESIMD_DNS::vector_type_t<RT, SZ>' to 'const void *' - This is because 'std::addressof()' from caller generates 'short * __attribute__((ext_vector_type(16)))' that cannot be converted into 'vector_type_t<>' while 'dpas.cpp' test is compiled.

- As template and runtime arguments are different (e.g. 'dpas_info' is not used in ESIMD_EMULATOR backend (hostmode))

pvchupin · 2022-07-27T21:58:51Z

Thanks @dongkyunahn-intel, do you know anything about SYCL :: unused_spec_const.cpp ?

dongkyunahn-intel · 2022-07-27T22:06:30Z

Thanks @dongkyunahn-intel, do you know anything about SYCL :: unused_spec_const.cpp ?

@pvchupin , as far as I know 'spec_const' tests are disabled for ESIMD_EMULATOR backend as they are failing with Unsupported 'online_compiler' error. It could be because of 'AOT' that is not supported by the ESIMD_EMULATOR backend. Currently, there is no plan for enabling those tests with the backend.

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

- Template arguments for dpas operations (repeat count, systolic depth, src1_precision, src2_precision) are converted into runtime arguments - __esimd_dpas* functions have same signature regardless of __SYCL_DEVICE_ONLY__

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

dongkyunahn-intel · 2022-07-30T01:51:10Z

Unrelated failure :

SYCL / Linux / CUDA LLVM Test Suite

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

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

dongkyunahn-intel · 2022-08-08T18:38:01Z

Function signatures for __esimd_dpas* will be changed to have precision, systolic depth, and repeat count (from dpas_info) as template arguments. Host implementation will be changed in this PR. Device implementation will be changed later.

- For precision, repeat, depth, and signedness info, runtime function arguments are changed to template argument for host mode compilation - dpas_info encoding/decoding, ARG_UNUSED removed as they are no longer used - TODO : Refactor intrinsic declarations for device mode

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

kbobrovs · 2022-08-12T17:18:22Z

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

-  __ESIMD_UNSUPPORTED_ON_HOST;
-  return __ESIMD_DNS::vector_type_t<T, N>();
-#endif // __SYCL_EXPLICIT_SIMD_PLUGIN__
+      (__ESIMD_DNS::vector_type_t<T, N> *)std::addressof(src0), src1, src2);


This cast seems very suspicious

This change is removed. dpas.cpp now fails as follows:

/export/users/dongkyun/github/opensource/intel_llvm/build/bin/../include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp:473:1: note: candidate function template not viable: no known conversion from 'short * __attribute__((ext_vector_type(16)))' to 'const sycl::ext::intel::esimd::detail::vector_type_t<float, 16U> *' (aka 'const cl::sycl::ext::intel::esimd::detail::raw_vector_type<float, 16>::type *') for 1st argument __esimd_dpas_inner(const __ESIMD_DNS::vector_type_t<RT, SZ> *src0,

kbobrovs · 2022-08-12T22:44:23Z

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

+#ifdef __SYCL_DEVICE_ONLY__
+
+// TODO: __esimd_dpas* should have single declaration for host and device:
+// void __esimd_dpas*(...)


it is not void actually Ret __esimd_dpas*(...). // void __esimd_dpas*(...) can be simply removed if complete pseudo-code skeleton of the definition is not shown.

kbobrovs · 2022-08-12T22:45:54Z

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

@@ -662,15 +671,13 @@ template <__ESIMD_ENS::argument_type src1_precision,
 inline __ESIMD_DNS::vector_type_t<T, N>
 __esimd_dpas(__ESIMD_DNS::vector_type_t<T0, N> src0,
             __ESIMD_DNS::vector_type_t<T1, N1> src1,
-             __ESIMD_DNS::vector_type_t<T2, N2> src2) {
-#ifdef __SYCL_EXPLICIT_SIMD_PLUGIN__
+             __ESIMD_DNS::vector_type_t<T2, N2> src2, int sign_res,


please remove the two unused arguments

kbobrovs · 2022-08-12T22:47:10Z

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

-                   T, T0, T1, T2, N, N1, N2>(src0.data(), src1.data(),
-                                             src2.data());
+                   T, T0, T1, T2, N, N1, N2>(
+          src0.data(), src1.data(), src2.data(), dst_signed, src0_signed);


please remove this change as well

dongkyunahn-intel · 2022-08-15T19:30:17Z

Unrelated failures

SYCL / Linux / HIP AMDGPU LLVM Test Suite 
- SYCL :: DeviceLib/imf_simd_emulate_test.cpp

SYCL / Linux / CUDA LLVM Test Suite
- SYCL :: DeviceLib/imf_simd_emulate_test.cpp

kbobrovs · 2022-08-15T19:35:49Z

Unrelated failures

What are the list of failing tests per each configuration?

dongkyunahn-intel · 2022-08-15T19:40:02Z

Unrelated failures

What are the list of failing tests per each configuration?

Updated the comment with failing tests.

- Bringing back of PR#1110 as all dpas tests are passing after PR - intel/llvm#6475 is merged in intel/llvm:sycl

…te#1151) - Bringing back of PR#1110 as all dpas tests are passing after PR - intel#6475 is merged in intel/llvm:sycl

[SYCL][ESIMD][EMU] Enabling dpas with ESIMD_EMULATOR backend

0bb5952

- +fix for 'saturate<>' for ESIMD_EMULATOR backend

dongkyunahn-intel requested a review from kbobrovs July 25, 2022 21:18

dongkyunahn-intel requested a review from a team as a code owner July 25, 2022 21:18

dongkyunahn-intel requested a review from sndmitriev July 25, 2022 21:22

v-klochkov previously approved these changes Jul 26, 2022

View reviewed changes

bader changed the title ~~[SYCL][ESIMD][EMU] Enabling dpas with ESIMD_EMULATOR backend~~ [SYCL][ESIMD][EMU] Enable dpas with ESIMD_EMULATOR backend Jul 26, 2022

dongkyunahn-intel added 2 commits July 26, 2022 10:43

Separating definitions of __esimd_dpas* per __SYCL_DEVICE_ONLY__

64d375b

- As template and runtime arguments are different (e.g. 'dpas_info' is not used in ESIMD_EMULATOR backend (hostmode))

dongkyunahn-intel requested a review from v-klochkov July 26, 2022 23:42

This was referenced Jul 27, 2022

[SYCL][ESIMD][EMU] Enabling dpas_test1/2/3 tests intel/llvm-test-suite#1110

Merged

Revert "[SYCL][ESIMD][EMU] Enabling dpas_test1/2/3 tests (#1110)" intel/llvm-test-suite#1115

Merged

kbobrovs reviewed Jul 27, 2022

View reviewed changes

sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp Show resolved Hide resolved

kbobrovs reviewed Jul 27, 2022

View reviewed changes

sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp Outdated Show resolved Hide resolved

dongkyunahn-intel added 5 commits July 28, 2022 15:11

Unifying __esimd_dpas* functions

fc3bd07

- Template arguments for dpas operations (repeat count, systolic depth, src1_precision, src2_precision) are converted into runtime arguments - __esimd_dpas* functions have same signature regardless of __SYCL_DEVICE_ONLY__

Missing changes in unifying __esimd_dpas* functions

a621dd8

systolic_depth decoding bug fix

b698423

Fix for passing dpas_test.cpp

a47893c

Fixing 'odr' test failure

0e220e9

kbobrovs reviewed Jul 29, 2022

View reviewed changes

sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp Outdated Show resolved Hide resolved

kbobrovs reviewed Jul 29, 2022

View reviewed changes

sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp Outdated Show resolved Hide resolved

kbobrovs reviewed Jul 29, 2022

View reviewed changes

sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp Outdated Show resolved Hide resolved

dongkyunahn-intel added 3 commits July 29, 2022 09:12

Reverting clang-formatting for '__esimd_dpas_inner'

1960d87

Reverting (void*) work-around

3907816

Bugfix for dpas_test.cpp test failure

27ab022

kbobrovs suggested changes Aug 1, 2022

View reviewed changes

dongkyunahn-intel requested a review from kbobrovs August 9, 2022 15:02

kbobrovs reviewed Aug 12, 2022

View reviewed changes

sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp Outdated Show resolved Hide resolved

kbobrovs reviewed Aug 12, 2022

View reviewed changes

sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp Outdated Show resolved Hide resolved

kbobrovs reviewed Aug 12, 2022

View reviewed changes

sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp Outdated Show resolved Hide resolved

dongkyunahn-intel commented Aug 12, 2022

View reviewed changes

sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp Outdated Show resolved Hide resolved

kbobrovs reviewed Aug 12, 2022

View reviewed changes

Removing unnecessary changes

d06821a

kbobrovs reviewed Aug 12, 2022

View reviewed changes

dongkyunahn-intel added 2 commits August 12, 2022 15:51

Applying suggestions

063b905

Working version

1615f45

kbobrovs approved these changes Aug 15, 2022

View reviewed changes

kbobrovs merged commit 3d506a3 into intel:sycl Aug 16, 2022

dongkyunahn-intel added a commit to dongkyunahn-intel/llvm-test-suite that referenced this pull request Aug 16, 2022

[SYCL][ESIMD][EMU] Enabling dpas_test1/2/3 tests

e9398a2

- Bringing back of PR#1110 as all dpas tests are passing after PR - intel/llvm#6475 is merged in intel/llvm:sycl

dongkyunahn-intel mentioned this pull request Aug 16, 2022

[SYCL][ESIMD][EMU] Enabling dpas_test1/2/3 tests intel/llvm-test-suite#1151

Merged

kbobrovs pushed a commit to intel/llvm-test-suite that referenced this pull request Aug 17, 2022

[SYCL][ESIMD][EMU] Enabling dpas_test1/2/3 tests (#1151)

03b449b

- Bringing back of PR#1110 as all dpas tests are passing after PR - intel/llvm#6475 is merged in intel/llvm:sycl

[SYCL][ESIMD][EMU] Enable dpas with ESIMD_EMULATOR backend #6475

[SYCL][ESIMD][EMU] Enable dpas with ESIMD_EMULATOR backend #6475

Uh oh!

Conversation

dongkyunahn-intel commented Jul 25, 2022

Uh oh!

dongkyunahn-intel commented Jul 25, 2022

Uh oh!

pvchupin commented Jul 27, 2022

Uh oh!

dongkyunahn-intel commented Jul 27, 2022

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

dongkyunahn-intel commented Jul 30, 2022

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

dongkyunahn-intel commented Aug 8, 2022

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

kbobrovs Aug 12, 2022

Choose a reason for hiding this comment

Uh oh!

dongkyunahn-intel Aug 12, 2022

Choose a reason for hiding this comment

Uh oh!

kbobrovs Aug 12, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kbobrovs Aug 12, 2022

Choose a reason for hiding this comment

Uh oh!

kbobrovs Aug 12, 2022

Choose a reason for hiding this comment

Uh oh!

dongkyunahn-intel commented Aug 15, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

kbobrovs commented Aug 15, 2022

Uh oh!

dongkyunahn-intel commented Aug 15, 2022

Uh oh!

Uh oh!

kbobrovs Aug 12, 2022 •

edited

Loading

dongkyunahn-intel commented Aug 15, 2022 •

edited

Loading