Skip to content

[SYCL] Optimize operator[] method of host accessor #6853

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Sep 23, 2022
Merged

[SYCL] Optimize operator[] method of host accessor #6853

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
baa67ab
[SYCL] Optimize operator[] method of host accessor
romanovvlad Sep 19, 2022
eaea384
clang-format
romanovvlad Sep 22, 2022
1ee0aa0
Adjusted test compile line
romanovvlad Sep 23, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
89 changes: 75 additions & 14 deletions sycl/include/sycl/accessor.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -227,6 +227,22 @@ template <typename DataT, int Dimensions = 1,
class accessor;

namespace detail {

// A helper structure which is shared between buffer accessor and accessor_impl
// TODO: Unify with AccessorImplDevice?
struct AccHostDataT {
AccHostDataT(const sycl::id<3> &Offset, const sycl::range<3> &Range,
const sycl::range<3> &MemoryRange, void *Data = nullptr)
: MOffset(Offset), MAccessRange(Range), MMemoryRange(MemoryRange),
MData(Data) {}

sycl::id<3> MOffset;
sycl::range<3> MAccessRange;
sycl::range<3> MMemoryRange;
void *MData = nullptr;
void *Reserved = nullptr;
};

// To ensure loop unrolling is done when processing dimensions.
template <size_t... Inds, class F>
void dim_loop_impl(std::integer_sequence<size_t, Inds...>, F &&f) {
Expand Down Expand Up @@ -474,6 +490,8 @@ class __SYCL_EXPORT AccessorBaseHost {
const range<3> &getMemoryRange() const;
void *getPtr() const;

detail::AccHostDataT &getAccData();

const property_list &getPropList() const;

void *getMemoryObject() const;
Expand Down Expand Up @@ -1106,21 +1124,42 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}

#else
id<3> &getOffset() { return AccessorBaseHost::getOffset(); }
id<3> &getOffset() {
if constexpr (IsHostBuf)
return MAccData->MOffset;
else
return AccessorBaseHost::getOffset();
}

range<3> &getAccessRange() { return AccessorBaseHost::getAccessRange(); }
range<3> &getMemoryRange() { return AccessorBaseHost::getMemoryRange(); }
range<3> &getMemoryRange() {
if constexpr (IsHostBuf)
return MAccData->MMemoryRange;
else
return AccessorBaseHost::getMemoryRange();
}
void *getPtr() { return AccessorBaseHost::getPtr(); }

const id<3> &getOffset() const { return AccessorBaseHost::getOffset(); }
const id<3> &getOffset() const {
if constexpr (IsHostBuf)
return MAccData->MOffset;
else
return AccessorBaseHost::getOffset();
}
const range<3> &getAccessRange() const {
return AccessorBaseHost::getAccessRange();
}
const range<3> &getMemoryRange() const {
return AccessorBaseHost::getMemoryRange();
if constexpr (IsHostBuf)
return MAccData->MMemoryRange;
else
return AccessorBaseHost::getMemoryRange();
}

void *getPtr() const { return AccessorBaseHost::getPtr(); }

void initHostAcc() { MAccData = &getAccData(); }

// The function references helper methods required by GDB pretty-printers
void GDBMethodsAnchor() {
#ifndef NDEBUG
Expand All @@ -1131,11 +1170,17 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
#endif
}

detail::AccHostDataT *MAccData = nullptr;

char padding[sizeof(detail::AccessorImplDevice<AdjustedDim>) +
sizeof(PtrType) - sizeof(detail::AccessorBaseHost)];
sizeof(PtrType) - sizeof(detail::AccessorBaseHost) -
sizeof(MAccData)];

PtrType getQualifiedPtr() const {
return reinterpret_cast<PtrType>(AccessorBaseHost::getPtr());
if constexpr (IsHostBuf)
return reinterpret_cast<PtrType>(MAccData->MData);
else
return reinterpret_cast<PtrType>(AccessorBaseHost::getPtr());
}

#endif // __SYCL_DEVICE_ONLY__
Expand Down Expand Up @@ -1196,9 +1241,11 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
preScreenAccessor(BufferRef.size(), PropertyList);
if (!IsPlaceH)
addHostAccessorAndWait(AccessorBaseHost::impl.get());
initHostAcc();
detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
detail::AccessorBaseHost::impl.get(),
AccessTarget, AccessMode, CodeLoc);
GDBMethodsAnchor();
#endif
}

