Hi Torsten,
Thanks for the bug report.  The STORE_FLAG_VALUE=1 patch was one of a series to dramatically improve the quality of nvptx code.  Alas not all of them have yet
been reviewed/approved, and it's likely these later improvements address the quality regression you're seeing.

The other patches in the "nvptx Boolean" series are:
patchq3: nvptx: Expand QI mode operations using SI mode instructions.
https://gcc.gnu.org/pipermail/gcc-patches/2022-January/587999.html

patchq4: nvptx: Fix and use BI mode logic instructions (e.g. and.pred).
https://gcc.gnu.org/pipermail/gcc-patches/2022-January/588555.html

[and purely for reference, my other outstanding nvptx patches are]
patchn: nvptx: Improved support for HFMode including neghf2 and abshf2.
https://gcc.gnu.org/pipermail/gcc-patches/2022-January/587949.html

patchw: nvptx: Add support for 64-bit mul.hi (and other) instructions.
https://gcc.gnu.org/pipermail/gcc-patches/2022-January/588453.html

And one other related patch is that there's also a middle-end SUBREG
patch intended to improve code generation on nvptx is also pending at:

patchs: Simplify paradoxical subreg extensions of TRUNCATE
https://gcc.gnu.org/pipermail/gcc-patches/2021-September/578848.html



My guess is that patchq3+patchq4 above should (hopefully) resolve this
particular regression.  If you could give them a spin on your system
to see if they reduce register pressure sufficiently for this case,
that would be greatly appreciated.  As you can read in the above postings,
the total number of instructions/registers (after all of these changes)
should be dramatically reduced.

I'll see what I can do from my end.

