Author: erichkeane Date: 2025-03-07T06:49:29-08:00 New Revision: 1c1140c4cf7d4fde5de4b20c6b993bab9060d9fc
URL: https://github.com/llvm/llvm-project/commit/1c1140c4cf7d4fde5de4b20c6b993bab9060d9fc DIFF: https://github.com/llvm/llvm-project/commit/1c1140c4cf7d4fde5de4b20c6b993bab9060d9fc.diff LOG: [OpenACC] Enable 'device_type' for 'routine' There is a slightly different list for routine on which clauses are permitted after it (like the rest of the constructs), but this implements and tests them to make sure we get them right. Added: Modified: clang/lib/Sema/SemaOpenACCClause.cpp clang/lib/Sema/SemaTemplateInstantiateDecl.cpp clang/test/AST/ast-print-openacc-routine-construct.cpp clang/test/SemaOpenACC/routine-construct-ast.cpp clang/test/SemaOpenACC/routine-construct-clauses.cpp Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp index 582681f247b31..1d70178d03611 100644 --- a/clang/lib/Sema/SemaOpenACCClause.cpp +++ b/clang/lib/Sema/SemaOpenACCClause.cpp @@ -509,11 +509,6 @@ bool checkAlreadyHasClauseOfKind( bool checkValidAfterDeviceType( SemaOpenACC &S, const OpenACCDeviceTypeClause &DeviceTypeClause, const SemaOpenACC::OpenACCParsedClause &NewClause) { - // This is implemented for everything but 'routine', so treat as 'fine' for - // that. - if (NewClause.getDirectiveKind() == OpenACCDirectiveKind::Routine) - return false; - // 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 @@ -601,6 +596,19 @@ bool checkValidAfterDeviceType( default: break; } + } else if (NewClause.getDirectiveKind() == OpenACCDirectiveKind::Routine) { + // OpenACC 3.3 section 2.15: Only the 'gang', 'worker', 'vector', 'seq', and + // 'bind' clauses may follow a device_type clause. + switch (NewClause.getClauseKind()) { + case OpenACCClauseKind::Gang: + case OpenACCClauseKind::Worker: + case OpenACCClauseKind::Vector: + case OpenACCClauseKind::Seq: + case OpenACCClauseKind::Bind: + return false; + default: + break; + } } S.Diag(NewClause.getBeginLoc(), diag::err_acc_clause_after_device_type) << NewClause.getClauseKind() << DeviceTypeClause.getClauseKind() @@ -1256,10 +1264,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWaitClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceTypeClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // Restrictions implemented properly on everything except 'routine'. - if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Routine) - return isNotImplemented(); - // OpenACC 3.3 2.14.3: Two instances of the same clause may not appear on the // same directive. if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Set && @@ -2056,11 +2060,8 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses, return nullptr; } - if (const auto *DevTypeClause = - llvm::find_if(ExistingClauses, - [&](const OpenACCClause *C) { - return isa<OpenACCDeviceTypeClause>(C); - }); + if (const auto *DevTypeClause = llvm::find_if( + ExistingClauses, llvm::IsaPred<OpenACCDeviceTypeClause>); DevTypeClause != ExistingClauses.end()) { if (checkValidAfterDeviceType( *this, *cast<OpenACCDeviceTypeClause>(*DevTypeClause), Clause)) diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index f3e33078c55c7..cb434dcf6eb9a 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -1056,7 +1056,6 @@ CLAUSE_NOT_ON_DECLS(Delete) CLAUSE_NOT_ON_DECLS(Detach) CLAUSE_NOT_ON_DECLS(Device) CLAUSE_NOT_ON_DECLS(DeviceNum) -CLAUSE_NOT_ON_DECLS(DeviceType) CLAUSE_NOT_ON_DECLS(Finalize) CLAUSE_NOT_ON_DECLS(FirstPrivate) CLAUSE_NOT_ON_DECLS(Host) @@ -1116,6 +1115,15 @@ void OpenACCDeclClauseInstantiator::VisitNoHostClause( ParsedClause.getEndLoc()); } +void OpenACCDeclClauseInstantiator::VisitDeviceTypeClause( + const OpenACCDeviceTypeClause &C) { + // Nothing to transform here, just create a new version of 'C'. + NewClause = OpenACCDeviceTypeClause::Create( + SemaRef.getASTContext(), C.getClauseKind(), ParsedClause.getBeginLoc(), + ParsedClause.getLParenLoc(), C.getArchitectures(), + ParsedClause.getEndLoc()); +} + void OpenACCDeclClauseInstantiator::VisitWorkerClause( const OpenACCWorkerClause &C) { assert(!C.hasIntExpr() && "Int Expr not allowed on routine 'worker' clause"); diff --git a/clang/test/AST/ast-print-openacc-routine-construct.cpp b/clang/test/AST/ast-print-openacc-routine-construct.cpp index 739718724f179..df41c61142d37 100644 --- a/clang/test/AST/ast-print-openacc-routine-construct.cpp +++ b/clang/test/AST/ast-print-openacc-routine-construct.cpp @@ -7,6 +7,11 @@ int function(); // CHECK: #pragma acc routine(function) vector nohost #pragma acc routine (function) vector nohost +// CHECK: #pragma acc routine(function) device_type(Something) seq +#pragma acc routine(function) device_type(Something) seq +// CHECK: #pragma acc routine(function) dtype(Something) seq +#pragma acc routine(function) dtype(Something) seq + namespace NS { int NSFunc(); auto Lambda = [](){}; @@ -70,6 +75,11 @@ struct DepS { #pragma acc routine(DepS<T>::MemFunc) seq nohost // CHECK: #pragma acc routine(DepS<T>::StaticMemFunc) nohost worker #pragma acc routine(DepS<T>::StaticMemFunc) nohost worker + +// CHECK: #pragma acc routine(MemFunc) worker dtype(*) +#pragma acc routine (MemFunc) worker dtype(*) +// CHECK: #pragma acc routine(MemFunc) device_type(Lambda) vector +#pragma acc routine (MemFunc) device_type(Lambda) vector }; // CHECK: #pragma acc routine(DepS<int>::Lambda) gang diff --git a/clang/test/SemaOpenACC/routine-construct-ast.cpp b/clang/test/SemaOpenACC/routine-construct-ast.cpp index 8fe2c3695826f..f9510eae31590 100644 --- a/clang/test/SemaOpenACC/routine-construct-ast.cpp +++ b/clang/test/SemaOpenACC/routine-construct-ast.cpp @@ -19,6 +19,18 @@ int function(); // CHECK-NEXT: nohost clause // CHECK-NEXT: vector clause +#pragma acc routine(function) device_type(Something) seq +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()' +// CHECK-NEXT: device_type(Something) +// CHECK-NEXT: seq clause +#pragma acc routine(function) nohost dtype(Something) vector +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()' +// CHECK-NEXT: nohost clause +// CHECK-NEXT: dtype(Something) +// CHECK-NEXT: vector clause + namespace NS { // CHECK-NEXT: NamespaceDecl int NSFunc(); @@ -77,6 +89,16 @@ struct S { // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified // CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at // CHECK-NEXT: worker clause +#pragma acc routine(Lambda) worker device_type(Lambda) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at +// CHECK-NEXT: worker clause +// CHECK-NEXT: device_type(Lambda) +#pragma acc routine(Lambda) dtype(Lambda) vector +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at +// CHECK-NEXT: dtype(Lambda) +// CHECK-NEXT: vector clause }; #pragma acc routine(S::MemFunc) gang(dim: 1) @@ -150,6 +172,12 @@ struct DepS { // CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()' // CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>' // CHECK-NEXT: worker clause +#pragma acc routine(DepS<T>::StaticMemFunc) worker device_type(T) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>' +// CHECK-NEXT: worker clause +// CHECK-NEXT: device_type(T) // Instantiation: // CHECK: ClassTemplateSpecializationDecl{{.*}}struct DepS @@ -200,6 +228,12 @@ struct DepS { // CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()' // CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>' // CHECK-NEXT: worker clause + +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>' +// CHECK-NEXT: worker clause +// CHECK-NEXT: device_type(T) }; #pragma acc routine(DepS<int>::Lambda) gang(dim:1) diff --git a/clang/test/SemaOpenACC/routine-construct-clauses.cpp b/clang/test/SemaOpenACC/routine-construct-clauses.cpp index ac9289b798a3e..5fea3d734a410 100644 --- a/clang/test/SemaOpenACC/routine-construct-clauses.cpp +++ b/clang/test/SemaOpenACC/routine-construct-clauses.cpp @@ -121,3 +121,24 @@ struct DependentT { void Inst() { DependentT<HasFuncs> T;// expected-note{{in instantiation of}} } + +#pragma acc routine(Func) gang device_type(Inst) +#pragma acc routine(Func) gang dtype(Inst) +#pragma acc routine(Func) device_type(*) worker +#pragma acc routine(Func) dtype(*) worker + +#pragma acc routine(Func) dtype(*) gang +#pragma acc routine(Func) device_type(*) worker +#pragma acc routine(Func) device_type(*) vector +#pragma acc routine(Func) dtype(*) seq +// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} +#pragma acc routine(Func) seq device_type(*) bind("asdf") +// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} +#pragma acc routine(Func) seq device_type(*) bind(getSomeInt) +#pragma acc routine(Func) seq dtype(*) device_type(*) +// expected-error@+2{{OpenACC clause 'nohost' may not follow a 'dtype' clause in a 'routine' construct}} +// expected-note@+1{{previous clause is here}} +#pragma acc routine(Func) seq dtype(*) nohost +// expected-error@+2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a 'routine' construct}} +// expected-note@+1{{previous clause is here}} +#pragma acc routine(Func) seq device_type(*) nohost _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
