Re: [RFC][nvptx, libgomp] Add 128-bit atomic support

[Jakub Jelinek <jakub@redhat.com>]

On Wed, Sep 02, 2020 at 12:22:28PM +0200, Tom de Vries wrote:
> And test-case passes on x86_64 with this patch (obviously, in
> combination with trigger patch above).
> 
> Jakub, WDYT?

I guess the normal answer would be use libatomic, but it isn't ported for
nvptx.
I guess at least temporarily this is ok, though I'm wondering why
you need __sync_*_16 rather than __atomic_*_16, or perhaps both __sync_* and
__atomic_*.

What happens if you try
unsigned __int128 v;
#pragma omp declare target (v)
int
main ()
{
  #pragma omp target
  {
    __atomic_add_fetch (&v, 1, __ATOMIC_RELAXED);
    __atomic_fetch_add (&v, 1, __ATOMIC_RELAXED);
    unsigned __int128v exp = 2;
    __atomic_compare_exchange_n (&v, &expected, 7, 0, __ATOMIC_RELEASE, __ATOMIC_ACQUIRE);
  }
}
etc. (see some gcc.dg/atomic* tests, ditto for __sync_*)?
I guess better not to throw everything into one test, because not every
target supports them all (e.g. I think x86_64 doesn't really do 128-bit
atomic loads because the cmpxchg16b insn are not appropriate for .rodata
locations).

	Jakub