Skip to content

Commit

Permalink
[OpenACC] implement 'device_type' for combined constructs
Browse files Browse the repository at this point in the history 
This clause is pretty small/doesn't do much semantic-analysis-wise, , other than
have two spellings and disallow certain clauses after it. However, as
most of those aren't implemented yet, the diagnostic is left as a TODO.
  • Loading branch information
@erichkeane
erichkeane committed Nov 13, 2024
1 parent 67fb268 commit 6b2de10
Show file tree
Hide file tree
Showing 11 changed files with 514 additions and 57 deletions.
2 changes: 1 addition & 1 deletion clang/include/clang/Basic/DiagnosticSemaKinds.td
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12686,7 +12686,7 @@ def err_acc_var_not_pointer_type
def note_acc_expected_pointer_var : Note<"expected variable of pointer type">;
def err_acc_clause_after_device_type
: Error<"OpenACC clause '%0' may not follow a '%1' clause in a "
"%select{'%3'|compute}2 construct">;
"'%2' construct">;
def err_acc_clause_cannot_combine
: Error<"OpenACC clause '%0' may not appear on the same construct as a "
"'%1' clause on a '%2' construct">;
Expand Down
12 changes: 6 additions & 6 deletions clang/lib/Sema/SemaOpenACC.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -502,7 +502,6 @@ bool checkValidAfterDeviceType(
}
S.Diag(NewClause.getBeginLoc(), diag::err_acc_clause_after_device_type)
<< NewClause.getClauseKind() << DeviceTypeClause.getClauseKind()
<< isOpenACCComputeDirectiveKind(NewClause.getDirectiveKind())
<< NewClause.getDirectiveKind();
S.Diag(DeviceTypeClause.getBeginLoc(), diag::note_acc_previous_clause_here);
return true;
Expand Down Expand Up @@ -999,12 +998,13 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWaitClause(

OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceTypeClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
// Restrictions only properly implemented on 'compute' and 'loop'
// constructs, and 'compute'/'loop' 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', and
// 'loop' constructs, and 'compute'/'combined'/'loop' constructs are the only
// construct that can do anything with this yet, so skip/treat as
// unimplemented in this case.
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop &&
!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
return isNotImplemented();

// TODO OpenACC: Once we get enough of the CodeGen implemented that we have
Expand Down
41 changes: 41 additions & 0 deletions clang/test/AST/ast-print-openacc-combined-construct.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -31,4 +31,45 @@ void foo() {
// CHECK-NEXT: ;
#pragma acc kernels loop independent
for(int i = 0;i<5;++i);

bool SomeB;
struct SomeStruct{} SomeStructImpl;

//CHECK: #pragma acc parallel loop dtype(SomeB)
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
// CHECK-NEXT: ;
#pragma acc parallel loop dtype(SomeB)
for(int i = 0;i<5;++i);

//CHECK: #pragma acc serial loop device_type(SomeStruct)
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
// CHECK-NEXT: ;
#pragma acc serial loop device_type(SomeStruct)
for(int i = 0;i<5;++i);

//CHECK: #pragma acc kernels loop device_type(int)
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
// CHECK-NEXT: ;
#pragma acc kernels loop device_type(int)
for(int i = 0;i<5;++i);

//CHECK: #pragma acc parallel loop dtype(bool)
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
// CHECK-NEXT: ;
#pragma acc parallel loop dtype(bool)
for(int i = 0;i<5;++i);

//CHECK: #pragma acc serial loop device_type(SomeStructImpl)
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
// CHECK-NEXT: ;
#pragma acc serial loop device_type (SomeStructImpl)
for(int i = 0;i<5;++i);

// CHECK: #pragma acc kernels loop dtype(AnotherIdent)
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
// CHECK-NEXT: ;
#pragma acc kernels loop dtype(AnotherIdent)
for(int i = 0;i<5;++i);


}
12 changes: 0 additions & 12 deletions clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -195,10 +195,8 @@ void uses() {
// expected-warning@+1{{OpenACC clause 'default_async' not yet implemented}}
#pragma acc parallel loop auto default_async(1)
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'device_type' not yet implemented}}
#pragma acc parallel loop auto device_type(*)
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'dtype' not yet implemented}}
#pragma acc parallel loop auto dtype(*)
for(unsigned i = 0; i < 5; ++i);
// TODOexpected-error@+1{{OpenACC 'async' clause is not valid on 'parallel loop' directive}}
Expand Down Expand Up @@ -371,10 +369,8 @@ void uses() {
// expected-warning@+1{{OpenACC clause 'default_async' not yet implemented}}
#pragma acc parallel loop default_async(1) auto
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'device_type' not yet implemented}}
#pragma acc parallel loop device_type(*) auto
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'dtype' not yet implemented}}
#pragma acc parallel loop dtype(*) auto
for(unsigned i = 0; i < 5; ++i);
// TODOexpected-error@+1{{OpenACC 'async' clause is not valid on 'parallel loop' directive}}
Expand Down Expand Up @@ -548,10 +544,8 @@ void uses() {
// expected-warning@+1{{OpenACC clause 'default_async' not yet implemented}}
#pragma acc parallel loop independent default_async(1)
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'device_type' not yet implemented}}
#pragma acc parallel loop independent device_type(*)
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'dtype' not yet implemented}}
#pragma acc parallel loop independent dtype(*)
for(unsigned i = 0; i < 5; ++i);
// TODOexpected-error@+1{{OpenACC 'async' clause is not valid on 'parallel loop' directive}}
Expand Down Expand Up @@ -724,10 +718,8 @@ void uses() {
// expected-warning@+1{{OpenACC clause 'default_async' not yet implemented}}
#pragma acc parallel loop default_async(1) independent
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'device_type' not yet implemented}}
#pragma acc parallel loop device_type(*) independent
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'dtype' not yet implemented}}
#pragma acc parallel loop dtype(*) independent
for(unsigned i = 0; i < 5; ++i);
// TODOexpected-error@+1{{OpenACC 'async' clause is not valid on 'parallel loop' directive}}
Expand Down Expand Up @@ -907,10 +899,8 @@ void uses() {
// expected-warning@+1{{OpenACC clause 'default_async' not yet implemented}}
#pragma acc parallel loop seq default_async(1)
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'device_type' not yet implemented}}
#pragma acc parallel loop seq device_type(*)
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'dtype' not yet implemented}}
#pragma acc parallel loop seq dtype(*)
for(unsigned i = 0; i < 5; ++i);
// TODOexpected-error@+1{{OpenACC 'async' clause is not valid on 'parallel loop' directive}}
Expand Down Expand Up @@ -1089,10 +1079,8 @@ void uses() {
// expected-warning@+1{{OpenACC clause 'default_async' not yet implemented}}
#pragma acc parallel loop default_async(1) seq
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'device_type' not yet implemented}}
#pragma acc parallel loop device_type(*) seq
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'dtype' not yet implemented}}
#pragma acc parallel loop dtype(*) seq
for(unsigned i = 0; i < 5; ++i);
// TODOexpected-error@+1{{OpenACC 'async' clause is not valid on 'parallel loop' directive}}
Expand Down
178 changes: 178 additions & 0 deletions clang/test/SemaOpenACC/combined-construct-device_type-ast.cpp
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,178 @@
// 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

