[clang] 1c1140c - [OpenACC] Enable 'device_type' for 'routine'
via cfe-commits
cfe-commits at lists.llvm.org
Fri Mar 7 06:49:39 PST 2025
Author: erichkeane
Date: 2025-03-07T06:49:29-08:00
New Revision: 1c1140c4cf7d4fde5de4b20c6b993bab9060d9fc
URL: https://github.com/llvm/llvm-project/commit/1c1140c4cf7d4fde5de4b20c6b993bab9060d9fc
DIFF: https://github.com/llvm/llvm-project/commit/1c1140c4cf7d4fde5de4b20c6b993bab9060d9fc.diff
LOG: [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.
Added:
Modified:
clang/lib/Sema/SemaOpenACCClause.cpp
clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
clang/test/AST/ast-print-openacc-routine-construct.cpp
clang/test/SemaOpenACC/routine-construct-ast.cpp
clang/test/SemaOpenACC/routine-construct-clauses.cpp
Removed:
################################################################################
diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp
index 582681f247b31..1d70178d03611 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/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))
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index f3e33078c55c7..cb434dcf6eb9a 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/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");
diff --git a/clang/test/AST/ast-print-openacc-routine-construct.cpp b/clang/test/AST/ast-print-openacc-routine-construct.cpp
index 739718724f179..df41c61142d37 100644
--- a/clang/test/AST/ast-print-openacc-routine-construct.cpp
+++ b/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
diff --git a/clang/test/SemaOpenACC/routine-construct-ast.cpp b/clang/test/SemaOpenACC/routine-construct-ast.cpp
index 8fe2c3695826f..f9510eae31590 100644
--- a/clang/test/SemaOpenACC/routine-construct-ast.cpp
+++ b/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)
diff --git a/clang/test/SemaOpenACC/routine-construct-clauses.cpp b/clang/test/SemaOpenACC/routine-construct-clauses.cpp
index ac9289b798a3e..5fea3d734a410 100644
--- a/clang/test/SemaOpenACC/routine-construct-clauses.cpp
+++ b/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 at +1{{OpenACC clause 'bind' not yet implemented}}
+#pragma acc routine(Func) seq device_type(*) bind("asdf")
+// expected-warning at +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 at +2{{OpenACC clause 'nohost' may not follow a 'dtype' clause in a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) seq dtype(*) nohost
+// expected-error at +2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) seq device_type(*) nohost
More information about the cfe-commits
mailing list