Hi!
On Wed, 4 Oct 2017 10:01:10 +0200, Jakub Jelinek <ja...@redhat.com> wrote:
> This patch propagates attributes, including opt and target nodes, from
> the containing function to the OMP/OACC outlined region functions.
>
> Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk.
Hmm, in your testing, why didn't you observe this breaking a few OpenACC
test cases? Committed to trunk r253402, as obvious:
commit da7a1a683a84d32e9d7be6a5d00925fdfa125430
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Oct 4 11:13:24 2017 +0000
Adjust test cases for attributes propagation changes for OMP outlined
regions
PR tree-optimization/82374
* c-c++-common/goacc/kernels-double-reduction-n.c: Adjust for
attributes propagation changes for OMP outlined regions.
* c-c++-common/goacc/kernels-double-reduction.c: Likewise.
* c-c++-common/goacc/kernels-reduction.c: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@253402
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/testsuite/ChangeLog | 8 ++++++++
gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c | 2 +-
gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c | 2 +-
gcc/testsuite/c-c++-common/goacc/kernels-reduction.c | 2 +-
4 files changed, 11 insertions(+), 3 deletions(-)
diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog
index 2c96f3a..4ac50d0 100644
--- gcc/testsuite/ChangeLog
+++ gcc/testsuite/ChangeLog
@@ -1,3 +1,11 @@
+2017-10-04 Thomas Schwinge <tho...@codesourcery.com>
+
+ PR tree-optimization/82374
+ * c-c++-common/goacc/kernels-double-reduction-n.c: Adjust for
+ attributes propagation changes for OMP outlined regions.
+ * c-c++-common/goacc/kernels-double-reduction.c: Likewise.
+ * c-c++-common/goacc/kernels-reduction.c: Likewise.
+
2017-10-04 Richard Sandiford <richard.sandif...@linaro.org>
PR tree-optimization/82413
diff --git gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c
gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c
index 27ea2e9..10b364b 100644
--- gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c
+++ gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c
@@ -27,7 +27,7 @@ foo (unsigned int n)
/* Check that only one loop is analyzed, and that it can be parallelized. */
/* { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1
"parloops1" } } */
-/* { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc kernels
parallelized, oacc function \\(, , \\), oacc kernels, omp target
entrypoint\\)\\)" 1 "parloops1" } } */
+/* { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc kernels
parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint,
noclone, noinline\\)\\)" 1 "parloops1" } } */
/* { dg-final { scan-tree-dump-not "FAILED:" "parloops1" } } */
/* { dg-final { scan-tree-dump-times "parallelizing outer loop" 1 "parloops1"
} } */
diff --git gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c
gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c
index 0841e90..c026346 100644
--- gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c
+++ gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c
@@ -27,7 +27,7 @@ foo (void)
/* Check that only one loop is analyzed, and that it can be parallelized. */
/* { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1
"parloops1" } } */
-/* { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc kernels
parallelized, oacc function \\(, , \\), oacc kernels, omp target
entrypoint\\)\\)" 1 "parloops1" } } */
+/* { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc kernels
parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint,
noclone, noinline\\)\\)" 1 "parloops1" } } */
/* { dg-final { scan-tree-dump-not "FAILED:" "parloops1" } } */
/* { dg-final { scan-tree-dump-times "parallelizing outer loop" 1 "parloops1"
} } */
diff --git gcc/testsuite/c-c++-common/goacc/kernels-reduction.c
gcc/testsuite/c-c++-common/goacc/kernels-reduction.c
index 4a18272..5921b88 100644
--- gcc/testsuite/c-c++-common/goacc/kernels-reduction.c
+++ gcc/testsuite/c-c++-common/goacc/kernels-reduction.c
@@ -26,7 +26,7 @@ foo (void)
/* Check that only one loop is analyzed, and that it can be parallelized. */
/* { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1
"parloops1" } } */
-/* { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc kernels
parallelized, oacc function \\(, , \\), oacc kernels, omp target
entrypoint\\)\\)" 1 "parloops1" } } */
+/* { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc kernels
parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint,
noclone, noinline\\)\\)" 1 "parloops1" } } */
/* { dg-final { scan-tree-dump-not "FAILED:" "parloops1" } } */
/* Check that the loop has been split off into a function. */
Grüße
Thomas
> 2017-10-04 Jakub Jelinek <ja...@redhat.com>
>
> PR tree-optimization/82374
> * omp-low.c (create_omp_child_function): Copy DECL_ATTRIBUTES,
> DECL_FUNCTION_SPECIFIC_OPTIMIZATION,
> DECL_FUNCTION_SPECIFIC_TARGET and DECL_FUNCTION_VERSIONED from
> current_function_decl to the new decl.
>
> * gcc.dg/gomp/pr82374.c: New test.
>
> --- gcc/omp-low.c.jj 2017-09-05 23:32:02.000000000 +0200
> +++ gcc/omp-low.c 2017-10-03 12:25:13.956261522 +0200
> @@ -1626,6 +1626,14 @@ create_omp_child_function (omp_context *
> DECL_CONTEXT (decl) = NULL_TREE;
> DECL_INITIAL (decl) = make_node (BLOCK);
> BLOCK_SUPERCONTEXT (DECL_INITIAL (decl)) = decl;
> + DECL_ATTRIBUTES (decl) = DECL_ATTRIBUTES (current_function_decl);
> + DECL_FUNCTION_SPECIFIC_OPTIMIZATION (decl)
> + = DECL_FUNCTION_SPECIFIC_OPTIMIZATION (current_function_decl);
> + DECL_FUNCTION_SPECIFIC_TARGET (decl)
> + = DECL_FUNCTION_SPECIFIC_TARGET (current_function_decl);
> + DECL_FUNCTION_VERSIONED (decl)
> + = DECL_FUNCTION_VERSIONED (current_function_decl);
> +
> if (omp_maybe_offloaded_ctx (ctx))
> {
> cgraph_node::get_create (decl)->offloadable = 1;
> --- gcc/testsuite/gcc.dg/gomp/pr82374.c.jj 2017-10-03 13:09:15.515879118
> +0200
> +++ gcc/testsuite/gcc.dg/gomp/pr82374.c 2017-10-03 13:08:56.000000000
> +0200
> @@ -0,0 +1,31 @@
> +/* PR tree-optimization/82374 */
> +/* { dg-do compile } */
> +/* { dg-options "-O2 -fno-tree-vectorize -fdump-tree-vect-details" } */
> +/* { dg-additional-options "-mavx -mno-avx2" { target i?86-*-* x86_64-*-* }
> } */
> +/* { dg-additional-options "-mvsx" { target powerpc_vsx_ok } } */
> +
> +#define SIZE (1024 * 1024 * 1024)
> +
> +float a[SIZE];
> +float b[SIZE];
> +float c[SIZE];
> +float d[SIZE];
> +
> +__attribute__((optimize ("O2", "tree-vectorize"))) void
> +foo (void)
> +{
> + int i;
> +#pragma omp parallel for
> + for (i = 0; i < SIZE; i++)
> + c[i] = a[i] + b[i];
> +}
> +
> +__attribute__((optimize ("O2", "tree-vectorize"))) void
> +bar (void)
> +{
> + int i;
> + for (i = 0; i < SIZE; i++)
> + d[i] = a[i] + b[i];
> +}
> +
> +/* { dg-final { scan-tree-dump-times "vectorized 1 loops" 2 "vect" { target
> { { i?86-*-* x86_64-*-* } || { powerpc_vsx_ok } } } } } */
>
> Jakub
