public inbox for gcc-bugs@sourceware.org help / color / mirror / Atom feed
From: "tschwinge at gcc dot gnu.org" <gcc-bugzilla@gcc.gnu.org> To: gcc-bugs@gcc.gnu.org Subject: [Bug tree-optimization/90591] Avoid unnecessary data transfer out of OMP construct Date: Mon, 20 Apr 2020 19:57:13 +0000 [thread overview] Message-ID: <bug-90591-4-4D71oO30WM@http.gcc.gnu.org/bugzilla/> (raw) In-Reply-To: <bug-90591-4@http.gcc.gnu.org/bugzilla/> https://gcc.gnu.org/bugzilla/show_bug.cgi?id=90591 Thomas Schwinge <tschwinge at gcc dot gnu.org> changed: What |Removed |Added ---------------------------------------------------------------------------- CC| |rguenth at gcc dot gnu.org Ever confirmed|0 |1 Last reconfirmed| |2020-04-20 Status|UNCONFIRMED |ASSIGNED Assignee|unassigned at gcc dot gnu.org |sandra at gcc dot gnu.org --- Comment #2 from Thomas Schwinge <tschwinge at gcc dot gnu.org> --- I'm not very familiar with IPA optimizations, I'm not yet clear on what can be done (optimized) with reasonable effort in GCC, and where/how, so will very much appreciate your input, Jakub, Richard, others. Sandra is going to look into this topic (who is not too familiar with OMP handling). As a first step, we shall not consider any offloading compilation specifics, but just the usual host-side compilation. Looking at the following simple example: int main(void) { int var; #pragma acc parallel copyout(var) #pragma omp target map(from:var) { var = 1; } // 'var' never read. return 0; } For reference, with '-fno-openacc -fno-openmp -O1' (so, OMP deactiveated, pragmas ignored), the 'var = 1' assignment disappears in the '030t.ccp1' dump file, and 'var' itself disappears in the '047t.release_ssa' dump file. With OMP enabled ('-fopenacc' shown in the following, but '-fopenmp' is very similar), it's more difficult to optimize 'var': the OMP region (here: just the 'var = 1' assignment, plus any set-up and tear-down code) is moved (outlined) into a separate function 'main._omp_fn.0', and the address of 'var' is taken, stored in an internal data structure '.omp_data_arr.1', and is dereferenced in 'main._omp_fn.0' to access the original 'var'. The outlined function 'main._omp_fn.0' is called via 'GOACC_parallel_keyed'. One step is to add logic so that in this example, we can optimize 'copyout' to 'create'. The 'copyout' clauses is encoded in '.omp_data_kinds.3[0]' as value '514', where with the upper data alignment byte stripped off, '514 & 255 = 2', which is 'GOMP_MAP_FROM'. This shall -- at a suitable point in the pass pipeline -- be optimized to 'GOMP_MAP_ALLOC'. Another step is to add logic so that the "dead"ness of 'var' after the outlined function 'main._omp_fn.0' called via 'GOACC_parallel_keyed' gets propagated into 'main._omp_fn.0', so that the 'var = 1' assignment can be eliminated. By the way, there already is some special IPA information handling for 'GOACC_parallel_keyed' ('BUILT_IN_GOACC_PARALLEL') in 'gcc/tree-ssa-structalias.c'. This is, if I remember correctly, to evaluate aliasing "transparently through" the outlined OMP function; see the PR46032 commit r231076 "Handle BUILT_IN_GOMP_PARALLEL in ipa-pta" and later commits, for reference. (And, there is a 'pass_oacc_ipa' to enable '-fipa-pta' for OpenACC OMP functions.) Working on eliminating 'var' completely from the internal '.omp_data_arr.1' etc. data structures shall be a separate step, for later. If that elimination happens when we're already in offloading compilation pipeline, this is expected to require more infrastructure to communicate that information back from the offloading compiler(s) to the host compiler/runtime. But first one step back -- I noticed that compiling the following: #pragma acc routine #pragma omp declare target static void __attribute__((noinline)) f(int *var) { *var = 1; } #pragma omp end declare target int main(void) { int var; #pragma acc parallel copyout(var) #pragma omp target map(from:var) { f(&var); } // 'var' never read. return 0; } ... with '-fno-openacc -fno-openmp -O1' (so, OMP deactiveated, pragmas ignored) also does *not* see any optimization of 'var'. Is that something that needs to be addressed first, before attempting the OMP case?
next parent reply other threads:[~2020-04-20 19:57 UTC|newest] Thread overview: 5+ messages / expand[flat|nested] mbox.gz Atom feed top [not found] <bug-90591-4@http.gcc.gnu.org/bugzilla/> 2020-04-20 19:57 ` tschwinge at gcc dot gnu.org [this message] 2020-04-21 6:42 ` rguenth at gcc dot gnu.org 2020-04-28 16:53 ` burnus at gcc dot gnu.org 2020-05-05 14:16 ` burnus at gcc dot gnu.org 2021-03-22 11:37 ` tschwinge at gcc dot gnu.org
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=bug-90591-4-4D71oO30WM@http.gcc.gnu.org/bugzilla/ \ --to=gcc-bugzilla@gcc.gnu.org \ --cc=gcc-bugs@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).