From 5c5cf8f721d2c8d90a2f4bac67936cf8295c84ed Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 18 Sep 2024 11:11:52 -0700 Subject: [PATCH 01/22] initial commit, retrieval not working yet. several todos pending --- sycl/source/detail/kernel_bundle_impl.hpp | 30 +++- .../detail/persistent_device_code_cache.cpp | 139 ++++++++++++++++++ .../detail/persistent_device_code_cache.hpp | 18 +++ 3 files changed, 186 insertions(+), 1 deletion(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 2a128ba9a901e..b630580a6ec8f 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,15 @@ class kernel_bundle_impl { return SS.str(); } + // TODO: remove duplication in kernel_compiler_sycl.cpp. + 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; + }); + } + std::shared_ptr build_from_source(const std::vector Devices, const std::vector &BuildOptions, @@ -415,6 +425,18 @@ class kernel_bundle_impl { DeviceVec.push_back(Dev); } + ur_program_handle_t UrProgram = nullptr; + // bool FetchedFromCache = false; + // if (Language == syclex::source_language::sycl){ + // auto BinProg = PersistentDeviceCodeCache::getItemFromDisc(Device, + // AllImages, SpecConsts, CompileAndLinkOptions); if (!BinProg.empty()) { + // FetchedFromCache = true; + // UrProgram= createBinaryProgram(getSyclObjImpl(Context), Device, + // (const unsigned char *)BinProg[0].data(), + // BinProg[0].size(), ProgMetadataVector); + // } + // } + const auto spirv = [&]() -> std::vector { if (Language == syclex::source_language::opencl) { // if successful, the log is empty. if failed, throws an error with the @@ -451,7 +473,7 @@ class kernel_bundle_impl { "OpenCL C and SPIR-V are the only supported languages at this time"); }(); - ur_program_handle_t UrProgram = nullptr; + // CP ur_program_handle_t UrProgram = nullptr; Plugin->call(ContextImpl->getHandleRef(), spirv.data(), spirv.size(), nullptr, &UrProgram); @@ -495,6 +517,12 @@ class kernel_bundle_impl { nullptr, MContext, MDevices, bundle_state::executable, KernelIDs, UrProgram); device_image_plain DevImg{DevImgImpl}; + + // if we didn't get this from cache then... + const auto &SourceStr = std::get(this->Source); + PersistentDeviceCodeCache::putCompiledKernelToDisc( + Devices[0], userArgsAsString(BuildOptions), SourceStr, UrProgram); + return std::make_shared(MContext, MDevices, DevImg, KernelNames, Language); } diff --git a/sycl/source/detail/persistent_device_code_cache.cpp b/sycl/source/detail/persistent_device_code_cache.cpp index 7484ed00150e2..a0848a702906e 100644 --- a/sycl/source/detail/persistent_device_code_cache.cpp +++ b/sycl/source/detail/persistent_device_code_cache.cpp @@ -180,6 +180,85 @@ void PersistentDeviceCodeCache::putItemToDisc( } } +// TODO: unify this with putItemToDisc. Too much code duplication. +void PersistentDeviceCodeCache::putCompiledKernelToDisc( + const device &Device, const std::string &BuildOptionsString, + const std::string SourceStr, const ur_program_handle_t &NativePrg) { + + // Directory + std::string DirName = + getCompiledKernelItemPath(Device, BuildOptionsString, SourceStr); + std::cout << "DirName: " << DirName << std::endl; + + // File + size_t i = 0; + std::string FileName; + do { + FileName = DirName + "/" + std::to_string(i++); + } while (OSUtil::isPathPresent(FileName + ".bin") || + OSUtil::isPathPresent(FileName + ".lock")); + + std::cout << "FileName: " << FileName << std::endl; + + // Number of Devices? + auto Plugin = detail::getSyclObjImpl(Device)->getPlugin(); + + std::vector> Result; + std::vector Pointers; + + try { + unsigned int DeviceNum = 0; + Plugin->call( + NativePrg, UR_PROGRAM_INFO_NUM_DEVICES, sizeof(DeviceNum), &DeviceNum, + nullptr); + std::cout << "DeviceNum: " << DeviceNum << std::endl; + + // Actual Data + std::vector BinarySizes(DeviceNum); + Plugin->call( + NativePrg, UR_PROGRAM_INFO_BINARY_SIZES, + sizeof(size_t) * BinarySizes.size(), BinarySizes.data(), nullptr); + + for (size_t I = 0; I < BinarySizes.size(); ++I) { + Result.emplace_back(BinarySizes[I]); + Pointers.push_back(Result[I].data()); + std::cout << "BinarySizes[" << I << "]: " << BinarySizes[I] << std::endl; + } + Plugin->call( + NativePrg, UR_PROGRAM_INFO_BINARIES, sizeof(char *) * Pointers.size(), + Pointers.data(), nullptr); + } catch (sycl::exception &e) { + PersistentDeviceCodeCache::trace( + std::string( + "exception when retrieving program info for persistent cache: ") + + e.what()); + return; + } + + // Write + try { + OSUtil::makeDir(DirName.c_str()); + LockCacheItem Lock{FileName}; + if (Lock.isOwned()) { + std::string FullFileName = FileName + ".bin"; + writeBinaryDataToFile(FullFileName, Result); + trace("device binary has been cached: " + FullFileName); + // writeSourceItem(FileName + ".src", Device, SortedImgs, SpecConsts, + // BuildOptionsString); + } else { + PersistentDeviceCodeCache::trace("cache lock not owned " + FileName); + } + } catch (std::exception &e) { + PersistentDeviceCodeCache::trace( + std::string("exception encountered making persistent cache: ") + + e.what()); + } catch (...) { + PersistentDeviceCodeCache::trace( + std::string("error outputting persistent 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 +301,47 @@ std::vector> PersistentDeviceCodeCache::getItemFromDisc( return {}; } +/* + */ +ur_program_handle_t PersistentDeviceCodeCache::getCompiledKernelFromDisc( + const device &Device, const std::string &BuildOptionsString, + const std::string SourceStr) { + std::cout << "getCompiledKernelFromDisc" << std::endl; + std::string DirName = + getCompiledKernelItemPath(Device, BuildOptionsString, SourceStr); + std::cout << " DirName: " << DirName << std::endl; + + /* + if (Path.empty() || !OSUtil::isPathPresent(Path)) + return {}; + + int i = 0; + + std::string FileName{Path + "/" + std::to_string(i)}; + while (OSUtil::isPathPresent(FileName + ".bin") || + OSUtil::isPathPresent(FileName + ".src")) { + + if (!LockCacheItem::isLocked(FileName) && + isCacheItemSrcEqual(FileName + ".src", Device, SortedImgs, SpecConsts, + BuildOptionsString)) { + try { + std::string FullFileName = FileName + ".bin"; + std::vector> res = + readBinaryDataFromFile(FullFileName); + trace("using cached device binary: " + FullFileName); + return res; // subject for NRVO + } catch (...) { + // If read was unsuccessfull try the next item + } + } + FileName = Path + "/" + 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 +516,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 e2b3c8f72c4da..605c94806670f 100644 --- a/sycl/source/detail/persistent_device_code_cache.hpp +++ b/sycl/source/detail/persistent_device_code_cache.hpp @@ -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 ur_program_handle_t + 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 char *TraceEnabled = SYCLConfig::get(); From 0d3ac02a470abba99d93ec35b0cf8921fe129b47 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 18 Sep 2024 18:32:46 -0700 Subject: [PATCH 02/22] s.b. working, but seeing an error ... investigating --- sycl/source/detail/kernel_bundle_impl.hpp | 157 ++++++++++-------- .../detail/persistent_device_code_cache.cpp | 44 +++-- .../detail/persistent_device_code_cache.hpp | 2 +- .../program_manager/program_manager.cpp | 4 +- 4 files changed, 123 insertions(+), 84 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index b630580a6ec8f..e37e513574881 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -426,71 +426,94 @@ class kernel_bundle_impl { } ur_program_handle_t UrProgram = nullptr; - // bool FetchedFromCache = false; - // if (Language == syclex::source_language::sycl){ - // auto BinProg = PersistentDeviceCodeCache::getItemFromDisc(Device, - // AllImages, SpecConsts, CompileAndLinkOptions); if (!BinProg.empty()) { - // FetchedFromCache = true; - // UrProgram= createBinaryProgram(getSyclObjImpl(Context), Device, - // (const unsigned char *)BinProg[0].data(), - // BinProg[0].size(), ProgMetadataVector); - // } - // } - - 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; - Plugin->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); + const auto &SourceStr = std::get(this->Source); + bool FetchedFromCache = false; + if (Language == syclex::source_language::sycl) { + auto BinProg = PersistentDeviceCodeCache::getCompiledKernelFromDisc( + Devices[0], userArgsAsString(BuildOptions), SourceStr); + if (!BinProg.empty()) { + ur_device_handle_t UrDevice = + getSyclObjImpl(Devices[0])->getHandleRef(); + ur_result_t BinaryStatus = UR_RESULT_SUCCESS; + ur_program_properties_t Properties = {}; + std::vector Metadata = {}; + Properties.stype = UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES; + Properties.pNext = nullptr; + Properties.count = Metadata.size(); + Properties.pMetadatas = Metadata.data(); + BinaryStatus = + Plugin->call_nocheck( + ContextImpl->getHandleRef(), UrDevice, BinProg[0].size(), + (const unsigned char *)BinProg[0].data(), &Properties, + &UrProgram); + + std::cout << " BinaryStatus: " << BinaryStatus << std::endl; + + if (BinaryStatus == UR_RESULT_SUCCESS) { + FetchedFromCache = true; + std::cout << "zOMG, fetched from cache!" << std::endl; + } } - throw sycl::exception( - make_error_code(errc::invalid), - "OpenCL C and SPIR-V are the only supported languages at this time"); - }(); - - // CP ur_program_handle_t UrProgram = nullptr; - Plugin->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."); - - std::string XsFlags = extractXsFlags(BuildOptions); - auto Res = Plugin->call_nocheck( - UrProgram, DeviceVec.size(), DeviceVec.data(), XsFlags.c_str()); - if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { - Res = Plugin->call_nocheck( - ContextImpl->getHandleRef(), UrProgram, XsFlags.c_str()); } - Plugin->checkUrResult(Res); + + 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. + 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; + Plugin->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); + } + throw sycl::exception(make_error_code(errc::invalid), + "OpenCL C and SPIR-V are the only supported " + "languages at this time"); + }(); + + // CP ur_program_handle_t UrProgram = nullptr; + Plugin->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."); + + std::string XsFlags = extractXsFlags(BuildOptions); + auto Res = Plugin->call_nocheck( + UrProgram, DeviceVec.size(), DeviceVec.data(), XsFlags.c_str()); + if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { + Res = Plugin->call_nocheck( + ContextImpl->getHandleRef(), UrProgram, XsFlags.c_str()); + } + Plugin->checkUrResult(Res); + + } // if(!FetchedFromCache) // Get the number of kernels in the program. size_t NumKernels; @@ -518,10 +541,10 @@ class kernel_bundle_impl { UrProgram); device_image_plain DevImg{DevImgImpl}; - // if we didn't get this from cache then... - const auto &SourceStr = std::get(this->Source); - PersistentDeviceCodeCache::putCompiledKernelToDisc( - Devices[0], userArgsAsString(BuildOptions), SourceStr, UrProgram); + if (!FetchedFromCache) { + PersistentDeviceCodeCache::putCompiledKernelToDisc( + Devices[0], userArgsAsString(BuildOptions), SourceStr, UrProgram); + } return std::make_shared(MContext, MDevices, DevImg, KernelNames, Language); diff --git a/sycl/source/detail/persistent_device_code_cache.cpp b/sycl/source/detail/persistent_device_code_cache.cpp index a0848a702906e..60cabc0f26eef 100644 --- a/sycl/source/detail/persistent_device_code_cache.cpp +++ b/sycl/source/detail/persistent_device_code_cache.cpp @@ -117,6 +117,8 @@ void PersistentDeviceCodeCache::putItemToDisc( const SerializedObj &SpecConsts, const std::string &BuildOptionsString, const ur_program_handle_t &NativePrg) { + std::cout << "putItemToDisc" << std::endl; + if (!areImagesCacheable(Imgs)) return; @@ -164,6 +166,8 @@ void PersistentDeviceCodeCache::putItemToDisc( std::string FullFileName = FileName + ".bin"; writeBinaryDataToFile(FullFileName, Result); trace("device binary has been cached: " + FullFileName); + std::cout << " device binary has been cached: " << FullFileName + << std::endl; writeSourceItem(FileName + ".src", Device, SortedImgs, SpecConsts, BuildOptionsString); } else { @@ -185,10 +189,12 @@ void PersistentDeviceCodeCache::putCompiledKernelToDisc( const device &Device, const std::string &BuildOptionsString, const std::string SourceStr, const ur_program_handle_t &NativePrg) { + std::cout << "putCompiledKernelToDisc" << std::endl; + // Directory std::string DirName = getCompiledKernelItemPath(Device, BuildOptionsString, SourceStr); - std::cout << "DirName: " << DirName << std::endl; + std::cout << " DirName: " << DirName << std::endl; // File size_t i = 0; @@ -198,7 +204,7 @@ void PersistentDeviceCodeCache::putCompiledKernelToDisc( } while (OSUtil::isPathPresent(FileName + ".bin") || OSUtil::isPathPresent(FileName + ".lock")); - std::cout << "FileName: " << FileName << std::endl; + std::cout << " FileName: " << FileName << std::endl; // Number of Devices? auto Plugin = detail::getSyclObjImpl(Device)->getPlugin(); @@ -211,7 +217,7 @@ void PersistentDeviceCodeCache::putCompiledKernelToDisc( Plugin->call( NativePrg, UR_PROGRAM_INFO_NUM_DEVICES, sizeof(DeviceNum), &DeviceNum, nullptr); - std::cout << "DeviceNum: " << DeviceNum << std::endl; + std::cout << " DeviceNum: " << DeviceNum << std::endl; // Actual Data std::vector BinarySizes(DeviceNum); @@ -222,7 +228,7 @@ void PersistentDeviceCodeCache::putCompiledKernelToDisc( for (size_t I = 0; I < BinarySizes.size(); ++I) { Result.emplace_back(BinarySizes[I]); Pointers.push_back(Result[I].data()); - std::cout << "BinarySizes[" << I << "]: " << BinarySizes[I] << std::endl; + std::cout << " BinarySizes[" << I << "]: " << BinarySizes[I] << std::endl; } Plugin->call( NativePrg, UR_PROGRAM_INFO_BINARIES, sizeof(char *) * Pointers.size(), @@ -243,6 +249,8 @@ void PersistentDeviceCodeCache::putCompiledKernelToDisc( std::string FullFileName = FileName + ".bin"; writeBinaryDataToFile(FullFileName, Result); trace("device binary has been cached: " + FullFileName); + std::cout << " kernel_compiler device binary has been cached: " + << FullFileName << std::endl; // writeSourceItem(FileName + ".src", Device, SortedImgs, SpecConsts, // BuildOptionsString); } else { @@ -267,6 +275,8 @@ std::vector> PersistentDeviceCodeCache::getItemFromDisc( const device &Device, const std::vector &Imgs, const SerializedObj &SpecConsts, const std::string &BuildOptionsString) { + std::cout << "getItemFromDisc" << std::endl; + if (!areImagesCacheable(Imgs)) return {}; @@ -291,6 +301,8 @@ std::vector> PersistentDeviceCodeCache::getItemFromDisc( std::vector> res = readBinaryDataFromFile(FullFileName); trace("using cached device binary: " + FullFileName); + std::cout << " using cached device binary: " << FullFileName + << std::endl; return res; // subject for NRVO } catch (...) { // If read was unsuccessfull try the next item @@ -303,7 +315,9 @@ std::vector> PersistentDeviceCodeCache::getItemFromDisc( /* */ -ur_program_handle_t PersistentDeviceCodeCache::getCompiledKernelFromDisc( +// ur_program_handle_t +std::vector> +PersistentDeviceCodeCache::getCompiledKernelFromDisc( const device &Device, const std::string &BuildOptionsString, const std::string SourceStr) { std::cout << "getCompiledKernelFromDisc" << std::endl; @@ -311,37 +325,37 @@ ur_program_handle_t PersistentDeviceCodeCache::getCompiledKernelFromDisc( getCompiledKernelItemPath(Device, BuildOptionsString, SourceStr); std::cout << " DirName: " << DirName << std::endl; - /* - if (Path.empty() || !OSUtil::isPathPresent(Path)) + if (DirName.empty() || !OSUtil::isPathPresent(DirName)) return {}; int i = 0; - std::string FileName{Path + "/" + std::to_string(i)}; + std::string FileName{DirName + "/" + std::to_string(i)}; while (OSUtil::isPathPresent(FileName + ".bin") || OSUtil::isPathPresent(FileName + ".src")) { - if (!LockCacheItem::isLocked(FileName) && - isCacheItemSrcEqual(FileName + ".src", Device, SortedImgs, SpecConsts, - BuildOptionsString)) { + if (!LockCacheItem::isLocked( + FileName)) //&& + // isCacheItemSrcEqual(FileName + ".src", Device, + // SortedImgs, SpecConsts, BuildOptionsString)) + { try { std::string FullFileName = FileName + ".bin"; std::vector> res = readBinaryDataFromFile(FullFileName); trace("using cached device binary: " + FullFileName); + std::cout << " kernel_compiler using cached device binary: " + << FullFileName << std::endl; return res; // subject for NRVO } catch (...) { // If read was unsuccessfull try the next item } } - FileName = Path + "/" + std::to_string(++i); + 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) { diff --git a/sycl/source/detail/persistent_device_code_cache.hpp b/sycl/source/detail/persistent_device_code_cache.hpp index 605c94806670f..ba50aeebf6cf0 100644 --- a/sycl/source/detail/persistent_device_code_cache.hpp +++ b/sycl/source/detail/persistent_device_code_cache.hpp @@ -187,7 +187,7 @@ class PersistentDeviceCodeCache { const SerializedObj &SpecConsts, const std::string &BuildOptionsString); - static ur_program_handle_t + static std::vector> // ur_program_handle_t getCompiledKernelFromDisc(const device &Device, const std::string &BuildOptionsString, const std::string SourceStr); diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index d6f063e5fada6..a86b2af5c38d7 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -95,9 +95,11 @@ createBinaryProgram(const ContextImplPtr Context, const device &Device, Properties.pNext = nullptr; Properties.count = Metadata.size(); Properties.pMetadatas = Metadata.data(); - Plugin->call( + BinaryStatus = Plugin->call_nocheck( Context->getHandleRef(), UrDevice, DataLen, Data, &Properties, &Program); + std::cout << "createBinaryProgram BinaryStatus: " << BinaryStatus + << std::endl; if (BinaryStatus != UR_RESULT_SUCCESS) { throw detail::set_ur_error( exception(make_error_code(errc::runtime), From e4da64dde970a2fa15d18d01fbcc4ac59ae2e41b Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 24 Sep 2024 10:38:54 -0700 Subject: [PATCH 03/22] guards --- sycl/source/detail/kernel_bundle_impl.hpp | 5 +++-- sycl/source/detail/persistent_device_code_cache.hpp | 8 ++++---- 2 files changed, 7 insertions(+), 6 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index e37e513574881..8e585e0425434 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -428,7 +428,7 @@ class kernel_bundle_impl { ur_program_handle_t UrProgram = nullptr; const auto &SourceStr = std::get(this->Source); bool FetchedFromCache = false; - if (Language == syclex::source_language::sycl) { + if (PersistentDeviceCodeCache::isEnabled()) { auto BinProg = PersistentDeviceCodeCache::getCompiledKernelFromDisc( Devices[0], userArgsAsString(BuildOptions), SourceStr); if (!BinProg.empty()) { @@ -541,7 +541,8 @@ class kernel_bundle_impl { UrProgram); device_image_plain DevImg{DevImgImpl}; - if (!FetchedFromCache) { + // If caching enabled and kernel not fetched from cache, cache. + if (PersistentDeviceCodeCache::isEnabled() && !FetchedFromCache) { PersistentDeviceCodeCache::putCompiledKernelToDisc( Devices[0], userArgsAsString(BuildOptions), SourceStr, UrProgram); } diff --git a/sycl/source/detail/persistent_device_code_cache.hpp b/sycl/source/detail/persistent_device_code_cache.hpp index ba50aeebf6cf0..1e275dce701bf 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 From 30a11bbc57cc2e4223cffc15774c75af124a143d Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 25 Sep 2024 16:46:42 -0700 Subject: [PATCH 04/22] zomg working --- sycl/source/detail/kernel_bundle_impl.hpp | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 8e585e0425434..1bdba817e5025 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -452,6 +452,14 @@ class kernel_bundle_impl { if (BinaryStatus == UR_RESULT_SUCCESS) { FetchedFromCache = true; std::cout << "zOMG, fetched from cache!" << std::endl; + + ur_result_t Error = + Plugin->call_nocheck( + UrProgram, + /*num devices =*/1, &UrDevice, + userArgsAsString(BuildOptions).c_str()); + + std::cout << "error? " << Error << std::endl; } } } @@ -534,6 +542,11 @@ class kernel_bundle_impl { std::vector KernelNames = detail::split_string(KernelNamesStr, ';'); + // CP + std::cout << "KernelNames: " << KernelNamesStr << std::endl; + + //} // if(!FetchedFromCache) + // make the device image and the kernel_bundle_impl auto KernelIDs = std::make_shared>(); auto DevImgImpl = std::make_shared( @@ -547,6 +560,8 @@ class kernel_bundle_impl { Devices[0], userArgsAsString(BuildOptions), SourceStr, UrProgram); } + // std::vector KernelNames = { "__sycl_kernel_ff_cp" }; + return std::make_shared(MContext, MDevices, DevImg, KernelNames, Language); } From 12e12a7f2f03a6ea4366287ac7e56fe3cfaa2e25 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 26 Sep 2024 12:04:03 -0700 Subject: [PATCH 05/22] code reorg, step 1 --- sycl/source/detail/kernel_bundle_impl.hpp | 86 ++++++++++--------- .../detail/persistent_device_code_cache.cpp | 21 ++--- 2 files changed, 51 insertions(+), 56 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 1bdba817e5025..7c2e128896733 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -406,6 +406,50 @@ class kernel_bundle_impl { }); } + 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 PluginPtr &Plugin = ContextImpl->getPlugin(); + + std::string UserArgs = userArgsAsString(BuildOptions); + auto BinProg = PersistentDeviceCodeCache::getCompiledKernelFromDisc( + Devices[0], UserArgs, SourceStr); + if (!BinProg.empty()) { + ur_device_handle_t UrDevice = getSyclObjImpl(Devices[0])->getHandleRef(); + ur_result_t BinaryStatus = UR_RESULT_SUCCESS; + ur_program_properties_t Properties = {}; + std::vector Metadata = {}; + Properties.stype = UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES; + Properties.pNext = nullptr; + Properties.count = Metadata.size(); + Properties.pMetadatas = Metadata.data(); + BinaryStatus = Plugin->call_nocheck( + ContextImpl->getHandleRef(), UrDevice, BinProg[0].size(), + (const unsigned char *)BinProg[0].data(), &Properties, &UrProgram); + + std::cout << " BinaryStatus: " << BinaryStatus << std::endl; + + if (BinaryStatus == UR_RESULT_SUCCESS) { + std::cout << "zOMG, fetched from cache!" << std::endl; + + ur_result_t Error = Plugin->call_nocheck( + UrProgram, + /*num devices =*/1, &UrDevice, UserArgs.c_str()); + + std::cout << "error? " << Error << std::endl; + if (Error == UR_RESULT_SUCCESS) { + + return true; + } + } + } + return false; + } + std::shared_ptr build_from_source(const std::vector Devices, const std::vector &BuildOptions, @@ -429,39 +473,8 @@ class kernel_bundle_impl { const auto &SourceStr = std::get(this->Source); bool FetchedFromCache = false; if (PersistentDeviceCodeCache::isEnabled()) { - auto BinProg = PersistentDeviceCodeCache::getCompiledKernelFromDisc( - Devices[0], userArgsAsString(BuildOptions), SourceStr); - if (!BinProg.empty()) { - ur_device_handle_t UrDevice = - getSyclObjImpl(Devices[0])->getHandleRef(); - ur_result_t BinaryStatus = UR_RESULT_SUCCESS; - ur_program_properties_t Properties = {}; - std::vector Metadata = {}; - Properties.stype = UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES; - Properties.pNext = nullptr; - Properties.count = Metadata.size(); - Properties.pMetadatas = Metadata.data(); - BinaryStatus = - Plugin->call_nocheck( - ContextImpl->getHandleRef(), UrDevice, BinProg[0].size(), - (const unsigned char *)BinProg[0].data(), &Properties, - &UrProgram); - - std::cout << " BinaryStatus: " << BinaryStatus << std::endl; - - if (BinaryStatus == UR_RESULT_SUCCESS) { - FetchedFromCache = true; - std::cout << "zOMG, fetched from cache!" << std::endl; - - ur_result_t Error = - Plugin->call_nocheck( - UrProgram, - /*num devices =*/1, &UrDevice, - userArgsAsString(BuildOptions).c_str()); - - std::cout << "error? " << Error << std::endl; - } - } + FetchedFromCache = extKernelCompilerFetchFromCache(Devices, BuildOptions, + SourceStr, UrProgram); } if (!FetchedFromCache) { @@ -542,11 +555,6 @@ class kernel_bundle_impl { std::vector KernelNames = detail::split_string(KernelNamesStr, ';'); - // CP - std::cout << "KernelNames: " << KernelNamesStr << std::endl; - - //} // if(!FetchedFromCache) - // make the device image and the kernel_bundle_impl auto KernelIDs = std::make_shared>(); auto DevImgImpl = std::make_shared( @@ -560,8 +568,6 @@ class kernel_bundle_impl { Devices[0], userArgsAsString(BuildOptions), SourceStr, UrProgram); } - // std::vector KernelNames = { "__sycl_kernel_ff_cp" }; - return std::make_shared(MContext, MDevices, DevImg, KernelNames, Language); } diff --git a/sycl/source/detail/persistent_device_code_cache.cpp b/sycl/source/detail/persistent_device_code_cache.cpp index 60cabc0f26eef..10d0c41cff740 100644 --- a/sycl/source/detail/persistent_device_code_cache.cpp +++ b/sycl/source/detail/persistent_device_code_cache.cpp @@ -275,8 +275,6 @@ std::vector> PersistentDeviceCodeCache::getItemFromDisc( const device &Device, const std::vector &Imgs, const SerializedObj &SpecConsts, const std::string &BuildOptionsString) { - std::cout << "getItemFromDisc" << std::endl; - if (!areImagesCacheable(Imgs)) return {}; @@ -301,8 +299,6 @@ std::vector> PersistentDeviceCodeCache::getItemFromDisc( std::vector> res = readBinaryDataFromFile(FullFileName); trace("using cached device binary: " + FullFileName); - std::cout << " using cached device binary: " << FullFileName - << std::endl; return res; // subject for NRVO } catch (...) { // If read was unsuccessfull try the next item @@ -313,17 +309,16 @@ std::vector> PersistentDeviceCodeCache::getItemFromDisc( return {}; } -/* +/* kernel_compiler extension uses slightly different format for path + and does not cache a .src separate from the binary. */ -// ur_program_handle_t std::vector> PersistentDeviceCodeCache::getCompiledKernelFromDisc( const device &Device, const std::string &BuildOptionsString, const std::string SourceStr) { - std::cout << "getCompiledKernelFromDisc" << std::endl; + std::string DirName = getCompiledKernelItemPath(Device, BuildOptionsString, SourceStr); - std::cout << " DirName: " << DirName << std::endl; if (DirName.empty() || !OSUtil::isPathPresent(DirName)) return {}; @@ -334,18 +329,12 @@ PersistentDeviceCodeCache::getCompiledKernelFromDisc( while (OSUtil::isPathPresent(FileName + ".bin") || OSUtil::isPathPresent(FileName + ".src")) { - if (!LockCacheItem::isLocked( - FileName)) //&& - // isCacheItemSrcEqual(FileName + ".src", Device, - // SortedImgs, SpecConsts, BuildOptionsString)) - { + if (!LockCacheItem::isLocked(FileName)) { try { std::string FullFileName = FileName + ".bin"; std::vector> res = readBinaryDataFromFile(FullFileName); - trace("using cached device binary: " + FullFileName); - std::cout << " kernel_compiler using cached device binary: " - << FullFileName << std::endl; + trace("kernel_compiler using cached binary: " + FullFileName); return res; // subject for NRVO } catch (...) { // If read was unsuccessfull try the next item From 194b5182c952a58ec15241600a39c0faad657072 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 26 Sep 2024 18:10:54 -0700 Subject: [PATCH 06/22] more reorg, testing --- sycl/source/detail/kernel_bundle_impl.hpp | 6 - .../detail/persistent_device_code_cache.cpp | 127 ++++++------------ .../detail/persistent_device_code_cache.hpp | 2 +- .../KernelCompiler/kernel_compiler_opencl.cpp | 17 ++- .../KernelCompiler/kernel_compiler_sycl.cpp | 41 +++++- 5 files changed, 92 insertions(+), 101 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 7c2e128896733..aabb3acd6adb4 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -431,18 +431,12 @@ class kernel_bundle_impl { ContextImpl->getHandleRef(), UrDevice, BinProg[0].size(), (const unsigned char *)BinProg[0].data(), &Properties, &UrProgram); - std::cout << " BinaryStatus: " << BinaryStatus << std::endl; - if (BinaryStatus == UR_RESULT_SUCCESS) { - std::cout << "zOMG, fetched from cache!" << std::endl; - ur_result_t Error = Plugin->call_nocheck( UrProgram, /*num devices =*/1, &UrDevice, UserArgs.c_str()); - std::cout << "error? " << Error << std::endl; if (Error == UR_RESULT_SUCCESS) { - return true; } } diff --git a/sycl/source/detail/persistent_device_code_cache.cpp b/sycl/source/detail/persistent_device_code_cache.cpp index 10d0c41cff740..212614c18aaea 100644 --- a/sycl/source/detail/persistent_device_code_cache.cpp +++ b/sycl/source/detail/persistent_device_code_cache.cpp @@ -110,29 +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) { - - std::cout << "putItemToDisc" << std::endl; - - 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 Plugin = detail::getSyclObjImpl(Device)->getPlugin(); - unsigned int DeviceNum = 0; - Plugin->call( NativePrg, UR_PROGRAM_INFO_NUM_DEVICES, sizeof(DeviceNum), &DeviceNum, nullptr); @@ -152,22 +145,35 @@ void PersistentDeviceCodeCache::putItemToDisc( Plugin->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); - std::cout << " device binary has been cached: " << FullFileName - << std::endl; writeSourceItem(FileName + ".src", Device, SortedImgs, SpecConsts, BuildOptionsString); } else { @@ -184,75 +190,22 @@ void PersistentDeviceCodeCache::putItemToDisc( } } -// TODO: unify this with putItemToDisc. Too much code duplication. void PersistentDeviceCodeCache::putCompiledKernelToDisc( const device &Device, const std::string &BuildOptionsString, - const std::string SourceStr, const ur_program_handle_t &NativePrg) { + const std::string &SourceStr, const ur_program_handle_t &NativePrg) { - std::cout << "putCompiledKernelToDisc" << std::endl; - - // Directory std::string DirName = getCompiledKernelItemPath(Device, BuildOptionsString, SourceStr); - std::cout << " DirName: " << DirName << std::endl; - - // File - size_t i = 0; - std::string FileName; - do { - FileName = DirName + "/" + std::to_string(i++); - } while (OSUtil::isPathPresent(FileName + ".bin") || - OSUtil::isPathPresent(FileName + ".lock")); - std::cout << " FileName: " << FileName << std::endl; - - // Number of Devices? - auto Plugin = detail::getSyclObjImpl(Device)->getPlugin(); - - std::vector> Result; - std::vector Pointers; - - try { - unsigned int DeviceNum = 0; - Plugin->call( - NativePrg, UR_PROGRAM_INFO_NUM_DEVICES, sizeof(DeviceNum), &DeviceNum, - nullptr); - std::cout << " DeviceNum: " << DeviceNum << std::endl; - - // Actual Data - std::vector BinarySizes(DeviceNum); - Plugin->call( - NativePrg, UR_PROGRAM_INFO_BINARY_SIZES, - sizeof(size_t) * BinarySizes.size(), BinarySizes.data(), nullptr); - - for (size_t I = 0; I < BinarySizes.size(); ++I) { - Result.emplace_back(BinarySizes[I]); - Pointers.push_back(Result[I].data()); - std::cout << " BinarySizes[" << I << "]: " << BinarySizes[I] << std::endl; - } - Plugin->call( - NativePrg, UR_PROGRAM_INFO_BINARIES, sizeof(char *) * Pointers.size(), - Pointers.data(), nullptr); - } catch (sycl::exception &e) { - PersistentDeviceCodeCache::trace( - std::string( - "exception when retrieving program info for persistent cache: ") + - e.what()); - return; - } - - // Write 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); - trace("device binary has been cached: " + FullFileName); - std::cout << " kernel_compiler device binary has been cached: " - << FullFileName << std::endl; - // writeSourceItem(FileName + ".src", Device, SortedImgs, SpecConsts, - // BuildOptionsString); + writeBinaryDataToFile(FullFileName, + getProgramBinaryData(NativePrg, Device)); + trace("kernel_compiler binary has been cached: " + FullFileName); } else { PersistentDeviceCodeCache::trace("cache lock not owned " + FileName); } diff --git a/sycl/source/detail/persistent_device_code_cache.hpp b/sycl/source/detail/persistent_device_code_cache.hpp index 1e275dce701bf..ddcd488e3f5ae 100644 --- a/sycl/source/detail/persistent_device_code_cache.hpp +++ b/sycl/source/detail/persistent_device_code_cache.hpp @@ -203,7 +203,7 @@ class PersistentDeviceCodeCache { static void putCompiledKernelToDisc(const device &Device, const std::string &BuildOptionsString, - const std::string SourceStr, + const std::string &SourceStr, const ur_program_handle_t &NativePrg); /* Sends message to std:cerr stream when SYCL_CACHE_TRACE environemnt is set*/ diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp index cfe2824ec0564..8a9b977e0b171 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp @@ -9,8 +9,23 @@ // REQUIRES: (opencl || level_zero) // UNSUPPORTED: accelerator +// The leak check env var is set even if it might be ignored by some backends. + // RUN: %{build} -o %t.out -// RUN: %{run} %t.out +// RUN: env UR_L0_LEAKS_DEBUG=1 %{run} %t.out + +// DEFINE: %{cache_vars} = env UR_L0_LEAKS_DEBUG=1 env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=1 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 + +// CHECK-WRITTEN-TO-CACHE: Code caching: enabled +// CHECK-WRITTEN-TO-CACHE-NOT: *** Code caching: kernel_compiler using cached binary +// CHECK-WRITTEN-TO-CACHE: *** Code caching: kernel_compiler binary has been cached + +// CHECK-READ-FROM-CACHE: *** Code caching: enabled +// CHECK-READ-FROM-CACHE-NOT: *** Code caching: kernel_compiler binary has been cached +// CHECK-READ-FROM-CACHE: *** Code caching: kernel_compiler 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..cb3d3d9e43918 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -9,8 +9,25 @@ // REQUIRES: (opencl || level_zero) // UNSUPPORTED: accelerator +// The leak check env var is set even if it might be ignored by some backends. + // RUN: %{build} -o %t.out -// RUN: %{run} %t.out +// RUN: env UR_L0_LEAKS_DEBUG=1 %{run} %t.out + +// 'reading-from-cache' is just a string we pass to differentiate between the +// two runs. +// DEFINE: %{cache_vars} = env UR_L0_LEAKS_DEBUG=1 env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=1 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 + +// CHECK-WRITTEN-TO-CACHE: Code caching: enabled +// CHECK-WRITTEN-TO-CACHE-NOT: *** Code caching: kernel_compiler using cached binary +// CHECK-WRITTEN-TO-CACHE: *** Code caching: kernel_compiler binary has been cached + +// CHECK-READ-FROM-CACHE: *** Code caching: enabled +// CHECK-READ-FROM-CACHE-NOT: *** Code caching: kernel_compiler binary has been cached +// CHECK-READ-FROM-CACHE: *** Code caching: kernel_compiler using cached binary #include #include @@ -115,7 +132,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 +174,11 @@ 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 (!readingFromCache) { + assert(log.find("warning: 'this_nd_item<1>' is deprecated") != + std::string::npos); + } // clang-format off @@ -271,10 +291,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 From 8844e203af21a214ee05040277094772b012cb31 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 27 Sep 2024 10:30:00 -0700 Subject: [PATCH 07/22] last minute clean up --- sycl/source/detail/kernel_bundle_impl.hpp | 15 +++------------ .../kernel_compiler/kernel_compiler_sycl.hpp | 2 ++ .../KernelCompiler/kernel_compiler_opencl.cpp | 2 ++ .../KernelCompiler/kernel_compiler_sycl.cpp | 3 +++ 4 files changed, 10 insertions(+), 12 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index aabb3acd6adb4..8c18984b73bb0 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -397,15 +397,6 @@ class kernel_bundle_impl { return SS.str(); } - // TODO: remove duplication in kernel_compiler_sycl.cpp. - 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; - }); - } - bool extKernelCompilerFetchFromCache(const std::vector Devices, const std::vector &BuildOptions, @@ -415,7 +406,7 @@ class kernel_bundle_impl { ContextImplPtr ContextImpl = getSyclObjImpl(MContext); const PluginPtr &Plugin = ContextImpl->getPlugin(); - std::string UserArgs = userArgsAsString(BuildOptions); + std::string UserArgs = syclex::detail::userArgsAsString(BuildOptions); auto BinProg = PersistentDeviceCodeCache::getCompiledKernelFromDisc( Devices[0], UserArgs, SourceStr); if (!BinProg.empty()) { @@ -509,7 +500,6 @@ class kernel_bundle_impl { "languages at this time"); }(); - // CP ur_program_handle_t UrProgram = nullptr; Plugin->call( ContextImpl->getHandleRef(), spirv.data(), spirv.size(), nullptr, &UrProgram); @@ -559,7 +549,8 @@ class kernel_bundle_impl { // If caching enabled and kernel not fetched from cache, cache. if (PersistentDeviceCodeCache::isEnabled() && !FetchedFromCache) { PersistentDeviceCodeCache::putCompiledKernelToDisc( - Devices[0], userArgsAsString(BuildOptions), SourceStr, UrProgram); + Devices[0], syclex::detail::userArgsAsString(BuildOptions), SourceStr, + UrProgram); } return std::make_shared(MContext, MDevices, DevImg, diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp index dfff9ac839e84..4c3c3822eea9d 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); + } // namespace detail } // namespace ext::oneapi::experimental diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp index 8a9b977e0b171..a1d40334db28a 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp @@ -11,9 +11,11 @@ // The leak check env var is set even if it might be ignored by some backends. +// -- Test the kernel_compiler with OpenCL source. // RUN: %{build} -o %t.out // RUN: env UR_L0_LEAKS_DEBUG=1 %{run} %t.out +// -- Test again, with caching. // DEFINE: %{cache_vars} = env UR_L0_LEAKS_DEBUG=1 env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=1 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 diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index cb3d3d9e43918..bfc7e9ba0813d 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -11,11 +11,14 @@ // The leak check env var is set even if it might be ignored by some backends. +// -- Test the kernel_compiler with SYCL source. // RUN: %{build} -o %t.out // RUN: env UR_L0_LEAKS_DEBUG=1 %{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} = env UR_L0_LEAKS_DEBUG=1 env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=1 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 From 4e4b93e8d9ab06a1d3943d8896f588a196af39b4 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Sat, 28 Sep 2024 10:12:51 -0700 Subject: [PATCH 08/22] fix for spirv --- sycl/source/detail/kernel_bundle_impl.hpp | 28 ++++++++++++----------- 1 file changed, 15 insertions(+), 13 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 33614f9c8b795..fbc3caac655db 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -455,11 +455,12 @@ class kernel_bundle_impl { } ur_program_handle_t UrProgram = nullptr; - const auto &SourceStr = std::get(this->Source); + // 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()) { - FetchedFromCache = extKernelCompilerFetchFromCache(Devices, BuildOptions, - SourceStr, UrProgram); + if (PersistentDeviceCodeCache::isEnabled() && SourceStrPtr) { + FetchedFromCache = extKernelCompilerFetchFromCache( + Devices, BuildOptions, *SourceStrPtr, UrProgram); } if (!FetchedFromCache) { @@ -467,7 +468,6 @@ class kernel_bundle_impl { 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) { @@ -477,7 +477,7 @@ class kernel_bundle_impl { &ipVersion, nullptr); return ipVersion; }); - return syclex::detail::OpenCLC_to_SPIRV(SourceStr, IPVersionVec, + return syclex::detail::OpenCLC_to_SPIRV(*SourceStrPtr, IPVersionVec, BuildOptions, LogPtr); } if (Language == syclex::source_language::spirv) { @@ -490,13 +490,14 @@ class kernel_bundle_impl { return Result; } if (Language == syclex::source_language::sycl) { - return syclex::detail::SYCL_to_SPIRV(SourceStr, IncludePairs, + return syclex::detail::SYCL_to_SPIRV(*SourceStrPtr, 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"); + 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(), @@ -546,10 +547,11 @@ class kernel_bundle_impl { device_image_plain DevImg{DevImgImpl}; // If caching enabled and kernel not fetched from cache, cache. - if (PersistentDeviceCodeCache::isEnabled() && !FetchedFromCache) { + if (PersistentDeviceCodeCache::isEnabled() && !FetchedFromCache && + SourceStrPtr) { PersistentDeviceCodeCache::putCompiledKernelToDisc( - Devices[0], syclex::detail::userArgsAsString(BuildOptions), SourceStr, - UrProgram); + Devices[0], syclex::detail::userArgsAsString(BuildOptions), + *SourceStrPtr, UrProgram); } return std::make_shared(MContext, MDevices, DevImg, From ae2e98ae3c8668be26a6a428770cb7c5050f6651 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Sat, 28 Sep 2024 12:40:54 -0700 Subject: [PATCH 09/22] forgot to qualify leak check for windows --- .../KernelCompiler/kernel_compiler_opencl.cpp | 12 ++++++++---- .../test-e2e/KernelCompiler/kernel_compiler_sycl.cpp | 12 ++++++++---- 2 files changed, 16 insertions(+), 8 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp index a1d40334db28a..02625971949ce 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp @@ -9,18 +9,22 @@ // REQUIRES: (opencl || level_zero) // UNSUPPORTED: accelerator -// The leak check env var is set even if it might be ignored by some backends. - // -- Test the kernel_compiler with OpenCL source. // RUN: %{build} -o %t.out -// RUN: env UR_L0_LEAKS_DEBUG=1 %{run} %t.out +// RUN: %{run} %t.out +// RUN: %{l0_leak_check} %{run} %t.out // -- Test again, with caching. -// DEFINE: %{cache_vars} = env UR_L0_LEAKS_DEBUG=1 env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=1 SYCL_CACHE_DIR=%t/cache_dir +// DEFINE: %{cache_vars} = env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=1 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: Code caching: enabled // CHECK-WRITTEN-TO-CACHE-NOT: *** Code caching: kernel_compiler using cached binary // CHECK-WRITTEN-TO-CACHE: *** Code caching: kernel_compiler binary has been cached diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index bfc7e9ba0813d..9699571e24104 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -9,21 +9,25 @@ // REQUIRES: (opencl || level_zero) // UNSUPPORTED: accelerator -// The leak check env var is set even if it might be ignored by some backends. - // -- Test the kernel_compiler with SYCL source. // RUN: %{build} -o %t.out -// RUN: env UR_L0_LEAKS_DEBUG=1 %{run} %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} = env UR_L0_LEAKS_DEBUG=1 env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=1 SYCL_CACHE_DIR=%t/cache_dir +// DEFINE: %{cache_vars} = %{l0_leak_check} env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=1 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: Code caching: enabled // CHECK-WRITTEN-TO-CACHE-NOT: *** Code caching: kernel_compiler using cached binary // CHECK-WRITTEN-TO-CACHE: *** Code caching: kernel_compiler binary has been cached From 126e39a3d5f93ffb24ffa40bbbcc5cc2b512478d Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Sat, 28 Sep 2024 13:00:53 -0700 Subject: [PATCH 10/22] all OCL tests should declare that ocloc dependency --- sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp index 02625971949ce..5841742f99ddb 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -// REQUIRES: (opencl || level_zero) +// REQUIRES: ocloc && (opencl || level_zero) // UNSUPPORTED: accelerator // -- Test the kernel_compiler with OpenCL source. From 01fe513a4add98de1767440dc25ccf306b8ccc8b Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 1 Oct 2024 11:45:34 -0700 Subject: [PATCH 11/22] GCC<8 support --- .../detail/kernel_compiler/kernel_compiler_sycl.cpp | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 1e3106255b9c5..d348f7f97ad5d 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 From 19c1f08504b7169439858c5c1a946564eb02d018 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 1 Oct 2024 16:26:08 -0700 Subject: [PATCH 12/22] reviewer feedback --- sycl/source/detail/kernel_bundle_impl.hpp | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index fbc3caac655db..1c19484b7a361 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -413,11 +413,10 @@ class kernel_bundle_impl { ur_device_handle_t UrDevice = getSyclObjImpl(Devices[0])->getHandleRef(); ur_result_t BinaryStatus = UR_RESULT_SUCCESS; ur_program_properties_t Properties = {}; - std::vector Metadata = {}; Properties.stype = UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES; Properties.pNext = nullptr; - Properties.count = Metadata.size(); - Properties.pMetadatas = Metadata.data(); + Properties.count = 0; + Properties.pMetadatas = nullptr; BinaryStatus = Adapter->call_nocheck( ContextImpl->getHandleRef(), UrDevice, BinProg[0].size(), (const unsigned char *)BinProg[0].data(), &Properties, &UrProgram); @@ -427,9 +426,7 @@ class kernel_bundle_impl { UrProgram, /*num devices =*/1, &UrDevice, UserArgs.c_str()); - if (Error == UR_RESULT_SUCCESS) { - return true; - } + return (Error == UR_RESULT_SUCCESS); } } return false; From a5711851e74b2bce832ea529426a0035ec7be9e5 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 23 Oct 2024 12:20:42 -0700 Subject: [PATCH 13/22] multiple devices supported for kernel_compiler caching --- sycl/source/detail/kernel_bundle_impl.hpp | 39 ++++++++++++++--------- 1 file changed, 24 insertions(+), 15 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 1c19484b7a361..d52472d327475 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -410,24 +410,31 @@ class kernel_bundle_impl { auto BinProg = PersistentDeviceCodeCache::getCompiledKernelFromDisc( Devices[0], UserArgs, SourceStr); if (!BinProg.empty()) { - ur_device_handle_t UrDevice = getSyclObjImpl(Devices[0])->getHandleRef(); - ur_result_t BinaryStatus = UR_RESULT_SUCCESS; + 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; + for (size_t i = 0; i < Devices.size(); i++) { + auto BinProg = PersistentDeviceCodeCache::getCompiledKernelFromDisc( + Devices[i], UserArgs, SourceStr); + Binaries.push_back((uint8_t *)(BinProg.data())); + Lengths.push_back(BinProg.size()); + } + ur_program_properties_t Properties = {}; Properties.stype = UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES; Properties.pNext = nullptr; Properties.count = 0; Properties.pMetadatas = nullptr; - BinaryStatus = Adapter->call_nocheck( - ContextImpl->getHandleRef(), UrDevice, BinProg[0].size(), - (const unsigned char *)BinProg[0].data(), &Properties, &UrProgram); - - if (BinaryStatus == UR_RESULT_SUCCESS) { - ur_result_t Error = Adapter->call_nocheck( - UrProgram, - /*num devices =*/1, &UrDevice, UserArgs.c_str()); - return (Error == UR_RESULT_SUCCESS); - } + Adapter->call( + ContextImpl->getHandleRef(), DeviceHandles.size(), + DeviceHandles.data(), Lengths.data(), Binaries.data(), &Properties, + &UrProgram); } return false; } @@ -546,9 +553,11 @@ class kernel_bundle_impl { // If caching enabled and kernel not fetched from cache, cache. if (PersistentDeviceCodeCache::isEnabled() && !FetchedFromCache && SourceStrPtr) { - PersistentDeviceCodeCache::putCompiledKernelToDisc( - Devices[0], syclex::detail::userArgsAsString(BuildOptions), - *SourceStrPtr, UrProgram); + for (const auto &Device : Devices) { + PersistentDeviceCodeCache::putCompiledKernelToDisc( + Device, syclex::detail::userArgsAsString(BuildOptions), + *SourceStrPtr, UrProgram); + } } return std::make_shared(MContext, MDevices, DevImg, From 49a3828a51bb242062825f70325c5442611f9df2 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 25 Oct 2024 10:34:21 -0700 Subject: [PATCH 14/22] multiple devices --- sycl/source/detail/kernel_bundle_impl.hpp | 25 +++++++++++++++-------- 1 file changed, 16 insertions(+), 9 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index d52472d327475..f6bcd148e5998 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -407,9 +407,7 @@ class kernel_bundle_impl { const AdapterPtr &Adapter = ContextImpl->getAdapter(); std::string UserArgs = syclex::detail::userArgsAsString(BuildOptions); - auto BinProg = PersistentDeviceCodeCache::getCompiledKernelFromDisc( - Devices[0], UserArgs, SourceStr); - if (!BinProg.empty()) { + std::vector DeviceHandles; std::transform(Devices.begin(), Devices.end(), std::back_inserter(DeviceHandles), [](const device &Dev) { @@ -418,11 +416,20 @@ class kernel_bundle_impl { std::vector Binaries; std::vector Lengths; + std::vector>> PersistentBinaries; for (size_t i = 0; i < Devices.size(); i++) { - auto BinProg = PersistentDeviceCodeCache::getCompiledKernelFromDisc( - Devices[i], UserArgs, SourceStr); - Binaries.push_back((uint8_t *)(BinProg.data())); - Lengths.push_back(BinProg.size()); + 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 *)(BinProg[0].data())); + Lengths.push_back(BinProg[0].size()); } ur_program_properties_t Properties = {}; @@ -435,8 +442,8 @@ class kernel_bundle_impl { ContextImpl->getHandleRef(), DeviceHandles.size(), DeviceHandles.data(), Lengths.data(), Binaries.data(), &Properties, &UrProgram); - } - return false; + + return true; } std::shared_ptr From 9841c43eb2748a2a7861a4b7377a3363600d61a1 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 25 Oct 2024 13:52:40 -0700 Subject: [PATCH 15/22] clang-format fighting with itself? --- sycl/source/detail/kernel_bundle_impl.hpp | 100 +++++++++--------- .../detail/persistent_device_code_cache.cpp | 6 +- 2 files changed, 52 insertions(+), 54 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index f6bcd148e5998..5f46885291a06 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -408,42 +408,40 @@ class kernel_bundle_impl { 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 *)(BinProg[0].data())); - Lengths.push_back(BinProg[0].size()); + 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); - ur_program_properties_t Properties = {}; - Properties.stype = UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES; - Properties.pNext = nullptr; - Properties.count = 0; - Properties.pMetadatas = nullptr; + Binaries.push_back((uint8_t *)(BinProg[0].data())); + Lengths.push_back(BinProg[0].size()); + } - Adapter->call( - ContextImpl->getHandleRef(), DeviceHandles.size(), - DeviceHandles.data(), Lengths.data(), Binaries.data(), &Properties, - &UrProgram); + ur_program_properties_t Properties = {}; + Properties.stype = UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES; + Properties.pNext = nullptr; + Properties.count = 0; + Properties.pMetadatas = nullptr; - return true; + Adapter->call( + ContextImpl->getHandleRef(), DeviceHandles.size(), DeviceHandles.data(), + Lengths.data(), Binaries.data(), &Properties, &UrProgram); + + return true; } std::shared_ptr @@ -480,8 +478,8 @@ class kernel_bundle_impl { // 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) { + 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), @@ -511,23 +509,23 @@ class kernel_bundle_impl { "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."); - - std::string XsFlags = extractXsFlags(BuildOptions); - auto Res = Adapter->call_nocheck( - UrProgram, DeviceVec.size(), DeviceVec.data(), XsFlags.c_str()); - if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { - Res = Adapter->call_nocheck( - ContextImpl->getHandleRef(), UrProgram, XsFlags.c_str()); - } - Adapter->checkUrResult(Res); + 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."); + + std::string XsFlags = extractXsFlags(BuildOptions); + auto Res = Adapter->call_nocheck( + UrProgram, DeviceVec.size(), DeviceVec.data(), XsFlags.c_str()); + if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { + Res = Adapter->call_nocheck( + ContextImpl->getHandleRef(), UrProgram, XsFlags.c_str()); + } + Adapter->checkUrResult(Res); } // if(!FetchedFromCache) diff --git a/sycl/source/detail/persistent_device_code_cache.cpp b/sycl/source/detail/persistent_device_code_cache.cpp index 406693574e359..fc4be8694275b 100644 --- a/sycl/source/detail/persistent_device_code_cache.cpp +++ b/sycl/source/detail/persistent_device_code_cache.cpp @@ -142,9 +142,9 @@ getProgramBinaryData(const ur_program_handle_t &NativePrg, Pointers.push_back(Result[I].data()); } - Adapter->call(NativePrg, UR_PROGRAM_INFO_BINARIES, - sizeof(char *) * Pointers.size(), - Pointers.data(), nullptr); + Adapter->call( + NativePrg, UR_PROGRAM_INFO_BINARIES, sizeof(char *) * Pointers.size(), + Pointers.data(), nullptr); return Result; } From 3340a2313f2c358b955d1cecfa33c756a642c846 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 29 Oct 2024 09:13:25 -0700 Subject: [PATCH 16/22] finish multi-device support --- sycl/source/detail/kernel_bundle_impl.hpp | 23 ++++++++++++----------- 1 file changed, 12 insertions(+), 11 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 5829075d3d3d2..c8cf4bab56acd 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -427,8 +427,8 @@ class kernel_bundle_impl { } PersistentBinaries.push_back(BinProg); - Binaries.push_back((uint8_t *)(BinProg[0].data())); - Lengths.push_back(BinProg[0].size()); + Binaries.push_back((uint8_t *)(PersistentBinaries[i][0].data())); + Lengths.push_back(PersistentBinaries[i][0].size()); } ur_program_properties_t Properties = {}; @@ -524,17 +524,18 @@ class kernel_bundle_impl { sycl::make_error_code(errc::invalid), "urProgramCreateWithIL resulted in a null program handle."); - std::string XsFlags = extractXsFlags(BuildOptions); - auto Res = Adapter->call_nocheck( - UrProgram, DeviceVec.size(), DeviceVec.data(), XsFlags.c_str()); - if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { - Res = Adapter->call_nocheck( - ContextImpl->getHandleRef(), UrProgram, XsFlags.c_str()); - } - Adapter->checkUrResult(Res); - } // if(!FetchedFromCache) + std::string XsFlags = extractXsFlags(BuildOptions); + auto Res = Adapter->call_nocheck( + UrProgram, DeviceVec.size(), DeviceVec.data(), XsFlags.c_str()); + if (Res == + UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { // (Res != UR_RESULT_SUCCESS) { + Res = Adapter->call_nocheck( + ContextImpl->getHandleRef(), UrProgram, XsFlags.c_str()); + } + Adapter->checkUrResult(Res); + // Get the number of kernels in the program. size_t NumKernels; Adapter->call( From e5af452e90d7883e84f0e141a30b2dce61120381 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 29 Oct 2024 10:12:00 -0700 Subject: [PATCH 17/22] updated for new cache_trace env, added cache testing to sycl_jit --- .../detail/persistent_device_code_cache.cpp | 19 +++++++++-------- .../detail/persistent_device_code_cache.hpp | 6 ++++++ .../KernelCompiler/kernel_compiler_opencl.cpp | 14 ++++++------- .../KernelCompiler/kernel_compiler_sycl.cpp | 14 ++++++------- .../kernel_compiler_sycl_jit.cpp | 21 +++++++++++++++++++ 5 files changed, 51 insertions(+), 23 deletions(-) diff --git a/sycl/source/detail/persistent_device_code_cache.cpp b/sycl/source/detail/persistent_device_code_cache.cpp index fc4be8694275b..469e4ceac2c8d 100644 --- a/sycl/source/detail/persistent_device_code_cache.cpp +++ b/sycl/source/detail/persistent_device_code_cache.cpp @@ -205,18 +205,18 @@ void PersistentDeviceCodeCache::putCompiledKernelToDisc( std::string FullFileName = FileName + ".bin"; writeBinaryDataToFile(FullFileName, getProgramBinaryData(NativePrg, Device)); - trace("kernel_compiler binary has been cached: " + FullFileName); + PersistentDeviceCodeCache::trace_KernelCompiler( + "binary has been cached: " + FullFileName); } else { - PersistentDeviceCodeCache::trace("cache lock not owned " + FileName); + PersistentDeviceCodeCache::trace_KernelCompiler("cache lock not owned " + + FileName); } } catch (std::exception &e) { - PersistentDeviceCodeCache::trace( - std::string("exception encountered making persistent cache: ") + - e.what()); + PersistentDeviceCodeCache::trace_KernelCompiler( + std::string("exception encountered making cache: ") + e.what()); } catch (...) { - PersistentDeviceCodeCache::trace( - std::string("error outputting persistent cache: ") + - std::strerror(errno)); + PersistentDeviceCodeCache::trace_KernelCompiler( + std::string("error outputting cache: ") + std::strerror(errno)); } } @@ -287,7 +287,8 @@ PersistentDeviceCodeCache::getCompiledKernelFromDisc( std::string FullFileName = FileName + ".bin"; std::vector> res = readBinaryDataFromFile(FullFileName); - trace("kernel_compiler using cached binary: " + FullFileName); + PersistentDeviceCodeCache::trace_KernelCompiler( + "using cached binary: " + FullFileName); return res; // subject for NRVO } catch (...) { // If read was unsuccessfull try the next item diff --git a/sycl/source/detail/persistent_device_code_cache.hpp b/sycl/source/detail/persistent_device_code_cache.hpp index 08b0611f62c54..19b145f6de895 100644 --- a/sycl/source/detail/persistent_device_code_cache.hpp +++ b/sycl/source/detail/persistent_device_code_cache.hpp @@ -213,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 5841742f99ddb..0fa13aece546f 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp @@ -15,7 +15,7 @@ // RUN: %{l0_leak_check} %{run} %t.out // -- Test again, with caching. -// DEFINE: %{cache_vars} = env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=1 SYCL_CACHE_DIR=%t/cache_dir +// 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 @@ -25,13 +25,13 @@ // 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: Code caching: enabled -// CHECK-WRITTEN-TO-CACHE-NOT: *** Code caching: kernel_compiler using cached binary -// CHECK-WRITTEN-TO-CACHE: *** Code caching: kernel_compiler binary has been cached +// 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: *** Code caching: enabled -// CHECK-READ-FROM-CACHE-NOT: *** Code caching: kernel_compiler binary has been cached -// CHECK-READ-FROM-CACHE: *** Code caching: kernel_compiler using cached binary +// 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 9699571e24104..5c08a904b97c6 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -18,7 +18,7 @@ // '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=1 SYCL_CACHE_DIR=%t/cache_dir +// 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 @@ -28,13 +28,13 @@ // 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: Code caching: enabled -// CHECK-WRITTEN-TO-CACHE-NOT: *** Code caching: kernel_compiler using cached binary -// CHECK-WRITTEN-TO-CACHE: *** Code caching: kernel_compiler binary has been cached +// 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: *** Code caching: enabled -// CHECK-READ-FROM-CACHE-NOT: *** Code caching: kernel_compiler binary has been cached -// CHECK-READ-FROM-CACHE: *** Code caching: kernel_compiler using cached binary +// 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 diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index 1491483781834..6b7b039a8f819 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -11,6 +11,27 @@ // 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 From a57c6650e3b66d7358cb3f0ac2aac9ac7e667ba1 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 29 Oct 2024 10:20:35 -0700 Subject: [PATCH 18/22] comments and cleanup --- sycl/source/detail/kernel_bundle_impl.hpp | 3 +-- sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp | 2 ++ 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index c8cf4bab56acd..f1355020738dc 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -529,8 +529,7 @@ class kernel_bundle_impl { std::string XsFlags = extractXsFlags(BuildOptions); auto Res = Adapter->call_nocheck( UrProgram, DeviceVec.size(), DeviceVec.data(), XsFlags.c_str()); - if (Res == - UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { // (Res != UR_RESULT_SUCCESS) { + if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { Res = Adapter->call_nocheck( ContextImpl->getHandleRef(), UrProgram, XsFlags.c_str()); } diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index 5c08a904b97c6..26ca820558f66 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -182,6 +182,8 @@ void test_build_and_run(bool readingFromCache) { syclex::properties{syclex::build_options{flags}, syclex::save_log{&log}, syclex::registered_kernel_names{"ff_templated"}}); + // 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); From 5bbceb72ae5fd1007b77238a816b82083a8d4345 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 29 Oct 2024 10:24:25 -0700 Subject: [PATCH 19/22] clang-format can go f..ree itself --- sycl/source/detail/kernel_bundle_impl.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index f1355020738dc..bf74779f69a77 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -506,8 +506,8 @@ class kernel_bundle_impl { 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); + BuildOptions, LogPtr, + RegisteredKernelNames); } throw sycl::exception( make_error_code(errc::invalid), From 10d918b6c947f4891d4338fed0bfc528aacf1e1d Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 30 Oct 2024 10:21:55 -0700 Subject: [PATCH 20/22] update unsupported test --- sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index 6b7b039a8f819..9de4c004baec6 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -9,6 +9,9 @@ // REQUIRES: (opencl || level_zero) // UNSUPPORTED: accelerator +// TODO: enable on Windows (via SYCL_ENABLE_EXTENSION_JIT in CMakeLists). +// UNSUPPORTED: windows + // RUN: %{build} -o %t.out // RUN: %{run} %t.out // RUN: %{l0_leak_check} %{run} %t.out From dc7d2a79467e8dd52186d859aa75e32b4e1437f0 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 30 Oct 2024 12:18:38 -0700 Subject: [PATCH 21/22] unsupported-INTENDED --- sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index 9de4c004baec6..f9f249713b11b 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -9,8 +9,9 @@ // REQUIRES: (opencl || level_zero) // UNSUPPORTED: accelerator -// TODO: enable on Windows (via SYCL_ENABLE_EXTENSION_JIT in CMakeLists). // UNSUPPORTED: windows +// UNSUPPORTED-INTENDED: To Do: enable on Windows (via SYCL_ENABLE_EXTENSION_JIT +// in CMakeLists). // RUN: %{build} -o %t.out // RUN: %{run} %t.out From a3d9f96d52f668839d44c788954dc192807e47ad Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 1 Nov 2024 08:36:16 -0700 Subject: [PATCH 22/22] committed-tracker and other sins --- sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index f9f249713b11b..cc45096b8564c 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -10,7 +10,7 @@ // UNSUPPORTED: accelerator // UNSUPPORTED: windows -// UNSUPPORTED-INTENDED: To Do: enable on Windows (via SYCL_ENABLE_EXTENSION_JIT +// UNSUPPORTED-TRACKER: CMPLRLLVM-63166 // in CMakeLists). // RUN: %{build} -o %t.out