[gomp4] fix c++ reference mappings in openacc

[gomp4] fix c++ reference mappings in openacc

[Cesar Philippidis]

I've applied this patch to gomp-4_0-branch which fixes of problems
involving reference type variables in openacc data clauses. The first
problem was, the c++ front end was incorrectly handling reference types
in general in openacc. Instead of mapping the variable, it would map the
pointer to the variable by itself. The second problem was, if the
gimplifier saw a pointer mapping for a data clause, it would propagate
it to omp-lower. That's bad because if you have something like this

  int &var = ...

  #pragma acc data copy (var)
  {
     ...var...
  }

where the var inside the data region would have some uninitialized value
because omplower installs a new variable for it. The gimpifier is
already handling openmp target data regions properly, so this patch
extends it to ignore pointer mappings in acc enter/exit and data constructs.

Ultimately this patch will need to go in trunk, but the c++ changes
don't apply cleanly. I'll need to work on that later.

Cesar

Re: [gomp4] fix c++ reference mappings in openacc

[Cesar Philippidis]

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

On 01/20/2016 07:46 PM, Cesar Philippidis wrote:
> I've applied this patch to gomp-4_0-branch which fixes of problems
> involving reference type variables in openacc data clauses. The first
> problem was, the c++ front end was incorrectly handling reference types
> in general in openacc. Instead of mapping the variable, it would map the
> pointer to the variable by itself. The second problem was, if the
> gimplifier saw a pointer mapping for a data clause, it would propagate
> it to omp-lower. That's bad because if you have something like this
> 
>   int &var = ...
> 
>   #pragma acc data copy (var)
>   {
>      ...var...
>   }
> 
> where the var inside the data region would have some uninitialized value
> because omplower installs a new variable for it. The gimpifier is
> already handling openmp target data regions properly, so this patch
> extends it to ignore pointer mappings in acc enter/exit and data constructs.
> 
> Ultimately this patch will need to go in trunk, but the c++ changes
> don't apply cleanly. I'll need to work on that later.

And here's the patch.

Cesar

[-- Attachment #2: cxx-data-20160120.diff --]
[-- Type: text/x-patch, Size: 6012 bytes --]

2016-01-20  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/cp/
	* parser.c (cp_parser_oacc_all_clauses): Call finish_omp_clauses
	with allow_fields set to true.
	(cp_parser_oacc_cache): Likewise.
	(cp_parser_oacc_loop): Likewise.
	* semantics.c (finish_omp_clauses): Ensure that is_oacc is properly
	set when calling hanlde_omp_array_sections.

	gcc/
	* gimplify.c (gimplify_scan_omp_clauses):  Consider OACC_{DATA,
	PARALLEL, KERNELS} when processing firstprivate pointers and
	references, and setting target_kind_firstprivatize_array_bases.

	libgomp/
	* testsuite/libgomp.oacc-c++/non-scalar-data.C: New test.


diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index d88877a..4882b19 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -32324,7 +32324,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, true, false);
+    return finish_omp_clauses (clauses, true, true);
 
   return clauses;
 }
@@ -35140,7 +35140,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
   tree stmt, clauses;
 
   clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_TREE);
-  clauses = finish_omp_clauses (clauses, true, false);
+  clauses = finish_omp_clauses (clauses, true, true);
 
   cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer));
 
@@ -35471,9 +35471,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
     {
       clauses = c_oacc_split_loop_clauses (clauses, cclauses);
       if (*cclauses)
-	finish_omp_clauses (*cclauses, true, false);
+	finish_omp_clauses (*cclauses, true, true);
       if (clauses)
-	finish_omp_clauses (clauses, true, false);
+	finish_omp_clauses (clauses, true, true);
     }
 
   tree block = begin_omp_structured_block ();
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 3ca6137..e161186 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5807,7 +5807,7 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c, allow_fields))
+	      if (handle_omp_array_sections (c, allow_fields && !is_oacc))
 		{
 		  remove = true;
 		  break;
@@ -6567,7 +6567,7 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
 	    }
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c, allow_fields))
+	      if (handle_omp_array_sections (c, allow_fields && !is_oacc))
 		remove = true;
 	      break;
 	    }
