From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 19 May 2017 13:32:48 +0000 (+0200)
Subject: OpenACC 2.5 default (present) clause
X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=7fd549d24fda05c859fb17697c51c16886902dad;p=gcc.git

OpenACC 2.5 default (present) clause

	gcc/c/
	* c-parser.c (c_parser_omp_clause_default): Handle
	"OMP_CLAUSE_DEFAULT_PRESENT".
	gcc/cp/
	* parser.c (cp_parser_omp_clause_default): Handle
	"OMP_CLAUSE_DEFAULT_PRESENT".
	gcc/fortran/
	* gfortran.h (enum gfc_omp_default_sharing): Add
	"OMP_DEFAULT_PRESENT".
	* dump-parse-tree.c (show_omp_clauses): Handle it.
	* openmp.c (gfc_match_omp_clauses): Likewise.
	* trans-openmp.c (gfc_trans_omp_clauses): Likewise.
	gcc/
	* tree-core.h (enum omp_clause_default_kind): Add
	"OMP_CLAUSE_DEFAULT_PRESENT".
	* tree-pretty-print.c (dump_omp_clause): Handle it.
	* gimplify.c (enum gimplify_omp_var_data): Add
	"GOVD_MAP_FORCE_PRESENT".
	(gimplify_adjust_omp_clauses_1): Map it to
	"GOMP_MAP_FORCE_PRESENT".
	(oacc_default_clause): Handle "OMP_CLAUSE_DEFAULT_PRESENT".
	gcc/testsuite/
	* c-c++-common/goacc/default-1.c: Update.
	* c-c++-common/goacc/default-2.c: Likewise.
	* c-c++-common/goacc/default-4.c: Likewise.
	* gfortran.dg/goacc/default-1.f95: Likewise.
	* gfortran.dg/goacc/default-4.f: Likewise.
	* c-c++-common/goacc/default-5.c: New file.
	* gfortran.dg/goacc/default-5.f: Likewise.
	libgomp/
	* testsuite/libgomp.oacc-c++/template-reduction.C: Update.
	* testsuite/libgomp.oacc-c-c++-common/nested-2.c: Update.
	* testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/default-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/non-scalar-data.f90: Likewise.

From-SVN: r248280
---

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 8f63902d064..c40af474ea3 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,5 +1,14 @@
 2017-05-19  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* tree-core.h (enum omp_clause_default_kind): Add
+	"OMP_CLAUSE_DEFAULT_PRESENT".
+	* tree-pretty-print.c (dump_omp_clause): Handle it.
+	* gimplify.c (enum gimplify_omp_var_data): Add
+	"GOVD_MAP_FORCE_PRESENT".
+	(gimplify_adjust_omp_clauses_1): Map it to
+	"GOMP_MAP_FORCE_PRESENT".
+	(oacc_default_clause): Handle "OMP_CLAUSE_DEFAULT_PRESENT".
+
 	* gimplify.c (oacc_default_clause): Clarify.
 
 2017-05-19  Nathan Sidwell  <nathan@acm.org>
diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 79643cc1b3b..a70eaf910b0 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,3 +1,8 @@
+2017-05-19  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-parser.c (c_parser_omp_clause_default): Handle
+	"OMP_CLAUSE_DEFAULT_PRESENT".
+
 2017-05-18  Bernd Edlinger  <bernd.edlinger@hotmail.de>
 
 	* config-lang.in (gtfiles): Add c-family/c-format.c.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 96c074961da..2e013161a9a 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11057,10 +11057,10 @@ c_parser_omp_clause_copyprivate (c_parser *parser, tree list)
 }
 
 /* OpenMP 2.5:
-   default ( shared | none )
+   default ( none | shared )
 
-   OpenACC 2.0:
-   default (none) */
+   OpenACC:
+   default ( none | present ) */
 
 static tree
 c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc)
@@ -11083,6 +11083,12 @@ c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc)
 	  kind = OMP_CLAUSE_DEFAULT_NONE;
 	  break;
 
