[Bug libgomp/83589] [nvptx] mode-transitions.c and private-variables.{c,f90} execution FAILs at GOMP_NVPTX_JIT=-O0

[jakub at gcc dot gnu.org]

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)?