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 cabb18d04f7f..33fe7b115ffa 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 7f646704d3f4..774706ac6079 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 234d896b2ce2..bee3015e484f 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 320920ed74c0..7549436944c3 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 ac169a903e98..73dbc9084758 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 867175d48473..e12cb86d097b 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 37e3c3555cde..59ac1d79668a 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 98ed34200c6c..c1edf4c81371 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 30f411f70ab6..4e89b6859bad 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 9dc83cbe601d..2cb07a8cbca2 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
--
GitLab