Skip to content

[SYCL] Fix kernel bundles don't really carry kernel IDs #5121

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 7 commits into from
Dec 22, 2021
Merged

[SYCL] Fix kernel bundles don't really carry kernel IDs #5121

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
0e2c6fa
[SYCL] Fix kernel bundles don't really carry kernel IDs
denis-kabanov Dec 9, 2021
ca5c3a4
Fix windows symbols
denis-kabanov Dec 10, 2021
ffa6a88
Add regression test
denis-kabanov Dec 10, 2021
b7a54f0
Fix clang-format
denis-kabanov Dec 10, 2021
99a71b2
Move test to llvm-test-suite
denis-kabanov Dec 13, 2021
1b42495
Update comment
denis-kabanov Dec 15, 2021
d4cb543
Merge branch 'sycl' of https://github.com/intel/llvm into fix_KernelId
denis-kabanov Dec 21, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
15 changes: 15 additions & 0 deletions sycl/include/CL/sycl/kernel_bundle.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -471,6 +471,21 @@ using DevImgSelectorImpl =
__SYCL_EXPORT detail::KernelBundleImplPtr
get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &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<device> &Devs);

/// make_kernel may need an empty interop kernel bundle. This function supplies
/// this.
template <bundle_state State>
kernel_bundle<State> 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<sycl::kernel_bundle<State>>(Impl);
}
} // namespace detail

/// A kernel bundle in state State which contains all of the device images for
Expand Down
8 changes: 4 additions & 4 deletions sycl/source/backend.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<bundle_state::executable>(
TargetContext, std::vector<kernel_id>{}),
NativeHandle, false, Backend);
return make_kernel(
TargetContext,
get_empty_interop_kernel_bundle<bundle_state::executable>(TargetContext),
NativeHandle, false, Backend);
}

} // namespace detail
Expand Down
16 changes: 12 additions & 4 deletions sycl/source/detail/kernel_bundle_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -87,19 +87,24 @@ class kernel_bundle_impl {
MContext, MDevices, State);
}

// Interop constructor
kernel_bundle_impl(context Ctx, std::vector<device> Devs,
device_image_plain &DevImage)
// Interop constructor used by make_kernel
kernel_bundle_impl(context Ctx, std::vector<device> 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<device> 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
Expand Down Expand Up @@ -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
Expand Down
19 changes: 10 additions & 9 deletions sycl/source/detail/scheduler/commands.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1922,15 +1922,16 @@ cl_int enqueueImpKernel(
std::shared_ptr<kernel_impl> SyclKernelImpl;
std::shared_ptr<device_image_impl> DeviceImageImpl;

// Use kernel_bundle is available
if (KernelBundleImplPtr) {

std::shared_ptr<kernel_id_impl> KernelIDImpl =
std::make_shared<kernel_id_impl>(KernelName);

kernel SyclKernel = KernelBundleImplPtr->get_kernel(
detail::createSyclObjFromImpl<kernel_id>(KernelIDImpl),
KernelBundleImplPtr);
// 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);
kernel SyclKernel =
KernelBundleImplPtr->get_kernel(KernelID, KernelBundleImplPtr);

SyclKernelImpl = detail::getSyclObjImpl(SyclKernel);

Expand Down
6 changes: 6 additions & 0 deletions sycl/source/kernel_bundle.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -138,6 +138,12 @@ get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
State);
}

detail::KernelBundleImplPtr
get_empty_interop_kernel_bundle_impl(const context &Ctx,
const std::vector<device> &Devs) {
return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs);
}

std::shared_ptr<detail::kernel_bundle_impl>
join_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles) {
return std::make_shared<detail::kernel_bundle_impl>(Bundles);
Expand Down
3 changes: 2 additions & 1 deletion sycl/test/abi/sycl_symbols_linux.dump
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3891,6 +3891,7 @@ _ZN2cl4sycl6detail2pi9assertionEbPKc
_ZN2cl4sycl6detail2pi9getPluginILNS0_7backendE1EEERKNS1_6pluginEv
_ZN2cl4sycl6detail2pi9getPluginILNS0_7backendE2EEERKNS1_6pluginEv
_ZN2cl4sycl6detail2pi9getPluginILNS0_7backendE5EEERKNS1_6pluginEv
_ZN2cl4sycl6detail36get_empty_interop_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EE
_ZN2cl4sycl6detail6OSUtil10getDirNameB5cxx11EPKc
_ZN2cl4sycl6detail6OSUtil11alignedFreeEPv
_ZN2cl4sycl6detail6OSUtil12alignedAllocEmm
Expand Down Expand Up @@ -4243,7 +4244,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
Expand Down Expand Up @@ -4272,6 +4272,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_
Expand Down
5 changes: 3 additions & 2 deletions sycl/test/abi/sycl_symbols_windows.dump
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1873,7 +1873,6 @@
?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
?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@PEAPEAU_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
Expand Down Expand Up @@ -2094,6 +2093,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
Expand All @@ -2116,13 +2116,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
Expand Down Expand Up @@ -2180,6 +2180,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
Expand Down