This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
OpenACC 2.5 default (present) clause
- From: Thomas Schwinge <thomas at codesourcery dot com>
- To: <gcc-patches at gcc dot gnu dot org>, Jakub Jelinek <jakub at redhat dot com>
- Date: Fri, 7 Apr 2017 17:08:55 +0200
- Subject: OpenACC 2.5 default (present) clause
- Authentication-results: sourceware.org; auth=none
Hi!
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". Preceded by the following cleanup patch (see
<https://gcc.gnu.org/ml/gcc-patches/2017-04/msg00377.html> for its
origin), OK for trunk in next stage 1?
commit 787fea9e71f693c1b629a699f8476f392c4bc55d
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Thu Apr 6 07:58:37 2017 +0200
Clarify gcc/gimplify.c:oacc_default_clause
gcc/
* gimplify.c (oacc_default_clause): Clarify.
---
gcc/gimplify.c | 40 ++++++++++++++++++++++------------------
1 file changed, 22 insertions(+), 18 deletions(-)
diff --git gcc/gimplify.c gcc/gimplify.c
index ff8d56b..3fcf9ab 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -6929,30 +6929,34 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
switch (ctx->region_type)
{
- default:
- gcc_unreachable ();
-
case ORT_ACC_KERNELS:
- /* Scalars are default 'copy' under kernels, non-scalars are default
- 'present_or_copy'. */
- flags |= GOVD_MAP;
- if (!AGGREGATE_TYPE_P (type))
- flags |= GOVD_MAP_FORCE;
-
rkind = "kernels";
+
+ if (AGGREGATE_TYPE_P (type))
+ /* Aggregates default to 'present_or_copy'. */
+ flags |= GOVD_MAP;
+ else
+ /* Scalars default to 'copy'. */
+ flags |= GOVD_MAP | GOVD_MAP_FORCE;
+
break;
case ORT_ACC_PARALLEL:
- {
- if (on_device || AGGREGATE_TYPE_P (type) || declared)
- /* Aggregates default to 'present_or_copy'. */
- flags |= GOVD_MAP;
- else
- /* Scalars default to 'firstprivate'. */
- flags |= GOVD_FIRSTPRIVATE;
- rkind = "parallel";
- }
+ rkind = "parallel";
+
+ if (on_device || declared)
+ flags |= GOVD_MAP;
+ else if (AGGREGATE_TYPE_P (type))
+ /* Aggregates default to 'present_or_copy'. */
+ flags |= GOVD_MAP;
+ else
+ /* Scalars default to 'firstprivate'. */
+ flags |= GOVD_FIRSTPRIVATE;
+
break;
+
+ default:
+ gcc_unreachable ();
}
if (DECL_ARTIFICIAL (decl))
commit 4f7643dd63177ba9fdb99063e34f8e287c2e30aa
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed Apr 5 15:33:17 2017 +0200
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".
(omp_default_clause, 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.
---
gcc/c/c-parser.c | 14 ++++--
gcc/cp/parser.c | 14 ++++--
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 | 53 ++++++++++++++++++----
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 | 5 +-
gcc/tree-pretty-print.c | 3 ++
.../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 +++++++++++++++++-
21 files changed, 310 insertions(+), 52 deletions(-)
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 5e841c4..9f276a1 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -11052,10 +11052,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 2.5:
+ default ( none | present ) */
static tree
c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc)
@@ -11078,6 +11078,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;
@@ -11094,7 +11100,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/parser.c gcc/cp/parser.c
index 0be7ab5..70c187d 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -31426,10 +31426,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 2.5
+ default ( none | present ) */
static tree
cp_parser_omp_clause_default (cp_parser *parser, tree list,
@@ -31453,6 +31453,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;
@@ -31469,7 +31475,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/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 d594fb8..e5c3da3 100644
--- gcc/fortran/gfortran.h
+++ gcc/fortran/gfortran.h
@@ -1255,7 +1255,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 3bbeced..e954332 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 b6957c8..ce78e74 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 3fcf9ab..811a9b2 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)
@@ -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 ();
}
@@ -6933,8 +6937,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;
@@ -6947,8 +6956,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;
@@ -6968,6 +6982,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);
@@ -8685,11 +8701,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 0 | 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/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 2b1759e..a747165 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -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). */
@@ -499,6 +501,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 d823c2e..f0457a3 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, int 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/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..e8ead3d 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
For now committed to gomp-4_0-branch in r246763:
commit ed4aefe0cd00111abc233151a97cd57ea30fa842
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri Apr 7 15:07:46 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_1): 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".
(omp_default_clause, 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/branches/gomp-4_0-branch@246763 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 10 ++++
gcc/c/ChangeLog.gomp | 5 ++
gcc/c/c-parser.c | 14 ++++--
gcc/cp/ChangeLog.gomp | 5 ++
gcc/cp/parser.c | 14 ++++--
gcc/fortran/ChangeLog.gomp | 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 | 56 +++++++++++++++++-----
gcc/testsuite/ChangeLog.gomp | 10 ++++
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 | 5 +-
gcc/tree-pretty-print.c | 3 ++
libgomp/ChangeLog.gomp | 8 ++++
.../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 | 22 ++++-----
.../testsuite/libgomp.oacc-fortran/default-1.f90 | 10 ++++
.../libgomp.oacc-fortran/non-scalar-data.f90 | 44 ++++++++++++++++-
27 files changed, 357 insertions(+), 55 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 811e190..0aefea6 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,15 @@
2017-04-07 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".
+ (omp_default_clause, oacc_default_clause): Handle
+ "OMP_CLAUSE_DEFAULT_PRESENT".
+
* gimplify.c (oacc_default_clause): Clarify.
2017-04-05 Cesar Philippidis <cesar@codesourcery.com>
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index f6bb468..7406783 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2017-04-07 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-parser.c (c_parser_omp_clause_default): Handle
+ "OMP_CLAUSE_DEFAULT_PRESENT".
+
2017-02-27 Chung-Lin Tang <cltang@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 728c31b..b3c4634 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -10928,10 +10928,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 2.5:
+ default ( none | present ) */
static tree
c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc)
@@ -10954,6 +10954,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;
@@ -10970,7 +10976,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.gomp gcc/cp/ChangeLog.gomp
index 99ea570..f4f1ecd 100644
--- gcc/cp/ChangeLog.gomp
+++ gcc/cp/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2017-04-07 Thomas Schwinge <thomas@codesourcery.com>
+
+ * parser.c (cp_parser_omp_clause_default): Handle
+ "OMP_CLAUSE_DEFAULT_PRESENT".
+
2017-03-28 Thomas Schwinge <thomas@codesourcery.com>
* parser.c (cp_parser_omp_clause_default): Avoid printing more
diff --git gcc/cp/parser.c gcc/cp/parser.c
index d80d056..dbe9b27 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -30674,10 +30674,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 2.5
+ default ( none | present ) */
static tree
cp_parser_omp_clause_default (cp_parser *parser, tree list,
@@ -30701,6 +30701,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;
@@ -30717,7 +30723,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.gomp gcc/fortran/ChangeLog.gomp
index f29c6e2..cd1f62b 100644
--- gcc/fortran/ChangeLog.gomp
+++ gcc/fortran/ChangeLog.gomp
@@ -1,3 +1,11 @@
+2017-04-07 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_1): Likewise.
+
2017-04-05 Cesar Philippidis <cesar@codesourcery.com>
* gfortran.h (enum gfc_omp_map_op): Add OMP_MAP_DECLARE_ALLOCATE,
diff --git gcc/fortran/dump-parse-tree.c gcc/fortran/dump-parse-tree.c
index 1054514..db0ef94 100644
--- gcc/fortran/dump-parse-tree.c
+++ gcc/fortran/dump-parse-tree.c
@@ -1270,6 +1270,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 75217c7..7913ac7 100644
--- gcc/fortran/gfortran.h
+++ gcc/fortran/gfortran.h
@@ -1237,7 +1237,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 31e4885..868ad73 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -1130,13 +1130,19 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, 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 ba738a9..238eebe 100644
--- gcc/fortran/trans-openmp.c
+++ gcc/fortran/trans-openmp.c
@@ -2624,6 +2624,9 @@ gfc_trans_omp_clauses_1 (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 f2bb814..ff7674a 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -93,6 +93,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)
@@ -6092,6 +6095,7 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl,
found_outer:
break;
+ case OMP_CLAUSE_DEFAULT_PRESENT:
default:
gcc_unreachable ();
}
@@ -6135,8 +6139,13 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
if (is_private)
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 'copy'. */
flags |= GOVD_MAP | GOVD_MAP_FORCE;
@@ -6151,8 +6160,13 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
else 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;
@@ -6172,6 +6186,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);
@@ -7987,14 +8003,32 @@ 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);
- tree c2 = NULL_TREE;
- 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 0 | 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);
- c2 = gomp_needs_data_present (decl);
+ tree c2 = gomp_needs_data_present (decl);
/* Handle OpenACC pointers that were declared inside acc data
regions. */
if (c2 != NULL && OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_POINTER)
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 3071ab5..04ec34f 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,13 @@
+2017-04-07 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.
+
2017-04-06 Cesar Philippidis <cesar@codesourcery.com>
* gfortran.dg/goacc/declare-allocatable-1.f90: Correct test.
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 a3c4b18..d0f490c 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -351,7 +351,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). */
@@ -498,6 +500,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 aae409d..a36c938 100644
--- gcc/tree-pretty-print.c
+++ gcc/tree-pretty-print.c
@@ -501,6 +501,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int 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.gomp libgomp/ChangeLog.gomp
index aca124a..873668e 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,11 @@
+2017-04-07 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.
+
2017-04-05 Cesar Philippidis <cesar@codesourcery.com>
* libgomp.h: Declare gomp_acc_declare_allocate.
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..e8ead3d 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 d1ecf0a..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 present (a(1:N), b(1:N), N)
+ !$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) present (a(1:N), b(1:N), N)
+ !$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) present (a(1:N), b(1:N), N)
+ !$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) present (a(1:N), c(1:N), N)
+ !$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) present (a(1:N), d(1:N), N)
+ !$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,26 +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) present (a(1:N), b(1:N), N)
+ !$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) present (a(1:N), c(1:N), N)
+ !$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) present (a(1:N), d(1:N), N)
+ !$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) present (a(1:N), b(1:N), c(1:N)) &
- !$acc& present (d(1:N), e(1:N), N)
+ !$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 7d031c6..76f7fc1 100644
--- libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
@@ -53,4 +53,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 ab5771e..7e822e5 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 }
! { dg-additional-options "-foffload-force" }
@@ -13,6 +14,7 @@ program main
n = size
!$acc data copy(array)
+
call kernels(array, n)
!$acc update host(array)
@@ -21,12 +23,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)
@@ -40,6 +59,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
@@ -51,3 +80,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