From d7724c8ea313e5d64ef35102958213dcd98f53ea Mon Sep 17 00:00:00 2001 From: erichkeane Date: Wed, 2 Apr 2025 08:06:21 -0700 Subject: [PATCH] [OpenACC] allow 'if' clause on 'atomic' construct This was added in OpenACC PR #511 in the 3.4 branch. From an AST/Sema perspective this is pretty trivial as the infrastructure for 'if' already exists, however the atomic construct needed to be taught to take clauses. This patch does that and adds some testing to do so. --- clang/include/clang/AST/StmtOpenACC.h | 32 ++- clang/lib/AST/StmtOpenACC.cpp | 20 +- clang/lib/AST/StmtPrinter.cpp | 1 + clang/lib/AST/StmtProfile.cpp | 2 + clang/lib/Sema/SemaOpenACC.cpp | 3 +- clang/lib/Sema/SemaOpenACCClause.cpp | 2 + clang/lib/Sema/TreeTransform.h | 9 +- clang/lib/Serialization/ASTReaderStmt.cpp | 7 +- clang/lib/Serialization/ASTWriterStmt.cpp | 4 +- .../ast-print-openacc-atomic-construct.cpp | 32 +++ clang/test/ParserOpenACC/parse-clauses.c | 24 ++ clang/test/ParserOpenACC/parse-clauses.cpp | 24 ++ .../test/SemaOpenACC/atomic-construct-ast.cpp | 229 ++++++++++++++++++ clang/test/SemaOpenACC/atomic-construct.cpp | 21 ++ clang/tools/libclang/CIndex.cpp | 2 + 15 files changed, 388 insertions(+), 24 deletions(-) diff --git a/clang/include/clang/AST/StmtOpenACC.h b/clang/include/clang/AST/StmtOpenACC.h index c2c74f5cf195..39c4c8184491 100644 --- a/clang/include/clang/AST/StmtOpenACC.h +++ b/clang/include/clang/AST/StmtOpenACC.h @@ -829,24 +829,42 @@ public: // This class represents the 'atomic' construct, which has an associated // statement, but no clauses. -class OpenACCAtomicConstruct final : public OpenACCAssociatedStmtConstruct { +class OpenACCAtomicConstruct final + : public OpenACCAssociatedStmtConstruct, + private llvm::TrailingObjects { friend class ASTStmtReader; + friend TrailingObjects; OpenACCAtomicKind AtomicKind = OpenACCAtomicKind::None; - OpenACCAtomicConstruct(EmptyShell) + OpenACCAtomicConstruct(unsigned NumClauses) : OpenACCAssociatedStmtConstruct( OpenACCAtomicConstructClass, OpenACCDirectiveKind::Atomic, SourceLocation{}, SourceLocation{}, SourceLocation{}, - /*AssociatedStmt=*/nullptr) {} + /*AssociatedStmt=*/nullptr) { + std::uninitialized_value_construct( + getTrailingObjects(), + getTrailingObjects() + NumClauses); + setClauseList(MutableArrayRef(getTrailingObjects(), + NumClauses)); + } OpenACCAtomicConstruct(SourceLocation Start, SourceLocation DirectiveLoc, OpenACCAtomicKind AtKind, SourceLocation End, + ArrayRef Clauses, Stmt *AssociatedStmt) : OpenACCAssociatedStmtConstruct(OpenACCAtomicConstructClass, OpenACCDirectiveKind::Atomic, Start, DirectiveLoc, End, AssociatedStmt), - AtomicKind(AtKind) {} + AtomicKind(AtKind) { + // Initialize the trailing storage. + std::uninitialized_copy(Clauses.begin(), Clauses.end(), + getTrailingObjects()); + + setClauseList(MutableArrayRef(getTrailingObjects(), + Clauses.size())); + } void setAssociatedStmt(Stmt *S) { OpenACCAssociatedStmtConstruct::setAssociatedStmt(S); @@ -857,10 +875,12 @@ public: return T->getStmtClass() == OpenACCAtomicConstructClass; } - static OpenACCAtomicConstruct *CreateEmpty(const ASTContext &C); + static OpenACCAtomicConstruct *CreateEmpty(const ASTContext &C, + unsigned NumClauses); static OpenACCAtomicConstruct * Create(const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc, - OpenACCAtomicKind AtKind, SourceLocation End, Stmt *AssociatedStmt); + OpenACCAtomicKind AtKind, SourceLocation End, + ArrayRef Clauses, Stmt *AssociatedStmt); OpenACCAtomicKind getAtomicKind() const { return AtomicKind; } const Stmt *getAssociatedStmt() const { diff --git a/clang/lib/AST/StmtOpenACC.cpp b/clang/lib/AST/StmtOpenACC.cpp index 8a86074fe68a..c45eca92dc87 100644 --- a/clang/lib/AST/StmtOpenACC.cpp +++ b/clang/lib/AST/StmtOpenACC.cpp @@ -307,20 +307,26 @@ OpenACCUpdateConstruct::Create(const ASTContext &C, SourceLocation Start, } OpenACCAtomicConstruct * -OpenACCAtomicConstruct::CreateEmpty(const ASTContext &C) { - void *Mem = C.Allocate(sizeof(OpenACCAtomicConstruct)); - auto *Inst = new (Mem) OpenACCAtomicConstruct(EmptyShell{}); +OpenACCAtomicConstruct::CreateEmpty(const ASTContext &C, unsigned NumClauses) { + void *Mem = C.Allocate( + OpenACCAtomicConstruct::totalSizeToAlloc( + NumClauses)); + auto *Inst = new (Mem) OpenACCAtomicConstruct(NumClauses); return Inst; } OpenACCAtomicConstruct *OpenACCAtomicConstruct::Create( const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc, - OpenACCAtomicKind AtKind, SourceLocation End, Stmt *AssociatedStmt) { - void *Mem = C.Allocate(sizeof(OpenACCAtomicConstruct)); - auto *Inst = new (Mem) - OpenACCAtomicConstruct(Start, DirectiveLoc, AtKind, End, AssociatedStmt); + OpenACCAtomicKind AtKind, SourceLocation End, + ArrayRef Clauses, Stmt *AssociatedStmt) { + void *Mem = C.Allocate( + OpenACCAtomicConstruct::totalSizeToAlloc( + Clauses.size())); + auto *Inst = new (Mem) OpenACCAtomicConstruct(Start, DirectiveLoc, AtKind, + End, Clauses, AssociatedStmt); return Inst; } + OpenACCCacheConstruct *OpenACCCacheConstruct::CreateEmpty(const ASTContext &C, unsigned NumVars) { void *Mem = diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp index fe9784a15b76..dbe2432d5c79 100644 --- a/clang/lib/AST/StmtPrinter.cpp +++ b/clang/lib/AST/StmtPrinter.cpp @@ -1258,6 +1258,7 @@ void StmtPrinter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) { if (S->getAtomicKind() != OpenACCAtomicKind::None) OS << " " << S->getAtomicKind(); + PrintOpenACCClauseList(S); OS << '\n'; PrintStmt(S->getAssociatedStmt()); } diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp index 96852858b3b9..83d54da9be7e 100644 --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -2839,6 +2839,8 @@ void StmtProfiler::VisitOpenACCUpdateConstruct( void StmtProfiler::VisitOpenACCAtomicConstruct( const OpenACCAtomicConstruct *S) { VisitStmt(S); + OpenACCClauseProfiler P{*this}; + P.VisitOpenACCClauseList(S->clauses()); } void StmtProfiler::VisitHLSLOutArgExpr(const HLSLOutArgExpr *S) { diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index fea334e9e80f..231e73fdbc9e 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -1832,9 +1832,8 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective( EndLoc, Clauses); } case OpenACCDirectiveKind::Atomic: { - assert(Clauses.empty() && "Atomic doesn't allow clauses"); return OpenACCAtomicConstruct::Create( - getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc, + getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc, Clauses, AssocStmt.isUsable() ? AssocStmt.get() : nullptr); } case OpenACCDirectiveKind::Cache: { diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp index 78e250b75950..a98b6712014c 100644 --- a/clang/lib/Sema/SemaOpenACCClause.cpp +++ b/clang/lib/Sema/SemaOpenACCClause.cpp @@ -55,6 +55,8 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind, case OpenACCDirectiveKind::ParallelLoop: case OpenACCDirectiveKind::SerialLoop: case OpenACCDirectiveKind::KernelsLoop: + // OpenACC 3.4(prerelease) PR #511 adds 'if' to atomic. + case OpenACCDirectiveKind::Atomic: return true; default: return false; diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 916b8e2735cd..524e73242a2d 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -4222,10 +4222,11 @@ public: SourceLocation DirLoc, OpenACCAtomicKind AtKind, SourceLocation EndLoc, + ArrayRef Clauses, StmtResult AssociatedStmt) { return getSema().OpenACC().ActOnEndStmtDirective( OpenACCDirectiveKind::Atomic, BeginLoc, DirLoc, SourceLocation{}, - SourceLocation{}, {}, AtKind, SourceLocation{}, EndLoc, {}, + SourceLocation{}, {}, AtKind, SourceLocation{}, EndLoc, Clauses, AssociatedStmt); } @@ -12744,6 +12745,10 @@ StmtResult TreeTransform::TransformOpenACCAtomicConstruct( OpenACCAtomicConstruct *C) { getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc()); + llvm::SmallVector TransformedClauses = + getDerived().TransformOpenACCClauseList(C->getDirectiveKind(), + C->clauses()); + if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(), C->getBeginLoc(), {})) return StmtError(); @@ -12759,7 +12764,7 @@ StmtResult TreeTransform::TransformOpenACCAtomicConstruct( return getDerived().RebuildOpenACCAtomicConstruct( C->getBeginLoc(), C->getDirectiveLoc(), C->getAtomicKind(), - C->getEndLoc(), AssocStmt); + C->getEndLoc(), TransformedClauses, AssocStmt); } template diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp index 8dceca6ff3db..7c7abcb2d49d 100644 --- a/clang/lib/Serialization/ASTReaderStmt.cpp +++ b/clang/lib/Serialization/ASTReaderStmt.cpp @@ -2940,9 +2940,7 @@ void ASTStmtReader::VisitOpenACCCacheConstruct(OpenACCCacheConstruct *S) { void ASTStmtReader::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) { VisitStmt(S); - S->Kind = Record.readEnum(); - S->Range = Record.readSourceRange(); - S->DirectiveLoc = Record.readSourceLocation(); + VisitOpenACCConstructStmt(S); S->AtomicKind = Record.readEnum(); S->setAssociatedStmt(Record.readSubStmt()); } @@ -4490,7 +4488,8 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) { break; } case STMT_OPENACC_ATOMIC_CONSTRUCT: { - S = OpenACCAtomicConstruct::CreateEmpty(Context); + unsigned NumClauses = Record[ASTStmtReader::NumStmtFields]; + S = OpenACCAtomicConstruct::CreateEmpty(Context, NumClauses); break; } case EXPR_REQUIRES: { diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp index 037eb13b1dc5..0860704368f3 100644 --- a/clang/lib/Serialization/ASTWriterStmt.cpp +++ b/clang/lib/Serialization/ASTWriterStmt.cpp @@ -3014,9 +3014,7 @@ void ASTStmtWriter::VisitOpenACCWaitConstruct(OpenACCWaitConstruct *S) { void ASTStmtWriter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) { VisitStmt(S); - Record.writeEnum(S->Kind); - Record.AddSourceRange(S->Range); - Record.AddSourceLocation(S->DirectiveLoc); + VisitOpenACCConstructStmt(S); Record.writeEnum(S->getAtomicKind()); Record.AddStmt(S->getAssociatedStmt()); diff --git a/clang/test/AST/ast-print-openacc-atomic-construct.cpp b/clang/test/AST/ast-print-openacc-atomic-construct.cpp index 572f2ea4842d..50e988be603d 100644 --- a/clang/test/AST/ast-print-openacc-atomic-construct.cpp +++ b/clang/test/AST/ast-print-openacc-atomic-construct.cpp @@ -31,3 +31,35 @@ void foo(int v, int x) { { x--; v = x; } } + +void foo2(int v, int x) { +// CHECK: #pragma acc atomic read if(v) +// CHECK-NEXT: v = x; +#pragma acc atomic read if (v) + v = x; +// CHECK-NEXT: pragma acc atomic write if(x) +// CHECK-NEXT: v = x + 1; +#pragma acc atomic write if (x) + v = x + 1; +// CHECK-NEXT: pragma acc atomic update if(true) +// CHECK-NEXT: x++; +#pragma acc atomic update if (true) + x++; +// CHECK-NEXT: pragma acc atomic if(false) +// CHECK-NEXT: x--; +#pragma acc atomic if (false) + x--; +// CHECK-NEXT: pragma acc atomic capture if(v < x) +// CHECK-NEXT: v = x++; +#pragma acc atomic capture if (v < x) + v = x++; + +// CHECK-NEXT: #pragma acc atomic capture if(x > v) +// CHECK-NEXT: { +// CHECK-NEXT: x--; +// CHECK-NEXT: v = x; +// CHECK-NEXT: } +#pragma acc atomic capture if (x > v) + { x--; v = x; } + +} diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c index 100dc07a9fd6..e31b7492dab2 100644 --- a/clang/test/ParserOpenACC/parse-clauses.c +++ b/clang/test/ParserOpenACC/parse-clauses.c @@ -1288,3 +1288,27 @@ void BCP1(); #pragma acc routine seq bind("ReductionClauseParsing") #pragma acc routine(BCP1) seq bind(unknown_thing) + +void AtomicIf() { + int i, j; + // expected-error@+1{{expected '('}} +#pragma acc atomic read if + i = j; +#pragma acc atomic read if (0) + i = j; +#pragma acc atomic write if (1) + i = j + 1; + +#pragma acc atomic update if (i) + ++i; +#pragma acc atomic if (j) + ++i; + +#pragma acc atomic capture if (0) + i = j++; +#pragma acc atomic capture if (i) + { + ++j; + i = j; + } +} diff --git a/clang/test/ParserOpenACC/parse-clauses.cpp b/clang/test/ParserOpenACC/parse-clauses.cpp index 6c8401806094..0ba70f5c6a1f 100644 --- a/clang/test/ParserOpenACC/parse-clauses.cpp +++ b/clang/test/ParserOpenACC/parse-clauses.cpp @@ -61,3 +61,27 @@ void use() { void another_func(); // expected-warning@+1{{encoding prefix 'U' on an unevaluated string literal has no effect}} #pragma acc routine(another_func) seq bind(U"32 bits") + +void AtomicIf() { + int i, j; + // expected-error@+1{{expected '('}} +#pragma acc atomic read if + i = j; +#pragma acc atomic read if (true) + i = j; +#pragma acc atomic write if (false) + i = j + 1; + +#pragma acc atomic update if (i) + ++i; +#pragma acc atomic if (j) + ++i; + +#pragma acc atomic capture if (true) + i = j++; +#pragma acc atomic capture if (i) + { + ++j; + i = j; + } +} diff --git a/clang/test/SemaOpenACC/atomic-construct-ast.cpp b/clang/test/SemaOpenACC/atomic-construct-ast.cpp index bfd7f32884f1..157532951a16 100644 --- a/clang/test/SemaOpenACC/atomic-construct-ast.cpp +++ b/clang/test/SemaOpenACC/atomic-construct-ast.cpp @@ -68,6 +68,90 @@ _Pragma("acc atomic read") #pragma acc atomic capture { x--; v = x; } +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic read +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '=' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +#pragma acc atomic read if (v < x) + v = x; + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic write +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '=' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: BinaryOperator{{.*}}'int' '+' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 +_Pragma("acc atomic write if (v < x)") + v = x + 1; + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic update +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '++' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +#pragma acc atomic update if (v < x) + x++; + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '--' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +#pragma acc atomic if (v < x) + x--; + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '=' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '++' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +#pragma acc atomic capture if (v < x) + v = x++; + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '--' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '=' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +#pragma acc atomic capture if (v < x) + { x--; v = x; } } template @@ -124,6 +208,76 @@ void templ_foo(T v, T x) { #pragma acc atomic capture { x--; v = x; } +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic read +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'' '<' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T' +// CHECK-NEXT: BinaryOperator{{.*}} '' '=' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T' +#pragma acc atomic read if (v < x) + v = x; + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic write +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'' '<' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T' +// CHECK-NEXT: BinaryOperator{{.*}} '' '=' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T' +// CHECK-NEXT: BinaryOperator{{.*}}'' '+' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T' +// CHECK-NEXT: DeclRefExpr{{.*}} 'I' 'int' +_Pragma("acc atomic write if (v < x)") + v = x + I; + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic update +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'' '<' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T' +// CHECK-NEXT: UnaryOperator{{.*}} '' postfix '++' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T' +#pragma acc atomic update if (v < x) + x++; + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'' '<' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T' +// CHECK-NEXT: UnaryOperator{{.*}} '' postfix '--' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T' +#pragma acc atomic if (v < x) + x--; + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'' '<' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T' +// CHECK-NEXT: BinaryOperator{{.*}} '' '=' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T' +// CHECK-NEXT: UnaryOperator{{.*}} '' postfix '++' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T' +#pragma acc atomic capture if (v < x) + v = x++; + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'' '<' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T' +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: UnaryOperator{{.*}} '' postfix '--' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T' +// CHECK-NEXT: BinaryOperator{{.*}} '' '=' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T' +#pragma acc atomic capture if (v < x) + { x--; v = x; } + // CHECK-NEXT: FunctionDecl{{.*}} templ_foo 'void (int, int)' implicit_instantiation // CHECK-NEXT: TemplateArgument type 'int' // CHECK-NEXT: BuiltinType{{.*}} 'int' @@ -170,6 +324,81 @@ void templ_foo(T v, T x) { // CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' // CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic read +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '=' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic write +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '=' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: BinaryOperator{{.*}}'int' '+' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'int' +// CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'int'{{.*}}I +// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5 + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic update +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '++' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '--' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '=' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '++' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' + +// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture +// CHECK-NEXT: if clause +// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '--' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' +// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '=' +// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int' +// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' +// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int' } void use() { diff --git a/clang/test/SemaOpenACC/atomic-construct.cpp b/clang/test/SemaOpenACC/atomic-construct.cpp index 7357d91d704f..83fbc4bff2ec 100644 --- a/clang/test/SemaOpenACC/atomic-construct.cpp +++ b/clang/test/SemaOpenACC/atomic-construct.cpp @@ -47,9 +47,30 @@ Struct &operator-(Struct&, Struct&); Struct S1, S2; +struct NotCondition{} NC; + template T &getRValue(); +void IfClause(int x, int v) { + // expected-error@+1{{OpenACC 'seq' clause is not valid on 'atomic' directive}} +#pragma acc atomic read seq + x = v; + + // expected-error@+1{{expected '('}} +#pragma acc atomic read if + x = v; + + // expected-error@+1{{value of type 'struct NotCondition' is not contextually convertible to 'bool'}} +#pragma acc atomic read if(NC) + x = v; + + // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'atomic' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc atomic read if(x) if (v) + x = v; +} + template void AtomicReadTemplate(T LHS, T RHS) { #pragma acc atomic read diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index 197ba2cd6856..6ea6447d1d59 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -3719,6 +3719,8 @@ void EnqueueVisitor::VisitOpenACCUpdateConstruct( void EnqueueVisitor::VisitOpenACCAtomicConstruct( const OpenACCAtomicConstruct *C) { EnqueueChildren(C); + for (auto *Clause : C->clauses()) + EnqueueChildren(Clause); } void EnqueueVisitor::VisitAnnotateAttr(const AnnotateAttr *A) {