libgomp: Structure element mapping for OpenMP 5.0

This patch implement OpenMP 5.0 requirements of incrementing/decrementing
the reference count of a mapped structure at most once (across all elements)
on a construct.

This is implemented by pulling in libgomp/hashtab.h and using htab_t as a
pointer set. Structure element list siblings also have pointers-to-refcounts
linked together, to naturally achieve uniform increment/decrement without
repeating.

There are still some questions on whether using such a htab_t based set is
faster/slower than using a sorted pointer array based implementation. This
is to be researched on later.

libgomp/ChangeLog:

	* hashtab.h (htab_clear): New function with initialization code
	factored out from...
	(htab_create): ...here, adjust to use htab_clear function.

	* libgomp.h (REFCOUNT_SPECIAL): New symbol to denote range of
	special refcount values, add comments.
	(REFCOUNT_INFINITY): Adjust definition to use REFCOUNT_SPECIAL.
	(REFCOUNT_LINK): Likewise.
	(REFCOUNT_STRUCTELEM): New special refcount range for structure
	element siblings.
	(REFCOUNT_STRUCTELEM_P): Macro for testing for structure element
	sibling maps.
	(REFCOUNT_STRUCTELEM_FLAG_FIRST): Flag to indicate first sibling.
	(REFCOUNT_STRUCTELEM_FLAG_LAST):  Flag to indicate last sibling.
	(REFCOUNT_STRUCTELEM_FIRST_P): Macro to test _FIRST flag.
	(REFCOUNT_STRUCTELEM_LAST_P): Macro to test _LAST flag.
	(struct splay_tree_key_s): Add structelem_refcount and
	structelem_refcount_ptr fields into a union with dynamic_refcount.
	Add comments.
	(gomp_map_vars): Delete declaration.
	(gomp_map_vars_async): Likewise.
	(gomp_unmap_vars): Likewise.
	(gomp_unmap_vars_async): Likewise.
	(goacc_map_vars): New declaration.
	(goacc_unmap_vars): Likewise.

	* oacc-mem.c (acc_map_data): Adjust to use goacc_map_vars.
	(goacc_enter_datum): Likewise.
	(goacc_enter_data_internal): Likewise.
	* oacc-parallel.c (GOACC_parallel_keyed): Adjust to use goacc_map_vars
	and goacc_unmap_vars.
	(GOACC_data_start): Adjust to use goacc_map_vars.
	(GOACC_data_end): Adjust to use goacc_unmap_vars.

	* target.c (hash_entry_type): New typedef.
	(htab_alloc): New function hook for hashtab.h.
	(htab_free): Likewise.
	(htab_hash): Likewise.
	(htab_eq): Likewise.
	(hashtab.h): Add file include.
	(gomp_increment_refcount): New function.
	(gomp_decrement_refcount): Likewise.
	(gomp_map_vars_existing): Add refcount_set parameter, adjust to use
	gomp_increment_refcount.
	(gomp_map_fields_existing): Add refcount_set parameter, adjust calls
	to gomp_map_vars_existing.

	(gomp_map_vars_internal): Add refcount_set parameter, add local openmp_p
	variable to guard OpenMP specific paths, adjust calls to
	gomp_map_vars_existing, add structure element sibling splay_tree_key
	sequence creation code, adjust Fortran map case to avoid increment
	under OpenMP.
	(gomp_map_vars): Adjust to static, add refcount_set parameter, manage
	local refcount_set if caller passed in NULL, adjust call to
	gomp_map_vars_internal.
	(gomp_map_vars_async): Adjust and rename into...
	(goacc_map_vars): ...this new function, adjust call to
	gomp_map_vars_internal.

	(gomp_remove_splay_tree_key): New function with code factored out from
	gomp_remove_var_internal.
	(gomp_remove_var_internal): Add code to handle removing multiple
	splay_tree_key sequence for structure elements, adjust code to use
	gomp_remove_splay_tree_key for splay-tree key removal.
	(gomp_unmap_vars_internal): Add refcount_set parameter, adjust to use
	gomp_decrement_refcount.
	(gomp_unmap_vars): Adjust to static, add refcount_set parameter, manage
	local refcount_set if caller passed in NULL, adjust call to
	gomp_unmap_vars_internal.
	(gomp_unmap_vars_async): Adjust and rename into...
	(goacc_unmap_vars): ...this new function, adjust call to
	gomp_unmap_vars_internal.
	(GOMP_target): Manage refcount_set and adjust calls to gomp_map_vars and
	gomp_unmap_vars.
	(GOMP_target_ext): Likewise.
	(gomp_target_data_fallback): Adjust call to gomp_map_vars.
	(GOMP_target_data): Likewise.
	(GOMP_target_data_ext): Likewise.
	(GOMP_target_end_data): Adjust call to gomp_unmap_vars.
	(gomp_exit_data): Add refcount_set parameter, adjust to use
	gomp_decrement_refcount, adjust to queue splay-tree keys for removal
	after main loop.
	(GOMP_target_enter_exit_data): Manage refcount_set and adjust calls to
	gomp_map_vars and gomp_exit_data.
	(gomp_target_task_fn): Likewise.

	* testsuite/libgomp.c-c++-common/refcount-1.c: New testcase.
	* testsuite/libgomp.c-c++-common/struct-elem-1.c: New testcase.
	* testsuite/libgomp.c-c++-common/struct-elem-2.c: New testcase.
	* testsuite/libgomp.c-c++-common/struct-elem-3.c: New testcase.
	* testsuite/libgomp.c-c++-common/struct-elem-4.c: New testcase.
	* testsuite/libgomp.c-c++-common/struct-elem-5.c: New testcase.