+	case 'p':
+	  if (strcmp ("present", p) != 0 || !is_oacc)
+	    goto invalid_kind;
+	  kind = OMP_CLAUSE_DEFAULT_PRESENT;
+	  break;
+
 	case 's':
 	  if (strcmp ("shared", p) != 0 || is_oacc)
 	    goto invalid_kind;
@@ -11099,7 +11105,7 @@ c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc)
     {
     invalid_kind:
       if (is_oacc)
-	c_parser_error (parser, "expected %<none%>");
+	c_parser_error (parser, "expected %<none%> or %<present%>");
       else
 	c_parser_error (parser, "expected %<none%> or %<shared%>");
     }
diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index c89f71981e3..f2dced41ddb 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,5 +1,8 @@
 2017-05-19  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* parser.c (cp_parser_omp_clause_default): Handle
+	"OMP_CLAUSE_DEFAULT_PRESENT".
+
 	* parser.c (cp_parser_omp_clause_default): Avoid printing more
 	than one syntax error message.
 
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 6453397b530..b345110fbc4 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -31456,10 +31456,10 @@ cp_parser_omp_clause_collapse (cp_parser *parser, tree list, location_t location
 }
 
 /* OpenMP 2.5:
-   default ( shared | none )
+   default ( none | shared )
 
-   OpenACC 2.0
-   default (none) */
+   OpenACC:
+   default ( none | present ) */
 
 static tree
 cp_parser_omp_clause_default (cp_parser *parser, tree list,
@@ -31483,6 +31483,12 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list,
 	  kind = OMP_CLAUSE_DEFAULT_NONE;
 	  break;
 
+	case 'p':
+	  if (strcmp ("present", p) != 0 || !is_oacc)
+	    goto invalid_kind;
+	  kind = OMP_CLAUSE_DEFAULT_PRESENT;
+	  break;
+
 	case 's':
 	  if (strcmp ("shared", p) != 0 || is_oacc)
 	    goto invalid_kind;
@@ -31499,7 +31505,7 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list,
     {
     invalid_kind:
       if (is_oacc)
-	cp_parser_error (parser, "expected %<none%>");
+	cp_parser_error (parser, "expected %<none%> or %<present%>");
       else
 	cp_parser_error (parser, "expected %<none%> or %<shared%>");
     }
diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog
index 9a8b3e1c396..4b7f1f4e4bf 100644
--- a/gcc/fortran/ChangeLog
+++ b/gcc/fortran/ChangeLog
@@ -1,3 +1,11 @@
+2017-05-19  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* gfortran.h (enum gfc_omp_default_sharing): Add
+	"OMP_DEFAULT_PRESENT".
+	* dump-parse-tree.c (show_omp_clauses): Handle it.
+	* openmp.c (gfc_match_omp_clauses): Likewise.
+	* trans-openmp.c (gfc_trans_omp_clauses): Likewise.
+
 2017-05-18  Fritz Reese <fritzoreese@gmail.com>
 
 	PR fortran/79968
diff --git a/gcc/fortran/dump-parse-tree.c b/gcc/fortran/dump-parse-tree.c
index 87a530458f9..49b23d8697e 100644
--- a/gcc/fortran/dump-parse-tree.c
+++ b/gcc/fortran/dump-parse-tree.c
@@ -1283,6 +1283,7 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses)
 	case OMP_DEFAULT_PRIVATE: type = "PRIVATE"; break;
 	case OMP_DEFAULT_SHARED: type = "SHARED"; break;
 	case OMP_DEFAULT_FIRSTPRIVATE: type = "FIRSTPRIVATE"; break;
+	case OMP_DEFAULT_PRESENT: type = "PRESENT"; break;
 	default:
 	  gcc_unreachable ();
 	}
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 293655078a4..26b89bee98e 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1241,7 +1241,8 @@ enum gfc_omp_default_sharing
   OMP_DEFAULT_NONE,
   OMP_DEFAULT_PRIVATE,
   OMP_DEFAULT_SHARED,
