From 524da07f8449154a40796c5734d674df64e3f9af Mon Sep 17 00:00:00 2001 From: "Sabianin, Maksim" 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> 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 Images) { return std::move(Buffers); } +Expected>> +bundleSYCL(ArrayRef Images) { + SmallVector> 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 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>> bundleCuda(ArrayRef Images, const ArgList &Args) { SmallVector, 4> InputFiles; @@ -800,8 +838,9 @@ bundleLinkedOutput(ArrayRef 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 + +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> 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 StringData; std::unique_ptr 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 +#include +#include + +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 getSizetConstPair(size_t First, size_t Second) { + IntegerType *SizeTTy = getSizeTTy(); + return SmallVector{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 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 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 + addArrayToModule(ArrayRef 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 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 + addStructArrayToModule(ArrayRef 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(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 + addOffloadEntriesToModule(StringRef Entries) { + if (Entries.empty()) { + auto *NullPtr = Constant::getNullValue(PointerType::getUnqual(C)); + return std::pair(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 EntriesInits; + std::unique_ptr 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(OI.TheOffloadKind)); + Constant *ImageKindConstant = ConstantInt::get( + Type::getInt8Ty(C), static_cast(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 PropertiesConstants = { + Constant::getNullValue(PointerType::getUnqual(C)), + Constant::getNullValue(PointerType::getUnqual(C))}; + + const MemoryBuffer &RawImage = *OI.Image; + std::pair Binary = addArrayToModule( + ArrayRef(RawImage.getBufferStart(), RawImage.getBufferEnd()), + Twine(OffloadKindTag) + ImageID + ".data", ".llvm.offloading"); + + // For SYCL images offload entries are defined here per image. + std::pair 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 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 Images) { + StringRef OffloadKindTag = ".sycl_offloading."; + SmallVector 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> Buffers, + SYCLWrappingOptions Options) { + Wrapper W(M, Options); + SmallVector> OffloadBinaries; + OffloadBinaries.reserve(Buffers.size()); + SmallVector 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 &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]); }