[SYCL] Extend eviction to kernel_compiler cache

sycl/source/detail/persistent_device_code_cache.cpp

@@ -49,13 +49,11 @@ LockCacheItem::LockCacheItem(const std::string &Path)
 
 LockCacheItem::~LockCacheItem() {
   if (Owned && std::remove(FileName.c_str()))
-    PersistentDeviceCodeCache::trace("Failed to release lock file: " +
-                                     FileName);
+    PersistentDeviceCodeCache::trace("Failed to release lock file: ", FileName);
 }
 
 // Returns true if the specified format is either SPIRV or a native binary.
-static bool
-IsSupportedImageFormat(ur::DeviceBinaryType Format) {
+static bool IsSupportedImageFormat(ur::DeviceBinaryType Format) {
   return Format == SYCL_DEVICE_BINARY_TYPE_SPIRV ||
          Format == SYCL_DEVICE_BINARY_TYPE_NATIVE;
 }
@@ -210,6 +208,16 @@ void PersistentDeviceCodeCache::repopulateCacheSizeFile(
   const std::string CacheSizeFileName = "cache_size.txt";
   const std::string CacheSizeFile = CacheRoot + "/" + CacheSizeFileName;
 
+  // Create cache root, if it does not exist.
+  try {
+    if (!OSUtil::isPathPresent(CacheRoot))
+      OSUtil::makeDir(CacheRoot.c_str());
+  } catch (...) {
+    throw sycl::exception(make_error_code(errc::runtime),
+                          "Failed to create cache root directory: " +
+                              CacheRoot);
+  }
+
   // If the cache size file is not present, calculate the size of the cache size
   // directory and write it to the file.
   if (!OSUtil::isPathPresent(CacheSizeFile)) {
@@ -316,6 +324,8 @@ void PersistentDeviceCodeCache::evictItemsFromCache(
         auto RemoveFileAndSubtractSize = [&CurrCacheSize](
                                              const std::string &FileName) {
           // If the file is not present, return.
+          // Src file is not present inj kernel_compiler cache, we will
+          // skip removing it.
           if (!OSUtil::isPathPresent(FileName))
             return;
 
@@ -324,7 +334,7 @@ void PersistentDeviceCodeCache::evictItemsFromCache(
             throw sycl::exception(make_error_code(errc::runtime),
                                   "Failed to evict cache entry: " + FileName);
           } else {
-            PersistentDeviceCodeCache::trace("File removed: " + FileName);
+            PersistentDeviceCodeCache::trace("File removed: ", FileName);
             CurrCacheSize -= FileSize;
           }
         };
@@ -464,7 +474,7 @@ void PersistentDeviceCodeCache::putItemToDisc(
       if (Lock.isOwned()) {
         std::string FullFileName = FileName + ".bin";
         writeBinaryDataToFile(FullFileName, BinaryData[DeviceIndex]);
-        trace("device binary has been cached: " + FullFileName);
+        trace("device binary has been cached: ", FullFileName);
         writeSourceItem(FileName + ".src", Devices[DeviceIndex], SortedImgs,
                         SpecConsts, BuildOptionsString);
 
@@ -474,7 +484,7 @@ void PersistentDeviceCodeCache::putItemToDisc(
 
         saveCurrentTimeInAFile(FileName + CacheEntryAccessTimeSuffix);
       } else {
-        PersistentDeviceCodeCache::trace("cache lock not owned " + FileName);
+        PersistentDeviceCodeCache::trace("cache lock not owned ", FileName);
       }
     } catch (std::exception &e) {
       PersistentDeviceCodeCache::trace(
@@ -495,7 +505,20 @@ void PersistentDeviceCodeCache::putItemToDisc(
 void PersistentDeviceCodeCache::putCompiledKernelToDisc(
     const std::vector<device> &Devices, const std::string &BuildOptionsString,
     const std::string &SourceStr, const ur_program_handle_t &NativePrg) {
+
+  repopulateCacheSizeFile(getRootDir());
+
+  // Do not insert any new item if eviction is in progress.
+  // Since evictions are rare, we can afford to spin lock here.
+  const std::string EvictionInProgressFile =
+      getRootDir() + EvictionInProgressFileSuffix;
+  // Stall until the other process finishes eviction.
+  while (OSUtil::isPathPresent(EvictionInProgressFile))
+    continue;
+
   auto BinaryData = getProgramBinaryData(NativePrg, Devices);
+  // Total size of the item that we are writing to the cache.
+  size_t TotalSize = 0;
 
   for (size_t DeviceIndex = 0; DeviceIndex < Devices.size(); DeviceIndex++) {
     // If we don't have binary for the device, skip it.
@@ -512,10 +535,13 @@ void PersistentDeviceCodeCache::putCompiledKernelToDisc(
         std::string FullFileName = FileName + ".bin";
         writeBinaryDataToFile(FullFileName, BinaryData[DeviceIndex]);
         PersistentDeviceCodeCache::trace_KernelCompiler(
-            "binary has been cached: " + FullFileName);
+            "binary has been cached: ", FullFileName);
+
+        TotalSize += getFileSize(FullFileName);
+        saveCurrentTimeInAFile(FileName + CacheEntryAccessTimeSuffix);
       } else {
-        PersistentDeviceCodeCache::trace_KernelCompiler(
-            "cache lock not owned " + FileName);
+        PersistentDeviceCodeCache::trace_KernelCompiler("cache lock not owned ",
+                                                        FileName);
       }
     } catch (std::exception &e) {
       PersistentDeviceCodeCache::trace_KernelCompiler(
@@ -525,6 +551,10 @@ void PersistentDeviceCodeCache::putCompiledKernelToDisc(
           std::string("error outputting cache: ") + std::strerror(errno));
     }
   }
+
+  // Update the cache size file and trigger cache eviction if needed.
+  if (TotalSize)
+    updateCacheFileSizeAndTriggerEviction(getRootDir(), TotalSize);
 }
 
 /* Program binaries built for one or more devices are read from persistent
@@ -581,7 +611,7 @@ std::vector<std::vector<char>> PersistentDeviceCodeCache::getItemFromDisc(
     if (Binaries[DeviceIndex].empty())
       return {};
   }
-  PersistentDeviceCodeCache::trace("using cached device binary: " + FileNames);
+  PersistentDeviceCodeCache::trace("using cached device binary: ", FileNames);
   return Binaries;
 }
 
@@ -611,6 +641,12 @@ PersistentDeviceCodeCache::getCompiledKernelFromDisc(
         try {
           std::string FullFileName = FileName + ".bin";
           Binaries[DeviceIndex] = readBinaryDataFromFile(FullFileName);
+
+          // Explicitly update the access time of the file. This is required for
+          // eviction.
+          if (isEvictionEnabled())
+            saveCurrentTimeInAFile(FileName + CacheEntryAccessTimeSuffix);
+
           FileNames += FullFileName + ";";
           break;
         } catch (...) {
@@ -623,7 +659,7 @@ PersistentDeviceCodeCache::getCompiledKernelFromDisc(
     if (Binaries[DeviceIndex].empty())
       return {};
   }
-  PersistentDeviceCodeCache::trace_KernelCompiler("using cached binary: " +
+  PersistentDeviceCodeCache::trace_KernelCompiler("using cached binary: ",
                                                   FileNames);
   return Binaries;
 }
@@ -654,7 +690,7 @@ void PersistentDeviceCodeCache::writeBinaryDataToFile(
   FileStream.write((char *)&Size, sizeof(Size));
   FileStream.write(Data.data(), Size);
   if (FileStream.fail())
-    trace("Failed to write to binary file " + FileName);
+    trace("Failed to write to binary file ", FileName);
 }
 
 /* Read built binary from persistent cache. Each persistent cache file contains
@@ -671,7 +707,7 @@ PersistentDeviceCodeCache::readBinaryDataFromFile(const std::string &FileName) {
   size_t NumBinaries = 0;
   FileStream.read((char *)&NumBinaries, sizeof(NumBinaries));
   if (FileStream.fail()) {
-    trace("Failed to read number of binaries from " + FileName);
+    trace("Failed to read number of binaries from ", FileName);
     return {};
   }
   // Even in the old implementation we could only put a single binary to the
@@ -686,7 +722,7 @@ PersistentDeviceCodeCache::readBinaryDataFromFile(const std::string &FileName) {
   FileStream.close();
 
   if (FileStream.fail()) {
-    trace("Failed to read binary file from " + FileName);
+    trace("Failed to read binary file from ", FileName);
     return {};
   }
 
@@ -726,7 +762,7 @@ void PersistentDeviceCodeCache::writeSourceItem(
   FileStream.close();
 
   if (FileStream.fail()) {
-    trace("Failed to write source file to " + FileName);
+    trace("Failed to write source file to ", FileName);
   }
 }
 
@@ -774,7 +810,7 @@ bool PersistentDeviceCodeCache::isCacheItemSrcEqual(
   FileStream.close();
 
   if (FileStream.fail()) {
-    trace("Failed to read source file from " + FileName);
+    trace("Failed to read source file from ", FileName);
   }
 
   return true;

sycl/source/detail/persistent_device_code_cache.hpp

@@ -208,17 +208,23 @@ class PersistentDeviceCodeCache {
                                       const ur_program_handle_t &NativePrg);
 
   /* Sends message to std:cerr stream when SYCL_CACHE_TRACE environemnt is set*/
-  static void trace(const std::string &msg) {
+  static void trace(const std::string &msg, std::string path = "") {
     static const bool traceEnabled =
         SYCLConfig<SYCL_CACHE_TRACE>::isTraceDiskCache();
-    if (traceEnabled)
-      std::cerr << "[Persistent Cache]: " << msg << std::endl;
+    if (traceEnabled) {
+      std::replace(path.begin(), path.end(), '\\', '/');
+      std::cerr << "[Persistent Cache]: " << msg << path << std::endl;
+    }
   }
-  static void trace_KernelCompiler(const std::string &msg) {
+  static void trace_KernelCompiler(const std::string &msg,
+                                   std::string path = "") {
     static const bool traceEnabled =
         SYCLConfig<SYCL_CACHE_TRACE>::isTraceKernelCompiler();
-    if (traceEnabled)
-      std::cerr << "[kernel_compiler Persistent Cache]: " << msg << std::endl;
+    if (traceEnabled) {
+      std::replace(path.begin(), path.end(), '\\', '/');
+      std::cerr << "[kernel_compiler Persistent Cache]: " << msg << path
+                << std::endl;
+    }
   }
 
 private:

sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp

@@ -0,0 +1,128 @@
+//==-kernel_compiler_cache_eviction.cpp -- kernel_compiler extension tests -==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// Tests on-disk cache and eviction with kernel_compiler.
+
+// REQUIRES: ocloc && (opencl || level_zero)
+// UNSUPPORTED: accelerator
+// UNSUPPORTED-INTENDED: kernel_compiler is not available for accelerator
+// devices.
+
+// -- Test the kernel_compiler with OpenCL source.
+// RUN: %{build} -o %t.out
+
+// -- Test again, with caching.
+// DEFINE: %{cache_vars} = env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=7 SYCL_CACHE_DIR=%t/cache_dir SYCL_CACHE_MAX_SIZE=30000
+// RUN: %if run-mode %{rm -rf %t/cache_dir%}
+// RUN: %{cache_vars} %{run-unfiltered-devices} %t.out 2>&1 | FileCheck %s --check-prefix=CHECK
+
+// CHECK: [Persistent Cache]: enabled
+
+#include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
+
+auto constexpr CLSource = R"===(
+__kernel void my_kernel(__global int *in, __global int *out) {
+  size_t i = get_global_id(0);
+  out[i] = in[i]*2 + 100;
+}
+__kernel void her_kernel(__global int *in, __global int *out) {
+  size_t i = get_global_id(0);
+  out[i] = in[i]*5 + 1000;
+}
+)===";
+
+using namespace sycl;
+
+void test_build_and_run() {
+  namespace syclex = sycl::ext::oneapi::experimental;
+  using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
+
+  // only one device is supported at this time, so we limit the queue and
+  // context to that
+  sycl::device d{sycl::default_selector_v};
+  sycl::context ctx{d};
+  sycl::queue q{ctx, d};
+
+  bool ok =
+      q.get_device().ext_oneapi_can_compile(syclex::source_language::opencl);
+  if (!ok) {
+    std::cout << "Apparently this device does not support OpenCL C source "
+                 "kernel bundle extension: "
+              << q.get_device().get_info<sycl::info::device::name>()
+              << std::endl;
+    return;
+  }
+
+  auto CreateAndVerifyKB = [](source_kb &kbSrc,
+                              std::vector<std::string> &&BuildFlags) {
+    std::string log;
+    std::vector<sycl::device> devs = kbSrc.get_devices();
+    sycl::context ctxRes = kbSrc.get_context();
+    sycl::backend beRes = kbSrc.get_backend();
+
+    auto kb =
+        syclex::build(kbSrc, devs,
+                      syclex::properties{syclex::build_options{BuildFlags},
+                                         syclex::save_log{&log}});
+
+    bool hasMyKernel = kb.ext_oneapi_has_kernel("my_kernel");
+    bool hasHerKernel = kb.ext_oneapi_has_kernel("her_kernel");
+    bool notExistKernel = kb.ext_oneapi_has_kernel("not_exist");
+    assert(hasMyKernel && "my_kernel should exist, but doesn't");
+    assert(hasHerKernel && "her_kernel should exist, but doesn't");
+    assert(!notExistKernel && "non-existing kernel should NOT exist.");
+  };
+
+  source_kb kbSrc = syclex::create_kernel_bundle_from_source(
+      ctx, syclex::source_language::opencl, CLSource);
+
+  // compilation with props and devices
+  std::vector<std::string> flags{"-cl-fast-relaxed-math",
+                                 "-cl-finite-math-only", "-cl-no-signed-zeros",
+                                 "-cl-unsafe-math-optimizations"};
+
+  // Device image #1
+  // CHECK: [Persistent Cache]: Cache size file not present. Creating one.
+  // CHECK-NEXT: [Persistent Cache]: Cache size file created.
+  // CHECK-NEXT: [kernel_compiler Persistent Cache]: binary has been cached: [[DEVIMG1:.*]]
+  // CHECK-NEXT: [Persistent Cache]: Updating the cache size file.
+  CreateAndVerifyKB(kbSrc, {});
+
+  // Device image #2
+  // CHECK-NEXT: [kernel_compiler Persistent Cache]: binary has been cached: [[DEVIMG2:.*]]
+  // CHECK-NEXT: [Persistent Cache]: Updating the cache size file.
+  CreateAndVerifyKB(kbSrc, {flags[0], flags[1], flags[2], flags[3]});
+
+  // Re-insert device image #1
+  // CHECK-NEXT: [kernel_compiler Persistent Cache]: using cached binary: [[DEVIMG1]]
+  CreateAndVerifyKB(kbSrc, {});
+
+  // Insert more unique device images to trigger cache eviction.
+  // Make sure Device image #2 is evicted before device image #1 as
+  // eviction is LRU-based.
+  // CHECK: [Persistent Cache]: Cache eviction triggered.
+  // CHECK-NEXT: [Persistent Cache]: File removed: [[DEVIMG2]]
+  // CHECK-NEXT: [Persistent Cache]: File removed: [[DEVIMG1]]
+  for (int i = 0; i < flags.size(); i++) {
+    CreateAndVerifyKB(kbSrc, {flags[i]});
+  }
+}
+
+int main() {
+#ifndef SYCL_EXT_ONEAPI_KERNEL_COMPILER_OPENCL
+  static_assert(false, "KernelCompiler OpenCL feature test macro undefined");
+#endif
+
+#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER
+  test_build_and_run();
+#else
+  static_assert(false, "Kernel Compiler feature test macro undefined");
+#endif
+  return 0;
+}