2016-02-09  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/c/
	* c-parser.c (c_parser_oacc_all_clauses): Update call to
	c_finish_omp_clauses.
	(c_parser_omp_all_clauses): Likewise.
	(c_parser_oacc_cache): Likewise.
	(c_parser_oacc_loop): Likewise.
	(omp_split_clauses): Likewise.
	(c_parser_omp_declare_target): Likewise.
	(c_parser_cilk_for): Likewise.
	* c-tree.h (c_finish_omp_clauses): Update prototype.
	* c-typeck.c (c_finish_omp_clauses): Add is_oacc argument.  Report
	OMP_CLAUSE_MAP errors as data clause errors for OpenACC.  Allow OpenACC
	reductions variables to appear in data clauses.

	gcc/cp/
	* cp-tree.h (finish_omp_clauses): Update prototype.
	* parser.c (cp_parser_oacc_all_clauses): Update call to
	finish_omp_clauses.
	(cp_parser_omp_all_clauses): Likewise.
	(cp_parser_omp_for_loop): Likewise.
	(cp_omp_split_clauses): Likewise.
	(cp_parser_oacc_cache): Likewise.
	(cp_parser_oacc_loop): Likewise.
	(cp_parser_omp_declare_target): Likewise.
	(cp_parser_cilk_for): Likewise.
	* pt.c (tsubst_attribute): Update call to tsubst_omp_clauses.
	(tsubst_omp_clauses): New is_oacc argument.  Use it when calling
	finish_omp_clauses.
	(tsubst_omp_for_iterator): Update call to finish_omp_clauses.
	(tsubst_expr): Update calls to tsubst_omp_clauses.
	* semantics.c (finish_omp_clauses): Add is_oacc argument.  Report
	OMP_CLAUSE_MAP errors as data clause errors for OpenACC.  Check for
	duplicate reference mappings.  Exclude "omp declare simd"-isms when
	processing OpenACC clauses.  Allow OpenACC reductions variables to
	appear in data clauses.
	(finish_omp_for): Update call to finish_omp_clauses.

	gcc/
	* gimplify.c (gimplify_scan_omp_clauses): Consider OACC_{DATA, PARALLEL,
	KERNELS} when setting target_firstprivatize_array_bases.  Consider
	OACC_{DATA, ENTER_DATA, EXIT_DATA, UPDATE} when filtering out pointer
	mappings.  Also filter out GOMP_MAP_POINTER.

	gcc/testsuite/
	* c-c++-common/goacc/data-clause-duplicate-1.c: Adjust test.
	* c-c++-common/goacc/declare-2.c: Likewise.
	* c-c++-common/goacc/deviceptr-1.c: Likewise.
	* c-c++-common/goacc/kernels-alias-ipa-pta-3.c: Likewise.
	* c-c++-common/goacc/parallel-reduction.c: New test.
	* c-c++-common/goacc/pcopy.c: Adjust test.
	* c-c++-common/goacc/pcopyin.c: Likewise.
	* c-c++-common/goacc/pcopyout.c: Likewise.
	* c-c++-common/goacc/pcreate.c: Likewise.
	* c-c++-common/goacc/present-1.c: Likewise.
	* c-c++-common/goacc/private-reduction-1.c: New test.
	* c-c++-common/goacc/reduction-5.c: New test.
	* g++.dg/goacc/data-1.C: New test.
	* g++.dg/goacc/data-2.C: New test.
	* g++.dg/goacc/reduction-1.C: New test.
	* g++.dg/goacc/update.C: New test.
	* g++.dg/gomp/template-data.C: New test.

	libgomp/
	* testsuite/libgomp.c++/non-scalar-data.C: New test.
	* testsuite/libgomp.oacc-c++/data-references.C: New test.
	* testsuite/libgomp.oacc-c++/data-templates.C: New test.
	* testsuite/libgomp.oacc-c++/non-scalar-data-templates.C: New test.
	* testsuite/libgomp.oacc-c++/update-reference.C: New test.
	* testsuite/libgomp.oacc-c++/update-template.C: New test.
	* testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Adjust test.
	* testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/if-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/present-2.c: Likewise.

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index eede3a7..20ff7da 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -13115,7 +13115,7 @@ 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, false);
+    return c_finish_omp_clauses (clauses, false, true, false);
 
   return clauses;
 }
@@ -13400,8 +13400,8 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
   if (finish_p)
     {
       if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_UNIFORM)) != 0)
-	return c_finish_omp_clauses (clauses, true, true);
-      return c_finish_omp_clauses (clauses, true);
+	return c_finish_omp_clauses (clauses, true, false, true);
+      return c_finish_omp_clauses (clauses, true, false);
     }
 
   return clauses;
@@ -13435,7 +13435,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser)
   tree stmt, clauses;
 
   clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL);
-  clauses = c_finish_omp_clauses (clauses, false);
+  clauses = c_finish_omp_clauses (clauses, false, true);
 
   c_parser_skip_to_pragma_eol (parser);
 
@@ -13769,9 +13769,9 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
     {
       clauses = c_oacc_split_loop_clauses (clauses, cclauses);
       if (*cclauses)
-	c_finish_omp_clauses (*cclauses, false);
+	c_finish_omp_clauses (*cclauses, false, true);
       if (clauses)
-	c_finish_omp_clauses (clauses, false);
+	c_finish_omp_clauses (clauses, false, true);
     }
 
   tree block = c_begin_compound_stmt (true);
@@ -14942,7 +14942,7 @@ 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] = c_finish_omp_clauses (cclauses[i], true);
+      cclauses[i] = c_finish_omp_clauses (cclauses[i], true, false);
 }
 
 /* OpenMP 4.0:
@@ -16455,7 +16455,7 @@ c_parser_omp_declare_target (c_parser *parser)
     {
       clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO_DECLARE,
 					      clauses);
-      clauses = c_finish_omp_clauses (clauses, true);
+      clauses = c_finish_omp_clauses (clauses, true, false);
       c_parser_skip_to_pragma_eol (parser);
     }
   else
@@ -17503,7 +17503,7 @@ c_parser_cilk_for (c_parser *parser, tree grain)
   tree clauses = build_omp_clause (EXPR_LOCATION (grain), OMP_CLAUSE_SCHEDULE);
   OMP_CLAUSE_SCHEDULE_KIND (clauses) = OMP_CLAUSE_SCHEDULE_CILKFOR;
   OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (clauses) = grain;
-  clauses = c_finish_omp_clauses (clauses, false);
+  clauses = c_finish_omp_clauses (clauses, false, false);
 
   tree block = c_begin_compound_stmt (true);
   tree sb = push_stmt_list ();
@@ -17568,7 +17568,7 @@ c_parser_cilk_for (c_parser *parser, tree grain)
       OMP_CLAUSE_OPERAND (c, 0)
 	= cilk_for_number_of_iterations (omp_for);
       OMP_CLAUSE_CHAIN (c) = clauses;
-      OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c, true);
+      OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c, true, false);
       add_stmt (omp_par);
     }
 
diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h
index cf79ba7..152e0d1 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -660,7 +660,7 @@ extern tree c_begin_omp_task (void);
 extern tree c_finish_omp_task (location_t, tree, tree);
 extern void c_finish_omp_cancel (location_t, tree);
 extern void c_finish_omp_cancellation_point (location_t, tree);
-extern tree c_finish_omp_clauses (tree, bool, bool = false);
+extern tree c_finish_omp_clauses (tree, bool, bool, bool = false);
 extern tree c_build_va_arg (location_t, tree, location_t, tree);
 extern tree c_finish_transaction (location_t, tree, int);
 extern bool c_tree_equal (tree, tree);
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 65925cb..149cdd5 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12526,10 +12526,10 @@ c_find_omp_placeholder_r (tree *tp, int *, void *data)
    Remove any elements from the list that are invalid.  */
 
 tree
-c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
+c_finish_omp_clauses (tree clauses, bool is_omp, bool is_oacc, bool declare_simd)
 {
   bitmap_head generic_head, firstprivate_head, lastprivate_head;
-  bitmap_head aligned_head, map_head, map_field_head;
+  bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head;
   tree c, t, type, *pc;
   tree simdlen = NULL_TREE, safelen = NULL_TREE;
   bool branch_seen = false;
@@ -12546,6 +12546,7 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
   bitmap_initialize (&aligned_head, &bitmap_default_obstack);
   bitmap_initialize (&map_head, &bitmap_default_obstack);
   bitmap_initialize (&map_field_head, &bitmap_default_obstack);
+  bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
 
   for (pc = &clauses, c = clauses; c ; c = *pc)
     {
@@ -12866,6 +12867,16 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
 			omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	      remove = true;
 	    }
+	  else if (is_oacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+	    {
+	      if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
+		{
+		  error ("%qD appears more than once in reduction clauses", t);
+		  remove = true;
+		}
+	      else
+		bitmap_set_bit (&oacc_reduction_head, DECL_UID (t));
+	    }
 	  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)))
