37 #include <type_traits>
215 namespace ext::intel::esimd::detail {
217 class AccessorPrivateProxy;
233 const sycl::range<3> &MemoryRange,
void *Data =
nullptr)
234 : MOffset(Offset), MAccessRange(Range), MMemoryRange(MemoryRange),
240 void *MData =
nullptr;
241 void *Reserved =
nullptr;
245 template <
size_t... Inds,
class F>
250 template <
size_t count,
class F>
void dim_loop(F &&f) {
251 dim_loop_impl(std::make_index_sequence<count>{}, std::forward<F>(f));
259 template <
typename T>
262 template <
typename T>
264 typename std::is_same<ext::oneapi::accessor_property_list<>, T>;
267 constexpr
static bool value =
false;
270 template <
typename... Props>
272 constexpr
static bool value =
true;
276 constexpr
static bool value =
false;
288 constexpr
static bool IsHostBuf = AccessTarget == access::target::host_buffer;
299 constexpr
static bool IsPlaceH = !IsHostBuf;
304 constexpr
static bool IsGlobalBuf =
305 AccessTarget == access::target::global_buffer;
307 constexpr
static bool IsConstantBuf =
308 AccessTarget == access::target::constant_buffer;
310 constexpr
static bool IsAccessAnyWrite =
314 AccessMode == access::mode::discard_read_write;
316 constexpr
static bool IsAccessReadOnly =
AccessMode == access::mode::read;
318 constexpr
static bool IsAccessReadWrite =
321 constexpr
static bool IsAccessAtomic =
AccessMode == access::mode::atomic;
329 template <
int SubDims,
341 : MIDs(IDs), MAccessor(Accessor) {}
349 template <
int CurDims = SubDims,
351 auto operator[](
size_t Index) {
352 MIDs[Dims - CurDims] = Index;
356 template <
int CurDims = SubDims,
358 CurDims == 1 && (IsAccessReadOnly || IsAccessAnyWrite)>>
359 typename AccType::reference
operator[](
size_t Index)
const {
360 MIDs[Dims - CurDims] = Index;
361 return MAccessor[MIDs];
364 template <
int CurDims = SubDims>
368 MIDs[Dims - CurDims] = Index;
369 return MAccessor[MIDs];
374 template <
typename MayBeTag1,
typename MayBeTag2>
379 if constexpr (std::is_same<MayBeTag1,
381 std::is_same<MayBeTag2,
383 return access::mode::read;
386 if constexpr (std::is_same<MayBeTag1,
388 std::is_same<MayBeTag2,
394 std::is_same<MayBeTag1,
396 access::target::constant_buffer>>::value ||
397 std::is_same<MayBeTag2,
399 access::target::constant_buffer>>::value) {
400 return access::mode::read;
406 template <
typename MayBeTag1,
typename MayBeTag2>
409 std::is_same<MayBeTag1,
411 access::target::constant_buffer>>::value ||
412 std::is_same<MayBeTag2,
414 access::target::constant_buffer>>::value) {
415 return access::target::constant_buffer;
418 return defaultTarget;
447 : Offset(Offset), AccessRange(AccessRange), MemRange(MemoryRange) {}
459 class AccessorImplHost;
474 int ElemSize,
int OffsetInBytes = 0,
475 bool IsSubBuffer =
false,
480 int ElemSize,
bool IsPlaceH,
int OffsetInBytes = 0,
481 bool IsSubBuffer =
false,
486 range<3> &getAccessRange();
487 range<3> &getMemoryRange();
489 unsigned int getElemSize()
const;
491 const id<3> &getOffset()
const;
492 const range<3> &getAccessRange()
const;
493 const range<3> &getMemoryRange()
const;
494 void *getPtr()
const;
495 bool isPlaceholder()
const;
497 detail::AccHostDataT &getAccData();
499 const property_list &getPropList()
const;
501 void *getMemoryObject()
const;
516 friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
529 sycl::range<3> &getSize();
530 const sycl::range<3> &getSize()
const;
532 void *getPtr()
const;
534 int getElementSize();
549 constexpr
static bool value =
553 constexpr
static bool value =
558 constexpr
static bool value =
571 #ifndef __SYCL_DEVICE_ONLY__
581 OCLImageTy MImageObj;
587 void imageAccessorInit(OCLImageTy Image) { MImageObj = Image; }
591 template <
typename T1,
int T2, access::mode T3, access::placeholder T4>
594 constexpr
static bool IsHostImageAcc =
595 (AccessTarget == access::target::host_image);
597 constexpr
static bool IsImageAcc = (AccessTarget == access::target::image);
599 constexpr
static bool IsImageArrayAcc =
600 (AccessTarget == access::target::image_array);
602 constexpr
static bool IsImageAccessWriteOnly =
606 constexpr
static bool IsImageAccessAnyWrite =
609 constexpr
static bool IsImageAccessReadOnly =
612 constexpr
static bool IsImageAccessAnyRead =
615 static_assert(std::is_same<DataT, cl_int4>::value ||
616 std::is_same<DataT, cl_uint4>::value ||
617 std::is_same<DataT, cl_float4>::value ||
618 std::is_same<DataT, cl_half4>::value,
619 "The data type of an image accessor must be only cl_int4, "
620 "cl_uint4, cl_float4 or cl_half4 from SYCL namespace");
622 static_assert(IsImageAcc || IsHostImageAcc || IsImageArrayAcc,
623 "Expected image type");
626 "Expected false as Placeholder value for image accessor.");
629 ((IsImageAcc || IsImageArrayAcc) &&
630 (IsImageAccessWriteOnly || IsImageAccessReadOnly)) ||
631 (IsHostImageAcc && (IsImageAccessAnyWrite || IsImageAccessAnyRead)),
632 "Access modes can be only read/write/discard_write for image/image_array "
633 "target accessor, or they can be only "
634 "read/write/discard_write/read_write for host_image target accessor.");
637 "Dimensions can be 1/2/3 for image accessor.");
639 template <
typename Param>
640 void checkDeviceFeatureSupported(
const device &
Device) {
643 PI_ERROR_INVALID_OPERATION);
646 #ifdef __SYCL_DEVICE_ONLY__
648 sycl::vec<int, Dimensions> getRangeInternal()
const {
649 return __invoke_ImageQuerySize<sycl::vec<int, Dimensions>, OCLImageTy>(
653 size_t getElementSize()
const {
654 int ChannelType = __invoke_ImageQueryFormat<int, OCLImageTy>(MImageObj);
655 int ChannelOrder = __invoke_ImageQueryOrder<int, OCLImageTy>(MImageObj);
656 int ElementSize = getSPIRVElementSize(ChannelType, ChannelOrder);
662 sycl::vec<int, Dimensions> getRangeInternal()
const {
664 throw runtime_error(
"image::getRangeInternal() is not implemented for host",
665 PI_ERROR_INVALID_OPERATION);
666 return sycl::vec<int, Dimensions>{1};
671 #ifndef __SYCL_DEVICE_ONLY__
677 friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
679 #ifdef __SYCL_DEVICE_ONLY__
680 const OCLImageTy getNativeImageObj()
const {
return MImageObj; }
690 #ifdef __SYCL_DEVICE_ONLY__
702 #ifdef __SYCL_DEVICE_ONLY__
705 (void)ImageElementSize;
711 detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
712 detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
715 MImageCount(ImageRef.
size()),
716 MImgChannelOrder(ImageRef.getChannelOrder()),
717 MImgChannelType(ImageRef.getChannelType()) {
730 handler &CommandGroupHandlerRef,
int ImageElementSize)
731 #ifdef __SYCL_DEVICE_ONLY__
734 (void)CommandGroupHandlerRef;
735 (void)ImageElementSize;
741 detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
742 detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
745 MImageCount(ImageRef.
size()),
746 MImgChannelOrder(ImageRef.getChannelOrder()),
747 MImgChannelType(ImageRef.getChannelType()) {
748 checkDeviceFeatureSupported<info::device::image_support>(
757 #ifndef __SYCL_DEVICE_ONLY__
776 #ifdef __SYCL_DEVICE_ONLY__
779 size_t get_count()
const {
return size(); }
780 size_t size() const noexcept {
return get_range<Dimensions>().size(); }
782 template <
int Dims = Dimensions,
typename = detail::enable_if_t<Dims == 1>>
783 range<1> get_range()
const {
784 int Range = getRangeInternal();
785 return range<1>(Range);
787 template <
int Dims = Dimensions,
typename = detail::enable_if_t<Dims == 2>>
788 range<2> get_range()
const {
789 int2 Range = getRangeInternal();
790 return range<2>(Range[0], Range[1]);
792 template <
int Dims = Dimensions,
typename = detail::enable_if_t<Dims == 3>>
793 range<3> get_range()
const {
794 int3 Range = getRangeInternal();
795 return range<3>(Range[0], Range[1], Range[2]);
800 size_t get_count()
const {
return size(); };
801 size_t size() const noexcept {
return MImageCount; };
803 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 0)>>
805 return detail::convertToArrayOfN<Dims, 1>(getAccessRange());
814 template <
typename CoordT,
int Dims =
Dimensions,
818 ((IsImageAcc && IsImageAccessReadOnly) ||
819 (IsHostImageAcc && IsImageAccessAnyRead))>>
820 DataT
read(
const CoordT &Coords)
const {
821 #ifdef __SYCL_DEVICE_ONLY__
822 return __invoke__ImageRead<DataT, OCLImageTy, CoordT>(MImageObj, Coords);
824 sampler Smpl(coordinate_normalization_mode::unnormalized,
825 addressing_mode::none, filtering_mode::nearest);
826 return read<CoordT, Dims>(Coords, Smpl);
834 template <
typename CoordT,
int Dims =
Dimensions,
837 ((IsImageAcc && IsImageAccessReadOnly) ||
838 (IsHostImageAcc && IsImageAccessAnyRead))>>
839 DataT
read(
const CoordT &Coords,
const sampler &Smpl)
const {
840 #ifdef __SYCL_DEVICE_ONLY__
841 return __invoke__ImageReadSampler<DataT, OCLImageTy, CoordT>(
842 MImageObj, Coords, Smpl.impl.m_Sampler);
844 return imageReadSamplerHostImpl<CoordT, DataT>(
845 Coords, Smpl, getAccessRange() ,
846 getOffset() , MImgChannelType, MImgChannelOrder,
847 AccessorBaseHost::getPtr() ,
848 AccessorBaseHost::getElemSize());
858 template <
typename CoordT,
int Dims =
Dimensions,
862 ((IsImageAcc && IsImageAccessWriteOnly) ||
863 (IsHostImageAcc && IsImageAccessAnyWrite))>>
864 void write(
const CoordT &Coords,
const DataT &Color)
const {
865 #ifdef __SYCL_DEVICE_ONLY__
866 __invoke__ImageWrite<OCLImageTy, CoordT, DataT>(MImageObj, Coords, Color);
869 AccessorBaseHost::getElemSize(), MImgChannelType,
871 AccessorBaseHost::getPtr() );
881 "Image slice cannot have more then 2 dimensions");
885 template <
typename CoordT,
886 typename CoordElemType =
888 sycl::vec<CoordElemType, AdjustedDims>
889 getAdjustedCoords(
const CoordT &Coords)
const {
890 CoordElemType LastCoord = 0;
892 if (std::is_same<float, CoordElemType>::value) {
893 sycl::vec<int, Dimensions + 1> Size = MBaseAcc.getRangeInternal();
895 MIdx /
static_cast<float>(Size.template swizzle<Dimensions>());
900 sycl::vec<CoordElemType, Dimensions> LeftoverCoords{LastCoord};
901 sycl::vec<CoordElemType, AdjustedDims> AdjustedCoords{Coords,
903 return AdjustedCoords;
912 : MBaseAcc(BaseAcc), MIdx(Idx) {}
914 template <
typename CoordT,
int Dims =
Dimensions,
917 DataT
read(
const CoordT &Coords)
const {
918 return MBaseAcc.read(getAdjustedCoords(Coords));
921 template <
typename CoordT,
int Dims =
Dimensions,
924 DataT
read(
const CoordT &Coords,
const sampler &Smpl)
const {
925 return MBaseAcc.read(getAdjustedCoords(Coords), Smpl);
928 template <
typename CoordT,
int Dims =
Dimensions,
931 void write(
const CoordT &Coords,
const DataT &Color)
const {
932 return MBaseAcc.write(getAdjustedCoords(Coords), Color);
935 #ifdef __SYCL_DEVICE_ONLY__
937 size_t get_count()
const {
return size(); }
938 size_t size() const noexcept {
return get_range<Dimensions>().size(); }
940 template <
int Dims = Dimensions,
typename = detail::enable_if_t<Dims == 1>>
941 range<1> get_range()
const {
942 int2 Count = MBaseAcc.getRangeInternal();
943 return range<1>(Count.x());
945 template <
int Dims = Dimensions,
typename = detail::enable_if_t<Dims == 2>>
946 range<2> get_range()
const {
947 int3 Count = MBaseAcc.getRangeInternal();
948 return range<2>(Count.x(), Count.y());
954 size_t get_count()
const {
return size(); }
956 return MBaseAcc.MImageCount / MBaseAcc.getAccessRange()[
Dimensions];
962 return detail::convertToArrayOfN<Dims, 1>(MBaseAcc.getAccessRange());
983 typename PropertyListT>
985 #ifndef __SYCL_DEVICE_ONLY__
986 public detail::AccessorBaseHost,
988 public detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
989 IsPlaceholder, PropertyListT>,
990 public detail::OwnerLessBase<
991 accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
994 static_assert((AccessTarget == access::target::global_buffer ||
995 AccessTarget == access::target::constant_buffer ||
996 AccessTarget == access::target::host_buffer),
997 "Expected buffer type");
999 static_assert((AccessTarget == access::target::global_buffer ||
1000 AccessTarget == access::target::host_buffer) ||
1001 (AccessTarget == access::target::constant_buffer &&
1003 "Access mode can be only read for constant buffers");
1005 static_assert(detail::IsPropertyListT<PropertyListT>::value,
1006 "PropertyListT must be accessor_property_list");
1008 using AccessorCommonT =
1014 using AccessorCommonT::AS;
1017 static constexpr
bool IsAccessAnyWrite = AccessorCommonT::IsAccessAnyWrite;
1018 static constexpr
bool IsAccessReadOnly = AccessorCommonT::IsAccessReadOnly;
1019 static constexpr
bool IsConstantBuf = AccessorCommonT::IsConstantBuf;
1020 static constexpr
bool IsGlobalBuf = AccessorCommonT::IsGlobalBuf;
1021 static constexpr
bool IsHostBuf = AccessorCommonT::IsHostBuf;
1022 static constexpr
bool IsPlaceH = AccessorCommonT::IsPlaceH;
1024 using AccessorSubscript =
1025 typename AccessorCommonT::template AccessorSubscript<Dims>;
1027 using ConcreteASPtrType =
typename detail::DecoratedType<DataT, AS>::type *;
1029 using RefType = detail::const_if_const_AS<AS, DataT> &;
1030 using ConstRefType =
const DataT &;
1031 using PtrType = detail::const_if_const_AS<AS, DataT> *;
1033 template <
int Dims = Dimensions>
size_t getLinearIndex(id<Dims> Id)
const {
1036 detail::dim_loop<Dims>([&,
this](
size_t I) {
1037 Result = Result * getMemoryRange()[I] + Id[I];
1040 #ifndef __SYCL_DEVICE_ONLY__
1043 Result += getOffset()[I];
1051 template <
typename T,
int Dims>
1052 struct IsSameAsBuffer
1054 (Dims == Dimensions)> {};
1056 static access::mode getAdjustedMode(
const PropertyListT &PropertyList) {
1059 if (PropertyList.template has_property<property::no_init>() ||
1060 PropertyList.template has_property<property::noinit>()) {
1062 AdjustedMode = access::mode::discard_write;
1064 AdjustedMode = access::mode::discard_read_write;
1068 return AdjustedMode;
1071 template <
typename TagT>
1074 std::is_same<TagT, mode_tag_t<AccessMode>>,
1075 std::is_same<TagT, mode_target_tag_t<AccessMode, AccessTarget>>> {};
1077 #ifdef __SYCL_DEVICE_ONLY__
1079 id<AdjustedDim> &getOffset() {
return impl.Offset; }
1080 range<AdjustedDim> &getAccessRange() {
return impl.AccessRange; }
1081 range<AdjustedDim> &getMemoryRange() {
return impl.MemRange; }
1083 const id<AdjustedDim> &getOffset()
const {
return impl.Offset; }
1084 const range<AdjustedDim> &getAccessRange()
const {
return impl.AccessRange; }
1085 const range<AdjustedDim> &getMemoryRange()
const {
return impl.MemRange; }
1087 detail::AccessorImplDevice<AdjustedDim> impl;
1090 ConcreteASPtrType MData;
1094 const ConcreteASPtrType getNativeImageObj()
const {
return MData; }
1096 void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
1097 range<AdjustedDim> MemRange, id<AdjustedDim> Offset) {
1099 detail::dim_loop<AdjustedDim>([&,
this](
size_t I) {
1102 getOffset()[I] = Offset[I];
1104 getAccessRange()[I] = AccessRange[I];
1105 getMemoryRange()[I] = MemRange[I];
1110 MData += getTotalOffset();
1116 void __init_esimd(ConcreteASPtrType Ptr) { MData = Ptr; }
1118 ConcreteASPtrType getQualifiedPtr() const noexcept {
return MData; }
1120 template <
typename DataT_,
int Dimensions_,
access::mode AccessMode_,
1122 typename PropertyListT_>
1125 #ifndef __SYCL_DEVICE_ONLY__
1126 using AccessorBaseHost::impl;
1132 : impl({}, detail::InitializedVal<AdjustedDim, range>::template get<0>(),
1133 detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
1137 : detail::AccessorBaseHost{Impl} {}
1139 id<3> &getOffset() {
1140 if constexpr (IsHostBuf)
1141 return MAccData->MOffset;
1143 return AccessorBaseHost::getOffset();
1146 range<3> &getAccessRange() {
return AccessorBaseHost::getAccessRange(); }
1147 range<3> &getMemoryRange() {
1148 if constexpr (IsHostBuf)
1149 return MAccData->MMemoryRange;
1151 return AccessorBaseHost::getMemoryRange();
1153 void *getPtr() {
return AccessorBaseHost::getPtr(); }
1155 const id<3> &getOffset()
const {
1156 if constexpr (IsHostBuf)
1157 return MAccData->MOffset;
1159 return AccessorBaseHost::getOffset();
1161 const range<3> &getAccessRange()
const {
1162 return AccessorBaseHost::getAccessRange();
1164 const range<3> &getMemoryRange()
const {
1165 if constexpr (IsHostBuf)
1166 return MAccData->MMemoryRange;
1168 return AccessorBaseHost::getMemoryRange();
1171 void *getPtr()
const {
return AccessorBaseHost::getPtr(); }
1173 void initHostAcc() { MAccData = &getAccData(); }
1176 void GDBMethodsAnchor() {
1178 const auto *this_const =
this;
1179 (void)getMemoryRange();
1180 (void)this_const->getMemoryRange();
1182 (void)this_const->getOffset();
1184 (void)this_const->getPtr();
1185 (void)getAccessRange();
1186 (void)this_const->getAccessRange();
1190 detail::AccHostDataT *MAccData =
nullptr;
1192 char padding[
sizeof(detail::AccessorImplDevice<AdjustedDim>) +
1193 sizeof(PtrType) -
sizeof(detail::AccessorBaseHost) -
1196 PtrType getQualifiedPtr() const noexcept {
1197 if constexpr (IsHostBuf)
1198 return reinterpret_cast<PtrType
>(MAccData->MData);
1200 return reinterpret_cast<PtrType
>(AccessorBaseHost::getPtr());
1206 {0, 0, 0}, {0, 0, 0},
1208 getAdjustedMode({}),
1216 friend class sycl::stream;
1217 friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
1226 using value_type =
typename std::conditional<
AccessMode == access_mode::read,
1227 const DataT, DataT>::type;
1228 using reference = value_type &;
1229 using const_reference =
const DataT &;
1231 template <access::decorated IsDecorated>
1232 using accessor_ptr =
1234 global_ptr<value_type, IsDecorated>, value_type *>;
1236 using iterator =
typename detail::accessor_iterator<value_type, Dimensions>;
1237 using const_iterator =
1238 typename detail::accessor_iterator<const value_type, Dimensions>;
1239 using reverse_iterator = std::reverse_iterator<iterator>;
1240 using const_reverse_iterator = std::reverse_iterator<const_iterator>;
1241 using difference_type =
1242 typename std::iterator_traits<iterator>::difference_type;
1268 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1270 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1271 std::is_same<T, DataT>::value && Dims == 0 &&
1272 (IsHostBuf || (IsGlobalBuf || IsConstantBuf))> * =
nullptr>
1274 buffer<T, 1, AllocatorT> &BufferRef,
1275 const property_list &PropertyList = {},
1276 const detail::code_location CodeLoc = detail::code_location::current())
1277 #ifdef __SYCL_DEVICE_ONLY__
1278 : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1282 {0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1283 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1284 getAdjustedMode(PropertyList),
1286 IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1288 preScreenAccessor(BufferRef.size(), PropertyList);
1289 if (!AccessorBaseHost::isPlaceholder())
1293 detail::AccessorBaseHost::impl.
get(),
1299 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1300 typename... PropTypes,
1302 detail::IsCxPropertyList<PropertyListT>::value &&
1303 std::is_same<T, DataT>::value && Dims == 0 &&
1304 (IsHostBuf || (IsGlobalBuf || IsConstantBuf))> * =
nullptr>
1306 buffer<T, 1, AllocatorT> &BufferRef,
1307 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1309 const detail::code_location CodeLoc = detail::code_location::current())
1310 #ifdef __SYCL_DEVICE_ONLY__
1311 : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1315 {0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1316 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1317 getAdjustedMode(PropertyList),
1319 IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1321 preScreenAccessor(BufferRef.size(), PropertyList);
1322 if (!AccessorBaseHost::isPlaceholder())
1326 detail::AccessorBaseHost::impl.
get(),
1332 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1334 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1335 std::is_same<T, DataT>::value && (Dims == 0) &&
1336 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1338 buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1339 const property_list &PropertyList = {},
1340 const detail::code_location CodeLoc = detail::code_location::current())
1341 #ifdef __SYCL_DEVICE_ONLY__
1342 : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1343 (void)CommandGroupHandler;
1348 {0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1349 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1350 getAdjustedMode(PropertyList),
1352 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1353 preScreenAccessor(BufferRef.size(), PropertyList);
1357 detail::AccessorBaseHost::impl.
get(),
1363 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1364 typename... PropTypes,
1366 detail::IsCxPropertyList<PropertyListT>::value &&
1367 std::is_same<T, DataT>::value && (Dims == 0) &&
1368 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1370 buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1371 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1373 const detail::code_location CodeLoc = detail::code_location::current())
1374 #ifdef __SYCL_DEVICE_ONLY__
1375 : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
1376 (void)CommandGroupHandler;
1381 {0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1382 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1383 getAdjustedMode(PropertyList),
1385 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1386 preScreenAccessor(BufferRef.size(), PropertyList);
1390 detail::AccessorBaseHost::impl.
get(),
1396 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1398 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1399 IsSameAsBuffer<T, Dims>::value &&
1400 (IsHostBuf || (IsGlobalBuf || IsConstantBuf))>>
1402 buffer<T, Dims, AllocatorT> &BufferRef,
1403 const property_list &PropertyList = {},
1404 const detail::code_location CodeLoc = detail::code_location::current())
1405 #ifdef __SYCL_DEVICE_ONLY__
1406 : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1412 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1413 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1414 getAdjustedMode(PropertyList),
1416 IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1418 preScreenAccessor(BufferRef.size(), PropertyList);
1419 if (!AccessorBaseHost::isPlaceholder())
1423 detail::AccessorBaseHost::impl.
get(),
1429 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1430 typename... PropTypes,
1432 detail::IsCxPropertyList<PropertyListT>::value &&
1433 IsSameAsBuffer<T, Dims>::value &&
1434 (IsHostBuf || (IsGlobalBuf || IsConstantBuf))>>
1436 buffer<T, Dims, AllocatorT> &BufferRef,
1437 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1439 const detail::code_location CodeLoc = detail::code_location::current())
1440 #ifdef __SYCL_DEVICE_ONLY__
1441 : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1447 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1448 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1449 getAdjustedMode(PropertyList),
1451 IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1453 preScreenAccessor(BufferRef.size(), PropertyList);
1454 if (!AccessorBaseHost::isPlaceholder())
1458 detail::AccessorBaseHost::impl.
get(),
1464 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1467 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1468 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1469 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1471 buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1472 const property_list &PropertyList = {},
1473 const detail::code_location CodeLoc = detail::code_location::current())
1474 :
accessor(BufferRef, PropertyList, CodeLoc) {
1475 adjustAccPropsInBuf(BufferRef);
1478 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1479 typename TagT,
typename... PropTypes,
1481 detail::IsCxPropertyList<PropertyListT>::value &&
1482 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1483 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1485 buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1486 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1488 const detail::code_location CodeLoc = detail::code_location::current())
1489 :
accessor(BufferRef, PropertyList, CodeLoc) {
1490 adjustAccPropsInBuf(BufferRef);
1493 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1495 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1496 IsSameAsBuffer<T, Dims>::value &&
1497 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1499 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1500 const property_list &PropertyList = {},
1501 const detail::code_location CodeLoc = detail::code_location::current())
1502 #ifdef __SYCL_DEVICE_ONLY__
1503 : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1504 (void)CommandGroupHandler;
1510 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1511 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1512 getAdjustedMode(PropertyList),
1514 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1515 preScreenAccessor(BufferRef.size(), PropertyList);
1519 detail::AccessorBaseHost::impl.
get(),
1525 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1526 typename... PropTypes,
1528 detail::IsCxPropertyList<PropertyListT>::value &&
1529 IsSameAsBuffer<T, Dims>::value &&
1530 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1532 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1533 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1535 const detail::code_location CodeLoc = detail::code_location::current())
1536 #ifdef __SYCL_DEVICE_ONLY__
1537 : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1538 (void)CommandGroupHandler;
1544 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1545 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1546 getAdjustedMode(PropertyList),
1548 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1549 preScreenAccessor(BufferRef.size(), PropertyList);
1553 detail::AccessorBaseHost::impl.
get(),
1559 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1562 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1563 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1564 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1566 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1567 TagT,
const property_list &PropertyList = {},
1568 const detail::code_location CodeLoc = detail::code_location::current())
1569 :
accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1570 adjustAccPropsInBuf(BufferRef);
1573 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1574 typename TagT,
typename... PropTypes,
1576 detail::IsCxPropertyList<PropertyListT>::value &&
1577 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1578 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1580 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1582 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1584 const detail::code_location CodeLoc = detail::code_location::current())
1585 :
accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1586 adjustAccPropsInBuf(BufferRef);
1589 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1591 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1592 IsSameAsBuffer<T, Dims>::value &&
1593 (IsHostBuf || (IsGlobalBuf || IsConstantBuf))>>
1595 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1596 const property_list &PropertyList = {},
1597 const detail::code_location CodeLoc = detail::code_location::current())
1598 :
accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1600 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1601 typename... PropTypes,
1603 detail::IsCxPropertyList<PropertyListT>::value &&
1604 IsSameAsBuffer<T, Dims>::value &&
1605 (IsHostBuf || (IsGlobalBuf || IsConstantBuf))>>
1607 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1608 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1610 const detail::code_location CodeLoc = detail::code_location::current())
1611 :
accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1613 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1616 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1617 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1618 (IsGlobalBuf || IsConstantBuf)>>
1620 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1621 TagT,
const property_list &PropertyList = {},
1622 const detail::code_location CodeLoc = detail::code_location::current())
1623 :
accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1624 adjustAccPropsInBuf(BufferRef);
1627 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1628 typename TagT,
typename... PropTypes,
1630 detail::IsCxPropertyList<PropertyListT>::value &&
1631 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1632 (IsGlobalBuf || IsConstantBuf)>>
1634 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1636 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1638 const detail::code_location CodeLoc = detail::code_location::current())
1639 :
accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1640 adjustAccPropsInBuf(BufferRef);
1643 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1645 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1646 IsSameAsBuffer<T, Dims>::value &&
1647 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1649 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1650 range<Dimensions> AccessRange,
const property_list &PropertyList = {},
1651 const detail::code_location CodeLoc = detail::code_location::current())
1652 :
accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1655 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1656 typename... PropTypes,
1658 detail::IsCxPropertyList<PropertyListT>::value &&
1659 IsSameAsBuffer<T, Dims>::value &&
1660 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1662 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1663 range<Dimensions> AccessRange,
1664 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1666 const detail::code_location CodeLoc = detail::code_location::current())
1667 :
accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1670 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1673 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1674 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1675 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1677 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1678 range<Dimensions> AccessRange, TagT,
1679 const property_list &PropertyList = {},
1680 const detail::code_location CodeLoc = detail::code_location::current())
1681 :
accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1683 adjustAccPropsInBuf(BufferRef);
1686 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1687 typename TagT,
typename... PropTypes,
1689 detail::IsCxPropertyList<PropertyListT>::value &&
1690 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1691 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1693 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1694 range<Dimensions> AccessRange, TagT,
1695 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1697 const detail::code_location CodeLoc = detail::code_location::current())
1698 :
accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1700 adjustAccPropsInBuf(BufferRef);
1703 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1705 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1706 IsSameAsBuffer<T, Dims>::value &&
1707 (IsHostBuf || (IsGlobalBuf || IsConstantBuf))>>
1709 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1710 id<Dimensions> AccessOffset,
const property_list &PropertyList = {},
1711 const detail::code_location CodeLoc = detail::code_location::current())
1712 #ifdef __SYCL_DEVICE_ONLY__
1713 : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1717 : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1718 detail::convertToArrayOfN<3, 1>(AccessRange),
1719 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1720 getAdjustedMode(PropertyList),
1722 sizeof(DataT), IsPlaceH, BufferRef.OffsetInBytes,
1723 BufferRef.IsSubBuffer, PropertyList) {
1724 preScreenAccessor(BufferRef.size(), PropertyList);
1725 if (!AccessorBaseHost::isPlaceholder())
1727 if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1728 BufferRef.get_range()))
1729 throw sycl::invalid_object_error(
1730 "accessor with requested offset and range would exceed the bounds of "
1732 PI_ERROR_INVALID_VALUE);
1736 detail::AccessorBaseHost::impl.
get(),
1742 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1743 typename... PropTypes,
1745 detail::IsCxPropertyList<PropertyListT>::value &&
1746 IsSameAsBuffer<T, Dims>::value &&
1747 (IsHostBuf || (IsGlobalBuf || IsConstantBuf))>>
1749 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1750 id<Dimensions> AccessOffset,
1751 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1753 const detail::code_location CodeLoc = detail::code_location::current())
1754 #ifdef __SYCL_DEVICE_ONLY__
1755 : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1759 : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1760 detail::convertToArrayOfN<3, 1>(AccessRange),
1761 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1762 getAdjustedMode(PropertyList),
1764 sizeof(DataT), IsPlaceH, BufferRef.OffsetInBytes,
1765 BufferRef.IsSubBuffer, PropertyList) {
1766 preScreenAccessor(BufferRef.size(), PropertyList);
1767 if (!AccessorBaseHost::isPlaceholder())
1769 if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1770 BufferRef.get_range()))
1771 throw sycl::invalid_object_error(
1772 "accessor with requested offset and range would exceed the bounds of "
1774 PI_ERROR_INVALID_VALUE);
1778 detail::AccessorBaseHost::impl.
get(),
1784 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1787 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1788 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1789 (IsGlobalBuf || IsConstantBuf)>>
1791 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1792 id<Dimensions> AccessOffset, TagT,
const property_list &PropertyList = {},
1793 const detail::code_location CodeLoc = detail::code_location::current())
1794 :
accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
1795 adjustAccPropsInBuf(BufferRef);
1798 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1799 typename TagT,
typename... PropTypes,
1801 detail::IsCxPropertyList<PropertyListT>::value &&
1802 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1803 (IsGlobalBuf || IsConstantBuf)>>
1805 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1806 id<Dimensions> AccessOffset, TagT,
1807 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1809 const detail::code_location CodeLoc = detail::code_location::current())
1810 :
accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
1811 adjustAccPropsInBuf(BufferRef);
1814 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1816 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1817 IsSameAsBuffer<T, Dims>::value &&
1818 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1820 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1821 range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
1822 const property_list &PropertyList = {},
1823 const detail::code_location CodeLoc = detail::code_location::current())
1824 #ifdef __SYCL_DEVICE_ONLY__
1825 : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1826 (void)CommandGroupHandler;
1830 : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1831 detail::convertToArrayOfN<3, 1>(AccessRange),
1832 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1833 getAdjustedMode(PropertyList),
1835 sizeof(DataT), BufferRef.OffsetInBytes,
1836 BufferRef.IsSubBuffer, PropertyList) {
1837 preScreenAccessor(BufferRef.size(), PropertyList);
1838 if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1839 BufferRef.get_range()))
1840 throw sycl::invalid_object_error(
1841 "accessor with requested offset and range would exceed the bounds of "
1843 PI_ERROR_INVALID_VALUE);
1848 detail::AccessorBaseHost::impl.
get(),
1854 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1855 typename... PropTypes,
1857 detail::IsCxPropertyList<PropertyListT>::value &&
1858 IsSameAsBuffer<T, Dims>::value &&
1859 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1861 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1862 range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
1863 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1865 const detail::code_location CodeLoc = detail::code_location::current())
1866 #ifdef __SYCL_DEVICE_ONLY__
1867 : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
1868 (void)CommandGroupHandler;
1872 : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1873 detail::convertToArrayOfN<3, 1>(AccessRange),
1874 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1875 getAdjustedMode(PropertyList),
1877 sizeof(DataT), BufferRef.OffsetInBytes,
1878 BufferRef.IsSubBuffer, PropertyList) {
1879 preScreenAccessor(BufferRef.size(), PropertyList);
1880 if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
1881 BufferRef.get_range()))
1882 throw sycl::invalid_object_error(
1883 "accessor with requested offset and range would exceed the bounds of "
1885 PI_ERROR_INVALID_VALUE);
1890 detail::AccessorBaseHost::impl.
get(),
1896 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1899 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1900 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1901 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1903 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1904 range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
1905 const property_list &PropertyList = {},
1906 const detail::code_location CodeLoc = detail::code_location::current())
1907 :
accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1908 PropertyList, CodeLoc) {
1909 adjustAccPropsInBuf(BufferRef);
1912 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1913 typename TagT,
typename... PropTypes,
1915 detail::IsCxPropertyList<PropertyListT>::value &&
1916 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1917 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1919 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1920 range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
1921 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1923 const detail::code_location CodeLoc = detail::code_location::current())
1924 :
accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1925 PropertyList, CodeLoc) {
1926 adjustAccPropsInBuf(BufferRef);
1929 template <
typename... NewPropsT>
1932 ext::oneapi::accessor_property_list<NewPropsT...>> &Other,
1933 const detail::code_location CodeLoc = detail::code_location::current())
1934 #ifdef __SYCL_DEVICE_ONLY__
1937 : detail::AccessorBaseHost(Other)
1940 static_assert(detail::IsCxPropertyList<PropertyListT>::value,
1941 "Conversion is only available for accessor_property_list");
1943 PropertyListT::template areSameCompileTimeProperties<NewPropsT...>(),
1944 "Compile-time-constant properties must be the same");
1945 #ifndef __SYCL_DEVICE_ONLY__
1951 void swap(accessor &other) {
1952 std::swap(impl, other.impl);
1953 #ifndef __SYCL_DEVICE_ONLY__
1954 std::swap(MAccData, other.MAccData);
1958 bool is_placeholder()
const {
1959 #ifdef __SYCL_DEVICE_ONLY__
1962 return detail::AccessorBaseHost::isPlaceholder();
1966 size_t get_size()
const {
return getAccessRange().size() *
sizeof(DataT); }
1969 size_t get_count()
const {
return size(); }
1970 size_t size() const noexcept {
return getAccessRange().size(); }
1972 size_t byte_size() const noexcept {
return size() *
sizeof(DataT); }
1974 size_t max_size() const noexcept {
1978 bool empty() const noexcept {
return size() == 0; }
1980 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 0)>>
1981 range<Dimensions> get_range()
const {
1982 return detail::convertToArrayOfN<Dimensions, 1>(getAccessRange());
1985 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 0)>>
1986 id<Dimensions> get_offset()
const {
1990 "Accessor has no_offset property, get_offset() can not be used");
1991 return detail::convertToArrayOfN<Dimensions, 0>(getOffset());
1994 template <
int Dims =
Dimensions,
typename RefT = RefType,
1997 operator reference()
const {
1999 return *(getQualifiedPtr() + LinearIndex);
2005 reference
operator[](id<Dimensions> Index)
const {
2007 return getQualifiedPtr()[LinearIndex];
2010 template <
int Dims = Dimensions>
2013 #ifdef __ENABLE_USM_ADDR_SPACE__
2020 return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2021 getQualifiedPtr() + LinearIndex));
2024 template <
int Dims = Dimensions>
2027 operator[](id<Dimensions> Index)
const {
2029 return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2030 getQualifiedPtr() + LinearIndex));
2033 template <
int Dims = Dimensions>
2037 const size_t LinearIndex =
getLinearIndex(id<AdjustedDim>(Index));
2038 return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2039 getQualifiedPtr() + LinearIndex));
2041 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 1)>>
2043 return AccessorSubscript<Dims - 1>(*
this, Index);
2048 access::target::host_buffer>>
2049 #if SYCL_LANGUAGE_VERSION >= 202001
2050 std::add_pointer_t<value_type> get_pointer() const noexcept
2052 DataT *get_pointer() const
2055 return getPointerAdjusted();
2060 typename = detail::enable_if_t<AccessTarget_ == access::target::device>>
2061 global_ptr<DataT> get_pointer()
const {
2062 return global_ptr<DataT>(getPointerAdjusted());
2067 access::target::constant_buffer>>
2068 constant_ptr<DataT> get_pointer()
const {
2069 return constant_ptr<DataT>(getPointerAdjusted());
2075 template <
typename Property>
2077 !ext::oneapi::is_compile_time_property<Property>::value,
bool>
2079 #ifndef __SYCL_DEVICE_ONLY__
2080 return getPropList().template has_property<Property>();
2089 template <
typename Property,
2091 !ext::oneapi::is_compile_time_property<Property>::value>>
2093 #ifndef __SYCL_DEVICE_ONLY__
2094 return getPropList().template get_property<Property>();
2100 template <
typename Property>
2103 ext::oneapi::is_compile_time_property<Property>::value> * = 0) {
2104 return PropertyListT::template has_property<Property>();
2107 template <
typename Property>
2110 ext::oneapi::is_compile_time_property<Property>::value> * = 0) {
2111 return PropertyListT::template get_property<Property>();
2114 bool operator==(
const accessor &Rhs)
const {
return impl == Rhs.impl; }
2115 bool operator!=(
const accessor &Rhs)
const {
return !(*
this == Rhs); }
2117 iterator begin() const noexcept {
2118 return iterator::getBegin(
2120 detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
2124 iterator end() const noexcept {
2125 return iterator::getEnd(
2127 detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
2131 const_iterator cbegin() const noexcept {
2132 return const_iterator::getBegin(
2134 detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
2138 const_iterator cend() const noexcept {
2139 return const_iterator::getEnd(
2141 detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
2145 reverse_iterator rbegin() const noexcept {
return reverse_iterator(end()); }
2146 reverse_iterator rend() const noexcept {
return reverse_iterator(begin()); }
2148 const_reverse_iterator crbegin() const noexcept {
2149 return const_reverse_iterator(cend());
2151 const_reverse_iterator crend() const noexcept {
2152 return const_reverse_iterator(cbegin());
2156 #ifdef __SYCL_DEVICE_ONLY__
2157 size_t getTotalOffset() const noexcept {
2158 size_t TotalOffset = 0;
2159 detail::dim_loop<Dimensions>([&,
this](
size_t I) {
2160 TotalOffset = TotalOffset * impl.MemRange[I];
2163 TotalOffset += impl.Offset[I];
2176 auto getPointerAdjusted() const noexcept {
2177 #ifdef __SYCL_DEVICE_ONLY__
2178 return getQualifiedPtr() - getTotalOffset();
2180 return getQualifiedPtr();
2184 void preScreenAccessor(
const size_t elemInBuffer,
2185 const PropertyListT &PropertyList) {
2187 if (!IsHostBuf && elemInBuffer == 0)
2188 throw sycl::invalid_object_error(
2189 "SYCL buffer size is zero. To create a device accessor, SYCL "
2190 "buffer size must be greater than zero.",
2191 PI_ERROR_INVALID_VALUE);
2194 if (PropertyList.template has_property<property::no_init>() &&
2196 throw sycl::invalid_object_error(
2197 "accessor would cannot be both read_only and no_init",
2198 PI_ERROR_INVALID_VALUE);
2202 template <
typename BufT,
typename... PropTypes>
2203 void adjustAccPropsInBuf(BufT &Buffer) {
2209 property_list PropList{
2211 Buffer.addOrReplaceAccessorProperties(PropList);
2213 deleteAccPropsFromBuf(Buffer);
2217 template <
typename BufT>
void deleteAccPropsFromBuf(BufT &Buffer) {
2218 Buffer.deleteAccProps(
2223 template <
typename DataT,
int Dimensions,
typename AllocatorT>
2226 access::placeholder::true_t>;
2228 template <
typename DataT,
int Dimensions,
typename AllocatorT,
2233 access::placeholder::true_t,
2236 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1>
2239 detail::deduceAccessTarget<Type1, Type1>(target::device),
2240 access::placeholder::true_t>;
2242 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2247 detail::deduceAccessTarget<Type1, Type1>(target::device),
2248 access::placeholder::true_t,
2251 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2255 detail::deduceAccessTarget<Type1, Type2>(target::device),
2256 access::placeholder::true_t>;
2258 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2259 typename Type2,
typename... PropsT>
2263 detail::deduceAccessTarget<Type1, Type2>(target::device),
2264 access::placeholder::true_t,
2267 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2268 typename Type2,
typename Type3>
2271 detail::deduceAccessTarget<Type2, Type3>(target::device),
2272 access::placeholder::true_t>;
2274 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2275 typename Type2,
typename Type3,
typename... PropsT>
2279 detail::deduceAccessTarget<Type2, Type3>(target::device),
2280 access::placeholder::true_t,
2283 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2284 typename Type2,
typename Type3,
typename Type4>
2287 detail::deduceAccessTarget<Type3, Type4>(target::device),
2288 access::placeholder::true_t>;
2290 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2291 typename Type2,
typename Type3,
typename Type4,
typename... PropsT>
2295 detail::deduceAccessTarget<Type3, Type4>(target::device),
2296 access::placeholder::true_t,
2299 template <
typename DataT,
int Dimensions,
typename AllocatorT>
2302 access::placeholder::false_t>;
2304 template <
typename DataT,
int Dimensions,
typename AllocatorT,
2309 access::placeholder::false_t,
2312 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1>
2315 detail::deduceAccessTarget<Type1, Type1>(target::device),
2316 access::placeholder::false_t>;
2318 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2323 detail::deduceAccessTarget<Type1, Type1>(target::device),
2324 access::placeholder::false_t,
2327 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2331 detail::deduceAccessTarget<Type1, Type2>(target::device),
2332 access::placeholder::false_t>;
2334 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2335 typename Type2,
typename... PropsT>
2339 detail::deduceAccessTarget<Type1, Type2>(target::device),
2340 access::placeholder::false_t,
2343 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2344 typename Type2,
typename Type3>
2347 detail::deduceAccessTarget<Type2, Type3>(target::device),
2348 access::placeholder::false_t>;
2350 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2351 typename Type2,
typename Type3,
typename... PropsT>
2355 detail::deduceAccessTarget<Type2, Type3>(target::device),
2356 access::placeholder::false_t,
2359 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2360 typename Type2,
typename Type3,
typename Type4>
2364 detail::deduceAccessTarget<Type3, Type4>(target::device),
2365 access::placeholder::false_t>;
2367 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2368 typename Type2,
typename Type3,
typename Type4,
typename... PropsT>
2372 detail::deduceAccessTarget<Type3, Type4>(target::device),
2373 access::placeholder::false_t,
2382 #ifndef __SYCL_DEVICE_ONLY__
2386 access::target::local, IsPlaceholder> {
2394 using AccessorCommonT::AS;
2395 using AccessorCommonT::IsAccessAnyWrite;
2407 #ifdef __SYCL_DEVICE_ONLY__
2410 sycl::range<AdjustedDim> &getSize() {
return impl.
MemRange; }
2411 const sycl::range<AdjustedDim> &getSize()
const {
return impl.
MemRange; }
2413 void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
2414 range<AdjustedDim>, id<AdjustedDim>) {
2416 detail::dim_loop<AdjustedDim>(
2417 [&,
this](
size_t I) { getSize()[I] = AccessRange[I]; });
2422 local_accessor_base()
2423 : impl(detail::InitializedVal<AdjustedDim, range>::template
get<0>()) {}
2426 ConcreteASPtrType getQualifiedPtr()
const {
return MData; }
2428 ConcreteASPtrType MData;
2433 : detail::LocalAccessorBaseHost{
sycl::
range<3>{0, 0, 0},
2434 0,
sizeof(DataT)} {}
2438 : detail::LocalAccessorBaseHost{Impl} {}
2442 using detail::LocalAccessorBaseHost::getSize;
2445 return reinterpret_cast<PtrType>(LocalAccessorBaseHost::getPtr());
2448 void *
getPtr() {
return detail::LocalAccessorBaseHost::getPtr(); }
2449 void *
getPtr()
const {
return detail::LocalAccessorBaseHost::getPtr(); }
2451 return detail::LocalAccessorBaseHost::getSize();
2458 const auto *this_const =
this;
2460 (void)this_const->getSize();
2462 (void)this_const->getPtr();
2471 for (
int I = 0; I < Dims; ++I)
2472 Result = Result * getSize()[I] + Id[I];
2484 template <
int Dims = Dimensions,
typename = detail::enable_if_t<Dims == 0>>
2486 detail::code_location::current())
2487 #ifdef __SYCL_DEVICE_ONLY__
2490 : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim,
sizeof(DataT)) {
2498 typename = detail::enable_if_t<Dims == 0>>
2501 detail::code_location::current())
2502 #ifdef __SYCL_DEVICE_ONLY__
2507 : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim,
sizeof(DataT),
2515 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 0)>>
2519 #ifdef __SYCL_DEVICE_ONLY__
2520 : impl(AllocationSize){}
2522 : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2523 AdjustedDim,
sizeof(DataT)) {
2535 detail::code_location::current())
2536 #ifdef __SYCL_DEVICE_ONLY__
2537 : impl(AllocationSize) {
2541 : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2542 AdjustedDim,
sizeof(DataT), propList) {
2549 size_t get_size()
const {
return getSize().size() *
sizeof(DataT); }
2552 size_t get_count()
const {
return size(); }
2553 size_t size() const noexcept {
return getSize().size(); }
2555 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 0)>>
2557 return detail::convertToArrayOfN<Dims, 1>(getSize());
2563 return *getQualifiedPtr();
2570 return getQualifiedPtr()[LinearIndex];
2576 return getQualifiedPtr()[Index];
2579 template <
int Dims = Dimensions>
2581 Dims == 0 &&
AccessMode == access::mode::atomic, atomic<DataT, AS>>()
2583 return atomic<DataT, AS>(
2587 template <
int Dims = Dimensions>
2593 getQualifiedPtr() + LinearIndex));
2596 template <
int Dims = Dimensions>
2601 getQualifiedPtr() + Index));
2604 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 1)>>
2605 typename AccessorCommonT::template AccessorSubscript<
2617 return impl == Rhs.
impl;
2620 return !(*
this == Rhs);
2631 accessor<DataT, Dimensions, AccessMode, access::target::local,
2638 using local_acc::local_acc;
2640 #ifdef __SYCL_DEVICE_ONLY__
2648 local_acc::__init(Ptr, AccessRange,
range,
id);
2655 range>::template get<0>();
2664 template <
typename DataT,
int Dimensions = 1>
2667 access::placeholder::false_t>,
2672 access::placeholder::false_t>;
2675 using local_acc::local_acc;
2677 #ifdef __SYCL_DEVICE_ONLY__
2681 void __init(
typename local_acc::ConcreteASPtrType Ptr,
2685 local_acc::__init(Ptr, AccessRange,
range,
id);
2691 local_acc::impl = detail::InitializedVal<local_acc::AdjustedDim,
2692 range>::template get<0>();
2700 using value_type = DataT;
2701 using iterator = value_type *;
2702 using const_iterator =
const value_type *;
2703 using reverse_iterator = std::reverse_iterator<iterator>;
2704 using const_reverse_iterator = std::reverse_iterator<const_iterator>;
2705 using difference_type =
2706 typename std::iterator_traits<iterator>::difference_type;
2708 template <access::decorated IsDecorated>
2709 using accessor_ptr = local_ptr<value_type, IsDecorated>;
2711 void swap(local_accessor &other) { std::swap(this->impl, other.impl); }
2713 size_t byte_size() const noexcept {
return this->size() *
sizeof(DataT); }
2715 size_t max_size() const noexcept {
2719 bool empty() const noexcept {
return this->size() == 0; }
2721 iterator begin() const noexcept {
2724 iterator end() const noexcept {
return begin() + this->size(); }
2726 const_iterator cbegin() const noexcept {
return const_iterator(begin()); }
2727 const_iterator cend() const noexcept {
return const_iterator(end()); }
2729 reverse_iterator rbegin() const noexcept {
return reverse_iterator(end()); }
2730 reverse_iterator rend() const noexcept {
return reverse_iterator(begin()); }
2732 const_reverse_iterator crbegin() const noexcept {
2733 return const_reverse_iterator(end());
2735 const_reverse_iterator crend() const noexcept {
2736 return const_reverse_iterator(begin());
2739 template <
typename Property>
bool has_property() const noexcept {
2740 #ifndef __SYCL_DEVICE_ONLY__
2741 return this->getPropList().template has_property<Property>();
2747 template <
typename Property> Property
get_property()
const {
2748 #ifndef __SYCL_DEVICE_ONLY__
2749 return this->getPropList().template get_property<Property>();
2766 access::target::image, IsPlaceholder>,
2768 accessor<DataT, Dimensions, AccessMode, access::target::image,
2776 template <
typename AllocatorT>
2777 accessor(sycl::image<Dimensions, AllocatorT> &Image,
2781 Image, CommandGroupHandler, Image.getElementSize()) {
2782 #ifndef __SYCL_DEVICE_ONLY__
2784 access::target::image);
2788 template <
typename AllocatorT>
2789 accessor(sycl::image<Dimensions, AllocatorT> &Image,
2793 Image, CommandGroupHandler, Image.getElementSize()) {
2795 #ifndef __SYCL_DEVICE_ONLY__
2797 access::target::image);
2800 #ifdef __SYCL_DEVICE_ONLY__
2804 access::target::image>::type;
2808 void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
2811 void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
2831 access::target::host_image, IsPlaceholder>,
2833 accessor<DataT, Dimensions, AccessMode, access::target::host_image,
2836 template <
typename AllocatorT>
2840 Image, Image.getElementSize()) {}
2842 template <
typename AllocatorT>
2847 Image, Image.getElementSize()) {
2865 access::target::image, IsPlaceholder>,
2867 accessor<DataT, Dimensions, AccessMode, access::target::image_array,
2869 #ifdef __SYCL_DEVICE_ONLY__
2873 access::target::image>::type;
2877 void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
2880 void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
2887 template <
typename AllocatorT>
2888 accessor(sycl::image<Dimensions + 1, AllocatorT> &Image,
2892 Image, CommandGroupHandler, Image.getElementSize()) {
2893 #ifndef __SYCL_DEVICE_ONLY__
2895 access::target::image_array);
2899 template <
typename AllocatorT>
2900 accessor(sycl::image<Dimensions + 1, AllocatorT> &Image,
2904 Image, CommandGroupHandler, Image.getElementSize()) {
2906 #ifndef __SYCL_DEVICE_ONLY__
2908 access::target::image_array);
2919 template <
typename DataT,
int Dimensions = 1,
2922 :
public accessor<DataT, Dimensions, AccessMode, target::host_buffer,
2923 access::placeholder::false_t>,
2925 host_accessor<DataT, Dimensions, AccessMode>> {
2928 access::placeholder::false_t>;
2932 template <
typename T,
int Dims>
2935 (Dims == Dimensions)> {};
2939 access::placeholder::false_t>::ConcreteASPtrType Ptr,
2942 AccessorT::__init(Ptr, AccessRange, MemRange, Offset);
2971 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2973 std::is_same<T, DataT>::value && Dims == 0>>
2978 : AccessorT(BufferRef, PropertyList, CodeLoc) {}
2980 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2981 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
2986 : AccessorT(BufferRef, PropertyList, CodeLoc) {}
2988 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2989 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
2996 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2997 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3002 : AccessorT(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
3004 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3005 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3010 :
host_accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
3012 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3013 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3018 : AccessorT(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
3020 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3021 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3026 :
host_accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
3028 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3029 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3034 : AccessorT(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
3037 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3038 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3044 :
host_accessor(BufferRef, CommandGroupHandler, AccessRange, {},
3045 PropertyList, CodeLoc) {}
3047 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3048 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3053 : AccessorT(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
3056 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3057 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3063 :
host_accessor(BufferRef, AccessRange, AccessOffset, PropertyList,
3066 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3067 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3073 : AccessorT(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
3074 PropertyList, CodeLoc) {}
3076 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3077 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3083 :
host_accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
3084 PropertyList, CodeLoc) {}
3087 template <
typename DataT,
int Dimensions,
typename AllocatorT>
3091 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1>
3094 detail::deduceAccessMode<Type1, Type1>()>;
3096 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
3100 detail::deduceAccessMode<Type1, Type2>()>;
3102 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
3103 typename Type2,
typename Type3>
3106 detail::deduceAccessMode<Type2, Type3>()>;
3108 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
3109 typename Type2,
typename Type3,
typename Type4>
3112 detail::deduceAccessMode<Type3, Type4>()>;
3114 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
3115 typename Type2,
typename Type3,
typename Type4,
typename Type5>
3118 detail::deduceAccessMode<Type4, Type5>()>;
3132 #ifdef __SYCL_DEVICE_ONLY__
3140 return hash<decltype(AccImplPtr)>()(AccImplPtr);
The file contains implementation of accessor iterator class.
accessor(sycl::image< Dimensions, AllocatorT > &Image)
accessor(sycl::image< Dimensions, AllocatorT > &Image, const property_list &propList)
Defines a shared array that can be used by kernels in queues.
AccessorBaseHost(const AccessorImplPtr &Impl)
decltype(Obj::impl) friend getSyclObjImpl(const Obj &SyclObject)
AccessorImplDevice(id< Dims > Offset, range< Dims > AccessRange, range< Dims > MemoryRange)
bool operator==(const AccessorImplDevice &Rhs) const
range< Dims > AccessRange
AccessorImplDevice()=default
range< Dims > AccessRange
LocalAccessorBaseDevice(sycl::range< Dims > Size)
bool operator==(const LocalAccessorBaseDevice &Rhs) const
decltype(Obj::impl) friend getSyclObjImpl(const Obj &SyclObject)
LocalAccessorBaseHost(const LocalAccessorImplPtr &Impl)
LocalAccessorImplPtr impl
decltype(check(T())) type
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
range< Dims > get_range() const
DataT read(const CoordT &Coords) const
size_t size() const noexcept
void write(const CoordT &Coords, const DataT &Color) const
DataT read(const CoordT &Coords, const sampler &Smpl) const
__image_array_slice__(accessor< DataT, Dimensions, AccessMode, access::target::image_array, IsPlaceholder, ext::oneapi::accessor_property_list<>> BaseAcc, size_t Idx)
detail::enable_if_t< CurDims==1 &&IsAccessAtomic, atomic< DataT, AS > > operator[](size_t Index) const
AccessorSubscript(AccType Accessor, size_t Index)
AccessorSubscript(AccType Accessor, id< Dims > IDs)
const DataT & ConstRefType
detail::const_if_const_AS< AS, DataT > * PtrType
detail::const_if_const_AS< AS, DataT > & RefType
size_t size() const noexcept
image_accessor(image< Dims, AllocatorT > &ImageRef, int ImageElementSize)
bool operator!=(const image_accessor &Rhs) const
const DataT & const_reference
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
void write(const CoordT &Coords, const DataT &Color) const
range< Dims > get_range() const
image_accessor(const AccessorImplPtr &Impl)
bool operator==(const image_accessor &Rhs) const
DataT read(const CoordT &Coords) const
DataT read(const CoordT &Coords, const sampler &Smpl) const
image_accessor(image< Dims, AllocatorT > &ImageRef, handler &CommandGroupHandlerRef, int ImageElementSize)
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Command group handler class.
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
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)
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, mode_tag_t< AccessMode >, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, mode_tag_t< AccessMode >, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, mode_tag_t< AccessMode >, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, mode_tag_t< AccessMode >, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, mode_tag_t< AccessMode >, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
host_accessor(buffer< T, Dims, AllocatorT > &BufferRef, handler &CommandGroupHandler, range< Dimensions > AccessRange, id< Dimensions > AccessOffset, mode_tag_t< AccessMode >, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
host_accessor(buffer< T, 1, AllocatorT > &BufferRef, const property_list &PropertyList={}, const detail::code_location CodeLoc=detail::code_location::current())
Defines a shared image data.
range< Dimensions > get_range() const
size_t size() const noexcept
detail::const_if_const_AS< AS, DataT > & RefType
const range< 3 > & getSize() const
bool operator!=(const local_accessor_base &Rhs) const
AccessorCommonT::template AccessorSubscript< Dims - 1, local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder > > operator[](size_t Index) const
local_accessor_base(range< Dimensions > AllocationSize, handler &, const detail::code_location CodeLoc=detail::code_location::current())
range< Dims > get_range() const
local_accessor_base(handler &, const detail::code_location CodeLoc=detail::code_location::current())
typename detail::DecoratedType< DataT, AS >::type * ConcreteASPtrType
typename AccessorCommonT::template AccessorSubscript< Dims, local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder > > AccessorSubscript
size_t getLinearIndex(id< Dims > Id) const
local_accessor_base(range< Dimensions > AllocationSize, handler &, const property_list &propList, const detail::code_location CodeLoc=detail::code_location::current())
local_accessor_base(handler &, const property_list &propList, const detail::code_location CodeLoc=detail::code_location::current())
bool operator==(const local_accessor_base &Rhs) const
RefType operator[](size_t Index) const
local_ptr< DataT > get_pointer() const
const DataT & const_reference
__SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") size_t get_count() const
detail::enable_if_t< Dims==1 &&AccessMode==access::mode::atomic, atomic< DataT, AS > > operator[](size_t Index) const
detail::const_if_const_AS< AS, DataT > * PtrType
size_t size() const noexcept
PtrType getQualifiedPtr() const
local_accessor_base(const detail::LocalAccessorImplPtr &Impl)
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Objects of the property_list class are containers for the SYCL properties.
Defines the iteration domain of either a single work-group in a parallel dispatch,...
#define __SYCL_SPECIAL_CLASS
#define __SYCL_INLINE_VER_NAMESPACE(X)
#define __SYCL2020_DEPRECATED(message)
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor accessor(buffer< DataT, Dimensions, AllocatorT >) -> accessor< DataT, Dimensions, access::mode::read_write, target::device, access::placeholder::true_t >
Buffer accessor.
constexpr tuple_element< I, tuple< Types... > >::type & get(sycl::detail::tuple< Types... > &Arg) noexcept
void addHostAccessorAndWait(AccessorImplHost *Req)
device getDeviceFromHandler(handler &CommandGroupHandlerRef)
size_t getLinearIndex(const T< Dims > &Index, const U< Dims > &Range)
void imageWriteHostImpl(const CoordT &Coords, const WriteDataT &Color, id< 3 > ImgPitch, uint8_t ElementSize, image_channel_type ImgChannelType, image_channel_order ImgChannelOrder, void *BasePtr)
typename std::conditional< B, T, F >::type conditional_t
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
typename std::is_same< ext::oneapi::accessor_property_list<>, T > IsRunTimePropertyListT
void constructorNotification(void *BufferObj, void *AccessorObj, access::target Target, access::mode Mode, const code_location &CodeLoc)
std::integral_constant< bool, V > bool_constant
T createSyclObjFromImpl(decltype(T::impl) ImplObj)
std::shared_ptr< LocalAccessorImplHost > LocalAccessorImplPtr
void dim_loop_impl(std::integer_sequence< size_t, Inds... >, F &&f)
constexpr access::target deduceAccessTarget(access::target defaultTarget)
typename std::is_base_of< PropertyListBase, T > IsPropertyListT
constexpr access::mode deduceAccessMode()
void associateWithHandler(handler &, AccessorBaseHost *, access::target)
typename std::enable_if< B, T >::type enable_if_t
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
std::shared_ptr< AccessorImplHost > AccessorImplPtr
static constexpr bool has_property()
static constexpr auto get_property()
constexpr std::enable_if_t< detail::IsCompileTimeProperty< PropertyT >::value, bool > operator!=(const property_value< PropertyT, A... > &, const property_value< PropertyT, B... > &)
T & operator[](std::ptrdiff_t idx) const noexcept
constexpr buffer_location_key::value_t< N > buffer_location
constexpr property::no_offset::instance no_offset
host_accessor(buffer< DataT, Dimensions, AllocatorT >, Type1, Type2, Type3, Type4, Type5) -> host_accessor< DataT, Dimensions, detail::deduceAccessMode< Type4, Type5 >()>
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS Dimensions
constexpr mode_tag_t< access_mode::read_write > read_write
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS IsPlaceholder
accessor(buffer< DataT, Dimensions, AllocatorT >, handler &, Type1, Type2, Type3, Type4, const ext::oneapi::accessor_property_list< PropsT... > &) -> accessor< DataT, Dimensions, detail::deduceAccessMode< Type3, Type4 >(), detail::deduceAccessTarget< Type3, Type4 >(target::device), access::placeholder::false_t, ext::oneapi::accessor_property_list< PropsT... >>
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS AccessMode
---— Error handling, matching OpenCL plugin semantics.
simd< _Tp, _Abi > max(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
sycl::accessor< DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder > AccType
size_t operator()(const AccType &A) const
sycl::range< 3 > MMemoryRange
sycl::range< 3 > MAccessRange
AccHostDataT(const sycl::id< 3 > &Offset, const sycl::range< 3 > &Range, const sycl::range< 3 > &MemoryRange, void *Data=nullptr)
bool operator==(const Slab &Lhs, const Slab &Rhs)