Fortran/openmp: Fix '!$omp end'
gcc/fortran/ChangeLog:
* parse.c (decode_omp_directive): Fix permitting 'nowait' for some
combined directives, add missing 'omp end ... loop'.
(gfc_ascii_statement): Fix ST_OMP_END_TEAMS_LOOP result.
* openmp.c (resolve_omp_clauses): Add missing combined loop constructs
case values to the 'if(directive-name: ...)' check.
* trans-openmp.c (gfc_split_omp_clauses): Put nowait on target if
first leaf construct accepting it.
(gfc_trans_omp_parallel_sections, gfc_trans_omp_parallel_workshare):
Unset nowait for parallel if set.
gcc/testsuite/ChangeLog:
* gfortran.dg/gomp/unexpected-end.f90: Update dg-error.
* gfortran.dg/gomp/clauses-1.f90: New test.
* gfortran.dg/gomp/nowait-2.f90: New test.
* gfortran.dg/gomp/nowait-3.f90: New test.
gcc/fortran/openmp.c | 3 +
gcc/fortran/parse.c | 49 +-
gcc/fortran/trans-openmp.c | 5 +
gcc/testsuite/gfortran.dg/gomp/clauses-1.f90 | 667 ++++++++++++++++++++++
gcc/testsuite/gfortran.dg/gomp/nowait-2.f90 | 240 ++++++++
gcc/testsuite/gfortran.dg/gomp/nowait-3.f90 | 151 +++++
gcc/testsuite/gfortran.dg/gomp/unexpected-end.f90 | 12 +-
7 files changed, 1102 insertions(+), 25 deletions(-)
@@ -6232,6 +6232,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
case EXEC_OMP_PARALLEL:
case EXEC_OMP_PARALLEL_DO:
+ case EXEC_OMP_PARALLEL_LOOP:
case EXEC_OMP_PARALLEL_MASKED:
case EXEC_OMP_PARALLEL_MASTER:
case EXEC_OMP_PARALLEL_SECTIONS:
@@ -6285,6 +6286,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
case EXEC_OMP_TARGET:
case EXEC_OMP_TARGET_TEAMS:
case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TARGET_TEAMS_LOOP:
ok = ifc == OMP_IF_TARGET;
break;
@@ -6312,6 +6314,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_TARGET_PARALLEL:
case EXEC_OMP_TARGET_PARALLEL_DO:
+ case EXEC_OMP_TARGET_PARALLEL_LOOP:
ok = ifc == OMP_IF_TARGET || ifc == OMP_IF_PARALLEL;
break;
@@ -915,15 +915,16 @@ decode_omp_directive (void)
matcho ("error", gfc_match_omp_error, ST_OMP_ERROR);
matcho ("end atomic", gfc_match_omp_eos_error, ST_OMP_END_ATOMIC);
matcho ("end critical", gfc_match_omp_end_critical, ST_OMP_END_CRITICAL);
- matchs ("end distribute parallel do simd", gfc_match_omp_eos_error,
+ matchs ("end distribute parallel do simd", gfc_match_omp_end_nowait,
ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD);
- matcho ("end distribute parallel do", gfc_match_omp_eos_error,
+ matcho ("end distribute parallel do", gfc_match_omp_end_nowait,
ST_OMP_END_DISTRIBUTE_PARALLEL_DO);
matchs ("end distribute simd", gfc_match_omp_eos_error,
ST_OMP_END_DISTRIBUTE_SIMD);
matcho ("end distribute", gfc_match_omp_eos_error, ST_OMP_END_DISTRIBUTE);
matchs ("end do simd", gfc_match_omp_end_nowait, ST_OMP_END_DO_SIMD);
matcho ("end do", gfc_match_omp_end_nowait, ST_OMP_END_DO);
+ matcho ("end loop", gfc_match_omp_eos_error, ST_OMP_END_LOOP);
matchs ("end simd", gfc_match_omp_eos_error, ST_OMP_END_SIMD);
matcho ("end masked taskloop simd", gfc_match_omp_eos_error,
ST_OMP_END_MASKED_TASKLOOP_SIMD);
@@ -936,9 +937,12 @@ decode_omp_directive (void)
ST_OMP_END_MASTER_TASKLOOP);
matcho ("end master", gfc_match_omp_eos_error, ST_OMP_END_MASTER);
matchs ("end ordered", gfc_match_omp_eos_error, ST_OMP_END_ORDERED);
- matchs ("end parallel do simd", gfc_match_omp_eos_error,
+ matchs ("end parallel do simd", gfc_match_omp_end_nowait,
ST_OMP_END_PARALLEL_DO_SIMD);
- matcho ("end parallel do", gfc_match_omp_eos_error, ST_OMP_END_PARALLEL_DO);
+ matcho ("end parallel do", gfc_match_omp_end_nowait,
+ ST_OMP_END_PARALLEL_DO);
+ matcho ("end parallel loop", gfc_match_omp_eos_error,
+ ST_OMP_END_PARALLEL_LOOP);
matcho ("end parallel masked taskloop simd", gfc_match_omp_eos_error,
ST_OMP_END_PARALLEL_MASKED_TASKLOOP_SIMD);
matcho ("end parallel masked taskloop", gfc_match_omp_eos_error,
@@ -951,46 +955,53 @@ decode_omp_directive (void)
ST_OMP_END_PARALLEL_MASTER_TASKLOOP);
matcho ("end parallel master", gfc_match_omp_eos_error,
ST_OMP_END_PARALLEL_MASTER);
- matcho ("end parallel sections", gfc_match_omp_eos_error,
+ matcho ("end parallel sections", gfc_match_omp_end_nowait,
ST_OMP_END_PARALLEL_SECTIONS);
- matcho ("end parallel workshare", gfc_match_omp_eos_error,
+ matcho ("end parallel workshare", gfc_match_omp_end_nowait,
ST_OMP_END_PARALLEL_WORKSHARE);
matcho ("end parallel", gfc_match_omp_eos_error, ST_OMP_END_PARALLEL);
matcho ("end scope", gfc_match_omp_end_nowait, ST_OMP_END_SCOPE);
matcho ("end sections", gfc_match_omp_end_nowait, ST_OMP_END_SECTIONS);
matcho ("end single", gfc_match_omp_end_single, ST_OMP_END_SINGLE);
matcho ("end target data", gfc_match_omp_eos_error, ST_OMP_END_TARGET_DATA);
- matchs ("end target parallel do simd", gfc_match_omp_eos_error,
+ matchs ("end target parallel do simd", gfc_match_omp_end_nowait,
ST_OMP_END_TARGET_PARALLEL_DO_SIMD);
- matcho ("end target parallel do", gfc_match_omp_eos_error,
+ matcho ("end target parallel do", gfc_match_omp_end_nowait,
ST_OMP_END_TARGET_PARALLEL_DO);
- matcho ("end target parallel", gfc_match_omp_eos_error,
+ matcho ("end target parallel loop", gfc_match_omp_end_nowait,
+ ST_OMP_END_TARGET_PARALLEL_LOOP);
+ matcho ("end target parallel", gfc_match_omp_end_nowait,
ST_OMP_END_TARGET_PARALLEL);
- matchs ("end target simd", gfc_match_omp_eos_error, ST_OMP_END_TARGET_SIMD);
+ matchs ("end target simd", gfc_match_omp_end_nowait, ST_OMP_END_TARGET_SIMD);
matchs ("end target teams distribute parallel do simd",
- gfc_match_omp_eos_error,
+ gfc_match_omp_end_nowait,
ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD);
- matcho ("end target teams distribute parallel do", gfc_match_omp_eos_error,
+ matcho ("end target teams distribute parallel do",
+ gfc_match_omp_end_nowait,
ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO);
- matchs ("end target teams distribute simd", gfc_match_omp_eos_error,
+ matchs ("end target teams distribute simd", gfc_match_omp_end_nowait,
ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD);
- matcho ("end target teams distribute", gfc_match_omp_eos_error,
+ matcho ("end target teams distribute", gfc_match_omp_end_nowait,
ST_OMP_END_TARGET_TEAMS_DISTRIBUTE);
- matcho ("end target teams", gfc_match_omp_eos_error, ST_OMP_END_TARGET_TEAMS);
- matcho ("end target", gfc_match_omp_eos_error, ST_OMP_END_TARGET);
+ matcho ("end target teams loop", gfc_match_omp_end_nowait,
+ ST_OMP_END_TARGET_TEAMS_LOOP);
+ matcho ("end target teams", gfc_match_omp_end_nowait,
+ ST_OMP_END_TARGET_TEAMS);
+ matcho ("end target", gfc_match_omp_end_nowait, ST_OMP_END_TARGET);
matcho ("end taskgroup", gfc_match_omp_eos_error, ST_OMP_END_TASKGROUP);
matchs ("end taskloop simd", gfc_match_omp_eos_error,
ST_OMP_END_TASKLOOP_SIMD);
matcho ("end taskloop", gfc_match_omp_eos_error, ST_OMP_END_TASKLOOP);
matcho ("end task", gfc_match_omp_eos_error, ST_OMP_END_TASK);
- matchs ("end teams distribute parallel do simd", gfc_match_omp_eos_error,
+ matchs ("end teams distribute parallel do simd", gfc_match_omp_end_nowait,
ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD);
- matcho ("end teams distribute parallel do", gfc_match_omp_eos_error,
+ matcho ("end teams distribute parallel do", gfc_match_omp_end_nowait,
ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO);
matchs ("end teams distribute simd", gfc_match_omp_eos_error,
ST_OMP_END_TEAMS_DISTRIBUTE_SIMD);
matcho ("end teams distribute", gfc_match_omp_eos_error,
ST_OMP_END_TEAMS_DISTRIBUTE);
+ matcho ("end teams loop", gfc_match_omp_eos_error, ST_OMP_END_TEAMS_LOOP);
matcho ("end teams", gfc_match_omp_eos_error, ST_OMP_END_TEAMS);
matcho ("end workshare", gfc_match_omp_end_nowait,
ST_OMP_END_WORKSHARE);
@@ -2553,7 +2564,7 @@ gfc_ascii_statement (gfc_statement st)
p = "!$OMP END TEAMS DISTRIBUTE SIMD";
break;
case ST_OMP_END_TEAMS_LOOP:
- p = "!$OMP END TEAMS LOP";
+ p = "!$OMP END TEAMS LOOP";
break;
case ST_OMP_END_WORKSHARE:
p = "!$OMP END WORKSHARE";
@@ -5878,6 +5878,9 @@ gfc_split_omp_clauses (gfc_code *code,
/* And this is copied to all. */
clausesa[GFC_OMP_SPLIT_TARGET].if_expr
= code->ext.omp_clauses->if_expr;
+ if (is_loop || !(mask & GFC_OMP_SPLIT_DO))
+ clausesa[GFC_OMP_SPLIT_TARGET].nowait
+ = code->ext.omp_clauses->nowait;
}
if (mask & GFC_OMP_MASK_TEAMS)
{
@@ -6296,6 +6299,7 @@ gfc_trans_omp_parallel_sections (gfc_code *code)
memset (§ion_clauses, 0, sizeof (section_clauses));
section_clauses.nowait = true;
+ code->ext.omp_clauses->nowait = false;
gfc_start_block (&block);
omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
@@ -6322,6 +6326,7 @@ gfc_trans_omp_parallel_workshare (gfc_code *code)
memset (&workshare_clauses, 0, sizeof (workshare_clauses));
workshare_clauses.nowait = true;
+ code->ext.omp_clauses->nowait = false;
gfc_start_block (&block);
omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
new file mode 100644
@@ -0,0 +1,667 @@
+! { dg-do compile }
+
+module m
+ use iso_c_binding, only: c_intptr_t
+ implicit none (external, type)
+
+ integer(c_intptr_t), parameter :: &
+ omp_null_allocator = 0, &
+ omp_default_mem_alloc = 1, &
+ omp_large_cap_mem_alloc = 2, &
+ omp_const_mem_alloc = 3, &
+ omp_high_bw_mem_alloc = 4, &
+ omp_low_lat_mem_alloc = 5, &
+ omp_cgroup_mem_alloc = 6, &
+ omp_pteam_mem_alloc = 7, &
+ omp_thread_mem_alloc = 8
+
+ integer, parameter :: &
+ omp_allocator_handle_kind = c_intptr_t
+
+ integer :: t
+ !$omp threadprivate (t)
+
+ integer :: f, l, ll, r, r2
+ !$omp declare target (f, l, ll, r, r2)
+
+contains
+
+subroutine foo (d, m, i1, i2, i3, p, idp, s, nte, tl, nth, g, nta, fi, pp, q, dd, ntm)
+ !$omp declare target (foo)
+ integer :: d, m, p, idp, s, nte, tl, nth, g, nta, pp, q, dd, ntm
+ logical :: i1, i2, i3, fi
+ pointer :: q
+ integer :: i
+
+ !$omp distribute parallel do &
+ !$omp& private (p) firstprivate (f) collapse(1) dist_schedule(static, 16) &
+ !$omp& if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4) order(concurrent)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc:f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp distribute parallel do simd &
+ !$omp& private (p) firstprivate (f) collapse(1) dist_schedule(static, 16) &
+ !$omp& if (parallel: i2) if(simd: i1) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4) nontemporal(ntm) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) order(concurrent)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc:f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp distribute simd &
+ !$omp& private (p) firstprivate (f) collapse(1) dist_schedule(static, 16) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) reduction(+:r) if(i1) nontemporal(ntm) &
+ !$omp& order(concurrent)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc:f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+end
+
+subroutine qux (p)
+ !$omp declare target (qux)
+ integer, value :: p
+
+ !$omp loop bind(teams) order(concurrent) &
+ !$omp& private (p) lastprivate (l) collapse(1) reduction(+:r)
+ do l = 1, 64
+ ll = ll + 1
+ end do
+end
+
+subroutine baz (d, m, i1, i2, i3, p, idp, s, nte, tl, nth, g, nta, fi, pp, q, dd, ntm)
+ integer :: d, m, p, idp, s, nte, tl, nth, g, nta, pp, q, dd, ntm
+ logical :: i1, i2, i3, fi
+ pointer :: q
+ integer :: i
+ !$omp distribute parallel do &
+ !$omp& private (p) firstprivate (f) collapse(1) dist_schedule(static, 16) &
+ !$omp& if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4) copyin(t)
+ ! FIXME/TODO: allocate (p)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp distribute parallel do &
+ !$omp& private (p) firstprivate (f) collapse(1) dist_schedule(static, 16) &
+ !$omp& if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4) order(concurrent)
+ ! FIXME/TODO: allocate (p)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp distribute parallel do simd &
+ !$omp& private (p) firstprivate (f) collapse(1) dist_schedule(static, 16) &
+ !$omp& if (parallel: i2) if(simd: i1) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4) nontemporal(ntm) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) copyin(t)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+
+ !$omp distribute parallel do simd &
+ !$omp& private (p) firstprivate (f) collapse(1) dist_schedule(static, 16) &
+ !$omp& if (parallel: i2) if(simd: i1) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4) nontemporal(ntm) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) order(concurrent)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+
+ !$omp distribute simd &
+ !$omp& private (p) firstprivate (f) collapse(1) dist_schedule(static, 16) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) reduction(+:r) if(i1) nontemporal(ntm) &
+ !$omp& order(concurrent)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+
+ !$omp loop bind(parallel) order(concurrent) &
+ !$omp& private (p) lastprivate (l) collapse(1) reduction(+:r)
+ do l = 1, 64
+ ll = ll + 1
+ end do
+end
+
+subroutine bar (d, m, i1, i2, i3, p, idp, s, nte, tl, nth, g, nta, fi, pp, q, dd, ntm)
+ integer :: d, m, p, idp, s, nte, tl, nth, g, nta, pp, q, dd(0:5), ntm
+ logical :: i1, i2, i3, fi
+ pointer :: q
+ integer :: i
+
+ !$omp do simd &
+ !$omp& private (p) firstprivate (f) lastprivate (l) linear (ll:1) reduction(+:r) schedule(static, 4) collapse(1) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) nontemporal(ntm) if(i1) order(concurrent)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+ !$omp end do simd
+
+ !$omp parallel do &
+ !$omp& private (p) firstprivate (f) if (parallel: i2) default(shared) shared(s) copyin(t) reduction(+:r) num_threads (nth) &
+ !$omp& proc_bind(spread) lastprivate (l) linear (ll:1) ordered schedule(static, 4) collapse(1)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+
+ !$omp parallel do &
+ !$omp& private (p) firstprivate (f) if (parallel: i2) default(shared) shared(s) copyin(t) reduction(+:r) num_threads (nth) &
+ !$omp& proc_bind(spread) lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) order(concurrent)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+
+ !$omp parallel do simd &
+ !$omp& private (p) firstprivate (f) if (i2) default(shared) shared(s) copyin(t) reduction(+:r) num_threads (nth) &
+ !$omp& proc_bind(spread) lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) nontemporal(ntm) order(concurrent)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+
+ !$omp parallel sections &
+ !$omp& private (p) firstprivate (f) if (parallel: i2) default(shared) shared(s) copyin(t) reduction(+:r) num_threads (nth) &
+ !$omp& proc_bind(spread) lastprivate (l)
+ ! FIXME/TODO: allocate (f)
+ !$omp section
+ block; end block
+ !$omp section
+ block; end block
+ !$omp end parallel sections
+
+ !$omp target parallel &
+ !$omp& device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) &
+ !$omp& if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) &
+ !$omp& depend(inout: dd(0)) in_reduction(+:r2)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc:f)
+ !$omp end target parallel nowait
+
+ !$omp target parallel do &
+ !$omp& device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) &
+ !$omp& if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) linear (ll:1) ordered schedule(static, 4) collapse(1) depend(inout: dd(0)) &
+ !$omp& in_reduction(+:r2)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc:f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+ !$omp end target parallel do nowait
+
+ !$omp target parallel do &
+ !$omp& device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) &
+ !$omp& if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) depend(inout: dd(0)) order(concurrent) &
+ !$omp& in_reduction(+:r2)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc:f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+ !$omp end target parallel do nowait
+
+ !$omp target parallel do simd &
+ !$omp& device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) &
+ !$omp& if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) depend(inout: dd(0)) nontemporal(ntm) if (simd: i3) order(concurrent) &
+ !$omp& in_reduction(+:r2)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc:f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+ !$omp end target parallel do simd nowait
+
+ !$omp target teams &
+ !$omp& device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) &
+ !$omp& shared(s) default(shared) reduction(+:r) num_teams(nte - 1:nte) thread_limit(tl) depend(inout: dd(0)) &
+ !$omp& in_reduction(+:r2)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc:f)
+ !$omp end target teams nowait
+
+ !$omp target teams distribute &
+ !$omp& device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) &
+ !$omp& shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) order(concurrent) &
+ !$omp& collapse(1) dist_schedule(static, 16) depend(inout: dd(0)) in_reduction(+:r2)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc:f)
+ do i = 1, 64
+ end do
+ !$omp end target teams distribute nowait
+
+ !$omp target teams distribute parallel do &
+ !$omp& device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) &
+ !$omp& shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) &
+ !$omp& collapse(1) dist_schedule(static, 16) &
+ !$omp& if (parallel: i2) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4) depend(inout: dd(0)) order(concurrent) &
+ !$omp& in_reduction(+:r2)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc:f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+ !$omp end target teams distribute parallel do nowait
+
+ !$omp target teams distribute parallel do simd &
+ !$omp& device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) &
+ !$omp& shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) &
+ !$omp& collapse(1) dist_schedule(static, 16) &
+ !$omp& if (parallel: i2) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4) order(concurrent) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) depend(inout: dd(0)) nontemporal(ntm) if (simd: i3) &
+ !$omp& in_reduction(+:r2)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc:f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+ !$omp end target teams distribute parallel do simd
+
+ !$omp target teams distribute simd &
+ !$omp& device(d) map (tofrom: m) if (i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) &
+ !$omp& shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) &
+ !$omp& collapse(1) dist_schedule(static, 16) order(concurrent) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) depend(inout: dd(0)) nontemporal(ntm) &
+ !$omp& in_reduction(+:r2)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc:f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+ !$omp end target teams distribute simd
+
+ !$omp target simd &
+ !$omp& device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) &
+ !$omp& safelen(8) simdlen(4) lastprivate (l) linear(ll: 1) aligned(q: 32) reduction(+:r) &
+ !$omp& depend(inout: dd(0)) nontemporal(ntm) if(simd:i3) order(concurrent) &
+ !$omp& in_reduction(+:r2)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc:f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+ !$omp end target simd
+
+ !$omp taskgroup task_reduction(+:r2)
+ ! FIXME/TODO: allocate (r2)
+ !$omp taskloop simd &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) grainsize (g) collapse(1) untied &
+ !$omp& if(taskloop: i1) if(simd: i2) final(fi) mergeable priority (pp) &
+ !$omp& safelen(8) simdlen(4) linear(ll: 1) aligned(q: 32) reduction(default, +:r) in_reduction(+:r2) nontemporal(ntm) &
+ !$omp& order(concurrent)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+ !$omp end taskgroup
+
+ !$omp taskgroup task_reduction(+:r)
+ ! FIXME/TODO: allocate (r)
+ !$omp taskloop simd &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) grainsize (g) collapse(1) untied if(i1) &
+ !$omp& final(fi) mergeable nogroup priority (pp) &
+ !$omp& safelen(8) simdlen(4) linear(ll: 1) aligned(q: 32) in_reduction(+:r) nontemporal(ntm) &
+ !$omp& order(concurrent)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+ !$omp end taskgroup
+
+ !$omp taskwait
+ !$omp taskloop simd &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) num_tasks (nta) collapse(1) if(taskloop: i1) &
+ !$omp& final(fi) priority (pp) safelen(8) simdlen(4) linear(ll: 1) aligned(q: 32) reduction(+:r) if (simd: i3) nontemporal(ntm) &
+ !$omp& order(concurrent)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll + 1
+ end do
+
+ !$omp target depend(inout: dd(0)) in_reduction(+:r2)
+ !$omp teams distribute &
+ !$omp& private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) &
+ !$omp& collapse(1) dist_schedule(static, 16) order(concurrent)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc: f)
+ do i = 1, 64
+ end do
+ !$omp end target nowait
+
+ !$omp target
+ !$omp teams distribute parallel do &
+ !$omp& private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) &
+ !$omp& collapse(1) dist_schedule(static, 16) &
+ !$omp& if (parallel: i2) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4) order(concurrent)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc: f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+ !$omp end target
+
+ !$omp target
+ !$omp teams distribute parallel do simd &
+ !$omp& private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) &
+ !$omp& collapse(1) dist_schedule(static, 16) &
+ !$omp& if (parallel: i2) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4) order(concurrent) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) if (simd: i3) nontemporal(ntm)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc: f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+ !$omp end target
+
+ !$omp target
+ !$omp teams distribute simd &
+ !$omp& private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) &
+ !$omp& collapse(1) dist_schedule(static, 16) order(concurrent) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) if(i3) nontemporal(ntm)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc: f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+ !$omp end target
+
+ !$omp teams distribute parallel do &
+ !$omp& private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) &
+ !$omp& collapse(1) dist_schedule(static, 16) &
+ !$omp& if (parallel: i2) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4) copyin(t)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp teams distribute parallel do &
+ !$omp& private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) &
+ !$omp& collapse(1) dist_schedule(static, 16) order(concurrent) &
+ !$omp& if (parallel: i2) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp teams distribute parallel do simd &
+ !$omp& private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) &
+ !$omp& collapse(1) dist_schedule(static, 16) &
+ !$omp& if (parallel: i2) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) if (simd: i3) nontemporal(ntm) copyin(t)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp teams distribute parallel do simd &
+ !$omp& private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) &
+ !$omp& collapse(1) dist_schedule(static, 16) &
+ !$omp& if (parallel: i2) num_threads (nth) proc_bind(spread) &
+ !$omp& lastprivate (l) schedule(static, 4) order(concurrent) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) if (simd: i3) nontemporal(ntm)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp teams distribute simd &
+ !$omp& private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) &
+ !$omp& collapse(1) dist_schedule(static, 16) order(concurrent) &
+ !$omp& safelen(8) simdlen(4) aligned(q: 32) if(i3) nontemporal(ntm)
+ ! FIXME/TODO: allocate(f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp parallel master &
+ !$omp& private (p) firstprivate (f) if (parallel: i2) default(shared) shared(s) reduction(+:r) &
+ !$omp& num_threads (nth) proc_bind(spread) copyin(t)
+ ! FIXME/TODO: allocate (f)
+ !$omp end parallel master
+
+ !$omp parallel masked &
+ !$omp& private (p) firstprivate (f) if (parallel: i2) default(shared) shared(s) reduction(+:r) &
+ !$omp& num_threads (nth) proc_bind(spread) copyin(t) filter (d)
+ ! FIXME/TODO: allocate (f)
+ !$omp end parallel masked
+
+ !$omp taskgroup task_reduction (+:r2)
+ ! FIXME/TODO: allocate (r2)
+ !$omp master taskloop &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) grainsize (g) collapse(1) untied &
+ !$omp& if(taskloop: i1) final(fi) mergeable priority (pp) &
+ !$omp& reduction(default, +:r) in_reduction(+:r2)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+ !$omp end taskgroup
+
+ !$omp taskgroup task_reduction (+:r2)
+ ! FIXME/TODO: allocate (r2)
+ !$omp masked taskloop &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) grainsize (g) collapse(1) untied &
+ !$omp& if(taskloop: i1) final(fi) mergeable priority (pp) reduction(default, +:r) in_reduction(+:r2) filter (d)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+ !$omp end taskgroup
+
+ !$omp taskgroup task_reduction (+:r2)
+ ! FIXME/TODO: allocate (r2)
+ !$omp master taskloop simd &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) grainsize (g) collapse(1) untied &
+ !$omp& if(taskloop: i1) if(simd: i2) final(fi) mergeable priority (pp) &
+ !$omp& safelen(8) simdlen(4) linear(ll: 1) aligned(q: 32) reduction(default, +:r) in_reduction(+:r2) nontemporal(ntm) &
+ !$omp& order(concurrent)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+ !$omp end taskgroup
+
+ !$omp taskgroup task_reduction (+:r2)
+ ! FIXME/TODO: allocate (r2)
+ !$omp masked taskloop simd &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) grainsize (g) collapse(1) untied &
+ !$omp& if(taskloop: i1) if(simd: i2) final(fi) mergeable priority (pp) &
+ !$omp& safelen(8) simdlen(4) linear(ll: 1) aligned(q: 32) reduction(default, +:r) in_reduction(+:r2) nontemporal(ntm) &
+ !$omp& order(concurrent) filter (d)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+ !$omp end taskgroup
+
+ !$omp parallel master taskloop &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) grainsize (g) collapse(1) untied &
+ !$omp& if(taskloop: i1) final(fi) mergeable priority (pp) &
+ !$omp& reduction(default, +:r) if (parallel: i2) num_threads (nth) proc_bind(spread) copyin(t)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp parallel masked taskloop &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) grainsize (g) collapse(1) untied &
+ !$omp& if(taskloop: i1) final(fi) mergeable priority (pp) &
+ !$omp& reduction(default, +:r) if (parallel: i2) num_threads (nth) proc_bind(spread) copyin(t) filter (d)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp parallel master taskloop simd &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) grainsize (g) collapse(1) untied &
+ !$omp& if(taskloop: i1) if(simd: i2) final(fi) mergeable priority (pp) &
+ !$omp& safelen(8) simdlen(4) linear(ll: 1) aligned(q: 32) reduction(default, +:r) nontemporal(ntm) if (parallel: i2) &
+ !$omp& num_threads (nth) proc_bind(spread) copyin(t) order(concurrent)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp parallel masked taskloop simd &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) grainsize (g) collapse(1) untied &
+ !$omp& if(taskloop: i1) if(simd: i2) final(fi) mergeable priority (pp) &
+ !$omp& safelen(8) simdlen(4) linear(ll: 1) aligned(q: 32) reduction(default, +:r) nontemporal(ntm) if (parallel: i2) &
+ !$omp& num_threads (nth) proc_bind(spread) copyin(t) order(concurrent) filter (d)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp taskgroup task_reduction (+:r2)
+ ! FIXME/TODO: allocate (r2)
+ !$omp master taskloop &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) num_tasks (nta) collapse(1) &
+ !$omp& untied if(i1) final(fi) mergeable priority (pp) reduction(default, +:r) in_reduction(+:r2)
+ do i = 1, 64
+ ll = ll +1
+ end do
+ !$omp end taskgroup
+
+ !$omp taskgroup task_reduction (+:r2)
+ ! FIXME/TODO: allocate (r2)
+ !$omp masked taskloop &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) num_tasks (nta) collapse(1) &
+ !$omp& untied if(i1) final(fi) mergeable priority (pp) reduction(default, +:r) in_reduction(+:r2) filter (d)
+ do i = 1, 64
+ ll = ll +1
+ end do
+ !$omp end taskgroup
+
+ !$omp taskgroup task_reduction (+:r2)
+ ! FIXME/TODO: allocate (r2)
+ !$omp master taskloop simd &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) num_tasks (nta) collapse(1) untied if(i1) &
+ !$omp& final(fi) mergeable priority (pp) safelen(8) simdlen(4) linear(ll: 1) aligned(q: 32) reduction(default, +:r) &
+ !$omp& in_reduction(+:r2) nontemporal(ntm) order(concurrent)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+ !$omp end taskgroup
+
+ !$omp taskgroup task_reduction (+:r2)
+ ! FIXME/TODO: allocate (r2)
+ !$omp masked taskloop simd &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) num_tasks (nta) collapse(1) untied &
+ !$omp& if(i1) final(fi) mergeable priority (pp) safelen(8) simdlen(4) linear(ll: 1) aligned(q: 32) reduction(default, +:r) &
+ !$omp& in_reduction(+:r2) nontemporal(ntm) order(concurrent) filter (d)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+ !$omp end taskgroup
+
+ !$omp parallel master taskloop &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) num_tasks (nta) collapse(1) untied &
+ !$omp& if(i1) final(fi) mergeable priority (pp) reduction(default, +:r) num_threads (nth) proc_bind(spread) copyin(t)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp parallel masked taskloop &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) num_tasks (nta) collapse(1) untied &
+ !$omp& if(i1) final(fi) mergeable priority (pp) reduction(default, +:r) num_threads (nth) proc_bind(spread) &
+ !$omp& copyin(t) filter (d)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp parallel master taskloop simd &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) num_tasks (nta) collapse(1) untied &
+ !$omp& if(i1) final(fi) mergeable priority (pp) safelen(8) simdlen(4) linear(ll: 1) aligned(q: 32) reduction(default, +:r) &
+ !$omp& nontemporal(ntm) num_threads (nth) proc_bind(spread)copyin(t) order(concurrent)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp parallel masked taskloop simd &
+ !$omp& private (p) firstprivate (f) lastprivate (l) shared (s) default(shared) num_tasks (nta) collapse(1) untied if(i1) &
+ !$omp& final(fi) mergeable priority (pp) safelen(8) simdlen(4) linear(ll: 1) aligned(q: 32) reduction(default, +:r) &
+ !$omp& nontemporal(ntm) num_threads (nth) proc_bind(spread) copyin(t) order(concurrent) filter (d)
+ ! FIXME/TODO: allocate (f)
+ do i = 1, 64
+ ll = ll +1
+ end do
+
+ !$omp loop bind(thread) order(concurrent) &
+ !$omp& private (p) lastprivate (l) collapse(1) reduction(+:r)
+ do l = 1, 64
+ ll = ll + 1
+ end do
+
+ !$omp parallel loop &
+ !$omp& private (p) firstprivate (f) default(shared) shared(s) copyin(t) reduction(+:r) num_threads (nth) &
+ !$omp& proc_bind(spread) lastprivate (l) collapse(1) bind(parallel) order(concurrent) if (parallel: i2)
+ ! FIXME/TODO: allocate (f)
+ do l = 1, 64
+ ll = ll + 1
+ end do
+
+ !$omp parallel loop &
+ !$omp& private (p) firstprivate (f) default(shared) shared(s) copyin(t) reduction(+:r) num_threads (nth) &
+ !$omp& proc_bind(spread) lastprivate (l) collapse(1) if (parallel: i2)
+ ! FIXME/TODO: allocate (f)
+ do l = 1, 64
+ ll = ll + 1
+ end do
+
+ !$omp teams loop &
+ !$omp& private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) &
+ !$omp& collapse(1) lastprivate (l) bind(teams)
+ ! FIXME/TODO: allocate (f)
+ do l = 1, 64
+ end do
+
+ !$omp teams loop &
+ !$omp& private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) &
+ !$omp& collapse(1) lastprivate (l) order(concurrent)
+ ! FIXME/TODO: allocate (f)
+ do l = 1, 64
+ end do
+
+ !$omp target parallel loop &
+ !$omp& device(d) map (tofrom: m) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) &
+ !$omp& default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) &
+ !$omp& depend(inout: dd(0)) lastprivate (l) order(concurrent) collapse(1) in_reduction(+:r2) &
+ !$omp& if (target: i1) if (parallel: i2)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc: f)
+ do l = 1, 64
+ end do
+ !$omp end target parallel loop
+
+ !$omp target teams loop &
+ !$omp& device(d) map (tofrom: m) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) &
+ !$omp& shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) depend(inout: dd(0)) &
+ !$omp& lastprivate (l) bind(teams) collapse(1) in_reduction(+:r2) if (target: i1)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc: f)
+ do l = 1, 64
+ end do
+ !$omp end target teams loop
+
+ !$omp target teams loop &
+ !$omp& device(d) map (tofrom: m) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) &
+ !$omp& shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) depend(inout: dd(0)) &
+ !$omp& lastprivate (l) order(concurrent) collapse(1) in_reduction(+:r2) if (target: i1)
+ ! FIXME/TODO: allocate (omp_default_mem_alloc: f)
+ do l = 1, 64
+ end do
+ !$omp end target teams loop
+
+end
+end module
new file mode 100644
@@ -0,0 +1,240 @@
+! Cross check that it is accepted without nowait
+subroutine bar()
+implicit none
+integer :: i
+!$omp atomic write
+i = 5
+!$omp end atomic
+
+!$omp critical
+!$omp end critical
+
+!$omp distribute
+do i = 1, 5
+end do
+!$omp end distribute
+
+!$omp distribute simd
+do i = 1, 5
+end do
+!$omp end distribute simd
+
+!$omp masked
+!$omp end masked
+
+!$omp masked taskloop
+do i = 1, 5
+end do
+!$omp end masked taskloop
+
+!$omp masked taskloop simd
+do i = 1, 5
+end do
+!$omp end masked taskloop simd
+
+!$omp master
+!$omp end master
+
+!$omp master taskloop
+do i = 1, 5
+end do
+!$omp end master taskloop
+
+!$omp master taskloop simd
+do i = 1, 5
+end do
+!$omp end master taskloop simd
+
+!$omp ordered
+!$omp end ordered
+
+!$omp parallel
+!$omp end parallel
+
+!$omp parallel masked
+!$omp end parallel masked
+
+!$omp parallel masked taskloop
+do i = 1, 5
+end do
+!$omp end parallel masked taskloop
+
+!$omp parallel masked taskloop simd
+do i = 1, 5
+end do
+!$omp end parallel masked taskloop simd
+
+!$omp parallel master
+!$omp end parallel master
+
+!$omp parallel master taskloop
+do i = 1, 5
+end do
+!$omp end parallel master taskloop
+
+!$omp parallel master taskloop simd
+do i = 1, 5
+end do
+!$omp end parallel master taskloop simd
+
+!$omp parallel sections
+!$omp end parallel sections
+
+!$omp simd
+do i = 1, 5
+end do
+!$omp end simd
+
+!$omp task
+!$omp end task
+
+!$omp taskgroup
+!$omp end taskgroup
+
+!$omp taskloop
+do i = 1, 5
+end do
+!$omp end taskloop
+
+!$omp taskloop simd
+do i = 1, 5
+end do
+!$omp end taskloop simd
+
+!$omp teams
+!$omp end teams
+
+!$omp teams distribute
+do i = 1, 5
+end do
+!$omp end teams distribute
+
+!$omp teams distribute simd
+do i = 1, 5
+end do
+!$omp end teams distribute simd
+
+!$omp target data map(tofrom:i)
+!$omp end target data
+
+end
+
+! invalid nowait
+
+subroutine foo
+implicit none
+integer :: i
+!$omp atomic write
+i = 5
+!$omp end atomic nowait ! { dg-error "Unexpected junk" }
+
+!$omp critical
+!$omp end critical nowait ! { dg-error "Unexpected junk" }
+
+!$omp distribute
+do i = 1, 5
+end do
+!$omp end distribute nowait ! { dg-error "Unexpected junk" }
+
+!$omp distribute simd
+do i = 1, 5
+end do
+!$omp end distribute simd nowait ! { dg-error "Unexpected junk" }
+
+!$omp masked
+!$omp end masked nowait ! { dg-error "Unexpected junk" }
+
+!$omp masked taskloop
+do i = 1, 5
+end do
+!$omp end masked taskloop nowait ! { dg-error "Unexpected junk" }
+
+!$omp masked taskloop simd
+do i = 1, 5
+end do
+!$omp end masked taskloop simd nowait ! { dg-error "Unexpected junk" }
+
+!$omp master
+!$omp end master nowait ! { dg-error "Unexpected junk" }
+
+!$omp master taskloop
+do i = 1, 5
+end do
+!$omp end master taskloop nowait ! { dg-error "Unexpected junk" }
+
+!$omp master taskloop simd
+do i = 1, 5
+end do
+!$omp end master taskloop simd nowait ! { dg-error "Unexpected junk" }
+
+!$omp ordered
+!$omp end ordered nowait ! { dg-error "Unexpected junk" }
+
+!$omp parallel
+!$omp end parallel nowait ! { dg-error "Unexpected junk" }
+
+!$omp parallel masked
+!$omp end parallel masked nowait ! { dg-error "Unexpected junk" }
+
+!$omp parallel masked taskloop
+do i = 1, 5
+end do
+!$omp end parallel masked taskloop nowait ! { dg-error "Unexpected junk" }
+
+!$omp parallel masked taskloop simd
+do i = 1, 5
+end do
+!$omp end parallel masked taskloop simd nowait ! { dg-error "Unexpected junk" }
+
+!$omp parallel master
+!$omp end parallel master nowait ! { dg-error "Unexpected junk" }
+
+!$omp parallel master taskloop
+do i = 1, 5
+end do
+!$omp end parallel master taskloop nowait ! { dg-error "Unexpected junk" }
+
+!$omp parallel master taskloop simd
+do i = 1, 5
+end do
+!$omp end parallel master taskloop simd nowait ! { dg-error "Unexpected junk" }
+
+!$omp simd
+do i = 1, 5
+end do
+!$omp end simd nowait ! { dg-error "Unexpected junk" }
+
+!$omp task
+!$omp end task nowait ! { dg-error "Unexpected junk" }
+
+!$omp taskgroup
+!$omp end taskgroup nowait ! { dg-error "Unexpected junk" }
+
+!$omp taskloop
+do i = 1, 5
+end do
+!$omp end taskloop nowait ! { dg-error "Unexpected junk" }
+
+!$omp taskloop simd
+do i = 1, 5
+end do
+!$omp end taskloop simd nowait ! { dg-error "Unexpected junk" }
+
+!$omp teams
+!$omp end teams nowait ! { dg-error "Unexpected junk" }
+
+!$omp teams distribute
+do i = 1, 5
+end do
+!$omp end teams distribute nowait ! { dg-error "Unexpected junk" }
+
+!$omp teams distribute simd
+do i = 1, 5
+end do
+!$omp end teams distribute simd
+
+!$omp target data map(tofrom:i)
+!$omp end target data nowait ! { dg-error "Unexpected junk" }
+
+end ! { dg-error "Unexpected END statement" }
+! { dg-prune-output "Unexpected end of file" }
new file mode 100644
@@ -0,0 +1,151 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+subroutine foo
+implicit none
+integer :: i, a(5)
+
+!$omp distribute parallel do
+do i = 1, 5
+end do
+!$omp end distribute parallel do nowait
+
+!$omp distribute parallel do simd
+do i = 1, 5
+end do
+!$omp end distribute parallel do simd nowait
+
+!$omp do
+do i = 1, 5
+end do
+!$omp end do nowait
+
+!$omp do simd
+do i = 1, 5
+end do
+!$omp end do simd nowait
+
+!$omp parallel do
+do i = 1, 5
+end do
+!$omp end parallel do nowait
+
+!$omp parallel sections
+ !$omp section
+ block; end block
+!$omp end parallel sections nowait
+
+!$omp parallel do simd
+do i = 1, 5
+end do
+!$omp end parallel do simd nowait
+
+!$omp parallel workshare
+a(:) = 5
+!$omp end parallel workshare nowait
+
+!$omp scope
+!$omp end scope nowait
+
+!$omp sections
+ !$omp section
+ block; end block
+!$omp end sections nowait
+
+!$omp single
+!$omp end single nowait
+
+!$omp target
+!$omp end target nowait
+
+!$omp target parallel
+!$omp end target parallel nowait
+
+!$omp target parallel do
+do i = 1, 5
+end do
+!$omp end target parallel do
+
+!$omp target parallel do simd
+do i = 1, 5
+end do
+!$omp end target parallel do simd nowait
+
+!$omp target parallel loop
+do i = 1, 5
+end do
+!$omp end target parallel loop nowait
+
+!$omp target simd
+do i = 1, 5
+end do
+!$omp end target simd nowait
+
+!$omp target teams
+!$omp end target teams nowait
+
+!$omp target teams distribute
+do i = 1, 5
+end do
+!$omp end target teams distribute nowait
+
+!$omp target teams distribute parallel do
+do i = 1, 5
+end do
+!$omp end target teams distribute parallel do nowait
+
+!$omp target teams distribute parallel do simd
+do i = 1, 5
+end do
+!$omp end target teams distribute parallel do simd nowait
+
+!$omp target teams distribute simd
+do i = 1, 5
+end do
+!$omp end target teams distribute simd nowait
+
+!$omp target teams loop
+do i = 1, 5
+end do
+!$omp end target teams loop nowait
+
+!$omp teams distribute parallel do
+do i = 1, 5
+end do
+!$omp end teams distribute parallel do nowait
+
+!$omp teams distribute parallel do simd
+do i = 1, 5
+end do
+!$omp end teams distribute parallel do simd nowait
+
+!$omp workshare
+A(:) = 5
+!$omp end workshare nowait
+end
+
+! Expected with 'nowait'
+
+! { dg-final { scan-tree-dump-times "#pragma omp for nowait" 12 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp for schedule\\(static\\) nowait" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp sections nowait" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp single nowait" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target nowait" 7 "original" } }
+
+! Never:
+
+! { dg-final { scan-tree-dump-not "#pragma omp distribute\[^\n\r]*nowait" "original" } }
+! { dg-final { scan-tree-dump-not "#pragma omp loop\[^\n\r]*nowait" "original" } }
+! { dg-final { scan-tree-dump-not "#pragma omp parallel\[^\n\r]*nowait" "original" } }
+! { dg-final { scan-tree-dump-not "#pragma omp section\[^s\]\[^\n\r]*nowait" "original" } }
+! { dg-final { scan-tree-dump-not "#pragma omp simd\[^\n\r]*nowait" "original" } }
+! { dg-final { scan-tree-dump-not "#pragma omp teams\[^\n\r]*nowait" "original" } }
+
+! Sometimes or never with nowait:
+
+! { dg-final { scan-tree-dump-times "#pragma omp distribute\[\n\r]" 8 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp loop\[\n\r]" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel\[\n\r]" 14 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp section\[\n\r]" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp simd linear\\(i:1\\)\[\n\r]" 8 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target\[\n\r]" 5 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams\[\n\r]" 8 "original" } }
@@ -16,12 +16,12 @@
!$omp end DO SIMD ! { dg-error "Unexpected !.OMP END DO SIMD" }
-!$omp end LOOP ! { dg-error "Unclassifiable OpenMP directive" }
+!$omp end LOOP ! { dg-error "Unexpected !.OMP END LOOP" }
!$omp parallel loop
do i = 1, 5
end do
-!$omp end LOOP ! { dg-error "Unclassifiable OpenMP directive" }
+!$omp end LOOP ! { dg-error "Unexpected !.OMP END LOOP" }
!$omp end MASKED ! { dg-error "Unexpected !.OMP END MASKED" }
@@ -44,7 +44,7 @@ end do
!$omp end PARALLEL DO SIMD ! { dg-error "Unexpected !.OMP END PARALLEL DO SIMD" }
!$omp loop
-!$omp end PARALLEL LOOP ! { dg-error "Unexpected junk" }
+!$omp end PARALLEL LOOP ! { dg-error "Unexpected !.OMP END PARALLEL LOOP" }
!$omp end PARALLEL MASKED ! { dg-error "Unexpected !.OMP END PARALLEL MASKED" }
@@ -80,7 +80,7 @@ end do
!$omp end TARGET PARALLEL DO SIMD ! { dg-error "Unexpected !.OMP END TARGET PARALLEL DO SIMD" }
-!$omp end TARGET PARALLEL LOOP ! { dg-error "Unexpected junk" }
+!$omp end TARGET PARALLEL LOOP ! { dg-error "Unexpected !.OMP END TARGET PARALLEL LOOP" }
!$omp end TARGET SIMD ! { dg-error "Unexpected !.OMP END TARGET SIMD" }
@@ -94,7 +94,7 @@ end do
!$omp end TARGET TEAMS DISTRIBUTE SIMD ! { dg-error "Unexpected !.OMP END TARGET TEAMS DISTRIBUTE SIMD" }
-!$omp end TARGET TEAMS LOOP ! { dg-error "Unexpected junk" }
+!$omp end TARGET TEAMS LOOP ! { dg-error "Unexpected !.OMP END TARGET TEAMS LOOP" }
!$omp end TASK ! { dg-error "Unexpected !.OMP END TASK" }
@@ -114,7 +114,7 @@ end do
!$omp end TEAMS DISTRIBUTE SIMD ! { dg-error "Unexpected !.OMP END TEAMS DISTRIBUTE SIMD" }
-!$omp end TEAMS LOOP ! { dg-error "Unexpected junk" }
+!$omp end TEAMS LOOP ! { dg-error "Unexpected !.OMP END TEAMS LOOP" }
!$omp end WORKSHARE ! { dg-error "Unexpected !.OMP END WORKSHARE" }