arpith-jacob created this revision.
arpith-jacob added reviewers: kkwli0, sfantao, fraggamuffin, hfinkel, ABataev,
carlo.bertolli.
arpith-jacob added a subscriber: cfe-commits.
This patch adds parsing and sema support for "#pragma omp target enter data"
and "#pragma omp target exit data" directives.
http://reviews.llvm.org/D15989
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/target_enter_data_ast_print.cpp
test/OpenMP/target_enter_data_device_messages.cpp
test/OpenMP/target_enter_data_if_messages.cpp
test/OpenMP/target_enter_data_map_messages.c
test/OpenMP/target_exit_data_ast_print.cpp
test/OpenMP/target_exit_data_device_messages.cpp
test/OpenMP/target_exit_data_if_messages.cpp
test/OpenMP/target_exit_data_map_messages.c
tools/libclang/CIndex.cpp
tools/libclang/CXCursor.cpp
Index: tools/libclang/CXCursor.cpp
===================================================================
--- tools/libclang/CXCursor.cpp
+++ tools/libclang/CXCursor.cpp
@@ -600,6 +600,12 @@
case Stmt::OMPTargetDataDirectiveClass:
K = CXCursor_OMPTargetDataDirective;
break;
+ case Stmt::OMPTargetEnterDataDirectiveClass:
+ K = CXCursor_OMPTargetEnterDataDirective;
+ break;
+ case Stmt::OMPTargetExitDataDirectiveClass:
+ K = CXCursor_OMPTargetExitDataDirective;
+ break;
case Stmt::OMPTeamsDirectiveClass:
K = CXCursor_OMPTeamsDirective;
break;
Index: tools/libclang/CIndex.cpp
===================================================================
--- tools/libclang/CIndex.cpp
+++ tools/libclang/CIndex.cpp
@@ -1949,6 +1949,8 @@
void VisitOMPAtomicDirective(const OMPAtomicDirective *D);
void VisitOMPTargetDirective(const OMPTargetDirective *D);
void VisitOMPTargetDataDirective(const OMPTargetDataDirective *D);
+ void VisitOMPTargetEnterDataDirective(const OMPTargetEnterDataDirective *D);
+ void VisitOMPTargetExitDataDirective(const OMPTargetExitDataDirective *D);
void VisitOMPTeamsDirective(const OMPTeamsDirective *D);
void VisitOMPTaskLoopDirective(const OMPTaskLoopDirective *D);
void VisitOMPTaskLoopSimdDirective(const OMPTaskLoopSimdDirective *D);
@@ -2624,6 +2626,16 @@
VisitOMPExecutableDirective(D);
}
+void EnqueueVisitor::VisitOMPTargetEnterDataDirective(
+ const OMPTargetEnterDataDirective *D) {
+ VisitOMPExecutableDirective(D);
+}
+
+void EnqueueVisitor::VisitOMPTargetExitDataDirective(
+ const OMPTargetExitDataDirective *D) {
+ VisitOMPExecutableDirective(D);
+}
+
void EnqueueVisitor::VisitOMPTeamsDirective(const OMPTeamsDirective *D) {
VisitOMPExecutableDirective(D);
}
@@ -4506,6 +4518,10 @@
return cxstring::createRef("OMPTargetDirective");
case CXCursor_OMPTargetDataDirective:
return cxstring::createRef("OMPTargetDataDirective");
+ case CXCursor_OMPTargetEnterDataDirective:
+ return cxstring::createRef("OMPTargetEnterDataDirective");
+ case CXCursor_OMPTargetExitDataDirective:
+ return cxstring::createRef("OMPTargetExitDataDirective");
case CXCursor_OMPTeamsDirective:
return cxstring::createRef("OMPTeamsDirective");
case CXCursor_OMPCancellationPointDirective:
Index: test/OpenMP/target_exit_data_map_messages.c
===================================================================
--- /dev/null
+++ test/OpenMP/target_exit_data_map_messages.c
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s
+
+int main(int argc, char **argv) {
+
+ int r;
+ #pragma omp target exit data map(r) // expected-error {{map type must be specified for '#pragma omp target exit data'}}
+
+ #pragma omp target exit data // expected-error {{expected at least one map clause for '#pragma omp target exit data'}}
+
+ #pragma omp target exit data map(tofrom: r) // expected-error {{map type 'tofrom' is not allowed for '#pragma omp target exit data'}}
+
+ #pragma omp target exit data map(always, from: r)
+ #pragma omp target exit data map(delete: r)
+ #pragma omp target exit data map(release: r)
+ #pragma omp target exit data map(always, alloc: r) // expected-error {{map type 'alloc' is not allowed for '#pragma omp target exit data'}}
+ #pragma omp target exit data map(to: r) // expected-error {{map type 'to' is not allowed for '#pragma omp target exit data'}}
+
+ return 0;
+}
Index: test/OpenMP/target_exit_data_if_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_exit_data_if_messages.cpp
@@ -0,0 +1,35 @@
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+ return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+int main(int argc, char **argv) {
+ int i;
+ #pragma omp target exit data map(from: i) if // expected-error {{expected '(' after 'if'}}
+ #pragma omp target exit data map(from: i) if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target exit data map(from: i) if () // expected-error {{expected expression}}
+ #pragma omp target exit data map(from: i) if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target exit data map(from: i) if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target exit data' are ignored}}
+ #pragma omp target exit data map(from: i) if (argc > 0 ? argv[1] : argv[2])
+ #pragma omp target exit data map(from: i) if (argc + argc)
+ #pragma omp target exit data map(from: i) if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target exit data' cannot contain more than one 'if' clause}}
+ #pragma omp target exit data map(from: i) if (S1) // expected-error {{'S1' does not refer to a value}}
+ #pragma omp target exit data map(from: i) if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target exit data map(from: i) if(target data : true) // expected-error {{directive name modifier 'target data' is not allowed for '#pragma omp target exit data'}}
+ #pragma omp target exit data map(from: i) if(target exit data : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target exit data map(from: i) if(target exit data : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target exit data map(from: i) if(target exit data : argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target exit data map(from: i) if(target exit data : argc)
+ #pragma omp target exit data map(from: i) if(target exit data : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target exit data'}}
+ #pragma omp target exit data map(from: i) if(target exit data : argc) if (target exit data:argc) // expected-error {{directive '#pragma omp target exit data' cannot contain more than one 'if' clause with 'target exit data' name modifier}}
+ #pragma omp target exit data map(from: i) if(target exit data : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}}
+ foo();
+
+ return 0;
+}
Index: test/OpenMP/target_exit_data_device_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_exit_data_device_messages.cpp
@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+ return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+int main(int argc, char **argv) {
+ int i;
+ #pragma omp target exit data map(from: i) device // expected-error {{expected '(' after 'device'}}
+ #pragma omp target exit data map(from: i) device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target exit data map(from: i) device () // expected-error {{expected expression}}
+ #pragma omp target exit data map(from: i) device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target exit data map(from: i) device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target exit data' are ignored}}
+#pragma omp target exit data map(from: i) device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+ #pragma omp target exit data map(from: i) device (argc + argc)
+ #pragma omp target exit data map(from: i) device (argc), device (argc+1) // expected-error {{directive '#pragma omp target exit data' cannot contain more than one 'device' clause}}
+ #pragma omp target exit data map(from: i) device (S1) // expected-error {{'S1' does not refer to a value}}
+ #pragma omp target exit data map(from: i) device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}}
+ #pragma omp target exit data map(from: i) device (-10u)
+ #pragma omp target exit data map(from: i) device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}}
+ foo();
+
+ return 0;
+}
Index: test/OpenMP/target_exit_data_ast_print.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_exit_data_ast_print.cpp
@@ -0,0 +1,100 @@
+// 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
+
+template <typename T, int C>
+T tmain(T argc, T *argv) {
+ T i, j, b, c, d, e, x[20];
+
+ i = argc;
+#pragma omp target exit data map(from: i)
+
+#pragma omp target exit data map(from: i) if (target exit data: j > 0)
+
+#pragma omp target exit data map(from: i) if (b)
+
+#pragma omp target exit data map(from: c)
+
+#pragma omp target exit data map(from: c) if(b>e)
+
+#pragma omp target exit data map(release: x[0:10], c)
+
+#pragma omp target exit data map(from: c) map(release: d)
+
+#pragma omp target exit data map(always,release: e)
+
+ return 0;
+}
+
+// CHECK: template <typename T = int, int C = 5> int tmain(int argc, int *argv) {
+// CHECK-NEXT: int i, j, b, c, d, e, x[20];
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target exit data map(from: i)
+// CHECK-NEXT: #pragma omp target exit data map(from: i) if(target exit data: j > 0)
+// CHECK-NEXT: #pragma omp target exit data map(from: i) if(b)
+// CHECK-NEXT: #pragma omp target exit data map(from: c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) if(b > e)
+// CHECK-NEXT: #pragma omp target exit data map(release: x[0:10],c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) map(release: d)
+// CHECK-NEXT: #pragma omp target exit data map(always,release: e)
+// CHECK: template <typename T = char, int C = 1> char tmain(char argc, char *argv) {
+// CHECK-NEXT: char i, j, b, c, d, e, x[20];
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target exit data map(from: i)
+// CHECK-NEXT: #pragma omp target exit data map(from: i) if(target exit data: j > 0)
+// CHECK-NEXT: #pragma omp target exit data map(from: i) if(b)
+// CHECK-NEXT: #pragma omp target exit data map(from: c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) if(b > e)
+// CHECK-NEXT: #pragma omp target exit data map(release: x[0:10],c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) map(release: d)
+// CHECK-NEXT: #pragma omp target exit data map(always,release: e)
+// CHECK: template <typename T, int C> T tmain(T argc, T *argv) {
+// CHECK-NEXT: T i, j, b, c, d, e, x[20];
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target exit data map(from: i)
+// CHECK-NEXT: #pragma omp target exit data map(from: i) if(target exit data: j > 0)
+// CHECK-NEXT: #pragma omp target exit data map(from: i) if(b)
+// CHECK-NEXT: #pragma omp target exit data map(from: c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) if(b > e)
+// CHECK-NEXT: #pragma omp target exit data map(release: x[0:10],c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) map(release: d)
+// CHECK-NEXT: #pragma omp target exit data map(always,release: e)
+
+int main (int argc, char **argv) {
+ int b = argc, c, d, e, f, g, x[20];
+ static int a;
+// CHECK: static int a;
+
+#pragma omp target exit data map(from: a)
+// CHECK: #pragma omp target exit data map(from: a)
+ a=2;
+// CHECK-NEXT: a = 2;
+#pragma omp target exit data map(from: a) if (target exit data: b)
+// CHECK: #pragma omp target exit data map(from: a) if(target exit data: b)
+
+#pragma omp target exit data map(from: a) if (b > g)
+// CHECK: #pragma omp target exit data map(from: a) if(b > g)
+
+#pragma omp target exit data map(from: c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c)
+
+#pragma omp target exit data map(release: c) if(b>g)
+// CHECK-NEXT: #pragma omp target exit data map(release: c) if(b > g)
+
+#pragma omp target exit data map(from: x[0:10], c)
+// CHECK-NEXT: #pragma omp target exit data map(from: x[0:10],c)
+
+#pragma omp target exit data map(from: c) map(release: d)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) map(release: d)
+
+#pragma omp target exit data map(always,release: e)
+// CHECK-NEXT: #pragma omp target exit data map(always,release: e)
+
+ return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);
+}
+
+#endif
Index: test/OpenMP/target_enter_data_map_messages.c
===================================================================
--- /dev/null
+++ test/OpenMP/target_enter_data_map_messages.c
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s
+
+int main(int argc, char **argv) {
+
+ int r;
+ #pragma omp target enter data map(r) // expected-error {{map type must be specified for '#pragma omp target enter data'}}
+
+ #pragma omp target enter data // expected-error {{expected at least one map clause for '#pragma omp target enter data'}}
+
+ #pragma omp target enter data map(tofrom: r) // expected-error {{map type 'tofrom' is not allowed for '#pragma omp target enter data'}}
+
+ #pragma omp target enter data map(always, to: r)
+ #pragma omp target enter data map(always, alloc: r)
+ #pragma omp target enter data map(always, from: r) // expected-error {{map type 'from' is not allowed for '#pragma omp target enter data'}}
+ #pragma omp target enter data map(release: r) // expected-error {{map type 'release' is not allowed for '#pragma omp target enter data'}}
+ #pragma omp target enter data map(delete: r) // expected-error {{map type 'delete' is not allowed for '#pragma omp target enter data'}}
+
+ return 0;
+}
Index: test/OpenMP/target_enter_data_if_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_enter_data_if_messages.cpp
@@ -0,0 +1,35 @@
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+ return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+int main(int argc, char **argv) {
+ int i;
+ #pragma omp target enter data map(to: i) if // expected-error {{expected '(' after 'if'}}
+ #pragma omp target enter data map(to: i) if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target enter data map(to: i) if () // expected-error {{expected expression}}
+ #pragma omp target enter data map(to: i) if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target enter data map(to: i) if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}}
+ #pragma omp target enter data map(to: i) if (argc > 0 ? argv[1] : argv[2])
+ #pragma omp target enter data map(to: i) if (argc + argc)
+ #pragma omp target enter data map(to: i) if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target enter data' cannot contain more than one 'if' clause}}
+ #pragma omp target enter data map(to: i) if (S1) // expected-error {{'S1' does not refer to a value}}
+ #pragma omp target enter data map(to: i) if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target enter data map(to: i) if(target data : true) // expected-error {{directive name modifier 'target data' is not allowed for '#pragma omp target enter data'}}
+ #pragma omp target enter data map(to: i) if(target enter data : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target enter data map(to: i) if(target enter data : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target enter data map(to: i) if(target enter data : argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target enter data map(to: i) if(target enter data : argc)
+ #pragma omp target enter data map(to: i) if(target enter data : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target enter data'}}
+ #pragma omp target enter data map(to: i) if(target enter data : argc) if (target enter data:argc) // expected-error {{directive '#pragma omp target enter data' cannot contain more than one 'if' clause with 'target enter data' name modifier}}
+ #pragma omp target enter data map(to: i) if(target enter data : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}}
+ foo();
+
+ return 0;
+}
Index: test/OpenMP/target_enter_data_device_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_enter_data_device_messages.cpp
@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+ return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+int main(int argc, char **argv) {
+ int i;
+ #pragma omp target enter data map(to: i) device // expected-error {{expected '(' after 'device'}}
+ #pragma omp target enter data map(to: i) device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target enter data map(to: i) device () // expected-error {{expected expression}}
+ #pragma omp target enter data map(to: i) device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target enter data map(to: i) device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}}
+#pragma omp target enter data map(to: i) device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+ #pragma omp target enter data map(to: i) device (argc + argc)
+ #pragma omp target enter data map(to: i) device (argc), device (argc+1) // expected-error {{directive '#pragma omp target enter data' cannot contain more than one 'device' clause}}
+ #pragma omp target enter data map(to: i) device (S1) // expected-error {{'S1' does not refer to a value}}
+ #pragma omp target enter data map(to: i) device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}}
+ #pragma omp target enter data map(to: i) device (-10u)
+ #pragma omp target enter data map(to: i) device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}}
+ foo();
+
+ return 0;
+}
Index: test/OpenMP/target_enter_data_ast_print.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_enter_data_ast_print.cpp
@@ -0,0 +1,100 @@
+// 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
+
+template <typename T, int C>
+T tmain(T argc, T *argv) {
+ T i, j, b, c, d, e, x[20];
+
+ i = argc;
+#pragma omp target enter data map(to: i)
+
+#pragma omp target enter data map(to: i) if (target enter data: j > 0)
+
+#pragma omp target enter data map(to: i) if (b)
+
+#pragma omp target enter data map(to: c)
+
+#pragma omp target enter data map(to: c) if(b>e)
+
+#pragma omp target enter data map(alloc: x[0:10], c)
+
+#pragma omp target enter data map(to: c) map(alloc: d)
+
+#pragma omp target enter data map(always,alloc: e)
+
+ return 0;
+}
+
+// CHECK: template <typename T = int, int C = 5> int tmain(int argc, int *argv) {
+// CHECK-NEXT: int i, j, b, c, d, e, x[20];
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target enter data map(to: i)
+// CHECK-NEXT: #pragma omp target enter data map(to: i) if(target enter data: j > 0)
+// CHECK-NEXT: #pragma omp target enter data map(to: i) if(b)
+// CHECK-NEXT: #pragma omp target enter data map(to: c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) if(b > e)
+// CHECK-NEXT: #pragma omp target enter data map(alloc: x[0:10],c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) map(alloc: d)
+// CHECK-NEXT: #pragma omp target enter data map(always,alloc: e)
+// CHECK: template <typename T = char, int C = 1> char tmain(char argc, char *argv) {
+// CHECK-NEXT: char i, j, b, c, d, e, x[20];
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target enter data map(to: i)
+// CHECK-NEXT: #pragma omp target enter data map(to: i) if(target enter data: j > 0)
+// CHECK-NEXT: #pragma omp target enter data map(to: i) if(b)
+// CHECK-NEXT: #pragma omp target enter data map(to: c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) if(b > e)
+// CHECK-NEXT: #pragma omp target enter data map(alloc: x[0:10],c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) map(alloc: d)
+// CHECK-NEXT: #pragma omp target enter data map(always,alloc: e)
+// CHECK: template <typename T, int C> T tmain(T argc, T *argv) {
+// CHECK-NEXT: T i, j, b, c, d, e, x[20];
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target enter data map(to: i)
+// CHECK-NEXT: #pragma omp target enter data map(to: i) if(target enter data: j > 0)
+// CHECK-NEXT: #pragma omp target enter data map(to: i) if(b)
+// CHECK-NEXT: #pragma omp target enter data map(to: c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) if(b > e)
+// CHECK-NEXT: #pragma omp target enter data map(alloc: x[0:10],c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) map(alloc: d)
+// CHECK-NEXT: #pragma omp target enter data map(always,alloc: e)
+
+int main (int argc, char **argv) {
+ int b = argc, c, d, e, f, g, x[20];
+ static int a;
+// CHECK: static int a;
+
+#pragma omp target enter data map(to: a)
+// CHECK: #pragma omp target enter data map(to: a)
+ a=2;
+// CHECK-NEXT: a = 2;
+#pragma omp target enter data map(to: a) if (target enter data: b)
+// CHECK: #pragma omp target enter data map(to: a) if(target enter data: b)
+
+#pragma omp target enter data map(to: a) if (b > g)
+// CHECK: #pragma omp target enter data map(to: a) if(b > g)
+
+#pragma omp target enter data map(to: c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c)
+
+#pragma omp target enter data map(alloc: c) if(b>g)
+// CHECK-NEXT: #pragma omp target enter data map(alloc: c) if(b > g)
+
+#pragma omp target enter data map(to: x[0:10], c)
+// CHECK-NEXT: #pragma omp target enter data map(to: x[0:10],c)
+
+#pragma omp target enter data map(to: c) map(alloc: d)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) map(alloc: d)
+
+#pragma omp target enter data map(always,alloc: e)
+// CHECK-NEXT: #pragma omp target enter data map(always,alloc: e)
+
+ return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);
+}
+
+#endif
Index: lib/StaticAnalyzer/Core/ExprEngine.cpp
===================================================================
--- lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -830,6 +830,8 @@
case Stmt::OMPAtomicDirectiveClass:
case Stmt::OMPTargetDirectiveClass:
case Stmt::OMPTargetDataDirectiveClass:
+ case Stmt::OMPTargetEnterDataDirectiveClass:
+ case Stmt::OMPTargetExitDataDirectiveClass:
case Stmt::OMPTeamsDirectiveClass:
case Stmt::OMPCancellationPointDirectiveClass:
case Stmt::OMPCancelDirectiveClass:
Index: lib/Serialization/ASTWriterStmt.cpp
===================================================================
--- lib/Serialization/ASTWriterStmt.cpp
+++ lib/Serialization/ASTWriterStmt.cpp
@@ -2212,6 +2212,22 @@
Code = serialization::STMT_OMP_TARGET_DATA_DIRECTIVE;
}
+void ASTStmtWriter::VisitOMPTargetEnterDataDirective(
+ OMPTargetEnterDataDirective *D) {
+ VisitStmt(D);
+ Record.push_back(D->getNumClauses());
+ VisitOMPExecutableDirective(D);
+ Code = serialization::STMT_OMP_TARGET_ENTER_DATA_DIRECTIVE;
+}
+
+void ASTStmtWriter::VisitOMPTargetExitDataDirective(
+ OMPTargetExitDataDirective *D) {
+ VisitStmt(D);
+ Record.push_back(D->getNumClauses());
+ VisitOMPExecutableDirective(D);
+ Code = serialization::STMT_OMP_TARGET_EXIT_DATA_DIRECTIVE;
+}
+
void ASTStmtWriter::VisitOMPTaskyieldDirective(OMPTaskyieldDirective *D) {
VisitStmt(D);
VisitOMPExecutableDirective(D);
Index: lib/Serialization/ASTReaderStmt.cpp
===================================================================
--- lib/Serialization/ASTReaderStmt.cpp
+++ lib/Serialization/ASTReaderStmt.cpp
@@ -2445,6 +2445,20 @@
VisitOMPExecutableDirective(D);
}
+void ASTStmtReader::VisitOMPTargetEnterDataDirective(
+ OMPTargetEnterDataDirective *D) {
+ VisitStmt(D);
+ ++Idx;
+ VisitOMPExecutableDirective(D);
+}
+
+void ASTStmtReader::VisitOMPTargetExitDataDirective(
+ OMPTargetExitDataDirective *D) {
+ VisitStmt(D);
+ ++Idx;
+ VisitOMPExecutableDirective(D);
+}
+
void ASTStmtReader::VisitOMPTeamsDirective(OMPTeamsDirective *D) {
VisitStmt(D);
// The NumClauses field was read in ReadStmtFromStream.
@@ -3088,6 +3102,16 @@
Context, Record[ASTStmtReader::NumStmtFields], Empty);
break;
+ case STMT_OMP_TARGET_ENTER_DATA_DIRECTIVE:
+ S = OMPTargetEnterDataDirective::CreateEmpty(
+ Context, Record[ASTStmtReader::NumStmtFields], Empty);
+ break;
+
+ case STMT_OMP_TARGET_EXIT_DATA_DIRECTIVE:
+ S = OMPTargetExitDataDirective::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
@@ -1657,10 +1657,11 @@
/// Subclasses may override this routine to provide different behavior.
OMPClause *RebuildOMPMapClause(
OpenMPMapClauseKind MapTypeModifier, OpenMPMapClauseKind MapType,
- SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
- SourceLocation StartLoc, SourceLocation LParenLoc,
- SourceLocation EndLoc) {
- return getSema().ActOnOpenMPMapClause(MapTypeModifier, MapType, MapLoc,
+ bool IsMapTypeImplicit, SourceLocation MapLoc, SourceLocation ColonLoc,
+ ArrayRef<Expr *> VarList, SourceLocation StartLoc,
+ SourceLocation LParenLoc, SourceLocation EndLoc) {
+ return getSema().ActOnOpenMPMapClause(MapTypeModifier, MapType,
+ IsMapTypeImplicit, MapLoc,
ColonLoc, VarList,StartLoc,
LParenLoc, EndLoc);
}
@@ -7353,6 +7354,28 @@
}
template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOMPTargetEnterDataDirective(
+ OMPTargetEnterDataDirective *D) {
+ DeclarationNameInfo DirName;
+ getDerived().getSema().StartOpenMPDSABlock(OMPD_target_enter_data, DirName,
+ nullptr, D->getLocStart());
+ StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
+ getDerived().getSema().EndOpenMPDSABlock(Res.get());
+ return Res;
+}
+
+template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOMPTargetExitDataDirective(
+ OMPTargetExitDataDirective *D) {
+ DeclarationNameInfo DirName;
+ getDerived().getSema().StartOpenMPDSABlock(OMPD_target_exit_data, 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;
@@ -7800,9 +7823,9 @@
Vars.push_back(EVar.get());
}
return getDerived().RebuildOMPMapClause(
- C->getMapTypeModifier(), C->getMapType(), C->getMapLoc(),
- C->getColonLoc(), Vars, C->getLocStart(), C->getLParenLoc(),
- C->getLocEnd());
+ C->getMapTypeModifier(), C->getMapType(), C->isImplicitMapType(),
+ C->getMapLoc(), C->getColonLoc(), Vars, C->getLocStart(),
+ C->getLParenLoc(), C->getLocEnd());
}
template <typename Derived>
Index: lib/Sema/SemaOpenMP.cpp
===================================================================
--- lib/Sema/SemaOpenMP.cpp
+++ lib/Sema/SemaOpenMP.cpp
@@ -1610,6 +1610,8 @@
case OMPD_cancellation_point:
case OMPD_cancel:
case OMPD_flush:
+ case OMPD_target_enter_data:
+ case OMPD_target_exit_data:
llvm_unreachable("OpenMP Directive is not allowed");
case OMPD_unknown:
llvm_unreachable("Unknown OpenMP directive");
@@ -2697,6 +2699,16 @@
EndLoc);
AllowedNameModifiers.push_back(OMPD_target_data);
break;
+ case OMPD_target_enter_data:
+ Res = ActOnOpenMPTargetEnterDataDirective(ClausesWithImplicit, StartLoc,
+ EndLoc);
+ AllowedNameModifiers.push_back(OMPD_target_enter_data);
+ break;
+ case OMPD_target_exit_data:
+ Res = ActOnOpenMPTargetExitDataDirective(ClausesWithImplicit, StartLoc,
+ EndLoc);
+ AllowedNameModifiers.push_back(OMPD_target_exit_data);
+ break;
case OMPD_taskloop:
Res = ActOnOpenMPTaskLoopDirective(ClausesWithImplicit, AStmt, StartLoc,
EndLoc, VarsWithInheritedDSA);
@@ -5448,6 +5460,34 @@
AStmt);
}
+StmtResult Sema::ActOnOpenMPTargetEnterDataDirective(
+ ArrayRef<OMPClause *> Clauses, SourceLocation StartLoc,
+ SourceLocation EndLoc) {
+ // OpenMP [2.10.2, Restrictions, p. 99]
+ // At least one map clause must appear on the directive.
+ if (!HasMapClause(Clauses)) {
+ Diag(StartLoc, diag::err_omp_no_map_for_directive) <<
+ getOpenMPDirectiveName(OMPD_target_enter_data);
+ return StmtError();
+ }
+
+ return OMPTargetEnterDataDirective::Create(Context, StartLoc, EndLoc, Clauses);
+}
+
+StmtResult Sema::ActOnOpenMPTargetExitDataDirective(
+ ArrayRef<OMPClause *> Clauses, SourceLocation StartLoc,
+ SourceLocation EndLoc) {
+ // OpenMP [2.10.3, Restrictions, p. 102]
+ // At least one map clause must appear on the directive.
+ if (!HasMapClause(Clauses)) {
+ Diag(StartLoc, diag::err_omp_no_map_for_directive) <<
+ getOpenMPDirectiveName(OMPD_target_exit_data);
+ return StmtError();
+ }
+
+ return OMPTargetExitDataDirective::Create(Context, StartLoc, EndLoc, Clauses);
+}
+
StmtResult Sema::ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
Stmt *AStmt, SourceLocation StartLoc,
SourceLocation EndLoc) {
@@ -5863,6 +5903,19 @@
return ICE;
}
+bool Sema::HasMapClause(ArrayRef<OMPClause *> Clauses) {
+ for (ArrayRef<OMPClause *>::iterator I = Clauses.begin(), E = Clauses.end();
+ I != E; ++I) {
+ if (*I) {
+ OMPClause *Clause = *I;
+ if (Clause->getClauseKind() == OMPC_map)
+ return true;
+ }
+ }
+
+ return false;
+}
+
OMPClause *Sema::ActOnOpenMPSafelenClause(Expr *Len, SourceLocation StartLoc,
SourceLocation LParenLoc,
SourceLocation EndLoc) {
@@ -6371,7 +6424,8 @@
SourceLocation EndLoc, CXXScopeSpec &ReductionIdScopeSpec,
const DeclarationNameInfo &ReductionId, OpenMPDependClauseKind DepKind,
OpenMPLinearClauseKind LinKind, OpenMPMapClauseKind MapTypeModifier,
- OpenMPMapClauseKind MapType, SourceLocation DepLinMapLoc) {
+ OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
+ SourceLocation DepLinMapLoc) {
OMPClause *Res = nullptr;
switch (Kind) {
case OMPC_private:
@@ -6412,8 +6466,9 @@
StartLoc, LParenLoc, EndLoc);
break;
case OMPC_map:
- Res = ActOnOpenMPMapClause(MapTypeModifier, MapType, DepLinMapLoc, ColonLoc,
- VarList, StartLoc, LParenLoc, EndLoc);
+ Res = ActOnOpenMPMapClause(MapTypeModifier, MapType, IsMapTypeImplicit,
+ DepLinMapLoc, ColonLoc, VarList, StartLoc,
+ LParenLoc, EndLoc);
break;
case OMPC_if:
case OMPC_final:
@@ -8332,8 +8387,9 @@
OMPClause *Sema::ActOnOpenMPMapClause(
OpenMPMapClauseKind MapTypeModifier, OpenMPMapClauseKind MapType,
- SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
- SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc) {
+ bool IsMapTypeImplicit, SourceLocation MapLoc, SourceLocation ColonLoc,
+ ArrayRef<Expr *> VarList, SourceLocation StartLoc, SourceLocation LParenLoc,
+ SourceLocation EndLoc) {
SmallVector<Expr *, 4> Vars;
for (auto &RE : VarList) {
@@ -8440,15 +8496,47 @@
DSAStack, 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) <<
+ static_cast<unsigned>(IsMapTypeImplicit) <<
+ getOpenMPSimpleClauseTypeName(OMPC_map, MapType) <<
+ getOpenMPDirectiveName(DKind);
+ // Proceed to add the variable in a map clause anyway, to prevent
+ // further spurious messages
+ }
+
+ // 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) <<
+ static_cast<unsigned>(IsMapTypeImplicit) <<
+ getOpenMPSimpleClauseTypeName(OMPC_map, MapType) <<
+ getOpenMPDirectiveName(DKind);
+ // Proceed to add the variable in a map clause anyway, to prevent
+ // further spurious messages
+ }
+
Vars.push_back(RE);
MI.RefExpr = RE;
DSAStack->addMapInfoForVar(VD, MI);
}
if (Vars.empty())
return nullptr;
return OMPMapClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars,
- MapTypeModifier, MapType, MapLoc);
+ MapTypeModifier, MapType, IsMapTypeImplicit,
+ MapLoc);
}
OMPClause *Sema::ActOnOpenMPNumTeamsClause(Expr *NumTeams,
Index: lib/Parse/ParseOpenMP.cpp
===================================================================
--- lib/Parse/ParseOpenMP.cpp
+++ lib/Parse/ParseOpenMP.cpp
@@ -34,6 +34,12 @@
{OMPD_unknown /*cancellation*/, OMPD_unknown /*point*/,
OMPD_cancellation_point},
{OMPD_target, OMPD_unknown /*data*/, OMPD_target_data},
+ {OMPD_target, OMPD_unknown /*enter/exit*/,
+ OMPD_unknown /*target enter/exit*/},
+ {OMPD_unknown /*target enter*/, OMPD_unknown /*data*/,
+ OMPD_target_enter_data},
+ {OMPD_unknown /*target exit*/, OMPD_unknown /*data*/,
+ OMPD_target_exit_data},
{OMPD_for, OMPD_simd, OMPD_for_simd},
{OMPD_parallel, OMPD_for, OMPD_parallel_for},
{OMPD_parallel_for, OMPD_simd, OMPD_parallel_for_simd},
@@ -49,8 +55,12 @@
for (unsigned i = 0; i < llvm::array_lengthof(F); ++i) {
if (!Tok.isAnnotation() && DKind == OMPD_unknown) {
TokenMatched =
- (i == 0) &&
- !P.getPreprocessor().getSpelling(Tok).compare("cancellation");
+ ((i == 0) &&
+ !P.getPreprocessor().getSpelling(Tok).compare("cancellation")) ||
+ ((i == 3) &&
+ !P.getPreprocessor().getSpelling(Tok).compare("enter")) ||
+ ((i == 4) &&
+ !P.getPreprocessor().getSpelling(Tok).compare("exit"));
} else {
TokenMatched = DKind == F[i][0] && DKind != OMPD_unknown;
}
@@ -67,7 +77,11 @@
TokenMatched =
((i == 0) &&
!P.getPreprocessor().getSpelling(Tok).compare("point")) ||
- ((i == 1) && !P.getPreprocessor().getSpelling(Tok).compare("data"));
+ ((i == 1 || i == 3 || i == 4) &&
+ !P.getPreprocessor().getSpelling(Tok).compare("data")) ||
+ ((i == 2) &&
+ (!P.getPreprocessor().getSpelling(Tok).compare("enter") ||
+ !P.getPreprocessor().getSpelling(Tok).compare("exit")));
} else {
TokenMatched = SDKind == F[i][1] && SDKind != OMPD_unknown;
}
@@ -138,6 +152,8 @@
case OMPD_cancellation_point:
case OMPD_cancel:
case OMPD_target_data:
+ case OMPD_target_enter_data:
+ case OMPD_target_exit_data:
case OMPD_taskloop:
case OMPD_taskloop_simd:
case OMPD_distribute:
@@ -162,7 +178,7 @@
/// 'barrier' | 'taskwait' | 'flush' | 'ordered' | 'atomic' |
/// 'for simd' | 'parallel for simd' | 'target' | 'target data' |
/// 'taskgroup' | 'teams' | 'taskloop' | 'taskloop simd' {clause} |
-/// 'distribute'
+/// 'distribute' | 'target enter data' | 'target exit data'
/// annot_pragma_openmp_end
///
StmtResult
@@ -213,6 +229,8 @@
case OMPD_taskwait:
case OMPD_cancellation_point:
case OMPD_cancel:
+ case OMPD_target_enter_data:
+ case OMPD_target_exit_data:
if (!StandAloneAllowed) {
Diag(Tok, diag::err_omp_immediate_directive)
<< getOpenMPDirectiveName(DKind) << 0;
@@ -866,6 +884,7 @@
OpenMPLinearClauseKind LinearModifier = OMPC_LINEAR_val;
OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown;
+ bool MapTypeIsImplicit = false;
bool MapTypeModifierSpecified = false;
bool UnexpectedId = false;
SourceLocation DepLinMapLoc;
@@ -916,7 +935,8 @@
Kind, llvm::None, /*TailExpr=*/nullptr, Loc, LOpen,
/*ColonLoc=*/SourceLocation(), Tok.getLocation(),
ReductionIdScopeSpec, DeclarationNameInfo(), DepKind,
- LinearModifier, MapTypeModifier, MapType, DepLinMapLoc);
+ LinearModifier, MapTypeModifier, MapType, MapTypeIsImplicit,
+ DepLinMapLoc);
}
}
if (Tok.is(tok::colon)) {
@@ -980,9 +1000,11 @@
ConsumeToken();
} else {
MapType = OMPC_MAP_tofrom;
+ MapTypeIsImplicit = true;
}
} else {
MapType = OMPC_MAP_tofrom;
+ MapTypeIsImplicit = true;
}
} else {
UnexpectedId = true;
@@ -1063,6 +1085,7 @@
ReductionIdScopeSpec,
ReductionId.isValid() ? Actions.GetNameFromUnqualifiedId(ReductionId)
: DeclarationNameInfo(),
- DepKind, LinearModifier, MapTypeModifier, MapType, DepLinMapLoc);
+ DepKind, LinearModifier, MapTypeModifier, MapType, MapTypeIsImplicit,
+ DepLinMapLoc);
}
Index: lib/CodeGen/CodeGenFunction.h
===================================================================
--- lib/CodeGen/CodeGenFunction.h
+++ lib/CodeGen/CodeGenFunction.h
@@ -2335,6 +2335,8 @@
void EmitOMPAtomicDirective(const OMPAtomicDirective &S);
void EmitOMPTargetDirective(const OMPTargetDirective &S);
void EmitOMPTargetDataDirective(const OMPTargetDataDirective &S);
+ void EmitOMPTargetEnterDataDirective(const OMPTargetEnterDataDirective &S);
+ void EmitOMPTargetExitDataDirective(const OMPTargetExitDataDirective &S);
void EmitOMPTeamsDirective(const OMPTeamsDirective &S);
void
EmitOMPCancellationPointDirective(const OMPCancellationPointDirective &S);
Index: lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- lib/CodeGen/CGStmtOpenMP.cpp
+++ lib/CodeGen/CGStmtOpenMP.cpp
@@ -2660,6 +2660,16 @@
[&CS](CodeGenFunction &CGF) { CGF.EmitStmt(CS->getCapturedStmt()); });
}
+void CodeGenFunction::EmitOMPTargetEnterDataDirective(
+ const OMPTargetEnterDataDirective &S) {
+ llvm_unreachable("CodeGen for 'omp target enter data' is not supported.");
+}
+
+void CodeGenFunction::EmitOMPTargetExitDataDirective(
+ const OMPTargetExitDataDirective &S) {
+ llvm_unreachable("CodeGen for 'omp target exit data' is not supported.");
+}
+
void CodeGenFunction::EmitOMPTaskLoopDirective(const OMPTaskLoopDirective &S) {
// emit the code inside the construct for now
auto CS = cast<CapturedStmt>(S.getAssociatedStmt());
Index: lib/CodeGen/CGStmt.cpp
===================================================================
--- lib/CodeGen/CGStmt.cpp
+++ lib/CodeGen/CGStmt.cpp
@@ -256,6 +256,12 @@
case Stmt::OMPTargetDataDirectiveClass:
EmitOMPTargetDataDirective(cast<OMPTargetDataDirective>(*S));
break;
+ case Stmt::OMPTargetEnterDataDirectiveClass:
+ EmitOMPTargetEnterDataDirective(cast<OMPTargetEnterDataDirective>(*S));
+ break;
+ case Stmt::OMPTargetExitDataDirectiveClass:
+ EmitOMPTargetExitDataDirective(cast<OMPTargetExitDataDirective>(*S));
+ break;
case Stmt::OMPTaskLoopDirectiveClass:
EmitOMPTaskLoopDirective(cast<OMPTaskLoopDirective>(*S));
break;
Index: lib/Basic/OpenMPKinds.cpp
===================================================================
--- lib/Basic/OpenMPKinds.cpp
+++ lib/Basic/OpenMPKinds.cpp
@@ -398,6 +398,26 @@
break;
}
break;
+ case OMPD_target_enter_data:
+ switch (CKind) {
+#define OPENMP_TARGET_ENTER_DATA_CLAUSE(Name) \
+ case OMPC_##Name: \
+ return true;
+#include "clang/Basic/OpenMPKinds.def"
+ default:
+ break;
+ }
+ break;
+ case OMPD_target_exit_data:
+ switch (CKind) {
+#define OPENMP_TARGET_EXIT_DATA_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) \
Index: lib/AST/StmtProfile.cpp
===================================================================
--- lib/AST/StmtProfile.cpp
+++ lib/AST/StmtProfile.cpp
@@ -584,6 +584,16 @@
VisitOMPExecutableDirective(S);
}
+void StmtProfiler::VisitOMPTargetEnterDataDirective(
+ const OMPTargetEnterDataDirective *S) {
+ VisitOMPExecutableDirective(S);
+}
+
+void StmtProfiler::VisitOMPTargetExitDataDirective(
+ const OMPTargetExitDataDirective *S) {
+ VisitOMPExecutableDirective(S);
+}
+
void StmtProfiler::VisitOMPTeamsDirective(const OMPTeamsDirective *S) {
VisitOMPExecutableDirective(S);
}
Index: lib/AST/StmtPrinter.cpp
===================================================================
--- lib/AST/StmtPrinter.cpp
+++ lib/AST/StmtPrinter.cpp
@@ -1051,6 +1051,18 @@
PrintOMPExecutableDirective(Node);
}
+void StmtPrinter::VisitOMPTargetEnterDataDirective(
+ OMPTargetEnterDataDirective *Node) {
+ Indent() << "#pragma omp target enter data ";
+ PrintOMPExecutableDirective(Node);
+}
+
+void StmtPrinter::VisitOMPTargetExitDataDirective(
+ OMPTargetExitDataDirective *Node) {
+ Indent() << "#pragma omp target exit data ";
+ PrintOMPExecutableDirective(Node);
+}
+
void StmtPrinter::VisitOMPTeamsDirective(OMPTeamsDirective *Node) {
Indent() << "#pragma omp teams ";
PrintOMPExecutableDirective(Node);
Index: lib/AST/StmtOpenMP.cpp
===================================================================
--- lib/AST/StmtOpenMP.cpp
+++ lib/AST/StmtOpenMP.cpp
@@ -718,6 +718,50 @@
return new (Mem) OMPTargetDataDirective(N);
}
+OMPTargetEnterDataDirective *OMPTargetEnterDataDirective::Create(
+ const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+ ArrayRef<OMPClause *> Clauses) {
+ void *Mem =
+ C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetEnterDataDirective),
+ llvm::alignOf<OMPClause *>()) +
+ sizeof(OMPClause *) * Clauses.size());
+ OMPTargetEnterDataDirective *Dir =
+ new (Mem) OMPTargetEnterDataDirective(StartLoc, EndLoc, Clauses.size());
+ Dir->setClauses(Clauses);
+ return Dir;
+}
+
+OMPTargetEnterDataDirective *OMPTargetEnterDataDirective::CreateEmpty(
+ const ASTContext &C, unsigned N, EmptyShell) {
+ void *Mem =
+ C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetEnterDataDirective),
+ llvm::alignOf<OMPClause *>()) +
+ sizeof(OMPClause *) * N);
+ return new (Mem) OMPTargetEnterDataDirective(N);
+}
+
+OMPTargetExitDataDirective *OMPTargetExitDataDirective::Create(
+ const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+ ArrayRef<OMPClause *> Clauses) {
+ void *Mem =
+ C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetExitDataDirective),
+ llvm::alignOf<OMPClause *>()) +
+ sizeof(OMPClause *) * Clauses.size());
+ OMPTargetExitDataDirective *Dir =
+ new (Mem) OMPTargetExitDataDirective(StartLoc, EndLoc, Clauses.size());
+ Dir->setClauses(Clauses);
+ return Dir;
+}
+
+OMPTargetExitDataDirective *OMPTargetExitDataDirective::CreateEmpty(
+ const ASTContext &C, unsigned N, EmptyShell) {
+ void *Mem =
+ C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetExitDataDirective),
+ llvm::alignOf<OMPClause *>()) +
+ sizeof(OMPClause *) * N);
+ return new (Mem) OMPTargetExitDataDirective(N);
+}
+
OMPTeamsDirective *OMPTeamsDirective::Create(const ASTContext &C,
SourceLocation StartLoc,
SourceLocation EndLoc,
Index: lib/AST/OpenMPClause.cpp
===================================================================
--- lib/AST/OpenMPClause.cpp
+++ lib/AST/OpenMPClause.cpp
@@ -400,10 +400,12 @@
SourceLocation EndLoc, ArrayRef<Expr *> VL,
OpenMPMapClauseKind TypeModifier,
OpenMPMapClauseKind Type,
+ bool TypeIsImplicit,
SourceLocation TypeLoc) {
void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size()));
OMPMapClause *Clause = new (Mem) OMPMapClause(
- TypeModifier, Type, TypeLoc, StartLoc, LParenLoc, EndLoc, VL.size());
+ TypeModifier, Type, TypeIsImplicit, TypeLoc, StartLoc, LParenLoc, EndLoc,
+ VL.size());
Clause->setVarRefs(VL);
Clause->setMapTypeModifier(TypeModifier);
Clause->setMapType(Type);
Index: include/clang/Serialization/ASTBitCodes.h
===================================================================
--- include/clang/Serialization/ASTBitCodes.h
+++ include/clang/Serialization/ASTBitCodes.h
@@ -1443,6 +1443,8 @@
STMT_OMP_ATOMIC_DIRECTIVE,
STMT_OMP_TARGET_DIRECTIVE,
STMT_OMP_TARGET_DATA_DIRECTIVE,
+ STMT_OMP_TARGET_ENTER_DATA_DIRECTIVE,
+ STMT_OMP_TARGET_EXIT_DATA_DIRECTIVE,
STMT_OMP_TEAMS_DIRECTIVE,
STMT_OMP_TASKGROUP_DIRECTIVE,
STMT_OMP_CANCELLATION_POINT_DIRECTIVE,
Index: include/clang/Sema/Sema.h
===================================================================
--- include/clang/Sema/Sema.h
+++ include/clang/Sema/Sema.h
@@ -7761,6 +7761,8 @@
ExprResult
VerifyPositiveIntegerConstantInClause(Expr *Op, OpenMPClauseKind CKind,
bool StrictlyPositive = true);
+ /// \brief Check for existence of a map clause
+ bool HasMapClause(ArrayRef<OMPClause *> Clauses);
public:
/// \brief Return true if the provided declaration \a VD should be captured by
@@ -7938,6 +7940,16 @@
StmtResult ActOnOpenMPTargetDataDirective(ArrayRef<OMPClause *> Clauses,
Stmt *AStmt, SourceLocation StartLoc,
SourceLocation EndLoc);
+ /// \brief Called on well-formed '\#pragma omp target enter data' after
+ /// parsing of the associated statement.
+ StmtResult ActOnOpenMPTargetEnterDataDirective(ArrayRef<OMPClause *> Clauses,
+ SourceLocation StartLoc,
+ SourceLocation EndLoc);
+ /// \brief Called on well-formed '\#pragma omp target exit data' after
+ /// parsing of the associated statement.
+ StmtResult ActOnOpenMPTargetExitDataDirective(ArrayRef<OMPClause *> Clauses,
+ SourceLocation StartLoc,
+ SourceLocation EndLoc);
/// \brief Called on well-formed '\#pragma omp teams' after parsing of the
/// associated statement.
StmtResult ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
@@ -8099,7 +8111,8 @@
CXXScopeSpec &ReductionIdScopeSpec,
const DeclarationNameInfo &ReductionId, OpenMPDependClauseKind DepKind,
OpenMPLinearClauseKind LinKind, OpenMPMapClauseKind MapTypeModifier,
- OpenMPMapClauseKind MapType, SourceLocation DepLinMapLoc);
+ OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
+ SourceLocation DepLinMapLoc);
/// \brief Called on well-formed 'private' clause.
OMPClause *ActOnOpenMPPrivateClause(ArrayRef<Expr *> VarList,
SourceLocation StartLoc,
@@ -8168,8 +8181,9 @@
/// \brief Called on well-formed 'map' clause.
OMPClause *ActOnOpenMPMapClause(
OpenMPMapClauseKind MapTypeModifier, OpenMPMapClauseKind MapType,
- SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
- SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc);
+ bool IsMapTypeImplicit, SourceLocation MapLoc, SourceLocation ColonLoc,
+ ArrayRef<Expr *> VarList, SourceLocation StartLoc, SourceLocation LParenLoc,
+ SourceLocation EndLoc);
/// \brief Called on well-formed 'num_teams' clause.
OMPClause *ActOnOpenMPNumTeamsClause(Expr *NumTeams, SourceLocation StartLoc,
SourceLocation LParenLoc,
Index: include/clang/Basic/StmtNodes.td
===================================================================
--- include/clang/Basic/StmtNodes.td
+++ include/clang/Basic/StmtNodes.td
@@ -216,6 +216,8 @@
def OMPAtomicDirective : DStmt<OMPExecutableDirective>;
def OMPTargetDirective : DStmt<OMPExecutableDirective>;
def OMPTargetDataDirective : DStmt<OMPExecutableDirective>;
+def OMPTargetEnterDataDirective : DStmt<OMPExecutableDirective>;
+def OMPTargetExitDataDirective : 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
@@ -60,6 +60,12 @@
#ifndef OPENMP_TARGET_DATA_CLAUSE
# define OPENMP_TARGET_DATA_CLAUSE(Name)
#endif
+#ifndef OPENMP_TARGET_ENTER_DATA_CLAUSE
+#define OPENMP_TARGET_ENTER_DATA_CLAUSE(Name)
+#endif
+#ifndef OPENMP_TARGET_EXIT_DATA_CLAUSE
+#define OPENMP_TARGET_EXIT_DATA_CLAUSE(Name)
+#endif
#ifndef OPENMP_TEAMS_CLAUSE
# define OPENMP_TEAMS_CLAUSE(Name)
#endif
@@ -125,6 +131,8 @@
OPENMP_DIRECTIVE(teams)
OPENMP_DIRECTIVE(cancel)
OPENMP_DIRECTIVE_EXT(target_data, "target data")
+OPENMP_DIRECTIVE_EXT(target_enter_data, "target enter data")
+OPENMP_DIRECTIVE_EXT(target_exit_data, "target exit data")
OPENMP_DIRECTIVE_EXT(parallel_for, "parallel for")
OPENMP_DIRECTIVE_EXT(parallel_for_simd, "parallel for simd")
OPENMP_DIRECTIVE_EXT(parallel_sections, "parallel sections")
@@ -349,6 +357,18 @@
OPENMP_TARGET_DATA_CLAUSE(device)
OPENMP_TARGET_DATA_CLAUSE(map)
+// Clauses allowed for OpenMP directive 'target enter data'.
+// TODO More clauses for 'target enter data' directive.
+OPENMP_TARGET_ENTER_DATA_CLAUSE(if)
+OPENMP_TARGET_ENTER_DATA_CLAUSE(device)
+OPENMP_TARGET_ENTER_DATA_CLAUSE(map)
+
+// Clauses allowed for OpenMP directive 'target exit data'.
+// TODO More clauses for 'target exit data' directive.
+OPENMP_TARGET_EXIT_DATA_CLAUSE(if)
+OPENMP_TARGET_EXIT_DATA_CLAUSE(device)
+OPENMP_TARGET_EXIT_DATA_CLAUSE(map)
+
// Clauses allowed for OpenMP directive 'teams'.
// TODO More clauses for 'teams' directive.
OPENMP_TEAMS_CLAUSE(default)
@@ -443,6 +463,8 @@
#undef OPENMP_ATOMIC_CLAUSE
#undef OPENMP_TARGET_CLAUSE
#undef OPENMP_TARGET_DATA_CLAUSE
+#undef OPENMP_TARGET_ENTER_DATA_CLAUSE
+#undef OPENMP_TARGET_EXIT_DATA_CLAUSE
#undef OPENMP_TEAMS_CLAUSE
#undef OPENMP_SIMD_CLAUSE
#undef OPENMP_FOR_CLAUSE
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -7927,6 +7927,10 @@
"variable already marked as mapped in current construct">;
def err_omp_not_mappable_type : Error<
"type %0 is not mappable to target">;
+def err_omp_invalid_map_type_for_directive : Error<
+ "%select{map type '%1' is not allowed|map type must be specified}0 for '#pragma omp %2'">;
+def err_omp_no_map_for_directive : Error<
+ "expected at least one map clause for '#pragma omp %0'">;
def note_omp_polymorphic_in_target : Note<
"mappable type cannot be polymorphic">;
def note_omp_static_member_in_target : Note<
Index: include/clang/AST/StmtOpenMP.h
===================================================================
--- include/clang/AST/StmtOpenMP.h
+++ include/clang/AST/StmtOpenMP.h
@@ -2039,6 +2039,124 @@
}
};
+/// \brief This represents '#pragma omp target enter data' directive.
+///
+/// \code
+/// #pragma omp target enter data device(0) if(a) map(b[:])
+/// \endcode
+/// In this example directive '#pragma omp target enter data' has clauses
+/// 'device' with the value '0', 'if' with condition 'a' and 'map' with array
+/// section 'b[:]'.
+///
+class OMPTargetEnterDataDirective : 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.
+ ///
+ OMPTargetEnterDataDirective(SourceLocation StartLoc, SourceLocation EndLoc,
+ unsigned NumClauses)
+ : OMPExecutableDirective(this, OMPTargetEnterDataDirectiveClass,
+ OMPD_target_enter_data, StartLoc, EndLoc,
+ NumClauses, /*NumChildren=*/0) {}
+
+ /// \brief Build an empty directive.
+ ///
+ /// \param NumClauses Number of clauses.
+ ///
+ explicit OMPTargetEnterDataDirective(unsigned NumClauses)
+ : OMPExecutableDirective(this, OMPTargetEnterDataDirectiveClass,
+ OMPD_target_enter_data, SourceLocation(),
+ SourceLocation(), NumClauses,
+ /*NumChildren=*/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.
+ /// \param AssociatedStmt Statement, associated with the directive.
+ ///
+ static OMPTargetEnterDataDirective *
+ Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+ ArrayRef<OMPClause *> Clauses);
+
+ /// \brief Creates an empty directive with the place for \a N clauses.
+ ///
+ /// \param C AST context.
+ /// \param N The number of clauses.
+ ///
+ static OMPTargetEnterDataDirective *CreateEmpty(const ASTContext &C,
+ unsigned N, EmptyShell);
+
+ static bool classof(const Stmt *T) {
+ return T->getStmtClass() == OMPTargetEnterDataDirectiveClass;
+ }
+};
+
+/// \brief This represents '#pragma omp target exit data' directive.
+///
+/// \code
+/// #pragma omp target exit data device(0) if(a) map(b[:])
+/// \endcode
+/// In this example directive '#pragma omp target exit data' has clauses
+/// 'device' with the value '0', 'if' with condition 'a' and 'map' with array
+/// section 'b[:]'.
+///
+class OMPTargetExitDataDirective : 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.
+ ///
+ OMPTargetExitDataDirective(SourceLocation StartLoc, SourceLocation EndLoc,
+ unsigned NumClauses)
+ : OMPExecutableDirective(this, OMPTargetExitDataDirectiveClass,
+ OMPD_target_exit_data, StartLoc, EndLoc,
+ NumClauses, /*NumChildren=*/0) {}
+
+ /// \brief Build an empty directive.
+ ///
+ /// \param NumClauses Number of clauses.
+ ///
+ explicit OMPTargetExitDataDirective(unsigned NumClauses)
+ : OMPExecutableDirective(this, OMPTargetExitDataDirectiveClass,
+ OMPD_target_exit_data, SourceLocation(),
+ SourceLocation(), NumClauses,
+ /*NumChildren=*/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.
+ /// \param AssociatedStmt Statement, associated with the directive.
+ ///
+ static OMPTargetExitDataDirective *
+ Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+ ArrayRef<OMPClause *> Clauses);
+
+ /// \brief Creates an empty directive with the place for \a N clauses.
+ ///
+ /// \param C AST context.
+ /// \param N The number of clauses.
+ ///
+ static OMPTargetExitDataDirective *CreateEmpty(const ASTContext &C,
+ unsigned N, EmptyShell);
+
+ static bool classof(const Stmt *T) {
+ return T->getStmtClass() == OMPTargetExitDataDirectiveClass;
+ }
+};
+
/// \brief This represents '#pragma omp teams' directive.
///
/// \code
Index: include/clang/AST/RecursiveASTVisitor.h
===================================================================
--- include/clang/AST/RecursiveASTVisitor.h
+++ include/clang/AST/RecursiveASTVisitor.h
@@ -2429,6 +2429,12 @@
DEF_TRAVERSE_STMT(OMPTargetDataDirective,
{ TRY_TO(TraverseOMPExecutableDirective(S)); })
+DEF_TRAVERSE_STMT(OMPTargetEnterDataDirective,
+ { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
+DEF_TRAVERSE_STMT(OMPTargetExitDataDirective,
+ { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
DEF_TRAVERSE_STMT(OMPTeamsDirective,
{ TRY_TO(TraverseOMPExecutableDirective(S)); })
Index: include/clang/AST/OpenMPClause.h
===================================================================
--- include/clang/AST/OpenMPClause.h
+++ include/clang/AST/OpenMPClause.h
@@ -2741,6 +2741,8 @@
OpenMPMapClauseKind MapTypeModifier;
/// \brief Map type for the 'map' clause.
OpenMPMapClauseKind MapType;
+ /// \brief Is this an implicit map type or not.
+ bool MapTypeIsImplicit;
/// \brief Location of the map type.
SourceLocation MapLoc;
/// \brief Colon location.
@@ -2777,20 +2779,23 @@
/// \param N Number of the variables in the clause.
///
explicit OMPMapClause(OpenMPMapClauseKind MapTypeModifier,
- OpenMPMapClauseKind MapType, SourceLocation MapLoc,
- SourceLocation StartLoc, SourceLocation LParenLoc,
- SourceLocation EndLoc, unsigned N)
+ OpenMPMapClauseKind MapType, bool MapTypeIsImplicit,
+ SourceLocation MapLoc, SourceLocation StartLoc,
+ SourceLocation LParenLoc, SourceLocation EndLoc,
+ unsigned N)
: OMPVarListClause<OMPMapClause>(OMPC_map, StartLoc, LParenLoc, EndLoc, N),
- MapTypeModifier(MapTypeModifier), MapType(MapType), MapLoc(MapLoc) {}
+ MapTypeModifier(MapTypeModifier), MapType(MapType),
+ MapTypeIsImplicit(MapTypeIsImplicit), MapLoc(MapLoc) {}
/// \brief Build an empty clause.
///
/// \param N Number of variables.
///
explicit OMPMapClause(unsigned N)
: OMPVarListClause<OMPMapClause>(OMPC_map, SourceLocation(),
SourceLocation(), SourceLocation(), N),
- MapTypeModifier(OMPC_MAP_unknown), MapType(OMPC_MAP_unknown), MapLoc() {}
+ MapTypeModifier(OMPC_MAP_unknown), MapType(OMPC_MAP_unknown),
+ MapTypeIsImplicit(false), MapLoc() {}
public:
/// \brief Creates clause with a list of variables \a VL.
@@ -2807,7 +2812,8 @@
SourceLocation LParenLoc,
SourceLocation EndLoc, ArrayRef<Expr *> VL,
OpenMPMapClauseKind TypeModifier,
- OpenMPMapClauseKind Type, SourceLocation TypeLoc);
+ OpenMPMapClauseKind Type, bool TypeIsImplicit,
+ SourceLocation TypeLoc);
/// \brief Creates an empty clause with the place for \a N variables.
///
/// \param C AST context.
@@ -2818,6 +2824,11 @@
/// \brief Fetches mapping kind for the clause.
OpenMPMapClauseKind getMapType() const LLVM_READONLY { return MapType; }
+ /// \brief Is this an implicit map type?
+ bool isImplicitMapType() const LLVM_READONLY {
+ return MapTypeIsImplicit;
+ }
+
/// \brief Fetches the map type modifier for the clause.
OpenMPMapClauseKind getMapTypeModifier() const LLVM_READONLY {
return MapTypeModifier;
Index: include/clang-c/Index.h
===================================================================
--- include/clang-c/Index.h
+++ include/clang-c/Index.h
@@ -2274,7 +2274,15 @@
*/
CXCursor_OMPDistributeDirective = 260,
- CXCursor_LastStmt = CXCursor_OMPDistributeDirective,
+ /** \brief OpenMP target enter data directive.
+ */
+ CXCursor_OMPTargetEnterDataDirective = 261,
+
+ /** \brief OpenMP target exit data directive.
+ */
+ CXCursor_OMPTargetExitDataDirective = 262,
+
+ CXCursor_LastStmt = CXCursor_OMPTargetExitDataDirective,
/**
* \brief Cursor that represents the translation unit itself.
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
