From 6456fe89646deb8bf30c0eb32827a62ff6e58ffb Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Thu, 14 Nov 2024 12:41:45 -0800 Subject: [PATCH] [SYCL] Fix multi-device support for persistent cache (#16056) Associated UR PR: https://github.com/oneapi-src/unified-runtime/pull/2313 --- sycl/source/detail/kernel_bundle_impl.hpp | 31 +- .../detail/persistent_device_code_cache.cpp | 321 ++++++++++-------- .../detail/persistent_device_code_cache.hpp | 17 +- .../program_manager/program_manager.cpp | 28 +- .../persistent-cache-multi-device.cpp | 38 +++ .../PersistentDeviceCodeCache.cpp | 75 ++-- 6 files changed, 300 insertions(+), 210 deletions(-) create mode 100644 sycl/test-e2e/KernelAndProgram/persistent-cache-multi-device.cpp diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index bf74779f69a77..7ce5971711291 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -415,20 +415,15 @@ class kernel_bundle_impl { 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()); + std::vector> BinProgs = + PersistentDeviceCodeCache::getCompiledKernelFromDisc(Devices, UserArgs, + SourceStr); + if (BinProgs.empty()) { + return false; + } + for (auto &BinProg : BinProgs) { + Binaries.push_back((uint8_t *)(BinProg.data())); + Lengths.push_back(BinProg.size()); } ur_program_properties_t Properties = {}; @@ -564,11 +559,9 @@ class kernel_bundle_impl { // 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); - } + PersistentDeviceCodeCache::putCompiledKernelToDisc( + Devices, syclex::detail::userArgsAsString(BuildOptions), + *SourceStrPtr, UrProgram); } return std::make_shared(MContext, MDevices, DevImg, diff --git a/sycl/source/detail/persistent_device_code_cache.cpp b/sycl/source/detail/persistent_device_code_cache.cpp index 469e4ceac2c8d..205ebd7d42d26 100644 --- a/sycl/source/detail/persistent_device_code_cache.cpp +++ b/sycl/source/detail/persistent_device_code_cache.cpp @@ -121,37 +121,69 @@ std::string getUniqueFilename(const std::string &base_name) { return filename; } +/* Returns binary data for the UR program. There is a one-to-one + * correspondence between the vector of programs returned from the function and + * the input vector of devices. + */ std::vector> getProgramBinaryData(const ur_program_handle_t &NativePrg, - const device &Device) { - auto Adapter = detail::getSyclObjImpl(Device)->getAdapter(); + const std::vector &Devices) { + assert(!Devices.empty() && "At least one device is expected"); + // We expect all devices to be from the same platform/adpater. + auto Adapter = detail::getSyclObjImpl(Devices[0])->getAdapter(); unsigned int DeviceNum = 0; Adapter->call( NativePrg, UR_PROGRAM_INFO_NUM_DEVICES, sizeof(DeviceNum), &DeviceNum, nullptr); + std::vector URDevices(DeviceNum); + Adapter->call( + NativePrg, UR_PROGRAM_INFO_DEVICES, + sizeof(ur_device_handle_t) * URDevices.size(), URDevices.data(), nullptr); + std::vector BinarySizes(DeviceNum); Adapter->call( NativePrg, UR_PROGRAM_INFO_BINARY_SIZES, sizeof(size_t) * BinarySizes.size(), BinarySizes.data(), nullptr); - std::vector> Result; + std::vector> Binaries; std::vector Pointers; for (size_t I = 0; I < BinarySizes.size(); ++I) { - Result.emplace_back(BinarySizes[I]); - Pointers.push_back(Result[I].data()); + Binaries.emplace_back(BinarySizes[I]); + Pointers.push_back(Binaries[I].data()); } Adapter->call( NativePrg, UR_PROGRAM_INFO_BINARIES, sizeof(char *) * Pointers.size(), Pointers.data(), nullptr); + + // Select only binaries for the input devices preserving one to one + // correpsondence. + std::vector> Result(Devices.size()); + for (size_t DeviceIndex = 0; DeviceIndex < Devices.size(); DeviceIndex++) { + auto DeviceIt = std::find_if( + URDevices.begin(), URDevices.end(), + [&Devices, &DeviceIndex](const ur_device_handle_t &URDevice) { + return URDevice == + detail::getSyclObjImpl(Devices[DeviceIndex])->getHandleRef(); + }); + assert(DeviceIt != URDevices.end() && + "Device is not associated with the program"); + auto URDeviceIndex = std::distance(URDevices.begin(), DeviceIt); + Result[DeviceIndex] = std::move(Binaries[URDeviceIndex]); + } + + // Return binaries correpsonding to the input devices. + return Result; } -/* Stores built program in persistent cache +/* Stores built program in persistent cache. We will put the binary for each + * device in the list to a separate file. */ void PersistentDeviceCodeCache::putItemToDisc( - const device &Device, const std::vector &Imgs, + const std::vector &Devices, + const std::vector &Imgs, const SerializedObj &SpecConsts, const std::string &BuildOptionsString, const ur_program_handle_t &NativePrg) { @@ -159,107 +191,127 @@ void PersistentDeviceCodeCache::putItemToDisc( 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, - getProgramBinaryData(NativePrg, Device)); - trace("device binary has been cached: " + FullFileName); - writeSourceItem(FileName + ".src", Device, SortedImgs, SpecConsts, - BuildOptionsString); - } else { - PersistentDeviceCodeCache::trace("cache lock not owned " + FileName); + auto BinaryData = getProgramBinaryData(NativePrg, Devices); + for (size_t DeviceIndex = 0; DeviceIndex < Devices.size(); DeviceIndex++) { + // If we don't have binary for the device, skip it. + if (BinaryData[DeviceIndex].empty()) + continue; + std::string DirName = getCacheItemPath(Devices[DeviceIndex], 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, BinaryData[DeviceIndex]); + trace("device binary has been cached: " + FullFileName); + writeSourceItem(FileName + ".src", Devices[DeviceIndex], 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)); } - } 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)); } } void PersistentDeviceCodeCache::putCompiledKernelToDisc( - const device &Device, const std::string &BuildOptionsString, + const std::vector &Devices, 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)); + auto BinaryData = getProgramBinaryData(NativePrg, Devices); + + for (size_t DeviceIndex = 0; DeviceIndex < Devices.size(); DeviceIndex++) { + // If we don't have binary for the device, skip it. + if (BinaryData[DeviceIndex].empty()) + continue; + std::string DirName = getCompiledKernelItemPath( + Devices[DeviceIndex], 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, BinaryData[DeviceIndex]); + 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( - "binary has been cached: " + FullFileName); - } else { - PersistentDeviceCodeCache::trace_KernelCompiler("cache lock not owned " + - FileName); + std::string("error outputting cache: ") + std::strerror(errno)); } - } 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. + * stored in vector of chars. There is a one-to-one correspondence between + * the vector of programs returned from the function and the input vector of + * devices. */ std::vector> PersistentDeviceCodeCache::getItemFromDisc( - const device &Device, const std::vector &Imgs, + const std::vector &Devices, + const std::vector &Imgs, const SerializedObj &SpecConsts, const std::string &BuildOptionsString) { - + assert(!Devices.empty()); if (!areImagesCacheable(Imgs)) return {}; std::vector SortedImgs = getSortedImages(Imgs); - std::string Path = - getCacheItemPath(Device, SortedImgs, SpecConsts, BuildOptionsString); - - 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 + std::vector> Binaries(Devices.size()); + std::string FileNames; + for (size_t DeviceIndex = 0; DeviceIndex < Devices.size(); DeviceIndex++) { + std::string Path = getCacheItemPath(Devices[DeviceIndex], SortedImgs, + SpecConsts, BuildOptionsString); + + 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", Devices[DeviceIndex], + SortedImgs, SpecConsts, BuildOptionsString)) { + try { + std::string FullFileName = FileName + ".bin"; + Binaries[DeviceIndex] = readBinaryDataFromFile(FullFileName); + FileNames += FullFileName + ";"; + break; + } catch (...) { + // If read was unsuccessfull try the next item + } } + FileName = Path + "/" + std::to_string(++i); } - FileName = Path + "/" + std::to_string(++i); + // If there is no binary for any device, return empty vector. + if (Binaries[DeviceIndex].empty()) + return {}; } - return {}; + PersistentDeviceCodeCache::trace("using cached device binary: " + FileNames); + return Binaries; } /* kernel_compiler extension uses slightly different format for path @@ -267,36 +319,42 @@ std::vector> PersistentDeviceCodeCache::getItemFromDisc( */ std::vector> PersistentDeviceCodeCache::getCompiledKernelFromDisc( - const device &Device, const std::string &BuildOptionsString, + const std::vector &Devices, 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 + assert(!Devices.empty()); + std::vector> Binaries(Devices.size()); + std::string FileNames; + for (size_t DeviceIndex = 0; DeviceIndex < Devices.size(); DeviceIndex++) { + std::string DirName = getCompiledKernelItemPath( + Devices[DeviceIndex], 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"; + Binaries[DeviceIndex] = readBinaryDataFromFile(FullFileName); + FileNames += FullFileName + ";"; + break; + } catch (...) { + // If read was unsuccessfull try the next item + } } + FileName = DirName + "/" + std::to_string(++i); } - FileName = DirName + "/" + std::to_string(++i); + // If there is no binary for any device, return empty vector. + if (Binaries[DeviceIndex].empty()) + return {}; } - return {}; + PersistentDeviceCodeCache::trace_KernelCompiler("using cached binary: " + + FileNames); + return Binaries; } /* Returns string value which can be used to identify different device @@ -309,44 +367,29 @@ std::string PersistentDeviceCodeCache::getDeviceIDString(const device &Device) { } /* Write built binary to persistent cache - * Format: numImages, 1stImageSize, Image[, NthImageSize, NthImage...] - * Return on first unsuccessfull file operation + * Format: BinarySize, Binary */ void PersistentDeviceCodeCache::writeBinaryDataToFile( - const std::string &FileName, const std::vector> &Data) { + const std::string &FileName, const std::vector &Data) { std::ofstream FileStream{FileName, std::ios::binary}; - - size_t Size = Data.size(); + auto Size = Data.size(); FileStream.write((char *)&Size, sizeof(Size)); - - for (size_t i = 0; i < Data.size(); ++i) { - Size = Data[i].size(); - FileStream.write((char *)&Size, sizeof(Size)); - FileStream.write(Data[i].data(), Size); - } - FileStream.close(); + FileStream.write(Data.data(), Size); if (FileStream.fail()) - trace("Failed to write binary file " + FileName); + trace("Failed to write to binary file " + FileName); } -/* Read built binary to persistent cache - * Format: numImages, 1stImageSize, Image[, NthImageSize, NthImage...] +/* Read built binary from persistent cache. Each persistent cache file contains + * binary for a single device. Format: BinarySize, Binary */ -std::vector> +std::vector PersistentDeviceCodeCache::readBinaryDataFromFile(const std::string &FileName) { std::ifstream FileStream{FileName, std::ios::binary}; - size_t ImgNum = 0, ImgSize = 0; - FileStream.read((char *)&ImgNum, sizeof(ImgNum)); - - std::vector> Res(ImgNum); - for (size_t i = 0; i < ImgNum; ++i) { - FileStream.read((char *)&ImgSize, sizeof(ImgSize)); + size_t BinarySize = 0; + FileStream.read((char *)&BinarySize, sizeof(BinarySize)); - std::vector ImgData(ImgSize); - FileStream.read(ImgData.data(), ImgSize); - - Res[i] = std::move(ImgData); - } + std::vector BinaryData(BinarySize); + FileStream.read(BinaryData.data(), BinarySize); FileStream.close(); if (FileStream.fail()) { @@ -354,7 +397,7 @@ PersistentDeviceCodeCache::readBinaryDataFromFile(const std::string &FileName) { return {}; } - return Res; + return BinaryData; } /* Writing cache item key sources to be used for reliable identification diff --git a/sycl/source/detail/persistent_device_code_cache.hpp b/sycl/source/detail/persistent_device_code_cache.hpp index 19b145f6de895..d2038aaa65969 100644 --- a/sycl/source/detail/persistent_device_code_cache.hpp +++ b/sycl/source/detail/persistent_device_code_cache.hpp @@ -95,16 +95,15 @@ class PersistentDeviceCodeCache { */ private: /* Write built binary to persistent cache - * Format: numImages, 1stImageSize, Image[, NthImageSize, NthImage...] + * Format: BinarySize, Binary */ static void writeBinaryDataToFile(const std::string &FileName, - const std::vector> &Data); + const std::vector &Data); /* Read built binary to persistent cache - * Format: numImages, 1stImageSize, Image[, NthImageSize, NthImage...] + * Format: BinarySize, Binary */ - static std::vector> - readBinaryDataFromFile(const std::string &FileName); + static std::vector readBinaryDataFromFile(const std::string &FileName); /* Writing cache item key sources to be used for reliable identification * Format: Four pairs of [size, value] for device, build options, @@ -182,26 +181,26 @@ class PersistentDeviceCodeCache { * stored in vector of chars. */ static std::vector> - getItemFromDisc(const device &Device, + getItemFromDisc(const std::vector &Devices, const std::vector &Imgs, const SerializedObj &SpecConsts, const std::string &BuildOptionsString); static std::vector> - getCompiledKernelFromDisc(const device &Device, + getCompiledKernelFromDisc(const std::vector &Devices, const std::string &BuildOptionsString, const std::string SourceStr); /* Stores build program in persistent cache */ static void - putItemToDisc(const device &Device, + putItemToDisc(const std::vector &Devices, const std::vector &Imgs, const SerializedObj &SpecConsts, const std::string &BuildOptionsString, const ur_program_handle_t &NativePrg); - static void putCompiledKernelToDisc(const device &Device, + static void putCompiledKernelToDisc(const std::vector &Devices, const std::string &BuildOptionsString, const std::string &SourceStr, const ur_program_handle_t &NativePrg); diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 276dc0795f003..d6602725663ff 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -511,9 +511,17 @@ std::pair ProgramManager::getOrCreateURProgram( const std::string &CompileAndLinkOptions, SerializedObj SpecConsts) { ur_program_handle_t NativePrg; // TODO: Or native? - auto BinProg = PersistentDeviceCodeCache::getItemFromDisc( - Devices[0], AllImages, SpecConsts, CompileAndLinkOptions); - if (BinProg.size()) { + // Get binaries for each device (1:1 correpsondence with input Devices). + auto Binaries = PersistentDeviceCodeCache::getItemFromDisc( + Devices, AllImages, SpecConsts, CompileAndLinkOptions); + if (!Binaries.empty()) { + std::vector BinPtrs; + std::vector Lengths; + for (auto &Bin : Binaries) { + Lengths.push_back(Bin.size()); + BinPtrs.push_back(reinterpret_cast(Bin.data())); + } + // Get program metadata from properties std::vector ProgMetadataVector; for (const RTDeviceBinaryImage *Img : AllImages) { @@ -521,16 +529,13 @@ std::pair ProgramManager::getOrCreateURProgram( ProgMetadataVector.insert(ProgMetadataVector.end(), ImgProgMetadata.begin(), ImgProgMetadata.end()); } - std::vector Binaries(Devices.size(), - (const uint8_t *)BinProg[0].data()); - std::vector Lengths(Devices.size(), BinProg[0].size()); NativePrg = - createBinaryProgram(getSyclObjImpl(Context), Devices, Binaries.data(), + createBinaryProgram(getSyclObjImpl(Context), Devices, BinPtrs.data(), Lengths.data(), ProgMetadataVector); } else { NativePrg = createURProgram(MainImg, Context, Devices); } - return {NativePrg, BinProg.size()}; + return {NativePrg, Binaries.size()}; } /// Emits information about built programs if the appropriate contitions are @@ -901,7 +906,7 @@ ur_program_handle_t ProgramManager::getBuiltURProgram( // Save program to persistent cache if it is not there if (!DeviceCodeWasInCache) { - PersistentDeviceCodeCache::putItemToDisc(Device, AllImages, SpecConsts, + PersistentDeviceCodeCache::putItemToDisc({Device}, AllImages, SpecConsts, CompileOpts + LinkOpts, BuiltProgram.get()); } @@ -2667,9 +2672,8 @@ device_image_plain ProgramManager::build(const device_image_plain &DeviceImage, // Save program to persistent cache if it is not there if (!DeviceCodeWasInCache) - PersistentDeviceCodeCache::putItemToDisc(Devs[0], {&Img}, SpecConsts, - CompileOpts + LinkOpts, - BuiltProgram.get()); + PersistentDeviceCodeCache::putItemToDisc( + Devs, {&Img}, SpecConsts, CompileOpts + LinkOpts, BuiltProgram.get()); return BuiltProgram.release(); }; diff --git a/sycl/test-e2e/KernelAndProgram/persistent-cache-multi-device.cpp b/sycl/test-e2e/KernelAndProgram/persistent-cache-multi-device.cpp new file mode 100644 index 0000000000000..4486a7822dbd5 --- /dev/null +++ b/sycl/test-e2e/KernelAndProgram/persistent-cache-multi-device.cpp @@ -0,0 +1,38 @@ +// REQUIRES: (level_zero || opencl) && linux && gpu + +// RUN: %{build} -o %t.out +// RUN: rm -rf %t/cache_dir +// RUN: env NEOReadDebugKeys=1 CreateMultipleRootDevices=3 SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=1 SYCL_CACHE_DIR=%t/cache_dir env -u XDG_CACHE_HOME env -u HOME %{run} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-BUILD +// RUN: env NEOReadDebugKeys=1 CreateMultipleRootDevices=3 SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=1 SYCL_CACHE_DIR=%t/cache_dir env -u XDG_CACHE_HOME env -u HOME %{run} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-CACHE + +// Test checks that persistent cache works correctly with multiple devices. + +#include + +using namespace sycl; + +class SimpleKernel; + +int main(void) { + platform plt; + auto devs = plt.get_devices(); + context ctx(devs); + assert(devs.size() >= 3); + + constexpr size_t sz = 1024; + sycl::buffer bufA(sz); + auto bundle = sycl::get_kernel_bundle(ctx); + // CHECK-BUILD: [Persistent Cache]: device binary has been cached + // CHECK-CACHE: [Persistent Cache]: using cached device binary + auto bundle_exe = sycl::build(bundle, {devs[0], devs[2]}); + auto kernel = bundle_exe.get_kernel(sycl::get_kernel_id()); + sycl::queue q(devs[2]); + q.submit([&](sycl::handler &cgh) { + sycl::accessor accA(bufA, cgh, sycl::write_only); + cgh.parallel_for(sycl::range<1>(sz), [=](sycl::item<1> item) { + accA[item] = item.get_linear_id(); + }); + }); + q.wait(); + return 0; +} diff --git a/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp b/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp index c1f4af826c8eb..1cd0fcee45dc7 100644 --- a/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp +++ b/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp @@ -61,6 +61,19 @@ static ur_result_t redefinedProgramGetInfoAfter(void *pParams) { *value = Progs[DeviceCodeID].size(); } + if (*params.ppropName == UR_PROGRAM_INFO_DEVICES) { + if (*params.ppPropValue) { + for (size_t i = 0; i < Progs[DeviceCodeID].size(); i++) { + auto devs = static_cast(*params.ppPropValue); + devs[i] = reinterpret_cast(i + 1); + } + } + if (*params.ppPropSizeRet) + **params.ppPropSizeRet = + sizeof(ur_device_handle_t) * Progs[DeviceCodeID].size(); + return UR_RESULT_SUCCESS; + } + if (*params.ppropName == UR_PROGRAM_INFO_BINARY_SIZES) { auto value = reinterpret_cast(*params.ppPropValue); for (size_t i = 0; i < Progs[DeviceCodeID].size(); ++i) @@ -174,8 +187,8 @@ class PersistentDeviceCodeCache &redefinedProgramGetInfoAfter); } - /* Helper function for concurent cache item read/write from diffrent number - * of threads with diffrent cache item sizes: + /* Helper function for concurent cache item read/write from different number + * of threads with different cache item sizes: * ProgramID - defines program parameters to be used for testing (see Progs * vector above. * ThreadCount - number of parallel executors used for the test*/ @@ -184,7 +197,7 @@ class PersistentDeviceCodeCache std::to_string(ThreadCount)}; DeviceCodeID = ProgramID; std::string ItemDir = detail::PersistentDeviceCodeCache::getCacheItemPath( - Dev, {&Img}, {'S', 'p', 'e', 'c', 'C', 'o', 'n', 's', 't', ProgramID}, + {Dev}, {&Img}, {'S', 'p', 'e', 'c', 'C', 'o', 'n', 's', 't', ProgramID}, BuildOptions); ASSERT_NO_ERROR(llvm::sys::fs::remove_directories(ItemDir)); @@ -193,12 +206,12 @@ class PersistentDeviceCodeCache auto testLambda = [&](std::size_t threadId) { b.wait(); detail::PersistentDeviceCodeCache::putItemToDisc( - Dev, {&Img}, + {Dev}, {&Img}, std::vector( {'S', 'p', 'e', 'c', 'C', 'o', 'n', 's', 't', ProgramID}), BuildOptions, NativeProg); auto Res = detail::PersistentDeviceCodeCache::getItemFromDisc( - Dev, {&Img}, + {Dev}, {&Img}, std::vector( {'S', 'p', 'e', 'c', 'C', 'o', 'n', 's', 't', ProgramID}), BuildOptions); @@ -251,9 +264,9 @@ TEST_P(PersistentDeviceCodeCache, KeysWithNullTermSymbol) { Dev, {&Img}, SpecConst, Key); ASSERT_NO_ERROR(llvm::sys::fs::remove_directories(ItemDir)); - detail::PersistentDeviceCodeCache::putItemToDisc(Dev, {&Img}, SpecConst, Key, - NativeProg); - auto Res = detail::PersistentDeviceCodeCache::getItemFromDisc(Dev, {&Img}, + detail::PersistentDeviceCodeCache::putItemToDisc({Dev}, {&Img}, SpecConst, + Key, NativeProg); + auto Res = detail::PersistentDeviceCodeCache::getItemFromDisc({Dev}, {&Img}, SpecConst, Key); EXPECT_NE(Res.size(), static_cast(0)) << "Failed to load cache item"; for (size_t i = 0; i < Res.size(); ++i) { @@ -303,11 +316,11 @@ TEST_P(PersistentDeviceCodeCache, MultipleImages) { Dev, Imgs, {}, BuildOptions); ASSERT_NO_ERROR(llvm::sys::fs::remove_directories(ItemDir)); - detail::PersistentDeviceCodeCache::putItemToDisc(Dev, Imgs, {}, BuildOptions, - NativeProg); + detail::PersistentDeviceCodeCache::putItemToDisc({Dev}, Imgs, {}, + BuildOptions, NativeProg); // Check that the order of images does not affect the result. std::reverse(Imgs.begin(), Imgs.end()); - auto Res = detail::PersistentDeviceCodeCache::getItemFromDisc(Dev, Imgs, {}, + auto Res = detail::PersistentDeviceCodeCache::getItemFromDisc({Dev}, Imgs, {}, BuildOptions); EXPECT_NE(Res.size(), static_cast(0)) << "Failed to load cache item"; for (size_t i = 0; i < Res.size(); ++i) { @@ -357,29 +370,29 @@ TEST_P(PersistentDeviceCodeCache, CorruptedCacheFiles) { ASSERT_NO_ERROR(llvm::sys::fs::remove_directories(ItemDir)); // Only source file is present - detail::PersistentDeviceCodeCache::putItemToDisc(Dev, {&Img}, {}, + detail::PersistentDeviceCodeCache::putItemToDisc({Dev}, {&Img}, {}, BuildOptions, NativeProg); EXPECT_FALSE(llvm::sys::fs::remove(ItemDir + "/0.bin")) << "Failed to remove binary file"; - auto Res = detail::PersistentDeviceCodeCache::getItemFromDisc(Dev, {&Img}, {}, - BuildOptions); + auto Res = detail::PersistentDeviceCodeCache::getItemFromDisc( + {Dev}, {&Img}, {}, BuildOptions); EXPECT_EQ(Res.size(), static_cast(0)) << "Item with missed binary file was read"; ASSERT_NO_ERROR(llvm::sys::fs::remove_directories(ItemDir)); // Only binary file is present - detail::PersistentDeviceCodeCache::putItemToDisc(Dev, {&Img}, {}, + detail::PersistentDeviceCodeCache::putItemToDisc({Dev}, {&Img}, {}, BuildOptions, NativeProg); EXPECT_FALSE(llvm::sys::fs::remove(ItemDir + "/0.src")) << "Failed to remove source file"; - Res = detail::PersistentDeviceCodeCache::getItemFromDisc(Dev, {&Img}, {}, + Res = detail::PersistentDeviceCodeCache::getItemFromDisc({Dev}, {&Img}, {}, BuildOptions); EXPECT_EQ(Res.size(), static_cast(0)) << "Item with missed source file was read"; ASSERT_NO_ERROR(llvm::sys::fs::remove_directories(ItemDir)); // Binary file is corrupted - detail::PersistentDeviceCodeCache::putItemToDisc(Dev, {&Img}, {}, + detail::PersistentDeviceCodeCache::putItemToDisc({Dev}, {&Img}, {}, BuildOptions, NativeProg); std::ofstream FileStream(ItemDir + "/0.bin", std::ofstream::out | std::ofstream::trunc); @@ -389,7 +402,7 @@ TEST_P(PersistentDeviceCodeCache, CorruptedCacheFiles) { FileStream << 2 << 12 << "123456789012" << 23 << "1234"; FileStream.close(); EXPECT_FALSE(FileStream.fail()) << "Failed to create trancated binary file"; - Res = detail::PersistentDeviceCodeCache::getItemFromDisc(Dev, {&Img}, {}, + Res = detail::PersistentDeviceCodeCache::getItemFromDisc({Dev}, {&Img}, {}, BuildOptions); EXPECT_EQ(Res.size(), static_cast(0)) << "Item with corrupted binary file was read"; @@ -397,13 +410,13 @@ TEST_P(PersistentDeviceCodeCache, CorruptedCacheFiles) { ASSERT_NO_ERROR(llvm::sys::fs::remove_directories(ItemDir)); // Source file is empty - detail::PersistentDeviceCodeCache::putItemToDisc(Dev, {&Img}, {}, + detail::PersistentDeviceCodeCache::putItemToDisc({Dev}, {&Img}, {}, BuildOptions, NativeProg); { std::ofstream FileStream(ItemDir + "/0.src", std::ofstream::out | std::ofstream::trunc); } - Res = detail::PersistentDeviceCodeCache::getItemFromDisc(Dev, {&Img}, {}, + Res = detail::PersistentDeviceCodeCache::getItemFromDisc({Dev}, {&Img}, {}, BuildOptions); EXPECT_EQ(Res.size(), static_cast(0)) << "Item with corrupted binary file was read"; @@ -421,7 +434,7 @@ TEST_P(PersistentDeviceCodeCache, LockFile) { ASSERT_NO_ERROR(llvm::sys::fs::remove_directories(ItemDir)); // Create 1st cahe item - detail::PersistentDeviceCodeCache::putItemToDisc(Dev, {&Img}, {}, + detail::PersistentDeviceCodeCache::putItemToDisc({Dev}, {&Img}, {}, BuildOptions, NativeProg); EXPECT_TRUE(llvm::sys::fs::exists(ItemDir + "/0.bin")) << "No file created"; std::string LockFile = ItemDir + "/0.lock"; @@ -431,25 +444,25 @@ TEST_P(PersistentDeviceCodeCache, LockFile) { { std::ofstream File{LockFile}; } // Cache item is locked, cache miss happens on read - auto Res = detail::PersistentDeviceCodeCache::getItemFromDisc(Dev, {&Img}, {}, - BuildOptions); + auto Res = detail::PersistentDeviceCodeCache::getItemFromDisc( + {Dev}, {&Img}, {}, BuildOptions); EXPECT_EQ(Res.size(), static_cast(0)) << "Locked item was read"; // Cache item is locked - new cache item to be created - detail::PersistentDeviceCodeCache::putItemToDisc(Dev, {&Img}, {}, + detail::PersistentDeviceCodeCache::putItemToDisc({Dev}, {&Img}, {}, BuildOptions, NativeProg); EXPECT_TRUE(llvm::sys::fs::exists(ItemDir + "/1.bin")) << "No file created"; // Second cache item is locked, cache miss happens on read { std::ofstream File{ItemDir + "/1.lock"}; } - Res = detail::PersistentDeviceCodeCache::getItemFromDisc(Dev, {&Img}, {}, + Res = detail::PersistentDeviceCodeCache::getItemFromDisc({Dev}, {&Img}, {}, BuildOptions); EXPECT_EQ(Res.size(), static_cast(0)) << "Locked item was read"; // First cache item was unlocked and successfully read std::remove(LockFile.c_str()); - Res = detail::PersistentDeviceCodeCache::getItemFromDisc(Dev, {&Img}, {}, + Res = detail::PersistentDeviceCodeCache::getItemFromDisc({Dev}, {&Img}, {}, BuildOptions); for (size_t i = 0; i < Res.size(); ++i) { for (size_t j = 0; j < Res[i].size(); ++j) { @@ -469,20 +482,20 @@ TEST_P(PersistentDeviceCodeCache, AccessDeniedForCacheDir) { std::string ItemDir = detail::PersistentDeviceCodeCache::getCacheItemPath( Dev, {&Img}, {}, BuildOptions); ASSERT_NO_ERROR(llvm::sys::fs::remove_directories(ItemDir)); - detail::PersistentDeviceCodeCache::putItemToDisc(Dev, {&Img}, {}, + detail::PersistentDeviceCodeCache::putItemToDisc({Dev}, {&Img}, {}, BuildOptions, NativeProg); EXPECT_TRUE(llvm::sys::fs::exists(ItemDir + "/0.bin")) << "No file created"; ASSERT_NO_ERROR(llvm::sys::fs::setPermissions(ItemDir + "/0.bin", llvm::sys::fs::no_perms)); // No access to binary file new cache item to be created - detail::PersistentDeviceCodeCache::putItemToDisc(Dev, {&Img}, {}, + detail::PersistentDeviceCodeCache::putItemToDisc({Dev}, {&Img}, {}, BuildOptions, NativeProg); EXPECT_TRUE(llvm::sys::fs::exists(ItemDir + "/1.bin")) << "No file created"; ASSERT_NO_ERROR(llvm::sys::fs::setPermissions(ItemDir + "/1.bin", llvm::sys::fs::no_perms)); - auto Res = detail::PersistentDeviceCodeCache::getItemFromDisc(Dev, {&Img}, {}, - BuildOptions); + auto Res = detail::PersistentDeviceCodeCache::getItemFromDisc( + {Dev}, {&Img}, {}, BuildOptions); // No image to be read due to lack of permissions from source file EXPECT_EQ(Res.size(), static_cast(0)) @@ -493,7 +506,7 @@ TEST_P(PersistentDeviceCodeCache, AccessDeniedForCacheDir) { ASSERT_NO_ERROR(llvm::sys::fs::setPermissions(ItemDir + "/1.bin", llvm::sys::fs::all_perms)); - Res = detail::PersistentDeviceCodeCache::getItemFromDisc(Dev, {&Img}, {}, + Res = detail::PersistentDeviceCodeCache::getItemFromDisc({Dev}, {&Img}, {}, BuildOptions); // Image should be successfully read for (size_t i = 0; i < Res.size(); ++i) {