[OpenMP] Implement the creation and registration of the offloading descriptor.

[sfantao]

Summary

All the offloading information is bundled into a descriptor that is passed to the runtime library so that it can extract all the binaries and map variables properly. This descriptor includes the range of the target binaries (that will be defined by the linker) for each device selected by the user, as well as table with the information about each function and variable that is to be mapped (given that we do not support 'declare target' directives yet this is only implemented for function, but the logic can/will be reused).

This patch adds support for the creation of the descriptor as well as the registration/unregistration of the descriptor with the runtime library. The registration is implemented in a high priority global initializer so that the registration happens always before any initializer (that can potentially include target regions) is run.

The frontend flag (it, or something similar, will have to be promoted to driver option in the future) -omptargets= was created to exercise the new functionality. It takes the list of triples of the devices the user wants to offload to.

This patch depends on http://reviews.llvm.org/D11361

[sfantao]

I am closing this revision as I added the functionality for the device codegen on top of this patch in http://reviews.llvm.org/D12614 because the two are tightly coupled.

Thanks,
Samuel

include/clang/Basic/DiagnosticDriverKinds.td

def err_drv_mg_requires_m_or_mm : Error<
   "option '-MG' requires '-M' or '-MM'">;
 def err_drv_unknown_objc_runtime : Error<
   "unknown or ill-formed Objective-C runtime '%0'">;
 def err_drv_emit_llvm_link : Error<
    "-emit-llvm cannot be used when linking">;
 def err_drv_optimization_remark_pattern : Error<
   "%0 in '%1'">;
 def err_drv_no_neon_modifier : Error<"[no]neon is not accepted as modifier, please use [no]simd instead">;
+def err_drv_invalid_omp_target : Error<"OpenMP target is invalid: '%0'">;
 
 def warn_O4_is_O3 : Warning<"-O4 is equivalent to -O3">, InGroup<Deprecated>;
 def warn_drv_optimization_value : Warning<"optimization level '%0' is not supported; using '%1%2' instead">,
   InGroup<InvalidCommandLineArgument>;
 def warn_ignored_gcc_optimization : Warning<"optimization flag '%0' is not supported">,
   InGroup<IgnoredOptimizationArgument>;
 def warn_c_kext : Warning<
   "ignoring -fapple-kext which is valid for C++ and Objective-C++ only">;

include/clang/Basic/LangOptions.h

public:
   /// \brief The names of any features to enable in module 'requires' decls
   /// in addition to the hard-coded list in Module.cpp and the target features.
   ///
   /// This list is sorted.
   std::vector<std::string> ModuleFeatures;
 
   /// \brief Options for parsing comments.
   CommentOptions CommentOpts;
-  
+  /// \brief Triples of the OpenMP targets that the host code codegen should
+  /// take into account in order to generate accurate offloading descriptors.
+  std::vector<llvm::Triple> OMPTargetTriples;
+
   LangOptions();
 
   // Define accessors/mutators for language options of enumeration type.
 #define LANGOPT(Name, Bits, Default, Description) 
 #define ENUM_LANGOPT(Name, Type, Bits, Default, Description) \
   Type get##Name() const { return static_cast<Type>(Name); } \
   void set##Name(Type Value) { Name = static_cast<unsigned>(Value); }  
 #include "clang/Basic/LangOptions.def"

include/clang/Driver/Options.td

 def nostdinc : Flag<["-"], "nostdinc">;
 def nostdlibinc : Flag<["-"], "nostdlibinc">;
 def nostdincxx : Flag<["-"], "nostdinc++">, Flags<[CC1Option]>,
   HelpText<"Disable standard #include directories for the C++ standard library">;
 def nostdlib : Flag<["-"], "nostdlib">;
 def object : Flag<["-"], "object">;
 def o : JoinedOrSeparate<["-"], "o">, Flags<[DriverOption, RenderAsInput, CC1Option, CC1AsOption]>,
   HelpText<"Write output to <file>">, MetaVarName<"<file>">;
+def omptargets_EQ : CommaJoined<["-"], "omptargets=">, Flags<[CC1Option]>;
 def pagezero__size : JoinedOrSeparate<["-"], "pagezero_size">;
 def pass_exit_codes : Flag<["-", "--"], "pass-exit-codes">, Flags<[Unsupported]>;
 def pedantic_errors : Flag<["-", "--"], "pedantic-errors">, Group<pedantic_Group>, Flags<[CC1Option]>;
 def pedantic : Flag<["-", "--"], "pedantic">, Group<pedantic_Group>, Flags<[CC1Option]>;
 def pg : Flag<["-"], "pg">, HelpText<"Enable mcount instrumentation">, Flags<[CC1Option]>;
 def pipe : Flag<["-", "--"], "pipe">,
   HelpText<"Use pipes between commands, when possible">;
 def prebind__all__twolevel__modules : Flag<["-"], "prebind_all_twolevel_modules">;

lib/CodeGen/CGOpenMPRuntime.h

enum OpenMPRTLFunction {
 
