From patchwork Tue Dec 1 14:25:42 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 550913 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 5F8F41401E7 for ; Wed, 2 Dec 2015 01:27:08 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=wyImOVeG; 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=i91487Nk+2GL/mvDA calTY5snA4OxRL9BIfmRUPgaN0FF0dB03L2FXPOs17jLJVMi/udqBpnMve76Ce+S /E2wc2YhGUQ+GnzaSm1gug7hIpVqjx9TMKeqKY2jfcagehNCr/zo2lTAUPwd3f+R X5B4vXiTI4w2ahTRh9EbUiUGcg= 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=JDQ9MWJM8R281wqNoQXvZwu GLqA=; b=wyImOVeGhvGuuyJS2FSPH1K7BYIB2eW2iAjT6HC1xgUfEYhWkybo8N+ G0a7nXyEci/gMsTdiYYiUS6dpqIYLq/hkFyEMIUPiVAhLhAd7sfhTp9pukdgHE2C CQ9RZSQTSj94UuPeGRUdiUiuPEYfYuJMN5kMjHLBb3nzF5zuD37g= Received: (qmail 3996 invoked by alias); 1 Dec 2015 14:27:00 -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 3979 invoked by uid 89); 1 Dec 2015 14:26:59 -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, SPF_PASS, T_RP_MATCHES_RCVD autolearn=ham version=3.3.2 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; Tue, 01 Dec 2015 14:26:57 +0000 Received: from eggs.gnu.org ([2001:4830:134:3::10]:46101) by fencepost.gnu.org with esmtps (TLS1.0:RSA_AES_256_CBC_SHA1:256) (Exim 4.82) (envelope-from ) id 1a3ltP-0004gz-2p for gcc-patches@gnu.org; Tue, 01 Dec 2015 09:26:55 -0500 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1a3ltK-0003gp-J7 for gcc-patches@gnu.org; Tue, 01 Dec 2015 09:26:54 -0500 Received: from relay1.mentorg.com ([192.94.38.131]:58831) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1a3ltK-0003gi-Ae for gcc-patches@gnu.org; Tue, 01 Dec 2015 09:26:50 -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 1a3ltG-0007Ht-52 from Tom_deVries@mentor.com ; Tue, 01 Dec 2015 06:26:46 -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; Tue, 1 Dec 2015 14:26:44 +0000 Subject: [PATCH] Handle BUILT_IN_GOACC_PARALLEL in ipa-pta To: Richard Biener References: <565C0F47.5020604@mentor.com> <565C3CEC.9040209@mentor.com> <565C7B09.6000206@mentor.com> CC: Jakub Jelinek , "gcc-patches@gnu.org" From: Tom de Vries Message-ID: <565DADE6.8020908@mentor.com> Date: Tue, 1 Dec 2015 15:25:42 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 In-Reply-To: <565C7B09.6000206@mentor.com> X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 [ was: Re: [PATCH, PR46032] Handle BUILT_IN_GOMP_PARALLEL in ipa-pta ] On 30/11/15 17:36, Tom de Vries wrote: > On 30/11/15 14:24, Richard Biener wrote: >> On Mon, 30 Nov 2015, Tom de Vries wrote: >> >>> On 30/11/15 10:16, Richard Biener wrote: >>>> On Mon, 30 Nov 2015, Tom de Vries wrote: >>>> >>>>> Hi, >>>>> >>>>> this patch fixes PR46032. >>>>> >>>>> It handles a call: >>>>> ... >>>>> __builtin_GOMP_parallel (fn, data, num_threads, flags) >>>>> ... >>>>> as: >>>>> ... >>>>> fn (data) >>>>> ... >>>>> in ipa-pta. >>>>> >>>>> This improves ipa-pta alias analysis in the parallelized function >>>>> fn, This follow-up patch does the same for BUILT_IN_GOACC_PARALLEL. Bootstrapped and reg-tested on x86_64. OK for stage3 trunk? Thanks, - Tom Handle BUILT_IN_GOACC_PARALLEL in ipa-pta 2015-12-01 Tom de Vries * tree-ssa-structalias.c (find_func_aliases_for_builtin_call) (find_func_clobbers, ipa_pta_execute): Handle BUILT_IN_GOACC_PARALLEL. * c-c++-common/goacc/kernels-alias-ipa-pta-2.c: New test. * c-c++-common/goacc/kernels-alias-ipa-pta-3.c: New test. * c-c++-common/goacc/kernels-alias-ipa-pta.c: New test. --- .../c-c++-common/goacc/kernels-alias-ipa-pta-2.c | 37 ++++++++++++++++++++++ .../c-c++-common/goacc/kernels-alias-ipa-pta-3.c | 36 +++++++++++++++++++++ .../c-c++-common/goacc/kernels-alias-ipa-pta.c | 23 ++++++++++++++ gcc/tree-ssa-structalias.c | 28 +++++++++++++--- .../kernels-alias-ipa-pta-2.c | 27 ++++++++++++++++ .../kernels-alias-ipa-pta-3.c | 26 +++++++++++++++ .../kernels-alias-ipa-pta.c | 26 +++++++++++++++ 7 files changed, 199 insertions(+), 4 deletions(-) diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-2.c new file mode 100644 index 0000000..f16d698 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-2.c @@ -0,0 +1,37 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fipa-pta -fdump-tree-optimized" } */ + +#ifdef __cplusplus +extern "C" { +#endif +typedef __SIZE_TYPE__ size_t; +void *malloc (size_t); +void free (void *); +#ifdef __cplusplus +} +#endif + +#define N 2 + +void +foo (void) +{ + unsigned int *a = (unsigned int *)malloc (N * sizeof (unsigned int)); + unsigned int *b = (unsigned int *)malloc (N * sizeof (unsigned int)); + unsigned int *c = (unsigned int *)malloc (N * sizeof (unsigned int)); + +#pragma acc kernels pcopyout (a[0:N], b[0:N], c[0:N]) + { + a[0] = 0; + b[0] = 1; + c[0] = a[0]; + } + + free (a); + free (b); + free (c); +} + +/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 2 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= \\*a" 0 "optimized" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c new file mode 100644 index 0000000..1eb56eb --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c @@ -0,0 +1,36 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fipa-pta -fdump-tree-optimized" } */ + +#ifdef __cplusplus +extern "C" { +#endif +typedef __SIZE_TYPE__ size_t; +void *malloc (size_t); +void free (void *); +#ifdef __cplusplus +} +#endif + +#define N 2 + +void +foo (void) +{ + unsigned int *a = (unsigned int *)malloc (N * sizeof (unsigned int)); + unsigned int *b = a; + unsigned int *c = (unsigned int *)malloc (N * sizeof (unsigned int)); + +#pragma acc kernels pcopyout (a[0:N], b[0:N], c[0:N]) + { + a[0] = 0; + b[0] = 1; + c[0] = a[0]; + } + + free (a); + free (c); +} + +/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 1 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= \\*a" 1 "optimized" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta.c new file mode 100644 index 0000000..969b466 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta.c @@ -0,0 +1,23 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fipa-pta -fdump-tree-optimized" } */ + +#define N 2 + +void +foo (void) +{ + unsigned int a[N]; + unsigned int b[N]; + unsigned int c[N]; + +#pragma acc kernels pcopyout (a, b, c) + { + a[0] = 0; + b[0] = 1; + c[0] = a[0]; + } +} + +/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 2 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= \\*_\[0-9\]\\\[0\\\];$" 0 "optimized" } } */ diff --git a/gcc/tree-ssa-structalias.c b/gcc/tree-ssa-structalias.c index 7f4a8ad..060ff3e 100644 --- a/gcc/tree-ssa-structalias.c +++ b/gcc/tree-ssa-structalias.c @@ -4507,15 +4507,32 @@ find_func_aliases_for_builtin_call (struct function *fn, gcall *t) return true; } case BUILT_IN_GOMP_PARALLEL: + case BUILT_IN_GOACC_PARALLEL: { - /* Handle __builtin_GOMP_parallel (fn, data, num_threads, flags) as - fn (data). */ if (in_ipa_mode) { - tree fnarg = gimple_call_arg (t, 0); + unsigned int fnpos, argpos; + switch (DECL_FUNCTION_CODE (fndecl)) + { + case BUILT_IN_GOMP_PARALLEL: + /* __builtin_GOMP_parallel (fn, data, num_threads, flags). */ + fnpos = 0; + argpos = 1; + break; + case BUILT_IN_GOACC_PARALLEL: + /* __builtin_GOACC_parallel (device, fn, mapnum, hostaddrs, + sizes, kinds, ...). */ + fnpos = 1; + argpos = 3; + break; + default: + gcc_unreachable (); + } + + tree fnarg = gimple_call_arg (t, fnpos); gcc_assert (TREE_CODE (fnarg) == ADDR_EXPR); tree fndecl = TREE_OPERAND (fnarg, 0); - tree arg = gimple_call_arg (t, 1); + tree arg = gimple_call_arg (t, argpos); gcc_assert (TREE_CODE (arg) == ADDR_EXPR); varinfo_t fi = get_vi_for_tree (fndecl); @@ -5064,6 +5081,7 @@ find_func_clobbers (struct function *fn, gimple *origt) case BUILT_IN_VA_END: return; case BUILT_IN_GOMP_PARALLEL: + case BUILT_IN_GOACC_PARALLEL: return; /* printf-style functions may have hooks to set pointers to point to somewhere into the generated string. Leave them @@ -7547,6 +7565,8 @@ ipa_pta_execute (void) /* Handle direct calls to functions with body. */ if (gimple_call_builtin_p (stmt, BUILT_IN_GOMP_PARALLEL)) decl = TREE_OPERAND (gimple_call_arg (stmt, 0), 0); + else if (gimple_call_builtin_p (stmt, BUILT_IN_GOACC_PARALLEL)) + decl = TREE_OPERAND (gimple_call_arg (stmt, 1), 0); else decl = gimple_call_fndecl (stmt); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c new file mode 100644 index 0000000..0f323c8 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c @@ -0,0 +1,27 @@ +/* { dg-additional-options "-O2 -fipa-pta" } */ + +#include + +#define N 2 + +int +main (void) +{ + unsigned int *a = (unsigned int *)malloc (N * sizeof (unsigned int)); + unsigned int *b = (unsigned int *)malloc (N * sizeof (unsigned int)); + unsigned int *c = (unsigned int *)malloc (N * sizeof (unsigned int)); + +#pragma acc kernels pcopyout (a[0:N], b[0:N], c[0:N]) + { + a[0] = 0; + b[0] = 1; + c[0] = a[0]; + } + + if (a[0] != 0 || b[0] != 1 || c[0] != 0) + abort (); + + free (a); + free (b); + free (c); +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c new file mode 100644 index 0000000..654e750 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c @@ -0,0 +1,26 @@ +/* { dg-additional-options "-O2 -fipa-pta" } */ + +#include + +#define N 2 + +int +main (void) +{ + unsigned int *a = (unsigned int *)malloc (N * sizeof (unsigned int)); + unsigned int *b = a; + unsigned int *c = (unsigned int *)malloc (N * sizeof (unsigned int)); + +#pragma acc kernels pcopyout (a[0:N], b[0:N], c[0:N]) + { + a[0] = 0; + b[0] = 1; + c[0] = a[0]; + } + + if (a[0] != 1 || b[0] != 1 || c[0] != 1) + abort (); + + free (a); + free (c); +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c new file mode 100644 index 0000000..44d4fd2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c @@ -0,0 +1,26 @@ +/* { dg-additional-options "-O2 -fipa-pta" } */ + +#include + +#define N 2 + +int +main (void) +{ + unsigned int a[N]; + unsigned int b[N]; + unsigned int c[N]; + +#pragma acc kernels pcopyout (a, b, c) + { + a[0] = 0; + b[0] = 1; + c[0] = a[0]; + } + + if (a[0] != 0 || b[0] != 1 || c[0] != 0) + abort (); + + return 0; +} +