From 2e4bfc92c0c019f2d086056f719f7665f7f6a4d3 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Thu, 8 Sep 2022 07:12:11 -0700 Subject: [PATCH 1/4] [SYCL][ABI-Break] Implement property interface for local_accessor & usm_allocator --- sycl/include/sycl/accessor.hpp | 24 ++++++++++++++++++---- sycl/include/sycl/detail/accessor_impl.hpp | 13 ++++++++---- sycl/include/sycl/usm/usm_allocator.hpp | 8 ++++++++ sycl/test/abi/layout_accessors_host.cpp | 18 ++++++++++++++-- sycl/test/abi/symbol_size_alignment.cpp | 2 +- 5 files changed, 54 insertions(+), 11 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index e6cf64c25a0d8..b485295abcf3b 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -2167,8 +2167,8 @@ class __SYCL_SPECIAL_CLASS local_accessor_base : (void)propList; } #else - : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT)) { - (void)propList; + : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT), + propList) { detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(), access::target::local, AccessMode, CodeLoc); } @@ -2200,8 +2200,7 @@ class __SYCL_SPECIAL_CLASS local_accessor_base : } #else : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize), - AdjustedDim, sizeof(DataT)) { - (void)propList; + AdjustedDim, sizeof(DataT), propList) { detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(), access::target::local, AccessMode, CodeLoc); } @@ -2345,6 +2344,23 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor } #endif + +public: + template bool has_property() const noexcept { +#ifndef __SYCL_DEVICE_ONLY__ + return this->getPropList().template has_property(); +#else + return false; +#endif + } + + template Property get_property() const { +#ifndef __SYCL_DEVICE_ONLY__ + return this->getPropList().template get_property(); +#else + return Property(); +#endif + } }; /// Image accessors. diff --git a/sycl/include/sycl/detail/accessor_impl.hpp b/sycl/include/sycl/detail/accessor_impl.hpp index a820b9dc478de..a248e1035de88 100644 --- a/sycl/include/sycl/detail/accessor_impl.hpp +++ b/sycl/include/sycl/detail/accessor_impl.hpp @@ -180,23 +180,27 @@ class __SYCL_EXPORT LocalAccessorImplHost { public: // Allocate ElemSize more data to have sufficient padding to enforce // alignment. - LocalAccessorImplHost(sycl::range<3> Size, int Dims, int ElemSize) + LocalAccessorImplHost(sycl::range<3> Size, int Dims, int ElemSize, + const property_list &PropertyList) : MSize(Size), MDims(Dims), MElemSize(ElemSize), - MMem(Size[0] * Size[1] * Size[2] * ElemSize + ElemSize) {} + MMem(Size[0] * Size[1] * Size[2] * ElemSize + ElemSize), + MPropertyList(PropertyList) {} sycl::range<3> MSize; int MDims; int MElemSize; std::vector MMem; + property_list MPropertyList; }; using LocalAccessorImplPtr = std::shared_ptr; class LocalAccessorBaseHost { public: - LocalAccessorBaseHost(sycl::range<3> Size, int Dims, int ElemSize) { + LocalAccessorBaseHost(sycl::range<3> Size, int Dims, int ElemSize, + const property_list &PropertyList = {}) { impl = std::shared_ptr( - new LocalAccessorImplHost(Size, Dims, ElemSize)); + new LocalAccessorImplHost(Size, Dims, ElemSize, PropertyList)); } sycl::range<3> &getSize() { return impl->MSize; } const sycl::range<3> &getSize() const { return impl->MSize; } @@ -218,6 +222,7 @@ class LocalAccessorBaseHost { int getNumOfDims() { return impl->MDims; } int getElementSize() { return impl->MElemSize; } + const property_list &getPropList() const { return impl->MPropertyList; } protected: template diff --git a/sycl/include/sycl/usm/usm_allocator.hpp b/sycl/include/sycl/usm/usm_allocator.hpp index 3f56b42e63fb0..233c92ab62f01 100644 --- a/sycl/include/sycl/usm/usm_allocator.hpp +++ b/sycl/include/sycl/usm/usm_allocator.hpp @@ -106,6 +106,14 @@ class usm_allocator { (One.MDevice == Two.MDevice)); } + template bool has_property() const noexcept { + return MPropList.has_property(); + } + + template Property get_property() const { + return MPropList.get_property(); + } + private: constexpr size_t getAlignment() const { return max(alignof(T), Alignment); } diff --git a/sycl/test/abi/layout_accessors_host.cpp b/sycl/test/abi/layout_accessors_host.cpp index 09f2a088bbb98..a463d96603d78 100644 --- a/sycl/test/abi/layout_accessors_host.cpp +++ b/sycl/test/abi/layout_accessors_host.cpp @@ -58,8 +58,22 @@ using namespace sycl; // CHECK: 32 | pointer _M_start // CHECK-NEXT: 40 | pointer _M_finish // CHECK-NEXT: 48 | pointer _M_end_of_storage -// CHECK-NEXT: | [sizeof=56, dsize=56, align=8, -// CHECK-NEXT: | nvsize=56, nvalign=8] +// CHECK-NEXT: 56 | class sycl::property_list MPropertyList +// CHECK-NEXT: 56 | class sycl::detail::PropertyListBase (base) +// CHECK-NEXT: 56 | class std::bitset<32> MDataLessProps +// CHECK-NEXT: 56 | struct std::_Base_bitset<1> (base) +// CHECK-NEXT: 56 | _WordT _M_w +// CHECK-NEXT: 64 | class std::vector > MPropsWithData +// CHECK-NEXT: 64 | struct std::_Vector_base, class std::allocator > > (base) +// CHECK-NEXT: 64 | struct std::_Vector_base, class std::allocator > >::_Vector_impl _M_impl +// CHECK-NEXT: 64 | class std::allocator > (base) (empty) +// CHECK-NEXT: 64 | class __gnu_cxx::new_allocator > (base) (empty) +// CHECK-NEXT: 64 | struct std::_Vector_base, class std::allocator > >::_Vector_impl_data (base) +// CHECK-NEXT: 64 | pointer _M_start +// CHECK-NEXT: 72 | pointer _M_finish +// CHECK-NEXT: 80 | pointer _M_end_of_storage +// CHECK-NEXT: | [sizeof=88, dsize=88, align=8, +// CHECK-NEXT: | nvsize=88, nvalign=8] //----------------------------------------------------------------------------// // Host buffer accessor. diff --git a/sycl/test/abi/symbol_size_alignment.cpp b/sycl/test/abi/symbol_size_alignment.cpp index c681df3434c90..6dfd2182d297b 100644 --- a/sycl/test/abi/symbol_size_alignment.cpp +++ b/sycl/test/abi/symbol_size_alignment.cpp @@ -45,7 +45,7 @@ int main() { check, 24, 8>(); check(); check(); - check(); + check(); check, 40, 8>(); check(); check(); From be7f99f594d57753e418db0261a5082ba46346f7 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Thu, 8 Sep 2022 09:49:37 -0700 Subject: [PATCH 2/4] Update SYCL_DEV_ABI_VERSION --- sycl/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/CMakeLists.txt b/sycl/CMakeLists.txt index 984bd649f1b9d..ae368fbce300e 100644 --- a/sycl/CMakeLists.txt +++ b/sycl/CMakeLists.txt @@ -30,7 +30,7 @@ set(SYCL_MINOR_VERSION 7) set(SYCL_PATCH_VERSION 0) # Don't forget to re-enable sycl_symbols_windows.dump once we leave ABI-breaking # window! -set(SYCL_DEV_ABI_VERSION 15) +set(SYCL_DEV_ABI_VERSION 16) if (SYCL_ADD_DEV_VERSION_POSTFIX) set(SYCL_VERSION_POSTFIX "-${SYCL_DEV_ABI_VERSION}") endif() From b5067a1e348687a8367e924561adbf932d502b38 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Wed, 14 Sep 2022 04:24:39 -0700 Subject: [PATCH 3/4] Update llvm/sycl/test/abi/sycl_symbols_linux.dump --- sycl/test/abi/sycl_symbols_linux.dump | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index ca6b049b964bf..eb5e391b40418 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3843,8 +3843,8 @@ _ZN4sycl3_V16detail21LocalAccessorBaseHost12getNumOfDimsEv _ZN4sycl3_V16detail21LocalAccessorBaseHost14getElementSizeEv _ZN4sycl3_V16detail21LocalAccessorBaseHost6getPtrEv _ZN4sycl3_V16detail21LocalAccessorBaseHost7getSizeEv -_ZN4sycl3_V16detail21LocalAccessorBaseHostC1ENS0_5rangeILi3EEEii -_ZN4sycl3_V16detail21LocalAccessorBaseHostC2ENS0_5rangeILi3EEEii +_ZN4sycl3_V16detail21LocalAccessorBaseHostC1ENS0_5rangeILi3EEEiiRKNS0_13property_listE +_ZN4sycl3_V16detail21LocalAccessorBaseHostC2ENS0_5rangeILi3EEEiiRKNS0_13property_listE _ZN4sycl3_V16detail22addHostAccessorAndWaitEPNS1_16AccessorImplHostE _ZN4sycl3_V16detail22getImageNumberChannelsENS0_19image_channel_orderE _ZN4sycl3_V16detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EENS0_12bundle_stateE @@ -4136,6 +4136,7 @@ _ZNK4sycl3_V16detail19kernel_bundle_plain33contains_specialization_constantsEv _ZNK4sycl3_V16detail19kernel_bundle_plain3endEv _ZNK4sycl3_V16detail19kernel_bundle_plain5beginEv _ZNK4sycl3_V16detail19kernel_bundle_plain5emptyEv +_ZNK4sycl3_V16detail21LocalAccessorBaseHost11getPropListEv _ZNK4sycl3_V16detail21LocalAccessorBaseHost6getPtrEv _ZNK4sycl3_V16detail21LocalAccessorBaseHost7getSizeEv _ZNK4sycl3_V16device11get_backendEv From 12f9b462b452f94dcf5f7174ede76eb19c30ad11 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Wed, 14 Sep 2022 04:35:51 -0700 Subject: [PATCH 4/4] Fix merge error & clang-format --- sycl/include/sycl/detail/accessor_impl.hpp | 291 --------------------- sycl/source/accessor.cpp | 6 +- 2 files changed, 3 insertions(+), 294 deletions(-) delete mode 100644 sycl/include/sycl/detail/accessor_impl.hpp diff --git a/sycl/include/sycl/detail/accessor_impl.hpp b/sycl/include/sycl/detail/accessor_impl.hpp deleted file mode 100644 index a248e1035de88..0000000000000 --- a/sycl/include/sycl/detail/accessor_impl.hpp +++ /dev/null @@ -1,291 +0,0 @@ -//==------------ accessor_impl.hpp - SYCL standard header file -------------==// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#pragma once - -#include -#include -#include -#include -#include -#include - -namespace sycl { -__SYCL_INLINE_VER_NAMESPACE(_V1) { -template -class accessor; - -namespace ext { -namespace intel { -namespace esimd { -namespace detail { -// Forward declare a "back-door" access class to support ESIMD. -class AccessorPrivateProxy; -} // namespace detail -} // namespace esimd -} // namespace intel -} // namespace ext - -namespace detail { - -class SYCLMemObjI; - -class Command; - -// The class describes a requirement to access a SYCL memory object such as -// sycl::buffer and sycl::image. For example, each accessor used in a kernel, -// except one with access target "local", adds such requirement for the command -// group. - -template class AccessorImplDevice { -public: - AccessorImplDevice() = default; - AccessorImplDevice(id Offset, range AccessRange, - range MemoryRange) - : Offset(Offset), AccessRange(AccessRange), MemRange(MemoryRange) {} - - id Offset; - range AccessRange; - range MemRange; - - bool operator==(const AccessorImplDevice &Rhs) const { - return (Offset == Rhs.Offset && AccessRange == Rhs.AccessRange && - MemRange == Rhs.MemRange); - } -}; - -template class LocalAccessorBaseDevice { -public: - LocalAccessorBaseDevice(sycl::range Size) - : AccessRange(Size), - MemRange(InitializedVal::template get<0>()) {} - // TODO: Actually we need only one field here, but currently compiler requires - // all of them. - range AccessRange; - range MemRange; - id Offset; - - bool operator==(const LocalAccessorBaseDevice &Rhs) const { - return (AccessRange == Rhs.AccessRange); - } -}; - -class __SYCL_EXPORT AccessorImplHost { -public: - AccessorImplHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange, - access::mode AccessMode, void *SYCLMemObject, int Dims, - int ElemSize, int OffsetInBytes = 0, - bool IsSubBuffer = false, - const property_list &PropertyList = {}) - : MOffset(Offset), MAccessRange(AccessRange), MMemoryRange(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), - MSYCLMemObj(Other.MSYCLMemObj), MDims(Other.MDims), - MElemSize(Other.MElemSize), MOffsetInBytes(Other.MOffsetInBytes), - MIsSubBuffer(Other.MIsSubBuffer) {} - - // The resize method provides a way to change the size of the - // allocated memory and corresponding properties for the accessor. - // These are normally fixed for the accessor, but this capability - // is needed to support the stream class. - // Stream implementation creates an accessor with initial size for - // work item. But the number of work items is not available during - // stream construction. The resize method allows to update the accessor - // as the information becomes available to the handler. - - void resize(size_t GlobalSize); - - id<3> MOffset; - // The size of accessing region. - range<3> MAccessRange; - // The size of memory object this requirement is created for. - range<3> MMemoryRange; - access::mode MAccessMode; - - detail::SYCLMemObjI *MSYCLMemObj; - - unsigned int MDims; - unsigned int MElemSize; - unsigned int MOffsetInBytes; - bool MIsSubBuffer; - - void *MData = nullptr; - - Command *MBlockedCmd = nullptr; - - bool PerWI = false; - - // To preserve runtime properties - property_list MPropertyList; -}; - -using AccessorImplPtr = std::shared_ptr; - -class AccessorBaseHost { -public: - template - AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange, - access::mode AccessMode, void *SYCLMemObject, int Dims, - int ElemSize, int OffsetInBytes = 0, - bool IsSubBuffer = false, - const PropertyListT &PropertyList = {}) { - impl = std::shared_ptr(new AccessorImplHost( - Offset, AccessRange, MemoryRange, AccessMode, - (detail::SYCLMemObjI *)SYCLMemObject, Dims, ElemSize, OffsetInBytes, - IsSubBuffer, PropertyList)); - } - -protected: - id<3> &getOffset() { return impl->MOffset; } - range<3> &getAccessRange() { return impl->MAccessRange; } - range<3> &getMemoryRange() { return impl->MMemoryRange; } - void *getPtr() { return impl->MData; } - unsigned int getElemSize() const { return impl->MElemSize; } - - const id<3> &getOffset() const { return impl->MOffset; } - const range<3> &getAccessRange() const { return impl->MAccessRange; } - const range<3> &getMemoryRange() const { return impl->MMemoryRange; } - void *getPtr() const { return const_cast(impl->MData); } - - template - friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject); - - template - friend class accessor; - - AccessorImplPtr impl; - -private: - friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy; -}; - -class __SYCL_EXPORT LocalAccessorImplHost { -public: - // Allocate ElemSize more data to have sufficient padding to enforce - // alignment. - LocalAccessorImplHost(sycl::range<3> Size, int Dims, int ElemSize, - const property_list &PropertyList) - : MSize(Size), MDims(Dims), MElemSize(ElemSize), - MMem(Size[0] * Size[1] * Size[2] * ElemSize + ElemSize), - MPropertyList(PropertyList) {} - - sycl::range<3> MSize; - int MDims; - int MElemSize; - std::vector MMem; - property_list MPropertyList; -}; - -using LocalAccessorImplPtr = std::shared_ptr; - -class LocalAccessorBaseHost { -public: - LocalAccessorBaseHost(sycl::range<3> Size, int Dims, int ElemSize, - const property_list &PropertyList = {}) { - impl = std::shared_ptr( - new LocalAccessorImplHost(Size, Dims, ElemSize, PropertyList)); - } - sycl::range<3> &getSize() { return impl->MSize; } - const sycl::range<3> &getSize() const { return impl->MSize; } - void *getPtr() { - // Const cast this in order to call the const getPtr. - return const_cast(this)->getPtr(); - } - void *getPtr() const { - char *ptr = impl->MMem.data(); - - // Align the pointer to MElemSize. - size_t val = reinterpret_cast(ptr); - if (val % impl->MElemSize != 0) { - ptr += impl->MElemSize - val % impl->MElemSize; - } - - return ptr; - } - - int getNumOfDims() { return impl->MDims; } - int getElementSize() { return impl->MElemSize; } - const property_list &getPropList() const { return impl->MPropertyList; } - -protected: - template - friend decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject); - - std::shared_ptr impl; -}; - -using Requirement = AccessorImplHost; - -void __SYCL_EXPORT addHostAccessorAndWait(Requirement *Req); - -#if __cplusplus >= 201703L - -template -constexpr access::mode deduceAccessMode() { - // property_list = {} is not properly detected by deduction guide, - // when parameter is passed without curly braces: access(buffer, no_init) - // thus simplest approach is to check 2 last arguments for being a tag - if constexpr (std::is_same>::value || - std::is_same>::value) { - return access::mode::read; - } - - if constexpr (std::is_same>::value || - std::is_same>::value) { - return access::mode::write; - } - - if constexpr ( - std::is_same>::value || - std::is_same>::value) { - return access::mode::read; - } - - return access::mode::read_write; -} - -template -constexpr access::target deduceAccessTarget(access::target defaultTarget) { - if constexpr ( - std::is_same>::value || - std::is_same>::value) { - return access::target::constant_buffer; - } - - return defaultTarget; -} - -#endif - -} // namespace detail -} // __SYCL_INLINE_VER_NAMESPACE(_V1) -} // namespace sycl diff --git a/sycl/source/accessor.cpp b/sycl/source/accessor.cpp index bfb9fd347d7f8..e9d6c6be901f3 100644 --- a/sycl/source/accessor.cpp +++ b/sycl/source/accessor.cpp @@ -52,9 +52,9 @@ void *AccessorBaseHost::getPtr() const { void *AccessorBaseHost::getMemoryObject() const { return impl->MSYCLMemObj; } -LocalAccessorBaseHost::LocalAccessorBaseHost(sycl::range<3> Size, int Dims, - int ElemSize, - const property_list &PropertyList) { +LocalAccessorBaseHost::LocalAccessorBaseHost( + sycl::range<3> Size, int Dims, int ElemSize, + const property_list &PropertyList) { impl = std::shared_ptr( new LocalAccessorImplHost(Size, Dims, ElemSize, PropertyList)); }