Re: [PATCH v5 2/4] OpenMP/OpenACC: Rework clause expansion and nested struct handling

[Julian Brown <julian@codesourcery.com>]

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

On Wed, 7 Dec 2022 17:13:42 +0100
Tobias Burnus <tobias@codesourcery.com> wrote:

> I think we need to distinguish:
> 
>    #pragma omp target enter data map(to: s.w[:10])
> 
> from
> 
>    #pragma omp target map(tofrom: s.arr[:20])
>      s.arr[0] = 5;
> 
> As in the latter case 's' gets implicitly mapped and then applies to
> the base pointer 's.arr' of 's.arr[:20]'. While in the former case,
> only the pointee gets mapped without the pointer 's.arr' (and, hence,
> there is also no pointer attachment).

Here's an incremental patch that fixes the mapping behaviour in that
case. This is to be applied on top of the approved (but not committed)
parent patch:

  https://gcc.gnu.org/pipermail/gcc-patches/2022-October/603792.html

and also the unreviewed patch posted here (ping?):

  https://gcc.gnu.org/pipermail/gcc-patches/2022-November/607543.html

though it might actually make more sense to commit the three patches
squashed together.

Tested with offloading to NVPTX. OK?

Thanks,

Julian

[-- Attachment #2: implicitly-map-base-pointer-1.diff --]
[-- Type: text/x-patch, Size: 7943 bytes --]

commit abb1e04f9ef93221ecd4292f43cc1ea901843766
Author: Julian Brown <julian@codesourcery.com>
Date:   Thu Dec 8 13:31:01 2022 +0000

    OpenMP: implicitly map base pointer for array-section pointer components
    
    Following from discussion in:
    
      https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570075.html
    
    and:
    
      https://gcc.gnu.org/pipermail/gcc-patches/2022-December/608100.html
    
    and also upstream OpenMP issue 342, this patch changes mapping for array
    sections of pointer components on compute regions like this:
    
      #pragma omp target map(s.ptr[0:10])
      {
        ...use of 's'...
      }
    
    so the base pointer 's.ptr' is implicitly mapped, and thus pointer
    attachment happens.  This is subtly different in the "enter data"
    case, e.g:
    
      #pragma omp target enter data map(s.ptr[0:10])
    
    if 's.ptr' (or the whole of 's') is not present on the target before
    the directive is executed, the array section is copied to the target
    but pointer attachment does *not* take place, since 's' (or 's.ptr')
    is not mapped implicitly for "enter data".
    
    To get a pointer attachment with "enter data", you can do, e.g:
    
      #pragma omp target enter data map(s.ptr, s.ptr[0:10])
    
      #pragma omp target
      {
        ...implicit use of 's'...
      }
    
    That is, once the attachment has happened, implicit mapping of 's'
    and uses of 's.ptr[...]' work correctly in the target region.
    
    ChangeLog
    
    2022-12-12  Julian Brown  <julian@codesourcery.com>
    
    gcc/
            * gimplify.cc (omp_accumulate_sibling_list): Don't require
            explicitly-mapped base pointer for compute regions.
    
    gcc/testsuite/
            * c-c++-comon/gomp/target-implicit-map-2.c: Update expected scan output.
    
    libgomp/
            * testsuite/libgomp.c-c++-common/target-implicit-map-2.c: Fix missing
            "free".
            * testsuite/libgomp.c-c++-common/target-implicit-map-3.c: New test.
            * testsuite/libgomp.c-c++-common/target-map-zlas-1.c: New test.
            * testsuite/libgomp.c/target-22.c: Remove explicit base pointer
            mappings.

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 34cac30d7d92..a8dd298559e8 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -10617,6 +10617,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
   poly_int64 cbitpos;
   tree ocd = OMP_CLAUSE_DECL (grp_end);
   bool openmp = !(region_type & ORT_ACC);
+  bool target = (region_type & ORT_TARGET) != 0;
   tree *continue_at = NULL;
 
   while (TREE_CODE (ocd) == ARRAY_REF)
@@ -10721,9 +10722,9 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
 	    }
 
 	  /* For OpenMP semantics, we don't want to implicitly allocate
-	     space for the pointer here.  A FRAGILE_P node is only being
-	     created so that omp-low.cc is able to rewrite the struct
-	     properly.
+	     space for the pointer here for non-compute regions (e.g. "enter
+	     data").  A FRAGILE_P node is only being created so that
+	     omp-low.cc is able to rewrite the struct properly.
 	     For references (to pointers), we want to actually allocate the
 	     space for the reference itself in the sorted list following the
 	     struct node.
@@ -10731,6 +10732,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
 	     mapping of the attachment point, but not otherwise.  */
 	  if (*fragile_p
 	      || (openmp
+		  && !target
 		  && attach_detach
 		  && TREE_CODE (TREE_TYPE (ocd)) == POINTER_TYPE
 		  && !OMP_CLAUSE_ATTACHMENT_MAPPING_ERASED (grp_end)))
