Re: OpenACC wait clause

Re: OpenACC wait clause

[Cesar Philippidis]

[-- Attachment #1: Type: text/plain, Size: 1564 bytes --]

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


[-- Attachment #2: fortran-wait-async.diff --]
[-- Type: text/x-patch, Size: 3287 bytes --]

2016-06-16  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/fortran/
	* openmp.c (gfc_match_omp_clauses): Update use a needs_space for the
	OpenACC wait and async clauses.
	(gfc_match_oacc_wait): Ensure that there is a space when the optional
	parenthesis is missing.

	gcc/testsuite/
	* gfortran.dg/goacc/asyncwait-2.f95: Add additional test coverage.
	* gfortran.dg/goacc/asyncwait-4.f95: Likewise.

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 2c92794..435c709 100644
--- 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)
@@ -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)
     {
diff --git a/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95 b/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95
index db0ce1f..7b2ae07 100644
--- a/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95
@@ -83,6 +83,18 @@ program asyncwait
   end do
   !$acc end parallel ! { dg-error "Unexpected \\\!\\\$ACC END PARALLEL" }
 
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) waitasync ! { dg-error "Unclassifiable OpenACC directive" }
+  do i = 1, N
+     b(i) = a(i)
+  end do
+  !$acc end parallel ! { dg-error "Unexpected \\\!\\\$ACC END PARALLEL" }
+
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) asyncwait ! { dg-error "Unclassifiable OpenACC directive" }
+  do i = 1, N
+     b(i) = a(i)
+  end do
+  !$acc end parallel ! { dg-error "Unexpected \\\!\\\$ACC END PARALLEL" }
+  
   !$acc parallel copyin (a(1:N)) copy (b(1:N)) wait
   do i = 1, N
      b(i) = a(i)
diff --git a/gcc/testsuite/gfortran.dg/goacc/asyncwait-4.f95 b/gcc/testsuite/gfortran.dg/goacc/asyncwait-4.f95
index cd64ef3..01349b0 100644
--- a/gcc/testsuite/gfortran.dg/goacc/asyncwait-4.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/asyncwait-4.f95
@@ -34,4 +34,6 @@ program asyncwait
   !$acc wait async (1.0) ! { dg-error "ASYNC clause at \\\(1\\\) requires a scalar INTEGER expression" }
 
   !$acc wait async 1 ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+
+  !$acc waitasync ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
 end program asyncwait

Re: OpenACC wait clause

[Jakub Jelinek]

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

Re: OpenACC wait clause

[Cesar Philippidis]

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

Re: OpenACC wait clause

[Jakub Jelinek]

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

Re: OpenACC wait clause

[Cesar Philippidis]

[-- Attachment #1: Type: text/plain, Size: 2788 bytes --]

On 06/24/2016 08:53 AM, Jakub Jelinek wrote:
> 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).

That's reasonable. I wasn't considering the gcc6 backport.

Improving the infrastructure is definitely a good ideal. I'll need to
clear that out with Thomas though. As a heads up, we've been getting a
lot of customer feedback to make the compiler emit more performance
diagnostics in general. Basically, they want to know what type of
parallelization the compiler detected and applied (e.g., which loops are
independent and were vectorized/partitioned across gang, workers and
vectors), along with missed optimizations.

With that aside, this patch should correct the issues you pointed out.
gfc_match_omp_clauses now uses a common function to parse the gang,
worker and vector clauses. Also, this patch takes extra care with
MATCH_ERRORs when parsing the wait and async clauses and the wait
directive. One oddity I noticed is how the fortran FE accepts directives
with clauses specified like this

  !$omp parallel if(.true.)private(x)

Shouldn't there be some sort of separator between if and private? I know
that it's easy to distinguish the clauses because of the parenthesis,
but it still looks weird.

With that oddity aside, is this patch OK for trunk and gcc6?

Cesar

[-- Attachment #2: fortran-parser-errors-20160627.diff --]
[-- Type: text/x-patch, Size: 13030 bytes --]

2016-06-27  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/fortran/
	* openmp.c (match_oacc_clause_gang): Rename to ...
	(match_oacc_clause_gwv): this.  Add support for OpenACC worker and
	vector clauses.
	(gfc_match_omp_clauses): Use match_oacc_clause_gwv for
	OMP_CLAUSE_{GANG,WORKER,VECTOR}.  Propagate any MATCH_ERRORs for
	invalid OMP_CLAUSE_{ASYNC,WAIT,GANG,WORKER,VECTOR} clauses.
	(gfc_match_oacc_wait): Propagate MATCH_ERROR for invalid
	oacc_expr_lists.  Adjust the first and needs_space arguments to
	gfc_match_omp_clauses.

	gcc/testsuite/
	* gfortran.dg/goacc/asyncwait-2.f95: Updated expected diagnostics.
	* gfortran.dg/goacc/asyncwait-3.f95: Likewise.
	* gfortran.dg/goacc/asyncwait-4.f95: Add test coverage.

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index f514866..fa6587e 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -396,43 +396,66 @@ cleanup:
 }
 
 static match
