[OpenMP] Fx link clause

[gtbercea]

Summary

Adjust linkage of variables under the link clause.

[ABataev]

The description does not explain anything.

lib/CodeGen/CGOpenMPRuntime.cpp

2578

Remove this line of code, it is not required

[gtbercea]

Not needed as a separate patch.

lib/CodeGen/CGOpenMPRuntime.cpp

if (Res && (*Res == OMPDeclareTargetDeclAttr::MT_Link ||
     }
     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);
         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 *

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
 // 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]
 // 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*)]

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]