On 22-04-15 09:42, Richard Biener wrote:
This patch adds pass_loop_ccp to pass group pass_oacc_kernels.
> > >
> > >We need this pass to simplify the loop body, and allow pass_parloops to
detect
> > >that loop iterations are independent.
> > >
> >
> >As suggested here (https://gcc.gnu.org/ml/gcc-patches/2014-11/msg02993.html
)
> >I've replaced the pass_ccp with pass_copyprop, which performs trivial
constant
> >propagation in addition to copy propagation.
> >
> >Bootstrapped and reg-tested as before.
> >
> >OK for trunk?
I've recently wondered why we do copy propagation after LIM and I don't
remember. Can you remind me? Can you add testcases that fail before
this kind of patches and pass afterwards?
For attached test-case, we manage to parallelize with pass_copy_prop (but then
run into an ICE):
...
PASS: c-c++-common/goacc/kernels-loop-reduction.c scan-tree-dump-not
parloops_oacc_kernels "FAILED:"
PASS: c-c++-common/goacc/kernels-loop-reduction.c scan-tree-dump-times
parloops_oacc_kernels "SUCCESS: may be parallelized" 1
FAIL: c-c++-common/goacc/kernels-loop-reduction.c (internal compiler error)
FAIL: c-c++-common/goacc/kernels-loop-reduction.c (test for excess errors)
...
Without pass_copy_prop we don't manage to parallelize:
...
FAIL: c-c++-common/goacc/kernels-loop-reduction.c scan-tree-dump-not
parloops_oacc_kernels "FAILED:"
FAIL: c-c++-common/goacc/kernels-loop-reduction.c scan-tree-dump-times
parloops_oacc_kernels "SUCCESS: may be parallelized" 1
PASS: c-c++-common/goacc/kernels-loop-reduction.c (test for excess errors)
...
In more detail, before pass_copy_prop, we have:
...
<bb 7>:
# D__lsm.14_3 = PHI <D__lsm.14_9(15), D__lsm.14_21(6)>
...
sum.3_39 = D__lsm.14_3;
sum.4_40 = _37 + sum.3_39;
D__lsm.14_9 = sum.4_40;
...
if (ii_43 <= 524287)
goto <bb 15>;
else
goto <bb 8>;
<bb 15>:
goto <bb 7>;
...
And after pass_copy_prop, we have:
...
<bb 7>:
# D__lsm.14_3 = PHI <sum.4_40(8), D__lsm.14_21(6)>
...
sum.4_40 = D__lsm.14_3 + _37;
...
if (ii_43 <= 524287)
goto <bb 8>;
else
goto <bb 9>;
<bb 8>:
goto <bb 7>;
...
The testcase is not committed yet, because reductions are not handled yet (which
explains the ICE).
Thanks,
- Tom
/* { dg-additional-options "-O2" } */
/* { dg-additional-options "-ftree-parallelize-loops=32" } */
/* { dg-additional-options "-fdump-tree-parloops_oacc_kernels-all" } */
#include <stdlib.h>
#define N (1024 * 512)
#define COUNTERTYPE unsigned int
int
main (void)
{
unsigned int *__restrict a;
unsigned int sum = 0;
unsigned int sum2 = 0;
a = (unsigned int *)malloc (N * sizeof (unsigned int));
for (COUNTERTYPE i = 0; i < N; i++)
a[i] = i * 2;
#pragma acc kernels copy (sum) copyin (a[0:N])
{
for (COUNTERTYPE ii = 0; ii < N; ii++)
sum += a[ii];
}
for (COUNTERTYPE i = 0; i < N; i++)
sum2 += a[i];
if (sum != sum2)
abort ();
free (a);
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" } } */
/* { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } } */
