[OpenACC] Enable 'device_type' for 'routine'

There is a slightly different list for routine on which clauses are
permitted after it (like the rest of the constructs), but this
implements and tests them to make sure we get them right.

Author: erichkeane

clang/lib/Sema/SemaOpenACCClause.cpp

@@ -509,11 +509,6 @@ bool checkAlreadyHasClauseOfKind(
 bool checkValidAfterDeviceType(
     SemaOpenACC &S, const OpenACCDeviceTypeClause &DeviceTypeClause,
     const SemaOpenACC::OpenACCParsedClause &NewClause) {
-  // This is implemented for everything but 'routine', so treat as 'fine' for
-  // that.
-  if (NewClause.getDirectiveKind() == OpenACCDirectiveKind::Routine)
-    return false;
-
   // OpenACC3.3: Section 2.4: Clauses that precede any device_type clause are
   // default clauses.  Clauses that follow a device_type clause up to the end of
   // the directive or up to the next device_type clause are device-specific
@@ -601,6 +596,19 @@ bool checkValidAfterDeviceType(
     default:
       break;
     }
+  } else if (NewClause.getDirectiveKind() == OpenACCDirectiveKind::Routine) {
+    // OpenACC 3.3 section 2.15: Only the 'gang', 'worker', 'vector', 'seq', and
+    // 'bind' clauses may follow a device_type clause.
+    switch (NewClause.getClauseKind()) {
+    case OpenACCClauseKind::Gang:
+    case OpenACCClauseKind::Worker:
+    case OpenACCClauseKind::Vector:
+    case OpenACCClauseKind::Seq:
+    case OpenACCClauseKind::Bind:
+      return false;
+    default:
+      break;
+    }
   }
   S.Diag(NewClause.getBeginLoc(), diag::err_acc_clause_after_device_type)
       << NewClause.getClauseKind() << DeviceTypeClause.getClauseKind()
@@ -1256,10 +1264,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWaitClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceTypeClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // Restrictions implemented properly on everything except 'routine'.
-  if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Routine)
-    return isNotImplemented();
-
   // OpenACC 3.3 2.14.3: Two instances of the same clause may not appear on the
   // same directive.
   if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Set &&
@@ -2056,11 +2060,8 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
     return nullptr;
   }
 