(In reply to Roger Sayle from comment #1)
> The other patches in the "nvptx Boolean" series are:
> patchq3: nvptx: Expand QI mode operations using SI mode instructions.
> https://gcc.gnu.org/pipermail/gcc-patches/2022-January/587999.html
> 
> patchq4: nvptx: Fix and use BI mode logic instructions (e.g. and.pred).
> https://gcc.gnu.org/pipermail/gcc-patches/2022-January/588555.html
> 
> [and purely for reference, my other outstanding nvptx patches are]
> patchn: nvptx: Improved support for HFMode including neghf2 and abshf2.
> https://gcc.gnu.org/pipermail/gcc-patches/2022-January/587949.html
> 
> patchw: nvptx: Add support for 64-bit mul.hi (and other) instructions.
> https://gcc.gnu.org/pipermail/gcc-patches/2022-January/588453.html

FYI, I'm currently testing these.

Additional patch proposed:
https://gcc.gnu.org/pipermail/gcc-patches/2022-February/589802.html
I need to figure out how to (re)produce Thomas' "used N registers" reports.
If someone could summarize the effect of this patch (and previous patches) on register usage, that would be much appreciated (and help reviewers).

Created attachment 52370 [details]
'_muldc3-good.o'

Created attachment 52371 [details]
'_muldc3-bad.o'

Created attachment 52372 [details]
'_muldc3-WIP.o'

All your three patches combined still doesn't help resolve the problem.
And, what I realized: they don't even change the Nvidia/CUDA Driver reported "used [...] registers".
Does that mean that the Driver "doesn't care" that we feed into it simple PTX code, using less PTX registers -- it seems able to do these optimizations all by itself?  :-O
(At least regarding number of registers used -- I didn't verify the SASS code generated.)
(Emitting cleaner, "de-cluttered" code in GCC/nvptx is still very much valuable, of course: if only for our own visual consumption!  ... at least as long as it doesn't make 'nvptx.md' etc. much more complex...)

For "good" vs. "bad"/"not-so-good" (before vs. after "nvptx: Transition nvptx backend to STORE_FLAG_VALUE = 1"), the only code generation difference is in the '__muldc3' function ('nvptx-none/libgcc/_muldc3.o'), and that one is the only '.extern' dependency aside from the '__reduction_lock' global variable ('nvptx-none/libgcc/reduction.o').
(In the following, working around <https://gcc.gnu.org/PR104416> "'lto-wrapper' invoking 'mkoffload's with duplicated command-line options".)

This means, I can conveniently manually create a minimal nvptx 'libgcc.a':

    $ cp build-gcc-offload-nvptx-none/nvptx-none/libgcc/_muldc3.o ./
    $ rm -f libgcc.a && ar q libgcc.a _muldc3.o build-gcc-offload-nvptx-none/nvptx-none/libgcc/reduction.o

..., and compile 'libgomp.oacc-c-c++-common/reduction-cplx-dbl.c' with '-foffload=nvptx-none=-Wl,-L.'.  (Via 'GOMP_DEBUG=1' verified that identical PTX code is loaded to GPU.)

Then, hand-modify '_muldc3.o', re-create 'libgcc.a', re-compile, re-execute.

Verified that before "nvptx: Transition nvptx backend to STORE_FLAG_VALUE = 1" '__muldc3' (attached '_muldc3-good.o') works fine, and after "nvptx: Transition nvptx backend to STORE_FLAG_VALUE = 1" '__muldc3' (attached '_muldc3-bad.o') does show the problem originally reported here.

I then gradually morphed the former into the latter (beginning with eliding simple changes like renumbered registers etc.), until only one last change was necessary to turn "good" into "bad" (attached '_muldc3-WIP.o'); showing the "still-good" state:

    [...]
    @@ -1716,8 +1718,16 @@
     cvt.rn.f64.s32 %r84,%r27;
     copysign.f64 %r61,%r61,%r84;
     .loc 2 1981 32
    +// Current/after "nvptx: Transition nvptx backend to STORE_FLAG_VALUE = 1":
    +/*
     selp.u32 %r86,1,0,%r138;
     mov.u32 %r85,%r86;
    +*/
    +// Before "nvptx: Transition nvptx backend to STORE_FLAG_VALUE = 1":
    +set.u32.leu.f64 %r86,%r57,0d7fefffffffffffff;
    +neg.s32 %r87,%r86;
    +mov.u32 %r85,%r87;
    +//
     cvt.u16.u8 %r140,%r85;
     mov.u16 %r89,%r140;
     xor.b16 %r88,%r89,1;
    @@ -1770,8 +1780,16 @@
    [...]

That is, things go "bad" if we here use the '%r138' that was computed (and used) earlier in the code, and things are "good" if we re-compute locally.  Same for '%r139' in the second code block.

(Interestingly, it is tolerated if one of the long-lived registers are used, but things go "bad" only if *both* are used.  But that's probably just a distraction?  And again, I have not inspected the actual SASS code, but just looked at the JIT-time "used [...] registers".)

Do we thus conclude that what happens is that the "nvptx: Transition nvptx backend to STORE_FLAG_VALUE = 1" changes here enable an "optimization" in GCC such that values that have previously been compute may be re-used later in the code, without re-computing them.  But on the flip side, of course, this means that the values have to kept live in (SASS) registers.  (That's just my theory; I haven't verified the actual SASS.)

In other words: at least in this case here, it seems preferrable to re-compute instead of keeping registers occupied.  (But I'm of course not claiming that to be a simple yes/no decision...)

It seem we're now in territory of tuning CPU vs. GPU code generation?

Certainly, GCC has not seen much care for the latter (GPU code generation).
I mean: verify GCC pass pipeline generally, and parameterization of individual passes for GPU code generation.
Impossible to get "right", of course, but maybe some heuristics for CPU vs. GPU may be discovered and implemented?
I'm sure there must be some literature on that topic?

All that complicated by the fact the with the (several different versions of the) Nvidia/CUDA Driver's PTX -> SASS translation/optimization we have another moving part...

Well.  Here's another problem.

Re-testing things using a "bad" '__muldc3' from a build with your three patches applied, I notice that my '_muldc3-WIP.o' "old"/replacement code uses a 'set.u32.leu.f64', 'neg.s32', 'mov.u32' sequence.  That works.
However, if I replace that with a localized (its own '.reg .pred %r138_;') 'setp.leu.f64', 'selp.u32' sequence (as per your code generation changes), then it doesn't in fact work!  Huh.

Will later continue looking into this.

OK!  Putting your "nvptx: Add support for 64-bit mul.hi (and other) instructions" on top, that considerably changes (simplifies!) the generated '__muldc3' PTX code; the regression disappears.  :-)

(I have, so far, only manually tested 'libgomp.oacc-c-c++-common/reduction-cplx-dbl.c'.  I'll report later in the unlikely case that any other/new issues should appear.)


(And, will later test your "nvptx: Tweak constraints on copysign instructions" on top, too.)

The master branch has been updated by Tom de Vries <vries@gcc.gnu.org>:

https://gcc.gnu.org/g:9bacd7af2e3bba9ddad17e7de4e2d299419d819d

commit r12-7167-g9bacd7af2e3bba9ddad17e7de4e2d299419d819d
Author: Roger Sayle <roger@nextmovesoftware.com>
Date:   Fri Feb 4 04:13:53 2022 +0100

    PR target/104345: Use nvptx "set" instruction for cond ? -1 : 0
    
    This patch addresses the "increased register pressure" regression on
    nvptx-none caused by my change to transition the backend to a
    STORE_FLAG_VALUE = 1 target.  This improved code generation for the
    more common case of producing 0/1 Boolean values, but unfortunately
    made things marginally worse when a 0/-1 mask value is desired.
    Unfortunately, nvptx kernels are extremely sensitive to changes in
    register usage, which was observable in the reported PR.
    
    This patch provides optimizations for -(cond ? 1 : 0), effectively
    simplify this into cond ? -1 : 0, where these ternary operators are
    provided by nvptx's selp instruction, and for the specific case of
    SImode, using (restoring) nvptx's "set" instruction (which avoids
    the need for a predicate register).
    
    This patch has been tested on nvptx-none hosted on x86_64-pc-linux-gnu
    with a "make" and "make -k check" with no new failures.  Unfortunately,
    the exact register usage of a nvptx kernel depends upon the version of
    the Cuda drivers being used (and the hardware), but I believe this
    change should resolve the PR (for Thomas) by improving code generation
    for the cases that regressed.
    
    gcc/ChangeLog:
    
            PR target/104345
            * config/nvptx/nvptx.md (sel_true<mode>): Fix indentation.
            (sel_false<mode>): Likewise.
            (define_code_iterator eqne): New code iterator for EQ and NE.
            (*selp<mode>_neg_<code>): New define_insn_and_split to optimize
            the negation of a selp instruction.
            (*selp<mode>_not_<code>): New define_insn_and_split to optimize
            the bitwise not of a selp instruction.
            (*setcc_int<mode>): Use set instruction for neg:SI of a selp.
    
    gcc/testsuite/ChangeLog:
    
            PR target/104345
            * gcc.target/nvptx/neg-selp.c: New test case.

The master branch has been updated by Tom de Vries <vries@gcc.gnu.org>:

https://gcc.gnu.org/g:6d98e83b2c919bd9fba2c61333d613bafc37357f

commit r12-7168-g6d98e83b2c919bd9fba2c61333d613bafc37357f
Author: Roger Sayle <roger@nextmovesoftware.com>
Date:   Tue Feb 8 20:56:55 2022 +0100

    nvptx: Tweak constraints on copysign instructions
    
    Many thanks to Thomas Schwinge for confirming my hypothesis that the register
    usage regression, PR target/104345, is solely due to libgcc's _muldc3 function.
    In addition to the isinf functionality in the previously proposed nvptx patch at
    https://gcc.gnu.org/pipermail/gcc-patches/2022-January/588453.html which
    significantly reduces the number of instructions in _muldc3, the patch below
    further reduces both the number of instructions and the number of explicitly
    declared registers, by permitting floating point constant immediate operands
    in nvptx's copysign instruction.
    
    Fingers-crossed, the combination with all of the previous proposed nvptx
    patches improves things.  Ultimately, increasing register usage from 50 to
    51 registers, reducing the number of concurrent threads by ~2%, can easily
    be countered if we're now executing significantly fewer instructions in each
    kernel, for a net performance win.
    
    This patch has been tested on nvptx-none hosted on x86_64-pc-linux-gnu
    with a "make" and "make -k check" with no new failures.
    
    gcc/ChangeLog:
    
            * config/nvptx/nvptx.md (copysign<mode>3): Allow immediate
            floating point constants as operands 1 and/or 2.

This should now be fixed on mainline.