     //
     // Offloading related calls
     //
     // Call to int32_t __tgt_target(int32_t device_id, void *host_ptr, int32_t
     // arg_num, void** args_base, void **args, size_t *arg_sizes, int32_t
     // *arg_types);
     OMPRTL__tgt_target,
+    // Call to void __tgt_register_lib(__tgt_bin_desc *desc);
+    OMPRTL__tgt_register_lib,
+    // Call to void __tgt_unregister_lib(__tgt_bin_desc *desc);
+    OMPRTL__tgt_unregister_lib,
   };
 
   /// \brief Values for bit flags used in the ident_t to describe the fields.
   /// All enumeric elements are named and described in accordance with the code
   /// from http://llvm.org/svn/llvm-project/openmp/trunk/runtime/src/kmp.h
   enum OpenMPLocationFlags {
     /// \brief Use trampoline for internal microtask.
     OMP_IDENT_IMD = 0x01,
private:
   ///    kmp_intptr_t               base_addr;
   ///    size_t                     len;
   ///    struct {
   ///             bool                   in:1;
   ///             bool                   out:1;
   ///    } flags;
   /// } kmp_depend_info_t;
   QualType KmpDependInfoTy;
+  /// \brief Type struct __tgt_offload_entry{
+  ///   void      *addr;       // Pointer to the offload entry info.
+  ///                          // (function or global)
+  ///   char      *name;       // Name of the function or global.
+  ///   size_t     size;       // Size of the entry info (0 if it a function).
+  /// };
+  llvm::StructType *TgtOffloadEntryTy;
+  /// struct __tgt_device_image{
+  /// void   *ImageStart;       // Pointer to the target code start.
+  /// void   *ImageEnd;         // Pointer to the target code end.
+  /// // We also add the host entries to the device image, as it may be useful
+  /// // for the target runtime to have access to that information.
+  /// __tgt_offload_entry  *EntriesBegin;   // Begin of the table with all
+  ///                                       // the entries.
+  /// __tgt_offload_entry  *EntriesEnd;     // End of the table with all the
+  ///                                       // entries (non inclusive).
+  /// };
+  llvm::StructType *TgtDeviceImageTy;
+  /// struct __tgt_bin_desc{
+  ///   int32_t              NumDevices;      // Number of devices supported.
+  ///   __tgt_device_image   *DeviceImages;   // Arrays of device images
+  ///                                         // (one per device).
+  ///   __tgt_offload_entry  *EntriesBegin;   // Begin of the table with all the
+  ///                                         // entries.
+  ///   __tgt_offload_entry  *EntriesEnd;     // End of the table with all the
+  ///                                         // entries (non inclusive).
+  /// };
+  llvm::StructType *TgtBinaryDescriptorTy;
+  /// \brief Array that registers the offloading constants that were emitted so
+  /// far. The order of the registration is the same as the order the entries
+  /// associated with these constants are emitted.
+  struct OffloadingEntryInfo {
+    llvm::Constant *Addr; // Address of the entity that has to be mapped for
+                          // offloading.
+    llvm::Constant *ID;   // Address that can be used as the ID of the entity.
+                          // This is different than Addr if the entity is a
+                          // target function.
+    OffloadingEntryInfo(llvm::Constant *_Addr, llvm::Constant *_ID)
+        : Addr(_Addr), ID(_ID) {}
+  };
+  llvm::SmallVector<OffloadingEntryInfo, 32> OffloadingEntriesInfo;
+
+  /// \brief Register the information required to create an entry for the
+  /// provided offloading entity.
+  void registerOffloadingEntryInfo(llvm::Function *F, llvm::Constant *ID);
+
+  /// \brief Creates and registers offloading binary descriptor for the current
+  /// compilation unit. The function that does the registration is returned.
+  llvm::Function *createOffloadingBinaryDescriptorRegistration();
+
+  /// \brief Creates offloading entry for the provided address \a Addr,
+  /// name \a Name and size \a Size.
+  void createOffloadEntry(llvm::Constant *Addr, StringRef Name, uint64_t Size);
+
+  /// \brief Returns __tgt_offload_entry type.
+  llvm::StructType *getTgtOffloadEntryTy();
 
+  /// \brief Returns __tgt_device_image type.
+  llvm::StructType *getTgtDeviceImageTy();
+
+  /// \brief Returns __tgt_bin_desc type.
+  llvm::StructType *getTgtBinaryDescriptorTy();
 
   /// \brief Build type kmp_routine_entry_t (if not built yet).
   void emitKmpRoutineEntryT(QualType KmpInt32Ty);
 
   /// \brief Emits object of ident_t type with info for source location.
   /// \param Flags Flags for OpenMP location.
   ///
   llvm::Value *emitUpdateLocation(CodeGenFunction &CGF, SourceLocation Loc,
public:
   /// \param CancelRegion Region kind for which the cancel must be emitted.
   ///
   virtual void emitCancelCall(CodeGenFunction &CGF, SourceLocation Loc,
                               OpenMPDirectiveKind CancelRegion);
 
   /// \brief Emit outilined function for 'target' directive.
   /// \param D Directive to emit.
   /// \param CodeGen Code generation sequence for the \a D directive.
-  virtual llvm::Value *
-  emitTargetOutlinedFunction(const OMPExecutableDirective &D,
-                             const RegionCodeGenTy &CodeGen);
+  /// \param OutlinedFn Outlined function value to be defined by this call.
+  /// \param OutlinedFnID Outlined function ID value to be defined by this call.
+  virtual void emitTargetOutlinedFunction(const OMPExecutableDirective &D,
+                                          const RegionCodeGenTy &CodeGen,
+                                          llvm::Function *&OutlinedFn,
+                                          llvm::Constant *&OutlinedFnID);
 
   /// \brief Emit the target offloading code associated with \a D. The emitted
   /// code attempts offloading the execution to the device, an the event of
   /// a failure it executes the host version outlined in \a OutlinedFn.
   /// \param D Directive to emit.
   /// \param OutlinedFn Host version of the code to be offloaded.
+  /// \param OutlinedFnID ID of host version of the code to be offloaded.
   /// \param IfCond Expression evaluated in if clause associated with the target
   /// directive, or null if no if clause is used.
   /// \param Device Expression evaluated in device clause associated with the
   /// target directive, or null if no device clause is used.
   virtual void emitTargetCall(CodeGenFunction &CGF,
                               const OMPExecutableDirective &D,
-                              llvm::Value *OutlinedFn, const Expr *IfCond,
+                              llvm::Value *OutlinedFn,
+                              llvm::Value *OutlinedFnID, const Expr *IfCond,
                               const Expr *Device,
                               ArrayRef<llvm::Value *> VLASizesInit);
+
+  /// \brief Creates the offloading descriptor in the event any target region
+  /// was emitted in the current module and return the function that registers
+  /// it.
+  virtual llvm::Function *emitRegistrationFunction();
 };
 
 } // namespace CodeGen
 } // namespace clang
 
 #endif

lib/CodeGen/CGOpenMPRuntime.cpp

 //===----- CGOpenMPRuntime.cpp - Interface to OpenMP Runtimes -------------===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 //
 // This provides a class for OpenMP runtime code generation.
 //
 //===----------------------------------------------------------------------===//
 
 #include "CGOpenMPRuntime.h"
 #include "CodeGenFunction.h"
 #include "CGCleanup.h"
+#include "CGCXXABI.h"
 #include "clang/AST/Decl.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "llvm/ADT/ArrayRef.h"
 #include "llvm/IR/CallSite.h"
 #include "llvm/IR/DerivedTypes.h"
 #include "llvm/IR/GlobalValue.h"
 #include "llvm/IR/Value.h"
