maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[OpenMP] Parsing + sema for "target enter data" directive. 

Patch by Arpith Jacob. Thanks!

llvm-svn: 258165
This commit is contained in:
Samuel Antao 2016-01-19 19:15:56 +00:00
parent 251cf1336a
commit df67fc468e
28 changed files with 666 additions and 10 deletions

6
clang/include/clang-c/Index.h
View File

@ -2274,7 +2274,11 @@ enum CXCursorKind {
*/
CXCursor_OMPDistributeDirective = 260,
CXCursor_LastStmt = CXCursor_OMPDistributeDirective,
/** \brief OpenMP target enter data directive.
*/
CXCursor_OMPTargetEnterDataDirective = 261,
CXCursor_LastStmt = CXCursor_OMPTargetEnterDataDirective,
/**
* \brief Cursor that represents the translation unit itself.

3
clang/include/clang/AST/RecursiveASTVisitor.h
View File

@ -2433,6 +2433,9 @@ DEF_TRAVERSE_STMT(OMPTargetDirective,
DEF_TRAVERSE_STMT(OMPTargetDataDirective,
{ TRY_TO(TraverseOMPExecutableDirective(S)); })
DEF_TRAVERSE_STMT(OMPTargetEnterDataDirective,
{ TRY_TO(TraverseOMPExecutableDirective(S)); })
DEF_TRAVERSE_STMT(OMPTeamsDirective,
{ TRY_TO(TraverseOMPExecutableDirective(S)); })

60
clang/include/clang/AST/StmtOpenMP.h
View File

@ -2038,6 +2038,66 @@ public:
}
};
/// \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 teams' directive.
///
/// \code

4
clang/include/clang/Basic/DiagnosticSemaKinds.td
View File

@ -7937,6 +7937,10 @@ def err_omp_map_shared_storage : Error<
"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<

11
clang/include/clang/Basic/OpenMPKinds.def
View File

@ -60,6 +60,9 @@
#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_TEAMS_CLAUSE
# define OPENMP_TEAMS_CLAUSE(Name)
#endif
@ -128,6 +131,7 @@ OPENMP_DIRECTIVE(target)
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(parallel_for, "parallel for")
OPENMP_DIRECTIVE_EXT(parallel_for_simd, "parallel for simd")
OPENMP_DIRECTIVE_EXT(parallel_sections, "parallel sections")
@ -354,6 +358,12 @@ OPENMP_TARGET_DATA_CLAUSE(if)
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 'teams'.
// TODO More clauses for 'teams' directive.
OPENMP_TEAMS_CLAUSE(default)
@ -452,6 +462,7 @@ OPENMP_DIST_SCHEDULE_KIND(static)
#undef OPENMP_ATOMIC_CLAUSE
#undef OPENMP_TARGET_CLAUSE
#undef OPENMP_TARGET_DATA_CLAUSE
#undef OPENMP_TARGET_ENTER_DATA_CLAUSE
#undef OPENMP_TEAMS_CLAUSE
#undef OPENMP_SIMD_CLAUSE
#undef OPENMP_FOR_CLAUSE

1
clang/include/clang/Basic/StmtNodes.td
View File

@ -216,6 +216,7 @@ def OMPOrderedDirective : DStmt<OMPExecutableDirective>;
def OMPAtomicDirective : DStmt<OMPExecutableDirective>;
def OMPTargetDirective : DStmt<OMPExecutableDirective>;
def OMPTargetDataDirective : DStmt<OMPExecutableDirective>;
def OMPTargetEnterDataDirective : DStmt<OMPExecutableDirective>;
def OMPTeamsDirective : DStmt<OMPExecutableDirective>;
def OMPCancellationPointDirective : DStmt<OMPExecutableDirective>;
def OMPCancelDirective : DStmt<OMPExecutableDirective>;

5
clang/include/clang/Sema/Sema.h
View File

@ -7950,6 +7950,11 @@ public:
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 teams' after parsing of the
/// associated statement.
StmtResult ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,

1
clang/include/clang/Serialization/ASTBitCodes.h
View File

@ -1445,6 +1445,7 @@ namespace clang {
STMT_OMP_ATOMIC_DIRECTIVE,
STMT_OMP_TARGET_DIRECTIVE,
STMT_OMP_TARGET_DATA_DIRECTIVE,
STMT_OMP_TARGET_ENTER_DATA_DIRECTIVE,
STMT_OMP_TEAMS_DIRECTIVE,
STMT_OMP_TASKGROUP_DIRECTIVE,
STMT_OMP_CANCELLATION_POINT_DIRECTIVE,

21
clang/lib/AST/StmtOpenMP.cpp
View File

@ -716,6 +716,27 @@ OMPTargetDataDirective *OMPTargetDataDirective::CreateEmpty(const ASTContext &C,
return new (Mem) OMPTargetDataDirective(N);
}
OMPTargetEnterDataDirective *OMPTargetEnterDataDirective::Create(
const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
ArrayRef<OMPClause *> Clauses) {
void *Mem = C.Allocate(llvm::alignTo(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::alignTo(sizeof(OMPTargetEnterDataDirective),
llvm::alignOf<OMPClause *>()) +
sizeof(OMPClause *) * N);
return new (Mem) OMPTargetEnterDataDirective(N);
}
OMPTeamsDirective *OMPTeamsDirective::Create(const ASTContext &C,
SourceLocation StartLoc,
SourceLocation EndLoc,

6
clang/lib/AST/StmtPrinter.cpp
View File

@ -1061,6 +1061,12 @@ void StmtPrinter::VisitOMPTargetDataDirective(OMPTargetDataDirective *Node) {
PrintOMPExecutableDirective(Node);
}
void StmtPrinter::VisitOMPTargetEnterDataDirective(
OMPTargetEnterDataDirective *Node) {
Indent() << "#pragma omp target enter data ";
PrintOMPExecutableDirective(Node);
}
void StmtPrinter::VisitOMPTeamsDirective(OMPTeamsDirective *Node) {
Indent() << "#pragma omp teams ";
PrintOMPExecutableDirective(Node);

5
clang/lib/AST/StmtProfile.cpp
View File

@ -584,6 +584,11 @@ void StmtProfiler::VisitOMPTargetDataDirective(const OMPTargetDataDirective *S)
VisitOMPExecutableDirective(S);
}
void StmtProfiler::VisitOMPTargetEnterDataDirective(
const OMPTargetEnterDataDirective *S) {
VisitOMPExecutableDirective(S);
}
void StmtProfiler::VisitOMPTeamsDirective(const OMPTeamsDirective *S) {
VisitOMPExecutableDirective(S);
}

10
clang/lib/Basic/OpenMPKinds.cpp
View File

@ -408,6 +408,16 @@ bool clang::isAllowedClauseForDirective(OpenMPDirectiveKind DKind,
#define OPENMP_TARGET_DATA_CLAUSE(Name) \
case OMPC_##Name: \
return true;
#include "clang/Basic/OpenMPKinds.def"
default:
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;

3
clang/lib/CodeGen/CGStmt.cpp
View File

@ -256,6 +256,9 @@ void CodeGenFunction::EmitStmt(const Stmt *S) {
case Stmt::OMPTargetDataDirectiveClass:
EmitOMPTargetDataDirective(cast<OMPTargetDataDirective>(*S));
break;
case Stmt::OMPTargetEnterDataDirectiveClass:
EmitOMPTargetEnterDataDirective(cast<OMPTargetEnterDataDirective>(*S));
break;
case Stmt::OMPTaskLoopDirectiveClass:
EmitOMPTaskLoopDirective(cast<OMPTaskLoopDirective>(*S));
break;

5
clang/lib/CodeGen/CGStmtOpenMP.cpp
View File

@ -2661,6 +2661,11 @@ void CodeGenFunction::EmitOMPTargetDataDirective(
[&CS](CodeGenFunction &CGF) { CGF.EmitStmt(CS->getCapturedStmt()); });
}
void CodeGenFunction::EmitOMPTargetEnterDataDirective(
const OMPTargetEnterDataDirective &S) {
// TODO: codegen for target enter data.
}
void CodeGenFunction::EmitOMPTaskLoopDirective(const OMPTaskLoopDirective &S) {
// emit the code inside the construct for now
auto CS = cast<CapturedStmt>(S.getAssociatedStmt());

1
clang/lib/CodeGen/CodeGenFunction.h
View File

@ -2336,6 +2336,7 @@ public:
void EmitOMPAtomicDirective(const OMPAtomicDirective &S);
void EmitOMPTargetDirective(const OMPTargetDirective &S);
void EmitOMPTargetDataDirective(const OMPTargetDataDirective &S);
void EmitOMPTargetEnterDataDirective(const OMPTargetEnterDataDirective &S);
void EmitOMPTeamsDirective(const OMPTeamsDirective &S);
void
EmitOMPCancellationPointDirective(const OMPCancellationPointDirective &S);

19
clang/lib/Parse/ParseOpenMP.cpp
View File

@ -34,6 +34,10 @@ static OpenMPDirectiveKind ParseOpenMPDirectiveKind(Parser &P) {
{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_for, OMPD_simd, OMPD_for_simd},
{OMPD_parallel, OMPD_for, OMPD_parallel_for},
{OMPD_parallel_for, OMPD_simd, OMPD_parallel_for_simd},
@ -49,8 +53,9 @@ static OpenMPDirectiveKind ParseOpenMPDirectiveKind(Parser &P) {
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"));
} else {
TokenMatched = DKind == F[i][0] && DKind != OMPD_unknown;
}
@ -67,7 +72,10 @@ static OpenMPDirectiveKind ParseOpenMPDirectiveKind(Parser &P) {
TokenMatched =
((i == 0) &&
!P.getPreprocessor().getSpelling(Tok).compare("point")) ||
((i == 1) && !P.getPreprocessor().getSpelling(Tok).compare("data"));
((i == 1 || i == 3) &&
!P.getPreprocessor().getSpelling(Tok).compare("data")) ||
((i == 2) &&
!P.getPreprocessor().getSpelling(Tok).compare("enter"));
} else {
TokenMatched = SDKind == F[i][1] && SDKind != OMPD_unknown;
}
@ -138,6 +146,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirective() {
case OMPD_cancellation_point:
case OMPD_cancel:
case OMPD_target_data:
case OMPD_target_enter_data:
case OMPD_taskloop:
case OMPD_taskloop_simd:
case OMPD_distribute:
@ -162,8 +171,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirective() {
/// 'barrier' | 'taskwait' | 'flush' | 'ordered' | 'atomic' |
/// 'for simd' | 'parallel for simd' | 'target' | 'target data' |
/// 'taskgroup' | 'teams' | 'taskloop' | 'taskloop simd' {clause} |
/// 'distribute'
/// annot_pragma_openmp_end
/// 'distribute' | 'target enter data' | annot_pragma_openmp_end
///
StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
AllowedContsructsKind Allowed) {
@ -217,6 +225,7 @@ StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
case OMPD_taskwait:
case OMPD_cancellation_point:
case OMPD_cancel:
case OMPD_target_enter_data:
if (Allowed == ACK_StatementsOpenMPNonStandalone) {
Diag(Tok, diag::err_omp_immediate_directive)
<< getOpenMPDirectiveName(DKind) << 0;

96
clang/lib/Sema/SemaOpenMP.cpp
View File

@ -1610,6 +1610,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
case OMPD_cancellation_point:
case OMPD_cancel:
case OMPD_flush:
case OMPD_target_enter_data:
llvm_unreachable("OpenMP Directive is not allowed");
case OMPD_unknown:
llvm_unreachable("Unknown OpenMP directive");
@ -1724,13 +1725,15 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | parallel | ordered | + |
// | parallel | atomic | * |
// | parallel | target | * |
// | parallel | target enter | * |
// | | data | |
// | parallel | teams | + |
// | parallel | cancellation | |
// | | point | ! |
// | parallel | cancel | ! |
// | parallel | taskloop | * |
// | parallel | taskloop simd | * |
// | parallel | distribute | |
// | parallel | distribute | |
// +------------------+-----------------+------------------------------------+
// | for | parallel | * |
// | for | for | + |
@ -1753,6 +1756,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | for | ordered | * (if construct is ordered) |
// | for | atomic | * |
// | for | target | * |
// | for | target enter | * |
// | | data | |
// | for | teams | + |
// | for | cancellation | |
// | | point | ! |
@ -1782,6 +1787,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | master | ordered | + |
// | master | atomic | * |
// | master | target | * |
// | master | target enter | * |
// | | data | |
// | master | teams | + |
// | master | cancellation | |
// | | point | |
@ -1810,6 +1817,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | critical | ordered | + |
// | critical | atomic | * |
// | critical | target | * |
// | critical | target enter | * |
// | | data | |
// | critical | teams | + |
// | critical | cancellation | |
// | | point | |
@ -1839,6 +1848,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | simd | ordered | + (with simd clause) |
// | simd | atomic | |
// | simd | target | |
// | simd | target enter | |
// | | data | |
// | simd | teams | |
// | simd | cancellation | |
// | | point | |
@ -1868,6 +1879,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | for simd | ordered | + (with simd clause) |
// | for simd | atomic | |
// | for simd | target | |
// | for simd | target enter | |
// | | data | |
// | for simd | teams | |
// | for simd | cancellation | |
// | | point | |
@ -1897,6 +1910,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | parallel for simd| ordered | + (with simd clause) |
// | parallel for simd| atomic | |
// | parallel for simd| target | |
// | parallel for simd| target enter | |
// | | data | |
// | parallel for simd| teams | |
// | parallel for simd| cancellation | |
// | | point | |
@ -1926,6 +1941,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | sections | ordered | + |
// | sections | atomic | * |
// | sections | target | * |
// | sections | target enter | * |
// | | data | |
// | sections | teams | + |
// | sections | cancellation | |
// | | point | ! |
@ -1955,6 +1972,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | section | ordered | + |
// | section | atomic | * |
// | section | target | * |
// | section | target enter | * |
// | | data | |
// | section | teams | + |
// | section | cancellation | |
// | | point | ! |
@ -1984,6 +2003,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | single | ordered | + |
// | single | atomic | * |
// | single | target | * |
// | single | target enter | * |
// | | data | |
// | single | teams | + |
// | single | cancellation | |
// | | point | |
@ -2013,6 +2034,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | parallel for | ordered | * (if construct is ordered) |
// | parallel for | atomic | * |
// | parallel for | target | * |
// | parallel for | target enter | * |
// | | data | |
// | parallel for | teams | + |
// | parallel for | cancellation | |
// | | point | ! |
@ -2042,13 +2065,15 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | parallel sections| ordered | + |
// | parallel sections| atomic | * |
// | parallel sections| target | * |
// | parallel sections| target enter | * |
// | | data | |
// | parallel sections| teams | + |
// | parallel sections| cancellation | |
// | | point | ! |
// | parallel sections| cancel | ! |
// | parallel sections| taskloop | * |
// | parallel sections| taskloop simd | * |
// | parallel sections| distribute | |
// | parallel sections| distribute | |
// +------------------+-----------------+------------------------------------+
// | task | parallel | * |
// | task | for | + |
@ -2071,6 +2096,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | task | ordered | + |
// | task | atomic | * |
// | task | target | * |
// | task | target enter | * |
// | | data | |
// | task | teams | + |
// | task | cancellation | |
// | | point | ! |
@ -2100,6 +2127,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | ordered | ordered | + |
// | ordered | atomic | * |
// | ordered | target | * |
// | ordered | target enter | * |
// | | data | |
// | ordered | teams | + |
// | ordered | cancellation | |
// | | point | |
@ -2129,13 +2158,15 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | atomic | ordered | |
// | atomic | atomic | |
// | atomic | target | |
// | atomic | target enter | |
// | | data | |
// | atomic | teams | |
// | atomic | cancellation | |
// | | point | |
// | atomic | cancel | |
// | atomic | taskloop | |
// | atomic | taskloop simd | |
// | atomic | distribute | |
// | atomic | distribute | |
// +------------------+-----------------+------------------------------------+
// | target | parallel | * |
// | target | for | * |
@ -2158,6 +2189,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | target | ordered | * |
// | target | atomic | * |
// | target | target | * |
// | target | target enter | * |
// | | data | |
// | target | teams | * |
// | target | cancellation | |
// | | point | |
@ -2187,6 +2220,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | teams | ordered | + |
// | teams | atomic | + |
// | teams | target | + |
// | teams | target enter | + |
// | | data | |
// | teams | teams | + |
// | teams | cancellation | |
// | | point | |
@ -2216,6 +2251,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | taskloop | ordered | + |
// | taskloop | atomic | * |
// | taskloop | target | * |
// | taskloop | target enter | * |
// | | data | |
// | taskloop | teams | + |
// | taskloop | cancellation | |
// | | point | |
@ -2244,6 +2281,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | taskloop simd | ordered | + (with simd clause) |
// | taskloop simd | atomic | |
// | taskloop simd | target | |
// | taskloop simd | target enter | |
// | | data | |
// | taskloop simd | teams | |
// | taskloop simd | cancellation | |
// | | point | |
@ -2273,6 +2312,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
// | distribute | ordered | + |
// | distribute | atomic | * |
// | distribute | target | |
// | distribute | target enter | |
// | | data | |
// | distribute | teams | |
// | distribute | cancellation | + |
// | | point | |
@ -2697,6 +2738,11 @@ StmtResult Sema::ActOnOpenMPExecutableDirective(
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_taskloop:
Res = ActOnOpenMPTaskLoopDirective(ClausesWithImplicit, AStmt, StartLoc,
EndLoc, VarsWithInheritedDSA);
@ -5433,6 +5479,18 @@ StmtResult Sema::ActOnOpenMPTargetDirective(ArrayRef<OMPClause *> Clauses,
return OMPTargetDirective::Create(Context, StartLoc, EndLoc, Clauses, AStmt);
}
/// \brief Check for existence of a map clause in the list of clauses.
static bool HasMapClause(ArrayRef<OMPClause *> Clauses) {
for (ArrayRef<OMPClause *>::iterator I = Clauses.begin(), E = Clauses.end();
I != E; ++I) {
if (*I != nullptr && (*I)->getClauseKind() == OMPC_map) {
return true;
}
}
return false;
}
StmtResult Sema::ActOnOpenMPTargetDataDirective(ArrayRef<OMPClause *> Clauses,
Stmt *AStmt,
SourceLocation StartLoc,
@ -5448,6 +5506,22 @@ StmtResult Sema::ActOnOpenMPTargetDataDirective(ArrayRef<OMPClause *> Clauses,
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::ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
Stmt *AStmt, SourceLocation StartLoc,
SourceLocation EndLoc) {
@ -8449,6 +8523,22 @@ OMPClause *Sema::ActOnOpenMPMapClause(
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)
<<
// TODO: Need to determine if map type is implicitly determined
0 << 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);

11
clang/lib/Sema/TreeTransform.h
View File

@ -7390,6 +7390,17 @@ StmtResult TreeTransform<Derived>::TransformOMPTargetDataDirective(
return Res;
}
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>::TransformOMPTeamsDirective(OMPTeamsDirective *D) {

12
clang/lib/Serialization/ASTReaderStmt.cpp
View File

@ -2455,6 +2455,13 @@ void ASTStmtReader::VisitOMPTargetDataDirective(OMPTargetDataDirective *D) {
VisitOMPExecutableDirective(D);
}
void ASTStmtReader::VisitOMPTargetEnterDataDirective(
OMPTargetEnterDataDirective *D) {
VisitStmt(D);
++Idx;
VisitOMPExecutableDirective(D);
}
void ASTStmtReader::VisitOMPTeamsDirective(OMPTeamsDirective *D) {
VisitStmt(D);
// The NumClauses field was read in ReadStmtFromStream.
@ -3098,6 +3105,11 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
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_TEAMS_DIRECTIVE:
S = OMPTeamsDirective::CreateEmpty(
Context, Record[ASTStmtReader::NumStmtFields], Empty);

8
clang/lib/Serialization/ASTWriterStmt.cpp
View File

@ -2214,6 +2214,14 @@ void ASTStmtWriter::VisitOMPTargetDataDirective(OMPTargetDataDirective *D) {
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::VisitOMPTaskyieldDirective(OMPTaskyieldDirective *D) {
VisitStmt(D);
VisitOMPExecutableDirective(D);

1
clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
View File

@ -830,6 +830,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
case Stmt::OMPAtomicDirectiveClass:
case Stmt::OMPTargetDirectiveClass:
case Stmt::OMPTargetDataDirectiveClass:
case Stmt::OMPTargetEnterDataDirectiveClass:
case Stmt::OMPTeamsDirectiveClass:
case Stmt::OMPCancellationPointDirectiveClass:
case Stmt::OMPCancelDirectiveClass:

195
clang/test/OpenMP/nesting_of_regions.cpp
View File

@ -97,6 +97,11 @@ void foo() {
}
#pragma omp parallel
{
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp parallel
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'parallel' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -236,6 +241,11 @@ void foo() {
}
#pragma omp simd
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside a simd region}}
++a;
}
#pragma omp simd
for (int i = 0; i < 10; ++i) {
#pragma omp teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
++a;
}
@ -398,6 +408,11 @@ void foo() {
}
#pragma omp for
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp for
for (int i = 0; i < 10; ++i) {
#pragma omp teams // expected-error {{region cannot be closely nested inside 'for' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -537,6 +552,11 @@ void foo() {
}
#pragma omp for simd
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside a simd region}}
++a;
}
#pragma omp for simd
for (int i = 0; i < 10; ++i) {
#pragma omp teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
++a;
}
@ -706,6 +726,10 @@ void foo() {
}
#pragma omp sections
{
#pragma omp target enter data map(to: a)
}
#pragma omp sections
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'sections' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -910,6 +934,14 @@ void foo() {
}
#pragma omp sections
{
#pragma omp section
{
#pragma omp target enter data map(to: a)
++a;
}
}
#pragma omp sections
{
#pragma omp section
#pragma omp teams // expected-error {{region cannot be closely nested inside 'section' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
@ -1065,6 +1097,11 @@ void foo() {
}
#pragma omp single
{
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp single
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'single' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -1217,6 +1254,11 @@ void foo() {
}
#pragma omp master
{
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp master
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'master' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -1383,6 +1425,11 @@ void foo() {
}
#pragma omp critical
{
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp critical
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'critical' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -1550,6 +1597,11 @@ void foo() {
}
#pragma omp parallel for
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp parallel for
for (int i = 0; i < 10; ++i) {
#pragma omp teams // expected-error {{region cannot be closely nested inside 'parallel for' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -1717,6 +1769,11 @@ void foo() {
}
#pragma omp parallel for simd
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside a simd region}}
++a;
}
#pragma omp parallel for simd
for (int i = 0; i < 10; ++i) {
#pragma omp teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
++a;
}
@ -1875,6 +1932,10 @@ void foo() {
}
#pragma omp parallel sections
{
#pragma omp target enter data map(to: a)
}
#pragma omp parallel sections
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'parallel sections' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -1979,6 +2040,11 @@ void foo() {
}
#pragma omp task
{
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp task
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'task' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -2141,6 +2207,11 @@ void foo() {
}
#pragma omp ordered
{
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp ordered
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'ordered' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -2322,6 +2393,13 @@ void foo() {
// 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 enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
++a;
}
#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 teams // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
++a;
}
@ -2458,6 +2536,10 @@ void foo() {
for (int i = 0; i < 10; ++i)
;
}
#pragma omp target
{
#pragma omp target enter data map(to: a)
}
// TEAMS DIRECTIVE
#pragma omp target
@ -2573,6 +2655,12 @@ void foo() {
++a;
}
#pragma omp target
#pragma omp teams
{
#pragma omp target enter data map(to: a) // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target enter data' directive into a parallel region?}}
++a;
}
#pragma omp target
#pragma omp teams
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
@ -2740,6 +2828,11 @@ void foo() {
}
#pragma omp taskloop
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp taskloop
for (int i = 0; i < 10; ++i) {
#pragma omp teams // expected-error {{region cannot be closely nested inside 'taskloop' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -2932,6 +3025,13 @@ void foo() {
}
#pragma omp target
#pragma omp teams
#pragma omp distribute
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp target
#pragma omp teams
#pragma omp distribute
for (int i = 0; i < 10; ++i) {
#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?}}
@ -3033,6 +3133,11 @@ void foo() {
}
#pragma omp parallel
{
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp parallel
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'parallel' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -3165,6 +3270,11 @@ void foo() {
}
#pragma omp simd
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside a simd region}}
++a;
}
#pragma omp simd
for (int i = 0; i < 10; ++i) {
#pragma omp teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
++a;
}
@ -3317,6 +3427,11 @@ void foo() {
}
#pragma omp for
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp for
for (int i = 0; i < 10; ++i) {
#pragma omp teams // expected-error {{region cannot be closely nested inside 'for' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -3449,6 +3564,11 @@ void foo() {
}
#pragma omp for simd
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside a simd region}}
++a;
}
#pragma omp for simd
for (int i = 0; i < 10; ++i) {
#pragma omp teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
++a;
}
@ -3593,6 +3713,10 @@ void foo() {
}
#pragma omp sections
{
#pragma omp target enter data map(to: a)
}
#pragma omp sections
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'sections' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -3803,6 +3927,14 @@ void foo() {
{
#pragma omp section
{
#pragma omp target enter data map(to: a)
++a;
}
}
#pragma omp sections
{
#pragma omp section
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'section' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -3950,6 +4082,11 @@ void foo() {
}
#pragma omp single
{
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp single
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'single' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -4102,6 +4239,11 @@ void foo() {
}
#pragma omp master
{
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp master
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'master' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -4273,6 +4415,11 @@ void foo() {
}
#pragma omp critical
{
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp critical
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'critical' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -4440,6 +4587,11 @@ void foo() {
}
#pragma omp parallel for
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp parallel for
for (int i = 0; i < 10; ++i) {
#pragma omp teams // expected-error {{region cannot be closely nested inside 'parallel for' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -4607,6 +4759,11 @@ void foo() {
}
#pragma omp parallel for simd
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside a simd region}}
++a;
}
#pragma omp parallel for simd
for (int i = 0; i < 10; ++i) {
#pragma omp teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
++a;
}
@ -4761,6 +4918,10 @@ void foo() {
}
#pragma omp parallel sections
{
#pragma omp target enter data map(to: a)
}
#pragma omp parallel sections
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'parallel sections' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -4864,6 +5025,11 @@ void foo() {
}
#pragma omp task
{
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp task
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'task' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -5045,6 +5211,13 @@ void foo() {
// 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 enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
++a;
}
#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 teams // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
++a;
}
@ -5160,6 +5333,10 @@ void foo() {
}
#pragma omp target
{
#pragma omp target enter data map(to: a)
}
#pragma omp target
{
#pragma omp teams
++a;
}
@ -5296,6 +5473,11 @@ void foo() {
++a;
}
#pragma omp target
#pragma omp teams
{
#pragma omp target enter data map(to: a) // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target enter data' directive into a parallel region?}}
}
#pragma omp target
#pragma omp teams
{
#pragma omp teams // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
@ -5463,6 +5645,11 @@ void foo() {
}
#pragma omp taskloop
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a)
++a;
}
#pragma omp taskloop
for (int i = 0; i < 10; ++i) {
#pragma omp teams // expected-error {{region cannot be closely nested inside 'taskloop' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
++a;
}
@ -5662,5 +5849,11 @@ void foo() {
++a;
}
return foo<int>();
#pragma omp target
#pragma omp teams
#pragma omp distribute
for (int i = 0; i < 10; ++i) {
#pragma omp target enter data map(to: a)
++a;
}
}

100
clang/test/OpenMP/target_enter_data_ast_print.cpp Normal file
View File

@ -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

29
clang/test/OpenMP/target_enter_data_device_messages.cpp Normal file
View File

@ -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;
}

35
clang/test/OpenMP/target_enter_data_if_messages.cpp Normal file
View File

@ -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;
}

17
clang/test/OpenMP/target_enter_data_map_messages.c Normal file
View File

@ -0,0 +1,17 @@
// 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 // 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;
}

8
clang/tools/libclang/CIndex.cpp
View File

@ -1953,6 +1953,7 @@ public:
void VisitOMPAtomicDirective(const OMPAtomicDirective *D);
void VisitOMPTargetDirective(const OMPTargetDirective *D);
void VisitOMPTargetDataDirective(const OMPTargetDataDirective *D);
void VisitOMPTargetEnterDataDirective(const OMPTargetEnterDataDirective *D);
void VisitOMPTeamsDirective(const OMPTeamsDirective *D);
void VisitOMPTaskLoopDirective(const OMPTaskLoopDirective *D);
void VisitOMPTaskLoopSimdDirective(const OMPTaskLoopSimdDirective *D);
@ -2633,6 +2634,11 @@ void EnqueueVisitor::VisitOMPTargetDataDirective(const
VisitOMPExecutableDirective(D);
}
void EnqueueVisitor::VisitOMPTargetEnterDataDirective(
const OMPTargetEnterDataDirective *D) {
VisitOMPExecutableDirective(D);
}
void EnqueueVisitor::VisitOMPTeamsDirective(const OMPTeamsDirective *D) {
VisitOMPExecutableDirective(D);
}
@ -4838,6 +4844,8 @@ CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
return cxstring::createRef("OMPTargetDirective");
case CXCursor_OMPTargetDataDirective:
return cxstring::createRef("OMPTargetDataDirective");
case CXCursor_OMPTargetEnterDataDirective:
return cxstring::createRef("OMPTargetEnterDataDirective");
case CXCursor_OMPTeamsDirective:
return cxstring::createRef("OMPTeamsDirective");
case CXCursor_OMPCancellationPointDirective:

3
clang/tools/libclang/CXCursor.cpp
View File

@ -600,6 +600,9 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent,
case Stmt::OMPTargetDataDirectiveClass:
K = CXCursor_OMPTargetDataDirective;
break;
case Stmt::OMPTargetEnterDataDirectiveClass:
K = CXCursor_OMPTargetEnterDataDirective;
break;
case Stmt::OMPTeamsDirectiveClass:
K = CXCursor_OMPTeamsDirective;
break;