DPC++ Runtime
Runtime libraries for oneAPI Data Parallel C++
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 <CL/sycl/atomic.hpp>
13 #include <CL/sycl/buffer.hpp>
21 #include <CL/sycl/exception.hpp>
22 #include <CL/sycl/id.hpp>
23 #include <CL/sycl/image.hpp>
24 #include <CL/sycl/pointers.hpp>
28 #include <CL/sycl/sampler.hpp>
30 
31 #include <type_traits>
32 
92 // +------------------+ +-----------------+ +-----------------------+
93 // | | | | | |
94 // | AccessorBaseHost | | accessor_common | | LocalAccessorBaseHost |
95 // | | | | | |
96 // +------------------+ +-----+-----------+ +--------+--------------+
97 // | | | | |
98 // | +-----------+ +----+ +---------+ +------+
99 // | | | | |
100 // v v v v v
101 // +----------------+ +-----------------+ +-------------+
102 // | | | accessor(1) | | accessor(3) |
103 // | image_accessor | +-----------------| +-------------+
104 // | | | for targets: | | for target: |
105 // +---+---+---+----+ | | | |
106 // | | | | host_buffer | | local |
107 // | | | | global_buffer | +-------------+
108 // | | | | constant_buffer |
109 // | | | +-----------------+
110 // | | |
111 // | | +------------------------------------+
112 // | | |
113 // | +----------------------+ |
114 // v v v
115 // +-----------------+ +--------------+ +-------------+
116 // | acessor(2) | | accessor(4) | | accessor(5) |
117 // +-----------------+ +--------------+ +-------------+
118 // | for targets: | | for targets: | | for target: |
119 // | | | | | |
120 // | host_image | | image | | image_array |
121 // +-----------------+ +--------------+ +-------------+
122 //
147 //
148 // +-----------------+
149 // | |
150 // | accessor_common |
151 // | |
152 // +-----+-------+---+
153 // | |
154 // +----+ +-----+
155 // | |
156 // v v
157 // +----------------+ +-----------------+ +-------------+
158 // | | | accessor(1) | | accessor(3) |
159 // | image_accessor | +-----------------| +-------------+
160 // | | | for targets: | | for target: |
161 // +---+---+---+----+ | | | |
162 // | | | | host_buffer | | local |
163 // | | | | global_buffer | +-------------+
164 // | | | | constant_buffer |
165 // | | | +-----------------+
166 // | | | |
167 // | | | v
168 // | | | +-----------------+
169 // | | | | |
170 // | | | | host_accessor |
171 // | | | | |
172 // | | | +-----------------+
173 // | | |
174 // | | +------------------------------------+
175 // | | |
176 // | +----------------------+ |
177 // v v v
178 // +-----------------+ +--------------+ +-------------+
179 // | acessor(2) | | accessor(4) | | accessor(5) |
180 // +-----------------+ +--------------+ +-------------+
181 // | for targets: | | for targets: | | for target: |
182 // | | | | | |
183 // | host_image | | image | | image_array |
184 // +-----------------+ +--------------+ +-------------+
185 //
203 
205 namespace sycl {
206 class stream;
207 namespace ext {
208 namespace intel {
209 namespace experimental {
210 namespace esimd {
211 namespace detail {
212 // Forward declare a "back-door" access class to support ESIMD.
214 } // namespace detail
215 } // namespace esimd
216 } // namespace experimental
217 } // namespace intel
218 } // namespace ext
219 
220 template <typename DataT, int Dimensions = 1,
221  access::mode AccessMode = access::mode::read_write,
222  access::target AccessTarget = access::target::device,
223  access::placeholder IsPlaceholder = access::placeholder::false_t,
224  typename PropertyListT = ext::oneapi::accessor_property_list<>>
225 class accessor;
226 
227 namespace detail {
228 template <typename T>
229 using IsPropertyListT = typename std::is_base_of<PropertyListBase, T>;
230 
231 template <typename T>
233  typename std::is_same<ext::oneapi::accessor_property_list<>, T>;
234 
235 template <typename T> struct IsCxPropertyList {
236  constexpr static bool value = false;
237 };
238 
239 template <typename... Props>
240 struct IsCxPropertyList<ext::oneapi::accessor_property_list<Props...>> {
241  constexpr static bool value = true;
242 };
243 
244 template <> struct IsCxPropertyList<ext::oneapi::accessor_property_list<>> {
245  constexpr static bool value = false;
246 };
247 
248 // The function extends or truncates number of dimensions of objects of id
249 // or ranges classes. When extending the new values are filled with
250 // DefaultValue, truncation just removes extra values.
251 template <int NewDim, int DefaultValue, template <int> class T, int OldDim>
254  const int CopyDims = NewDim > OldDim ? OldDim : NewDim;
255  for (int I = 0; I < CopyDims; ++I)
256  NewObj[I] = OldObj[I];
257  for (int I = CopyDims; I < NewDim; ++I)
258  NewObj[I] = DefaultValue;
259  return NewObj;
260 }
261 
262 __SYCL_EXPORT device getDeviceFromHandler(handler &CommandGroupHandlerRef);
263 
264 template <typename DataT, int Dimensions, access::mode AccessMode,
265  access::target AccessTarget, access::placeholder IsPlaceholder,
266  typename PropertyListT = ext::oneapi::accessor_property_list<>>
268 protected:
269  constexpr static bool IsPlaceH = IsPlaceholder == access::placeholder::true_t;
271 
272  constexpr static bool IsHostBuf = AccessTarget == access::target::host_buffer;
273 
274  // TODO: SYCL 2020 deprecates four of the target enum values
275  // and replaces them with 2 (device and host_task). May want
276  // to change these constexpr.
277  constexpr static bool IsGlobalBuf =
278  AccessTarget == access::target::global_buffer;
279 
280  constexpr static bool IsConstantBuf =
281  AccessTarget == access::target::constant_buffer;
282 
283  constexpr static bool IsAccessAnyWrite =
284  AccessMode == access::mode::write ||
285  AccessMode == access::mode::read_write ||
286  AccessMode == access::mode::discard_write ||
287  AccessMode == access::mode::discard_read_write;
288 
289  constexpr static bool IsAccessReadOnly = AccessMode == access::mode::read;
290 
291  constexpr static bool IsAccessReadWrite =
292  AccessMode == access::mode::read_write;
293 
294  constexpr static bool IsAccessAtomic = AccessMode == access::mode::atomic;
295 
297  using ConstRefType = const DataT &;
299 
300  using AccType = accessor<DataT, Dimensions, AccessMode, AccessTarget,
301  IsPlaceholder, PropertyListT>;
302 
303  // The class which allows to access value of N dimensional accessor using N
304  // subscript operators, e.g. accessor[2][2][3]
305  template <int SubDims> class AccessorSubscript {
306  static constexpr int Dims = Dimensions;
307 
308  mutable id<Dims> MIDs;
309  AccType MAccessor;
310 
311  public:
313  : MAccessor(Accessor), MIDs(IDs) {}
314 
315  // Only accessor class is supposed to use this c'tor for the first
316  // operator[].
317  AccessorSubscript(AccType Accessor, size_t Index) : MAccessor(Accessor) {
318  MIDs[0] = Index;
319  }
320 
321  template <int CurDims = SubDims>
322  typename detail::enable_if_t<(CurDims > 1), AccessorSubscript<CurDims - 1>>
323  operator[](size_t Index) {
324  MIDs[Dims - CurDims] = Index;
325  return AccessorSubscript<CurDims - 1>(MAccessor, MIDs);
326  }
327 
328  template <int CurDims = SubDims,
330  RefType operator[](size_t Index) const {
331  MIDs[Dims - CurDims] = Index;
332  return MAccessor[MIDs];
333  }
334 
335  template <int CurDims = SubDims>
336  typename detail::enable_if_t<CurDims == 1 && IsAccessAtomic,
338  operator[](size_t Index) const {
339  MIDs[Dims - CurDims] = Index;
340  return MAccessor[MIDs];
341  }
342 
343  template <int CurDims = SubDims,
345  ConstRefType operator[](size_t Index) const {
346  MIDs[Dims - SubDims] = Index;
347  return MAccessor[MIDs];
348  }
349  };
350 };
351 
352 template <int Dim, typename T> struct IsValidCoordDataT;
353 template <typename T> struct IsValidCoordDataT<1, T> {
354  constexpr static bool value =
356 };
357 template <typename T> struct IsValidCoordDataT<2, T> {
358  constexpr static bool value =
361 };
362 template <typename T> struct IsValidCoordDataT<3, T> {
363  constexpr static bool value =
366 };
367 
368 template <typename DataT, int Dimensions, access::mode AccessMode,
369  access::placeholder IsPlaceholder>
371 
372 // Image accessor
373 template <typename DataT, int Dimensions, access::mode AccessMode,
374  access::target AccessTarget, access::placeholder IsPlaceholder>
376 #ifndef __SYCL_DEVICE_ONLY__
377  : public detail::AccessorBaseHost {
378  size_t MImageCount;
379  image_channel_order MImgChannelOrder;
380  image_channel_type MImgChannelType;
381 #else
382 {
383 
384  using OCLImageTy = typename detail::opencl_image_type<Dimensions, AccessMode,
385  AccessTarget>::type;
386  OCLImageTy MImageObj;
387  char MPadding[sizeof(detail::AccessorBaseHost) +
388  sizeof(size_t /*MImageCount*/) + sizeof(image_channel_order) +
389  sizeof(image_channel_type) - sizeof(OCLImageTy)];
390 
391 protected:
392  void imageAccessorInit(OCLImageTy Image) { MImageObj = Image; }
393 
394 private:
395 #endif
396  template <typename T1, int T2, access::mode T3, access::placeholder T4>
397  friend class __image_array_slice__;
398 
399  constexpr static bool IsHostImageAcc =
400  (AccessTarget == access::target::host_image);
401 
402  constexpr static bool IsImageAcc = (AccessTarget == access::target::image);
403 
404  constexpr static bool IsImageArrayAcc =
405  (AccessTarget == access::target::image_array);
406 
407  constexpr static bool IsImageAccessWriteOnly =
408  (AccessMode == access::mode::write ||
409  AccessMode == access::mode::discard_write);
410 
411  constexpr static bool IsImageAccessAnyWrite =
412  (IsImageAccessWriteOnly || AccessMode == access::mode::read_write);
413 
414  constexpr static bool IsImageAccessReadOnly =
415  (AccessMode == access::mode::read);
416 
417  constexpr static bool IsImageAccessAnyRead =
418  (IsImageAccessReadOnly || AccessMode == access::mode::read_write);
419 
420  static_assert(std::is_same<DataT, cl_int4>::value ||
421  std::is_same<DataT, cl_uint4>::value ||
422  std::is_same<DataT, cl_float4>::value ||
423  std::is_same<DataT, cl_half4>::value,
424  "The data type of an image accessor must be only cl_int4, "
425  "cl_uint4, cl_float4 or cl_half4 from SYCL namespace");
426 
427  static_assert(IsImageAcc || IsHostImageAcc || IsImageArrayAcc,
428  "Expected image type");
429 
430  static_assert(IsPlaceholder == access::placeholder::false_t,
431  "Expected false as Placeholder value for image accessor.");
432 
433  static_assert(
434  ((IsImageAcc || IsImageArrayAcc) &&
435  (IsImageAccessWriteOnly || IsImageAccessReadOnly)) ||
436  (IsHostImageAcc && (IsImageAccessAnyWrite || IsImageAccessAnyRead)),
437  "Access modes can be only read/write/discard_write for image/image_array "
438  "target accessor, or they can be only "
439  "read/write/discard_write/read_write for host_image target accessor.");
440 
441  static_assert(Dimensions > 0 && Dimensions <= 3,
442  "Dimensions can be 1/2/3 for image accessor.");
443 
444  template <info::device param>
445  void checkDeviceFeatureSupported(const device &Device) {
446  if (!Device.get_info<param>())
447  throw feature_not_supported("Images are not supported by this device.",
449  }
450 
451 #ifdef __SYCL_DEVICE_ONLY__
452 
453  sycl::vec<int, Dimensions> getRangeInternal() const {
454  return __invoke_ImageQuerySize<sycl::vec<int, Dimensions>, OCLImageTy>(
455  MImageObj);
456  }
457 
458  size_t getElementSize() const {
459  int ChannelType = __invoke_ImageQueryFormat<int, OCLImageTy>(MImageObj);
460  int ChannelOrder = __invoke_ImageQueryOrder<int, OCLImageTy>(MImageObj);
461  int ElementSize = getSPIRVElementSize(ChannelType, ChannelOrder);
462  return ElementSize;
463  }
464 
465 #else
466 
467  sycl::vec<int, Dimensions> getRangeInternal() const {
468  // TODO: Implement for host.
469  throw runtime_error("image::getRangeInternal() is not implemented for host",
471  return sycl::vec<int, Dimensions>{1};
472  }
473 
474 #endif
475 
476 private:
477  friend class sycl::ext::intel::experimental::esimd::detail::
478  AccessorPrivateProxy;
479 
480 #ifdef __SYCL_DEVICE_ONLY__
481  const OCLImageTy getNativeImageObj() const { return MImageObj; }
482 #endif // __SYCL_DEVICE_ONLY__
483 
484 public:
485  using value_type = DataT;
486  using reference = DataT &;
487  using const_reference = const DataT &;
488 
489  // image_accessor Constructors.
490 
491 #ifdef __SYCL_DEVICE_ONLY__
492  // Default constructor for objects later initialized with __init member.
493  image_accessor() = default;
494 #endif
495 
496  // Available only when: accessTarget == access::target::host_image
497  // template <typename AllocatorT>
498  // accessor(image<dimensions, AllocatorT> &imageRef);
499  template <
500  typename AllocatorT, int Dims = Dimensions,
501  typename = detail::enable_if_t<(Dims > 0 && Dims <= 3) && IsHostImageAcc>>
502  image_accessor(image<Dims, AllocatorT> &ImageRef, int ImageElementSize)
503 #ifdef __SYCL_DEVICE_ONLY__
504  {
505  (void)ImageRef;
506  (void)ImageElementSize;
507  // No implementation needed for device. The constructor is only called by
508  // host.
509  }
510 #else
511  : AccessorBaseHost({detail::getSyclObjImpl(ImageRef)->getRowPitch(),
512  detail::getSyclObjImpl(ImageRef)->getSlicePitch(), 0},
513  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
514  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
515  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
516  Dimensions, ImageElementSize),
517  MImageCount(ImageRef.size()),
518  MImgChannelOrder(detail::getSyclObjImpl(ImageRef)->getChannelOrder()),
519  MImgChannelType(detail::getSyclObjImpl(ImageRef)->getChannelType()) {
520  addHostAccessorAndWait(AccessorBaseHost::impl.get());
521  }
522 #endif
523 
524  // Available only when: accessTarget == access::target::image
525  // template <typename AllocatorT>
526  // accessor(image<dimensions, AllocatorT> &imageRef,
527  // handler &commandGroupHandlerRef);
528  template <
529  typename AllocatorT, int Dims = Dimensions,
530  typename = detail::enable_if_t<(Dims > 0 && Dims <= 3) && IsImageAcc>>
532  handler &CommandGroupHandlerRef, int ImageElementSize)
533 #ifdef __SYCL_DEVICE_ONLY__
534  {
535  (void)ImageRef;
536  (void)CommandGroupHandlerRef;
537  (void)ImageElementSize;
538  // No implementation needed for device. The constructor is only called by
539  // host.
540  }
541 #else
542  : AccessorBaseHost({detail::getSyclObjImpl(ImageRef)->getRowPitch(),
543  detail::getSyclObjImpl(ImageRef)->getSlicePitch(), 0},
544  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
545  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
546  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
547  Dimensions, ImageElementSize),
548  MImageCount(ImageRef.size()),
549  MImgChannelOrder(detail::getSyclObjImpl(ImageRef)->getChannelOrder()),
550  MImgChannelType(detail::getSyclObjImpl(ImageRef)->getChannelType()) {
551  checkDeviceFeatureSupported<info::device::image_support>(
552  getDeviceFromHandler(CommandGroupHandlerRef));
553  }
554 #endif
555 
556  /* -- common interface members -- */
557 
558  // operator == and != need to be defined only for host application as per the
559  // SYCL spec 1.2.1
560 #ifndef __SYCL_DEVICE_ONLY__
561  bool operator==(const image_accessor &Rhs) const { return Rhs.impl == impl; }
562 #else
563  // The operator with __SYCL_DEVICE_ONLY__ need to be declared for compilation
564  // of host application with device compiler.
565  // Usage of this operator inside the kernel code will give a runtime failure.
566  bool operator==(const image_accessor &Rhs) const;
567 #endif
568 
569  bool operator!=(const image_accessor &Rhs) const { return !(Rhs == *this); }
570 
571  // get_count() method : Returns the number of elements of the SYCL image this
572  // SYCL accessor is accessing.
573  //
574  // get_range() method : Returns a range object which represents the number of
575  // elements of dataT per dimension that this accessor may access.
576  // The range object returned must equal to the range of the image this
577  // accessor is associated with.
578 
579 #ifdef __SYCL_DEVICE_ONLY__
580 
581  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
582  size_t get_count() const { return size(); }
583  size_t size() const noexcept { return get_range<Dimensions>().size(); }
584 
585  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 1>>
586  range<1> get_range() const {
587  cl_int Range = getRangeInternal();
588  return range<1>(Range);
589  }
590  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 2>>
591  range<2> get_range() const {
592  cl_int2 Range = getRangeInternal();
593  return range<2>(Range[0], Range[1]);
594  }
595  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 3>>
596  range<3> get_range() const {
597  cl_int3 Range = getRangeInternal();
598  return range<3>(Range[0], Range[1], Range[2]);
599  }
600 
601 #else
602  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
603  size_t get_count() const { return size(); };
604  size_t size() const noexcept { return MImageCount; };
605 
606  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
608  return detail::convertToArrayOfN<Dims, 1>(getAccessRange());
609  }
610 
611 #endif
612 
613  // Available only when:
614  // (accessTarget == access::target::image && accessMode == access::mode::read)
615  // || (accessTarget == access::target::host_image && ( accessMode ==
616  // access::mode::read || accessMode == access::mode::read_write))
617  template <typename CoordT, int Dims = Dimensions,
618  typename = detail::enable_if_t<
619  (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
621  ((IsImageAcc && IsImageAccessReadOnly) ||
622  (IsHostImageAcc && IsImageAccessAnyRead))>>
623  DataT read(const CoordT &Coords) const {
624 #ifdef __SYCL_DEVICE_ONLY__
625  return __invoke__ImageRead<DataT, OCLImageTy, CoordT>(MImageObj, Coords);
626 #else
627  sampler Smpl(coordinate_normalization_mode::unnormalized,
628  addressing_mode::none, filtering_mode::nearest);
629  return read<CoordT, Dims>(Coords, Smpl);
630 #endif
631  }
632 
633  // Available only when:
634  // (accessTarget == access::target::image && accessMode == access::mode::read)
635  // || (accessTarget == access::target::host_image && ( accessMode ==
636  // access::mode::read || accessMode == access::mode::read_write))
637  template <typename CoordT, int Dims = Dimensions,
638  typename = detail::enable_if_t<
639  (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
640  ((IsImageAcc && IsImageAccessReadOnly) ||
641  (IsHostImageAcc && IsImageAccessAnyRead))>>
642  DataT read(const CoordT &Coords, const sampler &Smpl) const {
643 #ifdef __SYCL_DEVICE_ONLY__
644  return __invoke__ImageReadSampler<DataT, OCLImageTy, CoordT>(
645  MImageObj, Coords, Smpl.impl.m_Sampler);
646 #else
647  return imageReadSamplerHostImpl<CoordT, DataT>(
648  Coords, Smpl, getAccessRange() /*Image Range*/,
649  getOffset() /*Image Pitch*/, MImgChannelType, MImgChannelOrder,
650  AccessorBaseHost::getPtr() /*ptr to image*/,
651  AccessorBaseHost::getElemSize());
652 #endif
653  }
654 
655  // Available only when:
656  // (accessTarget == access::target::image && (accessMode ==
657  // access::mode::write || accessMode == access::mode::discard_write)) ||
658  // (accessTarget == access::target::host_image && (accessMode ==
659  // access::mode::write || accessMode == access::mode::discard_write ||
660  // accessMode == access::mode::read_write))
661  template <typename CoordT, int Dims = Dimensions,
662  typename = detail::enable_if_t<
663  (Dims > 0) && (detail::is_genint<CoordT>::value) &&
665  ((IsImageAcc && IsImageAccessWriteOnly) ||
666  (IsHostImageAcc && IsImageAccessAnyWrite))>>
667  void write(const CoordT &Coords, const DataT &Color) const {
668 #ifdef __SYCL_DEVICE_ONLY__
669  __invoke__ImageWrite<OCLImageTy, CoordT, DataT>(MImageObj, Coords, Color);
670 #else
671  imageWriteHostImpl(Coords, Color, getOffset() /*ImagePitch*/,
672  AccessorBaseHost::getElemSize(), MImgChannelType,
673  MImgChannelOrder,
674  AccessorBaseHost::getPtr() /*Ptr to Image*/);
675 #endif
676  }
677 };
678 
679 template <typename DataT, int Dimensions, access::mode AccessMode,
680  access::placeholder IsPlaceholder>
681 class __image_array_slice__ {
682 
683  static_assert(Dimensions < 3,
684  "Image slice cannot have more then 2 dimensions");
685 
686  constexpr static int AdjustedDims = (Dimensions == 2) ? 4 : Dimensions + 1;
687 
688  template <typename CoordT,
689  typename CoordElemType =
692  getAdjustedCoords(const CoordT &Coords) const {
693  CoordElemType LastCoord = 0;
694 
695  if (std::is_same<float, CoordElemType>::value) {
696  sycl::vec<int, Dimensions + 1> Size = MBaseAcc.getRangeInternal();
697  LastCoord =
698  MIdx / static_cast<float>(Size.template swizzle<Dimensions>());
699  } else {
700  LastCoord = MIdx;
701  }
702 
703  sycl::vec<CoordElemType, Dimensions> LeftoverCoords{LastCoord};
704  sycl::vec<CoordElemType, AdjustedDims> AdjustedCoords{Coords,
705  LeftoverCoords};
706  return AdjustedCoords;
707  }
708 
709 public:
711  accessor<DataT, Dimensions, AccessMode, access::target::image_array,
713  BaseAcc,
714  size_t Idx)
715  : MBaseAcc(BaseAcc), MIdx(Idx) {}
716 
717  template <typename CoordT, int Dims = Dimensions,
718  typename = detail::enable_if_t<
720  DataT read(const CoordT &Coords) const {
721  return MBaseAcc.read(getAdjustedCoords(Coords));
722  }
723 
724  template <typename CoordT, int Dims = Dimensions,
725  typename = detail::enable_if_t<
727  DataT read(const CoordT &Coords, const sampler &Smpl) const {
728  return MBaseAcc.read(getAdjustedCoords(Coords), Smpl);
729  }
730 
731  template <typename CoordT, int Dims = Dimensions,
732  typename = detail::enable_if_t<
734  void write(const CoordT &Coords, const DataT &Color) const {
735  return MBaseAcc.write(getAdjustedCoords(Coords), Color);
736  }
737 
738 #ifdef __SYCL_DEVICE_ONLY__
739  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
740  size_t get_count() const { return size(); }
741  size_t size() const noexcept { return get_range<Dimensions>().size(); }
742 
743  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 1>>
744  range<1> get_range() const {
745  cl_int2 Count = MBaseAcc.getRangeInternal();
746  return range<1>(Count.x());
747  }
748  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 2>>
749  range<2> get_range() const {
750  cl_int3 Count = MBaseAcc.getRangeInternal();
751  return range<2>(Count.x(), Count.y());
752  }
753 
754 #else
755 
756  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
757  size_t get_count() const { return size(); }
758  size_t size() const noexcept {
759  return MBaseAcc.MImageCount / MBaseAcc.getAccessRange()[Dimensions];
760  }
761 
762  template <int Dims = Dimensions,
763  typename = detail::enable_if_t<(Dims == 1 || Dims == 2)>>
765  return detail::convertToArrayOfN<Dims, 1>(MBaseAcc.getAccessRange());
766  }
767 
768 #endif
769 
770 private:
771  size_t MIdx;
772  accessor<DataT, Dimensions, AccessMode, access::target::image_array,
774  MBaseAcc;
775 };
776 
777 } // namespace detail
778 
784 template <typename DataT, int Dimensions, access::mode AccessMode,
785  access::target AccessTarget, access::placeholder IsPlaceholder,
786  typename PropertyListT>
787 class __SYCL_SPECIAL_CLASS accessor :
788 #ifndef __SYCL_DEVICE_ONLY__
789  public detail::AccessorBaseHost,
790 #endif
791  public detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
792  IsPlaceholder, PropertyListT> {
793 protected:
794  static_assert((AccessTarget == access::target::global_buffer ||
795  AccessTarget == access::target::constant_buffer ||
796  AccessTarget == access::target::host_buffer),
797  "Expected buffer type");
798 
799  static_assert((AccessTarget == access::target::global_buffer ||
800  AccessTarget == access::target::host_buffer) ||
801  (AccessTarget == access::target::constant_buffer &&
802  AccessMode == access::mode::read),
803  "Access mode can be only read for constant buffers");
804 
805  static_assert(detail::IsPropertyListT<PropertyListT>::value,
806  "PropertyListT must be accessor_property_list");
807 
808  using AccessorCommonT =
809  detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
810  IsPlaceholder, PropertyListT>;
811 
812  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
813 
814  using AccessorCommonT::AS;
815  using AccessorCommonT::IsAccessAnyWrite;
816  using AccessorCommonT::IsAccessReadOnly;
817  using AccessorCommonT::IsConstantBuf;
818  using AccessorCommonT::IsGlobalBuf;
819  using AccessorCommonT::IsHostBuf;
820  using AccessorCommonT::IsPlaceH;
821  template <int Dims>
822  using AccessorSubscript =
823  typename AccessorCommonT::template AccessorSubscript<Dims>;
824 
826 
828  using ConstRefType = const DataT &;
830 
831  template <int Dims = Dimensions> size_t getLinearIndex(id<Dims> Id) const {
832 
833 #ifdef __SYCL_DEVICE_ONLY__
834  // Pointer is already adjusted for 1D case.
835  if (Dimensions == 1)
836  return Id[0];
837 #endif // __SYCL_DEVICE_ONLY__
838 
839  size_t Result = 0;
840  // Unroll the following loop for both host and device code
841  __SYCL_UNROLL(3)
842  for (int I = 0; I < Dims; ++I)
843  Result = Result * getMemoryRange()[I] + getOffset()[I] + Id[I];
844  return Result;
845  }
846 
847  template <typename T, int Dims> static constexpr bool IsSameAsBuffer() {
848  return std::is_same<T, DataT>::value && (Dims > 0) && (Dims == Dimensions);
849  }
850 
851  static access::mode getAdjustedMode(const PropertyListT &PropertyList) {
852  access::mode AdjustedMode = AccessMode;
853 
854  if (PropertyList.template has_property<property::no_init>() ||
855  PropertyList.template has_property<property::noinit>()) {
856  if (AdjustedMode == access::mode::write) {
857  AdjustedMode = access::mode::discard_write;
858  } else if (AdjustedMode == access::mode::read_write) {
859  AdjustedMode = access::mode::discard_read_write;
860  }
861  }
862 
863  return AdjustedMode;
864  }
865 
866 #if __cplusplus > 201402L
867 
868  template <typename TagT> static constexpr bool IsValidTag() {
869  return std::is_same<TagT, mode_tag_t<AccessMode>>::value ||
870  std::is_same<TagT,
872  }
873 
874 #endif
875 
876 #ifdef __SYCL_DEVICE_ONLY__
877 
878  id<AdjustedDim> &getOffset() { return impl.Offset; }
879  range<AdjustedDim> &getAccessRange() { return impl.AccessRange; }
880  range<AdjustedDim> &getMemoryRange() { return impl.MemRange; }
881 
882  const id<AdjustedDim> &getOffset() const { return impl.Offset; }
883  const range<AdjustedDim> &getAccessRange() const { return impl.AccessRange; }
884  const range<AdjustedDim> &getMemoryRange() const { return impl.MemRange; }
885 
886  detail::AccessorImplDevice<AdjustedDim> impl;
887 
888  union {
889  ConcreteASPtrType MData;
890  };
891 
892  // TODO replace usages with getQualifiedPtr
893  const ConcreteASPtrType getNativeImageObj() const { return MData; }
894 
895  void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
896  range<AdjustedDim> MemRange, id<AdjustedDim> Offset) {
897  MData = Ptr;
898 #pragma unroll
899  for (int I = 0; I < AdjustedDim; ++I) {
900  getOffset()[I] = Offset[I];
901  getAccessRange()[I] = AccessRange[I];
902  getMemoryRange()[I] = MemRange[I];
903  }
904  // In case of 1D buffer, adjust pointer during initialization rather
905  // then each time in operator[] or get_pointer functions.
906  if (1 == AdjustedDim)
907  MData += Offset[0];
908  }
909 
910  // __init variant used by the device compiler for ESIMD kernels.
911  // TODO In ESIMD accessors usage is limited for now - access range, mem
912  // range and offset are not supported.
913  void __init_esimd(ConcreteASPtrType Ptr) { MData = Ptr; }
914 
915  ConcreteASPtrType getQualifiedPtr() const { return MData; }
916 
917  template <typename, int, access::mode, access::target, access::placeholder,
918  typename>
919  friend class accessor;
920 
921 #ifndef __SYCL_DEVICE_ONLY__
922  using AccessorBaseHost::impl;
923 #endif
924 
925 public:
926  // Default constructor for objects later initialized with __init member.
927  accessor()
928  : impl({}, detail::InitializedVal<AdjustedDim, range>::template get<0>(),
929  detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
930 
931 #else
932  using AccessorBaseHost::getAccessRange;
933  using AccessorBaseHost::getMemoryRange;
934  using AccessorBaseHost::getOffset;
935 
936  char padding[sizeof(detail::AccessorImplDevice<AdjustedDim>) +
937  sizeof(PtrType) - sizeof(detail::AccessorBaseHost)];
938 
940  return reinterpret_cast<PtrType>(AccessorBaseHost::getPtr());
941  }
942 
943 #endif // __SYCL_DEVICE_ONLY__
944 
945 private:
946  friend class sycl::stream;
947  friend class sycl::ext::intel::experimental::esimd::detail::
948  AccessorPrivateProxy;
949 
950 public:
951  using value_type = DataT;
952  using reference = DataT &;
953  using const_reference = const DataT &;
954 
955  // The list of accessor constructors with their arguments
956  // -------+---------+-------+----+-----+--------------
957  // Dimensions = 0
958  // -------+---------+-------+----+-----+--------------
959  // buffer | | | | | property_list
960  // buffer | handler | | | | property_list
961  // -------+---------+-------+----+-----+--------------
962  // Dimensions >= 1
963  // -------+---------+-------+----+-----+--------------
964  // buffer | | | | | property_list
965  // buffer | | | | tag | property_list
966  // buffer | handler | | | | property_list
967  // buffer | handler | | | tag | property_list
968  // buffer | | range | | | property_list
969  // buffer | | range | | tag | property_list
970  // buffer | handler | range | | | property_list
971  // buffer | handler | range | | tag | property_list
972  // buffer | | range | id | | property_list
973  // buffer | | range | id | tag | property_list
974  // buffer | handler | range | id | | property_list
975  // buffer | handler | range | id | tag | property_list
976  // -------+---------+-------+----+-----+--------------
977 
978 public:
979  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
980  typename detail::enable_if_t<
981  detail::IsRunTimePropertyListT<PropertyListT>::value &&
982  std::is_same<T, DataT>::value && Dims == 0 &&
983  ((!IsPlaceH && IsHostBuf) ||
984  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> * = nullptr>
985  accessor(buffer<T, 1, AllocatorT> &BufferRef,
986  const property_list &PropertyList = {})
987 #ifdef __SYCL_DEVICE_ONLY__
988  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
989  (void)PropertyList;
990 #else
992  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
993  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
994  getAdjustedMode(PropertyList),
995  detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
996  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
997  checkDeviceAccessorBufferSize(BufferRef.size());
998  if (!IsPlaceH)
999  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1000 #endif
1001  }
1002 
1003  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1004  typename... PropTypes,
1005  typename detail::enable_if_t<
1007  std::is_same<T, DataT>::value && Dims == 0 &&
1008  ((!IsPlaceH && IsHostBuf) ||
1009  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> * = nullptr>
1012  &PropertyList = {})
1013 #ifdef __SYCL_DEVICE_ONLY__
1014  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1015  (void)PropertyList;
1016 #else
1017  : AccessorBaseHost(
1018  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1019  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1020  getAdjustedMode(PropertyList),
1021  detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
1022  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1023  checkDeviceAccessorBufferSize(BufferRef.size());
1024  if (!IsPlaceH)
1025  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1026 #endif
1027  }
1028 
1029  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1030  typename = typename detail::enable_if_t<
1031  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1032  std::is_same<T, DataT>::value && (Dims == 0) &&
1033  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1034  accessor(buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1035  const property_list &PropertyList = {})
1036 #ifdef __SYCL_DEVICE_ONLY__
1037  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1038  (void)CommandGroupHandler;
1039  (void)PropertyList;
1040  }
1041 #else
1042  : AccessorBaseHost(
1043  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1044  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1045  getAdjustedMode(PropertyList),
1046  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1047  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1048  checkDeviceAccessorBufferSize(BufferRef.size());
1049  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1050  }
1051 #endif
1052 
1053  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1054  typename... PropTypes,
1055  typename = typename detail::enable_if_t<
1056  detail::IsCxPropertyList<PropertyListT>::value &&
1057  std::is_same<T, DataT>::value && (Dims == 0) &&
1058  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1059  accessor(buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1061  &PropertyList = {})
1062 #ifdef __SYCL_DEVICE_ONLY__
1063  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1064  (void)CommandGroupHandler;
1065  (void)PropertyList;
1066  }
1067 #else
1068  : AccessorBaseHost(
1069  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1070  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1071  getAdjustedMode(PropertyList),
1072  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1073  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1074  checkDeviceAccessorBufferSize(BufferRef.size());
1075  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1076  }
1077 #endif
1078 
1079  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1080  typename = detail::enable_if_t<
1081  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1082  IsSameAsBuffer<T, Dims>() &&
1083  ((!IsPlaceH && IsHostBuf) ||
1084  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1086  const property_list &PropertyList = {})
1087 #ifdef __SYCL_DEVICE_ONLY__
1088  : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1089  (void)PropertyList;
1090  }
1091 #else
1092  : AccessorBaseHost(
1093  /*Offset=*/{0, 0, 0},
1094  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1095  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1096  getAdjustedMode(PropertyList),
1097  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1098  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1099  checkDeviceAccessorBufferSize(BufferRef.size());
1100  if (!IsPlaceH)
1101  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1102  }
1103 #endif
1104 
1105  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1106  typename... PropTypes,
1107  typename = detail::enable_if_t<
1108  detail::IsCxPropertyList<PropertyListT>::value &&
1109  IsSameAsBuffer<T, Dims>() &&
1110  ((!IsPlaceH && IsHostBuf) ||
1111  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1114  &PropertyList = {})
1115 #ifdef __SYCL_DEVICE_ONLY__
1116  : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1117  (void)PropertyList;
1118  }
1119 #else
1120  : AccessorBaseHost(
1121  /*Offset=*/{0, 0, 0},
1122  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1123  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1124  getAdjustedMode(PropertyList),
1125  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1126  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1127  checkDeviceAccessorBufferSize(BufferRef.size());
1128  if (!IsPlaceH)
1129  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1130  }
1131 #endif
1132 
1133 #if __cplusplus > 201402L
1134 
1135  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1136  typename TagT,
1137  typename = detail::enable_if_t<
1138  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1139  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1140  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1141  accessor(buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1142  const property_list &PropertyList = {})
1143  : accessor(BufferRef, PropertyList) {}
1144 
1145  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1146  typename TagT, typename... PropTypes,
1147  typename = detail::enable_if_t<
1148  detail::IsCxPropertyList<PropertyListT>::value &&
1149  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1150  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1151  accessor(buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1152  const ext::oneapi::accessor_property_list<PropTypes...>
1153  &PropertyList = {})
1154  : accessor(BufferRef, PropertyList) {}
1155 #endif
1156 
1157  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1158  typename = detail::enable_if_t<
1159  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1160  IsSameAsBuffer<T, Dims>() &&
1161  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1162  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1163  const property_list &PropertyList = {})
1164 #ifdef __SYCL_DEVICE_ONLY__
1165  : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1166  (void)CommandGroupHandler;
1167  (void)PropertyList;
1168  }
1169 #else
1170  : AccessorBaseHost(
1171  /*Offset=*/{0, 0, 0},
1172  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1173  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1174  getAdjustedMode(PropertyList),
1175  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1176  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1177  checkDeviceAccessorBufferSize(BufferRef.size());
1178  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1179  }
1180 #endif
1181 
1182  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1183  typename... PropTypes,
1184  typename = detail::enable_if_t<
1185  detail::IsCxPropertyList<PropertyListT>::value &&
1186  IsSameAsBuffer<T, Dims>() &&
1187  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1188  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1190  &PropertyList = {})
1191 #ifdef __SYCL_DEVICE_ONLY__
1192  : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1193  (void)CommandGroupHandler;
1194  (void)PropertyList;
1195  }
1196 #else
1197  : AccessorBaseHost(
1198  /*Offset=*/{0, 0, 0},
1199  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1200  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1201  getAdjustedMode(PropertyList),
1202  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1203  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1204  checkDeviceAccessorBufferSize(BufferRef.size());
1205  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1206  }
1207 #endif
1208 
1209 #if __cplusplus > 201402L
1210 
1211  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1212  typename TagT,
1213  typename = detail::enable_if_t<
1214  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1215  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1216  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1217  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1218  TagT, const property_list &PropertyList = {})
1219  : accessor(BufferRef, CommandGroupHandler, PropertyList) {}
1220 
1221  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1222  typename TagT, typename... PropTypes,
1223  typename = detail::enable_if_t<
1224  detail::IsCxPropertyList<PropertyListT>::value &&
1225  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1226  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1227  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1228  TagT,
1229  const ext::oneapi::accessor_property_list<PropTypes...>
1230  &PropertyList = {})
1231  : accessor(BufferRef, CommandGroupHandler, PropertyList) {}
1232 
1233 #endif
1234 
1235  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1236  typename = detail::enable_if_t<
1237  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1238  IsSameAsBuffer<T, Dims>() &&
1239  ((!IsPlaceH && IsHostBuf) ||
1240  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1242  range<Dimensions> AccessRange,
1243  const property_list &PropertyList = {})
1244  : accessor(BufferRef, AccessRange, {}, PropertyList) {}
1245 
1246  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1247  typename... PropTypes,
1248  typename = detail::enable_if_t<
1249  detail::IsCxPropertyList<PropertyListT>::value &&
1250  IsSameAsBuffer<T, Dims>() &&
1251  ((!IsPlaceH && IsHostBuf) ||
1252  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1254  range<Dimensions> AccessRange,
1256  &PropertyList = {})
1257  : accessor(BufferRef, AccessRange, {}, PropertyList) {}
1258 
1259 #if __cplusplus > 201402L
1260 
1261  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1262  typename TagT,
1263  typename = detail::enable_if_t<
1264  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1265  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1266  (IsGlobalBuf || IsConstantBuf)>>
1267  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
1268  range<Dimensions> AccessRange, TagT,
1269  const property_list &PropertyList = {})
1270  : accessor(BufferRef, AccessRange, {}, PropertyList) {}
1271 
1272  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1273  typename TagT, typename... PropTypes,
1274  typename = detail::enable_if_t<
1275  detail::IsCxPropertyList<PropertyListT>::value &&
1276  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1277  (IsGlobalBuf || IsConstantBuf)>>
1278  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
1279  range<Dimensions> AccessRange, TagT,
1280  const ext::oneapi::accessor_property_list<PropTypes...>
1281  &PropertyList = {})
1282  : accessor(BufferRef, AccessRange, {}, PropertyList) {}
1283 #endif
1284 
1285  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1286  typename = detail::enable_if_t<
1287  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1288  IsSameAsBuffer<T, Dims>() &&
1289  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1290  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1291  range<Dimensions> AccessRange,
1292  const property_list &PropertyList = {})
1293  : accessor(BufferRef, CommandGroupHandler, AccessRange, {},
1294  PropertyList) {}
1295 
1296  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1297  typename... PropTypes,
1298  typename = detail::enable_if_t<
1299  detail::IsCxPropertyList<PropertyListT>::value &&
1300  IsSameAsBuffer<T, Dims>() &&
1301  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1302  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1303  range<Dimensions> AccessRange,
1305  &PropertyList = {})
1306  : accessor(BufferRef, CommandGroupHandler, AccessRange, {},
1307  PropertyList) {}
1308 
1309 #if __cplusplus > 201402L
1310 
1311  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1312  typename TagT,
1313  typename = detail::enable_if_t<
1314  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1315  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1316  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1317  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1318  range<Dimensions> AccessRange, TagT,
1319  const property_list &PropertyList = {})
1320  : accessor(BufferRef, CommandGroupHandler, AccessRange, {},
1321  PropertyList) {}
1322 
1323  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1324  typename TagT, typename... PropTypes,
1325  typename = detail::enable_if_t<
1326  detail::IsCxPropertyList<PropertyListT>::value &&
1327  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1328  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1329  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1330  range<Dimensions> AccessRange, TagT,
1331  const ext::oneapi::accessor_property_list<PropTypes...>
1332  &PropertyList = {})
1333  : accessor(BufferRef, CommandGroupHandler, AccessRange, {},
1334  PropertyList) {}
1335 #endif
1336 
1337  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1338  typename = detail::enable_if_t<
1339  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1340  IsSameAsBuffer<T, Dims>() &&
1341  ((!IsPlaceH && IsHostBuf) ||
1342  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1344  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
1345  const property_list &PropertyList = {})
1346 #ifdef __SYCL_DEVICE_ONLY__
1347  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1348  (void)PropertyList;
1349  }
1350 #else
1351  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1352  detail::convertToArrayOfN<3, 1>(AccessRange),
1353  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1354  getAdjustedMode(PropertyList),
1355  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1356  sizeof(DataT), BufferRef.OffsetInBytes,
1357  BufferRef.IsSubBuffer) {
1358  checkDeviceAccessorBufferSize(BufferRef.size());
1359  if (!IsPlaceH)
1360  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1361  }
1362 #endif
1363 
1364  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1365  typename... PropTypes,
1366  typename = detail::enable_if_t<
1367  detail::IsCxPropertyList<PropertyListT>::value &&
1368  IsSameAsBuffer<T, Dims>() &&
1369  ((!IsPlaceH && IsHostBuf) ||
1370  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1372  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
1374  &PropertyList = {})
1375 #ifdef __SYCL_DEVICE_ONLY__
1376  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1377  (void)PropertyList;
1378  }
1379 #else
1380  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1381  detail::convertToArrayOfN<3, 1>(AccessRange),
1382  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1383  getAdjustedMode(PropertyList),
1384  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1385  sizeof(DataT), BufferRef.OffsetInBytes,
1386  BufferRef.IsSubBuffer) {
1387  checkDeviceAccessorBufferSize(BufferRef.size());
1388  if (!IsPlaceH)
1389  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1390  }
1391 #endif
1392 
1393 #if __cplusplus > 201402L
1394 
1395  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1396  typename TagT,
1397  typename = detail::enable_if_t<
1398  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1399  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1400  (IsGlobalBuf || IsConstantBuf)>>
1401  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
1402  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
1403  const property_list &PropertyList = {})
1404  : accessor(BufferRef, AccessRange, AccessOffset, PropertyList) {}
1405 
1406  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1407  typename TagT, typename... PropTypes,
1408  typename = detail::enable_if_t<
1409  detail::IsCxPropertyList<PropertyListT>::value &&
1410  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1411  (IsGlobalBuf || IsConstantBuf)>>
1412  accessor(buffer<T, Dims, AllocatorT> &BufferRef,
1413  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
1414  const ext::oneapi::accessor_property_list<PropTypes...>
1415  &PropertyList = {})
1416  : accessor(BufferRef, AccessRange, AccessOffset, PropertyList) {}
1417 #endif
1418 
1419  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1420  typename = detail::enable_if_t<
1421  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1422  IsSameAsBuffer<T, Dims>() &&
1423  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1424  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1425  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
1426  const property_list &PropertyList = {})
1427 #ifdef __SYCL_DEVICE_ONLY__
1428  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1429  (void)CommandGroupHandler;
1430  (void)PropertyList;
1431  }
1432 #else
1433  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1434  detail::convertToArrayOfN<3, 1>(AccessRange),
1435  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1436  getAdjustedMode(PropertyList),
1437  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1438  sizeof(DataT), BufferRef.OffsetInBytes,
1439  BufferRef.IsSubBuffer) {
1440  checkDeviceAccessorBufferSize(BufferRef.size());
1441  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1442  }
1443 #endif
1444 
1445  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1446  typename... PropTypes,
1447  typename = detail::enable_if_t<
1448  detail::IsCxPropertyList<PropertyListT>::value &&
1449  IsSameAsBuffer<T, Dims>() &&
1450  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1451  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1452  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
1454  &PropertyList = {})
1455 #ifdef __SYCL_DEVICE_ONLY__
1456  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1457  (void)CommandGroupHandler;
1458  (void)PropertyList;
1459  }
1460 #else
1461  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1462  detail::convertToArrayOfN<3, 1>(AccessRange),
1463  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1464  getAdjustedMode(PropertyList),
1465  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1466  sizeof(DataT), BufferRef.OffsetInBytes,
1467  BufferRef.IsSubBuffer) {
1468  checkDeviceAccessorBufferSize(BufferRef.size());
1469  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1470  }
1471 #endif
1472 
1473 #if __cplusplus > 201402L
1474 
1475  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1476  typename TagT,
1477  typename = detail::enable_if_t<
1478  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1479  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1480  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1481  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1482  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
1483  const property_list &PropertyList = {})
1484  : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1485  PropertyList) {}
1486 
1487  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1488  typename TagT, typename... PropTypes,
1489  typename = detail::enable_if_t<
1490  detail::IsCxPropertyList<PropertyListT>::value &&
1491  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1492  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1493  accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1494  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
1495  const ext::oneapi::accessor_property_list<PropTypes...>
1496  &PropertyList = {})
1497  : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1498  PropertyList) {}
1499 #endif
1500 
1501  template <typename... NewPropsT>
1503  const accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
1505 #ifdef __SYCL_DEVICE_ONLY__
1506  : impl(Other.impl)
1507 #else
1508  : detail::AccessorBaseHost(Other)
1509 #endif
1510  {
1512  "Conversion is only available for accessor_property_list");
1513  static_assert(
1514  PropertyListT::template areSameCompileTimeProperties<NewPropsT...>(),
1515  "Compile-time-constant properties must be the same");
1516  }
1517 
1518  constexpr bool is_placeholder() const { return IsPlaceH; }
1519 
1520  size_t get_size() const { return getAccessRange().size() * sizeof(DataT); }
1521 
1522  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1523  size_t get_count() const { return size(); }
1524  size_t size() const noexcept { return getAccessRange().size(); }
1525 
1526  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
1528  return detail::convertToArrayOfN<Dimensions, 1>(getAccessRange());
1529  }
1530 
1531  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
1533  return detail::convertToArrayOfN<Dimensions, 0>(getOffset());
1534  }
1535 
1536  template <int Dims = Dimensions, typename RefT = RefType,
1537  typename = detail::enable_if_t<Dims == 0 && IsAccessAnyWrite &&
1538  !std::is_const<RefT>::value>>
1539  operator RefType() const {
1540  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
1541  return *(getQualifiedPtr() + LinearIndex);
1542  }
1543 
1544  template <int Dims = Dimensions,
1546  operator ConstRefType() const {
1547  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
1548  return *(getQualifiedPtr() + LinearIndex);
1549  }
1550 
1551  template <int Dims = Dimensions,
1552  typename = detail::enable_if_t<(Dims > 0) && IsAccessAnyWrite>>
1553  RefType operator[](id<Dimensions> Index) const {
1554  const size_t LinearIndex = getLinearIndex(Index);
1555  return getQualifiedPtr()[LinearIndex];
1556  }
1557 
1558  template <int Dims = Dimensions>
1559  typename detail::enable_if_t<(Dims > 0) && IsAccessReadOnly, ConstRefType>
1560  operator[](id<Dimensions> Index) const {
1561  const size_t LinearIndex = getLinearIndex(Index);
1562  return getQualifiedPtr()[LinearIndex];
1563  }
1564 
1565  template <int Dims = Dimensions>
1566  operator typename detail::enable_if_t<Dims == 0 &&
1567 #ifdef __ENABLE_USM_ADDR_SPACE__
1568  AccessMode == access::mode::atomic,
1569  atomic<DataT>>() const {
1570 #else
1571  AccessMode == access::mode::atomic,
1573 #endif
1574  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
1575  return atomic<DataT, AS>(
1576  multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
1577  }
1578 
1579  template <int Dims = Dimensions>
1580  typename detail::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
1582  operator[](id<Dimensions> Index) const {
1583  const size_t LinearIndex = getLinearIndex(Index);
1584  return atomic<DataT, AS>(
1585  multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
1586  }
1587 
1588  template <int Dims = Dimensions>
1589  typename detail::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
1591  operator[](size_t Index) const {
1592  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
1593  return atomic<DataT, AS>(
1594  multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
1595  }
1596  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>
1597  typename AccessorCommonT::template AccessorSubscript<Dims - 1>
1598  operator[](size_t Index) const {
1599  return AccessorSubscript<Dims - 1>(*this, Index);
1600  }
1601 
1602  template <access::target AccessTarget_ = AccessTarget,
1603  typename = detail::enable_if_t<AccessTarget_ ==
1604  access::target::host_buffer>>
1605  DataT *get_pointer() const {
1606  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
1607  return getQualifiedPtr() + LinearIndex;
1608  }
1609 
1610  template <
1611  access::target AccessTarget_ = AccessTarget,
1614  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
1615  return global_ptr<DataT>(getQualifiedPtr() + LinearIndex);
1616  }
1617 
1618  template <access::target AccessTarget_ = AccessTarget,
1619  typename = detail::enable_if_t<AccessTarget_ ==
1620  access::target::constant_buffer>>
1622  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
1623  return constant_ptr<DataT>(getQualifiedPtr() + LinearIndex);
1624  }
1625 
1626  bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
1627  bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
1628 
1629 private:
1630  void checkDeviceAccessorBufferSize(const size_t elemInBuffer) {
1631  if (!IsHostBuf && elemInBuffer == 0)
1632  throw cl::sycl::invalid_object_error(
1633  "SYCL buffer size is zero. To create a device accessor, SYCL "
1634  "buffer size must be greater than zero.",
1636  }
1637 };
1638 
1639 #if __cplusplus > 201402L
1640 
1641 template <typename DataT, int Dimensions, typename AllocatorT>
1642 accessor(buffer<DataT, Dimensions, AllocatorT>)
1643  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
1644  access::placeholder::true_t>;
1645 
1646 template <typename DataT, int Dimensions, typename AllocatorT,
1647  typename... PropsT>
1648 accessor(buffer<DataT, Dimensions, AllocatorT>,
1649  const ext::oneapi::accessor_property_list<PropsT...> &)
1650  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
1651  access::placeholder::true_t,
1652  ext::oneapi::accessor_property_list<PropsT...>>;
1653 
1654 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
1655 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1)
1656  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
1657  detail::deduceAccessTarget<Type1, Type1>(target::device),
1658  access::placeholder::true_t>;
1659 
1660 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1661  typename... PropsT>
1662 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1,
1663  const ext::oneapi::accessor_property_list<PropsT...> &)
1664  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
1665  detail::deduceAccessTarget<Type1, Type1>(target::device),
1666  access::placeholder::true_t,
1667  ext::oneapi::accessor_property_list<PropsT...>>;
1668 
1669 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1670  typename Type2>
1671 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2)
1672  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
1673  detail::deduceAccessTarget<Type1, Type2>(target::device),
1674  access::placeholder::true_t>;
1675 
1676 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1677  typename Type2, typename... PropsT>
1678 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2,
1679  const ext::oneapi::accessor_property_list<PropsT...> &)
1680  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
1681  detail::deduceAccessTarget<Type1, Type2>(target::device),
1682  access::placeholder::true_t,
1683  ext::oneapi::accessor_property_list<PropsT...>>;
1684 
1685 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1686  typename Type2, typename Type3>
1687 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3)
1688  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
1689  detail::deduceAccessTarget<Type2, Type3>(target::device),
1690  access::placeholder::true_t>;
1691 
1692 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1693  typename Type2, typename Type3, typename... PropsT>
1694 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3,
1695  const ext::oneapi::accessor_property_list<PropsT...> &)
1696  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
1697  detail::deduceAccessTarget<Type2, Type3>(target::device),
1698  access::placeholder::true_t,
1699  ext::oneapi::accessor_property_list<PropsT...>>;
1700 
1701 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1702  typename Type2, typename Type3, typename Type4>
1703 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
1704  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
1705  detail::deduceAccessTarget<Type3, Type4>(target::device),
1706  access::placeholder::true_t>;
1707 
1708 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1709  typename Type2, typename Type3, typename Type4, typename... PropsT>
1710 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
1711  const ext::oneapi::accessor_property_list<PropsT...> &)
1712  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
1713  detail::deduceAccessTarget<Type3, Type4>(target::device),
1714  access::placeholder::true_t,
1715  ext::oneapi::accessor_property_list<PropsT...>>;
1716 
1717 template <typename DataT, int Dimensions, typename AllocatorT>
1718 accessor(buffer<DataT, Dimensions, AllocatorT>, handler)
1719  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
1720  access::placeholder::false_t>;
1721 
1722 template <typename DataT, int Dimensions, typename AllocatorT,
1723  typename... PropsT>
1724 accessor(buffer<DataT, Dimensions, AllocatorT>, handler,
1725  const ext::oneapi::accessor_property_list<PropsT...> &)
1726  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
1727  access::placeholder::false_t,
1728  ext::oneapi::accessor_property_list<PropsT...>>;
1729 
1730 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
1731 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1)
1732  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
1733  detail::deduceAccessTarget<Type1, Type1>(target::device),
1734  access::placeholder::false_t>;
1735 
1736 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1737  typename... PropsT>
1738 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1,
1739  const ext::oneapi::accessor_property_list<PropsT...> &)
1740  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
1741  detail::deduceAccessTarget<Type1, Type1>(target::device),
1742  access::placeholder::false_t,
1743  ext::oneapi::accessor_property_list<PropsT...>>;
1744 
1745 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1746  typename Type2>
1747 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2)
1748  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
1749  detail::deduceAccessTarget<Type1, Type2>(target::device),
1750  access::placeholder::false_t>;
1751 
1752 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1753  typename Type2, typename... PropsT>
1754 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2,
1755  const ext::oneapi::accessor_property_list<PropsT...> &)
1756  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
1757  detail::deduceAccessTarget<Type1, Type2>(target::device),
1758  access::placeholder::false_t,
1759  ext::oneapi::accessor_property_list<PropsT...>>;
1760 
1761 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1762  typename Type2, typename Type3>
1763 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3)
1764  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
1765  detail::deduceAccessTarget<Type2, Type3>(target::device),
1766  access::placeholder::false_t>;
1767 
1768 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1769  typename Type2, typename Type3, typename... PropsT>
1770 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
1771  const ext::oneapi::accessor_property_list<PropsT...> &)
1772  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
1773  detail::deduceAccessTarget<Type2, Type3>(target::device),
1774  access::placeholder::false_t,
1775  ext::oneapi::accessor_property_list<PropsT...>>;
1776 
1777 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1778  typename Type2, typename Type3, typename Type4>
1779 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
1780  Type4)
1781  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
1782  detail::deduceAccessTarget<Type3, Type4>(target::device),
1783  access::placeholder::false_t>;
1784 
1785 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1786  typename Type2, typename Type3, typename Type4, typename... PropsT>
1787 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
1788  Type4, const ext::oneapi::accessor_property_list<PropsT...> &)
1789  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
1790  detail::deduceAccessTarget<Type3, Type4>(target::device),
1791  access::placeholder::false_t,
1792  ext::oneapi::accessor_property_list<PropsT...>>;
1793 #endif
1794 
1798 template <typename DataT, int Dimensions, access::mode AccessMode,
1799  access::placeholder IsPlaceholder>
1800 class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
1801  access::target::local, IsPlaceholder> :
1802 #ifndef __SYCL_DEVICE_ONLY__
1804 #endif
1805  public detail::accessor_common<DataT, Dimensions, AccessMode,
1806  access::target::local, IsPlaceholder> {
1807 protected:
1808  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
1809 
1810  using AccessorCommonT =
1811  detail::accessor_common<DataT, Dimensions, AccessMode,
1812  access::target::local, IsPlaceholder>;
1813 
1814  using AccessorCommonT::AS;
1815  using AccessorCommonT::IsAccessAnyWrite;
1816  template <int Dims>
1817  using AccessorSubscript =
1818  typename AccessorCommonT::template AccessorSubscript<Dims>;
1819 
1821 
1824 
1825 #ifdef __SYCL_DEVICE_ONLY__
1827 
1828  sycl::range<AdjustedDim> &getSize() { return impl.MemRange; }
1829  const sycl::range<AdjustedDim> &getSize() const { return impl.MemRange; }
1830 
1831  void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
1832  range<AdjustedDim>, id<AdjustedDim>) {
1833  MData = Ptr;
1834 #pragma unroll
1835  for (int I = 0; I < AdjustedDim; ++I)
1836  getSize()[I] = AccessRange[I];
1837  }
1838 
1839 public:
1840  // Default constructor for objects later initialized with __init member.
1841  accessor()
1842  : impl(detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
1843 
1844 protected:
1845  ConcreteASPtrType getQualifiedPtr() const { return MData; }
1846 
1847  ConcreteASPtrType MData;
1848 
1849 #else
1850 
1851  char padding[sizeof(detail::LocalAccessorBaseDevice<AdjustedDim>) +
1853  using detail::LocalAccessorBaseHost::getSize;
1854 
1856  return reinterpret_cast<PtrType>(LocalAccessorBaseHost::getPtr());
1857  }
1858 
1859 #endif // __SYCL_DEVICE_ONLY__
1860 
1861  // Method which calculates linear offset for the ID using Range and Offset.
1862  template <int Dims = AdjustedDim> size_t getLinearIndex(id<Dims> Id) const {
1863  size_t Result = 0;
1864  for (int I = 0; I < Dims; ++I)
1865  Result = Result * getSize()[I] + Id[I];
1866  return Result;
1867  }
1868 
1869 public:
1870  using value_type = DataT;
1871  using reference = DataT &;
1872  using const_reference = const DataT &;
1873 
1874  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 0>>
1876 #ifdef __SYCL_DEVICE_ONLY__
1877  : impl(range<AdjustedDim>{1}) {
1878  }
1879 #else
1880  : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT)) {
1881  }
1882 #endif
1883 
1884  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 0>>
1885  accessor(handler &, const property_list &propList)
1886 #ifdef __SYCL_DEVICE_ONLY__
1887  : impl(range<AdjustedDim>{1}) {
1888  (void)propList;
1889  }
1890 #else
1891  : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT)) {
1892  (void)propList;
1893  }
1894 #endif
1895 
1896  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
1898 #ifdef __SYCL_DEVICE_ONLY__
1899  : impl(AllocationSize) {
1900  }
1901 #else
1902  : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
1903  AdjustedDim, sizeof(DataT)) {
1904  }
1905 #endif
1906 
1907  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
1909  const property_list &propList)
1910 #ifdef __SYCL_DEVICE_ONLY__
1911  : impl(AllocationSize) {
1912  (void)propList;
1913  }
1914 #else
1915  : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
1916  AdjustedDim, sizeof(DataT)) {
1917  (void)propList;
1918  }
1919 #endif
1920 
1921  size_t get_size() const { return getSize().size() * sizeof(DataT); }
1922 
1923  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1924  size_t get_count() const { return size(); }
1925  size_t size() const noexcept { return getSize().size(); }
1926 
1927  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
1929  return detail::convertToArrayOfN<Dims, 1>(getSize());
1930  }
1931 
1932  template <int Dims = Dimensions,
1934  operator RefType() const {
1935  return *getQualifiedPtr();
1936  }
1937 
1938  template <int Dims = Dimensions,
1939  typename = detail::enable_if_t<(Dims > 0) && IsAccessAnyWrite>>
1940  RefType operator[](id<Dimensions> Index) const {
1941  const size_t LinearIndex = getLinearIndex(Index);
1942  return getQualifiedPtr()[LinearIndex];
1943  }
1944 
1945  template <int Dims = Dimensions,
1947  RefType operator[](size_t Index) const {
1948  return getQualifiedPtr()[Index];
1949  }
1950 
1951  template <int Dims = Dimensions>
1952  operator typename detail::enable_if_t<
1953  Dims == 0 && AccessMode == access::mode::atomic, atomic<DataT, AS>>()
1954  const {
1955  return atomic<DataT, AS>(multi_ptr<DataT, AS>(getQualifiedPtr()));
1956  }
1957 
1958  template <int Dims = Dimensions>
1959  typename detail::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
1961  operator[](id<Dimensions> Index) const {
1962  const size_t LinearIndex = getLinearIndex(Index);
1963  return atomic<DataT, AS>(
1964  multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
1965  }
1966 
1967  template <int Dims = Dimensions>
1968  typename detail::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
1970  operator[](size_t Index) const {
1971  return atomic<DataT, AS>(multi_ptr<DataT, AS>(getQualifiedPtr() + Index));
1972  }
1973 
1974  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>
1975  typename AccessorCommonT::template AccessorSubscript<Dims - 1>
1976  operator[](size_t Index) const {
1977  return AccessorSubscript<Dims - 1>(*this, Index);
1978  }
1979 
1981  return local_ptr<DataT>(getQualifiedPtr());
1982  }
1983 
1984  bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
1985  bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
1986 };
1987 
1993 template <typename DataT, int Dimensions, access::mode AccessMode,
1994  access::placeholder IsPlaceholder>
1995 class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
1996  access::target::image, IsPlaceholder>
1997  : public detail::image_accessor<DataT, Dimensions, AccessMode,
1998  access::target::image, IsPlaceholder> {
1999 public:
2000  template <typename AllocatorT>
2002  handler &CommandGroupHandler)
2003  : detail::image_accessor<DataT, Dimensions, AccessMode,
2004  access::target::image, IsPlaceholder>(
2005  Image, CommandGroupHandler,
2006  (detail::getSyclObjImpl(Image))->getElementSize()) {
2007 #ifndef __SYCL_DEVICE_ONLY__
2008  detail::associateWithHandler(CommandGroupHandler, this,
2009  access::target::image);
2010 #endif
2011  }
2012 
2013  template <typename AllocatorT>
2015  handler &CommandGroupHandler, const property_list &propList)
2016  : detail::image_accessor<DataT, Dimensions, AccessMode,
2017  access::target::image, IsPlaceholder>(
2018  Image, CommandGroupHandler,
2019  (detail::getSyclObjImpl(Image))->getElementSize()) {
2020  (void)propList;
2021 #ifndef __SYCL_DEVICE_ONLY__
2022  detail::associateWithHandler(CommandGroupHandler, this,
2023  access::target::image);
2024 #endif
2025  }
2026 #ifdef __SYCL_DEVICE_ONLY__
2027 private:
2028  using OCLImageTy =
2029  typename detail::opencl_image_type<Dimensions, AccessMode,
2030  access::target::image>::type;
2031 
2032  // Front End requires this method to be defined in the accessor class.
2033  // It does not call the base class's init method.
2034  void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
2035 
2036  // __init variant used by the device compiler for ESIMD kernels.
2037  void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
2038 
2039 public:
2040  // Default constructor for objects later initialized with __init member.
2041  accessor() = default;
2042 #endif
2043 };
2044 
2052 template <typename DataT, int Dimensions, access::mode AccessMode,
2053  access::placeholder IsPlaceholder>
2054 class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
2055  access::target::host_image, IsPlaceholder>
2056  : public detail::image_accessor<DataT, Dimensions, AccessMode,
2057  access::target::host_image, IsPlaceholder> {
2058 public:
2059  template <typename AllocatorT>
2061  : detail::image_accessor<DataT, Dimensions, AccessMode,
2062  access::target::host_image, IsPlaceholder>(
2063  Image, (detail::getSyclObjImpl(Image))->getElementSize()) {}
2064 
2065  template <typename AllocatorT>
2067  const property_list &propList)
2068  : detail::image_accessor<DataT, Dimensions, AccessMode,
2069  access::target::host_image, IsPlaceholder>(
2070  Image, (detail::getSyclObjImpl(Image))->getElementSize()) {
2071  (void)propList;
2072  }
2073 };
2074 
2083 template <typename DataT, int Dimensions, access::mode AccessMode,
2084  access::placeholder IsPlaceholder>
2085 class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
2086  access::target::image_array, IsPlaceholder>
2087  : public detail::image_accessor<DataT, Dimensions + 1, AccessMode,
2088  access::target::image, IsPlaceholder> {
2089 #ifdef __SYCL_DEVICE_ONLY__
2090 private:
2091  using OCLImageTy =
2092  typename detail::opencl_image_type<Dimensions + 1, AccessMode,
2093  access::target::image>::type;
2094 
2095  // Front End requires this method to be defined in the accessor class.
2096  // It does not call the base class's init method.
2097  void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
2098 
2099  // __init variant used by the device compiler for ESIMD kernels.
2100  void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
2101 
2102 public:
2103  // Default constructor for objects later initialized with __init member.
2104  accessor() = default;
2105 #endif
2106 public:
2107  template <typename AllocatorT>
2109  handler &CommandGroupHandler)
2110  : detail::image_accessor<DataT, Dimensions + 1, AccessMode,
2111  access::target::image, IsPlaceholder>(
2112  Image, CommandGroupHandler,
2113  (detail::getSyclObjImpl(Image))->getElementSize()) {
2114 #ifndef __SYCL_DEVICE_ONLY__
2115  detail::associateWithHandler(CommandGroupHandler, this,
2116  access::target::image_array);
2117 #endif
2118  }
2119 
2120  template <typename AllocatorT>
2122  handler &CommandGroupHandler, const property_list &propList)
2123  : detail::image_accessor<DataT, Dimensions + 1, AccessMode,
2124  access::target::image, IsPlaceholder>(
2125  Image, CommandGroupHandler,
2126  (detail::getSyclObjImpl(Image))->getElementSize()) {
2127  (void)propList;
2128 #ifndef __SYCL_DEVICE_ONLY__
2129  detail::associateWithHandler(CommandGroupHandler, this,
2130  access::target::image_array);
2131 #endif
2132  }
2133 
2135  operator[](size_t Index) const {
2136  return detail::__image_array_slice__<DataT, Dimensions, AccessMode,
2137  IsPlaceholder>(*this, Index);
2138  }
2139 };
2140 
2141 template <typename DataT, int Dimensions = 1,
2142  access_mode AccessMode = access_mode::read_write>
2144  : public accessor<DataT, Dimensions, AccessMode, target::host_buffer,
2145  access::placeholder::false_t> {
2146 protected:
2147  using AccessorT = accessor<DataT, Dimensions, AccessMode, target::host_buffer,
2148  access::placeholder::false_t>;
2149 
2150  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
2151 
2152  template <typename T, int Dims> static constexpr bool IsSameAsBuffer() {
2153  return std::is_same<T, DataT>::value && (Dims > 0) && (Dims == Dimensions);
2154  }
2155 
2156 #if __cplusplus > 201402L
2157 
2158  template <typename TagT> static constexpr bool IsValidTag() {
2159  return std::is_same<TagT, mode_tag_t<AccessMode>>::value;
2160  }
2161 
2162 #endif
2163 
2164  void
2165  __init(typename accessor<DataT, Dimensions, AccessMode, target::host_buffer,
2166  access::placeholder::false_t>::ConcreteASPtrType Ptr,
2167  range<AdjustedDim> AccessRange, range<AdjustedDim> MemRange,
2168  id<AdjustedDim> Offset) {
2169  AccessorT::__init(Ptr, AccessRange, MemRange, Offset);
2170  }
2171 
2172 public:
2174 
2175  // The list of host_accessor constructors with their arguments
2176  // -------+---------+-------+----+----------+--------------
2177  // Dimensions = 0
2178  // -------+---------+-------+----+----------+--------------
2179  // buffer | | | | | property_list
2180  // buffer | handler | | | | property_list
2181  // -------+---------+-------+----+----------+--------------
2182  // Dimensions >= 1
2183  // -------+---------+-------+----+----------+--------------
2184  // buffer | | | | | property_list
2185  // buffer | | | | mode_tag | property_list
2186  // buffer | handler | | | | property_list
2187  // buffer | handler | | | mode_tag | property_list
2188  // buffer | | range | | | property_list
2189  // buffer | | range | | mode_tag | property_list
2190  // buffer | handler | range | | | property_list
2191  // buffer | handler | range | | mode_tag | property_list
2192  // buffer | | range | id | | property_list
2193  // buffer | | range | id | mode_tag | property_list
2194  // buffer | handler | range | id | | property_list
2195  // buffer | handler | range | id | mode_tag | property_list
2196  // -------+---------+-------+----+----------+--------------
2197 
2198  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2199  typename = typename detail::enable_if_t<
2200  std::is_same<T, DataT>::value && Dims == 0>>
2202  const property_list &PropertyList = {})
2203  : AccessorT(BufferRef, PropertyList) {}
2204 
2205  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2206  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2208  const property_list &PropertyList = {})
2209  : AccessorT(BufferRef, PropertyList) {}
2210 
2211 #if __cplusplus > 201402L
2212 
2213  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2214  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2215  host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2216  mode_tag_t<AccessMode>, const property_list &PropertyList = {})
2217  : host_accessor(BufferRef, PropertyList) {}
2218 
2219 #endif
2220 
2221  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2222  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2224  handler &CommandGroupHandler,
2225  const property_list &PropertyList = {})
2226  : AccessorT(BufferRef, CommandGroupHandler, PropertyList) {}
2227 
2228 #if __cplusplus > 201402L
2229 
2230  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2231  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2232  host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2233  handler &CommandGroupHandler, mode_tag_t<AccessMode>,
2234  const property_list &PropertyList = {})
2235  : host_accessor(BufferRef, CommandGroupHandler, PropertyList) {}
2236 
2237 #endif
2238 
2239  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2240  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2242  range<Dimensions> AccessRange,
2243  const property_list &PropertyList = {})
2244  : AccessorT(BufferRef, AccessRange, {}, PropertyList) {}
2245 
2246 #if __cplusplus > 201402L
2247 
2248  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2249  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2250  host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2251  range<Dimensions> AccessRange, mode_tag_t<AccessMode>,
2252  const property_list &PropertyList = {})
2253  : host_accessor(BufferRef, AccessRange, {}, PropertyList) {}
2254 
2255 #endif
2256 
2257  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2258  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2260  handler &CommandGroupHandler, range<Dimensions> AccessRange,
2261  const property_list &PropertyList = {})
2262  : AccessorT(BufferRef, CommandGroupHandler, AccessRange, {},
2263  PropertyList) {}
2264 
2265 #if __cplusplus > 201402L
2266 
2267  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2268  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2269  host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2270  handler &CommandGroupHandler, range<Dimensions> AccessRange,
2271  mode_tag_t<AccessMode>, const property_list &PropertyList = {})
2272  : host_accessor(BufferRef, CommandGroupHandler, AccessRange, {},
2273  PropertyList) {}
2274 
2275 #endif
2276 
2277  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2278  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2280  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
2281  const property_list &PropertyList = {})
2282  : AccessorT(BufferRef, AccessRange, AccessOffset, PropertyList) {}
2283 
2284 #if __cplusplus > 201402L
2285 
2286  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2287  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2288  host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2289  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
2290  mode_tag_t<AccessMode>, const property_list &PropertyList = {})
2291  : host_accessor(BufferRef, AccessRange, AccessOffset, PropertyList) {}
2292 
2293 #endif
2294 
2295  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2296  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2298  handler &CommandGroupHandler, range<Dimensions> AccessRange,
2299  id<Dimensions> AccessOffset,
2300  const property_list &PropertyList = {})
2301  : AccessorT(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2302  PropertyList) {}
2303 
2304 #if __cplusplus > 201402L
2305 
2306  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2307  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2308  host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2309  handler &CommandGroupHandler, range<Dimensions> AccessRange,
2310  id<Dimensions> AccessOffset, mode_tag_t<AccessMode>,
2311  const property_list &PropertyList = {})
2312  : host_accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2313  PropertyList) {}
2314 
2315 #endif
2316 };
2317 
2318 #if __cplusplus > 201402L
2319 
2320 template <typename DataT, int Dimensions, typename AllocatorT>
2321 host_accessor(buffer<DataT, Dimensions, AllocatorT>)
2322  -> host_accessor<DataT, Dimensions, access::mode::read_write>;
2323 
2324 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
2325 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1)
2326  -> host_accessor<DataT, Dimensions,
2327  detail::deduceAccessMode<Type1, Type1>()>;
2328 
2329 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2330  typename Type2>
2331 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2)
2332  -> host_accessor<DataT, Dimensions,
2333  detail::deduceAccessMode<Type1, Type2>()>;
2334 
2335 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2336  typename Type2, typename Type3>
2337 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3)
2338  -> host_accessor<DataT, Dimensions,
2339  detail::deduceAccessMode<Type2, Type3>()>;
2340 
2341 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2342  typename Type2, typename Type3, typename Type4>
2343 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
2344  -> host_accessor<DataT, Dimensions,
2345  detail::deduceAccessMode<Type3, Type4>()>;
2346 
2347 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2348  typename Type2, typename Type3, typename Type4, typename Type5>
2349 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
2350  Type5) -> host_accessor<DataT, Dimensions,
2351  detail::deduceAccessMode<Type4, Type5>()>;
2352 
2353 #endif
2354 
2355 } // namespace sycl
2356 } // __SYCL_INLINE_NAMESPACE(cl)
2357 
2358 namespace std {
2359 template <typename DataT, int Dimensions, cl::sycl::access::mode AccessMode,
2360  cl::sycl::access::target AccessTarget,
2361  cl::sycl::access::placeholder IsPlaceholder>
2362 struct hash<cl::sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
2363  IsPlaceholder>> {
2364  using AccType = cl::sycl::accessor<DataT, Dimensions, AccessMode,
2365  AccessTarget, IsPlaceholder>;
2366 
2367  size_t operator()(const AccType &A) const {
2368 #ifdef __SYCL_DEVICE_ONLY__
2369  // Hash is not supported on DEVICE. Just return 0 here.
2370  (void)A;
2371  return 0;
2372 #else
2373  // getSyclObjImpl() here returns a pointer to either AccessorImplHost
2374  // or LocalAccessorImplHost depending on the AccessTarget.
2375  auto AccImplPtr = cl::sycl::detail::getSyclObjImpl(A);
2376  return hash<decltype(AccImplPtr)>()(AccImplPtr);
2377 #endif
2378  }
2379 };
2380 
2381 } // namespace std
cl::sycl::ext::intel::experimental::esimd::detail::AccessorPrivateProxy
Definition: memory_intrin.hpp:31
cl::sycl::detail::associateWithHandler
void associateWithHandler(handler &, AccessorBaseHost *, access::target)
Definition: handler_proxy.cpp:17
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::image, IsPlaceholder >::accessor
accessor(cl::sycl::image< Dimensions, AllocatorT > &Image, handler &CommandGroupHandler)
Definition: accessor.hpp:2001
property_list.hpp
cl::sycl::buffer::get_range
range< dimensions > get_range() const
Definition: buffer.hpp:269
cl::sycl::accessor::get_pointer
constant_ptr< DataT > get_pointer() const
Definition: accessor.hpp:1621
type
cl::sycl::detail::getLinearIndex
size_t getLinearIndex(const T< Dims > &Index, const U< Dims > &Range)
Definition: common.hpp:314
cl::sycl::detail::get< 0 >
Definition: tuple.hpp:75
cl::sycl::detail::accessor_common< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::ConstRefType
const DataT & ConstRefType
Definition: accessor.hpp:297
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::image_array, IsPlaceholder >::accessor
accessor(cl::sycl::image< Dimensions+1, AllocatorT > &Image, handler &CommandGroupHandler, const property_list &propList)
Definition: accessor.hpp:2121
cl::sycl::detail::accessor_common::AccessorSubscript
Definition: accessor.hpp:305
T
cl::sycl::info::device
device
Definition: info_desc.hpp:49
cl::sycl::accessor::accessor
accessor(buffer< T, 1, AllocatorT > &BufferRef, handler &CommandGroupHandler, const property_list &PropertyList={})
Definition: accessor.hpp:1034
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={})
Definition: accessor.hpp:1188
cl::sycl::detail::image_accessor::read
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:623
cl::sycl::detail::__image_array_slice__::read
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:720
cl::sycl::host_accessor::host_accessor
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, const property_list &PropertyList={})
Definition: accessor.hpp:2223
PI_INVALID_OPERATION
@ PI_INVALID_OPERATION
Definition: pi.h:84
__SYCL2020_DEPRECATED
#define __SYCL2020_DEPRECATED(message)
Definition: defines_elementary.hpp:56
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::ConcreteASPtrType
typename detail::DecoratedType< DataT, AS >::type * ConcreteASPtrType
Definition: accessor.hpp:1820
cl::sycl::accessor::operator[]
detail::enable_if_t< Dims==1 &&AccessMode==access::mode::atomic, atomic< DataT, AS > > operator[](size_t Index) const
Definition: accessor.hpp:1591
cl::sycl::detail::__image_array_slice__::__SYCL2020_DEPRECATED
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:756
cl::sycl::accessor< DataT, 1, access_mode::read_write, target::host_buffer, access::placeholder::false_t >::reference
DataT & reference
Definition: accessor.hpp:952
cl::sycl::detail::__image_array_slice__::get_range
range< Dims > get_range() const
Definition: accessor.hpp:764
cl::sycl::access::placeholder
placeholder
Definition: access.hpp:43
cl::sycl::detail::type_list
Definition: type_list.hpp:23
cl::sycl::host_accessor::IsSameAsBuffer
static constexpr bool IsSameAsBuffer()
Definition: accessor.hpp:2152
cl::sycl::image::get_range
range< Dimensions > get_range() const
Definition: image.hpp:261
cl::sycl::accessor::accessor
accessor(const accessor< DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder, ext::oneapi::accessor_property_list< NewPropsT... >> &Other)
Definition: accessor.hpp:1502
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::operator!=
bool operator!=(const accessor &Rhs) const
Definition: accessor.hpp:1985
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const property_list &PropertyList={})
Definition: accessor.hpp:1424
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, const property_list &PropertyList={})
Definition: accessor.hpp:1085
cl::sycl::accessor::operator[]
AccessorCommonT::template AccessorSubscript< Dims - 1 > operator[](size_t Index) const
Definition: accessor.hpp:1598
cl::sycl::id< Dims >
cl::sycl::image::size
size_t size() const noexcept
Definition: image.hpp:275
cl::sycl::host_accessor::host_accessor
host_accessor(buffer< DataT, Dimensions, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, const property_list &PropertyList={})
Definition: accessor.hpp:2259
cl::sycl::detail::__image_array_slice__::write
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:734
cl::sycl::accessor< DataT, 1, access_mode::read_write, target::host_buffer, access::placeholder::false_t >::ConstRefType
const DataT & ConstRefType
Definition: accessor.hpp:828
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, const property_list &PropertyList={})
Definition: accessor.hpp:1241
cl::sycl::sampler
Encapsulates a configuration for sampling an image accessor.
Definition: sampler.hpp:65
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::get_size
size_t get_size() const
Definition: accessor.hpp:1921
handler_proxy.hpp
accessor_property_list.hpp
cl::sycl::detail::is_contained
Definition: type_list.hpp:54
cl::sycl::detail::InitializedVal
Definition: common.hpp:205
cl::sycl::errc::feature_not_supported
@ feature_not_supported
cl::sycl::accessor::IsSameAsBuffer
static constexpr bool IsSameAsBuffer()
Definition: accessor.hpp:847
cl::sycl::detail::IsRunTimePropertyListT
typename std::is_same< ext::oneapi::accessor_property_list<>, T > IsRunTimePropertyListT
Definition: accessor.hpp:233
cl::sycl::accessor::get_range
range< Dimensions > get_range() const
Definition: accessor.hpp:1527
cl::sycl::property_list
Objects of the property_list class are containers for the SYCL properties.
Definition: property_list.hpp:25
cl::sycl::multi_ptr
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: atomic.hpp:32
__SYCL_SPECIAL_CLASS
#define __SYCL_SPECIAL_CLASS
Definition: defines.hpp:27
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={})
Definition: accessor.hpp:1253
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::operator[]
RefType operator[](size_t Index) const
Definition: accessor.hpp:1947
cl::sycl::detail::LocalAccessorBaseDevice
Definition: accessor_impl.hpp:60
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::accessor
accessor(cl::sycl::image< Dimensions, AllocatorT > &Image)
Definition: accessor.hpp:2060
operator==
bool operator==(const Slab &Lhs, const Slab &Rhs)
Definition: usm_allocator.cpp:343
cl::sycl::detail::const_if_const_AS
DataT const_if_const_AS
Definition: type_traits.hpp:348
cl::sycl::detail::AccessorBaseHost::impl
AccessorImplPtr impl
Definition: accessor_impl.hpp:164
cl::sycl::buffer
Defines a shared array that can be used by kernels in queues.
Definition: buffer.hpp:46
cl::sycl::detail::write
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
Definition: stream.hpp:110
cl::sycl::detail::LocalAccessorBaseHost
Definition: accessor_impl.hpp:185
id.hpp
cl::sycl::range< Dims >
cl::sycl::accessor< DataT, 1, access_mode::read_write, target::host_buffer, access::placeholder::false_t >::value_type
DataT value_type
Definition: accessor.hpp:951
cl::sycl::host_accessor
Definition: accessor.hpp:2143
cl::sycl::detail::image_accessor::operator!=
bool operator!=(const image_accessor &Rhs) const
Definition: accessor.hpp:569
cl::sycl::accessor::accessor
accessor(buffer< T, 1, AllocatorT > &BufferRef, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={})
Definition: accessor.hpp:1010
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::operator==
bool operator==(const accessor &Rhs) const
Definition: accessor.hpp:1984
cl::sycl::buffer::size
size_t size() const noexcept
Definition: buffer.hpp:273
cl::sycl::access_mode
access::mode access_mode
Definition: access.hpp:63
cl::sycl::accessor::get_pointer
global_ptr< DataT > get_pointer() const
Definition: accessor.hpp:1613
cl::sycl::detail::TryToGetElementType::type
decltype(check(T())) type
Definition: generic_type_traits.hpp:308
image_accessor_util.hpp
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::operator[]
AccessorCommonT::template AccessorSubscript< Dims - 1 > operator[](size_t Index) const
Definition: accessor.hpp:1976
cl::sycl::accessor::getQualifiedPtr
PtrType getQualifiedPtr() const
Definition: accessor.hpp:939
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::const_reference
const DataT & const_reference
Definition: accessor.hpp:1872
cl::sycl::host_accessor::host_accessor
host_accessor(buffer< DataT, Dimensions, AllocatorT > &BufferRef, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const property_list &PropertyList={})
Definition: accessor.hpp:2279
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={})
Definition: accessor.hpp:1371
export.hpp
cl::sycl::detail::image_accessor::image_accessor
image_accessor(image< Dims, AllocatorT > &ImageRef, handler &CommandGroupHandlerRef, int ImageElementSize)
Definition: accessor.hpp:531
cl::sycl::detail::image_accessor::image_accessor
image_accessor(image< Dims, AllocatorT > &ImageRef, int ImageElementSize)
Definition: accessor.hpp:502
cl::sycl::detail::LocalAccessorBaseDevice::MemRange
range< Dims > MemRange
Definition: accessor_impl.hpp:68
cl::sycl::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
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::accessor
accessor(range< Dimensions > AllocationSize, handler &, const property_list &propList)
Definition: accessor.hpp:1908
cl::sycl::host_accessor::host_accessor
host_accessor(buffer< DataT, Dimensions, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const property_list &PropertyList={})
Definition: accessor.hpp:2297
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::image, IsPlaceholder >::accessor
accessor(cl::sycl::image< Dimensions, AllocatorT > &Image, handler &CommandGroupHandler, const property_list &propList)
Definition: accessor.hpp:2014
cl::sycl::detail::__image_array_slice__::read
DataT read(const CoordT &Coords, const sampler &Smpl) const
Definition: accessor.hpp:727
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={})
Definition: accessor.hpp:1302
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::PtrType
detail::const_if_const_AS< AS, DataT > * PtrType
Definition: accessor.hpp:1823
cl::sycl::device
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:34
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::reference
DataT & reference
Definition: accessor.hpp:1871
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::operator[]
detail::enable_if_t< Dims==1 &&AccessMode==access::mode::atomic, atomic< DataT, AS > > operator[](size_t Index) const
Definition: accessor.hpp:1970
cl::sycl::detail::IsValidCoordDataT
Definition: accessor.hpp:352
generic_type_traits.hpp
cl::sycl::detail::IsCxPropertyList
Definition: accessor.hpp:235
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, const property_list &PropertyList={})
Definition: accessor.hpp:1162
cl::sycl::detail::image_accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::reference
DataT & reference
Definition: accessor.hpp:486
cl::sycl::accessor< DataT, 1, access_mode::read_write, target::host_buffer, access::placeholder::false_t >::PtrType
detail::const_if_const_AS< AS, DataT > * PtrType
Definition: accessor.hpp:829
cl::sycl::detail::__image_array_slice__
Definition: accessor.hpp:370
cl::sycl::detail::accessor_common::AccessorSubscript::operator[]
ConstRefType operator[](size_t Index) const
Definition: accessor.hpp:345
cl::sycl::detail::accessor_common::AccessorSubscript::AccessorSubscript
AccessorSubscript(AccType Accessor, id< Dims > IDs)
Definition: accessor.hpp:312
cl::sycl::detail::TargetToAS
Definition: access.hpp:134
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::accessor
accessor(handler &, const property_list &propList)
Definition: accessor.hpp:1885
cl::sycl::accessor
Buffer accessor.
Definition: accessor.hpp:225
cl::sycl::access::target
target
Definition: access.hpp:17
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={})
Definition: accessor.hpp:1112
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::accessor< DataT, 1, access_mode::read_write, target::host_buffer, access::placeholder::false_t >::AccessorSubscript
typename AccessorCommonT::template AccessorSubscript< Dims > AccessorSubscript
Definition: accessor.hpp:823
cl::sycl::accessor::operator!=
bool operator!=(const accessor &Rhs) const
Definition: accessor.hpp:1627
cl::sycl::accessor::getAdjustedMode
static access::mode getAdjustedMode(const PropertyListT &PropertyList)
Definition: accessor.hpp:851
cl::sycl::detail::accessor_common
Definition: accessor.hpp:267
cl::sycl::accessor::get_offset
id< Dimensions > get_offset() const
Definition: accessor.hpp:1532
cl::sycl::detail::IsPropertyListT
typename std::is_base_of< PropertyListBase, T > IsPropertyListT
Definition: accessor.hpp:229
cl::sycl::detail::image_accessor::get_range
range< Dims > get_range() const
Definition: accessor.hpp:607
cl::sycl::host_accessor::host_accessor
host_accessor(buffer< T, 1, AllocatorT > &BufferRef, const property_list &PropertyList={})
Definition: accessor.hpp:2201
cl::sycl::detail::image_accessor::read
DataT read(const CoordT &Coords, const sampler &Smpl) const
Definition: accessor.hpp:642
cl::sycl::handler
Command group handler class.
Definition: handler.hpp:361
cl::sycl::ext::oneapi::accessor_property_list
Objects of the accessor_property_list class are containers for the SYCL properties.
Definition: property_list.hpp:18
cl::sycl::image
Defines a shared image data.
Definition: image_impl.hpp:29
cl::sycl::accessor::get_pointer
DataT * get_pointer() const
Definition: accessor.hpp:1605
image_ocl_types.hpp
accessor_impl.hpp
cl::sycl::detail::addHostAccessorAndWait
void addHostAccessorAndWait(Requirement *Req)
Definition: accessor_impl.cpp:35
cl::sycl::access::address_space
address_space
Definition: access.hpp:45
cl::sycl::detail::accessor_common< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::RefType
detail::const_if_const_AS< AS, DataT > & RefType
Definition: accessor.hpp:296
cl::sycl::accessor::__SYCL2020_DEPRECATED
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:1522
image.hpp
cl::sycl::detail::accessor_common::AccessorSubscript::AccessorSubscript
AccessorSubscript(AccType Accessor, size_t Index)
Definition: accessor.hpp:317
cl::sycl::detail::image_accessor::__SYCL2020_DEPRECATED
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:602
std::get
constexpr tuple_element< I, tuple< Types... > >::type & get(cl::sycl::detail::tuple< Types... > &Arg) noexcept
Definition: tuple.hpp:199
cl::sycl::device::get_info
info::param_traits< info::device, param >::return_type get_info() const
Queries this SYCL device for information requested by the template parameter param.
Definition: device.cpp:147
cl::sycl::accessor::get_size
size_t get_size() const
Definition: accessor.hpp:1520
cl::sycl::detail::image_accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::value_type
DataT value_type
Definition: accessor.hpp:485
cl::sycl::detail::image_accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::const_reference
const DataT & const_reference
Definition: accessor.hpp:487
cl::sycl::accessor::getLinearIndex
size_t getLinearIndex(id< Dims > Id) const
Definition: accessor.hpp:831
PI_INVALID_VALUE
@ PI_INVALID_VALUE
Definition: pi.h:87
atomic.hpp
cl::sycl::host_accessor::host_accessor
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, const property_list &PropertyList={})
Definition: accessor.hpp:2207
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::accessor
accessor(cl::sycl::image< Dimensions, AllocatorT > &Image, const property_list &propList)
Definition: accessor.hpp:2066
cl::sycl::detail::getSyclObjImpl
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: common.hpp:182
cl::sycl::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:710
cl::sycl::vec
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
Definition: aliases.hpp:18
cl::sycl::detail::image_accessor::operator==
bool operator==(const image_accessor &Rhs) const
Definition: accessor.hpp:561
cl::sycl::cl_int
std::int32_t cl_int
Definition: aliases.hpp:82
cl::sycl::accessor::size
size_t size() const noexcept
Definition: accessor.hpp:1524
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::value_type
DataT value_type
Definition: accessor.hpp:1870
std::hash< cl::sycl::accessor< DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:2367
exception.hpp
cl::sycl::detail::accessor_common::AccessorSubscript::operator[]
detail::enable_if_t< CurDims==1 &&IsAccessAtomic, atomic< DataT, AS > > operator[](size_t Index) const
Definition: accessor.hpp:338
cl::sycl::accessor< DataT, 1, access_mode::read_write, target::host_buffer, access::placeholder::false_t >::RefType
detail::const_if_const_AS< AS, DataT > & RefType
Definition: accessor.hpp:827
std
Definition: accessor.hpp:2358
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::get_range
range< Dims > get_range() const
Definition: accessor.hpp:1928
__SYCL_UNROLL
#define __SYCL_UNROLL(x)
Definition: defines_elementary.hpp:120
cl::sycl::detail::image_accessor
Definition: accessor.hpp:375
cl::sycl::detail::convertToArrayOfN
static T< NewDim > convertToArrayOfN(T< OldDim > OldObj)
Definition: accessor.hpp:252
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::getQualifiedPtr
PtrType getQualifiedPtr() const
Definition: accessor.hpp:1855
cl::sycl::detail::accessor_common::AccType
accessor< DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder, PropertyListT > AccType
Definition: accessor.hpp:301
sampler.hpp
cl::sycl::detail::getDeviceFromHandler
device getDeviceFromHandler(handler &CommandGroupHandlerRef)
Definition: accessor.cpp:15
cl::sycl::detail::accessor_common< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::PtrType
detail::const_if_const_AS< AS, DataT > * PtrType
Definition: accessor.hpp:298
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::accessor
accessor(handler &)
Definition: accessor.hpp:1875
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::accessor
accessor(range< Dimensions > AllocationSize, handler &)
Definition: accessor.hpp:1897
cl::sycl::mode_target_tag_t
Definition: access.hpp:69
cl::sycl::accessor< DataT, 1, access_mode::read_write, target::host_buffer, access::placeholder::false_t >::const_reference
const DataT & const_reference
Definition: accessor.hpp:953
buffer.hpp
cl::sycl::detail::AccessorBaseHost
Definition: accessor_impl.hpp:134
cl::sycl::host_accessor::host_accessor
host_accessor()
Definition: accessor.hpp:2173
common.hpp
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::AccessorSubscript
typename AccessorCommonT::template AccessorSubscript< Dims > AccessorSubscript
Definition: accessor.hpp:1818
cl::sycl::image_channel_type
image_channel_type
Definition: image.hpp:44
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::__SYCL2020_DEPRECATED
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:1923
cl::sycl::detail::__image_array_slice__::size
size_t size() const noexcept
Definition: accessor.hpp:758
cl::sycl::access::mode
mode
Definition: access.hpp:28
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::get_pointer
local_ptr< DataT > get_pointer() const
Definition: accessor.hpp:1980
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::image_array, IsPlaceholder >::accessor
accessor(cl::sycl::image< Dimensions+1, AllocatorT > &Image, handler &CommandGroupHandler)
Definition: accessor.hpp:2108
cl::sycl::stream
A buffered output stream that allows outputting the values of built-in, vector and SYCL types to the ...
Definition: stream.hpp:742
cl::sycl::accessor< DataT, 1, access_mode::read_write, target::host_buffer, access::placeholder::false_t >::ConcreteASPtrType
typename detail::DecoratedType< DataT, AS >::type * ConcreteASPtrType
Definition: accessor.hpp:825
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::RefType
detail::const_if_const_AS< AS, DataT > & RefType
Definition: accessor.hpp:1822
cl::sycl::detail::image_accessor::size
size_t size() const noexcept
Definition: accessor.hpp:604
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::image_array, IsPlaceholder >::operator[]
detail::__image_array_slice__< DataT, Dimensions, AccessMode, IsPlaceholder > operator[](size_t Index) const
Definition: accessor.hpp:2135
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::host_accessor::host_accessor
host_accessor(buffer< DataT, Dimensions, AllocatorT > &BufferRef, range< Dimensions > AccessRange, const property_list &PropertyList={})
Definition: accessor.hpp:2241
cl::sycl::detail::accessor_common::AccessorSubscript::operator[]
RefType operator[](size_t Index) const
Definition: accessor.hpp:330
pointers.hpp
cl::sycl::detail::image_accessor::write
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:667
cl::sycl::accessor::operator==
bool operator==(const accessor &Rhs) const
Definition: accessor.hpp:1626
cl::sycl::accessor::is_placeholder
constexpr bool is_placeholder() const
Definition: accessor.hpp:1518
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, const property_list &PropertyList={})
Definition: accessor.hpp:1290
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const property_list &PropertyList={})
Definition: accessor.hpp:1343
cl::sycl::image_channel_order
image_channel_order
Definition: image.hpp:26
cl::sycl::accessor::accessor
accessor(buffer< T, 1, AllocatorT > &BufferRef, handler &CommandGroupHandler, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={})
Definition: accessor.hpp:1059
accessor_properties.hpp
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::getLinearIndex
size_t getLinearIndex(id< Dims > Id) const
Definition: accessor.hpp:1862
cl::sycl::atomic
Definition: atomic.hpp:171
cl::sycl::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:2165
cl::sycl::errc::accessor
@ accessor
cl::sycl::detail::DecoratedType
Definition: access.hpp:157
spirv_types.hpp
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={})
Definition: accessor.hpp:1451
property_list_conversion.hpp
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::size
size_t size() const noexcept
Definition: accessor.hpp:1925