From 464c23edb52f5af669627ee26442c127b9662778 Mon Sep 17 00:00:00 2001
From: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
Date: Sat, 17 Apr 2021 02:17:57 -0700
Subject: [PATCH 1/7] [SYCL][ESIMD] Deprecate block_load/store, add
 simd::copy_from/to.

This patch:
1) Fixes the following TODO in esimd_memory.hpp:
  // TODO @rolandschulz
  // Should follow existing std::simd naming for similar APIs - "copy_from" and
  // "copy_to" to avoid confusion.

2) Adds type checks for the sycl accessor arguments in the added APIs.

Signed-off-by: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
---
 .../INTEL/esimd/detail/esimd_sycl_util.hpp    |  87 +++++++++++++
 sycl/include/CL/sycl/INTEL/esimd/esimd.hpp    | 122 ++++++++++++++++++
 .../CL/sycl/INTEL/esimd/esimd_memory.hpp      |  53 +++-----
 sycl/test/esimd/block_load_store.cpp          |  32 ++++-
 sycl/test/esimd/simd_copy_to_copy_from.cpp    |  78 +++++++++++
 5 files changed, 327 insertions(+), 45 deletions(-)
 create mode 100644 sycl/include/CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp
 create mode 100644 sycl/test/esimd/simd_copy_to_copy_from.cpp

diff --git a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp
new file mode 100644
index 0000000000000..ffed6b9fd7e61
--- /dev/null
+++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp
@@ -0,0 +1,87 @@
+//==------------- esimd_sycl_util.hpp - DPC++ Explicit SIMD API  -----------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// Utility functions related to interaction with generic SYCL and used for
+// implementing Explicit SIMD APIs.
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/accessor.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace INTEL {
+namespace gpu {
+namespace detail {
+
+// Checks that given type is a SYCL accessor type. Sets its static field
+// \c value accordingly. Also, if the check is succesful, sets \c mode and
+// \c target static fields to the accessor type's access mode and access target
+// respectively. Otherwise they are set to -1.
+template <typename T> struct is_sycl_accessor : public std::false_type {
+  static constexpr sycl::access::mode mode =
+      static_cast<sycl::access::mode>(-1);
+  static constexpr sycl::access::target target =
+      static_cast<sycl::access::target>(-1);
+};
+
+template <typename DataT, int Dimensions, sycl::access::mode AccessMode,
+          sycl::access::target AccessTarget,
+          sycl::access::placeholder IsPlaceholder, typename PropertyListT>
+struct is_sycl_accessor<sycl::accessor<
+    DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder, PropertyListT>>
+    : public std::true_type {
+  static constexpr sycl::access::mode mode = AccessMode;
+  static constexpr sycl::access::target target = AccessTarget;
+};
+
+using accessor_mode_cap_val_t = bool;
+
+// Denotes an accessor's capability - whether it can read or write.
+struct accessor_mode_cap {
+  static inline constexpr accessor_mode_cap_val_t can_read = false;
+  static inline constexpr accessor_mode_cap_val_t can_write = true;
+};
+
+template <sycl::access::mode Mode, accessor_mode_cap_val_t Cap>
+constexpr bool accessor_mode_has_capability() {
+  static_assert(Cap == accessor_mode_cap::can_read ||
+                    Cap == accessor_mode_cap::can_write,
+                "unsupported capability");
+
+  if constexpr (Mode == sycl::access::mode::atomic ||
+                Mode == sycl::access::mode::read_write ||
+                Mode == sycl::access::mode::discard_read_write)
+    return true; // atomic and *read_write accessors can read/write
+
+  return (Cap == accessor_mode_cap::can_read) ==
+         (Mode == sycl::access::mode::read);
+}
+
+// Checks that given type is a SYCL accessor type with given capability and
+// target.
+template <typename T, accessor_mode_cap_val_t Capability,
+          sycl::access::target AccessTarget>
+struct is_sycl_accessor_with
+    : public std::conditional_t<
+          accessor_mode_has_capability<is_sycl_accessor<T>::mode,
+                                       Capability>() &&
+              (is_sycl_accessor<T>::target == AccessTarget),
+          std::true_type, std::false_type> {};
+
+#define __ESIMD_ENABLE_IF_ACCESSOR(T, acc_capability, acc_target, ret_type)    \
+  sycl::detail::enable_if_t<detail::is_sycl_accessor_with<                     \
+                                T, detail::accessor_mode_cap::acc_capability,  \
+                                sycl::access::target::acc_target>::value,      \
+                            ret_type>
+
+} // namespace detail
+} // namespace gpu
+} // namespace INTEL
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
index d2e21bd77cc50..62ad569a6efd2 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
@@ -11,6 +11,8 @@
 #pragma once
 
 #include <CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp>
+#include <CL/sycl/INTEL/esimd/detail/esimd_memory_intrin.hpp>
+#include <CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp>
 #include <CL/sycl/INTEL/esimd/detail/esimd_types.hpp>
 
 __SYCL_INLINE_NAMESPACE(cl) {
@@ -480,6 +482,47 @@ template <typename Ty, int N> class simd {
     }
   }
 
