From patchwork Wed Feb 17 15:48:40 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 584205 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 175991402A1 for ; Thu, 18 Feb 2016 02:49:16 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=T9A1aKsG; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=HTQlbJUUZqZwrexhA r6a8943ijgF9/9LJ3m7vL8eWju7x62nzBDxVR0C1Ka51QBIF9edlzTN93WQoz3io PW8rz45QlSJzVgKe1TIJNS9RXyj2ZiVpWsoRJTgjeSXf0wq21pDLOGQ9LZ9LG/4h r5BEQ2Lb+LE19wN+syp/mXKIfs= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=WXZPV43Y5kg+p3qzdJ9zlXk WY/A=; b=T9A1aKsGTmvNQcfD4/SSUccTRGr3x6yHK9MTCGR6IqW5e+0cfOCYTIb 67ZR+2ENJSBDU4Mml9Ik+ZSmNr97ChoOZwBO5yiKkhks/Puu0hOrcfJ+TMJ4Dvdr A+15enmtJMo7pRVlzfM3BeP7d3xY8lYmzzviK/A4JyMU73ov4/do= Received: (qmail 79988 invoked by alias); 17 Feb 2016 15:49:07 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 79419 invoked by uid 89); 17 Feb 2016 15:49:06 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL, BAYES_00, RP_MATCHES_RCVD, SPF_PASS autolearn=ham version=3.3.2 spammy=capital, Capital, HX-detected-operating-system:Windows, 198 X-HELO: fencepost.gnu.org Received: from fencepost.gnu.org (HELO fencepost.gnu.org) (208.118.235.10) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-SHA encrypted) ESMTPS; Wed, 17 Feb 2016 15:49:05 +0000 Received: from eggs.gnu.org ([2001:4830:134:3::10]:59219) by fencepost.gnu.org with esmtps (TLS1.0:RSA_AES_256_CBC_SHA1:256) (Exim 4.82) (envelope-from ) id 1aW4Ld-0003Lm-Rb for gcc-patches@gnu.org; Wed, 17 Feb 2016 10:49:02 -0500 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1aW4LX-0005Mk-IO for gcc-patches@gnu.org; Wed, 17 Feb 2016 10:49:00 -0500 Received: from relay1.mentorg.com ([192.94.38.131]:33348) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1aW4LX-0005Mb-1O for gcc-patches@gnu.org; Wed, 17 Feb 2016 10:48:55 -0500 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1aW4LQ-0002Pj-Sr from Tom_deVries@mentor.com ; Wed, 17 Feb 2016 07:48:49 -0800 Received: from [127.0.0.1] (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Wed, 17 Feb 2016 15:48:47 +0000 Subject: Re: [PATCH, PR69607] Mark offload symbols as global in lto To: Jakub Jelinek , Jan Hubicka , Richard Biener References: <56B89150.7000209@mentor.com> <56C46149.3060503@mentor.com> <20160217123002.GC3017@tucnak.redhat.com> CC: "gcc-patches@gnu.org" , Ilya Verbin From: Tom de Vries Message-ID: <56C49658.1000908@mentor.com> Date: Wed, 17 Feb 2016 16:48:40 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.5.1 MIME-Version: 1.0 In-Reply-To: <20160217123002.GC3017@tucnak.redhat.com> X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 On 17/02/16 13:30, Jakub Jelinek wrote: > On Wed, Feb 17, 2016 at 01:02:17PM +0100, Tom de Vries wrote: >> Mark offload symbols as global in lto > > I'm really not familiar with that part of LTO, so I'm CCing Honza and > Richard here. >> 2016-02-08 Tom de Vries >> >> PR lto/69607 >> * lto-partition.c (promote_offload_tables): New function. >> * lto-partition.h (promote_offload_tables): Declare. > > Just one space instead of two after : > >> * lto.c (do_whole_program_analysis): call promote_offload_tables. > > Capital C in Call. > Done. >> diff --git a/libgomp/testsuite/libgomp.c/target-37.c b/libgomp/testsuite/libgomp.c/target-37.c >> new file mode 100644 >> index 0000000..1edb21e >> --- /dev/null >> +++ b/libgomp/testsuite/libgomp.c/target-37.c >> @@ -0,0 +1,98 @@ >> +/* { dg-do run { target lto } } */ >> +/* { dg-additional-sources "target-38.c" } */ >> +/* { dg-additional-options "-flto -flto-partition=1to1 -fno-toplevel-reorder" } */ >> + >> +extern >> +#ifdef __cplusplus >> +"C" >> +#endif >> +void abort (void); > > Why the C++ stuff in there? Do you intend to include the testcase > also in libgomp.c++? No, that's just there because I started both target-37.c and target-38.c by copying target-1.c. > If not, it is not needed. Removed. > Otherwise, the tests LGTM. > Updated patch attached. Thanks, - Tom Mark offload symbols as global in lto 2016-02-17 Tom de Vries PR lto/69607 * lto-partition.c (promote_offload_tables): New function. * lto-partition.h (promote_offload_tables): Declare. * lto.c (do_whole_program_analysis): Call promote_offload_tables. * testsuite/libgomp.c/target-36.c: New test. * testsuite/libgomp.c/target-37.c: New test. * testsuite/libgomp.c/target-38.c: New test. --- gcc/lto/lto-partition.c | 28 ++++++++++ gcc/lto/lto-partition.h | 1 + gcc/lto/lto.c | 2 + libgomp/testsuite/libgomp.c/target-36.c | 4 ++ libgomp/testsuite/libgomp.c/target-37.c | 94 +++++++++++++++++++++++++++++++++ libgomp/testsuite/libgomp.c/target-38.c | 91 +++++++++++++++++++++++++++++++ 6 files changed, 220 insertions(+) diff --git a/gcc/lto/lto-partition.c b/gcc/lto/lto-partition.c index 9eb63c2..56598d4 100644 --- a/gcc/lto/lto-partition.c +++ b/gcc/lto/lto-partition.c @@ -34,6 +34,7 @@ along with GCC; see the file COPYING3. If not see #include "ipa-prop.h" #include "ipa-inline.h" #include "lto-partition.h" +#include "omp-low.h" vec ltrans_partitions; @@ -1003,6 +1004,33 @@ promote_symbol (symtab_node *node) "Promoting as hidden: %s\n", node->name ()); } +/* Promote the symbols in the offload tables. */ + +void +promote_offload_tables (void) +{ + if (vec_safe_is_empty (offload_funcs) && vec_safe_is_empty (offload_vars)) + return; + + for (unsigned i = 0; i < vec_safe_length (offload_funcs); i++) + { + tree fn_decl = (*offload_funcs)[i]; + cgraph_node *node = cgraph_node::get (fn_decl); + if (node->externally_visible) + continue; + promote_symbol (node); + } + + for (unsigned i = 0; i < vec_safe_length (offload_vars); i++) + { + tree var_decl = (*offload_vars)[i]; + varpool_node *node = varpool_node::get (var_decl); + if (node->externally_visible) + continue; + promote_symbol (node); + } +} + /* Return true if NODE needs named section even if it won't land in the partition symbol table. FIXME: we should really not use named sections for inline clones and master diff --git a/gcc/lto/lto-partition.h b/gcc/lto/lto-partition.h index 31e3764..1a38126 100644 --- a/gcc/lto/lto-partition.h +++ b/gcc/lto/lto-partition.h @@ -36,6 +36,7 @@ extern vec ltrans_partitions; void lto_1_to_1_map (void); void lto_max_map (void); void lto_balanced_map (int); +extern void promote_offload_tables (void); void lto_promote_cross_file_statics (void); void free_ltrans_partitions (void); void lto_promote_statics_nonwpa (void); diff --git a/gcc/lto/lto.c b/gcc/lto/lto.c index 9dd513f..2736c5c 100644 --- a/gcc/lto/lto.c +++ b/gcc/lto/lto.c @@ -3138,6 +3138,8 @@ do_whole_program_analysis (void) to globals with hidden visibility because they are accessed from multiple partitions. */ lto_promote_cross_file_statics (); + /* Promote all the offload symbols. */ + promote_offload_tables (); timevar_pop (TV_WHOPR_PARTITIONING); timevar_stop (TV_PHASE_OPT_GEN); diff --git a/libgomp/testsuite/libgomp.c/target-36.c b/libgomp/testsuite/libgomp.c/target-36.c new file mode 100644 index 0000000..bafb718 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-36.c @@ -0,0 +1,4 @@ +/* { dg-do run { target lto } } */ +/* { dg-additional-options "-flto -flto-partition=1to1 -fno-toplevel-reorder" } */ + +#include "target-1.c" diff --git a/libgomp/testsuite/libgomp.c/target-37.c b/libgomp/testsuite/libgomp.c/target-37.c new file mode 100644 index 0000000..fe5b8ef --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-37.c @@ -0,0 +1,94 @@ +/* { dg-do run { target lto } } */ +/* { dg-additional-sources "target-38.c" } */ +/* { dg-additional-options "-flto -flto-partition=1to1 -fno-toplevel-reorder" } */ + +extern void abort (void); + +void +fn1 (double *x, double *y, int z) +{ + int i; + for (i = 0; i < z; i++) + { + x[i] = i & 31; + y[i] = (i & 63) - 30; + } +} + +#pragma omp declare target +static int tgtv = 6; +static int +tgt (void) +{ + #pragma omp atomic update + tgtv++; + return 0; +} +#pragma omp end declare target + +static double +fn2 (int x, int y, int z) +{ + double b[1024], c[1024], s = 0; + int i, j; + fn1 (b, c, x); + #pragma omp target data map(to: b) + { + #pragma omp target map(tofrom: c, s) + #pragma omp teams num_teams(y) thread_limit(z) reduction(+:s) firstprivate(x) + #pragma omp distribute dist_schedule(static, 4) collapse(1) + for (j=0; j < x; j += y) + #pragma omp parallel for reduction(+:s) + for (i = j; i < j + y; i++) + tgt (), s += b[i] * c[i]; + #pragma omp target update from(b, tgtv) + } + return s; +} + +static double +fn3 (int x) +{ + double b[1024], c[1024], s = 0; + int i; + fn1 (b, c, x); + #pragma omp target map(to: b, c) map(tofrom:s) + #pragma omp parallel for reduction(+:s) + for (i = 0; i < x; i++) + tgt (), s += b[i] * c[i]; + return s; +} + +static double +fn4 (int x, double *p) +{ + double b[1024], c[1024], d[1024], s = 0; + int i; + fn1 (b, c, x); + fn1 (d + x, p + x, x); + #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)]) \ + map(tofrom: s) + #pragma omp parallel for reduction(+:s) + for (i = 0; i < x; i++) + s += b[i] * c[i] + d[x + i] + p[x + i]; + return s; +} + +extern int other_main (void); + +int +main () +{ + double a = fn2 (128, 4, 6); + int b = tgtv; + double c = fn3 (61); + #pragma omp target update from(tgtv) + int d = tgtv; + double e[1024]; + double f = fn4 (64, e); + if (a != 13888.0 || b != 6 + 128 || c != 4062.0 || d != 6 + 128 + 61 + || f != 8032.0) + abort (); + other_main (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-38.c b/libgomp/testsuite/libgomp.c/target-38.c new file mode 100644 index 0000000..de2b4c8 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-38.c @@ -0,0 +1,91 @@ +/* { dg-skip-if "additional source" { *-*-* } } */ + +extern void abort (void); + +static void +fna1 (double *x, double *y, int z) +{ + int i; + for (i = 0; i < z; i++) + { + x[i] = i & 31; + y[i] = (i & 63) - 30; + } +} + +#pragma omp declare target +static int tgtva = 6; +static int +tgta (void) +{ + #pragma omp atomic update + tgtva++; + return 0; +} +#pragma omp end declare target + +double +fna2 (int x, int y, int z) +{ + double b[1024], c[1024], s = 0; + int i, j; + fna1 (b, c, x); + #pragma omp target data map(to: b) + { + #pragma omp target map(tofrom: c, s) + #pragma omp teams num_teams(y) thread_limit(z) reduction(+:s) firstprivate(x) + #pragma omp distribute dist_schedule(static, 4) collapse(1) + for (j=0; j < x; j += y) + #pragma omp parallel for reduction(+:s) + for (i = j; i < j + y; i++) + tgta (), s += b[i] * c[i]; + #pragma omp target update from(b, tgtva) + } + return s; +} + +static double +fna3 (int x) +{ + double b[1024], c[1024], s = 0; + int i; + fna1 (b, c, x); + #pragma omp target map(to: b, c) map(tofrom:s) + #pragma omp parallel for reduction(+:s) + for (i = 0; i < x; i++) + tgta (), s += b[i] * c[i]; + return s; +} + +static double +fna4 (int x, double *p) +{ + double b[1024], c[1024], d[1024], s = 0; + int i; + fna1 (b, c, x); + fna1 (d + x, p + x, x); + #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)]) \ + map(tofrom: s) + #pragma omp parallel for reduction(+:s) + for (i = 0; i < x; i++) + s += b[i] * c[i] + d[x + i] + p[x + i]; + return s; +} + +extern int other_main (void); + +int +other_main (void) +{ + double a = fna2 (128, 4, 6); + int b = tgtva; + double c = fna3 (61); + #pragma omp target update from(tgtva) + int d = tgtva; + double e[1024]; + double f = fna4 (64, e); + if (a != 13888.0 || b != 6 + 128 || c != 4062.0 || d != 6 + 128 + 61 + || f != 8032.0) + abort (); + return 0; +}