[OPENMP] ThreadId in serialized parallel regions is 0.

The first argument for the parallel outlined functions, called as
serialized parallel regions, should be a pointer to the global thread id
that always is 0.

llvm-svn: 337957

2 files changed, 12 insertions, 12 deletions

diff --git a/clang/test/OpenMP/nvptx_target_codegen.cpp b/clang/test/OpenMP/nvptx_target_codegen.cpp
index d4d9941bc43..5f9b3bd3280 100644
--- a/clang/test/OpenMP/nvptx_target_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_codegen.cpp
@@ -562,7 +562,6 @@ int baz(int f, double &a) {
   // CHECK: [[REC_ADDR:%.+]] = bitcast i8* [[PTR]] to %struct._globalized_locals_ty*
   // CHECK: [[F_PTR:%.+]] = getelementptr inbounds %struct._globalized_locals_ty, %struct._globalized_locals_ty* [[REC_ADDR]], i32 0, i32 0
   // CHECK: store i32 %{{.+}}, i32* [[F_PTR]],
-  // CHECK: store i32 [[GTID]], i32* [[GTID_ADDR]],
 
   // CHECK: [[RES:%.+]] = call i8 @__kmpc_is_spmd_exec_mode()
   // CHECK: icmp ne i8 [[RES]], 0
@@ -573,7 +572,7 @@ int baz(int f, double &a) {
   // CHECK: br i1
 
   // CHECK: call void @__kmpc_serialized_parallel(%struct.ident_t* @{{.+}}, i32 [[GTID]])
-  // CHECK: call void [[OUTLINED:@.+]](i32* [[GTID_ADDR]], i32* [[ZERO_ADDR]], i32* [[F_PTR]], double* %{{.+}})
+  // CHECK: call void [[OUTLINED:@.+]](i32* [[ZERO_ADDR]], i32* [[ZERO_ADDR]], i32* [[F_PTR]], double* %{{.+}})
   // CHECK: call void @__kmpc_end_serialized_parallel(%struct.ident_t* @{{.+}}, i32 [[GTID]])
   // CHECK: br label
 
@@ -591,6 +590,7 @@ int baz(int f, double &a) {
   // CHECK: call void @__kmpc_end_sharing_variables()
   // CHECK: br label
 
+  // CHECK: store i32 [[GTID]], i32* [[GTID_ADDR]],
   // CHECK: call void [[OUTLINED]](i32* [[GTID_ADDR]], i32* [[ZERO_ADDR]], i32* [[F_PTR]], double* %{{.+}})
   // CHECK: br label
 
diff --git a/clang/test/OpenMP/parallel_if_codegen.cpp b/clang/test/OpenMP/parallel_if_codegen.cpp
index e59415a9d29..49e4bd402ce 100644
--- a/clang/test/OpenMP/parallel_if_codegen.cpp
+++ b/clang/test/OpenMP/parallel_if_codegen.cpp
@@ -29,12 +29,12 @@ void gtid_test() {
 }
 
 // CHECK: define internal {{.*}}void [[GTID_TEST_REGION1]](i{{.+}}* noalias [[GTID_PARAM:%.+]], i32* noalias
+// CHECK: store i32 0, i32* [[ZERO_ADDR:%.+]],
 // CHECK: store i{{[0-9]+}}* [[GTID_PARAM]], i{{[0-9]+}}** [[GTID_ADDR_REF:%.+]],
 // CHECK: [[GTID_ADDR:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[GTID_ADDR_REF]]
 // CHECK: [[GTID:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[GTID_ADDR]]
 // CHECK: call {{.*}}void @__kmpc_serialized_parallel(%{{.+}}* @{{.+}}, i{{.+}} [[GTID]])
-// CHECK: [[GTID_ADDR:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[GTID_ADDR_REF]]
-// CHECK: call void [[GTID_TEST_REGION2:@.+]](i{{[0-9]+}}* [[GTID_ADDR]]
+// CHECK: call void [[GTID_TEST_REGION2:@.+]](i{{[0-9]+}}* [[ZERO_ADDR]]
 // CHECK: call {{.*}}void @__kmpc_end_serialized_parallel(%{{.+}}* @{{.+}}, i{{.+}} [[GTID]])
 // CHECK: ret void
 
@@ -56,12 +56,13 @@ int tmain(T Arg) {
 // CHECK-LABEL: define {{.*}}i{{[0-9]+}} @main()
 int main() {
 // CHECK: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num(
+// CHECK: store i32 0, i32* [[ZERO_ADDR2:%.+]],
+// CHECK: store i32 0, i32* [[ZERO_ADDR1:%.+]],
 // CHECK: call {{.*}}void {{.+}} @__kmpc_fork_call(%{{.+}}* @{{.+}}, i{{.+}} 0, void {{.+}}* [[CAP_FN4:@.+]] to void
 #pragma omp parallel if (true)
   fn4();
 // CHECK: call {{.*}}void @__kmpc_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]])
-// CHECK: store i32 [[GTID]], i32* [[GTID_ADDR:%.+]],
-// CHECK: call void [[CAP_FN5:@.+]](i32* [[GTID_ADDR]],
+// CHECK: call void [[CAP_FN5:@.+]](i32* [[ZERO_ADDR1]], i32* [[ZERO_ADDR1]])
 // CHECK: call {{.*}}void @__kmpc_end_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]])
 #pragma omp parallel if (false)
   fn5();
@@ -72,8 +73,7 @@ int main() {
 // CHECK: br label %[[OMP_END:.+]]
 // CHECK: [[OMP_ELSE]]
 // CHECK: call {{.*}}void @__kmpc_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]])
-// CHECK: store i32 [[GTID]], i32* [[GTID_ADDR:%.+]],
-// CHECK: call void [[CAP_FN6]](i32* [[GTID_ADDR]],
+// CHECK: call void [[CAP_FN6]](i32* [[ZERO_ADDR2]], i32* [[ZERO_ADDR2]])
 // CHECK: call {{.*}}void @__kmpc_end_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]])
 // CHECK: br label %[[OMP_END]]
 // CHECK: [[OMP_END]]
@@ -97,10 +97,11 @@ int main() {
 
 // CHECK-LABEL: define {{.+}} @{{.+}}tmain
 // CHECK: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num(
+// CHECK: store i32 0, i32* [[ZERO_ADDR2:%.+]],
+// CHECK: store i32 0, i32* [[ZERO_ADDR1:%.+]],
 // CHECK: call {{.*}}void {{.+}} @__kmpc_fork_call(%{{.+}}* @{{.+}}, i{{.+}} 0, void {{.+}}* [[CAP_FN1:@.+]] to void
 // CHECK: call {{.*}}void @__kmpc_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]])
-// CHECK: store i32 [[GTID]], i32* [[GTID_ADDR:%.+]],
-// CHECK: call void [[CAP_FN2:@.+]](i32* [[GTID_ADDR]],
+// CHECK: call void [[CAP_FN2:@.+]](i32* [[ZERO_ADDR1]], i32* [[ZERO_ADDR1]])
 // CHECK: call {{.*}}void @__kmpc_end_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]])
 // CHECK: br i1 %{{.+}}, label %[[OMP_THEN:.+]], label %[[OMP_ELSE:.+]]
 // CHECK: [[OMP_THEN]]
@@ -108,8 +109,7 @@ int main() {
 // CHECK: br label %[[OMP_END:.+]]
 // CHECK: [[OMP_ELSE]]
 // CHECK: call {{.*}}void @__kmpc_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]])
-// CHECK: store i32 [[GTID]], i32* [[GTID_ADDR:%.+]],
-// CHECK: call void [[CAP_FN3]](i32* [[GTID_ADDR]],
+// CHECK: call void [[CAP_FN3]](i32* [[ZERO_ADDR2]], i32* [[ZERO_ADDR2]])
 // CHECK: call {{.*}}void @__kmpc_end_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]])
 // CHECK: br label %[[OMP_END]]
 // CHECK: [[OMP_END]]