Skip to content

Commit f10e71f

Browse files
erichkeaneerichkeane
committed
[OpenACC] Implement 'device_type' sema for 'loop' construct
This clause is effectively identical to how this works on compute clauses, however the list of clauses allowed after it are slightly different. This enables the clause for the 'loop', and ensures we're permitting the correct list.
1 parent 7578c31 commit f10e71f

File tree

6 files changed

+443
-24
lines changed

6 files changed

+443
-24
lines changed

‎clang/include/clang/Basic/DiagnosticSemaKinds.td

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -12398,7 +12398,7 @@ def err_acc_var_not_pointer_type
1239812398
def note_acc_expected_pointer_var : Note<"expected variable of pointer type">;
1239912399
def err_acc_clause_after_device_type
1240012400
: Error<"OpenACC clause '%0' may not follow a '%1' clause in a "
12401-
"compute construct">;
12401+
"%select{'%3'|compute}2 construct">;
1240212402
def err_acc_reduction_num_gangs_conflict
1240312403
: Error<
1240412404
"OpenACC 'reduction' clause may not appear on a 'parallel' construct "

‎clang/lib/Sema/SemaOpenACC.cpp

Lines changed: 58 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -341,31 +341,64 @@ bool checkAlreadyHasClauseOfKind(
341341
return false;
342342
}
343343

