Hi Chung-Lin!
On Tue, 27 Nov 2018 22:41:54 +0800, Chung-Lin Tang <chunglin_t...@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
