From 1e0bd3a3452abb6d09115937d6fb25f781cdef86 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Fri, 20 Dec 2024 15:57:16 -0800 Subject: [PATCH 1/5] Extend eviction to kernel_compiler cache --- .../detail/persistent_device_code_cache.cpp | 28 +++ .../kernel_compiler_cache_eviction.cpp | 192 ++++++++++++++++++ 2 files changed, 220 insertions(+) create mode 100644 sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp diff --git a/sycl/source/detail/persistent_device_code_cache.cpp b/sycl/source/detail/persistent_device_code_cache.cpp index 47adbf133e85c..f5f9e926a007a 100644 --- a/sycl/source/detail/persistent_device_code_cache.cpp +++ b/sycl/source/detail/persistent_device_code_cache.cpp @@ -316,6 +316,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; @@ -495,7 +497,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. @@ -513,6 +528,9 @@ void PersistentDeviceCodeCache::putCompiledKernelToDisc( writeBinaryDataToFile(FullFileName, BinaryData[DeviceIndex]); PersistentDeviceCodeCache::trace_KernelCompiler( "binary has been cached: " + FullFileName); + + TotalSize += getFileSize(FullFileName); + saveCurrentTimeInAFile(FileName + CacheEntryAccessTimeSuffix); } else { PersistentDeviceCodeCache::trace_KernelCompiler( "cache lock not owned " + FileName); @@ -525,6 +543,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 @@ -611,6 +633,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 (...) { 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..2465f3a6f1581 --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp @@ -0,0 +1,192 @@ +//==-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 +// +//===----------------------------------------------------------------------===// + +// REQUIRES: ocloc && (opencl || level_zero) +// UNSUPPORTED: accelerator + +// -- Test the kernel_compiler with OpenCL source. +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// RUN: %{l0_leak_check} %{run} %t.out + +// -- Test again, with caching. +// DEFINE: %{cache_vars} = env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=5 SYCL_CACHE_DIR=%t/cache_dir +// RUN: rm -rf %t/cache_dir +// RUN: %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-WRITTEN-TO-CACHE +// RUN: %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-READ-FROM-CACHE + +// -- Add leak check. +// RUN: rm -rf %t/cache_dir +// RUN: %{l0_leak_check} %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-WRITTEN-TO-CACHE +// RUN: %{l0_leak_check} %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-READ-FROM-CACHE + +// CHECK-WRITTEN-TO-CACHE: [Persistent Cache]: enabled +// CHECK-WRITTEN-TO-CACHE-NOT: [kernel_compiler Persistent Cache]: using cached binary +// CHECK-WRITTEN-TO-CACHE: [kernel_compiler Persistent Cache]: binary has been cached + +// CHECK-READ-FROM-CACHE: [Persistent Cache]: enabled +// CHECK-READ-FROM-CACHE-NOT: [kernel_compiler Persistent Cache]: binary has been cached +// CHECK-READ-FROM-CACHE: [kernel_compiler Persistent Cache]: using cached binary + +#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; +} +)==="; + +auto constexpr BadCLSource = R"===( +__kernel void my_kernel(__global int *in, __global int *out) { + size_t i = get_global_id(0) + no semi-colon!! + out[i] = in[i]*2 + 100; +} +)==="; +/* +Compile Log: +1:3:34: error: use of undeclared identifier 'no' + size_t i = get_global_id(0) + no semi-colon!! + ^ +1:3:36: error: expected ';' at end of declaration + size_t i = get_global_id(0) + no semi-colon!! + ^ + ; + +Build failed with error code: -11 + +============= + +*/ + +using namespace sycl; + +void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel, int multiplier, + int added) { + constexpr int N = 4; + cl_int InputArray[N] = {0, 1, 2, 3}; + cl_int OutputArray[N] = {}; + + sycl::buffer InputBuf(InputArray, sycl::range<1>(N)); + sycl::buffer OutputBuf(OutputArray, sycl::range<1>(N)); + + Q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, InputBuf.get_access(CGH)); + CGH.set_arg(1, OutputBuf.get_access(CGH)); + CGH.parallel_for(sycl::range<1>{N}, Kernel); + }); + + sycl::host_accessor Out{OutputBuf}; + for (int I = 0; I < N; I++) + assert(Out[I] == ((I * multiplier) + added)); +} + +void test_build_and_run() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_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; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::opencl, CLSource); + // compilation of empty prop list, no devices + exe_kb kbExe1 = syclex::build(kbSrc); + + // compilation with props and devices + std::string log; + std::vector flags{"-cl-fast-relaxed-math", + "-cl-finite-math-only"}; + std::vector devs = kbSrc.get_devices(); + sycl::context ctxRes = kbSrc.get_context(); + assert(ctxRes == ctx); + sycl::backend beRes = kbSrc.get_backend(); + assert(beRes == ctx.get_backend()); + + exe_kb kbExe2 = syclex::build( + kbSrc, devs, + syclex::properties{syclex::build_options{flags}, syclex::save_log{&log}}); + + bool hasMyKernel = kbExe2.ext_oneapi_has_kernel("my_kernel"); + bool hasHerKernel = kbExe2.ext_oneapi_has_kernel("her_kernel"); + bool notExistKernel = kbExe2.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, but does?"); + + sycl::kernel my_kernel = kbExe2.ext_oneapi_get_kernel("my_kernel"); + sycl::kernel her_kernel = kbExe2.ext_oneapi_get_kernel("her_kernel"); + + auto my_num_args = my_kernel.get_info(); + assert(my_num_args == 2 && "my_kernel should take 2 args"); + + testSyclKernel(q, my_kernel, 2, 100); + testSyclKernel(q, her_kernel, 5, 1000); +} + +void test_error() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_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) { + return; + } + + try { + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::opencl, BadCLSource); + exe_kb kbExe1 = syclex::build(kbSrc); + assert(false && "we should not be here."); + } catch (sycl::exception &e) { + // nice! + assert(e.code() == sycl::errc::build); + } + // any other error will escape and cause the test to fail ( as it should ). +} + +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(); + test_error(); +#else + static_assert(false, "Kernel Compiler feature test macro undefined"); +#endif + return 0; +} From 2f9599cfd7244e293701faea30dd29d2006c2b14 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Sat, 21 Dec 2024 10:41:54 -0800 Subject: [PATCH 2/5] Fix E2E test for kernel compiler --- .../detail/persistent_device_code_cache.cpp | 13 +- .../kernel_compiler_cache_eviction.cpp | 175 ++++++------------ 2 files changed, 69 insertions(+), 119 deletions(-) diff --git a/sycl/source/detail/persistent_device_code_cache.cpp b/sycl/source/detail/persistent_device_code_cache.cpp index f5f9e926a007a..4ee02bdce18f5 100644 --- a/sycl/source/detail/persistent_device_code_cache.cpp +++ b/sycl/source/detail/persistent_device_code_cache.cpp @@ -54,8 +54,7 @@ LockCacheItem::~LockCacheItem() { } // 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 +209,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)) { diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp index 2465f3a6f1581..e8273dfda529b 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp @@ -6,32 +6,20 @@ // //===----------------------------------------------------------------------===// +// Tests on-disk cache and eviction with kernel_compiler. + // REQUIRES: ocloc && (opencl || level_zero) // UNSUPPORTED: accelerator // -- Test the kernel_compiler with OpenCL source. // RUN: %{build} -o %t.out -// RUN: %{run} %t.out -// RUN: %{l0_leak_check} %{run} %t.out // -- Test again, with caching. -// DEFINE: %{cache_vars} = env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=5 SYCL_CACHE_DIR=%t/cache_dir -// RUN: rm -rf %t/cache_dir -// RUN: %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-WRITTEN-TO-CACHE -// RUN: %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-READ-FROM-CACHE - -// -- Add leak check. +// DEFINE: %{cache_vars} = env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=7 SYCL_CACHE_DIR=%t/cache_dir SYCL_CACHE_MAX_SIZE=23000 // RUN: rm -rf %t/cache_dir -// RUN: %{l0_leak_check} %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-WRITTEN-TO-CACHE -// RUN: %{l0_leak_check} %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-READ-FROM-CACHE +// RUN: %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefix=CHECK -// CHECK-WRITTEN-TO-CACHE: [Persistent Cache]: enabled -// CHECK-WRITTEN-TO-CACHE-NOT: [kernel_compiler Persistent Cache]: using cached binary -// CHECK-WRITTEN-TO-CACHE: [kernel_compiler Persistent Cache]: binary has been cached - -// CHECK-READ-FROM-CACHE: [Persistent Cache]: enabled -// CHECK-READ-FROM-CACHE-NOT: [kernel_compiler Persistent Cache]: binary has been cached -// CHECK-READ-FROM-CACHE: [kernel_compiler Persistent Cache]: using cached binary +// CHECK: [Persistent Cache]: enabled #include #include @@ -47,54 +35,11 @@ __kernel void her_kernel(__global int *in, __global int *out) { } )==="; -auto constexpr BadCLSource = R"===( -__kernel void my_kernel(__global int *in, __global int *out) { - size_t i = get_global_id(0) + no semi-colon!! - out[i] = in[i]*2 + 100; -} -)==="; -/* -Compile Log: -1:3:34: error: use of undeclared identifier 'no' - size_t i = get_global_id(0) + no semi-colon!! - ^ -1:3:36: error: expected ';' at end of declaration - size_t i = get_global_id(0) + no semi-colon!! - ^ - ; - -Build failed with error code: -11 - -============= - -*/ - using namespace sycl; -void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel, int multiplier, - int added) { - constexpr int N = 4; - cl_int InputArray[N] = {0, 1, 2, 3}; - cl_int OutputArray[N] = {}; - - sycl::buffer InputBuf(InputArray, sycl::range<1>(N)); - sycl::buffer OutputBuf(OutputArray, sycl::range<1>(N)); - - Q.submit([&](sycl::handler &CGH) { - CGH.set_arg(0, InputBuf.get_access(CGH)); - CGH.set_arg(1, OutputBuf.get_access(CGH)); - CGH.parallel_for(sycl::range<1>{N}, Kernel); - }); - - sycl::host_accessor Out{OutputBuf}; - for (int I = 0; I < N; I++) - assert(Out[I] == ((I * multiplier) + added)); -} - void test_build_and_run() { namespace syclex = sycl::ext::oneapi::experimental; using source_kb = sycl::kernel_bundle; - using exe_kb = sycl::kernel_bundle; // only one device is supported at this time, so we limit the queue and // context to that @@ -112,69 +57,66 @@ void test_build_and_run() { 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 of empty prop list, no devices - exe_kb kbExe1 = syclex::build(kbSrc); // compilation with props and devices - std::string log; std::vector flags{"-cl-fast-relaxed-math", "-cl-finite-math-only"}; - std::vector devs = kbSrc.get_devices(); - sycl::context ctxRes = kbSrc.get_context(); - assert(ctxRes == ctx); - sycl::backend beRes = kbSrc.get_backend(); - assert(beRes == ctx.get_backend()); - - exe_kb kbExe2 = syclex::build( - kbSrc, devs, - syclex::properties{syclex::build_options{flags}, syclex::save_log{&log}}); - - bool hasMyKernel = kbExe2.ext_oneapi_has_kernel("my_kernel"); - bool hasHerKernel = kbExe2.ext_oneapi_has_kernel("her_kernel"); - bool notExistKernel = kbExe2.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, but does?"); - - sycl::kernel my_kernel = kbExe2.ext_oneapi_get_kernel("my_kernel"); - sycl::kernel her_kernel = kbExe2.ext_oneapi_get_kernel("her_kernel"); - - auto my_num_args = my_kernel.get_info(); - assert(my_num_args == 2 && "my_kernel should take 2 args"); - - testSyclKernel(q, my_kernel, 2, 100); - testSyclKernel(q, her_kernel, 5, 1000); -} -void test_error() { - namespace syclex = sycl::ext::oneapi::experimental; - using source_kb = sycl::kernel_bundle; - using exe_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) { - return; - } - - try { - source_kb kbSrc = syclex::create_kernel_bundle_from_source( - ctx, syclex::source_language::opencl, BadCLSource); - exe_kb kbExe1 = syclex::build(kbSrc); - assert(false && "we should not be here."); - } catch (sycl::exception &e) { - // nice! - assert(e.code() == sycl::errc::build); - } - // any other error will escape and cause the test to fail ( as it should ). + // 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]}); + + // Device image #3 + // CHECK: [kernel_compiler Persistent Cache]: binary has been cached: [[DEVIMG3:.*]] + // CHECK: [Persistent Cache]: Updating the cache size file. + CreateAndVerifyKB(kbSrc, {flags[1]}); + + // Re-insert device image #1 + // CHECK: [kernel_compiler Persistent Cache]: using cached binary: [[DEVIMG1]] + CreateAndVerifyKB(kbSrc, {}); + + // Device image #4 + // CHECK: [kernel_compiler Persistent Cache]: binary has been cached: [[DEVIMG4:.*]] + // CHECK: [Persistent Cache]: Updating the cache size file. + // CHECK: [Persistent Cache]: Cache eviction triggered. + // CHECK: [Persistent Cache]: File removed: [[DEVIMG2]] + // CHECK: [Persistent Cache]: File removed: [[DEVIMG3]] + // CHECK: [Persistent Cache]: File removed: [[DEVIMG1]] + CreateAndVerifyKB(kbSrc, {flags[0], flags[1]}); + + // Re-insert device image #4 + // CHECK: [kernel_compiler Persistent Cache]: using cached binary: [[DEVIMG4]] + CreateAndVerifyKB(kbSrc, {flags[0], flags[1]}); } int main() { @@ -184,7 +126,6 @@ int main() { #ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER test_build_and_run(); - test_error(); #else static_assert(false, "Kernel Compiler feature test macro undefined"); #endif From c3603da8da31729b92c6c58446b73d0ee4b88733 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Sat, 21 Dec 2024 11:13:06 -0800 Subject: [PATCH 3/5] Fix E2E test for seperate run and build mode --- .../KernelCompiler/kernel_compiler_cache_eviction.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp index e8273dfda529b..94bec1854bfd4 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp @@ -10,14 +10,16 @@ // 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=23000 -// RUN: rm -rf %t/cache_dir -// RUN: %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefix=CHECK +// RUN: %if run-mode %{rm -rf %t/cache_dir%} +// RUN: %{cache_vars} %{run} %t.out 2>&1 | FileCheck %s --check-prefix=CHECK // CHECK: [Persistent Cache]: enabled From 83bd015c0adc4c559241e00a8d0ac8f7b30d1a9f Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Mon, 30 Dec 2024 11:09:10 -0800 Subject: [PATCH 4/5] Fix test failure --- .../kernel_compiler_cache_eviction.cpp | 35 ++++++++----------- 1 file changed, 14 insertions(+), 21 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp index 94bec1854bfd4..ad9135225d361 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp @@ -17,9 +17,9 @@ // 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=23000 +// DEFINE: %{cache_vars} = env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=7 SYCL_CACHE_DIR=%t/cache_dir SYCL_CACHE_MAX_SIZE=27000 // RUN: %if run-mode %{rm -rf %t/cache_dir%} -// RUN: %{cache_vars} %{run} %t.out 2>&1 | FileCheck %s --check-prefix=CHECK +// RUN: %{cache_vars} %{run-unfiltered-devices} %t.out 2>&1 | FileCheck %s --check-prefix=CHECK // CHECK: [Persistent Cache]: enabled @@ -84,7 +84,8 @@ void test_build_and_run() { // compilation with props and devices std::vector flags{"-cl-fast-relaxed-math", - "-cl-finite-math-only"}; + "-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. @@ -96,29 +97,21 @@ void test_build_and_run() { // 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]}); - - // Device image #3 - // CHECK: [kernel_compiler Persistent Cache]: binary has been cached: [[DEVIMG3:.*]] - // CHECK: [Persistent Cache]: Updating the cache size file. - CreateAndVerifyKB(kbSrc, {flags[1]}); + CreateAndVerifyKB(kbSrc, {flags[0], flags[1], flags[2], flags[3]}); // Re-insert device image #1 - // CHECK: [kernel_compiler Persistent Cache]: using cached binary: [[DEVIMG1]] + // CHECK-NEXT: [kernel_compiler Persistent Cache]: using cached binary: [[DEVIMG1]] CreateAndVerifyKB(kbSrc, {}); - // Device image #4 - // CHECK: [kernel_compiler Persistent Cache]: binary has been cached: [[DEVIMG4:.*]] - // CHECK: [Persistent Cache]: Updating the cache size file. + // 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: [Persistent Cache]: File removed: [[DEVIMG2]] - // CHECK: [Persistent Cache]: File removed: [[DEVIMG3]] - // CHECK: [Persistent Cache]: File removed: [[DEVIMG1]] - CreateAndVerifyKB(kbSrc, {flags[0], flags[1]}); - - // Re-insert device image #4 - // CHECK: [kernel_compiler Persistent Cache]: using cached binary: [[DEVIMG4]] - CreateAndVerifyKB(kbSrc, {flags[0], flags[1]}); + // 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() { From 56bb3dc11859cdbbe778747aef5b36c2f1a3d536 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Tue, 31 Dec 2024 13:44:39 -0800 Subject: [PATCH 5/5] Fix formatting of trace on Widnows --- .../detail/persistent_device_code_cache.cpp | 29 +++++++++---------- .../detail/persistent_device_code_cache.hpp | 18 ++++++++---- .../kernel_compiler_cache_eviction.cpp | 2 +- 3 files changed, 27 insertions(+), 22 deletions(-) diff --git a/sycl/source/detail/persistent_device_code_cache.cpp b/sycl/source/detail/persistent_device_code_cache.cpp index 4ee02bdce18f5..b37f9100d0dbb 100644 --- a/sycl/source/detail/persistent_device_code_cache.cpp +++ b/sycl/source/detail/persistent_device_code_cache.cpp @@ -49,8 +49,7 @@ 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. @@ -335,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; } }; @@ -475,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); @@ -485,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( @@ -536,13 +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( @@ -612,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; } @@ -660,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; } @@ -691,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 @@ -708,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 @@ -723,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 {}; } @@ -763,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); } } @@ -811,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 index ad9135225d361..2340a6d96c06e 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_cache_eviction.cpp @@ -17,7 +17,7 @@ // 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=27000 +// 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