[clang-offload-wrapper] Add standard notes for ELF offload images

[vzakhari]

Summary

This patch is a piece of D99360.

The patch adds ELF notes into SHT_NOTE sections of ELF offload images passed to clang-offload-wrapper.

The new notes use a null-terminated "LLVMOMPOFFLOAD" note name. There are currently three types of notes:

VERSION: a string (not null-terminated) representing the ELF offload image structure. The current version '1.0' does not put any restrictions on the structure of the image. If we ever need to come up with a common structure for ELF offload images (e.g. to be able to analyze the images in libomptarget in some standard way), then we will introduce new versions.
PRODUCER: a vendor specific name of the producing toolchain. Upstream LLVM uses "LLVM" (not null-terminated).
PRODUCER_VERSION: a vendor specific version of the producing toolchain. Upstream LLVM uses LLVM_VERSION_STRING with optional <space> LLVM_REVISION.

All three notes are not mandatory currently.

[vzakhari]

I am not sure what is wrong with the uploaded ELF files. The test works locally. Let's try it with a thinner YAML template.

[vzakhari]

Updated revision fixes BUILD_SHARED_LIBS build.

[grokos]

Change looks good, so it's accepted on my end. I'll let the other reviewers have a look and post their comments. Please do not commit until we have reached an agreement for all 4 patches together (D99551, D99552, D99553, D99612).

[jhuber6]

This patch broke offloading on my machine. If I compile a basic offloading application I get an invalid device image.

#include <omp.h>

int main() {
  int isHost;
#pragma omp target map(from:isHost)
  { isHost = omp_is_initial_device(); }
  return isHost;
}

Compiling and running with clang -fopenmp -fopenmp-targets=nvptx64 gives me the following results.

CUDA error: Error returned from cuModuleLoadDataEx
CUDA error: device kernel image is invalid 
Libomptarget error: Unable to generate entries table for device id 0.
Libomptarget error: Failed to init globals on device 0
Libomptarget error: Run with LIBOMPTARGET_INFO=4 to dump host-target pointer mappings.
Libomptarget fatal error 1: failure of target construct while offloading is mandatory

Reverting 93d08acaacec951dbb302f77eeae51974985b6b2 fixes it.

[vzakhari]

@jhuber6, thank you for reporting this. I do not have a properly setup CUDA system currently. Can you please invoke clang-offload-wrapper with -save-temps and send the temporary files to me?

[jhuber6]

In D99551#2950281, @vzakhari wrote:

@jhuber6, thank you for reporting this. I do not have a properly setup CUDA system currently. Can you please invoke clang-offload-wrapper with -save-temps and send the temporary files to me?

I sent it to you, any luck or should we revert this upstream for the time being.

[vzakhari]

In D99551#2950464, @jhuber6 wrote:

In D99551#2950281, @vzakhari wrote:

@jhuber6, thank you for reporting this. I do not have a properly setup CUDA system currently. Can you please invoke clang-offload-wrapper with -save-temps and send the temporary files to me?

I sent it to you, any luck or should we revert this upstream for the time being.

Unfortunately, I did not receive it. Can you please check if the mail was blocked on your side? Can you please try to archive it with a password and send it again?
I would like to prepare a patch to put this functionality under a switch that is off by default. I think we want to have this working eventually (and it actually works with x86_64 offload), so it will be easier to just flip a switch, when I figure out what CUDA API does not like about the modified ELF image. Does it sound appropriate to you?

[jhuber6]

In D99551#2950515, @vzakhari wrote:

In D99551#2950464, @jhuber6 wrote:

In D99551#2950281, @vzakhari wrote:

@jhuber6, thank you for reporting this. I do not have a properly setup CUDA system currently. Can you please invoke clang-offload-wrapper with -save-temps and send the temporary files to me?

I sent it to you, any luck or should we revert this upstream for the time being.

Unfortunately, I did not receive it. Can you please check if the mail was blocked on your side? Can you please try to archive it with a password and send it again?
I would like to prepare a patch to put this functionality under a switch that is off by default. I think we want to have this working eventually (and it actually works with x86_64 offload), so it will be easier to just flip a switch, when I figure out what CUDA API does not like about the modified ELF image. Does it sound appropriate to you?

