[SYCL][ESIMD][EMU] Memory intrinsic updates for ESIMD_EMU plugin #4748

dongkyunahn-intel · 2021-10-11T23:30:47Z

Intrinsic implementations for host device are replaced with implementation for ESIMD_EMU

- Intrinsic implementations for host device are replaced with implementation for ESIMD_EMU

- media_ld/st host-mode implementation is replaced with ESIMD_EMU support implementation

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

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

- + Relocating SLM_BTI definition

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

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

llvm/lib/SYCLLowerIR/LowerESIMD.cpp

- and typo fix in memory.hpp

dongkyunahn-intel · 2021-10-25T23:41:42Z

Due to changes from memory intrinsic refactoring (#4720), the ESIMD Device interface used in this PR needs to be changed. I'll resume this PR as soon as my other high priority PRs are merged to the origin/sycl branch.

- intel#4748 (comment)

…emory_intrinsic

- In order to use surface-index including SLM_BTI

sycl/include/sycl/ext/intel/experimental/esimd/emu/detail/esimd_emulator_device_interface.hpp

…emory_intrinsic

- Kernel build fails because of assertion from 'translateSLMInit()' - +Typo fix in error log message

- With Version string info enclosed in device struct

- as ESIMD emulator is not productized yet

- lower_intrins.ll test fails because of '__esimd_slm_init' lowering - Failing lines are removed

…emory_intrinsic

kbobrovs · 2022-02-02T19:32:13Z

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

@@ -116,6 +116,18 @@ static bool PrintPiTrace = false;
 // Sycl RT calls piTearDown().
 static sycl::detail::ESIMDEmuPluginOpaqueData *PiESimdDeviceAccess;

+// Single-entry cache for piPlatformsGet call.
+static pi_platform PiPlatformCache;
+static sycl::detail::SpinLock *PiPlatformCacheMutex =


Please add a TODO comment that this is a memory leak, and then address in a separate PR.

Konst : Follow level_zero implementation

Please, fix this as soon as possible.

kbobrovs · 2022-02-02T19:32:32Z

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

+// Mapping between surface index and CM-managed surface
+static std::unordered_map<unsigned int, _pi_mem *> *PiESimdSurfaceMap =
+    new std::unordered_map<unsigned int, _pi_mem *>;
+static sycl::detail::SpinLock *PiESimdSurfaceMapLock =


kbobrovs · 2022-02-02T19:42:37Z

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

+static pi_platform PiPlatformCache;
+static sycl::detail::SpinLock *PiPlatformCacheMutex =
+    new sycl::detail::SpinLock;
+static bool PiPlatformCachePopulated = false;


Please make it local static in the function where it is used

- Removing unnecessary 'cm_rt.h' - TODO comments for memory leaks associated with SpinLocks - Making a static variable local one for minimized variable scope - Removing comments related to removed 'v2.h' file

…emory_intrinsic

lsatanov · 2022-02-10T13:49:54Z

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

+    new std::unordered_map<unsigned int, _pi_mem *>;
+// TODO/FIXME : Memory leak. Handle with 'piTearDown'.
+static sycl::detail::SpinLock *PiESimdSurfaceMapLock =
+    new sycl::detail::SpinLock;


Hi. Question: what's the purpose of using dynamic allocation here and above?

This part was imported from level_zero plug-in implementation. We have to follow level_zero implementation in general. I didn't have chance to ask why to level_zero implementation authors.

@smaslov-intel , could you comment on this - why level_zero use dynamic allocation for such global locks? -

llvm/sycl/plugins/level_zero/pi_level_zero.cpp

Line 394 in e6d0547

static sycl::detail::SpinLock *PiPlatformsCacheMutex =

I have a guess: it's one of ways to deal with static objects destruction order trouble (when you use a static/global object on destruction of other global/static objects and it may be already destroyed corresponding to the order of their initialization). (also this method doesn't help with static initialization order problem).

