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 
12 #include <sycl/atomic.hpp>
13 #include <sycl/buffer.hpp>
15 #include <sycl/detail/cl.h>
16 #include <sycl/detail/common.hpp>
17 #include <sycl/detail/export.hpp>
23 #include <sycl/device.hpp>
24 #include <sycl/exception.hpp>
27 #include <sycl/id.hpp>
28 #include <sycl/image.hpp>
29 #include <sycl/pointers.hpp>
32 #include <sycl/property_list.hpp>
34 #include <sycl/sampler.hpp>
35 
36 #include <iterator>
37 #include <type_traits>
38 
39 #include <utility>
40 
100 // +------------------+ +-----------------+ +-----------------------+
101 // | | | | | |
102 // | AccessorBaseHost | | accessor_common | | LocalAccessorBaseHost |
103 // | | | | | |
104 // +------------------+ +-----+-----------+ +--------+--------------+
105 // | | | | |
106 // | +-----------+ +----+ +---------+ +------+
107 // | | | | |
108 // v v v v v
109 // +----------------+ +-----------------+ +-------------+
110 // | | | accessor(1) | | accessor(3) |
111 // | image_accessor | +-----------------| +-------------+
112 // | | | for targets: | | for target: |
113 // +---+---+---+----+ | | | |
114 // | | | | host_buffer | | local |
115 // | | | | global_buffer | +-------------+
116 // | | | | constant_buffer |
117 // | | | +-----------------+
118 // | | |
119 // | | +------------------------------------+
120 // | | |
121 // | +----------------------+ |
122 // v v v
123 // +-----------------+ +--------------+ +-------------+
124 // | acessor(2) | | accessor(4) | | accessor(5) |
125 // +-----------------+ +--------------+ +-------------+
126 // | for targets: | | for targets: | | for target: |
127 // | | | | | |
128 // | host_image | | image | | image_array |
129 // +-----------------+ +--------------+ +-------------+
130 //
155 //
156 // +-----------------+
157 // | |
158 // | accessor_common |
159 // | |
160 // +-----+-------+---+
161 // | |
162 // +----+ +-----+
163 // | |
164 // v v
165 // +----------------+ +-----------------+ +-------------+
166 // | | | accessor(1) | | accessor(3) |
167 // | image_accessor | +-----------------| +-------------+
168 // | | | for targets: | | for target: |
169 // +---+---+---+----+ | | | |
170 // | | | | host_buffer | | local |
171 // | | | | global_buffer | +-------------+
172 // | | | | constant_buffer |
173 // | | | +-----------------+
174 // | | | |
175 // | | | v
176 // | | | +-----------------+
177 // | | | | |
178 // | | | | host_accessor |
179 // | | | | |
180 // | | | +-----------------+
181 // | | |
182 // | | +------------------------------------+
183 // | | |
184 // | +----------------------+ |
185 // v v v
186 // +-----------------+ +--------------+ +-------------+
187 // | acessor(2) | | accessor(4) | | accessor(5) |
188 // +-----------------+ +--------------+ +-------------+
189 // | for targets: | | for targets: | | for target: |
190 // | | | | | |
191 // | host_image | | image | | image_array |
192 // +-----------------+ +--------------+ +-------------+
193 //
211 
212 namespace sycl {
214 class stream;
215 namespace ext::intel::esimd::detail {
216 // Forward declare a "back-door" access class to support ESIMD.
217 class AccessorPrivateProxy;
218 } // namespace ext::intel::esimd::detail
219 
220 template <typename DataT, int Dimensions = 1,
222  access::target AccessTarget = access::target::device,
223  access::placeholder IsPlaceholder = access::placeholder::false_t,
224  typename PropertyListT = ext::oneapi::accessor_property_list<>>
225 class accessor;
226 
227 namespace detail {
228 
229 // A helper structure which is shared between buffer accessor and accessor_impl
230 // TODO: Unify with AccessorImplDevice?
231 struct AccHostDataT {
232  AccHostDataT(const sycl::id<3> &Offset, const sycl::range<3> &Range,
233  const sycl::range<3> &MemoryRange, void *Data = nullptr)
234  : MOffset(Offset), MAccessRange(Range), MMemoryRange(MemoryRange),
235  MData(Data) {}
236 
237  sycl::id<3> MOffset;
238  sycl::range<3> MAccessRange;
239  sycl::range<3> MMemoryRange;
240  void *MData = nullptr;
241  void *Reserved = nullptr;
242 };
243 
244 void __SYCL_EXPORT constructorNotification(void *BufferObj, void *AccessorObj,
245  access::target Target,
246  access::mode Mode,
247  const code_location &CodeLoc);
248 
249 template <typename T>
250 using IsPropertyListT = typename std::is_base_of<PropertyListBase, T>;
251 
252 template <typename T>
254  typename std::is_same<ext::oneapi::accessor_property_list<>, T>;
255 
256 template <typename T> struct IsCxPropertyList {
257  constexpr static bool value = false;
258 };
259 
260 template <typename... Props>
261 struct IsCxPropertyList<ext::oneapi::accessor_property_list<Props...>> {
262  constexpr static bool value = true;
263 };
264 
265 template <> struct IsCxPropertyList<ext::oneapi::accessor_property_list<>> {
266  constexpr static bool value = false;
267 };
268 
269 __SYCL_EXPORT device getDeviceFromHandler(handler &CommandGroupHandlerRef);
270 
271 template <typename DataT, int Dimensions, access::mode AccessMode,
273  typename PropertyListT = ext::oneapi::accessor_property_list<>>
275 protected:
277 
278  constexpr static bool IsHostBuf = AccessTarget == access::target::host_buffer;
279  constexpr static bool IsHostTask = AccessTarget == access::target::host_task;
280  // SYCL2020 4.7.6.9.4.3
281  // IsPlaceHolder template parameter has no bearing on whether the accessor
282  // instance is a placeholder. This is determined solely by the constructor.
283  // The rule seems to be: if the constructor receives a CommandGroupHandler
284  // it is NOT a placeholder. Otherwise, it is a placeholder.
285  // However, according to 4.7.6.9.4.6. accessor specialization with
286  // target::host_buffer is never a placeholder. So, if the constructor
287  // used receives a CommandGroupHandler, the accessor will never be a
288  // placeholder. If it doesn't, but IsHostBuf is true, it won't be a
289  // placeholder either. Otherwise, the accessor is a placeholder.
290  constexpr static bool IsPlaceH = !IsHostBuf;
291 
292  // TODO: SYCL 2020 deprecates four of the target enum values
293  // and replaces them with 2 (device and host_task). May want
294  // to change these constexpr.
295  constexpr static bool IsGlobalBuf =
296  AccessTarget == access::target::global_buffer;
297 
298  constexpr static bool IsConstantBuf =
299  AccessTarget == access::target::constant_buffer;
300 
301  constexpr static bool IsAccessAnyWrite =
304  AccessMode == access::mode::discard_write ||
305  AccessMode == access::mode::discard_read_write;
306 
307  constexpr static bool IsAccessReadOnly = AccessMode == access::mode::read;
308  static constexpr bool IsConst = std::is_const_v<DataT>;
309 
310  constexpr static bool IsAccessReadWrite =
312 
313  constexpr static bool IsAccessAtomic = AccessMode == access::mode::atomic;
314 
316  using ConstRefType = const DataT &;
318 
319  // The class which allows to access value of N dimensional accessor using N
320  // subscript operators, e.g. accessor[2][2][3]
321  template <int SubDims,
322  typename AccType =
323  accessor<DataT, Dimensions, AccessMode, AccessTarget,
324  IsPlaceholder, PropertyListT>>
326  static constexpr int Dims = Dimensions;
327 
328  mutable id<Dims> MIDs;
329  AccType MAccessor;
330 
331  public:
332  AccessorSubscript(AccType Accessor, id<Dims> IDs)
333  : MIDs(IDs), MAccessor(Accessor) {}
334 
335  // Only accessor class is supposed to use this c'tor for the first
336  // operator[].
337  AccessorSubscript(AccType Accessor, size_t Index) : MAccessor(Accessor) {
338  MIDs[0] = Index;
339  }
340 
341  template <int CurDims = SubDims, typename = std::enable_if_t<(CurDims > 1)>>
342  auto operator[](size_t Index) {
343  MIDs[Dims - CurDims] = Index;
344  return AccessorSubscript<CurDims - 1, AccType>(MAccessor, MIDs);
345  }
346 
347  template <int CurDims = SubDims,
348  typename = std::enable_if_t<CurDims == 1 && (IsAccessReadOnly ||
349  IsAccessAnyWrite)>>
350  typename AccType::reference operator[](size_t Index) const {
351  MIDs[Dims - CurDims] = Index;
352  return MAccessor[MIDs];
353  }
354 
355  template <int CurDims = SubDims>
356  typename std::enable_if_t<CurDims == 1 && IsAccessAtomic, atomic<DataT, AS>>
357  operator[](size_t Index) const {
358  MIDs[Dims - CurDims] = Index;
359  return MAccessor[MIDs];
360  }
361  };
362 };
363 
364 template <typename DataT> constexpr access::mode accessModeFromConstness() {
365  if constexpr (std::is_const_v<DataT>)
366  return access::mode::read;
367  else
369 }
370 
371 template <typename MayBeTag1, typename MayBeTag2>
373  // property_list = {} is not properly detected by deduction guide,
374  // when parameter is passed without curly braces: access(buffer, no_init)
375  // thus simplest approach is to check 2 last arguments for being a tag
376  if constexpr (std::is_same_v<MayBeTag1, mode_tag_t<access::mode::read>> ||
377  std::is_same_v<MayBeTag2, mode_tag_t<access::mode::read>>) {
378  return access::mode::read;
379  }
380 
381  if constexpr (std::is_same_v<MayBeTag1, mode_tag_t<access::mode::write>> ||
382  std::is_same_v<MayBeTag2, mode_tag_t<access::mode::write>>) {
383  return access::mode::write;
384  }
385 
386  if constexpr (std::is_same_v<
387  MayBeTag1,
388  mode_target_tag_t<access::mode::read,
389  access::target::constant_buffer>> ||
390  std::is_same_v<
391  MayBeTag2,
392  mode_target_tag_t<access::mode::read,
393  access::target::constant_buffer>>) {
394  return access::mode::read;
395  }
396 
397  if constexpr (std::is_same_v<MayBeTag1,
398  mode_target_tag_t<access::mode::read,
399  access::target::host_task>> ||
400  std::is_same_v<MayBeTag2,
401  mode_target_tag_t<access::mode::read,
402  access::target::host_task>>) {
403  return access::mode::read;
404  }
405 
406  if constexpr (std::is_same_v<MayBeTag1,
408  access::target::host_task>> ||
409  std::is_same_v<MayBeTag2,
411  access::target::host_task>>) {
412  return access::mode::write;
413  }
414 
416 }
417 
418 template <typename MayBeTag1, typename MayBeTag2>
420  if constexpr (std::is_same_v<
421  MayBeTag1,
422  mode_target_tag_t<access::mode::read,
423  access::target::constant_buffer>> ||
424  std::is_same_v<
425  MayBeTag2,
426  mode_target_tag_t<access::mode::read,
427  access::target::constant_buffer>>) {
428  return access::target::constant_buffer;
429  }
430 
431  if constexpr (
432  std::is_same_v<MayBeTag1, mode_target_tag_t<access::mode::read,
433  access::target::host_task>> ||
434  std::is_same_v<MayBeTag2, mode_target_tag_t<access::mode::read,
435  access::target::host_task>> ||
436  std::is_same_v<MayBeTag1, mode_target_tag_t<access::mode::write,
437  access::target::host_task>> ||
438  std::is_same_v<MayBeTag2, mode_target_tag_t<access::mode::write,
439  access::target::host_task>> ||
440  std::is_same_v<MayBeTag1, mode_target_tag_t<access::mode::read_write,
441  access::target::host_task>> ||
442  std::is_same_v<MayBeTag2, mode_target_tag_t<access::mode::read_write,
443  access::target::host_task>>) {
444  return access::target::host_task;
445  }
446 
447  return defaultTarget;
448 }
449 
450 template <int Dims> class LocalAccessorBaseDevice {
451 public:
452  LocalAccessorBaseDevice(sycl::range<Dims> Size)
453  : AccessRange(Size),
454  MemRange(InitializedVal<Dims, range>::template get<0>()) {}
455  // TODO: Actually we need only one field here, but currently compiler requires
456  // all of them.
460 
461  bool operator==(const LocalAccessorBaseDevice &Rhs) const {
462  return (AccessRange == Rhs.AccessRange);
463  }
464 };
465 
466 // The class describes a requirement to access a SYCL memory object such as
467 // sycl::buffer and sycl::image. For example, each accessor used in a kernel,
468 // except one with access target "local", adds such requirement for the command
469 // group.
470 
471 template <int Dims> class AccessorImplDevice {
472 public:
473  AccessorImplDevice() = default;
475  range<Dims> MemoryRange)
476  : Offset(Offset), AccessRange(AccessRange), MemRange(MemoryRange) {}
477 
481 
482  bool operator==(const AccessorImplDevice &Rhs) const {
483  return (Offset == Rhs.Offset && AccessRange == Rhs.AccessRange &&
484  MemRange == Rhs.MemRange);
485  }
486 };
487 
488 class AccessorImplHost;
489 
490 void __SYCL_EXPORT addHostAccessorAndWait(AccessorImplHost *Req);
491 
492 class SYCLMemObjI;
493 
494 using AccessorImplPtr = std::shared_ptr<AccessorImplHost>;
495 
496 class __SYCL_EXPORT AccessorBaseHost {
497 protected:
498  AccessorBaseHost(const AccessorImplPtr &Impl) : impl{Impl} {}
499 
500 public:
501  AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
502  access::mode AccessMode, void *SYCLMemObject, int Dims,
503  int ElemSize, int OffsetInBytes = 0,
504  bool IsSubBuffer = false,
505  const property_list &PropertyList = {});
506 
507  AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
508  access::mode AccessMode, void *SYCLMemObject, int Dims,
509  int ElemSize, bool IsPlaceH, int OffsetInBytes = 0,
510  bool IsSubBuffer = false,
511  const property_list &PropertyList = {});
512 
513 public:
514  id<3> &getOffset();
515  range<3> &getAccessRange();
516  range<3> &getMemoryRange();
517  void *getPtr();
518  unsigned int getElemSize() const;
519 
520  const id<3> &getOffset() const;
521  const range<3> &getAccessRange() const;
522  const range<3> &getMemoryRange() const;
523  void *getPtr() const;
524  bool isPlaceholder() const;
525 
526  detail::AccHostDataT &getAccData();
527 
528  const property_list &getPropList() const;
529 
530  void *getMemoryObject() const;
531 
532  template <class Obj>
533  friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject);
534 
535  template <class T>
536  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
537 
538  template <typename, int, access::mode, access::target, access::placeholder,
539  typename>
540  friend class accessor;
541 
543 
544 private:
545  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
546 };
547 
549 using LocalAccessorImplPtr = std::shared_ptr<LocalAccessorImplHost>;
550 
551 class __SYCL_EXPORT LocalAccessorBaseHost {
552 protected:
553  LocalAccessorBaseHost(const LocalAccessorImplPtr &Impl) : impl{Impl} {}
554 
555 public:
556  LocalAccessorBaseHost(sycl::range<3> Size, int Dims, int ElemSize,
557  const property_list &PropertyList = {});
558  sycl::range<3> &getSize();
559  const sycl::range<3> &getSize() const;
560  void *getPtr();
561  void *getPtr() const;
562  int getNumOfDims();
563  int getElementSize();
564  const property_list &getPropList() const;
565 
566 protected:
567  template <class Obj>
568  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
569 
570  template <class T>
571  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
572 
574 };
575 
579  std::shared_ptr<UnsampledImageAccessorImplHost>;
581  std::shared_ptr<SampledImageAccessorImplHost>;
582 
583 void __SYCL_EXPORT
585 void __SYCL_EXPORT
587 
588 class __SYCL_EXPORT UnsampledImageAccessorBaseHost {
589 protected:
591  : impl{Impl} {}
592 
593 public:
594  UnsampledImageAccessorBaseHost(sycl::range<3> Size, access_mode AccessMode,
595  void *SYCLMemObject, int Dims, int ElemSize,
596  id<3> Pitch, image_channel_type ChannelType,
597  image_channel_order ChannelOrder,
598  const property_list &PropertyList = {});
599  const sycl::range<3> &getSize() const;
600  void *getMemoryObject() const;
601  detail::AccHostDataT &getAccData();
602  void *getPtr();
603  void *getPtr() const;
604  int getNumOfDims() const;
605  int getElementSize() const;
606  id<3> getPitch() const;
607  image_channel_type getChannelType() const;
608  image_channel_order getChannelOrder() const;
609  const property_list &getPropList() const;
610 
611 protected:
612  template <class Obj>
613  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
614 
615  template <class T>
616  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
617 
619 };
620 
621 class __SYCL_EXPORT SampledImageAccessorBaseHost {
622 protected:
624  : impl{Impl} {}
625 
626 public:
627  SampledImageAccessorBaseHost(sycl::range<3> Size, void *SYCLMemObject,
628  int Dims, int ElemSize, id<3> Pitch,
629  image_channel_type ChannelType,
630  image_channel_order ChannelOrder,
631  image_sampler Sampler,
632  const property_list &PropertyList = {});
633  const sycl::range<3> &getSize() const;
634  void *getMemoryObject() const;
635  detail::AccHostDataT &getAccData();
636  void *getPtr();
637  void *getPtr() const;
638  int getNumOfDims() const;
639  int getElementSize() const;
640  id<3> getPitch() const;
641  image_channel_type getChannelType() const;
642  image_channel_order getChannelOrder() const;
643  image_sampler getSampler() const;
644  const property_list &getPropList() const;
645 
646 protected:
647  template <class Obj>
648  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
649 
650  template <class T>
651  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
652 
654 };
655 
656 template <int Dim, typename T> struct IsValidCoordDataT;
657 template <typename T> struct IsValidCoordDataT<1, T> {
658  constexpr static bool value = detail::is_contained<
660 };
661 template <typename T> struct IsValidCoordDataT<2, T> {
662  constexpr static bool value = detail::is_contained<
664  vec<opencl::cl_float, 2>>>::type::value;
665 };
666 template <typename T> struct IsValidCoordDataT<3, T> {
667  constexpr static bool value = detail::is_contained<
669  vec<opencl::cl_float, 4>>>::type::value;
670 };
671 
672 template <int Dim, typename T> struct IsValidUnsampledCoord2020DataT;
673 template <typename T> struct IsValidUnsampledCoord2020DataT<1, T> {
674  constexpr static bool value = std::is_same_v<T, int>;
675 };
676 template <typename T> struct IsValidUnsampledCoord2020DataT<2, T> {
677  constexpr static bool value = std::is_same_v<T, int2>;
678 };
679 template <typename T> struct IsValidUnsampledCoord2020DataT<3, T> {
680  constexpr static bool value = std::is_same_v<T, int4>;
681 };
682 
683 template <int Dim, typename T> struct IsValidSampledCoord2020DataT;
684 template <typename T> struct IsValidSampledCoord2020DataT<1, T> {
685  constexpr static bool value = std::is_same_v<T, float>;
686 };
687 template <typename T> struct IsValidSampledCoord2020DataT<2, T> {
688  constexpr static bool value = std::is_same_v<T, float2>;
689 };
690 template <typename T> struct IsValidSampledCoord2020DataT<3, T> {
691  constexpr static bool value = std::is_same_v<T, float4>;
692 };
693 
694 template <typename DataT, int Dimensions, access::mode AccessMode,
697 
698 // Image accessor
699 template <typename DataT, int Dimensions, access::mode AccessMode,
702 #ifndef __SYCL_DEVICE_ONLY__
703  : public detail::AccessorBaseHost {
704  size_t MImageCount;
705  image_channel_order MImgChannelOrder;
706  image_channel_type MImgChannelType;
707 #else
708 {
709 
710  using OCLImageTy = typename detail::opencl_image_type<Dimensions, AccessMode,
711  AccessTarget>::type;
712  OCLImageTy MImageObj;
713  char MPadding[sizeof(detail::AccessorBaseHost) +
714  sizeof(size_t /*MImageCount*/) + sizeof(image_channel_order) +
715  sizeof(image_channel_type) - sizeof(OCLImageTy)];
716 
717 protected:
718  void imageAccessorInit(OCLImageTy Image) { MImageObj = Image; }
719 
720 private:
721 #endif
722  template <typename T1, int T2, access::mode T3, access::placeholder T4>
723  friend class __image_array_slice__;
724 
725  constexpr static bool IsHostImageAcc =
726  (AccessTarget == access::target::host_image);
727 
728  constexpr static bool IsImageAcc = (AccessTarget == access::target::image);
729 
730  constexpr static bool IsImageArrayAcc =
731  (AccessTarget == access::target::image_array);
732 
733  constexpr static bool IsImageAccessWriteOnly =
735  AccessMode == access::mode::discard_write);
736 
737  constexpr static bool IsImageAccessAnyWrite =
738  (IsImageAccessWriteOnly || AccessMode == access::mode::read_write);
739 
740  constexpr static bool IsImageAccessReadOnly =
741  (AccessMode == access::mode::read);
742 
743  constexpr static bool IsImageAccessAnyRead =
744  (IsImageAccessReadOnly || AccessMode == access::mode::read_write);
745 
746  static_assert(std::is_same_v<DataT, vec<opencl::cl_int, 4>> ||
747  std::is_same_v<DataT, vec<opencl::cl_uint, 4>> ||
748  std::is_same_v<DataT, vec<opencl::cl_float, 4>> ||
749  std::is_same_v<DataT, vec<opencl::cl_half, 4>>,
750  "The data type of an image accessor must be only cl_int4, "
751  "cl_uint4, cl_float4 or cl_half4 from SYCL namespace");
752 
753  static_assert(IsImageAcc || IsHostImageAcc || IsImageArrayAcc,
754  "Expected image type");
755 
756  static_assert(IsPlaceholder == access::placeholder::false_t,
757  "Expected false as Placeholder value for image accessor.");
758 
759  static_assert(
760  ((IsImageAcc || IsImageArrayAcc) &&
761  (IsImageAccessWriteOnly || IsImageAccessReadOnly)) ||
762  (IsHostImageAcc && (IsImageAccessAnyWrite || IsImageAccessAnyRead)),
763  "Access modes can be only read/write/discard_write for image/image_array "
764  "target accessor, or they can be only "
765  "read/write/discard_write/read_write for host_image target accessor.");
766 
767  static_assert(Dimensions > 0 && Dimensions <= 3,
768  "Dimensions can be 1/2/3 for image accessor.");
769 
770 #ifdef __SYCL_DEVICE_ONLY__
771 
772  sycl::vec<int, Dimensions> getRangeInternal() const {
773  return __invoke_ImageQuerySize<sycl::vec<int, Dimensions>, OCLImageTy>(
774  MImageObj);
775  }
776 
777  size_t getElementSize() const {
778  int ChannelType = __invoke_ImageQueryFormat<int, OCLImageTy>(MImageObj);
779  int ChannelOrder = __invoke_ImageQueryOrder<int, OCLImageTy>(MImageObj);
780  int ElementSize = getSPIRVElementSize(ChannelType, ChannelOrder);
781  return ElementSize;
782  }
783 
784 #else
785 
786  sycl::vec<int, Dimensions> getRangeInternal() const {
787  // TODO: Implement for host.
788  throw runtime_error("image::getRangeInternal() is not implemented for host",
789  PI_ERROR_INVALID_OPERATION);
790  return sycl::vec<int, Dimensions>{1};
791  }
792 
793 #endif
794 
795 #ifndef __SYCL_DEVICE_ONLY__
796 protected:
798 #endif // __SYCL_DEVICE_ONLY__
799 
800 private:
801  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
802 
803 #ifdef __SYCL_DEVICE_ONLY__
804  const OCLImageTy getNativeImageObj() const { return MImageObj; }
805 #endif // __SYCL_DEVICE_ONLY__
806 
807 public:
808  using value_type = DataT;
809  using reference = DataT &;
810  using const_reference = const DataT &;
811 
812  // image_accessor Constructors.
813 
814 #ifdef __SYCL_DEVICE_ONLY__
815  // Default constructor for objects later initialized with __init member.
816  image_accessor() : MImageObj() {}
817 #endif
818 
819  // Available only when: accessTarget == access::target::host_image
820  // template <typename AllocatorT>
821  // accessor(image<dimensions, AllocatorT> &imageRef);
822  template <
823  typename AllocatorT, int Dims = Dimensions,
824  typename = std::enable_if_t<(Dims > 0 && Dims <= 3) && IsHostImageAcc>>
825  image_accessor(image<Dims, AllocatorT> &ImageRef, int ImageElementSize)
826 #ifdef __SYCL_DEVICE_ONLY__
827  {
828  (void)ImageRef;
829  (void)ImageElementSize;
830  // No implementation needed for device. The constructor is only called by
831  // host.
832  }
833 #else
834  : AccessorBaseHost({ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
835  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
836  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
837  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
838  Dimensions, ImageElementSize),
839  MImageCount(ImageRef.size()),
840  MImgChannelOrder(ImageRef.getChannelOrder()),
841  MImgChannelType(ImageRef.getChannelType()) {
842  addHostAccessorAndWait(AccessorBaseHost::impl.get());
843  }
844 #endif
845 
846  // Available only when: accessTarget == access::target::image
847  // template <typename AllocatorT>
848  // accessor(image<dimensions, AllocatorT> &imageRef,
849  // handler &commandGroupHandlerRef);
850  template <typename AllocatorT, int Dims = Dimensions,
851  typename = std::enable_if_t<(Dims > 0 && Dims <= 3) && IsImageAcc>>
853  handler &CommandGroupHandlerRef, int ImageElementSize)
854 #ifdef __SYCL_DEVICE_ONLY__
855  {
856  (void)ImageRef;
857  (void)CommandGroupHandlerRef;
858  (void)ImageElementSize;
859  // No implementation needed for device. The constructor is only called by
860  // host.
861  }
862 #else
863  : AccessorBaseHost({ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
864  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
865  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
866  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
867  Dimensions, ImageElementSize),
868  MImageCount(ImageRef.size()),
869  MImgChannelOrder(ImageRef.getChannelOrder()),
870  MImgChannelType(ImageRef.getChannelType()) {
871 
872  device Device = getDeviceFromHandler(CommandGroupHandlerRef);
873  if (!Device.has(aspect::ext_intel_legacy_image))
874  throw feature_not_supported(
875  "SYCL 1.2.1 images are not supported by this device.",
876  PI_ERROR_INVALID_OPERATION);
877  }
878 #endif
879 
880  /* -- common interface members -- */
881 
882  // operator == and != need to be defined only for host application as per the
883  // SYCL spec 1.2.1
884 #ifndef __SYCL_DEVICE_ONLY__
885  bool operator==(const image_accessor &Rhs) const { return Rhs.impl == impl; }
886 #else
887  // The operator with __SYCL_DEVICE_ONLY__ need to be declared for compilation
888  // of host application with device compiler.
889  // Usage of this operator inside the kernel code will give a runtime failure.
890  bool operator==(const image_accessor &Rhs) const;
891 #endif
892 
893  bool operator!=(const image_accessor &Rhs) const { return !(Rhs == *this); }
894 
895  // get_count() method : Returns the number of elements of the SYCL image this
896  // SYCL accessor is accessing.
897  //
898  // get_range() method : Returns a range object which represents the number of
899  // elements of dataT per dimension that this accessor may access.
900  // The range object returned must equal to the range of the image this
901  // accessor is associated with.
902 
903 #ifdef __SYCL_DEVICE_ONLY__
904 
905  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
906  size_t get_count() const { return size(); }
907  size_t size() const noexcept { return get_range<Dimensions>().size(); }
908 
909  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 1>>
910  range<1> get_range() const {
911  int Range = getRangeInternal();
912  return range<1>(Range);
913  }
914  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 2>>
915  range<2> get_range() const {
916  int2 Range = getRangeInternal();
917  return range<2>(Range[0], Range[1]);
918  }
919  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 3>>
920  range<3> get_range() const {
921  int3 Range = getRangeInternal();
922  return range<3>(Range[0], Range[1], Range[2]);
923  }
924 
925 #else
926  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
927  size_t get_count() const { return size(); };
928  size_t size() const noexcept { return MImageCount; };
929 
930  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
932  return detail::convertToArrayOfN<Dims, 1>(getAccessRange());
933  }
934 
935 #endif
936 
937  // Available only when:
938  // (accessTarget == access::target::image && accessMode == access::mode::read)
939  // || (accessTarget == access::target::host_image && ( accessMode ==
940  // access::mode::read || accessMode == access::mode::read_write))
941  template <typename CoordT, int Dims = Dimensions,
942  typename = std::enable_if_t<
943  (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
945  ((IsImageAcc && IsImageAccessReadOnly) ||
946  (IsHostImageAcc && IsImageAccessAnyRead))>>
947  DataT read(const CoordT &Coords) const {
948 #ifdef __SYCL_DEVICE_ONLY__
949  return __invoke__ImageRead<DataT, OCLImageTy, CoordT>(MImageObj, Coords);
950 #else
951  sampler Smpl(coordinate_normalization_mode::unnormalized,
952  addressing_mode::none, filtering_mode::nearest);
953  return read<CoordT, Dims>(Coords, Smpl);
954 #endif
955  }
956 
957  // Available only when:
958  // (accessTarget == access::target::image && accessMode == access::mode::read)
959  // || (accessTarget == access::target::host_image && ( accessMode ==
960  // access::mode::read || accessMode == access::mode::read_write))
961  template <typename CoordT, int Dims = Dimensions,
962  typename = std::enable_if_t<
963  (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
964  ((IsImageAcc && IsImageAccessReadOnly) ||
965  (IsHostImageAcc && IsImageAccessAnyRead))>>
966  DataT read(const CoordT &Coords, const sampler &Smpl) const {
967 #ifdef __SYCL_DEVICE_ONLY__
968  return __invoke__ImageReadSampler<DataT, OCLImageTy, CoordT>(
969  MImageObj, Coords, Smpl.impl.m_Sampler);
970 #else
971  return imageReadSamplerHostImpl<CoordT, DataT>(
972  Coords, Smpl, getAccessRange() /*Image Range*/,
973  getOffset() /*Image Pitch*/, MImgChannelType, MImgChannelOrder,
974  AccessorBaseHost::getPtr() /*ptr to image*/,
975  AccessorBaseHost::getElemSize());
976 #endif
977  }
978 
979  // Available only when:
980  // (accessTarget == access::target::image && (accessMode ==
981  // access::mode::write || accessMode == access::mode::discard_write)) ||
982  // (accessTarget == access::target::host_image && (accessMode ==
983  // access::mode::write || accessMode == access::mode::discard_write ||
984  // accessMode == access::mode::read_write))
985  template <typename CoordT, int Dims = Dimensions,
986  typename = std::enable_if_t<
987  (Dims > 0) && (detail::is_genint<CoordT>::value) &&
989  ((IsImageAcc && IsImageAccessWriteOnly) ||
990  (IsHostImageAcc && IsImageAccessAnyWrite))>>
991  void write(const CoordT &Coords, const DataT &Color) const {
992 #ifdef __SYCL_DEVICE_ONLY__
993  __invoke__ImageWrite<OCLImageTy, CoordT, DataT>(MImageObj, Coords, Color);
994 #else
995  imageWriteHostImpl(Coords, Color, getOffset() /*ImagePitch*/,
996  AccessorBaseHost::getElemSize(), MImgChannelType,
997  MImgChannelOrder,
998  AccessorBaseHost::getPtr() /*Ptr to Image*/);
999 #endif
1000  }
1001 };
1002 
1003 template <typename DataT, int Dimensions, access::mode AccessMode,
1005 class __image_array_slice__ {
1006 
1007  static_assert(Dimensions < 3,
1008  "Image slice cannot have more then 2 dimensions");
1009 
1010  constexpr static int AdjustedDims = (Dimensions == 2) ? 4 : Dimensions + 1;
1011 
1012  template <typename CoordT,
1013  typename CoordElemType =
1015  sycl::vec<CoordElemType, AdjustedDims>
1016  getAdjustedCoords(const CoordT &Coords) const {
1017  CoordElemType LastCoord = 0;
1018 
1019  if (std::is_same<float, CoordElemType>::value) {
1020  sycl::vec<int, Dimensions + 1> Size = MBaseAcc.getRangeInternal();
1021  LastCoord =
1022  MIdx / static_cast<float>(Size.template swizzle<Dimensions>());
1023  } else {
1024  LastCoord = MIdx;
1025  }
1026 
1027  sycl::vec<CoordElemType, Dimensions> LeftoverCoords{LastCoord};
1028  sycl::vec<CoordElemType, AdjustedDims> AdjustedCoords{Coords,
1029  LeftoverCoords};
1030  return AdjustedCoords;
1031  }
1032 
1033 public:
1035  accessor<DataT, Dimensions, AccessMode, access::target::image_array,
1037  BaseAcc,
1038  size_t Idx)
1039  : MBaseAcc(BaseAcc), MIdx(Idx) {}
1040 
1041  template <typename CoordT, int Dims = Dimensions,
1042  typename = std::enable_if_t<
1043  (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value)>>
1044  DataT read(const CoordT &Coords) const {
1045  return MBaseAcc.read(getAdjustedCoords(Coords));
1046  }
1047 
1048  template <typename CoordT, int Dims = Dimensions,
1049  typename = std::enable_if_t<(Dims > 0) &&
1051  DataT read(const CoordT &Coords, const sampler &Smpl) const {
1052  return MBaseAcc.read(getAdjustedCoords(Coords), Smpl);
1053  }
1054 
1055  template <typename CoordT, int Dims = Dimensions,
1056  typename = std::enable_if_t<(Dims > 0) &&
1058  void write(const CoordT &Coords, const DataT &Color) const {
1059  return MBaseAcc.write(getAdjustedCoords(Coords), Color);
1060  }
1061 
1062 #ifdef __SYCL_DEVICE_ONLY__
1063  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1064  size_t get_count() const { return size(); }
1065  size_t size() const noexcept { return get_range<Dimensions>().size(); }
1066 
1067  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 1>>
1068  range<1> get_range() const {
1069  int2 Count = MBaseAcc.getRangeInternal();
1070  return range<1>(Count.x());
1071  }
1072  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 2>>
1073  range<2> get_range() const {
1074  int3 Count = MBaseAcc.getRangeInternal();
1075  return range<2>(Count.x(), Count.y());
1076  }
1077 
1078 #else
1079 
1080  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1081  size_t get_count() const { return size(); }
1082  size_t size() const noexcept {
1083  return MBaseAcc.MImageCount / MBaseAcc.getAccessRange()[Dimensions];
1084  }
1085 
1086  template <int Dims = Dimensions,
1087  typename = std::enable_if_t<(Dims == 1 || Dims == 2)>>
1089  return detail::convertToArrayOfN<Dims, 1>(MBaseAcc.getAccessRange());
1090  }
1091 
1092 #endif
1093 
1094 private:
1095  size_t MIdx;
1096  accessor<DataT, Dimensions, AccessMode, access::target::image_array,
1098  MBaseAcc;
1099 };
1100 
1101 } // namespace detail
1102 
1108 template <typename DataT, int Dimensions, access::mode AccessMode,
1110  typename PropertyListT>
1112 #ifndef __SYCL_DEVICE_ONLY__
1113  public detail::AccessorBaseHost,
1114 #endif
1115  public detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
1116  IsPlaceholder, PropertyListT>,
1117  public detail::OwnerLessBase<
1118  accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
1119  PropertyListT>> {
1120 protected:
1121  static_assert((AccessTarget == access::target::global_buffer ||
1122  AccessTarget == access::target::constant_buffer ||
1123  AccessTarget == access::target::host_buffer ||
1124  AccessTarget == access::target::host_task),
1125  "Expected buffer type");
1126 
1127  static_assert((AccessTarget == access::target::global_buffer ||
1128  AccessTarget == access::target::host_buffer ||
1129  AccessTarget == access::target::host_task) ||
1130  (AccessTarget == access::target::constant_buffer &&
1131  AccessMode == access::mode::read),
1132  "Access mode can be only read for constant buffers");
1133 
1134  static_assert(detail::IsPropertyListT<PropertyListT>::value,
1135  "PropertyListT must be accessor_property_list");
1136 
1137  using AccessorCommonT =
1138  detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
1139  IsPlaceholder, PropertyListT>;
1140 
1141  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
1142 
1143  using AccessorCommonT::AS;
1144  // Cannot do "using AccessorCommonT::Flag" as it doesn't work with g++ as host
1145  // compiler, for some reason.
1146  static constexpr bool IsAccessAnyWrite = AccessorCommonT::IsAccessAnyWrite;
1147  static constexpr bool IsAccessReadOnly = AccessorCommonT::IsAccessReadOnly;
1148  static constexpr bool IsConstantBuf = AccessorCommonT::IsConstantBuf;
1149  static constexpr bool IsGlobalBuf = AccessorCommonT::IsGlobalBuf;
1150  static constexpr bool IsHostBuf = AccessorCommonT::IsHostBuf;
1151  static constexpr bool IsPlaceH = AccessorCommonT::IsPlaceH;
1152  static constexpr bool IsConst = AccessorCommonT::IsConst;
1153  static constexpr bool IsHostTask = AccessorCommonT::IsHostTask;
1154  template <int Dims>
1155  using AccessorSubscript =
1156  typename AccessorCommonT::template AccessorSubscript<Dims>;
1157 
1158  static_assert(
1159  !IsConst || IsAccessReadOnly,
1160  "A const qualified DataT is only allowed for a read-only accessor");
1161 
1162  using ConcreteASPtrType = typename detail::DecoratedType<DataT, AS>::type *;
1163 
1164  using RefType = detail::const_if_const_AS<AS, DataT> &;
1165  using ConstRefType = const DataT &;
1166  using PtrType = detail::const_if_const_AS<AS, DataT> *;
1167 
1168  template <int Dims = Dimensions> size_t getLinearIndex(id<Dims> Id) const {
1169 
1170  size_t Result = 0;
1171  detail::loop<Dims>([&, this](size_t I) {
1172  Result = Result * getMemoryRange()[I] + Id[I];
1173  // We've already adjusted for the accessor's offset in the __init, so
1174  // don't include it here in case of device.
1175 #ifndef __SYCL_DEVICE_ONLY__
1176  if constexpr (!(PropertyListT::template has_property<
1178  Result += getOffset()[I];
1179  }
1180 #endif // __SYCL_DEVICE_ONLY__
1181  });
1182 
1183  return Result;
1184  }
1185 
1186  template <typename T, int Dims>
1187  struct IsSameAsBuffer
1188  : std::bool_constant<std::is_same_v<T, DataT> && (Dims > 0) &&
1189  (Dims == Dimensions)> {};
1190 
1191  static access::mode getAdjustedMode(const PropertyListT &PropertyList) {
1192  access::mode AdjustedMode = AccessMode;
1193 
1194  if (PropertyList.template has_property<property::no_init>() ||
1195  PropertyList.template has_property<property::noinit>()) {
1196  if (AdjustedMode == access::mode::write) {
1197  AdjustedMode = access::mode::discard_write;
1198  } else if (AdjustedMode == access::mode::read_write) {
1199  AdjustedMode = access::mode::discard_read_write;
1200  }
1201  }
1202 
1203  return AdjustedMode;
1204  }
1205 
1206  template <typename TagT>
1207  struct IsValidTag
1208  : std::disjunction<
1209  std::is_same<TagT, mode_tag_t<AccessMode>>,
1210  std::is_same<TagT, mode_target_tag_t<AccessMode, AccessTarget>>> {};
1211 
1212  template <typename DataT_, int Dimensions_, access::mode AccessMode_,
1213  access::target AccessTarget_, access::placeholder IsPlaceholder_,
1214  typename PropertyListT_>
1215  friend class accessor;
1216 
1217 #ifdef __SYCL_DEVICE_ONLY__
1218 
1219  id<AdjustedDim> &getOffset() { return impl.Offset; }
1220  range<AdjustedDim> &getAccessRange() { return impl.AccessRange; }
1221  range<AdjustedDim> &getMemoryRange() { return impl.MemRange; }
1222 
1223  const id<AdjustedDim> &getOffset() const { return impl.Offset; }
1224  const range<AdjustedDim> &getAccessRange() const { return impl.AccessRange; }
1225  const range<AdjustedDim> &getMemoryRange() const { return impl.MemRange; }
1226 
1227  detail::AccessorImplDevice<AdjustedDim> impl;
1228 
1229  union {
1230  ConcreteASPtrType MData;
1231  };
1232 
1233  void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
1234  range<AdjustedDim> MemRange, id<AdjustedDim> Offset) {
1235  MData = Ptr;
1236  detail::loop<AdjustedDim>([&, this](size_t I) {
1237  if constexpr (!(PropertyListT::template has_property<
1239  getOffset()[I] = Offset[I];
1240  }
1241  getAccessRange()[I] = AccessRange[I];
1242  getMemoryRange()[I] = MemRange[I];
1243  });
1244 
1245  // Adjust for offsets as that part is invariant for all invocations of
1246  // operator[]. Will have to re-adjust in get_pointer.
1247  MData += getTotalOffset();
1248  }
1249 
1250  // __init variant used by the device compiler for ESIMD kernels.
1251  // TODO: In ESIMD accessors usage is limited for now - access range, mem
1252  // range and offset are not supported.
1253  void __init_esimd(ConcreteASPtrType Ptr) {
1254  MData = Ptr;
1255 #ifdef __ESIMD_FORCE_STATELESS_MEM
1256  detail::loop<AdjustedDim>([&, this](size_t I) {
1257  getOffset()[I] = 0;
1258  getAccessRange()[I] = 0;
1259  getMemoryRange()[I] = 0;
1260  });
1261 #endif
1262  }
1263 
1264  ConcreteASPtrType getQualifiedPtr() const noexcept { return MData; }
1265 
1266 #ifndef __SYCL_DEVICE_ONLY__
1267  using AccessorBaseHost::impl;
1268 #endif
1269 
1270 public:
1271  // Default constructor for objects later initialized with __init member.
1272  accessor()
1273  : impl({}, detail::InitializedVal<AdjustedDim, range>::template get<0>(),
1274  detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
1275 
1276 #else
1277  accessor(const detail::AccessorImplPtr &Impl)
1278  : detail::AccessorBaseHost{Impl} {}
1279 
1280  void *getPtr() { return AccessorBaseHost::getPtr(); }
1281 
1282  const id<3> getOffset() const {
1283  if constexpr (IsHostBuf)
1284  return MAccData ? MAccData->MOffset : id<3>();
1285  else
1286  return AccessorBaseHost::getOffset();
1287  }
1288  const range<3> &getAccessRange() const {
1289  return AccessorBaseHost::getAccessRange();
1290  }
1291  const range<3> getMemoryRange() const {
1292  if constexpr (IsHostBuf)
1293  return MAccData ? MAccData->MMemoryRange : range(0, 0, 0);
1294  else
1295  return AccessorBaseHost::getMemoryRange();
1296  }
1297 
1298  void *getPtr() const { return AccessorBaseHost::getPtr(); }
1299 
1300  void initHostAcc() { MAccData = &getAccData(); }
1301 
1302  // The function references helper methods required by GDB pretty-printers
1303  void GDBMethodsAnchor() {
1304 #ifndef NDEBUG
1305  const auto *this_const = this;
1306  (void)getMemoryRange();
1307  (void)this_const->getMemoryRange();
1308  (void)getOffset();
1309  (void)this_const->getOffset();
1310  (void)getPtr();
1311  (void)this_const->getPtr();
1312  (void)getAccessRange();
1313  (void)this_const->getAccessRange();
1314 #endif
1315  }
1316 
1317  detail::AccHostDataT *MAccData = nullptr;
1318 
1319  char padding[sizeof(detail::AccessorImplDevice<AdjustedDim>) +
1320  sizeof(PtrType) - sizeof(detail::AccessorBaseHost) -
1321  sizeof(MAccData)];
1322 
1323  PtrType getQualifiedPtr() const noexcept {
1324  if constexpr (IsHostBuf)
1325  return MAccData ? reinterpret_cast<PtrType>(MAccData->MData) : nullptr;
1326  else
1327  return reinterpret_cast<PtrType>(AccessorBaseHost::getPtr());
1328  }
1329 
1330 public:
1331  accessor()
1332  : AccessorBaseHost(
1333  /*Offset=*/{0, 0, 0}, /*AccessRange=*/{0, 0, 0},
1334  /*MemoryRange=*/{0, 0, 0},
1335  /*AccessMode=*/getAdjustedMode({}),
1336  /*SYCLMemObject=*/nullptr, /*Dims=*/0, /*ElemSize=*/0,
1337  /*IsPlaceH=*/true,
1338  /*OffsetInBytes=*/0, /*IsSubBuffer=*/false, /*PropertyList=*/{}){};
1339 
1340 #endif // __SYCL_DEVICE_ONLY__
1341 
1342 private:
1343  friend class sycl::stream;
1344  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
1345 
1346  template <class Obj>
1347  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
1348 
1349  template <class T>
1350  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
1351 
1352 public:
1353  // 4.7.6.9.1. Interface for buffer command accessors
1354  // value_type is defined as const DataT for read_only accessors, DataT
1355  // otherwise
1356  using value_type =
1357  std::conditional_t<AccessMode == access_mode::read, const DataT, DataT>;
1358  using reference = value_type &;
1359  using const_reference = const DataT &;
1360 
1361  template <access::decorated IsDecorated>
1362  using accessor_ptr =
1363  std::conditional_t<AccessTarget == access::target::device,
1364  global_ptr<value_type, IsDecorated>, value_type *>;
1365 
1366  using iterator = typename detail::accessor_iterator<value_type, AdjustedDim>;
1367  using const_iterator =
1368  typename detail::accessor_iterator<const value_type, AdjustedDim>;
1369  using reverse_iterator = std::reverse_iterator<iterator>;
1370  using const_reverse_iterator = std::reverse_iterator<const_iterator>;
1371  using difference_type =
1372  typename std::iterator_traits<iterator>::difference_type;
1373  using size_type = std::size_t;
1374 
1375  // The list of accessor constructors with their arguments
1376  // -------+---------+-------+----+-----+--------------
1377  // Dimensions = 0
1378  // -------+---------+-------+----+-----+--------------
1379  // buffer | | | | | property_list
1380  // buffer | handler | | | | property_list
1381  // -------+---------+-------+----+-----+--------------
1382  // Dimensions >= 1
1383  // -------+---------+-------+----+-----+--------------
1384  // buffer | | | | | property_list
1385  // buffer | | | | tag | property_list
1386  // buffer | handler | | | | property_list
1387  // buffer | handler | | | tag | property_list
1388  // buffer | | range | | | property_list
1389  // buffer | | range | | tag | property_list
1390  // buffer | handler | range | | | property_list
1391  // buffer | handler | range | | tag | property_list
1392  // buffer | | range | id | | property_list
1393  // buffer | | range | id | tag | property_list
1394  // buffer | handler | range | id | | property_list
1395  // buffer | handler | range | id | tag | property_list
1396  // -------+---------+-------+----+-----+--------------
1397 
1398 public:
1399  // implicit conversion between const / non-const types for read only accessors
1400  template <typename DataT_,
1401  typename = std::enable_if_t<
1402  IsAccessReadOnly && !std::is_same_v<DataT_, DataT> &&
1403  std::is_same_v<std::remove_const_t<DataT_>,
1404  std::remove_const_t<DataT>>>>
1405  accessor(const accessor<DataT_, Dimensions, AccessMode, AccessTarget,
1406  IsPlaceholder, PropertyListT> &other)
1407 #ifdef __SYCL_DEVICE_ONLY__
1408  : impl(other.impl), MData(other.MData) {
1409 #else
1410  : accessor(other.impl) {
1411 #endif // __SYCL_DEVICE_ONLY__
1412  }
1413 
1414  // implicit conversion from read_write T accessor to read only T (const)
1415  // accessor
1416  template <typename DataT_, access::mode AccessMode_,
1417  typename = std::enable_if_t<
1418  (AccessMode_ == access_mode::read_write) && IsAccessReadOnly &&
1419  std::is_same_v<std::remove_const_t<DataT_>,
1420  std::remove_const_t<DataT>>>>
1421  accessor(const accessor<DataT_, Dimensions, AccessMode_, AccessTarget,
1422  IsPlaceholder, PropertyListT> &other)
1423 #ifdef __SYCL_DEVICE_ONLY__
1424  : impl(other.impl), MData(other.MData) {
1425 #else
1426  : accessor(other.impl) {
1427 #endif // __SYCL_DEVICE_ONLY__
1428  }
1429 
1430  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1431  typename std::enable_if_t<
1432  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1433  std::is_same_v<T, DataT> && Dims == 0 &&
1434  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))> * =
1435  nullptr>
1436  accessor(
1437  buffer<T, 1, AllocatorT> &BufferRef,
1438  const property_list &PropertyList = {},
1439  const detail::code_location CodeLoc = detail::code_location::current())
1440 #ifdef __SYCL_DEVICE_ONLY__
1441  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1442  (void)PropertyList;
1443 #else
1444  : AccessorBaseHost(
1445  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1446  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1447  getAdjustedMode(PropertyList),
1448  detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
1449  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1450  PropertyList) {
1451  preScreenAccessor(BufferRef.size(), PropertyList);
1452  if (!AccessorBaseHost::isPlaceholder())
1453  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1454  initHostAcc();
1456  detail::AccessorBaseHost::impl.get(),
1457  AccessTarget, AccessMode, CodeLoc);
1458  GDBMethodsAnchor();
1459 #endif
1460  }
1461 
1462  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1463  typename... PropTypes,
1464  typename std::enable_if_t<
1465  detail::IsCxPropertyList<PropertyListT>::value &&
1466  // VS2019 can't compile sycl/test/regression/bit_cast_win.cpp
1467  // if std::is_same_v is used here.
1468  std::is_same<T, DataT>::value && Dims == 0 &&
1469  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))> * =
1470  nullptr>
1471  accessor(
1472  buffer<T, 1, AllocatorT> &BufferRef,
1473  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1474  {},
1475  const detail::code_location CodeLoc = detail::code_location::current())
1476 #ifdef __SYCL_DEVICE_ONLY__
1477  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1478  (void)PropertyList;
1479 #else
1480  : AccessorBaseHost(
1481  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1482  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1483  getAdjustedMode(PropertyList),
1484  detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
1485  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1486  PropertyList) {
1487  preScreenAccessor(BufferRef.size(), PropertyList);
1488  if (!AccessorBaseHost::isPlaceholder())
1489  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1490  initHostAcc();
1492  detail::AccessorBaseHost::impl.get(),
1493  AccessTarget, AccessMode, CodeLoc);
1494  GDBMethodsAnchor();
1495 #endif
1496  }
1497 
1498  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1499  typename = typename std::enable_if_t<
1500  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1501  std::is_same_v<T, DataT> && (Dims == 0) &&
1502  (IsGlobalBuf || IsHostBuf || IsConstantBuf || IsHostTask)>>
1503  accessor(
1504  buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1505  const property_list &PropertyList = {},
1506  const detail::code_location CodeLoc = detail::code_location::current())
1507 #ifdef __SYCL_DEVICE_ONLY__
1508  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1509  (void)CommandGroupHandler;
1510  (void)PropertyList;
1511  }
1512 #else
1513  : AccessorBaseHost(
1514  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1515  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1516  getAdjustedMode(PropertyList),
1517  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1518  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1519  preScreenAccessor(BufferRef.size(), PropertyList);
1520  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1521  initHostAcc();
1523  detail::AccessorBaseHost::impl.get(),
1524  AccessTarget, AccessMode, CodeLoc);
1525  GDBMethodsAnchor();
1526  }
1527 #endif
1528 
1529  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1530  typename... PropTypes,
1531  typename = typename std::enable_if_t<
1532  detail::IsCxPropertyList<PropertyListT>::value &&
1533  std::is_same_v<T, DataT> && (Dims == 0) &&
1534  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1535  accessor(
1536  buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1537  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1538  {},
1539  const detail::code_location CodeLoc = detail::code_location::current())
1540 #ifdef __SYCL_DEVICE_ONLY__
1541  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1542  (void)CommandGroupHandler;
1543  (void)PropertyList;
1544  }
1545 #else
1546  : AccessorBaseHost(
1547  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1548  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1549  getAdjustedMode(PropertyList),
1550  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1551  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1552  preScreenAccessor(BufferRef.size(), PropertyList);
1553  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1554  initHostAcc();
1556  detail::AccessorBaseHost::impl.get(),
1557  AccessTarget, AccessMode, CodeLoc);
1558  GDBMethodsAnchor();
1559  }
1560 #endif
1561 
1562  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1563  typename = std::enable_if_t<
1564  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1565  IsSameAsBuffer<T, Dims>::value &&
1566  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1567  accessor(
1568  buffer<T, Dims, AllocatorT> &BufferRef,
1569  const property_list &PropertyList = {},
1570  const detail::code_location CodeLoc = detail::code_location::current())
1571 #ifdef __SYCL_DEVICE_ONLY__
1572  : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1573  (void)PropertyList;
1574  }
1575 #else
1576  : AccessorBaseHost(
1577  /*Offset=*/{0, 0, 0},
1578  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1579  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1580  getAdjustedMode(PropertyList),
1581  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1582  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1583  PropertyList) {
1584  preScreenAccessor(BufferRef.size(), PropertyList);
1585  if (!AccessorBaseHost::isPlaceholder())
1586  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1587  initHostAcc();
1589  detail::AccessorBaseHost::impl.get(),
1590  AccessTarget, AccessMode, CodeLoc);
1591  GDBMethodsAnchor();
1592  }
1593 #endif
1594 
1595  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1596  typename... PropTypes,
1597  typename = std::enable_if_t<
1598  detail::IsCxPropertyList<PropertyListT>::value &&
1599  IsSameAsBuffer<T, Dims>::value &&
1600  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1601  accessor(
1602  buffer<T, Dims, AllocatorT> &BufferRef,
1603  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1604  {},
1605  const detail::code_location CodeLoc = detail::code_location::current())
1606 #ifdef __SYCL_DEVICE_ONLY__
1607  : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1608  (void)PropertyList;
1609  }
1610 #else
1611  : AccessorBaseHost(
1612  /*Offset=*/{0, 0, 0},
1613  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1614  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1615  getAdjustedMode(PropertyList),
1616  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1617  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1618  PropertyList) {
1619  preScreenAccessor(BufferRef.size(), PropertyList);
1620  if (!AccessorBaseHost::isPlaceholder())
1621  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1622  initHostAcc();
1624  detail::AccessorBaseHost::impl.get(),
1625  AccessTarget, AccessMode, CodeLoc);
1626  GDBMethodsAnchor();
1627  }
1628 #endif
1629 
1630  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1631  typename TagT,
1632  typename = std::enable_if_t<
1633  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1634  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1635  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1636  accessor(
1637  buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1638  const property_list &PropertyList = {},
1639  const detail::code_location CodeLoc = detail::code_location::current())
1640  : accessor(BufferRef, PropertyList, CodeLoc) {
1641  adjustAccPropsInBuf(BufferRef);
1642  }
1643 
1644  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1645  typename TagT, typename... PropTypes,
1646  typename = std::enable_if_t<
1647  detail::IsCxPropertyList<PropertyListT>::value &&
1648  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1649  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1650  accessor(
1651  buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1652  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1653  {},
1654  const detail::code_location CodeLoc = detail::code_location::current())
1655  : accessor(BufferRef, PropertyList, CodeLoc) {
1656  adjustAccPropsInBuf(BufferRef);
1657  }
1658 
1659  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1660  typename = std::enable_if_t<
1661  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1662  IsSameAsBuffer<T, Dims>::value &&
1663  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1664  accessor(
1665  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1666  const property_list &PropertyList = {},
1667  const detail::code_location CodeLoc = detail::code_location::current())
1668 #ifdef __SYCL_DEVICE_ONLY__
1669  : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1670  (void)CommandGroupHandler;
1671  (void)PropertyList;
1672  }
1673 #else
1674  : AccessorBaseHost(
1675  /*Offset=*/{0, 0, 0},
1676  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1677  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1678  getAdjustedMode(PropertyList),
1679  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1680  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1681  preScreenAccessor(BufferRef.size(), PropertyList);
1682  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1683  initHostAcc();
1685  detail::AccessorBaseHost::impl.get(),
1686  AccessTarget, AccessMode, CodeLoc);
1687  GDBMethodsAnchor();
1688  }
1689 #endif
1690 
1691  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1692  typename... PropTypes,
1693  typename = std::enable_if_t<
1694  detail::IsCxPropertyList<PropertyListT>::value &&
1695  IsSameAsBuffer<T, Dims>::value &&
1696  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1697  accessor(
1698  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1699  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1700  {},
1701  const detail::code_location CodeLoc = detail::code_location::current())
1702 #ifdef __SYCL_DEVICE_ONLY__
1703  : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1704  (void)CommandGroupHandler;
1705  (void)PropertyList;
1706  }
1707 #else
1708  : AccessorBaseHost(
1709  /*Offset=*/{0, 0, 0},
1710  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1711  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1712  getAdjustedMode(PropertyList),
1713  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1714  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1715  preScreenAccessor(BufferRef.size(), PropertyList);
1716  initHostAcc();
1717  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1719  detail::AccessorBaseHost::impl.get(),
1720  AccessTarget, AccessMode, CodeLoc);
1721  GDBMethodsAnchor();
1722  }
1723 #endif
1724 
1725  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1726  typename TagT,
1727  typename = std::enable_if_t<
1728  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1729  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1730  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1731  accessor(
1732  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1733  TagT, const property_list &PropertyList = {},
1734  const detail::code_location CodeLoc = detail::code_location::current())
1735  : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1736  adjustAccPropsInBuf(BufferRef);
1737  }
1738 
1739  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1740  typename TagT, typename... PropTypes,
1741  typename = std::enable_if_t<
1742  detail::IsCxPropertyList<PropertyListT>::value &&
1743  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1744  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1745  accessor(
1746  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1747  TagT,
1748  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1749  {},
1750  const detail::code_location CodeLoc = detail::code_location::current())
1751  : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1752  adjustAccPropsInBuf(BufferRef);
1753  }
1754 
1755  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1756  typename = std::enable_if_t<
1757  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1758  IsSameAsBuffer<T, Dims>::value &&
1759  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1760  accessor(
1761  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1762  const property_list &PropertyList = {},
1763  const detail::code_location CodeLoc = detail::code_location::current())
1764  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1765 
1766  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1767  typename... PropTypes,
1768  typename = std::enable_if_t<
1769  detail::IsCxPropertyList<PropertyListT>::value &&
1770  IsSameAsBuffer<T, Dims>::value &&
1771  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1772  accessor(
1773  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1774  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1775  {},
1776  const detail::code_location CodeLoc = detail::code_location::current())
1777  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1778 
1779  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1780  typename TagT,
1781  typename = std::enable_if_t<
1782  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1783  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1784  (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
1785  accessor(
1786  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1787  TagT, const property_list &PropertyList = {},
1788  const detail::code_location CodeLoc = detail::code_location::current())
1789  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1790  adjustAccPropsInBuf(BufferRef);
1791  }
1792 
1793  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1794  typename TagT, typename... PropTypes,
1795  typename = std::enable_if_t<
1796  detail::IsCxPropertyList<PropertyListT>::value &&
1797  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1798  (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
1799  accessor(
1800  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1801  TagT,
1802  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1803  {},
1804  const detail::code_location CodeLoc = detail::code_location::current())
1805  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1806  adjustAccPropsInBuf(BufferRef);
1807  }
1808 
1809  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1810  typename = std::enable_if_t<
1811  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1812  IsSameAsBuffer<T, Dims>::value &&
1813  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1814  accessor(
1815  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1816  range<Dimensions> AccessRange, const property_list &PropertyList = {},
1817  const detail::code_location CodeLoc = detail::code_location::current())
1818  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1819  CodeLoc) {}
1820 
1821  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1822  typename... PropTypes,
1823  typename = std::enable_if_t<
1824  detail::IsCxPropertyList<PropertyListT>::value &&
1825  IsSameAsBuffer<T, Dims>::value &&
1826  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1827  accessor(
1828  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1829  range<Dimensions> AccessRange,
1830  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1831  {},
1832  const detail::code_location CodeLoc = detail::code_location::current())
1833  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1834  CodeLoc) {}
1835 
1836  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1837  typename TagT,
1838  typename = std::enable_if_t<
1839  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1840  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1841  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1842  accessor(
1843  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1844  range<Dimensions> AccessRange, TagT,
1845  const property_list &PropertyList = {},
1846  const detail::code_location CodeLoc = detail::code_location::current())
1847  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1848  CodeLoc) {
1849  adjustAccPropsInBuf(BufferRef);
1850  }
1851 
1852  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1853  typename TagT, typename... PropTypes,
1854  typename = std::enable_if_t<
1855  detail::IsCxPropertyList<PropertyListT>::value &&
1856  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1857  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1858  accessor(
1859  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1860  range<Dimensions> AccessRange, TagT,
1861  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1862  {},
1863  const detail::code_location CodeLoc = detail::code_location::current())
1864  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1865  CodeLoc) {
1866  adjustAccPropsInBuf(BufferRef);
1867  }
1868 
1869  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1870  typename = std::enable_if_t<
1871  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1872  IsSameAsBuffer<T, Dims>::value &&
1873  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1874  accessor(
1875  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1876  id<Dimensions> AccessOffset, const property_list &PropertyList = {},
1877  const detail::code_location CodeLoc = detail::code_location::current())
1878 #ifdef __SYCL_DEVICE_ONLY__
1879  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1880  (void)PropertyList;
1881  }
1882 #else
1883  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1884  detail::convertToArrayOfN<3, 1>(AccessRange),
1885  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1886  getAdjustedMode(PropertyList),
1887  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1888  sizeof(DataT), IsPlaceH, BufferRef.OffsetInBytes,
1889  BufferRef.IsSubBuffer, PropertyList) {
1890  preScreenAccessor(BufferRef.size(), PropertyList);
1891  if (!AccessorBaseHost::isPlaceholder())
1892  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1893  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1894  BufferRef.get_range()))
1895  throw sycl::invalid_object_error(
1896  "accessor with requested offset and range would exceed the bounds of "
1897  "the buffer",
1898  PI_ERROR_INVALID_VALUE);
1899 
1900  initHostAcc();
1902  detail::AccessorBaseHost::impl.get(),
1903  AccessTarget, AccessMode, CodeLoc);
1904  GDBMethodsAnchor();
1905  }
1906 #endif
1907 
1908  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1909  typename... PropTypes,
1910  typename = std::enable_if_t<
1911  detail::IsCxPropertyList<PropertyListT>::value &&
1912  IsSameAsBuffer<T, Dims>::value &&
1913  (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1914  accessor(
1915  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1916  id<Dimensions> AccessOffset,
1917  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1918  {},
1919  const detail::code_location CodeLoc = detail::code_location::current())
1920 #ifdef __SYCL_DEVICE_ONLY__
1921  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1922  (void)PropertyList;
1923  }
1924 #else
1925  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1926  detail::convertToArrayOfN<3, 1>(AccessRange),
1927  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1928  getAdjustedMode(PropertyList),
1929  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1930  sizeof(DataT), IsPlaceH, BufferRef.OffsetInBytes,
1931  BufferRef.IsSubBuffer, PropertyList) {
1932  preScreenAccessor(BufferRef.size(), PropertyList);
1933  if (!AccessorBaseHost::isPlaceholder())
1934  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1935  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1936  BufferRef.get_range()))
1937  throw sycl::invalid_object_error(
1938  "accessor with requested offset and range would exceed the bounds of "
1939  "the buffer",
1940  PI_ERROR_INVALID_VALUE);
1941 
1942  initHostAcc();
1944  detail::AccessorBaseHost::impl.get(),
1945  AccessTarget, AccessMode, CodeLoc);
1946  GDBMethodsAnchor();
1947  }
1948 #endif
1949 
1950  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1951  typename TagT,
1952  typename = std::enable_if_t<
1953  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1954  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1955  (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
1956  accessor(
1957  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1958  id<Dimensions> AccessOffset, TagT, const property_list &PropertyList = {},
1959  const detail::code_location CodeLoc = detail::code_location::current())
1960  : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
1961  adjustAccPropsInBuf(BufferRef);
1962  }
1963 
1964  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1965  typename TagT, typename... PropTypes,
1966  typename = std::enable_if_t<
1967  detail::IsCxPropertyList<PropertyListT>::value &&
1968  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1969  (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
1970  accessor(
1971  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1972  id<Dimensions> AccessOffset, TagT,
1973  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1974  {},
1975  const detail::code_location CodeLoc = detail::code_location::current())
1976  : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
1977  adjustAccPropsInBuf(BufferRef);
1978  }
1979 
1980  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1981  typename = std::enable_if_t<
1982  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1983  IsSameAsBuffer<T, Dims>::value &&
1984  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1985  accessor(
1986  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1987  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
1988  const property_list &PropertyList = {},
1989  const detail::code_location CodeLoc = detail::code_location::current())
1990 #ifdef __SYCL_DEVICE_ONLY__
1991  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1992  (void)CommandGroupHandler;
1993  (void)PropertyList;
1994  }
1995 #else
1996  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1997  detail::convertToArrayOfN<3, 1>(AccessRange),
1998  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1999  getAdjustedMode(PropertyList),
2000  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
2001  sizeof(DataT), BufferRef.OffsetInBytes,
2002  BufferRef.IsSubBuffer, PropertyList) {
2003  preScreenAccessor(BufferRef.size(), PropertyList);
2004  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2005  BufferRef.get_range()))
2006  throw sycl::invalid_object_error(
2007  "accessor with requested offset and range would exceed the bounds of "
2008  "the buffer",
2009  PI_ERROR_INVALID_VALUE);
2010 
2011  initHostAcc();
2012  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
2014  detail::AccessorBaseHost::impl.get(),
2015  AccessTarget, AccessMode, CodeLoc);
2016  GDBMethodsAnchor();
2017  }
2018 #endif
2019 
2020  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2021  typename... PropTypes,
2022  typename = std::enable_if_t<
2023  detail::IsCxPropertyList<PropertyListT>::value &&
2024  IsSameAsBuffer<T, Dims>::value &&
2025  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2026  accessor(
2027  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2028  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
2029  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2030  {},
2031  const detail::code_location CodeLoc = detail::code_location::current())
2032 #ifdef __SYCL_DEVICE_ONLY__
2033  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
2034  (void)CommandGroupHandler;
2035  (void)PropertyList;
2036  }
2037 #else
2038  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
2039  detail::convertToArrayOfN<3, 1>(AccessRange),
2040  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
2041  getAdjustedMode(PropertyList),
2042  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
2043  sizeof(DataT), BufferRef.OffsetInBytes,
2044  BufferRef.IsSubBuffer, PropertyList) {
2045  preScreenAccessor(BufferRef.size(), PropertyList);
2046  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2047  BufferRef.get_range()))
2048  throw sycl::invalid_object_error(
2049  "accessor with requested offset and range would exceed the bounds of "
2050  "the buffer",
2051  PI_ERROR_INVALID_VALUE);
2052 
2053  initHostAcc();
2054  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
2056  detail::AccessorBaseHost::impl.get(),
2057  AccessTarget, AccessMode, CodeLoc);
2058  GDBMethodsAnchor();
2059  }
2060 #endif
2061 
2062  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2063  typename TagT,
2064  typename = std::enable_if_t<
2065  detail::IsRunTimePropertyListT<PropertyListT>::value &&
2066  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2067  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2068  accessor(
2069  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2070  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
2071  const property_list &PropertyList = {},
2072  const detail::code_location CodeLoc = detail::code_location::current())
2073  : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2074  PropertyList, CodeLoc) {
2075  adjustAccPropsInBuf(BufferRef);
2076  }
2077 
2078  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2079  typename TagT, typename... PropTypes,
2080  typename = std::enable_if_t<
2081  detail::IsCxPropertyList<PropertyListT>::value &&
2082  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2083  (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2084  accessor(
2085  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2086  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
2087  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2088  {},
2089  const detail::code_location CodeLoc = detail::code_location::current())
2090  : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2091  PropertyList, CodeLoc) {
2092  adjustAccPropsInBuf(BufferRef);
2093  }
2094 
2095  template <typename... NewPropsT>
2096  accessor(
2097  const accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
2098  ext::oneapi::accessor_property_list<NewPropsT...>> &Other,
2099  const detail::code_location CodeLoc = detail::code_location::current())
2100 #ifdef __SYCL_DEVICE_ONLY__
2101  : impl(Other.impl), MData(Other.MData)
2102 #else
2103  : detail::AccessorBaseHost(Other), MAccData(Other.MAccData)
2104 #endif
2105  {
2106  static_assert(detail::IsCxPropertyList<PropertyListT>::value,
2107  "Conversion is only available for accessor_property_list");
2108  static_assert(
2109  PropertyListT::template areSameCompileTimeProperties<NewPropsT...>(),
2110  "Compile-time-constant properties must be the same");
2111 #ifndef __SYCL_DEVICE_ONLY__
2112  detail::constructorNotification(getMemoryObject(), impl.get(), AccessTarget,
2113  AccessMode, CodeLoc);
2114 #endif
2115  }
2116 
2117  void swap(accessor &other) {
2118  std::swap(impl, other.impl);
2119 #ifdef __SYCL_DEVICE_ONLY__
2120  std::swap(MData, other.MData);
2121 #else
2122  std::swap(MAccData, other.MAccData);
2123 #endif
2124  }
2125 
2126  bool is_placeholder() const {
2127 #ifdef __SYCL_DEVICE_ONLY__
2128  return false;
2129 #else
2130  return detail::AccessorBaseHost::isPlaceholder();
2131 #endif
2132  }
2133 
2134  size_t get_size() const { return getAccessRange().size() * sizeof(DataT); }
2135 
2136  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
2137  size_t get_count() const { return size(); }
2138  size_type size() const noexcept { return getAccessRange().size(); }
2139 
2140  size_type byte_size() const noexcept { return size() * sizeof(DataT); }
2141 
2142  size_type max_size() const noexcept {
2143  return empty() ? 0 : (std::numeric_limits<difference_type>::max)();
2144  }
2145 
2146  bool empty() const noexcept { return size() == 0; }
2147 
2148  template <int Dims = Dimensions,
2149  typename = std::enable_if_t<Dims == Dimensions && (Dims > 0)>>
2150  range<Dimensions> get_range() const {
2151  return getRange<Dims>();
2152  }
2153 
2154  template <int Dims = Dimensions,
2155  typename = std::enable_if_t<Dims == Dimensions && (Dims > 0)>>
2156  id<Dimensions> get_offset() const {
2157  return getOffset<Dims>();
2158  }
2159 
2160  template <int Dims = Dimensions, typename RefT = RefType,
2161  typename = std::enable_if_t<Dims == 0 &&
2162  (IsAccessAnyWrite || IsAccessReadOnly)>>
2163  operator reference() const {
2164  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
2165  return *(getQualifiedPtr() + LinearIndex);
2166  }
2167 
2168  template <int Dims = Dimensions,
2169  typename = std::enable_if_t<AccessMode != access_mode::atomic &&
2170  !IsAccessReadOnly && Dims == 0>>
2171  const accessor &operator=(const value_type &Other) const {
2172  *getQualifiedPtr() = Other;
2173  return *this;
2174  }
2175 
2176  template <int Dims = Dimensions,
2177  typename = std::enable_if_t<AccessMode != access_mode::atomic &&
2178  !IsAccessReadOnly && Dims == 0>>
2179  const accessor &operator=(value_type &&Other) const {
2180  *getQualifiedPtr() = std::move(Other);
2181  return *this;
2182  }
2183 
2184  template <int Dims = Dimensions,
2185  typename = std::enable_if_t<(Dims > 0) &&
2186  (IsAccessAnyWrite || IsAccessReadOnly)>>
2187  reference operator[](id<Dimensions> Index) const {
2188  const size_t LinearIndex = getLinearIndex(Index);
2189  return getQualifiedPtr()[LinearIndex];
2190  }
2191 
2192  template <int Dims = Dimensions>
2193  operator typename std::enable_if_t<Dims == 0 &&
2194  AccessMode == access::mode::atomic,
2195 #ifdef __ENABLE_USM_ADDR_SPACE__
2196  atomic<DataT>
2197 #else
2198  atomic<DataT, AS>
2199 #endif
2200  >() const {
2201  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
2202  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2203  getQualifiedPtr() + LinearIndex));
2204  }
2205 
2206  template <int Dims = Dimensions>
2207  typename std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
2208  atomic<DataT, AS>>
2209  operator[](id<Dimensions> Index) const {
2210  const size_t LinearIndex = getLinearIndex(Index);
2211  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2212  getQualifiedPtr() + LinearIndex));
2213  }
2214 
2215  template <int Dims = Dimensions>
2216  typename std::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
2217  atomic<DataT, AS>>
2218  operator[](size_t Index) const {
2219  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
2220  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2221  getQualifiedPtr() + LinearIndex));
2222  }
2223  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 1)>>
2224  auto operator[](size_t Index) const {
2225  return AccessorSubscript<Dims - 1>(*this, Index);
2226  }
2227 
2228  template <access::target AccessTarget_ = AccessTarget,
2229  typename = std::enable_if_t<
2230  (AccessTarget_ == access::target::host_buffer) ||
2231  (AccessTarget_ == access::target::host_task) ||
2232  (AccessTarget_ == access::target::device)>>
2233  std::add_pointer_t<value_type> get_pointer() const noexcept {
2234  return getPointerAdjusted();
2235  }
2236 
2237  template <access::target AccessTarget_ = AccessTarget,
2238  typename = std::enable_if_t<AccessTarget_ ==
2239  access::target::constant_buffer>>
2240  constant_ptr<DataT> get_pointer() const {
2241  return constant_ptr<DataT>(getPointerAdjusted());
2242  }
2243 
2244  template <access::decorated IsDecorated>
2245  accessor_ptr<IsDecorated> get_multi_ptr() const noexcept {
2246  return accessor_ptr<IsDecorated>(getPointerAdjusted());
2247  }
2248 
2249  // accessor::has_property for runtime properties is only available in host
2250  // code. This restriction is not listed in the core spec and will be added in
2251  // future versions.
2252  template <typename Property>
2253  typename std::enable_if_t<
2254  !ext::oneapi::is_compile_time_property<Property>::value, bool>
2255  has_property() const noexcept {
2256 #ifndef __SYCL_DEVICE_ONLY__
2257  return getPropList().template has_property<Property>();
2258 #else
2259  return false;
2260 #endif
2261  }
2262 
2263  // accessor::get_property for runtime properties is only available in host
2264  // code. This restriction is not listed in the core spec and will be added in
2265  // future versions.
2266  template <typename Property,
2267  typename = typename std::enable_if_t<
2268  !ext::oneapi::is_compile_time_property<Property>::value>>
2269  Property get_property() const {
2270 #ifndef __SYCL_DEVICE_ONLY__
2271  return getPropList().template get_property<Property>();
2272 #else
2273  return Property();
2274 #endif
2275  }
2276 
2277  template <typename Property>
2278  static constexpr bool has_property(
2279  typename std::enable_if_t<
2280  ext::oneapi::is_compile_time_property<Property>::value> * = 0) {
2281  return PropertyListT::template has_property<Property>();
2282  }
2283 
2284  template <typename Property>
2285  static constexpr auto get_property(
2286  typename std::enable_if_t<
2287  ext::oneapi::is_compile_time_property<Property>::value> * = 0) {
2288  return PropertyListT::template get_property<Property>();
2289  }
2290 
2291  bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
2292  bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
2293 
2294  iterator begin() const noexcept {
2295  return iterator::getBegin(
2296  get_pointer(),
2297  detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2298  getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2299  }
2300 
2301  iterator end() const noexcept {
2302  return iterator::getEnd(
2303  get_pointer(),
2304  detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2305  getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2306  }
2307 
2308  const_iterator cbegin() const noexcept {
2309  return const_iterator::getBegin(
2310  get_pointer(),
2311  detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2312  getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2313  }
2314 
2315  const_iterator cend() const noexcept {
2316  return const_iterator::getEnd(
2317  get_pointer(),
2318  detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2319  getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2320  }
2321 
2322  reverse_iterator rbegin() const noexcept { return reverse_iterator(end()); }
2323  reverse_iterator rend() const noexcept { return reverse_iterator(begin()); }
2324 
2325  const_reverse_iterator crbegin() const noexcept {
2326  return const_reverse_iterator(cend());
2327  }
2328  const_reverse_iterator crend() const noexcept {
2329  return const_reverse_iterator(cbegin());
2330  }
2331 
2332 private:
2333  template <int Dims, typename = std::enable_if_t<(Dims > 0)>>
2334  range<Dims> getRange() const {
2335  if constexpr (Dimensions == 0)
2336  return range<1>{1};
2337  else
2338  return detail::convertToArrayOfN<Dims, 1>(getAccessRange());
2339  }
2340 
2341  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
2342  id<Dims> getOffset() const {
2343  static_assert(
2344  !(PropertyListT::template has_property<
2346  "Accessor has no_offset property, get_offset() can not be used");
2347  return detail::convertToArrayOfN<Dims, 0>(getOffset());
2348  }
2349 
2350 #ifdef __SYCL_DEVICE_ONLY__
2351  size_t getTotalOffset() const noexcept {
2352  size_t TotalOffset = 0;
2353  detail::loop<Dimensions>([&, this](size_t I) {
2354  TotalOffset = TotalOffset * impl.MemRange[I];
2355  if constexpr (!(PropertyListT::template has_property<
2357  TotalOffset += impl.Offset[I];
2358  }
2359  });
2360 
2361  return TotalOffset;
2362  }
2363 #endif
2364 
2365  // supporting function for get_pointer()
2366  // MData has been preadjusted with offset for faster access with []
2367  // but for get_pointer() we must return the original pointer.
2368  // On device, getQualifiedPtr() returns MData, so we need to backjust it.
2369  // On host, getQualifiedPtr() does not return MData, no need to adjust.
2370  auto getPointerAdjusted() const noexcept {
2371 #ifdef __SYCL_DEVICE_ONLY__
2372  return getQualifiedPtr() - getTotalOffset();
2373 #else
2374  return getQualifiedPtr();
2375 #endif
2376  }
2377 
2378  void preScreenAccessor(const size_t elemInBuffer,
2379  const PropertyListT &PropertyList) {
2380  // check device accessor buffer size
2381  if (!IsHostBuf && elemInBuffer == 0)
2382  throw sycl::invalid_object_error(
2383  "SYCL buffer size is zero. To create a device accessor, SYCL "
2384  "buffer size must be greater than zero.",
2385  PI_ERROR_INVALID_VALUE);
2386 
2387  // check that no_init property is compatible with access mode
2388  if (PropertyList.template has_property<property::no_init>() &&
2389  AccessMode == access::mode::read) {
2390  throw sycl::invalid_object_error(
2391  "accessor would cannot be both read_only and no_init",
2392  PI_ERROR_INVALID_VALUE);
2393  }
2394  }
2395 
2396  template <typename BufT, typename... PropTypes>
2397  void adjustAccPropsInBuf(BufT &Buffer) {
2398  if constexpr (PropertyListT::template has_property<
2400  auto location = (PropertyListT::template get_property<
2402  .get_location();
2403  property_list PropList{
2405  Buffer.addOrReplaceAccessorProperties(PropList);
2406  } else {
2407  deleteAccPropsFromBuf(Buffer);
2408  }
2409  }
2410 
2411  template <typename BufT> void deleteAccPropsFromBuf(BufT &Buffer) {
2412  Buffer.deleteAccProps(
2414  }
2415 };
2416 
2417 template <typename DataT, int Dimensions, typename AllocatorT>
2418 accessor(buffer<DataT, Dimensions, AllocatorT>)
2419  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2420  access::placeholder::true_t>;
2421 
2422 template <typename DataT, int Dimensions, typename AllocatorT,
2423  typename... PropsT>
2424 accessor(buffer<DataT, Dimensions, AllocatorT>,
2425  const ext::oneapi::accessor_property_list<PropsT...> &)
2426  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2427  access::placeholder::true_t,
2428  ext::oneapi::accessor_property_list<PropsT...>>;
2429 
2430 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
2431 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1)
2432  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
2433  detail::deduceAccessTarget<Type1, Type1>(target::device),
2434  access::placeholder::true_t>;
2435 
2436 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2437  typename... PropsT>
2438 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1,
2439  const ext::oneapi::accessor_property_list<PropsT...> &)
2440  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
2441  detail::deduceAccessTarget<Type1, Type1>(target::device),
2442  access::placeholder::true_t,
2443  ext::oneapi::accessor_property_list<PropsT...>>;
2444 
2445 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2446  typename Type2>
2447 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2)
2448  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
2449  detail::deduceAccessTarget<Type1, Type2>(target::device),
2450  access::placeholder::true_t>;
2451 
2452 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2453  typename Type2, typename... PropsT>
2454 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2,
2455  const ext::oneapi::accessor_property_list<PropsT...> &)
2456  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
2457  detail::deduceAccessTarget<Type1, Type2>(target::device),
2458  access::placeholder::true_t,
2459  ext::oneapi::accessor_property_list<PropsT...>>;
2460 
2461 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2462  typename Type2, typename Type3>
2463 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3)
2464  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
2465  detail::deduceAccessTarget<Type2, Type3>(target::device),
2466  access::placeholder::true_t>;
2467 
2468 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2469  typename Type2, typename Type3, typename... PropsT>
2470 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3,
2471  const ext::oneapi::accessor_property_list<PropsT...> &)
2472  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
2473  detail::deduceAccessTarget<Type2, Type3>(target::device),
2474  access::placeholder::true_t,
2475  ext::oneapi::accessor_property_list<PropsT...>>;
2476 
2477 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2478  typename Type2, typename Type3, typename Type4>
2479 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
2480  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2481  detail::deduceAccessTarget<Type3, Type4>(target::device),
2482  access::placeholder::true_t>;
2483 
2484 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2485  typename Type2, typename Type3, typename Type4, typename... PropsT>
2486 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
2487  const ext::oneapi::accessor_property_list<PropsT...> &)
2488  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2489  detail::deduceAccessTarget<Type3, Type4>(target::device),
2490  access::placeholder::true_t,
2491  ext::oneapi::accessor_property_list<PropsT...>>;
2492 
2493 template <typename DataT, int Dimensions, typename AllocatorT>
2494 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &)
2495  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2496  access::placeholder::false_t>;
2497 
2498 template <typename DataT, int Dimensions, typename AllocatorT,
2499  typename... PropsT>
2500 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &,
2501  const ext::oneapi::accessor_property_list<PropsT...> &)
2502  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2503  access::placeholder::false_t,
2504  ext::oneapi::accessor_property_list<PropsT...>>;
2505 
2506 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
2507 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1)
2508  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
2509  detail::deduceAccessTarget<Type1, Type1>(target::device),
2510  access::placeholder::false_t>;
2511 
2512 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2513  typename... PropsT>
2514 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1,
2515  const ext::oneapi::accessor_property_list<PropsT...> &)
2516  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
2517  detail::deduceAccessTarget<Type1, Type1>(target::device),
2518  access::placeholder::false_t,
2519  ext::oneapi::accessor_property_list<PropsT...>>;
2520 
2521 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2522  typename Type2>
2523 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2)
2524  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
2525  detail::deduceAccessTarget<Type1, Type2>(target::device),
2526  access::placeholder::false_t>;
2527 
2528 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2529  typename Type2, typename... PropsT>
2530 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2,
2531  const ext::oneapi::accessor_property_list<PropsT...> &)
2532  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
2533  detail::deduceAccessTarget<Type1, Type2>(target::device),
2534  access::placeholder::false_t,
2535  ext::oneapi::accessor_property_list<PropsT...>>;
2536 
2537 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2538  typename Type2, typename Type3>
2539 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2, Type3)
2540  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
2541  detail::deduceAccessTarget<Type2, Type3>(target::device),
2542  access::placeholder::false_t>;
2543 
2544 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2545  typename Type2, typename Type3, typename... PropsT>
2546 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2, Type3,
2547  const ext::oneapi::accessor_property_list<PropsT...> &)
2548  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
2549  detail::deduceAccessTarget<Type2, Type3>(target::device),
2550  access::placeholder::false_t,
2551  ext::oneapi::accessor_property_list<PropsT...>>;
2552 
2553 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2554  typename Type2, typename Type3, typename Type4>
2555 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2, Type3,
2556  Type4)
2557  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2558  detail::deduceAccessTarget<Type3, Type4>(target::device),
2559  access::placeholder::false_t>;
2560 
2561 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2562  typename Type2, typename Type3, typename Type4, typename... PropsT>
2563 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2, Type3,
2564  Type4, const ext::oneapi::accessor_property_list<PropsT...> &)
2565  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2566  detail::deduceAccessTarget<Type3, Type4>(target::device),
2567  access::placeholder::false_t,
2568  ext::oneapi::accessor_property_list<PropsT...>>;
2569 
2573 template <typename DataT, int Dimensions, access::mode AccessMode,
2576 #ifndef __SYCL_DEVICE_ONLY__
2578 #endif
2579  public detail::accessor_common<DataT, Dimensions, AccessMode,
2580  access::target::local, IsPlaceholder> {
2581 protected:
2582  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
2583 
2584  using AccessorCommonT =
2586  access::target::local, IsPlaceholder>;
2587 
2588  using AccessorCommonT::AS;
2589 
2590  // Cannot do "using AccessorCommonT::Flag" as it doesn't work with g++ as host
2591  // compiler, for some reason.
2592  static constexpr bool IsAccessAnyWrite = AccessorCommonT::IsAccessAnyWrite;
2593  static constexpr bool IsAccessReadOnly = AccessorCommonT::IsAccessReadOnly;
2594  static constexpr bool IsConst = AccessorCommonT::IsConst;
2595 
2596  template <int Dims>
2597  using AccessorSubscript =
2598  typename AccessorCommonT::template AccessorSubscript<
2599  Dims,
2601 
2603 
2606 
2607 #ifdef __SYCL_DEVICE_ONLY__
2609 
2610  sycl::range<AdjustedDim> &getSize() { return impl.MemRange; }
2611  const sycl::range<AdjustedDim> &getSize() const { return impl.MemRange; }
2612 
2613  void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
2614  range<AdjustedDim>, id<AdjustedDim>) {
2615  MData = Ptr;
2616  detail::loop<AdjustedDim>(
2617  [&, this](size_t I) { getSize()[I] = AccessRange[I]; });
2618  }
2619 
2620  // __init variant used by the device compiler for ESIMD kernels.
2621  // TODO: In ESIMD accessors usage is limited for now - access range, mem
2622  // range and offset are not supported.
2623  void __init_esimd(ConcreteASPtrType Ptr) {
2624  MData = Ptr;
2625  detail::loop<AdjustedDim>([&, this](size_t I) { getSize()[I] = 0; });
2626  }
2627 
2628 public:
2629  // Default constructor for objects later initialized with __init member.
2630  local_accessor_base()
2631  : impl(detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
2632 
2633 protected:
2634  ConcreteASPtrType getQualifiedPtr() const { return MData; }
2635 
2636  ConcreteASPtrType MData;
2637 
2638 #else
2639 public:
2641  : detail::LocalAccessorBaseHost{/*Size*/ sycl::range<3>{0, 0, 0},
2642  /*Dims*/ 0, /*ElemSize*/ sizeof(DataT)} {}
2643 
2644 protected:
2646  : detail::LocalAccessorBaseHost{Impl} {}
2647 
2648  char padding[sizeof(detail::LocalAccessorBaseDevice<AdjustedDim>) +
2650  using detail::LocalAccessorBaseHost::getSize;
2651 
2653  return reinterpret_cast<PtrType>(LocalAccessorBaseHost::getPtr());
2654  }
2655 
2656  void *getPtr() { return detail::LocalAccessorBaseHost::getPtr(); }
2657  void *getPtr() const { return detail::LocalAccessorBaseHost::getPtr(); }
2658  const range<3> &getSize() const {
2659  return detail::LocalAccessorBaseHost::getSize();
2660  }
2661  range<3> &getSize() { return detail::LocalAccessorBaseHost::getSize(); }
2662 
2663  // The function references helper methods required by GDB pretty-printers
2665 #ifndef NDEBUG
2666  const auto *this_const = this;
2667  (void)getSize();
2668  (void)this_const->getSize();
2669  (void)getPtr();
2670  (void)this_const->getPtr();
2671 #endif
2672  }
2673 
2674 #endif // __SYCL_DEVICE_ONLY__
2675 
2676  // Method which calculates linear offset for the ID using Range and Offset.
2677  template <int Dims = AdjustedDim> size_t getLinearIndex(id<Dims> Id) const {
2678  size_t Result = 0;
2679  detail::loop<Dims>(
2680  [&, this](size_t I) { Result = Result * getSize()[I] + Id[I]; });
2681  return Result;
2682  }
2683 
2684  template <class Obj>
2685  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
2686 
2687  template <class T>
2688  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
2689 
2690  template <typename DataT_, int Dimensions_> friend class local_accessor;
2691 
2692 public:
2693  using value_type = DataT;
2694  using reference = DataT &;
2695  using const_reference = const DataT &;
2696 
2697  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 0>>
2699  detail::code_location::current())
2700 #ifdef __SYCL_DEVICE_ONLY__
2701  : impl(range<AdjustedDim>{1}){}
2702 #else
2703  : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT)) {
2704  detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
2705  access::target::local, AccessMode, CodeLoc);
2706  GDBMethodsAnchor();
2707  }
2708 #endif
2709 
2710  template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 0>>
2712  const detail::code_location CodeLoc =
2713  detail::code_location::current())
2714 #ifdef __SYCL_DEVICE_ONLY__
2715  : impl(range<AdjustedDim>{1}) {
2716  (void)propList;
2717  }
2718 #else
2719  : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT),
2720  propList) {
2721  detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
2722  access::target::local, AccessMode, CodeLoc);
2723  GDBMethodsAnchor();
2724  }
2725 #endif
2726 
2727  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
2729  range<Dimensions> AllocationSize, handler &,
2730  const detail::code_location CodeLoc = detail::code_location::current())
2731 #ifdef __SYCL_DEVICE_ONLY__
2732  : impl(AllocationSize){}
2733 #else
2734  : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2735  AdjustedDim, sizeof(DataT)) {
2736  detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
2737  access::target::local, AccessMode, CodeLoc);
2738  GDBMethodsAnchor();
2739  }
2740 #endif
2741 
2742  template <int Dims = Dimensions,
2743  typename = std::enable_if_t<(Dims > 0)>>
2745  const property_list &propList,
2746  const detail::code_location CodeLoc =
2747  detail::code_location::current())
2748 #ifdef __SYCL_DEVICE_ONLY__
2749  : impl(AllocationSize) {
2750  (void)propList;
2751  }
2752 #else
2753  : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2754  AdjustedDim, sizeof(DataT), propList) {
2755  detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
2756  access::target::local, AccessMode, CodeLoc);
2757  GDBMethodsAnchor();
2758  }
2759 #endif
2760 
2761  size_t get_size() const { return getSize().size() * sizeof(DataT); }
2762 
2763  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
2764  size_t get_count() const { return size(); }
2765  size_t size() const noexcept { return getSize().size(); }
2766 
2767  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
2769  return detail::convertToArrayOfN<Dims, 1>(getSize());
2770  }
2771 
2772  template <int Dims = Dimensions,
2773  typename = std::enable_if_t<Dims == 0 &&
2774  (IsAccessAnyWrite || IsAccessReadOnly)>>
2775  operator RefType() const {
2776  return *getQualifiedPtr();
2777  }
2778 
2779  template <int Dims = Dimensions,
2780  typename = std::enable_if_t<(Dims > 0) &&
2781  (IsAccessAnyWrite || IsAccessReadOnly)>>
2783  const size_t LinearIndex = getLinearIndex(Index);
2784  return getQualifiedPtr()[LinearIndex];
2785  }
2786 
2787  template <int Dims = Dimensions,
2788  typename = std::enable_if_t<Dims == 1 &&
2789  (IsAccessAnyWrite || IsAccessReadOnly)>>
2790  RefType operator[](size_t Index) const {
2791  return getQualifiedPtr()[Index];
2792  }
2793 
2794  template <int Dims = Dimensions>
2795  operator typename std::enable_if_t<
2796  Dims == 0 && AccessMode == access::mode::atomic, atomic<DataT, AS>>()
2797  const {
2798  return atomic<DataT, AS>(
2800  }
2801 
2802  template <int Dims = Dimensions>
2803  typename std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
2804  atomic<DataT, AS>>
2805  operator[](id<Dimensions> Index) const {
2806  const size_t LinearIndex = getLinearIndex(Index);
2807  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2808  getQualifiedPtr() + LinearIndex));
2809  }
2810 
2811  template <int Dims = Dimensions>
2812  typename std::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
2813  atomic<DataT, AS>>
2814  operator[](size_t Index) const {
2815  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2816  getQualifiedPtr() + Index));
2817  }
2818 
2819  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 1)>>
2820  typename AccessorCommonT::template AccessorSubscript<
2821  Dims - 1,
2823  operator[](size_t Index) const {
2824  return AccessorSubscript<Dims - 1>(*this, Index);
2825  }
2826 
2827  bool operator==(const local_accessor_base &Rhs) const {
2828  return impl == Rhs.impl;
2829  }
2830  bool operator!=(const local_accessor_base &Rhs) const {
2831  return !(*this == Rhs);
2832  }
2833 };
2834 
2835 // TODO: Remove deprecated specialization once no longer needed
2836 template <typename DataT, int Dimensions, access::mode AccessMode,
2839  DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder>
2840  : public local_accessor_base<DataT, Dimensions, AccessMode, IsPlaceholder>,
2841  public detail::OwnerLessBase<
2842  accessor<DataT, Dimensions, AccessMode, access::target::local,
2843  IsPlaceholder>> {
2844 
2845  using local_acc =
2847 
2848  static_assert(
2849  !local_acc::IsConst || local_acc::IsAccessReadOnly,
2850  "A const qualified DataT is only allowed for a read-only accessor");
2851 
2852  // Use base classes constructors
2853  using local_acc::local_acc;
2854 
2855 public:
2857  return local_ptr<DataT>(local_acc::getQualifiedPtr());
2858  }
2859 
2860 #ifdef __SYCL_DEVICE_ONLY__
2861 
2862  // __init needs to be defined within the class not through inheritance.
2863  // Map this function to inherited func.
2864  void __init(typename local_acc::ConcreteASPtrType Ptr,
2865  range<local_acc::AdjustedDim> AccessRange,
2868  local_acc::__init(Ptr, AccessRange, range, id);
2869  }
2870 
2871  // __init variant used by the device compiler for ESIMD kernels.
2872  // TODO: In ESIMD accessors usage is limited for now - access range, mem
2873  // range and offset are not supported.
2874  void __init_esimd(typename local_acc::ConcreteASPtrType Ptr) {
2875  local_acc::__init_esimd(Ptr);
2876  }
2877 
2878 public:
2879  // Default constructor for objects later initialized with __init member.
2880  accessor() {
2881  local_acc::impl = detail::InitializedVal<local_acc::AdjustedDim,
2882  range>::template get<0>();
2883  }
2884 
2885 #else
2886 private:
2887  accessor(const detail::AccessorImplPtr &Impl) : local_acc{Impl} {}
2888 #endif
2889 };
2890 
2891 template <typename DataT, int Dimensions = 1>
2892 class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
2893  : public local_accessor_base<DataT, Dimensions,
2894  detail::accessModeFromConstness<DataT>(),
2895  access::placeholder::false_t>,
2896  public detail::OwnerLessBase<local_accessor<DataT, Dimensions>> {
2897 
2898  using local_acc =
2899  local_accessor_base<DataT, Dimensions,
2900  detail::accessModeFromConstness<DataT>(),
2901  access::placeholder::false_t>;
2902 
2903  static_assert(
2904  !local_acc::IsConst || local_acc::IsAccessReadOnly,
2905  "A const qualified DataT is only allowed for a read-only accessor");
2906 
2907  // Use base classes constructors
2908  using local_acc::local_acc;
2909 
2910 #ifdef __SYCL_DEVICE_ONLY__
2911 
2912  // __init needs to be defined within the class not through inheritance.
2913  // Map this function to inherited func.
2914  void __init(typename local_acc::ConcreteASPtrType Ptr,
2915  range<local_acc::AdjustedDim> AccessRange,
2916  range<local_acc::AdjustedDim> range,
2917  id<local_acc::AdjustedDim> id) {
2918  local_acc::__init(Ptr, AccessRange, range, id);
2919  }
2920 
2921  // __init variant used by the device compiler for ESIMD kernels.
2922  // TODO: In ESIMD accessors usage is limited for now - access range, mem
2923  // range and offset are not supported.
2924  void __init_esimd(typename local_acc::ConcreteASPtrType Ptr) {
2925  local_acc::__init_esimd(Ptr);
2926  }
2927 
2928 public:
2929  // Default constructor for objects later initialized with __init member.
2930  local_accessor() {
2931  local_acc::impl = detail::InitializedVal<local_acc::AdjustedDim,
2932  range>::template get<0>();
2933  }
2934 
2935 #else
2936  local_accessor(const detail::AccessorImplPtr &Impl) : local_acc{Impl} {}
2937 #endif
2938 
2939  // implicit conversion between non-const read-write accessor to const
2940  // read-only accessor
2941 public:
2942  template <typename DataT_,
2943  typename = std::enable_if_t<
2944  std::is_const_v<DataT> &&
2945  std::is_same_v<DataT_, std::remove_const_t<DataT>>>>
2946  local_accessor(const local_accessor<DataT_, Dimensions> &other) {
2947  local_acc::impl = other.impl;
2948  }
2949 
2950  using value_type = DataT;
2951  using iterator = value_type *;
2952  using const_iterator = const value_type *;
2953  using reverse_iterator = std::reverse_iterator<iterator>;
2954  using const_reverse_iterator = std::reverse_iterator<const_iterator>;
2955  using difference_type =
2956  typename std::iterator_traits<iterator>::difference_type;
2957  using size_type = std::size_t;
2958 
2959  template <access::decorated IsDecorated>
2960  using accessor_ptr = local_ptr<value_type, IsDecorated>;
2961 
2962  template <typename DataT_>
2963  bool operator==(const local_accessor<DataT_, Dimensions> &Rhs) const {
2964  return local_acc::impl == Rhs.impl;
2965  }
2966 
2967  template <typename DataT_>
2968  bool operator!=(const local_accessor<DataT_, Dimensions> &Rhs) const {
2969  return !(*this == Rhs);
2970  }
2971 
2972  void swap(local_accessor &other) { std::swap(this->impl, other.impl); }
2973 
2974  size_type byte_size() const noexcept { return this->size() * sizeof(DataT); }
2975 
2976  size_type max_size() const noexcept {
2977  return empty() ? 0 : (std::numeric_limits<difference_type>::max)();
2978  }
2979 
2980  bool empty() const noexcept { return this->size() == 0; }
2981 
2982  iterator begin() const noexcept {
2983  if constexpr (Dimensions == 0)
2984  return local_acc::getQualifiedPtr();
2985  else
2986  return &this->operator[](id<Dimensions>());
2987  }
2988  iterator end() const noexcept {
2989  if constexpr (Dimensions == 0)
2990  return begin() + 1;
2991  else
2992  return begin() + this->size();
2993  }
2994 
2995  const_iterator cbegin() const noexcept { return const_iterator(begin()); }
2996  const_iterator cend() const noexcept { return const_iterator(end()); }
2997 
2998  reverse_iterator rbegin() const noexcept { return reverse_iterator(end()); }
2999  reverse_iterator rend() const noexcept { return reverse_iterator(begin()); }
3000 
3001  const_reverse_iterator crbegin() const noexcept {
3002  return const_reverse_iterator(end());
3003  }
3004  const_reverse_iterator crend() const noexcept {
3005  return const_reverse_iterator(begin());
3006  }
3007 
3008  std::add_pointer_t<value_type> get_pointer() const noexcept {
3009  return std::add_pointer_t<value_type>(local_acc::getQualifiedPtr());
3010  }
3011 
3012  template <access::decorated IsDecorated>
3013  accessor_ptr<IsDecorated> get_multi_ptr() const noexcept {
3014  return accessor_ptr<IsDecorated>(local_acc::getQualifiedPtr());
3015  }
3016 
3017  template <typename Property> bool has_property() const noexcept {
3018 #ifndef __SYCL_DEVICE_ONLY__
3019  return this->getPropList().template has_property<Property>();
3020 #else
3021  return false;
3022 #endif
3023  }
3024 
3025  template <typename Property> Property get_property() const {
3026 #ifndef __SYCL_DEVICE_ONLY__
3027  return this->getPropList().template get_property<Property>();
3028 #else
3029  return Property();
3030 #endif
3031  }
3032 
3033  template <int Dims = Dimensions,
3034  typename = std::enable_if_t<!std::is_const_v<DataT> && Dims == 0>>
3035  const local_accessor &operator=(const value_type &Other) const {
3036  *local_acc::getQualifiedPtr() = Other;
3037  return *this;
3038  }
3039 
3040  template <int Dims = Dimensions,
3041  typename = std::enable_if_t<!std::is_const_v<DataT> && Dims == 0>>
3042  const local_accessor &operator=(value_type &&Other) const {
3043  *local_acc::getQualifiedPtr() = std::move(Other);
3044  return *this;
3045  }
3046 
3047 private:
3048  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
3049 };
3050 
3056 template <typename DataT, int Dimensions, access::mode AccessMode,
3059  DataT, Dimensions, AccessMode, access::target::image, IsPlaceholder>
3060  : public detail::image_accessor<DataT, Dimensions, AccessMode,
3061  access::target::image, IsPlaceholder>,
3062  public detail::OwnerLessBase<
3063  accessor<DataT, Dimensions, AccessMode, access::target::image,
3064  IsPlaceholder>> {
3065 private:
3066  accessor(const detail::AccessorImplPtr &Impl)
3068  access::target::image, IsPlaceholder>{Impl} {}
3069 
3070 public:
3071  template <typename AllocatorT>
3072  accessor(sycl::image<Dimensions, AllocatorT> &Image,
3073  handler &CommandGroupHandler)
3075  access::target::image, IsPlaceholder>(
3076  Image, CommandGroupHandler, Image.getElementSize()) {
3077 #ifndef __SYCL_DEVICE_ONLY__
3078  detail::associateWithHandler(CommandGroupHandler, this,
3079  access::target::image);
3080 #endif
3081  }
3082 
3083  template <typename AllocatorT>
3084  accessor(sycl::image<Dimensions, AllocatorT> &Image,
3085  handler &CommandGroupHandler, const property_list &propList)
3087  access::target::image, IsPlaceholder>(
3088  Image, CommandGroupHandler, Image.getElementSize()) {
3089  (void)propList;
3090 #ifndef __SYCL_DEVICE_ONLY__
3091  detail::associateWithHandler(CommandGroupHandler, this,
3092  access::target::image);
3093 #endif
3094  }
3095 #ifdef __SYCL_DEVICE_ONLY__
3096 private:
3097  using OCLImageTy =
3098  typename detail::opencl_image_type<Dimensions, AccessMode,
3099  access::target::image>::type;
3100 
3101  // Front End requires this method to be defined in the accessor class.
3102  // It does not call the base class's init method.
3103  void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
3104 
3105  // __init variant used by the device compiler for ESIMD kernels.
3106  void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
3107 
3108 public:
3109  // Default constructor for objects later initialized with __init member.
3110  accessor() = default;
3111 #endif
3112 };
3113 
3121 template <typename DataT, int Dimensions, access::mode AccessMode,
3124  access::target::host_image, IsPlaceholder>
3125  : public detail::image_accessor<DataT, Dimensions, AccessMode,
3126  access::target::host_image, IsPlaceholder>,
3127  public detail::OwnerLessBase<
3128  accessor<DataT, Dimensions, AccessMode, access::target::host_image,
3129  IsPlaceholder>> {
3130 public:
3131  template <typename AllocatorT>
3132  accessor(sycl::image<Dimensions, AllocatorT> &Image)
3133  : detail::image_accessor<DataT, Dimensions, AccessMode,
3134  access::target::host_image, IsPlaceholder>(
3135  Image, Image.getElementSize()) {}
3136 
3137  template <typename AllocatorT>
3138  accessor(sycl::image<Dimensions, AllocatorT> &Image,
3139  const property_list &propList)
3140  : detail::image_accessor<DataT, Dimensions, AccessMode,
3141  access::target::host_image, IsPlaceholder>(
3142  Image, Image.getElementSize()) {
3143  (void)propList;
3144  }
3145 };
3146 
3155 template <typename DataT, int Dimensions, access::mode AccessMode,
3158  DataT, Dimensions, AccessMode, access::target::image_array, IsPlaceholder>
3159  : public detail::image_accessor<DataT, Dimensions + 1, AccessMode,
3160  access::target::image, IsPlaceholder>,
3161  public detail::OwnerLessBase<
3162  accessor<DataT, Dimensions, AccessMode, access::target::image_array,
3163  IsPlaceholder>> {
3164 #ifdef __SYCL_DEVICE_ONLY__
3165 private:
3166  using OCLImageTy =
3167  typename detail::opencl_image_type<Dimensions + 1, AccessMode,
3168  access::target::image>::type;
3169 
3170  // Front End requires this method to be defined in the accessor class.
3171  // It does not call the base class's init method.
3172  void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
3173 
3174  // __init variant used by the device compiler for ESIMD kernels.
3175  void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
3176 
3177 public:
3178  // Default constructor for objects later initialized with __init member.
3179  accessor() = default;
3180 #endif
3181 public:
3182  template <typename AllocatorT>
3183  accessor(sycl::image<Dimensions + 1, AllocatorT> &Image,
3184  handler &CommandGroupHandler)
3186  access::target::image, IsPlaceholder>(
3187  Image, CommandGroupHandler, Image.getElementSize()) {
3188 #ifndef __SYCL_DEVICE_ONLY__
3189  detail::associateWithHandler(CommandGroupHandler, this,
3190  access::target::image_array);
3191 #endif
3192  }
3193 
3194  template <typename AllocatorT>
3195  accessor(sycl::image<Dimensions + 1, AllocatorT> &Image,
3196  handler &CommandGroupHandler, const property_list &propList)
3198  access::target::image, IsPlaceholder>(
3199  Image, CommandGroupHandler, Image.getElementSize()) {
3200  (void)propList;
3201 #ifndef __SYCL_DEVICE_ONLY__
3202  detail::associateWithHandler(CommandGroupHandler, this,
3203  access::target::image_array);
3204 #endif
3205  }
3206 
3208  operator[](size_t Index) const {
3210  IsPlaceholder>(*this, Index);
3211  }
3212 };
3213 
3214 template <typename DataT, int Dimensions = 1,
3217  : public accessor<DataT, Dimensions, AccessMode, target::host_buffer,
3218  access::placeholder::false_t> {
3219 protected:
3220  using AccessorT = accessor<DataT, Dimensions, AccessMode, target::host_buffer,
3221  access::placeholder::false_t>;
3222 
3223  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
3224  constexpr static bool IsAccessReadOnly = AccessMode == access::mode::read;
3225 
3226  template <typename T, int Dims>
3228  : std::bool_constant<std::is_same_v<T, DataT> && (Dims > 0) &&
3229  (Dims == Dimensions)> {};
3230 
3231  void
3232  __init(typename accessor<DataT, Dimensions, AccessMode, target::host_buffer,
3233  access::placeholder::false_t>::ConcreteASPtrType Ptr,
3234  range<AdjustedDim> AccessRange, range<AdjustedDim> MemRange,
3235  id<AdjustedDim> Offset) {
3236  AccessorT::__init(Ptr, AccessRange, MemRange, Offset);
3237  }
3238 
3239 #ifndef __SYCL_DEVICE_ONLY__
3241  : accessor<DataT, Dimensions, AccessMode, target::host_buffer,
3242  access::placeholder::false_t>{Impl} {}
3243 
3244  template <class Obj>
3245  friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject);
3246 
3247  template <class T>
3248  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
3249 #endif // __SYCL_DEVICE_ONLY__
3250 
3251 public:
3253 
3254  // The list of host_accessor constructors with their arguments
3255  // -------+---------+-------+----+----------+--------------
3256  // Dimensions = 0
3257  // -------+---------+-------+----+----------+--------------
3258  // buffer | | | | | property_list
3259  // buffer | handler | | | | property_list
3260  // -------+---------+-------+----+----------+--------------
3261  // Dimensions >= 1
3262  // -------+---------+-------+----+----------+--------------
3263  // buffer | | | | | property_list
3264  // buffer | | | | mode_tag | property_list
3265  // buffer | handler | | | | property_list
3266  // buffer | handler | | | mode_tag | property_list
3267  // buffer | | range | | | property_list
3268  // buffer | | range | | mode_tag | property_list
3269  // buffer | handler | range | | | property_list
3270  // buffer | handler | range | | mode_tag | property_list
3271  // buffer | | range | id | | property_list
3272  // buffer | | range | id | mode_tag | property_list
3273  // buffer | handler | range | id | | property_list
3274  // buffer | handler | range | id | mode_tag | property_list
3275  // -------+---------+-------+----+----------+--------------
3276 
3277  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3278  typename = typename std::enable_if_t<std::is_same_v<T, DataT> &&
3279  Dims == 0>>
3281  buffer<T, 1, AllocatorT> &BufferRef,
3282  const property_list &PropertyList = {},
3283  const detail::code_location CodeLoc = detail::code_location::current())
3284  : AccessorT(BufferRef, PropertyList, CodeLoc) {}
3285 
3286  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3287  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3289  buffer<T, Dims, AllocatorT> &BufferRef,
3290  const property_list &PropertyList = {},
3291  const detail::code_location CodeLoc = detail::code_location::current())
3292  : AccessorT(BufferRef, PropertyList, CodeLoc) {}
3293 
3294  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3295  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3298  const property_list &PropertyList = {},
3299  const detail::code_location CodeLoc = detail::code_location::current())
3300  : host_accessor(BufferRef, PropertyList, CodeLoc) {}
3301 
3302  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3303  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3305  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3306  const property_list &PropertyList = {},
3307  const detail::code_location CodeLoc = detail::code_location::current())
3308  : AccessorT(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
3309 
3310  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3311  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3313  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3314  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
3315  const detail::code_location CodeLoc = detail::code_location::current())
3316  : host_accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
3317 
3318  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3319  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3321  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3322  const property_list &PropertyList = {},
3323  const detail::code_location CodeLoc = detail::code_location::current())
3324  : AccessorT(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
3325 
3326  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3327  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3329  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3330  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
3331  const detail::code_location CodeLoc = detail::code_location::current())
3332  : host_accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
3333 
3334  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3335  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3337  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3338  range<Dimensions> AccessRange, const property_list &PropertyList = {},
3339  const detail::code_location CodeLoc = detail::code_location::current())
3340  : AccessorT(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
3341  CodeLoc) {}
3342 
3343  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3344  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3346  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3348  const property_list &PropertyList = {},
3349  const detail::code_location CodeLoc = detail::code_location::current())
3350  : host_accessor(BufferRef, CommandGroupHandler, AccessRange, {},
3351  PropertyList, CodeLoc) {}
3352 
3353  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3354  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3356  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3357  id<Dimensions> AccessOffset, const property_list &PropertyList = {},
3358  const detail::code_location CodeLoc = detail::code_location::current())
3359  : AccessorT(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
3360  }
3361 
3362  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3363  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3365  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3366  id<Dimensions> AccessOffset, mode_tag_t<AccessMode>,
3367  const property_list &PropertyList = {},
3368  const detail::code_location CodeLoc = detail::code_location::current())
3369  : host_accessor(BufferRef, AccessRange, AccessOffset, PropertyList,
3370  CodeLoc) {}
3371 
3372  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3373  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3375  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3376  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
3377  const property_list &PropertyList = {},
3378  const detail::code_location CodeLoc = detail::code_location::current())
3379  : AccessorT(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
3380  PropertyList, CodeLoc) {}
3381 
3382  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3383  typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3385  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3386  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
3387  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
3388  const detail::code_location CodeLoc = detail::code_location::current())
3389  : host_accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
3390  PropertyList, CodeLoc) {}
3391 
3392  template <int Dims = Dimensions,
3393  typename = std::enable_if_t<AccessMode != access_mode::atomic &&
3394  !IsAccessReadOnly && Dims == 0>>
3395  const host_accessor &
3396  operator=(const typename AccessorT::value_type &Other) const {
3397  *AccessorT::getQualifiedPtr() = Other;
3398  return *this;
3399  }
3400 
3401  template <int Dims = Dimensions,
3402  typename = std::enable_if_t<AccessMode != access_mode::atomic &&
3403  !IsAccessReadOnly && Dims == 0>>
3404  const host_accessor &operator=(typename AccessorT::value_type &&Other) const {
3405  *AccessorT::getQualifiedPtr() = std::move(Other);
3406  return *this;
3407  }
3408 
3409  // implicit conversion between const / non-const types for read only accessors
3410  template <typename DataT_,
3411  typename = std::enable_if_t<
3412  IsAccessReadOnly && !std::is_same_v<DataT_, DataT> &&
3413  std::is_same_v<std::remove_const_t<DataT_>,
3414  std::remove_const_t<DataT>>>>
3416 #ifndef __SYCL_DEVICE_ONLY__
3417  : host_accessor(other.impl)
3418 #endif // __SYCL_DEVICE_ONLY__
3419  {
3420  }
3421 
3422  // implicit conversion from read_write T accessor to read only T (const)
3423  // accessor
3424  template <typename DataT_, access::mode AccessMode_,
3425  typename = std::enable_if_t<
3426  (AccessMode_ == access_mode::read_write) && IsAccessReadOnly &&
3427  std::is_same_v<DataT_, std::remove_const_t<DataT>>>>
3429 #ifndef __SYCL_DEVICE_ONLY__
3430  : host_accessor(other.impl)
3431 #endif // __SYCL_DEVICE_ONLY__
3432  {
3433  }
3434 
3435  // host_accessor needs to explicitly define the owner_before member functions
3436  // as inheriting from OwnerLessBase causes base class conflicts.
3437  // TODO: Once host_accessor is detached from accessor, inherit from
3438  // OwnerLessBase instead.
3439 #ifndef __SYCL_DEVICE_ONLY__
3442  const noexcept {
3443  return this->impl.owner_before(
3445  }
3446 
3447  bool ext_oneapi_owner_before(const host_accessor &Other) const noexcept {
3448  return this->impl.owner_before(Other.impl);
3449  }
3450 #else
3451  bool ext_oneapi_owner_before(
3453  const noexcept;
3454  bool ext_oneapi_owner_before(const host_accessor &Other) const noexcept;
3455 #endif
3456 };
3457 
3458 template <typename DataT, int Dimensions, typename AllocatorT>
3459 host_accessor(buffer<DataT, Dimensions, AllocatorT>)
3460  -> host_accessor<DataT, Dimensions, access::mode::read_write>;
3461 
3462 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
3463 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1)
3464  -> host_accessor<DataT, Dimensions,
3465  detail::deduceAccessMode<Type1, Type1>()>;
3466 
3467 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3468  typename Type2>
3469 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2)
3470  -> host_accessor<DataT, Dimensions,
3471  detail::deduceAccessMode<Type1, Type2>()>;
3472 
3473 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3474  typename Type2, typename Type3>
3475 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3)
3476  -> host_accessor<DataT, Dimensions,
3477  detail::deduceAccessMode<Type2, Type3>()>;
3478 
3479 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3480  typename Type2, typename Type3, typename Type4>
3481 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
3482  -> host_accessor<DataT, Dimensions,
3483  detail::deduceAccessMode<Type3, Type4>()>;
3484 
3485 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3486  typename Type2, typename Type3, typename Type4, typename Type5>
3487 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
3488  Type5) -> host_accessor<DataT, Dimensions,
3489  detail::deduceAccessMode<Type4, Type5>()>;
3490 
3491 // SYCL 2020 image accessors
3492 
3493 template <typename DataT, int Dimensions, access_mode AccessMode,
3494  image_target AccessTarget = image_target::device>
3496  static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
3497  std::is_same_v<DataT, float4> ||
3498  std::is_same_v<DataT, half4>,
3499  "The data type of an image accessor must be only int4, "
3500  "uint4, float4 or half4 from SYCL namespace");
3501 
3502 public:
3503  using value_type = typename std::conditional<AccessMode == access_mode::read,
3504  const DataT, DataT>::type;
3506  using const_reference = const DataT &;
3507 
3508  template <typename AllocatorT>
3510  handler &CommandGroupHandlerRef,
3511  const property_list &PropList = {}) {
3512  device Device = detail::getDeviceFromHandler(CommandGroupHandlerRef);
3513  if (AccessTarget == image_target::device && !Device.has(aspect::image))
3514  throw sycl::exception(
3515  sycl::make_error_code(sycl::errc::feature_not_supported),
3516  "Device associated with command group handler does not have "
3517  "aspect::image.");
3518 
3519  std::ignore = ImageRef;
3520  std::ignore = PropList;
3521  throw sycl::exception(
3522  sycl::make_error_code(sycl::errc::feature_not_supported),
3523  "unsampled_image_accessor is not yet available.");
3524  }
3525 
3526  /* -- common interface members -- */
3527 
3528  unsampled_image_accessor(const unsampled_image_accessor &Rhs) = default;
3529 
3530  unsampled_image_accessor(unsampled_image_accessor &&Rhs) = default;
3531 
3532  unsampled_image_accessor &
3533  operator=(const unsampled_image_accessor &Rhs) = default;
3534 
3535  unsampled_image_accessor &operator=(unsampled_image_accessor &&Rhs) = default;
3536 
3537  ~unsampled_image_accessor() = default;
3538 
3539 #ifdef __SYCL_DEVICE_ONLY__
3540  bool operator==(const unsampled_image_accessor &Rhs) const;
3541 #else
3542  bool operator==(const unsampled_image_accessor &Rhs) const {
3543  std::ignore = Rhs;
3544  throw sycl::exception(
3545  sycl::make_error_code(sycl::errc::feature_not_supported),
3546  "operator== is not yet implemented.");
3547  }
3548 #endif // __SYCL_DEVICE_ONLY__
3549 
3550  bool operator!=(const unsampled_image_accessor &Rhs) const {
3551  return !(Rhs == *this);
3552  }
3553 
3554  /* -- property interface members -- */
3555 
3556  size_t size() const noexcept {
3557 #ifdef __SYCL_DEVICE_ONLY__
3558  return 0;
3559 #else
3560  throw sycl::exception(
3561  sycl::make_error_code(sycl::errc::feature_not_supported),
3562  "size() is not yet implemented.");
3563 #endif // __SYCL_DEVICE_ONLY__
3564  }
3565 
3566  /* Available only when: AccessMode == access_mode::read
3567  if Dimensions == 1, CoordT = int
3568  if Dimensions == 2, CoordT = int2
3569  if Dimensions == 3, CoordT = int4 */
3570  template <typename CoordT,
3571  typename = std::enable_if_t<AccessMode == access_mode::read &&
3573  Dimensions, CoordT>::value>>
3574  DataT read(const CoordT &Coords) const noexcept {
3575  std::ignore = Coords;
3576 #ifdef __SYCL_DEVICE_ONLY__
3577  return {0, 0, 0, 0};
3578 #else
3579  throw sycl::exception(
3580  sycl::make_error_code(sycl::errc::feature_not_supported),
3581  "read() is not yet implemented.");
3582 #endif // __SYCL_DEVICE_ONLY__
3583  }
3584 
3585  /* Available only when: AccessMode == access_mode::write
3586  if Dimensions == 1, CoordT = int
3587  if Dimensions == 2, CoordT = int2
3588  if Dimensions == 3, CoordT = int4 */
3589  template <typename CoordT,
3590  typename = std::enable_if_t<AccessMode == access_mode::write &&
3592  Dimensions, CoordT>::value>>
3593  void write(const CoordT &Coords, const DataT &Color) const {
3594  std::ignore = Coords;
3595  std::ignore = Color;
3596 #ifndef __SYCL_DEVICE_ONLY__
3597  throw sycl::exception(
3598  sycl::make_error_code(sycl::errc::feature_not_supported),
3599  "write() is not yet implemented.");
3600 #endif // __SYCL_DEVICE_ONLY__
3601  }
3602 };
3603 
3604 template <typename DataT, int Dimensions = 1,
3606  (std::is_const_v<DataT> ? access_mode::read
3610  public detail::OwnerLessBase<
3612  static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
3613  std::is_same_v<DataT, float4> ||
3614  std::is_same_v<DataT, half4>,
3615  "The data type of an image accessor must be only int4, "
3616  "uint4, float4 or half4 from SYCL namespace");
3617 
3619 
3620 public:
3621  using value_type = typename std::conditional<AccessMode == access_mode::read,
3622  const DataT, DataT>::type;
3624  using const_reference = const DataT &;
3625 
3626  template <typename AllocatorT>
3629  const property_list &PropList = {})
3630  : base_class(detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
3631  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
3632  Dimensions, ImageRef.getElementSize(),
3633  {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
3634  ImageRef.getChannelType(), ImageRef.getChannelOrder(),
3635  PropList) {
3636  addHostUnsampledImageAccessorAndWait(base_class::impl.get());
3637  }
3638 
3639  /* -- common interface members -- */
3640 
3641  host_unsampled_image_accessor(const host_unsampled_image_accessor &Rhs) =
3642  default;
3643 
3644  host_unsampled_image_accessor(host_unsampled_image_accessor &&Rhs) = default;
3645 
3646  host_unsampled_image_accessor &
3647  operator=(const host_unsampled_image_accessor &Rhs) = default;
3648 
3649  host_unsampled_image_accessor &
3650  operator=(host_unsampled_image_accessor &&Rhs) = default;
3651 
3652  ~host_unsampled_image_accessor() = default;
3653 
3655  return Rhs.impl == impl;
3656  }
3658  return !(Rhs == *this);
3659  }
3660 
3661  /* -- property interface members -- */
3662 
3663  size_t size() const noexcept { return base_class::getSize().size(); }
3664 
3665  /* Available only when: (AccessMode == access_mode::read ||
3666  AccessMode == access_mode::read_write)
3667  if Dimensions == 1, CoordT = int
3668  if Dimensions == 2, CoordT = int2
3669  if Dimensions == 3, CoordT = int4 */
3670  template <
3671  typename CoordT,
3672  typename = std::enable_if_t<
3673  (AccessMode == access_mode::read ||
3676  DataT read(const CoordT &Coords) const noexcept
3677 #ifdef __SYCL_DEVICE_ONLY__
3678  ;
3679 #else
3680  {
3681  // Host implementation is only available in host code. Device is not allowed
3682  // to use host_unsampled_image_accessor.
3683  image_sampler Smpl{addressing_mode::none,
3684  coordinate_normalization_mode::unnormalized,
3685  filtering_mode::nearest};
3686  return detail::imageReadSamplerHostImpl<CoordT, DataT>(
3687  Coords, Smpl, base_class::getSize(), base_class::getPitch(),
3688  base_class::getChannelType(), base_class::getChannelOrder(),
3689  base_class::getPtr(), base_class::getElementSize());
3690  }
3691 #endif
3692 
3693  /* Available only when: (AccessMode == access_mode::write ||
3694  AccessMode == access_mode::read_write)
3695  if Dimensions == 1, CoordT = int
3696  if Dimensions == 2, CoordT = int2
3697  if Dimensions == 3, CoordT = int4 */
3698  template <
3699  typename CoordT,
3700  typename = std::enable_if_t<
3704  void write(const CoordT &Coords, const DataT &Color) const
3705 #ifdef __SYCL_DEVICE_ONLY__
3706  ;
3707 #else
3708  {
3709  // Host implementation is only available in host code. Device is not allowed
3710  // to use host_unsampled_image_accessor.
3712  Coords, Color, base_class::getPitch(), base_class::getElementSize(),
3713  base_class::getChannelType(), base_class::getChannelOrder(),
3714  base_class::getPtr());
3715  }
3716 #endif
3717 
3718 private:
3721  : base_class{Impl} {}
3722 
3723  template <class Obj>
3724  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
3725 
3726  template <class T>
3727  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
3728 };
3729 
3730 template <typename DataT, int Dimensions,
3731  image_target AccessTarget = image_target::device>
3733  static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
3734  std::is_same_v<DataT, float4> ||
3735  std::is_same_v<DataT, half4>,
3736  "The data type of an image accessor must be only int4, "
3737  "uint4, float4 or half4 from SYCL namespace");
3738 
3739 public:
3740  using value_type = const DataT;
3741  using reference = const DataT &;
3742  using const_reference = const DataT &;
3743 
3744  template <typename AllocatorT>
3746  handler &CommandGroupHandlerRef,
3747  const property_list &PropList = {}) {
3748  device Device = detail::getDeviceFromHandler(CommandGroupHandlerRef);
3749  if (AccessTarget == image_target::device && !Device.has(aspect::image))
3750  throw sycl::exception(
3751  sycl::make_error_code(sycl::errc::feature_not_supported),
3752  "Device associated with command group handler does not have "
3753  "aspect::image.");
3754 
3755  std::ignore = ImageRef;
3756  std::ignore = PropList;
3757  throw sycl::exception(
3758  sycl::make_error_code(sycl::errc::feature_not_supported),
3759  "sampled_image_accessor is not yet available.");
3760  }
3761 
3762  /* -- common interface members -- */
3763 
3764  sampled_image_accessor(const sampled_image_accessor &Rhs) = default;
3765 
3766  sampled_image_accessor(sampled_image_accessor &&Rhs) = default;
3767 
3768  sampled_image_accessor &
3769  operator=(const sampled_image_accessor &Rhs) = default;
3770 
3771  sampled_image_accessor &operator=(sampled_image_accessor &&Rhs) = default;
3772 
3773  ~sampled_image_accessor() = default;
3774 
3775 #ifdef __SYCL_DEVICE_ONLY__
3776  bool operator==(const sampled_image_accessor &Rhs) const;
3777 #else
3778  bool operator==(const sampled_image_accessor &Rhs) const {
3779  std::ignore = Rhs;
3780  throw sycl::exception(
3781  sycl::make_error_code(sycl::errc::feature_not_supported),
3782  "operator== is not yet implemented.");
3783  }
3784 #endif // __SYCL_DEVICE_ONLY__
3785 
3786  bool operator!=(const sampled_image_accessor &Rhs) const {
3787  return !(Rhs == *this);
3788  }
3789 
3790  /* -- property interface members -- */
3791 
3792  size_t size() const noexcept {
3793 #ifdef __SYCL_DEVICE_ONLY__
3794  return 0;
3795 #else
3796  throw sycl::exception(
3797  sycl::make_error_code(sycl::errc::feature_not_supported),
3798  "size() is not yet implemented.");
3799 #endif // __SYCL_DEVICE_ONLY__
3800  }
3801 
3802  /* if Dimensions == 1, CoordT = float
3803  if Dimensions == 2, CoordT = float2
3804  if Dimensions == 3, CoordT = float4 */
3805  template <typename CoordT,
3806  typename = std::enable_if_t<detail::IsValidSampledCoord2020DataT<
3807  Dimensions, CoordT>::value>>
3808  DataT read(const CoordT &Coords) const noexcept {
3809  std::ignore = Coords;
3810 #ifdef __SYCL_DEVICE_ONLY__
3811  return {0, 0, 0, 0};
3812 #else
3813  throw sycl::exception(
3814  sycl::make_error_code(sycl::errc::feature_not_supported),
3815  "read() is not yet implemented.");
3816 #endif // __SYCL_DEVICE_ONLY__
3817  }
3818 };
3819 
3820 template <typename DataT, int Dimensions>
3823  public detail::OwnerLessBase<
3824  host_sampled_image_accessor<DataT, Dimensions>> {
3825  static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
3826  std::is_same_v<DataT, float4> ||
3827  std::is_same_v<DataT, half4>,
3828  "The data type of an image accessor must be only int4, "
3829  "uint4, float4 or half4 from SYCL namespace");
3830 
3832 
3833 public:
3834  using value_type = const DataT;
3835  using reference = const DataT &;
3836  using const_reference = const DataT &;
3837 
3838  template <typename AllocatorT>
3840  const property_list &PropList = {})
3841  : base_class(detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
3842  detail::getSyclObjImpl(ImageRef).get(), Dimensions,
3843  ImageRef.getElementSize(),
3844  {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
3845  ImageRef.getChannelType(), ImageRef.getChannelOrder(),
3846  ImageRef.getSampler(), PropList) {
3847  addHostSampledImageAccessorAndWait(base_class::impl.get());
3848  }
3849 
3850  /* -- common interface members -- */
3851 
3852  host_sampled_image_accessor(const host_sampled_image_accessor &Rhs) = default;
3853 
3854  host_sampled_image_accessor(host_sampled_image_accessor &&Rhs) = default;
3855 
3856  host_sampled_image_accessor &
3857  operator=(const host_sampled_image_accessor &Rhs) = default;
3858 
3859  host_sampled_image_accessor &
3860  operator=(host_sampled_image_accessor &&Rhs) = default;
3861 
3862  ~host_sampled_image_accessor() = default;
3863 
3864  bool operator==(const host_sampled_image_accessor &Rhs) const {
3865  return Rhs.impl == impl;
3866  }
3867  bool operator!=(const host_sampled_image_accessor &Rhs) const {
3868  return !(Rhs == *this);
3869  }
3870 
3871  /* -- property interface members -- */
3872 
3873  size_t size() const noexcept { return base_class::getSize().size(); }
3874 
3875  /* if Dimensions == 1, CoordT = float
3876  if Dimensions == 2, CoordT = float2
3877  if Dimensions == 3, CoordT = float4 */
3878  template <typename CoordT,
3879  typename = std::enable_if_t<detail::IsValidSampledCoord2020DataT<
3880  Dimensions, CoordT>::value>>
3881  DataT read(const CoordT &Coords) const
3882 #ifdef __SYCL_DEVICE_ONLY__
3883  ;
3884 #else
3885  {
3886  // Host implementation is only available in host code. Device is not allowed
3887  // to use host_sampled_image_accessor.
3888  return detail::imageReadSamplerHostImpl<CoordT, DataT>(
3889  Coords, base_class::getSampler(), base_class::getSize(),
3890  base_class::getPitch(), base_class::getChannelType(),
3891  base_class::getChannelOrder(), base_class::getPtr(),
3892  base_class::getElementSize());
3893  }
3894 #endif
3895 
3896 private:
3898  : base_class{Impl} {}
3899 
3900  template <class Obj>
3901  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
3902 
3903  template <class T>
3904  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
3905 };
3906 
3907 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
3908 } // namespace sycl
3909 
3910 namespace std {
3911 template <typename DataT, int Dimensions, sycl::access::mode AccessMode,
3912  sycl::access::target AccessTarget,
3914 struct hash<sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
3915  IsPlaceholder>> {
3916  using AccType = sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
3918 
3919  size_t operator()(const AccType &A) const {
3920 #ifdef __SYCL_DEVICE_ONLY__
3921  // Hash is not supported on DEVICE. Just return 0 here.
3922  (void)A;
3923  return 0;
3924 #else
3925  // getSyclObjImpl() here returns a pointer to either AccessorImplHost
3926  // or LocalAccessorImplHost depending on the AccessTarget.
3927  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
3928  return hash<decltype(AccImplPtr)>()(AccImplPtr);
3929 #endif
3930  }
3931 };
3932 
3933 template <typename DataT, int Dimensions, sycl::access_mode AccessMode>
3934 struct hash<sycl::host_accessor<DataT, Dimensions, AccessMode>> {
3935  using AccType = sycl::host_accessor<DataT, Dimensions, AccessMode>;
3936 
3937  size_t operator()(const AccType &A) const {
3938 #ifdef __SYCL_DEVICE_ONLY__
3939  // Hash is not supported on DEVICE. Just return 0 here.
3940  (void)A;
3941  return 0;
3942 #else
3943  // getSyclObjImpl() here returns a pointer to AccessorImplHost.
3944  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
3945  return hash<decltype(AccImplPtr)>()(AccImplPtr);
3946 #endif
3947  }
3948 };
3949 
3950 template <typename DataT, int Dimensions>
3951 struct hash<sycl::local_accessor<DataT, Dimensions>> {
3952  using AccType = sycl::local_accessor<DataT, Dimensions>;
3953 
3954  size_t operator()(const AccType &A) const {
3955 #ifdef __SYCL_DEVICE_ONLY__
3956  // Hash is not supported on DEVICE. Just return 0 here.
3957  (void)A;
3958  return 0;
3959 #else
3960  // getSyclObjImpl() here returns a pointer to LocalAccessorImplHost.
3961  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
3962  return hash<decltype(AccImplPtr)>()(AccImplPtr);
3963 #endif
3964  }
3965 };
3966 
3967 template <typename DataT, int Dimensions, sycl::access_mode AccessMode,
3968  sycl::image_target AccessTarget>
3969 struct hash<sycl::unsampled_image_accessor<DataT, Dimensions, AccessMode,
3970  AccessTarget>> {
3971  using AccType = sycl::unsampled_image_accessor<DataT, Dimensions, AccessMode,
3972  AccessTarget>;
3973 
3974  size_t operator()(const AccType &A) const {
3975  // TODO: Implement.
3976  (void)A;
3977  return 0;
3978  }
3979 };
3980 
3981 template <typename DataT, int Dimensions, sycl::access_mode AccessMode>
3982 struct hash<
3983  sycl::host_unsampled_image_accessor<DataT, Dimensions, AccessMode>> {
3984  using AccType =
3985  sycl::host_unsampled_image_accessor<DataT, Dimensions, AccessMode>;
3986 
3987  size_t operator()(const AccType &A) const {
3988  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
3989  return hash<decltype(AccImplPtr)>()(AccImplPtr);
3990  }
3991 };
3992 
3993 template <typename DataT, int Dimensions, sycl::image_target AccessTarget>
3994 struct hash<sycl::sampled_image_accessor<DataT, Dimensions, AccessTarget>> {
3995  using AccType = sycl::sampled_image_accessor<DataT, Dimensions, AccessTarget>;
3996 
3997  size_t operator()(const AccType &A) const {
3998  // TODO: Implement.
3999  (void)A;
4000  return 0;
4001  }
4002 };
4003 
4004 template <typename DataT, int Dimensions>
4005 struct hash<sycl::host_sampled_image_accessor<DataT, Dimensions>> {
4006  using AccType = sycl::host_sampled_image_accessor<DataT, Dimensions>;
4007 
4008  size_t operator()(const AccType &A) const {
4009  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
4010  return hash<decltype(AccImplPtr)>()(AccImplPtr);
4011  }
4012 };
4013 
4014 } // 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:3060
sycl::_V1::detail::LocalAccessorBaseDevice::Offset
id< Dims > Offset
Definition: accessor.hpp:459
sycl::_V1::property_list
Objects of the property_list class are containers for the SYCL properties.
Definition: property_list.hpp:24
sycl::_V1::detail::AccessorImplPtr
std::shared_ptr< AccessorImplHost > AccessorImplPtr
Definition: accessor.hpp:494
sycl::_V1::__SYCL2020_DEPRECATED
signed char __SYCL2020_DEPRECATED
Definition: aliases.hpp:96
sycl::_V1::detail::AccPropBufferLocation
@ AccPropBufferLocation
Definition: property_helper.hpp:59
sycl::_V1::detail::InitializedVal
Definition: common.hpp:326
sycl::_V1::detail::associateWithHandler
void associateWithHandler(handler &, AccessorBaseHost *, access::target)
Definition: handler_proxy.cpp:17
sycl::_V1::detail::LocalAccessorImplHost
Definition: accessor_impl.hpp:131
sycl::_V1::detail::SampledImageAccessorBaseHost::SampledImageAccessorBaseHost
SampledImageAccessorBaseHost(const SampledImageAccessorImplPtr &Impl)
Definition: accessor.hpp:623
sycl::_V1::host_unsampled_image_accessor::const_reference
const DataT & const_reference
Definition: accessor.hpp:3624
sycl::_V1::detail::UnsampledImageAccessorBaseHost::impl
UnsampledImageAccessorImplPtr impl
Definition: accessor.hpp:618
property_list.hpp
sycl::_V1::detail::image_accessor
Definition: accessor.hpp:701
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:3364
sycl::_V1::local_accessor_base::operator==
bool operator==(const local_accessor_base &Rhs) const
Definition: accessor.hpp:2827
sycl::_V1::image
Defines a shared image data.
Definition: image.hpp:30
sycl::_V1::access::mode
mode
Definition: access.hpp:30
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:3374
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:3328
sycl::_V1::host_unsampled_image_accessor::host_unsampled_image_accessor
host_unsampled_image_accessor(unsampled_image< Dimensions, AllocatorT > &ImageRef, const property_list &PropList={})
Definition: accessor.hpp:3627
sycl::_V1::image_sampler
Definition: sampler.hpp:129
T
sycl::_V1::host_unsampled_image_accessor::value_type
typename std::conditional< AccessMode==access_mode::read, const DataT, DataT >::type value_type
Definition: accessor.hpp:3622
sycl::_V1::host_accessor::host_accessor
host_accessor(const detail::AccessorImplPtr &Impl)
Definition: accessor.hpp:3240
sycl::_V1::mode_tag_t
Definition: access.hpp:67
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:3867
sycl::_V1::host_unsampled_image_accessor
Definition: accessor.hpp:3608
sycl::_V1::local_accessor_base::operator!=
bool operator!=(const local_accessor_base &Rhs) const
Definition: accessor.hpp:2830
sycl::_V1::ext::oneapi::experimental::buffer_location
constexpr buffer_location_key::value_t< N > buffer_location
Definition: properties.hpp:86
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:3138
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:3296
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:357
sycl::_V1::local_accessor_base::value_type
DataT value_type
Definition: accessor.hpp:2693
device.hpp
sycl::_V1::errc::feature_not_supported
@ feature_not_supported
sycl::_V1::local_accessor_base::getSize
const range< 3 > & getSize() const
Definition: accessor.hpp:2658
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:11
sycl::_V1::detail::accessor_common::AccessorSubscript::AccessorSubscript
AccessorSubscript(AccType Accessor, id< Dims > IDs)
Definition: accessor.hpp:332
sycl::_V1::buffer
Defines a shared array that can be used by kernels in queues.
Definition: buffer.hpp:37
sycl::_V1::local_accessor_base::local_accessor_base
local_accessor_base(const detail::LocalAccessorImplPtr &Impl)
Definition: accessor.hpp:2645
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:3355
sycl::_V1::detail::image_accessor::read
DataT read(const CoordT &Coords, const sampler &Smpl) const
Definition: accessor.hpp:966
sycl::_V1::local_accessor_base::const_reference
const DataT & const_reference
Definition: accessor.hpp:2695
sycl::_V1::local_accessor_base::get_size
size_t get_size() const
Definition: accessor.hpp:2761
sycl::_V1::detail::image_common< Dimensions, sycl::image_allocator >::get_range
range< Dimensions > get_range() const
Definition: image.hpp:293
handler_proxy.hpp
sycl::_V1::local_accessor_base::reference
DataT & reference
Definition: accessor.hpp:2694
sycl::_V1::detail::image_accessor::write
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:991
accessor_property_list.hpp
sycl::_V1::local_accessor_base::getPtr
void * getPtr() const
Definition: accessor.hpp:2657
std::hash< sycl::unsampled_image_accessor< DataT, Dimensions, AccessMode, AccessTarget > >::AccType
sycl::unsampled_image_accessor< DataT, Dimensions, AccessMode, AccessTarget > AccType
Definition: accessor.hpp:3972
sycl::_V1::detail::LocalAccessorBaseHost
Definition: accessor.hpp:551
sycl::_V1::image::get_range
range< Dimensions > get_range() const
Definition: image.hpp:599
sycl::_V1::detail::SampledImageAccessorBaseHost
Definition: accessor.hpp:621
std::hash< sycl::accessor< DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:3919
sycl::_V1::ext::oneapi::experimental::operator[]
T & operator[](std::ptrdiff_t idx) const noexcept
Definition: annotated_arg.hpp:160
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:232
sycl::_V1::detail::accessor_common::AccessorSubscript
Definition: accessor.hpp:325
sycl::_V1::host_accessor::host_accessor
host_accessor(const host_accessor< DataT_, Dimensions, AccessMode_ > &other)
Definition: accessor.hpp:3428
sycl::_V1::detail::IsPropertyListT
typename std::is_base_of< PropertyListBase, T > IsPropertyListT
Definition: accessor.hpp:250
sycl::_V1::detail::image_accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::reference
DataT & reference
Definition: accessor.hpp:809
sycl::_V1::detail::accessor_common::AccessorSubscript::operator[]
auto operator[](size_t Index)
Definition: accessor.hpp:342
std::hash< sycl::local_accessor< DataT, Dimensions > >::AccType
sycl::local_accessor< DataT, Dimensions > AccType
Definition: accessor.hpp:3952
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:693
sycl::_V1::detail::LocalAccessorBaseDevice
Definition: accessor.hpp:450
sycl::_V1::local_accessor_base::ConcreteASPtrType
typename detail::DecoratedType< DataT, AS >::type * ConcreteASPtrType
Definition: accessor.hpp:2602
sycl::_V1::detail::AccHostDataT::MOffset
sycl::id< 3 > MOffset
Definition: accessor.hpp:237
sycl::_V1::detail::constructorNotification
void constructorNotification(void *BufferObj, void *AccessorObj, access::target Target, access::mode Mode, const code_location &CodeLoc)
accessor_iterator.hpp
sycl::_V1::host_accessor
Definition: accessor.hpp:3216
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:2728
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:3345
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:3059
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:2856
__SYCL_EBO
#define __SYCL_EBO
Definition: common.hpp:283
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
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:71
sycl::_V1::detail::accessor_common::AccessorSubscript::AccessorSubscript
AccessorSubscript(AccType Accessor, size_t Index)
Definition: accessor.hpp:337
sycl::_V1::local_accessor_base::RefType
detail::const_if_const_AS< AS, DataT > & RefType
Definition: accessor.hpp:2604
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:1034
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:926
__SYCL_SPECIAL_CLASS
#define __SYCL_SPECIAL_CLASS
Definition: defines.hpp:30
sycl::_V1::detail::image_plain::getChannelOrder
image_channel_order getChannelOrder() const
Definition: image.cpp:197
sycl::_V1::detail::accessor_common
Definition: accessor.hpp:274
operator==
bool operator==(const Slab &Lhs, const Slab &Rhs)
Definition: usm_allocator.cpp:313
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:2744
sycl::_V1::id< Dims >
sycl::_V1::detail::LocalAccessorBaseDevice::AccessRange
range< Dims > AccessRange
Definition: accessor.hpp:457
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:254
sycl::_V1::local_accessor_base
Local accessor.
Definition: accessor.hpp:2575
sycl::_V1::detail::__image_array_slice__::size
size_t size() const noexcept
Definition: accessor.hpp:1082
owner_less_base.hpp
sycl::_V1::detail::SampledImageAccessorImplHost
Definition: accessor_impl.hpp:166
std::hash< sycl::unsampled_image_accessor< DataT, Dimensions, AccessMode, AccessTarget > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:3974
sycl::_V1::detail::accessModeFromConstness
constexpr access::mode accessModeFromConstness()
Definition: accessor.hpp:364
sycl::_V1::detail::AccessorImplDevice::AccessorImplDevice
AccessorImplDevice(id< Dims > Offset, range< Dims > AccessRange, range< Dims > MemoryRange)
Definition: accessor.hpp:474
sycl::_V1::host_accessor::IsSameAsBuffer
Definition: accessor.hpp:3227
sycl::_V1::host_sampled_image_accessor::size
size_t size() const noexcept
Definition: accessor.hpp:3873
sycl::_V1::detail::accessor_common< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::PtrType
detail::const_if_const_AS< AS, DataT > * PtrType
Definition: accessor.hpp:317
sycl::_V1::detail::__image_array_slice__::read
DataT read(const CoordT &Coords, const sampler &Smpl) const
Definition: accessor.hpp:1051
sycl::_V1::local_accessor_base::get_range
range< Dims > get_range() const
Definition: accessor.hpp:2768
image_accessor_util.hpp
sycl::_V1::detail::write
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
Definition: stream.hpp:110
sycl::_V1::detail::IsCxPropertyList
Definition: accessor.hpp:256
sycl::_V1::detail::LocalAccessorBaseDevice::LocalAccessorBaseDevice
LocalAccessorBaseDevice(sycl::range< Dims > Size)
Definition: accessor.hpp:452
sycl::_V1::range
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: buffer.hpp:28
std::hash< sycl::host_sampled_image_accessor< DataT, Dimensions > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:4008
sycl::_V1::local_accessor_base::local_accessor_base
local_accessor_base()
Definition: accessor.hpp:2640
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:3312
sycl::_V1::image_accessor
Definition: accessor_properties.hpp:104
sycl::_V1::access::placeholder
placeholder
Definition: access.hpp:45
sycl::_V1::host_sampled_image_accessor
Definition: accessor.hpp:3821
sycl::_V1::unsampled_image_accessor::value_type
typename std::conditional< AccessMode==access_mode::read, const DataT, DataT >::type value_type
Definition: accessor.hpp:3504
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::unsampled_image_accessor::unsampled_image_accessor
unsampled_image_accessor(unsampled_image< Dimensions, AllocatorT > &ImageRef, handler &CommandGroupHandlerRef, const property_list &PropList={})
Definition: accessor.hpp:3509
sycl::_V1::host_accessor::operator=
const host_accessor & operator=(typename AccessorT::value_type &&Other) const
Definition: accessor.hpp:3404
sycl::_V1::ext::oneapi::experimental::has_property
static constexpr bool has_property()
Definition: annotated_arg.hpp:162
std::get
constexpr tuple_element< I, tuple< Types... > >::type & get(sycl::detail::tuple< Types... > &Arg) noexcept
Definition: tuple.hpp:199
sycl::_V1::host_unsampled_image_accessor::operator==
bool operator==(const host_unsampled_image_accessor &Rhs) const
Definition: accessor.hpp:3654
sycl::_V1::detail::image_accessor::operator==
bool operator==(const image_accessor &Rhs) const
Definition: accessor.hpp:885
sycl::_V1::unsampled_image_accessor::size
size_t size() const noexcept
Definition: accessor.hpp:3556
sycl::_V1::local_accessor_base::getPtr
void * getPtr()
Definition: accessor.hpp:2656
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:2763
sycl::_V1::unsampled_image_accessor::write
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:3593
cl.h
sycl::_V1::ext::oneapi::detail::getSyclWeakObjImpl
decltype(weak_object_base< SYCLObjT >::MObjWeakPtr) getSyclWeakObjImpl(const weak_object_base< SYCLObjT > &WeakObj)
Definition: weak_object_base.hpp:21
sycl::_V1::detail::deduceAccessMode
constexpr access::mode deduceAccessMode()
Definition: accessor.hpp:372
std::hash< sycl::host_accessor< DataT, Dimensions, AccessMode > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:3937
std::hash< sycl::sampled_image_accessor< DataT, Dimensions, AccessTarget > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:3997
sycl::_V1::detail::image_accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::const_reference
const DataT & const_reference
Definition: accessor.hpp:810
sycl::_V1::sampled_image_accessor::value_type
const DataT value_type
Definition: accessor.hpp:3740
generic_type_traits.hpp
sycl::_V1::ext::oneapi::experimental::operator=
annotated_arg & operator=(annotated_arg &)=default
sycl::_V1::detail::SampledImageAccessorBaseHost::impl
SampledImageAccessorImplPtr impl
Definition: accessor.hpp:653
sycl::_V1::host_unsampled_image_accessor::write
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:3704
sycl::_V1::host_unsampled_image_accessor::size
size_t size() const noexcept
Definition: accessor.hpp:3663
sycl::_V1::detail::UnsampledImageAccessorImplPtr
std::shared_ptr< UnsampledImageAccessorImplHost > UnsampledImageAccessorImplPtr
Definition: accessor.hpp:579
sycl::_V1::detail::OwnerLessBase
Definition: owner_less_base.hpp:21
sycl::_V1::detail::SampledImageAccessorImplPtr
std::shared_ptr< SampledImageAccessorImplHost > SampledImageAccessorImplPtr
Definition: accessor.hpp:581
sycl::_V1::host_sampled_image_accessor::const_reference
const DataT & const_reference
Definition: accessor.hpp:3836
sycl::_V1::handler
Command group handler class.
Definition: handler.hpp:325
sycl::_V1::detail::__image_array_slice__
Definition: accessor.hpp:696
sycl::_V1::access_mode
access::mode access_mode
Definition: access.hpp:63
common.hpp
sycl::_V1::ext::oneapi::detail::weak_object_base
Definition: weak_object_base.hpp:16
std::hash< sycl::local_accessor< DataT, Dimensions > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:3954
sycl::_V1::sampled_image_accessor::size
size_t size() const noexcept
Definition: accessor.hpp:3792
sycl::_V1::detail::TargetToAS
Definition: access.hpp:131
sycl::_V1::local_accessor_base::size
size_t size() const noexcept
Definition: accessor.hpp:2765
sycl::_V1::image_channel_order
image_channel_order
Definition: image.hpp:38
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::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:3232
sycl::_V1::device
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:49
sycl::_V1::detail::createSyclObjFromImpl
T createSyclObjFromImpl(decltype(T::impl) ImplObj)
Definition: common.hpp:320
sycl::_V1::host_accessor::host_accessor
host_accessor()
Definition: accessor.hpp:3252
sycl::_V1::detail::__image_array_slice__::get_range
range< Dims > get_range() const
Definition: accessor.hpp:1088
sycl::_V1::ext::oneapi::accessor_property_list
Definition: property_list.hpp:18
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:2711
sycl::_V1::detail::AccessorBaseHost::impl
AccessorImplPtr impl
Definition: accessor.hpp:542
sycl::_V1::detail::UnsampledImageAccessorBaseHost
Definition: accessor.hpp:588
sycl::_V1::host_unsampled_image_accessor::read
DataT read(const CoordT &Coords) const noexcept
Definition: accessor.hpp:3676
sycl::_V1::local_accessor_base::AccessorSubscript
typename AccessorCommonT::template AccessorSubscript< Dims, local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder > > AccessorSubscript
Definition: accessor.hpp:2600
sycl::_V1::access::target
target
Definition: access.hpp:18
sycl::_V1::read_write
constexpr mode_tag_t< access_mode::read_write > read_write
Definition: access.hpp:76
sycl::_V1::detail::IsValidSampledCoord2020DataT
Definition: accessor.hpp:683
sycl::_V1::host_sampled_image_accessor::value_type
const DataT value_type
Definition: accessor.hpp:3834
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:3506
sycl::_V1::unsampled_image_accessor
Definition: accessor.hpp:3495
sycl::_V1::ext::oneapi::experimental::get_property
static constexpr auto get_property()
Definition: annotated_arg.hpp:166
sycl::_V1::host_accessor
host_accessor(buffer< DataT, Dimensions, AllocatorT >, Type1, Type2, Type3, Type4, Type5) -> host_accessor< DataT, Dimensions, detail::deduceAccessMode< Type4, Type5 >()>
sycl::_V1::detail::AccessorImplDevice
Definition: accessor.hpp:471
image_ocl_types.hpp
sycl::_V1::sampled_image_accessor::read
DataT read(const CoordT &Coords) const noexcept
Definition: accessor.hpp:3808
sycl::_V1::accessor
Definition: accessor.hpp:225
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:3384
sycl::_V1::detail::IsValidCoordDataT
Definition: accessor.hpp:656
sycl::_V1::detail::LocalAccessorBaseHost::LocalAccessorBaseHost
LocalAccessorBaseHost(const LocalAccessorImplPtr &Impl)
Definition: accessor.hpp:553
image.hpp
sycl::_V1::detail::LocalAccessorImplPtr
std::shared_ptr< LocalAccessorImplHost > LocalAccessorImplPtr
Definition: accessor.hpp:549
sycl::_V1::local_accessor_base::GDBMethodsAnchor
void GDBMethodsAnchor()
Definition: accessor.hpp:2664
sycl::_V1::detail::AccessorImplDevice::MemRange
range< Dims > MemRange
Definition: accessor.hpp:480
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:3288
std::hash< sycl::host_unsampled_image_accessor< DataT, Dimensions, AccessMode > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:3987
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:3505
sycl::_V1::detail::type_list
Definition: type_list.hpp:23
sycl::_V1::detail::accessor_common< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::ConstRefType
const DataT & ConstRefType
Definition: accessor.hpp:316
sycl::_V1::local_accessor_base::operator[]
RefType operator[](id< Dimensions > Index) const
Definition: accessor.hpp:2782
sycl::_V1::host_sampled_image_accessor::read
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:3881
sycl::_V1::detail::TryToGetElementType::type
decltype(check(T())) type
Definition: generic_type_traits.hpp:355
sycl::_V1::detail::UnsampledImageAccessorImplHost
Definition: accessor_impl.hpp:150
sycl::_V1::ext::oneapi::no_offset
constexpr property::no_offset::instance no_offset
Definition: accessor_properties.hpp:78
sycl::_V1::detail::AccHostDataT
Definition: accessor.hpp:231
sycl::_V1::detail::AccessorImplDevice::operator==
bool operator==(const AccessorImplDevice &Rhs) const
Definition: accessor.hpp:482
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:3304
sycl::_V1::detail::code_location
Definition: common.hpp:66
sycl::_V1::detail::deduceAccessTarget
constexpr access::target deduceAccessTarget(access::target defaultTarget)
Definition: accessor.hpp:419
sycl::_V1::detail::AccHostDataT::MAccessRange
sycl::range< 3 > MAccessRange
Definition: accessor.hpp:238
sycl::_V1::detail::AccHostDataT::MMemoryRange
sycl::range< 3 > MMemoryRange
Definition: accessor.hpp:239
sycl::_V1::host_unsampled_image_accessor::operator!=
bool operator!=(const host_unsampled_image_accessor &Rhs) const
Definition: accessor.hpp:3657
sycl::_V1::detail::__image_array_slice__::read
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:1044
sycl::_V1::accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::accessor
accessor(sycl::image< Dimensions, AllocatorT > &Image)
Definition: accessor.hpp:3132
sycl::_V1::detail::image_accessor::image_accessor
image_accessor(image< Dims, AllocatorT > &ImageRef, int ImageElementSize)
Definition: accessor.hpp:825
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:3440
sycl::_V1::local_accessor_base::PtrType
detail::const_if_const_AS< AS, DataT > * PtrType
Definition: accessor.hpp:2605
sycl::_V1::vec
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
Definition: aliases.hpp:20
sycl::_V1::detail::LocalAccessorBaseDevice::operator==
bool operator==(const LocalAccessorBaseDevice &Rhs) const
Definition: accessor.hpp:461
exception.hpp
sycl::_V1::host_sampled_image_accessor::host_sampled_image_accessor
host_sampled_image_accessor(sampled_image< Dimensions, AllocatorT > &ImageRef, const property_list &PropList={})
Definition: accessor.hpp:3839
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:1080
sycl::_V1::local_accessor
Definition: multi_ptr.hpp:68
std
Definition: accessor.hpp:3910
sycl::_V1::host_accessor::ext_oneapi_owner_before
bool ext_oneapi_owner_before(const host_accessor &Other) const noexcept
Definition: accessor.hpp:3447
sycl::_V1::detail::accessor_common< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::RefType
detail::const_if_const_AS< AS, DataT > & RefType
Definition: accessor.hpp:315
sycl::_V1::detail::getDeviceFromHandler
device getDeviceFromHandler(handler &CommandGroupHandlerRef)
Definition: accessor.cpp:15
sycl::_V1::detail::is_contained
Definition: type_list.hpp:55
sycl::_V1::local_accessor_base::local_accessor_base
local_accessor_base(handler &, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2698
sycl::_V1::detail::image_accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::value_type
DataT value_type
Definition: accessor.hpp:808
sycl::_V1::local_accessor_base::getQualifiedPtr
PtrType getQualifiedPtr() const
Definition: accessor.hpp:2652
__SYCL_TYPE
#define __SYCL_TYPE(x)
Definition: defines.hpp:40
sampler.hpp
sycl::_V1::detail::image_accessor::get_range
range< Dims > get_range() const
Definition: accessor.hpp:931
sycl::_V1::local_accessor_base::getLinearIndex
size_t getLinearIndex(id< Dims > Id) const
Definition: accessor.hpp:2677
sycl::_V1::detail::AccessorImplDevice::AccessRange
range< Dims > AccessRange
Definition: accessor.hpp:479
sycl::_V1::host_accessor::host_accessor
host_accessor(const host_accessor< DataT_, Dimensions, AccessMode > &other)
Definition: accessor.hpp:3415
buffer.hpp
sycl::_V1::sampled_image_accessor::operator!=
bool operator!=(const sampled_image_accessor &Rhs) const
Definition: accessor.hpp:3786
sycl::_V1::sampled_image_accessor::const_reference
const DataT & const_reference
Definition: accessor.hpp:3742
sycl::_V1::host_unsampled_image_accessor::reference
value_type & reference
Definition: accessor.hpp:3623
sycl::_V1::detail::image_accessor::read
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:947
std::hash< sycl::host_sampled_image_accessor< DataT, Dimensions > >::AccType
sycl::host_sampled_image_accessor< DataT, Dimensions > AccType
Definition: accessor.hpp:4006
sycl::_V1::detail::IsValidUnsampledCoord2020DataT
Definition: accessor.hpp:672
sycl::_V1::detail::getLinearIndex
size_t getLinearIndex(const T< Dims > &Index, const U< Dims > &Range)
Definition: common.hpp:435
sycl::_V1::detail::image_accessor::image_accessor
image_accessor(image< Dims, AllocatorT > &ImageRef, handler &CommandGroupHandlerRef, int ImageElementSize)
Definition: accessor.hpp:852
sycl::_V1::detail::DecoratedType
Definition: access.hpp:154
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:2814
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:3336
sycl::_V1::unsampled_image_accessor::read
DataT read(const CoordT &Coords) const noexcept
Definition: accessor.hpp:3574
std::hash< sycl::host_unsampled_image_accessor< DataT, Dimensions, AccessMode > >::AccType
sycl::host_unsampled_image_accessor< DataT, Dimensions, AccessMode > AccType
Definition: accessor.hpp:3985
sycl::_V1::ext::oneapi::experimental::reference
sycl::ext::oneapi::experimental::annotated_ref< T, property_list_t > reference
Definition: annotated_ptr.hpp:101
sycl::_V1::ext::intel::experimental::operator!=
bool operator!=(const cache_config &lhs, const cache_config &rhs)
Definition: kernel_execution_properties.hpp:37
sycl::_V1::unsampled_image_accessor::operator!=
bool operator!=(const unsampled_image_accessor &Rhs) const
Definition: accessor.hpp:3550
std::hash< sycl::accessor< DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder > >::AccType
sycl::accessor< DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder > AccType
Definition: accessor.hpp:3917
sycl::_V1::local_accessor_base::getSize
range< 3 > & getSize()
Definition: accessor.hpp:2661
sycl::_V1::image_channel_type
image_channel_type
Definition: image.hpp:56
sycl::_V1::sampled_image_accessor
Definition: accessor.hpp:3732
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:3059
weak_object_base.hpp
sycl::_V1::detail::image_accessor::image_accessor
image_accessor(const AccessorImplPtr &Impl)
Definition: accessor.hpp:797
pointers.hpp
sycl::_V1::detail::get
Definition: tuple.hpp:59
sycl::_V1::detail::image_accessor::operator!=
bool operator!=(const image_accessor &Rhs) const
Definition: accessor.hpp:893
sycl::_V1::detail::AccessorImplDevice::Offset
id< Dims > Offset
Definition: accessor.hpp:478
sycl::_V1::sampled_image_accessor::sampled_image_accessor
sampled_image_accessor(sampled_image< Dimensions, AllocatorT > &ImageRef, handler &CommandGroupHandlerRef, const property_list &PropList={})
Definition: accessor.hpp:3745
sycl::_V1::detail::__image_array_slice__::write
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:1058
sycl::_V1::host_accessor::operator=
const host_accessor & operator=(const typename AccessorT::value_type &Other) const
Definition: accessor.hpp:3396
std::hash< sycl::host_accessor< DataT, Dimensions, AccessMode > >::AccType
sycl::host_accessor< DataT, Dimensions, AccessMode > AccType
Definition: accessor.hpp:3935
sycl::_V1::sampled_image_accessor::operator==
bool operator==(const sampled_image_accessor &Rhs) const
Definition: accessor.hpp:3778
std::hash< sycl::sampled_image_accessor< DataT, Dimensions, AccessTarget > >::AccType
sycl::sampled_image_accessor< DataT, Dimensions, AccessTarget > AccType
Definition: accessor.hpp:3995
sycl::_V1::sampled_image
Definition: image.hpp:911
sycl::_V1::detail::LocalAccessorBaseHost::impl
LocalAccessorImplPtr impl
Definition: accessor.hpp:573
atomic.hpp
__SYCL2020_DEPRECATED
#define __SYCL2020_DEPRECATED(message)
Definition: defines_elementary.hpp:57
sycl::_V1::detail::AccessorBaseHost::AccessorBaseHost
AccessorBaseHost(const AccessorImplPtr &Impl)
Definition: accessor.hpp:498
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:2823
sycl::_V1::image_target
image_target
Definition: access.hpp:65
sycl::_V1::unsampled_image
Definition: image.hpp:696
sycl::_V1::unsampled_image_accessor::operator==
bool operator==(const unsampled_image_accessor &Rhs) const
Definition: accessor.hpp:3542
accessor_properties.hpp
buffer_properties.hpp
sycl::_V1::detail::getSyclObjImpl
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: common.hpp:302
sycl::_V1::host_sampled_image_accessor::operator==
bool operator==(const host_sampled_image_accessor &Rhs) const
Definition: accessor.hpp:3864
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:3320
sycl::_V1::detail::image_accessor::size
size_t size() const noexcept
Definition: accessor.hpp:928
sycl::_V1::detail::LocalAccessorBaseDevice::MemRange
range< Dims > MemRange
Definition: accessor.hpp:458
sycl::_V1::detail::const_if_const_AS
DataT const_if_const_AS
Definition: type_traits.hpp:462
sycl::_V1::detail::UnsampledImageAccessorBaseHost::UnsampledImageAccessorBaseHost
UnsampledImageAccessorBaseHost(const UnsampledImageAccessorImplPtr &Impl)
Definition: accessor.hpp:590
spirv_types.hpp
property_list_conversion.hpp
sycl::_V1::host_sampled_image_accessor::reference
const DataT & reference
Definition: accessor.hpp:3835
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:3280
sycl::_V1::access::address_space
address_space
Definition: access.hpp:47
sycl::_V1::detail::AccessorBaseHost
Definition: accessor.hpp:496
sycl::_V1::sampled_image_accessor::reference
const DataT & reference
Definition: accessor.hpp:3741