diff options
Diffstat (limited to 'clang/test/OpenMP/nvptx_target_codegen.cpp')
| -rw-r--r-- | clang/test/OpenMP/nvptx_target_codegen.cpp | 88 |
1 files changed, 70 insertions, 18 deletions
diff --git a/clang/test/OpenMP/nvptx_target_codegen.cpp b/clang/test/OpenMP/nvptx_target_codegen.cpp index 70f3973acd4..718c650bec6 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 @@ int foo(int n) { 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 @@ int foo(int n) { // 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 @@ int foo(int n) { { } - // 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 @@ int foo(int n) { // 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 @@ int foo(int n) { 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 @@ int foo(int n) { // 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 @@ struct S1 { { this->a = (double)b + 1.5; c[1][1] = ++a; + baz(a); } return c[1][1] + (int)b; @@ -364,7 +367,13 @@ int bar(int n){ 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 @@ int bar(int n){ // 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 @@ int bar(int n){ - // 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 @@ int bar(int n){ // 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 @@ int bar(int n){ // 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 @@ int bar(int n){ // 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 @@ int bar(int n){ // 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 @@ int bar(int n){ // 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]] |
