diff mbox series

OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065]

Message ID b915c835-71fb-ae37-487e-221a98ac22ec@codesourcery.com
State New
Headers show
Series OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065] | expand

Commit Message

Tobias Burnus July 24, 2023, 7:43 p.m. UTC 
This patch adds diagnostic for additional code alongside a nested teams
in a target region.

The diagnostic is happening soon after parsing such that expressions
in clauses are not yet expanded - those would end up before TEAMS
and can be very complicated (e.g. assume an allocatable-returning function).

(The patch diagnoses it in openmp.cc; after trans-openmp.cc it would
already be to late.)

Comments, remarks, suggestions?

Tobias

PS: Something similar is also needed for C/C++ but there templates
and lambda functions might make it harder to implement. In any case,
it has to be done in the FE. Tracked at PR71065
-----------------
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

Comments

Jakub Jelinek July 24, 2023, 7:49 p.m. UTC | #1 
On Mon, Jul 24, 2023 at 09:43:10PM +0200, Tobias Burnus wrote:
> This patch adds diagnostic for additional code alongside a nested teams
> in a target region.
> 
> The diagnostic is happening soon after parsing such that expressions
> in clauses are not yet expanded - those would end up before TEAMS
> and can be very complicated (e.g. assume an allocatable-returning function).
> 
> (The patch diagnoses it in openmp.cc; after trans-openmp.cc it would
> already be to late.)
> 
> Comments, remarks, suggestions?

Thanks for working on this.  The fuzzy thing on the Fortran side is
if e.g. multiple nested BLOCK statements can appear sandwiched in between
target and teams (of course without declarations in them), or if e.g.
extra empty BLOCK; END BLOCK could appear next to it etc.
And on C/C++ side similarly with {}s, ; is an empty statement, so
#pragma omp target
{
  ;
  #pragma omp teams
  ;
  ;
}
etc. would be invalid.

	Jakub
Tobias Burnus July 24, 2023, 8:05 p.m. UTC | #2 
On 24.07.23 21:49, Jakub Jelinek via Fortran wrote:
> Thanks for working on this.  The fuzzy thing on the Fortran side is
> if e.g. multiple nested BLOCK statements can appear sandwiched in between
> target and teams (of course without declarations in them), or if e.g.

The current patch rejects nested blocks, be it 'omp target; block;
block; omp teams;' or be it 'omp target; block; block;end block; omp teams'.

The current wording in the spec is also rather explicit as 'block' is a
statement.

(BTW: For 'block; block; omp teams', the simplistic search won't work
such that for those only the location of TARGET and not of TEAMS is
shown. I could try harder but as it is useful as is and such code should
be rare, I don't do it.)

Thus, I believe the patch should be fine.

Tobias

PS: I know that some regard {{{ }}} and block; block, ... end block; ...
as something to be ignored. Thus, for 'omp atomic', TR12 will allow any
number of curly braces and BLOCK/ENDBLOCK pairs. The wording there is
rather explicit but also localized, i.e. it won't affect other code
locations (for now at least).

-----------------
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
Tobias Burnus July 25, 2023, 7:37 a.m. UTC | #3 
Now committed as r14-2754-g2e31fe431b08b0 with a minor addition:

On 24.07.23 22:05, Tobias Burnus wrote:
> The current patch rejects nested blocks, be it 'omp target; block;
> block; omp teams;'
which was before in the testcase. But now also
> or be it 'omp target; block; block;end block; omp teams'.
is tested for.

Somehow, the second dg-error line in an modified testcase did not make
it in the first commit; now fixed in r14-2759-g50656980497d77

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
diff mbox series

Patch

OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065]

OpenMP requires: "If a teams region is nested inside a target region, the
corresponding target construct must not contain any statements, declarations
or directives outside of the corresponding teams construct."

Test for it!

	PR fortran/110725
	PR middle-end/71065

gcc/fortran/ChangeLog:

	* gfortran.h (gfc_omp_clauses): Add contains_teams_construct.
	* openmp.cc (resolve_omp_target): New; check for teams nesting.
	(gfc_resolve_omp_directive): Call it.
	* parse.cc (decode_omp_directive): Set contains_teams_construct
	on enclosing ST_OMP_TARGET.