@@ -13157,8 +13168,10 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
 	    {
 	      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 		error ("%qD appears more than once in motion clauses", t);
-	      else
+	      else if (is_omp)
 		error ("%qD appears more than once in map clauses", t);
+	      else
+		error ("%qD appears more than once in data clauses", t);
 	      remove = true;
 	    }
 	  else if (bitmap_bit_p (&generic_head, DECL_UID (t))
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index ead017e..2b70c96 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -6417,7 +6417,7 @@ extern tree omp_reduction_id			(enum tree_code, tree, tree);
 extern tree cp_remove_omp_priv_cleanup_stmt	(tree *, int *, void *);
 extern void cp_check_omp_declare_reduction	(tree);
 extern void finish_omp_declare_simd_methods	(tree);
-extern tree finish_omp_clauses			(tree, bool, bool = false);
+extern tree finish_omp_clauses			(tree, bool, bool, bool = false);
 extern tree push_omp_privatization_clauses	(bool);
 extern void pop_omp_privatization_clauses	(tree);
 extern void save_omp_privatization_clauses	(vec<tree> &);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index d03b0c9..458abbe 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -32168,7 +32168,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, false);
+    return finish_omp_clauses (clauses, true, true);
 
   return clauses;
 }
@@ -32487,9 +32487,9 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
   if (finish_p)
     {
       if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_UNIFORM)) != 0)
-	return finish_omp_clauses (clauses, false, true);
+	return finish_omp_clauses (clauses, false, false, true);
       else
-	return finish_omp_clauses (clauses, true);
+	return finish_omp_clauses (clauses, false, true);
     }
   return clauses;
 }
@@ -33564,7 +33564,7 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tree_code code, tree clauses,
 	      else
 		c = build_omp_clause (loc, OMP_CLAUSE_LASTPRIVATE);
 	      OMP_CLAUSE_DECL (c) = add_private_clause;
-	      c = finish_omp_clauses (c, true);
+	      c = finish_omp_clauses (c, false, true);
 	      if (c)
 		{
 		  OMP_CLAUSE_CHAIN (c) = clauses;
@@ -33713,7 +33713,7 @@ 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], true);
+      cclauses[i] = finish_omp_clauses (cclauses[i], false, true);
 }
 
 /* OpenMP 4.0:
@@ -34985,7 +34985,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, false);
+  clauses = finish_omp_clauses (clauses, true, true);
 
   cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer));
 
@@ -35310,9 +35310,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, false);
+	finish_omp_clauses (*cclauses, true, true);
       if (clauses)
-	finish_omp_clauses (clauses, false);
+	finish_omp_clauses (clauses, true, true);
     }
 
   tree block = begin_omp_structured_block ();
@@ -35678,7 +35678,7 @@ cp_parser_omp_declare_target (cp_parser *parser, cp_token *pragma_tok)
     {
       clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO_DECLARE,
 					clauses);
-      clauses = finish_omp_clauses (clauses, true);
+      clauses = finish_omp_clauses (clauses, false, true);
       cp_parser_require_pragma_eol (parser, pragma_tok);
     }
   else
@@ -37652,7 +37652,7 @@ cp_parser_cilk_for (cp_parser *parser, tree grain)
   tree clauses = build_omp_clause (EXPR_LOCATION (grain), OMP_CLAUSE_SCHEDULE);
   OMP_CLAUSE_SCHEDULE_KIND (clauses) = OMP_CLAUSE_SCHEDULE_CILKFOR;
   OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (clauses) = grain;
-  clauses = finish_omp_clauses (clauses, false);
+  clauses = finish_omp_clauses (clauses, false, false);
 
   tree ret = cp_parser_omp_for_loop (parser, CILK_FOR, clauses, NULL);
   if (ret)
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 76a6019..278c110 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -9537,7 +9537,8 @@ can_complete_type_without_circularity (tree type)
     return 1;
 }
 
-static tree tsubst_omp_clauses (tree, bool, bool, tree, tsubst_flags_t, tree);
+static tree tsubst_omp_clauses (tree, bool, bool, tree, tsubst_flags_t, tree,
+				bool);
 
 /* Instantiate a single dependent attribute T (a TREE_LIST), and return either
    T or a new TREE_LIST, possibly a chain in the case of a pack expansion.  */
@@ -9557,9 +9558,9 @@ tsubst_attribute (tree t, tree *decl_p, tree args,
     {
       tree clauses = TREE_VALUE (val);
       clauses = tsubst_omp_clauses (clauses, true, false, args,
-				    complain, in_decl);
+				    complain, in_decl, false);
       c_omp_declare_simd_clauses_to_decls (*decl_p, clauses);
-      clauses = finish_omp_clauses (clauses, false, true);
+      clauses = finish_omp_clauses (clauses, false, false, true);
       tree parms = DECL_ARGUMENTS (*decl_p);
       clauses
 	= c_omp_declare_simd_clauses_to_numbers (parms, clauses);
@@ -14456,7 +14457,8 @@ tsubst_omp_clause_decl (tree decl, tree args, tsubst_flags_t complain,
 
 static tree
 tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
-		    tree args, tsubst_flags_t complain, tree in_decl)
+		    tree args, tsubst_flags_t complain, tree in_decl,
+		    bool is_oacc)
 {
   tree new_clauses = NULL_TREE, nc, oc;
   tree linear_no_step = NULL_TREE;
@@ -14669,7 +14671,7 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
   new_clauses = nreverse (new_clauses);
   if (!declare_simd)
     {
-      new_clauses = finish_omp_clauses (new_clauses, allow_fields);
+      new_clauses = finish_omp_clauses (new_clauses, is_oacc, allow_fields);
       if (linear_no_step)
 	for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc))
 	  if (nc == linear_no_step)
@@ -14890,7 +14892,7 @@ tsubst_omp_for_iterator (tree t, int i, tree declv, tree orig_declv,
 	{
 	  tree c = build_omp_clause (input_location, OMP_CLAUSE_PRIVATE);
 	  OMP_CLAUSE_DECL (c) = decl;
-	  c = finish_omp_clauses (c, true);
+	  c = finish_omp_clauses (c, false, true);
 	  if (c)
 	    {
 	      OMP_CLAUSE_CHAIN (c) = *clauses;
@@ -15373,8 +15375,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
 
     case OACC_KERNELS:
     case OACC_PARALLEL:
-      tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, false, args, complain,
-				in_decl);
+      tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true, args, complain,
+				in_decl, true);
       stmt = begin_omp_parallel ();
       RECUR (OMP_BODY (t));
       finish_omp_construct (TREE_CODE (t), stmt, tmp);
@@ -15383,7 +15385,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
     case OMP_PARALLEL:
       r = push_omp_privatization_clauses (OMP_PARALLEL_COMBINED (t));
       tmp = tsubst_omp_clauses (OMP_PARALLEL_CLAUSES (t), false, true,
-				args, complain, in_decl);
+				args, complain, in_decl, false);
       if (OMP_PARALLEL_COMBINED (t))
 	omp_parallel_combined_clauses = &tmp;
       stmt = begin_omp_parallel ();
@@ -15397,7 +15399,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
     case OMP_TASK:
       r = push_omp_privatization_clauses (false);
       tmp = tsubst_omp_clauses (OMP_TASK_CLAUSES (t), false, true,
-				args, complain, in_decl);
+				args, complain, in_decl, false);
       stmt = begin_omp_task ();
       RECUR (OMP_TASK_BODY (t));
       finish_omp_task (tmp, stmt);
@@ -15419,9 +15421,9 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
 	int i;
 
 	r = push_omp_privatization_clauses (OMP_FOR_INIT (t) == NULL_TREE);
-	clauses = tsubst_omp_clauses (OMP_FOR_CLAUSES (t), false,
-				      TREE_CODE (t) != OACC_LOOP,
-				      args, complain, in_decl);
+	clauses = tsubst_omp_clauses (OMP_FOR_CLAUSES (t), false, true,
+				      args, complain, in_decl,
+				      TREE_CODE (t) == OACC_LOOP);
 	if (OMP_FOR_INIT (t) != NULL_TREE)
 	  {
 	    declv = make_tree_vec (TREE_VEC_LENGTH (OMP_FOR_INIT (t)));
@@ -15478,7 +15480,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
       r = push_omp_privatization_clauses (TREE_CODE (t) == OMP_TEAMS
 					  && OMP_TEAMS_COMBINED (t));
       tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true,
-				args, complain, in_decl);
+				args, complain, in_decl, false);
       stmt = push_stmt_list ();
       RECUR (OMP_BODY (t));
       stmt = pop_stmt_list (stmt);
@@ -15493,9 +15495,8 @@ 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), false,
-				TREE_CODE (t) != OACC_DATA,
-				args, complain, in_decl);
+      tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true, args, complain,
+				in_decl, TREE_CODE (t) == OACC_DATA);
       keep_next_level (true);
       stmt = begin_omp_structured_block ();
 
@@ -15540,8 +15541,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
 
     case OACC_DECLARE:
       t = copy_node (t);
-      tmp = tsubst_omp_clauses (OACC_DECLARE_CLAUSES (t), false, false,
-				args, complain, in_decl);
+      tmp = tsubst_omp_clauses (OACC_DECLARE_CLAUSES (t), false, true,
+				args, complain, in_decl, true);
       OACC_DECLARE_CLAUSES (t) = tmp;
       add_stmt (t);
       break;
@@ -15550,7 +15551,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
     case OMP_TARGET_ENTER_DATA:
     case OMP_TARGET_EXIT_DATA:
       tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, true,
-				args, complain, in_decl);
+				args, complain, in_decl, false);
       t = copy_node (t);
       OMP_STANDALONE_CLAUSES (t) = tmp;
       add_stmt (t);
@@ -15559,8 +15560,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
     case OACC_ENTER_DATA:
     case OACC_EXIT_DATA:
     case OACC_UPDATE:
-      tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, false,
-				args, complain, in_decl);
+      tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, true,
+				args, complain, in_decl, true);
       t = copy_node (t);
       OMP_STANDALONE_CLAUSES (t) = tmp;
       add_stmt (t);
@@ -15568,7 +15569,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
 
     case OMP_ORDERED:
       tmp = tsubst_omp_clauses (OMP_ORDERED_CLAUSES (t), false, true,
-				args, complain, in_decl);
+				args, complain, in_decl, false);
       stmt = push_stmt_list ();
       RECUR (OMP_BODY (t));
       stmt = pop_stmt_list (stmt);
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index c9f9db4..ae6ba0f 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5736,10 +5736,11 @@ cp_finish_omp_clause_depend_sink (tree sink_clause)
    Remove any elements from the list that are invalid.  */
 
 tree
-finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
+finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
+		    bool declare_simd)
 {
   bitmap_head generic_head, firstprivate_head, lastprivate_head;
-  bitmap_head aligned_head, map_head, map_field_head;
+  bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head;
   tree c, t, *pc;
   tree safelen = NULL_TREE;
   bool branch_seen = false;
@@ -5753,6 +5754,7 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
   bitmap_initialize (&aligned_head, &bitmap_default_obstack);
   bitmap_initialize (&map_head, &bitmap_default_obstack);
   bitmap_initialize (&map_field_head, &bitmap_default_obstack);
+  bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
 
   for (pc = &clauses, c = clauses; c ; c = *pc)
     {
@@ -5966,6 +5968,16 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 		       omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	      remove = true;
 	    }
+	  else if (is_oacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+	    {
+	      if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
+		{
+		  error ("%qD appears more than once in reduction clauses", t);
+		  remove = true;
+		}
+	      else
+		bitmap_set_bit (&oacc_reduction_head, DECL_UID (t));
+	    }
 	  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)))
@@ -5976,7 +5988,10 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 	  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
 		   && bitmap_bit_p (&map_head, DECL_UID (t)))
 	    {
-	      error ("%qD appears both in data and map clauses", t);
+	      if (is_oacc)
+		error ("%qD appears more than once in data clauses", t);
+	      else
+		error ("%qD appears both in data and map clauses", t);
 	      remove = true;
 	    }
 	  else
@@ -6002,7 +6017,8 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 	    omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
 	  else
 	    t = OMP_CLAUSE_DECL (c);
-	  if (t == current_class_ptr)
+	  if (!is_oacc
+	      && t == current_class_ptr)
 	    {
 	      error ("%<this%> allowed in OpenMP only in %<declare simd%>"
 		     " clauses");
@@ -6028,7 +6044,10 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 	    }
 	  else if (bitmap_bit_p (&map_head, DECL_UID (t)))
 	    {
-	      error ("%qD appears both in data and map clauses", t);
+	      if (is_oacc)
+		error ("%qD appears more than once in data clauses", t);
+	      else
+		error ("%qD appears both in data and map clauses", t);
 	      remove = true;
 	    }
 	  else
@@ -6538,6 +6557,9 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 			  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 			    error ("%qD appears more than once in motion"
 				   " clauses", t);
+			  else if (is_oacc)
+			    error ("%qD appears more than once in data"
+				   " clauses", t);
 			  else
 			    error ("%qD appears more than once in map"
 				   " clauses", t);
@@ -6549,6 +6571,27 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 			  bitmap_set_bit (&map_field_head, DECL_UID (t));
 			}
 		    }
+		  else if (TREE_CODE (t) == TREE_LIST)
+		    {
+		      while (TREE_CODE (t = TREE_CHAIN (t)) == TREE_LIST)
+			;
+
+		      if (DECL_P (t))
+			{
+			  if (bitmap_bit_p (&map_head, DECL_UID (t)))
+			    {
+			      if (is_oacc)
+				error ("%qD appears more than once in data "
+				       "clauses", t);
+			      else
+				error ("%qD appears more than once in map "
+				       "clauses", t);
+			      remove = true;
+			    }
+			  else
+			    bitmap_set_bit (&map_head, DECL_UID (t));
+			}
+		    }
 		}
 	      break;
 	    }
@@ -6625,7 +6668,8 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 		     omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	      remove = true;
 	    }
-	  else if (t == current_class_ptr)
+	  else if (!is_oacc
+		   && t == current_class_ptr)
 	    {
 	      error ("%<this%> allowed in OpenMP only in %<declare simd%>"
 		     " clauses");
@@ -6666,7 +6710,10 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 		}
 	      else if (bitmap_bit_p (&map_head, DECL_UID (t)))
 		{
-		  error ("%qD appears both in data and map clauses", t);
+		  if (is_oacc)
+		    error ("%qD appears more than once in data clauses", t);
+		  else
+		    error ("%qD appears both in data and map clauses", t);
 		  remove = true;
 		}
 	      else
@@ -6676,6 +6723,8 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 	    {
 	      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 		error ("%qD appears more than once in motion clauses", t);
+	      if (is_oacc)
+		error ("%qD appears more than once in data clauses", t);
 	      else
 		error ("%qD appears more than once in map clauses", t);
 	      remove = true;
@@ -6683,7 +6732,10 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 	  else if (bitmap_bit_p (&generic_head, DECL_UID (t))
 		   || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
 	    {
-	      error ("%qD appears both in data and map clauses", t);
+	      if (is_oacc)
+		error ("%qD appears more than once in data clauses", t);
+	      else
+		error ("%qD appears both in data and map clauses", t);
 	      remove = true;
 	    }
 	  else
@@ -8260,7 +8312,7 @@ finish_omp_for (location_t locus, enum tree_code code, tree declv,
       OMP_CLAUSE_OPERAND (c, 0)
 	= cilk_for_number_of_iterations (omp_for);
       OMP_CLAUSE_CHAIN (c) = clauses;
-      OMP_PARALLEL_CLAUSES (omp_par) = finish_omp_clauses (c, false);
+      OMP_PARALLEL_CLAUSES (omp_par) = finish_omp_clauses (c, false, false);
       add_stmt (omp_par);
       return omp_par;
     }
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index b0ee27e..6584bb7 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6495,7 +6495,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;
@@ -6761,10 +6764,16 @@ 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:
+	    case OACC_UPDATE:
 	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		  || (OMP_CLAUSE_MAP_KIND (c)
-		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+		  || (!lang_GNU_Fortran () &&
+		      OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER))
 		/* For target {,enter ,exit }data only the array slice is
 		   mapped, but not the pointer to it.  */
 		remove = true;
diff --git a/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c b/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
index 7a1cf68..6245beb 100644
--- a/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
@@ -2,12 +2,12 @@ void
 fun (void)
 {
   float *fp;
-#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */
   ;
-#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in data clauses" } */
   ;
-#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
   ;
-#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
   ;
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/declare-2.c b/gcc/testsuite/c-c++-common/goacc/declare-2.c
index d24cb22..1b1e60a 100644
--- a/gcc/testsuite/c-c++-common/goacc/declare-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c
@@ -77,3 +77,5 @@ f (void)
 
 #pragma acc declare present (v9) /* { dg-error "invalid use of" } */
 }
+
+/* { dg-error "at file scope" "" { target c++ } 10 } */
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
index 546fa82..bc10b82 100644
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
@@ -47,7 +47,7 @@ fun2 (void)
   /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 46 } */
   /* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 46 } */
   /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 46 } */
-  /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 46 } */
+  /* { dg-error "'fp' appears more than once in data clauses" "fp more than once" { target *-*-* } 46 } */
   ;
 }
 
@@ -55,11 +55,11 @@ void
 fun3 (void)
 {
   float *fp;
-#pragma acc data deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc data deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in data clauses" } */
   ;
-#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
   ;
-#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
   ;
 }
 
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c
index 1eb56eb..97d8f52 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c
@@ -31,6 +31,6 @@ foo (void)
   free (c);
 }
 
