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,
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,
267  access::mode Mode,
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,
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,
364  static constexpr int Dims = Dimensions;
365 
366  mutable id<Dims> MIDs;
367  AccType MAccessor;
368 
369  public:
370  AccessorSubscript(AccType Accessor, id<Dims> IDs)
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:
646  UnsampledImageAccessorBaseHost(sycl::range<3> Size, access_mode AccessMode,
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() : MImageObj() {}
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 feature_not_supported(
986  "SYCL 1.2.1 images are not supported by this device.",
987  PI_ERROR_INVALID_OPERATION);
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) &&
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 <typename CoordT, int Dims = Dimensions,
1097  typename = std::enable_if_t<
1098  (Dims > 0) && (detail::is_genint<CoordT>::value) &&
1100  ((IsImageAcc && IsImageAccessWriteOnly) ||
1101  (IsHostImageAcc && IsImageAccessAnyWrite))>>
1102  void write(const CoordT &Coords, const DataT &Color) const {
1103 #ifdef __SYCL_DEVICE_ONLY__
1104  __invoke__ImageWrite<OCLImageTy, CoordT, DataT>(MImageObj, Coords, Color);
1105 #else
1106  imageWriteHostImpl(Coords, Color, getOffset() /*ImagePitch*/,
1107  AccessorBaseHost::getElemSize(), MImgChannelType,
1108  MImgChannelOrder,
1109  AccessorBaseHost::getPtr() /*Ptr to Image*/);
1110 #endif
1111  }
1112 };
1113 
1114 template <typename DataT, int Dimensions, access::mode AccessMode,
1116 class __image_array_slice__ {
1117 
1118  static_assert(Dimensions < 3,
1119  "Image slice cannot have more then 2 dimensions");
1120 
1121  constexpr static int AdjustedDims = (Dimensions == 2) ? 4 : Dimensions + 1;
1122 
1123  template <typename CoordT,
1124  typename CoordElemType =
1127  getAdjustedCoords(const CoordT &Coords) const {
1128  CoordElemType LastCoord = 0;
1129 
1130  if (std::is_same<float, CoordElemType>::value) {
1131  sycl::vec<int, Dimensions + 1> Size = MBaseAcc.getRangeInternal();
1132  LastCoord =
1133  MIdx / static_cast<float>(Size.template swizzle<Dimensions>());
1134  } else {
1135  LastCoord = MIdx;
1136  }
1137 
1138  sycl::vec<CoordElemType, Dimensions> LeftoverCoords{LastCoord};
1139  sycl::vec<CoordElemType, AdjustedDims> AdjustedCoords{Coords,
1140  LeftoverCoords};
1141  return AdjustedCoords;
1142  }
1143 
1144 public:
1148  BaseAcc,
1149  size_t Idx)
1150  : MBaseAcc(BaseAcc), MIdx(Idx) {}
1151 
1152  template <typename CoordT, int Dims = Dimensions,
1153  typename = std::enable_if_t<
1154  (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value)>>
1155  DataT read(const CoordT &Coords) const {
1156  return MBaseAcc.read(getAdjustedCoords(Coords));
1157  }
1158 
1159  template <typename CoordT, int Dims = Dimensions,
1160  typename = std::enable_if_t<(Dims > 0) &&
1162  DataT read(const CoordT &Coords, const sampler &Smpl) const {
1163  return MBaseAcc.read(getAdjustedCoords(Coords), Smpl);
1164  }
1165 
1166  template <typename CoordT, int Dims = Dimensions,
1167  typename = std::enable_if_t<(Dims > 0) &&
1169  void write(const CoordT &Coords, const DataT &Color) const {
1170  return MBaseAcc.write(getAdjustedCoords(Coords), Color);
1171  }
1172 
1173 #ifdef __SYCL_DEVICE_ONLY__
1174  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1175  size_t get_count() const { return size(); }
1176  size_t size() const noexcept { return get_range<Dimensions>().size(); }
1177 
1178  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 1>>
1179  range<1> get_range() const {
1180  int2 Count = MBaseAcc.getRangeInternal();
1181  return range<1>(Count.x());
1182  }
1183  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 2>>
1184  range<2> get_range() const {
1185  int3 Count = MBaseAcc.getRangeInternal();
1186  return range<2>(Count.x(), Count.y());
1187  }
1188 
1189 #else
1190 
1191  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1192  size_t get_count() const { return size(); }
1193  size_t size() const noexcept {
1194  return MBaseAcc.MImageCount / MBaseAcc.getAccessRange()[Dimensions];
1195  }
1196 
1197  template <int Dims = Dimensions,
1198  typename = std::enable_if_t<(Dims == 1 || Dims == 2)>>
1200  return detail::convertToArrayOfN<Dims, 1>(MBaseAcc.getAccessRange());
1201  }
1202 
1203 #endif
1204 
1205 private:
1206  size_t MIdx;
1209  MBaseAcc;
1210 };
1211 
1212 } // namespace detail
1213 
1219 template <typename DataT, int Dimensions, access::mode AccessMode,
1221  typename PropertyListT>
1223 #ifndef __SYCL_DEVICE_ONLY__
1224  public detail::AccessorBaseHost,
1225 #endif
1226  public detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
1227  IsPlaceholder, PropertyListT>,
1228  public detail::OwnerLessBase<
1229  accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
1230  PropertyListT>> {
1231 protected:
1232  static_assert((AccessTarget == access::target::global_buffer ||
1233  AccessTarget == access::target::constant_buffer ||
1234  AccessTarget == access::target::host_buffer ||
1235  AccessTarget == access::target::host_task),
1236  "Expected buffer type");
1237 
1238  static_assert((AccessTarget == access::target::global_buffer ||
1239  AccessTarget == access::target::host_buffer ||
1240  AccessTarget == access::target::host_task) ||
1241  (AccessTarget == access::target::constant_buffer &&
1243  "Access mode can be only read for constant buffers");
1244 
1245  static_assert(detail::IsPropertyListT<PropertyListT>::value,
1246  "PropertyListT must be accessor_property_list");
1247 
1248  using AccessorCommonT =
1249  detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
1251 
1252  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
1253 
1254  using AccessorCommonT::AS;
1255  // Cannot do "using AccessorCommonT::Flag" as it doesn't work with g++ as host
1256  // compiler, for some reason.
1257  static constexpr bool IsAccessAnyWrite = AccessorCommonT::IsAccessAnyWrite;
1258  static constexpr bool IsAccessReadOnly = AccessorCommonT::IsAccessReadOnly;
1259  static constexpr bool IsConstantBuf = AccessorCommonT::IsConstantBuf;
1260  static constexpr bool IsGlobalBuf = AccessorCommonT::IsGlobalBuf;
1261  static constexpr bool IsHostBuf = AccessorCommonT::IsHostBuf;
1262  static constexpr bool IsPlaceH = AccessorCommonT::IsPlaceH;
1263  static constexpr bool IsConst = AccessorCommonT::IsConst;
1264  static constexpr bool IsHostTask = AccessorCommonT::IsHostTask;
1265  template <int Dims>
1266  using AccessorSubscript =
1267  typename AccessorCommonT::template AccessorSubscript<Dims>;
1268 
1269  static_assert(
1270  !IsConst || IsAccessReadOnly,
1271  "A const qualified DataT is only allowed for a read-only accessor");
1272 
1273  using ConcreteASPtrType = typename detail::DecoratedType<
1274  typename std::conditional_t<IsAccessReadOnly && !IsConstantBuf,
1275  const DataT, DataT>,
1276  AS>::type *;
1277 
1278  using RefType = detail::const_if_const_AS<AS, DataT> &;
1279  using ConstRefType = const DataT &;
1280  using PtrType = detail::const_if_const_AS<AS, DataT> *;
1281 
1282  template <int Dims = Dimensions> size_t getLinearIndex(id<Dims> Id) const {
1283 
1284  size_t Result = 0;
1285  detail::loop<Dims>([&, this](size_t I) {
1286  Result = Result * getMemoryRange()[I] + Id[I];
1287  // We've already adjusted for the accessor's offset in the __init, so
1288  // don't include it here in case of device.
1289 #ifndef __SYCL_DEVICE_ONLY__
1290  if constexpr (!(PropertyListT::template has_property<
1292  Result += getOffset()[I];
1293  }
1294 #endif // __SYCL_DEVICE_ONLY__
1295  });
1296 
1297  return Result;
1298  }
1299 
1300  template <typename T, int Dims>
1301  struct IsSameAsBuffer
1302  : std::bool_constant<std::is_same_v<T, DataT> && (Dims > 0) &&
1303  (Dims == Dimensions)> {};
1304 
1305  static access::mode getAdjustedMode(const PropertyListT &PropertyList) {
1306  access::mode AdjustedMode = AccessMode;
1307 
1308  if (PropertyList.template has_property<property::no_init>() ||
1309  PropertyList.template has_property<property::noinit>()) {
1310  if (AdjustedMode == access::mode::write) {
1311  AdjustedMode = access::mode::discard_write;
1312  } else if (AdjustedMode == access::mode::read_write) {
1313  AdjustedMode = access::mode::discard_read_write;
1314  }
1315  }
1316 
1317  return AdjustedMode;
1318  }
1319 
1320  template <typename TagT>
1321  struct IsValidTag
1322  : std::disjunction<
1323  std::is_same<TagT, mode_tag_t<AccessMode>>,
1324  std::is_same<TagT, mode_target_tag_t<AccessMode, AccessTarget>>> {};
1325 
1326  template <typename DataT_, int Dimensions_, access::mode AccessMode_,
1327  access::target AccessTarget_, access::placeholder IsPlaceholder_,
1328  typename PropertyListT_>
1329  friend class accessor;
1330 
1331 #ifdef __SYCL_DEVICE_ONLY__
1332 
1333  id<AdjustedDim> &getOffset() { return impl.Offset; }
1334  range<AdjustedDim> &getAccessRange() { return impl.AccessRange; }
1335  range<AdjustedDim> &getMemoryRange() { return impl.MemRange; }
1336 
1337  const id<AdjustedDim> &getOffset() const { return impl.Offset; }
1338  const range<AdjustedDim> &getAccessRange() const { return impl.AccessRange; }
1339  const range<AdjustedDim> &getMemoryRange() const { return impl.MemRange; }
1340 
1341  detail::AccessorImplDevice<AdjustedDim> impl;
1342 
1343  union {
1344  ConcreteASPtrType MData;
1345  };
1346 
1347  void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
1348  range<AdjustedDim> MemRange, id<AdjustedDim> Offset) {
1349  MData = Ptr;
1350  detail::loop<AdjustedDim>([&, this](size_t I) {
1351  if constexpr (!(PropertyListT::template has_property<
1353  getOffset()[I] = Offset[I];
1354  }
1355  getAccessRange()[I] = AccessRange[I];
1356  getMemoryRange()[I] = MemRange[I];
1357  });
1358 
1359  // Adjust for offsets as that part is invariant for all invocations of
1360  // operator[]. Will have to re-adjust in get_pointer.
1361  MData += getTotalOffset();
1362  }
1363 
1364  // __init variant used by the device compiler for ESIMD kernels.
1365  // TODO: In ESIMD accessors usage is limited for now - access range, mem
1366  // range and offset are not supported.
1367  void __init_esimd(ConcreteASPtrType Ptr) {
1368  MData = Ptr;
1369 #ifdef __ESIMD_FORCE_STATELESS_MEM
1370  detail::loop<AdjustedDim>([&, this](size_t I) {
1371  getOffset()[I] = 0;
1372  getAccessRange()[I] = 0;
1373  getMemoryRange()[I] = 0;
1374  });
1375 #endif
1376  }
1377 
1378  ConcreteASPtrType getQualifiedPtr() const noexcept { return MData; }
1379 
1380 #ifndef __SYCL_DEVICE_ONLY__
1381  using AccessorBaseHost::impl;
1382 #endif
1383 
1384 public:
1385  // Default constructor for objects later initialized with __init member.
1386  accessor()
1387  : impl({}, detail::InitializedVal<AdjustedDim, range>::template get<0>(),
1388  detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
1389 
1390 #else
1391  accessor(const detail::AccessorImplPtr &Impl)
1392  : detail::AccessorBaseHost{Impl} {}
1393 
1394  void *getPtr() { return AccessorBaseHost::getPtr(); }
1395 
1396  const id<3> getOffset() const {
1397  if constexpr (IsHostBuf)
1398  return MAccData ? MAccData->MOffset : id<3>();
1399  else
1400  return AccessorBaseHost::getOffset();
1401  }
1402  const range<3> &getAccessRange() const {
1404  }
1405  const range<3> getMemoryRange() const {
1406  if constexpr (IsHostBuf)
1407  return MAccData ? MAccData->MMemoryRange : range(0, 0, 0);
1408  else
1410  }
1411 
1412  void *getPtr() const { return AccessorBaseHost::getPtr(); }
1413 
1414  void initHostAcc() { MAccData = &getAccData(); }
1415 
1416  // The function references helper methods required by GDB pretty-printers
1417  void GDBMethodsAnchor() {
1418 #ifndef NDEBUG
1419  const auto *this_const = this;
1420  (void)getMemoryRange();
1421  (void)this_const->getMemoryRange();
1422  (void)getOffset();
1423  (void)this_const->getOffset();
1424  (void)getPtr();
1425  (void)this_const->getPtr();
1426  (void)getAccessRange();
1427  (void)this_const->getAccessRange();
1428 #endif
1429  }
1430 
1431  detail::AccHostDataT *MAccData = nullptr;
1432 
1433  char padding[sizeof(detail::AccessorImplDevice<AdjustedDim>) +
1434  sizeof(PtrType) - sizeof(detail::AccessorBaseHost) -
1435  sizeof(MAccData)];
1436 
1437  PtrType getQualifiedPtr() const noexcept {
1438  if constexpr (IsHostBuf)
1439  return MAccData ? reinterpret_cast<PtrType>(MAccData->MData) : nullptr;
1440  else
1441  return reinterpret_cast<PtrType>(AccessorBaseHost::getPtr());
1442  }
1443 
1444 public:
1445  accessor()
1446  : AccessorBaseHost(
1447  /*Offset=*/{0, 0, 0}, /*AccessRange=*/{0, 0, 0},
1448  /*MemoryRange=*/{0, 0, 0},
1449  /*AccessMode=*/getAdjustedMode({}),
1450  /*SYCLMemObject=*/nullptr, /*Dims=*/0, /*ElemSize=*/0,
1451  /*IsPlaceH=*/false,
1452  /*OffsetInBytes=*/0, /*IsSubBuffer=*/false, /*PropertyList=*/{}){};
1453 
1454  template <typename, int, access_mode> friend class host_accessor;
1455 
1456 #endif // __SYCL_DEVICE_ONLY__
1457 
1458 private:
1459  friend class sycl::stream;
1460  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
1461 
1462  template <class Obj>
1463  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
1464 
1465  template <class T>
1466  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
1467 
1468 public:
1469  // 4.7.6.9.1. Interface for buffer command accessors
1470  // value_type is defined as const DataT for read_only accessors, DataT
1471  // otherwise
1472  using value_type =
1473  std::conditional_t<AccessMode == access_mode::read, const DataT, DataT>;
1474  using reference = value_type &;
1475  using const_reference = const DataT &;
1476 
1477  template <access::decorated IsDecorated>
1478  using accessor_ptr =
1479  std::conditional_t<AccessTarget == access::target::device,
1480  global_ptr<value_type, IsDecorated>, value_type *>;
1481 
1482  using iterator = typename detail::accessor_iterator<value_type, AdjustedDim>;
1483  using const_iterator =
1484  typename detail::accessor_iterator<const value_type, AdjustedDim>;
1485  using reverse_iterator = std::reverse_iterator<iterator>;
1486  using const_reverse_iterator = std::reverse_iterator<const_iterator>;
1487  using difference_type =
1489  using size_type = std::size_t;
1490 
1493  void throwIfUsedByGraph() const {
1494 #ifndef __SYCL_DEVICE_ONLY__
1495  if (IsHostBuf && AccessorBaseHost::isMemoryObjectUsedByGraph()) {
1497  "Host accessors cannot be created for buffers "
1498  "which are currently in use by a command graph.");
1499  }
1500 #endif
1501  }
1502 
1503  // The list of accessor constructors with their arguments
1504  // -------+---------+-------+----+-----+--------------
1505  // Dimensions = 0
1506  // -------+---------+-------+----+-----+--------------
1507  // buffer | | | | | property_list
1508  // buffer | handler | | | | property_list
1509  // -------+---------+-------+----+-----+--------------
1510  // Dimensions >= 1
1511  // -------+---------+-------+----+-----+--------------
1512  // buffer | | | | | property_list
1513  // buffer | | | | tag | property_list
1514  // buffer | handler | | | | property_list
1515  // buffer | handler | | | tag | property_list
1516  // buffer | | range | | | property_list
1517  // buffer | | range | | tag | property_list
1518  // buffer | handler | range | | | property_list
1519  // buffer | handler | range | | tag | property_list
1520  // buffer | | range | id | | property_list
1521  // buffer | | range | id | tag | property_list
1522  // buffer | handler | range | id | | property_list
1523  // buffer | handler | range | id | tag | property_list
1524  // -------+---------+-------+----+-----+--------------
1525 
1526 public:
1527  // implicit conversion between const / non-const types for read only accessors
1528  template <typename DataT_,
1529  typename = std::enable_if_t<
1530  IsAccessReadOnly && !std::is_same_v<DataT_, DataT> &&
1531  std::is_same_v<std::remove_const_t<DataT_>,
1532  std::remove_const_t<DataT>>>>
1533  accessor(const accessor<DataT_, Dimensions, AccessMode, AccessTarget,
1534  IsPlaceholder, PropertyListT> &other)
1535 #ifdef __SYCL_DEVICE_ONLY__
1536  : impl(other.impl), MData(other.MData) {
1537 #else
1538  : accessor(other.impl) {
1539 #endif // __SYCL_DEVICE_ONLY__
1540  }
1541 
1542  // implicit conversion from read_write T accessor to read only T (const)
1543  // accessor
1544  template <typename DataT_, access::mode AccessMode_,
1545  typename = std::enable_if_t<
1546  (AccessMode_ == access_mode::read_write) && IsAccessReadOnly &&
1547  std::is_same_v<std::remove_const_t<DataT_>,
1548  std::remove_const_t<DataT>>>>
1549  accessor(const accessor<DataT_, Dimensions, AccessMode_, AccessTarget,
1550  IsPlaceholder, PropertyListT> &other)
1551 #ifdef __SYCL_DEVICE_ONLY__
1552  : impl(other.impl), MData(other.MData) {
1553 #else
1554  : accessor(other.impl) {
1555 #endif // __SYCL_DEVICE_ONLY__
1556  }
1557 
1558  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1559  typename std::enable_if_t<
1560  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1561  std::is_same_v<T, DataT> && Dims == 0 &&
1562  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))> * =
1563  nullptr>
1564  accessor(
1565  buffer<T, 1, AllocatorT> &BufferRef,
1566  const property_list &PropertyList = {},
1567  const detail::code_location CodeLoc = detail::code_location::current())
1568 #ifdef __SYCL_DEVICE_ONLY__
1569  : impl(id<AdjustedDim>(), detail::GetZeroDimAccessRange(BufferRef),
1570  BufferRef.get_range()) {
1571  (void)PropertyList;
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 #else
1613  : AccessorBaseHost(
1614  /*Offset=*/{0, 0, 0},
1615  detail::convertToArrayOfN<3, 1>(
1616  detail::GetZeroDimAccessRange(BufferRef)),
1617  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1618  getAdjustedMode(PropertyList),
1619  detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
1620  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1621  PropertyList) {
1622  throwIfUsedByGraph();
1623  preScreenAccessor(PropertyList);
1626  initHostAcc();
1629  AccessTarget, AccessMode, CodeLoc);
1630  GDBMethodsAnchor();
1631 #endif
1632  }
1633 
1634  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1635  typename = typename std::enable_if_t<
1636  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1637  std::is_same_v<T, DataT> && (Dims == 0) &&
1638  (IsGlobalBuf || IsHostBuf || IsConstantBuf || IsHostTask)>>
1639  accessor(
1640  buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1641  const property_list &PropertyList = {},
1642  const detail::code_location CodeLoc = detail::code_location::current())
1643 #ifdef __SYCL_DEVICE_ONLY__
1644  : impl(id<AdjustedDim>(), detail::GetZeroDimAccessRange(BufferRef),
1645  BufferRef.get_range()) {
1646  (void)CommandGroupHandler;
1647  (void)PropertyList;
1648  }
1649 #else
1650  : AccessorBaseHost(
1651  /*Offset=*/{0, 0, 0},
1652  detail::convertToArrayOfN<3, 1>(
1653  detail::GetZeroDimAccessRange(BufferRef)),
1654  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1655  getAdjustedMode(PropertyList),
1656  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1657  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1658  throwIfUsedByGraph();
1659  preScreenAccessor(PropertyList);
1660  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1661  initHostAcc();
1664  AccessTarget, AccessMode, CodeLoc);
1665  GDBMethodsAnchor();
1666  }
1667 #endif
1668 
1669  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1670  typename... PropTypes,
1671  typename = typename std::enable_if_t<
1673  std::is_same_v<T, DataT> && (Dims == 0) &&
1674  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1675  accessor(
1676  buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1677  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1678  {},
1679  const detail::code_location CodeLoc = detail::code_location::current())
1680 #ifdef __SYCL_DEVICE_ONLY__
1681  : impl(id<AdjustedDim>(), detail::GetZeroDimAccessRange(BufferRef),
1682  BufferRef.get_range()) {
1683  (void)CommandGroupHandler;
1684  (void)PropertyList;
1685  }
1686 #else
1687  : AccessorBaseHost(
1688  /*Offset=*/{0, 0, 0},
1689  detail::convertToArrayOfN<3, 1>(
1690  detail::GetZeroDimAccessRange(BufferRef)),
1691  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1692  getAdjustedMode(PropertyList),
1693  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1694  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1695  throwIfUsedByGraph();
1696  preScreenAccessor(PropertyList);
1697  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1698  initHostAcc();
1701  AccessTarget, AccessMode, CodeLoc);
1702  GDBMethodsAnchor();
1703  }
1704 #endif
1705 
1706  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1707  typename = std::enable_if_t<
1708  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1709  IsSameAsBuffer<T, Dims>::value &&
1710  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1711  accessor(
1712  buffer<T, Dims, AllocatorT> &BufferRef,
1713  const property_list &PropertyList = {},
1714  const detail::code_location CodeLoc = detail::code_location::current())
1715 #ifdef __SYCL_DEVICE_ONLY__
1716  : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1717  (void)PropertyList;
1718  }
1719 #else
1720  : AccessorBaseHost(
1721  /*Offset=*/{0, 0, 0},
1722  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1723  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1724  getAdjustedMode(PropertyList),
1725  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1726  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1727  PropertyList) {
1728  throwIfUsedByGraph();
1729  preScreenAccessor(PropertyList);
1732  initHostAcc();
1735  AccessTarget, AccessMode, CodeLoc);
1736  GDBMethodsAnchor();
1737  }
1738 #endif
1739 
1740  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1741  typename... PropTypes,
1742  typename = std::enable_if_t<
1744  IsSameAsBuffer<T, Dims>::value &&
1745  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1746  accessor(
1747  buffer<T, Dims, AllocatorT> &BufferRef,
1748  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1749  {},
1750  const detail::code_location CodeLoc = detail::code_location::current())
1751 #ifdef __SYCL_DEVICE_ONLY__
1752  : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1753  (void)PropertyList;
1754  }
1755 #else
1756  : AccessorBaseHost(
1757  /*Offset=*/{0, 0, 0},
1758  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1759  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1760  getAdjustedMode(PropertyList),
1761  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1762  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1763  PropertyList) {
1764  throwIfUsedByGraph();
1765  preScreenAccessor(PropertyList);
1768  initHostAcc();
1771  AccessTarget, AccessMode, CodeLoc);
1772  GDBMethodsAnchor();
1773  }
1774 #endif
1775 
1776  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1777  typename TagT,
1778  typename = std::enable_if_t<
1779  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1780  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1781  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1782  accessor(
1783  buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1784  const property_list &PropertyList = {},
1785  const detail::code_location CodeLoc = detail::code_location::current())
1786  : accessor(BufferRef, PropertyList, CodeLoc) {
1787  adjustAccPropsInBuf(BufferRef);
1788  }
1789 
1790  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1791  typename TagT, typename... PropTypes,
1792  typename = std::enable_if_t<
1794  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1795  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1796  accessor(
1797  buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1798  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1799  {},
1800  const detail::code_location CodeLoc = detail::code_location::current())
1801  : accessor(BufferRef, PropertyList, CodeLoc) {
1802  adjustAccPropsInBuf(BufferRef);
1803  }
1804 
1805  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1806  typename = std::enable_if_t<
1807  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1808  IsSameAsBuffer<T, Dims>::value &&
1809  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1810  accessor(
1811  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1812  const property_list &PropertyList = {},
1813  const detail::code_location CodeLoc = detail::code_location::current())
1814 #ifdef __SYCL_DEVICE_ONLY__
1815  : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1816  (void)CommandGroupHandler;
1817  (void)PropertyList;
1818  }
1819 #else
1820  : AccessorBaseHost(
1821  /*Offset=*/{0, 0, 0},
1822  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1823  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1824  getAdjustedMode(PropertyList),
1825  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1826  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1827  throwIfUsedByGraph();
1828  preScreenAccessor(PropertyList);
1829  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1830  initHostAcc();
1833  AccessTarget, AccessMode, CodeLoc);
1834  GDBMethodsAnchor();
1835  }
1836 #endif
1837 
1838  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1839  typename... PropTypes,
1840  typename = std::enable_if_t<
1842  IsSameAsBuffer<T, Dims>::value &&
1843  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1844  accessor(
1845  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1846  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1847  {},
1848  const detail::code_location CodeLoc = detail::code_location::current())
1849 #ifdef __SYCL_DEVICE_ONLY__
1850  : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1851  (void)CommandGroupHandler;
1852  (void)PropertyList;
1853  }
1854 #else
1855  : AccessorBaseHost(
1856  /*Offset=*/{0, 0, 0},
1857  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1858  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1859  getAdjustedMode(PropertyList),
1860  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1861  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1862  throwIfUsedByGraph();
1863  preScreenAccessor(PropertyList);
1864  initHostAcc();
1865  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1868  AccessTarget, AccessMode, CodeLoc);
1869  GDBMethodsAnchor();
1870  }
1871 #endif
1872 
1873  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1874  typename TagT,
1875  typename = std::enable_if_t<
1876  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1877  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1878  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1879  accessor(
1880  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1881  TagT, const property_list &PropertyList = {},
1882  const detail::code_location CodeLoc = detail::code_location::current())
1883  : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1884  adjustAccPropsInBuf(BufferRef);
1885  }
1886 
1887  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1888  typename TagT, typename... PropTypes,
1889  typename = std::enable_if_t<
1891  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1892  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1893  accessor(
1894  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1895  TagT,
1896  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1897  {},
1898  const detail::code_location CodeLoc = detail::code_location::current())
1899  : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1900  adjustAccPropsInBuf(BufferRef);
1901  }
1902 
1903  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1904  typename = std::enable_if_t<
1905  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1906  IsSameAsBuffer<T, Dims>::value &&
1907  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1908  accessor(
1909  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1910  const property_list &PropertyList = {},
1911  const detail::code_location CodeLoc = detail::code_location::current())
1912  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1913 
1914  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1915  typename... PropTypes,
1916  typename = std::enable_if_t<
1918  IsSameAsBuffer<T, Dims>::value &&
1919  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1920  accessor(
1921  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1922  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1923  {},
1924  const detail::code_location CodeLoc = detail::code_location::current())
1925  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1926 
1927  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1928  typename TagT,
1929  typename = std::enable_if_t<
1930  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1931  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1932  (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
1933  accessor(
1934  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1935  TagT, const property_list &PropertyList = {},
1936  const detail::code_location CodeLoc = detail::code_location::current())
1937  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1938  adjustAccPropsInBuf(BufferRef);
1939  }
1940 
1941  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1942  typename TagT, typename... PropTypes,
1943  typename = std::enable_if_t<
1945  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1946  (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
1947  accessor(
1948  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1949  TagT,
1950  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1951  {},
1952  const detail::code_location CodeLoc = detail::code_location::current())
1953  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1954  adjustAccPropsInBuf(BufferRef);
1955  }
1956 
1957  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1958  typename = std::enable_if_t<
1959  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1960  IsSameAsBuffer<T, Dims>::value &&
1961  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1962  accessor(
1963  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1964  range<Dimensions> AccessRange, const property_list &PropertyList = {},
1965  const detail::code_location CodeLoc = detail::code_location::current())
1966  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1967  CodeLoc) {}
1968 
1969  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1970  typename... PropTypes,
1971  typename = std::enable_if_t<
1973  IsSameAsBuffer<T, Dims>::value &&
1974  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1975  accessor(
1976  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1977  range<Dimensions> AccessRange,
1978  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1979  {},
1980  const detail::code_location CodeLoc = detail::code_location::current())
1981  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1982  CodeLoc) {}
1983 
1984  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1985  typename TagT,
1986  typename = std::enable_if_t<
1987  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1988  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1989  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1990  accessor(
1991  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1992  range<Dimensions> AccessRange, TagT,
1993  const property_list &PropertyList = {},
1994  const detail::code_location CodeLoc = detail::code_location::current())
1995  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1996  CodeLoc) {
1997  adjustAccPropsInBuf(BufferRef);
1998  }
1999 
2000  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2001  typename TagT, typename... PropTypes,
2002  typename = std::enable_if_t<
2004  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2005  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2006  accessor(
2007  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2008  range<Dimensions> AccessRange, TagT,
2009  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2010  {},
2011  const detail::code_location CodeLoc = detail::code_location::current())
2012  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
2013  CodeLoc) {
2014  adjustAccPropsInBuf(BufferRef);
2015  }
2016 
2017  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2018  typename = std::enable_if_t<
2019  detail::IsRunTimePropertyListT<PropertyListT>::value &&
2020  IsSameAsBuffer<T, Dims>::value &&
2021  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
2022  accessor(
2023  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
2024  id<Dimensions> AccessOffset, const property_list &PropertyList = {},
2025  const detail::code_location CodeLoc = detail::code_location::current())
2026 #ifdef __SYCL_DEVICE_ONLY__
2027  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
2028  (void)PropertyList;
2029  }
2030 #else
2031  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
2032  detail::convertToArrayOfN<3, 1>(AccessRange),
2033  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
2034  getAdjustedMode(PropertyList),
2035  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
2036  sizeof(DataT), IsPlaceH, BufferRef.OffsetInBytes,
2037  BufferRef.IsSubBuffer, PropertyList) {
2038  throwIfUsedByGraph();
2039  preScreenAccessor(PropertyList);
2042  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2043  BufferRef.get_range()))
2044  throw sycl::invalid_object_error(
2045  "accessor with requested offset and range would exceed the bounds of "
2046  "the buffer",
2047  PI_ERROR_INVALID_VALUE);
2048 
2049  initHostAcc();
2052  AccessTarget, AccessMode, CodeLoc);
2053  GDBMethodsAnchor();
2054  }
2055 #endif
2056 
2057  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2058  typename... PropTypes,
2059  typename = std::enable_if_t<
2061  IsSameAsBuffer<T, Dims>::value &&
2062  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
2063  accessor(
2064  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
2065  id<Dimensions> AccessOffset,
2066  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2067  {},
2068  const detail::code_location CodeLoc = detail::code_location::current())
2069 #ifdef __SYCL_DEVICE_ONLY__
2070  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
2071  (void)PropertyList;
2072  }
2073 #else
2074  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
2075  detail::convertToArrayOfN<3, 1>(AccessRange),
2076  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
2077  getAdjustedMode(PropertyList),
2078  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
2079  sizeof(DataT), IsPlaceH, BufferRef.OffsetInBytes,
2080  BufferRef.IsSubBuffer, PropertyList) {
2081  throwIfUsedByGraph();
2082  preScreenAccessor(PropertyList);
2085  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2086  BufferRef.get_range()))
2087  throw sycl::invalid_object_error(
2088  "accessor with requested offset and range would exceed the bounds of "
2089  "the buffer",
2090  PI_ERROR_INVALID_VALUE);
2091 
2092  initHostAcc();
2095  AccessTarget, AccessMode, CodeLoc);
2096  GDBMethodsAnchor();
2097  }
2098 #endif
2099 
2100  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2101  typename TagT,
2102  typename = std::enable_if_t<
2103  detail::IsRunTimePropertyListT<PropertyListT>::value &&
2104  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2105  (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
2106  accessor(
2107  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
2108  id<Dimensions> AccessOffset, TagT, const property_list &PropertyList = {},
2109  const detail::code_location CodeLoc = detail::code_location::current())
2110  : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
2111  adjustAccPropsInBuf(BufferRef);
2112  }
2113 
2114  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2115  typename TagT, typename... PropTypes,
2116  typename = std::enable_if_t<
2118  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2119  (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
2120  accessor(
2121  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
2122  id<Dimensions> AccessOffset, TagT,
2123  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2124  {},
2125  const detail::code_location CodeLoc = detail::code_location::current())
2126  : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
2127  adjustAccPropsInBuf(BufferRef);
2128  }
2129 
2130  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2131  typename = std::enable_if_t<
2132  detail::IsRunTimePropertyListT<PropertyListT>::value &&
2133  IsSameAsBuffer<T, Dims>::value &&
2134  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2135  accessor(
2136  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2137  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
2138  const property_list &PropertyList = {},
2139  const detail::code_location CodeLoc = detail::code_location::current())
2140 #ifdef __SYCL_DEVICE_ONLY__
2141  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
2142  (void)CommandGroupHandler;
2143  (void)PropertyList;
2144  }
2145 #else
2146  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
2147  detail::convertToArrayOfN<3, 1>(AccessRange),
2148  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
2149  getAdjustedMode(PropertyList),
2150  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
2151  sizeof(DataT), BufferRef.OffsetInBytes,
2152  BufferRef.IsSubBuffer, PropertyList) {
2153  throwIfUsedByGraph();
2154  preScreenAccessor(PropertyList);
2155  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2156  BufferRef.get_range()))
2157  throw sycl::invalid_object_error(
2158  "accessor with requested offset and range would exceed the bounds of "
2159  "the buffer",
2160  PI_ERROR_INVALID_VALUE);
2161 
2162  initHostAcc();
2163  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
2166  AccessTarget, AccessMode, CodeLoc);
2167  GDBMethodsAnchor();
2168  }
2169 #endif
2170 
2171  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2172  typename... PropTypes,
2173  typename = std::enable_if_t<
2175  IsSameAsBuffer<T, Dims>::value &&
2176  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2177  accessor(
2178  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2179  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
2180  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2181  {},
2182  const detail::code_location CodeLoc = detail::code_location::current())
2183 #ifdef __SYCL_DEVICE_ONLY__
2184  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
2185  (void)CommandGroupHandler;
2186  (void)PropertyList;
2187  }
2188 #else
2189  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
2190  detail::convertToArrayOfN<3, 1>(AccessRange),
2191  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
2192  getAdjustedMode(PropertyList),
2193  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
2194  sizeof(DataT), BufferRef.OffsetInBytes,
2195  BufferRef.IsSubBuffer, PropertyList) {
2196  throwIfUsedByGraph();
2197  preScreenAccessor(PropertyList);
2198  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2199  BufferRef.get_range()))
2200  throw sycl::invalid_object_error(
2201  "accessor with requested offset and range would exceed the bounds of "
2202  "the buffer",
2203  PI_ERROR_INVALID_VALUE);
2204 
2205  initHostAcc();
2206  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
2209  AccessTarget, AccessMode, CodeLoc);
2210  GDBMethodsAnchor();
2211  }
2212 #endif
2213 
2214  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2215  typename TagT,
2216  typename = std::enable_if_t<
2217  detail::IsRunTimePropertyListT<PropertyListT>::value &&
2218  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2219  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2220  accessor(
2221  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2222  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
2223  const property_list &PropertyList = {},
2224  const detail::code_location CodeLoc = detail::code_location::current())
2225  : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2226  PropertyList, CodeLoc) {
2227  adjustAccPropsInBuf(BufferRef);
2228  }
2229 
2230  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2231  typename TagT, typename... PropTypes,
2232  typename = std::enable_if_t<
2234  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2235  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2236  accessor(
2237  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2238  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
2239  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2240  {},
2241  const detail::code_location CodeLoc = detail::code_location::current())
2242  : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2243  PropertyList, CodeLoc) {
2244  adjustAccPropsInBuf(BufferRef);
2245  }
2246 
2247  template <typename... NewPropsT>
2248  accessor(
2249  const accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
2250  ext::oneapi::accessor_property_list<NewPropsT...>> &Other,
2251  const detail::code_location CodeLoc = detail::code_location::current())
2252 #ifdef __SYCL_DEVICE_ONLY__
2253  : impl(Other.impl), MData(Other.MData)
2254 #else
2255  : detail::AccessorBaseHost(Other), MAccData(Other.MAccData)
2256 #endif
2257  {
2259  "Conversion is only available for accessor_property_list");
2260  static_assert(
2261  PropertyListT::template areSameCompileTimeProperties<NewPropsT...>(),
2262  "Compile-time-constant properties must be the same");
2263 #ifndef __SYCL_DEVICE_ONLY__
2264  detail::constructorNotification(getMemoryObject(), impl.get(), AccessTarget,
2265  AccessMode, CodeLoc);
2266 #endif
2267  }
2268 
2269  void swap(accessor &other) {
2270  std::swap(impl, other.impl);
2271 #ifdef __SYCL_DEVICE_ONLY__
2272  std::swap(MData, other.MData);
2273 #else
2274  std::swap(MAccData, other.MAccData);
2275 #endif
2276  }
2277 
2278  bool is_placeholder() const {
2279 #ifdef __SYCL_DEVICE_ONLY__
2280  return false;
2281 #else
2283 #endif
2284  }
2285 
2286  size_t get_size() const { return getAccessRange().size() * sizeof(DataT); }
2287 
2288  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
2289  size_t get_count() const { return size(); }
2290  size_type size() const noexcept { return getAccessRange().size(); }
2291 
2292  size_type byte_size() const noexcept { return size() * sizeof(DataT); }
2293 
2294  size_type max_size() const noexcept {
2295  return empty() ? 0 : (std::numeric_limits<difference_type>::max)();
2296  }
2297 
2298  bool empty() const noexcept { return size() == 0; }
2299 
2300  template <int Dims = Dimensions,
2301  typename = std::enable_if_t<Dims == Dimensions && (Dims > 0)>>
2302  range<Dimensions> get_range() const {
2303  return getRange<Dims>();
2304  }
2305 
2306  template <int Dims = Dimensions,
2307  typename = std::enable_if_t<Dims == Dimensions && (Dims > 0)>>
2308  id<Dimensions> get_offset() const {
2309  return getOffset<Dims>();
2310  }
2311 
2312  template <int Dims = Dimensions, typename RefT = RefType,
2313  typename = std::enable_if_t<Dims == 0 &&
2314  (IsAccessAnyWrite || IsAccessReadOnly)>>
2315  operator reference() const {
2316  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
2317  return *(getQualifiedPtr() + LinearIndex);
2318  }
2319 
2320  template <int Dims = Dimensions,
2321  typename = std::enable_if_t<AccessMode != access_mode::atomic &&
2322  !IsAccessReadOnly && Dims == 0>>
2323  const accessor &operator=(const value_type &Other) const {
2324  *getQualifiedPtr() = Other;
2325  return *this;
2326  }
2327 
2328  template <int Dims = Dimensions,
2329  typename = std::enable_if_t<AccessMode != access_mode::atomic &&
2330  !IsAccessReadOnly && Dims == 0>>
2331  const accessor &operator=(value_type &&Other) const {
2332  *getQualifiedPtr() = std::move(Other);
2333  return *this;
2334  }
2335 
2336  template <int Dims = Dimensions,
2337  typename = std::enable_if_t<(Dims > 0) &&
2338  (IsAccessAnyWrite || IsAccessReadOnly)>>
2339  reference operator[](id<Dimensions> Index) const {
2340  const size_t LinearIndex = getLinearIndex(Index);
2341  return getQualifiedPtr()[LinearIndex];
2342  }
2343 
2344  template <int Dims = Dimensions>
2345  operator typename std::enable_if_t<Dims == 0 &&
2347 #ifdef __ENABLE_USM_ADDR_SPACE__
2348  atomic<DataT>
2349 #else
2350  atomic<DataT, AS>
2351 #endif
2352  >() const {
2353  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
2354  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2355  getQualifiedPtr() + LinearIndex));
2356  }
2357 
2358  template <int Dims = Dimensions>
2359  typename std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
2360  atomic<DataT, AS>>
2361  operator[](id<Dimensions> Index) const {
2362  const size_t LinearIndex = getLinearIndex(Index);
2363  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2364  getQualifiedPtr() + LinearIndex));
2365  }
2366 
2367  template <int Dims = Dimensions>
2368  typename std::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
2369  atomic<DataT, AS>>
2370  operator[](size_t Index) const {
2371  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
2372  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2373  getQualifiedPtr() + LinearIndex));
2374  }
2375  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 1)>>
2376  auto operator[](size_t Index) const {
2377  return AccessorSubscript<Dims - 1>(*this, Index);
2378  }
2379 
2380  template <access::target AccessTarget_ = AccessTarget,
2381  typename = std::enable_if_t<
2382  (AccessTarget_ == access::target::host_buffer) ||
2383  (AccessTarget_ == access::target::host_task)>>
2384  std::add_pointer_t<value_type> get_pointer() const noexcept {
2385  return getPointerAdjusted();
2386  }
2387 
2388  template <
2389  access::target AccessTarget_ = AccessTarget,
2390  typename = std::enable_if_t<(AccessTarget_ == access::target::device)>>
2392  "accessor::get_pointer() is deprecated, please use get_multi_ptr()")
2393  global_ptr<DataT> get_pointer() const noexcept {
2394  return global_ptr<DataT>(
2395  const_cast<typename detail::DecoratedType<DataT, AS>::type *>(
2396  getPointerAdjusted()));
2397  }
2398 
2399  template <access::target AccessTarget_ = AccessTarget,
2400  typename = std::enable_if_t<AccessTarget_ ==
2401  access::target::constant_buffer>>
2402  constant_ptr<DataT> get_pointer() const {
2403  return constant_ptr<DataT>(getPointerAdjusted());
2404  }
2405 
2406  template <access::decorated IsDecorated>
2407  accessor_ptr<IsDecorated> get_multi_ptr() const noexcept {
2408  return accessor_ptr<IsDecorated>(getPointerAdjusted());
2409  }
2410 
2411  // accessor::has_property for runtime properties is only available in host
2412  // code. This restriction is not listed in the core spec and will be added in
2413  // future versions.
2414  template <typename Property>
2415  typename std::enable_if_t<
2416  !ext::oneapi::is_compile_time_property<Property>::value, bool>
2417  has_property() const noexcept {
2418 #ifndef __SYCL_DEVICE_ONLY__
2419  return getPropList().template has_property<Property>();
2420 #else
2421  return false;
2422 #endif
2423  }
2424 
2425  // accessor::get_property for runtime properties is only available in host
2426  // code. This restriction is not listed in the core spec and will be added in
2427  // future versions.
2428  template <typename Property,
2429  typename = typename std::enable_if_t<
2430  !ext::oneapi::is_compile_time_property<Property>::value>>
2431  Property get_property() const {
2432 #ifndef __SYCL_DEVICE_ONLY__
2433  return getPropList().template get_property<Property>();
2434 #else
2435  return Property();
2436 #endif
2437  }
2438 
2439  template <typename Property>
2440  static constexpr bool has_property(
2441  typename std::enable_if_t<
2442  ext::oneapi::is_compile_time_property<Property>::value> * = 0) {
2443  return PropertyListT::template has_property<Property>();
2444  }
2445 
2446  template <typename Property>
2447  static constexpr auto get_property(
2448  typename std::enable_if_t<
2449  ext::oneapi::is_compile_time_property<Property>::value> * = 0) {
2450  return PropertyListT::template get_property<Property>();
2451  }
2452 
2453  bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
2454  bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
2455 
2456  iterator begin() const noexcept {
2457  return iterator::getBegin(
2458  get_pointer(),
2459  detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2460  getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2461  }
2462 
2463  iterator end() const noexcept {
2464  return iterator::getEnd(
2465  get_pointer(),
2466  detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2467  getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2468  }
2469 
2470  const_iterator cbegin() const noexcept {
2471  return const_iterator::getBegin(
2472  get_pointer(),
2473  detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2474  getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2475  }
2476 
2477  const_iterator cend() const noexcept {
2478  return const_iterator::getEnd(
2479  get_pointer(),
2480  detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2481  getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2482  }
2483 
2484  reverse_iterator rbegin() const noexcept { return reverse_iterator(end()); }
2485  reverse_iterator rend() const noexcept { return reverse_iterator(begin()); }
2486 
2487  const_reverse_iterator crbegin() const noexcept {
2488  return const_reverse_iterator(cend());
2489  }
2490  const_reverse_iterator crend() const noexcept {
2491  return const_reverse_iterator(cbegin());
2492  }
2493 
2494 private:
2495  template <int Dims, typename = std::enable_if_t<(Dims > 0)>>
2496  range<Dims> getRange() const {
2497  return detail::convertToArrayOfN<AdjustedDim, 1>(getAccessRange());
2498  }
2499 
2500  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
2501  id<Dims> getOffset() const {
2502  static_assert(
2503  !(PropertyListT::template has_property<
2505  "Accessor has no_offset property, get_offset() can not be used");
2506  return detail::convertToArrayOfN<Dims, 0>(getOffset());
2507  }
2508 
2509 #ifdef __SYCL_DEVICE_ONLY__
2510  size_t getTotalOffset() const noexcept {
2511  size_t TotalOffset = 0;
2512  detail::loop<Dimensions>([&, this](size_t I) {
2513  TotalOffset = TotalOffset * impl.MemRange[I];
2514  if constexpr (!(PropertyListT::template has_property<
2516  TotalOffset += impl.Offset[I];
2517  }
2518  });
2519 
2520  return TotalOffset;
2521  }
2522 #endif
2523 
2524  // supporting function for get_pointer()
2525  // MData has been preadjusted with offset for faster access with []
2526  // but for get_pointer() we must return the original pointer.
2527  // On device, getQualifiedPtr() returns MData, so we need to backjust it.
2528  // On host, getQualifiedPtr() does not return MData, no need to adjust.
2529  auto getPointerAdjusted() const noexcept {
2530 #ifdef __SYCL_DEVICE_ONLY__
2531  return getQualifiedPtr() - getTotalOffset();
2532 #else
2533  return getQualifiedPtr();
2534 #endif
2535  }
2536 
2537  void preScreenAccessor(const PropertyListT &PropertyList) {
2538  // check that no_init property is compatible with access mode
2539  if (PropertyList.template has_property<property::no_init>() &&
2541  throw sycl::invalid_object_error(
2542  "accessor would cannot be both read_only and no_init",
2543  PI_ERROR_INVALID_VALUE);
2544  }
2545  }
2546 
2547  template <typename BufT, typename... PropTypes>
2548  void adjustAccPropsInBuf(BufT &Buffer) {
2549  if constexpr (PropertyListT::template has_property<
2551  auto location = (PropertyListT::template get_property<
2553  .get_location();
2554  property_list PropList{
2556  Buffer.addOrReplaceAccessorProperties(PropList);
2557  } else {
2558  deleteAccPropsFromBuf(Buffer);
2559  }
2560  }
2561 
2562  template <typename BufT> void deleteAccPropsFromBuf(BufT &Buffer) {
2563  Buffer.deleteAccProps(
2565  }
2566 };
2567 
2568 template <typename DataT, int Dimensions, typename AllocatorT>
2569 accessor(buffer<DataT, Dimensions, AllocatorT>)
2570  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2572 
2573 template <typename DataT, int Dimensions, typename AllocatorT,
2574  typename... PropsT>
2575 accessor(buffer<DataT, Dimensions, AllocatorT>,
2576  const ext::oneapi::accessor_property_list<PropsT...> &)
2577  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2579  ext::oneapi::accessor_property_list<PropsT...>>;
2580 
2581 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
2582 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1)
2583  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
2584  detail::deduceAccessTarget<Type1, Type1>(target::device),
2586 
2587 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2588  typename... PropsT>
2589 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1,
2590  const ext::oneapi::accessor_property_list<PropsT...> &)
2591  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
2592  detail::deduceAccessTarget<Type1, Type1>(target::device),
2594  ext::oneapi::accessor_property_list<PropsT...>>;
2595 
2596 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2597  typename Type2>
2598 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2)
2599  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
2600  detail::deduceAccessTarget<Type1, Type2>(target::device),
2602 
2603 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2604  typename Type2, typename... PropsT>
2605 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2,
2606  const ext::oneapi::accessor_property_list<PropsT...> &)
2607  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
2608  detail::deduceAccessTarget<Type1, Type2>(target::device),
2610  ext::oneapi::accessor_property_list<PropsT...>>;
2611 
2612 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2613  typename Type2, typename Type3>
2614 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3)
2615  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
2616  detail::deduceAccessTarget<Type2, Type3>(target::device),
2618 
2619 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2620  typename Type2, typename Type3, typename... PropsT>
2621 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3,
2622  const ext::oneapi::accessor_property_list<PropsT...> &)
2623  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
2624  detail::deduceAccessTarget<Type2, Type3>(target::device),
2626  ext::oneapi::accessor_property_list<PropsT...>>;
2627 
2628 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2629  typename Type2, typename Type3, typename Type4>
2630 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
2631  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2632  detail::deduceAccessTarget<Type3, Type4>(target::device),
2634 
2635 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2636  typename Type2, typename Type3, typename Type4, typename... PropsT>
2637 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
2638  const ext::oneapi::accessor_property_list<PropsT...> &)
2639  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2640  detail::deduceAccessTarget<Type3, Type4>(target::device),
2642  ext::oneapi::accessor_property_list<PropsT...>>;
2643 
2644 template <typename DataT, int Dimensions, typename AllocatorT>
2645 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &)
2646  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2648 
2649 template <typename DataT, int Dimensions, typename AllocatorT,
2650  typename... PropsT>
2651 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &,
2652  const ext::oneapi::accessor_property_list<PropsT...> &)
2653  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2655  ext::oneapi::accessor_property_list<PropsT...>>;
2656 
2657 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
2658 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1)
2659  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
2660  detail::deduceAccessTarget<Type1, Type1>(target::device),
2662 
2663 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2664  typename... PropsT>
2665 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1,
2666  const ext::oneapi::accessor_property_list<PropsT...> &)
2667  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
2668  detail::deduceAccessTarget<Type1, Type1>(target::device),
2670  ext::oneapi::accessor_property_list<PropsT...>>;
2671 
2672 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2673  typename Type2>
2674 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2)
2675  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
2676  detail::deduceAccessTarget<Type1, Type2>(target::device),
2678 
2679 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2680  typename Type2, typename... PropsT>
2681 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2,
2682  const ext::oneapi::accessor_property_list<PropsT...> &)
2683  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
2684  detail::deduceAccessTarget<Type1, Type2>(target::device),
2686  ext::oneapi::accessor_property_list<PropsT...>>;
2687 
2688 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2689  typename Type2, typename Type3>
2690 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2, Type3)
2691  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
2692  detail::deduceAccessTarget<Type2, Type3>(target::device),
2694 
2695 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2696  typename Type2, typename Type3, typename... PropsT>
2697 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2, Type3,
2698  const ext::oneapi::accessor_property_list<PropsT...> &)
2699  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
2700  detail::deduceAccessTarget<Type2, Type3>(target::device),
2702  ext::oneapi::accessor_property_list<PropsT...>>;
2703 
2704 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2705  typename Type2, typename Type3, typename Type4>
2706 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2, Type3,
2707  Type4)
2708  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2709  detail::deduceAccessTarget<Type3, Type4>(target::device),
2711 
2712 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2713  typename Type2, typename Type3, typename Type4, typename... PropsT>
2714 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2, Type3,
2715  Type4, const ext::oneapi::accessor_property_list<PropsT...> &)
2716  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2717  detail::deduceAccessTarget<Type3, Type4>(target::device),
2719  ext::oneapi::accessor_property_list<PropsT...>>;
2720 
2724 template <typename DataT, int Dimensions, access::mode AccessMode,
2727 #ifndef __SYCL_DEVICE_ONLY__
2728  public detail::LocalAccessorBaseHost,
2729 #endif
2730  public detail::accessor_common<DataT, Dimensions, AccessMode,
2731  access::target::local, IsPlaceholder> {
2732 protected:
2733  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
2734 
2735  using AccessorCommonT =
2736  detail::accessor_common<DataT, Dimensions, AccessMode,
2737  access::target::local, IsPlaceholder>;
2738 
2739  using AccessorCommonT::AS;
2740 
2741  // Cannot do "using AccessorCommonT::Flag" as it doesn't work with g++ as host
2742  // compiler, for some reason.
2743  static constexpr bool IsAccessAnyWrite = AccessorCommonT::IsAccessAnyWrite;
2744  static constexpr bool IsAccessReadOnly = AccessorCommonT::IsAccessReadOnly;
2745  static constexpr bool IsConst = AccessorCommonT::IsConst;
2746 
2747  template <int Dims>
2748  using AccessorSubscript =
2749  typename AccessorCommonT::template AccessorSubscript<
2750  Dims,
2752 
2754 
2757 
2758 #ifdef __SYCL_DEVICE_ONLY__
2760 
2761  sycl::range<AdjustedDim> &getSize() { return impl.MemRange; }
2762  const sycl::range<AdjustedDim> &getSize() const { return impl.MemRange; }
2763 
2764  void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
2766  MData = Ptr;
2767  detail::loop<AdjustedDim>(
2768  [&, this](size_t I) { getSize()[I] = AccessRange[I]; });
2769  }
2770 
2771  // __init variant used by the device compiler for ESIMD kernels.
2772  // TODO: In ESIMD accessors usage is limited for now - access range, mem
2773  // range and offset are not supported.
2774  void __init_esimd(ConcreteASPtrType Ptr) {
2775  MData = Ptr;
2776  detail::loop<AdjustedDim>([&, this](size_t I) { getSize()[I] = 0; });
2777  }
2778 
2779 public:
2780  // Default constructor for objects later initialized with __init member.
2781  local_accessor_base()
2782  : impl(detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
2783 
2784 protected:
2785  ConcreteASPtrType getQualifiedPtr() const { return MData; }
2786 
2787  ConcreteASPtrType MData;
2788 
2789 #else
2790 public:
2792  : detail::LocalAccessorBaseHost{/*Size*/ sycl::range<3>{0, 0, 0},
2793  /*Dims*/ 0, /*ElemSize*/ sizeof(DataT)} {}
2794 
2795 protected:
2797  : detail::LocalAccessorBaseHost{Impl} {}
2798 
2799  char padding[sizeof(detail::LocalAccessorBaseDevice<AdjustedDim>) +
2800  sizeof(PtrType) - sizeof(detail::LocalAccessorBaseHost)];
2802 
2804  return reinterpret_cast<PtrType>(LocalAccessorBaseHost::getPtr());
2805  }
2806 
2808  void *getPtr() const { return detail::LocalAccessorBaseHost::getPtr(); }
2809  const range<3> &getSize() const {
2811  }
2813 
2814  // The function references helper methods required by GDB pretty-printers
2816 #ifndef NDEBUG
2817  const auto *this_const = this;
2818  (void)getSize();
2819  (void)this_const->getSize();
2820  (void)getPtr();
2821  (void)this_const->getPtr();
2822 #endif
2823  }
2824 
2825 #endif // __SYCL_DEVICE_ONLY__
2826 
2827  // Method which calculates linear offset for the ID using Range and Offset.
2828  template <int Dims = AdjustedDim> size_t getLinearIndex(id<Dims> Id) const {
2829  size_t Result = 0;
2830  detail::loop<Dims>(
2831  [&, this](size_t I) { Result = Result * getSize()[I] + Id[I]; });
2832  return Result;
2833  }
2834 
2835  template <class Obj>
2836  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
2837 
2838  template <class T>
2839  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
2840 
2841  template <typename DataT_, int Dimensions_> friend class local_accessor;
2842 
2843 public:
2844  using value_type = DataT;
2845  using reference = DataT &;
2846  using const_reference = const DataT &;
2847 
2848  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 0>>
2851 #ifdef __SYCL_DEVICE_ONLY__
2852  : impl(range<AdjustedDim>{1}){}
2853 #else
2854  : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT)) {
2856  access::target::local, AccessMode, CodeLoc);
2857  GDBMethodsAnchor();
2858  }
2859 #endif
2860 
2861  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 0>>
2863  const detail::code_location CodeLoc =
2865 #ifdef __SYCL_DEVICE_ONLY__
2866  : impl(range<AdjustedDim>{1}) {
2867  (void)propList;
2868  }
2869 #else
2870  : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT),
2871  propList) {
2873  access::target::local, AccessMode, CodeLoc);
2874  GDBMethodsAnchor();
2875  }
2876 #endif
2877 
2878  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
2880  range<Dimensions> AllocationSize, handler &,
2882 #ifdef __SYCL_DEVICE_ONLY__
2883  : impl(AllocationSize){}
2884 #else
2885  : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2886  AdjustedDim, sizeof(DataT)) {
2888  access::target::local, AccessMode, CodeLoc);
2889  GDBMethodsAnchor();
2890  }
2891 #endif
2892 
2893  template <int Dims = Dimensions,
2894  typename = std::enable_if_t<(Dims > 0)>>
2896  const property_list &propList,
2897  const detail::code_location CodeLoc =
2899 #ifdef __SYCL_DEVICE_ONLY__
2900  : impl(AllocationSize) {
2901  (void)propList;
2902  }
2903 #else
2904  : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2905  AdjustedDim, sizeof(DataT), propList) {
2907  access::target::local, AccessMode, CodeLoc);
2908  GDBMethodsAnchor();
2909  }
2910 #endif
2911 
2912  size_t get_size() const { return getSize().size() * sizeof(DataT); }
2913 
2914  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
2915  size_t get_count() const { return size(); }
2916  size_t size() const noexcept { return getSize().size(); }
2917 
2918  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
2920  return detail::convertToArrayOfN<Dims, 1>(getSize());
2921  }
2922 
2923  template <int Dims = Dimensions,
2924  typename = std::enable_if_t<Dims == 0 &&
2925  (IsAccessAnyWrite || IsAccessReadOnly)>>
2926  operator RefType() const {
2927  return *getQualifiedPtr();
2928  }
2929 
2930  template <int Dims = Dimensions,
2931  typename = std::enable_if_t<(Dims > 0) &&
2932  (IsAccessAnyWrite || IsAccessReadOnly)>>
2934  const size_t LinearIndex = getLinearIndex(Index);
2935  return getQualifiedPtr()[LinearIndex];
2936  }
2937 
2938  template <int Dims = Dimensions,
2939  typename = std::enable_if_t<Dims == 1 &&
2940  (IsAccessAnyWrite || IsAccessReadOnly)>>
2941  RefType operator[](size_t Index) const {
2942  return getQualifiedPtr()[Index];
2943  }
2944 
2945  template <int Dims = Dimensions>
2946  operator typename std::enable_if_t<
2947  Dims == 0 && AccessMode == access::mode::atomic, atomic<DataT, AS>>()
2948  const {
2949  return atomic<DataT, AS>(
2951  }
2952 
2953  template <int Dims = Dimensions>
2954  typename std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
2955  atomic<DataT, AS>>
2956  operator[](id<Dimensions> Index) const {
2957  const size_t LinearIndex = getLinearIndex(Index);
2958  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2959  getQualifiedPtr() + LinearIndex));
2960  }
2961 
2962  template <int Dims = Dimensions>
2963  typename std::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
2964  atomic<DataT, AS>>
2965  operator[](size_t Index) const {
2966  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2967  getQualifiedPtr() + Index));
2968  }
2969 
2970  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 1)>>
2971  typename AccessorCommonT::template AccessorSubscript<
2972  Dims - 1,
2974  operator[](size_t Index) const {
2975  return AccessorSubscript<Dims - 1>(*this, Index);
2976  }
2977 
2978  bool operator==(const local_accessor_base &Rhs) const {
2979  return impl == Rhs.impl;
2980  }
2981  bool operator!=(const local_accessor_base &Rhs) const {
2982  return !(*this == Rhs);
2983  }
2984 };
2985 
2986 // TODO: Remove deprecated specialization once no longer needed
2987 template <typename DataT, int Dimensions, access::mode AccessMode,
2990  DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder>
2991  : public local_accessor_base<DataT, Dimensions, AccessMode, IsPlaceholder>,
2992  public detail::OwnerLessBase<
2993  accessor<DataT, Dimensions, AccessMode, access::target::local,
2994  IsPlaceholder>> {
2995 
2996  using local_acc =
2998 
2999  static_assert(
3000  !local_acc::IsConst || local_acc::IsAccessReadOnly,
3001  "A const qualified DataT is only allowed for a read-only accessor");
3002 
3003  // Use base classes constructors
3004  using local_acc::local_acc;
3005 
3006 public:
3008  return local_ptr<DataT>(local_acc::getQualifiedPtr());
3009  }
3010 
3011 #ifdef __SYCL_DEVICE_ONLY__
3012 
3013  // __init needs to be defined within the class not through inheritance.
3014  // Map this function to inherited func.
3015  void __init(typename local_acc::ConcreteASPtrType Ptr,
3016  range<local_acc::AdjustedDim> AccessRange,
3019  local_acc::__init(Ptr, AccessRange, range, id);
3020  }
3021 
3022  // __init variant used by the device compiler for ESIMD kernels.
3023  // TODO: In ESIMD accessors usage is limited for now - access range, mem
3024  // range and offset are not supported.
3025  void __init_esimd(typename local_acc::ConcreteASPtrType Ptr) {
3026  local_acc::__init_esimd(Ptr);
3027  }
3028 
3029 public:
3030  // Default constructor for objects later initialized with __init member.
3031  accessor() {
3032  local_acc::impl = detail::InitializedVal<local_acc::AdjustedDim,
3033  range>::template get<0>();
3034  }
3035 
3036 #else
3037 private:
3038  accessor(const detail::AccessorImplPtr &Impl) : local_acc{Impl} {}
3039 #endif
3040 };
3041 
3042 template <typename DataT, int Dimensions = 1>
3043 class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
3044  : public local_accessor_base<DataT, Dimensions,
3045  detail::accessModeFromConstness<DataT>(),
3046  access::placeholder::false_t>,
3047  public detail::OwnerLessBase<local_accessor<DataT, Dimensions>> {
3048 
3049  using local_acc =
3050  local_accessor_base<DataT, Dimensions,
3051  detail::accessModeFromConstness<DataT>(),
3053 
3054  static_assert(
3055  !local_acc::IsConst || local_acc::IsAccessReadOnly,
3056  "A const qualified DataT is only allowed for a read-only accessor");
3057 
3058  // Use base classes constructors
3059  using local_acc::local_acc;
3060 
3061 #ifdef __SYCL_DEVICE_ONLY__
3062 
3063  // __init needs to be defined within the class not through inheritance.
3064  // Map this function to inherited func.
3065  void __init(typename local_acc::ConcreteASPtrType Ptr,
3066  range<local_acc::AdjustedDim> AccessRange,
3067  range<local_acc::AdjustedDim> range,
3068  id<local_acc::AdjustedDim> id) {
3069  local_acc::__init(Ptr, AccessRange, range, id);
3070  }
3071 
3072  // __init variant used by the device compiler for ESIMD kernels.
3073  // TODO: In ESIMD accessors usage is limited for now - access range, mem
3074  // range and offset are not supported.
3075  void __init_esimd(typename local_acc::ConcreteASPtrType Ptr) {
3076  local_acc::__init_esimd(Ptr);
3077  }
3078 
3079 public:
3080  // Default constructor for objects later initialized with __init member.
3081  local_accessor() {
3082  local_acc::impl = detail::InitializedVal<local_acc::AdjustedDim,
3083  range>::template get<0>();
3084  }
3085 
3086 #else
3087  local_accessor(const detail::AccessorImplPtr &Impl) : local_acc{Impl} {}
3088 #endif
3089 
3090  // implicit conversion between non-const read-write accessor to const
3091  // read-only accessor
3092 public:
3093  template <typename DataT_,
3094  typename = std::enable_if_t<
3095  std::is_const_v<DataT> &&
3096  std::is_same_v<DataT_, std::remove_const_t<DataT>>>>
3097  local_accessor(const local_accessor<DataT_, Dimensions> &other) {
3098  local_acc::impl = other.impl;
3099 #ifdef __SYCL_DEVICE_ONLY__
3100  local_acc::MData = other.MData;
3101 #endif
3102  }
3103 
3104  using value_type = DataT;
3105  using iterator = value_type *;
3106  using const_iterator = const value_type *;
3107  using reverse_iterator = std::reverse_iterator<iterator>;
3108  using const_reverse_iterator = std::reverse_iterator<const_iterator>;
3109  using difference_type =
3111  using size_type = std::size_t;
3112 
3113  template <access::decorated IsDecorated>
3114  using accessor_ptr = local_ptr<value_type, IsDecorated>;
3115 
3116  template <typename DataT_>
3117  bool operator==(const local_accessor<DataT_, Dimensions> &Rhs) const {
3118  return local_acc::impl == Rhs.impl;
3119  }
3120 
3121  template <typename DataT_>
3122  bool operator!=(const local_accessor<DataT_, Dimensions> &Rhs) const {
3123  return !(*this == Rhs);
3124  }
3125 
3126  void swap(local_accessor &other) { std::swap(this->impl, other.impl); }
3127 
3128  size_type byte_size() const noexcept { return this->size() * sizeof(DataT); }
3129 
3130  size_type max_size() const noexcept {
3131  return empty() ? 0 : (std::numeric_limits<difference_type>::max)();
3132  }
3133 
3134  bool empty() const noexcept { return this->size() == 0; }
3135 
3136  iterator begin() const noexcept {
3137  if constexpr (Dimensions == 0)
3138  return local_acc::getQualifiedPtr();
3139  else
3140  return &this->operator[](id<Dimensions>());
3141  }
3142  iterator end() const noexcept {
3143  if constexpr (Dimensions == 0)
3144  return begin() + 1;
3145  else
3146  return begin() + this->size();
3147  }
3148 
3149  const_iterator cbegin() const noexcept { return const_iterator(begin()); }
3150  const_iterator cend() const noexcept { return const_iterator(end()); }
3151 
3152  reverse_iterator rbegin() const noexcept { return reverse_iterator(end()); }
3153  reverse_iterator rend() const noexcept { return reverse_iterator(begin()); }
3154 
3155  const_reverse_iterator crbegin() const noexcept {
3156  return const_reverse_iterator(end());
3157  }
3158  const_reverse_iterator crend() const noexcept {
3159  return const_reverse_iterator(begin());
3160  }
3161 
3163  "local_accessor::get_pointer() is deprecated, please use get_multi_ptr()")
3164  local_ptr<DataT> get_pointer() const noexcept {
3165  return local_ptr<DataT>(local_acc::getQualifiedPtr());
3166  }
3167 
3168  template <access::decorated IsDecorated>
3169  accessor_ptr<IsDecorated> get_multi_ptr() const noexcept {
3170  return accessor_ptr<IsDecorated>(local_acc::getQualifiedPtr());
3171  }
3172 
3173  template <typename Property> bool has_property() const noexcept {
3174 #ifndef __SYCL_DEVICE_ONLY__
3175  return this->getPropList().template has_property<Property>();
3176 #else
3177  return false;
3178 #endif
3179  }
3180 
3181  template <typename Property> Property get_property() const {
3182 #ifndef __SYCL_DEVICE_ONLY__
3183  return this->getPropList().template get_property<Property>();
3184 #else
3185  return Property();
3186 #endif
3187  }
3188 
3189  template <int Dims = Dimensions,
3190  typename = std::enable_if_t<!std::is_const_v<DataT> && Dims == 0>>
3191  const local_accessor &operator=(const value_type &Other) const {
3192  *local_acc::getQualifiedPtr() = Other;
3193  return *this;
3194  }
3195 
3196  template <int Dims = Dimensions,
3197  typename = std::enable_if_t<!std::is_const_v<DataT> && Dims == 0>>
3198  const local_accessor &operator=(value_type &&Other) const {
3199  *local_acc::getQualifiedPtr() = std::move(Other);
3200  return *this;
3201  }
3202 
3203 private:
3204  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
3205 };
3206 
3212 template <typename DataT, int Dimensions, access::mode AccessMode,
3215  DataT, Dimensions, AccessMode, access::target::image, IsPlaceholder>
3216  : public detail::image_accessor<DataT, Dimensions, AccessMode,
3217  access::target::image, IsPlaceholder>,
3218  public detail::OwnerLessBase<
3219  accessor<DataT, Dimensions, AccessMode, access::target::image,
3220  IsPlaceholder>> {
3221 private:
3222  accessor(const detail::AccessorImplPtr &Impl)
3225 
3226 public:
3227  template <typename AllocatorT>
3229  handler &CommandGroupHandler)
3232  Image, CommandGroupHandler, Image.getElementSize()) {
3233 #ifndef __SYCL_DEVICE_ONLY__
3234  detail::associateWithHandler(CommandGroupHandler, this,
3236 #endif
3237  }
3238 
3239  template <typename AllocatorT>
3241  handler &CommandGroupHandler, const property_list &propList)
3244  Image, CommandGroupHandler, Image.getElementSize()) {
3245  (void)propList;
3246 #ifndef __SYCL_DEVICE_ONLY__
3247  detail::associateWithHandler(CommandGroupHandler, this,
3249 #endif
3250  }
3251 #ifdef __SYCL_DEVICE_ONLY__
3252 private:
3253  using OCLImageTy =
3254  typename detail::opencl_image_type<Dimensions, AccessMode,
3255  access::target::image>::type;
3256 
3257  // Front End requires this method to be defined in the accessor class.
3258  // It does not call the base class's init method.
3259  void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
3260 
3261  // __init variant used by the device compiler for ESIMD kernels.
3262  void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
3263 
3264 public:
3265  // Default constructor for objects later initialized with __init member.
3266  accessor() = default;
3267 #endif
3268 };
3269 
3277 template <typename DataT, int Dimensions, access::mode AccessMode,
3280  access::target::host_image, IsPlaceholder>
3281  : public detail::image_accessor<DataT, Dimensions, AccessMode,
3282  access::target::host_image, IsPlaceholder>,
3283  public detail::OwnerLessBase<
3284  accessor<DataT, Dimensions, AccessMode, access::target::host_image,
3285  IsPlaceholder>> {
3286 public:
3287  template <typename AllocatorT>
3290  access::target::host_image, IsPlaceholder>(
3291  Image, Image.getElementSize()) {}
3292 
3293  template <typename AllocatorT>
3295  const property_list &propList)
3297  access::target::host_image, IsPlaceholder>(
3298  Image, Image.getElementSize()) {
3299  (void)propList;
3300  }
3301 };
3302 
3311 template <typename DataT, int Dimensions, access::mode AccessMode,
3314  DataT, Dimensions, AccessMode, access::target::image_array, IsPlaceholder>
3315  : public detail::image_accessor<DataT, Dimensions + 1, AccessMode,
3316  access::target::image, IsPlaceholder>,
3317  public detail::OwnerLessBase<
3318  accessor<DataT, Dimensions, AccessMode, access::target::image_array,
3319  IsPlaceholder>> {
3320 #ifdef __SYCL_DEVICE_ONLY__
3321 private:
3322  using OCLImageTy =
3323  typename detail::opencl_image_type<Dimensions + 1, AccessMode,
3324  access::target::image>::type;
3325 
3326  // Front End requires this method to be defined in the accessor class.
3327  // It does not call the base class's init method.
3328  void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
3329 
3330  // __init variant used by the device compiler for ESIMD kernels.
3331  void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
3332 
3333 public:
3334  // Default constructor for objects later initialized with __init member.
3335  accessor() = default;
3336 #endif
3337 public:
3338  template <typename AllocatorT>
3340  handler &CommandGroupHandler)
3343  Image, CommandGroupHandler, Image.getElementSize()) {
3344 #ifndef __SYCL_DEVICE_ONLY__
3345  detail::associateWithHandler(CommandGroupHandler, this,
3347 #endif
3348  }
3349 
3350  template <typename AllocatorT>
3352  handler &CommandGroupHandler, const property_list &propList)
3355  Image, CommandGroupHandler, Image.getElementSize()) {
3356  (void)propList;
3357 #ifndef __SYCL_DEVICE_ONLY__
3358  detail::associateWithHandler(CommandGroupHandler, this,
3360 #endif
3361  }
3362 
3364  operator[](size_t Index) const {
3366  IsPlaceholder>(*this, Index);
3367  }
3368 };
3369 
3370 template <typename DataT, int Dimensions = 1,
3373  : public accessor<DataT, Dimensions, AccessMode, target::host_buffer,
3374  access::placeholder::false_t> {
3375 protected:
3376  using AccessorT = accessor<DataT, Dimensions, AccessMode, target::host_buffer,
3378 
3379  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
3380  constexpr static bool IsAccessReadOnly = AccessMode == access::mode::read;
3381 
3382  template <typename T, int Dims>
3384  : std::bool_constant<std::is_same_v<T, DataT> && (Dims > 0) &&
3385  (Dims == Dimensions)> {};
3386 
3387  void
3388  __init(typename accessor<DataT, Dimensions, AccessMode, target::host_buffer,
3389  access::placeholder::false_t>::ConcreteASPtrType Ptr,
3390  range<AdjustedDim> AccessRange, range<AdjustedDim> MemRange,
3391  id<AdjustedDim> Offset) {
3392  AccessorT::__init(Ptr, AccessRange, MemRange, Offset);
3393  }
3394 
3395 #ifndef __SYCL_DEVICE_ONLY__
3397  : accessor<DataT, Dimensions, AccessMode, target::host_buffer,
3398  access::placeholder::false_t>{Impl} {}
3399 
3400  template <class Obj>
3401  friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject);
3402 
3403  template <class T>
3404  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
3405 #endif // __SYCL_DEVICE_ONLY__
3406 
3407 public:
3408  host_accessor() : AccessorT() {}
3409 
3410  // The list of host_accessor constructors with their arguments
3411  // -------+---------+-------+----+----------+--------------
3412  // Dimensions = 0
3413  // -------+---------+-------+----+----------+--------------
3414  // buffer | | | | | property_list
3415  // buffer | handler | | | | property_list
3416  // -------+---------+-------+----+----------+--------------
3417  // Dimensions >= 1
3418  // -------+---------+-------+----+----------+--------------
3419  // buffer | | | | | property_list
3420  // buffer | | | | mode_tag | property_list
3421  // buffer | handler | | | | property_list
3422  // buffer | handler | | | mode_tag | property_list
3423  // buffer | | range | | | property_list
3424  // buffer | | range | | mode_tag | property_list
3425  // buffer | handler | range | | | property_list
3426  // buffer | handler | range | | mode_tag | property_list
3427  // buffer | | range | id | | property_list
3428  // buffer | | range | id | mode_tag | property_list
3429  // buffer | handler | range | id | | property_list
3430  // buffer | handler | range | id | mode_tag | property_list
3431  // -------+---------+-------+----+----------+--------------
3432 
3433  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3434  typename = typename std::enable_if_t<std::is_same_v<T, DataT> &&
3435  Dims == 0>>
3437  buffer<T, 1, AllocatorT> &BufferRef,
3438  const property_list &PropertyList = {},
3440  : AccessorT(BufferRef, PropertyList, CodeLoc) {}
3441 
3442  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3443  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3445  buffer<T, Dims, AllocatorT> &BufferRef,
3446  const property_list &PropertyList = {},
3448  : AccessorT(BufferRef, PropertyList, CodeLoc) {}
3449 
3450  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3451  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3454  const property_list &PropertyList = {},
3456  : host_accessor(BufferRef, PropertyList, CodeLoc) {}
3457 
3458  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3459  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3461  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3462  const property_list &PropertyList = {},
3464  : AccessorT(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
3465 
3466  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3467  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3469  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3470  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
3472  : host_accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
3473 
3474  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3475  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3477  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3478  const property_list &PropertyList = {},
3480  : AccessorT(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
3481 
3482  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3483  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3485  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3486  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
3488  : host_accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
3489 
3490  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3491  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3493  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3494  range<Dimensions> AccessRange, const property_list &PropertyList = {},
3496  : AccessorT(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
3497  CodeLoc) {}
3498 
3499  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3500  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3502  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3504  const property_list &PropertyList = {},
3506  : host_accessor(BufferRef, CommandGroupHandler, AccessRange, {},
3507  PropertyList, CodeLoc) {}
3508 
3509  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3510  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3512  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3513  id<Dimensions> AccessOffset, const property_list &PropertyList = {},
3515  : AccessorT(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
3516  }
3517 
3518  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3519  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3521  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3522  id<Dimensions> AccessOffset, mode_tag_t<AccessMode>,
3523  const property_list &PropertyList = {},
3525  : host_accessor(BufferRef, AccessRange, AccessOffset, PropertyList,
3526  CodeLoc) {}
3527 
3528  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3529  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3531  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3532  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
3533  const property_list &PropertyList = {},
3535  : AccessorT(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
3536  PropertyList, CodeLoc) {}
3537 
3538  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3539  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3541  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3542  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
3543  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
3545  : host_accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
3546  PropertyList, CodeLoc) {}
3547 
3548  template <int Dims = Dimensions,
3549  typename = std::enable_if_t<AccessMode != access_mode::atomic &&
3550  !IsAccessReadOnly && Dims == 0>>
3551  const host_accessor &
3552  operator=(const typename AccessorT::value_type &Other) const {
3553  *AccessorT::getQualifiedPtr() = Other;
3554  return *this;
3555  }
3556 
3557  template <int Dims = Dimensions,
3558  typename = std::enable_if_t<AccessMode != access_mode::atomic &&
3559  !IsAccessReadOnly && Dims == 0>>
3560  const host_accessor &operator=(typename AccessorT::value_type &&Other) const {
3561  *AccessorT::getQualifiedPtr() = std::move(Other);
3562  return *this;
3563  }
3564 
3565  // implicit conversion between const / non-const types for read only accessors
3566  template <typename DataT_,
3567  typename = std::enable_if_t<
3568  IsAccessReadOnly && !std::is_same_v<DataT_, DataT> &&
3569  std::is_same_v<std::remove_const_t<DataT_>,
3570  std::remove_const_t<DataT>>>>
3572 #ifndef __SYCL_DEVICE_ONLY__
3573  : host_accessor(other.impl) {
3574  AccessorT::MAccData = other.MAccData;
3575 #else
3576  {
3577 #endif // __SYCL_DEVICE_ONLY__
3578  }
3579 
3580  // implicit conversion from read_write T accessor to read only T (const)
3581  // accessor
3582  template <typename DataT_, access::mode AccessMode_,
3583  typename = std::enable_if_t<
3584  (AccessMode_ == access_mode::read_write) && IsAccessReadOnly &&
3585  std::is_same_v<DataT_, std::remove_const_t<DataT>>>>
3587 #ifndef __SYCL_DEVICE_ONLY__
3588  : host_accessor(other.impl) {
3589  AccessorT::MAccData = other.MAccData;
3590 #else
3591  {
3592 #endif // __SYCL_DEVICE_ONLY__
3593  }
3594 
3595  // host_accessor needs to explicitly define the owner_before member functions
3596  // as inheriting from OwnerLessBase causes base class conflicts.
3597  // TODO: Once host_accessor is detached from accessor, inherit from
3598  // OwnerLessBase instead.
3599 #ifndef __SYCL_DEVICE_ONLY__
3602  const noexcept {
3603  return this->impl.owner_before(
3605  }
3606 
3607  bool ext_oneapi_owner_before(const host_accessor &Other) const noexcept {
3608  return this->impl.owner_before(Other.impl);
3609  }
3610 #else
3611  bool ext_oneapi_owner_before(
3613  const noexcept;
3614  bool ext_oneapi_owner_before(const host_accessor &Other) const noexcept;
3615 #endif
3616 };
3617 
3618 template <typename DataT, int Dimensions, typename AllocatorT>
3621 
3622 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
3624  -> host_accessor<DataT, Dimensions,
3625  detail::deduceAccessMode<Type1, Type1>()>;
3626 
3627 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3628  typename Type2>
3630  -> host_accessor<DataT, Dimensions,
3631  detail::deduceAccessMode<Type1, Type2>()>;
3632 
3633 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3634  typename Type2, typename Type3>
3636  -> host_accessor<DataT, Dimensions,
3637  detail::deduceAccessMode<Type2, Type3>()>;
3638 
3639 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3640  typename Type2, typename Type3, typename Type4>
3641 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
3642  -> host_accessor<DataT, Dimensions,
3643  detail::deduceAccessMode<Type3, Type4>()>;
3644 
3645 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3646  typename Type2, typename Type3, typename Type4, typename Type5>
3647 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
3648  Type5) -> host_accessor<DataT, Dimensions,
3649  detail::deduceAccessMode<Type4, Type5>()>;
3650 
3651 // SYCL 2020 image accessors
3652 
3653 template <typename DataT, int Dimensions, access_mode AccessMode,
3654  image_target AccessTarget = image_target::device>
3656 #ifndef __SYCL_DEVICE_ONLY__
3657  private detail::UnsampledImageAccessorBaseHost,
3658 #endif // __SYCL_DEVICE_ONLY__
3659  public detail::OwnerLessBase<
3660  unsampled_image_accessor<DataT, Dimensions, AccessMode, AccessTarget>> {
3661  static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
3662  std::is_same_v<DataT, float4> ||
3663  std::is_same_v<DataT, half4>,
3664  "The data type of an image accessor must be only int4, "
3665  "uint4, float4 or half4 from SYCL namespace");
3666  static_assert(AccessMode == access_mode::read ||
3668  "Access mode must be either read or write.");
3669 
3670 #ifdef __SYCL_DEVICE_ONLY__
3671  char MPadding[sizeof(detail::UnsampledImageAccessorBaseHost)];
3672 #else
3673  using host_base_class = detail::UnsampledImageAccessorBaseHost;
3674 #endif // __SYCL_DEVICE_ONLY__
3675 
3676 public:
3677  using value_type = typename std::conditional<AccessMode == access_mode::read,
3678  const DataT, DataT>::type;
3679  using reference = value_type &;
3680  using const_reference = const DataT &;
3681 
3682  template <typename AllocatorT>
3685  handler &CommandGroupHandlerRef, const property_list &PropList = {},
3687 #ifdef __SYCL_DEVICE_ONLY__
3688  {}
3689 #else
3690  : host_base_class(detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
3691  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
3692  Dimensions, ImageRef.getElementSize(),
3693  {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
3694  ImageRef.getChannelType(), ImageRef.getChannelOrder(),
3695  PropList) {
3696  device Device = detail::getDeviceFromHandler(CommandGroupHandlerRef);
3697  // Avoid aspect::image warning.
3698  aspect ImageAspect = aspect::image;
3699  if (AccessTarget == image_target::device && !Device.has(ImageAspect))
3700  throw sycl::exception(
3701  sycl::make_error_code(sycl::errc::feature_not_supported),
3702  "Device associated with command group handler does not have "
3703  "aspect::image.");
3704 
3706  detail::getSyclObjImpl(ImageRef).get(), this->impl.get(), AccessTarget,
3707  AccessMode, (const void *)typeid(DataT).name(), sizeof(DataT), CodeLoc);
3708  detail::associateWithHandler(CommandGroupHandlerRef, this, AccessTarget);
3709  GDBMethodsAnchor();
3710  }
3711 #endif // __SYCL_DEVICE_ONLY__
3712 
3713  /* -- common interface members -- */
3714 
3715  unsampled_image_accessor(const unsampled_image_accessor &Rhs) = default;
3716 
3717  unsampled_image_accessor(unsampled_image_accessor &&Rhs) = default;
3718 
3719  unsampled_image_accessor &
3720  operator=(const unsampled_image_accessor &Rhs) = default;
3721 
3722  unsampled_image_accessor &operator=(unsampled_image_accessor &&Rhs) = default;
3723 
3724  ~unsampled_image_accessor() = default;
3725 
3726 #ifdef __SYCL_DEVICE_ONLY__
3727  bool operator==(const unsampled_image_accessor &Rhs) const;
3728 #else
3729  bool operator==(const unsampled_image_accessor &Rhs) const {
3730  return Rhs.impl == impl;
3731  }
3732 #endif // __SYCL_DEVICE_ONLY__
3733 
3734  bool operator!=(const unsampled_image_accessor &Rhs) const {
3735  return !(Rhs == *this);
3736  }
3737 
3738  /* -- property interface members -- */
3739 
3740  size_t size() const noexcept {
3741 #ifdef __SYCL_DEVICE_ONLY__
3742  // Currently not reachable on device.
3743  return 0;
3744 #else
3745  return host_base_class::getSize().size();
3746 #endif // __SYCL_DEVICE_ONLY__
3747  }
3748 
3749  /* Available only when: AccessMode == access_mode::read
3750  if Dimensions == 1, CoordT = int
3751  if Dimensions == 2, CoordT = int2
3752  if Dimensions == 3, CoordT = int4 */
3753  template <typename CoordT,
3754  typename = std::enable_if_t<AccessMode == access_mode::read &&
3756  Dimensions, CoordT>::value>>
3757  DataT read(const CoordT &Coords) const noexcept {
3758 #ifdef __SYCL_DEVICE_ONLY__
3759  // Currently not reachable on device.
3760  std::ignore = Coords;
3761  return {0, 0, 0, 0};
3762 #else
3763  return host_base_class::read<DataT>(Coords);
3764 #endif // __SYCL_DEVICE_ONLY__
3765  }
3766 
3767  /* Available only when: AccessMode == access_mode::write
3768  if Dimensions == 1, CoordT = int
3769  if Dimensions == 2, CoordT = int2
3770  if Dimensions == 3, CoordT = int4 */
3771  template <typename CoordT,
3772  typename = std::enable_if_t<AccessMode == access_mode::write &&
3774  Dimensions, CoordT>::value>>
3775  void write(const CoordT &Coords, const DataT &Color) const {
3776 #ifdef __SYCL_DEVICE_ONLY__
3777  // Currently not reachable on device.
3778  std::ignore = Coords;
3779  std::ignore = Color;
3780 #else
3781  host_base_class::write<DataT>(Coords, Color);
3782 #endif // __SYCL_DEVICE_ONLY__
3783  }
3784 
3785 private:
3787 #ifndef __SYCL_DEVICE_ONLY__
3788  : host_base_class{Impl}
3789 #endif // __SYCL_DEVICE_ONLY__
3790  {
3791  std::ignore = Impl;
3792  }
3793 
3794  template <class Obj>
3795  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
3796 
3797  template <class T>
3798  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
3799 };
3800 
3801 template <typename DataT, int Dimensions = 1,
3803  (std::is_const_v<DataT> ? access_mode::read
3806  : private detail::UnsampledImageAccessorBaseHost,
3807  public detail::OwnerLessBase<
3809  static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
3810  std::is_same_v<DataT, float4> ||
3811  std::is_same_v<DataT, half4>,
3812  "The data type of an image accessor must be only int4, "
3813  "uint4, float4 or half4 from SYCL namespace");
3814 
3815  using base_class = detail::UnsampledImageAccessorBaseHost;
3816 
3817 public:
3818  using value_type = typename std::conditional<AccessMode == access_mode::read,
3819  const DataT, DataT>::type;
3820  using reference = value_type &;
3821  using const_reference = const DataT &;
3822 
3823  template <typename AllocatorT>
3826  const property_list &PropList = {},
3828  : base_class(detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
3829  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
3830  Dimensions, ImageRef.getElementSize(),
3831  {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
3832  ImageRef.getChannelType(), ImageRef.getChannelOrder(),
3833  PropList) {
3834  addHostUnsampledImageAccessorAndWait(base_class::impl.get());
3835 
3837  detail::getSyclObjImpl(ImageRef).get(), this->impl.get(), std::nullopt,
3838  AccessMode, (const void *)typeid(DataT).name(), sizeof(DataT), CodeLoc);
3839  }
3840 
3841  /* -- common interface members -- */
3842 
3843  host_unsampled_image_accessor(const host_unsampled_image_accessor &Rhs) =
3844  default;
3845 
3846  host_unsampled_image_accessor(host_unsampled_image_accessor &&Rhs) = default;
3847 
3848  host_unsampled_image_accessor &
3849  operator=(const host_unsampled_image_accessor &Rhs) = default;
3850 
3851  host_unsampled_image_accessor &
3852  operator=(host_unsampled_image_accessor &&Rhs) = default;
3853 
3854  ~host_unsampled_image_accessor() = default;
3855 
3857  return Rhs.impl == impl;
3858  }
3860  return !(Rhs == *this);
3861  }
3862 
3863  /* -- property interface members -- */
3864 
3865  size_t size() const noexcept { return base_class::getSize().size(); }
3866 
3867  /* Available only when: (AccessMode == access_mode::read ||
3868  AccessMode == access_mode::read_write)
3869  if Dimensions == 1, CoordT = int
3870  if Dimensions == 2, CoordT = int2
3871  if Dimensions == 3, CoordT = int4 */
3872  template <
3873  typename CoordT,
3874  typename = std::enable_if_t<
3878  DataT read(const CoordT &Coords) const noexcept
3879 #ifdef __SYCL_DEVICE_ONLY__
3880  ;
3881 #else
3882  {
3883  // Host implementation is only available in host code. Device is not allowed
3884  // to use host_unsampled_image_accessor.
3885  return base_class::read<DataT>(Coords);
3886  }
3887 #endif
3888 
3889  /* Available only when: (AccessMode == access_mode::write ||
3890  AccessMode == access_mode::read_write)
3891  if Dimensions == 1, CoordT = int
3892  if Dimensions == 2, CoordT = int2
3893  if Dimensions == 3, CoordT = int4 */
3894  template <
3895  typename CoordT,
3896  typename = std::enable_if_t<
3900  void write(const CoordT &Coords, const DataT &Color) const
3901 #ifdef __SYCL_DEVICE_ONLY__
3902  ;
3903 #else
3904  {
3905  // Host implementation is only available in host code. Device is not allowed
3906  // to use host_unsampled_image_accessor.
3907  base_class::write<DataT>(Coords, Color);
3908  }
3909 #endif
3910 
3911 private:
3914  : base_class{Impl} {}
3915 
3916  template <class Obj>
3917  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
3918 
3919  template <class T>
3920  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
3921 };
3922 
3923 template <typename DataT, int Dimensions,
3924  image_target AccessTarget = image_target::device>
3926 #ifndef __SYCL_DEVICE_ONLY__
3928 #endif // __SYCL_DEVICE_ONLY__
3929  public detail::OwnerLessBase<
3930  sampled_image_accessor<DataT, Dimensions, AccessTarget>> {
3931  static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
3932  std::is_same_v<DataT, float4> ||
3933  std::is_same_v<DataT, half4>,
3934  "The data type of an image accessor must be only int4, "
3935  "uint4, float4 or half4 from SYCL namespace");
3936 
3937 #ifdef __SYCL_DEVICE_ONLY__
3938  char MPadding[sizeof(detail::SampledImageAccessorBaseHost)];
3939 #else
3941 #endif // __SYCL_DEVICE_ONLY__
3942 
3943 public:
3944  using value_type = const DataT;
3945  using reference = const DataT &;
3946  using const_reference = const DataT &;
3947 
3948  template <typename AllocatorT>
3951  handler &CommandGroupHandlerRef, const property_list &PropList = {},
3953 #ifdef __SYCL_DEVICE_ONLY__
3954  {}
3955 #else
3956  : host_base_class(detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
3957  detail::getSyclObjImpl(ImageRef).get(), Dimensions,
3958  ImageRef.getElementSize(),
3959  {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
3960  ImageRef.getChannelType(), ImageRef.getChannelOrder(),
3961  ImageRef.getSampler(), PropList) {
3962  device Device = detail::getDeviceFromHandler(CommandGroupHandlerRef);
3963  // Avoid aspect::image warning.
3964  aspect ImageAspect = aspect::image;
3965  if (AccessTarget == image_target::device && !Device.has(ImageAspect))
3966  throw sycl::exception(
3967  sycl::make_error_code(sycl::errc::feature_not_supported),
3968  "Device associated with command group handler does not have "
3969  "aspect::image.");
3970 
3972  detail::getSyclObjImpl(ImageRef).get(), this->impl.get(), AccessTarget,
3973  (const void *)typeid(DataT).name(), sizeof(DataT), CodeLoc);
3974  detail::associateWithHandler(CommandGroupHandlerRef, this, AccessTarget);
3975  GDBMethodsAnchor();
3976  }
3977 #endif // __SYCL_DEVICE_ONLY__
3978 
3979  /* -- common interface members -- */
3980 
3981  sampled_image_accessor(const sampled_image_accessor &Rhs) = default;
3982 
3983  sampled_image_accessor(sampled_image_accessor &&Rhs) = default;
3984 
3985  sampled_image_accessor &
3986  operator=(const sampled_image_accessor &Rhs) = default;
3987 
3988  sampled_image_accessor &operator=(sampled_image_accessor &&Rhs) = default;
3989 
3990  ~sampled_image_accessor() = default;
3991 
3992 #ifdef __SYCL_DEVICE_ONLY__
3993  bool operator==(const sampled_image_accessor &Rhs) const;
3994 #else
3995  bool operator==(const sampled_image_accessor &Rhs) const {
3996  return Rhs.impl == impl;
3997  }
3998 #endif // __SYCL_DEVICE_ONLY__
3999 
4000  bool operator!=(const sampled_image_accessor &Rhs) const {
4001  return !(Rhs == *this);
4002  }
4003 
4004  /* -- property interface members -- */
4005 
4006  size_t size() const noexcept {
4007 #ifdef __SYCL_DEVICE_ONLY__
4008  // Currently not reachable on device.
4009  return 0;
4010 #else
4011  return host_base_class::getSize().size();
4012 #endif // __SYCL_DEVICE_ONLY__
4013  }
4014 
4015  /* if Dimensions == 1, CoordT = float
4016  if Dimensions == 2, CoordT = float2
4017  if Dimensions == 3, CoordT = float4 */
4018  template <typename CoordT,
4019  typename = std::enable_if_t<detail::IsValidSampledCoord2020DataT<
4020  Dimensions, CoordT>::value>>
4021  DataT read(const CoordT &Coords) const noexcept {
4022 #ifdef __SYCL_DEVICE_ONLY__
4023  // Currently not reachable on device.
4024  std::ignore = Coords;
4025  return {0, 0, 0, 0};
4026 #else
4027  return host_base_class::read<DataT>(Coords);
4028 #endif // __SYCL_DEVICE_ONLY__
4029  }
4030 
4031 private:
4033 #ifndef __SYCL_DEVICE_ONLY__
4034  : host_base_class{Impl}
4035 #endif // __SYCL_DEVICE_ONLY__
4036  {
4037  std::ignore = Impl;
4038  }
4039 
4040  template <class Obj>
4041  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
4042 
4043  template <class T>
4044  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
4045 };
4046 
4047 template <typename DataT, int Dimensions>
4049  : private detail::SampledImageAccessorBaseHost,
4050  public detail::OwnerLessBase<
4051  host_sampled_image_accessor<DataT, Dimensions>> {
4052  static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
4053  std::is_same_v<DataT, float4> ||
4054  std::is_same_v<DataT, half4>,
4055  "The data type of an image accessor must be only int4, "
4056  "uint4, float4 or half4 from SYCL namespace");
4057 
4058  using base_class = detail::SampledImageAccessorBaseHost;
4059 
4060 public:
4061  using value_type = const DataT;
4062  using reference = const DataT &;
4063  using const_reference = const DataT &;
4064 
4065  template <typename AllocatorT>
4068  const property_list &PropList = {},
4070  : base_class(detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
4071  detail::getSyclObjImpl(ImageRef).get(), Dimensions,
4072  ImageRef.getElementSize(),
4073  {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
4074  ImageRef.getChannelType(), ImageRef.getChannelOrder(),
4075  ImageRef.getSampler(), PropList) {
4076  addHostSampledImageAccessorAndWait(base_class::impl.get());
4077 
4079  detail::getSyclObjImpl(ImageRef).get(), this->impl.get(), std::nullopt,
4080  (const void *)typeid(DataT).name(), sizeof(DataT), CodeLoc);
4081  }
4082 
4083  /* -- common interface members -- */
4084 
4085  host_sampled_image_accessor(const host_sampled_image_accessor &Rhs) = default;
4086 
4087  host_sampled_image_accessor(host_sampled_image_accessor &&Rhs) = default;
4088 
4089  host_sampled_image_accessor &
4090  operator=(const host_sampled_image_accessor &Rhs) = default;
4091 
4092  host_sampled_image_accessor &
4093  operator=(host_sampled_image_accessor &&Rhs) = default;
4094 
4095  ~host_sampled_image_accessor() = default;
4096 
4097  bool operator==(const host_sampled_image_accessor &Rhs) const {
4098  return Rhs.impl == impl;
4099  }
4100  bool operator!=(const host_sampled_image_accessor &Rhs) const {
4101  return !(Rhs == *this);
4102  }
4103 
4104  /* -- property interface members -- */
4105 
4106  size_t size() const noexcept { return base_class::getSize().size(); }
4107 
4108  /* if Dimensions == 1, CoordT = float
4109  if Dimensions == 2, CoordT = float2
4110  if Dimensions == 3, CoordT = float4 */
4111  template <typename CoordT,
4112  typename = std::enable_if_t<detail::IsValidSampledCoord2020DataT<
4113  Dimensions, CoordT>::value>>
4114  DataT read(const CoordT &Coords) const
4115 #ifdef __SYCL_DEVICE_ONLY__
4116  ;
4117 #else
4118  {
4119  // Host implementation is only available in host code. Device is not allowed
4120  // to use host_sampled_image_accessor.
4121  return base_class::read<DataT>(Coords);
4122  }
4123 #endif
4124 
4125 private:
4127  : base_class{Impl} {}
4128 
4129  template <class Obj>
4130  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
4131 
4132  template <class T>
4133  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
4134 };
4135 
4136 } // namespace _V1
4137 } // namespace sycl
4138 
4139 namespace std {
4140 template <typename DataT, int Dimensions, sycl::access::mode AccessMode,
4141  sycl::access::target AccessTarget,
4143 struct hash<sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
4144  IsPlaceholder>> {
4145  using AccType = sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
4147 
4148  size_t operator()(const AccType &A) const {
4149 #ifdef __SYCL_DEVICE_ONLY__
4150  // Hash is not supported on DEVICE. Just return 0 here.
4151  (void)A;
4152  return 0;
4153 #else
4154  // getSyclObjImpl() here returns a pointer to either AccessorImplHost
4155  // or LocalAccessorImplHost depending on the AccessTarget.
4156  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4157  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4158 #endif
4159  }
4160 };
4161 
4162 template <typename DataT, int Dimensions, sycl::access_mode AccessMode>
4163 struct hash<sycl::host_accessor<DataT, Dimensions, AccessMode>> {
4165 
4166  size_t operator()(const AccType &A) const {
4167 #ifdef __SYCL_DEVICE_ONLY__
4168  // Hash is not supported on DEVICE. Just return 0 here.
4169  (void)A;
4170  return 0;
4171 #else
4172  // getSyclObjImpl() here returns a pointer to AccessorImplHost.
4173  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4174  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4175 #endif
4176  }
4177 };
4178 
4179 template <typename DataT, int Dimensions>
4180 struct hash<sycl::local_accessor<DataT, Dimensions>> {
4182 
4183  size_t operator()(const AccType &A) const {
4184 #ifdef __SYCL_DEVICE_ONLY__
4185  // Hash is not supported on DEVICE. Just return 0 here.
4186  (void)A;
4187  return 0;
4188 #else
4189  // getSyclObjImpl() here returns a pointer to LocalAccessorImplHost.
4190  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4191  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4192 #endif
4193  }
4194 };
4195 
4196 template <typename DataT, int Dimensions, sycl::access_mode AccessMode,
4197  sycl::image_target AccessTarget>
4198 struct hash<sycl::unsampled_image_accessor<DataT, Dimensions, AccessMode,
4199  AccessTarget>> {
4201  AccessTarget>;
4202 
4203  size_t operator()(const AccType &A) const {
4204 #ifdef __SYCL_DEVICE_ONLY__
4205  // Hash is not supported on DEVICE. Just return 0 here.
4206  (void)A;
4207  return 0;
4208 #else
4209  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4210  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4211 #endif
4212  }
4213 };
4214 
4215 template <typename DataT, int Dimensions, sycl::access_mode AccessMode>
4216 struct hash<
4217  sycl::host_unsampled_image_accessor<DataT, Dimensions, AccessMode>> {
4218  using AccType =
4220 
4221  size_t operator()(const AccType &A) const {
4222  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4223  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4224  }
4225 };
4226 
4227 template <typename DataT, int Dimensions, sycl::image_target AccessTarget>
4228 struct hash<sycl::sampled_image_accessor<DataT, Dimensions, AccessTarget>> {
4230 
4231  size_t operator()(const AccType &A) const {
4232 #ifdef __SYCL_DEVICE_ONLY__
4233  // Hash is not supported on DEVICE. Just return 0 here.
4234  (void)A;
4235  return 0;
4236 #else
4237  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4238  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4239 #endif
4240  }
4241 };
4242 
4243 template <typename DataT, int Dimensions>
4244 struct hash<sycl::host_sampled_image_accessor<DataT, Dimensions>> {
4246 
4247  size_t operator()(const AccType &A) const {
4248  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4249  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4250  }
4251 };
4252 
4253 } // namespace std
sycl::_V1::IsPlaceholder
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS IsPlaceholder
Definition: accessor.hpp:3216
sycl::_V1::detail::LocalAccessorBaseDevice::Offset
id< Dims > Offset
Definition: accessor.hpp:497
sycl::_V1::property_list
Objects of the property_list class are containers for the SYCL properties.
Definition: property_list.hpp:31
sycl::_V1::detail::AccessorImplPtr
std::shared_ptr< AccessorImplHost > AccessorImplPtr
Definition: accessor.hpp:532
aspects.hpp
sycl::_V1::detail::AccPropBufferLocation
@ AccPropBufferLocation
Definition: property_helper.hpp:62
sycl::_V1::detail::InitializedVal
Definition: common.hpp:258
sycl::_V1::detail::associateWithHandler
void associateWithHandler(handler &, AccessorBaseHost *, access::target)
Definition: handler_proxy.cpp:17
sycl::_V1::detail::LocalAccessorBaseHost::getPtr
void * getPtr()
Definition: accessor.cpp:117
sycl::_V1::ext::oneapi::experimental::reference
sycl::ext::oneapi::experimental::annotated_ref< T, typename unpack< filtered_properties >::type > reference
Definition: annotated_ptr.hpp:174
sycl::_V1::operator=
multi_ptr & operator=(const multi_ptr &)=default
sycl::_V1::detail::LocalAccessorImplHost
Definition: accessor_impl.hpp:154
sycl::_V1::detail::SampledImageAccessorBaseHost::SampledImageAccessorBaseHost
SampledImageAccessorBaseHost(const SampledImageAccessorImplPtr &Impl)
Definition: accessor.hpp:710
sycl::_V1::host_unsampled_image_accessor::const_reference
const DataT & const_reference
Definition: accessor.hpp:3821
sycl::_V1::detail::UnsampledImageAccessorBaseHost::impl
UnsampledImageAccessorImplPtr impl
Definition: accessor.hpp:670
syclcompat::detail::pointer_access_attribute::end
@ end
property_list.hpp
sycl::_V1::detail::image_accessor
Definition: accessor.hpp:811
sycl::_V1::access::mode::discard_read_write
@ discard_read_write
sycl::_V1::host_accessor::host_accessor
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:3520
sycl::_V1::local_accessor_base::operator==
bool operator==(const local_accessor_base &Rhs) const
Definition: accessor.hpp:2978
sycl::_V1::image
Defines a shared image data.
Definition: image.hpp:49
sycl::_V1::detail::AccessorBaseHost::getOffset
id< 3 > & getOffset()
Definition: accessor.cpp:75
sycl::_V1::host_accessor
host_accessor(buffer< DataT, Dimensions, AllocatorT >) -> host_accessor< DataT, Dimensions, access::mode::read_write >
sycl::_V1::access::mode
mode
Definition: access.hpp:34
pi.h
sycl::_V1::host_accessor::host_accessor
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:3530
sycl::_V1::host_accessor::host_accessor
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:3484
sycl::_V1::detail::UnsampledImageAccessorBaseHost::read
DataT read(const CoordT &Coords) const noexcept
Definition: accessor.hpp:688
sycl::_V1::image_sampler
Definition: sampler.hpp:132
sycl::_V1::ext::intel::experimental::operator==
bool operator==(const cache_config &lhs, const cache_config &rhs)
Definition: kernel_execution_properties.hpp:36
sycl::_V1::detail::UnsampledImageAccessorBaseHost::write
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:701
sycl::_V1::host_accessor::host_accessor
host_accessor(const detail::AccessorImplPtr &Impl)
Definition: accessor.hpp:3396
sycl::_V1::mode_tag_t
Definition: access.hpp:76
type_traits.hpp
property_list_base.hpp
sycl::_V1::make_error_code
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:94
sycl::_V1::host_sampled_image_accessor::operator!=
bool operator!=(const host_sampled_image_accessor &Rhs) const
Definition: accessor.hpp:4100
sycl::_V1::host_unsampled_image_accessor
Definition: accessor.hpp:3805
sycl::_V1::local_accessor_base::operator!=
bool operator!=(const local_accessor_base &Rhs) const
Definition: accessor.hpp:2981
syclcompat::detail::get_offset
static size_t get_offset(sycl::id< 3 > id, size_t slice, size_t pitch)
Definition: memory.hpp:291
sycl::_V1::access::placeholder::false_t
@ false_t
sycl::_V1::accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::accessor
accessor(sycl::image< Dimensions, AllocatorT > &Image, const property_list &propList)
Definition: accessor.hpp:3294
sycl::_V1::host_accessor::host_accessor
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:3452
sycl::_V1::detail::accessor_common::AccessorSubscript::operator[]
std::enable_if_t< CurDims==1 &&IsAccessAtomic, atomic< DataT, AS > > operator[](size_t Index) const
Definition: accessor.hpp:395
sycl::_V1::local_accessor_base::value_type
DataT value_type
Definition: accessor.hpp:2844
sycl::_V1::errc::feature_not_supported
@ feature_not_supported
sycl::_V1::local_accessor_base::getSize
const range< 3 > & getSize() const
Definition: accessor.hpp:2809
sycl::_V1::detail::accessor_common::AccessorSubscript::AccessorSubscript
AccessorSubscript(AccType Accessor, id< Dims > IDs)
Definition: accessor.hpp:370
sycl::_V1::buffer
Defines a shared array that can be used by kernels in queues.
Definition: buffer.hpp:60
sycl::_V1::local_accessor_base::local_accessor_base
local_accessor_base(const detail::LocalAccessorImplPtr &Impl)
Definition: accessor.hpp:2796
sycl::_V1::host_accessor::host_accessor
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:3511
sycl::_V1::detail::image_accessor::read
DataT read(const CoordT &Coords, const sampler &Smpl) const
Definition: accessor.hpp:1077
sycl::_V1::local_accessor_base::const_reference
const DataT & const_reference
Definition: accessor.hpp:2846
sycl::_V1::local_accessor_base::get_size
size_t get_size() const
Definition: accessor.hpp:2912
types.hpp
sycl::_V1::detail::image_common< Dimensions, sycl::image_allocator >::get_range
range< Dimensions > get_range() const
Definition: image.hpp:324
sycl::_V1::get
pointer_t get() const
Definition: multi_ptr.hpp:974
sycl::_V1::device::has
bool has(aspect Aspect) const __SYCL_WARN_IMAGE_ASPECT(Aspect)
Indicates if the SYCL device has the given feature.
Definition: device.cpp:209
handler_proxy.hpp
sycl::_V1::local_accessor_base::reference
DataT & reference
Definition: accessor.hpp:2845
sycl::_V1::detail::image_accessor::write
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:1102
sycl::_V1::detail::accessor_common::IsAccessAnyWrite
constexpr static bool IsAccessAnyWrite
Definition: accessor.hpp:339
accessor_property_list.hpp
sycl::_V1::local_accessor_base::getPtr
void * getPtr() const
Definition: accessor.hpp:2808
sycl::_V1::detail::LocalAccessorBaseHost
Definition: accessor.hpp:603
sycl::_V1::image::get_range
range< Dimensions > get_range() const
Definition: image.hpp:630
sycl::_V1::detail::SampledImageAccessorBaseHost
Definition: accessor.hpp:708
common.hpp
sycl::_V1::detail::LocalAccessorBaseHost::getSize
sycl::range< 3 > & getSize()
Definition: accessor.cpp:113
sycl::_V1::detail::accessor_common::IsAccessReadOnly
constexpr static bool IsAccessReadOnly
Definition: accessor.hpp:345
std::hash< sycl::accessor< DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:4148
sycl::_V1::detail::image_plain::getSlicePitch
size_t getSlicePitch() const
Definition: image.cpp:191
sycl::_V1::detail::AccHostDataT::AccHostDataT
AccHostDataT(const sycl::id< 3 > &Offset, const sycl::range< 3 > &Range, const sycl::range< 3 > &MemoryRange, void *Data=nullptr)
Definition: accessor.hpp:253
sycl::_V1::detail::accessor_common::AccessorSubscript
Definition: accessor.hpp:363
sycl::_V1::host_accessor::host_accessor
host_accessor(const host_accessor< DataT_, Dimensions, AccessMode_ > &other)
Definition: accessor.hpp:3586
sycl::_V1::detail::IsPropertyListT
typename std::is_base_of< PropertyListBase, T > IsPropertyListT
Definition: accessor.hpp:281
sycl::_V1::detail::image_accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::reference
DataT & reference
Definition: accessor.hpp:920
sycl::_V1::detail::unsampledImageConstructorNotification
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)
Definition: accessor_impl.cpp:58
sycl::_V1::detail::accessor_common::AccessorSubscript::operator[]
auto operator[](size_t Index)
Definition: accessor.hpp:380
sycl::_V1::detail::imageWriteHostImpl
void imageWriteHostImpl(const CoordT &Coords, const WriteDataT &Color, id< 3 > ImgPitch, uint8_t ElementSize, image_channel_type ImgChannelType, image_channel_order ImgChannelOrder, void *BasePtr)
Definition: image_accessor_util.hpp:701
sycl::_V1::detail::LocalAccessorBaseDevice
Definition: accessor.hpp:488
sycl::_V1::access::target::host_image
@ host_image
sycl::_V1::detail::accessor_common::IsAccessAtomic
constexpr static bool IsAccessAtomic
Definition: accessor.hpp:351
detail
---— Error handling, matching OpenCL plugin semantics.
Definition: common.hpp:44
sycl::_V1::local_accessor_base::ConcreteASPtrType
typename detail::DecoratedType< DataT, AS >::type * ConcreteASPtrType
Definition: accessor.hpp:2753
sycl::_V1::detail::AccHostDataT::MOffset
sycl::id< 3 > MOffset
Definition: accessor.hpp:258
sycl::_V1::detail::constructorNotification
void constructorNotification(void *BufferObj, void *AccessorObj, access::target Target, access::mode Mode, const code_location &CodeLoc)
sycl::_V1::image_target::device
@ device
sycl::_V1::access::mode::atomic
@ atomic
accessor_iterator.hpp
sycl::_V1::host_accessor
Definition: accessor.hpp:3372
sycl::_V1::detail::UnsampledImageAccessorBaseHost::GDBMethodsAnchor
void GDBMethodsAnchor()
Definition: accessor.hpp:673
sycl::_V1::__SYCL_TYPE
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.
sycl::_V1::local_accessor_base::local_accessor_base
local_accessor_base(range< Dimensions > AllocationSize, handler &, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2879
sycl::_V1::host_accessor::host_accessor
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:3501
sycl::_V1::Dimensions
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS Dimensions
Definition: accessor.hpp:3215
sycl::_V1::detail::image_plain::getChannelType
image_channel_type getChannelType() const
Definition: image.cpp:201
sycl::_V1::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::get_pointer
local_ptr< DataT > get_pointer() const
Definition: accessor.hpp:3007
helpers.hpp
sycl
Definition: access.hpp:18
sycl::_V1::accessor
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.
sycl::_V1::mode_target_tag_t
Definition: access.hpp:80
sycl::_V1::detail::accessor_common::AccessorSubscript::AccessorSubscript
AccessorSubscript(AccType Accessor, size_t Index)
Definition: accessor.hpp:375
sycl::_V1::local_accessor_base::RefType
detail::const_if_const_AS< AS, DataT > & RefType
Definition: accessor.hpp:2755
sycl::_V1::detail::__image_array_slice__::__image_array_slice__
__image_array_slice__(accessor< DataT, Dimensions, AccessMode, access::target::image_array, IsPlaceholder, ext::oneapi::accessor_property_list<>> BaseAcc, size_t Idx)
Definition: accessor.hpp:1145
max
simd< _Tp, _Abi > max(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
sycl::_V1::detail::image_accessor::__SYCL2020_DEPRECATED
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:1037
__SYCL_SPECIAL_CLASS
#define __SYCL_SPECIAL_CLASS
Definition: defines.hpp:29
sycl::_V1::detail::IsCxPropertyList::value
constexpr static bool value
Definition: accessor.hpp:288
sycl::_V1::detail::image_plain::getChannelOrder
image_channel_order getChannelOrder() const
Definition: image.cpp:197
access.hpp
sycl::_V1::detail::accessor_common
Definition: accessor.hpp:312
sycl::_V1::detail::image_plain::getSampler
image_sampler getSampler() const noexcept
Definition: image.cpp:193
sycl::_V1::local_accessor_base::local_accessor_base
local_accessor_base(range< Dimensions > AllocationSize, handler &, const property_list &propList, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2895
sycl::_V1::id
A unique identifier of an item in an index space.
Definition: array.hpp:20
sycl::_V1::detail::LocalAccessorBaseDevice::AccessRange
range< Dims > AccessRange
Definition: accessor.hpp:495
sycl::_V1::detail::addHostAccessorAndWait
void addHostAccessorAndWait(AccessorImplHost *Req)
Definition: accessor_impl.cpp:36
id.hpp
sycl::_V1::detail::IsRunTimePropertyListT
typename std::is_same< ext::oneapi::accessor_property_list<>, T > IsRunTimePropertyListT
Definition: accessor.hpp:285
sycl::_V1::local_accessor_base
Local accessor.
Definition: accessor.hpp:2726
sycl::_V1::detail::__image_array_slice__::size
size_t size() const noexcept
Definition: accessor.hpp:1193
owner_less_base.hpp
sycl::_V1::detail::SampledImageAccessorImplHost
Definition: accessor_impl.hpp:189
std::hash< sycl::unsampled_image_accessor< DataT, Dimensions, AccessMode, AccessTarget > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:4203
sycl::_V1::detail::accessModeFromConstness
constexpr access::mode accessModeFromConstness()
Definition: accessor.hpp:402
sycl::_V1::detail::AccessorImplDevice::AccessorImplDevice
AccessorImplDevice(id< Dims > Offset, range< Dims > AccessRange, range< Dims > MemoryRange)
Definition: accessor.hpp:512
sycl::_V1::host_accessor::IsSameAsBuffer
Definition: accessor.hpp:3383
sycl::_V1::host_sampled_image_accessor::size
size_t size() const noexcept
Definition: accessor.hpp:4106
sycl::_V1::detail::accessor_common< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::PtrType
detail::const_if_const_AS< AS, DataT > * PtrType
Definition: accessor.hpp:355
sycl::_V1::coordinate_normalization_mode::unnormalized
@ unnormalized
sycl::_V1::detail::__image_array_slice__::read
DataT read(const CoordT &Coords, const sampler &Smpl) const
Definition: accessor.hpp:1162
sycl::_V1::local_accessor_base::get_range
range< Dims > get_range() const
Definition: accessor.hpp:2919
image_accessor_util.hpp
sycl::_V1::detail::SampledImageAccessorBaseHost::GDBMethodsAnchor
void GDBMethodsAnchor()
Definition: accessor.hpp:743
sycl::_V1::detail::IsCxPropertyList
Definition: accessor.hpp:287
sycl::_V1::detail::LocalAccessorBaseDevice::LocalAccessorBaseDevice
LocalAccessorBaseDevice(sycl::range< Dims > Size)
Definition: accessor.hpp:490
sycl::_V1::host_sampled_image_accessor::host_sampled_image_accessor
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:4066
sycl::_V1::range
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: buffer.hpp:51
std::hash< sycl::host_sampled_image_accessor< DataT, Dimensions > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:4247
sycl::_V1::local_accessor_base::local_accessor_base
local_accessor_base()
Definition: accessor.hpp:2791
export.hpp
sycl::_V1::host_accessor::host_accessor
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:3468
sycl::_V1::image_accessor
Definition: accessor_properties.hpp:109
sycl::_V1::access::placeholder
placeholder
Definition: access.hpp:49
sycl::_V1::host_sampled_image_accessor
Definition: accessor.hpp:4048
sycl::_V1::sampled_image_accessor::sampled_image_accessor
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:3949
sycl::_V1::multi_ptr
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: atomic.hpp:34
sycl::_V1::host_accessor::operator=
const host_accessor & operator=(typename AccessorT::value_type &&Other) const
Definition: accessor.hpp:3560
sycl::_V1::ext::oneapi::experimental::has_property
static constexpr bool has_property()
Definition: annotated_arg.hpp:169
sycl::_V1::host_unsampled_image_accessor::operator==
bool operator==(const host_unsampled_image_accessor &Rhs) const
Definition: accessor.hpp:3856
defines_elementary.hpp
sycl::_V1::detail::image_accessor::operator==
bool operator==(const image_accessor &Rhs) const
Definition: accessor.hpp:996
sycl::_V1::unsampled_image_accessor::size
size_t size() const noexcept
Definition: accessor.hpp:3740
sycl::_V1::local_accessor_base::getPtr
void * getPtr()
Definition: accessor.hpp:2807
sycl::_V1::local_accessor_base::__SYCL2020_DEPRECATED
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:2914
sycl::_V1::unsampled_image_accessor::write
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:3775
sycl::_V1::access::placeholder::true_t
@ true_t
sycl::_V1::ext::oneapi::detail::getSyclWeakObjImpl
decltype(weak_object_base< SYCLObjT >::MObjWeakPtr) getSyclWeakObjImpl(const weak_object_base< SYCLObjT > &WeakObj)
Definition: weak_object_base.hpp:23
sycl::_V1::detail::deduceAccessMode
constexpr access::mode deduceAccessMode()
Definition: accessor.hpp:410
std::hash< sycl::host_accessor< DataT, Dimensions, AccessMode > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:4166
std::hash< sycl::sampled_image_accessor< DataT, Dimensions, AccessTarget > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:4231
sycl::_V1::detail::image_accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::const_reference
const DataT & const_reference
Definition: accessor.hpp:921
sycl::_V1::sampled_image_accessor::value_type
const DataT value_type
Definition: accessor.hpp:3944
sycl::_V1::access::target::host_task
@ host_task
generic_type_traits.hpp
sycl::_V1::ext::oneapi::property::no_offset
Definition: accessor_properties.hpp:61
range.hpp
sycl::_V1::detail::SampledImageAccessorBaseHost::impl
SampledImageAccessorImplPtr impl
Definition: accessor.hpp:740
sycl::_V1::host_unsampled_image_accessor::write
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:3900
sycl::_V1::access::target::image
@ image
sycl::_V1::host_unsampled_image_accessor::size
size_t size() const noexcept
Definition: accessor.hpp:3865
sycl::_V1::detail::UnsampledImageAccessorImplPtr
std::shared_ptr< UnsampledImageAccessorImplHost > UnsampledImageAccessorImplPtr
Definition: accessor.hpp:631
sycl::_V1::detail::OwnerLessBase
Definition: owner_less_base.hpp:20
sycl::_V1::global_ptr
multi_ptr< ElementType, access::address_space::global_space, IsDecorated > global_ptr
Definition: pointers.hpp:30
multi_ptr.hpp
sycl::_V1::detail::SampledImageAccessorImplPtr
std::shared_ptr< SampledImageAccessorImplHost > SampledImageAccessorImplPtr
Definition: accessor.hpp:633
sycl::_V1::host_sampled_image_accessor::const_reference
const DataT & const_reference
Definition: accessor.hpp:4063
defines.hpp
sycl::_V1::unsampled_image_accessor::unsampled_image_accessor
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:3683
sycl::_V1::detail::AccHostDataT::Reserved
void * Reserved
Definition: accessor.hpp:262
sycl::_V1::handler
Command group handler class.
Definition: handler.hpp:454
sycl::_V1::detail::__image_array_slice__
Definition: accessor.hpp:806
sycl::_V1::exception
Definition: exception.hpp:68
sycl::_V1::detail::accessor_common::AS
constexpr static access::address_space AS
Definition: accessor.hpp:314
sycl::_V1::access_mode
access::mode access_mode
Definition: access.hpp:72
sycl::_V1::ext::oneapi::detail::weak_object_base
Definition: weak_object_base.hpp:18
std::hash< sycl::local_accessor< DataT, Dimensions > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:4183
sycl::_V1::sampled_image_accessor::size
size_t size() const noexcept
Definition: accessor.hpp:4006
sycl::_V1::detail::accessor_common::IsAccessReadWrite
constexpr static bool IsAccessReadWrite
Definition: accessor.hpp:348
sycl::_V1::detail::AccHostDataT::MData
void * MData
Definition: accessor.hpp:261
sycl::_V1::detail::TargetToAS
Definition: access.hpp:140
sycl::_V1::local_accessor_base::size
size_t size() const noexcept
Definition: accessor.hpp:2916
sycl::_V1::image_channel_order
image_channel_order
Definition: image.hpp:57
sycl::_V1::accessor
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... >>
sycl::_V1::detail::SampledImageAccessorBaseHost::read
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:758
sycl::_V1::host_accessor::__init
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:3388
sycl::_V1::device
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:59
device.hpp
sycl::_V1::detail::createSyclObjFromImpl
T createSyclObjFromImpl(decltype(T::impl) ImplObj)
Definition: impl_utils.hpp:48
sycl::_V1::detail::AccessorBaseHost::isPlaceholder
bool isPlaceholder() const
Definition: accessor.cpp:101
sycl::_V1::host_accessor::host_accessor
host_accessor()
Definition: accessor.hpp:3408
sycl::_V1::detail::__image_array_slice__::get_range
range< Dims > get_range() const
Definition: accessor.hpp:1199
sycl::_V1::ext::oneapi::accessor_property_list
Definition: property_list.hpp:25
sycl::_V1::detail::image_plain::getElementSize
size_t getElementSize() const
Definition: image.cpp:187
sycl::_V1::local_accessor_base::local_accessor_base
local_accessor_base(handler &, const property_list &propList, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2862
sycl::_V1::detail::AccessorBaseHost::impl
AccessorImplPtr impl
Definition: accessor.hpp:594
sycl::_V1::detail::UnsampledImageAccessorBaseHost
Definition: accessor.hpp:640
sycl::_V1::host_unsampled_image_accessor::read
DataT read(const CoordT &Coords) const noexcept
Definition: accessor.hpp:3878
sycl::_V1::local_accessor_base::AccessorSubscript
typename AccessorCommonT::template AccessorSubscript< Dims, local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder > > AccessorSubscript
Definition: accessor.hpp:2751
sycl::_V1::access::target
target
Definition: access.hpp:22
sycl::_V1::access::target::device
@ device
sycl::_V1::read_write
constexpr mode_tag_t< access_mode::read_write > read_write
Definition: access.hpp:85
sycl::_V1::detail::IsValidSampledCoord2020DataT
Definition: accessor.hpp:793
sycl::_V1::operator[]
ReturnRef operator[](difference_type index)
Definition: multi_ptr.hpp:840
sycl::_V1::host_sampled_image_accessor::value_type
const DataT value_type
Definition: accessor.hpp:4061
sycl::_V1::detail::addHostSampledImageAccessorAndWait
void addHostSampledImageAccessorAndWait(SampledImageAccessorImplHost *Req)
Definition: accessor_impl.cpp:46
sycl::_V1::unsampled_image_accessor::const_reference
const DataT & const_reference
Definition: accessor.hpp:3680
sycl::_V1::unsampled_image_accessor
Definition: accessor.hpp:3655
sycl::_V1::errc::invalid
@ invalid
sycl::_V1::ext::oneapi::experimental::get_property
static constexpr auto get_property()
Definition: annotated_arg.hpp:173
sycl::_V1::detail::sampledImageConstructorNotification
void sampledImageConstructorNotification(void *ImageObj, void *AccessorObj, const std::optional< image_target > &Target, const void *Type, uint32_t ElemSize, const code_location &CodeLoc)
Definition: accessor_impl.cpp:71
sycl::_V1::detail::image_plain::getRowPitch
size_t getRowPitch() const
Definition: image.cpp:189
sycl::_V1::detail::AccessorImplDevice
Definition: accessor.hpp:509
image_ocl_types.hpp
sycl::_V1::sampled_image_accessor::read
DataT read(const CoordT &Coords) const noexcept
Definition: accessor.hpp:4021
sycl::_V1::accessor
Definition: accessor.hpp:246
sycl::_V1::host_accessor::host_accessor
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:3540
sycl::_V1::detail::IsValidCoordDataT
Definition: accessor.hpp:766
sycl::_V1::detail::LocalAccessorBaseHost::LocalAccessorBaseHost
LocalAccessorBaseHost(const LocalAccessorImplPtr &Impl)
Definition: accessor.hpp:605
sycl::_V1::detail::LocalAccessorImplPtr
std::shared_ptr< LocalAccessorImplHost > LocalAccessorImplPtr
Definition: accessor.hpp:601
sycl::_V1::local_accessor_base::GDBMethodsAnchor
void GDBMethodsAnchor()
Definition: accessor.hpp:2815
property_helper.hpp
sycl::_V1::detail::AccessorImplDevice::MemRange
range< Dims > MemRange
Definition: accessor.hpp:518
sampler.hpp
sycl::_V1::detail::code_location::current
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
sycl::_V1::host_accessor::host_accessor
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3444
sycl::_V1::difference_type
std::ptrdiff_t difference_type
Definition: multi_ptr.hpp:753
std::hash< sycl::host_unsampled_image_accessor< DataT, Dimensions, AccessMode > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:4221
sycl::_V1::detail::AccessorBaseHost::getElemSize
unsigned int getElemSize() const
Definition: accessor.cpp:86
sycl::_V1::detail::addHostUnsampledImageAccessorAndWait
void addHostUnsampledImageAccessorAndWait(UnsampledImageAccessorImplHost *Req)
Definition: accessor_impl.cpp:42
sycl::_V1::unsampled_image_accessor::reference
value_type & reference
Definition: accessor.hpp:3679
sycl::_V1::detail::type_list
Definition: type_list.hpp:24
sycl::_V1::detail::accessor_common< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::ConstRefType
const DataT & ConstRefType
Definition: accessor.hpp:354
sycl::_V1::local_accessor_base::operator[]
RefType operator[](id< Dimensions > Index) const
Definition: accessor.hpp:2933
sycl::_V1::host_sampled_image_accessor::read
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:4114
sycl::_V1::local_ptr
multi_ptr< ElementType, access::address_space::local_space, IsDecorated > local_ptr
Definition: pointers.hpp:52
sycl::_V1::detail::TryToGetElementType::type
decltype(check(T())) type
Definition: generic_type_traits.hpp:357
sycl::_V1::detail::UnsampledImageAccessorImplHost
Definition: accessor_impl.hpp:173
sycl::_V1::detail::AccHostDataT
Definition: accessor.hpp:252
sycl::_V1::detail::AccessorImplDevice::AccessorImplDevice
AccessorImplDevice()=default
sycl::_V1::detail::AccessorImplDevice::operator==
bool operator==(const AccessorImplDevice &Rhs) const
Definition: accessor.hpp:520
sycl::_V1::host_accessor::host_accessor
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:3460
sycl::_V1::detail::code_location
Definition: common.hpp:66
sycl::_V1::detail::deduceAccessTarget
constexpr access::target deduceAccessTarget(access::target defaultTarget)
Definition: accessor.hpp:457
sycl::_V1::detail::AccHostDataT::MAccessRange
sycl::range< 3 > MAccessRange
Definition: accessor.hpp:259
sycl::_V1::detail::AccHostDataT::MMemoryRange
sycl::range< 3 > MMemoryRange
Definition: accessor.hpp:260
sycl::_V1::host_unsampled_image_accessor::operator!=
bool operator!=(const host_unsampled_image_accessor &Rhs) const
Definition: accessor.hpp:3859
sycl::_V1::ext::intel::experimental::buffer_location
constexpr buffer_location_key::value_t< N > buffer_location
Definition: fpga_annotated_properties.hpp:94
sycl::_V1::detail::__image_array_slice__::read
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:1155
sycl::_V1::accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::accessor
accessor(sycl::image< Dimensions, AllocatorT > &Image)
Definition: accessor.hpp:3288
sycl::_V1::detail::image_accessor::image_accessor
image_accessor(image< Dims, AllocatorT > &ImageRef, int ImageElementSize)
Definition: accessor.hpp:936
sycl::_V1::host_accessor::ext_oneapi_owner_before
bool ext_oneapi_owner_before(const ext::oneapi::detail::weak_object_base< host_accessor > &Other) const noexcept
Definition: accessor.hpp:3600
sycl::_V1::local_accessor_base::PtrType
detail::const_if_const_AS< AS, DataT > * PtrType
Definition: accessor.hpp:2756
sycl::_V1::vec
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
Definition: aliases.hpp:18
sycl::_V1::detail::LocalAccessorBaseDevice::operator==
bool operator==(const LocalAccessorBaseDevice &Rhs) const
Definition: accessor.hpp:499
sycl::_V1::detail::AccessorBaseHost::getMemoryRange
range< 3 > & getMemoryRange()
Definition: accessor.cpp:77
exception.hpp
sycl::_V1::detail::__image_array_slice__::__SYCL2020_DEPRECATED
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:1191
sycl::_V1::local_accessor
Definition: multi_ptr.hpp:73
std
Definition: accessor.hpp:4139
sycl::_V1::host_accessor::ext_oneapi_owner_before
bool ext_oneapi_owner_before(const host_accessor &Other) const noexcept
Definition: accessor.hpp:3607
sycl::_V1::detail::accessor_common< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::RefType
detail::const_if_const_AS< AS, DataT > & RefType
Definition: accessor.hpp:353
sycl::_V1::detail::getDeviceFromHandler
device getDeviceFromHandler(handler &CommandGroupHandlerRef)
Definition: accessor.cpp:16
sycl::_V1::detail::is_contained
Definition: type_list.hpp:56
sycl::_V1::local_accessor_base::local_accessor_base
local_accessor_base(handler &, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2849
sycl::_V1::detail::image_accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::value_type
DataT value_type
Definition: accessor.hpp:919
sycl::_V1::detail::accessor_common::IsPlaceH
constexpr static bool IsPlaceH
Definition: accessor.hpp:328
sycl::_V1::local_accessor_base::getQualifiedPtr
PtrType getQualifiedPtr() const
Definition: accessor.hpp:2803
image.hpp
sycl::_V1::detail::image_accessor::get_range
range< Dims > get_range() const
Definition: accessor.hpp:1042
aliases.hpp
sycl::_V1::local_accessor_base::getLinearIndex
size_t getLinearIndex(id< Dims > Id) const
Definition: accessor.hpp:2828
sycl::_V1::detail::accessor_common::IsGlobalBuf
constexpr static bool IsGlobalBuf
Definition: accessor.hpp:333
sycl::_V1::detail::AccessorImplDevice::AccessRange
range< Dims > AccessRange
Definition: accessor.hpp:517
sycl::_V1::host_accessor::host_accessor
host_accessor(const host_accessor< DataT_, Dimensions, AccessMode > &other)
Definition: accessor.hpp:3571
buffer.hpp
sycl::_V1::host_unsampled_image_accessor::host_unsampled_image_accessor
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:3824
sycl::_V1::sampled_image_accessor::operator!=
bool operator!=(const sampled_image_accessor &Rhs) const
Definition: accessor.hpp:4000
sycl::_V1::sampled_image_accessor::const_reference
const DataT & const_reference
Definition: accessor.hpp:3946
sycl::_V1::detail::accessor_common::IsHostBuf
constexpr static bool IsHostBuf
Definition: accessor.hpp:316
sycl::_V1::detail::__SYCL2020_DEPRECATED
struct __SYCL2020_DEPRECATED("This type isn't device copyable in SYCL 2020") IsDeprecatedDeviceCopyable< T
sycl::_V1::host_unsampled_image_accessor::reference
value_type & reference
Definition: accessor.hpp:3820
sycl::_V1::detail::image_accessor::read
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:1058
sycl::_V1::filtering_mode::nearest
@ nearest
sycl::_V1::detail::IsValidUnsampledCoord2020DataT
Definition: accessor.hpp:782
sycl::_V1::detail::getLinearIndex
size_t getLinearIndex(const T< Dims > &Index, const U< Dims > &Range)
Definition: common.hpp:367
sycl::_V1::detail::GetZeroDimAccessRange
sycl::range< 1 > GetZeroDimAccessRange(BufferT Buffer)
Definition: accessor.hpp:303
sycl::_V1::addressing_mode::none
@ none
sycl::_V1::detail::image_accessor::image_accessor
image_accessor(image< Dims, AllocatorT > &ImageRef, handler &CommandGroupHandlerRef, int ImageElementSize)
Definition: accessor.hpp:963
sycl::_V1::detail::DecoratedType
Definition: access.hpp:163
sycl::_V1::local_accessor_base::operator[]
std::enable_if_t< Dims==1 &&AccessMode==access::mode::atomic, atomic< DataT, AS > > operator[](size_t Index) const
Definition: accessor.hpp:2965
sycl::_V1::host_accessor::host_accessor
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:3492
sycl::_V1::unsampled_image_accessor::read
DataT read(const CoordT &Coords) const noexcept
Definition: accessor.hpp:3757
sycl::_V1::detail::accessor_common::IsConst
static constexpr bool IsConst
Definition: accessor.hpp:346
sycl::_V1::detail::accessor_common::IsConstantBuf
constexpr static bool IsConstantBuf
Definition: accessor.hpp:336
sycl::_V1::ext::intel::experimental::operator!=
bool operator!=(const cache_config &lhs, const cache_config &rhs)
Definition: kernel_execution_properties.hpp:40
sycl::_V1::unsampled_image_accessor::operator!=
bool operator!=(const unsampled_image_accessor &Rhs) const
Definition: accessor.hpp:3734
sycl::_V1::local_accessor_base::getSize
range< 3 > & getSize()
Definition: accessor.hpp:2812
sycl::_V1::image_channel_type
image_channel_type
Definition: image.hpp:75
sycl::_V1::sampled_image_accessor
Definition: accessor.hpp:3925
sycl::_V1::AccessMode
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS AccessMode
Definition: accessor.hpp:3215
sycl::_V1::access::mode::read_write
@ read_write
weak_object_base.hpp
sycl::_V1::detail::image_accessor::image_accessor
image_accessor(const AccessorImplPtr &Impl)
Definition: accessor.hpp:908
pointers.hpp
sycl::_V1::detail::get
Definition: tuple.hpp:58
sycl::_V1::detail::image_accessor::operator!=
bool operator!=(const image_accessor &Rhs) const
Definition: accessor.hpp:1004
sycl::_V1::detail::AccessorImplDevice::Offset
id< Dims > Offset
Definition: accessor.hpp:516
sycl::_V1::access::mode::write
@ write
sycl::_V1::detail::__image_array_slice__::write
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:1169
sycl::_V1::detail::AccessorBaseHost::isMemoryObjectUsedByGraph
bool isMemoryObjectUsedByGraph() const
Definition: accessor.cpp:103
PropertyListT
sycl::_V1::detail::AccessorBaseHost::getAccessRange
range< 3 > & getAccessRange()
Definition: accessor.cpp:76
sycl::_V1::host_accessor::operator=
const host_accessor & operator=(const typename AccessorT::value_type &Other) const
Definition: accessor.hpp:3552
sycl::_V1::detail::AccessorBaseHost::getPtr
void * getPtr() noexcept
Definition: accessor.cpp:78
sycl::_V1::sampled_image_accessor::operator==
bool operator==(const sampled_image_accessor &Rhs) const
Definition: accessor.hpp:3995
sycl::_V1::sampled_image
Definition: image.hpp:1009
sycl::_V1::detail::LocalAccessorBaseHost::impl
LocalAccessorImplPtr impl
Definition: accessor.hpp:625
atomic.hpp
sycl::_V1::access::mode::read
@ read
sycl::_V1::detail::AccessorBaseHost::AccessorBaseHost
AccessorBaseHost(const AccessorImplPtr &Impl)
Definition: accessor.hpp:536
sycl::_V1::local_accessor_base::operator[]
AccessorCommonT::template AccessorSubscript< Dims - 1, local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder > > operator[](size_t Index) const
Definition: accessor.hpp:2974
sycl::_V1::image_target
image_target
Definition: access.hpp:74
sycl::_V1::unsampled_image
Definition: image.hpp:727
sycl::_V1::unsampled_image_accessor::operator==
bool operator==(const unsampled_image_accessor &Rhs) const
Definition: accessor.hpp:3729
accessor_properties.hpp
sycl::_V1::access::mode::discard_write
@ discard_write
buffer_properties.hpp
sycl::_V1::detail::getSyclObjImpl
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: impl_utils.hpp:30
sycl::_V1::host_sampled_image_accessor::operator==
bool operator==(const host_sampled_image_accessor &Rhs) const
Definition: accessor.hpp:4097
sycl::_V1::host_accessor::host_accessor
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:3476
sycl::_V1::detail::image_accessor::size
size_t size() const noexcept
Definition: accessor.hpp:1039
sycl::_V1::detail::LocalAccessorBaseDevice::MemRange
range< Dims > MemRange
Definition: accessor.hpp:496
sycl::_V1::detail::const_if_const_AS
DataT const_if_const_AS
Definition: type_traits.hpp:480
sycl::_V1::detail::UnsampledImageAccessorBaseHost::UnsampledImageAccessorBaseHost
UnsampledImageAccessorBaseHost(const UnsampledImageAccessorImplPtr &Impl)
Definition: accessor.hpp:642
type_list.hpp
__SYCL_EBO
#define __SYCL_EBO
Definition: common.hpp:250
sycl::_V1::host_sampled_image_accessor::reference
const DataT & reference
Definition: accessor.hpp:4062
sycl::_V1::detail::accessor_common::IsHostTask
constexpr static bool IsHostTask
Definition: accessor.hpp:317
sycl::_V1::host_accessor::host_accessor
host_accessor(buffer< T, 1, AllocatorT > &BufferRef, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:3436
sycl::_V1::access::address_space
address_space
Definition: access.hpp:51
sycl::_V1::access::target::image_array
@ image_array
Definition: accessor.hpp:3314
sycl::_V1::detail::AccessorBaseHost
Definition: accessor.hpp:534
sycl::_V1::sampled_image_accessor::reference
const DataT & reference
Definition: accessor.hpp:3945