[OpenACC] enable 'attach' clause sema for 'data' and 'enter data'

This is very similar to deviceptr, and is the same implementation as for
combined/compute constructs, so this just enables that, and adds tests.

Author: erichkeane

clang/lib/Sema/SemaOpenACC.cpp

@@ -1013,13 +1013,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitCreateClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitAttachClause(
     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.
-  if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
-      !isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
-    return isNotImplemented();
-
   // ActOnVar ensured that everything is a valid variable reference, but we
   // still have to make sure it is a pointer type.
   llvm::SmallVector<Expr *> VarList{Clause.getVarList()};

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

@@ -106,5 +106,8 @@ void foo() {
 
 // CHECK: #pragma acc data default(none) deviceptr(iPtr, arrayPtr[0])
 #pragma acc data default(none) deviceptr(iPtr, arrayPtr[0])
+
+// CHECK: #pragma acc data default(none) attach(iPtr, arrayPtr[0])
+#pragma acc data default(none) attach(iPtr, arrayPtr[0])
   ;
 }

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

@@ -0,0 +1,66 @@
+// 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
+
+int Global;
+short GlobalArray[5];
+
+void NormalUses(float *PointerParam) {
+  // CHECK: FunctionDecl{{.*}}NormalUses
+  // CHECK: ParmVarDecl
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc data default(present) attach(PointerParam)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCDataConstruct{{.*}} data
+  // CHECK-NEXT: default(present)
+  // CHECK-NEXT: attach clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}} 'PointerParam' 'float *'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+}
+
+template<typename T>
+void TemplUses(T *t) {
+  // CHECK-NEXT: FunctionTemplateDecl
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (T *)'
+  // CHECK-NEXT: ParmVarDecl{{.*}} referenced t 'T *'
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc data default(present) attach(t)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCDataConstruct{{.*}} data
+  // CHECK-NEXT: default(present)
+  // CHECK-NEXT: attach clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'T *' lvalue ParmVar{{.*}} 't' 'T *'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+
+  // Check the instantiated versions of the above.
+  // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (int *)' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument type 'int'
+  // CHECK-NEXT: BuiltinType{{.*}} 'int'
+  // CHECK-NEXT: ParmVarDecl{{.*}} used t 'int *'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCDataConstruct{{.*}} data
+  // CHECK-NEXT: default(present)
+  // CHECK-NEXT: attach clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 't' 'int *'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+}
+
+void Inst() {
+  int i;
+  TemplUses(&i);
+}
+#endif

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

@@ -0,0 +1,65 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct S {
+  int IntMem;
+  int *PtrMem;
+};
+
+void uses() {
+  int LocalInt;
+  int *LocalPtr;
+  int Array[5];
+  int *PtrArray[5];
+  struct S s;
+
+  // expected-error@+1{{expected pointer in 'attach' clause, type is 'int'}}
+#pragma acc data default(none) attach(LocalInt)
+  ;
+
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+#pragma acc data default(none) attach(&LocalInt)
+  ;
+
+
+  // expected-error@+1{{expected pointer in 'attach' clause, type is 'int[5]'}}
+#pragma acc enter data copyin(LocalInt) attach(Array)
+
+  // expected-error@+1{{expected pointer in 'attach' clause, type is 'int'}}
+#pragma acc data default(none) attach(Array[0])
+  ;
+
+  // expected-error@+2{{OpenACC sub-array is not allowed here}}
+  // expected-note@+1{{expected variable of pointer type}}
+#pragma acc data default(none) attach(Array[0:1])
+  ;
+
+  // expected-error@+1{{expected pointer in 'attach' clause, type is 'int *[5]'}}
+#pragma acc data default(none) attach(PtrArray)
+  ;
+
+#pragma acc data default(none) attach(PtrArray[0])
+  ;
+
+  // expected-error@+2{{OpenACC sub-array is not allowed here}}
+  // expected-note@+1{{expected variable of pointer type}}
+#pragma acc data default(none) attach(PtrArray[0:1])
+  ;
+
+  // expected-error@+1{{expected pointer in 'attach' clause, type is 'struct S'}}
+#pragma acc data default(none) attach(s)
+  ;
+
+  // expected-error@+1{{expected pointer in 'attach' clause, type is 'int'}}
+#pragma acc data default(none) attach(s.IntMem)
+  ;
+
+#pragma acc data default(none) attach(s.PtrMem)
+  ;
+
+  // expected-error@+1{{OpenACC 'attach' clause is not valid on 'exit data' directive}}
+#pragma acc exit data copyout(LocalInt) attach(PtrArray[0])
+  // expected-warning@+2{{OpenACC clause 'use_device' not yet implemented}}
+  // expected-error@+1{{OpenACC 'attach' clause is not valid on 'host_data' directive}}
+#pragma acc host_data use_device(LocalInt) attach(PtrArray[0])
+  ;
+}

clang/test/SemaOpenACC/data-construct.cpp

@@ -36,7 +36,6 @@ void AtLeastOneOf() {
   ;
 #pragma acc data deviceptr(VarPtr)
   ;
-  // expected-warning@+1{{OpenACC clause 'attach' not yet implemented}}
 #pragma acc data attach(VarPtr)
   ;
 #pragma acc data default(none)
@@ -62,8 +61,7 @@ void AtLeastOneOf() {
   // Enter Data
 #pragma acc enter data copyin(Var)
 #pragma acc enter data create(Var)
-  // expected-warning@+1{{OpenACC clause 'attach' not yet implemented}}
-#pragma acc enter data attach(Var)
+#pragma acc enter data attach(VarPtr)
 
   // OpenACC TODO: The following 'enter data' directives should diagnose, since
   // they don't have at least one of the above clauses.