[ESIMD] Implement unified memory API - part2 - block_load(acc, ...) #11545

v-klochkov · 2023-10-14T06:21:49Z

No description provided.

Signed-off-by: Vyacheslav N Klochkov <[email protected]>

…thru Signed-off-by: Vyacheslav N Klochkov <[email protected]>

Signed-off-by: Vyacheslav N Klochkov <[email protected]>

v-klochkov · 2023-10-18T18:43:38Z

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

-    __ESIMD_DNS::lsc_data_size DS = __ESIMD_DNS::lsc_data_size::default_size,
-    cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none,
-    typename FlagsT>
+template <typename T, int NElts, cache_hint L1H = cache_hint::none,


this block_load_impl() function is actually the copy of the old experimental::lsc_block_load().
Please pay attention that this function here doesn't have the 'lsc_data_size' template parameter that was used in the old experimental::lsc_block_load().
The reason for that is that the data-size is always deduced from the loaded element type.

v-klochkov · 2023-10-18T18:52:08Z

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

-template <typename Tx, int N,
-          typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
-__ESIMD_API std::enable_if_t<is_simd_flag_type_v<Flags>, simd<Tx, N>>
-block_load(const Tx *addr, Flags) {


This function and the next 2 are not removed - they are just moved, grouped with other user-visible block_load() functions (and block_store() functions).

sarnex

LGTM! Only minor comments!

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

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

sarnex · 2023-10-18T19:01:44Z

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

@@ -1475,6 +1348,7 @@ lsc_block_load(AccessorTy acc, uint32_t offset, __ESIMD_NS::simd_mask<1> pred,
  return lsc_slm_block_load<T, NElts, DS>(
      offset + __ESIMD_DNS::localAccessorToOffset(acc), pred, old_values);
 }
+


Nit: Remove newline if no other changes

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

sarnex · 2023-10-18T20:12:21Z

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

+__ESIMD_API std::enable_if_t<
+    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
+        sycl::detail::acc_properties::is_accessor_v<AccessorT> &&
+        !sycl::detail::acc_properties::is_local_accessor_v<AccessorT>,


Maybe a dumb question but we don't allow local accessors for a couple of the functions here, why is that?

There is such function in experimental/memory.hpp. It is just not ported by this patch yet:

template <typename T, int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename AccessorTy, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag> __ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> && __ESIMD_NS::is_simd_flag_type_v<FlagsT>, __ESIMD_NS::simd<T, NElts>> lsc_block_load(AccessorTy acc, uint32_t offset, __ESIMD_NS::simd_mask<1> pred = 1, FlagsT flags = FlagsT{}) { return lsc_slm_block_load<T, NElts, DS>( offset + __ESIMD_DNS::localAccessorToOffset(acc), pred); }

sycl/test-e2e/ESIMD/unified_memory_api/block_load_acc_pvc.cpp

v-klochkov · 2023-10-19T01:14:50Z

@sarnex - Thank you for quick code-review. Because the comments required only few NFC changes - I will merge the current change and upload a separate NFC fix now.

Signed-off-by: Vyacheslav N Klochkov <[email protected]>

turinevgeny · 2023-10-18T23:43:50Z

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

-/// Allowed \c NElts values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256,
+/// Instruction can load max: DG2(64xD32 or 32xD64), PVC(64xD32 or 64xD64).
+///
+/// Accesses contiguous block of memory of `NElts * sizeof(T)` bytes  starting


Suggested change

/// Accesses contiguous block of memory of `NElts * sizeof(T)` bytes starting

/// Accesses contiguous block of memory of `NElts * sizeof(T)` bytes starting

…11594) Signed-off-by: Vyacheslav N Klochkov <[email protected]>

