Re: [PATCH, 10/16] Add pass_oacc_kernels pass group in passes.def

[Richard Biener <rguenther at suse dot de>]

On Thu, 19 Nov 2015, Tom de Vries wrote:

> On 17/11/15 15:53, Tom de Vries wrote:
> > > And the above LIM example
> > > is none for why you need two LIM passes...
> > 
> > Indeed. I'm planning a separate reply to explain in more detail the need
> > for the two pass_lims.
> 
> I.
> 
> I managed to get rid of the two pass_lims for the motivating example that I
> used until now (goacc/kernels-double-reduction.c). I found that by adding a
> pass_dominator instance after pass_ch, I could get rid of the second pass_lim
> (and pass_copyprop as well).
> 
> But... then I wrote a counter example (goacc/kernels-double-reduction-n.c),
> and I'm back at two pass_lims (and two pass_dominators).
> Also I've split the pass group into a bit before and after pass_fre.
> 
> So, the current pass group looks like:
> ...
> NEXT_PASS (pass_build_ealias);
> 
> /* Pass group that runs when the function is an offloaded function
>    containing oacc kernels loops.  Part 1.  */
> NEXT_PASS (pass_oacc_kernels);
> PUSH_INSERT_PASSES_WITHIN (pass_oacc_kernels)
>     /* We need pass_ch here, because pass_lim has no effect on
>        exit-first loops (PR65442).  Ideally we want to remove both
>        this pass instantiation, and the reverse transformation
>        transform_to_exit_first_loop_alt, which is done in
>        pass_parallelize_loops_oacc_kernels. */
>     NEXT_PASS (pass_ch);
> POP_INSERT_PASSES ()
> 
> NEXT_PASS (pass_fre);
> 
> /* Pass group that runs when the function is an offloaded function
>    containing oacc kernels loops.  Part 2.  */
> NEXT_PASS (pass_oacc_kernels2);
> PUSH_INSERT_PASSES_WITHIN (pass_oacc_kernels2)
>     /* We use pass_lim to rewrite in-memory iteration and reduction
>        variable accesses in loops into local variables accesses.  */
>     NEXT_PASS (pass_lim);
>     NEXT_PASS (pass_dominator, false /* may_peel_loop_headers_p */);
>     NEXT_PASS (pass_lim);
>     NEXT_PASS (pass_dominator, false /* may_peel_loop_headers_p */);
>     NEXT_PASS (pass_dce);
>     NEXT_PASS (pass_parallelize_loops_oacc_kernels);
>     NEXT_PASS (pass_expand_omp_ssa);
> POP_INSERT_PASSES ()
> NEXT_PASS (pass_merge_phi);
> ...
> 
> 
> II.
> 
> The motivating test-case kernels-double-reduction-n.c:
> ...
> #include <stdlib.h>
> 
> #define N 500
> 
> unsigned int a[N][N];
> 
> void  __attribute__((noinline,noclone))
> foo (unsigned int n)
> {
>   int i, j;
>   unsigned int sum = 1;
> 
> #pragma acc kernels copyin (a[0:n]) copy (sum)
>   {
>     for (i = 0; i < n; ++i)
>       for (j = 0; j < n; ++j)
>         sum += a[i][j];
>   }
> 
>   if (sum != 5001)
>     abort ();
> }
> ...
> 
> 
> III.
> 
> Before first pass_lim. Note no phis on inner or outer loop header for
> iteration varables or reduction variable:
> ...
>   <bb 2>:
>   _5 = *.omp_data_i_4(D).i;
>   *_5 = 0;
>   _44 = *.omp_data_i_4(D).n;
>   _45 = *_44;
>   if (_45 != 0)
>     goto <bb 4>;
>   else
>     goto <bb 3>;
> 
>   <bb 4>: outer loop header
>   _12 = *.omp_data_i_4(D).j;
>   *_12 = 0;
>   if (_45 != 0)
>     goto <bb 6>;
>   else
>     goto <bb 5>;
> 
>   <bb 6>: inner loop header, latch
>   _19 = *.omp_data_i_4(D).a;
>   _21 = *_5;
>   _23 = *_12;
>   _24 = *_19[_21][_23];
>   _25 = *.omp_data_i_4(D).sum;
>   sum.0_26 = *_25;
>   sum.1_27 = _24 + sum.0_26;
>   *_25 = sum.1_27;
>   _33 = _23 + 1;
>   *_12 = _33;
>   j.2_16 = (unsigned int) _33;
>   if (j.2_16 < _45)
>     goto <bb 6>;
>   else
>     goto <bb 5>;
> 
>   <bb 5>: outer loop latch
>   _36 = *_5;
>   _38 = _36 + 1;
>   *_5 = _38;
>   i.3_9 = (unsigned int) _38;
>   if (i.3_9 < _45)
>     goto <bb 4>;
>   else
>     goto <bb 3>;
> 
>   <bb 3>:
>   return;
> ...
> 
> 
> IV.
> 
> After first pass_lim/pass_dom pair. Note there are phis on the inner loop
> header for the reduction and the iteration variable, but not on the outer loop
> header:
> ...
>   <bb 2>:
>   _5 = *.omp_data_i_4(D).i;
>   *_5 = 0;
>   _44 = *.omp_data_i_4(D).n;
>   _45 = *_44;
>   if (_45 != 0)
>     goto <bb 4>;
>   else
>     goto <bb 3>;
> 
>   <bb 4>:
>   _12 = *.omp_data_i_4(D).j;
>   _19 = *.omp_data_i_4(D).a;
>   D__lsm.10_50 = *_12;
>   D__lsm.11_51 = 0;
>   _25 = *.omp_data_i_4(D).sum;
> 
>   <bb 5>: outer loop header
>   D__lsm.10_20 = 0;
>   D__lsm.11_22 = 1;
>   _21 = *_5;
>   D__lsm.12_28 = *_25;
>   D__lsm.13_30 = 0;
>   goto <bb 7>;
> 
>   <bb 7>: inner loop header, latch
>   # D__lsm.10_47 = PHI <0(5), _33(7)>
>   # D__lsm.12_49 = PHI <D__lsm.12_28(5), sum.1_27(7)>
>   _23 = D__lsm.10_47;
>   _24 = *_19[_21][D__lsm.10_47];
>   sum.0_26 = D__lsm.12_49;
>   sum.1_27 = _24 + D__lsm.12_49;
>   D__lsm.12_31 = sum.1_27;
>   D__lsm.13_32 = 1;
>   _33 = D__lsm.10_47 + 1;
>   D__lsm.10_14 = _33;
>   D__lsm.11_15 = 1;
>   j.2_16 = (unsigned int) _33;
>   if (j.2_16 < _45)
>     goto <bb 7>;
>   else
>     goto <bb 8>;
> 
>   <bb 8>: outer loop latch
>   # D__lsm.10_35 = PHI <_33(7)>
>   # D__lsm.11_37 = PHI <1(7)>
>   # D__lsm.12_7 = PHI <sum.1_27(7)>
>   # D__lsm.13_8 = PHI <1(7)>
>   *_25 = sum.1_27;
>   _36 = *_5;
>   _38 = _36 + 1;
>   *_5 = _38;
>   i.3_9 = (unsigned int) _38;
>   if (i.3_9 < _45)
>     goto <bb 5>;
>   else
>     goto <bb 6>;
> 
>   <bb 6>:
>   # D__lsm.10_10 = PHI <_33(8)>
>   # D__lsm.11_11 = PHI <1(8)>
>   *_12 = _33;
>   goto <bb 3>;
> 
>   <bb 3>:
>   return;
> ...
> 
> 
> V.
> 
> After second pass_lim/pass_dom pair. Note there are phis on the inner and
> outer loop header for the reduction and the iteration variables:
> ...
>   <bb 2>:
>   _5 = *.omp_data_i_4(D).i;
>   *_5 = 0;
>   _44 = *.omp_data_i_4(D).n;
>   _45 = *_44;
>   if (_45 != 0)
>     goto <bb 4>;
>   else
>     goto <bb 3>;
> 
>   <bb 4>:
>   _12 = *.omp_data_i_4(D).j;
>   _19 = *.omp_data_i_4(D).a;
>   D__lsm.10_50 = *_12;
>   D__lsm.11_51 = 0;
>   _25 = *.omp_data_i_4(D).sum;
>   D__lsm.14_40 = 0;
>   D__lsm.15_2 = 0;
>   D__lsm.16_1 = *_25;
>   D__lsm.17_46 = 0;
> 
>   <bb 5>: outer loop header
>   # D__lsm.14_13 = PHI <0(4), _38(8)>
>   # D__lsm.16_34 = PHI <D__lsm.16_1(4), sum.1_27(8)>
>   D__lsm.10_20 = 0;
>   D__lsm.11_22 = 1;
>   _21 = D__lsm.14_13;
>   D__lsm.12_28 = D__lsm.16_34;
>   D__lsm.13_30 = 0;
>   goto <bb 7>;
> 
>   <bb 7>: inner loop header, latch
>   # D__lsm.10_47 = PHI <0(5), _33(7)>
>   # D__lsm.12_49 = PHI <D__lsm.16_34(5), sum.1_27(7)>
>   _23 = D__lsm.10_47;
>   _24 = *_19[D__lsm.14_13][D__lsm.10_47];
>   sum.0_26 = D__lsm.12_49;
>   sum.1_27 = _24 + D__lsm.12_49;
>   D__lsm.12_31 = sum.1_27;
>   D__lsm.13_32 = 1;
>   _33 = D__lsm.10_47 + 1;
>   D__lsm.10_14 = _33;
>   D__lsm.11_15 = 1;
>   j.2_16 = (unsigned int) _33;
>   if (j.2_16 < _45)
>     goto <bb 7>;
>   else
>     goto <bb 8>;
> 
>   <bb 8>: outer loop latch
>   # D__lsm.10_35 = PHI <_33(7)>
>   # D__lsm.11_37 = PHI <1(7)>
>   # D__lsm.12_7 = PHI <sum.1_27(7)>
>   # D__lsm.13_8 = PHI <1(7)>
>   # sum.1_48 = PHI <sum.1_27(7)>
>   # _53 = PHI <_33(7)>
>   D__lsm.16_56 = sum.1_27;
>   D__lsm.17_57 = 1;
>   _36 = D__lsm.14_13;
>   _38 = D__lsm.14_13 + 1;
>   D__lsm.14_58 = _38;
>   D__lsm.15_59 = 1;
>   i.3_9 = (unsigned int) _38;
>   if (i.3_9 < _45)
>     goto <bb 5>;
>   else
>     goto <bb 6>;
> 
>   <bb 6>:
>   # D__lsm.10_10 = PHI <_33(8)>
>   # D__lsm.11_11 = PHI <1(8)>
>   # _43 = PHI <_33(8)>
>   # D__lsm.16_62 = PHI <sum.1_27(8)>
>   # D__lsm.17_63 = PHI <1(8)>
>   # D__lsm.14_64 = PHI <_38(8)>
>   # D__lsm.15_65 = PHI <1(8)>
>   *_5 = _38;
>   *_25 = sum.1_27;
>   *_12 = _33;
>   goto <bb 3>;
> 
>   <bb 3>:
>   return;
> ...

