https://gcc.gnu.org/bugzilla/show_bug.cgi?id=83589
--- Comment #6 from Jakub Jelinek <jakub at gcc dot gnu.org> --- (In reply to Tom de Vries from comment #4) > Using this rudimentary workaround, I got the failing tests of this PR > passing again: > ... > diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c > index afb0e4dd185..3ac28b3d903 100644 > --- a/gcc/config/nvptx/nvptx.c > +++ b/gcc/config/nvptx/nvptx.c > @@ -78,6 +78,7 @@ > #include "target-def.h" > > #define WORKAROUND_PTXJIT_BUG 1 > +#define WORKAROUND_PTXJIT_BUG_2 1 > > /* The various PTX memory areas an object might reside in. */ > enum nvptx_data_area > @@ -4431,6 +4432,12 @@ nvptx_reorg (void) > if (TARGET_UNIFORM_SIMT) > nvptx_reorg_uniform_simt (); > > +#if WORKAROUND_PTXJIT_BUG_2 > + for (rtx_insn *insn = get_insns (); insn; insn = NEXT_INSN (insn)) > + if (LABEL_P (insn)) > + emit_insn_before (gen_fake_nop (), insn); > +#endif > + > regstat_free_n_sets_and_refs (); > > df_finish_pass (true); > diff --git a/gcc/config/nvptx/nvptx.md b/gcc/config/nvptx/nvptx.md > index f9c087b6d22..909484c329a 100644 > --- a/gcc/config/nvptx/nvptx.md > +++ b/gcc/config/nvptx/nvptx.md > @@ -994,6 +994,15 @@ > "" > "") > > +(define_insn "fake_nop" > + [(const_int 1)] > + "" > + "{ > + .reg .u32 %%nop_src; > + .reg .u32 %%nop_dst; > + mov.u32 %%nop_dst, %%nop_src; > + }") > + > (define_insn "return" > [(return)] > "" > ... Shouldn't it be sufficient to emit this only for JUMP_Ps that jump to immediately following LABEL_Ps (without intervening non-note insns)?