From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 2153) id AABCE388A837; Tue, 20 Apr 2021 23:30:47 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org AABCE388A837 MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Jakub Jelinek To: gcc-cvs@gcc.gnu.org Subject: [gcc r9-9401] openmp, openacc: Fix up handling of data regions [PR98183] X-Act-Checkin: gcc X-Git-Author: Jakub Jelinek X-Git-Refname: refs/heads/releases/gcc-9 X-Git-Oldrev: 6f0a0d1c2bb7ad2020852ccf14ca86967ddb134a X-Git-Newrev: eca61f3b10c177258f09c28613062d2adb588984 Message-Id: <20210420233047.AABCE388A837@sourceware.org> Date: Tue, 20 Apr 2021 23:30:47 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Tue, 20 Apr 2021 23:30:47 -0000 https://gcc.gnu.org/g:eca61f3b10c177258f09c28613062d2adb588984 commit r9-9401-geca61f3b10c177258f09c28613062d2adb588984 Author: Jakub Jelinek Date: Sat Dec 12 08:36:02 2020 +0100 openmp, openacc: Fix up handling of data regions [PR98183] While the data regions (target data and OpenACC counterparts) aren't standalone directives, unlike most other OpenMP/OpenACC constructs we allow (apparently as an extension) exceptions and goto out of the block. During gimplification we place an *end* call into a finally block so that it is reached even on exceptions or goto out etc.). During omplower pass we then add paired #pragma omp return for them, but due to the exceptions because the region is not SESE we can end up with #pragma omp return appearing only conditionally in the CFG etc., which the ompexp pass can't handle. For the ompexp pass, we actually don't care about the end part or about target data nesting, so we can treat it as standalone directive. 2020-12-12 Jakub Jelinek PR middle-end/98183 * omp-low.c (lower_omp_target): Don't add OMP_RETURN for data regions. * omp-expand.c (expand_omp_target): Don't try to remove OMP_RETURN for data regions. (build_omp_regions_1, omp_make_gimple_edges): Don't expect OMP_RETURN for data regions. * gcc.dg/gomp/pr98183.c: New test. * gcc.dg/goacc/pr98183.c: New test. (cherry picked from commit 8c1ed7223ad1bc19ed9c936ba496220c8ef673bc) Diff: --- gcc/omp-expand.c | 24 +++++++----------------- gcc/omp-low.c | 7 ++++--- gcc/testsuite/gcc.dg/goacc/pr98183.c | 15 +++++++++++++++ gcc/testsuite/gcc.dg/gomp/pr98183.c | 15 +++++++++++++++ 4 files changed, 41 insertions(+), 20 deletions(-) diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index db3ede084aa..99ad994266c 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -7281,7 +7281,7 @@ expand_omp_target (struct omp_region *region) gomp_target *entry_stmt; gimple *stmt; edge e; - bool offloaded, data_region; + bool offloaded; entry_stmt = as_a (last_stmt (region->entry)); new_bb = region->entry; @@ -7298,12 +7298,9 @@ expand_omp_target (struct omp_region *region) case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: - data_region = false; - break; case GF_OMP_TARGET_KIND_DATA: case GF_OMP_TARGET_KIND_OACC_DATA: case GF_OMP_TARGET_KIND_OACC_HOST_DATA: - data_region = true; break; default: gcc_unreachable (); @@ -7822,13 +7819,6 @@ expand_omp_target (struct omp_region *region) gcc_assert (g && gimple_code (g) == GIMPLE_OMP_TARGET); gsi_remove (&gsi, true); } - if (data_region && region->exit) - { - gsi = gsi_last_nondebug_bb (region->exit); - g = gsi_stmt (gsi); - gcc_assert (g && gimple_code (g) == GIMPLE_OMP_RETURN); - gsi_remove (&gsi, true); - } } /* Expand KFOR loop as a HSA grifidied kernel, i.e. as a body only with @@ -8291,15 +8281,15 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent, switch (gimple_omp_target_kind (stmt)) { case GF_OMP_TARGET_KIND_REGION: - case GF_OMP_TARGET_KIND_DATA: case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: - case GF_OMP_TARGET_KIND_OACC_DATA: - case GF_OMP_TARGET_KIND_OACC_HOST_DATA: break; case GF_OMP_TARGET_KIND_UPDATE: case GF_OMP_TARGET_KIND_ENTER_DATA: case GF_OMP_TARGET_KIND_EXIT_DATA: + case GF_OMP_TARGET_KIND_DATA: + case GF_OMP_TARGET_KIND_OACC_DATA: + case GF_OMP_TARGET_KIND_OACC_HOST_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: @@ -8545,15 +8535,15 @@ omp_make_gimple_edges (basic_block bb, struct omp_region **region, switch (gimple_omp_target_kind (last)) { case GF_OMP_TARGET_KIND_REGION: - case GF_OMP_TARGET_KIND_DATA: case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: - case GF_OMP_TARGET_KIND_OACC_DATA: - case GF_OMP_TARGET_KIND_OACC_HOST_DATA: break; case GF_OMP_TARGET_KIND_UPDATE: case GF_OMP_TARGET_KIND_ENTER_DATA: case GF_OMP_TARGET_KIND_EXIT_DATA: + case GF_OMP_TARGET_KIND_DATA: + case GF_OMP_TARGET_KIND_OACC_DATA: + case GF_OMP_TARGET_KIND_OACC_HOST_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: diff --git a/gcc/omp-low.c b/gcc/omp-low.c index d3cb3ba5333..615f7f6327e 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -10176,9 +10176,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_seq_add_seq (&new_body, join_seq); if (offloaded) - new_body = maybe_catch_exception (new_body); - - gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false)); + { + new_body = maybe_catch_exception (new_body); + gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false)); + } gimple_omp_set_body (stmt, new_body); } diff --git a/gcc/testsuite/gcc.dg/goacc/pr98183.c b/gcc/testsuite/gcc.dg/goacc/pr98183.c new file mode 100644 index 00000000000..b0410117c84 --- /dev/null +++ b/gcc/testsuite/gcc.dg/goacc/pr98183.c @@ -0,0 +1,15 @@ +/* PR middle-end/98183 */ +/* { dg-additional-options "-fexceptions -O0" } */ + +void bar (void); +int x, y; + +void +foo (void) +{ +#pragma acc data copyout(x) + { +#pragma acc data copyout(y) + bar (); + } +} diff --git a/gcc/testsuite/gcc.dg/gomp/pr98183.c b/gcc/testsuite/gcc.dg/gomp/pr98183.c new file mode 100644 index 00000000000..dd114997920 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/pr98183.c @@ -0,0 +1,15 @@ +/* PR middle-end/98183 */ +/* { dg-additional-options "-fexceptions -O0" } */ + +void bar (void); +int x, y; + +void +foo (void) +{ +#pragma omp target data map(tofrom: x) + { +#pragma omp target data map(tofrom: y) + bar (); + } +}