diff mbox series

[v6,03/11] OpenMP/OpenACC: Refine condition for when map clause expansion happens

Message ID 44822b56c7eda01cea993e05c19aa7e881966b5a.1671796515.git.julian@codesourcery.com
State New
Headers show
Series OpenMP: C/C++ lvalue parsing, C/C++/Fortran "declare mapper" support | expand

Commit Message

Julian Brown Dec. 23, 2022, 12:12 p.m. UTC
This patch fixes some cases for OpenACC and OpenMP where map clauses were
being expanded (adding firstprivate_pointer, attach/detach nodes, and so
forth) unnecessarily, after the "OpenMP/OpenACC: Rework clause expansion
and nested struct handling" patch (approved but not yet committed):

  https://gcc.gnu.org/pipermail/gcc-patches/2022-October/603792.html

This is done by introducing a C_ORT_ACC_TARGET region type for OpenACC
compute regions to help distinguish them from non-compute regions that
need different handling, and by passing the region type through to the
clause expansion functions.

The patch also fixes clause expansion for OpenMP TO/FROM clauses, which
need to dereference references but not have any additional mapping nodes.

(These cases showed up due to the gimplification changes in the C++
"declare mapper" patch, but logically belong next to the earlier patch
named above.)

2022-11-30  Julian Brown  <julian@codesourcery.com>

gcc/
	* c-family/c-common.h (c_omp_region_type): Add C_ORT_ACC_TARGET.
	(c_omp_address_inspector): Pass c_omp_region_type instead of "target"
	bool.
	* c-family/c-omp.cc (c_omp_address_inspector::expand_array_base):
	Adjust clause expansion for OpenACC and non-map (OpenMP to/from)
	clauses.
	(c_omp_address_inspector::expand_component_selector): Use
	c_omp_region_type parameter.  Don't expand OpenMP to/from clauses.
	(c_omp_address_inspector::expand_map_clause): Take ORT parameter, pass
	to expand_array_base, etc.

gcc/c/
	* c-parser.cc (c_parser_oacc_all_clauses): Add TARGET parameter. Use
	to select region type for c_finish_omp_clauses call.
	(c_parser_oacc_loop): Update calls to c_parser_oacc_all_clauses.
	(c_parser_oacc_compute): Likewise.
	* c-typeck.cc (handle_omp_array_sctions_1): Update for C_ORT_ACC_TARGET
	addition and ai.expand_map_clause signature change.
	(c_finish_omp_clauses): Likewise.

gcc/cp/
	* parser.cc (cp_parser_oacc_all_clauses): Add TARGET parameter. Use
	to select region type for finish_omp_clauses call.
	(cp_parser_oacc_declare): Update call to cp_parser_oacc_all_clauses.
	(cp_parser_oacc_loop): Update calls to cp_parser_oacc_all_clauses.
	(cp_parser_oacc_compute): Likewise.
	* pt.cc (tsubst_expr): Use C_ORT_ACC_TARGET for call to
	tsubst_omp_clauses for compute regions.
	* semantics.cc (handle_omp_array_sections_1): Update for
	C_ORT_ACC_TARGET addition and ai.expand_map_clause signature change.
	(finish_omp_clauses): Likewise.
---
 gcc/c-family/c-common.h | 10 +++--
 gcc/c-family/c-omp.cc   | 90 ++++++++++++++++++++++++++++++++++++-----
 gcc/c/c-parser.cc       | 15 ++++---
 gcc/c/c-typeck.cc       | 39 ++++++++----------
 gcc/cp/parser.cc        | 15 ++++---
 gcc/cp/pt.cc            |  4 +-
 gcc/cp/semantics.cc     | 47 ++++++++++-----------
 7 files changed, 144 insertions(+), 76 deletions(-)
diff mbox series

Patch

diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index d935d4b3d7d9..06674e769bd4 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1245,7 +1245,8 @@  enum c_omp_region_type
   C_ORT_DECLARE_SIMD		= 1 << 2,
   C_ORT_TARGET			= 1 << 3,
   C_ORT_OMP_DECLARE_SIMD	= C_ORT_OMP | C_ORT_DECLARE_SIMD,
