[clang] 8d6c73c - [OpenACC] enable 'deviceptr' for combined constructs.
via cfe-commits
cfe-commits at lists.llvm.org
Wed Nov 27 07:04:21 PST 2024
Author: erichkeane
Date: 2024-11-27T07:04:15-08:00
New Revision: 8d6c73cbf53bd6eb410ac08836e7b128d4a99a16
URL: https://github.com/llvm/llvm-project/commit/8d6c73cbf53bd6eb410ac08836e7b128d4a99a16
DIFF: https://github.com/llvm/llvm-project/commit/8d6c73cbf53bd6eb410ac08836e7b128d4a99a16.diff
LOG: [OpenACC] enable 'deviceptr' for combined constructs.
This is another clause whose implementation is identical for combined
constructs as with compute constructs, so this adds tests and enables
it.
Added:
clang/test/SemaOpenACC/combined-construct-deviceptr-ast.cpp
clang/test/SemaOpenACC/combined-construct-deviceptr-clause.c
clang/test/SemaOpenACC/combined-construct-deviceptr-clause.cpp
Modified:
clang/lib/Sema/SemaOpenACC.cpp
clang/test/AST/ast-print-openacc-combined-construct.cpp
clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
Removed:
################################################################################
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index d146edeabab741..f9c002d116b20d 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -961,10 +961,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitAttachClause(
OpenACCClause *SemaOpenACCClauseVisitor::VisitDevicePtrClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
- // 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'/'combined' constructs,
+ // and 'compute'/'combined' constructs are the only construct that can do
+ // anything with this yet, so skip/treat as unimplemented in this case.
+ if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
+ !isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
return isNotImplemented();
// ActOnVar ensured that everything is a valid variable reference, but we
diff --git a/clang/test/AST/ast-print-openacc-combined-construct.cpp b/clang/test/AST/ast-print-openacc-combined-construct.cpp
index e04c39ac9bc5be..0d02d3daad162e 100644
--- a/clang/test/AST/ast-print-openacc-combined-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-combined-construct.cpp
@@ -131,4 +131,10 @@ void foo() {
// CHECK: #pragma acc kernels loop present(i, array[1], array, array[1:2])
#pragma acc kernels loop present(i, array[1], array, array[1:2])
for(int i = 0;i<5;++i);
+
+ float *arrayPtr[5];
+
+ // CHECK: #pragma acc kernels loop deviceptr(iPtr, arrayPtr[0])
+#pragma acc kernels loop deviceptr(iPtr, arrayPtr[0])
+ for(int i = 0;i<5;++i);
}
diff --git a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
index 69f93a6c605156..fe86960cac4ba1 100644
--- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
@@ -92,8 +92,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
#pragma acc parallel loop auto device(VarPtr)
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +1{{OpenACC clause 'deviceptr' not yet implemented}}
#pragma acc parallel loop auto deviceptr(VarPtr)
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
@@ -253,8 +251,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
#pragma acc parallel loop device(VarPtr) auto
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +1{{OpenACC clause 'deviceptr' not yet implemented}}
#pragma acc parallel loop deviceptr(VarPtr) auto
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
@@ -415,8 +411,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
#pragma acc parallel loop independent device(VarPtr)
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +1{{OpenACC clause 'deviceptr' not yet implemented}}
#pragma acc parallel loop independent deviceptr(VarPtr)
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
@@ -576,8 +570,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
#pragma acc parallel loop device(VarPtr) independent
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +1{{OpenACC clause 'deviceptr' not yet implemented}}
#pragma acc parallel loop deviceptr(VarPtr) independent
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
@@ -744,8 +736,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
#pragma acc parallel loop seq device(VarPtr)
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +1{{OpenACC clause 'deviceptr' not yet implemented}}
#pragma acc parallel loop seq deviceptr(VarPtr)
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
@@ -911,8 +901,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
#pragma acc parallel loop device(VarPtr) seq
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +1{{OpenACC clause 'deviceptr' not yet implemented}}
#pragma acc parallel loop deviceptr(VarPtr) seq
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
diff --git a/clang/test/SemaOpenACC/combined-construct-deviceptr-ast.cpp b/clang/test/SemaOpenACC/combined-construct-deviceptr-ast.cpp
new file mode 100644
index 00000000000000..5cfc9584016900
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-deviceptr-ast.cpp
@@ -0,0 +1,64 @@
+// 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
+
+int Global;
+short GlobalArray[5];
+
+void NormalUses(float *PointerParam) {
+ // CHECK: FunctionDecl{{.*}}NormalUses
+ // CHECK: ParmVarDecl
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel loop deviceptr(PointerParam)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: deviceptr clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}} 'PointerParam' 'float *'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+}
+
+template<typename T>
+void TemplUses(T *t) {
+ // CHECK-NEXT: FunctionTemplateDecl
+ // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T
+ // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (T *)'
+ // CHECK-NEXT: ParmVarDecl{{.*}} referenced t 'T *'
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel loop deviceptr(t)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: deviceptr clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'T *' lvalue ParmVar{{.*}} 't' 'T *'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+
+ // Check the instantiated versions of the above.
+ // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (int *)' implicit_instantiation
+ // CHECK-NEXT: TemplateArgument type 'int'
+ // CHECK-NEXT: BuiltinType{{.*}} 'int'
+ // CHECK-NEXT: ParmVarDecl{{.*}} used t 'int *'
+ // CHECK-NEXT: CompoundStmt
+
+// #pragma acc parallel loop deviceptr(t)
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: deviceptr clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 't' 'int *'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+}
+
+void Inst() {
+ int i;
+ TemplUses(&i);
+}
+#endif
diff --git a/clang/test/SemaOpenACC/combined-construct-deviceptr-clause.c b/clang/test/SemaOpenACC/combined-construct-deviceptr-clause.c
new file mode 100644
index 00000000000000..29b787cf934af3
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-deviceptr-clause.c
@@ -0,0 +1,69 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct S {
+ int IntMem;
+ int *PtrMem;
+};
+
+void uses() {
+ int LocalInt;
+ int *LocalPtr;
+ int Array[5];
+ int *PtrArray[5];
+ struct S s;
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel loop deviceptr(LocalInt)
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+#pragma acc parallel loop deviceptr(&LocalInt)
+ for (int i = 0; i < 5; ++i);
+
+#pragma acc serial loop deviceptr(LocalPtr)
+ for (int i = 0; i < 5; ++i);
+
+#pragma acc kernels loop deviceptr(LocalPtr)
+ for (int i = 0; i < 5; ++i);
+
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}}
+#pragma acc kernels loop deviceptr(Array)
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel loop deviceptr(Array[0])
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel loop deviceptr(Array[0:1])
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}}
+#pragma acc parallel loop deviceptr(PtrArray)
+ for (int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop deviceptr(PtrArray[0])
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel loop deviceptr(PtrArray[0:1])
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'struct S'}}
+#pragma acc parallel loop deviceptr(s)
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel loop deviceptr(s.IntMem)
+ for (int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop deviceptr(s.PtrMem)
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
+#pragma acc loop deviceptr(LocalInt)
+ for(int i = 5; i < 10;++i);
+}
diff --git a/clang/test/SemaOpenACC/combined-construct-deviceptr-clause.cpp b/clang/test/SemaOpenACC/combined-construct-deviceptr-clause.cpp
new file mode 100644
index 00000000000000..4baa4ce987405a
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-deviceptr-clause.cpp
@@ -0,0 +1,120 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct S {
+ int IntMem;
+ int *PtrMem;
+ operator int*();
+};
+
+void uses() {
+ int LocalInt;
+ int *LocalPtr;
+ int Array[5];
+ int *PtrArray[5];
+ struct S s;
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel loop deviceptr(LocalInt)
+ for (int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop deviceptr(LocalPtr)
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}}
+#pragma acc parallel loop deviceptr(Array)
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel loop deviceptr(Array[0])
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel loop deviceptr(Array[0:1])
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}}
+#pragma acc parallel loop deviceptr(PtrArray)
+ for (int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop deviceptr(PtrArray[0])
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel loop deviceptr(PtrArray[0:1])
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'struct S'}}
+#pragma acc parallel loop deviceptr(s)
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel loop deviceptr(s.IntMem)
+ for (int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop deviceptr(s.PtrMem)
+ for (int i = 0; i < 5; ++i);
+}
+
+template<typename T, typename TPtr, typename TStruct, auto &R1>
+void Templ() {
+ T SomeInt;
+ TPtr SomePtr;
+ T SomeIntArray[5];
+ TPtr SomeIntPtrArray[5];
+ TStruct SomeStruct;
+
+ // expected-error at +2{{expected pointer in 'deviceptr' clause, type is 'int'}}
+ // expected-note@#INST{{in instantiation of function template specialization}}
+#pragma acc parallel loop deviceptr(SomeInt)
+ for (int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop deviceptr(SomePtr)
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}}
+#pragma acc parallel loop deviceptr(SomeIntArray)
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel loop deviceptr(SomeIntArray[0])
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel loop deviceptr(SomeIntArray[0:1])
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}}
+#pragma acc parallel loop deviceptr(SomeIntPtrArray)
+ for (int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop deviceptr(SomeIntPtrArray[0])
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel loop deviceptr(SomeIntPtrArray[0:1])
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'S'}}
+#pragma acc parallel loop deviceptr(SomeStruct)
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel loop deviceptr(SomeStruct.IntMem)
+ for (int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop deviceptr(SomeStruct.PtrMem)
+ for (int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel loop deviceptr(R1)
+ for (int i = 0; i < 5; ++i);
+}
+
+void inst() {
+ static constexpr int CEVar = 1;
+ Templ<int, int*, S, CEVar>(); // #INST
+}
More information about the cfe-commits
mailing list