struct SomeS{};
void NormalUses() {
// CHECK: FunctionDecl{{.*}}NormalUses
// CHECK-NEXT: CompoundStmt

SomeS SomeImpl;
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl{{.*}} SomeImpl 'SomeS'
// CHECK-NEXT: CXXConstructExpr
bool SomeVar;
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl{{.*}} SomeVar 'bool'

#pragma acc parallel loop device_type(SomeS) dtype(SomeImpl)
for(int i = 0; i < 5; ++i){}
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
// CHECK-NEXT: device_type(SomeS)
// CHECK-NEXT: dtype(SomeImpl)
// CHECK-NEXT: ForStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl{{.*}} i 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
// CHECK-NEXT: <<<NULL>>>
// CHECK-NEXT: BinaryOperator{{.*}}'<'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
// CHECK-NEXT: UnaryOperator{{.*}}++
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: CompoundStmt
#pragma acc serial loop device_type(SomeVar) dtype(int)
for(int i = 0; i < 5; ++i){}
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
// CHECK-NEXT: device_type(SomeVar)
// CHECK-NEXT: dtype(int)
// CHECK-NEXT: ForStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl{{.*}} i 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
// CHECK-NEXT: <<<NULL>>>
// CHECK-NEXT: BinaryOperator{{.*}}'<'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
// CHECK-NEXT: UnaryOperator{{.*}}++
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: CompoundStmt
#pragma acc kernels loop device_type(private) dtype(struct)
for(int i = 0; i < 5; ++i){}
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
// CHECK-NEXT: device_type(private)
// CHECK-NEXT: dtype(struct)
// CHECK-NEXT: ForStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl{{.*}} i 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
// CHECK-NEXT: <<<NULL>>>
// CHECK-NEXT: BinaryOperator{{.*}}'<'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
// CHECK-NEXT: UnaryOperator{{.*}}++
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: CompoundStmt
#pragma acc parallel loop device_type(private) dtype(class)
for(int i = 0; i < 5; ++i){}
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
// CHECK-NEXT: device_type(private)
// CHECK-NEXT: dtype(class)
// CHECK-NEXT: ForStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl{{.*}} i 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
// CHECK-NEXT: <<<NULL>>>
// CHECK-NEXT: BinaryOperator{{.*}}'<'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
// CHECK-NEXT: UnaryOperator{{.*}}++
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: CompoundStmt
#pragma acc serial loop device_type(float) dtype(*)
for(int i = 0; i < 5; ++i){}
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
// CHECK-NEXT: device_type(float)
// CHECK-NEXT: dtype(*)
// CHECK-NEXT: ForStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl{{.*}} i 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
// CHECK-NEXT: <<<NULL>>>
// CHECK-NEXT: BinaryOperator{{.*}}'<'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
// CHECK-NEXT: UnaryOperator{{.*}}++
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: CompoundStmt
#pragma acc kernels loop device_type(float, int) dtype(*)
for(int i = 0; i < 5; ++i){}
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
// CHECK-NEXT: device_type(float, int)
// CHECK-NEXT: dtype(*)
// CHECK-NEXT: ForStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl{{.*}} i 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
// CHECK-NEXT: <<<NULL>>>
// CHECK-NEXT: BinaryOperator{{.*}}'<'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
// CHECK-NEXT: UnaryOperator{{.*}}++
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: CompoundStmt
}

