diff --git a/clang/test/Driver/clang-linker-wrapper-zstd.cpp b/clang/test/Driver/clang-linker-wrapper-zstd.cpp new file mode 100644 index 0000000000000..4d22ce735442b --- /dev/null +++ b/clang/test/Driver/clang-linker-wrapper-zstd.cpp @@ -0,0 +1,77 @@ +// REQUIRES: zstd && system-linux && x86-registered-target + +// clang-linker-wrapper compression test: checks that the wrapper compresses +// SYCL device images when --compress is set, tags them with +// SYCLBinaryImageFormat::BIF_Compressed (int8 value 4) in the emitted wrapper +// module, honors --compression-level=, and reports invalid values. + +// Prepare test data. The bitcode is over the 512-byte compression +// threshold, so the compression actually runs. +// RUN:%clang_cc1 -fsycl-is-device -disable-llvm-passes -triple=spir64-unknown-unknown %s -emit-llvm-bc -o %t.device.bc +// RUN: llvm-offload-binary -o %t.fat --image=file=%t.device.bc,kind=sycl,triple=spir64-unknown-unknown +// RUN:%clang_cc1 %s -triple=x86_64-unknown-linux-gnu -emit-obj -o %t.o -fembed-offload-object=%t.fat +// RUN: touch %t.devicelib.bc + +// With --compress and --wrapper-verbose the compression path fires and the +// wrapped image's Format slot in %__sycl.tgt_device_image is i8 4 +// (BIF_Compressed). +// RUN: clang-linker-wrapper --print-wrapped-module --host-triple=x86_64-unknown-linux-gnu \ +// RUN: --bitcode-library=spir64-unknown-unknown=%t.devicelib.bc \ +// RUN: -sycl-post-link-options="-split=auto -symbols -properties" \ +// RUN: --compress --compression-level=9 --wrapper-verbose \ +// RUN: --dry-run %t.o -o %t.out 2>&1 \ +// RUN: | FileCheck %s --check-prefix=CHECK-COMPRESS + +// Capture the original and compressed sizes. +// CHECK-COMPRESS: [Compression] Original image size: [[#ORIG:]] +// CHECK-COMPRESS: [Compression] Compressed image size: [[#COMP:]] +// CHECK-COMPRESS: [Compression] Compression level used: 9 +// +// Assert ORIG - COMP > 0. +// FileCheck's numeric matching form only supports the +// `==` constraint, but sub() rejects underflow at +// expression-evaluation time. We attach sub(ORIG, COMP) to a CHECK-NOT +// pattern that begins with a sentinel string never present in the output: +// * when ORIG > COMP, sub() succeeds; the substituted pattern is +// "COMPRESSION_SIZE_CHECK", which doesn't appear, so the +// CHECK-NOT is satisfied. +// * when COMP >= ORIG, sub() underflows and FileCheck fails the pattern. +// CHECK-COMPRESS-NOT: COMPRESSION_SIZE_CHECK[[#sub(ORIG, COMP)]] +// CHECK-COMPRESS: @.sycl_offloading.device_images = internal unnamed_addr constant [1 x %__sycl.tgt_device_image] [%__sycl.tgt_device_image { i16 {{[0-9]+}}, i8 4, i8 4, + +// Without --compress the image is left untagged (Format = BIF_None = 0) and +// no [Compression] verbose lines are emitted. +// RUN: clang-linker-wrapper --print-wrapped-module --host-triple=x86_64-unknown-linux-gnu \ +// RUN: --bitcode-library=spir64-unknown-unknown=%t.devicelib.bc \ +// RUN: -sycl-post-link-options="-split=auto -symbols -properties" \ +// RUN: --wrapper-verbose \ +// RUN: %t.o -o %t.out --linker-path="/usr/bin/ld" 2>&1 \ +// RUN: | FileCheck %s --check-prefix=CHECK-NO-COMPRESS + +// CHECK-NO-COMPRESS-NOT: [Compression] +// CHECK-NO-COMPRESS: @.sycl_offloading.device_images = internal unnamed_addr constant [1 x %__sycl.tgt_device_image] [%__sycl.tgt_device_image { i16 {{[0-9]+}}, i8 4, i8 0, + +// A non-integer --compression-level= is diagnosed. +// RUN: not clang-linker-wrapper --host-triple=x86_64-unknown-linux-gnu \ +// RUN: --bitcode-library=spir64-unknown-unknown=%t.devicelib.bc \ +// RUN: -sycl-post-link-options="-split=auto -symbols -properties" \ +// RUN: --compress --compression-level=notanumber \ +// RUN: %t.o -o %t.out --linker-path="/usr/bin/ld" 2>&1 \ +// RUN: | FileCheck %s --check-prefix=CHECK-BAD-LEVEL + +// CHECK-BAD-LEVEL: invalid value for --offload-compression-level=: 'notanumber' + +template +__attribute__((sycl_kernel)) void kernel(const Func &func) { + func(); +} + +extern "C" { +// Symbols so the linker doesn't fail resolving them. +void __sycl_register_lib(void *) {} +void __sycl_unregister_lib(void *) {} +} + +int main() { + kernel([](){}); +} diff --git a/clang/test/Driver/sycl-offload-new-driver-compression.cpp b/clang/test/Driver/sycl-offload-new-driver-compression.cpp new file mode 100644 index 0000000000000..244dcf9e048e1 --- /dev/null +++ b/clang/test/Driver/sycl-offload-new-driver-compression.cpp @@ -0,0 +1,26 @@ +/// Check that '--offload-compress' and '--offload-compression-level=' are +/// forwarded to clang-linker-wrapper under the new offloading driver via the +/// existing addOffloadCompressArgs() plumbing (as '--compress' and +/// '--compression-level='). The wrapper consumes them in +/// wrapSYCLBinariesFromFile to zstd-compress each SYCL device image and tag +/// it with BIF_Compressed. + +// RUN: %clangxx -### -fsycl --offload-new-driver --offload-compress \ +// RUN: --offload-compression-level=3 %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=CHECK-COMPRESS +// CHECK-COMPRESS: {{.*}}clang-linker-wrapper{{.*}}"--compress"{{.*}}"--compression-level=3" + +// RUN: %clangxx -### -fsycl --offload-new-driver %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=CHECK-NO-COMPRESS +// CHECK-NO-COMPRESS-NOT: {{.*}}clang-linker-wrapper{{.*}}"--compress" + +// --no-offload-compress overrides an earlier --offload-compress. +// RUN: %clangxx -### -fsycl --offload-new-driver --offload-compress \ +// RUN: --no-offload-compress %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=CHECK-NO-COMPRESS + +// ...and vice versa: a later --offload-compress wins over --no-offload-compress. +// RUN: %clangxx -### -fsycl --offload-new-driver --no-offload-compress \ +// RUN: --offload-compress %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=CHECK-COMPRESS-LAST +// CHECK-COMPRESS-LAST: {{.*}}clang-linker-wrapper{{.*}}"--compress" diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index bf1d3b19add9c..f3fb9ed92f3d7 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1107,6 +1107,41 @@ static Expected runAOTCompile(StringRef InputFile, "Unsupported SYCL Triple and Arch"); } +/// Compress each SYCL device image in-place when --compress is set, +/// tagging its Format as BIF_Compressed. --compress and --compression-level= +/// are the same flags HIP forwards to clang-offload-bundler; +static Error compressImages(SmallVectorImpl &Images, + const ArgList &Args) { + if (!Args.hasArg(OPT_compress)) + return Error::success(); + + int Level = llvm::offloading::DefaultSYCLCompressionLevel; + if (auto *A = Args.getLastArg(OPT_compression_level_eq)) + if (StringRef(A->getValue()).getAsInteger(10, Level)) + return createStringError( + "invalid value for --offload-compression-level=: '%s'", + A->getValue()); + + for (auto &Image : Images) { + SmallVector CompressedBytes; + Expected DidCompressOrErr = offloading::compressSYCLDeviceImage( + ArrayRef( + reinterpret_cast(Image.Image->getBufferStart()), + Image.Image->getBufferSize()), + CompressedBytes, Level, /*Threshold=*/512, Verbose); + if (!DidCompressOrErr) + return DidCompressOrErr.takeError(); + if (!*DidCompressOrErr) + continue; + Image.Image = MemoryBuffer::getMemBufferCopy( + StringRef(reinterpret_cast(CompressedBytes.data()), + CompressedBytes.size()), + Image.Image->getBufferIdentifier()); + Image.Format = offloading::SYCLBinaryImageFormat::BIF_Compressed; + } + return Error::success(); +} + /// Reads device images from the given \p InputFile and wraps them /// in one LLVM IR Module as a constant data. /// @@ -1178,6 +1213,9 @@ wrapSYCLBinariesFromFile(ArrayRef SplitModules, ImageTarget, SI.CompileOptions, SI.LinkOptions); } + if (Error E = compressImages(Images, Args)) + return std::move(E); + LLVMContext C; Module M("offload.wrapper.object", C); M.setTargetTriple(Triple( diff --git a/clang/tools/clang-offload-wrapper/CMakeLists.txt b/clang/tools/clang-offload-wrapper/CMakeLists.txt index 3195f18fe23cd..561ffc59a85aa 100644 --- a/clang/tools/clang-offload-wrapper/CMakeLists.txt +++ b/clang/tools/clang-offload-wrapper/CMakeLists.txt @@ -1,4 +1,4 @@ -set(LLVM_LINK_COMPONENTS BitWriter BitReader Core Object Support TransformUtils TargetParser) +set(LLVM_LINK_COMPONENTS BitWriter BitReader Core FrontendOffloading Object Support TransformUtils TargetParser) add_clang_tool(clang-offload-wrapper ClangOffloadWrapper.cpp diff --git a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp index 3129dcff42177..344f483e785ab 100644 --- a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp +++ b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp @@ -21,6 +21,7 @@ #include "llvm/ADT/Twine.h" #include "llvm/BinaryFormat/ELF.h" #include "llvm/Bitcode/BitcodeWriter.h" +#include "llvm/Frontend/Offloading/Utility.h" #include "llvm/IR/Constant.h" #include "llvm/IR/Constants.h" #include "llvm/IR/DerivedTypes.h" @@ -67,9 +68,6 @@ #include #include -// For device image compression. -#include - #define OPENMP_OFFLOAD_IMAGE_VERSION "1.0" using namespace llvm; @@ -84,8 +82,13 @@ static constexpr char COL_PROPS[] = "Properties"; // Offload models supported by this tool. The support basically means mapping // a string representation given at the command line to a value from this -// enum. -enum OffloadKind { +// enum. Distinct from llvm::object::OffloadKind (in Object/OffloadBinary.h), +// which is the 16-bit bitmask stored in the OffloadBinary header +// (OFK_None/OFK_OpenMP/... = powers of two). This enum is a sequential +// command-line-parsed tag with an extra Host slot and First/Last sentinels, +// serialized as a uint8_t into the SYCL runtime's sycl_device_binary_struct +// Kind field. +enum OldOffloadKind { Unknown = 0, Host, OpenMP, @@ -96,12 +99,12 @@ enum OffloadKind { }; namespace llvm { -template <> struct DenseMapInfo { - static unsigned getHashValue(const OffloadKind &Val) { +template <> struct DenseMapInfo { + static unsigned getHashValue(const OldOffloadKind &Val) { return DenseMapInfo::getHashValue(static_cast(Val)); } - static bool isEqual(const OffloadKind &LHS, const OffloadKind &RHS) { + static bool isEqual(const OldOffloadKind &LHS, const OldOffloadKind &RHS) { return LHS == RHS; } }; @@ -170,7 +173,7 @@ enum BinaryImageFormat { }; /// Sets offload kind. -static cl::list Kinds( +static cl::list Kinds( "kind", cl::desc("offload kind:"), cl::OneOrMore, cl::values(clEnumValN(Unknown, "unknown", "unknown"), clEnumValN(Host, "host", "host"), @@ -256,17 +259,17 @@ static cl::opt BatchMode( cl::cat(ClangOffloadWrapperCategory)); // clang-format on -static StringRef offloadKindToString(OffloadKind Kind) { +static StringRef offloadKindToString(OldOffloadKind Kind) { switch (Kind) { - case OffloadKind::Unknown: + case OldOffloadKind::Unknown: return "unknown"; - case OffloadKind::Host: + case OldOffloadKind::Host: return "host"; - case OffloadKind::OpenMP: + case OldOffloadKind::OpenMP: return "openmp"; - case OffloadKind::HIP: + case OldOffloadKind::HIP: return "hip"; - case OffloadKind::SYCL: + case OldOffloadKind::SYCL: return "sycl"; } llvm_unreachable("bad offload kind"); @@ -357,12 +360,12 @@ class BinaryWrapper { PointerType *PtrTy = nullptr; /// Records all added device binary images per offload kind. - llvm::DenseMap> Packs; + llvm::DenseMap> Packs; /// Records all created memory buffers for safe auto-gc llvm::SmallVector, 4> AutoGcBufs; public: - void addImage(const OffloadKind Kind, llvm::StringRef File, + void addImage(const OldOffloadKind Kind, llvm::StringRef File, llvm::StringRef Tgt, const BinaryImageFormat Fmt, llvm::StringRef CompileOpts, llvm::StringRef LinkOpts, llvm::StringRef EntriesFile, llvm::StringRef PropsFile) { @@ -524,7 +527,7 @@ class BinaryWrapper { // /// should increment the version. // uint16_t Version; // /// the kind of offload model the image employs. - // uint8_t OffloadKind; + // uint8_t OldOffloadKind; // /// format of the image data - SPIRV, LLVMIR bitcode,... // uint8_t Format; // /// null-terminated string representation of the device's target @@ -553,7 +556,7 @@ class BinaryWrapper { SyclImageTy = StructType::create( { Type::getInt16Ty(C), // Version - Type::getInt8Ty(C), // OffloadKind + Type::getInt8Ty(C), // OldOffloadKind Type::getInt8Ty(C), // Format getPtrTy(), // DeviceTargetSpec getPtrTy(), // CompileOptions @@ -722,7 +725,7 @@ class BinaryWrapper { // contains the image data. std::pair addDeviceImageToModule(ArrayRef Buf, const Twine &Name, - OffloadKind Kind, StringRef TargetTriple) { + OldOffloadKind Kind, StringRef TargetTriple) { // Create global variable for the image data. return addArrayToModule(Buf, Name, TargetTriple.empty() @@ -961,14 +964,14 @@ class BinaryWrapper { /// }; /// /// Global variable that represents BinDesc is returned. - Expected createBinDesc(OffloadKind Kind, + Expected createBinDesc(OldOffloadKind Kind, SameKindPack &Pack) { - const std::string OffloadKindTag = + const std::string OldOffloadKindTag = (Twine(".") + offloadKindToString(Kind) + Twine("_offloading.")).str(); Constant *EntriesB = nullptr, *EntriesE = nullptr; - if (Kind != OffloadKind::SYCL) { + if (Kind != OldOffloadKind::SYCL) { // Create external begin/end symbols for the offload entries table. auto *EntriesStart = new GlobalVariable( M, getEntryTy(), /*isConstant*/ true, GlobalValue::ExternalLinkage, @@ -1023,13 +1026,13 @@ class BinaryWrapper { auto *Fknd = ConstantInt::get(Type::getInt8Ty(C), Kind); auto *Ffmt = ConstantInt::get(Type::getInt8Ty(C), Img.Fmt); auto *Ftgt = addStringToModule( - Img.Tgt, Twine(OffloadKindTag) + Twine("target.") + Twine(ImgId)); + Img.Tgt, Twine(OldOffloadKindTag) + Twine("target.") + Twine(ImgId)); auto *Foptcompile = addStringToModule( Img.CompileOpts, - Twine(OffloadKindTag) + Twine("opts.compile.") + Twine(ImgId)); - auto *Foptlink = addStringToModule(Img.LinkOpts, Twine(OffloadKindTag) + - Twine("opts.link.") + - Twine(ImgId)); + Twine(OldOffloadKindTag) + Twine("opts.compile.") + Twine(ImgId)); + auto *Foptlink = addStringToModule( + Img.LinkOpts, + Twine(OldOffloadKindTag) + Twine("opts.link.") + Twine(ImgId)); if (MySymPropReader) MySymPropReader->getNextDeviceImageInitializer(); @@ -1046,7 +1049,7 @@ class BinaryWrapper { if (!BinOrErr) return BinOrErr.takeError(); MemoryBuffer *Bin = *BinOrErr; - if (Img.File != "-" && Kind == OffloadKind::OpenMP && + if (Img.File != "-" && Kind == OldOffloadKind::OpenMP && AddOpenMPOffloadNotes) { // Adding ELF notes for STDIN is not supported yet. Bin = addELFNotes(Bin, Img.File); @@ -1059,72 +1062,42 @@ class BinaryWrapper { Fbin = *FBinOrErr; } else { - // If '--offload-compress' option is specified and zstd is not - // available, throw an error. - if (OffloadCompressDevImgs && !llvm::compression::zstd::isAvailable()) { - return createStringError( - inconvertibleErrorCode(), - "'--offload-compress' is specified but the compiler is " - "built without zstd support.\n" - "If you are using a custom DPC++ build, please refer to " - "https://github.com/intel/llvm/blob/sycl/sycl/doc/" - "GetStartedGuide.md#build-dpc-toolchain-with-device-image-" - "compression-support" - " for more information on how to build with zstd support."); + // Compress only for SYCL images that don't have an explicitly-pinned + // binary image format. compressSYCLDeviceImage() surfaces the + // zstd-unavailable error when it actually attempts compression, so + // there's no separate up-front check. + SmallVector CompressedBuffer; + bool DidCompress = false; + if (OffloadCompressDevImgs && Kind == OldOffloadKind::SYCL && + Img.Fmt == BinaryImageFormat::none) { + Expected Compressed = offloading::compressSYCLDeviceImage( + ArrayRef( + reinterpret_cast(Bin->getBufferStart()), + Bin->getBufferSize()), + CompressedBuffer, OffloadCompressLevel, OffloadCompressThreshold, + Verbose); + if (!Compressed) + return Compressed.takeError(); + DidCompress = *Compressed; } - // Don't compress if the user explicitly specifies the binary image - // format or if the image is smaller than OffloadCompressThreshold - // bytes. - if (Kind != OffloadKind::SYCL || !OffloadCompressDevImgs || - Img.Fmt != BinaryImageFormat::none || - !llvm::compression::zstd::isAvailable() || - static_cast(Bin->getBufferSize()) < OffloadCompressThreshold) { - Fbin = addDeviceImageToModule( - ArrayRef(Bin->getBufferStart(), Bin->getBufferSize()), - Twine(OffloadKindTag) + Twine(ImgId) + Twine(".data"), Kind, - Img.Tgt); - } else { - - // Compress the image using zstd. - SmallVector CompressedBuffer; -#if LLVM_ENABLE_EXCEPTIONS - try { -#endif - llvm::compression::zstd::compress( - ArrayRef( - (const unsigned char *)(Bin->getBufferStart()), - Bin->getBufferSize()), - CompressedBuffer, OffloadCompressLevel); -#if LLVM_ENABLE_EXCEPTIONS - } catch (const std::exception &ex) { - return createStringError(inconvertibleErrorCode(), - std::string("Failed to compress the device image: \n") + - std::string(ex.what())); - } -#endif - if (Verbose) - errs() << "[Compression] Original image size: " - << Bin->getBufferSize() << "\n" - << "[Compression] Compressed image size: " - << CompressedBuffer.size() << "\n" - << "[Compression] Compression level used: " - << OffloadCompressLevel << "\n"; - - // Add the compressed image to the module. + if (DidCompress) { Fbin = addDeviceImageToModule( ArrayRef((const char *)CompressedBuffer.data(), CompressedBuffer.size()), - Twine(OffloadKindTag) + Twine(ImgId) + Twine(".data"), Kind, + Twine(OldOffloadKindTag) + Twine(ImgId) + Twine(".data"), Kind, Img.Tgt); - - // Change image format to compressed_none. Ffmt = ConstantInt::get(Type::getInt8Ty(C), BinaryImageFormat::compressed_none); + } else { + Fbin = addDeviceImageToModule( + ArrayRef(Bin->getBufferStart(), Bin->getBufferSize()), + Twine(OldOffloadKindTag) + Twine(ImgId) + Twine(".data"), Kind, + Img.Tgt); } } - if (Kind == OffloadKind::SYCL) { + if (Kind == OldOffloadKind::SYCL) { // For SYCL image offload entries are defined here, by wrapper, so // those are created per image Expected> EntriesOrErr = @@ -1154,7 +1127,7 @@ class BinaryWrapper { auto *ImgInfoVar = new GlobalVariable( M, ImgInfoArr->getType(), /*isConstant*/ true, GlobalVariable::InternalLinkage, ImgInfoArr, - Twine(OffloadKindTag) + Twine(ImgId) + Twine(".info")); + Twine(OldOffloadKindTag) + Twine(ImgId) + Twine(".info")); ImgInfoVar->setAlignment( MaybeAlign(M.getDataLayout().getTypeStoreSize(IntPtrTy) * 2u)); ImgInfoVar->setUnnamedAddr(GlobalValue::UnnamedAddr::Local); @@ -1170,7 +1143,7 @@ class BinaryWrapper { // Then create images array. auto *ImagesData = - Kind == OffloadKind::SYCL + Kind == OldOffloadKind::SYCL ? ConstantArray::get( ArrayType::get(getSyclDeviceImageTy(), ImagesInits.size()), ImagesInits) @@ -1181,7 +1154,7 @@ class BinaryWrapper { auto *Images = new GlobalVariable(M, ImagesData->getType(), /*isConstant*/ true, GlobalValue::InternalLinkage, ImagesData, - Twine(OffloadKindTag) + "device_images"); + Twine(OldOffloadKindTag) + "device_images"); if (Verbose) errs() << " global added: " << Images->getName() << "\n"; Images->setUnnamedAddr(GlobalValue::UnnamedAddr::Global); @@ -1191,7 +1164,7 @@ class BinaryWrapper { // And finally create the binary descriptor object. auto *DescInit = - Kind == OffloadKind::SYCL + Kind == OldOffloadKind::SYCL ? ConstantStruct::get( getSyclBinDescTy(), ConstantInt::get(Type::getInt16Ty(C), BinDescStructVersion), @@ -1207,13 +1180,13 @@ class BinaryWrapper { : GlobalValue::InternalLinkage; auto *Res = new GlobalVariable( M, DescInit->getType(), /*isConstant*/ true, Lnk, DescInit, - Twine(OffloadKindTag) + Twine(DescriptorName)); + Twine(OldOffloadKindTag) + Twine(DescriptorName)); if (Verbose) errs() << " global added: " << Res->getName() << "\n"; return Res; } - void createRegisterFunction(OffloadKind Kind, GlobalVariable *BinDesc) { + void createRegisterFunction(OldOffloadKind Kind, GlobalVariable *BinDesc) { auto *FuncTy = FunctionType::get(Type::getVoidTy(C), /*isVarArg*/ false); auto *Func = Function::Create(FuncTy, GlobalValue::InternalLinkage, @@ -1223,10 +1196,10 @@ class BinaryWrapper { // Get RegFuncName function declaration. auto *RegFuncTy = FunctionType::get(Type::getVoidTy(C), getPtrTy(), /*isVarArg=*/false); - FunctionCallee RegFuncC = - M.getOrInsertFunction(Kind == OffloadKind::SYCL ? "__sycl_register_lib" + FunctionCallee RegFuncC = M.getOrInsertFunction(Kind == OldOffloadKind::SYCL + ? "__sycl_register_lib" : "__tgt_register_lib", - RegFuncTy); + RegFuncTy); // Construct function body IRBuilder<> Builder(BasicBlock::Create(C, "entry", Func)); @@ -1242,7 +1215,7 @@ class BinaryWrapper { appendToGlobalCtors(M, Func, /*Priority*/ 1); } - void createUnregisterFunction(OffloadKind Kind, GlobalVariable *BinDesc) { + void createUnregisterFunction(OldOffloadKind Kind, GlobalVariable *BinDesc) { auto *FuncTy = FunctionType::get(Type::getVoidTy(C), /*isVarArg*/ false); auto *Func = Function::Create(FuncTy, GlobalValue::InternalLinkage, @@ -1253,8 +1226,8 @@ class BinaryWrapper { auto *UnRegFuncTy = FunctionType::get(Type::getVoidTy(C), getPtrTy(), /*isVarArg=*/false); FunctionCallee UnRegFuncC = M.getOrInsertFunction( - Kind == OffloadKind::SYCL ? "__sycl_unregister_lib" - : "__tgt_unregister_lib", + Kind == OldOffloadKind::SYCL ? "__sycl_unregister_lib" + : "__tgt_unregister_lib", UnRegFuncTy); // Construct function body @@ -1378,7 +1351,7 @@ class BinaryWrapper { Expected wrap() { for (auto &X : Packs) { - OffloadKind Kind = X.first; + OldOffloadKind Kind = X.first; SameKindPack *Pack = X.second.get(); Expected DescOrErr = createBinDesc(Kind, *Pack); if (!DescOrErr) @@ -1386,7 +1359,7 @@ class BinaryWrapper { if (EmitRegFuncs) { GlobalVariable *Desc = *DescOrErr; - if (Kind == OffloadKind::SYCL && + if (Kind == OldOffloadKind::SYCL && Triple(M.getTargetTriple()).isOSWindows()) { createSyclRegisterWithAtexitUnregister(Desc); } else { @@ -1841,7 +1814,7 @@ int main(int argc, const char **argv) { // add them to the wrapper BinaryWrapper Wr(Target, argv[0], SymPropBCFiles); - OffloadKind Knd = OffloadKind::Unknown; + OldOffloadKind Knd = OldOffloadKind::Unknown; llvm::StringRef Tgt = ""; BinaryImageFormat Fmt = BinaryImageFormat::none; llvm::StringRef CompileOpts = ""; @@ -1891,7 +1864,7 @@ int main(int argc, const char **argv) { Row.getCell(COL_PROPS, "")); } } else { - if (Knd == OffloadKind::Unknown) { + if (Knd == OldOffloadKind::Unknown) { reportError(createStringError(errc::invalid_argument, "offload model not set")); return 1; diff --git a/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h b/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h index bf39a7c0775d1..b99a6e31bd9da 100644 --- a/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h +++ b/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h @@ -23,11 +23,16 @@ namespace llvm { namespace offloading { // SYCL binary image formats supported. +// +// BIF_Compressed marks a zstd-compressed payload. The runtime re-derives the +// real format (SPIR-V / LLVM bitcode / native) post-decompression by reading +// the magic-number, so any pre-compression tag is not preserved. enum class SYCLBinaryImageFormat { - BIF_None, // Undetermined Image kind - BIF_Native, // Native Image kind - BIF_SPIRV, // SPIR-V - BIF_LLVMBC // LLVM bitcode + BIF_None, // Undetermined Image kind + BIF_Native, // Native Image kind + BIF_SPIRV, // SPIR-V + BIF_LLVMBC, // LLVM bitcode + BIF_Compressed // zstd-compressed image; format-of-original unknown }; struct SYCLImage { diff --git a/llvm/include/llvm/Frontend/Offloading/Utility.h b/llvm/include/llvm/Frontend/Offloading/Utility.h index 4c0bc87786dfb..1d151d35698bc 100644 --- a/llvm/include/llvm/Frontend/Offloading/Utility.h +++ b/llvm/include/llvm/Frontend/Offloading/Utility.h @@ -234,6 +234,23 @@ LLVM_ABI Error containerizeOpenMPSPIRVImage( std::unique_ptr &Binary, llvm::Triple Triple, StringRef CompileOpts = "", StringRef LinkOpts = ""); } // namespace intel + +// Default compression level for --compress. +constexpr int DefaultSYCLCompressionLevel = 10; + +/// zstd-compress a SYCL device image. +/// Errors if zstd is unavailable at build time. +/// \param Input The uncompressed image bytes. +/// \param Output Receives the compressed bytes on a hit; left untouched on a +/// skip. +/// \param Level zstd compression level. +/// \param Threshold Skip compression when Input.size() < Threshold. +/// \param Verbose Emit the "[Compression]" size/level lines to errs(). +/// \returns true if compression ran. +LLVM_ABI Expected +compressSYCLDeviceImage(ArrayRef Input, + SmallVectorImpl &Output, int Level = 10, + size_t Threshold = 512, bool Verbose = false); } // namespace offloading } // namespace llvm diff --git a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp index 0e21fe3143737..6118859ae63cc 100644 --- a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp +++ b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp @@ -61,6 +61,8 @@ int8_t binaryImageFormatToInt8(SYCLBinaryImageFormat Format) { return 2; case SYCLBinaryImageFormat::BIF_LLVMBC: return 3; + case SYCLBinaryImageFormat::BIF_Compressed: + return 4; default: llvm_unreachable("unexpected SYCLBinaryImageFormat"); } diff --git a/llvm/lib/Frontend/Offloading/Utility.cpp b/llvm/lib/Frontend/Offloading/Utility.cpp index c07d276244ee1..3688033a44b56 100644 --- a/llvm/lib/Frontend/Offloading/Utility.cpp +++ b/llvm/lib/Frontend/Offloading/Utility.cpp @@ -18,7 +18,9 @@ #include "llvm/Object/OffloadBinary.h" #include "llvm/ObjectYAML/ELFYAML.h" #include "llvm/ObjectYAML/yaml2obj.h" +#include "llvm/Support/Compression.h" #include "llvm/Support/MemoryBufferRef.h" +#include "llvm/Support/raw_ostream.h" #include "llvm/Transforms/Utils/ModuleUtils.h" using namespace llvm; @@ -506,3 +508,34 @@ void sycl::writeSymbolTable(ArrayRef Names, SmallString<0> &Out) { CurrentOffset += Names[I].size() + 1; } } + +Expected +offloading::compressSYCLDeviceImage(ArrayRef Input, + SmallVectorImpl &Output, int Level, + size_t Threshold, bool Verbose) { + if (!compression::zstd::isAvailable()) + return createStringError( + "Device image compression was requested, but LLVM was built without " + "zstd support. Rebuild with LLVM_ENABLE_ZSTD enabled to use this " + "feature."); + + if (Input.size() < Threshold) + return false; + +#if LLVM_ENABLE_EXCEPTIONS + try { +#endif + compression::zstd::compress(Input, Output, Level); +#if LLVM_ENABLE_EXCEPTIONS + } catch (const std::exception &ex) { + return createStringError(std::string("Failed to compress the device " + "image: \n") + + std::string(ex.what())); + } +#endif + if (Verbose) + errs() << "[Compression] Original image size: " << Input.size() << "\n" + << "[Compression] Compressed image size: " << Output.size() << "\n" + << "[Compression] Compression level used: " << Level << "\n"; + return true; +}