[PATCH] OpenMP: Duplicate checking for map clauses in Fortran (PR107214)

[PATCH] OpenMP: Duplicate checking for map clauses in Fortran (PR107214)

[Julian Brown]

This patch adds duplicate checking for OpenMP "map" clauses, taking some
cues from the implementation for C in c-typeck.cc:c_finish_omp_clauses
(and similar for C++).

In addition to the existing use of the "mark" and "comp_mark" bitfields
in the gfc_symbol structure, the patch adds several new bits handling
duplicate checking within various categories of clause types.  If "mark"
is being used for map clauses, we need to use different bits for other
clauses for cases where "map" and some other clause can refer to the
same symbol (e.g. "map(n) shared(n)").

Tested with offloading to NVPTX. OK?

2022-10-20  Julian Brown  <julian@codesourcery.com>

gcc/fortran/
	PR fortran/107214
	* gfortran.h (gfc_symbol): Add data_mark, dev_mark, gen_mark and
	reduc_mark bitfields.
	* openmp.cc (resolve_omp_clauses): Use above bitfields to improve
	duplicate clause detection.

gcc/testsuite/
	PR fortran/107214
	* gfortran.dg/gomp/pr107214.f90: New test.
---
 gcc/fortran/gfortran.h                      | 16 +++++-
 gcc/fortran/openmp.cc                       | 63 +++++++++++++++++----
 gcc/testsuite/gfortran.dg/gomp/pr107214.f90 |  7 +++
 3 files changed, 72 insertions(+), 14 deletions(-)
 create mode 100644 gcc/testsuite/gfortran.dg/gomp/pr107214.f90

diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index fe8c4e131f3..511a1ec3623 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1861,9 +1861,21 @@ typedef struct gfc_symbol
      the current statement.  Otherwise, old_symbol points to a copy of
      the old symbol. gfc_new is used in symbol.cc to flag new symbols.
      comp_mark is used to indicate variables which have component accesses
-     in OpenMP/OpenACC directive clauses.  */
+     in OpenMP/OpenACC directive clauses (cf. c-typeck.cc:c_finish_omp_clauses,
+     map_field_head).
+     data_mark is used to check duplicate mappings for OpenMP data-sharing
+     clauses (see firstprivate_head/lastprivate_head in the above function).
+     dev_mark is used to check duplicate mappings for OpenMP
+     is_device_ptr/has_device_addr clauses (see is_on_device_head in above
+     function).
+     gen_mark is used to check duplicate mappings for OpenMP
+     use_device_ptr/use_device_addr/private/shared clauses (see generic_head in
+     above functon).
+     reduc_mark is used to check duplicate mappings for OpenMP reduction
+     clauses.  */
   struct gfc_symbol *old_symbol;