-match_oacc_clause_gang (gfc_omp_clauses *cp)
+match_oacc_clause_gwv (gfc_omp_clauses *cp, unsigned gwv)
 {
   match ret = MATCH_YES;
 
   if (gfc_match (" ( ") != MATCH_YES)
     return MATCH_NO;
 
-  /* The gang clause accepts two optional arguments, num and static.
-     The num argument may either be explicit (num: <val>) or
-     implicit without (<val> without num:).  */
-
-  while (ret == MATCH_YES)
+  if (gwv == GOMP_DIM_GANG)
     {
-      if (gfc_match (" static :") == MATCH_YES)
+        /* The gang clause accepts two optional arguments, num and static.
+	 The num argument may either be explicit (num: <val>) or
+	 implicit without (<val> without num:).  */
+
+      while (ret == MATCH_YES)
 	{
-	  if (cp->gang_static)
-	    return MATCH_ERROR;
+	  if (gfc_match (" static :") == MATCH_YES)
+	    {
+	      if (cp->gang_static)
+		return MATCH_ERROR;
+	      else
+		cp->gang_static = true;
+	      if (gfc_match_char ('*') == MATCH_YES)
+		cp->gang_static_expr = NULL;
+	      else if (gfc_match (" %e ", &cp->gang_static_expr) != MATCH_YES)
+		return MATCH_ERROR;
+	    }
 	  else
-	    cp->gang_static = true;
-	  if (gfc_match_char ('*') == MATCH_YES)
-	    cp->gang_static_expr = NULL;
-	  else if (gfc_match (" %e ", &cp->gang_static_expr) != MATCH_YES)
-	    return MATCH_ERROR;
-	}
-      else
-	{
-	  /* This is optional.  */
-	  if (cp->gang_num_expr || gfc_match (" num :") == MATCH_ERROR)
-	    return MATCH_ERROR;
-	  else if (gfc_match (" %e ", &cp->gang_num_expr) != MATCH_YES)
-	    return MATCH_ERROR;
+	    {
+	      /* This is optional.  */
+	      if (cp->gang_num_expr || gfc_match (" num :") == MATCH_ERROR)
+		return MATCH_ERROR;
+	      else if (gfc_match (" %e ", &cp->gang_num_expr) != MATCH_YES)
+		return MATCH_ERROR;
+	    }
+
+	  ret = gfc_match (" , ");
 	}
+    }
+  else if (gwv == GOMP_DIM_WORKER)
+    {
+      /* The 'num' arugment is optional.  */
+      if (gfc_match (" num :") == MATCH_ERROR)
+	return MATCH_ERROR;
 
-      ret = gfc_match (" , ");
+      if (gfc_match (" %e ", &cp->worker_expr) != MATCH_YES)
+	return MATCH_ERROR;
     }
+  else if (gwv == GOMP_DIM_VECTOR)
+    {
+      /* The 'length' arugment is optional.  */
+      if (gfc_match (" length :") == MATCH_ERROR)
+	return MATCH_ERROR;
 
-  return gfc_match (" ) ");
+      if (gfc_match (" %e ", &cp->vector_expr) != MATCH_YES)
+	return MATCH_ERROR;
+    }
+  else
+    gfc_fatal_error ("Unexpected OpenACC parallelism.");
+
+  return gfc_match (" )");
 }
 
 static match
@@ -630,9 +653,10 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
 {
   gfc_omp_clauses *c = gfc_get_omp_clauses ();
   locus old_loc;
+  bool seen_error = false;
 
   *cp = NULL;
-  while (1)
+  while (!seen_error)
     {
       if ((first || gfc_match_char (',') != MATCH_YES)
 	  && (needs_space && gfc_match_space () != MATCH_YES))
@@ -677,14 +701,20 @@ 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)
+	      match m = gfc_match (" ( %e )", &c->async_expr);
+	      if (m == MATCH_ERROR)
+		{
+		  seen_error = true;
+		  break;
+		}
+	      else if (m == MATCH_NO)
 		{
 		  c->async_expr
 		    = gfc_get_constant_expr (BT_INTEGER,
 					     gfc_default_integer_kind,
 					     &gfc_current_locus);
 		  mpz_set_si (c->async_expr->value.integer, GOMP_ASYNC_NOVAL);
+		  needs_space = true;
 		}
 	      continue;
 	    }
@@ -877,9 +907,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
 	      && gfc_match ("gang") == MATCH_YES)
 	    {
 	      c->gang = true;
-	      if (match_oacc_clause_gang(c) == MATCH_YES)
-		needs_space = false;
-	      else
+	      match m = match_oacc_clause_gwv(c, GOMP_DIM_GANG);
+	      if (m == MATCH_ERROR)
+		{
+		  seen_error = true;
+		  break;
+		}
+	      else if (m == MATCH_NO)
 		needs_space = true;
 	      continue;
 	    }
