[Bug target/104364] New: [12 Regression] OpenMP/nvptx regressions after "[nvptx] Add some support for .local atomics"

[Bug target/104364] New: [12 Regression] OpenMP/nvptx regressions after "[nvptx] Add some support for .local atomics"

[tschwinge at gcc dot gnu.org]

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

            Bug ID: 104364
           Summary: [12 Regression] OpenMP/nvptx regressions after
                    "[nvptx] Add some support for .local atomics"
           Product: gcc
           Version: 12.0
            Status: UNCONFIRMED
          Keywords: openmp
          Severity: normal
          Priority: P3
         Component: target
          Assignee: unassigned at gcc dot gnu.org
          Reporter: tschwinge at gcc dot gnu.org
                CC: vries at gcc dot gnu.org
  Target Milestone: ---
            Target: nvptx

As we've (slightly) discussed before, I acknowledge that commit
r12-6966-ge0451f93d9faa13495132f4e246e9bef30b51417 "[nvptx] Add some support
for .local atomics" is a good idea, thanks!  (Though I wonder, wouldn't it be
possible for the Driver/JIT to handle that internally?  Aha, you had the same
comment; first sentence in
<https://gcc.gnu.org/bugzilla/show_bug.cgi?id=96494#c2>.)

For certain GPU/Driver combinations, this change however causes regressions in
OpenMP/nvptx testing:

    PASS: libgomp.c/../libgomp.c-c++-common/for-11.c (test for excess errors)
    [-PASS:-]{+WARNING: program timed out.+}
    {+FAIL:+} libgomp.c/../libgomp.c-c++-common/for-11.c execution test

    PASS: libgomp.c/../libgomp.c-c++-common/for-12.c (test for excess errors)
    [-PASS:-]{+WARNING: program timed out.+}
    {+FAIL:+} libgomp.c/../libgomp.c-c++-common/for-12.c execution test

    PASS: libgomp.c/../libgomp.c-c++-common/for-3.c (test for excess errors)
    [-PASS:-]{+WARNING: program timed out.+}
    {+FAIL:+} libgomp.c/../libgomp.c-c++-common/for-3.c execution test

    PASS: libgomp.c/../libgomp.c-c++-common/for-5.c (test for excess errors)
    [-PASS:-]{+FAIL:+} libgomp.c/../libgomp.c-c++-common/for-5.c execution test

    PASS: libgomp.c/../libgomp.c-c++-common/for-6.c (test for excess errors)
    [-PASS:-]{+WARNING: program timed out.+}
    {+FAIL:+} libgomp.c/../libgomp.c-c++-common/for-6.c execution test

    PASS: libgomp.c/../libgomp.c-c++-common/for-9.c (test for excess errors)
    [-PASS:-]{+WARNING: program timed out.+}
    {+FAIL:+} libgomp.c/../libgomp.c-c++-common/for-9.c execution test

But varying, not always all of these.

Sometimes with, sometimes without the timeout.

Similar for C++.

... but only seen regressing for:

  - Nvidia Tesla K20c, Driver Version: 346.46
  - Nvidia Tesla K20c, Driver Version: 455.38
  - Nvidia Tesla K40c, Driver Version: 455.38
  - Nvidia Tesla K80, Driver Version: 440.33.01

... and never seen regressing for:

  - Nvidia Quadro P1000, Driver Version: 450.119.03
  - Nvidia GeForce GTX 1080, Driver Version: 440.33.01
  - Nvidia TITAN V, Driver Version: 455.23.05

(What is the underlying characteristic here?)


I did notice that this change affects nvptx target libraries -- but OpenMP
'mgomp' multilib only -- as follows:

'nvptx-none/mgomp/libatomic/cas_1_.o' (complete diff):

    @@ -113,7 +113,7 @@
     .loc 3 80 9
     or.b64 %r61,%r60,%r39;
     .loc 3 82 11
    -@ %r71 atom.cas.b64 %r62,[%r35],%r29,%r61;
    +atom.cas.b64 %r62,[%r35],%r29,%r61;
     mov.b64 {%r69,%r70},%r62;
     shfl.idx.b32 %r69,%r69,%r68,31;
     shfl.idx.b32 %r70,%r70,%r68,31;

