[clang] f10e71f - [OpenACC] Implement 'device_type' sema for 'loop' construct
via cfe-commits
cfe-commits at lists.llvm.org
Wed Jun 5 08:01:41 PDT 2024
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 at +1{{expected identifier}}
+#pragma acc loop dtype(int, *)
+ for(;;);
+
+#pragma acc loop device_type(I, int)
+ for(;;);
+ // expected-error at +2{{expected ','}}
+ // expected-error at +1{{expected identifier}}
+#pragma acc loop dtype(int{})
+ for(;;);
+ // expected-error at +1{{expected identifier}}
+#pragma acc loop dtype(5)
+ for(;;);
+ // expected-error at +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 at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
+#pragma acc loop device_type(*) vector
+ for(;;);
+
+ // expected-error at +2{{OpenACC clause 'finalize' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) finalize
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'if_present' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) if_present
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
+#pragma acc loop device_type(*) seq
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'independent' not yet implemented, clause ignored}}
+#pragma acc loop device_type(*) independent
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'auto' not yet implemented, clause ignored}}
+#pragma acc loop device_type(*) auto
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
+#pragma acc loop device_type(*) worker
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) nohost
+ for(;;);
+ // expected-error at +1{{OpenACC 'default' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) default(none)
+ for(;;);
+ // expected-error at +1{{OpenACC 'if' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) if(1)
+ for(;;);
+ // expected-error at +1{{OpenACC 'self' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) self
+ for(;;);
+
+ int Var;
+ int *VarPtr;
+ // expected-error at +1{{OpenACC 'copy' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) copy(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'pcopy' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) pcopy(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'present_or_copy' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) present_or_copy(Var)
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'use_device' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) use_device(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'attach' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) attach(Var)
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'delete' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) delete(Var)
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'detach' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) detach(Var)
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'device' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) device(VarPtr)
+ for(;;);
+ // expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) deviceptr(VarPtr)
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'device_resident' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) device_resident(VarPtr)
+ for(;;);
+ // expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) firstprivate(Var)
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'host' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) host(Var)
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'link' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) link(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) no_create(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'present' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) present(Var)
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'private' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) private(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'copyout' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) copyout(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'pcopyout' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) pcopyout(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'present_or_copyout' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) present_or_copyout(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'copyin' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) copyin(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'pcopyin' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) pcopyin(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'present_or_copyin' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) present_or_copyin(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'create' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) create(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'pcreate' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) pcreate(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'present_or_create' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) present_or_create(Var)
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'reduction' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) reduction(+:Var)
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
+#pragma acc loop device_type(*) collapse(1)
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) bind(Var)
+ for(;;);
+ // expected-error at +1{{OpenACC 'vector_length' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) vector_length(1)
+ for(;;);
+ // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) num_gangs(1)
+ for(;;);
+ // expected-error at +1{{OpenACC 'num_workers' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) num_workers(1)
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'device_num' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) device_num(1)
+ for(;;);
+ // expected-error at +2{{OpenACC clause 'default_async' may not follow a 'device_type' clause in a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop device_type(*) default_async(1)
+ for(;;);
+ // expected-error at +1{{OpenACC 'async' clause is not valid on 'loop' directive}}
+#pragma acc loop device_type(*) async
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
+#pragma acc loop device_type(*) tile(Var, 1)
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+#pragma acc loop dtype(*) gang
+ for(;;);
+ // expected-error at +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 at +2{{expected ','}}
+ // expected-error at +1{{expected identifier}}
+#pragma acc loop device_type(T::value)
+ for(;;);
+}
+
+void Inst() {
+ TemplUses<int>(); // #INST
+}
More information about the cfe-commits
mailing list