template<typename T>
void TemplUses() {
// CHECK-NEXT: FunctionTemplateDecl{{.*}}TemplUses
// CHECK-NEXT: TemplateTypeParmDecl{{.*}}T
// CHECK-NEXT: FunctionDecl{{.*}}TemplUses
// CHECK-NEXT: CompoundStmt
#pragma acc parallel loop device_type(T) dtype(T)
for(int i = 0; i < 5; ++i){}
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
// CHECK-NEXT: device_type(T)
// CHECK-NEXT: dtype(T)
// CHECK-NEXT: ForStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl{{.*}} i 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
// CHECK-NEXT: <<<NULL>>>
// CHECK-NEXT: BinaryOperator{{.*}}'<'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
// CHECK-NEXT: UnaryOperator{{.*}}++
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: CompoundStmt


// Instantiations
// CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void ()' implicit_instantiation
// CHECK-NEXT: TemplateArgument type 'int'
// CHECK-NEXT: BuiltinType{{.*}} 'int'
// CHECK-NEXT: CompoundStmt

// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
// CHECK-NEXT: device_type(T)
// CHECK-NEXT: dtype(T)
// CHECK-NEXT: ForStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl{{.*}} i 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
// CHECK-NEXT: <<<NULL>>>
// CHECK-NEXT: BinaryOperator{{.*}}'<'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
// CHECK-NEXT: UnaryOperator{{.*}}++
// CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
// CHECK-NEXT: CompoundStmt
}

void Inst() {
TemplUses<int>();
}

#endif // PCH_HELPER
Loading

0 comments on commit 6b2de10

Please sign in to comment.