-/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 1 "optimized" } } */
-/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" } } */
-/* { dg-final { scan-tree-dump-times "(?n)= \\*a" 1 "optimized" } } */
+/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 1 "optimized" { target c } } } */
+/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" { target c }  } } */
+/* { dg-final { scan-tree-dump-times "(?n)= \\*a" 1 "optimized" { target c }  } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c
new file mode 100644
index 0000000..d7cc947
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c
@@ -0,0 +1,17 @@
+int
+main ()
+{
+  int sum = 0;
+  int dummy = 0;
+
+#pragma acc data copy (dummy)
+  {
+#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum)
+    {
+      int v = 5;
+      sum += 10 + v;
+    }
+  }
+
+  return sum;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/pcopy.c b/gcc/testsuite/c-c++-common/goacc/pcopy.c
index 02c4383..7e39e42 100644
--- a/gcc/testsuite/c-c++-common/goacc/pcopy.c
+++ b/gcc/testsuite/c-c++-common/goacc/pcopy.c
@@ -7,4 +7,6 @@ f (char *cp)
   ;
 }
 
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(alloc:cp \\\[pointer assign, bias: 3]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(alloc:cp \\\[pointer assign, bias: 3]\\)" 1 "original" { target c } } } */
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 3]\\)" 1 "original" { target c++ } } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/pcopyin.c b/gcc/testsuite/c-c++-common/goacc/pcopyin.c
index 10911fc..f2c0354 100644
--- a/gcc/testsuite/c-c++-common/goacc/pcopyin.c
+++ b/gcc/testsuite/c-c++-common/goacc/pcopyin.c
@@ -7,4 +7,6 @@ f (char *cp)
   ;
 }
 
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(alloc:cp \\\[pointer assign, bias: 4]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(alloc:cp \\\[pointer assign, bias: 4]\\)" 1 "original" { target c } } } */
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 4]\\)" 1 "original" { target c++ } } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/pcopyout.c b/gcc/testsuite/c-c++-common/goacc/pcopyout.c
index 703ac2f..555a0b6 100644
--- a/gcc/testsuite/c-c++-common/goacc/pcopyout.c
+++ b/gcc/testsuite/c-c++-common/goacc/pcopyout.c
@@ -7,4 +7,6 @@ f (char *cp)
   ;
 }
 
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(alloc:cp \\\[pointer assign, bias: 5]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(alloc:cp \\\[pointer assign, bias: 5]\\)" 1 "original" { target c } } } */
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 5]\\)" 1 "original" { target c++ } } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/pcreate.c b/gcc/testsuite/c-c++-common/goacc/pcreate.c
index 00bf155..fa1b985 100644
--- a/gcc/testsuite/c-c++-common/goacc/pcreate.c
+++ b/gcc/testsuite/c-c++-common/goacc/pcreate.c
@@ -7,4 +7,6 @@ f (char *cp)
   ;
 }
 
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(alloc:cp \\\[pointer assign, bias: 6]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(alloc:cp \\\[pointer assign, bias: 6]\\)" 1 "original" { target c } } } */
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 6]\\)" 1 "original" { target c++ } } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/present-1.c b/gcc/testsuite/c-c++-common/goacc/present-1.c
index 7537948..52ff4c6 100644
--- a/gcc/testsuite/c-c++-common/goacc/present-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/present-1.c
@@ -7,4 +7,6 @@ f (char *cp)
   ;
 }
 
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(alloc:cp \\\[pointer assign, bias: 7]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(alloc:cp \\\[pointer assign, bias: 7]\\)" 1 "original" { target c } } } */
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 7]\\)" 1 "original" { target c++ } } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/private-reduction-1.c b/gcc/testsuite/c-c++-common/goacc/private-reduction-1.c
new file mode 100644
index 0000000..d4e3995
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/private-reduction-1.c
@@ -0,0 +1,12 @@
+int
+reduction ()
+{
+  int i, r;
+
+  #pragma acc parallel
+  #pragma acc loop private (r) reduction (+:r)
+  for (i = 0; i < 100; i++)
+    r += 10;
+
+  return r;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-5.c b/gcc/testsuite/c-c++-common/goacc/reduction-5.c
new file mode 100644
index 0000000..42cf9c2
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-5.c
@@ -0,0 +1,21 @@
+int
+main ()
+{
+  int sum = 0;
+  int dummy = 0;
+
+#pragma acc data copy (dummy)
+  {
+#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) reduction (+:sum) /* { dg-error "appears more than once in reduction clauses" } */
+    {
+      int i;
+      int v = 5;
+#pragma acc loop reduction (+:sum, sum) /* { dg-error "appears more than once in reduction clauses" } */
+      for (i = 0; i < 100; i++)
+	sum += 10 + v;
+    }
+  }
+
+  return sum;
+}
+
diff --git a/gcc/testsuite/g++.dg/goacc/data-1.C b/gcc/testsuite/g++.dg/goacc/data-1.C
new file mode 100644
index 0000000..54676dc
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/data-1.C
@@ -0,0 +1,39 @@
+void
+foo (int &a, int (&b)[100], int &n)
+{
+#pragma acc enter data copyin (a, b) async wait
+#pragma acc enter data create (b[20:30]) async wait
+#pragma acc enter data (a) /* { dg-error "expected '#pragma acc' clause before '\\\(' token" } */
+#pragma acc enter data create (b(1:10)) /* { dg-error "expected '\\\)' before '\\\(' token" } */
+#pragma acc exit data delete (a) if (0)
+#pragma acc exit data copyout (b) if (a)
+#pragma acc exit data delete (b)
+#pragma acc enter /* { dg-error "expected 'data' in" } */
+#pragma acc exit /* { dg-error "expected 'data' in" } */
+#pragma acc enter data /* { dg-error "has no data movement clause" } */
+#pragma acc exit data /* { dg-error "has no data movement clause" } */
+#pragma acc enter Data /* { dg-error "invalid pragma before" } */
+#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */
+}
+
+template<typename T>
+void
+foo (T &a, T (&b)[100], T &n)
+{
+#pragma acc enter data copyin (a, b) async wait
+#pragma acc enter data create (b[20:30]) async wait
+#pragma acc enter data (a) /* { dg-error "expected '#pragma acc' clause before '\\\(' token" } */
+#pragma acc enter data create (b(1:10)) /* { dg-error "expected '\\\)' before '\\\(' token" } */
+#pragma acc exit data delete (a) if (0)
+#pragma acc exit data copyout (b) if (a)
+#pragma acc exit data delete (b)
+#pragma acc enter /* { dg-error "expected 'data' in" } */
+#pragma acc exit /* { dg-error "expected 'data' in" } */
+#pragma acc enter data /* { dg-error "has no data movement clause" } */
+#pragma acc exit data /* { dg-error "has no data movement clause" } */
+#pragma acc enter Data /* { dg-error "invalid pragma before" } */
+#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */
+}
+
+/* { dg-error "has no data movement clause" "" { target *-*-* } 6 } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } 25 } */
diff --git a/gcc/testsuite/g++.dg/goacc/data-2.C b/gcc/testsuite/g++.dg/goacc/data-2.C
new file mode 100644
index 0000000..efa002d
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/data-2.C
@@ -0,0 +1,30 @@
+void
+fun (float (&fp)[100])
+{
+  float *dptr = &fp[50];
+
+#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */
+  ;
+#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in data clauses" } */
+  ;
+#pragma acc data create(fp[:10]) deviceptr(dptr)
+  ;
+#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
+  ;
+}
+
+template<typename T>
+void
+fun (T (&fp)[100])
+{
+  T *dptr = &fp[50];
+
+#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */
+  ;
+#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in data clauses" } */
+  ;
+#pragma acc data create(fp[:10]) deviceptr(dptr)
+  ;
+#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
+  ;
+}
diff --git a/gcc/testsuite/g++.dg/goacc/reduction-1.C b/gcc/testsuite/g++.dg/goacc/reduction-1.C
new file mode 100644
index 0000000..de97125
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/reduction-1.C
@@ -0,0 +1,72 @@
+/* { dg-require-effective-target alloca } */
+/* Integer reductions.  */
+
+#define vl 32
+
+int
+main(void)
+{
+  const int n = 1000;
+  int i;
+  int result, array[n];
+  int lresult;
+
+  /* '+' reductions.  */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (+:result)
+  for (i = 0; i < n; i++)
+    result += array[i];
+
+  /* '*' reductions.  */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (*:result)
+  for (i = 0; i < n; i++)
+    result *= array[i];
+
+//   result = 0;
+//   vresult = 0;
+// 
+//   /* 'max' reductions.  */
+// #pragma acc parallel vector_length (vl)
+// #pragma acc loop reduction (+:result)
+//   for (i = 0; i < n; i++)
+//       result = result > array[i] ? result : array[i];
+//
+//   /* 'min' reductions.  */
+// #pragma acc parallel vector_length (vl)
+// #pragma acc loop reduction (+:result)
+//   for (i = 0; i < n; i++)
+//       result = result < array[i] ? result : array[i];
+
+  /* '&' reductions.  */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (&:result)
+  for (i = 0; i < n; i++)
+    result &= array[i];
+
+  /* '|' reductions.  */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (|:result)
+  for (i = 0; i < n; i++)
+    result |= array[i];
+
+  /* '^' reductions.  */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (^:result)
+  for (i = 0; i < n; i++)
+    result ^= array[i];
+
+  /* '&&' reductions.  */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (&&:lresult)
+  for (i = 0; i < n; i++)
+    lresult = lresult && (result > array[i]);
+
+  /* '||' reductions.  */
+#pragma acc parallel vector_length (vl)
+#pragma acc loop reduction (||:lresult)
+  for (i = 0; i < n; i++)
+    lresult = lresult || (result > array[i]);
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/goacc/update.C b/gcc/testsuite/g++.dg/goacc/update.C
new file mode 100644
index 0000000..18091c9
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/update.C
@@ -0,0 +1,22 @@
+void
+f (int &i, int (&a)[10])
+{
+#pragma acc update device(i)
+#pragma acc update host(i)
+#pragma acc update self(i)
+#pragma acc update device(a[1:3])
+#pragma acc update host(a[1:3])
+#pragma acc update self(a[1:3])
+}
+
+template<typename T>
+void
+f (T &i, T (&a)[10])
+{
+#pragma acc update device(i)
+#pragma acc update host(i)
+#pragma acc update self(i)
+#pragma acc update device(a[1:3])
+#pragma acc update host(a[1:3])
+#pragma acc update self(a[1:3])
+}
diff --git a/gcc/testsuite/g++.dg/gomp/template-data.C b/gcc/testsuite/g++.dg/gomp/template-data.C
new file mode 100644
index 0000000..0be14d4
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/template-data.C
@@ -0,0 +1,18 @@
+void
+fun (float (&fp)[100])
+{
+  float *dptr = &fp[50];
+
+#pragma omp target data map(tofrom:fp[0:2], fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */
+  ;
+}
+
+template<typename T>
+void
+fun (T (&fp)[100])
+{
+  T *dptr = &fp[50];
+
+#pragma omp target data map(tofrom:fp[0:2], fp[0:2]) /* { dg-error "'fp' appears more than once in map clauses" } */
+  ;
+}
diff --git a/libgomp/testsuite/libgomp.c++/non-scalar-data.C b/libgomp/testsuite/libgomp.c++/non-scalar-data.C
new file mode 100644
index 0000000..180e86f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c++/data-references.C b/libgomp/testsuite/libgomp.oacc-c++/data-references.C
new file mode 100644
index 0000000..1f2abf4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/data-references.C
@@ -0,0 +1,178 @@
+/* Test data mappings on variables with reference types.  */
+
+/* { dg-do run } */
+
+#include <stdlib.h>
+
+void
+test (int &N, float *(&a), float *(&b), float *(&c), float *(&d), float *(&e))
+{
+  int i;
+
+  for (i = 0; i < N; i++)
+    {
+      a[i] = 3.0;
+      b[i] = 0.0;
+    }
+
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
+#pragma acc parallel async wait present (N, a[0:N], b[0:N])
+#pragma acc loop
+  for (i = 0; i < N; i++)
+    b[i] = a[i];
+
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async
+#pragma acc wait
+
+  for (i = 0; i < N; i++)
+    {
+      if (a[i] != 3.0)
+	abort ();
+
+      if (b[i] != 3.0)
+	abort ();
+    }
+
+  for (i = 0; i < N; i++)
+    {
+      a[i] = 2.0;
+      b[i] = 0.0;
+    }
+
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1)
+#pragma acc parallel async (1) present (N, a[0:N], b[0:N])
+#pragma acc loop
+  for (i = 0; i < N; i++)
+    b[i] = a[i];
+
+#pragma acc update host (a[0:N], b[0:N]) wait (1) async (1)
+#pragma acc wait (1)
+
+  for (i = 0; i < N; i++)
+    {
+      if (a[i] != 2.0)
+	abort ();
+
+      if (b[i] != 2.0)
+	abort ();
+    }
+
+  for (i = 0; i < N; i++)
+    {
+      a[i] = 3.0;
+      b[i] = 0.0;
+      c[i] = 0.0;
+      d[i] = 0.0;
+    }
+
+#pragma acc update device (a[0:N], b[0:N])
+#pragma acc enter data copyin (c[0:N], d[0:N]) async (1)
+
+#pragma acc parallel present (N, a[0:N], b[0:N]) async (1) wait (1)
+#pragma acc loop 
+  for (i = 0; i < N; i++)
+    b[i] = (a[i] * a[i] * a[i]) / a[i];
+
+#pragma acc parallel present (N, a[0:N], c[0:N]) async (2) wait (1)
+#pragma acc loop 
+  for (i = 0; i < N; i++)
+    c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
+
+#pragma acc parallel present (N, a[0:N], d[0:N]) async (3) wait (1)
+#pragma acc loop 
+  for (i = 0; i < N; i++)
+    d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
+
+#pragma acc exit data copyout (a[0:N], b[0:N], c[0:N], d[0:N]) wait (1, 2, 3) async (1)
+#pragma acc wait (1)
+
+  for (i = 0; i < N; i++)
+    {
+      if (a[i] != 3.0)
+	abort ();
+
+      if (b[i] != 9.0)
+	abort ();
+
+      if (c[i] != 4.0)
+	abort ();
+
+      if (d[i] != 1.0)
+	abort ();
+    }
+
+  for (i = 0; i < N; i++)
+    {
+      a[i] = 2.0;
+      b[i] = 0.0;
+      c[i] = 0.0;
+      d[i] = 0.0;
+      e[i] = 0.0;
+    }
+
+#pragma acc enter data copyin (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) async (1)
+
+#pragma acc parallel loop present (N, a[0:N], b[0:N]) async (1) wait (1) 
+  for (int ii = 0; ii < N; ii++)
+    b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
+
+#pragma acc parallel loop present (N, a[0:N], c[0:N]) async (2) wait (1) 
+  for (int ii = 0; ii < N; ii++)
+    c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
+
+#pragma acc parallel loop present (N, a[0:N], d[0:N]) async (3) wait (1) 
+  for (int ii = 0; ii < N; ii++)
+    d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
+
+#pragma acc parallel loop present (N, a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) wait (1) async (4) 
+  for (int ii = 0; ii < N; ii++)
+    e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
+
+#pragma acc exit data copyout (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) wait (1, 2, 3, 4) async (1) delete (N)
+#pragma acc wait (1)
+
+  for (i = 0; i < N; i++)
+    {
+      if (a[i] != 2.0)
+	abort ();
+
+      if (b[i] != 4.0)
+	abort ();
+
+      if (c[i] != 4.0)
+	abort ();
+
+      if (d[i] != 1.0)
+	abort ();
+
+      if (e[i] != 11.0)
+	abort ();
+    }
+}
+
+int
+main (int argc, char **argv)
+{
+  int N = 128; //1024 * 1024;
+  float *a, *b, *c, *d, *e;
+  int i;
+  int nbytes;
+
+  nbytes = N * sizeof (float);
+
+  a = (float *) malloc (nbytes);
+  b = (float *) malloc (nbytes);
+  c = (float *) malloc (nbytes);
+  d = (float *) malloc (nbytes);
+  e = (float *) malloc (nbytes);
+
+  test (N, a, b, c, d, e);
+
+  free (a);
+  free (b);
+  free (c);
+  free (d);
+  free (e);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c++/data-templates.C b/libgomp/testsuite/libgomp.oacc-c++/data-templates.C
new file mode 100644
index 0000000..34bd556
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/data-templates.C
@@ -0,0 +1,179 @@
+/* Test data mappings on variables with template types.  */
+
+/* { dg-do run } */
+
+#include <stdlib.h>
+
+template<typename T1, typename T2>
+void
+test (T1 &N, T2 *(&a), T2 *(&b), T2 *(&c), T2 *(&d), T2 *(&e))
+{
+  int i;
+
+  for (i = 0; i < N; i++)
+    {
+      a[i] = 3.0;
+      b[i] = 0.0;
+    }
+
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
+#pragma acc parallel async wait present (N, a[0:N], b[0:N])
+#pragma acc loop
+  for (i = 0; i < N; i++)
+    b[i] = a[i];
+
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async
+#pragma acc wait
+
+  for (i = 0; i < N; i++)
+    {
+      if (a[i] != 3.0)
+	abort ();
+
+      if (b[i] != 3.0)
+	abort ();
+    }
+
+  for (i = 0; i < N; i++)
+    {
+      a[i] = 2.0;
+      b[i] = 0.0;
+    }
+
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1)
+#pragma acc parallel async (1) present (N, a[0:N], b[0:N])
+#pragma acc loop
+  for (i = 0; i < N; i++)
+    b[i] = a[i];
+
+#pragma acc update host (a[0:N], b[0:N]) wait (1) async (1)
+#pragma acc wait (1)
+
+  for (i = 0; i < N; i++)
+    {
+      if (a[i] != 2.0)
+	abort ();
+
+      if (b[i] != 2.0)
+	abort ();
+    }
+
+  for (i = 0; i < N; i++)
+    {
+      a[i] = 3.0;
+      b[i] = 0.0;
+      c[i] = 0.0;
+      d[i] = 0.0;
+    }
+
+#pragma acc update device (a[0:N], b[0:N])
+#pragma acc enter data copyin (c[0:N], d[0:N]) async (1)
+
+#pragma acc parallel present (N, a[0:N], b[0:N]) async (1) wait (1)
+#pragma acc loop 
+  for (i = 0; i < N; i++)
+    b[i] = (a[i] * a[i] * a[i]) / a[i];
+
+#pragma acc parallel present (N, a[0:N], c[0:N]) async (2) wait (1)
+#pragma acc loop 
+  for (i = 0; i < N; i++)
+    c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
+
+#pragma acc parallel present (N, a[0:N], d[0:N]) async (3) wait (1)
+#pragma acc loop 
+  for (i = 0; i < N; i++)
+    d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
+
+#pragma acc exit data copyout (a[0:N], b[0:N], c[0:N], d[0:N]) wait (1, 2, 3) async (1)
+#pragma acc wait (1)
+
+  for (i = 0; i < N; i++)
+    {
+      if (a[i] != 3.0)
+	abort ();
+
+      if (b[i] != 9.0)
+	abort ();
+
+      if (c[i] != 4.0)
+	abort ();
+
+      if (d[i] != 1.0)
+	abort ();
+    }
+
+  for (i = 0; i < N; i++)
+    {
+      a[i] = 2.0;
+      b[i] = 0.0;
+      c[i] = 0.0;
+      d[i] = 0.0;
+      e[i] = 0.0;
+    }
+
+#pragma acc enter data copyin (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) async (1)
+
+#pragma acc parallel loop present (N, a[0:N], b[0:N]) async (1) wait (1) 
+  for (int ii = 0; ii < N; ii++)
+    b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
+
+#pragma acc parallel loop present (N, a[0:N], c[0:N]) async (2) wait (1) 
+  for (int ii = 0; ii < N; ii++)
+    c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
+
+#pragma acc parallel loop present (N, a[0:N], d[0:N]) async (3) wait (1) 
+  for (int ii = 0; ii < N; ii++)
+    d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
+
+#pragma acc parallel loop present (N, a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) wait (1) async (4) 
+  for (int ii = 0; ii < N; ii++)
+    e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
+
+#pragma acc exit data copyout (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) wait (1, 2, 3, 4) async (1) delete (N)
+#pragma acc wait (1)
+
+  for (i = 0; i < N; i++)
+    {
+      if (a[i] != 2.0)
+	abort ();
+
+      if (b[i] != 4.0)
+	abort ();
+
+      if (c[i] != 4.0)
+	abort ();
+
+      if (d[i] != 1.0)
+	abort ();
+
+      if (e[i] != 11.0)
+	abort ();
+    }
+}
+
+int
+main (int argc, char **argv)
+{
+  int N = 128; //1024 * 1024;
+  float *a, *b, *c, *d, *e;
+  int i;
+  int nbytes;
+
+  nbytes = N * sizeof (float);
+
+  a = (float *) malloc (nbytes);
+  b = (float *) malloc (nbytes);
+  c = (float *) malloc (nbytes);
+  d = (float *) malloc (nbytes);
+  e = (float *) malloc (nbytes);
+
+  test (N, a, b, c, d, e);
+
+  free (a);
+  free (b);
+  free (c);
+  free (d);
+  free (e);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data-templates.C b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data-templates.C
new file mode 100644
index 0000000..03b2b01
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data-templates.C
@@ -0,0 +1,114 @@
+// 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;
+};
+
+template <class t1, class t2>
+void
+kernels_present (t1 &d, t2 &x)
+{
+#pragma acc kernels present (d, x) default (none)
+  {
+    d.v = x;
+  }
+}
+
+template <class t1, class t2>
+void
+parallel_present (t1 &d, t2 &x)
+{
+#pragma acc parallel present (d, x) default (none)
+  {
+    d.v = x;
+  }
+}
+
+template <class t1, class t2>
+void
+kernels_implicit (t1 &d, t2 &x)
+{
+#pragma acc kernels
+  {
+    d.v = x;
+  }
+}
+
+template <class t1, class t2>
+void
+parallel_implicit (t1 &d, t2 &x)
+{
+#pragma acc parallel
+  {
+    d.v = x;
+  }
+}
+
+template <class t1, class t2>
+void
+reference_data (t1 &d, t2 &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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c++/update-reference.C b/libgomp/testsuite/libgomp.oacc-c++/update-reference.C
new file mode 100644
index 0000000..bd85d13
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/update-reference.C
@@ -0,0 +1,300 @@
+/* Copy of update-1.c with self exchanged with host for #pragma acc update.
+   This exercises references types.  */
+
+/* { dg-do run } */
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+#include <string.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdbool.h>
+
+int
+test (int &N, float *(&a), float *(&b), float *(&c), float *(&d_a),
+      float *(&d_b), float *(&d_c), int &i)
+{
+
+    
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 3.0;
+        b[i] = 0.0;
+    }
+
+    acc_map_data (a, d_a, N * sizeof (float));
+    acc_map_data (b, d_b, N * sizeof (float));
+    acc_map_data (c, d_c, N * sizeof (float));
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != 3.0)
+            abort ();
+
+        if (b[i] != 3.0)
+            abort ();
+    }
+
+    if (!acc_is_present (&a[0], (N * sizeof (float))))
+      abort ();
+
+    if (!acc_is_present (&b[0], (N * sizeof (float))))
+      abort ();
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 5.0;
+        b[i] = 1.0;
+    }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != 5.0)
+            abort ();
+
+        if (b[i] != 5.0)
+            abort ();
+    }
+
+    if (!acc_is_present (&a[0], (N * sizeof (float))))
+      abort ();
+
+    if (!acc_is_present (&b[0], (N * sizeof (float))))
+      abort ();
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 5.0;
+        b[i] = 1.0;
+    }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update host (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != 5.0)
+            abort ();
+
+        if (b[i] != 5.0)
+            abort ();
+    }
+
+    if (!acc_is_present (&a[0], (N * sizeof (float))))
+      abort ();
+
+    if (!acc_is_present (&b[0], (N * sizeof (float))))
+      abort ();
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 6.0;
+        b[i] = 0.0;
+    }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 9.0;
+    }
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+
+        if (b[i] != 6.0)
+            abort ();
+    }
+
+    if (!acc_is_present (&a[0], (N * sizeof (float))))
+      abort ();
+
+    if (!acc_is_present (&b[0], (N * sizeof (float))))
+      abort ();
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 7.0;
+        b[i] = 2.0;
+    }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 9.0;
+    }
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != 7.0)
+            abort ();
+
+        if (b[i] != 7.0)
+            abort ();
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 9.0;
+    }
+
+#pragma acc update device (a[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != 9.0)
+            abort ();
+
+        if (b[i] != 9.0)
+            abort ();
+    }
+
+    if (!acc_is_present (&a[0], (N * sizeof (float))))
+      abort ();
+
+    if (!acc_is_present (&b[0], (N * sizeof (float))))
+      abort ();
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 5.0;
+    }
+
+#pragma acc update device (a[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 6.0;
+    }
+
+#pragma acc update device (a[0:N >> 1])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+    for (i = 0; i < (N >> 1); i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+
+        if (b[i] != 6.0)
+            abort ();
+    }
+
+    for (i = (N >> 1); i < N; i++)
+    {
+        if (a[i] != 5.0)
+            abort ();
+
+        if (b[i] != 5.0)
+            abort ();
+    }
+
+    if (!acc_is_present (&a[0], (N * sizeof (float))))
+      abort ();
+
+    if (!acc_is_present (&b[0], (N * sizeof (float))))
+      abort ();
+}
+
+int
+main (int argc, char **argv)
+{
+    int N = 8;
+    float *a, *b, *c;
+    float *d_a, *d_b, *d_c;
+    int i;
+
+    a = (float *) malloc (N * sizeof (float));
+    b = (float *) malloc (N * sizeof (float));
+    c = (float *) malloc (N * sizeof (float));
+
+    d_a = (float *) acc_malloc (N * sizeof (float));
+    d_b = (float *) acc_malloc (N * sizeof (float));
+    d_c = (float *) acc_malloc (N * sizeof (float));
+
+    test (N, a, b, c, d_a, d_b, d_c, i);
+
+    acc_free (d_a);
+    acc_free (d_b);
+    acc_free (d_c);
+
+    free (a);
+    free (b);
+    free (c);
+    
+    return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c++/update-template.C b/libgomp/testsuite/libgomp.oacc-c++/update-template.C
new file mode 100644
index 0000000..ae21e73
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/update-template.C
@@ -0,0 +1,301 @@
+/* Copy of update-1.c with self exchanged with host for #pragma acc update.
+   This exercises templates.  */
+
+/* { dg-do run } */
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+#include <string.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdbool.h>
+
+template<typename T1, typename T2>
+int
+test (T1 &N, T2 *(&a), T2 *(&b), T2 *(&c), T2 *(&d_a), T2 *(&d_b),
+      T2 *(&d_c), T1 &i)
+{
+
+    
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 3.0;
+        b[i] = 0.0;
+    }
+
+    acc_map_data (a, d_a, N * sizeof (float));
+    acc_map_data (b, d_b, N * sizeof (float));
+    acc_map_data (c, d_c, N * sizeof (float));
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != 3.0)
+            abort ();
+
+        if (b[i] != 3.0)
+            abort ();
+    }
+
+    if (!acc_is_present (&a[0], (N * sizeof (float))))
+      abort ();
+
+    if (!acc_is_present (&b[0], (N * sizeof (float))))
+      abort ();
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 5.0;
+        b[i] = 1.0;
+    }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != 5.0)
+            abort ();
+
+        if (b[i] != 5.0)
+            abort ();
+    }
+
+    if (!acc_is_present (&a[0], (N * sizeof (float))))
+      abort ();
+
+    if (!acc_is_present (&b[0], (N * sizeof (float))))
+      abort ();
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 5.0;
+        b[i] = 1.0;
+    }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update host (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != 5.0)
+            abort ();
+
+        if (b[i] != 5.0)
+            abort ();
+    }
+
+    if (!acc_is_present (&a[0], (N * sizeof (float))))
+      abort ();
+
+    if (!acc_is_present (&b[0], (N * sizeof (float))))
+      abort ();
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 6.0;
+        b[i] = 0.0;
+    }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 9.0;
+    }
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+
+        if (b[i] != 6.0)
+            abort ();
+    }
+
+    if (!acc_is_present (&a[0], (N * sizeof (float))))
+      abort ();
+
+    if (!acc_is_present (&b[0], (N * sizeof (float))))
+      abort ();
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 7.0;
+        b[i] = 2.0;
+    }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 9.0;
+    }
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != 7.0)
+            abort ();
+
+        if (b[i] != 7.0)
+            abort ();
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 9.0;
+    }
+
+#pragma acc update device (a[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != 9.0)
+            abort ();
+
+        if (b[i] != 9.0)
+            abort ();
+    }
+
+    if (!acc_is_present (&a[0], (N * sizeof (float))))
+      abort ();
+
+    if (!acc_is_present (&b[0], (N * sizeof (float))))
+      abort ();
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 5.0;
+    }
+
+#pragma acc update device (a[0:N])
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 6.0;
+    }
+
+#pragma acc update device (a[0:N >> 1])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            b[ii] = a[ii];
+    }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+    for (i = 0; i < (N >> 1); i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+
+        if (b[i] != 6.0)
+            abort ();
+    }
+
+    for (i = (N >> 1); i < N; i++)
+    {
+        if (a[i] != 5.0)
+            abort ();
+
+        if (b[i] != 5.0)
+            abort ();
+    }
+
+    if (!acc_is_present (&a[0], (N * sizeof (float))))
+      abort ();
+
+    if (!acc_is_present (&b[0], (N * sizeof (float))))
+      abort ();
+}
+
+int
+main (int argc, char **argv)
+{
+    int N = 8;
+    float *a, *b, *c;
+    float *d_a, *d_b, *d_c;
+    int i;
+
+    a = (float *) malloc (N * sizeof (float));
+    b = (float *) malloc (N * sizeof (float));
+    c = (float *) malloc (N * sizeof (float));
+
+    d_a = (float *) acc_malloc (N * sizeof (float));
+    d_b = (float *) acc_malloc (N * sizeof (float));
+    d_c = (float *) acc_malloc (N * sizeof (float));
+
+    test (N, a, b, c, d_a, d_b, d_c, i);
+
+    acc_free (d_a);
+    acc_free (d_b);
+    acc_free (d_c);
+
+    free (a);
+    free (b);
+    free (c);
+    
+    return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c
index 22cef6d..67c382c 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c
@@ -1,3 +1,5 @@
+/* Test acc_set_cuda_stream integration with async and wait.  */
+
 /* { dg-do run { target openacc_nvidia_accel_selected } } */
 /* { dg-additional-options "-lcuda" } */
 
@@ -37,7 +39,7 @@ main (int argc, char **argv)
 #pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N)
     {
 
-#pragma acc parallel async
+#pragma acc parallel async present (a[0:N], b[0:N], N)
     {
         int ii;
 
@@ -67,7 +69,7 @@ main (int argc, char **argv)
 #pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N)
     {
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
     {
         int ii;
 
@@ -99,7 +101,7 @@ main (int argc, char **argv)
 #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N)
     {
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
     {
         int ii;
 
@@ -107,7 +109,7 @@ main (int argc, char **argv)
             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
     }
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], c[0:N], N)
     {
         int ii;
 
@@ -116,7 +118,7 @@ main (int argc, char **argv)
     }
 
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], d[0:N], N)
     {
         int ii;
 
@@ -155,7 +157,7 @@ main (int argc, char **argv)
 #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N)
     {
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
     {
         int ii;
 
@@ -163,7 +165,7 @@ main (int argc, char **argv)
             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
     }
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], c[0:N], N)
     {
         int ii;
 
@@ -171,7 +173,7 @@ main (int argc, char **argv)
             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
     }
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], d[0:N], N)
     {
         int ii;
 
@@ -179,7 +181,8 @@ main (int argc, char **argv)
             d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
     }
 