344-
/// Implement check from OpenACC3.3: section 2.5.4:
345-
/// Only the async, wait, num_gangs, num_workers, and vector_length clauses may
346-
/// follow a device_type clause.
347344
bool checkValidAfterDeviceType(
348345
SemaOpenACC &S, const OpenACCDeviceTypeClause &DeviceTypeClause,
349346
const SemaOpenACC::OpenACCParsedClause &NewClause) {
350-
// This is only a requirement on compute constructs so far, so this is fine
351-
// otherwise.
352-
if (!isOpenACCComputeDirectiveKind(NewClause.getDirectiveKind()))
347+
// This is only a requirement on compute and loop constructs so far, so this
348+
// is fine otherwise.
349+
if (!isOpenACCComputeDirectiveKind(NewClause.getDirectiveKind()) &&
350+
NewClause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
353351
return false;
354-
switch (NewClause.getClauseKind()) {
355-
case OpenACCClauseKind::Async:
356-
case OpenACCClauseKind::Wait:
357-
case OpenACCClauseKind::NumGangs:
358-
case OpenACCClauseKind::NumWorkers:
359-
case OpenACCClauseKind::VectorLength:
360-
case OpenACCClauseKind::DType:
361-
case OpenACCClauseKind::DeviceType:
352+
353+
// OpenACC3.3: Section 2.4: Clauses that precede any device_type clause are
354+
// default clauses. Clauses that follow a device_type clause up to the end of
355+
// the directive or up to the next device_type clause are device-specific
356+
// clauses for the device types specified in the device_type argument.
357+
//
358+
// The above implies that despite what the individual text says, these are
359+
// valid.
360+
if (NewClause.getClauseKind() == OpenACCClauseKind::DType ||
361+
NewClause.getClauseKind() == OpenACCClauseKind::DeviceType)
362362
return false;
363-
default:
364-
S.Diag(NewClause.getBeginLoc(), diag::err_acc_clause_after_device_type)
365-
<< NewClause.getClauseKind() << DeviceTypeClause.getClauseKind();
366-
S.Diag(DeviceTypeClause.getBeginLoc(), diag::note_acc_previous_clause_here);
367-
return true;
363+
364+
// Implement check from OpenACC3.3: section 2.5.4:
365+
// Only the async, wait, num_gangs, num_workers, and vector_length clauses may
366+
// follow a device_type clause.
367+
if (isOpenACCComputeDirectiveKind(NewClause.getDirectiveKind())) {
368+
switch (NewClause.getClauseKind()) {
369+
case OpenACCClauseKind::Async:
370+
case OpenACCClauseKind::Wait:
371+
case OpenACCClauseKind::NumGangs:
372+
case OpenACCClauseKind::NumWorkers:
373+
case OpenACCClauseKind::VectorLength:
374+
return false;
375+
default:
376+
break;
377+
}
378+
} else if (NewClause.getDirectiveKind() == OpenACCDirectiveKind::Loop) {
379+
// Implement check from OpenACC3.3: section 2.9:
380+
// Only the collapse, gang, worker, vector, seq, independent, auto, and tile
381+
// clauses may follow a device_type clause.
382+
switch (NewClause.getClauseKind()) {
383+
case OpenACCClauseKind::Collapse:
384+
case OpenACCClauseKind::Gang:
385+
case OpenACCClauseKind::Worker:
386+
case OpenACCClauseKind::Vector:
387+
case OpenACCClauseKind::Seq:
388+
case OpenACCClauseKind::Independent:
389+
case OpenACCClauseKind::Auto:
390+
case OpenACCClauseKind::Tile:
391+
return false;
392+
default:
393+
break;
394+
}
368395
}
396+
S.Diag(NewClause.getBeginLoc(), diag::err_acc_clause_after_device_type)
397+
<< NewClause.getClauseKind() << DeviceTypeClause.getClauseKind()
398+
<< isOpenACCComputeDirectiveKind(NewClause.getDirectiveKind())
399+
<< NewClause.getDirectiveKind();
400+
S.Diag(DeviceTypeClause.getBeginLoc(), diag::note_acc_previous_clause_here);
401+
return true;
369402
}
370403
} // namespace
371404

@@ -816,10 +849,12 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
816849
}
817850
case OpenACCClauseKind::DType:
818851
case OpenACCClauseKind::DeviceType: {
819-
// Restrictions only properly implemented on 'compute' constructs, and
820-
// 'compute' constructs are the only construct that can do anything with
821-
// this yet, so skip/treat as unimplemented in this case.
822-
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
852+
// Restrictions only properly implemented on 'compute' and 'loop'
853+
// constructs, and 'compute'/'loop' constructs are the only construct that
854+
// can do anything with this yet, so skip/treat as unimplemented in this
855+
// case.
856+
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
857+
Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
823858
break;
824859

825860
// TODO OpenACC: Once we get enough of the CodeGen implemented that we have

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

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,35 @@
11
// RUN: %clang_cc1 -fopenacc -Wno-openacc-deprecated-clause-alias -ast-print %s -o - | FileCheck %s
22

3+
struct SomeStruct{};
4+
35
void foo() {
46
// CHECK: #pragma acc loop
57
// CHECK-NEXT: for (;;)
68
// CHECK-NEXT: ;
79
#pragma acc loop
810
for(;;);
11+
12+
// CHECK: #pragma acc loop device_type(SomeStruct)
13+
// CHECK-NEXT: for (;;)
14+
// CHECK-NEXT: ;
15+
#pragma acc loop device_type(SomeStruct)
16+
for(;;);
17+
18+
// CHECK: #pragma acc loop device_type(int)
19+
// CHECK-NEXT: for (;;)
20+
// CHECK-NEXT: ;
21+
#pragma acc loop device_type(int)
22+
for(;;);
23+
24+
// CHECK: #pragma acc loop dtype(bool)
25+
// CHECK-NEXT: for (;;)
26+
// CHECK-NEXT: ;
27+
#pragma acc loop dtype(bool)
28+
for(;;);
29+
30+
// CHECK: #pragma acc loop dtype(AnotherIdent)
31+
// CHECK-NEXT: for (;;)
32+
// CHECK-NEXT: ;
33+
#pragma acc loop dtype(AnotherIdent)
34+
for(;;);
935
}

‎clang/test/SemaOpenACC/loop-construct-device_type-ast.cpp

Lines changed: 129 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,129 @@
1+
// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
2+
3+
// Test this with PCH.
4+
// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
5+
// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
6+
#ifndef PCH_HELPER
7+
#define PCH_HELPER
8+
9+
struct SomeS{};
10+
void NormalUses() {
11+
// CHECK: FunctionDecl{{.*}}NormalUses
12+
// CHECK-NEXT: CompoundStmt
13+
14+
SomeS SomeImpl;
15+
// CHECK-NEXT: DeclStmt
16+
// CHECK-NEXT: VarDecl{{.*}} SomeImpl 'SomeS'
17+
// CHECK-NEXT: CXXConstructExpr
18+
bool SomeVar;
19+
// CHECK-NEXT: DeclStmt
20+
// CHECK-NEXT: VarDecl{{.*}} SomeVar 'bool'
21+
22+
#pragma acc loop device_type(SomeS) dtype(SomeImpl)
23+
for(;;){}
24+
// CHECK-NEXT: OpenACCLoopConstruct{{.*}}
25+
// CHECK-NEXT: device_type(SomeS)
26+
// CHECK-NEXT: dtype(SomeImpl)
27+
// CHECK-NEXT: ForStmt
28+
// CHECK-NEXT: <<<NULL>>>
29+
// CHECK-NEXT: <<<NULL>>>
30+
// CHECK-NEXT: <<<NULL>>>
31+
// CHECK-NEXT: <<<NULL>>>
32+
// CHECK-NEXT: CompoundStmt
33+
#pragma acc loop device_type(SomeVar) dtype(int)
34+
for(;;){}
35+
// CHECK-NEXT: OpenACCLoopConstruct{{.*}}
36+
// CHECK-NEXT: device_type(SomeVar)
37+
// CHECK-NEXT: dtype(int)
38+
// CHECK-NEXT: ForStmt
39+
// CHECK-NEXT: <<<NULL>>>
40+
// CHECK-NEXT: <<<NULL>>>
41+
// CHECK-NEXT: <<<NULL>>>
42+
// CHECK-NEXT: <<<NULL>>>
43+
// CHECK-NEXT: CompoundStmt
44+
#pragma acc loop device_type(private) dtype(struct)
45+
for(;;){}
46+
// CHECK-NEXT: OpenACCLoopConstruct{{.*}}
47+
// CHECK-NEXT: device_type(private)
48+
// CHECK-NEXT: dtype(struct)
49+
// CHECK-NEXT: ForStmt
50+
// CHECK-NEXT: <<<NULL>>>
51+
// CHECK-NEXT: <<<NULL>>>
52+
// CHECK-NEXT: <<<NULL>>>
53+
// CHECK-NEXT: <<<NULL>>>
54+
// CHECK-NEXT: CompoundStmt
55+
#pragma acc loop device_type(private) dtype(class)
56+
for(;;){}
57+
// CHECK-NEXT: OpenACCLoopConstruct{{.*}}
58+
// CHECK-NEXT: device_type(private)
59+
// CHECK-NEXT: dtype(class)
60+
// CHECK-NEXT: ForStmt
61+
// CHECK-NEXT: <<<NULL>>>
62+
// CHECK-NEXT: <<<NULL>>>
63+
// CHECK-NEXT: <<<NULL>>>
64+
// CHECK-NEXT: <<<NULL>>>
65+
// CHECK-NEXT: CompoundStmt
66+
#pragma acc loop device_type(float) dtype(*)
67+
for(;;){}
68+
// CHECK-NEXT: OpenACCLoopConstruct{{.*}}
69+
// CHECK-NEXT: device_type(float)
70+
// CHECK-NEXT: dtype(*)
71+
// CHECK-NEXT: ForStmt
72+
// CHECK-NEXT: <<<NULL>>>
73+
// CHECK-NEXT: <<<NULL>>>
74+
// CHECK-NEXT: <<<NULL>>>
75+
// CHECK-NEXT: <<<NULL>>>
76+
// CHECK-NEXT: CompoundStmt
77+
#pragma acc loop device_type(float, int) dtype(*)
78+
for(;;){}
79+
// CHECK-NEXT: OpenACCLoopConstruct{{.*}}
80+
// CHECK-NEXT: device_type(float, int)
81+
// CHECK-NEXT: dtype(*)
82+
// CHECK-NEXT: ForStmt
83+
// CHECK-NEXT: <<<NULL>>>
84+
// CHECK-NEXT: <<<NULL>>>
85+
// CHECK-NEXT: <<<NULL>>>
86+
// CHECK-NEXT: <<<NULL>>>
87+
// CHECK-NEXT: CompoundStmt
88+
}
89+
90+
template<typename T>
91+
void TemplUses() {
92+
// CHECK-NEXT: FunctionTemplateDecl{{.*}}TemplUses
93+
// CHECK-NEXT: TemplateTypeParmDecl{{.*}}T
94+
// CHECK-NEXT: FunctionDecl{{.*}}TemplUses
95+
// CHECK-NEXT: CompoundStmt
96+
#pragma acc loop device_type(T) dtype(T)
97+
for(;;){}
98+
// CHECK-NEXT: OpenACCLoopConstruct{{.*}}
99+
// CHECK-NEXT: device_type(T)
100+
// CHECK-NEXT: dtype(T)
101+
// CHECK-NEXT: ForStmt
102+
// CHECK-NEXT: <<<NULL>>>
103+
// CHECK-NEXT: <<<NULL>>>
104+
// CHECK-NEXT: <<<NULL>>>
105+
// CHECK-NEXT: <<<NULL>>>
106+
// CHECK-NEXT: CompoundStmt
107+
108+
109+
// Instantiations
110+
// CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void ()' implicit_instantiation
111+
// CHECK-NEXT: TemplateArgument type 'int'
112+
// CHECK-NEXT: BuiltinType{{.*}} 'int'
113+
// CHECK-NEXT: CompoundStmt
114+
115+
// CHECK-NEXT: OpenACCLoopConstruct{{.*}}
116+
// CHECK-NEXT: device_type(T)
117+
// CHECK-NEXT: dtype(T)
118+
// CHECK-NEXT: ForStmt
119+
// CHECK-NEXT: <<<NULL>>>
120+
// CHECK-NEXT: <<<NULL>>>
121+
// CHECK-NEXT: <<<NULL>>>
122+
// CHECK-NEXT: <<<NULL>>>
123+
// CHECK-NEXT: CompoundStmt
124+
}
125+
126+
void Inst() {
127+
TemplUses<int>();
128+
}
129+
#endif // PCH_HELPER

0 commit comments

Comments
 (0)