This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: OpenACC 2.5 default (present) clause
- From: Thomas Schwinge <thomas at codesourcery dot com>
- To: Jakub Jelinek <jakub at redhat dot com>, <gcc-patches at gcc dot gnu dot org>
- Date: Fri, 19 May 2017 15:46:53 +0200
- Subject: Re: OpenACC 2.5 default (present) clause
- Authentication-results: sourceware.org; auth=none
- References: <87a87s5k6w.fsf@hertz.schwinge.homeip.net> <20170510154836.GA1809@tucnak>
Hi!
On Wed, 10 May 2017 17:48:36 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Fri, Apr 07, 2017 at 05:08:55PM +0200, Thomas Schwinge wrote:
> > OpenACC 2.5 added a default (present) clause, which "causes all arrays or
> > variables of aggregate data type used in the compute construct that have
> > implicitly determined data attributes to be treated as if they appeared
> > in a present clause". [...] OK for trunk in next stage 1?
> > @@ -6897,6 +6900,7 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl,
> > found_outer:
> > break;
> >
> > + case OMP_CLAUSE_DEFAULT_PRESENT:
> > default:
> > gcc_unreachable ();
> > }
>
> What is the point of this hunk? Document that present clause is not in
> OpenMP?
Right.
> I don't think that is needed.
As you please.
> > + case 0 | GOVD_MAP_FORCE:
> > + kind = GOMP_MAP_TOFROM | GOMP_MAP_FLAG_FORCE;
> > + break;
>
> Please drop the "0 | ".
Removed; I thought that was a good way to relate this "case" to the
"case 0" just above.
> > @@ -357,7 +357,9 @@ enum omp_clause_code {
> > /* OpenMP clause: ordered [(constant-integer-expression)]. */
> > OMP_CLAUSE_ORDERED,
> >
> > - /* OpenMP clause: default. */
> > + /* OpenACC clause: default ( none | present ).
> > +
> > + OpenMP clause: default ( firstprivate | none | private | shared ). */
> > OMP_CLAUSE_DEFAULT,
> >
> > /* OpenACC/OpenMP clause: collapse (constant-integer-expression). */
>
> I think this hunk isn't needed (plus it is not accurate anyway).
Now you got me curious: why isn't it accurate? Anyway, I changed that to
just say "OpenACC/OpenMP clause: default".
> Otherwise LGTM.
Thanks. Committed to trunk in r248280:
commit 6acf639f20cb088fbc81b465c924c66d152af288
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri May 19 13:32:48 2017 +0000
OpenACC 2.5 default (present) clause
gcc/c/
* c-parser.c (c_parser_omp_clause_default): Handle
"OMP_CLAUSE_DEFAULT_PRESENT".
gcc/cp/
* parser.c (cp_parser_omp_clause_default): Handle
"OMP_CLAUSE_DEFAULT_PRESENT".
gcc/fortran/
* gfortran.h (enum gfc_omp_default_sharing): Add
"OMP_DEFAULT_PRESENT".
* dump-parse-tree.c (show_omp_clauses): Handle it.
* openmp.c (gfc_match_omp_clauses): Likewise.
* trans-openmp.c (gfc_trans_omp_clauses): Likewise.
gcc/
* tree-core.h (enum omp_clause_default_kind): Add
"OMP_CLAUSE_DEFAULT_PRESENT".
* tree-pretty-print.c (dump_omp_clause): Handle it.
* gimplify.c (enum gimplify_omp_var_data): Add
"GOVD_MAP_FORCE_PRESENT".
(gimplify_adjust_omp_clauses_1): Map it to
"GOMP_MAP_FORCE_PRESENT".
(oacc_default_clause): Handle "OMP_CLAUSE_DEFAULT_PRESENT".
gcc/testsuite/
* c-c++-common/goacc/default-1.c: Update.
* c-c++-common/goacc/default-2.c: Likewise.
* c-c++-common/goacc/default-4.c: Likewise.
* gfortran.dg/goacc/default-1.f95: Likewise.
* gfortran.dg/goacc/default-4.f: Likewise.
* c-c++-common/goacc/default-5.c: New file.
* gfortran.dg/goacc/default-5.f: Likewise.
libgomp/
* testsuite/libgomp.oacc-c++/template-reduction.C: Update.
* testsuite/libgomp.oacc-c-c++-common/nested-2.c: Update.
* testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise.
* testsuite/libgomp.oacc-fortran/default-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/non-scalar-data.f90: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@248280 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 9 ++++
gcc/c/ChangeLog | 5 +++
gcc/c/c-parser.c | 14 ++++--
gcc/cp/ChangeLog | 3 ++
gcc/cp/parser.c | 14 ++++--
gcc/fortran/ChangeLog | 8 ++++
gcc/fortran/dump-parse-tree.c | 1 +
gcc/fortran/gfortran.h | 3 +-
gcc/fortran/openmp.c | 20 ++++++---
gcc/fortran/trans-openmp.c | 3 ++
gcc/gimplify.c | 52 ++++++++++++++++++----
gcc/testsuite/ChangeLog | 8 ++++
gcc/testsuite/c-c++-common/goacc/default-1.c | 5 +++
gcc/testsuite/c-c++-common/goacc/default-2.c | 28 ++++++------
gcc/testsuite/c-c++-common/goacc/default-4.c | 21 +++++++++
gcc/testsuite/c-c++-common/goacc/default-5.c | 20 +++++++++
gcc/testsuite/gfortran.dg/goacc/default-1.f95 | 5 +++
gcc/testsuite/gfortran.dg/goacc/default-4.f | 18 ++++++++
gcc/testsuite/gfortran.dg/goacc/default-5.f | 18 ++++++++
gcc/tree-core.h | 3 +-
gcc/tree-pretty-print.c | 3 ++
libgomp/ChangeLog | 6 +++
.../libgomp.oacc-c++/template-reduction.C | 25 +++++++++++
.../testsuite/libgomp.oacc-c-c++-common/nested-2.c | 31 +++++++++++++
.../testsuite/libgomp.oacc-fortran/data-4-2.f90 | 21 ++++-----
.../testsuite/libgomp.oacc-fortran/default-1.f90 | 10 +++++
.../libgomp.oacc-fortran/non-scalar-data.f90 | 44 +++++++++++++++++-
27 files changed, 346 insertions(+), 52 deletions(-)
diff --git gcc/ChangeLog gcc/ChangeLog
index 8f63902..c40af47 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,5 +1,14 @@
2017-05-19 Thomas Schwinge <thomas@codesourcery.com>
+ * tree-core.h (enum omp_clause_default_kind): Add
+ "OMP_CLAUSE_DEFAULT_PRESENT".
+ * tree-pretty-print.c (dump_omp_clause): Handle it.
+ * gimplify.c (enum gimplify_omp_var_data): Add
+ "GOVD_MAP_FORCE_PRESENT".
+ (gimplify_adjust_omp_clauses_1): Map it to
+ "GOMP_MAP_FORCE_PRESENT".
+ (oacc_default_clause): Handle "OMP_CLAUSE_DEFAULT_PRESENT".
+
* gimplify.c (oacc_default_clause): Clarify.
2017-05-19 Nathan Sidwell <nathan@acm.org>
diff --git gcc/c/ChangeLog gcc/c/ChangeLog
index 79643cc..a70eaf9 100644
--- gcc/c/ChangeLog
+++ gcc/c/ChangeLog
@@ -1,3 +1,8 @@
+2017-05-19 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-parser.c (c_parser_omp_clause_default): Handle
+ "OMP_CLAUSE_DEFAULT_PRESENT".
+
2017-05-18 Bernd Edlinger <bernd.edlinger@hotmail.de>
* config-lang.in (gtfiles): Add c-family/c-format.c.
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 96c0749..2e01316 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -11057,10 +11057,10 @@ c_parser_omp_clause_copyprivate (c_parser *parser, tree list)
}
/* OpenMP 2.5:
- default ( shared | none )
+ default ( none | shared )
- OpenACC 2.0:
- default (none) */
+ OpenACC:
+ default ( none | present ) */
static tree
c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc)
@@ -11083,6 +11083,12 @@ c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc)
kind = OMP_CLAUSE_DEFAULT_NONE;
break;
+ case 'p':
+ if (strcmp ("present", p) != 0 || !is_oacc)
+ goto invalid_kind;
+ kind = OMP_CLAUSE_DEFAULT_PRESENT;
+ break;
+
case 's':
if (strcmp ("shared", p) != 0 || is_oacc)
goto invalid_kind;
@@ -11099,7 +11105,7 @@ c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc)
{
invalid_kind:
if (is_oacc)
- c_parser_error (parser, "expected %<none%>");
+ c_parser_error (parser, "expected %<none%> or %<present%>");
else
c_parser_error (parser, "expected %<none%> or %<shared%>");
}
diff --git gcc/cp/ChangeLog gcc/cp/ChangeLog
index c89f719..f2dced4 100644
--- gcc/cp/ChangeLog
+++ gcc/cp/ChangeLog
@@ -1,5 +1,8 @@
2017-05-19 Thomas Schwinge <thomas@codesourcery.com>
+ * parser.c (cp_parser_omp_clause_default): Handle
+ "OMP_CLAUSE_DEFAULT_PRESENT".
+
* parser.c (cp_parser_omp_clause_default): Avoid printing more
than one syntax error message.
diff --git gcc/cp/parser.c gcc/cp/parser.c
index 6453397..b345110 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -31456,10 +31456,10 @@ cp_parser_omp_clause_collapse (cp_parser *parser, tree list, location_t location
}
/* OpenMP 2.5:
- default ( shared | none )
+ default ( none | shared )
- OpenACC 2.0
- default (none) */
+ OpenACC:
+ default ( none | present ) */
static tree
cp_parser_omp_clause_default (cp_parser *parser, tree list,
@@ -31483,6 +31483,12 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list,
kind = OMP_CLAUSE_DEFAULT_NONE;
break;
+ case 'p':
+ if (strcmp ("present", p) != 0 || !is_oacc)
+ goto invalid_kind;
+ kind = OMP_CLAUSE_DEFAULT_PRESENT;
+ break;
+
case 's':
if (strcmp ("shared", p) != 0 || is_oacc)
goto invalid_kind;
@@ -31499,7 +31505,7 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list,
{
invalid_kind:
if (is_oacc)
- cp_parser_error (parser, "expected %<none%>");
+ cp_parser_error (parser, "expected %<none%> or %<present%>");
else
cp_parser_error (parser, "expected %<none%> or %<shared%>");
}
diff --git gcc/fortran/ChangeLog gcc/fortran/ChangeLog
index 9a8b3e1..4b7f1f4 100644
--- gcc/fortran/ChangeLog
+++ gcc/fortran/ChangeLog
@@ -1,3 +1,11 @@
+2017-05-19 Thomas Schwinge <thomas@codesourcery.com>
+
+ * gfortran.h (enum gfc_omp_default_sharing): Add
+ "OMP_DEFAULT_PRESENT".
+ * dump-parse-tree.c (show_omp_clauses): Handle it.
+ * openmp.c (gfc_match_omp_clauses): Likewise.
+ * trans-openmp.c (gfc_trans_omp_clauses): Likewise.
+
2017-05-18 Fritz Reese <fritzoreese@gmail.com>
PR fortran/79968
diff --git gcc/fortran/dump-parse-tree.c gcc/fortran/dump-parse-tree.c
index 87a5304..49b23d8 100644
--- gcc/fortran/dump-parse-tree.c
+++ gcc/fortran/dump-parse-tree.c
@@ -1283,6 +1283,7 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses)
case OMP_DEFAULT_PRIVATE: type = "PRIVATE"; break;
case OMP_DEFAULT_SHARED: type = "SHARED"; break;
case OMP_DEFAULT_FIRSTPRIVATE: type = "FIRSTPRIVATE"; break;
+ case OMP_DEFAULT_PRESENT: type = "PRESENT"; break;
default:
gcc_unreachable ();
}
diff --git gcc/fortran/gfortran.h gcc/fortran/gfortran.h
index 2936550..26b89be 100644
--- gcc/fortran/gfortran.h
+++ gcc/fortran/gfortran.h
@@ -1241,7 +1241,8 @@ enum gfc_omp_default_sharing
OMP_DEFAULT_NONE,
OMP_DEFAULT_PRIVATE,
OMP_DEFAULT_SHARED,
- OMP_DEFAULT_FIRSTPRIVATE
+ OMP_DEFAULT_FIRSTPRIVATE,
+ OMP_DEFAULT_PRESENT
};
enum gfc_omp_proc_bind_kind
diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
index 89eecfa..80146e2 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -1080,13 +1080,19 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if (gfc_match ("default ( none )") == MATCH_YES)
c->default_sharing = OMP_DEFAULT_NONE;
else if (openacc)
- /* c->default_sharing = OMP_DEFAULT_UNKNOWN */;
- else if (gfc_match ("default ( shared )") == MATCH_YES)
- c->default_sharing = OMP_DEFAULT_SHARED;
- else if (gfc_match ("default ( private )") == MATCH_YES)
- c->default_sharing = OMP_DEFAULT_PRIVATE;
- else if (gfc_match ("default ( firstprivate )") == MATCH_YES)
- c->default_sharing = OMP_DEFAULT_FIRSTPRIVATE;
+ {
+ if (gfc_match ("default ( present )") == MATCH_YES)
+ c->default_sharing = OMP_DEFAULT_PRESENT;
+ }
+ else
+ {
+ if (gfc_match ("default ( firstprivate )") == MATCH_YES)
+ c->default_sharing = OMP_DEFAULT_FIRSTPRIVATE;
+ else if (gfc_match ("default ( private )") == MATCH_YES)
+ c->default_sharing = OMP_DEFAULT_PRIVATE;
+ else if (gfc_match ("default ( shared )") == MATCH_YES)
+ c->default_sharing = OMP_DEFAULT_SHARED;
+ }
if (c->default_sharing != OMP_DEFAULT_UNKNOWN)
continue;
}
diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c
index 662036f..1d254c6 100644
--- gcc/fortran/trans-openmp.c
+++ gcc/fortran/trans-openmp.c
@@ -2564,6 +2564,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_DEFAULT_FIRSTPRIVATE:
OMP_CLAUSE_DEFAULT_KIND (c) = OMP_CLAUSE_DEFAULT_FIRSTPRIVATE;
break;
+ case OMP_DEFAULT_PRESENT:
+ OMP_CLAUSE_DEFAULT_KIND (c) = OMP_CLAUSE_DEFAULT_PRESENT;
+ break;
default:
gcc_unreachable ();
}
diff --git gcc/gimplify.c gcc/gimplify.c
index 0c02ee4..810d9f4 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -99,6 +99,9 @@ enum gimplify_omp_var_data
/* Flag for GOVD_MAP, if it is a forced mapping. */
GOVD_MAP_FORCE = 262144,
+ /* Flag for GOVD_MAP: must be present already. */
+ GOVD_MAP_FORCE_PRESENT = 524288,
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -6956,8 +6959,13 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
rkind = "kernels";
if (AGGREGATE_TYPE_P (type))
- /* Aggregates default to 'present_or_copy'. */
- flags |= GOVD_MAP;
+ {
+ /* Aggregates default to 'present_or_copy', or 'present'. */
+ if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
+ flags |= GOVD_MAP;
+ else
+ flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
+ }
else
/* Scalars default to 'copy'. */
flags |= GOVD_MAP | GOVD_MAP_FORCE;
@@ -6970,8 +6978,13 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
if (on_device || declared)
flags |= GOVD_MAP;
else if (AGGREGATE_TYPE_P (type))
- /* Aggregates default to 'present_or_copy'. */
- flags |= GOVD_MAP;
+ {
+ /* Aggregates default to 'present_or_copy', or 'present'. */
+ if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
+ flags |= GOVD_MAP;
+ else
+ flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
+ }
else
/* Scalars default to 'firstprivate'. */
flags |= GOVD_FIRSTPRIVATE;
@@ -6991,6 +7004,8 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
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)
+ ; /* Handled above. */
else
gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED);
@@ -8708,11 +8723,30 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
}
else if (code == OMP_CLAUSE_MAP)
{
- int kind = (flags & GOVD_MAP_TO_ONLY
- ? GOMP_MAP_TO
- : GOMP_MAP_TOFROM);
- if (flags & GOVD_MAP_FORCE)
- kind |= GOMP_MAP_FLAG_FORCE;
+ int kind;
+ /* Not all combinations of these GOVD_MAP flags are actually valid. */
+ switch (flags & (GOVD_MAP_TO_ONLY
+ | GOVD_MAP_FORCE
+ | GOVD_MAP_FORCE_PRESENT))
+ {
+ case 0:
+ kind = GOMP_MAP_TOFROM;
+ break;
+ case GOVD_MAP_FORCE:
+ kind = GOMP_MAP_TOFROM | GOMP_MAP_FLAG_FORCE;
+ break;
+ case GOVD_MAP_TO_ONLY:
+ kind = GOMP_MAP_TO;
+ break;
+ case GOVD_MAP_TO_ONLY | GOVD_MAP_FORCE:
+ kind = GOMP_MAP_TO | GOMP_MAP_FLAG_FORCE;
+ break;
+ case GOVD_MAP_FORCE_PRESENT:
+ kind = GOMP_MAP_FORCE_PRESENT;
+ break;
+ default:
+ gcc_unreachable ();
+ }
OMP_CLAUSE_SET_MAP_KIND (clause, kind);
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog
index ca901c2..e8e2df5 100644
--- gcc/testsuite/ChangeLog
+++ gcc/testsuite/ChangeLog
@@ -1,5 +1,13 @@
2017-05-19 Thomas Schwinge <thomas@codesourcery.com>
+ * c-c++-common/goacc/default-1.c: Update.
+ * c-c++-common/goacc/default-2.c: Likewise.
+ * c-c++-common/goacc/default-4.c: Likewise.
+ * gfortran.dg/goacc/default-1.f95: Likewise.
+ * gfortran.dg/goacc/default-4.f: Likewise.
+ * c-c++-common/goacc/default-5.c: New file.
+ * gfortran.dg/goacc/default-5.f: Likewise.
+
* c-c++-common/goacc/default-1.c: New file.
* c-c++-common/goacc/default-2.c: Likewise.
* c-c++-common/goacc/data-default-1.c: Remove file, including its
diff --git gcc/testsuite/c-c++-common/goacc/default-1.c gcc/testsuite/c-c++-common/goacc/default-1.c
index 4d31dbc..23c25ab 100644
--- gcc/testsuite/c-c++-common/goacc/default-1.c
+++ gcc/testsuite/c-c++-common/goacc/default-1.c
@@ -6,4 +6,9 @@ void f1 ()
;
#pragma acc parallel default (none)
;
+
+#pragma acc kernels default (present)
+ ;
+#pragma acc parallel default (present)
+ ;
}
diff --git gcc/testsuite/c-c++-common/goacc/default-2.c gcc/testsuite/c-c++-common/goacc/default-2.c
index d6b6cdc..c0cdd12 100644
--- gcc/testsuite/c-c++-common/goacc/default-2.c
+++ gcc/testsuite/c-c++-common/goacc/default-2.c
@@ -7,39 +7,39 @@ void f1 ()
#pragma acc parallel default /* { dg-error "expected .\\(. before end of line" } */
;
-#pragma acc kernels default ( /* { dg-error "expected .none. before end of line" } */
+#pragma acc kernels default ( /* { dg-error "expected .none. or .present. before end of line" } */
;
-#pragma acc parallel default ( /* { dg-error "expected .none. before end of line" } */
+#pragma acc parallel default ( /* { dg-error "expected .none. or .present. before end of line" } */
;
-#pragma acc kernels default (, /* { dg-error "expected .none. before .,. token" } */
+#pragma acc kernels default (, /* { dg-error "expected .none. or .present. before .,. token" } */
;
-#pragma acc parallel default (, /* { dg-error "expected .none. before .,. token" } */
+#pragma acc parallel default (, /* { dg-error "expected .none. or .present. before .,. token" } */
;
-#pragma acc kernels default () /* { dg-error "expected .none. before .\\). token" } */
+#pragma acc kernels default () /* { dg-error "expected .none. or .present. before .\\). token" } */
;
-#pragma acc parallel default () /* { dg-error "expected .none. before .\\). token" } */
+#pragma acc parallel default () /* { dg-error "expected .none. or .present. before .\\). token" } */
;
-#pragma acc kernels default (,) /* { dg-error "expected .none. before .,. token" } */
+#pragma acc kernels default (,) /* { dg-error "expected .none. or .present. before .,. token" } */
;
-#pragma acc parallel default (,) /* { dg-error "expected .none. before .,. token" } */
+#pragma acc parallel default (,) /* { dg-error "expected .none. or .present. before .,. token" } */
;
-#pragma acc kernels default (firstprivate) /* { dg-error "expected .none. before .firstprivate." } */
+#pragma acc kernels default (firstprivate) /* { dg-error "expected .none. or .present. before .firstprivate." } */
;
-#pragma acc parallel default (firstprivate) /* { dg-error "expected .none. before .firstprivate." } */
+#pragma acc parallel default (firstprivate) /* { dg-error "expected .none. or .present. before .firstprivate." } */
;
-#pragma acc kernels default (private) /* { dg-error "expected .none. before .private." } */
+#pragma acc kernels default (private) /* { dg-error "expected .none. or .present. before .private." } */
;
-#pragma acc parallel default (private) /* { dg-error "expected .none. before .private." } */
+#pragma acc parallel default (private) /* { dg-error "expected .none. or .present. before .private." } */
;
-#pragma acc kernels default (shared) /* { dg-error "expected .none. before .shared." } */
+#pragma acc kernels default (shared) /* { dg-error "expected .none. or .present. before .shared." } */
;
-#pragma acc parallel default (shared) /* { dg-error "expected .none. before .shared." } */
+#pragma acc parallel default (shared) /* { dg-error "expected .none. or .present. before .shared." } */
;
#pragma acc kernels default (none /* { dg-error "expected .\\). before end of line" } */
diff --git gcc/testsuite/c-c++-common/goacc/default-4.c gcc/testsuite/c-c++-common/goacc/default-4.c
index 787b352..dfa79bb 100644
--- gcc/testsuite/c-c++-common/goacc/default-4.c
+++ gcc/testsuite/c-c++-common/goacc/default-4.c
@@ -43,3 +43,24 @@ void f2 ()
}
}
}
+
+void f3 ()
+{
+ int f3_a = 2;
+ float f3_b[2];
+
+#pragma acc data copyin (f3_a) copyout (f3_b)
+ /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f3_b \[^\\)\]+\\) map\\(force_to:f3_a" 1 "gimple" } } */
+ {
+#pragma acc kernels default (present)
+ /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */
+ {
+ f3_b[0] = f3_a;
+ }
+#pragma acc parallel default (present)
+ /* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */
+ {
+ f3_b[0] = f3_a;
+ }
+ }
+}
diff --git gcc/testsuite/c-c++-common/goacc/default-5.c gcc/testsuite/c-c++-common/goacc/default-5.c
new file mode 100644
index 0000000..37e3c355
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/default-5.c
@@ -0,0 +1,20 @@
+/* OpenACC default (present) clause. */
+
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void f1 ()
+{
+ int f1_a = 2;
+ float f1_b[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" } } */
+ {
+ f1_b[0] = f1_a;
+ }
+#pragma acc parallel default (present)
+ /* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) firstprivate\\(f1_a\\)" 1 "gimple" } } */
+ {
+ f1_b[0] = f1_a;
+ }
+}
diff --git gcc/testsuite/gfortran.dg/goacc/default-1.f95 gcc/testsuite/gfortran.dg/goacc/default-1.f95
index a79b444..41fa944 100644
--- gcc/testsuite/gfortran.dg/goacc/default-1.f95
+++ gcc/testsuite/gfortran.dg/goacc/default-1.f95
@@ -7,4 +7,9 @@ subroutine f1
!$acc end kernels
!$acc parallel default (none)
!$acc end parallel
+
+ !$acc kernels default (present)
+ !$acc end kernels
+ !$acc parallel default (present)
+ !$acc end parallel
end subroutine f1
diff --git gcc/testsuite/gfortran.dg/goacc/default-4.f gcc/testsuite/gfortran.dg/goacc/default-4.f
index b2f73c3..77291f4 100644
--- gcc/testsuite/gfortran.dg/goacc/default-4.f
+++ gcc/testsuite/gfortran.dg/goacc/default-4.f
@@ -37,3 +37,21 @@
!$ACC END PARALLEL
!$ACC END DATA
END SUBROUTINE F2
+
+ SUBROUTINE F3
+ IMPLICIT NONE
+ INTEGER :: F3_A = 2
+ REAL, DIMENSION (2) :: F3_B
+
+!$ACC DATA COPYIN (F3_A) COPYOUT (F3_B)
+! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f3_a \[^\\)\]+\\) map\\(force_from:f3_b" 1 "gimple" } }
+!$ACC KERNELS DEFAULT (PRESENT)
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } }
+ F3_B(1) = F3_A;
+!$ACC END KERNELS
+!$ACC PARALLEL DEFAULT (PRESENT)
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) 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 gcc/testsuite/gfortran.dg/goacc/default-5.f gcc/testsuite/gfortran.dg/goacc/default-5.f
new file mode 100644
index 0000000..9dc83cb
--- /dev/null
+++ gcc/testsuite/gfortran.dg/goacc/default-5.f
@@ -0,0 +1,18 @@
+! OpenACC default (present) clause.
+
+! { dg-additional-options "-fdump-tree-gimple" }
+
+ SUBROUTINE F1
+ IMPLICIT NONE
+ INTEGER :: F1_A = 2
+ REAL, DIMENSION (2) :: F1_B
+
+!$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" } }
+ F1_B(1) = F1_A;
+!$ACC END KERNELS
+!$ACC PARALLEL DEFAULT (PRESENT)
+! { 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
+ END SUBROUTINE F1
diff --git gcc/tree-core.h gcc/tree-core.h
index c76fc7b..ea73477 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -357,7 +357,7 @@ enum omp_clause_code {
/* OpenMP clause: ordered [(constant-integer-expression)]. */
OMP_CLAUSE_ORDERED,
- /* OpenMP clause: default. */
+ /* OpenACC/OpenMP clause: default. */
OMP_CLAUSE_DEFAULT,
/* OpenACC/OpenMP clause: collapse (constant-integer-expression). */
@@ -499,6 +499,7 @@ enum omp_clause_default_kind {
OMP_CLAUSE_DEFAULT_NONE,
OMP_CLAUSE_DEFAULT_PRIVATE,
OMP_CLAUSE_DEFAULT_FIRSTPRIVATE,
+ OMP_CLAUSE_DEFAULT_PRESENT,
OMP_CLAUSE_DEFAULT_LAST
};
diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c
index ec28b1e..b70e325 100644
--- gcc/tree-pretty-print.c
+++ gcc/tree-pretty-print.c
@@ -502,6 +502,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case OMP_CLAUSE_DEFAULT_FIRSTPRIVATE:
pp_string (pp, "firstprivate");
break;
+ case OMP_CLAUSE_DEFAULT_PRESENT:
+ pp_string (pp, "present");
+ break;
default:
gcc_unreachable ();
}
diff --git libgomp/ChangeLog libgomp/ChangeLog
index 84d1c839..460605b 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,5 +1,11 @@
2017-05-19 Thomas Schwinge <thomas@codesourcery.com>
+ * testsuite/libgomp.oacc-c++/template-reduction.C: Update.
+ * testsuite/libgomp.oacc-c-c++-common/nested-2.c: Update.
+ * testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/default-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/non-scalar-data.f90: Likewise.
+
* plugin/plugin-hsa.c (DLSYM_FN, init_hsa_runtime_functions):
Debug output for failure.
diff --git libgomp/testsuite/libgomp.oacc-c++/template-reduction.C libgomp/testsuite/libgomp.oacc-c++/template-reduction.C
index 6c85fba..ede0aae 100644
--- libgomp/testsuite/libgomp.oacc-c++/template-reduction.C
+++ libgomp/testsuite/libgomp.oacc-c++/template-reduction.C
@@ -32,6 +32,28 @@ sum ()
return s;
}
+// Check template with default (present)
+
+template<typename T> T
+sum_default_present ()
+{
+ T s = 0;
+ T array[n];
+
+ for (int i = 0; i < n; i++)
+ array[i] = i+1;
+
+#pragma acc enter data copyin (array)
+
+#pragma acc parallel loop num_gangs (10) gang reduction (+:s) default (present)
+ for (int i = 0; i < n; i++)
+ s += array[i];
+
+#pragma acc exit data delete (array)
+
+ return s;
+}
+
// Check present and async
template<typename T> T
@@ -86,6 +108,9 @@ main()
if (sum<int> () != result)
__builtin_abort ();
+ if (sum_default_present<int> () != result)
+ __builtin_abort ();
+
#pragma acc enter data copyin (a)
if (async_sum (a) != result)
__builtin_abort ();
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
index c164598..51b3b18 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
@@ -137,5 +137,36 @@ main (int argc, char *argv[])
abort ();
}
+
+#pragma acc enter data create (a)
+
+#pragma acc parallel default (present)
+ {
+ for (int j = 0; j < N; ++j)
+ a[j] = j + 1;
+ }
+
+#pragma acc update host (a)
+
+ for (i = 0; i < N; ++i)
+ {
+ if (a[i] != i + 1)
+ abort ();
+ }
+
+#pragma acc kernels default (present)
+ {
+ for (int j = 0; j < N; ++j)
+ a[j] = j + 2;
+ }
+
+#pragma acc exit data copyout (a)
+
+ for (i = 0; i < N; ++i)
+ {
+ if (a[i] != i + 2)
+ abort ();
+ }
+
return 0;
}
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90 libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
index 16a8598..df4aca0 100644
--- libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
@@ -1,4 +1,5 @@
-! Copy of data-4.f90 with self exchanged with host for !acc update.
+! Copy of data-4.f90 with self exchanged with host for !acc update, and with
+! default (present) clauses added.
! { dg-do run }
@@ -19,7 +20,7 @@ program asyncwait
!$acc enter data copyin (a(1:N)) copyin (b(1:N)) copyin (N) async
- !$acc parallel async wait
+ !$acc parallel default (present) async wait
!$acc loop
do i = 1, N
b(i) = a(i)
@@ -39,7 +40,7 @@ program asyncwait
!$acc update device (a(1:N), b(1:N)) async (1)
- !$acc parallel async (1) wait (1)
+ !$acc parallel default (present) async (1) wait (1)
!$acc loop
do i = 1, N
b(i) = a(i)
@@ -62,19 +63,19 @@ program asyncwait
!$acc enter data copyin (c(1:N), d(1:N)) async (1)
!$acc update device (a(1:N), b(1:N)) async (1)
- !$acc parallel async (1)
+ !$acc parallel default (present) async (1)
do i = 1, N
b(i) = (a(i) * a(i) * a(i)) / a(i)
end do
!$acc end parallel
- !$acc parallel async (1)
+ !$acc parallel default (present) async (1)
do i = 1, N
c(i) = (a(i) * 4) / a(i)
end do
!$acc end parallel
- !$acc parallel async (1)
+ !$acc parallel default (present) async (1)
do i = 1, N
d(i) = ((a(i) * a(i) + a(i)) / a(i)) - a(i)
end do
@@ -100,25 +101,25 @@ program asyncwait
!$acc enter data copyin (e(1:N)) async (1)
!$acc update device (a(1:N), b(1:N), c(1:N), d(1:N)) async (1)
- !$acc parallel async (1)
+ !$acc parallel default (present) async (1)
do i = 1, N
b(i) = (a(i) * a(i) * a(i)) / a(i)
end do
!$acc end parallel
- !$acc parallel async (1)
+ !$acc parallel default (present) async (1)
do i = 1, N
c(i) = (a(i) * 4) / a(i)
end do
!$acc end parallel
- !$acc parallel async (1)
+ !$acc parallel default (present) async (1)
do i = 1, N
d(i) = ((a(i) * a(i) + a(i)) / a(i)) - a(i)
end do
!$acc end parallel
- !$acc parallel wait (1) async (1)
+ !$acc parallel default (present) wait (1) async (1)
do i = 1, N
e(i) = a(i) + b(i) + c(i) + d(i)
end do
diff --git libgomp/testsuite/libgomp.oacc-fortran/default-1.f90 libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
index 1059089..d8ceb60 100644
--- libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
@@ -51,4 +51,14 @@ program main
if (a .ne. 7.0) call abort
+ ! The default (present) clause doesn't affect scalar variables; these will
+ ! still get an implicit copy clause added.
+ !$acc kernels default (present)
+ c = a
+ a = 1.0
+ a = a + c
+ !$acc end kernels
+
+ if (a .ne. 8.0) call abort
+
end program main
diff --git libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90 libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
index 94e4228..53b2cd0 100644
--- libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
@@ -1,5 +1,6 @@
! Ensure that a non-scalar dummy arguments which are implicitly used inside
-! offloaded regions are properly mapped using present_or_copy.
+! offloaded regions are properly mapped using present_or_copy, or (default)
+! present.
! { dg-do run }
@@ -12,6 +13,7 @@ program main
n = size
!$acc data copy(array)
+
call kernels(array, n)
!$acc update host(array)
@@ -20,12 +22,29 @@ program main
if (array(i) .ne. i) call abort
end do
+ call kernels_default_present(array, n)
+
+ !$acc update host(array)
+
+ do i = 1, n
+ if (array(i) .ne. i+1) call abort
+ end do
+
call parallel(array, n)
- !$acc end data
+
+ !$acc update host(array)
do i = 1, n
if (array(i) .ne. i+i) call abort
end do
+
+ call parallel_default_present(array, n)
+
+ !$acc end data
+
+ do i = 1, n
+ if (array(i) .ne. i+i+1) call abort
+ end do
end program main
subroutine kernels (array, n)
@@ -39,6 +58,16 @@ subroutine kernels (array, n)
!$acc end kernels
end subroutine kernels
+subroutine kernels_default_present (array, n)
+ integer, dimension (n) :: array
+ integer :: n, i
+
+ !$acc kernels default(present)
+ do i = 1, n
+ array(i) = i+1
+ end do
+ !$acc end kernels
+end subroutine kernels_default_present
subroutine parallel (array, n)
integer, dimension (n) :: array
@@ -50,3 +79,14 @@ subroutine parallel (array, n)
end do
!$acc end parallel
end subroutine parallel
+
+subroutine parallel_default_present (array, n)
+ integer, dimension (n) :: array
+ integer :: n, i
+
+ !$acc parallel default(present)
+ do i = 1, n
+ array(i) = i+i+1
+ end do
+ !$acc end parallel
+end subroutine parallel_default_present
Grüße
Thomas