But. The big question is: is it really needed in this particular place: are those mutexes and data structures to be used on destruction of other static/global objects indirectly out of this program module (?).

@lsatanov, @smaslov-intel, @romanovvlad, @kbobrovs , I created an issue ticket and assigned it to me in order to keep track of this issue.

#5599

dongkyunahn-intel · 2022-02-15T06:36:24Z

/build

dongkyunahn-intel · 2022-02-15T06:40:33Z

/summary:run

…emory_intrinsic

dongkyunahn-intel · 2022-02-16T03:28:12Z

@romanovvlad / @intel/llvm-reviewers-runtime , would you review this PR?

romanovvlad

A couple of comments. I'm OK if they are resolved in a separate PR.

romanovvlad · 2022-02-16T16:26:08Z

sycl/source/esimd_emulator_device_interface.cpp

+    std::cerr << "ESIMD EMU plugin error or not loaded - try setting "
+                 "SYCL_DEVICE_FILTER=esimd_emulator:gpu environment variable"
+              << std::endl;


Please, do not print messages unconditionally. This message can be in the exception which is thrown below, so a user can read it by calling .what() . Or it can be done under pi trace check, see https://github.com/intel/llvm/blob/sycl/sycl/source/device_selector.cpp#L48

Will change.

romanovvlad · 2022-02-16T16:28:21Z

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

@@ -116,6 +116,18 @@ static bool PrintPiTrace = false;
 // Sycl RT calls piTearDown().
 static sycl::detail::ESIMDEmuPluginOpaqueData *PiESimdDeviceAccess;

+// Single-entry cache for piPlatformsGet call.
+static pi_platform PiPlatformCache;
+static sycl::detail::SpinLock *PiPlatformCacheMutex =


Please, fix this as soon as possible.

romanovvlad · 2022-02-16T16:33:57Z

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

@@ -401,9 +462,15 @@ pi_result piPlatformsGet(pi_uint32 NumEntries, pi_platform *Platforms,
    return PI_INVALID_VALUE;
  }

+  const std::lock_guard<sycl::detail::SpinLock> Lock{*PiPlatformCacheMutex};


This is not safe to access this var here as it might not be constructed yet. Hopefully it should be solved when all global variables are allocated(new) in piInitialize and deallocated in piTearDown.

Like most of atomicity controls in this file, PiPlatformCacheMutex was imported from level_zero implementation. I'll create a PR designated for atomicity issue.

@romanovvlad, I created an issue ticket and assigned it to me in order to keep track of this issue.

#5599

romanovvlad · 2022-02-16T16:34:41Z

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

@@ -401,9 +462,15 @@ pi_result piPlatformsGet(pi_uint32 NumEntries, pi_platform *Platforms,
    return PI_INVALID_VALUE;
  }

