openmp: in_reduction clause support on target construct

This patch adds support for in_reduction clause on target construct, though
for now only for synchronous targets (without nowait clause).
The encountering thread in that case runs the target task and blocks until
the target region ends, so it is implemented by remapping it before entering
the target, initializing the private copy if not yet initialized for the
current thread and then using the remapped addresses for the mapping
addresses.
For nowait combined with in_reduction the patch contains a hack where the
nowait clause is ignored.  To implement it correctly, I think we would need
to create a new private variable for the in_reduction and initialize it before
doing the async target and adjust the map addresses to that private variable
and then pass a function pointer to the library routine with code where the callback
would remap the address to the current threads private variable and use in_reduction
combiner to combine the private variable we've created into the thread's copy.
The library would then need to make sure that the routine is called in some thread
participating in the parallel (and not in an unshackeled thread).

2021-06-24  Jakub Jelinek  <jakub@redhat.com>

gcc/
	* tree.h (OMP_CLAUSE_MAP_IN_REDUCTION): Document meaning for OpenMP.
	* gimplify.c (gimplify_scan_omp_clauses): For OpenMP map clauses
	with OMP_CLAUSE_MAP_IN_REDUCTION flag partially defer gimplification
	of non-decl OMP_CLAUSE_DECL.  For OMP_CLAUSE_IN_REDUCTION on
	OMP_TARGET user outer_ctx instead of ctx for placeholders and
	initializer/combiner gimplification.
	* omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_MAP_IN_REDUCTION
	on target constructs.
	(lower_rec_input_clauses): Likewise.
	(lower_omp_target): Likewise.
	* omp-expand.c (expand_omp_target): Temporarily ignore nowait clause
	on target if in_reduction is present.
gcc/c-family/
	* c-common.h (enum c_omp_region_type): Add C_ORT_TARGET and
	C_ORT_OMP_TARGET.
	* c-omp.c (c_omp_split_clauses): For OMP_CLAUSE_IN_REDUCTION on
	combined target constructs also add map (always, tofrom:) clause.
gcc/c/
	* c-parser.c (omp_split_clauses): Pass C_ORT_OMP_TARGET instead of
	C_ORT_OMP for clauses on target construct.
	(OMP_TARGET_CLAUSE_MASK): Add in_reduction clause.
	(c_parser_omp_target): For non-combined target add
	map (always, tofrom:) clauses for OMP_CLAUSE_IN_REDUCTION.  Pass
	C_ORT_OMP_TARGET to c_finish_omp_clauses.
	* c-typeck.c (handle_omp_array_sections): Adjust ort handling
	for addition of C_ORT_OMP_TARGET and simplify, mapping clauses are
	never present on C_ORT_*DECLARE_SIMD.
	(c_finish_omp_clauses): Likewise.  Handle OMP_CLAUSE_IN_REDUCTION
	on C_ORT_OMP_TARGET, set OMP_CLAUSE_MAP_IN_REDUCTION on
	corresponding map clauses.
gcc/cp/
	* parser.c (cp_omp_split_clauses): Pass C_ORT_OMP_TARGET instead of
	C_ORT_OMP for clauses on target construct.
	(OMP_TARGET_CLAUSE_MASK): Add in_reduction clause.
	(cp_parser_omp_target): For non-combined target add
	map (always, tofrom:) clauses for OMP_CLAUSE_IN_REDUCTION.  Pass
	C_ORT_OMP_TARGET to finish_omp_clauses.
	* semantics.c (handle_omp_array_sections_1): Adjust ort handling
	for addition of C_ORT_OMP_TARGET and simplify, mapping clauses are
	never present on C_ORT_*DECLARE_SIMD.
	(handle_omp_array_sections): Likewise.
	(finish_omp_clauses): Likewise.  Handle OMP_CLAUSE_IN_REDUCTION
	on C_ORT_OMP_TARGET, set OMP_CLAUSE_MAP_IN_REDUCTION on
	corresponding map clauses.
	* pt.c (tsubst_expr): Pass C_ORT_OMP_TARGET instead of C_ORT_OMP for
	clauses on target construct.
