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 // To ensure loop unrolling is done when processing dimensions.
245 template <size_t... Inds, class F>
246 void dim_loop_impl(std::integer_sequence<size_t, Inds...>, F &&f) {
247  (f(Inds), ...);
248 }
249 
250 template <size_t count, class F> void dim_loop(F &&f) {
251  dim_loop_impl(std::make_index_sequence<count>{}, std::forward<F>(f));
252 }
253 
254 void __SYCL_EXPORT constructorNotification(void *BufferObj, void *AccessorObj,
255  access::target Target,
256  access::mode Mode,
257  const code_location &CodeLoc);
258 
259 template <typename T>
260 using IsPropertyListT = typename std::is_base_of<PropertyListBase, T>;
261 
262 template <typename T>
264  typename std::is_same<ext::oneapi::accessor_property_list<>, T>;
265 
266 template <typename T> struct IsCxPropertyList {
267  constexpr static bool value = false;
268 };
269 
270 template <typename... Props>
271 struct IsCxPropertyList<ext::oneapi::accessor_property_list<Props...>> {
272  constexpr static bool value = true;
273 };
274 
275 template <> struct IsCxPropertyList<ext::oneapi::accessor_property_list<>> {
276  constexpr static bool value = false;
277 };
278 
279 __SYCL_EXPORT device getDeviceFromHandler(handler &CommandGroupHandlerRef);
280 
281 template <typename DataT, int Dimensions, access::mode AccessMode,
283  typename PropertyListT = ext::oneapi::accessor_property_list<>>
285 protected:
287 
288  constexpr static bool IsHostBuf = AccessTarget == access::target::host_buffer;
289  // SYCL2020 4.7.6.9.4.3
290  // IsPlaceHolder template parameter has no bearing on whether the accessor
291  // instance is a placeholder. This is determined solely by the constructor.
292  // The rule seems to be: if the constructor receives a CommandGroupHandler
293  // it is NOT a placeholder. Otherwise, it is a placeholder.
294  // However, according to 4.7.6.9.4.6. accessor specialization with
295  // target::host_buffer is never a placeholder. So, if the constructor
296  // used receives a CommandGroupHandler, the accessor will never be a
297  // placeholder. If it doesn't, but IsHostBuf is true, it won't be a
298  // placeholder either. Otherwise, the accessor is a placeholder.
299  constexpr static bool IsPlaceH = !IsHostBuf;
300 
301  // TODO: SYCL 2020 deprecates four of the target enum values
302  // and replaces them with 2 (device and host_task). May want
303  // to change these constexpr.
304  constexpr static bool IsGlobalBuf =
305  AccessTarget == access::target::global_buffer;
306 
307  constexpr static bool IsConstantBuf =
308  AccessTarget == access::target::constant_buffer;
309 
310  constexpr static bool IsAccessAnyWrite =
313  AccessMode == access::mode::discard_write ||
314  AccessMode == access::mode::discard_read_write;
315 
316  constexpr static bool IsAccessReadOnly = AccessMode == access::mode::read;
317 
318  constexpr static bool IsAccessReadWrite =
320 
321  constexpr static bool IsAccessAtomic = AccessMode == access::mode::atomic;
322 
324  using ConstRefType = const DataT &;
326 
327  // The class which allows to access value of N dimensional accessor using N
328  // subscript operators, e.g. accessor[2][2][3]
329  template <int SubDims,
330  typename AccType =
331  accessor<DataT, Dimensions, AccessMode, AccessTarget,
332  IsPlaceholder, PropertyListT>>
334  static constexpr int Dims = Dimensions;
335 
336  mutable id<Dims> MIDs;
337  AccType MAccessor;
338 
339  public:
340  AccessorSubscript(AccType Accessor, id<Dims> IDs)
341  : MIDs(IDs), MAccessor(Accessor) {}
342 
343  // Only accessor class is supposed to use this c'tor for the first
344  // operator[].
345  AccessorSubscript(AccType Accessor, size_t Index) : MAccessor(Accessor) {
346  MIDs[0] = Index;
347  }
348 
349  template <int CurDims = SubDims,
350  typename = detail::enable_if_t<(CurDims > 1)>>
351  auto operator[](size_t Index) {
352  MIDs[Dims - CurDims] = Index;
353  return AccessorSubscript<CurDims - 1, AccType>(MAccessor, MIDs);
354  }
355 
356  template <int CurDims = SubDims,
357  typename = detail::enable_if_t<
358  CurDims == 1 && (IsAccessReadOnly || IsAccessAnyWrite)>>
359  typename AccType::reference operator[](size_t Index) const {
360  MIDs[Dims - CurDims] = Index;
361  return MAccessor[MIDs];
362  }
363 
364  template <int CurDims = SubDims>
365  typename detail::enable_if_t<CurDims == 1 && IsAccessAtomic,
366  atomic<DataT, AS>>
367  operator[](size_t Index) const {
368  MIDs[Dims - CurDims] = Index;
369  return MAccessor[MIDs];
370  }
371  };
372 };
373 
374 template <typename MayBeTag1, typename MayBeTag2>
376  // property_list = {} is not properly detected by deduction guide,
377  // when parameter is passed without curly braces: access(buffer, no_init)
378  // thus simplest approach is to check 2 last arguments for being a tag
379  if constexpr (std::is_same<MayBeTag1,
381  std::is_same<MayBeTag2,
383  return access::mode::read;
384  }
385 
386  if constexpr (std::is_same<MayBeTag1,
388  std::is_same<MayBeTag2,
390  return access::mode::write;
391  }
392 
393  if constexpr (
394  std::is_same<MayBeTag1,
395  mode_target_tag_t<access::mode::read,
396  access::target::constant_buffer>>::value ||
397  std::is_same<MayBeTag2,
398  mode_target_tag_t<access::mode::read,
399  access::target::constant_buffer>>::value) {
400  return access::mode::read;
401  }
402 
404 }
405 
406 template <typename MayBeTag1, typename MayBeTag2>
408  if constexpr (
409  std::is_same<MayBeTag1,
410  mode_target_tag_t<access::mode::read,
411  access::target::constant_buffer>>::value ||
412  std::is_same<MayBeTag2,
413  mode_target_tag_t<access::mode::read,
414  access::target::constant_buffer>>::value) {
415  return access::target::constant_buffer;
416  }
417 
418  return defaultTarget;
419 }
420 
421 template <int Dims> class LocalAccessorBaseDevice {
422 public:
423  LocalAccessorBaseDevice(sycl::range<Dims> Size)
424  : AccessRange(Size),
425  MemRange(InitializedVal<Dims, range>::template get<0>()) {}
426  // TODO: Actually we need only one field here, but currently compiler requires
427  // all of them.
431 
432  bool operator==(const LocalAccessorBaseDevice &Rhs) const {
433  return (AccessRange == Rhs.AccessRange);
434  }
435 };
436 
437 // The class describes a requirement to access a SYCL memory object such as
438 // sycl::buffer and sycl::image. For example, each accessor used in a kernel,
439 // except one with access target "local", adds such requirement for the command
440 // group.
441 
442 template <int Dims> class AccessorImplDevice {
443 public:
444  AccessorImplDevice() = default;
446  range<Dims> MemoryRange)
447  : Offset(Offset), AccessRange(AccessRange), MemRange(MemoryRange) {}
448 
452 
453  bool operator==(const AccessorImplDevice &Rhs) const {
454  return (Offset == Rhs.Offset && AccessRange == Rhs.AccessRange &&
455  MemRange == Rhs.MemRange);
456  }
457 };
458 
459 class AccessorImplHost;
460 
461 void __SYCL_EXPORT addHostAccessorAndWait(AccessorImplHost *Req);
462 
463 class SYCLMemObjI;
464 
465 using AccessorImplPtr = std::shared_ptr<AccessorImplHost>;
466 
467 class __SYCL_EXPORT AccessorBaseHost {
468 protected:
469  AccessorBaseHost(const AccessorImplPtr &Impl) : impl{Impl} {}
470 
471 public:
472  AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
473  access::mode AccessMode, void *SYCLMemObject, int Dims,
474  int ElemSize, int OffsetInBytes = 0,
475  bool IsSubBuffer = false,
476  const property_list &PropertyList = {});
477 
478  AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
479  access::mode AccessMode, void *SYCLMemObject, int Dims,
480  int ElemSize, bool IsPlaceH, int OffsetInBytes = 0,
481  bool IsSubBuffer = false,
482  const property_list &PropertyList = {});
483 
484 public:
485  id<3> &getOffset();
486  range<3> &getAccessRange();
487  range<3> &getMemoryRange();
488  void *getPtr();
489  unsigned int getElemSize() const;
490 
491  const id<3> &getOffset() const;
492  const range<3> &getAccessRange() const;
493  const range<3> &getMemoryRange() const;
494  void *getPtr() const;
495  bool isPlaceholder() const;
496 
497  detail::AccHostDataT &getAccData();
498 
499  const property_list &getPropList() const;
500 
501  void *getMemoryObject() const;
502 
503  template <class Obj>
504  friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject);
505 
506  template <class T>
507  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
508 
509  template <typename, int, access::mode, access::target, access::placeholder,
510  typename>
511  friend class accessor;
512 
514 
515 private:
516  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
517 };
518 
520 using LocalAccessorImplPtr = std::shared_ptr<LocalAccessorImplHost>;
521 
522 class __SYCL_EXPORT LocalAccessorBaseHost {
523 protected:
524  LocalAccessorBaseHost(const LocalAccessorImplPtr &Impl) : impl{Impl} {}
525 
526 public:
527  LocalAccessorBaseHost(sycl::range<3> Size, int Dims, int ElemSize,
528  const property_list &PropertyList = {});
529  sycl::range<3> &getSize();
530  const sycl::range<3> &getSize() const;
531  void *getPtr();
532  void *getPtr() const;
533  int getNumOfDims();
534  int getElementSize();
535  const property_list &getPropList() const;
536 
537 protected:
538  template <class Obj>
539  friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject);
540 
541  template <class T>
542  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
543 
545 };
546 
547 template <int Dim, typename T> struct IsValidCoordDataT;
548 template <typename T> struct IsValidCoordDataT<1, T> {
549  constexpr static bool value =
551 };
552 template <typename T> struct IsValidCoordDataT<2, T> {
553  constexpr static bool value =
556 };
557 template <typename T> struct IsValidCoordDataT<3, T> {
558  constexpr static bool value =
561 };
562 
563 template <typename DataT, int Dimensions, access::mode AccessMode,
566 
567 // Image accessor
568 template <typename DataT, int Dimensions, access::mode AccessMode,
571 #ifndef __SYCL_DEVICE_ONLY__
572  : public detail::AccessorBaseHost {
573  size_t MImageCount;
574  image_channel_order MImgChannelOrder;
575  image_channel_type MImgChannelType;
576 #else
577 {
578 
579  using OCLImageTy = typename detail::opencl_image_type<Dimensions, AccessMode,
580  AccessTarget>::type;
581  OCLImageTy MImageObj;
582  char MPadding[sizeof(detail::AccessorBaseHost) +
583  sizeof(size_t /*MImageCount*/) + sizeof(image_channel_order) +
584  sizeof(image_channel_type) - sizeof(OCLImageTy)];
585 
586 protected:
587  void imageAccessorInit(OCLImageTy Image) { MImageObj = Image; }
588 
589 private:
590 #endif
591  template <typename T1, int T2, access::mode T3, access::placeholder T4>
592  friend class __image_array_slice__;
593 
594  constexpr static bool IsHostImageAcc =
595  (AccessTarget == access::target::host_image);
596 
597  constexpr static bool IsImageAcc = (AccessTarget == access::target::image);
598 
599  constexpr static bool IsImageArrayAcc =
600  (AccessTarget == access::target::image_array);
601 
602  constexpr static bool IsImageAccessWriteOnly =
604  AccessMode == access::mode::discard_write);
605 
606  constexpr static bool IsImageAccessAnyWrite =
607  (IsImageAccessWriteOnly || AccessMode == access::mode::read_write);
608 
609  constexpr static bool IsImageAccessReadOnly =
610  (AccessMode == access::mode::read);
611 
612  constexpr static bool IsImageAccessAnyRead =
613  (IsImageAccessReadOnly || AccessMode == access::mode::read_write);
614 
615  static_assert(std::is_same<DataT, cl_int4>::value ||
616  std::is_same<DataT, cl_uint4>::value ||
617  std::is_same<DataT, cl_float4>::value ||
618  std::is_same<DataT, cl_half4>::value,
619  "The data type of an image accessor must be only cl_int4, "
620  "cl_uint4, cl_float4 or cl_half4 from SYCL namespace");
621 
622  static_assert(IsImageAcc || IsHostImageAcc || IsImageArrayAcc,
623  "Expected image type");
624 
625  static_assert(IsPlaceholder == access::placeholder::false_t,
626  "Expected false as Placeholder value for image accessor.");
627 
628  static_assert(
629  ((IsImageAcc || IsImageArrayAcc) &&
630  (IsImageAccessWriteOnly || IsImageAccessReadOnly)) ||
631  (IsHostImageAcc && (IsImageAccessAnyWrite || IsImageAccessAnyRead)),
632  "Access modes can be only read/write/discard_write for image/image_array "
633  "target accessor, or they can be only "
634  "read/write/discard_write/read_write for host_image target accessor.");
635 
636  static_assert(Dimensions > 0 && Dimensions <= 3,
637  "Dimensions can be 1/2/3 for image accessor.");
638 
639  template <typename Param>
640  void checkDeviceFeatureSupported(const device &Device) {
641  if (!Device.get_info<Param>())
642  throw feature_not_supported("Images are not supported by this device.",
643  PI_ERROR_INVALID_OPERATION);
644  }
645 
646 #ifdef __SYCL_DEVICE_ONLY__
647 
648  sycl::vec<int, Dimensions> getRangeInternal() const {
649  return __invoke_ImageQuerySize<sycl::vec<int, Dimensions>, OCLImageTy>(
650  MImageObj);
651  }
652 
653  size_t getElementSize() const {
654  int ChannelType = __invoke_ImageQueryFormat<int, OCLImageTy>(MImageObj);
655  int ChannelOrder = __invoke_ImageQueryOrder<int, OCLImageTy>(MImageObj);
656  int ElementSize = getSPIRVElementSize(ChannelType, ChannelOrder);
657  return ElementSize;
658  }
659 
660 #else
661 
662  sycl::vec<int, Dimensions> getRangeInternal() const {
663  // TODO: Implement for host.
664  throw runtime_error("image::getRangeInternal() is not implemented for host",
665  PI_ERROR_INVALID_OPERATION);
666  return sycl::vec<int, Dimensions>{1};
667  }
668 
669 #endif
670 
671 #ifndef __SYCL_DEVICE_ONLY__
672 protected:
674 #endif // __SYCL_DEVICE_ONLY__
675 
676 private:
677  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
678 
679 #ifdef __SYCL_DEVICE_ONLY__
680  const OCLImageTy getNativeImageObj() const { return MImageObj; }
681 #endif // __SYCL_DEVICE_ONLY__
682 
683 public:
684  using value_type = DataT;
685  using reference = DataT &;
686  using const_reference = const DataT &;
687 
688  // image_accessor Constructors.
689 
690 #ifdef __SYCL_DEVICE_ONLY__
691  // Default constructor for objects later initialized with __init member.
692  image_accessor() : MImageObj() {}
693 #endif
694 
695  // Available only when: accessTarget == access::target::host_image
696  // template <typename AllocatorT>
697  // accessor(image<dimensions, AllocatorT> &imageRef);
698  template <
699  typename AllocatorT, int Dims = Dimensions,
700  typename = detail::enable_if_t<(Dims > 0 && Dims <= 3) && IsHostImageAcc>>
701  image_accessor(image<Dims, AllocatorT> &ImageRef, int ImageElementSize)
702 #ifdef __SYCL_DEVICE_ONLY__
703  {
704  (void)ImageRef;
705  (void)ImageElementSize;
706  // No implementation needed for device. The constructor is only called by
707  // host.
708  }
709 #else
710  : AccessorBaseHost({ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
711  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
712  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
713  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
714  Dimensions, ImageElementSize),
715  MImageCount(ImageRef.size()),
716  MImgChannelOrder(ImageRef.getChannelOrder()),
717  MImgChannelType(ImageRef.getChannelType()) {
718  addHostAccessorAndWait(AccessorBaseHost::impl.get());
719  }
720 #endif
721 
722  // Available only when: accessTarget == access::target::image
723  // template <typename AllocatorT>
724  // accessor(image<dimensions, AllocatorT> &imageRef,
725  // handler &commandGroupHandlerRef);
726  template <
727  typename AllocatorT, int Dims = Dimensions,
728  typename = detail::enable_if_t<(Dims > 0 && Dims <= 3) && IsImageAcc>>
730  handler &CommandGroupHandlerRef, int ImageElementSize)
731 #ifdef __SYCL_DEVICE_ONLY__
732  {
733  (void)ImageRef;
734  (void)CommandGroupHandlerRef;
735  (void)ImageElementSize;
736  // No implementation needed for device. The constructor is only called by
737  // host.
738  }
739 #else
740  : AccessorBaseHost({ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
741  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
742  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
743  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
744  Dimensions, ImageElementSize),
745  MImageCount(ImageRef.size()),
746  MImgChannelOrder(ImageRef.getChannelOrder()),
747  MImgChannelType(ImageRef.getChannelType()) {
748  checkDeviceFeatureSupported<info::device::image_support>(
749  getDeviceFromHandler(CommandGroupHandlerRef));
750  }
751 #endif
752 
753  /* -- common interface members -- */
754 
755  // operator == and != need to be defined only for host application as per the
756  // SYCL spec 1.2.1
757 #ifndef __SYCL_DEVICE_ONLY__
758  bool operator==(const image_accessor &Rhs) const { return Rhs.impl == impl; }
759 #else
760  // The operator with __SYCL_DEVICE_ONLY__ need to be declared for compilation
761  // of host application with device compiler.
762  // Usage of this operator inside the kernel code will give a runtime failure.
763  bool operator==(const image_accessor &Rhs) const;
764 #endif
765 
766  bool operator!=(const image_accessor &Rhs) const { return !(Rhs == *this); }
767 
768  // get_count() method : Returns the number of elements of the SYCL image this
769  // SYCL accessor is accessing.
770  //
771  // get_range() method : Returns a range object which represents the number of
772  // elements of dataT per dimension that this accessor may access.
773  // The range object returned must equal to the range of the image this
774  // accessor is associated with.
775 
776 #ifdef __SYCL_DEVICE_ONLY__
777 
778  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
779  size_t get_count() const { return size(); }
780  size_t size() const noexcept { return get_range<Dimensions>().size(); }
781 
782  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 1>>
783  range<1> get_range() const {
784  int Range = getRangeInternal();
785  return range<1>(Range);
786  }
787  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 2>>
788  range<2> get_range() const {
789  int2 Range = getRangeInternal();
790  return range<2>(Range[0], Range[1]);
791  }
792  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 3>>
793  range<3> get_range() const {
794  int3 Range = getRangeInternal();
795  return range<3>(Range[0], Range[1], Range[2]);
796  }
797 
798 #else
799  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
800  size_t get_count() const { return size(); };
801  size_t size() const noexcept { return MImageCount; };
802 
803  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
805  return detail::convertToArrayOfN<Dims, 1>(getAccessRange());
806  }
807 
808 #endif
809 
810  // Available only when:
811  // (accessTarget == access::target::image && accessMode == access::mode::read)
812  // || (accessTarget == access::target::host_image && ( accessMode ==
813  // access::mode::read || accessMode == access::mode::read_write))
814  template <typename CoordT, int Dims = Dimensions,
815  typename = detail::enable_if_t<
816  (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
818  ((IsImageAcc && IsImageAccessReadOnly) ||
819  (IsHostImageAcc && IsImageAccessAnyRead))>>
820  DataT read(const CoordT &Coords) const {
821 #ifdef __SYCL_DEVICE_ONLY__
822  return __invoke__ImageRead<DataT, OCLImageTy, CoordT>(MImageObj, Coords);
823 #else
824  sampler Smpl(coordinate_normalization_mode::unnormalized,
825  addressing_mode::none, filtering_mode::nearest);
826  return read<CoordT, Dims>(Coords, Smpl);
827 #endif
828  }
829 
830  // Available only when:
831  // (accessTarget == access::target::image && accessMode == access::mode::read)
832  // || (accessTarget == access::target::host_image && ( accessMode ==
833  // access::mode::read || accessMode == access::mode::read_write))
834  template <typename CoordT, int Dims = Dimensions,
835  typename = detail::enable_if_t<
836  (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
837  ((IsImageAcc && IsImageAccessReadOnly) ||
838  (IsHostImageAcc && IsImageAccessAnyRead))>>
839  DataT read(const CoordT &Coords, const sampler &Smpl) const {
840 #ifdef __SYCL_DEVICE_ONLY__
841  return __invoke__ImageReadSampler<DataT, OCLImageTy, CoordT>(
842  MImageObj, Coords, Smpl.impl.m_Sampler);
843 #else
844  return imageReadSamplerHostImpl<CoordT, DataT>(
845  Coords, Smpl, getAccessRange() /*Image Range*/,
846  getOffset() /*Image Pitch*/, MImgChannelType, MImgChannelOrder,
847  AccessorBaseHost::getPtr() /*ptr to image*/,
848  AccessorBaseHost::getElemSize());
849 #endif
850  }
851 
852  // Available only when:
853  // (accessTarget == access::target::image && (accessMode ==
854  // access::mode::write || accessMode == access::mode::discard_write)) ||
855  // (accessTarget == access::target::host_image && (accessMode ==
856  // access::mode::write || accessMode == access::mode::discard_write ||
857  // accessMode == access::mode::read_write))
858  template <typename CoordT, int Dims = Dimensions,
859  typename = detail::enable_if_t<
860  (Dims > 0) && (detail::is_genint<CoordT>::value) &&
862  ((IsImageAcc && IsImageAccessWriteOnly) ||
863  (IsHostImageAcc && IsImageAccessAnyWrite))>>
864  void write(const CoordT &Coords, const DataT &Color) const {
865 #ifdef __SYCL_DEVICE_ONLY__
866  __invoke__ImageWrite<OCLImageTy, CoordT, DataT>(MImageObj, Coords, Color);
867 #else
868  imageWriteHostImpl(Coords, Color, getOffset() /*ImagePitch*/,
869  AccessorBaseHost::getElemSize(), MImgChannelType,
870  MImgChannelOrder,
871  AccessorBaseHost::getPtr() /*Ptr to Image*/);
872 #endif
873  }
874 };
875 
876 template <typename DataT, int Dimensions, access::mode AccessMode,
879 
880  static_assert(Dimensions < 3,
881  "Image slice cannot have more then 2 dimensions");
882 
883  constexpr static int AdjustedDims = (Dimensions == 2) ? 4 : Dimensions + 1;
884 
885  template <typename CoordT,
886  typename CoordElemType =
888  sycl::vec<CoordElemType, AdjustedDims>
889  getAdjustedCoords(const CoordT &Coords) const {
890  CoordElemType LastCoord = 0;
891 
892  if (std::is_same<float, CoordElemType>::value) {
893  sycl::vec<int, Dimensions + 1> Size = MBaseAcc.getRangeInternal();
894  LastCoord =
895  MIdx / static_cast<float>(Size.template swizzle<Dimensions>());
896  } else {
897  LastCoord = MIdx;
898  }
899 
900  sycl::vec<CoordElemType, Dimensions> LeftoverCoords{LastCoord};
901  sycl::vec<CoordElemType, AdjustedDims> AdjustedCoords{Coords,
902  LeftoverCoords};
903  return AdjustedCoords;
904  }
905 
906 public:
908  accessor<DataT, Dimensions, AccessMode, access::target::image_array,
910  BaseAcc,
911  size_t Idx)
912  : MBaseAcc(BaseAcc), MIdx(Idx) {}
913 
914  template <typename CoordT, int Dims = Dimensions,
915  typename = detail::enable_if_t<
917  DataT read(const CoordT &Coords) const {
918  return MBaseAcc.read(getAdjustedCoords(Coords));
919  }
920 
921  template <typename CoordT, int Dims = Dimensions,
922  typename = detail::enable_if_t<
924  DataT read(const CoordT &Coords, const sampler &Smpl) const {
925  return MBaseAcc.read(getAdjustedCoords(Coords), Smpl);
926  }
927 
928  template <typename CoordT, int Dims = Dimensions,
929  typename = detail::enable_if_t<
931  void write(const CoordT &Coords, const DataT &Color) const {
932  return MBaseAcc.write(getAdjustedCoords(Coords), Color);
933  }
934 
935 #ifdef __SYCL_DEVICE_ONLY__
936  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
937  size_t get_count() const { return size(); }
938  size_t size() const noexcept { return get_range<Dimensions>().size(); }
939 
940  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 1>>
941  range<1> get_range() const {
942  int2 Count = MBaseAcc.getRangeInternal();
943  return range<1>(Count.x());
944  }
945  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 2>>
946  range<2> get_range() const {
947  int3 Count = MBaseAcc.getRangeInternal();
948  return range<2>(Count.x(), Count.y());
949  }
950 
951 #else
952 
953  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
954  size_t get_count() const { return size(); }
955  size_t size() const noexcept {
956  return MBaseAcc.MImageCount / MBaseAcc.getAccessRange()[Dimensions];
957  }
958 
959  template <int Dims = Dimensions,
960  typename = detail::enable_if_t<(Dims == 1 || Dims == 2)>>
962  return detail::convertToArrayOfN<Dims, 1>(MBaseAcc.getAccessRange());
963  }
964 
965 #endif
966 
967 private:
968  size_t MIdx;
969  accessor<DataT, Dimensions, AccessMode, access::target::image_array,
971  MBaseAcc;
972 };
973 
974 } // namespace detail
975 
981 template <typename DataT, int Dimensions, access::mode AccessMode,
983  typename PropertyListT>
985 #ifndef __SYCL_DEVICE_ONLY__
986  public detail::AccessorBaseHost,
987 #endif
988  public detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
989  IsPlaceholder, PropertyListT>,
990  public detail::OwnerLessBase<
991  accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
992  PropertyListT>> {
993 protected:
994  static_assert((AccessTarget == access::target::global_buffer ||
995  AccessTarget == access::target::constant_buffer ||
996  AccessTarget == access::target::host_buffer),
997  "Expected buffer type");
998 
999  static_assert((AccessTarget == access::target::global_buffer ||
1000  AccessTarget == access::target::host_buffer) ||
1001  (AccessTarget == access::target::constant_buffer &&
1002  AccessMode == access::mode::read),
1003  "Access mode can be only read for constant buffers");
1004 
1005  static_assert(detail::IsPropertyListT<PropertyListT>::value,
1006  "PropertyListT must be accessor_property_list");
1007 
1008  using AccessorCommonT =
1009  detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
1010  IsPlaceholder, PropertyListT>;
1011 
1012  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
1013 
1014  using AccessorCommonT::AS;
1015  // Cannot do "using AccessorCommonT::Flag" as it doesn't work with g++ as host
1016  // compiler, for some reason.
1017  static constexpr bool IsAccessAnyWrite = AccessorCommonT::IsAccessAnyWrite;
1018  static constexpr bool IsAccessReadOnly = AccessorCommonT::IsAccessReadOnly;
1019  static constexpr bool IsConstantBuf = AccessorCommonT::IsConstantBuf;
1020  static constexpr bool IsGlobalBuf = AccessorCommonT::IsGlobalBuf;
1021  static constexpr bool IsHostBuf = AccessorCommonT::IsHostBuf;
1022  static constexpr bool IsPlaceH = AccessorCommonT::IsPlaceH;
1023  template <int Dims>
1024  using AccessorSubscript =
1025  typename AccessorCommonT::template AccessorSubscript<Dims>;
1026 
1027  using ConcreteASPtrType = typename detail::DecoratedType<DataT, AS>::type *;
1028 
1029  using RefType = detail::const_if_const_AS<AS, DataT> &;
1030  using ConstRefType = const DataT &;
1031  using PtrType = detail::const_if_const_AS<AS, DataT> *;
1032 
1033  template <int Dims = Dimensions> size_t getLinearIndex(id<Dims> Id) const {
1034 
1035  size_t Result = 0;
1036  detail::dim_loop<Dims>([&, this](size_t I) {
1037  Result = Result * getMemoryRange()[I] + Id[I];
1038  // We've already adjusted for the accessor's offset in the __init, so
1039  // don't include it here in case of device.
1040 #ifndef __SYCL_DEVICE_ONLY__
1041  if constexpr (!(PropertyListT::template has_property<
1043  Result += getOffset()[I];
1044  }
1045 #endif // __SYCL_DEVICE_ONLY__
1046  });
1047 
1048  return Result;
1049  }
1050 
1051  template <typename T, int Dims>
1052  struct IsSameAsBuffer
1053  : std::bool_constant<std::is_same<T, DataT>::value && (Dims > 0) &&
1054  (Dims == Dimensions)> {};
1055 
1056  static access::mode getAdjustedMode(const PropertyListT &PropertyList) {
1057  access::mode AdjustedMode = AccessMode;
1058 
1059  if (PropertyList.template has_property<property::no_init>() ||
1060  PropertyList.template has_property<property::noinit>()) {
1061  if (AdjustedMode == access::mode::write) {
1062  AdjustedMode = access::mode::discard_write;
1063  } else if (AdjustedMode == access::mode::read_write) {
1064  AdjustedMode = access::mode::discard_read_write;
1065  }
1066  }
1067 
1068  return AdjustedMode;
1069  }
1070 
1071  template <typename TagT>
1072  struct IsValidTag
1073  : std::disjunction<
1074  std::is_same<TagT, mode_tag_t<AccessMode>>,
1075  std::is_same<TagT, mode_target_tag_t<AccessMode, AccessTarget>>> {};
1076 
1077 #ifdef __SYCL_DEVICE_ONLY__
1078 
1079  id<AdjustedDim> &getOffset() { return impl.Offset; }
1080  range<AdjustedDim> &getAccessRange() { return impl.AccessRange; }
1081  range<AdjustedDim> &getMemoryRange() { return impl.MemRange; }
1082 
1083  const id<AdjustedDim> &getOffset() const { return impl.Offset; }
1084  const range<AdjustedDim> &getAccessRange() const { return impl.AccessRange; }
1085  const range<AdjustedDim> &getMemoryRange() const { return impl.MemRange; }
1086 
1087  detail::AccessorImplDevice<AdjustedDim> impl;
1088 
1089  union {
1090  ConcreteASPtrType MData;
1091  };
1092 
1093  // TODO replace usages with getQualifiedPtr
1094  const ConcreteASPtrType getNativeImageObj() const { return MData; }
1095 
1096  void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
1097  range<AdjustedDim> MemRange, id<AdjustedDim> Offset) {
1098  MData = Ptr;
1099  detail::dim_loop<AdjustedDim>([&, this](size_t I) {
1100  if constexpr (!(PropertyListT::template has_property<
1102  getOffset()[I] = Offset[I];
1103  }
1104  getAccessRange()[I] = AccessRange[I];
1105  getMemoryRange()[I] = MemRange[I];
1106  });
1107 
1108  // Adjust for offsets as that part is invariant for all invocations of
1109  // operator[]. Will have to re-adjust in get_pointer.
1110  MData += getTotalOffset();
1111  }
1112 
1113  // __init variant used by the device compiler for ESIMD kernels.
1114  // TODO In ESIMD accessors usage is limited for now - access range, mem
1115  // range and offset are not supported.
1116  void __init_esimd(ConcreteASPtrType Ptr) { MData = Ptr; }
1117 
1118  ConcreteASPtrType getQualifiedPtr() const noexcept { return MData; }
1119 
1120  template <typename DataT_, int Dimensions_, access::mode AccessMode_,
1121  access::target AccessTarget_, access::placeholder IsPlaceholder_,
1122  typename PropertyListT_>
1123  friend class accessor;
1124 
1125 #ifndef __SYCL_DEVICE_ONLY__
1126  using AccessorBaseHost::impl;
1127 #endif
1128 
1129 public:
1130  // Default constructor for objects later initialized with __init member.
1131  accessor()
1132  : impl({}, detail::InitializedVal<AdjustedDim, range>::template get<0>(),
1133  detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
1134 
1135 #else
1136  accessor(const detail::AccessorImplPtr &Impl)
1137  : detail::AccessorBaseHost{Impl} {}
1138 
1139  id<3> &getOffset() {
1140  if constexpr (IsHostBuf)
1141  return MAccData->MOffset;
1142  else
1143  return AccessorBaseHost::getOffset();
1144  }
1145 
1146  range<3> &getAccessRange() { return AccessorBaseHost::getAccessRange(); }
1147  range<3> &getMemoryRange() {
1148  if constexpr (IsHostBuf)
1149  return MAccData->MMemoryRange;
1150  else
1151  return AccessorBaseHost::getMemoryRange();
1152  }
1153  void *getPtr() { return AccessorBaseHost::getPtr(); }
1154 
1155  const id<3> &getOffset() const {
1156  if constexpr (IsHostBuf)
1157  return MAccData->MOffset;
1158  else
1159  return AccessorBaseHost::getOffset();
1160  }
1161  const range<3> &getAccessRange() const {
1162  return AccessorBaseHost::getAccessRange();
1163  }
1164  const range<3> &getMemoryRange() const {
1165  if constexpr (IsHostBuf)
1166  return MAccData->MMemoryRange;
1167  else
1168  return AccessorBaseHost::getMemoryRange();
1169  }
1170 
1171  void *getPtr() const { return AccessorBaseHost::getPtr(); }
1172 
1173  void initHostAcc() { MAccData = &getAccData(); }
1174 
1175  // The function references helper methods required by GDB pretty-printers
1176  void GDBMethodsAnchor() {
1177 #ifndef NDEBUG
1178  const auto *this_const = this;
1179  (void)getMemoryRange();
1180  (void)this_const->getMemoryRange();
1181  (void)getOffset();
1182  (void)this_const->getOffset();
1183  (void)getPtr();
1184  (void)this_const->getPtr();
1185  (void)getAccessRange();
1186  (void)this_const->getAccessRange();
1187 #endif
1188  }
1189 
1190  detail::AccHostDataT *MAccData = nullptr;
1191 
1192  char padding[sizeof(detail::AccessorImplDevice<AdjustedDim>) +
1193  sizeof(PtrType) - sizeof(detail::AccessorBaseHost) -
1194  sizeof(MAccData)];
1195 
1196  PtrType getQualifiedPtr() const noexcept {
1197  if constexpr (IsHostBuf)
1198  return reinterpret_cast<PtrType>(MAccData->MData);
1199  else
1200  return reinterpret_cast<PtrType>(AccessorBaseHost::getPtr());
1201  }
1202 
1203 public:
1204  accessor()
1205  : AccessorBaseHost(
1206  /*Offset=*/{0, 0, 0}, /*AccessRange=*/{0, 0, 0},
1207  /*MemoryRange=*/{0, 0, 0},
1208  /*AccessMode=*/getAdjustedMode({}),
1209  /*SYCLMemObject=*/nullptr, /*Dims=*/0, /*ElemSize=*/0,
1210  /*IsPlaceH=*/true,
1211  /*OffsetInBytes=*/0, /*IsSubBuffer=*/false, /*PropertyList=*/{}){};
1212 
1213 #endif // __SYCL_DEVICE_ONLY__
1214 
1215 private:
1216  friend class sycl::stream;
1217  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
1218 
1219  template <class T>
1220  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
1221 
1222 public:
1223  // 4.7.6.9.1. Interface for buffer command accessors
1224  // value_type is defined as const DataT for read_only accessors, DataT
1225  // otherwise
1226  using value_type = typename std::conditional<AccessMode == access_mode::read,
1227  const DataT, DataT>::type;
1228  using reference = value_type &;
1229  using const_reference = const DataT &;
1230 
1231  template <access::decorated IsDecorated>
1232  using accessor_ptr =
1233  std::conditional_t<AccessTarget == access::target::device,
1234  global_ptr<value_type, IsDecorated>, value_type *>;
1235 
1236  using iterator = typename detail::accessor_iterator<value_type, Dimensions>;
1237  using const_iterator =
1238  typename detail::accessor_iterator<const value_type, Dimensions>;
1239  using reverse_iterator = std::reverse_iterator<iterator>;
1240  using const_reverse_iterator = std::reverse_iterator<const_iterator>;
1241  using difference_type =
1242  typename std::iterator_traits<iterator>::difference_type;
1243 
1244  // The list of accessor constructors with their arguments
1245  // -------+---------+-------+----+-----+--------------
1246  // Dimensions = 0
1247  // -------+---------+-------+----+-----+--------------
1248  // buffer | | | | | property_list
1249  // buffer | handler | | | | property_list
1250  // -------+---------+-------+----+-----+--------------
1251  // Dimensions >= 1
1252  // -------+---------+-------+----+-----+--------------
1253  // buffer | | | | | property_list
1254  // buffer | | | | tag | property_list
1255  // buffer | handler | | | | property_list
1256  // buffer | handler | | | tag | property_list
1257  // buffer | | range | | | property_list
1258  // buffer | | range | | tag | property_list
1259  // buffer | handler | range | | | property_list
1260  // buffer | handler | range | | tag | property_list
1261  // buffer | | range | id | | property_list
1262  // buffer | | range | id | tag | property_list
1263  // buffer | handler | range | id | | property_list
1264  // buffer | handler | range | id | tag | property_list
1265  // -------+---------+-------+----+-----+--------------
1266 
1267 public:
1268  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1269  typename detail::enable_if_t<
1270  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1271  std::is_same<T, DataT>::value && Dims == 0 &&
1272  (IsHostBuf || (IsGlobalBuf || IsConstantBuf))> * = nullptr>
1273  accessor(
1274  buffer<T, 1, AllocatorT> &BufferRef,
1275  const property_list &PropertyList = {},
1276  const detail::code_location CodeLoc = detail::code_location::current())
1277 #ifdef __SYCL_DEVICE_ONLY__
1278  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1279  (void)PropertyList;
1280 #else
1281  : AccessorBaseHost(
1282  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1283  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1284  getAdjustedMode(PropertyList),
1285  detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
1286  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1287  PropertyList) {
1288  preScreenAccessor(BufferRef.size(), PropertyList);
1289  if (!AccessorBaseHost::isPlaceholder())
1290  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1291  initHostAcc();
1293  detail::AccessorBaseHost::impl.get(),
1294  AccessTarget, AccessMode, CodeLoc);
1295  GDBMethodsAnchor();
1296 #endif
1297  }
1298 
1299  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1300  typename... PropTypes,
1301  typename detail::enable_if_t<
1302  detail::IsCxPropertyList<PropertyListT>::value &&
1303  std::is_same<T, DataT>::value && Dims == 0 &&
1304  (IsHostBuf || (IsGlobalBuf || IsConstantBuf))> * = nullptr>
1305  accessor(
1306  buffer<T, 1, AllocatorT> &BufferRef,
1307  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1308  {},
1309  const detail::code_location CodeLoc = detail::code_location::current())
1310 #ifdef __SYCL_DEVICE_ONLY__
1311  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1312  (void)PropertyList;
1313 #else
1314  : AccessorBaseHost(
1315  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1316  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1317  getAdjustedMode(PropertyList),
1318  detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
1319  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1320  PropertyList) {
1321  preScreenAccessor(BufferRef.size(), PropertyList);
1322  if (!AccessorBaseHost::isPlaceholder())
1323  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1324  initHostAcc();
1326  detail::AccessorBaseHost::impl.get(),
1327  AccessTarget, AccessMode, CodeLoc);
1328  GDBMethodsAnchor();
1329 #endif
1330  }
1331 
1332  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1333  typename = typename detail::enable_if_t<
1334  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1335  std::is_same<T, DataT>::value && (Dims == 0) &&
1336  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1337  accessor(
1338  buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1339  const property_list &PropertyList = {},
1340  const detail::code_location CodeLoc = detail::code_location::current())
1341 #ifdef __SYCL_DEVICE_ONLY__
1342  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1343  (void)CommandGroupHandler;
1344  (void)PropertyList;
1345  }
1346 #else
1347  : AccessorBaseHost(
1348  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1349  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1350  getAdjustedMode(PropertyList),
1351  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1352  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1353  preScreenAccessor(BufferRef.size(), PropertyList);
1354  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1355  initHostAcc();
1357  detail::AccessorBaseHost::impl.get(),
1358  AccessTarget, AccessMode, CodeLoc);
1359  GDBMethodsAnchor();
1360  }
1361 #endif
1362 
1363  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1364  typename... PropTypes,
1365  typename = typename detail::enable_if_t<
1366  detail::IsCxPropertyList<PropertyListT>::value &&
1367  std::is_same<T, DataT>::value && (Dims == 0) &&
1368  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1369  accessor(
1370  buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1371  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1372  {},
1373  const detail::code_location CodeLoc = detail::code_location::current())
1374 #ifdef __SYCL_DEVICE_ONLY__
1375  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1376  (void)CommandGroupHandler;
1377  (void)PropertyList;
1378  }
1379 #else
1380  : AccessorBaseHost(
1381  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1382  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1383  getAdjustedMode(PropertyList),
1384  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1385  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1386  preScreenAccessor(BufferRef.size(), PropertyList);
1387  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1388  initHostAcc();
1390  detail::AccessorBaseHost::impl.get(),
1391  AccessTarget, AccessMode, CodeLoc);
1392  GDBMethodsAnchor();
1393  }
1394 #endif
1395 
1396  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1397  typename = detail::enable_if_t<
1398  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1399  IsSameAsBuffer<T, Dims>::value &&
1400  (IsHostBuf || (IsGlobalBuf || IsConstantBuf))>>
1401  accessor(
1402  buffer<T, Dims, AllocatorT> &BufferRef,
1403  const property_list &PropertyList = {},
1404  const detail::code_location CodeLoc = detail::code_location::current())
1405 #ifdef __SYCL_DEVICE_ONLY__
1406  : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1407  (void)PropertyList;
1408  }
1409 #else
1410  : AccessorBaseHost(
1411  /*Offset=*/{0, 0, 0},
1412  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1413  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1414  getAdjustedMode(PropertyList),
1415  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1416  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1417  PropertyList) {
1418  preScreenAccessor(BufferRef.size(), PropertyList);
1419  if (!AccessorBaseHost::isPlaceholder())
1420  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1421  initHostAcc();
1423  detail::AccessorBaseHost::impl.get(),
1424  AccessTarget, AccessMode, CodeLoc);
1425  GDBMethodsAnchor();
1426  }
1427 #endif
1428 
1429  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1430  typename... PropTypes,
1431  typename = detail::enable_if_t<
1432  detail::IsCxPropertyList<PropertyListT>::value &&
1433  IsSameAsBuffer<T, Dims>::value &&
1434  (IsHostBuf || (IsGlobalBuf || IsConstantBuf))>>
1435  accessor(
1436  buffer<T, Dims, AllocatorT> &BufferRef,
1437  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1438  {},
1439  const detail::code_location CodeLoc = detail::code_location::current())
1440 #ifdef __SYCL_DEVICE_ONLY__
1441  : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1442  (void)PropertyList;
1443  }
1444 #else
1445  : AccessorBaseHost(
1446  /*Offset=*/{0, 0, 0},
1447  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1448  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1449  getAdjustedMode(PropertyList),
1450  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1451  IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1452  PropertyList) {
1453  preScreenAccessor(BufferRef.size(), PropertyList);
1454  if (!AccessorBaseHost::isPlaceholder())
1455  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1456  initHostAcc();
1458  detail::AccessorBaseHost::impl.get(),
1459  AccessTarget, AccessMode, CodeLoc);
1460  GDBMethodsAnchor();
1461  }
1462 #endif
1463 
1464  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1465  typename TagT,
1466  typename = detail::enable_if_t<
1467  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1468  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1469  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1470  accessor(
1471  buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1472  const property_list &PropertyList = {},
1473  const detail::code_location CodeLoc = detail::code_location::current())
1474  : accessor(BufferRef, PropertyList, CodeLoc) {
1475  adjustAccPropsInBuf(BufferRef);
1476  }
1477 
1478  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1479  typename TagT, typename... PropTypes,
1480  typename = detail::enable_if_t<
1481  detail::IsCxPropertyList<PropertyListT>::value &&
1482  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1483  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1484  accessor(
1485  buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1486  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1487  {},
1488  const detail::code_location CodeLoc = detail::code_location::current())
1489  : accessor(BufferRef, PropertyList, CodeLoc) {
1490  adjustAccPropsInBuf(BufferRef);
1491  }
1492 
1493  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1494  typename = detail::enable_if_t<
1495  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1496  IsSameAsBuffer<T, Dims>::value &&
1497  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1498  accessor(
1499  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1500  const property_list &PropertyList = {},
1501  const detail::code_location CodeLoc = detail::code_location::current())
1502 #ifdef __SYCL_DEVICE_ONLY__
1503  : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1504  (void)CommandGroupHandler;
1505  (void)PropertyList;
1506  }
1507 #else
1508  : AccessorBaseHost(
1509  /*Offset=*/{0, 0, 0},
1510  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1511  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1512  getAdjustedMode(PropertyList),
1513  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1514  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1515  preScreenAccessor(BufferRef.size(), PropertyList);
1516  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1517  initHostAcc();
1519  detail::AccessorBaseHost::impl.get(),
1520  AccessTarget, AccessMode, CodeLoc);
1521  GDBMethodsAnchor();
1522  }
1523 #endif
1524 
1525  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1526  typename... PropTypes,
1527  typename = detail::enable_if_t<
1528  detail::IsCxPropertyList<PropertyListT>::value &&
1529  IsSameAsBuffer<T, Dims>::value &&
1530  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1531  accessor(
1532  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1533  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1534  {},
1535  const detail::code_location CodeLoc = detail::code_location::current())
1536 #ifdef __SYCL_DEVICE_ONLY__
1537  : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1538  (void)CommandGroupHandler;
1539  (void)PropertyList;
1540  }
1541 #else
1542  : AccessorBaseHost(
1543  /*Offset=*/{0, 0, 0},
1544  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1545  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1546  getAdjustedMode(PropertyList),
1547  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1548  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1549  preScreenAccessor(BufferRef.size(), PropertyList);
1550  initHostAcc();
1551  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1553  detail::AccessorBaseHost::impl.get(),
1554  AccessTarget, AccessMode, CodeLoc);
1555  GDBMethodsAnchor();
1556  }
1557 #endif
1558 
1559  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1560  typename TagT,
1561  typename = detail::enable_if_t<
1562  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1563  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1564  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1565  accessor(
1566  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1567  TagT, const property_list &PropertyList = {},
1568  const detail::code_location CodeLoc = detail::code_location::current())
1569  : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1570  adjustAccPropsInBuf(BufferRef);
1571  }
1572 
1573  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1574  typename TagT, typename... PropTypes,
1575  typename = detail::enable_if_t<
1576  detail::IsCxPropertyList<PropertyListT>::value &&
1577  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1578  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1579  accessor(
1580  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1581  TagT,
1582  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1583  {},
1584  const detail::code_location CodeLoc = detail::code_location::current())
1585  : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1586  adjustAccPropsInBuf(BufferRef);
1587  }
1588 
1589  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1590  typename = detail::enable_if_t<
1591  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1592  IsSameAsBuffer<T, Dims>::value &&
1593  (IsHostBuf || (IsGlobalBuf || IsConstantBuf))>>
1594  accessor(
1595  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1596  const property_list &PropertyList = {},
1597  const detail::code_location CodeLoc = detail::code_location::current())
1598  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1599 
1600  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1601  typename... PropTypes,
1602  typename = detail::enable_if_t<
1603  detail::IsCxPropertyList<PropertyListT>::value &&
1604  IsSameAsBuffer<T, Dims>::value &&
1605  (IsHostBuf || (IsGlobalBuf || IsConstantBuf))>>
1606  accessor(
1607  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1608  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1609  {},
1610  const detail::code_location CodeLoc = detail::code_location::current())
1611  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1612 
1613  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1614  typename TagT,
1615  typename = detail::enable_if_t<
1616  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1617  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1618  (IsGlobalBuf || IsConstantBuf)>>
1619  accessor(
1620  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1621  TagT, const property_list &PropertyList = {},
1622  const detail::code_location CodeLoc = detail::code_location::current())
1623  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1624  adjustAccPropsInBuf(BufferRef);
1625  }
1626 
1627  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1628  typename TagT, typename... PropTypes,
1629  typename = detail::enable_if_t<
1630  detail::IsCxPropertyList<PropertyListT>::value &&
1631  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1632  (IsGlobalBuf || IsConstantBuf)>>
1633  accessor(
1634  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1635  TagT,
1636  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1637  {},
1638  const detail::code_location CodeLoc = detail::code_location::current())
1639  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1640  adjustAccPropsInBuf(BufferRef);
1641  }
1642 
1643  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1644  typename = detail::enable_if_t<
1645  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1646  IsSameAsBuffer<T, Dims>::value &&
1647  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1648  accessor(
1649  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1650  range<Dimensions> AccessRange, const property_list &PropertyList = {},
1651  const detail::code_location CodeLoc = detail::code_location::current())
1652  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1653  CodeLoc) {}
1654 
1655  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1656  typename... PropTypes,
1657  typename = detail::enable_if_t<
1658  detail::IsCxPropertyList<PropertyListT>::value &&
1659  IsSameAsBuffer<T, Dims>::value &&
1660  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1661  accessor(
1662  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1663  range<Dimensions> AccessRange,
1664  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1665  {},
1666  const detail::code_location CodeLoc = detail::code_location::current())
1667  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1668  CodeLoc) {}
1669 
1670  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1671  typename TagT,
1672  typename = detail::enable_if_t<
1673  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1674  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1675  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1676  accessor(
1677  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1678  range<Dimensions> AccessRange, TagT,
1679  const property_list &PropertyList = {},
1680  const detail::code_location CodeLoc = detail::code_location::current())
1681  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1682  CodeLoc) {
1683  adjustAccPropsInBuf(BufferRef);
1684  }
1685 
1686  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1687  typename TagT, typename... PropTypes,
1688  typename = detail::enable_if_t<
1689  detail::IsCxPropertyList<PropertyListT>::value &&
1690  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1691  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1692  accessor(
1693  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1694  range<Dimensions> AccessRange, TagT,
1695  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1696  {},
1697  const detail::code_location CodeLoc = detail::code_location::current())
1698  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1699  CodeLoc) {
1700  adjustAccPropsInBuf(BufferRef);
1701  }
1702 
1703  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1704  typename = detail::enable_if_t<
1705  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1706  IsSameAsBuffer<T, Dims>::value &&
1707  (IsHostBuf || (IsGlobalBuf || IsConstantBuf))>>
1708  accessor(
1709  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1710  id<Dimensions> AccessOffset, const property_list &PropertyList = {},
1711  const detail::code_location CodeLoc = detail::code_location::current())
1712 #ifdef __SYCL_DEVICE_ONLY__
1713  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1714  (void)PropertyList;
1715  }
1716 #else
1717  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1718  detail::convertToArrayOfN<3, 1>(AccessRange),
1719  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1720  getAdjustedMode(PropertyList),
1721  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1722  sizeof(DataT), IsPlaceH, BufferRef.OffsetInBytes,
1723  BufferRef.IsSubBuffer, PropertyList) {
1724  preScreenAccessor(BufferRef.size(), PropertyList);
1725  if (!AccessorBaseHost::isPlaceholder())
1726  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1727  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1728  BufferRef.get_range()))
1729  throw sycl::invalid_object_error(
1730  "accessor with requested offset and range would exceed the bounds of "
1731  "the buffer",
1732  PI_ERROR_INVALID_VALUE);
1733 
1734  initHostAcc();
1736  detail::AccessorBaseHost::impl.get(),
1737  AccessTarget, AccessMode, CodeLoc);
1738  GDBMethodsAnchor();
1739  }
1740 #endif
1741 
1742  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1743  typename... PropTypes,
1744  typename = detail::enable_if_t<
1745  detail::IsCxPropertyList<PropertyListT>::value &&
1746  IsSameAsBuffer<T, Dims>::value &&
1747  (IsHostBuf || (IsGlobalBuf || IsConstantBuf))>>
1748  accessor(
1749  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1750  id<Dimensions> AccessOffset,
1751  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1752  {},
1753  const detail::code_location CodeLoc = detail::code_location::current())
1754 #ifdef __SYCL_DEVICE_ONLY__
1755  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1756  (void)PropertyList;
1757  }
1758 #else
1759  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1760  detail::convertToArrayOfN<3, 1>(AccessRange),
1761  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1762  getAdjustedMode(PropertyList),
1763  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1764  sizeof(DataT), IsPlaceH, BufferRef.OffsetInBytes,
1765  BufferRef.IsSubBuffer, PropertyList) {
1766  preScreenAccessor(BufferRef.size(), PropertyList);
1767  if (!AccessorBaseHost::isPlaceholder())
1768  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1769  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1770  BufferRef.get_range()))
1771  throw sycl::invalid_object_error(
1772  "accessor with requested offset and range would exceed the bounds of "
1773  "the buffer",
1774  PI_ERROR_INVALID_VALUE);
1775 
1776  initHostAcc();
1778  detail::AccessorBaseHost::impl.get(),
1779  AccessTarget, AccessMode, CodeLoc);
1780  GDBMethodsAnchor();
1781  }
1782 #endif
1783 
1784  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1785  typename TagT,
1786  typename = detail::enable_if_t<
1787  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1788  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1789  (IsGlobalBuf || IsConstantBuf)>>
1790  accessor(
1791  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1792  id<Dimensions> AccessOffset, TagT, const property_list &PropertyList = {},
1793  const detail::code_location CodeLoc = detail::code_location::current())
1794  : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
1795  adjustAccPropsInBuf(BufferRef);
1796  }
1797 
1798  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1799  typename TagT, typename... PropTypes,
1800  typename = detail::enable_if_t<
1801  detail::IsCxPropertyList<PropertyListT>::value &&
1802  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1803  (IsGlobalBuf || IsConstantBuf)>>
1804  accessor(
1805  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1806  id<Dimensions> AccessOffset, TagT,
1807  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1808  {},
1809  const detail::code_location CodeLoc = detail::code_location::current())
1810  : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
1811  adjustAccPropsInBuf(BufferRef);
1812  }
1813 
1814  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1815  typename = detail::enable_if_t<
1816  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1817  IsSameAsBuffer<T, Dims>::value &&
1818  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1819  accessor(
1820  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1821  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
1822  const property_list &PropertyList = {},
1823  const detail::code_location CodeLoc = detail::code_location::current())
1824 #ifdef __SYCL_DEVICE_ONLY__
1825  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1826  (void)CommandGroupHandler;
1827  (void)PropertyList;
1828  }
1829 #else
1830  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1831  detail::convertToArrayOfN<3, 1>(AccessRange),
1832  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1833  getAdjustedMode(PropertyList),
1834  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1835  sizeof(DataT), BufferRef.OffsetInBytes,
1836  BufferRef.IsSubBuffer, PropertyList) {
1837  preScreenAccessor(BufferRef.size(), PropertyList);
1838  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1839  BufferRef.get_range()))
1840  throw sycl::invalid_object_error(
1841  "accessor with requested offset and range would exceed the bounds of "
1842  "the buffer",
1843  PI_ERROR_INVALID_VALUE);
1844 
1845  initHostAcc();
1846  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1848  detail::AccessorBaseHost::impl.get(),
1849  AccessTarget, AccessMode, CodeLoc);
1850  GDBMethodsAnchor();
1851  }
1852 #endif
1853 
1854  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1855  typename... PropTypes,
1856  typename = detail::enable_if_t<
1857  detail::IsCxPropertyList<PropertyListT>::value &&
1858  IsSameAsBuffer<T, Dims>::value &&
1859  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1860  accessor(
1861  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1862  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
1863  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1864  {},
1865  const detail::code_location CodeLoc = detail::code_location::current())
1866 #ifdef __SYCL_DEVICE_ONLY__
1867  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1868  (void)CommandGroupHandler;
1869  (void)PropertyList;
1870  }
1871 #else
1872  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1873  detail::convertToArrayOfN<3, 1>(AccessRange),
1874  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1875  getAdjustedMode(PropertyList),
1876  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1877  sizeof(DataT), BufferRef.OffsetInBytes,
1878  BufferRef.IsSubBuffer, PropertyList) {
1879  preScreenAccessor(BufferRef.size(), PropertyList);
1880  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1881  BufferRef.get_range()))
1882  throw sycl::invalid_object_error(
1883  "accessor with requested offset and range would exceed the bounds of "
1884  "the buffer",
1885  PI_ERROR_INVALID_VALUE);
1886 
1887  initHostAcc();
1888  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1890  detail::AccessorBaseHost::impl.get(),
1891  AccessTarget, AccessMode, CodeLoc);
1892  GDBMethodsAnchor();
1893  }
1894 #endif
1895 
1896  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1897  typename TagT,
1898  typename = detail::enable_if_t<
1899  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1900  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1901  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1902  accessor(
1903  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1904  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
1905  const property_list &PropertyList = {},
1906  const detail::code_location CodeLoc = detail::code_location::current())
1907  : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1908  PropertyList, CodeLoc) {
1909  adjustAccPropsInBuf(BufferRef);
1910  }
1911 
1912  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1913  typename TagT, typename... PropTypes,
1914  typename = detail::enable_if_t<
1915  detail::IsCxPropertyList<PropertyListT>::value &&
1916  IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1917  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1918  accessor(
1919  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1920  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
1921  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1922  {},
1923  const detail::code_location CodeLoc = detail::code_location::current())
1924  : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1925  PropertyList, CodeLoc) {
1926  adjustAccPropsInBuf(BufferRef);
1927  }
1928 
1929  template <typename... NewPropsT>
1930  accessor(
1931  const accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
1932  ext::oneapi::accessor_property_list<NewPropsT...>> &Other,
1933  const detail::code_location CodeLoc = detail::code_location::current())
1934 #ifdef __SYCL_DEVICE_ONLY__
1935  : impl(Other.impl)
1936 #else
1937  : detail::AccessorBaseHost(Other)
1938 #endif
1939  {
1940  static_assert(detail::IsCxPropertyList<PropertyListT>::value,
1941  "Conversion is only available for accessor_property_list");
1942  static_assert(
1943  PropertyListT::template areSameCompileTimeProperties<NewPropsT...>(),
1944  "Compile-time-constant properties must be the same");
1945 #ifndef __SYCL_DEVICE_ONLY__
1946  detail::constructorNotification(getMemoryObject(), impl.get(), AccessTarget,
1947  AccessMode, CodeLoc);
1948 #endif
1949  }
1950 
1951  void swap(accessor &other) {
1952  std::swap(impl, other.impl);
1953 #ifndef __SYCL_DEVICE_ONLY__
1954  std::swap(MAccData, other.MAccData);
1955 #endif
1956  }
1957 
1958  bool is_placeholder() const {
1959 #ifdef __SYCL_DEVICE_ONLY__
1960  return false;
1961 #else
1962  return detail::AccessorBaseHost::isPlaceholder();
1963 #endif
1964  }
1965 
1966  size_t get_size() const { return getAccessRange().size() * sizeof(DataT); }
1967 
1968  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1969  size_t get_count() const { return size(); }
1970  size_t size() const noexcept { return getAccessRange().size(); }
1971 
1972  size_t byte_size() const noexcept { return size() * sizeof(DataT); }
1973 
1974  size_t max_size() const noexcept {
1975  return empty() ? 0 : (std::numeric_limits<difference_type>::max)();
1976  }
1977 
1978  bool empty() const noexcept { return size() == 0; }
1979 
1980  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
1981  range<Dimensions> get_range() const {
1982  return detail::convertToArrayOfN<Dimensions, 1>(getAccessRange());
1983  }
1984 
1985  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
1986  id<Dimensions> get_offset() const {
1987  static_assert(
1988  !(PropertyListT::template has_property<
1990  "Accessor has no_offset property, get_offset() can not be used");
1991  return detail::convertToArrayOfN<Dimensions, 0>(getOffset());
1992  }
1993 
1994  template <int Dims = Dimensions, typename RefT = RefType,
1995  typename = detail::enable_if_t<Dims == 0 && (IsAccessAnyWrite ||
1996  IsAccessReadOnly)>>
1997  operator reference() const {
1998  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
1999  return *(getQualifiedPtr() + LinearIndex);
2000  }
2001 
2002  template <int Dims = Dimensions,
2003  typename = detail::enable_if_t<(Dims > 0) && (IsAccessAnyWrite ||
2004  IsAccessReadOnly)>>
2005  reference operator[](id<Dimensions> Index) const {
2006  const size_t LinearIndex = getLinearIndex(Index);
2007  return getQualifiedPtr()[LinearIndex];
2008  }
2009 
2010  template <int Dims = Dimensions>
2011  operator typename detail::enable_if_t<Dims == 0 &&
2012  AccessMode == access::mode::atomic,
2013 #ifdef __ENABLE_USM_ADDR_SPACE__
2014  atomic<DataT>
2015 #else
2016  atomic<DataT, AS>
2017 #endif
2018  >() const {
2019  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
2020  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2021  getQualifiedPtr() + LinearIndex));
2022  }
2023 
2024  template <int Dims = Dimensions>
2025  typename detail::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
2026  atomic<DataT, AS>>
2027  operator[](id<Dimensions> Index) const {
2028  const size_t LinearIndex = getLinearIndex(Index);
2029  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2030  getQualifiedPtr() + LinearIndex));
2031  }
2032 
2033  template <int Dims = Dimensions>
2034  typename detail::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
2035  atomic<DataT, AS>>
2036  operator[](size_t Index) const {
2037  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
2038  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2039  getQualifiedPtr() + LinearIndex));
2040  }
2041  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>
2042  auto operator[](size_t Index) const {
2043  return AccessorSubscript<Dims - 1>(*this, Index);
2044  }
2045 
2046  template <access::target AccessTarget_ = AccessTarget,
2047  typename = detail::enable_if_t<AccessTarget_ ==
2048  access::target::host_buffer>>
2049 #if SYCL_LANGUAGE_VERSION >= 202001
2050  std::add_pointer_t<value_type> get_pointer() const noexcept
2051 #else
2052  DataT *get_pointer() const
2053 #endif
2054  {
2055  return getPointerAdjusted();
2056  }
2057 
2058  template <
2059  access::target AccessTarget_ = AccessTarget,
2060  typename = detail::enable_if_t<AccessTarget_ == access::target::device>>
2061  global_ptr<DataT> get_pointer() const {
2062  return global_ptr<DataT>(getPointerAdjusted());
2063  }
2064 
2065  template <access::target AccessTarget_ = AccessTarget,
2066  typename = detail::enable_if_t<AccessTarget_ ==
2067  access::target::constant_buffer>>
2068  constant_ptr<DataT> get_pointer() const {
2069  return constant_ptr<DataT>(getPointerAdjusted());
2070  }
2071 
2072  // accessor::has_property for runtime properties is only available in host
2073  // code. This restriction is not listed in the core spec and will be added in
2074  // future versions.
2075  template <typename Property>
2076  typename sycl::detail::enable_if_t<
2077  !ext::oneapi::is_compile_time_property<Property>::value, bool>
2078  has_property() const noexcept {
2079 #ifndef __SYCL_DEVICE_ONLY__
2080  return getPropList().template has_property<Property>();
2081 #else
2082  return false;
2083 #endif
2084  }
2085 
2086  // accessor::get_property for runtime properties is only available in host
2087  // code. This restriction is not listed in the core spec and will be added in
2088  // future versions.
2089  template <typename Property,
2090  typename = typename sycl::detail::enable_if_t<
2091  !ext::oneapi::is_compile_time_property<Property>::value>>
2092  Property get_property() const {
2093 #ifndef __SYCL_DEVICE_ONLY__
2094  return getPropList().template get_property<Property>();
2095 #else
2096  return Property();
2097 #endif
2098  }
2099 
2100  template <typename Property>
2101  static constexpr bool has_property(
2102  typename std::enable_if_t<
2103  ext::oneapi::is_compile_time_property<Property>::value> * = 0) {
2104  return PropertyListT::template has_property<Property>();
2105  }
2106 
2107  template <typename Property>
2108  static constexpr auto get_property(
2109  typename std::enable_if_t<
2110  ext::oneapi::is_compile_time_property<Property>::value> * = 0) {
2111  return PropertyListT::template get_property<Property>();
2112  }
2113 
2114  bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
2115  bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
2116 
2117  iterator begin() const noexcept {
2118  return iterator::getBegin(
2119  get_pointer(),
2120  detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
2121  get_offset());
2122  }
2123 
2124  iterator end() const noexcept {
2125  return iterator::getEnd(
2126  get_pointer(),
2127  detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
2128  get_offset());
2129  }
2130 
2131  const_iterator cbegin() const noexcept {
2132  return const_iterator::getBegin(
2133  get_pointer(),
2134  detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
2135  get_offset());
2136  }
2137 
2138  const_iterator cend() const noexcept {
2139  return const_iterator::getEnd(
2140  get_pointer(),
2141  detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
2142  get_offset());
2143  }
2144 
2145  reverse_iterator rbegin() const noexcept { return reverse_iterator(end()); }
2146  reverse_iterator rend() const noexcept { return reverse_iterator(begin()); }
2147 
2148  const_reverse_iterator crbegin() const noexcept {
2149  return const_reverse_iterator(cend());
2150  }
2151  const_reverse_iterator crend() const noexcept {
2152  return const_reverse_iterator(cbegin());
2153  }
2154 
2155 private:
2156 #ifdef __SYCL_DEVICE_ONLY__
2157  size_t getTotalOffset() const noexcept {
2158  size_t TotalOffset = 0;
2159  detail::dim_loop<Dimensions>([&, this](size_t I) {
2160  TotalOffset = TotalOffset * impl.MemRange[I];
2161  if constexpr (!(PropertyListT::template has_property<
2163  TotalOffset += impl.Offset[I];
2164  }
2165  });
2166 
2167  return TotalOffset;
2168  }
2169 #endif
2170 
2171  // supporting function for get_pointer()
2172  // MData has been preadjusted with offset for faster access with []
2173  // but for get_pointer() we must return the original pointer.
2174  // On device, getQualifiedPtr() returns MData, so we need to backjust it.
2175  // On host, getQualifiedPtr() does not return MData, no need to adjust.
2176  auto getPointerAdjusted() const noexcept {
2177 #ifdef __SYCL_DEVICE_ONLY__
2178  return getQualifiedPtr() - getTotalOffset();
2179 #else
2180  return getQualifiedPtr();
2181 #endif
2182  }
2183 
2184  void preScreenAccessor(const size_t elemInBuffer,
2185  const PropertyListT &PropertyList) {
2186  // check device accessor buffer size
2187  if (!IsHostBuf && elemInBuffer == 0)
2188  throw sycl::invalid_object_error(
2189  "SYCL buffer size is zero. To create a device accessor, SYCL "
2190  "buffer size must be greater than zero.",
2191  PI_ERROR_INVALID_VALUE);
2192 
2193  // check that no_init property is compatible with access mode
2194  if (PropertyList.template has_property<property::no_init>() &&
2195  AccessMode == access::mode::read) {
2196  throw sycl::invalid_object_error(
2197  "accessor would cannot be both read_only and no_init",
2198  PI_ERROR_INVALID_VALUE);
2199  }
2200  }
2201 
2202  template <typename BufT, typename... PropTypes>
2203  void adjustAccPropsInBuf(BufT &Buffer) {
2204  if constexpr (PropertyListT::template has_property<
2206  auto location = (PropertyListT::template get_property<
2208  .get_location();
2209  property_list PropList{
2211  Buffer.addOrReplaceAccessorProperties(PropList);
2212  } else {
2213  deleteAccPropsFromBuf(Buffer);
2214  }
2215  }
2216 
2217  template <typename BufT> void deleteAccPropsFromBuf(BufT &Buffer) {
2218  Buffer.deleteAccProps(
2220  }
2221 };
2222 
2223 template <typename DataT, int Dimensions, typename AllocatorT>
2225  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2226  access::placeholder::true_t>;
2227 
2228 template <typename DataT, int Dimensions, typename AllocatorT,
2229  typename... PropsT>
2232  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2233  access::placeholder::true_t,
2235 
2236 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
2239  detail::deduceAccessTarget<Type1, Type1>(target::device),
2240  access::placeholder::true_t>;
2241 
2242 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2243  typename... PropsT>
2247  detail::deduceAccessTarget<Type1, Type1>(target::device),
2248  access::placeholder::true_t,
2250 
2251 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2252  typename Type2>
2255  detail::deduceAccessTarget<Type1, Type2>(target::device),
2256  access::placeholder::true_t>;
2257 
2258 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2259  typename Type2, typename... PropsT>
2263  detail::deduceAccessTarget<Type1, Type2>(target::device),
2264  access::placeholder::true_t,
2266 
2267 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2268  typename Type2, typename Type3>
2271  detail::deduceAccessTarget<Type2, Type3>(target::device),
2272  access::placeholder::true_t>;
2273 
2274 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2275  typename Type2, typename Type3, typename... PropsT>
2279  detail::deduceAccessTarget<Type2, Type3>(target::device),
2280  access::placeholder::true_t,
2282 
2283 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2284  typename Type2, typename Type3, typename Type4>
2287  detail::deduceAccessTarget<Type3, Type4>(target::device),
2288  access::placeholder::true_t>;
2289 
2290 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2291  typename Type2, typename Type3, typename Type4, typename... PropsT>
2295  detail::deduceAccessTarget<Type3, Type4>(target::device),
2296  access::placeholder::true_t,
2298 
2299 template <typename DataT, int Dimensions, typename AllocatorT>
2301  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2302  access::placeholder::false_t>;
2303 
2304 template <typename DataT, int Dimensions, typename AllocatorT,
2305  typename... PropsT>
2308  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
2309  access::placeholder::false_t,
2311 
2312 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
2315  detail::deduceAccessTarget<Type1, Type1>(target::device),
2316  access::placeholder::false_t>;
2317 
2318 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2319  typename... PropsT>
2323  detail::deduceAccessTarget<Type1, Type1>(target::device),
2324  access::placeholder::false_t,
2326 
2327 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2328  typename Type2>
2331  detail::deduceAccessTarget<Type1, Type2>(target::device),
2332  access::placeholder::false_t>;
2333 
2334 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2335  typename Type2, typename... PropsT>
2339  detail::deduceAccessTarget<Type1, Type2>(target::device),
2340  access::placeholder::false_t,
2342 
2343 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2344  typename Type2, typename Type3>
2347  detail::deduceAccessTarget<Type2, Type3>(target::device),
2348  access::placeholder::false_t>;
2349 
2350 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2351  typename Type2, typename Type3, typename... PropsT>
2355  detail::deduceAccessTarget<Type2, Type3>(target::device),
2356  access::placeholder::false_t,
2358 
2359 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2360  typename Type2, typename Type3, typename Type4>
2362  Type4)
2364  detail::deduceAccessTarget<Type3, Type4>(target::device),
2365  access::placeholder::false_t>;
2366 
2367 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2368  typename Type2, typename Type3, typename Type4, typename... PropsT>
2372  detail::deduceAccessTarget<Type3, Type4>(target::device),
2373  access::placeholder::false_t,
2375 
2379 template <typename DataT, int Dimensions, access::mode AccessMode,
2382 #ifndef __SYCL_DEVICE_ONLY__
2384 #endif
2385  public detail::accessor_common<DataT, Dimensions, AccessMode,
2386  access::target::local, IsPlaceholder> {
2387 protected:
2388  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
2389 
2392  access::target::local, IsPlaceholder>;
2393 
2394  using AccessorCommonT::AS;
2395  using AccessorCommonT::IsAccessAnyWrite;
2396  template <int Dims>
2398  typename AccessorCommonT::template AccessorSubscript<
2399  Dims,
2401 
2403 
2406 
2407 #ifdef __SYCL_DEVICE_ONLY__
2409 
2410  sycl::range<AdjustedDim> &getSize() { return impl.MemRange; }
2411  const sycl::range<AdjustedDim> &getSize() const { return impl.MemRange; }
2412 
2413  void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
2414  range<AdjustedDim>, id<AdjustedDim>) {
2415  MData = Ptr;
2416  detail::dim_loop<AdjustedDim>(
2417  [&, this](size_t I) { getSize()[I] = AccessRange[I]; });
2418  }
2419 
2420 public:
2421  // Default constructor for objects later initialized with __init member.
2422  local_accessor_base()
2423  : impl(detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
2424 
2425 protected:
2426  ConcreteASPtrType getQualifiedPtr() const { return MData; }
2427 
2428  ConcreteASPtrType MData;
2429 
2430 #else
2431 public:
2433  : detail::LocalAccessorBaseHost{/*Size*/ sycl::range<3>{0, 0, 0},
2434  /*Dims*/ 0, /*ElemSize*/ sizeof(DataT)} {}
2435 
2436 protected:
2438  : detail::LocalAccessorBaseHost{Impl} {}
2439 
2441  sizeof(PtrType) - sizeof(detail::LocalAccessorBaseHost)];
2442  using detail::LocalAccessorBaseHost::getSize;
2443 
2445  return reinterpret_cast<PtrType>(LocalAccessorBaseHost::getPtr());
2446  }
2447 
2448  void *getPtr() { return detail::LocalAccessorBaseHost::getPtr(); }
2449  void *getPtr() const { return detail::LocalAccessorBaseHost::getPtr(); }
2450  const range<3> &getSize() const {
2451  return detail::LocalAccessorBaseHost::getSize();
2452  }
2453  range<3> &getSize() { return detail::LocalAccessorBaseHost::getSize(); }
2454 
2455  // The function references helper methods required by GDB pretty-printers
2457 #ifndef NDEBUG
2458  const auto *this_const = this;
2459  (void)getSize();
2460  (void)this_const->getSize();
2461  (void)getPtr();
2462  (void)this_const->getPtr();
2463 #endif
2464  }
2465 
2466 #endif // __SYCL_DEVICE_ONLY__
2467 
2468  // Method which calculates linear offset for the ID using Range and Offset.
2469  template <int Dims = AdjustedDim> size_t getLinearIndex(id<Dims> Id) const {
2470  size_t Result = 0;
2471  for (int I = 0; I < Dims; ++I)
2472  Result = Result * getSize()[I] + Id[I];
2473  return Result;
2474  }
2475 
2476  template <class T>
2477  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
2478 
2479 public:
2480  using value_type = DataT;
2481  using reference = DataT &;
2482  using const_reference = const DataT &;
2483 
2484  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 0>>
2486  detail::code_location::current())
2487 #ifdef __SYCL_DEVICE_ONLY__
2488  : impl(range<AdjustedDim>{1}){}
2489 #else
2490  : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT)) {
2491  detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
2492  access::target::local, AccessMode, CodeLoc);
2493  GDBMethodsAnchor();
2494  }
2495 #endif
2496 
2497  template <int Dims = Dimensions,
2498  typename = detail::enable_if_t<Dims == 0>>
2500  const detail::code_location CodeLoc =
2501  detail::code_location::current())
2502 #ifdef __SYCL_DEVICE_ONLY__
2503  : impl(range<AdjustedDim>{1}) {
2504  (void)propList;
2505  }
2506 #else
2507  : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT),
2508  propList) {
2509  detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
2510  access::target::local, AccessMode, CodeLoc);
2511  GDBMethodsAnchor();
2512  }
2513 #endif
2514 
2515  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
2517  range<Dimensions> AllocationSize, handler &,
2518  const detail::code_location CodeLoc = detail::code_location::current())
2519 #ifdef __SYCL_DEVICE_ONLY__
2520  : impl(AllocationSize){}
2521 #else
2522  : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2523  AdjustedDim, sizeof(DataT)) {
2524  detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
2525  access::target::local, AccessMode, CodeLoc);
2526  GDBMethodsAnchor();
2527  }
2528 #endif
2529 
2530  template <int Dims = Dimensions,
2531  typename = detail::enable_if_t<(Dims > 0)>>
2533  const property_list &propList,
2534  const detail::code_location CodeLoc =
2535  detail::code_location::current())
2536 #ifdef __SYCL_DEVICE_ONLY__
2537  : impl(AllocationSize) {
2538  (void)propList;
2539  }
2540 #else
2541  : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2542  AdjustedDim, sizeof(DataT), propList) {
2543  detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
2544  access::target::local, AccessMode, CodeLoc);
2545  GDBMethodsAnchor();
2546  }
2547 #endif
2548 
2549  size_t get_size() const { return getSize().size() * sizeof(DataT); }
2550 
2551  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
2552  size_t get_count() const { return size(); }
2553  size_t size() const noexcept { return getSize().size(); }
2554 
2555  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
2557  return detail::convertToArrayOfN<Dims, 1>(getSize());
2558  }
2559 
2560  template <int Dims = Dimensions,
2562  operator RefType() const {
2563  return *getQualifiedPtr();
2564  }
2565 
2566  template <int Dims = Dimensions,
2567  typename = detail::enable_if_t<(Dims > 0) && IsAccessAnyWrite>>
2568  RefType operator[](id<Dimensions> Index) const {
2569  const size_t LinearIndex = getLinearIndex(Index);
2570  return getQualifiedPtr()[LinearIndex];
2571  }
2572 
2573  template <int Dims = Dimensions,
2575  RefType operator[](size_t Index) const {
2576  return getQualifiedPtr()[Index];
2577  }
2578 
2579  template <int Dims = Dimensions>
2580  operator typename detail::enable_if_t<
2581  Dims == 0 && AccessMode == access::mode::atomic, atomic<DataT, AS>>()
2582  const {
2583  return atomic<DataT, AS>(
2585  }
2586 
2587  template <int Dims = Dimensions>
2588  typename detail::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
2589  atomic<DataT, AS>>
2590  operator[](id<Dimensions> Index) const {
2591  const size_t LinearIndex = getLinearIndex(Index);
2592  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2593  getQualifiedPtr() + LinearIndex));
2594  }
2595 
2596  template <int Dims = Dimensions>
2597  typename detail::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
2598  atomic<DataT, AS>>
2599  operator[](size_t Index) const {
2600  return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2601  getQualifiedPtr() + Index));
2602  }
2603 
2604  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>
2605  typename AccessorCommonT::template AccessorSubscript<
2606  Dims - 1,
2608  operator[](size_t Index) const {
2609  return AccessorSubscript<Dims - 1>(*this, Index);
2610  }
2611 
2613  return local_ptr<DataT>(getQualifiedPtr());
2614  }
2615 
2616  bool operator==(const local_accessor_base &Rhs) const {
2617  return impl == Rhs.impl;
2618  }
2619  bool operator!=(const local_accessor_base &Rhs) const {
2620  return !(*this == Rhs);
2621  }
2622 };
2623 
2624 // TODO: Remove deprecated specialization once no longer needed
2625 template <typename DataT, int Dimensions, access::mode AccessMode,
2628  DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder>
2629  : public local_accessor_base<DataT, Dimensions, AccessMode, IsPlaceholder>,
2630  public detail::OwnerLessBase<
2631  accessor<DataT, Dimensions, AccessMode, access::target::local,
2632  IsPlaceholder>> {
2633 
2634  using local_acc =
2636 
2637  // Use base classes constructors
2638  using local_acc::local_acc;
2639 
2640 #ifdef __SYCL_DEVICE_ONLY__
2641 
2642  // __init needs to be defined within the class not through inheritance.
2643  // Map this function to inherited func.
2644  void __init(typename local_acc::ConcreteASPtrType Ptr,
2645  range<local_acc::AdjustedDim> AccessRange,
2648  local_acc::__init(Ptr, AccessRange, range, id);
2649  }
2650 
2651 public:
2652  // Default constructor for objects later initialized with __init member.
2653  accessor() {
2654  local_acc::impl = detail::InitializedVal<local_acc::AdjustedDim,
2655  range>::template get<0>();
2656  }
2657 
2658 #else
2659 private:
2660  accessor(const detail::AccessorImplPtr &Impl) : local_acc{Impl} {}
2661 #endif
2662 };
2663 
2664 template <typename DataT, int Dimensions = 1>
2666  : public local_accessor_base<DataT, Dimensions, access::mode::read_write,
2667  access::placeholder::false_t>,
2668  public detail::OwnerLessBase<local_accessor<DataT, Dimensions>> {
2669 
2670  using local_acc =
2672  access::placeholder::false_t>;
2673 
2674  // Use base classes constructors
2675  using local_acc::local_acc;
2676 
2677 #ifdef __SYCL_DEVICE_ONLY__
2678 
2679  // __init needs to be defined within the class not through inheritance.
2680  // Map this function to inherited func.
2681  void __init(typename local_acc::ConcreteASPtrType Ptr,
2682  range<local_acc::AdjustedDim> AccessRange,
2685  local_acc::__init(Ptr, AccessRange, range, id);
2686  }
2687 
2688 public:
2689  // Default constructor for objects later initialized with __init member.
2690  local_accessor() {
2691  local_acc::impl = detail::InitializedVal<local_acc::AdjustedDim,
2692  range>::template get<0>();
2693  }
2694 
2695 #else
2696  local_accessor(const detail::AccessorImplPtr &Impl) : local_acc{Impl} {}
2697 #endif
2698 
2699 public:
2700  using value_type = DataT;
2701  using iterator = value_type *;
2702  using const_iterator = const value_type *;
2703  using reverse_iterator = std::reverse_iterator<iterator>;
2704  using const_reverse_iterator = std::reverse_iterator<const_iterator>;
2705  using difference_type =
2706  typename std::iterator_traits<iterator>::difference_type;
2707 
2708  template <access::decorated IsDecorated>
2709  using accessor_ptr = local_ptr<value_type, IsDecorated>;
2710 
2711  void swap(local_accessor &other) { std::swap(this->impl, other.impl); }
2712 
2713  size_t byte_size() const noexcept { return this->size() * sizeof(DataT); }
2714 
2715  size_t max_size() const noexcept {
2716  return empty() ? 0 : (std::numeric_limits<difference_type>::max)();
2717  }
2718 
2719  bool empty() const noexcept { return this->size() == 0; }
2720 
2721  iterator begin() const noexcept {
2722  return &this->operator[](id<Dimensions>());
2723  }
2724  iterator end() const noexcept { return begin() + this->size(); }
2725 
2726  const_iterator cbegin() const noexcept { return const_iterator(begin()); }
2727  const_iterator cend() const noexcept { return const_iterator(end()); }
2728 
2729  reverse_iterator rbegin() const noexcept { return reverse_iterator(end()); }
2730  reverse_iterator rend() const noexcept { return reverse_iterator(begin()); }
2731 
2732  const_reverse_iterator crbegin() const noexcept {
2733  return const_reverse_iterator(end());
2734  }
2735  const_reverse_iterator crend() const noexcept {
2736  return const_reverse_iterator(begin());
2737  }
2738 
2739  template <typename Property> bool has_property() const noexcept {
2740 #ifndef __SYCL_DEVICE_ONLY__
2741  return this->getPropList().template has_property<Property>();
2742 #else
2743  return false;
2744 #endif
2745  }
2746 
2747  template <typename Property> Property get_property() const {
2748 #ifndef __SYCL_DEVICE_ONLY__
2749  return this->getPropList().template get_property<Property>();
2750 #else
2751  return Property();
2752 #endif
2753  }
2754 };
2755 
2761 template <typename DataT, int Dimensions, access::mode AccessMode,
2764  DataT, Dimensions, AccessMode, access::target::image, IsPlaceholder>
2765  : public detail::image_accessor<DataT, Dimensions, AccessMode,
2766  access::target::image, IsPlaceholder>,
2768  accessor<DataT, Dimensions, AccessMode, access::target::image,
2769  IsPlaceholder>> {
2770 private:
2771  accessor(const detail::AccessorImplPtr &Impl)
2773  access::target::image, IsPlaceholder>{Impl} {}
2774 
2775 public:
2776  template <typename AllocatorT>
2777  accessor(sycl::image<Dimensions, AllocatorT> &Image,
2778  handler &CommandGroupHandler)
2780  access::target::image, IsPlaceholder>(
2781  Image, CommandGroupHandler, Image.getElementSize()) {
2782 #ifndef __SYCL_DEVICE_ONLY__
2783  detail::associateWithHandler(CommandGroupHandler, this,
2784  access::target::image);
2785 #endif
2786  }
2787 
2788  template <typename AllocatorT>
2789  accessor(sycl::image<Dimensions, AllocatorT> &Image,
2790  handler &CommandGroupHandler, const property_list &propList)
2792  access::target::image, IsPlaceholder>(
2793  Image, CommandGroupHandler, Image.getElementSize()) {
2794  (void)propList;
2795 #ifndef __SYCL_DEVICE_ONLY__
2796  detail::associateWithHandler(CommandGroupHandler, this,
2797  access::target::image);
2798 #endif
2799  }
2800 #ifdef __SYCL_DEVICE_ONLY__
2801 private:
2802  using OCLImageTy =
2803  typename detail::opencl_image_type<Dimensions, AccessMode,
2804  access::target::image>::type;
2805 
2806  // Front End requires this method to be defined in the accessor class.
2807  // It does not call the base class's init method.
2808  void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
2809 
2810  // __init variant used by the device compiler for ESIMD kernels.
2811  void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
2812 
2813 public:
2814  // Default constructor for objects later initialized with __init member.
2815  accessor() = default;
2816 #endif
2817 };
2818 
2826 template <typename DataT, int Dimensions, access::mode AccessMode,
2829  access::target::host_image, IsPlaceholder>
2830  : public detail::image_accessor<DataT, Dimensions, AccessMode,
2831  access::target::host_image, IsPlaceholder>,
2832  public detail::OwnerLessBase<
2833  accessor<DataT, Dimensions, AccessMode, access::target::host_image,
2834  IsPlaceholder>> {
2835 public:
2836  template <typename AllocatorT>
2837  accessor(sycl::image<Dimensions, AllocatorT> &Image)
2838  : detail::image_accessor<DataT, Dimensions, AccessMode,
2839  access::target::host_image, IsPlaceholder>(
2840  Image, Image.getElementSize()) {}
2841 
2842  template <typename AllocatorT>
2843  accessor(sycl::image<Dimensions, AllocatorT> &Image,
2844  const property_list &propList)
2845  : detail::image_accessor<DataT, Dimensions, AccessMode,
2846  access::target::host_image, IsPlaceholder>(
2847  Image, Image.getElementSize()) {
2848  (void)propList;
2849  }
2850 };
2851 
2860 template <typename DataT, int Dimensions, access::mode AccessMode,
2863  DataT, Dimensions, AccessMode, access::target::image_array, IsPlaceholder>
2864  : public detail::image_accessor<DataT, Dimensions + 1, AccessMode,
2865  access::target::image, IsPlaceholder>,
2866  public detail::OwnerLessBase<
2867  accessor<DataT, Dimensions, AccessMode, access::target::image_array,
2868  IsPlaceholder>> {
2869 #ifdef __SYCL_DEVICE_ONLY__
2870 private:
2871  using OCLImageTy =
2872  typename detail::opencl_image_type<Dimensions + 1, AccessMode,
2873  access::target::image>::type;
2874 
2875  // Front End requires this method to be defined in the accessor class.
2876  // It does not call the base class's init method.
2877  void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
2878 
2879  // __init variant used by the device compiler for ESIMD kernels.
2880  void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
2881 
2882 public:
2883  // Default constructor for objects later initialized with __init member.
2884  accessor() = default;
2885 #endif
2886 public:
2887  template <typename AllocatorT>
2888  accessor(sycl::image<Dimensions + 1, AllocatorT> &Image,
2889  handler &CommandGroupHandler)
2891  access::target::image, IsPlaceholder>(
2892  Image, CommandGroupHandler, Image.getElementSize()) {
2893 #ifndef __SYCL_DEVICE_ONLY__
2894  detail::associateWithHandler(CommandGroupHandler, this,
2895  access::target::image_array);
2896 #endif
2897  }
2898 
2899  template <typename AllocatorT>
2900  accessor(sycl::image<Dimensions + 1, AllocatorT> &Image,
2901  handler &CommandGroupHandler, const property_list &propList)
2903  access::target::image, IsPlaceholder>(
2904  Image, CommandGroupHandler, Image.getElementSize()) {
2905  (void)propList;
2906 #ifndef __SYCL_DEVICE_ONLY__
2907  detail::associateWithHandler(CommandGroupHandler, this,
2908  access::target::image_array);
2909 #endif
2910  }
2911 
2913  operator[](size_t Index) const {
2915  IsPlaceholder>(*this, Index);
2916  }
2917 };
2918 
2919 template <typename DataT, int Dimensions = 1,
2922  : public accessor<DataT, Dimensions, AccessMode, target::host_buffer,
2923  access::placeholder::false_t>,
2924  public detail::OwnerLessBase<
2925  host_accessor<DataT, Dimensions, AccessMode>> {
2926 protected:
2927  using AccessorT = accessor<DataT, Dimensions, AccessMode, target::host_buffer,
2928  access::placeholder::false_t>;
2929 
2930  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
2931 
2932  template <typename T, int Dims>
2934  : std::bool_constant<std::is_same<T, DataT>::value && (Dims > 0) &&
2935  (Dims == Dimensions)> {};
2936 
2937  void
2938  __init(typename accessor<DataT, Dimensions, AccessMode, target::host_buffer,
2939  access::placeholder::false_t>::ConcreteASPtrType Ptr,
2940  range<AdjustedDim> AccessRange, range<AdjustedDim> MemRange,
2941  id<AdjustedDim> Offset) {
2942  AccessorT::__init(Ptr, AccessRange, MemRange, Offset);
2943  }
2944 
2945 public:
2947 
2948  // The list of host_accessor constructors with their arguments
2949  // -------+---------+-------+----+----------+--------------
2950  // Dimensions = 0
2951  // -------+---------+-------+----+----------+--------------
2952  // buffer | | | | | property_list
2953  // buffer | handler | | | | property_list
2954  // -------+---------+-------+----+----------+--------------
2955  // Dimensions >= 1
2956  // -------+---------+-------+----+----------+--------------
2957  // buffer | | | | | property_list
2958  // buffer | | | | mode_tag | property_list
2959  // buffer | handler | | | | property_list
2960  // buffer | handler | | | mode_tag | property_list
2961  // buffer | | range | | | property_list
2962  // buffer | | range | | mode_tag | property_list
2963  // buffer | handler | range | | | property_list
2964  // buffer | handler | range | | mode_tag | property_list
2965  // buffer | | range | id | | property_list
2966  // buffer | | range | id | mode_tag | property_list
2967  // buffer | handler | range | id | | property_list
2968  // buffer | handler | range | id | mode_tag | property_list
2969  // -------+---------+-------+----+----------+--------------
2970 
2971  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2972  typename = typename detail::enable_if_t<
2973  std::is_same<T, DataT>::value && Dims == 0>>
2975  buffer<T, 1, AllocatorT> &BufferRef,
2976  const property_list &PropertyList = {},
2977  const detail::code_location CodeLoc = detail::code_location::current())
2978  : AccessorT(BufferRef, PropertyList, CodeLoc) {}
2979 
2980  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2981  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
2983  buffer<T, Dims, AllocatorT> &BufferRef,
2984  const property_list &PropertyList = {},
2985  const detail::code_location CodeLoc = detail::code_location::current())
2986  : AccessorT(BufferRef, PropertyList, CodeLoc) {}
2987 
2988  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2989  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
2992  const property_list &PropertyList = {},
2993  const detail::code_location CodeLoc = detail::code_location::current())
2994  : host_accessor(BufferRef, PropertyList, CodeLoc) {}
2995 
2996  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2997  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
2999  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3000  const property_list &PropertyList = {},
3001  const detail::code_location CodeLoc = detail::code_location::current())
3002  : AccessorT(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
3003 
3004  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3005  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3007  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3008  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
3009  const detail::code_location CodeLoc = detail::code_location::current())
3010  : host_accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
3011 
3012  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3013  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3015  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3016  const property_list &PropertyList = {},
3017  const detail::code_location CodeLoc = detail::code_location::current())
3018  : AccessorT(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
3019 
3020  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3021  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3023  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3024  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
3025  const detail::code_location CodeLoc = detail::code_location::current())
3026  : host_accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
3027 
3028  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3029  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3031  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3032  range<Dimensions> AccessRange, const property_list &PropertyList = {},
3033  const detail::code_location CodeLoc = detail::code_location::current())
3034  : AccessorT(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
3035  CodeLoc) {}
3036 
3037  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3038  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3040  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3042  const property_list &PropertyList = {},
3043  const detail::code_location CodeLoc = detail::code_location::current())
3044  : host_accessor(BufferRef, CommandGroupHandler, AccessRange, {},
3045  PropertyList, CodeLoc) {}
3046 
3047  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3048  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3050  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3051  id<Dimensions> AccessOffset, const property_list &PropertyList = {},
3052  const detail::code_location CodeLoc = detail::code_location::current())
3053  : AccessorT(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
3054  }
3055 
3056  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3057  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3059  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
3060  id<Dimensions> AccessOffset, mode_tag_t<AccessMode>,
3061  const property_list &PropertyList = {},
3062  const detail::code_location CodeLoc = detail::code_location::current())
3063  : host_accessor(BufferRef, AccessRange, AccessOffset, PropertyList,
3064  CodeLoc) {}
3065 
3066  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3067  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3069  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3070  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
3071  const property_list &PropertyList = {},
3072  const detail::code_location CodeLoc = detail::code_location::current())
3073  : AccessorT(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
3074  PropertyList, CodeLoc) {}
3075 
3076  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
3077  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3079  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
3080  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
3081  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
3082  const detail::code_location CodeLoc = detail::code_location::current())
3083  : host_accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
3084  PropertyList, CodeLoc) {}
3085 };
3086 
3087 template <typename DataT, int Dimensions, typename AllocatorT>
3090 
3091 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
3093  -> host_accessor<DataT, Dimensions,
3094  detail::deduceAccessMode<Type1, Type1>()>;
3095 
3096 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3097  typename Type2>
3099  -> host_accessor<DataT, Dimensions,
3100  detail::deduceAccessMode<Type1, Type2>()>;
3101 
3102 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3103  typename Type2, typename Type3>
3105  -> host_accessor<DataT, Dimensions,
3106  detail::deduceAccessMode<Type2, Type3>()>;
3107 
3108 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3109  typename Type2, typename Type3, typename Type4>
3111  -> host_accessor<DataT, Dimensions,
3112  detail::deduceAccessMode<Type3, Type4>()>;
3113 
3114 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
3115  typename Type2, typename Type3, typename Type4, typename Type5>
3117  Type5) -> host_accessor<DataT, Dimensions,
3118  detail::deduceAccessMode<Type4, Type5>()>;
3119 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
3120 } // namespace sycl
3121 
3122 namespace std {
3123 template <typename DataT, int Dimensions, sycl::access::mode AccessMode,
3124  sycl::access::target AccessTarget,
3126 struct hash<sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
3127  IsPlaceholder>> {
3128  using AccType = sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
3129  IsPlaceholder>;
3130 
3131  size_t operator()(const AccType &A) const {
3132 #ifdef __SYCL_DEVICE_ONLY__
3133  // Hash is not supported on DEVICE. Just return 0 here.
3134  (void)A;
3135  return 0;
3136 #else
3137  // getSyclObjImpl() here returns a pointer to either AccessorImplHost
3138  // or LocalAccessorImplHost depending on the AccessTarget.
3139  auto AccImplPtr = sycl::detail::getSyclObjImpl(A);
3140  return hash<decltype(AccImplPtr)>()(AccImplPtr);
3141 #endif
3142  }
3143 };
3144 
3145 } // namespace std
The file contains implementation of accessor iterator class.
accessor(sycl::image< Dimensions, AllocatorT > &Image, const property_list &propList)
Definition: accessor.hpp:2843
Defines a shared array that can be used by kernels in queues.
Definition: buffer.hpp:146
AccessorBaseHost(const AccessorImplPtr &Impl)
Definition: accessor.hpp:469
decltype(Obj::impl) friend getSyclObjImpl(const Obj &SyclObject)
AccessorImplDevice(id< Dims > Offset, range< Dims > AccessRange, range< Dims > MemoryRange)
Definition: accessor.hpp:445
bool operator==(const AccessorImplDevice &Rhs) const
Definition: accessor.hpp:453
LocalAccessorBaseDevice(sycl::range< Dims > Size)
Definition: accessor.hpp:423
bool operator==(const LocalAccessorBaseDevice &Rhs) const
Definition: accessor.hpp:432
decltype(Obj::impl) friend getSyclObjImpl(const Obj &SyclObject)
LocalAccessorBaseHost(const LocalAccessorImplPtr &Impl)
Definition: accessor.hpp:524
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:953
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:917
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:931
DataT read(const CoordT &Coords, const sampler &Smpl) const
Definition: accessor.hpp:924
__image_array_slice__(accessor< DataT, Dimensions, AccessMode, access::target::image_array, IsPlaceholder, ext::oneapi::accessor_property_list<>> BaseAcc, size_t Idx)
Definition: accessor.hpp:907
detail::enable_if_t< CurDims==1 &&IsAccessAtomic, atomic< DataT, AS > > operator[](size_t Index) const
Definition: accessor.hpp:367
AccessorSubscript(AccType Accessor, size_t Index)
Definition: accessor.hpp:345
AccessorSubscript(AccType Accessor, id< Dims > IDs)
Definition: accessor.hpp:340
detail::const_if_const_AS< AS, DataT > * PtrType
Definition: accessor.hpp:325
detail::const_if_const_AS< AS, DataT > & RefType
Definition: accessor.hpp:323
size_t size() const noexcept
Definition: accessor.hpp:801
image_accessor(image< Dims, AllocatorT > &ImageRef, int ImageElementSize)
Definition: accessor.hpp:701
bool operator!=(const image_accessor &Rhs) const
Definition: accessor.hpp:766
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:799
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:864
range< Dims > get_range() const
Definition: accessor.hpp:804
image_accessor(const AccessorImplPtr &Impl)
Definition: accessor.hpp:673
bool operator==(const image_accessor &Rhs) const
Definition: accessor.hpp:758
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:820
DataT read(const CoordT &Coords, const sampler &Smpl) const
Definition: accessor.hpp:839
image_accessor(image< Dims, AllocatorT > &ImageRef, handler &CommandGroupHandlerRef, int ImageElementSize)
Definition: accessor.hpp:729
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:49
Command group handler class.
Definition: handler.hpp:312
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:3068
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:2998
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:2938
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:3006
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:3030
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:3014
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2982
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:3022
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:2990
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:3058
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:3049
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:3039
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:3078
host_accessor(buffer< T, 1, AllocatorT > &BufferRef, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2974
Defines a shared image data.
Definition: image.hpp:181
range< Dimensions > get_range() const
Definition: image.hpp:364
size_t size() const noexcept
Definition: image.hpp:381
detail::const_if_const_AS< AS, DataT > & RefType
Definition: accessor.hpp:2404
const range< 3 > & getSize() const
Definition: accessor.hpp:2450
bool operator!=(const local_accessor_base &Rhs) const
Definition: accessor.hpp:2619
AccessorCommonT::template AccessorSubscript< Dims - 1, local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder > > operator[](size_t Index) const
Definition: accessor.hpp:2608
local_accessor_base(range< Dimensions > AllocationSize, handler &, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2516
range< Dims > get_range() const
Definition: accessor.hpp:2556
local_accessor_base(handler &, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2485
typename detail::DecoratedType< DataT, AS >::type * ConcreteASPtrType
Definition: accessor.hpp:2402
typename AccessorCommonT::template AccessorSubscript< Dims, local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder > > AccessorSubscript
Definition: accessor.hpp:2400
size_t getLinearIndex(id< Dims > Id) const
Definition: accessor.hpp:2469
local_accessor_base(range< Dimensions > AllocationSize, handler &, const property_list &propList, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2532
local_accessor_base(handler &, const property_list &propList, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2499
bool operator==(const local_accessor_base &Rhs) const
Definition: accessor.hpp:2616
RefType operator[](size_t Index) const
Definition: accessor.hpp:2575
local_ptr< DataT > get_pointer() const
Definition: accessor.hpp:2612
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:2551
detail::enable_if_t< Dims==1 &&AccessMode==access::mode::atomic, atomic< DataT, AS > > operator[](size_t Index) const
Definition: accessor.hpp:2599
detail::const_if_const_AS< AS, DataT > * PtrType
Definition: accessor.hpp:2405
size_t size() const noexcept
Definition: accessor.hpp:2553
PtrType getQualifiedPtr() const
Definition: accessor.hpp:2444
local_accessor_base(const detail::LocalAccessorImplPtr &Impl)
Definition: accessor.hpp:2437
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: multi_ptr.hpp:78
Objects of the property_list class are containers for the SYCL properties.
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: range.hpp:24
#define __SYCL_SPECIAL_CLASS
Definition: defines.hpp:30
#define __SYCL_TYPE(x)
Definition: defines.hpp:40
#define __SYCL_EBO
Definition: common.hpp:281
#define __SYCL_INLINE_VER_NAMESPACE(X)
#define __SYCL2020_DEPRECATED(message)
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.
constexpr tuple_element< I, tuple< Types... > >::type & get(sycl::detail::tuple< Types... > &Arg) noexcept
Definition: tuple.hpp:199
void addHostAccessorAndWait(AccessorImplHost *Req)
device getDeviceFromHandler(handler &CommandGroupHandlerRef)
Definition: accessor.cpp:15
size_t getLinearIndex(const T< Dims > &Index, const U< Dims > &Range)
Definition: common.hpp:433
void imageWriteHostImpl(const CoordT &Coords, const WriteDataT &Color, id< 3 > ImgPitch, uint8_t ElementSize, image_channel_type ImgChannelType, image_channel_order ImgChannelOrder, void *BasePtr)
void dim_loop(F &&f)
Definition: accessor.hpp:250
typename std::conditional< B, T, F >::type conditional_t
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: common.hpp:300
typename std::is_same< ext::oneapi::accessor_property_list<>, T > IsRunTimePropertyListT
Definition: accessor.hpp:264
void constructorNotification(void *BufferObj, void *AccessorObj, access::target Target, access::mode Mode, const code_location &CodeLoc)
std::integral_constant< bool, V > bool_constant
T createSyclObjFromImpl(decltype(T::impl) ImplObj)
Definition: common.hpp:318
std::shared_ptr< LocalAccessorImplHost > LocalAccessorImplPtr
Definition: accessor.hpp:520
void dim_loop_impl(std::integer_sequence< size_t, Inds... >, F &&f)
Definition: accessor.hpp:246
constexpr access::target deduceAccessTarget(access::target defaultTarget)
Definition: accessor.hpp:407
typename std::is_base_of< PropertyListBase, T > IsPropertyListT
Definition: accessor.hpp:260
constexpr access::mode deduceAccessMode()
Definition: accessor.hpp:375
void associateWithHandler(handler &, AccessorBaseHost *, access::target)
typename std::enable_if< B, T >::type enable_if_t
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
Definition: stream.hpp:111
std::shared_ptr< AccessorImplHost > AccessorImplPtr
Definition: accessor.hpp:465
static constexpr bool has_property()
static constexpr auto get_property()
constexpr std::enable_if_t< detail::IsCompileTimeProperty< PropertyT >::value, bool > operator!=(const property_value< PropertyT, A... > &, const property_value< PropertyT, B... > &)
T & operator[](std::ptrdiff_t idx) const noexcept
constexpr buffer_location_key::value_t< N > buffer_location
Definition: properties.hpp:85
constexpr property::no_offset::instance no_offset
host_accessor(buffer< DataT, Dimensions, AllocatorT >, Type1, Type2, Type3, Type4, Type5) -> host_accessor< DataT, Dimensions, detail::deduceAccessMode< Type4, Type5 >()>
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS Dimensions
Definition: accessor.hpp:2764
constexpr mode_tag_t< access_mode::read_write > read_write
Definition: access.hpp:73
image_channel_order
Definition: image.hpp:27
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS IsPlaceholder
Definition: accessor.hpp:2765
accessor(buffer< DataT, Dimensions, AllocatorT >, handler &, Type1, Type2, Type3, Type4, const ext::oneapi::accessor_property_list< PropsT... > &) -> accessor< DataT, Dimensions, detail::deduceAccessMode< Type3, Type4 >(), detail::deduceAccessTarget< Type3, Type4 >(target::device), access::placeholder::false_t, ext::oneapi::accessor_property_list< PropsT... >>
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS AccessMode
Definition: accessor.hpp:2764
image_channel_type
Definition: image.hpp:45
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
@ Device
simd< _Tp, _Abi > max(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
sycl::accessor< DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder > AccType
Definition: accessor.hpp:3129
sycl::range< 3 > MMemoryRange
Definition: accessor.hpp:239
sycl::range< 3 > MAccessRange
Definition: accessor.hpp:238
AccHostDataT(const sycl::id< 3 > &Offset, const sycl::range< 3 > &Range, const sycl::range< 3 > &MemoryRange, void *Data=nullptr)
Definition: accessor.hpp:232
bool operator==(const Slab &Lhs, const Slab &Rhs)