I put it in a .tar.gz file, I think your domain is blocking it. Anything else I can use?

[vzakhari]

In D99551#2950522, @jhuber6 wrote:

In D99551#2950515, @vzakhari wrote:

In D99551#2950464, @jhuber6 wrote:

In D99551#2950281, @vzakhari wrote:

@jhuber6, thank you for reporting this. I do not have a properly setup CUDA system currently. Can you please invoke clang-offload-wrapper with -save-temps and send the temporary files to me?

I sent it to you, any luck or should we revert this upstream for the time being.

Unfortunately, I did not receive it. Can you please check if the mail was blocked on your side? Can you please try to archive it with a password and send it again?
I would like to prepare a patch to put this functionality under a switch that is off by default. I think we want to have this working eventually (and it actually works with x86_64 offload), so it will be easier to just flip a switch, when I figure out what CUDA API does not like about the modified ELF image. Does it sound appropriate to you?

I put it in a .tar.gz file, I think your domain is blocking it. Anything else I can use?

Can you please try sharing it via Google drive?

[JonChesterfield]

Nvptx broken here too, amdgpu is fine. I'm guessing one of the cuda tools does some overly aggressive input validation that we're running afoul of.

There was a discussion about this on the call today - plan was to put it behind a disabled boolean argument while fixing to avoid downstream churn. Sadly the original patch was not authored with that in mind. I suggest if we can't get that patch together asap we revert this and fix it offline (even if the fix is adding said flag)

[vzakhari]

In D99551#2952336, @JonChesterfield wrote:

Nvptx broken here too, amdgpu is fine. I'm guessing one of the cuda tools does some overly aggressive input validation that we're running afoul of.

There was a discussion about this on the call today - plan was to put it behind a disabled boolean argument while fixing to avoid downstream churn. Sadly the original patch was not authored with that in mind. I suggest if we can't get that patch together asap we revert this and fix it offline (even if the fix is adding said flag)

I am about to merge D108246 that is adding the switch.

[JonChesterfield]

It's a hack, but D108303 will unblock nvptx offloading. Alternative to reverting. Suggest we go with that then revisit in a couple of weeks.

clang/test/Driver/Inputs/empty-elf-template.yaml

+--- !ELF
+FileHeader:
+  Class: ELFCLASS[[BITS]]
+  Data:  ELFDATA2[[ENCODING]]
+  Type:  ET_REL

clang/test/Driver/clang-offload-wrapper.c

 //
 // Generate a file to wrap.
 //
 // RUN: echo 'Content of device file' > %t.tgt
 
 //
 // Check bitcode produced by the wrapper tool.
 //
-// RUN: clang-offload-wrapper -target=x86_64-pc-linux-gnu -o %t.wrapper.bc %t.tgt
+// RUN: clang-offload-wrapper -target=x86_64-pc-linux-gnu -o %t.wrapper.bc %t.tgt 2>&1 | FileCheck %s --check-prefix ELF-WARNING
 // RUN: llvm-dis %t.wrapper.bc -o - | FileCheck %s --check-prefix CHECK-IR
 
+// ELF-WARNING: is not an ELF image, so notes cannot be added to it.
 // CHECK-IR: target triple = "x86_64-pc-linux-gnu"
 
 // CHECK-IR-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 }
 // CHECK-IR-DAG: [[IMAGETY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* }
 // CHECK-IR-DAG: [[DESCTY:%.+]] = type { i32, [[IMAGETY]]*, [[ENTTY]]*, [[ENTTY]]* }
 
 // CHECK-IR: [[ENTBEGIN:@.+]] = external hidden constant [[ENTTY]]
 // CHECK-IR: [[ENTEND:@.+]] = external hidden constant [[ENTTY]]

 
 // CHECK-IR: declare void @__tgt_register_lib([[DESCTY]]*)
 
 // CHECK-IR: define internal void [[UNREGFN]]()
 // CHECK-IR:   call void @__tgt_unregister_lib([[DESCTY]]* [[DESC]])
 // CHECK-IR:   ret void
 
 // CHECK-IR: declare void @__tgt_unregister_lib([[DESCTY]]*)
