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

This is an issue complaining about that return from getSyclObjImpl can be null.
I talked to @steffenlarsen off line and although this is considered to be impossible (every constructor should create for itself and impl), it doesn't seem to be enforced anywhere in the code. We decided that would be a good compromise to make sure that impl is always non-null.

Signed-off-by: Zahira Ammarguellat <[email protected]>

Author: zahiraam

sycl/include/sycl/detail/common.hpp

@@ -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;
 }
 

sycl/include/sycl/handler.hpp

@@ -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);
 
@@ -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});
@@ -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());
@@ -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());
@@ -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());
@@ -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());
@@ -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());
@@ -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());

sycl/source/handler.cpp

@@ -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()).
+  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.

sycl/test/abi/sycl_symbols_linux.dump

@@ -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_