-  OMP_DEFAULT_FIRSTPRIVATE
+  OMP_DEFAULT_FIRSTPRIVATE,
+  OMP_DEFAULT_PRESENT
 };
 
 enum gfc_omp_proc_bind_kind
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 89eecfa2ed1..80146e2f1dc 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -1080,13 +1080,19 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	      if (gfc_match ("default ( none )") == MATCH_YES)
 		c->default_sharing = OMP_DEFAULT_NONE;
 	      else if (openacc)
-		/* c->default_sharing = OMP_DEFAULT_UNKNOWN */;
-	      else if (gfc_match ("default ( shared )") == MATCH_YES)
-		c->default_sharing = OMP_DEFAULT_SHARED;
-	      else if (gfc_match ("default ( private )") == MATCH_YES)
-		c->default_sharing = OMP_DEFAULT_PRIVATE;
-	      else if (gfc_match ("default ( firstprivate )") == MATCH_YES)
-		c->default_sharing = OMP_DEFAULT_FIRSTPRIVATE;
+		{
+		  if (gfc_match ("default ( present )") == MATCH_YES)
+		    c->default_sharing = OMP_DEFAULT_PRESENT;
+		}
+	      else
+		{
+		  if (gfc_match ("default ( firstprivate )") == MATCH_YES)
+		    c->default_sharing = OMP_DEFAULT_FIRSTPRIVATE;
+		  else if (gfc_match ("default ( private )") == MATCH_YES)
+		    c->default_sharing = OMP_DEFAULT_PRIVATE;
+		  else if (gfc_match ("default ( shared )") == MATCH_YES)
+		    c->default_sharing = OMP_DEFAULT_SHARED;
+		}
 	      if (c->default_sharing != OMP_DEFAULT_UNKNOWN)
 		continue;
 	    }
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 662036f514d..1d254c6904d 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2564,6 +2564,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	case OMP_DEFAULT_FIRSTPRIVATE:
 	  OMP_CLAUSE_DEFAULT_KIND (c) = OMP_CLAUSE_DEFAULT_FIRSTPRIVATE;
 	  break;
+	case OMP_DEFAULT_PRESENT:
+	  OMP_CLAUSE_DEFAULT_KIND (c) = OMP_CLAUSE_DEFAULT_PRESENT;
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 0c02ee4371e..810d9f4a3e1 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -99,6 +99,9 @@ enum gimplify_omp_var_data
   /* Flag for GOVD_MAP, if it is a forced mapping.  */
   GOVD_MAP_FORCE = 262144,
 
+  /* Flag for GOVD_MAP: must be present already.  */
+  GOVD_MAP_FORCE_PRESENT = 524288,
+
   GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
 			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
 			   | GOVD_LOCAL)
@@ -6956,8 +6959,13 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
       rkind = "kernels";
 
       if (AGGREGATE_TYPE_P (type))
-	/* Aggregates default to 'present_or_copy'.  */
-	flags |= GOVD_MAP;
+	{
+	  /* Aggregates default to 'present_or_copy', or 'present'.  */
+	  if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
+	    flags |= GOVD_MAP;
+	  else
+	    flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
+	}
       else
 	/* Scalars default to 'copy'.  */
 	flags |= GOVD_MAP | GOVD_MAP_FORCE;
@@ -6970,8 +6978,13 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
       if (on_device || declared)
 	flags |= GOVD_MAP;
       else if (AGGREGATE_TYPE_P (type))
-	/* Aggregates default to 'present_or_copy'.  */
-	flags |= GOVD_MAP;
+	{
+	  /* Aggregates default to 'present_or_copy', or 'present'.  */
+	  if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
+	    flags |= GOVD_MAP;
+	  else
+	    flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
+	}
       else
 	/* Scalars default to 'firstprivate'.  */
 	flags |= GOVD_FIRSTPRIVATE;
