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 <[email protected]>
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(-)
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
--
2.27.0
