fix test failures, remove duplicated string from deprecation message

Signed-off-by: kbobrovs <[email protected]>

Author: kbobrovs

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");

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);