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