Re: [gomp4, committed] Fix parallelization for fortran oacc kernels tests

[Tom de Vries <Tom_deVries@mentor.com>]

On 18/06/15 12:48, Richard Biener wrote:
> On Thu, 18 Jun 2015, Tom de Vries wrote:
>
>> 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.
>
> You may now figure out that LIM needs FRE to detect equal memory
> references to apply store-motion.  But maybe the issues oacc
> lowering introduces are limited and under your control.
>

To show the context of the pass group, after this commit the pass group 
looks like this:
...
           NEXT_PASS (pass_sra_early);
           NEXT_PASS (pass_build_ealias);
           NEXT_PASS (pass_fre);
           NEXT_PASS (pass_oacc_kernels);
           PUSH_INSERT_PASSES_WITHIN (pass_oacc_kernels)
               NEXT_PASS (pass_ch_oacc_kernels);
               NEXT_PASS (pass_tree_loop_init);
               NEXT_PASS (pass_lim);
               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);
               NEXT_PASS (pass_tree_loop_done);
           POP_INSERT_PASSES ()
           NEXT_PASS (pass_merge_phi);
           NEXT_PASS (pass_dse);
           NEXT_PASS (pass_cd_dce);
...
In other words, the pass group is run directly after pass_fre.

When I move pass_fre before the pass group to directly after the pass 
group, I start seeing the failure mode you describe.

Thanks,
- Tom