[SYCL] Disable fallback assert for interop kernels (#4712)

Signed-off-by: Sergey Kanaev <[email protected]>

Author: sergei

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
@@ -482,13 +483,16 @@ class kernel_bundle_impl {
     return MSpecConstValues;
   }
 
+  bool isInterop() const { return MIsInterop; }
+
 private:
   context MContext;
   std::vector<device> MDevices;
   std::vector<device_image_plain> MDeviceImages;
   // 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,
@@ -47,6 +49,8 @@ kernel_impl::kernel_impl(RT::PiKernel Kernel, ContextImplPtr ContextImpl,
     throw cl::sycl::invalid_parameter_error(
         "Input context must be the same as the context of cl_kernel",
         PI_INVALID_CONTEXT);
+
+  MIsInterop = MProgramImpl->isInterop();
 }
 
 kernel_impl::kernel_impl(RT::PiKernel Kernel, ContextImplPtr ContextImpl,
@@ -60,6 +64,8 @@ kernel_impl::kernel_impl(RT::PiKernel Kernel, ContextImplPtr ContextImpl,
   if (!is_host()) {
     getPlugin().call<PiApiKind::piKernelRetain>(MKernel);
   }
+
+  MIsInterop = MKernelBundleImpl->isInterop();
 }
 
 kernel_impl::kernel_impl(ContextImplPtr Context,

sycl/source/detail/kernel_impl.hpp

@@ -200,13 +200,16 @@ class kernel_impl {
 
   KernelBundleImplPtr get_kernel_bundle() const { return MKernelBundleImpl; }
 
+  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

@@ -117,7 +117,9 @@ program_impl::program_impl(
 
 program_impl::program_impl(ContextImplPtr Context,
                            pi_native_handle InteropProgram)
-    : program_impl(Context, InteropProgram, nullptr) {}
+    : program_impl(Context, InteropProgram, nullptr) {
+  MIsInterop = true;
+}
 
 program_impl::program_impl(ContextImplPtr Context,
                            pi_native_handle InteropProgram,
@@ -198,7 +200,9 @@ program_impl::program_impl(ContextImplPtr Context,
 program_impl::program_impl(ContextImplPtr Context, RT::PiKernel Kernel)
     : program_impl(Context, reinterpret_cast<pi_native_handle>(nullptr),
                    ProgramManager::getInstance().getPiProgramFromPiKernel(
-                       Kernel, Context)) {}
+                       Kernel, Context)) {
+  MIsInterop = true;
+}
 
 program_impl::~program_impl() {
   // TODO catch an exception and put it to list of asynchronous exceptions
@@ -244,6 +248,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 +280,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) {

sycl/source/detail/program_impl.hpp

@@ -346,6 +346,8 @@ class program_impl {
   /// Returns the native plugin handle.
   pi_native_handle getNative() const;
 
+  bool isInterop() const { return MIsInterop; }
+
 private:
   // Deligating Constructor used in Implementation.
   program_impl(ContextImplPtr Context, pi_native_handle InteropProgram,
@@ -448,6 +450,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>
@@ -470,11 +471,12 @@ class queue_impl {
     if (PostProcess) {
       bool IsKernel = Type == CG::Kernel;
       bool KernelUsesAssert = false;
+
       if (IsKernel)
-        KernelUsesAssert =
-            Handler.MKernel ? true
-                            : ProgramManager::getInstance().kernelUsesAssert(
-                                  Handler.MOSModuleHandle, Handler.MKernelName);
+        // Kernel only uses assert if it's non interop one
+        KernelUsesAssert = !(Handler.MKernel && Handler.MKernel->isInterop()) &&
+                           ProgramManager::getInstance().kernelUsesAssert(
+                               Handler.MOSModuleHandle, Handler.MKernelName);
 
       finalizeHandler(Handler, NeedSeparateDependencyMgmt, Event);