[SYCL][ESIMD][EMU] ESIMD_CPU Kernel launch and ESIMD_EMU backend loading #4020

dongkyunahn-intel · 2021-06-29T05:41:50Z

This PR is for enabling kernel launching for ESIMD_CPU
Also contains emulated intrinsics for memory operations
esimd_cpu backend is loaded in SYCL runtime
Base PR : [SYCL][ESIMD][EMU] pi_esimd_cpu bringing up with CM library #4011

* This PR is for enabling kernel launching for ESIMD_CPU * Also contains emulated intrinsics for memory operations * esimd_cpu backend is loaded in SYCL runtime * Base PR : intel#4011

- Replacing deprecated ATOMIC_* with atomic_op - Relocating atomic_add implementation for flat_atomic1 from slm_atomic1

- _pi_buffer and _pi_image are defined in 'cmrt_if_defs.hpp' which is used for both PI and kernel compilation - Dependency is resolved by adding another device interface call (sycl_get_cm_buffer_params, sycl_get_cm_image_params)

smaslov-intel

PI-related changes LGTM

dongkyunahn-intel · 2021-07-09T05:37:37Z

I'll revert some changes (relocation + update) in following files and apply them in PR#4011 due to changes in Device interface.

esimd_emu_functions_v1.h
esimdcpu_device_interface.hpp

- These changes are already applied in PR#4011

- esimd_emu_functions_v1.h / esimdcpu_device_interface.hpp - Reverted file changes are to be relocated in PR#4011

dongkyunahn-intel · 2021-07-09T05:51:34Z

Revert is complete.

- Address of mutex is used for fetching mutex - Interface change is from PR#4011

smaslov-intel

PI related changes LGTM

dongkyunahn-intel · 2021-07-13T03:10:41Z

Jenkins and buildbot jobs are failing due to rebase failure. I'll make single commit using these change sets and merge it to the origin/sycl branch after PR#4011 is merged to origin/sycl.

- __esimd_surf_read/write

sycl/include/CL/sycl/accessor.hpp

romanovvlad