-  unsigned mark:1, comp_mark:1, gfc_new:1;
+  unsigned mark:1, comp_mark:1, data_mark:1, dev_mark:1, gen_mark:1;
+  unsigned reduc_mark:1, gfc_new:1;
 
   /* The tlink field is used in the front end to carry the module
      declaration of separate module procedures so that the characteristics
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index ce719bd5d92..d4595aae23e 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -6738,6 +6738,10 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	  continue;
 	n->sym->mark = 0;
 	n->sym->comp_mark = 0;
+	n->sym->data_mark = 0;
+	n->sym->dev_mark = 0;
+	n->sym->gen_mark = 0;
+	n->sym->reduc_mark = 0;
 	if (n->sym->attr.flavor == FL_VARIABLE
 	    || n->sym->attr.proc_pointer
 	    || (!code && (!n->sym->attr.dummy || n->sym->ns != ns)))
@@ -6806,7 +6810,6 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	&& list != OMP_LIST_LASTPRIVATE
 	&& list != OMP_LIST_ALIGNED
 	&& list != OMP_LIST_DEPEND
-	&& (list != OMP_LIST_MAP || openacc)
 	&& list != OMP_LIST_FROM
 	&& list != OMP_LIST_TO
 	&& (list != OMP_LIST_REDUCTION || !openacc)
@@ -6825,10 +6828,43 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	    for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next)
 	      if (ref->type == REF_COMPONENT)
 		component_ref_p = true;
-	  if ((!component_ref_p && n->sym->comp_mark)
-	      || (component_ref_p && n->sym->mark))
-	    gfc_error ("Symbol %qs has mixed component and non-component "
-		       "accesses at %L", n->sym->name, &n->where);
+	  if ((list == OMP_LIST_IS_DEVICE_PTR
+	       || list == OMP_LIST_HAS_DEVICE_ADDR)
+	      && !component_ref_p)
+	    {
+	      if (n->sym->gen_mark || n->sym->dev_mark || n->sym->reduc_mark)
+		gfc_error ("Symbol %qs present on multiple clauses at %L",
+			   n->sym->name, &n->where);
+	      else
+		n->sym->dev_mark = 1;
+	    }
+	  else if ((list == OMP_LIST_USE_DEVICE_PTR
+		    || list == OMP_LIST_USE_DEVICE_ADDR
+		    || list == OMP_LIST_PRIVATE
+		    || list == OMP_LIST_SHARED)
+		   && !component_ref_p)
+	    {
+	      if (n->sym->gen_mark || n->sym->dev_mark || n->sym->reduc_mark)
+		gfc_error ("Symbol %qs present on multiple clauses at %L",
+			   n->sym->name, &n->where);
+	      else
+		n->sym->gen_mark = 1;
+	    }
+	  else if (list == OMP_LIST_REDUCTION && !component_ref_p)
+	    {
+	      if (n->sym->gen_mark || n->sym->dev_mark || n->sym->reduc_mark)
+		gfc_error ("Symbol %qs present on multiple clauses at %L",
+			   n->sym->name, &n->where);
+	      else
+		n->sym->reduc_mark = 1;
+	    }
+	  else if ((!component_ref_p && n->sym->comp_mark)
+		   || (component_ref_p && n->sym->mark))
+	    {
+	      if (openacc)
+		gfc_error ("Symbol %qs has mixed component and non-component "
+			   "accesses at %L", n->sym->name, &n->where);
+	    }
 	  else if (n->sym->mark)
 	    gfc_error ("Symbol %qs present on multiple clauses at %L",
 		       n->sym->name, &n->where);
@@ -6844,31 +6880,34 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
   gcc_assert (OMP_LIST_LASTPRIVATE == OMP_LIST_FIRSTPRIVATE + 1);
   for (list = OMP_LIST_FIRSTPRIVATE; list <= OMP_LIST_LASTPRIVATE; list++)
     for (n = omp_clauses->lists[list]; n; n = n->next)
-      if (n->sym->mark)
+      if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark)
 	{
 	  gfc_error ("Symbol %qs present on multiple clauses at %L",
 		     n->sym->name, &n->where);
-	  n->sym->mark = 0;
+	  n->sym->data_mark = 0;
 	}
+      else if (n->sym->mark)
+	gfc_error ("Symbol %qs present on both data and map clauses "
+		   "at %L", n->sym->name, &n->where);
 
   for (n = omp_clauses->lists[OMP_LIST_FIRSTPRIVATE]; n; n = n->next)
     {
-      if (n->sym->mark)
+      if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark)
 	gfc_error ("Symbol %qs present on multiple clauses at %L",
 		   n->sym->name, &n->where);
       else
-	n->sym->mark = 1;
+	n->sym->data_mark = 1;
     }
   for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next)
-    n->sym->mark = 0;
+    n->sym->data_mark = 0;
 
   for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next)
     {
-      if (n->sym->mark)
+      if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark)
 	gfc_error ("Symbol %qs present on multiple clauses at %L",
 		   n->sym->name, &n->where);
       else
-	n->sym->mark = 1;
+	n->sym->data_mark = 1;
     }
 
   for (n = omp_clauses->lists[OMP_LIST_ALIGNED]; n; n = n->next)
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214.f90
new file mode 100644
index 00000000000..25949934e84
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr107214.f90
@@ -0,0 +1,7 @@
+! { dg-do compile }
+
+program p
+   integer, allocatable :: a
+   !$omp target map(tofrom: a, a) ! { dg-error "Symbol 'a' present on multiple clauses" }
+   !$omp end target
+end
-- 
2.29.2

Re: [PATCH] OpenMP: Duplicate checking for map clauses in Fortran (PR107214)

[Tobias Burnus]

Hi Julian,

I had a first quick lock at this patch, I should have a closer look
later. However, I stumbled over the following:

On 20.10.22 18:14, Julian Brown wrote:
> typedef struct gfc_symbol
> {
> ...
>    struct gfc_symbol *old_symbol;
>
>    unsigned mark:1, comp_mark:1, data_mark:1, dev_mark:1, gen_mark:1;
>    unsigned reduc_mark:1, gfc_new:1;
>
>    struct gfc_symbol *tlink;
>
>    unsigned equiv_built:1;
>    ...
I know that this was the case before, but can you move the mark:1 etc.
after 'tlink'? In that case all bitfields are grouped together. If I
have not miscounted, we have currently 7 bits before and 9 bits after
'tlink' and grouping them together reduced pointless padding.

* * *
> +      else if (n->sym->mark)
> +     gfc_error ("Symbol %qs present on both data and map clauses "
> +                "at %L", n->sym->name, &n->where);

I wonder whether that also rejects the following – which seems to be
valid. The 'map' goes to 'target' and the 'firstprivate' to 'parallel',
cf. OpenMP 5.2, "17.2 Clauses on Combined and Composite Constructs",
[340:3-4 & 12-14]. (BTW: While some fixes went into 5.1 regarding this section,
a likewise wording is already in 5.0.)

(Testing showed: it give an ICE without the patch and an error with.)

module m
   integer :: a = 1
end module m

module m2
contains
subroutine bar()
   use m
   !$omp declare target
   a = a + 5
end subroutine bar
end

program p
   use m
   !$omp target parallel do map(a) firstprivate(a)
     do i = 1, 1
        a = 7
       call bar()
        if (a /= 7) error stop 1
        a = a + 8

    end do
   if (a /= 6) error stop
end

  * * *

The ICE seems to be because gcc/fortran/trans-openmp.cc's gfc_split_omp_clauses
mishandles this as the dump shows the following:

   #pragma omp target firstprivate(a) map(tofrom:a)
         #pragma omp parallel firstprivate(a)

  * * *

In contrast, for the C testcase:

void foo(int x) {
#pragma omp target parallel for simd map(x) firstprivate(x)
for (int k = 0; k < 1; ++k)
   x = 1;
}

the dump is as follows, which seems to be sensible:

   #pragma omp target map(tofrom:x)
         #pragma omp parallel firstprivate(x)
               #pragma omp for nowait
                     #pragma omp simd

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

[PATCH 1/2] OpenMP/Fortran: Combined directives with map/firstprivate of same symbol

[Julian Brown]

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

On Wed, 26 Oct 2022 12:39:39 +0200
Tobias Burnus <tobias@codesourcery.com> wrote:

> The ICE seems to be because gcc/fortran/trans-openmp.cc's
> gfc_split_omp_clauses mishandles this as the dump shows the following:
> 
>    #pragma omp target firstprivate(a) map(tofrom:a)
>          #pragma omp parallel firstprivate(a)
> 
>   * * *
> 
> In contrast, for the C testcase:
> 
> void foo(int x) {
> #pragma omp target parallel for simd map(x) firstprivate(x)
> for (int k = 0; k < 1; ++k)
>    x = 1;
> }
> 
> the dump is as follows, which seems to be sensible:
> 
>    #pragma omp target map(tofrom:x)
>          #pragma omp parallel firstprivate(x)
>                #pragma omp for nowait
>                      #pragma omp simd

First, here's a patch to address this bit...

This patch fixes a case where a combined directive (e.g. "!$omp target
parallel ...") contains both a map and a firstprivate clause for the
same variable.  When the combined directive is split into two nested
directives, the outer "target" gets the "map" clause, and the inner
"parallel" gets the "firstprivate" clause, like so:
  
  !$omp target parallel map(x) firstprivate(x)
  
  -->
  
  !$omp target map(x)
    !$omp parallel firstprivate(x)
      ...

When there is no map of the same variable, the firstprivate is distributed
to both directives, e.g. for 'y' in:
  
  !$omp target parallel map(x) firstprivate(y)
  
  -->
  
  !$omp target map(x) firstprivate(y)
    !$omp parallel firstprivate(y)
      ...

This is not a recent regression, but appears to fix a long-standing ICE.
(The included testcase is based on one by Tobias.)

Tested with offloading to NVPTX, alongside previously-posted patches
(in review or approved but waiting for other patches), i.e.:
  
  OpenMP/OpenACC: Rework clause expansion and nested struct handling
  OpenMP/OpenACC: Refine condition for when map clause expansion happens
  OpenMP: Pointers and member mappings

and the patch following.  OK?

2022-12-06  Julian Brown  <julian@codesourcery.com>

gcc/fortran/
        * trans-openmp.cc (gfc_add_firstprivate_if_unmapped): New function.
        (gfc_split_omp_clauses): Call above.

libgomp/
        * testsuite/libgomp.fortran/combined-directive-splitting-1.f90: New
        test.

[-- Attachment #2: 0001-OpenMP-Fortran-Combined-directives-with-map-firstpri.patch --]
[-- Type: text/x-patch, Size: 4936 bytes --]

From c66db363066913ae4939f2aa706427338b109d71 Mon Sep 17 00:00:00 2001
Message-Id: <c66db363066913ae4939f2aa706427338b109d71.1670438768.git.julian@codesourcery.com>
From: Julian Brown <julian@codesourcery.com>
Date: Tue, 6 Dec 2022 12:18:33 +0000
Subject: [PATCH 1/2] OpenMP/Fortran: Combined directives with map/firstprivate
 of same symbol

This patch fixes a case where a combined directive (e.g. "!$omp target
parallel ...") contains both a map and a firstprivate clause for the
same variable.  When the combined directive is split into two nested
directives, the outer "target" gets the "map" clause, and the inner
"parallel" gets the "firstprivate" clause, like so:

  !$omp target parallel map(x) firstprivate(x)

  -->

  !$omp target map(x)
    !$omp parallel firstprivate(x)
      ...

When there is no map of the same variable, the firstprivate is distributed
to both directives, e.g. for 'y' in:

  !$omp target parallel map(x) firstprivate(y)

  -->

  !$omp target map(x) firstprivate(y)
    !$omp parallel firstprivate(y)
      ...

This is not a recent regression, but appears to fix a long-standing ICE.
(The included testcase is based on one by Tobias.)

Tested with offloading to NVPTX, alongside previously-posted patches
(in review or approved but waiting for other patches), i.e.:

  OpenMP/OpenACC: Rework clause expansion and nested struct handling
  OpenMP/OpenACC: Refine condition for when map clause expansion happens
  OpenMP: Pointers and member mappings

and the patch following.  OK?

2022-12-06  Julian Brown  <julian@codesourcery.com>

gcc/fortran/
	* trans-openmp.cc (gfc_add_firstprivate_if_unmapped): New function.
	(gfc_split_omp_clauses): Call above.

libgomp/
	* testsuite/libgomp.fortran/combined-directive-splitting-1.f90: New
	test.
---
 gcc/fortran/trans-openmp.cc                   | 37 ++++++++++++++++-
 .../combined-directive-splitting-1.f90        | 41 +++++++++++++++++++
 2 files changed, 76 insertions(+), 2 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.fortran/combined-directive-splitting-1.f90

diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index e39f7b1cb273..c61cd1bf55de 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -6121,6 +6121,39 @@ gfc_add_clause_implicitly (gfc_omp_clauses *clauses_out,
     }
 }
 
+/* Kind of opposite to above, add firstprivate to CLAUSES_OUT if it is mapped
+   in CLAUSES_IN's FIRSTPRIVATE list but not its MAP list.  */
+
+static void
+gfc_add_firstprivate_if_unmapped (gfc_omp_clauses *clauses_out,
+				  gfc_omp_clauses *clauses_in)
+{
+  gfc_omp_namelist *n = clauses_in->lists[OMP_LIST_FIRSTPRIVATE];
+  gfc_omp_namelist **tail = NULL;
+
+  for (; n != NULL; n = n->next)
+    {
+      gfc_omp_namelist *n2 = clauses_out->lists[OMP_LIST_MAP];
+      for (; n2 != NULL; n2 = n2->next)
+	if (n->sym == n2->sym)
+	  break;
+      if (n2 == NULL)
+	{
+	  gfc_omp_namelist *dup = gfc_get_omp_namelist ();
+	  *dup = *n;
+	  dup->next = NULL;
+	  if (!tail)
+	    {
+	      tail = &clauses_out->lists[OMP_LIST_FIRSTPRIVATE];
+	      while (*tail && (*tail)->next)
+		tail = &(*tail)->next;
+	    }
+	  *tail = dup;
+	  tail = &(*tail)->next;
+	}
+    }
+}
+
 static void
 gfc_free_split_omp_clauses (gfc_code *code, gfc_omp_clauses *clausesa)
 {
@@ -6504,8 +6537,8 @@ gfc_split_omp_clauses (gfc_code *code,
 	 simd and masked/master.  Put it on the outermost of those and duplicate
 	 on parallel and teams.  */
       if (mask & GFC_OMP_MASK_TARGET)
-	clausesa[GFC_OMP_SPLIT_TARGET].lists[OMP_LIST_FIRSTPRIVATE]
-	  = code->ext.omp_clauses->lists[OMP_LIST_FIRSTPRIVATE];
+	gfc_add_firstprivate_if_unmapped (&clausesa[GFC_OMP_SPLIT_TARGET],
+					  code->ext.omp_clauses);
       if (mask & GFC_OMP_MASK_TEAMS)
 	clausesa[GFC_OMP_SPLIT_TEAMS].lists[OMP_LIST_FIRSTPRIVATE]
 	  = code->ext.omp_clauses->lists[OMP_LIST_FIRSTPRIVATE];
diff --git a/libgomp/testsuite/libgomp.fortran/combined-directive-splitting-1.f90 b/libgomp/testsuite/libgomp.fortran/combined-directive-splitting-1.f90
new file mode 100644
index 000000000000..e662a2bd3b20
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/combined-directive-splitting-1.f90
@@ -0,0 +1,41 @@
+module m
+  integer :: a = 1
+  !$omp declare target enter(a)
+end module m
+
+module m2
+contains
+subroutine bar()
+  use m
+  implicit none
+  !$omp declare target
+  a = a + 5
+end subroutine bar
+end module m2
+
+program p
+  use m
+  use m2
+  implicit none
+  integer :: b, i
+
+  !$omp target parallel do map(always, tofrom: a) firstprivate(a)
+    do i = 1, 1
+      a = 7
+      call bar()
+      if (a /= 7) error stop 1
+      a = a + 8
+    end do
+  if (a /= 6) error stop 2
+
+  b = 3
+  !$omp target parallel do map(always, tofrom: a) firstprivate(b)
+    do i = 1, 1
+      a = 3
+      call bar ()
+      if (a /= 8) error stop 3
+      a = a + b
+    end do
+  if (a /= 11) error stop 4
+end program p
+
-- 
2.29.2

[PATCH 2/2] OpenMP: Duplicate checking for map clauses in Fortran (PR107214)

[Julian Brown]

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

> Hi Julian,
> 
> I had a first quick lock at this patch, I should have a closer look
> later. However, I stumbled over the following:
> 
> On 20.10.22 18:14, Julian Brown wrote:
> > typedef struct gfc_symbol
> > {
> > ...
> >    struct gfc_symbol *old_symbol;
> >
> >    unsigned mark:1, comp_mark:1, data_mark:1, dev_mark:1,
> > gen_mark:1; unsigned reduc_mark:1, gfc_new:1;
> >
> >    struct gfc_symbol *tlink;
> >
> >    unsigned equiv_built:1;
> >    ...  
> I know that this was the case before, but can you move the mark:1 etc.
> after 'tlink'? In that case all bitfields are grouped together. If I
> have not miscounted, we have currently 7 bits before and 9 bits after
> 'tlink' and grouping them together reduced pointless padding.
> 
> * * *
> > +      else if (n->sym->mark)
> > +     gfc_error ("Symbol %qs present on both data and map clauses "
> > +                "at %L", n->sym->name, &n->where);  
> 
> I wonder whether that also rejects the following – which seems to be
> valid. The 'map' goes to 'target' and the 'firstprivate' to
> 'parallel', cf. OpenMP 5.2, "17.2 Clauses on Combined and Composite
> Constructs", [340:3-4 & 12-14]. (BTW: While some fixes went into 5.1
> regarding this section, a likewise wording is already in 5.0.)
> 
> (Testing showed: it give an ICE without the patch and an error with.)

...and this patch avoids the error for combined directives, and
reorders the gfc_symbol bitfields.

--

This patch adds duplicate checking for OpenMP "map" clauses, taking some
cues from the implementation for C in c-typeck.cc:c_finish_omp_clauses
(and similar for C++).

In addition to the existing use of the "mark" and "comp_mark" bitfields
in the gfc_symbol structure, the patch adds several new bits handling
duplicate checking within various categories of clause types.  If "mark"
is being used for map clauses, we need to use different bits for other
clauses for cases where "map" and some other clause can refer to the
same symbol (e.g. "map(n) shared(n)").

This version of the patch avoids flagging variables that are listed on
both map and firstprivate clauses when they are on a combined directive,
as they get moved to separate nested directives later (see previous
patch in series).

Tested with offloading to NVPTX alongside previous patch (and
dependencies).  OK?

2022-12-06  Julian Brown  <julian@codesourcery.com>

gcc/fortran/
        PR fortran/107214
        * gfortran.h (gfc_symbol): Add data_mark, dev_mark, gen_mark and
        reduc_mark bitfields.
        * openmp.cc (resolve_omp_clauses): Use above bitfields to improve
        duplicate clause detection.

gcc/testsuite/
        PR fortran/107214
        * gfortran.dg/gomp/pr107214.f90: New test.

[-- Attachment #2: 0002-OpenMP-Duplicate-checking-for-map-clauses-in-Fortran.patch --]
[-- Type: text/x-patch, Size: 9807 bytes --]

From fa6d1e273449aff61833064027fed3787c13121f Mon Sep 17 00:00:00 2001
Message-Id: <fa6d1e273449aff61833064027fed3787c13121f.1670438768.git.julian@codesourcery.com>
In-Reply-To: <c66db363066913ae4939f2aa706427338b109d71.1670438768.git.julian@codesourcery.com>
References: <c66db363066913ae4939f2aa706427338b109d71.1670438768.git.julian@codesourcery.com>
From: Julian Brown <julian@codesourcery.com>
Date: Tue, 6 Dec 2022 23:10:58 +0000
Subject: [PATCH 2/2] OpenMP: Duplicate checking for map clauses in Fortran
 (PR107214)

This patch adds duplicate checking for OpenMP "map" clauses, taking some
cues from the implementation for C in c-typeck.cc:c_finish_omp_clauses
(and similar for C++).

In addition to the existing use of the "mark" and "comp_mark" bitfields
in the gfc_symbol structure, the patch adds several new bits handling
duplicate checking within various categories of clause types.  If "mark"
is being used for map clauses, we need to use different bits for other
clauses for cases where "map" and some other clause can refer to the
same symbol (e.g. "map(n) shared(n)").

This version of the patch avoids flagging variables that are listed on
both map and firstprivate clauses when they are on a combined directive,
as they get moved to separate nested directives later (see previous
patch in series).

Tested with offloading to NVPTX alongside previous patch (and
dependencies).  OK?

2022-12-06  Julian Brown  <julian@codesourcery.com>

gcc/fortran/
	PR fortran/107214
	* gfortran.h (gfc_symbol): Add data_mark, dev_mark, gen_mark and
	reduc_mark bitfields.
	* openmp.cc (resolve_omp_clauses): Use above bitfields to improve
	duplicate clause detection.

gcc/testsuite/
	PR fortran/107214
	* gfortran.dg/gomp/pr107214.f90: New test.
---
 gcc/fortran/gfortran.h                      | 32 ++++++---
 gcc/fortran/openmp.cc                       | 73 +++++++++++++++++----
 gcc/testsuite/gfortran.dg/gomp/pr107214.f90 |  7 ++
 3 files changed, 90 insertions(+), 22 deletions(-)
 create mode 100644 gcc/testsuite/gfortran.dg/gomp/pr107214.f90

diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index df90ed39bea7..47a7f5552385 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1871,16 +1871,6 @@ typedef struct gfc_symbol
 
   gfc_namelist *namelist, *namelist_tail;
 
-  /* Change management fields.  Symbols that might be modified by the
-     current statement have the mark member nonzero.  Of these symbols,
-     symbols with old_symbol equal to NULL are symbols created within
-     the current statement.  Otherwise, old_symbol points to a copy of
-     the old symbol. gfc_new is used in symbol.cc to flag new symbols.
-     comp_mark is used to indicate variables which have component accesses
-     in OpenMP/OpenACC directive clauses.  */
-  struct gfc_symbol *old_symbol;
-  unsigned mark:1, comp_mark:1, gfc_new:1;
-
   /* The tlink field is used in the front end to carry the module
      declaration of separate module procedures so that the characteristics
      can be compared with the corresponding declaration in a submodule. In
@@ -1888,6 +1878,28 @@ typedef struct gfc_symbol
      deferred initialization.  */
   struct gfc_symbol *tlink;
 
+  /* Change management fields.  Symbols that might be modified by the
+     current statement have the mark member nonzero.  Of these symbols,
+     symbols with old_symbol equal to NULL are symbols created within
+     the current statement.  Otherwise, old_symbol points to a copy of
+     the old symbol. gfc_new is used in symbol.cc to flag new symbols.
+     comp_mark is used to indicate variables which have component accesses
+     in OpenMP/OpenACC directive clauses (cf. c-typeck.cc:c_finish_omp_clauses,
+     map_field_head).
+     data_mark is used to check duplicate mappings for OpenMP data-sharing
+     clauses (see firstprivate_head/lastprivate_head in the above function).
+     dev_mark is used to check duplicate mappings for OpenMP
+     is_device_ptr/has_device_addr clauses (see is_on_device_head in above
+     function).
+     gen_mark is used to check duplicate mappings for OpenMP
+     use_device_ptr/use_device_addr/private/shared clauses (see generic_head in
+     above functon).
+     reduc_mark is used to check duplicate mappings for OpenMP reduction
+     clauses.  */
+  struct gfc_symbol *old_symbol;
+  unsigned mark:1, comp_mark:1, data_mark:1, dev_mark:1, gen_mark:1;
+  unsigned reduc_mark:1, gfc_new:1;
+
   /* Nonzero if all equivalences associated with this symbol have been
      processed.  */
   unsigned equiv_built:1;
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 653c43f79ffb..63a14daa6d7b 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -7135,6 +7135,10 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	  continue;
 	n->sym->mark = 0;
 	n->sym->comp_mark = 0;
+	n->sym->data_mark = 0;
+	n->sym->dev_mark = 0;
+	n->sym->gen_mark = 0;
+	n->sym->reduc_mark = 0;
 	if (n->sym->attr.flavor == FL_VARIABLE
 	    || n->sym->attr.proc_pointer
 	    || (!code && (!n->sym->attr.dummy || n->sym->ns != ns)))
@@ -7203,7 +7207,6 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	&& list != OMP_LIST_LASTPRIVATE
 	&& list != OMP_LIST_ALIGNED
 	&& list != OMP_LIST_DEPEND
-	&& (list != OMP_LIST_MAP || openacc)
 	&& list != OMP_LIST_FROM
 	&& list != OMP_LIST_TO
 	&& (list != OMP_LIST_REDUCTION || !openacc)
@@ -7222,10 +7225,43 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	    for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next)
 	      if (ref->type == REF_COMPONENT)
 		component_ref_p = true;
-	  if ((!component_ref_p && n->sym->comp_mark)
-	      || (component_ref_p && n->sym->mark))
-	    gfc_error ("Symbol %qs has mixed component and non-component "
-		       "accesses at %L", n->sym->name, &n->where);
+	  if ((list == OMP_LIST_IS_DEVICE_PTR
+	       || list == OMP_LIST_HAS_DEVICE_ADDR)
+	      && !component_ref_p)
+	    {
+	      if (n->sym->gen_mark || n->sym->dev_mark || n->sym->reduc_mark)
+		gfc_error ("Symbol %qs present on multiple clauses at %L",
+			   n->sym->name, &n->where);
+	      else
+		n->sym->dev_mark = 1;
+	    }
+	  else if ((list == OMP_LIST_USE_DEVICE_PTR
+		    || list == OMP_LIST_USE_DEVICE_ADDR
+		    || list == OMP_LIST_PRIVATE
+		    || list == OMP_LIST_SHARED)
+		   && !component_ref_p)
+	    {
+	      if (n->sym->gen_mark || n->sym->dev_mark || n->sym->reduc_mark)
+		gfc_error ("Symbol %qs present on multiple clauses at %L",
+			   n->sym->name, &n->where);
+	      else
+		n->sym->gen_mark = 1;
+	    }
+	  else if (list == OMP_LIST_REDUCTION && !component_ref_p)
+	    {
+	      if (n->sym->gen_mark || n->sym->dev_mark || n->sym->reduc_mark)
+		gfc_error ("Symbol %qs present on multiple clauses at %L",
+			   n->sym->name, &n->where);
+	      else
+		n->sym->reduc_mark = 1;
+	    }
+	  else if ((!component_ref_p && n->sym->comp_mark)
+		   || (component_ref_p && n->sym->mark))
+	    {
+	      if (openacc)
+		gfc_error ("Symbol %qs has mixed component and non-component "
+			   "accesses at %L", n->sym->name, &n->where);
+	    }
 	  else if (n->sym->mark)
 	    gfc_error ("Symbol %qs present on multiple clauses at %L",
 		       n->sym->name, &n->where);
@@ -7241,31 +7277,44 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
   gcc_assert (OMP_LIST_LASTPRIVATE == OMP_LIST_FIRSTPRIVATE + 1);
   for (list = OMP_LIST_FIRSTPRIVATE; list <= OMP_LIST_LASTPRIVATE; list++)
     for (n = omp_clauses->lists[list]; n; n = n->next)
-      if (n->sym->mark)
+      if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark)
 	{
 	  gfc_error ("Symbol %qs present on multiple clauses at %L",
 		     n->sym->name, &n->where);
-	  n->sym->mark = 0;
+	  n->sym->data_mark = 0;
 	}
+      else if (n->sym->mark
+	       && code->op != EXEC_OMP_TARGET_TEAMS
+	       && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE
+	       && code->op != EXEC_OMP_TARGET_TEAMS_LOOP
+	       && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD
+	       && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO
+	       && code->op != EXEC_OMP_TARGET_PARALLEL
+	       && code->op != EXEC_OMP_TARGET_PARALLEL_DO
+	       && code->op != EXEC_OMP_TARGET_PARALLEL_LOOP
+	       && code->op != EXEC_OMP_TARGET_PARALLEL_DO_SIMD
+	       && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD)
+	gfc_error ("Symbol %qs present on both data and map clauses "
+		   "at %L", n->sym->name, &n->where);
 
   for (n = omp_clauses->lists[OMP_LIST_FIRSTPRIVATE]; n; n = n->next)
     {
-      if (n->sym->mark)
+      if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark)
 	gfc_error ("Symbol %qs present on multiple clauses at %L",
 		   n->sym->name, &n->where);
       else
-	n->sym->mark = 1;
+	n->sym->data_mark = 1;
     }
   for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next)
