diff options
| -rw-r--r-- | gcc/omp-expand.c | 24 | ||||
| -rw-r--r-- | gcc/omp-low.c | 7 | ||||
| -rw-r--r-- | gcc/testsuite/gcc.dg/goacc/pr98183.c | 15 | ||||
| -rw-r--r-- | 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 68f899a..90df5c9 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -7920,7 +7920,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 <gomp_target *> (last_stmt (region->entry)); @@ -7940,12 +7940,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 (); @@ -8499,13 +8496,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 @@ -8968,16 +8958,16 @@ 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: 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: @@ -9223,16 +9213,16 @@ 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: 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 d951c54..625566a 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -12600,9 +12600,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 0000000..b041011 --- /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 0000000..dd11499 --- /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 (); + } +} |