public inbox for gcc-cvs@sourceware.org help / color / mirror / Atom feed
From: Jakub Jelinek <jakub@gcc.gnu.org> To: gcc-cvs@gcc.gnu.org Subject: [gcc r12-7459] openmp: Disable SSA form during gimplification on OMP_SIMD clauses and body [PR104757] Date: Thu, 3 Mar 2022 08:22:28 +0000 (GMT) [thread overview] Message-ID: <20220303082228.01B9E3858D39@sourceware.org> (raw) https://gcc.gnu.org/g:431414b5d934866af3f6415a56c35bb57b928fef commit r12-7459-g431414b5d934866af3f6415a56c35bb57b928fef Author: Jakub Jelinek <jakub@redhat.com> Date: Thu Mar 3 09:13:32 2022 +0100 openmp: Disable SSA form during gimplification on OMP_SIMD clauses and body [PR104757] When offloading to nvptx is enabled, scan_omp_simd duplicates the simd region including its clauses and body using inliner's copy_gimple_seq_and_replace_locals. That works nicely for decls, remaps only those that are seen in the nested bind expr vars (i.e. local variables) and doesn't remap other vars. But for SSA_NAMEs it remaps them always, doesn't know if their def stmt is outside of the simd (then it better shouldn't be remapped) or inside of it (then it should) and without cfg/dominators that is pretty hard to figure out (well, we could walk the region twice, once note SSA_NAMEs defined by each stmt seen there and once do the remapping of only those visited SSA_NAMEs). This patch uses a simpler way, disables temporarily into_ssa for the clauses and body of each simd region; we already disable into_ssa e.g. in parallel/target/task etc. regions through push_gimplify_context () but for simd we don't push any gimplification context and appart from into_ssa I think we don't need it. 2022-03-03 Jakub Jelinek <jakub@redhat.com> PR middle-end/104757 * gimplify.cc (gimplify_omp_loop): Call gimplify_expr rather than gimplify_omp_for. (gimplify_expr) <case OMP_SIMD>: Temporarily disable gimplify_ctxp->into_ssa around call to gimplify_omp_for. * gfortran.dg/gomp/pr104757.f90: New test. * gcc.dg/gomp/pr104757.c: New test. Diff: --- gcc/gimplify.cc | 15 +++++++++++++-- gcc/testsuite/gcc.dg/gomp/pr104757.c | 14 ++++++++++++++ gcc/testsuite/gfortran.dg/gomp/pr104757.f90 | 19 +++++++++++++++++++ 3 files changed, 46 insertions(+), 2 deletions(-) diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 2364d2e5182..be8f0d3c76f 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -13786,7 +13786,7 @@ gimplify_omp_loop (tree *expr_p, gimple_seq *pre_p) *pc = NULL_TREE; *expr_p = t; } - return gimplify_omp_for (expr_p, pre_p); + return gimplify_expr (expr_p, pre_p, NULL, is_gimple_stmt, fb_none); } @@ -15540,8 +15540,19 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, ret = GS_ALL_DONE; break; - case OMP_FOR: case OMP_SIMD: + { + /* Temporarily disable into_ssa, as scan_omp_simd + which calls copy_gimple_seq_and_replace_locals can't deal + with SSA_NAMEs defined outside of the body properly. */ + bool saved_into_ssa = gimplify_ctxp->into_ssa; + gimplify_ctxp->into_ssa = false; + ret = gimplify_omp_for (expr_p, pre_p); + gimplify_ctxp->into_ssa = saved_into_ssa; + break; + } + + case OMP_FOR: case OMP_DISTRIBUTE: case OMP_TASKLOOP: case OACC_LOOP: diff --git a/gcc/testsuite/gcc.dg/gomp/pr104757.c b/gcc/testsuite/gcc.dg/gomp/pr104757.c new file mode 100644 index 00000000000..738396523af --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/pr104757.c @@ -0,0 +1,14 @@ +/* PR middle-end/104757 */ +/* { dg-do compile } */ +/* { dg-options "-O2 -fopenmp" } */ + +#pragma omp declare target +void +foo (int x, int y, int *z) +{ + int j = 0; + #pragma omp simd linear(j:x + y) + for (int i = 0; i < 64; i++) + j += x + y; +} +#pragma omp end declare target diff --git a/gcc/testsuite/gfortran.dg/gomp/pr104757.f90 b/gcc/testsuite/gfortran.dg/gomp/pr104757.f90 new file mode 100644 index 00000000000..6f1aef54ed0 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr104757.f90 @@ -0,0 +1,19 @@ +! PR middle-end/104757 +! { dg-do compile } +! { dg-options "-O -fopenmp" } + +module pr104757 + implicit none (external, type) + integer :: ll + !$omp declare target (ll) +contains + subroutine foo (i1) + !$omp declare target (foo) + logical :: i1 + integer :: i + !$omp distribute simd if(i1) + do i = 1, 64 + ll = ll + 1 + end do + end +end module
reply other threads:[~2022-03-03 8:22 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=20220303082228.01B9E3858D39@sourceware.org \ --to=jakub@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).