* [gomp4, WIP] Implement -foffload-alias
@ 2015-09-28 16:15 Tom de Vries
2015-11-03 14:19 ` [gomp4, committed] " Tom de Vries
0 siblings, 1 reply; 5+ messages in thread
From: Tom de Vries @ 2015-09-28 16:15 UTC (permalink / raw)
To: Nathan Sidwell; +Cc: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 942 bytes --]
Hi,
this work-in-progress patch implements a new option
-foffload-alias=<none|pointer|all>.
The option -foffload-alias=none instructs the compiler to assume that
objects references and pointer dereferences in an offload region do not
alias.
The option -foffload-alias=pointer instructs the compiler to assume that
objects references in an offload region do not alias.
The option -foffload-alias=all instructs the compiler to make no
assumptions about aliasing in offload regions.
The default value is -foffload-alias=none.
The patch works by adding restrict to the types of the fields used to
pass data to an offloading region.
Atm, the kernels-loop-offload-alias-ptr.c test-case passes, but the
kernels-loop-offload-alias-none.c test-case fails. For the latter, the
required amount of restrict is added, but it has no effect. I've
reported this in a more basic form in PR67742: "3rd-level restrict ignored".
Thanks,
- Tom
[-- Attachment #2: 0001-Implement-foffload-alias.patch --]
[-- Type: text/x-patch, Size: 10058 bytes --]
Implement -foffload-alias
2015-09-28 Tom de Vries <tom@codesourcery.com>
* common.opt (foffload-alias): New option.
* flag-types.h (enum offload_alias): New enum.
* omp-low.c (is_gimple_oacc_offload): New function.
(install_var_field): Handle flag_offload_alias.
* doc/invoke.texi (@item Code Generation Options): Add -foffload-alias.
(@item -foffload-alias): New item.
* c-c++-common/goacc/kernels-loop-offload-alias-none.c: New test.
* c-c++-common/goacc/kernels-loop-offload-alias-ptr.c: New test.
---
gcc/common.opt | 16 ++++++
gcc/doc/invoke.texi | 11 ++++
gcc/flag-types.h | 7 +++
gcc/omp-low.c | 24 ++++++++-
.../goacc/kernels-loop-offload-alias-none.c | 63 ++++++++++++++++++++++
.../goacc/kernels-loop-offload-alias-ptr.c | 52 ++++++++++++++++++
6 files changed, 172 insertions(+), 1 deletion(-)
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/common.opt b/gcc/common.opt
index 290b6b3..28977a4 100644
--- a/gcc/common.opt
+++ b/gcc/common.opt
@@ -1730,6 +1730,22 @@ Enum(offload_abi) String(ilp32) Value(OFFLOAD_ABI_ILP32)
EnumValue
Enum(offload_abi) String(lp64) Value(OFFLOAD_ABI_LP64)
+foffload-alias=
+Common Joined RejectNegative Enum(offload_alias) Var(flag_offload_alias) Init(OFFLOAD_ALIAS_NONE)
+-foffload-alias=[all|pointer|none] Assume non-aliasing in an offload region
+
+Enum
+Name(offload_alias) Type(enum offload_alias) UnknownError(unknown offload aliasing %qs)
+
+EnumValue
+Enum(offload_alias) String(all) Value(OFFLOAD_ALIAS_ALL)
+
+EnumValue
+Enum(offload_alias) String(pointer) Value(OFFLOAD_ALIAS_POINTER)
+
+EnumValue
+Enum(offload_alias) String(none) Value(OFFLOAD_ALIAS_NONE)
+
fomit-frame-pointer
Common Report Var(flag_omit_frame_pointer) Optimization
When possible do not generate stack frames
diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi
index 909a453..a5ab785 100644
--- a/gcc/doc/invoke.texi
+++ b/gcc/doc/invoke.texi
@@ -1136,6 +1136,7 @@ See S/390 and zSeries Options.
-finstrument-functions-exclude-function-list=@var{sym},@var{sym},@dots{} @gol
-finstrument-functions-exclude-file-list=@var{file},@var{file},@dots{} @gol
-fno-common -fno-ident @gol
+-foffload-alias=@r{[}none@r{|}pointer@r{|}all@r{]} @gol
-fpcc-struct-return -fpic -fPIC -fpie -fPIE -fno-plt @gol
-fno-jump-tables @gol
-frecord-gcc-switches @gol
@@ -23695,6 +23696,16 @@ The options @option{-ftrapv} and @option{-fwrapv} override each other, so using
using @option{-ftrapv} @option{-fwrapv} @option{-fno-wrapv} on the command-line
results in @option{-ftrapv} being effective.
+@item -foffload-alias=@r{[}none@r{|}pointer@r{|}all@r{]}
+@opindex -foffload-alias
+The option @option{-foffload-alias=none} instructs the compiler to assume that
+objects references and pointer dereferences in an offload region do not alias.
+The option @option{-foffload-alias=pointer} instruct the compiler to assume that
+objects references in an offload region are presumed unaliased. The option
+@option{-foffload-alias=all} instructs the compiler to make no assumtions about
+aliasing in offload regions. The default value is
+@option{-foffload-alias=none}.
+
@item -fexceptions
@opindex fexceptions
Enable exception handling. Generates extra code needed to propagate
diff --git a/gcc/flag-types.h b/gcc/flag-types.h
index ac9ca0b..e8e672d 100644
--- a/gcc/flag-types.h
+++ b/gcc/flag-types.h
@@ -286,5 +286,12 @@ enum gfc_convert
GFC_FLAG_CONVERT_LITTLE
};
+enum offload_alias
+{
+ OFFLOAD_ALIAS_ALL,
+ OFFLOAD_ALIAS_POINTER,
+ OFFLOAD_ALIAS_NONE
+};
+
#endif /* ! GCC_FLAG_TYPES_H */
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index a5904eb..9cbba1f 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1140,6 +1140,16 @@ use_pointer_for_field (tree decl, omp_context *shared_ctx)
return false;
}
+static bool
+is_gimple_oacc_offload (const gimple *stmt)
+{
+ return (gimple_code (stmt) == GIMPLE_OMP_TARGET
+ && (((gimple_omp_target_kind (stmt)
+ == GF_OMP_TARGET_KIND_OACC_PARALLEL)
+ || (gimple_omp_target_kind (stmt)
+ == GF_OMP_TARGET_KIND_OACC_KERNELS))));
+}
+
/* Construct a new automatic decl similar to VAR. */
static tree
@@ -1283,6 +1293,7 @@ static void
install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
{
tree field, type, sfield = NULL_TREE;
+ bool in_oacc_offload = is_gimple_oacc_offload (ctx->stmt);
gcc_assert ((mask & 1) == 0
|| !splay_tree_lookup (ctx->field_map, (splay_tree_key) var));
@@ -1298,7 +1309,18 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
type = build_pointer_type (build_pointer_type (type));
}
else if (by_ref)
- type = build_pointer_type (type);
+ {
+ if (in_oacc_offload
+ && flag_offload_alias == OFFLOAD_ALIAS_NONE
+ && POINTER_TYPE_P (type))
+ type = build_qualified_type (type, TYPE_QUAL_RESTRICT);
+
+ type = build_pointer_type (type);
+
+ if (in_oacc_offload
+ && flag_offload_alias != OFFLOAD_ALIAS_ALL)
+ type = build_qualified_type (type, TYPE_QUAL_RESTRICT);
+ }
else if ((mask & 3) == 1 && is_reference (var))
type = TREE_TYPE (type);
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..918bced
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-offload-alias-none.c
@@ -0,0 +1,63 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+/* { dg-additional-options "-fdump-tree-parloops_oacc_kernels-all" } */
+/* { 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
+
+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));
+
+ 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 ();
+
+ free (a);
+ free (b);
+ free (c);
+
+ return 0;
+}
+
+/* Check that only one loop is analyzed, and that it can be parallelized. */
+/* { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1 "parloops_oacc_kernels" } } */
+/* { dg-final { scan-tree-dump-not "FAILED:" "parloops_oacc_kernels" } } */
+
+/* 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 "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 1 "parloops_oacc_kernels" } } */
+
+/* { 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..029d77c
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-offload-alias-ptr.c
@@ -0,0 +1,52 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+/* { dg-additional-options "-fdump-tree-parloops_oacc_kernels-all" } */
+/* { 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 only one loop is analyzed, and that it can be parallelized. */
+/* { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1 "parloops_oacc_kernels" } } */
+/* { dg-final { scan-tree-dump-not "FAILED:" "parloops_oacc_kernels" } } */
+
+/* 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 "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 1 "parloops_oacc_kernels" } } */
+
+/* { 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
^ permalink raw reply [flat|nested] 5+ messages in thread
* [gomp4, committed] Implement -foffload-alias
2015-09-28 16:15 [gomp4, WIP] Implement -foffload-alias Tom de Vries
@ 2015-11-03 14:19 ` Tom de Vries
2015-11-03 14:33 ` Tom de Vries
0 siblings, 1 reply; 5+ messages in thread
From: Tom de Vries @ 2015-11-03 14:19 UTC (permalink / raw)
To: Nathan Sidwell; +Cc: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 1477 bytes --]
[ was: Re: [gomp4, WIP] Implement -foffload-alias ]
On 28/09/15 17:38, Tom de Vries wrote:
> Hi,
>
> this work-in-progress patch implements a new option
> -foffload-alias=<none|pointer|all>.
>
>
> The option -foffload-alias=none instructs the compiler to assume that
> objects references and pointer dereferences in an offload region do not
> alias.
>
> The option -foffload-alias=pointer instructs the compiler to assume that
> objects references in an offload region do not alias.
>
> The option -foffload-alias=all instructs the compiler to make no
> assumptions about aliasing in offload regions.
>
> The default value is -foffload-alias=none.
>
>
> The patch works by adding restrict to the types of the fields used to
> pass data to an offloading region.
>
Updated patch attached, committed to gomp-4_0-branch.
> Atm, the kernels-loop-offload-alias-ptr.c test-case passes, but the
> kernels-loop-offload-alias-none.c test-case fails.
I've dropped the two testcases from this patch, I'll commit in a
follow-up patch.
> For the latter, the
> required amount of restrict is added, but it has no effect. I've
> reported this in a more basic form in PR67742: "3rd-level restrict
> ignored".
I've committed a fix for that PR as reported here:
https://gcc.gnu.org/ml/gcc-patches/2015-11/msg00204.html .
Furthermore, I've added support for the option in the 'mask & 4' case in
install_var_field, I ran into this when trying out some Fortran test-cases.
Thanks,
- Tom
[-- Attachment #2: 0002-Implement-foffload-alias.patch --]
[-- Type: text/x-patch, Size: 4857 bytes --]
Implement -foffload-alias
2015-09-28 Tom de Vries <tom@codesourcery.com>
* common.opt (foffload-alias): New option.
* flag-types.h (enum offload_alias): New enum.
* omp-low.c (install_var_field): Handle flag_offload_alias.
* doc/invoke.texi (@item Code Generation Options): Add -foffload-alias.
(@item -foffload-alias): New item.
---
gcc/common.opt | 16 ++++++++++++++++
gcc/doc/invoke.texi | 11 +++++++++++
gcc/flag-types.h | 7 +++++++
gcc/omp-low.c | 28 ++++++++++++++++++++++++++--
4 files changed, 60 insertions(+), 2 deletions(-)
diff --git a/gcc/common.opt b/gcc/common.opt
index c85ab49..135e777 100644
--- a/gcc/common.opt
+++ b/gcc/common.opt
@@ -1738,6 +1738,22 @@ Enum(offload_abi) String(ilp32) Value(OFFLOAD_ABI_ILP32)
EnumValue
Enum(offload_abi) String(lp64) Value(OFFLOAD_ABI_LP64)
+foffload-alias=
+Common Joined RejectNegative Enum(offload_alias) Var(flag_offload_alias) Init(OFFLOAD_ALIAS_NONE)
+-foffload-alias=[all|pointer|none] Assume non-aliasing in an offload region
+
+Enum
+Name(offload_alias) Type(enum offload_alias) UnknownError(unknown offload aliasing %qs)
+
+EnumValue
+Enum(offload_alias) String(all) Value(OFFLOAD_ALIAS_ALL)
+
+EnumValue
+Enum(offload_alias) String(pointer) Value(OFFLOAD_ALIAS_POINTER)
+
+EnumValue
+Enum(offload_alias) String(none) Value(OFFLOAD_ALIAS_NONE)
+
fomit-frame-pointer
Common Report Var(flag_omit_frame_pointer) Optimization
When possible do not generate stack frames.
diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi
index 5a07512..8967f88 100644
--- a/gcc/doc/invoke.texi
+++ b/gcc/doc/invoke.texi
@@ -1142,6 +1142,7 @@ See S/390 and zSeries Options.
-finstrument-functions-exclude-function-list=@var{sym},@var{sym},@dots{} @gol
-finstrument-functions-exclude-file-list=@var{file},@var{file},@dots{} @gol
-fno-common -fno-ident @gol
+-foffload-alias=@r{[}none@r{|}pointer@r{|}all@r{]} @gol
-fpcc-struct-return -fpic -fPIC -fpie -fPIE -fno-plt @gol
-fno-jump-tables @gol
-frecord-gcc-switches @gol
@@ -23842,6 +23843,16 @@ The options @option{-ftrapv} and @option{-fwrapv} override each other, so using
using @option{-ftrapv} @option{-fwrapv} @option{-fno-wrapv} on the command-line
results in @option{-ftrapv} being effective.
+@item -foffload-alias=@r{[}none@r{|}pointer@r{|}all@r{]}
+@opindex -foffload-alias
+The option @option{-foffload-alias=none} instructs the compiler to assume that
+objects references and pointer dereferences in an offload region do not alias.
+The option @option{-foffload-alias=pointer} instruct the compiler to assume that
+objects references in an offload region do not alias. The option
+@option{-foffload-alias=all} instructs the compiler to make no assumptions about
+aliasing in offload regions. The default value is
+@option{-foffload-alias=none}.
+
@item -fexceptions
@opindex fexceptions
Enable exception handling. Generates extra code needed to propagate
diff --git a/gcc/flag-types.h b/gcc/flag-types.h
index 6301cea..87b1677 100644
--- a/gcc/flag-types.h
+++ b/gcc/flag-types.h
@@ -293,5 +293,12 @@ enum gfc_convert
GFC_FLAG_CONVERT_LITTLE
};
+enum offload_alias
+{
+ OFFLOAD_ALIAS_ALL,
+ OFFLOAD_ALIAS_POINTER,
+ OFFLOAD_ALIAS_NONE
+};
+
#endif /* ! GCC_FLAG_TYPES_H */
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 3543785..6bac074 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1441,6 +1441,14 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
tree field, type, sfield = NULL_TREE;
splay_tree_key key = (splay_tree_key) var;
+ /* We use flag_offload_alias only for the oacc kernels region for the
+ moment. */
+ bool offload_alias_p = is_oacc_kernels (ctx);
+ bool no_alias_var_p
+ = offload_alias_p && flag_offload_alias != OFFLOAD_ALIAS_ALL;
+ bool no_alias_ptr_p
+ = offload_alias_p && flag_offload_alias == OFFLOAD_ALIAS_NONE;
+
if ((mask & 8) != 0)
{
key = (splay_tree_key) &DECL_UID (var);
@@ -1457,10 +1465,26 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
if (mask & 4)
{
gcc_assert (TREE_CODE (type) == ARRAY_TYPE);
- type = build_pointer_type (build_pointer_type (type));
+
+ type = build_pointer_type (type);
+ if (no_alias_var_p)
+ type = build_qualified_type (type, TYPE_QUAL_RESTRICT);
+
+ type = build_pointer_type (type);
+ if (no_alias_var_p)
+ type = build_qualified_type (type, TYPE_QUAL_RESTRICT);
}
else if (by_ref)
- type = build_pointer_type (type);
+ {
+ if (no_alias_ptr_p
+ && POINTER_TYPE_P (type))
+ type = build_qualified_type (type, TYPE_QUAL_RESTRICT);
+
+ type = build_pointer_type (type);
+
+ if (no_alias_var_p)
+ type = build_qualified_type (type, TYPE_QUAL_RESTRICT);
+ }
else if ((mask & 3) == 1 && is_reference (var))
type = TREE_TYPE (type);
--
1.9.1
^ permalink raw reply [flat|nested] 5+ messages in thread
* Re: [gomp4, committed] Implement -foffload-alias
2015-11-03 14:19 ` [gomp4, committed] " Tom de Vries
@ 2015-11-03 14:33 ` Tom de Vries
2015-11-04 8:47 ` Thomas Schwinge
0 siblings, 1 reply; 5+ messages in thread
From: Tom de Vries @ 2015-11-03 14:33 UTC (permalink / raw)
To: Nathan Sidwell; +Cc: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 184 bytes --]
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
[-- Attachment #2: 0001-Add-goacc-kernels-loop-offload-alias-none-ptr-.c.patch --]
[-- Type: text/x-patch, Size: 4305 bytes --]
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
^ permalink raw reply [flat|nested] 5+ messages in thread
* Re: [gomp4, committed] Implement -foffload-alias
2015-11-03 14:33 ` Tom de Vries
@ 2015-11-04 8:47 ` Thomas Schwinge
2015-11-04 14:55 ` Tom de Vries
0 siblings, 1 reply; 5+ messages in thread
From: Thomas Schwinge @ 2015-11-04 8:47 UTC (permalink / raw)
To: Tom de Vries; +Cc: gcc-patches, Nathan Sidwell
[-- Attachment #1: Type: text/plain, Size: 4383 bytes --]
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
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 5+ messages in thread
* Re: [gomp4, committed] Implement -foffload-alias
2015-11-04 8:47 ` Thomas Schwinge
@ 2015-11-04 14:55 ` Tom de Vries
0 siblings, 0 replies; 5+ messages in thread
From: Tom de Vries @ 2015-11-04 14:55 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches, Nathan Sidwell
[-- Attachment #1: Type: text/plain, Size: 715 bytes --]
On 04/11/15 09:47, Thomas Schwinge wrote:
>> +/* 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.
>
Hi Thomas,
thanks for noticing.
Fixed as attached.
Committed to gomp-4_0-branch.
Thanks,
- Tom
[-- Attachment #2: 0001-Fixup-goacc-kernels-loop-offload-alias-none.c.patch --]
[-- Type: text/x-patch, Size: 1099 bytes --]
Fixup goacc/kernels-loop-offload-alias-none.c
2015-11-04 Tom de Vries <tom@codesourcery.com>
* c-c++-common/goacc/kernels-loop-offload-alias-none.c: Fix
foo._omp_fn.0 function name scanning.
---
gcc/testsuite/c-c++-common/goacc/kernels-loop-offload-alias-none.c | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
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
index bb96330..79d8daa 100644
--- 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
@@ -49,7 +49,7 @@ main (void)
}
/* 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 "(?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" } } */
--
1.9.1
^ permalink raw reply [flat|nested] 5+ messages in thread
end of thread, other threads:[~2015-11-04 14:55 UTC | newest]
Thread overview: 5+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-09-28 16:15 [gomp4, WIP] Implement -foffload-alias Tom de Vries
2015-11-03 14:19 ` [gomp4, committed] " Tom de Vries
2015-11-03 14:33 ` Tom de Vries
2015-11-04 8:47 ` Thomas Schwinge
2015-11-04 14:55 ` Tom de Vries
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).