@@ -6991,6 +7004,8 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
 	     DECL_NAME (lang_hooks.decls.omp_report_decl (decl)), rkind);
       inform (ctx->location, "enclosing OpenACC %qs construct", rkind);
     }
+  else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_PRESENT)
+    ; /* Handled above.  */
   else
     gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED);
 
@@ -8708,11 +8723,30 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
     }
   else if (code == OMP_CLAUSE_MAP)
     {
-      int kind = (flags & GOVD_MAP_TO_ONLY
-		  ? GOMP_MAP_TO
-		  : GOMP_MAP_TOFROM);
-      if (flags & GOVD_MAP_FORCE)
-	kind |= GOMP_MAP_FLAG_FORCE;
+      int kind;
+      /* Not all combinations of these GOVD_MAP flags are actually valid.  */
+      switch (flags & (GOVD_MAP_TO_ONLY
+		       | GOVD_MAP_FORCE
+		       | GOVD_MAP_FORCE_PRESENT))
+	{
+	case 0:
+	  kind = GOMP_MAP_TOFROM;
+	  break;
+	case GOVD_MAP_FORCE:
+	  kind = GOMP_MAP_TOFROM | GOMP_MAP_FLAG_FORCE;
+	  break;
+	case GOVD_MAP_TO_ONLY:
+	  kind = GOMP_MAP_TO;
+	  break;
+	case GOVD_MAP_TO_ONLY | GOVD_MAP_FORCE:
+	  kind = GOMP_MAP_TO | GOMP_MAP_FLAG_FORCE;
+	  break;
+	case GOVD_MAP_FORCE_PRESENT:
+	  kind = GOMP_MAP_FORCE_PRESENT;
+	  break;
+	default:
+	  gcc_unreachable ();
+	}
       OMP_CLAUSE_SET_MAP_KIND (clause, kind);
       if (DECL_SIZE (decl)
 	  && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index ca901c29a66..e8e2df5c9a9 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,5 +1,13 @@
 2017-05-19  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* c-c++-common/goacc/default-1.c: Update.
+	* c-c++-common/goacc/default-2.c: Likewise.
+	* c-c++-common/goacc/default-4.c: Likewise.
+	* gfortran.dg/goacc/default-1.f95: Likewise.
+	* gfortran.dg/goacc/default-4.f: Likewise.
+	* c-c++-common/goacc/default-5.c: New file.
+	* gfortran.dg/goacc/default-5.f: Likewise.
+
 	* c-c++-common/goacc/default-1.c: New file.
 	* c-c++-common/goacc/default-2.c: Likewise.
 	* c-c++-common/goacc/data-default-1.c: Remove file, including its
diff --git a/gcc/testsuite/c-c++-common/goacc/default-1.c b/gcc/testsuite/c-c++-common/goacc/default-1.c
index 4d31dbc981b..23c25ab45a9 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/default-1.c
@@ -6,4 +6,9 @@ void f1 ()
   ;
 #pragma acc parallel default (none)
   ;
+
+#pragma acc kernels default (present)
+  ;
+#pragma acc parallel default (present)
+  ;
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/default-2.c b/gcc/testsuite/c-c++-common/goacc/default-2.c
index d6b6cdce4d7..c0cdd1263bd 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/default-2.c
@@ -7,39 +7,39 @@ void f1 ()
 #pragma acc parallel default /* { dg-error "expected .\\(. before end of line" } */
   ;
 
-#pragma acc kernels default ( /* { dg-error "expected .none. before end of line" } */
+#pragma acc kernels default ( /* { dg-error "expected .none. or .present. before end of line" } */
   ;
-#pragma acc parallel default ( /* { dg-error "expected .none. before end of line" } */
+#pragma acc parallel default ( /* { dg-error "expected .none. or .present. before end of line" } */
   ;
 
-#pragma acc kernels default (, /* { dg-error "expected .none. before .,. token" } */
+#pragma acc kernels default (, /* { dg-error "expected .none. or .present. before .,. token" } */
   ;
-#pragma acc parallel default (, /* { dg-error "expected .none. before .,. token" } */
+#pragma acc parallel default (, /* { dg-error "expected .none. or .present. before .,. token" } */
   ;
 
-#pragma acc kernels default () /* { dg-error "expected .none. before .\\). token" } */
+#pragma acc kernels default () /* { dg-error "expected .none. or .present. before .\\). token" } */
   ;
-#pragma acc parallel default () /* { dg-error "expected .none. before .\\). token" } */
+#pragma acc parallel default () /* { dg-error "expected .none. or .present. before .\\). token" } */
   ;
 
-#pragma acc kernels default (,) /* { dg-error "expected .none. before .,. token" } */
+#pragma acc kernels default (,) /* { dg-error "expected .none. or .present. before .,. token" } */
   ;
-#pragma acc parallel default (,) /* { dg-error "expected .none. before .,. token" } */
+#pragma acc parallel default (,) /* { dg-error "expected .none. or .present. before .,. token" } */
   ;
 
-#pragma acc kernels default (firstprivate) /* { dg-error "expected .none. before .firstprivate." } */
+#pragma acc kernels default (firstprivate) /* { dg-error "expected .none. or .present. before .firstprivate." } */
   ;
-#pragma acc parallel default (firstprivate) /* { dg-error "expected .none. before .firstprivate." } */
+#pragma acc parallel default (firstprivate) /* { dg-error "expected .none. or .present. before .firstprivate." } */
   ;
 
-#pragma acc kernels default (private) /* { dg-error "expected .none. before .private." } */
+#pragma acc kernels default (private) /* { dg-error "expected .none. or .present. before .private." } */
   ;
-#pragma acc parallel default (private) /* { dg-error "expected .none. before .private." } */
+#pragma acc parallel default (private) /* { dg-error "expected .none. or .present. before .private." } */
   ;
 
-#pragma acc kernels default (shared) /* { dg-error "expected .none. before .shared." } */
+#pragma acc kernels default (shared) /* { dg-error "expected .none. or .present. before .shared." } */
   ;
-#pragma acc parallel default (shared) /* { dg-error "expected .none. before .shared." } */
+#pragma acc parallel default (shared) /* { dg-error "expected .none. or .present. before .shared." } */
   ;
 
 #pragma acc kernels default (none /* { dg-error "expected .\\). before end of line" } */
diff --git a/gcc/testsuite/c-c++-common/goacc/default-4.c b/gcc/testsuite/c-c++-common/goacc/default-4.c
index 787b35256db..dfa79bbbe6e 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/default-4.c
@@ -43,3 +43,24 @@ void f2 ()
     }
   }
 }
+
+void f3 ()
+{
+  int f3_a = 2;
+  float f3_b[2];
+
+#pragma acc data copyin (f3_a) copyout (f3_b)
+  /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f3_b \[^\\)\]+\\) map\\(force_to:f3_a" 1 "gimple" } } */
+  {
+#pragma acc kernels default (present)
+    /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */
+    {
+      f3_b[0] = f3_a;
+    }
+#pragma acc parallel default (present)
+    /* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */
+    {
+      f3_b[0] = f3_a;
+    }
+  }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/default-5.c b/gcc/testsuite/c-c++-common/goacc/default-5.c
new file mode 100644
index 00000000000..37e3c3555cd
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/default-5.c
@@ -0,0 +1,20 @@
+/* OpenACC default (present) clause.  */
+
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void f1 ()
+{
+  int f1_a = 2;
+  float f1_b[2];
+
+#pragma acc kernels default (present)
+  /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_tofrom:f1_a" 1 "gimple" } } */
+  {
+    f1_b[0] = f1_a;
+  }
+#pragma acc parallel default (present)
+  /* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) firstprivate\\(f1_a\\)" 1 "gimple" } } */
+  {
+    f1_b[0] = f1_a;
+  }
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/default-1.f95 b/gcc/testsuite/gfortran.dg/goacc/default-1.f95
index a79b444f6f4..41fa9446b91 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/default-1.f95
@@ -7,4 +7,9 @@ subroutine f1
   !$acc end kernels
   !$acc parallel default (none)
   !$acc end parallel
+
+  !$acc kernels default (present)
+  !$acc end kernels
+  !$acc parallel default (present)
+  !$acc end parallel
 end subroutine f1
diff --git a/gcc/testsuite/gfortran.dg/goacc/default-4.f b/gcc/testsuite/gfortran.dg/goacc/default-4.f
index b2f73c3bcfb..77291f43eff 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-4.f
+++ b/gcc/testsuite/gfortran.dg/goacc/default-4.f
@@ -37,3 +37,21 @@
 !$ACC END PARALLEL
 !$ACC END DATA
       END SUBROUTINE F2
+
+      SUBROUTINE F3
+      IMPLICIT NONE
+      INTEGER :: F3_A = 2
+      REAL, DIMENSION (2) :: F3_B
+
+!$ACC DATA COPYIN (F3_A) COPYOUT (F3_B)
+! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f3_a \[^\\)\]+\\) map\\(force_from:f3_b" 1 "gimple" } }
+!$ACC KERNELS DEFAULT (PRESENT)
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } }
+      F3_B(1) = F3_A;
+!$ACC END KERNELS
+!$ACC PARALLEL DEFAULT (PRESENT)
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } }
+      F3_B(1) = F3_A;
+!$ACC END PARALLEL
+!$ACC END DATA
+      END SUBROUTINE F3
diff --git a/gcc/testsuite/gfortran.dg/goacc/default-5.f b/gcc/testsuite/gfortran.dg/goacc/default-5.f
new file mode 100644
index 00000000000..9dc83cbe601
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/default-5.f
@@ -0,0 +1,18 @@
+! OpenACC default (present) clause.
+
+! { dg-additional-options "-fdump-tree-gimple" } 
+
+      SUBROUTINE F1
+      IMPLICIT NONE
+      INTEGER :: F1_A = 2
+      REAL, DIMENSION (2) :: F1_B
+
+!$ACC KERNELS DEFAULT (PRESENT)
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_tofrom:f1_a" 1 "gimple" } }
+      F1_B(1) = F1_A;
+!$ACC END KERNELS
+!$ACC PARALLEL DEFAULT (PRESENT)
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) firstprivate\\(f1_a\\)" 1 "gimple" } }
+      F1_B(1) = F1_A;
+!$ACC END PARALLEL
+      END SUBROUTINE F1
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index c76fc7bc8ef..ea73477163d 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -357,7 +357,7 @@ enum omp_clause_code {
   /* OpenMP clause: ordered [(constant-integer-expression)].  */
   OMP_CLAUSE_ORDERED,
 
-  /* OpenMP clause: default.  */
+  /* OpenACC/OpenMP clause: default.  */
   OMP_CLAUSE_DEFAULT,
 
   /* OpenACC/OpenMP clause: collapse (constant-integer-expression).  */
@@ -499,6 +499,7 @@ enum omp_clause_default_kind {
   OMP_CLAUSE_DEFAULT_NONE,
   OMP_CLAUSE_DEFAULT_PRIVATE,
   OMP_CLAUSE_DEFAULT_FIRSTPRIVATE,
+  OMP_CLAUSE_DEFAULT_PRESENT,
   OMP_CLAUSE_DEFAULT_LAST
 };
 
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index ec28b1e3c95..b70e32573ee 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -502,6 +502,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	case OMP_CLAUSE_DEFAULT_FIRSTPRIVATE:
 	  pp_string (pp, "firstprivate");
 	  break;
+	case OMP_CLAUSE_DEFAULT_PRESENT:
+	  pp_string (pp, "present");
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 84d1c839d70..460605b1b8a 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,5 +1,11 @@
 2017-05-19  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* testsuite/libgomp.oacc-c++/template-reduction.C: Update.
+	* testsuite/libgomp.oacc-c-c++-common/nested-2.c: Update.
+	* testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise.
+	* testsuite/libgomp.oacc-fortran/default-1.f90: Likewise.
+	* testsuite/libgomp.oacc-fortran/non-scalar-data.f90: Likewise.
+
 	* plugin/plugin-hsa.c (DLSYM_FN, init_hsa_runtime_functions):
 	Debug output for failure.
 
diff --git a/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C b/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C
index 6c85fba7a92..ede0aae14c6 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C
@@ -32,6 +32,28 @@ sum ()
   return s;
 }
 
