From e03436e7ac2ddbbf397a6d64309b01ad37cfcadf Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Fri, 17 May 2019 21:13:04 +0200 Subject: [PATCH] [PR89433] Refer to OpenACC 'routine' clauses from "omp declare target" attribute gcc/c-family/ PR c/89433 * c-attribs.c (c_common_attribute_table): Set min_len to -1 for "omp declare target". gcc/c/ PR c/89433 * c-parser.c (c_finish_oacc_routine): Refer to OpenACC 'routine' clauses from "omp declare target" attribute. gcc/cp/ PR c++/89433 * parser.c (cp_finalize_oacc_routine): Refer to OpenACC 'routine' clauses from "omp declare target" attribute. gcc/fortran/ PR fortran/89433 * f95-lang.c (gfc_attribute_table): Set min_len to -1 for "omp declare target". * trans-decl.c (add_attributes_to_decl): Refer to OpenACC 'routine' clauses from "omp declare target" attribute. gcc/testsuite/ PR testsuite/89433 * c-c++-common/goacc/classify-routine.c: Update. * gfortran.dg/goacc/classify-routine.f95: Likewise. From-SVN: r271343 --- gcc/c-family/ChangeLog | 6 ++++++ gcc/c-family/c-attribs.c | 2 +- gcc/c/ChangeLog | 6 ++++++ gcc/c/c-parser.c | 2 +- gcc/cp/ChangeLog | 6 ++++++ gcc/cp/parser.c | 2 +- gcc/fortran/ChangeLog | 8 ++++++++ gcc/fortran/f95-lang.c | 2 +- gcc/fortran/trans-decl.c | 18 +++++++++++------- gcc/testsuite/ChangeLog | 6 ++++++ .../c-c++-common/goacc/classify-routine.c | 4 ++-- .../gfortran.dg/goacc/classify-routine.f95 | 4 ++-- 12 files changed, 51 insertions(+), 15 deletions(-) diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog index 47c1d3d51f5..50687764221 100644 --- a/gcc/c-family/ChangeLog +++ b/gcc/c-family/ChangeLog @@ -1,3 +1,9 @@ +2019-05-17 Thomas Schwinge + + PR c/89433 + * c-attribs.c (c_common_attribute_table): Set min_len to -1 for + "omp declare target". + 2019-05-16 Martin Sebor * c-attribs.c (handle_no_sanitize_attribute): Quote identifiers, diff --git a/gcc/c-family/c-attribs.c b/gcc/c-family/c-attribs.c index 12c0b9bfb54..03203470955 100644 --- a/gcc/c-family/c-attribs.c +++ b/gcc/c-family/c-attribs.c @@ -437,7 +437,7 @@ const struct attribute_spec c_common_attribute_table[] = handle_omp_declare_simd_attribute, NULL }, { "simd", 0, 1, true, false, false, false, handle_simd_attribute, NULL }, - { "omp declare target", 0, 0, true, false, false, false, + { "omp declare target", 0, -1, true, false, false, false, handle_omp_declare_target_attribute, NULL }, { "omp declare target link", 0, 0, true, false, false, false, handle_omp_declare_target_attribute, NULL }, diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index f558d28dc25..f0cab2e6592 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,9 @@ +2019-05-17 Thomas Schwinge + + PR c/89433 + * c-parser.c (c_finish_oacc_routine): Refer to OpenACC 'routine' + clauses from "omp declare target" attribute. + 2019-05-16 Martin Sebor * c-decl.c (start_decl): Quote keywords, operators, and diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 993cfe05ecb..3cbbb199bdd 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -15904,7 +15904,7 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl, /* Add an "omp declare target" attribute. */ DECL_ATTRIBUTES (fndecl) = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, DECL_ATTRIBUTES (fndecl)); + data->clauses, DECL_ATTRIBUTES (fndecl)); /* Remember that we've used this "#pragma acc routine". */ data->fndecl_seen = true; diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 08b7d533707..40622acc0ff 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,9 @@ +2019-05-17 Thomas Schwinge + + PR c++/89433 + * parser.c (cp_finalize_oacc_routine): Refer to OpenACC 'routine' + clauses from "omp declare target" attribute. + 2019-05-16 Martin Sebor * call.c (print_z_candidate): Wrap diagnostic text in a gettext diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index e6ef5a9bc00..15424b6b633 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40292,7 +40292,7 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn) /* Add an "omp declare target" attribute. */ DECL_ATTRIBUTES (fndecl) = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, DECL_ATTRIBUTES (fndecl)); + parser->oacc_routine->clauses, DECL_ATTRIBUTES (fndecl)); /* Don't unset parser->oacc_routine here: we may still need it to diagnose wrong usage. But, remember that we've used this "#pragma acc diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog index b8851067cf3..f09e715353b 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,11 @@ +2019-05-17 Thomas Schwinge + + PR fortran/89433 + * f95-lang.c (gfc_attribute_table): Set min_len to -1 for "omp + declare target". + * trans-decl.c (add_attributes_to_decl): Refer to OpenACC + 'routine' clauses from "omp declare target" attribute. + 2019-05-16 Martin Sebor * gfortranspec.c (append_arg): Spell out the word "argument." diff --git a/gcc/fortran/f95-lang.c b/gcc/fortran/f95-lang.c index 3e3d3046bdb..6b9f490d2bb 100644 --- a/gcc/fortran/f95-lang.c +++ b/gcc/fortran/f95-lang.c @@ -91,7 +91,7 @@ static const struct attribute_spec gfc_attribute_table[] = { /* { name, min_len, max_len, decl_req, type_req, fn_type_req, affects_type_identity, handler, exclude } */ - { "omp declare target", 0, 0, true, false, false, false, + { "omp declare target", 0, -1, true, false, false, false, gfc_handle_omp_declare_target_attribute, NULL }, { "omp declare target link", 0, 0, true, false, false, false, gfc_handle_omp_declare_target_attribute, NULL }, diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c index 07d1c33af72..8420870a6b7 100644 --- a/gcc/fortran/trans-decl.c +++ b/gcc/fortran/trans-decl.c @@ -1400,12 +1400,7 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list) list = chainon (list, attr); } - if (sym_attr.omp_declare_target_link) - list = tree_cons (get_identifier ("omp declare target link"), - NULL_TREE, list); - else if (sym_attr.omp_declare_target) - list = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, list); + tree clauses = NULL_TREE; if (sym_attr.oacc_routine_lop != OACC_ROUTINE_LOP_NONE) { @@ -1430,11 +1425,20 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list) gcc_unreachable (); } tree c = build_omp_clause (UNKNOWN_LOCATION, code); + OMP_CLAUSE_CHAIN (c) = clauses; + clauses = c; - tree dims = oacc_build_routine_dims (c); + tree dims = oacc_build_routine_dims (clauses); list = oacc_replace_fn_attrib_attr (list, dims); } + if (sym_attr.omp_declare_target_link) + list = tree_cons (get_identifier ("omp declare target link"), + NULL_TREE, list); + else if (sym_attr.omp_declare_target) + list = tree_cons (get_identifier ("omp declare target"), + clauses, list); + return list; } diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index e1aa25736ef..0f44c6d4ab0 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,9 @@ +2019-05-17 Thomas Schwinge + + PR testsuite/89433 + * c-c++-common/goacc/classify-routine.c: Update. + * gfortran.dg/goacc/classify-routine.f95: Likewise. + 2019-05-16 Martin Sebor * c-c++-common/Wbool-operation-1.c: Adjust text of expected diagnostics. diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c index a723d2cdf51..0b9ba6ea69f 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c @@ -22,10 +22,10 @@ void ROUTINE () } /* Check the offloaded function's attributes. - { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */ + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */ /* Check the offloaded function's classification and compute dimensions (will always be 1 x 1 x 1 for non-offloading compilation). { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } } { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } } */ + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 index e435f5d7eae..401d5270391 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 @@ -21,10 +21,10 @@ subroutine ROUTINE end subroutine ROUTINE ! Check the offloaded function's attributes. -! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 0, 1 0\\), omp declare target \\(worker\\)\\)\\)" 1 "ompexp" } } ! Check the offloaded function's classification and compute dimensions (will ! always be 1 x 1 x 1 for non-offloading compilation). ! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } } ! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\)\\)\\)" 1 "oaccdevlow" } } -- 2.30.2