[SYCL] Implement eviction for in-memory program cache

sycl/source/detail/kernel_program_cache.hpp

@@ -218,14 +218,15 @@ class KernelProgramCache {
         MProgramEvictionList.splice(MProgramEvictionList.end(),
                                     MProgramEvictionList, It->second);
         traceProgram("Program moved to the end of eviction list.", CacheKey);
-      } else
-        // This should never happen.
-        throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
-                              "Program not found in the eviction list.");
+      }
+      // else: This can happen if concurrently the program is removed from
+      // eviction list by another thread.
     }
 
     bool empty() { return MProgramEvictionList.empty(); }
 
+    size_t size() { return MProgramEvictionList.size(); }
+
     void popFront() {
       if (!MProgramEvictionList.empty()) {
         MProgramToEvictionListMap.erase(MProgramEvictionList.front());
@@ -294,6 +295,10 @@ class KernelProgramCache {
     return {MKernelsPerProgramCache, MKernelsPerProgramCacheMutex};
   }
 
+  Locked<EvictionListT> acquireEvictionList() {
+    return {MEvictionList, MProgramEvictionListMutex};
+  }
+
   std::pair<ProgramBuildResultPtr, bool>
   getOrInsertProgram(const ProgramCacheKeyT &CacheKey) {
     auto LockedCache = acquireCachedPrograms();
@@ -368,13 +373,11 @@ class KernelProgramCache {
       ur_program_handle_t Program = std::get<3>(CacheVal);
       // Save kernel in fast cache only if the corresponding program is also
       // in the cache.
-      {
-        auto LockedCache = acquireCachedPrograms();
-        auto &ProgCache = LockedCache.get();
-        if (ProgCache.ProgramSizeMap.find(Program) ==
-            ProgCache.ProgramSizeMap.end())
-          return;
-      }
+      auto LockedCache = acquireCachedPrograms();
+      auto &ProgCache = LockedCache.get();
+      if (ProgCache.ProgramSizeMap.find(Program) ==
+          ProgCache.ProgramSizeMap.end())
+        return;
 
       // Save reference between the program and the fast cache key.
       MProgramToKernelFastCacheKeyMap[Program].emplace_back(CacheKey);
@@ -483,9 +486,8 @@ class KernelProgramCache {
                             const ur_program_handle_t &Program,
                             const bool IsBuilt) {
 
-    static size_t ProgramCacheEvictionThreshold = static_cast<size_t>(
-        SYCLConfig<
-            SYCL_IN_MEM_CACHE_EVICTION_THRESHOLD>::getProgramCacheSize());
+    size_t ProgramCacheEvictionThreshold =
+        SYCLConfig<SYCL_IN_MEM_CACHE_EVICTION_THRESHOLD>::getProgramCacheSize();
 
     // No need to populate the eviction list if eviction is disabled.
     if (ProgramCacheEvictionThreshold == 0)
@@ -561,7 +563,7 @@ class KernelProgramCache {
     MProgramToKernelFastCacheKeyMap.clear();
 
     // Clear the eviction lists and its mutexes.
-    std::lock_guard<std::mutex> L4(MProgramEvictionListMutex);
+    std::lock_guard<std::mutex> EvictionListLock(MProgramEvictionListMutex);
     MEvictionList.clear();
   }
 

sycl/unittests/kernel-and-program/CMakeLists.txt

@@ -7,5 +7,6 @@ add_sycl_unittest(KernelAndProgramTests OBJECT
   PersistentDeviceCodeCache.cpp
   KernelBuildOptions.cpp
   OutOfResources.cpp
+  InMemCacheEviction.cpp
 )
 target_compile_definitions(KernelAndProgramTests PRIVATE -D__SYCL_INTERNAL_API)

sycl/unittests/kernel-and-program/InMemCacheEviction.cpp

@@ -0,0 +1,225 @@
+//==----- InMemCacheEviction.cpp --- In-memory cache eviction 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
+//
+//===----------------------------------------------------------------------===//
+// This file contains tests covering eviction in in-memory program cache.
+
+#define SYCL2020_DISABLE_DEPRECATION_WARNINGS
+
+#include "../thread_safety/ThreadUtils.h"
+#include "detail/context_impl.hpp"
+#include "detail/kernel_program_cache.hpp"
+#include <detail/config.hpp>
+#include <helpers/MockDeviceImage.hpp>
+#include <helpers/MockKernelInfo.hpp>
+#include <helpers/UrMock.hpp>
+
+#include <gtest/gtest.h>
+
+#include <iostream>
+
+using namespace sycl;
+
+class Kernel1;
+class Kernel2;
+class Kernel3;
+
+MOCK_INTEGRATION_HEADER(Kernel1)
+MOCK_INTEGRATION_HEADER(Kernel2)
+MOCK_INTEGRATION_HEADER(Kernel3)
+
+static sycl::unittest::MockDeviceImage Img[] = {
+    sycl::unittest::generateDefaultImage({"Kernel1"}),
+    sycl::unittest::generateDefaultImage({"Kernel2"}),
+    sycl::unittest::generateDefaultImage({"Kernel3"})};
+
+static sycl::unittest::MockDeviceImageArray<3> ImgArray{Img};
+
+// Number of times urProgramCreateWithIL is called. This is used to check
+// if the program is created or fetched from the cache.
+static int NumProgramBuild = 0;
+
+constexpr int ProgramSize = 10000;
+
+static ur_result_t redefinedProgramCreateWithIL(void *) {
+  ++NumProgramBuild;
+  return UR_RESULT_SUCCESS;
+}
+
+static ur_result_t redefinedProgramGetInfoAfter(void *pParams) {
+  auto params = *static_cast<ur_program_get_info_params_t *>(pParams);
+  if (*params.ppropName == UR_PROGRAM_INFO_NUM_DEVICES) {
+    auto value = reinterpret_cast<unsigned int *>(*params.ppPropValue);
+    *value = 1;
+  }
+
+  if (*params.ppropName == UR_PROGRAM_INFO_BINARY_SIZES) {
+    auto value = reinterpret_cast<size_t *>(*params.ppPropValue);
+    value[0] = ProgramSize;
+  }
+
+  if (*params.ppropName == UR_PROGRAM_INFO_BINARIES) {
+    auto value = reinterpret_cast<unsigned char **>(*params.ppPropValue);
+    value[0] = 0;
+  }
+
+  return UR_RESULT_SUCCESS;
+}
+
+// Function to set SYCL_IN_MEM_CACHE_EVICTION_THRESHOLD.
+static void setCacheEvictionEnv(const char *value) {
+#ifdef _WIN32
+  _putenv_s("SYCL_IN_MEM_CACHE_EVICTION_THRESHOLD", value);
+#else
+  if (value)
+    setenv("SYCL_IN_MEM_CACHE_EVICTION_THRESHOLD", value, 1);
+  else
+    (void)unsetenv("SYCL_IN_MEM_CACHE_EVICTION_THRESHOLD");
+#endif
+
+  sycl::detail::readConfig(true);
+  sycl::detail::SYCLConfig<
+      sycl::detail::SYCL_IN_MEM_CACHE_EVICTION_THRESHOLD>::reset();
+}
+
+// Function to check number of entries in the cache and eviction list.
+static inline void
+CheckNumberOfEntriesInCacheAndEvictionList(detail::context_impl &CtxImpl,
+                                           size_t ExpectedNumEntries) {
+  auto &KPCache = CtxImpl.getKernelProgramCache();
+  EXPECT_EQ(KPCache.acquireCachedPrograms().get().size(), ExpectedNumEntries)
+      << "Unexpected number of entries in the cache";
+  auto EvcList = KPCache.acquireEvictionList();
+  EXPECT_EQ(EvcList.get().size(), ExpectedNumEntries)
+      << "Unexpected number of entries in the eviction list";
+}
+
+class InMemCacheEvictionTests : public ::testing::Test {
+protected:
+  void TearDown() override { setCacheEvictionEnv(""); }
+};
+
+TEST(InMemCacheEvictionTests, TestBasicEvictionAndLRU) {
+  NumProgramBuild = 0;
+  sycl::unittest::UrMock<> Mock;
+  mock::getCallbacks().set_before_callback("urProgramCreateWithIL",
+                                           &redefinedProgramCreateWithIL);
+  mock::getCallbacks().set_after_callback("urProgramGetInfo",
+                                          &redefinedProgramGetInfoAfter);
+
+  sycl::platform Plt{sycl::platform()};
+  sycl::context Ctx{Plt};
+  auto CtxImpl = detail::getSyclObjImpl(Ctx);
+  queue q(Ctx, default_selector_v);
+
+  // One program is of 10000 bytes, so 20005 eviction threshold can
+  // accommodate two program.
+  setCacheEvictionEnv("20005");
+
+  // Cache is empty, so one urProgramCreateWithIL call.
+  q.single_task<class Kernel1>([] {});
+  EXPECT_EQ(NumProgramBuild, 1);
+  CheckNumberOfEntriesInCacheAndEvictionList(*CtxImpl, 1);
+
+  q.single_task<class Kernel2>([] {});
+  EXPECT_EQ(NumProgramBuild, 2);
+  CheckNumberOfEntriesInCacheAndEvictionList(*CtxImpl, 2);
+
+  // Move first program to end of eviction list.
+  q.single_task<class Kernel1>([] {});
+  EXPECT_EQ(NumProgramBuild, 2);
+
+  // Calling Kernel3, Kernel2, and Kernel1 in a cyclic manner to
+  // verify LRU's working.
+
+  // Kernel2's program should have been evicted.
+  q.single_task<class Kernel3>([] {});
+  EXPECT_EQ(NumProgramBuild, 3);
+  CheckNumberOfEntriesInCacheAndEvictionList(*CtxImpl, 2);
+
+  // Calling Kernel2 again should trigger urProgramCreateWithIL and
+  // should evict Kernel1's program.
+  q.single_task<class Kernel2>([] {});
+  EXPECT_EQ(NumProgramBuild, 3);
+  CheckNumberOfEntriesInCacheAndEvictionList(*CtxImpl, 2);
+
+  // Calling Kernel1 again should trigger urProgramCreateWithIL and
+  // should evict Kernel3's program.
+  q.single_task<class Kernel1>([] {});
+  EXPECT_EQ(NumProgramBuild, 4);
+  CheckNumberOfEntriesInCacheAndEvictionList(*CtxImpl, 2);
+}
+
+// Test to verify eviction using concurrent kernel invocation.
+TEST(InMemCacheEvictionTests, TestConcurrentEvictionDifferentQueue) {
+  NumProgramBuild = 0;
+  sycl::unittest::UrMock<> Mock;
+  mock::getCallbacks().set_before_callback("urProgramCreateWithIL",
+                                           &redefinedProgramCreateWithIL);
+  mock::getCallbacks().set_after_callback("urProgramGetInfo",
+                                          &redefinedProgramGetInfoAfter);
+
+  sycl::platform Plt{sycl::platform()};
+  context Ctx{Plt};
+  auto CtxImpl = detail::getSyclObjImpl(Ctx);
+
+  // One program is of 10000 bytes, so 20005 eviction threshold can
+  // accommodate two program.
+  setCacheEvictionEnv("20005");
+
+  constexpr size_t ThreadCount = 100;
+  Barrier barrier(ThreadCount);
+  {
+    auto ConcurrentInvokeKernels = [&](std::size_t threadId) {
+      queue q(Ctx, default_selector_v);
+      barrier.wait();
+      q.single_task<class Kernel1>([] {});
+      q.single_task<class Kernel2>([] {});
+      q.single_task<class Kernel3>([] {});
+      q.wait_and_throw();
+    };
+
+    ThreadPool MPool(ThreadCount, ConcurrentInvokeKernels);
+  }
+
+  CheckNumberOfEntriesInCacheAndEvictionList(*CtxImpl, 2);
+}
+
+// Test to verify eviction using concurrent kernel invocation when
+// cache size is very less so as to trigger immediate eviction.
+TEST(InMemCacheEvictionTests, TestConcurrentEvictionSmallCache) {
+  NumProgramBuild = 0;
+  sycl::unittest::UrMock<> Mock;
+  mock::getCallbacks().set_before_callback("urProgramCreateWithIL",
+                                           &redefinedProgramCreateWithIL);
+  mock::getCallbacks().set_after_callback("urProgramGetInfo",
+                                          &redefinedProgramGetInfoAfter);
+
+  context Ctx{platform()};
+  auto CtxImpl = detail::getSyclObjImpl(Ctx);
+
+  // One program is of 10000 bytes, so 100 eviction threshold will
+  // trigger immediate eviction.
+  setCacheEvictionEnv("100");
+
+  // Fetch the same kernel concurrently from multiple threads.
+  // This should cause some threads to insert a program and other
+  // threads to evict the same program.
+  constexpr size_t ThreadCount = 300;
+  Barrier barrier(ThreadCount);
+  {
+    auto ConcurrentInvokeKernels = [&](std::size_t threadId) {
+      queue q(Ctx, default_selector_v);
+      barrier.wait();
+      q.single_task<class Kernel1>([] {});
+      q.wait_and_throw();
+    };
+
+    ThreadPool MPool(ThreadCount, ConcurrentInvokeKernels);
+  }
+
+  CheckNumberOfEntriesInCacheAndEvictionList(*CtxImpl, 0);
+}