On Fri, Mar 31, 2017 at 06:38:09PM +0300, Alexander Monakov wrote: > I've noticed while re-reading that this patch incorrectly handled SIMT case > in lower_lastprivate_clauses. The code was changed to look for variables > with "omp simt private" attribute, and was left under > 'simduid && DECL_HAS_VALUE_EXPR_P (new_var)' condition. This effectively > constrained processing to privatized (i.e. addressable) variables, as > non-addressable variables now have neither the value-expr nor the attribute.
Sorry for the review delay. > This wasn't caught in testing, as apparently all testcases that have target > simd loops with linear/lastprivate clauses also have the corresponding > variables > mentioned in target map clause, which makes them addressable (is that > necessary?), Yes, in order to map something you need to map its address (+ size) on the host to its address on the device, so it needs to be addressable. Compared to that, firstprivate on target should not make it addressable. > and I didn't think to check something like that manually. > > The following patch fixes it by adjusting the if's in > lower_lastprivate_clauses; > alternatively it may be possible to keep that code as-is, and instead set up > value-expr and "omp simt private" attributes for all formally private > variables, > not just addressable ones, but to me that sounds less preferable. OK for > trunk? > > I think it's possible to improve test coverage for NVPTX by running all OpenMP > testcases via nvptx-none-run (it'll need some changes, but shouldn't be hard). > > gcc/ > * omp-low.c (lower_lastprivate_clauses): Correct handling of linear and > lastprivate clauses in SIMT case. > > libgomp/ > * testsuite/libgomp.c/target-36.c: New testcase. Ok for trunk/gcc-7-branch, thanks. > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c/target-36.c > @@ -0,0 +1,18 @@ > +int > +main () > +{ > + int ah, bh, n = 1024; > +#pragma omp target map(from: ah, bh) > + { > + int a, b; > +#pragma omp simd lastprivate(b) > + for (a = 0; a < n; a++) > + { > + b = a + n + 1; > + asm volatile ("" : "+r"(b)); > + } > + ah = a, bh = b; > + } > + if (ah != n || bh != 2 * n) > + __builtin_abort (); > +} Would be nice to also test explicit linear, perhaps in the same testcase, just add ch and c and use say linear(c:2). Jakub