Expand Down Expand Up @@ -1227,9 +1274,11 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
preScreenAccessor(BufferRef.size(), PropertyList);
if (!IsPlaceH)
addHostAccessorAndWait(AccessorBaseHost::impl.get());
initHostAcc();
detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
detail::AccessorBaseHost::impl.get(),
AccessTarget, AccessMode, CodeLoc);
GDBMethodsAnchor();
#endif
}

Expand All @@ -1256,9 +1305,11 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
preScreenAccessor(BufferRef.size(), PropertyList);
detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
initHostAcc();
detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
detail::AccessorBaseHost::impl.get(),
AccessTarget, AccessMode, CodeLoc);
GDBMethodsAnchor();
}
#endif

Expand Down Expand Up @@ -1287,9 +1338,11 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
preScreenAccessor(BufferRef.size(), PropertyList);
detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
initHostAcc();
detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
detail::AccessorBaseHost::impl.get(),
AccessTarget, AccessMode, CodeLoc);
GDBMethodsAnchor();
}
#endif

Expand All @@ -1315,13 +1368,14 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
getAdjustedMode(PropertyList),
detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
GDBMethodsAnchor();
preScreenAccessor(BufferRef.size(), PropertyList);
if (!IsPlaceH)
addHostAccessorAndWait(AccessorBaseHost::impl.get());
initHostAcc();
detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
detail::AccessorBaseHost::impl.get(),
AccessTarget, AccessMode, CodeLoc);
GDBMethodsAnchor();
}
#endif

Expand Down Expand Up @@ -1349,13 +1403,14 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
getAdjustedMode(PropertyList),
detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
GDBMethodsAnchor();
preScreenAccessor(BufferRef.size(), PropertyList);
if (!IsPlaceH)
addHostAccessorAndWait(AccessorBaseHost::impl.get());
initHostAcc();
detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
detail::AccessorBaseHost::impl.get(),
AccessTarget, AccessMode, CodeLoc);
GDBMethodsAnchor();
}
#endif

Expand Down Expand Up @@ -1413,12 +1468,13 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
getAdjustedMode(PropertyList),
detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
GDBMethodsAnchor();
preScreenAccessor(BufferRef.size(), PropertyList);
detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
initHostAcc();
detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
detail::AccessorBaseHost::impl.get(),
AccessTarget, AccessMode, CodeLoc);
GDBMethodsAnchor();
}
#endif

Expand Down Expand Up @@ -1446,12 +1502,13 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
getAdjustedMode(PropertyList),
detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
BufferRef.OffsetInBytes, BufferRef.IsSubBuffer, PropertyList) {
GDBMethodsAnchor();
preScreenAccessor(BufferRef.size(), PropertyList);
initHostAcc();
detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
detail::AccessorBaseHost::impl.get(),
AccessTarget, AccessMode, CodeLoc);
GDBMethodsAnchor();
}
#endif

Expand Down Expand Up @@ -1633,7 +1690,6 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
detail::getSyclObjImpl(BufferRef).get(), Dimensions,
sizeof(DataT), BufferRef.OffsetInBytes,
BufferRef.IsSubBuffer, PropertyList) {
GDBMethodsAnchor();
preScreenAccessor(BufferRef.size(), PropertyList);
if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
BufferRef.get_range()))
Expand All @@ -1644,9 +1700,11 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :

if (!IsPlaceH)
addHostAccessorAndWait(AccessorBaseHost::impl.get());
initHostAcc();
detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
detail::AccessorBaseHost::impl.get(),
AccessTarget, AccessMode, CodeLoc);
GDBMethodsAnchor();
}
#endif