gcc/testsuite/
	* c-c++-common/gomp/target-in-reduction-1.c: New test.
	* c-c++-common/gomp/clauses-1.c: Add in_reduction clauses on
	target or combined target constructs.
libgomp/
	* testsuite/libgomp.c-c++-common/target-in-reduction-1.c: New test.
	* testsuite/libgomp.c-c++-common/target-in-reduction-2.c: New test.
	* testsuite/libgomp.c++/target-in-reduction-1.C: New test.
	* testsuite/libgomp.c++/target-in-reduction-2.C: New test.

3 files changed, 90 insertions, 48 deletions

diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index d57ddc4..b7a4298 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40877,7 +40877,9 @@ cp_omp_split_clauses (location_t loc, enum tree_code code,
   c_omp_split_clauses (loc, code, mask, clauses, cclauses);
   for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++)
     if (cclauses[i])
-      cclauses[i] = finish_omp_clauses (cclauses[i], C_ORT_OMP);
+      cclauses[i] = finish_omp_clauses (cclauses[i],
+					i == C_OMP_CLAUSE_SPLIT_TARGET
+					? C_ORT_OMP_TARGET : C_ORT_OMP);
 }
 
 /* OpenMP 5.0:
@@ -42219,6 +42221,7 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ALLOCATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR))
 
 static bool
@@ -42381,7 +42384,18 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
 
   OMP_TARGET_CLAUSES (stmt)
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
-				 "#pragma omp target", pragma_tok);
+				 "#pragma omp target", pragma_tok, false);
+  for (tree c = OMP_TARGET_CLAUSES (stmt); c; c = OMP_CLAUSE_CHAIN (c))
+    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION)
+      {
+	tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+	OMP_CLAUSE_DECL (nc) = OMP_CLAUSE_DECL (c);
+	OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_ALWAYS_TOFROM);
+	OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
+	OMP_CLAUSE_CHAIN (c) = nc;
+      }
+  OMP_TARGET_CLAUSES (stmt)
+    = finish_omp_clauses (OMP_TARGET_CLAUSES (stmt), C_ORT_OMP_TARGET);
   c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
 
   pc = &OMP_TARGET_CLAUSES (stmt);
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 5c55507..1af8120 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -18886,9 +18886,12 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
     case OACC_DATA:
     case OMP_TARGET_DATA:
     case OMP_TARGET:
-      tmp = tsubst_omp_clauses (OMP_CLAUSES (t), (TREE_CODE (t) == OACC_DATA)
-				? C_ORT_ACC : C_ORT_OMP, args, complain,
-				in_decl);
+      tmp = tsubst_omp_clauses (OMP_CLAUSES (t),
+				TREE_CODE (t) == OACC_DATA
+				? C_ORT_ACC
+				: TREE_CODE (t) == OMP_TARGET
+				? C_ORT_OMP_TARGET : C_ORT_OMP,
+				args, complain, in_decl);
       keep_next_level (true);
       stmt = begin_omp_structured_block ();
 
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 384c54b..fbaabf6 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5042,7 +5042,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		      omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	  return error_mark_node;
 	}
-      else if (ort == C_ORT_OMP
+      else if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
 	       && TREE_CODE (t) == PARM_DECL
 	       && DECL_ARTIFICIAL (t)
 	       && DECL_NAME (t) == this_identifier
@@ -5069,7 +5069,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
       return ret;
     }
 
-  if (ort == C_ORT_OMP
+  if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
       && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
 	  || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION
 	  || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TASK_REDUCTION)
@@ -5571,33 +5571,30 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	      || (TREE_CODE (t) == COMPONENT_REF
 		  && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
 	    return false;
-	  if (ort == C_ORT_OMP || ort == C_ORT_ACC)
-	    switch (OMP_CLAUSE_MAP_KIND (c))
-	      {
-	      case GOMP_MAP_ALLOC:
-	      case GOMP_MAP_IF_PRESENT:
-	      case GOMP_MAP_TO:
-	      case GOMP_MAP_FROM:
-	      case GOMP_MAP_TOFROM:
-	      case GOMP_MAP_ALWAYS_TO:
-	      case GOMP_MAP_ALWAYS_FROM:
-	      case GOMP_MAP_ALWAYS_TOFROM:
-	      case GOMP_MAP_RELEASE:
-	      case GOMP_MAP_DELETE:
-	      case GOMP_MAP_FORCE_TO:
-	      case GOMP_MAP_FORCE_FROM:
-	      case GOMP_MAP_FORCE_TOFROM:
-	      case GOMP_MAP_FORCE_PRESENT:
-		OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
-		break;
-	      default:
-		break;
-	      }
+	  switch (OMP_CLAUSE_MAP_KIND (c))
+	    {
+	    case GOMP_MAP_ALLOC:
+	    case GOMP_MAP_IF_PRESENT:
+	    case GOMP_MAP_TO:
+	    case GOMP_MAP_FROM:
+	    case GOMP_MAP_TOFROM:
+	    case GOMP_MAP_ALWAYS_TO:
+	    case GOMP_MAP_ALWAYS_FROM:
+	    case GOMP_MAP_ALWAYS_TOFROM:
+	    case GOMP_MAP_RELEASE:
+	    case GOMP_MAP_DELETE:
+	    case GOMP_MAP_FORCE_TO:
+	    case GOMP_MAP_FORCE_FROM:
+	    case GOMP_MAP_FORCE_TOFROM:
+	    case GOMP_MAP_FORCE_PRESENT:
+	      OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+	      break;
+	    default:
+	      break;
+	    }
 	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 				      OMP_CLAUSE_MAP);
-	  if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC)
-	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
-	  else if (TREE_CODE (t) == COMPONENT_REF)
+	  if (TREE_CODE (t) == COMPONENT_REF)
 	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
 	  else if (REFERENCE_REF_P (t)
 		   && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
@@ -6592,6 +6589,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   tree detach_seen = NULL_TREE;
   bool mergeable_seen = false;
   bool implicit_moved = false;
+  bool target_in_reduction_seen = false;
 
   bitmap_obstack_initialize (NULL);
   bitmap_initialize (&generic_head, &bitmap_default_obstack);
@@ -6603,7 +6601,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   bitmap_initialize (&map_field_head, &bitmap_default_obstack);
   bitmap_initialize (&map_firstprivate_head, &bitmap_default_obstack);
   /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
-     instead.  */
+     instead and for ort == C_ORT_OMP_TARGET used as in_reduction_head.  */
   bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
 
   if (ort & C_ORT_ACC)
