https://gcc.gnu.org/bugzilla/show_bug.cgi?id=83589
--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
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)]
""
...
