Raw implementation

Author: vladimirlaz

sycl/doc/EnvironmentVariables.md

@@ -35,8 +35,8 @@ subject to change. Do not rely on these variables in production code.
 | SYCL_ENABLE_PCI | Integer | When set to 1, enables obtaining the GPU PCI address when using the Level Zero backend. The default is 0. |
 | SYCL_HOST_UNIFIED_MEMORY | Integer | Enforce host unified memory support or lack of it for the execution graph builder. If set to 0, it is enforced as not supported by all devices. If set to 1, it is enforced as supported by all devices. |
 | SYCL_CACHE_DIR | Path | Path to persistent cache root directory. Default values are `%AppData%\Intel\sycl_program_cache` for Windows and `$HOME/intel/sycl_program_cache` on Linux. |
-| SYCL_CACHE_ENABLED | ON, OFF | Switches persistent cache switch on or off. Default value is ON. |
-| SYCL_CACHE_EVICTION | ON, OFF | Switches cache eviction on or off. Default value is ON. |
+| SYCL_CACHE_DISABLE_PERSISTENT | Any(\*) | Switches persistent cache switch off. Default value is ON. |
+| SYCL_CACHE_EVICTION_DISABLE | Any(\*) | Switches cache eviction off. Default value is ON. |
 | SYCL_CACHE_MAX_SIZE | Positive integer | Cache eviction is triggered once total size of cached images exceeds the value in megabytes (default - 8 192 for 8 GB). Set to 0 to disable size-based cache eviction. |
 | SYCL_CACHE_THRESHOLD | Positive integer | Cache eviction threshold in days (default value is 7 for 1 week). Set to 0 for disabling time-based cache eviction. |
 | SYCL_CACHE_MIN_DEVICE_IMAGE_SIZE | Positive integer | Minimum size of device code image in kilobytes which is reasonable to cache on disk because disk access operation may take more time than do JIT compilation for it. Default value is 0 to cache all images. |

sycl/include/CL/sycl/detail/os_util.hpp

@@ -80,6 +80,9 @@ class __SYCL_EXPORT OSUtil {
 
   /// Deallocates the memory referenced by \p Ptr.
   static void alignedFree(void *Ptr);
+
+  /// Returns the path to directory storing on-disk SYCL program cache.
+  static std::string getCacheRoot();
 };
 
 } // namespace detail

sycl/source/detail/config.def

@@ -21,7 +21,8 @@ CONFIG(SYCL_PROGRAM_LINK_OPTIONS, 64, __SYCL_PROGRAM_LINK_OPTIONS)
 CONFIG(SYCL_PROGRAM_COMPILE_OPTIONS, 64, __SYCL_PROGRAM_COMPILE_OPTIONS)
 CONFIG(SYCL_HOST_UNIFIED_MEMORY, 1, __SYCL_HOST_UNIFIED_MEMORY)
 CONFIG(SYCL_CACHE_DIR, 1024, __SYCL_CACHE_DIR)
-CONFIG(SYCL_CACHE_ENABLED, 3, __SYCL_CACHE_ENABLED)
+CONFIG(SYCL_CACHE_DISABLE_PERSISTENT, 1, __SYCL_CACHE_DISABLE_PERSISTENT)
+CONFIG(SYCL_CACHE_EVICTION_DISABLE, 1, __SYCL_CACHE_EVICTION_DISABLE)
 CONFIG(SYCL_CACHE_MAX_SIZE, 16, __SYCL_CACHE_MAX_SIZE)
 CONFIG(SYCL_CACHE_THRESHOLD, 16, __SYCL_CACHE_THRESHOLD)
 CONFIG(SYCL_CACHE_MIN_DEVICE_IMAGE_SIZE, 16, __SYCL_CACHE_MIN_DEVICE_IMAGE_SIZE)

sycl/source/detail/os_util.cpp

@@ -8,6 +8,7 @@
 
 #include <CL/sycl/detail/os_util.hpp>
 #include <CL/sycl/exception.hpp>
+#include <detail/config.hpp>
 
 #include <cassert>
 
