llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-flang-parser @llvm/pr-subscribers-flang-fir-hlfir Author: Shraiysh (shraiysh) <details> <summary>Changes</summary> This patch adds the following check from OpenMP 5.2. ``` If the directive has a clause, it must contain at least one enter clause or at least one link clause. ``` Also added a warning for the deprication of `TO` clause on `DECLARE TARGET` construct. ``` The clause-name to may be used as a synonym for the clause-name enter. This use has been deprecated. ``` Last attempt to add this caused issues in buildbots - #<!-- -->71861. It should pass now. --- Patch is 67.44 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/72770.diff 24 Files Affected: - (modified) clang/lib/Parse/ParseOpenMP.cpp (+1) - (modified) clang/test/OpenMP/target_enter_data_nowait_messages.cpp (+13-1) - (modified) flang/lib/Lower/OpenMP.cpp (+15) - (modified) flang/lib/Parser/openmp-parsers.cpp (+2) - (modified) flang/lib/Semantics/check-omp-structure.cpp (+39-12) - (modified) flang/lib/Semantics/check-omp-structure.h (+2) - (modified) flang/lib/Semantics/resolve-directives.cpp (+3) - (modified) flang/test/Lower/OpenMP/FIR/declare-target-data.f90 (+16) - (modified) flang/test/Lower/OpenMP/FIR/declare-target-func-and-subr.f90 (+68) - (added) flang/test/Lower/OpenMP/FIR/declare-target-implicit-func-and-subr-cap-enter.f90 (+192) - (modified) flang/test/Lower/OpenMP/declare-target-data.f90 (+16) - (modified) flang/test/Lower/OpenMP/declare-target-func-and-subr.f90 (+68) - (added) flang/test/Lower/OpenMP/declare-target-implicit-func-and-subr-cap-enter.f90 (+192) - (modified) flang/test/Lower/OpenMP/declare-target-implicit-tarop-cap.f90 (+14) - (modified) flang/test/Lower/OpenMP/function-filtering-2.f90 (+8) - (modified) flang/test/Lower/OpenMP/function-filtering.f90 (+18) - (modified) flang/test/Parser/OpenMP/declare_target-device_type.f90 (+21-6) - (modified) flang/test/Semantics/OpenMP/declarative-directive.f90 (+11-1) - (modified) flang/test/Semantics/OpenMP/declare-target01.f90 (+48) - (modified) flang/test/Semantics/OpenMP/declare-target02.f90 (+54) - (modified) flang/test/Semantics/OpenMP/declare-target06.f90 (+5) - (modified) flang/test/Semantics/OpenMP/requires04.f90 (+4-1) - (modified) flang/test/Semantics/OpenMP/requires05.f90 (+2) - (modified) llvm/include/llvm/Frontend/OpenMP/OMP.td (+5-1) ``````````diff diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp index 3e7d8274aeefc52..ff3417f5a251ade 100644 --- a/clang/lib/Parse/ParseOpenMP.cpp +++ b/clang/lib/Parse/ParseOpenMP.cpp @@ -3369,6 +3369,7 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind, case OMPC_exclusive: case OMPC_affinity: case OMPC_doacross: + case OMPC_enter: if (getLangOpts().OpenMP >= 52 && DKind == OMPD_ordered && CKind == OMPC_depend) Diag(Tok, diag::warn_omp_depend_in_ordered_deprecated); diff --git a/clang/test/OpenMP/target_enter_data_nowait_messages.cpp b/clang/test/OpenMP/target_enter_data_nowait_messages.cpp index 3f5dde00b814ca4..ba5eaf1d2214a07 100644 --- a/clang/test/OpenMP/target_enter_data_nowait_messages.cpp +++ b/clang/test/OpenMP/target_enter_data_nowait_messages.cpp @@ -6,14 +6,26 @@ int main(int argc, char **argv) { int i; #pragma omp nowait target enter data map(to: i) // expected-error {{expected an OpenMP directive}} - #pragma omp target nowait enter data map(to: i) // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}} + {} + #pragma omp target nowait foo data map(to: i) // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}} + {} + #pragma omp target nowait enter data map(to: i) // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}} expected-error {{unexpected OpenMP clause 'enter' in directive '#pragma omp target'}} expected-error {{expected '(' after 'enter'}} + {} #pragma omp target enter nowait data map(to: i) // expected-error {{expected an OpenMP directive}} + {} #pragma omp target enter data nowait() map(to: i) // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}} expected-error {{expected at least one 'map' clause for '#pragma omp target enter data'}} + {} #pragma omp target enter data map(to: i) nowait( // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}} + {} #pragma omp target enter data map(to: i) nowait (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}} + {} #pragma omp target enter data map(to: i) nowait device (-10u) + {} #pragma omp target enter data map(to: i) nowait (3.14) device (-10u) // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}} + {} #pragma omp target enter data map(to: i) nowait nowait // expected-error {{directive '#pragma omp target enter data' cannot contain more than one 'nowait' clause}} + {} #pragma omp target enter data nowait map(to: i) nowait // expected-error {{directive '#pragma omp target enter data' cannot contain more than one 'nowait' clause}} + {} return 0; } diff --git a/flang/lib/Lower/OpenMP.cpp b/flang/lib/Lower/OpenMP.cpp index 00f16026d9b4bb8..6657f46ba6325ca 100644 --- a/flang/lib/Lower/OpenMP.cpp +++ b/flang/lib/Lower/OpenMP.cpp @@ -590,6 +590,8 @@ class ClauseProcessor { bool processSectionsReduction(mlir::Location currentLocation) const; bool processTo(llvm::SmallVectorImpl<DeclareTargetCapturePair> &result) const; bool + processEnter(llvm::SmallVectorImpl<DeclareTargetCapturePair> &result) const; + bool processUseDeviceAddr(llvm::SmallVectorImpl<mlir::Value> &operands, llvm::SmallVectorImpl<mlir::Type> &useDeviceTypes, llvm::SmallVectorImpl<mlir::Location> &useDeviceLocs, @@ -1851,6 +1853,18 @@ bool ClauseProcessor::processTo( }); } +bool ClauseProcessor::processEnter( + llvm::SmallVectorImpl<DeclareTargetCapturePair> &result) const { + return findRepeatableClause<ClauseTy::Enter>( + [&](const ClauseTy::Enter *enterClause, + const Fortran::parser::CharBlock &) { + // Case: declare target to(func, var1, var2)... + gatherFuncAndVarSyms(enterClause->v, + mlir::omp::DeclareTargetCaptureClause::enter, + result); + }); +} + bool ClauseProcessor::processUseDeviceAddr( llvm::SmallVectorImpl<mlir::Value> &operands, llvm::SmallVectorImpl<mlir::Type> &useDeviceTypes, @@ -2792,6 +2806,7 @@ static mlir::omp::DeclareTargetDeviceType getDeclareTargetInfo( ClauseProcessor cp(converter, *clauseList); cp.processTo(symbolAndClause); + cp.processEnter(symbolAndClause); cp.processLink(symbolAndClause); cp.processDeviceType(deviceType); cp.processTODO<Fortran::parser::OmpClause::Indirect>( diff --git a/flang/lib/Parser/openmp-parsers.cpp b/flang/lib/Parser/openmp-parsers.cpp index 0271f699a687ff0..bba1be27158ce7e 100644 --- a/flang/lib/Parser/openmp-parsers.cpp +++ b/flang/lib/Parser/openmp-parsers.cpp @@ -259,6 +259,8 @@ TYPE_PARSER( parenthesized("STATIC" >> maybe("," >> scalarIntExpr)))) || "DYNAMIC_ALLOCATORS" >> construct<OmpClause>(construct<OmpClause::DynamicAllocators>()) || + "ENTER" >> construct<OmpClause>(construct<OmpClause::Enter>( + parenthesized(Parser<OmpObjectList>{}))) || "FINAL" >> construct<OmpClause>(construct<OmpClause::Final>( parenthesized(scalarLogicalExpr))) || "FULL" >> construct<OmpClause>(construct<OmpClause::Full>()) || diff --git a/flang/lib/Semantics/check-omp-structure.cpp b/flang/lib/Semantics/check-omp-structure.cpp index 0b1a581bf298196..4dbbf1260f3ea31 100644 --- a/flang/lib/Semantics/check-omp-structure.cpp +++ b/flang/lib/Semantics/check-omp-structure.cpp @@ -1165,13 +1165,31 @@ void OmpStructureChecker::Enter(const parser::OmpClause::Allocate &x) { } } +void OmpStructureChecker::Enter(const parser::OmpDeclareTargetWithClause &x) { + SetClauseSets(llvm::omp::Directive::OMPD_declare_target); +} + +void OmpStructureChecker::Leave(const parser::OmpDeclareTargetWithClause &x) { + if (x.v.v.size() > 0) { + const parser::OmpClause *enterClause = + FindClause(llvm::omp::Clause::OMPC_enter); + const parser::OmpClause *toClause = FindClause(llvm::omp::Clause::OMPC_to); + const parser::OmpClause *linkClause = + FindClause(llvm::omp::Clause::OMPC_link); + if (!enterClause && !toClause && !linkClause) { + context_.Say(x.source, + "If the DECLARE TARGET directive has a clause, it must contain at lease one ENTER clause or LINK clause"_err_en_US); + } + if (toClause) { + context_.Say(toClause->source, + "The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead."_warn_en_US); + } + } +} + void OmpStructureChecker::Enter(const parser::OpenMPDeclareTargetConstruct &x) { const auto &dir{std::get<parser::Verbatim>(x.t)}; PushContext(dir.source, llvm::omp::Directive::OMPD_declare_target); - const auto &spec{std::get<parser::OmpDeclareTargetSpecifier>(x.t)}; - if (std::holds_alternative<parser::OmpDeclareTargetWithClause>(spec.u)) { - SetClauseSets(llvm::omp::Directive::OMPD_declare_target); - } } void OmpStructureChecker::Enter(const parser::OmpDeclareTargetWithList &x) { @@ -1245,7 +1263,8 @@ void OmpStructureChecker::Leave(const parser::OpenMPDeclareTargetConstruct &x) { CheckThreadprivateOrDeclareTargetVar(*objectList); } else if (const auto *clauseList{ parser::Unwrap<parser::OmpClauseList>(spec.u)}) { - bool toClauseFound{false}, deviceTypeClauseFound{false}; + bool toClauseFound{false}, deviceTypeClauseFound{false}, + enterClauseFound{false}; for (const auto &clause : clauseList->v) { common::visit( common::visitors{ @@ -1260,6 +1279,12 @@ void OmpStructureChecker::Leave(const parser::OpenMPDeclareTargetConstruct &x) { CheckIsVarPartOfAnotherVar(dir.source, linkClause.v); CheckThreadprivateOrDeclareTargetVar(linkClause.v); }, + [&](const parser::OmpClause::Enter &enterClause) { + enterClauseFound = true; + CheckSymbolNames(dir.source, enterClause.v); + CheckIsVarPartOfAnotherVar(dir.source, enterClause.v); + CheckThreadprivateOrDeclareTargetVar(enterClause.v); + }, [&](const parser::OmpClause::DeviceType &deviceTypeClause) { deviceTypeClauseFound = true; if (deviceTypeClause.v.v != @@ -1273,7 +1298,7 @@ void OmpStructureChecker::Leave(const parser::OpenMPDeclareTargetConstruct &x) { }, clause.u); - if (toClauseFound && !deviceTypeClauseFound) { + if ((toClauseFound || enterClauseFound) && !deviceTypeClauseFound) { deviceConstructFound_ = true; } } @@ -2228,6 +2253,7 @@ CHECK_SIMPLE_CLAUSE(CancellationConstructType, OMPC_cancellation_construct_type) CHECK_SIMPLE_CLAUSE(Doacross, OMPC_doacross) CHECK_SIMPLE_CLAUSE(OmpxAttribute, OMPC_ompx_attribute) CHECK_SIMPLE_CLAUSE(OmpxBare, OMPC_ompx_bare) +CHECK_SIMPLE_CLAUSE(Enter, OMPC_enter) CHECK_REQ_SCALAR_INT_CLAUSE(Grainsize, OMPC_grainsize) CHECK_REQ_SCALAR_INT_CLAUSE(NumTasks, OMPC_num_tasks) @@ -3229,12 +3255,13 @@ const parser::OmpObjectList *OmpStructureChecker::GetOmpObjectList( const parser::OmpClause &clause) { // Clauses with OmpObjectList as its data member - using MemberObjectListClauses = std::tuple<parser::OmpClause::Copyprivate, - parser::OmpClause::Copyin, parser::OmpClause::Firstprivate, - parser::OmpClause::From, parser::OmpClause::Lastprivate, - parser::OmpClause::Link, parser::OmpClause::Private, - parser::OmpClause::Shared, parser::OmpClause::To, - parser::OmpClause::UseDevicePtr, parser::OmpClause::UseDeviceAddr>; + using MemberObjectListClauses = + std::tuple<parser::OmpClause::Copyprivate, parser::OmpClause::Copyin, + parser::OmpClause::Firstprivate, parser::OmpClause::From, + parser::OmpClause::Lastprivate, parser::OmpClause::Link, + parser::OmpClause::Private, parser::OmpClause::Shared, + parser::OmpClause::To, parser::OmpClause::Enter, + parser::OmpClause::UseDevicePtr, parser::OmpClause::UseDeviceAddr>; // Clauses with OmpObjectList in the tuple using TupleObjectListClauses = diff --git a/flang/lib/Semantics/check-omp-structure.h b/flang/lib/Semantics/check-omp-structure.h index d35602cca75d549..90e5c9f19127750 100644 --- a/flang/lib/Semantics/check-omp-structure.h +++ b/flang/lib/Semantics/check-omp-structure.h @@ -80,6 +80,8 @@ class OmpStructureChecker void Enter(const parser::OpenMPDeclareTargetConstruct &); void Leave(const parser::OpenMPDeclareTargetConstruct &); void Enter(const parser::OmpDeclareTargetWithList &); + void Enter(const parser::OmpDeclareTargetWithClause &); + void Leave(const parser::OmpDeclareTargetWithClause &); void Enter(const parser::OpenMPExecutableAllocate &); void Leave(const parser::OpenMPExecutableAllocate &); void Enter(const parser::OpenMPAllocatorsConstruct &); diff --git a/flang/lib/Semantics/resolve-directives.cpp b/flang/lib/Semantics/resolve-directives.cpp index bbb105e3516da18..882ecad1f7ed023 100644 --- a/flang/lib/Semantics/resolve-directives.cpp +++ b/flang/lib/Semantics/resolve-directives.cpp @@ -1744,6 +1744,9 @@ bool OmpAttributeVisitor::Pre(const parser::OpenMPDeclareTargetConstruct &x) { } else if (const auto *linkClause{ std::get_if<parser::OmpClause::Link>(&clause.u)}) { ResolveOmpObjectList(linkClause->v, Symbol::Flag::OmpDeclareTarget); + } else if (const auto *enterClause{ + std::get_if<parser::OmpClause::Enter>(&clause.u)}) { + ResolveOmpObjectList(enterClause->v, Symbol::Flag::OmpDeclareTarget); } } } diff --git a/flang/test/Lower/OpenMP/FIR/declare-target-data.f90 b/flang/test/Lower/OpenMP/FIR/declare-target-data.f90 index e57d928f5497a17..bb3bbc8dfa83769 100644 --- a/flang/test/Lower/OpenMP/FIR/declare-target-data.f90 +++ b/flang/test/Lower/OpenMP/FIR/declare-target-data.f90 @@ -32,6 +32,10 @@ module test_0 INTEGER :: data_int_to = 5 !$omp declare target to(data_int_to) +!CHECK-DAG: fir.global @_QMtest_0Edata_int_enter {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : i32 +INTEGER :: data_int_enter = 5 +!$omp declare target enter(data_int_enter) + !CHECK-DAG: fir.global @_QMtest_0Edata_int_clauseless {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>} : i32 INTEGER :: data_int_clauseless = 1 !$omp declare target(data_int_clauseless) @@ -42,6 +46,12 @@ module test_0 REAL :: data_extended_to_2 = 3 !$omp declare target to(data_extended_to_1, data_extended_to_2) +!CHECK-DAG: fir.global @_QMtest_0Edata_extended_enter_1 {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : f32 +!CHECK-DAG: fir.global @_QMtest_0Edata_extended_enter_2 {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : f32 +REAL :: data_extended_enter_1 = 2 +REAL :: data_extended_enter_2 = 3 +!$omp declare target enter(data_extended_enter_1, data_extended_enter_2) + !CHECK-DAG: fir.global @_QMtest_0Edata_extended_link_1 {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (link)>} : f32 !CHECK-DAG: fir.global @_QMtest_0Edata_extended_link_2 {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (link)>} : f32 REAL :: data_extended_link_1 = 2 @@ -69,4 +79,10 @@ PROGRAM commons REAL :: two_to = 2 COMMON /numbers_to/ one_to, two_to !$omp declare target to(/numbers_to/) + + !CHECK-DAG: fir.global @numbers_enter_ {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : tuple<f32, f32> { + REAL :: one_enter = 1 + REAL :: two_enter = 2 + COMMON /numbers_enter/ one_enter, two_enter + !$omp declare target enter(/numbers_enter/) END diff --git a/flang/test/Lower/OpenMP/FIR/declare-target-func-and-subr.f90 b/flang/test/Lower/OpenMP/FIR/declare-target-func-and-subr.f90 index 26741c6795a6ce5..36d4d7db64e5f25 100644 --- a/flang/test/Lower/OpenMP/FIR/declare-target-func-and-subr.f90 +++ b/flang/test/Lower/OpenMP/FIR/declare-target-func-and-subr.f90 @@ -13,6 +13,14 @@ FUNCTION FUNC_T_DEVICE() RESULT(I) I = 1 END FUNCTION FUNC_T_DEVICE +! DEVICE-LABEL: func.func @_QPfunc_enter_device() +! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}} +FUNCTION FUNC_ENTER_DEVICE() RESULT(I) +!$omp declare target enter(FUNC_ENTER_DEVICE) device_type(nohost) + INTEGER :: I + I = 1 +END FUNCTION FUNC_ENTER_DEVICE + ! HOST-LABEL: func.func @_QPfunc_t_host() ! HOST-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (to)>{{.*}} FUNCTION FUNC_T_HOST() RESULT(I) @@ -21,6 +29,14 @@ FUNCTION FUNC_T_HOST() RESULT(I) I = 1 END FUNCTION FUNC_T_HOST +! HOST-LABEL: func.func @_QPfunc_enter_host() +! HOST-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (enter)>{{.*}} +FUNCTION FUNC_ENTER_HOST() RESULT(I) +!$omp declare target enter(FUNC_ENTER_HOST) device_type(host) + INTEGER :: I + I = 1 +END FUNCTION FUNC_ENTER_HOST + ! ALL-LABEL: func.func @_QPfunc_t_any() ! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}} FUNCTION FUNC_T_ANY() RESULT(I) @@ -29,6 +45,14 @@ FUNCTION FUNC_T_ANY() RESULT(I) I = 1 END FUNCTION FUNC_T_ANY +! ALL-LABEL: func.func @_QPfunc_enter_any() +! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}} +FUNCTION FUNC_ENTER_ANY() RESULT(I) +!$omp declare target enter(FUNC_ENTER_ANY) device_type(any) + INTEGER :: I + I = 1 +END FUNCTION FUNC_ENTER_ANY + ! ALL-LABEL: func.func @_QPfunc_default_t_any() ! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}} FUNCTION FUNC_DEFAULT_T_ANY() RESULT(I) @@ -37,6 +61,14 @@ FUNCTION FUNC_DEFAULT_T_ANY() RESULT(I) I = 1 END FUNCTION FUNC_DEFAULT_T_ANY +! ALL-LABEL: func.func @_QPfunc_default_enter_any() +! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}} +FUNCTION FUNC_DEFAULT_ENTER_ANY() RESULT(I) +!$omp declare target enter(FUNC_DEFAULT_ENTER_ANY) + INTEGER :: I + I = 1 +END FUNCTION FUNC_DEFAULT_ENTER_ANY + ! ALL-LABEL: func.func @_QPfunc_default_any() ! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}} FUNCTION FUNC_DEFAULT_ANY() RESULT(I) @@ -65,24 +97,48 @@ SUBROUTINE SUBR_T_DEVICE() !$omp declare target to(SUBR_T_DEVICE) device_type(nohost) END +! DEVICE-LABEL: func.func @_QPsubr_enter_device() +! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}} +SUBROUTINE SUBR_ENTER_DEVICE() +!$omp declare target enter(SUBR_ENTER_DEVICE) device_type(nohost) +END + ! HOST-LABEL: func.func @_QPsubr_t_host() ! HOST-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (to)>{{.*}} SUBROUTINE SUBR_T_HOST() !$omp declare target to(SUBR_T_HOST) device_type(host) END +! HOST-LABEL: func.func @_QPsubr_enter_host() +! HOST-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (enter)>{{.*}} +SUBROUTINE SUBR_ENTER_HOST() +!$omp declare target enter(SUBR_ENTER_HOST) device_type(host) +END + ! ALL-LABEL: func.func @_QPsubr_t_any() ! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}} SUBROUTINE SUBR_T_ANY() !$omp declare target to(SUBR_T_ANY) device_type(any) END +! ALL-LABEL: func.func @_QPsubr_enter_any() +! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}} +SUBROUTINE SUBR_ENTER_ANY() +!$omp declare target enter(SUBR_ENTER_ANY) device_type(any) +END + ! ALL-LABEL: func.func @_QPsubr_default_t_any() ! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}} SUBROUTINE SUBR_DEFAULT_T_ANY() !$omp declare target to(SUBR_DEFAULT_T_ANY) END +! ALL-LABEL: func.func @_QPsubr_default_enter_any() +! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}} +SUBROUTINE SUBR_DEFAULT_ENTER_ANY() +!$omp declare target enter(SUBR_DEFAULT_ENTER_ANY) +END + ! ALL-LABEL: func.func @_QPsubr_default_any() ! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}} SUBROUTINE SUBR_DEFAULT_ANY() @@ -108,3 +164,15 @@ RECURSIVE FUNCTION RECURSIVE_DECLARE_TARGET(INCREMENT) RESULT(K) K = RECURSIVE_DECLARE_TARGET(INCREMENT + 1) END IF END FUNCTION RECURSIVE_DECLARE_TARGET + +! DEVICE-LABEL: func.func @_QPrecursive_declare_target_enter +! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}} +RECURSIVE FUNCTION RECURSIVE_DECLARE_TARGET_ENTER(INCREMENT) RESULT(K) +!$omp declare target enter(RECURSIVE_DECLARE_TARGET_ENTER) device_type(nohost) + INTEGER :: INCREMENT, K + IF (INCREMENT == 10) THEN + K = INCREMENT + ELSE + K = RECURSIVE_DECLARE_TARGET_ENTER(INCREMENT + 1) + END IF +END FUNCTION RECURSIVE_DECLARE_TARGET_ENTER diff --git a/flang/test/Lower/OpenMP/FIR/declare-target-implicit-func-and-subr-cap-enter.f90 b/flang/test/Lower/OpenMP/FIR/declare-target-implicit-func-and-subr-cap-enter.f90 new file mode 100644 index 000000000000000..8e88d1b0f52a95f --- /dev/null +++ b/flang/test/Lower/OpenMP/FIR/declare-target-implicit-func-and-subr-cap-enter.f90 @@ -0,0 +1,192 @@ +!RUN: %flang_fc1 -emit-fir -fopenmp %s -o - | FileCheck %s +!RUN: %flang_fc1 -emit-fir -fopenmp -fopenmp-is-target-device %s -o - | FileCheck %s --check-prefix=DEVICE +!RUN: bbc -emit-fir -fopenmp %s -o - | FileCheck %s +!RUN: bbc -emit-fir -fopenmp -fopenmp-is-target-device %s -o - | FileCheck %s --check-prefix=DEVICE + +! CHECK-LABEL: func.func @_QPimplicitly_c... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/72770 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
