[OpenACC] Implement 'if_present' clause sema

The 'if_present' clause controls the replacement of addresses in the
var-list in current device memory.  This clause can only go on
'host_device'.  From a Sema perspective, there isn't anything to do
beyond add this to AST and pass it on.

Author: erichkeane

clang/include/clang/AST/OpenACCClause.h

@@ -98,6 +98,28 @@ class OpenACCFinalizeClause : public OpenACCClause {
   }
 };
 
+// Represents the 'if_present' clause.
+class OpenACCIfPresentClause : public OpenACCClause {
+protected:
+  OpenACCIfPresentClause(SourceLocation BeginLoc, SourceLocation EndLoc)
+      : OpenACCClause(OpenACCClauseKind::IfPresent, BeginLoc, EndLoc) {}
+
+public:
+  static bool classof(const OpenACCClause *C) {
+    return C->getClauseKind() == OpenACCClauseKind::IfPresent;
+  }
+
+  static OpenACCIfPresentClause *
+  Create(const ASTContext &Ctx, SourceLocation BeginLoc, SourceLocation EndLoc);
+
+  child_range children() {
+    return child_range(child_iterator(), child_iterator());
+  }
+  const_child_range children() const {
+    return const_child_range(const_child_iterator(), const_child_iterator());
+  }
+};
+
 // Represents the 'independent' clause.
 class OpenACCIndependentClause : public OpenACCClause {
 protected:

clang/include/clang/Basic/OpenACCClauses.def

@@ -45,6 +45,7 @@ VISIT_CLAUSE(Finalize)
 VISIT_CLAUSE(FirstPrivate)
 VISIT_CLAUSE(Gang)
 VISIT_CLAUSE(If)
+VISIT_CLAUSE(IfPresent)
 VISIT_CLAUSE(Independent)
 VISIT_CLAUSE(NoCreate)
 VISIT_CLAUSE(NumGangs)

clang/lib/AST/OpenACCClause.cpp

@@ -452,6 +452,14 @@ OpenACCFinalizeClause *OpenACCFinalizeClause::Create(const ASTContext &C,
   return new (Mem) OpenACCFinalizeClause(BeginLoc, EndLoc);
 }
 
+OpenACCIfPresentClause *OpenACCIfPresentClause::Create(const ASTContext &C,
+                                                       SourceLocation BeginLoc,
+                                                       SourceLocation EndLoc) {
+  void *Mem = C.Allocate(sizeof(OpenACCIfPresentClause),
+                         alignof(OpenACCIfPresentClause));
+  return new (Mem) OpenACCIfPresentClause(BeginLoc, EndLoc);
+}
+
 //===----------------------------------------------------------------------===//
 //  OpenACC clauses printing methods
 //===----------------------------------------------------------------------===//
@@ -697,3 +705,8 @@ void OpenACCClausePrinter::VisitVectorClause(const OpenACCVectorClause &C) {
 void OpenACCClausePrinter::VisitFinalizeClause(const OpenACCFinalizeClause &C) {
   OS << "finalize";
 }
+
+void OpenACCClausePrinter::VisitIfPresentClause(
+    const OpenACCIfPresentClause &C) {
+  OS << "if_present";
+}

clang/lib/AST/StmtProfile.cpp

@@ -2561,6 +2561,9 @@ void OpenACCClauseProfiler::VisitSelfClause(const OpenACCSelfClause &Clause) {
 void OpenACCClauseProfiler::VisitFinalizeClause(
     const OpenACCFinalizeClause &Clause) {}
 
+void OpenACCClauseProfiler::VisitIfPresentClause(
+    const OpenACCIfPresentClause &Clause) {}
+
 void OpenACCClauseProfiler::VisitNumGangsClause(
     const OpenACCNumGangsClause &Clause) {
   for (auto *E : Clause.getIntExprs())

clang/lib/AST/TextNodeDumper.cpp

@@ -409,6 +409,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
     case OpenACCClauseKind::PCopy:
     case OpenACCClauseKind::PresentOrCopy:
     case OpenACCClauseKind::If:
+    case OpenACCClauseKind::IfPresent:
     case OpenACCClauseKind::Independent:
     case OpenACCClauseKind::DevicePtr:
     case OpenACCClauseKind::Finalize:

clang/lib/Sema/SemaOpenACC.cpp

@@ -416,6 +416,15 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
       return false;
     }
   }
+  case OpenACCClauseKind::IfPresent: {
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::HostData:
+    case OpenACCDirectiveKind::Update:
+      return true;
+    default:
+      return false;
+    }
+  }
   }
 
   default:
@@ -1620,6 +1629,16 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitFinalizeClause(
                                        Clause.getEndLoc());
 }
 
+OpenACCClause *SemaOpenACCClauseVisitor::VisitIfPresentClause(
+    SemaOpenACC::OpenACCParsedClause &Clause) {
+  if (Clause.getDirectiveKind() != OpenACCDirectiveKind::HostData)
+    return isNotImplemented();
+  // There isn't anything to do here, this is only valid on one construct, and
+  // has no associated rules.
+  return OpenACCIfPresentClause::Create(Ctx, Clause.getBeginLoc(),
+                                        Clause.getEndLoc());
+}
+
 OpenACCClause *SemaOpenACCClauseVisitor::VisitSeqClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
   // Restrictions only properly implemented on 'loop' constructs and combined ,

clang/lib/Sema/TreeTransform.h

@@ -11986,6 +11986,14 @@ void OpenACCClauseTransform<Derived>::VisitFinalizeClause(
                                             ParsedClause.getEndLoc());
 }
 
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitIfPresentClause(
+    const OpenACCIfPresentClause &C) {
+  NewClause = OpenACCIfPresentClause::Create(Self.getSema().getASTContext(),
+                                             ParsedClause.getBeginLoc(),
+                                             ParsedClause.getEndLoc());
+}
+
 template <typename Derived>
 void OpenACCClauseTransform<Derived>::VisitReductionClause(
     const OpenACCReductionClause &C) {

clang/lib/Serialization/ASTReader.cpp

@@ -12536,6 +12536,8 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
     return OpenACCSeqClause::Create(getContext(), BeginLoc, EndLoc);
   case OpenACCClauseKind::Finalize:
     return OpenACCFinalizeClause::Create(getContext(), BeginLoc, EndLoc);
+  case OpenACCClauseKind::IfPresent:
+    return OpenACCIfPresentClause::Create(getContext(), BeginLoc, EndLoc);
   case OpenACCClauseKind::Independent:
     return OpenACCIndependentClause::Create(getContext(), BeginLoc, EndLoc);
   case OpenACCClauseKind::Auto:
@@ -12581,7 +12583,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
                                        VectorExpr, EndLoc);
   }
 
-  case OpenACCClauseKind::IfPresent:
   case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::UseDevice:
   case OpenACCClauseKind::Delete:

clang/lib/Serialization/ASTWriter.cpp

@@ -8452,6 +8452,7 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
   case OpenACCClauseKind::Independent:
   case OpenACCClauseKind::Auto:
   case OpenACCClauseKind::Finalize:
+  case OpenACCClauseKind::IfPresent:
     // Nothing to do here, there is no additional information beyond the
     // begin/end loc and clause kind.
     return;
@@ -8497,7 +8498,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
     return;
   }
 
-  case OpenACCClauseKind::IfPresent:
   case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::UseDevice:
   case OpenACCClauseKind::Delete:

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

@@ -113,4 +113,8 @@ void foo() {
 
 // CHECK: #pragma acc exit data copyout(i) finalize
 #pragma acc exit data copyout(i) finalize
+
+// CHECK: #pragma acc host_data if_present
+#pragma acc host_data use_device(i) if_present
+  ;
 }

clang/test/ParserOpenACC/parse-clauses.c

@@ -16,11 +16,8 @@ void func() {
 
 #pragma acc exit data wait finalize
 
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented, clause ignored}}
 #pragma acc host_data if_present
 
-  // expected-warning@+2{{OpenACC clause 'if_present' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented, clause ignored}}
 #pragma acc host_data if_present, if_present
 
   // expected-error@+4{{OpenACC clause 'independent' on 'loop' construct conflicts with previous data dependence clause}}

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

@@ -40,7 +40,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop auto finalize
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop auto if_present
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop auto worker
@@ -157,7 +157,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop finalize auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop if_present auto
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop worker auto
@@ -275,7 +275,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop independent finalize
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop independent if_present
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop independent worker
@@ -392,7 +392,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop finalize independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop if_present independent
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop worker independent
@@ -522,7 +522,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop seq finalize
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop seq if_present
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
@@ -645,7 +645,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop finalize seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop if_present seq
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}

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

@@ -45,8 +45,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'serial loop' directive}}
 #pragma acc serial loop device_type(*) finalize
   for(int i = 0; i < 5; ++i);
