@@ -314,12 +314,12 @@ class simd_obj_impl {
314
314
// / argument.
315
315
// / @param acc The accessor to read from.
316
316
// / @param offset 32-bit offset in bytes of the first element.
317
- template <typename AccessorT, typename Flags = element_aligned_tag,
318
- typename = std:: enable_if_t <
319
- detail::is_sycl_accessor_with <
320
- AccessorT, accessor_mode_cap::can_read,
321
- sycl::access::target::global_buffer >::value &&
322
- is_simd_flag_type_v<Flags>>>
317
+ template <
318
+ typename AccessorT, typename Flags = element_aligned_tag,
319
+ typename = std:: enable_if_t <
320
+ detail::is_sycl_accessor_with< AccessorT, accessor_mode_cap::can_read,
321
+ sycl::access::target::device >::value &&
322
+ is_simd_flag_type_v<Flags>>>
323
323
simd_obj_impl (AccessorT acc, uint32_t offset, Flags = {}) noexcept {
324
324
__esimd_dbg_print (simd_obj_impl (AccessorT acc, uint32_t offset, Flags));
325
325
copy_from (acc, offset, Flags{});
@@ -713,7 +713,7 @@ class simd_obj_impl {
713
713
int ChunkSize = 32 ,
714
714
typename = std::enable_if_t <is_simd_flag_type_v<Flags>>>
715
715
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read,
716
- sycl::access::target::global_buffer , void >
716
+ sycl::access::target::device , void >
717
717
copy_from (AccessorT acc, uint32_t offset, Flags = {}) SYCL_ESIMD_FUNCTION;
718
718
719
719
// / Copy all vector elements of this object into a contiguous block in memory.
@@ -739,7 +739,7 @@ class simd_obj_impl {
739
739
int ChunkSize = 32 ,
740
740
typename = std::enable_if_t <is_simd_flag_type_v<Flags>>>
741
741
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write,
742
- sycl::access::target::global_buffer , void >
742
+ sycl::access::target::device , void >
743
743
copy_to (AccessorT acc, uint32_t offset, Flags = {}) const SYCL_ESIMD_FUNCTION;
744
744
745
745
// Unary operations.
0 commit comments