[OpenACC] Implement 'if' clause sema for 'data' constructs

This is another one that has no additional sema work other than enabling
it, so this patch does just that.

clang/include/clang/Basic/OpenACCKinds.h

@@ -158,6 +158,14 @@ inline bool isOpenACCCombinedDirectiveKind(OpenACCDirectiveKind K) {
          K == OpenACCDirectiveKind::KernelsLoop;
 }
 
+// Tests 'K' to see if it is 'data', 'host_data', 'enter data', or 'exit data'.
+inline bool isOpenACCDataDirectiveKind(OpenACCDirectiveKind K) {
+  return K == OpenACCDirectiveKind::Data ||
+         K == OpenACCDirectiveKind::EnterData ||
+         K == OpenACCDirectiveKind::ExitData ||
+         K == OpenACCDirectiveKind::HostData;
+}
+
 enum class OpenACCAtomicKind : uint8_t {
   Read,
   Write,

clang/lib/Sema/SemaOpenACC.cpp

@@ -641,16 +641,20 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitTileClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitIfClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // Restrictions only properly implemented on 'compute'/'combined' constructs,
-  // and 'compute'/'combined' constructs are the only construct that can do
-  // anything with this yet, so skip/treat as unimplemented in this case.
+  // Restrictions only properly implemented on 'compute'/'combined'/'data'
+  // constructs, and 'compute'/'combined'/'data' constructs are the only
+  // constructs that can do anything with this yet, so skip/treat as
+  // unimplemented in this case.
   if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
-      !isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
+      !isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()) &&
+      !isOpenACCDataDirectiveKind(Clause.getDirectiveKind()))
     return isNotImplemented();
 
   // There is no prose in the standard that says duplicates aren't allowed,
   // but this diagnostic is present in other compilers, as well as makes
-  // sense.
+  // sense. Prose DOES exist for 'data' and 'host_data', 'enter data' and 'exit
+  // data' both don't, but other implmementations do this.  OpenACC issue 519
+  // filed for the latter two.
   if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause))
     return nullptr;
 

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

@@ -27,4 +27,20 @@ void foo() {
 // CHECK-NOT: use_device(Var)
 #pragma acc host_data use_device(Var)
   ;
+
+  int i;
+  int array[5];
+
+// CHECK: #pragma acc data if(i == array[1])
+#pragma acc data default(none) if(i == array[1])
+  ;
+// CHECK: #pragma acc enter data if(i == array[1])
+#pragma acc enter data copyin(Var) if(i == array[1])
+  ;
+// CHECK: #pragma acc exit data if(i == array[1])
+#pragma acc exit data copyout(Var) if(i == array[1])
+  ;
+// CHECK: #pragma acc host_data if(i == array[1])
+#pragma acc host_data use_device(Var) if(i == array[1])
+  ;
 }

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