gcc/testsuite/ChangeLog:

	* gfortran.dg/gomp/pr99226.f90: Update dg-error.
	* gfortran.dg/gomp/teams-5.f90: New test.

 gcc/fortran/gfortran.h                     |   1 +
 gcc/fortran/openmp.cc                      |  39 ++++++++-
 gcc/fortran/parse.cc                       |  33 ++++++++
 gcc/testsuite/gfortran.dg/gomp/pr99226.f90 |   2 +-
 gcc/testsuite/gfortran.dg/gomp/teams-5.f90 | 127 +++++++++++++++++++++++++++++
 5 files changed, 200 insertions(+), 2 deletions(-)

diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 6482a885211..577ef807af7 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1575,6 +1575,7 @@  typedef struct gfc_omp_clauses
   unsigned order_unconstrained:1, order_reproducible:1, capture:1;
   unsigned grainsize_strict:1, num_tasks_strict:1, compare:1, weak:1;
   unsigned non_rectangular:1, order_concurrent:1;
+  unsigned contains_teams_construct:1;
   ENUM_BITFIELD (gfc_omp_sched_kind) sched_kind:3;
   ENUM_BITFIELD (gfc_omp_device_type) device_type:2;
   ENUM_BITFIELD (gfc_omp_memorder) memorder:3;
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 05a697da071..675011a18ce 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -10653,6 +10653,41 @@  gfc_resolve_oacc_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED)
 }
 
 
+static void
+resolve_omp_target (gfc_code *code)
+{
+#define GFC_IS_TEAMS_CONSTRUCT(op)			\
+  (op == EXEC_OMP_TEAMS					\
+   || op == EXEC_OMP_TEAMS_DISTRIBUTE			\
+   || op == EXEC_OMP_TEAMS_DISTRIBUTE_SIMD		\
+   || op == EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO	\
+   || op == EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD	\
+   || op == EXEC_OMP_TEAMS_LOOP)
+
+  if (!code->ext.omp_clauses->contains_teams_construct)
+    return;
+  if ((GFC_IS_TEAMS_CONSTRUCT (code->block->next->op)
+       && code->block->next->next == NULL)
+      || (code->block->next->op == EXEC_BLOCK
+	  && code->block->next->next
+	  && GFC_IS_TEAMS_CONSTRUCT (code->block->next->next->op)
+	  && code->block->next->next->next == NULL))
+    return;
+  gfc_code *c = code->block->next;
+  while (c && !GFC_IS_TEAMS_CONSTRUCT (c->op))
+    c = c->next;
+  if (c)
+    gfc_error ("!$OMP TARGET region at %L with a nested TEAMS at %L may not "
+	       "contain any other statement, declaration or directive outside "
+	       "of the single TEAMS construct", &c->loc, &code->loc);
+  else
+    gfc_error ("!$OMP TARGET region at %L with a nested TEAMS may not "
+	       "contain any other statement, declaration or directive outside "
+	       "of the single TEAMS construct", &code->loc);
+#undef GFC_IS_TEAMS_CONSTRUCT
+}
+
+
 /* Resolve OpenMP directive clauses and check various requirements
    of each directive.  */
 
@@ -10703,6 +10738,9 @@  gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns)
     case EXEC_OMP_TEAMS_LOOP:
       resolve_omp_do (code);
       break;
+    case EXEC_OMP_TARGET:
+      resolve_omp_target (code);
+      gcc_fallthrough ();
     case EXEC_OMP_ALLOCATE:
     case EXEC_OMP_ALLOCATORS:
     case EXEC_OMP_ASSUME:
@@ -10718,7 +10756,6 @@  gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns)
     case EXEC_OMP_SCOPE:
     case EXEC_OMP_SECTIONS:
     case EXEC_OMP_SINGLE:
-    case EXEC_OMP_TARGET:
     case EXEC_OMP_TARGET_DATA:
     case EXEC_OMP_TARGET_ENTER_DATA:
     case EXEC_OMP_TARGET_EXIT_DATA:
diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc
index e53b7a42e92..011a39c3d04 100644
--- a/gcc/fortran/parse.cc
+++ b/gcc/fortran/parse.cc
@@ -1312,6 +1312,39 @@  decode_omp_directive (void)
 	  prog_unit->omp_target_seen = true;
 	break;
       }
