kkwli0 updated this revision to Diff 29959.
kkwli0 added a comment.
- this is the full diff
- sync up the code in ParseOpenMPDirectiveKind
- use the existing infrastructure for generating if clause on target data
Thanks.
http://reviews.llvm.org/D10765
Files:
include/clang-c/Index.h
include/clang/AST/DataRecursiveASTVisitor.h
include/clang/AST/RecursiveASTVisitor.h
include/clang/AST/StmtOpenMP.h
include/clang/Basic/OpenMPKinds.def
include/clang/Basic/StmtNodes.td
include/clang/Sema/Sema.h
include/clang/Serialization/ASTBitCodes.h
lib/AST/Stmt.cpp
lib/AST/StmtPrinter.cpp
lib/AST/StmtProfile.cpp
lib/Basic/OpenMPKinds.cpp
lib/CodeGen/CGExpr.cpp
lib/CodeGen/CGOpenMPRuntime.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_data_ast_print.cpp
test/OpenMP/target_data_if_messages.cpp
test/OpenMP/target_data_messages.c
tools/libclang/CIndex.cpp
tools/libclang/CXCursor.cpp
Index: tools/libclang/CXCursor.cpp
===================================================================
--- tools/libclang/CXCursor.cpp
+++ tools/libclang/CXCursor.cpp
@@ -585,6 +585,9 @@
case Stmt::OMPTargetDirectiveClass:
K = CXCursor_OMPTargetDirective;
break;
+ case Stmt::OMPTargetDataDirectiveClass:
+ K = CXCursor_OMPTargetDataDirective;
+ break;
case Stmt::OMPTeamsDirectiveClass:
K = CXCursor_OMPTeamsDirective;
break;
Index: tools/libclang/CIndex.cpp
===================================================================
--- tools/libclang/CIndex.cpp
+++ tools/libclang/CIndex.cpp
@@ -1930,6 +1930,7 @@
void VisitOMPOrderedDirective(const OMPOrderedDirective *D);
void VisitOMPAtomicDirective(const OMPAtomicDirective *D);
void VisitOMPTargetDirective(const OMPTargetDirective *D);
+ void VisitOMPTargetDataDirective(const OMPTargetDataDirective *D);
void VisitOMPTeamsDirective(const OMPTeamsDirective *D);
private:
@@ -2549,6 +2550,11 @@
VisitOMPExecutableDirective(D);
}
+void EnqueueVisitor::VisitOMPTargetDataDirective(const
+ OMPTargetDataDirective *D) {
+ VisitOMPExecutableDirective(D);
+}
+
void EnqueueVisitor::VisitOMPTeamsDirective(const OMPTeamsDirective *D) {
VisitOMPExecutableDirective(D);
}
@@ -4349,6 +4355,8 @@
return cxstring::createRef("OMPAtomicDirective");
case CXCursor_OMPTargetDirective:
return cxstring::createRef("OMPTargetDirective");
+ case CXCursor_OMPTargetDataDirective:
+ return cxstring::createRef("OMPTargetDataDirective");
case CXCursor_OMPTeamsDirective:
return cxstring::createRef("OMPTeamsDirective");
case CXCursor_OMPCancellationPointDirective:
Index: test/OpenMP/target_data_messages.c
===================================================================
--- /dev/null
+++ test/OpenMP/target_data_messages.c
@@ -0,0 +1,27 @@
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s
+
+void foo() { }
+
+int main(int argc, char **argv) {
+ L1:
+ foo();
+ #pragma omp target data
+ {
+ foo();
+ goto L1; // expected-error {{use of undeclared label 'L1'}}
+ }
+ goto L2; // expected-error {{use of undeclared label 'L2'}}
+ #pragma omp target data
+ L2:
+ foo();
+
+ #pragma omp target data(i) // expected-warning {{extra tokens at the end of '#pragma omp target data' are ignored}}
+ {
+ foo();
+ }
+ #pragma omp target unknown // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}
+ {
+ foo();
+ }
+ return 0;
+}
Index: test/OpenMP/target_data_if_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_data_if_messages.cpp
@@ -0,0 +1,26 @@
+// 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) {
+ #pragma omp target data if // expected-error {{expected '(' after 'if'}}
+ #pragma omp target data if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target data if () // expected-error {{expected expression}}
+ #pragma omp target data if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ #pragma omp target data if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target data' are ignored}}
+ #pragma omp target data if (argc > 0 ? argv[1] : argv[2])
+ #pragma omp target data if (argc + argc)
+ #pragma omp target data if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target data' cannot contain more than one 'if' clause}}
+ #pragma omp target data if (S1) // expected-error {{'S1' does not refer to a value}}
+ #pragma omp target data if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+
+ return 0;
+}
Index: test/OpenMP/target_data_ast_print.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_data_ast_print.cpp
@@ -0,0 +1,33 @@
+// 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() {}
+
+int main (int argc, char **argv) {
+ int b = argc, c, d, e, f, g;
+ static int a;
+// CHECK: static int a;
+
+#pragma omp target data
+// CHECK: #pragma omp target data
+ a=2;
+// CHECK-NEXT: a = 2;
+#pragma omp target data if (b)
+// CHECK: #pragma omp target data if(b)
+ foo();
+// CHECK-NEXT: foo();
+
+#pragma omp target data if (b > g)
+// CHECK: #pragma omp target data if(b > g)
+ foo();
+// CHECK-NEXT: foo();
+
+ return (0);
+}
+
+#endif
Index: lib/StaticAnalyzer/Core/ExprEngine.cpp
===================================================================
--- lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -821,6 +821,7 @@
case Stmt::OMPOrderedDirectiveClass:
case Stmt::OMPAtomicDirectiveClass:
case Stmt::OMPTargetDirectiveClass:
+ case Stmt::OMPTargetDataDirectiveClass:
case Stmt::OMPTeamsDirectiveClass:
case Stmt::OMPCancellationPointDirectiveClass:
case Stmt::OMPCancelDirectiveClass:
Index: lib/Serialization/ASTWriterStmt.cpp
===================================================================
--- lib/Serialization/ASTWriterStmt.cpp
+++ lib/Serialization/ASTWriterStmt.cpp
@@ -2063,6 +2063,13 @@
Code = serialization::STMT_OMP_TARGET_DIRECTIVE;
}
+void ASTStmtWriter::VisitOMPTargetDataDirective(OMPTargetDataDirective *D) {
+ VisitStmt(D);
+ Record.push_back(D->getNumClauses());
+ VisitOMPExecutableDirective(D);
+ Code = serialization::STMT_OMP_TARGET_DATA_DIRECTIVE;
+}
+
void ASTStmtWriter::VisitOMPTaskyieldDirective(OMPTaskyieldDirective *D) {
VisitStmt(D);
VisitOMPExecutableDirective(D);
Index: lib/Serialization/ASTReaderStmt.cpp
===================================================================
--- lib/Serialization/ASTReaderStmt.cpp
+++ lib/Serialization/ASTReaderStmt.cpp
@@ -2242,6 +2242,12 @@
VisitOMPExecutableDirective(D);
}
+void ASTStmtReader::VisitOMPTargetDataDirective(OMPTargetDataDirective *D) {
+ VisitStmt(D);
+ ++Idx;
+ VisitOMPExecutableDirective(D);
+}
+
void ASTStmtReader::VisitOMPTeamsDirective(OMPTeamsDirective *D) {
VisitStmt(D);
// The NumClauses field was read in ReadStmtFromStream.
@@ -2860,6 +2866,11 @@
Context, Record[ASTStmtReader::NumStmtFields], Empty);
break;
+ case STMT_OMP_TARGET_DATA_DIRECTIVE:
+ S = OMPTargetDataDirective::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
@@ -7115,6 +7115,17 @@
}
template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOMPTargetDataDirective(
+ OMPTargetDataDirective *D) {
+ DeclarationNameInfo DirName;
+ getDerived().getSema().StartOpenMPDSABlock(OMPD_target_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;
Index: lib/Sema/SemaOpenMP.cpp
===================================================================
--- lib/Sema/SemaOpenMP.cpp
+++ lib/Sema/SemaOpenMP.cpp
@@ -1300,6 +1300,7 @@
Params);
break;
}
+ case OMPD_target_data:
case OMPD_target: {
Sema::CapturedParamNameType Params[] = {
std::make_pair(StringRef(), QualType()) // __context with shared vars
@@ -2124,6 +2125,10 @@
"No associated statement allowed for 'omp cancel' directive");
Res = ActOnOpenMPCancelDirective(StartLoc, EndLoc, CancelRegion);
break;
+ case OMPD_target_data:
+ Res = ActOnOpenMPTargetDataDirective(ClausesWithImplicit, AStmt, StartLoc,
+ EndLoc);
+ break;
case OMPD_threadprivate:
llvm_unreachable("OpenMP Directive is not allowed");
case OMPD_unknown:
@@ -4320,6 +4325,16 @@
return OMPTargetDirective::Create(Context, StartLoc, EndLoc, Clauses, AStmt);
}
+StmtResult Sema::ActOnOpenMPTargetDataDirective(ArrayRef<OMPClause *> Clauses,
+ Stmt *AStmt,
+ SourceLocation StartLoc,
+ SourceLocation EndLoc) {
+ getCurFunction()->setHasBranchProtectedScope();
+
+ return OMPTargetDataDirective::Create(Context, StartLoc, EndLoc, Clauses,
+ AStmt);
+}
+
StmtResult Sema::ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
Stmt *AStmt, SourceLocation StartLoc,
SourceLocation EndLoc) {
@@ -6593,4 +6608,3 @@
return OMPDependClause::Create(Context, StartLoc, LParenLoc, EndLoc, DepKind,
DepLoc, ColonLoc, Vars);
}
-
Index: lib/Parse/ParseOpenMP.cpp
===================================================================
--- lib/Parse/ParseOpenMP.cpp
+++ lib/Parse/ParseOpenMP.cpp
@@ -19,6 +19,7 @@
#include "clang/Parse/Parser.h"
#include "clang/Sema/Scope.h"
#include "llvm/ADT/PointerIntPair.h"
+
using namespace clang;
//===----------------------------------------------------------------------===//
@@ -32,6 +33,7 @@
const OpenMPDirectiveKind F[][3] = {
{OMPD_unknown /*cancellation*/, OMPD_unknown /*point*/,
OMPD_cancellation_point},
+ {OMPD_target, OMPD_unknown /*data*/, OMPD_target_data},
{OMPD_for, OMPD_simd, OMPD_for_simd},
{OMPD_parallel, OMPD_for, OMPD_parallel_for},
{OMPD_parallel_for, OMPD_simd, OMPD_parallel_for_simd},
@@ -41,6 +43,7 @@
Tok.isAnnotation()
? OMPD_unknown
: getOpenMPDirectiveKind(P.getPreprocessor().getSpelling(Tok));
+
bool TokenMatched = false;
for (unsigned i = 0; i < llvm::array_lengthof(F); ++i) {
if (!Tok.isAnnotation() && DKind == OMPD_unknown) {
@@ -50,18 +53,24 @@
} else {
TokenMatched = DKind == F[i][0] && DKind != OMPD_unknown;
}
+
if (TokenMatched) {
Tok = P.getPreprocessor().LookAhead(0);
+ auto TokenIsAnnotation = Tok.isAnnotation();
auto SDKind =
- Tok.isAnnotation()
+ TokenIsAnnotation
? OMPD_unknown
: getOpenMPDirectiveKind(P.getPreprocessor().getSpelling(Tok));
- if (!Tok.isAnnotation() && DKind == OMPD_unknown) {
+
+ if (!TokenIsAnnotation && DKind == OMPD_unknown) {
TokenMatched =
(i == 0) && !P.getPreprocessor().getSpelling(Tok).compare("point");
+ } else if (!TokenIsAnnotation && DKind == OMPD_target) {
+ TokenMatched = !P.getPreprocessor().getSpelling(Tok).compare("data");
} else {
TokenMatched = SDKind == F[i][1] && SDKind != OMPD_unknown;
}
+
if (TokenMatched) {
P.ConsumeToken();
DKind = F[i][2];
@@ -146,8 +155,8 @@
/// 'section' | 'single' | 'master' | 'critical' [ '(' <name> ')' ] |
/// 'parallel for' | 'parallel sections' | 'task' | 'taskyield' |
/// 'barrier' | 'taskwait' | 'flush' | 'ordered' | 'atomic' |
-/// 'for simd' | 'parallel for simd' | 'target' | 'teams' | 'taskgroup'
-/// {clause}
+/// 'for simd' | 'parallel for simd' | 'target' | 'target data' |
+/// 'taskgroup' | 'teams' {clause}
/// annot_pragma_openmp_end
///
StmtResult
@@ -221,7 +230,8 @@
case OMPD_atomic:
case OMPD_target:
case OMPD_teams:
- case OMPD_taskgroup: {
+ case OMPD_taskgroup:
+ case OMPD_target_data: {
ConsumeToken();
// Parse directive name of the 'critical' directive if any.
if (DKind == OMPD_critical) {
@@ -385,7 +395,7 @@
/// lastprivate-clause | reduction-clause | proc_bind-clause |
/// schedule-clause | copyin-clause | copyprivate-clause | untied-clause |
/// mergeable-clause | flush-clause | read-clause | write-clause |
-/// update-clause | capture-clause | seq_cst-clause
+/// update-clause | capture-clause | seq_cst-clause
///
OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
OpenMPClauseKind CKind, bool FirstClause) {
Index: lib/CodeGen/CodeGenFunction.h
===================================================================
--- lib/CodeGen/CodeGenFunction.h
+++ lib/CodeGen/CodeGenFunction.h
@@ -2213,11 +2213,17 @@
void EmitOMPOrderedDirective(const OMPOrderedDirective &S);
void EmitOMPAtomicDirective(const OMPAtomicDirective &S);
void EmitOMPTargetDirective(const OMPTargetDirective &S);
+ void EmitOMPTargetDataDirective(const OMPTargetDataDirective &S);
void EmitOMPTeamsDirective(const OMPTeamsDirective &S);
void
EmitOMPCancellationPointDirective(const OMPCancellationPointDirective &S);
void EmitOMPCancelDirective(const OMPCancelDirective &S);
+ void EmitOMPIfClause(
+ CodeGenFunction &CGF, const Expr *Cond,
+ const llvm::function_ref<void(CodeGenFunction &)> &ThenGen,
+ const llvm::function_ref<void(CodeGenFunction &)> &ElseGen);
+
/// \brief Emit inner loop of the worksharing/simd construct.
///
/// \param S Directive, for which the inner loop must be emitted.
Index: lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- lib/CodeGen/CGStmtOpenMP.cpp
+++ lib/CodeGen/CGStmtOpenMP.cpp
@@ -2121,3 +2121,32 @@
: BreakContinueStack.back().BreakBlock;
return JumpDest();
}
+
+// Generate the instructions for '#pragma omp target data' directive.
+void CodeGenFunction::EmitOMPTargetDataDirective(
+ const OMPTargetDataDirective &S) {
+
+ // if clause condition
+ const Expr *IfCond = nullptr;
+ if (auto C = S.getSingleClause(OMPC_if)) {
+ IfCond = cast<OMPIfClause>(C)->getCondition();
+ }
+
+ if (IfCond) {
+ auto &&ThenGen = [](CodeGenFunction &CGF) { /*code gen for data mapping*/ };
+ auto &&ElseGen = [](CodeGenFunction &CGF) {};
+ CodeGenFunction::EmitOMPIfClause(*this, IfCond, ThenGen, ElseGen);
+ }
+
+ auto CS = cast<CapturedStmt>(S.getAssociatedStmt());
+
+ CGM.getOpenMPRuntime().emitInlinedDirective(
+ *this, OMPD_target_data,
+ [&CS](CodeGenFunction &CGF) { CGF.EmitStmt(CS->getCapturedStmt()); });
+
+ if (IfCond) {
+ auto &&ThenGen = [](CodeGenFunction &CGF) { /*code gen for data mapping*/ };
+ auto &&ElseGen = [](CodeGenFunction &CGF) {};
+ CodeGenFunction::EmitOMPIfClause(*this, IfCond, ThenGen, ElseGen);
+ }
+}
Index: lib/CodeGen/CGStmt.cpp
===================================================================
--- lib/CodeGen/CGStmt.cpp
+++ lib/CodeGen/CGStmt.cpp
@@ -246,6 +246,9 @@
case Stmt::OMPCancelDirectiveClass:
EmitOMPCancelDirective(cast<OMPCancelDirective>(*S));
break;
+ case Stmt::OMPTargetDataDirectiveClass:
+ EmitOMPTargetDataDirective(cast<OMPTargetDataDirective>(*S));
+ break;
}
}
Index: lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- lib/CodeGen/CGOpenMPRuntime.cpp
+++ lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1096,9 +1096,9 @@
/// } else {
/// ElseGen();
/// }
-static void emitOMPIfClause(CodeGenFunction &CGF, const Expr *Cond,
- const RegionCodeGenTy &ThenGen,
- const RegionCodeGenTy &ElseGen) {
+void CodeGenFunction::EmitOMPIfClause(CodeGenFunction &CGF, const Expr *Cond,
+ const RegionCodeGenTy &ThenGen,
+ const RegionCodeGenTy &ElseGen) {
CodeGenFunction::LexicalScope ConditionScope(CGF, Cond->getSourceRange());
// If the condition constant folds and can be elided, try to avoid emitting
@@ -1190,7 +1190,7 @@
createRuntimeFunction(OMPRTL__kmpc_end_serialized_parallel), EndArgs);
};
if (IfCond) {
- emitOMPIfClause(CGF, IfCond, ThenGen, ElseGen);
+ CGF.EmitOMPIfClause(CGF, IfCond, ThenGen, ElseGen);
} else {
CodeGenFunction::RunCleanupsScope Scope(CGF);
ThenGen(CGF);
@@ -2430,7 +2430,7 @@
CGF.EmitCallOrInvoke(TaskEntry, OutlinedFnArgs);
};
if (IfCond) {
- emitOMPIfClause(CGF, IfCond, ThenCodeGen, ElseCodeGen);
+ CGF.EmitOMPIfClause(CGF, IfCond, ThenCodeGen, ElseCodeGen);
} else {
CodeGenFunction::RunCleanupsScope Scope(CGF);
ThenCodeGen(CGF);
Index: lib/CodeGen/CGExpr.cpp
===================================================================
--- lib/CodeGen/CGExpr.cpp
+++ lib/CodeGen/CGExpr.cpp
@@ -1959,7 +1959,8 @@
else
return EmitCapturedFieldLValue(*this, CapturedStmtInfo->lookup(VD),
CapturedStmtInfo->getContextValue());
- }
+ }
+
assert(isa<BlockDecl>(CurCodeDecl));
return MakeAddrLValue(GetAddrOfBlockDecl(VD, VD->hasAttr<BlocksAttr>()),
T, Alignment);
Index: lib/Basic/OpenMPKinds.cpp
===================================================================
--- lib/Basic/OpenMPKinds.cpp
+++ lib/Basic/OpenMPKinds.cpp
@@ -328,6 +328,16 @@
break;
}
break;
+ case OMPD_target_data:
+ switch (CKind) {
+#define OPENMP_TARGET_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
@@ -533,6 +533,10 @@
VisitOMPExecutableDirective(S);
}
+void StmtProfiler::VisitOMPTargetDataDirective(const OMPTargetDataDirective *S) {
+ VisitOMPExecutableDirective(S);
+}
+
void StmtProfiler::VisitOMPTeamsDirective(const OMPTeamsDirective *S) {
VisitOMPExecutableDirective(S);
}
Index: lib/AST/StmtPrinter.cpp
===================================================================
--- lib/AST/StmtPrinter.cpp
+++ lib/AST/StmtPrinter.cpp
@@ -946,6 +946,11 @@
PrintOMPExecutableDirective(Node);
}
+void StmtPrinter::VisitOMPTargetDataDirective(OMPTargetDataDirective *Node) {
+ Indent() << "#pragma omp target data ";
+ PrintOMPExecutableDirective(Node);
+}
+
void StmtPrinter::VisitOMPTeamsDirective(OMPTeamsDirective *Node) {
Indent() << "#pragma omp teams ";
PrintOMPExecutableDirective(Node);
Index: lib/AST/Stmt.cpp
===================================================================
--- lib/AST/Stmt.cpp
+++ lib/AST/Stmt.cpp
@@ -2225,6 +2225,30 @@
return new (Mem) OMPTargetDirective(NumClauses);
}
+OMPTargetDataDirective *OMPTargetDataDirective::Create(
+ const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+ ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt) {
+ void *Mem =
+ C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetDataDirective),
+ llvm::alignOf<OMPClause *>()) +
+ sizeof(OMPClause *) * Clauses.size() + sizeof(Stmt *));
+ OMPTargetDataDirective *Dir =
+ new (Mem) OMPTargetDataDirective(StartLoc, EndLoc, Clauses.size());
+ Dir->setClauses(Clauses);
+ Dir->setAssociatedStmt(AssociatedStmt);
+ return Dir;
+}
+
+OMPTargetDataDirective *OMPTargetDataDirective::CreateEmpty(const ASTContext &C,
+ unsigned N,
+ EmptyShell) {
+ void *Mem =
+ C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetDataDirective),
+ llvm::alignOf<OMPClause *>()) +
+ sizeof(OMPClause *) * N + sizeof(Stmt *));
+ return new (Mem) OMPTargetDataDirective(N);
+}
+
OMPTeamsDirective *OMPTeamsDirective::Create(const ASTContext &C,
SourceLocation StartLoc,
SourceLocation EndLoc,
Index: include/clang/Serialization/ASTBitCodes.h
===================================================================
--- include/clang/Serialization/ASTBitCodes.h
+++ include/clang/Serialization/ASTBitCodes.h
@@ -1397,6 +1397,7 @@
STMT_OMP_ORDERED_DIRECTIVE,
STMT_OMP_ATOMIC_DIRECTIVE,
STMT_OMP_TARGET_DIRECTIVE,
+ STMT_OMP_TARGET_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
@@ -7855,6 +7855,11 @@
StmtResult ActOnOpenMPTargetDirective(ArrayRef<OMPClause *> Clauses,
Stmt *AStmt, SourceLocation StartLoc,
SourceLocation EndLoc);
+ /// \brief Called on well-formed '\#pragma omp target data' after parsing of
+ /// the associated statement.
+ StmtResult ActOnOpenMPTargetDataDirective(ArrayRef<OMPClause *> Clauses,
+ Stmt *AStmt, SourceLocation StartLoc,
+ SourceLocation EndLoc);
/// \brief Called on well-formed '\#pragma omp teams' after parsing of the
/// associated statement.
StmtResult ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
Index: include/clang/Basic/StmtNodes.td
===================================================================
--- include/clang/Basic/StmtNodes.td
+++ include/clang/Basic/StmtNodes.td
@@ -204,6 +204,7 @@
def OMPOrderedDirective : DStmt<OMPExecutableDirective>;
def OMPAtomicDirective : DStmt<OMPExecutableDirective>;
def OMPTargetDirective : DStmt<OMPExecutableDirective>;
+def OMPTargetDataDirective : 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
@@ -57,6 +57,9 @@
#ifndef OPENMP_TARGET_CLAUSE
# define OPENMP_TARGET_CLAUSE(Name)
#endif
+#ifndef OPENMP_TARGET_DATA_CLAUSE
+# define OPENMP_TARGET_DATA_CLAUSE(Name)
+#endif
#ifndef OPENMP_TEAMS_CLAUSE
# define OPENMP_TEAMS_CLAUSE(Name)
#endif
@@ -94,6 +97,7 @@
OPENMP_DIRECTIVE(target)
OPENMP_DIRECTIVE(teams)
OPENMP_DIRECTIVE(cancel)
+OPENMP_DIRECTIVE_EXT(target_data, "target data")
OPENMP_DIRECTIVE_EXT(parallel_for, "parallel for")
OPENMP_DIRECTIVE_EXT(parallel_for_simd, "parallel for simd")
OPENMP_DIRECTIVE_EXT(parallel_sections, "parallel sections")
@@ -272,6 +276,10 @@
// TODO More clauses for 'target' directive.
OPENMP_TARGET_CLAUSE(if)
+// Clauses allowed for OpenMP directive 'target data'.
+// TODO More clauses for 'target data' directive.
+OPENMP_TARGET_DATA_CLAUSE(if)
+
// Clauses allowed for OpenMP directive 'teams'.
// TODO More clauses for 'teams' directive.
OPENMP_TEAMS_CLAUSE(default)
@@ -296,6 +304,7 @@
#undef OPENMP_TASK_CLAUSE
#undef OPENMP_ATOMIC_CLAUSE
#undef OPENMP_TARGET_CLAUSE
+#undef OPENMP_TARGET_DATA_CLAUSE
#undef OPENMP_TEAMS_CLAUSE
#undef OPENMP_SIMD_CLAUSE
#undef OPENMP_FOR_CLAUSE
Index: include/clang/AST/StmtOpenMP.h
===================================================================
--- include/clang/AST/StmtOpenMP.h
+++ include/clang/AST/StmtOpenMP.h
@@ -1798,6 +1798,64 @@
}
};
+/// \brief This represents '#pragma omp target data' directive.
+///
+/// \code
+/// #pragma omp target data device(0) if(a) map(b[:])
+/// \endcode
+/// In this example directive '#pragma omp target data' has clauses 'device'
+/// with the value '0', 'if' with condition 'a' and 'map' with array
+/// section 'b[:]'.
+///
+class OMPTargetDataDirective : 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 N The number of clauses.
+ ///
+ OMPTargetDataDirective(SourceLocation StartLoc, SourceLocation EndLoc,
+ unsigned NumClauses)
+ : OMPExecutableDirective(this, OMPTargetDataDirectiveClass,
+ OMPD_target_data, StartLoc, EndLoc, NumClauses,
+ 1) {}
+
+ /// \brief Build an empty directive.
+ ///
+ /// \param N Number of clauses.
+ ///
+ explicit OMPTargetDataDirective(unsigned NumClauses)
+ : OMPExecutableDirective(this, OMPTargetDataDirectiveClass,
+ OMPD_target_data, SourceLocation(),
+ SourceLocation(), NumClauses, 1) {}
+
+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 OMPTargetDataDirective *
+ Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+ ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt);
+
+ /// \brief Creates an empty directive with the place for \a N clauses.
+ ///
+ /// \param C AST context.
+ /// \param N The number of clauses.
+ ///
+ static OMPTargetDataDirective *CreateEmpty(const ASTContext &C, unsigned N,
+ EmptyShell);
+
+ static bool classof(const Stmt *T) {
+ return T->getStmtClass() == OMPTargetDataDirectiveClass;
+ }
+};
+
/// \brief This represents '#pragma omp teams' directive.
///
/// \code
Index: include/clang/AST/RecursiveASTVisitor.h
===================================================================
--- include/clang/AST/RecursiveASTVisitor.h
+++ include/clang/AST/RecursiveASTVisitor.h
@@ -2437,6 +2437,9 @@
DEF_TRAVERSE_STMT(OMPTargetDirective,
{ TRY_TO(TraverseOMPExecutableDirective(S)); })
+DEF_TRAVERSE_STMT(OMPTargetDataDirective,
+ { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
DEF_TRAVERSE_STMT(OMPTeamsDirective,
{ TRY_TO(TraverseOMPExecutableDirective(S)); })
Index: include/clang/AST/DataRecursiveASTVisitor.h
===================================================================
--- include/clang/AST/DataRecursiveASTVisitor.h
+++ include/clang/AST/DataRecursiveASTVisitor.h
@@ -2405,6 +2405,9 @@
DEF_TRAVERSE_STMT(OMPTargetDirective,
{ TRY_TO(TraverseOMPExecutableDirective(S)); })
+DEF_TRAVERSE_STMT(OMPTargetDataDirective,
+ { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
DEF_TRAVERSE_STMT(OMPTeamsDirective,
{ TRY_TO(TraverseOMPExecutableDirective(S)); })
Index: include/clang-c/Index.h
===================================================================
--- include/clang-c/Index.h
+++ include/clang-c/Index.h
@@ -2227,17 +2227,21 @@
/** \brief OpenMP taskgroup directive.
*/
- CXCursor_OMPTaskgroupDirective = 254,
+ CXCursor_OMPTaskgroupDirective = 254,
/** \brief OpenMP cancellation point directive.
*/
- CXCursor_OMPCancellationPointDirective = 255,
+ CXCursor_OMPCancellationPointDirective = 255,
/** \brief OpenMP cancel directive.
*/
- CXCursor_OMPCancelDirective = 256,
+ CXCursor_OMPCancelDirective = 256,
- CXCursor_LastStmt = CXCursor_OMPCancelDirective,
+ /** \brief OpenMP target data directive.
+ */
+ CXCursor_OMPTargetDataDirective = 257,
+
+ CXCursor_LastStmt = CXCursor_OMPTargetDataDirective,
/**
* \brief Cursor that represents the translation unit itself.
_______________________________________________
cfe-commits mailing list
cfe-commits@cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
