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