Re: [PATCH, OpenACC] Properly handle wait clause with no arguments

[Thomas Schwinge <thomas at codesourcery dot com>]

Hi Chung-Lin!

On Tue, 27 Nov 2018 22:41:54 +0800, Chung-Lin Tang <chunglin_tang@mentor.com> wrote:
> On 2018/11/8 3:13 AM, Thomas Schwinge wrote:
> > Now, why do we need the following changes, in this rather "convoluted"
> > form:
> 
|    --- gcc/omp-expand.c	(revision 263981)
|    +++ gcc/omp-expand.c	(working copy)
|    @@ -7381,16 +7381,32 @@ expand_omp_target (struct omp_region *region)
|     	  /* ... push a placeholder.  */
|     	  args.safe_push (integer_zero_node);
|     
|    +	bool noval_seen = false;
|    +	tree noval = build_int_cst (integer_type_node, GOMP_ASYNC_NOVAL);
|    +
|     	for (; c; c = OMP_CLAUSE_CHAIN (c))
|     	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_WAIT)
|     	    {
> >> +	      tree wait_expr = OMP_CLAUSE_WAIT_EXPR (c);
> >> +
> >> +	      if (TREE_CODE (wait_expr) == INTEGER_CST
> >> +		  && tree_int_cst_compare (wait_expr, noval) == 0)
> >> +		{
> >> +		  noval_seen = true;
> >> +		  continue;
> >> +		}
> >> +
> >>   	      args.safe_push (fold_convert_loc (OMP_CLAUSE_LOCATION (c),
> >> -						integer_type_node,
> >> -						OMP_CLAUSE_WAIT_EXPR (c)));
> >> +						integer_type_node, wait_expr));
> >>   	      num_waits++;
> >>   	    }
> >>   
> >> -	if (!tagging || num_waits)
> >> +	if (noval_seen && num_waits == 0)
> >> +	  args[t_wait_idx] =
> >> +	    (tagging
> >> +	     ? oacc_launch_pack (GOMP_LAUNCH_WAIT, NULL_TREE, GOMP_ASYNC_NOVAL)
> >> +	     : noval);
> >> +	else if (!tagging || num_waits)
> >>   	  {
> >>   	    tree len;
> 
> >>   	case GOMP_LAUNCH_WAIT:
> >>   	  {
> >> -	    unsigned num_waits = GOMP_LAUNCH_OP (tag);
> >> +	    /* Be careful to cast the op field as a signed 16-bit, and
> >> +	       sign-extend to full integer.  */
> >> +	    int num_waits = ((signed short) GOMP_LAUNCH_OP (tag));
> >>   
> >> -	    if (num_waits)
> >> +	    if (num_waits > 0)
> >>   	      goacc_wait (async, num_waits, &ap);
> >> +	    else if (num_waits == acc_async_noval)
> >> +	      acc_wait_all_async (async);
> 
> > Why can't we just pass "GOMP_ASYNC_NOVAL" through like any other
> > async-argument (that is, map a single "wait" clause to "num_waits == 1,
> > *ap == GOMP_ASYNC_NOVAL"), and then handle that case in "goacc_wait",
> > avoiding all these interface changes and special casing in different
> > functions?
> > 
> > Or am I not understanding correctly what the purpose of this is?
> 
> I think the original intention was that wait(acc_async_noval) should
> correspond to "wait all" semantics, hence we should be able to ignore
> and discard other wait(<arg>) clauses if they exist.

Ah, I see.  But, I'm not sure whether an optimization for such "strange"
user code ("#pragma acc [...] wait(0, 1, acc_async_noval, 5, 0, [...])")
really warrants any such GCC code complications.  ;-)

> Having that said, I think there is some incorrect code in my patch wrt
> this intended behavior, which I'll revise.

(OK, still waiting for that.)

> (The assumption of an argument-less wait clause to mean "wait all" is
> derived from the closely documented OpenACC wait *directive* specification.
> Frankly speaking, the prior section on the wait *clause* is not explicitly
> clear on this, though 'wait all' is a reasonable assumption. It would still
> be helpful if we asked the OpenACC SC to clarify)

(We're discussing that with them, but what you describe indeed I also
would agree to be what's intended, so OK to proceed assuming that.)

> As for the idea on stuffing more code into goacc_wait(), I think that's
> a pretty good suggestion, since all uses of it in oacc-parallel.c are
> actually quite similar; re-factoring this part should make things more elegant.

ACK.

> > My understanding is that before, GCC never generates "negative
> > async-arguments" (now used for "GOMP_ASYNC_NOVAL"), but only non-negative
> > ones (real "async-arguments"), which we continue to handle, as before.
> 
> > Isn't that sufficient for the ABI compatibility that we promise, which is
> > (unless I'm confused now?) that old (existing) executables continue to
> > run correctly when dynamically linking against a new libgomp.  Or do we
> > also have to care about the case that an executable built with a new
> > version of GCC has to work when dynamically linked against an old
> > libgomp?
> 
> I think either way, encoding GOMP_ASYNC_NOVAL in num_waits or as an argument
> should be okay for backward compatibility, i.e. old binaries should still
> work with new libgomp with this modification.
> 
> As for new binaries vs old libgomp, I believe with the original libgomp
> oacc-parallel.c code, it's not quite possible to achieve the intended wait all
> behavior by playing with num_waits or arguments.
> 
> I'll revise the patch and re-submit later.

OK, thanks!


Grüße
 Thomas