OpenACC wait clause

On 06/07/2016 08:02 AM, Jakub Jelinek wrote:
> On Tue, Jun 07, 2016 at 08:01:10AM -0700, Cesar Philippidis wrote:
>> On 06/07/2016 04:13 AM, Jakub Jelinek wrote:
>>
>>> I've noticed
>>>           if ((mask & OMP_CLAUSE_WAIT)
>>>               && !c->wait
>>>               && gfc_match ("wait") == MATCH_YES)
>>>             {
>>>               c->wait = true; 
>>>               match_oacc_expr_list (" (", &c->wait_list, false);
>>>               continue;
>>>             }
>>> which looks just weird and confusing.  Why isn't this instead:
>>>           if ((mask & OMP_CLAUSE_WAIT)
>>>               && !c->wait
>>> 	      && (match_oacc_expr_list ("wait (", &c->wait_list, false)
>>> 		  == MATCH_YES))
>>>             {
>>>               c->wait = true; 
>>>               continue;
>>>             }
>>> ?  Otherwise you happily accept wait without following (, perhaps even
>>> combined with another clause without any space in between etc.
>>
>> Both acc wait and async accept optional parenthesis arguments. E.g.,
>>
>>   #pragma acc wait
>>
>> blocks for all of the async streams to complete before proceeding, whereas
>>
>>   #pragma acc wait (1, 5)
>>
>> only blocks for async streams 1 and 5.
> 
> But then you need to set need_space = true; if it doesn't have the ( after
> it.

I was distracted with acc routine stuff, so this took me a little longer
to get around to this. In addition to that problem with the wait clause,
I discovered a similar problem with the async clause and the wait
directive. It this patch ok for trunk and gcc-6?

Cesar

On Thu, Jun 16, 2016 at 08:22:29PM -0700, Cesar Philippidis wrote:
> --- a/gcc/fortran/openmp.c
> +++ b/gcc/fortran/openmp.c
> @@ -677,7 +677,6 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>  	      && gfc_match ("async") == MATCH_YES)
>  	    {
>  	      c->async = true;
> -	      needs_space = false;
>  	      if (gfc_match (" ( %e )", &c->async_expr) != MATCH_YES)
>  		{
>  		  c->async_expr
> @@ -685,6 +684,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>  					     gfc_default_integer_kind,
>  					     &gfc_current_locus);
>  		  mpz_set_si (c->async_expr->value.integer, GOMP_ASYNC_NOVAL);
> +		  needs_space = true;
>  		}
>  	      continue;
>  	    }
> @@ -1328,7 +1328,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>  	      && gfc_match ("wait") == MATCH_YES)
>  	    {
>  	      c->wait = true;
> -	      match_oacc_expr_list (" (", &c->wait_list, false);
> +	      if (match_oacc_expr_list (" (", &c->wait_list, false) == MATCH_NO)
> +		needs_space = true;
>  	      continue;
>  	    }
>  	  if ((mask & OMP_CLAUSE_WORKER)

I think it is still problematic.  Most of the parsing fortran FE errors are deferred,
meaning that if you don't reject the whole gfc_match_omp_clauses, then no
diagnostics is actually emitted.  Both
gfc_match (" ( %e )", &c->async_expr) and match_oacc_expr_list (" (", &c->wait_list, false)
IMHO can return MATCH_YES, MATCH_NO and MATCH_ERROR, and I believe you need
to do different actions in each case.
In particular, if something is optional, then for MATCH_YES you should
accept it (continue) and not set needs_space, because after ) you don't need
space.  If MATCH_NO, then you should accept it too (because it is optional),
and set needs_space = true; first and perhaps do whatever else you need to
do.  If MATCH_ERROR, then you should make sure not to accept it, e.g. by
doing break; or making sure continue will not be done (which one depends on
whether it might be validly parsed as some other clause, which is very
likely not the case).  In the above changes, you do it all except for the
MATCH_ERROR handling, where you still do continue; and thus I bet
diagnostics for it won't be reported.
E.g. for
!$omp acc parallel async(&abc)
!$omp acc end parallel
end
no diagnostics is reported.  Looking around, there are many more issues like
that, e.g. match_oacc_clause_gang(c) (note, wrong formatting) also ignores
MATCH_ERROR, etc.

