* [PATCH, OpenACC 2.7] Implement default clause support for data constructs
@ 2023-06-06 15:11 Chung-Lin Tang
2023-06-23 10:47 ` Thomas Schwinge
0 siblings, 1 reply; 5+ messages in thread
From: Chung-Lin Tang @ 2023-06-06 15:11 UTC (permalink / raw)
To: gcc-patches, Thomas Schwinge, Catherine Moore
[-- Attachment #1: Type: text/plain, Size: 1175 bytes --]
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.
[-- Attachment #2: 0002-OpenACC-2.7-default-clause-support-for-data-construc.patch --]
[-- Type: text/plain, Size: 11311 bytes --]
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(-)
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
^ permalink raw reply [flat|nested] 5+ messages in thread
* Re: [PATCH, OpenACC 2.7] Implement default clause support for data constructs
2023-06-06 15:11 [PATCH, OpenACC 2.7] Implement default clause support for data constructs Chung-Lin Tang
@ 2023-06-23 10:47 ` Thomas Schwinge
2023-07-14 10:34 ` Chung-Lin Tang
0 siblings, 1 reply; 5+ messages in thread
From: Thomas Schwinge @ 2023-06-23 10:47 UTC (permalink / raw)
To: Chung-Lin Tang; +Cc: gcc-patches, Catherine Moore
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
^ permalink raw reply [flat|nested] 5+ messages in thread
* Re: [PATCH, OpenACC 2.7] Implement default clause support for data constructs
2023-06-23 10:47 ` Thomas Schwinge
@ 2023-07-14 10:34 ` Chung-Lin Tang
2023-08-15 14:35 ` [v3] OpenACC 2.7: default clause support for data constructs (was: [PATCH, OpenACC 2.7, v2] Implement default clause support for data constructs) Thomas Schwinge
0 siblings, 1 reply; 5+ messages in thread
From: Chung-Lin Tang @ 2023-07-14 10:34 UTC (permalink / raw)
To: Thomas Schwinge, Chung-Lin Tang; +Cc: gcc-patches, Catherine Moore
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
^ permalink raw reply [flat|nested] 5+ messages in thread
* [PATCH, OpenACC 2.7, v2] Implement default clause support for data constructs
@ 2023-08-01 15:35 Chung-Lin Tang
0 siblings, 0 replies; 5+ messages in thread
From: Chung-Lin Tang @ 2023-08-01 15:35 UTC (permalink / raw)
To: gcc-patches, Thomas Schwinge, Catherine Moore
[-- Attachment #1: Type: text/plain, Size: 1883 bytes --]
Hi Thomas,
this is v2 of the patch for implementing the OpenACC 2.7 addition of
default(none|present) support for data constructs.
Instead of propagating an additional 'oacc_default_kind' for OpenACC,
this patch does it in a more complete way: it directly propagates the
gimplify_omp_ctx* pointer of the inner most context where we found
a default-clause. This supports displaying the location/type of OpenACC
construct where the default-clause is in the error messages.
The testcases also have the multiple nested data construct testing added,
where we can now have messages referring precisely to the exact innermost
default clause that was active at that program point.
Note, I got rid of the dummy OMP_CLAUSE_DEFAULT creation in this version,
since it seemed not really needed.
Re-tested on master on powerpc64le-linux/nvptx. Okay to commit?
Thanks,
Chung-Lin
2023-08-01 Chung-Lin Tang <cltang@codesourcery.com>
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_default_clause_ctx
field.
(new_omp_context): Initialize oacc_default_clause_ctx field.
(oacc_region_type_name): New function.
(oacc_default_clause): Lookup current default_kind value from
ctx->oacc_default_clause_ctx, adjust default(none) error and inform
message dumping.
(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-5.c: Adjust testcase.
* gfortran.dg/goacc/default-3.f95: Adjust testcase.
* gfortran.dg/goacc/default-5.f: Adjust testcase.
[-- Attachment #2: oaccdata-default-clause-v2.patch --]
[-- Type: text/plain, Size: 15769 bytes --]
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 24a6eb6e459..974f0132787 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -18196,6 +18196,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 d7ef5b34d42..bc59fbeac20 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -45860,6 +45860,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 2952cd300ac..c37f843ec3b 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..ec0ccc67da8 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;
+ struct gimplify_omp_ctx *oacc_default_clause_ctx;
enum omp_region_type region_type;
enum tree_code code;
bool combined_loop;
@@ -459,6 +460,10 @@ 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_default_clause_ctx = gimplify_omp_ctxp->oacc_default_clause_ctx;
+ else
+ c->oacc_default_clause_ctx = c;
c->defaultmap[GDMK_SCALAR] = GOVD_MAP;
c->defaultmap[GDMK_SCALAR_TARGET] = GOVD_MAP;
c->defaultmap[GDMK_AGGREGATE] = GOVD_MAP;
@@ -7699,6 +7704,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 +7730,6 @@ 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;
bool on_device = false;
bool is_private = false;
bool declared = is_oacc_declared (decl);
@@ -7735,17 +7758,20 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
flags |= GOVD_MAP_TO_ONLY;
}
+ /* Use the enclosing construct with a default clause to set the current
+ default kind. */
+ enum omp_clause_default_kind default_kind
+ = ctx->oacc_default_clause_ctx->default_kind;
+
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 (default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
flags |= GOVD_MAP;
else
flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
@@ -7758,8 +7784,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 +7791,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 (default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
flags |= GOVD_MAP;
else
flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
@@ -7785,16 +7809,20 @@ 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 (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));
+ inform (ctx->oacc_default_clause_ctx->location,
+ "enclosing OpenACC %qs construct",
+ oacc_region_type_name
+ (ctx->oacc_default_clause_ctx->region_type));
+ }
+ else if (default_kind == OMP_CLAUSE_DEFAULT_PRESENT)
; /* Handled above. */
else
- gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED);
+ gcc_checking_assert (default_kind == OMP_CLAUSE_DEFAULT_SHARED);
return flags;
}
@@ -12124,6 +12152,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_DEFAULT:
ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
+ if (flag_openacc)
+ /* For OpenACC, set current context to the 'active' one for
+ default-clause lookup. */
+ ctx->oacc_default_clause_ctx = ctx;
break;
case OMP_CLAUSE_INCLUSIVE:
diff --git a/gcc/testsuite/c-c++-common/goacc/default-3.c b/gcc/testsuite/c-c++-common/goacc/default-3.c
index ac169a903e9..e68f33606fd 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,49 @@ 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) /* { dg-message "enclosing OpenACC .data. construct" } */
+#pragma acc kernels
+ {
+ 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) /* { dg-message "enclosing OpenACC .data. construct" } */
+#pragma acc parallel
+ {
+ 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-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" } */
+ }
+
+#pragma acc data default (none) /* { dg-message "enclosing OpenACC .data. construct" } */
+#pragma acc data
+#pragma acc data
+#pragma acc parallel
+ {
+ 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-message "enclosing OpenACC .data. construct" } */
+#pragma acc data
+#pragma acc parallel
+ {
+ 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-message "enclosing OpenACC .data. construct" } */
+#pragma acc parallel
+ {
+ 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..76526396124 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 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 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..09b8c1bbd87 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/default-3.f95
@@ -15,4 +15,65 @@ 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) ! { dg-message "enclosing OpenACC .data. construct" }
+ !$acc kernels
+ 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-message "enclosing OpenACC .data. construct" }
+ !$acc parallel
+ 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-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
+
+ !$acc data default (none) ! { dg-message "enclosing OpenACC .data. construct" }
+ !$acc data
+ !$acc data
+ !$acc parallel
+ 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-message "enclosing OpenACC .data. construct" }
+ !$acc data
+ !$acc parallel
+ 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-message "enclosing OpenACC .data. construct" }
+ !$acc parallel
+ 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-5.f b/gcc/testsuite/gfortran.dg/goacc/default-5.f
index 9dc83cbe601..9260db07a1a 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 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
^ permalink raw reply [flat|nested] 5+ messages in thread
* [v3] OpenACC 2.7: default clause support for data constructs (was: [PATCH, OpenACC 2.7, v2] Implement default clause support for data constructs)
2023-07-14 10:34 ` Chung-Lin Tang
@ 2023-08-15 14:35 ` Thomas Schwinge
0 siblings, 0 replies; 5+ messages in thread
From: Thomas Schwinge @ 2023-08-15 14:35 UTC (permalink / raw)
To: Chung-Lin Tang, gcc-patches; +Cc: Catherine Moore
[-- Attachment #1: Type: text/plain, Size: 3004 bytes --]
Hi!
On 2023-08-01T23:35:16+0800, Chung-Lin Tang <chunglin.tang@siemens.com> wrote:
> this is v2 of the patch for implementing the OpenACC 2.7 addition of
> default(none|present) support for data constructs.
Thanks!
> Instead of propagating an additional 'oacc_default_kind' for OpenACC,
> this patch does it in a more complete way: it directly propagates the
> gimplify_omp_ctx* pointer of the inner most context where we found
> a default-clause.
Right -- but reviewing this, it came upon me that we don't need any such
new code at all, and instead may in 'gcc/gimplify.cc:oacc_default_clause'
simply look through the 'ctx's to find the 'default' clause information.
This centralizes the logic in the one place where it's relevant.
> This supports displaying the location/type of OpenACC
> construct where the default-clause is in the error messages.
This is preserved...
> The testcases also have the multiple nested data construct testing added,
> where we can now have messages referring precisely to the exact innermost
> default clause that was active at that program point.
..., but we should also still 'inform' about the compute construct, where
the user is expected to add explicit data clauses (if not adding to the
'data' construct where the 'default(none)' clause appears):
> --- a/gcc/gimplify.cc
> +++ b/gcc/gimplify.cc
> @@ -7785,16 +7809,20 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
> - else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_NONE)
> + else if (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));
> + inform (ctx->oacc_default_clause_ctx->location,
> + "enclosing OpenACC %qs construct",
> + oacc_region_type_name
> + (ctx->oacc_default_clause_ctx->region_type));
> + }
That is, we should keep here the original 'inform' for 'ctx->location',
and *add another* 'inform' for 'ctx->oacc_default_clause_ctx->location'.
Otherwise that's confusing to users.
Instead of requiring another iteration through you, I've now implemented
that, and with test cases enhanced some more, pushed to master branch
commit bed993884b149851fe930b43cf11cbcdf05f1578
"OpenACC 2.7: default clause support for data constructs", see attached.
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
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-OpenACC-2.7-default-clause-support-for-data-construc.patch --]
[-- Type: text/x-diff, Size: 24541 bytes --]
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
^ permalink raw reply [flat|nested] 5+ messages in thread
end of thread, other threads:[~2023-08-15 14:37 UTC | newest]
Thread overview: 5+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-06-06 15:11 [PATCH, OpenACC 2.7] Implement default clause support for data constructs Chung-Lin Tang
2023-06-23 10:47 ` Thomas Schwinge
2023-07-14 10:34 ` Chung-Lin Tang
2023-08-15 14:35 ` [v3] OpenACC 2.7: default clause support for data constructs (was: [PATCH, OpenACC 2.7, v2] Implement default clause support for data constructs) Thomas Schwinge
2023-08-01 15:35 [PATCH, OpenACC 2.7, v2] Implement default clause support for data constructs Chung-Lin Tang
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).