+    case ST_OMP_TEAMS:
+    case ST_OMP_TEAMS_DISTRIBUTE:
+    case ST_OMP_TEAMS_DISTRIBUTE_SIMD:
+    case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+    case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case ST_OMP_TEAMS_LOOP:
+      if (gfc_state_stack->previous && gfc_state_stack->previous->tail)
+	{
+	  gfc_state_data *stk = gfc_state_stack;
+	  do {
+	       stk = stk->previous;
+	     } while (stk && stk->tail && stk->tail->op == EXEC_BLOCK);
+	  if (stk && stk->tail)
+	    switch (stk->tail->op)
+	      {
+	      case EXEC_OMP_TARGET:
+	      case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+	      case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+	      case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	      case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	      case EXEC_OMP_TARGET_TEAMS_LOOP:
+	      case EXEC_OMP_TARGET_PARALLEL:
+	      case EXEC_OMP_TARGET_PARALLEL_DO:
+	      case EXEC_OMP_TARGET_PARALLEL_DO_SIMD:
+	      case EXEC_OMP_TARGET_PARALLEL_LOOP:
+	      case EXEC_OMP_TARGET_SIMD:
+		stk->tail->ext.omp_clauses->contains_teams_construct = 1;
+		break;
+	  default:
+	    break;
+	  }
+	}
+      break;
     case ST_OMP_ERROR:
       if (new_st.ext.omp_clauses->at != OMP_AT_EXECUTION)
 	return ST_NONE;
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr99226.f90 b/gcc/testsuite/gfortran.dg/gomp/pr99226.f90
index 72dbdde2e28..2aea0c15585 100644
--- a/gcc/testsuite/gfortran.dg/gomp/pr99226.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/pr99226.f90
@@ -2,7 +2,7 @@ 
 
 subroutine sub (n)
    integer :: n, i
-   !$omp target	! { dg-error "construct with nested 'teams' construct contains directives outside of the 'teams' construct" }
+   !$omp target	! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
    !$omp teams distribute dist_schedule (static,n+4)
    do i = 1, 8
    end do
diff --git a/gcc/testsuite/gfortran.dg/gomp/teams-5.f90 b/gcc/testsuite/gfortran.dg/gomp/teams-5.f90
new file mode 100644
index 00000000000..bf5461b87c8
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/teams-5.f90
@@ -0,0 +1,127 @@ 
+! { dg-do compile }
+
+! PR fortran/110725
+! PR middle-end/71065
+
+implicit none
+integer :: x
+!$omp target device(1)
+  block
+    !$omp teams num_teams(f())
+    !$omp end teams
+  end block
+!!$omp end target
+
+!$omp target device(1)
+  !$omp teams num_teams(f())
+  !$omp end teams
+!$omp end target
+
+!$omp target device(1)  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+  x = 5
+  !$omp teams num_teams(f())  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+  !$omp end teams
+!$omp end target
+
+!$omp target device(1)  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+  !$omp teams num_teams(f())  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+  !$omp end teams
+  x = 5
+!$omp end target
+
+!$omp target  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+  block
+    block
+    !$omp teams num_teams(f())
+    !$omp end teams
+    end block
+  end block
+!$omp end target
+
+!$omp target device(1)  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+  block
+    x = 5
+    !$omp teams num_teams(f())  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+    !$omp end teams
+  end block
+!$omp end target
+
+!$omp target device(1)  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+  block
+    !$omp teams num_teams(f())  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+    !$omp end teams
+    x = 5
+  end block
+!$omp end target
+
+!$omp target device(1)  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+  block
+    !$omp teams num_teams(f())  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+    !$omp end teams
+    x = 5
+  end block
+!$omp end target
+
+contains
+
+function f()
+  !$omp declare target
+  integer, allocatable :: f
+  f = 5
+end
+end
+
+subroutine sub1
+  implicit none
+  integer :: x,i
+
+  !$omp target device(1)  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+  block
+    !$omp teams distribute num_teams(f())  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+    do i = 1, 5
+    end do
+    x = 7
+  end block
+  !$omp end target
+
+  !$omp target device(1)  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+  block
+    !$omp teams loop num_teams(f())  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+    do i = 1, 5
+    end do
+    x = 7
+  end block
+  !$omp end target
+
+  !$omp target device(1)  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+    !$omp teams distribute simd num_teams(f())  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+    do i = 1, 5
+    end do
+    x = 7
+  !$omp end target
+
+  !$omp target device(1)  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+    !$omp teams distribute parallel do num_teams(f())  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+    do i = 1, 5
+    end do
+    x = 7
+  !$omp end target
+
+  !$omp target device(1)  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+  block
+    x = 7
+    !$omp teams distribute parallel do simd num_teams(f())  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+    do i = 1, 5
+    end do
+  end block
+  !$omp end target
+
+contains
+
+function f()
+  !$omp declare target
+  integer, allocatable :: f
+  f = 5
+end
+
+end