[OpenACC] allow 'if' clause on 'atomic' construct

erichkeane · erichkeane · commit d7724c8ea313 · 2025-04-02T10:03:24.000-07:00
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.
diff --git a/clang/include/clang/AST/StmtOpenACC.h b/clang/include/clang/AST/StmtOpenACC.h
@@ -829,24 +829,42 @@ class OpenACCUpdateConstruct final
 
 // 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<OpenACCAtomicConstruct,
+                                    const OpenACCClause *> {
 
   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<const OpenACCClause *>(),
+        getTrailingObjects<const OpenACCClause *>() + NumClauses);
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  NumClauses));
+  }
 
   OpenACCAtomicConstruct(SourceLocation Start, SourceLocation DirectiveLoc,
                          OpenACCAtomicKind AtKind, SourceLocation End,
+                         ArrayRef<const OpenACCClause *> 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<const OpenACCClause *>());
+
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  Clauses.size()));
+  }
 
   void setAssociatedStmt(Stmt *S) {
     OpenACCAssociatedStmtConstruct::setAssociatedStmt(S);
@@ -857,10 +875,12 @@ class OpenACCAtomicConstruct final : public OpenACCAssociatedStmtConstruct {
     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<const OpenACCClause *> 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
@@ -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<const OpenACCClause *>(
+          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<const OpenACCClause *> Clauses, Stmt *AssociatedStmt) {
+  void *Mem = C.Allocate(
+      OpenACCAtomicConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -4222,10 +4222,11 @@ class TreeTransform {
                                            SourceLocation DirLoc,
                                            OpenACCAtomicKind AtKind,
                                            SourceLocation EndLoc,
+                                           ArrayRef<OpenACCClause *> 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<Derived>::TransformOpenACCAtomicConstruct(
     OpenACCAtomicConstruct *C) {
   getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());
 
+  llvm::SmallVector<OpenACCClause *> TransformedClauses =
+      getDerived().TransformOpenACCClauseList(C->getDirectiveKind(),
+                                              C->clauses());
+
   if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(),
                                                   C->getBeginLoc(), {}))
     return StmtError();
@@ -12759,7 +12764,7 @@ StmtResult TreeTransform<Derived>::TransformOpenACCAtomicConstruct(
 
   return getDerived().RebuildOpenACCAtomicConstruct(
       C->getBeginLoc(), C->getDirectiveLoc(), C->getAtomicKind(),
-      C->getEndLoc(), AssocStmt);
+      C->getEndLoc(), TransformedClauses, AssocStmt);
 }
 
 template <typename Derived>
diff --git 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<OpenACCDirectiveKind>();
-  S->Range = Record.readSourceRange();
-  S->DirectiveLoc = Record.readSourceLocation();
+  VisitOpenACCConstructStmt(S);
   S->AtomicKind = Record.readEnum<OpenACCAtomicKind>();
   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
@@ -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
@@ -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
@@ -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
@@ -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
diff --git a/clang/test/SemaOpenACC/atomic-construct.cpp b/clang/test/SemaOpenACC/atomic-construct.cpp
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp

Original file line number	Diff line number	Diff line change
`@@ -1258,6 +1258,7 @@ void StmtPrinter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {`
`1258`	`1258`	`if (S->getAtomicKind() != OpenACCAtomicKind::None)`
`1259`	`1259`	`OS << " " << S->getAtomicKind();`
`1260`	`1260`
	`1261`	`+ PrintOpenACCClauseList(S);`
`1261`	`1262`	`OS << '\n';`
`1262`	`1263`	`PrintStmt(S->getAssociatedStmt());`
`1263`	`1264`	`}`
Original file line number	Diff line number	Diff line change
`@@ -2839,6 +2839,8 @@ void StmtProfiler::VisitOpenACCUpdateConstruct(`
`2839`	`2839`	`void StmtProfiler::VisitOpenACCAtomicConstruct(`
`2840`	`2840`	`const OpenACCAtomicConstruct *S) {`
`2841`	`2841`	`VisitStmt(S);`
	`2842`	`+ OpenACCClauseProfiler P{*this};`
	`2843`	`+ P.VisitOpenACCClauseList(S->clauses());`
`2842`	`2844`	`}`
`2843`	`2845`
`2844`	`2846`	`void StmtProfiler::VisitHLSLOutArgExpr(const HLSLOutArgExpr *S) {`
Original file line number	Diff line number	Diff line change
`@@ -1832,9 +1832,8 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(`
`1832`	`1832`	`EndLoc, Clauses);`
`1833`	`1833`	`}`
`1834`	`1834`	`case OpenACCDirectiveKind::Atomic: {`
`1835`		`- assert(Clauses.empty() && "Atomic doesn't allow clauses");`
`1836`	`1835`	`return OpenACCAtomicConstruct::Create(`
`1837`		`- getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc,`
	`1836`	`+ getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc, Clauses,`
`1838`	`1837`	`AssocStmt.isUsable() ? AssocStmt.get() : nullptr);`
`1839`	`1838`	`}`
`1840`	`1839`	`case OpenACCDirectiveKind::Cache: {`
Original file line number	Diff line number	Diff line change
`@@ -2940,9 +2940,7 @@ void ASTStmtReader::VisitOpenACCCacheConstruct(OpenACCCacheConstruct *S) {`
`2940`	`2940`
`2941`	`2941`	`void ASTStmtReader::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {`
`2942`	`2942`	`VisitStmt(S);`
`2943`		`- S->Kind = Record.readEnum<OpenACCDirectiveKind>();`
`2944`		`- S->Range = Record.readSourceRange();`
`2945`		`- S->DirectiveLoc = Record.readSourceLocation();`
	`2943`	`+ VisitOpenACCConstructStmt(S);`
`2946`	`2944`	`S->AtomicKind = Record.readEnum<OpenACCAtomicKind>();`
`2947`	`2945`	`S->setAssociatedStmt(Record.readSubStmt());`
`2948`	`2946`	`}`
`@@ -4490,7 +4488,8 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {`
`4490`	`4488`	`break;`
`4491`	`4489`	`}`
`4492`	`4490`	`case STMT_OPENACC_ATOMIC_CONSTRUCT: {`
`4493`		`- S = OpenACCAtomicConstruct::CreateEmpty(Context);`
	`4491`	`+ unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];`
	`4492`	`+ S = OpenACCAtomicConstruct::CreateEmpty(Context, NumClauses);`
`4494`	`4493`	`break;`
`4495`	`4494`	`}`
`4496`	`4495`	`case EXPR_REQUIRES: {`