Handle oacc region in oacc routine

2016-03-05  Tom de Vries  <tom@codesourcery.com>

	* omp-low.c (check_omp_nesting_restrictions): Check for non-oacc
	construct in oacc routine.  Check for oacc region in oacc routine.

	* c-c++-common/goacc/nesting-fail-1.c (f_acc_routine): New function.
	* c-c++-common/goacc-gomp/nesting-fail-1.c (f_acc_routine): New
	function.

From-SVN: r233998

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index ea4e2d52491d3cf94b3e91a6122004519a592e4b..5c238366354c90414fadf0d6ddbe6db5e98ea1dc 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,8 @@
+2016-03-05  Tom de Vries  <tom@codesourcery.com>
+
+       * omp-low.c (check_omp_nesting_restrictions): Check for non-oacc
+       construct in oacc routine.  Check for oacc region in oacc routine.
+
 2016-03-04  Jakub Jelinek  <jakub@redhat.com>
 
        PR target/70062

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ecbf74ad71bf7201c0e4a5fe4aabea8d1cf895d1..c69fe448da62f320fc91e95ce171d80f437575ef 100644 (file)
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -3236,19 +3236,26 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
   /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP builtin)
      inside an OpenACC CTX.  */
   if (!(is_gimple_omp (stmt)
-       && is_gimple_omp_oacc (stmt)))
-    {
-      for (omp_context *octx = ctx; octx != NULL; octx = octx->outer)
-       if (is_gimple_omp (octx->stmt)
-           && is_gimple_omp_oacc (octx->stmt)
-           /* Except for atomic codes that we share with OpenMP.  */
-           && ! (gimple_code (stmt) == GIMPLE_OMP_ATOMIC_LOAD
-                 || gimple_code (stmt) == GIMPLE_OMP_ATOMIC_STORE))
-         {
-           error_at (gimple_location (stmt),
-                     "non-OpenACC construct inside of OpenACC region");
-           return false;
-         }
+       && is_gimple_omp_oacc (stmt))
+      /* Except for atomic codes that we share with OpenMP.  */
+      && !(gimple_code (stmt) == GIMPLE_OMP_ATOMIC_LOAD
+          || gimple_code (stmt) == GIMPLE_OMP_ATOMIC_STORE))
+    {
+      if (get_oacc_fn_attrib (cfun->decl) != NULL)
+       {
+         error_at (gimple_location (stmt),
+                   "non-OpenACC construct inside of OpenACC routine");
+         return false;
+       }
+      else
+       for (omp_context *octx = ctx; octx != NULL; octx = octx->outer)
+         if (is_gimple_omp (octx->stmt)
+             && is_gimple_omp_oacc (octx->stmt))
+           {
+             error_at (gimple_location (stmt),
+                       "non-OpenACC construct inside of OpenACC region");
+             return false;
+           }
     }
 
   if (ctx != NULL)
@@ -3715,6 +3722,14 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
                      kind == OMP_CLAUSE_DEPEND_SOURCE ? "source" : "sink");
            return false;
          }
+      if (is_gimple_omp_offloaded (stmt)
+         && get_oacc_fn_attrib (cfun->decl) != NULL)
+       {
+         error_at (gimple_location (stmt),
+                   "OpenACC region inside of OpenACC routine, nested "
+                   "parallelism not supported yet");
+         return false;
+       }
       for (; ctx != NULL; ctx = ctx->outer)
        {
          if (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET)

diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index ca65abf63389416fb7b69ffed45beb7034aa1a9c..51469eaee8111628f2254b23e324a5787a041c1a 100644 (file)
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,9 @@
+2016-03-05  Tom de Vries  <tom@codesourcery.com>
+
+       * c-c++-common/goacc/nesting-fail-1.c (f_acc_routine): New function.
+       * c-c++-common/goacc-gomp/nesting-fail-1.c (f_acc_routine): New
+       function.
+
 2016-03-05  Patrick Palka  <ppalka@gcc.gnu.org>
 
        PR c++/66786

diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
index 1a4472107c676c3f503cecf888b1946ce3b1f9a0..5e3f183998a3a42fee1867ffa811dd2548e1ae30 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
@@ -439,3 +439,11 @@ f_acc_loop (void)
 #pragma omp target update to(i) /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
     }
 }
+
+#pragma acc routine
+void
+f_acc_routine (void)
+{
+#pragma omp target /* { dg-error "non-OpenACC construct inside of OpenACC routine" } */
+  ;
+}

diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
index 7a36074ae384fd78e4a69e07e21e58911da54a63..506a1aeaa6a9391b54740fc4dfa4ad04292a8cb3 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
@@ -37,3 +37,11 @@ f_acc_kernels (void)
 #pragma acc exit data delete(i) /* { dg-error ".enter/exit data. construct inside of .kernels. region" } */
   }
 }
+
+#pragma acc routine
+void
+f_acc_routine (void)
+{
+#pragma acc parallel /* { dg-error "OpenACC region inside of OpenACC routine, nested parallelism not supported yet" } */
+  ;
+}