[SYCL] adding binary caching support to kernel_compiler extension

sycl/source/detail/kernel_bundle_impl.hpp

@@ -12,6 +12,7 @@
 #include <detail/kernel_compiler/kernel_compiler_opencl.hpp>
 #include <detail/kernel_compiler/kernel_compiler_sycl.hpp>
 #include <detail/kernel_impl.hpp>
+#include <detail/persistent_device_code_cache.hpp>
 #include <detail/program_manager/program_manager.hpp>
 #include <sycl/backend_types.hpp>
 #include <sycl/context.hpp>
@@ -396,6 +397,44 @@ class kernel_bundle_impl {
     return SS.str();
   }
 
+  bool
+  extKernelCompilerFetchFromCache(const std::vector<device> Devices,
+                                  const std::vector<std::string> &BuildOptions,
+                                  const std::string &SourceStr,
+                                  ur_program_handle_t &UrProgram) {
+    using ContextImplPtr = std::shared_ptr<sycl::detail::context_impl>;
+    ContextImplPtr ContextImpl = getSyclObjImpl(MContext);
+    const AdapterPtr &Adapter = ContextImpl->getAdapter();
+
+    std::string UserArgs = syclex::detail::userArgsAsString(BuildOptions);
+    auto BinProg = PersistentDeviceCodeCache::getCompiledKernelFromDisc(
+        Devices[0], UserArgs, SourceStr);
+    if (!BinProg.empty()) {
+      ur_device_handle_t UrDevice = getSyclObjImpl(Devices[0])->getHandleRef();
+      ur_result_t BinaryStatus = UR_RESULT_SUCCESS;
+      ur_program_properties_t Properties = {};
+      std::vector<ur_program_metadata_t> Metadata = {};
+      Properties.stype = UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES;
+      Properties.pNext = nullptr;
+      Properties.count = Metadata.size();
+      Properties.pMetadatas = Metadata.data();
+      BinaryStatus = Adapter->call_nocheck<UrApiKind::urProgramCreateWithBinary>(
+          ContextImpl->getHandleRef(), UrDevice, BinProg[0].size(),
+          (const unsigned char *)BinProg[0].data(), &Properties, &UrProgram);
+
+      if (BinaryStatus == UR_RESULT_SUCCESS) {
+        ur_result_t Error = Adapter->call_nocheck<UrApiKind::urProgramBuildExp>(
+            UrProgram,
+            /*num devices =*/1, &UrDevice, UserArgs.c_str());
+
+        if (Error == UR_RESULT_SUCCESS) {
+          return true;
+        }
+      }
+    }
+    return false;
+  }
+
   std::shared_ptr<kernel_bundle_impl>
   build_from_source(const std::vector<device> Devices,
                     const std::vector<std::string> &BuildOptions,
@@ -415,46 +454,55 @@ class kernel_bundle_impl {
       DeviceVec.push_back(Dev);
     }
 
-    const auto spirv = [&]() -> std::vector<uint8_t> {
-      if (Language == syclex::source_language::opencl) {
-        // if successful, the log is empty. if failed, throws an error with the
-        // compilation log.
-        const auto &SourceStr = std::get<std::string>(this->Source);
-        std::vector<uint32_t> IPVersionVec(Devices.size());
-        std::transform(DeviceVec.begin(), DeviceVec.end(), IPVersionVec.begin(),
-                       [&](ur_device_handle_t d) {
-                         uint32_t ipVersion = 0;
-                         Adapter->call<UrApiKind::urDeviceGetInfo>(
-                             d, UR_DEVICE_INFO_IP_VERSION, sizeof(uint32_t),
-                             &ipVersion, nullptr);
-                         return ipVersion;
-                       });
-        return syclex::detail::OpenCLC_to_SPIRV(SourceStr, IPVersionVec,
-                                                BuildOptions, LogPtr);
-      }
-      if (Language == syclex::source_language::spirv) {
-        const auto &SourceBytes =
-            std::get<std::vector<std::byte>>(this->Source);
-        std::vector<uint8_t> Result(SourceBytes.size());
-        std::transform(SourceBytes.cbegin(), SourceBytes.cend(), Result.begin(),
-                       [](std::byte B) { return static_cast<uint8_t>(B); });
-        return Result;
-      }
-      if (Language == syclex::source_language::sycl) {
-        const auto &SourceStr = std::get<std::string>(this->Source);
-        return syclex::detail::SYCL_to_SPIRV(SourceStr, IncludePairs,
-                                             BuildOptions, LogPtr,
-                                             RegisteredKernelNames);
-      }
-      throw sycl::exception(
-          make_error_code(errc::invalid),
-          "OpenCL C and SPIR-V are the only supported languages at this time");
-    }();
-
     ur_program_handle_t UrProgram = nullptr;
+    // SourceStrPtr will be null when source is Spir-V bytes.
+    const std::string *SourceStrPtr = std::get_if<std::string>(&this->Source);
+    bool FetchedFromCache = false;
+    if (PersistentDeviceCodeCache::isEnabled() && SourceStrPtr) {
+      FetchedFromCache = extKernelCompilerFetchFromCache(
+          Devices, BuildOptions, *SourceStrPtr, UrProgram);
+    }
+
+    if (!FetchedFromCache) {
+      const auto spirv = [&]() -> std::vector<uint8_t> {
+        if (Language == syclex::source_language::opencl) {
+          // if successful, the log is empty. if failed, throws an error with
+          // the compilation log.
+          std::vector<uint32_t> IPVersionVec(Devices.size());
+          std::transform(DeviceVec.begin(), DeviceVec.end(), IPVersionVec.begin(),
+           [&](ur_device_handle_t d) {
+                           uint32_t ipVersion = 0;
+                           Adapter->call<UrApiKind::urDeviceGetInfo>(
+                               d, UR_DEVICE_INFO_IP_VERSION, sizeof(uint32_t),
+                               &ipVersion, nullptr);
+                           return ipVersion;
+                         });
+          return syclex::detail::OpenCLC_to_SPIRV(*SourceStrPtr, IPVersionVec,
+                                                  BuildOptions, LogPtr);
+        }
+        if (Language == syclex::source_language::spirv) {
+          const auto &SourceBytes =
+              std::get<std::vector<std::byte>>(this->Source);
+          std::vector<uint8_t> Result(SourceBytes.size());
+          std::transform(SourceBytes.cbegin(), SourceBytes.cend(),
+                         Result.begin(),
+                         [](std::byte B) { return static_cast<uint8_t>(B); });
+          return Result;
+        }
+        if (Language == syclex::source_language::sycl) {
+          return syclex::detail::SYCL_to_SPIRV(*SourceStrPtr, IncludePairs,
+                                               BuildOptions, LogPtr,
+                                               RegisteredKernelNames);
+        }
+        throw sycl::exception(
+            make_error_code(errc::invalid),
+            "SYCL C++, OpenCL C and SPIR-V are the only supported "
+            "languages at this time");
+      }();
+
     Adapter->call<UrApiKind::urProgramCreateWithIL>(ContextImpl->getHandleRef(),
-                                                    spirv.data(), spirv.size(),
-                                                    nullptr, &UrProgram);
+                                                   spirv.data(), spirv.size(),
+                                                   nullptr, &UrProgram);
     // program created by urProgramCreateWithIL is implicitly retained.
     if (UrProgram == nullptr)
       throw sycl::exception(
@@ -470,6 +518,8 @@ class kernel_bundle_impl {
     }
     Adapter->checkUrResult<errc::build>(Res);
 
+    } // if(!FetchedFromCache)
+
     // Get the number of kernels in the program.
     size_t NumKernels;
     Adapter->call<UrApiKind::urProgramGetInfo>(
@@ -495,6 +545,15 @@ class kernel_bundle_impl {
         nullptr, MContext, MDevices, bundle_state::executable, KernelIDs,
         UrProgram);
     device_image_plain DevImg{DevImgImpl};
+
+    // If caching enabled and kernel not fetched from cache, cache.
+    if (PersistentDeviceCodeCache::isEnabled() && !FetchedFromCache &&
+        SourceStrPtr) {
+      PersistentDeviceCodeCache::putCompiledKernelToDisc(
+          Devices[0], syclex::detail::userArgsAsString(BuildOptions),
+          *SourceStrPtr, UrProgram);
+    }
+
     return std::make_shared<kernel_bundle_impl>(MContext, MDevices, DevImg,
                                                 KernelNames, Language);
   }

sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp

@@ -31,6 +31,8 @@ SYCL_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs,
 
 bool SYCL_Compilation_Available();
 
+std::string userArgsAsString(const std::vector<std::string> &UserArguments);
+
 } // namespace detail
 } // namespace ext::oneapi::experimental
 