@@ -43,7 +43,6 @@ void BoolExpr(int *I, float *F) {
 #pragma acc kernels loop if (*I < *F)
   for (unsigned i = 0; i < 5; ++i);
 
-  // expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
 #pragma acc data if (*I < *F)
   for (unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop if (*I < *F)

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

@@ -43,7 +43,6 @@ void BoolExpr(int *I, float *F) {
 #pragma acc kernels if (*I < *F)
   while(0);
 
-  // expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
 #pragma acc data if (*I < *F)
   while(0);
 #pragma acc parallel loop if (*I < *F)

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

@@ -0,0 +1,128 @@
+// 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 NormalFunc(int j, float f) {
+  // CHECK: FunctionDecl{{.*}}NormalFunc
+  // CHECK-NEXT: ParmVarDecl
+  // CHECK-NEXT: ParmVarDecl
+  // CHECK-NEXT: CompoundStmt
+#pragma acc data if( j < f) default(none)
+  ;
+  // CHECK-NEXT: OpenACCDataConstruct{{.*}}data
+  // CHECK-NEXT: if clause
+  // CHECK-NEXT: BinaryOperator{{.*}} 'bool' '<'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'float' <IntegralToFloating>
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'j' 'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'float' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'float' lvalue ParmVar{{.*}} 'f' 'float'
+  // CHECK-NEXT: NullStmt
+
+}
+
+int Global;
+
+template<typename T>
+void TemplFunc() {
+  // CHECK: FunctionTemplateDecl{{.*}}TemplFunc
+  // CHECK-NEXT: TemplateTypeParmDecl
+
+  // Match the prototype:
+  // CHECK-NEXT: FunctionDecl{{.*}}TemplFunc
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc data default(none) if(T::SomeFloat < typename T::IntTy{})
+  ;
+  // CHECK-NEXT: OpenACCDataConstruct{{.*}}data
+  // CHECK-NEXT: if clause
+  // CHECK-NEXT: BinaryOperator{{.*}} '<dependent type>' '<'
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+  // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'typename T::IntTy' 'typename T::IntTy'
+  // CHECK-NEXT: InitListExpr{{.*}} 'void'
+  // CHECK-NEXT: NullStmt
+
+#pragma acc enter data copyin(Global) if(typename T::IntTy{})
+  ;
+  // CHECK-NEXT: OpenACCEnterDataConstruct{{.*}}enter data
+  // CHECK-NEXT: if clause
+  // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'typename T::IntTy' 'typename T::IntTy'
+  // CHECK-NEXT: InitListExpr{{.*}} 'void'
+  // CHECK-NEXT: NullStmt
+
+#pragma acc exit data copyout(Global) if(T::SomeFloat)
+  ;
+  // CHECK-NEXT: OpenACCExitDataConstruct{{.*}}exit data
+  // CHECK-NEXT: if clause
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+  // CHECK-NEXT: NullStmt
+
+#pragma acc host_data use_device(Global) if(T::BC)
+  ;
+  // CHECK-NEXT: OpenACCHostDataConstruct{{.*}}host_data
+  // CHECK-NEXT: if clause
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+  // CHECK-NEXT: NullStmt
+
+  // Match the instantiation:
+  // CHECK: FunctionDecl{{.*}}TemplFunc{{.*}}implicit_instantiation
+  // CHECK-NEXT: TemplateArgument type 'InstTy'
+  // CHECK-NEXT: RecordType{{.*}} 'InstTy'
+  // CHECK-NEXT: CXXRecord{{.*}} 'InstTy'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCDataConstruct{{.*}}data
+  // CHECK-NEXT: if clause
+  // CHECK-NEXT: BinaryOperator{{.*}} 'bool' '<'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'float' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'float' <IntegralToFloating>
+  // CHECK-NEXT: CXXFunctionalCastExpr{{.*}}'typename InstTy::IntTy':'int' functional cast to typename struct InstTy::IntTy <NoOp>
+  // CHECK-NEXT: InitListExpr {{.*}}'typename InstTy::IntTy':'int'
+  // CHECK-NEXT: NullStmt
+
+  // CHECK-NEXT: OpenACCEnterDataConstruct{{.*}}enter data
+  // CHECK-NEXT: if clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <IntegralToBoolean>
+  // CHECK-NEXT: CXXFunctionalCastExpr{{.*}}'typename InstTy::IntTy':'int' functional cast to typename struct InstTy::IntTy <NoOp>
+  // CHECK-NEXT: InitListExpr {{.*}}'typename InstTy::IntTy':'int'
+  // CHECK-NEXT: NullStmt
+
+  // CHECK-NEXT: OpenACCExitDataConstruct{{.*}}exit data
+  // CHECK-NEXT: if clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <FloatingToBoolean>
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'float' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+  // CHECK-NEXT: NullStmt
+
+  // CHECK-NEXT: OpenACCHostDataConstruct{{.*}}host_data
+  // CHECK-NEXT: if clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'bool' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'bool'
+  // CHECK-NEXT: MemberExpr{{.*}} .operator bool
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const BoolConversion' lvalue Var{{.*}} 'BC' 'const BoolConversion'
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+  // CHECK-NEXT: NullStmt
+
+}
+
+struct BoolConversion{ operator bool() const;};
+struct InstTy {
+  using IntTy = int;
+  static constexpr float SomeFloat = 5.0;
+  static constexpr BoolConversion BC;
+};
+
+void Instantiate() {
+  TemplFunc<InstTy>();
+}
+#endif

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

@@ -0,0 +1,37 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+void Foo() {
+  int Var;
+  // expected-warning@+1{{OpenACC clause 'default' not yet implemented}}
+#pragma acc data default(present) if(1)
+  ;
+  // expected-warning@+3{{OpenACC clause 'default' not yet implemented}}
+  // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'data' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc data default(present) if(1) if (2)
+  ;
+
+  // expected-warning@+1{{OpenACC clause 'copyin' not yet implemented}}
+#pragma acc enter data copyin(Var) if(1)
+
+  // expected-warning@+3{{OpenACC clause 'copyin' not yet implemented}}
+  // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'enter data' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc enter data copyin(Var) if(1) if (2)
+
+  // expected-warning@+1{{OpenACC clause 'copyout' not yet implemented}}
+#pragma acc exit data copyout(Var) if(1)
+  // expected-warning@+3{{OpenACC clause 'copyout' not yet implemented}}
+  // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'exit data' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc exit data copyout(Var) if(1) if (2)
+
+  // expected-warning@+1{{OpenACC clause 'use_device' not yet implemented}}
+#pragma acc host_data use_device(Var) if(1)
+  ;
+  // expected-warning@+3{{OpenACC clause 'use_device' not yet implemented}}
+  // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'host_data' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc host_data use_device(Var) if(1) if (2)
+  ;
+}

clang/test/SemaOpenACC/data-construct.cpp

@@ -52,7 +52,6 @@ void AtLeastOneOf() {
   // OpenACC TODO: The following 'data' directives should diagnose, since they
   // don't have at least one of the above clauses.
 
-  // expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
 #pragma acc data if(Var)
   ;
 
@@ -80,7 +79,6 @@ void AtLeastOneOf() {
   // OpenACC TODO: The following 'enter data' directives should diagnose, since
   // they don't have at least one of the above clauses.
 
-  // expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
 #pragma acc enter data if(Var)
   // expected-warning@+1{{OpenACC clause 'async' not yet implemented}}
 #pragma acc enter data async
@@ -99,7 +97,6 @@ void AtLeastOneOf() {
   // OpenACC TODO: The following 'exit data' directives should diagnose, since
   // they don't have at least one of the above clauses.
 
-  // expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
 #pragma acc exit data if(Var)
   // expected-warning@+1{{OpenACC clause 'async' not yet implemented}}
 #pragma acc exit data async
@@ -116,7 +113,6 @@ void AtLeastOneOf() {
   // OpenACC TODO: The following 'host_data' directives should diagnose, since
   // they don't have at least one of the above clauses.
 
-  // expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
 #pragma acc host_data if(Var)
   ;
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
@@ -200,10 +196,8 @@ struct HasMembers {
 
 void HostDataRules() {
   int Var, Var2;
-  // TODO OpenACC: The following line should diagnose, since only 1 'if' is
-  // allowed per directive on host_data.
-  // expected-warning@+2{{OpenACC clause 'if' not yet implemented}}
-  // expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
+  // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'host_data' directive}}
+  // expected-note@+1{{previous clause is here}}
 #pragma acc host_data if(Var) if (Var2)
   ;