aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authortschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>2017-05-19 13:32:48 +0000
committertschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>2017-05-19 13:32:48 +0000
commit6acf639f20cb088fbc81b465c924c66d152af288 (patch)
treea1ddeff9f474c8969d68b1b0b726f4ca763411f3
parentf7c10d53cac5a46d06d23177526aa36c021e135c (diff)
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
-rw-r--r--gcc/ChangeLog9
-rw-r--r--gcc/c/ChangeLog5
-rw-r--r--gcc/c/c-parser.c14
-rw-r--r--gcc/cp/ChangeLog3
-rw-r--r--gcc/cp/parser.c14
-rw-r--r--gcc/fortran/ChangeLog8
-rw-r--r--gcc/fortran/dump-parse-tree.c1
-rw-r--r--gcc/fortran/gfortran.h3
-rw-r--r--gcc/fortran/openmp.c20
-rw-r--r--gcc/fortran/trans-openmp.c3
-rw-r--r--gcc/gimplify.c52
-rw-r--r--gcc/testsuite/ChangeLog8
-rw-r--r--gcc/testsuite/c-c++-common/goacc/default-1.c5
-rw-r--r--gcc/testsuite/c-c++-common/goacc/default-2.c28
-rw-r--r--gcc/testsuite/c-c++-common/goacc/default-4.c21
-rw-r--r--gcc/testsuite/c-c++-common/goacc/default-5.c20
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/default-1.f955
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/default-4.f18
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/default-5.f18
-rw-r--r--gcc/tree-core.h3
-rw-r--r--gcc/tree-pretty-print.c3
-rw-r--r--libgomp/ChangeLog6
-rw-r--r--libgomp/testsuite/libgomp.oacc-c++/template-reduction.C25
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c31
-rw-r--r--libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f9021
-rw-r--r--libgomp/testsuite/libgomp.oacc-fortran/default-1.f9010
-rw-r--r--libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f9044
27 files changed, 346 insertions, 52 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 8f63902d064..c40af474ea3 100644
--- a/gcc/ChangeLog
+++ b/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 a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 79643cc1b3b..a70eaf910b0 100644
--- a/gcc/c/ChangeLog
+++ b/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 a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 96c074961da..2e013161a9a 100644
--- a/gcc/c/c-parser.c
+++ b/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 a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index c89f71981e3..f2dced41ddb 100644
--- a/gcc/cp/ChangeLog
+++ b/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 a/gcc/cp/parser.c b/gcc/cp/parser.c
index 6453397b530..b345110fbc4 100644
--- a/gcc/cp/parser.c
+++ b/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 a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog
index 9a8b3e1c396..4b7f1f4e4bf 100644
--- a/gcc/fortran/ChangeLog
+++ b/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 a/gcc/fortran/dump-parse-tree.c b/gcc/fortran/dump-parse-tree.c
index 87a530458f9..49b23d8697e 100644
--- a/gcc/fortran/dump-parse-tree.c
+++ b/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 a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 293655078a4..26b89bee98e 100644
--- a/gcc/fortran/gfortran.h
+++ b/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 a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 89eecfa2ed1..80146e2f1dc 100644
--- a/gcc/fortran/openmp.c
+++ b/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 a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 662036f514d..1d254c6904d 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/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 a/gcc/gimplify.c b/gcc/gimplify.c
index 0c02ee4371e..810d9f4a3e1 100644
--- a/gcc/gimplify.c
+++ b/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 a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index ca901c29a66..e8e2df5c9a9 100644
--- a/gcc/testsuite/ChangeLog
+++ b/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 a/gcc/testsuite/c-c++-common/goacc/default-1.c b/gcc/testsuite/c-c++-common/goacc/default-1.c
index 4d31dbc981b..23c25ab45a9 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-1.c
+++ b/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 a/gcc/testsuite/c-c++-common/goacc/default-2.c b/gcc/testsuite/c-c++-common/goacc/default-2.c
index d6b6cdce4d7..c0cdd1263bd 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-2.c
+++ b/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 a/gcc/testsuite/c-c++-common/goacc/default-4.c b/gcc/testsuite/c-c++-common/goacc/default-4.c
index 787b35256db..dfa79bbbe6e 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-4.c
+++ b/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 a/gcc/testsuite/c-c++-common/goacc/default-5.c b/gcc/testsuite/c-c++-common/goacc/default-5.c
new file mode 100644
index 00000000000..37e3c3555cd
--- /dev/null
+++ b/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 a/gcc/testsuite/gfortran.dg/goacc/default-1.f95 b/gcc/testsuite/gfortran.dg/goacc/default-1.f95
index a79b444f6f4..41fa9446b91 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-1.f95
+++ b/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 a/gcc/testsuite/gfortran.dg/goacc/default-4.f b/gcc/testsuite/gfortran.dg/goacc/default-4.f
index b2f73c3bcfb..77291f43eff 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-4.f
+++ b/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 a/gcc/testsuite/gfortran.dg/goacc/default-5.f b/gcc/testsuite/gfortran.dg/goacc/default-5.f
new file mode 100644
index 00000000000..9dc83cbe601
--- /dev/null
+++ b/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 a/gcc/tree-core.h b/gcc/tree-core.h
index c76fc7bc8ef..ea73477163d 100644
--- a/gcc/tree-core.h
+++ b/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 a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index ec28b1e3c95..b70e32573ee 100644
--- a/gcc/tree-pretty-print.c
+++ b/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 a/libgomp/ChangeLog b/libgomp/ChangeLog
index 84d1c839d70..460605b1b8a 100644
--- a/libgomp/ChangeLog
+++ b/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 a/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C b/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C
index 6c85fba7a92..ede0aae14c6 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/template-reduction.C
+++ b/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 a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
index c16459826af..51b3b18ad66 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
+++ b/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 a/libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
index 16a85980b4d..df4aca06a70 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
+++ b/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 a/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
index 10590892458..d8ceb60b493 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
+++ b/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 a/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90 b/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
index 94e4228e7a9..53b2cd060f2 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
+++ b/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