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  <jakub@redhat.com>

	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.

diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index 23a13d82808b50fdd6f2033acbd9c1e2f9f75291..4baec32e84353a3e8f8f5e7d6858f15c25c4bee9 100644 (file)
--- 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 <gomp_target *> (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 09a8cbdc433ab9ec6be9dc2d728470cf941c9ba0..91a5e3d143160c21ee2f4aa2b5f5a6a4683c31aa 100644 (file)
--- 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 (file)
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 (file)
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 ();
+  }
+}