+#include "llvm/Support/Format.h"
 #include "llvm/Support/raw_ostream.h"
 #include <cassert>
 
 using namespace clang;
 using namespace CodeGen;
 
 namespace {
 /// \brief Base class for handling code generation inside OpenMP regions.

 LValue CGOpenMPTaskOutlinedRegionInfo::getThreadIDVariableLValue(
     CodeGenFunction &CGF) {
   return CGF.MakeNaturalAlignAddrLValue(
       CGF.GetAddrOfLocalVar(getThreadIDVariable()),
       getThreadIDVariable()->getType());
 }
 
 CGOpenMPRuntime::CGOpenMPRuntime(CodeGenModule &CGM)
-    : CGM(CGM), DefaultOpenMPPSource(nullptr), KmpRoutineEntryPtrTy(nullptr) {
+    : CGM(CGM), DefaultOpenMPPSource(nullptr), KmpRoutineEntryPtrTy(nullptr),
+      TgtOffloadEntryTy(nullptr), TgtDeviceImageTy(nullptr),
+      TgtBinaryDescriptorTy(nullptr) {
   IdentTy = llvm::StructType::create(
       "ident_t", CGM.Int32Ty /* reserved_1 */, CGM.Int32Ty /* flags */,
       CGM.Int32Ty /* reserved_2 */, CGM.Int32Ty /* reserved_3 */,
       CGM.Int8PtrTy /* psource */, nullptr);
   // Build void (*kmpc_micro)(kmp_int32 *global_tid, kmp_int32 *bound_tid,...)
   llvm::Type *MicroParams[] = {llvm::PointerType::getUnqual(CGM.Int32Ty),
                                llvm::PointerType::getUnqual(CGM.Int32Ty)};
   Kmpc_MicroTy = llvm::FunctionType::get(CGM.VoidTy, MicroParams, true);
llvm::Type *TypeParams[] = {CGM.Int32Ty,
                                 CGM.VoidPtrPtrTy,
                                 CGM.SizeTy->getPointerTo(),
                                 CGM.Int32Ty->getPointerTo()};
     llvm::FunctionType *FnTy =
         llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
     RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target");
     break;
   }
+  case OMPRTL__tgt_register_lib: {
+    // Build void __tgt_register_lib(__tgt_bin_desc *desc);
+    llvm::Type *TypeParams[] = {getTgtBinaryDescriptorTy()->getPointerTo()};
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_register_lib");
+    break;
+  }
+  case OMPRTL__tgt_unregister_lib: {
+    // Build void __tgt_unregister_lib(__tgt_bin_desc *desc);
+    llvm::Type *TypeParams[] = {getTgtBinaryDescriptorTy()->getPointerTo()};
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_unregister_lib");
+    break;
+  }
   }
   return RTLFn;
 }
 
 llvm::Constant *CGOpenMPRuntime::createForStaticInitFunction(unsigned IVSize,
                                                              bool IVSigned) {
   assert((IVSize == 32 || IVSize == 64) &&
          "IV size is not compatible with the omp runtime");
enum KmpTaskTFields {
   KmpTaskTRoutine,
   /// \brief Partition id for the untied tasks.
   KmpTaskTPartId,
   /// \brief Function with call of destructors for private variables.
   KmpTaskTDestructors,
 };
 } // namespace
 