-    n->sym->mark = 0;
+    n->sym->data_mark = 0;
 
   for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next)
     {
-      if (n->sym->mark)
+      if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark)
 	gfc_error ("Symbol %qs present on multiple clauses at %L",
 		   n->sym->name, &n->where);
       else
-	n->sym->mark = 1;
+	n->sym->data_mark = 1;
     }
 
   for (n = omp_clauses->lists[OMP_LIST_ALIGNED]; n; n = n->next)
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214.f90
new file mode 100644
index 000000000000..25949934e840
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr107214.f90
@@ -0,0 +1,7 @@
+! { dg-do compile }
+
+program p
+   integer, allocatable :: a
+   !$omp target map(tofrom: a, a) ! { dg-error "Symbol 'a' present on multiple clauses" }
+   !$omp end target
+end
-- 
2.29.2

Re: [PATCH 1/2] OpenMP/Fortran: Combined directives with map/firstprivate of same symbol

[Tobias Burnus]

On 07.12.22 20:09, Julian Brown wrote:
> On Wed, 26 Oct 2022 12:39:39 +0200
> Tobias Burnus <tobias@codesourcery.com> wrote:
>> The ICE seems to be because gcc/fortran/trans-openmp.cc's
>> gfc_split_omp_clauses mishandles this as the dump shows the following:
>>
>>     #pragma omp target firstprivate(a) map(tofrom:a)
>>           #pragma omp parallel firstprivate(a)
>>
>> In contrast, for the C testcase:
>>
>> #pragma omp target parallel for simd map(x) firstprivate(x)
>>
>> the dump is as follows, which seems to be sensible:
>>
>>     #pragma omp target map(tofrom:x)
>>           #pragma omp parallel firstprivate(x)
> This patch fixes a case where a combined directive (e.g. "!$omp target
> parallel ...") contains both a map and a firstprivate clause for the
> same variable.  When the combined directive is split into two nested
> directives, the outer "target" gets the "map" clause, and the inner
> "parallel" gets the "firstprivate" clause, like so:
...
> This is not a recent regression, but appears to fix a long-standing ICE.
...
> gcc/fortran/
>          * trans-openmp.cc (gfc_add_firstprivate_if_unmapped): New function.
>          (gfc_split_omp_clauses): Call above.
>
> libgomp/
>          * testsuite/libgomp.fortran/combined-directive-splitting-1.f90: New
>          test.

LGTM – thanks!

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

Re: [PATCH 2/2] OpenMP: Duplicate checking for map clauses in Fortran (PR107214)

[Tobias Burnus]

Hi Julian:

On 07.12.22 20:13, Julian Brown wrote:
>> I know that this was the case before, but can you move the mark:1 etc.
>> after 'tlink'? In that case all bitfields are grouped together.
Thanks for doing so.
>> I wonder whether that also rejects the following – which seems to be
>> valid. The 'map' goes to 'target' and the 'firstprivate' to
>> 'parallel', cf. OpenMP 5.2, "17.2 Clauses on Combined and Composite
>> Constructs", [340:3-4 & 12-14]. (BTW: While some fixes went into 5.1
>> regarding this section, a likewise wording is already in 5.0.)
>>
>> (Testing showed: it give an ICE without the patch and an error with.)
> ...and this patch avoids the error for combined directives, and
> reorders the gfc_symbol bitfields.

