[OpenACC,2.7] Implement default clause support for data constructs

Message ID f2b7a229-0ec8-5404-fea2-4612ffb73bf2@siemens.com
State Accepted
Headers
Series [OpenACC,2.7] Implement default clause support for data constructs |

Checks

Context Check Description
snail/gcc-patch-check success Github commit url

Commit Message

Chung-Lin Tang June 6, 2023, 3:11 p.m. UTC
  Hi Thomas,
this patch implements the OpenACC 2.7 addition of default(none|present) support
for data constructs.

Apart from adjusting the front-ends for allowed clauses masks (for acc data),
mostly implemented in gimplify.

Tested on powerpc64le-linux/nvptx, x86_64-linux/amdgcn tests in progress (expect
no surprises). Is this okay for trunk?

Thanks,
Chung-Lin
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 (struct gimplify_omp_ctx): Add oacc_data_default_kind
	field.
	(new_omp_context): Initialize oacc_data_default_kind field.
	(gimplify_scan_omp_clauses): Set oacc_data_default_kind for data
	constructs. Set ctx->default_kind for compute constructs from
	ctx->oacc_data_default_kind.

gcc/testsuite/ChangeLog:

	* c-c++-common/goacc/default-3.c: Adjust testcase.
	* c-c++-common/goacc/default-5.c: Adjust testcase.
	* gfortran.dg/goacc/default-3.f95: Adjust testcase.
	* gfortran.dg/goacc/default-5.f: Adjust testcase.
From 101305aee9b27c6df00d7c403e469bdf8d7f45a4 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 2/2] 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 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 (struct gimplify_omp_ctx): Add oacc_data_default_kind
	field.
	(new_omp_context): Initialize oacc_data_default_kind field.
	(gimplify_scan_omp_clauses): Set oacc_data_default_kind for data
	constructs. Set ctx->default_kind for compute constructs from
	ctx->oacc_data_default_kind.

gcc/testsuite/ChangeLog:

	* c-c++-common/goacc/default-3.c: Adjust testcase.
	* c-c++-common/goacc/default-5.c: Adjust testcase.
	* gfortran.dg/goacc/default-3.f95: Adjust testcase.
	* gfortran.dg/goacc/default-5.f: Adjust testcase.
---
 gcc/c/c-parser.cc                             |  1 +
 gcc/cp/parser.cc                              |  1 +
 gcc/fortran/openmp.cc                         |  3 ++-
 gcc/gimplify.cc                               | 20 +++++++++++++++++++
 gcc/testsuite/c-c++-common/goacc/default-3.c  | 15 +++++++++++++-
 gcc/testsuite/c-c++-common/goacc/default-5.c  | 18 +++++++++++++++--
 gcc/testsuite/gfortran.dg/goacc/default-3.f95 | 15 ++++++++++++++
 gcc/testsuite/gfortran.dg/goacc/default-5.f   | 17 ++++++++++++++--
 8 files changed, 84 insertions(+), 6 deletions(-)
  

Comments

Thomas Schwinge June 23, 2023, 10:47 a.m. UTC | #1
Hi Chung-Lin!

On 2023-06-06T23:11:55+0800, Chung-Lin Tang <chunglin.tang@siemens.com> wrote:
> this patch implements the OpenACC 2.7 addition of default(none|present) support
> for data constructs.

Thanks!

It wasn't clear to me what is supposed to happen, for example, for:

    #pragma acc data default(none)
    {
      #pragma acc data // no 'default' clause
      {
        #pragma acc parallel

Specifically, does the "no 'default' clause" inner 'data' construct
invalidate the 'default(none)' clause of the outer 'data' construct?

In later revisions of the OpenACC specification, wording for 'default'
clause etc. generally has been changed; for example, OpenACC 3.3,
2.6.2 "Variables with Implicitly Determined Data Attributes" defines:

    *Visible 'default' clause*: The nearest 'default' clause appearing on the compute construct or a lexically containing 'data' construct.

Therefore, in the example above, the 'default(none)' still holds.

> Apart from adjusting the front-ends for allowed clauses masks (for acc data),
> mostly implemented in gimplify.

ACK ('s%mostly%%') -- but a little bit differently, please:

> From 101305aee9b27c6df00d7c403e469bdf8d7f45a4 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 2/2] 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 enclosed compute constructs (which don't
> already themselves have a default clause).

Please say "lexically enclosed" -- it's that only, not any dynamic
extent.

> --- a/gcc/gimplify.cc
> +++ b/gcc/gimplify.cc
> @@ -225,6 +225,7 @@ struct gimplify_omp_ctx
>    vec<tree> loop_iter_var;
>    location_t location;
>    enum omp_clause_default_kind default_kind;
> +  enum omp_clause_default_kind oacc_data_default_kind;
>    enum omp_region_type region_type;
>    enum tree_code code;
>    bool combined_loop;
> @@ -459,6 +460,8 @@ new_omp_context (enum omp_region_type region_type)
>      c->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
>    else
>      c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
> +  if (gimplify_omp_ctxp)
> +    c->oacc_data_default_kind = gimplify_omp_ctxp->oacc_data_default_kind;
>    c->defaultmap[GDMK_SCALAR] = GOVD_MAP;
>    c->defaultmap[GDMK_SCALAR_TARGET] = GOVD_MAP;
>    c->defaultmap[GDMK_AGGREGATE] = GOVD_MAP;

Instead of adding a new 'oacc_data_default_kind' to 'gimplify_omp_ctx',
let's please do this the other way round:

> @@ -12050,6 +12053,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>
>       case OMP_CLAUSE_DEFAULT:
>         ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);