+// Check template with default (present)
+
+template<typename T> T
+sum_default_present ()
+{
+  T s = 0;
+  T array[n];
+
+  for (int i = 0; i < n; i++)
+    array[i] = i+1;
+
+#pragma acc enter data copyin (array)
+
+#pragma acc parallel loop num_gangs (10) gang reduction (+:s) default (present)
+  for (int i = 0; i < n; i++)
+    s += array[i];
+
+#pragma acc exit data delete (array)
+
+  return s;
+}
+
 // Check present and async
 
 template<typename T> T
@@ -86,6 +108,9 @@ main()
   if (sum<int> () != result)
     __builtin_abort ();
 
+  if (sum_default_present<int> () != result)
+    __builtin_abort ();
+
 #pragma acc enter data copyin (a)
   if (async_sum (a) != result)
     __builtin_abort ();
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
index c16459826af..51b3b18ad66 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
@@ -137,5 +137,36 @@ main (int argc, char *argv[])
             abort ();
     }
 
+
+#pragma acc enter data create (a)
+
+#pragma acc parallel default (present)
+    {
+      for (int j = 0; j < N; ++j)
+	a[j] = j + 1;
+    }
+
+#pragma acc update host (a)
+
+    for (i = 0; i < N; ++i)
+      {
+        if (a[i] != i + 1)
+	  abort ();
+      }
+
+#pragma acc kernels default (present)
+    {
+      for (int j = 0; j < N; ++j)
+	a[j] = j + 2;
+    }
+
+#pragma acc exit data copyout (a)
+
+    for (i = 0; i < N; ++i)
+      {
+        if (a[i] != i + 2)
+	  abort ();
+      }
+
     return 0;
 }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
