openmp: Implement OpenMP 5.0 base-pointer attachement and clause ordering

This patch implements some parts of the target variable mapping changes
specified in OpenMP 5.0, including base-pointer attachment/detachment
behavior for array section list-items in map clauses, and ordering of
map clauses according to map kind.

2020-11-10  Chung-Lin Tang  <cltang@codesourcery.com>

gcc/c-family/ChangeLog:

	* c-common.h (c_omp_adjust_map_clauses): New declaration.
	* c-omp.c (struct map_clause): Helper type for c_omp_adjust_map_clauses.
	(c_omp_adjust_map_clauses): New function.

gcc/c/ChangeLog:

	* c-parser.c (c_parser_omp_target_data): Add use of
	new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as
	handled map clause kind.
	(c_parser_omp_target_enter_data): Likewise.
	(c_parser_omp_target_exit_data): Likewise.
	(c_parser_omp_target): Likewise.
	* c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to
	use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type.
	(c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and
	same struct field access to co-exist on OpenMP construct.

gcc/cp/ChangeLog:

	* parser.c (cp_parser_omp_target_data): Add use of
	new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as
	handled map clause kind.
	(cp_parser_omp_target_enter_data): Likewise.
	(cp_parser_omp_target_exit_data): Likewise.
	(cp_parser_omp_target): Likewise.
	* semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to
	use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix
	interaction between reference case and attach/detach.
	(finish_omp_clauses): Adjust bitmap checks to allow struct decl and
	same struct field access to co-exist on OpenMP construct.

gcc/ChangeLog:

	* gimplify.c (is_or_contains_p): New static helper function.
	(omp_target_reorder_clauses): New function.
	(gimplify_scan_omp_clauses): Add use of omp_target_reorder_clauses to
	reorder clause list according to OpenMP 5.0 rules. Add handling of
	GOMP_MAP_ATTACH_DETACH for OpenMP cases.
	* omp-low.c (is_omp_target): New static helper function.
	(scan_sharing_clauses): Add scan phase handling of GOMP_MAP_ATTACH/DETACH
	for OpenMP cases.
	(lower_omp_target): Add lowering handling of GOMP_MAP_ATTACH/DETACH for
	OpenMP cases.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/clauses-2.c: Remove dg-error cases now valid.
	* gfortran.dg/gomp/map-2.f90: Likewise.
	* c-c++-common/gomp/map-5.c: New testcase.

libgomp/ChangeLog:

	* libgomp.h (enum gomp_map_vars_kind): Adjust enum values to be bit-flag
	usable.
	* oacc-mem.c (acc_map_data): Adjust gomp_map_vars argument flags to
	'GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA'.
	(goacc_enter_datum): Likewise for call to gomp_map_vars_async.
	(goacc_enter_data_internal): Likewise.
	* target.c (gomp_map_vars_internal):
	Change checks of GOMP_MAP_VARS_ENTER_DATA to use bit-and (&). Adjust use
	of gomp_attach_pointer for OpenMP cases.
	(gomp_exit_data): Add handling of GOMP_MAP_DETACH.
	(GOMP_target_enter_exit_data): Add handling of GOMP_MAP_ATTACH.
	* testsuite/libgomp.c-c++-common/ptr-attach-1.c: New testcase.

4 files changed, 122 insertions, 13 deletions

diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index da7ac03..0cc3f4d 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1162,10 +1162,10 @@ struct gomp_device_descr
 /* Kind of the pragma, for which gomp_map_vars () is called.  */
 enum gomp_map_vars_kind
 {
-  GOMP_MAP_VARS_OPENACC,
-  GOMP_MAP_VARS_TARGET,
-  GOMP_MAP_VARS_DATA,
-  GOMP_MAP_VARS_ENTER_DATA
+  GOMP_MAP_VARS_OPENACC    = 1,
+  GOMP_MAP_VARS_TARGET     = 2,
+  GOMP_MAP_VARS_DATA       = 4,
+  GOMP_MAP_VARS_ENTER_DATA = 8
 };
 
 extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *,
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 65757ab..4c8f0e0 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -403,7 +403,8 @@ acc_map_data (void *h, void *d, size_t s)
 
       struct target_mem_desc *tgt
 	= gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes,
-			 &kinds, true, GOMP_MAP_VARS_ENTER_DATA);
+			 &kinds, true,
+			 GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA);
       assert (tgt);
       assert (tgt->list_count == 1);
       splay_tree_key n = tgt->list[0].key;
@@ -572,7 +573,8 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async)
 
       struct target_mem_desc *tgt
 	= gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes,
-			       kinds, true, GOMP_MAP_VARS_ENTER_DATA);
+			       kinds, true, (GOMP_MAP_VARS_OPENACC
+					     | GOMP_MAP_VARS_ENTER_DATA));
       assert (tgt);
       assert (tgt->list_count == 1);
       n = tgt->list[0].key;
@@ -1202,7 +1204,8 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 	  struct target_mem_desc *tgt
 	    = gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL,
 				   &sizes[i], &kinds[i], true,
-				   GOMP_MAP_VARS_ENTER_DATA);
+				   (GOMP_MAP_VARS_OPENACC
+				    | GOMP_MAP_VARS_ENTER_DATA));
 	  assert (tgt);
 
 	  gomp_mutex_lock (&acc_dev->lock);