Here, we already preserve 'default' for whichever OMP construct.

> +       if (code == OACC_DATA)
> +         ctx->oacc_data_default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
>         break;
>
>       case OMP_CLAUSE_INCLUSIVE:
> @@ -12098,6 +12103,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>       list_p = &OMP_CLAUSE_CHAIN (c);
>      }
>
> +  if ((code == OACC_PARALLEL
> +       || code == OACC_KERNELS
> +       || code == OACC_SERIAL)
> +      && ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED
> +      && ctx->oacc_data_default_kind != OMP_CLAUSE_DEFAULT_UNSPECIFIED)
> +    {
> +      ctx->default_kind = ctx->oacc_data_default_kind;
> +
> +      /* Append actual default clause on compute construct. Not really needed
> +      for omp_notice_variable to work properly, but for debug dump files.  */
> +      c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_DEFAULT);
> +      OMP_CLAUSE_DEFAULT_KIND (c) = ctx->oacc_data_default_kind;
> +      *list_p = c;
> +    }
> +
>    ctx->clauses = *orig_list_p;
>    gimplify_omp_ctxp = ctx;
>  }

Instead of this, in 'gimplify_omp_workshare', before the
'gimplify_scan_omp_clauses' call, do something like:

    if ((ort & ORT_ACC)
        && !omp_find_clause (OMP_CLAUSES (expr), OMP_CLAUSE_DEFAULT))
      {
        /* Determine effective 'default' clause for OpenACC compute construct.  */
        for (struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context)
          {
            if (ctx->region_type == ORT_ACC_DATA
                && ctx->default_kind != OMP_CLAUSE_DEFAULT_SHARED)
              {
                [Append actual default clause on compute construct.]
                break;
              }
          }
      }

That seems conceptually simpler to me?

For the 'build_omp_clause', does using 'ctx->location' instead of
'UNKNOWN_LOCATION' help diagnostics in any way?  Like if we add in
'gcc/gimplify.cc:oacc_default_clause',
'if (ctx->default_kind == OMP_CLAUSE_DEFAULT_NONE)' another 'inform' to
point to the 'data' construct's 'default' clause?  (But not sure if
that's easily done; otherwise don't.)

Similar to the ones you've already got, please also add a few test cases
for nested 'default' clauses, like:

    #pragma acc data // no vs. 'default(none)' vs. 'default(present)'
    {
      #pragma acc data // no vs. same vs. different 'default' clause
      {
        #pragma acc data // no vs. same vs. different 'default' clause
        {
          #pragma acc parallel

Similarly, test cases where 'default' on the compute construct overrides
'default' of an outer 'data' construct.


Grüße
 Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
  
Chung-Lin Tang July 14, 2023, 10:34 a.m. UTC | #2
Hi Thomas,

On 2023/6/23 6:47 PM, Thomas Schwinge wrote:
>> +
>>    ctx->clauses = *orig_list_p;
>>    gimplify_omp_ctxp = ctx;
>>  }
> Instead of this, in 'gimplify_omp_workshare', before the
> 'gimplify_scan_omp_clauses' call, do something like:
> 
>     if ((ort & ORT_ACC)
>         && !omp_find_clause (OMP_CLAUSES (expr), OMP_CLAUSE_DEFAULT))
>       {
>         /* Determine effective 'default' clause for OpenACC compute construct.  */
>         for (struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context)
>           {
>             if (ctx->region_type == ORT_ACC_DATA
>                 && ctx->default_kind != OMP_CLAUSE_DEFAULT_SHARED)
>               {
>                 [Append actual default clause on compute construct.]
>                 break;
>               }
>           }
>       }
> 
> That seems conceptually simpler to me?

I'm not sure if this is conceptually simpler, but using 'oacc_default_kind'
is definitely faster computationally :)

However, as you mention below...

> For the 'build_omp_clause', does using 'ctx->location' instead of
> 'UNKNOWN_LOCATION' help diagnostics in any way?  Like if we add in
> 'gcc/gimplify.cc:oacc_default_clause',
> 'if (ctx->default_kind == OMP_CLAUSE_DEFAULT_NONE)' another 'inform' to
> point to the 'data' construct's 'default' clause?  (But not sure if
> that's easily done; otherwise don't.)

