https://github.com/maksimsab created 
https://github.com/llvm/llvm-project/pull/147508

None

>From 524da07f8449154a40796c5734d674df64e3f9af Mon Sep 17 00:00:00 2001
From: "Sabianin, Maksim" <maksim.sabia...@intel.com>
Date: Mon, 7 Jul 2025 08:30:12 -0700
Subject: [PATCH] [SYCL] Add offload wrapping for SYCL kind.

---
 clang/test/Driver/linker-wrapper-image.c      |  35 ++
 clang/test/Driver/linker-wrapper.c            |   2 +-
 .../tools/clang-linker-wrapper/CMakeLists.txt |   1 +
 .../ClangLinkerWrapper.cpp                    |  41 +-
 .../llvm/Frontend/SYCL/OffloadWrapper.h       |  44 ++
 llvm/include/llvm/Object/OffloadBinary.h      |   9 +-
 llvm/lib/Frontend/CMakeLists.txt              |   1 +
 llvm/lib/Frontend/SYCL/CMakeLists.txt         |  14 +
 llvm/lib/Frontend/SYCL/OffloadWrapper.cpp     | 513 ++++++++++++++++++
 llvm/lib/Object/OffloadBinary.cpp             |  11 +
 llvm/unittests/Object/OffloadingTest.cpp      |   9 +
 11 files changed, 675 insertions(+), 5 deletions(-)
 create mode 100644 llvm/include/llvm/Frontend/SYCL/OffloadWrapper.h
 create mode 100644 llvm/lib/Frontend/SYCL/CMakeLists.txt
 create mode 100644 llvm/lib/Frontend/SYCL/OffloadWrapper.cpp

diff --git a/clang/test/Driver/linker-wrapper-image.c 
b/clang/test/Driver/linker-wrapper-image.c
index c0de56d58196a..67bb21bfe49b4 100644
--- a/clang/test/Driver/linker-wrapper-image.c
+++ b/clang/test/Driver/linker-wrapper-image.c
@@ -1,6 +1,7 @@
 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 // REQUIRES: amdgpu-registered-target
+// REQUIRES: spirv-registered-target
 
 // RUN: %clang -cc1 %s -triple x86_64-unknown-linux-gnu -emit-obj -o %t.elf.o
 
@@ -263,3 +264,37 @@
 //      HIP: while.end:
 // HIP-NEXT:   ret void
 // HIP-NEXT: }
+
+// RUN: clang-offload-packager -o %t.out 
--image=file=%t.elf.o,kind=sycl,triple=spirv64-unknown-unknown,arch=generic
+// RUN: %clang -cc1 %s -triple x86_64-unknown-linux-gnu -emit-obj -o %t.o \
+// RUN:   -fembed-offload-object=%t.out
+// RUN: clang-linker-wrapper --print-wrapped-module --dry-run 
--host-triple=x86_64-unknown-linux-gnu \
+// RUN:   --linker-path=/usr/bin/ld %t.o -o a.out 2>&1 | FileCheck %s 
--check-prefixes=SYCL
+// RUN: clang-linker-wrapper --print-wrapped-module --dry-run 
--host-triple=x86_64-unknown-linux-gnu -r \
+// RUN:   --linker-path=/usr/bin/ld %t.o -o a.out 2>&1 | FileCheck %s 
--check-prefixes=SYCL
+
+//      SYCL: %__sycl.tgt_device_image = type { i16, i8, i8, ptr, ptr, ptr, 
ptr, ptr, ptr, ptr, ptr, ptr }
+// SYCL-NEXT: %__sycl.tgt_bin_desc = type { i16, i16, ptr, ptr, ptr }
+
+//      SYCL: @.sycl_offloading.target.0 = internal unnamed_addr constant [1 x 
i8] zeroinitializer
+// SYCL-NEXT: @.sycl_offloading.opts.compile.0 = internal unnamed_addr 
constant [1 x i8] zeroinitializer
+// SYCL-NEXT: @.sycl_offloading.opts.link.0 = internal unnamed_addr constant 
[1 x i8] zeroinitializer
+// SYCL-NEXT: @.sycl_offloading.0.data = internal unnamed_addr constant [0 x 
i8] zeroinitializer
+// SYCL-NEXT: @.sycl_offloading.0.info = internal local_unnamed_addr constant 
[2 x i64] [i64 ptrtoint (ptr @.sycl_offloading.0.data to i64), i64 0], section 
".tgtimg", align 16
+// SYCL-NEXT: @llvm.used = appending global [1 x ptr] [ptr 
@.sycl_offloading.0.info], section "llvm.metadata"
+// SYCL-NEXT: @.sycl_offloading.device_images = internal unnamed_addr constant 
[1 x %__sycl.tgt_device_image] [%__sycl.tgt_device_image { i16 3, i8 8, i8 0, 
ptr @.sycl_offloading.target.0, ptr @.sycl_offloading.opts.compile.0, ptr 
@.sycl_offloading.opts.link.0, ptr @.sycl_offloading.0.data, ptr 
@.sycl_offloading.0.data, ptr null, ptr null, ptr null, ptr null }]
+// SYCL-NEXT: @.sycl_offloading.descriptor = internal constant 
%__sycl.tgt_bin_desc { i16 1, i16 1, ptr @.sycl_offloading.device_images, ptr 
null, ptr null }
+// SYCL-NEXT: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ 
i32, ptr, ptr } { i32 1, ptr @sycl.descriptor_reg, ptr null }]
+// SYCL-NEXT: @llvm.global_dtors = appending global [1 x { i32, ptr, ptr }] [{ 
i32, ptr, ptr } { i32 1, ptr @sycl.descriptor_unreg, ptr null }]
+
+//      SYCL: define internal void @sycl.descriptor_reg() section 
".text.startup" {
+// SYCL-NEXT: entry:
+// SYCL-NEXT:   call void @__sycl_register_lib(ptr 
@.sycl_offloading.descriptor)
+// SYCL-NEXT:   ret void
+// SYCL-NEXT: }
+
+//      SYCL: define internal void @sycl.descriptor_unreg() section 
".text.startup" {
+// SYCL-NEXT: entry:
+// SYCL-NEXT:   call void @__sycl_unregister_lib(ptr 
@.sycl_offloading.descriptor)
+// SYCL-NEXT:   ret void
+// SYCL-NEXT: }
diff --git a/clang/test/Driver/linker-wrapper.c 
b/clang/test/Driver/linker-wrapper.c
index 80b1a5745a123..5ab8a09660e57 100644
--- a/clang/test/Driver/linker-wrapper.c
+++ b/clang/test/Driver/linker-wrapper.c
@@ -54,7 +54,7 @@ __attribute__((visibility("protected"), used)) int x;
 // RUN: clang-offload-packager -o %t.out \
 // RUN:   