@@ -121,7 +122,7 @@ std::string OSUtil::getCurrentDSODir() {
   //
   //  4) Extract an absolute path to a filename and get a dirname from it.
   //
-  uintptr_t CurrentFunc = (uintptr_t) &getCurrentDSODir;
+  uintptr_t CurrentFunc = (uintptr_t)&getCurrentDSODir;
   std::ifstream Stream("/proc/self/maps");
   Stream >> std::hex;
   while (!Stream.eof()) {
@@ -166,7 +167,7 @@ std::string OSUtil::getCurrentDSODir() {
   return "";
 }
 
-std::string OSUtil::getDirName(const char* Path) {
+std::string OSUtil::getDirName(const char *Path) {
   std::string Tmp(Path);
   // dirname(3) needs a writable C string: a null-terminator is written where a
   // path should split.
@@ -258,6 +259,22 @@ void OSUtil::alignedFree(void *Ptr) {
 #endif
 }
 
+std::string OSUtil::getCacheRoot() {
+  static const char *PersistenCacheRoot = SYCLConfig<SYCL_CACHE_DIR>::get();
+  if (PersistenCacheRoot)
+    return PersistenCacheRoot;
+
+#if defined(__SYCL_RT_OS_LINUX)
+  static const char *RootDir = std::getenv("HOME");
+#else
+  static const char *RootDir = std::getenv("AppData");
+#endif
+  std::string Root{RootDir ? RootDir : "."};
+
+  Root += "/intel/sycl_cache";
+  return Root;
+}
+
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/source/detail/program_manager/program_manager.cpp

@@ -30,18 +30,26 @@
 #include <cstdlib>
 #include <cstring>
 #include <fstream>
+#include <functional>
+#include <libgen.h>
 #include <memory>
 #include <mutex>
 #include <sstream>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
 #include <string>
+#include <sys/stat.h>
+#include <sys/types.h>
+#include <unistd.h>
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
 
 using ContextImplPtr = std::shared_ptr<cl::sycl::detail::context_impl>;
 
-static constexpr int DbgProgMgr = 1;
+static constexpr int DbgProgMgr = 2;
 
 enum BuildState { BS_InProgress, BS_Done, BS_Failed };
 
@@ -182,6 +190,7 @@ getOrBuild(KernelProgramCache &KPCache, KeyT &&CacheKey, AcquireFT &&Acquire,
 
   // only the building thread will run this
   try {
+
     RetT *Desired = Build();
 
 #ifndef NDEBUG
@@ -346,6 +355,204 @@ RT::PiProgram ProgramManager::createPIProgram(const RTDeviceBinaryImage &Img,
   return Res;
 }
 
+long GetFileSize(const char *FileName) {
+  struct stat Stat;
+  if (!stat(FileName, &Stat))
+    return Stat.st_size;
+  return -1;
+}
+
+inline bool IsFSEntryPresent(const char *Path) {
+  struct stat Stat;
+  return !stat(Path, &Stat);
+}
+
+int MakePathRecur(const char *Dir, mode_t Mode) {
+  assert((Dir != nullptr) && "Passed null-pointer as directory name.");
+
+  // Directory is present - do nothing
+  if (IsFSEntryPresent(Dir))
+    return 0;
+
+  char *CurDir = strdup(Dir);
+  MakePathRecur(dirname(CurDir), Mode);
+  if (DbgProgMgr > 1)
+    std::cerr << "Created directory: " << CurDir << std::endl;
+
+  free(CurDir);
+  return mkdir(Dir, Mode);
+}
+
+void WriteCacheItem(const std::string &FileName,
+                    const std::vector<std::vector<char>> &Data) {
+  std::ofstream FileStream{FileName, std::ios::binary};
+  if (DbgProgMgr > 1) {
+    std::cerr << "####Writing programs built for " << std::dec << Data.size()
+              << " devices:\n";
+  }
+
+  size_t Size = Data.size();
+  FileStream.write((char *)&Size, sizeof(Size));
+  for (size_t i = 0; i < Data.size(); ++i) {
+    if (DbgProgMgr > 1) {
+      std::cerr << "\tWrite " << i << "-th image of size " << std::dec
+                << Data[i].size() << "\n";
+    }
+    Size = Data[i].size();
+    FileStream.write((char *)&Size, sizeof(Size));
+    FileStream.write(Data[i].data(), Size);
+  }
+  FileStream.close();
+}
+
+std::vector<std::vector<char>> ReadCacheItem(const std::string &FileName) {
+  std::vector<std::vector<char>> Res;
+  std::ifstream FileStream{FileName, std::ios::binary};
+  size_t ImgNum, ImgSize;
+  FileStream.read((char *)&ImgNum, sizeof(ImgNum));
+  if (DbgProgMgr > 1) {
+    std::cerr << "####Reading programs built for " << std::dec << ImgNum
+              << " devices:\n";
+  }
+
+  Res.resize(ImgNum);
+
+  for (size_t i = 0; i < ImgNum; ++i) {
+    FileStream.read((char *)&ImgSize, sizeof(ImgSize));
+    if (DbgProgMgr > 1) {
+      std::cerr << "\tRead " << i << "-th image of size " << std::dec << ImgSize
+                << "\n";
+    }
+
+    Res[i].resize(ImgSize);
+    FileStream.read(Res[i].data(), ImgSize);
+  }
+
+  return Res;
+}
+
+std::string getDeviceString(const device &Device) {
+  return {Device.get_platform().get_info<sycl::info::platform::name>() +
+          Device.get_info<sycl::info::device::name>() +
+          Device.get_info<sycl::info::device::version>() +
+          Device.get_info<sycl::info::device::driver_version>()};
+}
+
+std::string DumpBinData(const unsigned char *Data, size_t Size) {
+  if (!Size)
+    return "NONE";
+  std::stringstream ss;
+  for (size_t i = 0; i < Size; i++) {
+    ss << std::hex << (int)Data[i];
+  }
+  return ss.str();
+}
+
+std::string GetCacheItemDirName(const device &Device,
+                                const RTDeviceBinaryImage &Img,
+                                const SerializedObj SpecConsts,
+                                const std::string &BuildOptionsString) {
+  static std::string cache_root{detail::OSUtil::getCacheRoot()};
+
+  std::string ImgString{
+      DumpBinData(Img.getRawData().BinaryStart, Img.getSize())};
+  std::string DeviceString{getDeviceString(Device)};
+  std::string SpecConstsString{
+      DumpBinData(SpecConsts.data(), SpecConsts.size())};
+  std::hash<std::string> StringHasher{};
+  return {cache_root + "/" + std::to_string(StringHasher(DeviceString)) + "/" +
+          std::to_string(StringHasher(ImgString)) + "/" +
+          std::to_string(StringHasher(SpecConstsString)) + "/" +
+          std::to_string(StringHasher(BuildOptionsString))};
+}
+
+bool IsPersistentCacheEnabled() {
+  static const char *PersistenCacheDisabled =
+      SYCLConfig<SYCL_CACHE_DISABLE_PERSISTENT>::get();
+
+  if (DbgProgMgr > 0)
+    std::cerr << "Persistent cache "
+              << (PersistenCacheDisabled ? "disabled." : "enabled.")
+              << std::endl;
+  return !PersistenCacheDisabled;
+}
+
+void ProgramManager::putPIProgramToDisc(const detail::plugin &Plugin,
+                                        const device &Device,
+                                        const RTDeviceBinaryImage &Img,
+                                        const SerializedObj SpecConsts,
+                                        const std::string &BuildOptionsString,
+                                        const RT::PiProgram &Program) {
+  if (!IsPersistentCacheEnabled()) {
+    return;
+  }
+
+  static std::string DirName =
+      GetCacheItemDirName(Device, Img, SpecConsts, BuildOptionsString);
+
+  size_t i = 0;
+  std::string FileName;
+  do {
+    FileName = DirName + "/" + std::to_string(i++) + ".bin";
+  } while (IsFSEntryPresent(FileName.c_str()));
+
+  size_t DeviceNum;
+  Plugin.call<PiApiKind::piProgramGetInfo>(Program, PI_PROGRAM_INFO_NUM_DEVICES,
+                                           sizeof(DeviceNum), &DeviceNum,
+                                           nullptr);
+  std::vector<size_t> BinarySizes(DeviceNum);
+  Plugin.call<PiApiKind::piProgramGetInfo>(
+      Program, PI_PROGRAM_INFO_BINARY_SIZES,
+      sizeof(size_t) * BinarySizes.size(), BinarySizes.data(), nullptr);
+
+  std::vector<std::vector<char>> Result;
+  std::vector<char *> Pointers;
+  for (size_t I = 0; I < BinarySizes.size(); ++I) {
+    Result.emplace_back(BinarySizes[I]);
+    Pointers.push_back(Result[I].data());
+  }
+
+  Plugin.call<PiApiKind::piProgramGetInfo>(Program, PI_PROGRAM_INFO_BINARIES,
+                                           sizeof(char *) * Pointers.size(),
+                                           Pointers.data(), nullptr);
+
+  MakePathRecur(DirName.c_str(), 0777);
+  WriteCacheItem(FileName, Result);
+}
+
+bool ProgramManager::getPIProgramFromDisc(ContextImplPtr ContextImpl,
+                                          const device &Device,
+                                          const RTDeviceBinaryImage &Img,
+                                          const SerializedObj SpecConsts,
+                                          const std::string &BuildOptionsString,
+                                          RT::PiProgram &NativePrg) {
+
+  if (!IsPersistentCacheEnabled())
+    return false;
+
+  std::string Path{
+      GetCacheItemDirName(Device, Img, SpecConsts, BuildOptionsString)};
+
+  if (!IsFSEntryPresent(Path.c_str()))
+    return false;
+
+  int i = 0;
+  std::string BinFileName{Path + "/" + std::to_string(i) + ".bin"};
+  while (IsFSEntryPresent(BinFileName.c_str())) {
+    auto BinDataItem = ReadCacheItem(BinFileName);
+    if (BinDataItem.size()) {
+      // TODO: Build for multiple devices once supported by program manager
+      NativePrg = createBinaryProgram(
+          ContextImpl, Device, (const unsigned char *)BinDataItem[0].data(),
+          BinDataItem[0].size());
+      return true;
+    }
+    BinFileName = Path + "/" + std::to_string(++i) + ".bin";
+  }
+
+  return false;
+}
+
 RT::PiProgram ProgramManager::getBuiltPIProgram(OSModuleHandle M,
                                                 const context &Context,
                                                 const device &Device,
@@ -390,9 +597,12 @@ RT::PiProgram ProgramManager::getBuiltPIProgram(OSModuleHandle M,
   if (LinkOptsEnv) {
     LinkOpts = LinkOptsEnv;
   }
+  SerializedObj SpecConsts;
+  if (Prg)
+    Prg->stableSerializeSpecConstRegistry(SpecConsts);
 
   auto BuildF = [this, &M, &KSId, &Context, &Device, Prg, &CompileOpts,
-                 &LinkOpts, &JITCompilationIsRequired] {
+                 &LinkOpts, &JITCompilationIsRequired, SpecConsts] {
     const RTDeviceBinaryImage &Img =
         getDeviceImage(M, KSId, Context, Device, JITCompilationIsRequired);
     // Update only if compile options are not overwritten by environment
@@ -413,19 +623,28 @@ RT::PiProgram ProgramManager::getBuiltPIProgram(OSModuleHandle M,
       LinkOpts += Img.getLinkOptions();
     ContextImplPtr ContextImpl = getSyclObjImpl(Context);
     const detail::plugin &Plugin = ContextImpl->getPlugin();
-    RT::PiProgram NativePrg = createPIProgram(Img, Context, Device);
-    if (Prg)
-      flushSpecConstants(*Prg, NativePrg, &Img);
+    RT::PiProgram NativePrg;
+    bool LoadedFromDiskCache =
+        getPIProgramFromDisc(ContextImpl, Device, Img, SpecConsts,
+                             CompileOpts + LinkOpts, NativePrg);
+    if (!LoadedFromDiskCache) {
+      NativePrg = createPIProgram(Img, Context, Device);
+      if (Prg)
+        flushSpecConstants(*Prg, NativePrg, &Img);
+    }
+
     ProgramPtr ProgramManaged(
         NativePrg, Plugin.getPiPlugin().PiFunctionTable.piProgramRelease);
 
     // Link a fallback implementation of device libraries if they are not
     // supported by a device compiler.
-    // Pre-compiled programs are supposed to be already linked.
+    // Pre-compiled programs (after AOT compilation or read from persitent
+    // cache) are supposed to be already linked.
     // If device image is not SPIR-V, DeviceLibReqMask will be 0 which means
     // no fallback device library will be linked.
     uint32_t DeviceLibReqMask = 0;
-    if (Img.getFormat() == PI_DEVICE_BINARY_TYPE_SPIRV &&
+    if (!LoadedFromDiskCache &&
+        Img.getFormat() == PI_DEVICE_BINARY_TYPE_SPIRV &&
         !SYCLConfig<SYCL_DEVICELIB_NO_FALLBACK>::get())
       DeviceLibReqMask = getDeviceLibReqMask(Img);
 
@@ -438,13 +657,12 @@ RT::PiProgram ProgramManager::getBuiltPIProgram(OSModuleHandle M,
       std::lock_guard<std::mutex> Lock(MNativeProgramsMutex);
       NativePrograms[BuiltProgram.get()] = &Img;
     }
+    if (!LoadedFromDiskCache)
+      putPIProgramToDisc(Plugin, Device, Img, SpecConsts,
+                         CompileOpts + LinkOpts, BuiltProgram.get());
     return BuiltProgram.release();
   };
 
-  SerializedObj SpecConsts;
-  if (Prg)
-    Prg->stableSerializeSpecConstRegistry(SpecConsts);
-
   const RT::PiDevice PiDevice = getRawSyclObjImpl(Device)->getHandleRef();
   auto BuildResult = getOrBuild<PiProgramT, compile_program_error>(
       Cache,