https://gcc.gnu.org/bugzilla/show_bug.cgi?id=126015
Bug ID: 126015
Summary: [OpenMP] target [nvptx]: reduction(+:arr[0:2]) of
doubles triggers missing __atomic_compare_exchange_16
Product: gcc
Version: 16.1.1
Status: UNCONFIRMED
Severity: normal
Priority: P3
Component: libgomp
Assignee: unassigned at gcc dot gnu.org
Reporter: schulz.benjamin at googlemail dot com
CC: jakub at gcc dot gnu.org
Target Milestone: ---
Created attachment 64883
--> https://gcc.gnu.org/bugzilla/attachment.cgi?id=64883&action=edit
main.cpp
Hi there, i wrote a library for linear algebra, which works with templates.
I wanted to know what happens if one uses std::complex. It turns out, that
almost everything works. std:: complex is completely mappable,
except reductions involving complex variables.
If one uses a reduction over a complex variable on nvptx target in an openmp
simd loop, gcc 16 would always return: unresolved symbol
__atomic_compare_exchange_16
of course, one may resolve this by including libatomic. On my system, gentoo,
this is not there in nvptx.
But using atomics on gpu is also disadvised because they lead to a severe
performance drop with many threads.
But why does gcc want these atomics for reductions over complex numbers at all?
I simplified my code further, to the point of removing std::complex.
And it turns out that even this here fails without atomics:
A reduction over a native [2] array, which is allowed by the openmp standard
and where the array members are totally independent and atomics should not be
needed.
int main(int argc, char** argv)
{
// A 2-element double array is exactly 16 bytes.
double res_view[2] = {0.0, 0.0};
// OpenMP array section reduction on a 16-byte buffer triggers the bug
#pragma omp target teams distribute parallel for simd
reduction(+:res_view[0:2])
for (int i = 0; i < 100; ++i) {
res_view[0] += 1.0;
res_view[1] += 2.0;
}
std::cout << "Real Sum: " << res_view[0] << ", Imag Sum: " << res_view[1]
<< "\n";
return 0;
}
compile with
g++ -march=native -fopenmp -foffload=nvptx-none -fno-stack-protector
-U_FORTIFY_SOURCE -O3 -Wall ./main.cpp
will result in
/usr/bin/cmake -E cmake_link_script CMakeFiles/gpu_compiler_test.dir/link.txt
--verbose=1
unresolved symbol __atomic_compare_exchange_16
collect2: error: ld gab 1 als Ende-Status zurück
There should be no atomic needed since we have two entirely separate iterations
clang will compile this example on gpu of course...
googles ai describes the problem as follows:
instead of lowering a 2-element double array reduction into independent 8-byte
scalar registers or sequential steps,
GCC's intermediate code pipeline treats the 16-byte array section as an opaque
block.
It emits a host-style __atomic_compare_exchange_16 loop during the final
inter-team reduction rollup, which the nvptx target backend cannot resolve.
That is probably why gcc can't make reductions with complex variables easily.
I have attached a cpp file to this bug with this example.
In the attached file, there is also code commented out.
If removes the comments, then it shows that one can easily define 2 complex
matrices, offload them on nvptx gpu, add them together in parallel, and even
print them! on gpu!, with print commants issued on the device in a kernel.
All this works...
what does not work is reductions over complex variables,
because of a problem in the way how gcc treats independend reductions over
array members...
A workaround for reductions over complex numbers, or any other variable that is
an array of fixed size is, until now, to make two loops, one for ar[0] and the
other for ar[1]...
Especially, if the array is larger than 2, this is highly cumbersome.
Using atomics would slow the system down severely here.
I would like this to be fixed... Especially since clang does not have this
problem.
Best regards,
Benjamin
