[gomp4,committed] Implement -foffload-alias

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.

Thanks,
- Tom

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