+  const std::lock_guard<sycl::detail::SpinLock> Lock{*PiPlatformCacheMutex};
+  if (!PiPlatformCachePopulated) {


Probably, we could check PiPlatformCache == nullptr instead.

Will apply in the PR designated for atomicity as mentioned above.

…el#4748) Intrinsic implementations for host device are replaced with implementation for ESIMD_EMU

* upstream/sycl: (2757 commits) [SYCL][Doc] Fixing incorrect merge of community Readme.md with our version (intel#5636) [SYCL] Change USM pooling parameters. (intel#5457) [CI] Fix cache location on Windows (intel#5603) [SYCL][NFC] Fix a warning about uninitialized struct members (intel#5610) [Buildbot] Update Windows GPU version to 101.1340 (intel#5620) Fix SPIRV -> OCL barrier call argument attributes Move SPV_INTEL_memory_access_aliasing tokens from spirv_internal [SYCL][ESIMD] Add support for named barrier APIs (intel#5583) [SYCL][L0] Remove ZeModule when program build failed (intel#5541) [SYCL] Silence "unknown attribute" warning for `device_indirectly_callable` (intel#5591) [SYCL][DOC] Introductory material for extensions (intel#5605) [SYCL][DOC] Change extension names to lower case (intel#5607) [SYCL] Improve get_kernel_bundle performance (intel#5496) [SYCL] Do not build device code for sub-devices (intel#5240) [sycl-post-link] Fix a crash during spec-constant properties generation (intel#5538) [SYCL][DOC] Move SPIR-V and OpenCL extensions (intel#5578) [SYCL][ESIMD][EMU] Update memory intrinsics for ESIMD_EMU plugin (intel#4748) [CI] Allow stale issue bot to analyze more issues (intel#5602) [SYCL][L0] Honor property::queue::enable_profiling (intel#5543) [OpenMP] Properly save strings when doing LTO ...

dongkyunahn-intel added 2 commits October 11, 2021 11:44

[SYCL][ESIMD][EMU] Memory intrinsic updates for ESIMD_EMU plugin

77e5c8d

- Intrinsic implementations for host device are replaced with implementation for ESIMD_EMU

Intrinsic implementations for ESIMD_EMU

bff7edd

- media_ld/st host-mode implementation is replaced with ESIMD_EMU support implementation

dongkyunahn-intel requested review from kbobrovs, kychendev and v-klochkov as code owners October 11, 2021 23:30

dongkyunahn-intel mentioned this pull request Oct 11, 2021

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

Merged

kbobrovs reviewed Oct 11, 2021

View reviewed changes

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

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

slm_init() refactoring

54da408

dongkyunahn-intel requested a review from sndmitriev as a code owner October 12, 2021 01:05

romanovvlad reviewed Oct 12, 2021

View reviewed changes

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

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

Comments for if(SLM_BTI) and otherwise case

ab85953

- + Relocating SLM_BTI definition

dongkyunahn-intel mentioned this pull request Oct 12, 2021

[ESIMD] Refactor esimd intrinsic mapping to BE intrinsics. #4720

Merged

SLM_BTI fixes / Function argument reference fixes

0e6c1ee

kbobrovs reviewed Oct 13, 2021

View reviewed changes

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

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

kbobrovs reviewed Oct 13, 2021

View reviewed changes

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

llvm/lib/SYCLLowerIR/LowerESIMD.cpp Outdated Show resolved Hide resolved

dongkyunahn-intel added 2 commits October 13, 2021 11:46

Removing unnecessary comment

e3a94bc

Sanity check for whether ESIMD_EMU pi module is loaded or not

8cb68ec

- and typo fix in memory.hpp

dongkyunahn-intel added a commit to dongkyunahn-intel/llvm that referenced this pull request Nov 1, 2021

Partially reverting slm_init() refactoring

5609718

- intel#4748 (comment)

dongkyunahn-intel added 3 commits November 18, 2021 14:24

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

7b6629b

…emory_intrinsic

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

fa4afaf

…emory_intrinsic

Memory intrinsic update for ESIMD_EMULATOR

692f8c8

- In order to use surface-index including SLM_BTI

dongkyunahn-intel requested a review from smaslov-intel as a code owner December 1, 2021 06:08

romanovvlad reviewed Dec 6, 2021

View reviewed changes

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

dongkyunahn-intel added 5 commits December 6, 2021 08:09

Version info fix for interface version upgrade

808d0c4

Removing outdated LowerESIMD.cpp for pulling sycl

09ec042

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

edb02da

…emory_intrinsic

Build error fix after merging sycl

a7e346e

Reverting __esimd_slm_init lowering

7be1079

- Kernel build fails because of assertion from 'translateSLMInit()' - +Typo fix in error log message

dongkyunahn-intel requested a review from a team as a code owner January 24, 2022 18:49

dongkyunahn-intel added 6 commits January 26, 2022 19:37

PlatformCache for piPlatformsGet PI_API

faa7bda

Enabling single-entry device cache

e0cf9ed

- With Version string info enclosed in device struct

Merging ESIMDDeviceInterface version 2 table to version 1

3863705

- as ESIMD emulator is not productized yet

Fixing lower_intrins.ll failure

263e831

- lower_intrins.ll test fails because of '__esimd_slm_init' lowering - Failing lines are removed

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

f00480f

…emory_intrinsic

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

0c6f70b

…emory_intrinsic

kbobrovs reviewed Feb 2, 2022

View reviewed changes

dongkyunahn-intel added 5 commits February 2, 2022 12:30

Addressing Konst's comments

251fe21

- Removing unnecessary 'cm_rt.h' - TODO comments for memory leaks associated with SpinLocks - Making a static variable local one for minimized variable scope - Removing comments related to removed 'v2.h' file

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

16f7c3f

…emory_intrinsic

Missing change from pulling sycl branch

3ae4be4

Excluding 'PromoT' work-around for ESIMD_EMULATOR

60d1107

Initializing return vector in 'gather' memory intrinsic

d4edc16

lsatanov reviewed Feb 10, 2022

View reviewed changes

kbobrovs previously approved these changes Feb 15, 2022

View reviewed changes

dongkyunahn-intel added 2 commits February 15, 2022 10:26

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

d7ec3ca

…emory_intrinsic

Recovering 'PromoT' for ESIMD_EMULATOR for half_type support

69f01cf

dongkyunahn-intel dismissed kbobrovs’s stale review via 69f01cf February 15, 2022 22:52

dongkyunahn-intel requested a review from romanovvlad February 16, 2022 02:28

kbobrovs approved these changes Feb 16, 2022

View reviewed changes

romanovvlad approved these changes Feb 16, 2022

View reviewed changes

dongkyunahn-intel mentioned this pull request Feb 16, 2022

[SYCL][ESIMD][EMU] Atomicity control for ESIMD_EMULATOR plug-in #5599

Open

romanovvlad merged commit 1a8f501 into intel:sycl Feb 17, 2022

maximdimakov pushed a commit to maximdimakov/llvm that referenced this pull request Feb 17, 2022

[SYCL][ESIMD][EMU] Update memory intrinsics for ESIMD_EMU plugin (int…

8358b81

…el#4748) Intrinsic implementations for host device are replaced with implementation for ESIMD_EMU

smaslov-intel pushed a commit to smaslov-intel/llvm that referenced this pull request Feb 19, 2022

[SYCL][ESIMD][EMU] Update memory intrinsics for ESIMD_EMU plugin (int…

d7a50a8

…el#4748) Intrinsic implementations for host device are replaced with implementation for ESIMD_EMU

[SYCL][ESIMD][EMU] Memory intrinsic updates for ESIMD_EMU plugin #4748

[SYCL][ESIMD][EMU] Memory intrinsic updates for ESIMD_EMU plugin #4748

Uh oh!

Conversation

dongkyunahn-intel commented Oct 11, 2021

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

dongkyunahn-intel commented Oct 25, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

kbobrovs Feb 2, 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!

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!

lsatanov Feb 10, 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!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

lsatanov Feb 11, 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!

Choose a reason for hiding this comment

Uh oh!

dongkyunahn-intel commented Feb 15, 2022

Uh oh!

dongkyunahn-intel commented Feb 15, 2022

Uh oh!

dongkyunahn-intel commented Feb 16, 2022

Uh oh!

romanovvlad left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

dongkyunahn-intel Feb 16, 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!

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!

Uh oh!

dongkyunahn-intel commented Oct 25, 2021 •

edited

Loading

kbobrovs Feb 2, 2022 •

edited

Loading

lsatanov Feb 10, 2022 •

edited

Loading

lsatanov Feb 11, 2022 •

edited

Loading

dongkyunahn-intel Feb 16, 2022 •

edited

Loading