@@ -6601,7 +6601,7 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c, allow_fields))
+	      if (handle_omp_array_sections (c, allow_fields && !is_oacc))
 		remove = true;
 	      else
 		{
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index cdb5b96..152942f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6092,7 +6092,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
 	{
 	  unsigned nflags = flags;
 	  if (ctx->target_map_pointers_as_0len_arrays
-	      || ctx->target_map_scalars_firstprivate)
+	       || ctx->target_map_scalars_firstprivate)
 	    {
 	      bool is_declare_target = false;
 	      bool is_scalar = false;
@@ -6456,7 +6456,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
       case OMP_TARGET_DATA:
       case OMP_TARGET_ENTER_DATA:
       case OMP_TARGET_EXIT_DATA:
+      case OACC_DATA:
       case OACC_HOST_DATA:
+      case OACC_PARALLEL:
+      case OACC_KERNELS:
 	ctx->target_firstprivatize_array_bases = true;
       default:
 	break;
@@ -6726,7 +6729,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    case OMP_TARGET_DATA:
 	    case OMP_TARGET_ENTER_DATA:
 	    case OMP_TARGET_EXIT_DATA:
+	    case OACC_DATA:
 	    case OACC_HOST_DATA:
+	    case OACC_ENTER_DATA:
+	    case OACC_EXIT_DATA:
 	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		  || (OMP_CLAUSE_MAP_KIND (c)
 		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
diff --git a/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C
new file mode 100644
index 0000000..180e86f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C
@@ -0,0 +1,109 @@
+// Ensure that a non-scalar dummy arguments which are implicitly used inside
+// offloaded regions are properly mapped using present_or_copy.
+
+// { dg-do run }
+
+#include <cassert>
+
+const int n = 100;
+
+struct data {
+  int v;
+};
+
+void
+kernels_present (data &d, int &x)
+{
+#pragma acc kernels present (d, x) default (none)
+  {
+    d.v = x;
+  }
+}
+
+void
+parallel_present (data &d, int &x)
+{
+#pragma acc parallel present (d, x) default (none)
+  {
+    d.v = x;
+  }
+}
+
+void
+kernels_implicit (data &d, int &x)
+{
+#pragma acc kernels
+  {
+    d.v = x;
+  }
+}
+
+void
+parallel_implicit (data &d, int &x)
+{
+#pragma acc parallel
+  {
+    d.v = x;
+  }
+}
+
+void
+reference_data (data &d, int &x)
+{
+#pragma acc data copy(d, x)
+  {
+    kernels_present (d, x);
+
+#pragma acc update host(d)
+    assert (d.v == x);
+
+    x = 200;
+#pragma acc update device(x)
+    
+    parallel_present (d, x);
+  }
+
+  assert (d.v = x);
+
+  x = 300;
+  kernels_implicit (d, x);
+  assert (d.v = x);
+
+  x = 400;
+  parallel_implicit (d, x);
+  assert (d.v = x);
+}
+
+int
+main ()
+{
+  data d;
+  int x = 100;
+
+#pragma acc data copy(d, x)
+  {
+    kernels_present (d, x);
+
+#pragma acc update host(d)
+    assert (d.v == x);
+
+    x = 200;
+#pragma acc update device(x)
+    
+    parallel_present (d, x);
+  }
+
+  assert (d.v = x);
+
+  x = 300;
+  kernels_implicit (d, x);
+  assert (d.v = x);
+
+  x = 400;
+  parallel_implicit (d, x);
+  assert (d.v = x);
+
+  reference_data (d, x);
+
+  return 0;
+}

[og7, openacc, libgomp, testsuite] Fix asserts in non-scalar-data.C

[Tom de Vries]

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

[ was: Re: [gomp4] fix c++ reference mappings in openacc ]


On 01/21/2016 04:47 AM, Cesar Philippidis wrote:
> diff --git a/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C
> new file mode 100644
> index 0000000..180e86f
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C
> @@ -0,0 +1,109 @@
> +// Ensure that a non-scalar dummy arguments which are implicitly used inside
> +// offloaded regions are properly mapped using present_or_copy.
> +
> +// { dg-do run }
> +
> +#include <cassert>
> +
> +const int n = 100;
> +
> +struct data {
> +  int v;
> +};
> +
> +void
> +kernels_present (data &d, int &x)
> +{
> +#pragma acc kernels present (d, x) default (none)
> +  {
> +    d.v = x;
> +  }
> +}
> +
> +void
> +parallel_present (data &d, int &x)
> +{
> +#pragma acc parallel present (d, x) default (none)
> +  {
> +    d.v = x;
> +  }
> +}
> +
> +void
> +kernels_implicit (data &d, int &x)
> +{
> +#pragma acc kernels
> +  {
> +    d.v = x;
> +  }
> +}
> +
> +void
> +parallel_implicit (data &d, int &x)
> +{
> +#pragma acc parallel
> +  {
> +    d.v = x;
> +  }
> +}
> +
> +void
> +reference_data (data &d, int &x)
> +{
> +#pragma acc data copy(d, x)
> +  {
> +    kernels_present (d, x);
> +
> +#pragma acc update host(d)
> +    assert (d.v == x);
> +
> +    x = 200;
> +#pragma acc update device(x)
> +
> +    parallel_present (d, x);
> +  }
> +
> +  assert (d.v = x);
> +
> +  x = 300;
> +  kernels_implicit (d, x);
> +  assert (d.v = x);
> +
> +  x = 400;
> +  parallel_implicit (d, x);
> +  assert (d.v = x);
> +}
> +
> +int
> +main ()
> +{
> +  data d;
> +  int x = 100;
> +
> +#pragma acc data copy(d, x)
> +  {
> +    kernels_present (d, x);
> +
> +#pragma acc update host(d)
> +    assert (d.v == x);
> +
> +    x = 200;
> +#pragma acc update device(x)
> +
> +    parallel_present (d, x);
> +  }
> +
> +  assert (d.v = x);
> +
> +  x = 300;
> +  kernels_implicit (d, x);
> +  assert (d.v = x);
> +
> +  x = 400;
> +  parallel_implicit (d, x);
> +  assert (d.v = x);
> +
> +  reference_data (d, x);
> +
> +  return 0;
> +}
> 

Some of these assert have assigns in them.

Fixed in attached patch, committed.

Thanks,
- Tom

[-- Attachment #2: 0002-openacc-libgomp-testsuite-Fix-asserts-in-non-scalar-data.C.patch --]
[-- Type: text/x-patch, Size: 1219 bytes --]

[openacc, libgomp, testsuite] Fix asserts in non-scalar-data.C

2018-04-18  Tom de Vries  <tom@codesourcery.com>

	* testsuite/libgomp.oacc-c++/non-scalar-data.C (reference_data, main):
	Fix asserts.

---
 libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C | 12 ++++++------
 1 file changed, 6 insertions(+), 6 deletions(-)

diff --git a/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C
index f24e31e..8143533 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C
@@ -69,15 +69,15 @@ reference_data (data &d, int &x)
     parallel_present (d, x);
   }
 
-  assert (d.v = x);
+  assert (d.v == x);
 
   x = 300;
   kernels_implicit (d, x);
-  assert (d.v = x);
+  assert (d.v == x);
 
   x = 400;
   parallel_implicit (d, x);
-  assert (d.v = x);
+  assert (d.v == x);
 }
 
 int
@@ -99,15 +99,15 @@ main ()
     parallel_present (d, x);
   }
 
-  assert (d.v = x);
+  assert (d.v == x);
 
   x = 300;
   kernels_implicit (d, x);
-  assert (d.v = x);
+  assert (d.v == x);
 
   x = 400;
   parallel_implicit (d, x);
-  assert (d.v = x);
+  assert (d.v == x);
 
   reference_data (d, x);