diff --git a/sycl/source/detail/persistent_device_code_cache.cpp b/sycl/source/detail/persistent_device_code_cache.cpp index 47adbf133e85c..b37f9100d0dbb 100644 --- a/sycl/source/detail/persistent_device_code_cache.cpp +++ b/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 &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> 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; diff --git a/sycl/source/detail/persistent_device_code_cache.hpp b/sycl/source/detail/persistent_device_code_cache.hpp index c51e5e55bc22b..48ef6e15b6fce 100644 --- a/sycl/source/detail/persistent_device_code_cache.hpp +++ b/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::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::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: diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp new file mode 100644 index 0000000000000..2340a6d96c06e --- /dev/null +++ b/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 +#include + +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; + + // 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() + << std::endl; + return; + } + + auto CreateAndVerifyKB = [](source_kb &kbSrc, + std::vector &&BuildFlags) { + std::string log; + std::vector 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 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; +}