[PATCH 2/5] omp-low: implement SIMT privatization, part 1

[Jakub Jelinek]

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