[Bug target/85653] [nvptx] Work around subsequent bar.sync JIT/ptxas bug

[vries at gcc dot gnu.org]

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

--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
Author: vries
Date: Sat May  5 07:56:21 2018
New Revision: 259967

URL: https://gcc.gnu.org/viewcvs?rev=259967&root=gcc&view=rev
Log:
[nvptx] Add workaround for subsequent bar.syncs

2018-05-05  Tom de Vries  <tom@codesourcery.com>

        PR target/85653
        * config/nvptx/nvptx.c (WORKAROUND_PTXJIT_BUG_3): Define.
        (workaround_barsyncs): New function.
        (nvptx_reorg): Use workaround_barsyncs.
        * config/nvptx/nvptx.md (define_c_enum "unspecv"): Add UNSPECV_MEMBAR.
        (define_expand "nvptx_membar_cta"): New define_expand.
        (define_insn "*nvptx_membar_cta"): New insn.

Modified:
    trunk/gcc/ChangeLog
    trunk/gcc/config/nvptx/nvptx.c
    trunk/gcc/config/nvptx/nvptx.md