sycl/source/detail/persistent_device_code_cache.cpp

@@ -110,27 +110,22 @@ getSortedImages(const std::vector<const RTDeviceBinaryImage *> &Imgs) {
   return SortedImgs;
 }
 
-/* Stores built program in persistent cache
- */
-void PersistentDeviceCodeCache::putItemToDisc(
-    const device &Device, const std::vector<const RTDeviceBinaryImage *> &Imgs,
-    const SerializedObj &SpecConsts, const std::string &BuildOptionsString,
-    const ur_program_handle_t &NativePrg) {
-
-  if (!areImagesCacheable(Imgs))
-    return;
-
-  std::vector<const RTDeviceBinaryImage *> SortedImgs = getSortedImages(Imgs);
-  std::string DirName =
-      getCacheItemPath(Device, SortedImgs, SpecConsts, BuildOptionsString);
-
-  if (DirName.empty())
-    return;
+// Utility function to get a non-yet-existing unique filename.
+std::string getUniqueFilename(const std::string &base_name) {
+  size_t i = 0;
+  std::string filename = base_name + "/" + std::to_string(i++);
+  while (OSUtil::isPathPresent(filename + ".bin") ||
+         OSUtil::isPathPresent(filename + ".lock")) {
+    filename = base_name + "/" + std::to_string(i++);
+  }
+  return filename;
+}
 
+std::vector<std::vector<char>>
+getProgramBinaryData(const ur_program_handle_t &NativePrg,
+                     const device &Device) {
   auto Adapter = detail::getSyclObjImpl(Device)->getAdapter();
-
   unsigned int DeviceNum = 0;
-
   Adapter->call<UrApiKind::urProgramGetInfo>(
       NativePrg, UR_PROGRAM_INFO_NUM_DEVICES, sizeof(DeviceNum), &DeviceNum,
       nullptr);
@@ -147,22 +142,37 @@ void PersistentDeviceCodeCache::putItemToDisc(
     Pointers.push_back(Result[I].data());
   }
 
-  Adapter->call<UrApiKind::urProgramGetInfo>(
-      NativePrg, UR_PROGRAM_INFO_BINARIES, sizeof(char *) * Pointers.size(),
-      Pointers.data(), nullptr);
-  size_t i = 0;
-  std::string FileName;
-  do {
-    FileName = DirName + "/" + std::to_string(i++);
-  } while (OSUtil::isPathPresent(FileName + ".bin") ||
-           OSUtil::isPathPresent(FileName + ".lock"));
+  Adapter->call<UrApiKind::urProgramGetInfo>(NativePrg, UR_PROGRAM_INFO_BINARIES,
+                                            sizeof(char *) * Pointers.size(),
+                                            Pointers.data(), nullptr);
+  return Result;
+}
+
+/* Stores built program in persistent cache
+ */
+void PersistentDeviceCodeCache::putItemToDisc(
+    const device &Device, const std::vector<const RTDeviceBinaryImage *> &Imgs,
+    const SerializedObj &SpecConsts, const std::string &BuildOptionsString,
+    const ur_program_handle_t &NativePrg) {
+
+  if (!areImagesCacheable(Imgs))
+    return;
+
+  std::vector<const RTDeviceBinaryImage *> SortedImgs = getSortedImages(Imgs);
+  std::string DirName =
+      getCacheItemPath(Device, SortedImgs, SpecConsts, BuildOptionsString);
+
+  if (DirName.empty())
+    return;
 
   try {
     OSUtil::makeDir(DirName.c_str());
+    std::string FileName = getUniqueFilename(DirName);
     LockCacheItem Lock{FileName};
     if (Lock.isOwned()) {
       std::string FullFileName = FileName + ".bin";
-      writeBinaryDataToFile(FullFileName, Result);
+      writeBinaryDataToFile(FullFileName,
+                            getProgramBinaryData(NativePrg, Device));
       trace("device binary has been cached: " + FullFileName);
       writeSourceItem(FileName + ".src", Device, SortedImgs, SpecConsts,
                       BuildOptionsString);
@@ -180,6 +190,36 @@ void PersistentDeviceCodeCache::putItemToDisc(
   }
 }
 
+void PersistentDeviceCodeCache::putCompiledKernelToDisc(
+    const device &Device, const std::string &BuildOptionsString,
+    const std::string &SourceStr, const ur_program_handle_t &NativePrg) {
+
+  std::string DirName =
+      getCompiledKernelItemPath(Device, BuildOptionsString, SourceStr);
+
+  try {
+    OSUtil::makeDir(DirName.c_str());
+    std::string FileName = getUniqueFilename(DirName);
+    LockCacheItem Lock{FileName};
+    if (Lock.isOwned()) {
+      std::string FullFileName = FileName + ".bin";
+      writeBinaryDataToFile(FullFileName,
+                            getProgramBinaryData(NativePrg, Device));
+      trace("kernel_compiler binary has been cached: " + FullFileName);
+    } else {
+      PersistentDeviceCodeCache::trace("cache lock not owned " + FileName);
+    }
+  } catch (std::exception &e) {
+    PersistentDeviceCodeCache::trace(
+        std::string("exception encountered making persistent cache: ") +
+        e.what());
+  } catch (...) {
+    PersistentDeviceCodeCache::trace(
+        std::string("error outputting persistent cache: ") +
+        std::strerror(errno));
+  }
+}
+
 /* Program binaries built for one or more devices are read from persistent
  * cache and returned in form of vector of programs. Each binary program is
  * stored in vector of chars.
@@ -222,6 +262,42 @@ std::vector<std::vector<char>> PersistentDeviceCodeCache::getItemFromDisc(
   return {};
 }
 
+/*  kernel_compiler extension uses slightly different format for path
+    and does not cache a .src separate from the binary.
+ */
+std::vector<std::vector<char>>
+PersistentDeviceCodeCache::getCompiledKernelFromDisc(
+    const device &Device, const std::string &BuildOptionsString,
+    const std::string SourceStr) {
+
+  std::string DirName =
+      getCompiledKernelItemPath(Device, BuildOptionsString, SourceStr);
+
+  if (DirName.empty() || !OSUtil::isPathPresent(DirName))
+    return {};
+
+  int i = 0;
+
+  std::string FileName{DirName + "/" + std::to_string(i)};
+  while (OSUtil::isPathPresent(FileName + ".bin") ||
+         OSUtil::isPathPresent(FileName + ".src")) {
+
+    if (!LockCacheItem::isLocked(FileName)) {
+      try {
+        std::string FullFileName = FileName + ".bin";
+        std::vector<std::vector<char>> res =
+            readBinaryDataFromFile(FullFileName);
+        trace("kernel_compiler using cached binary: " + FullFileName);
+        return res; // subject for NRVO
+      } catch (...) {
+        // If read was unsuccessfull try the next item
+      }
+    }
+    FileName = DirName + "/" + std::to_string(++i);
+  }
+  return {};
+}
+
 /* Returns string value which can be used to identify different device
  */
 std::string PersistentDeviceCodeCache::getDeviceIDString(const device &Device) {
@@ -396,6 +472,25 @@ std::string PersistentDeviceCodeCache::getCacheItemPath(
          std::to_string(StringHasher(BuildOptionsString));
 }
 
+std::string PersistentDeviceCodeCache::getCompiledKernelItemPath(
+    const device &Device, const std::string &BuildOptionsString,
+    const std::string SourceString) {
+
+  std::string cache_root{getRootDir()};
+  if (cache_root.empty()) {
+    trace("Disable persistent cache due to unconfigured cache root.");
+    return {};
+  }
+
+  std::string DeviceString{getDeviceIDString(Device)};
+  std::hash<std::string> StringHasher{};
+
+  return cache_root + "/ext_kernel_compiler" + "/" +
+         std::to_string(StringHasher(DeviceString)) + "/" +
+         std::to_string(StringHasher(BuildOptionsString)) + "/" +
+         std::to_string(StringHasher(SourceString));
+}
+
 /* Returns true if persistent cache is enabled.
  */
 bool PersistentDeviceCodeCache::isEnabled() {