Noticed that we will need to track the actually lexically enclosing OpenACC construct
with the user set default-clause somewhere in 'ctx', in order to satisfy the current
diagnostics in oacc_default_clause().

(the UNKNOWN_LOCATION for the internally created default-clause probably doesn't
matter, that one is just for reminder in internal dumps, probably never plays role
in user diagnostics)

> Similar to the ones you've already got, please also add a few test cases
> for nested 'default' clauses, like:
> 
>     #pragma acc data // no vs. 'default(none)' vs. 'default(present)'
>     {
>       #pragma acc data // no vs. same vs. different 'default' clause
>       {
>         #pragma acc data // no vs. same vs. different 'default' clause
>         {
>           #pragma acc parallel
> 
> Similarly, test cases where 'default' on the compute construct overrides
> 'default' of an outer 'data' construct.

Okay, will add more testcases.

Thanks,
Chung-Lin
  

Patch

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index b61aef8b1a2..645d28b320d 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -18133,6 +18133,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 dd7638f1c93..4b4df29a406 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -45759,6 +45759,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 4c30548567f..b785e71f20f 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -3645,7 +3645,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 4aa6229fc74..368f5fd7ec8 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -225,6 +225,7 @@  struct gimplify_omp_ctx
   vec<tree> loop_iter_var;
   location_t location;
   enum omp_clause_default_kind default_kind;
+  enum omp_clause_default_kind oacc_data_default_kind;
   enum omp_region_type region_type;
   enum tree_code code;
   bool combined_loop;
@@ -459,6 +460,8 @@  new_omp_context (enum omp_region_type region_type)
     c->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
   else
     c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
+  if (gimplify_omp_ctxp)
+    c->oacc_data_default_kind = gimplify_omp_ctxp->oacc_data_default_kind;
   c->defaultmap[GDMK_SCALAR] = GOVD_MAP;
   c->defaultmap[GDMK_SCALAR_TARGET] = GOVD_MAP;
   c->defaultmap[GDMK_AGGREGATE] = GOVD_MAP;
@@ -12050,6 +12053,8 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 
 	case OMP_CLAUSE_DEFAULT:
 	  ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
+	  if (code == OACC_DATA)
+	    ctx->oacc_data_default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
 	  break;
 
 	case OMP_CLAUSE_INCLUSIVE:
@@ -12098,6 +12103,21 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	list_p = &OMP_CLAUSE_CHAIN (c);
     }
 
+  if ((code == OACC_PARALLEL
+       || code == OACC_KERNELS
+       || code == OACC_SERIAL)
+      && ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED
+      && ctx->oacc_data_default_kind != OMP_CLAUSE_DEFAULT_UNSPECIFIED)
+    {
+      ctx->default_kind = ctx->oacc_data_default_kind;
+
+      /* Append actual default clause on compute construct. Not really needed
+	 for omp_notice_variable to work properly, but for debug dump files.  */
+      c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_DEFAULT);
+      OMP_CLAUSE_DEFAULT_KIND (c) = ctx->oacc_data_default_kind;
+      *list_p = c;
+    }
+
   ctx->clauses = *orig_list_p;
   gimplify_omp_ctxp = ctx;
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/default-3.c b/gcc/testsuite/c-c++-common/goacc/default-3.c
index ac169a903e9..060629057c1 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/default-3.c
@@ -4,7 +4,7 @@  void f1 ()
 {
   int f1_a = 2;
   float f1_b[2];
-  
+
 #pragma acc kernels default (none) /* { dg-message "enclosing OpenACC .kernels. construct" } */
   {
     f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" } */
@@ -15,4 +15,17 @@  void f1 ()
     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 kernels /* { dg-message "enclosing OpenACC .kernels. construct" } */
+  {
+    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 data default (none)
+#pragma acc parallel /* { dg-message "enclosing OpenACC .parallel. construct" } */
+  {
+    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-5.c b/gcc/testsuite/c-c++-common/goacc/default-5.c
index 37e3c3555cd..e44a0dd49d7 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,18 @@  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 default\\(present\\) map\\(force_present:f1_d \[^\\)\]+\\) map\\(force_tofrom:f1_c" 1 "gimple" } } */
+  {
+    f1_d[0] = f1_c;
+  }
+#pragma acc data default (present)
+#pragma acc parallel
+  /* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) 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..13ccce30fd7 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/default-3.f95
@@ -15,4 +15,19 @@  subroutine f1
        = 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)
+  !$acc kernels ! { dg-message "enclosing OpenACC .kernels. construct" }
+  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)
+  !$acc parallel ! { dg-message "enclosing OpenACC .parallel. construct" }
+  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
 end subroutine f1
diff --git a/gcc/testsuite/gfortran.dg/goacc/default-5.f b/gcc/testsuite/gfortran.dg/goacc/default-5.f
index 9dc83cbe601..4424f9e7523 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,17 @@ 
 ! { 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 default\\(present\\) 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 (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
       END SUBROUTINE F1