--image=file=%t.spirv.bc,kind=sycl,triple=spirv64-unknown-unknown,arch=generic
 // RUN: %clang -cc1 %s -triple x86_64-unknown-linux-gnu -emit-obj -o %t.o 
-fembed-offload-object=%t.out
-// RUN: not clang-linker-wrapper --host-triple=x86_64-unknown-linux-gnu 
--dry-run \
+// RUN: clang-linker-wrapper --host-triple=x86_64-unknown-linux-gnu --dry-run \
 // RUN:   --linker-path=/usr/bin/ld %t.o -o a.out 2>&1 | FileCheck %s 
--check-prefix=SPIRV-LINK
 
 // SPIRV-LINK: clang{{.*}} -o {{.*}}.img --target=spirv64-unknown-unknown 
{{.*}}.o --sycl-link -Xlinker -triple=spirv64-unknown-unknown -Xlinker -arch=
diff --git a/clang/tools/clang-linker-wrapper/CMakeLists.txt 
b/clang/tools/clang-linker-wrapper/CMakeLists.txt
index bf37d8031025e..741e3fbbefb74 100644
--- a/clang/tools/clang-linker-wrapper/CMakeLists.txt
+++ b/clang/tools/clang-linker-wrapper/CMakeLists.txt
@@ -16,6 +16,7 @@ set(LLVM_LINK_COMPONENTS
   CodeGen
   LTO
   FrontendOffloading
+  FrontendSYCL
   )
 
 set(LLVM_TARGET_DEFINITIONS LinkerWrapperOpts.td)
diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp 
b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
index 0f1fa8b329fd6..9a466d6e69c31 100644
--- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
+++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
@@ -22,6 +22,7 @@
 #include "llvm/CodeGen/CommandFlags.h"
 #include "llvm/Frontend/Offloading/OffloadWrapper.h"
 #include "llvm/Frontend/Offloading/Utility.h"
+#include "llvm/Frontend/SYCL/OffloadWrapper.h"
 #include "llvm/IR/Constants.h"
 #include "llvm/IR/DiagnosticPrinter.h"
 #include "llvm/IR/Module.h"
@@ -711,6 +712,13 @@ wrapDeviceImages(ArrayRef<std::unique_ptr<MemoryBuffer>> 
Buffers,
             M, BuffersToWrap.front(), offloading::getOffloadEntryArray(M)))
       return std::move(Err);
     break;
+  case OFK_SYCL: {
+    offloading::sycl::SYCLWrappingOptions WrappingOptions;
+    if (Error Err = offloading::sycl::wrapSYCLBinaries(M, BuffersToWrap,
+                                                       WrappingOptions))
+      return Err;
+    break;
+  }
   default:
     return createStringError(getOffloadKindName(Kind) +
                              " wrapping is not supported");
@@ -748,6 +756,36 @@ bundleOpenMP(ArrayRef<OffloadingImage> Images) {
   return std::move(Buffers);
 }
 
