From ec93747fd6aabf9663a9cf7cd253c98dd1febd60 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Thu, 1 Apr 2021 14:17:34 +0300 Subject: [PATCH 01/21] [SYCL] Add support for set(get)_specialization_constant --- sycl/include/CL/sycl/detail/kernel_desc.hpp | 5 ++ sycl/include/CL/sycl/kernel_bundle.hpp | 44 ++++++++---- sycl/source/detail/device_image_impl.hpp | 80 +++++++++++++++------ sycl/source/detail/kernel_bundle_impl.hpp | 9 +-- sycl/source/kernel_bundle.cpp | 19 +++++ 5 files changed, 118 insertions(+), 39 deletions(-) diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp index 7fbe0d1dffe76..0204ec057ee34 100644 --- a/sycl/include/CL/sycl/detail/kernel_desc.hpp +++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp @@ -49,6 +49,11 @@ template struct SpecConstantInfo { static constexpr const char *getName() { return ""; } }; +// Translates SYCL 2020 specialization constant type to its name. +template const char *get_spec_constant_symbolic_ID() { + return __builtin_unique_stable_name(SpecName); +} + #ifndef __SYCL_UNNAMED_LAMBDA__ template struct KernelInfo { static constexpr unsigned getNumParams() { return 0; } diff --git a/sycl/include/CL/sycl/kernel_bundle.hpp b/sycl/include/CL/sycl/kernel_bundle.hpp index 41ab9c3405899..dd421e4e53ee7 100644 --- a/sycl/include/CL/sycl/kernel_bundle.hpp +++ b/sycl/include/CL/sycl/kernel_bundle.hpp @@ -168,6 +168,14 @@ class __SYCL_EXPORT kernel_bundle_plain { // \returns an iterator to the last device image kernel_bundle contains const device_image_plain *end() const; + bool has_specialization_constant_impl(const char *SpecName) const noexcept; + + void set_specialization_constant_impl(const char *SpecName, void *Value, + size_t Size); + + void get_specialization_constant_impl(const char *SpecName, void *Value, + size_t Size) const; + detail::KernelBundleImplPtr impl; }; @@ -247,9 +255,8 @@ class kernel_bundle : public detail::kernel_bundle_plain { /// \returns true if any device image in the kernel_bundle uses specialization /// constant whose address is SpecName template bool has_specialization_constant() const noexcept { - throw sycl::runtime_error( - "kernel_bundle::has_specialization_constant is not implemented yet", - PI_INVALID_OPERATION); + const char *SpecSymName = detail::get_spec_constant_symbolic_ID(); + return has_specialization_constant(SpecSymName); } /// Sets the value of the specialization constant whose address is SpecName @@ -259,20 +266,33 @@ class kernel_bundle : public detail::kernel_bundle_plain { typename = detail::enable_if_t<_State == bundle_state::input>> void set_specialization_constant( typename std::remove_reference_t::value_type Value) { - (void)Value; - throw sycl::runtime_error( - "kernel_bundle::set_specialization_constant is not implemented yet", - PI_INVALID_OPERATION); + + using SCType = std::remove_reference_t::value_type; + static_assert(std::is_trivially_copyable_v); + // TODO can this be simply default constructible + static_assert(std::is_trivially_default_constructible_v); + + const char *SpecSymName = detail::get_spec_constant_symbolic_ID(); + set_specialization_constant_impl(SpecSymName, &Value, sizeof(SCType)); } - /// The value of the specialization constant whose address is SpecName for - /// this kernel bundle. + /// \returns the value of the specialization constant whose address is + /// SpecName for this kernel bundle. template typename std::remove_reference_t::value_type get_specialization_constant() const { - throw sycl::runtime_error( - "kernel_bundle::get_specialization_constant is not implemented yet", - PI_INVALID_OPERATION); + + using SCType = std::remove_reference_t::value_type; + static_assert(std::is_trivially_copyable_v); + // TODO can this be simply default constructible + static_assert(std::is_trivially_default_constructible_v); + + SCType RetValue; + + const char *SpecSymName = detail::get_spec_constant_symbolic_ID(); + get_specialization_constant_impl(SpecSymName, &RetValue, sizeof(SCType)); + + return RetValue; } #endif diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index dbfc545bcea3d..dbfbba3f95db7 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -39,7 +39,9 @@ class device_image_impl { std::vector KernelIDs, RT::PiProgram Program) : MBinImage(BinImage), MContext(std::move(Context)), MDevices(std::move(Devices)), MState(State), MProgram(Program), - MKernelIDs(std::move(KernelIDs)) {} + MKernelIDs(std::move(KernelIDs)) { + updateSpecConstSymMap(); + } bool has_kernel(const kernel_id &KernelIDCand) const noexcept { return std::binary_search(MKernelIDs.begin(), MKernelIDs.end(), @@ -76,16 +78,21 @@ class device_image_impl { bool IsSet = false; }; - bool has_specialization_constant(unsigned int SpecID) const noexcept { + bool has_specialization_constant(const char *SpecName) const noexcept { + if (SpecID.count(SpecID) == 0) + return false; + + unsigned SpecID = MSpecConstSymMap.at(SpecName); return std::any_of(MSpecConstDescs.begin(), MSpecConstDescs.end(), [SpecID](const SpecConstDescT &SpecConstDesc) { return SpecConstDesc.ID == SpecID; }); } - void set_specialization_constant_raw_value(unsigned int SpecID, + void set_specialization_constant_raw_value(const char *SpecName, const void *Value, size_t ValueSize) noexcept { + unsigned SpecID = MSpecConstSymMap[SpecName]; for (const SpecConstDescT &SpecConstDesc : MSpecConstDescs) if (SpecConstDesc.ID == SpecID) { // Lock the mutex to prevent when one thread in the middle of writing a @@ -98,9 +105,10 @@ class device_image_impl { } } - void get_specialization_constant_raw_value(unsigned int SpecID, + void get_specialization_constant_raw_value(const char *SpecName, void *ValueRet, size_t ValueSize) const noexcept { + unsigned SpecID = MSpecConstSymMap[SpecName]; for (const SpecConstDescT &SpecConstDesc : MSpecConstDescs) if (SpecConstDesc.ID == SpecID) { // Lock the mutex to prevent when one thread in the middle of writing a @@ -150,25 +158,51 @@ class device_image_impl { } private: - const RTDeviceBinaryImage *MBinImage = nullptr; - context MContext; - std::vector MDevices; - bundle_state MState; - // Native program handler which this device image represents - RT::PiProgram MProgram = nullptr; - // List of kernel ids available in this image, elements should be sorted - // according to LessByNameComp - std::vector MKernelIDs; - - // A mutex for sycnhronizing access to spec constants blob. Mutable because - // needs to be locked in the const method for getting spec constant value. - mutable std::mutex MSpecConstAccessMtx; - // Binary blob which can have values of all specialization constants in the - // image - std::vector MSpecConstsBlob; - // Contains list of spec ID + their offsets in the MSpecConstsBlob - std::vector MSpecConstDescs; -}; + void updateSpecConstSymMap() { + if (MBinImage) { + const pi_device_binary_struct &RawImg = MBinImage->getRawData(); + const auto &PropSetsBegin = RawImg->PropertySetsBegin; + const auto &PropSetsEnd = RawImg->PropertySetsEnd; + const auto &SpecConstMap = std::find_if( + PropSetsBegin, PropSetsEnd, + [](const _pi_device_binary_property_set_struct &Set) { + return strcmp(Set->Name, __SYCL_PI_PROPERTY_SET_SPEC_CONST_MAP) == + 0; + }); + + if (SpecConstMap != PropSetsEnd) { + const auto &PropsBegin = SpecConstMap->PropertiesBegin; + const auto &PropsEnd = SpecConstMap->PropertiesEnd; + + std::for_each(PropsBegin, PropsEnd, + [&](const _pi_device_binary_property_struct &Prop) { + MSpecConstSymMap[Prop.Name] = + *static_cast(Prop.ValAddr); + }); + } + } + + const RTDeviceBinaryImage *MBinImage = nullptr; + context MContext; + std::vector MDevices; + bundle_state MState; + // Native program handler which this device image represents + RT::PiProgram MProgram = nullptr; + // List of kernel ids available in this image, elements should be sorted + // according to LessByNameComp + std::vector MKernelIDs; + + // A mutex for sycnhronizing access to spec constants blob. Mutable because + // needs to be locked in the const method for getting spec constant value. + mutable std::mutex MSpecConstAccessMtx; + // Binary blob which can have values of all specialization constants in the + // image + std::vector MSpecConstsBlob; + // Contains list of spec ID + their offsets in the MSpecConstsBlob + std::vector MSpecConstDescs; + + std::map MSpecConstSymMap; + }; } // namespace detail } // namespace sycl diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index a7a5a93990725..c920a47e84e5d 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -254,9 +254,9 @@ class kernel_bundle_impl { ->set_specialization_constant_raw_value(SpecID, Value, ValueSize); } - const void *get_specialization_constant_raw_value(unsigned int SpecID, - void *ValueRet, - size_t ValueSize) const { + void get_specialization_constant_raw_value(unsigned int SpecID, + void *ValueRet, + size_t ValueSize) const { for (const device_image_plain &DeviceImage : MDeviceImages) if (getSyclObjImpl(DeviceImage)->has_specialization_constant(SpecID)) { getSyclObjImpl(DeviceImage) @@ -264,7 +264,8 @@ class kernel_bundle_impl { ValueSize); } - return nullptr; + throw sycl::runtime_error("Specialization constant not found", + PI_INVALID_VALUE); } const device_image_plain *begin() const { return &MDeviceImages.front(); } diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 7f35903792bf7..733ab23e4f5f2 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -83,6 +83,25 @@ bool kernel_bundle_plain::has_kernel(const kernel_id &KernelID, return impl->has_kernel(KernelID, Dev); } +bool kernel_bundle_plain::has_specialization_constant_impl( + const char *SpecName) const noexcept { + return impl->has_specialization_constant(SpecName); +} + +void kernel_bundle_plain::set_specialization_constant_impl(const char *SpecName, + void *Value, + size_t Size) { + unsigned SpecID = 0; + impl->set_specialization_constant_raw_value(SpecName, Value, Size); +} + +void kernel_bundle_plain::get_specialization_constant_impl(const char *SpecID, + void *Value, + size_t Size) const { + unsigned SpecID = 0; + impl->get_specialization_constant_raw_value(SpecName, Value, Size); +} + //////////////////////////// ///// free functions /////////////////////////// From 3495c0e38003aed70bf807cbf5ed8a35bf77a68c Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Fri, 2 Apr 2021 14:18:27 +0300 Subject: [PATCH 02/21] Address reviewer feedback --- sycl/include/CL/sycl/detail/kernel_desc.hpp | 6 +++++- sycl/source/detail/device_image_impl.hpp | 12 ++++++------ sycl/source/detail/kernel_bundle_impl.hpp | 16 ++++++++-------- sycl/source/kernel_bundle.cpp | 4 +--- 4 files changed, 20 insertions(+), 18 deletions(-) diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp index 0204ec057ee34..052b59cec82ce 100644 --- a/sycl/include/CL/sycl/detail/kernel_desc.hpp +++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp @@ -12,6 +12,7 @@ #include #include +#include __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { @@ -49,10 +50,13 @@ template struct SpecConstantInfo { static constexpr const char *getName() { return ""; } }; +#if __cplusplus >= 201703L // Translates SYCL 2020 specialization constant type to its name. template const char *get_spec_constant_symbolic_ID() { - return __builtin_unique_stable_name(SpecName); + return __builtin_unique_stable_name( + specialization_id_name_generator); } +#endif #ifndef __SYCL_UNNAMED_LAMBDA__ template struct KernelInfo { diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index dbfbba3f95db7..beef8f7488fb2 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -79,7 +79,7 @@ class device_image_impl { }; bool has_specialization_constant(const char *SpecName) const noexcept { - if (SpecID.count(SpecID) == 0) + if (MSpecConstSymMap.count(SpecName) == 0) return false; unsigned SpecID = MSpecConstSymMap.at(SpecName); @@ -108,7 +108,7 @@ class device_image_impl { void get_specialization_constant_raw_value(const char *SpecName, void *ValueRet, size_t ValueSize) const noexcept { - unsigned SpecID = MSpecConstSymMap[SpecName]; + unsigned SpecID = MSpecConstSymMap.at(SpecName); for (const SpecConstDescT &SpecConstDesc : MSpecConstDescs) if (SpecConstDesc.ID == SpecID) { // Lock the mutex to prevent when one thread in the middle of writing a @@ -161,13 +161,12 @@ class device_image_impl { void updateSpecConstSymMap() { if (MBinImage) { const pi_device_binary_struct &RawImg = MBinImage->getRawData(); - const auto &PropSetsBegin = RawImg->PropertySetsBegin; - const auto &PropSetsEnd = RawImg->PropertySetsEnd; + const auto &PropSetsBegin = RawImg.PropertySetsBegin; + const auto &PropSetsEnd = RawImg.PropertySetsEnd; const auto &SpecConstMap = std::find_if( PropSetsBegin, PropSetsEnd, [](const _pi_device_binary_property_set_struct &Set) { - return strcmp(Set->Name, __SYCL_PI_PROPERTY_SET_SPEC_CONST_MAP) == - 0; + return strcmp(Set.Name, __SYCL_PI_PROPERTY_SET_SPEC_CONST_MAP) == 0; }); if (SpecConstMap != PropSetsEnd) { @@ -181,6 +180,7 @@ class device_image_impl { }); } } + } const RTDeviceBinaryImage *MBinImage = nullptr; context MContext; diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index c920a47e84e5d..2cea9e9ff6733 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -238,29 +238,29 @@ class kernel_bundle_impl { }); } - bool has_specialization_constant(unsigned int SpecID) const noexcept { + bool has_specialization_constant(const char *SpecName) const noexcept { return std::any_of(MDeviceImages.begin(), MDeviceImages.end(), - [SpecID](const device_image_plain &DeviceImage) { + [SpecName](const device_image_plain &DeviceImage) { return getSyclObjImpl(DeviceImage) - ->has_specialization_constant(SpecID); + ->has_specialization_constant(SpecName); }); } - void set_specialization_constant_raw_value(unsigned int SpecID, + void set_specialization_constant_raw_value(const char *SpecName, const void *Value, size_t ValueSize) { for (const device_image_plain &DeviceImage : MDeviceImages) getSyclObjImpl(DeviceImage) - ->set_specialization_constant_raw_value(SpecID, Value, ValueSize); + ->set_specialization_constant_raw_value(SpecName, Value, ValueSize); } - void get_specialization_constant_raw_value(unsigned int SpecID, + void get_specialization_constant_raw_value(const char *SpecName, void *ValueRet, size_t ValueSize) const { for (const device_image_plain &DeviceImage : MDeviceImages) - if (getSyclObjImpl(DeviceImage)->has_specialization_constant(SpecID)) { + if (getSyclObjImpl(DeviceImage)->has_specialization_constant(SpecName)) { getSyclObjImpl(DeviceImage) - ->get_specialization_constant_raw_value(SpecID, ValueRet, + ->get_specialization_constant_raw_value(SpecName, ValueRet, ValueSize); } diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 733ab23e4f5f2..abe3400cb64d3 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -91,14 +91,12 @@ bool kernel_bundle_plain::has_specialization_constant_impl( void kernel_bundle_plain::set_specialization_constant_impl(const char *SpecName, void *Value, size_t Size) { - unsigned SpecID = 0; impl->set_specialization_constant_raw_value(SpecName, Value, Size); } -void kernel_bundle_plain::get_specialization_constant_impl(const char *SpecID, +void kernel_bundle_plain::get_specialization_constant_impl(const char *SpecName, void *Value, size_t Size) const { - unsigned SpecID = 0; impl->get_specialization_constant_raw_value(SpecName, Value, Size); } From 0f22c62aef75920633752402477c8a167036e4fc Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Mon, 5 Apr 2021 14:21:13 +0300 Subject: [PATCH 03/21] More fixes --- sycl/source/detail/device_image_impl.hpp | 8 +++++++- sycl/source/detail/kernel_bundle_impl.hpp | 1 + 2 files changed, 8 insertions(+), 1 deletion(-) diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index beef8f7488fb2..f9c982bb1c9a1 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -93,8 +93,9 @@ class device_image_impl { const void *Value, size_t ValueSize) noexcept { unsigned SpecID = MSpecConstSymMap[SpecName]; - for (const SpecConstDescT &SpecConstDesc : MSpecConstDescs) + for (SpecConstDescT &SpecConstDesc : MSpecConstDescs) if (SpecConstDesc.ID == SpecID) { + SpecConstDesc.IsSet = true; // Lock the mutex to prevent when one thread in the middle of writing a // new value while another thread is reading the value to pass it to // JIT compiler. @@ -177,6 +178,11 @@ class device_image_impl { [&](const _pi_device_binary_property_struct &Prop) { MSpecConstSymMap[Prop.Name] = *static_cast(Prop.ValAddr); + MSpecConstDescs.emplace_back( + *static_cast(Prop.ValAddr), // ID + *(static_cast(Prop.ValAddr) + 1), // Offset + false + ); }); } } diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 2cea9e9ff6733..bc3a85bcb445c 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -262,6 +262,7 @@ class kernel_bundle_impl { getSyclObjImpl(DeviceImage) ->get_specialization_constant_raw_value(SpecName, ValueRet, ValueSize); + return; } throw sycl::runtime_error("Specialization constant not found", From f7672b7ba036c615bf69374612e41c604744dbb0 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Mon, 5 Apr 2021 15:16:24 +0300 Subject: [PATCH 04/21] Fix minor issues --- sycl/include/CL/sycl/handler.hpp | 2 +- sycl/include/CL/sycl/kernel_bundle.hpp | 6 +-- sycl/include/CL/sycl/kernel_handler.hpp | 3 ++ sycl/source/detail/device_image_impl.hpp | 4 +- .../specialization_constants/host_apis.cpp | 41 +++++++++++++++++++ 5 files changed, 50 insertions(+), 6 deletions(-) create mode 100644 sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 2523e917890e1..dc1a6a2d072dc 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -1048,7 +1048,7 @@ class __SYCL_EXPORT handler { } template - typename std::remove_reference_t::type + typename std::remove_reference_t::value_type get_specialization_constant() const { std::shared_ptr KernelBundleImplPtr = diff --git a/sycl/include/CL/sycl/kernel_bundle.hpp b/sycl/include/CL/sycl/kernel_bundle.hpp index dd421e4e53ee7..7168c5f46105a 100644 --- a/sycl/include/CL/sycl/kernel_bundle.hpp +++ b/sycl/include/CL/sycl/kernel_bundle.hpp @@ -256,7 +256,7 @@ class kernel_bundle : public detail::kernel_bundle_plain { /// constant whose address is SpecName template bool has_specialization_constant() const noexcept { const char *SpecSymName = detail::get_spec_constant_symbolic_ID(); - return has_specialization_constant(SpecSymName); + return has_specialization_constant_impl(SpecSymName); } /// Sets the value of the specialization constant whose address is SpecName @@ -267,7 +267,7 @@ class kernel_bundle : public detail::kernel_bundle_plain { void set_specialization_constant( typename std::remove_reference_t::value_type Value) { - using SCType = std::remove_reference_t::value_type; + using SCType = typename std::remove_reference_t::value_type; static_assert(std::is_trivially_copyable_v); // TODO can this be simply default constructible static_assert(std::is_trivially_default_constructible_v); @@ -282,7 +282,7 @@ class kernel_bundle : public detail::kernel_bundle_plain { typename std::remove_reference_t::value_type get_specialization_constant() const { - using SCType = std::remove_reference_t::value_type; + using SCType = typename std::remove_reference_t::value_type; static_assert(std::is_trivially_copyable_v); // TODO can this be simply default constructible static_assert(std::is_trivially_default_constructible_v); diff --git a/sycl/include/CL/sycl/kernel_handler.hpp b/sycl/include/CL/sycl/kernel_handler.hpp index 0a3f77381b364..b6a2e96d20313 100644 --- a/sycl/include/CL/sycl/kernel_handler.hpp +++ b/sycl/include/CL/sycl/kernel_handler.hpp @@ -8,6 +8,9 @@ #pragma once +#include +#include + __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { namespace detail { diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index f9c982bb1c9a1..2bfa3a6dc7551 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -178,11 +178,11 @@ class device_image_impl { [&](const _pi_device_binary_property_struct &Prop) { MSpecConstSymMap[Prop.Name] = *static_cast(Prop.ValAddr); - MSpecConstDescs.emplace_back( + MSpecConstDescs.push_back(SpecConstDescT{ *static_cast(Prop.ValAddr), // ID *(static_cast(Prop.ValAddr) + 1), // Offset false - ); + }); }); } } diff --git a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp new file mode 100644 index 0000000000000..6ea665205b06f --- /dev/null +++ b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp @@ -0,0 +1,41 @@ +#include + +class Kernel1Name; +class Kernel2Name; + +static sycl::specialization_id SpecConst1{42}; +static sycl::specialization_id SpecConst2{42.f}; +static sycl::specialization_id SpecConst3{42.f}; +static sycl::specialization_id SpecConst4{42}; + +int main() { + sycl::queue Q; + + // No support for host device so far + if (Q.is_host()) + return 0; + + // The code is needed to just have device images in the executable + if (0) { + Q.submit([](sycl::handler &CGH) { CGH.single_task([]{}); }); + Q.submit([](sycl::handler &CGH) { CGH.single_task([]{}); }); + } + + const sycl::context Ctx = Q.get_context(); + const sycl::device Dev = Q.get_device(); + + sycl::kernel_bundle KernelBundle = + sycl::get_kernel_bundle(Ctx, {Dev}); + + assert(KernelBundle.has_specialization_constant() == true); + KernelBundle.set_specialization_constant(1.f); + const auto SC = KernelBundle.get_specialization_constant(); + + Q.submit([](sycl::handler &CGH) { + CGH.set_specialization_constant(0.f); + const auto SC = CGH.get_specialization_constant(); + CGH.single_task([]{}); + }); + + return 0; +} From 7c138e629086bd0b0ed4705f3522d97b942fe656 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Tue, 6 Apr 2021 13:45:34 +0300 Subject: [PATCH 05/21] Improve tests and fix some errors --- .../CL/sycl/detail/sycl_fe_intrins.hpp | 4 +- sycl/include/CL/sycl/kernel_handler.hpp | 9 ++-- .../specialization_constants/host_apis.cpp | 49 ++++++++++++++----- 3 files changed, 45 insertions(+), 17 deletions(-) diff --git a/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp b/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp index 461b1b9b52d2a..403b418f92f2f 100644 --- a/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp +++ b/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp @@ -37,12 +37,12 @@ SYCL_EXTERNAL T __sycl_getCompositeSpecConstantValue(const char *ID); // are not available. template SYCL_EXTERNAL T __sycl_getScalar2020SpecConstantValue(const char *SymbolicID, - void *DefaultValue, + const void *DefaultValue, void *RTBuffer); template SYCL_EXTERNAL T __sycl_getComposite2020SpecConstantValue(const char *SymbolicID, - void *DefaultValue, + const void *DefaultValue, void *RTBuffer); // Request a fixed-size allocation in local address space at kernel scope. diff --git a/sycl/include/CL/sycl/kernel_handler.hpp b/sycl/include/CL/sycl/kernel_handler.hpp index b6a2e96d20313..eecaa0a3919d5 100644 --- a/sycl/include/CL/sycl/kernel_handler.hpp +++ b/sycl/include/CL/sycl/kernel_handler.hpp @@ -8,9 +8,10 @@ #pragma once -#include #include +#include + __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { namespace detail { @@ -31,7 +32,7 @@ class kernel_handler { public: #if __cplusplus > 201402L template - typename std::remove_reference_t get_specialization_constant() { + typename std::remove_reference_t::value_type get_specialization_constant() { #ifdef __SYCL_DEVICE_ONLY__ return getSpecializationConstantOnDevice(); #else @@ -51,7 +52,7 @@ class kernel_handler { } #ifdef __SYCL_DEVICE_ONLY__ - template , + template ::value_type, std::enable_if_t> * = nullptr> T getSpecializationConstantOnDevice() { const char *SymbolicID = __builtin_unique_stable_name( @@ -59,7 +60,7 @@ class kernel_handler { return __sycl_getScalar2020SpecConstantValue( SymbolicID, &S, MSpecializationConstantsBuffer); } - template , + template ::value_type, std::enable_if_t> * = nullptr> T getSpecializationConstantOnDevice() { const char *SymbolicID = __builtin_unique_stable_name( diff --git a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp index 6ea665205b06f..5589abeadba1c 100644 --- a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp +++ b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp @@ -1,12 +1,13 @@ +#include #include class Kernel1Name; class Kernel2Name; -static sycl::specialization_id SpecConst1{42}; -static sycl::specialization_id SpecConst2{42.f}; -static sycl::specialization_id SpecConst3{42.f}; -static sycl::specialization_id SpecConst4{42}; +const static sycl::specialization_id SpecConst1{42}; +const static sycl::specialization_id SpecConst2{42.f}; +const static sycl::specialization_id SpecConst3{42.f}; +const static sycl::specialization_id SpecConst4{42}; int main() { sycl::queue Q; @@ -27,15 +28,41 @@ int main() { sycl::kernel_bundle KernelBundle = sycl::get_kernel_bundle(Ctx, {Dev}); - assert(KernelBundle.has_specialization_constant() == true); + assert(KernelBundle.has_specialization_constant() == false); KernelBundle.set_specialization_constant(1.f); - const auto SC = KernelBundle.get_specialization_constant(); + // TODO uncomment once spec constants work correctly. + /* + { + auto ExecBundle = sycl::build(KernelBundle); + sycl::buffer Buf{sycl::range{1}}; + sycl::event Evt = Q.submit([&](sycl::handler &CGH) { + CGH.use_kernel_bundle(ExecBundle); + auto Acc = Buf.get_access(CGH); + CGH.single_task([=](sycl::kernel_handler KH) { + Acc[0] = KH.get_specialization_constant(); + }); + }); + Evt.wait(); + auto Acc = Buf.get_access(); + assert(std::fabs(Acc[0] - 1.f) <= 0.01); + } + assert(KernelBundle.has_specialization_constant() == true); - Q.submit([](sycl::handler &CGH) { - CGH.set_specialization_constant(0.f); - const auto SC = CGH.get_specialization_constant(); - CGH.single_task([]{}); - }); + { + sycl::buffer Buf{sycl::range{1}}; + sycl::event Evt = Q.submit([](sycl::handler &CGH) { + auto Acc = Buf.get_access(CGH); + CGH.set_specialization_constant(0.f); + const auto SC = CGH.get_specialization_constant(); + CGH.single_task([](sycl::kernel_handler KH) { + Acc[0] = KH.get_specialization_constant(); + }); + Evt.wait(); + auto Acc = Buf.get_access(); + assert(std::fabs(Acc[0]) <= 0.01); + }); + } + */ return 0; } From 0cddba7b6579e9068c4d527bba5598cf536a60b0 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Tue, 6 Apr 2021 13:49:54 +0300 Subject: [PATCH 06/21] More fixes --- sycl/include/CL/sycl/kernel_bundle.hpp | 5 ++++- sycl/source/detail/device_image_impl.hpp | 1 + 2 files changed, 5 insertions(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/kernel_bundle.hpp b/sycl/include/CL/sycl/kernel_bundle.hpp index 7168c5f46105a..8ce7b6ca5037f 100644 --- a/sycl/include/CL/sycl/kernel_bundle.hpp +++ b/sycl/include/CL/sycl/kernel_bundle.hpp @@ -281,12 +281,15 @@ class kernel_bundle : public detail::kernel_bundle_plain { template typename std::remove_reference_t::value_type get_specialization_constant() const { - using SCType = typename std::remove_reference_t::value_type; static_assert(std::is_trivially_copyable_v); // TODO can this be simply default constructible static_assert(std::is_trivially_default_constructible_v); + if (!has_specialization_constant()) + // TODO replace with SYCL 2020 exception + throw sycl::runtime_error("Unknown specialization constant", PI_INVALID_VALUE); + SCType RetValue; const char *SpecSymName = detail::get_spec_constant_symbolic_ID(); diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index 2bfa3a6dc7551..236fc85b4f2b0 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -109,6 +109,7 @@ class device_image_impl { void get_specialization_constant_raw_value(const char *SpecName, void *ValueRet, size_t ValueSize) const noexcept { + // operator[] can't be used here, since it's not marked as const unsigned SpecID = MSpecConstSymMap.at(SpecName); for (const SpecConstDescT &SpecConstDesc : MSpecConstDescs) if (SpecConstDesc.ID == SpecID) { From e8bfcc7d9f767410c60b3d1dc086307314a888b2 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Wed, 7 Apr 2021 09:11:16 +0300 Subject: [PATCH 07/21] add run line --- .../basic_tests/specialization_constants/host_apis.cpp | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp index 5589abeadba1c..086ab8c636da5 100644 --- a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp +++ b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp @@ -1,6 +1,12 @@ -#include +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %t.out + +// UNSUPPORTED: cuda + #include +#include + class Kernel1Name; class Kernel2Name; From 4d05216ae71be066a02de5022b5a44568dc86c73 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Wed, 7 Apr 2021 09:13:52 +0300 Subject: [PATCH 08/21] clang-format --- .../CL/sycl/detail/sycl_fe_intrins.hpp | 5 +- sycl/include/CL/sycl/kernel_bundle.hpp | 9 ++- sycl/include/CL/sycl/kernel_handler.hpp | 15 +++-- sycl/source/detail/device_image_impl.hpp | 62 +++++++++---------- .../specialization_constants/host_apis.cpp | 6 +- 5 files changed, 52 insertions(+), 45 deletions(-) diff --git a/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp b/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp index 403b418f92f2f..0043dcf291bee 100644 --- a/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp +++ b/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp @@ -41,9 +41,8 @@ SYCL_EXTERNAL T __sycl_getScalar2020SpecConstantValue(const char *SymbolicID, void *RTBuffer); template -SYCL_EXTERNAL T __sycl_getComposite2020SpecConstantValue(const char *SymbolicID, - const void *DefaultValue, - void *RTBuffer); +SYCL_EXTERNAL T __sycl_getComposite2020SpecConstantValue( + const char *SymbolicID, const void *DefaultValue, void *RTBuffer); // Request a fixed-size allocation in local address space at kernel scope. extern "C" SYCL_EXTERNAL __attribute__((opencl_local)) std::uint8_t * diff --git a/sycl/include/CL/sycl/kernel_bundle.hpp b/sycl/include/CL/sycl/kernel_bundle.hpp index 8ce7b6ca5037f..bb8dccdcc788f 100644 --- a/sycl/include/CL/sycl/kernel_bundle.hpp +++ b/sycl/include/CL/sycl/kernel_bundle.hpp @@ -267,7 +267,8 @@ class kernel_bundle : public detail::kernel_bundle_plain { void set_specialization_constant( typename std::remove_reference_t::value_type Value) { - using SCType = typename std::remove_reference_t::value_type; + using SCType = + typename std::remove_reference_t::value_type; static_assert(std::is_trivially_copyable_v); // TODO can this be simply default constructible static_assert(std::is_trivially_default_constructible_v); @@ -281,14 +282,16 @@ class kernel_bundle : public detail::kernel_bundle_plain { template typename std::remove_reference_t::value_type get_specialization_constant() const { - using SCType = typename std::remove_reference_t::value_type; + using SCType = + typename std::remove_reference_t::value_type; static_assert(std::is_trivially_copyable_v); // TODO can this be simply default constructible static_assert(std::is_trivially_default_constructible_v); if (!has_specialization_constant()) // TODO replace with SYCL 2020 exception - throw sycl::runtime_error("Unknown specialization constant", PI_INVALID_VALUE); + throw sycl::runtime_error("Unknown specialization constant", + PI_INVALID_VALUE); SCType RetValue; diff --git a/sycl/include/CL/sycl/kernel_handler.hpp b/sycl/include/CL/sycl/kernel_handler.hpp index eecaa0a3919d5..9a508d5611b3c 100644 --- a/sycl/include/CL/sycl/kernel_handler.hpp +++ b/sycl/include/CL/sycl/kernel_handler.hpp @@ -32,7 +32,8 @@ class kernel_handler { public: #if __cplusplus > 201402L template - typename std::remove_reference_t::value_type get_specialization_constant() { + typename std::remove_reference_t::value_type + get_specialization_constant() { #ifdef __SYCL_DEVICE_ONLY__ return getSpecializationConstantOnDevice(); #else @@ -52,16 +53,20 @@ class kernel_handler { } #ifdef __SYCL_DEVICE_ONLY__ - template ::value_type, - std::enable_if_t> * = nullptr> + template < + auto &S, + typename T = typename std::remove_reference_t::value_type, + std::enable_if_t> * = nullptr> T getSpecializationConstantOnDevice() { const char *SymbolicID = __builtin_unique_stable_name( detail::specialization_id_name_generator); return __sycl_getScalar2020SpecConstantValue( SymbolicID, &S, MSpecializationConstantsBuffer); } - template ::value_type, - std::enable_if_t> * = nullptr> + template < + auto &S, + typename T = typename std::remove_reference_t::value_type, + std::enable_if_t> * = nullptr> T getSpecializationConstantOnDevice() { const char *SymbolicID = __builtin_unique_stable_name( detail::specialization_id_name_generator); diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index 236fc85b4f2b0..e9c8c99813919 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -175,41 +175,41 @@ class device_image_impl { const auto &PropsBegin = SpecConstMap->PropertiesBegin; const auto &PropsEnd = SpecConstMap->PropertiesEnd; - std::for_each(PropsBegin, PropsEnd, - [&](const _pi_device_binary_property_struct &Prop) { - MSpecConstSymMap[Prop.Name] = - *static_cast(Prop.ValAddr); - MSpecConstDescs.push_back(SpecConstDescT{ - *static_cast(Prop.ValAddr), // ID - *(static_cast(Prop.ValAddr) + 1), // Offset - false - }); - }); + std::for_each( + PropsBegin, PropsEnd, + [&](const _pi_device_binary_property_struct &Prop) { + MSpecConstSymMap[Prop.Name] = + *static_cast(Prop.ValAddr); + MSpecConstDescs.push_back(SpecConstDescT{ + *static_cast(Prop.ValAddr), // ID + *(static_cast(Prop.ValAddr) + 1), // Offset + false}); + }); } } } - const RTDeviceBinaryImage *MBinImage = nullptr; - context MContext; - std::vector MDevices; - bundle_state MState; - // Native program handler which this device image represents - RT::PiProgram MProgram = nullptr; - // List of kernel ids available in this image, elements should be sorted - // according to LessByNameComp - std::vector MKernelIDs; - - // A mutex for sycnhronizing access to spec constants blob. Mutable because - // needs to be locked in the const method for getting spec constant value. - mutable std::mutex MSpecConstAccessMtx; - // Binary blob which can have values of all specialization constants in the - // image - std::vector MSpecConstsBlob; - // Contains list of spec ID + their offsets in the MSpecConstsBlob - std::vector MSpecConstDescs; - - std::map MSpecConstSymMap; - }; + const RTDeviceBinaryImage *MBinImage = nullptr; + context MContext; + std::vector MDevices; + bundle_state MState; + // Native program handler which this device image represents + RT::PiProgram MProgram = nullptr; + // List of kernel ids available in this image, elements should be sorted + // according to LessByNameComp + std::vector MKernelIDs; + + // A mutex for sycnhronizing access to spec constants blob. Mutable because + // needs to be locked in the const method for getting spec constant value. + mutable std::mutex MSpecConstAccessMtx; + // Binary blob which can have values of all specialization constants in the + // image + std::vector MSpecConstsBlob; + // Contains list of spec ID + their offsets in the MSpecConstsBlob + std::vector MSpecConstDescs; + + std::map MSpecConstSymMap; +}; } // namespace detail } // namespace sycl diff --git a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp index 086ab8c636da5..f78aac0cca077 100644 --- a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp +++ b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp @@ -24,10 +24,10 @@ int main() { // The code is needed to just have device images in the executable if (0) { - Q.submit([](sycl::handler &CGH) { CGH.single_task([]{}); }); - Q.submit([](sycl::handler &CGH) { CGH.single_task([]{}); }); + Q.submit([](sycl::handler &CGH) { CGH.single_task([] {}); }); + Q.submit([](sycl::handler &CGH) { CGH.single_task([] {}); }); } - + const sycl::context Ctx = Q.get_context(); const sycl::device Dev = Q.get_device(); From f3e011e17c043a0a8ef7722860291455717b092e Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Wed, 7 Apr 2021 10:29:33 +0300 Subject: [PATCH 09/21] Update symbol table --- sycl/test/abi/sycl_symbols_linux.dump | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 4601549761311..e58303e8a0d17 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3786,6 +3786,7 @@ _ZN2cl4sycl6detail18stringifyErrorCodeEi _ZN2cl4sycl6detail19convertChannelOrderE23_pi_image_channel_order _ZN2cl4sycl6detail19convertChannelOrderENS0_19image_channel_orderE _ZN2cl4sycl6detail19getImageElementSizeEhNS0_18image_channel_typeE +_ZN2cl4sycl6detail19kernel_bundle_plain32set_specialization_constant_implEPKcPvm _ZN2cl4sycl6detail20associateWithHandlerERNS0_7handlerEPNS1_16AccessorBaseHostENS0_6access6targetE _ZN2cl4sycl6detail20getDeviceFromHandlerERNS0_7handlerE _ZN2cl4sycl6detail22addHostAccessorAndWaitEPNS1_16AccessorImplHostE @@ -3977,6 +3978,8 @@ _ZNK2cl4sycl6detail19kernel_bundle_plain11get_contextEv _ZNK2cl4sycl6detail19kernel_bundle_plain11get_devicesEv _ZNK2cl4sycl6detail19kernel_bundle_plain14get_kernel_idsEv _ZNK2cl4sycl6detail19kernel_bundle_plain30native_specialization_constantEv +_ZNK2cl4sycl6detail19kernel_bundle_plain32get_specialization_constant_implEPKcPvm +_ZNK2cl4sycl6detail19kernel_bundle_plain32has_specialization_constant_implEPKc _ZNK2cl4sycl6detail19kernel_bundle_plain33contains_specialization_constantsEv _ZNK2cl4sycl6detail19kernel_bundle_plain3endEv _ZNK2cl4sycl6detail19kernel_bundle_plain5beginEv From af901b2273a25dd7c1fc682f528fe547d462276a Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Thu, 8 Apr 2021 10:38:42 +0300 Subject: [PATCH 10/21] Address feedback --- sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp | 4 ++-- sycl/include/CL/sycl/kernel_bundle.hpp | 12 ------------ 2 files changed, 2 insertions(+), 14 deletions(-) diff --git a/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp b/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp index 0043dcf291bee..a79c496c3a225 100644 --- a/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp +++ b/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp @@ -38,11 +38,11 @@ SYCL_EXTERNAL T __sycl_getCompositeSpecConstantValue(const char *ID); template SYCL_EXTERNAL T __sycl_getScalar2020SpecConstantValue(const char *SymbolicID, const void *DefaultValue, - void *RTBuffer); + const void *RTBuffer); template SYCL_EXTERNAL T __sycl_getComposite2020SpecConstantValue( - const char *SymbolicID, const void *DefaultValue, void *RTBuffer); + const char *SymbolicID, const void *DefaultValue, const void *RTBuffer); // Request a fixed-size allocation in local address space at kernel scope. extern "C" SYCL_EXTERNAL __attribute__((opencl_local)) std::uint8_t * diff --git a/sycl/include/CL/sycl/kernel_bundle.hpp b/sycl/include/CL/sycl/kernel_bundle.hpp index bb8dccdcc788f..e90704fab47fa 100644 --- a/sycl/include/CL/sycl/kernel_bundle.hpp +++ b/sycl/include/CL/sycl/kernel_bundle.hpp @@ -266,12 +266,8 @@ class kernel_bundle : public detail::kernel_bundle_plain { typename = detail::enable_if_t<_State == bundle_state::input>> void set_specialization_constant( typename std::remove_reference_t::value_type Value) { - using SCType = typename std::remove_reference_t::value_type; - static_assert(std::is_trivially_copyable_v); - // TODO can this be simply default constructible - static_assert(std::is_trivially_default_constructible_v); const char *SpecSymName = detail::get_spec_constant_symbolic_ID(); set_specialization_constant_impl(SpecSymName, &Value, sizeof(SCType)); @@ -284,14 +280,6 @@ class kernel_bundle : public detail::kernel_bundle_plain { get_specialization_constant() const { using SCType = typename std::remove_reference_t::value_type; - static_assert(std::is_trivially_copyable_v); - // TODO can this be simply default constructible - static_assert(std::is_trivially_default_constructible_v); - - if (!has_specialization_constant()) - // TODO replace with SYCL 2020 exception - throw sycl::runtime_error("Unknown specialization constant", - PI_INVALID_VALUE); SCType RetValue; From 203c17b6474857ab6f33a8f2ed99ed50b5b3d2be Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Thu, 8 Apr 2021 10:51:35 +0300 Subject: [PATCH 11/21] Fix kernel_desc --- sycl/include/CL/sycl/detail/kernel_desc.hpp | 8 +++++++- sycl/include/CL/sycl/kernel_handler.hpp | 12 +----------- 2 files changed, 8 insertions(+), 12 deletions(-) diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp index 052b59cec82ce..3851651cf4a8a 100644 --- a/sycl/include/CL/sycl/detail/kernel_desc.hpp +++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp @@ -12,12 +12,18 @@ #include #include -#include __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { namespace detail { +// This guard is needed because the libsycl.so can compiled with C++ <=14 +// while the code requires C++17. This code is not supposed to be used by the +// libsycl.so so it should not be a problem. +#if __cplusplus > 201402L +template struct specialization_id_name_generator {}; +#endif + #ifndef __SYCL_DEVICE_ONLY__ #define _Bool bool #endif diff --git a/sycl/include/CL/sycl/kernel_handler.hpp b/sycl/include/CL/sycl/kernel_handler.hpp index 9a508d5611b3c..fd5415b802e28 100644 --- a/sycl/include/CL/sycl/kernel_handler.hpp +++ b/sycl/include/CL/sycl/kernel_handler.hpp @@ -8,23 +8,13 @@ #pragma once +#include #include #include __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { -namespace detail { - -// This guard is needed because the libsycl.so can compiled with C++ <=14 -// while the code requires C++17. This code is not supposed to be used by the -// libsycl.so so it should not be a problem. -#if __cplusplus > 201402L -template struct specialization_id_name_generator {}; -#endif - -} // namespace detail - /// Reading the value of a specialization constant /// /// \ingroup sycl_api From 80ad75cc5aaccab6c7e60b4063cc99bbf7e99488 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Mon, 12 Apr 2021 11:10:44 +0300 Subject: [PATCH 12/21] Address feedback --- sycl/include/CL/sycl/kernel_bundle.hpp | 26 ++-- sycl/include/CL/sycl/specialization_id.hpp | 3 + sycl/source/detail/device_image_impl.hpp | 134 ++++++++++-------- sycl/source/detail/kernel_bundle_impl.hpp | 20 +-- .../program_manager/program_manager.cpp | 20 +-- sycl/source/kernel_bundle.cpp | 19 +-- 6 files changed, 122 insertions(+), 100 deletions(-) diff --git a/sycl/include/CL/sycl/kernel_bundle.hpp b/sycl/include/CL/sycl/kernel_bundle.hpp index e90704fab47fa..c8524492c8eb7 100644 --- a/sycl/include/CL/sycl/kernel_bundle.hpp +++ b/sycl/include/CL/sycl/kernel_bundle.hpp @@ -170,11 +170,13 @@ class __SYCL_EXPORT kernel_bundle_plain { bool has_specialization_constant_impl(const char *SpecName) const noexcept; - void set_specialization_constant_impl(const char *SpecName, void *Value, - size_t Size); + void set_specialization_constant_impl(const char *SpecName, + void *Value) noexcept; - void get_specialization_constant_impl(const char *SpecName, void *Value, - size_t Size) const; + void get_specialization_constant_impl(const char *SpecName, + void *Value) const noexcept; + + bool is_specialization_constant_set(const char *SpecName) const noexcept; detail::KernelBundleImplPtr impl; }; @@ -266,11 +268,8 @@ class kernel_bundle : public detail::kernel_bundle_plain { typename = detail::enable_if_t<_State == bundle_state::input>> void set_specialization_constant( typename std::remove_reference_t::value_type Value) { - using SCType = - typename std::remove_reference_t::value_type; - const char *SpecSymName = detail::get_spec_constant_symbolic_ID(); - set_specialization_constant_impl(SpecSymName, &Value, sizeof(SCType)); + set_specialization_constant_impl(SpecSymName, &Value); } /// \returns the value of the specialization constant whose address is @@ -278,15 +277,18 @@ class kernel_bundle : public detail::kernel_bundle_plain { template typename std::remove_reference_t::value_type get_specialization_constant() const { + const char *SpecSymName = detail::get_spec_constant_symbolic_ID(); + if (!is_specialization_constant_set(SpecSymName)) + return SpecName.getDefaultValue(); + using SCType = typename std::remove_reference_t::value_type; - SCType RetValue; + std::array RetValue; - const char *SpecSymName = detail::get_spec_constant_symbolic_ID(); - get_specialization_constant_impl(SpecSymName, &RetValue, sizeof(SCType)); + get_specialization_constant_impl(SpecSymName, RetValue.data()); - return RetValue; + return *reinterpret_cast(RetValue.data()); } #endif diff --git a/sycl/include/CL/sycl/specialization_id.hpp b/sycl/include/CL/sycl/specialization_id.hpp index cc4f25d8ebad6..0023507a7702a 100644 --- a/sycl/include/CL/sycl/specialization_id.hpp +++ b/sycl/include/CL/sycl/specialization_id.hpp @@ -28,6 +28,9 @@ template class specialization_id { specialization_id &operator=(specialization_id &&rhs) = delete; private: + template friend class kernel_bundle; + T getDefaultValue() const noexcept { return MDefaultValue; } + T MDefaultValue; }; diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index e9c8c99813919..26e392e591c47 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -74,53 +74,58 @@ class device_image_impl { // for this spec const should be. struct SpecConstDescT { unsigned int ID = 0; + unsigned int BlobOffset = 0; unsigned int Offset = 0; + unsigned int Size = 0; bool IsSet = false; }; bool has_specialization_constant(const char *SpecName) const noexcept { - if (MSpecConstSymMap.count(SpecName) == 0) - return false; - - unsigned SpecID = MSpecConstSymMap.at(SpecName); - return std::any_of(MSpecConstDescs.begin(), MSpecConstDescs.end(), - [SpecID](const SpecConstDescT &SpecConstDesc) { - return SpecConstDesc.ID == SpecID; - }); + return MSpecConstSymMap.count(SpecName) == 0; } void set_specialization_constant_raw_value(const char *SpecName, - const void *Value, - size_t ValueSize) noexcept { - unsigned SpecID = MSpecConstSymMap[SpecName]; - for (SpecConstDescT &SpecConstDesc : MSpecConstDescs) - if (SpecConstDesc.ID == SpecID) { - SpecConstDesc.IsSet = true; - // Lock the mutex to prevent when one thread in the middle of writing a - // new value while another thread is reading the value to pass it to - // JIT compiler. - const std::lock_guard SpecConstLock(MSpecConstAccessMtx); - std::memcpy(MSpecConstsBlob.data() + SpecConstDesc.Offset, Value, - ValueSize); - return; - } + const void *Value) noexcept { + // Lock the mutex to prevent when one thread in the middle of writing a + // new value while another thread is reading the value to pass it to + // JIT compiler. + const std::lock_guard SpecConstLock(MSpecConstAccessMtx); + + std::vector &Descs = MSpecConstSymMap[SpecName]; + for (SpecConstDescT &Desc : Descs) { + Desc.IsSet = true; + std::memcpy(MSpecConstsBlob.data() + Desc.BlobOffset, + static_cast(Value) + Desc.Offset, Desc.Size); + } } void get_specialization_constant_raw_value(const char *SpecName, - void *ValueRet, - size_t ValueSize) const noexcept { + void *ValueRet) const noexcept { + // Lock the mutex to prevent when one thread in the middle of writing a + // new value while another thread is reading the value to pass it to + // JIT compiler. + const std::lock_guard SpecConstLock(MSpecConstAccessMtx); + // operator[] can't be used here, since it's not marked as const - unsigned SpecID = MSpecConstSymMap.at(SpecName); - for (const SpecConstDescT &SpecConstDesc : MSpecConstDescs) - if (SpecConstDesc.ID == SpecID) { - // Lock the mutex to prevent when one thread in the middle of writing a - // new value while another thread is reading the value to pass it to - // JIT compiler. - const std::lock_guard SpecConstLock(MSpecConstAccessMtx); - std::memcpy(ValueRet, MSpecConstsBlob.data() + SpecConstDesc.Offset, - ValueSize); - return; - } + const std::vector &Descs = MSpecConstSymMap.at(SpecName); + for (const SpecConstDescT &Desc : Descs) { + + std::memcpy(static_cast(ValueRet) + Desc.Offset, + MSpecConstsBlob.data() + Desc.BlobOffset, Desc.Size); + } + } + + bool is_specialization_constant_set(const char *SpecName) const noexcept { + if (MSpecConstSymMap.count(SpecName) == 0) + return false; + + // Lock the mutex to prevent when one thread in the middle of writing a + // new value while another thread is reading the value to pass it to + // JIT compiler. + const std::lock_guard SpecConstLock(MSpecConstAccessMtx); + + const std::vector &Descs = MSpecConstSymMap.at(SpecName); + return Descs.front().IsSet; } bundle_state get_state() const noexcept { return MState; } @@ -147,8 +152,9 @@ class device_image_impl { return MSpecConstsBlob; } - std::vector &get_spec_const_offsets_ref() noexcept { - return MSpecConstDescs; + const std::map> & + get_spec_const_data_ref() const noexcept { + return MSpecConstSymMap; } ~device_image_impl() { @@ -162,29 +168,33 @@ class device_image_impl { private: void updateSpecConstSymMap() { if (MBinImage) { - const pi_device_binary_struct &RawImg = MBinImage->getRawData(); - const auto &PropSetsBegin = RawImg.PropertySetsBegin; - const auto &PropSetsEnd = RawImg.PropertySetsEnd; - const auto &SpecConstMap = std::find_if( - PropSetsBegin, PropSetsEnd, - [](const _pi_device_binary_property_set_struct &Set) { - return strcmp(Set.Name, __SYCL_PI_PROPERTY_SET_SPEC_CONST_MAP) == 0; - }); - - if (SpecConstMap != PropSetsEnd) { - const auto &PropsBegin = SpecConstMap->PropertiesBegin; - const auto &PropsEnd = SpecConstMap->PropertiesEnd; - - std::for_each( - PropsBegin, PropsEnd, - [&](const _pi_device_binary_property_struct &Prop) { - MSpecConstSymMap[Prop.Name] = - *static_cast(Prop.ValAddr); - MSpecConstDescs.push_back(SpecConstDescT{ - *static_cast(Prop.ValAddr), // ID - *(static_cast(Prop.ValAddr) + 1), // Offset - false}); - }); + const pi::DeviceBinaryImage::PropertyRange &SCRange = + MBinImage->getSpecConstants(); + using SCItTy = pi::DeviceBinaryImage::PropertyRange::ConstIterator; + + for (SCItTy SCIt : SCRange) { + const char *SCName = (*SCIt)->Name; + + pi::ByteArray Descriptors = + pi::DeviceBinaryProperty(*SCIt).asByteArray(); + assert(Descriptors.size() > 8 && "Unexpected property size"); + + // Expected layout is vector of 3-component tuples (flattened into a + // vector of scalars), where each tuple consists of: ID of a scalar spec + // constant, (which might be a member of the composite); offset, which + // is used to calculate location of scalar member within the composite + // or zero for scalar spec constants; size of a spec constant + assert(((Descriptors.size() - 8) / sizeof(std::uint32_t)) % 3 == 0 && + "unexpected layout of composite spec const descriptors"); + auto *It = reinterpret_cast(&Descriptors[8]); + auto *End = reinterpret_cast(&Descriptors[0] + + Descriptors.size()); + unsigned BlobOffset = 0; + while (It != End) { + MSpecConstSymMap[SCName].push_back(SpecConstDescT{ + /*ID*/ It[0], BlobOffset, /*Offset*/ It[1], It[2]}); + BlobOffset += /*Size*/ It[2]; + } } } } @@ -206,9 +216,9 @@ class device_image_impl { // image std::vector MSpecConstsBlob; // Contains list of spec ID + their offsets in the MSpecConstsBlob - std::vector MSpecConstDescs; + // std::vector MSpecConstDescs; - std::map MSpecConstSymMap; + std::map> MSpecConstSymMap; }; } // namespace detail diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index bc3a85bcb445c..20840a2b1c5ba 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -247,26 +247,28 @@ class kernel_bundle_impl { } void set_specialization_constant_raw_value(const char *SpecName, - const void *Value, - size_t ValueSize) { + const void *Value) noexcept { for (const device_image_plain &DeviceImage : MDeviceImages) getSyclObjImpl(DeviceImage) - ->set_specialization_constant_raw_value(SpecName, Value, ValueSize); + ->set_specialization_constant_raw_value(SpecName, Value); } void get_specialization_constant_raw_value(const char *SpecName, - void *ValueRet, - size_t ValueSize) const { + void *ValueRet) const noexcept { for (const device_image_plain &DeviceImage : MDeviceImages) if (getSyclObjImpl(DeviceImage)->has_specialization_constant(SpecName)) { getSyclObjImpl(DeviceImage) - ->get_specialization_constant_raw_value(SpecName, ValueRet, - ValueSize); + ->get_specialization_constant_raw_value(SpecName, ValueRet); return; } + } - throw sycl::runtime_error("Specialization constant not found", - PI_INVALID_VALUE); + bool is_specialization_constant_set(const char *SpecName) const noexcept { + return std::all_of(MDeviceImages.begin(), MDeviceImages.end(), + [SpecName](const device_image_plain &DeviceImage) { + return getSyclObjImpl(DeviceImage) + ->is_specialization_constant_set(SpecName); + }); } const device_image_plain *begin() const { return &MDeviceImages.front(); } diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index dd00f357032a2..ffdecbbf3f458 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -1540,17 +1540,19 @@ device_image_plain ProgramManager::build(const device_image_plain &DeviceImage, const std::vector &SpecConstsBlob = InputImpl->get_spec_const_blob_ref(); - std::vector &SpecConstOffsets = - InputImpl->get_spec_const_offsets_ref(); + const std::map> + &SpecConstData = InputImpl->get_spec_const_data_ref(); unsigned int PrevOffset = 0; - for (const device_image_impl::SpecConstDescT &SpecIDDesc : - SpecConstOffsets) { - - Plugin.call( - NativePrg, SpecIDDesc.ID, SpecIDDesc.Offset - PrevOffset, - SpecConstsBlob.data() + SpecIDDesc.Offset); - PrevOffset = SpecIDDesc.Offset; + for (const auto &DescPair : SpecConstData) { + for (const device_image_impl::SpecConstDescT &SpecIDDesc : + DescPair.second) { + + Plugin.call( + NativePrg, SpecIDDesc.ID, SpecIDDesc.BlobOffset - PrevOffset, + SpecConstsBlob.data() + SpecIDDesc.BlobOffset); + PrevOffset = SpecIDDesc.BlobOffset; + } } ProgramPtr ProgramManaged( diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index abe3400cb64d3..eab2731f137b6 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -88,16 +88,19 @@ bool kernel_bundle_plain::has_specialization_constant_impl( return impl->has_specialization_constant(SpecName); } -void kernel_bundle_plain::set_specialization_constant_impl(const char *SpecName, - void *Value, - size_t Size) { - impl->set_specialization_constant_raw_value(SpecName, Value, Size); +void kernel_bundle_plain::set_specialization_constant_impl( + const char *SpecName, void *Value) noexcept { + impl->set_specialization_constant_raw_value(SpecName, Value); } -void kernel_bundle_plain::get_specialization_constant_impl(const char *SpecName, - void *Value, - size_t Size) const { - impl->get_specialization_constant_raw_value(SpecName, Value, Size); +void kernel_bundle_plain::get_specialization_constant_impl( + const char *SpecName, void *Value) const noexcept { + impl->get_specialization_constant_raw_value(SpecName, Value); +} + +bool kernel_bundle_plain::is_specialization_constant_set( + const char *SpecName) const noexcept { + return impl->is_specialization_constant_set(SpecName); } //////////////////////////// From c0196ef65b09e874660916968cc394256e9db4e1 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Mon, 12 Apr 2021 11:12:25 +0300 Subject: [PATCH 13/21] update abi table --- 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 e58303e8a0d17..a21cab277a49b 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3786,7 +3786,7 @@ _ZN2cl4sycl6detail18stringifyErrorCodeEi _ZN2cl4sycl6detail19convertChannelOrderE23_pi_image_channel_order _ZN2cl4sycl6detail19convertChannelOrderENS0_19image_channel_orderE _ZN2cl4sycl6detail19getImageElementSizeEhNS0_18image_channel_typeE -_ZN2cl4sycl6detail19kernel_bundle_plain32set_specialization_constant_implEPKcPvm +_ZN2cl4sycl6detail19kernel_bundle_plain32set_specialization_constant_implEPKcPv _ZN2cl4sycl6detail20associateWithHandlerERNS0_7handlerEPNS1_16AccessorBaseHostENS0_6access6targetE _ZN2cl4sycl6detail20getDeviceFromHandlerERNS0_7handlerE _ZN2cl4sycl6detail22addHostAccessorAndWaitEPNS1_16AccessorImplHostE @@ -3977,8 +3977,9 @@ _ZNK2cl4sycl6detail19kernel_bundle_plain11get_backendEv _ZNK2cl4sycl6detail19kernel_bundle_plain11get_contextEv _ZNK2cl4sycl6detail19kernel_bundle_plain11get_devicesEv _ZNK2cl4sycl6detail19kernel_bundle_plain14get_kernel_idsEv +_ZNK2cl4sycl6detail19kernel_bundle_plain30is_specialization_constant_setEPKc _ZNK2cl4sycl6detail19kernel_bundle_plain30native_specialization_constantEv -_ZNK2cl4sycl6detail19kernel_bundle_plain32get_specialization_constant_implEPKcPvm +_ZNK2cl4sycl6detail19kernel_bundle_plain32get_specialization_constant_implEPKcPv _ZNK2cl4sycl6detail19kernel_bundle_plain32has_specialization_constant_implEPKc _ZNK2cl4sycl6detail19kernel_bundle_plain33contains_specialization_constantsEv _ZNK2cl4sycl6detail19kernel_bundle_plain3endEv From d55533cf05d89905f281d0c266eddb9146b908a6 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Mon, 12 Apr 2021 11:23:47 +0300 Subject: [PATCH 14/21] clang-format --- sycl/include/CL/sycl/kernel_bundle.hpp | 4 ++-- sycl/source/kernel_bundle.cpp | 5 +++-- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/sycl/include/CL/sycl/kernel_bundle.hpp b/sycl/include/CL/sycl/kernel_bundle.hpp index c8524492c8eb7..e8c93a02f079a 100644 --- a/sycl/include/CL/sycl/kernel_bundle.hpp +++ b/sycl/include/CL/sycl/kernel_bundle.hpp @@ -173,8 +173,8 @@ class __SYCL_EXPORT kernel_bundle_plain { void set_specialization_constant_impl(const char *SpecName, void *Value) noexcept; - void get_specialization_constant_impl(const char *SpecName, - void *Value) const noexcept; + void get_specialization_constant_impl(const char *SpecName, void *Value) const + noexcept; bool is_specialization_constant_set(const char *SpecName) const noexcept; diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index eab2731f137b6..b5ea314428d38 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -93,8 +93,9 @@ void kernel_bundle_plain::set_specialization_constant_impl( impl->set_specialization_constant_raw_value(SpecName, Value); } -void kernel_bundle_plain::get_specialization_constant_impl( - const char *SpecName, void *Value) const noexcept { +void kernel_bundle_plain::get_specialization_constant_impl(const char *SpecName, + void *Value) const + noexcept { impl->get_specialization_constant_raw_value(SpecName, Value); } From 6b4ce8aea024a9915c3324bc452123fa805d10dc Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Mon, 12 Apr 2021 13:07:18 +0300 Subject: [PATCH 15/21] small test fix --- sycl/source/detail/device_image_impl.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index 26e392e591c47..729746634c186 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -81,7 +81,7 @@ class device_image_impl { }; bool has_specialization_constant(const char *SpecName) const noexcept { - return MSpecConstSymMap.count(SpecName) == 0; + return MSpecConstSymMap.count(SpecName) != 0; } void set_specialization_constant_raw_value(const char *SpecName, @@ -94,6 +94,7 @@ class device_image_impl { std::vector &Descs = MSpecConstSymMap[SpecName]; for (SpecConstDescT &Desc : Descs) { Desc.IsSet = true; + MSpecConstsBlob.reserve(MSpecConstsBlob.size() + Desc.Size); std::memcpy(MSpecConstsBlob.data() + Desc.BlobOffset, static_cast(Value) + Desc.Offset, Desc.Size); } From d600ebfef11eae9d0eca346a6f7bcc8808cb62c9 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Mon, 12 Apr 2021 15:01:13 +0300 Subject: [PATCH 16/21] Address feedback --- sycl/source/detail/device_image_impl.hpp | 49 +++++++++++++------ .../program_manager/program_manager.cpp | 13 +++-- .../specialization_constants/host_apis.cpp | 15 +++--- 3 files changed, 45 insertions(+), 32 deletions(-) diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index 729746634c186..7772e873df47c 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -74,13 +74,17 @@ class device_image_impl { // for this spec const should be. struct SpecConstDescT { unsigned int ID = 0; - unsigned int BlobOffset = 0; - unsigned int Offset = 0; + unsigned int CompositeOffset = 0; unsigned int Size = 0; + unsigned int BlobOffset = 0; bool IsSet = false; }; bool has_specialization_constant(const char *SpecName) const noexcept { + // Lock the mutex to prevent when one thread in the middle of writing a + // new value while another thread is reading the value to pass it to + // JIT compiler. + const std::lock_guard SpecConstLock(MSpecConstAccessMtx); return MSpecConstSymMap.count(SpecName) != 0; } @@ -91,12 +95,17 @@ class device_image_impl { // JIT compiler. const std::lock_guard SpecConstLock(MSpecConstAccessMtx); - std::vector &Descs = MSpecConstSymMap[SpecName]; + if (MSpecConstSymMap.count(std::string{SpecName}) == 0) + return; + + std::vector &Descs = + MSpecConstSymMap[std::string{SpecName}]; for (SpecConstDescT &Desc : Descs) { Desc.IsSet = true; MSpecConstsBlob.reserve(MSpecConstsBlob.size() + Desc.Size); std::memcpy(MSpecConstsBlob.data() + Desc.BlobOffset, - static_cast(Value) + Desc.Offset, Desc.Size); + static_cast(Value) + Desc.CompositeOffset, + Desc.Size); } } @@ -108,16 +117,17 @@ class device_image_impl { const std::lock_guard SpecConstLock(MSpecConstAccessMtx); // operator[] can't be used here, since it's not marked as const - const std::vector &Descs = MSpecConstSymMap.at(SpecName); + const std::vector &Descs = + MSpecConstSymMap.at(std::string{SpecName}); for (const SpecConstDescT &Desc : Descs) { - std::memcpy(static_cast(ValueRet) + Desc.Offset, + std::memcpy(static_cast(ValueRet) + Desc.CompositeOffset, MSpecConstsBlob.data() + Desc.BlobOffset, Desc.Size); } } bool is_specialization_constant_set(const char *SpecName) const noexcept { - if (MSpecConstSymMap.count(SpecName) == 0) + if (MSpecConstSymMap.count(std::string{SpecName}) == 0) return false; // Lock the mutex to prevent when one thread in the middle of writing a @@ -125,7 +135,8 @@ class device_image_impl { // JIT compiler. const std::lock_guard SpecConstLock(MSpecConstAccessMtx); - const std::vector &Descs = MSpecConstSymMap.at(SpecName); + const std::vector &Descs = + MSpecConstSymMap.at(std::string{SpecName}); return Descs.front().IsSet; } @@ -153,7 +164,7 @@ class device_image_impl { return MSpecConstsBlob; } - const std::map> & + const std::map> & get_spec_const_data_ref() const noexcept { return MSpecConstSymMap; } @@ -185,16 +196,23 @@ class device_image_impl { // constant, (which might be a member of the composite); offset, which // is used to calculate location of scalar member within the composite // or zero for scalar spec constants; size of a spec constant - assert(((Descriptors.size() - 8) / sizeof(std::uint32_t)) % 3 == 0 && + constexpr size_t NumElements = 3; + assert(((Descriptors.size() - 8) / sizeof(std::uint32_t)) % + NumElements == + 0 && "unexpected layout of composite spec const descriptors"); auto *It = reinterpret_cast(&Descriptors[8]); auto *End = reinterpret_cast(&Descriptors[0] + Descriptors.size()); unsigned BlobOffset = 0; while (It != End) { - MSpecConstSymMap[SCName].push_back(SpecConstDescT{ - /*ID*/ It[0], BlobOffset, /*Offset*/ It[1], It[2]}); + // The map is not locked here because updateSpecConstSymMap() is only + // supposed to be called from c'tor. + MSpecConstSymMap[std::string{SCName}].push_back( + SpecConstDescT{/*ID*/ It[0], /*CompositeOffset*/ It[1], + /*Size*/ It[2], BlobOffset}); BlobOffset += /*Size*/ It[2]; + It += NumElements; } } } @@ -216,10 +234,9 @@ class device_image_impl { // Binary blob which can have values of all specialization constants in the // image std::vector MSpecConstsBlob; - // Contains list of spec ID + their offsets in the MSpecConstsBlob - // std::vector MSpecConstDescs; - - std::map> MSpecConstSymMap; + // Contains map of spec const names to their descriptions + offsets in + // the MSpecConstsBlob + std::map> MSpecConstSymMap; }; } // namespace detail diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index ffdecbbf3f458..6d2e67cf9db45 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -1540,18 +1540,17 @@ device_image_plain ProgramManager::build(const device_image_plain &DeviceImage, const std::vector &SpecConstsBlob = InputImpl->get_spec_const_blob_ref(); - const std::map> + const std::map> &SpecConstData = InputImpl->get_spec_const_data_ref(); - unsigned int PrevOffset = 0; for (const auto &DescPair : SpecConstData) { for (const device_image_impl::SpecConstDescT &SpecIDDesc : DescPair.second) { - - Plugin.call( - NativePrg, SpecIDDesc.ID, SpecIDDesc.BlobOffset - PrevOffset, - SpecConstsBlob.data() + SpecIDDesc.BlobOffset); - PrevOffset = SpecIDDesc.BlobOffset; + if (SpecIDDesc.IsSet) { + Plugin.call( + NativePrg, SpecIDDesc.ID, SpecIDDesc.Size, + SpecConstsBlob.data() + SpecIDDesc.BlobOffset); + } } } diff --git a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp index f78aac0cca077..c0976232c651f 100644 --- a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp +++ b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp @@ -35,9 +35,8 @@ int main() { sycl::get_kernel_bundle(Ctx, {Dev}); assert(KernelBundle.has_specialization_constant() == false); + assert(KernelBundle.has_specialization_constant() == true); KernelBundle.set_specialization_constant(1.f); - // TODO uncomment once spec constants work correctly. - /* { auto ExecBundle = sycl::build(KernelBundle); sycl::buffer Buf{sycl::range{1}}; @@ -52,23 +51,21 @@ int main() { auto Acc = Buf.get_access(); assert(std::fabs(Acc[0] - 1.f) <= 0.01); } - assert(KernelBundle.has_specialization_constant() == true); { sycl::buffer Buf{sycl::range{1}}; - sycl::event Evt = Q.submit([](sycl::handler &CGH) { + sycl::event Evt = Q.submit([&](sycl::handler &CGH) { auto Acc = Buf.get_access(CGH); CGH.set_specialization_constant(0.f); const auto SC = CGH.get_specialization_constant(); - CGH.single_task([](sycl::kernel_handler KH) { + CGH.single_task([=](sycl::kernel_handler KH) { Acc[0] = KH.get_specialization_constant(); }); - Evt.wait(); - auto Acc = Buf.get_access(); - assert(std::fabs(Acc[0]) <= 0.01); }); + Evt.wait(); + auto Acc = Buf.get_access(); + assert(std::fabs(Acc[0]) <= 0.01); } - */ return 0; } From 164ac17f30b8d1e7b6a435b5495f4769e18e2f4c Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Mon, 12 Apr 2021 15:02:11 +0300 Subject: [PATCH 17/21] Update sycl/source/detail/kernel_bundle_impl.hpp Co-authored-by: Romanov Vlad <17316488+romanovvlad@users.noreply.github.com> --- sycl/source/detail/kernel_bundle_impl.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 20840a2b1c5ba..ccfcd5134b585 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -264,7 +264,7 @@ class kernel_bundle_impl { } bool is_specialization_constant_set(const char *SpecName) const noexcept { - return std::all_of(MDeviceImages.begin(), MDeviceImages.end(), + return std::any_of(MDeviceImages.begin(), MDeviceImages.end(), [SpecName](const device_image_plain &DeviceImage) { return getSyclObjImpl(DeviceImage) ->is_specialization_constant_set(SpecName); From 6d1283cd43f0da7c054512f49d4fe9473196c392 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Mon, 12 Apr 2021 15:10:50 +0300 Subject: [PATCH 18/21] Last minute bugfix --- sycl/source/detail/device_image_impl.hpp | 6 +++++- .../basic_tests/specialization_constants/host_apis.cpp | 1 + 2 files changed, 6 insertions(+), 1 deletion(-) diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index 7772e873df47c..85cc884d2de0c 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -62,7 +62,11 @@ class device_image_impl { } bool has_specialization_constants() const noexcept { - return !MSpecConstsBlob.empty(); + // Lock the mutex to prevent when one thread in the middle of writing a + // new value while another thread is reading the value to pass it to + // JIT compiler. + const std::lock_guard SpecConstLock(MSpecConstAccessMtx); + return !MSpecConstSymMap.empty(); } bool all_specialization_constant_native() const noexcept { diff --git a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp index c0976232c651f..163dde5697a09 100644 --- a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp +++ b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp @@ -34,6 +34,7 @@ int main() { sycl::kernel_bundle KernelBundle = sycl::get_kernel_bundle(Ctx, {Dev}); + assert(KernelBundle.contains_specialization_constants() == true); assert(KernelBundle.has_specialization_constant() == false); assert(KernelBundle.has_specialization_constant() == true); KernelBundle.set_specialization_constant(1.f); From 7dc0f98465c877e0af7aca3c929c17dc10a6330b Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Tue, 13 Apr 2021 09:28:40 +0300 Subject: [PATCH 19/21] More feedback --- sycl/source/detail/device_image_impl.hpp | 12 +++++--- .../program_manager/program_manager.cpp | 24 +++++++++------- .../specialization_constants/host_apis.cpp | 28 +++++++++++-------- 3 files changed, 38 insertions(+), 26 deletions(-) diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index 85cc884d2de0c..a09d1d6403a9a 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -106,7 +106,6 @@ class device_image_impl { MSpecConstSymMap[std::string{SpecName}]; for (SpecConstDescT &Desc : Descs) { Desc.IsSet = true; - MSpecConstsBlob.reserve(MSpecConstsBlob.size() + Desc.Size); std::memcpy(MSpecConstsBlob.data() + Desc.BlobOffset, static_cast(Value) + Desc.CompositeOffset, Desc.Size); @@ -115,6 +114,7 @@ class device_image_impl { void get_specialization_constant_raw_value(const char *SpecName, void *ValueRet) const noexcept { + assert(is_specialization_constant_set(SpecName)); // Lock the mutex to prevent when one thread in the middle of writing a // new value while another thread is reading the value to pass it to // JIT compiler. @@ -131,13 +131,12 @@ class device_image_impl { } bool is_specialization_constant_set(const char *SpecName) const noexcept { - if (MSpecConstSymMap.count(std::string{SpecName}) == 0) - return false; - // Lock the mutex to prevent when one thread in the middle of writing a // new value while another thread is reading the value to pass it to // JIT compiler. const std::lock_guard SpecConstLock(MSpecConstAccessMtx); + if (MSpecConstSymMap.count(std::string{SpecName}) == 0) + return false; const std::vector &Descs = MSpecConstSymMap.at(std::string{SpecName}); @@ -173,6 +172,10 @@ class device_image_impl { return MSpecConstSymMap; } + std::mutex &get_spec_const_data_lock() noexcept { + return MSpecConstAccessMtx; + } + ~device_image_impl() { if (MProgram) { @@ -218,6 +221,7 @@ class device_image_impl { BlobOffset += /*Size*/ It[2]; It += NumElements; } + MSpecConstsBlob.resize(BlobOffset); } } } diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 6d2e67cf9db45..f04b6cc2aa74e 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -1540,16 +1540,20 @@ device_image_plain ProgramManager::build(const device_image_plain &DeviceImage, const std::vector &SpecConstsBlob = InputImpl->get_spec_const_blob_ref(); - const std::map> - &SpecConstData = InputImpl->get_spec_const_data_ref(); - - for (const auto &DescPair : SpecConstData) { - for (const device_image_impl::SpecConstDescT &SpecIDDesc : - DescPair.second) { - if (SpecIDDesc.IsSet) { - Plugin.call( - NativePrg, SpecIDDesc.ID, SpecIDDesc.Size, - SpecConstsBlob.data() + SpecIDDesc.BlobOffset); + { + std::lock_guard Lock{InputImpl->get_spec_const_data_lock()}; + const std::map> + &SpecConstData = InputImpl->get_spec_const_data_ref(); + + for (const auto &DescPair : SpecConstData) { + for (const device_image_impl::SpecConstDescT &SpecIDDesc : + DescPair.second) { + if (SpecIDDesc.IsSet) { + Plugin.call( + NativePrg, SpecIDDesc.ID, SpecIDDesc.Size, + SpecConstsBlob.data() + SpecIDDesc.BlobOffset); + } } } } diff --git a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp index 163dde5697a09..e73d2944dba66 100644 --- a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp +++ b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp @@ -10,9 +10,14 @@ class Kernel1Name; class Kernel2Name; +struct TestStruct { + int a; + int b; +}; + const static sycl::specialization_id SpecConst1{42}; -const static sycl::specialization_id SpecConst2{42.f}; -const static sycl::specialization_id SpecConst3{42.f}; +const static sycl::specialization_id SpecConst2{42}; +const static sycl::specialization_id SpecConst3{TestStruct{42, 42}}; const static sycl::specialization_id SpecConst4{42}; int main() { @@ -37,35 +42,34 @@ int main() { assert(KernelBundle.contains_specialization_constants() == true); assert(KernelBundle.has_specialization_constant() == false); assert(KernelBundle.has_specialization_constant() == true); - KernelBundle.set_specialization_constant(1.f); + KernelBundle.set_specialization_constant(1); { auto ExecBundle = sycl::build(KernelBundle); - sycl::buffer Buf{sycl::range{1}}; - sycl::event Evt = Q.submit([&](sycl::handler &CGH) { + sycl::buffer Buf{sycl::range{1}}; + Q.submit([&](sycl::handler &CGH) { CGH.use_kernel_bundle(ExecBundle); auto Acc = Buf.get_access(CGH); CGH.single_task([=](sycl::kernel_handler KH) { Acc[0] = KH.get_specialization_constant(); }); }); - Evt.wait(); auto Acc = Buf.get_access(); - assert(std::fabs(Acc[0] - 1.f) <= 0.01); + assert(Acc[0] == 1); } { - sycl::buffer Buf{sycl::range{1}}; - sycl::event Evt = Q.submit([&](sycl::handler &CGH) { + sycl::buffer Buf{sycl::range{1}}; + Q.submit([&](sycl::handler &CGH) { auto Acc = Buf.get_access(CGH); - CGH.set_specialization_constant(0.f); + CGH.set_specialization_constant(TestStruct{1, 2}); const auto SC = CGH.get_specialization_constant(); + assert(SC == 42); CGH.single_task([=](sycl::kernel_handler KH) { Acc[0] = KH.get_specialization_constant(); }); }); - Evt.wait(); auto Acc = Buf.get_access(); - assert(std::fabs(Acc[0]) <= 0.01); + assert(Acc[0].a == 1 && Acc[0].b == 2); } return 0; From 3dce625ed0b5e3440f4d679f0cc38d9e14b11dbe Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Tue, 13 Apr 2021 09:33:35 +0300 Subject: [PATCH 20/21] bug --- sycl/source/detail/device_image_impl.hpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index a09d1d6403a9a..9fa83db6aa150 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -191,6 +191,9 @@ class device_image_impl { MBinImage->getSpecConstants(); using SCItTy = pi::DeviceBinaryImage::PropertyRange::ConstIterator; + // This variable is used to calculate spec constant value offset in a + // flat byte array. + unsigned BlobOffset = 0; for (SCItTy SCIt : SCRange) { const char *SCName = (*SCIt)->Name; @@ -211,7 +214,6 @@ class device_image_impl { auto *It = reinterpret_cast(&Descriptors[8]); auto *End = reinterpret_cast(&Descriptors[0] + Descriptors.size()); - unsigned BlobOffset = 0; while (It != End) { // The map is not locked here because updateSpecConstSymMap() is only // supposed to be called from c'tor. @@ -221,8 +223,8 @@ class device_image_impl { BlobOffset += /*Size*/ It[2]; It += NumElements; } - MSpecConstsBlob.resize(BlobOffset); } + MSpecConstsBlob.resize(BlobOffset); } } From ea525f0b8205c6c724a27d5a8fe0a3ca7ca8dbf4 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Tue, 13 Apr 2021 17:34:08 +0300 Subject: [PATCH 21/21] minor update --- sycl/include/CL/sycl/detail/kernel_desc.hpp | 2 +- sycl/source/detail/kernel_bundle_impl.hpp | 2 ++ 2 files changed, 3 insertions(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp index 3851651cf4a8a..87359ff6c779b 100644 --- a/sycl/include/CL/sycl/detail/kernel_desc.hpp +++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp @@ -17,7 +17,7 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { namespace detail { -// This guard is needed because the libsycl.so can compiled with C++ <=14 +// This guard is needed because the libsycl.so can be compiled with C++ <=14 // while the code requires C++17. This code is not supposed to be used by the // libsycl.so so it should not be a problem. #if __cplusplus > 201402L diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index ccfcd5134b585..984fd6651189e 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -248,6 +248,8 @@ class kernel_bundle_impl { void set_specialization_constant_raw_value(const char *SpecName, const void *Value) noexcept { + // TODO add support for specialization constants, that are missing in any + // device image. for (const device_image_plain &DeviceImage : MDeviceImages) getSyclObjImpl(DeviceImage) ->set_specialization_constant_raw_value(SpecName, Value);