@@ -6866,8 +6864,22 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		   || (ort == C_ORT_OMP
 		       && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
 			   || (OMP_CLAUSE_CODE (c)
-			       == OMP_CLAUSE_USE_DEVICE_ADDR))))
+			       == OMP_CLAUSE_USE_DEVICE_ADDR)))
+		   || (ort == C_ORT_OMP_TARGET
+		       && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION))
 	    {
+	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION
+		  && (bitmap_bit_p (&generic_head, DECL_UID (t))
+		      || bitmap_bit_p (&firstprivate_head, DECL_UID (t))))
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "%qD appears more than once in data-sharing "
+			    "clauses", t);
+		  remove = true;
+		  break;
+		}
+	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION)
+		target_in_reduction_seen = true;
 	      if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
 		{
 		  error_at (OMP_CLAUSE_LOCATION (c),
@@ -6882,7 +6894,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    }
 	  else if (bitmap_bit_p (&generic_head, DECL_UID (t))
 		   || bitmap_bit_p (&firstprivate_head, DECL_UID (t))
-		   || bitmap_bit_p (&lastprivate_head, DECL_UID (t)))
+		   || bitmap_bit_p (&lastprivate_head, DECL_UID (t))
+		   || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t)))
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
 			"%qD appears more than once in data clauses", t);
