From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id 86E62383A606; Thu, 9 Jun 2022 12:48:39 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 86E62383A606 MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc r13-1024] OpenMP: Handle ancestor:1 with discover_declare_target X-Act-Checkin: gcc X-Git-Author: Tobias Burnus X-Git-Refname: refs/heads/master X-Git-Oldrev: 2dc19a1b590b9afb58fd5f779e87ae8fc1998ae7 X-Git-Newrev: 209de00fdb9da90a727337d6e752ea583a11d982 Message-Id: <20220609124839.86E62383A606@sourceware.org> Date: Thu, 9 Jun 2022 12:48:39 +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: Thu, 09 Jun 2022 12:48:39 -0000 https://gcc.gnu.org/g:209de00fdb9da90a727337d6e752ea583a11d982 commit r13-1024-g209de00fdb9da90a727337d6e752ea583a11d982 Author: Tobias Burnus Date: Thu Jun 9 14:48:24 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. Diff: --- gcc/omp-offload.cc | 21 ++--- .../c-c++-common/gomp/reverse-offload-1.c | 93 ++++++++++++++++++++++ 2 files changed, 104 insertions(+), 10 deletions(-) diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc index ad4e772015e..fcbe6cf83d8 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/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; +}