From 9659b8895e8de00d911a849f68f86a5be9d6ece4 Mon Sep 17 00:00:00 2001 From: Vlad Romanov Date: Mon, 31 May 2021 20:23:00 +0300 Subject: [PATCH 1/6] [SYCL] Add sycl::kernel::get_kernel_bundle method --- sycl/include/CL/sycl/kernel.hpp | 7 +++ sycl/include/CL/sycl/kernel_bundle.hpp | 3 +- sycl/include/CL/sycl/kernel_bundle_enums.hpp | 17 +++++++ sycl/source/detail/kernel_impl.hpp | 2 + sycl/source/kernel.cpp | 7 +++ sycl/test/abi/sycl_symbols_linux.dump | 1 + sycl/unittests/CMakeLists.txt | 2 +- .../CMakeLists.txt | 4 +- .../SYCL2020APITests.cpp} | 44 ++++++++++++++++++- 9 files changed, 80 insertions(+), 7 deletions(-) create mode 100644 sycl/include/CL/sycl/kernel_bundle_enums.hpp rename sycl/unittests/{spec_constants => SYCL2020API}/CMakeLists.txt (62%) rename sycl/unittests/{spec_constants/DefaultValues.cpp => SYCL2020API/SYCL2020APITests.cpp} (87%) diff --git a/sycl/include/CL/sycl/kernel.hpp b/sycl/include/CL/sycl/kernel.hpp index 12ad8cd1ecb31..7e19b2c669122 100644 --- a/sycl/include/CL/sycl/kernel.hpp +++ b/sycl/include/CL/sycl/kernel.hpp @@ -12,6 +12,7 @@ #include #include #include +#include #include #include @@ -22,6 +23,7 @@ namespace sycl { class program; class context; template class backend_traits; +template class kernel_bundle; namespace detail { class kernel_impl; @@ -100,6 +102,11 @@ class __SYCL_EXPORT kernel { /// \return a valid SYCL context context get_context() const; + /// Get the kernel_bundle that this kernel is associated with. + /// + /// \return a valid kernel_bundle + kernel_bundle get_kernel_bundle() const; + /// Get the program that this kernel is defined for. /// /// The value returned must be equal to that returned by diff --git a/sycl/include/CL/sycl/kernel_bundle.hpp b/sycl/include/CL/sycl/kernel_bundle.hpp index 9eb256b63d2ab..4669b2163d45c 100644 --- a/sycl/include/CL/sycl/kernel_bundle.hpp +++ b/sycl/include/CL/sycl/kernel_bundle.hpp @@ -15,6 +15,7 @@ #include #include #include +#include #include #include @@ -25,8 +26,6 @@ namespace sycl { // Forward declaration template class backend_traits; -enum class bundle_state : char { input = 0, object = 1, executable = 2 }; - namespace detail { class kernel_id_impl; } diff --git a/sycl/include/CL/sycl/kernel_bundle_enums.hpp b/sycl/include/CL/sycl/kernel_bundle_enums.hpp new file mode 100644 index 0000000000000..89bb0be1ca6c7 --- /dev/null +++ b/sycl/include/CL/sycl/kernel_bundle_enums.hpp @@ -0,0 +1,17 @@ +//==------- kernel_bundle_enums.hpp - SYCL kernel_bundle related enums -----==// +// +// 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 + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { + +enum class bundle_state : char { input = 0, object = 1, executable = 2 }; + +} +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp index a7c66e8f0e69b..f4b734d37ebe8 100644 --- a/sycl/source/detail/kernel_impl.hpp +++ b/sycl/source/detail/kernel_impl.hpp @@ -198,6 +198,8 @@ class kernel_impl { return NativeKernel; } + KernelBundleImplPtr get_kernel_bundle() const { return MKernelBundleImpl; } + private: RT::PiKernel MKernel; const ContextImplPtr MContext; diff --git a/sycl/source/kernel.cpp b/sycl/source/kernel.cpp index 79ab595a1b845..f36fbe3e7b0b2 100644 --- a/sycl/source/kernel.cpp +++ b/sycl/source/kernel.cpp @@ -10,6 +10,7 @@ #include #include #include +#include #include __SYCL_INLINE_NAMESPACE(cl) { @@ -28,6 +29,12 @@ context kernel::get_context() const { return impl->get_info(); } +kernel_bundle +kernel::get_kernel_bundle() const { + return detail::createSyclObjFromImpl< + kernel_bundle>(impl->get_kernel_bundle()); +} + program kernel::get_program() const { return impl->get_info(); } diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 011113631ec29..b89f8137df7d7 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -4100,6 +4100,7 @@ _ZNK2cl4sycl6device9getNativeEv _ZNK2cl4sycl6kernel11get_contextEv _ZNK2cl4sycl6kernel11get_programEv _ZNK2cl4sycl6kernel13getNativeImplEv +_ZNK2cl4sycl6kernel17get_kernel_bundleEv _ZNK2cl4sycl6kernel18get_sub_group_infoILNS0_4info16kernel_sub_groupE16650EEENS3_12param_traitsIS4_XT_EE11return_typeERKNS0_6deviceE _ZNK2cl4sycl6kernel18get_sub_group_infoILNS0_4info16kernel_sub_groupE4537EEENS3_12param_traitsIS4_XT_EE11return_typeERKNS0_6deviceE _ZNK2cl4sycl6kernel18get_sub_group_infoILNS0_4info16kernel_sub_groupE4538EEENS3_12param_traitsIS4_XT_EE11return_typeERKNS0_6deviceE diff --git a/sycl/unittests/CMakeLists.txt b/sycl/unittests/CMakeLists.txt index 5da1955d9e5ff..cbfe46f354aeb 100644 --- a/sycl/unittests/CMakeLists.txt +++ b/sycl/unittests/CMakeLists.txt @@ -15,5 +15,5 @@ add_subdirectory(pi) add_subdirectory(kernel-and-program) add_subdirectory(queue) add_subdirectory(scheduler) -add_subdirectory(spec_constants) +add_subdirectory(SYCL2020API) add_subdirectory(thread_safety) diff --git a/sycl/unittests/spec_constants/CMakeLists.txt b/sycl/unittests/SYCL2020API/CMakeLists.txt similarity index 62% rename from sycl/unittests/spec_constants/CMakeLists.txt rename to sycl/unittests/SYCL2020API/CMakeLists.txt index cff537cd2963a..93981d20c36bf 100644 --- a/sycl/unittests/spec_constants/CMakeLists.txt +++ b/sycl/unittests/SYCL2020API/CMakeLists.txt @@ -2,7 +2,7 @@ set(CMAKE_CXX_EXTENSIONS OFF) # Enable exception handling for these unit tests set(LLVM_REQUIRES_EH 1) -add_sycl_unittest(SpecConstantsTests OBJECT - DefaultValues.cpp +add_sycl_unittest(SYCL2020APITests OBJECT + SYCL2020APITests.cpp ) diff --git a/sycl/unittests/spec_constants/DefaultValues.cpp b/sycl/unittests/SYCL2020API/SYCL2020APITests.cpp similarity index 87% rename from sycl/unittests/spec_constants/DefaultValues.cpp rename to sycl/unittests/SYCL2020API/SYCL2020APITests.cpp index a5605124a533a..7090fe573f198 100644 --- a/sycl/unittests/spec_constants/DefaultValues.cpp +++ b/sycl/unittests/SYCL2020API/SYCL2020APITests.cpp @@ -205,7 +205,7 @@ static sycl::unittest::PiImage generateDefaultImage() { sycl::unittest::PiImage Img = generateDefaultImage(); sycl::unittest::PiImageArray ImgArray{Img}; -TEST(DefaultValues, DISABLED_DefaultValuesAreSet) { +TEST(SpecConstants, DISABLED_DefaultValuesAreSet) { sycl::platform Plt{sycl::default_selector()}; if (Plt.is_host()) { std::cerr << "Test is not supported on host, skipping\n"; @@ -238,7 +238,7 @@ TEST(DefaultValues, DISABLED_DefaultValuesAreSet) { EXPECT_EQ(SpecConstVal1, 8); } -TEST(DefaultValues, DISABLED_DefaultValuesAreOverriden) { +TEST(SpecConstants, DISABLED_DefaultValuesAreOverriden) { sycl::platform Plt{sycl::default_selector()}; if (Plt.is_host()) { std::cerr << "Test is not supported on host, skipping\n"; @@ -271,3 +271,43 @@ TEST(DefaultValues, DISABLED_DefaultValuesAreOverriden) { EXPECT_EQ(SpecConstVal0, 80); EXPECT_EQ(SpecConstVal1, 8); } + +TEST(KernelBundle, GetKernelBundleFromKernel) { + 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; + } + + 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 Kernel = + KernelBundle.get_kernel(sycl::get_kernel_id()); + + sycl::kernel_bundle RetKernelBundle = + Kernel.get_kernel_bundle(); + + EXPECT_EQ(KernelBundle, RetKernelBundle); + + //KernelBundle.set_specialization_constant(80); + //auto ExecBundle = sycl::build(KernelBundle); + //if (0) + //Queue.submit([&](sycl::handler &CGH) { + //CGH.single_task([] {}); // Actual kernel does not matter + //}); +} From 70012263df4b99021e6e1b8693bb8474e0cfb58d Mon Sep 17 00:00:00 2001 From: Vlad Romanov Date: Tue, 1 Jun 2021 11:12:40 +0300 Subject: [PATCH 2/6] upply comments --- sycl/unittests/SYCL2020API/SYCL2020APITests.cpp | 7 ------- 1 file changed, 7 deletions(-) diff --git a/sycl/unittests/SYCL2020API/SYCL2020APITests.cpp b/sycl/unittests/SYCL2020API/SYCL2020APITests.cpp index 7090fe573f198..a1cc8e84958b1 100644 --- a/sycl/unittests/SYCL2020API/SYCL2020APITests.cpp +++ b/sycl/unittests/SYCL2020API/SYCL2020APITests.cpp @@ -303,11 +303,4 @@ TEST(KernelBundle, GetKernelBundleFromKernel) { Kernel.get_kernel_bundle(); EXPECT_EQ(KernelBundle, RetKernelBundle); - - //KernelBundle.set_specialization_constant(80); - //auto ExecBundle = sycl::build(KernelBundle); - //if (0) - //Queue.submit([&](sycl::handler &CGH) { - //CGH.single_task([] {}); // Actual kernel does not matter - //}); } From 0409c2d91ae31f47c16996766d737110dcb559a1 Mon Sep 17 00:00:00 2001 From: Vlad Romanov Date: Thu, 1 Jul 2021 13:41:09 +0300 Subject: [PATCH 3/6] Finish merge --- sycl/unittests/SYCL2020/CMakeLists.txt | 1 + sycl/unittests/SYCL2020/KernelBundle.cpp | 94 +++++++++++++++++++ .../SYCL2020/SpecConstDefaultValues.cpp | 4 +- 3 files changed, 97 insertions(+), 2 deletions(-) create mode 100644 sycl/unittests/SYCL2020/KernelBundle.cpp diff --git a/sycl/unittests/SYCL2020/CMakeLists.txt b/sycl/unittests/SYCL2020/CMakeLists.txt index c284fabd5f09f..7d8be599bfb5e 100644 --- a/sycl/unittests/SYCL2020/CMakeLists.txt +++ b/sycl/unittests/SYCL2020/CMakeLists.txt @@ -5,5 +5,6 @@ set(LLVM_REQUIRES_EH 1) add_sycl_unittest(SYCL2020Tests OBJECT GetNativeOpenCL.cpp SpecConstDefaultValues.cpp + KernelBundle.cpp ) diff --git a/sycl/unittests/SYCL2020/KernelBundle.cpp b/sycl/unittests/SYCL2020/KernelBundle.cpp new file mode 100644 index 0000000000000..d620f8a851e5b --- /dev/null +++ b/sycl/unittests/SYCL2020/KernelBundle.cpp @@ -0,0 +1,94 @@ +//==---- 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 + +class TestKernel; + +__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 "TestKernel"; } + static constexpr bool isESIMD() { return false; } + static constexpr bool callsThisItem() { return false; } + static constexpr bool callsAnyThisFreeFunction() { return false; } +}; + +} // namespace detail +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) + +static sycl::unittest::PiImage generateDefaultImage() { + using namespace sycl::unittest; + + PiPropertySet PropSet; + + std::vector Bin{0, 1, 2, 3, 4, 5}; // Random data + + PiArray Entries = makeEmptyKernels({"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 = generateDefaultImage(); +static sycl::unittest::PiImageArray ImgArray{Img}; + +TEST(KernelBundle, GetKernelBundleFromKernel) { + 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; + } + + 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 Kernel = + KernelBundle.get_kernel(sycl::get_kernel_id()); + + sycl::kernel_bundle RetKernelBundle = + Kernel.get_kernel_bundle(); + + EXPECT_EQ(KernelBundle, RetKernelBundle); +} diff --git a/sycl/unittests/SYCL2020/SpecConstDefaultValues.cpp b/sycl/unittests/SYCL2020/SpecConstDefaultValues.cpp index 83c18e1ec72b1..7707096411d0b 100644 --- a/sycl/unittests/SYCL2020/SpecConstDefaultValues.cpp +++ b/sycl/unittests/SYCL2020/SpecConstDefaultValues.cpp @@ -81,8 +81,8 @@ static sycl::unittest::PiImage generateImageWithSpecConsts() { return Img; } -sycl::unittest::PiImage Img = generateImageWithSpecConsts(); -sycl::unittest::PiImageArray ImgArray{Img}; +static sycl::unittest::PiImage Img = generateImageWithSpecConsts(); +static sycl::unittest::PiImageArray ImgArray{Img}; TEST(SpecConstDefaultValues, DISABLED_DefaultValuesAreSet) { sycl::platform Plt{sycl::default_selector()}; From 9f37b630e8b88ce2dfbe2e652ad33f6cf3ff966c Mon Sep 17 00:00:00 2001 From: Vlad Romanov Date: Thu, 1 Jul 2021 13:47:54 +0300 Subject: [PATCH 4/6] Address comments --- sycl/include/CL/sycl/kernel.hpp | 2 +- sycl/unittests/SYCL2020/KernelBundle.cpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/include/CL/sycl/kernel.hpp b/sycl/include/CL/sycl/kernel.hpp index 5c1df9814b669..acc607fbaa9e0 100644 --- a/sycl/include/CL/sycl/kernel.hpp +++ b/sycl/include/CL/sycl/kernel.hpp @@ -102,7 +102,7 @@ class __SYCL_EXPORT kernel { /// \return a valid SYCL context context get_context() const; - /// Get the kernel_bundle that this kernel is associated with. + /// Get the kernel_bundle associated with this kernel. /// /// \return a valid kernel_bundle kernel_bundle get_kernel_bundle() const; diff --git a/sycl/unittests/SYCL2020/KernelBundle.cpp b/sycl/unittests/SYCL2020/KernelBundle.cpp index d620f8a851e5b..ea22dd8fd0456 100644 --- a/sycl/unittests/SYCL2020/KernelBundle.cpp +++ b/sycl/unittests/SYCL2020/KernelBundle.cpp @@ -63,12 +63,12 @@ static sycl::unittest::PiImageArray ImgArray{Img}; TEST(KernelBundle, GetKernelBundleFromKernel) { sycl::platform Plt{sycl::default_selector()}; if (Plt.is_host()) { - std::cerr << "Test is not supported on host, skipping\n"; + std::cout << "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"; + std::cout << "Test is not supported on CUDA platform, skipping\n"; return; } From a1bbab22124cf61c702ba479f32fb4a6fd3baf7c Mon Sep 17 00:00:00 2001 From: Vlad Romanov Date: Thu, 1 Jul 2021 19:12:55 +0300 Subject: [PATCH 5/6] Remove deprecation warning suppression macro --- sycl/unittests/SYCL2020/KernelBundle.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/sycl/unittests/SYCL2020/KernelBundle.cpp b/sycl/unittests/SYCL2020/KernelBundle.cpp index ea22dd8fd0456..83e57a85b5ccd 100644 --- a/sycl/unittests/SYCL2020/KernelBundle.cpp +++ b/sycl/unittests/SYCL2020/KernelBundle.cpp @@ -6,8 +6,6 @@ // //===----------------------------------------------------------------------===// -#define SYCL2020_DISABLE_DEPRECATION_WARNINGS - #include #include From 9e0f78c79aadeab1ff9d769ed2de7a6486675908 Mon Sep 17 00:00:00 2001 From: Vlad Romanov Date: Thu, 1 Jul 2021 19:57:14 +0300 Subject: [PATCH 6/6] Adapt to new reality --- sycl/unittests/SYCL2020/KernelBundle.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/unittests/SYCL2020/KernelBundle.cpp b/sycl/unittests/SYCL2020/KernelBundle.cpp index 83e57a85b5ccd..c8c2edb48f0ba 100644 --- a/sycl/unittests/SYCL2020/KernelBundle.cpp +++ b/sycl/unittests/SYCL2020/KernelBundle.cpp @@ -56,7 +56,7 @@ static sycl::unittest::PiImage generateDefaultImage() { } static sycl::unittest::PiImage Img = generateDefaultImage(); -static sycl::unittest::PiImageArray ImgArray{Img}; +static sycl::unittest::PiImageArray<1> ImgArray{&Img}; TEST(KernelBundle, GetKernelBundleFromKernel) { sycl::platform Plt{sycl::default_selector()};