DPC++ Runtime
Runtime libraries for oneAPI DPC++
accessor.hpp
Go to the documentation of this file.
1 //==------------ accessor.hpp - SYCL standard header file ------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
12 #include <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 esimd {
210 namespace detail {
211 // Forward declare a "back-door" access class to support ESIMD.
212 class AccessorPrivateProxy;
213 } // namespace detail
214 } // namespace esimd
215 } // namespace intel
216 } // namespace ext
217 
218 template <typename DataT, int Dimensions = 1,
219  access::mode AccessMode = access::mode::read_write,
220  access::target AccessTarget = access::target::device,
221  access::placeholder IsPlaceholder = access::placeholder::false_t,
222  typename PropertyListT = ext::oneapi::accessor_property_list<>>
223 class accessor;
224 
225 namespace detail {
226 void __SYCL_EXPORT constructorNotification(void *BufferObj, void *AccessorObj,
227  access::target Target,
228  access::mode Mode,
229  const code_location &CodeLoc);
230 template <typename T>
231 using IsPropertyListT = typename std::is_base_of<PropertyListBase, T>;
232 
233 template <typename T>
235  typename std::is_same<ext::oneapi::accessor_property_list<>, T>;
236 
237 template <typename T> struct IsCxPropertyList {
238  constexpr static bool value = false;
239 };
240 
241 template <typename... Props>
242 struct IsCxPropertyList<ext::oneapi::accessor_property_list<Props...>> {
243  constexpr static bool value = true;
244 };
245 
246 template <> struct IsCxPropertyList<ext::oneapi::accessor_property_list<>> {
247  constexpr static bool value = false;
248 };
249 
250 // The function extends or truncates number of dimensions of objects of id
251 // or ranges classes. When extending the new values are filled with
252 // DefaultValue, truncation just removes extra values.
253 template <int NewDim, int DefaultValue, template <int> class T, int OldDim>
254 static T<NewDim> convertToArrayOfN(T<OldDim> OldObj) {
255  T<NewDim> NewObj = InitializedVal<NewDim, T>::template get<0>();
256  const int CopyDims = NewDim > OldDim ? OldDim : NewDim;
257  for (int I = 0; I < CopyDims; ++I)
258  NewObj[I] = OldObj[I];
259  for (int I = CopyDims; I < NewDim; ++I)
260  NewObj[I] = DefaultValue;
261  return NewObj;
262 }
263 
264 __SYCL_EXPORT device getDeviceFromHandler(handler &CommandGroupHandlerRef);
265 
266 template <typename DataT, int Dimensions, access::mode AccessMode,
267  access::target AccessTarget, access::placeholder IsPlaceholder,
268  typename PropertyListT = ext::oneapi::accessor_property_list<>>
270 protected:
271  constexpr static bool IsPlaceH = IsPlaceholder == access::placeholder::true_t;
273 
274  constexpr static bool IsHostBuf = AccessTarget == access::target::host_buffer;
275 
276  // TODO: SYCL 2020 deprecates four of the target enum values
277  // and replaces them with 2 (device and host_task). May want
278  // to change these constexpr.
279  constexpr static bool IsGlobalBuf =
280  AccessTarget == access::target::global_buffer;
281 
282  constexpr static bool IsConstantBuf =
283  AccessTarget == access::target::constant_buffer;
284 
285  constexpr static bool IsAccessAnyWrite =
286  AccessMode == access::mode::write ||
287  AccessMode == access::mode::read_write ||
288  AccessMode == access::mode::discard_write ||
289  AccessMode == access::mode::discard_read_write;
290 
291  constexpr static bool IsAccessReadOnly = AccessMode == access::mode::read;
292 
293  constexpr static bool IsAccessReadWrite =
294  AccessMode == access::mode::read_write;
295 
296  constexpr static bool IsAccessAtomic = AccessMode == access::mode::atomic;
297 
299  using ConstRefType = const DataT &;
301 
302  using AccType = accessor<DataT, Dimensions, AccessMode, AccessTarget,
303  IsPlaceholder, PropertyListT>;
304 
305  // The class which allows to access value of N dimensional accessor using N
306  // subscript operators, e.g. accessor[2][2][3]
307  template <int SubDims> class AccessorSubscript {
308  static constexpr int Dims = Dimensions;
309 
310  mutable id<Dims> MIDs;
311  AccType MAccessor;
312 
313  public:
315  : MAccessor(Accessor), MIDs(IDs) {}
316 
317  // Only accessor class is supposed to use this c'tor for the first
318  // operator[].
319  AccessorSubscript(AccType Accessor, size_t Index) : MAccessor(Accessor) {
320  MIDs[0] = Index;
321  }
322 
323  template <int CurDims = SubDims>
324  typename detail::enable_if_t<(CurDims > 1), AccessorSubscript<CurDims - 1>>
325  operator[](size_t Index) {
326  MIDs[Dims - CurDims] = Index;
327  return AccessorSubscript<CurDims - 1>(MAccessor, MIDs);
328  }
329 
330  template <int CurDims = SubDims,
332  RefType operator[](size_t Index) const {
333  MIDs[Dims - CurDims] = Index;
334  return MAccessor[MIDs];
335  }
336 
337  template <int CurDims = SubDims>
338  typename detail::enable_if_t<CurDims == 1 && IsAccessAtomic,
339  atomic<DataT, AS>>
340  operator[](size_t Index) const {
341  MIDs[Dims - CurDims] = Index;
342  return MAccessor[MIDs];
343  }
344 
345  template <int CurDims = SubDims,
347  ConstRefType operator[](size_t Index) const {
348  MIDs[Dims - SubDims] = Index;
349  return MAccessor[MIDs];
350  }
351  };
352 };
353 
354 template <int Dim, typename T> struct IsValidCoordDataT;
355 template <typename T> struct IsValidCoordDataT<1, T> {
356  constexpr static bool value =
358 };
359 template <typename T> struct IsValidCoordDataT<2, T> {
360  constexpr static bool value =
363 };
364 template <typename T> struct IsValidCoordDataT<3, T> {
365  constexpr static bool value =
368 };
369 
370 template <typename DataT, int Dimensions, access::mode AccessMode,
371  access::placeholder IsPlaceholder>
373 
374 // Image accessor
375 template <typename DataT, int Dimensions, access::mode AccessMode,
376  access::target AccessTarget, access::placeholder IsPlaceholder>
378 #ifndef __SYCL_DEVICE_ONLY__
379  : public detail::AccessorBaseHost {
380  size_t MImageCount;
381  image_channel_order MImgChannelOrder;
382  image_channel_type MImgChannelType;
383 #else
384 {
385 
386  using OCLImageTy = typename detail::opencl_image_type<Dimensions, AccessMode,
387  AccessTarget>::type;
388  OCLImageTy MImageObj;
389  char MPadding[sizeof(detail::AccessorBaseHost) +
390  sizeof(size_t /*MImageCount*/) + sizeof(image_channel_order) +
391  sizeof(image_channel_type) - sizeof(OCLImageTy)];
392 
393 protected:
394  void imageAccessorInit(OCLImageTy Image) { MImageObj = Image; }
395 
396 private:
397 #endif
398  template <typename T1, int T2, access::mode T3, access::placeholder T4>
399  friend class __image_array_slice__;
400 
401  constexpr static bool IsHostImageAcc =
402  (AccessTarget == access::target::host_image);
403 
404  constexpr static bool IsImageAcc = (AccessTarget == access::target::image);
405 
406  constexpr static bool IsImageArrayAcc =
407  (AccessTarget == access::target::image_array);
408 
409  constexpr static bool IsImageAccessWriteOnly =
410  (AccessMode == access::mode::write ||
411  AccessMode == access::mode::discard_write);
412 
413  constexpr static bool IsImageAccessAnyWrite =
414  (IsImageAccessWriteOnly || AccessMode == access::mode::read_write);
415 
416  constexpr static bool IsImageAccessReadOnly =
417  (AccessMode == access::mode::read);
418 
419  constexpr static bool IsImageAccessAnyRead =
420  (IsImageAccessReadOnly || AccessMode == access::mode::read_write);
421 
422  static_assert(std::is_same<DataT, cl_int4>::value ||
423  std::is_same<DataT, cl_uint4>::value ||
424  std::is_same<DataT, cl_float4>::value ||
425  std::is_same<DataT, cl_half4>::value,
426  "The data type of an image accessor must be only cl_int4, "
427  "cl_uint4, cl_float4 or cl_half4 from SYCL namespace");
428 
429  static_assert(IsImageAcc || IsHostImageAcc || IsImageArrayAcc,
430  "Expected image type");
431 
432  static_assert(IsPlaceholder == access::placeholder::false_t,
433  "Expected false as Placeholder value for image accessor.");
434 
435  static_assert(
436  ((IsImageAcc || IsImageArrayAcc) &&
437  (IsImageAccessWriteOnly || IsImageAccessReadOnly)) ||
438  (IsHostImageAcc && (IsImageAccessAnyWrite || IsImageAccessAnyRead)),
439  "Access modes can be only read/write/discard_write for image/image_array "
440  "target accessor, or they can be only "
441  "read/write/discard_write/read_write for host_image target accessor.");
442 
443  static_assert(Dimensions > 0 && Dimensions <= 3,
444  "Dimensions can be 1/2/3 for image accessor.");
445 
446  template <info::device param>
447  void checkDeviceFeatureSupported(const device &Device) {
448  if (!Device.get_info<param>())
449  throw feature_not_supported("Images are not supported by this device.",
451  }
452 
453 #ifdef __SYCL_DEVICE_ONLY__
454 
455  sycl::vec<int, Dimensions> getRangeInternal() const {
456  return __invoke_ImageQuerySize<sycl::vec<int, Dimensions>, OCLImageTy>(
457  MImageObj);
458  }
459 
460  size_t getElementSize() const {
461  int ChannelType = __invoke_ImageQueryFormat<int, OCLImageTy>(MImageObj);
462  int ChannelOrder = __invoke_ImageQueryOrder<int, OCLImageTy>(MImageObj);
463  int ElementSize = getSPIRVElementSize(ChannelType, ChannelOrder);
464  return ElementSize;
465  }
466 
467 #else
468 
469  sycl::vec<int, Dimensions> getRangeInternal() const {
470  // TODO: Implement for host.
471  throw runtime_error("image::getRangeInternal() is not implemented for host",
473  return sycl::vec<int, Dimensions>{1};
474  }
475 
476 #endif
477 
478 private:
479  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
480 
481 #ifdef __SYCL_DEVICE_ONLY__
482  const OCLImageTy getNativeImageObj() const { return MImageObj; }
483 #endif // __SYCL_DEVICE_ONLY__
484 
485 public:
486  using value_type = DataT;
487  using reference = DataT &;
488  using const_reference = const DataT &;
489 
490  // image_accessor Constructors.
491 
492 #ifdef __SYCL_DEVICE_ONLY__
493  // Default constructor for objects later initialized with __init member.
494  image_accessor() : MImageObj() {}
495 #endif
496 
497  // Available only when: accessTarget == access::target::host_image
498  // template <typename AllocatorT>
499  // accessor(image<dimensions, AllocatorT> &imageRef);
500  template <
501  typename AllocatorT, int Dims = Dimensions,
502  typename = detail::enable_if_t<(Dims > 0 && Dims <= 3) && IsHostImageAcc>>
503  image_accessor(image<Dims, AllocatorT> &ImageRef, int ImageElementSize)
504 #ifdef __SYCL_DEVICE_ONLY__
505  {
506  (void)ImageRef;
507  (void)ImageElementSize;
508  // No implementation needed for device. The constructor is only called by
509  // host.
510  }
511 #else
512  : AccessorBaseHost({detail::getSyclObjImpl(ImageRef)->getRowPitch(),
513  detail::getSyclObjImpl(ImageRef)->getSlicePitch(), 0},
514  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
515  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
516  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
517  Dimensions, ImageElementSize),
518  MImageCount(ImageRef.size()),
519  MImgChannelOrder(detail::getSyclObjImpl(ImageRef)->getChannelOrder()),
520  MImgChannelType(detail::getSyclObjImpl(ImageRef)->getChannelType()) {
521  addHostAccessorAndWait(AccessorBaseHost::impl.get());
522  }
523 #endif
524 
525  // Available only when: accessTarget == access::target::image
526  // template <typename AllocatorT>
527  // accessor(image<dimensions, AllocatorT> &imageRef,
528  // handler &commandGroupHandlerRef);
529  template <
530  typename AllocatorT, int Dims = Dimensions,
531  typename = detail::enable_if_t<(Dims > 0 && Dims <= 3) && IsImageAcc>>
533  handler &CommandGroupHandlerRef, int ImageElementSize)
534 #ifdef __SYCL_DEVICE_ONLY__
535  {
536  (void)ImageRef;
537  (void)CommandGroupHandlerRef;
538  (void)ImageElementSize;
539  // No implementation needed for device. The constructor is only called by
540  // host.
541  }
542 #else
543  : AccessorBaseHost({detail::getSyclObjImpl(ImageRef)->getRowPitch(),
544  detail::getSyclObjImpl(ImageRef)->getSlicePitch(), 0},
545  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
546  detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
547  AccessMode, detail::getSyclObjImpl(ImageRef).get(),
548  Dimensions, ImageElementSize),
549  MImageCount(ImageRef.size()),
550  MImgChannelOrder(detail::getSyclObjImpl(ImageRef)->getChannelOrder()),
551  MImgChannelType(detail::getSyclObjImpl(ImageRef)->getChannelType()) {
552  checkDeviceFeatureSupported<info::device::image_support>(
553  getDeviceFromHandler(CommandGroupHandlerRef));
554  }
555 #endif
556 
557  /* -- common interface members -- */
558 
559  // operator == and != need to be defined only for host application as per the
560  // SYCL spec 1.2.1
561 #ifndef __SYCL_DEVICE_ONLY__
562  bool operator==(const image_accessor &Rhs) const { return Rhs.impl == impl; }
563 #else
564  // The operator with __SYCL_DEVICE_ONLY__ need to be declared for compilation
565  // of host application with device compiler.
566  // Usage of this operator inside the kernel code will give a runtime failure.
567  bool operator==(const image_accessor &Rhs) const;
568 #endif
569 
570  bool operator!=(const image_accessor &Rhs) const { return !(Rhs == *this); }
571 
572  // get_count() method : Returns the number of elements of the SYCL image this
573  // SYCL accessor is accessing.
574  //
575  // get_range() method : Returns a range object which represents the number of
576  // elements of dataT per dimension that this accessor may access.
577  // The range object returned must equal to the range of the image this
578  // accessor is associated with.
579 
580 #ifdef __SYCL_DEVICE_ONLY__
581 
582  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
583  size_t get_count() const { return size(); }
584  size_t size() const noexcept { return get_range<Dimensions>().size(); }
585 
586  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 1>>
587  range<1> get_range() const {
588  cl_int Range = getRangeInternal();
589  return range<1>(Range);
590  }
591  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 2>>
592  range<2> get_range() const {
593  cl_int2 Range = getRangeInternal();
594  return range<2>(Range[0], Range[1]);
595  }
596  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 3>>
597  range<3> get_range() const {
598  cl_int3 Range = getRangeInternal();
599  return range<3>(Range[0], Range[1], Range[2]);
600  }
601 
602 #else
603  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
604  size_t get_count() const { return size(); };
605  size_t size() const noexcept { return MImageCount; };
606 
607  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
609  return detail::convertToArrayOfN<Dims, 1>(getAccessRange());
610  }
611 
612 #endif
613 
614  // Available only when:
615  // (accessTarget == access::target::image && accessMode == access::mode::read)
616  // || (accessTarget == access::target::host_image && ( accessMode ==
617  // access::mode::read || accessMode == access::mode::read_write))
618  template <typename CoordT, int Dims = Dimensions,
619  typename = detail::enable_if_t<
620  (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
622  ((IsImageAcc && IsImageAccessReadOnly) ||
623  (IsHostImageAcc && IsImageAccessAnyRead))>>
624  DataT read(const CoordT &Coords) const {
625 #ifdef __SYCL_DEVICE_ONLY__
626  return __invoke__ImageRead<DataT, OCLImageTy, CoordT>(MImageObj, Coords);
627 #else
628  sampler Smpl(coordinate_normalization_mode::unnormalized,
629  addressing_mode::none, filtering_mode::nearest);
630  return read<CoordT, Dims>(Coords, Smpl);
631 #endif
632  }
633 
634  // Available only when:
635  // (accessTarget == access::target::image && accessMode == access::mode::read)
636  // || (accessTarget == access::target::host_image && ( accessMode ==
637  // access::mode::read || accessMode == access::mode::read_write))
638  template <typename CoordT, int Dims = Dimensions,
639  typename = detail::enable_if_t<
640  (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
641  ((IsImageAcc && IsImageAccessReadOnly) ||
642  (IsHostImageAcc && IsImageAccessAnyRead))>>
643  DataT read(const CoordT &Coords, const sampler &Smpl) const {
644 #ifdef __SYCL_DEVICE_ONLY__
645  return __invoke__ImageReadSampler<DataT, OCLImageTy, CoordT>(
646  MImageObj, Coords, Smpl.impl.m_Sampler);
647 #else
648  return imageReadSamplerHostImpl<CoordT, DataT>(
649  Coords, Smpl, getAccessRange() /*Image Range*/,
650  getOffset() /*Image Pitch*/, MImgChannelType, MImgChannelOrder,
651  AccessorBaseHost::getPtr() /*ptr to image*/,
652  AccessorBaseHost::getElemSize());
653 #endif
654  }
655 
656  // Available only when:
657  // (accessTarget == access::target::image && (accessMode ==
658  // access::mode::write || accessMode == access::mode::discard_write)) ||
659  // (accessTarget == access::target::host_image && (accessMode ==
660  // access::mode::write || accessMode == access::mode::discard_write ||
661  // accessMode == access::mode::read_write))
662  template <typename CoordT, int Dims = Dimensions,
663  typename = detail::enable_if_t<
664  (Dims > 0) && (detail::is_genint<CoordT>::value) &&
666  ((IsImageAcc && IsImageAccessWriteOnly) ||
667  (IsHostImageAcc && IsImageAccessAnyWrite))>>
668  void write(const CoordT &Coords, const DataT &Color) const {
669 #ifdef __SYCL_DEVICE_ONLY__
670  __invoke__ImageWrite<OCLImageTy, CoordT, DataT>(MImageObj, Coords, Color);
671 #else
672  imageWriteHostImpl(Coords, Color, getOffset() /*ImagePitch*/,
673  AccessorBaseHost::getElemSize(), MImgChannelType,
674  MImgChannelOrder,
675  AccessorBaseHost::getPtr() /*Ptr to Image*/);
676 #endif
677  }
678 };
679 
680 template <typename DataT, int Dimensions, access::mode AccessMode,
681  access::placeholder IsPlaceholder>
682 class __image_array_slice__ {
683 
684  static_assert(Dimensions < 3,
685  "Image slice cannot have more then 2 dimensions");
686 
687  constexpr static int AdjustedDims = (Dimensions == 2) ? 4 : Dimensions + 1;
688 
689  template <typename CoordT,
690  typename CoordElemType =
693  getAdjustedCoords(const CoordT &Coords) const {
694  CoordElemType LastCoord = 0;
695 
696  if (std::is_same<float, CoordElemType>::value) {
697  sycl::vec<int, Dimensions + 1> Size = MBaseAcc.getRangeInternal();
698  LastCoord =
699  MIdx / static_cast<float>(Size.template swizzle<Dimensions>());
700  } else {
701  LastCoord = MIdx;
702  }
703 
704  sycl::vec<CoordElemType, Dimensions> LeftoverCoords{LastCoord};
705  sycl::vec<CoordElemType, AdjustedDims> AdjustedCoords{Coords,
706  LeftoverCoords};
707  return AdjustedCoords;
708  }
709 
710 public:
712  accessor<DataT, Dimensions, AccessMode, access::target::image_array,
714  BaseAcc,
715  size_t Idx)
716  : MBaseAcc(BaseAcc), MIdx(Idx) {}
717 
718  template <typename CoordT, int Dims = Dimensions,
719  typename = detail::enable_if_t<
721  DataT read(const CoordT &Coords) const {
722  return MBaseAcc.read(getAdjustedCoords(Coords));
723  }
724 
725  template <typename CoordT, int Dims = Dimensions,
726  typename = detail::enable_if_t<
728  DataT read(const CoordT &Coords, const sampler &Smpl) const {
729  return MBaseAcc.read(getAdjustedCoords(Coords), Smpl);
730  }
731 
732  template <typename CoordT, int Dims = Dimensions,
733  typename = detail::enable_if_t<
735  void write(const CoordT &Coords, const DataT &Color) const {
736  return MBaseAcc.write(getAdjustedCoords(Coords), Color);
737  }
738 
739 #ifdef __SYCL_DEVICE_ONLY__
740  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
741  size_t get_count() const { return size(); }
742  size_t size() const noexcept { return get_range<Dimensions>().size(); }
743 
744  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 1>>
745  range<1> get_range() const {
746  cl_int2 Count = MBaseAcc.getRangeInternal();
747  return range<1>(Count.x());
748  }
749  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 2>>
750  range<2> get_range() const {
751  cl_int3 Count = MBaseAcc.getRangeInternal();
752  return range<2>(Count.x(), Count.y());
753  }
754 
755 #else
756 
757  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
758  size_t get_count() const { return size(); }
759  size_t size() const noexcept {
760  return MBaseAcc.MImageCount / MBaseAcc.getAccessRange()[Dimensions];
761  }
762 
763  template <int Dims = Dimensions,
764  typename = detail::enable_if_t<(Dims == 1 || Dims == 2)>>
766  return detail::convertToArrayOfN<Dims, 1>(MBaseAcc.getAccessRange());
767  }
768 
769 #endif
770 
771 private:
772  size_t MIdx;
773  accessor<DataT, Dimensions, AccessMode, access::target::image_array,
775  MBaseAcc;
776 };
777 
778 } // namespace detail
779 
785 template <typename DataT, int Dimensions, access::mode AccessMode,
786  access::target AccessTarget, access::placeholder IsPlaceholder,
787  typename PropertyListT>
788 class __SYCL_SPECIAL_CLASS accessor :
789 #ifndef __SYCL_DEVICE_ONLY__
790  public detail::AccessorBaseHost,
791 #endif
792  public detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
793  IsPlaceholder, PropertyListT> {
794 protected:
795  static_assert((AccessTarget == access::target::global_buffer ||
796  AccessTarget == access::target::constant_buffer ||
797  AccessTarget == access::target::host_buffer),
798  "Expected buffer type");
799 
800  static_assert((AccessTarget == access::target::global_buffer ||
801  AccessTarget == access::target::host_buffer) ||
802  (AccessTarget == access::target::constant_buffer &&
803  AccessMode == access::mode::read),
804  "Access mode can be only read for constant buffers");
805 
806  static_assert(detail::IsPropertyListT<PropertyListT>::value,
807  "PropertyListT must be accessor_property_list");
808 
809  using AccessorCommonT =
810  detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
811  IsPlaceholder, PropertyListT>;
812 
813  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
814 
815  using AccessorCommonT::AS;
816  using AccessorCommonT::IsAccessAnyWrite;
817  using AccessorCommonT::IsAccessReadOnly;
818  using AccessorCommonT::IsConstantBuf;
819  using AccessorCommonT::IsGlobalBuf;
820  using AccessorCommonT::IsHostBuf;
821  using AccessorCommonT::IsPlaceH;
822  template <int Dims>
823  using AccessorSubscript =
824  typename AccessorCommonT::template AccessorSubscript<Dims>;
825 
827 
829  using ConstRefType = const DataT &;
831 
832  template <int Dims = Dimensions> size_t getLinearIndex(id<Dims> Id) const {
833 
834 #ifdef __SYCL_DEVICE_ONLY__
835  // Pointer is already adjusted for 1D case.
836  if (Dimensions == 1)
837  return Id[0];
838 #endif // __SYCL_DEVICE_ONLY__
839 
840  size_t Result = 0;
841  // Unroll the following loop for both host and device code
842  __SYCL_UNROLL(3)
843  for (int I = 0; I < Dims; ++I) {
844  Result = Result * getMemoryRange()[I] + Id[I];
845 #if __cplusplus >= 201703L
846  if constexpr (!(PropertyListT::template has_property<
848  Result += getOffset()[I];
849  }
850 #else
851  Result += getOffset()[I];
852 #endif
853  }
854  return Result;
855  }
856 
857  template <typename T, int Dims> static constexpr bool IsSameAsBuffer() {
858  return std::is_same<T, DataT>::value && (Dims > 0) && (Dims == Dimensions);
859  }
860 
861  static access::mode getAdjustedMode(const PropertyListT &PropertyList) {
862  access::mode AdjustedMode = AccessMode;
863 
864  if (PropertyList.template has_property<property::no_init>() ||
865  PropertyList.template has_property<property::noinit>()) {
866  if (AdjustedMode == access::mode::write) {
867  AdjustedMode = access::mode::discard_write;
868  } else if (AdjustedMode == access::mode::read_write) {
869  AdjustedMode = access::mode::discard_read_write;
870  }
871  }
872 
873  return AdjustedMode;
874  }
875 
876 #if __cplusplus >= 201703L
877 
878  template <typename TagT> static constexpr bool IsValidTag() {
879  return std::is_same<TagT, mode_tag_t<AccessMode>>::value ||
880  std::is_same<TagT,
882  }
883 
884 #endif
885 
886 #ifdef __SYCL_DEVICE_ONLY__
887 
888  id<AdjustedDim> &getOffset() { return impl.Offset; }
889  range<AdjustedDim> &getAccessRange() { return impl.AccessRange; }
890  range<AdjustedDim> &getMemoryRange() { return impl.MemRange; }
891 
892  const id<AdjustedDim> &getOffset() const { return impl.Offset; }
893  const range<AdjustedDim> &getAccessRange() const { return impl.AccessRange; }
894  const range<AdjustedDim> &getMemoryRange() const { return impl.MemRange; }
895 
896  detail::AccessorImplDevice<AdjustedDim> impl;
897 
898  union {
899  ConcreteASPtrType MData;
900  };
901 
902  // TODO replace usages with getQualifiedPtr
903  const ConcreteASPtrType getNativeImageObj() const { return MData; }
904 
905  void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
906  range<AdjustedDim> MemRange, id<AdjustedDim> Offset) {
907  MData = Ptr;
908 #pragma unroll
909  for (int I = 0; I < AdjustedDim; ++I) {
910 #if __cplusplus >= 201703L
911  if constexpr (!(PropertyListT::template has_property<
912  sycl::ext::oneapi::property::no_offset>())) {
913  getOffset()[I] = Offset[I];
914  }
915 #else
916  getOffset()[I] = Offset[I];
917 #endif
918  getAccessRange()[I] = AccessRange[I];
919  getMemoryRange()[I] = MemRange[I];
920  }
921  // In case of 1D buffer, adjust pointer during initialization rather
922  // then each time in operator[]. Will have to re-adjust in get_pointer
923  if (1 == AdjustedDim)
924 #if __cplusplus >= 201703L
925  if constexpr (!(PropertyListT::template has_property<
926  sycl::ext::oneapi::property::no_offset>())) {
927  MData += Offset[0];
928  }
929 #else
930  MData += Offset[0];
931 #endif
932  }
933 
934  // __init variant used by the device compiler for ESIMD kernels.
935  // TODO In ESIMD accessors usage is limited for now - access range, mem
936  // range and offset are not supported.
937  void __init_esimd(ConcreteASPtrType Ptr) { MData = Ptr; }
938 
939  ConcreteASPtrType getQualifiedPtr() const { return MData; }
940 
941  template <typename, int, access::mode, access::target, access::placeholder,
942  typename>
943  friend class accessor;
944 
945 #ifndef __SYCL_DEVICE_ONLY__
946  using AccessorBaseHost::impl;
947 #endif
948 
949 public:
950  // Default constructor for objects later initialized with __init member.
951  accessor()
952  : impl({}, detail::InitializedVal<AdjustedDim, range>::template get<0>(),
953  detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
954 
955 #else
956  using AccessorBaseHost::getAccessRange;
957  using AccessorBaseHost::getMemoryRange;
958  using AccessorBaseHost::getOffset;
959 
960  char padding[sizeof(detail::AccessorImplDevice<AdjustedDim>) +
961  sizeof(PtrType) - sizeof(detail::AccessorBaseHost)];
962 
964  return reinterpret_cast<PtrType>(AccessorBaseHost::getPtr());
965  }
966 
967 #endif // __SYCL_DEVICE_ONLY__
968 
969 private:
970  friend class sycl::stream;
971  friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
972 
973 public:
974  using value_type = DataT;
975  using reference = DataT &;
976  using const_reference = const DataT &;
977 
978  // The list of accessor constructors with their arguments
979  // -------+---------+-------+----+-----+--------------
980  // Dimensions = 0
981  // -------+---------+-------+----+-----+--------------
982  // buffer | | | | | property_list
983  // buffer | handler | | | | property_list
984  // -------+---------+-------+----+-----+--------------
985  // Dimensions >= 1
986  // -------+---------+-------+----+-----+--------------
987  // buffer | | | | | property_list
988  // buffer | | | | tag | property_list
989  // buffer | handler | | | | property_list
990  // buffer | handler | | | tag | property_list
991  // buffer | | range | | | property_list
992  // buffer | | range | | tag | property_list
993  // buffer | handler | range | | | property_list
994  // buffer | handler | range | | tag | property_list
995  // buffer | | range | id | | property_list
996  // buffer | | range | id | tag | property_list
997  // buffer | handler | range | id | | property_list
998  // buffer | handler | range | id | tag | property_list
999  // -------+---------+-------+----+-----+--------------
1000 
1001 public:
1002  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1003  typename detail::enable_if_t<
1005  std::is_same<T, DataT>::value && Dims == 0 &&
1006  ((!IsPlaceH && IsHostBuf) ||
1007  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> * = nullptr>
1009  buffer<T, 1, AllocatorT> &BufferRef,
1010  const property_list &PropertyList = {},
1011  const detail::code_location CodeLoc = detail::code_location::current())
1012 #ifdef __SYCL_DEVICE_ONLY__
1013  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1014  (void)PropertyList;
1015 #else
1016  : AccessorBaseHost(
1017  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1018  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1019  getAdjustedMode(PropertyList),
1020  detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
1021  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1022  preScreenAccessor(BufferRef.size(), PropertyList);
1023  if (!IsPlaceH)
1024  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1026  detail::AccessorBaseHost::impl.get(),
1027  AccessTarget, AccessMode, CodeLoc);
1028 #endif
1029  }
1030 
1031  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1032  typename... PropTypes,
1033  typename detail::enable_if_t<
1034  detail::IsCxPropertyList<PropertyListT>::value &&
1035  std::is_same<T, DataT>::value && Dims == 0 &&
1036  ((!IsPlaceH && IsHostBuf) ||
1037  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> * = nullptr>
1039  buffer<T, 1, AllocatorT> &BufferRef,
1041  {},
1042  const detail::code_location CodeLoc = detail::code_location::current())
1043 #ifdef __SYCL_DEVICE_ONLY__
1044  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1045  (void)PropertyList;
1046 #else
1047  : AccessorBaseHost(
1048  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1049  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1050  getAdjustedMode(PropertyList),
1051  detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
1052  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1053  preScreenAccessor(BufferRef.size(), PropertyList);
1054  if (!IsPlaceH)
1055  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1057  detail::AccessorBaseHost::impl.get(),
1058  AccessTarget, AccessMode, CodeLoc);
1059 #endif
1060  }
1061 
1062  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1063  typename = typename detail::enable_if_t<
1064  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1065  std::is_same<T, DataT>::value && (Dims == 0) &&
1066  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1068  buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1069  const property_list &PropertyList = {},
1070  const detail::code_location CodeLoc = detail::code_location::current())
1071 #ifdef __SYCL_DEVICE_ONLY__
1072  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1073  (void)CommandGroupHandler;
1074  (void)PropertyList;
1075  }
1076 #else
1077  : AccessorBaseHost(
1078  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1079  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1080  getAdjustedMode(PropertyList),
1081  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1082  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1083  preScreenAccessor(BufferRef.size(), PropertyList);
1084  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1086  detail::AccessorBaseHost::impl.get(),
1087  AccessTarget, AccessMode, CodeLoc);
1088  }
1089 #endif
1090 
1091  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1092  typename... PropTypes,
1093  typename = typename detail::enable_if_t<
1094  detail::IsCxPropertyList<PropertyListT>::value &&
1095  std::is_same<T, DataT>::value && (Dims == 0) &&
1096  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1098  buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1100  {},
1101  const detail::code_location CodeLoc = detail::code_location::current())
1102 #ifdef __SYCL_DEVICE_ONLY__
1103  : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1104  (void)CommandGroupHandler;
1105  (void)PropertyList;
1106  }
1107 #else
1108  : AccessorBaseHost(
1109  /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1110  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1111  getAdjustedMode(PropertyList),
1112  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1113  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1114  preScreenAccessor(BufferRef.size(), PropertyList);
1115  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1117  detail::AccessorBaseHost::impl.get(),
1118  AccessTarget, AccessMode, CodeLoc);
1119  }
1120 #endif
1121 
1122  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1123  typename = detail::enable_if_t<
1124  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1125  IsSameAsBuffer<T, Dims>() &&
1126  ((!IsPlaceH && IsHostBuf) ||
1127  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1129  buffer<T, Dims, AllocatorT> &BufferRef,
1130  const property_list &PropertyList = {},
1131  const detail::code_location CodeLoc = detail::code_location::current())
1132 #ifdef __SYCL_DEVICE_ONLY__
1133  : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1134  (void)PropertyList;
1135  }
1136 #else
1137  : AccessorBaseHost(
1138  /*Offset=*/{0, 0, 0},
1139  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1140  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1141  getAdjustedMode(PropertyList),
1142  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1143  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1144  preScreenAccessor(BufferRef.size(), PropertyList);
1145  if (!IsPlaceH)
1146  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1148  detail::AccessorBaseHost::impl.get(),
1149  AccessTarget, AccessMode, CodeLoc);
1150  }
1151 #endif
1152 
1153  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1154  typename... PropTypes,
1155  typename = detail::enable_if_t<
1156  detail::IsCxPropertyList<PropertyListT>::value &&
1157  IsSameAsBuffer<T, Dims>() &&
1158  ((!IsPlaceH && IsHostBuf) ||
1159  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1161  buffer<T, Dims, AllocatorT> &BufferRef,
1163  {},
1164  const detail::code_location CodeLoc = detail::code_location::current())
1165 #ifdef __SYCL_DEVICE_ONLY__
1166  : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
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  preScreenAccessor(BufferRef.size(), PropertyList);
1178  if (!IsPlaceH)
1179  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1181  detail::AccessorBaseHost::impl.get(),
1182  AccessTarget, AccessMode, CodeLoc);
1183  }
1184 #endif
1185 
1186 #if __cplusplus >= 201703L
1187 
1188  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1189  typename TagT,
1190  typename = detail::enable_if_t<
1191  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1192  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1193  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1194  accessor(
1195  buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1196  const property_list &PropertyList = {},
1197  const detail::code_location CodeLoc = detail::code_location::current())
1198  : accessor(BufferRef, PropertyList, CodeLoc) {
1199  adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1200  }
1201 
1202  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1203  typename TagT, typename... PropTypes,
1204  typename = detail::enable_if_t<
1205  detail::IsCxPropertyList<PropertyListT>::value &&
1206  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1207  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1208  accessor(
1209  buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1210  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1211  {},
1212  const detail::code_location CodeLoc = detail::code_location::current())
1213  : accessor(BufferRef, PropertyList, CodeLoc) {
1214  adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1215  }
1216 #endif
1217 
1218  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1219  typename = detail::enable_if_t<
1220  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1221  IsSameAsBuffer<T, Dims>() &&
1222  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1224  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1225  const property_list &PropertyList = {},
1226  const detail::code_location CodeLoc = detail::code_location::current())
1227 #ifdef __SYCL_DEVICE_ONLY__
1228  : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1229  (void)CommandGroupHandler;
1230  (void)PropertyList;
1231  }
1232 #else
1233  : AccessorBaseHost(
1234  /*Offset=*/{0, 0, 0},
1235  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1236  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1237  getAdjustedMode(PropertyList),
1238  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1239  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1240  preScreenAccessor(BufferRef.size(), PropertyList);
1241  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1243  detail::AccessorBaseHost::impl.get(),
1244  AccessTarget, AccessMode, CodeLoc);
1245  }
1246 #endif
1247 
1248  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1249  typename... PropTypes,
1250  typename = detail::enable_if_t<
1251  detail::IsCxPropertyList<PropertyListT>::value &&
1252  IsSameAsBuffer<T, Dims>() &&
1253  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1255  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1257  {},
1258  const detail::code_location CodeLoc = detail::code_location::current())
1259 #ifdef __SYCL_DEVICE_ONLY__
1260  : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1261  (void)CommandGroupHandler;
1262  (void)PropertyList;
1263  }
1264 #else
1265  : AccessorBaseHost(
1266  /*Offset=*/{0, 0, 0},
1267  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1268  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1269  getAdjustedMode(PropertyList),
1270  detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
1271  BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1272  preScreenAccessor(BufferRef.size(), PropertyList);
1273  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1275  detail::AccessorBaseHost::impl.get(),
1276  AccessTarget, AccessMode, CodeLoc);
1277  }
1278 #endif
1279 
1280 #if __cplusplus >= 201703L
1281 
1282  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1283  typename TagT,
1284  typename = detail::enable_if_t<
1285  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1286  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1287  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1288  accessor(
1289  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1290  TagT, const property_list &PropertyList = {},
1291  const detail::code_location CodeLoc = detail::code_location::current())
1292  : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1293  adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1294  }
1295 
1296  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1297  typename TagT, typename... PropTypes,
1298  typename = detail::enable_if_t<
1299  detail::IsCxPropertyList<PropertyListT>::value &&
1300  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1301  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1302  accessor(
1303  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1304  TagT,
1305  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1306  {},
1307  const detail::code_location CodeLoc = detail::code_location::current())
1308  : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1309  adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1310  }
1311 
1312 #endif
1313 
1314  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1315  typename = detail::enable_if_t<
1316  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1317  IsSameAsBuffer<T, Dims>() &&
1318  ((!IsPlaceH && IsHostBuf) ||
1319  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1321  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1322  const property_list &PropertyList = {},
1323  const detail::code_location CodeLoc = detail::code_location::current())
1324  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1325 
1326  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1327  typename... PropTypes,
1328  typename = detail::enable_if_t<
1329  detail::IsCxPropertyList<PropertyListT>::value &&
1330  IsSameAsBuffer<T, Dims>() &&
1331  ((!IsPlaceH && IsHostBuf) ||
1332  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1334  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1336  {},
1337  const detail::code_location CodeLoc = detail::code_location::current())
1338  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1339 
1340 #if __cplusplus >= 201703L
1341 
1342  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1343  typename TagT,
1344  typename = detail::enable_if_t<
1345  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1346  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1347  (IsGlobalBuf || IsConstantBuf)>>
1348  accessor(
1349  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1350  TagT, const property_list &PropertyList = {},
1351  const detail::code_location CodeLoc = detail::code_location::current())
1352  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1353  adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1354  }
1355 
1356  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1357  typename TagT, typename... PropTypes,
1358  typename = detail::enable_if_t<
1359  detail::IsCxPropertyList<PropertyListT>::value &&
1360  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1361  (IsGlobalBuf || IsConstantBuf)>>
1362  accessor(
1363  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1364  TagT,
1365  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1366  {},
1367  const detail::code_location CodeLoc = detail::code_location::current())
1368  : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1369  adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1370  }
1371 #endif
1372 
1373  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1374  typename = detail::enable_if_t<
1375  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1376  IsSameAsBuffer<T, Dims>() &&
1377  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1379  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1380  range<Dimensions> AccessRange, const property_list &PropertyList = {},
1381  const detail::code_location CodeLoc = detail::code_location::current())
1382  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1383  CodeLoc) {}
1384 
1385  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1386  typename... PropTypes,
1387  typename = detail::enable_if_t<
1388  detail::IsCxPropertyList<PropertyListT>::value &&
1389  IsSameAsBuffer<T, Dims>() &&
1390  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1392  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1393  range<Dimensions> AccessRange,
1395  {},
1396  const detail::code_location CodeLoc = detail::code_location::current())
1397  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1398  CodeLoc) {}
1399 
1400 #if __cplusplus >= 201703L
1401 
1402  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1403  typename TagT,
1404  typename = detail::enable_if_t<
1405  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1406  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1407  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1408  accessor(
1409  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1410  range<Dimensions> AccessRange, TagT,
1411  const property_list &PropertyList = {},
1412  const detail::code_location CodeLoc = detail::code_location::current())
1413  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1414  CodeLoc) {
1415  adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1416  }
1417 
1418  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1419  typename TagT, typename... PropTypes,
1420  typename = detail::enable_if_t<
1421  detail::IsCxPropertyList<PropertyListT>::value &&
1422  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1423  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1424  accessor(
1425  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1426  range<Dimensions> AccessRange, TagT,
1427  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1428  {},
1429  const detail::code_location CodeLoc = detail::code_location::current())
1430  : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1431  CodeLoc) {
1432  adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1433  }
1434 #endif
1435 
1436  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1437  typename = detail::enable_if_t<
1438  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1439  IsSameAsBuffer<T, Dims>() &&
1440  ((!IsPlaceH && IsHostBuf) ||
1441  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1443  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1444  id<Dimensions> AccessOffset, const property_list &PropertyList = {},
1445  const detail::code_location CodeLoc = detail::code_location::current())
1446 #ifdef __SYCL_DEVICE_ONLY__
1447  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1448  (void)PropertyList;
1449  }
1450 #else
1451  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1452  detail::convertToArrayOfN<3, 1>(AccessRange),
1453  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1454  getAdjustedMode(PropertyList),
1455  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1456  sizeof(DataT), BufferRef.OffsetInBytes,
1457  BufferRef.IsSubBuffer) {
1458  preScreenAccessor(BufferRef.size(), PropertyList);
1459  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1460  BufferRef.get_range()))
1461  throw sycl::invalid_object_error(
1462  "accessor with requested offset and range would exceed the bounds of "
1463  "the buffer",
1465 
1466  if (!IsPlaceH)
1467  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1469  detail::AccessorBaseHost::impl.get(),
1470  AccessTarget, AccessMode, CodeLoc);
1471  }
1472 #endif
1473 
1474  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1475  typename... PropTypes,
1476  typename = detail::enable_if_t<
1477  detail::IsCxPropertyList<PropertyListT>::value &&
1478  IsSameAsBuffer<T, Dims>() &&
1479  ((!IsPlaceH && IsHostBuf) ||
1480  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1482  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1483  id<Dimensions> AccessOffset,
1485  {},
1486  const detail::code_location CodeLoc = detail::code_location::current())
1487 #ifdef __SYCL_DEVICE_ONLY__
1488  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1489  (void)PropertyList;
1490  }
1491 #else
1492  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1493  detail::convertToArrayOfN<3, 1>(AccessRange),
1494  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1495  getAdjustedMode(PropertyList),
1496  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1497  sizeof(DataT), BufferRef.OffsetInBytes,
1498  BufferRef.IsSubBuffer) {
1499  preScreenAccessor(BufferRef.size(), PropertyList);
1500  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1501  BufferRef.get_range()))
1502  throw sycl::invalid_object_error(
1503  "accessor with requested offset and range would exceed the bounds of "
1504  "the buffer",
1506 
1507  if (!IsPlaceH)
1508  addHostAccessorAndWait(AccessorBaseHost::impl.get());
1510  detail::AccessorBaseHost::impl.get(),
1511  AccessTarget, AccessMode, CodeLoc);
1512  }
1513 #endif
1514 
1515 #if __cplusplus >= 201703L
1516 
1517  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1518  typename TagT,
1519  typename = detail::enable_if_t<
1520  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1521  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1522  (IsGlobalBuf || IsConstantBuf)>>
1523  accessor(
1524  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1525  id<Dimensions> AccessOffset, TagT, const property_list &PropertyList = {},
1526  const detail::code_location CodeLoc = detail::code_location::current())
1527  : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
1528  adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1529  }
1530 
1531  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1532  typename TagT, typename... PropTypes,
1533  typename = detail::enable_if_t<
1534  detail::IsCxPropertyList<PropertyListT>::value &&
1535  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1536  (IsGlobalBuf || IsConstantBuf)>>
1537  accessor(
1538  buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1539  id<Dimensions> AccessOffset, TagT,
1540  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1541  {},
1542  const detail::code_location CodeLoc = detail::code_location::current())
1543  : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
1544  adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1545  }
1546 #endif
1547 
1548  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1549  typename = detail::enable_if_t<
1550  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1551  IsSameAsBuffer<T, Dims>() &&
1552  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1554  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1555  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
1556  const property_list &PropertyList = {},
1557  const detail::code_location CodeLoc = detail::code_location::current())
1558 #ifdef __SYCL_DEVICE_ONLY__
1559  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1560  (void)CommandGroupHandler;
1561  (void)PropertyList;
1562  }
1563 #else
1564  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1565  detail::convertToArrayOfN<3, 1>(AccessRange),
1566  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1567  getAdjustedMode(PropertyList),
1568  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1569  sizeof(DataT), BufferRef.OffsetInBytes,
1570  BufferRef.IsSubBuffer) {
1571  preScreenAccessor(BufferRef.size(), PropertyList);
1572  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1573  BufferRef.get_range()))
1574  throw sycl::invalid_object_error(
1575  "accessor with requested offset and range would exceed the bounds of "
1576  "the buffer",
1578 
1579  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1581  detail::AccessorBaseHost::impl.get(),
1582  AccessTarget, AccessMode, CodeLoc);
1583  }
1584 #endif
1585 
1586  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1587  typename... PropTypes,
1588  typename = detail::enable_if_t<
1589  detail::IsCxPropertyList<PropertyListT>::value &&
1590  IsSameAsBuffer<T, Dims>() &&
1591  (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1593  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1594  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
1596  {},
1597  const detail::code_location CodeLoc = detail::code_location::current())
1598 #ifdef __SYCL_DEVICE_ONLY__
1599  : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1600  (void)CommandGroupHandler;
1601  (void)PropertyList;
1602  }
1603 #else
1604  : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1605  detail::convertToArrayOfN<3, 1>(AccessRange),
1606  detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1607  getAdjustedMode(PropertyList),
1608  detail::getSyclObjImpl(BufferRef).get(), Dimensions,
1609  sizeof(DataT), BufferRef.OffsetInBytes,
1610  BufferRef.IsSubBuffer) {
1611  preScreenAccessor(BufferRef.size(), PropertyList);
1612  if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1613  BufferRef.get_range()))
1614  throw sycl::invalid_object_error(
1615  "accessor with requested offset and range would exceed the bounds of "
1616  "the buffer",
1618 
1619  detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
1621  detail::AccessorBaseHost::impl.get(),
1622  AccessTarget, AccessMode, CodeLoc);
1623  }
1624 #endif
1625 
1626 #if __cplusplus >= 201703L
1627 
1628  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1629  typename TagT,
1630  typename = detail::enable_if_t<
1631  detail::IsRunTimePropertyListT<PropertyListT>::value &&
1632  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1633  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1634  accessor(
1635  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1636  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
1637  const property_list &PropertyList = {},
1638  const detail::code_location CodeLoc = detail::code_location::current())
1639  : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1640  PropertyList, CodeLoc) {
1641  adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1642  }
1643 
1644  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1645  typename TagT, typename... PropTypes,
1646  typename = detail::enable_if_t<
1647  detail::IsCxPropertyList<PropertyListT>::value &&
1648  IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1649  (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1650  accessor(
1651  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1652  range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
1653  const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1654  {},
1655  const detail::code_location CodeLoc = detail::code_location::current())
1656  : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1657  PropertyList, CodeLoc) {
1658  adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1659  }
1660 #endif
1661 
1662  template <typename... NewPropsT>
1664  const accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
1666  const detail::code_location CodeLoc = detail::code_location::current())
1667 #ifdef __SYCL_DEVICE_ONLY__
1668  : impl(Other.impl)
1669 #else
1670  : detail::AccessorBaseHost(Other)
1671 #endif
1672  {
1674  "Conversion is only available for accessor_property_list");
1675  static_assert(
1676  PropertyListT::template areSameCompileTimeProperties<NewPropsT...>(),
1677  "Compile-time-constant properties must be the same");
1678 #ifndef __SYCL_DEVICE_ONLY__
1679  detail::constructorNotification(impl.get()->MSYCLMemObj, impl.get(),
1680  AccessTarget, AccessMode, CodeLoc);
1681 #endif
1682  }
1683 
1684  constexpr bool is_placeholder() const { return IsPlaceH; }
1685 
1686  size_t get_size() const { return getAccessRange().size() * sizeof(DataT); }
1687 
1688  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1689  size_t get_count() const { return size(); }
1690  size_t size() const noexcept { return getAccessRange().size(); }
1691 
1692  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
1694  return detail::convertToArrayOfN<Dimensions, 1>(getAccessRange());
1695  }
1696 
1697  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
1699 #if __cplusplus >= 201703L
1700  static_assert(
1701  !(PropertyListT::template has_property<
1703  "Accessor has no_offset property, get_offset() can not be used");
1704 #endif
1705  return detail::convertToArrayOfN<Dimensions, 0>(getOffset());
1706  }
1707 
1708  template <int Dims = Dimensions, typename RefT = RefType,
1709  typename = detail::enable_if_t<Dims == 0 && IsAccessAnyWrite &&
1710  !std::is_const<RefT>::value>>
1711  operator RefType() const {
1712  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
1713  return *(getQualifiedPtr() + LinearIndex);
1714  }
1715 
1716  template <int Dims = Dimensions,
1718  operator ConstRefType() const {
1719  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
1720  return *(getQualifiedPtr() + LinearIndex);
1721  }
1722 
1723  template <int Dims = Dimensions,
1724  typename = detail::enable_if_t<(Dims > 0) && IsAccessAnyWrite>>
1725  RefType operator[](id<Dimensions> Index) const {
1726  const size_t LinearIndex = getLinearIndex(Index);
1727  return getQualifiedPtr()[LinearIndex];
1728  }
1729 
1730  template <int Dims = Dimensions>
1731  typename detail::enable_if_t<(Dims > 0) && IsAccessReadOnly, ConstRefType>
1732  operator[](id<Dimensions> Index) const {
1733  const size_t LinearIndex = getLinearIndex(Index);
1734  return getQualifiedPtr()[LinearIndex];
1735  }
1736 
1737  template <int Dims = Dimensions>
1738  operator typename detail::enable_if_t<Dims == 0 &&
1739  AccessMode == access::mode::atomic,
1740 #ifdef __ENABLE_USM_ADDR_SPACE__
1741  atomic<DataT>
1742 #else
1743  atomic<DataT, AS>
1744 #endif
1745  >() const {
1746  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
1747  return atomic<DataT, AS>(
1748  multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
1749  }
1750 
1751  template <int Dims = Dimensions>
1752  typename detail::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
1753  atomic<DataT, AS>>
1754  operator[](id<Dimensions> Index) const {
1755  const size_t LinearIndex = getLinearIndex(Index);
1756  return atomic<DataT, AS>(
1757  multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
1758  }
1759 
1760  template <int Dims = Dimensions>
1761  typename detail::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
1762  atomic<DataT, AS>>
1763  operator[](size_t Index) const {
1764  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
1765  return atomic<DataT, AS>(
1766  multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
1767  }
1768  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>
1769  typename AccessorCommonT::template AccessorSubscript<Dims - 1>
1770  operator[](size_t Index) const {
1771  return AccessorSubscript<Dims - 1>(*this, Index);
1772  }
1773 
1774  template <access::target AccessTarget_ = AccessTarget,
1775  typename = detail::enable_if_t<AccessTarget_ ==
1776  access::target::host_buffer>>
1777  DataT *get_pointer() const {
1778  return getPointerAdjusted();
1779  }
1780 
1781  template <
1782  access::target AccessTarget_ = AccessTarget,
1785  return global_ptr<DataT>(getPointerAdjusted());
1786  }
1787 
1788  template <access::target AccessTarget_ = AccessTarget,
1789  typename = detail::enable_if_t<AccessTarget_ ==
1790  access::target::constant_buffer>>
1792  return constant_ptr<DataT>(getPointerAdjusted());
1793  }
1794 
1795  bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
1796  bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
1797 
1798 private:
1799  // supporting function for get_pointer()
1800  // when dim==1, MData will have been preadjusted for faster access with []
1801  // but for get_pointer() we must return the original pointer.
1802  // On device, getQualifiedPtr() returns MData, so we need to backjust it.
1803  // On host, getQualifiedPtr() does not return MData, no need to adjust.
1804  PtrType getPointerAdjusted() const {
1805 #ifdef __SYCL_DEVICE_ONLY__
1806  if (1 == AdjustedDim)
1807  return getQualifiedPtr() - impl.Offset[0];
1808 #endif
1809  return getQualifiedPtr();
1810  }
1811 
1812  void preScreenAccessor(const size_t elemInBuffer,
1813  const PropertyListT &PropertyList) {
1814  // check device accessor buffer size
1815  if (!IsHostBuf && elemInBuffer == 0)
1816  throw sycl::invalid_object_error(
1817  "SYCL buffer size is zero. To create a device accessor, SYCL "
1818  "buffer size must be greater than zero.",
1820 
1821  // check that no_init property is compatible with access mode
1822  if (PropertyList.template has_property<property::no_init>() &&
1823  AccessMode == access::mode::read) {
1824  throw sycl::invalid_object_error(
1825  "accessor would cannot be both read_only and no_init",
1827  }
1828  }
1829 
1830 #if __cplusplus >= 201703L
1831  template <typename... PropTypes>
1832  void adjustAccPropsInBuf(detail::SYCLMemObjI *SYCLMemObject) {
1833  if constexpr (PropertyListT::template has_property<
1834  sycl::ext::intel::property::buffer_location>()) {
1835  auto location = (PropertyListT::template get_property<
1836  sycl::ext::intel::property::buffer_location>())
1837  .get_location();
1838  property_list PropList{
1839  sycl::property::buffer::detail::buffer_location(location)};
1840  detail::SYCLMemObjT *SYCLMemObjectT =
1841  dynamic_cast<detail::SYCLMemObjT *>(SYCLMemObject);
1842  SYCLMemObjectT->addOrReplaceAccessorProperties(PropList);
1843  } else {
1844  deleteAccPropsFromBuf(SYCLMemObject);
1845  }
1846  }
1847 
1848  void deleteAccPropsFromBuf(detail::SYCLMemObjI *SYCLMemObject) {
1849  detail::SYCLMemObjT *SYCLMemObjectT =
1850  dynamic_cast<detail::SYCLMemObjT *>(SYCLMemObject);
1851  SYCLMemObjectT->deleteAccessorProperty(
1853  }
1854 #endif
1855 };
1856 
1857 #if __cplusplus >= 201703L
1858 
1859 template <typename DataT, int Dimensions, typename AllocatorT>
1860 accessor(buffer<DataT, Dimensions, AllocatorT>)
1861  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
1862  access::placeholder::true_t>;
1863 
1864 template <typename DataT, int Dimensions, typename AllocatorT,
1865  typename... PropsT>
1866 accessor(buffer<DataT, Dimensions, AllocatorT>,
1867  const ext::oneapi::accessor_property_list<PropsT...> &)
1868  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
1869  access::placeholder::true_t,
1870  ext::oneapi::accessor_property_list<PropsT...>>;
1871 
1872 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
1873 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1)
1874  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
1875  detail::deduceAccessTarget<Type1, Type1>(target::device),
1876  access::placeholder::true_t>;
1877 
1878 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1879  typename... PropsT>
1880 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1,
1881  const ext::oneapi::accessor_property_list<PropsT...> &)
1882  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
1883  detail::deduceAccessTarget<Type1, Type1>(target::device),
1884  access::placeholder::true_t,
1885  ext::oneapi::accessor_property_list<PropsT...>>;
1886 
1887 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1888  typename Type2>
1889 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2)
1890  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
1891  detail::deduceAccessTarget<Type1, Type2>(target::device),
1892  access::placeholder::true_t>;
1893 
1894 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1895  typename Type2, typename... PropsT>
1896 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2,
1897  const ext::oneapi::accessor_property_list<PropsT...> &)
1898  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
1899  detail::deduceAccessTarget<Type1, Type2>(target::device),
1900  access::placeholder::true_t,
1901  ext::oneapi::accessor_property_list<PropsT...>>;
1902 
1903 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1904  typename Type2, typename Type3>
1905 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3)
1906  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
1907  detail::deduceAccessTarget<Type2, Type3>(target::device),
1908  access::placeholder::true_t>;
1909 
1910 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1911  typename Type2, typename Type3, typename... PropsT>
1912 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3,
1913  const ext::oneapi::accessor_property_list<PropsT...> &)
1914  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
1915  detail::deduceAccessTarget<Type2, Type3>(target::device),
1916  access::placeholder::true_t,
1917  ext::oneapi::accessor_property_list<PropsT...>>;
1918 
1919 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1920  typename Type2, typename Type3, typename Type4>
1921 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
1922  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
1923  detail::deduceAccessTarget<Type3, Type4>(target::device),
1924  access::placeholder::true_t>;
1925 
1926 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1927  typename Type2, typename Type3, typename Type4, typename... PropsT>
1928 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
1929  const ext::oneapi::accessor_property_list<PropsT...> &)
1930  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
1931  detail::deduceAccessTarget<Type3, Type4>(target::device),
1932  access::placeholder::true_t,
1933  ext::oneapi::accessor_property_list<PropsT...>>;
1934 
1935 template <typename DataT, int Dimensions, typename AllocatorT>
1936 accessor(buffer<DataT, Dimensions, AllocatorT>, handler)
1937  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
1938  access::placeholder::false_t>;
1939 
1940 template <typename DataT, int Dimensions, typename AllocatorT,
1941  typename... PropsT>
1942 accessor(buffer<DataT, Dimensions, AllocatorT>, handler,
1943  const ext::oneapi::accessor_property_list<PropsT...> &)
1944  -> accessor<DataT, Dimensions, access::mode::read_write, target::device,
1945  access::placeholder::false_t,
1946  ext::oneapi::accessor_property_list<PropsT...>>;
1947 
1948 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
1949 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1)
1950  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
1951  detail::deduceAccessTarget<Type1, Type1>(target::device),
1952  access::placeholder::false_t>;
1953 
1954 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1955  typename... PropsT>
1956 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1,
1957  const ext::oneapi::accessor_property_list<PropsT...> &)
1958  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
1959  detail::deduceAccessTarget<Type1, Type1>(target::device),
1960  access::placeholder::false_t,
1961  ext::oneapi::accessor_property_list<PropsT...>>;
1962 
1963 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1964  typename Type2>
1965 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2)
1966  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
1967  detail::deduceAccessTarget<Type1, Type2>(target::device),
1968  access::placeholder::false_t>;
1969 
1970 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1971  typename Type2, typename... PropsT>
1972 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2,
1973  const ext::oneapi::accessor_property_list<PropsT...> &)
1974  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
1975  detail::deduceAccessTarget<Type1, Type2>(target::device),
1976  access::placeholder::false_t,
1977  ext::oneapi::accessor_property_list<PropsT...>>;
1978 
1979 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1980  typename Type2, typename Type3>
1981 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3)
1982  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
1983  detail::deduceAccessTarget<Type2, Type3>(target::device),
1984  access::placeholder::false_t>;
1985 
1986 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1987  typename Type2, typename Type3, typename... PropsT>
1988 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
1989  const ext::oneapi::accessor_property_list<PropsT...> &)
1990  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
1991  detail::deduceAccessTarget<Type2, Type3>(target::device),
1992  access::placeholder::false_t,
1993  ext::oneapi::accessor_property_list<PropsT...>>;
1994 
1995 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1996  typename Type2, typename Type3, typename Type4>
1997 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
1998  Type4)
1999  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2000  detail::deduceAccessTarget<Type3, Type4>(target::device),
2001  access::placeholder::false_t>;
2002 
2003 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2004  typename Type2, typename Type3, typename Type4, typename... PropsT>
2005 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
2006  Type4, const ext::oneapi::accessor_property_list<PropsT...> &)
2007  -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2008  detail::deduceAccessTarget<Type3, Type4>(target::device),
2009  access::placeholder::false_t,
2010  ext::oneapi::accessor_property_list<PropsT...>>;
2011 #endif
2012 
2016 template <typename DataT, int Dimensions, access::mode AccessMode,
2017  access::placeholder IsPlaceholder>
2018 class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
2019  access::target::local, IsPlaceholder> :
2020 #ifndef __SYCL_DEVICE_ONLY__
2022 #endif
2023  public detail::accessor_common<DataT, Dimensions, AccessMode,
2024  access::target::local, IsPlaceholder> {
2025 protected:
2026  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
2027 
2028  using AccessorCommonT =
2029  detail::accessor_common<DataT, Dimensions, AccessMode,
2030  access::target::local, IsPlaceholder>;
2031 
2032  using AccessorCommonT::AS;
2033  using AccessorCommonT::IsAccessAnyWrite;
2034  template <int Dims>
2035  using AccessorSubscript =
2036  typename AccessorCommonT::template AccessorSubscript<Dims>;
2037 
2039 
2042 
2043 #ifdef __SYCL_DEVICE_ONLY__
2045 
2046  sycl::range<AdjustedDim> &getSize() { return impl.MemRange; }
2047  const sycl::range<AdjustedDim> &getSize() const { return impl.MemRange; }
2048 
2049  void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
2050  range<AdjustedDim>, id<AdjustedDim>) {
2051  MData = Ptr;
2052 #pragma unroll
2053  for (int I = 0; I < AdjustedDim; ++I)
2054  getSize()[I] = AccessRange[I];
2055  }
2056 
2057 public:
2058  // Default constructor for objects later initialized with __init member.
2059  accessor()
2060  : impl(detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
2061 
2062 protected:
2063  ConcreteASPtrType getQualifiedPtr() const { return MData; }
2064 
2065  ConcreteASPtrType MData;
2066 
2067 #else
2068 
2069  char padding[sizeof(detail::LocalAccessorBaseDevice<AdjustedDim>) +
2071  using detail::LocalAccessorBaseHost::getSize;
2072 
2074  return reinterpret_cast<PtrType>(LocalAccessorBaseHost::getPtr());
2075  }
2076 
2077 #endif // __SYCL_DEVICE_ONLY__
2078 
2079  // Method which calculates linear offset for the ID using Range and Offset.
2080  template <int Dims = AdjustedDim> size_t getLinearIndex(id<Dims> Id) const {
2081  size_t Result = 0;
2082  for (int I = 0; I < Dims; ++I)
2083  Result = Result * getSize()[I] + Id[I];
2084  return Result;
2085  }
2086 
2087 public:
2088  using value_type = DataT;
2089  using reference = DataT &;
2090  using const_reference = const DataT &;
2091 
2092  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 0>>
2094  detail::code_location::current())
2095 #ifdef __SYCL_DEVICE_ONLY__
2096  : impl(range<AdjustedDim>{1}){}
2097 #else
2098  : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT)) {
2099  detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
2100  access::target::local, AccessMode, CodeLoc);
2101  }
2102 #endif
2103 
2104  template <int Dims = Dimensions,
2105  typename = detail::enable_if_t<Dims == 0>>
2106  accessor(handler &, const property_list &propList,
2107  const detail::code_location CodeLoc =
2108  detail::code_location::current())
2109 #ifdef __SYCL_DEVICE_ONLY__
2110  : impl(range<AdjustedDim>{1}) {
2111  (void)propList;
2112  }
2113 #else
2114  : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT)) {
2115  (void)propList;
2116  detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
2117  access::target::local, AccessMode, CodeLoc);
2118  }
2119 #endif
2120 
2121  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
2123  range<Dimensions> AllocationSize, handler &,
2124  const detail::code_location CodeLoc = detail::code_location::current())
2125 #ifdef __SYCL_DEVICE_ONLY__
2126  : impl(AllocationSize){}
2127 #else
2128  : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2129  AdjustedDim, sizeof(DataT)) {
2130  detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
2131  access::target::local, AccessMode, CodeLoc);
2132  }
2133 #endif
2134 
2135  template <int Dims = Dimensions,
2136  typename = detail::enable_if_t<(Dims > 0)>>
2138  const property_list &propList,
2139  const detail::code_location CodeLoc =
2140  detail::code_location::current())
2141 #ifdef __SYCL_DEVICE_ONLY__
2142  : impl(AllocationSize) {
2143  (void)propList;
2144  }
2145 #else
2146  : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2147  AdjustedDim, sizeof(DataT)) {
2148  (void)propList;
2149  detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
2150  access::target::local, AccessMode, CodeLoc);
2151  }
2152 #endif
2153 
2154  size_t get_size() const { return getSize().size() * sizeof(DataT); }
2155 
2156  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
2157  size_t get_count() const { return size(); }
2158  size_t size() const noexcept { return getSize().size(); }
2159 
2160  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
2162  return detail::convertToArrayOfN<Dims, 1>(getSize());
2163  }
2164 
2165  template <int Dims = Dimensions,
2167  operator RefType() const {
2168  return *getQualifiedPtr();
2169  }
2170 
2171  template <int Dims = Dimensions,
2172  typename = detail::enable_if_t<(Dims > 0) && IsAccessAnyWrite>>
2173  RefType operator[](id<Dimensions> Index) const {
2174  const size_t LinearIndex = getLinearIndex(Index);
2175  return getQualifiedPtr()[LinearIndex];
2176  }
2177 
2178  template <int Dims = Dimensions,
2180  RefType operator[](size_t Index) const {
2181  return getQualifiedPtr()[Index];
2182  }
2183 
2184  template <int Dims = Dimensions>
2185  operator typename detail::enable_if_t<
2186  Dims == 0 && AccessMode == access::mode::atomic, atomic<DataT, AS>>()
2187  const {
2188  return atomic<DataT, AS>(multi_ptr<DataT, AS>(getQualifiedPtr()));
2189  }
2190 
2191  template <int Dims = Dimensions>
2192  typename detail::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
2193  atomic<DataT, AS>>
2194  operator[](id<Dimensions> Index) const {
2195  const size_t LinearIndex = getLinearIndex(Index);
2196  return atomic<DataT, AS>(
2197  multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
2198  }
2199 
2200  template <int Dims = Dimensions>
2201  typename detail::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
2202  atomic<DataT, AS>>
2203  operator[](size_t Index) const {
2204  return atomic<DataT, AS>(multi_ptr<DataT, AS>(getQualifiedPtr() + Index));
2205  }
2206 
2207  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>
2208  typename AccessorCommonT::template AccessorSubscript<Dims - 1>
2209  operator[](size_t Index) const {
2210  return AccessorSubscript<Dims - 1>(*this, Index);
2211  }
2212 
2214  return local_ptr<DataT>(getQualifiedPtr());
2215  }
2216 
2217  bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
2218  bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
2219 };
2220 
2226 template <typename DataT, int Dimensions, access::mode AccessMode,
2227  access::placeholder IsPlaceholder>
2228 class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
2229  access::target::image, IsPlaceholder>
2230  : public detail::image_accessor<DataT, Dimensions, AccessMode,
2231  access::target::image, IsPlaceholder> {
2232 public:
2233  template <typename AllocatorT>
2235  handler &CommandGroupHandler)
2236  : detail::image_accessor<DataT, Dimensions, AccessMode,
2237  access::target::image, IsPlaceholder>(
2238  Image, CommandGroupHandler,
2239  (detail::getSyclObjImpl(Image))->getElementSize()) {
2240 #ifndef __SYCL_DEVICE_ONLY__
2241  detail::associateWithHandler(CommandGroupHandler, this,
2242  access::target::image);
2243 #endif
2244  }
2245 
2246  template <typename AllocatorT>
2248  handler &CommandGroupHandler, const property_list &propList)
2249  : detail::image_accessor<DataT, Dimensions, AccessMode,
2250  access::target::image, IsPlaceholder>(
2251  Image, CommandGroupHandler,
2252  (detail::getSyclObjImpl(Image))->getElementSize()) {
2253  (void)propList;
2254 #ifndef __SYCL_DEVICE_ONLY__
2255  detail::associateWithHandler(CommandGroupHandler, this,
2256  access::target::image);
2257 #endif
2258  }
2259 #ifdef __SYCL_DEVICE_ONLY__
2260 private:
2261  using OCLImageTy =
2262  typename detail::opencl_image_type<Dimensions, AccessMode,
2263  access::target::image>::type;
2264 
2265  // Front End requires this method to be defined in the accessor class.
2266  // It does not call the base class's init method.
2267  void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
2268 
2269  // __init variant used by the device compiler for ESIMD kernels.
2270  void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
2271 
2272 public:
2273  // Default constructor for objects later initialized with __init member.
2274  accessor() = default;
2275 #endif
2276 };
2277 
2285 template <typename DataT, int Dimensions, access::mode AccessMode,
2286  access::placeholder IsPlaceholder>
2287 class accessor<DataT, Dimensions, AccessMode, access::target::host_image,
2288  IsPlaceholder>
2289  : public detail::image_accessor<DataT, Dimensions, AccessMode,
2290  access::target::host_image, IsPlaceholder> {
2291 public:
2292  template <typename AllocatorT>
2294  : detail::image_accessor<DataT, Dimensions, AccessMode,
2295  access::target::host_image, IsPlaceholder>(
2296  Image, (detail::getSyclObjImpl(Image))->getElementSize()) {}
2297 
2298  template <typename AllocatorT>
2300  const property_list &propList)
2301  : detail::image_accessor<DataT, Dimensions, AccessMode,
2302  access::target::host_image, IsPlaceholder>(
2303  Image, (detail::getSyclObjImpl(Image))->getElementSize()) {
2304  (void)propList;
2305  }
2306 };
2307 
2316 template <typename DataT, int Dimensions, access::mode AccessMode,
2317  access::placeholder IsPlaceholder>
2318 class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
2319  access::target::image_array, IsPlaceholder>
2320  : public detail::image_accessor<DataT, Dimensions + 1, AccessMode,
2321  access::target::image, IsPlaceholder> {
2322 #ifdef __SYCL_DEVICE_ONLY__
2323 private:
2324  using OCLImageTy =
2325  typename detail::opencl_image_type<Dimensions + 1, AccessMode,
2326  access::target::image>::type;
2327 
2328  // Front End requires this method to be defined in the accessor class.
2329  // It does not call the base class's init method.
2330  void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
2331 
2332  // __init variant used by the device compiler for ESIMD kernels.
2333  void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
2334 
2335 public:
2336  // Default constructor for objects later initialized with __init member.
2337  accessor() = default;
2338 #endif
2339 public:
2340  template <typename AllocatorT>
2342  handler &CommandGroupHandler)
2343  : detail::image_accessor<DataT, Dimensions + 1, AccessMode,
2344  access::target::image, IsPlaceholder>(
2345  Image, CommandGroupHandler,
2346  (detail::getSyclObjImpl(Image))->getElementSize()) {
2347 #ifndef __SYCL_DEVICE_ONLY__
2348  detail::associateWithHandler(CommandGroupHandler, this,
2349  access::target::image_array);
2350 #endif
2351  }
2352 
2353  template <typename AllocatorT>
2355  handler &CommandGroupHandler, const property_list &propList)
2356  : detail::image_accessor<DataT, Dimensions + 1, AccessMode,
2357  access::target::image, IsPlaceholder>(
2358  Image, CommandGroupHandler,
2359  (detail::getSyclObjImpl(Image))->getElementSize()) {
2360  (void)propList;
2361 #ifndef __SYCL_DEVICE_ONLY__
2362  detail::associateWithHandler(CommandGroupHandler, this,
2363  access::target::image_array);
2364 #endif
2365  }
2366 
2368  operator[](size_t Index) const {
2369  return detail::__image_array_slice__<DataT, Dimensions, AccessMode,
2370  IsPlaceholder>(*this, Index);
2371  }
2372 };
2373 
2374 template <typename DataT, int Dimensions = 1,
2375  access_mode AccessMode = access_mode::read_write>
2377  : public accessor<DataT, Dimensions, AccessMode, target::host_buffer,
2378  access::placeholder::false_t> {
2379 protected:
2380  using AccessorT = accessor<DataT, Dimensions, AccessMode, target::host_buffer,
2381  access::placeholder::false_t>;
2382 
2383  constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
2384 
2385  template <typename T, int Dims> static constexpr bool IsSameAsBuffer() {
2386  return std::is_same<T, DataT>::value && (Dims > 0) && (Dims == Dimensions);
2387  }
2388 
2389 #if __cplusplus >= 201703L
2390 
2391  template <typename TagT> static constexpr bool IsValidTag() {
2392  return std::is_same<TagT, mode_tag_t<AccessMode>>::value;
2393  }
2394 
2395 #endif
2396 
2397  void
2398  __init(typename accessor<DataT, Dimensions, AccessMode, target::host_buffer,
2399  access::placeholder::false_t>::ConcreteASPtrType Ptr,
2400  range<AdjustedDim> AccessRange, range<AdjustedDim> MemRange,
2401  id<AdjustedDim> Offset) {
2402  AccessorT::__init(Ptr, AccessRange, MemRange, Offset);
2403  }
2404 
2405 public:
2407 
2408  // The list of host_accessor constructors with their arguments
2409  // -------+---------+-------+----+----------+--------------
2410  // Dimensions = 0
2411  // -------+---------+-------+----+----------+--------------
2412  // buffer | | | | | property_list
2413  // buffer | handler | | | | property_list
2414  // -------+---------+-------+----+----------+--------------
2415  // Dimensions >= 1
2416  // -------+---------+-------+----+----------+--------------
2417  // buffer | | | | | property_list
2418  // buffer | | | | mode_tag | property_list
2419  // buffer | handler | | | | property_list
2420  // buffer | handler | | | mode_tag | property_list
2421  // buffer | | range | | | property_list
2422  // buffer | | range | | mode_tag | property_list
2423  // buffer | handler | range | | | property_list
2424  // buffer | handler | range | | mode_tag | property_list
2425  // buffer | | range | id | | property_list
2426  // buffer | | range | id | mode_tag | property_list
2427  // buffer | handler | range | id | | property_list
2428  // buffer | handler | range | id | mode_tag | property_list
2429  // -------+---------+-------+----+----------+--------------
2430 
2431  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2432  typename = typename detail::enable_if_t<
2433  std::is_same<T, DataT>::value && Dims == 0>>
2435  buffer<T, 1, AllocatorT> &BufferRef,
2436  const property_list &PropertyList = {},
2437  const detail::code_location CodeLoc = detail::code_location::current())
2438  : AccessorT(BufferRef, PropertyList, CodeLoc) {}
2439 
2440  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2441  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2443  buffer<T, Dims, AllocatorT> &BufferRef,
2444  const property_list &PropertyList = {},
2445  const detail::code_location CodeLoc = detail::code_location::current())
2446  : AccessorT(BufferRef, PropertyList, CodeLoc) {}
2447 
2448 #if __cplusplus >= 201703L
2449 
2450  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2451  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2452  host_accessor(
2453  buffer<DataT, Dimensions, AllocatorT> &BufferRef, mode_tag_t<AccessMode>,
2454  const property_list &PropertyList = {},
2455  const detail::code_location CodeLoc = detail::code_location::current())
2456  : host_accessor(BufferRef, PropertyList, CodeLoc) {}
2457 
2458 #endif
2459 
2460  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2461  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2463  buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2464  const property_list &PropertyList = {},
2465  const detail::code_location CodeLoc = detail::code_location::current())
2466  : AccessorT(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
2467 
2468 #if __cplusplus >= 201703L
2469 
2470  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2471  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2472  host_accessor(
2473  buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2474  handler &CommandGroupHandler, mode_tag_t<AccessMode>,
2475  const property_list &PropertyList = {},
2476  const detail::code_location CodeLoc = detail::code_location::current())
2477  : host_accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
2478 
2479 #endif
2480 
2481  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2482  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2485  range<Dimensions> AccessRange, const property_list &PropertyList = {},
2486  const detail::code_location CodeLoc = detail::code_location::current())
2487  : AccessorT(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
2488 
2489 #if __cplusplus >= 201703L
2490 
2491  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2492  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2493  host_accessor(
2494  buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2495  range<Dimensions> AccessRange, mode_tag_t<AccessMode>,
2496  const property_list &PropertyList = {},
2497  const detail::code_location CodeLoc = detail::code_location::current())
2498  : host_accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
2499 
2500 #endif
2501 
2502  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2503  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2506  handler &CommandGroupHandler, range<Dimensions> AccessRange,
2507  const property_list &PropertyList = {},
2508  const detail::code_location CodeLoc = detail::code_location::current())
2509  : AccessorT(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
2510  CodeLoc) {}
2511 
2512 #if __cplusplus >= 201703L
2513 
2514  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2515  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2516  host_accessor(
2517  buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2518  handler &CommandGroupHandler, range<Dimensions> AccessRange,
2519  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
2520  const detail::code_location CodeLoc = detail::code_location::current())
2521  : host_accessor(BufferRef, CommandGroupHandler, AccessRange, {},
2522  PropertyList, CodeLoc) {}
2523 
2524 #endif
2525 
2526  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2527  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2530  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
2531  const property_list &PropertyList = {},
2532  const detail::code_location CodeLoc = detail::code_location::current())
2533  : AccessorT(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
2534  }
2535 
2536 #if __cplusplus >= 201703L
2537 
2538  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2539  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2540  host_accessor(
2541  buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2542  range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
2543  mode_tag_t<AccessMode>, const property_list &PropertyList = {},
2544  const detail::code_location CodeLoc = detail::code_location::current())
2545  : host_accessor(BufferRef, AccessRange, AccessOffset, PropertyList,
2546  CodeLoc) {}
2547 
2548 #endif
2549 
2550  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2551  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2554  handler &CommandGroupHandler, range<Dimensions> AccessRange,
2555  id<Dimensions> AccessOffset, const property_list &PropertyList = {},
2556  const detail::code_location CodeLoc = detail::code_location::current())
2557  : AccessorT(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2558  PropertyList, CodeLoc) {}
2559 
2560 #if __cplusplus >= 201703L
2561 
2562  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
2563  typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2564  host_accessor(
2565  buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2566  handler &CommandGroupHandler, range<Dimensions> AccessRange,
2567  id<Dimensions> AccessOffset, mode_tag_t<AccessMode>,
2568  const property_list &PropertyList = {},
2569  const detail::code_location CodeLoc = detail::code_location::current())
2570  : host_accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2571  PropertyList, CodeLoc) {}
2572 
2573 #endif
2574 };
2575 
2576 #if __cplusplus >= 201703L
2577 
2578 template <typename DataT, int Dimensions, typename AllocatorT>
2579 host_accessor(buffer<DataT, Dimensions, AllocatorT>)
2580  -> host_accessor<DataT, Dimensions, access::mode::read_write>;
2581 
2582 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1>
2583 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1)
2584  -> host_accessor<DataT, Dimensions,
2585  detail::deduceAccessMode<Type1, Type1>()>;
2586 
2587 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2588  typename Type2>
2589 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2)
2590  -> host_accessor<DataT, Dimensions,
2591  detail::deduceAccessMode<Type1, Type2>()>;
2592 
2593 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2594  typename Type2, typename Type3>
2595 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3)
2596  -> host_accessor<DataT, Dimensions,
2597  detail::deduceAccessMode<Type2, Type3>()>;
2598 
2599 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2600  typename Type2, typename Type3, typename Type4>
2601 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
2602  -> host_accessor<DataT, Dimensions,
2603  detail::deduceAccessMode<Type3, Type4>()>;
2604 
2605 template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
2606  typename Type2, typename Type3, typename Type4, typename Type5>
2607 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
2608  Type5) -> host_accessor<DataT, Dimensions,
2609  detail::deduceAccessMode<Type4, Type5>()>;
2610 
2611 #endif
2612 
2613 } // namespace sycl
2614 } // __SYCL_INLINE_NAMESPACE(cl)
2615 
2616 namespace std {
2617 template <typename DataT, int Dimensions, cl::sycl::access::mode AccessMode,
2618  cl::sycl::access::target AccessTarget,
2619  cl::sycl::access::placeholder IsPlaceholder>
2620 struct hash<cl::sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
2621  IsPlaceholder>> {
2622  using AccType = cl::sycl::accessor<DataT, Dimensions, AccessMode,
2623  AccessTarget, IsPlaceholder>;
2624 
2625  size_t operator()(const AccType &A) const {
2626 #ifdef __SYCL_DEVICE_ONLY__
2627  // Hash is not supported on DEVICE. Just return 0 here.
2628  (void)A;
2629  return 0;
2630 #else
2631  // getSyclObjImpl() here returns a pointer to either AccessorImplHost
2632  // or LocalAccessorImplHost depending on the AccessTarget.
2633  auto AccImplPtr = cl::sycl::detail::getSyclObjImpl(A);
2634  return hash<decltype(AccImplPtr)>()(AccImplPtr);
2635 #endif
2636  }
2637 };
2638 
2639 } // namespace std
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:2234
cl::sycl::ext::oneapi::property::no_offset
Definition: accessor_properties.hpp:75
property_list.hpp
cl::sycl::buffer::get_range
range< dimensions > get_range() const
Definition: buffer.hpp:372
cl::sycl::accessor::get_pointer
constant_ptr< DataT > get_pointer() const
Definition: accessor.hpp:1791
cl::sycl::detail::AccPropBufferLocation
@ AccPropBufferLocation
Definition: property_helper.hpp:51
cl::sycl::detail::getLinearIndex
size_t getLinearIndex(const T< Dims > &Index, const U< Dims > &Range)
Definition: common.hpp:336
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:299
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:2354
cl::sycl::accessor::accessor
accessor(buffer< T, 1, AllocatorT > &BufferRef, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1038
cl::sycl::accessor::accessor
accessor(buffer< T, 1, AllocatorT > &BufferRef, handler &CommandGroupHandler, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1097
cl::sycl::detail::accessor_common::AccessorSubscript
Definition: accessor.hpp:307
T
cl::sycl::info::device
device
Definition: info_desc.hpp:53
cl::sycl::detail::image_accessor::read
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:624
cl::sycl::detail::__image_array_slice__::read
DataT read(const CoordT &Coords) const
Definition: accessor.hpp:721
PI_INVALID_OPERATION
@ PI_INVALID_OPERATION
Definition: pi.h:88
__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:2038
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:1763
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::accessor
accessor(handler &, const property_list &propList, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2106
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:757
cl::sycl::accessor< DataT, 1, access_mode::read_write, target::host_buffer, access::placeholder::false_t >::reference
DataT & reference
Definition: accessor.hpp:975
cl::sycl::detail::__image_array_slice__::get_range
range< Dims > get_range() const
Definition: accessor.hpp:765
cl::sycl::access::placeholder
placeholder
Definition: access.hpp:43
cl::sycl::detail::type_list
Definition: type_list.hpp:23
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1254
cl::sycl::host_accessor::IsSameAsBuffer
static constexpr bool IsSameAsBuffer()
Definition: accessor.hpp:2385
cl::sycl::image::get_range
range< Dimensions > get_range() const
Definition: image.hpp:258
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::operator!=
bool operator!=(const accessor &Rhs) const
Definition: accessor.hpp:2218
cl::sycl::accessor::operator[]
AccessorCommonT::template AccessorSubscript< Dims - 1 > operator[](size_t Index) const
Definition: accessor.hpp:1770
cl::sycl::id< Dims >
cl::sycl::image::size
size_t size() const noexcept
Definition: image.hpp:272
cl::sycl::detail::__image_array_slice__::write
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:735
cl::sycl::accessor< DataT, 1, access_mode::read_write, target::host_buffer, access::placeholder::false_t >::ConstRefType
const DataT & ConstRefType
Definition: accessor.hpp:829
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:2154
handler_proxy.hpp
accessor_property_list.hpp
cl::sycl::detail::is_contained
Definition: type_list.hpp:54
cl::sycl::detail::InitializedVal
Definition: common.hpp:227
cl::sycl::errc::feature_not_supported
@ feature_not_supported
cl::sycl::accessor::IsSameAsBuffer
static constexpr bool IsSameAsBuffer()
Definition: accessor.hpp:857
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1223
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={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2552
cl::sycl::detail::IsRunTimePropertyListT
typename std::is_same< ext::oneapi::accessor_property_list<>, T > IsRunTimePropertyListT
Definition: accessor.hpp:235
cl::sycl::accessor::get_range
range< Dimensions > get_range() const
Definition: accessor.hpp:1693
sycl
Definition: invoke_simd.hpp:68
cl::sycl::property_list
Objects of the property_list class are containers for the SYCL properties.
Definition: property_list.hpp:26
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={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1333
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::operator[]
RefType operator[](size_t Index) const
Definition: accessor.hpp:2180
cl::sycl::detail::LocalAccessorBaseDevice
Definition: accessor_impl.hpp:58
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::accessor
accessor(range< Dimensions > AllocationSize, handler &, const property_list &propList, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2137
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::accessor
accessor(cl::sycl::image< Dimensions, AllocatorT > &Image)
Definition: accessor.hpp:2293
operator==
bool operator==(const Slab &Lhs, const Slab &Rhs)
Definition: usm_allocator.cpp:568
cl::sycl::detail::const_if_const_AS
DataT const_if_const_AS
Definition: type_traits.hpp:348
cl::sycl::detail::code_location
Definition: common.hpp:54
cl::sycl::detail::AccessorBaseHost::impl
AccessorImplPtr impl
Definition: accessor_impl.hpp:162
cl::sycl::buffer
Defines a shared array that can be used by kernels in queues.
Definition: buffer.hpp:58
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:184
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:974
cl::sycl::host_accessor
Definition: accessor.hpp:2376
cl::sycl::detail::image_accessor::operator!=
bool operator!=(const image_accessor &Rhs) const
Definition: accessor.hpp:570
cl::sycl::host_accessor::host_accessor
host_accessor(buffer< DataT, Dimensions, AllocatorT > &BufferRef, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2528
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1553
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::operator==
bool operator==(const accessor &Rhs) const
Definition: accessor.hpp:2217
cl::sycl::buffer::size
size_t size() const noexcept
Definition: buffer.hpp:376
cl::sycl::access_mode
access::mode access_mode
Definition: access.hpp:65
cl::sycl::accessor::get_pointer
global_ptr< DataT > get_pointer() const
Definition: accessor.hpp:1784
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:2209
cl::sycl::accessor::getQualifiedPtr
PtrType getQualifiedPtr() const
Definition: accessor.hpp:963
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1442
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::const_reference
const DataT & const_reference
Definition: accessor.hpp:2090
export.hpp
cl::sycl::detail::image_accessor::image_accessor
image_accessor(image< Dims, AllocatorT > &ImageRef, handler &CommandGroupHandlerRef, int ImageElementSize)
Definition: accessor.hpp:532
cl::sycl::detail::image_accessor::image_accessor
image_accessor(image< Dims, AllocatorT > &ImageRef, int ImageElementSize)
Definition: accessor.hpp:503
cl::sycl::detail::LocalAccessorBaseDevice::MemRange
range< Dims > MemRange
Definition: accessor_impl.hpp:66
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::image, IsPlaceholder >::accessor
accessor(cl::sycl::image< Dimensions, AllocatorT > &Image, handler &CommandGroupHandler, const property_list &propList)
Definition: accessor.hpp:2247
cl::sycl::detail::__image_array_slice__::read
DataT read(const CoordT &Coords, const sampler &Smpl) const
Definition: accessor.hpp:728
cl::sycl::accessor::accessor
accessor(const accessor< DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder, ext::oneapi::accessor_property_list< NewPropsT... >> &Other, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1663
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::PtrType
detail::const_if_const_AS< AS, DataT > * PtrType
Definition: accessor.hpp:2041
cl::sycl::accessor::accessor
accessor(buffer< T, 1, AllocatorT > &BufferRef, handler &CommandGroupHandler, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1067
cl::sycl::device
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:35
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::reference
DataT & reference
Definition: accessor.hpp:2089
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:2203
cl::sycl::detail::IsValidCoordDataT
Definition: accessor.hpp:354
generic_type_traits.hpp
cl::sycl::detail::IsCxPropertyList
Definition: accessor.hpp:237
cl::sycl::detail::image_accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::reference
DataT & reference
Definition: accessor.hpp:487
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:830
cl::sycl::detail::__image_array_slice__
Definition: accessor.hpp:372
cl::sycl::detail::accessor_common::AccessorSubscript::operator[]
ConstRefType operator[](size_t Index) const
Definition: accessor.hpp:347
cl::sycl::detail::accessor_common::AccessorSubscript::AccessorSubscript
AccessorSubscript(AccType Accessor, id< Dims > IDs)
Definition: accessor.hpp:314
cl::sycl::detail::TargetToAS
Definition: access.hpp:136
cl::sycl::host_accessor::host_accessor
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2442
cl::sycl::accessor
Buffer accessor.
Definition: accessor.hpp:223
cl::sycl::access::target
target
Definition: access.hpp:17
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:824
cl::sycl::accessor::operator!=
bool operator!=(const accessor &Rhs) const
Definition: accessor.hpp:1796
cl::sycl::accessor::getAdjustedMode
static access::mode getAdjustedMode(const PropertyListT &PropertyList)
Definition: accessor.hpp:861
cl::sycl::detail::accessor_common
Definition: accessor.hpp:269
cl::sycl::accessor::get_offset
id< Dimensions > get_offset() const
Definition: accessor.hpp:1698
cl::sycl::host_accessor::host_accessor
host_accessor(buffer< T, 1, AllocatorT > &BufferRef, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2434
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1481
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1128
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1378
cl::sycl::detail::IsPropertyListT
typename std::is_base_of< PropertyListBase, T > IsPropertyListT
Definition: accessor.hpp:231
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={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1592
cl::sycl::detail::image_accessor::get_range
range< Dims > get_range() const
Definition: accessor.hpp:608
cl::sycl::detail::image_accessor::read
DataT read(const CoordT &Coords, const sampler &Smpl) const
Definition: accessor.hpp:643
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:19
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:1777
image_ocl_types.hpp
accessor_impl.hpp
cl::sycl::detail::addHostAccessorAndWait
void addHostAccessorAndWait(Requirement *Req)
Definition: accessor_impl.cpp:36
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:298
cl::sycl::accessor::__SYCL2020_DEPRECATED
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
Definition: accessor.hpp:1688
image.hpp
cl::sycl::detail::accessor_common::AccessorSubscript::AccessorSubscript
AccessorSubscript(AccType Accessor, size_t Index)
Definition: accessor.hpp:319
cl::sycl::host_accessor::host_accessor
host_accessor(buffer< DataT, Dimensions, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2504
cl::sycl::host_accessor::host_accessor
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2462
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:603
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::accessor
accessor(buffer< T, 1, AllocatorT > &BufferRef, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1008
cl::sycl::accessor::get_size
size_t get_size() const
Definition: accessor.hpp:1686
cl::sycl::detail::image_accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::value_type
DataT value_type
Definition: accessor.hpp:486
cl::sycl::detail::image_accessor< DataT, Dimensions, AccessMode, access::target::host_image, IsPlaceholder >::const_reference
const DataT & const_reference
Definition: accessor.hpp:488
cl::sycl::accessor::getLinearIndex
size_t getLinearIndex(id< Dims > Id) const
Definition: accessor.hpp:832
PI_INVALID_VALUE
@ PI_INVALID_VALUE
Definition: pi.h:91
atomic.hpp
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::accessor
accessor(handler &, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2093
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:2299
cl::sycl::detail::getSyclObjImpl
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: common.hpp:204
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:711
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:562
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:1690
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::value_type
DataT value_type
Definition: accessor.hpp:2088
std::hash< cl::sycl::accessor< DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder > >::operator()
size_t operator()(const AccType &A) const
Definition: accessor.hpp:2625
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:340
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:828
std
Definition: accessor.hpp:2616
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::get_range
range< Dims > get_range() const
Definition: accessor.hpp:2161
cl::sycl::host_accessor::host_accessor
host_accessor(buffer< DataT, Dimensions, AllocatorT > &BufferRef, range< Dimensions > AccessRange, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2483
cl::sycl::ext::oneapi::experimental::detail::get_property
static constexpr std::enable_if_t< HasProperty, typename FindCompileTimePropertyValueType< CTPropertyT, PropertiesT >::type > get_property()
Definition: properties.hpp:65
__SYCL_UNROLL
#define __SYCL_UNROLL(x)
Definition: defines_elementary.hpp:120
cl::sycl::detail::image_accessor
Definition: accessor.hpp:377
cl::sycl::detail::convertToArrayOfN
static T< NewDim > convertToArrayOfN(T< OldDim > OldObj)
Definition: accessor.hpp:254
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::getQualifiedPtr
PtrType getQualifiedPtr() const
Definition: accessor.hpp:2073
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1391
cl::sycl::detail::accessor_common::AccType
accessor< DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder, PropertyListT > AccType
Definition: accessor.hpp:303
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:300
cl::sycl::mode_target_tag_t
Definition: access.hpp:71
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:976
buffer.hpp
cl::sycl::detail::AccessorBaseHost
Definition: accessor_impl.hpp:132
cl::sycl::host_accessor::host_accessor
host_accessor()
Definition: accessor.hpp:2406
common.hpp
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::AccessorSubscript
typename AccessorCommonT::template AccessorSubscript< Dims > AccessorSubscript
Definition: accessor.hpp:2036
cl::sycl::image_channel_type
image_channel_type
Definition: image.hpp:41
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:2156
cl::sycl::detail::__image_array_slice__::size
size_t size() const noexcept
Definition: accessor.hpp:759
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:2213
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:2341
cl::sycl::stream
A buffered output stream that allows outputting the values of built-in, vector and SYCL types to the ...
Definition: stream.hpp:743
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:826
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::RefType
detail::const_if_const_AS< AS, DataT > & RefType
Definition: accessor.hpp:2040
cl::sycl::detail::image_accessor::size
size_t size() const noexcept
Definition: accessor.hpp:605
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:2368
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::detail::accessor_common::AccessorSubscript::operator[]
RefType operator[](size_t Index) const
Definition: accessor.hpp:332
pointers.hpp
cl::sycl::detail::image_accessor::write
void write(const CoordT &Coords, const DataT &Color) const
Definition: accessor.hpp:668
cl::sycl::Dimensions
Dimensions
Definition: backend.hpp:138
cl::sycl::accessor::operator==
bool operator==(const accessor &Rhs) const
Definition: accessor.hpp:1795
cl::sycl::buffer::isOutOfBounds
bool isOutOfBounds(const id< dimensions > &offset, const range< dimensions > &newRange, const range< dimensions > &parentRange)
Definition: buffer.hpp:522
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, const ext::oneapi::accessor_property_list< PropTypes... > &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1160
cl::sycl::accessor::is_placeholder
constexpr bool is_placeholder() const
Definition: accessor.hpp:1684
cl::sycl::image_channel_order
image_channel_order
Definition: image.hpp:23
cl::sycl::accessor::accessor
accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:1320
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::accessor
accessor(range< Dimensions > AllocationSize, handler &, const detail::code_location CodeLoc=detail::code_location::current())
Definition: accessor.hpp:2122
accessor_properties.hpp
cl::sycl::accessor< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >::getLinearIndex
size_t getLinearIndex(id< Dims > Id) const
Definition: accessor.hpp:2080
cl::sycl::detail::constructorNotification
void constructorNotification(void *BufferObj, void *AccessorObj, access::target Target, access::mode Mode, const code_location &CodeLoc)
Definition: accessor_impl.cpp:42
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:2398
cl::sycl::errc::accessor
@ accessor
cl::sycl::detail::DecoratedType
Definition: access.hpp:159
spirv_types.hpp
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:2158