diff --git a/sycl/CMakeLists.txt b/sycl/CMakeLists.txt index 04830be9cbc82..23bc873d63c03 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 16) +set(SYCL_DEV_ABI_VERSION 17) if (SYCL_ADD_DEV_VERSION_POSTFIX) set(SYCL_VERSION_POSTFIX "-${SYCL_DEV_ABI_VERSION}") endif() diff --git a/sycl/include/sycl/detail/device_binary_image.hpp b/sycl/include/sycl/detail/device_binary_image.hpp deleted file mode 100644 index 61070124d33b9..0000000000000 --- a/sycl/include/sycl/detail/device_binary_image.hpp +++ /dev/null @@ -1,72 +0,0 @@ -//==----- device_binary_image.hpp --- SYCL device binary image abstraction -==// -// -// 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 - -namespace sycl { -__SYCL_INLINE_VER_NAMESPACE(_V1) { -namespace detail { - -// SYCL RT wrapper over PI binary image. -class RTDeviceBinaryImage : public pi::DeviceBinaryImage { -public: - RTDeviceBinaryImage(OSModuleHandle ModuleHandle) - : pi::DeviceBinaryImage(), ModuleHandle(ModuleHandle) {} - RTDeviceBinaryImage(pi_device_binary Bin, OSModuleHandle ModuleHandle) - : pi::DeviceBinaryImage(Bin), ModuleHandle(ModuleHandle) {} - // Explicitly delete copy constructor/operator= to avoid unintentional copies - RTDeviceBinaryImage(const RTDeviceBinaryImage &) = delete; - RTDeviceBinaryImage &operator=(const RTDeviceBinaryImage &) = delete; - // Explicitly retain move constructors to facilitate potential moves across - // collections - RTDeviceBinaryImage(RTDeviceBinaryImage &&) = default; - RTDeviceBinaryImage &operator=(RTDeviceBinaryImage &&) = default; - - OSModuleHandle getOSModuleHandle() const { return ModuleHandle; } - - ~RTDeviceBinaryImage() override {} - - bool supportsSpecConstants() const { - return getFormat() == PI_DEVICE_BINARY_TYPE_SPIRV; - } - - const pi_device_binary_struct &getRawData() const { return *get(); } - - void print() const override { - pi::DeviceBinaryImage::print(); - std::cerr << " OSModuleHandle=" << ModuleHandle << "\n"; - } - -protected: - OSModuleHandle ModuleHandle; -}; - -// Dynamically allocated device binary image, which de-allocates its binary -// data in destructor. -class DynRTDeviceBinaryImage : public RTDeviceBinaryImage { -public: - DynRTDeviceBinaryImage(std::unique_ptr &&DataPtr, size_t DataSize, - OSModuleHandle M); - ~DynRTDeviceBinaryImage() override; - - void print() const override { - RTDeviceBinaryImage::print(); - std::cerr << " DYNAMICALLY CREATED\n"; - } - -protected: - std::unique_ptr Data; -}; - -} // namespace detail -} // __SYCL_INLINE_VER_NAMESPACE(_V1) -} // namespace sycl diff --git a/sycl/include/sycl/detail/pi.hpp b/sycl/include/sycl/detail/pi.hpp index d666d6aacc8b6..4c05c5d52aad5 100644 --- a/sycl/include/sycl/detail/pi.hpp +++ b/sycl/include/sycl/detail/pi.hpp @@ -207,196 +207,6 @@ void emitFunctionWithArgsEndTrace(uint64_t CorrelationID, uint32_t FuncID, const char *FName, unsigned char *ArgsData, pi_result Result, pi_plugin Plugin); -// A wrapper for passing around byte array properties -class ByteArray { -public: - using ConstIterator = const std::uint8_t *; - - ByteArray(const std::uint8_t *Ptr, std::size_t Size) : Ptr{Ptr}, Size{Size} {} - const std::uint8_t &operator[](std::size_t Idx) const { return Ptr[Idx]; } - std::size_t size() const { return Size; } - ConstIterator begin() const { return Ptr; } - ConstIterator end() const { return Ptr + Size; } - -private: - const std::uint8_t *Ptr; - const std::size_t Size; -}; - -// C++ wrapper over the _pi_device_binary_property_struct structure. -class DeviceBinaryProperty { -public: - DeviceBinaryProperty(const _pi_device_binary_property_struct *Prop) - : Prop(Prop) {} - - pi_uint32 asUint32() const; - ByteArray asByteArray() const; - const char *asCString() const; - -protected: - friend std::ostream &operator<<(std::ostream &Out, - const DeviceBinaryProperty &P); - const _pi_device_binary_property_struct *Prop; -}; - -std::ostream &operator<<(std::ostream &Out, const DeviceBinaryProperty &P); - -// C++ convenience wrapper over the pi_device_binary_struct structure. -class DeviceBinaryImage { -public: - // Represents a range of properties to enable iteration over them. - // Implements the standard C++ STL input iterator interface. - class PropertyRange { - public: - using ValTy = std::remove_pointer::type; - - class ConstIterator { - pi_device_binary_property Cur; - - public: - using iterator_category = std::input_iterator_tag; - using value_type = ValTy; - using difference_type = ptrdiff_t; - using pointer = const pi_device_binary_property; - using reference = pi_device_binary_property; - - ConstIterator(pi_device_binary_property Cur = nullptr) : Cur(Cur) {} - ConstIterator &operator++() { - Cur++; - return *this; - } - ConstIterator operator++(int) { - ConstIterator Ret = *this; - ++(*this); - return Ret; - } - bool operator==(ConstIterator Other) const { return Cur == Other.Cur; } - bool operator!=(ConstIterator Other) const { return !(*this == Other); } - reference operator*() const { return Cur; } - }; - ConstIterator begin() const { return ConstIterator(Begin); } - ConstIterator end() const { return ConstIterator(End); } - friend class DeviceBinaryImage; - bool isAvailable() const { return !(Begin == nullptr); } - - private: - PropertyRange() : Begin(nullptr), End(nullptr) {} - // Searches for a property set with given name and constructs a - // PropertyRange spanning all its elements. If property set is not found, - // the range will span zero elements. - PropertyRange(pi_device_binary Bin, const char *PropSetName) - : PropertyRange() { - init(Bin, PropSetName); - }; - void init(pi_device_binary Bin, const char *PropSetName); - pi_device_binary_property Begin; - pi_device_binary_property End; - }; - -public: - DeviceBinaryImage(pi_device_binary Bin) { init(Bin); } - DeviceBinaryImage() : Bin(nullptr){}; - - virtual void print() const; - virtual void dump(std::ostream &Out) const; - - size_t getSize() const { - assert(Bin && "binary image data not set"); - return static_cast(Bin->BinaryEnd - Bin->BinaryStart); - } - - const char *getCompileOptions() const { - assert(Bin && "binary image data not set"); - return Bin->CompileOptions; - } - - const char *getLinkOptions() const { - assert(Bin && "binary image data not set"); - return Bin->LinkOptions; - } - - /// Returns the format of the binary image - pi::PiDeviceBinaryType getFormat() const { - assert(Bin && "binary image data not set"); - return Format; - } - - /// Returns a single property from SYCL_MISC_PROP category. - pi_device_binary_property getProperty(const char *PropName) const; - - /// Gets the iterator range over specialization constants in this binary - /// image. For each property pointed to by an iterator within the - /// range, the name of the property is the specialization constant symbolic ID - /// and the value is a list of 3-element tuples of 32-bit unsigned integers, - /// describing the specialization constant. - /// This is done in order to unify representation of both scalar and composite - /// specialization constants: composite specialization constant is represented - /// by its leaf elements, so for scalars the list contains only a single - /// tuple, while for composite there might be more of them. - /// Each tuple consists of ID of scalar specialization constant, its location - /// within a composite (offset in bytes from the beginning or 0 if it is not - /// an element of a composite specialization constant) and its size. - /// For example, for the following structure: - /// struct A { int a; float b; }; - /// struct POD { A a[2]; int b; }; - /// List of tuples will look like: - /// { ID0, 0, 4 }, // .a[0].a - /// { ID1, 4, 4 }, // .a[0].b - /// { ID2, 8, 4 }, // .a[1].a - /// { ID3, 12, 4 }, // .a[1].b - /// { ID4, 16, 4 }, // .b - /// And for an interger specialization constant, the list of tuples will look - /// like: - /// { ID5, 0, 4 } - const PropertyRange &getSpecConstants() const { return SpecConstIDMap; } - const PropertyRange getSpecConstantsDefaultValues() const { - // We can't have this variable as a class member, since it would break - // the ABI backwards compatibility. - DeviceBinaryImage::PropertyRange SpecConstDefaultValuesMap; - SpecConstDefaultValuesMap.init( - Bin, __SYCL_PI_PROPERTY_SET_SPEC_CONST_DEFAULT_VALUES_MAP); - return SpecConstDefaultValuesMap; - } - const PropertyRange &getDeviceLibReqMask() const { return DeviceLibReqMask; } - const PropertyRange &getKernelParamOptInfo() const { - return KernelParamOptInfo; - } - const PropertyRange getAssertUsed() const { - // We can't have this variable as a class member, since it would break - // the ABI backwards compatibility. - PropertyRange AssertUsed; - AssertUsed.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_ASSERT_USED); - return AssertUsed; - } - const PropertyRange &getProgramMetadata() const { return ProgramMetadata; } - const PropertyRange getExportedSymbols() const { - // We can't have this variable as a class member, since it would break - // the ABI backwards compatibility. - DeviceBinaryImage::PropertyRange ExportedSymbols; - ExportedSymbols.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_EXPORTED_SYMBOLS); - return ExportedSymbols; - } - const PropertyRange getDeviceGlobals() const { - // We can't have this variable as a class member, since it would break - // the ABI backwards compatibility. - DeviceBinaryImage::PropertyRange DeviceGlobals; - DeviceGlobals.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_GLOBALS); - return DeviceGlobals; - } - virtual ~DeviceBinaryImage() {} - -protected: - void init(pi_device_binary Bin); - pi_device_binary get() const { return Bin; } - - pi_device_binary Bin; - pi::PiDeviceBinaryType Format = PI_DEVICE_BINARY_TYPE_NONE; - DeviceBinaryImage::PropertyRange SpecConstIDMap; - DeviceBinaryImage::PropertyRange DeviceLibReqMask; - DeviceBinaryImage::PropertyRange KernelParamOptInfo; - DeviceBinaryImage::PropertyRange ProgramMetadata; -}; - /// Tries to determine the device binary image foramat. Returns /// PI_DEVICE_BINARY_TYPE_NONE if unsuccessful. PiDeviceBinaryType getBinaryImageFormat(const unsigned char *ImgData, diff --git a/sycl/source/detail/device_binary_image.cpp b/sycl/source/detail/device_binary_image.cpp index 281c64cfd0644..708e5ec44b8ec 100644 --- a/sycl/source/detail/device_binary_image.cpp +++ b/sycl/source/detail/device_binary_image.cpp @@ -6,16 +6,177 @@ // //===----------------------------------------------------------------------===// +#include #include +#include +#include #include -#include - namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { namespace detail { +std::ostream &operator<<(std::ostream &Out, const DeviceBinaryProperty &P) { + switch (P.Prop->Type) { + case PI_PROPERTY_TYPE_UINT32: + Out << "[UINT32] "; + break; + case PI_PROPERTY_TYPE_BYTE_ARRAY: + Out << "[Byte array] "; + break; + case PI_PROPERTY_TYPE_STRING: + Out << "[String] "; + break; + default: + assert(false && "unsupported property"); + return Out; + } + Out << P.Prop->Name << "="; + + switch (P.Prop->Type) { + case PI_PROPERTY_TYPE_UINT32: + Out << P.asUint32(); + break; + case PI_PROPERTY_TYPE_BYTE_ARRAY: { + ByteArray BA = P.asByteArray(); + std::ios_base::fmtflags FlagsBackup = Out.flags(); + Out << std::hex; + for (const auto &Byte : BA) { + Out << "0x" << static_cast(Byte) << " "; + } + Out.flags(FlagsBackup); + break; + } + case PI_PROPERTY_TYPE_STRING: + Out << P.asCString(); + break; + default: + assert(false && "Unsupported property"); + return Out; + } + return Out; +} + +pi_uint32 DeviceBinaryProperty::asUint32() const { + assert(Prop->Type == PI_PROPERTY_TYPE_UINT32 && "property type mismatch"); + // if type fits into the ValSize - it is used to store the property value + assert(Prop->ValAddr == nullptr && "primitive types must be stored inline"); + const auto *P = reinterpret_cast(&Prop->ValSize); + return (*P) | (*(P + 1) << 8) | (*(P + 2) << 16) | (*(P + 3) << 24); +} + +ByteArray DeviceBinaryProperty::asByteArray() const { + assert(Prop->Type == PI_PROPERTY_TYPE_BYTE_ARRAY && "property type mismatch"); + assert(Prop->ValSize > 0 && "property size mismatch"); + const auto *Data = pi::cast(Prop->ValAddr); + return {Data, Prop->ValSize}; +} + +const char *DeviceBinaryProperty::asCString() const { + assert(Prop->Type == PI_PROPERTY_TYPE_STRING && "property type mismatch"); + assert(Prop->ValSize > 0 && "property size mismatch"); + return pi::cast(Prop->ValAddr); +} + +void RTDeviceBinaryImage::PropertyRange::init(pi_device_binary Bin, + const char *PropSetName) { + assert(!this->Begin && !this->End && "already initialized"); + pi_device_binary_property_set PS = nullptr; + + for (PS = Bin->PropertySetsBegin; PS != Bin->PropertySetsEnd; ++PS) { + assert(PS->Name && "nameless property set - bug in the offload wrapper?"); + if (!strcmp(PropSetName, PS->Name)) + break; + } + if (PS == Bin->PropertySetsEnd) { + Begin = End = nullptr; + return; + } + Begin = PS->PropertiesBegin; + End = Begin ? PS->PropertiesEnd : nullptr; +} + +void RTDeviceBinaryImage::print() const { + std::cerr << " --- Image " << Bin << "\n"; + if (!Bin) + return; + std::cerr << " Version : " << (int)Bin->Version << "\n"; + std::cerr << " Kind : " << (int)Bin->Kind << "\n"; + std::cerr << " Format : " << (int)Bin->Format << "\n"; + std::cerr << " Target : " << Bin->DeviceTargetSpec << "\n"; + std::cerr << " Bin size : " + << ((intptr_t)Bin->BinaryEnd - (intptr_t)Bin->BinaryStart) << "\n"; + std::cerr << " OSModuleHandle : " << ModuleHandle << "\n"; + std::cerr << " Compile options : " + << (Bin->CompileOptions ? Bin->CompileOptions : "NULL") << "\n"; + std::cerr << " Link options : " + << (Bin->LinkOptions ? Bin->LinkOptions : "NULL") << "\n"; + std::cerr << " Entries : "; + for (_pi_offload_entry EntriesIt = Bin->EntriesBegin; + EntriesIt != Bin->EntriesEnd; ++EntriesIt) + std::cerr << EntriesIt->name << " "; + std::cerr << "\n"; + std::cerr << " Properties [" << Bin->PropertySetsBegin << "-" + << Bin->PropertySetsEnd << "]:\n"; + + for (pi_device_binary_property_set PS = Bin->PropertySetsBegin; + PS != Bin->PropertySetsEnd; ++PS) { + std::cerr << " Category " << PS->Name << " [" << PS->PropertiesBegin + << "-" << PS->PropertiesEnd << "]:\n"; + + for (pi_device_binary_property P = PS->PropertiesBegin; + P != PS->PropertiesEnd; ++P) { + std::cerr << " " << DeviceBinaryProperty(P) << "\n"; + } + } +} + +void RTDeviceBinaryImage::dump(std::ostream &Out) const { + size_t ImgSize = getSize(); + Out.write(reinterpret_cast(Bin->BinaryStart), ImgSize); +} + +pi_device_binary_property +RTDeviceBinaryImage::getProperty(const char *PropName) const { + RTDeviceBinaryImage::PropertyRange BoolProp; + BoolProp.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_MISC_PROP); + if (!BoolProp.isAvailable()) + return nullptr; + auto It = std::find_if(BoolProp.begin(), BoolProp.end(), + [=](pi_device_binary_property Prop) { + return !strcmp(PropName, Prop->Name); + }); + if (It == BoolProp.end()) + return nullptr; + + return *It; +} + +void RTDeviceBinaryImage::init(pi_device_binary Bin) { + this->Bin = Bin; + // If device binary image format wasn't set by its producer, then can't change + // now, because 'Bin' data is part of the executable image loaded into memory + // which can't be modified (easily). + // TODO clang driver + ClangOffloadWrapper can figure out the format and set + // it when invoking the offload wrapper job + Format = static_cast(Bin->Format); + + if (Format == PI_DEVICE_BINARY_TYPE_NONE) + // try to determine the format; may remain "NONE" + Format = pi::getBinaryImageFormat(Bin->BinaryStart, getSize()); + + SpecConstIDMap.init(Bin, __SYCL_PI_PROPERTY_SET_SPEC_CONST_MAP); + SpecConstDefaultValuesMap.init( + Bin, __SYCL_PI_PROPERTY_SET_SPEC_CONST_DEFAULT_VALUES_MAP); + DeviceLibReqMask.init(Bin, __SYCL_PI_PROPERTY_SET_DEVICELIB_REQ_MASK); + KernelParamOptInfo.init(Bin, __SYCL_PI_PROPERTY_SET_KERNEL_PARAM_OPT_INFO); + AssertUsed.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_ASSERT_USED); + ProgramMetadata.init(Bin, __SYCL_PI_PROPERTY_SET_PROGRAM_METADATA); + ExportedSymbols.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_EXPORTED_SYMBOLS); + DeviceGlobals.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_GLOBALS); +} + DynRTDeviceBinaryImage::DynRTDeviceBinaryImage( std::unique_ptr &&DataPtr, size_t DataSize, OSModuleHandle M) : RTDeviceBinaryImage(M) { diff --git a/sycl/source/detail/device_binary_image.hpp b/sycl/source/detail/device_binary_image.hpp new file mode 100644 index 0000000000000..8d0251b5bac36 --- /dev/null +++ b/sycl/source/detail/device_binary_image.hpp @@ -0,0 +1,231 @@ +//==----- device_binary_image.hpp --- SYCL device binary image abstraction -==// +// +// 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 + +namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { +namespace detail { + +// A wrapper for passing around byte array properties +class ByteArray { +public: + using ConstIterator = const std::uint8_t *; + + ByteArray(const std::uint8_t *Ptr, std::size_t Size) : Ptr{Ptr}, Size{Size} {} + const std::uint8_t &operator[](std::size_t Idx) const { return Ptr[Idx]; } + std::size_t size() const { return Size; } + ConstIterator begin() const { return Ptr; } + ConstIterator end() const { return Ptr + Size; } + +private: + const std::uint8_t *Ptr; + const std::size_t Size; +}; + +// C++ wrapper over the _pi_device_binary_property_struct structure. +class DeviceBinaryProperty { +public: + DeviceBinaryProperty(const _pi_device_binary_property_struct *Prop) + : Prop(Prop) {} + + pi_uint32 asUint32() const; + ByteArray asByteArray() const; + const char *asCString() const; + +protected: + friend std::ostream &operator<<(std::ostream &Out, + const DeviceBinaryProperty &P); + const _pi_device_binary_property_struct *Prop; +}; + +std::ostream &operator<<(std::ostream &Out, const DeviceBinaryProperty &P); + +// SYCL RT wrapper over PI binary image. +class RTDeviceBinaryImage { +public: + // Represents a range of properties to enable iteration over them. + // Implements the standard C++ STL input iterator interface. + class PropertyRange { + public: + using ValTy = std::remove_pointer::type; + + class ConstIterator { + pi_device_binary_property Cur; + + public: + using iterator_category = std::input_iterator_tag; + using value_type = ValTy; + using difference_type = ptrdiff_t; + using pointer = const pi_device_binary_property; + using reference = pi_device_binary_property; + + ConstIterator(pi_device_binary_property Cur = nullptr) : Cur(Cur) {} + ConstIterator &operator++() { + Cur++; + return *this; + } + ConstIterator operator++(int) { + ConstIterator Ret = *this; + ++(*this); + return Ret; + } + bool operator==(ConstIterator Other) const { return Cur == Other.Cur; } + bool operator!=(ConstIterator Other) const { return !(*this == Other); } + reference operator*() const { return Cur; } + }; + ConstIterator begin() const { return ConstIterator(Begin); } + ConstIterator end() const { return ConstIterator(End); } + friend class RTDeviceBinaryImage; + bool isAvailable() const { return !(Begin == nullptr); } + + private: + PropertyRange() : Begin(nullptr), End(nullptr) {} + // Searches for a property set with given name and constructs a + // PropertyRange spanning all its elements. If property set is not found, + // the range will span zero elements. + PropertyRange(pi_device_binary Bin, const char *PropSetName) + : PropertyRange() { + init(Bin, PropSetName); + }; + void init(pi_device_binary Bin, const char *PropSetName); + pi_device_binary_property Begin; + pi_device_binary_property End; + }; + +public: + RTDeviceBinaryImage(OSModuleHandle ModuleHandle) + : Bin(nullptr), ModuleHandle(ModuleHandle) {} + RTDeviceBinaryImage(pi_device_binary Bin, OSModuleHandle ModuleHandle) + : ModuleHandle(ModuleHandle) { + init(Bin); + } + // Explicitly delete copy constructor/operator= to avoid unintentional copies + RTDeviceBinaryImage(const RTDeviceBinaryImage &) = delete; + RTDeviceBinaryImage &operator=(const RTDeviceBinaryImage &) = delete; + // Explicitly retain move constructors to facilitate potential moves across + // collections + RTDeviceBinaryImage(RTDeviceBinaryImage &&) = default; + RTDeviceBinaryImage &operator=(RTDeviceBinaryImage &&) = default; + + OSModuleHandle getOSModuleHandle() const { return ModuleHandle; } + + virtual ~RTDeviceBinaryImage() {} + + bool supportsSpecConstants() const { + return getFormat() == PI_DEVICE_BINARY_TYPE_SPIRV; + } + + const pi_device_binary_struct &getRawData() const { return *get(); } + + virtual void print() const; + virtual void dump(std::ostream &Out) const; + + size_t getSize() const { + assert(Bin && "binary image data not set"); + return static_cast(Bin->BinaryEnd - Bin->BinaryStart); + } + + const char *getCompileOptions() const { + assert(Bin && "binary image data not set"); + return Bin->CompileOptions; + } + + const char *getLinkOptions() const { + assert(Bin && "binary image data not set"); + return Bin->LinkOptions; + } + + /// Returns the format of the binary image + pi::PiDeviceBinaryType getFormat() const { + assert(Bin && "binary image data not set"); + return Format; + } + + /// Returns a single property from SYCL_MISC_PROP category. + pi_device_binary_property getProperty(const char *PropName) const; + + /// Gets the iterator range over specialization constants in this binary + /// image. For each property pointed to by an iterator within the + /// range, the name of the property is the specialization constant symbolic ID + /// and the value is a list of 3-element tuples of 32-bit unsigned integers, + /// describing the specialization constant. + /// This is done in order to unify representation of both scalar and composite + /// specialization constants: composite specialization constant is represented + /// by its leaf elements, so for scalars the list contains only a single + /// tuple, while for composite there might be more of them. + /// Each tuple consists of ID of scalar specialization constant, its location + /// within a composite (offset in bytes from the beginning or 0 if it is not + /// an element of a composite specialization constant) and its size. + /// For example, for the following structure: + /// struct A { int a; float b; }; + /// struct POD { A a[2]; int b; }; + /// List of tuples will look like: + /// { ID0, 0, 4 }, // .a[0].a + /// { ID1, 4, 4 }, // .a[0].b + /// { ID2, 8, 4 }, // .a[1].a + /// { ID3, 12, 4 }, // .a[1].b + /// { ID4, 16, 4 }, // .b + /// And for an interger specialization constant, the list of tuples will look + /// like: + /// { ID5, 0, 4 } + const PropertyRange &getSpecConstants() const { return SpecConstIDMap; } + const PropertyRange &getSpecConstantsDefaultValues() const { + return SpecConstDefaultValuesMap; + } + const PropertyRange &getDeviceLibReqMask() const { return DeviceLibReqMask; } + const PropertyRange &getKernelParamOptInfo() const { + return KernelParamOptInfo; + } + const PropertyRange &getAssertUsed() const { return AssertUsed; } + const PropertyRange &getProgramMetadata() const { return ProgramMetadata; } + const PropertyRange &getExportedSymbols() const { return ExportedSymbols; } + const PropertyRange &getDeviceGlobals() const { return DeviceGlobals; } + +protected: + void init(pi_device_binary Bin); + pi_device_binary get() const { return Bin; } + + pi_device_binary Bin; + OSModuleHandle ModuleHandle; + + pi::PiDeviceBinaryType Format = PI_DEVICE_BINARY_TYPE_NONE; + RTDeviceBinaryImage::PropertyRange SpecConstIDMap; + RTDeviceBinaryImage::PropertyRange SpecConstDefaultValuesMap; + RTDeviceBinaryImage::PropertyRange DeviceLibReqMask; + RTDeviceBinaryImage::PropertyRange KernelParamOptInfo; + RTDeviceBinaryImage::PropertyRange AssertUsed; + RTDeviceBinaryImage::PropertyRange ProgramMetadata; + RTDeviceBinaryImage::PropertyRange ExportedSymbols; + RTDeviceBinaryImage::PropertyRange DeviceGlobals; +}; + +// Dynamically allocated device binary image, which de-allocates its binary +// data in destructor. +class DynRTDeviceBinaryImage : public RTDeviceBinaryImage { +public: + DynRTDeviceBinaryImage(std::unique_ptr &&DataPtr, size_t DataSize, + OSModuleHandle M); + ~DynRTDeviceBinaryImage() override; + + void print() const override { + RTDeviceBinaryImage::print(); + std::cerr << " DYNAMICALLY CREATED\n"; + } + +protected: + std::unique_ptr Data; +}; + +} // namespace detail +} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace sycl diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index 9f438cf169243..f155106dd6e97 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -245,12 +245,12 @@ class device_image_impl { private: void updateSpecConstSymMap() { if (MBinImage) { - const pi::DeviceBinaryImage::PropertyRange &SCRange = + const RTDeviceBinaryImage::PropertyRange &SCRange = MBinImage->getSpecConstants(); - using SCItTy = pi::DeviceBinaryImage::PropertyRange::ConstIterator; + using SCItTy = RTDeviceBinaryImage::PropertyRange::ConstIterator; // get default values for specialization constants - const pi::DeviceBinaryImage::PropertyRange &SCDefValRange = + const RTDeviceBinaryImage::PropertyRange &SCDefValRange = MBinImage->getSpecConstantsDefaultValues(); // This variable is used to calculate spec constant value offset in a @@ -259,8 +259,7 @@ class device_image_impl { for (SCItTy SCIt : SCRange) { const char *SCName = (*SCIt)->Name; - pi::ByteArray Descriptors = - pi::DeviceBinaryProperty(*SCIt).asByteArray(); + ByteArray Descriptors = DeviceBinaryProperty(*SCIt).asByteArray(); assert(Descriptors.size() > 8 && "Unexpected property size"); // Expected layout is vector of 3-component tuples (flattened into a @@ -301,8 +300,8 @@ class device_image_impl { bool HasDefaultValues = SCDefValRange.begin() != SCDefValRange.end(); if (HasDefaultValues) { - pi::ByteArray DefValDescriptors = - pi::DeviceBinaryProperty(*SCDefValRange.begin()).asByteArray(); + ByteArray DefValDescriptors = + DeviceBinaryProperty(*SCDefValRange.begin()).asByteArray(); assert(DefValDescriptors.size() - 8 == MSpecConstsBlob.size() && "Specialization constant default value blob do not have the " "expected size."); diff --git a/sycl/source/detail/persistent_device_code_cache.hpp b/sycl/source/detail/persistent_device_code_cache.hpp index e5aadeebabf62..b5ebc17a3638c 100644 --- a/sycl/source/detail/persistent_device_code_cache.hpp +++ b/sycl/source/detail/persistent_device_code_cache.hpp @@ -9,9 +9,9 @@ #pragma once #include +#include #include #include -#include #include #include #include diff --git a/sycl/source/detail/pi.cpp b/sycl/source/detail/pi.cpp index c5a6d7c0d3f24..f81b20bb9d92d 100644 --- a/sycl/source/detail/pi.cpp +++ b/sycl/source/detail/pi.cpp @@ -549,146 +549,6 @@ void assertion(bool Condition, const char *Message) { die(Message); } -std::ostream &operator<<(std::ostream &Out, const DeviceBinaryProperty &P) { - switch (P.Prop->Type) { - case PI_PROPERTY_TYPE_UINT32: - Out << "[UINT32] "; - break; - case PI_PROPERTY_TYPE_BYTE_ARRAY: - Out << "[Byte array] "; - break; - case PI_PROPERTY_TYPE_STRING: - Out << "[String] "; - break; - default: - assert(false && "unsupported property"); - return Out; - } - Out << P.Prop->Name << "="; - - switch (P.Prop->Type) { - case PI_PROPERTY_TYPE_UINT32: - Out << P.asUint32(); - break; - case PI_PROPERTY_TYPE_BYTE_ARRAY: { - ByteArray BA = P.asByteArray(); - std::ios_base::fmtflags FlagsBackup = Out.flags(); - Out << std::hex; - for (const auto &Byte : BA) { - Out << "0x" << static_cast(Byte) << " "; - } - Out.flags(FlagsBackup); - break; - } - case PI_PROPERTY_TYPE_STRING: - Out << P.asCString(); - break; - default: - assert(false && "Unsupported property"); - return Out; - } - return Out; -} - -void DeviceBinaryImage::print() const { - std::cerr << " --- Image " << Bin << "\n"; - if (!Bin) - return; - std::cerr << " Version : " << (int)Bin->Version << "\n"; - std::cerr << " Kind : " << (int)Bin->Kind << "\n"; - std::cerr << " Format : " << (int)Bin->Format << "\n"; - std::cerr << " Target : " << Bin->DeviceTargetSpec << "\n"; - std::cerr << " Bin size : " - << ((intptr_t)Bin->BinaryEnd - (intptr_t)Bin->BinaryStart) << "\n"; - std::cerr << " Compile options : " - << (Bin->CompileOptions ? Bin->CompileOptions : "NULL") << "\n"; - std::cerr << " Link options : " - << (Bin->LinkOptions ? Bin->LinkOptions : "NULL") << "\n"; - std::cerr << " Entries : "; - for (_pi_offload_entry EntriesIt = Bin->EntriesBegin; - EntriesIt != Bin->EntriesEnd; ++EntriesIt) - std::cerr << EntriesIt->name << " "; - std::cerr << "\n"; - std::cerr << " Properties [" << Bin->PropertySetsBegin << "-" - << Bin->PropertySetsEnd << "]:\n"; - - for (pi_device_binary_property_set PS = Bin->PropertySetsBegin; - PS != Bin->PropertySetsEnd; ++PS) { - std::cerr << " Category " << PS->Name << " [" << PS->PropertiesBegin - << "-" << PS->PropertiesEnd << "]:\n"; - - for (pi_device_binary_property P = PS->PropertiesBegin; - P != PS->PropertiesEnd; ++P) { - std::cerr << " " << DeviceBinaryProperty(P) << "\n"; - } - } -} - -void DeviceBinaryImage::dump(std::ostream &Out) const { - size_t ImgSize = getSize(); - Out.write(reinterpret_cast(Bin->BinaryStart), ImgSize); -} - -static pi_uint32 asUint32(const void *Addr) { - assert(Addr && "Addr is NULL"); - const auto *P = reinterpret_cast(Addr); - return (*P) | (*(P + 1) << 8) | (*(P + 2) << 16) | (*(P + 3) << 24); -} - -pi_uint32 DeviceBinaryProperty::asUint32() const { - assert(Prop->Type == PI_PROPERTY_TYPE_UINT32 && "property type mismatch"); - // if type fits into the ValSize - it is used to store the property value - assert(Prop->ValAddr == nullptr && "primitive types must be stored inline"); - return sycl::detail::pi::asUint32(&Prop->ValSize); -} - -ByteArray DeviceBinaryProperty::asByteArray() const { - assert(Prop->Type == PI_PROPERTY_TYPE_BYTE_ARRAY && "property type mismatch"); - assert(Prop->ValSize > 0 && "property size mismatch"); - const auto *Data = pi::cast(Prop->ValAddr); - return {Data, Prop->ValSize}; -} - -const char *DeviceBinaryProperty::asCString() const { - assert(Prop->Type == PI_PROPERTY_TYPE_STRING && "property type mismatch"); - assert(Prop->ValSize > 0 && "property size mismatch"); - return pi::cast(Prop->ValAddr); -} - -void DeviceBinaryImage::PropertyRange::init(pi_device_binary Bin, - const char *PropSetName) { - assert(!this->Begin && !this->End && "already initialized"); - pi_device_binary_property_set PS = nullptr; - - for (PS = Bin->PropertySetsBegin; PS != Bin->PropertySetsEnd; ++PS) { - assert(PS->Name && "nameless property set - bug in the offload wrapper?"); - if (!strcmp(PropSetName, PS->Name)) - break; - } - if (PS == Bin->PropertySetsEnd) { - Begin = End = nullptr; - return; - } - Begin = PS->PropertiesBegin; - End = Begin ? PS->PropertiesEnd : nullptr; -} - -pi_device_binary_property -DeviceBinaryImage::getProperty(const char *PropName) const { - DeviceBinaryImage::PropertyRange BoolProp; - BoolProp.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_MISC_PROP); - if (!BoolProp.isAvailable()) - return nullptr; - auto It = std::find_if(BoolProp.begin(), BoolProp.end(), - [=](pi_device_binary_property Prop) { - return !strcmp(PropName, Prop->Name); - }); - if (It == BoolProp.end()) - return nullptr; - - return *It; -} - // Reads an integer value from ELF data. template static ResT readELFValue(const unsigned char *Data, size_t NumBytes, @@ -823,25 +683,6 @@ RT::PiDeviceBinaryType getBinaryImageFormat(const unsigned char *ImgData, return PI_DEVICE_BINARY_TYPE_NONE; } -void DeviceBinaryImage::init(pi_device_binary Bin) { - this->Bin = Bin; - // If device binary image format wasn't set by its producer, then can't change - // now, because 'Bin' data is part of the executable image loaded into memory - // which can't be modified (easily). - // TODO clang driver + ClangOffloadWrapper can figure out the format and set - // it when invoking the offload wrapper job - Format = static_cast(Bin->Format); - - if (Format == PI_DEVICE_BINARY_TYPE_NONE) - // try to determine the format; may remain "NONE" - Format = getBinaryImageFormat(Bin->BinaryStart, getSize()); - - SpecConstIDMap.init(Bin, __SYCL_PI_PROPERTY_SET_SPEC_CONST_MAP); - DeviceLibReqMask.init(Bin, __SYCL_PI_PROPERTY_SET_DEVICELIB_REQ_MASK); - KernelParamOptInfo.init(Bin, __SYCL_PI_PROPERTY_SET_KERNEL_PARAM_OPT_INFO); - ProgramMetadata.init(Bin, __SYCL_PI_PROPERTY_SET_PROGRAM_METADATA); -} - } // namespace pi } // namespace detail } // __SYCL_INLINE_VER_NAMESPACE(_V1) diff --git a/sycl/source/detail/program_impl.cpp b/sycl/source/detail/program_impl.cpp index 505efbacf9ef8..a0458c5fcc1bb 100644 --- a/sycl/source/detail/program_impl.cpp +++ b/sycl/source/detail/program_impl.cpp @@ -526,9 +526,9 @@ void program_impl::flush_spec_constants(const RTDeviceBinaryImage &Img, RT::PiProgram NativePrg) const { // iterate via all specialization constants the program's image depends on, // and set each to current runtime value (if any) - const pi::DeviceBinaryImage::PropertyRange &SCRange = Img.getSpecConstants(); + const RTDeviceBinaryImage::PropertyRange &SCRange = Img.getSpecConstants(); ContextImplPtr Ctx = getSyclObjImpl(get_context()); - using SCItTy = pi::DeviceBinaryImage::PropertyRange::ConstIterator; + using SCItTy = RTDeviceBinaryImage::PropertyRange::ConstIterator; auto LockGuard = Ctx->getKernelProgramCache().acquireCachedPrograms(); NativePrg = NativePrg ? NativePrg : getHandleRef(); @@ -540,7 +540,7 @@ void program_impl::flush_spec_constants(const RTDeviceBinaryImage &Img, continue; const spec_constant_impl &SC = SCEntry->second; assert(SC.isSet() && "uninitialized spec constant"); - pi::ByteArray Descriptors = pi::DeviceBinaryProperty(*SCIt).asByteArray(); + ByteArray Descriptors = DeviceBinaryProperty(*SCIt).asByteArray(); // First 8 bytes are consumed by size of the property assert(Descriptors.size() > 8 && "Unexpected property size"); // Expected layout is vector of 3-component tuples (flattened into a vector diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index f1b08b27cad79..7ad1c1b649a63 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -383,7 +383,7 @@ static void appendLinkOptionsFromImage(std::string &LinkOpts, static bool getUint32PropAsBool(const RTDeviceBinaryImage &Img, const char *PropName) { pi_device_binary_property Prop = Img.getProperty(PropName); - return Prop && (pi::DeviceBinaryProperty(Prop).asUint32() != 0); + return Prop && (DeviceBinaryProperty(Prop).asUint32() != 0); } static void appendCompileOptionsFromImage(std::string &CompileOpts, @@ -1084,7 +1084,7 @@ ProgramManager::build(ProgramPtr Program, const ContextImplPtr Context, } static ProgramManager::KernelArgMask -createKernelArgMask(const pi::ByteArray &Bytes) { +createKernelArgMask(const ByteArray &Bytes) { const int NBytesForSize = 8; const int NBitsInElement = 8; std::uint64_t SizeInBits = 0; @@ -1102,7 +1102,7 @@ createKernelArgMask(const pi::ByteArray &Bytes) { void ProgramManager::cacheKernelUsesAssertInfo(OSModuleHandle M, RTDeviceBinaryImage &Img) { - const pi::DeviceBinaryImage::PropertyRange &AssertUsedRange = + const RTDeviceBinaryImage::PropertyRange &AssertUsedRange = Img.getAssertUsed(); if (AssertUsedRange.isAvailable()) for (const auto &Prop : AssertUsedRange) { @@ -1129,14 +1129,14 @@ void ProgramManager::addImages(pi_device_binaries DeviceBinary) { auto Img = make_unique_ptr(RawImg, M); // Fill the kernel argument mask map - const pi::DeviceBinaryImage::PropertyRange &KPOIRange = + const RTDeviceBinaryImage::PropertyRange &KPOIRange = Img->getKernelParamOptInfo(); if (KPOIRange.isAvailable()) { KernelNameToArgMaskMap &ArgMaskMap = m_EliminatedKernelArgMasks[Img.get()]; for (const auto &Info : KPOIRange) ArgMaskMap[Info->Name] = - createKernelArgMask(pi::DeviceBinaryProperty(Info).asByteArray()); + createKernelArgMask(DeviceBinaryProperty(Info).asByteArray()); } // Fill maps for kernel bundles @@ -1226,8 +1226,8 @@ void ProgramManager::addImages(pi_device_binaries DeviceBinary) { auto DeviceGlobals = Img->getDeviceGlobals(); for (const pi_device_binary_property &DeviceGlobal : DeviceGlobals) { - pi::ByteArray DeviceGlobalInfo = - pi::DeviceBinaryProperty(DeviceGlobal).asByteArray(); + ByteArray DeviceGlobalInfo = + DeviceBinaryProperty(DeviceGlobal).asByteArray(); // The supplied device_global info property is expected to contain: // * 8 bytes - Size of the property. @@ -1391,10 +1391,10 @@ void ProgramManager::flushSpecConstants(const program_impl &Prg, // mask, sycl runtime won't know which fallback device libraries are needed. In // such case, the safest way is to load all fallback device libraries. uint32_t ProgramManager::getDeviceLibReqMask(const RTDeviceBinaryImage &Img) { - const pi::DeviceBinaryImage::PropertyRange &DLMRange = + const RTDeviceBinaryImage::PropertyRange &DLMRange = Img.getDeviceLibReqMask(); if (DLMRange.isAvailable()) - return pi::DeviceBinaryProperty(*(DLMRange.begin())).asUint32(); + return DeviceBinaryProperty(*(DLMRange.begin())).asUint32(); else return 0xFFFFFFFF; } diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp index 25efbc1525d0a..0a66a576abed7 100644 --- a/sycl/source/detail/program_manager/program_manager.hpp +++ b/sycl/source/detail/program_manager/program_manager.hpp @@ -7,10 +7,10 @@ //===----------------------------------------------------------------------===// #pragma once +#include #include #include #include -#include #include #include #include diff --git a/sycl/test/abi/vtable.cpp b/sycl/test/abi/vtable.cpp index 008a510496db7..11a4189ebaa00 100644 --- a/sycl/test/abi/vtable.cpp +++ b/sycl/test/abi/vtable.cpp @@ -23,16 +23,6 @@ void foo(sycl::detail::HostKernelBase &HKB) { // CHECK-NEXT: 4 | sycl::detail::HostKernelBase::~HostKernelBase() [complete] // CHECK-NEXT: 5 | sycl::detail::HostKernelBase::~HostKernelBase() [deleting] -void foo(sycl::detail::pi::DeviceBinaryImage &Img) { Img.print(); } -// CHECK: Vtable for 'sycl::detail::pi::DeviceBinaryImage' (6 entries). -// CHECK-NEXT: 0 | offset_to_top (0) -// CHECK-NEXT: 1 | sycl::detail::pi::DeviceBinaryImage RTTI -// CHECK-NEXT: -- (sycl::detail::pi::DeviceBinaryImage, 0) vtable address -- -// CHECK-NEXT: 2 | void sycl::detail::pi::DeviceBinaryImage::print() const -// CHECK-NEXT: 3 | void sycl::detail::pi::DeviceBinaryImage::dump(std::ostream &) const -// CHECK-NEXT: 4 | sycl::detail::pi::DeviceBinaryImage::~DeviceBinaryImage() [complete] -// CHECK-NEXT: 5 | sycl::detail::pi::DeviceBinaryImage::~DeviceBinaryImage() [deleting] - void foo(sycl::detail::CG *CG) { delete CG; } // CHECK: Vtable for 'sycl::detail::CG' (4 entries). // CHECK-NEXT: 0 | offset_to_top (0) diff --git a/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp b/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp index bb36fccf7f422..bedbb3d4e85f2 100644 --- a/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp +++ b/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp @@ -10,11 +10,11 @@ #include "../thread_safety/ThreadUtils.h" #include "detail/persistent_device_code_cache.hpp" #include +#include #include #include #include #include -#include #include #include #include