address review comments

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

Author: kbobrovs

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

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,

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 {{}}