Re: [PATCH] omp-low: fix lastprivate/linear lowering for SIMT

[Alexander Monakov <amonakov@ispras.ru>]

Ping - as this patch addresses a wrong-code issue in new functionality, I'd like
to ask if it may be applied to gcc-7 branch too.

On Fri, 7 Apr 2017, Alexander Monakov wrote:

> Ping.
> 
> > 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.
> > 
> > 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?),
> > 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.
> > 
> > Thanks.
> > Alexander
> > 	
> > diff --git a/gcc/omp-low.c b/gcc/omp-low.c
> > index 253dc85..02b681c 100644
> > --- a/gcc/omp-low.c
> > +++ b/gcc/omp-low.c
> > @@ -4768,11 +4768,10 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
> >  		TREE_NO_WARNING (new_var) = 1;
> >  	    }
> >  
> > -	  if (simduid && DECL_HAS_VALUE_EXPR_P (new_var))
> > +	  if (!maybe_simt && simduid && DECL_HAS_VALUE_EXPR_P (new_var))
> >  	    {
> >  	      tree val = DECL_VALUE_EXPR (new_var);
> > -	      if (!maybe_simt
> > -		  && TREE_CODE (val) == ARRAY_REF
> > +	      if (TREE_CODE (val) == ARRAY_REF
> >  		  && VAR_P (TREE_OPERAND (val, 0))
> >  		  && lookup_attribute ("omp simd array",
> >  				       DECL_ATTRIBUTES (TREE_OPERAND (val,
> > @@ -4792,26 +4791,26 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
> >  				    TREE_OPERAND (val, 0), lastlane,
> >  				    NULL_TREE, NULL_TREE);
> >  		}
> > -	      else if (maybe_simt
> > -		       && VAR_P (val)
> > -		       && lookup_attribute ("omp simt private",
> > -					    DECL_ATTRIBUTES (val)))
> > +	    }
> > +	  else if (maybe_simt)
> > +	    {
> > +	      tree val = (DECL_HAS_VALUE_EXPR_P (new_var)
> > +			  ? DECL_VALUE_EXPR (new_var)
> > +			  : new_var);
> > +	      if (simtlast == NULL)
> >  		{
> > -		  if (simtlast == NULL)
> > -		    {
> > -		      simtlast = create_tmp_var (unsigned_type_node);
> > -		      gcall *g = gimple_build_call_internal
> > -			(IFN_GOMP_SIMT_LAST_LANE, 1, simtcond);
> > -		      gimple_call_set_lhs (g, simtlast);
> > -		      gimple_seq_add_stmt (stmt_list, g);
> > -		    }
> > -		  x = build_call_expr_internal_loc
> > -		    (UNKNOWN_LOCATION, IFN_GOMP_SIMT_XCHG_IDX,
> > -		     TREE_TYPE (val), 2, val, simtlast);
> > -		  new_var = unshare_expr (new_var);
> > -		  gimplify_assign (new_var, x, stmt_list);
> > -		  new_var = unshare_expr (new_var);
> > +		  simtlast = create_tmp_var (unsigned_type_node);
> > +		  gcall *g = gimple_build_call_internal
> > +		    (IFN_GOMP_SIMT_LAST_LANE, 1, simtcond);
> > +		  gimple_call_set_lhs (g, simtlast);
> > +		  gimple_seq_add_stmt (stmt_list, g);
> >  		}
> > +	      x = build_call_expr_internal_loc
> > +		(UNKNOWN_LOCATION, IFN_GOMP_SIMT_XCHG_IDX,
> > +		 TREE_TYPE (val), 2, val, simtlast);
> > +	      new_var = unshare_expr (new_var);
> > +	      gimplify_assign (new_var, x, stmt_list);
> > +	      new_var = unshare_expr (new_var);
> >  	    }
> >  
> >  	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE
> > diff --git a/libgomp/testsuite/libgomp.c/target-36.c b/libgomp/testsuite/libgomp.c/target-36.c
> > new file mode 100644
> > index 0000000..6925a2a
> > --- /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 ();
> > +}
> > 
>