[PATCH, OpenMP 5.0] Target mapping C++ members inside member functions

[PATCH, OpenMP 5.0] Target mapping C++ members inside member functions

[Chung-Lin Tang]

[-- Attachment #1: Type: text/plain, Size: 1319 bytes --]

Hi Jakub,
this patch is to fix some of the problems with target mapping when inside
C++ member functions:

1. Allow deref '->' in map clauses.
2. Allow this[X] in map clauses.
3. Create map(this->member) from map(member), when encountering member's FIELD_DECL.

This actually may not be the last related patch, since this patch is only for C++,
while the deref feature parity probably is also needed in C. Also the convention
required "map(struct:) map(<element>)..." sequence still needs further middle-end
gimplify modifications to generate. However this is already useful for compiling
some programs, and enter/exit data usage.

Tested C++ and libgomp without regressions, is this okay? Probably will need to wait
till stage1 to commit.

Thanks,
Chung-Lin

2020-12-23  Chung-Lin Tang  <cltang@codesourcery.com>

         gcc/cp/
         * parser.c (cp_parser_omp_clause_map): Adjust call to
         cp_parser_omp_var_list_no_open to set 'allow_deref' argument to true.
         * semantics.c (handle_omp_array_sections_1): Add handling to create
         'this->member' from 'member' FIELD_DECL.
         (finish_omp_clauses): Likewise. Adjust to allow 'this[]' in OpenMP
         map clauses.

         gcc/testsuite/
         * g++.dg/gomp/target-3.C: New test.
         * g++.dg/gomp/this-2.C: Adjust testcase.

[-- Attachment #2: ext-map-clause.patch --]
[-- Type: text/plain, Size: 6901 bytes --]

diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 7ea8c28830e..1d5a508dd71 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -37720,7 +37720,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
     }
 
   nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list,
-					  NULL);
+					  NULL, true);
 
   for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
     OMP_CLAUSE_SET_MAP_KIND (c, kind);
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 92e32c8e0ad..d21b58df954 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4808,6 +4808,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
       if (REFERENCE_REF_P (t)
 	  && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
 	t = TREE_OPERAND (t, 0);
+      if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+	  && TREE_CODE (t) == FIELD_DECL)
+	t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
       ret = t;
       if (TREE_CODE (t) == COMPONENT_REF
 	  && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -4833,8 +4836,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		  return error_mark_node;
 		}
 	      t = TREE_OPERAND (t, 0);
-	      if (ort == C_ORT_ACC && TREE_CODE (t) == INDIRECT_REF)
-		t = TREE_OPERAND (t, 0);
+	      if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+		  && TREE_CODE (t) == INDIRECT_REF)
+		{
+		  t = TREE_OPERAND (t, 0);
+		  STRIP_NOPS (t);
+		}
 	    }
 	  if (REFERENCE_REF_P (t))
 	    t = TREE_OPERAND (t, 0);
@@ -4854,6 +4861,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	  return error_mark_node;
 	}
       else if (ort == C_ORT_OMP
+	       && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+	       && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TO
+	       && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FROM
 	       && TREE_CODE (t) == PARM_DECL
 	       && DECL_ARTIFICIAL (t)
 	       && DECL_NAME (t) == this_identifier)
@@ -7441,6 +7451,11 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			t = TREE_OPERAND (t, 0);
 		      if (REFERENCE_REF_P (t))
 			t = TREE_OPERAND (t, 0);
+		      if (TREE_CODE (t) == INDIRECT_REF)
+			{
+			  t = TREE_OPERAND (t, 0);
+			  STRIP_NOPS (t);
+			}
 		      if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
 			break;
 		      if (bitmap_bit_p (&map_head, DECL_UID (t)))
@@ -7555,6 +7570,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		    goto handle_map_references;
 		}
 	    }
+	  if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+	      && !processing_template_decl
+	      && TREE_CODE (t) == FIELD_DECL)
+	    {
+	      OMP_CLAUSE_DECL (c) = finish_non_static_data_member (t, NULL_TREE,
+								   NULL_TREE);
+	      break;
+	    }
 	  if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
 	    {
 	      if (processing_template_decl && TREE_CODE (t) != OVERLOAD)
@@ -7581,7 +7604,9 @@ 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 && t == current_class_ptr)
+	  else if (ort != C_ORT_ACC
+		   && ort != C_ORT_OMP
+		   && t == current_class_ptr)
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
 			"%<this%> allowed in OpenMP only in %<declare simd%>"
