[SYCL] Don't set PI_USM_INDIRECT_ACCESS if platform don't support it

sycl/source/detail/kernel_impl.cpp

@@ -27,8 +27,9 @@ kernel_impl::kernel_impl(sycl::detail::pi::PiKernel Kernel,
   // Enable USM indirect access for interoperability kernels.
   // Some PI Plugins (like OpenCL) require this call to enable USM
   // For others, PI will turn this into a NOP.
-  getPlugin()->call<PiApiKind::piKernelSetExecInfo>(
-      MKernel, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE);
+  if (Context->getPlatformImpl()->supports_usm())
+    getPlugin()->call<PiApiKind::piKernelSetExecInfo>(
+        MKernel, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE);
 
   // This constructor is only called in the interoperability kernel constructor.
   MIsInterop = true;

sycl/source/detail/platform_impl.cpp

@@ -589,6 +589,11 @@ bool platform_impl::has_extension(const std::string &ExtensionName) const {
   return (AllExtensionNames.find(ExtensionName) != std::string::npos);
 }
 
+bool platform_impl::supports_usm() const {
+  return getBackend() != backend::opencl ||
+         has_extension("cl_intel_unified_shared_memory");
+}
+
 pi_native_handle platform_impl::getNative() const {
   const auto &Plugin = getPlugin();
   pi_native_handle Handle;

sycl/source/detail/platform_impl.hpp

@@ -61,6 +61,12 @@ class platform_impl {
   /// \return true if platform supports specified extension.
   bool has_extension(const std::string &ExtensionName) const;
 
+  /// Checks if this platform supports usm.
+  /// Non opencl backends are assumed to support it.
+  ///
+  /// \return true if platform supports usm.
+  bool supports_usm() const;
+
   /// Returns all SYCL devices associated with this platform.
   ///
   /// If this platform is a host platform and device type requested is either

sycl/source/detail/program_impl.cpp

@@ -402,8 +402,9 @@ program_impl::get_pi_kernel_arg_mask_pair(const std::string &KernelName) const {
 
     // Some PI Plugins (like OpenCL) require this call to enable USM
     // For others, PI will turn this into a NOP.
-    Plugin->call<PiApiKind::piKernelSetExecInfo>(
-        Result.first, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE);
+    if (getContextImplPtr()->getPlatformImpl()->supports_usm())
+      Plugin->call<PiApiKind::piKernelSetExecInfo>(
+          Result.first, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE);
 
     return Result;
 }

sycl/source/detail/program_manager/program_manager.cpp

@@ -689,10 +689,13 @@ ProgramManager::getOrCreateKernel(const ContextImplPtr &ContextImpl,
     Plugin->call<errc::kernel_not_supported, PiApiKind::piKernelCreate>(
         Program, KernelName.c_str(), &Kernel);
 
-    // Some PI Plugins (like OpenCL) require this call to enable USM
-    // For others, PI will turn this into a NOP.
-    Plugin->call<PiApiKind::piKernelSetExecInfo>(Kernel, PI_USM_INDIRECT_ACCESS,
-                                                 sizeof(pi_bool), &PI_TRUE);
+    // Only set PI_USM_INDIRECT_ACCESS if the platform can handle it.
+    if (ContextImpl->getPlatformImpl()->supports_usm()) {
+      // Some PI Plugins (like OpenCL) require this call to enable USM
+      // For others, PI will turn this into a NOP.
+      Plugin->call<PiApiKind::piKernelSetExecInfo>(
+          Kernel, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE);
+    }
 
     const KernelArgMask *ArgMask = nullptr;
     if (!m_UseSpvFile)
@@ -2361,8 +2364,10 @@ ProgramManager::getOrCreateKernel(const context &Context,
     Plugin->call<PiApiKind::piKernelCreate>(Program, KernelName.c_str(),
                                             &Kernel);
 
-    Plugin->call<PiApiKind::piKernelSetExecInfo>(Kernel, PI_USM_INDIRECT_ACCESS,
-                                                 sizeof(pi_bool), &PI_TRUE);
+    // Only set PI_USM_INDIRECT_ACCESS if the platform can handle it.
+    if (Ctx->getPlatformImpl()->supports_usm())
+      Plugin->call<PiApiKind::piKernelSetExecInfo>(
+          Kernel, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE);
 
     // Ignore possible m_UseSpvFile for now.
     // TODO consider making m_UseSpvFile interact with kernel bundles as well.

sycl/test-e2e/XPTI/basic_event_collection_linux.cpp

@@ -17,11 +17,15 @@
 // CHECK:      PI Call Begin : piextQueueCreate
 // CHECK:      PI Call Begin : piextDeviceSelectBinary
 // CHECK:      PI Call Begin : piKernelCreate
+// CHECK-NEXT: PI Call Begin : piPlatformGetInfo
+// CHECK-NEXT: PI Call Begin : piPlatformGetInfo
 // CHECK-NEXT: PI Call Begin : piKernelSetExecInfo
 // CHECK:      PI Call Begin : piextKernelSetArgPointer
 // CHECK-NEXT: PI Call Begin : piKernelGetGroupInfo
 // CHECK-NEXT: PI Call Begin : piEnqueueKernelLaunch
 // CHECK:      PI Call Begin : piKernelCreate
+// CHECK-NEXT: PI Call Begin : piPlatformGetInfo
+// CHECK-NEXT: PI Call Begin : piPlatformGetInfo
 // CHECK-NEXT: PI Call Begin : piKernelSetExecInfo
 // CHECK:      Node create
 // CHECK-DAG:    sym_line_no : {{.*}}