@@ -1275,9 +1309,16 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
 	    continue;
 	  if ((mask & OMP_CLAUSE_TILE)
 	      && !c->tile_list
-	      && match_oacc_expr_list ("tile (", &c->tile_list,
-				       true) == MATCH_YES)
-	    continue;
+	      && gfc_match ("tile") == MATCH_YES)
+	    {
+	      if (match_oacc_expr_list (" (", &c->tile_list, true) != MATCH_YES)
+		{
+		  seen_error = true;
+		  break;
+		}
+	      needs_space = true;
+	      continue;
+	    }
 	  if ((mask & OMP_CLAUSE_TO)
 	      && gfc_match_omp_variable_list ("to (",
 					      &c->lists[OMP_LIST_TO], false,
@@ -1309,10 +1350,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
 	      && gfc_match ("vector") == MATCH_YES)
 	    {
 	      c->vector = true;
-	      if (gfc_match (" ( length : %e )", &c->vector_expr) == MATCH_YES
-		  || gfc_match (" ( %e )", &c->vector_expr) == MATCH_YES)
-		needs_space = false;
-	      else
+	      match m = match_oacc_clause_gwv(c, GOMP_DIM_VECTOR);
+	      if (m == MATCH_ERROR)
+		{
+		  seen_error = true;
+		  break;
+		}
+	      if (m == MATCH_NO)
 		needs_space = true;
 	      continue;
 	    }
@@ -1328,7 +1372,14 @@ 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);
+	      match m = match_oacc_expr_list (" (", &c->wait_list, false);
+	      if (m == MATCH_ERROR)
+		{
+		  seen_error = true;
+		  break;
+		}
+	      else if (m == MATCH_NO)
+		needs_space = true;
 	      continue;
 	    }
 	  if ((mask & OMP_CLAUSE_WORKER)
@@ -1336,10 +1387,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
 	      && gfc_match ("worker") == MATCH_YES)
 	    {
 	      c->worker = true;
-	      if (gfc_match (" ( num : %e )", &c->worker_expr) == MATCH_YES
-		  || gfc_match (" ( %e )", &c->worker_expr) == MATCH_YES)
-		needs_space = false;
-	      else
+	      match m = match_oacc_clause_gwv(c, GOMP_DIM_WORKER);
+	      if (m == MATCH_ERROR)
+		{
+		  seen_error = true;
+		  break;
+		}
+	      else if (m == MATCH_NO)
 		needs_space = true;
 	      continue;
 	    }
@@ -1348,7 +1402,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
       break;
     }
 
-  if (gfc_match_omp_eos () != MATCH_YES)
+  if (seen_error || gfc_match_omp_eos () != MATCH_YES)
     {
       gfc_free_omp_clauses (c);
       return MATCH_ERROR;
@@ -1595,15 +1649,18 @@ gfc_match_oacc_wait (void)
 {
   gfc_omp_clauses *c = gfc_get_omp_clauses ();
   gfc_expr_list *wait_list = NULL, *el;
+  bool space = true;
+  match m;
 
-  match_oacc_expr_list (" (", &wait_list, true);
-  gfc_match_omp_clauses (&c, OACC_WAIT_CLAUSES, false, false, true);
+  m = match_oacc_expr_list (" (", &wait_list, true);
+  if (m == MATCH_ERROR)
+    return m;
+  else if (m == MATCH_YES)
+    space = false;
 
-  if (gfc_match_omp_eos () != MATCH_YES)
-    {
-      gfc_error ("Unexpected junk in !$ACC WAIT at %C");
-      return MATCH_ERROR;
-    }
+  if (gfc_match_omp_clauses (&c, OACC_WAIT_CLAUSES, space, space, true)
+      == MATCH_ERROR)
+    return MATCH_ERROR;
 
   if (wait_list)
     for (el = wait_list; el; el = el->next)
diff --git a/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95 b/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95
index db0ce1f..7b2ae07 100644
--- a/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95
@@ -83,6 +83,18 @@ program asyncwait
   end do
   !$acc end parallel ! { dg-error "Unexpected \\\!\\\$ACC END PARALLEL" }
 
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) waitasync ! { dg-error "Unclassifiable OpenACC directive" }
+  do i = 1, N
+     b(i) = a(i)
+  end do
+  !$acc end parallel ! { dg-error "Unexpected \\\!\\\$ACC END PARALLEL" }
+
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) asyncwait ! { dg-error "Unclassifiable OpenACC directive" }
+  do i = 1, N
+     b(i) = a(i)
+  end do
+  !$acc end parallel ! { dg-error "Unexpected \\\!\\\$ACC END PARALLEL" }
+  
   !$acc parallel copyin (a(1:N)) copy (b(1:N)) wait
   do i = 1, N
      b(i) = a(i)
diff --git a/gcc/testsuite/gfortran.dg/goacc/asyncwait-3.f95 b/gcc/testsuite/gfortran.dg/goacc/asyncwait-3.f95
index 32c11de..ed72a9b 100644
--- a/gcc/testsuite/gfortran.dg/goacc/asyncwait-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/asyncwait-3.f95
@@ -11,17 +11,17 @@ program asyncwait
   a(:) = 3.0
   b(:) = 0.0
 
