[Patch] Fortran/openmp: Fix '!$omp end'

[Tobias Burnus <tobias@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 390 bytes --]

Found this when looking at the num_teams patch – and when
converting clauses-1.c to clauses-1.f90.

OK?

Tobias
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

[-- Attachment #2: omp-end.diff --]
[-- Type: text/x-patch, Size: 49614 bytes --]

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(-)

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 7b2df0d0be3..2893ab2befb 100644
--- 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 12aa80ec45c..d4b985e75eb 100644
--- a/gcc/fortran/parse.c
+++ b/gcc/fortran/parse.c
@@ -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";
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 6bc7e9a6017..928d205a5aa 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -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 (&section_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,
diff --git a/gcc/testsuite/gfortran.dg/gomp/clauses-1.f90 b/gcc/testsuite/gfortran.dg/gomp/clauses-1.f90
new file mode 100644
index 00000000000..d22274d7033
--- /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
+
+  !$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
diff --git a/gcc/testsuite/gfortran.dg/gomp/nowait-2.f90 b/gcc/testsuite/gfortran.dg/gomp/nowait-2.f90
new file mode 100644
index 00000000000..a1a3e86f6b0
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/nowait-2.f90
@@ -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" }
diff --git a/gcc/testsuite/gfortran.dg/gomp/nowait-3.f90 b/gcc/testsuite/gfortran.dg/gomp/nowait-3.f90
new file mode 100644
index 00000000000..94d95ba6dc9
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/nowait-3.f90
@@ -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" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/unexpected-end.f90 b/gcc/testsuite/gfortran.dg/gomp/unexpected-end.f90
index 96f10b594cf..70f54f9be5e 100644
--- 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" }