[flang-commits] [llvm] [flang] [clang] [flang][OpenMP] Add semantic check for declare target (PR #72770)
via flang-commits
flang-commits at lists.llvm.org
Sat Nov 18 18:05:37 PST 2023
https://github.com/shraiysh created https://github.com/llvm/llvm-project/pull/72770
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.
>From 70e5a05ad2a65f824fa1c8248513e402e8af210f Mon Sep 17 00:00:00 2001
From: Shraiysh Vaishay <shraiysh.vaishay at amd.com>
Date: Sat, 18 Nov 2023 20:02:30 -0600
Subject: [PATCH] [flang][OpenMP] Add semantic check for declare target
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.
```
---
clang/lib/Parse/ParseOpenMP.cpp | 1 +
.../target_enter_data_nowait_messages.cpp | 14 +-
flang/lib/Lower/OpenMP.cpp | 15 ++
flang/lib/Parser/openmp-parsers.cpp | 2 +
flang/lib/Semantics/check-omp-structure.cpp | 51 +++--
flang/lib/Semantics/check-omp-structure.h | 2 +
flang/lib/Semantics/resolve-directives.cpp | 3 +
.../Lower/OpenMP/FIR/declare-target-data.f90 | 16 ++
.../FIR/declare-target-func-and-subr.f90 | 68 +++++++
...arget-implicit-func-and-subr-cap-enter.f90 | 192 ++++++++++++++++++
.../test/Lower/OpenMP/declare-target-data.f90 | 16 ++
.../OpenMP/declare-target-func-and-subr.f90 | 68 +++++++
...arget-implicit-func-and-subr-cap-enter.f90 | 192 ++++++++++++++++++
.../declare-target-implicit-tarop-cap.f90 | 14 ++
.../Lower/OpenMP/function-filtering-2.f90 | 8 +
.../test/Lower/OpenMP/function-filtering.f90 | 18 ++
.../OpenMP/declare_target-device_type.f90 | 27 ++-
.../OpenMP/declarative-directive.f90 | 12 +-
.../Semantics/OpenMP/declare-target01.f90 | 48 +++++
.../Semantics/OpenMP/declare-target02.f90 | 54 +++++
.../Semantics/OpenMP/declare-target06.f90 | 5 +
flang/test/Semantics/OpenMP/requires04.f90 | 5 +-
flang/test/Semantics/OpenMP/requires05.f90 | 2 +
llvm/include/llvm/Frontend/OpenMP/OMP.td | 6 +-
24 files changed, 817 insertions(+), 22 deletions(-)
create mode 100644 flang/test/Lower/OpenMP/FIR/declare-target-implicit-func-and-subr-cap-enter.f90
create mode 100644 flang/test/Lower/OpenMP/declare-target-implicit-func-and-subr-cap-enter.f90
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_captured_twice
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_twice() result(k)
+ integer :: i
+ i = 10
+ k = i
+end function implicitly_captured_twice
+
+! CHECK-LABEL: func.func @_QPtarget_function_twice_host
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (enter)>{{.*}}}
+function target_function_twice_host() result(i)
+!$omp declare target enter(target_function_twice_host) device_type(host)
+ integer :: i
+ i = implicitly_captured_twice()
+end function target_function_twice_host
+
+! DEVICE-LABEL: func.func @_QPtarget_function_twice_device
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+function target_function_twice_device() result(i)
+!$omp declare target enter(target_function_twice_device) device_type(nohost)
+ integer :: i
+ i = implicitly_captured_twice()
+end function target_function_twice_device
+
+!! -----
+
+! DEVICE-LABEL: func.func @_QPimplicitly_captured_nest
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_nest() result(k)
+ integer :: i
+ i = 10
+ k = i
+end function implicitly_captured_nest
+
+! DEVICE-LABEL: func.func @_QPimplicitly_captured_one
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter){{.*}}}
+function implicitly_captured_one() result(k)
+ k = implicitly_captured_nest()
+end function implicitly_captured_one
+
+! DEVICE-LABEL: func.func @_QPimplicitly_captured_two
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_two() result(k)
+ integer :: i
+ i = 10
+ k = i
+end function implicitly_captured_two
+
+! DEVICE-LABEL: func.func @_QPtarget_function_test
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+function target_function_test() result(j)
+!$omp declare target enter(target_function_test) device_type(nohost)
+ integer :: i, j
+ i = implicitly_captured_one()
+ j = implicitly_captured_two() + i
+end function target_function_test
+
+!! -----
+
+! CHECK-LABEL: func.func @_QPimplicitly_captured_nest_twice
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_nest_twice() result(k)
+ integer :: i
+ i = 10
+ k = i
+end function implicitly_captured_nest_twice
+
+! CHECK-LABEL: func.func @_QPimplicitly_captured_one_twice
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_one_twice() result(k)
+ k = implicitly_captured_nest_twice()
+end function implicitly_captured_one_twice
+
+! CHECK-LABEL: func.func @_QPimplicitly_captured_two_twice
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_two_twice() result(k)
+ integer :: i
+ i = 10
+ k = i
+end function implicitly_captured_two_twice
+
+! DEVICE-LABEL: func.func @_QPtarget_function_test_device
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+function target_function_test_device() result(j)
+ !$omp declare target enter(target_function_test_device) device_type(nohost)
+ integer :: i, j
+ i = implicitly_captured_one_twice()
+ j = implicitly_captured_two_twice() + i
+end function target_function_test_device
+
+! CHECK-LABEL: func.func @_QPtarget_function_test_host
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (enter)>{{.*}}}
+function target_function_test_host() result(j)
+ !$omp declare target enter(target_function_test_host) device_type(host)
+ integer :: i, j
+ i = implicitly_captured_one_twice()
+ j = implicitly_captured_two_twice() + i
+end function target_function_test_host
+
+!! -----
+
+! DEVICE-LABEL: func.func @_QPimplicitly_captured_with_dev_type_recursive
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+recursive function implicitly_captured_with_dev_type_recursive(increment) result(k)
+!$omp declare target enter(implicitly_captured_with_dev_type_recursive) device_type(host)
+ integer :: increment, k
+ if (increment == 10) then
+ k = increment
+ else
+ k = implicitly_captured_with_dev_type_recursive(increment + 1)
+ end if
+end function implicitly_captured_with_dev_type_recursive
+
+! DEVICE-LABEL: func.func @_QPtarget_function_with_dev_type_recurse
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+function target_function_with_dev_type_recurse() result(i)
+!$omp declare target enter(target_function_with_dev_type_recurse) device_type(nohost)
+ integer :: i
+ i = implicitly_captured_with_dev_type_recursive(0)
+end function target_function_with_dev_type_recurse
+
+!! -----
+
+module test_module
+contains
+! CHECK-LABEL: func.func @_QMtest_modulePimplicitly_captured_nest_twice
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+ function implicitly_captured_nest_twice() result(i)
+ integer :: i
+ i = 10
+ end function implicitly_captured_nest_twice
+
+! CHECK-LABEL: func.func @_QMtest_modulePimplicitly_captured_one_twice
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+ function implicitly_captured_one_twice() result(k)
+ !$omp declare target enter(implicitly_captured_one_twice) device_type(host)
+ k = implicitly_captured_nest_twice()
+ end function implicitly_captured_one_twice
+
+! DEVICE-LABEL: func.func @_QMtest_modulePimplicitly_captured_two_twice
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+ function implicitly_captured_two_twice() result(y)
+ integer :: y
+ y = 5
+ end function implicitly_captured_two_twice
+
+! DEVICE-LABEL: func.func @_QMtest_modulePtarget_function_test_device
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+ function target_function_test_device() result(j)
+ !$omp declare target enter(target_function_test_device) device_type(nohost)
+ integer :: i, j
+ i = implicitly_captured_one_twice()
+ j = implicitly_captured_two_twice() + i
+ end function target_function_test_device
+end module test_module
+
+!! -----
+
+program mb
+ interface
+ subroutine caller_recursive
+ !$omp declare target enter(caller_recursive) device_type(nohost)
+ end subroutine
+
+ recursive subroutine implicitly_captured_recursive(increment)
+ integer :: increment
+ end subroutine
+ end interface
+end program
+
+! DEVICE-LABEL: func.func @_QPimplicitly_captured_recursive
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+recursive subroutine implicitly_captured_recursive(increment)
+ integer :: increment
+ if (increment == 10) then
+ return
+ else
+ call implicitly_captured_recursive(increment + 1)
+ end if
+end subroutine
+
+! DEVICE-LABEL: func.func @_QPcaller_recursive
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+subroutine caller_recursive
+!$omp declare target enter(caller_recursive) device_type(nohost)
+ call implicitly_captured_recursive(0)
+end subroutine
diff --git a/flang/test/Lower/OpenMP/declare-target-data.f90 b/flang/test/Lower/OpenMP/declare-target-data.f90
index e74f84719b71299..4568810d8d07d4a 100644
--- a/flang/test/Lower/OpenMP/declare-target-data.f90
+++ b/flang/test/Lower/OpenMP/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/declare-target-func-and-subr.f90 b/flang/test/Lower/OpenMP/declare-target-func-and-subr.f90
index a5cdfca844807a5..3d2c4067dab7161 100644
--- a/flang/test/Lower/OpenMP/declare-target-func-and-subr.f90
+++ b/flang/test/Lower/OpenMP/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/declare-target-implicit-func-and-subr-cap-enter.f90 b/flang/test/Lower/OpenMP/declare-target-implicit-func-and-subr-cap-enter.f90
new file mode 100644
index 000000000000000..ed718a485e3ddce
--- /dev/null
+++ b/flang/test/Lower/OpenMP/declare-target-implicit-func-and-subr-cap-enter.f90
@@ -0,0 +1,192 @@
+!RUN: %flang_fc1 -emit-hlfir -fopenmp %s -o - | FileCheck %s
+!RUN: %flang_fc1 -emit-hlfir -fopenmp -fopenmp-is-target-device %s -o - | FileCheck %s --check-prefix=DEVICE
+!RUN: bbc -emit-hlfir -fopenmp %s -o - | FileCheck %s
+!RUN: bbc -emit-hlfir -fopenmp -fopenmp-is-target-device %s -o - | FileCheck %s --check-prefix=DEVICE
+
+! CHECK-LABEL: func.func @_QPimplicitly_captured_twice
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_twice() result(k)
+ integer :: i
+ i = 10
+ k = i
+end function implicitly_captured_twice
+
+! CHECK-LABEL: func.func @_QPtarget_function_twice_host
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (enter)>{{.*}}}
+function target_function_twice_host() result(i)
+!$omp declare target enter(target_function_twice_host) device_type(host)
+ integer :: i
+ i = implicitly_captured_twice()
+end function target_function_twice_host
+
+! DEVICE-LABEL: func.func @_QPtarget_function_twice_device
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+function target_function_twice_device() result(i)
+!$omp declare target enter(target_function_twice_device) device_type(nohost)
+ integer :: i
+ i = implicitly_captured_twice()
+end function target_function_twice_device
+
+!! -----
+
+! DEVICE-LABEL: func.func @_QPimplicitly_captured_nest
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_nest() result(k)
+ integer :: i
+ i = 10
+ k = i
+end function implicitly_captured_nest
+
+! DEVICE-LABEL: func.func @_QPimplicitly_captured_one
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter){{.*}}}
+function implicitly_captured_one() result(k)
+ k = implicitly_captured_nest()
+end function implicitly_captured_one
+
+! DEVICE-LABEL: func.func @_QPimplicitly_captured_two
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_two() result(k)
+ integer :: i
+ i = 10
+ k = i
+end function implicitly_captured_two
+
+! DEVICE-LABEL: func.func @_QPtarget_function_test
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+function target_function_test() result(j)
+!$omp declare target enter(target_function_test) device_type(nohost)
+ integer :: i, j
+ i = implicitly_captured_one()
+ j = implicitly_captured_two() + i
+end function target_function_test
+
+!! -----
+
+! CHECK-LABEL: func.func @_QPimplicitly_captured_nest_twice
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_nest_twice() result(k)
+ integer :: i
+ i = 10
+ k = i
+end function implicitly_captured_nest_twice
+
+! CHECK-LABEL: func.func @_QPimplicitly_captured_one_twice
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_one_twice() result(k)
+ k = implicitly_captured_nest_twice()
+end function implicitly_captured_one_twice
+
+! CHECK-LABEL: func.func @_QPimplicitly_captured_two_twice
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_two_twice() result(k)
+ integer :: i
+ i = 10
+ k = i
+end function implicitly_captured_two_twice
+
+! DEVICE-LABEL: func.func @_QPtarget_function_test_device
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+function target_function_test_device() result(j)
+ !$omp declare target enter(target_function_test_device) device_type(nohost)
+ integer :: i, j
+ i = implicitly_captured_one_twice()
+ j = implicitly_captured_two_twice() + i
+end function target_function_test_device
+
+! CHECK-LABEL: func.func @_QPtarget_function_test_host
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (enter)>{{.*}}}
+function target_function_test_host() result(j)
+ !$omp declare target enter(target_function_test_host) device_type(host)
+ integer :: i, j
+ i = implicitly_captured_one_twice()
+ j = implicitly_captured_two_twice() + i
+end function target_function_test_host
+
+!! -----
+
+! DEVICE-LABEL: func.func @_QPimplicitly_captured_with_dev_type_recursive
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+recursive function implicitly_captured_with_dev_type_recursive(increment) result(k)
+!$omp declare target enter(implicitly_captured_with_dev_type_recursive) device_type(host)
+ integer :: increment, k
+ if (increment == 10) then
+ k = increment
+ else
+ k = implicitly_captured_with_dev_type_recursive(increment + 1)
+ end if
+end function implicitly_captured_with_dev_type_recursive
+
+! DEVICE-LABEL: func.func @_QPtarget_function_with_dev_type_recurse
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+function target_function_with_dev_type_recurse() result(i)
+!$omp declare target enter(target_function_with_dev_type_recurse) device_type(nohost)
+ integer :: i
+ i = implicitly_captured_with_dev_type_recursive(0)
+end function target_function_with_dev_type_recurse
+
+!! -----
+
+module test_module
+contains
+! CHECK-LABEL: func.func @_QMtest_modulePimplicitly_captured_nest_twice
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+ function implicitly_captured_nest_twice() result(i)
+ integer :: i
+ i = 10
+ end function implicitly_captured_nest_twice
+
+! CHECK-LABEL: func.func @_QMtest_modulePimplicitly_captured_one_twice
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+ function implicitly_captured_one_twice() result(k)
+ !$omp declare target enter(implicitly_captured_one_twice) device_type(host)
+ k = implicitly_captured_nest_twice()
+ end function implicitly_captured_one_twice
+
+! DEVICE-LABEL: func.func @_QMtest_modulePimplicitly_captured_two_twice
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+ function implicitly_captured_two_twice() result(y)
+ integer :: y
+ y = 5
+ end function implicitly_captured_two_twice
+
+! DEVICE-LABEL: func.func @_QMtest_modulePtarget_function_test_device
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+ function target_function_test_device() result(j)
+ !$omp declare target enter(target_function_test_device) device_type(nohost)
+ integer :: i, j
+ i = implicitly_captured_one_twice()
+ j = implicitly_captured_two_twice() + i
+ end function target_function_test_device
+end module test_module
+
+!! -----
+
+program mb
+ interface
+ subroutine caller_recursive
+ !$omp declare target enter(caller_recursive) device_type(nohost)
+ end subroutine
+
+ recursive subroutine implicitly_captured_recursive(increment)
+ integer :: increment
+ end subroutine
+ end interface
+end program
+
+! DEVICE-LABEL: func.func @_QPimplicitly_captured_recursive
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+recursive subroutine implicitly_captured_recursive(increment)
+ integer :: increment
+ if (increment == 10) then
+ return
+ else
+ call implicitly_captured_recursive(increment + 1)
+ end if
+end subroutine
+
+! DEVICE-LABEL: func.func @_QPcaller_recursive
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+subroutine caller_recursive
+!$omp declare target enter(caller_recursive) device_type(nohost)
+ call implicitly_captured_recursive(0)
+end subroutine
diff --git a/flang/test/Lower/OpenMP/declare-target-implicit-tarop-cap.f90 b/flang/test/Lower/OpenMP/declare-target-implicit-tarop-cap.f90
index 0f548600f760178..7d1ae06c80561d5 100644
--- a/flang/test/Lower/OpenMP/declare-target-implicit-tarop-cap.f90
+++ b/flang/test/Lower/OpenMP/declare-target-implicit-tarop-cap.f90
@@ -34,6 +34,20 @@ function implicitly_captured_one_twice() result(k)
k = implicitly_captured_nest_twice()
end function implicitly_captured_one_twice
+! CHECK-LABEL: func.func @_QPimplicitly_captured_nest_twice_enter
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_nest_twice_enter() result(i)
+ integer :: i
+ i = 10
+end function implicitly_captured_nest_twice_enter
+
+! CHECK-LABEL: func.func @_QPimplicitly_captured_one_twice_enter
+! CHECK-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (enter)>{{.*}}}
+function implicitly_captured_one_twice_enter() result(k)
+!$omp declare target enter(implicitly_captured_one_twice_enter) device_type(host)
+ k = implicitly_captured_nest_twice_enter()
+end function implicitly_captured_one_twice_enter
+
! DEVICE-LABEL: func.func @_QPimplicitly_captured_two_twice
! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (to)>{{.*}}}
function implicitly_captured_two_twice() result(y)
diff --git a/flang/test/Lower/OpenMP/function-filtering-2.f90 b/flang/test/Lower/OpenMP/function-filtering-2.f90
index 17cd0d44c01b4bc..e1d0f72b9f525a9 100644
--- a/flang/test/Lower/OpenMP/function-filtering-2.f90
+++ b/flang/test/Lower/OpenMP/function-filtering-2.f90
@@ -19,6 +19,14 @@ subroutine declaretarget()
call implicit_invocation()
end subroutine declaretarget
+! MLIR: func.func @{{.*}}declaretarget_enter() attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>}
+! MLIR: return
+! LLVM: define {{.*}} @{{.*}}declaretarget_enter{{.*}}(
+subroutine declaretarget_enter()
+!$omp declare target enter(declaretarget_enter) device_type(nohost)
+ call implicit_invocation()
+end subroutine declaretarget_enter
+
! MLIR: func.func @{{.*}}no_declaretarget() attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (to)>}
! MLIR: return
! LLVM: define {{.*}} @{{.*}}no_declaretarget{{.*}}(
diff --git a/flang/test/Lower/OpenMP/function-filtering.f90 b/flang/test/Lower/OpenMP/function-filtering.f90
index 45d8c2e2533d07a..c473b9961907d16 100644
--- a/flang/test/Lower/OpenMP/function-filtering.f90
+++ b/flang/test/Lower/OpenMP/function-filtering.f90
@@ -19,6 +19,16 @@ function device_fn() result(x)
x = 10
end function device_fn
+! MLIR-ALL: func.func @{{.*}}device_fn_enter(
+! MLIR-ALL: return
+
+! LLVM-ALL: define {{.*}} @{{.*}}device_fn_enter{{.*}}(
+function device_fn_enter() result(x)
+ !$omp declare target enter(device_fn_enter) device_type(nohost)
+ integer :: x
+ x = 10
+end function device_fn_enter
+
! MLIR-HOST: func.func @{{.*}}host_fn(
! MLIR-HOST: return
! MLIR-DEVICE-NOT: func.func {{.*}}host_fn(
@@ -31,6 +41,14 @@ function host_fn() result(x)
x = 10
end function host_fn
+! LLVM-HOST: define {{.*}} @{{.*}}host_fn_enter{{.*}}(
+! LLVM-DEVICE-NOT: {{.*}} @{{.*}}host_fn_enter{{.*}}(
+function host_fn_enter() result(x)
+ !$omp declare target enter(host_fn_enter) device_type(host)
+ integer :: x
+ x = 10
+end function host_fn_enter
+
! MLIR-ALL: func.func @{{.*}}target_subr(
! MLIR-ALL: return
diff --git a/flang/test/Parser/OpenMP/declare_target-device_type.f90 b/flang/test/Parser/OpenMP/declare_target-device_type.f90
index 2f6b68ab30f6e91..0b4f75e7ddccbf8 100644
--- a/flang/test/Parser/OpenMP/declare_target-device_type.f90
+++ b/flang/test/Parser/OpenMP/declare_target-device_type.f90
@@ -2,12 +2,27 @@
! RUN: %flang_fc1 -fdebug-dump-parse-tree -fopenmp %s | FileCheck --check-prefix="PARSE-TREE" %s
subroutine openmp_declare_target
- !CHECK: !$omp declare target device_type(host)
- !$omp declare target device_type(host)
- !CHECK: !$omp declare target device_type(nohost)
- !$omp declare target device_type(nohost)
- !CHECK: !$omp declare target device_type(any)
- !$omp declare target device_type(any)
+ integer, save :: x, y
+ !CHECK: !$omp declare target device_type(host) enter(x)
+ !$omp declare target device_type(host) enter(x)
+ !CHECK: !$omp declare target device_type(nohost) enter(x)
+ !$omp declare target device_type(nohost) enter(x)
+ !CHECK: !$omp declare target device_type(any) enter(x)
+ !$omp declare target device_type(any) enter(x)
+
+ !CHECK: !$omp declare target device_type(host) to(x)
+ !$omp declare target device_type(host) to(x)
+ !CHECK: !$omp declare target device_type(nohost) to(x)
+ !$omp declare target device_type(nohost) to(x)
+ !CHECK: !$omp declare target device_type(any) to(x)
+ !$omp declare target device_type(any) to(x)
+
+ !CHECK: !$omp declare target device_type(host) enter(y) to(x)
+ !$omp declare target device_type(host) enter(y) to(x)
+ !CHECK: !$omp declare target device_type(nohost) enter(y) to(x)
+ !$omp declare target device_type(nohost) enter(y) to(x)
+ !CHECK: !$omp declare target device_type(any) enter(y) to(x)
+ !$omp declare target device_type(any) enter(y) to(x)
integer :: a(1024), i
!CHECK: do
do i = 1, 1024
diff --git a/flang/test/Semantics/OpenMP/declarative-directive.f90 b/flang/test/Semantics/OpenMP/declarative-directive.f90
index e48a682288583f7..4d10dc2d1b123e3 100644
--- a/flang/test/Semantics/OpenMP/declarative-directive.f90
+++ b/flang/test/Semantics/OpenMP/declarative-directive.f90
@@ -61,12 +61,22 @@ subroutine foo
!WARNING: The entity with PARAMETER attribute is used in a DECLARE TARGET directive
!WARNING: The entity with PARAMETER attribute is used in a DECLARE TARGET directive
!$omp declare target (foo, N, M)
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
!ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
!ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
!$omp declare target to(Q, S) link(R)
+ !ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
+ !ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
+ !ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
+ !$omp declare target enter(Q, S) link(R)
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
+ !ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
+ !ERROR: MAP clause is not allowed on the DECLARE TARGET directive
+ !$omp declare target to(Q) map(from:Q)
+ !ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
!ERROR: MAP clause is not allowed on the DECLARE TARGET directive
- !$omp declare target map(from:Q)
+ !$omp declare target enter(Q) map(from:Q)
integer, parameter :: N=10000, M=1024
integer :: i
real :: Q(N, N), R(N,M), S(M,M)
diff --git a/flang/test/Semantics/OpenMP/declare-target01.f90 b/flang/test/Semantics/OpenMP/declare-target01.f90
index 862b4c5866f3fb6..3168781ba383917 100644
--- a/flang/test/Semantics/OpenMP/declare-target01.f90
+++ b/flang/test/Semantics/OpenMP/declare-target01.f90
@@ -49,41 +49,89 @@ module declare_target01
!ERROR: A type parameter inquiry cannot appear on the DECLARE TARGET directive
!$omp declare target (y%KIND)
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (my_var)
+ !$omp declare target enter (my_var)
+
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (my_var) device_type(host)
+
+ !$omp declare target enter (my_var) device_type(host)
!ERROR: A variable that is part of another variable (as an array or structure element) cannot appear on the DECLARE TARGET directive
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (my_var%t_i)
!ERROR: A variable that is part of another variable (as an array or structure element) cannot appear on the DECLARE TARGET directive
+ !$omp declare target enter (my_var%t_i)
+
+ !ERROR: A variable that is part of another variable (as an array or structure element) cannot appear on the DECLARE TARGET directive
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (my_var%t_arr)
+ !ERROR: A variable that is part of another variable (as an array or structure element) cannot appear on the DECLARE TARGET directive
+ !$omp declare target enter (my_var%t_arr)
+
!ERROR: A type parameter inquiry cannot appear on the DECLARE TARGET directive
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (my_var%kind_param)
!ERROR: A type parameter inquiry cannot appear on the DECLARE TARGET directive
+ !$omp declare target enter (my_var%kind_param)
+
+ !ERROR: A type parameter inquiry cannot appear on the DECLARE TARGET directive
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (my_var%len_param)
+ !ERROR: A type parameter inquiry cannot appear on the DECLARE TARGET directive
+ !$omp declare target enter (my_var%len_param)
+
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (arr)
+ !$omp declare target enter (arr)
+
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (arr) device_type(nohost)
+ !$omp declare target enter (arr) device_type(nohost)
+
!ERROR: A variable that is part of another variable (as an array or structure element) cannot appear on the DECLARE TARGET directive
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (arr(1))
!ERROR: A variable that is part of another variable (as an array or structure element) cannot appear on the DECLARE TARGET directive
+ !$omp declare target enter (arr(1))
+
+ !ERROR: A variable that is part of another variable (as an array or structure element) cannot appear on the DECLARE TARGET directive
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (arr(1:2))
+ !ERROR: A variable that is part of another variable (as an array or structure element) cannot appear on the DECLARE TARGET directive
+ !$omp declare target enter (arr(1:2))
+
!ERROR: A type parameter inquiry cannot appear on the DECLARE TARGET directive
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (x%KIND)
!ERROR: A type parameter inquiry cannot appear on the DECLARE TARGET directive
+ !$omp declare target enter (x%KIND)
+
+ !ERROR: A type parameter inquiry cannot appear on the DECLARE TARGET directive
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (w%LEN)
!ERROR: A type parameter inquiry cannot appear on the DECLARE TARGET directive
+ !$omp declare target enter (w%LEN)
+
+ !ERROR: A type parameter inquiry cannot appear on the DECLARE TARGET directive
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (y%KIND)
+ !ERROR: A type parameter inquiry cannot appear on the DECLARE TARGET directive
+ !$omp declare target enter (y%KIND)
+
!$omp declare target link (my_var2)
!$omp declare target link (my_var2) device_type(any)
diff --git a/flang/test/Semantics/OpenMP/declare-target02.f90 b/flang/test/Semantics/OpenMP/declare-target02.f90
index 57a2776ff934054..8166e10d702b899 100644
--- a/flang/test/Semantics/OpenMP/declare-target02.f90
+++ b/flang/test/Semantics/OpenMP/declare-target02.f90
@@ -16,13 +16,23 @@ program declare_target02
!ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
!$omp declare target (a1)
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (arr1_to)
+ !$omp declare target enter (arr1_to)
+
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (blk1_to)
+ !$omp declare target enter (blk1_to)
+
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
!$omp declare target to (a1_to)
+ !ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
+ !$omp declare target enter (a1_to)
+
!$omp declare target link (arr1_link)
!$omp declare target link (blk1_link)
@@ -34,18 +44,26 @@ program declare_target02
!ERROR: A variable in a DECLARE TARGET directive cannot appear in an EQUIVALENCE statement
!$omp declare target (eq_a)
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!ERROR: A variable in a DECLARE TARGET directive cannot appear in an EQUIVALENCE statement
!$omp declare target to (eq_a)
+ !ERROR: A variable in a DECLARE TARGET directive cannot appear in an EQUIVALENCE statement
+ !$omp declare target enter (eq_a)
+
!ERROR: A variable in a DECLARE TARGET directive cannot appear in an EQUIVALENCE statement
!$omp declare target link (eq_b)
!ERROR: A variable in a DECLARE TARGET directive cannot appear in an EQUIVALENCE statement
!$omp declare target (eq_c)
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!ERROR: A variable in a DECLARE TARGET directive cannot appear in an EQUIVALENCE statement
!$omp declare target to (eq_c)
+ !ERROR: A variable in a DECLARE TARGET directive cannot appear in an EQUIVALENCE statement
+ !$omp declare target enter (eq_c)
+
!ERROR: A variable in a DECLARE TARGET directive cannot appear in an EQUIVALENCE statement
!$omp declare target link (eq_d)
equivalence(eq_c, eq_d)
@@ -69,17 +87,32 @@ subroutine func()
!ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
!$omp declare target (a3)
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
!$omp declare target to (arr2_to)
+ !ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
+ !$omp declare target enter (arr2_to)
+
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (arr3_to)
+ !$omp declare target enter (arr3_to)
+
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
!$omp declare target to (a2_to)
+ !ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
+ !$omp declare target enter (a2_to)
+
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
!$omp declare target to (a3_to)
+ !ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
+ !$omp declare target enter (a3_to)
+
!ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
!$omp declare target link (arr2_link)
@@ -104,13 +137,22 @@ module mod4
!ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
!$omp declare target (a4)
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (arr4_to)
+ !$omp declare target enter (arr4_to)
+
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to (blk4_to)
+ !$omp declare target enter (blk4_to)
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
!$omp declare target to (a4_to)
+ !ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
+ !$omp declare target enter (a4_to)
+
!$omp declare target link (arr4_link)
!$omp declare target link (blk4_link)
@@ -132,15 +174,27 @@ subroutine func5()
!ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
!$omp declare target (a5)
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
!$omp declare target to (arr5_to)
+
+ !ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
+ !$omp declare target enter (arr5_to)
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
!$omp declare target to (blk5_to)
+ !ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
+ !$omp declare target enter (blk5_to)
+
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
!$omp declare target to (a5_to)
+ !ERROR: A variable in a DECLARE TARGET directive cannot be an element of a common block
+ !$omp declare target enter (a5_to)
+
!ERROR: A variable that appears in a DECLARE TARGET directive must be declared in the scope of a module or have the SAVE attribute, either explicitly or implicitly
!$omp declare target link (arr5_link)
diff --git a/flang/test/Semantics/OpenMP/declare-target06.f90 b/flang/test/Semantics/OpenMP/declare-target06.f90
index a2d8263a60294ba..a1c55d39e1b684c 100644
--- a/flang/test/Semantics/OpenMP/declare-target06.f90
+++ b/flang/test/Semantics/OpenMP/declare-target06.f90
@@ -15,9 +15,14 @@ module test_0
!$omp declare target link(no_implicit_materialization_2)
!ERROR: The given DECLARE TARGET directive clause has an invalid argument
+!WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!ERROR: No explicit type declared for 'no_implicit_materialization_3'
!$omp declare target to(no_implicit_materialization_3)
+!ERROR: The given DECLARE TARGET directive clause has an invalid argument
+!ERROR: No explicit type declared for 'no_implicit_materialization_3'
+!$omp declare target enter(no_implicit_materialization_3)
+
INTEGER :: data_int = 10
!$omp declare target(data_int)
end module test_0
diff --git a/flang/test/Semantics/OpenMP/requires04.f90 b/flang/test/Semantics/OpenMP/requires04.f90
index ae1c2c08aa1c7ec..bb4101c1cbd6c4e 100644
--- a/flang/test/Semantics/OpenMP/requires04.f90
+++ b/flang/test/Semantics/OpenMP/requires04.f90
@@ -5,7 +5,10 @@
! device constructs, such as declare target with device_type=nohost|any.
subroutine f
- !$omp declare target device_type(nohost)
+ integer, save :: x
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
+ !$omp declare target to(x) device_type(nohost)
+ !$omp declare target enter(x) device_type(nohost)
end subroutine f
subroutine g
diff --git a/flang/test/Semantics/OpenMP/requires05.f90 b/flang/test/Semantics/OpenMP/requires05.f90
index 9281f8bab4a1694..dd27e3895e394cb 100644
--- a/flang/test/Semantics/OpenMP/requires05.f90
+++ b/flang/test/Semantics/OpenMP/requires05.f90
@@ -5,7 +5,9 @@
! device constructs, such as declare target with 'to' clause and no device_type.
subroutine f
+ !WARNING: The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead.
!$omp declare target to(f)
+ !$omp declare target enter(f)
end subroutine f
subroutine g
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMP.td b/llvm/include/llvm/Frontend/OpenMP/OMP.td
index f8b3b0c7524979b..8bfacd85864c062 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMP.td
+++ b/llvm/include/llvm/Frontend/OpenMP/OMP.td
@@ -289,6 +289,9 @@ def OMPC_To : Clause<"to"> {
let clangClass = "OMPToClause";
let flangClass = "OmpObjectList";
}
+def OMPC_Enter : Clause<"enter"> {
+ let flangClass = "OmpObjectList";
+}
def OMPC_From : Clause<"from"> {
let clangClass = "OMPFromClause";
let flangClass = "OmpObjectList";
@@ -1124,7 +1127,8 @@ def OMP_DeclareTarget : Directive<"declare target"> {
let allowedClauses = [
VersionedClause<OMPC_To>,
VersionedClause<OMPC_Link>,
- VersionedClause<OMPC_Indirect>
+ VersionedClause<OMPC_Indirect>,
+ VersionedClause<OMPC_Enter, 52>
];
let allowedOnceClauses = [
VersionedClause<OMPC_DeviceType, 50>
More information about the flang-commits
mailing list