From bed993884b149851fe930b43cf11cbcdf05f1578 Mon Sep 17 00:00:00 2001
From: Chung-Lin Tang <cltang@codesourcery.com>
Date: Tue, 6 Jun 2023 03:46:29 -0700
Subject: [PATCH] OpenACC 2.7: default clause support for data constructs

This patch implements the OpenACC 2.7 addition of default(none|present) support
for data constructs.

Now, specifying "default(none|present)" on a data construct turns on same
default clause behavior for all lexically enclosed compute constructs (which
don't already themselves have a default clause).

gcc/c/ChangeLog:
	* c-parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.

gcc/cp/ChangeLog:
	* parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.

gcc/fortran/ChangeLog:
	* openmp.cc (OACC_DATA_CLAUSES): Add OMP_CLAUSE_DEFAULT.

gcc/ChangeLog:
	* gimplify.cc (oacc_region_type_name): New function.
	(oacc_default_clause): If no 'default' clause appears on this
	compute construct, see if one appears on a lexically containing
	'data' construct.
	(gimplify_scan_omp_clauses): Upon OMP_CLAUSE_DEFAULT case, set
	ctx->oacc_default_clause_ctx to current context.

gcc/testsuite/ChangeLog:
	* c-c++-common/goacc/default-3.c: Adjust testcase.
	* c-c++-common/goacc/default-4.c: Adjust testcase.
	* c-c++-common/goacc/default-5.c: Adjust testcase.
	* gfortran.dg/goacc/default-3.f95: Adjust testcase.
	* gfortran.dg/goacc/default-4.f: Adjust testcase.
	* gfortran.dg/goacc/default-5.f: Adjust testcase.

Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
---
 gcc/c/c-parser.cc                             |  1 +
 gcc/cp/parser.cc                              |  1 +
 gcc/fortran/openmp.cc                         |  3 +-
 gcc/gimplify.cc                               | 64 +++++++++++----
 gcc/testsuite/c-c++-common/goacc/default-3.c  | 59 +++++++++++++-
 gcc/testsuite/c-c++-common/goacc/default-4.c  | 42 ++++++++++
 gcc/testsuite/c-c++-common/goacc/default-5.c  | 19 ++++-
 gcc/testsuite/gfortran.dg/goacc/default-3.f95 | 77 ++++++++++++++++++-
 gcc/testsuite/gfortran.dg/goacc/default-4.f   | 36 +++++++++
 gcc/testsuite/gfortran.dg/goacc/default-5.f   | 19 ++++-
 10 files changed, 298 insertions(+), 23 deletions(-)

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index cabb18d04f7..33fe7b115ff 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -18284,6 +18284,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser)
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE)		\
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 7f646704d3f..774706ac607 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -45854,6 +45854,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 234d896b2ce..bee3015e484 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -3802,7 +3802,8 @@ error:
 #define OACC_DATA_CLAUSES \
   (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR  | OMP_CLAUSE_COPY	      \
    | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE		      \
-   | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_ATTACH)
+   | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_ATTACH	      \
+   | OMP_CLAUSE_DEFAULT)
 #define OACC_LOOP_CLAUSES \
   (omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER	      \
    | OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT	      \
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 320920ed74c..7549436944c 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -7699,6 +7699,25 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl,
   return flags;
 }
 
+/* Return string name for types of OpenACC constructs from ORT_* values.  */
+
+static const char *
+oacc_region_type_name (enum omp_region_type region_type)
+{
+  switch (region_type)
+    {
+    case ORT_ACC_DATA:
+      return "data";
+    case ORT_ACC_PARALLEL:
+      return "parallel";
+    case ORT_ACC_KERNELS:
+      return "kernels";
+    case ORT_ACC_SERIAL:
+      return "serial";
+    default:
+      gcc_unreachable ();
+    }
+}
 
 /* Determine outer default flags for DECL mentioned in an OACC region
    but not declared in an enclosing clause.  */
