[OpenMP] Fix declare target link implementation

[gtbercea]

Summary

This patch fixes the case where variables in different compilation units or the same compilation unit are under the declare target link clause AND have the same name.
This also fixes the name clash error that occurs when unified memory is activated.
The changes in this patch include:

• Pointers to internal variables are given unique names.
• Externally visible variables are given the same name as before.
• All pointer variables (external or internal) are weakly linked.

[lebedev.ri]

Description lacking.
Why is that bad?
How does this play wrt reproducibility of the output?
Normally value names in IR are completely discarded in Release build mode, why do they matter here?

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

2567 ↗

(On Diff #209319)

1. Capitalize the first letter in tbe variable name.
2. Why do you need to use hash? Could you use variable sourcelocation instead?

Also, what if the variable is defined in one file but declared in another one file. Will it be linked correctly?

lib/Sema/SemaOpenMP.cpp

2625 ↗

(On Diff #209319)

I assume the changes in this file are from the different fix.

[gtbercea]

• Fix tests.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

2610 ↗

(On Diff #210386)

Better to fix the link clause processing in a different patch, it has nothing to do with the unified memory.

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.cpp

2610 ↗

(On Diff #210386)

Sure I can split this patch.

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.cpp

2610 ↗

(On Diff #210386)

It turns out that the patch cannot be split. Splitting the patch would lead to an intermediate state which produced incorrect results for certain corner cases.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

2589 ↗

(On Diff #210386)

Currently, I think, it would better to use isExternallyVisible() member function of VD variable.

2612–2614 ↗

(On Diff #210386)

Remove braces

2615 ↗

(On Diff #210386)

Remove this, it is not required anymore since the pointers have weakany linkage

test/OpenMP/declare_target_codegen.cpp

23 ↗

(On Diff #210386)

I would not rely on the fact that the FileId is always 0

[gtbercea]

• Address comments.

[ABataev]

Fix the name of the patch. It is the patch for declare target, not target.

lib/CodeGen/CGOpenMPRuntime.cpp

2598–2599 ↗

(On Diff #210645)

Why do you want to add to zeroed FileId for externally visible vars? Leave them as-is.

[gtbercea]

• Fix tests. Address comments.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

2594–2597 ↗

(On Diff #210650)

Merge these 2 checks into one, please.

OS << CGM.getMangledName(GlobalDecl(VD));
if (!VD->isExternallyVisible()) {
      unsigned DeviceID, Line;
      unsigned FileID = 0;
        getTargetEntryUniqueInfo(CGM.getContext(),
                                 VD->getCanonicalDecl()->getBeginLoc(),
                                 DeviceID, FileID, Line);
       OS << llvm::format("_%x", FileID);
}
OS << "_decl_tgt_ref_ptr";

[gtbercea]

• Merge conditionals.

[ABataev]

I don't see a single test for static variables. Do we have at least one? If not, add it.

[gtbercea]

• Add static variable.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

2592–2593 ↗

(On Diff #210673)

Move those 3 vars to if statement, they are not needed here.

test/OpenMP/declare_target_link_codegen.cpp

22–23 ↗

(On Diff #210673)

Better to define the variable with the name of the original d variable for better and stable check

[gtbercea]

test/OpenMP/declare_target_link_codegen.cpp

22–23 ↗

(On Diff #210673)

Can you rephrase your comment please? I don't understand what you'd like me to do here.

Do you want me to use [[...]] to define the variable?

[gtbercea]

• Address comments.

[ABataev]

LG with a nit.

lib/CodeGen/CGOpenMPRuntime.cpp

2595 ↗

(On Diff #210854)

No need to initialize FileID.

[gtbercea]

• Address comments.

cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp

llvm::Type *TypeParams[] = {
     PtrTy,                                     // p_upper
     PtrTy                                      // p_stride
   };
   auto *FnTy =
       llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
   return CGM.CreateRuntimeFunction(FnTy, Name);
 }
 
+/// Obtain information that uniquely identifies a target entry. This
+/// consists of the file and device IDs as well as line number associated with
+/// the relevant entry source location.
+static void getTargetEntryUniqueInfo(ASTContext &C, SourceLocation Loc,
+                                     unsigned &DeviceID, unsigned &FileID,
+                                     unsigned &LineNum) {
+  SourceManager &SM = C.getSourceManager();
+
+  // 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.");
+
+  PresumedLoc PLoc = SM.getPresumedLoc(Loc);
+  assert(PLoc.isValid() && "Source location is expected to be always valid.");
+
+  llvm::sys::fs::UniqueID ID;
+  if (auto EC = llvm::sys::fs::getUniqueID(PLoc.getFilename(), ID))
+    SM.getDiagnostics().Report(diag::err_cannot_open_file)
+        << PLoc.getFilename() << EC.message();
+
+  DeviceID = ID.getDevice();
+  FileID = ID.getFile();
+  LineNum = PLoc.getLine();
+}
+
 Address CGOpenMPRuntime::getAddrOfDeclareTargetVar(const VarDecl *VD) {
   if (CGM.getLangOpts().OpenMPSimd)
     return Address::invalid();
   llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
       OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
   if (Res && (*Res == OMPDeclareTargetDeclAttr::MT_Link ||
               (*Res == OMPDeclareTargetDeclAttr::MT_To &&
                HasRequiresUnifiedSharedMemory))) {
     SmallString<64> PtrName;
     {
       llvm::raw_svector_ostream OS(PtrName);
-      OS << CGM.getMangledName(GlobalDecl(VD)) << "_decl_tgt_ref_ptr";
+      OS << CGM.getMangledName(GlobalDecl(VD));
+      if (!VD->isExternallyVisible()) {
+        unsigned DeviceID, FileID, Line;
+        getTargetEntryUniqueInfo(CGM.getContext(),
+                                 VD->getCanonicalDecl()->getBeginLoc(),
+                                 DeviceID, FileID, Line);
+        OS << llvm::format("_%x", FileID);
+      }
+      OS << "_decl_tgt_ref_ptr";
     }
     llvm::Value *Ptr = CGM.getModule().getNamedValue(PtrName);
     if (!Ptr) {
       QualType PtrTy = CGM.getContext().getPointerType(VD->getType());
       Ptr = getOrCreateInternalVariable(CGM.getTypes().ConvertTypeForMem(PtrTy),
                                         PtrName);
-      if (!CGM.getLangOpts().OpenMPIsDevice) {
-        auto *GV = cast<llvm::GlobalVariable>(Ptr);
-        GV->setLinkage(llvm::GlobalValue::ExternalLinkage);
+
+      GV->setLinkage(llvm::GlobalValue::WeakAnyLinkage);
+
+      if (!CGM.getLangOpts().OpenMPIsDevice)
         GV->setInitializer(CGM.GetAddrOfGlobal(VD));
-      }
-      CGM.addUsedGlobal(cast<llvm::GlobalValue>(Ptr));
       registerTargetGlobalVariable(VD, cast<llvm::Constant>(Ptr));
     }
     return Address(Ptr, CGM.getContext().getDeclAlign(VD));
   }
   return Address::invalid();
 }
 
 llvm::Constant *
if (!CGF) {
       InitCGF.FinishFunction();
       return InitFunction;
     }
     emitThreadPrivateVarInit(*CGF, VDAddr, Ctor, CopyCtor, Dtor, Loc);
   }
   return nullptr;
 }
 
-/// Obtain information that uniquely identifies a target entry. This
-/// consists of the file and device IDs as well as line number associated with
-/// the relevant entry source location.
-static void getTargetEntryUniqueInfo(ASTContext &C, SourceLocation Loc,
-                                     unsigned &DeviceID, unsigned &FileID,
-                                     unsigned &LineNum) {
-  SourceManager &SM = C.getSourceManager();
-
-  // 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.");
-
-  PresumedLoc PLoc = SM.getPresumedLoc(Loc);
-  assert(PLoc.isValid() && "Source location is expected to be always valid.");
-
-  llvm::sys::fs::UniqueID ID;
-  if (auto EC = llvm::sys::fs::getUniqueID(PLoc.getFilename(), ID))
-    SM.getDiagnostics().Report(diag::err_cannot_open_file)
-        << PLoc.getFilename() << EC.message();
-
-  DeviceID = ID.getDevice();
-  FileID = ID.getFile();
-  LineNum = PLoc.getLine();
-}
-
 bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD,
                                                      llvm::GlobalVariable *Addr,
                                                      bool PerformInit) {
   Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
       OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
   if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link ||
       (*Res == OMPDeclareTargetDeclAttr::MT_To &&
        HasRequiresUnifiedSharedMemory))

cfe/trunk/test/OpenMP/declare_target_codegen.cpp

 // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}}
 // CHECK-DAG: Bake
 // CHECK-NOT: @{{hhh|ggg|fff|eee}} =
 // CHECK-DAG: @aaa = external global i32,
 // CHECK-DAG: @bbb ={{ dso_local | }}global i32 0,
 // CHECK-DAG: weak constant %struct.__tgt_offload_entry { i8* bitcast (i32* @bbb to i8*),
 // CHECK-DAG: @ccc = external global i32,
 // CHECK-DAG: @ddd ={{ dso_local | }}global i32 0,
-// CHECK-DAG: @hhh_decl_tgt_ref_ptr = common global i32* null
-// CHECK-DAG: @ggg_decl_tgt_ref_ptr = common global i32* null
-// CHECK-DAG: @fff_decl_tgt_ref_ptr = common global i32* null
-// CHECK-DAG: @eee_decl_tgt_ref_ptr = common global i32* null
+// CHECK-DAG: @hhh_decl_tgt_ref_ptr = weak global i32* null
+// CHECK-DAG: @ggg_decl_tgt_ref_ptr = weak global i32* null
+// CHECK-DAG: @fff_decl_tgt_ref_ptr = weak global i32* null
+// CHECK-DAG: @eee_decl_tgt_ref_ptr = weak global i32* null
 // CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23,
 // CHECK-DAG: @b ={{ dso_local | }}global i32 15,
 // CHECK-DAG: @d ={{ dso_local | }}global i32 0,
 // CHECK-DAG: @c = external global i32,
 // CHECK-DAG: @globals ={{ dso_local | }}global %struct.S zeroinitializer,
 // CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer,
 // CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]]
 // CHECK-DAG: @out_decl_target ={{ dso_local | }}global i32 0,
-// CHECK-DAG: @llvm.used = appending global [6 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+80]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+81]]_ctor to i8*),
+// CHECK-DAG: @llvm.used = appending global [2 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+80]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+81]]_ctor to i8*)],
 // CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)],
 
 // CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA|f_method}}{{.*}}()
 // CHECK-DAG: define {{.*}}void @{{.*}}TemplateClass{{.*}}(%class.TemplateClass* %{{.*}})
 // CHECK-DAG: define {{.*}}i32 @{{.*}}TemplateClass{{.*}}f_method{{.*}}(%class.TemplateClass* %{{.*}})
 // CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+74]]_ctor()
 
 #ifndef HEADER

cfe/trunk/test/OpenMP/declare_target_link_codegen.cpp

 // expected-no-diagnostics
 
 // SIMD-ONLY-NOT: {{__kmpc|__tgt}}
 
 #ifndef HEADER
 #define HEADER
 
 // HOST-DAG: @c = external global i32,
-// HOST-DAG: @c_decl_tgt_ref_ptr = global i32* @c
+// HOST-DAG: @c_decl_tgt_ref_ptr = weak global i32* @c
+// HOST-DAG: @[[D:.+]] = internal global i32 2
+// HOST-DAG: @[[D_PTR:.+]] = weak global i32* @[[D]]
 // DEVICE-NOT: @c =
-// DEVICE: @c_decl_tgt_ref_ptr = common global i32* null
-// HOST: [[SIZES:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 4]
-// HOST: [[MAPTYPES:@.+]] = private unnamed_addr constant [2 x i64] [i64 35, i64 531]
+// DEVICE: @c_decl_tgt_ref_ptr = weak global i32* null
+// HOST: [[SIZES:@.+]] = private unnamed_addr constant [3 x i64] [i64 4, i64 4, i64 4]
+// HOST: [[MAPTYPES:@.+]] = private unnamed_addr constant [3 x i64] [i64 35, i64 531, i64 531]
 // HOST: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"c_decl_tgt_ref_ptr\00"
 // HOST: @.omp_offloading.entry.c_decl_tgt_ref_ptr = weak constant %struct.__tgt_offload_entry { i8* bitcast (i32** @c_decl_tgt_ref_ptr to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name, i32 0, i32 0), i64 8, i32 1, i32 0 }, section ".omp_offloading.entries", align 1
-// DEVICE-NOT: internal unnamed_addr constant [{{[0-9]+}} x i8] c"c_decl_tgt_ref_ptr\00"
-// CHECK: @llvm.used = appending global [1 x i8*] [i8* bitcast (i32** @c_decl_tgt_ref_ptr to i8*)]
+// DEVICE-NOT: internal unnamed_addr constant [{{[0-9]+}} x i8] c"c_{{.*}}_decl_tgt_ref_ptr\00"
+// HOST: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"_{{.*}}d_{{.*}}_decl_tgt_ref_ptr\00"
+// HOST: @.omp_offloading.entry.[[D_PTR]] = weak constant %struct.__tgt_offload_entry { i8* bitcast (i32** @[[D_PTR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0
 
 extern int c;
 #pragma omp declare target link(c)
 
+static int d = 2;
+#pragma omp declare target link(d)
+
 int maini1() {
   int a;
 #pragma omp target map(tofrom : a)
   {
     a = c;
+    d++;
   }
 #pragma omp target
 #pragma omp teams
   c = a;
   return 0;
 }
 
-// DEVICE: define weak void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-10]](i32* dereferenceable{{[^,]*}}
+// DEVICE: define weak void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* dereferenceable{{[^,]*}}
 // DEVICE: [[C_REF:%.+]] = load i32*, i32** @c_decl_tgt_ref_ptr,
 // DEVICE: [[C:%.+]] = load i32, i32* [[C_REF]],
 // DEVICE: store i32 [[C]], i32* %
 
 // HOST: define {{.*}}i32 @{{.*}}maini1{{.*}}()
-// HOST: [[BASEPTRS:%.+]] = alloca [2 x i8*],
-// HOST: [[PTRS:%.+]] = alloca [2 x i8*],
-// HOST: getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
-// HOST: getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
-// HOST: [[BP1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+// HOST: [[BASEPTRS:%.+]] = alloca [3 x i8*],
+// HOST: [[PTRS:%.+]] = alloca [3 x i8*],
+// HOST: getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+// HOST: getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+
+// HOST: [[BP1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
 // HOST: [[BP1_CAST:%.+]] = bitcast i8** [[BP1]] to i32***
 // HOST: store i32** @c_decl_tgt_ref_ptr, i32*** [[BP1_CAST]],
-// HOST: [[P1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+// HOST: [[P1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
 // HOST: [[P1_CAST:%.+]] = bitcast i8** [[P1]] to i32**
 // HOST: store i32* @c, i32** [[P1_CAST]],
-// HOST: [[BP0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
-// HOST: [[P0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
-// HOST: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 2, i8** [[BP0]], i8** [[P0]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[SIZES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPTYPES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0))
-// HOST: call void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-29]](i32* %{{[^,]+}})
-// HOST: call i32 @__tgt_target_teams(i64 -1, i8* @.__omp_offloading_{{.+}}_l40.region_id, i32 2, {{.+}})
 
-// HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-32]](i32* dereferenceable{{.*}})
+// HOST: [[BP2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// HOST: [[BP2_CAST:%.+]] = bitcast i8** [[BP2]] to i32***
+// HOST: store i32** @[[D_PTR]], i32*** [[BP2_CAST]],
+// HOST: [[P2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// HOST: [[P2_CAST:%.+]] = bitcast i8** [[P2]] to i32**
+// HOST: store i32* @[[D]], i32** [[P2_CAST]],
+
+// HOST: [[BP0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+// HOST: [[P0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+// HOST: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP0]], i8** [[P0]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[SIZES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPTYPES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0))
+// HOST: call void @__omp_offloading_{{.*}}_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* %{{[^,]+}})
+// HOST: call i32 @__tgt_target_teams(i64 -1, i8* @.__omp_offloading_{{.+}}_l47.region_id, i32 2, {{.+}})
+
+// HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* dereferenceable{{.*}})
 // HOST: [[C:%.*]] = load i32, i32* @c,
 // HOST: store i32 [[C]], i32* %
 
 // CHECK: !{i32 1, !"c_decl_tgt_ref_ptr", i32 1, i32 {{[0-9]+}}}
 #endif // HEADER

cfe/trunk/test/OpenMP/nvptx_target_requires_unified_shared_memory.cpp

#pragma omp target
   for(int i = 0; i < n; i++) {
     sum += var + to_var;
   }
 
   return sum;
 }
 
 // CHECK-HOST: [[VAR:@.+]] = global double 1.000000e+01
-// CHECK-HOST: [[VAR_DECL_TGT_LINK_PTR:@.+]] = global double* [[VAR]]
+// CHECK-HOST: [[VAR_DECL_TGT_LINK_PTR:@.+]] = weak global double* [[VAR]]
 
 // CHECK-HOST: [[TO_VAR:@.+]] = global double 2.000000e+01
-// CHECK-HOST: [[VAR_DECL_TGT_TO_PTR:@.+]] = global double* [[TO_VAR]]
+// CHECK-HOST: [[VAR_DECL_TGT_TO_PTR:@.+]] = weak global double* [[TO_VAR]]
 
 // CHECK-HOST: [[OFFLOAD_SIZES:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 8]
 // CHECK-HOST: [[OFFLOAD_MAPTYPES:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800]
 
 // CHECK-HOST: [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR_NAME:@.+]] = internal unnamed_addr constant [21 x i8]
 // CHECK-HOST: [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR:@.+]] = weak constant %struct.__tgt_offload_entry { i8* bitcast (double** [[VAR_DECL_TGT_LINK_PTR]] to i8*), i8* getelementptr inbounds ([21 x i8], [21 x i8]* [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR_NAME]], i32 0, i32 0), i64 8, i32 1, i32 0 }, section ".omp_offloading.entries"
 
 // CHECK-HOST: [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR_NAME:@.+]] = internal unnamed_addr constant [24 x i8]
 // CHECK-HOST: [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR:@.+]] = weak constant %struct.__tgt_offload_entry { i8* bitcast (double** [[VAR_DECL_TGT_TO_PTR]] to i8*), i8* getelementptr inbounds ([24 x i8], [24 x i8]* [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR_NAME]], i32 0, i32 0), i64 8, i32 0, i32 0 }, section ".omp_offloading.entries"
 
-// CHECK-HOST: @llvm.used = appending global [2 x i8*] [i8* bitcast (double** [[VAR_DECL_TGT_LINK_PTR]] to i8*), i8* bitcast (double** [[VAR_DECL_TGT_TO_PTR]] to i8*)], section "llvm.metadata"
-
 // CHECK-HOST: [[N_CASTED:%.+]] = alloca i64
 // CHECK-HOST: [[SUM_CASTED:%.+]] = alloca i64
 
 // CHECK-HOST: [[OFFLOAD_BASEPTRS:%.+]] = alloca [2 x i8*]
 // CHECK-HOST: [[OFFLOAD_PTRS:%.+]] = alloca [2 x i8*]
 
 // CHECK-HOST: [[LOAD1:%.+]] = load i64, i64* [[N_CASTED]]
 // CHECK-HOST: [[LOAD2:%.+]] = load i64, i64* [[SUM_CASTED]]

 // CHECK-HOST: [[BCAST4:%.+]] = bitcast i8** [[BPTR4]] to i64*
 // CHECK-HOST: store i64 [[LOAD2]], i64* [[BCAST4]]
 
 // CHECK-HOST: [[BPTR7:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 0
 // CHECK-HOST: [[BPTR8:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0
 
 // CHECK-HOST: call i32 @__tgt_target(i64 -1, i8* @{{.*}}.region_id, i32 2, i8** [[BPTR7]], i8** [[BPTR8]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[OFFLOAD_SIZES]], i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[OFFLOAD_MAPTYPES]], i32 0, i32 0))
 
-// CHECK-DEVICE: [[VAR_LINK:@.+]] = common global double* null
-// CHECK-DEVICE: [[VAR_TO:@.+]] = common global double* null
-
-// CHECK-DEVICE: @llvm.used = appending global [2 x i8*] [i8* bitcast (double** [[VAR_LINK]] to i8*), i8* bitcast (double** [[VAR_TO]] to i8*)], section "llvm.metadata"
+// CHECK-DEVICE: [[VAR_LINK:@.+]] = weak global double* null
+// CHECK-DEVICE: [[VAR_TO:@.+]] = weak global double* null
 
 // CHECK-DEVICE: [[VAR_TO_PTR:%.+]] = load double*, double** [[VAR_TO]]
 // CHECK-DEVICE: load double, double* [[VAR_TO_PTR]]
 
 #endif