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

…12780)

If the OpenCL platform doesn't support USM, don't set
PI_USM_INDIRECT_ACCESS exec info. This will avoid SYCL program to fail
when they don't use USM. If the program do need USM support, the runtime
will fail on other API calls (like memory allocation).

---------

Signed-off-by: Victor Lomuller <victor@codeplay.com>

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 : {{.*}}