diff options
author | Alexey Bataev <a.bataev@hotmail.com> | 2017-12-13 21:04:20 +0000 |
---|---|---|
committer | Alexey Bataev <a.bataev@hotmail.com> | 2017-12-13 21:04:20 +0000 |
commit | e8ca256e5e3f0705a447f60d297e97457fd28438 (patch) | |
tree | e1e8820ecb468dd8333c625ba61848faa12c37da /test/OpenMP | |
parent | 95387d6350f7fda33efed7b94757d580c60bd92d (diff) |
[OPENMP] Add codegen for `nowait` clause in target directives.
Added basic codegen for `nowait` clauses in target-based directives.
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@320613 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/OpenMP')
-rw-r--r-- | test/OpenMP/target_codegen.cpp | 4 | ||||
-rw-r--r-- | test/OpenMP/target_parallel_codegen.cpp | 4 | ||||
-rw-r--r-- | test/OpenMP/target_parallel_for_codegen.cpp | 4 | ||||
-rw-r--r-- | test/OpenMP/target_parallel_for_simd_codegen.cpp | 4 | ||||
-rw-r--r-- | test/OpenMP/target_simd_codegen.cpp | 4 | ||||
-rw-r--r-- | test/OpenMP/target_teams_codegen.cpp | 4 | ||||
-rw-r--r-- | test/OpenMP/target_teams_distribute_codegen.cpp | 4 | ||||
-rw-r--r-- | test/OpenMP/target_teams_distribute_simd_codegen.cpp | 4 |
8 files changed, 16 insertions, 16 deletions
diff --git a/test/OpenMP/target_codegen.cpp b/test/OpenMP/target_codegen.cpp index 91a0b83e6d..20bc96b84c 100644 --- a/test/OpenMP/target_codegen.cpp +++ b/test/OpenMP/target_codegen.cpp @@ -111,7 +111,7 @@ int foo(int n) { // CHECK-DAG: [[ADD:%.+]] = add nsw i32 // CHECK-DAG: [[DEVICE:%.+]] = sext i32 [[ADD]] to i64 - // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target(i64 [[DEVICE]], i8* @{{[^,]+}}, i32 2, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* getelementptr inbounds ([2 x i[[SZ]]], [2 x i[[SZ]]]* [[SIZET]], i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPT]], i32 0, i32 0) + // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_nowait(i64 [[DEVICE]], i8* @{{[^,]+}}, i32 2, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* getelementptr inbounds ([2 x i[[SZ]]], [2 x i[[SZ]]]* [[SIZET]], i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPT]], i32 0, i32 0) // CHECK-DAG: [[BPR]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP:%[^,]+]], i32 0, i32 0 // CHECK-DAG: [[PR]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P:%[^,]+]], i32 0, i32 0 @@ -134,7 +134,7 @@ int foo(int n) { // CHECK: call void [[HVT0_:@.+]](i[[SZ]]* [[BP0]], i[[SZ]] [[BP1]]) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] - #pragma omp target device(global + a) + #pragma omp target device(global + a) nowait { static int local1; *plocal = global; diff --git a/test/OpenMP/target_parallel_codegen.cpp b/test/OpenMP/target_parallel_codegen.cpp index ce280b4549..b9d00da02f 100644 --- a/test/OpenMP/target_parallel_codegen.cpp +++ b/test/OpenMP/target_parallel_codegen.cpp @@ -93,14 +93,14 @@ int foo(int n) { double cn[5][n]; TT<long long, char> d; - // CHECK: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i64* null, i32 1, i32 0) + // CHECK: [[RET:%.+]] = call i32 @__tgt_target_teams_nowait(i64 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i64* null, i32 1, i32 0) // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT0:@.+]]() // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] - #pragma omp target parallel + #pragma omp target parallel nowait { } diff --git a/test/OpenMP/target_parallel_for_codegen.cpp b/test/OpenMP/target_parallel_for_codegen.cpp index 55dd4d2a57..58ac10bf9b 100644 --- a/test/OpenMP/target_parallel_for_codegen.cpp +++ b/test/OpenMP/target_parallel_for_codegen.cpp @@ -116,7 +116,7 @@ int foo(int n) { a += 1; } - // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET2]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT2]], i32 0, i32 0), i32 1, i32 0) + // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_nowait(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET2]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT2]], i32 0, i32 0), i32 1, i32 0) // CHECK-DAG: [[BP]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR:%[^,]+]], i32 0, i32 0 // CHECK-DAG: [[P]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR:%[^,]+]], i32 0, i32 0 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR]], i32 0, i32 0 @@ -145,7 +145,7 @@ int foo(int n) { // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] int lin = 12; - #pragma omp target parallel for if(target: 1) linear(lin, a : get_val()) + #pragma omp target parallel for if(target: 1) linear(lin, a : get_val()) nowait for (unsigned long long it = 2000; it >= 600; it-=400) { aa += 1; } diff --git a/test/OpenMP/target_parallel_for_simd_codegen.cpp b/test/OpenMP/target_parallel_for_simd_codegen.cpp index 1ed8d3a8bf..d168214df3 100644 --- a/test/OpenMP/target_parallel_for_simd_codegen.cpp +++ b/test/OpenMP/target_parallel_for_simd_codegen.cpp @@ -95,14 +95,14 @@ int foo(int n) { double cn[5][n]; TT<long long, char> d; - // CHECK: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i64* null, i32 1, i32 0) + // CHECK: [[RET:%.+]] = call i32 @__tgt_target_teams_nowait(i64 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i64* null, i32 1, i32 0) // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT0:@.+]]() // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] - #pragma omp target parallel for simd + #pragma omp target parallel for simd nowait for (int i = 3; i < 32; i += 5) { } diff --git a/test/OpenMP/target_simd_codegen.cpp b/test/OpenMP/target_simd_codegen.cpp index edabfc5eaa..d57e381855 100644 --- a/test/OpenMP/target_simd_codegen.cpp +++ b/test/OpenMP/target_simd_codegen.cpp @@ -92,14 +92,14 @@ int foo(int n) { double cn[5][n]; TT<long long, char> d; - // CHECK: [[RET:%.+]] = call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i64* null) + // CHECK: [[RET:%.+]] = call i32 @__tgt_target_nowait(i64 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i64* null) // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT0:@.+]]() // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] - #pragma omp target simd + #pragma omp target simd nowait for (int i = 3; i < 32; i += 5) { } diff --git a/test/OpenMP/target_teams_codegen.cpp b/test/OpenMP/target_teams_codegen.cpp index 284a147f28..e6c250178d 100644 --- a/test/OpenMP/target_teams_codegen.cpp +++ b/test/OpenMP/target_teams_codegen.cpp @@ -97,7 +97,7 @@ int foo(int n) { double cn[5][n]; TT<long long, char> d; - // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT]], i32 0, i32 0), i32 {{[^,]+}}, i32 {{[^)]+}}) + // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_nowait(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT]], i32 0, i32 0), i32 {{[^,]+}}, i32 {{[^)]+}}) // CHECK-DAG: [[BP]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR:%[^,]+]], i32 0, i32 0 // CHECK-DAG: [[P]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR:%[^,]+]], i32 0, i32 0 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR]], i32 0, i32 [[IDX0:[0-9]+]] @@ -124,7 +124,7 @@ int foo(int n) { // CHECK: call void [[HVT0:@.+]](i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^)]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] - #pragma omp target teams num_teams(a) thread_limit(a) firstprivate(aa) + #pragma omp target teams num_teams(a) thread_limit(a) firstprivate(aa) nowait { } diff --git a/test/OpenMP/target_teams_distribute_codegen.cpp b/test/OpenMP/target_teams_distribute_codegen.cpp index b3eae6c4bd..17baa603d1 100644 --- a/test/OpenMP/target_teams_distribute_codegen.cpp +++ b/test/OpenMP/target_teams_distribute_codegen.cpp @@ -97,7 +97,7 @@ int foo(int n) { double cn[5][n]; TT<long long, char> d; - // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT]], i32 0, i32 0), i32 {{[^,]+}}, i32 {{[^)]+}}) + // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_nowait(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT]], i32 0, i32 0), i32 {{[^,]+}}, i32 {{[^)]+}}) // CHECK-DAG: [[BP]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR:%[^,]+]], i32 0, i32 0 // CHECK-DAG: [[P]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR:%[^,]+]], i32 0, i32 0 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR]], i32 0, i32 [[IDX0:[0-9]+]] @@ -124,7 +124,7 @@ int foo(int n) { // CHECK: call void [[HVT0:@.+]](i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^)]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] - #pragma omp target teams distribute num_teams(a) thread_limit(a) firstprivate(aa) + #pragma omp target teams distribute num_teams(a) thread_limit(a) firstprivate(aa) nowait for (int i = 0; i < 10; ++i) { } diff --git a/test/OpenMP/target_teams_distribute_simd_codegen.cpp b/test/OpenMP/target_teams_distribute_simd_codegen.cpp index b7f031fce7..7fb76d9198 100644 --- a/test/OpenMP/target_teams_distribute_simd_codegen.cpp +++ b/test/OpenMP/target_teams_distribute_simd_codegen.cpp @@ -97,7 +97,7 @@ int foo(int n) { double cn[5][n]; TT<long long, char> d; - // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT]], i32 0, i32 0), i32 {{[^,]+}}, i32 {{[^)]+}}) + // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_nowait(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT]], i32 0, i32 0), i32 {{[^,]+}}, i32 {{[^)]+}}) // CHECK-DAG: [[BP]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR:%[^,]+]], i32 0, i32 0 // CHECK-DAG: [[P]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR:%[^,]+]], i32 0, i32 0 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR]], i32 0, i32 [[IDX0:[0-9]+]] @@ -124,7 +124,7 @@ int foo(int n) { // CHECK: call void [[HVT0:@.+]](i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^)]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] - #pragma omp target teams distribute simd num_teams(a) thread_limit(a) firstprivate(aa) simdlen(16) + #pragma omp target teams distribute simd num_teams(a) thread_limit(a) firstprivate(aa) simdlen(16) nowait for (int i = 0; i < 10; ++i) { } |