Hi,
I ran into a problem with fortran loops in oacc kernels regions not
being parallelized, after introducting transform_to_exit_first_loop_alt.
For gfortran.dg/goacc/kernels-loop.f95, we get:
...
#pragma omp target oacc_parallel num_gangs(1)
...
instead of the desired num_gangs (32).
transform_to_exit_first_loop_alt fails because nit is _135, where nit is
defined by:
...
*_105 = 0;
D__lsm.27_50 = *_105;
_32 = (unsigned int) D__lsm.27_50;
_135 = 1023 - _32;
...
pass_fre would manage to propagate the '*105 = 0' assignment. But in the
current pass order, pass_fre is run before pass_lim, where this pattern
is introduced:
...
NEXT_PASS (pass_ch_oacc_kernels);
NEXT_PASS (pass_fre);
NEXT_PASS (pass_tree_loop_init);
NEXT_PASS (pass_lim);
NEXT_PASS (pass_copy_prop);
NEXT_PASS (pass_scev_cprop);
NEXT_PASS (pass_parallelize_loops_oacc_kernels);
NEXT_PASS (pass_expand_omp_ssa);
NEXT_PASS (pass_tree_loop_done);
...
The patch moves pass_fre to the location of pass_copy_prop, and replaces
it. Furthermore, it adds scans to the fortran test-cases to make sure
they get properly parallelized.
Bootstrapped and reg-tested on x86_64.
Committed to gomp-4_0-branch.
Thanks,
- Tom
Fix parallelization for fortran oacc kernels tests
2015-06-18 Tom de Vries <t...@codesourcery.com>
* passes.def: Move pass_fre later in pass group pass_oacc_kernels,
replacing pass_copy_prop.
* gfortran.dg/goacc/kernels-loop-2.f95: Add check for num_gangs (32).
* gfortran.dg/goacc/kernels-loop-data-2.f95: Same.
* gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95: Same.
* gfortran.dg/goacc/kernels-loop-data-enter-exit.f95: Same.
* gfortran.dg/goacc/kernels-loop-data-update.f95: Same.
* gfortran.dg/goacc/kernels-loop-data.f95: Same.
* gfortran.dg/goacc/kernels-loop.f95: Same.
* gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95: Same.
---
gcc/passes.def | 5 +++--
gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95 | 2 ++
gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95 | 2 ++
gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95 | 2 ++
gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95 | 2 ++
gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95 | 2 ++
gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95 | 2 ++
gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95 | 2 ++
.../gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95 | 2 ++
9 files changed, 19 insertions(+), 2 deletions(-)
diff --git a/gcc/passes.def b/gcc/passes.def
index da497ed..8b00f17 100644
--- a/gcc/passes.def
+++ b/gcc/passes.def
@@ -93,10 +93,11 @@ along with GCC; see the file COPYING3. If not see
NEXT_PASS (pass_oacc_kernels);
PUSH_INSERT_PASSES_WITHIN (pass_oacc_kernels)
NEXT_PASS (pass_ch_oacc_kernels);
- NEXT_PASS (pass_fre);
NEXT_PASS (pass_tree_loop_init);
NEXT_PASS (pass_lim);
- NEXT_PASS (pass_copy_prop);
+ NEXT_PASS (pass_tree_loop_done);
+ NEXT_PASS (pass_fre);
+ NEXT_PASS (pass_tree_loop_init);
NEXT_PASS (pass_scev_cprop);
NEXT_PASS (pass_parallelize_loops_oacc_kernels);
NEXT_PASS (pass_expand_omp_ssa);
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95
index bef69f8..634c670 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95
@@ -42,5 +42,7 @@ end program main
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.1 " 1 "optimized" } }
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.2 " 1 "optimized" } }
+! { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 3 "parloops_oacc_kernels" } }
+
! { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } }
! { dg-final { cleanup-tree-dump "optimized" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95
index 1b75a23..58d92ff 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95
@@ -48,5 +48,7 @@ end program main
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.1 " 1 "optimized" } }
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.2 " 1 "optimized" } }
+! { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 3 "parloops_oacc_kernels" } }
+
! { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } }
! { dg-final { cleanup-tree-dump "optimized" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95
index 4ba83b6..689002d 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95
@@ -48,5 +48,7 @@ end program main
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.1 " 1 "optimized" } }
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.2 " 1 "optimized" } }
+! { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 3 "parloops_oacc_kernels" } }
+
! { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } }
! { dg-final { cleanup-tree-dump "optimized" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95
index 2b05b33..0145754 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95
@@ -46,5 +46,7 @@ end program main
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.1 " 1 "optimized" } }
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.2 " 1 "optimized" } }
+! { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 3 "parloops_oacc_kernels" } }
+
! { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } }
! { dg-final { cleanup-tree-dump "optimized" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95
index b3c80dc..2c2da27 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95
@@ -45,5 +45,7 @@ end program main
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.0 " 1 "optimized" } }
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.1 " 1 "optimized" } }
+! { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 2 "parloops_oacc_kernels" } }
+
! { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } }
! { dg-final { cleanup-tree-dump "optimized" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95
index 98c5e7a..f61fe7e 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95
@@ -46,5 +46,7 @@ end program main
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.1 " 1 "optimized" } }
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.2 " 1 "optimized" } }
+! { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 3 "parloops_oacc_kernels" } }
+
! { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } }
! { dg-final { cleanup-tree-dump "optimized" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95
index be5f26d..f9f8bbd 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95
@@ -36,5 +36,7 @@ end program main
! 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 { cleanup-tree-dump "parloops_oacc_kernels" } }
! { dg-final { cleanup-tree-dump "optimized" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95
index 7ea2b49..c6ab436 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95
@@ -47,5 +47,7 @@ end program main
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.1 " 1 "optimized" } }
! { dg-final { scan-tree-dump-times "(?n);; Function MAIN__._omp_fn.2 " 1 "optimized" } }
+! { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 2 "parloops_oacc_kernels" } }
+
! { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } }
! { dg-final { cleanup-tree-dump "optimized" } }
--
1.9.1
