DPC++ Runtime
Runtime libraries for oneAPI DPC++
accessor.hpp
Go to the documentation of this file.
1 //==------------ accessor.hpp - SYCL standard header file ------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
11 #include <sycl/access/access.hpp> // for target, mode
12 #include <sycl/aliases.hpp> // for float4, int4
13 #include <sycl/aspects.hpp> // for aspect
14 #include <sycl/atomic.hpp> // for atomic
15 #include <sycl/buffer.hpp> // for range
16 #include <sycl/detail/accessor_iterator.hpp> // for accessor_iterator
17 #include <sycl/detail/common.hpp> // for code_location
18 #include <sycl/detail/defines.hpp> // for __SYCL_SPECIAL...
19 #include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEP...
20 #include <sycl/detail/export.hpp> // for __SYCL_EXPORT
21 #include <sycl/detail/generic_type_traits.hpp> // for is_genint, Try...
22 #include <sycl/detail/handler_proxy.hpp> // for associateWithH...
23 #include <sycl/detail/helpers.hpp> // for loop
24 #include <sycl/detail/image_accessor_util.hpp> // for imageReadSampl...
25 #include <sycl/detail/owner_less_base.hpp> // for OwnerLessBase
26 #include <sycl/detail/pi.h> // for PI_ERROR_INVAL...
27 #include <sycl/detail/property_helper.hpp> // for PropWithDataKind
28 #include <sycl/detail/property_list_base.hpp> // for PropertyListBase
29 #include <sycl/detail/type_list.hpp> // for is_contained
30 #include <sycl/detail/type_traits.hpp> // for const_if_const_AS
31 #include <sycl/device.hpp> // for device
32 #include <sycl/exception.hpp> // for make_error_code
33 #include <sycl/ext/oneapi/accessor_property_list.hpp> // for accessor_prope...
34 #include <sycl/ext/oneapi/weak_object_base.hpp> // for getSyclWeakObj...
35 #include <sycl/id.hpp> // for id
36 #include <sycl/image.hpp> // for image, image_c...
37 #include <sycl/multi_ptr.hpp> // for multi_ptr
38 #include <sycl/pointers.hpp> // for local_ptr, glo...
39 #include <sycl/properties/accessor_properties.hpp> // for buffer_location
40 #include <sycl/properties/buffer_properties.hpp> // for buffer, buffer...
41 #include <sycl/property_list.hpp> // for property_list
42 #include <sycl/range.hpp> // for range
43 #include <sycl/sampler.hpp> // for addressing_mode
44 #include <sycl/types.hpp> // for vec
45 
46 #ifdef __SYCL_DEVICE_ONLY__
48 #endif
49 
50 #include <cstddef> // for size_t
51 #include <functional> // for hash
52 #include <iterator> // for reverse_iterator
53 #include <limits> // for numeric_limits
54 #include <memory> // for shared_ptr
55 #include <optional> // for nullopt, optional
56 #include <stdint.h> // for uint32_t
57 #include <tuple> // for _Swallow_assign
58 #include <type_traits> // for enable_if_t
59 #include <typeinfo> // for type_info
60 #include <variant> // for hash
61 
121 // +------------------+ +-----------------+ +-----------------------+
122 // | | | | | |
123 // | AccessorBaseHost | | accessor_common | | LocalAccessorBaseHost |
124 // | | | | | |
125 // +------------------+ +-----+-----------+ +--------+--------------+
126 // | | | | |
127 // | +-----------+ +----+ +---------+ +------+
128 // | | | | |
129 // v v v v v
130 // +----------------+ +-----------------+ +-------------+
131 // | | | accessor(1) | | accessor(3) |
132 // | image_accessor | +-----------------| +-------------+
133 // | | | for targets: | | for target: |
134 // +---+---+---+----+ | | | |
135 // | | | | host_buffer | | local |
136 // | | | | global_buffer | +-------------+
137 // | | | | constant_buffer |
138 // | | | +-----------------+
139 // | | |
140 // | | +------------------------------------+
141 // | | |
142 // | +----------------------+ |
143 // v v v
144 // +-----------------+ +--------------+ +-------------+
145 // | acessor(2) | | accessor(4) | | accessor(5) |
146 // +-----------------+ +--------------+ +-------------+
147 // | for targets: | | for targets: | | for target: |
148 // | | | | | |
149 // | host_image | | image | | image_array |
150 // +-----------------+ +--------------+ +-------------+
151 //
176 //
177 // +-----------------+
178 // | |
179 // | accessor_common |
180 // | |
181 // +-----+-------+---+
182 // | |
183 // +----+ +-----+
184 // | |
185 // v v
186 // +----------------+ +-----------------+ +-------------+
187 // | | | accessor(1) | | accessor(3) |
188 // | image_accessor | +-----------------| +-------------+
189 // | | | for targets: | | for target: |
190 // +---+---+---+----+ | | | |
191 // | | | | host_buffer | | local |
192 // | | | | global_buffer | +-------------+
193 // | | | | constant_buffer |
194 // | | | +-----------------+
195 // | | | |
196 // | | | v
197 // | | | +-----------------+
198 // | | | | |
199 // | | | | host_accessor |
200 // | | | | |
201 // | | | +-----------------+
202 // | | |
203 // | | +------------------------------------+
204 // | | |
205 // | +----------------------+ |
206 // v v v
207 // +-----------------+ +--------------+ +-------------+
208 // | acessor(2) | | accessor(4) | | accessor(5) |
209 // +-----------------+ +--------------+ +-------------+
210 // | for targets: | | for targets: | | for target: |
211 // | | | | | |
212 // | host_image | | image | | image_array |
213 // +-----------------+ +--------------+ +-------------+
214 //
232 
233 namespace sycl {
234 inline namespace _V1 {
235 class stream;
236 namespace ext::intel::esimd::detail {
237 // Forward declare a "back-door" access class to support ESIMD.
238 class AccessorPrivateProxy;
239 } // namespace ext::intel::esimd::detail
240 
241 template <typename DataT, int Dimensions = 1,
243  access::target AccessTarget = access::target::device,
245  typename PropertyListT = ext::oneapi::accessor_property_list<>>
246 class accessor;
247 
248 namespace detail {
249 
250 // A helper structure which is shared between buffer accessor and accessor_impl
251 // TODO: Unify with AccessorImplDevice?
252 struct AccHostDataT {
253  AccHostDataT(const sycl::id<3> &Offset, const sycl::range<3> &Range,
254  const sycl::range<3> &MemoryRange, void *Data = nullptr)
255  : MOffset(Offset), MAccessRange(Range), MMemoryRange(MemoryRange),
256  MData(Data) {}
257 
261  void *MData = nullptr;
262  void *Reserved = nullptr;
263 };
264 
265 void __SYCL_EXPORT constructorNotification(void *BufferObj, void *AccessorObj,
266  access::target Target,
268  const code_location &CodeLoc);
269 
270 void __SYCL_EXPORT unsampledImageConstructorNotification(
271  void *ImageObj, void *AccessorObj,
272  const std::optional<image_target> &Target, access::mode Mode,
273  const void *Type, uint32_t ElemSize, const code_location &CodeLoc);
274 
275 void __SYCL_EXPORT sampledImageConstructorNotification(
276  void *ImageObj, void *AccessorObj,
277  const std::optional<image_target> &Target, const void *Type,
278  uint32_t ElemSize, const code_location &CodeLoc);
279 
280 template <typename T>
281 using IsPropertyListT = typename std::is_base_of<PropertyListBase, T>;
282 
283 template <typename T>
285  typename std::is_same<ext::oneapi::accessor_property_list<>, T>;
286 
287 template <typename T> struct IsCxPropertyList {
288  constexpr static bool value = false;
289 };
290 
291 template <typename... Props>
292 struct IsCxPropertyList<ext::oneapi::accessor_property_list<Props...>> {
293  constexpr static bool value = true;
294 };
295 
296 template <> struct IsCxPropertyList<ext::oneapi::accessor_property_list<>> {
297  constexpr static bool value = false;
298 };
299 
300 // Zero-dimensional accessors references at-most a single element, so the range
301 // is either 0 if the associated buffer is empty or 1 otherwise.
302 template <typename BufferT>
304  return std::min(Buffer.size(), size_t{1});
305 }
306 
307 __SYCL_EXPORT device getDeviceFromHandler(handler &CommandGroupHandlerRef);
308 
309 template <typename DataT, int Dimensions, access::mode AccessMode,
311  typename PropertyListT = ext::oneapi::accessor_property_list<>>
313 protected:
315 
316  constexpr static bool IsHostBuf = AccessTarget == access::target::host_buffer;
317  constexpr static bool IsHostTask = AccessTarget == access::target::host_task;
318  // SYCL2020 4.7.6.9.4.3
319  // IsPlaceHolder template parameter has no bearing on whether the accessor
320  // instance is a placeholder. This is determined solely by the constructor.
321  // The rule seems to be: if the constructor receives a CommandGroupHandler
322  // it is NOT a placeholder. Otherwise, it is a placeholder.
323  // However, according to 4.7.6.9.4.6. accessor specialization with
324  // target::host_buffer is never a placeholder. So, if the constructor
325  // used receives a CommandGroupHandler, the accessor will never be a
326  // placeholder. If it doesn't, but IsHostBuf is true, it won't be a
327  // placeholder either. Otherwise, the accessor is a placeholder.
328  constexpr static bool IsPlaceH = !IsHostBuf;
329 
330  // TODO: SYCL 2020 deprecates four of the target enum values
331  // and replaces them with 2 (device and host_task). May want
332  // to change these constexpr.
333  constexpr static bool IsGlobalBuf =
334  AccessTarget == access::target::global_buffer;
335 
336  constexpr static bool IsConstantBuf =
337  AccessTarget == access::target::constant_buffer;
338 
339  constexpr static bool IsAccessAnyWrite =
344 
345  constexpr static bool IsAccessReadOnly = AccessMode == access::mode::read;
346  static constexpr bool IsConst = std::is_const_v<DataT>;
347 
348  constexpr static bool IsAccessReadWrite =
350 
351  constexpr static bool IsAccessAtomic = AccessMode == access::mode::atomic;
352 
354  using ConstRefType = const DataT &;
356 
357  // The class which allows to access value of N dimensional accessor using N
358  // subscript operators, e.g. accessor[2][2][3]
359  template <int SubDims,
360  typename AccType =
361  accessor<DataT, Dimensions, AccessMode, AccessTarget,
362  IsPlaceholder, PropertyListT>>
364  static constexpr int Dims = Dimensions;
365 
366  mutable id<Dims> MIDs;
367  AccType MAccessor;
368 
369  public:
371  : MIDs(IDs), MAccessor(Accessor) {}
372 
373  // Only accessor class is supposed to use this c'tor for the first
374  // operator[].
375  AccessorSubscript(AccType Accessor, size_t Index) : MAccessor(Accessor) {
376  MIDs[0] = Index;
377  }
378 
379  template <int CurDims = SubDims, typename = std::enable_if_t<(CurDims > 1)>>
380  auto operator[](size_t Index) {
381  MIDs[Dims - CurDims] = Index;
382  return AccessorSubscript<CurDims - 1, AccType>(MAccessor, MIDs);
383  }
384 
385  template <int CurDims = SubDims,
386  typename = std::enable_if_t<CurDims == 1 && (IsAccessReadOnly ||
388  typename AccType::reference operator[](size_t Index) const {
389  MIDs[Dims - CurDims] = Index;
390  return MAccessor[MIDs];
391  }
392 
393  template <int CurDims = SubDims>
394  typename std::enable_if_t<CurDims == 1 && IsAccessAtomic, atomic<DataT, AS>>
395  operator[](size_t Index) const {
396  MIDs[Dims - CurDims] = Index;
397  return MAccessor[MIDs];
398  }
399  };
400 };
401 
402 template <typename DataT> constexpr access::mode accessModeFromConstness() {
403  if constexpr (std::is_const_v<DataT>)
404  return access::mode::read;
405  else
407 }
408 
409 template <typename MayBeTag1, typename MayBeTag2>
411  // property_list = {} is not properly detected by deduction guide,
412  // when parameter is passed without curly braces: access(buffer, no_init)
413  // thus simplest approach is to check 2 last arguments for being a tag
414  if constexpr (std::is_same_v<MayBeTag1, mode_tag_t<access::mode::read>> ||
415  std::is_same_v<MayBeTag2, mode_tag_t<access::mode::read>>) {
416  return access::mode::read;
417  }
418 
419  if constexpr (std::is_same_v<MayBeTag1, mode_tag_t<access::mode::write>> ||
420  std::is_same_v<MayBeTag2, mode_tag_t<access::mode::write>>) {
421  return access::mode::write;
422  }
423 
424  if constexpr (std::is_same_v<
425  MayBeTag1,
427  access::target::constant_buffer>> ||
428  std::is_same_v<
429  MayBeTag2,
431  access::target::constant_buffer>>) {
432  return access::mode::read;
433  }
434 
435  if constexpr (std::is_same_v<MayBeTag1,
438  std::is_same_v<MayBeTag2,
441  return access::mode::read;
442  }
443 
444  if constexpr (std::is_same_v<MayBeTag1,
447  std::is_same_v<MayBeTag2,
450  return access::mode::write;
451  }
452 
454 }
455 
456 template <typename MayBeTag1, typename MayBeTag2>
458  if constexpr (std::is_same_v<
459  MayBeTag1,
461  access::target::constant_buffer>> ||
462  std::is_same_v<
463  MayBeTag2,
465  access::target::constant_buffer>>) {
466  return access::target::constant_buffer;
467  }
468 
469  if constexpr (
470  std::is_same_v<MayBeTag1, mode_target_tag_t<access::mode::read,
472  std::is_same_v<MayBeTag2, mode_target_tag_t<access::mode::read,
474  std::is_same_v<MayBeTag1, mode_target_tag_t<access::mode::write,
476  std::is_same_v<MayBeTag2, mode_target_tag_t<access::mode::write,
478  std::is_same_v<MayBeTag1, mode_target_tag_t<access::mode::read_write,
480  std::is_same_v<MayBeTag2, mode_target_tag_t<access::mode::read_write,
483  }
484 
485  return defaultTarget;
486 }
487 
488 template <int Dims> class LocalAccessorBaseDevice {
489 public:
491  : AccessRange(Size),
492  MemRange(InitializedVal<Dims, range>::template get<0>()) {}
493  // TODO: Actually we need only one field here, but currently compiler requires
494  // all of them.
498 
499  bool operator==(const LocalAccessorBaseDevice &Rhs) const {
500  return (AccessRange == Rhs.AccessRange);
501  }
502 };
503 
504 // The class describes a requirement to access a SYCL memory object such as
505 // sycl::buffer and sycl::image. For example, each accessor used in a kernel,
506 // except one with access target "local", adds such requirement for the command
507 // group.
508 
509 template <int Dims> class AccessorImplDevice {
510 public:
511  AccessorImplDevice() = default;
513  range<Dims> MemoryRange)
514  : Offset(Offset), AccessRange(AccessRange), MemRange(MemoryRange) {}
515 
519 
520  bool operator==(const AccessorImplDevice &Rhs) const {
521  return (Offset == Rhs.Offset && AccessRange == Rhs.AccessRange &&
522  MemRange == Rhs.MemRange);
523  }
524 };
525 
526 class AccessorImplHost;
527 
528 void __SYCL_EXPORT addHostAccessorAndWait(AccessorImplHost *Req);
529 
530 class SYCLMemObjI;
531 
532 using AccessorImplPtr = std::shared_ptr<AccessorImplHost>;
533 
534 class __SYCL_EXPORT AccessorBaseHost {
535 protected:
536  AccessorBaseHost(const AccessorImplPtr &Impl) : impl{Impl} {}
537 
538 public:
539  // TODO: the following function to be removed during next ABI break window
540  AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
541  access::mode AccessMode, void *SYCLMemObject, int Dims,
542  int ElemSize, int OffsetInBytes = 0,
543  bool IsSubBuffer = false,
544  const property_list &PropertyList = {});
545  // TODO: the following function to be removed during next ABI break window
546  AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
547  access::mode AccessMode, void *SYCLMemObject, int Dims,
548  int ElemSize, bool IsPlaceH, int OffsetInBytes = 0,
549  bool IsSubBuffer = false,
550  const property_list &PropertyList = {});
551 
552  AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
553  access::mode AccessMode, void *SYCLMemObject, int Dims,
554  int ElemSize, size_t OffsetInBytes = 0,
555  bool IsSubBuffer = false,
556  const property_list &PropertyList = {});
557 
558  AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
559  access::mode AccessMode, void *SYCLMemObject, int Dims,
560  int ElemSize, bool IsPlaceH, size_t OffsetInBytes = 0,
561  bool IsSubBuffer = false,
562  const property_list &PropertyList = {});
563 
564 public:
565  id<3> &getOffset();
566  range<3> &getAccessRange();
567  range<3> &getMemoryRange();
568  void *getPtr() noexcept;
569  unsigned int getElemSize() const;
570 
571  const id<3> &getOffset() const;
572  const range<3> &getAccessRange() const;
573  const range<3> &getMemoryRange() const;
574  void *getPtr() const noexcept;
575  bool isPlaceholder() const;
576  bool isMemoryObjectUsedByGraph() const;
577 
578  detail::AccHostDataT &getAccData();
579 
580  const property_list &getPropList() const;
581 
582  void *getMemoryObject() const;
583 
584  template <class Obj>
585  friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject);
586 
587  template <class T>
588  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
589 
590  template <typename, int, access::mode, access::target, access::placeholder,
591  typename>
592  friend class accessor;
593 
595 
596 private:
597  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
598 };
599 
602 
603 class __SYCL_EXPORT LocalAccessorBaseHost {
604 protected:
605  LocalAccessorBaseHost(const LocalAccessorImplPtr &Impl) : impl{Impl} {}
606 
607 public:
608  LocalAccessorBaseHost(sycl::range<3> Size, int Dims, int ElemSize,
609  const property_list &PropertyList = {});
610  sycl::range<3> &getSize();
611  const sycl::range<3> &getSize() const;
612  void *getPtr();
613  void *getPtr() const;
614  int getNumOfDims();
615  int getElementSize();
616  const property_list &getPropList() const;
617 
618 protected:
619  template <class Obj>
620  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
621 
622  template <class T>
623  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
624 
626 };
627 
631  std::shared_ptr<UnsampledImageAccessorImplHost>;
633  std::shared_ptr<SampledImageAccessorImplHost>;
634 
635 void __SYCL_EXPORT
637 void __SYCL_EXPORT
639 
640 class __SYCL_EXPORT UnsampledImageAccessorBaseHost {
641 protected:
643  : impl{Impl} {}
644 
645 public:
647  void *SYCLMemObject, int Dims, int ElemSize,
648  id<3> Pitch, image_channel_type ChannelType,
649  image_channel_order ChannelOrder,
650  const property_list &PropertyList = {});
651  const sycl::range<3> &getSize() const;
652  void *getMemoryObject() const;
653  detail::AccHostDataT &getAccData();
654  void *getPtr();
655  void *getPtr() const;
656  int getNumOfDims() const;
657  int getElementSize() const;
658  id<3> getPitch() const;
659  image_channel_type getChannelType() const;
660  image_channel_order getChannelOrder() const;
661  const property_list &getPropList() const;
662 
663 protected:
664  template <class Obj>
665  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
666 
667  template <class T>
668  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
669 
671 
672  // The function references helper methods required by GDB pretty-printers
674 #ifndef NDEBUG
675  const auto *this_const = this;
676  (void)getSize();
677  (void)this_const->getSize();
678  (void)getPtr();
679  (void)this_const->getPtr();
680 #endif
681  }
682 
683 #ifndef __SYCL_DEVICE_ONLY__
684  // Reads a pixel of the underlying image at the specified coordinate. It is
685  // the responsibility of the caller to ensure that the coordinate type is
686  // valid.
687  template <typename DataT, typename CoordT>
688  DataT read(const CoordT &Coords) const noexcept {
692  return imageReadSamplerHostImpl<CoordT, DataT>(
693  Coords, Smpl, getSize(), getPitch(), getChannelType(),
694  getChannelOrder(), getPtr(), getElementSize());
695  }
696 
697  // Writes to a pixel of the underlying image at the specified coordinate. It
698  // is the responsibility of the caller to ensure that the coordinate type is
699  // valid.
700  template <typename DataT, typename CoordT>
701  void write(const CoordT &Coords, const DataT &Color) const {
702  imageWriteHostImpl(Coords, Color, getPitch(), getElementSize(),
703  getChannelType(), getChannelOrder(), getPtr());
704  }
705 #endif
706 };
707 
708 class __SYCL_EXPORT SampledImageAccessorBaseHost {
709 protected:
711  : impl{Impl} {}
712 
713 public:
714  SampledImageAccessorBaseHost(sycl::range<3> Size, void *SYCLMemObject,
715  int Dims, int ElemSize, id<3> Pitch,
716  image_channel_type ChannelType,
717  image_channel_order ChannelOrder,
718  image_sampler Sampler,
719  const property_list &PropertyList = {});
720  const sycl::range<3> &getSize() const;
721  void *getMemoryObject() const;
722  detail::AccHostDataT &getAccData();
723  void *getPtr();
724  void *getPtr() const;
725  int getNumOfDims() const;
726  int getElementSize() const;
727  id<3> getPitch() const;
728  image_channel_type getChannelType() const;
729  image_channel_order getChannelOrder() const;
730  image_sampler getSampler() const;
731  const property_list &getPropList() const;
732 
733 protected:
734  template <class Obj>
735  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
736 
737  template <class T>
738  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
739 
741 
742  // The function references helper methods required by GDB pretty-printers
744 #ifndef NDEBUG
745  const auto *this_const = this;
746  (void)getSize();
747  (void)this_const->getSize();
748  (void)getPtr();
749  (void)this_const->getPtr();
750 #endif
751  }
752 
753 #ifndef __SYCL_DEVICE_ONLY__
754  // Reads a pixel of the underlying image at the specified coordinate. It is
755  // the responsibility of the caller to ensure that the coordinate type is
756  // valid.
757  template <typename DataT, typename CoordT>
758  DataT read(const CoordT &Coords) const {
759  return imageReadSamplerHostImpl<CoordT, DataT>(
760  Coords, getSampler(), getSize(), getPitch(), getChannelType(),
761  getChannelOrder(), getPtr(), getElementSize());
762  }
763 #endif
764 };
765 
766 template <int Dim, typename T> struct IsValidCoordDataT;
767 template <typename T> struct IsValidCoordDataT<1, T> {
768  constexpr static bool value = detail::is_contained<
770 };
771 template <typename T> struct IsValidCoordDataT<2, T> {
772  constexpr static bool value = detail::is_contained<
774  vec<opencl::cl_float, 2>>>::type::value;
775 };
776 template <typename T> struct IsValidCoordDataT<3, T> {
777  constexpr static bool value = detail::is_contained<
779  vec<opencl::cl_float, 4>>>::type::value;
780 };
781 
782 template <int Dim, typename T> struct IsValidUnsampledCoord2020DataT;
783 template <typename T> struct IsValidUnsampledCoord2020DataT<1, T> {
784  constexpr static bool value = std::is_same_v<T, int>;
785 };
786 template <typename T> struct IsValidUnsampledCoord2020DataT<2, T> {
787  constexpr static bool value = std::is_same_v<T, int2>;
788 };
789 template <typename T> struct IsValidUnsampledCoord2020DataT<3, T> {
790  constexpr static bool value = std::is_same_v<T, int4>;
791 };
792 
793 template <int Dim, typename T> struct IsValidSampledCoord2020DataT;
794 template <typename T> struct IsValidSampledCoord2020DataT<1, T> {
795  constexpr static bool value = std::is_same_v<T, float>;
796 };
797 template <typename T> struct IsValidSampledCoord2020DataT<2, T> {
798  constexpr static bool value = std::is_same_v<T, float2>;
799 };
800 template <typename T> struct IsValidSampledCoord2020DataT<3, T> {
801  constexpr static bool value = std::is_same_v<T, float4>;
802 };
803 
804 template <typename DataT, int Dimensions, access::mode AccessMode,
807 
808 // Image accessor
809 template <typename DataT, int Dimensions, access::mode AccessMode,
812 #ifndef __SYCL_DEVICE_ONLY__
813  : public detail::AccessorBaseHost {
814  size_t MImageCount;
815  image_channel_order MImgChannelOrder;
816  image_channel_type MImgChannelType;
817 #else
818 {
819 
820  using OCLImageTy = typename detail::opencl_image_type<Dimensions, AccessMode,
821  AccessTarget>::type;
822  OCLImageTy MImageObj;
823  char MPadding[sizeof(detail::AccessorBaseHost) +
824  sizeof(size_t /*MImageCount*/) + sizeof(image_channel_order) +
825  sizeof(image_channel_type) - sizeof(OCLImageTy)];
826 
827 protected:
828  void imageAccessorInit(OCLImageTy Image) { MImageObj = Image; }
829 
830 private:
831 #endif
832  template <typename T1, int T2, access::mode T3, access::placeholder T4>
833  friend class __image_array_slice__;
834 
835  constexpr static bool IsHostImageAcc =
836  (AccessTarget == access::target::host_image);
837 
838  constexpr static bool IsImageAcc = (AccessTarget == access::target::image);
839 
840  constexpr static bool IsImageArrayAcc =
841  (AccessTarget == access::target::image_array);
842 
843  constexpr static bool IsImageAccessWriteOnly =
846 
847  constexpr static bool IsImageAccessAnyWrite =
848  (IsImageAccessWriteOnly || AccessMode == access::mode::read_write);
849 
850  constexpr static bool IsImageAccessReadOnly =
852 
853  constexpr static bool IsImageAccessAnyRead =
854  (IsImageAccessReadOnly || AccessMode == access::mode::read_write);
855 
856  static_assert(std::is_same_v<DataT, vec<opencl::cl_int, 4>> ||
857  std::is_same_v<DataT, vec<opencl::cl_uint, 4>> ||
858  std::is_same_v<DataT, vec<opencl::cl_float, 4>> ||
859  std::is_same_v<DataT, vec<opencl::cl_half, 4>>,
860  "The data type of an image accessor must be only cl_int4, "
861  "cl_uint4, cl_float4 or cl_half4 from SYCL namespace");
862 
863  static_assert(IsImageAcc || IsHostImageAcc || IsImageArrayAcc,
864  "Expected image type");
865 
867  "Expected false as Placeholder value for image accessor.");
868 
869  static_assert(
870  ((IsImageAcc || IsImageArrayAcc) &&
871  (IsImageAccessWriteOnly || IsImageAccessReadOnly)) ||
872  (IsHostImageAcc && (IsImageAccessAnyWrite || IsImageAccessAnyRead)),
873  "Access modes can be only read/write/discard_write for image/image_array "
874  "target accessor, or they can be only "
875  "read/write/discard_write/read_write for host_image target accessor.");
876 
877  static_assert(Dimensions > 0 && Dimensions <= 3,
878  "Dimensions can be 1/2/3 for image accessor.");
879 
880 #ifdef __SYCL_DEVICE_ONLY__
881 
882  sycl::vec<int, Dimensions> getRangeInternal() const {
883  return __invoke_ImageQuerySize<sycl::vec<int, Dimensions>, OCLImageTy>(
884  MImageObj);
885  }
886 
887  size_t getElementSize() const {
888  int ChannelType = __invoke_ImageQueryFormat<int, OCLImageTy>(MImageObj);
889  int ChannelOrder = __invoke_ImageQueryOrder<int, OCLImageTy>(MImageObj);
890  int ElementSize = getSPIRVElementSize(ChannelType, ChannelOrder);
891  return ElementSize;
892  }
893 
894 #else
895 
896  sycl::vec<int, Dimensions> getRangeInternal() const {
897  // TODO: Implement for host.
898  throw sycl::exception(
900  "image::getRangeInternal() is not implemented for host");
901  return sycl::vec<int, Dimensions>{1};
902  }
903 
904 #endif
905 
906 #ifndef __SYCL_DEVICE_ONLY__
907 protected:
909 #endif // __SYCL_DEVICE_ONLY__
910 
911 private:
912  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
913 
914 #ifdef __SYCL_DEVICE_ONLY__
915  const OCLImageTy getNativeImageObj() const { return MImageObj; }
916 #endif // __SYCL_DEVICE_ONLY__
917 
918 public:
919  using value_type = DataT;
920  using reference = DataT &;
921  using const_reference = const DataT &;
922 
923  // image_accessor Constructors.
924 
925 #ifdef __SYCL_DEVICE_ONLY__
926  // Default constructor for objects later initialized with __init member.
927  image_accessor() {}
928 #endif
929 
930  // Available only when: accessTarget == access::target::host_image
931  // template <typename AllocatorT>
932  // accessor(image<dimensions, AllocatorT> &imageRef);
933  template <
934  typename AllocatorT, int Dims = Dimensions,
935  typename = std::enable_if_t<(Dims > 0 && Dims <= 3) && IsHostImageAcc>>
936  image_accessor(image<Dims, AllocatorT> &ImageRef, int ImageElementSize)
937 #ifdef __SYCL_DEVICE_ONLY__
938  {
939  (void)ImageRef;
940  (void)ImageElementSize;
941  // No implementation needed for device. The constructor is only called by
942  // host.
943  }
944 #else
945  : AccessorBaseHost({ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
946  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
947  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
948  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
949  Dimensions, ImageElementSize, size_t(0)),
950  MImageCount(ImageRef.size()),
951  MImgChannelOrder(ImageRef.getChannelOrder()),
952  MImgChannelType(ImageRef.getChannelType()) {
954  }
955 #endif
956 
957  // Available only when: accessTarget == access::target::image
958  // template <typename AllocatorT>
959  // accessor(image<dimensions, AllocatorT> &imageRef,
960  // handler &commandGroupHandlerRef);
961  template <typename AllocatorT, int Dims = Dimensions,
962  typename = std::enable_if_t<(Dims > 0 && Dims <= 3) && IsImageAcc>>
964  handler &CommandGroupHandlerRef, int ImageElementSize)
965 #ifdef __SYCL_DEVICE_ONLY__
966  {
967  (void)ImageRef;
968  (void)CommandGroupHandlerRef;
969  (void)ImageElementSize;
970  // No implementation needed for device. The constructor is only called by
971  // host.
972  }
973 #else
974  : AccessorBaseHost({ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
975  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
976  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
977  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
978  Dimensions, ImageElementSize, size_t(0)),
979  MImageCount(ImageRef.size()),
980  MImgChannelOrder(ImageRef.getChannelOrder()),
981  MImgChannelType(ImageRef.getChannelType()) {
982 
983  device Device = getDeviceFromHandler(CommandGroupHandlerRef);
984  if (!Device.has(aspect::ext_intel_legacy_image))
985  throw sycl::exception(
986  sycl::errc::feature_not_supported,
987  "SYCL 1.2.1 images are not supported by this device.");
988  }
989 #endif
990 
991  /* -- common interface members -- */
992 
993  // operator == and != need to be defined only for host application as per the
994  // SYCL spec 1.2.1
995 #ifndef __SYCL_DEVICE_ONLY__
996  bool operator==(const image_accessor &Rhs) const { return Rhs.impl == impl; }
997 #else
998  // The operator with __SYCL_DEVICE_ONLY__ need to be declared for compilation
999  // of host application with device compiler.
1000  // Usage of this operator inside the kernel code will give a runtime failure.
1001  bool operator==(const image_accessor &Rhs) const;
1002 #endif
1003 
1004  bool operator!=(const image_accessor &Rhs) const { return !(Rhs == *this); }
1005 
1006  // get_count() method : Returns the number of elements of the SYCL image this
1007  // SYCL accessor is accessing.
1008  //
1009  // get_range() method : Returns a range object which represents the number of
1010  // elements of dataT per dimension that this accessor may access.
1011  // The range object returned must equal to the range of the image this
1012  // accessor is associated with.
1013 
1014 #ifdef __SYCL_DEVICE_ONLY__
1015 
1016  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1017  size_t get_count() const { return size(); }
1018  size_t size() const noexcept { return get_range<Dimensions>().size(); }
1019 
1020  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 1>>
1021  range<1> get_range() const {
1022  int Range = getRangeInternal();
1023  return range<1>(Range);
1024  }
1025  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 2>>
1026  range<2> get_range() const {
1027  int2 Range = getRangeInternal();
1028  return range<2>(Range[0], Range[1]);
1029  }
1030  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 3>>
1031  range<3> get_range() const {
1032  int3 Range = getRangeInternal();
1033  return range<3>(Range[0], Range[1], Range[2]);
1034  }
1035 
1036 #else
1037  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1038  size_t get_count() const { return size(); };
1039  size_t size() const noexcept { return MImageCount; };
1040 
1041  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
1043  return detail::convertToArrayOfN<Dims, 1>(getAccessRange());
1044  }
1045 
1046 #endif
1047 
1048  // Available only when:
1049  // (accessTarget == access::target::image && accessMode == access::mode::read)
1050  // || (accessTarget == access::target::host_image && ( accessMode ==
1051  // access::mode::read || accessMode == access::mode::read_write))
1052  template <typename CoordT, int Dims = Dimensions,
1053  typename = std::enable_if_t<
1054  (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
1055  (detail::is_genint_v<CoordT>)&&(
1056  (IsImageAcc && IsImageAccessReadOnly) ||
1057  (IsHostImageAcc && IsImageAccessAnyRead))>>
1058  DataT read(const CoordT &Coords) const {
1059 #ifdef __SYCL_DEVICE_ONLY__
1060  return __invoke__ImageRead<DataT, OCLImageTy, CoordT>(MImageObj, Coords);
1061 #else
1064  return read<CoordT, Dims>(Coords, Smpl);
1065 #endif
1066  }
1067 
1068  // Available only when:
1069  // (accessTarget == access::target::image && accessMode == access::mode::read)
1070  // || (accessTarget == access::target::host_image && ( accessMode ==
1071  // access::mode::read || accessMode == access::mode::read_write))
1072  template <typename CoordT, int Dims = Dimensions,
1073  typename = std::enable_if_t<
1074  (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
1075  ((IsImageAcc && IsImageAccessReadOnly) ||
1076  (IsHostImageAcc && IsImageAccessAnyRead))>>
1077  DataT read(const CoordT &Coords, const sampler &Smpl) const {
1078 #ifdef __SYCL_DEVICE_ONLY__
1079  return __invoke__ImageReadSampler<DataT, OCLImageTy, CoordT>(
1080  MImageObj, Coords, Smpl.impl.m_Sampler);
1081 #else
1082  return imageReadSamplerHostImpl<CoordT, DataT>(
1083  Coords, Smpl, getAccessRange() /*Image Range*/,
1084  getOffset() /*Image Pitch*/, MImgChannelType, MImgChannelOrder,
1085  AccessorBaseHost::getPtr() /*ptr to image*/,
1087 #endif
1088  }
1089 
1090  // Available only when:
1091  // (accessTarget == access::target::image && (accessMode ==
1092  // access::mode::write || accessMode == access::mode::discard_write)) ||
1093  // (accessTarget == access::target::host_image && (accessMode ==
1094  // access::mode::write || accessMode == access::mode::discard_write ||
1095  // accessMode == access::mode::read_write))
1096  template <
1097  typename CoordT, int Dims = Dimensions,
1098  typename = std::enable_if_t<(Dims > 0) &&
1099  (detail::is_genint_v<CoordT>)&&(
1101  ((IsImageAcc && IsImageAccessWriteOnly) ||
1102  (IsHostImageAcc && IsImageAccessAnyWrite))>>
1103  void write(const CoordT &Coords, const DataT &Color) const {
1104 #ifdef __SYCL_DEVICE_ONLY__
1105  __invoke__ImageWrite<OCLImageTy, CoordT, DataT>(MImageObj, Coords, Color);
1106 #else
1107  imageWriteHostImpl(Coords, Color, getOffset() /*ImagePitch*/,
1108  AccessorBaseHost::getElemSize(), MImgChannelType,
1109  MImgChannelOrder,
1110  AccessorBaseHost::getPtr() /*Ptr to Image*/);
1111 #endif
1112  }
1113 };
1114 
1115 template <typename DataT, int Dimensions, access::mode AccessMode,
1118 
1119  static_assert(Dimensions < 3,
1120  "Image slice cannot have more then 2 dimensions");
1121 
1122  constexpr static int AdjustedDims = (Dimensions == 2) ? 4 : Dimensions + 1;
1123 
1124  template <typename CoordT, typename CoordElemType = get_elem_type_t<CoordT>>
1126  getAdjustedCoords(const CoordT &Coords) const {
1127  CoordElemType LastCoord = 0;
1128 
1129  if (std::is_same<float, CoordElemType>::value) {
1130  sycl::vec<int, Dimensions + 1> Size = MBaseAcc.getRangeInternal();
1131  LastCoord =
1132  MIdx / static_cast<float>(Size.template swizzle<Dimensions>());
1133  } else {
1134  LastCoord = MIdx;
1135  }
1136 
1137  sycl::vec<CoordElemType, Dimensions> LeftoverCoords{LastCoord};
1138  sycl::vec<CoordElemType, AdjustedDims> AdjustedCoords{Coords,
1139  LeftoverCoords};
1140  return AdjustedCoords;
1141  }
1142 
1143 public:
1147  BaseAcc,
1148  size_t Idx)
1149  : MBaseAcc(BaseAcc), MIdx(Idx) {}
1150 
1151  template <typename CoordT, int Dims = Dimensions,
1152  typename = std::enable_if_t<
1153  (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value)>>
1154  DataT read(const CoordT &Coords) const {
1155  return MBaseAcc.read(getAdjustedCoords(Coords));
1156  }
1157 
1158  template <typename CoordT, int Dims = Dimensions,
1159  typename = std::enable_if_t<(Dims > 0) &&
1161  DataT read(const CoordT &Coords, const sampler &Smpl) const {
1162  return MBaseAcc.read(getAdjustedCoords(Coords), Smpl);
1163  }
1164 
1165  template <typename CoordT, int Dims = Dimensions,
1166  typename = std::enable_if_t<(Dims > 0) &&
1168  void write(const CoordT &Coords, const DataT &Color) const {
1169  return MBaseAcc.write(getAdjustedCoords(Coords), Color);
1170  }
1171 
1172 #ifdef __SYCL_DEVICE_ONLY__
1173  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1174  size_t get_count() const { return size(); }
1175  size_t size() const noexcept { return get_range<Dimensions>().size(); }
1176 
1177  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 1>>
1178  range<1> get_range() const {
1179  int2 Count = MBaseAcc.getRangeInternal();
1180  return range<1>(Count.x());
1181  }
1182  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 2>>
1183  range<2> get_range() const {
1184  int3 Count = MBaseAcc.getRangeInternal();
1185  return range<2>(Count.x(), Count.y());
1186  }
1187 
1188 #else
1189 
1190  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1191  size_t get_count() const { return size(); }
1192  size_t size() const noexcept {
1193  return MBaseAcc.MImageCount / MBaseAcc.getAccessRange()[Dimensions];
1194  }
1195 
1196  template <int Dims = Dimensions,
1197  typename = std::enable_if_t<(Dims == 1 || Dims == 2)>>
1199  return detail::convertToArrayOfN<Dims, 1>(MBaseAcc.getAccessRange());
1200  }
1201 
1202 #endif
1203 
1204 private:
1205  size_t MIdx;
1208  MBaseAcc;
1209 };
1210 
1211 } // namespace detail
1212 
1218 template <typename DataT, int Dimensions, access::mode AccessMode,
1220  typename PropertyListT>
1222 #ifndef __SYCL_DEVICE_ONLY__
1223  public detail::AccessorBaseHost,
1224 #endif
1225  public detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
1226  IsPlaceholder, PropertyListT>,
1227  public detail::OwnerLessBase<
1228  accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
1229  PropertyListT>> {
1230 protected:
1231  static_assert((AccessTarget == access::target::global_buffer ||
1232  AccessTarget == access::target::constant_buffer ||
1233  AccessTarget == access::target::host_buffer ||
1234  AccessTarget == access::target::host_task),
1235  "Expected buffer type");
1236 
1237  static_assert((AccessTarget == access::target::global_buffer ||
1238  AccessTarget == access::target::host_buffer ||
1239  AccessTarget == access::target::host_task) ||
1240  (AccessTarget == access::target::constant_buffer &&
1242  "Access mode can be only read for constant buffers");
1243 
1244  static_assert(detail::IsPropertyListT<PropertyListT>::value,
1245  "PropertyListT must be accessor_property_list");
1246 
1247  using AccessorCommonT =
1248  detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
1249  IsPlaceholder, PropertyListT>;
1250 
1251  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
1252 
1253  using AccessorCommonT::AS;
1254  // Cannot do "using AccessorCommonT::Flag" as it doesn't work with g++ as host
1255  // compiler, for some reason.
1256  static constexpr bool IsAccessAnyWrite = AccessorCommonT::IsAccessAnyWrite;
1257  static constexpr bool IsAccessReadOnly = AccessorCommonT::IsAccessReadOnly;
1258  static constexpr bool IsConstantBuf = AccessorCommonT::IsConstantBuf;
1259  static constexpr bool IsGlobalBuf = AccessorCommonT::IsGlobalBuf;
1260  static constexpr bool IsHostBuf = AccessorCommonT::IsHostBuf;
1261  static constexpr bool IsPlaceH = AccessorCommonT::IsPlaceH;
1262  static constexpr bool IsConst = AccessorCommonT::IsConst;
1263  static constexpr bool IsHostTask = AccessorCommonT::IsHostTask;
1264  template <int Dims>
1265  using AccessorSubscript =
1266  typename AccessorCommonT::template AccessorSubscript<Dims>;
1267 
1268  static_assert(
1269  !IsConst || IsAccessReadOnly,
1270  "A const qualified DataT is only allowed for a read-only accessor");
1271 
1272  using ConcreteASPtrType = typename detail::DecoratedType<
1273  typename std::conditional_t<IsAccessReadOnly && !IsConstantBuf,
1274  const DataT, DataT>,
1275  AS>::type *;
1276 
1277  using RefType = detail::const_if_const_AS<AS, DataT> &;
1278  using ConstRefType = const DataT &;
1279  using PtrType = detail::const_if_const_AS<AS, DataT> *;
1280 
1281  template <int Dims = Dimensions> size_t getLinearIndex(id<Dims> Id) const {
1282 
1283  size_t Result = 0;
1284  detail::loop<Dims>([&, this](size_t I) {
1285  Result = Result * getMemoryRange()[I] + Id[I];
1286  // We've already adjusted for the accessor's offset in the __init, so
1287  // don't include it here in case of device.
1288 #ifndef __SYCL_DEVICE_ONLY__
1289  if constexpr (!(PropertyListT::template has_property<
1291  Result += getOffset()[I];
1292  }
1293 #endif // __SYCL_DEVICE_ONLY__
1294  });
1295 
1296  return Result;
1297  }
1298 
1299  template <typename T, int Dims>
1300  struct IsSameAsBuffer
1301  : std::bool_constant<std::is_same_v<T, DataT> && (Dims > 0) &&
1302  (Dims == Dimensions)> {};
1303 
1304  static access::mode getAdjustedMode(const PropertyListT &PropertyList) {
1305  access::mode AdjustedMode = AccessMode;
1306 
1307  if (PropertyList.template has_property<property::no_init>() ||
1308  PropertyList.template has_property<property::noinit>()) {
1309  if (AdjustedMode == access::mode::write) {
1310  AdjustedMode = access::mode::discard_write;
1311  } else if (AdjustedMode == access::mode::read_write) {
1312  AdjustedMode = access::mode::discard_read_write;
1313  }
1314  }
1315 
1316  return AdjustedMode;
1317  }
1318 
1319  template <typename TagT>
1320  struct IsValidTag
1321  : std::disjunction<
1322  std::is_same<TagT, mode_tag_t<AccessMode>>,
1323  std::is_same<TagT, mode_target_tag_t<AccessMode, AccessTarget>>> {};
1324 
1325  template <typename DataT_, int Dimensions_, access::mode AccessMode_,
1326  access::target AccessTarget_, access::placeholder IsPlaceholder_,
1327  typename PropertyListT_>
1328  friend class accessor;
1329 
1330 #ifdef __SYCL_DEVICE_ONLY__
1331 
1332  id<AdjustedDim> &getOffset() { return impl.Offset; }
1333  range<AdjustedDim> &getAccessRange() { return impl.AccessRange; }
1334  range<AdjustedDim> &getMemoryRange() { return impl.MemRange; }
1335 
1336  const id<AdjustedDim> &getOffset() const { return impl.Offset; }
1337  const range<AdjustedDim> &getAccessRange() const { return impl.AccessRange; }
1338  const range<AdjustedDim> &getMemoryRange() const { return impl.MemRange; }
1339 
1340  detail::AccessorImplDevice<AdjustedDim> impl;
1341 
1342  union {
1343  ConcreteASPtrType MData;
1344  };
1345 
1346  void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
1347  range<AdjustedDim> MemRange, id<AdjustedDim> Offset) {
1348  MData = Ptr;
1349  detail::loop<AdjustedDim>([&, this](size_t I) {
1350  if constexpr (!(PropertyListT::template has_property<
1352  getOffset()[I] = Offset[I];
1353  }
1354  getAccessRange()[I] = AccessRange[I];
1355  getMemoryRange()[I] = MemRange[I];
1356  });
1357 
1358  // Adjust for offsets as that part is invariant for all invocations of
1359  // operator[]. Will have to re-adjust in get_pointer.
1360  MData += getTotalOffset();
1361  }
1362 
1363  // __init variant used by the device compiler for ESIMD kernels.
1364  // TODO: In ESIMD accessors usage is limited for now - access range, mem
1365  // range and offset are not supported.
1366  void __init_esimd(ConcreteASPtrType Ptr) {
1367  MData = Ptr;
1368 #ifdef __ESIMD_FORCE_STATELESS_MEM
1369  detail::loop<AdjustedDim>([&, this](size_t I) {
1370  getOffset()[I] = 0;
1371  getAccessRange()[I] = 0;
1372  getMemoryRange()[I] = 0;
1373  });
1374 #endif
1375  }
1376 
1377  ConcreteASPtrType getQualifiedPtr() const noexcept { return MData; }
1378 
1379 #ifndef __SYCL_DEVICE_ONLY__
1380  using AccessorBaseHost::impl;
1381 #endif
1382 
1383 public:
1384  // Default constructor for objects later initialized with __init member.
1385  accessor()
1386  : impl({}, detail::InitializedVal<AdjustedDim, range>::template get<0>(),
1387  detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
1388 
1389 #else
1390  accessor(const detail::AccessorImplPtr &Impl)
1391  : detail::AccessorBaseHost{Impl} {}
1392 
1393  void *getPtr() { return AccessorBaseHost::getPtr(); }
1394 
1395  const id<3> getOffset() const {
1396  if constexpr (IsHostBuf)
1397  return MAccData ? MAccData->MOffset : id<3>();
1398  else
1399  return AccessorBaseHost::getOffset();
1400  }
1401  const range<3> &getAccessRange() const {
1403  }
1404  const range<3> getMemoryRange() const {
1405  if constexpr (IsHostBuf)
1406  return MAccData ? MAccData->MMemoryRange : range(0, 0, 0);
1407  else
1409  }
1410 
1411  void *getPtr() const { return AccessorBaseHost::getPtr(); }
1412 
1413  void initHostAcc() { MAccData = &getAccData(); }
1414 
1415  // The function references helper methods required by GDB pretty-printers
1416  void GDBMethodsAnchor() {
1417 #ifndef NDEBUG
1418  const auto *this_const = this;
1419  (void)getMemoryRange();
1420  (void)this_const->getMemoryRange();
1421  (void)getOffset();
1422  (void)this_const->getOffset();
1423  (void)getPtr();
1424  (void)this_const->getPtr();
1425  (void)getAccessRange();
1426  (void)this_const->getAccessRange();
1427 #endif
1428  }
1429 
1430  detail::AccHostDataT *MAccData = nullptr;
1431 
1432  char padding[sizeof(detail::AccessorImplDevice<AdjustedDim>) +
1433  sizeof(PtrType) - sizeof(detail::AccessorBaseHost) -
1434  sizeof(MAccData)];
1435 
1436  PtrType getQualifiedPtr() const noexcept {
1437  if constexpr (IsHostBuf)
1438  return MAccData ? reinterpret_cast<PtrType>(MAccData->MData) : nullptr;
1439  else
1440  return reinterpret_cast<PtrType>(AccessorBaseHost::getPtr());
1441  }
1442 
1443 public:
1444  accessor()
1445  : AccessorBaseHost(
1446  /*Offset=*/{0, 0, 0}, /*AccessRange=*/{0, 0, 0},
1447  /*MemoryRange=*/{0, 0, 0},
1448  /*AccessMode=*/getAdjustedMode({}),
1449  /*SYCLMemObject=*/nullptr, /*Dims=*/0, /*ElemSize=*/0,
1450  /*IsPlaceH=*/false,
1451  /*OffsetInBytes=*/0, /*IsSubBuffer=*/false, /*PropertyList=*/{}){};
1452 
1453  template <typename, int, access_mode> friend class host_accessor;
1454 
1455 #endif // __SYCL_DEVICE_ONLY__
1456 
1457 private:
1458  friend class sycl::stream;
1459  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
1460 
1461  template <class Obj>
1462  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
1463 
1464  template <class T>
1465  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
1466 
1467 public:
1468  // 4.7.6.9.1. Interface for buffer command accessors
1469  // value_type is defined as const DataT for read_only accessors, DataT
1470  // otherwise
1471  using value_type =
1472  std::conditional_t<AccessMode == access_mode::read, const DataT, DataT>;
1473  using reference = value_type &;
1474  using const_reference = const DataT &;
1475 
1476  template <access::decorated IsDecorated>
1477  using accessor_ptr =
1478  std::conditional_t<AccessTarget == access::target::device,
1479  global_ptr<value_type, IsDecorated>, value_type *>;
1480 
1481  using iterator = typename detail::accessor_iterator<value_type, AdjustedDim>;
1482  using const_iterator =
1483  typename detail::accessor_iterator<const value_type, AdjustedDim>;
1484  using reverse_iterator = std::reverse_iterator<iterator>;
1485  using const_reverse_iterator = std::reverse_iterator<const_iterator>;
1486  using difference_type =
1488  using size_type = std::size_t;
1489 
1492  void throwIfUsedByGraph() const {
1493 #ifndef __SYCL_DEVICE_ONLY__
1494  if (IsHostBuf && AccessorBaseHost::isMemoryObjectUsedByGraph()) {
1496  "Host accessors cannot be created for buffers "
1497  "which are currently in use by a command graph.");
1498  }
1499 #endif
1500  }
1501 
1502  // The list of accessor constructors with their arguments
1503  // -------+---------+-------+----+-----+--------------
1504  // Dimensions = 0
1505  // -------+---------+-------+----+-----+--------------
1506  // buffer | | | | | property_list
1507  // buffer | handler | | | | property_list
1508  // -------+---------+-------+----+-----+--------------
1509  // Dimensions >= 1
1510  // -------+---------+-------+----+-----+--------------
1511  // buffer | | | | | property_list
1512  // buffer | | | | tag | property_list
1513  // buffer | handler | | | | property_list
1514  // buffer | handler | | | tag | property_list
1515  // buffer | | range | | | property_list
1516  // buffer | | range | | tag | property_list
1517  // buffer | handler | range | | | property_list
1518  // buffer | handler | range | | tag | property_list
1519  // buffer | | range | id | | property_list
1520  // buffer | | range | id | tag | property_list
1521  // buffer | handler | range | id | | property_list
1522  // buffer | handler | range | id | tag | property_list
1523  // -------+---------+-------+----+-----+--------------
1524 
1525 public:
1526  // implicit conversion between const / non-const types for read only accessors
1527  template <typename DataT_,
1528  typename = std::enable_if_t<
1529  IsAccessReadOnly && !std::is_same_v<DataT_, DataT> &&
1530  std::is_same_v<std::remove_const_t<DataT_>,
1531  std::remove_const_t<DataT>>>>
1532  accessor(const accessor<DataT_, Dimensions, AccessMode, AccessTarget,
1533  IsPlaceholder, PropertyListT> &other)
1534 #ifdef __SYCL_DEVICE_ONLY__
1535  : impl(other.impl), MData(other.MData) {
1536 #else
1537  : accessor(other.impl) {
1538 #endif // __SYCL_DEVICE_ONLY__
1539  }
1540 
1541  // implicit conversion from read_write T accessor to read only T (const)
1542  // accessor
1543  template <typename DataT_, access::mode AccessMode_,
1544  typename = std::enable_if_t<
1545  (AccessMode_ == access_mode::read_write) && IsAccessReadOnly &&
1546  std::is_same_v<std::remove_const_t<DataT_>,
1547  std::remove_const_t<DataT>>>>
1548  accessor(const accessor<DataT_, Dimensions, AccessMode_, AccessTarget,
1549  IsPlaceholder, PropertyListT> &other)
1550 #ifdef __SYCL_DEVICE_ONLY__
1551  : impl(other.impl), MData(other.MData) {
1552 #else
1553  : accessor(other.impl) {
1554 #endif // __SYCL_DEVICE_ONLY__
1555  }
1556 
1557  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1558  typename std::enable_if_t<
1559  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1560  std::is_same_v<T, DataT> && Dims == 0 &&
1561  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))> * =
1562  nullptr>
1563  accessor(
1564  buffer<T, 1, AllocatorT> &BufferRef,
1565  const property_list &PropertyList = {},
1566  const detail::code_location CodeLoc = detail::code_location::current())
1567 #ifdef __SYCL_DEVICE_ONLY__
1568  : impl(id<AdjustedDim>(), detail::GetZeroDimAccessRange(BufferRef),
1569  BufferRef.get_range()) {
1570  (void)PropertyList;
1571  (void)CodeLoc;
1572 #else
1573  : AccessorBaseHost(
1574  /*Offset=*/{0, 0, 0},
1575  detail::convertToArrayOfN<3, 1>(
1576  detail::GetZeroDimAccessRange(BufferRef)),
1577  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1578  getAdjustedMode(PropertyList),
1579  detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
1580  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1581  PropertyList) {
1582  throwIfUsedByGraph();
1583  preScreenAccessor(PropertyList);
1586  initHostAcc();
1589  AccessTarget, AccessMode, CodeLoc);
1590  GDBMethodsAnchor();
1591 #endif
1592  }
1593 
1594  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1595  typename... PropTypes,
1596  typename std::enable_if_t<
1598  // VS2019 can't compile sycl/test/regression/bit_cast_win.cpp
1599  // if std::is_same_v is used here.
1600  std::is_same<T, DataT>::value && Dims == 0 &&
1601  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))> * =
1602  nullptr>
1603  accessor(
1604  buffer<T, 1, AllocatorT> &BufferRef,
1605  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1606  {},
1607  const detail::code_location CodeLoc = detail::code_location::current())
1608 #ifdef __SYCL_DEVICE_ONLY__
1609  : impl(id<AdjustedDim>(), detail::GetZeroDimAccessRange(BufferRef),
1610  BufferRef.get_range()) {
1611  (void)PropertyList;
1612  (void)CodeLoc;
1613 #else
1614  : AccessorBaseHost(
1615  /*Offset=*/{0, 0, 0},
1616  detail::convertToArrayOfN<3, 1>(
1617  detail::GetZeroDimAccessRange(BufferRef)),
1618  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1619  getAdjustedMode(PropertyList),
1620  detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
1621  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1622  PropertyList) {
1623  throwIfUsedByGraph();
1624  preScreenAccessor(PropertyList);
1627  initHostAcc();
1630  AccessTarget, AccessMode, CodeLoc);
1631  GDBMethodsAnchor();
1632 #endif
1633  }
1634 
1635  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1636  typename = typename std::enable_if_t<
1637  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1638  std::is_same_v<T, DataT> && (Dims == 0) &&
1639  (IsGlobalBuf || IsHostBuf || IsConstantBuf || IsHostTask)>>
1640  accessor(
1641  buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1642  const property_list &PropertyList = {},
1643  const detail::code_location CodeLoc = detail::code_location::current())
1644 #ifdef __SYCL_DEVICE_ONLY__
1645  : impl(id<AdjustedDim>(), detail::GetZeroDimAccessRange(BufferRef),
1646  BufferRef.get_range()) {
1647  (void)CommandGroupHandler;
1648  (void)PropertyList;
1649  (void)CodeLoc;
1650  }
1651 #else
1652  : AccessorBaseHost(
1653  /*Offset=*/{0, 0, 0},
1654  detail::convertToArrayOfN<3, 1>(
1655  detail::GetZeroDimAccessRange(BufferRef)),
1656  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1657  getAdjustedMode(PropertyList),
1658  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1659  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1660  throwIfUsedByGraph();
1661  preScreenAccessor(PropertyList);
1662  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1663  initHostAcc();
1666  AccessTarget, AccessMode, CodeLoc);
1667  GDBMethodsAnchor();
1668  }
1669 #endif
1670 
1671  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1672  typename... PropTypes,
1673  typename = typename std::enable_if_t<
1675  std::is_same_v<T, DataT> && (Dims == 0) &&
1676  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1677  accessor(
1678  buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1679  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1680  {},
1681  const detail::code_location CodeLoc = detail::code_location::current())
1682 #ifdef __SYCL_DEVICE_ONLY__
1683  : impl(id<AdjustedDim>(), detail::GetZeroDimAccessRange(BufferRef),
1684  BufferRef.get_range()) {
1685  (void)CommandGroupHandler;
1686  (void)PropertyList;
1687  (void)CodeLoc;
1688  }
1689 #else
1690  : AccessorBaseHost(
1691  /*Offset=*/{0, 0, 0},
1692  detail::convertToArrayOfN<3, 1>(
1693  detail::GetZeroDimAccessRange(BufferRef)),
1694  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1695  getAdjustedMode(PropertyList),
1696  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1697  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1698  throwIfUsedByGraph();
1699  preScreenAccessor(PropertyList);
1700  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1701  initHostAcc();
1704  AccessTarget, AccessMode, CodeLoc);
1705  GDBMethodsAnchor();
1706  }
1707 #endif
1708 
1709  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1710  typename = std::enable_if_t<
1711  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1712  IsSameAsBuffer<T, Dims>::value &&
1713  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1714  accessor(
1715  buffer<T, Dims, AllocatorT> &BufferRef,
1716  const property_list &PropertyList = {},
1717  const detail::code_location CodeLoc = detail::code_location::current())
1718 #ifdef __SYCL_DEVICE_ONLY__
1719  : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1720  (void)PropertyList;
1721  (void)CodeLoc;
1722  }
1723 #else
1724  : AccessorBaseHost(
1725  /*Offset=*/{0, 0, 0},
1726  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1727  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1728  getAdjustedMode(PropertyList),
1729  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1730  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1731  PropertyList) {
1732  throwIfUsedByGraph();
1733  preScreenAccessor(PropertyList);
1736  initHostAcc();
1739  AccessTarget, AccessMode, CodeLoc);
1740  GDBMethodsAnchor();
1741  }
1742 #endif
1743 
1744  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1745  typename... PropTypes,
1746  typename = std::enable_if_t<
1748  IsSameAsBuffer<T, Dims>::value &&
1749  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1750  accessor(
1751  buffer<T, Dims, AllocatorT> &BufferRef,
1752  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1753  {},
1754  const detail::code_location CodeLoc = detail::code_location::current())
1755 #ifdef __SYCL_DEVICE_ONLY__
1756  : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1757  (void)PropertyList;
1758  (void)CodeLoc;
1759  }
1760 #else
1761  : AccessorBaseHost(
1762  /*Offset=*/{0, 0, 0},
1763  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1764  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1765  getAdjustedMode(PropertyList),
1766  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1767  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1768  PropertyList) {
1769  throwIfUsedByGraph();
1770  preScreenAccessor(PropertyList);
1773  initHostAcc();
1776  AccessTarget, AccessMode, CodeLoc);
1777  GDBMethodsAnchor();
1778  }
1779 #endif
1780 
1781  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1782  typename TagT,
1783  typename = std::enable_if_t<
1784  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1785  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1786  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1787  accessor(
1788  buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1789  const property_list &PropertyList = {},
1790  const detail::code_location CodeLoc = detail::code_location::current())
1791  : accessor(BufferRef, PropertyList, CodeLoc) {
1792  adjustAccPropsInBuf(BufferRef);
1793  }
1794 
1795  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1796  typename TagT, typename... PropTypes,
1797  typename = std::enable_if_t<
1799  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1800  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1801  accessor(
1802  buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1803  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1804  {},
1805  const detail::code_location CodeLoc = detail::code_location::current())
1806  : accessor(BufferRef, PropertyList, CodeLoc) {
1807  adjustAccPropsInBuf(BufferRef);
1808  }
1809 
1810  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1811  typename = std::enable_if_t<
1812  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1813  IsSameAsBuffer<T, Dims>::value &&
1814  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1815  accessor(
1816  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1817  const property_list &PropertyList = {},
1818  const detail::code_location CodeLoc = detail::code_location::current())
1819 #ifdef __SYCL_DEVICE_ONLY__
1820  : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1821  (void)CommandGroupHandler;
1822  (void)PropertyList;
1823  (void)CodeLoc;
1824  }
1825 #else
1826  : AccessorBaseHost(
1827  /*Offset=*/{0, 0, 0},
1828  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1829  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1830  getAdjustedMode(PropertyList),
1831  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1832  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1833  throwIfUsedByGraph();
1834  preScreenAccessor(PropertyList);
1835  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1836  initHostAcc();
1839  AccessTarget, AccessMode, CodeLoc);
1840  GDBMethodsAnchor();
1841  }
1842 #endif
1843 
1844  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1845  typename... PropTypes,
1846  typename = std::enable_if_t<
1848  IsSameAsBuffer<T, Dims>::value &&
1849  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1850  accessor(
1851  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1852  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1853  {},
1854  const detail::code_location CodeLoc = detail::code_location::current())
1855 #ifdef __SYCL_DEVICE_ONLY__
1856  : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1857  (void)CommandGroupHandler;
1858  (void)PropertyList;
1859  (void)CodeLoc;
1860  }
1861 #else
1862  : AccessorBaseHost(
1863  /*Offset=*/{0, 0, 0},
1864  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1865  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1866  getAdjustedMode(PropertyList),
1867  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1868  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1869  throwIfUsedByGraph();
1870  preScreenAccessor(PropertyList);
1871  initHostAcc();
1872  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1875  AccessTarget, AccessMode, CodeLoc);
1876  GDBMethodsAnchor();
1877  }
1878 #endif
1879 
1880  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1881  typename TagT,
1882  typename = std::enable_if_t<
1883  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1884  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1885  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1886  accessor(
1887  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1888  TagT, const property_list &PropertyList = {},
1889  const detail::code_location CodeLoc = detail::code_location::current())
1890  : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1891  adjustAccPropsInBuf(BufferRef);
1892  }
1893 
1894  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1895  typename TagT, typename... PropTypes,
1896  typename = std::enable_if_t<
1898  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1899  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1900  accessor(
1901  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1902  TagT,
1903  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1904  {},
1905  const detail::code_location CodeLoc = detail::code_location::current())
1906  : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1907  adjustAccPropsInBuf(BufferRef);
1908  }
1909 
1910  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1911  typename = std::enable_if_t<
1912  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1913  IsSameAsBuffer<T, Dims>::value &&
1914  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1915  accessor(
1916  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1917  const property_list &PropertyList = {},
1918  const detail::code_location CodeLoc = detail::code_location::current())
1919  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1920 
1921  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1922  typename... PropTypes,
1923  typename = std::enable_if_t<
1925  IsSameAsBuffer<T, Dims>::value &&
1926  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1927  accessor(
1928  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1929  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1930  {},
1931  const detail::code_location CodeLoc = detail::code_location::current())
1932  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1933 
1934  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1935  typename TagT,
1936  typename = std::enable_if_t<
1937  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1938  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1939  (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
1940  accessor(
1941  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1942  TagT, const property_list &PropertyList = {},
1943  const detail::code_location CodeLoc = detail::code_location::current())
1944  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1945  adjustAccPropsInBuf(BufferRef);
1946  }
1947 
1948  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1949  typename TagT, typename... PropTypes,
1950  typename = std::enable_if_t<
1952  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1953  (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
1954  accessor(
1955  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1956  TagT,
1957  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1958  {},
1959  const detail::code_location CodeLoc = detail::code_location::current())
1960  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1961  adjustAccPropsInBuf(BufferRef);
1962  }
1963 
1964  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1965  typename = std::enable_if_t<
1966  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1967  IsSameAsBuffer<T, Dims>::value &&
1968  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1969  accessor(
1970  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1971  range<Dimensions> AccessRange, const property_list &PropertyList = {},
1972  const detail::code_location CodeLoc = detail::code_location::current())
1973  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1974  CodeLoc) {}
1975 
1976  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1977  typename... PropTypes,
1978  typename = std::enable_if_t<
1980  IsSameAsBuffer<T, Dims>::value &&
1981  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1982  accessor(
1983  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1984  range<Dimensions> AccessRange,
1985  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1986  {},
1987  const detail::code_location CodeLoc = detail::code_location::current())
1988  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1989  CodeLoc) {}
1990 
1991  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1992  typename TagT,
1993  typename = std::enable_if_t<
1994  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1995  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1996  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1997  accessor(
1998  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1999  range<Dimensions> AccessRange, TagT,
2000  const property_list &PropertyList = {},
2001  const detail::code_location CodeLoc = detail::code_location::current())
2002  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
2003  CodeLoc) {
2004  adjustAccPropsInBuf(BufferRef);
2005  }
2006 
2007  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2008  typename TagT, typename... PropTypes,
2009  typename = std::enable_if_t<
2011  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2012  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2013  accessor(
2014  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2015  range<Dimensions> AccessRange, TagT,
2016  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2017  {},
2018  const detail::code_location CodeLoc = detail::code_location::current())
2019  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
2020  CodeLoc) {
2021  adjustAccPropsInBuf(BufferRef);
2022  }
2023 
2024  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2025  typename = std::enable_if_t<
2026  detail::IsRunTimePropertyListT<PropertyListT>::value &&
2027  IsSameAsBuffer<T, Dims>::value &&
2028  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
2029  accessor(
2030  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
2031  id<Dimensions> AccessOffset, const property_list &PropertyList = {},
2032  const detail::code_location CodeLoc = detail::code_location::current())
2033 #ifdef __SYCL_DEVICE_ONLY__
2034  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
2035  (void)PropertyList;
2036  (void)CodeLoc;
2037  }
2038 #else
2039  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
2040  detail::convertToArrayOfN<3, 1>(AccessRange),
2041  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
2042  getAdjustedMode(PropertyList),
2043  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
2044  sizeof(DataT), IsPlaceH, BufferRef.OffsetInBytes,
2045  BufferRef.IsSubBuffer, PropertyList) {
2046  throwIfUsedByGraph();
2047  preScreenAccessor(PropertyList);
2050  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2051  BufferRef.get_range()))
2052  throw sycl::invalid_object_error(
2053  "accessor with requested offset and range would exceed the bounds of "
2054  "the buffer",
2055  PI_ERROR_INVALID_VALUE);
2056 
2057  initHostAcc();
2060  AccessTarget, AccessMode, CodeLoc);
2061  GDBMethodsAnchor();
2062  }
2063 #endif
2064 
2065  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2066  typename... PropTypes,
2067  typename = std::enable_if_t<
2069  IsSameAsBuffer<T, Dims>::value &&
2070  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
2071  accessor(
2072  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
2073  id<Dimensions> AccessOffset,
2074  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2075  {},
2076  const detail::code_location CodeLoc = detail::code_location::current())
2077 #ifdef __SYCL_DEVICE_ONLY__
2078  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
2079  (void)PropertyList;
2080  (void)CodeLoc;
2081  }
2082 #else
2083  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
2084  detail::convertToArrayOfN<3, 1>(AccessRange),
2085  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
2086  getAdjustedMode(PropertyList),
2087  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
2088  sizeof(DataT), IsPlaceH, BufferRef.OffsetInBytes,
2089  BufferRef.IsSubBuffer, PropertyList) {
2090  throwIfUsedByGraph();
2091  preScreenAccessor(PropertyList);
2094  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2095  BufferRef.get_range()))
2096  throw sycl::invalid_object_error(
2097  "accessor with requested offset and range would exceed the bounds of "
2098  "the buffer",
2099  PI_ERROR_INVALID_VALUE);
2100 
2101  initHostAcc();
2104  AccessTarget, AccessMode, CodeLoc);
2105  GDBMethodsAnchor();
2106  }
2107 #endif
2108 
2109  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2110  typename TagT,
2111  typename = std::enable_if_t<
2112  detail::IsRunTimePropertyListT<PropertyListT>::value &&
2113  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2114  (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
2115  accessor(
2116  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
2117  id<Dimensions> AccessOffset, TagT, const property_list &PropertyList = {},
2118  const detail::code_location CodeLoc = detail::code_location::current())
2119  : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
2120  adjustAccPropsInBuf(BufferRef);
2121  }
2122 
2123  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2124  typename TagT, typename... PropTypes,
2125  typename = std::enable_if_t<
2127  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2128  (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
2129  accessor(
2130  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
2131  id<Dimensions> AccessOffset, TagT,
2132  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2133  {},
2134  const detail::code_location CodeLoc = detail::code_location::current())
2135  : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
2136  adjustAccPropsInBuf(BufferRef);
2137  }
2138 
2139  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2140  typename = std::enable_if_t<
2141  detail::IsRunTimePropertyListT<PropertyListT>::value &&
2142  IsSameAsBuffer<T, Dims>::value &&
2143  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2144  accessor(
2145  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2146  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
2147  const property_list &PropertyList = {},
2148  const detail::code_location CodeLoc = detail::code_location::current())
2149 #ifdef __SYCL_DEVICE_ONLY__
2150  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
2151  (void)CommandGroupHandler;
2152  (void)PropertyList;
2153  (void)CodeLoc;
2154  }
2155 #else
2156  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
2157  detail::convertToArrayOfN<3, 1>(AccessRange),
2158  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
2159  getAdjustedMode(PropertyList),
2160  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
2161  sizeof(DataT), BufferRef.OffsetInBytes,
2162  BufferRef.IsSubBuffer, PropertyList) {
2163  throwIfUsedByGraph();
2164  preScreenAccessor(PropertyList);
2165  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2166  BufferRef.get_range()))
2167  throw sycl::invalid_object_error(
2168  "accessor with requested offset and range would exceed the bounds of "
2169  "the buffer",
2170  PI_ERROR_INVALID_VALUE);
2171 
2172  initHostAcc();
2173  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
2176  AccessTarget, AccessMode, CodeLoc);
2177  GDBMethodsAnchor();
2178  }
2179 #endif
2180 
2181  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2182  typename... PropTypes,
2183  typename = std::enable_if_t<
2185  IsSameAsBuffer<T, Dims>::value &&
2186  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2187  accessor(
2188  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2189  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
2190  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2191  {},
2192  const detail::code_location CodeLoc = detail::code_location::current())
2193 #ifdef __SYCL_DEVICE_ONLY__
2194  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
2195  (void)CommandGroupHandler;
2196  (void)PropertyList;
2197  (void)CodeLoc;
2198  }
2199 #else
2200  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
2201  detail::convertToArrayOfN<3, 1>(AccessRange),
2202  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
2203  getAdjustedMode(PropertyList),
2204  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
2205  sizeof(DataT), BufferRef.OffsetInBytes,
2206  BufferRef.IsSubBuffer, PropertyList) {
2207  throwIfUsedByGraph();
2208  preScreenAccessor(PropertyList);
2209  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2210  BufferRef.get_range()))
2211  throw sycl::invalid_object_error(
2212  "accessor with requested offset and range would exceed the bounds of "
2213  "the buffer",
2214  PI_ERROR_INVALID_VALUE);
2215 
2216  initHostAcc();
2217  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
2220  AccessTarget, AccessMode, CodeLoc);
2221  GDBMethodsAnchor();
2222  }
2223 #endif
2224 
2225  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2226  typename TagT,
2227  typename = std::enable_if_t<
2228  detail::IsRunTimePropertyListT<PropertyListT>::value &&
2229  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2230  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2231  accessor(
2232  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2233  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
2234  const property_list &PropertyList = {},
2235  const detail::code_location CodeLoc = detail::code_location::current())
2236  : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2237  PropertyList, CodeLoc) {
2238  adjustAccPropsInBuf(BufferRef);
2239  }
2240 
2241  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2242  typename TagT, typename... PropTypes,
2243  typename = std::enable_if_t<
2245  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2246  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2247  accessor(
2248  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2249  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
2250  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2251  {},
2252  const detail::code_location CodeLoc = detail::code_location::current())
2253  : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2254  PropertyList, CodeLoc) {
2255  adjustAccPropsInBuf(BufferRef);
2256  }
2257 
2258  template <typename... NewPropsT>
2259  accessor(
2260  const accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
2261  ext::oneapi::accessor_property_list<NewPropsT...>> &Other,
2262  const detail::code_location CodeLoc = detail::code_location::current())
2263 #ifdef __SYCL_DEVICE_ONLY__
2264  : impl(Other.impl), MData(Other.MData)
2265 #else
2266  : detail::AccessorBaseHost(Other), MAccData(Other.MAccData)
2267 #endif
2268  {
2270  "Conversion is only available for accessor_property_list");
2271  static_assert(
2272  PropertyListT::template areSameCompileTimeProperties<NewPropsT...>(),
2273  "Compile-time-constant properties must be the same");
2274  (void)CodeLoc;
2275 #ifndef __SYCL_DEVICE_ONLY__
2276  detail::constructorNotification(getMemoryObject(), impl.get(), AccessTarget,
2277  AccessMode, CodeLoc);
2278 #endif
2279  }
2280 
2281  void swap(accessor &other) {
2282  std::swap(impl, other.impl);
2283 #ifdef __SYCL_DEVICE_ONLY__
2284  std::swap(MData, other.MData);
2285 #else
2286  std::swap(MAccData, other.MAccData);
2287 #endif
2288  }
2289 
2290  bool is_placeholder() const {
2291 #ifdef __SYCL_DEVICE_ONLY__
2292  return false;
2293 #else
2295 #endif
2296  }
2297 
2298  size_t get_size() const { return getAccessRange().size() * sizeof(DataT); }
2299 
2300  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
2301  size_t get_count() const { return size(); }
2302  size_type size() const noexcept { return getAccessRange().size(); }
2303 
2304  size_type byte_size() const noexcept { return size() * sizeof(DataT); }
2305 
2306  size_type max_size() const noexcept {
2308  }
2309 
2310  bool empty() const noexcept { return size() == 0; }
2311 
2312  template <int Dims = Dimensions,
2313  typename = std::enable_if_t<Dims == Dimensions && (Dims > 0)>>
2314  range<Dimensions> get_range() const {
2315  return getRange<Dims>();
2316  }
2317 
2318  template <int Dims = Dimensions,
2319  typename = std::enable_if_t<Dims == Dimensions && (Dims > 0)>>
2320  id<Dimensions> get_offset() const {
2321  return getOffset<Dims>();
2322  }
2323 
2324  template <int Dims = Dimensions, typename RefT = RefType,
2325  typename = std::enable_if_t<Dims == 0 &&
2326  (IsAccessAnyWrite || IsAccessReadOnly)>>
2327  operator reference() const {
2328  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
2329  return *(getQualifiedPtr() + LinearIndex);
2330  }
2331 
2332  template <int Dims = Dimensions,
2333  typename = std::enable_if_t<AccessMode != access_mode::atomic &&
2334  !IsAccessReadOnly && Dims == 0>>
2335  const accessor &operator=(const value_type &Other) const {
2336  *getQualifiedPtr() = Other;
2337  return *this;
2338  }
2339 
2340  template <int Dims = Dimensions,
2341  typename = std::enable_if_t<AccessMode != access_mode::atomic &&
2342  !IsAccessReadOnly && Dims == 0>>
2343  const accessor &operator=(value_type &&Other) const {
2344  *getQualifiedPtr() = std::move(Other);
2345  return *this;
2346  }
2347 
2348  template <int Dims = Dimensions,
2349  typename = std::enable_if_t<(Dims > 0) &&
2350  (IsAccessAnyWrite || IsAccessReadOnly)>>
2351  reference operator[](id<Dimensions> Index) const {
2352  const size_t LinearIndex = getLinearIndex(Index);
2353  return getQualifiedPtr()[LinearIndex];
2354  }
2355 
2356  template <int Dims = Dimensions>
2357  operator typename std::enable_if_t<Dims == 0 &&
2359 #ifdef __ENABLE_USM_ADDR_SPACE__
2360  atomic<DataT>
2361 #else
2362  atomic<DataT, AS>
2363 #endif
2364  >() const {
2365  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
2366  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2367  getQualifiedPtr() + LinearIndex));
2368  }
2369 
2370  template <int Dims = Dimensions>
2371  typename std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
2372  atomic<DataT, AS>>
2373  operator[](id<Dimensions> Index) const {
2374  const size_t LinearIndex = getLinearIndex(Index);
2375  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2376  getQualifiedPtr() + LinearIndex));
2377  }
2378 
2379  template <int Dims = Dimensions>
2380  typename std::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
2381  atomic<DataT, AS>>
2382  operator[](size_t Index) const {
2383  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
2384  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2385  getQualifiedPtr() + LinearIndex));
2386  }
2387  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 1)>>
2388  auto operator[](size_t Index) const {
2389  return AccessorSubscript<Dims - 1>(*this, Index);
2390  }
2391 
2392  template <access::target AccessTarget_ = AccessTarget,
2393  typename = std::enable_if_t<
2394  (AccessTarget_ == access::target::host_buffer) ||
2395  (AccessTarget_ == access::target::host_task)>>
2396  std::add_pointer_t<value_type> get_pointer() const noexcept {
2397  return getPointerAdjusted();
2398  }
2399 
2400  template <
2401  access::target AccessTarget_ = AccessTarget,
2402  typename = std::enable_if_t<(AccessTarget_ == access::target::device)>>
2404  "accessor::get_pointer() is deprecated, please use get_multi_ptr()")
2405  global_ptr<DataT> get_pointer() const noexcept {
2406  return global_ptr<DataT>(
2407  const_cast<typename detail::DecoratedType<DataT, AS>::type *>(
2408  getPointerAdjusted()));
2409  }
2410 
2411  template <access::target AccessTarget_ = AccessTarget,
2412  typename = std::enable_if_t<AccessTarget_ ==
2413  access::target::constant_buffer>>
2414  constant_ptr<DataT> get_pointer() const {
2415  return constant_ptr<DataT>(getPointerAdjusted());
2416  }
2417 
2418  template <access::decorated IsDecorated>
2419  accessor_ptr<IsDecorated> get_multi_ptr() const noexcept {
2420  return accessor_ptr<IsDecorated>(getPointerAdjusted());
2421  }
2422 
2423  // accessor::has_property for runtime properties is only available in host
2424  // code. This restriction is not listed in the core spec and will be added in
2425  // future versions.
2426  template <typename Property>
2427  typename std::enable_if_t<
2428  !ext::oneapi::is_compile_time_property<Property>::value, bool>
2429  has_property() const noexcept {
2430 #ifndef __SYCL_DEVICE_ONLY__
2431  return getPropList().template has_property<Property>();
2432 #else
2433  return false;
2434 #endif
2435  }
2436 
2437  // accessor::get_property for runtime properties is only available in host
2438  // code. This restriction is not listed in the core spec and will be added in
2439  // future versions.
2440  template <typename Property,
2441  typename = typename std::enable_if_t<
2442  !ext::oneapi::is_compile_time_property<Property>::value>>
2443  Property get_property() const {
2444 #ifndef __SYCL_DEVICE_ONLY__
2445  return getPropList().template get_property<Property>();
2446 #else
2447  return Property();
2448 #endif
2449  }
2450 
2451  template <typename Property>
2452  static constexpr bool has_property(
2453  typename std::enable_if_t<
2454  ext::oneapi::is_compile_time_property<Property>::value> * = 0) {
2455  return PropertyListT::template has_property<Property>();
2456  }
2457 
2458  template <typename Property>
2459  static constexpr auto get_property(
2460  typename std::enable_if_t<
2461  ext::oneapi::is_compile_time_property<Property>::value> * = 0) {
2462  return PropertyListT::template get_property<Property>();
2463  }
2464 
2465  bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
2466  bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
2467 
2468  iterator begin() const noexcept {
2469  return iterator::getBegin(
2470  get_pointer(),
2471  detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2472  getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2473  }
2474 
2475  iterator end() const noexcept {
2476  return iterator::getEnd(
2477  get_pointer(),
2478  detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2479  getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2480  }
2481 
2482  const_iterator cbegin() const noexcept {
2483  return const_iterator::getBegin(
2484  get_pointer(),
2485  detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2486  getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2487  }
2488 
2489  const_iterator cend() const noexcept {
2490  return const_iterator::getEnd(
2491  get_pointer(),
2492  detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2493  getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2494  }
2495 
2496  reverse_iterator rbegin() const noexcept { return reverse_iterator(end()); }
2497  reverse_iterator rend() const noexcept { return reverse_iterator(begin()); }
2498 
2499  const_reverse_iterator crbegin() const noexcept {
2500  return const_reverse_iterator(cend());
2501  }
2502  const_reverse_iterator crend() const noexcept {
2503  return const_reverse_iterator(cbegin());
2504  }
2505 
2506 private:
2507  template <int Dims, typename = std::enable_if_t<(Dims > 0)>>
2508  range<Dims> getRange() const {
2509  return detail::convertToArrayOfN<AdjustedDim, 1>(getAccessRange());
2510  }
2511 
2512  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
2513  id<Dims> getOffset() const {
2514  static_assert(
2515  !(PropertyListT::template has_property<
2517  "Accessor has no_offset property, get_offset() can not be used");
2518  return detail::convertToArrayOfN<Dims, 0>(getOffset());
2519  }
2520 
2521 #ifdef __SYCL_DEVICE_ONLY__
2522  size_t getTotalOffset() const noexcept {
2523  size_t TotalOffset = 0;
2524  detail::loop<Dimensions>([&, this](size_t I) {
2525  TotalOffset = TotalOffset * impl.MemRange[I];
2526  if constexpr (!(PropertyListT::template has_property<
2528  TotalOffset += impl.Offset[I];
2529  }
2530  });
2531 
2532  return TotalOffset;
2533  }
2534 #endif
2535 
2536  // supporting function for get_pointer()
2537  // MData has been preadjusted with offset for faster access with []
2538  // but for get_pointer() we must return the original pointer.
2539  // On device, getQualifiedPtr() returns MData, so we need to backjust it.
2540  // On host, getQualifiedPtr() does not return MData, no need to adjust.
2541  auto getPointerAdjusted() const noexcept {
2542 #ifdef __SYCL_DEVICE_ONLY__
2543  return getQualifiedPtr() - getTotalOffset();
2544 #else
2545  return getQualifiedPtr();
2546 #endif
2547  }
2548 
2549  void preScreenAccessor(const PropertyListT &PropertyList) {
2550  // check that no_init property is compatible with access mode
2551  if (PropertyList.template has_property<property::no_init>() &&
2553  throw sycl::invalid_object_error(
2554  "accessor would cannot be both read_only and no_init",
2555  PI_ERROR_INVALID_VALUE);
2556  }
2557  }
2558 
2559  template <typename BufT, typename... PropTypes>
2560  void adjustAccPropsInBuf(BufT &Buffer) {
2561  if constexpr (PropertyListT::template has_property<
2563  auto location = (PropertyListT::template get_property<
2565  .get_location();
2566  property_list PropList{
2568  Buffer.addOrReplaceAccessorProperties(PropList);
2569  } else {
2570  deleteAccPropsFromBuf(Buffer);
2571  }
2572  }
2573 
2574  template <typename BufT> void deleteAccPropsFromBuf(BufT &Buffer) {
2575  Buffer.deleteAccProps(
2577  }
2578 };
2579 
2580 template <typename DataT, int Dimensions, typename AllocatorT>
2582  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2584 
2585 template <typename DataT, int Dimensions, typename AllocatorT,
2586  typename... PropsT>
2589  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2592 
2593 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
2596  detail::deduceAccessTarget<Type1, Type1>(target::device),
2598 
2599 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2600  typename... PropsT>
2604  detail::deduceAccessTarget<Type1, Type1>(target::device),
2607 
2608 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2609  typename Type2>
2612  detail::deduceAccessTarget<Type1, Type2>(target::device),
2614 
2615 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2616  typename Type2, typename... PropsT>
2620  detail::deduceAccessTarget<Type1, Type2>(target::device),
2623 
2624 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2625  typename Type2, typename Type3>
2628  detail::deduceAccessTarget<Type2, Type3>(target::device),
2630 
2631 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2632  typename Type2, typename Type3, typename... PropsT>
2636  detail::deduceAccessTarget<Type2, Type3>(target::device),
2639 
2640 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2641  typename Type2, typename Type3, typename Type4>
2644  detail::deduceAccessTarget<Type3, Type4>(target::device),
2646 
2647 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2648  typename Type2, typename Type3, typename Type4, typename... PropsT>
2652  detail::deduceAccessTarget<Type3, Type4>(target::device),
2655 
2656 template <typename DataT, int Dimensions, typename AllocatorT>
2658  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2660 
2661 template <typename DataT, int Dimensions, typename AllocatorT,
2662  typename... PropsT>
2665  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2668 
2669 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
2672  detail::deduceAccessTarget<Type1, Type1>(target::device),
2674 
2675 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2676  typename... PropsT>
2680  detail::deduceAccessTarget<Type1, Type1>(target::device),
2683 
2684 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2685  typename Type2>
2688  detail::deduceAccessTarget<Type1, Type2>(target::device),
2690 
2691 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2692  typename Type2, typename... PropsT>
2696  detail::deduceAccessTarget<Type1, Type2>(target::device),
2699 
2700 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2701  typename Type2, typename Type3>
2704  detail::deduceAccessTarget<Type2, Type3>(target::device),
2706 
2707 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2708  typename Type2, typename Type3, typename... PropsT>
2712  detail::deduceAccessTarget<Type2, Type3>(target::device),
2715 
2716 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2717  typename Type2, typename Type3, typename Type4>
2719  Type4)
2721  detail::deduceAccessTarget<Type3, Type4>(target::device),
2723 
2724 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2725  typename Type2, typename Type3, typename Type4, typename... PropsT>
2729  detail::deduceAccessTarget<Type3, Type4>(target::device),
2732 
2736 template <typename DataT, int Dimensions, access::mode AccessMode,
2739 #ifndef __SYCL_DEVICE_ONLY__
2741 #endif
2742  public detail::accessor_common<DataT, Dimensions, AccessMode,
2743  access::target::local, IsPlaceholder> {
2744 protected:
2745  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
2746 
2749  access::target::local, IsPlaceholder>;
2750 
2751  using AccessorCommonT::AS;
2752 
2753  // Cannot do "using AccessorCommonT::Flag" as it doesn't work with g++ as host
2754  // compiler, for some reason.
2755  static constexpr bool IsAccessAnyWrite = AccessorCommonT::IsAccessAnyWrite;
2756  static constexpr bool IsAccessReadOnly = AccessorCommonT::IsAccessReadOnly;
2757  static constexpr bool IsConst = AccessorCommonT::IsConst;
2758 
2759  template <int Dims>
2761  typename AccessorCommonT::template AccessorSubscript<
2762  Dims,
2764 
2766 
2769 
2770 #ifdef __SYCL_DEVICE_ONLY__
2772 
2773  sycl::range<AdjustedDim> &getSize() { return impl.MemRange; }
2774  const sycl::range<AdjustedDim> &getSize() const { return impl.MemRange; }
2775 
2776  void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
2778  MData = Ptr;
2779  detail::loop<AdjustedDim>(
2780  [&, this](size_t I) { getSize()[I] = AccessRange[I]; });
2781  }
2782 
2783  // __init variant used by the device compiler for ESIMD kernels.
2784  // TODO: In ESIMD accessors usage is limited for now - access range, mem
2785  // range and offset are not supported.
2786  void __init_esimd(ConcreteASPtrType Ptr) {
2787  MData = Ptr;
2788  detail::loop<AdjustedDim>([&, this](size_t I) { getSize()[I] = 0; });
2789  }
2790 
2791 public:
2792  // Default constructor for objects later initialized with __init member.
2793  local_accessor_base()
2794  : impl(detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
2795 
2796 protected:
2797  ConcreteASPtrType getQualifiedPtr() const { return MData; }
2798 
2799  ConcreteASPtrType MData;
2800 
2801 #else
2802 public:
2804  : detail::LocalAccessorBaseHost{/*Size*/ sycl::range<3>{0, 0, 0},
2805  /*Dims*/ 0, /*ElemSize*/ sizeof(DataT)} {}
2806 
2807 protected:
2809  : detail::LocalAccessorBaseHost{Impl} {}
2810 
2812  sizeof(PtrType) - sizeof(detail::LocalAccessorBaseHost)];
2814 
2816  return reinterpret_cast<PtrType>(LocalAccessorBaseHost::getPtr());
2817  }
2818 
2820  void *getPtr() const { return detail::LocalAccessorBaseHost::getPtr(); }
2821  const range<3> &getSize() const {
2823  }
2825 
2826  // The function references helper methods required by GDB pretty-printers
2828 #ifndef NDEBUG
2829  const auto *this_const = this;
2830  (void)getSize();
2831  (void)this_const->getSize();
2832  (void)getPtr();
2833  (void)this_const->getPtr();
2834 #endif
2835  }
2836 
2837 #endif // __SYCL_DEVICE_ONLY__
2838 
2839  // Method which calculates linear offset for the ID using Range and Offset.
2840  template <int Dims = AdjustedDim> size_t getLinearIndex(id<Dims> Id) const {
2841  size_t Result = 0;
2842  detail::loop<Dims>(
2843  [&, this](size_t I) { Result = Result * getSize()[I] + Id[I]; });
2844  return Result;
2845  }
2846 
2847  template <class Obj>
2848  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
2849 
2850  template <class T>
2851  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
2852 
2853  template <typename DataT_, int Dimensions_> friend class local_accessor;
2854 
2855 public:
2856  using value_type = DataT;
2857  using reference = DataT &;
2858  using const_reference = const DataT &;
2859 
2860  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 0>>
2863 #ifdef __SYCL_DEVICE_ONLY__
2864  : impl(range<AdjustedDim>{1}) {
2865  (void)CodeLoc;
2866  }
2867 #else
2868  : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT)) {
2870  access::target::local, AccessMode, CodeLoc);
2871  GDBMethodsAnchor();
2872  }
2873 #endif
2874 
2875  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 0>>
2877  const detail::code_location CodeLoc =
2879 #ifdef __SYCL_DEVICE_ONLY__
2880  : impl(range<AdjustedDim>{1}) {
2881  (void)propList;
2882  (void)CodeLoc;
2883  }
2884 #else
2885  : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT),
2886  propList) {
2888  access::target::local, AccessMode, CodeLoc);
2889  GDBMethodsAnchor();
2890  }
2891 #endif
2892 
2893  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
2895  range<Dimensions> AllocationSize, handler &,
2897 #ifdef __SYCL_DEVICE_ONLY__
2898  : impl(AllocationSize) {
2899  (void)CodeLoc;
2900  }
2901 #else
2902  : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2903  AdjustedDim, sizeof(DataT)) {
2905  access::target::local, AccessMode, CodeLoc);
2906  GDBMethodsAnchor();
2907  }
2908 #endif
2909 
2910  template <int Dims = Dimensions,
2911  typename = std::enable_if_t<(Dims > 0)>>
2913  const property_list &propList,
2914  const detail::code_location CodeLoc =
2916 #ifdef __SYCL_DEVICE_ONLY__
2917  : impl(AllocationSize) {
2918  (void)propList;
2919  (void)CodeLoc;
2920  }
2921 #else
2922  : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2923  AdjustedDim, sizeof(DataT), propList) {
2925  access::target::local, AccessMode, CodeLoc);
2926  GDBMethodsAnchor();
2927  }
2928 #endif
2929 
2930  size_t get_size() const { return getSize().size() * sizeof(DataT); }
2931 
2932  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
2933  size_t get_count() const { return size(); }
2934  size_t size() const noexcept { return getSize().size(); }
2935 
2936  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
2938  return detail::convertToArrayOfN<Dims, 1>(getSize());
2939  }
2940 
2941  template <int Dims = Dimensions,
2942  typename = std::enable_if_t<Dims == 0 &&
2943  (IsAccessAnyWrite || IsAccessReadOnly)>>
2944  operator RefType() const {
2945  return *getQualifiedPtr();
2946  }
2947 
2948  template <int Dims = Dimensions,
2949  typename = std::enable_if_t<(Dims > 0) &&
2950  (IsAccessAnyWrite || IsAccessReadOnly)>>
2952  const size_t LinearIndex = getLinearIndex(Index);
2953  return getQualifiedPtr()[LinearIndex];
2954  }
2955 
2956  template <int Dims = Dimensions,
2957  typename = std::enable_if_t<Dims == 1 &&
2958  (IsAccessAnyWrite || IsAccessReadOnly)>>
2959  RefType operator[](size_t Index) const {
2960  return getQualifiedPtr()[Index];
2961  }
2962 
2963  template <int Dims = Dimensions>
2964  operator typename std::enable_if_t<
2965  Dims == 0 && AccessMode == access::mode::atomic, atomic<DataT, AS>>()
2966  const {
2967  return atomic<DataT, AS>(
2969  }
2970 
2971  template <int Dims = Dimensions>
2972  typename std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
2973  atomic<DataT, AS>>
2974  operator[](id<Dimensions> Index) const {
2975  const size_t LinearIndex = getLinearIndex(Index);
2976  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2977  getQualifiedPtr() + LinearIndex));
2978  }
2979 
2980  template <int Dims = Dimensions>
2981  typename std::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
2982  atomic<DataT, AS>>
2983  operator[](size_t Index) const {
2984  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2985  getQualifiedPtr() + Index));
2986  }
2987 
2988  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 1)>>
2989  typename AccessorCommonT::template AccessorSubscript<
2990  Dims - 1,
2992  operator[](size_t Index) const {
2993  return AccessorSubscript<Dims - 1>(*this, Index);
2994  }
2995 
2996  bool operator==(const local_accessor_base &Rhs) const {
2997  return impl == Rhs.impl;
2998  }
2999  bool operator!=(const local_accessor_base &Rhs) const {
3000  return !(*this == Rhs);
3001  }
3002 };
3003 
3004 // TODO: Remove deprecated specialization once no longer needed
3005 template <typename DataT, int Dimensions, access::mode AccessMode,
3008  DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder>
3009  : public local_accessor_base<DataT, Dimensions, AccessMode, IsPlaceholder>,
3010  public detail::OwnerLessBase<
3011  accessor<DataT, Dimensions, AccessMode, access::target::local,
3012  IsPlaceholder>> {
3013 
3014  using local_acc =
3016 
3017  static_assert(
3018  !local_acc::IsConst || local_acc::IsAccessReadOnly,
3019  "A const qualified DataT is only allowed for a read-only accessor");
3020 
3021  // Use base classes constructors
3022  using local_acc::local_acc;
3023 
3024 public:
3026  return local_ptr<DataT>(local_acc::getQualifiedPtr());
3027  }
3028 
3029 #ifdef __SYCL_DEVICE_ONLY__
3030 
3031  // __init needs to be defined within the class not through inheritance.
3032  // Map this function to inherited func.
3033  void __init(typename local_acc::ConcreteASPtrType Ptr,
3034  range<local_acc::AdjustedDim> AccessRange,
3037  local_acc::__init(Ptr, AccessRange, range, id);
3038  }
3039 
3040  // __init variant used by the device compiler for ESIMD kernels.
3041  // TODO: In ESIMD accessors usage is limited for now - access range, mem
3042  // range and offset are not supported.
3043  void __init_esimd(typename local_acc::ConcreteASPtrType Ptr) {
3044  local_acc::__init_esimd(Ptr);
3045  }
3046 
3047 public:
3048  // Default constructor for objects later initialized with __init member.
3049  accessor() {
3050  local_acc::impl = detail::InitializedVal<local_acc::AdjustedDim,
3051  range>::template get<0>();
3052  }
3053 
3054 #else
3055 private:
3056  accessor(const detail::AccessorImplPtr &Impl) : local_acc{Impl} {}
3057 #endif
3058 };
3059 
3060 template <typename DataT, int Dimensions = 1>
3061 class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
3062  : public local_accessor_base<DataT, Dimensions,
3063  detail::accessModeFromConstness<DataT>(),
3064  access::placeholder::false_t>,
3065  public detail::OwnerLessBase<local_accessor<DataT, Dimensions>> {
3066 
3067  using local_acc =
3068  local_accessor_base<DataT, Dimensions,
3069  detail::accessModeFromConstness<DataT>(),
3071 
3072  static_assert(
3073  !local_acc::IsConst || local_acc::IsAccessReadOnly,
3074  "A const qualified DataT is only allowed for a read-only accessor");
3075 
3076  // Use base classes constructors
3077  using local_acc::local_acc;
3078 
3079 #ifdef __SYCL_DEVICE_ONLY__
3080 
3081  // __init needs to be defined within the class not through inheritance.
3082  // Map this function to inherited func.
3083  void __init(typename local_acc::ConcreteASPtrType Ptr,
3084  range<local_acc::AdjustedDim> AccessRange,
3085  range<local_acc::AdjustedDim> range,
3086  id<local_acc::AdjustedDim> id) {
3087  local_acc::__init(Ptr, AccessRange, range, id);
3088  }
3089 
3090  // __init variant used by the device compiler for ESIMD kernels.
3091  // TODO: In ESIMD accessors usage is limited for now - access range, mem
3092  // range and offset are not supported.
3093  void __init_esimd(typename local_acc::ConcreteASPtrType Ptr) {
3094  local_acc::__init_esimd(Ptr);
3095  }
3096 
3097 public:
3098  // Default constructor for objects later initialized with __init member.
3099  local_accessor() {
3100  local_acc::impl = detail::InitializedVal<local_acc::AdjustedDim,
3101  range>::template get<0>();
3102  }
3103 
3104 #else
3105  local_accessor(const detail::AccessorImplPtr &Impl) : local_acc{Impl} {}
3106 #endif
3107 
3108  // implicit conversion between non-const read-write accessor to const
3109  // read-only accessor
3110 public:
3111  template <typename DataT_,
3112  typename = std::enable_if_t<
3113  std::is_const_v<DataT> &&
3114  std::is_same_v<DataT_, std::remove_const_t<DataT>>>>
3115  local_accessor(const local_accessor<DataT_, Dimensions> &other) {
3116  local_acc::impl = other.impl;
3117 #ifdef __SYCL_DEVICE_ONLY__
3118  local_acc::MData = other.MData;
3119 #endif
3120  }
3121 
3122  using value_type = DataT;
3123  using iterator = value_type *;
3124  using const_iterator = const value_type *;
3125  using reverse_iterator = std::reverse_iterator<iterator>;
3126  using const_reverse_iterator = std::reverse_iterator<const_iterator>;
3127  using difference_type =
3129  using size_type = std::size_t;
3130 
3131  template <access::decorated IsDecorated>
3132  using accessor_ptr = local_ptr<value_type, IsDecorated>;
3133 
3134  template <typename DataT_>
3135  bool operator==(const local_accessor<DataT_, Dimensions> &Rhs) const {
3136  return local_acc::impl == Rhs.impl;
3137  }
3138 
3139  template <typename DataT_>
3140  bool operator!=(const local_accessor<DataT_, Dimensions> &Rhs) const {
3141  return !(*this == Rhs);
3142  }
3143 
3144  void swap(local_accessor &other) { std::swap(this->impl, other.impl); }
3145 
3146  size_type byte_size() const noexcept { return this->size() * sizeof(DataT); }
3147 
3148  size_type max_size() const noexcept {
3150  }
3151 
3152  bool empty() const noexcept { return this->size() == 0; }
3153 
3154  iterator begin() const noexcept {
3155  if constexpr (Dimensions == 0)
3156  return local_acc::getQualifiedPtr();
3157  else
3158  return &this->operator[](id<Dimensions>());
3159  }
3160  iterator end() const noexcept {
3161  if constexpr (Dimensions == 0)
3162  return begin() + 1;
3163  else
3164  return begin() + this->size();
3165  }
3166 
3167  const_iterator cbegin() const noexcept { return const_iterator(begin()); }
3168  const_iterator cend() const noexcept { return const_iterator(end()); }
3169 
3170  reverse_iterator rbegin() const noexcept { return reverse_iterator(end()); }
3171  reverse_iterator rend() const noexcept { return reverse_iterator(begin()); }
3172 
3173  const_reverse_iterator crbegin() const noexcept {
3174  return const_reverse_iterator(end());
3175  }
3176  const_reverse_iterator crend() const noexcept {
3177  return const_reverse_iterator(begin());
3178  }
3179 
3181  "local_accessor::get_pointer() is deprecated, please use get_multi_ptr()")
3182  local_ptr<DataT> get_pointer() const noexcept {
3183  return local_ptr<DataT>(local_acc::getQualifiedPtr());
3184  }
3185 
3186  template <access::decorated IsDecorated>
3187  accessor_ptr<IsDecorated> get_multi_ptr() const noexcept {
3188  return accessor_ptr<IsDecorated>(local_acc::getQualifiedPtr());
3189  }
3190 
3191  template <typename Property> bool has_property() const noexcept {
3192 #ifndef __SYCL_DEVICE_ONLY__
3193  return this->getPropList().template has_property<Property>();
3194 #else
3195  return false;
3196 #endif
3197  }
3198 
3199  template <typename Property> Property get_property() const {
3200 #ifndef __SYCL_DEVICE_ONLY__
3201  return this->getPropList().template get_property<Property>();
3202 #else
3203  return Property();
3204 #endif
3205  }
3206 
3207  template <int Dims = Dimensions,
3208  typename = std::enable_if_t<!std::is_const_v<DataT> && Dims == 0>>
3209  const local_accessor &operator=(const value_type &Other) const {
3210  *local_acc::getQualifiedPtr() = Other;
3211  return *this;
3212  }
3213 
3214  template <int Dims = Dimensions,
3215  typename = std::enable_if_t<!std::is_const_v<DataT> && Dims == 0>>
3216  const local_accessor &operator=(value_type &&Other) const {
3217  *local_acc::getQualifiedPtr() = std::move(Other);
3218  return *this;
3219  }
3220 
3221 private:
3222  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
3223 };
3224 
3230 template <typename DataT, int Dimensions, access::mode AccessMode,
3233  DataT, Dimensions, AccessMode, access::target::image, IsPlaceholder>
3234  : public detail::image_accessor<DataT, Dimensions, AccessMode,
3235  access::target::image, IsPlaceholder>,
3236  public detail::OwnerLessBase<
3237  accessor<DataT, Dimensions, AccessMode, access::target::image,
3238  IsPlaceholder>> {
3239 private:
3240  accessor(const detail::AccessorImplPtr &Impl)
3241  : detail::image_accessor<DataT, Dimensions, AccessMode,
3243 
3244 public:
3245  template <typename AllocatorT>
3247  handler &CommandGroupHandler)
3248  : detail::image_accessor<DataT, Dimensions, AccessMode,
3250  Image, CommandGroupHandler, Image.getElementSize()) {
3251 #ifndef __SYCL_DEVICE_ONLY__
3252  detail::associateWithHandler(CommandGroupHandler, this,
3254 #endif
3255  }
3256 
3257  template <typename AllocatorT>
3259  handler &CommandGroupHandler, const property_list &propList)
3260  : detail::image_accessor<DataT, Dimensions, AccessMode,
3262  Image, CommandGroupHandler, Image.getElementSize()) {
3263  (void)propList;
3264 #ifndef __SYCL_DEVICE_ONLY__
3265  detail::associateWithHandler(CommandGroupHandler, this,
3267 #endif
3268  }
3269 #ifdef __SYCL_DEVICE_ONLY__
3270 private:
3271  using OCLImageTy =
3272  typename detail::opencl_image_type<Dimensions, AccessMode,
3273  access::target::image>::type;
3274 
3275  // Front End requires this method to be defined in the accessor class.
3276  // It does not call the base class's init method.
3277  void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
3278 
3279  // __init variant used by the device compiler for ESIMD kernels.
3280  void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
3281 
3282 public:
3283  // Default constructor for objects later initialized with __init member.
3284  accessor() = default;
3285 #endif
3286 };
3287 
3295 template <typename DataT, int Dimensions, access::mode AccessMode,
3298  access::target::host_image, IsPlaceholder>
3299  : public detail::image_accessor<DataT, Dimensions, AccessMode,
3300  access::target::host_image, IsPlaceholder>,
3301  public detail::OwnerLessBase<
3302  accessor<DataT, Dimensions, AccessMode, access::target::host_image,
3303  IsPlaceholder>> {
3304 public:
3305  template <typename AllocatorT>
3307  : detail::image_accessor<DataT, Dimensions, AccessMode,
3308  access::target::host_image, IsPlaceholder>(
3309  Image, Image.getElementSize()) {}
3310 
3311  template <typename AllocatorT>
3313  const property_list &propList)
3314  : detail::image_accessor<DataT, Dimensions, AccessMode,
3315  access::target::host_image, IsPlaceholder>(
3316  Image, Image.getElementSize()) {
3317  (void)propList;
3318  }
3319 };
3320 
3329 template <typename DataT, int Dimensions, access::mode AccessMode,
3332  DataT, Dimensions, AccessMode, access::target::image_array, IsPlaceholder>
3333  : public detail::image_accessor<DataT, Dimensions + 1, AccessMode,
3334  access::target::image, IsPlaceholder>,
3335  public detail::OwnerLessBase<
3336  accessor<DataT, Dimensions, AccessMode, access::target::image_array,
3337  IsPlaceholder>> {
3338 #ifdef __SYCL_DEVICE_ONLY__
3339 private:
3340  using OCLImageTy =
3341  typename detail::opencl_image_type<Dimensions + 1, AccessMode,
3342  access::target::image>::type;
3343 
3344  // Front End requires this method to be defined in the accessor class.
3345  // It does not call the base class's init method.
3346  void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
3347 
3348  // __init variant used by the device compiler for ESIMD kernels.
3349  void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
3350 
3351 public:
3352  // Default constructor for objects later initialized with __init member.
3353  accessor() = default;
3354 #endif
3355 public:
3356  template <typename AllocatorT>
3358  handler &CommandGroupHandler)
3359  : detail::image_accessor<DataT, Dimensions + 1, AccessMode,
3360  access::target::image, IsPlaceholder>(
3361  Image, CommandGroupHandler, Image.getElementSize()) {
3362 #ifndef __SYCL_DEVICE_ONLY__
3363  detail::associateWithHandler(CommandGroupHandler, this,
3365 #endif
3366  }
3367 
3368  template <typename AllocatorT>
3370  handler &CommandGroupHandler, const property_list &propList)
3371  : detail::image_accessor<DataT, Dimensions + 1, AccessMode,
3372  access::target::image, IsPlaceholder>(
3373  Image, CommandGroupHandler, Image.getElementSize()) {
3374  (void)propList;
3375 #ifndef __SYCL_DEVICE_ONLY__
3376  detail::associateWithHandler(CommandGroupHandler, this,
3378 #endif
3379  }
3380 
3381  detail::__image_array_slice__<DataT, Dimensions, AccessMode, IsPlaceholder>
3382  operator[](size_t Index) const {
3383  return detail::__image_array_slice__<DataT, Dimensions, AccessMode,
3384  IsPlaceholder>(*this, Index);
3385  }
3386 };
3387 
3388 template <typename DataT, int Dimensions = 1,
3391  : public accessor<DataT, Dimensions, AccessMode, target::host_buffer,
3392  access::placeholder::false_t> {
3393 protected:
3394  using AccessorT = accessor<DataT, Dimensions, AccessMode, target::host_buffer,
3396 
3397  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
3398  constexpr static bool IsAccessReadOnly = AccessMode == access::mode::read;
3399 
3400  template <typename T, int Dims>
3402  : std::bool_constant<std::is_same_v<T, DataT> && (Dims > 0) &&
3403  (Dims == Dimensions)> {};
3404 
3405  void
3406  __init(typename accessor<DataT, Dimensions, AccessMode, target::host_buffer,
3407  access::placeholder::false_t>::ConcreteASPtrType Ptr,
3408  range<AdjustedDim> AccessRange, range<AdjustedDim> MemRange,
3409  id<AdjustedDim> Offset) {
3410  AccessorT::__init(Ptr, AccessRange, MemRange, Offset);
3411  }
3412 
3413 #ifndef __SYCL_DEVICE_ONLY__
3415  : accessor<DataT, Dimensions, AccessMode, target::host_buffer,
3416  access::placeholder::false_t>{Impl} {}
3417 
3418  template <class Obj>
3419  friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject);
3420 
3421  template <class T>
3422  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
3423 #endif // __SYCL_DEVICE_ONLY__
3424 
3425 public:
3427 
3428  // The list of host_accessor constructors with their arguments
3429  // -------+---------+-------+----+----------+--------------
3430  // Dimensions = 0
3431  // -------+---------+-------+----+----------+--------------
3432  // buffer | | | | | property_list
3433  // buffer | handler | | | | property_list
3434  // -------+---------+-------+----+----------+--------------
3435  // Dimensions >= 1
3436  // -------+---------+-------+----+----------+--------------
3437  // buffer | | | | | property_list
3438  // buffer | | | | mode_tag | property_list
3439  // buffer | handler | | | | property_list
3440  // buffer | handler | | | mode_tag | property_list
3441  // buffer | | range | | | property_list
3442  // buffer | | range | | mode_tag | property_list
3443  // buffer | handler | range | | | property_list
3444  // buffer | handler | range | | mode_tag | property_list
3445  // buffer | | range | id | | property_list
3446  // buffer | | range | id | mode_tag | property_list
3447  // buffer | handler | range | id | | property_list
3448  // buffer | handler | range | id | mode_tag | property_list
3449  // -------+---------+-------+----+----------+--------------
3450 
3451  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3452  typename = typename std::enable_if_t<std::is_same_v<T, DataT> &&
3453  Dims == 0>>
3455  buffer<T, 1, AllocatorT> &BufferRef,
3456  const property_list &PropertyList = {},
3458  : AccessorT(BufferRef, PropertyList, CodeLoc) {}
3459 
3460  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3461  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3463  buffer<T, Dims, AllocatorT> &BufferRef,
3464  const property_list &PropertyList = {},
3466  : AccessorT(BufferRef, PropertyList, CodeLoc) {}
3467 
3468  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3469  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3472  const property_list &PropertyList = {},
3474  : host_accessor(BufferRef, PropertyList, CodeLoc) {}
3475 
3476  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3477  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3479  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3480  const property_list &PropertyList = {},
3482  : AccessorT(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
3483 
3484  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3485  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3487  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3488  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
3490  : host_accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
3491 
3492  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3493  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3495  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3496  const property_list &PropertyList = {},
3498  : AccessorT(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
3499 
3500  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3501  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3503  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3504  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
3506  : host_accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
3507 
3508  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3509  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3511  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3512  range<Dimensions> AccessRange, const property_list &PropertyList = {},
3514  : AccessorT(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
3515  CodeLoc) {}
3516 
3517  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3518  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3520  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3522  const property_list &PropertyList = {},
3524  : host_accessor(BufferRef, CommandGroupHandler, AccessRange, {},
3525  PropertyList, CodeLoc) {}
3526 
3527  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3528  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3530  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3531  id<Dimensions> AccessOffset, const property_list &PropertyList = {},
3533  : AccessorT(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
3534  }
3535 
3536  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3537  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3539  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3540  id<Dimensions> AccessOffset, mode_tag_t<AccessMode>,
3541  const property_list &PropertyList = {},
3543  : host_accessor(BufferRef, AccessRange, AccessOffset, PropertyList,
3544  CodeLoc) {}
3545 
3546  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3547  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3549  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3550  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
3551  const property_list &PropertyList = {},
3553  : AccessorT(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
3554  PropertyList, CodeLoc) {}
3555 
3556  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3557  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3559  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3560  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
3561  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
3563  : host_accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
3564  PropertyList, CodeLoc) {}
3565 
3566  template <int Dims = Dimensions,
3567  typename = std::enable_if_t<AccessMode != access_mode::atomic &&
3568  !IsAccessReadOnly && Dims == 0>>
3569  const host_accessor &
3570  operator=(const typename AccessorT::value_type &Other) const {
3571  *AccessorT::getQualifiedPtr() = Other;
3572  return *this;
3573  }
3574 
3575  template <int Dims = Dimensions,
3576  typename = std::enable_if_t<AccessMode != access_mode::atomic &&
3577  !IsAccessReadOnly && Dims == 0>>
3578  const host_accessor &operator=(typename AccessorT::value_type &&Other) const {
3579  *AccessorT::getQualifiedPtr() = std::move(Other);
3580  return *this;
3581  }
3582 
3583  // implicit conversion between const / non-const types for read only accessors
3584  template <typename DataT_,
3585  typename = std::enable_if_t<
3586  IsAccessReadOnly && !std::is_same_v<DataT_, DataT> &&
3587  std::is_same_v<std::remove_const_t<DataT_>,
3588  std::remove_const_t<DataT>>>>
3590 #ifndef __SYCL_DEVICE_ONLY__
3591  : host_accessor(other.impl) {
3592  AccessorT::MAccData = other.MAccData;
3593 #else
3594  {
3595  (void)other;
3596 #endif // __SYCL_DEVICE_ONLY__
3597  }
3598 
3599  // implicit conversion from read_write T accessor to read only T (const)
3600  // accessor
3601  template <typename DataT_, access::mode AccessMode_,
3602  typename = std::enable_if_t<
3603  (AccessMode_ == access_mode::read_write) && IsAccessReadOnly &&
3604  std::is_same_v<DataT_, std::remove_const_t<DataT>>>>
3606 #ifndef __SYCL_DEVICE_ONLY__
3607  : host_accessor(other.impl) {
3608  AccessorT::MAccData = other.MAccData;
3609 #else
3610  {
3611  (void)other;
3612 #endif // __SYCL_DEVICE_ONLY__
3613  }
3614 
3615  // host_accessor needs to explicitly define the owner_before member functions
3616  // as inheriting from OwnerLessBase causes base class conflicts.
3617  // TODO: Once host_accessor is detached from accessor, inherit from
3618  // OwnerLessBase instead.
3619 #ifndef __SYCL_DEVICE_ONLY__
3622  const noexcept {
3623  return this->impl.owner_before(
3625  }
3626 
3628  return this->impl.owner_before(Other.impl);
3629  }
3630 #else
3631  bool ext_oneapi_owner_before(
3633  const noexcept;
3634  bool ext_oneapi_owner_before(const host_accessor &Other) const noexcept;
3635 #endif
3636 };
3637 
3638 template <typename DataT, int Dimensions, typename AllocatorT>
3641 
3642 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
3644  -> host_accessor<DataT, Dimensions,
3645  detail::deduceAccessMode<Type1, Type1>()>;
3646 
3647 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3648  typename Type2>
3650  -> host_accessor<DataT, Dimensions,
3651  detail::deduceAccessMode<Type1, Type2>()>;
3652 
3653 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3654  typename Type2, typename Type3>
3656  -> host_accessor<DataT, Dimensions,
3657  detail::deduceAccessMode<Type2, Type3>()>;
3658 
3659 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3660  typename Type2, typename Type3, typename Type4>
3662  -> host_accessor<DataT, Dimensions,
3663  detail::deduceAccessMode<Type3, Type4>()>;
3664 
3665 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3666  typename Type2, typename Type3, typename Type4, typename Type5>
3668  Type5) -> host_accessor<DataT, Dimensions,
3669  detail::deduceAccessMode<Type4, Type5>()>;
3670 
3671 // SYCL 2020 image accessors
3672 
3673 template <typename DataT, int Dimensions, access_mode AccessMode,
3674  image_target AccessTarget = image_target::device>
3676 #ifndef __SYCL_DEVICE_ONLY__
3678 #endif // __SYCL_DEVICE_ONLY__
3679  public detail::OwnerLessBase<
3680  unsampled_image_accessor<DataT, Dimensions, AccessMode, AccessTarget>> {
3681  static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
3682  std::is_same_v<DataT, float4> ||
3683  std::is_same_v<DataT, half4>,
3684  "The data type of an image accessor must be only int4, "
3685  "uint4, float4 or half4 from SYCL namespace");
3686  static_assert(AccessMode == access_mode::read ||
3688  "Access mode must be either read or write.");
3689 
3690 #ifdef __SYCL_DEVICE_ONLY__
3691  char MPadding[sizeof(detail::UnsampledImageAccessorBaseHost)];
3692 #else
3694 #endif // __SYCL_DEVICE_ONLY__
3695 
3696 public:
3697  using value_type = typename std::conditional<AccessMode == access_mode::read,
3698  const DataT, DataT>::type;
3700  using const_reference = const DataT &;
3701 
3702  template <typename AllocatorT>
3705  handler &CommandGroupHandlerRef, const property_list &PropList = {},
3707 #ifdef __SYCL_DEVICE_ONLY__
3708  {
3709  (void)ImageRef;
3710  (void)CommandGroupHandlerRef;
3711  (void)PropList;
3712  (void)CodeLoc;
3713  }
3714 #else
3715  : host_base_class(detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
3716  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
3717  Dimensions, ImageRef.getElementSize(),
3718  {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
3719  ImageRef.getChannelType(), ImageRef.getChannelOrder(),
3720  PropList) {
3721  device Device = detail::getDeviceFromHandler(CommandGroupHandlerRef);
3722  // Avoid aspect::image warning.
3723  aspect ImageAspect = aspect::image;
3724  if (AccessTarget == image_target::device && !Device.has(ImageAspect))
3725  throw sycl::exception(
3726  sycl::make_error_code(sycl::errc::feature_not_supported),
3727  "Device associated with command group handler does not have "
3728  "aspect::image.");
3729 
3731  detail::getSyclObjImpl(ImageRef).get(), this->impl.get(), AccessTarget,
3732  AccessMode, (const void *)typeid(DataT).name(), sizeof(DataT), CodeLoc);
3733  detail::associateWithHandler(CommandGroupHandlerRef, this, AccessTarget);
3734  GDBMethodsAnchor();
3735  }
3736 #endif // __SYCL_DEVICE_ONLY__
3737 
3738  /* -- common interface members -- */
3739 
3741 
3743 
3745  operator=(const unsampled_image_accessor &Rhs) = default;
3746 
3748 
3750 
3751 #ifdef __SYCL_DEVICE_ONLY__
3752  bool operator==(const unsampled_image_accessor &Rhs) const;
3753 #else
3754  bool operator==(const unsampled_image_accessor &Rhs) const {
3755  return Rhs.impl == impl;
3756  }
3757 #endif // __SYCL_DEVICE_ONLY__
3758 
3759  bool operator!=(const unsampled_image_accessor &Rhs) const {
3760  return !(Rhs == *this);
3761  }
3762 
3763  /* -- property interface members -- */
3764  template <typename Property> bool has_property() const noexcept {
3765 #ifndef __SYCL_DEVICE_ONLY__
3766  return getPropList().template has_property<Property>();
3767 #else
3768  return false;
3769 #endif
3770  }
3771  template <typename Property> Property get_property() const {
3772 #ifndef __SYCL_DEVICE_ONLY__
3773  return getPropList().template get_property<Property>();
3774 #else
3775  return Property();
3776 #endif
3777  }
3778 
3779  size_t size() const noexcept {
3780 #ifdef __SYCL_DEVICE_ONLY__
3781  // Currently not reachable on device.
3782  return 0;
3783 #else
3784  return host_base_class::getSize().size();
3785 #endif // __SYCL_DEVICE_ONLY__
3786  }
3787 
3788  /* Available only when: AccessMode == access_mode::read
3789  if Dimensions == 1, CoordT = int
3790  if Dimensions == 2, CoordT = int2
3791  if Dimensions == 3, CoordT = int4 */
3792  template <typename CoordT,
3793  typename = std::enable_if_t<AccessMode == access_mode::read &&
3795  Dimensions, CoordT>::value>>
3796  DataT read(const CoordT &Coords) const noexcept {
3797 #ifdef __SYCL_DEVICE_ONLY__
3798  // Currently not reachable on device.
3799  std::ignore = Coords;
3800  return {0, 0, 0, 0};
3801 #else
3802  return host_base_class::read<DataT>(Coords);
3803 #endif // __SYCL_DEVICE_ONLY__
3804  }
3805 
3806  /* Available only when: AccessMode == access_mode::write
3807  if Dimensions == 1, CoordT = int
3808  if Dimensions == 2, CoordT = int2
3809  if Dimensions == 3, CoordT = int4 */
3810  template <typename CoordT,
3811  typename = std::enable_if_t<AccessMode == access_mode::write &&
3813  Dimensions, CoordT>::value>>
3814  void write(const CoordT &Coords, const DataT &Color) const {
3815 #ifdef __SYCL_DEVICE_ONLY__
3816  // Currently not reachable on device.
3817  std::ignore = Coords;
3818  std::ignore = Color;
3819 #else
3820  host_base_class::write<DataT>(Coords, Color);
3821 #endif // __SYCL_DEVICE_ONLY__
3822  }
3823 
3824 private:
3826 #ifndef __SYCL_DEVICE_ONLY__
3827  : host_base_class{Impl}
3828 #endif // __SYCL_DEVICE_ONLY__
3829  {
3830  std::ignore = Impl;
3831  }
3832 
3833  template <class Obj>
3834  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
3835 
3836  template <class T>
3837  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
3838 };
3839 
3840 template <typename DataT, int Dimensions = 1,
3842  (std::is_const_v<DataT> ? access_mode::read
3846  public detail::OwnerLessBase<
3848  static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
3849  std::is_same_v<DataT, float4> ||
3850  std::is_same_v<DataT, half4>,
3851  "The data type of an image accessor must be only int4, "
3852  "uint4, float4 or half4 from SYCL namespace");
3853 
3855 
3856 public:
3857  using value_type = typename std::conditional<AccessMode == access_mode::read,
3858  const DataT, DataT>::type;
3860  using const_reference = const DataT &;
3861 
3862  template <typename AllocatorT>
3865  const property_list &PropList = {},
3867  : base_class(detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
3868  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
3869  Dimensions, ImageRef.getElementSize(),
3870  {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
3871  ImageRef.getChannelType(), ImageRef.getChannelOrder(),
3872  PropList) {
3873  addHostUnsampledImageAccessorAndWait(base_class::impl.get());
3874 
3876  detail::getSyclObjImpl(ImageRef).get(), this->impl.get(), std::nullopt,
3877  AccessMode, (const void *)typeid(DataT).name(), sizeof(DataT), CodeLoc);
3878  }
3879 
3880  /* -- common interface members -- */
3881 
3883  default;
3884 
3886 
3889 
3892 
3894 
3896  return Rhs.impl == impl;
3897  }
3899  return !(Rhs == *this);
3900  }
3901 
3902  /* -- property interface members -- */
3903  template <typename Property> bool has_property() const noexcept {
3904 #ifndef __SYCL_DEVICE_ONLY__
3905  return getPropList().template has_property<Property>();
3906 #else
3907  return false;
3908 #endif
3909  }
3910  template <typename Property> Property get_property() const {
3911 #ifndef __SYCL_DEVICE_ONLY__
3912  return getPropList().template get_property<Property>();
3913 #else
3914  return Property();
3915 #endif
3916  }
3917 
3918  size_t size() const noexcept { return base_class::getSize().size(); }
3919 
3920  /* Available only when: (AccessMode == access_mode::read ||
3921  AccessMode == access_mode::read_write)
3922  if Dimensions == 1, CoordT = int
3923  if Dimensions == 2, CoordT = int2
3924  if Dimensions == 3, CoordT = int4 */
3925  template <
3926  typename CoordT,
3927  typename = std::enable_if_t<
3931  DataT read(const CoordT &Coords) const noexcept
3932 #ifdef __SYCL_DEVICE_ONLY__
3933  ;
3934 #else
3935  {
3936  // Host implementation is only available in host code. Device is not allowed
3937  // to use host_unsampled_image_accessor.
3938  return base_class::read<DataT>(Coords);
3939  }
3940 #endif
3941 
3942  /* Available only when: (AccessMode == access_mode::write ||
3943  AccessMode == access_mode::read_write)
3944  if Dimensions == 1, CoordT = int
3945  if Dimensions == 2, CoordT = int2
3946  if Dimensions == 3, CoordT = int4 */
3947  template <
3948  typename CoordT,
3949  typename = std::enable_if_t<
3953  void write(const CoordT &Coords, const DataT &Color) const
3954 #ifdef __SYCL_DEVICE_ONLY__
3955  ;
3956 #else
3957  {
3958  // Host implementation is only available in host code. Device is not allowed
3959  // to use host_unsampled_image_accessor.
3960  base_class::write<DataT>(Coords, Color);
3961  }
3962 #endif
3963 
3964 private:
3967  : base_class{Impl} {}
3968 
3969  template <class Obj>
3970  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
3971 
3972  template <class T>
3973  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
3974 };
3975 
3976 template <typename DataT, int Dimensions,
3977  image_target AccessTarget = image_target::device>
3979 #ifndef __SYCL_DEVICE_ONLY__
3981 #endif // __SYCL_DEVICE_ONLY__
3982  public detail::OwnerLessBase<
3983  sampled_image_accessor<DataT, Dimensions, AccessTarget>> {
3984  static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
3985  std::is_same_v<DataT, float4> ||
3986  std::is_same_v<DataT, half4>,
3987  "The data type of an image accessor must be only int4, "
3988  "uint4, float4 or half4 from SYCL namespace");
3989 
3990 #ifdef __SYCL_DEVICE_ONLY__
3991  char MPadding[sizeof(detail::SampledImageAccessorBaseHost)];
3992 #else
3994 #endif // __SYCL_DEVICE_ONLY__
3995 
3996 public:
3997  using value_type = const DataT;
3998  using reference = const DataT &;
3999  using const_reference = const DataT &;
4000 
4001  template <typename AllocatorT>
4004  handler &CommandGroupHandlerRef, const property_list &PropList = {},
4006 #ifdef __SYCL_DEVICE_ONLY__
4007  {
4008  (void)ImageRef;
4009  (void)CommandGroupHandlerRef;
4010  (void)PropList;
4011  (void)CodeLoc;
4012  }
4013 #else
4014  : host_base_class(detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
4015  detail::getSyclObjImpl(ImageRef).get(), Dimensions,
4016  ImageRef.getElementSize(),
4017  {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
4018  ImageRef.getChannelType(), ImageRef.getChannelOrder(),
4019  ImageRef.getSampler(), PropList) {
4020  device Device = detail::getDeviceFromHandler(CommandGroupHandlerRef);
4021  // Avoid aspect::image warning.
4022  aspect ImageAspect = aspect::image;
4023  if (AccessTarget == image_target::device && !Device.has(ImageAspect))
4024  throw sycl::exception(
4025  sycl::make_error_code(sycl::errc::feature_not_supported),
4026  "Device associated with command group handler does not have "
4027  "aspect::image.");
4028 
4030  detail::getSyclObjImpl(ImageRef).get(), this->impl.get(), AccessTarget,
4031  (const void *)typeid(DataT).name(), sizeof(DataT), CodeLoc);
4032  detail::associateWithHandler(CommandGroupHandlerRef, this, AccessTarget);
4033  GDBMethodsAnchor();
4034  }
4035 #endif // __SYCL_DEVICE_ONLY__
4036 
4037  /* -- common interface members -- */
4038 
4040 
4042 
4044  operator=(const sampled_image_accessor &Rhs) = default;
4045 
4047 
4049 
4050 #ifdef __SYCL_DEVICE_ONLY__
4051  bool operator==(const sampled_image_accessor &Rhs) const;
4052 #else
4053  bool operator==(const sampled_image_accessor &Rhs) const {
4054  return Rhs.impl == impl;
4055  }
4056 #endif // __SYCL_DEVICE_ONLY__
4057 
4058  bool operator!=(const sampled_image_accessor &Rhs) const {
4059  return !(Rhs == *this);
4060  }
4061 
4062  /* -- property interface members -- */
4063  template <typename Property> bool has_property() const noexcept {
4064 #ifndef __SYCL_DEVICE_ONLY__
4065  return getPropList().template has_property<Property>();
4066 #else
4067  return false;
4068 #endif
4069  }
4070  template <typename Property> Property get_property() const {
4071 #ifndef __SYCL_DEVICE_ONLY__
4072  return getPropList().template get_property<Property>();
4073 #else
4074  return Property();
4075 #endif
4076  }
4077 
4078  size_t size() const noexcept {
4079 #ifdef __SYCL_DEVICE_ONLY__
4080  // Currently not reachable on device.
4081  return 0;
4082 #else
4083  return host_base_class::getSize().size();
4084 #endif // __SYCL_DEVICE_ONLY__
4085  }
4086 
4087  /* if Dimensions == 1, CoordT = float
4088  if Dimensions == 2, CoordT = float2
4089  if Dimensions == 3, CoordT = float4 */
4090  template <typename CoordT,
4091  typename = std::enable_if_t<detail::IsValidSampledCoord2020DataT<
4092  Dimensions, CoordT>::value>>
4093  DataT read(const CoordT &Coords) const noexcept {
4094 #ifdef __SYCL_DEVICE_ONLY__
4095  // Currently not reachable on device.
4096  std::ignore = Coords;
4097  return {0, 0, 0, 0};
4098 #else
4099  return host_base_class::read<DataT>(Coords);
4100 #endif // __SYCL_DEVICE_ONLY__
4101  }
4102 
4103 private:
4105 #ifndef __SYCL_DEVICE_ONLY__
4106  : host_base_class{Impl}
4107 #endif // __SYCL_DEVICE_ONLY__
4108  {
4109  std::ignore = Impl;
4110  }
4111 
4112  template <class Obj>
4113  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
4114 
4115  template <class T>
4116  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
4117 };
4118 
4119 template <typename DataT, int Dimensions>
4122  public detail::OwnerLessBase<
4123  host_sampled_image_accessor<DataT, Dimensions>> {
4124  static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
4125  std::is_same_v<DataT, float4> ||
4126  std::is_same_v<DataT, half4>,
4127  "The data type of an image accessor must be only int4, "
4128  "uint4, float4 or half4 from SYCL namespace");
4129 
4131 
4132 public:
4133  using value_type = const DataT;
4134  using reference = const DataT &;
4135  using const_reference = const DataT &;
4136 
4137  template <typename AllocatorT>
4140  const property_list &PropList = {},
4142  : base_class(detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
4143  detail::getSyclObjImpl(ImageRef).get(), Dimensions,
4144  ImageRef.getElementSize(),
4145  {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
4146  ImageRef.getChannelType(), ImageRef.getChannelOrder(),
4147  ImageRef.getSampler(), PropList) {
4148  addHostSampledImageAccessorAndWait(base_class::impl.get());
4149 
4151  detail::getSyclObjImpl(ImageRef).get(), this->impl.get(), std::nullopt,
4152  (const void *)typeid(DataT).name(), sizeof(DataT), CodeLoc);
4153  }
4154 
4155  /* -- common interface members -- */
4156 
4158 
4160 
4163 
4166 
4168 
4169  bool operator==(const host_sampled_image_accessor &Rhs) const {
4170  return Rhs.impl == impl;
4171  }
4172  bool operator!=(const host_sampled_image_accessor &Rhs) const {
4173  return !(Rhs == *this);
4174  }
4175 
4176  /* -- property interface members -- */
4177  template <typename Property> bool has_property() const noexcept {
4178 #ifndef __SYCL_DEVICE_ONLY__
4179  return getPropList().template has_property<Property>();
4180 #else
4181  return false;
4182 #endif
4183  }
4184  template <typename Property> Property get_property() const {
4185 #ifndef __SYCL_DEVICE_ONLY__
4186  return getPropList().template get_property<Property>();
4187 #else
4188  return Property();
4189 #endif
4190  }
4191 
4192  size_t size() const noexcept { return base_class::getSize().size(); }
4193 
4194  /* if Dimensions == 1, CoordT = float
4195  if Dimensions == 2, CoordT = float2
4196  if Dimensions == 3, CoordT = float4 */
4197  template <typename CoordT,
4198  typename = std::enable_if_t<detail::IsValidSampledCoord2020DataT<
4199  Dimensions, CoordT>::value>>
4200  DataT read(const CoordT &Coords) const
4201 #ifdef __SYCL_DEVICE_ONLY__
4202  ;
4203 #else
4204  {
4205  // Host implementation is only available in host code. Device is not allowed
4206  // to use host_sampled_image_accessor.
4207  return base_class::read<DataT>(Coords);
4208  }
4209 #endif
4210 
4211 private:
4213  : base_class{Impl} {}
4214 
4215  template <class Obj>
4216  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
4217 
4218  template <class T>
4219  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
4220 };
4221 
4222 } // namespace _V1
4223 } // namespace sycl
4224 
4225 namespace std {
4226 template <typename DataT, int Dimensions, sycl::access::mode AccessMode,
4227  sycl::access::target AccessTarget,
4229 struct hash<sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
4230  IsPlaceholder>> {
4231  using AccType = sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
4232  IsPlaceholder>;
4233 
4234  size_t operator()(const AccType &A) const {
4235 #ifdef __SYCL_DEVICE_ONLY__
4236  // Hash is not supported on DEVICE. Just return 0 here.
4237  (void)A;
4238  return 0;
4239 #else
4240  // getSyclObjImpl() here returns a pointer to either AccessorImplHost
4241  // or LocalAccessorImplHost depending on the AccessTarget.
4242  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4243  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4244 #endif
4245  }
4246 };
4247 
4248 template <typename DataT, int Dimensions, sycl::access_mode AccessMode>
4249 struct hash<sycl::host_accessor<DataT, Dimensions, AccessMode>> {
4251 
4252  size_t operator()(const AccType &A) const {
4253 #ifdef __SYCL_DEVICE_ONLY__
4254  // Hash is not supported on DEVICE. Just return 0 here.
4255  (void)A;
4256  return 0;
4257 #else
4258  // getSyclObjImpl() here returns a pointer to AccessorImplHost.
4259  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4260  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4261 #endif
4262  }
4263 };
4264 
4265 template <typename DataT, int Dimensions>
4266 struct hash<sycl::local_accessor<DataT, Dimensions>> {
4268 
4269  size_t operator()(const AccType &A) const {
4270 #ifdef __SYCL_DEVICE_ONLY__
4271  // Hash is not supported on DEVICE. Just return 0 here.
4272  (void)A;
4273  return 0;
4274 #else
4275  // getSyclObjImpl() here returns a pointer to LocalAccessorImplHost.
4276  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4277  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4278 #endif
4279  }
4280 };
4281 
4282 template <typename DataT, int Dimensions, sycl::access_mode AccessMode,
4283  sycl::image_target AccessTarget>
4284 struct hash<sycl::unsampled_image_accessor<DataT, Dimensions, AccessMode,
4285  AccessTarget>> {
4287  AccessTarget>;
4288 
4289  size_t operator()(const AccType &A) const {
4290 #ifdef __SYCL_DEVICE_ONLY__
4291  // Hash is not supported on DEVICE. Just return 0 here.
4292  (void)A;
4293  return 0;
4294 #else
4295  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4296  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4297 #endif
4298  }
4299 };
4300 
4301 template <typename DataT, int Dimensions, sycl::access_mode AccessMode>
4302 struct hash<
4303  sycl::host_unsampled_image_accessor<DataT, Dimensions, AccessMode>> {
4304  using AccType =
4306 
4307  size_t operator()(const AccType &A) const {
4308  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4309  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4310  }
4311 };
4312 
4313 template <typename DataT, int Dimensions, sycl::image_target AccessTarget>
4314 struct hash<sycl::sampled_image_accessor<DataT, Dimensions, AccessTarget>> {
4316 
4317  size_t operator()(const AccType &A) const {
4318 #ifdef __SYCL_DEVICE_ONLY__
4319  // Hash is not supported on DEVICE. Just return 0 here.
4320  (void)A;
4321  return 0;
4322 #else
4323  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4324  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4325 #endif
4326  }
4327 };
4328 
4329 template <typename DataT, int Dimensions>
4330 struct hash<sycl::host_sampled_image_accessor<DataT, Dimensions>> {
4332 
4333  size_t operator()(const AccType &A) const {
4334  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4335  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4336  }
4337 };
4338 
4339 } // namespace std
The file contains implementation of accessor iterator class.
accessor(sycl::image< Dimensions, AllocatorT > &Image, const property_list &propList)
Definition: accessor.hpp:3312
Defines a shared array that can be used by kernels in queues.
Definition: buffer.hpp:169
unsigned int getElemSize() const
Definition: accessor.cpp:86
AccessorBaseHost(const AccessorImplPtr &Impl)
Definition: accessor.hpp:536
AccessorImplDevice(id< Dims > Offset, range< Dims > AccessRange, range< Dims > MemoryRange)
Definition: accessor.hpp:512
bool operator==(const AccessorImplDevice &Rhs) const
Definition: accessor.hpp:520
LocalAccessorBaseDevice(sycl::range< Dims > Size)
Definition: accessor.hpp:490
bool operator==(const LocalAccessorBaseDevice &Rhs) const
Definition: accessor.hpp:499
LocalAccessorBaseHost(const LocalAccessorImplPtr &Impl)
Definition: accessor.hpp:605
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:758
SampledImageAccessorBaseHost(const SampledImageAccessorImplPtr &Impl)
Definition: accessor.hpp:710
UnsampledImageAccessorBaseHost(const UnsampledImageAccessorImplPtr &Impl)
Definition: accessor.hpp:642
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:701
DataT read(const CoordT &Coords) const noexcept
Definition: accessor.hpp:688
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:1190
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:1154
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:1168
DataT read(const CoordT &Coords, const sampler &Smpl) const
Definition: accessor.hpp:1161
__image_array_slice__(accessor< DataT, Dimensions, AccessMode, access::target::image_array, IsPlaceholder, ext::oneapi::accessor_property_list<>> BaseAcc, size_t Idx)
Definition: accessor.hpp:1144
std::enable_if_t< CurDims==1 &&IsAccessAtomic, atomic< DataT, AS > > operator[](size_t Index) const
Definition: accessor.hpp:395
AccessorSubscript(AccType Accessor, size_t Index)
Definition: accessor.hpp:375
AccessorSubscript(AccType Accessor, id< Dims > IDs)
Definition: accessor.hpp:370
constexpr static access::address_space AS
Definition: accessor.hpp:314
constexpr static bool IsAccessReadWrite
Definition: accessor.hpp:348
static constexpr bool IsConst
Definition: accessor.hpp:346
constexpr static bool IsAccessAtomic
Definition: accessor.hpp:351
constexpr static bool IsHostBuf
Definition: accessor.hpp:316
constexpr static bool IsHostTask
Definition: accessor.hpp:317
constexpr static bool IsAccessReadOnly
Definition: accessor.hpp:345
constexpr static bool IsAccessAnyWrite
Definition: accessor.hpp:339
constexpr static bool IsPlaceH
Definition: accessor.hpp:328
detail::const_if_const_AS< AS, DataT > * PtrType
Definition: accessor.hpp:355
detail::const_if_const_AS< AS, DataT > & RefType
Definition: accessor.hpp:353
constexpr static bool IsConstantBuf
Definition: accessor.hpp:336
constexpr static bool IsGlobalBuf
Definition: accessor.hpp:333
size_t size() const noexcept
Definition: accessor.hpp:1039
image_accessor(image< Dims, AllocatorT > &ImageRef, int ImageElementSize)
Definition: accessor.hpp:936
bool operator!=(const image_accessor &Rhs) const
Definition: accessor.hpp:1004
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:1037
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:1103
range< Dims > get_range() const
Definition: accessor.hpp:1042
image_accessor(const AccessorImplPtr &Impl)
Definition: accessor.hpp:908
bool operator==(const image_accessor &Rhs) const
Definition: accessor.hpp:996
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:1058
DataT read(const CoordT &Coords, const sampler &Smpl) const
Definition: accessor.hpp:1077
image_accessor(image< Dims, AllocatorT > &ImageRef, handler &CommandGroupHandlerRef, int ImageElementSize)
Definition: accessor.hpp:963
range< Dimensions > get_range() const
Definition: image.hpp:323
size_t size() const noexcept
Definition: image.hpp:335
size_t getElementSize() const
Definition: image.cpp:187
image_sampler getSampler() const noexcept
Definition: image.cpp:193
size_t getRowPitch() const
Definition: image.cpp:189
image_channel_order getChannelOrder() const
Definition: image.cpp:197
size_t getSlicePitch() const
Definition: image.cpp:191
image_channel_type getChannelType() const
Definition: image.cpp:201
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:64
bool has(aspect Aspect) const __SYCL_WARN_IMAGE_ASPECT(Aspect)
Indicates if the SYCL device has the given feature.
Definition: device.cpp:219
Command group handler class.
Definition: handler.hpp:458
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3548
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3478
void __init(typename accessor< DataT, Dimensions, AccessMode, target::host_buffer, access::placeholder::false_t >::ConcreteASPtrType Ptr, range< AdjustedDim > AccessRange, range< AdjustedDim > MemRange, id< AdjustedDim > Offset)
Definition: accessor.hpp:3406
host_accessor(const detail::AccessorImplPtr &Impl)
Definition: accessor.hpp:3414
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, mode_tag_t< AccessMode >, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3486
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3510
host_accessor(const host_accessor< DataT_, Dimensions, AccessMode_ > &other)
Definition: accessor.hpp:3605
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3494
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3462
const host_accessor & operator=(const typename AccessorT::value_type &Other) const
Definition: accessor.hpp:3570
host_accessor(const host_accessor< DataT_, Dimensions, AccessMode > &other)
Definition: accessor.hpp:3589
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, mode_tag_t< AccessMode >, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3502
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, mode_tag_t< AccessMode >, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3470
decltype(Obj::impl) friend getSyclObjImpl(const Obj &SyclObject)
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, mode_tag_t< AccessMode >, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3538
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3529
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, mode_tag_t< AccessMode >, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3519
const host_accessor & operator=(typename AccessorT::value_type &&Other) const
Definition: accessor.hpp:3578
bool ext_oneapi_owner_before(const ext::oneapi::detail::weak_object_base< host_accessor > &Other) const noexcept
Definition: accessor.hpp:3620
bool ext_oneapi_owner_before(const host_accessor &Other) const noexcept
Definition: accessor.hpp:3627
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, mode_tag_t< AccessMode >, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3558
host_accessor(buffer< T, 1, AllocatorT > &BufferRef, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3454
bool has_property() const noexcept
Definition: accessor.hpp:4177
bool operator!=(const host_sampled_image_accessor &Rhs) const
Definition: accessor.hpp:4172
host_sampled_image_accessor(const host_sampled_image_accessor &Rhs)=default
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:4200
host_sampled_image_accessor & operator=(const host_sampled_image_accessor &Rhs)=default
host_sampled_image_accessor(host_sampled_image_accessor &&Rhs)=default
bool operator==(const host_sampled_image_accessor &Rhs) const
Definition: accessor.hpp:4169
host_sampled_image_accessor(sampled_image< Dimensions, AllocatorT > &ImageRef, const property_list &PropList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:4138
host_sampled_image_accessor & operator=(host_sampled_image_accessor &&Rhs)=default
host_unsampled_image_accessor(unsampled_image< Dimensions, AllocatorT > &ImageRef, const property_list &PropList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3863
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:3953
host_unsampled_image_accessor(const host_unsampled_image_accessor &Rhs)=default
host_unsampled_image_accessor(host_unsampled_image_accessor &&Rhs)=default
DataT read(const CoordT &Coords) const noexcept
Definition: accessor.hpp:3931
bool operator==(const host_unsampled_image_accessor &Rhs) const
Definition: accessor.hpp:3895
typename std::conditional< AccessMode==access_mode::read, const DataT, DataT >::type value_type
Definition: accessor.hpp:3858
host_unsampled_image_accessor & operator=(host_unsampled_image_accessor &&Rhs)=default
host_unsampled_image_accessor & operator=(const host_unsampled_image_accessor &Rhs)=default
bool operator!=(const host_unsampled_image_accessor &Rhs) const
Definition: accessor.hpp:3898
A unique identifier of an item in an index space.
Definition: id.hpp:36
Defines a shared image data.
Definition: image.hpp:443
range< Dimensions > get_range() const
Definition: image.hpp:629
detail::const_if_const_AS< AS, DataT > & RefType
Definition: accessor.hpp:2767
const range< 3 > & getSize() const
Definition: accessor.hpp:2821
bool operator!=(const local_accessor_base &Rhs) const
Definition: accessor.hpp:2999
AccessorCommonT::template AccessorSubscript< Dims - 1, local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder > > operator[](size_t Index) const
Definition: accessor.hpp:2992
local_accessor_base(range< Dimensions > AllocationSize, handler &, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2894
range< Dims > get_range() const
Definition: accessor.hpp:2937
local_accessor_base(handler &, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2861
typename detail::DecoratedType< DataT, AS >::type * ConcreteASPtrType
Definition: accessor.hpp:2765
typename AccessorCommonT::template AccessorSubscript< Dims, local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder > > AccessorSubscript
Definition: accessor.hpp:2763
size_t getLinearIndex(id< Dims > Id) const
Definition: accessor.hpp:2840
local_accessor_base(range< Dimensions > AllocationSize, handler &, const property_list &propList, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2912
local_accessor_base(handler &, const property_list &propList, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2876
bool operator==(const local_accessor_base &Rhs) const
Definition: accessor.hpp:2996
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:2932
detail::const_if_const_AS< AS, DataT > * PtrType
Definition: accessor.hpp:2768
size_t size() const noexcept
Definition: accessor.hpp:2934
std::enable_if_t< Dims==1 &&AccessMode==access::mode::atomic, atomic< DataT, AS > > operator[](size_t Index) const
Definition: accessor.hpp:2983
RefType operator[](id< Dimensions > Index) const
Definition: accessor.hpp:2951
PtrType getQualifiedPtr() const
Definition: accessor.hpp:2815
local_accessor_base(const detail::LocalAccessorImplPtr &Impl)
Definition: accessor.hpp:2808
Objects of the property_list class are containers for the SYCL properties.
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: range.hpp:26
bool has_property() const noexcept
Definition: accessor.hpp:4063
sampled_image_accessor(const sampled_image_accessor &Rhs)=default
bool operator==(const sampled_image_accessor &Rhs) const
Definition: accessor.hpp:4053
sampled_image_accessor(sampled_image_accessor &&Rhs)=default
sampled_image_accessor & operator=(sampled_image_accessor &&Rhs)=default
sampled_image_accessor(sampled_image< Dimensions, AllocatorT > &ImageRef, handler &CommandGroupHandlerRef, const property_list &PropList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:4002
size_t size() const noexcept
Definition: accessor.hpp:4078
sampled_image_accessor & operator=(const sampled_image_accessor &Rhs)=default
DataT read(const CoordT &Coords) const noexcept
Definition: accessor.hpp:4093
bool operator!=(const sampled_image_accessor &Rhs) const
Definition: accessor.hpp:4058
DataT read(const CoordT &Coords) const noexcept
Definition: accessor.hpp:3796
bool operator!=(const unsampled_image_accessor &Rhs) const
Definition: accessor.hpp:3759
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:3814
bool has_property() const noexcept
Definition: accessor.hpp:3764
unsampled_image_accessor & operator=(unsampled_image_accessor &&Rhs)=default
size_t size() const noexcept
Definition: accessor.hpp:3779
bool operator==(const unsampled_image_accessor &Rhs) const
Definition: accessor.hpp:3754
unsampled_image_accessor & operator=(const unsampled_image_accessor &Rhs)=default
typename std::conditional< AccessMode==access_mode::read, const DataT, DataT >::type value_type
Definition: accessor.hpp:3698
unsampled_image_accessor(const unsampled_image_accessor &Rhs)=default
unsampled_image_accessor(unsampled_image_accessor &&Rhs)=default
unsampled_image_accessor(unsampled_image< Dimensions, AllocatorT > &ImageRef, handler &CommandGroupHandlerRef, const property_list &PropList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3703
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
Definition: types.hpp:284
#define __SYCL_SPECIAL_CLASS
Definition: defines.hpp:29
#define __SYCL_EBO
Definition: common.hpp:250
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor< DataT
Image accessors.
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor accessor(buffer< DataT, Dimensions, AllocatorT >) -> accessor< DataT, Dimensions, access::mode::read_write, target::device, access::placeholder::true_t >
Buffer accessor.
void unsampledImageConstructorNotification(void *ImageObj, void *AccessorObj, const std::optional< image_target > &Target, access::mode Mode, const void *Type, uint32_t ElemSize, const code_location &CodeLoc)
void addHostUnsampledImageAccessorAndWait(UnsampledImageAccessorImplHost *Req)
void addHostAccessorAndWait(AccessorImplHost *Req)
device getDeviceFromHandler(handler &CommandGroupHandlerRef)
Definition: accessor.cpp:16
size_t getLinearIndex(const T< Dims > &Index, const U< Dims > &Range)
Definition: common.hpp:367
sycl::range< 1 > GetZeroDimAccessRange(BufferT Buffer)
Definition: accessor.hpp:303
std::shared_ptr< UnsampledImageAccessorImplHost > UnsampledImageAccessorImplPtr
Definition: accessor.hpp:631
void imageWriteHostImpl(const CoordT &Coords, const WriteDataT &Color, id< 3 > ImgPitch, uint8_t ElementSize, image_channel_type ImgChannelType, image_channel_order ImgChannelOrder, void *BasePtr)
constexpr access::mode accessModeFromConstness()
Definition: accessor.hpp:402
void addHostSampledImageAccessorAndWait(SampledImageAccessorImplHost *Req)
boost::mp11::mp_set_contains< TypeList, std::remove_cv_t< T > > is_contained
Definition: type_list.hpp:32
std::shared_ptr< SampledImageAccessorImplHost > SampledImageAccessorImplPtr
Definition: accessor.hpp:633
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: impl_utils.hpp:30
typename std::is_same< ext::oneapi::accessor_property_list<>, T > IsRunTimePropertyListT
Definition: accessor.hpp:285
void constructorNotification(void *BufferObj, void *AccessorObj, access::target Target, access::mode Mode, const code_location &CodeLoc)
boost::mp11::mp_list< T... > type_list
Definition: type_list.hpp:22
T createSyclObjFromImpl(decltype(T::impl) ImplObj)
Definition: impl_utils.hpp:48
std::shared_ptr< LocalAccessorImplHost > LocalAccessorImplPtr
Definition: accessor.hpp:601
struct __SYCL2020_DEPRECATED("This type isn't device copyable in SYCL 2020") IsDeprecatedDeviceCopyable< T
constexpr access::target deduceAccessTarget(access::target defaultTarget)
Definition: accessor.hpp:457
typename std::is_base_of< PropertyListBase, T > IsPropertyListT
Definition: accessor.hpp:281
void sampledImageConstructorNotification(void *ImageObj, void *AccessorObj, const std::optional< image_target > &Target, const void *Type, uint32_t ElemSize, const code_location &CodeLoc)
constexpr access::mode deduceAccessMode()
Definition: accessor.hpp:410
void associateWithHandler(handler &, AccessorBaseHost *, access::target)
std::shared_ptr< AccessorImplHost > AccessorImplPtr
Definition: accessor.hpp:532
constexpr buffer_location_key::value_t< N > buffer_location
bool operator==(const cache_config &lhs, const cache_config &rhs)
bool operator!=(const cache_config &lhs, const cache_config &rhs)
decltype(weak_object_base< SYCLObjT >::MObjWeakPtr) getSyclWeakObjImpl(const weak_object_base< SYCLObjT > &WeakObj)
static constexpr bool has_property()
sycl::ext::oneapi::experimental::annotated_ref< T, property_list_t > reference
static constexpr auto get_property()
T & operator[](std::ptrdiff_t idx) const noexcept
access::mode access_mode
Definition: access.hpp:72
host_accessor(buffer< DataT, Dimensions, AllocatorT >) -> host_accessor< DataT, Dimensions, access::mode::read_write >
image_target
Definition: access.hpp:74
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS Dimensions
Definition: accessor.hpp:3233
multi_ptr< ElementType, access::address_space::global_space, IsDecorated > global_ptr
Definition: pointers.hpp:29
return(x >> one)+(y >> one)+((y &x) &one)
constexpr mode_tag_t< access_mode::read_write > read_write
Definition: access.hpp:85
pointer get() const
Definition: multi_ptr.hpp:544
std::ptrdiff_t difference_type
Definition: multi_ptr.hpp:460
image_channel_order
Definition: image.hpp:56
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS IsPlaceholder
Definition: accessor.hpp:3234
PropertyListT int access::address_space multi_ptr & operator=(multi_ptr &&)=default
PropertyListT Accessor
Definition: multi_ptr.hpp:510
accessor(buffer< DataT, Dimensions, AllocatorT >, handler &, Type1, Type2, Type3, Type4, const ext::oneapi::accessor_property_list< PropsT... > &) -> accessor< DataT, Dimensions, detail::deduceAccessMode< Type3, Type4 >(), detail::deduceAccessTarget< Type3, Type4 >(target::device), access::placeholder::false_t, ext::oneapi::accessor_property_list< PropsT... >>
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS AccessMode
Definition: accessor.hpp:3233
image_channel_type
Definition: image.hpp:74
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:87
multi_ptr< ElementType, access::address_space::local_space, IsDecorated > local_ptr
Definition: pointers.hpp:34
const void value_type
Definition: multi_ptr.hpp:457
Definition: access.hpp:18
static size_t get_offset(sycl::id< 3 > id, size_t slice, size_t pitch)
Definition: memory.hpp:297
_Abi const simd< _Tp, _Abi > & noexcept
Definition: simd.hpp:1324
sycl::range< 3 > MMemoryRange
Definition: accessor.hpp:260
sycl::range< 3 > MAccessRange
Definition: accessor.hpp:259
AccHostDataT(const sycl::id< 3 > &Offset, const sycl::range< 3 > &Range, const sycl::range< 3 > &MemoryRange, void *Data=nullptr)
Definition: accessor.hpp:253
constexpr static bool value
Definition: accessor.hpp:288
static constexpr code_location current(const char *fileName=__CODELOC_FILE_NAME, const char *funcName=__CODELOC_FUNCTION, unsigned long lineNo=__CODELOC_LINE, unsigned long columnNo=__CODELOC_COLUMN) noexcept
Definition: common.hpp:68