[OPENMP, NVPTX] Fix linkage of the global entries.

The linkage of the global entries must be weak to enable support of redefinition of the same target regions in multiple compilation units. llvm-svn: 331768
2018-05-08 14:16:57 +00:00 · 2018-05-08 14:16:57 +00:00 · 9a70017537
parent 2580554333
commit 9a70017537
51 changed files with 377 additions and 377 deletions
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@ -3863,9 +3863,9 @@ void CGOpenMPRuntime::createOffloadEntry(
                            llvm::ConstantInt::get(CGM.Int32Ty, Flags),
                            llvm::ConstantInt::get(CGM.Int32Ty, 0)};
  std::string EntryName = getName({"omp_offloading", "entry", ""});
-  llvm::GlobalVariable *Entry =
-      createConstantGlobalStruct(CGM, getTgtOffloadEntryQTy(), Data,
-                                 Twine(EntryName).concat(Name), Linkage);
+  llvm::GlobalVariable *Entry = createConstantGlobalStruct(
+      CGM, getTgtOffloadEntryQTy(), Data, Twine(EntryName).concat(Name),
+      llvm::GlobalValue::WeakAnyLinkage);

  // The entry has to be created in the section the linker expects it to be.
  std::string Section = getName({"omp_offloading", "entries"});
