diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index 23a13d82808b..4baec32e8435 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -9260,7 +9260,7 @@ expand_omp_target (struct omp_region *region) gomp_target *entry_stmt; gimple *stmt; edge e; - bool offloaded, data_region; + bool offloaded; int target_kind; entry_stmt = as_a (last_stmt (region->entry)); @@ -9282,13 +9282,10 @@ expand_omp_target (struct omp_region *region) case GF_OMP_TARGET_KIND_OACC_DECLARE: case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED: case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE: - 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: case GF_OMP_TARGET_KIND_OACC_DATA_KERNELS: - data_region = true; break; default: gcc_unreachable (); @@ -9855,13 +9852,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 the parallel region tree rooted at REGION. Expansion @@ -10026,19 +10016,19 @@ 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_SERIAL: - case GF_OMP_TARGET_KIND_OACC_DATA: - case GF_OMP_TARGET_KIND_OACC_HOST_DATA: case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED: case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE: - case GF_OMP_TARGET_KIND_OACC_DATA_KERNELS: 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_DATA_KERNELS: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: @@ -10283,19 +10273,19 @@ 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_SERIAL: - case GF_OMP_TARGET_KIND_OACC_DATA: - case GF_OMP_TARGET_KIND_OACC_HOST_DATA: case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED: case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE: - case GF_OMP_TARGET_KIND_OACC_DATA_KERNELS: 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_DATA_KERNELS: 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 09a8cbdc433a..91a5e3d14316 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -12979,9 +12979,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 000000000000..b0410117c84f --- /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 000000000000..dd1149979204 --- /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 (); + } +}