Expand Down Expand Up @@ -1675,7 +1733,6 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
detail::getSyclObjImpl(BufferRef).get(), Dimensions,
sizeof(DataT), BufferRef.OffsetInBytes,
BufferRef.IsSubBuffer, PropertyList) {
GDBMethodsAnchor();
preScreenAccessor(BufferRef.size(), PropertyList);
if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
BufferRef.get_range()))
Expand All @@ -1686,9 +1743,11 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :

if (!IsPlaceH)
addHostAccessorAndWait(AccessorBaseHost::impl.get());
initHostAcc();
detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
detail::AccessorBaseHost::impl.get(),
AccessTarget, AccessMode, CodeLoc);
GDBMethodsAnchor();
}
#endif

Expand Down Expand Up @@ -1748,7 +1807,6 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
detail::getSyclObjImpl(BufferRef).get(), Dimensions,
sizeof(DataT), BufferRef.OffsetInBytes,
BufferRef.IsSubBuffer, PropertyList) {
GDBMethodsAnchor();
preScreenAccessor(BufferRef.size(), PropertyList);
if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
BufferRef.get_range()))
Expand All @@ -1757,10 +1815,12 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
"the buffer",
PI_ERROR_INVALID_VALUE);

initHostAcc();
detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
detail::AccessorBaseHost::impl.get(),
AccessTarget, AccessMode, CodeLoc);
GDBMethodsAnchor();
}
#endif

Expand Down Expand Up @@ -1789,7 +1849,6 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
detail::getSyclObjImpl(BufferRef).get(), Dimensions,
sizeof(DataT), BufferRef.OffsetInBytes,
BufferRef.IsSubBuffer, PropertyList) {
GDBMethodsAnchor();
preScreenAccessor(BufferRef.size(), PropertyList);
if (BufferRef.isOutOfBounds(AccessOffset, AccessRange,
BufferRef.get_range()))
Expand All @@ -1798,10 +1857,12 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
"the buffer",
PI_ERROR_INVALID_VALUE);

initHostAcc();
detail::associateWithHandler(CommandGroupHandler, this, AccessTarget);
detail::constructorNotification(detail::getSyclObjImpl(BufferRef).get(),
detail::AccessorBaseHost::impl.get(),
AccessTarget, AccessMode, CodeLoc);
GDBMethodsAnchor();
}
#endif

Expand Down
2 changes: 2 additions & 0 deletions sycl/source/accessor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,8 @@ range<3> &AccessorBaseHost::getAccessRange() { return impl->MAccessRange; }
range<3> &AccessorBaseHost::getMemoryRange() { return impl->MMemoryRange; }
void *AccessorBaseHost::getPtr() { return impl->MData; }

detail::AccHostDataT &AccessorBaseHost::getAccData() { return impl->MAccData; }

const property_list &AccessorBaseHost::getPropList() const {
return impl->MPropertyList;
}
Expand Down
17 changes: 9 additions & 8 deletions sycl/source/detail/accessor_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,6 +9,7 @@
#pragma once

#include <sycl/access/access.hpp>
#include <sycl/accessor.hpp>
#include <sycl/detail/export.hpp>
#include <sycl/id.hpp>
#include <sycl/property_list.hpp>
Expand Down Expand Up @@ -45,17 +46,15 @@ class __SYCL_EXPORT AccessorImplHost {
int ElemSize, int OffsetInBytes = 0,
bool IsSubBuffer = false,
const property_list &PropertyList = {})
: MOffset(Offset), MAccessRange(AccessRange), MMemoryRange(MemoryRange),
MAccessMode(AccessMode),
: MAccData(Offset, AccessRange, MemoryRange), MAccessMode(AccessMode),
MSYCLMemObj((detail::SYCLMemObjI *)SYCLMemObject), MDims(Dims),
MElemSize(ElemSize), MOffsetInBytes(OffsetInBytes),
MIsSubBuffer(IsSubBuffer), MPropertyList(PropertyList) {}

