From d89bf164516d06684f120a42d9f8d351f8e5dd2f Mon Sep 17 00:00:00 2001 From: Tobias Burnus Date: Tue, 5 Jul 2022 08:36:42 +0200 Subject: [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. (cherry picked from commit 47554478a13f64bff1ee4b9bb0319ae63d42ca52) --- gcc/ChangeLog.omp | 8 ++++++ gcc/omp-low.cc | 10 +++++++ gcc/testsuite/ChangeLog.omp | 7 +++++ .../gomp/target-device-ancestor-5.c | 28 +++++++++++++++++++ 4 files changed, 53 insertions(+) create mode 100644 gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index e146e14e6f8e..06a194e744c4 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,11 @@ +2022-07-05 Tobias Burnus + + Backport from mainline: + 2022-05-17 Tobias Burnus + + * omp-low.cc (check_omp_nesting_restrictions): Skip warning for + target inside target if inner is reverse offload. + 2022-07-04 Tobias Burnus Backport from mainline: diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index c366fb55c806..adefaa527a72 100644 --- 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 0bfd4f2a3b72..4bd6d97f8a0c 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,10 @@ +2022-07-05 Tobias Burnus + + Backport from mainline: + 2022-05-17 Tobias Burnus + + * c-c++-common/gomp/target-device-ancestor-5.c: New test. + 2022-07-04 Tobias Burnus 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 index 000000000000..b6ff84bcdab9 --- /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'" } */ + { } + +} -- 2.47.2