From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id 2813B385AE4F; Tue, 5 Jul 2022 10:32:04 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 2813B385AE4F Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-12] OpenMP: Handle ancestor:1 with discover_declare_target X-Act-Checkin: gcc X-Git-Author: Tobias Burnus X-Git-Refname: refs/heads/devel/omp/gcc-12 X-Git-Oldrev: 6c0936c4c9a91cd56be2fa08a0f9b5f78e05d25e X-Git-Newrev: 0bded9db16c93b1a713733c56de806a7a12858fe Message-Id: <20220705103204.2813B385AE4F@sourceware.org> Date: Tue, 5 Jul 2022 10:32:04 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Tue, 05 Jul 2022 10:32:04 -0000 https://gcc.gnu.org/g:0bded9db16c93b1a713733c56de806a7a12858fe commit 0bded9db16c93b1a713733c56de806a7a12858fe Author: Tobias Burnus 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 + + Backport from mainline: + 2022-06-09 Jakub Jelinek + + * 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 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 + + Backport from mainline: + 2022-06-08 Tobias Burnus + + * openmp.cc (gfc_match_omp_clauses): Check also parent namespace + for 'requires reverse_offload'. + 2022-07-05 Tobias Burnus 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 + + Backport from mainline: + 2022-06-09 Jakub Jelinek + + * c-c++-common/gomp/reverse-offload-1.c: New. + +2022-07-05 Tobias Burnus + + Backport from mainline: + 2022-06-08 Tobias Burnus + + * gfortran.dg/gomp/target-device-ancestor-5.f90: New test. + 2022-07-05 Tobias Burnus 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; +}