[SYCL] Initial implementation of dynamic linking support in runtime (#14587)

This patch provides an initial implementation for supporting the dynamic
linking feature. Current known limitations are: lack of kernel bundle
and AOT support.

Author: sergey-semenov

sycl/include/sycl/detail/pi.h

@@ -1083,6 +1083,8 @@ static const uint8_t PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL = 4;
 #define __SYCL_PI_PROPERTY_SET_SYCL_ASSERT_USED "SYCL/assert used"
 /// PropertySetRegistry::SYCL_EXPORTED_SYMBOLS defined in PropertySetIO.h
 #define __SYCL_PI_PROPERTY_SET_SYCL_EXPORTED_SYMBOLS "SYCL/exported symbols"
+/// PropertySetRegistry::SYCL_IMPORTED_SYMBOLS defined in PropertySetIO.h
+#define __SYCL_PI_PROPERTY_SET_SYCL_IMPORTED_SYMBOLS "SYCL/imported symbols"
 /// PropertySetRegistry::SYCL_DEVICE_GLOBALS defined in PropertySetIO.h
 #define __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_GLOBALS "SYCL/device globals"
 /// PropertySetRegistry::SYCL_DEVICE_REQUIREMENTS defined in PropertySetIO.h

sycl/source/detail/device_binary_image.cpp

@@ -178,6 +178,7 @@ void RTDeviceBinaryImage::init(pi_device_binary Bin) {
   AssertUsed.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_ASSERT_USED);
   ProgramMetadata.init(Bin, __SYCL_PI_PROPERTY_SET_PROGRAM_METADATA);
   ExportedSymbols.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_EXPORTED_SYMBOLS);
+  ImportedSymbols.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_IMPORTED_SYMBOLS);
   DeviceGlobals.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_GLOBALS);
   DeviceRequirements.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_REQUIREMENTS);
   HostPipes.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_HOST_PIPES);

sycl/source/detail/device_binary_image.hpp