> @@ -1649,7 +1650,7 @@ gfc_match_oacc_wait (void)
>    gfc_expr_list *wait_list = NULL, *el;
>  
>    match_oacc_expr_list (" (", &wait_list, true);
> -  gfc_match_omp_clauses (&c, OACC_WAIT_CLAUSES, false, false, true);
> +  gfc_match_omp_clauses (&c, OACC_WAIT_CLAUSES, false, true, true);
>  
>    if (gfc_match_omp_eos () != MATCH_YES)
>      {

Can you explain this change?  I bet it again suffers from the above
mentioned issue.  If match_oacc_expr_list returns MATCH_YES, I believe you
want false, false, true as you don't need space in between the closing
) of the wait_list and name of next clause.  Note, does OpenACC allow also comma
in that case?
!$acc wait (whatever),async
?
If match_oacc_expr_list returns MATCH_NO, then IMHO it should be
true, true, true, because you don't want to accept
!$acc waitasync
and also don't want to accept
!$acc wait,async
And if match_oacc_expr_list returns MATCH_ERROR, you should reject it, so
that diagnostics is emitted.

	Jakub

On 06/17/2016 07:34 AM, Jakub Jelinek wrote:
> On Thu, Jun 16, 2016 at 08:22:29PM -0700, Cesar Philippidis wrote:
>> --- a/gcc/fortran/openmp.c
>> +++ b/gcc/fortran/openmp.c
>> @@ -677,7 +677,6 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>>  	      && gfc_match ("async") == MATCH_YES)
>>  	    {
>>  	      c->async = true;
>> -	      needs_space = false;
>>  	      if (gfc_match (" ( %e )", &c->async_expr) != MATCH_YES)
>>  		{
>>  		  c->async_expr
>> @@ -685,6 +684,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>>  					     gfc_default_integer_kind,
>>  					     &gfc_current_locus);
>>  		  mpz_set_si (c->async_expr->value.integer, GOMP_ASYNC_NOVAL);
>> +		  needs_space = true;
>>  		}
>>  	      continue;
>>  	    }
>> @@ -1328,7 +1328,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>>  	      && gfc_match ("wait") == MATCH_YES)
>>  	    {
>>  	      c->wait = true;
>> -	      match_oacc_expr_list (" (", &c->wait_list, false);
>> +	      if (match_oacc_expr_list (" (", &c->wait_list, false) == MATCH_NO)
>> +		needs_space = true;
>>  	      continue;
>>  	    }
>>  	  if ((mask & OMP_CLAUSE_WORKER)
> 
> I think it is still problematic.  Most of the parsing fortran FE errors are deferred,
> meaning that if you don't reject the whole gfc_match_omp_clauses, then no
> diagnostics is actually emitted.

What exactly is the problem here? Do you want more precise errors or do
you want to keep the errors generic and deferred?

E.g. Should

  !$acc loop gang(bogus_arg)

report that bogus_arg is an invalid gang argument, or would you rather
see this being reported as an "Unclassifiable OpenACC directive"? I'm
not sure if it's intentional, but the unclassifiable omp/acc directive
error won't show of if there were any error messages were deferred.
Maybe that error message should be unconditional. Also, if you want to
go with more precise errors, should gfc_match_omp_clauses abort early as
soon as it detect the first bogus clause, or should it try to work
through the errors?

I'd personally favor making error messages more precise like the c and
c++ FEs, but everything mostly works as-is right now in the fortran FE.

One thing slightly unrelated is the annoying "Error: Unexpected
!${OMP/ACC} END PARALLEL statement at (1)" message if there is a bogus
omp/acc directive. Should we leave this alone or create a dummy omp/acc
region so that this error gets ignored?

>  Both
> gfc_match (" ( %e )", &c->async_expr) and match_oacc_expr_list (" (", &c->wait_list, false)
> IMHO can return MATCH_YES, MATCH_NO and MATCH_ERROR, and I believe you need
> to do different actions in each case.
> In particular, if something is optional, then for MATCH_YES you should
> accept it (continue) and not set needs_space, because after ) you don't need
> space.  If MATCH_NO, then you should accept it too (because it is optional),
> and set needs_space = true; first and perhaps do whatever else you need to
> do.  If MATCH_ERROR, then you should make sure not to accept it, e.g. by
> doing break; or making sure continue will not be done (which one depends on
> whether it might be validly parsed as some other clause, which is very
> likely not the case).  In the above changes, you do it all except for the
> MATCH_ERROR handling, where you still do continue; and thus I bet
> diagnostics for it won't be reported.
> E.g. for
> !$omp acc parallel async(&abc)
> !$omp acc end parallel
> end
> no diagnostics is reported.  Looking around, there are many more issues like
> that, e.g. match_oacc_clause_gang(c) (note, wrong formatting) also ignores
> MATCH_ERROR, etc.

