[SYCL][ESIMD] Add more stringent compile time checks for accessor versions of block_load/block_store, gather/scatter API

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

@@ -330,13 +330,11 @@ class [[__sycl_detail__::__uses_aspects__(
   ///   argument.
   /// @param acc The accessor to read from.
   /// @param offset offset in bytes of the first element.
-  template <typename AccessorT, typename Flags = element_aligned_tag,
-            typename = std::enable_if_t<
-                (sycl::detail::acc_properties::is_local_accessor_v<AccessorT> ||
-                 detail::is_sycl_accessor_with<
-                     AccessorT, accessor_mode_cap::can_read,
-                     sycl::access::target::device>::value) &&
-                is_simd_flag_type_v<Flags>>>
+  template <
+      typename AccessorT, typename Flags = element_aligned_tag,
+      typename = std::enable_if_t<
+          detail::is_accessor_with_v<AccessorT, accessor_mode_cap::can_read> &&
+          is_simd_flag_type_v<Flags>>>
   simd_obj_impl(AccessorT acc,
 #ifdef __ESIMD_FORCE_STATELESS_MEM
                 uint64_t offset,
@@ -744,8 +742,7 @@ class [[__sycl_detail__::__uses_aspects__(
   template <typename AccessorT, typename Flags = element_aligned_tag,
             int ChunkSize = 32,
             typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
-  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read,
-                                sycl::access::target::device, void>
+  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read, void>
   copy_from(AccessorT acc,
 #ifdef __ESIMD_FORCE_STATELESS_MEM
             uint64_t offset,
@@ -754,6 +751,25 @@ class [[__sycl_detail__::__uses_aspects__(
 #endif
             Flags = {}) SYCL_ESIMD_FUNCTION;
 
+  /// Copy a contiguous block of data from memory into this simd_obj_impl
+  /// object. The amount of memory copied equals the total size of vector
+  /// elements in this object. Source memory location is represented via a
+  /// local accessor and offset.
+  /// None of the template parameters except documented ones can/should be
+  /// specified by callers.
+  /// @tparam AccessorT Type of the accessor (auto-deduced).
+  /// @tparam Flags Alignment control for the copy operation.
+  ///   See @ref sycl_esimd_core_align for more info.
+  /// @param acc accessor to copy from.
+  /// @param offset offset to copy from (in bytes).
+  template <typename AccessorT, typename Flags = element_aligned_tag,
+            int ChunkSize = 32,
+            typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
+  ESIMD_INLINE std::enable_if_t<
+      detail::is_local_accessor_with_v<AccessorT, accessor_mode_cap::can_read>,
+      void>
+  copy_from(AccessorT acc, uint32_t offset, Flags = {}) SYCL_ESIMD_FUNCTION;
+
   /// Copy all vector elements of this object into a contiguous block in memory.
   /// None of the template parameters should be be specified by callers.
   /// @tparam Flags Alignment control for the copy operation.
@@ -776,8 +792,7 @@ class [[__sycl_detail__::__uses_aspects__(
   template <typename AccessorT, typename Flags = element_aligned_tag,
             int ChunkSize = 32,
             typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
-  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write,
-                                sycl::access::target::device, void>
+  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write, void>
   copy_to(AccessorT acc,
 #ifdef __ESIMD_FORCE_STATELESS_MEM
           uint64_t offset,
@@ -786,6 +801,23 @@ class [[__sycl_detail__::__uses_aspects__(
 #endif
           Flags = {}) const SYCL_ESIMD_FUNCTION;
 
+  /// Copy all vector elements of this object into a contiguous block in memory.
+  /// Destination memory location is represented via a local accessor and
+  /// offset.
+  /// None of the template parameters should be be specified by callers.
+  /// @tparam AccessorT Type of the accessor (auto-deduced).
+  /// @tparam Flags Alignment control for the copy operation.
+  ///   See @ref sycl_esimd_core_align for more info.
+  /// @param acc accessor to copy from.
+  /// @param offset offset to copy from.
+  template <typename AccessorT, typename Flags = element_aligned_tag,
+            int ChunkSize = 32,
+            typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
+  ESIMD_INLINE std::enable_if_t<
+      detail::is_local_accessor_with_v<AccessorT, accessor_mode_cap::can_write>,
+      void>
+  copy_to(AccessorT acc, uint32_t offset, Flags = {}) const SYCL_ESIMD_FUNCTION;
+
   // Unary operations.
 
   /// Per-element bitwise inversion, available in all subclasses, but only for
@@ -916,6 +948,13 @@ class [[__sycl_detail__::__uses_aspects__(
   // The underlying data for this vector.
   raw_vector_type M_data;
 
+  template <int ChunkSize, typename Flags, typename AccessorT, typename TOffset>
+  ESIMD_INLINE void copy_to_impl(AccessorT acc,
+                                 TOffset offset) const SYCL_ESIMD_FUNCTION;
+  template <int ChunkSize, typename Flags, typename AccessorT, typename TOffset>
+  ESIMD_INLINE void copy_from_impl(AccessorT acc,
+                                   TOffset offset) SYCL_ESIMD_FUNCTION;
+
 protected:
   // The test proxy if enabled
   __ESIMD_DECLARE_TEST_PROXY

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

@@ -63,23 +63,51 @@ constexpr bool accessor_mode_has_capability() {
          (Mode == sycl::access::mode::read);
 }
 
+template <typename T> struct local_accessor_access_mode {
+  static constexpr sycl::access::mode mode =
+      static_cast<sycl::access::mode>(-1);
+};
+
+template <typename DataT, int Dimensions>
+struct local_accessor_access_mode<local_accessor<DataT, Dimensions>> {
+  static constexpr sycl::access::mode mode =
+      sycl::detail::accessModeFromConstness<DataT>();
+};
+
 // 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
+template <typename T, accessor_mode_cap_val_t Capability>
+struct is_device_accessor_with
     : public std::conditional_t<
           accessor_mode_has_capability<is_sycl_accessor<T>::mode,
                                        Capability>() &&
-              (is_sycl_accessor<T>::target == AccessTarget),
+              (is_sycl_accessor<T>::target == sycl::access::target::device),
           std::true_type, std::false_type> {};
 
-template <typename T, accessor_mode_cap_val_t Capability,
-          sycl::access::target AccessTarget, typename RetT>
-using EnableIfAccessor = std::enable_if_t<
-    detail::is_sycl_accessor_with<T, Capability, AccessTarget>::value ||
-        sycl::detail::acc_properties::is_local_accessor_v<T>,
-    RetT>;
+template <typename T, accessor_mode_cap_val_t Capability>
+struct is_local_accessor_with
+    : public std::conditional_t<
+          sycl::detail::acc_properties::is_local_accessor_v<T> &&
+              accessor_mode_has_capability<local_accessor_access_mode<T>::mode,
+                                           Capability>(),
+          std::true_type, std::false_type> {};
+
+template <typename T, accessor_mode_cap_val_t Capability>
+inline constexpr bool is_local_accessor_with_v =
+    is_local_accessor_with<T, Capability>::value;
+
+template <typename T, accessor_mode_cap_val_t Capability>
+inline constexpr bool is_device_accessor_with_v =
+    is_device_accessor_with<T, Capability>::value;
+
+template <typename T, accessor_mode_cap_val_t Capability>
+inline constexpr bool is_accessor_with_v =
+    is_device_accessor_with_v<T, Capability> ||
+    is_local_accessor_with_v<T, Capability>;
+
+template <typename T, accessor_mode_cap_val_t Capability, typename RetT>
+using EnableIfAccessor =
+    std::enable_if_t<detail::is_device_accessor_with_v<T, Capability>, RetT>;
 
 template <typename T, int Dimensions>
 __ESIMD_API uint32_t localAccessorToOffset(local_accessor<T, Dimensions> acc) {