-  C_ORT_OMP_TARGET		= C_ORT_OMP | C_ORT_TARGET
+  C_ORT_OMP_TARGET		= C_ORT_OMP | C_ORT_TARGET,
+  C_ORT_ACC_TARGET		= C_ORT_ACC | C_ORT_TARGET
 };
 
 extern tree c_finish_omp_master (location_t, tree);
@@ -1345,10 +1346,11 @@  public:
   bool maybe_zero_length_array_section (tree);
 
   tree expand_array_base (tree, vec<omp_addr_token *> &, tree, unsigned *,
-			  bool, bool);
+			  c_omp_region_type, bool);
   tree expand_component_selector (tree, vec<omp_addr_token *> &, tree,
-				  unsigned *, bool);
-  tree expand_map_clause (tree, tree, vec<omp_addr_token *> &, bool);
+				  unsigned *, c_omp_region_type);
+  tree expand_map_clause (tree, tree, vec<omp_addr_token *> &,
+			  c_omp_region_type);
 };
 
 enum c_omp_directive_kind {
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index d32c2a977304..74c01d8f2a52 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -3370,7 +3370,8 @@  tree
 c_omp_address_inspector::expand_array_base (tree c,
 					    vec<omp_addr_token *> &addr_tokens,
 					    tree expr, unsigned *idx,
-					    bool target, bool decl_p)
+					    c_omp_region_type ort,
+					    bool decl_p)
 {
   using namespace omp_addr_tokenizer;
   location_t loc = OMP_CLAUSE_LOCATION (c);
@@ -3380,14 +3381,26 @@  c_omp_address_inspector::expand_array_base (tree c,
 			   && is_global_var (decl)
 			   && lookup_attribute ("omp declare target",
 						DECL_ATTRIBUTES (decl)));
+  bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP;
   bool implicit_p = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		     && OMP_CLAUSE_MAP_IMPLICIT (c));
   bool chain_p = omp_access_chain_p (addr_tokens, i + 1);
   tree c2 = NULL_TREE, c3 = NULL_TREE;
   unsigned consume_tokens = 2;
+  bool target = (ort & C_ORT_TARGET) != 0;
+  bool openmp = (ort & C_ORT_OMP) != 0;
 
   gcc_assert (i == 0);
 
