[OpenMP] Add support for registering requires directives with the runtime

[gtbercea]

Summary

This patch adds support for the registration of the requires directives with the runtime.

Each requires directive clause will enable a particular flag to be set.

The set of flags is passed to the runtime to be checked for compatibility with other such flags coming from other object files.

The registration function is called whenever OpenMP is present even if a requires directive is not present. This helps detect cases in which requires directives are used inconsistently.

[ABataev]

lib/CodeGen/CGDecl.cpp

2576 ↗

(On Diff #194700)

You don't need to pass the reference to CodeGenModule to CGOpenMPRuntime class, it handles a reference to the CodeGenMode already.

lib/CodeGen/CGOpenMPRuntime.cpp

746

You must use LLVM_ENABLE_BITMASK_ENUMS_IN_NAMESPACE(); and LLVM_MARK_AS_BITMASK_ENUM(/* LargestFlag = */) from clang/Basic/BitmaskEnum.h to enable bit operations on your new bit-arithmetic based enumeric.

746

Add a comment for the whole new type

748

You should not handle all the possible flags for the requires directives, since you try to support only target-specific flags.

2323

Function return type must be void

3871

Why do you need a new member function? Can you make a static local function?

3883

Use getTypes().arrangeNullaryFunction()

3888

Use OpenMPOffloadingRequiresDirFlags instead of int64_t

7955

Seems to me, this must be in another patch, has nothing to do with this patch

9037

Why do you need the second function?

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

4946

Call this function after the target-specific processing.

lib/CodeGen/CodeGenModule.h

294 ↗

(On Diff #194700)

Why public? Must be private. Also, add comments for this new data member.
Also, I don't think this must be in CGM. It must be a member of CGOpenMPRuntime class.

[gtbercea]

• Address comments.

[gtbercea]

• Address comments.

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.cpp

3871

Same as for register lib.

3888

Leads to error if I do that. This enum behaves like OpenMPLocationFlags.

7955

I will split it.

9037

Second function eliminated.

[gtbercea]

• Fix enum.

[gtbercea]

• Split patch.

[gtbercea]

• Handle OpenMP simd case.

[AlexEichenberger]

see note above; apologies if it is already done and hiding somewhere I did not see

lib/CodeGen/CGOpenMPRuntime.h

641

Don't you need a bool for each characteristics? Your intention is to have one bit vector for each characteristics that matter to the compiler?

Also, it is my belief that you need to record 2 states:
unmaterialized (meaning I have not processed any target yet, so I should record any requires as they arrive)
finalized (I am processing a target, so the state is now fixed)

you need this in case you have an input like this:

declare target
int x
end declare target

requires unified memory

which is illegal

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

473

YOu don't need al these flags, add only target-specific.

1250

No need to do this initialization here

8973

Just HasRequiresUnifiedSharedMemory = true;

9040

Just add a check for OpenMPSimd mode here and do not emit anything in this case.

9056

No need for braces

9060

CGF.CGM.Int64Ty->CGM.Int64Ty

lib/CodeGen/CGOpenMPRuntime.h

1445

Why do you need this vertual funtion? I think static local is going to be enough

lib/CodeGen/CodeGenModule.cpp

415

You can remove the third nullptr argument

[gtbercea]

• Address comments.

[gtbercea]

• Remove atomic flags.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

8973

Not done

8973

You can do break; here, no need for further analysis

9055

Change the type from int64_t to OpenMPOffloadingRequiresDirFlags

lib/CodeGen/CGOpenMPRuntime.h

641

What about this comment?

1445

Can you make it const member function?

[gtbercea]

• Fix type.

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.cpp

8973

Ah I misunderstood your initial comment. I thought you were complaining about the break being in there.

8973

Because the member is non-static I have to provide the CGM.getOpenMPRuntime() part.

[gtbercea]

• Add break.

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.h

1445

Cannot add that due to createRuntimeFunction not being const.

[gtbercea]

• Remove const.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

8973

Because the member function is non-static and the member data is non-static, you can just use HasRequiresUnifiedSharedMemory = true;

test/OpenMP/openmp_offload_registration.cpp

38

Add a test where the unified memory flag is used

[gtbercea]

test/OpenMP/openmp_offload_registration.cpp

38

Agreed. More test will be coming anyway. There are a lot of regression tests which need to be updated.

[gtbercea]

• Fix checks for use of requires directive.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.h

644

Why do you need this? I think your function should be called without any conditions. It does not depend on the presence of the target regions or not. Plus, your check is not consistent if you're calling the function from another module, which has target region inside.

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.h

644

This does not determine if the function is called or not. This helps determine the flags with which the function is called.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.h

644

It does not matter, it still does not work correctly if the target region is called in the function from another module

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.h

644

If the target is in another compilation unit, that unit will need to have a requires directive.

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.h

644

If you can come up with an example which you think doesn't work I'm happy to try it.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.h

644

If the module without target directives was compiled with unified memory and the module with target directives compiled without unfied memory support? Is this a problem? Shall we diagnose it?

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.h

644

The requires directives in a module without targets are just not taken into consideration. In general, a target region is needed before the requires flags are checked for compatibility with flags in other modules.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.h

644

You're saying that we're going to ignore the directive completely if the module does not have target regions. I'd suggest to discuss it with Alex if this is ok per standard.

[AlexEichenberger]

fantastic

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.h

644

I discussed with Alex and we agreed on this being the most practical solution. This will enable users to ensure consistency of requires flags across relevant compilation units only.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

8975

The message speaks about target region but points to the clause. You need to correct the message.

9062

Hmm, according to the standard A requires directive with any of the following clauses must appear in all compilation units of a program that contain device constructs or device routines or in none of them. You can not detect use of the device routines in the code, can you?

lib/CodeGen/CGOpenMPRuntime.h

644

Must start from capital letter

[gtbercea]

• Add support for declare target routines.

[gtbercea]

• Update error message.

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.cpp

8975

"A target region was emitted before this requires directive." ?

[ABataev]

Other tests in presence of the requires directive with the clause?

lib/CodeGen/CGOpenMPRuntime.cpp

7956

Restore the original, this must be changed in a separate patch.

8975

Still does not look good, better to add previously emitted target regions in the diagnostic somehow. Also, it would be good to mention the clause in the message.

test/OpenMP/openmp_offload_registration.cpp

40

Why do you need this check?

[gtbercea]

• Update tests.
• Move error check in sema.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

7956

Restore

lib/Sema/SemaOpenMP.cpp

196 ↗

(On Diff #195644)

I would suggest to split this patch into 2 parts: the sema checking part and the codegen part.

468 ↗

(On Diff #195644)

Make function const

469 ↗

(On Diff #195644)

const OMPClause *

476 ↗

(On Diff #195644)

DSAStack usually is not used for diagnostic emission, it serves as the container. Better to move the diagnostics to Act... functions.

[gtbercea]

lib/Sema/SemaOpenMP.cpp

476 ↗

(On Diff #195644)

The "bool hasDuplicateRequiresClause(ArrayRef<OMPClause *> ClauseList) const {" function works in a similar function as this one and does emit diagnostics. See a few lines above.

I need to iterate through a private member of DSAStack so I don't see how that can be in any other place than in DSAStack.

[ABataev]

lib/Sema/SemaOpenMP.cpp

476 ↗

(On Diff #195644)

1. hasDuplicateRequiresClause must be fixed.
2. You can return ArrayRef<SourceLocation> and iterate over it in the Act.. function

[gtbercea]

• Add const.
• Address comments.

[ABataev]

include/clang/Basic/DiagnosticSemaKinds.td

9135 ↗

(On Diff #195739)

Split the patch, sema part must be separate.

[gtbercea]

• Move error check in sema.

[ABataev]

test/OpenMP/openmp_offload_registration.cpp

40

Again, this check is not required

test/OpenMP/target_codegen.cpp

99

Add a check for function too

[gtbercea]

• Add tests.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

9797

Bad formatting.

9799

No need for the braces

9799

What if declare target is used only for variabes but not for the functions?

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.cpp

9799

Even more reason to error in that case since it may contain clauses like link or to which need for requires directives to be used consistently.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

9799

But I don't see that your patch works in this situation. Currently, it will emit the error only if the declare target function is found, no?

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.cpp

9799

Actually it will work even when just variables are used in the declare target region. There is another problem which I have a fix for. I will update the patch in a bit.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

9799

Why will it work in this case? If you have just a declare target region in the code for the variables and nothing else. You don't have target regions, declare target functions etc. It won't work in this case.

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.cpp

9799

It will work because the OffloadEntriesInfoManager.empty() will return false in that case.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

9041

Missed check for HasEmittedTargetRegion here

9062

The first part of the condition is excessive. You have early exit from the function. Instead of the condition, use assert((!OffloadEntriesInfoManager.empty() || HasEmittedDeclareTargetRegion || HasEmittedTargetRegion) && "Expected bla-bla-bla");

9799

But you don't have a check for OffloadEntriesInfoManager.empty() when you set Flags for __tgt_register_requires function call.

[gtbercea]

• Fix conditions.

[gtbercea]

• Fix format.

[gtbercea]

• Fix braces.

[gtbercea]

ping

[ABataev]

Looks good, in general. Just one more question: will tgt_register_requires be emitted if only -fopenmp is used? If -fopenmp-targets is not provided, this function should not be called, I think, because we're not going to use offloading at all in this case and tgt_register_requires should not be emitted. Otherwise, we may end up with the linker error that tgt_register_requires is not found. Would be good to have a test with -fopenmp and #pragma omp requires that check that no tgt_register_requires is emitted in this case.

[gtbercea]

• Avoid requires registration when no targets are specified.

[gtbercea]

In D60568#1508909, @ABataev wrote:

Looks good, in general. Just one more question: will tgt_register_requires be emitted if only -fopenmp is used? If -fopenmp-targets is not provided, this function should not be called, I think, because we're not going to use offloading at all in this case and tgt_register_requires should not be emitted. Otherwise, we may end up with the linker error that tgt_register_requires is not found. Would be good to have a test with -fopenmp and #pragma omp requires that check that no tgt_register_requires is emitted in this case.

Done with test included.

[ABataev]

lib/Frontend/CompilerInvocation.cpp

2840 ↗

(On Diff #200355)

We already have Opts.OMPTargetTripples. We use it currently to check if the offloading enabled !CGM.getLangOpts().OMPTargetTriples.empty(), so, probably, we don't need this new flag.

[gtbercea]

lib/Frontend/CompilerInvocation.cpp

2840 ↗

(On Diff #200355)

Even better!! I'll make the change!

[gtbercea]

• Remove new option.

[ABataev]

LG. Please, commit it only after the runtime part is committed.

lib/CodeGen/CGOpenMPRuntime.h

private:
   /// List of the global variables with their addresses that should not be
   /// emitted for the target.
   llvm::StringMap<llvm::WeakTrackingVH> EmittedNonTargetVariables;
 
   /// List of variables that can become declare target implicitly and, thus,
   /// must be emitted.
   llvm::SmallDenseSet<const VarDecl *> DeferredGlobalVariables;
 
+  /// Flag for keeping track of weather a requires unified_shared_memory
+  /// directive is present.
+  bool HasRequiresUnifiedSharedMemory = false;
+
+  /// Flag for keeping track of weather a target region has been emitted.
+  bool HasEmittedTargetRegion = false;
+
+  /// Flag for keeping track of weather a device routine has been emitted.
+  /// Device routines are specific to the
+  bool HasEmittedDeclareTargetRegion = false;
+
   /// Creates and registers offloading binary descriptor for the current
   /// compilation unit. The function that does the registration is returned.
   llvm::Function *createOffloadingBinaryDescriptorRegistration();
 
   /// Creates all the offload entries in the current compilation unit
   /// along with the associated metadata.
   void createOffloadEntriesAndInfoMetadata();
 
public:
   llvm::Constant *registerTargetFirstprivateCopy(CodeGenFunction &CGF,
                                                  const VarDecl *VD);
 
   /// Emit the global \a GD if it is meaningful for the target. Returns
   /// if it was emitted successfully.
   /// \param GD Global to scan.
   virtual bool emitTargetGlobal(GlobalDecl GD);
 
+  /// Creates and returns a registration function for when at least one
+  /// requires directives was used in the current module.
+  llvm::Function *emitRequiresDirectiveRegFun();
+
   /// 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();
 
   /// Emits code for teams call of the \a OutlinedFn with
   /// variables captured in a record which address is stored in \a
   /// CapturedStruct.
public:
   /// translating these arguments to correct target-specific arguments.
   virtual void
   emitOutlinedFunctionCall(CodeGenFunction &CGF, SourceLocation Loc,
                            llvm::FunctionCallee OutlinedFn,
                            ArrayRef<llvm::Value *> Args = llvm::None) const;
 
   /// Emits OpenMP-specific function prolog.
   /// Required for device constructs.
-  virtual void emitFunctionProlog(CodeGenFunction &CGF, const Decl *D) {}
+  virtual void emitFunctionProlog(CodeGenFunction &CGF, const Decl *D);
 
   /// Gets the OpenMP-specific address of the local variable.
   virtual Address getAddressOfLocalVariable(CodeGenFunction &CGF,
                                             const VarDecl *VD);
 
   /// Marks the declaration as already emitted for the device code and returns
   /// true, if it was marked already, and false, otherwise.
   bool markAsGlobalTarget(GlobalDecl GD);
 
   /// Emit deferred declare target variables marked for deferred emission.
   void emitDeferredTargetDecls() const;
 
   /// Adjust some parameters for the target-based directives, like addresses of
   /// the variables captured by reference in lambdas.
   virtual void
   adjustTargetSpecificDataForLambdas(CodeGenFunction &CGF,
                                      const OMPExecutableDirective &D) const;
 
   /// Perform check on requires decl to ensure that target architecture
   /// supports unified addressing
-  virtual void checkArchForUnifiedAddressing(const OMPRequiresDecl *D) const {}
+  virtual void checkArchForUnifiedAddressing(const OMPRequiresDecl *D);
 
   /// Checks if the variable has associated OMPAllocateDeclAttr attribute with
   /// the predefined allocator and translates it into the corresponding address
   /// space.
   virtual bool hasAllocateAttributeForGlobalVar(const VarDecl *VD, LangAS &AS);
 };
 
 /// Class supports emissionof SIMD-only code.

lib/CodeGen/CGOpenMPRuntime.cpp

enum OpenMPLocationFlags : unsigned {
   OMP_IDENT_WORK_LOOP = 0x200,
   /// Call of __kmp_for_static_init for sections.
   OMP_IDENT_WORK_SECTIONS = 0x400,
   /// Call of __kmp_for_static_init for distribute.
   OMP_IDENT_WORK_DISTRIBUTE = 0x800,
   LLVM_MARK_AS_BITMASK_ENUM(/*LargestValue=*/OMP_IDENT_WORK_DISTRIBUTE)
 };
 
+namespace {
+LLVM_ENABLE_BITMASK_ENUMS_IN_NAMESPACE();
+/// Values for bit flags for marking which requires clauses have been used.
+enum OpenMPOffloadingRequiresDirFlags : int64_t {
+  /// flag undefined.
+  OMP_REQ_UNDEFINED               = 0x000,
+  /// no requires clause present.
+  OMP_REQ_NONE                    = 0x001,
+  /// reverse_offload clause.
+  OMP_REQ_REVERSE_OFFLOAD         = 0x002,
+  /// unified_address clause.
+  OMP_REQ_UNIFIED_ADDRESS         = 0x004,
+  /// unified_shared_memory clause.
+  OMP_REQ_UNIFIED_SHARED_MEMORY   = 0x008,
+  /// dynamic_allocators clause.
+  OMP_REQ_DYNAMIC_ALLOCATORS      = 0x010,
+  LLVM_MARK_AS_BITMASK_ENUM(/*LargestValue=*/OMP_REQ_DYNAMIC_ALLOCATORS)
+};
+} // anonymous namespace
+
 /// Describes ident structure that describes a source location.
 /// All descriptions are taken from
 /// https://github.com/llvm/llvm-project/blob/master/openmp/runtime/src/kmp.h
 /// Original structure:
 /// typedef struct ident {
 ///    kmp_int32 reserved_1;   /**<  might be used in Fortran;
 ///                                  see above  */
 ///    kmp_int32 flags;        /**<  also f.flags; KMP_IDENT_xxx flags;
enum OpenMPRTLFunction {
   // Call to int32_t __tgt_target_teams(int64_t device_id, void *host_ptr,
   // int32_t arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t
   // *arg_types, int32_t num_teams, int32_t thread_limit);
   OMPRTL__tgt_target_teams,
   // Call to int32_t __tgt_target_teams_nowait(int64_t device_id, void
   // *host_ptr, int32_t arg_num, void** args_base, void **args, size_t
   // *arg_sizes, int64_t *arg_types, int32_t num_teams, int32_t thread_limit);
   OMPRTL__tgt_target_teams_nowait,
+  // Call to void __tgt_register_requires(int64_t flags);
+  OMPRTL__tgt_register_requires,
   // 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,
   // Call to void __tgt_target_data_begin(int64_t device_id, int32_t arg_num,
   // void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types);
   OMPRTL__tgt_target_data_begin,
   // Call to void __tgt_target_data_begin_nowait(int64_t device_id, int32_t
enum OpenMPRTLFunction {
   // void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types);
   OMPRTL__tgt_target_data_update,
   // Call to void __tgt_target_data_update_nowait(int64_t device_id, int32_t
   // arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t
   // *arg_types);
   OMPRTL__tgt_target_data_update_nowait,
 };
 
 /// A basic class for pre|post-action for advanced codegen sequence for OpenMP
 /// region.
 class CleanupTy final : public EHScopeStack::Cleanup {
   PrePostActionTy *Action;
 
 public:
   explicit CleanupTy(PrePostActionTy *Action) : Action(Action) {}
   void Emit(CodeGenFunction &CGF, Flags /*flags*/) override {
     if (!CGF.HaveInsertPoint())
       return;
     Action->Exit(CGF);
CGOpenMPRuntime::CGOpenMPRuntime(CodeGenModule &CGM, StringRef FirstSeparator,
   addFieldToRecordDecl(C, RD, C.VoidPtrTy);
   RD->completeDefinition();
   IdentQTy = C.getRecordType(RD);
   IdentTy = CGM.getTypes().ConvertRecordDeclType(RD);
   KmpCriticalNameTy = llvm::ArrayType::get(CGM.Int32Ty, /*NumElements*/ 8);
 
   loadOffloadInfoMetadata();
 }
 
 void CGOpenMPRuntime::clear() {
   InternalVars.clear();
   // Clean non-target variable declarations possibly used only in debug info.
   for (const auto &Data : EmittedNonTargetVariables) {
     if (!Data.getValue().pointsToAliveValue())
       continue;
     auto *GV = dyn_cast<llvm::GlobalVariable>(Data.getValue());
     if (!GV)
llvm::Type *TypeParams[] = {CGM.Int64Ty,
                                 CGM.Int64Ty->getPointerTo(),
                                 CGM.Int32Ty,
                                 CGM.Int32Ty};
     auto *FnTy =
         llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
     RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_teams_nowait");
     break;
   }
+  case OMPRTL__tgt_register_requires: {
+    // Build void __tgt_register_requires(int64_t flags);
+    llvm::Type *TypeParams[] = {CGM.Int64Ty};
+    auto *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_register_requires");
+    break;
+  }
   case OMPRTL__tgt_register_lib: {
     // Build void __tgt_register_lib(__tgt_bin_desc *desc);
     QualType ParamTy =
         CGM.getContext().getPointerType(getTgtBinaryDescriptorQTy());
     llvm::Type *TypeParams[] = {CGM.getTypes().ConvertTypeForMem(ParamTy)};
     auto *FnTy =
         llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
     RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_register_lib");
void CGOpenMPRuntime::OffloadEntriesInfoManagerTy::
     actOnDeviceGlobalVarEntriesInfo(
         const OffloadDeviceGlobalVarEntryInfoActTy &Action) {
   // Scan all target region entries and perform the provided action.
   for (const auto &E : OffloadEntriesDeviceGlobalVar)
     Action(E.getKey(), E.getValue());
 }
 
 llvm::Function *
 CGOpenMPRuntime::createOffloadingBinaryDescriptorRegistration() {
   // If we don't have entries or if we are emitting code for the device, we
   // don't need to do anything.
   if (CGM.getLangOpts().OpenMPIsDevice || OffloadEntriesInfoManager.empty())
     return nullptr;
 
   llvm::Module &M = CGM.getModule();
   ASTContext &C = CGM.getContext();
 
   // Get list of devices we care about
   const std::vector<llvm::Triple> &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::Type *OffloadEntryTy =
       CGM.getTypes().ConvertTypeForMem(getTgtOffloadEntryQTy());
   std::string EntriesBeginName = getName({"omp_offloading", "entries_begin"});
   auto *HostEntriesBegin = new llvm::GlobalVariable(
       M, OffloadEntryTy, /*isConstant=*/true,
       llvm::GlobalValue::ExternalLinkage, /*Initializer=*/nullptr,
       EntriesBeginName);
   std::string EntriesEndName = getName({"omp_offloading", "entries_end"});
void CGOpenMPRuntime::emitCancelCall(CodeGenFunction &CGF, SourceLocation Loc,
   }
 }
 
 void CGOpenMPRuntime::emitTargetOutlinedFunction(
     const OMPExecutableDirective &D, StringRef ParentName,
     llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
     bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
   assert(!ParentName.empty() && "Invalid target region parent name!");
+  HasEmittedTargetRegion = true;
   emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
                                    IsOffloadEntry, CodeGen);
 }
 
 void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
     const OMPExecutableDirective &D, StringRef ParentName,
     llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
     bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
public:
   }
 
   /// Generate the base pointers, section pointers, sizes and map types
   /// associated with the declare target link variables.
   void generateInfoForDeclareTargetLink(MapBaseValuesArrayTy &BasePointers,
                                         MapValuesArrayTy &Pointers,
                                         MapValuesArrayTy &Sizes,
                                         MapFlagsArrayTy &Types) const {
     // Map other list items in the map clause which are not captured variables
-    // but "declare target link" global variables.,
+    // but "declare target link" global variables.
     for (const auto *C : this->CurDir.getClausesOfKind<OMPMapClause>()) {
       for (const auto &L : C->component_lists()) {
         if (!L.first)
           continue;
         const auto *VD = dyn_cast<VarDecl>(L.first);
         if (!VD)
           continue;
         llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =

 }
 
 void CGOpenMPRuntime::adjustTargetSpecificDataForLambdas(
     CodeGenFunction &CGF, const OMPExecutableDirective &D) const {
   assert(isOpenMPTargetExecutionDirective(D.getDirectiveKind()) &&
          " Expected target-based directive.");
 }
 
+void CGOpenMPRuntime::checkArchForUnifiedAddressing(
+    const OMPRequiresDecl *D) {
+  for (const OMPClause *Clause : D->clauselists()) {
+    if (Clause->getClauseKind() == OMPC_unified_shared_memory) {
+      HasRequiresUnifiedSharedMemory = true;
+      break;
+    }
+  }
+}
+
 bool CGOpenMPRuntime::hasAllocateAttributeForGlobalVar(const VarDecl *VD,
                                                        LangAS &AS) {
   if (!VD || !VD->hasAttr<OMPAllocateDeclAttr>())
     return false;
   const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
   switch(A->getAllocatorType()) {
   case OMPAllocateDeclAttr::OMPDefaultMemAlloc:
   // Not supported, fallback to the default mem space.
if (D->hasBody() && AlreadyEmittedTargetFunctions.count(Name) == 0) {
       return false;
     }
     return true;
   }
 
   return !AlreadyEmittedTargetFunctions.insert(Name).second;
 }
 
+llvm::Function *CGOpenMPRuntime::emitRequiresDirectiveRegFun() {
+  // If we don't have entries or if we are emitting code for the device, we
+  // don't need to do anything.
+  if (CGM.getLangOpts().OpenMPSimd || CGM.getLangOpts().OpenMPIsDevice ||
+      (OffloadEntriesInfoManager.empty() && !HasEmittedDeclareTargetRegion))
+    return nullptr;
+
+  // Create and register the function that handles the requires directives.
+  ASTContext &C = CGM.getContext();
+
+  llvm::Function *RequiresRegFn;
+  {
+    CodeGenFunction CGF(CGM);
+    const auto &FI = CGM.getTypes().arrangeNullaryFunction();
+    llvm::FunctionType *FTy = CGM.getTypes().GetFunctionType(FI);
+    std::string ReqName = getName({"omp_offloading", "requires_reg"});
+    RequiresRegFn = CGM.CreateGlobalInitOrDestructFunction(FTy, ReqName, FI);
+    CGF.StartFunction(GlobalDecl(), C.VoidTy, RequiresRegFn, FI, {});
+    OpenMPOffloadingRequiresDirFlags Flags = OMP_REQ_NONE;
+    // TODO: check for other requires clauses.
+    // The requires directive takes effect only when a target region is
+    // present in the compilation unit. Otherwise it is ignored and not
+    // passed to the runtime. This avoids the runtime from throwing an error
+    // for mismatching requires clauses across compilation units that don't
+    // contain at least 1 target region.
+    if ((HasEmittedTargetRegion || HasEmittedDeclareTargetRegion) &&
+        HasRequiresUnifiedSharedMemory)
+      Flags = OMP_REQ_UNIFIED_SHARED_MEMORY;
+    CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_register_requires),
+        llvm::ConstantInt::get(CGM.Int64Ty, Flags));
+    CGF.FinishFunction();
+  }
+  return RequiresRegFn;
+}
+
 llvm::Function *CGOpenMPRuntime::emitRegistrationFunction() {
   // If we have offloading in the current module, we need to emit the entries
   // now and register the offloading descriptor.
   createOffloadEntriesAndInfoMetadata();
 
   // Create and register the offloading binary descriptors. This is the main
   // entity that captures all the information about offloading in the current
   // compilation unit.

 }
 
 void CGOpenMPRuntime::emitOutlinedFunctionCall(
     CodeGenFunction &CGF, SourceLocation Loc, llvm::FunctionCallee OutlinedFn,
     ArrayRef<llvm::Value *> Args) const {
   emitCall(CGF, Loc, OutlinedFn, Args);
 }
 
+void CGOpenMPRuntime::emitFunctionProlog(CodeGenFunction &CGF, const Decl *D){
+  if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
+    if (OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(FD)) {
+      HasEmittedDeclareTargetRegion = true;
+    }
+  }
+}
+
 Address CGOpenMPRuntime::getParameterAddress(CodeGenFunction &CGF,
                                              const VarDecl *NativeParam,
                                              const VarDecl *TargetParam) const {
   return CGF.GetAddrOfLocalVar(NativeParam);
 }
 
 namespace {
 /// Cleanup action for allocate support.

lib/CodeGen/CGOpenMPRuntimeNVPTX.h

public:
 
   /// Adjust some parameters for the target-based directives, like addresses of
   /// the variables captured by reference in lambdas.
   void adjustTargetSpecificDataForLambdas(
       CodeGenFunction &CGF, const OMPExecutableDirective &D) const override;
 
   /// Perform check on requires decl to ensure that target architecture
   /// supports unified addressing
-  void checkArchForUnifiedAddressing(const OMPRequiresDecl *D) const override;
+  void checkArchForUnifiedAddressing(const OMPRequiresDecl *D) override;
 
   /// Returns default address space for the constant firstprivates, __constant__
   /// address space by default.
   unsigned getDefaultFirstprivateAddressSpace() const override;
 
   /// Checks if the variable has associated OMPAllocateDeclAttr attribute with
   /// the predefined allocator and translates it into the corresponding address
   /// space.

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

for (const auto &Feature : Features) {
     }
   }
   return CudaArch::UNKNOWN;
 }
 
 /// Check to see if target architecture supports unified addressing which is
 /// a restriction for OpenMP requires clause "unified_shared_memory".
 void CGOpenMPRuntimeNVPTX::checkArchForUnifiedAddressing(
-    const OMPRequiresDecl *D) const {
+    const OMPRequiresDecl *D) {
   for (const OMPClause *Clause : D->clauselists()) {
     if (Clause->getClauseKind() == OMPC_unified_shared_memory) {
       switch (getCudaArch(CGM)) {
       case CudaArch::SM_20:
       case CudaArch::SM_21:
       case CudaArch::SM_30:
       case CudaArch::SM_32:
       case CudaArch::SM_35:
       case CudaArch::SM_37:
if (Clause->getClauseKind() == OMPC_unified_shared_memory) {
       case CudaArch::GFX909:
       case CudaArch::UNKNOWN:
         break;
       case CudaArch::LAST:
         llvm_unreachable("Unexpected Cuda arch.");
       }
     }
   }
+  CGOpenMPRuntime::checkArchForUnifiedAddressing(D);
 }
 
 /// Get number of SMs and number of blocks per SM.
 static std::pair<unsigned, unsigned> getSMsBlocksPerSM(CodeGenModule &CGM) {
   std::pair<unsigned, unsigned> Data;
   if (CGM.getLangOpts().OpenMPCUDANumSMs)
     Data.first = CGM.getLangOpts().OpenMPCUDANumSMs;
   if (CGM.getLangOpts().OpenMPCUDABlocksPerSM)

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 (OpenMPRuntime) {
+    if (llvm::Function *OpenMPRequiresDirectiveRegFun =
+            OpenMPRuntime->emitRequiresDirectiveRegFun()) {
+      AddGlobalCtor(OpenMPRequiresDirectiveRegFun, 0);
+    }
     if (llvm::Function *OpenMPRegistrationFunction =
             OpenMPRuntime->emitRegistrationFunction()) {
       auto ComdatKey = OpenMPRegistrationFunction->hasComdat() ?
         OpenMPRegistrationFunction : nullptr;
       AddGlobalCtor(OpenMPRegistrationFunction, 0, ComdatKey);
     }
     OpenMPRuntime->clear();
   }

test/OpenMP/openmp_offload_registration.cpp

 // CHECK: [[DEV1BEGIN:@.+]] = extern_weak constant i8
 // CHECK: [[DEV1END:@.+]] = extern_weak constant i8
 // CHECK: [[DEV2BEGIN:@.+]] = extern_weak constant i8
 // CHECK: [[DEV2END:@.+]] = extern_weak constant i8
 // CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [2 x [[DEVTY]]] [{{.+}} { i8* [[DEV1BEGIN]], i8* [[DEV1END]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }, {{.+}} { i8* [[DEV2BEGIN]], i8* [[DEV2END]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }], comdat($[[REGFN]])
 // CHECK: [[DESC:@.+]] = internal constant [[DSCTY]] { i32 2, [[DEVTY]]* getelementptr inbounds ([2 x [[DEVTY]]], [2 x [[DEVTY]]]* [[IMAGES]], i32 0, i32 0), [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }, comdat($[[REGFN]])
 
 // Check target registration is registered as a Ctor.
-// CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* @[[REGFN]], i8* bitcast (void ()* @[[REGFN]] to i8*) }]
+// CHECK: appending global [2 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* @.omp_offloading.requires_reg, i8* null }, { i32, void ()*, i8* } { i32 0, void ()* @[[REGFN]], i8* bitcast (void ()* @[[REGFN]] to i8*) }]
 
 // Check presence of foo() and the outlined target region
 // CHECK: define void [[FOO:@.+]]()
 // CHECK: define internal void [[OUTLINEDTARGET:@.+]]() 
 
 // Check registration and unregistration code.
 
+// CHECK:     define internal void @.omp_offloading.requires_reg()
+// CHECK:     call void @__tgt_register_requires(i64 1)
+// CHECK:     ret void
+// CHECK:     declare void @__tgt_register_requires(i64)
+
 // CHECK:     define internal void @[[UNREGFN:.+]](i8*)
 // CHECK-SAME: comdat($[[REGFN]]) {
 // CHECK:     call i32 @__tgt_unregister_lib([[DSCTY]]* [[DESC]])
 // CHECK:     ret void
 // CHECK:     declare i32 @__tgt_unregister_lib([[DSCTY]]*)
 
 // CHECK:     define linkonce hidden void @[[REGFN]]()
 // CHECK-SAME: comdat {
 // 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]]*)
 

test/OpenMP/target_codegen.cpp

 // CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]]
 // CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]]
 // CHECK: [[DEVBEGIN:@.+]] = extern_weak constant i8
 // CHECK: [[DEVEND:@.+]] = extern_weak constant i8
 // CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [1 x [[DEVTY]]] [{{.+}} { i8* [[DEVBEGIN]], i8* [[DEVEND]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }], comdat($[[REGFN]])
 // 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]] }, comdat($[[REGFN]])
 
 // Check target registration is registered as a Ctor.
-// CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* @[[REGFN]], i8* bitcast (void ()* @[[REGFN]] to i8*) }]
+// CHECK: appending global [2 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* @.omp_offloading.requires_reg, i8* null }, { i32, void ()*, i8* } { i32 0, void ()* @[[REGFN]], i8* bitcast (void ()* @[[REGFN]] to i8*) }]
 
 
 template<typename tx, typename ty>
 struct TT{
   tx X;
   ty Y;
 };
 

test/OpenMP/target_codegen_registration.cpp

 // CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]]
 // CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]]
 // CHECK: [[DEVBEGIN:@.+]] = extern_weak constant i8
 // CHECK: [[DEVEND:@.+]] = extern_weak constant i8
 // CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [1 x [[DEVTY]]] [{{.+}} { i8* [[DEVBEGIN]], i8* [[DEVEND]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }], comdat($[[REGFN]])
 // 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]] }, comdat($[[REGFN]])
 
 // 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: @llvm.global_ctors = appending global [5 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 ()* @.omp_offloading.requires_reg, i8* null },
 // CHECK-SAME: { i32, void ()*, i8* } { i32 0, void ()* @[[REGFN]], i8* bitcast (void ()* @[[REGFN]] to i8*) }]
 
 // CHECK-NTARGET: @llvm.global_ctors = appending global [3   x { i32, void ()*, i8* }] [
 
 extern int *R;
 
 struct SA {
   int arr[4];

 //CHECK:     call void @{{.+}}()
 //CHECK:     call void @{{.+}}()
 //CHECK:     call void @{{.+}}()
 //CHECK-NOT: call void @{{.+}}()
 //CHECK:     ret void
 
 // Check registration and unregistration
 
+//CHECK:     define internal void @.omp_offloading.requires_reg()
+//CHECK:     call void @__tgt_register_requires(i64 1)
+//CHECK:     ret void
+
 //CHECK:     define internal void @[[UNREGFN:.+]](i8*)
 //CHECK-SAME: comdat($[[REGFN]]) {
 //CHECK:     call i32 @__tgt_unregister_lib([[DSCTY]]* [[DESC]])
 //CHECK:     ret void
 //CHECK:     declare i32 @__tgt_unregister_lib([[DSCTY]]*)
 
 //CHECK:     define linkonce hidden void @[[REGFN]]()
 //CHECK-SAME: comdat {
int bar(int a){
   #pragma omp target
   ++r;
 
   return r + *R;
 }
 
 // Check metadata is properly generated:
 // CHECK:     !omp_offload.info = !{!{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}}
-// CHECK-DAG: = !{i32 0, i32 [[DEVID:-?[0-9]+]], i32 [[FILEID:-?[0-9]+]], !"_ZN2SB3fooEv", i32 216, i32 {{[0-9]+}}}
-// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SDD1Ev", i32 266, i32 {{[0-9]+}}}
-// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SEC1Ev", i32 282, i32 {{[0-9]+}}}
-// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SED1Ev", i32 288, i32 {{[0-9]+}}}
-// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EE3fooEv", i32 299, i32 {{[0-9]+}}}
-// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EEC1Ev", i32 305, i32 {{[0-9]+}}}
-// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_Z3bari", i32 427, i32 {{[0-9]+}}}
-// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EED1Ev", i32 311, i32 {{[0-9]+}}}
-// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EEC1Ev", i32 305, i32 {{[0-9]+}}}
-// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EED1Ev", i32 311, i32 {{[0-9]+}}}
-// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EE3fooEv", i32 299, i32 {{[0-9]+}}}
-// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SCC1Ev", i32 241, i32 {{[0-9]+}}}
+// CHECK-DAG: = !{i32 0, i32 [[DEVID:-?[0-9]+]], i32 [[FILEID:-?[0-9]+]], !"_ZN2SB3fooEv", i32 217, i32 {{[0-9]+}}}
+// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SDD1Ev", i32 267, i32 {{[0-9]+}}}
+// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SEC1Ev", i32 283, i32 {{[0-9]+}}}
+// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SED1Ev", i32 289, i32 {{[0-9]+}}}
+// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EE3fooEv", i32 300, i32 {{[0-9]+}}}
+// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EEC1Ev", i32 306, i32 {{[0-9]+}}}
+// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_Z3bari", i32 432, i32 {{[0-9]+}}}
+// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EED1Ev", i32 312, i32 {{[0-9]+}}}
+// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EEC1Ev", i32 306, i32 {{[0-9]+}}}
+// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EED1Ev", i32 312, i32 {{[0-9]+}}}
+// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EE3fooEv", i32 300, i32 {{[0-9]+}}}
+// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SCC1Ev", i32 242, i32 {{[0-9]+}}}
 
 // TCHECK:     !omp_offload.info = !{!{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}}
-// TCHECK-DAG: = !{i32 0, i32 [[DEVID:-?[0-9]+]], i32 [[FILEID:-?[0-9]+]], !"_ZN2SB3fooEv", i32 216, i32 {{[0-9]+}}}
-// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SDD1Ev", i32 266, i32 {{[0-9]+}}}
-// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SEC1Ev", i32 282, i32 {{[0-9]+}}}
-// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SED1Ev", i32 288, i32 {{[0-9]+}}}
-// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EE3fooEv", i32 299, i32 {{[0-9]+}}}
-// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EEC1Ev", i32 305, i32 {{[0-9]+}}}
-// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_Z3bari", i32 427, i32 {{[0-9]+}}}
-// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EED1Ev", i32 311, i32 {{[0-9]+}}}
-// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EEC1Ev", i32 305, i32 {{[0-9]+}}}
-// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EED1Ev", i32 311, i32 {{[0-9]+}}}
-// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EE3fooEv", i32 299, i32 {{[0-9]+}}}
-// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SCC1Ev", i32 241, i32 {{[0-9]+}}}
+// TCHECK-DAG: = !{i32 0, i32 [[DEVID:-?[0-9]+]], i32 [[FILEID:-?[0-9]+]], !"_ZN2SB3fooEv", i32 217, i32 {{[0-9]+}}}
+// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SDD1Ev", i32 267, i32 {{[0-9]+}}}
+// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SEC1Ev", i32 283, i32 {{[0-9]+}}}
+// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SED1Ev", i32 289, i32 {{[0-9]+}}}
+// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EE3fooEv", i32 300, i32 {{[0-9]+}}}
+// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EEC1Ev", i32 306, i32 {{[0-9]+}}}
+// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_Z3bari", i32 432, i32 {{[0-9]+}}}
+// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EED1Ev", i32 312, i32 {{[0-9]+}}}
+// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EEC1Ev", i32 306, i32 {{[0-9]+}}}
+// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EED1Ev", i32 312, i32 {{[0-9]+}}}
+// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EE3fooEv", i32 300, i32 {{[0-9]+}}}
+// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SCC1Ev", i32 242, i32 {{[0-9]+}}}
 
 #endif