All in all, I am fine with the patch - but I spotted some issues.

First, I think you need to set for some error cases mark = 0 to avoid duplicated errors.
Namely:

   ! Outputs the error twice ('Symbol ‘y’ present on multiple clauses')
   !$omp target has_device_addr(y) firstprivate(y)
   block; end block

  * * *

Additionally, I think it would be good to have besides 'target' + map/firstprivate (→ error)
also a testcase for 'target simd' + map/firstprivate → error

And I think also gives-no-error checks all combined 'target ...' that take firstprivate
should be added, cf. your own patch - possibly with looking at the original dump (scan-tree-dump)
to see that the clause is properly attached correctly. Example for 'target teams':

   !$omp target teams map(x) firstprivate(x)
   block; end block

(Works but no testcase.)

  * * *

The following is not diagnosed and gives an ICE:

!$omp target in_reduction(+: x) private(x)
   block; end block
end

The C testcase properly has:
   error: ‘x’ appears more than once in data-sharing clauses

Note: Using 'firstprivate' instead of 'private' shows the proper error also in Fortran.


The following does not ICE but does not make sense (and is rejected in C):

     4 | #pragma omp target private(x) map(x)

vs.

   !$omp target map(x) private(x)
   block; end block

(The latter produces "#pragma omp target private(x.0) map(tofrom:*x.0)", ups!)

  * * *