6 files changed, 282 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.c-c++-common/refcount-1.c b/libgomp/testsuite/libgomp.c-c++-common/refcount-1.c
new file mode 100644
index 0000000..5ccd908
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/refcount-1.c
@@ -0,0 +1,61 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int main (void)
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  unsigned int a = 0xcdcdcdcd;
+  #pragma omp target enter data map (to:a)
+
+  a = 0xabababab;
+  unsigned char *p = (unsigned char *) &a;
+  unsigned char *q = p + 2;
+
+  #pragma omp target enter data map (alloc:p[:1], q[:1])
+
+  if (d != id)
+    {
+      if (!omp_target_is_present (&a, d))
+	abort ();
+      if (!omp_target_is_present (&p[0], d))
+	abort ();
+      if (!omp_target_is_present (&q[0], d))
+	abort ();
+    }
+
+  #pragma omp target exit data map (release:a)
+
+  if (d != id)
+    {
+      if (!omp_target_is_present (&a, d))
+	abort ();
+      if (!omp_target_is_present (&p[0], d))
+	abort ();
+      if (!omp_target_is_present (&q[0], d))
+	abort ();
+    }
+
+  #pragma omp target exit data map (from:q[:1])
+
+  if (d != id)
+    {
+      if (omp_target_is_present (&a, d))
+	abort ();
+      if (omp_target_is_present (&p[0], d))
+	abort ();
+      if (omp_target_is_present (&q[0], d))
+	abort ();
+
+      if (q[0] != 0xcd)
+	abort ();
+      if (p[0] != 0xab)
+	abort ();
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c
new file mode 100644
index 0000000..5f40fd7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c
@@ -0,0 +1,29 @@
+#include <omp.h>
+#include <stdlib.h>
+
+struct S
+{
+  int a, b;
+};
+typedef struct S S;
+
+int main (void)
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  S s;
+  #pragma omp target enter data map (alloc: s.a, s.b)
+  #pragma omp target exit data map (release: s.b)
+
+  /* OpenMP 5.0 structure element mapping rules describe that elements of same
+     structure variable should allocate/deallocate in a uniform fashion, so
+     "s.a" should be removed together by above 'exit data'.  */
+  if (d != id && omp_target_is_present (&s.a, d))
+    abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-2.c b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-2.c
new file mode 100644
index 0000000..c50b299
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-2.c
@@ -0,0 +1,47 @@
+#include <omp.h>
+#include <stdlib.h>
+
+struct S
+{
+  int a, b, c, d;
+};
+typedef struct S S;
+
+int main (void)
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  S s;
+  #pragma omp target enter data map (alloc: s.a, s.b, s.c, s.d)
+  #pragma omp target enter data map (alloc: s.c)
+  #pragma omp target enter data map (alloc: s.b, s.d)
+  #pragma omp target enter data map (alloc: s.a, s.c, s.b)
+
+  #pragma omp target exit data map (release: s.a)
+  #pragma omp target exit data map (release: s.d)
+  #pragma omp target exit data map (release: s.c)
+  #pragma omp target exit data map (release: s.b)
+
+  /* OpenMP 5.0 structure element mapping rules describe that elements of same
+     structure variable should allocate/deallocate in a uniform fashion, so
+     all elements of 's' should be removed together by above 'exit data's.  */
+  if (d != id)
+    {
+      if (omp_target_is_present (&s, d))
+	abort ();
+      if (omp_target_is_present (&s.a, d))
+	abort ();
+      if (omp_target_is_present (&s.b, d))
+	abort ();
+      if (omp_target_is_present (&s.c, d))
+	abort ();
+      if (omp_target_is_present (&s.d, d))
+	abort ();
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-3.c b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-3.c
new file mode 100644
index 0000000..e2b6a6a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-3.c
@@ -0,0 +1,69 @@
+#include <omp.h>
+#include <stdlib.h>
+
+struct S
+{
+  int a, b, c, d;
+};
+typedef struct S S;
+
+int main (void)
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  S s;
+
+  #pragma omp target enter data map (alloc: s)
+  #pragma omp target enter data map (alloc: s)
+
+  #pragma omp target exit data map (release: s.a)
+  #pragma omp target exit data map (release: s.b)
+
+  /* OpenMP 5.0 structure element mapping rules describe that elements of same
+     structure variable should allocate/deallocate in a uniform fashion, so
+     all elements of 's' should be removed together by above 'exit data's.  */
+  if (d != id)
+    {
+      if (omp_target_is_present (&s, d))
+	abort ();
+      if (omp_target_is_present (&s.a, d))
+	abort ();
+      if (omp_target_is_present (&s.b, d))
+	abort ();
+      if (omp_target_is_present (&s.c, d))
+	abort ();
+      if (omp_target_is_present (&s.d, d))
+	abort ();
+    }
+
+  #pragma omp target enter data map (alloc: s.a, s.b)
+  #pragma omp target enter data map (alloc: s.a)
+  #pragma omp target enter data map (alloc: s.b)
+
+  #pragma omp target exit data map (release: s)
+  #pragma omp target exit data map (release: s)
+  #pragma omp target exit data map (release: s)
+
+  /* OpenMP 5.0 structure element mapping rules describe that elements of same
+     structure variable should allocate/deallocate in a uniform fashion, so
+     all elements of 's' should be removed together by above 'exit data's.  */
+  if (d != id)
+    {
+      if (omp_target_is_present (&s, d))
+	abort ();
+      if (omp_target_is_present (&s.a, d))
+	abort ();
+      if (omp_target_is_present (&s.b, d))
+	abort ();
+      if (omp_target_is_present (&s.c, d))
+	abort ();
+      if (omp_target_is_present (&s.d, d))
+	abort ();
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-4.c b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-4.c
new file mode 100644
index 0000000..9a23b4f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-4.c
@@ -0,0 +1,56 @@
+#include <omp.h>
+#include <stdlib.h>
+
+struct S
+{
+  int a, b, c, d, e;
+};
+typedef struct S S;
+
+int main (void)
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  S s = { 1, 2, 3, 4, 5 };
+  #pragma omp target enter data map (to:s)
+
+  int *p = &s.b;
+  int *q = &s.d;
+  #pragma omp target enter data map (alloc: p[:1], q[:1])
+
+  s.b = 88;
+  s.d = 99;
+
+  #pragma omp target exit data map (release: s)
+  if (d != id)
+    {
+      if (!omp_target_is_present (&s, d))
+	abort ();
+      if (!omp_target_is_present (&p[0], d))
+	abort ();
+      if (!omp_target_is_present (&q[0], d))
+	abort ();
+    }
+
+  #pragma omp target exit data map (from: q[:1])
+  if (d != id)
+    {
+      if (omp_target_is_present (&s, d))
+	abort ();
+      if (omp_target_is_present (&p[0], d))
+	abort ();
+      if (omp_target_is_present (&q[0], d))
+	abort ();
+
+      if (q[0] != 4)
+	abort ();
+      if (p[0] != 88)
+	abort ();
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-5.c b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-5.c
new file mode 100644
index 0000000..814c301
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-5.c
@@ -0,0 +1,20 @@
+/* { dg-do run } */
+
+struct S
+{
+  int a, b, c;
+};
+typedef struct S S;
+
+int main (void)
+{
+  S s;
+  #pragma omp target data map (alloc: s.a, s.c)
+  {
+    #pragma omp target enter data map (alloc: s.b)
+  }
+
+  return 0;
+}
+/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) structure element when other mapped elements from the same structure weren't mapped together with it" } */
+/* { dg-shouldfail "" } */