+
+// Check that clang-offload-wrapper adds LLVMOMPOFFLOAD notes
+// into the ELF offload images:
+// RUN: yaml2obj %S/Inputs/empty-elf-template.yaml -o %t.64le -DBITS=64 -DENCODING=LSB
+// RUN: clang-offload-wrapper -target=x86_64-pc-linux-gnu -o %t.wrapper.elf64le.bc %t.64le
+// RUN: llvm-dis %t.wrapper.elf64le.bc -o - | FileCheck %s --check-prefix OMPNOTES
+// RUN: yaml2obj %S/Inputs/empty-elf-template.yaml -o %t.64be -DBITS=64 -DENCODING=MSB
+// RUN: clang-offload-wrapper -target=x86_64-pc-linux-gnu -o %t.wrapper.elf64be.bc %t.64be
+// RUN: llvm-dis %t.wrapper.elf64be.bc -o - | FileCheck %s --check-prefix OMPNOTES
+// RUN: yaml2obj %S/Inputs/empty-elf-template.yaml -o %t.32le -DBITS=32 -DENCODING=LSB
+// RUN: clang-offload-wrapper -target=x86_64-pc-linux-gnu -o %t.wrapper.elf32le.bc %t.32le
+// RUN: llvm-dis %t.wrapper.elf32le.bc -o - | FileCheck %s --check-prefix OMPNOTES
+// RUN: yaml2obj %S/Inputs/empty-elf-template.yaml -o %t.32be -DBITS=32 -DENCODING=MSB
+// RUN: clang-offload-wrapper -target=x86_64-pc-linux-gnu -o %t.wrapper.elf32be.bc %t.32be
+// RUN: llvm-dis %t.wrapper.elf32be.bc -o - | FileCheck %s --check-prefix OMPNOTES
+
+// There is no clean way for extracting the offload image
+// from the object file currently, so try to find
+// the inserted ELF notes in the device image variable's
+// initializer:
+// OMPNOTES: @{{.+}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"{{.*}}LLVMOMPOFFLOAD{{.*}}LLVMOMPOFFLOAD{{.*}}LLVMOMPOFFLOAD{{.*}}"

clang/tools/clang-offload-wrapper/CMakeLists.txt

-set(LLVM_LINK_COMPONENTS BitWriter Core Support TransformUtils)
+set(LLVM_LINK_COMPONENTS BitWriter Core Object Support TransformUtils)
 
 add_clang_tool(clang-offload-wrapper
   ClangOffloadWrapper.cpp
 
   DEPENDS
   intrinsics_gen
   )
 

clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp

 /// packaged as data. Wrapper bitcode also includes initialization code which
 /// registers target binaries in offloading runtime at program startup.
 ///
 //===----------------------------------------------------------------------===//
 
 #include "clang/Basic/Version.h"
 #include "llvm/ADT/ArrayRef.h"
 #include "llvm/ADT/Triple.h"
+#include "llvm/BinaryFormat/ELF.h"
 #include "llvm/Bitcode/BitcodeWriter.h"
 #include "llvm/IR/Constants.h"
 #include "llvm/IR/GlobalVariable.h"
 #include "llvm/IR/IRBuilder.h"
 #include "llvm/IR/LLVMContext.h"
 #include "llvm/IR/Module.h"
+#include "llvm/Object/ELFObjectFile.h"
+#include "llvm/Object/ObjectFile.h"
 #include "llvm/Support/CommandLine.h"
+#include "llvm/Support/EndianStream.h"
 #include "llvm/Support/Errc.h"
 #include "llvm/Support/Error.h"
 #include "llvm/Support/ErrorOr.h"
 #include "llvm/Support/FileSystem.h"
 #include "llvm/Support/MemoryBuffer.h"
+#include "llvm/Support/Path.h"
+#include "llvm/Support/Program.h"
 #include "llvm/Support/Signals.h"
 #include "llvm/Support/ToolOutputFile.h"
+#include "llvm/Support/VCSRevision.h"
 #include "llvm/Support/WithColor.h"
 #include "llvm/Support/raw_ostream.h"
 #include "llvm/Transforms/Utils/ModuleUtils.h"
 #include <cassert>
 #include <cstdint>
 