Similar for 'nvptx-none/mgomp/libatomic/cas_2_.o'.

'nvptx-none/mgomp/libatomic/cas_4_.o' (complete diff):

    @@ -50,9 +50,9 @@
     mov.u32 %r25,%ar2;
     .loc 2 41 12
     ld.u32 %r28,[%r24];
    -@ %r37 membar.sys;
    -@ %r37 atom.cas.b32 %r29,[%r23],%r28,%r25;
    -@ %r37 membar.sys;
    +membar.sys;
    +atom.cas.b32 %r29,[%r23],%r28,%r25;
    +membar.sys;
     shfl.idx.b32 %r29,%r29,%r36,31;
     setp.eq.u32 %r31,%r29,%r28;
     selp.u32 %r30,1,0,%r31;

Similar for 'nvptx-none/mgomp/libatomic/cas_8_.o'.

Simlarly for several others in 'nvptx-none/mgomp/libatomic/',
'nvptx-none/mgomp/libbacktrace/dwarf.o', 'nvptx-none/mgomp/libgcc/atomic.o',
'nvptx-none/mgomp/libgfortran/single.o', several in
'nvptx-none/mgomp/libgomp/'.

That is, remove predication from all the "atomic" instructions.  I of course do
see the code changes related to that in the commit, but I do wonder whether
that's all intentional/correct?  (Or, is the 'atom' provably a no-op if not '@
%r37'?)

And, in 'nvptx-none/mgomp/libgomp/task.o' we've got several things like:

    -@ %r330 atom.cas.b32 %r208,[%frame+16],%r80,%r207;
    -@ %r330 membar.sys;
    +{
    +.reg .pred %eq_p;
    +.reg .u32 %val;
    +ld.u32 %val,[%frame+16];
    +setp.eq.u32 %eq_p,%val,%r80;
    +@ %eq_p st.u32 [%frame+16],%r207;
    +mov.u32 %r208,%val;
    +}

So these are the non-'atom' replacement sequences that now get synthesized
(haven't verified) -- but again this one lost the '@ %r330' predication from
all the "atomic" instructions?

And, is it correct here to use the non-'atom' replacement, though?  '%frame'
comes from:

    .visible .func GOMP_taskwait
    {
    .reg .u64 %stack;
    .reg .u64 %frame;
    .reg .u64 %sspslot;
    .reg .u64 %sspprev;
    {
    .reg .u32 %fstmp0;
    .reg .u64 %fstmp1;
    .reg .u64 %fstmp2;
    mov.u32 %fstmp0,%tid.y;
    mul.wide.u32 %fstmp1,%fstmp0,8;
    mov.u64 %fstmp2,__nvptx_stacks;
    add.u64 %sspslot,%fstmp2,%fstmp1;
    ld.shared.u64 %sspprev,[%sspslot];
    sub.u64 %frame,%sspprev,32;
    sub.u64 %stack,%frame,16;
    st.shared.u64 [%sspslot],%stack;
    }
    [...]

... which -- at least from my superficial look -- seems to be some kind of
'shared' location, thus mandating 'atom' access?  (But haven't looked in
detail, so I may certainly be wrong, of course.)

Are these changes and/or similar changes in nvptx code generation responsible
for the regressing test cases?  (Again, I have not looked.)


I'll report in case anything here changes due to Tom's further recent GCC/nvptx
back end commits.

[Bug target/104364] [12 Regression] OpenMP/nvptx regressions after "[nvptx] Add some support for .local atomics"

[vries at gcc dot gnu.org]

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

--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Thomas Schwinge from comment #0)
> 'nvptx-none/mgomp/libatomic/cas_1_.o' (complete diff):
>     
>     @@ -113,7 +113,7 @@
>      .loc 3 80 9
>      or.b64 %r61,%r60,%r39;
>      .loc 3 82 11
>     -@ %r71 atom.cas.b64 %r62,[%r35],%r29,%r61;
>     +atom.cas.b64 %r62,[%r35],%r29,%r61;
>      mov.b64 {%r69,%r70},%r62;
>      shfl.idx.b32 %r69,%r69,%r68,31;
>      shfl.idx.b32 %r70,%r70,%r68,31;
> 

Interesting.  This could be fallout from nvptx_reorg_uniform_simt.

Basically, the function assumes that atom.cas is predicable, which it's not
(anymore).

What is odd is that the resulting insn is still validated, I would have
expected that to fail.

[Bug target/104364] [12 Regression] OpenMP/nvptx regressions after "[nvptx] Add some support for .local atomics"

[rguenth at gcc dot gnu.org]

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

Richard Biener <rguenth at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
   Target Milestone|---                         |12.0

[Bug target/104364] [12 Regression] OpenMP/nvptx regressions after "[nvptx] Add some support for .local atomics"

[vries at gcc dot gnu.org]

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

--- Comment #2 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #1)
> What is odd is that the resulting insn is still validated, I would have
> expected that to fail.

Ah, the change is just silently rejected, this makes the problem visible:
...
diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc
index b3bb97c3c14d..9d083f46ce0a 100644
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
@@ -3150,7 +3150,8 @@ nvptx_reorg_uniform_simt ()
       rtx pred = nvptx_get_unisimt_predicate ();
       pred = gen_rtx_NE (BImode, pred, const0_rtx);
       pat = gen_rtx_COND_EXEC (VOIDmode, pred, pat);
-      validate_change (insn, &PATTERN (insn), pat, false);
+      bool changed_p = validate_change (insn, &PATTERN (insn), pat, false);
+      gcc_assert (changed_p);
     }
 }
