llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang-driver @llvm/pr-subscribers-clang Author: Maksim Sabianin (maksimsab) <details> <summary>Changes</summary> --- Patch is 34.84 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/147508.diff 11 Files Affected: - (modified) clang/test/Driver/linker-wrapper-image.c (+35) - (modified) clang/test/Driver/linker-wrapper.c (+1-1) - (modified) clang/tools/clang-linker-wrapper/CMakeLists.txt (+1) - (modified) clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp (+40-1) - (added) llvm/include/llvm/Frontend/SYCL/OffloadWrapper.h (+44) - (modified) llvm/include/llvm/Object/OffloadBinary.h (+6-3) - (modified) llvm/lib/Frontend/CMakeLists.txt (+1) - (added) llvm/lib/Frontend/SYCL/CMakeLists.txt (+14) - (added) llvm/lib/Frontend/SYCL/OffloadWrapper.cpp (+513) - (modified) llvm/lib/Object/OffloadBinary.cpp (+11) - (modified) llvm/unittests/Object/OffloadingTest.cpp (+9) ``````````diff 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)... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/147508 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
