46 #ifdef __SYCL_DEVICE_ONLY__
58 #include <type_traits>
234 inline namespace _V1 {
236 namespace ext::intel::esimd::detail {
238 class AccessorPrivateProxy;
271 void *ImageObj,
void *AccessorObj,
272 const std::optional<image_target> &Target,
access::mode Mode,
273 const void *Type, uint32_t ElemSize,
const code_location &CodeLoc);
276 void *ImageObj,
void *AccessorObj,
277 const std::optional<image_target> &Target,
const void *Type,
280 template <
typename T>
283 template <
typename T>
285 typename std::is_same<ext::oneapi::accessor_property_list<>, T>;
288 constexpr
static bool value =
false;
291 template <
typename... Props>
297 constexpr
static bool value =
false;
302 template <
typename BufferT>
304 return std::min(Buffer.size(),
size_t{1});
316 constexpr
static bool IsHostBuf = AccessTarget == access::target::host_buffer;
334 AccessTarget == access::target::global_buffer;
337 AccessTarget == access::target::constant_buffer;
346 static constexpr
bool IsConst = std::is_const_v<DataT>;
359 template <
int SubDims,
371 : MIDs(IDs), MAccessor(Accessor) {}
379 template <
int CurDims = SubDims,
typename = std::enable_if_t<(CurDims > 1)>>
381 MIDs[Dims - CurDims] = Index;
385 template <
int CurDims = SubDims,
389 MIDs[Dims - CurDims] = Index;
390 return MAccessor[MIDs];
393 template <
int CurDims = SubDims>
394 typename std::enable_if_t<CurDims == 1 && IsAccessAtomic, atomic<DataT, AS>>
396 MIDs[Dims - CurDims] = Index;
397 return MAccessor[MIDs];
403 if constexpr (std::is_const_v<DataT>)
409 template <
typename MayBeTag1,
typename MayBeTag2>
424 if constexpr (std::is_same_v<
427 access::target::constant_buffer>> ||
431 access::target::constant_buffer>>) {
435 if constexpr (std::is_same_v<MayBeTag1,
438 std::is_same_v<MayBeTag2,
444 if constexpr (std::is_same_v<MayBeTag1,
447 std::is_same_v<MayBeTag2,
456 template <
typename MayBeTag1,
typename MayBeTag2>
458 if constexpr (std::is_same_v<
461 access::target::constant_buffer>> ||
465 access::target::constant_buffer>>) {
466 return access::target::constant_buffer;
485 return defaultTarget;
526 class AccessorImplHost;
542 int ElemSize,
int OffsetInBytes = 0,
543 bool IsSubBuffer =
false,
548 int ElemSize,
bool IsPlaceH,
int OffsetInBytes = 0,
549 bool IsSubBuffer =
false,
552 AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
554 int ElemSize,
size_t OffsetInBytes = 0,
555 bool IsSubBuffer =
false,
556 const property_list &PropertyList = {});
558 AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
560 int ElemSize,
bool IsPlaceH,
size_t OffsetInBytes = 0,
561 bool IsSubBuffer =
false,
562 const property_list &PropertyList = {});
566 range<3> &getAccessRange();
567 range<3> &getMemoryRange();
568 void *getPtr() noexcept;
569 unsigned int getElemSize() const;
571 const
id<3> &getOffset() const;
572 const range<3> &getAccessRange() const;
573 const range<3> &getMemoryRange() const;
574 void *getPtr() const noexcept;
575 bool isPlaceholder() const;
576 bool isMemoryObjectUsedByGraph() const;
578 detail::AccHostDataT &getAccData();
580 const property_list &getPropList() const;
582 void *getMemoryObject() const;
597 friend class
sycl::ext::intel::esimd::
detail::AccessorPrivateProxy;
608 LocalAccessorBaseHost(
sycl::range<3> Size,
int Dims,
int ElemSize,
613 void *getPtr()
const;
615 int getElementSize();
631 std::shared_ptr<UnsampledImageAccessorImplHost>;
633 std::shared_ptr<SampledImageAccessorImplHost>;
647 void *SYCLMemObject,
int Dims,
int ElemSize,
652 void *getMemoryObject()
const;
653 detail::AccHostDataT &getAccData();
655 void *getPtr()
const;
656 int getNumOfDims()
const;
657 int getElementSize()
const;
658 id<3> getPitch()
const;
675 const auto *this_const =
this;
677 (void)this_const->getSize();
679 (void)this_const->getPtr();
683 #ifndef __SYCL_DEVICE_ONLY__
687 template <
typename DataT,
typename CoordT>
688 DataT
read(
const CoordT &Coords)
const noexcept {
692 return imageReadSamplerHostImpl<CoordT, DataT>(
693 Coords, Smpl, getSize(), getPitch(), getChannelType(),
694 getChannelOrder(), getPtr(), getElementSize());
700 template <
typename DataT,
typename CoordT>
701 void write(
const CoordT &Coords,
const DataT &Color)
const {
703 getChannelType(), getChannelOrder(), getPtr());
714 SampledImageAccessorBaseHost(
sycl::range<3> Size,
void *SYCLMemObject,
715 int Dims,
int ElemSize,
id<3> Pitch,
721 void *getMemoryObject()
const;
722 detail::AccHostDataT &getAccData();
724 void *getPtr()
const;
725 int getNumOfDims()
const;
726 int getElementSize()
const;
727 id<3> getPitch()
const;
745 const auto *this_const =
this;
747 (void)this_const->getSize();
749 (void)this_const->getPtr();
753 #ifndef __SYCL_DEVICE_ONLY__
757 template <
typename DataT,
typename CoordT>
758 DataT
read(
const CoordT &Coords)
const {
759 return imageReadSamplerHostImpl<CoordT, DataT>(
760 Coords, getSampler(), getSize(), getPitch(), getChannelType(),
761 getChannelOrder(), getPtr(), getElementSize());
784 constexpr
static bool value = std::is_same_v<T, int>;
787 constexpr
static bool value = std::is_same_v<T, int2>;
790 constexpr
static bool value = std::is_same_v<T, int4>;
795 constexpr
static bool value = std::is_same_v<T, float>;
798 constexpr
static bool value = std::is_same_v<T, float2>;
801 constexpr
static bool value = std::is_same_v<T, float4>;
812 #ifndef __SYCL_DEVICE_ONLY__
822 OCLImageTy MImageObj;
828 void imageAccessorInit(OCLImageTy Image) { MImageObj = Image; }
832 template <
typename T1,
int T2, access::mode T3, access::placeholder T4>
835 constexpr
static bool IsHostImageAcc =
840 constexpr
static bool IsImageArrayAcc =
843 constexpr
static bool IsImageAccessWriteOnly =
847 constexpr
static bool IsImageAccessAnyWrite =
850 constexpr
static bool IsImageAccessReadOnly =
853 constexpr
static bool IsImageAccessAnyRead =
860 "The data type of an image accessor must be only cl_int4, "
861 "cl_uint4, cl_float4 or cl_half4 from SYCL namespace");
863 static_assert(IsImageAcc || IsHostImageAcc || IsImageArrayAcc,
864 "Expected image type");
867 "Expected false as Placeholder value for image accessor.");
870 ((IsImageAcc || IsImageArrayAcc) &&
871 (IsImageAccessWriteOnly || IsImageAccessReadOnly)) ||
872 (IsHostImageAcc && (IsImageAccessAnyWrite || IsImageAccessAnyRead)),
873 "Access modes can be only read/write/discard_write for image/image_array "
874 "target accessor, or they can be only "
875 "read/write/discard_write/read_write for host_image target accessor.");
878 "Dimensions can be 1/2/3 for image accessor.");
880 #ifdef __SYCL_DEVICE_ONLY__
883 return __invoke_ImageQuerySize<sycl::vec<int, Dimensions>, OCLImageTy>(
887 size_t getElementSize()
const {
888 int ChannelType = __invoke_ImageQueryFormat<int, OCLImageTy>(MImageObj);
889 int ChannelOrder = __invoke_ImageQueryOrder<int, OCLImageTy>(MImageObj);
890 int ElementSize = getSPIRVElementSize(ChannelType, ChannelOrder);
900 "image::getRangeInternal() is not implemented for host");
906 #ifndef __SYCL_DEVICE_ONLY__
909 #endif // __SYCL_DEVICE_ONLY__
912 friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
914 #ifdef __SYCL_DEVICE_ONLY__
915 const OCLImageTy getNativeImageObj()
const {
return MImageObj; }
916 #endif // __SYCL_DEVICE_ONLY__
925 #ifdef __SYCL_DEVICE_ONLY__
935 typename = std::enable_if_t<(Dims > 0 && Dims <= 3) && IsHostImageAcc>>
937 #ifdef __SYCL_DEVICE_ONLY__
940 (void)ImageElementSize;
946 detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
947 detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
950 MImageCount(ImageRef.size()),
951 MImgChannelOrder(ImageRef.getChannelOrder()),
952 MImgChannelType(ImageRef.getChannelType()) {
961 template <
typename AllocatorT,
int Dims =
Dimensions,
962 typename = std::enable_if_t<(Dims > 0 && Dims <= 3) && IsImageAcc>>
964 handler &CommandGroupHandlerRef,
int ImageElementSize)
965 #ifdef __SYCL_DEVICE_ONLY__
968 (void)CommandGroupHandlerRef;
969 (void)ImageElementSize;
975 detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
976 detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
979 MImageCount(ImageRef.size()),
980 MImgChannelOrder(ImageRef.getChannelOrder()),
981 MImgChannelType(ImageRef.getChannelType()) {
984 if (!Device.
has(aspect::ext_intel_legacy_image))
986 "SYCL 1.2.1 images are not supported by this device.",
987 PI_ERROR_INVALID_OPERATION);
995 #ifndef __SYCL_DEVICE_ONLY__
1014 #ifdef __SYCL_DEVICE_ONLY__
1017 size_t get_count()
const {
return size(); }
1018 size_t size() const noexcept {
return get_range<Dimensions>().size(); }
1020 template <
int Dims = Dimensions,
typename = std::enable_if_t<Dims == 1>>
1022 int Range = getRangeInternal();
1025 template <
int Dims = Dimensions,
typename = std::enable_if_t<Dims == 2>>
1026 range<2> get_range()
const {
1027 int2 Range = getRangeInternal();
1028 return range<2>(Range[0], Range[1]);
1030 template <
int Dims = Dimensions,
typename = std::enable_if_t<Dims == 3>>
1031 range<3> get_range()
const {
1032 int3 Range = getRangeInternal();
1033 return range<3>(Range[0], Range[1], Range[2]);
1038 size_t get_count()
const {
return size(); };
1039 size_t size() const noexcept {
return MImageCount; };
1041 template <
int Dims = Dimensions,
typename = std::enable_if_t<(Dims > 0)>>
1043 return detail::convertToArrayOfN<Dims, 1>(getAccessRange());
1052 template <
typename CoordT,
int Dims =
Dimensions,
1053 typename = std::enable_if_t<
1056 ((IsImageAcc && IsImageAccessReadOnly) ||
1057 (IsHostImageAcc && IsImageAccessAnyRead))>>
1058 DataT
read(
const CoordT &Coords)
const {
1059 #ifdef __SYCL_DEVICE_ONLY__
1060 return __invoke__ImageRead<DataT, OCLImageTy, CoordT>(MImageObj, Coords);
1064 return read<CoordT, Dims>(Coords, Smpl);
1072 template <
typename CoordT,
int Dims =
Dimensions,
1073 typename = std::enable_if_t<
1075 ((IsImageAcc && IsImageAccessReadOnly) ||
1076 (IsHostImageAcc && IsImageAccessAnyRead))>>
1077 DataT
read(
const CoordT &Coords,
const sampler &Smpl)
const {
1078 #ifdef __SYCL_DEVICE_ONLY__
1079 return __invoke__ImageReadSampler<DataT, OCLImageTy, CoordT>(
1080 MImageObj, Coords, Smpl.impl.m_Sampler);
1082 return imageReadSamplerHostImpl<CoordT, DataT>(
1083 Coords, Smpl, getAccessRange() ,
1084 getOffset() , MImgChannelType, MImgChannelOrder,
1096 template <
typename CoordT,
int Dims =
Dimensions,
1097 typename = std::enable_if_t<
1100 ((IsImageAcc && IsImageAccessWriteOnly) ||
1101 (IsHostImageAcc && IsImageAccessAnyWrite))>>
1102 void write(
const CoordT &Coords,
const DataT &Color)
const {
1103 #ifdef __SYCL_DEVICE_ONLY__
1104 __invoke__ImageWrite<OCLImageTy, CoordT, DataT>(MImageObj, Coords, Color);
1116 class __image_array_slice__ {
1119 "Image slice cannot have more then 2 dimensions");
1123 template <
typename CoordT,
1124 typename CoordElemType =
1127 getAdjustedCoords(
const CoordT &Coords)
const {
1128 CoordElemType LastCoord = 0;
1130 if (std::is_same<float, CoordElemType>::value) {
1133 MIdx /
static_cast<float>(Size.template swizzle<Dimensions>());
1141 return AdjustedCoords;
1150 : MBaseAcc(BaseAcc), MIdx(Idx) {}
1152 template <
typename CoordT,
int Dims =
Dimensions,
1153 typename = std::enable_if_t<
1155 DataT
read(
const CoordT &Coords)
const {
1156 return MBaseAcc.read(getAdjustedCoords(Coords));
1159 template <
typename CoordT,
int Dims =
Dimensions,
1160 typename = std::enable_if_t<(Dims > 0) &&
1162 DataT
read(
const CoordT &Coords,
const sampler &Smpl)
const {
1163 return MBaseAcc.read(getAdjustedCoords(Coords), Smpl);
1166 template <
typename CoordT,
int Dims =
Dimensions,
1167 typename = std::enable_if_t<(Dims > 0) &&
1169 void write(
const CoordT &Coords,
const DataT &Color)
const {
1170 return MBaseAcc.write(getAdjustedCoords(Coords), Color);
1173 #ifdef __SYCL_DEVICE_ONLY__
1175 size_t get_count()
const {
return size(); }
1176 size_t size() const noexcept {
return get_range<Dimensions>().size(); }
1178 template <
int Dims = Dimensions,
typename = std::enable_if_t<Dims == 1>>
1180 int2 Count = MBaseAcc.getRangeInternal();
1183 template <
int Dims = Dimensions,
typename = std::enable_if_t<Dims == 2>>
1184 range<2> get_range()
const {
1185 int3 Count = MBaseAcc.getRangeInternal();
1186 return range<2>(Count.x(), Count.y());
1192 size_t get_count()
const {
return size(); }
1194 return MBaseAcc.MImageCount / MBaseAcc.getAccessRange()[
Dimensions];
1198 typename = std::enable_if_t<(Dims == 1 || Dims == 2)>>
1200 return detail::convertToArrayOfN<Dims, 1>(MBaseAcc.getAccessRange());
1223 #ifndef __SYCL_DEVICE_ONLY__
1224 public detail::AccessorBaseHost,
1226 public detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
1227 IsPlaceholder, PropertyListT>,
1228 public detail::OwnerLessBase<
1229 accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
1232 static_assert((AccessTarget == access::target::global_buffer ||
1233 AccessTarget == access::target::constant_buffer ||
1234 AccessTarget == access::target::host_buffer ||
1236 "Expected buffer type");
1238 static_assert((AccessTarget == access::target::global_buffer ||
1239 AccessTarget == access::target::host_buffer ||
1241 (AccessTarget == access::target::constant_buffer &&
1243 "Access mode can be only read for constant buffers");
1245 static_assert(detail::IsPropertyListT<PropertyListT>::value,
1246 "PropertyListT must be accessor_property_list");
1248 using AccessorCommonT =
1254 using AccessorCommonT::AS;
1257 static constexpr
bool IsAccessAnyWrite = AccessorCommonT::IsAccessAnyWrite;
1258 static constexpr
bool IsAccessReadOnly = AccessorCommonT::IsAccessReadOnly;
1259 static constexpr
bool IsConstantBuf = AccessorCommonT::IsConstantBuf;
1260 static constexpr
bool IsGlobalBuf = AccessorCommonT::IsGlobalBuf;
1261 static constexpr
bool IsHostBuf = AccessorCommonT::IsHostBuf;
1262 static constexpr
bool IsPlaceH = AccessorCommonT::IsPlaceH;
1263 static constexpr
bool IsConst = AccessorCommonT::IsConst;
1264 static constexpr
bool IsHostTask = AccessorCommonT::IsHostTask;
1266 using AccessorSubscript =
1267 typename AccessorCommonT::template AccessorSubscript<Dims>;
1270 !IsConst || IsAccessReadOnly,
1271 "A const qualified DataT is only allowed for a read-only accessor");
1273 using ConcreteASPtrType =
typename detail::DecoratedType<
1274 typename std::conditional_t<IsAccessReadOnly && !IsConstantBuf,
1275 const DataT, DataT>,
1278 using RefType = detail::const_if_const_AS<AS, DataT> &;
1279 using ConstRefType =
const DataT &;
1280 using PtrType = detail::const_if_const_AS<AS, DataT> *;
1285 detail::loop<Dims>([&,
this](
size_t I) {
1286 Result = Result * getMemoryRange()[I] + Id[I];
1289 #ifndef __SYCL_DEVICE_ONLY__
1292 Result += getOffset()[I];
1294 #endif // __SYCL_DEVICE_ONLY__
1300 template <
typename T,
int Dims>
1301 struct IsSameAsBuffer
1302 : std::bool_constant<std::is_same_v<T, DataT> && (Dims > 0) &&
1303 (Dims == Dimensions)> {};
1308 if (PropertyList.template has_property<property::no_init>() ||
1309 PropertyList.template has_property<property::noinit>()) {
1317 return AdjustedMode;
1320 template <
typename TagT>
1323 std::is_same<TagT, mode_tag_t<AccessMode>>,
1324 std::is_same<TagT, mode_target_tag_t<AccessMode, AccessTarget>>> {};
1326 template <
typename DataT_,
int Dimensions_,
access::mode AccessMode_,
1328 typename PropertyListT_>
1331 #ifdef __SYCL_DEVICE_ONLY__
1333 id<AdjustedDim> &getOffset() {
return impl.Offset; }
1334 range<AdjustedDim> &getAccessRange() {
return impl.AccessRange; }
1335 range<AdjustedDim> &getMemoryRange() {
return impl.MemRange; }
1337 const id<AdjustedDim> &getOffset()
const {
return impl.Offset; }
1338 const range<AdjustedDim> &getAccessRange()
const {
return impl.AccessRange; }
1339 const range<AdjustedDim> &getMemoryRange()
const {
return impl.MemRange; }
1341 detail::AccessorImplDevice<AdjustedDim> impl;
1344 ConcreteASPtrType MData;
1347 void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
1348 range<AdjustedDim> MemRange, id<AdjustedDim> Offset) {
1350 detail::loop<AdjustedDim>([&,
this](
size_t I) {
1353 getOffset()[I] = Offset[I];
1355 getAccessRange()[I] = AccessRange[I];
1356 getMemoryRange()[I] = MemRange[I];
1361 MData += getTotalOffset();
1367 void __init_esimd(ConcreteASPtrType Ptr) {
1369 #ifdef __ESIMD_FORCE_STATELESS_MEM
1370 detail::loop<AdjustedDim>([&,
this](
size_t I) {
1372 getAccessRange()[I] = 0;
1373 getMemoryRange()[I] = 0;
1378 ConcreteASPtrType getQualifiedPtr() const noexcept {
return MData; }
1380 #ifndef __SYCL_DEVICE_ONLY__
1387 : impl({}, detail::InitializedVal<AdjustedDim, range>::template get<0>(),
1388 detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
1392 :
detail::AccessorBaseHost{Impl} {}
1396 const id<3> getOffset()
const {
1397 if constexpr (IsHostBuf)
1398 return MAccData ? MAccData->MOffset : id<3>();
1402 const range<3> &getAccessRange()
const {
1405 const range<3> getMemoryRange()
const {
1406 if constexpr (IsHostBuf)
1407 return MAccData ? MAccData->MMemoryRange : range(0, 0, 0);
1414 void initHostAcc() { MAccData = &getAccData(); }
1417 void GDBMethodsAnchor() {
1419 const auto *this_const =
this;
1420 (void)getMemoryRange();
1421 (void)this_const->getMemoryRange();
1423 (void)this_const->getOffset();
1425 (void)this_const->getPtr();
1426 (void)getAccessRange();
1427 (void)this_const->getAccessRange();
1431 detail::AccHostDataT *MAccData =
nullptr;
1433 char padding[
sizeof(detail::AccessorImplDevice<AdjustedDim>) +
1434 sizeof(PtrType) -
sizeof(detail::AccessorBaseHost) -
1437 PtrType getQualifiedPtr() const noexcept {
1438 if constexpr (IsHostBuf)
1439 return MAccData ?
reinterpret_cast<PtrType
>(MAccData->MData) :
nullptr;
1447 {0, 0, 0}, {0, 0, 0},
1449 getAdjustedMode({}),
1454 template <
typename,
int, access_mode>
friend class host_accessor;
1456 #endif // __SYCL_DEVICE_ONLY__
1459 friend class sycl::stream;
1460 friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
1462 template <
class Obj>
1473 std::conditional_t<AccessMode == access_mode::read, const DataT, DataT>;
1475 using const_reference =
const DataT &;
1477 template <access::decorated IsDecorated>
1478 using accessor_ptr =
1480 global_ptr<value_type, IsDecorated>, value_type *>;
1482 using iterator =
typename detail::accessor_iterator<value_type, AdjustedDim>;
1483 using const_iterator =
1484 typename detail::accessor_iterator<const value_type, AdjustedDim>;
1485 using reverse_iterator = std::reverse_iterator<iterator>;
1486 using const_reverse_iterator = std::reverse_iterator<const_iterator>;
1489 using size_type = std::size_t;
1493 void throwIfUsedByGraph()
const {
1494 #ifndef __SYCL_DEVICE_ONLY__
1497 "Host accessors cannot be created for buffers "
1498 "which are currently in use by a command graph.");
1528 template <
typename DataT_,
1529 typename = std::enable_if_t<
1530 IsAccessReadOnly && !std::is_same_v<DataT_, DataT> &&
1531 std::is_same_v<std::remove_const_t<DataT_>,
1532 std::remove_const_t<DataT>>>>
1535 #ifdef __SYCL_DEVICE_ONLY__
1536 : impl(other.impl), MData(other.MData) {
1539 #endif // __SYCL_DEVICE_ONLY__
1545 typename = std::enable_if_t<
1547 std::is_same_v<std::remove_const_t<DataT_>,
1548 std::remove_const_t<DataT>>>>
1551 #ifdef __SYCL_DEVICE_ONLY__
1552 : impl(other.impl), MData(other.MData) {
1555 #endif // __SYCL_DEVICE_ONLY__
1558 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1559 typename std::enable_if_t<
1560 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1561 std::is_same_v<T, DataT> && Dims == 0 &&
1562 (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))> * =
1565 buffer<T, 1, AllocatorT> &BufferRef,
1566 const property_list &PropertyList = {},
1568 #ifdef __SYCL_DEVICE_ONLY__
1570 BufferRef.get_range()) {
1575 detail::convertToArrayOfN<3, 1>(
1577 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1578 getAdjustedMode(PropertyList),
1580 IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1582 throwIfUsedByGraph();
1583 preScreenAccessor(PropertyList);
1594 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1595 typename... PropTypes,
1596 typename std::enable_if_t<
1600 std::is_same<T, DataT>::value && Dims == 0 &&
1601 (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))> * =
1604 buffer<T, 1, AllocatorT> &BufferRef,
1605 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1608 #ifdef __SYCL_DEVICE_ONLY__
1610 BufferRef.get_range()) {
1615 detail::convertToArrayOfN<3, 1>(
1617 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1618 getAdjustedMode(PropertyList),
1620 IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1622 throwIfUsedByGraph();
1623 preScreenAccessor(PropertyList);
1634 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1635 typename =
typename std::enable_if_t<
1636 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1637 std::is_same_v<T, DataT> && (Dims == 0) &&
1638 (IsGlobalBuf || IsHostBuf || IsConstantBuf || IsHostTask)>>
1640 buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1641 const property_list &PropertyList = {},
1643 #ifdef __SYCL_DEVICE_ONLY__
1645 BufferRef.get_range()) {
1646 (void)CommandGroupHandler;
1652 detail::convertToArrayOfN<3, 1>(
1654 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1655 getAdjustedMode(PropertyList),
1657 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1658 throwIfUsedByGraph();
1659 preScreenAccessor(PropertyList);
1669 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1670 typename... PropTypes,
1671 typename =
typename std::enable_if_t<
1673 std::is_same_v<T, DataT> && (Dims == 0) &&
1674 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1676 buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1677 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1680 #ifdef __SYCL_DEVICE_ONLY__
1682 BufferRef.get_range()) {
1683 (void)CommandGroupHandler;
1689 detail::convertToArrayOfN<3, 1>(
1691 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1692 getAdjustedMode(PropertyList),
1694 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1695 throwIfUsedByGraph();
1696 preScreenAccessor(PropertyList);
1706 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1707 typename = std::enable_if_t<
1708 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1709 IsSameAsBuffer<T, Dims>::value &&
1710 (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1712 buffer<T, Dims, AllocatorT> &BufferRef,
1713 const property_list &PropertyList = {},
1715 #ifdef __SYCL_DEVICE_ONLY__
1716 : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1722 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1723 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1724 getAdjustedMode(PropertyList),
1726 IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1728 throwIfUsedByGraph();
1729 preScreenAccessor(PropertyList);
1740 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1741 typename... PropTypes,
1742 typename = std::enable_if_t<
1744 IsSameAsBuffer<T, Dims>::value &&
1745 (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1747 buffer<T, Dims, AllocatorT> &BufferRef,
1748 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1751 #ifdef __SYCL_DEVICE_ONLY__
1752 : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
1758 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1759 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1760 getAdjustedMode(PropertyList),
1762 IsPlaceH, BufferRef.OffsetInBytes, BufferRef.IsSubBuffer,
1764 throwIfUsedByGraph();
1765 preScreenAccessor(PropertyList);
1776 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1778 typename = std::enable_if_t<
1779 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1780 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1781 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1783 buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1784 const property_list &PropertyList = {},
1786 :
accessor(BufferRef, PropertyList, CodeLoc) {
1787 adjustAccPropsInBuf(BufferRef);
1790 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1791 typename TagT,
typename... PropTypes,
1792 typename = std::enable_if_t<
1794 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1795 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1797 buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1798 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1801 :
accessor(BufferRef, PropertyList, CodeLoc) {
1802 adjustAccPropsInBuf(BufferRef);
1805 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1806 typename = std::enable_if_t<
1807 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1808 IsSameAsBuffer<T, Dims>::value &&
1809 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1811 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1812 const property_list &PropertyList = {},
1814 #ifdef __SYCL_DEVICE_ONLY__
1815 : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1816 (void)CommandGroupHandler;
1822 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1823 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1824 getAdjustedMode(PropertyList),
1826 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1827 throwIfUsedByGraph();
1828 preScreenAccessor(PropertyList);
1838 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1839 typename... PropTypes,
1840 typename = std::enable_if_t<
1842 IsSameAsBuffer<T, Dims>::value &&
1843 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1845 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1846 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1849 #ifdef __SYCL_DEVICE_ONLY__
1850 : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
1851 (void)CommandGroupHandler;
1857 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1858 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
1859 getAdjustedMode(PropertyList),
1861 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
1862 throwIfUsedByGraph();
1863 preScreenAccessor(PropertyList);
1873 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1875 typename = std::enable_if_t<
1876 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1877 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1878 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1880 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1881 TagT,
const property_list &PropertyList = {},
1883 :
accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1884 adjustAccPropsInBuf(BufferRef);
1887 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1888 typename TagT,
typename... PropTypes,
1889 typename = std::enable_if_t<
1891 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1892 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1894 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1896 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1899 :
accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1900 adjustAccPropsInBuf(BufferRef);
1903 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1904 typename = std::enable_if_t<
1905 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1906 IsSameAsBuffer<T, Dims>::value &&
1907 (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1909 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1910 const property_list &PropertyList = {},
1912 :
accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1914 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1915 typename... PropTypes,
1916 typename = std::enable_if_t<
1918 IsSameAsBuffer<T, Dims>::value &&
1919 (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
1921 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1922 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1925 :
accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1927 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1929 typename = std::enable_if_t<
1930 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1931 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1932 (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
1934 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1935 TagT,
const property_list &PropertyList = {},
1937 :
accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1938 adjustAccPropsInBuf(BufferRef);
1941 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1942 typename TagT,
typename... PropTypes,
1943 typename = std::enable_if_t<
1945 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1946 (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
1948 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1950 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1953 :
accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1954 adjustAccPropsInBuf(BufferRef);
1957 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1958 typename = std::enable_if_t<
1959 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1960 IsSameAsBuffer<T, Dims>::value &&
1961 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1963 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1964 range<Dimensions> AccessRange,
const property_list &PropertyList = {},
1966 :
accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1969 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1970 typename... PropTypes,
1971 typename = std::enable_if_t<
1973 IsSameAsBuffer<T, Dims>::value &&
1974 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1976 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1977 range<Dimensions> AccessRange,
1978 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1981 :
accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1984 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1986 typename = std::enable_if_t<
1987 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1988 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
1989 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
1991 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1992 range<Dimensions> AccessRange, TagT,
1993 const property_list &PropertyList = {},
1995 :
accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1997 adjustAccPropsInBuf(BufferRef);
2000 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2001 typename TagT,
typename... PropTypes,
2002 typename = std::enable_if_t<
2004 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2005 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2007 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2008 range<Dimensions> AccessRange, TagT,
2009 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2012 :
accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
2014 adjustAccPropsInBuf(BufferRef);
2017 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2018 typename = std::enable_if_t<
2019 detail::IsRunTimePropertyListT<PropertyListT>::value &&
2020 IsSameAsBuffer<T, Dims>::value &&
2021 (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
2023 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
2024 id<Dimensions> AccessOffset,
const property_list &PropertyList = {},
2026 #ifdef __SYCL_DEVICE_ONLY__
2027 : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
2031 : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
2032 detail::convertToArrayOfN<3, 1>(AccessRange),
2033 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
2034 getAdjustedMode(PropertyList),
2036 sizeof(DataT), IsPlaceH, BufferRef.OffsetInBytes,
2037 BufferRef.IsSubBuffer, PropertyList) {
2038 throwIfUsedByGraph();
2039 preScreenAccessor(PropertyList);
2042 if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2043 BufferRef.get_range()))
2044 throw sycl::invalid_object_error(
2045 "accessor with requested offset and range would exceed the bounds of "
2047 PI_ERROR_INVALID_VALUE);
2057 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2058 typename... PropTypes,
2059 typename = std::enable_if_t<
2061 IsSameAsBuffer<T, Dims>::value &&
2062 (IsHostBuf || IsHostTask || (IsGlobalBuf || IsConstantBuf))>>
2064 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
2065 id<Dimensions> AccessOffset,
2066 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2069 #ifdef __SYCL_DEVICE_ONLY__
2070 : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
2074 : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
2075 detail::convertToArrayOfN<3, 1>(AccessRange),
2076 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
2077 getAdjustedMode(PropertyList),
2079 sizeof(DataT), IsPlaceH, BufferRef.OffsetInBytes,
2080 BufferRef.IsSubBuffer, PropertyList) {
2081 throwIfUsedByGraph();
2082 preScreenAccessor(PropertyList);
2085 if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2086 BufferRef.get_range()))
2087 throw sycl::invalid_object_error(
2088 "accessor with requested offset and range would exceed the bounds of "
2090 PI_ERROR_INVALID_VALUE);
2100 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2102 typename = std::enable_if_t<
2103 detail::IsRunTimePropertyListT<PropertyListT>::value &&
2104 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2105 (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
2107 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
2108 id<Dimensions> AccessOffset, TagT,
const property_list &PropertyList = {},
2110 :
accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
2111 adjustAccPropsInBuf(BufferRef);
2114 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2115 typename TagT,
typename... PropTypes,
2116 typename = std::enable_if_t<
2118 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2119 (IsGlobalBuf || IsConstantBuf || IsHostTask)>>
2121 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
2122 id<Dimensions> AccessOffset, TagT,
2123 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2126 :
accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
2127 adjustAccPropsInBuf(BufferRef);
2130 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2131 typename = std::enable_if_t<
2132 detail::IsRunTimePropertyListT<PropertyListT>::value &&
2133 IsSameAsBuffer<T, Dims>::value &&
2134 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2136 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2137 range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
2138 const property_list &PropertyList = {},
2140 #ifdef __SYCL_DEVICE_ONLY__
2141 : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
2142 (void)CommandGroupHandler;
2146 : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
2147 detail::convertToArrayOfN<3, 1>(AccessRange),
2148 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
2149 getAdjustedMode(PropertyList),
2151 sizeof(DataT), BufferRef.OffsetInBytes,
2152 BufferRef.IsSubBuffer, PropertyList) {
2153 throwIfUsedByGraph();
2154 preScreenAccessor(PropertyList);
2155 if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2156 BufferRef.get_range()))
2157 throw sycl::invalid_object_error(
2158 "accessor with requested offset and range would exceed the bounds of "
2160 PI_ERROR_INVALID_VALUE);
2171 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2172 typename... PropTypes,
2173 typename = std::enable_if_t<
2175 IsSameAsBuffer<T, Dims>::value &&
2176 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2178 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2179 range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
2180 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2183 #ifdef __SYCL_DEVICE_ONLY__
2184 : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
2185 (void)CommandGroupHandler;
2189 : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
2190 detail::convertToArrayOfN<3, 1>(AccessRange),
2191 detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
2192 getAdjustedMode(PropertyList),
2194 sizeof(DataT), BufferRef.OffsetInBytes,
2195 BufferRef.IsSubBuffer, PropertyList) {
2196 throwIfUsedByGraph();
2197 preScreenAccessor(PropertyList);
2198 if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
2199 BufferRef.get_range()))
2200 throw sycl::invalid_object_error(
2201 "accessor with requested offset and range would exceed the bounds of "
2203 PI_ERROR_INVALID_VALUE);
2214 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2216 typename = std::enable_if_t<
2217 detail::IsRunTimePropertyListT<PropertyListT>::value &&
2218 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2219 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2221 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2222 range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
2223 const property_list &PropertyList = {},
2225 :
accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2226 PropertyList, CodeLoc) {
2227 adjustAccPropsInBuf(BufferRef);
2230 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2231 typename TagT,
typename... PropTypes,
2232 typename = std::enable_if_t<
2234 IsSameAsBuffer<T, Dims>::value && IsValidTag<TagT>::value &&
2235 (IsGlobalBuf || IsConstantBuf || IsHostBuf || IsHostTask)>>
2237 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
2238 range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
2239 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
2242 :
accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2243 PropertyList, CodeLoc) {
2244 adjustAccPropsInBuf(BufferRef);
2247 template <
typename... NewPropsT>
2250 ext::oneapi::accessor_property_list<NewPropsT...>> &Other,
2252 #ifdef __SYCL_DEVICE_ONLY__
2253 : impl(Other.impl), MData(Other.MData)
2255 : detail::AccessorBaseHost(Other), MAccData(Other.MAccData)
2259 "Conversion is only available for accessor_property_list");
2261 PropertyListT::template areSameCompileTimeProperties<NewPropsT...>(),
2262 "Compile-time-constant properties must be the same");
2263 #ifndef __SYCL_DEVICE_ONLY__
2270 std::swap(impl, other.impl);
2271 #ifdef __SYCL_DEVICE_ONLY__
2272 std::swap(MData, other.MData);
2274 std::swap(MAccData, other.MAccData);
2278 bool is_placeholder()
const {
2279 #ifdef __SYCL_DEVICE_ONLY__
2286 size_t get_size()
const {
return getAccessRange().size() *
sizeof(DataT); }
2289 size_t get_count()
const {
return size(); }
2290 size_type size() const noexcept {
return getAccessRange().size(); }
2292 size_type byte_size() const noexcept {
return size() *
sizeof(DataT); }
2294 size_type max_size() const noexcept {
2298 bool empty() const noexcept {
return size() == 0; }
2301 typename = std::enable_if_t<Dims ==
Dimensions && (Dims > 0)>>
2302 range<Dimensions> get_range()
const {
2303 return getRange<Dims>();
2307 typename = std::enable_if_t<Dims ==
Dimensions && (Dims > 0)>>
2309 return getOffset<Dims>();
2312 template <
int Dims =
Dimensions,
typename RefT = RefType,
2313 typename = std::enable_if_t<Dims == 0 &&
2314 (IsAccessAnyWrite || IsAccessReadOnly)>>
2317 return *(getQualifiedPtr() + LinearIndex);
2322 !IsAccessReadOnly && Dims == 0>>
2324 *getQualifiedPtr() = Other;
2330 !IsAccessReadOnly && Dims == 0>>
2332 *getQualifiedPtr() = std::move(Other);
2337 typename = std::enable_if_t<(Dims > 0) &&
2338 (IsAccessAnyWrite || IsAccessReadOnly)>>
2341 return getQualifiedPtr()[LinearIndex];
2344 template <
int Dims = Dimensions>
2345 operator typename std::enable_if_t<Dims == 0 &&
2347 #ifdef __ENABLE_USM_ADDR_SPACE__
2354 return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2355 getQualifiedPtr() + LinearIndex));
2358 template <
int Dims = Dimensions>
2361 operator[](id<Dimensions> Index)
const {
2363 return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2364 getQualifiedPtr() + LinearIndex));
2367 template <
int Dims = Dimensions>
2371 const size_t LinearIndex =
getLinearIndex(id<AdjustedDim>(Index));
2372 return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
2373 getQualifiedPtr() + LinearIndex));
2375 template <
int Dims = Dimensions,
typename = std::enable_if_t<(Dims > 1)>>
2377 return AccessorSubscript<Dims - 1>(*
this, Index);
2381 typename = std::enable_if_t<
2382 (AccessTarget_ == access::target::host_buffer) ||
2384 std::add_pointer_t<value_type> get_pointer() const noexcept {
2385 return getPointerAdjusted();
2392 "accessor::get_pointer() is deprecated, please use get_multi_ptr()")
2393 global_ptr<DataT> get_pointer() const noexcept {
2394 return global_ptr<DataT>(
2395 const_cast<typename detail::DecoratedType<DataT, AS>::type *
>(
2396 getPointerAdjusted()));
2400 typename = std::enable_if_t<AccessTarget_ ==
2401 access::target::constant_buffer>>
2402 constant_ptr<DataT> get_pointer()
const {
2403 return constant_ptr<DataT>(getPointerAdjusted());
2406 template <access::decorated IsDecorated>
2407 accessor_ptr<IsDecorated> get_multi_ptr() const noexcept {
2408 return accessor_ptr<IsDecorated>(getPointerAdjusted());
2414 template <
typename Property>
2415 typename std::enable_if_t<
2416 !ext::oneapi::is_compile_time_property<Property>::value,
bool>
2418 #ifndef __SYCL_DEVICE_ONLY__
2419 return getPropList().template has_property<Property>();
2428 template <
typename Property,
2429 typename =
typename std::enable_if_t<
2430 !ext::oneapi::is_compile_time_property<Property>::value>>
2432 #ifndef __SYCL_DEVICE_ONLY__
2433 return getPropList().template get_property<Property>();
2439 template <
typename Property>
2441 typename std::enable_if_t<
2442 ext::oneapi::is_compile_time_property<Property>::value> * = 0) {
2443 return PropertyListT::template has_property<Property>();
2446 template <
typename Property>
2448 typename std::enable_if_t<
2449 ext::oneapi::is_compile_time_property<Property>::value> * = 0) {
2450 return PropertyListT::template get_property<Property>();
2456 iterator begin() const noexcept {
2457 return iterator::getBegin(
2459 detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2460 getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2463 iterator
end() const noexcept {
2464 return iterator::getEnd(
2466 detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2467 getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2470 const_iterator cbegin() const noexcept {
2471 return const_iterator::getBegin(
2473 detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2474 getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2477 const_iterator cend() const noexcept {
2478 return const_iterator::getEnd(
2480 detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
2481 getRange<AdjustedDim>(), getOffset<AdjustedDim>());
2484 reverse_iterator rbegin() const noexcept {
return reverse_iterator(
end()); }
2485 reverse_iterator rend() const noexcept {
return reverse_iterator(begin()); }
2487 const_reverse_iterator crbegin() const noexcept {
2488 return const_reverse_iterator(cend());
2490 const_reverse_iterator crend() const noexcept {
2491 return const_reverse_iterator(cbegin());
2495 template <
int Dims,
typename = std::enable_if_t<(Dims > 0)>>
2496 range<Dims> getRange()
const {
2497 return detail::convertToArrayOfN<AdjustedDim, 1>(getAccessRange());
2500 template <
int Dims = Dimensions,
typename = std::enable_if_t<(Dims > 0)>>
2501 id<Dims> getOffset()
const {
2505 "Accessor has no_offset property, get_offset() can not be used");
2506 return detail::convertToArrayOfN<Dims, 0>(getOffset());
2509 #ifdef __SYCL_DEVICE_ONLY__
2510 size_t getTotalOffset() const noexcept {
2511 size_t TotalOffset = 0;
2512 detail::loop<Dimensions>([&,
this](
size_t I) {
2513 TotalOffset = TotalOffset * impl.MemRange[I];
2516 TotalOffset += impl.Offset[I];
2529 auto getPointerAdjusted() const noexcept {
2530 #ifdef __SYCL_DEVICE_ONLY__
2531 return getQualifiedPtr() - getTotalOffset();
2533 return getQualifiedPtr();
2539 if (PropertyList.template has_property<property::no_init>() &&
2541 throw sycl::invalid_object_error(
2542 "accessor would cannot be both read_only and no_init",
2543 PI_ERROR_INVALID_VALUE);
2547 template <
typename BufT,
typename... PropTypes>
2548 void adjustAccPropsInBuf(BufT &Buffer) {
2554 property_list PropList{
2556 Buffer.addOrReplaceAccessorProperties(PropList);
2558 deleteAccPropsFromBuf(Buffer);
2562 template <
typename BufT>
void deleteAccPropsFromBuf(BufT &Buffer) {
2563 Buffer.deleteAccProps(
2568 template <
typename DataT,
int Dimensions,
typename AllocatorT>
2569 accessor(buffer<DataT, Dimensions, AllocatorT>)
2573 template <
typename DataT,
int Dimensions,
typename AllocatorT,
2575 accessor(buffer<DataT, Dimensions, AllocatorT>,
2576 const ext::oneapi::accessor_property_list<PropsT...> &)
2579 ext::oneapi::accessor_property_list<PropsT...>>;
2581 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1>
2582 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1)
2583 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
2584 detail::deduceAccessTarget<Type1, Type1>(target::device),
2587 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2589 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1,
2590 const ext::oneapi::accessor_property_list<PropsT...> &)
2591 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
2592 detail::deduceAccessTarget<Type1, Type1>(target::device),
2594 ext::oneapi::accessor_property_list<PropsT...>>;
2596 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2598 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2)
2599 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
2600 detail::deduceAccessTarget<Type1, Type2>(target::device),
2603 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2604 typename Type2,
typename... PropsT>
2605 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2,
2606 const ext::oneapi::accessor_property_list<PropsT...> &)
2607 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
2608 detail::deduceAccessTarget<Type1, Type2>(target::device),
2610 ext::oneapi::accessor_property_list<PropsT...>>;
2612 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2613 typename Type2,
typename Type3>
2614 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3)
2615 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
2616 detail::deduceAccessTarget<Type2, Type3>(target::device),
2619 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2620 typename Type2,
typename Type3,
typename... PropsT>
2621 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3,
2622 const ext::oneapi::accessor_property_list<PropsT...> &)
2623 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
2624 detail::deduceAccessTarget<Type2, Type3>(target::device),
2626 ext::oneapi::accessor_property_list<PropsT...>>;
2628 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2629 typename Type2,
typename Type3,
typename Type4>
2630 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
2631 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2632 detail::deduceAccessTarget<Type3, Type4>(target::device),
2635 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2636 typename Type2,
typename Type3,
typename Type4,
typename... PropsT>
2637 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
2638 const ext::oneapi::accessor_property_list<PropsT...> &)
2639 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2640 detail::deduceAccessTarget<Type3, Type4>(target::device),
2642 ext::oneapi::accessor_property_list<PropsT...>>;
2644 template <
typename DataT,
int Dimensions,
typename AllocatorT>
2645 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &)
2649 template <
typename DataT,
int Dimensions,
typename AllocatorT,
2651 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &,
2652 const ext::oneapi::accessor_property_list<PropsT...> &)
2655 ext::oneapi::accessor_property_list<PropsT...>>;
2657 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1>
2658 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1)
2659 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
2660 detail::deduceAccessTarget<Type1, Type1>(target::device),
2663 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2665 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1,
2666 const ext::oneapi::accessor_property_list<PropsT...> &)
2667 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
2668 detail::deduceAccessTarget<Type1, Type1>(target::device),
2670 ext::oneapi::accessor_property_list<PropsT...>>;
2672 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2674 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2)
2675 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
2676 detail::deduceAccessTarget<Type1, Type2>(target::device),
2679 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2680 typename Type2,
typename... PropsT>
2681 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2,
2682 const ext::oneapi::accessor_property_list<PropsT...> &)
2683 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
2684 detail::deduceAccessTarget<Type1, Type2>(target::device),
2686 ext::oneapi::accessor_property_list<PropsT...>>;
2688 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2689 typename Type2,
typename Type3>
2690 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2, Type3)
2691 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
2692 detail::deduceAccessTarget<Type2, Type3>(target::device),
2695 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2696 typename Type2,
typename Type3,
typename... PropsT>
2697 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2, Type3,
2698 const ext::oneapi::accessor_property_list<PropsT...> &)
2699 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
2700 detail::deduceAccessTarget<Type2, Type3>(target::device),
2702 ext::oneapi::accessor_property_list<PropsT...>>;
2704 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2705 typename Type2,
typename Type3,
typename Type4>
2706 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2, Type3,
2708 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2709 detail::deduceAccessTarget<Type3, Type4>(target::device),
2712 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2713 typename Type2,
typename Type3,
typename Type4,
typename... PropsT>
2714 accessor(buffer<DataT, Dimensions, AllocatorT>, handler &, Type1, Type2, Type3,
2715 Type4,
const ext::oneapi::accessor_property_list<PropsT...> &)
2716 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2717 detail::deduceAccessTarget<Type3, Type4>(target::device),
2719 ext::oneapi::accessor_property_list<PropsT...>>;
2727 #ifndef __SYCL_DEVICE_ONLY__
2728 public detail::LocalAccessorBaseHost,
2730 public detail::accessor_common<DataT, Dimensions, AccessMode,
2731 access::target::local, IsPlaceholder> {
2739 using AccessorCommonT::AS;
2743 static constexpr
bool IsAccessAnyWrite = AccessorCommonT::IsAccessAnyWrite;
2744 static constexpr
bool IsAccessReadOnly = AccessorCommonT::IsAccessReadOnly;
2745 static constexpr
bool IsConst = AccessorCommonT::IsConst;
2758 #ifdef __SYCL_DEVICE_ONLY__
2767 detail::loop<AdjustedDim>(
2768 [&,
this](
size_t I) { getSize()[I] = AccessRange[I]; });
2774 void __init_esimd(ConcreteASPtrType Ptr) {
2776 detail::loop<AdjustedDim>([&,
this](
size_t I) { getSize()[I] = 0; });
2781 local_accessor_base()
2782 : impl(
detail::InitializedVal<AdjustedDim, range>::template
get<0>()) {}
2785 ConcreteASPtrType getQualifiedPtr()
const {
return MData; }
2787 ConcreteASPtrType MData;
2793 0,
sizeof(DataT)} {}
2797 :
detail::LocalAccessorBaseHost{Impl} {}
2799 char padding[
sizeof(detail::LocalAccessorBaseDevice<AdjustedDim>) +
2800 sizeof(
PtrType) -
sizeof(detail::LocalAccessorBaseHost)];
2817 const auto *this_const =
this;
2819 (void)this_const->getSize();
2821 (void)this_const->getPtr();
2825 #endif // __SYCL_DEVICE_ONLY__
2831 [&,
this](
size_t I) { Result = Result * getSize()[I] + Id[I]; });
2835 template <
class Obj>
2848 template <
int Dims = Dimensions,
typename = std::enable_if_t<Dims == 0>>
2851 #ifdef __SYCL_DEVICE_ONLY__
2854 : LocalAccessorBaseHost(
range<3>{1, 1, 1}, AdjustedDim,
sizeof(DataT)) {
2861 template <
int Dims = Dimensions,
typename = std::enable_if_t<Dims == 0>>
2865 #ifdef __SYCL_DEVICE_ONLY__
2870 : LocalAccessorBaseHost(
range<3>{1, 1, 1}, AdjustedDim,
sizeof(DataT),
2878 template <
int Dims = Dimensions,
typename = std::enable_if_t<(Dims > 0)>>
2882 #ifdef __SYCL_DEVICE_ONLY__
2883 : impl(AllocationSize){}
2885 : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2886 AdjustedDim,
sizeof(DataT)) {
2894 typename = std::enable_if_t<(Dims > 0)>>
2899 #ifdef __SYCL_DEVICE_ONLY__
2900 : impl(AllocationSize) {
2904 : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
2905 AdjustedDim,
sizeof(DataT), propList) {
2912 size_t get_size()
const {
return getSize().size() *
sizeof(DataT); }
2915 size_t get_count()
const {
return size(); }
2916 size_t size() const noexcept {
return getSize().size(); }
2918 template <
int Dims = Dimensions,
typename = std::enable_if_t<(Dims > 0)>>
2920 return detail::convertToArrayOfN<Dims, 1>(getSize());
2924 typename = std::enable_if_t<Dims == 0 &&
2925 (IsAccessAnyWrite || IsAccessReadOnly)>>
2927 return *getQualifiedPtr();
2931 typename = std::enable_if_t<(Dims > 0) &&
2932 (IsAccessAnyWrite || IsAccessReadOnly)>>
2935 return getQualifiedPtr()[LinearIndex];
2939 typename = std::enable_if_t<Dims == 1 &&
2940 (IsAccessAnyWrite || IsAccessReadOnly)>>
2942 return getQualifiedPtr()[Index];
2945 template <
int Dims = Dimensions>
2946 operator typename std::enable_if_t<
2949 return atomic<DataT, AS>(
2953 template <
int Dims = Dimensions>
2959 getQualifiedPtr() + LinearIndex));
2962 template <
int Dims = Dimensions>
2967 getQualifiedPtr() + Index));
2970 template <
int Dims = Dimensions,
typename = std::enable_if_t<(Dims > 1)>>
2971 typename AccessorCommonT::template AccessorSubscript<
2979 return impl == Rhs.
impl;
2982 return !(*
this == Rhs);
2992 public detail::OwnerLessBase<
2993 accessor<DataT, Dimensions, AccessMode, access::target::local,
3000 !local_acc::IsConst || local_acc::IsAccessReadOnly,
3001 "A const qualified DataT is only allowed for a read-only accessor");
3004 using local_acc::local_acc;
3011 #ifdef __SYCL_DEVICE_ONLY__
3015 void __init(
typename local_acc::ConcreteASPtrType Ptr,
3019 local_acc::__init(Ptr, AccessRange,
range,
id);
3025 void __init_esimd(
typename local_acc::ConcreteASPtrType Ptr) {
3026 local_acc::__init_esimd(Ptr);
3032 local_acc::impl = detail::InitializedVal<local_acc::AdjustedDim,
3033 range>::template get<0>();
3042 template <
typename DataT,
int Dimensions = 1>
3044 :
public local_accessor_base<DataT, Dimensions,
3045 detail::accessModeFromConstness<DataT>(),
3046 access::placeholder::false_t>,
3047 public detail::OwnerLessBase<local_accessor<DataT, Dimensions>> {
3051 detail::accessModeFromConstness<DataT>(),
3055 !local_acc::IsConst || local_acc::IsAccessReadOnly,
3056 "A const qualified DataT is only allowed for a read-only accessor");
3059 using local_acc::local_acc;
3061 #ifdef __SYCL_DEVICE_ONLY__
3065 void __init(
typename local_acc::ConcreteASPtrType Ptr,
3066 range<local_acc::AdjustedDim> AccessRange,
3067 range<local_acc::AdjustedDim> range,
3068 id<local_acc::AdjustedDim>
id) {
3069 local_acc::__init(Ptr, AccessRange, range,
id);
3075 void __init_esimd(
typename local_acc::ConcreteASPtrType Ptr) {
3076 local_acc::__init_esimd(Ptr);
3082 local_acc::impl = detail::InitializedVal<local_acc::AdjustedDim,
3083 range>::template get<0>();
3093 template <
typename DataT_,
3094 typename = std::enable_if_t<
3095 std::is_const_v<DataT> &&
3096 std::is_same_v<DataT_, std::remove_const_t<DataT>>>>
3097 local_accessor(
const local_accessor<DataT_, Dimensions> &other) {
3098 local_acc::impl = other.impl;
3099 #ifdef __SYCL_DEVICE_ONLY__
3100 local_acc::MData = other.MData;
3104 using value_type = DataT;
3105 using iterator = value_type *;
3106 using const_iterator =
const value_type *;
3107 using reverse_iterator = std::reverse_iterator<iterator>;
3108 using const_reverse_iterator = std::reverse_iterator<const_iterator>;
3111 using size_type = std::size_t;
3113 template <access::decorated IsDecorated>
3114 using accessor_ptr = local_ptr<value_type, IsDecorated>;
3116 template <
typename DataT_>
3117 bool operator==(
const local_accessor<DataT_, Dimensions> &Rhs)
const {
3118 return local_acc::impl == Rhs.impl;
3121 template <
typename DataT_>
3122 bool operator!=(
const local_accessor<DataT_, Dimensions> &Rhs)
const {
3123 return !(*
this == Rhs);
3126 void swap(local_accessor &other) { std::swap(this->impl, other.impl); }
3128 size_type byte_size() const noexcept {
return this->size() *
sizeof(DataT); }
3130 size_type max_size() const noexcept {
3134 bool empty() const noexcept {
return this->size() == 0; }
3136 iterator begin() const noexcept {
3138 return local_acc::getQualifiedPtr();
3142 iterator end() const noexcept {
3146 return begin() + this->size();
3149 const_iterator cbegin() const noexcept {
return const_iterator(begin()); }
3150 const_iterator cend() const noexcept {
return const_iterator(
end()); }
3152 reverse_iterator rbegin() const noexcept {
return reverse_iterator(
end()); }
3153 reverse_iterator rend() const noexcept {
return reverse_iterator(begin()); }
3155 const_reverse_iterator crbegin() const noexcept {
3156 return const_reverse_iterator(
end());
3158 const_reverse_iterator crend() const noexcept {
3159 return const_reverse_iterator(begin());
3163 "local_accessor::get_pointer() is deprecated, please use get_multi_ptr()")
3164 local_ptr<DataT> get_pointer() const noexcept {
3165 return local_ptr<DataT>(local_acc::getQualifiedPtr());
3168 template <access::decorated IsDecorated>
3169 accessor_ptr<IsDecorated> get_multi_ptr() const noexcept {
3170 return accessor_ptr<IsDecorated>(local_acc::getQualifiedPtr());
3173 template <
typename Property>
bool has_property() const noexcept {
3174 #ifndef __SYCL_DEVICE_ONLY__
3175 return this->getPropList().template has_property<Property>();
3181 template <
typename Property> Property
get_property()
const {
3182 #ifndef __SYCL_DEVICE_ONLY__
3183 return this->getPropList().template get_property<Property>();
3190 typename = std::enable_if_t<!std::is_const_v<DataT> && Dims == 0>>
3191 const local_accessor &
operator=(
const value_type &Other)
const {
3192 *local_acc::getQualifiedPtr() = Other;
3197 typename = std::enable_if_t<!std::is_const_v<DataT> && Dims == 0>>
3198 const local_accessor &
operator=(value_type &&Other)
const {
3199 *local_acc::getQualifiedPtr() = std::move(Other);
3204 friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
3217 access::target::image, IsPlaceholder>,
3219 accessor<DataT, Dimensions, AccessMode, access::target::image,
3227 template <
typename AllocatorT>
3232 Image, CommandGroupHandler, Image.getElementSize()) {
3233 #ifndef __SYCL_DEVICE_ONLY__
3239 template <
typename AllocatorT>
3244 Image, CommandGroupHandler, Image.getElementSize()) {
3246 #ifndef __SYCL_DEVICE_ONLY__
3251 #ifdef __SYCL_DEVICE_ONLY__
3259 void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
3262 void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
3281 :
public detail::image_accessor<DataT, Dimensions, AccessMode,
3282 access::target::host_image, IsPlaceholder>,
3283 public detail::OwnerLessBase<
3284 accessor<DataT, Dimensions, AccessMode, access::target::host_image,
3287 template <
typename AllocatorT>
3291 Image, Image.getElementSize()) {}
3293 template <
typename AllocatorT>
3298 Image, Image.getElementSize()) {
3316 access::target::image, IsPlaceholder>,
3318 accessor<DataT, Dimensions, AccessMode, access::target::image_array,
3320 #ifdef __SYCL_DEVICE_ONLY__
3328 void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
3331 void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
3338 template <
typename AllocatorT>
3343 Image, CommandGroupHandler, Image.getElementSize()) {
3344 #ifndef __SYCL_DEVICE_ONLY__
3350 template <
typename AllocatorT>
3355 Image, CommandGroupHandler, Image.getElementSize()) {
3357 #ifndef __SYCL_DEVICE_ONLY__
3370 template <
typename DataT,
int Dimensions = 1,
3373 :
public accessor<DataT, Dimensions, AccessMode, target::host_buffer,
3374 access::placeholder::false_t> {
3382 template <
typename T,
int Dims>
3384 : std::bool_constant<std::is_same_v<T, DataT> && (Dims > 0) &&
3385 (Dims == Dimensions)> {};
3392 AccessorT::__init(Ptr, AccessRange, MemRange, Offset);
3395 #ifndef __SYCL_DEVICE_ONLY__
3400 template <
class Obj>
3401 friend decltype(Obj::impl)
getSyclObjImpl(
const Obj &SyclObject);
3405 #endif // __SYCL_DEVICE_ONLY__
3433 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3434 typename =
typename std::enable_if_t<std::is_same_v<T, DataT> &&
3440 : AccessorT(BufferRef, PropertyList, CodeLoc) {}
3442 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3443 typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3448 : AccessorT(BufferRef, PropertyList, CodeLoc) {}
3450 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3451 typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3458 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3459 typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3464 : AccessorT(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
3466 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3467 typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3472 :
host_accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
3474 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3475 typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3480 : AccessorT(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
3482 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3483 typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3488 :
host_accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
3490 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3491 typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3496 : AccessorT(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
3499 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3500 typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3506 :
host_accessor(BufferRef, CommandGroupHandler, AccessRange, {},
3507 PropertyList, CodeLoc) {}
3509 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3510 typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3515 : AccessorT(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
3518 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3519 typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3525 :
host_accessor(BufferRef, AccessRange, AccessOffset, PropertyList,
3528 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3529 typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3535 : AccessorT(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
3536 PropertyList, CodeLoc) {}
3538 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
3539 typename = std::enable_if_t<IsSameAsBuffer<T, Dims>::value>>
3545 :
host_accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
3546 PropertyList, CodeLoc) {}
3550 !IsAccessReadOnly && Dims == 0>>
3552 operator=(
const typename AccessorT::value_type &Other)
const {
3553 *AccessorT::getQualifiedPtr() = Other;
3559 !IsAccessReadOnly && Dims == 0>>
3561 *AccessorT::getQualifiedPtr() = std::move(Other);
3566 template <
typename DataT_,
3567 typename = std::enable_if_t<
3568 IsAccessReadOnly && !std::is_same_v<DataT_, DataT> &&
3569 std::is_same_v<std::remove_const_t<DataT_>,
3570 std::remove_const_t<DataT>>>>
3572 #ifndef __SYCL_DEVICE_ONLY__
3574 AccessorT::MAccData = other.MAccData;
3577 #endif // __SYCL_DEVICE_ONLY__
3583 typename = std::enable_if_t<
3585 std::is_same_v<DataT_, std::remove_const_t<DataT>>>>
3587 #ifndef __SYCL_DEVICE_ONLY__
3589 AccessorT::MAccData = other.MAccData;
3592 #endif // __SYCL_DEVICE_ONLY__
3599 #ifndef __SYCL_DEVICE_ONLY__
3603 return this->impl.owner_before(
3608 return this->impl.owner_before(Other.impl);
3611 bool ext_oneapi_owner_before(
3614 bool ext_oneapi_owner_before(
const host_accessor &Other)
const noexcept;
3618 template <
typename DataT,
int Dimensions,
typename AllocatorT>
3622 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1>
3625 detail::deduceAccessMode<Type1, Type1>()>;
3627 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
3631 detail::deduceAccessMode<Type1, Type2>()>;
3633 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
3634 typename Type2,
typename Type3>
3637 detail::deduceAccessMode<Type2, Type3>()>;
3639 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
3640 typename Type2,
typename Type3,
typename Type4>
3643 detail::deduceAccessMode<Type3, Type4>()>;
3645 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
3646 typename Type2,
typename Type3,
typename Type4,
typename Type5>
3649 detail::deduceAccessMode<Type4, Type5>()>;
3656 #ifndef __SYCL_DEVICE_ONLY__
3657 private detail::UnsampledImageAccessorBaseHost,
3658 #endif // __SYCL_DEVICE_ONLY__
3659 public detail::OwnerLessBase<
3660 unsampled_image_accessor<DataT, Dimensions, AccessMode, AccessTarget>> {
3661 static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
3662 std::is_same_v<DataT, float4> ||
3663 std::is_same_v<DataT, half4>,
3664 "The data type of an image accessor must be only int4, "
3665 "uint4, float4 or half4 from SYCL namespace");
3668 "Access mode must be either read or write.");
3670 #ifdef __SYCL_DEVICE_ONLY__
3671 char MPadding[
sizeof(detail::UnsampledImageAccessorBaseHost)];
3673 using host_base_class = detail::UnsampledImageAccessorBaseHost;
3674 #endif // __SYCL_DEVICE_ONLY__
3678 const DataT, DataT>::type;
3682 template <
typename AllocatorT>
3687 #ifdef __SYCL_DEVICE_ONLY__
3690 : host_base_class(detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
3693 {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
3698 aspect ImageAspect = aspect::image;
3702 "Device associated with command group handler does not have "
3707 AccessMode, (
const void *)
typeid(DataT).name(),
sizeof(DataT), CodeLoc);
3711 #endif // __SYCL_DEVICE_ONLY__
3715 unsampled_image_accessor(
const unsampled_image_accessor &Rhs) =
default;
3717 unsampled_image_accessor(unsampled_image_accessor &&Rhs) =
default;
3719 unsampled_image_accessor &
3720 operator=(
const unsampled_image_accessor &Rhs) =
default;
3722 unsampled_image_accessor &
operator=(unsampled_image_accessor &&Rhs) =
default;
3724 ~unsampled_image_accessor() =
default;
3726 #ifdef __SYCL_DEVICE_ONLY__
3727 bool operator==(
const unsampled_image_accessor &Rhs)
const;
3730 return Rhs.impl == impl;
3732 #endif // __SYCL_DEVICE_ONLY__
3735 return !(Rhs == *
this);
3741 #ifdef __SYCL_DEVICE_ONLY__
3745 return host_base_class::getSize().size();
3746 #endif // __SYCL_DEVICE_ONLY__
3753 template <
typename CoordT,
3757 DataT
read(
const CoordT &Coords)
const noexcept {
3758 #ifdef __SYCL_DEVICE_ONLY__
3760 std::ignore = Coords;
3761 return {0, 0, 0, 0};
3763 return host_base_class::read<DataT>(Coords);
3764 #endif // __SYCL_DEVICE_ONLY__
3771 template <
typename CoordT,
3775 void write(
const CoordT &Coords,
const DataT &Color)
const {
3776 #ifdef __SYCL_DEVICE_ONLY__
3778 std::ignore = Coords;
3779 std::ignore = Color;
3781 host_base_class::write<DataT>(Coords, Color);
3782 #endif // __SYCL_DEVICE_ONLY__
3787 #ifndef __SYCL_DEVICE_ONLY__
3788 : host_base_class{Impl}
3789 #endif // __SYCL_DEVICE_ONLY__
3794 template <
class Obj>
3801 template <
typename DataT,
int Dimensions = 1,
3806 :
private detail::UnsampledImageAccessorBaseHost,
3807 public detail::OwnerLessBase<
3809 static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
3810 std::is_same_v<DataT, float4> ||
3811 std::is_same_v<DataT, half4>,
3812 "The data type of an image accessor must be only int4, "
3813 "uint4, float4 or half4 from SYCL namespace");
3815 using base_class = detail::UnsampledImageAccessorBaseHost;
3819 const DataT, DataT>::type;
3823 template <
typename AllocatorT>
3828 : base_class(detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
3838 AccessMode, (
const void *)
typeid(DataT).name(),
sizeof(DataT), CodeLoc);
3843 host_unsampled_image_accessor(
const host_unsampled_image_accessor &Rhs) =
3846 host_unsampled_image_accessor(host_unsampled_image_accessor &&Rhs) =
default;
3848 host_unsampled_image_accessor &
3849 operator=(
const host_unsampled_image_accessor &Rhs) =
default;
3851 host_unsampled_image_accessor &
3852 operator=(host_unsampled_image_accessor &&Rhs) =
default;
3854 ~host_unsampled_image_accessor() =
default;
3857 return Rhs.impl == impl;
3860 return !(Rhs == *
this);
3865 size_t size() const noexcept {
return base_class::getSize().size(); }
3874 typename = std::enable_if_t<
3878 DataT
read(
const CoordT &Coords)
const noexcept
3879 #ifdef __SYCL_DEVICE_ONLY__
3885 return base_class::read<DataT>(Coords);
3896 typename = std::enable_if_t<
3900 void write(
const CoordT &Coords,
const DataT &Color)
const
3901 #ifdef __SYCL_DEVICE_ONLY__
3907 base_class::write<DataT>(Coords, Color);
3914 : base_class{Impl} {}
3916 template <
class Obj>
3926 #ifndef __SYCL_DEVICE_ONLY__
3928 #endif // __SYCL_DEVICE_ONLY__
3930 sampled_image_accessor<DataT, Dimensions, AccessTarget>> {
3931 static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
3932 std::is_same_v<DataT, float4> ||
3933 std::is_same_v<DataT, half4>,
3934 "The data type of an image accessor must be only int4, "
3935 "uint4, float4 or half4 from SYCL namespace");
3937 #ifdef __SYCL_DEVICE_ONLY__
3941 #endif // __SYCL_DEVICE_ONLY__
3948 template <
typename AllocatorT>
3953 #ifdef __SYCL_DEVICE_ONLY__
3956 : host_base_class(detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
3959 {ImageRef.getRowPitch(), ImageRef.getSlicePitch(), 0},
3964 aspect ImageAspect = aspect::image;
3968 "Device associated with command group handler does not have "
3973 (
const void *)
typeid(DataT).name(),
sizeof(DataT), CodeLoc);
3977 #endif // __SYCL_DEVICE_ONLY__
3981 sampled_image_accessor(
const sampled_image_accessor &Rhs) =
default;
3983 sampled_image_accessor(sampled_image_accessor &&Rhs) =
default;
3985 sampled_image_accessor &
3986 operator=(
const sampled_image_accessor &Rhs) =
default;
3988 sampled_image_accessor &
operator=(sampled_image_accessor &&Rhs) =
default;
3990 ~sampled_image_accessor() =
default;
3992 #ifdef __SYCL_DEVICE_ONLY__
3993 bool operator==(
const sampled_image_accessor &Rhs)
const;
3996 return Rhs.impl == impl;
3998 #endif // __SYCL_DEVICE_ONLY__
4001 return !(Rhs == *
this);
4007 #ifdef __SYCL_DEVICE_ONLY__
4011 return host_base_class::getSize().size();
4012 #endif // __SYCL_DEVICE_ONLY__
4018 template <
typename CoordT,
4021 DataT
read(
const CoordT &Coords)
const noexcept {
4022 #ifdef __SYCL_DEVICE_ONLY__
4024 std::ignore = Coords;
4025 return {0, 0, 0, 0};
4027 return host_base_class::read<DataT>(Coords);
4028 #endif // __SYCL_DEVICE_ONLY__
4033 #ifndef __SYCL_DEVICE_ONLY__
4034 : host_base_class{Impl}
4035 #endif // __SYCL_DEVICE_ONLY__
4040 template <
class Obj>
4047 template <
typename DataT,
int Dimensions>
4049 :
private detail::SampledImageAccessorBaseHost,
4050 public detail::OwnerLessBase<
4051 host_sampled_image_accessor<DataT, Dimensions>> {
4052 static_assert(std::is_same_v<DataT, int4> || std::is_same_v<DataT, uint4> ||
4053 std::is_same_v<DataT, float4> ||
4054 std::is_same_v<DataT, half4>,
4055 "The data type of an image accessor must be only int4, "
4056 "uint4, float4 or half4 from SYCL namespace");
4058 using base_class = detail::SampledImageAccessorBaseHost;
4065 template <
typename AllocatorT>
4070 : base_class(detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
4080 (
const void *)
typeid(DataT).name(),
sizeof(DataT), CodeLoc);
4085 host_sampled_image_accessor(
const host_sampled_image_accessor &Rhs) =
default;
4087 host_sampled_image_accessor(host_sampled_image_accessor &&Rhs) =
default;
4089 host_sampled_image_accessor &
4090 operator=(
const host_sampled_image_accessor &Rhs) =
default;
4092 host_sampled_image_accessor &
4093 operator=(host_sampled_image_accessor &&Rhs) =
default;
4095 ~host_sampled_image_accessor() =
default;
4098 return Rhs.impl == impl;
4101 return !(Rhs == *
this);
4106 size_t size() const noexcept {
return base_class::getSize().size(); }
4111 template <
typename CoordT,
4114 DataT
read(
const CoordT &Coords)
const
4115 #ifdef __SYCL_DEVICE_ONLY__
4121 return base_class::read<DataT>(Coords);
4127 : base_class{Impl} {}
4129 template <
class Obj>
4149 #ifdef __SYCL_DEVICE_ONLY__
4157 return hash<decltype(AccImplPtr)>()(AccImplPtr);
4162 template <
typename DataT,
int Dimensions, sycl::access_mode AccessMode>
4167 #ifdef __SYCL_DEVICE_ONLY__
4174 return hash<decltype(AccImplPtr)>()(AccImplPtr);
4179 template <
typename DataT,
int Dimensions>
4180 struct hash<
sycl::local_accessor<DataT, Dimensions>> {
4184 #ifdef __SYCL_DEVICE_ONLY__
4191 return hash<decltype(AccImplPtr)>()(AccImplPtr);
4198 struct hash<
sycl::unsampled_image_accessor<DataT, Dimensions, AccessMode,
4204 #ifdef __SYCL_DEVICE_ONLY__
4210 return hash<decltype(AccImplPtr)>()(AccImplPtr);
4215 template <
typename DataT,
int Dimensions, sycl::access_mode AccessMode>
4217 sycl::host_unsampled_image_accessor<DataT, Dimensions, AccessMode>> {
4223 return hash<decltype(AccImplPtr)>()(AccImplPtr);
4227 template <
typename DataT,
int Dimensions, sycl::image_target AccessTarget>
4228 struct hash<
sycl::sampled_image_accessor<DataT, Dimensions, AccessTarget>> {
4232 #ifdef __SYCL_DEVICE_ONLY__
4238 return hash<decltype(AccImplPtr)>()(AccImplPtr);
4243 template <
typename DataT,
int Dimensions>
4244 struct hash<
sycl::host_sampled_image_accessor<DataT, Dimensions>> {
4249 return hash<decltype(AccImplPtr)>()(AccImplPtr);