[OPENMP, NVPTX] Added support for L2 parallelism.
Added initial codegen for level 2, 3 etc. parallelism. Currently, all
the second, the third etc. parallel regions will run sequentially.
llvm-svn: 331642
diff --git a/clang/test/OpenMP/nvptx_target_codegen.cpp b/clang/test/OpenMP/nvptx_target_codegen.cpp
index 70f3973..718c650 100644
--- a/clang/test/OpenMP/nvptx_target_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_codegen.cpp
@@ -9,15 +9,17 @@
#define HEADER
// Check that the execution mode of all 6 target regions is set to Generic Mode.
-// CHECK-DAG: {{@__omp_offloading_.+l100}}_exec_mode = weak constant i8 1
-// CHECK-DAG: {{@__omp_offloading_.+l177}}_exec_mode = weak constant i8 1
-// CHECK-DAG: {{@__omp_offloading_.+l287}}_exec_mode = weak constant i8 1
-// CHECK-DAG: {{@__omp_offloading_.+l324}}_exec_mode = weak constant i8 1
-// CHECK-DAG: {{@__omp_offloading_.+l342}}_exec_mode = weak constant i8 1
-// CHECK-DAG: {{@__omp_offloading_.+l307}}_exec_mode = weak constant i8 1
+// CHECK-DAG: {{@__omp_offloading_.+l102}}_exec_mode = weak constant i8 1
+// CHECK-DAG: {{@__omp_offloading_.+l179}}_exec_mode = weak constant i8 1
+// CHECK-DAG: {{@__omp_offloading_.+l289}}_exec_mode = weak constant i8 1
+// CHECK-DAG: {{@__omp_offloading_.+l326}}_exec_mode = weak constant i8 1
+// CHECK-DAG: {{@__omp_offloading_.+l344}}_exec_mode = weak constant i8 1
+// CHECK-DAG: {{@__omp_offloading_.+l309}}_exec_mode = weak constant i8 1
__thread int id;
+int baz(int f);
+
template<typename tx, typename ty>
struct TT{
tx X;
@@ -33,7 +35,7 @@
double cn[5][n];
TT<long long, char> d;
- // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l100}}_worker()
+ // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l102}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -64,7 +66,7 @@
// CHECK: [[EXIT]]
// CHECK: ret void
- // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+foo.+l100]]()
+ // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+foo.+l102]]()
// CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
// CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
// CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
@@ -106,7 +108,7 @@
{
}
- // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l177}}_worker()
+ // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l179}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -137,7 +139,7 @@
// CHECK: [[EXIT]]
// CHECK: ret void
- // CHECK: define {{.*}}void [[T2:@__omp_offloading_.+foo.+l177]](i[[SZ:32|64]] [[ARG1:%[a-zA-Z_]+]], i[[SZ:32|64]] [[ID:%[a-zA-Z_]+]])
+ // CHECK: define {{.*}}void [[T2:@__omp_offloading_.+foo.+l179]](i[[SZ:32|64]] [[ARG1:%[a-zA-Z_]+]], i[[SZ:32|64]] [[ID:%[a-zA-Z_]+]])
// CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]],
// CHECK: store i[[SZ]] [[ARG1]], i[[SZ]]* [[AA_ADDR]],
// CHECK: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16*
@@ -180,7 +182,7 @@
id = aa;
}
- // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l287}}_worker()
+ // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l289}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -211,7 +213,7 @@
// CHECK: [[EXIT]]
// CHECK: ret void
- // CHECK: define {{.*}}void [[T3:@__omp_offloading_.+foo.+l287]](i[[SZ]]
+ // CHECK: define {{.*}}void [[T3:@__omp_offloading_.+foo.+l289]](i[[SZ]]
// Create local storage for each capture.
// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
// CHECK: [[LOCAL_B:%.+]] = alloca [10 x float]*
@@ -343,6 +345,7 @@
{
this->a = (double)b + 1.5;
c[1][1] = ++a;
+ baz(a);
}
return c[1][1] + (int)b;
@@ -364,7 +367,13 @@
return a;
}
- // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+static.+324}}_worker()
+int baz(int f) {
+#pragma omp parallel
+ f = 2;
+ return f;
+}
+
+ // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+static.+326}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -395,7 +404,7 @@
// CHECK: [[EXIT]]
// CHECK: ret void
- // CHECK: define {{.*}}void [[T4:@__omp_offloading_.+static.+l324]](i[[SZ]]
+ // CHECK: define {{.*}}void [[T4:@__omp_offloading_.+static.+l326]](i[[SZ]]
// Create local storage for each capture.
// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
// CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]
@@ -450,9 +459,10 @@
- // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+S1.+l342}}_worker()
+ // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+S1.+l344}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
+ // CHECK: [[GTID:%.+]] = call i32 @__kmpc_global_thread_num(%struct.ident_t*
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
// CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
// CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
@@ -469,6 +479,8 @@
// CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
//
// CHECK: [[EXEC_PARALLEL]]
+ // CHECK: [[WORK_FN:%.+]] = bitcast i8* [[WORK]] to void (i16, i32)*
+ // CHECK: call void [[WORK_FN]](i16 0, i32 [[GTID]])
// CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
//
// CHECK: [[TERM_PARALLEL]]
@@ -481,7 +493,7 @@
// CHECK: [[EXIT]]
// CHECK: ret void
- // CHECK: define {{.*}}void [[T5:@__omp_offloading_.+S1.+l342]](
+ // CHECK: define {{.*}}void [[T5:@__omp_offloading_.+S1.+l344]](
// Create local storage for each capture.
// CHECK: [[LOCAL_THIS:%.+]] = alloca [[S1:%struct.*]]*
// CHECK: [[LOCAL_B:%.+]] = alloca i[[SZ]]
@@ -528,6 +540,7 @@
// CHECK-64-DAG:load i32, i32* [[REF_B]]
// CHECK-32-DAG:load i32, i32* [[LOCAL_B]]
// CHECK-DAG: getelementptr inbounds i16, i16* [[REF_C]], i[[SZ]] %{{.+}}
+ // CHECK: call i32 [[BAZ:@.*baz.*]](i32 %
// CHECK: br label {{%?}}[[TERMINATE:.+]]
//
// CHECK: [[TERMINATE]]
@@ -538,9 +551,48 @@
// CHECK: [[EXIT]]
// CHECK: ret void
+ // CHECK: define i32 [[BAZ]](i32 [[F:%.*]])
+ // CHECK: [[ZERO_ADDR:%.+]] = alloca i32,
+ // CHECK: [[GTID:%.+]] = call i32 @__kmpc_global_thread_num(%struct.ident_t*
+ // CHECK: [[GTID_ADDR:%.+]] = alloca i32,
+ // CHECK: store i32 0, i32* [[ZERO_ADDR]]
+ // CHECK: [[PTR:%.+]] = call i8* @__kmpc_data_sharing_push_stack(i{{64|32}} 4, i16 0)
+ // 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 [[F]], i32* [[F_PTR]],
+ // CHECK: store i32 [[GTID]], i32* [[GTID_ADDR]],
+ // CHECK: icmp eq i32
+ // CHECK: br i1
+
+ // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* @{{.+}} to i8*), i16 1)
+ // CHECK: call void @__kmpc_begin_sharing_variables(i8*** [[SHARED_PTR:%.+]], i{{64|32}} 1)
+ // CHECK: [[SHARED:%.+]] = load i8**, i8*** [[SHARED_PTR]],
+ // CHECK: [[REF:%.+]] = getelementptr inbounds i8*, i8** [[SHARED]], i{{64|32}} 0
+ // CHECK: [[F_REF:%.+]] = bitcast i32* [[F_PTR]] to i8*
+ // CHECK: store i8* [[F_REF]], i8** [[REF]],
+ // CHECK: call void @llvm.nvvm.barrier0()
+ // CHECK: call void @llvm.nvvm.barrier0()
+ // CHECK: call void @__kmpc_end_sharing_variables()
+ // CHECK: br label
+
+ // CHECK: [[RES:%.+]] = call i16 @__kmpc_parallel_level(%struct.ident_t* @{{.+}}, i32 [[GTID]])
+ // CHECK: icmp ne i16 [[RES]], 0
+ // CHECK: br i1
+
+ // CHECK: call void @__kmpc_serialized_parallel(%struct.ident_t* @{{.+}}, i32 [[GTID]])
+ // CHECK: call void @__omp_outlined__(i32* [[GTID_ADDR]], i32* [[ZERO_ADDR]], i32* [[F_PTR]])
+ // CHECK: call void @__kmpc_end_serialized_parallel(%struct.ident_t* @{{.+}}, i32 [[GTID]])
+ // CHECK: br label
+
+ // CHECK: call void @__omp_outlined__(i32* [[GTID_ADDR]], i32* [[ZERO_ADDR]], i32* [[F_PTR]])
+ // CHECK: br label
+
+ // CHECK: [[RES:%.+]] = load i32, i32* [[F_PTR]],
+ // CHECK: call void @__kmpc_data_sharing_pop_stack(i8* [[PTR]])
+ // CHECK: ret i32 [[RES]]
- // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l307}}_worker()
+ // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l309}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -571,7 +623,7 @@
// CHECK: [[EXIT]]
// CHECK: ret void
- // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l307]](i[[SZ]]
+ // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l309]](i[[SZ]]
// Create local storage for each capture.
// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
// CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]