On 9/24/20 8:40 AM, Tom de Vries wrote:
Maybe we could require building libatomic for nvptx.
If one does not explicitly disable it, it builds,* which
I think is a good default. Additionally, libatomic is
only very rarely needed on x86-64 + nvptx and only
somewhat regularly on PowerPC + nvptx. Hence, I am not
sure whether nvptx's libgomp build should unconditionally
require it.
Tobias
* configure(.ac): If no --(enable|disable)-libatomic was
used, libatomic/configure.tgt is checked.
For nvptx*-*-* it is supported, for *), which includes gcn,
it is UNSUPPORTED=1.
PS: Besides the indirect dependency of nvptx on
__atomic_compare_exchange_16, libatomic is unsurprisingly
also required for code explicitly using atomics. (I wonder
whether we should add some libgomp/testsuite/ testcase like
the following with an appropriate effective-target.)
__uint128_t 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);
__uint128_t exp = 2;
__atomic_compare_exchange_n (&v, &exp, 7, 0, __ATOMIC_RELEASE,
__ATOMIC_ACQUIRE);
}
}
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
