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, 9 insertions, 1 deletions

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index ecc3d21..377914c 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -19511,6 +19511,7 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
 				"#pragma omp target data");
+  c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
     {
@@ -19528,6 +19529,7 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
 	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
 	    break;
 	  default:
 	    map_seen |= 1;
@@ -19651,6 +19653,7 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
 				"#pragma omp target enter data");
+  c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
     {
@@ -19664,6 +19667,7 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
 	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
 	    break;
 	  default:
 	    map_seen |= 1;
@@ -19735,7 +19739,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
 				"#pragma omp target exit data");
-
+  c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
     {
@@ -19750,6 +19754,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
 	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
 	    break;
 	  default:
 	    map_seen |= 1;
@@ -19960,6 +19965,8 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
   OMP_TARGET_CLAUSES (stmt)
     = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
 				"#pragma omp target");
+  c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
+
   pc = &OMP_TARGET_CLAUSES (stmt);
   keep_next_level ();
   block = c_begin_compound_stmt (true);
@@ -19984,6 +19991,7 @@ check_clauses:
 	  case GOMP_MAP_ALLOC:
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
 	    break;
 	  default:
 	    error_at (OMP_CLAUSE_LOCATION (*pc),