[hsa] Produce naked kernels for simple omp loops

Hi,

this patch removes omp-expansion generated control flow from simple
omp parallel loops (identified by the previous patch) and makes the
functions generated from these loops with attribute "hsakernel" which
means that it also adds two parameters denoting the iteration size and
group size (which I take from chunk size, if there is any).

Before explaining this in more detail, let me clarify two things.
First, for various reasons, we do not attempt to make any provisions
for OMP 4.0 yet, so do not look for it in the code.  Second, this is
just a proof-of concept patch, I understand that it is quite horrible
in a number of ways which is a result of my lack of experience in this
area and fairly unpleasant time pressure that was caused mainly by
things well beyond our control.  I will wholeheartedly appreciate any
guidance in reimplementing this properly.

Basically, for expansion into HSA, we want omp parallel loop to only
contain the body of the loop, it should not contain the computation of
the portion of workload the particular thread or the loop over that
portion.  Each iteration is one thread for us and the size of the
iteration space must be set by code invoking the kernel.

For example, in

#pragma omp parallel for shared(a,b) firstprivate(n) private(i)
for (i = 0; i < n; i++)
  a[i] = b[i] * b[i];

We want the generated omp child function to only contain:

i = omp_get_thread_num();
a[i] = b[i] * b[i];

and nothing more (modulo loading from omp_data structure and other
unrelated stuff).  This is implemented by ignoring large parts of
expand_omp_for_static_nochunk when such simple loop is being expanded.
A far bigger obstacle was that the code invoking the kernel (which
represents the whole parallel construct) and thus code outside of the
parallel construct must calculate the iteration size in order to
verify that the loop should be run at all and so that it can provide
it as a kernel parameter.  Here I ran into severe problems caused by
variable remapping, because when I attempted to just move the
conditions above the parallel statement, the variables were already
remapped, resulting in undefined loads, and I did not find any way of
mapping them back.  Eventually I resorted to hiding away the loop
parameters one more time in the gimple statement itself but that is of
course a terrible hack.

Another problem which I have not attempted to solve in this patch is
how to generate both code for the host and the accelerator.  Basically
we would want OMP expansion to generate two very different child
functions for OMP parallel loops which we want to turn into kernels
but the bottom-up structure of OMP expansion makes this very
difficult.

I have not been able to find a public branch for offloading to Nvidia
PTX but I assume it faces the same problem.  Have you guys attempted
to tackle them somehow?  Or am I just completely misguided in my
thoughts?

Anyway, as with the previous patches, I have bootstrapped this just to
catch errors and tested it on a number of OMP testcases and it did not
introduce any new failures.  Committed to the hsa branch.

Thanks,

Martin

2014-09-26  Martin Jambor  <mjambor@suse.cz>

	* gimple.c (gimple_build_omp_for): Allocate prev_first_iter.
	* gimple.h (gimple_statement_omp_for): New field orig_first_iter.
	(gimple_omp_for_set_orig_first_iter): New function.
	* gimplify.c (gimplify_omp_for): Use it.
	* omp-low.c (omp_region): New fields req_group_size and orig_first_iter.
	(adjust_for_condition): New function.
	(get_omp_for_step_from_incr): Likewise.
	(extract_omp_for_data): Moved some functionality to the above two new
	functions.
	(create_omp_child_function): Do not append hsa attributes to child_fn.
	(expand_parallel_call): Handle kernelized parallel regions.
	(expand_omp_for_static_nochunk): Kernelize marked loops.
	(expand_omp_for): Copy prev_first_iter of to-be-kernelized omp for
	statements to the region structure, mark requested group size.  Always
	expand these loops with expand_omp_for_static_nochunk.
	* tree-sra.c (ipa_sra_preliminary_function_checks): Test TREE_USED.
---
 gcc/gimple.c   |   1 +
 gcc/gimple.h   |  20 +++
 gcc/gimplify.c |  13 ++
 gcc/omp-low.c  | 462 +++++++++++++++++++++++++++++++++++++--------------------
 gcc/tree-sra.c |   2 +-
 5 files changed, 338 insertions(+), 160 deletions(-)

[hsa] Produce naked kernels for simple omp loops

Commit Message

Patch