diff mbox

Handle oacc region in oacc routine

Message ID 56D5D06A.8020607@mentor.com
State New
Headers show

Commit Message

Tom de Vries March 1, 2016, 5:24 p.m. UTC 
Hi,

this patch fixes an ICE in an openacc testcase. The patch fixes it by emitting 
an 'unsupported' error.

We've been carrying this patch for a while in the gomp-4_0-branch ( 
https://gcc.gnu.org/ml/gcc-patches/2015-10/msg01545.html ).

Build for C-only on top of trunk, ran goacc.exp regression test.

OK for stage4 trunk, if complete bootstrap/reg-test succeeds?

Thanks,
- Tom

Comments

Jakub Jelinek March 1, 2016, 5:37 p.m. UTC | #1 
On Tue, Mar 01, 2016 at 06:24:58PM +0100, Tom de Vries wrote:
> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -3715,6 +3715,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;
> +	}

Won't this emit the same error message even for
#pragma omp target
inside of #pragma acc routine function?  That would be misleading...

	Jakub
diff mbox

Patch

Handle oacc region in oacc routine

2015-10-16  Tom de Vries  <tom@codesourcery.com>

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

	* c-c++-common/goacc/nesting-fail-1.c (f_acc_routine): Add oacc region
	in oacc routine test.

---
 gcc/omp-low.c                                     | 8 ++++++++
 gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c | 8 ++++++++
 2 files changed, 16 insertions(+)

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 989d03e..e84277b 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -3715,6 +3715,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/c-c++-common/goacc/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
index 7a36074..506a1ae 100644
--- 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" } */
+  ;
+}