public inbox for gcc-cvs@sourceware.org help / color / mirror / Atom feed
From: Tobias Burnus <burnus@gcc.gnu.org> To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-11] openmp: Fix up declare target handling for vars with DECL_LOCAL_DECL_ALIAS [PR102640] Date: Fri, 8 Oct 2021 11:35:25 +0000 (GMT) [thread overview] Message-ID: <20211008113525.4AEBD385803F@sourceware.org> (raw) https://gcc.gnu.org/g:cb639f7bf2e56acdb3c7832fad5cc15ded7db441 commit cb639f7bf2e56acdb3c7832fad5cc15ded7db441 Author: Jakub Jelinek <jakub@redhat.com> Date: Fri Oct 8 13:30:26 2021 +0200 openmp: Fix up declare target handling for vars with DECL_LOCAL_DECL_ALIAS [PR102640] The introduction of DECL_LOCAL_DECL_ALIAS and push_local_extern_decl_alias in r11-3699-g4e62aca0e0520e4ed2532f2d8153581190621c1a broke the following testcase. The following patch fixes it by treating similarly not just the variable to or link clause is put on, but also its DECL_LOCAL_DECL_ALIAS if any. If it hasn't been created yet, when it is created it will copy attributes and therefore should get it for free, and as it is an extern, nothing more than attributes is needed for it. 2021-10-08 Jakub Jelinek <jakub@redhat.com> PR c++/102640 gcc/cp/ * parser.c (handle_omp_declare_target_clause): New function. (cp_parser_omp_declare_target): Use it. gcc/testsuite/ * c-c++-common/gomp/pr102640.c: New test. (cherry picked from commit db3d7270b42fe27fb05664c4fdf524ab7ad13a75) Diff: --- gcc/c-family/ChangeLog.omp | 11 +++ gcc/cp/ChangeLog.omp | 9 ++ gcc/cp/parser.c | 132 ++++++++++++++++------------- gcc/testsuite/ChangeLog.omp | 8 ++ gcc/testsuite/c-c++-common/gomp/pr102640.c | 44 ++++++++++ 5 files changed, 145 insertions(+), 59 deletions(-) diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp index afb871f6d82..4d19a6e56c7 100644 --- a/gcc/c-family/ChangeLog.omp +++ b/gcc/c-family/ChangeLog.omp @@ -1,3 +1,14 @@ +2021-10-06 Tobias Burnus <tobias@codesourcery.com> + + Backported from master: + 2021-10-06 Jakub Jelinek <jakub@redhat.com> + + PR tree-optimization/102571 + * c-omp.c (c_finish_omp_atomic): Optimize the case where type has + padding, but the non-padding bits are contiguous set of bytes + by adjusting the memcmp call arguments instead of emitting + __builtin_clear_padding and then comparing all the type's bytes. + 2021-10-01 Tobias Burnus <tobias@codesourcery.com> Backported from master: diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp index 06b5dc93302..b5b92bc6a1c 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,12 @@ +2021-10-08 Tobias Burnus <tobias@codesourcery.com> + + Backported from master: + 2021-10-08 Jakub Jelinek <jakub@redhat.com> + + PR c++/102640 + * parser.c (handle_omp_declare_target_clause): New function. + (cp_parser_omp_declare_target): Use it. + 2021-10-01 Tobias Burnus <tobias@codesourcery.com> Backported from master: diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 954f85d4b2d..6694349adb7 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -45241,6 +45241,71 @@ cp_parser_late_parsing_omp_declare_simd (cp_parser *parser, tree attrs) return attrs; } +/* Helper for cp_parser_omp_declare_target, handle one to or link clause + on #pragma omp declare target. Return false if errors were reported. */ + +static bool +handle_omp_declare_target_clause (tree c, tree t, int device_type) +{ + tree at1 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (t)); + tree at2 = lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (t)); + tree id; + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINK) + { + id = get_identifier ("omp declare target link"); + std::swap (at1, at2); + } + else + id = get_identifier ("omp declare target"); + if (at2) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD specified both in declare target %<link%> and %<to%>" + " clauses", t); + return false; + } + if (!at1) + { + DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t)); + if (TREE_CODE (t) != FUNCTION_DECL && !is_global_var (t)) + return true; + + symtab_node *node = symtab_node::get (t); + if (node != NULL) + { + node->offloadable = 1; + if (ENABLE_OFFLOADING) + { + g->have_offload = true; + if (is_a <varpool_node *> (node)) + vec_safe_push (offload_vars, t); + } + } + } + if (TREE_CODE (t) != FUNCTION_DECL) + return true; + if ((device_type & OMP_CLAUSE_DEVICE_TYPE_HOST) != 0) + { + tree at3 = lookup_attribute ("omp declare target host", + DECL_ATTRIBUTES (t)); + if (at3 == NULL_TREE) + { + id = get_identifier ("omp declare target host"); + DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t)); + } + } + if ((device_type & OMP_CLAUSE_DEVICE_TYPE_NOHOST) != 0) + { + tree at3 = lookup_attribute ("omp declare target nohost", + DECL_ATTRIBUTES (t)); + if (at3 == NULL_TREE) + { + id = get_identifier ("omp declare target nohost"); + DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t)); + } + } + return true; +} /* OpenMP 4.0: # pragma omp declare target new-line @@ -45293,67 +45358,16 @@ cp_parser_omp_declare_target (cp_parser *parser, cp_token *pragma_tok) { if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE) continue; - tree t = OMP_CLAUSE_DECL (c), id; - tree at1 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (t)); - tree at2 = lookup_attribute ("omp declare target link", - DECL_ATTRIBUTES (t)); + tree t = OMP_CLAUSE_DECL (c); only_device_type = false; - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINK) - { - id = get_identifier ("omp declare target link"); - std::swap (at1, at2); - } - else - id = get_identifier ("omp declare target"); - if (at2) - { - error_at (OMP_CLAUSE_LOCATION (c), - "%qD specified both in declare target %<link%> and %<to%>" - " clauses", t); - continue; - } - if (!at1) - { - DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t)); - if (TREE_CODE (t) != FUNCTION_DECL && !is_global_var (t)) - continue; - - symtab_node *node = symtab_node::get (t); - if (node != NULL) - { - node->offloadable = 1; - if (ENABLE_OFFLOADING) - { - g->have_offload = true; - if (is_a <varpool_node *> (node)) - vec_safe_push (offload_vars, t); - } - } - } - if (TREE_CODE (t) != FUNCTION_DECL) + if (!handle_omp_declare_target_clause (c, t, device_type)) continue; - if ((device_type & OMP_CLAUSE_DEVICE_TYPE_HOST) != 0) - { - tree at3 = lookup_attribute ("omp declare target host", - DECL_ATTRIBUTES (t)); - if (at3 == NULL_TREE) - { - id = get_identifier ("omp declare target host"); - DECL_ATTRIBUTES (t) - = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t)); - } - } - if ((device_type & OMP_CLAUSE_DEVICE_TYPE_NOHOST) != 0) - { - tree at3 = lookup_attribute ("omp declare target nohost", - DECL_ATTRIBUTES (t)); - if (at3 == NULL_TREE) - { - id = get_identifier ("omp declare target nohost"); - DECL_ATTRIBUTES (t) - = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t)); - } - } + if (VAR_OR_FUNCTION_DECL_P (t) + && DECL_LOCAL_DECL_P (t) + && DECL_LANG_SPECIFIC (t) + && DECL_LOCAL_DECL_ALIAS (t)) + handle_omp_declare_target_clause (c, DECL_LOCAL_DECL_ALIAS (t), + device_type); } if (device_type && only_device_type) warning_at (OMP_CLAUSE_LOCATION (clauses), 0, diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 1027065ddf7..e806233da0a 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,11 @@ +2021-10-08 Tobias Burnus <tobias@codesourcery.com> + + Backported from master: + 2021-10-08 Jakub Jelinek <jakub@redhat.com> + + PR c++/102640 + * c-c++-common/gomp/pr102640.c: New test. + 2021-10-06 Tobias Burnus <tobias@codesourcery.com> Backported from master: diff --git a/gcc/testsuite/c-c++-common/gomp/pr102640.c b/gcc/testsuite/c-c++-common/gomp/pr102640.c new file mode 100644 index 00000000000..00ebab9c628 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/pr102640.c @@ -0,0 +1,44 @@ +/* PR c++/102640 */ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-omplower" } */ +/* Verify var[123] are mapped without any copying, because they are + mentioned in declare target directive to clauses. */ +/* { dg-final { scan-tree-dump-not "firstprivate\\\(var\[123]\\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump-not ".omp_data_arr.\[0-9]*.var" "omplower" } } */ +/* { dg-final { scan-tree-dump-not ".omp_data_i->var" "omplower" } } */ + +void +foo (void) +{ + extern int var1; + #pragma omp declare target to (var1) + + #pragma omp target + var1++; +} + +int +bar (int x) +{ + extern int var2; + #pragma omp declare target to (var2) + if (x) + return var2; + #pragma omp target + var2++; + return -1; +} +#pragma omp declare target to (bar) + +#pragma omp declare target +int +baz (int x) +{ + extern int var3; + if (x) + return var3; + #pragma omp target + var3++; + return -1; +} +#pragma omp end declare target
reply other threads:[~2021-10-08 11:35 UTC|newest] Thread overview: [no followups] expand[flat|nested] mbox.gz Atom feed
Reply instructions: You may reply publicly to this message via plain-text email using any one of the following methods: * Save the following mbox file, import it into your mail client, and reply-to-all from there: mbox Avoid top-posting and favor interleaved quoting: https://en.wikipedia.org/wiki/Posting_style#Interleaved_style * Reply using the --to, --cc, and --in-reply-to switches of git-send-email(1): git send-email \ --in-reply-to=20211008113525.4AEBD385803F@sourceware.org \ --to=burnus@gcc.gnu.org \ --cc=gcc-cvs@gcc.gnu.org \ /path/to/YOUR_REPLY https://kernel.org/pub/software/scm/git/docs/git-send-email.html * If your mail client supports setting the In-Reply-To header via mailto: links, try the mailto: linkBe sure your reply has a Subject: header at the top and a blank line before the message body.
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).