index 16a85980b4d..df4aca06a70 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
@@ -1,4 +1,5 @@
-! Copy of data-4.f90 with self exchanged with host for !acc update.
+! Copy of data-4.f90 with self exchanged with host for !acc update, and with
+! default (present) clauses added.
 
 ! { dg-do run }
 
@@ -19,7 +20,7 @@ program asyncwait
 
   !$acc enter data copyin (a(1:N)) copyin (b(1:N)) copyin (N) async
 
-  !$acc parallel async wait
+  !$acc parallel default (present) async wait
   !$acc loop
   do i = 1, N
      b(i) = a(i)
@@ -39,7 +40,7 @@ program asyncwait
 
   !$acc update device (a(1:N), b(1:N)) async (1)
 
-  !$acc parallel async (1) wait (1)
+  !$acc parallel default (present) async (1) wait (1)
   !$acc loop
   do i = 1, N
      b(i) = a(i)
@@ -62,19 +63,19 @@ program asyncwait
   !$acc enter data copyin (c(1:N), d(1:N)) async (1)
   !$acc update device (a(1:N), b(1:N)) async (1)
 
-  !$acc parallel async (1)
+  !$acc parallel default (present) async (1)
   do i = 1, N
      b(i) = (a(i) * a(i) * a(i)) / a(i)
   end do
   !$acc end parallel
 