...

[Bug target/104364] [12 Regression] OpenMP/nvptx regressions after "[nvptx] Add some support for .local atomics"

[vries at gcc dot gnu.org]

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

--- Comment #3 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Thomas Schwinge from comment #0)

> And, is it correct here to use the non-'atom' replacement, though?  '%frame'
> comes from:
> 
>     .visible .func GOMP_taskwait
>     {
>     .reg .u64 %stack;
>     .reg .u64 %frame;
>     .reg .u64 %sspslot;
>     .reg .u64 %sspprev;
>     {
>     .reg .u32 %fstmp0;
>     .reg .u64 %fstmp1;
>     .reg .u64 %fstmp2;
>     mov.u32 %fstmp0,%tid.y;
>     mul.wide.u32 %fstmp1,%fstmp0,8;
>     mov.u64 %fstmp2,__nvptx_stacks;
>     add.u64 %sspslot,%fstmp2,%fstmp1;
>     ld.shared.u64 %sspprev,[%sspslot];
>     sub.u64 %frame,%sspprev,32;
>     sub.u64 %stack,%frame,16;
>     st.shared.u64 [%sspslot],%stack;
>     }
>     [...]
> 
> ... which -- at least from my superficial look -- seems to be some kind of
> 'shared' location, thus mandating 'atom' access?  (But haven't looked in
> detail, so I may certainly be wrong, of course.)

Yeah, that's a mistake, when using -msoft-stack, we can't assume that
stack-related means local.

[Bug target/104364] [12 Regression] OpenMP/nvptx regressions after "[nvptx] Add some support for .local atomics"

[vries at gcc dot gnu.org]

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

--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
Created attachment 52341
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=52341&action=edit
Tentative patch

[Bug target/104364] [12 Regression] OpenMP/nvptx regressions after "[nvptx] Add some support for .local atomics"

[vries at gcc dot gnu.org]

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

--- Comment #5 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Thomas Schwinge from comment #0)
> ... but only seen regressing for:
> 
>   - Nvidia Tesla K20c, Driver Version: 346.46
>   - Nvidia Tesla K20c, Driver Version: 455.38
>   - Nvidia Tesla K40c, Driver Version: 455.38
>   - Nvidia Tesla K80, Driver Version: 440.33.01
> 
> ... and never seen regressing for:
> 
>   - Nvidia Quadro P1000, Driver Version: 450.119.03
>   - Nvidia GeForce GTX 1080, Driver Version: 440.33.01
>   - Nvidia TITAN V, Driver Version: 455.23.05
> 
> (What is the underlying characteristic here?)

Good question.

I've tested this using (recommended) driver 470.94 on boards:
- Kepler (Geforce GT 710)
- Maxwell (Quadro K620)
- Pascal (Geforce GT 1030)
- Turing (T400)
while iterating over dimensions { -mptx=3.1 , -mptx=6.3 } x {
GOMP_NVPTX_JIT=-O0, <default> }.

So I'm slightly surprised that I didn't see any regressions.

[Bug target/104364] [12 Regression] OpenMP/nvptx regressions after "[nvptx] Add some support for .local atomics"

[tschwinge at gcc dot gnu.org]

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

--- Comment #6 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
Thanks for having confirmed my findings and doubts -- seems I did correctly
understand a thing or two.  ;-)

