[Bug target/96932] New: [nvptx] atomic_exchange missing barrier

[Bug target/96932] New: [nvptx] atomic_exchange missing barrier

[vries at gcc dot gnu.org]

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

            Bug ID: 96932
           Summary: [nvptx] atomic_exchange missing barrier
           Product: gcc
           Version: 11.0
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: target
          Assignee: unassigned at gcc dot gnu.org
          Reporter: vries at gcc dot gnu.org
  Target Milestone: ---

After digging into GOMP_atomic_start/end I realized these also imply barrier
semantics.

And looking at the source code used for nvptx in libgomp/config/accel/mutex.h,
that should be fine:
...
static inline void
gomp_mutex_lock (gomp_mutex_t *mutex)
{
  while (__sync_lock_test_and_set (mutex, 1))
    /* spin */ ;
}

static inline void
gomp_mutex_unlock (gomp_mutex_t *mutex)
{
  __sync_lock_release (mutex);
}
...

However, when looking at the resulting code in libgomp.a we see there's no
barrier for GOMP_atomic_start:
...
.visible .func GOMP_atomic_start
{
.reg .u32 %r22;
.reg .pred %r23;
$L2:
.loc 1 51 10
atom.global.exch.b32 %r22,[atomic_lock],1;
.loc 1 51 9
setp.ne.u32 %r23,%r22,0;
@ %r23 bra $L2;
.loc 2 43 1
ret;
}
...

While there is for GOMP_atomic_end:
...
.visible .func GOMP_atomic_end
{
.reg .u32 %r22;
.loc 1 58 3
membar.sys;
mov.u32 %r22,0;
st.global.u32 [atomic_lock],%r22;
.loc 2 49 1
ret;
}
...

[Bug target/96932] [nvptx] atomic_exchange missing barrier

[vries at gcc dot gnu.org]

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

--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
FWIW, I've tried this test-case to trigger the problem, but it runs fine:
...
$ cat libgomp/testsuite/libgomp.oacc-c-c++-common/test.c  
/* { dg-do run } */

#include <stdlib.h>
#include <openacc.h>

#define assert(COND) \
  do { \
    if (!(COND)) \
      abort ();  \
  } while (0) 

int c;
#pragma acc declare copyin (c)

int
main (void)
{
  int *p = acc_malloc (4);

#pragma acc parallel num_gangs (1) deviceptr (p)
  {
    *p = 0;
  }

#pragma acc parallel num_gangs (100) deviceptr (p)
  {
    int i;
    for (i = 0; i < 10000; ++i)
      {
        int res;
        while (1)
          {
            res = __sync_lock_test_and_set (p, 1);
            if (res == 1)
              continue;
            break;
          }
        assert (res == 0);
        assert (*p == 1);

        assert (c == 0);
        c += 1;
        assert (c == 1);
        c -= 1;
        assert (c == 0);

        __sync_lock_release (p);
      }
  }
  return 0;
}
...

[Bug target/96932] [nvptx] atomic_exchange missing barrier

[vries at gcc dot gnu.org]

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

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

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |rguenth at gcc dot gnu.org

--- Comment #2 from Tom de Vries <vries at gcc dot gnu.org> ---
It could be interesting to use a ptx simulator, perhaps that one could trigger
this.

[Bug target/96932] [nvptx] atomic_exchange missing barrier

[burnus at gcc dot gnu.org]

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

Tobias Burnus <burnus at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |burnus at gcc dot gnu.org

--- Comment #3 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Crossref: PR100497 - fails on Volta without
  membar.sys;
before
  atom.global.exch.b32

Unfortunately, compared to pre-Volta, it is very slow - membar.gl is still slow
but a bit less.  Using (→ sm_70) fence.sys / fence.gnu instead of
fence.sc.{sys,gnu} (= membar.{sys,gl} on >= sm_70) does not seem to make a
performance difference for PR100497.

[Bug target/96932] [nvptx] atomic_exchange missing barrier

[vries at gcc dot gnu.org]

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

--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #3)
> Crossref: PR100497 - fails on Volta without
>   membar.sys;
> before
>   atom.global.exch.b32
> 
> Unfortunately, compared to pre-Volta, it is very slow - membar.gl is still
> slow but a bit less.  Using (→ sm_70) fence.sys / fence.gnu instead of
> fence.sc.{sys,gnu} (= membar.{sys,gl} on >= sm_70) does not seem to make a

fence.sc.gpu, funny typo :)

> performance difference for PR100497.

The GOMP_atomic_start/GOMP_atomic_end are fallbacks, and unfortunately cannot
be expected to be too optimal.

Following the introduction of -mptx=6.3 we can add support for atom.cas.b16
(well, once we also introduce misa=sm_70), and that should be the optimal
solution.

[Bug target/96932] [nvptx] atomic_exchange missing barrier

[vries at gcc dot gnu.org]

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

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

[Bug target/96932] [nvptx] atomic_exchange missing barrier

[vries at gcc dot gnu.org]

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

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

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

--- Comment #6 from Tom de Vries <vries at gcc dot gnu.org> ---
Duplicate.

*** This bug has been marked as a duplicate of bug 100497 ***