Message ID | 5638C5AD.5090208@mentor.com |
---|---|
State | New |
Headers | show |
Hi Tom! On Tue, 3 Nov 2015 15:33:17 +0100, Tom de Vries <Tom_deVries@mentor.com> wrote: > On 03/11/15 15:19, Tom de Vries wrote: > > I've dropped the two testcases from this patch, I'll commit in a > > follow-up patch. > > Committed to gomp-4_0-branch, as attached. > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-offload-alias-none.c > @@ -0,0 +1,61 @@ > +/* { dg-additional-options "-O2" } */ > +/* { dg-additional-options "-fdump-tree-optimized" } */ > +/* { dg-additional-options "-fdump-tree-alias-all" } */ > +/* { dg-additional-options "-foffload-alias=none" } */ > + > +#include <stdlib.h> > + > +#define N (1024 * 512) > +#define COUNTERTYPE unsigned int > + > +static void > +foo (unsigned int *a, unsigned int *b, unsigned int *c) > +{ > + for (COUNTERTYPE i = 0; i < N; i++) > + a[i] = i * 2; > + > + for (COUNTERTYPE i = 0; i < N; i++) > + b[i] = i * 4; > + > +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) > + { > + for (COUNTERTYPE ii = 0; ii < N; ii++) > + c[ii] = a[ii] + b[ii]; > + } > + > + for (COUNTERTYPE i = 0; i < N; i++) > + if (c[i] != a[i] + b[i]) > + abort (); > +} > + > +int > +main (void) > +{ > + unsigned int *a; > + unsigned int *b; > + unsigned int *c; > + > + a = (unsigned int *)malloc (N * sizeof (unsigned int)); > + b = (unsigned int *)malloc (N * sizeof (unsigned int)); > + c = (unsigned int *)malloc (N * sizeof (unsigned int)); > + > + foo (a, b, c); > + > + free (a); > + free (b); > + free (c); > + > + return 0; > +} > + > +/* Check that the loop has been split off into a function. */ > +/* { dg-final { scan-tree-dump-times "(?n);; Function .*foo._omp_fn.0" 1 "optimized" } } */ For C we get: ;; Function foo._omp_fn.0 (foo._omp_fn.0, funcdef_no=12, decl_uid=2534, cgraph_uid=14, symbol_order=14) ..., so that matches, but for C++ we get: ;; Function foo(unsigned int*, unsigned int*, unsigned int*) [clone ._omp_fn.0] (_ZL3fooPjS_S_._omp_fn.0, funcdef_no=12, decl_uid=2416, cgraph_uid=14, symbol_order=14) ..., which doesn't match, so this directive FAILs. > + > +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 3 "alias" } } */ > +/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "alias" } } */ > +/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "alias" } } */ > +/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "alias" } } */ > +/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "alias" } } */ > +/* { dg-final { scan-tree-dump-times "clique 1 base 6" 1 "alias" } } */ > +/* { dg-final { scan-tree-dump-times "clique 1 base 7" 1 "alias" } } */ > +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 9 "alias" } } */ > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-offload-alias-ptr.c > @@ -0,0 +1,44 @@ > +/* { dg-additional-options "-O2" } */ > +/* { dg-additional-options "-fdump-tree-optimized" } */ > +/* { dg-additional-options "-fdump-tree-alias-all" } */ > +/* { dg-additional-options "-foffload-alias=pointer" } */ > + > +#include <stdlib.h> > + > +#define N (1024 * 512) > +#define COUNTERTYPE unsigned int > + > +unsigned int a[N]; > +unsigned int b[N]; > +unsigned int c[N]; > + > +int > +main (void) > +{ > + for (COUNTERTYPE i = 0; i < N; i++) > + a[i] = i * 2; > + > + for (COUNTERTYPE i = 0; i < N; i++) > + b[i] = i * 4; > + > +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) > + { > + for (COUNTERTYPE ii = 0; ii < N; ii++) > + c[ii] = a[ii] + b[ii]; > + } > + > + for (COUNTERTYPE i = 0; i < N; i++) > + if (c[i] != a[i] + b[i]) > + abort (); > + > + return 0; > +} > + > +/* Check that the loop has been split off into a function. */ > +/* { dg-final { scan-tree-dump-times "(?n);; Function .*main._omp_fn.0" 1 "optimized" } } */ This works for both C and C++. > + > +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 3 "alias" } } */ > +/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "alias" } } */ > +/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "alias" } } */ > +/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "alias" } } */ > +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 6 "alias" } } */ Grüße Thomas
Add goacc/kernels-loop-offload-alias-{none,ptr}.c 2015-11-03 Tom de Vries <tom@codesourcery.com> * c-c++-common/goacc/kernels-loop-offload-alias-none.c: New test. * c-c++-common/goacc/kernels-loop-offload-alias-ptr.c: New test. --- .../goacc/kernels-loop-offload-alias-none.c | 61 ++++++++++++++++++++++ .../goacc/kernels-loop-offload-alias-ptr.c | 44 ++++++++++++++++ 2 files changed, 105 insertions(+) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-offload-alias-none.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-offload-alias-ptr.c diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-offload-alias-none.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-offload-alias-none.c new file mode 100644 index 0000000..bb96330 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-offload-alias-none.c @@ -0,0 +1,61 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-optimized" } */ +/* { dg-additional-options "-fdump-tree-alias-all" } */ +/* { dg-additional-options "-foffload-alias=none" } */ + +#include <stdlib.h> + +#define N (1024 * 512) +#define COUNTERTYPE unsigned int + +static void +foo (unsigned int *a, unsigned int *b, unsigned int *c) +{ + for (COUNTERTYPE i = 0; i < N; i++) + a[i] = i * 2; + + for (COUNTERTYPE i = 0; i < N; i++) + b[i] = i * 4; + +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) + { + for (COUNTERTYPE ii = 0; ii < N; ii++) + c[ii] = a[ii] + b[ii]; + } + + for (COUNTERTYPE i = 0; i < N; i++) + if (c[i] != a[i] + b[i]) + abort (); +} + +int +main (void) +{ + unsigned int *a; + unsigned int *b; + unsigned int *c; + + a = (unsigned int *)malloc (N * sizeof (unsigned int)); + b = (unsigned int *)malloc (N * sizeof (unsigned int)); + c = (unsigned int *)malloc (N * sizeof (unsigned int)); + + foo (a, b, c); + + free (a); + free (b); + free (c); + + return 0; +} + +/* Check that the loop has been split off into a function. */ +/* { dg-final { scan-tree-dump-times "(?n);; Function .*foo._omp_fn.0" 1 "optimized" } } */ + +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 3 "alias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "alias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "alias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "alias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "alias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 6" 1 "alias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 7" 1 "alias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 9 "alias" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-offload-alias-ptr.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-offload-alias-ptr.c new file mode 100644 index 0000000..de4f45a --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-offload-alias-ptr.c @@ -0,0 +1,44 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-optimized" } */ +/* { dg-additional-options "-fdump-tree-alias-all" } */ +/* { dg-additional-options "-foffload-alias=pointer" } */ + +#include <stdlib.h> + +#define N (1024 * 512) +#define COUNTERTYPE unsigned int + +unsigned int a[N]; +unsigned int b[N]; +unsigned int c[N]; + +int +main (void) +{ + for (COUNTERTYPE i = 0; i < N; i++) + a[i] = i * 2; + + for (COUNTERTYPE i = 0; i < N; i++) + b[i] = i * 4; + +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) + { + for (COUNTERTYPE ii = 0; ii < N; ii++) + c[ii] = a[ii] + b[ii]; + } + + for (COUNTERTYPE i = 0; i < N; i++) + if (c[i] != a[i] + b[i]) + abort (); + + return 0; +} + +/* Check that the loop has been split off into a function. */ +/* { dg-final { scan-tree-dump-times "(?n);; Function .*main._omp_fn.0" 1 "optimized" } } */ + +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 3 "alias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "alias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "alias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "alias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 6 "alias" } } */ -- 1.9.1