On Thu, Jan 14, 2016 at 07:08:23PM +0300, Alexander Monakov wrote: > On Thu, 14 Jan 2016, Jakub Jelinek wrote: > > That is weird, because c is also a firstprivate var in target and > > (implicitly) shared in teams, so if omp lowering/expansion is not buggy, > > you should see the exactly same problem with that. > > Wonder if we use GOMP_MAP_FIRSTPRIVATE_INT when we shouldn't (the var really > > should be addressable because of the atomics and atomics would be the only > > way to observe some changes on it (or teams reduction?)). > > I see that there's stack var 'c' initialized from omp_data_i->c in > main._omp_fn.0, which makes a per-team copy of 'c'. Thus I think it's a > lowering/expansion issue. I didn't dig further. > > > That said, for the testcase, as it is num_teams(4), I think I'd prefer > > #pragma omp atomic read c, g[0] and g[1] into temporaries and verify if they > > are one of the possible values. > > Done in the below patch. > > > > Second, there are failures on libgomp.c/examples-4/teams-{3,4}.c and their > > > Fortran counterparts. The issue is that 'sum' is not reduced across all > > > teams, but only across loop iterations within each team. I'm using the > > > following patch to add the missing reduction. Is that correct? > > > > Yes, but we need to report it upstream, the bug is on the OpenMP Examples > > side. > > Let me do it now. > > Thanks! Do you want me to repost that patch with the Fortran bits updated?
I've submitted https://github.com/OpenMP/Examples/pull/17 but it hasn't been applied. If you could post a patch that matches that formatting (in the hope that it will be accepted that way), it is preapproved. > diff --git a/libgomp/testsuite/libgomp.c/target-31.c > b/libgomp/testsuite/libgomp.c/target-31.c > index 255327c..27f9410 100644 > --- a/libgomp/testsuite/libgomp.c/target-31.c > +++ b/libgomp/testsuite/libgomp.c/target-31.c > @@ -36,7 +36,18 @@ main () > u3[i] = k + i; > #pragma omp parallel num_threads (1) > { > - if (c != 3 || d != 4 || g[0] != 9 || g[1] != 10 || h[0] != 11 || h[1] > != 12 || k != 14 || m[0] != 17 || m[1] != 18) > + int v1, v2, v3; > + #pragma omp atomic read > + v1 = c; > + #pragma omp atomic read > + v2 = g[0]; > + #pragma omp atomic read > + v3 = g[1]; > + if ((v1 != 3 && v1 != 4 && v1 != 5 && v1 != 6) I'd use just (v1 < 3 || v1 > 6) > + || d != 4 > + || (v2 != 9 && v2 != 11 && v2 != 13 && v2 != 15) and (v2 < 9 || v2 > 15 || (v2 & 1) == 0) > + || (v3 != 10 && v3 != 13 && v3 != 16 && v3 != 19) and (v3 < 10 || v3 > 19 || ((v3 - 10) % 3) != 0) > + || h[0] != 11 || h[1] != 12 || k != 14 || m[0] != 17 || m[1] != 18) > #pragma omp atomic write > err = 1; > b = omp_get_team_num (); Ok for trunk with those changes, thanks. Jakub