diff --git a/sycl/include/sycl/ext/intel/esimd/memory.hpp b/sycl/include/sycl/ext/intel/esimd/memory.hpp index 77a00108f0e2..36773ed1355f 100644 --- a/sycl/include/sycl/ext/intel/esimd/memory.hpp +++ b/sycl/include/sycl/ext/intel/esimd/memory.hpp @@ -3057,11 +3057,12 @@ __ESIMD_API void media_block_store(AccessorTy acc, unsigned x, unsigned y, /// template > -__ESIMD_API std::enable_if_t< - sycl::detail::acc_properties::is_local_accessor_v && - is_simd_flag_type_v, - simd> -block_load(AccessorTy acc, uint32_t offset, Flags = {}) { +__ESIMD_API + std::enable_if_t && + is_simd_flag_type_v, + simd> + block_load(AccessorTy acc, uint32_t offset, Flags = {}) { return slm_block_load(offset + __ESIMD_DNS::localAccessorToOffset(acc)); } @@ -3085,10 +3086,11 @@ block_load(AccessorTy acc, uint32_t offset, Flags = {}) { /// template > -__ESIMD_API std::enable_if_t< - sycl::detail::acc_properties::is_local_accessor_v && - is_simd_flag_type_v> -block_store(AccessorTy acc, uint32_t offset, simd vals, Flags = {}) { +__ESIMD_API + std::enable_if_t && + is_simd_flag_type_v> + block_store(AccessorTy acc, uint32_t offset, simd vals, Flags = {}) { slm_block_store( offset + __ESIMD_DNS::localAccessorToOffset(acc), vals); } @@ -3111,10 +3113,12 @@ block_store(AccessorTy acc, uint32_t offset, simd vals, Flags = {}) { /// undefined. /// template -__ESIMD_API std::enable_if_t< - sycl::detail::acc_properties::is_local_accessor_v, simd> -gather(AccessorTy acc, simd offsets, uint32_t glob_offset = 0, - simd_mask mask = 1) { +__ESIMD_API + std::enable_if_t, + simd> + gather(AccessorTy acc, simd offsets, uint32_t glob_offset = 0, + simd_mask mask = 1) { return slm_gather( offsets + glob_offset + __ESIMD_DNS::localAccessorToOffset(acc), mask); } @@ -3138,8 +3142,8 @@ gather(AccessorTy acc, simd offsets, uint32_t glob_offset = 0, /// /// template -__ESIMD_API std::enable_if_t< - sycl::detail::acc_properties::is_local_accessor_v> +__ESIMD_API std::enable_if_t> scatter(AccessorTy acc, simd offsets, simd vals, uint32_t glob_offset = 0, simd_mask mask = 1) { slm_scatter(offsets + glob_offset + @@ -3174,11 +3178,12 @@ scatter(AccessorTy acc, simd offsets, simd vals, template -__ESIMD_API std::enable_if_t< - sycl::detail::acc_properties::is_local_accessor_v, - simd> -gather_rgba(AccessorT acc, simd offsets, - uint32_t global_offset = 0, simd_mask mask = 1) { +__ESIMD_API + std::enable_if_t, + simd> + gather_rgba(AccessorT acc, simd offsets, + uint32_t global_offset = 0, simd_mask mask = 1) { return slm_gather_rgba( offsets + global_offset + __ESIMD_DNS::localAccessorToOffset(acc), mask); } @@ -3202,8 +3207,8 @@ gather_rgba(AccessorT acc, simd offsets, template -__ESIMD_API std::enable_if_t< - sycl::detail::acc_properties::is_local_accessor_v> +__ESIMD_API std::enable_if_t> scatter_rgba(AccessorT acc, simd offsets, simd vals, uint32_t global_offset = 0, simd_mask mask = 1) { diff --git a/sycl/test/esimd/block_load_store.cpp b/sycl/test/esimd/block_load_store.cpp index 0fbe46b9235a..79b9ab73f80b 100644 --- a/sycl/test/esimd/block_load_store.cpp +++ b/sycl/test/esimd/block_load_store.cpp @@ -30,7 +30,7 @@ SYCL_EXTERNAL void kernel2(int *ptr) SYCL_ESIMD_FUNCTION { // Incompatible mode (write). SYCL_EXTERNAL void -kernel3(accessor &buf) +kernel4(accessor &buf) SYCL_ESIMD_FUNCTION { simd v; // CHECK: block_load_store.cpp:38{{.*}}error: no matching function @@ -40,10 +40,19 @@ kernel3(accessor &buf) // Incompatible mode (read). SYCL_EXTERNAL void -kernel4(accessor &buf) +kernel5(accessor &buf) SYCL_ESIMD_FUNCTION { simd v(0, 1); // CHECK: block_load_store.cpp:48{{.*}}error: no matching function // function for call to 'block_store' block_store(buf, 0, v); } + +// Incompatible mode (read). +SYCL_EXTERNAL void +kernel6(local_accessor &buf) SYCL_ESIMD_FUNCTION { + simd v(0, 1); + // CHECK: block_load_store.cpp:57{{.*}}error: no matching function + // function for call to 'block_store' + block_store(buf, 0, v); +} diff --git a/sycl/test/esimd/gather_scatter.cpp b/sycl/test/esimd/gather_scatter.cpp index bc029dc017e1..80cd11697744 100644 --- a/sycl/test/esimd/gather_scatter.cpp +++ b/sycl/test/esimd/gather_scatter.cpp @@ -104,3 +104,23 @@ kernel5(accessor &buf) // function for call to 'scatter' scatter_rgba(buf, offset, v); } + +// Incompatible mode (read). +SYCL_EXTERNAL void +kernel6(local_accessor &buf) SYCL_ESIMD_FUNCTION { + simd v(0, 1); + simd offset(0, 1); + // CHECK: gather_scatter.cpp:115{{.*}}error: no matching function + // function for call to 'scatter' + scatter(buf, offset, v); +} + +// Incompatible mode (read). +SYCL_EXTERNAL void +kernel7(local_accessor &buf) SYCL_ESIMD_FUNCTION { + simd v(0, 1); + simd offset(0, sizeof(int) * 4); + // CHECK: gather_scatter.cpp:125{{.*}}error: no matching function + // function for call to 'scatter' + scatter_rgba(buf, offset, v); +}