From patchwork Thu Jun 9 10:47:17 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 1641157 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Received: from sourceware.org (ip-8-43-85-97.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (2048 bits) server-digest SHA256) (No client certificate requested) by bilbo.ozlabs.org (Postfix) with ESMTPS id 4LJgly412jz9sFk for ; Thu, 9 Jun 2022 20:47:41 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 96048383D809 for ; Thu, 9 Jun 2022 10:47:37 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 6957838582A5 for ; Thu, 9 Jun 2022 10:47:24 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 6957838582A5 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.91,287,1647331200"; d="diff'?scan'208";a="79681100" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 09 Jun 2022 02:47:23 -0800 IronPort-SDR: HCzK1y/1/DntoOSxUMac/0eA5OmmAqZKJurMlt91fD7I4PyEyI+9Z2Fow/941j49O7t8NivOrq rshT9XHMIF7ooV8qzNc+g7bEm8PZgtMWpnMMFNxBTostJO5df+57akKLhoHL+A36VmJscIpyyl sqemRO5YlJ9VJ7EhuJdQOhizhZt8cmvvT1RIIbsJb7QnYUmao9LgqUPmUH99NBM3XVTEaO9dN+ YPyLP+K1gBJpbGNdUTwk9aEiuVqSocyEgmN6BhUgu835VepY+pujoX+rH4Uok3J+8t02ODnCM1 /7k= Message-ID: Date: Thu, 9 Jun 2022 12:47:17 +0200 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:91.0) Gecko/20100101 Thunderbird/91.10.0 Content-Language: en-US To: gcc-patches , Jakub Jelinek From: Tobias Burnus Subject: [Patch] OpenMP: Handle ancestor:1 with discover_declare_target X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-06.mgc.mentorg.com (139.181.222.6) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-11.4 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" Another minor step to getting reverse offloading to work ... OK for mainline? Tobias PS: As attached, this patch is a stand-alone patch, which fails due to the requires sorry (see dg-prune-output). With the requires patch, it should fail with the next sorry: 'ancestor' not yet supported" (For Fortran, it currently fails already in the FE as the sorry prevents generating tree code. Thus, there cannot be a Fortran check until the requires patch is in.) ----------------- 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 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 --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; +}