LGTM except some "non-functional" comments".
(haven't reviewed logic of __esimd* builtins)

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

- Kernel build failure fix from fully-working branch - Changing Header file inclusion for ESIMD_CPU

romanovvlad

LGTM

dongkyunahn-intel · 2021-07-26T18:54:44Z

@romanovvlad , Thanks for approval. I'll resolve conflicts and resolve build failures.

dongkyunahn-intel · 2021-07-27T17:56:33Z

Buildbot job failed because of missing implementation of PI_API. PR containing PI_API implementation is under review (#4011). I'll try to merge this change set after the PR is merged into the SYCL branch.

romanovvlad · 2021-10-22T09:43:50Z

/summary:run

…rnel_launch_memory_intrinsic

dongkyunahn-intel · 2021-10-28T07:33:40Z

/summary:run

dm-vodopyanov · 2021-11-08T11:20:49Z

@dongkyunahn-intel, can you please resolve the CI failures in Jenkins/Summary?

dongkyunahn-intel · 2021-11-08T16:52:07Z

@dm-vodopyanov ,

I've been trying to resolve the failure, but is not successful yet.

The failure seems to be segmentation fault (memory access violation) under windows environment and occurs only for one precision - system::double16 (https://github.com/KhronosGroup/SYCL-CTS/blob/65856bfb5844fc73f1c4e0841bf4be4c6028eef9/tests/handler/handler_copy.cpp#L948). Other precisions like int and double are fine.

dongkyunahn-intel · 2021-11-12T22:22:38Z

It is found that the segmentation fault is because of violation against 16-byte memory access alignment required by XMM access. Currently, only warning message is printed out during compilation for cases where the requirement cannot be satisfied -

llvm/sycl/include/CL/sycl/types.hpp

Line 545 in 6761e73

#pragma message ("Alignment of class vec is not in accordance with SYCL \

.

This issue has to be addressed in a separate PR.

dongkyunahn-intel · 2021-11-13T00:34:34Z

PR for memory access alignment is posted : #4953

sycl/include/CL/sycl/detail/cg_types.hpp

sycl/include/CL/sycl/handler.hpp

- Allowing access 'MKernel' from sycl::handler class - +Reverting unnecessary std:move for Kernel function

…rnel_launch_memory_intrinsic

dongkyunahn-intel · 2021-11-19T03:58:29Z

/summary:run

dongkyunahn-intel · 2021-11-19T20:01:51Z

@romanovvlad , @kbobrovs , @s-kanaev , please review and approve this PR.

Cause of failures from 'test_handler' is identified, and corresponding ticket was created & shared internally.

dongkyunahn-intel · 2021-11-23T15:51:44Z

Second review/approval request : @romanovvlad , @kbobrovs , @s-kanaev.

romanovvlad · 2021-11-23T18:07:03Z

/summary:run

…rnel_launch_memory_intrinsic

dongkyunahn-intel · 2021-11-23T19:41:14Z

/summary:run

romanovvlad · 2021-11-24T07:27:12Z

@kbobrovs, @v-klochkov , @smaslov-intel Could you please review and approve if the patch is OK for you?

smaslov-intel

PI related changes LGTM

dongkyunahn-intel · 2021-11-29T19:41:58Z

@kbobrovs , @v-klochkov , Please review and approve if this patch looks okay to you.

v-klochkov · 2022-04-06T13:46:44Z

sycl/source/detail/device_filter.cpp

@@ -91,7 +91,7 @@ device_filter::device_filter(const std::string &FilterString) {
      std::string Message =
          std::string("Invalid device filter: ") + FilterString +
          "\nPossible backend values are "
-          "{host,opencl,level_zero,cuda,hip,*}.\n"
+          "{host,opencl,level_zero,cuda,hip,esimd_emulator*}.\n"


Late comment: This line would be better as below:
"{host,opencl,level_zero,cuda,hip,esimd_emulator,*}.\n"

Will apply this change as piggyback in another PR later.

intel#4020 introduced some kind of wrapper for SYCL kernels on host to be able to launch them on ESIMD emulator backend. Even though we had dropped that feature since then, we didn't remove corresponding `handler.hpp` modifications. This PR removes them. The main effect expected from this PR is compilation time improvement: those kernel wrappers are specialized by kernel name, meaning that there will be plenty of extra useless functions emitted during host compilation pass for every kernel in a program. This patch also uncovered a missing case in `HostKernel::InstantiateKernelOnHost` which we couldn't ever encounter because we transformed a host kernel to always accept `nd_item`, thus always skipping problematic `item` code path. This PR is not expected to introduce any functional changes, but since I'm not very familiar with the SYCL RT, I'm not entirely sure of that.

#4020 introduced some kind of wrapper for SYCL kernels on host to be able to launch them on ESIMD emulator backend. Even though we had dropped that feature since then, we didn't remove corresponding `handler.hpp` modifications. This PR removes them. The main effect expected from this PR is compilation time improvement: those kernel wrappers are specialized by kernel name, meaning that there will be plenty of extra useless functions emitted during host compilation pass for every kernel in a program. This patch also uncovered a missing case in `HostKernel::InstantiateKernelOnHost` which we couldn't ever encounter because we transformed a host kernel to always accept `nd_item`, thus always skipping problematic `item` code path. This PR is not expected to introduce any functional changes, but since I'm not very familiar with the SYCL RT, I'm not entirely sure of that.

[SYCL][ESIMD][EMU] ESIMD_CPU Kernel launch and Emulated Intrinsics

ee6e9e8

* This PR is for enabling kernel launching for ESIMD_CPU * Also contains emulated intrinsics for memory operations * esimd_cpu backend is loaded in SYCL runtime * Base PR : intel#4011

dongkyunahn-intel requested review from DenisBakhvalov, kbobrovs, smaslov-intel and a team as code owners June 29, 2021 05:41

dongkyunahn-intel added 2 commits June 29, 2021 13:57

Atomic fix

6b15320

- Replacing deprecated ATOMIC_* with atomic_op - Relocating atomic_add implementation for flat_atomic1 from slm_atomic1

dongkyunahn-intel mentioned this pull request Jul 8, 2021

[SYCL][ESIMD][EMU] pi_esimd_cpu bringing up with CM library #4011

Merged

smaslov-intel previously approved these changes Jul 8, 2021

View reviewed changes

dongkyunahn-intel added 2 commits July 8, 2021 22:39

Revert changes in ESIMD_CPU device interface

7b163d6

- These changes are already applied in PR#4011

Revert file relocation for future merging

0151dc7

- esimd_emu_functions_v1.h / esimdcpu_device_interface.hpp - Reverted file changes are to be relocated in PR#4011

dongkyunahn-intel dismissed smaslov-intel’s stale review via 0151dc7 July 9, 2021 05:51

Changes in ESIMD_CPU device interface definition

7533a83

- Address of mutex is used for fetching mutex - Interface change is from PR#4011

smaslov-intel previously approved these changes Jul 12, 2021

View reviewed changes

Removing interleaved '__SYCL_DEVICE_ONLY__' in memory_intrin.hpp

c96efe1

- __esimd_surf_read/write

dongkyunahn-intel dismissed smaslov-intel’s stale review via c96efe1 July 15, 2021 19:12

kbobrovs reviewed Jul 17, 2021

View reviewed changes

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

romanovvlad reviewed Jul 19, 2021

View reviewed changes

Removing 'get_pointer()' and changing space for 'raw_send'

1266ceb

- Kernel build failure fix from fully-working branch - Changing Header file inclusion for ESIMD_CPU

romanovvlad previously approved these changes Jul 26, 2021

View reviewed changes

Merge branch 'sycl' into esimdcpu_kernel_launch_memory_intrinsic

986a4f8

dongkyunahn-intel dismissed romanovvlad’s stale review via 986a4f8 July 26, 2021 18:53

Build failure fix after merging sycl branch

c14a755

kbobrovs previously approved these changes Jul 27, 2021

View reviewed changes

Merge branch 'sycl' into esimdcpu_kernel_launch_memory_intrinsic

33c3645

Merge branch 'sycl' of https://github.com/intel/llvm into esimdcpu_ke…

6985a64

…rnel_launch_memory_intrinsic

s-kanaev reviewed Nov 15, 2021

View reviewed changes

sycl/include/CL/sycl/detail/cg_types.hpp Outdated Show resolved Hide resolved

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

Re-privatizing 'MKernel' for backward compatibility

8ee0bd8

- Allowing access 'MKernel' from sycl::handler class - +Reverting unnecessary std:move for Kernel function

dongkyunahn-intel dismissed stale reviews from againull and kbobrovs via 8ee0bd8 November 16, 2021 19:22

dongkyunahn-intel added 3 commits November 18, 2021 10:44

Merge branch 'sycl' of https://github.com/intel/llvm into esimdcpu_ke…

2bf84de

…rnel_launch_memory_intrinsic

ParseAllowList failure fix

b85ba9f

clang-format fix

7a3e968

Merge branch 'sycl' of https://github.com/intel/llvm into esimdcpu_ke…

270763f

…rnel_launch_memory_intrinsic

romanovvlad approved these changes Nov 24, 2021

View reviewed changes

smaslov-intel approved these changes Nov 24, 2021

View reviewed changes

kbobrovs approved these changes Nov 29, 2021

View reviewed changes

kbobrovs merged commit f4ad3c1 into intel:sycl Nov 30, 2021

v-klochkov reviewed Apr 6, 2022

View reviewed changes

AlexeySachkov mentioned this pull request Sep 24, 2024

[SYCL][NFCI] Drop ESIMD emulator leftovers #15495

Merged

[SYCL][ESIMD][EMU] ESIMD_CPU Kernel launch and ESIMD_EMU backend loading #4020

[SYCL][ESIMD][EMU] ESIMD_CPU Kernel launch and ESIMD_EMU backend loading #4020

Uh oh!

Conversation

dongkyunahn-intel commented Jun 29, 2021

Uh oh!

smaslov-intel left a comment

Choose a reason for hiding this comment

Uh oh!

dongkyunahn-intel commented Jul 9, 2021

Uh oh!

dongkyunahn-intel commented Jul 9, 2021

Uh oh!

smaslov-intel left a comment

Choose a reason for hiding this comment

Uh oh!

dongkyunahn-intel commented Jul 13, 2021

Uh oh!

Uh oh!

romanovvlad left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

romanovvlad left a comment

Choose a reason for hiding this comment

Uh oh!

dongkyunahn-intel commented Jul 26, 2021

Uh oh!

dongkyunahn-intel commented Jul 27, 2021

Uh oh!

romanovvlad commented Oct 22, 2021

Uh oh!

dongkyunahn-intel commented Oct 28, 2021

Uh oh!

dm-vodopyanov commented Nov 8, 2021

Uh oh!

dongkyunahn-intel commented Nov 8, 2021

Uh oh!

dongkyunahn-intel commented Nov 12, 2021

Uh oh!

dongkyunahn-intel commented Nov 13, 2021

Uh oh!

Uh oh!

Uh oh!

dongkyunahn-intel commented Nov 19, 2021

Uh oh!

dongkyunahn-intel commented Nov 19, 2021

Uh oh!

dongkyunahn-intel commented Nov 23, 2021

Uh oh!

romanovvlad commented Nov 23, 2021

Uh oh!

dongkyunahn-intel commented Nov 23, 2021

Uh oh!

romanovvlad commented Nov 24, 2021

Uh oh!

smaslov-intel left a comment

Choose a reason for hiding this comment

Uh oh!

dongkyunahn-intel commented Nov 29, 2021

Uh oh!

v-klochkov Apr 6, 2022

Choose a reason for hiding this comment

Uh oh!

dongkyunahn-intel Apr 6, 2022

Choose a reason for hiding this comment

Uh oh!

Uh oh!