diff --git a/gcc/testsuite/g++.dg/gomp/target-3.C b/gcc/testsuite/g++.dg/gomp/target-3.C
new file mode 100644
index 00000000000..fe2e38b46a3
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-3.C
@@ -0,0 +1,36 @@
+// { dg-do compile }
+// { dg-options "-fopenmp -fdump-tree-gimple" }
+
+struct S
+{
+  int a, b;
+  void bar (int);
+};
+
+void
+S::bar (int x)
+{
+  #pragma omp target map (alloc: a, b)
+    ;
+  #pragma omp target enter data map (alloc: a, b)
+}
+
+template <int N>
+struct T
+{
+  int a, b;
+  void bar (int);
+};
+
+template <int N>
+void
+T<N>::bar (int x)
+{
+  #pragma omp target map (alloc: a, b)
+    ;
+  #pragma omp target enter data map (alloc: a, b)
+}
+
+template struct T<0>;
+
+/* { dg-final { scan-tree-dump-times "map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/this-2.C b/gcc/testsuite/g++.dg/gomp/this-2.C
index d03b8a0728e..b521a4faf5e 100644
--- a/gcc/testsuite/g++.dg/gomp/this-2.C
+++ b/gcc/testsuite/g++.dg/gomp/this-2.C
@@ -9,14 +9,14 @@ struct S
 void
 S::bar (int x)
 {
-  #pragma omp target map (this, x)		// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+  #pragma omp target map (this, x)		// { dg-error "cannot take the address of .this., which is an rvalue expression" }
     ;
-  #pragma omp target map (this[0], x)		// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+  #pragma omp target map (this[0], x)
     ;
-  #pragma omp target update to (this, x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
-  #pragma omp target update to (this[0], x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
-  #pragma omp target update from (this, x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
-  #pragma omp target update from (this[1], x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+  #pragma omp target update to (this, x)	// { dg-error "cannot take the address of .this., which is an rvalue expression" }
+  #pragma omp target update to (this[0], x)
+  #pragma omp target update from (this, x)	// { dg-error "cannot take the address of .this., which is an rvalue expression" }
+  #pragma omp target update from (this[1], x)
 }
 
 template <int N>
@@ -29,14 +29,14 @@ template <int N>
 void
 T<N>::bar (int x)
 {
-  #pragma omp target map (this, x)		// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+  #pragma omp target map (this, x)		// { dg-error "cannot take the address of .this., which is an rvalue expression" }
     ;
-  #pragma omp target map (this[0], x)		// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+  #pragma omp target map (this[0], x)
     ;
-  #pragma omp target update to (this, x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
-  #pragma omp target update to (this[0], x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
-  #pragma omp target update from (this, x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
-  #pragma omp target update from (this[1], x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+  #pragma omp target update to (this, x)	// { dg-error "cannot take the address of .this., which is an rvalue expression" }
+  #pragma omp target update to (this[0], x)
+  #pragma omp target update from (this, x)	// { dg-error "cannot take the address of .this., which is an rvalue expression" }
+  #pragma omp target update from (this[1], x)
 }
 
 template struct T<0>;

Re: [PATCH, OpenMP 5.0] Target mapping C++ members inside member functions

[Jakub Jelinek]

On Wed, Dec 23, 2020 at 07:02:09PM +0800, Chung-Lin Tang wrote:
> Hi Jakub,
> this patch is to fix some of the problems with target mapping when inside
> C++ member functions:
> 
> 1. Allow deref '->' in map clauses.
> 2. Allow this[X] in map clauses.

I must say I'd prefer not to go this way by adding further and further
exceptions on what we allow during parsing.
The OpenMP 5.0 wording basically says that we should parse (for map and
to/from) arbitrary base language expressions with the OpenMP additional
extensions (e.g. the array section syntax and array shaping), which I think
means we need to set some flag in the parse when we start parsing the
expression in the clause, which will tell the parser that array sections
and/or array shaping is allowed, add tree codes for ARRAY_SECTION_REF and
the array shaping, parse it (ideally with as few foldings as possible if the
foldings would prevent us from behaving correctly) and when parsed (probably
in the *finish_omp_clauses helper routines) go over those expressions,
diagnose ARRAY_SECTION_REF and/or array shaping used in inappropriate
locations in the expression and then go through what the 5.0 spec says on
how the expressions should be mapped, what to do on ARRAY_REF, what to do on
ARRAY_SECTION_REF, what to do on COMPONENT_REF, what to do on INDIRECT_REF,
...
Because I'm afraid we can't avoid doing that and when we do, the -> and
this[X] etc. handling would be thrown away again.
The hack to just parse decl or optional . member after it several times)
followed by [ with array section syntax in there was possibly only for 4.5
and earlier.

	Jakub