diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp
index 12ae796b2e2e2..89b240b816ff9 100644
--- a/sycl/source/detail/kernel_impl.cpp
+++ b/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;
diff --git a/sycl/source/detail/platform_impl.cpp b/sycl/source/detail/platform_impl.cpp
index 57b4a2f48030b..eb28ffaf3819a 100644
--- a/sycl/source/detail/platform_impl.cpp
+++ b/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;
diff --git a/sycl/source/detail/platform_impl.hpp b/sycl/source/detail/platform_impl.hpp
index 3b08d39ad3738..efa8d8b29c372 100644
--- a/sycl/source/detail/platform_impl.hpp
+++ b/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
diff --git a/sycl/source/detail/program_impl.cpp b/sycl/source/detail/program_impl.cpp
index f3c42e2adcd97..d65f3163b961f 100644
--- a/sycl/source/detail/program_impl.cpp
+++ b/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;
 }
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 961a2ae394f94..87b9da025c1a7 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/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.
diff --git a/sycl/test-e2e/XPTI/basic_event_collection_linux.cpp b/sycl/test-e2e/XPTI/basic_event_collection_linux.cpp
index 0af0121c9313a..61b53feed0622 100644
--- a/sycl/test-e2e/XPTI/basic_event_collection_linux.cpp
+++ b/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 : {{.*}}