https://gcc.gnu.org/bugzilla/show_bug.cgi?id=120173
Bug ID: 120173
Summary: [OpenACC][gcn-offload] wrong 'firstprivate' with 'acc
parallel async' [modified
libgomp.oacc-fortran/lib-13.f90]
Product: gcc
Version: 16.0
Status: UNCONFIRMED
Keywords: openacc, wrong-code
Severity: normal
Priority: P3
Component: fortran
Assignee: unassigned at gcc dot gnu.org
Reporter: burnus at gcc dot gnu.org
CC: ams at gcc dot gnu.org, sandra at gcc dot gnu.org,
tschwinge at gcc dot gnu.org
Target Milestone: ---
Created attachment 61366
--> https://gcc.gnu.org/bugzilla/attachment.cgi?id=61366&action=edit
Full test case (i.e. patch applied to libgomp.oacc-fortran/lib-13.f90)
The following is a modified version of
libgomp.oacc-fortran/lib-13.f90
which looks as follows (current version, might have changed by the time you
look at it):
https://gcc.gnu.org/git/?p=gcc.git;a=blob;f=libgomp/testsuite/libgomp.oacc-fortran/lib-13.f90
I decided to check the value of the 'k' variable:
k(:) = 7
...
! print *, k
if (any (k /= [(7 + N, i = 1, nprocs)])) stop 4
And that fails with GCN but not with Nvptx. Namely:
1000007 1000007 ! nvptx
vs.
7 1005444 ! gcn
7 2000007 ! gcn with '!$acc atomic update'
* * *
In particular, I do not understand why k(j==1) is not increment at all while
k(j==2) is too much incremented.
* * *
Changes to that program - in an attempt to fix the issue:
(A) Use atomic update:
!$acc atomic update
k(j) = k(j) + 1
→ Yields consistent results → 7 2000007
(B) Use explicit clauses:
!$acc parallel firstprivate(j) private(i) async (j)
→ no effect, unsurprisingly as 'firstprivate' is the default.
BTW: If I look at the dump, I see:
j.16 = j;
.omp_data_arr.13.j = &j.16;
#pragma omp target oacc_parallel private(i) firstprivate(j)
D.4787 = .omp_data_i->j;
j = *D.4787;
[I wonder whether the address of 'j.16' gets too quickly reused?]
* * *
Otherwise:
!$acc end data
should probably be moved after 'call acc_wait' to avoid races, even though
I failed to observe any issues for those.
See OG14 commit https://gcc.gnu.org/g:8998cc60c1cbbdef603090eab63e5835bf071d41
8998cc60c1c Fix libgomp.oacc-fortran/lib-13.f90 async bug
* * *
The changed testcase (added atomic update, value check for 'k', moved 'end
data'):
--- a/libgomp/testsuite/libgomp.oacc-fortran/lib-13.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/lib-13.f90
@@ -11,3 +11,3 @@ program main
- k(:) = 0
+ k(:) = 7
@@ -17,2 +17,3 @@ program main
do i = 1, N
+ !$acc atomic update
k(j) = k(j) + 1
@@ -21,7 +22,6 @@ program main
end do
- !$acc end data
call acc_wait_all_async (nprocs + 1)
-
call acc_wait (nprocs + 1)
+ !$acc end data
@@ -30,2 +30,3 @@ program main
if (acc_async_test (nprocs + 1) .neqv. .TRUE.) stop 3
+ if (any (k /= [(7 + N, i = 1, nprocs)])) stop 4