+void CGOpenMPRuntime::registerOffloadingEntryInfo(llvm::Function *F,
+                                                  llvm::Constant *ID) {
+  OffloadingEntriesInfo.push_back(OffloadingEntryInfo(F, ID));
+}
+
+/// \brief Create a Ctor/Dtor-like function whose body is emitted through
+/// \a Codegen. This is used to emit the two functions that register and
+/// unregister the descriptor of the current compilation unit.
+static llvm::Function *
+createOffloadingBinaryDescriptorFunction(CodeGenModule &CGM, StringRef Name,
+                                         const RegionCodeGenTy &Codegen) {
+  auto &C = CGM.getContext();
+  FunctionArgList Args;
+  ImplicitParamDecl DummyPtr(C, /*DC=*/nullptr, SourceLocation(),
+                             /*Id=*/nullptr, C.VoidPtrTy);
+  Args.push_back(&DummyPtr);
+
+  CodeGenFunction CGF(CGM);
+  auto &FI = CGM.getTypes().arrangeFreeFunctionDeclaration(
+      C.VoidTy, Args, FunctionType::ExtInfo(),
+      /*isVariadic=*/false);
+  auto FTy = CGM.getTypes().GetFunctionType(FI);
+  auto *Fn =
+      CGM.CreateGlobalInitOrDestructFunction(FTy, Name, SourceLocation());
+  CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, FI, Args, SourceLocation());
+  Codegen(CGF);
+  CGF.FinishFunction();
+  return Fn;
+}
+
+llvm::Function *
+CGOpenMPRuntime::createOffloadingBinaryDescriptorRegistration() {
+  auto &M = CGM.getModule();
+  auto &C = CGM.getContext();
+
+  // Get list of devices we care about
+  auto &Devices = CGM.getLangOpts().OMPTargetTriples;
+
+  // We should be creating an offloading descriptor only if there are devices
+  // specified.
+  assert(!Devices.empty() && "No OpenMP offloading devices??");
+
+  // Create the external variables that will point to the begin and end of the
+  // host entries section. These will be defined by the linker.
+
+  llvm::GlobalVariable *HostEntriesBegin = new llvm::GlobalVariable(
+      M, getTgtOffloadEntryTy(), /*isConstant=*/true,
+      llvm::GlobalValue::ExternalLinkage, /*Initializer=*/0,
+      ".omp_offloading.entries_begin");
+  llvm::GlobalVariable *HostEntriesEnd = new llvm::GlobalVariable(
+      M, getTgtOffloadEntryTy(), /*isConstant=*/true,
+      llvm::GlobalValue::ExternalLinkage, /*Initializer=*/0,
+      ".omp_offloading.entries_end");
+
+  // Create all device images
+  llvm::SmallVector<llvm::Constant *, 4> DeviceImagesEntires;
+
+  for (unsigned i = 0; i < Devices.size(); ++i) {
+    StringRef T = Devices[i].getTriple();
+    auto *ImgBegin = new llvm::GlobalVariable(
+        M, CGM.Int8Ty, /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage,
+        /*Initializer=*/0, Twine(".omp_offloading.img_start.") + Twine(T));
+    auto *ImgEnd = new llvm::GlobalVariable(
+        M, CGM.Int8Ty, /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage,
+        /*Initializer=*/0, Twine(".omp_offloading.img_end.") + Twine(T));
+
+    llvm::Constant *Dev =
+        llvm::ConstantStruct::get(getTgtDeviceImageTy(), ImgBegin, ImgEnd,
+                                  HostEntriesBegin, HostEntriesEnd, nullptr);
+    DeviceImagesEntires.push_back(Dev);
+  }
+
+  // Create device images global array.
+  llvm::ArrayType *DeviceImagesInitTy =
+      llvm::ArrayType::get(getTgtDeviceImageTy(), DeviceImagesEntires.size());
+  llvm::Constant *DeviceImagesInit =
+      llvm::ConstantArray::get(DeviceImagesInitTy, DeviceImagesEntires);
+
+  llvm::GlobalVariable *DeviceImages = new llvm::GlobalVariable(
+      M, DeviceImagesInitTy, /*isConstant=*/true,
+      llvm::GlobalValue::InternalLinkage, DeviceImagesInit,
+      ".omp_offloading.device_images");
+  DeviceImages->setUnnamedAddr(true);
+
+  // This is a Zero array to be used in the creation of the constant expressions
+  llvm::Constant *Index[] = {llvm::Constant::getNullValue(CGM.Int32Ty),
+                             llvm::Constant::getNullValue(CGM.Int32Ty)};
+
+  // Create the target region descriptor.
+  llvm::Constant *TargetRegionsDescriptorInit = llvm::ConstantStruct::get(
+      getTgtBinaryDescriptorTy(),
+      llvm::ConstantInt::get(CGM.Int32Ty, Devices.size()),
+      llvm::ConstantExpr::getGetElementPtr(DeviceImagesInitTy, DeviceImages,
+                                           Index),
+      HostEntriesBegin, HostEntriesEnd, nullptr);
+
+  auto *Desc = new llvm::GlobalVariable(
+      M, getTgtBinaryDescriptorTy(), /*isConstant=*/true,
+      llvm::GlobalValue::InternalLinkage, TargetRegionsDescriptorInit,
+      ".omp_offloading.descriptor");
+
+  // Emit code to register or unregister the descriptor at execution
+  // startup or closing, respectively.
+
+  // Create a variable to drive the registration and unregistration of the
+  // descriptor, so we can reuse the logic that emits Ctors and Dtors.
+  auto *IdentInfo = &C.Idents.get(".omp_offloading.reg_unreg_var");
+  ImplicitParamDecl RegUnregVar(C, C.getTranslationUnitDecl(), SourceLocation(),
+                                IdentInfo, C.CharTy);
+
+  auto *UnRegFn = createOffloadingBinaryDescriptorFunction(
+      CGM, ".omp_offloading.descriptor_unreg", [&](CodeGenFunction &CGF) {
+        CGF.EmitCallOrInvoke(createRuntimeFunction(OMPRTL__tgt_unregister_lib),
+                             Desc);
+      });
+  auto *RegFn = createOffloadingBinaryDescriptorFunction(
+      CGM, ".omp_offloading.descriptor_reg", [&](CodeGenFunction &CGF) {
+        CGF.EmitCallOrInvoke(createRuntimeFunction(OMPRTL__tgt_register_lib),
+                             Desc);
+        CGM.getCXXABI().registerGlobalDtor(CGF, RegUnregVar, UnRegFn, Desc);
+      });
+  return RegFn;
+}
+
+void CGOpenMPRuntime::createOffloadEntry(llvm::Constant *Addr, StringRef Name,
+                                         uint64_t Size) {
+  auto *TgtOffloadEntryType = getTgtOffloadEntryTy();
+  llvm::LLVMContext &C = CGM.getModule().getContext();
+  llvm::Module &M = CGM.getModule();
+
+  // Make sure the address has the right type.
+  llvm::Constant *AddrPtr = llvm::ConstantExpr::getBitCast(Addr, CGM.VoidPtrTy);
+
+  // Create constant string with the name.
+  llvm::Constant *StrPtrInit = llvm::ConstantDataArray::getString(C, Name);
+
+  llvm::GlobalVariable *Str =
+      new llvm::GlobalVariable(M, StrPtrInit->getType(), /*isConstant=*/true,
+                               llvm::GlobalValue::InternalLinkage, StrPtrInit,
+                               ".omp_offloading.entry_name");
+  Str->setUnnamedAddr(true);
+  llvm::Constant *StrPtr = llvm::ConstantExpr::getBitCast(Str, CGM.Int8PtrTy);
+
+  // Create the entry struct.
+  llvm::Constant *EntryInit = llvm::ConstantStruct::get(
+      TgtOffloadEntryType, AddrPtr, StrPtr,
+      llvm::ConstantInt::get(CGM.SizeTy, Size), nullptr);
+  llvm::GlobalVariable *Entry = new llvm::GlobalVariable(
+      M, TgtOffloadEntryType, true, llvm::GlobalValue::ExternalLinkage,
+      EntryInit, ".omp_offloading.entry");
+
+  // The entry has to be created in the section the linker expects it to be.
+  Entry->setSection(".omp_offloading.entries");
+  // We can't have any padding between symbols, so we need to have 1-byte
+  // alignment.
+  Entry->setAlignment(1);
+  return;
+}
+
+llvm::StructType *CGOpenMPRuntime::getTgtOffloadEntryTy() {
+
+  // Make sure the type of the entry is already created. This is the type we
+  // have to create:
+  // struct __tgt_offload_entry{
+  //   void      *addr;       // Pointer to the offload entry info.
+  //                          // (function or global)
+  //   char      *name;       // Name of the function or global.
+  //   size_t     size;       // Size of the entry info (0 if it a function).
+  // };
+  if (!TgtOffloadEntryTy)
+    TgtOffloadEntryTy = llvm::StructType::create(
+        "tgt_offload_entry", CGM.VoidPtrTy, CGM.Int8PtrTy, CGM.SizeTy, nullptr);
+  return TgtOffloadEntryTy;
+}
+
+llvm::StructType *CGOpenMPRuntime::getTgtDeviceImageTy() {
+  // These are the types we need to build:
+  // struct __tgt_device_image{
+  // void   *ImageStart;       // Pointer to the target code start.
+  // void   *ImageEnd;         // Pointer to the target code end.
+  // // We also add the host entries to the device image, as it may be useful
+  // // for the target runtime to have access to that information.
+  // __tgt_offload_entry  *EntriesBegin;   // Begin of the table with all
+  //                                       // the entries.
+  // __tgt_offload_entry  *EntriesEnd;     // End of the table with all the
+  //                                       // entries (non inclusive).
+  // };
+  if (!TgtDeviceImageTy)
+    TgtDeviceImageTy = llvm::StructType::create(
+        "tgt_device_image", CGM.VoidPtrTy, CGM.VoidPtrTy,
+        getTgtOffloadEntryTy()->getPointerTo(),
+        getTgtOffloadEntryTy()->getPointerTo(), nullptr);
+  return TgtDeviceImageTy;
+}
+
+llvm::StructType *CGOpenMPRuntime::getTgtBinaryDescriptorTy() {
+  // struct __tgt_bin_desc{
+  //   int32_t              NumDevices;      // Number of devices supported.
+  //   __tgt_device_image   *DeviceImages;   // Arrays of device images
+  //                                         // (one per device).
+  //   __tgt_offload_entry  *EntriesBegin;   // Begin of the table with all the
+  //                                         // entries.
+  //   __tgt_offload_entry  *EntriesEnd;     // End of the table with all the
+  //                                         // entries (non inclusive).
+  // };
+  if (!TgtBinaryDescriptorTy)
+    TgtBinaryDescriptorTy = llvm::StructType::create(
+        "tgt_bin_desc", CGM.Int32Ty, getTgtDeviceImageTy()->getPointerTo(),
+        getTgtOffloadEntryTy()->getPointerTo(),
+        getTgtOffloadEntryTy()->getPointerTo(), nullptr);
+  return TgtBinaryDescriptorTy;
+}
+
 void CGOpenMPRuntime::emitKmpRoutineEntryT(QualType KmpInt32Ty) {
   if (!KmpRoutineEntryPtrTy) {
     // Build typedef kmp_int32 (* kmp_routine_entry_t)(kmp_int32, void *); type.
     auto &C = CGM.getContext();
     QualType KmpRoutineEntryTyArgs[] = {KmpInt32Ty, C.VoidPtrTy};
     FunctionProtoType::ExtProtoInfo EPI;
     KmpRoutineEntryPtrQTy = C.getPointerType(
         C.getFunctionType(KmpInt32Ty, KmpRoutineEntryTyArgs, EPI));

 /// \code
 /// void .omp_offloading_entry.(ty1 capture1, ..., tyn capture n) {
 ///   Context ctx;
 ///   // fillup the context with all the arguments.
 ///   .omp_offloading.(ctx);
 ///   return;
 /// }
 /// \endcode
-static llvm::Value *emitProxyTargetFunction(CodeGenModule &CGM,
+static llvm::Function *emitProxyTargetFunction(CodeGenModule &CGM,
                                                const CapturedStmt &CS,
                                                SourceLocation Loc,
                                                llvm::Value *TargetFunction) {
   auto &C = CGM.getContext();
+  auto &SM = C.getSourceManager();
+
+  // Create a unique name for the proxy/entry function that using the source
+  // location information of the current target region. The name will be
+  // something like:
+  //
+  // .omp_offloading.AAAA_AA.lBB.cCC
+  //
+  // where AAAA_AA is an ID unique to the file, BB is the line number of the
+  // target region and CC is the column number of the target region.
+  SmallString<64> EntryFnName;
+
+  // The loc should be always valid and have a file ID (the user cannot use
+  // #pragma directives in macros)
+
+  assert(Loc.isValid() && "Source location is expected to be always valid.");
+  assert(Loc.isFileID() && "Source location is expected to refer to a file.");
+
+  PresumedLoc PLoc = SM.getPresumedLoc(Loc);
+  assert(PLoc.isValid() && "Source location is expected to be always valid.");
+
+  llvm::sys::fs::UniqueID ID;
+  if (llvm::sys::fs::getUniqueID(PLoc.getFilename(), ID))
+    llvm_unreachable("Source file with target region no longer exists!");
+
+  {
+    llvm::raw_svector_ostream OS(EntryFnName);
+    OS << ".omp_offloading" << llvm::format(".%llx", ID.getFile())
+       << llvm::format(".%llx", ID.getDevice()) << ".l" << PLoc.getLine()
+       << ".c" << PLoc.getColumn();
+  }
 
   // Collect the arguments of the main function.
   FunctionArgList Args;
   auto *RDecl = CS.getCapturedRecordDecl();
   auto RI = RDecl->field_begin();
   for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(),
                                             CE = CS.capture_end();
        CI != CE; ++CI, ++RI) {
static llvm::Function *emitProxyTargetFunction(CodeGenModule &CGM,
 
   // Create the main function.
   FunctionType::ExtInfo Info;
   auto &FnInfo =
       CGM.getTypes().arrangeFreeFunctionDeclaration(C.VoidTy, Args, Info,
                                                     /*isVariadic=*/false);
   auto *FnTy = CGM.getTypes().GetFunctionType(FnInfo);
   auto *Fn = llvm::Function::Create(FnTy, llvm::GlobalValue::InternalLinkage,
-                                    ".omp_offloading_entry.", &CGM.getModule());
+                                    EntryFnName, &CGM.getModule());
   CGM.SetLLVMFunctionAttributes(/*D=*/nullptr, FnInfo, Fn);
   CodeGenFunction CGF(CGM);
   CGF.disableDebugInfo();
   CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, FnInfo, Args);
 
   // Create the context.
   auto ContextType = C.getTagDeclType(RDecl);
   auto *Context = CGF.CreateMemTemp(ContextType, ".omp_offloading.context");
for (RecordDecl::field_iterator RI = RDecl->field_begin(),
     }
     CGF.EmitStoreOfScalar(Arg.getScalarVal(), Dst);
   }
   CGF.EmitCallOrInvoke(TargetFunction, Context);
   CGF.FinishFunction();
   return Fn;
 }
 
-llvm::Value *
-CGOpenMPRuntime::emitTargetOutlinedFunction(const OMPExecutableDirective &D,
-                                            const RegionCodeGenTy &CodeGen) {
+void CGOpenMPRuntime::emitTargetOutlinedFunction(
+    const OMPExecutableDirective &D, const RegionCodeGenTy &CodeGen,
+    llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID) {
+
   const CapturedStmt &CS = *cast<CapturedStmt>(D.getAssociatedStmt());
 
   CodeGenFunction CGF(CGM, true);
   CGOpenMPTargetRegionInfo CGInfo(CS, CodeGen);
   CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
   auto *Fn = CGF.GenerateCapturedStmtFunction(CS);
   Fn->addFnAttr(llvm::Attribute::AlwaysInline);
 
-  return emitProxyTargetFunction(CGM, CS, D.getLocStart(), Fn);
+  OutlinedFn = emitProxyTargetFunction(CGM, CS, D.getLocStart(), Fn);
+
+  // If we don't have any devices specified by the user, we don't need to bother
+  // registering the target region.
+  if (CGM.getLangOpts().OMPTargetTriples.empty())
+    return;
+
+  // The target region ID is used by the runtime library to identify the current
+  // target region, so it only has to be unique and not necessarily point to
+  // anything. It could be the pointer to the outlined function that implements
+  // the target region, but we aren't using that so that the compiler doesn't
+  // need to keep that, and could therefore inline the host function if proven
+  // worthwhile during optimization.
+
+  OutlinedFnID = new llvm::GlobalVariable(
+      CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
+      llvm::GlobalValue::PrivateLinkage,
+      llvm::Constant::getNullValue(CGM.Int8Ty), ".omp_offload.region_id");
+
+  // Register the information for the entry associated with this target region.
+  registerOffloadingEntryInfo(OutlinedFn, OutlinedFnID);
+  return;
 }
 
 void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
                                      const OMPExecutableDirective &D,
                                      llvm::Value *OutlinedFn,
+                                     llvm::Value *OutlinedFnID,
                                      const Expr *IfCond, const Expr *Device,
                                      ArrayRef<llvm::Value *> VLASizesInit) {
 
   // Fill up the arrays with the all the captured variables.
   SmallVector<llvm::Value *, 16> BasePointers;
   SmallVector<llvm::Value *, 16> Pointers;
   SmallVector<llvm::Value *, 16> Sizes;
   SmallVector<unsigned, 16> MapTypes;
for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(),
     }
 
     BasePointers.push_back(BasePointer);
     Pointers.push_back(Pointer);
     Sizes.push_back(Size);
     MapTypes.push_back(MapType);
   }
 
