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"});
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);
}
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)
{
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();
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() {
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* %
/// ========= 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
// 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"
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* %
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]],
// 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* %{{.*}})
// 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(
// 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]]
// 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"
//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
// 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]]
// 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]],
// 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]],
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]+}}]*,
// 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]+}},
// 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]+}},
// 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]]
// 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"
//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
// 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)* {{[^)]+}})
// 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)* {{[^)]+}})
// 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)*
// 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]]
// 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]]
// 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"
//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
// 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)* {{[^)]+}})
// 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)* {{[^)]+}})
// 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)*
// 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]]
// 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]]
// 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"
//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
// 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]]
// 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]]
// 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]]
{
}
- // 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]],
}
// 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]+}},
// 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]],
// 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]],
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]+}},
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]+}},
}
// 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]+}}],
{
}
- // 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]],
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]],
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]],
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]+}}*,
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]+}},
}
// 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]+}}]*,
// 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]]
// 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"
//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
// 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]]
// 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]]
// 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"
//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
// 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]]
// 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]]
// 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"
//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
// 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]]
// 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}},
// 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
// 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]]
#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]+}},
// 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
// 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}},
// 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
// 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"
//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
// 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]]
#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]+}},
// 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
// 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]]
// 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"
//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
// 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]]
// 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]]
// 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]]