+  /// @name Memory operations
+  /// TODO NOTE: These APIs do not support cache hint specification yet, as this
+  /// is WIP. Later addition of hints is not expected to break code using these
+  /// APIs.
+  ///
+  /// @{
+
+  /// Copy a contiguous block of data from memory into this simd object.
+  /// The amount of memory copied equals the total size of vector elements in
+  /// this object.
+  /// @param addr the memory address to copy from. Must be a pointer to the
+  /// global address space, otherwise behavior is undefined.
+  ESIMD_INLINE void copy_from(const Ty *const addr) SYCL_ESIMD_FUNCTION;
+
+  /// Copy a contiguous block of data from memory into this simd object.
+  /// The amount of memory copied equals the total size of vector elements in
+  /// this object.
+  /// Source memory location is represented via a global accessor and offset.
+  /// @param acc accessor to copy from.
+  /// @param offset offset to copy from.
+  template <typename AccessorT>
+  ESIMD_INLINE __ESIMD_ENABLE_IF_ACCESSOR(AccessorT, can_read, global_buffer,
+                                          void)
+      copy_from(AccessorT acc, uint32_t offset) SYCL_ESIMD_FUNCTION;
+
+  /// Copy all vector elements of this object into a contiguous block in memory.
+  /// @param addr the memory address to copy to. Must be a pointer to the
+  /// global address space, otherwise behavior is undefined.
+  ESIMD_INLINE void copy_to(Ty *addr) SYCL_ESIMD_FUNCTION;
+
+  /// Copy all vector elements of this object into a contiguous block in memory.
+  /// Destination memory location is represented via a global accessor and
+  /// offset.
+  /// @param acc accessor to copy from.
+  /// @param offset offset to copy from.
+  template <typename AccessorT>
+  ESIMD_INLINE __ESIMD_ENABLE_IF_ACCESSOR(AccessorT, can_write, global_buffer,
+                                          void)
+      copy_to(AccessorT acc, uint32_t offset) SYCL_ESIMD_FUNCTION;
+
+  /// @} // Memory operations
 private:
   // The underlying data for this vector.
   vector_type M_data;
@@ -498,6 +541,84 @@ ESIMD_INLINE simd<U, n> convert(simd<T, n> val) {
   return __builtin_convertvector(val.data(), detail::vector_type_t<U, n>);
 }
 
+// ----------- Outlined implementations of esimd class APIs.
+
+template <typename T, int N> void simd<T, N>::copy_from(const T *const addr) {
+  constexpr unsigned Sz = sizeof(T) * N;
+  static_assert(Sz >= detail::OperandSize::OWORD,
+                "block size must be at least 1 oword");
+  static_assert(Sz % detail::OperandSize::OWORD == 0,
+                "block size must be whole number of owords");
+  static_assert(detail::isPowerOf2(Sz / detail::OperandSize::OWORD),
+                "block must be 1, 2, 4 or 8 owords long");
+  static_assert(Sz <= 8 * detail::OperandSize::OWORD,
+                "block size must be at most 8 owords");
+
+  uintptr_t AddrVal = reinterpret_cast<uintptr_t>(addr);
+  *this =
+      __esimd_flat_block_read_unaligned<T, N, CacheHint::None, CacheHint::None>(
+          AddrVal);
+}
+
+template <typename T, int N>
+template <typename AccessorT>
+__ESIMD_ENABLE_IF_ACCESSOR(AccessorT, can_read, global_buffer, void)
+simd<T, N>::copy_from(AccessorT acc, uint32_t offset) {
+  constexpr unsigned Sz = sizeof(T) * N;
+  static_assert(Sz >= detail::OperandSize::OWORD,
+                "block size must be at least 1 oword");
+  static_assert(Sz % detail::OperandSize::OWORD == 0,
+                "block size must be whole number of owords");
+  static_assert(detail::isPowerOf2(Sz / detail::OperandSize::OWORD),
+                "block must be 1, 2, 4 or 8 owords long");
+  static_assert(Sz <= 8 * detail::OperandSize::OWORD,
+                "block size must be at most 8 owords");
+#if defined(__SYCL_DEVICE_ONLY__)
+  auto surf_ind = detail::AccessorPrivateProxy::getNativeImageObj(acc);
+  *this = __esimd_block_read<T, N>(surf_ind, offset);
+#else
+  *this = __esimd_block_read<T, N>(acc, offset);
+#endif // __SYCL_DEVICE_ONLY__
+}
+
+template <typename T, int N> void simd<T, N>::copy_to(T *addr) {
+  constexpr unsigned Sz = sizeof(T) * N;
+  static_assert(Sz >= detail::OperandSize::OWORD,
+                "block size must be at least 1 oword");
+  static_assert(Sz % detail::OperandSize::OWORD == 0,
+                "block size must be whole number of owords");
+  static_assert(detail::isPowerOf2(Sz / detail::OperandSize::OWORD),
+                "block must be 1, 2, 4 or 8 owords long");
+  static_assert(Sz <= 8 * detail::OperandSize::OWORD,
+                "block size must be at most 8 owords");
+
+  uintptr_t AddrVal = reinterpret_cast<uintptr_t>(addr);
+  __esimd_flat_block_write<T, N, CacheHint::None, CacheHint::None>(AddrVal,
+                                                                   data());
+}
+
+template <typename T, int N>
+template <typename AccessorT>
+__ESIMD_ENABLE_IF_ACCESSOR(AccessorT, can_write, global_buffer, void)
+simd<T, N>::copy_to(AccessorT acc, uint32_t offset) {
+  constexpr unsigned Sz = sizeof(T) * N;
+  static_assert(Sz >= detail::OperandSize::OWORD,
+                "block size must be at least 1 oword");
+  static_assert(Sz % detail::OperandSize::OWORD == 0,
+                "block size must be whole number of owords");
+  static_assert(detail::isPowerOf2(Sz / detail::OperandSize::OWORD),
+                "block must be 1, 2, 4 or 8 owords long");
+  static_assert(Sz <= 8 * detail::OperandSize::OWORD,
+                "block size must be at most 8 owords");
+
+#if defined(__SYCL_DEVICE_ONLY__)
+  auto surf_ind = detail::AccessorPrivateProxy::getNativeImageObj(acc);
+  __esimd_block_write<T, N>(surf_ind, offset >> 4, data());
+#else
+  __esimd_block_write<T, N>(acc, offset >> 4, data());
+#endif // __SYCL_DEVICE_ONLY__
+}
+
 } // namespace gpu
 } // namespace INTEL
 } // namespace sycl
@@ -516,4 +637,5 @@ std::ostream &operator<<(std::ostream &OS,
   OS << "}";
   return OS;
 }
+
 #endif
diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp
index a4f072c317af9..9160323adbed7 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp
@@ -158,14 +158,15 @@ scatter(T *p, simd<T, n * ElemsPerAddr> vals, simd<uint32_t, n> offsets,
                             pred.data());
 }
 
-// TODO @rolandschulz
-// Should follow existing std::simd naming for similar APIs - "copy_from" and
-// "copy_to" to avoid confusion.
-//
 /// Flat-address block-load.
 /// \ingroup sycl_esimd
