diff options
Diffstat (limited to 'test/OpenMP/nvptx_target_codegen.cpp')
| -rw-r--r-- | test/OpenMP/nvptx_target_codegen.cpp | 39 |
1 files changed, 21 insertions, 18 deletions
diff --git a/test/OpenMP/nvptx_target_codegen.cpp b/test/OpenMP/nvptx_target_codegen.cpp index dab0c5a082a52..d7ce7dc7d1060 100644 --- a/test/OpenMP/nvptx_target_codegen.cpp +++ b/test/OpenMP/nvptx_target_codegen.cpp @@ -9,12 +9,14 @@ #define HEADER // Check that the execution mode of all 6 target regions is set to Generic Mode. -// CHECK-DAG: {{@__omp_offloading_.+l98}}_exec_mode = weak constant i8 1 -// CHECK-DAG: {{@__omp_offloading_.+l175}}_exec_mode = weak constant i8 1 -// CHECK-DAG: {{@__omp_offloading_.+l284}}_exec_mode = weak constant i8 1 -// CHECK-DAG: {{@__omp_offloading_.+l321}}_exec_mode = weak constant i8 1 -// CHECK-DAG: {{@__omp_offloading_.+l339}}_exec_mode = weak constant i8 1 -// CHECK-DAG: {{@__omp_offloading_.+l304}}_exec_mode = weak constant i8 1 +// 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 + +__thread int id; template<typename tx, typename ty> struct TT{ @@ -31,7 +33,7 @@ int foo(int n) { double cn[5][n]; TT<long long, char> d; - // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l98}}_worker() + // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l100}}_worker() // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8, // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*, // CHECK: store i8* null, i8** [[OMP_WORK_FN]], @@ -62,7 +64,7 @@ int foo(int n) { // CHECK: [[EXIT]] // CHECK: ret void - // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+foo.+l98]]() + // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+foo.+l100]]() // 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() @@ -104,7 +106,7 @@ int foo(int n) { { } - // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l175}}_worker() + // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l177}}_worker() // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8, // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*, // CHECK: store i8* null, i8** [[OMP_WORK_FN]], @@ -135,7 +137,7 @@ int foo(int n) { // CHECK: [[EXIT]] // CHECK: ret void - // CHECK: define {{.*}}void [[T2:@__omp_offloading_.+foo.+l175]](i[[SZ:32|64]] [[ARG1:%[a-zA-Z_]+]]) + // 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: [[AA_ADDR:%.+]] = alloca i[[SZ]], // CHECK: store i[[SZ]] [[ARG1]], i[[SZ]]* [[AA_ADDR]], // CHECK: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16* @@ -175,9 +177,10 @@ int foo(int n) { #pragma omp target if(1) { aa += 1; + id = aa; } - // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l284}}_worker() + // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l287}}_worker() // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8, // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*, // CHECK: store i8* null, i8** [[OMP_WORK_FN]], @@ -208,7 +211,7 @@ int foo(int n) { // CHECK: [[EXIT]] // CHECK: ret void - // CHECK: define {{.*}}void [[T3:@__omp_offloading_.+foo.+l284]](i[[SZ]] + // CHECK: define {{.*}}void [[T3:@__omp_offloading_.+foo.+l287]](i[[SZ]] // Create local storage for each capture. // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]] // CHECK: [[LOCAL_B:%.+]] = alloca [10 x float]* @@ -361,7 +364,7 @@ int bar(int n){ return a; } - // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+static.+321}}_worker() + // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+static.+324}}_worker() // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8, // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*, // CHECK: store i8* null, i8** [[OMP_WORK_FN]], @@ -392,7 +395,7 @@ int bar(int n){ // CHECK: [[EXIT]] // CHECK: ret void - // CHECK: define {{.*}}void [[T4:@__omp_offloading_.+static.+l321]](i[[SZ]] + // CHECK: define {{.*}}void [[T4:@__omp_offloading_.+static.+l324]](i[[SZ]] // Create local storage for each capture. // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]] // CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]] @@ -447,7 +450,7 @@ int bar(int n){ - // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+S1.+l339}}_worker() + // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+S1.+l342}}_worker() // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8, // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*, // CHECK: store i8* null, i8** [[OMP_WORK_FN]], @@ -478,7 +481,7 @@ int bar(int n){ // CHECK: [[EXIT]] // CHECK: ret void - // CHECK: define {{.*}}void [[T5:@__omp_offloading_.+S1.+l339]]( + // CHECK: define {{.*}}void [[T5:@__omp_offloading_.+S1.+l342]]( // Create local storage for each capture. // CHECK: [[LOCAL_THIS:%.+]] = alloca [[S1:%struct.*]]* // CHECK: [[LOCAL_B:%.+]] = alloca i[[SZ]] @@ -537,7 +540,7 @@ int bar(int n){ - // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l304}}_worker() + // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l307}}_worker() // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8, // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*, // CHECK: store i8* null, i8** [[OMP_WORK_FN]], @@ -568,7 +571,7 @@ int bar(int n){ // CHECK: [[EXIT]] // CHECK: ret void - // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l304]](i[[SZ]] + // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l307]](i[[SZ]] // Create local storage for each capture. // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]] // CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]] |