Sorry but staring at dumps doesn't make me understand the issue you
run into.  Where can I reproduce this if I have time to look at this?

>From the dump below I understand you want no memory references in
the outer loop?  So the issue seems to be that store motion fails
to insert the preheader load / exit store to the outermost loop
possible and thus another LIM pass is needed to "store motion" those
again?  But a simple testcase

int a;
int *p = &a;
int foo (int n)
{
  for (int i = 0; i < n; ++i)
    for (int j = 0; j < 100; ++j)
      *p += j + i;
  return a;
}

shows that LIM can do this in one step.  Which means it should
be investigated why it doesn't do this properly for your testcase
(store motion of *_25).

Simply adding two LIM passes either papers over a wrong-code
bug (in LIM or in DOM) or over a missed-optimization in LIM.

Richard.
 
> 
> VI.
> 
> After pass_dce, so before parloops-oacc-kernels:
> ...
>   <bb 2>:
>   _5 = *.omp_data_i_4(D).i;
>   *_5 = 0;
>   _44 = *.omp_data_i_4(D).n;
>   _45 = *_44;
>   if (_45 != 0)
>     goto <bb 4>;
>   else
>     goto <bb 3>;
> 
>   <bb 4>:
>   _12 = *.omp_data_i_4(D).j;
>   _19 = *.omp_data_i_4(D).a;
>   _25 = *.omp_data_i_4(D).sum;
>   D__lsm.16_1 = *_25;
> 
>   <bb 5>: outer loop header
>   # D__lsm.14_13 = PHI <0(4), _38(8)>
>   # D__lsm.16_34 = PHI <D__lsm.16_1(4), sum.1_27(8)>
>   goto <bb 7>;
> 
>   <bb 7>: inner loop header, latch
>   # D__lsm.10_47 = PHI <0(5), _33(7)>
>   # D__lsm.12_49 = PHI <D__lsm.16_34(5), sum.1_27(7)>
>   _24 = *_19[D__lsm.14_13][D__lsm.10_47];
>   sum.1_27 = _24 + D__lsm.12_49;
>   _33 = D__lsm.10_47 + 1;
>   j.2_16 = (unsigned int) _33;
>   if (j.2_16 < _45)
>     goto <bb 7>;
>   else
>     goto <bb 8>;
> 
>   <bb 8>: outer loop latch
>   _38 = D__lsm.14_13 + 1;
>   i.3_9 = (unsigned int) _38;
>   if (i.3_9 < _45)
>     goto <bb 5>;
>   else
>     goto <bb 6>;
> 
>   <bb 6>:
>   *_5 = _38;
>   *_25 = sum.1_27;
>   *_12 = _33;
>   goto <bb 3>;
> 
>   <bb 3>:
>   return;
> ...
> 
> Thanks,
> - Tom
> 
> 

-- 
Richard Biener <rguenther@suse.de>
SUSE LINUX GmbH, GF: Felix Imendoerffer, Jane Smithard, Graham Norton, HRB 21284 (AG Nuernberg)