[OpenACC] Implement 'tile' for combined constructs

Like collapse, this clause has only mild changes that need to be made.
Most of the associated work for the RAII container was already done
previously, so this is mostly just updating the diagnostics to properly
emit the right construct.

Author: erichkeane

clang/include/clang/Sema/SemaOpenACC.h

@@ -107,6 +107,10 @@ class SemaOpenACC : public SemaBase {
     /// which allows us to diagnose if the number of arguments is too large for
     /// the current number of 'for' loops.
     bool TileDepthSatisfied = true;
+
+    /// Records the kind of the directive that this clause is attached to, which
+    /// allows us to use it in diagnostics.
+    OpenACCDirectiveKind DirectiveKind = OpenACCDirectiveKind::Invalid;
   } TileInfo;
 
   /// A list of the active reduction clauses, which allows us to check that all

clang/lib/Sema/SemaOpenACC.cpp

@@ -598,9 +598,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDefaultClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitTileClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // TODO OpenACC: Remove this when we get combined construct impl for this.
-  if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
-    return isNotImplemented();
 
   // Duplicates here are not really sensible.  We could possible permit
   // multiples if they all had the same value, but there isn't really a good
@@ -1718,6 +1715,7 @@ void SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt(
   SemaRef.TileInfo.ActiveTile = TileClause;
   SemaRef.TileInfo.TileDepthSatisfied = false;
   SemaRef.TileInfo.CurTileCount = TileClause->getSizeExprs().size();
+  SemaRef.TileInfo.DirectiveKind = DirKind;
 }
 
 SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
@@ -2608,7 +2606,7 @@ void SemaOpenACC::ActOnWhileStmt(SourceLocation WhileLoc) {
 
   if (TileInfo.CurTileCount && *TileInfo.CurTileCount > 0) {
     Diag(WhileLoc, diag::err_acc_invalid_in_loop)
-        << /*while loop*/ 1 << OpenACCDirectiveKind::Loop
+        << /*while loop*/ 1 << TileInfo.DirectiveKind
         << OpenACCClauseKind::Tile;
     assert(TileInfo.ActiveTile && "tile count without object?");
     Diag(TileInfo.ActiveTile->getBeginLoc(), diag::note_acc_active_clause_here)
@@ -2643,8 +2641,7 @@ void SemaOpenACC::ActOnDoStmt(SourceLocation DoLoc) {
 
   if (TileInfo.CurTileCount && *TileInfo.CurTileCount > 0) {
     Diag(DoLoc, diag::err_acc_invalid_in_loop)
-        << /*do loop*/ 2 << OpenACCDirectiveKind::Loop
-        << OpenACCClauseKind::Tile;
+        << /*do loop*/ 2 << TileInfo.DirectiveKind << OpenACCClauseKind::Tile;
     assert(TileInfo.ActiveTile && "tile count without object?");
     Diag(TileInfo.ActiveTile->getBeginLoc(), diag::note_acc_active_clause_here)
         << OpenACCClauseKind::Tile;
@@ -2692,7 +2689,7 @@ void SemaOpenACC::ForStmtBeginHelper(SourceLocation ForLoc,
 
     if (LoopInfo.CurLevelHasLoopAlready) {
       Diag(ForLoc, diag::err_acc_clause_multiple_loops)
-          << OpenACCDirectiveKind::Loop << OpenACCClauseKind::Tile;
+          << TileInfo.DirectiveKind << OpenACCClauseKind::Tile;
       assert(TileInfo.ActiveTile && "No tile object?");
       Diag(TileInfo.ActiveTile->getBeginLoc(),
            diag::note_acc_active_clause_here)
@@ -3203,7 +3200,7 @@ void SemaOpenACC::ActOnForStmtEnd(SourceLocation ForLoc, StmtResult Body) {
 
     if (OtherStmtLoc.isValid() && IsActiveTile) {
       Diag(OtherStmtLoc, diag::err_acc_intervening_code)
-          << OpenACCClauseKind::Tile << OpenACCDirectiveKind::Loop;
+          << OpenACCClauseKind::Tile << TileInfo.DirectiveKind;
       Diag(TileInfo.ActiveTile->getBeginLoc(),
            diag::note_acc_active_clause_here)
           << OpenACCClauseKind::Tile;
@@ -3232,7 +3229,7 @@ bool SemaOpenACC::ActOnStartStmtDirective(OpenACCDirectiveKind K,
   }
   if (TileInfo.CurTileCount && *TileInfo.CurTileCount > 0) {
     Diag(StartLoc, diag::err_acc_invalid_in_loop)
-        << /*OpenACC Construct*/ 0 << OpenACCDirectiveKind::Loop
+        << /*OpenACC Construct*/ 0 << TileInfo.DirectiveKind
         << OpenACCClauseKind::Tile << K;
     assert(TileInfo.ActiveTile && "Tile count without object?");
     Diag(TileInfo.ActiveTile->getBeginLoc(), diag::note_acc_active_clause_here)

clang/test/AST/ast-print-openacc-combined-construct.cpp

@@ -1,5 +1,6 @@
 // RUN: %clang_cc1 -fopenacc -Wno-openacc-deprecated-clause-alias -ast-print %s -o - | FileCheck %s
 
+constexpr int get_value() { return 1; }
 void foo() {
   int *iPtr;
 // CHECK: #pragma acc parallel loop
@@ -210,4 +211,17 @@ void foo() {
 #pragma acc parallel loop collapse(force:2)
   for(int i = 0;i<5;++i)
     for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc serial loop tile(1, 3, *, get_value())
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc serial loop tile(1, 3, *, get_value())
+  for(int i = 0;i<5;++i)
+    for(int i = 0;i<5;++i)
+      for(int i = 0;i<5;++i)
+        for(int i = 0;i<5;++i);
+
 }

clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c

@@ -158,7 +158,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop auto async
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}}
 #pragma acc parallel loop auto tile(1+2, 1)
   for(unsigned j = 0; j < 5; ++j)
     for(unsigned i = 0; i < 5; ++i);
@@ -286,7 +285,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop async auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}}
 #pragma acc parallel loop tile(1+2, 1) auto
   for(unsigned j = 0; j < 5; ++j)
     for(unsigned i = 0; i < 5; ++i);
@@ -415,7 +413,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop independent async
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}}
 #pragma acc parallel loop independent tile(1+2, 1)
   for(unsigned j = 0; j < 5; ++j)
     for(unsigned i = 0; i < 5; ++i);
@@ -543,7 +540,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop async independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}}
 #pragma acc parallel loop tile(1+2, 1) independent
   for(unsigned j = 0; j < 5; ++j)
     for(unsigned i = 0; i < 5; ++i);
@@ -678,7 +674,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop seq async
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}}
 #pragma acc parallel loop seq tile(1+2, 1)
   for(;;)
     for(unsigned i = 0; i < 5; ++i);
@@ -812,7 +807,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop async seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}}
 #pragma acc parallel loop tile(1+2, 1) seq
   for(;;)
     for(unsigned i = 0; i < 5; ++i);

clang/test/SemaOpenACC/combined-construct-device_type-clause.c

@@ -209,7 +209,6 @@ void uses() {
 #pragma acc parallel loop device_type(*) async
   for(int i = 0; i < 5; ++i);
 
-  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
 #pragma acc serial loop device_type(*) tile(*, 1)
   for(int j = 0; j < 5; ++j)
     for(int i = 0; i < 5; ++i);