Add new ESIMD/unified_memory_api/block_load_acc.cpp test (added by intel#11545) to XFAIL list. Signed-off-by: Victor Perez <[email protected]>

v-klochkov force-pushed the esimd_block_load_acc branch from fc989de to 1ec774a Compare October 14, 2023 07:35

v-klochkov had a problem deploying to WindowsCILock October 14, 2023 07:37 — with GitHub Actions Failure

v-klochkov had a problem deploying to WindowsCILock October 14, 2023 07:58 — with GitHub Actions Failure

v-klochkov force-pushed the esimd_block_load_acc branch from 1ec774a to 21f7f62 Compare October 16, 2023 18:18

v-klochkov had a problem deploying to WindowsCILock October 16, 2023 18:41 — with GitHub Actions Failure

v-klochkov had a problem deploying to WindowsCILock October 16, 2023 20:15 — with GitHub Actions Failure

v-klochkov force-pushed the esimd_block_load_acc branch from 21f7f62 to 004697d Compare October 16, 2023 21:29

v-klochkov temporarily deployed to WindowsCILock October 16, 2023 21:31 — with GitHub Actions Inactive

v-klochkov had a problem deploying to WindowsCILock October 16, 2023 22:11 — with GitHub Actions Failure

v-klochkov force-pushed the esimd_block_load_acc branch 2 times, most recently from 7dc87a3 to d3da039 Compare October 17, 2023 01:56

v-klochkov had a problem deploying to WindowsCILock October 17, 2023 01:57 — with GitHub Actions Failure

v-klochkov temporarily deployed to WindowsCILock October 17, 2023 02:17 — with GitHub Actions Inactive

EwanC mentioned this pull request Oct 17, 2023

[SYCL][Graph][L0] Reset L0 events for UR sync-points #11553

Merged

v-klochkov force-pushed the esimd_block_load_acc branch from d3da039 to f9750d1 Compare October 17, 2023 17:40

v-klochkov had a problem deploying to WindowsCILock October 17, 2023 17:46 — with GitHub Actions Failure

v-klochkov had a problem deploying to WindowsCILock October 17, 2023 18:25 — with GitHub Actions Failure

v-klochkov force-pushed the esimd_block_load_acc branch from f9750d1 to 60cf644 Compare October 18, 2023 01:23

v-klochkov temporarily deployed to WindowsCILock October 18, 2023 01:24 — with GitHub Actions Inactive

v-klochkov had a problem deploying to WindowsCILock October 18, 2023 02:01 — with GitHub Actions Failure

[ESIMD] Unified memory API - part2 - block_load(acc, ...)

91e8d45

Signed-off-by: Vyacheslav N Klochkov <[email protected]>

v-klochkov force-pushed the esimd_block_load_acc branch from 60cf644 to 91e8d45 Compare October 18, 2023 15:24

v-klochkov marked this pull request as ready for review October 18, 2023 15:26

v-klochkov requested a review from a team as a code owner October 18, 2023 15:26

v-klochkov changed the title ~~[ESIMD] Unified memory API - part2 - block_load(acc, ...)~~ [ESIMD] Implement unified memory API - part2 - block_load(acc, ...) Oct 18, 2023

v-klochkov temporarily deployed to WindowsCILock October 18, 2023 15:35 — with GitHub Actions Inactive

v-klochkov added 2 commits October 18, 2023 08:50

[NFC] removed forgotten comment/mark, var rename: old_values -> pass_…

ba98c74

…thru Signed-off-by: Vyacheslav N Klochkov <[email protected]>

clang-format

4c8366d

Signed-off-by: Vyacheslav N Klochkov <[email protected]>

v-klochkov requested review from sarnex and turinevgeny October 18, 2023 15:55

v-klochkov temporarily deployed to WindowsCILock October 18, 2023 16:03 — with GitHub Actions Inactive

v-klochkov temporarily deployed to WindowsCILock October 18, 2023 16:53 — with GitHub Actions Inactive

v-klochkov commented Oct 18, 2023

View reviewed changes

sarnex approved these changes Oct 18, 2023

View reviewed changes

v-klochkov merged commit 6a98330 into intel:sycl Oct 19, 2023

v-klochkov deleted the esimd_block_load_acc branch October 19, 2023 01:14

v-klochkov added a commit to v-klochkov/llvm that referenced this pull request Oct 19, 2023

[ESIMD][NFC] Follow-up changes for intel#11545

7282511

Signed-off-by: Vyacheslav N Klochkov <[email protected]>

turinevgeny reviewed Oct 19, 2023

View reviewed changes

againull pushed a commit that referenced this pull request Oct 19, 2023

[ESIMD][NFC] Remove few unused intrinsics, follow-up NFC for #11545 (#…

a0f688f

…11594) Signed-off-by: Vyacheslav N Klochkov <[email protected]>

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[ESIMD] Implement unified memory API - part2 - block_load(acc, ...) #11545

[ESIMD] Implement unified memory API - part2 - block_load(acc, ...) #11545

Uh oh!

v-klochkov commented Oct 14, 2023

Uh oh!

v-klochkov Oct 18, 2023

Uh oh!

v-klochkov Oct 18, 2023

Uh oh!

sarnex left a comment

Uh oh!

Uh oh!

Uh oh!

sarnex Oct 18, 2023

Uh oh!

Uh oh!

sarnex Oct 18, 2023

Uh oh!

v-klochkov Oct 19, 2023

Uh oh!

Uh oh!

v-klochkov commented Oct 19, 2023

Uh oh!

turinevgeny Oct 18, 2023

Uh oh!

Uh oh!

	/// Accesses contiguous block of memory of `NElts * sizeof(T)` bytes starting
	/// Accesses contiguous block of memory of `NElts * sizeof(T)` bytes starting

[ESIMD] Implement unified memory API - part2 - block_load(acc, ...) #11545

[ESIMD] Implement unified memory API - part2 - block_load(acc, ...) #11545

Uh oh!

Conversation

v-klochkov commented Oct 14, 2023

Uh oh!

v-klochkov Oct 18, 2023

Choose a reason for hiding this comment

Uh oh!

v-klochkov Oct 18, 2023

Choose a reason for hiding this comment

Uh oh!

sarnex left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

sarnex Oct 18, 2023

Choose a reason for hiding this comment

Uh oh!

Uh oh!

sarnex Oct 18, 2023

Choose a reason for hiding this comment

Uh oh!

v-klochkov Oct 19, 2023

Choose a reason for hiding this comment

Uh oh!

Uh oh!

v-klochkov commented Oct 19, 2023

Uh oh!

turinevgeny Oct 18, 2023

Choose a reason for hiding this comment

Uh oh!

Uh oh!