diff --git a/libgomp/target.c b/libgomp/target.c
index 3432a83..6152f58 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -683,7 +683,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
   struct target_mem_desc *tgt
     = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
   tgt->list_count = mapnum;
-  tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1;
+  tgt->refcount = (pragma_kind & GOMP_MAP_VARS_ENTER_DATA) ? 0 : 1;
   tgt->device_descr = devicep;
   tgt->prev = NULL;
   struct gomp_coalesce_buf cbuf, *cbufp = NULL;
@@ -1212,15 +1212,16 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 		      /* OpenACC 'attach'/'detach' doesn't affect
 			 structured/dynamic reference counts ('n->refcount',
 			 'n->dynamic_refcount').  */
+
+		      gomp_attach_pointer (devicep, aq, mem_map, n,
+					   (uintptr_t) hostaddrs[i], sizes[i],
+					   cbufp);
 		    }
-		  else
+		  else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0)
 		    {
 		      gomp_mutex_unlock (&devicep->lock);
 		      gomp_fatal ("outer struct not mapped for attach");
 		    }
-		  gomp_attach_pointer (devicep, aq, mem_map, n,
-				       (uintptr_t) hostaddrs[i], sizes[i],
-				       cbufp);
 		  continue;
 		}
 	      default:
@@ -1415,7 +1416,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
   /* If the variable from "omp target enter data" map-list was already mapped,
      tgt is not needed.  Otherwise tgt will be freed by gomp_unmap_vars or
      gomp_exit_data.  */
-  if (pragma_kind == GOMP_MAP_VARS_ENTER_DATA && tgt->refcount == 0)
+  if ((pragma_kind & GOMP_MAP_VARS_ENTER_DATA) && tgt->refcount == 0)
     {
       free (tgt);
       tgt = NULL;
@@ -2476,6 +2477,19 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
     }
 
   for (i = 0; i < mapnum; i++)
+    if ((kinds[i] & typemask) == GOMP_MAP_DETACH)
+      {
+	struct splay_tree_key_s cur_node;
+	cur_node.host_start = (uintptr_t) hostaddrs[i];
+	cur_node.host_end = cur_node.host_start + sizeof (void *);
+	splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
+
+	if (n)
+	  gomp_detach_pointer (devicep, NULL, n, (uintptr_t) hostaddrs[i],
+			       false, NULL);
+      }
+
+  for (i = 0; i < mapnum; i++)
     {
       struct splay_tree_key_s cur_node;
       unsigned char kind = kinds[i] & typemask;
@@ -2512,7 +2526,9 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
 				cur_node.host_end - cur_node.host_start);
 	  if (k->refcount == 0)
 	    gomp_remove_var (devicep, k);
+	  break;
 
+	case GOMP_MAP_DETACH:
 	  break;
 	default:
 	  gomp_mutex_unlock (&devicep->lock);
@@ -2621,6 +2637,14 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
 			 &kinds[i], true, GOMP_MAP_VARS_ENTER_DATA);
 	  i += j - i - 1;
 	}
+      else if (i + 1 < mapnum && (kinds[i + 1] & 0xff) == GOMP_MAP_ATTACH)
+	{
+	  /* An attach operation must be processed together with the mapped
+	     base-pointer list item.  */
+	  gomp_map_vars (devicep, 2, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
+			 true, GOMP_MAP_VARS_ENTER_DATA);
+	  i += 1;
+	}
       else
 	gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
 		       true, GOMP_MAP_VARS_ENTER_DATA);
diff --git a/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-1.c b/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-1.c
new file mode 100644
index 0000000..e7deec6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-1.c
@@ -0,0 +1,82 @@
+#include <stdlib.h>
+
+struct S
+{
+  int a, b;
+  int *ptr;
+  int c, d;
+};
+typedef struct S S;
+
+#pragma omp declare target
+int *gp;
+#pragma omp end declare target
+
+#define N 10
+int main (void)
+{
+  /* Test to see if pointer attachment works, for scalar pointers,
+     and pointer fields in structures.  */
+
+  int *ptr = (int *) malloc (sizeof (int) * N);
+  int *orig_ptr = ptr;
+
+  #pragma omp target map (ptr, ptr[:N])
+  {
+    for (int i = 0; i < N; i++)
+      ptr[i] = N - i;
+  }
+
+  if (ptr != orig_ptr)
+    abort ();
+
+  for (int i = 0; i < N; i++)
+    if (ptr[i] != N - i)
+      abort ();
+
+  S s = { 0 };
+  s.ptr = ptr;
+  #pragma omp target map (s, s.ptr[:N])
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i] = i;
+
+    s.a = 1;
+    s.b = 2;
+  }
+
+  if (s.ptr != ptr)
+    abort ();
+
+  for (int i = 0; i < N; i++)
+    if (s.ptr[i] != i)
+      abort ();
+
+  if (s.a != 1 || s.b != 2 || s.c != 0 || s.d != 0)
+    abort ();
+
+  gp = (int *) malloc (sizeof (int) * N);
+  orig_ptr = gp;
+
+  for (int i = 0; i < N; i++)
+    gp[i] = i - 1;
+
+  #pragma omp target map (gp[:N])
+  {
+    for (int i = 0; i < N; i++)
+      gp[i] += 1;
+  }
+
+  if (gp != orig_ptr)
+    abort ();
+
+  for (int i = 0; i < N; i++)
+    if (gp[i] != i)
+      abort ();
+
+  free (ptr);
+  free (gp);
+
+  return 0;
+}
+