Message ID | 564254FC.5030002@acm.org |
---|---|
State | New |
Headers | show |
On 11/10/2015 12:35 PM, Nathan Sidwell wrote: > I've committed this to gomp4. In preparing the reworked firstprivate > patch changes for gomp4's gimplify.c I discovered these testcases were > passing by accident, and lacked a data clause. It used to be if a reduction was on a parallel construct, the gimplifier would introduce a pcopy clause for the reduction variable if it was not associated with any data clause. Is that not the case anymore? Cesar
On 11/10/15 18:08, Cesar Philippidis wrote: > On 11/10/2015 12:35 PM, Nathan Sidwell wrote: >> I've committed this to gomp4. In preparing the reworked firstprivate >> patch changes for gomp4's gimplify.c I discovered these testcases were >> passing by accident, and lacked a data clause. > > It used to be if a reduction was on a parallel construct, the gimplifier > would introduce a pcopy clause for the reduction variable if it was not > associated with any data clause. Is that not the case anymore? AFAICT, the std doesn't specify that behaviour. 2.6 'Data Environment' doesn't mention reductions as a modifier for implicitly determined data attributes. nathan
On 11/11/2015 05:40 AM, Nathan Sidwell wrote: > On 11/10/15 18:08, Cesar Philippidis wrote: >> On 11/10/2015 12:35 PM, Nathan Sidwell wrote: >>> I've committed this to gomp4. In preparing the reworked firstprivate >>> patch changes for gomp4's gimplify.c I discovered these testcases were >>> passing by accident, and lacked a data clause. >> >> It used to be if a reduction was on a parallel construct, the gimplifier >> would introduce a pcopy clause for the reduction variable if it was not >> associated with any data clause. Is that not the case anymore? > > AFAICT, the std doesn't specify that behaviour. 2.6 'Data Environment' > doesn't mention reductions as a modifier for implicitly determined data > attributes. I guess I was confused because the reduction section in 2.5.11 mentions something about updating the original reduction variable after the parallel region. Cesar
On 11/11/15 09:50, Cesar Philippidis wrote: > On 11/11/2015 05:40 AM, Nathan Sidwell wrote: >> On 11/10/15 18:08, Cesar Philippidis wrote: >>> On 11/10/2015 12:35 PM, Nathan Sidwell wrote: >>>> I've committed this to gomp4. In preparing the reworked firstprivate >>>> patch changes for gomp4's gimplify.c I discovered these testcases were >>>> passing by accident, and lacked a data clause. >>> >>> It used to be if a reduction was on a parallel construct, the gimplifier >>> would introduce a pcopy clause for the reduction variable if it was not >>> associated with any data clause. Is that not the case anymore? >> >> AFAICT, the std doesn't specify that behaviour. 2.6 'Data Environment' >> doesn't mention reductions as a modifier for implicitly determined data >> attributes. > > I guess I was confused because the reduction section in 2.5.11 mentions > something about updating the original reduction variable after the > parallel region. I think that still relies on a copy clause to transfer the liveness of the original variable into and out of the region. (that's the implication of what 2.6 says) nathan
2015-11-10 Nathan Sidwell <nathan@codesourcery.com> * testsuite/libgomp.oacc-fortran/parallel-reduction.f90: Fix data missing data clause. * testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c: Likewise. Index: libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90 =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90 (revision 230116) +++ libgomp/testsuite/libgomp.oacc-fortran/parallel-reduction.f90 (working copy) @@ -7,7 +7,7 @@ program reduction sum = 0 - !$acc parallel reduction(+:sum) num_gangs (n) + !$acc parallel reduction(+:sum) num_gangs (n) copy(sum) sum = sum + 1 !$acc end parallel @@ -32,7 +32,7 @@ end program reduction subroutine redsub(sum, n) integer :: sum, n - !$acc parallel reduction(+:sum) num_gangs (10) + !$acc parallel reduction(+:sum) num_gangs (10) copy(sum) sum = sum + 1 !$acc end parallel end subroutine redsub Index: libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c (revision 230116) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c (working copy) @@ -14,7 +14,7 @@ main () #pragma acc data copy (dummy) { -#pragma acc parallel num_gangs (N) reduction (+:s1) +#pragma acc parallel num_gangs (N) reduction (+:s1) copy(s1) { s1++; } @@ -34,7 +34,7 @@ main () s1 = 0; s2 = 0; -#pragma acc parallel num_gangs (10) reduction (+:s1, s2) +#pragma acc parallel num_gangs (10) reduction (+:s1, s2) copy(s1, s2) { s1++; s2 += N; @@ -57,7 +57,7 @@ main () s1 = 0; -#pragma acc parallel num_gangs (10) reduction (+:s1) +#pragma acc parallel num_gangs (10) reduction (+:s1) copy(s1) { #pragma acc loop gang reduction (+:s1) for (i = 0; i < 10; i++) Index: libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c (revision 230116) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c (working copy) @@ -15,7 +15,7 @@ main (int argc, char *argv[]) # define GANGS 256 #endif #pragma acc parallel num_gangs(GANGS) num_workers(32) vector_length(32) \ - reduction(+:res1) copy(res2) + reduction(+:res1) copy(res2, res1) { res1 += 5; @@ -36,7 +36,7 @@ main (int argc, char *argv[]) # define GANGS 8 #endif #pragma acc parallel num_gangs(GANGS) num_workers(32) vector_length(32) \ - reduction(*:res1) copy(res2) + reduction(*:res1) copy(res1, res2) { res1 *= 5; Index: libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c (revision 230116) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c (working copy) @@ -13,7 +13,7 @@ main (int argc, char *argv[]) arr[i] = i; #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - reduction(+:res) + reduction(+:res) copy(res) { #pragma acc loop gang for (j = 0; j < 32; j++) Index: libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c (revision 230116) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c (working copy) @@ -14,7 +14,7 @@ main (int argc, char *argv[]) # define GANGS 256 #endif #pragma acc parallel num_gangs(GANGS) num_workers(32) vector_length(32) \ - reduction(+:res1) copy(res2) async(1) + reduction(+:res1) copy(res1, res2) async(1) { res1 += 5; @@ -37,7 +37,7 @@ main (int argc, char *argv[]) # define GANGS 8 #endif #pragma acc parallel num_gangs(GANGS) num_workers(32) vector_length(32) \ - reduction(*:res1) copy(res2) async(1) + reduction(*:res1) copy(res1, res2) async(1) { res1 *= 5; Index: libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c (revision 230116) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c (working copy) @@ -14,7 +14,7 @@ main (int argc, char *argv[]) arr[i] = i; #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - reduction(+:res) + reduction(+:res) copy(res) { #pragma acc loop gang for (j = 0; j < 32; j++) Index: libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c (revision 230116) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c (working copy) @@ -14,7 +14,7 @@ main (int argc, char *argv[]) arr[i] = i; #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - reduction(+:res) reduction(max:mres) + reduction(+:res) reduction(max:mres) copy(res, mres) { #pragma acc loop gang for (j = 0; j < 32; j++)