https://gcc.gnu.org/bugzilla/show_bug.cgi?id=85527
Bug ID: 85527
Summary: [openacc] atomic_capture-1.{c,f90} undefined behaviour
Product: gcc
Version: unknown
Status: UNCONFIRMED
Severity: normal
Priority: P3
Component: testsuite
Assignee: unassigned at gcc dot gnu.org
Reporter: vries at gcc dot gnu.org
Target Milestone: ---
I.
Consider this test-case, minimized from atomic_capture-1.f90:
...
program main
real fgot, fexp, ftmp
integer, parameter :: N = 32
fgot = 1234.0
fexp = 1266.0
!$acc parallel loop copy (fgot, ftmp)
do i = 1, N
!$acc atomic capture
ftmp = fgot
fgot = fgot + 1.0
!$acc end atomic
end do
!$acc end parallel loop
if (ftmp /= fexp - 1.0) call abort
end program main
...
We write different values to the scalar ftmp in a parallel loop.
So, roughtly equivalent to:
...
!$acc parallel loop copy (fgot, ftmp)
do i = 1, N
ftmp = i
end do
!$acc end parallel loop
...
This is undefined behaviour.
If we look in atomic_capture-1.c, the equivalent loop looks like this:
...
#pragma acc data copy (fgot, fdata[0:N])
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
#pragma acc atomic capture
fdata[i] = fgot++;
}
}
if (fexp != fgot)
abort ();
...
So, the solution here is to write to different entries in an array.
II.
atomic_capture-1.c contains only a single case of checking the contents of
fdata/ldata/idata:
...
fgot = 1.0;
fexp = 0.0;
#pragma acc data copy (fgot, fdata[0:N])
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 32.0;
#pragma acc atomic capture
fdata[i] = fgot = expr - fgot;
}
}
for (i = 0; i < N; i++)
if (i % 2 == 0)
{
if (fdata[i] != 31.0)
abort ();
}
else
{
if (fdata[i] != 1.0)
abort ();
}
...
That was already observed to be causing problem here
(https://gcc.gnu.org/ml/gcc-patches/2016-05/msg00173.html ), and was removed in
this patch:
...
commit 39be832ac8fb62836b54b98c118ce713a323bb2e
Author: nathan <nathan@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri Jan 22 22:01:33 2016 +0000
* testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c: Don't
assume atomic op ordering.
* testsuite/libgomp.oacc-fortran/atomic_capture-1.f90: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@232749
138bc75d-0d04-0410-961f-82ee72b054a4
...
but that's still missing on trunk.
That patch also removes some of the ftmp/itmp/ltmp checks in
atomic_capture-1.f90, but not the one mentioned in I.
III.
It may still be a good idea to check the result of fdata/idata/ldata array, but
not in the order-dependent fashion that's used in the example listed at II.