~AccessorImplHost();

AccessorImplHost(const AccessorImplHost &Other)
: MOffset(Other.MOffset), MAccessRange(Other.MAccessRange),
MMemoryRange(Other.MMemoryRange), MAccessMode(Other.MAccessMode),
: MAccData(Other.MAccData), MAccessMode(Other.MAccessMode),
MSYCLMemObj(Other.MSYCLMemObj), MDims(Other.MDims),
MElemSize(Other.MElemSize), MOffsetInBytes(Other.MOffsetInBytes),
MIsSubBuffer(Other.MIsSubBuffer), MPropertyList(Other.MPropertyList) {}
Expand All @@ -71,11 +70,13 @@ class __SYCL_EXPORT AccessorImplHost {

void resize(size_t GlobalSize);

id<3> MOffset;
detail::AccHostDataT MAccData;

id<3> &MOffset = MAccData.MOffset;
// The size of accessing region.
range<3> MAccessRange;
range<3> &MAccessRange = MAccData.MAccessRange;
// The size of memory object this requirement is created for.
range<3> MMemoryRange;
range<3> &MMemoryRange = MAccData.MMemoryRange;
access::mode MAccessMode;

detail::SYCLMemObjI *MSYCLMemObj;
Expand All @@ -85,7 +86,7 @@ class __SYCL_EXPORT AccessorImplHost {
unsigned int MOffsetInBytes;
bool MIsSubBuffer;

void *MData = nullptr;
void *&MData = MAccData.MData;

Command *MBlockedCmd = nullptr;

Expand Down
6 changes: 4 additions & 2 deletions sycl/test/abi/layout_accessors_host.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,8 @@ void hostAcc(accessor<int, 1, access::mode::read, access::target::host_buffer> A
// CHECK-NEXT: 8 | class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
// CHECK-NEXT: 8 | _Sp_counted_base<(_Lock_policy)2U> * _M_pi
// CHECK-NEXT: 0 | class sycl::detail::accessor_common<int, 1, sycl::access::mode::read, sycl::access::target::host_buffer, sycl::access::placeholder::false_t> (base) (empty)
// CHECK-NEXT: 16 | char[16] padding
// CHECK-NEXT: 16 | detail::AccHostDataT * MAccData
// CHECK-NEXT: 24 | char[8] padding
// CHECK-NEXT: [sizeof=32, dsize=32, align=8,
// CHECK-NEXT: nvsize=32, nvalign=8]

Expand All @@ -46,7 +47,8 @@ void hostAcc(accessor<int, 1, access::mode::read, access::target::global_buffer>
// CHECK-NEXT: 8 | class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
// CHECK-NEXT: 8 | _Sp_counted_base<(_Lock_policy)2U> * _M_pi
// CHECK-NEXT: 0 | class sycl::detail::accessor_common<int, 1, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::false_t> (base) (empty)
// CHECK-NEXT: 16 | char[16] padding
// CHECK-NEXT: 16 | detail::AccHostDataT * MAccData
// CHECK-NEXT: 24 | char[8] padding
// CHECK-NEXT: [sizeof=32, dsize=32, align=8,
// CHECK-NEXT: nvsize=32, nvalign=8]

Expand Down
1 change: 1 addition & 0 deletions sycl/test/abi/sycl_symbols_linux.dump
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3812,6 +3812,7 @@ _ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEE
_ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEERKNS0_7contextE
_ZN4sycl3_V16detail14getBorderColorENS0_19image_channel_orderE
_ZN4sycl3_V16detail15getOrWaitEventsESt6vectorINS0_5eventESaIS3_EESt10shared_ptrINS1_12context_implEE
_ZN4sycl3_V16detail16AccessorBaseHost10getAccDataEv
_ZN4sycl3_V16detail16AccessorBaseHost14getAccessRangeEv
_ZN4sycl3_V16detail16AccessorBaseHost14getMemoryRangeEv
_ZN4sycl3_V16detail16AccessorBaseHost6getPtrEv
Expand Down
Loading