I also note that 'simd' accepts private such that

#pragma omp target simd private(x) map(x)
  for (int i=0; i < 0; i++)
  ;

!$omp target simd map(x) private(x)
do i = 1, 0; end do

is valid. (It is accepted by gcc and gfortran, i.e. it just needs to be added as testcase.)

  * * *

I note that C rejects {map(x),firstprivate(x)} + {has_device_addr(x),is_device_ptr(x)}',
but gfortran + your patch accepts:

   !$omp target map(x) has_device_addr(x)
   !$omp target map(x) is_device_ptr(x)

while

   !$omp target firstprivate(x) has_device_addr(x)
   !$omp target firstprivate(x) is_device_ptr(x)

is rejected – showing the error message twice.

Expected: I think it should show an error in all four cases - but only once.

> 2022-12-06  Julian Brown  <julian@codesourcery.com>
>
> gcc/fortran/
>          PR fortran/107214
>          * gfortran.h (gfc_symbol): Add data_mark, dev_mark, gen_mark and
>          reduc_mark bitfields.
>          * openmp.cc (resolve_omp_clauses): Use above bitfields to improve
>          duplicate clause detection.
>
> gcc/testsuite/
>          PR fortran/107214
>          * gfortran.dg/gomp/pr107214.f90: New test.

Thanks,

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

Re: [PATCH 2/2] OpenMP: Duplicate checking for map clauses in Fortran (PR107214)

[Julian Brown]

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

On Thu, 8 Dec 2022 13:04:20 +0100
Tobias Burnus <tobias@codesourcery.com> wrote:

> All in all, I am fine with the patch - but I spotted some issues.
> 
> First, I think you need to set for some error cases mark = 0 to avoid
> duplicated errors. Namely:
> 
>    ! Outputs the error twice ('Symbol ‘y’ present on multiple
> clauses') !$omp target has_device_addr(y) firstprivate(y)
>    block; end block
> 
>   * * *
> 
> Additionally, I think it would be good to have besides 'target' +
> map/firstprivate (→ error) also a testcase for 'target simd' +
> map/firstprivate → error
> 
> And I think also gives-no-error checks all combined 'target ...' that
> take firstprivate should be added, cf. your own patch - possibly with
> looking at the original dump (scan-tree-dump) to see that the clause
> is properly attached correctly. Example for 'target teams':
> 
>    !$omp target teams map(x) firstprivate(x)
>    block; end block
> 
> (Works but no testcase.)
> 
>   * * *
> 
> The following is not diagnosed and gives an ICE:
> 
> !$omp target in_reduction(+: x) private(x)
>    block; end block
> end
> 
> The C testcase properly has:
>    error: ‘x’ appears more than once in data-sharing clauses
> 
> Note: Using 'firstprivate' instead of 'private' shows the proper
> error also in Fortran.
> 
> 
> The following does not ICE but does not make sense (and is rejected
> in C):
> 
>      4 | #pragma omp target private(x) map(x)
> 
> vs.
> 
>    !$omp target map(x) private(x)
>    block; end block
> 
> (The latter produces "#pragma omp target private(x.0)
> map(tofrom:*x.0)", ups!)
> 
>   * * *
> 
> I also note that 'simd' accepts private such that
> 
> #pragma omp target simd private(x) map(x)
>   for (int i=0; i < 0; i++)
>   ;
> 
> !$omp target simd map(x) private(x)
> do i = 1, 0; end do
> 
> is valid. (It is accepted by gcc and gfortran, i.e. it just needs to
> be added as testcase.)
> 
>   * * *
> 
> I note that C rejects {map(x),firstprivate(x)} +
> {has_device_addr(x),is_device_ptr(x)}', but gfortran + your patch
> accepts:
> 
>    !$omp target map(x) has_device_addr(x)
>    !$omp target map(x) is_device_ptr(x)
> 
> while
> 
>    !$omp target firstprivate(x) has_device_addr(x)
>    !$omp target firstprivate(x) is_device_ptr(x)
> 
> is rejected – showing the error message twice.
> 
> Expected: I think it should show an error in all four cases - but
> only once.

I believe this patch covers all the above cases (hopefully
appropriately generalised), at least for Fortran. I haven't attempted
to fix any missing cases for C, for now.

Re-tested with offloading to NVPTX (with a few supporting patches, as
before).

Does this look OK now?

Thanks,

Julian

[-- Attachment #2: pr107214-3.diff --]
[-- Type: text/x-patch, Size: 18217 bytes --]

commit c5e0cb20a07fcce4822dba2731b7af7b372b3376
Author: Julian Brown <julian@codesourcery.com>
Date:   Tue Dec 6 23:10:58 2022 +0000

    OpenMP: Duplicate checking for map clauses in Fortran (PR107214)
    
    This patch adds duplicate checking for OpenMP "map" clauses, taking some
    cues from the implementation for C in c-typeck.cc:c_finish_omp_clauses
    (and similar for C++).
    
    In addition to the existing use of the "mark" and "comp_mark" bitfields
    in the gfc_symbol structure, the patch adds several new bits handling
    duplicate checking within various categories of clause types.  If "mark"
    is being used for map clauses, we need to use different bits for other
    clauses for cases where "map" and some other clause can refer to the
    same symbol (e.g. "map(n) shared(n)").
    
    2022-12-06  Julian Brown  <julian@codesourcery.com>
    
    gcc/fortran/
            PR fortran/107214
            * gfortran.h (gfc_symbol): Add data_mark, dev_mark, gen_mark and
            reduc_mark bitfields.
            * openmp.cc (resolve_omp_clauses): Use above bitfields to improve
            duplicate clause detection.
    
    gcc/testsuite/
            PR fortran/107214
            * gfortran.dg/gomp/pr107214.f90: New test.
            * gfortran.dg/gomp/pr107214-2.f90: New test.
            * gfortran.dg/gomp/pr107214-3.f90: New test.
            * gfortran.dg/gomp/pr107214-4.f90: New test.
            * gfortran.dg/gomp/pr107214-5.f90: New test.
            * gfortran.dg/gomp/pr107214-6.f90: New test.
            * gfortran.dg/gomp/pr107214-7.f90: New test.
            * gfortran.dg/gomp/pr107214-8.f90: New test.

diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 612a54f4cc49..990ad377e5b2 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1871,16 +1871,6 @@ typedef struct gfc_symbol
 
   gfc_namelist *namelist, *namelist_tail;
 
-  /* Change management fields.  Symbols that might be modified by the
-     current statement have the mark member nonzero.  Of these symbols,
-     symbols with old_symbol equal to NULL are symbols created within
-     the current statement.  Otherwise, old_symbol points to a copy of
-     the old symbol. gfc_new is used in symbol.cc to flag new symbols.
-     comp_mark is used to indicate variables which have component accesses
-     in OpenMP/OpenACC directive clauses.  */
-  struct gfc_symbol *old_symbol;
-  unsigned mark:1, comp_mark:1, gfc_new:1;
-
   /* The tlink field is used in the front end to carry the module
      declaration of separate module procedures so that the characteristics
      can be compared with the corresponding declaration in a submodule. In
@@ -1888,6 +1878,28 @@ typedef struct gfc_symbol
      deferred initialization.  */
   struct gfc_symbol *tlink;
 
+  /* Change management fields.  Symbols that might be modified by the
+     current statement have the mark member nonzero.  Of these symbols,
+     symbols with old_symbol equal to NULL are symbols created within
+     the current statement.  Otherwise, old_symbol points to a copy of
+     the old symbol. gfc_new is used in symbol.cc to flag new symbols.
+     comp_mark is used to indicate variables which have component accesses
+     in OpenMP/OpenACC directive clauses (cf. c-typeck.cc:c_finish_omp_clauses,
+     map_field_head).
+     data_mark is used to check duplicate mappings for OpenMP data-sharing
+     clauses (see firstprivate_head/lastprivate_head in the above function).
+     dev_mark is used to check duplicate mappings for OpenMP
+     is_device_ptr/has_device_addr clauses (see is_on_device_head in above
+     function).
+     gen_mark is used to check duplicate mappings for OpenMP
+     use_device_ptr/use_device_addr/private/shared clauses (see generic_head in
+     above functon).
+     reduc_mark is used to check duplicate mappings for OpenMP reduction
+     clauses.  */
+  struct gfc_symbol *old_symbol;
+  unsigned mark:1, comp_mark:1, data_mark:1, dev_mark:1, gen_mark:1;
+  unsigned reduc_mark:1, gfc_new:1;
+
   /* Nonzero if all equivalences associated with this symbol have been
      processed.  */
   unsigned equiv_built:1;
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 862c649b0b62..70edbe7595d1 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -7131,6 +7131,10 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	  continue;
 	n->sym->mark = 0;
 	n->sym->comp_mark = 0;
+	n->sym->data_mark = 0;
+	n->sym->dev_mark = 0;
+	n->sym->gen_mark = 0;
+	n->sym->reduc_mark = 0;
 	if (n->sym->attr.flavor == FL_VARIABLE
 	    || n->sym->attr.proc_pointer
 	    || (!code && (!n->sym->attr.dummy || n->sym->ns != ns)))
@@ -7199,14 +7203,9 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	&& list != OMP_LIST_LASTPRIVATE
 	&& list != OMP_LIST_ALIGNED
 	&& list != OMP_LIST_DEPEND
-	&& (list != OMP_LIST_MAP || openacc)
 	&& list != OMP_LIST_FROM
 	&& list != OMP_LIST_TO
 	&& (list != OMP_LIST_REDUCTION || !openacc)
-	&& list != OMP_LIST_REDUCTION_INSCAN
-	&& list != OMP_LIST_REDUCTION_TASK
-	&& list != OMP_LIST_IN_REDUCTION
-	&& list != OMP_LIST_TASK_REDUCTION
 	&& list != OMP_LIST_ALLOCATE)
       for (n = omp_clauses->lists[list]; n; n = n->next)
 	{
@@ -7218,10 +7217,58 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	    for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next)
 	      if (ref->type == REF_COMPONENT)
 		component_ref_p = true;
-	  if ((!component_ref_p && n->sym->comp_mark)
-	      || (component_ref_p && n->sym->mark))
-	    gfc_error ("Symbol %qs has mixed component and non-component "
-		       "accesses at %L", n->sym->name, &n->where);
+	  if ((list == OMP_LIST_IS_DEVICE_PTR
+	       || list == OMP_LIST_HAS_DEVICE_ADDR)
+	      && !component_ref_p)
+	    {
+	      if (n->sym->gen_mark
+		  || n->sym->dev_mark
+		  || n->sym->reduc_mark
+		  || n->sym->mark)
+		gfc_error ("Symbol %qs present on multiple clauses at %L",
+			   n->sym->name, &n->where);
+	      else
+		n->sym->dev_mark = 1;
+	    }
+	  else if ((list == OMP_LIST_USE_DEVICE_PTR
+		    || list == OMP_LIST_USE_DEVICE_ADDR
+		    || list == OMP_LIST_PRIVATE
+		    || list == OMP_LIST_SHARED)
+		   && !component_ref_p)
+	    {
+	      if (n->sym->gen_mark || n->sym->dev_mark || n->sym->reduc_mark)
+		gfc_error ("Symbol %qs present on multiple clauses at %L",
+			   n->sym->name, &n->where);
+	      else
+		{
+		  n->sym->gen_mark = 1;
+		  /* Set both generic and device bits if we have
+		     use_device_*(x) or shared(x).  This allows us to diagnose
+		     "map(x) private(x)" below.  */
+		  if (list != OMP_LIST_PRIVATE)
+		    n->sym->dev_mark = 1;
+		}
+	    }
+	  else if ((list == OMP_LIST_REDUCTION
+		    || list == OMP_LIST_REDUCTION_TASK
+		    || list == OMP_LIST_REDUCTION_INSCAN
+		    || list == OMP_LIST_IN_REDUCTION
+		    || list == OMP_LIST_TASK_REDUCTION)
+		   && !component_ref_p)
+	    {
+	      /* Attempts to mix reduction types are diagnosed below.  */
+	      if (n->sym->gen_mark || n->sym->dev_mark)
+		gfc_error ("Symbol %qs present on multiple clauses at %L",
+			   n->sym->name, &n->where);
+	      n->sym->reduc_mark = 1;
+	    }
+	  else if ((!component_ref_p && n->sym->comp_mark)
+		   || (component_ref_p && n->sym->mark))
+	    {
+	      if (openacc)
+		gfc_error ("Symbol %qs has mixed component and non-component "
+			   "accesses at %L", n->sym->name, &n->where);
+	    }
 	  else if (n->sym->mark)
 	    gfc_error ("Symbol %qs present on multiple clauses at %L",
 		       n->sym->name, &n->where);
@@ -7234,34 +7281,62 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	    }
 	}
 