@@ -7706,7 +7725,23 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl,
 static unsigned
 oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
 {
-  const char *rkind;
+  struct gimplify_omp_ctx *ctx_default = ctx;
+  /* If no 'default' clause appears on this compute construct...  */
+  if (ctx_default->default_kind == OMP_CLAUSE_DEFAULT_SHARED)
+    {
+      /* ..., see if one appears on a lexically containing 'data'
+	 construct.  */
+      while ((ctx_default = ctx_default->outer_context))
+	{
+	  if (ctx_default->region_type == ORT_ACC_DATA
+	      && ctx_default->default_kind != OMP_CLAUSE_DEFAULT_SHARED)
+	    break;
+	}
+      /* If not, reset.  */
+      if (!ctx_default)
+	ctx_default = ctx;
+    }
+
   bool on_device = false;
   bool is_private = false;
   bool declared = is_oacc_declared (decl);
@@ -7738,14 +7773,12 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
   switch (ctx->region_type)
     {
     case ORT_ACC_KERNELS:
-      rkind = "kernels";
-
       if (is_private)
 	flags |= GOVD_FIRSTPRIVATE;
       else if (AGGREGATE_TYPE_P (type))
 	{
 	  /* Aggregates default to 'present_or_copy', or 'present'.  */
-	  if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
+	  if (ctx_default->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
 	    flags |= GOVD_MAP;
 	  else
 	    flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
@@ -7758,8 +7791,6 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
 
     case ORT_ACC_PARALLEL:
     case ORT_ACC_SERIAL:
-      rkind = ctx->region_type == ORT_ACC_PARALLEL ? "parallel" : "serial";
-
       if (is_private)
 	flags |= GOVD_FIRSTPRIVATE;
       else if (on_device || declared)
@@ -7767,7 +7798,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
       else if (AGGREGATE_TYPE_P (type))
 	{
 	  /* Aggregates default to 'present_or_copy', or 'present'.  */
-	  if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
+	  if (ctx_default->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
 	    flags |= GOVD_MAP;
 	  else
 	    flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
@@ -7785,16 +7816,23 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
   if (DECL_ARTIFICIAL (decl))
     ; /* We can get compiler-generated decls, and should not complain
 	 about them.  */
-  else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_NONE)
+  else if (ctx_default->default_kind == OMP_CLAUSE_DEFAULT_NONE)
     {
       error ("%qE not specified in enclosing OpenACC %qs construct",
-	     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)
+	     DECL_NAME (lang_hooks.decls.omp_report_decl (decl)),
+	     oacc_region_type_name (ctx->region_type));
+      if (ctx_default != ctx)
+	inform (ctx->location, "enclosing OpenACC %qs construct and",
+		oacc_region_type_name (ctx->region_type));
+      inform (ctx_default->location,
+	      "enclosing OpenACC %qs construct with %qs clause",
+	      oacc_region_type_name (ctx_default->region_type),
+	      "default(none)");
+    }
+  else if (ctx_default->default_kind == OMP_CLAUSE_DEFAULT_PRESENT)
     ; /* Handled above.  */
   else
-    gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED);
+    gcc_checking_assert (ctx_default->default_kind == OMP_CLAUSE_DEFAULT_SHARED);
 
   return flags;
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/default-3.c b/gcc/testsuite/c-c++-common/goacc/default-3.c
index ac169a903e9..73dbc908475 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/default-3.c
@@ -4,13 +4,66 @@ void f1 ()
 {
   int f1_a = 2;
   float f1_b[2];
-  
-#pragma acc kernels default (none) /* { dg-message "enclosing OpenACC .kernels. construct" } */
+
+#pragma acc kernels default (none) /* { dg-note "enclosing OpenACC 'kernels' construct with 'default\\\(none\\\)' clause" } */
+  {
+    f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" } */
+      = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .kernels. construct" } */
+  }
+#pragma acc parallel default (none) /* { dg-note "enclosing OpenACC 'parallel' construct with 'default\\\(none\\\)' clause" } */
+  {
+    f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
+      = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
+  }
+
+#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */
+#pragma acc kernels /* { dg-note "enclosing OpenACC 'kernels' construct and" } */
   {
     f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" } */
       = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .kernels. construct" } */
   }
-#pragma acc parallel default (none) /* { dg-message "enclosing OpenACC .parallel. construct" } */
+#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */
+#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */
+  {
+    f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
+      = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
+  }
+
+#pragma acc data default (none)
+#pragma acc parallel default (none) /* { dg-note "enclosing OpenACC 'parallel' construct with 'default\\\(none\\\)' clause" } */
+  {
+    f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
+      = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
+  }
+
+#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */
+#pragma acc data
+#pragma acc data
+#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */
+  {
+    f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
+      = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
+  }
+#pragma acc data
+#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */
+#pragma acc data
+#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */
+  {
+    f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
+      = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
+  }
+#pragma acc data
+#pragma acc data
+#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */
+#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */
+  {
+    f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
+      = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
+  }
+#pragma acc data
+#pragma acc data default (none)
+#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */
+#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */
   {
     f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
       = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
diff --git a/gcc/testsuite/c-c++-common/goacc/default-4.c b/gcc/testsuite/c-c++-common/goacc/default-4.c
index 867175d4847..e12cb86d097 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/default-4.c
@@ -44,6 +44,27 @@ void f2 ()
   }
 }
 
+void f2_ ()
+{
+  int f2__a = 2;
+  float f2__b[2];
+
+#pragma acc data default (none) copyin (f2__a) copyout (f2__b)
+  /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f2__b \[^\\)\]+\\) map\\(to:f2__a \[^\\)\]+\\) default\\(none\\)" 1 "gimple" } } */
+  {
+#pragma acc kernels
+    /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f2__b \[^\\)\]+\\) map\\(tofrom:f2__a" 1 "gimple" } } */
+    {
+      f2__b[0] = f2__a;
+    }
+#pragma acc parallel
+    /* { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(tofrom:f2__b \[^\\)\]+\\) map\\(tofrom:f2__a" 1 "gimple" } } */
+    {
+      f2__b[0] = f2__a;
+    }
+  }
+}
+
 void f3 ()
 {
   int f3_a = 2;
@@ -64,3 +85,24 @@ void f3 ()
     }
   }
 }
