@@ -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(-)
@@ -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%>");
}
@@ -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%>");
}
@@ -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 ();
}
@@ -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
@@ -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;
}
@@ -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 ();
}
@@ -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)
@@ -6,4 +6,9 @@ void f1 ()
;
#pragma acc parallel default (none)
;
+
+#pragma acc kernels default (present)
+ ;
+#pragma acc parallel default (present)
+ ;
}
@@ -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" } */
@@ -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;
+ }
+ }
+}
new file mode 100644
@@ -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;
+ }
+}
@@ -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
@@ -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
new file mode 100644
@@ -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
@@ -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
};
@@ -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 ();
}
@@ -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 ();
@@ -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;
}
@@ -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
@@ -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
@@ -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(-)
@@ -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>
@@ -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>
@@ -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%>");
}
@@ -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
@@ -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%>");
}
@@ -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,
@@ -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 ();
}
@@ -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
@@ -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;
}
@@ -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 ();
}
@@ -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)
@@ -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.
@@ -6,4 +6,9 @@ void f1 ()
;
#pragma acc parallel default (none)
;
+
+#pragma acc kernels default (present)
+ ;
+#pragma acc parallel default (present)
+ ;
}
@@ -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" } */
@@ -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;
+ }
+ }
+}
new file mode 100644
@@ -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;
+ }
+}
@@ -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
@@ -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
new file mode 100644
@@ -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
@@ -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
};
@@ -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 ();
}
@@ -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.
@@ -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 ();
@@ -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;
}
@@ -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
@@ -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
@@ -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