+  /* Detect specifically the case where we have "map(x) private(x)" and raise
+     an error.  If we have "...simd" combined directives though, the "private"
+     applies to the simd part, so this is permitted though.  */
+  for (n = omp_clauses->lists[OMP_LIST_PRIVATE]; n; n = n->next)
+    if (n->sym->mark
+	&& n->sym->gen_mark
+	&& !n->sym->dev_mark
+	&& !n->sym->reduc_mark
+	&& code->op != EXEC_OMP_TARGET_SIMD
+	&& code->op != EXEC_OMP_TARGET_PARALLEL_DO_SIMD
+	&& code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD
+	&& code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD)
+      gfc_error ("Symbol %qs present on multiple clauses at %L",
+		 n->sym->name, &n->where);
+
   gcc_assert (OMP_LIST_LASTPRIVATE == OMP_LIST_FIRSTPRIVATE + 1);
   for (list = OMP_LIST_FIRSTPRIVATE; list <= OMP_LIST_LASTPRIVATE; list++)
     for (n = omp_clauses->lists[list]; n; n = n->next)
-      if (n->sym->mark)
+      if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark)
 	{
 	  gfc_error ("Symbol %qs present on multiple clauses at %L",
 		     n->sym->name, &n->where);
-	  n->sym->mark = 0;
+	  n->sym->data_mark = n->sym->gen_mark = n->sym->dev_mark = 0;
 	}
+      else if (n->sym->mark
+	       && code->op != EXEC_OMP_TARGET_TEAMS
+	       && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE
+	       && code->op != EXEC_OMP_TARGET_TEAMS_LOOP
+	       && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD
+	       && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO
+	       && code->op != EXEC_OMP_TARGET_PARALLEL
+	       && code->op != EXEC_OMP_TARGET_PARALLEL_DO
+	       && code->op != EXEC_OMP_TARGET_PARALLEL_LOOP
+	       && code->op != EXEC_OMP_TARGET_PARALLEL_DO_SIMD
+	       && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD)
+	gfc_error ("Symbol %qs present on both data and map clauses "
+		   "at %L", n->sym->name, &n->where);
 
   for (n = omp_clauses->lists[OMP_LIST_FIRSTPRIVATE]; n; n = n->next)
     {
-      if (n->sym->mark)
+      if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark)
 	gfc_error ("Symbol %qs present on multiple clauses at %L",
 		   n->sym->name, &n->where);
       else
-	n->sym->mark = 1;
+	n->sym->data_mark = 1;
     }
   for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next)
-    n->sym->mark = 0;
+    n->sym->data_mark = 0;
 
   for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next)
     {
-      if (n->sym->mark)
+      if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark)
 	gfc_error ("Symbol %qs present on multiple clauses at %L",
 		   n->sym->name, &n->where);
       else