+
+void f3_ ()
+{
+  int f3__a = 2;
+  float f3__b[2];
+
+#pragma acc data default (present) copyin (f3__a) copyout (f3__b)
+  /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f3__b \[^\\)\]+\\) map\\(to:f3__a \[^\\)\]+\\) default\\(present\\)" 1 "gimple" } } */
+  {
+#pragma acc kernels
+    /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f3__b \[^\\)\]+\\) map\\(tofrom:f3__a" 1 "gimple" } } */
+    {
+      f3__b[0] = f3__a;
+    }
+#pragma acc parallel
+    /* { dg-final { scan-tree-dump-times "omp target oacc_parallel 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
index 37e3c3555cd..59ac1d79668 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/default-5.c
@@ -4,8 +4,8 @@
 
 void f1 ()
 {
-  int f1_a = 2;
-  float f1_b[2];
+  int f1_a = 2, f1_c = 3;
+  float f1_b[2], f1_d[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" } } */
@@ -17,4 +17,19 @@ void f1 ()
   {
     f1_b[0] = f1_a;
   }
+
+  /* { dg-final { scan-tree-dump-times "omp target oacc_data default\\(present\\)" 2 "gimple" } } */
+#pragma acc data default (present)
+#pragma acc kernels
+  /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(force_present:f1_d \[^\\)\]+\\) map\\(force_tofrom:f1_c" 1 "gimple" } } */
+  {
+    f1_d[0] = f1_c;
+  }
+#pragma acc data default (none)
+#pragma acc data default (present)
+#pragma acc parallel
+  /* { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:f1_d \[^\\)\]+\\) firstprivate\\(f1_c\\)" 1 "gimple" } } */
+  {
+    f1_d[0] = f1_c;
+  }
 }
diff --git a/gcc/testsuite/gfortran.dg/goacc/default-3.f95 b/gcc/testsuite/gfortran.dg/goacc/default-3.f95
index 98ed34200c6..c1edf4c8137 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/default-3.f95
@@ -5,14 +5,87 @@ subroutine f1
   integer :: f1_a = 2
   real, dimension (2) :: f1_b
 
-  !$acc kernels default (none) ! { dg-message "enclosing OpenACC .kernels. construct" }
+  !$acc kernels default (none) ! { dg-note "enclosing OpenACC .kernels. construct with 'default\\\(none\\\)' clause" }
   f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" "" { xfail *-*-* } }
        = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .kernels. construct" }
   ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .kernels. construct" "" { xfail *-*-* } .-1 }
   !$acc end kernels
-  !$acc parallel default (none) ! { dg-message "enclosing OpenACC .parallel. construct" }
+  !$acc parallel default (none) ! { dg-note "enclosing OpenACC .parallel. construct with 'default\\\(none\\\)' clause" }
   f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } }
        = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
   ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
   !$acc end parallel
