Skip to content

Commit d7724c8

Browse files
erichkeaneerichkeane
committed
[OpenACC] allow 'if' clause on 'atomic' construct
This was added in OpenACC PR llvm#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.
1 parent b6e2df5 commit d7724c8

File tree

15 files changed

+388
-24
lines changed

15 files changed

+388
-24
lines changed

‎clang/include/clang/AST/StmtOpenACC.h‎

Lines changed: 26 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -829,24 +829,42 @@ class OpenACCUpdateConstruct final
829829

830830
// This class represents the 'atomic' construct, which has an associated
831831
// statement, but no clauses.
832-
class OpenACCAtomicConstruct final : public OpenACCAssociatedStmtConstruct {
832+
class OpenACCAtomicConstruct final
833+
: public OpenACCAssociatedStmtConstruct,
834+
private llvm::TrailingObjects<OpenACCAtomicConstruct,
835+
const OpenACCClause *> {
833836

834837
friend class ASTStmtReader;
838+
friend TrailingObjects;
835839
OpenACCAtomicKind AtomicKind = OpenACCAtomicKind::None;
836840

837-
OpenACCAtomicConstruct(EmptyShell)
841+
OpenACCAtomicConstruct(unsigned NumClauses)
838842
: OpenACCAssociatedStmtConstruct(
839843
OpenACCAtomicConstructClass, OpenACCDirectiveKind::Atomic,
840844
SourceLocation{}, SourceLocation{}, SourceLocation{},
841-
/*AssociatedStmt=*/nullptr) {}
845+
/*AssociatedStmt=*/nullptr) {
846+
std::uninitialized_value_construct(
847+
getTrailingObjects<const OpenACCClause *>(),
848+
getTrailingObjects<const OpenACCClause *>() + NumClauses);
849+
setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
850+
NumClauses));
851+
}
842852

843853
OpenACCAtomicConstruct(SourceLocation Start, SourceLocation DirectiveLoc,
844854
OpenACCAtomicKind AtKind, SourceLocation End,
855+
ArrayRef<const OpenACCClause *> Clauses,
845856
Stmt *AssociatedStmt)
846857
: OpenACCAssociatedStmtConstruct(OpenACCAtomicConstructClass,
847858
OpenACCDirectiveKind::Atomic, Start,
848859
DirectiveLoc, End, AssociatedStmt),
849-
AtomicKind(AtKind) {}
860+
AtomicKind(AtKind) {
861+
// Initialize the trailing storage.
862+
std::uninitialized_copy(Clauses.begin(), Clauses.end(),
863+
getTrailingObjects<const OpenACCClause *>());
864+
865+
setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
866+
Clauses.size()));
867+
}
850868

851869
void setAssociatedStmt(Stmt *S) {
852870
OpenACCAssociatedStmtConstruct::setAssociatedStmt(S);
@@ -857,10 +875,12 @@ class OpenACCAtomicConstruct final : public OpenACCAssociatedStmtConstruct {
857875
return T->getStmtClass() == OpenACCAtomicConstructClass;
858876
}
859877

860-
static OpenACCAtomicConstruct *CreateEmpty(const ASTContext &C);
878+
static OpenACCAtomicConstruct *CreateEmpty(const ASTContext &C,
879+
unsigned NumClauses);
861880
static OpenACCAtomicConstruct *
862881
Create(const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
863-
OpenACCAtomicKind AtKind, SourceLocation End, Stmt *AssociatedStmt);
882+
OpenACCAtomicKind AtKind, SourceLocation End,
883+
ArrayRef<const OpenACCClause *> Clauses, Stmt *AssociatedStmt);
864884