-#pragma acc parallel wait (1) async (1)
+#pragma acc parallel wait (1) async (1) present (a[0:N], b[0:N], c[0:N]) \
+  present (d[0:N], e[0:N], N)
     {
         int ii;
 
@@ -228,7 +231,7 @@ main (int argc, char **argv)
 #pragma acc data copy (a[0:N], b[0:N]) copyin (N)
     {
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
     {
         int ii;
 
@@ -257,10 +260,11 @@ main (int argc, char **argv)
         d[i] = 0.0;
     }
 
-#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N)
+#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) \
+  copyin (N)
     {
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
     {
         int ii;
 
@@ -268,7 +272,7 @@ main (int argc, char **argv)
             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
     }
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], c[0:N], N)
     {
         int ii;
 
@@ -276,7 +280,7 @@ main (int argc, char **argv)
             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
     }
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], d[0:N], N)
     {
         int ii;
 
@@ -315,7 +319,7 @@ main (int argc, char **argv)
 #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N)
     {
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
     {
         int ii;
 
@@ -323,7 +327,7 @@ main (int argc, char **argv)
             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
     }
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], c[0:N], N)
     {
         int ii;
 
@@ -331,7 +335,7 @@ main (int argc, char **argv)
             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
     }
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], d[0:N], N)
     {
         int ii;
 
@@ -339,7 +343,8 @@ main (int argc, char **argv)
             d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
     }
 