-  !$acc wait (1 2) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait (1 2) ! { dg-error "Syntax error in OpenACC expression list at" }
 
-  !$acc wait (1,) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait (1,) ! { dg-error "Syntax error in OpenACC expression list at" }
 
-  !$acc wait (,1) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait (,1) ! { dg-error "Syntax error in OpenACC expression list at" }
 
-  !$acc wait (1, 2, ) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait (1, 2, ) ! { dg-error "Syntax error in OpenACC expression list at" }
 
-  !$acc wait (1, 2, ,) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait (1, 2, ,) ! { dg-error "Syntax error in OpenACC expression list at" }
 
-  !$acc wait (1 ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait (1 ! { dg-error "Syntax error in OpenACC expression list at" }
 
   !$acc wait (1, *) ! { dg-error "Invalid argument to \\\$\\\!ACC WAIT" }
 
@@ -33,9 +33,9 @@ program asyncwait
 
   !$acc wait (1.0) ! { dg-error "WAIT clause at \\\(1\\\) requires a scalar INTEGER expression" }
 
-  !$acc wait 1 ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait 1 ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait N ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait N ! { dg-error "Unclassifiable OpenACC directive" }
 
   !$acc wait (1)
 end program asyncwait
diff --git a/gcc/testsuite/gfortran.dg/goacc/asyncwait-4.f95 b/gcc/testsuite/gfortran.dg/goacc/asyncwait-4.f95
index cd64ef3..df31154 100644
--- a/gcc/testsuite/gfortran.dg/goacc/asyncwait-4.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/asyncwait-4.f95
@@ -11,21 +11,21 @@ program asyncwait
   a(:) = 3.0
   b(:) = 0.0
 
-  !$acc wait async (1 2) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1 2) ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait async (1,) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1,) ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait async (,1) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (,1) ! { dg-error "Invalid character in name" }
 
-  !$acc wait async (1, 2, ) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1, 2, ) ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait async (1, 2, ,) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1, 2, ,) ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait async (1 ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1 ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait async (1, *) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1, *) ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait async (1, a) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1, a) ! { dg-error "Unclassifiable OpenACC directive" }
 
   !$acc wait async (a) ! { dg-error "ASYNC clause at \\\(1\\\) requires a scalar INTEGER expression" }
 
@@ -33,5 +33,9 @@ program asyncwait
 
   !$acc wait async (1.0) ! { dg-error "ASYNC clause at \\\(1\\\) requires a scalar INTEGER expression" }
 
-  !$acc wait async 1 ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async 1 ! { dg-error "Unclassifiable OpenACC directive" }
+
+  !$acc waitasync ! { dg-error "Unclassifiable OpenACC directive" }
+
+  !$acc wait,async ! { dg-error "Unclassifiable OpenACC directive" }
 end program asyncwait

Re: OpenACC wait clause

[Jakub Jelinek]

On Mon, Jun 27, 2016 at 11:36:26AM -0700, Cesar Philippidis wrote:
> With that aside, this patch should correct the issues you pointed out.
> gfc_match_omp_clauses now uses a common function to parse the gang,
> worker and vector clauses. Also, this patch takes extra care with
> MATCH_ERRORs when parsing the wait and async clauses and the wait
> directive. One oddity I noticed is how the fortran FE accepts directives
> with clauses specified like this
> 
>   !$omp parallel if(.true.)private(x)

That is completely intentional, in that case the whitespace isn't needed to
separate tokens.  It is only a matter of coding style then.
Somebody will only use
!$omp parallel if(.true.) private(x)
other people
!$omp parallel if ( .true. ) , private ( x )
etc.

> Shouldn't there be some sort of separator between if and private? I know
> that it's easy to distinguish the clauses because of the parenthesis,
> but it still looks weird.

If user wants, it can, but it isn't required by the standard.
In C you can also write if(cond)if(cond2)if(cond3){foo(x,y,z);}else{bar(x);}
if you want, similarly Fortran.  OpenMP/Fortran even allows
!$omp endparalleldosimd
and similar, you don't have to use spaces in between end parallel do simd.
For fixed form of course even more whitespace is optional.

