public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] OpenMP: Handle ancestor:1 with discover_declare_target
@ 2022-07-05 10:32 Tobias Burnus
0 siblings, 0 replies; only message in thread
From: Tobias Burnus @ 2022-07-05 10:32 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:0bded9db16c93b1a713733c56de806a7a12858fe
commit 0bded9db16c93b1a713733c56de806a7a12858fe
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Tue Jul 5 11:18:23 2022 +0200
OpenMP: Handle ancestor:1 with discover_declare_target
gcc/
* omp-offload.cc (omp_discover_declare_target_tgt_fn_r,
omp_discover_declare_target_fn_r): Don't walk reverse-offload
target regions.
gcc/testsuite/
* c-c++-common/gomp/reverse-offload-1.c: New.
(cherry picked from commit 209de00fdb9da90a727337d6e752ea583a11d982)
Diff:
---
gcc/ChangeLog.omp | 9 +++
gcc/fortran/ChangeLog.omp | 8 ++
gcc/omp-offload.cc | 21 ++---
gcc/testsuite/ChangeLog.omp | 14 ++++
.../c-c++-common/gomp/reverse-offload-1.c | 93 ++++++++++++++++++++++
5 files changed, 135 insertions(+), 10 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 800b902fcb2..a39216bb7d4 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,12 @@
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-06-09 Jakub Jelinek <jakub@redhat.com>
+
+ * omp-offload.cc (omp_discover_declare_target_tgt_fn_r,
+ omp_discover_declare_target_fn_r): Don't walk reverse-offload
+ target regions.
+
2022-07-05 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index 39d2143b0cf..37a18320bbc 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,11 @@
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-06-08 Tobias Burnus <tobias@codesourcery.com>
+
+ * openmp.cc (gfc_match_omp_clauses): Check also parent namespace
+ for 'requires reverse_offload'.
+
2022-07-05 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc
index 6656bc48e95..2b9aa87dbbf 100644
--- a/gcc/omp-offload.cc
+++ b/gcc/omp-offload.cc
@@ -268,12 +268,12 @@ omp_discover_declare_target_tgt_fn_r (tree *tp, int *walk_subtrees, void *data)
}
else if (TYPE_P (*tp))
*walk_subtrees = 0;
- /* else if (TREE_CODE (*tp) == OMP_TARGET)
- {
- if (tree dev = omp_find_clause (OMP_TARGET_CLAUSES (*tp)))
- if (OMP_DEVICE_ANCESTOR (dev))
- *walk_subtrees = 0;
- } */
+ else if (TREE_CODE (*tp) == OMP_TARGET)
+ {
+ tree c = omp_find_clause (OMP_CLAUSES (*tp), OMP_CLAUSE_DEVICE);
+ if (c && OMP_CLAUSE_DEVICE_ANCESTOR (c))
+ *walk_subtrees = 0;
+ }
return NULL_TREE;
}
@@ -284,10 +284,11 @@ omp_discover_declare_target_fn_r (tree *tp, int *walk_subtrees, void *data)
{
if (TREE_CODE (*tp) == OMP_TARGET)
{
- /* And not OMP_DEVICE_ANCESTOR. */
- walk_tree_without_duplicates (&OMP_TARGET_BODY (*tp),
- omp_discover_declare_target_tgt_fn_r,
- data);
+ tree c = omp_find_clause (OMP_CLAUSES (*tp), OMP_CLAUSE_DEVICE);
+ if (!c || !OMP_CLAUSE_DEVICE_ANCESTOR (c))
+ walk_tree_without_duplicates (&OMP_TARGET_BODY (*tp),
+ omp_discover_declare_target_tgt_fn_r,
+ data);
*walk_subtrees = 0;
}
else if (TYPE_P (*tp))
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index bec3740ca7d..00f06f3fba5 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,17 @@
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-06-09 Jakub Jelinek <jakub@redhat.com>
+
+ * c-c++-common/gomp/reverse-offload-1.c: New.
+
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-06-08 Tobias Burnus <tobias@codesourcery.com>
+
+ * gfortran.dg/gomp/target-device-ancestor-5.f90: New test.
+
2022-07-05 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
diff --git a/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c b/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c
new file mode 100644
index 00000000000..9a3fa5230f8
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c
@@ -0,0 +1,93 @@
+/* { dg-additional-options "-fdump-tree-omplower" } */
+
+/* { dg-final { scan-tree-dump-times "omp declare target\[^ \]" 3 "omplower" } } */
+
+/* { dg-final { scan-tree-dump-times "__attribute__\\(\\(omp declare target\\)\\)\[\n\r\]*int called_in_target1" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "__attribute__\\(\\(omp declare target\\)\\)\[\n\r\]*int called_in_target2" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "__attribute__\\(\\(omp declare target, omp declare target block\\)\\)\[\n\r\]*void tg_fn" 1 "omplower" } } */
+
+/* { dg-prune-output "'reverse_offload' clause on 'requires' directive not supported yet" } */
+
+#pragma omp requires reverse_offload
+
+extern int add_3 (int);
+
+static int global_var = 5;
+
+void
+check_offload (int *x, int *y)
+{
+ *x = add_3 (*x);
+ *y = add_3 (*y);
+}
+
+int
+called_in_target1 ()
+{
+ return 42;
+}
+
+int
+called_in_target2 ()
+{
+ return -6;
+}
+
+#pragma omp declare target
+void
+tg_fn (int *x, int *y)
+{
+ int x2 = *x, y2 = *y;
+ if (x2 != 2 || y2 != 3)
+ __builtin_abort ();
+ x2 = x2 + 2 + called_in_target1 ();
+ y2 = y2 + 7;
+
+ #pragma omp target device(ancestor : 1) map(tofrom: x2)
+ check_offload(&x2, &y2);
+
+ if (x2 != 2+2+3+42 || y2 != 3 + 7)
+ __builtin_abort ();
+ *x = x2, *y = y2;
+}
+#pragma omp end declare target
+
+void
+my_func (int *x, int *y)
+{
+ if (global_var != 5)
+ __builtin_abort ();
+ global_var = 242;
+ *x = 2*add_3(*x);
+ *y = 3*add_3(*y);
+}
+
+int
+main ()
+{
+ #pragma omp target
+ {
+ int x = 2, y = 3;
+ tg_fn (&x, &y);
+ }
+
+ #pragma omp target
+ {
+ int x = -2, y = -1;
+ x += called_in_target2 ();
+ #pragma omp target device ( ancestor:1 ) firstprivate(y) map(tofrom:x)
+ {
+ if (x != -2-6 || y != -1)
+ __builtin_abort ();
+ my_func (&x, &y);
+ if (x != 2*(3-2) || y != 3*(3-1))
+ __builtin_abort ();
+ }
+ if (x != 2*(3-2) || y != -1)
+ __builtin_abort ();
+ }
+
+ if (global_var != 242)
+ __builtin_abort ();
+ return 0;
+}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-07-05 10:32 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-07-05 10:32 [gcc/devel/omp/gcc-12] OpenMP: Handle ancestor:1 with discover_declare_target Tobias Burnus
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).