-#pragma acc parallel wait (1) async (1)
+#pragma acc parallel wait (1) async (1) present (a[0:N], b[0:N], c[0:N]) \
+  present (d[0:N], e[0:N], N)
     {
         int ii;
 
@@ -381,7 +386,7 @@ main (int argc, char **argv)
 #pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N)
     {
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
     {
         int ii;
 
@@ -389,7 +394,7 @@ main (int argc, char **argv)
             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
     }
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], c[0:N], N)
     {
         int ii;
 
@@ -426,7 +431,7 @@ main (int argc, char **argv)
 #pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N)
     {
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
     {
         int ii;
 
@@ -434,7 +439,7 @@ main (int argc, char **argv)
             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
     }
 
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], c[0:N], N)
     {
         int ii;
 
@@ -462,5 +467,11 @@ main (int argc, char **argv)
 
     acc_shutdown (acc_device_nvidia);
 
+    free (a);
+    free (b);
+    free (c);
+    free (d);
+    free (e);
+
     return 0;
 }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
index f867a66..bbdaabe 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
@@ -1,3 +1,5 @@
+/* Test "acc enter data" and "acc exit data".  */
+
 /* { dg-do run } */
 
 #include <stdlib.h>
@@ -25,7 +27,7 @@ main (int argc, char **argv)
     }
 
 #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
