diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index d6930633c2cd1..bf74779f69a77 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -12,6 +12,7 @@ #include #include #include +#include #include #include #include @@ -396,6 +397,53 @@ class kernel_bundle_impl { return SS.str(); } + bool + extKernelCompilerFetchFromCache(const std::vector Devices, + const std::vector &BuildOptions, + const std::string &SourceStr, + ur_program_handle_t &UrProgram) { + using ContextImplPtr = std::shared_ptr; + ContextImplPtr ContextImpl = getSyclObjImpl(MContext); + const AdapterPtr &Adapter = ContextImpl->getAdapter(); + + std::string UserArgs = syclex::detail::userArgsAsString(BuildOptions); + + std::vector DeviceHandles; + std::transform( + Devices.begin(), Devices.end(), std::back_inserter(DeviceHandles), + [](const device &Dev) { return getSyclObjImpl(Dev)->getHandleRef(); }); + + std::vector Binaries; + std::vector Lengths; + std::vector>> PersistentBinaries; + for (size_t i = 0; i < Devices.size(); i++) { + std::vector> BinProg = + PersistentDeviceCodeCache::getCompiledKernelFromDisc( + Devices[i], UserArgs, SourceStr); + + // exit if any device binary is missing + if (BinProg.empty()) { + return false; + } + PersistentBinaries.push_back(BinProg); + + Binaries.push_back((uint8_t *)(PersistentBinaries[i][0].data())); + Lengths.push_back(PersistentBinaries[i][0].size()); + } + + ur_program_properties_t Properties = {}; + Properties.stype = UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES; + Properties.pNext = nullptr; + Properties.count = 0; + Properties.pMetadatas = nullptr; + + Adapter->call( + ContextImpl->getHandleRef(), DeviceHandles.size(), DeviceHandles.data(), + Lengths.data(), Binaries.data(), &Properties, &UrProgram); + + return true; + } + std::shared_ptr build_from_source(const std::vector Devices, const std::vector &BuildOptions, @@ -415,57 +463,68 @@ class kernel_bundle_impl { DeviceVec.push_back(Dev); } - const auto spirv = [&]() -> std::vector { - 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(this->Source); - std::vector IPVersionVec(Devices.size()); - std::transform(DeviceVec.begin(), DeviceVec.end(), IPVersionVec.begin(), - [&](ur_device_handle_t d) { - uint32_t ipVersion = 0; - Adapter->call( - 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>(this->Source); - std::vector Result(SourceBytes.size()); - std::transform(SourceBytes.cbegin(), SourceBytes.cend(), Result.begin(), - [](std::byte B) { return static_cast(B); }); - return Result; - } - if (Language == syclex::source_language::sycl) { - const auto &SourceStr = std::get(this->Source); - return syclex::detail::SYCL_to_SPIRV(SourceStr, IncludePairs, - BuildOptions, LogPtr, - RegisteredKernelNames); - } - if (Language == syclex::source_language::sycl_jit) { - const auto &SourceStr = std::get(this->Source); - return syclex::detail::SYCL_JIT_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; - Adapter->call(ContextImpl->getHandleRef(), - spirv.data(), spirv.size(), - nullptr, &UrProgram); - // program created by urProgramCreateWithIL is implicitly retained. - if (UrProgram == nullptr) - throw sycl::exception( - sycl::make_error_code(errc::invalid), - "urProgramCreateWithIL resulted in a null program handle."); + // SourceStrPtr will be null when source is Spir-V bytes. + const std::string *SourceStrPtr = std::get_if(&this->Source); + bool FetchedFromCache = false; + if (PersistentDeviceCodeCache::isEnabled() && SourceStrPtr) { + FetchedFromCache = extKernelCompilerFetchFromCache( + Devices, BuildOptions, *SourceStrPtr, UrProgram); + } + + if (!FetchedFromCache) { + const auto spirv = [&]() -> std::vector { + if (Language == syclex::source_language::opencl) { + // if successful, the log is empty. if failed, throws an error with + // the compilation log. + std::vector IPVersionVec(Devices.size()); + std::transform(DeviceVec.begin(), DeviceVec.end(), + IPVersionVec.begin(), [&](ur_device_handle_t d) { + uint32_t ipVersion = 0; + Adapter->call( + 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>(this->Source); + std::vector Result(SourceBytes.size()); + std::transform(SourceBytes.cbegin(), SourceBytes.cend(), + Result.begin(), + [](std::byte B) { return static_cast(B); }); + return Result; + } + if (Language == syclex::source_language::sycl) { + return syclex::detail::SYCL_to_SPIRV(*SourceStrPtr, IncludePairs, + BuildOptions, LogPtr, + RegisteredKernelNames); + } + if (Language == syclex::source_language::sycl_jit) { + const auto &SourceStr = std::get(this->Source); + return syclex::detail::SYCL_JIT_to_SPIRV(SourceStr, 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( + ContextImpl->getHandleRef(), spirv.data(), spirv.size(), nullptr, + &UrProgram); + // program created by urProgramCreateWithIL is implicitly retained. + if (UrProgram == nullptr) + throw sycl::exception( + sycl::make_error_code(errc::invalid), + "urProgramCreateWithIL resulted in a null program handle."); + + } // if(!FetchedFromCache) std::string XsFlags = extractXsFlags(BuildOptions); auto Res = Adapter->call_nocheck( @@ -501,6 +560,17 @@ 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) { + for (const auto &Device : Devices) { + PersistentDeviceCodeCache::putCompiledKernelToDisc( + Device, syclex::detail::userArgsAsString(BuildOptions), + *SourceStrPtr, UrProgram); + } + } + return std::make_shared(MContext, MDevices, DevImg, KernelNames, Language); } diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 195ed3f1987d8..89a8a548a6d8a 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -36,6 +36,15 @@ SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs, throw sycl::exception(sycl::errc::build, "kernel_compiler does not support GCC<8"); } + +std::string userArgsAsString(const std::vector &UserArguments) { + return std::accumulate(UserArguments.begin(), UserArguments.end(), + std::string(""), + [](const std::string &A, const std::string &B) { + return A.empty() ? B : A + " " + B; + }); +} + } // namespace detail } // namespace ext::oneapi::experimental } // namespace _V1 diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp index 2d591cfb0913a..948f199c1b3cc 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp +++ b/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 &UserArguments); + spirv_vec_t SYCL_JIT_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs, const std::vector &UserArgs, std::string *LogPtr, diff --git a/sycl/source/detail/persistent_device_code_cache.cpp b/sycl/source/detail/persistent_device_code_cache.cpp index 0023b47c80174..469e4ceac2c8d 100644 --- a/sycl/source/detail/persistent_device_code_cache.cpp +++ b/sycl/source/detail/persistent_device_code_cache.cpp @@ -110,27 +110,22 @@ getSortedImages(const std::vector &Imgs) { return SortedImgs; } -/* Stores built program in persistent cache - */ -void PersistentDeviceCodeCache::putItemToDisc( - const device &Device, const std::vector &Imgs, - const SerializedObj &SpecConsts, const std::string &BuildOptionsString, - const ur_program_handle_t &NativePrg) { - - if (!areImagesCacheable(Imgs)) - return; - - std::vector 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> +getProgramBinaryData(const ur_program_handle_t &NativePrg, + const device &Device) { auto Adapter = detail::getSyclObjImpl(Device)->getAdapter(); - unsigned int DeviceNum = 0; - Adapter->call( NativePrg, UR_PROGRAM_INFO_NUM_DEVICES, sizeof(DeviceNum), &DeviceNum, nullptr); @@ -150,19 +145,34 @@ void PersistentDeviceCodeCache::putItemToDisc( Adapter->call( 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")); + return Result; +} + +/* Stores built program in persistent cache + */ +void PersistentDeviceCodeCache::putItemToDisc( + const device &Device, const std::vector &Imgs, + const SerializedObj &SpecConsts, const std::string &BuildOptionsString, + const ur_program_handle_t &NativePrg) { + + if (!areImagesCacheable(Imgs)) + return; + + std::vector 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)); + PersistentDeviceCodeCache::trace_KernelCompiler( + "binary has been cached: " + FullFileName); + } else { + PersistentDeviceCodeCache::trace_KernelCompiler("cache lock not owned " + + FileName); + } + } catch (std::exception &e) { + PersistentDeviceCodeCache::trace_KernelCompiler( + std::string("exception encountered making cache: ") + e.what()); + } catch (...) { + PersistentDeviceCodeCache::trace_KernelCompiler( + std::string("error outputting 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,43 @@ std::vector> PersistentDeviceCodeCache::getItemFromDisc( return {}; } +/* kernel_compiler extension uses slightly different format for path + and does not cache a .src separate from the binary. + */ +std::vector> +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> res = + readBinaryDataFromFile(FullFileName); + PersistentDeviceCodeCache::trace_KernelCompiler( + "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 +473,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 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() { diff --git a/sycl/source/detail/persistent_device_code_cache.hpp b/sycl/source/detail/persistent_device_code_cache.hpp index 868c247f28903..19b145f6de895 100644 --- a/sycl/source/detail/persistent_device_code_cache.hpp +++ b/sycl/source/detail/persistent_device_code_cache.hpp @@ -123,10 +123,6 @@ class PersistentDeviceCodeCache { const std::vector &SortedImgs, const SerializedObj &SpecConsts, const std::string &BuildOptionsString); - /* Check if on-disk cache enabled. - */ - static bool isEnabled(); - /* Returns the path to directory storing persistent device code cache.*/ static std::string getRootDir(); @@ -161,6 +157,10 @@ class PersistentDeviceCodeCache { 1024 * 1024 * 1024; public: + /* Check if on-disk cache enabled. + */ + static bool isEnabled(); + /* Get directory name for storing current cache item */ static std::string @@ -169,6 +169,14 @@ class PersistentDeviceCodeCache { const SerializedObj &SpecConsts, const std::string &BuildOptionsString); + /* Get directory name when storing runtime compiled kernels ( via + * kernel_compiler ). + */ + static std::string + getCompiledKernelItemPath(const device &Device, + const std::string &BuildOptionsString, + const std::string SourceString); + /* 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. @@ -179,6 +187,11 @@ class PersistentDeviceCodeCache { const SerializedObj &SpecConsts, const std::string &BuildOptionsString); + static std::vector> + getCompiledKernelFromDisc(const device &Device, + const std::string &BuildOptionsString, + const std::string SourceStr); + /* Stores build program in persistent cache */ static void @@ -188,6 +201,11 @@ class PersistentDeviceCodeCache { const std::string &BuildOptionsString, const ur_program_handle_t &NativePrg); + static void putCompiledKernelToDisc(const device &Device, + const std::string &BuildOptionsString, + const std::string &SourceStr, + 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 const bool traceEnabled = @@ -195,6 +213,12 @@ class PersistentDeviceCodeCache { if (traceEnabled) std::cerr << "[Persistent Cache]: " << msg << std::endl; } + static void trace_KernelCompiler(const std::string &msg) { + static const bool traceEnabled = + SYCLConfig::isTraceKernelCompiler(); + if (traceEnabled) + std::cerr << "[kernel_compiler Persistent Cache]: " << msg << std::endl; + } }; } // namespace detail } // namespace _V1 diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp index cfe2824ec0564..0fa13aece546f 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp @@ -6,11 +6,32 @@ // //===----------------------------------------------------------------------===// -// REQUIRES: (opencl || level_zero) +// 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 diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index 511f713b7c95c..26ca820558f66 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -9,8 +9,32 @@ // REQUIRES: (opencl || level_zero) // UNSUPPORTED: accelerator +// -- Test the kernel_compiler with SYCL source. // RUN: %{build} -o %t.out // RUN: %{run} %t.out +// RUN: %{l0_leak_check} %{run} %t.out + +// -- Test again, with caching. +// 'reading-from-cache' is just a string we pass to differentiate between the +// two runs. + +// DEFINE: %{cache_vars} = %{l0_leak_check} 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 reading-from-cache 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 reading-from-cache 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 @@ -115,7 +139,7 @@ void test_1(sycl::queue &Queue, sycl::kernel &Kernel, int seed) { sycl::free(usmPtr, Queue); } -void test_build_and_run() { +void test_build_and_run(bool readingFromCache) { namespace syclex = sycl::ext::oneapi::experimental; using source_kb = sycl::kernel_bundle; using exe_kb = sycl::kernel_bundle; @@ -157,8 +181,13 @@ void test_build_and_run() { kbSrc, devs, syclex::properties{syclex::build_options{flags}, syclex::save_log{&log}, syclex::registered_kernel_names{"ff_templated"}}); - assert(log.find("warning: 'this_nd_item<1>' is deprecated") != - std::string::npos); + + // If the kernel was restored from cache, there will not have been + // any warning issued by the compilation of the kernel. + if (!readingFromCache) { + assert(log.find("warning: 'this_nd_item<1>' is deprecated") != + std::string::npos); + } // clang-format off @@ -271,10 +300,19 @@ void test_esimd() { sycl::free(C, q); } -int main() { +int main(int argc, char *argv[]) { + bool readingFromCache = false; + + // Check if the argument is present + if (argc > 1) { + std::string argument(argv[1]); + if (argument == "reading-from-cache") { + readingFromCache = true; + } + } #ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER - test_build_and_run(); + test_build_and_run(readingFromCache); test_error(); test_esimd(); #else diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index 1491483781834..cc45096b8564c 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -9,8 +9,33 @@ // REQUIRES: (opencl || level_zero) // UNSUPPORTED: accelerator +// UNSUPPORTED: windows +// UNSUPPORTED-TRACKER: CMPLRLLVM-63166 +// in CMakeLists). + // RUN: %{build} -o %t.out // RUN: %{run} %t.out +// RUN: %{l0_leak_check} %{run} %t.out + +// -- Test again, with caching. + +// DEFINE: %{cache_vars} = %{l0_leak_check} 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