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.

(cherry picked from commit 47554478a13f64bff1ee4b9bb0319ae63d42ca52)

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index e146e14e6f8e3960e35189744346e81fca4be9c8..06a194e744c47f9eec1522f508fd1532c3618c9a 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,11 @@
+2022-07-05  Tobias Burnus  <tobias@codesourcery.com>
+
+       Backport from mainline:
+       2022-05-17  Tobias Burnus  <tobias@codesourcery.com>
+
+       * omp-low.cc (check_omp_nesting_restrictions): Skip warning for
+       target inside target if inner is reverse offload.
+
 2022-07-04  Tobias Burnus  <tobias@codesourcery.com>
 
        Backport from mainline:

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index c366fb55c806b240df8c4d4c3b8099ab3365ca2a..adefaa527a72114433a48a03b112319e05fec57d 100644 (file)
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -4377,6 +4377,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);
+                     if (c && 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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 0bfd4f2a3b728c58c0c1e093072be583fc287541..4bd6d97f8a0c83d101ed92e5a5edd63b0cac8a27 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,10 @@
+2022-07-05  Tobias Burnus  <tobias@codesourcery.com>
+
+       Backport from mainline:
+       2022-05-17  Tobias Burnus  <tobias@codesourcery.com>
+
+       * c-c++-common/gomp/target-device-ancestor-5.c: New test.
+
 2022-07-04  Tobias Burnus  <tobias@codesourcery.com>
 
        Backport from mainline:

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 (file)
index 0000000..b6ff84b
--- /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'" }  */
+       { }
+
+}