+#define OPENMP_OFFLOAD_IMAGE_VERSION "1.0"
+
 using namespace llvm;
+using namespace llvm::object;
 
 static cl::opt<bool> Help("h", cl::desc("Alias for -help"), cl::Hidden);
 
 // Mark all our options with this category, everything else (except for -version
 // and -help) will be hidden.
 static cl::OptionCategory
     ClangOffloadWrapperCategory("clang-offload-wrapper options");
 
 static cl::opt<std::string> Output("o", cl::Required,
                                    cl::desc("Output filename"),
                                    cl::value_desc("filename"),
                                    cl::cat(ClangOffloadWrapperCategory));
 
 static cl::list<std::string> Inputs(cl::Positional, cl::OneOrMore,
                                     cl::desc("<input files>"),
                                     cl::cat(ClangOffloadWrapperCategory));
 
 static cl::opt<std::string>
     Target("target", cl::Required,
            cl::desc("Target triple for the output module"),
            cl::value_desc("triple"), cl::cat(ClangOffloadWrapperCategory));
 
+static cl::opt<bool> SaveTemps(
+    "save-temps",
+    cl::desc("Save temporary files that may be produced by the tool. "
+             "This option forces print-out of the temporary files' names."),
+    cl::Hidden);
+
 namespace {
 
 class BinaryWrapper {
   LLVMContext C;
   Module M;
 
   StructType *EntryTy = nullptr;
   StructType *ImageTy = nullptr;
   StructType *DescTy = nullptr;
 
+  std::string ToolName;
+  std::string ObjcopyPath;
+  // Temporary file names that may be created during adding notes
+  // to ELF offload images. Use -save-temps to keep them and also
+  // see their names. A temporary file's name includes the name
+  // of the original input ELF image, so you can easily match
+  // them, if you have multiple inputs.
+  std::vector<std::string> TempFiles;
+
 private:
   IntegerType *getSizeTTy() {
     switch (M.getDataLayout().getPointerTypeSize(Type::getInt8PtrTy(C))) {
     case 4u:
       return Type::getInt32Ty(C);
     case 8u:
       return Type::getInt64Ty(C);
     }
void createUnregisterFunction(GlobalVariable *BinDesc) {
     Builder.CreateRetVoid();
 
     // Add this function to global destructors.
     // Match priority of __tgt_register_lib
     appendToGlobalDtors(M, Func, /*Priority*/ 1);
   }
 
 public:
-  BinaryWrapper(StringRef Target) : M("offload.wrapper.object", C) {
+  BinaryWrapper(StringRef Target, StringRef ToolName)
+      : M("offload.wrapper.object", C), ToolName(ToolName) {
     M.setTargetTriple(Target);
+    // Look for llvm-objcopy in the same directory, from which
+    // clang-offload-wrapper is invoked. This helps OpenMP offload
+    // LIT tests.
+
+    // This just needs to be some symbol in the binary; C++ doesn't
+    // allow taking the address of ::main however.
+    void *P = (void *)(intptr_t)&Help;
+    std::string COWPath = sys::fs::getMainExecutable(ToolName.str().c_str(), P);
+    if (!COWPath.empty()) {
+      auto COWDir = sys::path::parent_path(COWPath);
+      ErrorOr<std::string> ObjcopyPathOrErr =
+          sys::findProgramByName("llvm-objcopy", {COWDir});
+      if (ObjcopyPathOrErr) {
+        ObjcopyPath = *ObjcopyPathOrErr;
+        return;
+      }
+
+      // Otherwise, look through PATH environment.
+    }
+
+    ErrorOr<std::string> ObjcopyPathOrErr =
+        sys::findProgramByName("llvm-objcopy");
+    if (!ObjcopyPathOrErr) {
+      WithColor::warning(errs(), ToolName)
+          << "cannot find llvm-objcopy[.exe] in PATH; ELF notes cannot be "
+             "added.\n";
+      return;
+    }
+
+    ObjcopyPath = *ObjcopyPathOrErr;
+  }
+
+  ~BinaryWrapper() {
+    if (TempFiles.empty())
+      return;
+
+    StringRef ToolNameRef(ToolName);
+    auto warningOS = [ToolNameRef]() -> raw_ostream & {
+      return WithColor::warning(errs(), ToolNameRef);
+    };
+
+    for (auto &F : TempFiles) {
+      if (SaveTemps) {
+        warningOS() << "keeping temporary file " << F << "\n";
+        continue;
+      }
+
+      auto EC = sys::fs::remove(F, false);
+      if (EC)
+        warningOS() << "cannot remove temporary file " << F << ": "
+                    << EC.message().c_str() << "\n";
+    }
   }
 
   const Module &wrapBinaries(ArrayRef<ArrayRef<char>> Binaries) {
     GlobalVariable *Desc = createBinDesc(Binaries);
     assert(Desc && "no binary descriptor");
     createRegisterFunction(Desc);
     createUnregisterFunction(Desc);
     return M;
   }
+
+  std::unique_ptr<MemoryBuffer> addELFNotes(std::unique_ptr<MemoryBuffer> Buf,
+                                            StringRef OriginalFileName) {
+    // Cannot add notes, if llvm-objcopy is not available.
+    //
+    // I did not find a clean way to add a new notes section into an existing
+    // ELF file. llvm-objcopy seems to recreate a new ELF from scratch,
+    // and we just try to use llvm-objcopy here.
+    if (ObjcopyPath.empty())
+      return Buf;
+
+    StringRef ToolNameRef(ToolName);
+
+    // Helpers to emit warnings.
+    auto warningOS = [ToolNameRef]() -> raw_ostream & {
+      return WithColor::warning(errs(), ToolNameRef);
+    };
+    auto handleErrorAsWarning = [&warningOS](Error E) {
+      logAllUnhandledErrors(std::move(E), warningOS());
+    };
+
+    Expected<std::unique_ptr<ObjectFile>> BinOrErr =
+        ObjectFile::createELFObjectFile(Buf->getMemBufferRef(),
+                                        /*InitContent=*/false);
+    if (Error E = BinOrErr.takeError()) {
+      consumeError(std::move(E));
+      // This warning is questionable, but let it be here,
+      // assuming that most OpenMP offload models use ELF offload images.
+      warningOS() << OriginalFileName
+                  << " is not an ELF image, so notes cannot be added to it.\n";
+      return Buf;
+    }
+
+    // If we fail to add the note section, we just pass through the original
+    // ELF image for wrapping. At some point we should enforce the note section
+    // and start emitting errors vs warnings.
+    support::endianness Endianness;
+    if (isa<ELF64LEObjectFile>(BinOrErr->get()) ||
+        isa<ELF32LEObjectFile>(BinOrErr->get())) {
+      Endianness = support::little;
+    } else if (isa<ELF64BEObjectFile>(BinOrErr->get()) ||
+               isa<ELF32BEObjectFile>(BinOrErr->get())) {
+      Endianness = support::big;
+    } else {
+      warningOS() << OriginalFileName
+                  << " is an ELF image of unrecognized format.\n";
+      return Buf;
+    }
+
+    // Create temporary file for the data of a new SHT_NOTE section.
+    // We fill it in with data and then pass to llvm-objcopy invocation
+    // for reading.
+    Twine NotesFileModel = OriginalFileName + Twine(".elfnotes.%%%%%%%.tmp");
+    Expected<sys::fs::TempFile> NotesTemp =
+        sys::fs::TempFile::create(NotesFileModel);
+    if (Error E = NotesTemp.takeError()) {
+      handleErrorAsWarning(createFileError(NotesFileModel, std::move(E)));
+      return Buf;
+    }
+    TempFiles.push_back(NotesTemp->TmpName);
+
+    // Create temporary file for the updated ELF image.
+    // This is an empty file that we pass to llvm-objcopy invocation
+    // for writing.
+    Twine ELFFileModel = OriginalFileName + Twine(".elfwithnotes.%%%%%%%.tmp");
+    Expected<sys::fs::TempFile> ELFTemp =
+        sys::fs::TempFile::create(ELFFileModel);
+    if (Error E = ELFTemp.takeError()) {
+      handleErrorAsWarning(createFileError(ELFFileModel, std::move(E)));
+      return Buf;
+    }
+    TempFiles.push_back(ELFTemp->TmpName);
+
+    // Keep the new ELF image file to reserve the name for the future
+    // llvm-objcopy invocation.
+    std::string ELFTmpFileName = ELFTemp->TmpName;
+    if (Error E = ELFTemp->keep(ELFTmpFileName)) {
+      handleErrorAsWarning(createFileError(ELFTmpFileName, std::move(E)));
+      return Buf;
+    }
+
+    // Write notes to the *elfnotes*.tmp file.
+    raw_fd_ostream NotesOS(NotesTemp->FD, false);
+
+    struct NoteTy {
+      // Note name is a null-terminated "LLVMOMPOFFLOAD".
+      std::string Name;
+      // Note type defined in llvm/include/llvm/BinaryFormat/ELF.h.
+      uint32_t Type = 0;
+      // Each note has type-specific associated data.
+      std::string Desc;
+
+      NoteTy(std::string &&Name, uint32_t Type, std::string &&Desc)
+          : Name(std::move(Name)), Type(Type), Desc(std::move(Desc)) {}
+    };
+
+    // So far we emit just three notes.
+    SmallVector<NoteTy, 3> Notes;
+    // Version of the offload image identifying the structure of the ELF image.
+    // Version 1.0 does not have any specific requirements.
+    // We may come up with some structure that has to be honored by all
+    // offload implementations in future (e.g. to let libomptarget
+    // get some information from the offload image).
+    Notes.emplace_back("LLVMOMPOFFLOAD", ELF::NT_LLVM_OPENMP_OFFLOAD_VERSION,
+                       OPENMP_OFFLOAD_IMAGE_VERSION);
+    // This is a producer identification string. We are LLVM!
+    Notes.emplace_back("LLVMOMPOFFLOAD", ELF::NT_LLVM_OPENMP_OFFLOAD_PRODUCER,
+                       "LLVM");
+    // This is a producer version. Use the same format that is used
+    // by clang to report the LLVM version.
+    Notes.emplace_back("LLVMOMPOFFLOAD",
+                       ELF::NT_LLVM_OPENMP_OFFLOAD_PRODUCER_VERSION,
+                       LLVM_VERSION_STRING
+#ifdef LLVM_REVISION
+                       " " LLVM_REVISION
+#endif
+    );
+
+    // Return the amount of padding required for a blob of N bytes
+    // to be aligned to Alignment bytes.
+    auto getPadAmount = [](uint32_t N, uint32_t Alignment) -> uint32_t {
+      uint32_t Mod = (N % Alignment);
+      if (Mod == 0)
+        return 0;
+      return Alignment - Mod;
+    };
+    auto emitPadding = [&getPadAmount](raw_ostream &OS, uint32_t Size) {
+      for (uint32_t I = 0; I < getPadAmount(Size, 4); ++I)
+        OS << '\0';
+    };
+
+    // Put notes into the file.
+    for (auto &N : Notes) {
+      assert(!N.Name.empty() && "We should not create notes with empty names.");
+      // Name must be null-terminated.
+      if (N.Name.back() != '\0')
+        N.Name += '\0';
+      uint32_t NameSz = N.Name.size();
+      uint32_t DescSz = N.Desc.size();
+      // A note starts with three 4-byte values:
+      //   NameSz
+      //   DescSz
+      //   Type
+      // These three fields are endian-sensitive.
+      support::endian::write<uint32_t>(NotesOS, NameSz, Endianness);
+      support::endian::write<uint32_t>(NotesOS, DescSz, Endianness);
+      support::endian::write<uint32_t>(NotesOS, N.Type, Endianness);
+      // Next, we have a null-terminated Name padded to a 4-byte boundary.
+      NotesOS << N.Name;
+      emitPadding(NotesOS, NameSz);
+      if (DescSz == 0)
+        continue;
+      // Finally, we have a descriptor, which is an arbitrary flow of bytes.
+      NotesOS << N.Desc;
+      emitPadding(NotesOS, DescSz);
+    }
+    NotesOS.flush();
+
+    // Keep the notes file.
+    std::string NotesTmpFileName = NotesTemp->TmpName;
+    if (Error E = NotesTemp->keep(NotesTmpFileName)) {
+      handleErrorAsWarning(createFileError(NotesTmpFileName, std::move(E)));
+      return Buf;
+    }
+
+    // Run llvm-objcopy like this:
+    //   llvm-objcopy --add-section=.note.openmp=<notes-tmp-file-name> \
+    //       <orig-file-name> <elf-tmp-file-name>
+    //
+    // This will add a SHT_NOTE section on top of the original ELF.
+    std::vector<StringRef> Args;
+    Args.push_back(ObjcopyPath);
+    std::string Option("--add-section=.note.openmp=" + NotesTmpFileName);
+    Args.push_back(Option);
+    Args.push_back(OriginalFileName);
+    Args.push_back(ELFTmpFileName);
+    bool ExecutionFailed = false;
+    std::string ErrMsg;
+    (void)sys::ExecuteAndWait(ObjcopyPath, Args,
+                              /*Env=*/llvm::None, /*Redirects=*/{},
+                              /*SecondsToWait=*/0,
+                              /*MemoryLimit=*/0, &ErrMsg, &ExecutionFailed);
+
+    if (ExecutionFailed) {
+      warningOS() << ErrMsg << "\n";
+      return Buf;
+    }
+
+    // Substitute the original ELF with new one.
+    ErrorOr<std::unique_ptr<MemoryBuffer>> BufOrErr =
+        MemoryBuffer::getFile(ELFTmpFileName);
+    if (!BufOrErr) {
+      handleErrorAsWarning(
+          createFileError(ELFTmpFileName, BufOrErr.getError()));
+      return Buf;
+    }
+
+    return std::move(*BufOrErr);
+  }
 };
 
 } // anonymous namespace
 
 int main(int argc, const char **argv) {
   sys::PrintStackTraceOnErrorSignal(argv[0]);
 
   cl::HideUnrelatedOptions(ClangOffloadWrapperCategory);
int main(int argc, const char **argv) {
   };
 
   if (Triple(Target).getArch() == Triple::UnknownArch) {
     reportError(createStringError(
         errc::invalid_argument, "'" + Target + "': unsupported target triple"));
     return 1;
   }
 
+  BinaryWrapper Wrapper(Target, argv[0]);
+
   // Read device binaries.
   SmallVector<std::unique_ptr<MemoryBuffer>, 4u> Buffers;
   SmallVector<ArrayRef<char>, 4u> Images;
   Buffers.reserve(Inputs.size());
   Images.reserve(Inputs.size());
   for (const std::string &File : Inputs) {
     ErrorOr<std::unique_ptr<MemoryBuffer>> BufOrErr =
         MemoryBuffer::getFileOrSTDIN(File);
     if (!BufOrErr) {
       reportError(createFileError(File, BufOrErr.getError()));
       return 1;
     }
+    std::unique_ptr<MemoryBuffer> Buffer(std::move(*BufOrErr));
+    if (File != "-") {
+      // Adding ELF notes for STDIN is not supported yet.
+      Buffer = Wrapper.addELFNotes(std::move(Buffer), File);
+    }
     const std::unique_ptr<MemoryBuffer> &Buf =
-        Buffers.emplace_back(std::move(*BufOrErr));
+        Buffers.emplace_back(std::move(Buffer));
     Images.emplace_back(Buf->getBufferStart(), Buf->getBufferSize());
   }
 
   // Create the output file to write the resulting bitcode to.
   std::error_code EC;
   ToolOutputFile Out(Output, EC, sys::fs::OF_None);
   if (EC) {
     reportError(createFileError(Output, EC));
     return 1;
   }
 
   // Create a wrapper for device binaries and write its bitcode to the file.
-  WriteBitcodeToFile(BinaryWrapper(Target).wrapBinaries(
-                         makeArrayRef(Images.data(), Images.size())),
-                     Out.os());
+  WriteBitcodeToFile(
+      Wrapper.wrapBinaries(makeArrayRef(Images.data(), Images.size())),
   if (Out.os().has_error()) {
     reportError(createFileError(Output, Out.os().error()));
     return 1;
   }
 
   // Success.
   Out.keep();
   return 0;
 }