The nvptx port [1/11+] indirect jumps

[Jakub Jelinek]

On Tue, Oct 21, 2014 at 11:00:35PM +0200, Bernd Schmidt wrote:
> On 10/21/2014 08:26 PM, Jeff Law wrote:
> >>    * optabs.c (emit_indirect_jump): Test HAVE_indirect_jump and emit a
> >>    sorry if necessary.
> >So doesn't this imply no hot-cold partitioning since we use indirect
> >jumps to get across the partition?  Similarly doesn't this imply other
> >missing features (setjmp/longjmp, nonlocal gotos, computed jumps?
> 
> Pretty much yes to all.
> 
> >Do you need some mechanism to ensure that hot/cold partitioning isn't
> >enabled?
> 
> I guess I could clear flag_reorder_blocks_and_partition in
> nvptx_option_override. The problem hasn't come up so far.
> 
> >Do you need some kind of message specific to the other
> >features, or are we going to assume that the user will map from the
> >indirect jump message back to the use of setjmp/longjmp or something
> >similar?
> 
> I have some sorry calls in things like a dummy nonlocal_goto pattern. It
> doesn't quite manage to catch everything without an ICE yet though.

With all the sorry additions, what is actually the plan for OpenMP (dunno how
OpenACC is different in this regard)?
At least for OpenMP, the best would be if the #pragma omp target regions
and/or #pragma omp declare target functions contain anything a particular
offloading accelerator can't handle, instead of failing the whole
compilation perhaps just emit some at least by default non-fatal warning
and not emit anything for the particular offloading target, which would mean
either host fallback, or, if some other offloading target succeeded, just
that target.
The unsupported stuff can be machine dependent builtins that can't be
transformed, or e.g. the various things you've listed as unsupportable by
the PTX backend right now.

	Jakub