+  // If we don't have any devices specified by the user, we don't need to bother
+  // about emitting any runtime calls. Instead, we just call the host version.
+  if (CGM.getLangOpts().OMPTargetTriples.empty()) {
+    CGF.Builder.CreateCall(OutlinedFn, BasePointers);
+    return;
+  }
+
   if (IfCond) {
     // Check if the if clause conditional always evaluates to true or false.
     // If it evaluates to false, we only need to emit the host version of the
     // target region. If it evaluates to true, we can proceed with the codegen
     // as if no if clause was provided.
     bool CondConstant;
     if (CGF.ConstantFoldsToSimpleInteger(IfCond, CondConstant)) {
       if (CondConstant) {
void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
   } else {
     BasePointersArray = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
     PointersArray = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
     SizesArray = llvm::ConstantPointerNull::get(CGM.SizeTy->getPointerTo());
     MapTypesArray = llvm::ConstantPointerNull::get(CGM.Int32Ty->getPointerTo());
   }
 
   // On top of the arrays that were filled up, the target offloading call takes
-  // as arguments the device id as well as the host pointer. The host pointer
-  // is used by the runtime library to identify the current target region, so
-  // it only has to be unique and not necessarily point to anything. It could be
-  // the pointer to the outlined function that implements the target region, but
-  // we aren't using that so that the compiler doesn't need to keep that, and
-  // could therefore inline the host function if proven worthwhile during
-  // optimization.
-
-  llvm::Value *HostPtr = new llvm::GlobalVariable(
-      CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
-      llvm::GlobalValue::PrivateLinkage,
-      llvm::Constant::getNullValue(CGM.Int8Ty), ".offload_hstptr");
+  // as arguments the device id as well as the ID of the target region.
 
   // Emit device ID if any.
   llvm::Value *DeviceID;
   if (Device)
     DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
                                          CGM.Int32Ty, /*isSigned=*/true);
   else
     DeviceID = CGF.Builder.getInt32(OMP_DEVICEID_UNDEF);
 
-  llvm::Value *OffloadingArgs[] = {DeviceID,          HostPtr,       PointerNum,
+  llvm::Value *OffloadingArgs[] = {DeviceID,          OutlinedFnID,  PointerNum,
                                    BasePointersArray, PointersArray, SizesArray,
                                    MapTypesArray};
   auto Return = CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_target),
                                     OffloadingArgs);
   auto Error = CGF.Builder.CreateIsNotNull(Return);
   CGF.Builder.CreateCondBr(Error, OffloadFailBlock, ContBlock);
 
   CGF.EmitBlock(OffloadFailBlock);
   CGF.Builder.CreateCall(OutlinedFn, BasePointers);
   CGF.EmitBranch(ContBlock);
 
   CGF.EmitBlock(ContBlock, /*IsFinished=*/true);
   return;
 }