865885
OpenACCAtomicKind getAtomicKind() const { return AtomicKind; }
866886
const Stmt *getAssociatedStmt() const {

‎clang/lib/AST/StmtOpenACC.cpp‎

Lines changed: 13 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -307,20 +307,26 @@ OpenACCUpdateConstruct::Create(const ASTContext &C, SourceLocation Start,
307307
}
308308

309309
OpenACCAtomicConstruct *
310-
OpenACCAtomicConstruct::CreateEmpty(const ASTContext &C) {
311-
void *Mem = C.Allocate(sizeof(OpenACCAtomicConstruct));
312-
auto *Inst = new (Mem) OpenACCAtomicConstruct(EmptyShell{});
310+
OpenACCAtomicConstruct::CreateEmpty(const ASTContext &C, unsigned NumClauses) {
311+
void *Mem = C.Allocate(
312+
OpenACCAtomicConstruct::totalSizeToAlloc<const OpenACCClause *>(
313+
NumClauses));
314+
auto *Inst = new (Mem) OpenACCAtomicConstruct(NumClauses);
313315
return Inst;
314316
}
315317

316318
OpenACCAtomicConstruct *OpenACCAtomicConstruct::Create(
317319
const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
318-
OpenACCAtomicKind AtKind, SourceLocation End, Stmt *AssociatedStmt) {
319-
void *Mem = C.Allocate(sizeof(OpenACCAtomicConstruct));
320-
auto *Inst = new (Mem)
321-
OpenACCAtomicConstruct(Start, DirectiveLoc, AtKind, End, AssociatedStmt);
320+
OpenACCAtomicKind AtKind, SourceLocation End,
321+
ArrayRef<const OpenACCClause *> Clauses, Stmt *AssociatedStmt) {
322+
void *Mem = C.Allocate(
323+
OpenACCAtomicConstruct::totalSizeToAlloc<const OpenACCClause *>(
324+
Clauses.size()));
325+
auto *Inst = new (Mem) OpenACCAtomicConstruct(Start, DirectiveLoc, AtKind,
326+
End, Clauses, AssociatedStmt);
322327
return Inst;
323328
}
329+
324330
OpenACCCacheConstruct *OpenACCCacheConstruct::CreateEmpty(const ASTContext &C,
325331
unsigned NumVars) {
326332
void *Mem =

‎clang/lib/AST/StmtPrinter.cpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1258,6 +1258,7 @@ void StmtPrinter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {
12581258
if (S->getAtomicKind() != OpenACCAtomicKind::None)
12591259
OS << " " << S->getAtomicKind();
12601260

1261+
PrintOpenACCClauseList(S);
12611262
OS << '\n';
12621263
PrintStmt(S->getAssociatedStmt());
12631264
}

‎clang/lib/AST/StmtProfile.cpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2839,6 +2839,8 @@ void StmtProfiler::VisitOpenACCUpdateConstruct(
28392839
void StmtProfiler::VisitOpenACCAtomicConstruct(
28402840
const OpenACCAtomicConstruct *S) {
28412841
VisitStmt(S);
2842+
OpenACCClauseProfiler P{*this};
2843+
P.VisitOpenACCClauseList(S->clauses());
28422844
}
28432845

28442846
void StmtProfiler::VisitHLSLOutArgExpr(const HLSLOutArgExpr *S) {

‎clang/lib/Sema/SemaOpenACC.cpp‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1832,9 +1832,8 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(
18321832
EndLoc, Clauses);
18331833
}
18341834
case OpenACCDirectiveKind::Atomic: {
1835-
assert(Clauses.empty() && "Atomic doesn't allow clauses");
18361835
return OpenACCAtomicConstruct::Create(
1837-
getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc,
1836+
getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc, Clauses,
18381837
AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
18391838
}
18401839
case OpenACCDirectiveKind::Cache: {

‎clang/lib/Sema/SemaOpenACCClause.cpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -55,6 +55,8 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
5555
case OpenACCDirectiveKind::ParallelLoop:
5656
case OpenACCDirectiveKind::SerialLoop:
5757
case OpenACCDirectiveKind::KernelsLoop:
58+
// OpenACC 3.4(prerelease) PR #511 adds 'if' to atomic.
59+
case OpenACCDirectiveKind::Atomic:
5860
return true;
5961
default:
6062
return false;

‎clang/lib/Sema/TreeTransform.h‎

Lines changed: 7 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4222,10 +4222,11 @@ class TreeTransform {
42224222
SourceLocation DirLoc,
42234223
OpenACCAtomicKind AtKind,
42244224
SourceLocation EndLoc,
4225+
ArrayRef<OpenACCClause *> Clauses,
42254226
StmtResult AssociatedStmt) {
42264227
return getSema().OpenACC().ActOnEndStmtDirective(
42274228
OpenACCDirectiveKind::Atomic, BeginLoc, DirLoc, SourceLocation{},
4228-
SourceLocation{}, {}, AtKind, SourceLocation{}, EndLoc, {},
4229+
SourceLocation{}, {}, AtKind, SourceLocation{}, EndLoc, Clauses,
42294230
AssociatedStmt);
42304231
}
42314232

@@ -12744,6 +12745,10 @@ StmtResult TreeTransform<Derived>::TransformOpenACCAtomicConstruct(
1274412745
OpenACCAtomicConstruct *C) {
1274512746
getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());
1274612747

12748+
llvm::SmallVector<OpenACCClause *> TransformedClauses =
12749+
getDerived().TransformOpenACCClauseList(C->getDirectiveKind(),
12750+
C->clauses());
12751+
1274712752
if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(),
1274812753
C->getBeginLoc(), {}))
1274912754
return StmtError();
@@ -12759,7 +12764,7 @@ StmtResult TreeTransform<Derived>::TransformOpenACCAtomicConstruct(
1275912764

1276012765
return getDerived().RebuildOpenACCAtomicConstruct(
1276112766
C->getBeginLoc(), C->getDirectiveLoc(), C->getAtomicKind(),
12762-
C->getEndLoc(), AssocStmt);
12767+
C->getEndLoc(), TransformedClauses, AssocStmt);
1276312768
}
1276412769

1276512770
template <typename Derived>

‎clang/lib/Serialization/ASTReaderStmt.cpp‎

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2940,9 +2940,7 @@ void ASTStmtReader::VisitOpenACCCacheConstruct(OpenACCCacheConstruct *S) {
29402940

29412941
void ASTStmtReader::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {
29422942
VisitStmt(S);
2943-
S->Kind = Record.readEnum<OpenACCDirectiveKind>();
2944-
S->Range = Record.readSourceRange();
2945-
S->DirectiveLoc = Record.readSourceLocation();
2943+
VisitOpenACCConstructStmt(S);
29462944
S->AtomicKind = Record.readEnum<OpenACCAtomicKind>();
29472945
S->setAssociatedStmt(Record.readSubStmt());
29482946
}
@@ -4490,7 +4488,8 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
44904488
break;
44914489
}
44924490
case STMT_OPENACC_ATOMIC_CONSTRUCT: {
4493-
S = OpenACCAtomicConstruct::CreateEmpty(Context);
4491+
unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
4492+
S = OpenACCAtomicConstruct::CreateEmpty(Context, NumClauses);
44944493
break;
44954494
}
44964495
case EXPR_REQUIRES: {

‎clang/lib/Serialization/ASTWriterStmt.cpp‎

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -3014,9 +3014,7 @@ void ASTStmtWriter::VisitOpenACCWaitConstruct(OpenACCWaitConstruct *S) {
30143014

30153015
void ASTStmtWriter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {
30163016
VisitStmt(S);
3017-
Record.writeEnum(S->Kind);
3018-
Record.AddSourceRange(S->Range);
3019-
Record.AddSourceLocation(S->DirectiveLoc);
3017+
VisitOpenACCConstructStmt(S);
30203018
Record.writeEnum(S->getAtomicKind());
30213019
Record.AddStmt(S->getAssociatedStmt());
30223020

‎clang/test/AST/ast-print-openacc-atomic-construct.cpp‎

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -31,3 +31,35 @@ void foo(int v, int x) {
3131
{ x--; v = x; }
3232

3333
}
34+
35+
void foo2(int v, int x) {
36+
// CHECK: #pragma acc atomic read if(v)
37+
// CHECK-NEXT: v = x;
38+
#pragma acc atomic read if (v)
39+
v = x;
40+
// CHECK-NEXT: pragma acc atomic write if(x)
41+
// CHECK-NEXT: v = x + 1;
42+
#pragma acc atomic write if (x)
43+
v = x + 1;
44+
// CHECK-NEXT: pragma acc atomic update if(true)
45+
// CHECK-NEXT: x++;
46+
#pragma acc atomic update if (true)
47+
x++;
48+
// CHECK-NEXT: pragma acc atomic if(false)
49+
// CHECK-NEXT: x--;
50+
#pragma acc atomic if (false)
51+
x--;
52+
// CHECK-NEXT: pragma acc atomic capture if(v < x)
53+
// CHECK-NEXT: v = x++;
54+
#pragma acc atomic capture if (v < x)
55+
v = x++;
56+
57+
// CHECK-NEXT: #pragma acc atomic capture if(x > v)
58+
// CHECK-NEXT: {
59+
// CHECK-NEXT: x--;
60+
// CHECK-NEXT: v = x;
61+
// CHECK-NEXT: }
62+
#pragma acc atomic capture if (x > v)
63+
{ x--; v = x; }
64+
65+
}

0 commit comments

Comments
 (0)