Skip to content

[SYCL] sycl::kernel::get_kernel_bundle() may return a kernel_bundle with a null impl. #6598

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 1 commit into from
Aug 24, 2022
Merged

[SYCL] sycl::kernel::get_kernel_bundle() may return a kernel_bundle with a null impl. #6598

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
1 change: 1 addition & 0 deletions sycl/include/sycl/detail/common.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -226,6 +226,7 @@ namespace detail {
// template <class T>
// friend decltype(T::impl) detail::getSyclObjImpl(const T &SyclObject);
template <class Obj> decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject) {
assert(SyclObject.impl && "every constructor should create an impl");
return SyclObject.impl;
}

Expand Down
16 changes: 9 additions & 7 deletions sycl/include/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1290,6 +1290,8 @@ class __SYCL_EXPORT handler {
std::shared_ptr<detail::kernel_bundle_impl>
getOrInsertHandlerKernelBundle(bool Insert) const;

void setHandlerKernelBundle(kernel Kernel);

void setHandlerKernelBundle(
const std::shared_ptr<detail::kernel_bundle_impl> &NewKernelBundleImpPtr);

Expand Down Expand Up @@ -1918,7 +1920,7 @@ class __SYCL_EXPORT handler {
throwIfActionIsCreated();
verifyKernelInvoc(Kernel);
// Ignore any set kernel bundles and use the one associated with the kernel
setHandlerKernelBundle(detail::getSyclObjImpl(Kernel.get_kernel_bundle()));
setHandlerKernelBundle(Kernel);
// No need to check if range is out of INT_MAX limits as it's compile-time
// known constant
MNDRDesc.set(range<1>{1});
Expand Down Expand Up @@ -1991,7 +1993,7 @@ class __SYCL_EXPORT handler {
void single_task(kernel Kernel, _KERNELFUNCPARAM(KernelFunc)) {
throwIfActionIsCreated();
// Ignore any set kernel bundles and use the one associated with the kernel
setHandlerKernelBundle(detail::getSyclObjImpl(Kernel.get_kernel_bundle()));
setHandlerKernelBundle(Kernel);
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
verifyUsedKernelBundle(detail::KernelInfo<NameT>::getName());
Expand Down Expand Up @@ -2037,7 +2039,7 @@ class __SYCL_EXPORT handler {
_KERNELFUNCPARAM(KernelFunc)) {
throwIfActionIsCreated();
// Ignore any set kernel bundles and use the one associated with the kernel
setHandlerKernelBundle(detail::getSyclObjImpl(Kernel.get_kernel_bundle()));
setHandlerKernelBundle(Kernel);
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
verifyUsedKernelBundle(detail::KernelInfo<NameT>::getName());
Expand Down Expand Up @@ -2075,7 +2077,7 @@ class __SYCL_EXPORT handler {
id<Dims> WorkItemOffset, _KERNELFUNCPARAM(KernelFunc)) {
throwIfActionIsCreated();
// Ignore any set kernel bundles and use the one associated with the kernel
setHandlerKernelBundle(detail::getSyclObjImpl(Kernel.get_kernel_bundle()));
setHandlerKernelBundle(Kernel);
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
verifyUsedKernelBundle(detail::KernelInfo<NameT>::getName());
Expand Down Expand Up @@ -2113,7 +2115,7 @@ class __SYCL_EXPORT handler {
_KERNELFUNCPARAM(KernelFunc)) {
throwIfActionIsCreated();
// Ignore any set kernel bundles and use the one associated with the kernel
setHandlerKernelBundle(detail::getSyclObjImpl(Kernel.get_kernel_bundle()));
setHandlerKernelBundle(Kernel);
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
verifyUsedKernelBundle(detail::KernelInfo<NameT>::getName());
Expand Down Expand Up @@ -2155,7 +2157,7 @@ class __SYCL_EXPORT handler {
_KERNELFUNCPARAM(KernelFunc)) {
throwIfActionIsCreated();
// Ignore any set kernel bundles and use the one associated with the kernel
setHandlerKernelBundle(detail::getSyclObjImpl(Kernel.get_kernel_bundle()));
setHandlerKernelBundle(Kernel);
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
verifyUsedKernelBundle(detail::KernelInfo<NameT>::getName());
Expand Down Expand Up @@ -2195,7 +2197,7 @@ class __SYCL_EXPORT handler {
_KERNELFUNCPARAM(KernelFunc)) {
throwIfActionIsCreated();
// Ignore any set kernel bundles and use the one associated with the kernel
setHandlerKernelBundle(detail::getSyclObjImpl(Kernel.get_kernel_bundle()));
setHandlerKernelBundle(Kernel);
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
verifyUsedKernelBundle(detail::KernelInfo<NameT>::getName());
Expand Down
9 changes: 9 additions & 0 deletions sycl/source/handler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -82,6 +82,15 @@ void handler::setHandlerKernelBundle(
MImpl->MKernelBundle = NewKernelBundleImpPtr;
}

void handler::setHandlerKernelBundle(kernel Kernel) {
// Kernel may not have an associated kernel bundle if it is created from a
// program. As such, apply getSyclObjImpl directly on the kernel, i.e. not
// the other way around: getSyclObjImp(Kernel->get_kernel_bundle()).
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think this is not directly related to your change, but this comment makes me think there is a problem someplace. Why is there no kernel bundle for a kernel created from a program? The SYCL spec defines kernel::get_kernel_bundle() for all kernels, even those created from backend interop.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Arguably there is, but since program is being removed I don't think it is worth the worry.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

OK, I was wondering if the issue about having no kernel bundle is limited to just the deprecated program case, or if it also included other cases like constructing a kernel from backend interop. If it is only the program case, I agree there is no need to fix it now. We should instead wait until program is removed.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This problem should be unique to kernel bundles in kernels created from sycl::program. For sycl::make_kernel we either use a supplied kernel bundle (part of the L0 input type) or an empty "interop" kernel bundle.

std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImpl =
detail::getSyclObjImpl(Kernel)->get_kernel_bundle();
setHandlerKernelBundle(KernelBundleImpl);
}

event handler::finalize() {
// This block of code is needed only for reduction implementation.
// It is harmless (does nothing) for everything else.
Expand Down
1 change: 1 addition & 0 deletions sycl/test/abi/sycl_symbols_linux.dump
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3957,6 +3957,7 @@ _ZN4sycl3_V17handler18extractArgsAndReqsEv
_ZN4sycl3_V17handler20DisableRangeRoundingEv
_ZN4sycl3_V17handler20associateWithHandlerEPNS0_6detail16AccessorBaseHostENS0_6access6targetE
_ZN4sycl3_V17handler20setStateSpecConstSetEv
_ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE
_ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE
_ZN4sycl3_V17handler22verifyUsedKernelBundleERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
_ZN4sycl3_V17handler24GetRangeRoundingSettingsERmS2_S2_
Expand Down