-  !$acc parallel async (1)
+  !$acc parallel default (present) async (1)
   do i = 1, N
      c(i) = (a(i) * 4) / a(i)
   end do
   !$acc end parallel
 
-  !$acc parallel async (1)
+  !$acc parallel default (present) async (1)
   do i = 1, N
      d(i) = ((a(i) * a(i)  + a(i)) / a(i)) - a(i)
   end do
@@ -100,25 +101,25 @@ program asyncwait
   !$acc enter data copyin (e(1:N)) async (1)
   !$acc update device (a(1:N), b(1:N), c(1:N), d(1:N)) async (1)
 
-  !$acc parallel async (1)
+  !$acc parallel default (present) async (1)
   do i = 1, N
      b(i) = (a(i) * a(i) * a(i)) / a(i)
   end do
   !$acc end parallel
 
-  !$acc parallel async (1)
+  !$acc parallel default (present) async (1)
   do i = 1, N
      c(i) = (a(i) * 4) / a(i)
   end do
   !$acc end parallel
 
-  !$acc parallel async (1)
+  !$acc parallel default (present) async (1)
   do i = 1, N
      d(i) = ((a(i) * a(i) + a(i)) / a(i)) - a(i)
   end do
   !$acc end parallel
 
-  !$acc parallel wait (1) async (1)
+  !$acc parallel default (present) wait (1) async (1)
   do i = 1, N
      e(i) = a(i) + b(i) + c(i) + d(i)
   end do
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
index 10590892458..d8ceb60b493 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
@@ -51,4 +51,14 @@ program main
 
   if (a .ne. 7.0) call abort
 
+  ! The default (present) clause doesn't affect scalar variables; these will
+  ! still get an implicit copy clause added.
+  !$acc kernels default (present)
+    c = a
+    a = 1.0
+    a = a + c
+  !$acc end kernels
+
+  if (a .ne. 8.0) call abort
+
 end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90 b/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
index 94e4228e7a9..53b2cd060f2 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
@@ -1,5 +1,6 @@
 ! Ensure that a non-scalar dummy arguments which are implicitly used inside
-! offloaded regions are properly mapped using present_or_copy.
+! offloaded regions are properly mapped using present_or_copy, or (default)
+! present.
 
 ! { dg-do run }
 
@@ -12,6 +13,7 @@ program main
   n = size
 
   !$acc data copy(array)
+
   call kernels(array, n)
 
   !$acc update host(array)
@@ -20,12 +22,29 @@ program main
      if (array(i) .ne. i) call abort
   end do
 
+  call kernels_default_present(array, n)
+
+  !$acc update host(array)
+
+  do i = 1, n
+     if (array(i) .ne. i+1) call abort
+  end do
+
   call parallel(array, n)
-  !$acc end data
+
+  !$acc update host(array)
 
   do i = 1, n
      if (array(i) .ne. i+i) call abort
   end do
+
+  call parallel_default_present(array, n)
+
+  !$acc end data
+
+  do i = 1, n
+     if (array(i) .ne. i+i+1) call abort
+  end do
 end program main
 
 subroutine kernels (array, n)
@@ -39,6 +58,16 @@ subroutine kernels (array, n)
   !$acc end kernels
 end subroutine kernels
 
+subroutine kernels_default_present (array, n)
+  integer, dimension (n) :: array
+  integer :: n, i
+
+  !$acc kernels default(present)
+  do i = 1, n
+     array(i) = i+1
+  end do
+  !$acc end kernels
+end subroutine kernels_default_present
 
 subroutine parallel (array, n)
   integer, dimension (n) :: array
@@ -50,3 +79,14 @@ subroutine parallel (array, n)
   end do
   !$acc end parallel
 end subroutine parallel
+
+subroutine parallel_default_present (array, n)
+  integer, dimension (n) :: array
+  integer :: n, i
+
+  !$acc parallel default(present)
+  do i = 1, n
+     array(i) = i+i+1
+  end do
+  !$acc end parallel
+end subroutine parallel_default_present