@ -6281,13 +6281,13 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(

  if (CGM.getLangOpts().OpenMPIsDevice) {
    OutlinedFnID = llvm::ConstantExpr::getBitCast(OutlinedFn, CGM.Int8PtrTy);
-    OutlinedFn->setLinkage(llvm::GlobalValue::ExternalLinkage);
+    OutlinedFn->setLinkage(llvm::GlobalValue::WeakAnyLinkage);
    OutlinedFn->setDSOLocal(false);
  } else {
    std::string Name = getName({"omp_offload", "region_id"});
    OutlinedFnID = new llvm::GlobalVariable(
        CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
-        llvm::GlobalValue::PrivateLinkage,
+        llvm::GlobalValue::WeakAnyLinkage,
        llvm::Constant::getNullValue(CGM.Int8Ty), Name);
  }

--- a/clang/test/OpenMP/declare_target_codegen.cpp
+++ b/clang/test/OpenMP/declare_target_codegen.cpp
@ -64,7 +64,7 @@ int bar() { return 1 + foo() + bar() + baz1() + baz2(); }
 int maini1() {
  int a;
  static long aa = 32;
-// CHECK-DAG: define void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+1]](i32* dereferenceable{{.*}}, i64 {{.*}}, i64 {{.*}})
+// CHECK-DAG: define weak void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+1]](i32* dereferenceable{{.*}}, i64 {{.*}}, i64 {{.*}})
 #pragma omp target map(tofrom \
                       : a, b)
  {
@ -77,7 +77,7 @@ int maini1() {

 int baz3() { return 2 + baz2(); }
 int baz2() {
-// CHECK-DAG: define void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](i64 {{.*}})
+// CHECK-DAG: define weak void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](i64 {{.*}})
 #pragma omp target
  ++c;
  return 2 + baz3();
--- a/clang/test/OpenMP/declare_target_codegen_globalization.cpp
+++ b/clang/test/OpenMP/declare_target_codegen_globalization.cpp
@ -9,7 +9,7 @@ int bar() {
  return foo(a);
 }

-// CHECK: define void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+5]](i32* dereferenceable{{.*}})
+// CHECK: define weak void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+5]](i32* dereferenceable{{.*}})
 // CHECK-NOT: @__kmpc_data_sharing_push_stack

 int maini1() {
--- a/clang/test/OpenMP/declare_target_link_codegen.cpp
+++ b/clang/test/OpenMP/declare_target_link_codegen.cpp
@ -39,7 +39,7 @@ int maini1() {
  return 0;
 }

-// DEVICE: define void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-7]](i32* dereferenceable{{[^,]*}}
+// DEVICE: define weak void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-7]](i32* dereferenceable{{[^,]*}}
 // DEVICE: [[C_REF:%.+]] = load i32*, i32** @c_decl_tgt_link_ptr,
 // DEVICE: [[C:%.+]] = load i32, i32* [[C_REF]],
 // DEVICE: store i32 [[C]], i32* %
--- a/clang/test/OpenMP/nvptx_data_sharing.cpp
+++ b/clang/test/OpenMP/nvptx_data_sharing.cpp
@ -34,7 +34,7 @@ void test_ds(){

 /// ========= In the kernel function ========= ///

-// CK1: {{.*}}define void @__omp_offloading{{.*}}test_ds{{.*}}()
+// CK1: {{.*}}define weak void @__omp_offloading{{.*}}test_ds{{.*}}()
 // CK1: [[SHAREDARGS1:%.+]] = alloca i8**
 // CK1: [[SHAREDARGS2:%.+]] = alloca i8**
 // CK1: call void @__kmpc_kernel_init
--- a/clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp
@ -73,9 +73,9 @@ S cd = doo() + car() + caz() + baz();
 // DEVICE-DAG: ret void

 // HOST-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[C_ADDR]]\00"
-// HOST-DAG: @.omp_offloading.entry.[[C_ADDR]] = internal constant %struct.__tgt_offload_entry { i8* bitcast (i32* @[[C_ADDR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 4, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
+// HOST-DAG: @.omp_offloading.entry.[[C_ADDR]] = weak constant %struct.__tgt_offload_entry { i8* bitcast (i32* @[[C_ADDR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 4, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
 // HOST-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[CD_ADDR]]\00"
-// HOST-DAG: @.omp_offloading.entry.[[CD_ADDR]] = constant %struct.__tgt_offload_entry { i8* bitcast (%struct.S* @[[CD_ADDR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 4, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
+// HOST-DAG: @.omp_offloading.entry.[[CD_ADDR]] = weak constant %struct.__tgt_offload_entry { i8* bitcast (%struct.S* @[[CD_ADDR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 4, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
 // HOST-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[C_CTOR]]\00"
 // HOST-DAG: @.omp_offloading.entry.[[C_CTOR]] = weak constant %struct.__tgt_offload_entry { i8* @[[C_CTOR]], i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 0, i32 2, i32 0 }, section ".omp_offloading.entries", align 1
 // HOST-DAG: @.omp_offloading.entry_name{{.*}}= internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[CD_CTOR]]\00"
@ -91,7 +91,7 @@ int maini1() {
  return 0;
 }

-// DEVICE: define void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-7]](i32* dereferenceable{{[^,]*}}
+// DEVICE: define weak void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-7]](i32* dereferenceable{{[^,]*}}
 // DEVICE: [[C:%.+]] = load i32, i32* [[C_ADDR]],
 // DEVICE: store i32 [[C]], i32* %

--- a/clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp
@ -91,7 +91,7 @@ int foo(int n, double *ptr) {
    ptr[0]++;
  }

-  // TCHECK:  define void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]])
+  // TCHECK:  define weak void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]])
  // TCHECK:  [[PTR_ADDR:%.+]] = alloca double*,
  // TCHECK-NOT: alloca double*,
  // TCHECK:  store double* [[PTR_IN]], double** [[PTR_ADDR]],
--- a/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp
@ -125,7 +125,7 @@ int bar(int n){
 // CHECK: call void @__kmpc_for_static_fini(
 // CHECK: ret void

-// CHECK: define void @__omp_offloading_{{.*}}_l56(i[[SZ:64|32]] %{{[^,]+}}, [1000 x i32]* dereferenceable{{.*}}, i32* %{{[^)]+}})
+// CHECK: define weak void @__omp_offloading_{{.*}}_l56(i[[SZ:64|32]] %{{[^,]+}}, [1000 x i32]* dereferenceable{{.*}}, i32* %{{[^)]+}})
 // CHECK: call void [[OUTLINED:@__omp_outlined.*]](i32* %{{.+}}, i32* %{{.+}}, i[[SZ]] %{{.*}}, i[[SZ]] %{{.*}}, i[[SZ]] %{{.*}}, [1000 x i32]* %{{.*}}, i32* %{{.*}})
 // CHECK: define internal void [[OUTLINED]](i32* noalias %{{.*}}, i32* noalias %{{.*}} i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, [1000 x i32]* dereferenceable{{.*}}, i32* %{{.*}})

--- a/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_generic_mode_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_generic_mode_codegen.cpp
@ -21,7 +21,7 @@ int main(int argc, char **argv) {

 // CHECK: @__omp_offloading_{{.*}}_main_l16_exec_mode = weak constant i8 0

-// CHECK: define void @__omp_offloading_{{.*}}_main_l16(i{{64|32}} %{{[^,].*}}, i32* dereferenceable{{[^,]*}}, i{{64|32}} %{{[^,)]*}})
+// CHECK: define weak void @__omp_offloading_{{.*}}_main_l16(i{{64|32}} %{{[^,].*}}, i32* dereferenceable{{[^,]*}}, i{{64|32}} %{{[^,)]*}})
 // CHECK: [[TID:%.+]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @
 // CHECK: call void @__kmpc_spmd_kernel_init(
 // CHECK: call void @__kmpc_for_static_init_4(
--- a/clang/test/OpenMP/target_codegen.cpp
+++ b/clang/test/OpenMP/target_codegen.cpp
@ -64,13 +64,13 @@
 // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40]
 // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 288, i64 288, i64 288, i64 547]
 // CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 288, i64 288, i64 288, i64 547]
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
--- a/clang/test/OpenMP/target_codegen_registration.cpp
+++ b/clang/test/OpenMP/target_codegen_registration.cpp
@ -81,45 +81,45 @@

 // We have 7 target regions

-// CHECK-DAG: {{@.+}} = private constant i8 0
-// TCHECK-NOT: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
+// TCHECK-NOT: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]

-// CHECK-NTARGET-NOT: private constant i8 0
+// CHECK-NTARGET-NOT: weak constant i8 0
 // CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i

 // CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
@ -340,18 +340,18 @@ struct ST {
 //CHECK-DAG: define internal void @[[NAME12]](
 //CHECK-DAG: call void @[[NAME12]](

-//TCHECK-DAG: define void @[[NAME1]](
-//TCHECK-DAG: define void @[[NAME2]](
-//TCHECK-DAG: define void @[[NAME3]](
-//TCHECK-DAG: define void @[[NAME4]](
-//TCHECK-DAG: define void @[[NAME5]](
-//TCHECK-DAG: define void @[[NAME6]](
-//TCHECK-DAG: define void @[[NAME7]](
-//TCHECK-DAG: define void @[[NAME8]](
-//TCHECK-DAG: define void @[[NAME9]](
-//TCHECK-DAG: define void @[[NAME10]](
-//TCHECK-DAG: define void @[[NAME11]](
-//TCHECK-DAG: define void @[[NAME12]](
+//TCHECK-DAG: define weak void @[[NAME1]](
+//TCHECK-DAG: define weak void @[[NAME2]](
+//TCHECK-DAG: define weak void @[[NAME3]](
+//TCHECK-DAG: define weak void @[[NAME4]](
+//TCHECK-DAG: define weak void @[[NAME5]](
+//TCHECK-DAG: define weak void @[[NAME6]](
+//TCHECK-DAG: define weak void @[[NAME7]](
+//TCHECK-DAG: define weak void @[[NAME8]](
+//TCHECK-DAG: define weak void @[[NAME9]](
+//TCHECK-DAG: define weak void @[[NAME10]](
+//TCHECK-DAG: define weak void @[[NAME11]](
+//TCHECK-DAG: define weak void @[[NAME12]](

 // CHECK-NTARGET-NOT: __tgt_target
 // CHECK-NTARGET-NOT: __tgt_register_lib
--- a/clang/test/OpenMP/target_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_depend_codegen.cpp
@ -49,7 +49,7 @@

 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 0, i[[SZ]] 4]
 // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 288]
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
--- a/clang/test/OpenMP/target_firstprivate_codegen.cpp
+++ b/clang/test/OpenMP/target_firstprivate_codegen.cpp
@ -120,7 +120,7 @@ int foo(int n, double *ptr) {
  // CHECK:  [[PTR_GEP_ARG:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
  // CHECK:  {{.+}} = call i32 @__tgt_target(i64 -1, {{.+}}, i32 1, i8** [[BASE_PTR_GEP_ARG]], i8** [[PTR_GEP_ARG]], i[[SZ]]* getelementptr inbounds ([1 x i[[SZ]]], [1 x i[[SZ]]]* [[SIZET]], i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[MAPT]], i32 0, i32 0))
  
-  // TCHECK:  define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]])
+  // TCHECK:  define weak void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]])
  // TCHECK:  [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
  // TCHECK-NOT: alloca i{{[0-9]+}},
  // TCHECK:  store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
@ -330,7 +330,7 @@ int foo(int n, double *ptr) {
  // CHECK:  [[PTR_GEP_ARG3:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
  // CHECK: {{.+}} = call i32 @__tgt_target(i64 -1, {{.+}}, i32 1, i8** [[BASE_PTR_GEP_ARG3]], i8** [[PTR_GEP_ARG3]], i[[SZ]]* getelementptr inbounds ([1 x i[[SZ]]], [1 x i[[SZ]]]* [[SIZET3]], i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[MAPT3]], i32 0, i32 0))

-  // TCHECK:  define void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]])
+  // TCHECK:  define weak void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]])
  // TCHECK:  [[PTR_ADDR:%.+]] = alloca double*,
  // TCHECK-NOT: alloca double*,
  // TCHECK:  store double* [[PTR_IN]], double** [[PTR_ADDR]],
@ -370,7 +370,7 @@ int fstatic(int n) {
  return a;
 }

-// TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: define weak void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
 // TCHECK:  [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
 // TCHECK:  [[A3_ADDR:%.+]] = alloca i{{[0-9]+}},
 // TCHECK:  [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
@ -466,7 +466,7 @@ struct S1 {
  // only check that we use the map types stored in the global variable
  // CHECK:  call i32 @__tgt_target(i64 -1, {{.+}}, i32 5, i8** {{.+}}, i8** {{.+}}, i{{[0-9]+}}* {{.+}}, i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPT4]], i32 0, i32 0))
  
-  // TCHECK: define void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[B_IN:%.+]], i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i{{[0-9]+}}{{.+}} [[C_IN:%.+]])
+  // TCHECK: define weak void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[B_IN:%.+]], i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i{{[0-9]+}}{{.+}} [[C_IN:%.+]])
  // TCHECK:  [[TH_ADDR:%.+]] = alloca [[S1]]*,
  // TCHECK:  [[B_ADDR:%.+]] = alloca i{{[0-9]+}},
  // TCHECK:  [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
@ -578,7 +578,7 @@ int bar(int n, double *ptr){
 // CHECK:  call i32 @__tgt_target(i64 -1, {{.+}}, i32 2, i8** {{.+}}, i8** {{.+}}, i[[SZ]]* getelementptr inbounds ([2 x i[[SZ]]], [2 x i[[SZ]]]* [[SIZET6]], i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPT6]], i32 0, i32 0))


-// TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: define weak void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
 // TCHECK:  [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
 // TCHECK:  [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
 // TCHECK-NOT: alloca i{{[0-9]+}},
--- a/clang/test/OpenMP/target_parallel_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_codegen.cpp
@ -66,13 +66,13 @@
 // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40]
 // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 288, i64 288, i64 288, i64 547]
 // CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 288, i64 288, i64 288, i64 547]
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
--- a/clang/test/OpenMP/target_parallel_codegen_registration.cpp
+++ b/clang/test/OpenMP/target_parallel_codegen_registration.cpp
@ -81,45 +81,45 @@

 // We have 7 target regions

-// CHECK-DAG: {{@.+}} = private constant i8 0
-// TCHECK-NOT: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
+// TCHECK-NOT: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]

-// CHECK-NTARGET-NOT: private constant i8 0
+// CHECK-NTARGET-NOT: weak constant i8 0
 // CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i

 // CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
@ -340,18 +340,18 @@ struct ST {
 //CHECK-DAG: define internal void @[[NAME12]](
 //CHECK-DAG: call void @[[NAME12]](

-//TCHECK-DAG: define void @[[NAME1]](
-//TCHECK-DAG: define void @[[NAME2]](
-//TCHECK-DAG: define void @[[NAME3]](
-//TCHECK-DAG: define void @[[NAME4]](
-//TCHECK-DAG: define void @[[NAME5]](
-//TCHECK-DAG: define void @[[NAME6]](
-//TCHECK-DAG: define void @[[NAME7]](
-//TCHECK-DAG: define void @[[NAME8]](
-//TCHECK-DAG: define void @[[NAME9]](
-//TCHECK-DAG: define void @[[NAME10]](
-//TCHECK-DAG: define void @[[NAME11]](
-//TCHECK-DAG: define void @[[NAME12]](
+//TCHECK-DAG: define weak void @[[NAME1]](
+//TCHECK-DAG: define weak void @[[NAME2]](
+//TCHECK-DAG: define weak void @[[NAME3]](
+//TCHECK-DAG: define weak void @[[NAME4]](
+//TCHECK-DAG: define weak void @[[NAME5]](
+//TCHECK-DAG: define weak void @[[NAME6]](
+//TCHECK-DAG: define weak void @[[NAME7]](
+//TCHECK-DAG: define weak void @[[NAME8]](
+//TCHECK-DAG: define weak void @[[NAME9]](
+//TCHECK-DAG: define weak void @[[NAME10]](
+//TCHECK-DAG: define weak void @[[NAME11]](
+//TCHECK-DAG: define weak void @[[NAME12]](

 // CHECK-NTARGET-NOT: __tgt_target
 // CHECK-NTARGET-NOT: __tgt_register_lib
--- a/clang/test/OpenMP/target_parallel_debug_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_debug_codegen.cpp
@ -65,7 +65,7 @@ int main() {
 // CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
 // CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})

-// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
+// CHECK: define weak void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
 // CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
 // CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})

@ -83,7 +83,7 @@ int main() {
 // CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
 // CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})

-// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
+// CHECK: define weak void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
 // CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
 // CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
 // CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
@ -105,7 +105,7 @@ int main() {
 // CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
 // CHECK: call void @[[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 addrspace(1)* {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})

-// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i32* dereferenceable{{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
+// CHECK: define weak void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i32* dereferenceable{{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
 // CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
 // CHECK: addrspacecast i32* %{{.+}} to i32 addrspace(1)*
 // CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
--- a/clang/test/OpenMP/target_parallel_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_depend_codegen.cpp
@ -49,7 +49,7 @@

 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 0, i[[SZ]] 4]
 // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 288]
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
--- a/clang/test/OpenMP/target_parallel_for_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_codegen.cpp
@ -66,13 +66,13 @@
 // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40]
 // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 288, i64 288, i64 288, i64 547]
 // CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 288, i64 288, i64 288, i64 547]
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
--- a/clang/test/OpenMP/target_parallel_for_codegen_registration.cpp
+++ b/clang/test/OpenMP/target_parallel_for_codegen_registration.cpp
@ -81,45 +81,45 @@

 // We have 7 target regions

-// CHECK-DAG: {{@.+}} = private constant i8 0
-// TCHECK-NOT: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
+// TCHECK-NOT: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]

-// CHECK-NTARGET-NOT: private constant i8 0
+// CHECK-NTARGET-NOT: weak constant i8 0
 // CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i

 // CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
@ -349,18 +349,18 @@ struct ST {
 //CHECK-DAG: define internal void @[[NAME12]](
 //CHECK-DAG: call void @[[NAME12]](

-//TCHECK-DAG: define void @[[NAME1]](
-//TCHECK-DAG: define void @[[NAME2]](
-//TCHECK-DAG: define void @[[NAME3]](
-//TCHECK-DAG: define void @[[NAME4]](
-//TCHECK-DAG: define void @[[NAME5]](
-//TCHECK-DAG: define void @[[NAME6]](
-//TCHECK-DAG: define void @[[NAME7]](
-//TCHECK-DAG: define void @[[NAME8]](
-//TCHECK-DAG: define void @[[NAME9]](
-//TCHECK-DAG: define void @[[NAME10]](
-//TCHECK-DAG: define void @[[NAME11]](
-//TCHECK-DAG: define void @[[NAME12]](
+//TCHECK-DAG: define weak void @[[NAME1]](
+//TCHECK-DAG: define weak void @[[NAME2]](
+//TCHECK-DAG: define weak void @[[NAME3]](
+//TCHECK-DAG: define weak void @[[NAME4]](
+//TCHECK-DAG: define weak void @[[NAME5]](
+//TCHECK-DAG: define weak void @[[NAME6]](
+//TCHECK-DAG: define weak void @[[NAME7]](
+//TCHECK-DAG: define weak void @[[NAME8]](
+//TCHECK-DAG: define weak void @[[NAME9]](
+//TCHECK-DAG: define weak void @[[NAME10]](
+//TCHECK-DAG: define weak void @[[NAME11]](
+//TCHECK-DAG: define weak void @[[NAME12]](

 // CHECK-NTARGET-NOT: __tgt_target
 // CHECK-NTARGET-NOT: __tgt_register_lib
--- a/clang/test/OpenMP/target_parallel_for_debug_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_debug_codegen.cpp
@ -65,7 +65,7 @@ int main() {
 // CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
 // CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})

-// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
+// CHECK: define weak void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
 // CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
 // CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})

@ -83,7 +83,7 @@ int main() {
 // CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
 // CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})

-// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
+// CHECK: define weak void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
 // CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
 // CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
 // CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
@ -105,7 +105,7 @@ int main() {
 // CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
 // CHECK: call void @[[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 addrspace(1)* {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})

-// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i32* dereferenceable{{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
+// CHECK: define weak void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i32* dereferenceable{{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
 // CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
 // CHECK: addrspacecast i32* %{{.+}} to i32 addrspace(1)*
 // CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
--- a/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp
@ -49,7 +49,7 @@

 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 0, i[[SZ]] 4]
 // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 288]
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
--- a/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp
@ -65,13 +65,13 @@
 // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40]
 // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 288, i64 288, i64 288, i64 547]
 // CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 288, i64 288, i64 288, i64 547]
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
--- a/clang/test/OpenMP/target_parallel_for_simd_codegen_registration.cpp
+++ b/clang/test/OpenMP/target_parallel_for_simd_codegen_registration.cpp
@ -81,45 +81,45 @@

 // We have 7 target regions

-// CHECK-DAG: {{@.+}} = private constant i8 0
-// TCHECK-NOT: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
+// TCHECK-NOT: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]

-// CHECK-NTARGET-NOT: private constant i8 0
+// CHECK-NTARGET-NOT: weak constant i8 0
 // CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i

 // CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
@ -349,18 +349,18 @@ struct ST {
 //CHECK-DAG: define internal void @[[NAME12]](
 //CHECK-DAG: call void @[[NAME12]](

-//TCHECK-DAG: define void @[[NAME1]](
-//TCHECK-DAG: define void @[[NAME2]](
-//TCHECK-DAG: define void @[[NAME3]](
-//TCHECK-DAG: define void @[[NAME4]](
-//TCHECK-DAG: define void @[[NAME5]](
-//TCHECK-DAG: define void @[[NAME6]](
-//TCHECK-DAG: define void @[[NAME7]](
-//TCHECK-DAG: define void @[[NAME8]](
-//TCHECK-DAG: define void @[[NAME9]](
-//TCHECK-DAG: define void @[[NAME10]](
-//TCHECK-DAG: define void @[[NAME11]](
-//TCHECK-DAG: define void @[[NAME12]](
+//TCHECK-DAG: define weak void @[[NAME1]](
+//TCHECK-DAG: define weak void @[[NAME2]](
+//TCHECK-DAG: define weak void @[[NAME3]](
+//TCHECK-DAG: define weak void @[[NAME4]](
+//TCHECK-DAG: define weak void @[[NAME5]](
+//TCHECK-DAG: define weak void @[[NAME6]](
+//TCHECK-DAG: define weak void @[[NAME7]](
+//TCHECK-DAG: define weak void @[[NAME8]](
+//TCHECK-DAG: define weak void @[[NAME9]](
+//TCHECK-DAG: define weak void @[[NAME10]](
+//TCHECK-DAG: define weak void @[[NAME11]](
+//TCHECK-DAG: define weak void @[[NAME12]](

 // CHECK-NTARGET-NOT: __tgt_target
 // CHECK-NTARGET-NOT: __tgt_register_lib
--- a/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp
@ -49,7 +49,7 @@

 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 0, i[[SZ]] 4]
 // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 288]
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
--- a/clang/test/OpenMP/target_parallel_if_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_if_codegen.cpp
@ -53,12 +53,12 @@

 // We have 6 target regions

-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
--- a/clang/test/OpenMP/target_parallel_num_threads_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_num_threads_codegen.cpp
@ -53,12 +53,12 @@

 // We have 6 target regions

-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
--- a/clang/test/OpenMP/target_private_codegen.cpp
+++ b/clang/test/OpenMP/target_private_codegen.cpp
@ -45,28 +45,28 @@ int foo(int n) {
  {
  }

-  // TCHECK:  define void @__omp_offloading_{{.+}}()
+  // TCHECK:  define weak void @__omp_offloading_{{.+}}()
  // TCHECK:  [[A:%.+]] = alloca i{{[0-9]+}},
  // TCHECK-NOT: store {{.+}}, {{.+}} [[A]],
-  // TCHECK:  ret void  
+  // TCHECK:  ret void

 #pragma omp target private(a)
  {
    a = 1;
  }

-  // TCHECK:  define void @__omp_offloading_{{.+}}()
+  // TCHECK:  define weak void @__omp_offloading_{{.+}}()
  // TCHECK:  [[A:%.+]] = alloca i{{[0-9]+}},
  // TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
  // TCHECK:  ret void
-  
-  #pragma omp target private(a, aa)
+
+#pragma omp target private(a, aa)
  {
    a = 1;
    aa = 1;
  }

-  // TCHECK:  define void @__omp_offloading_{{.+}}()
+  // TCHECK:  define weak void @__omp_offloading_{{.+}}()
  // TCHECK:  [[A:%.+]] = alloca i{{[0-9]+}},
  // TCHECK:  [[A2:%.+]] = alloca i{{[0-9]+}},
  // TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
@ -85,7 +85,7 @@ int foo(int n) {
  }
  // make sure that private variables are generated in all cases and that we use those instances for operations inside the
  // target region
-  // TCHECK:  define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i{{[0-9]+}} [[VLA3:%.+]])
+  // TCHECK:  define weak void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i{{[0-9]+}} [[VLA3:%.+]])
  // TCHECK:  [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
  // TCHECK:  [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},
  // TCHECK:  [[VLA_ADDR4:%.+]] = alloca i{{[0-9]+}},
@ -112,7 +112,7 @@ int foo(int n) {
  // b[2] = 1.0
  // TCHECK:  [[B_GEP:%.+]] = getelementptr inbounds [10 x float], [10 x float]* [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
  // TCHECK:  store float 1.0{{.*}}, float* [[B_GEP]],
-  
+
  // bn[3] = 1.0
  // TCHECK:  [[BN_GEP:%.+]] = getelementptr inbounds float, float* [[VLA5]], i{{[0-9]+}} 3
  // TCHECK:  store float 1.0{{.*}}, float* [[BN_GEP]],
@ -131,7 +131,7 @@ int foo(int n) {
  // d.X = 1
  // [[X_FIELD:%.+]] = getelementptr inbounds [[TT]] [[TT]]* [[D]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
  // store i{{[0-9]+}} 1, i{{[0-9]+}}* [[X_FIELD]],
-  
+
  // d.Y = 1
  // [[Y_FIELD:%.+]] = getelementptr inbounds [[TT]] [[TT]]* [[D]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
  // store i{{[0-9]+}} 1, i{{[0-9]+}}* [[Y_FIELD]],
@ -179,7 +179,7 @@ int fstatic(int n) {
  return a;
 }

-// TCHECK: define void @__omp_offloading_{{.+}}()
+// TCHECK: define weak void @__omp_offloading_{{.+}}()
 // TCHECK:  [[A:%.+]] = alloca i{{[0-9]+}},
 // TCHECK:  [[A2:%.+]] = alloca i{{[0-9]+}},
 // TCHECK:  [[A3:%.+]] = alloca i{{[0-9]+}},
@ -207,7 +207,7 @@ struct S1 {
    return c[1][1] + (int)b;
  }

-  // TCHECK: define void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]])
+  // TCHECK: define weak void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]])
  // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
  // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
  // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},
@ -261,7 +261,7 @@ int bar(int n){
 }

 // template
-// TCHECK: define void @__omp_offloading_{{.+}}()
+// TCHECK: define weak void @__omp_offloading_{{.+}}()
 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
 // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
 // TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}],
--- a/clang/test/OpenMP/target_reduction_codegen.cpp
+++ b/clang/test/OpenMP/target_reduction_codegen.cpp
@ -45,7 +45,7 @@ int foo(int n) {
  {
  }

-  // TCHECK: define void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
+  // TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
  // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store {{.+}}, {{.+}} [[A]],
  // TCHECK: load i32*, i32** [[A]],
@ -56,7 +56,7 @@ int foo(int n) {
    a = 1;
  }

-  // TCHECK:  define void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
+  // TCHECK:  define weak void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
  // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store {{.+}}, {{.+}} [[A]],
  // TCHECK: [[REF:%.+]] = load i32*, i32** [[A]],
@ -69,7 +69,7 @@ int foo(int n) {
    aa = 1;
  }

-  // TCHECK:  define void @__omp_offloading_{{.+}}(i32*{{.+}} [[A:%.+]], i16*{{.+}} [[AA:%.+]])
+  // TCHECK:  define weak void @__omp_offloading_{{.+}}(i32*{{.+}} [[A:%.+]], i16*{{.+}} [[AA:%.+]])
  // TCHECK:  [[A:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK:  [[AA:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store {{.+}}, {{.+}} [[A]],
@ -118,7 +118,7 @@ int fstatic(int n) {
  return a;
 }

-// TCHECK: define void @__omp_offloading_{{.+}}(i32*{{.+}}, i16*{{.+}}, i8*{{.+}}, [10 x i32]*{{.+}})
+// TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}}, i16*{{.+}}, i8*{{.+}}, [10 x i32]*{{.+}})
 // TCHECK:  [[A:%.+]] = alloca i{{[0-9]+}}*,
 // TCHECK:  [[A2:%.+]] = alloca i{{[0-9]+}}*,
 // TCHECK:  [[A3:%.+]] = alloca i{{[0-9]+}}*,
@ -154,7 +154,7 @@ struct S1 {
    return c[1][1] + (int)b;
  }

-  // TCHECK: define void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i32*{{.+}}, i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i16*{{.+}})
+  // TCHECK: define weak void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i32*{{.+}}, i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i16*{{.+}})
  // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
  // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
@ -206,7 +206,7 @@ int bar(int n){
 }

 // template
-// TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, [10 x i32]*{{.+}})
+// TCHECK: define weak void @__omp_offloading_{{.+}}(i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, [10 x i32]*{{.+}})
 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
 // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}}*,
 // TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}]*,
--- a/clang/test/OpenMP/target_simd_codegen.cpp
+++ b/clang/test/OpenMP/target_simd_codegen.cpp
@ -62,13 +62,13 @@
 // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40]
 // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 288, i64 288, i64 288, i64 547]
 // CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 288, i64 288, i64 288, i64 547]
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
--- a/clang/test/OpenMP/target_simd_codegen_registration.cpp
+++ b/clang/test/OpenMP/target_simd_codegen_registration.cpp
@ -81,45 +81,45 @@

 // We have 7 target regions

-// CHECK-DAG: {{@.+}} = private constant i8 0
-// TCHECK-NOT: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
+// TCHECK-NOT: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]

-// CHECK-NTARGET-NOT: private constant i8 0
+// CHECK-NTARGET-NOT: weak constant i8 0
 // CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i

 // CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
@ -349,18 +349,18 @@ struct ST {
 //CHECK-DAG: define internal void @[[NAME12]](
 //CHECK-DAG: call void @[[NAME12]](

-//TCHECK-DAG: define void @[[NAME1]](
-//TCHECK-DAG: define void @[[NAME2]](
-//TCHECK-DAG: define void @[[NAME3]](
-//TCHECK-DAG: define void @[[NAME4]](
-//TCHECK-DAG: define void @[[NAME5]](
-//TCHECK-DAG: define void @[[NAME6]](
-//TCHECK-DAG: define void @[[NAME7]](
-//TCHECK-DAG: define void @[[NAME8]](
-//TCHECK-DAG: define void @[[NAME9]](
-//TCHECK-DAG: define void @[[NAME10]](
-//TCHECK-DAG: define void @[[NAME11]](
-//TCHECK-DAG: define void @[[NAME12]](
+//TCHECK-DAG: define weak void @[[NAME1]](
+//TCHECK-DAG: define weak void @[[NAME2]](
+//TCHECK-DAG: define weak void @[[NAME3]](
+//TCHECK-DAG: define weak void @[[NAME4]](
+//TCHECK-DAG: define weak void @[[NAME5]](
+//TCHECK-DAG: define weak void @[[NAME6]](
+//TCHECK-DAG: define weak void @[[NAME7]](
+//TCHECK-DAG: define weak void @[[NAME8]](
+//TCHECK-DAG: define weak void @[[NAME9]](
+//TCHECK-DAG: define weak void @[[NAME10]](
+//TCHECK-DAG: define weak void @[[NAME11]](
+//TCHECK-DAG: define weak void @[[NAME12]](

 // CHECK-NTARGET-NOT: __tgt_target
 // CHECK-NTARGET-NOT: __tgt_register_lib
--- a/clang/test/OpenMP/target_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_simd_depend_codegen.cpp
@ -49,7 +49,7 @@

 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 0, i[[SZ]] 4]
 // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 288]
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
--- a/clang/test/OpenMP/target_teams_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_codegen.cpp
@ -68,13 +68,13 @@
 // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40]
 // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 288, i64 288, i64 288, i64 547]
 // CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 288, i64 288, i64 288, i64 547]
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
--- a/clang/test/OpenMP/target_teams_codegen_registration.cpp
+++ b/clang/test/OpenMP/target_teams_codegen_registration.cpp
@ -81,45 +81,45 @@

 // We have 7 target regions

-// CHECK-DAG: {{@.+}} = private constant i8 0
-// TCHECK-NOT: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
+// TCHECK-NOT: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]

-// CHECK-NTARGET-NOT: private constant i8 0
+// CHECK-NTARGET-NOT: weak constant i8 0
 // CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i

 // CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
@ -340,18 +340,18 @@ struct ST {
 //CHECK-DAG: define internal void @[[NAME12]](
 //CHECK-DAG: call void @[[NAME12]](

-//TCHECK-DAG: define void @[[NAME1]](
-//TCHECK-DAG: define void @[[NAME2]](
-//TCHECK-DAG: define void @[[NAME3]](
-//TCHECK-DAG: define void @[[NAME4]](
-//TCHECK-DAG: define void @[[NAME5]](
-//TCHECK-DAG: define void @[[NAME6]](
-//TCHECK-DAG: define void @[[NAME7]](
-//TCHECK-DAG: define void @[[NAME8]](
-//TCHECK-DAG: define void @[[NAME9]](
-//TCHECK-DAG: define void @[[NAME10]](
-//TCHECK-DAG: define void @[[NAME11]](
-//TCHECK-DAG: define void @[[NAME12]](
+//TCHECK-DAG: define weak void @[[NAME1]](
+//TCHECK-DAG: define weak void @[[NAME2]](
+//TCHECK-DAG: define weak void @[[NAME3]](
+//TCHECK-DAG: define weak void @[[NAME4]](
+//TCHECK-DAG: define weak void @[[NAME5]](
+//TCHECK-DAG: define weak void @[[NAME6]](
+//TCHECK-DAG: define weak void @[[NAME7]](
+//TCHECK-DAG: define weak void @[[NAME8]](
+//TCHECK-DAG: define weak void @[[NAME9]](
+//TCHECK-DAG: define weak void @[[NAME10]](
+//TCHECK-DAG: define weak void @[[NAME11]](
+//TCHECK-DAG: define weak void @[[NAME12]](

 // CHECK-NTARGET-NOT: __tgt_target
 // CHECK-NTARGET-NOT: __tgt_register_lib
--- a/clang/test/OpenMP/target_teams_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_depend_codegen.cpp
@ -49,7 +49,7 @@

 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 0, i[[SZ]] 4]
 // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 288]
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
--- a/clang/test/OpenMP/target_teams_distribute_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_codegen.cpp
@ -68,13 +68,13 @@
 // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [5 x i64] [i64 288, i64 288, i64 288, i64 288, i64 547]
 // CHECK-DAG: [[SIZET7:@.+]] = private unnamed_addr constant [3 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 40]
 // CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [3 x i64] [i64 288, i64 288, i64 547]
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
--- a/clang/test/OpenMP/target_teams_distribute_codegen_registration.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_codegen_registration.cpp
@ -81,45 +81,45 @@

 // We have 7 target regions

-// CHECK-DAG: {{@.+}} = private constant i8 0
-// TCHECK-NOT: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
+// TCHECK-NOT: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]

-// CHECK-NTARGET-NOT: private constant i8 0
+// CHECK-NTARGET-NOT: weak constant i8 0
 // CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i

 // CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
@ -349,18 +349,18 @@ struct ST {
 //CHECK-DAG: define internal void @[[NAME12]](
 //CHECK-DAG: call void @[[NAME12]](

-//TCHECK-DAG: define void @[[NAME1]](
-//TCHECK-DAG: define void @[[NAME2]](
-//TCHECK-DAG: define void @[[NAME3]](
-//TCHECK-DAG: define void @[[NAME4]](
-//TCHECK-DAG: define void @[[NAME5]](
-//TCHECK-DAG: define void @[[NAME6]](
-//TCHECK-DAG: define void @[[NAME7]](
-//TCHECK-DAG: define void @[[NAME8]](
-//TCHECK-DAG: define void @[[NAME9]](
-//TCHECK-DAG: define void @[[NAME10]](
-//TCHECK-DAG: define void @[[NAME11]](
-//TCHECK-DAG: define void @[[NAME12]](
+//TCHECK-DAG: define weak void @[[NAME1]](
+//TCHECK-DAG: define weak void @[[NAME2]](
+//TCHECK-DAG: define weak void @[[NAME3]](
+//TCHECK-DAG: define weak void @[[NAME4]](
+//TCHECK-DAG: define weak void @[[NAME5]](
+//TCHECK-DAG: define weak void @[[NAME6]](
+//TCHECK-DAG: define weak void @[[NAME7]](
+//TCHECK-DAG: define weak void @[[NAME8]](
+//TCHECK-DAG: define weak void @[[NAME9]](
+//TCHECK-DAG: define weak void @[[NAME10]](
+//TCHECK-DAG: define weak void @[[NAME11]](
+//TCHECK-DAG: define weak void @[[NAME12]](

 // CHECK-NTARGET-NOT: __tgt_target
 // CHECK-NTARGET-NOT: __tgt_register_lib
--- a/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp
@ -49,7 +49,7 @@

 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 0, i[[SZ]] 4]
 // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 288]
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_codegen.cpp
@ -77,7 +77,7 @@ int target_teams_fun(int *g){

  // outlined target regions
  // HCK1: define internal void @[[OFFL1]](i{{32|64}} [[N_ARG:%.+]], i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]])
-  // TCK1: define void @{{.+}}target_teams_fun{{.*}}(i{{32|64}} [[N_ARG:%.+]], {{.+}}, i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]])
+  // TCK1: define weak void @{{.+}}target_teams_fun{{.*}}(i{{32|64}} [[N_ARG:%.+]], {{.+}}, i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]])
  // CK1: [[N_ADDR:%.+]] = alloca i{{32|64}},
  // CK1: [[TE_ADDR:%.+]] = alloca i{{32|64}},
  // CK1: [[TH_ADDR:%.+]] = alloca i{{32|64}},
@ -101,7 +101,7 @@ int target_teams_fun(int *g){
  // CK1: ret void

  // HCK1: define internal void @[[OFFL2]](
-  // TCK1: define void @{{.+}}target_teams_fun{{.+}}(
+  // TCK1: define weak void @{{.+}}target_teams_fun{{.+}}(
  // CK1: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 3, {{.+}} @[[OUTL2:.+]] to {{.+}}, {{.+}}, {{.+}})
  // CK1: ret void

--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp
@ -49,7 +49,7 @@

 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 0, i[[SZ]] 4]
 // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 288]
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_firstprivate_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_firstprivate_codegen.cpp
@ -115,7 +115,7 @@ int main() {
 #pragma omp target teams distribute parallel for firstprivate(g, g1, sivar)
  for (int i = 0; i < 2; ++i) {
    // HLAMBDA: define{{.*}} internal{{.*}} void @[[LOFFL1]](i{{64|32}} {{%.+}}, i{{64|32}} {{%.+}})
-    // TLAMBDA: define void @[[LOFFL1:.+]](i{{64|32}} {{%.+}}, i{{64|32}} {{%.+}})
+    // TLAMBDA: define weak void @[[LOFFL1:.+]](i{{64|32}} {{%.+}}, i{{64|32}} {{%.+}})
    // LAMBDA: {{%.+}} = alloca i{{[0-9]+}},
    // LAMBDA: {{%.+}} = alloca i{{[0-9]+}},
    // LAMBDA: {{%.+}} = alloca i{{[0-9]+}},
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_private_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_private_codegen.cpp
@ -282,7 +282,7 @@ int main() {
 // HCHECK: ret

 // HCHECK: define {{.*}}void @[[TOFFL1]]()
-// TCHECK: define void @[[TOFFL1:.+]]()
+// TCHECK: define weak void @[[TOFFL1:.+]]()
 // CHECK: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 0, {{.+}} @[[TOUTL1:.+]] to {{.+}})
 // CHECK: ret void

--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_codegen.cpp
@ -80,7 +80,7 @@ int target_teams_fun(int *g){

  // outlined target regions
  // HCK1: define internal void @[[OFFL1]](i{{32|64}} [[I_ARG:%.+]], i{{32|64}} [[N_ARG:%.+]], {{.+}}, i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]])
-  // TCK1: define void @{{.+}}target_teams_fun{{.*}}(i{{32|64}} [[I_ARG:%.+]], i{{32|64}} [[N_ARG:%.+]], {{.+}}, i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]])
+  // TCK1: define weak void @{{.+}}target_teams_fun{{.*}}(i{{32|64}} [[I_ARG:%.+]], i{{32|64}} [[N_ARG:%.+]], {{.+}}, i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]])
  // CK1: [[I_ADDR:%.+]] = alloca i{{32|64}},
  // CK1: [[N_ADDR:%.+]] = alloca i{{32|64}},
  // CK1: [[TE_ADDR:%.+]] = alloca i{{32|64}},
@ -110,7 +110,7 @@ int target_teams_fun(int *g){
  // CK1: ret void

  // HCK1: define internal void @[[OFFL2]](
-  // TCK1: define void @{{.+}}target_teams_fun{{.+}}(
+  // TCK1: define weak void @{{.+}}target_teams_fun{{.+}}(
  // CK1: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 3, {{.+}} @[[OUTL2:.+]] to {{.+}}, {{.+}}, {{.+}})
  // CK1: ret void

--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_codegen_registration.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_codegen_registration.cpp
@ -81,45 +81,45 @@

 // We have 7 target regions

-// CHECK-DAG: {{@.+}} = private constant i8 0
-// TCHECK-NOT: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
+// TCHECK-NOT: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]

-// CHECK-NTARGET-NOT: private constant i8 0
+// CHECK-NTARGET-NOT: weak constant i8 0
 // CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i

 // CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
@ -349,18 +349,18 @@ struct ST {
 //CHECK-DAG: define internal void @[[NAME12]](
 //CHECK-DAG: call void @[[NAME12]](

-//TCHECK-DAG: define void @[[NAME1]](
-//TCHECK-DAG: define void @[[NAME2]](
-//TCHECK-DAG: define void @[[NAME3]](
-//TCHECK-DAG: define void @[[NAME4]](
-//TCHECK-DAG: define void @[[NAME5]](
-//TCHECK-DAG: define void @[[NAME6]](
-//TCHECK-DAG: define void @[[NAME7]](
-//TCHECK-DAG: define void @[[NAME8]](
-//TCHECK-DAG: define void @[[NAME9]](
-//TCHECK-DAG: define void @[[NAME10]](
-//TCHECK-DAG: define void @[[NAME11]](
-//TCHECK-DAG: define void @[[NAME12]](
+//TCHECK-DAG: define weak void @[[NAME1]](
+//TCHECK-DAG: define weak void @[[NAME2]](
+//TCHECK-DAG: define weak void @[[NAME3]](
+//TCHECK-DAG: define weak void @[[NAME4]](
+//TCHECK-DAG: define weak void @[[NAME5]](
+//TCHECK-DAG: define weak void @[[NAME6]](
+//TCHECK-DAG: define weak void @[[NAME7]](
+//TCHECK-DAG: define weak void @[[NAME8]](
+//TCHECK-DAG: define weak void @[[NAME9]](
+//TCHECK-DAG: define weak void @[[NAME10]](
+//TCHECK-DAG: define weak void @[[NAME11]](
+//TCHECK-DAG: define weak void @[[NAME12]](

 // CHECK-NTARGET-NOT: __tgt_target
 // CHECK-NTARGET-NOT: __tgt_register_lib
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp
@ -49,7 +49,7 @@

 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 0, i[[SZ]] 4]
 // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 288]
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_firstprivate_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_firstprivate_codegen.cpp
@ -115,7 +115,7 @@ int main() {
 #pragma omp target teams distribute parallel for simd firstprivate(g, g1, sivar)
  for (int i = 0; i < 2; ++i) {
    // HLAMBDA: define{{.*}} internal{{.*}} void @[[LOFFL1]](i{{64|32}} {{%.+}}, i{{64|32}} {{%.+}})
-    // TLAMBDA: define void @[[LOFFL1:.+]](i{{64|32}} {{%.+}}, i{{64|32}} {{%.+}})
+    // TLAMBDA: define weak void @[[LOFFL1:.+]](i{{64|32}} {{%.+}}, i{{64|32}} {{%.+}})
    // LAMBDA: {{%.+}} = alloca i{{[0-9]+}},
    // LAMBDA: {{%.+}} = alloca i{{[0-9]+}},
    // LAMBDA: {{%.+}} = alloca i{{[0-9]+}},
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_private_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_private_codegen.cpp
@ -282,7 +282,7 @@ int main() {
 // HCHECK: ret

 // HCHECK: define {{.*}}void @[[TOFFL1]]()
-// TCHECK: define void @[[TOFFL1:.+]]()
+// TCHECK: define weak void @[[TOFFL1:.+]]()
 // CHECK: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 0, {{.+}} @[[TOUTL1:.+]] to {{.+}})
 // CHECK: ret void

--- a/clang/test/OpenMP/target_teams_distribute_simd_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_simd_codegen.cpp
@ -68,13 +68,13 @@
 // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [5 x i64] [i64 288, i64 288, i64 288, i64 288, i64 547]
 // CHECK-DAG: [[SIZET7:@.+]] = private unnamed_addr constant [3 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 40]
 // CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [3 x i64] [i64 288, i64 288, i64 547]
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
--- a/clang/test/OpenMP/target_teams_distribute_simd_codegen_registration.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_simd_codegen_registration.cpp
@ -81,45 +81,45 @@

 // We have 7 target regions

-// CHECK-DAG: {{@.+}} = private constant i8 0
-// TCHECK-NOT: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
+// TCHECK-NOT: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]
-// CHECK-DAG: {{@.+}} = private constant i8 0
+// CHECK-DAG: {{@.+}} = weak constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 288]

-// CHECK-NTARGET-NOT: private constant i8 0
+// CHECK-NTARGET-NOT: weak constant i8 0
 // CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i

 // CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
@ -349,18 +349,18 @@ struct ST {
 //CHECK-DAG: define internal void @[[NAME12]](
 //CHECK-DAG: call void @[[NAME12]](

-//TCHECK-DAG: define void @[[NAME1]](
-//TCHECK-DAG: define void @[[NAME2]](
-//TCHECK-DAG: define void @[[NAME3]](
-//TCHECK-DAG: define void @[[NAME4]](
-//TCHECK-DAG: define void @[[NAME5]](
-//TCHECK-DAG: define void @[[NAME6]](
-//TCHECK-DAG: define void @[[NAME7]](
-//TCHECK-DAG: define void @[[NAME8]](
-//TCHECK-DAG: define void @[[NAME9]](
-//TCHECK-DAG: define void @[[NAME10]](
-//TCHECK-DAG: define void @[[NAME11]](
-//TCHECK-DAG: define void @[[NAME12]](
+//TCHECK-DAG: define weak void @[[NAME1]](
+//TCHECK-DAG: define weak void @[[NAME2]](
+//TCHECK-DAG: define weak void @[[NAME3]](
+//TCHECK-DAG: define weak void @[[NAME4]](
+//TCHECK-DAG: define weak void @[[NAME5]](
+//TCHECK-DAG: define weak void @[[NAME6]](
+//TCHECK-DAG: define weak void @[[NAME7]](
+//TCHECK-DAG: define weak void @[[NAME8]](
+//TCHECK-DAG: define weak void @[[NAME9]](
+//TCHECK-DAG: define weak void @[[NAME10]](
+//TCHECK-DAG: define weak void @[[NAME11]](
+//TCHECK-DAG: define weak void @[[NAME12]](

 // CHECK-NTARGET-NOT: __tgt_target
 // CHECK-NTARGET-NOT: __tgt_register_lib
--- a/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp
@ -49,7 +49,7 @@

 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 0, i[[SZ]] 4]
 // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 288]
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
--- a/clang/test/OpenMP/target_teams_num_teams_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_num_teams_codegen.cpp
@ -53,12 +53,12 @@

 // We have 6 target regions

-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
--- a/clang/test/OpenMP/target_teams_thread_limit_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_thread_limit_codegen.cpp
@ -53,12 +53,12 @@

 // We have 6 target regions

-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
-// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0
+// CHECK-DAG: @{{.*}} = weak constant i8 0

 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]