@@ -11043,6 +11045,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
 
 	  if (*fragile_p
 	      || (openmp
+		  && !target
 		  && attach_detach
 		  && TREE_CODE (TREE_TYPE (ocd)) == POINTER_TYPE
 		  && !OMP_CLAUSE_ATTACHMENT_MAPPING_ERASED (grp_end)))
diff --git a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c
index 5ba1d7efe08d..222272df5b1e 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c
@@ -49,4 +49,5 @@ main (void)
 
 /* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:a \[len: [0-9]+\]\[implicit\]\)} "gimple" } } */
 
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(struct:a \[len: 1\]\) map\(alloc:a\.ptr \[len: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:a\.ptr \[bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(struct:a \[len: 1\]\) map\(alloc:a\.ptr \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:a\.ptr \[bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump-not {map\(struct:a \[len: 1\]\) map\(alloc:a\.ptr \[len: 0\]\)} "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c
index 974a9786c3f6..4c49cd091c38 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c
@@ -42,5 +42,7 @@ main (void)
 
   #pragma omp target exit data map(from:a.ptr, a.ptr[:N])
 
+  free (a.ptr);
+
   return 0;
 }
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c
new file mode 100644
index 000000000000..81a7752685c5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c
@@ -0,0 +1,50 @@
+#include <stdlib.h>
+
+#define N 10
+
+struct S
+{
+  int a, b;
+  int *ptr;
+  int c, d;
+};
+
+int
+main (void)
+{
+  struct S a;
+  a.ptr = (int *) malloc (sizeof (int) * N);
+
+  for (int i = 0; i < N; i++)
+    a.ptr[i] = 0;
+
+  #pragma omp target enter data map(to: a.ptr)
+  #pragma omp target enter data map(to: a.ptr[:N])
+
+  #pragma omp target
+  for (int i = 0; i < N; i++)
+    a.ptr[i] += 1;
+
+  #pragma omp target update from(a.ptr[:N])
+
+  for (int i = 0; i < N; i++)
+    if (a.ptr[i] != 1)
+      abort ();
+
+  #pragma omp target map(a.ptr[:N])
+  for (int i = 0; i < N; i++)
+    a.ptr[i] += 1;
+
+  #pragma omp target update from(a.ptr[:N])
+
+  for (int i = 0; i < N; i++)
+    if (a.ptr[i] != 2)
+      abort ();
+
+  #pragma omp target exit data map(release: a.ptr[:N])
+  #pragma omp target exit data map(release: a.ptr)
+
+  free (a.ptr);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-zlas-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-zlas-1.c
new file mode 100644
index 000000000000..1ec0c9a0d5f9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-zlas-1.c
@@ -0,0 +1,36 @@
+#include <stdlib.h>
+
+#define N 10
+
+struct S
+{
+  int a, b;
+  int *ptr;
+  int c, d;
+};
+
+int
+main (void)
+{
+  struct S a;
+  a.ptr = (int *) malloc (sizeof (int) * N);
+
+  for (int i = 0; i < N; i++)
+    a.ptr[i] = 0;
+
+  #pragma omp target enter data map(to: a.ptr[:N])
+
+  #pragma omp target map(a, a.ptr[:0])
+  for (int i = 0; i < N; i++)
+    a.ptr[i] += 1;
+
+  #pragma omp target exit data map(from: a.ptr[:N])
+
+  for (int i = 0; i < N; i++)
+    if (a.ptr[i] != 1)
+      abort ();
+
+  free (a.ptr);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-22.c b/libgomp/testsuite/libgomp.c/target-22.c
index 492744ad0efd..aad8a0a09df7 100644
--- a/libgomp/testsuite/libgomp.c/target-22.c
+++ b/libgomp/testsuite/libgomp.c/target-22.c
@@ -21,8 +21,7 @@ main ()
   s.v.b = a + 16;
   s.w = c + 3;
   int err = 0;
-  #pragma omp target map (to: s.w, s.v.b, s.u, s.s) \
-		     map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \
+  #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \
 		     map (tofrom:s.s[3:3]) \
 		     map (from: s.w[z:4], err) private (i)
   {