@@ -214,6 +214,7 @@ class RTDeviceBinaryImage {
   const PropertyRange &getAssertUsed() const { return AssertUsed; }
   const PropertyRange &getProgramMetadata() const { return ProgramMetadata; }
   const PropertyRange &getExportedSymbols() const { return ExportedSymbols; }
+  const PropertyRange &getImportedSymbols() const { return ImportedSymbols; }
   const PropertyRange &getDeviceGlobals() const { return DeviceGlobals; }
   const PropertyRange &getDeviceRequirements() const {
     return DeviceRequirements;
@@ -240,6 +241,7 @@ class RTDeviceBinaryImage {
   RTDeviceBinaryImage::PropertyRange AssertUsed;
   RTDeviceBinaryImage::PropertyRange ProgramMetadata;
   RTDeviceBinaryImage::PropertyRange ExportedSymbols;
+  RTDeviceBinaryImage::PropertyRange ImportedSymbols;
   RTDeviceBinaryImage::PropertyRange DeviceGlobals;
   RTDeviceBinaryImage::PropertyRange DeviceRequirements;
   RTDeviceBinaryImage::PropertyRange HostPipes;

sycl/source/detail/persistent_device_code_cache.cpp

@@ -59,11 +59,19 @@ IsSupportedImageFormat(sycl::detail::pi::PiDeviceBinaryType Format) {
          Format == PI_DEVICE_BINARY_TYPE_NATIVE;
 }
 
-/* Returns true if specified image should be cached on disk. It checks if
- * cache is enabled, image has supported format and matches thresholds. */
-bool PersistentDeviceCodeCache::isImageCached(const RTDeviceBinaryImage &Img) {
+/* Returns true if specified images should be cached on disk. It checks if
+ * cache is enabled, images have supported format and match thresholds. */
+bool PersistentDeviceCodeCache::areImagesCacheable(
+    const std::vector<const RTDeviceBinaryImage *> &Imgs) {
+  assert(!Imgs.empty());
+  auto Format = Imgs[0]->getFormat();
+  assert(std::all_of(Imgs.begin(), Imgs.end(),
+                     [&Format](const RTDeviceBinaryImage *Img) {
+                       return Img->getFormat() == Format;
+                     }) &&
+         "All images are expected to have the same format");
   // Cache should be enabled and image type is one of the supported formats.
-  if (!isEnabled() || !IsSupportedImageFormat(Img.getFormat()))
+  if (!isEnabled() || !IsSupportedImageFormat(Format))
     return false;
 
   // Disable cache for ITT-profiled images.
@@ -79,25 +87,42 @@ bool PersistentDeviceCodeCache::isImageCached(const RTDeviceBinaryImage &Img) {
 
   // Make sure that image size is between caching thresholds if they are set.
   // Zero values for threshold is treated as disabled threshold.
-  if ((MaxImgSize && (Img.getSize() > MaxImgSize)) ||
-      (MinImgSize && (Img.getSize() < MinImgSize)))
+  size_t TotalSize = 0;
+  for (const RTDeviceBinaryImage *Img : Imgs)
+    TotalSize += Img->getSize();
+  if ((MaxImgSize && (TotalSize > MaxImgSize)) ||
+      (MinImgSize && (TotalSize < MinImgSize)))
     return false;
 
   return true;
 }
 
-/* Stores built program in persisten cache
+static std::vector<const RTDeviceBinaryImage *>
+getSortedImages(const std::vector<const RTDeviceBinaryImage *> &Imgs) {
+  std::vector<const RTDeviceBinaryImage *> SortedImgs = Imgs;
+  std::sort(SortedImgs.begin(), SortedImgs.end(),
+            [](const RTDeviceBinaryImage *A, const RTDeviceBinaryImage *B) {
+              // All entry names are unique among these images, so comparing the
+              // first ones is enough.
+              return std::strcmp(A->getRawData().EntriesBegin->name,
+                                 B->getRawData().EntriesBegin->name) < 0;
+            });
+  return SortedImgs;
+}
+
+/* Stores built program in persistent cache
  */
 void PersistentDeviceCodeCache::putItemToDisc(
-    const device &Device, const RTDeviceBinaryImage &Img,
+    const device &Device, const std::vector<const RTDeviceBinaryImage *> &Imgs,
     const SerializedObj &SpecConsts, const std::string &BuildOptionsString,
     const sycl::detail::pi::PiProgram &NativePrg) {
 
-  if (!isImageCached(Img))
+  if (!areImagesCacheable(Imgs))
     return;
 
+  std::vector<const RTDeviceBinaryImage *> SortedImgs = getSortedImages(Imgs);
   std::string DirName =
-      getCacheItemPath(Device, Img, SpecConsts, BuildOptionsString);
+      getCacheItemPath(Device, SortedImgs, SpecConsts, BuildOptionsString);
 
   if (DirName.empty())
     return;
@@ -139,7 +164,7 @@ void PersistentDeviceCodeCache::putItemToDisc(
       std::string FullFileName = FileName + ".bin";
       writeBinaryDataToFile(FullFileName, Result);
       trace("device binary has been cached: " + FullFileName);
-      writeSourceItem(FileName + ".src", Device, Img, SpecConsts,
+      writeSourceItem(FileName + ".src", Device, SortedImgs, SpecConsts,
                       BuildOptionsString);
     } else {
       PersistentDeviceCodeCache::trace("cache lock not owned " + FileName);
@@ -160,14 +185,15 @@ void PersistentDeviceCodeCache::putItemToDisc(
  * stored in vector of chars.
  */
 std::vector<std::vector<char>> PersistentDeviceCodeCache::getItemFromDisc(
-    const device &Device, const RTDeviceBinaryImage &Img,
+    const device &Device, const std::vector<const RTDeviceBinaryImage *> &Imgs,
     const SerializedObj &SpecConsts, const std::string &BuildOptionsString) {
 
-  if (!isImageCached(Img))
+  if (!areImagesCacheable(Imgs))
     return {};
 
+  std::vector<const RTDeviceBinaryImage *> SortedImgs = getSortedImages(Imgs);
   std::string Path =
-      getCacheItemPath(Device, Img, SpecConsts, BuildOptionsString);
+      getCacheItemPath(Device, SortedImgs, SpecConsts, BuildOptionsString);
 
   if (Path.empty() || !OSUtil::isPathPresent(Path))
     return {};
@@ -179,7 +205,7 @@ std::vector<std::vector<char>> PersistentDeviceCodeCache::getItemFromDisc(
          OSUtil::isPathPresent(FileName + ".src")) {
 
     if (!LockCacheItem::isLocked(FileName) &&
-        isCacheItemSrcEqual(FileName + ".src", Device, Img, SpecConsts,
+        isCacheItemSrcEqual(FileName + ".src", Device, SortedImgs, SpecConsts,
                             BuildOptionsString)) {
       try {
         std::string FullFileName = FileName + ".bin";
@@ -256,12 +282,12 @@ PersistentDeviceCodeCache::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, specialization
- * constant values, device code SPIR-V image.
+ * constant values, device code SPIR-V images.
  */
 void PersistentDeviceCodeCache::writeSourceItem(
     const std::string &FileName, const device &Device,
-    const RTDeviceBinaryImage &Img, const SerializedObj &SpecConsts,
-    const std::string &BuildOptionsString) {
+    const std::vector<const RTDeviceBinaryImage *> &SortedImgs,
+    const SerializedObj &SpecConsts, const std::string &BuildOptionsString) {
   std::ofstream FileStream{FileName, std::ios::binary};
 
   std::string DeviceString{getDeviceIDString(Device)};
@@ -277,9 +303,13 @@ void PersistentDeviceCodeCache::writeSourceItem(
   FileStream.write((char *)&Size, sizeof(Size));
   FileStream.write((const char *)SpecConsts.data(), Size);
 
-  Size = Img.getSize();
+  Size = 0;
+  for (const RTDeviceBinaryImage *Img : SortedImgs)
+    Size += Img->getSize();
   FileStream.write((char *)&Size, sizeof(Size));
-  FileStream.write((const char *)Img.getRawData().BinaryStart, Size);
+  for (const RTDeviceBinaryImage *Img : SortedImgs)
+    FileStream.write((const char *)Img->getRawData().BinaryStart,
+                     Img->getSize());
   FileStream.close();
 
   if (FileStream.fail()) {
@@ -292,12 +322,14 @@ void PersistentDeviceCodeCache::writeSourceItem(
  */
 bool PersistentDeviceCodeCache::isCacheItemSrcEqual(
     const std::string &FileName, const device &Device,
-    const RTDeviceBinaryImage &Img, const SerializedObj &SpecConsts,
-    const std::string &BuildOptionsString) {
+    const std::vector<const RTDeviceBinaryImage *> &SortedImgs,
+    const SerializedObj &SpecConsts, const std::string &BuildOptionsString) {
   std::ifstream FileStream{FileName, std::ios::binary};
 
-  std::string ImgString{(const char *)Img.getRawData().BinaryStart,
-                        Img.getSize()};
+  std::string ImgsString;
+  for (const RTDeviceBinaryImage *Img : SortedImgs)
+    ImgsString.append((const char *)Img->getRawData().BinaryStart,
+                      Img->getSize());
   std::string SpecConstsString{(const char *)SpecConsts.data(),
                                SpecConsts.size()};
 
@@ -323,7 +355,7 @@ bool PersistentDeviceCodeCache::isCacheItemSrcEqual(
   FileStream.read((char *)&Size, sizeof(Size));
   res.resize(Size);
   FileStream.read(&res[0], Size);
-  if (ImgString.compare(res))
+  if (ImgsString.compare(res))
     return false;
 
   FileStream.close();
@@ -335,29 +367,31 @@ bool PersistentDeviceCodeCache::isCacheItemSrcEqual(
   return true;
 }
 
-/* Returns directory name to store specific kernel image for specified
+/* Returns directory name to store specific kernel images for specified
  * device, build options and specialization constants values.
  */
 std::string PersistentDeviceCodeCache::getCacheItemPath(
-    const device &Device, const RTDeviceBinaryImage &Img,
+    const device &Device, const std::vector<const RTDeviceBinaryImage *> &Imgs,
     const SerializedObj &SpecConsts, const std::string &BuildOptionsString) {
   std::string cache_root{getRootDir()};
   if (cache_root.empty()) {
     trace("Disable persistent cache due to unconfigured cache root.");
     return {};
   }
 
-  std::string ImgString = "";
-  if (Img.getRawData().BinaryStart)
-    ImgString.assign((const char *)Img.getRawData().BinaryStart, Img.getSize());
+  std::string ImgsString;
+  for (const RTDeviceBinaryImage *Img : Imgs)
+    if (Img->getRawData().BinaryStart)
+      ImgsString.append((const char *)Img->getRawData().BinaryStart,
+                        Img->getSize());
 
   std::string DeviceString{getDeviceIDString(Device)};
   std::string SpecConstsString{(const char *)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(ImgsString)) + "/" +
          std::to_string(StringHasher(SpecConstsString)) + "/" +
          std::to_string(StringHasher(BuildOptionsString));
 }

sycl/source/detail/persistent_device_code_cache.hpp

@@ -69,7 +69,7 @@ class PersistentDeviceCodeCache {
    *   <cache_root>                 - root directory storing cache files;
    *   <device_hash>                - hash out of device information used to
    *                                  identify target device;
-   *   <device_image_hash>          - hash made out of device image used as
+   *   <device_image_hash>          - hash made out of device images used as
    *                                  input for the JIT compilation;
    *   <spec_constants_values_hash> - hash for specialization constants values;
    *   <build_options_hash>         - hash for all build options;
@@ -80,7 +80,7 @@ class PersistentDeviceCodeCache {
    *                                  started from 0).
    * Two files per cache item are stored on disk:
    *   <n>.src  - contains full values for build parameters (device information,
-   *              specialization constant values, build options, device image)
+   *              specialization constant values, build options, device images)
    *              which is used to resolve hash collisions and analysis of
    *              cached items.
    *   <n>.bin  - contains built device code.
@@ -108,20 +108,20 @@ class PersistentDeviceCodeCache {
 
   /* Writing cache item key sources to be used for reliable identification
    * Format: Four pairs of [size, value] for device, build options,
-   * specialization constant values, device code SPIR-V image.
+   * specialization constant values, device code SPIR-V images.
    */
-  static void writeSourceItem(const std::string &FileName, const device &Device,
-                              const RTDeviceBinaryImage &Img,
-                              const SerializedObj &SpecConsts,
-                              const std::string &BuildOptionsString);
+  static void
+  writeSourceItem(const std::string &FileName, const device &Device,
+                  const std::vector<const RTDeviceBinaryImage *> &SortedImgs,
+                  const SerializedObj &SpecConsts,
+                  const std::string &BuildOptionsString);
 
   /* Check that cache item key sources are equal to the current program
    */
-  static bool isCacheItemSrcEqual(const std::string &FileName,
-                                  const device &Device,
-                                  const RTDeviceBinaryImage &Img,
-                                  const SerializedObj &SpecConsts,
-                                  const std::string &BuildOptionsString);
+  static bool isCacheItemSrcEqual(
+      const std::string &FileName, const device &Device,
+      const std::vector<const RTDeviceBinaryImage *> &SortedImgs,
+      const SerializedObj &SpecConsts, const std::string &BuildOptionsString);
 
   /* Check if on-disk cache enabled.
    */
@@ -133,9 +133,10 @@ class PersistentDeviceCodeCache {
   /* Form string representing device version */
   static std::string getDeviceIDString(const device &Device);
 
-  /* Returns true if specified image should be cached on disk. It checks if
-   * cache is enabled, image has SPIRV type and matches thresholds. */
-  static bool isImageCached(const RTDeviceBinaryImage &Img);
+  /* Returns true if specified images should be cached on disk. It checks if
+   * cache is enabled, images have SPIRV type and match thresholds. */
+  static bool areImagesCacheable(
+      const std::vector<const RTDeviceBinaryImage *> &SortedImgs);
 
   /* Returns value of specified parameter. Default value is used if failure
    * happens during obtaining value. */
@@ -162,27 +163,30 @@ class PersistentDeviceCodeCache {
 public:
   /* Get directory name for storing current cache item
    */
-  static std::string getCacheItemPath(const device &Device,
-                                      const RTDeviceBinaryImage &Img,
-                                      const SerializedObj &SpecConsts,
-                                      const std::string &BuildOptionsString);
+  static std::string
+  getCacheItemPath(const device &Device,
+                   const std::vector<const RTDeviceBinaryImage *> &SortedImgs,
+                   const SerializedObj &SpecConsts,
+                   const std::string &BuildOptionsString);
 
   /* 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.
    */
   static std::vector<std::vector<char>>
-  getItemFromDisc(const device &Device, const RTDeviceBinaryImage &Img,
+  getItemFromDisc(const device &Device,
+                  const std::vector<const RTDeviceBinaryImage *> &Imgs,
                   const SerializedObj &SpecConsts,
                   const std::string &BuildOptionsString);
 
-  /* Stores build program in persisten cache
+  /* Stores build program in persistent cache
    */
-  static void putItemToDisc(const device &Device,
-                            const RTDeviceBinaryImage &Img,
-                            const SerializedObj &SpecConsts,
-                            const std::string &BuildOptionsString,
-                            const sycl::detail::pi::PiProgram &NativePrg);
+  static void
+  putItemToDisc(const device &Device,
+                const std::vector<const RTDeviceBinaryImage *> &Imgs,
+                const SerializedObj &SpecConsts,
+                const std::string &BuildOptionsString,
+                const sycl::detail::pi::PiProgram &NativePrg);
 
   /* Sends message to std:cerr stream when SYCL_CACHE_TRACE environemnt is set*/
   static void trace(const std::string &msg) {