Skip to content

Commit 8840eb3

Browse files
shraiyshshraiysh
authored
[flang][OpenMP] Add semantic check for declare target (#72770)
1 parent 7f215b1 commit 8840eb3

24 files changed

+817
-22
lines changed

‎clang/lib/Parse/ParseOpenMP.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3369,6 +3369,7 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
33693369
case OMPC_exclusive:
33703370
case OMPC_affinity:
33713371
case OMPC_doacross:
3372+
case OMPC_enter:
33723373
if (getLangOpts().OpenMP >= 52 && DKind == OMPD_ordered &&
33733374
CKind == OMPC_depend)
33743375
Diag(Tok, diag::warn_omp_depend_in_ordered_deprecated);

‎clang/test/OpenMP/target_enter_data_nowait_messages.cpp

Lines changed: 13 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,14 +6,26 @@ int main(int argc, char **argv) {
66
int i;
77

88
#pragma omp nowait target enter data map(to: i) // expected-error {{expected an OpenMP directive}}
9-
#pragma omp target nowait enter data map(to: i) // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}
9+
{}
10+
#pragma omp target nowait foo data map(to: i) // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}
11+
{}
12+
#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'}}
13+
{}
1014
#pragma omp target enter nowait data map(to: i) // expected-error {{expected an OpenMP directive}}
15+
{}
1116
#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'}}
17+
{}
1218
#pragma omp target enter data map(to: i) nowait( // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}}
19+
{}
1320
#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}}
21+
{}
1422
#pragma omp target enter data map(to: i) nowait device (-10u)
23+
{}
1524
#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}}
25+
{}
1626
#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}}
27+
{}
1728
#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}}
29+
{}
1830
return 0;
1931
}

‎flang/lib/Lower/OpenMP.cpp

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -564,6 +564,8 @@ class ClauseProcessor {
564564
bool processDepend(llvm::SmallVectorImpl<mlir::Attribute> &dependTypeOperands,
565565
llvm::SmallVectorImpl<mlir::Value> &dependOperands) const;
566566
bool
567+
processEnter(llvm::SmallVectorImpl<DeclareTargetCapturePair> &result) const;
568+
bool
567569
processIf(Fortran::parser::OmpIfClause::DirectiveNameModifier directiveName,
568570
mlir::Value &result) const;
569571
bool
@@ -1851,6 +1853,18 @@ bool ClauseProcessor::processTo(
18511853
});
18521854
}
18531855

