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.

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.

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 7b2df0d0be36630c3a6d2d803dfe6d5be3d9add0..2893ab2befb3301b16c8e3f4a02af09f3abeda42 100644 (file)
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -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;
 

diff --git a/gcc/fortran/parse.c b/gcc/fortran/parse.c
index 12aa80ec45ca91c7a2e4cc01edf244c765ac26e2..94b677f2a7087d5e049adb194a256564a5eb0e1c 100644 (file)
--- a/gcc/fortran/parse.c
+++ b/gcc/fortran/parse.c
@@ -924,6 +924,7 @@ decode_omp_directive (void)
       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);
@@ -939,6 +940,8 @@ decode_omp_directive (void)
       matchs ("end parallel do simd", gfc_match_omp_eos_error,
              ST_OMP_END_PARALLEL_DO_SIMD);
       matcho ("end parallel do", gfc_match_omp_eos_error, 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,
@@ -960,24 +963,29 @@ decode_omp_directive (void)
       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);
@@ -991,6 +999,7 @@ decode_omp_directive (void)
              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 +2562,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";

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 6bc7e9a6017ed43305a4e1880c176d85f4c9b338..b86c7cf98330ac2484095432230ba71696eb4c7a 100644 (file)
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -5878,6 +5878,8 @@ 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;
+         clausesa[GFC_OMP_SPLIT_TARGET].nowait
+           = code->ext.omp_clauses->nowait;
        }
       if (mask & GFC_OMP_MASK_TEAMS)
        {

diff --git a/gcc/testsuite/gfortran.dg/gomp/clauses-1.f90 b/gcc/testsuite/gfortran.dg/gomp/clauses-1.f90
new file mode 100644 (file)
index 0000000..639f5d1
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/clauses-1.f90
@@ -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 nowait
+
+  !$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 nowait
+
+  !$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 nowait
+
+  !$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 nowait
+
+  !$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 nowait
+
+  !$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 nowait
+
+  !$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 nowait
+
+end
+end module

diff --git a/gcc/testsuite/gfortran.dg/gomp/nowait-2.f90 b/gcc/testsuite/gfortran.dg/gomp/nowait-2.f90
new file mode 100644 (file)
index 0000000..d18459b
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/nowait-2.f90
@@ -0,0 +1,315 @@
+! Cross check that it is accepted without nowait
+subroutine bar()
+implicit none
+integer :: i, a(5)
+!$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 parallel do
+do i = 1, 5
+end do
+!$omp end distribute parallel do
+
+!$omp distribute parallel do simd
+do i = 1, 5
+end do
+!$omp end distribute parallel do simd
+
+!$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 workshare
+a(:) = 5
+!$omp end parallel workshare
+
+!$omp parallel do
+do i = 1, 5
+end do
+!$omp end parallel do
+
+!$omp parallel do simd
+do i = 1, 5
+end do
+!$omp end parallel do simd
+
+!$omp parallel sections
+  !$omp section
+  block; end block
+!$omp end parallel sections
+
+!$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 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 parallel do
+do i = 1, 5
+end do
+!$omp end teams distribute parallel do
+
+!$omp teams distribute parallel do simd
+do i = 1, 5
+end do
+!$omp end teams distribute parallel do simd
+
+!$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, a(5)
+!$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 parallel do
+do i = 1, 5
+end do
+!$omp end distribute parallel do nowait  ! { dg-error "Unexpected junk" }
+
+!$omp distribute parallel do simd
+do i = 1, 5
+end do
+!$omp end distribute parallel do simd nowait  ! { dg-error "Unexpected junk" }
+
+!$omp parallel sections
+  !$omp section
+  block; end block
+!$omp end parallel sections 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 workshare
+a(:) = 5
+!$omp end parallel workshare nowait  ! { dg-error "Unexpected junk" }
+
+!$omp parallel do
+do i = 1, 5
+end do
+!$omp end parallel do nowait  ! { dg-error "Unexpected junk" }
+
+!$omp parallel do simd
+do i = 1, 5
+end do
+!$omp end parallel do simd 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 parallel do
+do i = 1, 5
+end do
+!$omp end teams distribute parallel do nowait  ! { dg-error "Unexpected junk" }
+
+!$omp teams distribute parallel do simd
+do i = 1, 5
+end do
+!$omp end teams distribute parallel do simd nowait  ! { dg-error "Unexpected junk" }
+
+!$omp teams distribute simd
+do i = 1, 5
+end do
+!$omp end teams distribute simd nowait  ! { dg-error "Unexpected junk" }
+
+!$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" }

diff --git a/gcc/testsuite/gfortran.dg/gomp/nowait-3.f90 b/gcc/testsuite/gfortran.dg/gomp/nowait-3.f90
new file mode 100644 (file)
index 0000000..a58b8fd
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/nowait-3.f90
@@ -0,0 +1,118 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+subroutine foo
+implicit none
+integer :: i, a(5)
+
+!$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 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 nowait
+
+!$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 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 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 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 workshare
+A(:) = 5
+!$omp end workshare nowait
+end
+
+! Note: internally, for '... parallel do ...', 'nowait' is always added
+! such that for 'omp end target parallel do nowait', 'nowait' is on both
+! 'target' as specified in the OpenMP spec and and on 'do' due to internal usage.
+
+! Expected with 'nowait'
+
+! { dg-final { scan-tree-dump-times "#pragma omp for nowait" 6 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp for schedule\\(static\\) nowait" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp sections nowait" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp single nowait" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target nowait" 12 "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]" 4 "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]" 6 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp section\[\n\r]" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp simd linear\\(i:1\\)\[\n\r]" 5 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams\[\n\r]" 6 "original" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp target\[\n\r]" 0 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp for\[\n\r]" 0 "original" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/unexpected-end.f90 b/gcc/testsuite/gfortran.dg/gomp/unexpected-end.f90
index 96f10b594cf960bae45affb4b4705732924c8e56..70f54f9be5e584377c779af3660fab1bbeed7981 100644 (file)
--- a/gcc/testsuite/gfortran.dg/gomp/unexpected-end.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/unexpected-end.f90
@@ -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" }