Skip to content

Commit 1c1140c

Browse files
erichkeaneerichkeane
committed
[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.
1 parent d19218e commit 1c1140c

File tree

5 files changed

+89
-15
lines changed

5 files changed

+89
-15
lines changed

‎clang/lib/Sema/SemaOpenACCClause.cpp‎

Lines changed: 15 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -509,11 +509,6 @@ bool checkAlreadyHasClauseOfKind(
509509
bool checkValidAfterDeviceType(
510510
SemaOpenACC &S, const OpenACCDeviceTypeClause &DeviceTypeClause,
511511
const SemaOpenACC::OpenACCParsedClause &NewClause) {
512-
// This is implemented for everything but 'routine', so treat as 'fine' for
513-
// that.
514-
if (NewClause.getDirectiveKind() == OpenACCDirectiveKind::Routine)
515-
return false;
516-
517512
// OpenACC3.3: Section 2.4: Clauses that precede any device_type clause are
518513
// default clauses. Clauses that follow a device_type clause up to the end of
519514
// the directive or up to the next device_type clause are device-specific
@@ -601,6 +596,19 @@ bool checkValidAfterDeviceType(
601596
default:
602597
break;
603598
}
599+
} else if (NewClause.getDirectiveKind() == OpenACCDirectiveKind::Routine) {
600+
// OpenACC 3.3 section 2.15: Only the 'gang', 'worker', 'vector', 'seq', and
601+
// 'bind' clauses may follow a device_type clause.
602+
switch (NewClause.getClauseKind()) {
603+
case OpenACCClauseKind::Gang:
604+
case OpenACCClauseKind::Worker:
605+
case OpenACCClauseKind::Vector:
606+
case OpenACCClauseKind::Seq:
607+
case OpenACCClauseKind::Bind:
608+
return false;
609+
default:
610+
break;
611+
}
604612
}
605613
S.Diag(NewClause.getBeginLoc(), diag::err_acc_clause_after_device_type)
606614
<< NewClause.getClauseKind() << DeviceTypeClause.getClauseKind()
@@ -1256,10 +1264,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWaitClause(
12561264

12571265
OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceTypeClause(
12581266
SemaOpenACC::OpenACCParsedClause &Clause) {
1259-
// Restrictions implemented properly on everything except 'routine'.
1260-
if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Routine)
1261-
return isNotImplemented();
1262-
12631267
// OpenACC 3.3 2.14.3: Two instances of the same clause may not appear on the
12641268
// same directive.
12651269
if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Set &&
@@ -2056,11 +2060,8 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
20562060
return nullptr;
20572061
}
20582062

2059-
if (const auto *DevTypeClause =
2060-
llvm::find_if(ExistingClauses,
2061-
[&](const OpenACCClause *C) {
2062-
return isa<OpenACCDeviceTypeClause>(C);
2063-
});
2063+
if (const auto *DevTypeClause = llvm::find_if(
2064+
ExistingClauses, llvm::IsaPred<OpenACCDeviceTypeClause>);
20642065
DevTypeClause != ExistingClauses.end()) {
20652066
if (checkValidAfterDeviceType(
20662067
*this, *cast<OpenACCDeviceTypeClause>(*DevTypeClause), Clause))

‎clang/lib/Sema/SemaTemplateInstantiateDecl.cpp‎

Lines changed: 9 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1056,7 +1056,6 @@ CLAUSE_NOT_ON_DECLS(Delete)
10561056
CLAUSE_NOT_ON_DECLS(Detach)
10571057
CLAUSE_NOT_ON_DECLS(Device)
10581058
CLAUSE_NOT_ON_DECLS(DeviceNum)
1059-
CLAUSE_NOT_ON_DECLS(DeviceType)
10601059
CLAUSE_NOT_ON_DECLS(Finalize)
10611060
CLAUSE_NOT_ON_DECLS(FirstPrivate)
10621061
CLAUSE_NOT_ON_DECLS(Host)
@@ -1116,6 +1115,15 @@ void OpenACCDeclClauseInstantiator::VisitNoHostClause(
11161115
ParsedClause.getEndLoc());
11171116
}
11181117

1118+
void OpenACCDeclClauseInstantiator::VisitDeviceTypeClause(
1119+
const OpenACCDeviceTypeClause &C) {
1120+
// Nothing to transform here, just create a new version of 'C'.
1121+
NewClause = OpenACCDeviceTypeClause::Create(
1122+
SemaRef.getASTContext(), C.getClauseKind(), ParsedClause.getBeginLoc(),
1123+
ParsedClause.getLParenLoc(), C.getArchitectures(),
1124+
ParsedClause.getEndLoc());
1125+
}
1126+
11191127
void OpenACCDeclClauseInstantiator::VisitWorkerClause(
11201128
const OpenACCWorkerClause &C) {
11211129
assert(!C.hasIntExpr() && "Int Expr not allowed on routine 'worker' clause");

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

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,11 @@ int function();
77
// CHECK: #pragma acc routine(function) vector nohost
88
#pragma acc routine (function) vector nohost
99

10+
// CHECK: #pragma acc routine(function) device_type(Something) seq
11+
#pragma acc routine(function) device_type(Something) seq
12+
// CHECK: #pragma acc routine(function) dtype(Something) seq
13+
#pragma acc routine(function) dtype(Something) seq
14+
1015
namespace NS {
1116
int NSFunc();
1217
auto Lambda = [](){};
@@ -70,6 +75,11 @@ struct DepS {
7075
#pragma acc routine(DepS<T>::MemFunc) seq nohost
7176
// CHECK: #pragma acc routine(DepS<T>::StaticMemFunc) nohost worker
7277
#pragma acc routine(DepS<T>::StaticMemFunc) nohost worker
78+
79+
// CHECK: #pragma acc routine(MemFunc) worker dtype(*)
80+
#pragma acc routine (MemFunc) worker dtype(*)
81+
// CHECK: #pragma acc routine(MemFunc) device_type(Lambda) vector
82+
#pragma acc routine (MemFunc) device_type(Lambda) vector
7383
};
7484

7585
// CHECK: #pragma acc routine(DepS<int>::Lambda) gang

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

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19,6 +19,18 @@ int function();
1919
// CHECK-NEXT: nohost clause
2020
// CHECK-NEXT: vector clause
2121

22+
#pragma acc routine(function) device_type(Something) seq
23+
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
24+
// CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()'
25+
// CHECK-NEXT: device_type(Something)
26+
// CHECK-NEXT: seq clause
27+
#pragma acc routine(function) nohost dtype(Something) vector
28+
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
29+
// CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()'
30+
// CHECK-NEXT: nohost clause
31+
// CHECK-NEXT: dtype(Something)
32+
// CHECK-NEXT: vector clause
33+
2234
namespace NS {
2335
// CHECK-NEXT: NamespaceDecl
2436
int NSFunc();
@@ -77,6 +89,16 @@ struct S {
7789
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
7890
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
7991
// CHECK-NEXT: worker clause
92+
#pragma acc routine(Lambda) worker device_type(Lambda)
93+
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
94+
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
95+
// CHECK-NEXT: worker clause
96+
// CHECK-NEXT: device_type(Lambda)
97+
#pragma acc routine(Lambda) dtype(Lambda) vector
98+
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
99+
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
100+
// CHECK-NEXT: dtype(Lambda)
101+
// CHECK-NEXT: vector clause
80102
};
81103

82104
#pragma acc routine(S::MemFunc) gang(dim: 1)
@@ -150,6 +172,12 @@ struct DepS {
150172
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()'
151173
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
152174
// CHECK-NEXT: worker clause
175+
#pragma acc routine(DepS<T>::StaticMemFunc) worker device_type(T)
176+
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
177+
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()'
178+
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
179+
// CHECK-NEXT: worker clause
180+
// CHECK-NEXT: device_type(T)
153181

154182
// Instantiation:
155183
// CHECK: ClassTemplateSpecializationDecl{{.*}}struct DepS
@@ -200,6 +228,12 @@ struct DepS {
200228
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
201229
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
202230
// CHECK-NEXT: worker clause
231+
232+
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
233+
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
234+
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
235+
// CHECK-NEXT: worker clause
236+
// CHECK-NEXT: device_type(T)
203237
};
204238

205239
#pragma acc routine(DepS<int>::Lambda) gang(dim:1)

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

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -121,3 +121,24 @@ struct DependentT {
121121
void Inst() {
122122
DependentT<HasFuncs> T;// expected-note{{in instantiation of}}
123123
}
124+
125+
#pragma acc routine(Func) gang device_type(Inst)
126+
#pragma acc routine(Func) gang dtype(Inst)
127+
#pragma acc routine(Func) device_type(*) worker
128+
#pragma acc routine(Func) dtype(*) worker
129+
130+
#pragma acc routine(Func) dtype(*) gang
131+
#pragma acc routine(Func) device_type(*) worker
132+
#pragma acc routine(Func) device_type(*) vector
133+
#pragma acc routine(Func) dtype(*) seq
134+
// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
135+
#pragma acc routine(Func) seq device_type(*) bind("asdf")
136+
// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
137+
#pragma acc routine(Func) seq device_type(*) bind(getSomeInt)
138+
#pragma acc routine(Func) seq dtype(*) device_type(*)
139+
// expected-error@+2{{OpenACC clause 'nohost' may not follow a 'dtype' clause in a 'routine' construct}}
140+
// expected-note@+1{{previous clause is here}}
141+
#pragma acc routine(Func) seq dtype(*) nohost
142+
// expected-error@+2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a 'routine' construct}}
143+
// expected-note@+1{{previous clause is here}}
144+
#pragma acc routine(Func) seq device_type(*) nohost

0 commit comments

Comments
 (0)