* [Patch] OpenMP: Skip target-nesting warning for reverse offload
@ 2022-05-16 15:14 Tobias Burnus
2022-05-17 18:31 ` Jakub Jelinek
0 siblings, 1 reply; 2+ messages in thread
From: Tobias Burnus @ 2022-05-16 15:14 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 482 bytes --]
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
[-- Attachment #2: omp-ancestor-nest-warn.diff --]
[-- Type: text/x-patch, Size: 2714 bytes --]
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'" } */
+ { }
+
+}
^ permalink raw reply [flat|nested] 2+ messages in thread
* Re: [Patch] OpenMP: Skip target-nesting warning for reverse offload
2022-05-16 15:14 [Patch] OpenMP: Skip target-nesting warning for reverse offload Tobias Burnus
@ 2022-05-17 18:31 ` Jakub Jelinek
0 siblings, 0 replies; 2+ messages in thread
From: Jakub Jelinek @ 2022-05-17 18:31 UTC (permalink / raw)
To: Tobias Burnus; +Cc: gcc-patches
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
^ permalink raw reply [flat|nested] 2+ messages in thread
end of thread, other threads:[~2022-05-17 18:31 UTC | newest]
Thread overview: 2+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-05-16 15:14 [Patch] OpenMP: Skip target-nesting warning for reverse offload Tobias Burnus
2022-05-17 18:31 ` Jakub Jelinek
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).