diff options
author | Thomas Schwinge <thomas@codesourcery.com> | 2017-05-19 15:32:48 +0200 |
---|---|---|
committer | Thomas Schwinge <tschwinge@gcc.gnu.org> | 2017-05-19 15:32:48 +0200 |
commit | 7fd549d24fda05c859fb17697c51c16886902dad (patch) | |
tree | a1ddeff9f474c8969d68b1b0b726f4ca763411f3 /gcc | |
parent | 0d0afa9fafec1711b52259b9b8caebf51380216d (diff) | |
download | gcc-7fd549d24fda05c859fb17697c51c16886902dad.zip gcc-7fd549d24fda05c859fb17697c51c16886902dad.tar.gz gcc-7fd549d24fda05c859fb17697c51c16886902dad.tar.bz2 |
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.
From-SVN: r248280
Diffstat (limited to 'gcc')
-rw-r--r-- | gcc/ChangeLog | 9 | ||||
-rw-r--r-- | gcc/c/ChangeLog | 5 | ||||
-rw-r--r-- | gcc/c/c-parser.c | 14 | ||||
-rw-r--r-- | gcc/cp/ChangeLog | 3 | ||||
-rw-r--r-- | gcc/cp/parser.c | 14 | ||||
-rw-r--r-- | gcc/fortran/ChangeLog | 8 | ||||
-rw-r--r-- | gcc/fortran/dump-parse-tree.c | 1 | ||||
-rw-r--r-- | gcc/fortran/gfortran.h | 3 | ||||
-rw-r--r-- | gcc/fortran/openmp.c | 20 | ||||
-rw-r--r-- | gcc/fortran/trans-openmp.c | 3 | ||||
-rw-r--r-- | gcc/gimplify.c | 52 | ||||
-rw-r--r-- | gcc/testsuite/ChangeLog | 8 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/default-1.c | 5 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/default-2.c | 28 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/default-4.c | 21 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/default-5.c | 20 | ||||
-rw-r--r-- | gcc/testsuite/gfortran.dg/goacc/default-1.f95 | 5 | ||||
-rw-r--r-- | gcc/testsuite/gfortran.dg/goacc/default-4.f | 18 | ||||
-rw-r--r-- | gcc/testsuite/gfortran.dg/goacc/default-5.f | 18 | ||||
-rw-r--r-- | gcc/tree-core.h | 3 | ||||
-rw-r--r-- | gcc/tree-pretty-print.c | 3 |
21 files changed, 221 insertions, 40 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 8f63902..c40af47 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,5 +1,14 @@ 2017-05-19 Thomas Schwinge <thomas@codesourcery.com> + * tree-core.h (enum omp_clause_default_kind): Add + "OMP_CLAUSE_DEFAULT_PRESENT". + * tree-pretty-print.c (dump_omp_clause): Handle it. + * gimplify.c (enum gimplify_omp_var_data): Add + "GOVD_MAP_FORCE_PRESENT". + (gimplify_adjust_omp_clauses_1): Map it to + "GOMP_MAP_FORCE_PRESENT". + (oacc_default_clause): Handle "OMP_CLAUSE_DEFAULT_PRESENT". + * gimplify.c (oacc_default_clause): Clarify. 2017-05-19 Nathan Sidwell <nathan@acm.org> diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index 79643cc..a70eaf9 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,8 @@ +2017-05-19 Thomas Schwinge <thomas@codesourcery.com> + + * c-parser.c (c_parser_omp_clause_default): Handle + "OMP_CLAUSE_DEFAULT_PRESENT". + 2017-05-18 Bernd Edlinger <bernd.edlinger@hotmail.de> * config-lang.in (gtfiles): Add c-family/c-format.c. diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 96c0749..2e01316 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -11057,10 +11057,10 @@ c_parser_omp_clause_copyprivate (c_parser *parser, tree list) } /* OpenMP 2.5: - default ( shared | none ) + default ( none | shared ) - OpenACC 2.0: - default (none) */ + OpenACC: + default ( none | present ) */ static tree c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc) @@ -11083,6 +11083,12 @@ c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc) kind = OMP_CLAUSE_DEFAULT_NONE; break; + case 'p': + if (strcmp ("present", p) != 0 || !is_oacc) + goto invalid_kind; + kind = OMP_CLAUSE_DEFAULT_PRESENT; + break; + case 's': if (strcmp ("shared", p) != 0 || is_oacc) goto invalid_kind; @@ -11099,7 +11105,7 @@ c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc) { invalid_kind: if (is_oacc) - c_parser_error (parser, "expected %<none%>"); + c_parser_error (parser, "expected %<none%> or %<present%>"); else c_parser_error (parser, "expected %<none%> or %<shared%>"); } diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index c89f719..f2dced4 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,5 +1,8 @@ 2017-05-19 Thomas Schwinge <thomas@codesourcery.com> + * parser.c (cp_parser_omp_clause_default): Handle + "OMP_CLAUSE_DEFAULT_PRESENT". + * parser.c (cp_parser_omp_clause_default): Avoid printing more than one syntax error message. diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 6453397..b345110 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -31456,10 +31456,10 @@ cp_parser_omp_clause_collapse (cp_parser *parser, tree list, location_t location } /* OpenMP 2.5: - default ( shared | none ) + default ( none | shared ) - OpenACC 2.0 - default (none) */ + OpenACC: + default ( none | present ) */ static tree cp_parser_omp_clause_default (cp_parser *parser, tree list, @@ -31483,6 +31483,12 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list, kind = OMP_CLAUSE_DEFAULT_NONE; break; + case 'p': + if (strcmp ("present", p) != 0 || !is_oacc) + goto invalid_kind; + kind = OMP_CLAUSE_DEFAULT_PRESENT; + break; + case 's': if (strcmp ("shared", p) != 0 || is_oacc) goto invalid_kind; @@ -31499,7 +31505,7 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list, { invalid_kind: if (is_oacc) - cp_parser_error (parser, "expected %<none%>"); + cp_parser_error (parser, "expected %<none%> or %<present%>"); else cp_parser_error (parser, "expected %<none%> or %<shared%>"); } diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog index 9a8b3e1..4b7f1f4 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,11 @@ +2017-05-19 Thomas Schwinge <thomas@codesourcery.com> + + * gfortran.h (enum gfc_omp_default_sharing): Add + "OMP_DEFAULT_PRESENT". + * dump-parse-tree.c (show_omp_clauses): Handle it. + * openmp.c (gfc_match_omp_clauses): Likewise. + * trans-openmp.c (gfc_trans_omp_clauses): Likewise. + 2017-05-18 Fritz Reese <fritzoreese@gmail.com> PR fortran/79968 diff --git a/gcc/fortran/dump-parse-tree.c b/gcc/fortran/dump-parse-tree.c index 87a5304..49b23d8 100644 --- a/gcc/fortran/dump-parse-tree.c +++ b/gcc/fortran/dump-parse-tree.c @@ -1283,6 +1283,7 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses) case OMP_DEFAULT_PRIVATE: type = "PRIVATE"; break; case OMP_DEFAULT_SHARED: type = "SHARED"; break; case OMP_DEFAULT_FIRSTPRIVATE: type = "FIRSTPRIVATE"; break; + case OMP_DEFAULT_PRESENT: type = "PRESENT"; break; default: gcc_unreachable (); } diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 2936550..26b89be 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1241,7 +1241,8 @@ enum gfc_omp_default_sharing OMP_DEFAULT_NONE, OMP_DEFAULT_PRIVATE, OMP_DEFAULT_SHARED, - OMP_DEFAULT_FIRSTPRIVATE + OMP_DEFAULT_FIRSTPRIVATE, + OMP_DEFAULT_PRESENT }; enum gfc_omp_proc_bind_kind diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 89eecfa..80146e2 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -1080,13 +1080,19 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if (gfc_match ("default ( none )") == MATCH_YES) c->default_sharing = OMP_DEFAULT_NONE; else if (openacc) - /* c->default_sharing = OMP_DEFAULT_UNKNOWN */; - else if (gfc_match ("default ( shared )") == MATCH_YES) - c->default_sharing = OMP_DEFAULT_SHARED; - else if (gfc_match ("default ( private )") == MATCH_YES) - c->default_sharing = OMP_DEFAULT_PRIVATE; - else if (gfc_match ("default ( firstprivate )") == MATCH_YES) - c->default_sharing = OMP_DEFAULT_FIRSTPRIVATE; + { + if (gfc_match ("default ( present )") == MATCH_YES) + c->default_sharing = OMP_DEFAULT_PRESENT; + } + else + { + if (gfc_match ("default ( firstprivate )") == MATCH_YES) + c->default_sharing = OMP_DEFAULT_FIRSTPRIVATE; + else if (gfc_match ("default ( private )") == MATCH_YES) + c->default_sharing = OMP_DEFAULT_PRIVATE; + else if (gfc_match ("default ( shared )") == MATCH_YES) + c->default_sharing = OMP_DEFAULT_SHARED; + } if (c->default_sharing != OMP_DEFAULT_UNKNOWN) continue; } diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 662036f..1d254c6 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2564,6 +2564,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, case OMP_DEFAULT_FIRSTPRIVATE: OMP_CLAUSE_DEFAULT_KIND (c) = OMP_CLAUSE_DEFAULT_FIRSTPRIVATE; break; + case OMP_DEFAULT_PRESENT: + OMP_CLAUSE_DEFAULT_KIND (c) = OMP_CLAUSE_DEFAULT_PRESENT; + break; default: gcc_unreachable (); } diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 0c02ee4..810d9f4 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -99,6 +99,9 @@ enum gimplify_omp_var_data /* Flag for GOVD_MAP, if it is a forced mapping. */ GOVD_MAP_FORCE = 262144, + /* Flag for GOVD_MAP: must be present already. */ + GOVD_MAP_FORCE_PRESENT = 524288, + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -6956,8 +6959,13 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) rkind = "kernels"; if (AGGREGATE_TYPE_P (type)) - /* Aggregates default to 'present_or_copy'. */ - flags |= GOVD_MAP; + { + /* Aggregates default to 'present_or_copy', or 'present'. */ + if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT) + flags |= GOVD_MAP; + else + flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT; + } else /* Scalars default to 'copy'. */ flags |= GOVD_MAP | GOVD_MAP_FORCE; @@ -6970,8 +6978,13 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) if (on_device || declared) flags |= GOVD_MAP; else if (AGGREGATE_TYPE_P (type)) - /* Aggregates default to 'present_or_copy'. */ - flags |= GOVD_MAP; + { + /* Aggregates default to 'present_or_copy', or 'present'. */ + if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT) + flags |= GOVD_MAP; + else + flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT; + } else /* Scalars default to 'firstprivate'. */ flags |= GOVD_FIRSTPRIVATE; @@ -6991,6 +7004,8 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) DECL_NAME (lang_hooks.decls.omp_report_decl (decl)), rkind); inform (ctx->location, "enclosing OpenACC %qs construct", rkind); } + else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_PRESENT) + ; /* Handled above. */ else gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED); @@ -8708,11 +8723,30 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) } else if (code == OMP_CLAUSE_MAP) { - int kind = (flags & GOVD_MAP_TO_ONLY - ? GOMP_MAP_TO - : GOMP_MAP_TOFROM); - if (flags & GOVD_MAP_FORCE) - kind |= GOMP_MAP_FLAG_FORCE; + int kind; + /* Not all combinations of these GOVD_MAP flags are actually valid. */ + switch (flags & (GOVD_MAP_TO_ONLY + | GOVD_MAP_FORCE + | GOVD_MAP_FORCE_PRESENT)) + { + case 0: + kind = GOMP_MAP_TOFROM; + break; + case GOVD_MAP_FORCE: + kind = GOMP_MAP_TOFROM | GOMP_MAP_FLAG_FORCE; + break; + case GOVD_MAP_TO_ONLY: + kind = GOMP_MAP_TO; + break; + case GOVD_MAP_TO_ONLY | GOVD_MAP_FORCE: + kind = GOMP_MAP_TO | GOMP_MAP_FLAG_FORCE; + break; + case GOVD_MAP_FORCE_PRESENT: + kind = GOMP_MAP_FORCE_PRESENT; + break; + default: + gcc_unreachable (); + } OMP_CLAUSE_SET_MAP_KIND (clause, kind); if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index ca901c2..e8e2df5 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,5 +1,13 @@ 2017-05-19 Thomas Schwinge <thomas@codesourcery.com> + * c-c++-common/goacc/default-1.c: Update. + * c-c++-common/goacc/default-2.c: Likewise. + * c-c++-common/goacc/default-4.c: Likewise. + * gfortran.dg/goacc/default-1.f95: Likewise. + * gfortran.dg/goacc/default-4.f: Likewise. + * c-c++-common/goacc/default-5.c: New file. + * gfortran.dg/goacc/default-5.f: Likewise. + * c-c++-common/goacc/default-1.c: New file. * c-c++-common/goacc/default-2.c: Likewise. * c-c++-common/goacc/data-default-1.c: Remove file, including its diff --git a/gcc/testsuite/c-c++-common/goacc/default-1.c b/gcc/testsuite/c-c++-common/goacc/default-1.c index 4d31dbc..23c25ab 100644 --- a/gcc/testsuite/c-c++-common/goacc/default-1.c +++ b/gcc/testsuite/c-c++-common/goacc/default-1.c @@ -6,4 +6,9 @@ void f1 () ; #pragma acc parallel default (none) ; + +#pragma acc kernels default (present) + ; +#pragma acc parallel default (present) + ; } diff --git a/gcc/testsuite/c-c++-common/goacc/default-2.c b/gcc/testsuite/c-c++-common/goacc/default-2.c index d6b6cdc..c0cdd12 100644 --- a/gcc/testsuite/c-c++-common/goacc/default-2.c +++ b/gcc/testsuite/c-c++-common/goacc/default-2.c @@ -7,39 +7,39 @@ void f1 () #pragma acc parallel default /* { dg-error "expected .\\(. before end of line" } */ ; -#pragma acc kernels default ( /* { dg-error "expected .none. before end of line" } */ +#pragma acc kernels default ( /* { dg-error "expected .none. or .present. before end of line" } */ ; -#pragma acc parallel default ( /* { dg-error "expected .none. before end of line" } */ +#pragma acc parallel default ( /* { dg-error "expected .none. or .present. before end of line" } */ ; -#pragma acc kernels default (, /* { dg-error "expected .none. before .,. token" } */ +#pragma acc kernels default (, /* { dg-error "expected .none. or .present. before .,. token" } */ ; -#pragma acc parallel default (, /* { dg-error "expected .none. before .,. token" } */ +#pragma acc parallel default (, /* { dg-error "expected .none. or .present. before .,. token" } */ ; -#pragma acc kernels default () /* { dg-error "expected .none. before .\\). token" } */ +#pragma acc kernels default () /* { dg-error "expected .none. or .present. before .\\). token" } */ ; -#pragma acc parallel default () /* { dg-error "expected .none. before .\\). token" } */ +#pragma acc parallel default () /* { dg-error "expected .none. or .present. before .\\). token" } */ ; -#pragma acc kernels default (,) /* { dg-error "expected .none. before .,. token" } */ +#pragma acc kernels default (,) /* { dg-error "expected .none. or .present. before .,. token" } */ ; -#pragma acc parallel default (,) /* { dg-error "expected .none. before .,. token" } */ +#pragma acc parallel default (,) /* { dg-error "expected .none. or .present. before .,. token" } */ ; -#pragma acc kernels default (firstprivate) /* { dg-error "expected .none. before .firstprivate." } */ +#pragma acc kernels default (firstprivate) /* { dg-error "expected .none. or .present. before .firstprivate." } */ ; -#pragma acc parallel default (firstprivate) /* { dg-error "expected .none. before .firstprivate." } */ +#pragma acc parallel default (firstprivate) /* { dg-error "expected .none. or .present. before .firstprivate." } */ ; -#pragma acc kernels default (private) /* { dg-error "expected .none. before .private." } */ +#pragma acc kernels default (private) /* { dg-error "expected .none. or .present. before .private." } */ ; -#pragma acc parallel default (private) /* { dg-error "expected .none. before .private." } */ +#pragma acc parallel default (private) /* { dg-error "expected .none. or .present. before .private." } */ ; -#pragma acc kernels default (shared) /* { dg-error "expected .none. before .shared." } */ +#pragma acc kernels default (shared) /* { dg-error "expected .none. or .present. before .shared." } */ ; -#pragma acc parallel default (shared) /* { dg-error "expected .none. before .shared." } */ +#pragma acc parallel default (shared) /* { dg-error "expected .none. or .present. before .shared." } */ ; #pragma acc kernels default (none /* { dg-error "expected .\\). before end of line" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/default-4.c b/gcc/testsuite/c-c++-common/goacc/default-4.c index 787b352..dfa79bb 100644 --- a/gcc/testsuite/c-c++-common/goacc/default-4.c +++ b/gcc/testsuite/c-c++-common/goacc/default-4.c @@ -43,3 +43,24 @@ void f2 () } } } + +void f3 () +{ + int f3_a = 2; + float f3_b[2]; + +#pragma acc data copyin (f3_a) copyout (f3_b) + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f3_b \[^\\)\]+\\) map\\(force_to:f3_a" 1 "gimple" } } */ + { +#pragma acc kernels default (present) + /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */ + { + f3_b[0] = f3_a; + } +#pragma acc parallel default (present) + /* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */ + { + f3_b[0] = f3_a; + } + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/default-5.c b/gcc/testsuite/c-c++-common/goacc/default-5.c new file mode 100644 index 0000000..37e3c355 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/default-5.c @@ -0,0 +1,20 @@ +/* OpenACC default (present) clause. */ + +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void f1 () +{ + int f1_a = 2; + float f1_b[2]; + +#pragma acc kernels default (present) + /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_tofrom:f1_a" 1 "gimple" } } */ + { + f1_b[0] = f1_a; + } +#pragma acc parallel default (present) + /* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) firstprivate\\(f1_a\\)" 1 "gimple" } } */ + { + f1_b[0] = f1_a; + } +} diff --git a/gcc/testsuite/gfortran.dg/goacc/default-1.f95 b/gcc/testsuite/gfortran.dg/goacc/default-1.f95 index a79b444..41fa944 100644 --- a/gcc/testsuite/gfortran.dg/goacc/default-1.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/default-1.f95 @@ -7,4 +7,9 @@ subroutine f1 !$acc end kernels !$acc parallel default (none) !$acc end parallel + + !$acc kernels default (present) + !$acc end kernels + !$acc parallel default (present) + !$acc end parallel end subroutine f1 diff --git a/gcc/testsuite/gfortran.dg/goacc/default-4.f b/gcc/testsuite/gfortran.dg/goacc/default-4.f index b2f73c3..77291f4 100644 --- a/gcc/testsuite/gfortran.dg/goacc/default-4.f +++ b/gcc/testsuite/gfortran.dg/goacc/default-4.f @@ -37,3 +37,21 @@ !$ACC END PARALLEL !$ACC END DATA END SUBROUTINE F2 + + SUBROUTINE F3 + IMPLICIT NONE + INTEGER :: F3_A = 2 + REAL, DIMENSION (2) :: F3_B + +!$ACC DATA COPYIN (F3_A) COPYOUT (F3_B) +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f3_a \[^\\)\]+\\) map\\(force_from:f3_b" 1 "gimple" } } +!$ACC KERNELS DEFAULT (PRESENT) +! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } + F3_B(1) = F3_A; +!$ACC END KERNELS +!$ACC PARALLEL DEFAULT (PRESENT) +! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } + F3_B(1) = F3_A; +!$ACC END PARALLEL +!$ACC END DATA + END SUBROUTINE F3 diff --git a/gcc/testsuite/gfortran.dg/goacc/default-5.f b/gcc/testsuite/gfortran.dg/goacc/default-5.f new file mode 100644 index 0000000..9dc83cb --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/default-5.f @@ -0,0 +1,18 @@ +! OpenACC default (present) clause. + +! { dg-additional-options "-fdump-tree-gimple" } + + SUBROUTINE F1 + IMPLICIT NONE + INTEGER :: F1_A = 2 + REAL, DIMENSION (2) :: F1_B + +!$ACC KERNELS DEFAULT (PRESENT) +! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_tofrom:f1_a" 1 "gimple" } } + F1_B(1) = F1_A; +!$ACC END KERNELS +!$ACC PARALLEL DEFAULT (PRESENT) +! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) firstprivate\\(f1_a\\)" 1 "gimple" } } + F1_B(1) = F1_A; +!$ACC END PARALLEL + END SUBROUTINE F1 diff --git a/gcc/tree-core.h b/gcc/tree-core.h index c76fc7b..ea73477 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -357,7 +357,7 @@ enum omp_clause_code { /* OpenMP clause: ordered [(constant-integer-expression)]. */ OMP_CLAUSE_ORDERED, - /* OpenMP clause: default. */ + /* OpenACC/OpenMP clause: default. */ OMP_CLAUSE_DEFAULT, /* OpenACC/OpenMP clause: collapse (constant-integer-expression). */ @@ -499,6 +499,7 @@ enum omp_clause_default_kind { OMP_CLAUSE_DEFAULT_NONE, OMP_CLAUSE_DEFAULT_PRIVATE, OMP_CLAUSE_DEFAULT_FIRSTPRIVATE, + OMP_CLAUSE_DEFAULT_PRESENT, OMP_CLAUSE_DEFAULT_LAST }; diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index ec28b1e..b70e325 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -502,6 +502,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case OMP_CLAUSE_DEFAULT_FIRSTPRIVATE: pp_string (pp, "firstprivate"); break; + case OMP_CLAUSE_DEFAULT_PRESENT: + pp_string (pp, "present"); + break; default: gcc_unreachable (); } |