You're example does gets flagged with an invalid name error. Also, in
the case of, say, gang(bogus_arg), in the first pass,
gfc_match_omp_clauses would detect gang (by itself) as a valid clause.
Then in the second pass, it would detect (bogus_arg) as an invalid
clause. Therefore, I'd still expect this directive to be reported as an
unclassifiable directive.

>> @@ -1649,7 +1650,7 @@ gfc_match_oacc_wait (void)
>>    gfc_expr_list *wait_list = NULL, *el;
>>  
>>    match_oacc_expr_list (" (", &wait_list, true);
>> -  gfc_match_omp_clauses (&c, OACC_WAIT_CLAUSES, false, false, true);
>> +  gfc_match_omp_clauses (&c, OACC_WAIT_CLAUSES, false, true, true);
>>  
>>    if (gfc_match_omp_eos () != MATCH_YES)
>>      {
> 
> Can you explain this change?  I bet it again suffers from the above
> mentioned issue. If match_oacc_expr_list returns MATCH_YES, I believe you
> want false, false, true as you don't need space in between the closing
> ) of the wait_list and name of next clause.  Note, does OpenACC allow also comma
> in that case?
> !$acc wait (whatever),async
> ?
>
> If match_oacc_expr_list returns MATCH_NO, then IMHO it should be
> true, true, true, because you don't want to accept
> !$acc waitasync
> and also don't want to accept
> !$acc wait,async
> And if match_oacc_expr_list returns MATCH_ERROR, you should reject it, so
> that diagnostics is emitted.

Yeah, it should be true, true, true. I overlooked the comma search. I
fix that once I get feedback on how you want the errors to get reported.

Cesar

On Fri, Jun 24, 2016 at 08:42:49AM -0700, Cesar Philippidis wrote:
> >> @@ -1328,7 +1328,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
> >>  	      && gfc_match ("wait") == MATCH_YES)
> >>  	    {
> >>  	      c->wait = true;
> >> -	      match_oacc_expr_list (" (", &c->wait_list, false);
> >> +	      if (match_oacc_expr_list (" (", &c->wait_list, false) == MATCH_NO)
> >> +		needs_space = true;
> >>  	      continue;
> >>  	    }
> >>  	  if ((mask & OMP_CLAUSE_WORKER)
> > 
> > I think it is still problematic.  Most of the parsing fortran FE errors are deferred,
> > meaning that if you don't reject the whole gfc_match_omp_clauses, then no
> > diagnostics is actually emitted.
> 
> What exactly is the problem here? Do you want more precise errors or do
> you want to keep the errors generic and deferred?

The problem is that if you ignore MATCH_ERROR and don't reject the stmt,
then invalid source will be accepted, which is not acceptable.

I admit the Fortran OpenMP/OpenACC error recovery and diagnostics is not
very good, but it generally follows the model done elsewhere in the Fortran
FE.  So, I think before changing anything substantial first the bugs in
there (where MATCH_ERROR is ignored) should be fixed (something that can be
also backported), and only afterwards we should be discussing some plan to
improve the infrastructure (like for clauses if there are any non-fatal
errors in the parsing emit the diagnostics immediately, don't add the
clauses to the FE IL and don't reject the whole stmt).

	Jakub