* OpenACC 2.5 default (present) clause
@ 2017-04-07 15:09 Thomas Schwinge
2017-04-12 13:03 ` Thomas Schwinge
2017-05-10 15:49 ` Jakub Jelinek
0 siblings, 2 replies; 7+ messages in thread
From: Thomas Schwinge @ 2017-04-07 15:09 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek
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
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: OpenACC 2.5 default (present) clause
2017-04-07 15:09 OpenACC 2.5 default (present) clause Thomas Schwinge
@ 2017-04-12 13:03 ` Thomas Schwinge
2017-05-10 15:49 ` Jakub Jelinek
1 sibling, 0 replies; 7+ messages in thread
From: Thomas Schwinge @ 2017-04-12 13:03 UTC (permalink / raw)
To: gcc-patches
Hi!
On Fri, 07 Apr 2017 17:08:55 +0200, I 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". 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?
(Jakub asked me to ping this in next stage 1.)
> For now committed to gomp-4_0-branch in r246763:
> --- libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
> +++ libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
> +#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 ();
> + }
In our PowerPC testing, this change triggered linking failures (for host
comilation) and ICEs (for offloading compilation).
<https://gcc.gnu.org/PR80411> "DCE vs. offloading" filed (here, we got:
"char a[N]", causing GCC to figure out that "(char) -1 != (int) -1", thus
unconditionally "abort", thus the following OpenACC kernels getting DCEd,
thus PR80411), and committed to gomp-4_0-branch in r246871:
commit 6cf38f5ba09e4152336cf94d0eb9c80db9ffc024
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Apr 12 12:54:10 2017 +0000
Fix libgomp.oacc-c-c++-common/nested-2.c: "char" might mean "unsigned char"
libgomp/
* testsuite/libgomp.oacc-c-c++-common/nested-2.c: Respect that
"char" might mean "unsigned char".
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@246871 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog.gomp | 5 +++++
libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c | 8 ++++----
2 files changed, 9 insertions(+), 4 deletions(-)
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 873668e..32cb7e8 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2017-04-12 Thomas Schwinge <thomas@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/nested-2.c: Respect that
+ "char" might mean "unsigned char".
+
2017-04-07 Thomas Schwinge <thomas@codesourcery.com>
* testsuite/libgomp.oacc-c++/template-reduction.C: Update.
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
index e8ead3d..51b3b18 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
@@ -143,28 +143,28 @@ main (int argc, char *argv[])
#pragma acc parallel default (present)
{
for (int j = 0; j < N; ++j)
- a[j] = j - 1;
+ a[j] = j + 1;
}
#pragma acc update host (a)
for (i = 0; i < N; ++i)
{
- if (a[i] != i - 1)
+ if (a[i] != i + 1)
abort ();
}
#pragma acc kernels default (present)
{
for (int j = 0; j < N; ++j)
- a[j] = j - 2;
+ a[j] = j + 2;
}
#pragma acc exit data copyout (a)
for (i = 0; i < N; ++i)
{
- if (a[i] != i - 2)
+ if (a[i] != i + 2)
abort ();
}
Grüße
Thomas
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: OpenACC 2.5 default (present) clause
2017-04-07 15:09 OpenACC 2.5 default (present) clause Thomas Schwinge
2017-04-12 13:03 ` Thomas Schwinge
@ 2017-05-10 15:49 ` Jakub Jelinek
2017-05-19 13:47 ` Thomas Schwinge
2017-05-19 14:05 ` Thomas Schwinge
1 sibling, 2 replies; 7+ messages in thread
From: Jakub Jelinek @ 2017-05-10 15:49 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches
On Fri, Apr 07, 2017 at 05:08:55PM +0200, Thomas Schwinge wrote:
> 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(-)
>
> @@ -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? I don't think that is needed.
> + case 0 | GOVD_MAP_FORCE:
> + kind = GOMP_MAP_TOFROM | GOMP_MAP_FLAG_FORCE;
> + break;
Please drop the "0 | ".
> @@ -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).
Otherwise LGTM.
Jakub
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: OpenACC 2.5 default (present) clause
2017-05-10 15:49 ` Jakub Jelinek
@ 2017-05-19 13:47 ` Thomas Schwinge
2017-05-19 14:05 ` Thomas Schwinge
1 sibling, 0 replies; 7+ messages in thread
From: Thomas Schwinge @ 2017-05-19 13:47 UTC (permalink / raw)
To: Jakub Jelinek, gcc-patches
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". 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?
> > Clarify gcc/gimplify.c:oacc_default_clause
> [...] LGTM.
As posted, committed to trunk in r248279:
commit f7c10d53cac5a46d06d23177526aa36c021e135c
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri May 19 13:32:30 2017 +0000
Clarify gcc/gimplify.c:oacc_default_clause
gcc/
* gimplify.c (oacc_default_clause): Clarify.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@248279 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 4 ++++
gcc/gimplify.c | 40 ++++++++++++++++++++++------------------
2 files changed, 26 insertions(+), 18 deletions(-)
diff --git gcc/ChangeLog gcc/ChangeLog
index 44ff617..8f63902 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,3 +1,7 @@
+2017-05-19 Thomas Schwinge <thomas@codesourcery.com>
+
+ * gimplify.c (oacc_default_clause): Clarify.
+
2017-05-19 Nathan Sidwell <nathan@acm.org>
LANG_HOOK_REGISTER_DUMPS
diff --git gcc/gimplify.c gcc/gimplify.c
index acaab8b..0c02ee4 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -6952,30 +6952,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))
Grüße
Thomas
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: OpenACC 2.5 default (present) clause
2017-05-10 15:49 ` Jakub Jelinek
2017-05-19 13:47 ` Thomas Schwinge
@ 2017-05-19 14:05 ` Thomas Schwinge
2017-05-23 10:41 ` Jakub Jelinek
1 sibling, 1 reply; 7+ messages in thread
From: Thomas Schwinge @ 2017-05-19 14:05 UTC (permalink / raw)
To: Jakub Jelinek, gcc-patches
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
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: OpenACC 2.5 default (present) clause
2017-05-19 14:05 ` Thomas Schwinge
@ 2017-05-23 10:41 ` Jakub Jelinek
2017-05-23 16:08 ` Thomas Schwinge
0 siblings, 1 reply; 7+ messages in thread
From: Jakub Jelinek @ 2017-05-23 10:41 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches
On Fri, May 19, 2017 at 03:46:53PM +0200, Thomas Schwinge wrote:
> > > - /* 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?
Because the clause in OpenMP is default ( shared | none ) for C/C++
and default ( private | firstprivate | shared | none ) for Fortran.
Jakub
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: OpenACC 2.5 default (present) clause
2017-05-23 10:41 ` Jakub Jelinek
@ 2017-05-23 16:08 ` Thomas Schwinge
0 siblings, 0 replies; 7+ messages in thread
From: Thomas Schwinge @ 2017-05-23 16:08 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: gcc-patches
Hi!
On Tue, 23 May 2017 12:40:06 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Fri, May 19, 2017 at 03:46:53PM +0200, Thomas Schwinge wrote:
> > > > - /* 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?
>
> Because the clause in OpenMP is default ( shared | none ) for C/C++
> and default ( private | firstprivate | shared | none ) for Fortran.
So, the union of the C/C++ and Fortran variants is "private |
firstprivate | shared | none", and sorting these alphabetically, you'd
get "firstprivate | none | private | shared" as I posted? Anyway, no
point on spending more time on this. ;-)
Grüße
Thomas
^ permalink raw reply [flat|nested] 7+ messages in thread
end of thread, other threads:[~2017-05-23 16:06 UTC | newest]
Thread overview: 7+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2017-04-07 15:09 OpenACC 2.5 default (present) clause Thomas Schwinge
2017-04-12 13:03 ` Thomas Schwinge
2017-05-10 15:49 ` Jakub Jelinek
2017-05-19 13:47 ` Thomas Schwinge
2017-05-19 14:05 ` Thomas Schwinge
2017-05-23 10:41 ` Jakub Jelinek
2017-05-23 16:08 ` Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).