[gomp4] fix c++ reference mappings in openacc

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;
+}