kkwli0 updated this revision to Diff 50774.
kkwli0 marked 5 inline comments as done.
kkwli0 added a comment.
Addressed the comments from the last review: added assert calls and outline the
common code in ActOnOpenMPToClause, ActOnOpenMPFromClause and
ActOnOpenMPMapClause to a static function.
http://reviews.llvm.org/D15944
Files:
include/clang-c/Index.h
include/clang/AST/OpenMPClause.h
include/clang/AST/RecursiveASTVisitor.h
include/clang/AST/StmtOpenMP.h
include/clang/Basic/DiagnosticSemaKinds.td
include/clang/Basic/OpenMPKinds.def
include/clang/Basic/StmtNodes.td
include/clang/Sema/Sema.h
include/clang/Serialization/ASTBitCodes.h
lib/AST/OpenMPClause.cpp
lib/AST/StmtOpenMP.cpp
lib/AST/StmtPrinter.cpp
lib/AST/StmtProfile.cpp
lib/Basic/OpenMPKinds.cpp
lib/CodeGen/CGStmt.cpp
lib/CodeGen/CGStmtOpenMP.cpp
lib/CodeGen/CodeGenFunction.h
lib/Parse/ParseOpenMP.cpp
lib/Sema/SemaOpenMP.cpp
lib/Sema/TreeTransform.h
lib/Serialization/ASTReaderStmt.cpp
lib/Serialization/ASTWriterStmt.cpp
lib/StaticAnalyzer/Core/ExprEngine.cpp
test/OpenMP/nesting_of_regions.cpp
test/OpenMP/target_map_messages.cpp
test/OpenMP/target_parallel_for_map_messages.cpp
test/OpenMP/target_parallel_map_messages.cpp
test/OpenMP/target_update_ast_print.cpp
test/OpenMP/target_update_device_messages.cpp
test/OpenMP/target_update_from_messages.cpp
test/OpenMP/target_update_if_messages.cpp
test/OpenMP/target_update_messages.cpp
test/OpenMP/target_update_to_messages.cpp
tools/libclang/CIndex.cpp
tools/libclang/CXCursor.cpp
Index: tools/libclang/CXCursor.cpp
===================================================================
--- tools/libclang/CXCursor.cpp
+++ tools/libclang/CXCursor.cpp
@@ -612,6 +612,9 @@
case Stmt::OMPTargetParallelForDirectiveClass:
K = CXCursor_OMPTargetParallelForDirective;
break;
+ case Stmt::OMPTargetUpdateDirectiveClass:
+ K = CXCursor_OMPTargetUpdateDirective;
+ break;
case Stmt::OMPTeamsDirectiveClass:
K = CXCursor_OMPTeamsDirective;
break;
Index: tools/libclang/CIndex.cpp
===================================================================
--- tools/libclang/CIndex.cpp
+++ tools/libclang/CIndex.cpp
@@ -2255,6 +2255,12 @@
}
void OMPClauseEnqueue::VisitOMPDefaultmapClause(
const OMPDefaultmapClause * /*C*/) {}
+void OMPClauseEnqueue::VisitOMPFromClause(const OMPFromClause *C) {
+ VisitOMPClauseList(C);
+}
+void OMPClauseEnqueue::VisitOMPToClause(const OMPToClause *C) {
+ VisitOMPClauseList(C);
+}
}
void EnqueueVisitor::EnqueueChildren(const OMPClause *S) {
@@ -4825,6 +4831,8 @@
return cxstring::createRef("OMPTargetParallelDirective");
case CXCursor_OMPTargetParallelForDirective:
return cxstring::createRef("OMPTargetParallelForDirective");
+ case CXCursor_OMPTargetUpdateDirective:
+ return cxstring::createRef("OMPTargetUpdateDirective");
case CXCursor_OMPTeamsDirective:
return cxstring::createRef("OMPTeamsDirective");
case CXCursor_OMPCancellationPointDirective:
Index: test/OpenMP/target_update_to_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_update_to_messages.cpp
@@ -0,0 +1,175 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+ return argc;
+}
+
+struct S1; // expected-note 2 {{declared here}}
+extern S1 a;
+class S2 {
+ mutable int a;
+public:
+ S2():a(0) { }
+ S2(S2 &s2):a(s2.a) { }
+ static float S2s; // expected-note 4 {{mappable type cannot contain static members}}
+ static const float S2sc; // expected-note 4 {{mappable type cannot contain static members}}
+};
+const float S2::S2sc = 0;
+const S2 b;
+const S2 ba[5];
+class S3 {
+ int a;
+public:
+ S3():a(0) { }
+ S3(S3 &s3):a(s3.a) { }
+};
+const S3 c;
+const S3 ca[5];
+extern const int f;
+class S4 {
+ int a;
+ S4();
+ S4(const S4 &s4);
+public:
+ S4(int v):a(v) { }
+};
+class S5 {
+ int a;
+ S5():a(0) {}
+ S5(const S5 &s5):a(s5.a) { }
+public:
+ S5(int v):a(v) { }
+};
+struct S6 {
+ int ii;
+ int aa[30];
+ float xx;
+ double *pp;
+};
+struct S7 {
+ int i;
+ int a[50];
+ float x;
+ S6 s6[5];
+ double *p;
+ unsigned bfa : 4;
+};
+
+S3 h;
+#pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
+
+typedef int from;
+
+template <typename T, int I> // expected-note {{declared here}}
+T tmain(T argc) {
+ const T d = 5;
+ const T da[5] = { 0 };
+ S4 e(4);
+ S5 g(5);
+ T *m;
+ T i, t[20];
+ T &j = i;
+ T *k = &j;
+ T x;
+ T y;
+ T to;
+ const T (&l)[5] = da;
+ S7 s7;
+
+#pragma omp target update to // expected-error {{expected '(' after 'to'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to() // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update() // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(x)
+#pragma omp target update to(t[:I])
+#pragma omp target update to(T) // expected-error {{'T' does not refer to a value}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(I) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update to(S2::S2s)
+#pragma omp target update to(S2::S2sc)
+#pragma omp target update to(to)
+#pragma omp target update to(y x) // expected-error {{expected ',' or ')' in 'to' clause}}
+#pragma omp target update to(argc > 0 ? x : y) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update to(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target update to(ba) // expected-error 2 {{type 'S2' is not mappable to target}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(h) // expected-error {{threadprivate variables are not allowed in 'to' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(k), from(k) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}}
+#pragma omp target update to(t), to(t[:5]) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}}
+#pragma omp target update to(da)
+#pragma omp target update to(da[:4])
+
+#pragma omp target update to(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update to(x, c[:]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update to(x, (m+1)[2]) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update to(s7.i, s7.a[:3])
+#pragma omp target update to(s7.s6[1].aa[0:5])
+#pragma omp target update to(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update to(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update to(s7.p[:10])
+#pragma omp target update to(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'to' clause}}
+#pragma omp target update to(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+#pragma omp target data map(to: s7.i)
+ {
+#pragma omp target update to(s7.x)
+ }
+ return 0;
+}
+
+int main(int argc, char **argv) {
+ const int d = 5;
+ const int da[5] = { 0 };
+ S4 e(4);
+ S5 g(5);
+ int i, t[20];
+ int &j = i;
+ int *k = &j;
+ int x;
+ int y;
+ int to;
+ const int (&l)[5] = da;
+ S7 s7;
+ int *m;
+
+#pragma omp target update to // expected-error {{expected '(' after 'to'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to() // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update() // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(x)
+#pragma omp target update to(t[:i])
+#pragma omp target update to(S2::S2s)
+#pragma omp target update to(S2::S2sc)
+#pragma omp target update to(to)
+#pragma omp target update to(y x) // expected-error {{expected ',' or ')' in 'to' clause}}
+#pragma omp target update to(argc > 0 ? x : y) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target update to(ba) // expected-error 2 {{type 'S2' is not mappable to target}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(h) // expected-error {{threadprivate variables are not allowed in 'to' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(k), from(k) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}}
+#pragma omp target update to(t), to(t[:5]) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}}
+#pragma omp target update to(da)
+#pragma omp target update to(da[:4])
+
+#pragma omp target update to(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update to(x, c[:]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update to(x, (m+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update to(s7.i, s7.a[:3])
+#pragma omp target update to(s7.s6[1].aa[0:5])
+#pragma omp target update to(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update to(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update to(s7.p[:10])
+#pragma omp target update to(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'to' clause}}
+#pragma omp target update to(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+#pragma omp target data map(to: s7.i)
+ {
+#pragma omp target update to(s7.x)
+ }
+
+ return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
+}
+
Index: test/OpenMP/target_update_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_update_messages.cpp
@@ -0,0 +1,31 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+ return argc;
+}
+
+struct S1; // Aexpected-note {{declared here}}
+
+template <class T, class S> // Aexpected-note {{declared here}}
+int tmain(T argc, S **argv) {
+ int n;
+ return 0;
+}
+
+int main(int argc, char **argv) {
+ int m;
+ #pragma omp target update to(m) { // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+ #pragma omp target update to(m) ( // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+ #pragma omp target update to(m) [ // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+ #pragma omp target update to(m) ] // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+ #pragma omp target update to(m) ) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+
+ #pragma omp target update from(m) // OK
+ {
+ foo();
+ }
+ return tmain(argc, argv);
+}
Index: test/OpenMP/target_update_if_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_update_if_messages.cpp
@@ -0,0 +1,58 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+ return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+template <class T, class S> // expected-note {{declared here}}
+int tmain(T argc, S **argv) {
+ int n;
+#pragma omp target update to(n) if // expected-error {{expected '(' after 'if'}}
+#pragma omp target update from(n) if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(n) if () // expected-error {{expected expression}}
+#pragma omp target update from(n) if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(n) if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+#pragma omp target update from(n) if (argc > 0 ? argv[1] : argv[2])
+#pragma omp target update to(n) if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause}}
+#pragma omp target update from(n) if (S) // expected-error {{'S' does not refer to a value}}
+#pragma omp target update to(n) if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update from(n) if (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(n) if(argc)
+#pragma omp target update from(n) if(target update // expected-warning {{missing ':' after directive name modifier - ignoring}} expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(n) if(target update : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update from(n) if(target update : argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(n) if(target update : argc)
+#pragma omp target update from(n) if(target update : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target update'}}
+#pragma omp target update to(n) if(target update : argc) if (target update:argc) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause with 'target update' name modifier}}
+#pragma omp target update from(n) if(target update : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}}
+ return 0;
+}
+
+int main(int argc, char **argv) {
+ int m;
+#pragma omp target update to(m) if // expected-error {{expected '(' after 'if'}}
+#pragma omp target update from(m) if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(m) if () // expected-error {{expected expression}}
+#pragma omp target update from(m) if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(m) if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+#pragma omp target update from(m) if (argc > 0 ? argv[1] : argv[2])
+#pragma omp target update to(m) if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause}}
+#pragma omp target update from(m) if (S1) // expected-error {{'S1' does not refer to a value}}
+#pragma omp target update to(m) if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update from(m) if (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(m) if (1 0) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update from(m) if(if(tmain(argc, argv) // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(m) if(target update // expected-warning {{missing ':' after directive name modifier - ignoring}} expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update from(m) if(target update : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(m) if(target update : argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update from(m) if(target update : argc)
+#pragma omp target update to(m) if(target update : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target update'}}
+#pragma omp target update from(m) if(target update : argc) if (target update:argc) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause with 'target update' name modifier}}
+#pragma omp target update to(m) if(target update : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}}
+ return tmain(argc, argv);
+}
Index: test/OpenMP/target_update_from_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_update_from_messages.cpp
@@ -0,0 +1,176 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+ return argc;
+}
+
+struct S1; // expected-note 2 {{declared here}}
+extern S1 a;
+class S2 {
+ mutable int a;
+public:
+ S2():a(0) { }
+ S2(S2 &s2):a(s2.a) { }
+ static float S2s; // expected-note 4 {{mappable type cannot contain static members}}
+ static const float S2sc; // expected-note 4 {{mappable type cannot contain static members}}
+};
+const float S2::S2sc = 0;
+const S2 b;
+const S2 ba[5];
+class S3 {
+ int a;
+public:
+ S3():a(0) { }
+ S3(S3 &s3):a(s3.a) { }
+};
+const S3 c;
+const S3 ca[5];
+extern const int f;
+class S4 {
+ int a;
+ S4();
+ S4(const S4 &s4);
+public:
+ S4(int v):a(v) { }
+};
+class S5 {
+ int a;
+ S5():a(0) {}
+ S5(const S5 &s5):a(s5.a) { }
+public:
+ S5(int v):a(v) { }
+};
+struct S6 {
+ int ii;
+ int aa[30];
+ float xx;
+ double *pp;
+};
+struct S7 {
+ int i;
+ int a[50];
+ float x;
+ S6 s6[5];
+ double *p;
+ unsigned bfa : 4;
+};
+
+S3 h;
+#pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
+
+typedef int to;
+
+template <typename T, int I> // expected-note {{declared here}}
+T tmain(T argc) {
+ const T d = 5;
+ const T da[5] = { 0 };
+ S4 e(4);
+ S5 g(5);
+ T i, t[20];
+ T &j = i;
+ T *k = &j;
+ T x;
+ T y;
+ T from;
+ const T (&l)[5] = da;
+ T *m;
+ S7 s7;
+
+#pragma omp target update from // expected-error {{expected '(' after 'from'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from() // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update() // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(x)
+#pragma omp target update from(t[:I])
+#pragma omp target update from(T) // expected-error {{'T' does not refer to a value}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(I) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update from(S2::S2s)
+#pragma omp target update from(S2::S2sc)
+#pragma omp target update from(from)
+#pragma omp target update from(y x) // expected-error {{expected ',' or ')' in 'from' clause}}
+#pragma omp target update from(argc > 0 ? x : y) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update from(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target update from(ba) // expected-error 2 {{type 'S2' is not mappable to target}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(h) // expected-error {{threadprivate variables are not allowed in 'from' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(k), to(k) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}}
+#pragma omp target update from(t), from(t[:5]) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}}
+#pragma omp target update from(da)
+#pragma omp target update from(da[:4])
+
+#pragma omp target update from(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update from(x, c[:]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update from(x, (m+1)[2]) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update from(s7.i, s7.a[:3])
+#pragma omp target update from(s7.s6[1].aa[0:5])
+#pragma omp target update from(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update from(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update from(s7.p[:10])
+#pragma omp target update from(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'from' clause}}
+#pragma omp target update from(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+#pragma omp target data map(to: s7.i)
+ {
+#pragma omp target update from(s7.x)
+ }
+
+ return 0;
+}
+
+int main(int argc, char **argv) {
+ const int d = 5;
+ const int da[5] = { 0 };
+ S4 e(4);
+ S5 g(5);
+ int i, t[20];
+ int &j = i;
+ int *k = &j;
+ int x;
+ int y;
+ int from;
+ const int (&l)[5] = da;
+ int *m;
+ S7 s7;
+
+#pragma omp target update from // expected-error {{expected '(' after 'from'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from() // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update() // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(x)
+#pragma omp target update from(t[:i])
+#pragma omp target update from(S2::S2s)
+#pragma omp target update from(S2::S2sc)
+#pragma omp target update from(from)
+#pragma omp target update from(y x) // expected-error {{expected ',' or ')' in 'from' clause}}
+#pragma omp target update from(argc > 0 ? x : y) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target update from(ba) // expected-error 2 {{type 'S2' is not mappable to target}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(h) // expected-error {{threadprivate variables are not allowed in 'from' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(k), to(k) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}}
+#pragma omp target update from(t), from(t[:5]) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}}
+#pragma omp target update from(da)
+#pragma omp target update from(da[:4])
+
+#pragma omp target update from(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update from(x, c[:]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update from(x, (m+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update from(s7.i, s7.a[:3])
+#pragma omp target update from(s7.s6[1].aa[0:5])
+#pragma omp target update from(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update from(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update from(s7.p[:10])
+#pragma omp target update from(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'from' clause}}
+#pragma omp target update from(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+#pragma omp target data map(to: s7.i)
+ {
+#pragma omp target update from(s7.x)
+ }
+
+ return tmain<int, 3>(argc)+tmain<to, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
+}
+
Index: test/OpenMP/target_update_device_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_update_device_messages.cpp
@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+ return argc;
+}
+
+struct S1; // expected-note 2 {{declared here}}
+
+template <class T, class S>
+int tmain(T argc, S **argv) {
+ int i;
+#pragma omp target update to(i) device // expected-error {{expected '(' after 'device'}}
+#pragma omp target update to(i) device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(i) device () // expected-error {{expected expression}}
+#pragma omp target update to(i) device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(i) device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+#pragma omp target update from(i) device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+#pragma omp target update from(i) device (argc + argc)
+#pragma omp target update from(i) device (argc), device (argc+1) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'device' clause}}
+#pragma omp target update from(i) device (S1) // expected-error {{'S1' does not refer to a value}}
+#pragma omp target update from(i) device (3.14) // expected-error 2 {{expression must have integral or unscoped enumeration type, not 'double'}}
+#pragma omp target update from(i) device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}}
+}
+
+int main(int argc, char **argv) {
+ int j;
+#pragma omp target update to(j) device // expected-error {{expected '(' after 'device'}}
+#pragma omp target update from(j) device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(j) device () // expected-error {{expected expression}}
+#pragma omp target update from(j) device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(j) device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+#pragma omp target update from(j) device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+#pragma omp target update to(j) device (argc + argc)
+#pragma omp target update from(j) device (argc), device (argc+1) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'device' clause}}
+#pragma omp target update to(j) device (S1) // expected-error {{'S1' does not refer to a value}}
+#pragma omp target update from(j) device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}}
+#pragma omp target update to(j) device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}}
+
+ return tmain(argc, argv); // expected-note {{in instantiation of function template specialization 'tmain<int, char>' requested here}}
+}
Index: test/OpenMP/target_update_ast_print.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_update_ast_print.cpp
@@ -0,0 +1,52 @@
+// RUN: %clang_cc1 -verify -fopenmp -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+void foo() {}
+
+template <class T, class U>
+T foo(T targ, U uarg) {
+ static T a;
+ U b;
+ int l;
+#pragma omp target update to(a) if(l>5) device(l)
+
+#pragma omp target update from(b) if(l<5) device(l-1)
+ return a + targ + (T)b;
+}
+// CHECK: static int a;
+// CHECK-NEXT: float b;
+// CHECK-NEXT: int l;
+// CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l)
+// CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1)
+// CHECK: static char a;
+// CHECK-NEXT: float b;
+// CHECK-NEXT: int l;
+// CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l)
+// CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1)
+// CHECK: static T a;
+// CHECK-NEXT: U b;
+// CHECK-NEXT: int l;
+// CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l)
+// CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1)
+
+int main(int argc, char **argv) {
+ static int a;
+ int n;
+ float f;
+
+// CHECK: static int a;
+// CHECK-NEXT: int n;
+// CHECK-NEXT: float f;
+#pragma omp target update to(a) if(f>0.0) device(n)
+ // CHECK-NEXT: #pragma omp target update to(a) if(f > 0.) device(n)
+#pragma omp target update from(f) if(f<0.0) device(n+1)
+ // CHECK-NEXT: #pragma omp target update from(f) if(f < 0.) device(n + 1)
+ return foo(argc, f) + foo(argv[0][0], f) + a;
+}
+
+#endif
Index: test/OpenMP/target_parallel_map_messages.cpp
===================================================================
--- test/OpenMP/target_parallel_map_messages.cpp
+++ test/OpenMP/target_parallel_map_messages.cpp
@@ -126,7 +126,7 @@
foo();
#pragma omp target parallel map(e, g)
foo();
-#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
foo();
#pragma omp target parallel map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
foo();
@@ -229,7 +229,7 @@
foo();
#pragma omp target parallel map(e, g)
foo();
-#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
foo();
#pragma omp target parallel map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
foo();
Index: test/OpenMP/target_parallel_for_map_messages.cpp
===================================================================
--- test/OpenMP/target_parallel_for_map_messages.cpp
+++ test/OpenMP/target_parallel_for_map_messages.cpp
@@ -126,7 +126,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(e, g)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
@@ -230,7 +230,7 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(e, g)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
for (i = 0; i < argc; ++i) foo();
Index: test/OpenMP/target_map_messages.cpp
===================================================================
--- test/OpenMP/target_map_messages.cpp
+++ test/OpenMP/target_map_messages.cpp
@@ -40,7 +40,7 @@
#pragma omp target map(arg,a,d[:2]) // expected-error {{subscripted value is not an array or pointer}}
{}
- #pragma omp target map(to:ss) // expected-error {{threadprivate variables are not allowed in map clause}}
+ #pragma omp target map(to:ss) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
{}
#pragma omp target map(to:b,e)
@@ -239,7 +239,7 @@
{}
#pragma omp target map(r.C, t.C)
{}
- #pragma omp target map(r.A) // expected-error {{bit fields cannot be used to specify storage in a map clause}}
+ #pragma omp target map(r.A) // expected-error {{bit fields cannot be used to specify storage in a 'map' clause}}
{}
#pragma omp target map(r.Arr)
{}
@@ -407,7 +407,7 @@
#pragma omp target data map(S2::S2s)
#pragma omp target data map(S2::S2sc)
#pragma omp target data map(e, g)
-#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
#pragma omp target data map(k) map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
#pragma omp target map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
foo();
@@ -476,7 +476,7 @@
#pragma omp target data map(S2::S2s)
#pragma omp target data map(S2::S2sc)
#pragma omp target data map(e, g)
-#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
#pragma omp target data map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
#pragma omp target map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
foo();
Index: test/OpenMP/nesting_of_regions.cpp
===================================================================
--- test/OpenMP/nesting_of_regions.cpp
+++ test/OpenMP/nesting_of_regions.cpp
@@ -133,6 +133,10 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp parallel
+ {
+#pragma omp target update to(a)
+ }
// SIMD DIRECTIVE
#pragma omp simd
@@ -293,6 +297,10 @@
for (int j = 0; j < 10; ++j)
;
}
+#pragma omp simd
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update to(a) // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+ }
// FOR DIRECTIVE
#pragma omp for
@@ -476,6 +484,10 @@
for (int j = 0; j < 10; ++j)
;
}
+#pragma omp for
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update to(a)
+ }
// FOR SIMD DIRECTIVE
#pragma omp for simd
@@ -636,6 +648,11 @@
for (int j = 0; j < 10; ++j)
;
}
+#pragma omp for simd
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update to(a) // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+ bar();
+ }
// SECTIONS DIRECTIVE
#pragma omp sections
@@ -824,6 +841,10 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp sections
+ {
+#pragma omp target update to(a)
+ }
// SECTION DIRECTIVE
#pragma omp section // expected-error {{orphaned 'omp section' directives are prohibited, it must be closely nested to a sections region}}
@@ -1062,6 +1083,14 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp sections
+ {
+#pragma omp section
+ {
+ bar();
+#pragma omp target update to(a)
+ }
+ }
// SINGLE DIRECTIVE
#pragma omp single
@@ -1235,6 +1264,11 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp single
+ {
+#pragma omp target update to(a)
+ bar();
+ }
// MASTER DIRECTIVE
#pragma omp master
@@ -1408,6 +1442,11 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp master
+ {
+#pragma omp target update to(a)
+ bar();
+ }
// CRITICAL DIRECTIVE
#pragma omp critical
@@ -1595,6 +1634,11 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp critical
+ {
+#pragma omp target update to(a)
+ bar();
+ }
// PARALLEL FOR DIRECTIVE
#pragma omp parallel for
@@ -1783,6 +1827,10 @@
for (int j = 0; j < 10; ++j)
;
}
+#pragma omp parallel for
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update from(a)
+ }
// PARALLEL FOR SIMD DIRECTIVE
#pragma omp parallel for simd
@@ -1971,6 +2019,11 @@
for (int j = 0; j < 10; ++j)
;
}
+#pragma omp parallel for simd
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update to(a) // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+ bar();
+ }
// PARALLEL SECTIONS DIRECTIVE
#pragma omp parallel sections
@@ -2148,6 +2201,10 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp parallel sections
+ {
+#pragma omp target update to(a)
+ }
// TASK DIRECTIVE
#pragma omp task
@@ -2271,6 +2328,11 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp task
+ {
+#pragma omp target update to(a)
+ bar();
+ }
// ORDERED DIRECTIVE
#pragma omp ordered
@@ -2464,6 +2526,12 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp ordered
+ {
+ bar();
+#pragma omp target update to(a)
+ bar();
+ }
// ATOMIC DIRECTIVE
#pragma omp atomic
@@ -2678,6 +2746,13 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp atomic
+ // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
+ // expected-note@+1 {{expected an expression statement}}
+ {
+#pragma omp target update to(a) // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
+ bar();
+ }
// TARGET DIRECTIVE
#pragma omp target
@@ -2812,6 +2887,10 @@
{
#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target' region}}
}
+#pragma omp target
+ {
+#pragma omp target update to(a) // expected-error {{region cannot be nested inside 'target' region}}
+ }
// TARGET PARALLEL DIRECTIVE
#pragma omp target parallel
@@ -2946,6 +3025,10 @@
{
#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target parallel' region}}
}
+#pragma omp target parallel
+ {
+#pragma omp target update from(a) // expected-error {{region cannot be nested inside 'target parallel' region}}
+ }
// TARGET PARALLEL FOR DIRECTIVE
#pragma omp target parallel for
@@ -3134,6 +3217,10 @@
for (int j = 0; j < 10; ++j)
;
}
+#pragma omp target parallel for
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update from(a) // expected-error {{region cannot be nested inside 'target parallel for' region}}
+ }
// TEAMS DIRECTIVE
#pragma omp target
@@ -3297,6 +3384,11 @@
#pragma omp distribute
for (int j = 0; j < 10; ++j)
;
+#pragma omp target
+#pragma omp teams
+ {
+#pragma omp target update to(a) // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target update' directive into a parallel region?}}
+ }
// TASKLOOP DIRECTIVE
#pragma omp taskloop
@@ -3469,6 +3561,13 @@
for (int i = 0; i < 10; ++i)
++a;
}
+#pragma omp taskloop
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update to(a)
+ bar();
+ }
+
+
// DISTRIBUTE DIRECTIVE
#pragma omp target
#pragma omp teams
@@ -3686,6 +3785,13 @@
#pragma omp teams // expected-error {{region cannot be closely nested inside 'distribute' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update to(a) // expected-error {{region cannot be nested inside 'target' region}}
+ ++a;
+ }
}
void foo() {
@@ -3816,6 +3922,11 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp parallel
+ {
+#pragma omp target update to(a)
+ a++;
+ }
// SIMD DIRECTIVE
#pragma omp simd
@@ -3969,6 +4080,11 @@
for (int j = 0; j < 10; ++j)
;
}
+#pragma omp simd
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update from(a) // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+ a++;
+ }
// FOR DIRECTIVE
#pragma omp for
@@ -4142,6 +4258,11 @@
for (int j = 0; j < 10; ++j)
;
}
+#pragma omp for
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update from(a)
+ ++a;
+ }
// FOR SIMD DIRECTIVE
#pragma omp for simd
@@ -4295,6 +4416,11 @@
for (int j = 0; j < 10; ++j)
;
}
+#pragma omp for simd
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update from(a) // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+ ++a;
+ }
// SECTIONS DIRECTIVE
#pragma omp sections
@@ -4458,6 +4584,10 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp sections
+ {
+#pragma omp target update from(a)
+ }
// SECTION DIRECTIVE
#pragma omp section // expected-error {{orphaned 'omp section' directives are prohibited, it must be closely nested to a sections region}}
@@ -4706,6 +4836,14 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp sections
+ {
+#pragma omp section
+ {
+#pragma omp target update from(a)
+ a++;
+ }
+ }
// SINGLE DIRECTIVE
#pragma omp single
@@ -4869,6 +5007,11 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp single
+ {
+#pragma omp target update from(a)
+ a++;
+ }
// MASTER DIRECTIVE
#pragma omp master
@@ -5042,6 +5185,11 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp master
+ {
+#pragma omp target update from(a)
+ ++a;
+ }
// CRITICAL DIRECTIVE
#pragma omp critical
@@ -5234,6 +5382,11 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp critical
+ {
+#pragma omp target update from(a)
+ a++;
+ }
// PARALLEL FOR DIRECTIVE
#pragma omp parallel for
@@ -5422,6 +5575,11 @@
for (int j = 0; j < 10; ++j)
;
}
+#pragma omp parallel for
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update from(a)
+ a++;
+ }
// PARALLEL FOR SIMD DIRECTIVE
#pragma omp parallel for simd
@@ -5610,6 +5768,11 @@
for (int j = 0; j < 10; ++j)
;
}
+#pragma omp parallel for simd
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update from(a) // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+ a++;
+ }
// PARALLEL SECTIONS DIRECTIVE
#pragma omp parallel sections
@@ -5783,6 +5946,10 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp parallel sections
+ {
+#pragma omp target update from(a)
+ }
// TASK DIRECTIVE
#pragma omp task
@@ -5905,6 +6072,11 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp task
+ {
+#pragma omp target update from(a)
+ a++;
+ }
// ATOMIC DIRECTIVE
#pragma omp atomic
@@ -6119,6 +6291,12 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp atomic
+ // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
+ // expected-note@+1 {{expected an expression statement}}
+ {
+#pragma omp target update // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
+ }
// TARGET DIRECTIVE
#pragma omp target
@@ -6253,6 +6431,13 @@
for (int i = 0; i < 10; ++i)
;
}
+#pragma omp atomic
+ // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
+ // expected-note@+1 {{expected an expression statement}}
+ {
+#pragma omp target update from(a) // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
+ a++;
+ }
// TARGET PARALLEL DIRECTIVE
#pragma omp target parallel
@@ -6387,6 +6572,11 @@
{
#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target parallel' region}}
}
+#pragma omp target parallel
+ {
+#pragma omp target update from(a) // expected-error {{region cannot be nested inside 'target parallel' region}}
+ }
+
// TARGET PARALLEL FOR DIRECTIVE
#pragma omp target parallel for
@@ -6575,6 +6765,11 @@
for (int j = 0; j < 10; ++j)
;
}
+#pragma omp target parallel for
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update from(a) // expected-error {{region cannot be nested inside 'target parallel for' region}}
+ a++;
+ }
// TEAMS DIRECTIVE
#pragma omp target
@@ -6736,6 +6931,12 @@
#pragma omp distribute
for (int j = 0; j < 10; ++j)
;
+#pragma omp target
+#pragma omp teams
+ {
+#pragma omp target update from(a) // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target update' directive into a parallel region?}}
+ ++a;
+ }
// TASKLOOP DIRECTIVE
#pragma omp taskloop
@@ -6908,6 +7109,11 @@
for (int i = 0; i < 10; ++i)
++a;
}
+#pragma omp taskloop
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update from(a)
+ ++a;
+ }
// DISTRIBUTE DIRECTIVE
#pragma omp target
@@ -7127,4 +7333,11 @@
#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target' region}}
++a;
}
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target update from(a) // expected-error {{region cannot be nested inside 'target' region}}
+ ++a;
+ }
}
Index: lib/StaticAnalyzer/Core/ExprEngine.cpp
===================================================================
--- lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -835,6 +835,7 @@
case Stmt::OMPTargetExitDataDirectiveClass:
case Stmt::OMPTargetParallelDirectiveClass:
case Stmt::OMPTargetParallelForDirectiveClass:
+ case Stmt::OMPTargetUpdateDirectiveClass:
case Stmt::OMPTeamsDirectiveClass:
case Stmt::OMPCancellationPointDirectiveClass:
case Stmt::OMPCancelDirectiveClass:
Index: lib/Serialization/ASTWriterStmt.cpp
===================================================================
--- lib/Serialization/ASTWriterStmt.cpp
+++ lib/Serialization/ASTWriterStmt.cpp
@@ -2069,6 +2069,20 @@
Writer->Writer.AddSourceLocation(C->getDefaultmapKindLoc(), Record);
}
+void OMPClauseWriter::VisitOMPToClause(OMPToClause *C) {
+ Record.push_back(C->varlist_size());
+ Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record);
+ for (auto *VE : C->varlists())
+ Writer->Writer.AddStmt(VE);
+}
+
+void OMPClauseWriter::VisitOMPFromClause(OMPFromClause *C) {
+ Record.push_back(C->varlist_size());
+ Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record);
+ for (auto *VE : C->varlists())
+ Writer->Writer.AddStmt(VE);
+}
+
//===----------------------------------------------------------------------===//
// OpenMP Directives.
//===----------------------------------------------------------------------===//
@@ -2346,6 +2360,13 @@
Code = serialization::STMT_OMP_DISTRIBUTE_DIRECTIVE;
}
+void ASTStmtWriter::VisitOMPTargetUpdateDirective(OMPTargetUpdateDirective *D) {
+ VisitStmt(D);
+ Record.push_back(D->getNumClauses());
+ VisitOMPExecutableDirective(D);
+ Code = serialization::STMT_OMP_TARGET_UPDATE_DIRECTIVE;
+}
+
//===----------------------------------------------------------------------===//
// ASTWriter Implementation
//===----------------------------------------------------------------------===//
Index: lib/Serialization/ASTReaderStmt.cpp
===================================================================
--- lib/Serialization/ASTReaderStmt.cpp
+++ lib/Serialization/ASTReaderStmt.cpp
@@ -1885,6 +1885,12 @@
case OMPC_defaultmap:
C = new (Context) OMPDefaultmapClause();
break;
+ case OMPC_to:
+ C = OMPToClause::CreateEmpty(Context, Record[Idx++]);
+ break;
+ case OMPC_from:
+ C = OMPFromClause::CreateEmpty(Context, Record[Idx++]);
+ break;
}
Visit(C);
C->setLocStart(Reader->ReadSourceLocation(Record, Idx));
@@ -2280,6 +2286,28 @@
C->setDefaultmapKindLoc(Reader->ReadSourceLocation(Record, Idx));
}
+void OMPClauseReader::VisitOMPToClause(OMPToClause *C) {
+ C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx));
+ unsigned NumVars = C->varlist_size();
+ SmallVector<Expr *, 16> Vars;
+ Vars.reserve(NumVars);
+ for (unsigned i = 0; i != NumVars; ++i) {
+ Vars.push_back(Reader->Reader.ReadSubExpr());
+ }
+ C->setVarRefs(Vars);
+}
+
+void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) {
+ C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx));
+ unsigned NumVars = C->varlist_size();
+ SmallVector<Expr *, 16> Vars;
+ Vars.reserve(NumVars);
+ for (unsigned i = 0; i != NumVars; ++i) {
+ Vars.push_back(Reader->Reader.ReadSubExpr());
+ }
+ C->setVarRefs(Vars);
+}
+
//===----------------------------------------------------------------------===//
// OpenMP Directives.
//===----------------------------------------------------------------------===//
@@ -2545,6 +2573,12 @@
VisitOMPLoopDirective(D);
}
+void ASTStmtReader::VisitOMPTargetUpdateDirective(OMPTargetUpdateDirective *D) {
+ VisitStmt(D);
+ ++Idx;
+ VisitOMPExecutableDirective(D);
+}
+
//===----------------------------------------------------------------------===//
// ASTReader Implementation
//===----------------------------------------------------------------------===//
@@ -3177,6 +3211,11 @@
break;
}
+ case STMT_OMP_TARGET_UPDATE_DIRECTIVE:
+ S = OMPTargetUpdateDirective::CreateEmpty(
+ Context, Record[ASTStmtReader::NumStmtFields], Empty);
+ break;
+
case STMT_OMP_TEAMS_DIRECTIVE:
S = OMPTeamsDirective::CreateEmpty(
Context, Record[ASTStmtReader::NumStmtFields], Empty);
Index: lib/Sema/TreeTransform.h
===================================================================
--- lib/Sema/TreeTransform.h
+++ lib/Sema/TreeTransform.h
@@ -1750,6 +1750,29 @@
Kind, ChunkSize, StartLoc, LParenLoc, KindLoc, CommaLoc, EndLoc);
}
+ /// \brief Build a new OpenMP 'to' clause.
+ ///
+ /// By default, performs semantic analysis to build the new statement.
+ /// Subclasses may override this routine to provide different behavior.
+ OMPClause *RebuildOMPToClause(ArrayRef<Expr *> VarList,
+ SourceLocation StartLoc,
+ SourceLocation LParenLoc,
+ SourceLocation EndLoc) {
+ return getSema().ActOnOpenMPToClause(VarList, StartLoc, LParenLoc, EndLoc);
+ }
+
+ /// \brief Build a new OpenMP 'from' clause.
+ ///
+ /// By default, performs semantic analysis to build the new statement.
+ /// Subclasses may override this routine to provide different behavior.
+ OMPClause *RebuildOMPFromClause(ArrayRef<Expr *> VarList,
+ SourceLocation StartLoc,
+ SourceLocation LParenLoc,
+ SourceLocation EndLoc) {
+ return getSema().ActOnOpenMPFromClause(VarList, StartLoc, LParenLoc,
+ EndLoc);
+ }
+
/// \brief Rebuild the operand to an Objective-C \@synchronized statement.
///
/// By default, performs semantic analysis to build the new statement.
@@ -7468,6 +7491,17 @@
}
template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOMPTargetUpdateDirective(
+ OMPTargetUpdateDirective *D) {
+ DeclarationNameInfo DirName;
+ getDerived().getSema().StartOpenMPDSABlock(OMPD_target_update, DirName,
+ nullptr, D->getLocStart());
+ StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
+ getDerived().getSema().EndOpenMPDSABlock(Res.get());
+ return Res;
+}
+
+template <typename Derived>
StmtResult
TreeTransform<Derived>::TransformOMPTeamsDirective(OMPTeamsDirective *D) {
DeclarationNameInfo DirName;
@@ -7996,6 +8030,34 @@
return C;
}
+template <typename Derived>
+OMPClause *TreeTransform<Derived>::TransformOMPToClause(OMPToClause *C) {
+ llvm::SmallVector<Expr *, 16> Vars;
+ Vars.reserve(C->varlist_size());
+ for (auto *VE : C->varlists()) {
+ ExprResult EVar = getDerived().TransformExpr(cast<Expr>(VE));
+ if (EVar.isInvalid())
+ return 0;
+ Vars.push_back(EVar.get());
+ }
+ return getDerived().RebuildOMPToClause(Vars, C->getLocStart(),
+ C->getLParenLoc(), C->getLocEnd());
+}
+
+template <typename Derived>
+OMPClause *TreeTransform<Derived>::TransformOMPFromClause(OMPFromClause *C) {
+ llvm::SmallVector<Expr *, 16> Vars;
+ Vars.reserve(C->varlist_size());
+ for (auto *VE : C->varlists()) {
+ ExprResult EVar = getDerived().TransformExpr(cast<Expr>(VE));
+ if (EVar.isInvalid())
+ return 0;
+ Vars.push_back(EVar.get());
+ }
+ return getDerived().RebuildOMPFromClause(Vars, C->getLocStart(),
+ C->getLParenLoc(), C->getLocEnd());
+}
+
//===----------------------------------------------------------------------===//
// Expression transformation
//===----------------------------------------------------------------------===//
Index: lib/Sema/SemaOpenMP.cpp
===================================================================
--- lib/Sema/SemaOpenMP.cpp
+++ lib/Sema/SemaOpenMP.cpp
@@ -1700,6 +1700,7 @@
case OMPD_target_enter_data:
case OMPD_target_exit_data:
case OMPD_declare_reduction:
+ case OMPD_target_update:
llvm_unreachable("OpenMP Directive is not allowed");
case OMPD_unknown:
llvm_unreachable("Unknown OpenMP directive");
@@ -3140,6 +3141,12 @@
Res = ActOnOpenMPDistributeDirective(ClausesWithImplicit, AStmt, StartLoc,
EndLoc, VarsWithInheritedDSA);
break;
+ case OMPD_target_update:
+ assert(!AStmt && "Statement is not allowed for target update");
+ Res =
+ ActOnOpenMPTargetUpdateDirective(ClausesWithImplicit, StartLoc, EndLoc);
+ AllowedNameModifiers.push_back(OMPD_target_update);
+ break;
case OMPD_threadprivate:
case OMPD_declare_reduction:
llvm_unreachable("OpenMP Directive is not allowed");
@@ -6059,6 +6066,21 @@
return OMPTargetExitDataDirective::Create(Context, StartLoc, EndLoc, Clauses);
}
+StmtResult Sema::ActOnOpenMPTargetUpdateDirective(ArrayRef<OMPClause *> Clauses,
+ SourceLocation StartLoc,
+ SourceLocation EndLoc) {
+ bool seenMotionClause = false;
+ for (auto *C : Clauses) {
+ if (C->getClauseKind() == OMPC_to || C->getClauseKind() == OMPC_from)
+ seenMotionClause = true;
+ }
+ if (!seenMotionClause) {
+ Diag(StartLoc, diag::err_omp_at_least_one_motion_clause_required);
+ return StmtError();
+ }
+ return OMPTargetUpdateDirective::Create(Context, StartLoc, EndLoc, Clauses);
+}
+
StmtResult Sema::ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
Stmt *AStmt, SourceLocation StartLoc,
SourceLocation EndLoc) {
@@ -6312,6 +6334,8 @@
case OMPC_nogroup:
case OMPC_dist_schedule:
case OMPC_defaultmap:
+ case OMPC_to:
+ case OMPC_from:
case OMPC_unknown:
llvm_unreachable("Clause is not allowed.");
}
@@ -6597,6 +6621,8 @@
case OMPC_hint:
case OMPC_dist_schedule:
case OMPC_defaultmap:
+ case OMPC_to:
+ case OMPC_from:
case OMPC_unknown:
llvm_unreachable("Clause is not allowed.");
}
@@ -6747,6 +6773,8 @@
case OMPC_nogroup:
case OMPC_num_tasks:
case OMPC_hint:
+ case OMPC_to:
+ case OMPC_from:
case OMPC_unknown:
llvm_unreachable("Clause is not allowed.");
}
@@ -6934,6 +6962,8 @@
case OMPC_hint:
case OMPC_dist_schedule:
case OMPC_defaultmap:
+ case OMPC_to:
+ case OMPC_from:
case OMPC_unknown:
llvm_unreachable("Clause is not allowed.");
}
@@ -7048,6 +7078,12 @@
DepLinMapLoc, ColonLoc, VarList, StartLoc,
LParenLoc, EndLoc);
break;
+ case OMPC_to:
+ Res = ActOnOpenMPToClause(VarList, StartLoc, LParenLoc, EndLoc);
+ break;
+ case OMPC_from:
+ Res = ActOnOpenMPFromClause(VarList, StartLoc, LParenLoc, EndLoc);
+ break;
case OMPC_if:
case OMPC_final:
case OMPC_num_threads:
@@ -9195,7 +9231,8 @@
// Return the expression of the base of the map clause or null if it cannot
// be determined and do all the necessary checks to see if the expression is
// valid as a standalone map clause expression.
-static Expr *CheckMapClauseExpressionBase(Sema &SemaRef, Expr *E) {
+static Expr *CheckMapClauseExpressionBase(Sema &SemaRef, Expr *E,
+ OpenMPClauseKind CKind) {
SourceLocation ELoc = E->getExprLoc();
SourceRange ERange = E->getSourceRange();
@@ -9277,8 +9314,8 @@
// A bit-field cannot appear in a map clause.
//
if (FD->isBitField()) {
- SemaRef.Diag(ELoc, diag::err_omp_bit_fields_forbidden_in_map_clause)
- << CurE->getSourceRange();
+ SemaRef.Diag(ELoc, diag::err_omp_bit_fields_forbidden_in_clause) <<
+ CurE->getSourceRange() << getOpenMPClauseName(CKind);
break;
}
@@ -9386,7 +9423,8 @@
// Return true if expression E associated with value VD has conflicts with other
// map information.
static bool CheckMapConflicts(Sema &SemaRef, DSAStackTy *DSAS, ValueDecl *VD,
- Expr *E, bool CurrentRegionOnly) {
+ Expr *E, bool CurrentRegionOnly,
+ OpenMPClauseKind CKind) {
assert(VD && E);
// Types used to organize the components of a valid map clause.
@@ -9506,7 +9544,14 @@
// other, it means they are sharing storage.
if (CI == CE && SI == SE) {
if (CurrentRegionOnly) {
- SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange;
+ if (CKind == OMPC_map)
+ SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange;
+ else {
+ assert(CKind == OMPC_to || CKind == OMPC_from);
+ SemaRef.Diag(ELoc,
+ diag::err_omp_once_referenced_in_target_update)
+ << ERange;
+ }
SemaRef.Diag(RE->getExprLoc(), diag::note_used_here)
<< RE->getSourceRange();
return true;
@@ -9560,9 +9605,16 @@
//
// An expression is a subset of the other.
if (CurrentRegionOnly && (CI == CE || SI == SE)) {
- SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange;
+ if (CKind == OMPC_map)
+ SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange;
+ else {
+ assert(CKind == OMPC_to || CKind == OMPC_from);
+ SemaRef.Diag(ELoc,
+ diag::err_omp_once_referenced_in_target_update)
+ << ERange;
+ }
SemaRef.Diag(RE->getExprLoc(), diag::note_used_here)
- << RE->getSourceRange();
+ << RE->getSourceRange();
return true;
}
@@ -9605,16 +9657,15 @@
return FoundError;
}
-OMPClause *
-Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
- OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
- SourceLocation MapLoc, SourceLocation ColonLoc,
- ArrayRef<Expr *> VarList, SourceLocation StartLoc,
- SourceLocation LParenLoc, SourceLocation EndLoc) {
- SmallVector<Expr *, 4> Vars;
-
+static void checkOpenMPToFromMapVars(Sema &SemaRef, DSAStackTy *DSAS,
+ ArrayRef<Expr *> VarList,
+ SmallVector<Expr *, 4> &Vars,
+ SourceLocation StartLoc,
+ OpenMPClauseKind CKind,
+ OpenMPMapClauseKind MapType,
+ bool IsMapTypeImplicit) {
for (auto &RE : VarList) {
- assert(RE && "Null expr in omp map");
+ assert(RE && "Null expr in omp to/from/map clause");
if (isa<DependentScopeDeclRefExpr>(RE)) {
// It will be analyzed later.
Vars.push_back(RE);
@@ -9636,106 +9687,126 @@
auto *SimpleExpr = RE->IgnoreParenCasts();
if (!RE->IgnoreParenImpCasts()->isLValue()) {
- Diag(ELoc, diag::err_omp_expected_named_var_member_or_array_expression)
+ SemaRef.Diag(ELoc,
+ diag::err_omp_expected_named_var_member_or_array_expression)
<< RE->getSourceRange();
continue;
}
// Obtain the array or member expression bases if required.
- auto *BE = CheckMapClauseExpressionBase(*this, SimpleExpr);
+ auto *BE = CheckMapClauseExpressionBase(SemaRef, SimpleExpr, CKind);
if (!BE)
continue;
// If the base is a reference to a variable, we rely on that variable for
- // the following checks. If it is a 'this' expression we rely on the field.
+ // the following checks. if it is a 'this' expression we rely on the field.
ValueDecl *D = nullptr;
if (auto *DRE = dyn_cast<DeclRefExpr>(BE)) {
D = DRE->getDecl();
} else {
auto *ME = cast<MemberExpr>(BE);
assert(isa<CXXThisExpr>(ME->getBase()) && "Unexpected expression!");
D = ME->getMemberDecl();
}
- assert(D && "Null decl on map clause.");
+ assert(D && "Null decl on to/from/map clause.");
auto *VD = dyn_cast<VarDecl>(D);
auto *FD = dyn_cast<FieldDecl>(D);
assert((VD || FD) && "Only variables or fields are expected here!");
(void)FD;
// OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.10]
- // threadprivate variables cannot appear in a map clause.
- if (VD && DSAStack->isThreadPrivate(VD)) {
- auto DVar = DSAStack->getTopDSA(VD, false);
- Diag(ELoc, diag::err_omp_threadprivate_in_map);
- ReportOriginalDSA(*this, DSAStack, VD, DVar);
+ // threadprivate variables cannot appear in a map clause.
+ // OpenMP 4.5 [2.10.5, target update Construct]
+ // threadprivate variables cannot appear in a from clause.
+ if (VD && DSAS->isThreadPrivate(VD)) {
+ auto DVar = DSAS->getTopDSA(VD, false);
+ SemaRef.Diag(ELoc, diag::err_omp_threadprivate_in_clause)
+ << getOpenMPClauseName(CKind);
+ ReportOriginalDSA(SemaRef, DSAS, VD, DVar);
continue;
}
- // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9]
// A list item cannot appear in both a map clause and a data-sharing
// attribute clause on the same construct.
//
// TODO: Implement this check - it cannot currently be tested because of
// missing implementation of the other data sharing clauses in target
// directives.
- // Check conflicts with other map clause expressions. We check the conflicts
- // with the current construct separately from the enclosing data
- // environment, because the restrictions are different.
- if (CheckMapConflicts(*this, DSAStack, D, SimpleExpr,
- /*CurrentRegionOnly=*/true))
+ // Check conflicts with other map, to or from clause expressions. We check
+ // the conflicts with the current construct separately from the enclosing
+ // data environment, because the restrictions are different.
+ if (CheckMapConflicts(SemaRef, DSAS, D, SimpleExpr,
+ /*CurrentRegionOnly=*/true, CKind))
break;
- if (CheckMapConflicts(*this, DSAStack, D, SimpleExpr,
- /*CurrentRegionOnly=*/false))
+ if (CKind == OMPC_map &&
+ CheckMapConflicts(SemaRef, DSAS, D, SimpleExpr,
+ /*CurrentRegionOnly=*/false, CKind))
break;
+ // OpenMP 4.5 [2.10.5, target update Construct]
// OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1]
- // If the type of a list item is a reference to a type T then the type will
- // be considered to be T for all purposes of this clause.
- QualType Type = D->getType();
+ // If the type of a list item is a reference to a type T then the type will
+ // be considered to be T for all purpose of this clause.
+ auto Type = D->getType();
if (Type->isReferenceType())
Type = Type->getPointeeType();
+ // OpenMP 4.5 [2.10.5, target update Construct, Restrictions, p.4]
+ // A list item in a to or from clause must have a mappable type.
// OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9]
- // A list item must have a mappable type.
- if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), *this,
- DSAStack, Type))
+ // A list item must have a mappable type.
+ if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), SemaRef,
+ DSAS, Type))
continue;
- // target enter data
- // OpenMP [2.10.2, Restrictions, p. 99]
- // A map-type must be specified in all map clauses and must be either
- // to or alloc.
- OpenMPDirectiveKind DKind = DSAStack->getCurrentDirective();
- if (DKind == OMPD_target_enter_data &&
- !(MapType == OMPC_MAP_to || MapType == OMPC_MAP_alloc)) {
- Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive)
- << (IsMapTypeImplicit ? 1 : 0)
- << getOpenMPSimpleClauseTypeName(OMPC_map, MapType)
- << getOpenMPDirectiveName(DKind);
- continue;
- }
+ if (CKind == OMPC_map) {
+ // target enter data
+ // OpenMP [2.10.2, Restrictions, p. 99]
+ // A map-type must be specified in all map clauses and must be either
+ // to or alloc.
+ OpenMPDirectiveKind DKind = DSAS->getCurrentDirective();
+ if (DKind == OMPD_target_enter_data &&
+ !(MapType == OMPC_MAP_to || MapType == OMPC_MAP_alloc)) {
+ SemaRef.Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive)
+ << (IsMapTypeImplicit ? 1 : 0)
+ << getOpenMPSimpleClauseTypeName(OMPC_map, MapType)
+ << getOpenMPDirectiveName(DKind);
+ continue;
+ }
- // target exit_data
- // OpenMP [2.10.3, Restrictions, p. 102]
- // A map-type must be specified in all map clauses and must be either
- // from, release, or delete.
- DKind = DSAStack->getCurrentDirective();
- if (DKind == OMPD_target_exit_data &&
- !(MapType == OMPC_MAP_from || MapType == OMPC_MAP_release ||
- MapType == OMPC_MAP_delete)) {
- Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive)
- << (IsMapTypeImplicit ? 1 : 0)
- << getOpenMPSimpleClauseTypeName(OMPC_map, MapType)
- << getOpenMPDirectiveName(DKind);
- continue;
+ // target exit_data
+ // OpenMP [2.10.3, Restrictions, p. 102]
+ // A map-type must be specified in all map clauses and must be either
+ // from, release, or delete.
+ DKind = DSAS->getCurrentDirective();
+ if (DKind == OMPD_target_exit_data &&
+ !(MapType == OMPC_MAP_from || MapType == OMPC_MAP_release ||
+ MapType == OMPC_MAP_delete)) {
+ SemaRef.Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive)
+ << (IsMapTypeImplicit ? 1 : 0)
+ << getOpenMPSimpleClauseTypeName(OMPC_map, MapType)
+ << getOpenMPDirectiveName(DKind);
+ continue;
+ }
}
Vars.push_back(RE);
- DSAStack->addExprToVarMapInfo(D, RE);
+ DSAS->addExprToVarMapInfo(D, RE);
}
+}
+
+OMPClause *
+Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
+ OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
+ SourceLocation MapLoc, SourceLocation ColonLoc,
+ ArrayRef<Expr *> VarList, SourceLocation StartLoc,
+ SourceLocation LParenLoc, SourceLocation EndLoc) {
+ SmallVector<Expr *, 4> Vars;
+ checkOpenMPToFromMapVars(*this, DSAStack, VarList, Vars, StartLoc, OMPC_map,
+ MapType, IsMapTypeImplicit);
// We need to produce a map clause even if we don't have variables so that
// other diagnostics related with non-existing map clauses are accurate.
@@ -10146,3 +10217,29 @@
return new (Context)
OMPDefaultmapClause(StartLoc, LParenLoc, MLoc, KindLoc, EndLoc, Kind, M);
}
+
+OMPClause *Sema::ActOnOpenMPToClause(ArrayRef<Expr *> VarList,
+ SourceLocation StartLoc,
+ SourceLocation LParenLoc,
+ SourceLocation EndLoc) {
+ SmallVector<Expr *, 4> Vars;
+ checkOpenMPToFromMapVars(*this, DSAStack, VarList, Vars, StartLoc, OMPC_to,
+ OMPC_MAP_unknown, /*IsMapTypeImplicit=*/false);
+ if (Vars.empty())
+ return nullptr;
+
+ return OMPToClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars);
+}
+
+OMPClause *Sema::ActOnOpenMPFromClause(ArrayRef<Expr *> VarList,
+ SourceLocation StartLoc,
+ SourceLocation LParenLoc,
+ SourceLocation EndLoc) {
+ SmallVector<Expr *, 4> Vars;
+ checkOpenMPToFromMapVars(*this, DSAStack, VarList, Vars, StartLoc, OMPC_from,
+ OMPC_MAP_unknown, /*IsMapTypeImplicit=*/false);
+ if (Vars.empty())
+ return nullptr;
+
+ return OMPFromClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars);
+}
Index: lib/Parse/ParseOpenMP.cpp
===================================================================
--- lib/Parse/ParseOpenMP.cpp
+++ lib/Parse/ParseOpenMP.cpp
@@ -36,7 +36,8 @@
OMPD_point,
OMPD_reduction,
OMPD_target_enter,
- OMPD_target_exit
+ OMPD_target_exit,
+ OMPD_update
};
} // namespace
@@ -55,6 +56,7 @@
.Case("exit", OMPD_exit)
.Case("point", OMPD_point)
.Case("reduction", OMPD_reduction)
+ .Case("update", OMPD_update)
.Default(OMPD_unknown);
}
@@ -70,6 +72,7 @@
{ OMPD_target, OMPD_exit, OMPD_target_exit },
{ OMPD_target_enter, OMPD_data, OMPD_target_enter_data },
{ OMPD_target_exit, OMPD_data, OMPD_target_exit_data },
+ { OMPD_target, OMPD_update, OMPD_target_update },
{ OMPD_for, OMPD_simd, OMPD_for_simd },
{ OMPD_parallel, OMPD_for, OMPD_parallel_for },
{ OMPD_parallel_for, OMPD_simd, OMPD_parallel_for_simd },
@@ -398,6 +401,7 @@
case OMPD_taskloop:
case OMPD_taskloop_simd:
case OMPD_distribute:
+ case OMPD_target_update:
Diag(Tok, diag::err_omp_unexpected_directive)
<< getOpenMPDirectiveName(DKind);
break;
@@ -426,9 +430,10 @@
/// 'parallel for' | 'parallel sections' | 'task' | 'taskyield' |
/// 'barrier' | 'taskwait' | 'flush' | 'ordered' | 'atomic' |
/// 'for simd' | 'parallel for simd' | 'target' | 'target data' |
-/// 'taskgroup' | 'teams' | 'taskloop' | 'taskloop simd' |
+/// 'taskgroup' | 'teams' | 'taskloop' | 'taskloop simd' |
/// 'distribute' | 'target enter data' | 'target exit data' |
-/// 'target parallel' | 'target parallel for' {clause}
+/// 'target parallel' | 'target parallel for' |
+/// 'target update' {clause}
/// annot_pragma_openmp_end
///
StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
@@ -501,6 +506,7 @@
case OMPD_cancel:
case OMPD_target_enter_data:
case OMPD_target_exit_data:
+ case OMPD_target_update:
if (Allowed == ACK_StatementsOpenMPNonStandalone) {
Diag(Tok, diag::err_omp_immediate_directive)
<< getOpenMPDirectiveName(DKind) << 0;
@@ -706,7 +712,8 @@
/// update-clause | capture-clause | seq_cst-clause | device-clause |
/// simdlen-clause | threads-clause | simd-clause | num_teams-clause |
/// thread_limit-clause | priority-clause | grainsize-clause |
-/// nogroup-clause | num_tasks-clause | hint-clause
+/// nogroup-clause | num_tasks-clause | hint-clause | to-clause |
+/// from-clause
///
OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
OpenMPClauseKind CKind, bool FirstClause) {
@@ -830,6 +837,8 @@
case OMPC_flush:
case OMPC_depend:
case OMPC_map:
+ case OMPC_to:
+ case OMPC_from:
Clause = ParseOpenMPVarListClause(DKind, CKind);
break;
case OMPC_unknown:
@@ -1147,8 +1156,7 @@
TemplateKWLoc, ReductionId);
}
-/// \brief Parsing of OpenMP clause 'private', 'firstprivate', 'lastprivate',
-/// 'shared', 'copyin', 'copyprivate', 'flush' or 'reduction'.
+/// \brief Parsing of OpenMP clauses
///
/// private-clause:
/// 'private' '(' list ')'
@@ -1173,6 +1181,10 @@
/// map-clause:
/// 'map' '(' [ [ always , ]
/// to | from | tofrom | alloc | release | delete ':' ] list ')';
+/// to-clause:
+/// 'to' '(' list ')'
+/// from-clause:
+/// 'from' '(' list ')'
///
/// For 'linear' clause linear-list may have the following forms:
/// list
Index: lib/CodeGen/CodeGenFunction.h
===================================================================
--- lib/CodeGen/CodeGenFunction.h
+++ lib/CodeGen/CodeGenFunction.h
@@ -2354,6 +2354,7 @@
void EmitOMPTargetDataDirective(const OMPTargetDataDirective &S);
void EmitOMPTargetEnterDataDirective(const OMPTargetEnterDataDirective &S);
void EmitOMPTargetExitDataDirective(const OMPTargetExitDataDirective &S);
+ void EmitOMPTargetUpdateDirective(const OMPTargetUpdateDirective &S);
void EmitOMPTargetParallelDirective(const OMPTargetParallelDirective &S);
void
EmitOMPTargetParallelForDirective(const OMPTargetParallelForDirective &S);
Index: lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- lib/CodeGen/CGStmtOpenMP.cpp
+++ lib/CodeGen/CGStmtOpenMP.cpp
@@ -2824,6 +2824,8 @@
case OMPC_hint:
case OMPC_dist_schedule:
case OMPC_defaultmap:
+ case OMPC_to:
+ case OMPC_from:
llvm_unreachable("Clause is not allowed in 'omp atomic'.");
}
}
@@ -3038,3 +3040,9 @@
[&CS](CodeGenFunction &CGF) { CGF.EmitStmt(CS->getCapturedStmt()); });
}
+// Generate the instructions for '#pragma omp target update' directive.
+void CodeGenFunction::EmitOMPTargetUpdateDirective(
+ const OMPTargetUpdateDirective &S) {
+ // TODO: codegen for target update
+}
+
Index: lib/CodeGen/CGStmt.cpp
===================================================================
--- lib/CodeGen/CGStmt.cpp
+++ lib/CodeGen/CGStmt.cpp
@@ -277,6 +277,9 @@
case Stmt::OMPDistributeDirectiveClass:
EmitOMPDistributeDirective(cast<OMPDistributeDirective>(*S));
break;
+ case Stmt::OMPTargetUpdateDirectiveClass:
+ EmitOMPTargetUpdateDirective(cast<OMPTargetUpdateDirective>(*S));
+ break;
}
}
Index: lib/Basic/OpenMPKinds.cpp
===================================================================
--- lib/Basic/OpenMPKinds.cpp
+++ lib/Basic/OpenMPKinds.cpp
@@ -158,6 +158,8 @@
case OMPC_nogroup:
case OMPC_num_tasks:
case OMPC_hint:
+ case OMPC_to:
+ case OMPC_from:
break;
}
llvm_unreachable("Invalid OpenMP simple clause kind");
@@ -292,6 +294,8 @@
case OMPC_nogroup:
case OMPC_num_tasks:
case OMPC_hint:
+ case OMPC_to:
+ case OMPC_from:
break;
}
llvm_unreachable("Invalid OpenMP simple clause kind");
@@ -475,6 +479,16 @@
break;
}
break;
+ case OMPD_target_update:
+ switch (CKind) {
+#define OPENMP_TARGET_UPDATE_CLAUSE(Name) \
+ case OMPC_##Name: \
+ return true;
+#include "clang/Basic/OpenMPKinds.def"
+ default:
+ break;
+ }
+ break;
case OMPD_teams:
switch (CKind) {
#define OPENMP_TEAMS_CLAUSE(Name) \
@@ -596,7 +610,7 @@
bool clang::isOpenMPTargetDataManagementDirective(OpenMPDirectiveKind DKind) {
// TODO add target update directive check.
return DKind == OMPD_target_data || DKind == OMPD_target_enter_data ||
- DKind == OMPD_target_exit_data;
+ DKind == OMPD_target_exit_data || DKind == OMPD_target_update;
}
bool clang::isOpenMPTeamsDirective(OpenMPDirectiveKind DKind) {
Index: lib/AST/StmtProfile.cpp
===================================================================
--- lib/AST/StmtProfile.cpp
+++ lib/AST/StmtProfile.cpp
@@ -491,6 +491,12 @@
void OMPClauseProfiler::VisitOMPHintClause(const OMPHintClause *C) {
Profiler->VisitStmt(C->getHint());
}
+void OMPClauseProfiler::VisitOMPToClause(const OMPToClause *C) {
+ VisitOMPClauseList(C);
+}
+void OMPClauseProfiler::VisitOMPFromClause(const OMPFromClause *C) {
+ VisitOMPClauseList(C);
+}
}
void
@@ -656,6 +662,11 @@
void OMPClauseProfiler::VisitOMPDefaultmapClause(const OMPDefaultmapClause *) {}
+void StmtProfiler::VisitOMPTargetUpdateDirective(
+ const OMPTargetUpdateDirective *S) {
+ VisitOMPExecutableDirective(S);
+}
+
void StmtProfiler::VisitExpr(const Expr *S) {
VisitStmt(S);
}
Index: lib/AST/StmtPrinter.cpp
===================================================================
--- lib/AST/StmtPrinter.cpp
+++ lib/AST/StmtPrinter.cpp
@@ -912,6 +912,22 @@
}
}
+void OMPClausePrinter::VisitOMPToClause(OMPToClause *Node) {
+ if (!Node->varlist_empty()) {
+ OS << "to";
+ VisitOMPClauseList(Node, '(');
+ OS << ")";
+ }
+}
+
+void OMPClausePrinter::VisitOMPFromClause(OMPFromClause *Node) {
+ if (!Node->varlist_empty()) {
+ OS << "from";
+ VisitOMPClauseList(Node, '(');
+ OS << ")";
+ }
+}
+
void OMPClausePrinter::VisitOMPDistScheduleClause(OMPDistScheduleClause *Node) {
OS << "dist_schedule(" << getOpenMPSimpleClauseTypeName(
OMPC_dist_schedule, Node->getDistScheduleKind());
@@ -1131,6 +1147,11 @@
PrintOMPExecutableDirective(Node);
}
+void StmtPrinter::VisitOMPTargetUpdateDirective(OMPTargetUpdateDirective *Node) {
+ Indent() << "#pragma omp target update ";
+ PrintOMPExecutableDirective(Node);
+}
+
//===----------------------------------------------------------------------===//
// Expr printing methods.
//===----------------------------------------------------------------------===//
Index: lib/AST/StmtOpenMP.cpp
===================================================================
--- lib/AST/StmtOpenMP.cpp
+++ lib/AST/StmtOpenMP.cpp
@@ -995,3 +995,25 @@
numLoopChildren(CollapsedNum, OMPD_distribute));
return new (Mem) OMPDistributeDirective(CollapsedNum, NumClauses);
}
+
+OMPTargetUpdateDirective *
+OMPTargetUpdateDirective::Create(const ASTContext &C, SourceLocation StartLoc,
+ SourceLocation EndLoc,
+ ArrayRef<OMPClause *> Clauses) {
+ unsigned Size = llvm::alignTo(sizeof(OMPTargetUpdateDirective),
+ llvm::alignOf<OMPClause *>());
+ void *Mem = C.Allocate(Size + sizeof(OMPClause *) * Clauses.size());
+ OMPTargetUpdateDirective *Dir =
+ new (Mem) OMPTargetUpdateDirective(StartLoc, EndLoc, Clauses.size());
+ Dir->setClauses(Clauses);
+ return Dir;
+}
+
+OMPTargetUpdateDirective *
+OMPTargetUpdateDirective::CreateEmpty(const ASTContext &C, unsigned NumClauses,
+ EmptyShell) {
+ unsigned Size = llvm::alignTo(sizeof(OMPTargetUpdateDirective),
+ llvm::alignOf<OMPClause *>());
+ void *Mem = C.Allocate(Size + sizeof(OMPClause *) * NumClauses);
+ return new (Mem) OMPTargetUpdateDirective(NumClauses);
+}
Index: lib/AST/OpenMPClause.cpp
===================================================================
--- lib/AST/OpenMPClause.cpp
+++ lib/AST/OpenMPClause.cpp
@@ -85,6 +85,8 @@
case OMPC_num_tasks:
case OMPC_hint:
case OMPC_defaultmap:
+ case OMPC_to:
+ case OMPC_from:
case OMPC_unknown:
break;
}
@@ -145,6 +147,8 @@
case OMPC_num_tasks:
case OMPC_hint:
case OMPC_defaultmap:
+ case OMPC_to:
+ case OMPC_from:
case OMPC_unknown:
break;
}
@@ -550,3 +554,34 @@
void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(N));
return new (Mem) OMPMapClause(N);
}
+
+OMPToClause *OMPToClause::Create(const ASTContext &C, SourceLocation StartLoc,
+ SourceLocation LParenLoc, SourceLocation EndLoc,
+ ArrayRef<Expr *> VL) {
+ void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size()));
+ OMPToClause *Clause = new (Mem) OMPToClause(StartLoc, LParenLoc, EndLoc,
+ VL.size());
+ Clause->setVarRefs(VL);
+ return Clause;
+}
+
+OMPToClause *OMPToClause::CreateEmpty(const ASTContext &C, unsigned N) {
+ void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(N));
+ return new (Mem) OMPToClause(N);
+}
+
+OMPFromClause *
+OMPFromClause::Create(const ASTContext &C, SourceLocation StartLoc,
+ SourceLocation LParenLoc, SourceLocation EndLoc,
+ ArrayRef<Expr *> VL) {
+ void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size()));
+ OMPFromClause *Clause = new (Mem) OMPFromClause(StartLoc, LParenLoc, EndLoc,
+ VL.size());
+ Clause->setVarRefs(VL);
+ return Clause;
+}
+
+OMPFromClause *OMPFromClause::CreateEmpty(const ASTContext &C, unsigned N) {
+ void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(N));
+ return new (Mem) OMPFromClause(N);
+}
Index: include/clang/Serialization/ASTBitCodes.h
===================================================================
--- include/clang/Serialization/ASTBitCodes.h
+++ include/clang/Serialization/ASTBitCodes.h
@@ -1479,6 +1479,7 @@
STMT_OMP_TASKLOOP_DIRECTIVE,
STMT_OMP_TASKLOOP_SIMD_DIRECTIVE,
STMT_OMP_DISTRIBUTE_DIRECTIVE,
+ STMT_OMP_TARGET_UPDATE_DIRECTIVE,
EXPR_OMP_ARRAY_SECTION,
// ARC
Index: include/clang/Sema/Sema.h
===================================================================
--- include/clang/Sema/Sema.h
+++ include/clang/Sema/Sema.h
@@ -8074,6 +8074,10 @@
ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
SourceLocation EndLoc,
llvm::DenseMap<ValueDecl *, Expr *> &VarsWithImplicitDSA);
+ /// \brief Called on well-formed '\#pragma omp target update'.
+ StmtResult ActOnOpenMPTargetUpdateDirective(ArrayRef<OMPClause *> Clauses,
+ SourceLocation StartLoc,
+ SourceLocation EndLoc);
OMPClause *ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind,
Expr *Expr,
@@ -8299,6 +8303,16 @@
OpenMPDefaultmapClauseModifier M, OpenMPDefaultmapClauseKind Kind,
SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation MLoc,
SourceLocation KindLoc, SourceLocation EndLoc);
+ /// \brief Called on well-formed 'to' clause.
+ OMPClause *ActOnOpenMPToClause(ArrayRef<Expr *> VarList,
+ SourceLocation StartLoc,
+ SourceLocation LParenLoc,
+ SourceLocation EndLoc);
+ /// \brief Called on well-formed 'from' clause.
+ OMPClause *ActOnOpenMPFromClause(ArrayRef<Expr *> VarList,
+ SourceLocation StartLoc,
+ SourceLocation LParenLoc,
+ SourceLocation EndLoc);
/// \brief The kind of conversion being performed.
enum CheckedConversionKind {
Index: include/clang/Basic/StmtNodes.td
===================================================================
--- include/clang/Basic/StmtNodes.td
+++ include/clang/Basic/StmtNodes.td
@@ -220,6 +220,7 @@
def OMPTargetExitDataDirective : DStmt<OMPExecutableDirective>;
def OMPTargetParallelDirective : DStmt<OMPExecutableDirective>;
def OMPTargetParallelForDirective : DStmt<OMPExecutableDirective>;
+def OMPTargetUpdateDirective : DStmt<OMPExecutableDirective>;
def OMPTeamsDirective : DStmt<OMPExecutableDirective>;
def OMPCancellationPointDirective : DStmt<OMPExecutableDirective>;
def OMPCancelDirective : DStmt<OMPExecutableDirective>;
Index: include/clang/Basic/OpenMPKinds.def
===================================================================
--- include/clang/Basic/OpenMPKinds.def
+++ include/clang/Basic/OpenMPKinds.def
@@ -72,6 +72,9 @@
#ifndef OPENMP_TARGET_PARALLEL_FOR_CLAUSE
# define OPENMP_TARGET_PARALLEL_FOR_CLAUSE(Name)
#endif
+#ifndef OPENMP_TARGET_UPDATE_CLAUSE
+# define OPENMP_TARGET_UPDATE_CLAUSE(Name)
+#endif
#ifndef OPENMP_TEAMS_CLAUSE
# define OPENMP_TEAMS_CLAUSE(Name)
#endif
@@ -150,6 +153,7 @@
OPENMP_DIRECTIVE_EXT(target_exit_data, "target exit data")
OPENMP_DIRECTIVE_EXT(target_parallel, "target parallel")
OPENMP_DIRECTIVE_EXT(target_parallel_for, "target parallel for")
+OPENMP_DIRECTIVE_EXT(target_update, "target update")
OPENMP_DIRECTIVE_EXT(parallel_for, "parallel for")
OPENMP_DIRECTIVE_EXT(parallel_for_simd, "parallel for simd")
OPENMP_DIRECTIVE_EXT(parallel_sections, "parallel sections")
@@ -203,6 +207,8 @@
OPENMP_CLAUSE(hint, OMPHintClause)
OPENMP_CLAUSE(dist_schedule, OMPDistScheduleClause)
OPENMP_CLAUSE(defaultmap, OMPDefaultmapClause)
+OPENMP_CLAUSE(to, OMPToClause)
+OPENMP_CLAUSE(from, OMPFromClause)
// Clauses allowed for OpenMP directive 'parallel'.
OPENMP_PARALLEL_CLAUSE(if)
@@ -439,6 +445,13 @@
OPENMP_TARGET_PARALLEL_FOR_CLAUSE(ordered)
OPENMP_TARGET_PARALLEL_FOR_CLAUSE(linear)
+// Clauses allowed for OpenMP directive 'target update'.
+// TODO More clauses for 'target update' directive.
+OPENMP_TARGET_UPDATE_CLAUSE(if)
+OPENMP_TARGET_UPDATE_CLAUSE(device)
+OPENMP_TARGET_UPDATE_CLAUSE(to)
+OPENMP_TARGET_UPDATE_CLAUSE(from)
+
// Clauses allowed for OpenMP directive 'teams'.
// TODO More clauses for 'teams' directive.
OPENMP_TEAMS_CLAUSE(default)
@@ -550,3 +563,4 @@
#undef OPENMP_DIST_SCHEDULE_KIND
#undef OPENMP_DEFAULTMAP_KIND
#undef OPENMP_DEFAULTMAP_MODIFIER
+#undef OPENMP_TARGET_UPDATE_CLAUSE
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -7823,8 +7823,8 @@
"expected variable name%select{|, data member of current class}0, array element or array section">;
def err_omp_expected_named_var_member_or_array_expression: Error<
"expected expression containing only member accesses and/or array sections based on named variables">;
-def err_omp_bit_fields_forbidden_in_map_clause : Error<
- "bit fields cannot be used to specify storage in a map clause">;
+def err_omp_bit_fields_forbidden_in_clause : Error<
+ "bit fields cannot be used to specify storage in a '%0' clause">;
def err_array_section_does_not_specify_contiguous_storage : Error<
"array section does not specify contiguous storage">;
def err_omp_union_type_not_allowed : Error<
@@ -7949,6 +7949,8 @@
"arguments of OpenMP clause 'reduction' with bitwise operators cannot be of floating type">;
def err_omp_once_referenced : Error<
"variable can appear only once in OpenMP '%0' clause">;
+def err_omp_once_referenced_in_target_update : Error<
+ "variable can appear only once in OpenMP 'target update' construct">;
def note_omp_referenced : Note<
"previously referenced here">;
def err_omp_reduction_in_task : Error<
@@ -8082,8 +8084,8 @@
"mappable type cannot be polymorphic">;
def note_omp_static_member_in_target : Note<
"mappable type cannot contain static members">;
-def err_omp_threadprivate_in_map : Error<
- "threadprivate variables are not allowed in map clause">;
+def err_omp_threadprivate_in_clause : Error<
+ "threadprivate variables are not allowed in '%0' clause">;
def err_omp_wrong_ordered_loop_count : Error<
"the parameter of the 'ordered' clause must be greater than or equal to the parameter of the 'collapse' clause">;
def note_collapse_loop_count : Note<
@@ -8128,6 +8130,8 @@
"'schedule' clause with 'nonmonotonic' modifier cannot be specified if an 'ordered' clause is specified">;
def err_omp_ordered_simd : Error<
"'ordered' clause with a parameter can not be specified in '#pragma omp %0' directive">;
+def err_omp_at_least_one_motion_clause_required : Error<
+ "expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'">;
} // end of OpenMP category
let CategoryName = "Related Result Type Issue" in {
Index: include/clang/AST/StmtOpenMP.h
===================================================================
--- include/clang/AST/StmtOpenMP.h
+++ include/clang/AST/StmtOpenMP.h
@@ -2660,7 +2660,7 @@
/// \param Clauses List of clauses.
/// \param AssociatedStmt Statement, associated with the directive.
/// \param Exprs Helper expressions for CodeGen.
- ///
+ ///
static OMPDistributeDirective *
Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses,
@@ -2682,6 +2682,64 @@
}
};
+/// \brief This represents '#pragma omp target update' directive.
+///
+/// \code
+/// #pragma omp target update to(a) from(b) device(1)
+/// \endcode
+/// In this example directive '#pragma omp target update' has clause 'to' with
+/// argument 'a', clause 'from' with argument 'b' and clause 'device' with
+/// argument '1'.
+///
+class OMPTargetUpdateDirective : public OMPExecutableDirective {
+ friend class ASTStmtReader;
+ /// \brief Build directive with the given start and end location.
+ ///
+ /// \param StartLoc Starting location of the directive kind.
+ /// \param EndLoc Ending Location of the directive.
+ /// \param NumClauses The number of clauses.
+ ///
+ OMPTargetUpdateDirective(SourceLocation StartLoc, SourceLocation EndLoc,
+ unsigned NumClauses)
+ : OMPExecutableDirective(this, OMPTargetUpdateDirectiveClass,
+ OMPD_target_update, StartLoc, EndLoc, NumClauses,
+ 0) {}
+
+ /// \brief Build an empty directive.
+ ///
+ /// \param NumClauses Number of clauses.
+ ///
+ explicit OMPTargetUpdateDirective(unsigned NumClauses)
+ : OMPExecutableDirective(this, OMPTargetUpdateDirectiveClass,
+ OMPD_target_update, SourceLocation(),
+ SourceLocation(), NumClauses, 0) {}
+public:
+ /// \brief Creates directive with a list of \a Clauses.
+ ///
+ /// \param C AST context.
+ /// \param StartLoc Starting location of the directive kind.
+ /// \param EndLoc Ending Location of the directive.
+ /// \param Clauses List of clauses.
+ ///
+ static OMPTargetUpdateDirective *Create(const ASTContext &C,
+ SourceLocation StartLoc,
+ SourceLocation EndLoc,
+ ArrayRef<OMPClause *> Clauses);
+
+ /// \brief Creates an empty directive with the place for \a NumClauses clauses.
+ ///
+ /// \param C AST context.
+ /// \param NumClauses The number of clauses.
+ ///
+ static OMPTargetUpdateDirective *CreateEmpty(const ASTContext &C,
+ unsigned NumClauses,
+ EmptyShell);
+
+ static bool classof(const Stmt *T) {
+ return T->getStmtClass() == OMPTargetUpdateDirectiveClass;
+ }
+};
+
} // end namespace clang
#endif
Index: include/clang/AST/RecursiveASTVisitor.h
===================================================================
--- include/clang/AST/RecursiveASTVisitor.h
+++ include/clang/AST/RecursiveASTVisitor.h
@@ -2496,6 +2496,9 @@
DEF_TRAVERSE_STMT(OMPTeamsDirective,
{ TRY_TO(TraverseOMPExecutableDirective(S)); })
+DEF_TRAVERSE_STMT(OMPTargetUpdateDirective,
+ { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
DEF_TRAVERSE_STMT(OMPTaskLoopDirective,
{ TRY_TO(TraverseOMPExecutableDirective(S)); })
@@ -2874,6 +2877,18 @@
return true;
}
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPFromClause(OMPFromClause *C) {
+ TRY_TO(VisitOMPClauseList(C));
+ return true;
+}
+
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPToClause(OMPToClause *C) {
+ TRY_TO(VisitOMPClauseList(C));
+ return true;
+}
+
// FIXME: look at the following tricky-seeming exprs to see if we
// need to recurse on anything. These are ones that have methods
// returning decls or qualtypes or nestednamespecifier -- though I'm
Index: include/clang/AST/OpenMPClause.h
===================================================================
--- include/clang/AST/OpenMPClause.h
+++ include/clang/AST/OpenMPClause.h
@@ -3471,6 +3471,128 @@
return child_range(child_iterator(), child_iterator());
}
};
+
+/// \brief This represents clause 'from' in the '#pragma omp ...'
+/// directives.
+///
+/// \code
+/// #pragma omp target update from(a,b)
+/// \endcode
+/// In this example directive '#pragma omp target update' has clause 'from'
+/// with the variables 'a' and 'b'.
+///
+class OMPFromClause final : public OMPVarListClause<OMPFromClause>,
+ private llvm::TrailingObjects<OMPFromClause, Expr *> {
+ friend TrailingObjects;
+ friend OMPVarListClause;
+ friend class OMPClauseReader;
+
+ /// \brief Build clause with number of variables \a N.
+ ///
+ /// \param StartLoc Starting location of the clause.
+ /// \param EndLoc Ending location of the clause.
+ /// \param N Number of the variables in the clause.
+ ///
+ explicit OMPFromClause(SourceLocation StartLoc, SourceLocation LParenLoc,
+ SourceLocation EndLoc, unsigned N)
+ : OMPVarListClause<OMPFromClause>(OMPC_from, StartLoc, LParenLoc, EndLoc, N)
+ {}
+ /// \brief Build an empty clause.
+ ///
+ /// \param N Number of variables.
+ ///
+ explicit OMPFromClause(unsigned N)
+ : OMPVarListClause<OMPFromClause>(OMPC_from, SourceLocation(),
+ SourceLocation(), SourceLocation(), N) {}
+
+public:
+ /// \brief Creates clause with a list of variables \a VL.
+ ///
+ /// \param C AST context.
+ /// \brief StartLoc Starting location of the clause.
+ /// \brief EndLoc Ending location of the clause.
+ /// \param VL List of references to the variables.
+ ///
+ static OMPFromClause *Create(const ASTContext &C, SourceLocation StartLoc,
+ SourceLocation LParenLoc, SourceLocation EndLoc,
+ ArrayRef<Expr *> VL);
+ /// \brief Creates an empty clause with the place for \a N variables.
+ ///
+ /// \param C AST context.
+ /// \param N The number of variables.
+ ///
+ static OMPFromClause *CreateEmpty(const ASTContext &C, unsigned N);
+
+ static bool classof(const OMPClause *T) {
+ return T->getClauseKind() == OMPC_from;
+ }
+
+ child_range children() {
+ return child_range(reinterpret_cast<Stmt **>(varlist_begin()),
+ reinterpret_cast<Stmt **>(varlist_end()));
+ }
+};
+
+/// \brief This represents clause 'to' in the '#pragma omp ...'
+/// directives.
+///
+/// \code
+/// #pragma omp target update to(a,b)
+/// \endcode
+/// In this example directive '#pragma omp target update' has clause 'to'
+/// with the variables 'a' and 'b'.
+///
+class OMPToClause final : public OMPVarListClause<OMPToClause>,
+ private llvm::TrailingObjects<OMPToClause, Expr *> {
+ friend TrailingObjects;
+ friend OMPVarListClause;
+ friend class OMPClauseReader;
+
+ /// \brief Build clause with number of variables \a N.
+ ///
+ /// \param StartLoc Starting location of the clause.
+ /// \param EndLoc Ending location of the clause.
+ /// \param N Number of the variables in the clause.
+ ///
+ explicit OMPToClause(SourceLocation StartLoc, SourceLocation LParenLoc,
+ SourceLocation EndLoc, unsigned N)
+ : OMPVarListClause<OMPToClause>(OMPC_to, StartLoc, LParenLoc, EndLoc, N) {}
+ /// \brief Build an empty clause.
+ ///
+ /// \param N Number of variables.
+ ///
+ explicit OMPToClause(unsigned N)
+ : OMPVarListClause<OMPToClause>(OMPC_to, SourceLocation(), SourceLocation(),
+ SourceLocation(), N) {}
+
+public:
+ /// \brief Creates clause with a list of variables \a VL.
+ ///
+ /// \param C AST context.
+ /// \brief StartLoc Starting location of the clause.
+ /// \brief EndLoc Ending location of the clause.
+ /// \param VL List of references to the variables.
+ ///
+ static OMPToClause *Create(const ASTContext &C, SourceLocation StartLoc,
+ SourceLocation LParenLoc, SourceLocation EndLoc,
+ ArrayRef<Expr *> VL);
+ /// \brief Creates an empty clause with the place for \a N variables.
+ ///
+ /// \param C AST context.
+ /// \param N The number of variables.
+ ///
+ static OMPToClause *CreateEmpty(const ASTContext &C, unsigned N);
+
+ static bool classof(const OMPClause *T) {
+ return T->getClauseKind() == OMPC_to;
+ }
+
+ child_range children() {
+ return child_range(reinterpret_cast<Stmt **>(varlist_begin()),
+ reinterpret_cast<Stmt **>(varlist_end()));
+ }
+};
+
} // end namespace clang
#endif // LLVM_CLANG_AST_OPENMPCLAUSE_H
Index: include/clang-c/Index.h
===================================================================
--- include/clang-c/Index.h
+++ include/clang-c/Index.h
@@ -2281,7 +2281,7 @@
*/
CXCursor_OMPTaskLoopSimdDirective = 259,
- /** \brief OpenMP distribute directive.
+ /** \brief OpenMP distribute directive.
*/
CXCursor_OMPDistributeDirective = 260,
@@ -2301,7 +2301,11 @@
*/
CXCursor_OMPTargetParallelForDirective = 264,
- CXCursor_LastStmt = CXCursor_OMPTargetParallelForDirective,
+ /** \brief OpenMP target update directive.
+ */
+ CXCursor_OMPTargetUpdateDirective = 265,
+
+ CXCursor_LastStmt = CXCursor_OMPTargetUpdateDirective,
/**
* \brief Cursor that represents the translation unit itself.
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
