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.

1 files changed, 82 insertions, 0 deletions

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;
+}
+