[OpenACC] Enable 'deviceptr' clause sema on data construct

Another simple implementation, as it uses the same work as the previous
implementation, just enabling it for this construct.

clang/lib/Sema/SemaOpenACC.cpp

@@ -1035,11 +1035,13 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitAttachClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitDevicePtrClause(
     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
+  // construct 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();
 
   // ActOnVar ensured that everything is a valid variable reference, but we

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

@@ -101,4 +101,10 @@ void foo() {
 
 // CHECK: #pragma acc enter data create(i, array[1], array, array[1:2]) pcreate(zero: i, array[1], array, array[1:2]) present_or_create(i, array[1], array, array[1:2])
 #pragma acc enter data create(i, array[1], array, array[1:2]) pcreate(zero: i, array[1], array, array[1:2]) present_or_create(i, array[1], array, array[1:2])
+
+  float *arrayPtr[5];
+
+// CHECK: #pragma acc data default(none) deviceptr(iPtr, arrayPtr[0])
+#pragma acc data default(none) deviceptr(iPtr, arrayPtr[0])
+  ;
 }

clang/test/SemaOpenACC/data-construct-deviceptr-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) deviceptr(PointerParam)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCDataConstruct{{.*}} data
+  // CHECK-NEXT: default(present)
+  // CHECK-NEXT: deviceptr 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) deviceptr(t)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCDataConstruct{{.*}} data
+  // CHECK-NEXT: default(present)
+  // CHECK-NEXT: deviceptr 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: deviceptr 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-deviceptr-clause.c

@@ -0,0 +1,70 @@
+// 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 'deviceptr' clause, type is 'int'}}
+#pragma acc data default(none) deviceptr(LocalInt)
+  for (int i = 0; i < 5; ++i);
+
+  // 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) deviceptr(&LocalInt)
+  for (int i = 0; i < 5; ++i);
+
+
+  // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}}
+#pragma acc kernels loop deviceptr(Array)
+  for (int i = 0; i < 5; ++i);
+
+  // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc data default(none) deviceptr(Array[0])
+  for (int i = 0; i < 5; ++i);
+
+  // expected-error@+2{{OpenACC sub-array is not allowed here}}
+  // expected-note@+1{{expected variable of pointer type}}
+#pragma acc data default(none) deviceptr(Array[0:1])
+  for (int i = 0; i < 5; ++i);
+
+  // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}}
+#pragma acc data default(none) deviceptr(PtrArray)
+  for (int i = 0; i < 5; ++i);
+
+#pragma acc data default(none) deviceptr(PtrArray[0])
+  for (int i = 0; i < 5; ++i);
+
+  // expected-error@+2{{OpenACC sub-array is not allowed here}}
+  // expected-note@+1{{expected variable of pointer type}}
+#pragma acc data default(none) deviceptr(PtrArray[0:1])
+  for (int i = 0; i < 5; ++i);
+
+  // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'struct S'}}
+#pragma acc data default(none) deviceptr(s)
+  for (int i = 0; i < 5; ++i);
+
+  // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc data default(none) deviceptr(s.IntMem)
+  for (int i = 0; i < 5; ++i);
+
+#pragma acc data default(none) deviceptr(s.PtrMem)
+  for (int i = 0; i < 5; ++i);
+
+  // expected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'enter data' directive}}
+#pragma acc enter data copyin(LocalInt) deviceptr(LocalInt)
+  for(int i = 5; i < 10;++i);
+  // expected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'exit data' directive}}
+#pragma acc exit data copyout(LocalInt) deviceptr(LocalInt)
+  for(int i = 5; i < 10;++i);
+  // expected-warning@+2{{OpenACC clause 'use_device' not yet implemented}}
+  // expected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'host_data' directive}}
+#pragma acc host_data use_device(LocalInt) deviceptr(LocalInt)
+  for(int i = 5; i < 10;++i);
+}

clang/test/SemaOpenACC/data-construct.cpp

@@ -20,6 +20,7 @@ void HasStmt() {
 
 void AtLeastOneOf() {
   int Var;
+  int *VarPtr = &Var;
 // Data
 #pragma acc data copy(Var)
   ;
@@ -33,11 +34,10 @@ void AtLeastOneOf() {
   ;
 #pragma acc data present(Var)
   ;
-  // expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}}
-#pragma acc data deviceptr(Var)
+#pragma acc data deviceptr(VarPtr)
   ;
   // expected-warning@+1{{OpenACC clause 'attach' not yet implemented}}
-#pragma acc data attach(Var)
+#pragma acc data attach(VarPtr)
   ;
 #pragma acc data default(none)
   ;