(In reply to Tom de Vries from comment #5)
> (In reply to Thomas Schwinge from comment #0)
> > ... but only seen regressing for: [...]
> > 
> > ... and never seen regressing for: [...]
> > 
> > (What is the underlying characteristic here?)
> 
> Good question.
> 
> I've tested this using (recommended) driver 470.94 on boards: [...]
> while iterating over dimensions { -mptx=3.1 , -mptx=6.3 } x { GOMP_NVPTX_JIT=-O0, <default> }.
> 
> So I'm slightly surprised that I didn't see any regressions.

If indeed we're now generating some bad 'atom' code, it sure is confusing why
execution anyway PASSes for quite a number of configurations?  Are we just
"lucky", or is there some more fundamental issue that we're not even properly
using the concurrency in these configurations (and thus don't notice the 'atom'
issues)?

[Bug target/104364] [12 Regression] OpenMP/nvptx regressions after "[nvptx] Add some support for .local atomics"

[cvs-commit at gcc dot gnu.org]

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

--- Comment #7 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Tom de Vries <vries@gcc.gnu.org>:

https://gcc.gnu.org/g:04b54cc486cc6fcc40380445e500eaf46d7901dc

commit r12-7092-g04b54cc486cc6fcc40380445e500eaf46d7901dc
Author: Tom de Vries <tdevries@suse.de>
Date:   Thu Feb 3 14:00:02 2022 +0100

    [nvptx] Fix .local atomic regressions

    In PR target/104364, two problems were reported:
    - in muniform-simt mode, an atom.cas insn is no longer executed in the
      "master lane" only.
    - in msoft-stack mode, an __atomic_compare_exchange_n on stack memory is
      translated assuming it accesses local memory, while that's not the case.

    Fix these by:
    - ensuring that all insns with atomic attribute are also predicable, such
      that the validate_change in nvptx_reorg_uniform_simt will succeed, and
      asserting that it does, and
    - guarding the local atomics implementation with a new function
      nvptx_mem_local_p that correctly handles msoft-stack.

    Tested on x86_64 with nvptx accelerator.

    gcc/ChangeLog:

    2022-02-04  Tom de Vries  <tdevries@suse.de>

            PR target/104364
            * config/nvptx/nvptx-protos.h (nvptx_mem_local_p): Declare.
            * config/nvptx/nvptx.cc (nvptx_reorg_uniform_simt): Assert that
            change is validated.
            (nvptx_mem_local_p): New function.
            * config/nvptx/nvptx.md: Use nvptx_mem_local_p.
            (define_c_enum "unspecv"): Add UNSPECV_CAS_LOCAL.
            (define_insn "atomic_compare_and_swap<mode>_1_local"): New
            non-atomic, non-predicable define_insn, factored out of ...
            (define_insn "atomic_compare_and_swap<mode>_1"): ... here.
            Make predicable again.
            (define_expand "atomic_compare_and_swap<mode>"): Use
            atomic_compare_and_swap<mode>_1_local.

    gcc/testsuite/ChangeLog:

    2022-02-04  Tom de Vries  <tdevries@suse.de>

            PR target/104364
            * gcc.target/nvptx/softstack-2.c: New test.
            * gcc.target/nvptx/uniform-simt-1.c: New test.

[Bug target/104364] [12 Regression] OpenMP/nvptx regressions after "[nvptx] Add some support for .local atomics"

[vries at gcc dot gnu.org]

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

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
         Resolution|---                         |FIXED
             Status|UNCONFIRMED                 |RESOLVED

--- Comment #8 from Tom de Vries <vries at gcc dot gnu.org> ---
Fixed by "[nvptx] Fix .local atomic regressions".