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@+1{{OpenACC clause 'device' not yet implemented}} #pragma acc parallel loop auto device(VarPtr) for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}} #pragma acc parallel loop auto deviceptr(VarPtr) for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}} @@ -253,8 +251,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'device' not yet implemented}} #pragma acc parallel loop device(VarPtr) auto for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}} #pragma acc parallel loop deviceptr(VarPtr) auto for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}} @@ -415,8 +411,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'device' not yet implemented}} #pragma acc parallel loop independent device(VarPtr) for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}} #pragma acc parallel loop independent deviceptr(VarPtr) for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}} @@ -576,8 +570,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'device' not yet implemented}} #pragma acc parallel loop device(VarPtr) independent for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}} #pragma acc parallel loop deviceptr(VarPtr) independent for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}} @@ -744,8 +736,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'device' not yet implemented}} #pragma acc parallel loop seq device(VarPtr) for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}} #pragma acc parallel loop seq deviceptr(VarPtr) for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}} @@ -911,8 +901,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'device' not yet implemented}} #pragma acc parallel loop device(VarPtr) seq for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}} #pragma acc parallel loop deviceptr(VarPtr) seq for(unsigned i = 0; i < 5; ++i); // expected-warning@+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@+1{{expected pointer in 'deviceptr' clause, type is 'int'}} +#pragma acc parallel loop deviceptr(LocalInt) + for (int i = 0; i < 5; ++i); + + // expected-error@+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@+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@+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@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel loop deviceptr(Array[0:1]) + for (int i = 0; i < 5; ++i); + + // expected-error@+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@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel loop deviceptr(PtrArray[0:1]) + for (int i = 0; i < 5; ++i); + + // expected-error@+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@+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@+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@+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@+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@+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@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel loop deviceptr(Array[0:1]) + for (int i = 0; i < 5; ++i); + + // expected-error@+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@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel loop deviceptr(PtrArray[0:1]) + for (int i = 0; i < 5; ++i); + + // expected-error@+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@+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@+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@+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@+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@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel loop deviceptr(SomeIntArray[0:1]) + for (int i = 0; i < 5; ++i); + + // expected-error@+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@+2{{OpenACC sub-array is not allowed here}} + // expected-note@+1{{expected variable of pointer type}} +#pragma acc parallel loop deviceptr(SomeIntPtrArray[0:1]) + for (int i = 0; i < 5; ++i); + + // expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'S'}} +#pragma acc parallel loop deviceptr(SomeStruct) + for (int i = 0; i < 5; ++i); + + // expected-error@+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@+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 +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
