[SYCL] Kernel Compiler OpenCL ext_can_compile rejects older ocloc

sycl/include/sycl/kernel_bundle.hpp

@@ -1101,8 +1101,6 @@ struct is_property_key_of<registered_names_key,
 
 namespace detail {
 // forward decls
-__SYCL_EXPORT bool is_source_kernel_bundle_supported(backend BE,
-                                                     source_language Language);
 
 __SYCL_EXPORT kernel_bundle<bundle_state::ext_oneapi_source>
 make_kernel_bundle_from_source(

sycl/source/detail/device_impl.cpp

@@ -891,8 +891,13 @@ bool device_impl::isGetDeviceAndHostTimerSupported() {
 bool device_impl::extOneapiCanCompile(
     ext::oneapi::experimental::source_language Language) {
   try {
+    // Get the shared_ptr to this object from the platform that owns it.
+    std::shared_ptr<device_impl> Self =
+        MPlatform->getOrMakeDeviceImpl(MDevice, MPlatform);
     return sycl::ext::oneapi::experimental::detail::
-        is_source_kernel_bundle_supported(getBackend(), Language);
+        is_source_kernel_bundle_supported(Language,
+                                          std::vector<DeviceImplPtr>{Self});
+
   } catch (sycl::exception &) {
     return false;
   }

sycl/source/detail/kernel_bundle_impl.hpp

@@ -9,6 +9,7 @@
 #pragma once
 
 #include <detail/device_image_impl.hpp>
+#include <detail/device_impl.hpp>
 #include <detail/kernel_impl.hpp>
 #include <detail/program_manager/program_manager.hpp>
 #include <sycl/backend_types.hpp>
@@ -30,6 +31,18 @@
 
 namespace sycl {
 inline namespace _V1 {
+
+namespace ext::oneapi::experimental::detail {
+using DeviceImplPtr = std::shared_ptr<sycl::detail::device_impl>;
+bool is_source_kernel_bundle_supported(
+    sycl::ext::oneapi::experimental::source_language Language,
+    const context &Ctx);
+
+bool is_source_kernel_bundle_supported(
+    sycl::ext::oneapi::experimental::source_language Language,
+    const std::vector<DeviceImplPtr> &Devices);
+} // namespace ext::oneapi::experimental::detail
+
 namespace detail {
 
 static bool checkAllDevicesAreInContext(const std::vector<device> &Devices,

sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp

@@ -25,6 +25,10 @@ inline namespace _V1 {
 namespace ext::oneapi::experimental {
 namespace detail {
 
+// forward declaration
+std::string InvokeOclocQuery(const std::vector<uint32_t> &IPVersionVec,
+                             const char *identifier);
+
 // ensures the OclocLibrary has the right version, etc.
 void checkOclocLibrary(void *OclocLibrary) {
   void *OclocVersionHandle =
@@ -64,35 +68,60 @@ static std::unique_ptr<void, std::function<void(void *)>>
       std::ignore = sycl::detail::ur::unloadOsLibrary(StoredPtr);
     });
 
-// load the ocloc shared library, check it.
-void loadOclocLibrary() {
+// Load first compatible ocloc shared library.
+void loadOclocLibrary(const std::vector<uint32_t> &IPVersionVec) {
 #ifdef __SYCL_RT_OS_WINDOWS
-  static const std::string OclocLibraryName = "ocloc64.dll";
+  // first the environment, if not compatible will move on to absolute path.
+  static const std::vector<std::string> OclocPaths = {
+      "ocloc64.dll",
+      "C:\\Program Files (x86)\\Intel\\oneAPI\\ocloc\\latest\\ocloc64.dll"};
 #else
-  static const std::string OclocLibraryName = "libocloc.so";
+  // linux always uses the environment.
+  static const std::vector<std::string> OclocPaths = {"libocloc.so"};
 #endif
-  void *tempPtr = OclocLibrary.get();
-  if (tempPtr == nullptr) {
-    tempPtr = sycl::detail::ur::loadOsLibrary(OclocLibraryName);
 
-    if (tempPtr == nullptr)
-      throw sycl::exception(make_error_code(errc::build),
-                            "Unable to load ocloc library " + OclocLibraryName);
+  // attemptLoad() sets OclocLibrary value by side effect.
+  auto attemptLoad = [&](std::string path) {
+    void *tempPtr;
+    try {
+      // Load then perform checks. Each check throws.
+      tempPtr = sycl::detail::ur::loadOsLibrary(path);
+      OclocLibrary.reset(tempPtr);
+
+      if (tempPtr == nullptr)
+        throw sycl::exception(make_error_code(errc::build),
+                              "Unable to load ocloc from " + path);
+
+      checkOclocLibrary(tempPtr);
 
-    checkOclocLibrary(tempPtr);
+      InvokeOclocQuery(IPVersionVec, "CL_DEVICE_OPENCL_C_ALL_VERSIONS");
+    } catch (const sycl::exception &) {
+      tempPtr = nullptr;
+      OclocLibrary.reset(tempPtr);
+    }
 
-    OclocLibrary.reset(tempPtr);
+    return tempPtr;
+  };
+
+  // Attempt to load each, exiting as soon as we find compatible ocloc.
+  for (auto path : OclocPaths) {
+    void *tempPtr = attemptLoad(path);
+    if (tempPtr != nullptr)
+      return;
   }
+
+  // If we haven't exited yet, then throw to indicate failure.
+  throw sycl::exception(make_error_code(errc::build), "Unable to load ocloc");
 }
 
-bool OpenCLC_Compilation_Available() {
+bool OpenCLC_Compilation_Available(const std::vector<uint32_t> &IPVersionVec) {
   // Already loaded?
   if (OclocLibrary != nullptr)
     return true;
 
   try {
     // loads and checks version
-    loadOclocLibrary();
+    loadOclocLibrary(IPVersionVec);
     return true;
   } catch (...) {
     return false;
@@ -102,11 +131,12 @@ bool OpenCLC_Compilation_Available() {
 using voidPtr = void *;
 
 void SetupLibrary(voidPtr &oclocInvokeHandle, voidPtr &oclocFreeOutputHandle,
-                  std::error_code the_errc) {
-  if (!oclocInvokeHandle) {
-    if (OclocLibrary == nullptr)
-      loadOclocLibrary();
+                  std::error_code the_errc,
+                  const std::vector<uint32_t> &IPVersionVec) {
+  if (OclocLibrary == nullptr)
+    loadOclocLibrary(IPVersionVec);
 
+  if (!oclocInvokeHandle) {
     oclocInvokeHandle = sycl::detail::ur::getOsLibraryFuncAddress(
         OclocLibrary.get(), "oclocInvoke");
     if (!oclocInvokeHandle)
@@ -145,7 +175,8 @@ std::string InvokeOclocQuery(const std::vector<uint32_t> &IPVersionVec,
   static void *oclocFreeOutputHandle = nullptr;
   std::error_code the_errc = make_error_code(errc::runtime);
 
-  SetupLibrary(oclocInvokeHandle, oclocFreeOutputHandle, the_errc);
+  SetupLibrary(oclocInvokeHandle, oclocFreeOutputHandle, the_errc,
+               IPVersionVec);
 
   uint32_t NumOutputs = 0;
   uint8_t **Outputs = nullptr;
@@ -205,7 +236,8 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source,
   static void *oclocFreeOutputHandle = nullptr;
   std::error_code build_errc = make_error_code(errc::build);
 
-  SetupLibrary(oclocInvokeHandle, oclocFreeOutputHandle, build_errc);
+  SetupLibrary(oclocInvokeHandle, oclocFreeOutputHandle, build_errc,
+               IPVersionVec);
 
   // assemble ocloc args
   std::string CombinedUserArgs =

sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp

@@ -24,8 +24,8 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source,
                              const std::vector<uint32_t> &IPVersionVec,
                              const std::vector<std::string> &UserArgs,
                              std::string *LogPtr);
-
-bool OpenCLC_Compilation_Available();
+// IPVersionVec gets flattened and passed to ocloc as the -dev flag.
+bool OpenCLC_Compilation_Available(const std::vector<uint32_t> &IPVersionVec);
 
 bool OpenCLC_Feature_Available(const std::string &Feature, uint32_t IPVersion);
 

sycl/source/kernel_bundle.cpp

@@ -387,24 +387,62 @@ namespace detail {
 /////////////////////////
 // syclex::detail::is_source_kernel_bundle_supported
 /////////////////////////
-bool is_source_kernel_bundle_supported(backend BE, source_language Language) {
+
+bool is_source_kernel_bundle_supported(
+    sycl::ext::oneapi::experimental::source_language Language,
+    const std::vector<DeviceImplPtr> &DeviceImplVec) {
+  backend BE = DeviceImplVec[0]->getBackend();
   // Support is limited to the opencl and level_zero backends.
   bool BE_Acceptable = (BE == sycl::backend::ext_oneapi_level_zero) ||
                        (BE == sycl::backend::opencl);
-  if (BE_Acceptable) {
-    if (Language == source_language::opencl) {
-      return detail::OpenCLC_Compilation_Available();
-    } else if (Language == source_language::spirv) {
-      return true;
-    } else if (Language == source_language::sycl) {
-      return detail::SYCL_JIT_Compilation_Available();
-    }
+  if (!BE_Acceptable)
+    return false;
+
+  if (Language == source_language::spirv) {
+    return true;
+  } else if (Language == source_language::sycl) {
+    return detail::SYCL_JIT_Compilation_Available();
+  } else if (Language == source_language::opencl) {
+    if (DeviceImplVec.empty())
+      return false;
+
+    const AdapterPtr &Adapter = DeviceImplVec[0]->getAdapter();
+    std::vector<uint32_t> IPVersionVec;
+    IPVersionVec.reserve(DeviceImplVec.size());
+
+    std::transform(DeviceImplVec.begin(), DeviceImplVec.end(),
+                   std::back_inserter(IPVersionVec),
+                   [&](const DeviceImplPtr &Impl) {
+                     uint32_t ipVersion = 0;
+                     ur_device_handle_t DeviceHandle = Impl->getHandleRef();
+                     Adapter->call<UrApiKind::urDeviceGetInfo>(
+                         DeviceHandle, UR_DEVICE_INFO_IP_VERSION,
+                         sizeof(uint32_t), &ipVersion, nullptr);
+                     return ipVersion;
+                   });
+
+    return detail::OpenCLC_Compilation_Available(IPVersionVec);
   }
 
   // otherwise
   return false;
 }
 
+bool is_source_kernel_bundle_supported(
+    sycl::ext::oneapi::experimental::source_language Language,
+    const context &Ctx) {
+  const std::vector<sycl::device> Devices = Ctx.get_devices();
+  std::vector<DeviceImplPtr> DeviceImplVec;
+  DeviceImplVec.reserve(Devices.size());
+  std::transform(Devices.begin(), Devices.end(),
+                 std::back_inserter(DeviceImplVec),
+                 [](const sycl::device &dev) {
+                   return sycl::detail::getSyclObjImpl(dev);
+                 });
+
+  return is_source_kernel_bundle_supported(Language, DeviceImplVec);
+}
+
 /////////////////////////
 // syclex::detail::create_kernel_bundle_from_source
 /////////////////////////
@@ -428,8 +466,7 @@ make_kernel_bundle_from_source(const context &SyclContext,
   for (auto &p : IncludePairViews)
     IncludePairs.push_back({p.first.data(), p.second.data()});
 
-  backend BE = SyclContext.get_backend();
-  if (!is_source_kernel_bundle_supported(BE, Language))
+  if (!is_source_kernel_bundle_supported(Language, SyclContext))
     throw sycl::exception(make_error_code(errc::invalid),
                           "kernel_bundle creation from source not supported");
 
@@ -448,8 +485,7 @@ source_kb make_kernel_bundle_from_source(const context &SyclContext,
                                          const std::vector<std::byte> &Bytes,
                                          include_pairs_view_t IncludePairs) {
   (void)IncludePairs;
-  backend BE = SyclContext.get_backend();
-  if (!is_source_kernel_bundle_supported(BE, Language))
+  if (!is_source_kernel_bundle_supported(Language, SyclContext))
     throw sycl::exception(make_error_code(errc::invalid),
                           "kernel_bundle creation from source not supported");
 

sycl/test/abi/sycl_symbols_linux.dump

@@ -3105,7 +3105,6 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC2ERKNS0_5
 _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC2ERKNS0_7contextERKNS0_6deviceERKNS0_13property_listE
 _ZN4sycl3_V13ext6oneapi12experimental6detail30make_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageENS0_6detail11string_viewESt6vectorISt4pairISA_SA_ESaISD_EE
 _ZN4sycl3_V13ext6oneapi12experimental6detail30make_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageERKSt6vectorISt4byteSaISA_EES9_ISt4pairINS0_6detail11string_viewESH_ESaISI_EE
-_ZN4sycl3_V13ext6oneapi12experimental6detail33is_source_kernel_bundle_supportedENS0_7backendENS3_15source_languageE
 _ZN4sycl3_V13ext6oneapi12experimental6memcpyENS0_5queueEPvPKvmRKNS0_6detail13code_locationE
 _ZN4sycl3_V13ext6oneapi12experimental6memsetENS0_5queueEPvimRKNS0_6detail13code_locationE
 _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_5queueE

sycl/test/abi/sycl_symbols_windows.dump

@@ -4228,7 +4228,6 @@
 ?is_gpu@device@_V1@sycl@@QEBA_NXZ
 ?is_in_fusion_mode@fusion_wrapper@experimental@codeplay@ext@_V1@sycl@@QEBA_NXZ
 ?is_in_order@queue@_V1@sycl@@QEBA_NXZ
-?is_source_kernel_bundle_supported@detail@experimental@oneapi@ext@_V1@sycl@@YA_NW4backend@56@W4source_language@23456@@Z
 ?is_specialization_constant_set@kernel_bundle_plain@detail@_V1@sycl@@IEBA_NPEBD@Z
 ?join_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@AEBV?$vector@V?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@V?$allocator@V?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@@2@@5@W4bundle_state@23@@Z
 ?khr_get_default_context@platform@_V1@sycl@@QEBA?AVcontext@23@XZ