> @@ -630,9 +653,10 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>  {
>    gfc_omp_clauses *c = gfc_get_omp_clauses ();
>    locus old_loc;
> +  bool seen_error = false;
>  
>    *cp = NULL;
> -  while (1)
> +  while (!seen_error)
>      {
>        if ((first || gfc_match_char (',') != MATCH_YES)
>  	  && (needs_space && gfc_match_space () != MATCH_YES))

Why?
The main loop is while (1) which has break; as the last statement.
Instead of setting seen_error, just set
gfc_current_locus = old_loc;
if you have already matched successfully something, and then break;
as you already do.

> @@ -1275,9 +1309,16 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>  	    continue;
>  	  if ((mask & OMP_CLAUSE_TILE)
>  	      && !c->tile_list
> -	      && match_oacc_expr_list ("tile (", &c->tile_list,
> -				       true) == MATCH_YES)
> -	    continue;
> +	      && gfc_match ("tile") == MATCH_YES)
> +	    {
> +	      if (match_oacc_expr_list (" (", &c->tile_list, true) != MATCH_YES)
> +		{
> +		  seen_error = true;
> +		  break;
> +		}
> +	      needs_space = true;
> +	      continue;
> +	    }
>  	  if ((mask & OMP_CLAUSE_TO)
>  	      && gfc_match_omp_variable_list ("to (",
>  					      &c->lists[OMP_LIST_TO], false,

So, tile without ()s is also a valid clause in OpenACC?
If yes, what do you set in c structure for the existence of the clause?
If it is not valid, then the above change looks wrong.

> @@ -1309,10 +1350,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>  	      && gfc_match ("vector") == MATCH_YES)
>  	    {
>  	      c->vector = true;
> -	      if (gfc_match (" ( length : %e )", &c->vector_expr) == MATCH_YES
> -		  || gfc_match (" ( %e )", &c->vector_expr) == MATCH_YES)
> -		needs_space = false;
> -	      else
> +	      match m = match_oacc_clause_gwv(c, GOMP_DIM_VECTOR);

Formatting, space before (.

> @@ -1348,7 +1402,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>        break;
>      }
>  
> -  if (gfc_match_omp_eos () != MATCH_YES)
> +  if (seen_error || gfc_match_omp_eos () != MATCH_YES)
>      {
>        gfc_free_omp_clauses (c);
>        return MATCH_ERROR;

Again, IMHO not needed, if you restore the old_loc into gfc_current_loc,
then gfc_match_omp_eos () will surely fail.

	Jakub

Re: OpenACC wait clause

[Cesar Philippidis]

[-- Attachment #1: Type: text/plain, Size: 2721 bytes --]

On 06/27/2016 12:23 PM, Jakub Jelinek wrote:
> On Mon, Jun 27, 2016 at 11:36:26AM -0700, Cesar Philippidis wrote:

>> @@ -630,9 +653,10 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>>  {
>>    gfc_omp_clauses *c = gfc_get_omp_clauses ();
>>    locus old_loc;
>> +  bool seen_error = false;
>>  
>>    *cp = NULL;
>> -  while (1)
>> +  while (!seen_error)
>>      {
>>        if ((first || gfc_match_char (',') != MATCH_YES)
>>  	  && (needs_space && gfc_match_space () != MATCH_YES))
> 
> Why?
> The main loop is while (1) which has break; as the last statement.
> Instead of setting seen_error, just set
> gfc_current_locus = old_loc;
> if you have already matched successfully something, and then break;
> as you already do.

I made that change.

>> @@ -1275,9 +1309,16 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>>  	    continue;
>>  	  if ((mask & OMP_CLAUSE_TILE)
>>  	      && !c->tile_list
>> -	      && match_oacc_expr_list ("tile (", &c->tile_list,
>> -				       true) == MATCH_YES)
>> -	    continue;
>> +	      && gfc_match ("tile") == MATCH_YES)
>> +	    {
>> +	      if (match_oacc_expr_list (" (", &c->tile_list, true) != MATCH_YES)
>> +		{
>> +		  seen_error = true;
>> +		  break;
>> +		}
>> +	      needs_space = true;
>> +	      continue;
>> +	    }
>>  	  if ((mask & OMP_CLAUSE_TO)
>>  	      && gfc_match_omp_variable_list ("to (",
>>  					      &c->lists[OMP_LIST_TO], false,
> 
> So, tile without ()s is also a valid clause in OpenACC?
> If yes, what do you set in c structure for the existence of the clause?
> If it is not valid, then the above change looks wrong.

No. The tile clause always was a () argument. So I should have used
gfc_match_omp_variable_list ("tile (", &c->tile_list) instead. This
patch fixes that.

>> @@ -1309,10 +1350,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>>  	      && gfc_match ("vector") == MATCH_YES)
>>  	    {
>>  	      c->vector = true;
>> -	      if (gfc_match (" ( length : %e )", &c->vector_expr) == MATCH_YES
>> -		  || gfc_match (" ( %e )", &c->vector_expr) == MATCH_YES)
>> -		needs_space = false;
>> -	      else
>> +	      match m = match_oacc_clause_gwv(c, GOMP_DIM_VECTOR);
> 
> Formatting, space before (.

Fixed.

>> @@ -1348,7 +1402,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>>        break;
>>      }
>>  
>> -  if (gfc_match_omp_eos () != MATCH_YES)
>> +  if (seen_error || gfc_match_omp_eos () != MATCH_YES)
>>      {
>>        gfc_free_omp_clauses (c);
>>        return MATCH_ERROR;
> 
> Again, IMHO not needed, if you restore the old_loc into gfc_current_loc,
> then gfc_match_omp_eos () will surely fail.

Fixed.

Is this ok for trunk and gcc6?

Cesar


[-- Attachment #2: fortran-parser-errors-20160627a.diff --]
[-- Type: text/x-patch, Size: 12280 bytes --]

2016-06-27  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/fortran/
	* openmp.c (match_oacc_clause_gang): Rename to ...
	(match_oacc_clause_gwv): this.  Add support for OpenACC worker and
	vector clauses.
	(gfc_match_omp_clauses): Use match_oacc_clause_gwv for
	OMP_CLAUSE_{GANG,WORKER,VECTOR}.  Propagate any MATCH_ERRORs for
	invalid OMP_CLAUSE_{ASYNC,WAIT,GANG,WORKER,VECTOR} clauses.
	(gfc_match_oacc_wait): Propagate MATCH_ERROR for invalid
	oacc_expr_lists.  Adjust the first and needs_space arguments to
	gfc_match_omp_clauses.

	gcc/testsuite/
	* gfortran.dg/goacc/asyncwait-2.f95: Updated expected diagnostics.
	* gfortran.dg/goacc/asyncwait-3.f95: Likewise.
	* gfortran.dg/goacc/asyncwait-4.f95: Add test coverage.

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index f514866..a691af9 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -396,43 +396,66 @@ cleanup:
 }
 
 static match
-match_oacc_clause_gang (gfc_omp_clauses *cp)
+match_oacc_clause_gwv (gfc_omp_clauses *cp, unsigned gwv)
 {
   match ret = MATCH_YES;
 
   if (gfc_match (" ( ") != MATCH_YES)
     return MATCH_NO;
 
-  /* The gang clause accepts two optional arguments, num and static.
-     The num argument may either be explicit (num: <val>) or
-     implicit without (<val> without num:).  */
-
-  while (ret == MATCH_YES)
+  if (gwv == GOMP_DIM_GANG)
     {
-      if (gfc_match (" static :") == MATCH_YES)
+        /* The gang clause accepts two optional arguments, num and static.
+	 The num argument may either be explicit (num: <val>) or
+	 implicit without (<val> without num:).  */
+
+      while (ret == MATCH_YES)
 	{
-	  if (cp->gang_static)
-	    return MATCH_ERROR;
+	  if (gfc_match (" static :") == MATCH_YES)
+	    {
+	      if (cp->gang_static)
+		return MATCH_ERROR;
+	      else
+		cp->gang_static = true;
+	      if (gfc_match_char ('*') == MATCH_YES)
+		cp->gang_static_expr = NULL;
+	      else if (gfc_match (" %e ", &cp->gang_static_expr) != MATCH_YES)
+		return MATCH_ERROR;
+	    }
 	  else
-	    cp->gang_static = true;
-	  if (gfc_match_char ('*') == MATCH_YES)
-	    cp->gang_static_expr = NULL;
-	  else if (gfc_match (" %e ", &cp->gang_static_expr) != MATCH_YES)
-	    return MATCH_ERROR;
-	}
-      else
-	{
-	  /* This is optional.  */
-	  if (cp->gang_num_expr || gfc_match (" num :") == MATCH_ERROR)
-	    return MATCH_ERROR;
-	  else if (gfc_match (" %e ", &cp->gang_num_expr) != MATCH_YES)
-	    return MATCH_ERROR;
+	    {
+	      /* This is optional.  */
+	      if (cp->gang_num_expr || gfc_match (" num :") == MATCH_ERROR)
+		return MATCH_ERROR;
+	      else if (gfc_match (" %e ", &cp->gang_num_expr) != MATCH_YES)
+		return MATCH_ERROR;
+	    }
+
+	  ret = gfc_match (" , ");
 	}
+    }
+  else if (gwv == GOMP_DIM_WORKER)
+    {
+      /* The 'num' arugment is optional.  */
+      if (gfc_match (" num :") == MATCH_ERROR)
+	return MATCH_ERROR;
+
+      if (gfc_match (" %e ", &cp->worker_expr) != MATCH_YES)
+	return MATCH_ERROR;
+    }
+  else if (gwv == GOMP_DIM_VECTOR)
+    {
+      /* The 'length' arugment is optional.  */
+      if (gfc_match (" length :") == MATCH_ERROR)
+	return MATCH_ERROR;
 
-      ret = gfc_match (" , ");
+      if (gfc_match (" %e ", &cp->vector_expr) != MATCH_YES)
+	return MATCH_ERROR;
     }
+  else
+    gfc_fatal_error ("Unexpected OpenACC parallelism.");
 
-  return gfc_match (" ) ");
+  return gfc_match (" )");
 }
 
 static match
@@ -677,14 +700,20 @@ 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)
+	      match m = gfc_match (" ( %e )", &c->async_expr);
+	      if (m == MATCH_ERROR)
+		{
+		  gfc_current_locus = old_loc;
+		  break;
+		}
+	      else if (m == MATCH_NO)
 		{
 		  c->async_expr
 		    = gfc_get_constant_expr (BT_INTEGER,
 					     gfc_default_integer_kind,
 					     &gfc_current_locus);
 		  mpz_set_si (c->async_expr->value.integer, GOMP_ASYNC_NOVAL);
+		  needs_space = true;
 		}
 	      continue;
 	    }
@@ -877,9 +906,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
 	      && gfc_match ("gang") == MATCH_YES)
 	    {
 	      c->gang = true;
-	      if (match_oacc_clause_gang(c) == MATCH_YES)
-		needs_space = false;
-	      else
+	      match m = match_oacc_clause_gwv (c, GOMP_DIM_GANG);
+	      if (m == MATCH_ERROR)
+		{
+		  gfc_current_locus = old_loc;
+		  break;
+		}
+	      else if (m == MATCH_NO)
 		needs_space = true;
 	      continue;
 	    }
@@ -1275,8 +1308,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
 	    continue;
 	  if ((mask & OMP_CLAUSE_TILE)
 	      && !c->tile_list
-	      && match_oacc_expr_list ("tile (", &c->tile_list,
-				       true) == MATCH_YES)
+	      && match_oacc_expr_list ("tile (", &c->tile_list, true)
+	         == MATCH_YES)
 	    continue;
 	  if ((mask & OMP_CLAUSE_TO)
 	      && gfc_match_omp_variable_list ("to (",
@@ -1309,10 +1342,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
 	      && gfc_match ("vector") == MATCH_YES)
 	    {
 	      c->vector = true;
-	      if (gfc_match (" ( length : %e )", &c->vector_expr) == MATCH_YES
-		  || gfc_match (" ( %e )", &c->vector_expr) == MATCH_YES)
-		needs_space = false;
-	      else
+	      match m = match_oacc_clause_gwv (c, GOMP_DIM_VECTOR);
+	      if (m == MATCH_ERROR)
+		{
+		  gfc_current_locus = old_loc;
+		  break;
+		}
+	      if (m == MATCH_NO)
 		needs_space = true;
 	      continue;
 	    }
@@ -1328,7 +1364,14 @@ 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);
+	      match m = match_oacc_expr_list (" (", &c->wait_list, false);
+	      if (m == MATCH_ERROR)
+		{
+		  gfc_current_locus = old_loc;
+		  break;
+		}
+	      else if (m == MATCH_NO)
+		needs_space = true;
 	      continue;
 	    }
 	  if ((mask & OMP_CLAUSE_WORKER)
@@ -1336,10 +1379,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
 	      && gfc_match ("worker") == MATCH_YES)
 	    {
 	      c->worker = true;
-	      if (gfc_match (" ( num : %e )", &c->worker_expr) == MATCH_YES
-		  || gfc_match (" ( %e )", &c->worker_expr) == MATCH_YES)
-		needs_space = false;
-	      else
+	      match m = match_oacc_clause_gwv (c, GOMP_DIM_WORKER);
+	      if (m == MATCH_ERROR)
+		{
+		  gfc_current_locus = old_loc;
+		  break;
+		}
+	      else if (m == MATCH_NO)
 		needs_space = true;
 	      continue;
 	    }
@@ -1595,15 +1641,18 @@ gfc_match_oacc_wait (void)
 {
   gfc_omp_clauses *c = gfc_get_omp_clauses ();
   gfc_expr_list *wait_list = NULL, *el;
+  bool space = true;
+  match m;
 
-  match_oacc_expr_list (" (", &wait_list, true);
-  gfc_match_omp_clauses (&c, OACC_WAIT_CLAUSES, false, false, true);
+  m = match_oacc_expr_list (" (", &wait_list, true);
+  if (m == MATCH_ERROR)
+    return m;
+  else if (m == MATCH_YES)
+    space = false;
 
-  if (gfc_match_omp_eos () != MATCH_YES)
-    {
-      gfc_error ("Unexpected junk in !$ACC WAIT at %C");
-      return MATCH_ERROR;
-    }
+  if (gfc_match_omp_clauses (&c, OACC_WAIT_CLAUSES, space, space, true)
+      == MATCH_ERROR)
+    return MATCH_ERROR;
 
   if (wait_list)
     for (el = wait_list; el; el = el->next)
diff --git a/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95 b/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95
index db0ce1f..7b2ae07 100644
--- a/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/asyncwait-2.f95
@@ -83,6 +83,18 @@ program asyncwait
   end do
   !$acc end parallel ! { dg-error "Unexpected \\\!\\\$ACC END PARALLEL" }
 
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) waitasync ! { dg-error "Unclassifiable OpenACC directive" }
+  do i = 1, N
+     b(i) = a(i)
+  end do
+  !$acc end parallel ! { dg-error "Unexpected \\\!\\\$ACC END PARALLEL" }
+
+  !$acc parallel copyin (a(1:N)) copy (b(1:N)) asyncwait ! { dg-error "Unclassifiable OpenACC directive" }
+  do i = 1, N
+     b(i) = a(i)
+  end do
+  !$acc end parallel ! { dg-error "Unexpected \\\!\\\$ACC END PARALLEL" }
+  
   !$acc parallel copyin (a(1:N)) copy (b(1:N)) wait
   do i = 1, N
      b(i) = a(i)
diff --git a/gcc/testsuite/gfortran.dg/goacc/asyncwait-3.f95 b/gcc/testsuite/gfortran.dg/goacc/asyncwait-3.f95
index 32c11de..ed72a9b 100644
--- a/gcc/testsuite/gfortran.dg/goacc/asyncwait-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/asyncwait-3.f95
@@ -11,17 +11,17 @@ program asyncwait
   a(:) = 3.0
   b(:) = 0.0
 
-  !$acc wait (1 2) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait (1 2) ! { dg-error "Syntax error in OpenACC expression list at" }
 
-  !$acc wait (1,) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait (1,) ! { dg-error "Syntax error in OpenACC expression list at" }
 
-  !$acc wait (,1) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait (,1) ! { dg-error "Syntax error in OpenACC expression list at" }
 
-  !$acc wait (1, 2, ) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait (1, 2, ) ! { dg-error "Syntax error in OpenACC expression list at" }
 
-  !$acc wait (1, 2, ,) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait (1, 2, ,) ! { dg-error "Syntax error in OpenACC expression list at" }
 
-  !$acc wait (1 ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait (1 ! { dg-error "Syntax error in OpenACC expression list at" }
 
   !$acc wait (1, *) ! { dg-error "Invalid argument to \\\$\\\!ACC WAIT" }
 
@@ -33,9 +33,9 @@ program asyncwait
 
   !$acc wait (1.0) ! { dg-error "WAIT clause at \\\(1\\\) requires a scalar INTEGER expression" }
 
-  !$acc wait 1 ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait 1 ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait N ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait N ! { dg-error "Unclassifiable OpenACC directive" }
 
   !$acc wait (1)
 end program asyncwait
diff --git a/gcc/testsuite/gfortran.dg/goacc/asyncwait-4.f95 b/gcc/testsuite/gfortran.dg/goacc/asyncwait-4.f95
index cd64ef3..df31154 100644
--- a/gcc/testsuite/gfortran.dg/goacc/asyncwait-4.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/asyncwait-4.f95
@@ -11,21 +11,21 @@ program asyncwait
   a(:) = 3.0
   b(:) = 0.0
 
-  !$acc wait async (1 2) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1 2) ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait async (1,) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1,) ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait async (,1) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (,1) ! { dg-error "Invalid character in name" }
 
-  !$acc wait async (1, 2, ) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1, 2, ) ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait async (1, 2, ,) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1, 2, ,) ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait async (1 ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1 ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait async (1, *) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1, *) ! { dg-error "Unclassifiable OpenACC directive" }
 
-  !$acc wait async (1, a) ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async (1, a) ! { dg-error "Unclassifiable OpenACC directive" }
 
   !$acc wait async (a) ! { dg-error "ASYNC clause at \\\(1\\\) requires a scalar INTEGER expression" }
 
@@ -33,5 +33,9 @@ program asyncwait
 
   !$acc wait async (1.0) ! { dg-error "ASYNC clause at \\\(1\\\) requires a scalar INTEGER expression" }
 
-  !$acc wait async 1 ! { dg-error "Unexpected junk in \\\!\\\$ACC WAIT at" }
+  !$acc wait async 1 ! { dg-error "Unclassifiable OpenACC directive" }
+
+  !$acc waitasync ! { dg-error "Unclassifiable OpenACC directive" }
+
+  !$acc wait,async ! { dg-error "Unclassifiable OpenACC directive" }
 end program asyncwait

Re: OpenACC wait clause

[Jakub Jelinek]

On Mon, Jun 27, 2016 at 04:22:01PM -0700, Cesar Philippidis wrote:
> +      while (ret == MATCH_YES)
>  	{
> -	  if (cp->gang_static)
> -	    return MATCH_ERROR;
> +	  if (gfc_match (" static :") == MATCH_YES)
> +	    {
> +	      if (cp->gang_static)
> +		return MATCH_ERROR;

It might be useful to gfc_error before this (i.e. follow a rule,
if you return MATCH_ERROR where MATCH_ERROR has not been returned before,
you emit gfc_error, if it has been returned from nested calls, you don't),
but it is not a big deal, worse case you get the cryptic generic error.
Let's keep it as is for now.

> +      /* The 'num' arugment is optional.  */

Typo, argument.

> +      if (gfc_match (" num :") == MATCH_ERROR)
> +	return MATCH_ERROR;

This is unnecessary (I mean the == MATCH_ERROR check), gfc_match for
no % operands in it will never return MATCH_ERROR (why would it?),
just MATCH_YES or MATCH_NO.

> +  else if (gwv == GOMP_DIM_VECTOR)
> +    {
> +      /* The 'length' arugment is optional.  */

See above.

> +      if (gfc_match (" length :") == MATCH_ERROR)

And once more.

> @@ -1275,8 +1308,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
>  	    continue;
>  	  if ((mask & OMP_CLAUSE_TILE)
>  	      && !c->tile_list
> -	      && match_oacc_expr_list ("tile (", &c->tile_list,
> -				       true) == MATCH_YES)
> +	      && match_oacc_expr_list ("tile (", &c->tile_list, true)
> +	         == MATCH_YES)

Why this hunk?  I think it is better to put the line break before the last
argument here, you don't have to decide if it shouldn't be surrounded by
	      && (match_oacc... (...)
		  == MATCH_YES))

Otherwise LGTM.

	Jakub