This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: The nvptx port [1/11+] indirect jumps
- From: Jakub Jelinek <jakub at redhat dot com>
- To: Bernd Schmidt <bernds at codesourcery dot com>
- Cc: Jeff Law <law at redhat dot com>, GCC Patches <gcc-patches at gcc dot gnu dot org>
- Date: Tue, 21 Oct 2014 23:30:25 +0200
- Subject: Re: The nvptx port [1/11+] indirect jumps
- Authentication-results: sourceware.org; auth=none
- References: <54451994 dot 9070209 at codesourcery dot com> <544519D8 dot 70606 at codesourcery dot com> <5446A55A dot 6060807 at redhat dot com> <5446C973 dot 3040102 at codesourcery dot com>
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
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