-  if (const auto *DevTypeClause =
-          llvm::find_if(ExistingClauses,
-                        [&](const OpenACCClause *C) {
-                          return isa<OpenACCDeviceTypeClause>(C);
-                        });
+  if (const auto *DevTypeClause = llvm::find_if(
+          ExistingClauses, llvm::IsaPred<OpenACCDeviceTypeClause>);
       DevTypeClause != ExistingClauses.end()) {
     if (checkValidAfterDeviceType(
             *this, *cast<OpenACCDeviceTypeClause>(*DevTypeClause), Clause))

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

@@ -1056,7 +1056,6 @@ CLAUSE_NOT_ON_DECLS(Delete)
 CLAUSE_NOT_ON_DECLS(Detach)
 CLAUSE_NOT_ON_DECLS(Device)
 CLAUSE_NOT_ON_DECLS(DeviceNum)
-CLAUSE_NOT_ON_DECLS(DeviceType)
 CLAUSE_NOT_ON_DECLS(Finalize)
 CLAUSE_NOT_ON_DECLS(FirstPrivate)
 CLAUSE_NOT_ON_DECLS(Host)
@@ -1116,6 +1115,15 @@ void OpenACCDeclClauseInstantiator::VisitNoHostClause(
                                           ParsedClause.getEndLoc());
 }
 
+void OpenACCDeclClauseInstantiator::VisitDeviceTypeClause(
+    const OpenACCDeviceTypeClause &C) {
+  // Nothing to transform here, just create a new version of 'C'.
+  NewClause = OpenACCDeviceTypeClause::Create(
+      SemaRef.getASTContext(), C.getClauseKind(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), C.getArchitectures(),
+      ParsedClause.getEndLoc());
+}
+
 void OpenACCDeclClauseInstantiator::VisitWorkerClause(
     const OpenACCWorkerClause &C) {
   assert(!C.hasIntExpr() && "Int Expr not allowed on routine 'worker' clause");

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

@@ -7,6 +7,11 @@ int function();
 // CHECK: #pragma acc routine(function) vector nohost
 #pragma acc routine (function) vector nohost
 
+// CHECK: #pragma acc routine(function) device_type(Something) seq
+#pragma acc routine(function) device_type(Something) seq
+// CHECK: #pragma acc routine(function) dtype(Something) seq
+#pragma acc routine(function) dtype(Something) seq
+
 namespace NS {
   int NSFunc();
 auto Lambda = [](){};
@@ -70,6 +75,11 @@ struct DepS {
 #pragma acc routine(DepS<T>::MemFunc) seq nohost
 // CHECK: #pragma acc routine(DepS<T>::StaticMemFunc) nohost worker
 #pragma acc routine(DepS<T>::StaticMemFunc) nohost worker
+
+// CHECK: #pragma acc routine(MemFunc) worker dtype(*)
+#pragma acc routine (MemFunc) worker dtype(*)
+// CHECK: #pragma acc routine(MemFunc) device_type(Lambda) vector
+#pragma acc routine (MemFunc) device_type(Lambda) vector
 };
 
 // CHECK: #pragma acc routine(DepS<int>::Lambda) gang

clang/test/SemaOpenACC/routine-construct-ast.cpp

@@ -19,6 +19,18 @@ int function();
 // CHECK-NEXT: nohost clause
 // CHECK-NEXT: vector clause
 
+#pragma acc routine(function) device_type(Something) seq
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()'
+// CHECK-NEXT: device_type(Something)
+// CHECK-NEXT: seq clause
+#pragma acc routine(function) nohost dtype(Something) vector
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()'
+// CHECK-NEXT: nohost clause
+// CHECK-NEXT: dtype(Something)
+// CHECK-NEXT: vector clause
+
 namespace NS {
   // CHECK-NEXT: NamespaceDecl
   int NSFunc();
@@ -77,6 +89,16 @@ struct S {
 // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
 // CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
 // CHECK-NEXT: worker clause
+#pragma acc routine(Lambda) worker device_type(Lambda)
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
+// CHECK-NEXT: worker clause
+// CHECK-NEXT: device_type(Lambda)
+#pragma acc routine(Lambda) dtype(Lambda) vector
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
+// CHECK-NEXT: dtype(Lambda)
+// CHECK-NEXT: vector clause
 };
 
 #pragma acc routine(S::MemFunc) gang(dim: 1)
@@ -150,6 +172,12 @@ struct DepS {
 // CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()'
 // CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
 // CHECK-NEXT: worker clause
+#pragma acc routine(DepS<T>::StaticMemFunc) worker device_type(T)
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()'
+// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
+// CHECK-NEXT: worker clause
+// CHECK-NEXT: device_type(T)
 
 // Instantiation:
 // CHECK: ClassTemplateSpecializationDecl{{.*}}struct DepS
@@ -200,6 +228,12 @@ struct DepS {
 // CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
 // CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
 // CHECK-NEXT: worker clause
+
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
+// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
+// CHECK-NEXT: worker clause
+// CHECK-NEXT: device_type(T)
 };
 
 #pragma acc routine(DepS<int>::Lambda) gang(dim:1)

clang/test/SemaOpenACC/routine-construct-clauses.cpp

@@ -121,3 +121,24 @@ struct DependentT {
 void Inst() {
   DependentT<HasFuncs> T;// expected-note{{in instantiation of}}
 }
+
+#pragma acc routine(Func) gang device_type(Inst)
+#pragma acc routine(Func) gang dtype(Inst)
+#pragma acc routine(Func) device_type(*) worker
+#pragma acc routine(Func) dtype(*) worker
+
+#pragma acc routine(Func) dtype(*) gang
+#pragma acc routine(Func) device_type(*) worker
+#pragma acc routine(Func) device_type(*) vector
+#pragma acc routine(Func) dtype(*) seq
+// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
+#pragma acc routine(Func) seq device_type(*) bind("asdf")
+// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
+#pragma acc routine(Func) seq device_type(*) bind(getSomeInt)
+#pragma acc routine(Func) seq dtype(*) device_type(*)
+// expected-error@+2{{OpenACC clause 'nohost' may not follow a 'dtype' clause in a 'routine' construct}}
+// expected-note@+1{{previous clause is here}}
+#pragma acc routine(Func) seq dtype(*) nohost
+// expected-error@+2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a 'routine' construct}}
+// expected-note@+1{{previous clause is here}}
+#pragma acc routine(Func) seq device_type(*) nohost