* [Patch] OpenMP: Fix 'omp declare target' handling for vars [PR99509] @ 2021-03-10 14:20 Tobias Burnus 2021-03-12 12:20 ` Jakub Jelinek 0 siblings, 1 reply; 4+ messages in thread From: Tobias Burnus @ 2021-03-10 14:20 UTC (permalink / raw) To: gcc-patches, fortran, Jakub Jelinek [-- Attachment #1: Type: text/plain, Size: 668 bytes --] The C/C++ FE sets for an 'omp declare target' ... 'omp end declare target' the attribute 'omp declare target implicit'. That's later processed (for C++) in decl.c - which remove that attribute and either keeps and explicit 'omp declare target' or 'omp declare target link' attribute. Unfortunately, adding 'omp declare target' comes too late as the varpool has been generated. I am not sure what's the best way, but the following seems to work. OK for mainline? Tobias ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf [-- Attachment #2: omp-decl-tgt.diff --] [-- Type: text/x-patch, Size: 2107 bytes --] OpenMP: Fix 'omp declare target' handling for vars [PR99509] C/C++ use 'omp declare target implicit' when they call varpool_node::get_create; make sure that a later 'omp declare target' (but not 'omp declare target link') is honored. gcc/ChangeLog: PR c++/99509 * varpool.c (varpool_node::get_create): Re-check whether 'omp declare target'. libgomp/ChangeLog: PR c++/99509 * testsuite/libgomp.c-c++-common/declare_target-1.c: New test. gcc/varpool.c | 18 +++++++++++++++++- .../libgomp.c-c++-common/declare_target-1.c | 22 ++++++++++++++++++++++ 2 files changed, 39 insertions(+), 1 deletion(-) diff --git a/gcc/varpool.c b/gcc/varpool.c index 4830df5..5ae6778 100644 --- a/gcc/varpool.c +++ b/gcc/varpool.c @@ -145,7 +145,23 @@ varpool_node::get_create (tree decl) varpool_node *node = varpool_node::get (decl); gcc_checking_assert (VAR_P (decl)); if (node) - return node; + { + /* This happens with C/C++ when 'omp declare target implicit' is set + which is only later replaced by 'omp declate target' or '... link'. */ + if (!node->offloadable + && (flag_openacc || flag_openmp) + && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))) + { + node->offloadable = 1; + if (ENABLE_OFFLOADING && !DECL_EXTERNAL (decl)) + { + g->have_offload = true; + if (!in_lto_p) + vec_safe_push (offload_vars, decl); + } + } + return node; + } node = varpool_node::create_empty (); node->decl = decl; diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c b/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c new file mode 100644 index 0000000..c5670df --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c @@ -0,0 +1,22 @@ +/* PR c++/99509 */ + +#pragma omp declare target +int data[] = {5}; +#pragma omp end declare target + +static inline int +foo (int idx) +{ + return data[idx]; +} + +int +main () +{ + int i = -1; + #pragma omp target map(from:i) + i = foo(0); + if (i != 5) + __builtin_abort (); + return 0; +} ^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [Patch] OpenMP: Fix 'omp declare target' handling for vars [PR99509] 2021-03-10 14:20 [Patch] OpenMP: Fix 'omp declare target' handling for vars [PR99509] Tobias Burnus @ 2021-03-12 12:20 ` Jakub Jelinek 2021-03-15 8:17 ` Tobias Burnus 0 siblings, 1 reply; 4+ messages in thread From: Jakub Jelinek @ 2021-03-12 12:20 UTC (permalink / raw) To: Tobias Burnus; +Cc: gcc-patches, fortran On Wed, Mar 10, 2021 at 03:20:42PM +0100, Tobias Burnus wrote: > The C/C++ FE sets for an 'omp declare target' ... 'omp end declare target' > the attribute 'omp declare target implicit'. > > That's later processed (for C++) in decl.c - which remove that attribute > and either keeps and explicit 'omp declare target' or 'omp declare target link' > attribute. > > Unfortunately, adding 'omp declare target' comes too late as the varpool > has been generated. Looking attributes on every get rather than just once is IMHO too expensive. For explicit declare target, we use: 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); } } (e.g. from cp/parser.c). So, I think it would be better to do the same thing when we turn an "omp declare target implicit" into "omp declare target", i.e. in cp/decl.c (cp_finish_decl) and in c/c-decl.c (finish_decl) - right after adding the "omp declare target" attribute in there. Jakub ^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [Patch] OpenMP: Fix 'omp declare target' handling for vars [PR99509] 2021-03-12 12:20 ` Jakub Jelinek @ 2021-03-15 8:17 ` Tobias Burnus 2021-03-15 8:20 ` Jakub Jelinek 0 siblings, 1 reply; 4+ messages in thread From: Tobias Burnus @ 2021-03-15 8:17 UTC (permalink / raw) To: Jakub Jelinek; +Cc: gcc-patches, fortran [-- Attachment #1: Type: text/plain, Size: 1187 bytes --] On 12.03.21 13:20, Jakub Jelinek wrote: > On Wed, Mar 10, 2021 at 03:20:42PM +0100, Tobias Burnus wrote: >> The C/C++ FE sets for an 'omp declare target' ... 'omp end declare target' >> the attribute 'omp declare target implicit'. >> >> That's later processed (for C++) in decl.c - which remove that attribute >> and either keeps and explicit 'omp declare target' or 'omp declare target link' >> attribute. >> >> Unfortunately, adding 'omp declare target' comes too late as the varpool >> has been generated. >> >> Looking attributes on every get rather than just once is IMHO too expensive. >> For explicit declare target, we use: >> ... >> (e.g. from cp/parser.c). >> So, I think it would be better to do the same thing I concur. I was already wondering whether it should be in the FE or not, thinking about some separation between FE and ME. But given the precedent and the performance issue, it surely makes sense to solve it in c*/*decl.c. Done now. OK for mainline? Tobias ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf [-- Attachment #2: omp-decl-tgt-v2.diff --] [-- Type: text/x-patch, Size: 4284 bytes --] OpenMP: Fix 'omp declare target' handling for vars [PR99509] For variables with 'declare target' attribute, varpool_node::get_create marks variables as offload; however, if the node already exists, it is not updated. C/C++ may tag decl with 'declare target implicit', which may only be after varpool creation turned into 'declare target' or 'declare target link'; in this case, the tagging has to happen in the FE. gcc/c/ChangeLog: PR c++/99509 * c-decl.c (finish_decl): For 'omp declare target implicit' vars, ensure that the varpool node is marked as offloadable. gcc/cp/ChangeLog: PR c++/99509 * decl.c (cp_finish_decl): For 'omp declare target implicit' vars, ensure that the varpool node is marked as offloadable. libgomp/ChangeLog: PR c++/99509 * testsuite/libgomp.c-c++-common/declare_target-1.c: New test. gcc/c/c-decl.c | 22 +++++++++++++++++++--- gcc/cp/decl.c | 21 ++++++++++++++++++--- .../libgomp.c-c++-common/declare_target-1.c | 22 ++++++++++++++++++++++ 3 files changed, 59 insertions(+), 6 deletions(-) diff --git a/gcc/c/c-decl.c b/gcc/c/c-decl.c index b559ed5d76a..3b2241bfd97 100644 --- a/gcc/c/c-decl.c +++ b/gcc/c/c-decl.c @@ -58,6 +58,9 @@ along with GCC; see the file COPYING3. If not see #include "c-family/name-hint.h" #include "c-family/known-headers.h" #include "c-family/c-spellcheck.h" +#include "context.h" /* For 'g'. */ +#include "omp-general.h" +#include "omp-offload.h" /* For offload_vars. */ #include "tree-pretty-print.h" @@ -5658,9 +5661,22 @@ finish_decl (tree decl, location_t init_loc, tree init, DECL_ATTRIBUTES (decl)) && !lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (decl))) - DECL_ATTRIBUTES (decl) - = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, DECL_ATTRIBUTES (decl)); + { + DECL_ATTRIBUTES (decl) + = tree_cons (get_identifier ("omp declare target"), + NULL_TREE, DECL_ATTRIBUTES (decl)); + symtab_node *node = symtab_node::get (decl); + if (node != NULL) + { + node->offloadable = 1; + if (ENABLE_OFFLOADING) + { + g->have_offload = true; + if (is_a <varpool_node *> (node)) + vec_safe_push (offload_vars, decl); + } + } + } } /* This is the last point we can lower alignment so give the target the diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c index 9c7f6e59bbb..1b671cec9a3 100644 --- a/gcc/cp/decl.c +++ b/gcc/cp/decl.c @@ -53,7 +53,9 @@ along with GCC; see the file COPYING3. If not see #include "asan.h" #include "gcc-rich-location.h" #include "langhooks.h" +#include "context.h" /* For 'g'. */ #include "omp-general.h" +#include "omp-offload.h" /* For offload_vars. */ /* Possible cases of bad specifiers type used by bad_specifiers. */ enum bad_spec_place { @@ -8176,9 +8178,22 @@ cp_finish_decl (tree decl, tree init, bool init_const_expr_p, DECL_ATTRIBUTES (decl)) && !lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (decl))) - DECL_ATTRIBUTES (decl) - = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, DECL_ATTRIBUTES (decl)); + { + DECL_ATTRIBUTES (decl) + = tree_cons (get_identifier ("omp declare target"), + NULL_TREE, DECL_ATTRIBUTES (decl)); + symtab_node *node = symtab_node::get (decl); + if (node != NULL) + { + node->offloadable = 1; + if (ENABLE_OFFLOADING) + { + g->have_offload = true; + if (is_a <varpool_node *> (node)) + vec_safe_push (offload_vars, decl); + } + } + } } /* This is the last point we can lower alignment so give the target the diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c b/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c new file mode 100644 index 00000000000..c5670dfb7db --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c @@ -0,0 +1,22 @@ +/* PR c++/99509 */ + +#pragma omp declare target +int data[] = {5}; +#pragma omp end declare target + +static inline int +foo (int idx) +{ + return data[idx]; +} + +int +main () +{ + int i = -1; + #pragma omp target map(from:i) + i = foo(0); + if (i != 5) + __builtin_abort (); + return 0; +} ^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [Patch] OpenMP: Fix 'omp declare target' handling for vars [PR99509] 2021-03-15 8:17 ` Tobias Burnus @ 2021-03-15 8:20 ` Jakub Jelinek 0 siblings, 0 replies; 4+ messages in thread From: Jakub Jelinek @ 2021-03-15 8:20 UTC (permalink / raw) To: Tobias Burnus; +Cc: gcc-patches, fortran On Mon, Mar 15, 2021 at 09:17:47AM +0100, Tobias Burnus wrote: > gcc/c/ChangeLog: > > PR c++/99509 > * c-decl.c (finish_decl): For 'omp declare target implicit' vars, > ensure that the varpool node is marked as offloadable. > > gcc/cp/ChangeLog: > > PR c++/99509 > * decl.c (cp_finish_decl): For 'omp declare target implicit' vars, > ensure that the varpool node is marked as offloadable. > > libgomp/ChangeLog: > > PR c++/99509 > * testsuite/libgomp.c-c++-common/declare_target-1.c: New test. Ok, thanks. Jakub ^ permalink raw reply [flat|nested] 4+ messages in thread
end of thread, other threads:[~2021-03-15 8:21 UTC | newest] Thread overview: 4+ messages (download: mbox.gz / follow: Atom feed) -- links below jump to the message on this page -- 2021-03-10 14:20 [Patch] OpenMP: Fix 'omp declare target' handling for vars [PR99509] Tobias Burnus 2021-03-12 12:20 ` Jakub Jelinek 2021-03-15 8:17 ` Tobias Burnus 2021-03-15 8:20 ` Jakub Jelinek
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).