+  if (!openmp
+      && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+	  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+    {
+      *idx = ++i;
+      return c;
+    }
+
   switch (addr_tokens[i + 1]->u.access_kind)
     {
     case ACCESS_DIRECT:
@@ -3397,11 +3410,19 @@  c_omp_address_inspector::expand_array_base (tree c,
 
     case ACCESS_REF:
       {
-	/* Copy the referenced object.  */
+	/* Copy the referenced object.  Note that we do this even for !MAP_P
+	   clauses.  */
 	tree obj = convert_from_reference (addr_tokens[i + 1]->expr);
 	OMP_CLAUSE_DECL (c) = obj;
 	OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (obj));
 
+	if (!map_p)
+	  {
+	    if (decl_p)
+	      c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
+	    break;
+	  }
+
 	/* If we have a reference to a pointer, avoid using
 	   FIRSTPRIVATE_REFERENCE here in case the pointer is modified in the
 	   offload region (we can only do that if the pointer does not point
@@ -3441,6 +3462,13 @@  c_omp_address_inspector::expand_array_base (tree c,
 
     case ACCESS_INDEXED_REF_TO_ARRAY:
       {
+	if (!map_p)
+	  {
+	    if (decl_p)
+	      c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
+	    break;
+	  }
+
 	tree virtual_origin
 	  = convert_from_reference (addr_tokens[i + 1]->expr);
 	virtual_origin = build_fold_addr_expr (virtual_origin);
@@ -3467,6 +3495,13 @@  c_omp_address_inspector::expand_array_base (tree c,
 
     case ACCESS_INDEXED_ARRAY:
       {
+	if (!map_p)
+	  {
+	    if (decl_p)
+	      c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
+	    break;
+	  }
+
 	/* The code handling "firstprivatize_array_bases" in gimplify.cc is
 	   relevant here.  What do we need to create for arrays at this
 	   stage?  (This condition doesn't feel quite right.  FIXME?)  */
@@ -3501,6 +3536,13 @@  c_omp_address_inspector::expand_array_base (tree c,
     case ACCESS_POINTER:
     case ACCESS_POINTER_OFFSET:
       {
+	if (!map_p)
+	  {
+	    if (decl_p)
+	      c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
+	    break;
+	  }
+
 	unsigned last_access = i + 1;
 	tree virtual_origin;
 
@@ -3524,7 +3566,11 @@  c_omp_address_inspector::expand_array_base (tree c,
 					     addr_tokens[last_access]->expr);
 	tree data_addr = omp_accessed_addr (addr_tokens, last_access, expr);
 	c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
-	if (decl_p && target && !chain_p && !declare_target_p)
+	/* For OpenACC, use FIRSTPRIVATE_POINTER for decls even on non-compute
+	   regions (e.g. "acc data" constructs).  It'll be removed anyway in
+	   gimplify.cc, but doing it this way maintains diagnostic
+	   behaviour.  */
+	if (decl_p && (target || !openmp) && !chain_p && !declare_target_p)
 	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
 	else
 	  {
@@ -3544,6 +3590,13 @@  c_omp_address_inspector::expand_array_base (tree c,
     case ACCESS_REF_TO_POINTER:
     case ACCESS_REF_TO_POINTER_OFFSET:
       {
+	if (!map_p)
+	  {
+	    if (decl_p)
+	      c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
+	    break;
+	  }
+
 	unsigned last_access = i + 1;
 	tree virtual_origin;
 
@@ -3618,7 +3671,7 @@  c_omp_address_inspector::expand_array_base (tree c,
   i += consume_tokens;
   *idx = i;
 
-  if (target && chain_p)
+  if (target && chain_p && map_p)
     return omp_expand_access_chain (c, expr, addr_tokens, idx);
   else if (chain_p)
     while (*idx < addr_tokens.length ()
@@ -3635,13 +3688,15 @@  c_omp_address_inspector::expand_component_selector (tree c,
 						    vec<omp_addr_token *>
 						      &addr_tokens,
 						    tree expr, unsigned *idx,
-						    bool target)
+						    c_omp_region_type ort)
 {
   using namespace omp_addr_tokenizer;
   location_t loc = OMP_CLAUSE_LOCATION (c);
   unsigned i = *idx;
   tree c2 = NULL_TREE, c3 = NULL_TREE;
   bool chain_p = omp_access_chain_p (addr_tokens, i + 1);
+  bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP;
+  bool target = (ort & C_ORT_TARGET) != 0;
 
   switch (addr_tokens[i + 1]->u.access_kind)
     {
@@ -3651,11 +3706,15 @@  c_omp_address_inspector::expand_component_selector (tree c,
 
     case ACCESS_REF:
       {
-	/* Copy the referenced object.  */
+	/* Copy the referenced object.  Note that we also do this for !MAP_P
+	   clauses.  */
 	tree obj = convert_from_reference (addr_tokens[i + 1]->expr);
 	OMP_CLAUSE_DECL (c) = obj;
 	OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (obj));
 
+	if (!map_p)
+	  break;
+
 	c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
 	OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
@@ -3665,6 +3724,9 @@  c_omp_address_inspector::expand_component_selector (tree c,
 
     case ACCESS_INDEXED_REF_TO_ARRAY:
       {
+	if (!map_p)
+	  break;
+
 	tree virtual_origin
 	  = convert_from_reference (addr_tokens[i + 1]->expr);
 	virtual_origin = build_fold_addr_expr (virtual_origin);
@@ -3686,6 +3748,9 @@  c_omp_address_inspector::expand_component_selector (tree c,
     case ACCESS_POINTER:
     case ACCESS_POINTER_OFFSET:
       {
+	if (!map_p)
+	  break;
+
 	tree virtual_origin
 	  = fold_convert_loc (loc, ptrdiff_type_node,
 			      addr_tokens[i + 1]->expr);
@@ -3705,6 +3770,9 @@  c_omp_address_inspector::expand_component_selector (tree c,
     case ACCESS_REF_TO_POINTER:
     case ACCESS_REF_TO_POINTER_OFFSET:
       {
+	if (!map_p)
+	  break;
+
 	tree ptr = convert_from_reference (addr_tokens[i + 1]->expr);
 	tree virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
 						ptr);
@@ -3750,7 +3818,7 @@  c_omp_address_inspector::expand_component_selector (tree c,
   i += 2;
   *idx = i;
 
-  if (target && chain_p)
+  if (target && chain_p && map_p)
     return omp_expand_access_chain (c, expr, addr_tokens, idx);
   else if (chain_p)
     while (*idx < addr_tokens.length ()
@@ -3766,7 +3834,7 @@  c_omp_address_inspector::expand_component_selector (tree c,
 tree
 c_omp_address_inspector::expand_map_clause (tree c, tree expr,
 					    vec<omp_addr_token *> &addr_tokens,
-					    bool target)
+					    c_omp_region_type ort)
 {
   using namespace omp_addr_tokenizer;
   unsigned i, length = addr_tokens.length ();
@@ -3780,7 +3848,7 @@  c_omp_address_inspector::expand_map_clause (tree c, tree expr,
 	  && addr_tokens[i]->u.structure_base_kind == BASE_DECL
 	  && addr_tokens[i + 1]->type == ACCESS_METHOD)
 	{
-	  c = expand_array_base (c, addr_tokens, expr, &i, target, true);
+	  c = expand_array_base (c, addr_tokens, expr, &i, ort, true);
 	  if (c == error_mark_node)
 	    return error_mark_node;
 	}
@@ -3789,7 +3857,7 @@  c_omp_address_inspector::expand_map_clause (tree c, tree expr,
 	       && addr_tokens[i]->u.structure_base_kind == BASE_ARBITRARY_EXPR
 	       && addr_tokens[i + 1]->type == ACCESS_METHOD)
 	{
-	  c = expand_array_base (c, addr_tokens, expr, &i, target, false);
+	  c = expand_array_base (c, addr_tokens, expr, &i, ort, false);
 	  if (c == error_mark_node)
 	    return error_mark_node;
 	}
@@ -3825,7 +3893,7 @@  c_omp_address_inspector::expand_map_clause (tree c, tree expr,
 	       && addr_tokens[i]->type == COMPONENT_SELECTOR
 	       && addr_tokens[i + 1]->type == ACCESS_METHOD)
 	{
-	  c = expand_component_selector (c, addr_tokens, expr, &i, target);
+	  c = expand_component_selector (c, addr_tokens, expr, &i, ort);
 	  /* We used 'expr', so these must have been the last tokens.  */
 	  gcc_assert (i == length);
 	  if (c == error_mark_node)
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 7d6960fffbb6..15e2fd7962af 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -17514,7 +17514,8 @@  c_parser_omp_clause_detach (c_parser *parser, tree list)
 
 static tree
 c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
-			   const char *where, bool finish_p = true)
+			   const char *where, bool finish_p = true,
+			   bool target = false)
 {
   tree clauses = NULL;
   bool first = true;
@@ -17715,7 +17716,8 @@  c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
   c_parser_skip_to_pragma_eol (parser);
 
   if (finish_p)
-    return c_finish_omp_clauses (clauses, C_ORT_ACC);
+    return c_finish_omp_clauses (clauses, target ? C_ORT_ACC_TARGET
+						 : C_ORT_ACC);
 
   return clauses;
 }
@@ -18440,12 +18442,13 @@  c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
   mask |= OACC_LOOP_CLAUSE_MASK;
 
   tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name,
-					    cclauses == NULL);
+					    /*finish_p=*/cclauses == NULL,
+					    /*target=*/is_parallel);
   if (cclauses)
     {
       clauses = c_oacc_split_loop_clauses (clauses, cclauses, is_parallel);
       if (*cclauses)
-	*cclauses = c_finish_omp_clauses (*cclauses, C_ORT_ACC);
+	*cclauses = c_finish_omp_clauses (*cclauses, C_ORT_ACC_TARGET);
       if (clauses)
 	clauses = c_finish_omp_clauses (clauses, C_ORT_ACC);
     }
@@ -18570,7 +18573,9 @@  c_parser_oacc_compute (location_t loc, c_parser *parser,
 	}
     }
 
-  tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name);
+  tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name,
+					    /*finish_p=*/true,
+					    /*target=*/true);
 
   tree block = c_begin_omp_parallel ();
   add_stmt (c_parser_omp_structured_block (parser, if_p));
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index ca32d8cd5d21..a4f3b855f6e7 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -13634,6 +13634,7 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 			     enum c_omp_region_type ort)
 {
   tree ret, low_bound, length, type;
+  bool openacc = (ort & C_ORT_ACC) != 0;
   if (TREE_CODE (t) != TREE_LIST)
     {
       if (error_operand_p (t))
@@ -13753,7 +13754,7 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	{
 	  error_at (OMP_CLAUSE_LOCATION (c),
 		    "expected single pointer in %qs clause",
-		    user_omp_clause_code_name (c, ort == C_ORT_ACC));
+		    user_omp_clause_code_name (c, openacc));
 	  return error_mark_node;
 	}
     }
@@ -14201,9 +14202,7 @@  handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 
       c_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t);
 
-      tree nc = ai.expand_map_clause (c, first, addr_tokens,
-				      (ort == C_ORT_OMP_TARGET
-				       || ort == C_ORT_ACC));
+      tree nc = ai.expand_map_clause (c, first, addr_tokens, ort);
       if (nc != error_mark_node)
 	{
 	  if (ai.maybe_zero_length_array_section (c))
@@ -14486,6 +14485,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   bool allocate_seen = false;
   bool implicit_moved = false;
   bool target_in_reduction_seen = false;
+  bool openacc = (ort & C_ORT_ACC) != 0;
 
   bitmap_obstack_initialize (NULL);
   bitmap_initialize (&generic_head, &bitmap_default_obstack);
@@ -14501,7 +14501,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
   bitmap_initialize (&is_on_device_head, &bitmap_default_obstack);
 
-  if (ort & C_ORT_ACC)
+  if (openacc)
     for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
       if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC)
 	{
@@ -14895,8 +14895,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	      remove = true;
 	    }
-	  else if ((ort == C_ORT_ACC
-		    && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+	  else if ((openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
 		   || (ort == C_ORT_OMP
 		       && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
 			   || (OMP_CLAUSE_CODE (c)
@@ -14919,7 +14918,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
 		{
 		  error_at (OMP_CLAUSE_LOCATION (c),
-			    ort == C_ORT_ACC
+			    openacc
 			    ? "%qD appears more than once in reduction clauses"
 			    : "%qD appears more than once in data clauses",
 			    t);
@@ -14942,7 +14941,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		    || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
 		   && bitmap_bit_p (&map_head, DECL_UID (t)))
 	    {
-	      if (ort == C_ORT_ACC)
+	      if (openacc)
 		error_at (OMP_CLAUSE_LOCATION (c),
 			  "%qD appears more than once in data clauses", t);
 	      else
@@ -15010,7 +15009,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  else if (bitmap_bit_p (&map_head, DECL_UID (t))
 		   || bitmap_bit_p (&map_field_head, DECL_UID (t)))
 	    {
-	      if (ort == C_ORT_ACC)
+	      if (openacc)
 		error_at (OMP_CLAUSE_LOCATION (c),
 			  "%qD appears more than once in data clauses", t);
 	      else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
@@ -15357,7 +15356,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			      error_at (OMP_CLAUSE_LOCATION (c),
 					"%qD appears more than once in motion "
 					"clauses", rt);
-			    else if (ort == C_ORT_ACC)
+			    else if (openacc)
 			      error_at (OMP_CLAUSE_LOCATION (c),
 					"%qD appears more than once in data "
 					"clauses", rt);
@@ -15445,8 +15444,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	       to be written.  */
 	    if (addr_tokens[0]->type == STRUCTURE_BASE
 		&& (bitmap_bit_p (&map_field_head, DECL_UID (t))
-		    || (ort != C_ORT_ACC
-			&& bitmap_bit_p (&map_head, DECL_UID (t)))))
+		    || (!openacc && bitmap_bit_p (&map_head, DECL_UID (t)))))
 	      goto skip_decl_checks;
 
 	    if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
@@ -15512,7 +15510,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  }
 		else if (bitmap_bit_p (&map_head, DECL_UID (t))
 			 && !bitmap_bit_p (&map_field_head, DECL_UID (t))
-			 && ort == C_ORT_ACC)
+			 && openacc)
 		  {
 		    error_at (OMP_CLAUSE_LOCATION (c),
 			      "%qD appears more than once in data clauses", t);
@@ -15527,7 +15525,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 		  error_at (OMP_CLAUSE_LOCATION (c),
 			    "%qD appears more than once in motion clauses", t);
-		else if (ort == C_ORT_ACC)
+		else if (openacc)
 		  error_at (OMP_CLAUSE_LOCATION (c),
 			    "%qD appears more than once in data clauses", t);
 		else
@@ -15535,8 +15533,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			    "%qD appears more than once in map clauses", t);
 		remove = true;
 	      }
-	    else if (ort == C_ORT_ACC
-		     && bitmap_bit_p (&generic_head, DECL_UID (t)))
+	    else if (openacc && bitmap_bit_p (&generic_head, DECL_UID (t)))
 	      {
 		error_at (OMP_CLAUSE_LOCATION (c),
 			  "%qD appears more than once in data clauses", t);
@@ -15545,7 +15542,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))
 		     || bitmap_bit_p (&is_on_device_head, DECL_UID (t)))
 	      {
-		if (ort == C_ORT_ACC)
+		if (openacc)
 		  error_at (OMP_CLAUSE_LOCATION (c),
 			    "%qD appears more than once in data clauses", t);
 		else
@@ -15578,9 +15575,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		grp_start_p = pc;
 		grp_sentinel = OMP_CLAUSE_CHAIN (c);
 		tree nc = ai.expand_map_clause (c, OMP_CLAUSE_DECL (c),
-						addr_tokens,
-						(ort == C_ORT_OMP_TARGET
-						 || ort == C_ORT_ACC));
+						addr_tokens, ort);
 		if (nc != error_mark_node)
 		  c = nc;
 	      }
@@ -15661,7 +15656,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
 	    {
 	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
-		  && ort != C_ORT_ACC)
+		  && !openacc)
 		{
 		  error_at (OMP_CLAUSE_LOCATION (c),
 			    "%qs variable is not a pointer",
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index bfd8aeae39f6..3905d561152b 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -40764,7 +40764,7 @@  cp_parser_oacc_clause_async (cp_parser *parser, tree list)
 static tree
 cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 			    const char *where, cp_token *pragma_tok,
-			    bool finish_p = true)
+			    bool finish_p = true, bool target = false)
 {
   tree clauses = NULL;
   bool first = true;
@@ -40974,7 +40974,7 @@  cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
   cp_parser_skip_to_pragma_eol (parser, pragma_tok);
 
   if (finish_p)
-    return finish_omp_clauses (clauses, C_ORT_ACC);
+    return finish_omp_clauses (clauses, target ? C_ORT_ACC_TARGET : C_ORT_ACC);
 
   return clauses;
 }
@@ -45587,7 +45587,7 @@  cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
   bool found_in_scope = global_bindings_p ();
 
   clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK,
-					"#pragma acc declare", pragma_tok, true);
+					"#pragma acc declare", pragma_tok);
 
 
   if (omp_find_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
@@ -45835,12 +45835,13 @@  cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
   mask |= OACC_LOOP_CLAUSE_MASK;
 
   tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok,
-					     cclauses == NULL);
+					     /*finish_p=*/cclauses == NULL,
+					     /*target=*/is_parallel);
   if (cclauses)
     {
       clauses = c_oacc_split_loop_clauses (clauses, cclauses, is_parallel);
       if (*cclauses)
-	*cclauses = finish_omp_clauses (*cclauses, C_ORT_ACC);
+	*cclauses = finish_omp_clauses (*cclauses, C_ORT_ACC_TARGET);
       if (clauses)
 	clauses = finish_omp_clauses (clauses, C_ORT_ACC);
     }
@@ -45965,7 +45966,9 @@  cp_parser_oacc_compute (cp_parser *parser, cp_token *pragma_tok,
 	}
     }
 
-  tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok);
+  tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok,
+					     /*finish_p=*/true,
+					     /*target=*/true);
 
   tree block = begin_omp_parallel ();
   unsigned int save = cp_parser_begin_omp_structured_block (parser);
diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc
index e68c74913f5d..63772bd64e73 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -19258,8 +19258,8 @@  tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl)
     case OACC_KERNELS:
     case OACC_PARALLEL:
     case OACC_SERIAL:
-      tmp = tsubst_omp_clauses (OMP_CLAUSES (t), C_ORT_ACC, args, complain,
-				in_decl);
+      tmp = tsubst_omp_clauses (OMP_CLAUSES (t), C_ORT_ACC_TARGET, args,
+				complain, in_decl);
       stmt = begin_omp_parallel ();
       RECUR (OMP_BODY (t));
       finish_omp_construct (TREE_CODE (t), stmt, tmp);
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index d6cf1cee6be2..f8f8299bbf64 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -5167,6 +5167,7 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 			     enum c_omp_region_type ort)
 {
   tree ret, low_bound, length, type;
+  bool openacc = (ort & C_ORT_ACC) != 0;
   if (TREE_CODE (t) != TREE_LIST)
     {
       if (error_operand_p (t))
@@ -5283,7 +5284,7 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	{
 	  error_at (OMP_CLAUSE_LOCATION (c),
 		    "expected single pointer in %qs clause",
-		    user_omp_clause_code_name (c, ort == C_ORT_ACC));
+		    user_omp_clause_code_name (c, openacc));
 	  return error_mark_node;
 	}
     }
@@ -5772,9 +5773,7 @@  handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
 
 	  cp_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t);
 
-	  tree nc = ai.expand_map_clause (c, first, addr_tokens,
-					  (ort == C_ORT_OMP_TARGET
-					   || ort == C_ORT_ACC));
+	  tree nc = ai.expand_map_clause (c, first, addr_tokens, ort);
 	  if (nc != error_mark_node)
 	    {
 	      using namespace omp_addr_tokenizer;
@@ -6721,6 +6720,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   bitmap_head oacc_reduction_head, is_on_device_head;
   tree c, t, *pc;
   tree safelen = NULL_TREE;
+  bool openacc = (ort & C_ORT_ACC) != 0;
   bool branch_seen = false;
   bool copyprivate_seen = false;
   bool ordered_seen = false;
@@ -6753,7 +6753,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
   bitmap_initialize (&is_on_device_head, &bitmap_default_obstack);
 
-  if (ort & C_ORT_ACC)
+  if (openacc)
     for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
       if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC)
 	{
@@ -7002,7 +7002,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    t = OMP_CLAUSE_DECL (c);
 	check_dup_generic_t:
 	  if (t == current_class_ptr
-	      && ((ort != C_ORT_OMP_DECLARE_SIMD && ort != C_ORT_ACC)
+	      && ((ort != C_ORT_OMP_DECLARE_SIMD && !openacc)
 		  || (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_LINEAR
 		      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_UNIFORM)))
 	    {
@@ -7027,7 +7027,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			  omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	      remove = true;
 	    }
-	  else if ((ort == C_ORT_ACC
+	  else if ((openacc
 		    && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
 		   || (ort == C_ORT_OMP
 		       && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
@@ -7051,7 +7051,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
 		{
 		  error_at (OMP_CLAUSE_LOCATION (c),
-			    ort == C_ORT_ACC
+			    openacc
 			    ? "%qD appears more than once in reduction clauses"
 			    : "%qD appears more than once in data clauses",
 			    t);
@@ -7074,7 +7074,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		    || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
 		   && bitmap_bit_p (&map_head, DECL_UID (t)))
 	    {
-	      if (ort == C_ORT_ACC)
+	      if (openacc)
 		error_at (OMP_CLAUSE_LOCATION (c),
 			  "%qD appears more than once in data clauses", t);
 	      else
@@ -7136,7 +7136,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
 	  else
 	    t = OMP_CLAUSE_DECL (c);
-	  if (ort != C_ORT_ACC && t == current_class_ptr)
+	  if (!openacc && t == current_class_ptr)
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
 			"%<this%> allowed in OpenMP only in %<declare simd%>"
@@ -7175,7 +7175,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  else if (bitmap_bit_p (&map_head, DECL_UID (t))
 		   || bitmap_bit_p (&map_field_head, DECL_UID (t)))
 	    {
-	      if (ort == C_ORT_ACC)
+	      if (openacc)
 		error_at (OMP_CLAUSE_LOCATION (c),
 			  "%qD appears more than once in data clauses", t);
 	      else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
@@ -7196,7 +7196,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
 	  else
 	    t = OMP_CLAUSE_DECL (c);
-	  if (ort != C_ORT_ACC && t == current_class_ptr)
+	  if (!openacc && t == current_class_ptr)
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
 			"%<this%> allowed in OpenMP only in %<declare simd%>"
@@ -8076,7 +8076,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			      error_at (OMP_CLAUSE_LOCATION (c),
 					"%qD appears more than once in motion"
 					" clauses", rt);
-			    else if (ort == C_ORT_ACC)
+			    else if (openacc)
 			      error_at (OMP_CLAUSE_LOCATION (c),
 					"%qD appears more than once in data"
 					" clauses", rt);
@@ -8173,8 +8173,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	       to be written.  */
 	    if (addr_tokens[0]->type == STRUCTURE_BASE
 		&& (bitmap_bit_p (&map_field_head, DECL_UID (t))
-		    || (ort != C_ORT_ACC
-			&& bitmap_bit_p (&map_head, DECL_UID (t)))))
+		    || (!openacc && bitmap_bit_p (&map_head, DECL_UID (t)))))
 	      goto skip_decl_checks;
 
 	    if (!processing_template_decl && TREE_CODE (t) == FIELD_DECL)
@@ -8270,7 +8269,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  }
 		else if (bitmap_bit_p (&map_head, DECL_UID (t))
 			 && !bitmap_bit_p (&map_field_head, DECL_UID (t))
-			 && ort == C_ORT_ACC)
+			 && openacc)
 		  {
 		    error_at (OMP_CLAUSE_LOCATION (c),
 			      "%qD appears more than once in data clauses", t);
@@ -8289,7 +8288,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 		  error_at (OMP_CLAUSE_LOCATION (c),
 			    "%qD appears more than once in motion clauses", t);
-		else if (ort == C_ORT_ACC)
+		else if (openacc)
 		  error_at (OMP_CLAUSE_LOCATION (c),
 			    "%qD appears more than once in data clauses", t);
 		else
@@ -8297,8 +8296,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			    "%qD appears more than once in map clauses", t);
 		remove = true;
 	      }
-	    else if (ort == C_ORT_ACC
-		     && bitmap_bit_p (&generic_head, DECL_UID (t)))
+	    else if (openacc && bitmap_bit_p (&generic_head, DECL_UID (t)))
 	      {
 		error_at (OMP_CLAUSE_LOCATION (c),
 			  "%qD appears more than once in data clauses", t);
@@ -8307,7 +8305,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))
 		     || bitmap_bit_p (&is_on_device_head, DECL_UID (t)))
 	      {
-		if (ort == C_ORT_ACC)
+		if (openacc)
 		  error_at (OMP_CLAUSE_LOCATION (c),
 			    "%qD appears more than once in data clauses", t);
 		else
@@ -8331,7 +8329,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      }
 
 	  skip_decl_checks:
-	    /* If we call omp_expand_map_clause in handle_omp_array_sections,
+	    /* If we call ai.expand_map_clause in handle_omp_array_sections,
 	       the containing loop (here) iterates through the new nodes
 	       created by that expansion.  Avoid expanding those again (just
 	       by checking the node type).  */
@@ -8339,8 +8337,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		&& !processing_template_decl
 		&& ort != C_ORT_DECLARE_SIMD
 		&& (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
-		    || ((OMP_CLAUSE_MAP_KIND (c)
-			 != GOMP_MAP_FIRSTPRIVATE_POINTER)
+		    || (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
 			&& (OMP_CLAUSE_MAP_KIND (c)
 			    != GOMP_MAP_FIRSTPRIVATE_REFERENCE)
 			&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
@@ -8349,9 +8346,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		grp_start_p = pc;
 		grp_sentinel = OMP_CLAUSE_CHAIN (c);
 		tree nc = ai.expand_map_clause (c, OMP_CLAUSE_DECL (c),
-						addr_tokens,
-						(ort == C_ORT_OMP_TARGET
-						 || ort == C_ORT_ACC));
+						addr_tokens, ort);
 		if (nc != error_mark_node)
 		  c = nc;
 	      }