+
+llvm::Function *CGOpenMPRuntime::emitRegistrationFunction() {
+  // If we have offloading in the current module, we need to emit the entries
+  // now and register the offloading descriptor.
+  if (!OffloadingEntriesInfo.empty()) {
+    for (auto &Info : OffloadingEntriesInfo) {
+      // For the moment all entries are target regions, so they do not have a
+      // size associated as global variables do.
+      createOffloadEntry(Info.ID, Info.Addr->getName(), 0);
+    }
+
+    /// Create and register the offloading binary descriptors. This is the main
+    /// entity that captures all the information about offloading in the current
+    /// compilation unit.
+    return createOffloadingBinaryDescriptorRegistration();
+  }
+  return nullptr;
+}

lib/CodeGen/CGStmtOpenMP.cpp

void CodeGenFunction::EmitOMPTargetDirective(const OMPTargetDirective &S) {
   LexicalScope Scope(*this, S.getSourceRange());
   const CapturedStmt &CS = *cast<CapturedStmt>(S.getAssociatedStmt());
 
   // Emit target region as a standalone region.
   auto &&CodeGen = [&CS](CodeGenFunction &CGF) {
     CGF.EmitStmt(CS.getCapturedStmt());
   };
 
+  llvm::Function *Fn;
+  llvm::Constant *FnID;
+
   // Obtain the target region outlined function.
-  llvm::Value *Fn =
-      CGM.getOpenMPRuntime().emitTargetOutlinedFunction(S, CodeGen);
+  CGM.getOpenMPRuntime().emitTargetOutlinedFunction(S, CodeGen, Fn, FnID);
 
   // Check if we have any if clause associated with the directive.
   const Expr *IfCond = nullptr;
   if (auto C = S.getSingleClause(OMPC_if)) {
     IfCond = cast<OMPIfClause>(C)->getCondition();
   }
 
   // Check if we have any device clause associated with the directive.
   const Expr *Device = nullptr;
   if (auto C = S.getSingleClause(OMPC_device)) {
     Device = cast<OMPDeviceClause>(C)->getDevice();
   }
 
   SmallVector<llvm::Value *, 8> VLASizesInit;
   for (auto F : CS.getCapturedRecordDecl()->fields())
     if (F->hasCapturedVLAType()) {
       auto *V = VLASizeMap[F->getCapturedVLAType()->getSizeExpr()];
       assert(V && "VLA size must exist!");
       VLASizesInit.push_back(V);
     }
 
-  CGM.getOpenMPRuntime().emitTargetCall(*this, S, Fn, IfCond, Device,
+  CGM.getOpenMPRuntime().emitTargetCall(*this, S, Fn, FnID, IfCond, Device,
                                         VLASizesInit);
 }
 
 void CodeGenFunction::EmitOMPTeamsDirective(const OMPTeamsDirective &) {
   llvm_unreachable("CodeGen for 'omp teams' is not supported yet.");
 }
 
 void CodeGenFunction::EmitOMPCancellationPointDirective(

lib/CodeGen/CodeGenModule.cpp

if (llvm::Function *ObjCInitFunction = ObjCRuntime->ModuleInitFunction())
       AddGlobalCtor(ObjCInitFunction);
   if (Context.getLangOpts().CUDA && !Context.getLangOpts().CUDAIsDevice &&
       CUDARuntime) {
     if (llvm::Function *CudaCtorFunction = CUDARuntime->makeModuleCtorFunction())
       AddGlobalCtor(CudaCtorFunction);
     if (llvm::Function *CudaDtorFunction = CUDARuntime->makeModuleDtorFunction())
       AddGlobalDtor(CudaDtorFunction);
   }
+  if (OpenMPRuntime)
+    if (llvm::Function *OpenMPRegistrationFunction =
+            OpenMPRuntime->emitRegistrationFunction())
+      AddGlobalCtor(OpenMPRegistrationFunction, 0);
   if (PGOReader && PGOStats.hasDiagnostics())
     PGOStats.reportDiagnostics(getDiags(), getCodeGenOpts().MainFileName);
   EmitCtorList(GlobalCtors, "llvm.global_ctors");
   EmitCtorList(GlobalDtors, "llvm.global_dtors");
   EmitGlobalAnnotations();
   EmitStaticExternCAliases();
   EmitDeferredUnusedCoverageMappings();
   if (CoverageMapping)

lib/Frontend/CompilerInvocation.cpp

if (Arg *A = Args.getLastArg(OPT_fms_memptr_rep_EQ)) {
     Opts.setMSPointerToMemberRepresentationMethod(InheritanceModel);
   }
 
   // Check if -fopenmp is specified.
   Opts.OpenMP = Args.hasArg(options::OPT_fopenmp);
   Opts.OpenMPUseTLS =
       Opts.OpenMP && !Args.hasArg(options::OPT_fnoopenmp_use_tls);
 
+  // Get the OpenMP target triples if any
+  if (Arg *A = Args.getLastArg(options::OPT_omptargets_EQ)) {
+
+    for (unsigned i = 0; i < A->getNumValues(); ++i) {
+      llvm::Triple TT(A->getValue(i));
+
+      if (TT.getArch() == llvm::Triple::UnknownArch)
+        Diags.Report(clang::diag::err_drv_invalid_omp_target) << A->getValue(i);
+      else
+        Opts.OMPTargetTriples.push_back(TT);
+    }
+  }
+
   // Record whether the __DEPRECATED define was requested.
   Opts.Deprecated = Args.hasFlag(OPT_fdeprecated_macro,
                                  OPT_fno_deprecated_macro,
                                  Opts.Deprecated);
 
   // FIXME: Eliminate this dependency.
   unsigned Opt = getOptimizationLevel(Args, IK, Diags),
        OptSize = getOptimizationLevelSize(Args);

lib/Serialization/ASTReader.cpp

#include "clang/Basic/Sanitizers.def"
 
   // Comment options.
   for (unsigned N = Record[Idx++]; N; --N) {
     LangOpts.CommentOpts.BlockCommandNames.push_back(
       ReadString(Record, Idx));
   }
   LangOpts.CommentOpts.ParseAllComments = Record[Idx++];
 
+  // OpenMP offloading options.
+  for (unsigned N = Record[Idx++]; N; --N) {
+    LangOpts.OMPTargetTriples.push_back(llvm::Triple(ReadString(Record, Idx)));
+  }
+
   return Listener.ReadLanguageOptions(LangOpts, Complain,
                                       AllowCompatibleDifferences);
 }
 
 bool ASTReader::ParseTargetOptions(const RecordData &Record, bool Complain,
                                    ASTReaderListener &Listener,
                                    bool AllowCompatibleDifferences) {
   unsigned Idx = 0;

lib/Serialization/ASTWriter.cpp

#include "clang/Basic/Sanitizers.def"
   for (CommentOptions::BlockCommandNamesTy::const_iterator
            I = LangOpts.CommentOpts.BlockCommandNames.begin(),
            IEnd = LangOpts.CommentOpts.BlockCommandNames.end();
        I != IEnd; ++I) {
     AddString(*I, Record);
   }
   Record.push_back(LangOpts.CommentOpts.ParseAllComments);
 
+  // OpenMP offloading options.
+  Record.push_back(LangOpts.OMPTargetTriples.size());
+  for (auto &T : LangOpts.OMPTargetTriples)
+    AddString(T.getTriple(), Record);
+
   Stream.EmitRecord(LANGUAGE_OPTIONS, Record);
 
   // Target options.
   Record.clear();
   const TargetInfo &Target = Context.getTargetInfo();
   const TargetOptions &TargetOpts = Target.getTargetOpts();
   AddString(TargetOpts.Triple, Record);
   AddString(TargetOpts.CPU, Record);

test/OpenMP/target_codegen.cpp

-// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
-// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -omptargets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -omptargets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -omptargets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -omptargets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -omptargets=i386-pc-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -omptargets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
 // expected-no-diagnostics
 // REQUIRES: powerpc-registered-target
 #ifndef HEADER
 #define HEADER
 
 // CHECK-DAG: [[TT:%.+]] = type { i64, i8 }
 // CHECK-DAG: [[S1:%.+]] = type { double }
 

test/OpenMP/target_codegen_registration.cpp

+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -omptargets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -omptargets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -omptargets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -omptargets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -omptargets=i386-pc-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -omptargets=i386-pc-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+
+// Check that no target code is emmitted if no omptests flag was provided.
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s -check-prefix=CHECK-NTARGET
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+// CHECK-DAG: [[SA:%.+]] = type { [4 x i32] }
+// CHECK-DAG: [[SB:%.+]] = type { [8 x i32] }
+// CHECK-DAG: [[SC:%.+]] = type { [16 x i32] }
+// CHECK-DAG: [[SD:%.+]] = type { [32 x i32] }
+// CHECK-DAG: [[SE:%.+]] = type { [64 x i32] }
+// CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]] }
+// CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* }
+// CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* }
+
+// CHECK-DAG: [[A1:@.+]] = internal global [[SA]]
+// CHECK-DAG: [[A2:@.+]] = global [[SA]]
+// CHECK-DAG: [[B1:@.+]] = global [[SB]]
+// CHECK-DAG: [[B2:@.+]] = global [[SB]]
+// CHECK-DAG: [[C1:@.+]] = internal global [[SC]]
+// CHECK-DAG: [[D1:@.+]] = global [[SD]]
+// CHECK-DAG: [[E1:@.+]] = global [[SE]]
+
+// CHECK-NTARGET-DAG: [[SA:%.+]] = type { [4 x i32] }
+// CHECK-NTARGET-DAG: [[SB:%.+]] = type { [8 x i32] }
+// CHECK-NTARGET-DAG: [[SC:%.+]] = type { [16 x i32] }
+// CHECK-NTARGET-DAG: [[SD:%.+]] = type { [32 x i32] }
+// CHECK-NTARGET-DAG: [[SE:%.+]] = type { [64 x i32] }
+// CHECK-NTARGET-NOT: type { i8*,
+// CHECK-NTARGET-NOT: type { i32,
+
+// We have 7 target regions
+
+// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 3]
+// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 3]
+// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 3]
+// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 3]
+// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 3]
+// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 3]
+// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 3]
+
+// CHECK-NTARGET-NOT: private constant i8 0
+// CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i
+
+// CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:.+]]\00"
+// CHECK-DAG: [[ENTRY1:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR1]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
+// CHECK-DAG: [[NAMEPTR2:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME2:.+]]\00"
+// CHECK-DAG: [[ENTRY2:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR2]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
+// CHECK-DAG: [[NAMEPTR3:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME3:.+]]\00"
+// CHECK-DAG: [[ENTRY3:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR3]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
+// CHECK-DAG: [[NAMEPTR4:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME4:.+]]\00"
+// CHECK-DAG: [[ENTRY4:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR4]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
+// CHECK-DAG: [[NAMEPTR5:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME5:.+]]\00"
+// CHECK-DAG: [[ENTRY5:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR5]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
+// CHECK-DAG: [[NAMEPTR6:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME6:.+]]\00"
+// CHECK-DAG: [[ENTRY6:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR6]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
+// CHECK-DAG: [[NAMEPTR7:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME7:.+]]\00"
+// CHECK-DAG: [[ENTRY7:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR7]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
+
+// CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]]
+// CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]]
+// CHECK: [[DEVBEGIN:@.+]] = external constant i8
+// CHECK: [[DEVEND:@.+]] = external constant i8
+// CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [1 x [[DEVTY]]] [{{.+}} { i8* [[DEVBEGIN]], i8* [[DEVEND]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }]
+// CHECK: [[DESC:@.+]] = internal constant [[DSCTY]] { i32 1, [[DEVTY]]* getelementptr inbounds ([1 x [[DEVTY]]], [1 x [[DEVTY]]]* [[IMAGES]], i32 0, i32 0), [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }
+
+// We have 4 initializers, one for the 500 priority, another one for 501, or more for the default priority, and the last one for the offloading registration function.
+// CHECK: @llvm.global_ctors = appending global [4 x { i32, void ()*, i8* }] [
+// CHECK-SAME: { i32, void ()*, i8* } { i32 500, void ()* [[P500:@[^,]+]], i8* null },
+// CHECK-SAME: { i32, void ()*, i8* } { i32 501, void ()* [[P501:@[^,]+]], i8* null },
+// CHECK-SAME: { i32, void ()*, i8* } { i32 65535, void ()* [[PMAX:@[^,]+]], i8* null },
+// CHECK-SAME: { i32, void ()*, i8* } { i32 0, void ()* bitcast (void (i8*)* [[REGFN:@.+]] to void ()*), i8* null }]
+
+// CHECK-NTARGET: @llvm.global_ctors = appending global [3   x { i32, void ()*, i8* }] [
+
+extern int *R;
+
+struct SA {
+  int arr[4];
+  void foo() {
+    int a = *R;
+    a += 1;
+    *R = a;
+  }
+  SA() {
+    int a = *R;
+    a += 2;
+    *R = a;
+  }
+  ~SA() {
+    int a = *R;
+    a += 3;
+    *R = a;
+  }
+};
+
+struct SB {
+  int arr[8];
+  void foo() {
+    int a = *R;
+    #pragma omp target
+    a += 4;
+    *R = a;
+  }
+  SB() {
+    int a = *R;
+    a += 5;
+    *R = a;
+  }
+  ~SB() {
+    int a = *R;
+    a += 6;
+    *R = a;
+  }
+};
+
+struct SC {
+  int arr[16];
+  void foo() {
+    int a = *R;
+    a += 7;
+    *R = a;
+  }
+  SC() {
+    int a = *R;
+    #pragma omp target
+    a += 8;
+    *R = a;
+  }
+  ~SC() {
+    int a = *R;
+    a += 9;
+    *R = a;
+  }
+};
+
+struct SD {
+  int arr[32];
+  void foo() {
+    int a = *R;
+    a += 10;
+    *R = a;
+  }
+  SD() {
+    int a = *R;
+    a += 11;
+    *R = a;
+  }
+  ~SD() {
+    int a = *R;
+    #pragma omp target
+    a += 12;
+    *R = a;
+  }
+};
+
+struct SE {
+  int arr[64];
+  void foo() {
+    int a = *R;
+    #pragma omp target
+    a += 13;
+    *R = a;
+  }
+  SE() {
+    int a = *R;
+    #pragma omp target
+    a += 14;
+    *R = a;
+  }
+  ~SE() {
+    int a = *R;
+    #pragma omp target
+    a += 15;
+    *R = a;
+  }
+};
+
+// We have to make sure we us all the target regions:
+//CHECK-DAG: define internal void @[[NAME1]](
+//CHECK-DAG: call void @[[NAME1]](
+//CHECK-DAG: define internal void @[[NAME2]](
+//CHECK-DAG: call void @[[NAME2]](
+//CHECK-DAG: define internal void @[[NAME3]](
+//CHECK-DAG: call void @[[NAME3]](
+//CHECK-DAG: define internal void @[[NAME4]](
+//CHECK-DAG: call void @[[NAME4]](
+//CHECK-DAG: define internal void @[[NAME5]](
+//CHECK-DAG: call void @[[NAME5]](
+//CHECK-DAG: define internal void @[[NAME6]](
+//CHECK-DAG: call void @[[NAME6]](
+//CHECK-DAG: define internal void @[[NAME7]](
+//CHECK-DAG: call void @[[NAME7]](
+
+// CHECK-NTARGET-NOT: __tgt_target
+// CHECK-NTARGET-NOT: __tgt_register_lib
+// CHECK-NTARGET-NOT: __tgt_unregister_lib
+
+// We have 2 initializers with priority 500
+//CHECK: define internal void [[P500]](
+//CHECK:     call void @{{.+}}()
+//CHECK:     call void @{{.+}}()
+//CHECK-NOT: call void @{{.+}}()
+//CHECK:     ret void
+
+// We have 1 initializers with priority 501
+//CHECK: define internal void [[P501]](
+//CHECK:     call void @{{.+}}()
+//CHECK-NOT: call void @{{.+}}()
+//CHECK:     ret void
+
+// We have 4 initializers with default priority
+//CHECK: define internal void [[PMAX]](
+//CHECK:     call void @{{.+}}()
+//CHECK:     call void @{{.+}}()
+//CHECK:     call void @{{.+}}()
+//CHECK:     call void @{{.+}}()
+//CHECK-NOT: call void @{{.+}}()
+//CHECK:     ret void
+
+// Check registration and unregistration
+
+//CHECK:     define internal void [[UNREGFN:@.+]](i8*)
+//CHECK:     call i32 @__tgt_unregister_lib([[DSCTY]]* [[DESC]])
+//CHECK:     ret void
+//CHECK:     declare i32 @__tgt_unregister_lib([[DSCTY]]*)
+
+//CHECK:     define internal void [[REGFN]](i8*)
+//CHECK:     call i32 @__tgt_register_lib([[DSCTY]]* [[DESC]])
+//CHECK:     call i32 @__cxa_atexit(void (i8*)* [[UNREGFN]], i8* bitcast ([[DSCTY]]* [[DESC]] to i8*),
+//CHECK:     ret void
+//CHECK:     declare i32 @__tgt_register_lib([[DSCTY]]*)
+
+static __attribute__((init_priority(500))) SA a1;
+SA a2;
+SB __attribute__((init_priority(500))) b1;
+SB __attribute__((init_priority(501))) b2;
+static SC c1;
+SD d1;
+SE e1;
+
+
+int bar(int a){
+  int r = a;
+
+  a1.foo();
+  a2.foo();
+  b1.foo();
+  b2.foo();
+  c1.foo();
+  d1.foo();
+  e1.foo();
+
+  #pragma omp target
+  ++r;
+
+  return r + *R;
+}
+
+#endif

test/OpenMP/target_messages.cpp

 // RUN: %clang_cc1 -verify -fopenmp -std=c++11 -o - %s
+// RUN: not %clang_cc1 -fopenmp -std=c++11 -omptargets=aaa-bbb-ccc-ddd -o - %s 2>&1 | FileCheck %s
+// CHECK: error: OpenMP target is invalid: 'aaa-bbb-ccc-ddd'
 
 void foo() {
 }
 
 #pragma omp target // expected-error {{unexpected OpenMP directive '#pragma omp target'}}
 
 int main(int argc, char **argv) {
   #pragma omp target { // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}