1856+
bool ClauseProcessor::processEnter(
1857+
llvm::SmallVectorImpl<DeclareTargetCapturePair> &result) const {
1858+
return findRepeatableClause<ClauseTy::Enter>(
1859+
[&](const ClauseTy::Enter *enterClause,
1860+
const Fortran::parser::CharBlock &) {
1861+
// Case: declare target enter(func, var1, var2)...
1862+
gatherFuncAndVarSyms(enterClause->v,
1863+
mlir::omp::DeclareTargetCaptureClause::enter,
1864+
result);
1865+
});
1866+
}
1867+
18541868
bool ClauseProcessor::processUseDeviceAddr(
18551869
llvm::SmallVectorImpl<mlir::Value> &operands,
18561870
llvm::SmallVectorImpl<mlir::Type> &useDeviceTypes,
@@ -2792,6 +2806,7 @@ static mlir::omp::DeclareTargetDeviceType getDeclareTargetInfo(
27922806

27932807
ClauseProcessor cp(converter, *clauseList);
27942808
cp.processTo(symbolAndClause);
2809+
cp.processEnter(symbolAndClause);
27952810
cp.processLink(symbolAndClause);
27962811
cp.processDeviceType(deviceType);
27972812
cp.processTODO<Fortran::parser::OmpClause::Indirect>(

‎flang/lib/Parser/openmp-parsers.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -259,6 +259,8 @@ TYPE_PARSER(
259259
parenthesized("STATIC" >> maybe("," >> scalarIntExpr)))) ||
260260
"DYNAMIC_ALLOCATORS" >>
261261
construct<OmpClause>(construct<OmpClause::DynamicAllocators>()) ||
262+
"ENTER" >> construct<OmpClause>(construct<OmpClause::Enter>(
263+
parenthesized(Parser<OmpObjectList>{}))) ||
262264
"FINAL" >> construct<OmpClause>(construct<OmpClause::Final>(
263265
parenthesized(scalarLogicalExpr))) ||
264266
"FULL" >> construct<OmpClause>(construct<OmpClause::Full>()) ||

‎flang/lib/Semantics/check-omp-structure.cpp

Lines changed: 39 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -1165,13 +1165,31 @@ void OmpStructureChecker::Enter(const parser::OmpClause::Allocate &x) {
11651165
}
11661166
}
11671167

1168+
void OmpStructureChecker::Enter(const parser::OmpDeclareTargetWithClause &x) {
1169+
SetClauseSets(llvm::omp::Directive::OMPD_declare_target);
1170+
}
1171+
1172+
void OmpStructureChecker::Leave(const parser::OmpDeclareTargetWithClause &x) {
1173+
if (x.v.v.size() > 0) {
1174+
const parser::OmpClause *enterClause =
1175+
FindClause(llvm::omp::Clause::OMPC_enter);
1176+
const parser::OmpClause *toClause = FindClause(llvm::omp::Clause::OMPC_to);
1177+
const parser::OmpClause *linkClause =
1178+
FindClause(llvm::omp::Clause::OMPC_link);
1179+
if (!enterClause && !toClause && !linkClause) {
1180+
context_.Say(x.source,
1181+
"If the DECLARE TARGET directive has a clause, it must contain at lease one ENTER clause or LINK clause"_err_en_US);
1182+
}
1183+
if (toClause) {
1184+
context_.Say(toClause->source,
1185+
"The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead."_warn_en_US);
1186+
}
1187+
}
1188+
}
1189+
11681190
void OmpStructureChecker::Enter(const parser::OpenMPDeclareTargetConstruct &x) {
11691191
const auto &dir{std::get<parser::Verbatim>(x.t)};
11701192
PushContext(dir.source, llvm::omp::Directive::OMPD_declare_target);
1171-
const auto &spec{std::get<parser::OmpDeclareTargetSpecifier>(x.t)};
1172-
if (std::holds_alternative<parser::OmpDeclareTargetWithClause>(spec.u)) {
1173-
SetClauseSets(llvm::omp::Directive::OMPD_declare_target);
1174-
}
11751193
}
11761194

11771195
void OmpStructureChecker::Enter(const parser::OmpDeclareTargetWithList &x) {
@@ -1245,7 +1263,8 @@ void OmpStructureChecker::Leave(const parser::OpenMPDeclareTargetConstruct &x) {
12451263
CheckThreadprivateOrDeclareTargetVar(*objectList);
12461264
} else if (const auto *clauseList{
12471265
parser::Unwrap<parser::OmpClauseList>(spec.u)}) {
1248-
bool toClauseFound{false}, deviceTypeClauseFound{false};
1266+
bool toClauseFound{false}, deviceTypeClauseFound{false},
1267+
enterClauseFound{false};
12491268
for (const auto &clause : clauseList->v) {
12501269
common::visit(
12511270
common::visitors{
@@ -1260,6 +1279,12 @@ void OmpStructureChecker::Leave(const parser::OpenMPDeclareTargetConstruct &x) {
12601279
CheckIsVarPartOfAnotherVar(dir.source, linkClause.v);
12611280
CheckThreadprivateOrDeclareTargetVar(linkClause.v);
12621281
},
1282+
[&](const parser::OmpClause::Enter &enterClause) {
1283+
enterClauseFound = true;
1284+
CheckSymbolNames(dir.source, enterClause.v);
1285+
CheckIsVarPartOfAnotherVar(dir.source, enterClause.v);
1286+
CheckThreadprivateOrDeclareTargetVar(enterClause.v);
1287+
},
12631288
[&](const parser::OmpClause::DeviceType &deviceTypeClause) {
12641289
deviceTypeClauseFound = true;
12651290
if (deviceTypeClause.v.v !=
@@ -1273,7 +1298,7 @@ void OmpStructureChecker::Leave(const parser::OpenMPDeclareTargetConstruct &x) {
12731298
},
12741299
clause.u);
12751300

1276-
if (toClauseFound && !deviceTypeClauseFound) {
1301+
if ((toClauseFound || enterClauseFound) && !deviceTypeClauseFound) {
12771302
deviceConstructFound_ = true;
12781303
}
12791304
}
@@ -2228,6 +2253,7 @@ CHECK_SIMPLE_CLAUSE(CancellationConstructType, OMPC_cancellation_construct_type)
22282253
CHECK_SIMPLE_CLAUSE(Doacross, OMPC_doacross)
22292254
CHECK_SIMPLE_CLAUSE(OmpxAttribute, OMPC_ompx_attribute)
22302255
CHECK_SIMPLE_CLAUSE(OmpxBare, OMPC_ompx_bare)
2256+
CHECK_SIMPLE_CLAUSE(Enter, OMPC_enter)
22312257

22322258
CHECK_REQ_SCALAR_INT_CLAUSE(Grainsize, OMPC_grainsize)
22332259
CHECK_REQ_SCALAR_INT_CLAUSE(NumTasks, OMPC_num_tasks)
@@ -3240,12 +3266,13 @@ const parser::OmpObjectList *OmpStructureChecker::GetOmpObjectList(
32403266
const parser::OmpClause &clause) {
32413267

32423268
// Clauses with OmpObjectList as its data member
3243-
using MemberObjectListClauses = std::tuple<parser::OmpClause::Copyprivate,
3244-
parser::OmpClause::Copyin, parser::OmpClause::Firstprivate,
3245-
parser::OmpClause::From, parser::OmpClause::Lastprivate,
3246-
parser::OmpClause::Link, parser::OmpClause::Private,
3247-
parser::OmpClause::Shared, parser::OmpClause::To,
3248-
parser::OmpClause::UseDevicePtr, parser::OmpClause::UseDeviceAddr>;
3269+
using MemberObjectListClauses =
3270+
std::tuple<parser::OmpClause::Copyprivate, parser::OmpClause::Copyin,
3271+
parser::OmpClause::Firstprivate, parser::OmpClause::From,
3272+
parser::OmpClause::Lastprivate, parser::OmpClause::Link,
3273+
parser::OmpClause::Private, parser::OmpClause::Shared,
3274+
parser::OmpClause::To, parser::OmpClause::Enter,
3275+
parser::OmpClause::UseDevicePtr, parser::OmpClause::UseDeviceAddr>;
32493276

32503277
// Clauses with OmpObjectList in the tuple
32513278
using TupleObjectListClauses =

‎flang/lib/Semantics/check-omp-structure.h

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -80,6 +80,8 @@ class OmpStructureChecker
8080
void Enter(const parser::OpenMPDeclareTargetConstruct &);
8181
void Leave(const parser::OpenMPDeclareTargetConstruct &);
8282
void Enter(const parser::OmpDeclareTargetWithList &);
83+
void Enter(const parser::OmpDeclareTargetWithClause &);
84+
void Leave(const parser::OmpDeclareTargetWithClause &);
8385
void Enter(const parser::OpenMPExecutableAllocate &);
8486
void Leave(const parser::OpenMPExecutableAllocate &);
8587
void Enter(const parser::OpenMPAllocatorsConstruct &);

‎flang/lib/Semantics/resolve-directives.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1755,6 +1755,9 @@ bool OmpAttributeVisitor::Pre(const parser::OpenMPDeclareTargetConstruct &x) {
17551755
} else if (const auto *linkClause{
17561756
std::get_if<parser::OmpClause::Link>(&clause.u)}) {
17571757
ResolveOmpObjectList(linkClause->v, Symbol::Flag::OmpDeclareTarget);
1758+
} else if (const auto *enterClause{
1759+
std::get_if<parser::OmpClause::Enter>(&clause.u)}) {
1760+
ResolveOmpObjectList(enterClause->v, Symbol::Flag::OmpDeclareTarget);
17581761
}
17591762
}
17601763
}

‎flang/test/Lower/OpenMP/FIR/declare-target-data.f90

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -32,6 +32,10 @@ module test_0
3232
INTEGER :: data_int_to = 5
3333
!$omp declare target to(data_int_to)
3434

35+
!CHECK-DAG: fir.global @_QMtest_0Edata_int_enter {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : i32
36+
INTEGER :: data_int_enter = 5
37+
!$omp declare target enter(data_int_enter)
38+
3539
!CHECK-DAG: fir.global @_QMtest_0Edata_int_clauseless {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>} : i32
3640
INTEGER :: data_int_clauseless = 1
3741
!$omp declare target(data_int_clauseless)
@@ -42,6 +46,12 @@ module test_0
4246
REAL :: data_extended_to_2 = 3
4347
!$omp declare target to(data_extended_to_1, data_extended_to_2)
4448

49+
!CHECK-DAG: fir.global @_QMtest_0Edata_extended_enter_1 {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : f32
50+
!CHECK-DAG: fir.global @_QMtest_0Edata_extended_enter_2 {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : f32
51+
REAL :: data_extended_enter_1 = 2
52+
REAL :: data_extended_enter_2 = 3
53+
!$omp declare target enter(data_extended_enter_1, data_extended_enter_2)
54+
4555
!CHECK-DAG: fir.global @_QMtest_0Edata_extended_link_1 {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (link)>} : f32
4656
!CHECK-DAG: fir.global @_QMtest_0Edata_extended_link_2 {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (link)>} : f32
4757
REAL :: data_extended_link_1 = 2
@@ -69,4 +79,10 @@ PROGRAM commons
6979
REAL :: two_to = 2
7080
COMMON /numbers_to/ one_to, two_to
7181
!$omp declare target to(/numbers_to/)
82+
83+
!CHECK-DAG: fir.global @numbers_enter_ {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : tuple<f32, f32> {
84+
REAL :: one_enter = 1
85+
REAL :: two_enter = 2
86+
COMMON /numbers_enter/ one_enter, two_enter
87+
!$omp declare target enter(/numbers_enter/)
7288
END

‎flang/test/Lower/OpenMP/FIR/declare-target-func-and-subr.f90

Lines changed: 68 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -13,6 +13,14 @@ FUNCTION FUNC_T_DEVICE() RESULT(I)
1313
I = 1
1414
END FUNCTION FUNC_T_DEVICE
1515

16+
! DEVICE-LABEL: func.func @_QPfunc_enter_device()
17+
! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}
18+
FUNCTION FUNC_ENTER_DEVICE() RESULT(I)
19+
!$omp declare target enter(FUNC_ENTER_DEVICE) device_type(nohost)
20+
INTEGER :: I
21+
I = 1
22+
END FUNCTION FUNC_ENTER_DEVICE
23+
1624
! HOST-LABEL: func.func @_QPfunc_t_host()
1725
! HOST-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (to)>{{.*}}
1826
FUNCTION FUNC_T_HOST() RESULT(I)
@@ -21,6 +29,14 @@ FUNCTION FUNC_T_HOST() RESULT(I)
2129
I = 1
2230
END FUNCTION FUNC_T_HOST
2331

32+
! HOST-LABEL: func.func @_QPfunc_enter_host()
33+
! HOST-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (enter)>{{.*}}
34+
FUNCTION FUNC_ENTER_HOST() RESULT(I)
35+
!$omp declare target enter(FUNC_ENTER_HOST) device_type(host)
36+
INTEGER :: I
37+
I = 1
38+
END FUNCTION FUNC_ENTER_HOST
39+
2440
! ALL-LABEL: func.func @_QPfunc_t_any()
2541
! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}
2642
FUNCTION FUNC_T_ANY() RESULT(I)
@@ -29,6 +45,14 @@ FUNCTION FUNC_T_ANY() RESULT(I)
2945
I = 1
3046
END FUNCTION FUNC_T_ANY
3147

48+
! ALL-LABEL: func.func @_QPfunc_enter_any()
49+
! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}
50+
FUNCTION FUNC_ENTER_ANY() RESULT(I)
51+
!$omp declare target enter(FUNC_ENTER_ANY) device_type(any)
52+
INTEGER :: I
53+
I = 1
54+
END FUNCTION FUNC_ENTER_ANY
55+
3256
! ALL-LABEL: func.func @_QPfunc_default_t_any()
3357
! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}
3458
FUNCTION FUNC_DEFAULT_T_ANY() RESULT(I)
@@ -37,6 +61,14 @@ FUNCTION FUNC_DEFAULT_T_ANY() RESULT(I)
3761
I = 1
3862
END FUNCTION FUNC_DEFAULT_T_ANY
3963

64+
! ALL-LABEL: func.func @_QPfunc_default_enter_any()
65+
! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}
66+
FUNCTION FUNC_DEFAULT_ENTER_ANY() RESULT(I)
67+
!$omp declare target enter(FUNC_DEFAULT_ENTER_ANY)
68+
INTEGER :: I
69+
I = 1
70+
END FUNCTION FUNC_DEFAULT_ENTER_ANY
71+
4072
! ALL-LABEL: func.func @_QPfunc_default_any()
4173
! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}
4274
FUNCTION FUNC_DEFAULT_ANY() RESULT(I)
@@ -65,24 +97,48 @@ SUBROUTINE SUBR_T_DEVICE()
6597
!$omp declare target to(SUBR_T_DEVICE) device_type(nohost)
6698
END
6799

100+
! DEVICE-LABEL: func.func @_QPsubr_enter_device()
101+
! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}
102+
SUBROUTINE SUBR_ENTER_DEVICE()
103+
!$omp declare target enter(SUBR_ENTER_DEVICE) device_type(nohost)
104+
END
105+
68106
! HOST-LABEL: func.func @_QPsubr_t_host()
69107
! HOST-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (to)>{{.*}}
70108
SUBROUTINE SUBR_T_HOST()
71109
!$omp declare target to(SUBR_T_HOST) device_type(host)
72110
END
73111

112+
! HOST-LABEL: func.func @_QPsubr_enter_host()
113+
! HOST-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (enter)>{{.*}}
114+
SUBROUTINE SUBR_ENTER_HOST()
115+
!$omp declare target enter(SUBR_ENTER_HOST) device_type(host)
116+
END
117+
74118
! ALL-LABEL: func.func @_QPsubr_t_any()
75119
! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}
76120
SUBROUTINE SUBR_T_ANY()
77121
!$omp declare target to(SUBR_T_ANY) device_type(any)
78122
END
79123

124+
! ALL-LABEL: func.func @_QPsubr_enter_any()
125+
! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}
126+
SUBROUTINE SUBR_ENTER_ANY()
127+
!$omp declare target enter(SUBR_ENTER_ANY) device_type(any)
128+
END
129+
80130
! ALL-LABEL: func.func @_QPsubr_default_t_any()
81131
! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}
82132
SUBROUTINE SUBR_DEFAULT_T_ANY()
83133
!$omp declare target to(SUBR_DEFAULT_T_ANY)
84134
END
85135

136+
! ALL-LABEL: func.func @_QPsubr_default_enter_any()
137+
! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}
138+
SUBROUTINE SUBR_DEFAULT_ENTER_ANY()
139+
!$omp declare target enter(SUBR_DEFAULT_ENTER_ANY)
140+
END
141+
86142
! ALL-LABEL: func.func @_QPsubr_default_any()
87143
! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}
88144
SUBROUTINE SUBR_DEFAULT_ANY()
@@ -108,3 +164,15 @@ RECURSIVE FUNCTION RECURSIVE_DECLARE_TARGET(INCREMENT) RESULT(K)
108164
K = RECURSIVE_DECLARE_TARGET(INCREMENT + 1)
109165
END IF
110166
END FUNCTION RECURSIVE_DECLARE_TARGET
167+
168+
! DEVICE-LABEL: func.func @_QPrecursive_declare_target_enter
169+
! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}
170+
RECURSIVE FUNCTION RECURSIVE_DECLARE_TARGET_ENTER(INCREMENT) RESULT(K)
171+
!$omp declare target enter(RECURSIVE_DECLARE_TARGET_ENTER) device_type(nohost)
172+
INTEGER :: INCREMENT, K
173+
IF (INCREMENT == 10) THEN
174+
K = INCREMENT
175+
ELSE
176+
K = RECURSIVE_DECLARE_TARGET_ENTER(INCREMENT + 1)
177+
END IF
178+
END FUNCTION RECURSIVE_DECLARE_TARGET_ENTER

0 commit comments

Comments
 (0)