From 0e2c6fa040607ca9050eb22e9660b0a897487b10 Mon Sep 17 00:00:00 2001 From: Denis Kabanov Date: Thu, 9 Dec 2021 17:27:22 +0300 Subject: [PATCH 1/6] [SYCL] Fix kernel bundles don't really carry kernel IDs --- sycl/include/CL/sycl/kernel_bundle.hpp | 15 +++++++++++++++ sycl/source/backend.cpp | 8 ++++---- sycl/source/detail/kernel_bundle_impl.hpp | 16 ++++++++++++---- sycl/source/detail/scheduler/commands.cpp | 15 ++++++--------- sycl/source/kernel_bundle.cpp | 6 ++++++ sycl/test/abi/sycl_symbols_linux.dump | 3 ++- 6 files changed, 45 insertions(+), 18 deletions(-) diff --git a/sycl/include/CL/sycl/kernel_bundle.hpp b/sycl/include/CL/sycl/kernel_bundle.hpp index 854b76aef2b5e..81226a7af7dad 100644 --- a/sycl/include/CL/sycl/kernel_bundle.hpp +++ b/sycl/include/CL/sycl/kernel_bundle.hpp @@ -471,6 +471,21 @@ using DevImgSelectorImpl = __SYCL_EXPORT detail::KernelBundleImplPtr get_kernel_bundle_impl(const context &Ctx, const std::vector &Devs, bundle_state State, const DevImgSelectorImpl &Selector); + +// Internal non-template versions of get_empty_interop_kernel_bundle API which +// is used by public onces +__SYCL_EXPORT detail::KernelBundleImplPtr +get_empty_interop_kernel_bundle_impl(const context &Ctx, + const std::vector &Devs); + +/// make_kernel may need an empty interop kernel bundle. This function supplies +/// this. +template +kernel_bundle get_empty_interop_kernel_bundle(const context &Ctx) { + detail::KernelBundleImplPtr Impl = + detail::get_empty_interop_kernel_bundle_impl(Ctx, Ctx.get_devices()); + return detail::createSyclObjFromImpl>(Impl); +} } // namespace detail /// A kernel bundle in state State which contains all of the device images for diff --git a/sycl/source/backend.cpp b/sycl/source/backend.cpp index a2a6a7a1561a4..691bdbf0ab91c 100644 --- a/sycl/source/backend.cpp +++ b/sycl/source/backend.cpp @@ -246,10 +246,10 @@ kernel make_kernel(const context &TargetContext, kernel make_kernel(pi_native_handle NativeHandle, const context &TargetContext, backend Backend) { - return make_kernel(TargetContext, - get_kernel_bundle( - TargetContext, std::vector{}), - NativeHandle, false, Backend); + return make_kernel( + TargetContext, + get_empty_interop_kernel_bundle(TargetContext), + NativeHandle, false, Backend); } } // namespace detail diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index b70910c88ca5b..fa4a336aa8a3f 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -87,19 +87,24 @@ class kernel_bundle_impl { MContext, MDevices, State); } - // Interop constructor - kernel_bundle_impl(context Ctx, std::vector Devs, - device_image_plain &DevImage) + // Interop constructor used by make_kernel + kernel_bundle_impl(context Ctx, std::vector Devs) : MContext(Ctx), MDevices(Devs) { if (!checkAllDevicesAreInContext(Devs, Ctx)) throw sycl::exception( make_error_code(errc::invalid), "Not all devices are associated with the context or " "vector of devices is empty"); - MDeviceImages.push_back(DevImage); MIsInterop = true; } + // Interop constructor + kernel_bundle_impl(context Ctx, std::vector Devs, + device_image_plain &DevImage) + : kernel_bundle_impl(Ctx, Devs) { + MDeviceImages.push_back(DevImage); + } + // Matches sycl::build and sycl::compile // Have one constructor because sycl::build and sycl::compile have the same // signature @@ -476,6 +481,9 @@ class kernel_bundle_impl { size_t size() const noexcept { return MDeviceImages.size(); } bundle_state get_bundle_state() const { + // Interop kernel-bundles are always in executable state + if (MIsInterop) + return bundle_state::executable; // All device images are expected to have the same state return MDeviceImages.empty() ? bundle_state::input diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index f74b85e229942..8b277d4bda297 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -1911,15 +1911,12 @@ cl_int enqueueImpKernel( std::shared_ptr SyclKernelImpl; std::shared_ptr DeviceImageImpl; - // Use kernel_bundle is available - if (KernelBundleImplPtr) { - - std::shared_ptr KernelIDImpl = - std::make_shared(KernelName); - - kernel SyclKernel = KernelBundleImplPtr->get_kernel( - detail::createSyclObjFromImpl(KernelIDImpl), - KernelBundleImplPtr); + // Use kernel_bundle is available unless it is interop + if (KernelBundleImplPtr && !KernelBundleImplPtr->isInterop()) { + kernel_id KernelID = + detail::ProgramManager::getInstance().getSYCLKernelID(KernelName); + kernel SyclKernel = + KernelBundleImplPtr->get_kernel(KernelID, KernelBundleImplPtr); SyclKernelImpl = detail::getSyclObjImpl(SyclKernel); diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 1e52a424e5088..a0aeb84ddb3ff 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -138,6 +138,12 @@ get_kernel_bundle_impl(const context &Ctx, const std::vector &Devs, State); } +detail::KernelBundleImplPtr +get_empty_interop_kernel_bundle_impl(const context &Ctx, + const std::vector &Devs) { + return std::make_shared(Ctx, Devs); +} + std::shared_ptr join_impl(const std::vector &Bundles) { return std::make_shared(Bundles); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 12a2de86c252a..6cee9bf59f23f 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3865,6 +3865,7 @@ _ZN2cl4sycl6detail2pi9assertionEbPKc _ZN2cl4sycl6detail2pi9getPluginILNS0_7backendE1EEERKNS1_6pluginEv _ZN2cl4sycl6detail2pi9getPluginILNS0_7backendE2EEERKNS1_6pluginEv _ZN2cl4sycl6detail2pi9getPluginILNS0_7backendE5EEERKNS1_6pluginEv +_ZN2cl4sycl6detail36get_empty_interop_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EE _ZN2cl4sycl6detail6OSUtil10getDirNameB5cxx11EPKc _ZN2cl4sycl6detail6OSUtil11alignedFreeEPv _ZN2cl4sycl6detail6OSUtil12alignedAllocEmm @@ -4215,7 +4216,6 @@ _ZNK2cl4sycl6kernel11get_backendEv _ZNK2cl4sycl6kernel11get_contextEv _ZNK2cl4sycl6kernel11get_programEv _ZNK2cl4sycl6kernel13getNativeImplEv -_ZNK2cl4sycl6kernel9getNativeEv _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 @@ -4244,6 +4244,7 @@ _ZNK2cl4sycl6kernel8get_infoILNS0_4info6kernelE4498EEENS3_12param_traitsIS4_XT_E _ZNK2cl4sycl6kernel8get_infoILNS0_4info6kernelE4499EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl6kernel8get_infoILNS0_4info6kernelE4500EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl6kernel8get_infoILNS0_4info6kernelE4501EEENS3_12param_traitsIS4_XT_EE11return_typeEv +_ZNK2cl4sycl6kernel9getNativeEv _ZNK2cl4sycl6stream22get_max_statement_sizeEv _ZNK2cl4sycl6stream8get_sizeEv _ZNK2cl4sycl6streameqERKS1_ From ca5c3a4584f287d9242cfe1f84cf2e19aa744304 Mon Sep 17 00:00:00 2001 From: Denis Kabanov Date: Fri, 10 Dec 2021 14:08:12 +0300 Subject: [PATCH 2/6] Fix windows symbols --- sycl/test/abi/sycl_symbols_windows.dump | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index e33424c33e2ef..2e4881ce9a63e 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -1847,7 +1847,6 @@ ?fill@MemoryManager@detail@sycl@cl@@SAXPEAVSYCLMemObjI@234@PEAXV?$shared_ptr@Vqueue_impl@detail@sycl@cl@@@std@@_KPEBDIV?$range@$02@34@5V?$id@$02@34@IV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@7@AEAPEAU_pi_event@@@Z ?fill_usm@MemoryManager@detail@sycl@cl@@SAXPEAXV?$shared_ptr@Vqueue_impl@detail@sycl@cl@@@std@@_KHV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@6@AEAPEAU_pi_event@@@Z ?finalize@handler@sycl@cl@@AEAA?AVevent@23@XZ -?getESIMDDeviceInterface@detail@sycl@cl@@YAPEAUESIMDDeviceInterface@123@XZ ?find_device_intersection@detail@sycl@cl@@YA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@AEBV?$vector@V?$kernel_bundle@$00@sycl@cl@@V?$allocator@V?$kernel_bundle@$00@sycl@cl@@@std@@@5@@Z ?floor@__host_std@cl@@YA?AV?$vec@M$00@sycl@2@V342@@Z ?floor@__host_std@cl@@YA?AV?$vec@M$01@sycl@2@V342@@Z @@ -2068,6 +2067,7 @@ ?getDevices@?$image_impl@$01@detail@sycl@cl@@AEAA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@V?$shared_ptr@Vcontext_impl@detail@sycl@cl@@@6@@Z ?getDevices@?$image_impl@$02@detail@sycl@cl@@AEAA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@V?$shared_ptr@Vcontext_impl@detail@sycl@cl@@@6@@Z ?getDirName@OSUtil@detail@sycl@cl@@SA?AV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@PEBD@Z +?getESIMDDeviceInterface@detail@sycl@cl@@YAPEAUESIMDDeviceInterface@123@XZ ?getElementSize@?$image_impl@$00@detail@sycl@cl@@QEBA_KXZ ?getElementSize@?$image_impl@$01@detail@sycl@cl@@QEBA_KXZ ?getElementSize@?$image_impl@$02@detail@sycl@cl@@QEBA_KXZ @@ -2090,13 +2090,13 @@ ?getNative@device@sycl@cl@@AEBA_KXZ ?getNative@device_image_plain@detail@sycl@cl@@QEBA_KXZ ?getNative@event@sycl@cl@@AEBA_KXZ +?getNative@kernel@sycl@cl@@AEBA_KXZ ?getNative@platform@sycl@cl@@AEBA_KXZ ?getNative@program@sycl@cl@@AEBA_KXZ ?getNative@queue@sycl@cl@@AEBA_KXZ ?getNativeContext@interop_handle@sycl@cl@@AEBA_KXZ ?getNativeDevice@interop_handle@sycl@cl@@AEBA_KXZ ?getNativeImpl@kernel@sycl@cl@@AEBA_KXZ -?getNative@kernel@sycl@cl@@AEBA_KXZ ?getNativeMem@interop_handle@sycl@cl@@AEBA_KPEAVAccessorImplHost@detail@23@@Z ?getNativeQueue@interop_handle@sycl@cl@@AEBA_KXZ ?getOSMemSize@OSUtil@detail@sycl@cl@@SA_KXZ @@ -2154,6 +2154,7 @@ ?get_devices@kernel_bundle_plain@detail@sycl@cl@@QEBA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@XZ ?get_devices@platform@sycl@cl@@QEBA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@W4device_type@info@23@@Z ?get_devices@program@sycl@cl@@QEBA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@XZ +?get_empty_interop_kernel_bundle_impl@detail@sycl@cl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@sycl@cl@@@std@@AEBVcontext@23@AEBV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@5@@Z ?get_filtering_mode@sampler@sycl@cl@@QEBA?AW4filtering_mode@23@XZ ?get_filtering_mode@sampler_impl@detail@sycl@cl@@QEBA?AW4filtering_mode@34@XZ ?get_flags@stream@sycl@cl@@AEBAIXZ From ffa6a88432c7f16b34f7845e042d6392412c108f Mon Sep 17 00:00:00 2001 From: Denis Kabanov Date: Fri, 10 Dec 2021 20:45:45 +0300 Subject: [PATCH 3/6] Add regression test --- .../check_carrying_real_kernel_IDs.cpp | 37 +++++++++++++++++++ 1 file changed, 37 insertions(+) create mode 100644 sycl/test/regression/check_carrying_real_kernel_IDs.cpp diff --git a/sycl/test/regression/check_carrying_real_kernel_IDs.cpp b/sycl/test/regression/check_carrying_real_kernel_IDs.cpp new file mode 100644 index 0000000000000..b415f5bdbeb16 --- /dev/null +++ b/sycl/test/regression/check_carrying_real_kernel_IDs.cpp @@ -0,0 +1,37 @@ +// RUN: %clangxx -fsycl -lOpenCL -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %t.out + +#include +#include +#include + +using namespace cl::sycl; + +int main() { + queue Queue{}; + + const char KernelCode[] = "__kernel void foo() { }\n"; + const size_t KernelCodeSize = sizeof(KernelCode); + const char *CLCode[1] = {KernelCode}; + + auto Context = Queue.get_info(); + auto Device = Queue.get_info(); + cl_context CLContext = get_native(Context); + cl_device_id CLDevice = get_native(Device); + + cl_int Err; + + cl_program CLProgram = + clCreateProgramWithSource(CLContext, 1, CLCode, &KernelCodeSize, &Err); + assert(Err == CL_SUCCESS); + Err = clBuildProgram(CLProgram, 1, &CLDevice, "", nullptr, nullptr); + assert(Err == CL_SUCCESS); + + cl_kernel CLKernel = clCreateKernel(CLProgram, "foo", &Err); + assert(Err == CL_SUCCESS); + kernel SYCLKernel = sycl::make_kernel(CLKernel, Context); + + Queue.submit( + [&](handler &commandgroup) { commandgroup.single_task(SYCLKernel); }); + return 0; +} From b7a54f0ec02a4ecd25bbb3df3419623f3118c29a Mon Sep 17 00:00:00 2001 From: Denis Kabanov Date: Fri, 10 Dec 2021 20:51:58 +0300 Subject: [PATCH 4/6] Fix clang-format --- sycl/test/regression/check_carrying_real_kernel_IDs.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/regression/check_carrying_real_kernel_IDs.cpp b/sycl/test/regression/check_carrying_real_kernel_IDs.cpp index b415f5bdbeb16..6fa348fe79e35 100644 --- a/sycl/test/regression/check_carrying_real_kernel_IDs.cpp +++ b/sycl/test/regression/check_carrying_real_kernel_IDs.cpp @@ -13,7 +13,7 @@ int main() { const char KernelCode[] = "__kernel void foo() { }\n"; const size_t KernelCodeSize = sizeof(KernelCode); const char *CLCode[1] = {KernelCode}; - + auto Context = Queue.get_info(); auto Device = Queue.get_info(); cl_context CLContext = get_native(Context); From 99a71b2953638e27668e80bc9a7a84683bdf87e2 Mon Sep 17 00:00:00 2001 From: Denis Kabanov Date: Mon, 13 Dec 2021 18:02:29 +0300 Subject: [PATCH 5/6] Move test to llvm-test-suite --- .../check_carrying_real_kernel_IDs.cpp | 37 ------------------- 1 file changed, 37 deletions(-) delete mode 100644 sycl/test/regression/check_carrying_real_kernel_IDs.cpp diff --git a/sycl/test/regression/check_carrying_real_kernel_IDs.cpp b/sycl/test/regression/check_carrying_real_kernel_IDs.cpp deleted file mode 100644 index 6fa348fe79e35..0000000000000 --- a/sycl/test/regression/check_carrying_real_kernel_IDs.cpp +++ /dev/null @@ -1,37 +0,0 @@ -// RUN: %clangxx -fsycl -lOpenCL -fsycl-targets=%sycl_triple %s -o %t.out -// RUN: %t.out - -#include -#include -#include - -using namespace cl::sycl; - -int main() { - queue Queue{}; - - const char KernelCode[] = "__kernel void foo() { }\n"; - const size_t KernelCodeSize = sizeof(KernelCode); - const char *CLCode[1] = {KernelCode}; - - auto Context = Queue.get_info(); - auto Device = Queue.get_info(); - cl_context CLContext = get_native(Context); - cl_device_id CLDevice = get_native(Device); - - cl_int Err; - - cl_program CLProgram = - clCreateProgramWithSource(CLContext, 1, CLCode, &KernelCodeSize, &Err); - assert(Err == CL_SUCCESS); - Err = clBuildProgram(CLProgram, 1, &CLDevice, "", nullptr, nullptr); - assert(Err == CL_SUCCESS); - - cl_kernel CLKernel = clCreateKernel(CLProgram, "foo", &Err); - assert(Err == CL_SUCCESS); - kernel SYCLKernel = sycl::make_kernel(CLKernel, Context); - - Queue.submit( - [&](handler &commandgroup) { commandgroup.single_task(SYCLKernel); }); - return 0; -} From 1b42495b489fe8a106841490ed0f7bb854982adb Mon Sep 17 00:00:00 2001 From: Denis Kabanov Date: Wed, 15 Dec 2021 11:29:22 +0300 Subject: [PATCH 6/6] Update comment --- sycl/source/detail/scheduler/commands.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 8b277d4bda297..74195290428d3 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -1911,7 +1911,11 @@ cl_int enqueueImpKernel( std::shared_ptr SyclKernelImpl; std::shared_ptr DeviceImageImpl; - // Use kernel_bundle is available unless it is interop + // Use kernel_bundle if available unless it is interop. + // Interop bundles can't be used in the first branch, because the kernels + // in interop kernel bundles (if any) do not have kernel_id + // and can therefore not be looked up, but since they are self-contained + // they can simply be launched directly. if (KernelBundleImplPtr && !KernelBundleImplPtr->isInterop()) { kernel_id KernelID = detail::ProgramManager::getInstance().getSYCLKernelID(KernelName);