OpenMP: Expand "declare mapper" mappers for target {enter,exit,} data directives

This patch allows 'declare mapper' mappers to be used on 'omp target
data', 'omp target enter data' and 'omp target exit data' directives.
For each of these, only explicit mappings are supported, unlike for
'omp target' directives where implicit uses of variables inside an
offload region might trigger mappers also.

Each of C, C++ and Fortran are supported.

The patch also adjusts 'map kind decay' to match OpenMP 5.2 semantics,
which is particularly important with regard to 'exit data' operations.

2023-07-06  Julian Brown  <julian@codesourcery.com>

gcc/c-family/
	* c-common.h (c_omp_region_type): Add C_ORT_EXIT_DATA,
	C_ORT_OMP_EXIT_DATA.
	(c_omp_instantiate_mappers): Add region type parameter.
	* c-omp.cc (omp_split_map_kind, omp_join_map_kind,
	omp_map_decayed_kind): New functions.
	(omp_instantiate_mapper): Add ORT parameter.  Implement map kind decay
	for instantiated mapper clauses.
	(c_omp_instantiate_mappers): Add ORT parameter, pass to
	omp_instantiate_mapper.

gcc/c/
	* c-parser.cc (c_parser_omp_target_data): Instantiate mappers for
	'omp target data'.
	(c_parser_omp_target_enter_data): Instantiate mappers for 'omp target
	enter data'.
	(c_parser_omp_target_exit_data): Instantiate mappers for 'omp target
	exit data'.
	(c_parser_omp_target): Add c_omp_region_type argument to
	c_omp_instantiate_mappers call.
	* c-tree.h (c_omp_instantiate_mappers): Remove spurious prototype.

gcc/cp/
	* parser.cc (cp_parser_omp_target_data): Instantiate mappers for 'omp
	target data'.
	(cp_parser_omp_target_enter_data): Instantiate mappers for 'omp target
	enter data'.
	(cp_parser_omp_target_exit_data): Instantiate mappers for 'omp target
	exit data'.
	(cp_parser_omp_target): Add c_omp_region_type argument to
	c_omp_instantiate_mappers call.
	* pt.cc (tsubst_omp_clauses): Instantiate mappers for OMP regions other
	than just C_ORT_OMP_TARGET.
	(tsubst_expr): Update call to tsubst_omp_clauses for OMP_TARGET_UPDATE,
	OMP_TARGET_ENTER_DATA, OMP_TARGET_EXIT_DATA stanza.
	* semantics.cc (cxx_omp_map_array_section): Avoid calling
	build_array_ref for non-array/non-pointer bases (error reported
	already).

gcc/fortran/
	* trans-openmp.cc (omp_split_map_op, omp_join_map_op,
	omp_map_decayed_kind): New functions.
	(gfc_trans_omp_instantiate_mapper): Add CD parameter.  Implement map
	kind decay.
	(gfc_trans_omp_instantiate_mappers): Add CD parameter.  Pass to above
	function.
	(gfc_trans_omp_target_data): Instantiate mappers for 'omp target data'.
	(gfc_trans_omp_target_enter_data): Instantiate mappers for 'omp target
	enter data'.
	(gfc_trans_omp_target_exit_data): Instantiate mappers for 'omp target
	exit data'.

gcc/testsuite/
	* c-c++-common/gomp/declare-mapper-15.c: New test.
	* c-c++-common/gomp/declare-mapper-16.c: New test.
	* g++.dg/gomp/declare-mapper-1.C: Adjust expected scan output.
	* gfortran.dg/gomp/declare-mapper-22.f90: New test.
	* gfortran.dg/gomp/declare-mapper-23.f90: New test.

1 files changed, 8 insertions, 3 deletions

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index ee2d69e..78a2f18 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -26199,7 +26199,9 @@ 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");
+				"#pragma omp target data", false);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);
+  clauses = c_finish_omp_clauses (clauses, C_ORT_OMP);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
@@ -26357,7 +26359,9 @@ 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");
+				"#pragma omp target enter data", false);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);
+  clauses = c_finish_omp_clauses (clauses, C_ORT_OMP);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
@@ -26468,6 +26472,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", false);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_EXIT_DATA);
   clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
@@ -26725,7 +26730,7 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
 	OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
 	OMP_CLAUSE_CHAIN (c) = nc;
       }
-  clauses = c_omp_instantiate_mappers (clauses);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_TARGET);
   clauses  = c_finish_omp_clauses (clauses, C_ORT_OMP_TARGET);
   c_omp_adjust_map_clauses (clauses, true);