-#pragma acc parallel async wait
+#pragma acc parallel async wait present (a[0:N], b[0:N], N)
 #pragma acc loop
   for (i = 0; i < N; i++)
     b[i] = a[i];
@@ -41,7 +43,7 @@ main (int argc, char **argv)
       if (b[i] != 3.0)
 	abort ();
     }
-
+  return 0;
   for (i = 0; i < N; i++)
     {
       a[i] = 2.0;
@@ -49,7 +51,7 @@ main (int argc, char **argv)
     }
 
 #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1)
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
 #pragma acc loop
   for (i = 0; i < N; i++)
     b[i] = a[i];
@@ -74,24 +76,26 @@ main (int argc, char **argv)
       d[i] = 0.0;
     }
 
-#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (N) async (1)
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) \
+  copyin (d[0:N]) copyin (N) async (1)
 
-#pragma acc parallel async (1) wait (1)
+#pragma acc parallel async (1) wait (1) present (a[0:N], b[0:N], N)
 #pragma acc loop
   for (i = 0; i < N; i++)
     b[i] = (a[i] * a[i] * a[i]) / a[i];
 
-#pragma acc parallel async (2) wait (1)
+#pragma acc parallel async (2) wait (1) present (a[0:N], c[0:N], N)
 #pragma acc loop
   for (i = 0; i < N; i++)
     c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
 
-#pragma acc parallel async (3) wait (1)
+#pragma acc parallel async (3) wait (1) present (a[0:N], d[0:N], N)
 #pragma acc loop
   for (i = 0; i < N; i++)
     d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
 
-#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) wait (1, 2, 3) async (1)
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \
+  copyout (d[0:N]) wait (1, 2, 3) async (1)
 #pragma acc wait (1)
 
   for (i = 0; i < N; i++)
@@ -118,28 +122,30 @@ main (int argc, char **argv)
       e[i] = 0.0;
     }
 
-#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (e[0:N]) copyin (N) async (1)
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) \
+  copyin (d[0:N]) copyin (e[0:N]) copyin (N) async (1)
 
-#pragma acc parallel async (1) wait (1)
+#pragma acc parallel async (1) wait (1) present (a[0:N], b[0:N], N)
   for (int ii = 0; ii < N; ii++)
     b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 
-#pragma acc parallel async (2) wait (1)
+#pragma acc parallel async (2) wait (1) present (a[0:N], c[0:N], N)
   for (int ii = 0; ii < N; ii++)
     c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 
-#pragma acc parallel async (3) wait (1)
+#pragma acc parallel async (3) wait (1) present (a[0:N], d[0:N], N)
   for (int ii = 0; ii < N; ii++)
     d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
 
-#pragma acc parallel wait (1) async (4)
+#pragma acc parallel wait (1) async (4) present (a[0:N], b[0:N], c[0:N]) \
+  present (d[0:N], e[0:N], N)
   for (int ii = 0; ii < N; ii++)
     e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
 
-#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \
+  copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) delete (N)
 #pragma acc wait (1)
 
-
   for (i = 0; i < N; i++)
     {
       if (a[i] != 2.0)
@@ -158,5 +164,11 @@ main (int argc, char **argv)
 	abort ();
     }
 
+  free (a);
+  free (b);
+  free (c);
+  free (d);
+  free (e);
+
   return 0;
 }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
index 747109f..014eb85 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
@@ -1,3 +1,5 @@
+/* Test "acc enter data", "acc exit data" and "acc update".  */
+
 /* { dg-do run } */
 
 #include <stdlib.h>
@@ -25,7 +27,7 @@ main (int argc, char **argv)
     }
 
 #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
-#pragma acc parallel async wait
+#pragma acc parallel async wait present (a[0:N], b[0:N], N)
 #pragma acc loop
   for (i = 0; i < N; i++)
     b[i] = a[i];
@@ -49,7 +51,7 @@ main (int argc, char **argv)
     }
 
 #pragma acc update device (a[0:N], b[0:N]) async (1)
-#pragma acc parallel async (1)
+#pragma acc parallel async (1) present (a[0:N], b[0:N], N)
 #pragma acc loop
   for (i = 0; i < N; i++)
     b[i] = a[i];
