[PATCH] Handle oacc region in oacc routine

[PATCH] Handle oacc region in oacc routine

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 383 bytes --]

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

[-- Attachment #2: 0001-Handle-oacc-region-in-oacc-routine.patch --]
[-- Type: text/x-patch, Size: 1683 bytes --]

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" } */
+  ;
+}

Re: [PATCH] Handle oacc region in oacc routine

[Jakub Jelinek]

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

Re: [PATCH] Handle oacc region in oacc routine

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1096 bytes --]

On 01-03-16 12:37, Jakub Jelinek wrote:
> 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...

Hi,

indeed, thanks for pointing that out.

I've added:
- detection of that case earlier in the same function,
- an appropriate error message and
- a corresponding test-case.

Build c compiler, ran goacc.exp, gomp.exp, goacc-gomp.exp and libgomp c.exp.

OK for stage4 trunk if full bootstrap and reg-test succeeds?

Thanks,
- Tom


[-- Attachment #2: 0001-Handle-oacc-region-in-oacc-routine.patch --]
[-- Type: text/x-patch, Size: 3897 bytes --]

Handle oacc region in oacc routine

2015-10-16  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.

---
 gcc/omp-low.c                                      | 41 +++++++++++++++-------
 .../c-c++-common/goacc-gomp/nesting-fail-1.c       |  9 +++++
 gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c  |  8 +++++
 3 files changed, 45 insertions(+), 13 deletions(-)

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 989d03e..487eccc 100644
--- 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/c-c++-common/goacc-gomp/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
index 1a44721..728daf1 100644
--- 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,12 @@ 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 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" } */
+  ;
+}

Re: [PATCH] Handle oacc region in oacc routine

[Jakub Jelinek]

On Wed, Mar 02, 2016 at 12:03:12PM -0500, Tom de Vries wrote:
> 2015-10-16  Tom de Vries  <tom@codesourcery.com>

Please adjust the date ;)

> 	* 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.

Ok for trunk.

	Jakub