31 #include <type_traits>
212 class AccessorPrivateProxy;
230 template <
typename T>
233 template <
typename T>
235 typename std::is_same<ext::oneapi::accessor_property_list<>, T>;
238 constexpr
static bool value =
false;
241 template <
typename... Props>
243 constexpr
static bool value =
true;
247 constexpr
static bool value =
false;
253 template <
int NewDim,
int DefaultValue,
template <
int>
class T,
int OldDim>
256 const int CopyDims = NewDim > OldDim ? OldDim : NewDim;
257 for (
int I = 0; I < CopyDims; ++I)
258 NewObj[I] = OldObj[I];
259 for (
int I = CopyDims; I < NewDim; ++I)
260 NewObj[I] = DefaultValue;
271 constexpr
static bool IsPlaceH = IsPlaceholder == access::placeholder::true_t;
274 constexpr
static bool IsHostBuf = AccessTarget == access::target::host_buffer;
279 constexpr
static bool IsGlobalBuf =
280 AccessTarget == access::target::global_buffer;
282 constexpr
static bool IsConstantBuf =
283 AccessTarget == access::target::constant_buffer;
285 constexpr
static bool IsAccessAnyWrite =
287 AccessMode == access::mode::read_write ||
288 AccessMode == access::mode::discard_write ||
289 AccessMode == access::mode::discard_read_write;
291 constexpr
static bool IsAccessReadOnly = AccessMode == access::mode::read;
293 constexpr
static bool IsAccessReadWrite =
294 AccessMode == access::mode::read_write;
296 constexpr
static bool IsAccessAtomic = AccessMode == access::mode::atomic;
303 IsPlaceholder, PropertyListT>;
315 : MAccessor(Accessor), MIDs(IDs) {}
323 template <
int CurDims = SubDims>
325 operator[](
size_t Index) {
326 MIDs[Dims - CurDims] = Index;
330 template <
int CurDims = SubDims,
333 MIDs[Dims - CurDims] = Index;
334 return MAccessor[MIDs];
337 template <
int CurDims = SubDims>
341 MIDs[Dims - CurDims] = Index;
342 return MAccessor[MIDs];
345 template <
int CurDims = SubDims,
348 MIDs[Dims - SubDims] = Index;
349 return MAccessor[MIDs];
356 constexpr
static bool value =
360 constexpr
static bool value =
365 constexpr
static bool value =
378 #ifndef __SYCL_DEVICE_ONLY__
386 using OCLImageTy =
typename detail::opencl_image_type<
Dimensions, AccessMode,
388 OCLImageTy MImageObj;
394 void imageAccessorInit(OCLImageTy Image) { MImageObj = Image; }
398 template <
typename T1,
int T2, access::mode T3, access::placeholder T4>
401 constexpr
static bool IsHostImageAcc =
402 (AccessTarget == access::target::host_image);
404 constexpr
static bool IsImageAcc = (AccessTarget == access::target::image);
406 constexpr
static bool IsImageArrayAcc =
407 (AccessTarget == access::target::image_array);
409 constexpr
static bool IsImageAccessWriteOnly =
411 AccessMode == access::mode::discard_write);
413 constexpr
static bool IsImageAccessAnyWrite =
414 (IsImageAccessWriteOnly || AccessMode == access::mode::read_write);
416 constexpr
static bool IsImageAccessReadOnly =
417 (AccessMode == access::mode::read);
419 constexpr
static bool IsImageAccessAnyRead =
420 (IsImageAccessReadOnly || AccessMode == access::mode::read_write);
422 static_assert(std::is_same<DataT, cl_int4>::value ||
423 std::is_same<DataT, cl_uint4>::value ||
424 std::is_same<DataT, cl_float4>::value ||
425 std::is_same<DataT, cl_half4>::value,
426 "The data type of an image accessor must be only cl_int4, "
427 "cl_uint4, cl_float4 or cl_half4 from SYCL namespace");
429 static_assert(IsImageAcc || IsHostImageAcc || IsImageArrayAcc,
430 "Expected image type");
432 static_assert(IsPlaceholder == access::placeholder::false_t,
433 "Expected false as Placeholder value for image accessor.");
436 ((IsImageAcc || IsImageArrayAcc) &&
437 (IsImageAccessWriteOnly || IsImageAccessReadOnly)) ||
438 (IsHostImageAcc && (IsImageAccessAnyWrite || IsImageAccessAnyRead)),
439 "Access modes can be only read/write/discard_write for image/image_array "
440 "target accessor, or they can be only "
441 "read/write/discard_write/read_write for host_image target accessor.");
444 "Dimensions can be 1/2/3 for image accessor.");
446 template <info::device param>
447 void checkDeviceFeatureSupported(
const device &Device) {
453 #ifdef __SYCL_DEVICE_ONLY__
456 return __invoke_ImageQuerySize<sycl::vec<int, Dimensions>, OCLImageTy>(
460 size_t getElementSize()
const {
461 int ChannelType = __invoke_ImageQueryFormat<int, OCLImageTy>(MImageObj);
462 int ChannelOrder = __invoke_ImageQueryOrder<int, OCLImageTy>(MImageObj);
463 int ElementSize = getSPIRVElementSize(ChannelType, ChannelOrder);
471 throw runtime_error(
"image::getRangeInternal() is not implemented for host",
479 friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
481 #ifdef __SYCL_DEVICE_ONLY__
482 const OCLImageTy getNativeImageObj()
const {
return MImageObj; }
483 #endif // __SYCL_DEVICE_ONLY__
492 #ifdef __SYCL_DEVICE_ONLY__
504 #ifdef __SYCL_DEVICE_ONLY__
507 (void)ImageElementSize;
514 detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
515 detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
518 MImageCount(ImageRef.
size()),
533 handler &CommandGroupHandlerRef,
int ImageElementSize)
534 #ifdef __SYCL_DEVICE_ONLY__
537 (void)CommandGroupHandlerRef;
538 (void)ImageElementSize;
545 detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
546 detail::convertToArrayOfN<3, 1>(ImageRef.
get_range()),
549 MImageCount(ImageRef.
size()),
552 checkDeviceFeatureSupported<info::device::image_support>(
561 #ifndef __SYCL_DEVICE_ONLY__
580 #ifdef __SYCL_DEVICE_ONLY__
583 size_t get_count()
const {
return size(); }
584 size_t size() const noexcept {
return get_range<Dimensions>().size(); }
586 template <
int Dims = Dimensions,
typename = detail::enable_if_t<Dims == 1>>
587 range<1> get_range()
const {
588 cl_int Range = getRangeInternal();
589 return range<1>(Range);
591 template <
int Dims = Dimensions,
typename = detail::enable_if_t<Dims == 2>>
592 range<2> get_range()
const {
593 cl_int2 Range = getRangeInternal();
594 return range<2>(Range[0], Range[1]);
596 template <
int Dims = Dimensions,
typename = detail::enable_if_t<Dims == 3>>
597 range<3> get_range()
const {
598 cl_int3 Range = getRangeInternal();
599 return range<3>(Range[0], Range[1], Range[2]);
604 size_t get_count()
const {
return size(); };
605 size_t size() const noexcept {
return MImageCount; };
607 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 0)>>
609 return detail::convertToArrayOfN<Dims, 1>(getAccessRange());
618 template <
typename CoordT,
int Dims =
Dimensions,
622 ((IsImageAcc && IsImageAccessReadOnly) ||
623 (IsHostImageAcc && IsImageAccessAnyRead))>>
624 DataT
read(
const CoordT &Coords)
const {
625 #ifdef __SYCL_DEVICE_ONLY__
626 return __invoke__ImageRead<DataT, OCLImageTy, CoordT>(MImageObj, Coords);
628 sampler Smpl(coordinate_normalization_mode::unnormalized,
629 addressing_mode::none, filtering_mode::nearest);
630 return read<CoordT, Dims>(Coords, Smpl);
638 template <
typename CoordT,
int Dims =
Dimensions,
641 ((IsImageAcc && IsImageAccessReadOnly) ||
642 (IsHostImageAcc && IsImageAccessAnyRead))>>
644 #ifdef __SYCL_DEVICE_ONLY__
645 return __invoke__ImageReadSampler<DataT, OCLImageTy, CoordT>(
646 MImageObj, Coords, Smpl.impl.m_Sampler);
648 return imageReadSamplerHostImpl<CoordT, DataT>(
649 Coords, Smpl, getAccessRange() ,
650 getOffset() , MImgChannelType, MImgChannelOrder,
651 AccessorBaseHost::getPtr() ,
652 AccessorBaseHost::getElemSize());
662 template <
typename CoordT,
int Dims =
Dimensions,
666 ((IsImageAcc && IsImageAccessWriteOnly) ||
667 (IsHostImageAcc && IsImageAccessAnyWrite))>>
668 void write(
const CoordT &Coords,
const DataT &Color)
const {
669 #ifdef __SYCL_DEVICE_ONLY__
670 __invoke__ImageWrite<OCLImageTy, CoordT, DataT>(MImageObj, Coords, Color);
673 AccessorBaseHost::getElemSize(), MImgChannelType,
675 AccessorBaseHost::getPtr() );
685 "Image slice cannot have more then 2 dimensions");
689 template <
typename CoordT,
690 typename CoordElemType =
693 getAdjustedCoords(
const CoordT &Coords)
const {
694 CoordElemType LastCoord = 0;
696 if (std::is_same<float, CoordElemType>::value) {
699 MIdx /
static_cast<float>(Size.template swizzle<Dimensions>());
704 sycl::vec<CoordElemType, Dimensions> LeftoverCoords{LastCoord};
705 sycl::vec<CoordElemType, AdjustedDims> AdjustedCoords{Coords,
707 return AdjustedCoords;
716 : MBaseAcc(BaseAcc), MIdx(Idx) {}
718 template <
typename CoordT,
int Dims =
Dimensions,
721 DataT
read(
const CoordT &Coords)
const {
722 return MBaseAcc.read(getAdjustedCoords(Coords));
725 template <
typename CoordT,
int Dims =
Dimensions,
729 return MBaseAcc.read(getAdjustedCoords(Coords), Smpl);
732 template <
typename CoordT,
int Dims =
Dimensions,
735 void write(
const CoordT &Coords,
const DataT &Color)
const {
736 return MBaseAcc.write(getAdjustedCoords(Coords), Color);
739 #ifdef __SYCL_DEVICE_ONLY__
741 size_t get_count()
const {
return size(); }
742 size_t size() const noexcept {
return get_range<Dimensions>().size(); }
744 template <
int Dims = Dimensions,
typename = detail::enable_if_t<Dims == 1>>
745 range<1> get_range()
const {
746 cl_int2 Count = MBaseAcc.getRangeInternal();
747 return range<1>(Count.x());
749 template <
int Dims = Dimensions,
typename = detail::enable_if_t<Dims == 2>>
750 range<2> get_range()
const {
751 cl_int3 Count = MBaseAcc.getRangeInternal();
752 return range<2>(Count.x(), Count.y());
758 size_t get_count()
const {
return size(); }
760 return MBaseAcc.MImageCount / MBaseAcc.getAccessRange()[
Dimensions];
766 return detail::convertToArrayOfN<Dims, 1>(MBaseAcc.getAccessRange());
787 typename PropertyListT>
789 #ifndef __SYCL_DEVICE_ONLY__
790 public detail::AccessorBaseHost,
792 public detail::accessor_common<DataT, Dimensions, AccessMode, AccessTarget,
793 IsPlaceholder, PropertyListT> {
795 static_assert((AccessTarget == access::target::global_buffer ||
796 AccessTarget == access::target::constant_buffer ||
797 AccessTarget == access::target::host_buffer),
798 "Expected buffer type");
800 static_assert((AccessTarget == access::target::global_buffer ||
801 AccessTarget == access::target::host_buffer) ||
802 (AccessTarget == access::target::constant_buffer &&
803 AccessMode == access::mode::read),
804 "Access mode can be only read for constant buffers");
806 static_assert(detail::IsPropertyListT<PropertyListT>::value,
807 "PropertyListT must be accessor_property_list");
809 using AccessorCommonT =
810 detail::accessor_common<DataT,
Dimensions, AccessMode, AccessTarget,
811 IsPlaceholder, PropertyListT>;
815 using AccessorCommonT::AS;
816 using AccessorCommonT::IsAccessAnyWrite;
817 using AccessorCommonT::IsAccessReadOnly;
818 using AccessorCommonT::IsConstantBuf;
819 using AccessorCommonT::IsGlobalBuf;
820 using AccessorCommonT::IsHostBuf;
821 using AccessorCommonT::IsPlaceH;
834 #ifdef __SYCL_DEVICE_ONLY__
838 #endif // __SYCL_DEVICE_ONLY__
843 for (
int I = 0; I < Dims; ++I) {
844 Result = Result * getMemoryRange()[I] + Id[I];
845 #if __cplusplus >= 201703L
846 if constexpr (!(PropertyListT::template has_property<
848 Result += getOffset()[I];
851 Result += getOffset()[I];
858 return std::is_same<T, DataT>::value && (Dims > 0) && (Dims ==
Dimensions);
864 if (PropertyList.template has_property<property::no_init>() ||
865 PropertyList.template has_property<property::noinit>()) {
867 AdjustedMode = access::mode::discard_write;
868 }
else if (AdjustedMode == access::mode::read_write) {
869 AdjustedMode = access::mode::discard_read_write;
876 #if __cplusplus >= 201703L
878 template <
typename TagT>
static constexpr
bool IsValidTag() {
879 return std::is_same<TagT, mode_tag_t<AccessMode>>::value ||
886 #ifdef __SYCL_DEVICE_ONLY__
888 id<AdjustedDim> &getOffset() {
return impl.Offset; }
889 range<AdjustedDim> &getAccessRange() {
return impl.AccessRange; }
890 range<AdjustedDim> &getMemoryRange() {
return impl.MemRange; }
892 const id<AdjustedDim> &getOffset()
const {
return impl.Offset; }
893 const range<AdjustedDim> &getAccessRange()
const {
return impl.AccessRange; }
894 const range<AdjustedDim> &getMemoryRange()
const {
return impl.MemRange; }
896 detail::AccessorImplDevice<AdjustedDim> impl;
899 ConcreteASPtrType MData;
903 const ConcreteASPtrType getNativeImageObj()
const {
return MData; }
905 void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
906 range<AdjustedDim> MemRange, id<AdjustedDim> Offset) {
909 for (
int I = 0; I < AdjustedDim; ++I) {
910 #if __cplusplus >= 201703L
911 if constexpr (!(PropertyListT::template has_property<
912 sycl::ext::oneapi::property::no_offset>())) {
913 getOffset()[I] = Offset[I];
916 getOffset()[I] = Offset[I];
918 getAccessRange()[I] = AccessRange[I];
919 getMemoryRange()[I] = MemRange[I];
923 if (1 == AdjustedDim)
924 #if __cplusplus >= 201703L
925 if constexpr (!(PropertyListT::template has_property<
926 sycl::ext::oneapi::property::no_offset>())) {
937 void __init_esimd(ConcreteASPtrType Ptr) { MData = Ptr; }
939 ConcreteASPtrType getQualifiedPtr()
const {
return MData; }
945 #ifndef __SYCL_DEVICE_ONLY__
946 using AccessorBaseHost::impl;
952 : impl({}, detail::InitializedVal<AdjustedDim, range>::template get<0>(),
953 detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
956 using AccessorBaseHost::getAccessRange;
957 using AccessorBaseHost::getMemoryRange;
958 using AccessorBaseHost::getOffset;
960 char padding[
sizeof(detail::AccessorImplDevice<AdjustedDim>) +
964 return reinterpret_cast<PtrType>(AccessorBaseHost::getPtr());
967 #endif // __SYCL_DEVICE_ONLY__
971 friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
1002 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1005 std::is_same<T, DataT>::value && Dims == 0 &&
1006 ((!IsPlaceH && IsHostBuf) ||
1007 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> * =
nullptr>
1012 #ifdef __SYCL_DEVICE_ONLY__
1017 {0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1018 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1019 getAdjustedMode(PropertyList),
1021 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1022 preScreenAccessor(BufferRef.
size(), PropertyList);
1026 detail::AccessorBaseHost::impl.
get(),
1027 AccessTarget, AccessMode, CodeLoc);
1031 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1032 typename... PropTypes,
1034 detail::IsCxPropertyList<PropertyListT>::value &&
1035 std::is_same<T, DataT>::value && Dims == 0 &&
1036 ((!IsPlaceH && IsHostBuf) ||
1037 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> * =
nullptr>
1043 #ifdef __SYCL_DEVICE_ONLY__
1048 {0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1049 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1050 getAdjustedMode(PropertyList),
1052 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1053 preScreenAccessor(BufferRef.
size(), PropertyList);
1057 detail::AccessorBaseHost::impl.
get(),
1058 AccessTarget, AccessMode, CodeLoc);
1062 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1064 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1065 std::is_same<T, DataT>::value && (Dims == 0) &&
1066 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1071 #ifdef __SYCL_DEVICE_ONLY__
1073 (void)CommandGroupHandler;
1078 {0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1079 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1080 getAdjustedMode(PropertyList),
1082 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1083 preScreenAccessor(BufferRef.
size(), PropertyList);
1086 detail::AccessorBaseHost::impl.
get(),
1087 AccessTarget, AccessMode, CodeLoc);
1091 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1092 typename... PropTypes,
1094 detail::IsCxPropertyList<PropertyListT>::value &&
1095 std::is_same<T, DataT>::value && (Dims == 0) &&
1096 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1102 #ifdef __SYCL_DEVICE_ONLY__
1104 (void)CommandGroupHandler;
1109 {0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
1110 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1111 getAdjustedMode(PropertyList),
1113 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1114 preScreenAccessor(BufferRef.
size(), PropertyList);
1117 detail::AccessorBaseHost::impl.
get(),
1118 AccessTarget, AccessMode, CodeLoc);
1122 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1124 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1125 IsSameAsBuffer<T, Dims>() &&
1126 ((!IsPlaceH && IsHostBuf) ||
1127 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1132 #ifdef __SYCL_DEVICE_ONLY__
1139 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1140 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1141 getAdjustedMode(PropertyList),
1143 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1144 preScreenAccessor(BufferRef.
size(), PropertyList);
1148 detail::AccessorBaseHost::impl.
get(),
1149 AccessTarget, AccessMode, CodeLoc);
1153 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1154 typename... PropTypes,
1156 detail::IsCxPropertyList<PropertyListT>::value &&
1157 IsSameAsBuffer<T, Dims>() &&
1158 ((!IsPlaceH && IsHostBuf) ||
1159 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1165 #ifdef __SYCL_DEVICE_ONLY__
1172 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1173 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1174 getAdjustedMode(PropertyList),
1176 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1177 preScreenAccessor(BufferRef.
size(), PropertyList);
1181 detail::AccessorBaseHost::impl.
get(),
1182 AccessTarget, AccessMode, CodeLoc);
1186 #if __cplusplus >= 201703L
1188 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1191 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1192 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1193 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1195 buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1196 const property_list &PropertyList = {},
1197 const detail::code_location CodeLoc = detail::code_location::current())
1198 :
accessor(BufferRef, PropertyList, CodeLoc) {
1202 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1203 typename TagT,
typename... PropTypes,
1205 detail::IsCxPropertyList<PropertyListT>::value &&
1206 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1207 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1209 buffer<T, Dims, AllocatorT> &BufferRef, TagT,
1210 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1212 const detail::code_location CodeLoc = detail::code_location::current())
1213 :
accessor(BufferRef, PropertyList, CodeLoc) {
1218 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1220 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1221 IsSameAsBuffer<T, Dims>() &&
1222 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1227 #ifdef __SYCL_DEVICE_ONLY__
1229 (void)CommandGroupHandler;
1235 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1236 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1237 getAdjustedMode(PropertyList),
1239 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1240 preScreenAccessor(BufferRef.
size(), PropertyList);
1243 detail::AccessorBaseHost::impl.
get(),
1244 AccessTarget, AccessMode, CodeLoc);
1248 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1249 typename... PropTypes,
1251 detail::IsCxPropertyList<PropertyListT>::value &&
1252 IsSameAsBuffer<T, Dims>() &&
1253 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1259 #ifdef __SYCL_DEVICE_ONLY__
1261 (void)CommandGroupHandler;
1267 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1268 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1269 getAdjustedMode(PropertyList),
1271 BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
1272 preScreenAccessor(BufferRef.
size(), PropertyList);
1275 detail::AccessorBaseHost::impl.
get(),
1276 AccessTarget, AccessMode, CodeLoc);
1280 #if __cplusplus >= 201703L
1282 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1285 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1286 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1287 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1289 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1290 TagT,
const property_list &PropertyList = {},
1291 const detail::code_location CodeLoc = detail::code_location::current())
1292 :
accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1296 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1297 typename TagT,
typename... PropTypes,
1299 detail::IsCxPropertyList<PropertyListT>::value &&
1300 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1301 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1303 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1305 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1307 const detail::code_location CodeLoc = detail::code_location::current())
1308 :
accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1314 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1316 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1317 IsSameAsBuffer<T, Dims>() &&
1318 ((!IsPlaceH && IsHostBuf) ||
1319 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1324 :
accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1326 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1327 typename... PropTypes,
1329 detail::IsCxPropertyList<PropertyListT>::value &&
1330 IsSameAsBuffer<T, Dims>() &&
1331 ((!IsPlaceH && IsHostBuf) ||
1332 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1338 :
accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1340 #if __cplusplus >= 201703L
1342 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1345 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1346 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1347 (IsGlobalBuf || IsConstantBuf)>>
1349 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1350 TagT,
const property_list &PropertyList = {},
1351 const detail::code_location CodeLoc = detail::code_location::current())
1352 :
accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1356 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1357 typename TagT,
typename... PropTypes,
1359 detail::IsCxPropertyList<PropertyListT>::value &&
1360 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1361 (IsGlobalBuf || IsConstantBuf)>>
1363 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1365 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1367 const detail::code_location CodeLoc = detail::code_location::current())
1368 :
accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1373 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1375 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1376 IsSameAsBuffer<T, Dims>() &&
1377 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1382 :
accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1385 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1386 typename... PropTypes,
1388 detail::IsCxPropertyList<PropertyListT>::value &&
1389 IsSameAsBuffer<T, Dims>() &&
1390 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1397 :
accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1400 #if __cplusplus >= 201703L
1402 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1405 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1406 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1407 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1409 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1410 range<Dimensions> AccessRange, TagT,
1411 const property_list &PropertyList = {},
1412 const detail::code_location CodeLoc = detail::code_location::current())
1413 :
accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1418 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1419 typename TagT,
typename... PropTypes,
1421 detail::IsCxPropertyList<PropertyListT>::value &&
1422 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1423 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1425 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1426 range<Dimensions> AccessRange, TagT,
1427 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1429 const detail::code_location CodeLoc = detail::code_location::current())
1430 :
accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1436 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1438 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1439 IsSameAsBuffer<T, Dims>() &&
1440 ((!IsPlaceH && IsHostBuf) ||
1441 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1446 #ifdef __SYCL_DEVICE_ONLY__
1447 : impl(AccessOffset, AccessRange, BufferRef.
get_range()) {
1451 : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1452 detail::convertToArrayOfN<3, 1>(AccessRange),
1453 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1454 getAdjustedMode(PropertyList),
1456 sizeof(DataT), BufferRef.OffsetInBytes,
1457 BufferRef.IsSubBuffer) {
1458 preScreenAccessor(BufferRef.
size(), PropertyList);
1461 throw sycl::invalid_object_error(
1462 "accessor with requested offset and range would exceed the bounds of "
1469 detail::AccessorBaseHost::impl.
get(),
1470 AccessTarget, AccessMode, CodeLoc);
1474 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1475 typename... PropTypes,
1477 detail::IsCxPropertyList<PropertyListT>::value &&
1478 IsSameAsBuffer<T, Dims>() &&
1479 ((!IsPlaceH && IsHostBuf) ||
1480 (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
1487 #ifdef __SYCL_DEVICE_ONLY__
1488 : impl(AccessOffset, AccessRange, BufferRef.
get_range()) {
1492 : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1493 detail::convertToArrayOfN<3, 1>(AccessRange),
1494 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1495 getAdjustedMode(PropertyList),
1497 sizeof(DataT), BufferRef.OffsetInBytes,
1498 BufferRef.IsSubBuffer) {
1499 preScreenAccessor(BufferRef.
size(), PropertyList);
1502 throw sycl::invalid_object_error(
1503 "accessor with requested offset and range would exceed the bounds of "
1510 detail::AccessorBaseHost::impl.
get(),
1511 AccessTarget, AccessMode, CodeLoc);
1515 #if __cplusplus >= 201703L
1517 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1520 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1521 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1522 (IsGlobalBuf || IsConstantBuf)>>
1524 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1525 id<Dimensions> AccessOffset, TagT,
const property_list &PropertyList = {},
1526 const detail::code_location CodeLoc = detail::code_location::current())
1527 :
accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
1531 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1532 typename TagT,
typename... PropTypes,
1534 detail::IsCxPropertyList<PropertyListT>::value &&
1535 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
1536 (IsGlobalBuf || IsConstantBuf)>>
1538 buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
1539 id<Dimensions> AccessOffset, TagT,
1540 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1542 const detail::code_location CodeLoc = detail::code_location::current())
1543 :
accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
1548 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1550 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1551 IsSameAsBuffer<T, Dims>() &&
1552 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1558 #ifdef __SYCL_DEVICE_ONLY__
1559 : impl(AccessOffset, AccessRange, BufferRef.
get_range()) {
1560 (void)CommandGroupHandler;
1564 : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1565 detail::convertToArrayOfN<3, 1>(AccessRange),
1566 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1567 getAdjustedMode(PropertyList),
1569 sizeof(DataT), BufferRef.OffsetInBytes,
1570 BufferRef.IsSubBuffer) {
1571 preScreenAccessor(BufferRef.
size(), PropertyList);
1574 throw sycl::invalid_object_error(
1575 "accessor with requested offset and range would exceed the bounds of "
1581 detail::AccessorBaseHost::impl.
get(),
1582 AccessTarget, AccessMode, CodeLoc);
1586 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1587 typename... PropTypes,
1589 detail::IsCxPropertyList<PropertyListT>::value &&
1590 IsSameAsBuffer<T, Dims>() &&
1591 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
1598 #ifdef __SYCL_DEVICE_ONLY__
1599 : impl(AccessOffset, AccessRange, BufferRef.
get_range()) {
1600 (void)CommandGroupHandler;
1604 : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
1605 detail::convertToArrayOfN<3, 1>(AccessRange),
1606 detail::convertToArrayOfN<3, 1>(BufferRef.
get_range()),
1607 getAdjustedMode(PropertyList),
1609 sizeof(DataT), BufferRef.OffsetInBytes,
1610 BufferRef.IsSubBuffer) {
1611 preScreenAccessor(BufferRef.
size(), PropertyList);
1614 throw sycl::invalid_object_error(
1615 "accessor with requested offset and range would exceed the bounds of "
1621 detail::AccessorBaseHost::impl.
get(),
1622 AccessTarget, AccessMode, CodeLoc);
1626 #if __cplusplus >= 201703L
1628 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1631 detail::IsRunTimePropertyListT<PropertyListT>::value &&
1632 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1633 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1635 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1636 range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
1637 const property_list &PropertyList = {},
1638 const detail::code_location CodeLoc = detail::code_location::current())
1639 :
accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1640 PropertyList, CodeLoc) {
1644 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
1645 typename TagT,
typename... PropTypes,
1647 detail::IsCxPropertyList<PropertyListT>::value &&
1648 IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1649 (IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
1651 buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
1652 range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
1653 const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
1655 const detail::code_location CodeLoc = detail::code_location::current())
1656 :
accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1657 PropertyList, CodeLoc) {
1662 template <
typename... NewPropsT>
1667 #ifdef __SYCL_DEVICE_ONLY__
1674 "Conversion is only available for accessor_property_list");
1676 PropertyListT::template areSameCompileTimeProperties<NewPropsT...>(),
1677 "Compile-time-constant properties must be the same");
1678 #ifndef __SYCL_DEVICE_ONLY__
1680 AccessTarget, AccessMode, CodeLoc);
1686 size_t get_size()
const {
return getAccessRange().size() *
sizeof(DataT); }
1689 size_t get_count()
const {
return size(); }
1690 size_t size() const noexcept {
return getAccessRange().size(); }
1692 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 0)>>
1694 return detail::convertToArrayOfN<Dimensions, 1>(getAccessRange());
1697 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 0)>>
1699 #if __cplusplus >= 201703L
1701 !(PropertyListT::template has_property<
1703 "Accessor has no_offset property, get_offset() can not be used");
1705 return detail::convertToArrayOfN<Dimensions, 0>(getOffset());
1708 template <
int Dims =
Dimensions,
typename RefT = RefType,
1710 !std::is_const<RefT>::value>>
1713 return *(getQualifiedPtr() + LinearIndex);
1720 return *(getQualifiedPtr() + LinearIndex);
1727 return getQualifiedPtr()[LinearIndex];
1730 template <
int Dims = Dimensions>
1734 return getQualifiedPtr()[LinearIndex];
1737 template <
int Dims = Dimensions>
1739 AccessMode == access::mode::atomic,
1740 #ifdef __ENABLE_USM_ADDR_SPACE__
1747 return atomic<DataT, AS>(
1751 template <
int Dims = Dimensions>
1756 return atomic<DataT, AS>(
1760 template <
int Dims = Dimensions>
1765 return atomic<DataT, AS>(
1768 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 1)>>
1769 typename AccessorCommonT::template AccessorSubscript<Dims - 1>
1776 access::target::host_buffer>>
1778 return getPointerAdjusted();
1790 access::target::constant_buffer>>
1804 PtrType getPointerAdjusted()
const {
1805 #ifdef __SYCL_DEVICE_ONLY__
1806 if (1 == AdjustedDim)
1807 return getQualifiedPtr() - impl.Offset[0];
1809 return getQualifiedPtr();
1812 void preScreenAccessor(
const size_t elemInBuffer,
1813 const PropertyListT &PropertyList) {
1815 if (!IsHostBuf && elemInBuffer == 0)
1816 throw sycl::invalid_object_error(
1817 "SYCL buffer size is zero. To create a device accessor, SYCL "
1818 "buffer size must be greater than zero.",
1822 if (PropertyList.template has_property<property::no_init>() &&
1823 AccessMode == access::mode::read) {
1824 throw sycl::invalid_object_error(
1825 "accessor would cannot be both read_only and no_init",
1830 #if __cplusplus >= 201703L
1831 template <
typename... PropTypes>
1832 void adjustAccPropsInBuf(detail::SYCLMemObjI *SYCLMemObject) {
1833 if constexpr (PropertyListT::template has_property<
1834 sycl::ext::intel::property::buffer_location>()) {
1836 sycl::ext::intel::property::buffer_location>())
1838 property_list PropList{
1839 sycl::property::buffer::detail::buffer_location(location)};
1840 detail::SYCLMemObjT *SYCLMemObjectT =
1841 dynamic_cast<detail::SYCLMemObjT *
>(SYCLMemObject);
1842 SYCLMemObjectT->addOrReplaceAccessorProperties(PropList);
1844 deleteAccPropsFromBuf(SYCLMemObject);
1848 void deleteAccPropsFromBuf(detail::SYCLMemObjI *SYCLMemObject) {
1849 detail::SYCLMemObjT *SYCLMemObjectT =
1850 dynamic_cast<detail::SYCLMemObjT *
>(SYCLMemObject);
1851 SYCLMemObjectT->deleteAccessorProperty(
1857 #if __cplusplus >= 201703L
1859 template <
typename DataT,
int Dimensions,
typename AllocatorT>
1860 accessor(buffer<DataT, Dimensions, AllocatorT>)
1862 access::placeholder::true_t>;
1864 template <
typename DataT,
int Dimensions,
typename AllocatorT,
1866 accessor(buffer<DataT, Dimensions, AllocatorT>,
1867 const ext::oneapi::accessor_property_list<PropsT...> &)
1869 access::placeholder::true_t,
1870 ext::oneapi::accessor_property_list<PropsT...>>;
1872 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1>
1873 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1)
1874 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
1876 access::placeholder::true_t>;
1878 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
1880 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1,
1881 const ext::oneapi::accessor_property_list<PropsT...> &)
1882 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
1884 access::placeholder::true_t,
1885 ext::oneapi::accessor_property_list<PropsT...>>;
1887 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
1889 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2)
1890 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
1892 access::placeholder::true_t>;
1894 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
1895 typename Type2,
typename... PropsT>
1896 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2,
1897 const ext::oneapi::accessor_property_list<PropsT...> &)
1898 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
1900 access::placeholder::true_t,
1901 ext::oneapi::accessor_property_list<PropsT...>>;
1903 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
1904 typename Type2,
typename Type3>
1905 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3)
1906 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
1908 access::placeholder::true_t>;
1910 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
1911 typename Type2,
typename Type3,
typename... PropsT>
1912 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3,
1913 const ext::oneapi::accessor_property_list<PropsT...> &)
1914 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
1916 access::placeholder::true_t,
1917 ext::oneapi::accessor_property_list<PropsT...>>;
1919 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
1920 typename Type2,
typename Type3,
typename Type4>
1921 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
1922 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
1924 access::placeholder::true_t>;
1926 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
1927 typename Type2,
typename Type3,
typename Type4,
typename... PropsT>
1928 accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
1929 const ext::oneapi::accessor_property_list<PropsT...> &)
1930 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
1932 access::placeholder::true_t,
1933 ext::oneapi::accessor_property_list<PropsT...>>;
1935 template <
typename DataT,
int Dimensions,
typename AllocatorT>
1936 accessor(buffer<DataT, Dimensions, AllocatorT>, handler)
1938 access::placeholder::false_t>;
1940 template <
typename DataT,
int Dimensions,
typename AllocatorT,
1942 accessor(buffer<DataT, Dimensions, AllocatorT>, handler,
1943 const ext::oneapi::accessor_property_list<PropsT...> &)
1945 access::placeholder::false_t,
1946 ext::oneapi::accessor_property_list<PropsT...>>;
1948 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1>
1949 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1)
1950 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
1952 access::placeholder::false_t>;
1954 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
1956 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1,
1957 const ext::oneapi::accessor_property_list<PropsT...> &)
1958 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type1>(),
1960 access::placeholder::false_t,
1961 ext::oneapi::accessor_property_list<PropsT...>>;
1963 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
1965 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2)
1966 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
1968 access::placeholder::false_t>;
1970 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
1971 typename Type2,
typename... PropsT>
1972 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2,
1973 const ext::oneapi::accessor_property_list<PropsT...> &)
1974 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type1, Type2>(),
1976 access::placeholder::false_t,
1977 ext::oneapi::accessor_property_list<PropsT...>>;
1979 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
1980 typename Type2,
typename Type3>
1981 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3)
1982 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
1984 access::placeholder::false_t>;
1986 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
1987 typename Type2,
typename Type3,
typename... PropsT>
1988 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
1989 const ext::oneapi::accessor_property_list<PropsT...> &)
1990 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type2, Type3>(),
1992 access::placeholder::false_t,
1993 ext::oneapi::accessor_property_list<PropsT...>>;
1995 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
1996 typename Type2,
typename Type3,
typename Type4>
1997 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
1999 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2001 access::placeholder::false_t>;
2003 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2004 typename Type2,
typename Type3,
typename Type4,
typename... PropsT>
2005 accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
2006 Type4,
const ext::oneapi::accessor_property_list<PropsT...> &)
2007 -> accessor<DataT, Dimensions, detail::deduceAccessMode<Type3, Type4>(),
2009 access::placeholder::false_t,
2010 ext::oneapi::accessor_property_list<PropsT...>>;
2019 access::target::local, IsPlaceholder> :
2020 #ifndef __SYCL_DEVICE_ONLY__
2024 access::target::local, IsPlaceholder> {
2030 access::target::local, IsPlaceholder>;
2032 using AccessorCommonT::AS;
2033 using AccessorCommonT::IsAccessAnyWrite;
2043 #ifdef __SYCL_DEVICE_ONLY__
2049 void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
2050 range<AdjustedDim>, id<AdjustedDim>) {
2053 for (
int I = 0; I < AdjustedDim; ++I)
2054 getSize()[I] = AccessRange[I];
2060 : impl(detail::InitializedVal<AdjustedDim, range>::template
get<0>()) {}
2063 ConcreteASPtrType getQualifiedPtr()
const {
return MData; }
2065 ConcreteASPtrType MData;
2069 char padding[
sizeof(detail::LocalAccessorBaseDevice<AdjustedDim>) +
2071 using detail::LocalAccessorBaseHost::getSize;
2074 return reinterpret_cast<PtrType>(LocalAccessorBaseHost::getPtr());
2077 #endif // __SYCL_DEVICE_ONLY__
2082 for (
int I = 0; I < Dims; ++I)
2083 Result = Result * getSize()[I] + Id[I];
2092 template <
int Dims = Dimensions,
typename = detail::enable_if_t<Dims == 0>>
2094 detail::code_location::current())
2095 #ifdef __SYCL_DEVICE_ONLY__
2098 : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim,
sizeof(DataT)) {
2100 access::target::local, AccessMode, CodeLoc);
2105 typename = detail::enable_if_t<Dims == 0>>
2108 detail::code_location::current())
2109 #ifdef __SYCL_DEVICE_ONLY__
2114 : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim,
sizeof(DataT)) {
2117 access::target::local, AccessMode, CodeLoc);
2121 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 0)>>
2125 #ifdef __SYCL_DEVICE_ONLY__
2126 : impl(AllocationSize){}
2129 AdjustedDim,
sizeof(DataT)) {
2131 access::target::local, AccessMode, CodeLoc);
2140 detail::code_location::current())
2141 #ifdef __SYCL_DEVICE_ONLY__
2142 : impl(AllocationSize) {
2147 AdjustedDim,
sizeof(DataT)) {
2150 access::target::local, AccessMode, CodeLoc);
2154 size_t get_size()
const {
return getSize().size() *
sizeof(DataT); }
2157 size_t get_count()
const {
return size(); }
2158 size_t size() const noexcept {
return getSize().size(); }
2160 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 0)>>
2162 return detail::convertToArrayOfN<Dims, 1>(getSize());
2168 return *getQualifiedPtr();
2175 return getQualifiedPtr()[LinearIndex];
2181 return getQualifiedPtr()[Index];
2184 template <
int Dims = Dimensions>
2186 Dims == 0 && AccessMode == access::mode::atomic, atomic<DataT, AS>>()
2191 template <
int Dims = Dimensions>
2196 return atomic<DataT, AS>(
2200 template <
int Dims = Dimensions>
2207 template <
int Dims = Dimensions,
typename = detail::enable_if_t<(Dims > 1)>>
2208 typename AccessorCommonT::template AccessorSubscript<Dims - 1>
2229 access::target::image, IsPlaceholder>
2231 access::target::image, IsPlaceholder> {
2233 template <
typename AllocatorT>
2238 Image, CommandGroupHandler,
2240 #ifndef __SYCL_DEVICE_ONLY__
2242 access::target::image);
2246 template <
typename AllocatorT>
2251 Image, CommandGroupHandler,
2254 #ifndef __SYCL_DEVICE_ONLY__
2256 access::target::image);
2259 #ifdef __SYCL_DEVICE_ONLY__
2262 typename detail::opencl_image_type<
Dimensions, AccessMode,
2263 access::target::image>::type;
2267 void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
2270 void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
2274 accessor() =
default;
2290 access::target::host_image, IsPlaceholder> {
2292 template <
typename AllocatorT>
2295 access::
target::host_image, IsPlaceholder>(
2298 template <
typename AllocatorT>
2302 access::
target::host_image, IsPlaceholder>(
2319 access::target::image_array, IsPlaceholder>
2321 access::target::image, IsPlaceholder> {
2322 #ifdef __SYCL_DEVICE_ONLY__
2325 typename detail::opencl_image_type<
Dimensions + 1, AccessMode,
2326 access::target::image>::type;
2330 void __init(OCLImageTy Image) { this->imageAccessorInit(Image); }
2333 void __init_esimd(OCLImageTy Image) { this->imageAccessorInit(Image); }
2340 template <
typename AllocatorT>
2345 Image, CommandGroupHandler,
2347 #ifndef __SYCL_DEVICE_ONLY__
2349 access::target::image_array);
2353 template <
typename AllocatorT>
2358 Image, CommandGroupHandler,
2361 #ifndef __SYCL_DEVICE_ONLY__
2363 access::target::image_array);
2370 IsPlaceholder>(*
this, Index);
2374 template <
typename DataT,
int Dimensions = 1,
2377 :
public accessor<DataT, Dimensions, AccessMode, target::host_buffer,
2378 access::placeholder::false_t> {
2381 access::placeholder::false_t>;
2386 return std::is_same<T, DataT>::value && (Dims > 0) && (Dims ==
Dimensions);
2389 #if __cplusplus >= 201703L
2391 template <
typename TagT>
static constexpr
bool IsValidTag() {
2392 return std::is_same<TagT, mode_tag_t<AccessMode>>::value;
2402 AccessorT::__init(Ptr, AccessRange, MemRange, Offset);
2431 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2433 std::is_same<T, DataT>::value && Dims == 0>>
2438 : AccessorT(BufferRef, PropertyList, CodeLoc) {}
2440 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2441 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2446 : AccessorT(BufferRef, PropertyList, CodeLoc) {}
2448 #if __cplusplus >= 201703L
2450 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2451 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2453 buffer<DataT, Dimensions, AllocatorT> &BufferRef, mode_tag_t<AccessMode>,
2454 const property_list &PropertyList = {},
2455 const detail::code_location CodeLoc = detail::code_location::current())
2456 : host_accessor(BufferRef, PropertyList, CodeLoc) {}
2460 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2461 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2466 : AccessorT(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
2468 #if __cplusplus >= 201703L
2470 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2471 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2473 buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2474 handler &CommandGroupHandler, mode_tag_t<AccessMode>,
2475 const property_list &PropertyList = {},
2476 const detail::code_location CodeLoc = detail::code_location::current())
2477 : host_accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
2481 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2482 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2487 : AccessorT(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
2489 #if __cplusplus >= 201703L
2491 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2492 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2494 buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2495 range<Dimensions> AccessRange, mode_tag_t<AccessMode>,
2496 const property_list &PropertyList = {},
2497 const detail::code_location CodeLoc = detail::code_location::current())
2498 : host_accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
2502 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2503 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2509 : AccessorT(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
2512 #if __cplusplus >= 201703L
2514 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2515 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2517 buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2518 handler &CommandGroupHandler, range<Dimensions> AccessRange,
2519 mode_tag_t<AccessMode>,
const property_list &PropertyList = {},
2520 const detail::code_location CodeLoc = detail::code_location::current())
2521 : host_accessor(BufferRef, CommandGroupHandler, AccessRange, {},
2522 PropertyList, CodeLoc) {}
2526 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2527 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2533 : AccessorT(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
2536 #if __cplusplus >= 201703L
2538 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2539 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2541 buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2542 range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
2543 mode_tag_t<AccessMode>,
const property_list &PropertyList = {},
2544 const detail::code_location CodeLoc = detail::code_location::current())
2545 : host_accessor(BufferRef, AccessRange, AccessOffset, PropertyList,
2550 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2551 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2557 : AccessorT(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2558 PropertyList, CodeLoc) {}
2560 #if __cplusplus >= 201703L
2562 template <
typename T = DataT,
int Dims =
Dimensions,
typename AllocatorT,
2563 typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
2565 buffer<DataT, Dimensions, AllocatorT> &BufferRef,
2566 handler &CommandGroupHandler, range<Dimensions> AccessRange,
2567 id<Dimensions> AccessOffset, mode_tag_t<AccessMode>,
2568 const property_list &PropertyList = {},
2569 const detail::code_location CodeLoc = detail::code_location::current())
2570 : host_accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
2571 PropertyList, CodeLoc) {}
2576 #if __cplusplus >= 201703L
2578 template <
typename DataT,
int Dimensions,
typename AllocatorT>
2579 host_accessor(buffer<DataT, Dimensions, AllocatorT>)
2580 -> host_accessor<DataT, Dimensions, access::mode::read_write>;
2582 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1>
2583 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1)
2585 detail::deduceAccessMode<Type1, Type1>()>;
2587 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2589 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2)
2591 detail::deduceAccessMode<Type1, Type2>()>;
2593 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2594 typename Type2,
typename Type3>
2595 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3)
2597 detail::deduceAccessMode<Type2, Type3>()>;
2599 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2600 typename Type2,
typename Type3,
typename Type4>
2601 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
2603 detail::deduceAccessMode<Type3, Type4>()>;
2605 template <
typename DataT,
int Dimensions,
typename AllocatorT,
typename Type1,
2606 typename Type2,
typename Type3,
typename Type4,
typename Type5>
2607 host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
2609 detail::deduceAccessMode<Type4, Type5>()>;
2620 struct hash<
cl::sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
2623 AccessTarget, IsPlaceholder>;
2626 #ifdef __SYCL_DEVICE_ONLY__
2634 return hash<decltype(AccImplPtr)>()(AccImplPtr);