@@ -78,17 +80,17 @@ main (int argc, char **argv)
 #pragma acc update device (b[0:N]) async (2)
 #pragma acc enter data copyin (c[0:N], d[0:N]) async (3)
 
-#pragma acc parallel async (1) wait (1,2)
+#pragma acc parallel async (1) wait (1,2) present (a[0:N], b[0:N], N)
 #pragma acc loop
   for (i = 0; i < N; i++)
     b[i] = (a[i] * a[i] * a[i]) / a[i];
 
-#pragma acc parallel async (2) wait (1,3)
+#pragma acc parallel async (2) wait (1,3) present (a[0:N], c[0:N], N)
 #pragma acc loop
   for (i = 0; i < N; i++)
     c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
 
-#pragma acc parallel async (3) wait (1,3)
+#pragma acc parallel async (3) wait (1,3) present (a[0:N], d[0:N], N)
 #pragma acc loop
   for (i = 0; i < N; i++)
     d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
@@ -123,27 +125,28 @@ main (int argc, char **argv)
 #pragma acc update device (a[0:N], b[0:N], c[0:N], d[0:N]) async (1)
 #pragma acc enter data copyin (e[0:N]) async (5)
 
-#pragma acc parallel async (1) wait (1)
+#pragma acc parallel async (1) wait (1) present (a[0:N], b[0:N], N)
   for (int ii = 0; ii < N; ii++)
     b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 
-#pragma acc parallel async (2) wait (1)
+#pragma acc parallel async (2) wait (1) present (a[0:N], c[0:N], N)
   for (int ii = 0; ii < N; ii++)
     c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 
-#pragma acc parallel async (3) wait (1)
+#pragma acc parallel async (3) wait (1) present (a[0:N], d[0:N], N)
   for (int ii = 0; ii < N; ii++)
     d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
 
-#pragma acc parallel wait (1,5) async (4)
+#pragma acc parallel wait (1,5) async (4) present (a[0:N], b[0:N], c[0:N]) \
+  present (d[0:N], e[0:N], N)
   for (int ii = 0; ii < N; ii++)
     e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
 
-#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \
+  copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
 #pragma acc exit data delete (N)
 #pragma acc wait (1)
 
-
   for (i = 0; i < N; i++)
     {
       if (a[i] != 2.0)
@@ -162,5 +165,11 @@ main (int argc, char **argv)
 	abort ();
     }
 
+  free (a);
+  free (b);
+  free (c);
+  free (d);
+  free (e);
+
   return 0;
 }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
index 6aa3bb7..cb6e59c 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
@@ -457,7 +457,7 @@ main(int argc, char **argv)
 
 #pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(1)
 {
-#pragma acc parallel present(a[0:N])
+#pragma acc parallel present(a[0:N], b[0:N])
     {
         int ii;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c
index ededf2b..86dd491 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c
@@ -1,3 +1,5 @@
+/* Test implicit data clauses inside data regions.  */
+
 /* { dg-do run } */
 /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
 
@@ -11,20 +13,16 @@ int
 main (int argc, char **argv)
 {
     int N = 8;
-    float *a, *b, *c, *d;
+    float a[N], b[N], c[N], *d;
     int i;
 
-    a = (float *) malloc (N * sizeof (float));
-    b = (float *) malloc (N * sizeof (float));
-    c = (float *) malloc (N * sizeof (float));
-
     for (i = 0; i < N; i++)
     {
         a[i] = 3.0;
         b[i] = 0.0;
     }
 
-#pragma acc data copyin (a[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) copyout (b)
     {
 #pragma acc parallel
         {
@@ -53,7 +51,7 @@ main (int argc, char **argv)
         b[i] = 1.0;
     }
 
-#pragma acc data copyin (a[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) copyout (b)
     {
 #pragma acc parallel
         {
@@ -89,7 +87,7 @@ main (int argc, char **argv)
         a[i] = 9.0;
     }
 
-#pragma acc data present_or_copyin (a[0:N]) copyout (b[0:N])
+#pragma acc data present_or_copyin (a) copyout (b)
     {
 #pragma acc parallel
         {
@@ -120,7 +118,7 @@ main (int argc, char **argv)
         b[i] = 0.0;
     }
 
-#pragma acc data copyin (a[0:N]) present_or_copyout (b[0:N])
+#pragma acc data copyin (a) present_or_copyout (b)
     {
 #pragma acc parallel
         {
@@ -151,7 +149,7 @@ main (int argc, char **argv)
 
     d = (float *) acc_copyin (&b[0], N * sizeof (float));
 
-#pragma acc data copyin (a[0:N]) present_or_copyout (b[0:N])
+#pragma acc data copyin (a) present_or_copyout (b)
     {
 #pragma acc parallel
         {
@@ -188,7 +186,7 @@ main (int argc, char **argv)
         b[i] = 4.0;
     }
 
-#pragma acc data copy (a[0:N]) copyout (b[0:N])
+#pragma acc data copy (a) copyout (b)
     {
 #pragma acc parallel
         {
@@ -223,7 +221,7 @@ main (int argc, char **argv)
         b[i] = 7.0;
     }
 
-#pragma acc data present_or_copy (a[0:N]) present_or_copy (b[0:N])
+#pragma acc data present_or_copy (a) present_or_copy (b)
     {
 #pragma acc parallel
         {
@@ -261,7 +259,7 @@ main (int argc, char **argv)
     d = (float *) acc_copyin (&a[0], N * sizeof (float));
     d = (float *) acc_copyin (&b[0], N * sizeof (float));
 
-#pragma acc data present_or_copy (a[0:N]) present_or_copy (b[0:N])
+#pragma acc data present_or_copy (a) present_or_copy (b)
     {
 #pragma acc parallel
         {
@@ -305,7 +303,7 @@ main (int argc, char **argv)
         b[i] = 7.0;
     }
 
-#pragma acc data copyin (a[0:N]) create (c[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) create (c) copyout (b)
     {
 #pragma acc parallel
         {
@@ -343,7 +341,7 @@ main (int argc, char **argv)
         b[i] = 8.0;
     }
 
-#pragma acc data copyin (a[0:N]) present_or_create (c[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) present_or_create (c) copyout (b)
     {
 #pragma acc parallel
         {
@@ -384,7 +382,7 @@ main (int argc, char **argv)
     d = (float *) acc_malloc (N * sizeof (float));
     acc_map_data (c, d, N * sizeof (float));
 
-#pragma acc data copyin (a[0:N]) present_or_create (c[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) present_or_create (c) copyout (b)
     {
 #pragma acc parallel
         {
@@ -431,7 +429,7 @@ main (int argc, char **argv)
     d = (float *) acc_malloc (N * sizeof (float));
     acc_map_data (c, d, N * sizeof (float));
 
-#pragma acc data copyin (a[0:N]) present (c[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) present (c) copyout (b)
     {
 #pragma acc parallel
         {
@@ -488,7 +486,7 @@ main (int argc, char **argv)
     if (!acc_is_present (a, (N * sizeof (float))))
       abort ();
 
-#pragma acc data present (a[0:N]) present (c[0:N]) present (b[0:N])
+#pragma acc data present (a) present (c) present (b)
     {
 #pragma acc parallel
         {
@@ -543,7 +541,7 @@ main (int argc, char **argv)
 
     d = (float *) acc_malloc (N * sizeof (float));
 
-#pragma acc parallel copyin (a[0:N]) deviceptr (d) copyout (b[0:N])
+#pragma acc parallel copyin (a) deviceptr (d) copyout (b)
     {
         int ii;
 
@@ -584,7 +582,7 @@ main (int argc, char **argv)
         a[i] = 9.0;
     }
 
-#pragma acc data pcopyin (a[0:N]) copyout (b[0:N])
+#pragma acc data pcopyin (a) copyout (b)
     {
 #pragma acc parallel
         {
@@ -615,7 +613,7 @@ main (int argc, char **argv)
         b[i] = 0.0;
     }
 
-#pragma acc data copyin (a[0:N]) pcopyout (b[0:N])
+#pragma acc data copyin (a) pcopyout (b)
     {
 #pragma acc parallel
         {
@@ -644,7 +642,7 @@ main (int argc, char **argv)
         b[i] = 7.0;
     }
 
-#pragma acc data copyin (a[0:N]) pcreate (c[0:N]) copyout (b[0:N])
+#pragma acc data copyin (a) pcreate (c) copyout (b)
     {
 #pragma acc parallel
         {
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/present-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/present-2.c
index 41efa70..0c03bad 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/present-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/present-2.c
@@ -1,3 +1,5 @@
+/* Test both explicit and implicitly present data.  */
+
 /* { dg-do run } */
 /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
 
@@ -8,11 +10,10 @@ int
 main (int argc, char **argv)
 {
   int N = 8;
-  float *a, *b;
+  float *a, b[N];
   int i;
 
   a = (float *) malloc (N * sizeof (float));
-  b = (float *) malloc (N * sizeof (float));
 
   for (i = 0; i < N; i++)
     {
@@ -20,7 +21,7 @@ main (int argc, char **argv)
       b[i] = 0.0;
     }
 
-#pragma acc data copyin(a[0:N]) copyout(b[0:N])
+#pragma acc data copyin(a[0:N]) copy (b)
   {
 
 #pragma acc parallel present(a[0:N])
@@ -44,5 +45,7 @@ main (int argc, char **argv)
 	abort ();
     }
 
+  free (a);
+
   return 0;
 }
