diff mbox

[gomp4,committed] Handle oacc region in oacc routine

Message ID 5620C0A5.20506@mentor.com
State New
Headers show

Commit Message

Tom de Vries Oct. 16, 2015, 9:17 a.m. UTC
Hi,

this patch checks for occurance of oacc offload regions in oacc routines 
(which means nested parallelism, which is currently not supported) and 
gives an appropriate error message.

Committed to gomp-4_0-branch.

Thanks,
- Tom

Comments

Thomas Schwinge Oct. 22, 2015, 9:10 a.m. UTC | #1
Hi Tom!

On Fri, 16 Oct 2015 11:17:25 +0200, Tom de Vries <Tom_deVries@mentor.com> wrote:
> this patch checks for occurance of oacc offload regions in oacc routines 
> (which means nested parallelism, which is currently not supported) and 
> gives an appropriate error message.

> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/parallel-in-routine.c
> @@ -0,0 +1,8 @@
> +#pragma acc routine
> +void
> +foo (void)
> +{
> +#pragma acc parallel /* { dg-error "OpenACC region inside of OpenACC routine, nested parallelism not supported yet" } */
> +  ;
> +}

Please move such tests into
gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c and/or
gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c, and nesting tests
that are expected to succeed into
gcc/testsuite/c-c++-common/goacc/nesting-1.c and/or
gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c.  (I have not looked up
the corresponding Fortran testsuite files.)


When fixing that, please also add back (checking for appropriate
diagnostics) the tests that Nathan apparently didn't know where to put,
and thus removed, in
<http://news.gmane.org/find-root.php?message_id=%3C56192805.6090200%40acm.org%3E>
(gomp-4_0-branch r228678), and
<http://news.gmane.org/find-root.php?message_id=%3C561AE027.9040601%40acm.org%3E>
(gomp-4_0-branch r228694).


Grüße
 Thomas
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/parallel-in-routine.c: New test.
---
 gcc/omp-low.c                                          | 9 +++++++++
 gcc/testsuite/c-c++-common/goacc/parallel-in-routine.c | 8 ++++++++
 2 files changed, 17 insertions(+)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/parallel-in-routine.c

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index f27bde7..f7e4afc 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -3204,6 +3204,15 @@  check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
 	}
       break;
     case GIMPLE_OMP_TARGET:
+      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/parallel-in-routine.c b/gcc/testsuite/c-c++-common/goacc/parallel-in-routine.c
new file mode 100644
index 0000000..b93d63b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/parallel-in-routine.c
@@ -0,0 +1,8 @@ 
+#pragma acc routine
+void
+foo (void)
+{
+#pragma acc parallel /* { dg-error "OpenACC region inside of OpenACC routine, nested parallelism not supported yet" } */
+  ;
+}
+
-- 
1.9.1