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.

2 files changed, 23 insertions, 9 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),
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 9684037..df1dad4 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13584,11 +13584,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
       if (ort != C_ORT_OMP && ort != C_ORT_ACC)
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
       else if (TREE_CODE (t) == COMPONENT_REF)
-	{
-	  gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
-					       : GOMP_MAP_ALWAYS_POINTER;
-	  OMP_CLAUSE_SET_MAP_KIND (c2, k);
-	}
+	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
       else
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
       if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
@@ -14711,7 +14707,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		break;
 	      if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
 		{
-		  if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+		  if (bitmap_bit_p (&map_field_head, DECL_UID (t))
+		      || (ort == C_ORT_OMP
+			  && bitmap_bit_p (&map_head, DECL_UID (t))))
 		    break;
 		}
 	    }
@@ -14780,7 +14778,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      else
 		bitmap_set_bit (&generic_head, DECL_UID (t));
 	    }
-	  else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+	  else if (bitmap_bit_p (&map_head, DECL_UID (t))
+		   && (ort != C_ORT_OMP
+		       || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
 	    {
 	      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 		error_at (OMP_CLAUSE_LOCATION (c),
@@ -14794,7 +14794,13 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      remove = true;
 	    }
 	  else if (bitmap_bit_p (&generic_head, DECL_UID (t))
-		   || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+		   && ort == C_ORT_ACC)
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"%qD appears more than once in data clauses", t);
+	      remove = true;
+	    }
+	  else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
 	    {
 	      if (ort == C_ORT_ACC)
 		error_at (OMP_CLAUSE_LOCATION (c),