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

["vries at gcc dot gnu.org" <gcc-bugzilla@gcc.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;
}
...