-  // expected-error@+2{{OpenACC clause 'if_present' may not follow a 'device_type' clause in a 'kernels loop' construct}}
-  // expected-note@+1{{previous clause is here}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'kernels loop' directive}}
 #pragma acc kernels loop device_type(*) if_present
   for(int i = 0; i < 5; ++i);
 #pragma acc parallel loop device_type(*) seq

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

@@ -45,8 +45,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'kernels' directive}}
 #pragma acc kernels device_type(*) finalize
   while(1);
-  // expected-error@+2{{OpenACC clause 'if_present' may not follow a 'device_type' clause in a 'kernels' construct}}
-  // expected-note@+1{{previous clause is here}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'kernels' directive}}
 #pragma acc kernels device_type(*) if_present
   while(1);
   // expected-error@+1{{OpenACC 'seq' clause is not valid on 'kernels' directive}}

clang/test/SemaOpenACC/data-construct-if_present-ast.cpp

@@ -0,0 +1,59 @@
+// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
+
+// Test this with PCH.
+// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
+// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
+#ifndef PCH_HELPER
+#define PCH_HELPER
+
+void Uses() {
+  // CHECK: FunctionDecl{{.*}}Uses
+  // CHECK-NEXT: CompoundStmt
+
+  int I;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+
+#pragma acc host_data use_device(I) if_present
+  ;
+  // CHECK-NEXT: OpenACCHostDataConstruct{{.*}}host_data
+  // CHECK-NEXT: if_present clause
+  // CHECK-NEXT: NullStmt
+}
+
+template<typename T>
+void TemplUses() {
+  // CHECK: FunctionTemplateDecl{{.*}}TemplUses
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}T
+  // CHECK-NEXT: FunctionDecl{{.*}}TemplUses
+  // CHECK-NEXT: CompoundStmt
+
+  T I;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+
+#pragma acc host_data use_device(I) if_present
+  ;
+  // CHECK-NEXT: OpenACCHostDataConstruct{{.*}}host_data
+  // CHECK-NEXT: if_present clause
+  // CHECK-NEXT: NullStmt
+
+  // Instantiations
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void ()' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument type 'int'
+  // CHECK-NEXT: BuiltinType{{.*}} 'int'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+
+  // CHECK-NEXT: OpenACCHostDataConstruct{{.*}}host_data
+  // CHECK-NEXT: if_present clause
+  // CHECK-NEXT: NullStmt
+}
+void Inst() {
+  TemplUses<int>();
+}
+
+
+#endif // PCH_HELPER

clang/test/SemaOpenACC/data-construct-if_present-clause.c

@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+void Test() {
+  int I;
+
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'data' directive}}
+#pragma acc data copyin(I) if_present
+  ;
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'enter data' directive}}
+#pragma acc enter data copyin(I) if_present
+  ;
+
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'exit data' directive}}
+#pragma acc exit data copyout(I) if_present
+  ;
+  // expected-warning@+1{{OpenACC clause 'use_device' not yet implemented}}
+#pragma acc host_data use_device(I) if_present
+  ;
+}

clang/test/SemaOpenACC/data-construct.cpp

@@ -96,7 +96,6 @@ void AtLeastOneOf() {
 
 #pragma acc host_data if(Var)
   ;
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
 #pragma acc host_data if_present
   ;
 #pragma acc host_data

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

@@ -40,7 +40,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'loop' directive}}
 #pragma acc loop auto finalize
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'loop' directive}}
 #pragma acc loop auto if_present
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc loop auto worker
@@ -174,7 +174,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'loop' directive}}
 #pragma acc loop finalize auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'loop' directive}}
 #pragma acc loop if_present auto
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc loop worker auto
@@ -309,7 +309,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'loop' directive}}
 #pragma acc loop independent finalize
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'loop' directive}}
 #pragma acc loop independent if_present
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc loop independent worker
@@ -443,7 +443,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'loop' directive}}
 #pragma acc loop finalize independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'loop' directive}}
 #pragma acc loop if_present independent
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc loop worker independent
@@ -590,7 +590,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'loop' directive}}
 #pragma acc loop seq finalize
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'loop' directive}}
 #pragma acc loop seq if_present
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
@@ -730,7 +730,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'loop' directive}}
 #pragma acc loop finalize seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
+  // expected-error@+1{{OpenACC 'if_present' clause is not valid on 'loop' directive}}
 #pragma acc loop if_present seq
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}