@@ -6982,7 +6995,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		   && bitmap_bit_p (&map_firstprivate_head, DECL_UID (t)))
 	    remove = true;
 	  else if (bitmap_bit_p (&generic_head, DECL_UID (t))
-		   || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+		   || bitmap_bit_p (&firstprivate_head, DECL_UID (t))
+		   || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t)))
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
 			"%qD appears more than once in data clauses", t);
@@ -7795,13 +7809,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      t = TREE_OPERAND (t, 0);
 	      OMP_CLAUSE_DECL (c) = t;
 	    }
-	  if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
-	      && TREE_CODE (t) == COMPONENT_REF
+	  if (TREE_CODE (t) == COMPONENT_REF
 	      && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
 	    t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
 	  if (TREE_CODE (t) == COMPONENT_REF
-	      && ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
-		  || ort == C_ORT_ACC)
 	      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
 	    {
 	      if (type_dependent_expression_p (t))
@@ -7842,7 +7853,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
 		{
 		  if (bitmap_bit_p (&map_field_head, DECL_UID (t))
-		      || (ort == C_ORT_OMP
+		      || (ort != C_ORT_ACC
 			  && bitmap_bit_p (&map_head, DECL_UID (t))))
 		    goto handle_map_references;
 		}
@@ -7924,7 +7935,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		   && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
 	    {
 	      if (bitmap_bit_p (&generic_head, DECL_UID (t))
-		  || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+		  || bitmap_bit_p (&firstprivate_head, DECL_UID (t))
+		  || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t)))
 		{
 		  error_at (OMP_CLAUSE_LOCATION (c),
 			    "%qD appears more than once in data clauses", t);
@@ -7941,10 +7953,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  remove = true;
 		}
 	      else
-		{
-		  bitmap_set_bit (&generic_head, DECL_UID (t));
-		  bitmap_set_bit (&map_firstprivate_head, DECL_UID (t));
-		}
+		bitmap_set_bit (&map_firstprivate_head, DECL_UID (t));
 	    }
 	  else if (bitmap_bit_p (&map_head, DECL_UID (t))
 		   && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
@@ -7960,8 +7969,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			  "%qD appears more than once in map clauses", t);
 	      remove = true;
 	    }
-	  else if (bitmap_bit_p (&generic_head, DECL_UID (t))
-		   && ort == C_ORT_ACC)
+	  else if (ort == C_ORT_ACC
+		   && bitmap_bit_p (&generic_head, DECL_UID (t)))
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
 			"%qD appears more than once in data clauses", t);
@@ -8511,6 +8520,22 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    }
 	  pc = &OMP_CLAUSE_CHAIN (c);
 	  continue;
+	case OMP_CLAUSE_MAP:
+	  if (target_in_reduction_seen && !processing_template_decl)
+	    {
+	      t = OMP_CLAUSE_DECL (c);
+	      while (handled_component_p (t)
+		     || TREE_CODE (t) == INDIRECT_REF
+		     || TREE_CODE (t) == ADDR_EXPR
+		     || TREE_CODE (t) == MEM_REF
+		     || TREE_CODE (t) == NON_LVALUE_EXPR)
+		t = TREE_OPERAND (t, 0);
+	      if (DECL_P (t)
+		  && bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
+		OMP_CLAUSE_MAP_IN_REDUCTION (c) = 1;
+	    }
+	  pc = &OMP_CLAUSE_CHAIN (c);
+	  continue;
 	case OMP_CLAUSE_NOWAIT:
 	  if (copyprivate_seen)
 	    {