jdenny updated this revision to Diff 278658.
jdenny added a comment.
Rebased onto latest D83922 <https://reviews.llvm.org/D83922>.
Implemented rejection of `present` if not OpenMP >= 5.1.
Cleaned up tests some.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D83061/new/
https://reviews.llvm.org/D83061
Files:
clang/include/clang/AST/OpenMPClause.h
clang/include/clang/Basic/DiagnosticParseKinds.td
clang/include/clang/Basic/OpenMPKinds.def
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/Parse/ParseOpenMP.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/target_data_ast_print.cpp
clang/test/OpenMP/target_data_codegen.cpp
clang/test/OpenMP/target_enter_data_codegen.cpp
clang/test/OpenMP/target_map_codegen.cpp
clang/test/OpenMP/target_map_messages.cpp
clang/test/OpenMP/target_parallel_for_map_messages.cpp
clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
clang/test/OpenMP/target_parallel_map_messages.cpp
clang/test/OpenMP/target_simd_map_messages.cpp
clang/test/OpenMP/target_teams_distribute_map_messages.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
clang/test/OpenMP/target_teams_map_messages.cpp
Index: clang/test/OpenMP/target_teams_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_map_messages.cpp
+++ clang/test/OpenMP/target_teams_map_messages.cpp
@@ -468,7 +468,7 @@
#pragma omp target data map(always, tofrom: x)
#pragma omp target data map(always: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
#pragma omp target data map(always, tofrom: always, tofrom, x)
#pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
foo();
@@ -543,7 +543,7 @@
#pragma omp target data map(always, tofrom: x)
#pragma omp target data map(always: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
#pragma omp target data map(always, tofrom: always, tofrom, x)
#pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
foo();
Index: clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
+++ clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
@@ -175,7 +175,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
@@ -287,7 +287,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
+++ clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
@@ -175,7 +175,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
@@ -286,7 +286,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
+++ clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
@@ -174,7 +174,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
@@ -285,7 +285,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_teams_distribute_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_map_messages.cpp
+++ clang/test/OpenMP/target_teams_distribute_map_messages.cpp
@@ -175,7 +175,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
@@ -287,7 +287,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_simd_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_simd_map_messages.cpp
+++ clang/test/OpenMP/target_simd_map_messages.cpp
@@ -169,7 +169,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
@@ -275,7 +275,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_parallel_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_parallel_map_messages.cpp
+++ clang/test/OpenMP/target_parallel_map_messages.cpp
@@ -175,7 +175,7 @@
foo();
#pragma omp target parallel map(always: x) // expected-error {{missing map type}}
foo();
-#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
foo();
#pragma omp target parallel map(always, tofrom: always, tofrom, x)
foo();
@@ -285,7 +285,7 @@
foo();
#pragma omp target parallel map(always: x) // expected-error {{missing map type}}
foo();
-#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
foo();
#pragma omp target parallel map(always, tofrom: always, tofrom, x)
foo();
Index: clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
+++ clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
@@ -175,7 +175,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
@@ -287,7 +287,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_parallel_for_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_parallel_for_map_messages.cpp
+++ clang/test/OpenMP/target_parallel_for_map_messages.cpp
@@ -175,7 +175,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
@@ -287,7 +287,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(always, tofrom: always, tofrom, x)
for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_map_messages.cpp
+++ clang/test/OpenMP/target_map_messages.cpp
@@ -1,13 +1,15 @@
-// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
// RUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s -Wno-openmp -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
// RUN: %clang_cc1 -DCCODE -verify -fopenmp-simd -ferror-limit 200 -x c %s -Wno-openmp-mapping -Wuninitialized
#ifdef CCODE
void foo(int arg) {
@@ -63,9 +65,9 @@
{}
#pragma omp target map(arg[2:2],a,d) // expected-error {{subscripted value is not an array or pointer}}
{}
- #pragma omp target map(arg,a*2) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error {{expected addressable lvalue in 'map' clause}}
+ #pragma omp target map(arg,a*2) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}}
{}
- #pragma omp target map(arg,(c+1)[2]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+ #pragma omp target map(arg,(c+1)[2]) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
{}
#pragma omp target map(arg,a[:2],d) // expected-error {{subscripted value is not an array or pointer}}
{}
@@ -116,9 +118,38 @@
{}
#pragma omp target map(close) // expected-error {{use of undeclared identifier 'close'}}
{}
+ // lt51-error@+1 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+ #pragma omp target map(present, tofrom: c,f)
+ {}
+ // lt51-error@+1 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+ #pragma omp target map(present, tofrom: c[1:2],f)
+ {}
+ // lt51-error@+1 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+ #pragma omp target map(present, tofrom: c,f[1:2])
+ {}
+ // expected-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+ // lt51-error@+1 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+ #pragma omp target map(present, tofrom: c[:],f)
+ {}
+ // expected-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+ // lt51-error@+1 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+ #pragma omp target map(present, tofrom: c,f[:])
+ {}
+ // expected-error@+1 {{use of undeclared identifier 'present'}}
+ #pragma omp target map(present)
+ {}
#pragma omp target map(close, close, tofrom: a) // expected-error {{same map type modifier has been specified more than once}}
{}
- #pragma omp target map(always, close, always, close, tofrom: a) // expected-error {{same map type modifier has been specified more than once}} expected-error {{same map type modifier has been specified more than once}}
+ #pragma omp target map(always, close, always, close, tofrom: a) // expected-error 2 {{same map type modifier has been specified more than once}}
+ {}
+ // ge51-error@+2 {{same map type modifier has been specified more than once}}
+ // lt51-error@+1 2 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+ #pragma omp target map(present, present, tofrom: a)
+ {}
+ // expected-error@+3 2 {{same map type modifier has been specified more than once}}
+ // ge51-error@+2 1 {{same map type modifier has been specified more than once}}
+ // lt51-error@+1 2 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+ #pragma omp target map(always, close, present, always, close, present, tofrom: a)
{}
#pragma omp target map( , tofrom: a) // expected-error {{missing map type modifier}}
{}
@@ -126,21 +157,21 @@
{}
#pragma omp target map( , , : a) // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}} expected-error {{missing map type}}
{}
- #pragma omp target map( d, f, bf: a) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+ #pragma omp target map( d, f, bf: a) // expected-error 2 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
{}
- #pragma omp target map( , f, : a) // expected-error {{missing map type modifier}} expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+ #pragma omp target map( , f, : a) // expected-error {{missing map type modifier}} expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper'}} expected-error {{missing map type}}
{}
#pragma omp target map(always close: a) // expected-error {{missing map type}}
{}
#pragma omp target map(always close bf: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
{}
- #pragma omp target map(always tofrom close: a) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+ #pragma omp target map(always tofrom close: a) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
{}
- #pragma omp target map(tofrom from: a) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+ #pragma omp target map(tofrom from: a) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
{}
#pragma omp target map(close bf: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
{}
- #pragma omp target map(([b[I]][bf])f) // le45-error {{expected ',' or ']' in lambda capture list}} le45-error {{expected ')'}} le45-note {{to match this '('}}
+ #pragma omp target map(([b[I]][bf])f) // lt50-error {{expected ',' or ']' in lambda capture list}} lt50-error {{expected ')'}} lt50-note {{to match this '('}}
{}
return;
}
@@ -313,7 +344,7 @@
{}
#pragma omp target map(r.S.Arr[:12])
{}
-#pragma omp target map(r.S.foo() [:12]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error {{expected addressable lvalue in 'map' clause}}
+#pragma omp target map(r.S.foo() [:12]) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}}
{}
#pragma omp target map(r.C, r.D)
{}
@@ -347,7 +378,7 @@
{}
#pragma omp target map(r.S.Ptr[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
{}
-#pragma omp target map((p + 1)->A) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map((p + 1)->A) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
{}
#pragma omp target map(u.B) // expected-error {{mapping of union members is not allowed}}
{}
@@ -467,7 +498,7 @@
T *k = &j;
T x;
T y;
- T to, tofrom, always, close;
+ T to, tofrom, always, close, present;
const T (&l)[5] = da;
#pragma omp target map // expected-error {{expected '(' after 'map'}}
{}
@@ -493,7 +524,7 @@
foo();
#pragma omp target map(T) // expected-error {{'T' does not refer to a value}}
foo();
-#pragma omp target map(I) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error 2 {{expected addressable lvalue in 'map' clause}}
+#pragma omp target map(I) // lt50-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error 2 {{expected addressable lvalue in 'map' clause}}
foo();
#pragma omp target map(S2::S2s)
foo();
@@ -510,7 +541,7 @@
#pragma omp target map(to, x)
foo();
#pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
-#pragma omp target data map(tofrom: argc > 0 ? x : y) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error 2 {{expected addressable lvalue in 'map' clause}}
+#pragma omp target data map(tofrom: argc > 0 ? x : y) // lt50-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error 2 {{expected addressable lvalue in 'map' clause}}
#pragma omp target data map(argc)
#pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}}
#pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} warn-warning 2 {{Type 'const S2' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning 2 {{Type 'const S3' is not trivially copyable and not guaranteed to be mapped correctly}}
@@ -540,17 +571,30 @@
#pragma omp target data map(always, tofrom: x)
#pragma omp target data map(always: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
#pragma omp target data map(always, tofrom: always, tofrom, x)
#pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
foo();
#pragma omp target data map(close, tofrom: x)
#pragma omp target data map(close: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
#pragma omp target data map(close, tofrom: close, tofrom, x)
foo();
+// lt51-error@+1 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+#pragma omp target data map(present, tofrom: x)
+// expected-error@+2 {{missing map type}}
+// lt51-error@+1 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+#pragma omp target data map(present: x)
+// expected-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// expected-error@+2 {{missing map type}}
+// lt51-error@+1 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+#pragma omp target data map(tofrom, present: x)
+// lt51-error@+1 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+#pragma omp target data map(present, tofrom: present, tofrom, x)
+ foo();
+
T marr[10][10], iarr[5];
#pragma omp target data map(marr[10][0:2:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
{}
@@ -600,11 +644,11 @@
S6<int> m;
int x;
int y;
- int to, tofrom, always, close;
+ int to, tofrom, always, close, present;
const int (&l)[5] = da;
SC1 s;
SC1 *p;
-#pragma omp target data map // expected-error {{expected '(' after 'map'}} le45-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}} le50-error {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
+#pragma omp target data map // expected-error {{expected '(' after 'map'}} lt50-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}} ge50-error {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
#pragma omp target data map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
#pragma omp target data map() // expected-error {{expected expression}}
#pragma omp target data map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
@@ -623,7 +667,7 @@
#pragma omp target map(to, x)
foo();
#pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
-#pragma omp target data map(tofrom: argc > 0 ? argv[1] : argv[2]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error {{expected addressable lvalue in 'map' clause}}
+#pragma omp target data map(tofrom: argc > 0 ? argv[1] : argv[2]) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}}
#pragma omp target data map(argc)
#pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}}
#pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} warn-warning {{Type 'const S2' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning {{Type 'const S3' is not trivially copyable and not guaranteed to be mapped correctly}}
@@ -654,13 +698,23 @@
#pragma omp target data map(always, tofrom: x)
#pragma omp target data map(always: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
#pragma omp target data map(always, tofrom: always, tofrom, x)
#pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
foo();
#pragma omp target data map(close, tofrom: x)
#pragma omp target data map(close: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
+ foo();
+// lt51-error@+1 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+#pragma omp target data map(present, tofrom: x)
+// expected-error@+2 {{missing map type}}
+// lt51-error@+1 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+#pragma omp target data map(present: x)
+// expected-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// expected-error@+2 {{missing map type}}
+// lt51-error@+1 {{map type modifier 'present' requires OpenMP version 5.1 or above}}
+#pragma omp target data map(tofrom, present: x)
foo();
#pragma omp target private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target' directive}} expected-note {{defined as private}}
{}
@@ -717,21 +771,21 @@
int **BB, *offset, *a;
-#pragma omp target map(**(BB+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(**(BB+*offset)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
{}
-#pragma omp target map(**(BB+y)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(**(BB+y)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
{}
-#pragma omp target map(*(a+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(*(a+*offset)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
{}
-#pragma omp target map(**(*offset+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(**(*offset+BB)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
{}
-#pragma omp target map(**(y+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(**(y+BB)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
{}
-#pragma omp target map(*(*offset+a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(*(*offset+a)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
{}
-#pragma omp target map(**(*offset+BB+*a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(**(*offset+BB+*a)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
{}
-#pragma omp target map(**(*(*(&offset))+BB+*a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(**(*(*(&offset))+BB+*a)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
{}
#pragma omp target map(*(a+(a))) // expected-error {{invalid operands to binary expression ('int *' and 'int *')}}
{}
Index: clang/test/OpenMP/target_map_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_map_codegen.cpp
+++ clang/test/OpenMP/target_map_codegen.cpp
@@ -5702,6 +5702,288 @@
// CK31: define {{.+}}[[CALL00]]
// CK31: define {{.+}}[[CALL01]]
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DUSE -DCK31A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-64,CK31A-USE
+// RUN: %clang_cc1 -DUSE -DCK31A -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-64,CK31A-USE
+// RUN: %clang_cc1 -DUSE -DCK31A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-32,CK31A-USE
+// RUN: %clang_cc1 -DUSE -DCK31A -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-32,CK31A-USE
+
+// RUN: %clang_cc1 -DUSE -DCK31A -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DUSE -DCK31A -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DUSE -DCK31A -verify -fopenmp-version=51 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DUSE -DCK31A -fopenmp-version=51 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+
+// RUN: %clang_cc1 -DCK31A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-64,CK31A-NOUSE
+// RUN: %clang_cc1 -DCK31A -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-64,CK31A-NOUSE
+// RUN: %clang_cc1 -DCK31A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-32,CK31A-NOUSE
+// RUN: %clang_cc1 -DCK31A -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-32,CK31A-NOUSE
+
+// RUN: %clang_cc1 -DCK31A -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK31A -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK31A -verify -fopenmp-version=51 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK31A -fopenmp-version=51 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
+#ifdef CK31A
+
+// CK31A: [[ST:%.+]] = type { i32, i32 }
+struct ST {
+ int i;
+ int j;
+};
+
+// CK31A-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
+// CK31A: [[MTYPE00:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x820]],
+// CK31A-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000000803]], i64 [[#0x823]],
+// CK31A-SAME: {{^}} i64 [[#0x820]], i64 [[#0x5000000000803]], i64 [[#0x5000000000003]]]
+
+// CK31A-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
+// CK31A: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
+// CK31A: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 [[#0xC27]]]
+
+// CK31A-LABEL: explicit_maps_single{{.*}}(
+void explicit_maps_single (int ii){
+ // CK31A: alloca i32
+
+ // Map of a scalar.
+ // CK31A: [[A:%.+]] = alloca i32
+ int a = ii;
+
+ // CK31A: [[ST1:%.+]] = alloca [[ST]]
+ // CK31A: [[ST2:%.+]] = alloca [[ST]]
+ struct ST st1;
+ struct ST st2;
+
+ // Make sure the struct picks up present even if another element of the struct
+ // doesn't have present.
+ // Region 00
+ // CK31A: [[ST1_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 0
+ // CK31A: [[ST1_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 1
+ // CK31A: [[ST2_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 0
+ // CK31A: [[ST2_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 1
+ // CK31A-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 7, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[7 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK31A-DAG: [[GEPS]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0
+ // CK31A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK31A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // st1
+ // CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK31A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+ // CK31A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
+ // CK31A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+ // CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP0]]
+ // CK31A-DAG: store i32* [[ST1_I]], i32** [[CP0]]
+ // CK31A-DAG: store i64 %{{.+}}, i64* [[S0]]
+
+ // st1.i
+ // CK31A-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK31A-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK31A-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+ // CK31A-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
+ // CK31A-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
+ // CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP1]]
+ // CK31A-DAG: store i32* [[ST1_I]], i32** [[CP1]]
+ // CK31A-DAG: store i64 4, i64* [[S1]]
+
+ // st1.j
+ // CK31A-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+ // CK31A-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+ // CK31A-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
+ // CK31A-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
+ // CK31A-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32**
+ // CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP2]]
+ // CK31A-DAG: store i32* [[ST1_J]], i32** [[CP2]]
+ // CK31A-DAG: store i64 4, i64* [[S2]]
+
+ // a
+ // CK31A-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
+ // CK31A-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
+ // CK31A-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
+ // CK31A-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to i32**
+ // CK31A-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32**
+ // CK31A-DAG: store i32* [[A]], i32** [[CBP3]]
+ // CK31A-DAG: store i32* [[A]], i32** [[CP3]]
+ // CK31A-DAG: store i64 4, i64* [[S3]]
+
+ // st2
+ // CK31A-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 4
+ // CK31A-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 4
+ // CK31A-DAG: [[S4:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 4
+ // CK31A-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to [[ST]]**
+ // CK31A-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i32**
+ // CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP4]]
+ // CK31A-DAG: store i32* [[ST2_I]], i32** [[CP4]]
+ // CK31A-DAG: store i64 %{{.+}}, i64* [[S4]]
+
+ // st2.i
+ // CK31A-DAG: [[BP5:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 5
+ // CK31A-DAG: [[P5:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 5
+ // CK31A-DAG: [[S5:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 5
+ // CK31A-DAG: [[CBP5:%.+]] = bitcast i8** [[BP5]] to [[ST]]**
+ // CK31A-DAG: [[CP5:%.+]] = bitcast i8** [[P5]] to i32**
+ // CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP5]]
+ // CK31A-DAG: store i32* [[ST2_I]], i32** [[CP5]]
+ // CK31A-DAG: store i64 4, i64* [[S5]]
+
+ // st2.j
+ // CK31A-DAG: [[BP6:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 6
+ // CK31A-DAG: [[P6:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 6
+ // CK31A-DAG: [[S6:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 6
+ // CK31A-DAG: [[CBP6:%.+]] = bitcast i8** [[BP6]] to [[ST]]**
+ // CK31A-DAG: [[CP6:%.+]] = bitcast i8** [[P6]] to i32**
+ // CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP6]]
+ // CK31A-DAG: store i32* [[ST2_J]], i32** [[CP6]]
+ // CK31A-DAG: store i64 4, i64* [[S6]]
+
+ // CK31A-USE: call void [[CALL00:@.+]]([[ST]]* [[ST1]], i32* [[A]], [[ST]]* [[ST2]])
+ // CK31A-NOUSE: call void [[CALL00:@.+]]()
+ #pragma omp target map(tofrom: st1.i) map(present, tofrom: a, st1.j, st2.i) map(tofrom: st2.j)
+ {
+#ifdef USE
+ st1.i++;
+ a++;
+ st1.j++;
+ st2.i++;
+ st2.j++;
+#endif
+ }
+
+ // Always Close Present.
+ // Region 01
+ // CK31A-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+ // CK31A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK31A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK31A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+ // CK31A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+ // CK31A-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]]
+ // CK31A-DAG: store i32* [[VAR0]], i32** [[CP0]]
+
+ // CK31A-USE: call void [[CALL01:@.+]](i32* {{[^,]+}})
+ // CK31A-NOUSE: call void [[CALL01:@.+]]()
+ #pragma omp target map(always close present tofrom: a)
+ {
+#ifdef USE
+ a++;
+#endif
+ }
+}
+// CK31A: define {{.+}}[[CALL00]]
+// CK31A: define {{.+}}[[CALL01]]
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DUSE -DCK31B -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31B,CK31B-64,CK31B-USE
+// RUN: %clang_cc1 -DUSE -DCK31B -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31B,CK31B-64,CK31B-USE
+// RUN: %clang_cc1 -DUSE -DCK31B -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31B,CK31B-32,CK31B-USE
+// RUN: %clang_cc1 -DUSE -DCK31B -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31B,CK31B-32,CK31B-USE
+
+// RUN: %clang_cc1 -DUSE -DCK31B -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DUSE -DCK31B -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DUSE -DCK31B -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DUSE -DCK31B -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+
+// RUN: %clang_cc1 -DCK31B -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31B,CK31B-64,CK31B-NOUSE
+// RUN: %clang_cc1 -DCK31B -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31B,CK31B-64,CK31B-NOUSE
+// RUN: %clang_cc1 -DCK31B -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31B,CK31B-32,CK31B-NOUSE
+// RUN: %clang_cc1 -DCK31B -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31B,CK31B-32,CK31B-NOUSE
+
+// RUN: %clang_cc1 -DCK31B -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK31B -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK31B -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK31B -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
+#ifdef CK31B
+
+// CK31B: [[ST:%.+]] = type { i32, i32 }
+
+// CK31B-LABEL: @.__omp_offloading_{{.*}}test_present_members{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
+// CK31B: [[MTYPE00:@.+]] = private {{.*}}constant [3 x i64] [i64 [[#0x820]],
+// CK31B-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000000803]]]
+
+struct ST {
+ int i;
+ int j;
+ // CK31B-LABEL: define {{.*}}test_present_members{{.*}}(
+ void test_present_members() {
+ // Make sure the struct picks up present even if another element of the
+ // struct doesn't have present.
+ // Region 00
+ // CK31B: [[I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS:%.+]], i{{.+}} 0, i{{.+}} 0
+ // CK31B: [[J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS]], i{{.+}} 0, i{{.+}} 1
+ // CK31B-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK31B-DAG: [[GEPS]] = getelementptr inbounds [3 x i64], [3 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0
+ // CK31B-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK31B-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // this
+ // CK31B-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK31B-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK31B-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+ // CK31B-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
+ // CK31B-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+ // CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP0]]
+ // CK31B-DAG: store i32* [[I]], i32** [[CP0]]
+ // CK31B-DAG: store i64 %{{.+}}, i64* [[S0]]
+
+ // i
+ // CK31B-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK31B-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK31B-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+ // CK31B-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
+ // CK31B-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
+ // CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP1]]
+ // CK31B-DAG: store i32* [[I]], i32** [[CP1]]
+ // CK31B-DAG: store i64 4, i64* [[S1]]
+
+ // j
+ // CK31B-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+ // CK31B-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+ // CK31B-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
+ // CK31B-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
+ // CK31B-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32**
+ // CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP2]]
+ // CK31B-DAG: store i32* [[J]], i32** [[CP2]]
+ // CK31B-DAG: store i64 4, i64* [[S2]]
+
+ // CK31B-USE: call void [[CALL00:@.+]]([[ST]]* [[THIS]])
+ // CK31B-NOUSE: call void [[CALL00:@.+]]()
+ #pragma omp target map(tofrom: i) map(present, tofrom: j)
+ {
+#ifdef USE
+ i++;
+ j++;
+#endif
+ }
+ }
+};
+
+void test() {
+ ST s;
+ s.test_present_members();
+}
+
+// CK31B: define {{.+}}[[CALL00]]
+
#endif
///==========================================================================///
// RUN: %clang_cc1 -DCK32 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK32 --check-prefix CK32-64
Index: clang/test/OpenMP/target_enter_data_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_enter_data_codegen.cpp
+++ clang/test/OpenMP/target_enter_data_codegen.cpp
@@ -203,6 +203,90 @@
}
#endif
///==========================================================================///
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1A --check-prefix CK1A-64
+// RUN: %clang_cc1 -DCK1A -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1A --check-prefix CK1A-64
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1A --check-prefix CK1A-32
+// RUN: %clang_cc1 -DCK1A -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1A --check-prefix CK1A-32
+
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1A -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1A -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK1A
+
+// CK1A: [[ST:%.+]] = type { i32, double* }
+template <typename T>
+struct ST {
+ T a;
+ double *b;
+};
+
+ST<int> gb;
+double gc[100];
+
+// CK1A: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x821]]]
+
+// CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0xC25]]]
+
+// CK1A-LABEL: _Z3fooi
+void foo(int arg) {
+ int la;
+ float lb[arg];
+
+ // Region 00
+ // CK1A-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:32|64]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // CK1A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+ // CK1A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+ // CK1A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+ // CK1A-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
+ // CK1A-DAG: store float* [[VAR0]], float** [[CP0]]
+ // CK1A-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
+ // CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
+ // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
+ // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
+ // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+ // CK1A-NOT: __tgt_target_data_end
+ #pragma omp target enter data map(present, to: lb)
+ {++arg;}
+
+ // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+ {++arg;}
+
+ // Region 01
+ // CK1A-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+ // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // CK1A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+ // CK1A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+ // CK1A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+ // CK1A-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
+ // CK1A-DAG: store float* [[VAR0]], float** [[CP0]]
+ // CK1A-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
+ // CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
+ // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
+ // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
+ // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+ // CK1A-NOT: __tgt_target_data_end
+ #pragma omp target enter data map(always close present, to: lb)
+ {++arg;}
+}
+#endif
+///==========================================================================///
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
Index: clang/test/OpenMP/target_data_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_data_codegen.cpp
+++ clang/test/OpenMP/target_data_codegen.cpp
@@ -225,6 +225,97 @@
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
#pragma omp target data map(always close, to: lb)
{++arg;}
+
+}
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1A --check-prefix CK1A-64
+// RUN: %clang_cc1 -DCK1A -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1A --check-prefix CK1A-64
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1A --check-prefix CK1A-32
+// RUN: %clang_cc1 -DCK1A -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1A --check-prefix CK1A-32
+
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1A -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1A -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK1A
+
+// CK1A: [[ST:%.+]] = type { i32, double* }
+template <typename T>
+struct ST {
+ T a;
+ double *b;
+};
+
+ST<int> gb;
+double gc[100];
+
+// CK1A: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x821]]]
+
+// CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0xC25]]]
+
+// CK1A-LABEL: _Z3fooi
+void foo(int arg) {
+ int la;
+ float lb[arg];
+
+ // Region 00
+ // CK1A-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:32|64]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // CK1A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+ // CK1A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+ // CK1A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+ // CK1A-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
+ // CK1A-DAG: store float* [[VAR0]], float** [[CP0]]
+ // CK1A-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
+ // CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
+ // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
+ // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
+ // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+
+ // CK1A-DAG: call void @__tgt_target_data_end(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+ // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
+ // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
+ #pragma omp target data map(present, to: lb)
+ {++arg;}
+
+ // Region 01
+ // CK1A-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+ // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // CK1A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+ // CK1A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+ // CK1A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+ // CK1A-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
+ // CK1A-DAG: store float* [[VAR0]], float** [[CP0]]
+ // CK1A-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
+ // CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
+ // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
+ // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
+ // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+
+ // CK1A-DAG: call void @__tgt_target_data_end(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+ // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+ // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
+ // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
+ #pragma omp target data map(always close present, to: lb)
+ {++arg;}
+
}
#endif
///==========================================================================///
@@ -510,4 +601,77 @@
{ ++arg, ++(*p); }
}
#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK8 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK8 --check-prefix CK8-64
+// RUN: %clang_cc1 -DCK8 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK8 --check-prefix CK8-64
+// RUN: %clang_cc1 -DCK8 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK8 --check-prefix CK8-32
+// RUN: %clang_cc1 -DCK8 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK8 --check-prefix CK8-32
+
+// RUN: %clang_cc1 -DCK8 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK8 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK8 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK8 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}#ifdef CK8
+#ifdef CK8
+struct S1 {
+ int i;
+};
+struct S2 {
+ S1 s;
+ struct S2 *ps;
+};
+
+void test_present_modifier(int arg) {
+ S2 *ps1;
+ S2 *ps2;
+
+ // Make sure the struct picks up present even if another element of the struct
+ // doesn't have present.
+
+ // CK8: private unnamed_addr constant [11 x i64]
+
+ // ps1
+ // CK8-SAME: {{^}} [i64 [[#0x820]], i64 [[#0x1000000000003]],
+ // CK8-SAME: {{^}} i64 [[#0x1000000000810]], i64 [[#0x810]], i64 [[#0x813]],
+
+ // arg
+ // CK8-SAME: {{^}} i64 [[#0x823]],
+
+ // ps2
+ // CK8-SAME: {{^}} i64 [[#0x820]], i64 [[#0x7000000000803]],
+ // CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
+ #pragma omp target data map(tofrom: ps1->s) \
+ map(present,tofrom: arg, ps1->ps->ps->ps->s, ps2->s) \
+ map(tofrom: ps2->ps->ps->ps->s)
+ {
+ ++(arg);
+ }
+}
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK9 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK9 --check-prefix CK9-64
+// RUN: %clang_cc1 -DCK9 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK9 --check-prefix CK9-64
+// RUN: %clang_cc1 -DCK9 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK9 --check-prefix CK9-32
+// RUN: %clang_cc1 -DCK9 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK9 --check-prefix CK9-32
+
+// RUN: %clang_cc1 -DCK9 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK9 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK9 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK9 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
+#ifdef CK9
+void test_present_modifier(int arg) {
+ // CK9: private unnamed_addr constant [1 x i64] [i64 [[#0x823]]]
+ #pragma omp target data map(present,tofrom: arg)
+ {++arg;}
+}
+#endif
#endif
Index: clang/test/OpenMP/target_data_ast_print.cpp
===================================================================
--- clang/test/OpenMP/target_data_ast_print.cpp
+++ clang/test/OpenMP/target_data_ast_print.cpp
@@ -5,6 +5,14 @@
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -DOMP51 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
+
+// RUN: %clang_cc1 -DOMP51 -verify -fopenmp-simd -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
// expected-no-diagnostics
#ifndef HEADER
@@ -43,6 +51,11 @@
#pragma omp target data map(close,alloc: e)
foo();
+#ifdef OMP51
+#pragma omp target data map(present,alloc: e)
+ foo();
+#endif
+
// nesting a target region
#pragma omp target data map(e)
{
@@ -50,6 +63,10 @@
foo();
#pragma omp target map(close, alloc: e)
foo();
+#ifdef OMP51
+ #pragma omp target map(present, alloc: e)
+ foo();
+#endif
}
return 0;
@@ -75,12 +92,16 @@
// CHECK-NEXT: foo();
// CHECK-NEXT: #pragma omp target data map(close,alloc: e)
// CHECK-NEXT: foo();
+// OMP51-NEXT: #pragma omp target data map(present,alloc: e)
+// OMP51-NEXT: foo();
// CHECK-NEXT: #pragma omp target data map(tofrom: e)
// CHECK-NEXT: {
// CHECK-NEXT: #pragma omp target map(always,alloc: e)
// CHECK-NEXT: foo();
// CHECK-NEXT: #pragma omp target map(close,alloc: e)
// CHECK-NEXT: foo();
+// OMP51-NEXT: #pragma omp target map(present,alloc: e)
+// OMP51-NEXT: foo();
// CHECK: template<> int tmain<int, 5>(int argc, int *argv) {
// CHECK-NEXT: int i, j, b, c, d, e, x[20];
// CHECK-NEXT: #pragma omp target data map(to: c)
@@ -101,12 +122,16 @@
// CHECK-NEXT: foo();
// CHECK-NEXT: #pragma omp target data map(close,alloc: e)
// CHECK-NEXT: foo();
+// OMP51-NEXT: #pragma omp target data map(present,alloc: e)
+// OMP51-NEXT: foo();
// CHECK-NEXT: #pragma omp target data map(tofrom: e)
// CHECK-NEXT: {
// CHECK-NEXT: #pragma omp target map(always,alloc: e)
// CHECK-NEXT: foo();
// CHECK-NEXT: #pragma omp target map(close,alloc: e)
// CHECK-NEXT: foo();
+// OMP51-NEXT: #pragma omp target map(present,alloc: e)
+// OMP51-NEXT: foo();
// CHECK: template<> char tmain<char, 1>(char argc, char *argv) {
// CHECK-NEXT: char i, j, b, c, d, e, x[20];
// CHECK-NEXT: #pragma omp target data map(to: c)
@@ -127,12 +152,16 @@
// CHECK-NEXT: foo();
// CHECK-NEXT: #pragma omp target data map(close,alloc: e)
// CHECK-NEXT: foo();
+// OMP51-NEXT: #pragma omp target data map(present,alloc: e)
+// OMP51-NEXT: foo();
// CHECK-NEXT: #pragma omp target data map(tofrom: e)
// CHECK-NEXT: {
// CHECK-NEXT: #pragma omp target map(always,alloc: e)
// CHECK-NEXT: foo();
// CHECK-NEXT: #pragma omp target map(close,alloc: e)
// CHECK-NEXT: foo();
+// OMP51-NEXT: #pragma omp target map(present,alloc: e)
+// OMP51-NEXT: foo();
int main (int argc, char **argv) {
int b = argc, c, d, e, f, g, x[20];
@@ -185,6 +214,13 @@
foo();
// CHECK-NEXT: foo();
+// OMP51-NEXT: #pragma omp target data map(present,alloc: e)
+// OMP51-NEXT: foo();
+#ifdef OMP51
+#pragma omp target data map(present,alloc: e)
+ foo();
+#endif
+
// nesting a target region
#pragma omp target data map(e)
// CHECK-NEXT: #pragma omp target data map(tofrom: e)
@@ -197,6 +233,10 @@
#pragma omp target map(close, alloc: e)
// CHECK-NEXT: #pragma omp target map(close,alloc: e)
foo();
+// CHECK-NEXT: foo();
+#pragma omp target map(always, alloc: e)
+// CHECK-NEXT: #pragma omp target map(always,alloc: e)
+ foo();
}
return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -17623,9 +17623,9 @@
OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, SourceLocation MapLoc,
SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
const OMPVarListLocTy &Locs, ArrayRef<Expr *> UnresolvedMappers) {
- OpenMPMapModifierKind Modifiers[] = {OMPC_MAP_MODIFIER_unknown,
- OMPC_MAP_MODIFIER_unknown,
- OMPC_MAP_MODIFIER_unknown};
+ OpenMPMapModifierKind Modifiers[] = {
+ OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
+ OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown};
SourceLocation ModifiersLoc[NumberOfOMPMapClauseModifiers];
// Process map-type-modifiers, flag errors for duplicate modifiers.
Index: clang/lib/Parse/ParseOpenMP.cpp
===================================================================
--- clang/lib/Parse/ParseOpenMP.cpp
+++ clang/lib/Parse/ParseOpenMP.cpp
@@ -3104,9 +3104,16 @@
while (getCurToken().isNot(tok::colon)) {
OpenMPMapModifierKind TypeModifier = isMapModifier(*this);
if (TypeModifier == OMPC_MAP_MODIFIER_always ||
- TypeModifier == OMPC_MAP_MODIFIER_close) {
- Data.MapTypeModifiers.push_back(TypeModifier);
- Data.MapTypeModifiersLoc.push_back(Tok.getLocation());
+ TypeModifier == OMPC_MAP_MODIFIER_close ||
+ TypeModifier == OMPC_MAP_MODIFIER_present) {
+ if (getLangOpts().OpenMP < 51 &&
+ TypeModifier == OMPC_MAP_MODIFIER_present) {
+ Diag(Tok, diag::err_omp_map_type_modifier_wrong_version)
+ << PP.getSpelling(Tok) << "5.1";
+ } else {
+ Data.MapTypeModifiers.push_back(TypeModifier);
+ Data.MapTypeModifiersLoc.push_back(Tok.getLocation());
+ }
ConsumeToken();
} else if (TypeModifier == OMPC_MAP_MODIFIER_mapper) {
Data.MapTypeModifiers.push_back(TypeModifier);
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7040,6 +7040,8 @@
/// Close is a hint to the runtime to allocate memory close to
/// the target device.
OMP_MAP_CLOSE = 0x400,
+ /// Produce a runtime error if the data is not already allocated.
+ OMP_MAP_PRESENT = 0x800,
/// The 16 MSBs of the flags indicate whether the entry is member of some
/// struct/class.
OMP_MAP_MEMBER_OF = 0xffff000000000000,
@@ -7260,6 +7262,9 @@
if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_close)
!= MapModifiers.end())
Bits |= OMP_MAP_CLOSE;
+ if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_present)
+ != MapModifiers.end())
+ Bits |= OMP_MAP_PRESENT;
return Bits;
}
@@ -7926,6 +7931,13 @@
Sizes.push_back(Size);
// Map type is always TARGET_PARAM
Types.push_back(OMP_MAP_TARGET_PARAM);
+ // If any element has the present modifier, then make sure the runtime
+ // doesn't attempt to allocate the struct.
+ if (CurTypes.end() !=
+ llvm::find_if(CurTypes, [](OpenMPOffloadMappingFlags Type) {
+ return Type & OMP_MAP_PRESENT;
+ }))
+ Types.back() |= OMP_MAP_PRESENT;
// Remove TARGET_PARAM flag from the first element
(*CurTypes.begin()) &= ~OMP_MAP_TARGET_PARAM;
Index: clang/include/clang/Basic/OpenMPKinds.def
===================================================================
--- clang/include/clang/Basic/OpenMPKinds.def
+++ clang/include/clang/Basic/OpenMPKinds.def
@@ -124,6 +124,7 @@
OPENMP_MAP_MODIFIER_KIND(always)
OPENMP_MAP_MODIFIER_KIND(close)
OPENMP_MAP_MODIFIER_KIND(mapper)
+OPENMP_MAP_MODIFIER_KIND(present)
// Modifiers for 'to' clause.
OPENMP_TO_MODIFIER_KIND(mapper)
Index: clang/include/clang/Basic/DiagnosticParseKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticParseKinds.td
+++ clang/include/clang/Basic/DiagnosticParseKinds.td
@@ -1248,7 +1248,9 @@
def err_omp_unknown_map_type : Error<
"incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'">;
def err_omp_unknown_map_type_modifier : Error<
- "incorrect map type modifier, expected 'always', 'close', or 'mapper'">;
+ "incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'">;
+def err_omp_map_type_modifier_wrong_version : Error<
+ "map type modifier '%0' requires OpenMP version %1 or above">;
def err_omp_map_type_missing : Error<
"missing map type">;
def err_omp_map_type_modifier_missing : Error<
Index: clang/include/clang/AST/OpenMPClause.h
===================================================================
--- clang/include/clang/AST/OpenMPClause.h
+++ clang/include/clang/AST/OpenMPClause.h
@@ -5354,7 +5354,7 @@
/// Map-type-modifiers for the 'map' clause.
OpenMPMapModifierKind MapTypeModifiers[NumberOfOMPMapClauseModifiers] = {
OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
- OMPC_MAP_MODIFIER_unknown};
+ OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown};
/// Location of map-type-modifiers for the 'map' clause.
SourceLocation MapTypeModifiersLoc[NumberOfOMPMapClauseModifiers];
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