-	n->sym->mark = 1;
+	n->sym->data_mark = 1;
     }
 
   for (n = omp_clauses->lists[OMP_LIST_ALIGNED]; n; n = n->next)
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-2.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-2.f90
new file mode 100644
index 000000000000..da47e40f359b
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-2.f90
@@ -0,0 +1,6 @@
+integer :: y
+
+!$omp target has_device_addr(y) firstprivate(y)  ! { dg-error "Symbol 'y' present on multiple clauses" }
+!$omp end target
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-3.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-3.f90
new file mode 100644
index 000000000000..526152e11018
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-3.f90
@@ -0,0 +1,14 @@
+program p
+integer :: y
+
+!$omp target map(y) firstprivate(y)  ! { dg-error "Symbol 'y' present on both data and map clauses" }
+y = y + 1
+!$omp end target
+
+!$omp target simd map(y) firstprivate(y)  ! { dg-error "Symbol 'y' present on both data and map clauses" }
+do i=1,1
+  y = y + 1
+end do
+!$omp end target simd
+
+end program p
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-4.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-4.f90
new file mode 100644
index 000000000000..b4f343a17acc
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-4.f90
@@ -0,0 +1,147 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+
+integer :: x, y
+
+! EXEC_OMP_TARGET_TEAMS
+
+!$omp target teams map(x) firstprivate(x)
+x = x + 1
+!$omp end target teams
+
+!$omp target teams map(x) firstprivate(y)
+x = y + 1
+!$omp end target teams
+
+! EXEC_OMP_TARGET_TEAMS_DISTRIBUTE
+
+!$omp target teams distribute map(x) firstprivate(x)
+do i=1,1
+  x = x + 1
+end do
+!$omp end target teams distribute
+
+!$omp target teams distribute map(x) firstprivate(y)
+do i=1,1
+  x = y + 1
+end do
+!$omp end target teams distribute
+
+! EXEC_OMP_TARGET_TEAMS_LOOP
+
+!$omp target teams loop map(x) firstprivate(x)
+do i=1,1
+  x = x + 1
+end do
+!$omp end target teams loop
+
+!$omp target teams loop map(x) firstprivate(y)
+do i=1,1
+  x = y + 1
+end do
+!$omp end target teams loop
+
+! EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD
+
+!$omp target teams distribute simd map(x) firstprivate(x)
+do i=1,1
+  x = x + 1
+end do
+!$omp end target teams distribute simd
+
+!$omp target teams distribute simd map(x) firstprivate(y)
+do i=1,1
+  x = y + 1
+end do
+!$omp end target teams distribute simd
+
+! EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO
+
+!$omp target teams distribute parallel do map(x) firstprivate(x)
+do i=1,1
+  x = x + 1
+end do
+!$omp end target teams distribute parallel do
+
+!$omp target teams distribute parallel do map(x) firstprivate(y)
+do i=1,1
+  x = y + 1
+end do
+!$omp end target teams distribute parallel do
+
+! EXEC_OMP_TARGET_PARALLEL
+
+!$omp target parallel map(x) firstprivate(x)
+x = x + 1
+!$omp end target parallel
+
+!$omp target parallel map(x) firstprivate(y)
+x = y + 1
+!$omp end target parallel
+
+! EXEC_OMP_TARGET_PARALLEL_DO
+
+!$omp target parallel do map(x) firstprivate(x)
+do i=1,1
+  x = x + 1
+end do
+!$omp end target parallel do
+
+!$omp target parallel do map(x) firstprivate(y)
+do i=1,1
+  x = y + 1
+end do
+!$omp end target parallel do
+
+! EXEC_OMP_TARGET_PARALLEL_LOOP
+
+!$omp target parallel loop map(x) firstprivate(x)
+do i=1,1
+  x = x + 1
+end do
+!$omp end target parallel loop
+
+!$omp target parallel loop map(x) firstprivate(y)
+do i=1,1
+  x = y + 1
+end do
+!$omp end target parallel loop
+
+! EXEC_OMP_TARGET_PARALLEL_DO_SIMD
+
+!$omp target parallel do simd map(x) firstprivate(x)
+do i=1,1
+  x = x + 1
+end do
+!$omp end target parallel do simd
+
+!$omp target parallel do simd map(x) firstprivate(y)
+do i=1,1
+  x = y + 1
+end do
+!$omp end target parallel do simd
+
+! EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD
+
+!$omp target teams distribute parallel do simd map(x) firstprivate(x)
+do i=1,1
+  x = x + 1
+end do
+!$omp end target teams distribute parallel do simd
+
+!$omp target teams distribute parallel do simd map(x) firstprivate(y)
+do i=1,1
+  x = y + 1
+end do
+!$omp end target teams distribute parallel do simd
+
+! { dg-final { scan-tree-dump-times {omp target map\(tofrom:x\)} 10 "original" } }
+! { dg-final { scan-tree-dump-times {omp target firstprivate\(y\) map\(tofrom:x\)} 10 "original" } }
+
+! { dg-final { scan-tree-dump-times {omp teams firstprivate\(x\)} 6 "original" } }
+! { dg-final { scan-tree-dump-times {omp teams firstprivate\(y\)} 6 "original" } }
+
+! { dg-final { scan-tree-dump-times {omp parallel firstprivate\(x\)} 6 "original" } }
+! { dg-final { scan-tree-dump-times {omp parallel firstprivate\(y\)} 6 "original" } }
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-5.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-5.f90
new file mode 100644
index 000000000000..08a9f62b0881
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-5.f90
@@ -0,0 +1,11 @@
+integer :: x, y
+
+!$omp target in_reduction(+: x) private(x)  ! { dg-error "Symbol 'x' present on multiple clauses" }
+x = x + 1
+!$omp end target
+
+!$omp target in_reduction(+: y) firstprivate(y)  ! { dg-error "Symbol 'y' present on both data and map clauses" }
+y = y + 1
+!$omp end target
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-6.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-6.f90
new file mode 100644
index 000000000000..0a1270645518
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-6.f90
@@ -0,0 +1,9 @@
+! { dg-do compile }
+
+integer :: x
+
+!$omp target map(x) private(x)  ! { dg-error "Symbol 'x' present on multiple clauses" }
+x = x + 1
+!$omp end target
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-7.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-7.f90
new file mode 100644
index 000000000000..125d1bc4fed4
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-7.f90
@@ -0,0 +1,33 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+
+integer :: x
+
+!$omp target simd map(x) private(x)
+do i=1,1
+  x = x + 1
+end do
+!$omp end target simd
+
+!$omp target teams distribute simd map(x) private(x)
+do i=1,1
+  x = x + 1
+end do
+!$omp end target teams distribute simd
+
+!$omp target parallel do simd map(x) private(x)
+do i=1,1
+  x = x + 1
+end do
+!$omp end target parallel do simd
+
+!$omp target teams distribute parallel do simd map(x) private(x)
+do i=1,1
+  x = x + 1
+end do
+!$omp end target teams distribute parallel do simd
+
+! { dg-final { scan-tree-dump-times {omp target map\(tofrom:x\)} 4 "original" } }
+! { dg-final { scan-tree-dump-times {(?n)omp simd.* private\(x\)} 4 "original" } }
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-8.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-8.f90
new file mode 100644
index 000000000000..192c97a33e5c
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-8.f90
@@ -0,0 +1,18 @@
+! { dg-do compile }
+
+integer, allocatable :: x
+integer, pointer :: y
+
+!$omp target map(x) has_device_addr(x)  ! { dg-error "Symbol 'x' present on multiple clauses" }
+!$omp end target
+
+!$omp target map(y) is_device_ptr(y)  ! { dg-error "Symbol 'y' present on multiple clauses" }
+!$omp end target
+
+!$omp target firstprivate(x) has_device_addr(x)  ! { dg-error "Symbol 'x' present on multiple clauses" }
+!$omp end target
+
+!$omp target firstprivate(y) is_device_ptr(y)  ! { dg-error "Symbol 'y' present on multiple clauses" }
+!$omp end target
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214.f90
new file mode 100644
index 000000000000..25949934e840
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr107214.f90
@@ -0,0 +1,7 @@
+! { dg-do compile }
+
+program p
+   integer, allocatable :: a
+   !$omp target map(tofrom: a, a) ! { dg-error "Symbol 'a' present on multiple clauses" }
+   !$omp end target
+end

Re: [PATCH 2/2] OpenMP: Duplicate checking for map clauses in Fortran (PR107214)

[Tobias Burnus]

Hi Julian,

On 10.12.22 13:10, Julian Brown wrote:
> On Thu, 8 Dec 2022 13:04:20 +0100
> Tobias Burnus <tobias@codesourcery.com> wrote:
>> All in all, I am fine with the patch - but I spotted some issues.
...
> I believe this patch covers all the above cases (hopefully
> appropriately generalised), at least for Fortran. I haven't attempted
> to fix any missing cases for C, for now.
>
> Re-tested with offloading to NVPTX (with a few supporting patches, as
> before).
>
> Does this look OK now?

Yes, LGTM.

Thanks!

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