+
+  !$acc data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" }
+  !$acc kernels ! { dg-note "enclosing OpenACC 'kernels' construct and" }
+  f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" "" { xfail *-*-* } }
+       = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .kernels. construct" }
+  ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .kernels. construct" "" { xfail *-*-* } .-1 }
+  !$acc end kernels
+  !$acc end data
+
+  !$acc data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" }
+  !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" }
+  f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } }
+       = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
+  ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
+  !$acc end parallel
+  !$acc end data
+
+  !$acc data default (none)
+  !$acc parallel default (none) ! { dg-note "enclosing OpenACC .parallel. construct with 'default\\\(none\\\)' clause" }
+  f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } }
+       = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
+  ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
+  !$acc end parallel
+  !$acc end data
+
+  !$acc data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" }
+  !$acc data
+  !$acc data
+  !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" }
+  f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } }
+       = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
+  ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
+  !$acc end parallel
+  !$acc end data
+  !$acc end data
+  !$acc end data
+
+  !$acc data
+  !$acc data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" }
+  !$acc data
+  !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" }
+  f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } }
+       = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
+  ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
+  !$acc end parallel
+  !$acc end data
+  !$acc end data
+  !$acc end data
+
+  !$acc data
+  !$acc data
+  !$acc data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" }
+  !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" }
+  f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } }
+       = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
+  ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
+  !$acc end parallel
+  !$acc end data
+  !$acc end data
+  !$acc end data
+
+  !$acc data
+  !$acc data default (none)
+  !$acc data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" }
+  !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" }
+  f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } }
+       = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
+  ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
+  !$acc end parallel
+  !$acc end data
+  !$acc end data
+  !$acc end data
+
 end subroutine f1
diff --git a/gcc/testsuite/gfortran.dg/goacc/default-4.f b/gcc/testsuite/gfortran.dg/goacc/default-4.f
index 30f411f70ab..4e89b6859ba 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-4.f
+++ b/gcc/testsuite/gfortran.dg/goacc/default-4.f
@@ -38,6 +38,24 @@
 !$ACC END DATA
       END SUBROUTINE F2
 
+      SUBROUTINE F2_
+      IMPLICIT NONE
+      INTEGER :: F2__A = 2
+      REAL, DIMENSION (2) :: F2__B
+
+!$ACC DATA DEFAULT (NONE) COPYIN (F2__A) COPYOUT (F2__B)
+! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f2__a \[^\\)\]+\\) map\\(from:f2__b \[^\\)\]+\\) default\\(none\\)" 1 "gimple" } }
+!$ACC KERNELS
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f2__b \[^\\)\]+\\) map\\(tofrom:f2__a" 1 "gimple" } }
+      F2__B(1) = F2__A;
+!$ACC END KERNELS
+!$ACC PARALLEL
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(tofrom:f2__b \[^\\)\]+\\) map\\(tofrom:f2__a" 1 "gimple" } }
+      F2__B(1) = F2__A;
+!$ACC END PARALLEL
+!$ACC END DATA
+      END SUBROUTINE F2_
+
       SUBROUTINE F3
       IMPLICIT NONE
       INTEGER :: F3_A = 2
@@ -55,3 +73,21 @@
 !$ACC END PARALLEL
 !$ACC END DATA
       END SUBROUTINE F3
+
+      SUBROUTINE F3_
+      IMPLICIT NONE
+      INTEGER :: F3__A = 2
+      REAL, DIMENSION (2) :: F3__B
+
+!$ACC DATA DEFAULT (PRESENT) COPYIN (F3__A) COPYOUT (F3__B)
+! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f3__a \[^\\)\]+\\) map\\(from:f3__b \[^\\)\]+\\) default\\(present\\)" 1 "gimple" } }
+!$ACC KERNELS
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f3__b \[^\\)\]+\\) map\\(tofrom:f3__a" 1 "gimple" } }
+      F3__B(1) = F3__A;
+!$ACC END KERNELS
+!$ACC PARALLEL
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel 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
index 9dc83cbe601..2cb07a8cbca 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-5.f
+++ b/gcc/testsuite/gfortran.dg/goacc/default-5.f
@@ -4,8 +4,8 @@
 
       SUBROUTINE F1
       IMPLICIT NONE
-      INTEGER :: F1_A = 2
-      REAL, DIMENSION (2) :: F1_B
+      INTEGER :: F1_A = 2, F1_C = 3
+      REAL, DIMENSION (2) :: F1_B, F1_D
 
 !$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" } }
@@ -15,4 +15,19 @@
 ! { 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
+
+!$ACC DATA DEFAULT (PRESENT)
+!$ACC KERNELS
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(force_present:f1_d \[^\\)\]+\\) map\\(force_tofrom:f1_c" 1 "gimple" } }
+      F1_D(1) = F1_C;
+!$ACC END KERNELS
+!$ACC END DATA
+!$ACC DATA DEFAULT (NONE)
+!$ACC DATA DEFAULT (PRESENT)
+!$ACC PARALLEL DEFAULT (PRESENT)
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_d \[^\\)\]+\\) firstprivate\\(f1_c\\)" 1 "gimple" } }
+      F1_D(1) = F1_C;
+!$ACC END PARALLEL
+!$ACC END DATA
+!$ACC END DATA
       END SUBROUTINE F1
-- 
2.34.1

