This is the mail archive of the
mailing list for the GCC project.
IPA/cgraph: propagating node frequencies to offloaded functions
- From: Thomas Schwinge <thomas at codesourcery dot com>
- To: <gcc at gcc dot gnu dot org>, Jan Hubicka <hubicka at ucw dot cz>, Bernd Schmidt <bschmidt at redhat dot com>, Jakub Jelinek <jakub at redhat dot com>, Tom de Vries <vries at codesourcery dot com>
- Date: Fri, 12 Feb 2016 17:40:58 +0100
- Subject: IPA/cgraph: propagating node frequencies to offloaded functions
- Authentication-results: sourceware.org; auth=none
As I'm touching areas of GCC here, that I have no noteworthy experience
with (IPA optimizations, cgraph), I'm asking for your help. Thanks!
This is primarily to implement a better "avoid offloading" policy for
un-parallelized OpenACC kernels constructs with nvptx offloading,
but I imagine that potentially also any non-nvptx offloading can benefit
from this, by better parameterization of GCC's optimization passes.
For this consideration, simplified, "offloading" means that:
#pragma omp target
... is re-written into something like:
..., and the code of main_offloaded is then not run on the CPU but is in
GOMP_target launched to execute on an offloading device (not actually
important for this consideration).
GOMP_target, described by gcc/omp-builtins.def:BUILT_IN_GOMP_TARGET, is
implemented externally to the compiler (in libgomp), but we "control" its
implementation, and so we're free to have the compiler make certain
assumptions about its behavior.
(I guess) due to it being passed to the intermediary GOMP_target call,
the address of main_offloaded "escapes". But, as we know how/what for
GOMP_target is using it, we should be able to (teach the compiler to)
handle it like a "static" function -- I hope.
For reference, I'm assuming the example above should look/work a bit like
the following example:
static __attribute__((noinline)) int fun_s(int x)
return x + 1;
..., where I do observe that, for example, main's
NODE_FREQUENCY_EXECUTED_ONCE is propagated to fun_s. (For main
initialized in gcc/predict.c:compute_function_frequency, and the
propagation to fun_s then happens in
test.c.067i.profile_estimate: "Node foo_s promoted to executed once".
However, this node frequency propagation does not currently happen for
main_offloaded. Looking into this for a bit, I think I have identified
IPA/cgraph does not consider main_offloaded to be "local" (as in struct
cgraph_local_info; handled/set in
gcc/ipa-visibility.c:cgraph_node::local_p), so does not even attempt to
do such optimizations. (But, per my comment above, we should actually be
able to treat main_offloaded at least similar to a "static" function.)
If I hack gcc/ipa-visibility.c:cgraph_node::local_p to forcefully mark
main_offloaded as "local", the node frequency propagation still doesn't
work correctly, because no callers (cgraph_edge) have been registered for
the respective cgraph_node, so gcc/ipa-profile.c:ipa_propagate_frequency
doesn't have the data available that it needs in order to set the
ipa_propagate_frequency_data, which later will be used to set the node
As it already has some support for gomp_parallel and gomp_task (which I
understand to have similar semantics in that regard, comparing to
gomp_target), maybe extending gcc/cgraphbuild.c will help with the latter
issue? (That bit of code handling gomp_* looks like its purpose is to
"pass through" function addresses, such as main_offloaded, which are
called through intermediary functions, such as GOMP_target, but I have
not yet confirmed if that's exactly what's being done there.)
What to do about the cgraph_node "local" bit for main_offloaded?
Am I generally on the right track?