diff mbox series

OpenMP: Skip target-nesting warning for reverse offload

Message ID 9d5cb5b1-6b58-86b6-403f-e2ab9dd38d2e@codesourcery.com
State New
Headers
Series OpenMP: Skip target-nesting warning for reverse offload |

Commit Message

Tobias Burnus May 16, 2022, 3:14 p.m. UTC 
  A warning about target-nesting inside target makes sense,
but not if the inner target is one for reverse offload ("device(ancestor:1)").

Thus, silence the warning in this case.

OK for mainline?

Tobias
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

Comments

Jakub Jelinek May 17, 2022, 6:31 p.m. UTC | #1 
On Mon, May 16, 2022 at 05:14:12PM +0200, Tobias Burnus wrote:
> --- a/gcc/omp-low.cc
> +++ b/gcc/omp-low.cc
> @@ -3883,6 +3883,16 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
>  		}
>  	      else
>  		{
> +		  if ((gimple_omp_target_kind (ctx->stmt)
> +		       == GF_OMP_TARGET_KIND_REGION)
> +		      && (gimple_omp_target_kind (stmt)
> +			  == GF_OMP_TARGET_KIND_REGION)
> +		      && ((c = omp_find_clause (
> +				 gimple_omp_target_clauses (stmt),
> +				 OMP_CLAUSE_DEVICE))
> +			  != NULL_TREE)
> +		      && OMP_CLAUSE_DEVICE_ANCESTOR (c))
> +		    break;

The ( at the end of line is too ugly for me.
Can't you write it as:
		  if ((gimple_omp_target_kind (ctx->stmt)
		       == GF_OMP_TARGET_KIND_REGION)
		      && (gimple_omp_target_kind (stmt)
			  == GF_OMP_TARGET_KIND_REGION))
		    {
		      c = omp_find_clause (gimple_omp_target_clauses (stmt),
					   OMP_CLAUSE_DEVICE);
		      if (c && OMP_CLAUSE_DEVICE_ANCESTOR (c))
			break;
		    }
?

Otherwise LGTM.

	Jakub
diff mbox series

Patch

OpenMP: Skip target-nesting warning for reverse offload

gcc/ChangeLog:

	* omp-low.cc (check_omp_nesting_restrictions): Skip warning for
	target inside target if inner is reverse offload.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/target-device-ancestor-5.c: New test.

 omp-low.cc                                             |   10 ++++++++++
 testsuite/c-c++-common/gomp/target-device-ancestor-5.c |   28 ++++++++++++++++++++++++++++
 2 files changed, 38 insertions(+)

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 8aebaeecd42..f7f44fca450 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -3883,6 +3883,16 @@  check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
 		}
 	      else
 		{
+		  if ((gimple_omp_target_kind (ctx->stmt)
+		       == GF_OMP_TARGET_KIND_REGION)
+		      && (gimple_omp_target_kind (stmt)
+			  == GF_OMP_TARGET_KIND_REGION)
+		      && ((c = omp_find_clause (
+				 gimple_omp_target_clauses (stmt),
+				 OMP_CLAUSE_DEVICE))
+			  != NULL_TREE)
+		      && OMP_CLAUSE_DEVICE_ANCESTOR (c))
+		    break;
 		  warning_at (gimple_location (stmt), 0,
 			      "%qs construct inside of %qs region",
 			      stmt_name, ctx_stmt_name);
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c
new file mode 100644
index 00000000000..b6ff84bcdab
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c
@@ -0,0 +1,28 @@ 
+#pragma omp requires reverse_offload  /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+
+void
+foo ()
+{
+  /* Good nesting - as reverse offload */
+  #pragma omp target
+   #pragma omp target device(ancestor:1)  /* valid -> no warning */   /* { dg-bogus "'target' construct inside of 'target' region" }  */
+    { }
+
+  /* Bad nesting */
+  #pragma omp target
+   #pragma omp target  /* { dg-warning "'target' construct inside of 'target' region" }  */
+     #pragma omp target  /* { dg-warning "'target' construct inside of 'target' region" }  */
+    { }
+
+  /* Good nesting - as reverse offload */
+  #pragma omp target
+   #pragma omp target  /* { dg-warning "'target' construct inside of 'target' region" }  */
+     #pragma omp target device(ancestor:1)  /* valid -> no warning */   /* { dg-bogus "'target' construct inside of 'target' region" }  */
+      { }
+
+  #pragma omp target
+   #pragma omp target device(ancestor:1)  /* valid -> no warning */   /* { dg-bogus "'target' construct inside of 'target' region" }  */
+     #pragma omp target device(ancestor:1) /* { dg-error "OpenMP constructs are not allowed in target region with 'ancestor'" }  */
+       { }
+
+}