[SYCL] Disable fallback assert for interop kernels

sycl/source/detail/kernel_bundle_impl.hpp

@@ -97,6 +97,7 @@ class kernel_bundle_impl {
           "Not all devices are associated with the context or "
           "vector of devices is empty");
     MDeviceImages.push_back(DevImage);
+    MIsInterop = true;
   }
 
   // Matches sycl::build and sycl::compile
@@ -373,6 +374,8 @@ class kernel_bundle_impl {
     std::shared_ptr<kernel_impl> KernelImpl = std::make_shared<kernel_impl>(
         Kernel, detail::getSyclObjImpl(MContext), DeviceImageImpl, Self);
 
+    KernelImpl->setInterop(MIsInterop);
+
     return detail::createSyclObjFromImpl<kernel>(KernelImpl);
   }
 
@@ -489,6 +492,7 @@ class kernel_bundle_impl {
   // This map stores values for specialization constants, that are missing
   // from any device image.
   SpecConstMapT MSpecConstValues;
+  bool MIsInterop = false;
 };
 
 } // namespace detail

sycl/source/detail/kernel_impl.cpp

@@ -30,6 +30,8 @@ kernel_impl::kernel_impl(RT::PiKernel Kernel, ContextImplPtr Context)
   // For others, PI will turn this into a NOP.
   getPlugin().call<PiApiKind::piKernelSetExecInfo>(
       MKernel, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE);
+
+  MIsInterop = true;
 }
 
 kernel_impl::kernel_impl(RT::PiKernel Kernel, ContextImplPtr ContextImpl,

sycl/source/detail/kernel_impl.hpp

@@ -200,13 +200,17 @@ class kernel_impl {
 
   KernelBundleImplPtr get_kernel_bundle() const { return MKernelBundleImpl; }
 
+  void setInterop(bool V) { MIsInterop = V; }
+  bool isInterop() const { return MIsInterop; }
+
 private:
   RT::PiKernel MKernel;
   const ContextImplPtr MContext;
   const ProgramImplPtr MProgramImpl;
   bool MCreatedFromSource = true;
   const DeviceImageImplPtr MDeviceImageImpl;
   const KernelBundleImplPtr MKernelBundleImpl;
+  bool MIsInterop = false;
 };
 
 template <info::kernel param>

sycl/source/detail/program_impl.cpp

@@ -244,6 +244,7 @@ void program_impl::compile_with_source(std::string KernelSource,
     compile(CompileOptions);
   }
   MState = program_state::compiled;
+  MIsInterop = true;
 }
 
 void program_impl::build_with_kernel_name(std::string KernelName,
@@ -275,6 +276,7 @@ void program_impl::build_with_source(std::string KernelSource,
     build(BuildOptions);
   }
   MState = program_state::linked;
+  MIsInterop = true;
 }
 
 void program_impl::link(std::string LinkOptions) {
@@ -311,17 +313,21 @@ kernel program_impl::get_kernel(std::string KernelName,
                                 std::shared_ptr<program_impl> PtrToSelf,
                                 bool IsCreatedFromSource) const {
   throw_if_state_is(program_state::none);
+  std::shared_ptr<kernel_impl> KImpl;
   if (is_host()) {
     if (IsCreatedFromSource)
       throw invalid_object_error("This instance of program is a host instance",
                                  PI_INVALID_PROGRAM);
 
-    return createSyclObjFromImpl<kernel>(
-        std::make_shared<kernel_impl>(MContext, PtrToSelf));
-  }
-  return createSyclObjFromImpl<kernel>(std::make_shared<kernel_impl>(
-      get_pi_kernel(KernelName), MContext, PtrToSelf,
-      /*IsCreatedFromSource*/ IsCreatedFromSource));
+    KImpl = std::make_shared<kernel_impl>(MContext, PtrToSelf);
+
+  } else
+    KImpl = std::make_shared<kernel_impl>(
+        get_pi_kernel(KernelName), MContext, PtrToSelf,
+        /*IsCreatedFromSource*/ IsCreatedFromSource);
+
+  KImpl->setInterop(MIsInterop);
+  return createSyclObjFromImpl<kernel>(KImpl);
 }
 
 std::vector<std::vector<char>> program_impl::get_binaries() const {

sycl/source/detail/program_impl.hpp

@@ -448,6 +448,8 @@ class program_impl {
   /// device list and context) and built with build_with_kernel_type with
   /// default build options
   bool MProgramAndKernelCachingAllowed = false;
+
+  bool MIsInterop = false;
 };
 
 template <>

sycl/source/detail/queue_impl.hpp

@@ -24,6 +24,7 @@
 #include <detail/context_impl.hpp>
 #include <detail/device_impl.hpp>
 #include <detail/event_impl.hpp>
+#include <detail/kernel_impl.hpp>
 #include <detail/plugin.hpp>
 #include <detail/scheduler/scheduler.hpp>
 #include <detail/thread_pool.hpp>
@@ -448,7 +449,7 @@ class queue_impl {
       bool KernelUsesAssert = false;
       if (IsKernel)
         KernelUsesAssert =
-            Handler.MKernel ? true
+            Handler.MKernel ? Handler.MKernel->isInterop()
                             : ProgramManager::getInstance().kernelUsesAssert(
                                   Handler.MOSModuleHandle, Handler.MKernelName);