Handle BUILT_IN_GOACC_PARALLEL in ipa-pta

Message ID	565DADE6.8020908@mentor.com
State	New
Headers	show Return-Path: <gcc-patches-return-415919-incoming=patchwork.ozlabs.org@gcc.gnu.org> 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= Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk Sender: gcc-patches-owner@gcc.gnu.org Subject: [PATCH] Handle BUILT_IN_GOACC_PARALLEL in ipa-pta To: Richard Biener <rguenther@suse.de> References: <565C0F47.5020604@mentor.com> <alpine.LSU.2.11.1511301010570.4884@t29.fhfr.qr> <565C3CEC.9040209@mentor.com> <alpine.LSU.2.11.1511301423530.4884@t29.fhfr.qr> <565C7B09.6000206@mentor.com> CC: Jakub Jelinek <jakub@redhat.com>, "gcc-patches@gnu.org" <gcc-patches@gnu.org> From: Tom de Vries <Tom_deVries@mentor.com> 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> Content-Type: multipart/mixed; boundary="------------080808060301020900070402"

Handle BUILT_IN_GOACC_PARALLEL in ipa-pta 2015-12-01 Tom de Vries <tom@codesourcery.com> * 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 <stdlib.h> + +#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 <stdlib.h> + +#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 <stdlib.h> + +#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; +} +

Handle BUILT_IN_GOACC_PARALLEL in ipa-pta

Commit Message

Comments

Patch