[Bug tree-optimization/89499] [7/8/9 Regression] ICE in expand_UNIQUE, at internal-fn.c:2605

[tschwinge at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=89499

--- Comment #5 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
Created attachment 46094
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=46094&action=edit
[WIP] [PR89499] CIF_OPENACC_ROUTINE_MISMATCH

(In reply to Thomas Schwinge from comment #4)
> Unless there is an OpenACC 'bind' clause involved (also 'nohost' clause?)
> (which are not yet implemented, so not relevant in this discussion), it is
> permissible (and worthwhile for the usual reasons) to inline such functions
> into one another, given proper nesting of OpenACC levels of parallelism. 
> The latter is enforced by construction, by detecting improper caller/callee
> combinations (for example, trying to call a 'gang' routine from a 'vector'
> context).
> 
> By construction, 'IFN_UNIQUE' etc. can only appear in 'oacc function'
> functions.
> 
> There are 'oacc function' functions that do not contain 'IFN_UNIQUE' etc.
> (for example, 'seq' functions, such as math library functions), but which
> might still benefit from inlining.

And, in particular, also the C++ 'acc_on_device' forwarding function with
correctly typed argument.  (See 'openacc.h'.)

> I suppose it is fine to do inlining if the outer function will then be
> handled by 'oaccdevlow'.
> 
> And, I suppose it will be reasonable to forbid inlining of 'oacc function'
> functions into non-'oacc function' functions

..., and this breaks said 'acc_on_device' case...  :-/

This is: the '#pragma acc routine' 'inline' function 'acc_on_device' is
regularely called from non-'#pragma acc routine' functions.

It becomes worse if the C++ forwarding function with correctly typed argument
is actually marked 'always_inline' -- as it probably should be?

So, this has to be permitted.

> because in that case you're
> not applying any OpenACC parallelism anyway, at least for a backportable ICE
> fix, then later possibly more logic added to allow that.
> 
> That will probably be reasonably simple to implement; I'll give it some
> further thought, and testing.

Attaching my WIP patch for that, anyway.

But, as discussed, it seems that it'll have to be a bit more elaborate: do
allow such inlining after all.  I'm still pondering how to best do that.  (a)
Clean up the unexpected stuff at the place where it currently ICEs (by
refactoring stuff currently done in 'oaccdevlow'?).  (b) When the problematic
inlining happens, tag the inlined-into function such that 'oaccdevlow' will
then clean it up in there.  (c) Go as far as cloning the '#pragma acc routine'
functions early, so that the "host" function doesn't get the OpenACC magic
applied, but only the "OMP" one does.  (If this has other benefits maybe, by
not cluttering the "hose" code with OpenACC magic?)

Option (b) would be something similar to what Jakub had suggested:

(In reply to Jakub Jelinek from comment #3)
> Or it is inlinable, but we need some cleanup, in that case perhaps have some
> cfun->* flag that would be initially set to whether the function has
> oacc_get_fn_attrib and would be ored into functions into which those
> functions were inlined, and then the oaccdevlower pass would clean that
> stuff up or whatever.

I think I'll look into that option first, try to locate the places where such
processing is to happen, and leave option (c) for later.