From ef79ff6f4ac5d5fbade7d343df7d298d94c1dd8a Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Mon, 27 Sep 2021 15:24:44 +0300 Subject: [PATCH 1/6] [SYCL] Handler exceptions on mutually exclusive operations Calling handler::set_specialization_constant after or before calling handler::use_kernel_bundle should cause the latter operation to throw a SYCL exception with error code errc::invalid. These changes enforces this behavior. Signed-off-by: Steffen Larsen --- sycl/include/CL/sycl/detail/cg.hpp | 1 + sycl/include/CL/sycl/handler.hpp | 16 +- sycl/source/handler.cpp | 129 +++++-- sycl/test/abi/sycl_symbols_linux.dump | 1 + sycl/unittests/SYCL2020/CMakeLists.txt | 2 +- .../SYCL2020/SpecConstDefaultValues.cpp | 171 --------- .../SYCL2020/SpecializationConstant.cpp | 351 ++++++++++++++++++ 7 files changed, 475 insertions(+), 196 deletions(-) delete mode 100644 sycl/unittests/SYCL2020/SpecConstDefaultValues.cpp create mode 100644 sycl/unittests/SYCL2020/SpecializationConstant.cpp diff --git a/sycl/include/CL/sycl/detail/cg.hpp b/sycl/include/CL/sycl/detail/cg.hpp index fc6042bf69183..71c22d37334a7 100644 --- a/sycl/include/CL/sycl/detail/cg.hpp +++ b/sycl/include/CL/sycl/detail/cg.hpp @@ -94,6 +94,7 @@ namespace detail { enum class ExtendedMembersType : unsigned int { HANDLER_KERNEL_BUNDLE = 0, HANDLER_MEM_ADVICE, + HANDLER_KERNEL_BUNDLE_FLAGS, }; // Holds a pointer to an object of an arbitrary type and an ID value which diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index cbe982540e07b..3390bf925103a 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -234,6 +234,12 @@ class RoundedRangeKernelWithKH { KernelType KernelFunc; }; +// Kernel bundle flags used to identify previous uses of use_kernel_bundle and +// set_specialization_constant. These are used to detect invalid command group +// operation ordering. +constexpr std::uint8_t EXPLICIT_KERNEL_BUNDLE_FLAG = 1; +constexpr std::uint8_t SPEC_CONST_SET_FLAG = 2; + } // namespace detail namespace ext { @@ -1119,6 +1125,10 @@ class __SYCL_EXPORT handler { std::shared_ptr getOrInsertHandlerKernelBundle(bool Insert) const; + std::shared_ptr + getOrInsertNonExplicitHandlerKernelBundle(bool Insert, + bool MarkSpecConstSet) const; + void setHandlerKernelBundle( const std::shared_ptr &NewKernelBundleImpPtr); @@ -1151,7 +1161,8 @@ class __SYCL_EXPORT handler { typename std::remove_reference_t::value_type Value) { std::shared_ptr KernelBundleImplPtr = - getOrInsertHandlerKernelBundle(/*Insert=*/true); + getOrInsertNonExplicitHandlerKernelBundle(/*Insert=*/true, + /*MarkSpecConstSet=*/true); detail::createSyclObjFromImpl>( KernelBundleImplPtr) @@ -1163,7 +1174,8 @@ class __SYCL_EXPORT handler { get_specialization_constant() const { std::shared_ptr KernelBundleImplPtr = - getOrInsertHandlerKernelBundle(/*Insert=*/true); + getOrInsertNonExplicitHandlerKernelBundle(/*Insert=*/true, + /*MarkSpecConstSet=*/false); return detail::createSyclObjFromImpl>( KernelBundleImplPtr) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 7f2b0f14dc7da..e25443dc417f9 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -31,21 +31,14 @@ handler::handler(std::shared_ptr Queue, bool IsHost) std::make_shared>()); } -// Returns a shared_ptr to kernel_bundle stored in the extended members vector. -// If there is no kernel_bundle created: -// returns newly created kernel_bundle if Insert is true -// returns shared_ptr(nullptr) if Insert is false +// Common implementation for getting/inserting handler kernel bundle. +// detail::GlobalHandler::instance().getHandlerExtendedMembersMutex() must be +// held when calling this function. std::shared_ptr -handler::getOrInsertHandlerKernelBundle(bool Insert) const { - - std::lock_guard Lock( - detail::GlobalHandler::instance().getHandlerExtendedMembersMutex()); - - assert(!MSharedPtrStorage.empty()); - - std::shared_ptr> ExendedMembersVec = - detail::convertToExtendedMembers(MSharedPtrStorage[0]); - +getOrInsertHandlerKernelBundleCommon( + const std::shared_ptr> + &ExendedMembersVec, + const std::shared_ptr &Queue, bool Insert) { // Look for the kernel bundle in extended members std::shared_ptr KernelBundleImpPtr; for (const detail::ExtendedMemberT &EMember : *ExendedMembersVec) @@ -58,10 +51,10 @@ handler::getOrInsertHandlerKernelBundle(bool Insert) const { // No kernel bundle yet, create one if (!KernelBundleImpPtr && Insert) { KernelBundleImpPtr = detail::getSyclObjImpl( - get_kernel_bundle(MQueue->get_context())); + get_kernel_bundle(Queue->get_context())); if (KernelBundleImpPtr->empty()) { KernelBundleImpPtr = detail::getSyclObjImpl( - get_kernel_bundle(MQueue->get_context())); + get_kernel_bundle(Queue->get_context())); } detail::ExtendedMemberT EMember = { @@ -73,6 +66,70 @@ handler::getOrInsertHandlerKernelBundle(bool Insert) const { return KernelBundleImpPtr; } +// Returns a shared_ptr to kernel_bundle stored in the extended members vector. +// If there is no kernel_bundle created: +// returns newly created kernel_bundle if Insert is true +// returns shared_ptr(nullptr) if Insert is false +std::shared_ptr +handler::getOrInsertHandlerKernelBundle(bool Insert) const { + + std::lock_guard Lock( + detail::GlobalHandler::instance().getHandlerExtendedMembersMutex()); + + assert(!MSharedPtrStorage.empty()); + + std::shared_ptr> ExendedMembersVec = + detail::convertToExtendedMembers(MSharedPtrStorage[0]); + + return getOrInsertHandlerKernelBundleCommon(ExendedMembersVec, MQueue, + Insert); +} + +// This function exhibits the same behavior as getOrInsertHandlerKernelBundle +// but throws an exception with errc::invalid if a kernel bundle has been set +// in the command group by a call to use_kernel_bundle. +// If MarkSpecConstSet is true the command group is flagged as having had set +// a specialization constant. +std::shared_ptr +handler::getOrInsertNonExplicitHandlerKernelBundle( + bool Insert, bool MarkSpecConstSet) const { + + std::lock_guard Lock( + detail::GlobalHandler::instance().getHandlerExtendedMembersMutex()); + + assert(!MSharedPtrStorage.empty()); + + std::shared_ptr> ExendedMembersVec = + detail::convertToExtendedMembers(MSharedPtrStorage[0]); + + // If kernel was explicitly set through use_kernel_bundle then throw exception + bool KernelBundleFlagsExist = false; + for (detail::ExtendedMemberT &EMember : *ExendedMembersVec) { + if (detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE_FLAGS == + EMember.MType) { + auto Flags = std::static_pointer_cast(EMember.MData); + if (*Flags & detail::EXPLICIT_KERNEL_BUNDLE_FLAG) + throw sycl::exception( + make_error_code(errc::invalid), + "Specialization constants cannot be accessed after explicitly " + "setting the used kernel bundle"); + if (MarkSpecConstSet) + *Flags |= detail::SPEC_CONST_SET_FLAG; + KernelBundleFlagsExist = true; + } + } + + if (!KernelBundleFlagsExist && MarkSpecConstSet) { + detail::ExtendedMemberT EMember = { + detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE_FLAGS, + std::make_shared(detail::SPEC_CONST_SET_FLAG)}; + ExendedMembersVec->push_back(EMember); + } + + return getOrInsertHandlerKernelBundleCommon(ExendedMembersVec, MQueue, + Insert); +} + // Sets kernel bundle to the provided one. Either replaces existing one or // create a new entry in the extended members vector. void handler::setHandlerKernelBundle( @@ -85,17 +142,45 @@ void handler::setHandlerKernelBundle( std::shared_ptr> ExendedMembersVec = detail::convertToExtendedMembers(MSharedPtrStorage[0]); - for (detail::ExtendedMemberT &EMember : *ExendedMembersVec) + // Find the handler kernel bundle flags if they have been set. Throw exception + // if the explicit kernel bundle flag is set. + bool KernelBundleFlagsExist = false; + for (detail::ExtendedMemberT &EMember : *ExendedMembersVec) { + if (detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE_FLAGS == + EMember.MType) { + auto Flags = std::static_pointer_cast(EMember.MData); + if (*Flags & detail::SPEC_CONST_SET_FLAG) + throw sycl::exception(make_error_code(errc::invalid), + "Kernel bundle cannot be explicitly set after a " + "specialization constant has been set"); + *Flags |= detail::EXPLICIT_KERNEL_BUNDLE_FLAG; + KernelBundleFlagsExist = true; + break; + } + } + + bool KernelBundleExist = false; + for (detail::ExtendedMemberT &EMember : *ExendedMembersVec) { if (detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE == EMember.MType) { EMember.MData = NewKernelBundleImpPtr; - return; + KernelBundleExist = true; + break; } + } - detail::ExtendedMemberT EMember = { - detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE, - NewKernelBundleImpPtr}; + if (!KernelBundleFlagsExist) { + detail::ExtendedMemberT EMember = { + detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE_FLAGS, + std::make_shared(detail::EXPLICIT_KERNEL_BUNDLE_FLAG)}; + ExendedMembersVec->push_back(EMember); + } - ExendedMembersVec->push_back(EMember); + if (!KernelBundleExist) { + detail::ExtendedMemberT EMember = { + detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE, + NewKernelBundleImpPtr}; + ExendedMembersVec->push_back(EMember); + } } event handler::finalize() { diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 7bbd5f580fa52..674fecdc62629 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -4264,6 +4264,7 @@ _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE4228EEENS3_12param_traitsIS4_XT _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE65552EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl7context9getNativeEv _ZNK2cl4sycl7handler30getOrInsertHandlerKernelBundleEb +_ZNK2cl4sycl7handler41getOrInsertNonExplicitHandlerKernelBundleEbb _ZNK2cl4sycl7program10get_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE _ZNK2cl4sycl7program10get_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEb _ZNK2cl4sycl7program10has_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE diff --git a/sycl/unittests/SYCL2020/CMakeLists.txt b/sycl/unittests/SYCL2020/CMakeLists.txt index 247adc4b7e11d..f4b64df26afb0 100644 --- a/sycl/unittests/SYCL2020/CMakeLists.txt +++ b/sycl/unittests/SYCL2020/CMakeLists.txt @@ -4,7 +4,7 @@ set(CMAKE_CXX_EXTENSIONS OFF) set(LLVM_REQUIRES_EH 1) add_sycl_unittest(SYCL2020Tests OBJECT GetNativeOpenCL.cpp - SpecConstDefaultValues.cpp + SpecializationConstant.cpp KernelBundle.cpp KernelID.cpp ) diff --git a/sycl/unittests/SYCL2020/SpecConstDefaultValues.cpp b/sycl/unittests/SYCL2020/SpecConstDefaultValues.cpp deleted file mode 100644 index 8bdf85b162c9d..0000000000000 --- a/sycl/unittests/SYCL2020/SpecConstDefaultValues.cpp +++ /dev/null @@ -1,171 +0,0 @@ -//==---- DefaultValues.cpp --- Spec constants default values unit test -----==// -// -// 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 -// -//===----------------------------------------------------------------------===// - -#define SYCL2020_DISABLE_DEPRECATION_WARNINGS - -#include -#include - -#include -#include -#include - -#include - -class TestKernel; -const static sycl::specialization_id SpecConst1{42}; - -__SYCL_INLINE_NAMESPACE(cl) { -namespace sycl { -namespace detail { -template <> struct KernelInfo { - static constexpr unsigned getNumParams() { return 0; } - static const kernel_param_desc_t &getParamDesc(int) { - static kernel_param_desc_t Dummy; - return Dummy; - } - static constexpr const char *getName() { - return "SpecConstDefaultValues_TestKernel"; - } - static constexpr bool isESIMD() { return false; } - static constexpr bool callsThisItem() { return false; } - static constexpr bool callsAnyThisFreeFunction() { return false; } -}; - -template <> const char *get_spec_constant_symbolic_ID() { - return "SC1"; -} -} // namespace detail -} // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) - -static sycl::unittest::PiImage generateImageWithSpecConsts() { - using namespace sycl::unittest; - - std::vector SpecConstData; - PiProperty SC1 = makeSpecConstant(SpecConstData, "SC1", {0}, {0}, {42}); - PiProperty SC2 = makeSpecConstant(SpecConstData, "SC2", {1}, {0}, {8}); - - PiPropertySet PropSet; - addSpecConstants({SC1, SC2}, std::move(SpecConstData), PropSet); - - std::vector Bin{0, 1, 2, 3, 4, 5}; // Random data - - PiArray Entries = - makeEmptyKernels({"SpecConstDefaultValues_TestKernel"}); - - PiImage Img{PI_DEVICE_BINARY_TYPE_SPIRV, // Format - __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64, // DeviceTargetSpec - "", // Compile options - "", // Link options - std::move(Bin), - std::move(Entries), - std::move(PropSet)}; - - return Img; -} - -static sycl::unittest::PiImage Img = generateImageWithSpecConsts(); -static sycl::unittest::PiImageArray<1> ImgArray{&Img}; - -TEST(SpecConstDefaultValues, DefaultValuesAreSet) { - sycl::platform Plt{sycl::default_selector()}; - if (Plt.is_host()) { - std::cerr << "Test is not supported on host, skipping\n"; - return; // test is not supported on host. - } - - if (Plt.get_backend() == sycl::backend::cuda) { - std::cerr << "Test is not supported on CUDA platform, skipping\n"; - return; - } - - if (Plt.get_backend() == sycl::backend::hip) { - std::cerr << "Test is not supported on HIP platform, skipping\n"; - return; - } - - sycl::unittest::PiMock Mock{Plt}; - setupDefaultMockAPIs(Mock); - - const sycl::device Dev = Plt.get_devices()[0]; - - sycl::queue Queue{Dev}; - - const sycl::context Ctx = Queue.get_context(); - - sycl::kernel_bundle KernelBundle = - sycl::get_kernel_bundle(Ctx, {Dev}); - - sycl::kernel_id TestKernelID = sycl::get_kernel_id(); - auto DevImage = - std::find_if(KernelBundle.begin(), KernelBundle.end(), - [&](auto Image) { return Image.has_kernel(TestKernelID); }); - EXPECT_NE(DevImage, KernelBundle.end()); - - auto DevImageImpl = sycl::detail::getSyclObjImpl(*DevImage); - const auto &Blob = DevImageImpl->get_spec_const_blob_ref(); - - int SpecConstVal1 = *reinterpret_cast(Blob.data()); - int SpecConstVal2 = *(reinterpret_cast(Blob.data()) + 1); - - EXPECT_EQ(SpecConstVal1, 42); - EXPECT_EQ(SpecConstVal2, 8); -} - -TEST(SpecConstDefaultValues, DefaultValuesAreOverriden) { - sycl::platform Plt{sycl::default_selector()}; - if (Plt.is_host()) { - std::cerr << "Test is not supported on host, skipping\n"; - return; // test is not supported on host. - } - - if (Plt.get_backend() == sycl::backend::cuda) { - std::cerr << "Test is not supported on CUDA platform, skipping\n"; - return; - } - - if (Plt.get_backend() == sycl::backend::hip) { - std::cerr << "Test is not supported on HIP platform, skipping\n"; - return; - } - - sycl::unittest::PiMock Mock{Plt}; - setupDefaultMockAPIs(Mock); - - const sycl::device Dev = Plt.get_devices()[0]; - - sycl::queue Queue{Dev}; - - const sycl::context Ctx = Queue.get_context(); - - sycl::kernel_bundle KernelBundle = - sycl::get_kernel_bundle(Ctx, {Dev}); - - sycl::kernel_id TestKernelID = sycl::get_kernel_id(); - auto DevImage = - std::find_if(KernelBundle.begin(), KernelBundle.end(), - [&](auto Image) { return Image.has_kernel(TestKernelID); }); - EXPECT_NE(DevImage, KernelBundle.end()); - - auto DevImageImpl = sycl::detail::getSyclObjImpl(*DevImage); - auto &Blob = DevImageImpl->get_spec_const_blob_ref(); - int SpecConstVal1 = *reinterpret_cast(Blob.data()); - int SpecConstVal2 = *(reinterpret_cast(Blob.data()) + 1); - - EXPECT_EQ(SpecConstVal1, 42); - EXPECT_EQ(SpecConstVal2, 8); - - KernelBundle.set_specialization_constant(80); - - SpecConstVal1 = *reinterpret_cast(Blob.data()); - SpecConstVal2 = *(reinterpret_cast(Blob.data()) + 1); - - EXPECT_EQ(SpecConstVal1, 80); - EXPECT_EQ(SpecConstVal2, 8); -} diff --git a/sycl/unittests/SYCL2020/SpecializationConstant.cpp b/sycl/unittests/SYCL2020/SpecializationConstant.cpp new file mode 100644 index 0000000000000..8d89113049aa5 --- /dev/null +++ b/sycl/unittests/SYCL2020/SpecializationConstant.cpp @@ -0,0 +1,351 @@ +//==------ SpecializationConstant.cpp --- Spec constants unit tests --------==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#define SYCL2020_DISABLE_DEPRECATION_WARNINGS + +#include +#include + +#include +#include +#include + +#include + +class TestKernel; +const static sycl::specialization_id SpecConst1{42}; + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace detail { +template <> struct KernelInfo { + static constexpr unsigned getNumParams() { return 0; } + static const kernel_param_desc_t &getParamDesc(int) { + static kernel_param_desc_t Dummy; + return Dummy; + } + static constexpr const char *getName() { + return "SpecializationConstant_TestKernel"; + } + static constexpr bool isESIMD() { return false; } + static constexpr bool callsThisItem() { return false; } + static constexpr bool callsAnyThisFreeFunction() { return false; } +}; + +template <> const char *get_spec_constant_symbolic_ID() { + return "SC1"; +} +} // namespace detail +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) + +static sycl::unittest::PiImage generateImageWithSpecConsts() { + using namespace sycl::unittest; + + std::vector SpecConstData; + PiProperty SC1 = makeSpecConstant(SpecConstData, "SC1", {0}, {0}, {42}); + PiProperty SC2 = makeSpecConstant(SpecConstData, "SC2", {1}, {0}, {8}); + + PiPropertySet PropSet; + addSpecConstants({SC1, SC2}, std::move(SpecConstData), PropSet); + + std::vector Bin{0, 1, 2, 3, 4, 5}; // Random data + + PiArray Entries = + makeEmptyKernels({"SpecializationConstant_TestKernel"}); + + PiImage Img{PI_DEVICE_BINARY_TYPE_SPIRV, // Format + __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64, // DeviceTargetSpec + "", // Compile options + "", // Link options + std::move(Bin), + std::move(Entries), + std::move(PropSet)}; + + return Img; +} + +static sycl::unittest::PiImage Img = generateImageWithSpecConsts(); +static sycl::unittest::PiImageArray<1> ImgArray{&Img}; + +TEST(SpecializationConstant, DefaultValuesAreSet) { + sycl::platform Plt{sycl::default_selector()}; + if (Plt.is_host()) { + std::cerr << "Test is not supported on host, skipping\n"; + return; // test is not supported on host. + } + + if (Plt.get_backend() == sycl::backend::cuda) { + std::cerr << "Test is not supported on CUDA platform, skipping\n"; + return; + } + + if (Plt.get_backend() == sycl::backend::hip) { + std::cerr << "Test is not supported on HIP platform, skipping\n"; + return; + } + + sycl::unittest::PiMock Mock{Plt}; + setupDefaultMockAPIs(Mock); + + const sycl::device Dev = Plt.get_devices()[0]; + + sycl::queue Queue{Dev}; + + const sycl::context Ctx = Queue.get_context(); + + sycl::kernel_bundle KernelBundle = + sycl::get_kernel_bundle(Ctx, {Dev}); + + sycl::kernel_id TestKernelID = sycl::get_kernel_id(); + auto DevImage = + std::find_if(KernelBundle.begin(), KernelBundle.end(), + [&](auto Image) { return Image.has_kernel(TestKernelID); }); + EXPECT_NE(DevImage, KernelBundle.end()); + + auto DevImageImpl = sycl::detail::getSyclObjImpl(*DevImage); + const auto &Blob = DevImageImpl->get_spec_const_blob_ref(); + + int SpecConstVal1 = *reinterpret_cast(Blob.data()); + int SpecConstVal2 = *(reinterpret_cast(Blob.data()) + 1); + + EXPECT_EQ(SpecConstVal1, 42); + EXPECT_EQ(SpecConstVal2, 8); +} + +TEST(SpecializationConstant, DefaultValuesAreOverriden) { + sycl::platform Plt{sycl::default_selector()}; + if (Plt.is_host()) { + std::cerr << "Test is not supported on host, skipping\n"; + return; // test is not supported on host. + } + + if (Plt.get_backend() == sycl::backend::cuda) { + std::cerr << "Test is not supported on CUDA platform, skipping\n"; + return; + } + + if (Plt.get_backend() == sycl::backend::hip) { + std::cerr << "Test is not supported on HIP platform, skipping\n"; + return; + } + + sycl::unittest::PiMock Mock{Plt}; + setupDefaultMockAPIs(Mock); + + const sycl::device Dev = Plt.get_devices()[0]; + + sycl::queue Queue{Dev}; + + const sycl::context Ctx = Queue.get_context(); + + sycl::kernel_bundle KernelBundle = + sycl::get_kernel_bundle(Ctx, {Dev}); + + sycl::kernel_id TestKernelID = sycl::get_kernel_id(); + auto DevImage = + std::find_if(KernelBundle.begin(), KernelBundle.end(), + [&](auto Image) { return Image.has_kernel(TestKernelID); }); + EXPECT_NE(DevImage, KernelBundle.end()); + + auto DevImageImpl = sycl::detail::getSyclObjImpl(*DevImage); + auto &Blob = DevImageImpl->get_spec_const_blob_ref(); + int SpecConstVal1 = *reinterpret_cast(Blob.data()); + int SpecConstVal2 = *(reinterpret_cast(Blob.data()) + 1); + + EXPECT_EQ(SpecConstVal1, 42); + EXPECT_EQ(SpecConstVal2, 8); + + KernelBundle.set_specialization_constant(80); + + SpecConstVal1 = *reinterpret_cast(Blob.data()); + SpecConstVal2 = *(reinterpret_cast(Blob.data()) + 1); + + EXPECT_EQ(SpecConstVal1, 80); + EXPECT_EQ(SpecConstVal2, 8); +} + +TEST(SpecializationConstant, SetSpecConstAfterUseKernelBundle) { + sycl::platform Plt{sycl::default_selector()}; + if (Plt.is_host()) { + std::cerr << "Test is not supported on host, skipping\n"; + return; // test is not supported on host. + } + + if (Plt.get_backend() == sycl::backend::cuda) { + std::cerr << "Test is not supported on CUDA platform, skipping\n"; + return; + } + + if (Plt.get_backend() == sycl::backend::hip) { + std::cerr << "Test is not supported on HIP platform, skipping\n"; + return; + } + + sycl::unittest::PiMock Mock{Plt}; + setupDefaultMockAPIs(Mock); + + const sycl::device Dev = Plt.get_devices()[0]; + + sycl::queue Queue{Dev}; + + const sycl::context Ctx = Queue.get_context(); + + sycl::kernel_bundle KernelBundle = + sycl::get_kernel_bundle(Ctx, {Dev}); + + // Create uniquely identifyable class to throw on expected exception + class UniqueException {}; + + try { + Queue.submit([&](sycl::handler &CGH) { + CGH.use_kernel_bundle(KernelBundle); + try { + CGH.set_specialization_constant(80); + FAIL() << "No exception was thrown."; + } catch (const sycl::exception &e) { + if (static_cast(e.code().value()) != sycl::errc::invalid) { + FAIL() << "Unexpected SYCL exception was thrown."; + throw; + } + throw UniqueException{}; + } catch (...) { + FAIL() << "Unexpected non-SYCL exception was thrown."; + throw; + } + CGH.single_task([]() {}); + }); + } catch (const sycl::exception &e) { + if (static_cast(e.code().value()) == sycl::errc::invalid) { + FAIL() << "SYCL exception with error code sycl::errc::invalid was " + "thrown at the wrong level."; + } + throw; + } catch (const UniqueException &) { + // Expected path + } +} + +TEST(SpecializationConstant, GetSpecConstAfterUseKernelBundle) { + sycl::platform Plt{sycl::default_selector()}; + if (Plt.is_host()) { + std::cerr << "Test is not supported on host, skipping\n"; + return; // test is not supported on host. + } + + if (Plt.get_backend() == sycl::backend::cuda) { + std::cerr << "Test is not supported on CUDA platform, skipping\n"; + return; + } + + if (Plt.get_backend() == sycl::backend::hip) { + std::cerr << "Test is not supported on HIP platform, skipping\n"; + return; + } + + sycl::unittest::PiMock Mock{Plt}; + setupDefaultMockAPIs(Mock); + + const sycl::device Dev = Plt.get_devices()[0]; + sycl::queue Queue{Dev}; + const sycl::context Ctx = Queue.get_context(); + + sycl::kernel_bundle KernelBundle = + sycl::get_kernel_bundle(Ctx, {Dev}); + + // Create uniquely identifyable class to throw on expected exception + class UniqueException {}; + + try { + Queue.submit([&](sycl::handler &CGH) { + CGH.use_kernel_bundle(KernelBundle); + try { + auto SpecConst1Val = CGH.get_specialization_constant(); + (void)SpecConst1Val; + FAIL() << "No exception was thrown."; + } catch (const sycl::exception &e) { + if (static_cast(e.code().value()) != sycl::errc::invalid) { + FAIL() << "Unexpected SYCL exception was thrown."; + throw; + } + throw UniqueException{}; + } catch (...) { + FAIL() << "Unexpected non-SYCL exception was thrown."; + throw; + } + CGH.single_task([]() {}); + }); + } catch (const sycl::exception &e) { + if (static_cast(e.code().value()) == sycl::errc::invalid) { + FAIL() << "SYCL exception with error code sycl::errc::invalid was " + "thrown at the wrong level."; + } + throw; + } catch (const UniqueException &) { + // Expected path + } +} + +TEST(SpecializationConstant, UseKernelBundleAfterSetSpecConst) { + sycl::platform Plt{sycl::default_selector()}; + if (Plt.is_host()) { + std::cerr << "Test is not supported on host, skipping\n"; + return; // test is not supported on host. + } + + if (Plt.get_backend() == sycl::backend::cuda) { + std::cerr << "Test is not supported on CUDA platform, skipping\n"; + return; + } + + if (Plt.get_backend() == sycl::backend::hip) { + std::cerr << "Test is not supported on HIP platform, skipping\n"; + return; + } + + sycl::unittest::PiMock Mock{Plt}; + setupDefaultMockAPIs(Mock); + + const sycl::device Dev = Plt.get_devices()[0]; + sycl::queue Queue{Dev}; + const sycl::context Ctx = Queue.get_context(); + + sycl::kernel_bundle KernelBundle = + sycl::get_kernel_bundle(Ctx, {Dev}); + + // Create uniquely identifyable class to throw on expected exception + class UniqueException {}; + + try { + Queue.submit([&](sycl::handler &CGH) { + CGH.set_specialization_constant(80); + try { + CGH.use_kernel_bundle(KernelBundle); + FAIL() << "No exception was thrown."; + } catch (const sycl::exception &e) { + if (static_cast(e.code().value()) != sycl::errc::invalid) { + FAIL() << "Unexpected SYCL exception was thrown."; + throw; + } + throw UniqueException{}; + } catch (...) { + FAIL() << "Unexpected non-SYCL exception was thrown."; + throw; + } + CGH.single_task([]() {}); + }); + } catch (const sycl::exception &e) { + if (static_cast(e.code().value()) == sycl::errc::invalid) { + FAIL() << "SYCL exception with error code sycl::errc::invalid was " + "thrown at the wrong level."; + } + throw; + } catch (const UniqueException &) { + // Expected path + } +} From 4a25092df32900aee46a6eb1394c6ea246e48b58 Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Tue, 28 Sep 2021 18:03:53 +0300 Subject: [PATCH 2/6] Introduce handler_impl and refactoring Signed-off-by: Steffen Larsen --- sycl/include/CL/sycl/detail/cg.hpp | 5 +- sycl/include/CL/sycl/handler.hpp | 37 +++-- sycl/source/detail/handler_impl.hpp | 36 +++++ sycl/source/handler.cpp | 188 +++++++++++++++----------- sycl/test/abi/sycl_symbols_linux.dump | 4 +- 5 files changed, 176 insertions(+), 94 deletions(-) create mode 100644 sycl/source/detail/handler_impl.hpp diff --git a/sycl/include/CL/sycl/detail/cg.hpp b/sycl/include/CL/sycl/detail/cg.hpp index 71c22d37334a7..170c0f39906c3 100644 --- a/sycl/include/CL/sycl/detail/cg.hpp +++ b/sycl/include/CL/sycl/detail/cg.hpp @@ -94,7 +94,10 @@ namespace detail { enum class ExtendedMembersType : unsigned int { HANDLER_KERNEL_BUNDLE = 0, HANDLER_MEM_ADVICE, - HANDLER_KERNEL_BUNDLE_FLAGS, + // handler_impl is stored in the exended members to avoid breaking ABI. + // TODO: This should be made a member of the handler class once ABI can be + // broken. + HANDLER_IMPL, }; // Holds a pointer to an object of an arbitrary type and an ID value which diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 3390bf925103a..5ecb27e36215f 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -234,12 +234,6 @@ class RoundedRangeKernelWithKH { KernelType KernelFunc; }; -// Kernel bundle flags used to identify previous uses of use_kernel_bundle and -// set_specialization_constant. These are used to detect invalid command group -// operation ordering. -constexpr std::uint8_t EXPLICIT_KERNEL_BUNDLE_FLAG = 1; -constexpr std::uint8_t SPEC_CONST_SET_FLAG = 2; - } // namespace detail namespace ext { @@ -1122,12 +1116,12 @@ class __SYCL_EXPORT handler { kernel_parallel_for_work_group(KernelFunc); } - std::shared_ptr - getOrInsertHandlerKernelBundle(bool Insert) const; + bool setStateExplicitKernel(); + bool setStateSpecConstSet(); + bool isStateExplicitKernel() const; std::shared_ptr - getOrInsertNonExplicitHandlerKernelBundle(bool Insert, - bool MarkSpecConstSet) const; + getOrInsertHandlerKernelBundle(bool Insert) const; void setHandlerKernelBundle( const std::shared_ptr &NewKernelBundleImpPtr); @@ -1160,9 +1154,13 @@ class __SYCL_EXPORT handler { void set_specialization_constant( typename std::remove_reference_t::value_type Value) { + if (!setStateSpecConstSet()) + throw sycl::exception(make_error_code(errc::invalid), + "Specialization constants cannot be set after " + "explicitly setting the used kernel bundle"); + std::shared_ptr KernelBundleImplPtr = - getOrInsertNonExplicitHandlerKernelBundle(/*Insert=*/true, - /*MarkSpecConstSet=*/true); + getOrInsertHandlerKernelBundle(/*Insert=*/true); detail::createSyclObjFromImpl>( KernelBundleImplPtr) @@ -1173,9 +1171,13 @@ class __SYCL_EXPORT handler { typename std::remove_reference_t::value_type get_specialization_constant() const { + if (isStateExplicitKernel()) + throw sycl::exception(make_error_code(errc::invalid), + "Specialization constants cannot be read after " + "explicitly setting the used kernel bundle"); + std::shared_ptr KernelBundleImplPtr = - getOrInsertNonExplicitHandlerKernelBundle(/*Insert=*/true, - /*MarkSpecConstSet=*/false); + getOrInsertHandlerKernelBundle(/*Insert=*/true); return detail::createSyclObjFromImpl>( KernelBundleImplPtr) @@ -1186,6 +1188,13 @@ class __SYCL_EXPORT handler { void use_kernel_bundle(const kernel_bundle &ExecBundle) { + + if (!setStateExplicitKernel()) + throw sycl::exception( + make_error_code(errc::invalid), + "Kernel bundle cannot be explicitly set after a specialization " + "constant has been set"); + setHandlerKernelBundle(detail::getSyclObjImpl(ExecBundle)); } diff --git a/sycl/source/detail/handler_impl.hpp b/sycl/source/detail/handler_impl.hpp new file mode 100644 index 0000000000000..55427ab1d6493 --- /dev/null +++ b/sycl/source/detail/handler_impl.hpp @@ -0,0 +1,36 @@ +//==---------------- handler_impl.hpp - SYCL handler -----------------------==// +// +// 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 + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace detail { + +using KernelBundleImplPtr = std::shared_ptr; + +enum class HandlerSubmissionState : std::uint8_t { + NO_STATE = 0, + EXPLICIT_KERNEL_BUNDLE_STATE, + SPEC_CONST_SET_STATE, +}; + +class handler_impl { +public: + handler_impl() = default; + + /// Registers mutually exclusive submission states. + HandlerSubmissionState MSubmissionState = + detail::HandlerSubmissionState::NO_STATE; +}; + +} // namespace detail +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index e25443dc417f9..4a79544a540c1 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -17,6 +17,7 @@ #include #include #include +#include #include #include #include @@ -27,21 +28,45 @@ namespace sycl { handler::handler(std::shared_ptr Queue, bool IsHost) : MQueue(std::move(Queue)), MIsHost(IsHost) { - MSharedPtrStorage.emplace_back( - std::make_shared>()); + // Create extended members and insert handler_impl + // TODO: When allowed to break ABI the handler_impl should be made a member + // of the handler class. + auto ExtendedMembers = + std::make_shared>(); + detail::ExtendedMemberT HandlerImplMember = { + detail::ExtendedMembersType::HANDLER_IMPL, + std::make_shared()}; + ExtendedMembers->push_back(std::move(HandlerImplMember)); + MSharedPtrStorage.push_back(std::move(ExtendedMembers)); +} + +/// Gets the handler_impl at the start of the extended members. +// detail::GlobalHandler::instance().getHandlerExtendedMembersMutex() must +// be held when calling this function. +std::shared_ptr +getHandlerImpl(const std::shared_ptr> + &ExtendedMembersVec) { + assert(ExtendedMembersVec->size() > 0); + + auto HandlerImplMember = (*ExtendedMembersVec)[0]; + + assert(detail::ExtendedMembersType::HANDLER_IMPL == HandlerImplMember.MType); + + return std::static_pointer_cast( + HandlerImplMember.MData); } // Common implementation for getting/inserting handler kernel bundle. -// detail::GlobalHandler::instance().getHandlerExtendedMembersMutex() must be -// held when calling this function. +// detail::GlobalHandler::instance().getHandlerExtendedMembersMutex() must +// be held when calling this function. std::shared_ptr getOrInsertHandlerKernelBundleCommon( const std::shared_ptr> - &ExendedMembersVec, + &ExtendedMembersVec, const std::shared_ptr &Queue, bool Insert) { // Look for the kernel bundle in extended members std::shared_ptr KernelBundleImpPtr; - for (const detail::ExtendedMemberT &EMember : *ExendedMembersVec) + for (const detail::ExtendedMemberT &EMember : *ExtendedMembersVec) if (detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE == EMember.MType) { KernelBundleImpPtr = std::static_pointer_cast(EMember.MData); @@ -59,20 +84,57 @@ getOrInsertHandlerKernelBundleCommon( detail::ExtendedMemberT EMember = { detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE, KernelBundleImpPtr}; - - ExendedMembersVec->push_back(EMember); + ExtendedMembersVec->push_back(EMember); } return KernelBundleImpPtr; } -// Returns a shared_ptr to kernel_bundle stored in the extended members vector. -// If there is no kernel_bundle created: -// returns newly created kernel_bundle if Insert is true -// returns shared_ptr(nullptr) if Insert is false -std::shared_ptr -handler::getOrInsertHandlerKernelBundle(bool Insert) const { +// If the submission state is SPEC_CONST_SET_STATE this function returns false. +// Otherwise it sets the submission state to EXPLICIT_KERNEL_BUNDLE_STATE and +// returns true. +bool handler::setStateExplicitKernel() { + std::lock_guard Lock( + detail::GlobalHandler::instance().getHandlerExtendedMembersMutex()); + + assert(!MSharedPtrStorage.empty()); + std::shared_ptr> ExendedMembersVec = + detail::convertToExtendedMembers(MSharedPtrStorage[0]); + + auto HandlerImpl = getHandlerImpl(ExendedMembersVec); + if (HandlerImpl->MSubmissionState == + detail::HandlerSubmissionState::SPEC_CONST_SET_STATE) + return false; + HandlerImpl->MSubmissionState = + detail::HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE; + return true; +} + +// If the submission state is EXPLICIT_KERNEL_BUNDLE_STATE this function returns +// false. Otherwise it sets the submission state to SPEC_CONST_SET_STATE and +// returns true. +bool handler::setStateSpecConstSet() { + std::lock_guard Lock( + detail::GlobalHandler::instance().getHandlerExtendedMembersMutex()); + + assert(!MSharedPtrStorage.empty()); + + std::shared_ptr> ExendedMembersVec = + detail::convertToExtendedMembers(MSharedPtrStorage[0]); + + auto HandlerImpl = getHandlerImpl(ExendedMembersVec); + if (HandlerImpl->MSubmissionState == + detail::HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE) + return false; + HandlerImpl->MSubmissionState = + detail::HandlerSubmissionState::SPEC_CONST_SET_STATE; + return true; +} + +// Returns true if the submission state is EXPLICIT_KERNEL_BUNDLE_STATE and +// false otherwise. +bool handler::isStateExplicitKernel() const { std::lock_guard Lock( detail::GlobalHandler::instance().getHandlerExtendedMembersMutex()); @@ -81,53 +143,49 @@ handler::getOrInsertHandlerKernelBundle(bool Insert) const { std::shared_ptr> ExendedMembersVec = detail::convertToExtendedMembers(MSharedPtrStorage[0]); - return getOrInsertHandlerKernelBundleCommon(ExendedMembersVec, MQueue, - Insert); + auto HandlerImpl = getHandlerImpl(ExendedMembersVec); + return HandlerImpl->MSubmissionState == + detail::HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE; } -// This function exhibits the same behavior as getOrInsertHandlerKernelBundle -// but throws an exception with errc::invalid if a kernel bundle has been set -// in the command group by a call to use_kernel_bundle. -// If MarkSpecConstSet is true the command group is flagged as having had set -// a specialization constant. +// Returns a shared_ptr to kernel_bundle stored in the extended members vector. +// If there is no kernel_bundle created: +// returns newly created kernel_bundle if Insert is true +// returns shared_ptr(nullptr) if Insert is false std::shared_ptr -handler::getOrInsertNonExplicitHandlerKernelBundle( - bool Insert, bool MarkSpecConstSet) const { +handler::getOrInsertHandlerKernelBundle(bool Insert) const { std::lock_guard Lock( detail::GlobalHandler::instance().getHandlerExtendedMembersMutex()); assert(!MSharedPtrStorage.empty()); - std::shared_ptr> ExendedMembersVec = + std::shared_ptr> ExtendedMembersVec = detail::convertToExtendedMembers(MSharedPtrStorage[0]); + // Look for the kernel bundle in extended members + std::shared_ptr KernelBundleImpPtr; + for (const detail::ExtendedMemberT &EMember : *ExtendedMembersVec) + if (detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE == EMember.MType) { + KernelBundleImpPtr = + std::static_pointer_cast(EMember.MData); + break; + } - // If kernel was explicitly set through use_kernel_bundle then throw exception - bool KernelBundleFlagsExist = false; - for (detail::ExtendedMemberT &EMember : *ExendedMembersVec) { - if (detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE_FLAGS == - EMember.MType) { - auto Flags = std::static_pointer_cast(EMember.MData); - if (*Flags & detail::EXPLICIT_KERNEL_BUNDLE_FLAG) - throw sycl::exception( - make_error_code(errc::invalid), - "Specialization constants cannot be accessed after explicitly " - "setting the used kernel bundle"); - if (MarkSpecConstSet) - *Flags |= detail::SPEC_CONST_SET_FLAG; - KernelBundleFlagsExist = true; + // No kernel bundle yet, create one + if (!KernelBundleImpPtr && Insert) { + KernelBundleImpPtr = detail::getSyclObjImpl( + get_kernel_bundle(MQueue->get_context())); + if (KernelBundleImpPtr->empty()) { + KernelBundleImpPtr = detail::getSyclObjImpl( + get_kernel_bundle(MQueue->get_context())); } - } - if (!KernelBundleFlagsExist && MarkSpecConstSet) { detail::ExtendedMemberT EMember = { - detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE_FLAGS, - std::make_shared(detail::SPEC_CONST_SET_FLAG)}; - ExendedMembersVec->push_back(EMember); + detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE, KernelBundleImpPtr}; + ExtendedMembersVec->push_back(EMember); } - return getOrInsertHandlerKernelBundleCommon(ExendedMembersVec, MQueue, - Insert); + return KernelBundleImpPtr; } // Sets kernel bundle to the provided one. Either replaces existing one or @@ -142,45 +200,19 @@ void handler::setHandlerKernelBundle( std::shared_ptr> ExendedMembersVec = detail::convertToExtendedMembers(MSharedPtrStorage[0]); - // Find the handler kernel bundle flags if they have been set. Throw exception - // if the explicit kernel bundle flag is set. - bool KernelBundleFlagsExist = false; - for (detail::ExtendedMemberT &EMember : *ExendedMembersVec) { - if (detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE_FLAGS == - EMember.MType) { - auto Flags = std::static_pointer_cast(EMember.MData); - if (*Flags & detail::SPEC_CONST_SET_FLAG) - throw sycl::exception(make_error_code(errc::invalid), - "Kernel bundle cannot be explicitly set after a " - "specialization constant has been set"); - *Flags |= detail::EXPLICIT_KERNEL_BUNDLE_FLAG; - KernelBundleFlagsExist = true; - break; - } - } - - bool KernelBundleExist = false; + // Look for kernel bundle in extended members and overwrite it. for (detail::ExtendedMemberT &EMember : *ExendedMembersVec) { if (detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE == EMember.MType) { EMember.MData = NewKernelBundleImpPtr; - KernelBundleExist = true; - break; + return; } } - if (!KernelBundleFlagsExist) { - detail::ExtendedMemberT EMember = { - detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE_FLAGS, - std::make_shared(detail::EXPLICIT_KERNEL_BUNDLE_FLAG)}; - ExendedMembersVec->push_back(EMember); - } - - if (!KernelBundleExist) { - detail::ExtendedMemberT EMember = { - detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE, - NewKernelBundleImpPtr}; - ExendedMembersVec->push_back(EMember); - } + // Kernel bundle was set found so we add it. + detail::ExtendedMemberT EMember = { + detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE, + NewKernelBundleImpPtr}; + ExendedMembersVec->push_back(EMember); } event handler::finalize() { diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 674fecdc62629..45b5cba6bda85 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3920,7 +3920,9 @@ _ZN2cl4sycl7handler18ext_oneapi_barrierERKSt6vectorINS0_5eventESaIS3_EE _ZN2cl4sycl7handler18extractArgsAndReqsEv _ZN2cl4sycl7handler20DisableRangeRoundingEv _ZN2cl4sycl7handler20associateWithHandlerEPNS0_6detail16AccessorBaseHostENS0_6access6targetE +_ZN2cl4sycl7handler20setStateSpecConstSetEv _ZN2cl4sycl7handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE +_ZN2cl4sycl7handler22setStateExplicitKernelEv _ZN2cl4sycl7handler24GetRangeRoundingSettingsERmS2_S2_ _ZN2cl4sycl7handler28extractArgsAndReqsFromLambdaEPcmPKNS0_6detail19kernel_param_desc_tE _ZN2cl4sycl7handler28extractArgsAndReqsFromLambdaEPcmPKNS0_6detail19kernel_param_desc_tEb @@ -4263,8 +4265,8 @@ _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE4225EEENS3_12param_traitsIS4_XT _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE4228EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE65552EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl7context9getNativeEv +_ZNK2cl4sycl7handler21isStateExplicitKernelEv _ZNK2cl4sycl7handler30getOrInsertHandlerKernelBundleEb -_ZNK2cl4sycl7handler41getOrInsertNonExplicitHandlerKernelBundleEbb _ZNK2cl4sycl7program10get_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE _ZNK2cl4sycl7program10get_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEb _ZNK2cl4sycl7program10has_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE From 6626bd31dd5ed257eb2e9aedb32324d50184cb7d Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Tue, 28 Sep 2021 18:22:08 +0300 Subject: [PATCH 3/6] Fix formatting Signed-off-by: Steffen Larsen --- sycl/source/handler.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 4a79544a540c1..a4b5e3557cc86 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -105,7 +105,7 @@ bool handler::setStateExplicitKernel() { auto HandlerImpl = getHandlerImpl(ExendedMembersVec); if (HandlerImpl->MSubmissionState == detail::HandlerSubmissionState::SPEC_CONST_SET_STATE) - return false; + return false; HandlerImpl->MSubmissionState = detail::HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE; return true; @@ -126,7 +126,7 @@ bool handler::setStateSpecConstSet() { auto HandlerImpl = getHandlerImpl(ExendedMembersVec); if (HandlerImpl->MSubmissionState == detail::HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE) - return false; + return false; HandlerImpl->MSubmissionState = detail::HandlerSubmissionState::SPEC_CONST_SET_STATE; return true; From 99c4a0ead54663056201ad858462c0e1a516195f Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Tue, 28 Sep 2021 20:15:33 +0300 Subject: [PATCH 4/6] Add new functions to Windows dump Signed-off-by: Steffen Larsen --- sycl/test/abi/sycl_symbols_windows.dump | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 9b347ceb7b47a..24bb99151c880 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -2331,6 +2331,7 @@ ?isInterop@SYCLMemObjT@detail@sycl@cl@@QEBA_NXZ ?isOutOfRange@detail@sycl@cl@@YA_NV?$vec@H$03@23@W4addressing_mode@23@V?$range@$02@23@@Z ?isPathPresent@OSUtil@detail@sycl@cl@@SA_NAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z +?isStateExplicitKernel@handler@sycl@cl@@AEBA_NXZ ?isValidModeForDestinationAccessor@handler@sycl@cl@@CA_NW4mode@access@23@@Z ?isValidModeForSourceAccessor@handler@sycl@cl@@CA_NW4mode@access@23@@Z ?isValidTargetForExplicitOp@handler@sycl@cl@@CA_NW4target@access@23@@Z @@ -3812,6 +3813,8 @@ ?setPitches@?$image_impl@$00@detail@sycl@cl@@AEAAXXZ ?setPitches@?$image_impl@$01@detail@sycl@cl@@AEAAXXZ ?setPitches@?$image_impl@$02@detail@sycl@cl@@AEAAXXZ +?setStateExplicitKernel@handler@sycl@cl@@AEAA_NXZ +?setStateSpecConstSet@handler@sycl@cl@@AEAA_NXZ ?setType@handler@sycl@cl@@AEAAXW4CGTYPE@CG@detail@23@@Z ?set_final_data@SYCLMemObjT@detail@sycl@cl@@QEAAX$$T@Z ?set_final_data_from_storage@SYCLMemObjT@detail@sycl@cl@@QEAAXXZ From 8391e73cafbc224060416eb416e0d66b4a639a6f Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Wed, 29 Sep 2021 12:39:05 +0300 Subject: [PATCH 5/6] Move logic and exception to handler_impl Signed-off-by: Steffen Larsen --- sycl/include/CL/sycl/handler.hpp | 24 ++---- sycl/source/detail/handler_impl.hpp | 26 +++++- sycl/source/handler.cpp | 110 +++++------------------- sycl/test/abi/sycl_symbols_linux.dump | 5 +- sycl/test/abi/sycl_symbols_windows.dump | 4 +- 5 files changed, 59 insertions(+), 110 deletions(-) diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 5ecb27e36215f..e64bc4d77eea8 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -80,6 +80,7 @@ template class buffer; namespace detail { +class handler_impl; class kernel_impl; class queue_impl; class stream_impl; @@ -1116,9 +1117,11 @@ class __SYCL_EXPORT handler { kernel_parallel_for_work_group(KernelFunc); } - bool setStateExplicitKernel(); - bool setStateSpecConstSet(); - bool isStateExplicitKernel() const; + std::shared_ptr getHandlerImpl() const; + + void setStateExplicitKernelBundle(); + void setStateSpecConstSet(); + bool isStateExplicitKernelBundle() const; std::shared_ptr getOrInsertHandlerKernelBundle(bool Insert) const; @@ -1154,10 +1157,7 @@ class __SYCL_EXPORT handler { void set_specialization_constant( typename std::remove_reference_t::value_type Value) { - if (!setStateSpecConstSet()) - throw sycl::exception(make_error_code(errc::invalid), - "Specialization constants cannot be set after " - "explicitly setting the used kernel bundle"); + setStateSpecConstSet(); std::shared_ptr KernelBundleImplPtr = getOrInsertHandlerKernelBundle(/*Insert=*/true); @@ -1171,7 +1171,7 @@ class __SYCL_EXPORT handler { typename std::remove_reference_t::value_type get_specialization_constant() const { - if (isStateExplicitKernel()) + if (isStateExplicitKernelBundle()) throw sycl::exception(make_error_code(errc::invalid), "Specialization constants cannot be read after " "explicitly setting the used kernel bundle"); @@ -1188,13 +1188,7 @@ class __SYCL_EXPORT handler { void use_kernel_bundle(const kernel_bundle &ExecBundle) { - - if (!setStateExplicitKernel()) - throw sycl::exception( - make_error_code(errc::invalid), - "Kernel bundle cannot be explicitly set after a specialization " - "constant has been set"); - + setStateExplicitKernelBundle(); setHandlerKernelBundle(detail::getSyclObjImpl(ExecBundle)); } diff --git a/sycl/source/detail/handler_impl.hpp b/sycl/source/detail/handler_impl.hpp index 55427ab1d6493..96f1621d28d34 100644 --- a/sycl/source/detail/handler_impl.hpp +++ b/sycl/source/detail/handler_impl.hpp @@ -26,9 +26,31 @@ class handler_impl { public: handler_impl() = default; + void setStateExplicitKernelBundle() { + if (MSubmissionState == HandlerSubmissionState::SPEC_CONST_SET_STATE) + throw sycl::exception( + make_error_code(errc::invalid), + "Kernel bundle cannot be explicitly set after a specialization " + "constant has been set"); + MSubmissionState = HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE; + } + + void setStateSpecConstSet() { + if (MSubmissionState == + HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE) + throw sycl::exception(make_error_code(errc::invalid), + "Specialization constants cannot be set after " + "explicitly setting the used kernel bundle"); + MSubmissionState = HandlerSubmissionState::SPEC_CONST_SET_STATE; + } + + bool isStateExplicitKernelBundle() const { + return MSubmissionState == + HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE; + } + /// Registers mutually exclusive submission states. - HandlerSubmissionState MSubmissionState = - detail::HandlerSubmissionState::NO_STATE; + HandlerSubmissionState MSubmissionState = HandlerSubmissionState::NO_STATE; }; } // namespace detail diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index a4b5e3557cc86..b311146a62917 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -41,11 +41,15 @@ handler::handler(std::shared_ptr Queue, bool IsHost) } /// Gets the handler_impl at the start of the extended members. -// detail::GlobalHandler::instance().getHandlerExtendedMembersMutex() must -// be held when calling this function. -std::shared_ptr -getHandlerImpl(const std::shared_ptr> - &ExtendedMembersVec) { +std::shared_ptr handler::getHandlerImpl() const { + std::lock_guard Lock( + detail::GlobalHandler::instance().getHandlerExtendedMembersMutex()); + + assert(!MSharedPtrStorage.empty()); + + std::shared_ptr> ExtendedMembersVec = + detail::convertToExtendedMembers(MSharedPtrStorage[0]); + assert(ExtendedMembersVec->size() > 0); auto HandlerImplMember = (*ExtendedMembersVec)[0]; @@ -56,96 +60,24 @@ getHandlerImpl(const std::shared_ptr> HandlerImplMember.MData); } -// Common implementation for getting/inserting handler kernel bundle. -// detail::GlobalHandler::instance().getHandlerExtendedMembersMutex() must -// be held when calling this function. -std::shared_ptr -getOrInsertHandlerKernelBundleCommon( - const std::shared_ptr> - &ExtendedMembersVec, - const std::shared_ptr &Queue, bool Insert) { - // Look for the kernel bundle in extended members - std::shared_ptr KernelBundleImpPtr; - for (const detail::ExtendedMemberT &EMember : *ExtendedMembersVec) - if (detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE == EMember.MType) { - KernelBundleImpPtr = - std::static_pointer_cast(EMember.MData); - break; - } - - // No kernel bundle yet, create one - if (!KernelBundleImpPtr && Insert) { - KernelBundleImpPtr = detail::getSyclObjImpl( - get_kernel_bundle(Queue->get_context())); - if (KernelBundleImpPtr->empty()) { - KernelBundleImpPtr = detail::getSyclObjImpl( - get_kernel_bundle(Queue->get_context())); - } - - detail::ExtendedMemberT EMember = { - detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE, KernelBundleImpPtr}; - ExtendedMembersVec->push_back(EMember); - } - - return KernelBundleImpPtr; -} - -// If the submission state is SPEC_CONST_SET_STATE this function returns false. -// Otherwise it sets the submission state to EXPLICIT_KERNEL_BUNDLE_STATE and -// returns true. -bool handler::setStateExplicitKernel() { - std::lock_guard Lock( - detail::GlobalHandler::instance().getHandlerExtendedMembersMutex()); - - assert(!MSharedPtrStorage.empty()); - - std::shared_ptr> ExendedMembersVec = - detail::convertToExtendedMembers(MSharedPtrStorage[0]); - - auto HandlerImpl = getHandlerImpl(ExendedMembersVec); - if (HandlerImpl->MSubmissionState == - detail::HandlerSubmissionState::SPEC_CONST_SET_STATE) - return false; - HandlerImpl->MSubmissionState = - detail::HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE; - return true; +// Sets the submission state to indicate that an explicit kernel bundle has been +// set. This returns a sycl::exception with errc::invalid if the current state +// indicates that a specialization constant has been set. +void handler::setStateExplicitKernelBundle() { + getHandlerImpl()->setStateExplicitKernelBundle(); } -// If the submission state is EXPLICIT_KERNEL_BUNDLE_STATE this function returns -// false. Otherwise it sets the submission state to SPEC_CONST_SET_STATE and -// returns true. -bool handler::setStateSpecConstSet() { - std::lock_guard Lock( - detail::GlobalHandler::instance().getHandlerExtendedMembersMutex()); - - assert(!MSharedPtrStorage.empty()); - - std::shared_ptr> ExendedMembersVec = - detail::convertToExtendedMembers(MSharedPtrStorage[0]); - - auto HandlerImpl = getHandlerImpl(ExendedMembersVec); - if (HandlerImpl->MSubmissionState == - detail::HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE) - return false; - HandlerImpl->MSubmissionState = - detail::HandlerSubmissionState::SPEC_CONST_SET_STATE; - return true; +// Sets the submission state to indicate that a specialization constant has been +// set. This returns a sycl::exception with errc::invalid if the current state +// indicates that an explicit kernel bundle has been set. +void handler::setStateSpecConstSet() { + getHandlerImpl()->setStateSpecConstSet(); } // Returns true if the submission state is EXPLICIT_KERNEL_BUNDLE_STATE and // false otherwise. -bool handler::isStateExplicitKernel() const { - std::lock_guard Lock( - detail::GlobalHandler::instance().getHandlerExtendedMembersMutex()); - - assert(!MSharedPtrStorage.empty()); - - std::shared_ptr> ExendedMembersVec = - detail::convertToExtendedMembers(MSharedPtrStorage[0]); - - auto HandlerImpl = getHandlerImpl(ExendedMembersVec); - return HandlerImpl->MSubmissionState == - detail::HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE; +bool handler::isStateExplicitKernelBundle() const { + return getHandlerImpl()->isStateExplicitKernelBundle(); } // Returns a shared_ptr to kernel_bundle stored in the extended members vector. diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 45b5cba6bda85..3e891d3977e2d 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3922,7 +3922,7 @@ _ZN2cl4sycl7handler20DisableRangeRoundingEv _ZN2cl4sycl7handler20associateWithHandlerEPNS0_6detail16AccessorBaseHostENS0_6access6targetE _ZN2cl4sycl7handler20setStateSpecConstSetEv _ZN2cl4sycl7handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE -_ZN2cl4sycl7handler22setStateExplicitKernelEv +_ZN2cl4sycl7handler28setStateExplicitKernelBundleEv _ZN2cl4sycl7handler24GetRangeRoundingSettingsERmS2_S2_ _ZN2cl4sycl7handler28extractArgsAndReqsFromLambdaEPcmPKNS0_6detail19kernel_param_desc_tE _ZN2cl4sycl7handler28extractArgsAndReqsFromLambdaEPcmPKNS0_6detail19kernel_param_desc_tEb @@ -4265,7 +4265,8 @@ _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE4225EEENS3_12param_traitsIS4_XT _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE4228EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE65552EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl7context9getNativeEv -_ZNK2cl4sycl7handler21isStateExplicitKernelEv +_ZNK2cl4sycl7handler14getHandlerImplEv +_ZNK2cl4sycl7handler27isStateExplicitKernelBundleEv _ZNK2cl4sycl7handler30getOrInsertHandlerKernelBundleEb _ZNK2cl4sycl7program10get_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE _ZNK2cl4sycl7program10get_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEb diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 24bb99151c880..88a3c69efbffe 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -2331,7 +2331,7 @@ ?isInterop@SYCLMemObjT@detail@sycl@cl@@QEBA_NXZ ?isOutOfRange@detail@sycl@cl@@YA_NV?$vec@H$03@23@W4addressing_mode@23@V?$range@$02@23@@Z ?isPathPresent@OSUtil@detail@sycl@cl@@SA_NAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z -?isStateExplicitKernel@handler@sycl@cl@@AEBA_NXZ +?isStateExplicitKernelBundle@handler@sycl@cl@@AEBA_NXZ ?isValidModeForDestinationAccessor@handler@sycl@cl@@CA_NW4mode@access@23@@Z ?isValidModeForSourceAccessor@handler@sycl@cl@@CA_NW4mode@access@23@@Z ?isValidTargetForExplicitOp@handler@sycl@cl@@CA_NW4target@access@23@@Z @@ -3813,7 +3813,7 @@ ?setPitches@?$image_impl@$00@detail@sycl@cl@@AEAAXXZ ?setPitches@?$image_impl@$01@detail@sycl@cl@@AEAAXXZ ?setPitches@?$image_impl@$02@detail@sycl@cl@@AEAAXXZ -?setStateExplicitKernel@handler@sycl@cl@@AEAA_NXZ +?setStateExplicitKernelBundle@handler@sycl@cl@@AEAA_NXZ ?setStateSpecConstSet@handler@sycl@cl@@AEAA_NXZ ?setType@handler@sycl@cl@@AEAAXW4CGTYPE@CG@detail@23@@Z ?set_final_data@SYCLMemObjT@detail@sycl@cl@@QEAAX$$T@Z From 16880f2284575fbad18907e9c35b181bbcf93052 Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Wed, 29 Sep 2021 14:01:44 +0300 Subject: [PATCH 6/6] Change comments and fix windows dump Signed-off-by: Steffen Larsen --- sycl/source/handler.cpp | 4 ++-- sycl/test/abi/sycl_symbols_windows.dump | 5 +++-- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index b311146a62917..8f12abaee8a1f 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -61,14 +61,14 @@ std::shared_ptr handler::getHandlerImpl() const { } // Sets the submission state to indicate that an explicit kernel bundle has been -// set. This returns a sycl::exception with errc::invalid if the current state +// set. Throws a sycl::exception with errc::invalid if the current state // indicates that a specialization constant has been set. void handler::setStateExplicitKernelBundle() { getHandlerImpl()->setStateExplicitKernelBundle(); } // Sets the submission state to indicate that a specialization constant has been -// set. This returns a sycl::exception with errc::invalid if the current state +// set. Throws a sycl::exception with errc::invalid if the current state // indicates that an explicit kernel bundle has been set. void handler::setStateSpecConstSet() { getHandlerImpl()->setStateSpecConstSet(); diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 88a3c69efbffe..0fab219338d04 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -2056,6 +2056,7 @@ ?getElementSize@?$image_impl@$01@detail@sycl@cl@@QEBA_KXZ ?getElementSize@?$image_impl@$02@detail@sycl@cl@@QEBA_KXZ ?getEndTime@HostProfilingInfo@detail@sycl@cl@@QEBA_KXZ +?getHandlerImpl@handler@sycl@cl@@AEBA?AV?$shared_ptr@Vhandler_impl@detail@sycl@cl@@@std@@XZ ?getImageDesc@?$image_impl@$00@detail@sycl@cl@@AEAA?AU_pi_image_desc@@_N@Z ?getImageDesc@?$image_impl@$01@detail@sycl@cl@@AEAA?AU_pi_image_desc@@_N@Z ?getImageDesc@?$image_impl@$02@detail@sycl@cl@@AEAA?AU_pi_image_desc@@_N@Z @@ -3813,8 +3814,8 @@ ?setPitches@?$image_impl@$00@detail@sycl@cl@@AEAAXXZ ?setPitches@?$image_impl@$01@detail@sycl@cl@@AEAAXXZ ?setPitches@?$image_impl@$02@detail@sycl@cl@@AEAAXXZ -?setStateExplicitKernelBundle@handler@sycl@cl@@AEAA_NXZ -?setStateSpecConstSet@handler@sycl@cl@@AEAA_NXZ +?setStateExplicitKernelBundle@handler@sycl@cl@@AEAAXXZ +?setStateSpecConstSet@handler@sycl@cl@@AEAAXXZ ?setType@handler@sycl@cl@@AEAAXW4CGTYPE@CG@detail@23@@Z ?set_final_data@SYCLMemObjT@detail@sycl@cl@@QEAAX$$T@Z ?set_final_data_from_storage@SYCLMemObjT@detail@sycl@cl@@QEAAXXZ