+Expected<SmallVector<std::unique_ptr<MemoryBuffer>>>
+bundleSYCL(ArrayRef<OffloadingImage> Images) {
+  SmallVector<std::unique_ptr<MemoryBuffer>> Buffers;
+  if (DryRun) {
+    // In dry-run mode there is an empty input which is insufficient for
+    // the testing. Therefore, we insert a stub value.
+    OffloadBinary::OffloadingImage Image;
+    Image.TheOffloadKind = OffloadKind::OFK_SYCL;
+    Image.Image = MemoryBuffer::getMemBufferCopy("");
+    SmallString<0> SerializedImage = OffloadBinary::write(Image);
+    Buffers.emplace_back(MemoryBuffer::getMemBufferCopy(SerializedImage));
+    return Buffers;
+  }
+
+  for (const OffloadingImage &TheImage : Images) {
+    SmallVector<OffloadFile> OffloadBinaries;
+    if (Error E = extractOffloadBinaries(*TheImage.Image, OffloadBinaries))
+      return E;
+
+    for (const OffloadFile &File : OffloadBinaries) {
+      const OffloadBinary &Binary = *File.getBinary();
+      SmallString<0> SerializedImage =
+          OffloadBinary::write(Binary.getOffloadingImage());
+      Buffers.emplace_back(MemoryBuffer::getMemBufferCopy(SerializedImage));
+    }
+  }
+
+  return Buffers;
+}
+
 Expected<SmallVector<std::unique_ptr<MemoryBuffer>>>
 bundleCuda(ArrayRef<OffloadingImage> Images, const ArgList &Args) {
   SmallVector<std::pair<StringRef, StringRef>, 4> InputFiles;
@@ -800,8 +838,9 @@ bundleLinkedOutput(ArrayRef<OffloadingImage> Images, const 
ArgList &Args,
   llvm::TimeTraceScope TimeScope("Bundle linked output");
   switch (Kind) {
   case OFK_OpenMP:
-  case OFK_SYCL:
     return bundleOpenMP(Images);
+  case OFK_SYCL:
+    return bundleSYCL(Images);
   case OFK_Cuda:
     return bundleCuda(Images, Args);
   case OFK_HIP:
diff --git a/llvm/include/llvm/Frontend/SYCL/OffloadWrapper.h 
b/llvm/include/llvm/Frontend/SYCL/OffloadWrapper.h
new file mode 100644
index 0000000000000..f89411c86984d
--- /dev/null
+++ b/llvm/include/llvm/Frontend/SYCL/OffloadWrapper.h
@@ -0,0 +1,44 @@
+//===----- OffloadWrapper.h -------------------------------------*- C++ 
-*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM 
Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_FRONTEND_SYCL_OFFLOAD_WRAPPER_H
+#define LLVM_FRONTEND_SYCL_OFFLOAD_WRAPPER_H
+
+#include "llvm/ADT/ArrayRef.h"
+#include "llvm/Object/OffloadBinary.h"
+
+#include <string>
+
+namespace llvm {
+
+class Module;
+
+namespace offloading {
+namespace sycl {
+
+struct SYCLWrappingOptions {
+  // target/compiler specific options what are suggested to use to "compile"
+  // program at runtime.
+  std::string CompileOptions;
+  // Target/Compiler specific options that are suggested to use to "link"
+  // program at runtime.
+  std::string LinkOptions;
+};
+
+/// Wraps OffloadBinaries in the given \p Buffers into the module \p M
+/// as global symbols and registers the images with the SYCL Runtime.
+/// \param Options Settings that allows to turn on optional data and settings.
+llvm::Error
+wrapSYCLBinaries(llvm::Module &M, llvm::ArrayRef<llvm::ArrayRef<char>> Buffers,
+                 SYCLWrappingOptions Options = SYCLWrappingOptions());
+
+} // namespace sycl
+} // namespace offloading
+} // namespace llvm
+
+#endif // LLVM_FRONTEND_SYCL_OFFLOAD_WRAPPER_H
diff --git a/llvm/include/llvm/Object/OffloadBinary.h 
b/llvm/include/llvm/Object/OffloadBinary.h
index b5c845fa8eb70..9d137db834f08 100644
--- a/llvm/include/llvm/Object/OffloadBinary.h
+++ b/llvm/include/llvm/Object/OffloadBinary.h
@@ -48,6 +48,7 @@ enum ImageKind : uint16_t {
   IMG_Cubin,
   IMG_Fatbinary,
   IMG_PTX,
+  IMG_SPIRV,
   IMG_LAST,
 };
 
@@ -70,9 +71,9 @@ class OffloadBinary : public Binary {
 
   /// The offloading metadata that will be serialized to a memory buffer.
   struct OffloadingImage {
-    ImageKind TheImageKind;
-    OffloadKind TheOffloadKind;
-    uint32_t Flags;
+    ImageKind TheImageKind = ImageKind::IMG_None;
+    OffloadKind TheOffloadKind = OffloadKind::OFK_None;
+    uint32_t Flags = 0;
     MapVector<StringRef, StringRef> StringData;
     std::unique_ptr<MemoryBuffer> Image;
   };
@@ -84,6 +85,8 @@ class OffloadBinary : public Binary {
   /// Serialize the contents of \p File to a binary buffer to be read later.
   LLVM_ABI static SmallString<0> write(const OffloadingImage &);
 
+  OffloadingImage getOffloadingImage() const;
+
   static uint64_t getAlignment() { return 8; }
 
   ImageKind getImageKind() const { return TheEntry->TheImageKind; }
diff --git a/llvm/lib/Frontend/CMakeLists.txt b/llvm/lib/Frontend/CMakeLists.txt
index 3b31e6f8dec96..6c4b8362c04fd 100644
--- a/llvm/lib/Frontend/CMakeLists.txt
+++ b/llvm/lib/Frontend/CMakeLists.txt
@@ -5,3 +5,4 @@ add_subdirectory(HLSL)
 add_subdirectory(OpenACC)
 add_subdirectory(OpenMP)
 add_subdirectory(Offloading)
+add_subdirectory(SYCL)
diff --git a/llvm/lib/Frontend/SYCL/CMakeLists.txt 
b/llvm/lib/Frontend/SYCL/CMakeLists.txt
new file mode 100644
index 0000000000000..355ae5f7955a8
--- /dev/null
+++ b/llvm/lib/Frontend/SYCL/CMakeLists.txt
@@ -0,0 +1,14 @@
+add_llvm_component_library(LLVMFrontendSYCL
+  OffloadWrapper.cpp
+
+  ADDITIONAL_HEADER_DIRS
+  ${LLVM_MAIN_INCLUDE_DIR}/llvm/Frontend
+  ${LLVM_MAIN_INCLUDE_DIR}/llvm/Frontend/SYCL
+
+  LINK_COMPONENTS
+  Core
+  FrontendOffloading
+  Object
+  Support
+  TransformUtils
+  )
diff --git a/llvm/lib/Frontend/SYCL/OffloadWrapper.cpp 
b/llvm/lib/Frontend/SYCL/OffloadWrapper.cpp
new file mode 100644
index 0000000000000..c0d160c39e93d
--- /dev/null
+++ b/llvm/lib/Frontend/SYCL/OffloadWrapper.cpp
@@ -0,0 +1,513 @@
+//===- SYCLOffloadWrapper.cpp -----------------------------------*- C++ 
-*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM 
Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "llvm/Frontend/SYCL/OffloadWrapper.h"
+#include "llvm/ADT/ArrayRef.h"
+#include "llvm/ADT/SmallVector.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/ADT/Twine.h"
+#include "llvm/Frontend/Offloading/Utility.h"
+#include "llvm/IR/Constants.h"
+#include "llvm/IR/DerivedTypes.h"
+#include "llvm/IR/GlobalVariable.h"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/LLVMContext.h"
+#include "llvm/IR/Module.h"
+#include "llvm/IR/Type.h"
+#include "llvm/Object/OffloadBinary.h"
+#include "llvm/Support/Error.h"
+#include "llvm/Support/ErrorHandling.h"
+#include "llvm/Support/LineIterator.h"
+#include "llvm/Support/MemoryBufferRef.h"
+#include "llvm/Transforms/Utils/ModuleUtils.h"
+
+#include <memory>
+#include <string>
+#include <utility>
+
+using namespace llvm;
+using namespace llvm::object;
+using namespace llvm::offloading;
+using namespace llvm::offloading::sycl;
+
+using OffloadingImage = OffloadBinary::OffloadingImage;
+
+namespace {
+
+/// Wrapper helper class that creates all LLVM IRs wrapping given images.
+struct Wrapper {
+  Module &M;
+  LLVMContext &C;
+  SYCLWrappingOptions Options;
+
+  StructType *EntryTy = nullptr;
+  StructType *SyclDeviceImageTy = nullptr;
+  StructType *SyclBinDescTy = nullptr;
+
+  Wrapper(Module &M, const SYCLWrappingOptions &Options)
+      : M(M), C(M.getContext()), Options(Options) {
+
+    EntryTy = offloading::getEntryTy(M);
+    SyclDeviceImageTy = getSyclDeviceImageTy();
+    SyclBinDescTy = getSyclBinDescTy();
+  }
+
+  IntegerType *getSizeTTy() {
+    switch (M.getDataLayout().getPointerSize()) {
+    case 4:
+      return Type::getInt32Ty(C);
+    case 8:
+      return Type::getInt64Ty(C);
+    }
+    llvm_unreachable("unsupported pointer type size");
+  }
+
+  SmallVector<Constant *, 2> getSizetConstPair(size_t First, size_t Second) {
+    IntegerType *SizeTTy = getSizeTTy();
+    return SmallVector<Constant *, 2>{ConstantInt::get(SizeTTy, First),
+                                      ConstantInt::get(SizeTTy, Second)};
+  }
+
+  /// Note: Properties aren't supported and the support is going
+  /// to be added later.
+  /// Creates a structure corresponding to:
+  /// SYCL specific image descriptor type.
+  /// \code
+  /// struct __sycl.tgt_device_image {
+  ///   // version of this structure - for backward compatibility;
+  ///   // all modifications which change order/type/offsets of existing fields
+  ///   // should increment the version.
+  ///   uint16_t Version;
+  ///   // the kind of offload model the image employs.
+  ///   uint8_t OffloadKind;
+  ///   // format of the image data - SPIRV, LLVMIR bitcode, etc
+  ///   uint8_t Format;
+  ///   // null-terminated string representation of the device's target
+  ///   // architecture
+  ///   const char *Arch;
+  ///   // a null-terminated string; target- and compiler-specific options
+  ///   // which are suggested to use to "compile" program at runtime
+  ///   const char *CompileOptions;
+  ///   // a null-terminated string; target- and compiler-specific options
+  ///   // which are suggested to use to "link" program at runtime
+  ///   const char *LinkOptions;
+  ///   // Pointer to the device binary image start
+  ///   void *ImageStart;
+  ///   // Pointer to the device binary image end
+  ///   void *ImageEnd;
+  ///   // the entry table
+  ///   __tgt_offload_entry *EntriesBegin;
+  ///   __tgt_offload_entry *EntriesEnd;
+  ///   const char *PropertiesBegin;
+  ///   const char *PropertiesEnd;
+  /// };
+  /// \endcode
+  StructType *getSyclDeviceImageTy() {
+    return StructType::create(
+        {
+            Type::getInt16Ty(C),       // Version
+            Type::getInt8Ty(C),        // OffloadKind
+            Type::getInt8Ty(C),        // Format
+            PointerType::getUnqual(C), // Arch
+            PointerType::getUnqual(C), // CompileOptions
+            PointerType::getUnqual(C), // LinkOptions
+            PointerType::getUnqual(C), // ImageStart
+            PointerType::getUnqual(C), // ImageEnd
+            PointerType::getUnqual(C), // EntriesBegin
+            PointerType::getUnqual(C), // EntriesEnd
+            PointerType::getUnqual(C), // PropertiesBegin
+            PointerType::getUnqual(C)  // PropertiesEnd
+        },
+        "__sycl.tgt_device_image");
+  }
+
+  /// Creates a structure for SYCL specific binary descriptor type. Corresponds
+  /// to:
+  ///
+  /// \code
+  ///  struct __sycl.tgt_bin_desc {
+  ///    // version of this structure - for backward compatibility;
+  ///    // all modifications which change order/type/offsets of existing 
fields
+  ///    // should increment the version.
+  ///    uint16_t Version;
+  ///    uint16_t NumDeviceImages;
+  ///    __sycl.tgt_device_image *DeviceImages;
+  ///    // the offload entry table
+  ///    __tgt_offload_entry *HostEntriesBegin;
+  ///    __tgt_offload_entry *HostEntriesEnd;
+  ///  };
+  /// \endcode
+  StructType *getSyclBinDescTy() {
+    return StructType::create(
+        {Type::getInt16Ty(C), Type::getInt16Ty(C), PointerType::getUnqual(C),
+         PointerType::getUnqual(C), PointerType::getUnqual(C)},
+        "__sycl.tgt_bin_desc");
+  }
+
+  /// Adds a global readonly variable that is initialized by given
+  /// \p Initializer to the module.
+  GlobalVariable *addGlobalArrayVariable(const Twine &Name,
+                                         ArrayRef<char> Initializer,
+                                         const Twine &Section = "") {
+    auto *Arr = ConstantDataArray::get(M.getContext(), Initializer);
+    auto *Var = new GlobalVariable(M, Arr->getType(), /*isConstant*/ true,
+                                   GlobalVariable::InternalLinkage, Arr, Name);
+    Var->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
+
+    SmallVector<char, 32> NameBuf;
+    auto SectionName = Section.toStringRef(NameBuf);
+    if (!SectionName.empty())
+      Var->setSection(SectionName);
+    return Var;
+  }
+
+  /// Adds given \p Buf as a global variable into the module.
+  /// \returns Pair of pointers that point at the beginning and the end of the
+  /// variable.
+  std::pair<Constant *, Constant *>
+  addArrayToModule(ArrayRef<char> Buf, const Twine &Name,
+                   const Twine &Section = "") {
+    auto *Var = addGlobalArrayVariable(Name, Buf, Section);
+    auto *ImageB = ConstantExpr::getGetElementPtr(Var->getValueType(), Var,
+                                                  getSizetConstPair(0, 0));
+    auto *ImageE = ConstantExpr::getGetElementPtr(
+        Var->getValueType(), Var, getSizetConstPair(0, Buf.size()));
+    return std::make_pair(ImageB, ImageE);
+  }
+
+  /// Adds given \p Data as constant byte array in the module.
+  /// \returns Constant pointer to the added data. The pointer type does not
+  /// carry size information.
+  Constant *addRawDataToModule(ArrayRef<char> Data, const Twine &Name) {
+    auto *Var = addGlobalArrayVariable(Name, Data);
+    auto *DataPtr = ConstantExpr::getGetElementPtr(Var->getValueType(), Var,
+                                                   getSizetConstPair(0, 0));
+    return DataPtr;
+  }
+
+  /// Creates a global variable of const char* type and creates an
+  /// initializer that initializes it with \p Str.
+  ///
+  /// \returns Link-time constant pointer (constant expr) to that
+  /// variable.
+  Constant *addStringToModule(StringRef Str, const Twine &Name) {
+    auto *Arr = ConstantDataArray::getString(C, Str);
+    auto *Var = new GlobalVariable(M, Arr->getType(), /*isConstant*/ true,
+                                   GlobalVariable::InternalLinkage, Arr, Name);
+    Var->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
+    auto *Zero = ConstantInt::get(getSizeTTy(), 0);
+    Constant *ZeroZero[] = {Zero, Zero};
+    return ConstantExpr::getGetElementPtr(Var->getValueType(), Var, ZeroZero);
+  }
+
+  /// Creates a global variable of array of structs and initializes
+  /// it with the given values in \p ArrayData.
+  ///
+  /// \returns Pair of Constants that point at array content.
+  /// If \p ArrayData is empty then a returned pair contains nullptrs.
+  std::pair<Constant *, Constant *>
+  addStructArrayToModule(ArrayRef<Constant *> ArrayData, Type *ElemTy) {
+    if (ArrayData.empty()) {
+      auto *PtrTy = llvm::PointerType::getUnqual(ElemTy->getContext());
+      auto *NullPtr = Constant::getNullValue(PtrTy);
+      return std::make_pair(NullPtr, NullPtr);
+    }
+
+    assert(ElemTy == ArrayData[0]->getType() && "elem type mismatch");
+    auto *Arr =
+        ConstantArray::get(ArrayType::get(ElemTy, ArrayData.size()), 
ArrayData);
+    auto *ArrGlob = new GlobalVariable(M, Arr->getType(), /*isConstant*/ true,
+                                       GlobalVariable::InternalLinkage, Arr,
+                                       "__sycl_offload_prop_sets_arr");
+    auto *ArrB = ConstantExpr::getGetElementPtr(
+        ArrGlob->getValueType(), ArrGlob, getSizetConstPair(0, 0));
+    auto *ArrE =
+        ConstantExpr::getGetElementPtr(ArrGlob->getValueType(), ArrGlob,
+                                       getSizetConstPair(0, ArrayData.size()));
+    return std::pair<Constant *, Constant *>(ArrB, ArrE);
+  }
+
+  /// Creates a global variable that is initiazed with the given \p Entries.
+  ///
+  /// \returns Pair of Constants that point at entries content.
+  std::pair<Constant *, Constant *>
+  addOffloadEntriesToModule(StringRef Entries) {
+    if (Entries.empty()) {
+      auto *NullPtr = Constant::getNullValue(PointerType::getUnqual(C));
+      return std::pair<Constant *, Constant *>(NullPtr, NullPtr);
+    }
+
+    auto *I64Zero = ConstantInt::get(Type::getInt64Ty(C), 0);
+    auto *I32Zero = ConstantInt::get(Type::getInt32Ty(C), 0);
+    auto *NullPtr = Constant::getNullValue(PointerType::getUnqual(C));
+
+    SmallVector<Constant *> EntriesInits;
+    std::unique_ptr<MemoryBuffer> MB = MemoryBuffer::getMemBuffer(Entries);
+    for (line_iterator LI(*MB); !LI.is_at_eof(); ++LI) {
+      Constant *EntryData[] = {
+          ConstantExpr::getNullValue(Type::getInt64Ty(C)),
+          ConstantInt::get(Type::getInt16Ty(C), 1),
+          ConstantInt::get(Type::getInt16Ty(C), object::OffloadKind::OFK_SYCL),
+          I32Zero,
+          NullPtr,
+          addStringToModule(*LI, "__sycl_offload_entry_name"),
+          I64Zero,
+          I64Zero,
+          NullPtr};
+
+      EntriesInits.push_back(ConstantStruct::get(EntryTy, EntryData));
+    }
+
+    auto *Arr = ConstantArray::get(ArrayType::get(EntryTy, 
EntriesInits.size()),
+                                   EntriesInits);
+    auto *EntriesGV = new GlobalVariable(M, Arr->getType(), /*isConstant*/ 
true,
+                                         GlobalVariable::InternalLinkage, Arr,
+                                         "__sycl_offload_entries_arr");
+
+    auto *EntriesB = ConstantExpr::getGetElementPtr(
+        EntriesGV->getValueType(), EntriesGV, getSizetConstPair(0, 0));
+    auto *EntriesE = ConstantExpr::getGetElementPtr(
+        EntriesGV->getValueType(), EntriesGV,
+        getSizetConstPair(0, EntriesInits.size()));
+    return std::make_pair(EntriesB, EntriesE);
+  }
+
+  /// Emits a global array that contains \p Address and \P Size. Also add
+  /// it into llvm.used to force it to be emitted in the object file.
+  void emitRegistrationFunctions(Constant *Address, size_t Size,
+                                 const Twine &ImageID,
+                                 StringRef OffloadKindTag) {
+    Type *IntPtrTy = M.getDataLayout().getIntPtrType(C);
+    auto *ImgInfoArr =
+        ConstantArray::get(ArrayType::get(IntPtrTy, 2),
+                           {ConstantExpr::getPointerCast(Address, IntPtrTy),
+                            ConstantInt::get(IntPtrTy, Size)});
+    auto *ImgInfoVar = new GlobalVariable(
+        M, ImgInfoArr->getType(), true, GlobalVariable::InternalLinkage,
+        ImgInfoArr, Twine(OffloadKindTag) + ImageID + ".info");
+    ImgInfoVar->setAlignment(
+        MaybeAlign(M.getDataLayout().getTypeStoreSize(IntPtrTy) * 2u));
+    ImgInfoVar->setUnnamedAddr(GlobalValue::UnnamedAddr::Local);
+    ImgInfoVar->setSection(".tgtimg");
+
+    // Add image info to the used list to force it to be emitted to the
+    // object.
+    appendToUsed(M, ImgInfoVar);
+  }
+
+  Constant *wrapImage(const OffloadingImage &OI, const Twine &ImageID,
+                      StringRef OffloadKindTag) {
+    // Note: Intel DPC++ compiler had 2 versions of this structure
+    // and clang++ has a third different structure. To avoid ABI 
incompatibility
+    // between generated device images the Version here starts from 3.
+    constexpr uint16_t DeviceImageStructVersion = 3;
+    Constant *Version =
+        ConstantInt::get(Type::getInt16Ty(C), DeviceImageStructVersion);
+    Constant *OffloadKindConstant = ConstantInt::get(
+        Type::getInt8Ty(C), static_cast<uint8_t>(OI.TheOffloadKind));
+    Constant *ImageKindConstant = ConstantInt::get(
+        Type::getInt8Ty(C), static_cast<uint8_t>(OI.TheImageKind));
+    StringRef Triple = OI.StringData.lookup("triple");
+    Constant *TripleConstant =
+        addStringToModule(Triple, Twine(OffloadKindTag) + "target." + ImageID);
+    Constant *CompileOptions =
+        addStringToModule(Options.CompileOptions,
+                          Twine(OffloadKindTag) + "opts.compile." + ImageID);
+    Constant *LinkOptions = addStringToModule(
+        Options.LinkOptions, Twine(OffloadKindTag) + "opts.link." + ImageID);
+
+    // Note: NULL for now.
+    std::pair<Constant *, Constant *> PropertiesConstants = {
+        Constant::getNullValue(PointerType::getUnqual(C)),
+        Constant::getNullValue(PointerType::getUnqual(C))};
+
+    const MemoryBuffer &RawImage = *OI.Image;
+    std::pair<Constant *, Constant *> Binary = addArrayToModule(
+        ArrayRef<char>(RawImage.getBufferStart(), RawImage.getBufferEnd()),
+        Twine(OffloadKindTag) + ImageID + ".data", ".llvm.offloading");
+
+    // For SYCL images offload entries are defined here per image.
+    std::pair<Constant *, Constant *> ImageEntriesPtrs =
+        addOffloadEntriesToModule(OI.StringData.lookup("symbols"));
+    Constant *WrappedBinary = ConstantStruct::get(
+        SyclDeviceImageTy, Version, OffloadKindConstant, ImageKindConstant,
+        TripleConstant, CompileOptions, LinkOptions, Binary.first,
+        Binary.second, ImageEntriesPtrs.first, ImageEntriesPtrs.second,
+        PropertiesConstants.first, PropertiesConstants.second);
+
+    emitRegistrationFunctions(Binary.first, RawImage.getBufferSize(), ImageID,
+                              OffloadKindTag);
+
+    return WrappedBinary;
+  }
+
+  GlobalVariable *combineWrappedImages(ArrayRef<Constant *> WrappedImages,
+                                       StringRef OffloadKindTag) {
+    auto *ImagesData = ConstantArray::get(
+        ArrayType::get(SyclDeviceImageTy, WrappedImages.size()), 
WrappedImages);
+    auto *ImagesGV =
+        new GlobalVariable(M, ImagesData->getType(), /*isConstant*/ true,
+                           GlobalValue::InternalLinkage, ImagesData,
+                           Twine(OffloadKindTag) + "device_images");
+    ImagesGV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
+
+    auto *Zero = ConstantInt::get(getSizeTTy(), 0);
+    Constant *ZeroZero[] = {Zero, Zero};
+    auto *ImagesB = ConstantExpr::getGetElementPtr(ImagesGV->getValueType(),
+                                                   ImagesGV, ZeroZero);
+
+    Constant *EntriesB = Constant::getNullValue(PointerType::getUnqual(C));
+    Constant *EntriesE = Constant::getNullValue(PointerType::getUnqual(C));
+    static constexpr uint16_t BinDescStructVersion = 1;
+    auto *DescInit = ConstantStruct::get(
+        SyclBinDescTy,
+        ConstantInt::get(Type::getInt16Ty(C), BinDescStructVersion),
+        ConstantInt::get(Type::getInt16Ty(C), WrappedImages.size()), ImagesB,
+        EntriesB, EntriesE);
+
+    return new GlobalVariable(M, DescInit->getType(), /*isConstant*/ true,
+                              GlobalValue::InternalLinkage, DescInit,
+                              Twine(OffloadKindTag) + "descriptor");
+  }
+
+  /// Creates binary descriptor for the given device images. Binary descriptor
+  /// is an object that is passed to the offloading runtime at program startup
+  /// and it describes all device images available in the executable or shared
+  /// library. It is defined as follows:
+  ///
+  /// \code
+  /// __attribute__((visibility("hidden")))
+  /// extern __tgt_offload_entry *__start_offloading_entries0;
+  /// __attribute__((visibility("hidden")))
+  /// extern __tgt_offload_entry *__stop_offloading_entries0;
+  /// ...
+  ///
+  /// __attribute__((visibility("hidden")))
+  /// extern const char *CompileOptions = "...";
+  /// ...
+  /// __attribute__((visibility("hidden")))
+  /// extern const char *LinkOptions = "...";
+  /// ...
+  ///
+  /// static const char Image0[] = { ... };
+  ///  ...
+  /// static const char ImageN[] = { ... };
+  ///
+  /// static const __sycl.tgt_device_image Images[] = {
+  ///   {
+  ///     Version,                      // Version
+  ///     OffloadKind,                  // OffloadKind
+  ///     Format,                       // format of the image - SPIRV, LLVMIR
+  ///                                   // bc, etc
+  //      TripleString,                 // Arch
+  ///     CompileOptions0,              // CompileOptions
+  ///     LinkOptions0,                 // LinkOptions
+  ///     Image0,                       // ImageStart
+  ///     Image0 + N,                   // ImageEnd
+  ///     __start_offloading_entries0,  // EntriesBegin
+  ///     __stop_offloading_entries0,   // EntriesEnd
+  ///     NULL,                         // PropertiesBegin
+  ///     NULL,                         // PropertiesEnd
+  ///   },
+  ///   ...
+  /// };
+  ///
+  /// static const __sycl.tgt_bin_desc FatbinDesc = {
+  ///   Version,                             //Version
+  ///   sizeof(Images) / sizeof(Images[0]),  //NumDeviceImages
+  ///   Images,                              //DeviceImages
+  ///   NULL,                                //HostEntriesBegin
+  ///   NULL                                 //HostEntriesEnd
+  /// };
+  /// \endcode
+  ///
+  /// \returns Global variable that represents FatbinDesc.
+  GlobalVariable *createFatbinDesc(ArrayRef<OffloadingImage> Images) {
+    StringRef OffloadKindTag = ".sycl_offloading.";
+    SmallVector<Constant *> WrappedImages;
+    WrappedImages.reserve(Images.size());
+    for (size_t I = 0, E = Images.size(); I != E; ++I)
+      WrappedImages.push_back(wrapImage(Images[I], Twine(I), OffloadKindTag));
+
+    return combineWrappedImages(WrappedImages, OffloadKindTag);
+  }
+
+  void createRegisterFatbinFunction(GlobalVariable *FatbinDesc) {
+    auto *FuncTy = FunctionType::get(Type::getVoidTy(C), /*isVarArg*/ false);
+    auto *Func = Function::Create(FuncTy, GlobalValue::InternalLinkage,
+                                  Twine("sycl") + ".descriptor_reg", &M);
+    Func->setSection(".text.startup");
+
+    // Get RegFuncName function declaration.
+    auto *RegFuncTy =
+        FunctionType::get(Type::getVoidTy(C), PointerType::getUnqual(C),
+                          /*isVarArg=*/false);
+    FunctionCallee RegFuncC =
+        M.getOrInsertFunction("__sycl_register_lib", RegFuncTy);
+
+    // Construct function body
+    IRBuilder Builder(BasicBlock::Create(C, "entry", Func));
+    Builder.CreateCall(RegFuncC, FatbinDesc);
+    Builder.CreateRetVoid();
+
+    // Add this function to constructors.
+    appendToGlobalCtors(M, Func, /*Priority*/ 1);
+  }
+
+  void createUnregisterFunction(GlobalVariable *FatbinDesc) {
+    auto *FuncTy = FunctionType::get(Type::getVoidTy(C), /*isVarArg*/ false);
+    auto *Func = Function::Create(FuncTy, GlobalValue::InternalLinkage,
+                                  "sycl.descriptor_unreg", &M);
+    Func->setSection(".text.startup");
+
+    // Get UnregFuncName function declaration.
+    auto *UnRegFuncTy =
+        FunctionType::get(Type::getVoidTy(C), PointerType::getUnqual(C),
+                          /*isVarArg=*/false);
+    FunctionCallee UnRegFuncC =
+        M.getOrInsertFunction("__sycl_unregister_lib", UnRegFuncTy);
+
+    // Construct function body
+    IRBuilder<> Builder(BasicBlock::Create(C, "entry", Func));
+    Builder.CreateCall(UnRegFuncC, FatbinDesc);
+    Builder.CreateRetVoid();
+
+    // Add this function to global destructors.
+    appendToGlobalDtors(M, Func, /*Priority*/ 1);
+  }
+}; // end of Wrapper
+
+} // anonymous namespace
+
+Error llvm::offloading::sycl::wrapSYCLBinaries(llvm::Module &M,
+                                               ArrayRef<ArrayRef<char>> 
Buffers,
+                                               SYCLWrappingOptions Options) {
+  Wrapper W(M, Options);
+  SmallVector<std::unique_ptr<OffloadBinary>> OffloadBinaries;
+  OffloadBinaries.reserve(Buffers.size());
+  SmallVector<OffloadingImage> Images;
+  Images.reserve(Buffers.size());
+  for (auto Buf : Buffers) {
+    MemoryBufferRef MBR(StringRef(Buf.begin(), Buf.size()), /*Identifier*/ "");
+    auto OffloadBinaryOrErr = OffloadBinary::create(MBR);
+    if (!OffloadBinaryOrErr)
+      return OffloadBinaryOrErr.takeError();
+
+    OffloadBinaries.emplace_back(std::move(*OffloadBinaryOrErr));
+    Images.emplace_back(OffloadBinaries.back()->getOffloadingImage());
+  }
+
+  GlobalVariable *Desc = W.createFatbinDesc(Images);
+  if (!Desc)
+    return createStringError(inconvertibleErrorCode(),
+                             "No binary descriptors created.");
+
+  W.createRegisterFatbinFunction(Desc);
+  W.createUnregisterFunction(Desc);
+  return Error::success();
+}
diff --git a/llvm/lib/Object/OffloadBinary.cpp 
b/llvm/lib/Object/OffloadBinary.cpp
index 3fff6b6a09e08..ac0a74b34f9ee 100644
--- a/llvm/lib/Object/OffloadBinary.cpp
+++ b/llvm/lib/Object/OffloadBinary.cpp
@@ -266,6 +266,17 @@ SmallString<0> OffloadBinary::write(const OffloadingImage 
&OffloadingData) {
   return Data;
 }
 
+OffloadBinary::OffloadingImage OffloadBinary::getOffloadingImage() const {
+  OffloadingImage OI;
+  OI.TheImageKind = getImageKind();
+  OI.TheOffloadKind = getOffloadKind();
+  OI.Flags = getFlags();
+  OI.StringData = StringData;
+  OI.Image = MemoryBuffer::getMemBuffer(
+      MemoryBufferRef(getImage(), /*Identifier*/ ""));
+  return OI;
+}
+
 Error object::extractOffloadBinaries(MemoryBufferRef Buffer,
                                      SmallVectorImpl<OffloadFile> &Binaries) {
   file_magic Type = identify_magic(Buffer.getBuffer());
diff --git a/llvm/unittests/Object/OffloadingTest.cpp 
b/llvm/unittests/Object/OffloadingTest.cpp
index 18c9efaceed06..bae09028dba77 100644
--- a/llvm/unittests/Object/OffloadingTest.cpp
+++ b/llvm/unittests/Object/OffloadingTest.cpp
@@ -64,4 +64,13 @@ TEST(OffloadingTest, checkOffloadingBinary) {
   // Ensure the size and alignment of the data is correct.
   EXPECT_TRUE(Binary.getSize() % OffloadBinary::getAlignment() == 0);
   EXPECT_TRUE(Binary.getSize() == BinaryBuffer->getBuffer().size());
+
+  OffloadBinary::OffloadingImage OI = Binary.getOffloadingImage();
+  ASSERT_EQ(Data.TheImageKind, OI.TheImageKind);
+  ASSERT_EQ(Data.TheOffloadKind, OI.TheOffloadKind);
+  ASSERT_EQ(Data.Flags, OI.Flags);
+  ASSERT_EQ(Data.Image->getBuffer(), OI.Image->getBuffer());
+  for (const auto &KeyAndValue : Data.StringData)
+    ASSERT_EQ(Data.StringData[KeyAndValue.first],
+              OI.StringData[KeyAndValue.first]);
 }

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to