+// TODO normally, this function should just delegate to
+// simd::copy_from for the deprecation period, but separate implementations are
+// needed for now, as simd::copy_from does not support cache hints yet.
+// This API, even though deprecated, can't be removed until then.
 template <typename T, int n, CacheHint L1H = CacheHint::None,
           CacheHint L3H = CacheHint::None>
+__SYCL_DEPRECATED("Replaced by simd::copy_from")
 ESIMD_INLINE ESIMD_NODEBUG simd<T, n> block_load(const T *const addr) {
   constexpr unsigned Sz = sizeof(T) * n;
   static_assert(Sz >= detail::OperandSize::OWORD,
@@ -184,30 +185,20 @@ ESIMD_INLINE ESIMD_NODEBUG simd<T, n> block_load(const T *const addr) {
 /// Accessor-based block-load.
 /// \ingroup sycl_esimd
 template <typename T, int n, typename AccessorTy>
+__SYCL_DEPRECATED("Replaced by simd::copy_from")
 ESIMD_INLINE ESIMD_NODEBUG simd<T, n> block_load(AccessorTy acc,
                                                  uint32_t offset) {
-  constexpr unsigned Sz = sizeof(T) * n;
-  static_assert(Sz >= detail::OperandSize::OWORD,
-                "block size must be at least 1 oword");
-  static_assert(Sz % detail::OperandSize::OWORD == 0,
-                "block size must be whole number of owords");
-  static_assert(detail::isPowerOf2(Sz / detail::OperandSize::OWORD),
-                "block must be 1, 2, 4 or 8 owords long");
-  static_assert(Sz <= 8 * detail::OperandSize::OWORD,
-                "block size must be at most 8 owords");
-
-#if defined(__SYCL_DEVICE_ONLY__)
-  auto surf_ind = detail::AccessorPrivateProxy::getNativeImageObj(acc);
-  return __esimd_block_read<T, n>(surf_ind, offset);
-#else
-  return __esimd_block_read<T, n>(acc, offset);
-#endif // __SYCL_DEVICE_ONLY__
+  simd<T, n> Res;
+  Res.copy_from(acc, offset);
+  return Res;
 }
 
 /// Flat-address block-store.
 /// \ingroup sycl_esimd
+// TODO the above note about cache hints applies to this API as well.
 template <typename T, int n, CacheHint L1H = CacheHint::None,
           CacheHint L3H = CacheHint::None>
+__SYCL_DEPRECATED("Replaced by simd::copy_to")
 ESIMD_INLINE ESIMD_NODEBUG void block_store(T *p, simd<T, n> vals) {
   constexpr unsigned Sz = sizeof(T) * n;
   static_assert(Sz >= detail::OperandSize::OWORD,
@@ -226,24 +217,10 @@ ESIMD_INLINE ESIMD_NODEBUG void block_store(T *p, simd<T, n> vals) {
 /// Accessor-based block-store.
 /// \ingroup sycl_esimd
 template <typename T, int n, typename AccessorTy>
-ESIMD_INLINE ESIMD_NODEBUG void block_store(AccessorTy acc, uint32_t offset,
-                                            simd<T, n> vals) {
-  constexpr unsigned Sz = sizeof(T) * n;
-  static_assert(Sz >= detail::OperandSize::OWORD,
-                "block size must be at least 1 oword");
-  static_assert(Sz % detail::OperandSize::OWORD == 0,
-                "block size must be whole number of owords");
-  static_assert(detail::isPowerOf2(Sz / detail::OperandSize::OWORD),
-                "block must be 1, 2, 4 or 8 owords long");
-  static_assert(Sz <= 8 * detail::OperandSize::OWORD,
-                "block size must be at most 8 owords");
-
-#if defined(__SYCL_DEVICE_ONLY__)
-  auto surf_ind = detail::AccessorPrivateProxy::getNativeImageObj(acc);
-  __esimd_block_write<T, n>(surf_ind, offset >> 4, vals.data());
-#else
-  __esimd_block_write<T, n>(acc, offset >> 4, vals.data());
-#endif // __SYCL_DEVICE_ONLY__
+__SYCL_DEPRECATED("Replaced by simd::copy_to")
+ESIMD_INLINE ESIMD_NODEBUG
+    void block_store(AccessorTy acc, uint32_t offset, simd<T, n> vals) {
+  vals.copy_to(acc, offset);
 }
 
 /// Accessor-based gather.
diff --git a/sycl/test/esimd/block_load_store.cpp b/sycl/test/esimd/block_load_store.cpp
index 8399eccfe2f85..ff35e72ab7dad 100644
--- a/sycl/test/esimd/block_load_store.cpp
+++ b/sycl/test/esimd/block_load_store.cpp
@@ -1,5 +1,4 @@
-// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
-// expected-no-diagnostics
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
 
 #include <CL/sycl.hpp>
 #include <CL/sycl/INTEL/esimd.hpp>
@@ -9,12 +8,31 @@
 using namespace sycl::INTEL::gpu;
 using namespace cl::sycl;
 
-void kernel(accessor<int, 1, access::mode::read_write, access::target::global_buffer> &buf) __attribute__((sycl_device)) {
-  simd<int, 32> v1(0, 1);
-
-  auto v0 = block_load<int, 32>(buf.get_pointer());
+#ifdef __SYCL_DEVICE_ONLY__
+#define __SYCL_DEVICE_ATTR __attribute__((sycl_device))
+#else
+#define __SYCL_DEVICE_ATTR
+#endif // __SYCL_DEVICE_ONLY__
 
+void kernel1(accessor<int, 1, access::mode::read_write,
+                      access::target::global_buffer> &buf) __SYCL_DEVICE_ATTR {
+  simd<int, 32> v1(0, 1);
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd_memory.hpp:188 {{}}
+  auto v0 = block_load<int, 32>(buf, 0);
   v0 = v0 + v1;
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd_memory.hpp:220 {{}}
+  block_store<int, 32>(buf, 0, v0);
+}
 
-  block_store<int, 32>(buf.get_pointer(), v0);
+void kernel2(int *ptr) __SYCL_DEVICE_ATTR {
+  simd<int, 32> v1(0, 1);
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd_memory.hpp:169 {{}}
+  auto v0 = block_load<int, 32>(ptr);
+  v0 = v0 + v1;
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd_memory.hpp:201 {{}}
+  block_store<int, 32>(ptr, v0);
 }
diff --git a/sycl/test/esimd/simd_copy_to_copy_from.cpp b/sycl/test/esimd/simd_copy_to_copy_from.cpp
new file mode 100644
index 0000000000000..63c39241c12c2
--- /dev/null
+++ b/sycl/test/esimd/simd_copy_to_copy_from.cpp
@@ -0,0 +1,78 @@
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
+
+// This test checks that both host and device compilers can:
+// - successfully compile simd::copy_to and simd::copy_from APIs
+// - emit an error if argument of an incompatible type is used
+//   in place of the accessor argument
+
+#include <CL/sycl.hpp>
+#include <CL/sycl/INTEL/esimd.hpp>
+#include <limits>
+#include <utility>
+
+using namespace sycl::INTEL::gpu;
+using namespace cl::sycl;
+
+#ifdef __SYCL_DEVICE_ONLY__
+#define __SYCL_DEVICE_ATTR __attribute__((sycl_device))
+#else
+#define __SYCL_DEVICE_ATTR
+#endif // __SYCL_DEVICE_ONLY__
+
+// --- Postive tests.
+
+void kernel1(accessor<int, 1, access::mode::read_write,
+                      access::target::global_buffer> &buf) __SYCL_DEVICE_ATTR {
+  simd<int, 32> v1(0, 1);
+  simd<int, 32> v0;
+  v0.copy_from(buf, 0);
+  v0 = v0 + v1;
+  v0.copy_to(buf, 0);
+}
+
+void kernel2(int *ptr) __SYCL_DEVICE_ATTR {
+  simd<int, 32> v1(0, 1);
+  simd<int, 32> v0;
+  v0.copy_from(ptr);
+  v0 = v0 + v1;
+  v0.copy_to(ptr);
+}
+
+// --- Negative tests.
+
+// Incompatible target.
+void kernel3(accessor<int, 1, access::mode::read_write, access::target::local>
+                 &buf) __SYCL_DEVICE_ATTR {
+  simd<int, 32> v1(0, 1);
+  simd<int, 32> v0;
+  // expected-error@+3 {{no matching member function for call to 'copy_from'}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:513 {{}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:508 {{}}
+  v0.copy_from(buf, 0);
+  v0 = v0 + v1;
+  // expected-error@+3 {{no matching member function for call to 'copy_to'}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:497 {{}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:523 {{}}
+  v0.copy_to(buf, 0);
+}
+
+// Incompatible mode (write).
+void kernel4(
+    accessor<int, 1, access::mode::write, access::target::global_buffer> &buf)
+    __SYCL_DEVICE_ATTR {
+  simd<int, 32> v;
+  // expected-error@+3 {{no matching member function for call to 'copy_from'}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:513 {{}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:508 {{}}
+  v.copy_from(buf, 0);
+}
+
+// Incompatible mode (read).
+void kernel5(accessor<int, 1, access::mode::read, access::target::global_buffer>
+                 &buf) __SYCL_DEVICE_ATTR {
+  simd<int, 32> v(0, 1);
+  // expected-error@+3 {{no matching member function for call to 'copy_to'}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:497 {{}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:523 {{}}
+  v.copy_to(buf, 0);
+}

From 8ec7dea9574d8eda28b3553e08b9a751682e1fb2 Mon Sep 17 00:00:00 2001
From: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
Date: Tue, 20 Apr 2021 00:16:48 -0700
Subject: [PATCH 2/7] address review comments

Signed-off-by: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
---
 .../CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp   |  8 +++-----
 sycl/include/CL/sycl/INTEL/esimd/esimd.hpp           | 12 +++++-------
 sycl/test/esimd/block_load_store.cpp                 | 12 +++---------
 3 files changed, 11 insertions(+), 21 deletions(-)

diff --git a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp
index ffed6b9fd7e61..972a39c6d63f6 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp
@@ -74,11 +74,9 @@ struct is_sycl_accessor_with
               (is_sycl_accessor<T>::target == AccessTarget),
           std::true_type, std::false_type> {};
 
-#define __ESIMD_ENABLE_IF_ACCESSOR(T, acc_capability, acc_target, ret_type)    \
-  sycl::detail::enable_if_t<detail::is_sycl_accessor_with<                     \
-                                T, detail::accessor_mode_cap::acc_capability,  \
-                                sycl::access::target::acc_target>::value,      \
-                            ret_type>
+template <typename T, accessor_mode_cap_val_t Capability, sycl::access::target AccessTarget, typename RetT>
+using EnableIfAccessor =
+sycl::detail::enable_if_t<detail::is_sycl_accessor_with<T, Capability, AccessTarget>::value, RetT>;
 
 } // namespace detail
 } // namespace gpu
diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
index 62ad569a6efd2..5e3267a269d61 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
@@ -503,8 +503,7 @@ template <typename Ty, int N> class simd {
   /// @param acc accessor to copy from.
   /// @param offset offset to copy from.
   template <typename AccessorT>
-  ESIMD_INLINE __ESIMD_ENABLE_IF_ACCESSOR(AccessorT, can_read, global_buffer,
-                                          void)
+  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read, sycl::access::target::global_buffer, void>
       copy_from(AccessorT acc, uint32_t offset) SYCL_ESIMD_FUNCTION;
 
   /// Copy all vector elements of this object into a contiguous block in memory.
@@ -518,9 +517,8 @@ template <typename Ty, int N> class simd {
   /// @param acc accessor to copy from.
   /// @param offset offset to copy from.
   template <typename AccessorT>
-  ESIMD_INLINE __ESIMD_ENABLE_IF_ACCESSOR(AccessorT, can_write, global_buffer,
-                                          void)
-      copy_to(AccessorT acc, uint32_t offset) SYCL_ESIMD_FUNCTION;
+  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write, sycl::access::target::global_buffer, void>
+    copy_to(AccessorT acc, uint32_t offset) SYCL_ESIMD_FUNCTION;
 
   /// @} // Memory operations
 private:
@@ -562,7 +560,7 @@ template <typename T, int N> void simd<T, N>::copy_from(const T *const addr) {
 
 template <typename T, int N>
 template <typename AccessorT>
-__ESIMD_ENABLE_IF_ACCESSOR(AccessorT, can_read, global_buffer, void)
+ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read, sycl::access::target::global_buffer, void>
 simd<T, N>::copy_from(AccessorT acc, uint32_t offset) {
   constexpr unsigned Sz = sizeof(T) * N;
   static_assert(Sz >= detail::OperandSize::OWORD,
@@ -599,7 +597,7 @@ template <typename T, int N> void simd<T, N>::copy_to(T *addr) {
 
 template <typename T, int N>
 template <typename AccessorT>
-__ESIMD_ENABLE_IF_ACCESSOR(AccessorT, can_write, global_buffer, void)
+ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write, sycl::access::target::global_buffer, void>
 simd<T, N>::copy_to(AccessorT acc, uint32_t offset) {
   constexpr unsigned Sz = sizeof(T) * N;
   static_assert(Sz >= detail::OperandSize::OWORD,
diff --git a/sycl/test/esimd/block_load_store.cpp b/sycl/test/esimd/block_load_store.cpp
index ff35e72ab7dad..7cd20484af42e 100644
--- a/sycl/test/esimd/block_load_store.cpp
+++ b/sycl/test/esimd/block_load_store.cpp
@@ -8,14 +8,8 @@
 using namespace sycl::INTEL::gpu;
 using namespace cl::sycl;
 
-#ifdef __SYCL_DEVICE_ONLY__
-#define __SYCL_DEVICE_ATTR __attribute__((sycl_device))
-#else
-#define __SYCL_DEVICE_ATTR
-#endif // __SYCL_DEVICE_ONLY__
-
-void kernel1(accessor<int, 1, access::mode::read_write,
-                      access::target::global_buffer> &buf) __SYCL_DEVICE_ATTR {
+SYCL_EXTERNAL void kernel1(accessor<int, 1, access::mode::read_write,
+                      access::target::global_buffer> &buf) SYCL_ESIMD_FUNCTION {
   simd<int, 32> v1(0, 1);
   // expected-warning@+2 {{deprecated}}
   // expected-note@CL/sycl/INTEL/esimd/esimd_memory.hpp:188 {{}}
@@ -26,7 +20,7 @@ void kernel1(accessor<int, 1, access::mode::read_write,
   block_store<int, 32>(buf, 0, v0);
 }
 
-void kernel2(int *ptr) __SYCL_DEVICE_ATTR {
+SYCL_EXTERNAL void kernel2(int *ptr) SYCL_ESIMD_FUNCTION {
   simd<int, 32> v1(0, 1);
   // expected-warning@+2 {{deprecated}}
   // expected-note@CL/sycl/INTEL/esimd/esimd_memory.hpp:169 {{}}

From 5f0e1d0ab20c9dbab7542d71eebc93652de84add Mon Sep 17 00:00:00 2001
From: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
Date: Tue, 20 Apr 2021 00:22:49 -0700
Subject: [PATCH 3/7] clang format

Signed-off-by: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
---
 .../sycl/INTEL/esimd/detail/esimd_sycl_util.hpp  |  7 ++++---
 sycl/include/CL/sycl/INTEL/esimd/esimd.hpp       | 16 ++++++++++------
 sycl/test/esimd/block_load_store.cpp             |  5 +++--
 3 files changed, 17 insertions(+), 11 deletions(-)

diff --git a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp
index 972a39c6d63f6..043b40cfa08a7 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_sycl_util.hpp
@@ -74,9 +74,10 @@ struct is_sycl_accessor_with
               (is_sycl_accessor<T>::target == AccessTarget),
           std::true_type, std::false_type> {};
 
-template <typename T, accessor_mode_cap_val_t Capability, sycl::access::target AccessTarget, typename RetT>
-using EnableIfAccessor =
-sycl::detail::enable_if_t<detail::is_sycl_accessor_with<T, Capability, AccessTarget>::value, RetT>;
+template <typename T, accessor_mode_cap_val_t Capability,
+          sycl::access::target AccessTarget, typename RetT>
+using EnableIfAccessor = sycl::detail::enable_if_t<
+    detail::is_sycl_accessor_with<T, Capability, AccessTarget>::value, RetT>;
 
 } // namespace detail
 } // namespace gpu
diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
index 5e3267a269d61..5970e622917a5 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
@@ -503,8 +503,9 @@ template <typename Ty, int N> class simd {
   /// @param acc accessor to copy from.
   /// @param offset offset to copy from.
   template <typename AccessorT>
-  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read, sycl::access::target::global_buffer, void>
-      copy_from(AccessorT acc, uint32_t offset) SYCL_ESIMD_FUNCTION;
+  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read,
+                                sycl::access::target::global_buffer, void>
+  copy_from(AccessorT acc, uint32_t offset) SYCL_ESIMD_FUNCTION;
 
   /// Copy all vector elements of this object into a contiguous block in memory.
   /// @param addr the memory address to copy to. Must be a pointer to the
@@ -517,8 +518,9 @@ template <typename Ty, int N> class simd {
   /// @param acc accessor to copy from.
   /// @param offset offset to copy from.
   template <typename AccessorT>
-  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write, sycl::access::target::global_buffer, void>
-    copy_to(AccessorT acc, uint32_t offset) SYCL_ESIMD_FUNCTION;
+  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write,
+                                sycl::access::target::global_buffer, void>
+  copy_to(AccessorT acc, uint32_t offset) SYCL_ESIMD_FUNCTION;
 
   /// @} // Memory operations
 private:
@@ -560,7 +562,8 @@ template <typename T, int N> void simd<T, N>::copy_from(const T *const addr) {
 
 template <typename T, int N>
 template <typename AccessorT>
-ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read, sycl::access::target::global_buffer, void>
+ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read,
+                              sycl::access::target::global_buffer, void>
 simd<T, N>::copy_from(AccessorT acc, uint32_t offset) {
   constexpr unsigned Sz = sizeof(T) * N;
   static_assert(Sz >= detail::OperandSize::OWORD,
@@ -597,7 +600,8 @@ template <typename T, int N> void simd<T, N>::copy_to(T *addr) {
 
 template <typename T, int N>
 template <typename AccessorT>
-ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write, sycl::access::target::global_buffer, void>
+ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write,
+                              sycl::access::target::global_buffer, void>
 simd<T, N>::copy_to(AccessorT acc, uint32_t offset) {
   constexpr unsigned Sz = sizeof(T) * N;
   static_assert(Sz >= detail::OperandSize::OWORD,
diff --git a/sycl/test/esimd/block_load_store.cpp b/sycl/test/esimd/block_load_store.cpp
index 7cd20484af42e..1b0da0341b2a7 100644
--- a/sycl/test/esimd/block_load_store.cpp
+++ b/sycl/test/esimd/block_load_store.cpp
@@ -8,8 +8,9 @@
 using namespace sycl::INTEL::gpu;
 using namespace cl::sycl;
 
-SYCL_EXTERNAL void kernel1(accessor<int, 1, access::mode::read_write,
-                      access::target::global_buffer> &buf) SYCL_ESIMD_FUNCTION {
+SYCL_EXTERNAL void kernel1(
+    accessor<int, 1, access::mode::read_write, access::target::global_buffer>
+        &buf) SYCL_ESIMD_FUNCTION {
   simd<int, 32> v1(0, 1);
   // expected-warning@+2 {{deprecated}}
   // expected-note@CL/sycl/INTEL/esimd/esimd_memory.hpp:188 {{}}

From 00bf14f12cac3df8aa4ee172cd0a405789e38181 Mon Sep 17 00:00:00 2001
From: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
Date: Tue, 20 Apr 2021 00:29:11 -0700
Subject: [PATCH 4/7] more review comments

Signed-off-by: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
---
 sycl/test/esimd/simd_copy_to_copy_from.cpp | 27 ++++++++++------------
 1 file changed, 12 insertions(+), 15 deletions(-)

diff --git a/sycl/test/esimd/simd_copy_to_copy_from.cpp b/sycl/test/esimd/simd_copy_to_copy_from.cpp
index 63c39241c12c2..90531f1c73203 100644
--- a/sycl/test/esimd/simd_copy_to_copy_from.cpp
+++ b/sycl/test/esimd/simd_copy_to_copy_from.cpp
@@ -13,16 +13,11 @@
 using namespace sycl::INTEL::gpu;
 using namespace cl::sycl;
 
-#ifdef __SYCL_DEVICE_ONLY__
-#define __SYCL_DEVICE_ATTR __attribute__((sycl_device))
-#else
-#define __SYCL_DEVICE_ATTR
-#endif // __SYCL_DEVICE_ONLY__
-
 // --- Postive tests.
 
-void kernel1(accessor<int, 1, access::mode::read_write,
-                      access::target::global_buffer> &buf) __SYCL_DEVICE_ATTR {
+SYCL_EXTERNAL void kernel1(
+    accessor<int, 1, access::mode::read_write, access::target::global_buffer>
+        &buf) SYCL_ESIMD_FUNCTION {
   simd<int, 32> v1(0, 1);
   simd<int, 32> v0;
   v0.copy_from(buf, 0);
@@ -30,7 +25,7 @@ void kernel1(accessor<int, 1, access::mode::read_write,
   v0.copy_to(buf, 0);
 }
 
-void kernel2(int *ptr) __SYCL_DEVICE_ATTR {
+SYCL_EXTERNAL void kernel2(int *ptr) SYCL_ESIMD_FUNCTION {
   simd<int, 32> v1(0, 1);
   simd<int, 32> v0;
   v0.copy_from(ptr);
@@ -41,8 +36,9 @@ void kernel2(int *ptr) __SYCL_DEVICE_ATTR {
 // --- Negative tests.
 
 // Incompatible target.
-void kernel3(accessor<int, 1, access::mode::read_write, access::target::local>
-                 &buf) __SYCL_DEVICE_ATTR {
+SYCL_EXTERNAL void
+kernel3(accessor<int, 1, access::mode::read_write, access::target::local> &buf)
+    SYCL_ESIMD_FUNCTION {
   simd<int, 32> v1(0, 1);
   simd<int, 32> v0;
   // expected-error@+3 {{no matching member function for call to 'copy_from'}}
@@ -57,9 +53,9 @@ void kernel3(accessor<int, 1, access::mode::read_write, access::target::local>
 }
 
 // Incompatible mode (write).
-void kernel4(
+SYCL_EXTERNAL void kernel4(
     accessor<int, 1, access::mode::write, access::target::global_buffer> &buf)
-    __SYCL_DEVICE_ATTR {
+    SYCL_ESIMD_FUNCTION {
   simd<int, 32> v;
   // expected-error@+3 {{no matching member function for call to 'copy_from'}}
   // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:513 {{}}
@@ -68,8 +64,9 @@ void kernel4(
 }
 
 // Incompatible mode (read).
-void kernel5(accessor<int, 1, access::mode::read, access::target::global_buffer>
-                 &buf) __SYCL_DEVICE_ATTR {
+SYCL_EXTERNAL void kernel5(
+    accessor<int, 1, access::mode::read, access::target::global_buffer> &buf)
+    SYCL_ESIMD_FUNCTION {
   simd<int, 32> v(0, 1);
   // expected-error@+3 {{no matching member function for call to 'copy_to'}}
   // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:497 {{}}

From 6f6752cf127002c6707175ca2fc385638cf0c072 Mon Sep 17 00:00:00 2001
From: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
Date: Mon, 3 May 2021 23:02:05 -0700
Subject: [PATCH 5/7] Review comments.

Signed-off-by: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
---
 sycl/include/CL/sycl/INTEL/esimd/esimd.hpp        | 4 ++--
 sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp | 8 ++++----
 2 files changed, 6 insertions(+), 6 deletions(-)

diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
index 5970e622917a5..91a3d1301d2c2 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
@@ -543,7 +543,7 @@ ESIMD_INLINE simd<U, n> convert(simd<T, n> val) {
 
 // ----------- Outlined implementations of esimd class APIs.
 
-template <typename T, int N> void simd<T, N>::copy_from(const T *const addr) {
+template <typename T, int N> void simd<T, N>::copy_from(const T *const Addr) {
   constexpr unsigned Sz = sizeof(T) * N;
   static_assert(Sz >= detail::OperandSize::OWORD,
                 "block size must be at least 1 oword");
@@ -554,7 +554,7 @@ template <typename T, int N> void simd<T, N>::copy_from(const T *const addr) {
   static_assert(Sz <= 8 * detail::OperandSize::OWORD,
                 "block size must be at most 8 owords");
 
-  uintptr_t AddrVal = reinterpret_cast<uintptr_t>(addr);
+  uintptr_t AddrVal = reinterpret_cast<uintptr_t>(Addr);
   *this =
       __esimd_flat_block_read_unaligned<T, N, CacheHint::None, CacheHint::None>(
           AddrVal);
diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp
index 9160323adbed7..f3ae3430d2d63 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp
@@ -166,7 +166,7 @@ scatter(T *p, simd<T, n * ElemsPerAddr> vals, simd<uint32_t, n> offsets,
 // This API, even though deprecated, can't be removed until then.
 template <typename T, int n, CacheHint L1H = CacheHint::None,
           CacheHint L3H = CacheHint::None>
-__SYCL_DEPRECATED("Replaced by simd::copy_from")
+__SYCL_DEPRECATED("block_load is deprecated, use simd::copy_from.")
 ESIMD_INLINE ESIMD_NODEBUG simd<T, n> block_load(const T *const addr) {
   constexpr unsigned Sz = sizeof(T) * n;
   static_assert(Sz >= detail::OperandSize::OWORD,
@@ -185,7 +185,7 @@ ESIMD_INLINE ESIMD_NODEBUG simd<T, n> block_load(const T *const addr) {
 /// Accessor-based block-load.
 /// \ingroup sycl_esimd
 template <typename T, int n, typename AccessorTy>
-__SYCL_DEPRECATED("Replaced by simd::copy_from")
+__SYCL_DEPRECATED("block_load is deprecated, use simd::copy_from.")
 ESIMD_INLINE ESIMD_NODEBUG simd<T, n> block_load(AccessorTy acc,
                                                  uint32_t offset) {
   simd<T, n> Res;
@@ -198,7 +198,7 @@ ESIMD_INLINE ESIMD_NODEBUG simd<T, n> block_load(AccessorTy acc,
 // TODO the above note about cache hints applies to this API as well.
 template <typename T, int n, CacheHint L1H = CacheHint::None,
           CacheHint L3H = CacheHint::None>
-__SYCL_DEPRECATED("Replaced by simd::copy_to")
+__SYCL_DEPRECATED("block_store is deprecated, use simd::copy_to.")
 ESIMD_INLINE ESIMD_NODEBUG void block_store(T *p, simd<T, n> vals) {
   constexpr unsigned Sz = sizeof(T) * n;
   static_assert(Sz >= detail::OperandSize::OWORD,
@@ -217,7 +217,7 @@ ESIMD_INLINE ESIMD_NODEBUG void block_store(T *p, simd<T, n> vals) {
 /// Accessor-based block-store.
 /// \ingroup sycl_esimd
 template <typename T, int n, typename AccessorTy>
-__SYCL_DEPRECATED("Replaced by simd::copy_to")
+__SYCL_DEPRECATED("block_store is deprecated, use simd::copy_to.")
 ESIMD_INLINE ESIMD_NODEBUG
     void block_store(AccessorTy acc, uint32_t offset, simd<T, n> vals) {
   vals.copy_to(acc, offset);

From 9a4c3b63f3da7376f0edc1fb81a9fac4e5a87ccf Mon Sep 17 00:00:00 2001
From: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
Date: Tue, 4 May 2021 10:49:08 -0700
Subject: [PATCH 6/7] fix test failures, remove duplicated string from
 deprecation message

Signed-off-by: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
---
 sycl/include/CL/sycl/INTEL/esimd/esimd.hpp    | 28 +++++++++++--------
 .../CL/sycl/INTEL/esimd/esimd_memory.hpp      |  8 +++---
 2 files changed, 20 insertions(+), 16 deletions(-)

diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
index 91a3d1301d2c2..ca33a35949f96 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
@@ -503,9 +503,10 @@ template <typename Ty, int N> class simd {
   /// @param acc accessor to copy from.
   /// @param offset offset to copy from.
   template <typename AccessorT>
-  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read,
-                                sycl::access::target::global_buffer, void>
-  copy_from(AccessorT acc, uint32_t offset) SYCL_ESIMD_FUNCTION;
+  ESIMD_INLINE
+      detail::EnableIfAccessor<AccessorT, detail::accessor_mode_cap::can_read,
+                               sycl::access::target::global_buffer, void>
+      copy_from(AccessorT acc, uint32_t offset) SYCL_ESIMD_FUNCTION;
 
   /// Copy all vector elements of this object into a contiguous block in memory.
   /// @param addr the memory address to copy to. Must be a pointer to the
@@ -518,9 +519,10 @@ template <typename Ty, int N> class simd {
   /// @param acc accessor to copy from.
   /// @param offset offset to copy from.
   template <typename AccessorT>
-  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write,
-                                sycl::access::target::global_buffer, void>
-  copy_to(AccessorT acc, uint32_t offset) SYCL_ESIMD_FUNCTION;
+  ESIMD_INLINE
+      detail::EnableIfAccessor<AccessorT, detail::accessor_mode_cap::can_write,
+                               sycl::access::target::global_buffer, void>
+      copy_to(AccessorT acc, uint32_t offset) SYCL_ESIMD_FUNCTION;
 
   /// @} // Memory operations
 private:
@@ -562,9 +564,10 @@ template <typename T, int N> void simd<T, N>::copy_from(const T *const Addr) {
 
 template <typename T, int N>
 template <typename AccessorT>
-ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read,
-                              sycl::access::target::global_buffer, void>
-simd<T, N>::copy_from(AccessorT acc, uint32_t offset) {
+ESIMD_INLINE
+    detail::EnableIfAccessor<AccessorT, detail::accessor_mode_cap::can_read,
+                             sycl::access::target::global_buffer, void>
+    simd<T, N>::copy_from(AccessorT acc, uint32_t offset) {
   constexpr unsigned Sz = sizeof(T) * N;
   static_assert(Sz >= detail::OperandSize::OWORD,
                 "block size must be at least 1 oword");
@@ -600,9 +603,10 @@ template <typename T, int N> void simd<T, N>::copy_to(T *addr) {
 
 template <typename T, int N>
 template <typename AccessorT>
-ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write,
-                              sycl::access::target::global_buffer, void>
-simd<T, N>::copy_to(AccessorT acc, uint32_t offset) {
+ESIMD_INLINE
+    detail::EnableIfAccessor<AccessorT, detail::accessor_mode_cap::can_write,
+                             sycl::access::target::global_buffer, void>
+    simd<T, N>::copy_to(AccessorT acc, uint32_t offset) {
   constexpr unsigned Sz = sizeof(T) * N;
   static_assert(Sz >= detail::OperandSize::OWORD,
                 "block size must be at least 1 oword");
diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp
index f3ae3430d2d63..1c9e991a46fe4 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp
@@ -166,7 +166,7 @@ scatter(T *p, simd<T, n * ElemsPerAddr> vals, simd<uint32_t, n> offsets,
 // This API, even though deprecated, can't be removed until then.
 template <typename T, int n, CacheHint L1H = CacheHint::None,
           CacheHint L3H = CacheHint::None>
-__SYCL_DEPRECATED("block_load is deprecated, use simd::copy_from.")
+__SYCL_DEPRECATED("use simd::copy_from.")
 ESIMD_INLINE ESIMD_NODEBUG simd<T, n> block_load(const T *const addr) {
   constexpr unsigned Sz = sizeof(T) * n;
   static_assert(Sz >= detail::OperandSize::OWORD,
@@ -185,7 +185,7 @@ ESIMD_INLINE ESIMD_NODEBUG simd<T, n> block_load(const T *const addr) {
 /// Accessor-based block-load.
 /// \ingroup sycl_esimd
 template <typename T, int n, typename AccessorTy>
-__SYCL_DEPRECATED("block_load is deprecated, use simd::copy_from.")
+__SYCL_DEPRECATED("use simd::copy_from.")
 ESIMD_INLINE ESIMD_NODEBUG simd<T, n> block_load(AccessorTy acc,
                                                  uint32_t offset) {
   simd<T, n> Res;
@@ -198,7 +198,7 @@ ESIMD_INLINE ESIMD_NODEBUG simd<T, n> block_load(AccessorTy acc,
 // TODO the above note about cache hints applies to this API as well.
 template <typename T, int n, CacheHint L1H = CacheHint::None,
           CacheHint L3H = CacheHint::None>
-__SYCL_DEPRECATED("block_store is deprecated, use simd::copy_to.")
+__SYCL_DEPRECATED("use simd::copy_to.")
 ESIMD_INLINE ESIMD_NODEBUG void block_store(T *p, simd<T, n> vals) {
   constexpr unsigned Sz = sizeof(T) * n;
   static_assert(Sz >= detail::OperandSize::OWORD,
@@ -217,7 +217,7 @@ ESIMD_INLINE ESIMD_NODEBUG void block_store(T *p, simd<T, n> vals) {
 /// Accessor-based block-store.
 /// \ingroup sycl_esimd
 template <typename T, int n, typename AccessorTy>
-__SYCL_DEPRECATED("block_store is deprecated, use simd::copy_to.")
+__SYCL_DEPRECATED("use simd::copy_to.")
 ESIMD_INLINE ESIMD_NODEBUG
     void block_store(AccessorTy acc, uint32_t offset, simd<T, n> vals) {
   vals.copy_to(acc, offset);

From 72cddd14c4a79a1ac42a4ddc9c223c314b3abe09 Mon Sep 17 00:00:00 2001
From: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
Date: Tue, 4 May 2021 11:29:24 -0700
Subject: [PATCH 7/7] Fix line numbers in negative tests.

Signed-off-by: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
---
 sycl/test/esimd/simd_copy_to_copy_from.cpp | 12 ++++++------
 1 file changed, 6 insertions(+), 6 deletions(-)

diff --git a/sycl/test/esimd/simd_copy_to_copy_from.cpp b/sycl/test/esimd/simd_copy_to_copy_from.cpp
index 90531f1c73203..fd465ff068c0e 100644
--- a/sycl/test/esimd/simd_copy_to_copy_from.cpp
+++ b/sycl/test/esimd/simd_copy_to_copy_from.cpp
@@ -42,13 +42,13 @@ kernel3(accessor<int, 1, access::mode::read_write, access::target::local> &buf)
   simd<int, 32> v1(0, 1);
   simd<int, 32> v0;
   // expected-error@+3 {{no matching member function for call to 'copy_from'}}
-  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:513 {{}}
-  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:508 {{}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:514 {{}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:509 {{}}
   v0.copy_from(buf, 0);
   v0 = v0 + v1;
   // expected-error@+3 {{no matching member function for call to 'copy_to'}}
   // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:497 {{}}
-  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:523 {{}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:525 {{}}
   v0.copy_to(buf, 0);
 }
 
@@ -58,8 +58,8 @@ SYCL_EXTERNAL void kernel4(
     SYCL_ESIMD_FUNCTION {
   simd<int, 32> v;
   // expected-error@+3 {{no matching member function for call to 'copy_from'}}
-  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:513 {{}}
-  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:508 {{}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:514 {{}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:509 {{}}
   v.copy_from(buf, 0);
 }
 
@@ -70,6 +70,6 @@ SYCL_EXTERNAL void kernel5(
   simd<int, 32> v(0, 1);
   // expected-error@+3 {{no matching member function for call to 'copy_to'}}
   // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:497 {{}}
-  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:523 {{}}
+  // expected-note@CL/sycl/INTEL/esimd/esimd.hpp:525 {{}}
   v.copy_to(buf, 0);
 }