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(+)
@@ -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);
new file mode 100644
@@ -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'" } */
+ { }
+
+}