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
