Author: erichkeane Date: 2024-06-05T08:01:35-07:00 New Revision: f10e71f6d80719c47f3eed117120e74d9d3858c1
URL: https://github.com/llvm/llvm-project/commit/f10e71f6d80719c47f3eed117120e74d9d3858c1 DIFF: https://github.com/llvm/llvm-project/commit/f10e71f6d80719c47f3eed117120e74d9d3858c1.diff LOG: [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. Added: clang/test/SemaOpenACC/loop-construct-device_type-ast.cpp clang/test/SemaOpenACC/loop-construct-device_type-clause.c clang/test/SemaOpenACC/loop-construct-device_type-clause.cpp Modified: clang/include/clang/Basic/DiagnosticSemaKinds.td clang/lib/Sema/SemaOpenACC.cpp clang/test/AST/ast-print-openacc-loop-construct.cpp Removed: ################################################################################ diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 50332966a7e37..1a117245a325e 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12398,7 +12398,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 " - "compute construct">; + "%select{'%3'|compute}2 construct">; def err_acc_reduction_num_gangs_conflict : Error< "OpenACC 'reduction' clause may not appear on a 'parallel' construct " diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 92da218010c94..1f99543248f27 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -341,31 +341,64 @@ bool checkAlreadyHasClauseOfKind( return false; } -/// Implement check from OpenACC3.3: section 2.5.4: -/// Only the async, wait, num_gangs, num_workers, and vector_length clauses may -/// follow a device_type clause. bool checkValidAfterDeviceType( SemaOpenACC &S, const OpenACCDeviceTypeClause &DeviceTypeClause, const SemaOpenACC::OpenACCParsedClause &NewClause) { - // This is only a requirement on compute constructs so far, so this is fine - // otherwise. - if (!isOpenACCComputeDirectiveKind(NewClause.getDirectiveKind())) + // This is only a requirement on compute and loop constructs so far, so this + // is fine otherwise. + if (!isOpenACCComputeDirectiveKind(NewClause.getDirectiveKind()) && + NewClause.getDirectiveKind() != OpenACCDirectiveKind::Loop) return false; - switch (NewClause.getClauseKind()) { - case OpenACCClauseKind::Async: - case OpenACCClauseKind::Wait: - case OpenACCClauseKind::NumGangs: - case OpenACCClauseKind::NumWorkers: - case OpenACCClauseKind::VectorLength: - case OpenACCClauseKind::DType: - case OpenACCClauseKind::DeviceType: + + // 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 + // clauses for the device types specified in the device_type argument. + // + // The above implies that despite what the individual text says, these are + // valid. + if (NewClause.getClauseKind() == OpenACCClauseKind::DType || + NewClause.getClauseKind() == OpenACCClauseKind::DeviceType) return false; - default: - S.Diag(NewClause.getBeginLoc(), diag::err_acc_clause_after_device_type) - << NewClause.getClauseKind() << DeviceTypeClause.getClauseKind(); - S.Diag(DeviceTypeClause.getBeginLoc(), diag::note_acc_previous_clause_here); - return true; + + // Implement check from OpenACC3.3: section 2.5.4: + // Only the async, wait, num_gangs, num_workers, and vector_length clauses may + // follow a device_type clause. + if (isOpenACCComputeDirectiveKind(NewClause.getDirectiveKind())) { + switch (NewClause.getClauseKind()) { + case OpenACCClauseKind::Async: + case OpenACCClauseKind::Wait: + case OpenACCClauseKind::NumGangs: + case OpenACCClauseKind::NumWorkers: + case OpenACCClauseKind::VectorLength: + return false; + default: + break; + } + } else if (NewClause.getDirectiveKind() == OpenACCDirectiveKind::Loop) { + // Implement check from OpenACC3.3: section 2.9: + // Only the collapse, gang, worker, vector, seq, independent, auto, and tile + // clauses may follow a device_type clause. + switch (NewClause.getClauseKind()) { + case OpenACCClauseKind::Collapse: + case OpenACCClauseKind::Gang: + case OpenACCClauseKind::Worker: + case OpenACCClauseKind::Vector: + case OpenACCClauseKind::Seq: + case OpenACCClauseKind::Independent: + case OpenACCClauseKind::Auto: + case OpenACCClauseKind::Tile: + return false; + default: + break; + } } + 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; } } // namespace @@ -816,10 +849,12 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses, } case OpenACCClauseKind::DType: case OpenACCClauseKind::DeviceType: { - // Restrictions only properly implemented on 'compute' constructs, and - // 'compute' constructs are the only construct that can do anything with - // this yet, so skip/treat as unimplemented in this case. - if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind())) + // 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. + if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) && + Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop) break; // TODO OpenACC: Once we get enough of the CodeGen implemented that we have diff --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp b/clang/test/AST/ast-print-openacc-loop-construct.cpp index 21c92b17317ef..f6d7b54251db0 100644 --- a/clang/test/AST/ast-print-openacc-loop-construct.cpp +++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp @@ -1,9 +1,35 @@ // RUN: %clang_cc1 -fopenacc -Wno-openacc-deprecated-clause-alias -ast-print %s -o - | FileCheck %s +struct SomeStruct{}; + void foo() { // CHECK: #pragma acc loop // CHECK-NEXT: for (;;) // CHECK-NEXT: ; #pragma acc loop for(;;); + +// CHECK: #pragma acc loop device_type(SomeStruct) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc loop device_type(SomeStruct) + for(;;); + +// CHECK: #pragma acc loop device_type(int) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc loop device_type(int) + for(;;); + +// CHECK: #pragma acc loop dtype(bool) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc loop dtype(bool) + for(;;); + +// CHECK: #pragma acc loop dtype(AnotherIdent) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc loop dtype(AnotherIdent) + for(;;); } diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-ast.cpp b/clang/test/SemaOpenACC/loop-construct-device_type-ast.cpp new file mode 100644 index 0000000000000..537eae0908358 --- /dev/null +++ b/clang/test/SemaOpenACC/loop-construct-device_type-ast.cpp @@ -0,0 +1,129 @@ +// 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 loop device_type(SomeS) dtype(SomeImpl) + for(;;){} + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} + // CHECK-NEXT: device_type(SomeS) + // CHECK-NEXT: dtype(SomeImpl) + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: CompoundStmt +#pragma acc loop device_type(SomeVar) dtype(int) + for(;;){} + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} + // CHECK-NEXT: device_type(SomeVar) + // CHECK-NEXT: dtype(int) + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: CompoundStmt +#pragma acc loop device_type(private) dtype(struct) + for(;;){} + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} + // CHECK-NEXT: device_type(private) + // CHECK-NEXT: dtype(struct) + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: CompoundStmt +#pragma acc loop device_type(private) dtype(class) + for(;;){} + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} + // CHECK-NEXT: device_type(private) + // CHECK-NEXT: dtype(class) + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: CompoundStmt +#pragma acc loop device_type(float) dtype(*) + for(;;){} + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} + // CHECK-NEXT: device_type(float) + // CHECK-NEXT: dtype(*) + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: CompoundStmt +#pragma acc loop device_type(float, int) dtype(*) + for(;;){} + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} + // CHECK-NEXT: device_type(float, int) + // CHECK-NEXT: dtype(*) + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // 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 loop device_type(T) dtype(T) + for(;;){} + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} + // CHECK-NEXT: device_type(T) + // CHECK-NEXT: dtype(T) + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // 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: OpenACCLoopConstruct{{.*}} + // CHECK-NEXT: device_type(T) + // CHECK-NEXT: dtype(T) + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: CompoundStmt +} + +void Inst() { + TemplUses<int>(); +} +#endif // PCH_HELPER diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c new file mode 100644 index 0000000000000..3e3a85e4498b7 --- /dev/null +++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c @@ -0,0 +1,204 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +#define MACRO +FOO + +void uses() { + typedef struct S{} STy; + STy SImpl; + +#pragma acc loop device_type(I) + for(;;); +#pragma acc loop device_type(S) dtype(STy) + for(;;); +#pragma acc loop dtype(SImpl) + for(;;); +#pragma acc loop dtype(int) device_type(*) + for(;;); +#pragma acc loop dtype(true) device_type(false) + for(;;); + + // expected-error@+1{{expected identifier}} +#pragma acc loop dtype(int, *) + for(;;); + +#pragma acc loop device_type(I, int) + for(;;); + // expected-error@+2{{expected ','}} + // expected-error@+1{{expected identifier}} +#pragma acc loop dtype(int{}) + for(;;); + // expected-error@+1{{expected identifier}} +#pragma acc loop dtype(5) + for(;;); + // expected-error@+1{{expected identifier}} +#pragma acc loop dtype(MACRO) + for(;;); + + + // Only 'collapse', 'gang', 'worker', 'vector', 'seq', 'independent', 'auto', + // and 'tile' allowed after 'device_type'. + + // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}} +#pragma acc loop device_type(*) vector + for(;;); + + // expected-error@+2{{OpenACC clause 'finalize' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) finalize + for(;;); + // expected-error@+2{{OpenACC clause 'if_present' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) if_present + for(;;); + // expected-warning@+1{{OpenACC clause 'seq' not yet implemented, clause ignored}} +#pragma acc loop device_type(*) seq + for(;;); + // expected-warning@+1{{OpenACC clause 'independent' not yet implemented, clause ignored}} +#pragma acc loop device_type(*) independent + for(;;); + // expected-warning@+1{{OpenACC clause 'auto' not yet implemented, clause ignored}} +#pragma acc loop device_type(*) auto + for(;;); + // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause ignored}} +#pragma acc loop device_type(*) worker + for(;;); + // expected-error@+2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) nohost + for(;;); + // expected-error@+1{{OpenACC 'default' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) default(none) + for(;;); + // expected-error@+1{{OpenACC 'if' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) if(1) + for(;;); + // expected-error@+1{{OpenACC 'self' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) self + for(;;); + + int Var; + int *VarPtr; + // expected-error@+1{{OpenACC 'copy' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) copy(Var) + for(;;); + // expected-error@+1{{OpenACC 'pcopy' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) pcopy(Var) + for(;;); + // expected-error@+1{{OpenACC 'present_or_copy' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) present_or_copy(Var) + for(;;); + // expected-error@+2{{OpenACC clause 'use_device' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) use_device(Var) + for(;;); + // expected-error@+1{{OpenACC 'attach' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) attach(Var) + for(;;); + // expected-error@+2{{OpenACC clause 'delete' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) delete(Var) + for(;;); + // expected-error@+2{{OpenACC clause 'detach' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) detach(Var) + for(;;); + // expected-error@+2{{OpenACC clause 'device' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) device(VarPtr) + for(;;); + // expected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) deviceptr(VarPtr) + for(;;); + // expected-error@+2{{OpenACC clause 'device_resident' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) device_resident(VarPtr) + for(;;); + // expected-error@+1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) firstprivate(Var) + for(;;); + // expected-error@+2{{OpenACC clause 'host' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) host(Var) + for(;;); + // expected-error@+2{{OpenACC clause 'link' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) link(Var) + for(;;); + // expected-error@+1{{OpenACC 'no_create' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) no_create(Var) + for(;;); + // expected-error@+1{{OpenACC 'present' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) present(Var) + for(;;); + // expected-error@+2{{OpenACC clause 'private' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) private(Var) + for(;;); + // expected-error@+1{{OpenACC 'copyout' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) copyout(Var) + for(;;); + // expected-error@+1{{OpenACC 'pcopyout' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) pcopyout(Var) + for(;;); + // expected-error@+1{{OpenACC 'present_or_copyout' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) present_or_copyout(Var) + for(;;); + // expected-error@+1{{OpenACC 'copyin' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) copyin(Var) + for(;;); + // expected-error@+1{{OpenACC 'pcopyin' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) pcopyin(Var) + for(;;); + // expected-error@+1{{OpenACC 'present_or_copyin' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) present_or_copyin(Var) + for(;;); + // expected-error@+1{{OpenACC 'create' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) create(Var) + for(;;); + // expected-error@+1{{OpenACC 'pcreate' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) pcreate(Var) + for(;;); + // expected-error@+1{{OpenACC 'present_or_create' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) present_or_create(Var) + for(;;); + // expected-error@+2{{OpenACC clause 'reduction' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) reduction(+:Var) + for(;;); + // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented, clause ignored}} +#pragma acc loop device_type(*) collapse(1) + for(;;); + // expected-error@+2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) bind(Var) + for(;;); + // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) vector_length(1) + for(;;); + // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) num_gangs(1) + for(;;); + // expected-error@+1{{OpenACC 'num_workers' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) num_workers(1) + for(;;); + // expected-error@+2{{OpenACC clause 'device_num' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) device_num(1) + for(;;); + // expected-error@+2{{OpenACC clause 'default_async' may not follow a 'device_type' clause in a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop device_type(*) default_async(1) + for(;;); + // expected-error@+1{{OpenACC 'async' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) async + for(;;); + // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}} +#pragma acc loop device_type(*) tile(Var, 1) + for(;;); + // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} +#pragma acc loop dtype(*) gang + for(;;); + // expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop' directive}} +#pragma acc loop device_type(*) wait + for(;;); +} diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.cpp b/clang/test/SemaOpenACC/loop-construct-device_type-clause.cpp new file mode 100644 index 0000000000000..b8d8d6de4dcf6 --- /dev/null +++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.cpp @@ -0,0 +1,25 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +template<typename T> +void TemplUses() { +#pragma acc loop device_type(I) + for(;;); +#pragma acc loop dtype(*) + for(;;); +#pragma acc loop device_type(class) + for(;;); +#pragma acc loop device_type(private) + for(;;); +#pragma acc loop device_type(bool) + for(;;); +#pragma acc kernels dtype(true) device_type(false) + for(;;); + // expected-error@+2{{expected ','}} + // expected-error@+1{{expected identifier}} +#pragma acc loop device_type(T::value) + for(;;); +} + +void Inst() { + TemplUses<int>(); // #INST +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
