[OpenACC] Properly handle wait clause with no arguments

Hi, this patch properly handles OpenACC 'wait' clauses without arguments, making it an equivalent of "wait all".
(current trunk basically discards and ignores such argument-less wait clauses)  This adds additional handling in
the pack/unpack of the wait argument across the compiler/libgomp interface, but is done in a matter that
doesn't affect binary compatibility.

This patch was part of the OpenACC async re-work that was done on the gomp4 branch (later merged to OG7/OG8), see [1].
I'm separating this part out and submitting it first because it's logically independent.

[1] https://gcc.gnu.org/ml/gcc-patches/2017-06/msg01842.html

Re-tested with offloading to ensure no regressions, is this okay for trunk?

Thanks,
Chung-Lin

2018-08-30  Chung-Lin Tang  <cltang@codesourcery.com>

         gcc/c/
         * c-parser.c (c_parser_oacc_clause_wait): Add representation of wait
         clause without argument as 'wait (GOMP_ASYNC_NOVAL)', adjust comments.

         gcc/cp/
         * parser.c (cp_parser_oacc_clause_wait): Add representation of wait
         clause without argument as 'wait (GOMP_ASYNC_NOVAL)', adjust comments.

         gcc/fortran/
         * trans-openmp.c (gfc_trans_omp_clauses_1): Add representation of wait
         clause without argument as 'wait (GOMP_ASYNC_NOVAL)'.

         gcc/
         * omp-low.c (expand_omp_target): Add middle-end support for handling
         OMP_CLAUSE_WAIT clause with a GOMP_ASYNC_NOVAL(-1) as the argument.

         include/
         * gomp-constants.h (GOMP_LAUNCH_OP_MASK): Define.
         (GOMP_LAUNCH_PACK): Add bitwise-and of GOMP_LAUNCH_OP_MASK.
         (GOMP_LAUNCH_OP): Likewise.

         libgomp/
         * oacc-parallel.c (GOACC_parallel_keyed): Interpret launch op as
         signed 16-bit field, adjust num_waits handling.
         (GOACC_enter_exit_data): Adjust num_waits handling.
         (GOACC_update): Adjust num_waits handling.

Ping (adding Thomas to CC as OpenACC maintainer)

On 2018/8/30 9:27 PM, Chung-Lin Tang wrote:
> Hi, this patch properly handles OpenACC 'wait' clauses without arguments, making it an equivalent of "wait all".
> (current trunk basically discards and ignores such argument-less wait clauses)  This adds additional handling in
> the pack/unpack of the wait argument across the compiler/libgomp interface, but is done in a matter that
> doesn't affect binary compatibility.
> 
> This patch was part of the OpenACC async re-work that was done on the gomp4 branch (later merged to OG7/OG8), see [1].
> I'm separating this part out and submitting it first because it's logically independent.
> 
> [1] https://gcc.gnu.org/ml/gcc-patches/2017-06/msg01842.html
> 
> Re-tested with offloading to ensure no regressions, is this okay for trunk?
> 
> Thanks,
> Chung-Lin
> 
> 2018-08-30  Chung-Lin Tang  <cltang@codesourcery.com>
> 
>          gcc/c/
>          * c-parser.c (c_parser_oacc_clause_wait): Add representation of wait
>          clause without argument as 'wait (GOMP_ASYNC_NOVAL)', adjust comments.
> 
>          gcc/cp/
>          * parser.c (cp_parser_oacc_clause_wait): Add representation of wait
>          clause without argument as 'wait (GOMP_ASYNC_NOVAL)', adjust comments.
> 
>          gcc/fortran/
>          * trans-openmp.c (gfc_trans_omp_clauses_1): Add representation of wait
>          clause without argument as 'wait (GOMP_ASYNC_NOVAL)'.
> 
>          gcc/
>          * omp-low.c (expand_omp_target): Add middle-end support for handling
>          OMP_CLAUSE_WAIT clause with a GOMP_ASYNC_NOVAL(-1) as the argument.
> 
>          include/
>          * gomp-constants.h (GOMP_LAUNCH_OP_MASK): Define.
>          (GOMP_LAUNCH_PACK): Add bitwise-and of GOMP_LAUNCH_OP_MASK.
>          (GOMP_LAUNCH_OP): Likewise.
> 
>          libgomp/
>          * oacc-parallel.c (GOACC_parallel_keyed): Interpret launch op as
>          signed 16-bit field, adjust num_waits handling.
>          (GOACC_enter_exit_data): Adjust num_waits handling.
>          (GOACC_update): Adjust num_waits handling.

Hi Chung-Lin!

On Thu, 30 Aug 2018 21:27:22 +0800, Chung-Lin Tang <chunglin_tang@mentor.com> wrote:
> Hi, this patch properly handles OpenACC 'wait' clauses without arguments, making it an equivalent of "wait all".

Thanks!

> (current trunk basically discards and ignores such argument-less wait
> clauses)

Bugs should be filed, for later reference.  Now done:
<https://gcc.gnu.org/PR87924> "OpenACC wait clauses without
async-arguments".  (I couldn't put you in CC because "cltang@gcc.gnu.org
did not match anything"?)

> This adds additional handling in
> the pack/unpack of the wait argument across the compiler/libgomp interface, but is done in a matter that
> doesn't affect binary compatibility.

Hmm.  See below.  (Jakub, could you please review the last paragraph of
this email?)

> This patch was part of the OpenACC async re-work that was done on the gomp4 branch (later merged to OG7/OG8), see [1].
> I'm separating this part out and submitting it first because it's logically independent.
> 
> [1] https://gcc.gnu.org/ml/gcc-patches/2017-06/msg01842.html

Thanks for splitting it out!

> Re-tested with offloading to ensure no regressions, is this okay for trunk?

A few comments.

No test cases included.  I'm working on a few, will post/commit later.

>          gcc/c/
>          * c-parser.c (c_parser_oacc_clause_wait): Add representation of wait
>          clause without argument as 'wait (GOMP_ASYNC_NOVAL)', adjust comments.
> 
>          gcc/cp/
>          * parser.c (cp_parser_oacc_clause_wait): Add representation of wait
>          clause without argument as 'wait (GOMP_ASYNC_NOVAL)', adjust comments.
> 
>          gcc/fortran/
>          * trans-openmp.c (gfc_trans_omp_clauses_1): Add representation of wait
>          clause without argument as 'wait (GOMP_ASYNC_NOVAL)'.
> 
>          gcc/
>          * omp-low.c (expand_omp_target): Add middle-end support for handling
>          OMP_CLAUSE_WAIT clause with a GOMP_ASYNC_NOVAL(-1) as the argument.
> 
>          include/
>          * gomp-constants.h (GOMP_LAUNCH_OP_MASK): Define.
>          (GOMP_LAUNCH_PACK): Add bitwise-and of GOMP_LAUNCH_OP_MASK.
>          (GOMP_LAUNCH_OP): Likewise.
> 
>          libgomp/
>          * oacc-parallel.c (GOACC_parallel_keyed): Interpret launch op as
>          signed 16-bit field, adjust num_waits handling.
>          (GOACC_enter_exit_data): Adjust num_waits handling.
>          (GOACC_update): Adjust num_waits handling.

> --- gcc/c/c-parser.c	(revision 263981)
> +++ gcc/c/c-parser.c	(working copy)
> @@ -12719,7 +12719,7 @@ c_parser_oacc_clause_tile (c_parser *parser, tree
>  }
>  
>  /* OpenACC:
> -   wait ( int-expr-list ) */
> +   wait [( int-expr-list )] */
>  
>  static tree
>  c_parser_oacc_clause_wait (c_parser *parser, tree list)
> @@ -12728,7 +12728,15 @@ c_parser_oacc_clause_wait (c_parser *parser, tree
>  
>    if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
>      list = c_parser_oacc_wait_list (parser, clause_loc, list);
> +  else
> +    {
> +      tree c = build_omp_clause (clause_loc, OMP_CLAUSE_WAIT);
>  
> +      OMP_CLAUSE_DECL (c) = build_int_cst (integer_type_node, GOMP_ASYNC_NOVAL);
> +      OMP_CLAUSE_CHAIN (c) = list;
> +      list = c;
> +    }
> +
>    return list;
>  }

ACK.

> --- gcc/cp/parser.c	(revision 263981)
> +++ gcc/cp/parser.c	(working copy)
> @@ -32137,7 +32137,7 @@ cp_parser_oacc_wait_list (cp_parser *parser, locat
>  }
>  
>  /* OpenACC:
> -   wait ( int-expr-list ) */
> +   wait [( int-expr-list )] */
>  
>  static tree
>  cp_parser_oacc_clause_wait (cp_parser *parser, tree list)
> @@ -32144,10 +32144,16 @@ cp_parser_oacc_clause_wait (cp_parser *parser, tre
>  {
>    location_t location = cp_lexer_peek_token (parser->lexer)->location;
>  
> -  if (cp_lexer_peek_token (parser->lexer)->type != CPP_OPEN_PAREN)
> -    return list;
> +  if (cp_lexer_peek_token (parser->lexer)->type == CPP_OPEN_PAREN)
> +    list = cp_parser_oacc_wait_list (parser, location, list);
> +  else
> +    {
> +      tree c = build_omp_clause (location, OMP_CLAUSE_WAIT);
>  
> -  list = cp_parser_oacc_wait_list (parser, location, list);
> +      OMP_CLAUSE_DECL (c) = build_int_cst (integer_type_node, GOMP_ASYNC_NOVAL);
> +      OMP_CLAUSE_CHAIN (c) = list;
> +      list = c;
> +    }
>  
>    return list;
>  }

ACK.

> --- gcc/fortran/trans-openmp.c	(revision 263981)
> +++ gcc/fortran/trans-openmp.c	(working copy)
> @@ -2922,6 +2922,13 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp
>  	  omp_clauses = c;
>  	}
>      }
> +  else if (clauses->wait)
> +    {
> +      c = build_omp_clause (where.lb->location, OMP_CLAUSE_WAIT);
> +      OMP_CLAUSE_DECL (c) = build_int_cst (integer_type_node, GOMP_ASYNC_NOVAL);
> +      OMP_CLAUSE_CHAIN (c) = omp_clauses;
> +      omp_clauses = c;
> +    }
>    if (clauses->num_gangs_expr)
>      {
>        tree num_gangs_var

NACK.  Instead let's do the following, similar to C, C++, and also
similar to Fortran's OpenACC async clause handling without explicit
async-argument:

--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -1885,7 +1885,19 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 		  break;
 		}
 	      else if (m == MATCH_NO)
-		needs_space = true;
+		{
+		  gfc_expr *expr
+		    = gfc_get_constant_expr (BT_INTEGER,
+					     gfc_default_integer_kind,
+					     &gfc_current_locus);
+		  mpz_set_si (expr->value.integer, GOMP_ASYNC_NOVAL);
+		  gfc_expr_list **expr_list = &c->wait_list;
+		  while (*expr_list)
+		    expr_list = &(*expr_list)->next;
+		  *expr_list = gfc_get_expr_list ();
+		  (*expr_list)->expr = expr;
+		  needs_space = true;
+		}
 	      continue;
 	    }
 	  if ((mask & OMP_CLAUSE_WORKER)

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;

> --- include/gomp-constants.h	(revision 263981)
> +++ include/gomp-constants.h	(working copy)
> @@ -221,13 +221,14 @@ enum gomp_map_kind
>  #define GOMP_LAUNCH_CODE_SHIFT	28
>  #define GOMP_LAUNCH_DEVICE_SHIFT 16
>  #define GOMP_LAUNCH_OP_SHIFT 0
> +#define GOMP_LAUNCH_OP_MASK 0xffff
>  #define GOMP_LAUNCH_PACK(CODE,DEVICE,OP)	\
>    (((CODE) << GOMP_LAUNCH_CODE_SHIFT)		\
>     | ((DEVICE) << GOMP_LAUNCH_DEVICE_SHIFT)	\
> -   | ((OP) << GOMP_LAUNCH_OP_SHIFT))
> +   | (((OP) & GOMP_LAUNCH_OP_MASK) << GOMP_LAUNCH_OP_SHIFT))
>  #define GOMP_LAUNCH_CODE(X) (((X) >> GOMP_LAUNCH_CODE_SHIFT) & 0xf)
>  #define GOMP_LAUNCH_DEVICE(X) (((X) >> GOMP_LAUNCH_DEVICE_SHIFT) & 0xfff)
> -#define GOMP_LAUNCH_OP(X) (((X) >> GOMP_LAUNCH_OP_SHIFT) & 0xffff)
> +#define GOMP_LAUNCH_OP(X) (((X) >> GOMP_LAUNCH_OP_SHIFT) & GOMP_LAUNCH_OP_MASK)
>  #define GOMP_LAUNCH_OP_MAX 0xffff
>  
>  /* Bitmask to apply in order to find out the intended device of a target

> --- libgomp/oacc-parallel.c	(revision 263981)
> +++ libgomp/oacc-parallel.c	(working copy)
> @@ -194,10 +194,14 @@ GOACC_parallel_keyed (int device, void (*fn) (void
>  
>  	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);
>  	    break;
>  	  }
>  
> @@ -351,7 +355,7 @@ GOACC_enter_exit_data (int device, size_t mapnum,
>        || host_fallback)
>      return;
>  
> -  if (num_waits)
> +  if (num_waits > 0)
>      {
>        va_list ap;
>  
> @@ -359,6 +363,8 @@ GOACC_enter_exit_data (int device, size_t mapnum,
>        goacc_wait (async, num_waits, &ap);
>        va_end (ap);
>      }
> +  else if (num_waits == acc_async_noval)
> +    acc_wait_all_async (async);
>  
>    /* Determine whether "finalize" semantics apply to all mappings of this
>       OpenACC directive.  */
> @@ -542,7 +548,7 @@ GOACC_update (int device, size_t mapnum,
>        || host_fallback)
>      return;
>  
> -  if (num_waits)
> +  if (num_waits > 0)
>      {
>        va_list ap;
>  
> @@ -550,6 +556,8 @@ GOACC_update (int device, size_t mapnum,
>        goacc_wait (async, num_waits, &ap);
>        va_end (ap);
>      }
> +  else if (num_waits == acc_async_noval)
> +    acc_wait_all_async (async);
>  
>    acc_dev->openacc.async_set_async_func (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?


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?


Grüße
 Thomas

On Wed, Nov 07, 2018 at 08:13:29PM +0100, Thomas Schwinge wrote:
> 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?

Only old executables/libraries need to continue running correctly when
linking against new libgomp.  New programs against old libgomp might work,
or might not.

	Jakub

On 2018/11/8 3:13 AM, Thomas Schwinge wrote:
> NACK.  Instead let's do the following, similar to C, C++, and also
> similar to Fortran's OpenACC async clause handling without explicit
> async-argument:
> 
> --- gcc/fortran/openmp.c
> +++ gcc/fortran/openmp.c
> @@ -1885,7 +1885,19 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
>   		  break;
>   		}
>   	      else if (m == MATCH_NO)
> -		needs_space = true;
> +		{
> +		  gfc_expr *expr
> +		    = gfc_get_constant_expr (BT_INTEGER,
> +					     gfc_default_integer_kind,
> +					     &gfc_current_locus);
> +		  mpz_set_si (expr->value.integer, GOMP_ASYNC_NOVAL);
> +		  gfc_expr_list **expr_list = &c->wait_list;
> +		  while (*expr_list)
> +		    expr_list = &(*expr_list)->next;
> +		  *expr_list = gfc_get_expr_list ();
> +		  (*expr_list)->expr = expr;
> +		  needs_space = true;
> +		}
>   	      continue;
>   	    }
>   	  if ((mask & OMP_CLAUSE_WORKER)

Okay, I see what you mean.


> Now, why do we need the following changes, in this rather "convoluted"
> form:

>> +	      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.

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

(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)

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.

> 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.

Thanks,
Chung-Lin

Hi Chung-Lin!

On Wed, 07 Nov 2018 20:13:29 +0100, Thomas Schwinge <thomas@codesourcery.com> wrote:
> On Thu, 30 Aug 2018 21:27:22 +0800, Chung-Lin Tang <chunglin_tang@mentor.com> wrote:
> > Hi, this patch properly handles OpenACC 'wait' clauses without arguments, making it an equivalent of "wait all".

> > (current trunk basically discards and ignores such argument-less wait
> > clauses)
> 
> Bugs should be filed, for later reference.  Now done:
> <https://gcc.gnu.org/PR87924> "OpenACC wait clauses without
> async-arguments".  (I couldn't put you in CC because "cltang@gcc.gnu.org
> did not match anything"?)

This will, by the way, need to be fixed on all active release branches.


> No test cases included.  I'm working on a few, will post/commit later.

I thought I had also written a libgomp execution test case, during
travel/attending the SuperComputing 2018 conference, but I can't find it
right now...  ;-|

Anyway, with XFAILs (which you then please remove as part of your patch),
at least the following compile-time test cases now committed to trunk in
r266686:

commit 1d89613e77d7db420b13ce3ad8b98f07aaf474e8
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Fri Nov 30 20:39:30 2018 +0000

    [PR87924] Add (XFAILed) test cases for OpenACC wait clauses without async-arguments
    
            gcc/testsuite/
            PR c/87924
            * c-c++-common/goacc/asyncwait-5.c: Update.
            * gfortran.dg/goacc/asyncwait-5.f: Likewise.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@266686 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/testsuite/ChangeLog                        |  4 ++++
 gcc/testsuite/c-c++-common/goacc/asyncwait-5.c | 21 +++++++++++++++++++++
 gcc/testsuite/gfortran.dg/goacc/asyncwait-5.f  | 20 ++++++++++++++++++++
 3 files changed, 45 insertions(+)

diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog
index 75ca70b4af28..68186d8ab837 100644
--- gcc/testsuite/ChangeLog
+++ gcc/testsuite/ChangeLog
@@ -1,5 +1,9 @@
 2018-11-30  Thomas Schwinge  <thomas@codesourcery.com>
 
+	PR c/87924
+	* c-c++-common/goacc/asyncwait-5.c: Update.
+	* gfortran.dg/goacc/asyncwait-5.f: Likewise.
+
 	* c-c++-common/goacc/asyncwait-5.c: New file.
 	* gfortran.dg/goacc/asyncwait-5.f: Likewise.
 
diff --git gcc/testsuite/c-c++-common/goacc/asyncwait-5.c gcc/testsuite/c-c++-common/goacc/asyncwait-5.c
index fe6f8a0cf2da..80d4a8477b93 100644
--- gcc/testsuite/c-c++-common/goacc/asyncwait-5.c
+++ gcc/testsuite/c-c++-common/goacc/asyncwait-5.c
@@ -11,4 +11,25 @@ void f()
 #pragma acc parallel async (2) wait (11, 12) wait (13)
   ;
   /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(13\\) wait\\(12\\) wait\\(11\\) async\\(2\\)\$" 1 "original" } } */
+
+
+#pragma acc parallel async (3) wait
+  ;
+  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(-1\\) async\\(3\\)$" 1 "original" { xfail *-*-* } } } */
+
+#pragma acc parallel async (4) wait (100) wait
+  ;
+  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(-1\\) wait\\(100\\) async\\(4\\)$" 1 "original" { xfail *-*-* } } } */
+
+#pragma acc parallel async (5) wait wait (101)
+  ;
+  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(101\\) wait\\(-1\\) async\\(5\\)$" 1 "original" { xfail *-*-* } } } */
+
+#pragma acc parallel async (6) wait wait (102, 103) wait wait
+  ;
+  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(-1\\) wait\\(-1\\) wait\\(103\\) wait\\(102\\) wait\\(-1\\) async\\(6\\)$" 1 "original" { xfail *-*-* } } } */
+
+#pragma acc parallel async (7) wait (104) wait wait (105, 106)
+  ;
+  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(106\\) wait\\(105\\) wait\\(-1\\) wait\\(104\\) async\\(7\\)$" 1 "original" { xfail *-*-* } } } */
 }
diff --git gcc/testsuite/gfortran.dg/goacc/asyncwait-5.f gcc/testsuite/gfortran.dg/goacc/asyncwait-5.f
index 59b886343af6..7ad5813b8a03 100644
--- gcc/testsuite/gfortran.dg/goacc/asyncwait-5.f
+++ gcc/testsuite/gfortran.dg/goacc/asyncwait-5.f
@@ -10,4 +10,24 @@
 !$ACC END PARALLEL
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(2\\) wait\\(11\\) wait\\(12\\) wait\\(13\\)$" 1 "original" } }
 
+!$ACC PARALLEL ASYNC (3) WAIT
+!$ACC END PARALLEL
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(3\\) wait\\(-1\\)$" 1 "original" { xfail *-*-* } } }
+
+!$ACC PARALLEL ASYNC (4) WAIT (100) WAIT
+!$ACC END PARALLEL
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(4\\) wait\\(100\\) wait\\(-1\\)$" 1 "original" { xfail *-*-* } } }
+
+!$ACC PARALLEL ASYNC (5) WAIT WAIT (101)
+!$ACC END PARALLEL
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(5\\) wait\\(-1\\) wait\\(101\\)$" 1 "original" { xfail *-*-* } } }
+
+!$ACC PARALLEL ASYNC (6) WAIT WAIT (102, 103) WAIT WAIT
+!$ACC END PARALLEL
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(6\\) wait\\(-1\\) wait\\(102\\) wait\\(103\\) wait\\(-1\\) wait\\(-1\\)$" 1 "original" { xfail *-*-* } } }
+
+!$ACC PARALLEL ASYNC (7) WAIT (104) WAIT WAIT (105, 106)
+!$ACC END PARALLEL
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(7\\) wait\\(104\\) wait\\(-1\\) wait\\(105\\) wait\\(106\\)$" 1 "original" { xfail *-*-* } } }
+
       END


Grüße
 Thomas

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

Hi Thomas,
this version of the wait-clause-with-no-args patch revises the following:

(1) The way the Fortran FE parts are implemented, which essentially is your code.
(I'll reflect that in the final ChangeLog)

(2) Instead of trying to encode ACC_ASYNC_NOVAL into num_waits, I've followed
your suggestion to just treat it as a normal async. This means the gcc/omp-expand.c
parts in the last patch are discarded.

(3) Things in oacc-parallel.c have been mostly adjusted to only handle the wait(ACC_ASYNC_NOVAL)
case inside goacc_wait().

Hope this is now okay for trunk when appropriate.

Thanks,
Chung-Lin
Index: gcc/c/c-parser.c
===================================================================
--- gcc/c/c-parser.c	(revision 267913)
+++ gcc/c/c-parser.c	(working copy)
@@ -13410,7 +13410,7 @@ c_parser_oacc_clause_tile (c_parser *parser, tree
 }
 
 /* OpenACC:
-   wait ( int-expr-list ) */
+   wait [( int-expr-list )] */
 
 static tree
 c_parser_oacc_clause_wait (c_parser *parser, tree list)
@@ -13419,7 +13419,15 @@ c_parser_oacc_clause_wait (c_parser *parser, tree
 
   if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
     list = c_parser_oacc_wait_list (parser, clause_loc, list);
+  else
+    {
+      tree c = build_omp_clause (clause_loc, OMP_CLAUSE_WAIT);
 
+      OMP_CLAUSE_DECL (c) = build_int_cst (integer_type_node, GOMP_ASYNC_NOVAL);
+      OMP_CLAUSE_CHAIN (c) = list;
+      list = c;
+    }
+
   return list;
 }
 
Index: gcc/cp/parser.c
===================================================================
--- gcc/cp/parser.c	(revision 267913)
+++ gcc/cp/parser.c	(working copy)
@@ -32815,7 +32815,7 @@ cp_parser_oacc_wait_list (cp_parser *parser, locat
 }
 
 /* OpenACC:
-   wait ( int-expr-list ) */
+   wait [( int-expr-list )] */
 
 static tree
 cp_parser_oacc_clause_wait (cp_parser *parser, tree list)
@@ -32822,10 +32822,16 @@ cp_parser_oacc_clause_wait (cp_parser *parser, tre
 {
   location_t location = cp_lexer_peek_token (parser->lexer)->location;
 
-  if (cp_lexer_peek_token (parser->lexer)->type != CPP_OPEN_PAREN)
-    return list;
+  if (cp_lexer_peek_token (parser->lexer)->type == CPP_OPEN_PAREN)
+    list = cp_parser_oacc_wait_list (parser, location, list);
+  else
+    {
+      tree c = build_omp_clause (location, OMP_CLAUSE_WAIT);
 
-  list = cp_parser_oacc_wait_list (parser, location, list);
+      OMP_CLAUSE_DECL (c) = build_int_cst (integer_type_node, GOMP_ASYNC_NOVAL);
+      OMP_CLAUSE_CHAIN (c) = list;
+      list = c;
+    }
 
   return list;
 }
Index: gcc/fortran/openmp.c
===================================================================
--- gcc/fortran/openmp.c	(revision 267913)
+++ gcc/fortran/openmp.c	(working copy)
@@ -1885,7 +1885,19 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const
 		  break;
 		}
 	      else if (m == MATCH_NO)
-		needs_space = true;
+		{
+		  gfc_expr *expr
+		    = gfc_get_constant_expr (BT_INTEGER,
+					     gfc_default_integer_kind,
+					     &gfc_current_locus);
+		  mpz_set_si (expr->value.integer, GOMP_ASYNC_NOVAL);
+		  gfc_expr_list **expr_list = &c->wait_list;
+		  while (*expr_list)
+		    expr_list = &(*expr_list)->next;
+		  *expr_list = gfc_get_expr_list ();
+		  (*expr_list)->expr = expr;
+		  needs_space = true;
+		}
 	      continue;
 	    }
 	  if ((mask & OMP_CLAUSE_WORKER)
Index: libgomp/oacc-parallel.c
===================================================================
--- libgomp/oacc-parallel.c	(revision 267913)
+++ libgomp/oacc-parallel.c	(working copy)
@@ -206,9 +206,7 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (voi
 	case GOMP_LAUNCH_WAIT:
 	  {
 	    unsigned num_waits = GOMP_LAUNCH_OP (tag);
-
-	    if (num_waits)
-	      goacc_wait (async, num_waits, &ap);
+	    goacc_wait (async, num_waits, &ap);
 	    break;
 	  }
 
@@ -514,13 +512,20 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum,
 static void
 goacc_wait (int async, int num_waits, va_list *ap)
 {
-  struct goacc_thread *thr = goacc_thread ();
-  struct gomp_device_descr *acc_dev = thr->dev;
-
   while (num_waits--)
     {
       int qid = va_arg (*ap, int);
-      
+
+      /* Waiting on ACC_ASYNC_NOVAL maps to 'wait all'.  */
+      if (qid == acc_async_noval)
+	{
+	  if (async == acc_async_sync)
+	    acc_wait_all ();
+	  else
+	    acc_wait_all_async (async);
+	  break;
+	}
+
       if (acc_async_test (qid))
 	continue;
 
@@ -531,7 +536,7 @@ goacc_wait (int async, int num_waits, va_list *ap)
 	    launching on, the queue itself will order work as
 	    required, so there's no need to wait explicitly.  */
       else
-	acc_dev->openacc.async_wait_async_func (qid, async);
+	acc_wait_async (qid, async);
     }
 }

Hi Chung-Lin!

On Mon, 14 Jan 2019 21:56:05 +0800, Chung-Lin Tang <chunglin_tang@mentor.com> wrote:
> this version of the wait-clause-with-no-args patch revises the following:
> 
> (1) The way the Fortran FE parts are implemented, which essentially is your code.
> (I'll reflect that in the final ChangeLog)
> 
> (2) Instead of trying to encode ACC_ASYNC_NOVAL into num_waits, I've followed
> your suggestion to just treat it as a normal async. This means the gcc/omp-expand.c
> parts in the last patch are discarded.
> 
> (3) Things in oacc-parallel.c have been mostly adjusted to only handle the wait(ACC_ASYNC_NOVAL)
> case inside goacc_wait().

Thanks.

> Hope this is now okay for trunk when appropriate.

Given that this is a wrong-code generation bug fix, I do approve this in
the current GCC development stage (and will later backport to the other
release branches), but please include the following changes as discussed
before: mention PR87924 "OpenACC wait clauses without async-arguments" in
the ChangeLogs and commit log, and remove the XFAILs in
"c-c++-common/goacc/asyncwait-5.c" and "gfortran.dg/goacc/asyncwait-5.f".

To record the review effort, please include "Reviewed-by: Thomas Schwinge
<thomas@codesourcery.com>" in the commit log, see
<https://gcc.gnu.org/wiki/Reviewed-by>.


Grüße
 Thomas


> Index: gcc/c/c-parser.c
> ===================================================================
> --- gcc/c/c-parser.c	(revision 267913)
> +++ gcc/c/c-parser.c	(working copy)
> @@ -13410,7 +13410,7 @@ c_parser_oacc_clause_tile (c_parser *parser, tree
>  }
>  
>  /* OpenACC:
> -   wait ( int-expr-list ) */
> +   wait [( int-expr-list )] */
>  
>  static tree
>  c_parser_oacc_clause_wait (c_parser *parser, tree list)
> @@ -13419,7 +13419,15 @@ c_parser_oacc_clause_wait (c_parser *parser, tree
>  
>    if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
>      list = c_parser_oacc_wait_list (parser, clause_loc, list);
> +  else
> +    {
> +      tree c = build_omp_clause (clause_loc, OMP_CLAUSE_WAIT);
>  
> +      OMP_CLAUSE_DECL (c) = build_int_cst (integer_type_node, GOMP_ASYNC_NOVAL);
> +      OMP_CLAUSE_CHAIN (c) = list;
> +      list = c;
> +    }
> +
>    return list;
>  }
>  
> Index: gcc/cp/parser.c
> ===================================================================
> --- gcc/cp/parser.c	(revision 267913)
> +++ gcc/cp/parser.c	(working copy)
> @@ -32815,7 +32815,7 @@ cp_parser_oacc_wait_list (cp_parser *parser, locat
>  }
>  
>  /* OpenACC:
> -   wait ( int-expr-list ) */
> +   wait [( int-expr-list )] */
>  
>  static tree
>  cp_parser_oacc_clause_wait (cp_parser *parser, tree list)
> @@ -32822,10 +32822,16 @@ cp_parser_oacc_clause_wait (cp_parser *parser, tre
>  {
>    location_t location = cp_lexer_peek_token (parser->lexer)->location;
>  
> -  if (cp_lexer_peek_token (parser->lexer)->type != CPP_OPEN_PAREN)
> -    return list;
> +  if (cp_lexer_peek_token (parser->lexer)->type == CPP_OPEN_PAREN)
> +    list = cp_parser_oacc_wait_list (parser, location, list);
> +  else
> +    {
> +      tree c = build_omp_clause (location, OMP_CLAUSE_WAIT);
>  
> -  list = cp_parser_oacc_wait_list (parser, location, list);
> +      OMP_CLAUSE_DECL (c) = build_int_cst (integer_type_node, GOMP_ASYNC_NOVAL);
> +      OMP_CLAUSE_CHAIN (c) = list;
> +      list = c;
> +    }
>  
>    return list;
>  }
> Index: gcc/fortran/openmp.c
> ===================================================================
> --- gcc/fortran/openmp.c	(revision 267913)
> +++ gcc/fortran/openmp.c	(working copy)
> @@ -1885,7 +1885,19 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const
>  		  break;
>  		}
>  	      else if (m == MATCH_NO)
> -		needs_space = true;
> +		{
> +		  gfc_expr *expr
> +		    = gfc_get_constant_expr (BT_INTEGER,
> +					     gfc_default_integer_kind,
> +					     &gfc_current_locus);
> +		  mpz_set_si (expr->value.integer, GOMP_ASYNC_NOVAL);
> +		  gfc_expr_list **expr_list = &c->wait_list;
> +		  while (*expr_list)
> +		    expr_list = &(*expr_list)->next;
> +		  *expr_list = gfc_get_expr_list ();
> +		  (*expr_list)->expr = expr;
> +		  needs_space = true;
> +		}
>  	      continue;
>  	    }
>  	  if ((mask & OMP_CLAUSE_WORKER)
> Index: libgomp/oacc-parallel.c
> ===================================================================
> --- libgomp/oacc-parallel.c	(revision 267913)
> +++ libgomp/oacc-parallel.c	(working copy)
> @@ -206,9 +206,7 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (voi
>  	case GOMP_LAUNCH_WAIT:
>  	  {
>  	    unsigned num_waits = GOMP_LAUNCH_OP (tag);
> -
> -	    if (num_waits)
> -	      goacc_wait (async, num_waits, &ap);
> +	    goacc_wait (async, num_waits, &ap);
>  	    break;
>  	  }
>  
> @@ -514,13 +512,20 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum,
>  static void
>  goacc_wait (int async, int num_waits, va_list *ap)
>  {
> -  struct goacc_thread *thr = goacc_thread ();
> -  struct gomp_device_descr *acc_dev = thr->dev;
> -
>    while (num_waits--)
>      {
>        int qid = va_arg (*ap, int);
> -      
> +
> +      /* Waiting on ACC_ASYNC_NOVAL maps to 'wait all'.  */
> +      if (qid == acc_async_noval)
> +	{
> +	  if (async == acc_async_sync)
> +	    acc_wait_all ();
> +	  else
> +	    acc_wait_all_async (async);
> +	  break;
> +	}
> +
>        if (acc_async_test (qid))
>  	continue;
>  
> @@ -531,7 +536,7 @@ goacc_wait (int async, int num_waits, va_list *ap)
>  	    launching on, the queue itself will order work as
>  	    required, so there's no need to wait explicitly.  */
>        else
> -	acc_dev->openacc.async_wait_async_func (qid, async);
> +	acc_wait_async (qid, async);
>      }
>  }
>

Hi Chung-Lin!

On Tue, 22 Jan 2019 13:16:37 +0100, I wrote:
> On Mon, 14 Jan 2019 21:56:05 +0800, Chung-Lin Tang <chunglin_tang@mentor.com> wrote:
> > this version of the wait-clause-with-no-args patch revises the following:

> Thanks.
> 
> > Hope this is now okay for trunk when appropriate.
> 
> Given that this is a wrong-code generation bug fix, I do approve this in
> the current GCC development stage (and will later backport to the other
> release branches), but please include the following changes as discussed
> before: mention PR87924 "OpenACC wait clauses without async-arguments" in
> the ChangeLogs and commit log

Thanks for committing this, trunk r269016.

> and remove the XFAILs in
> "c-c++-common/goacc/asyncwait-5.c" and "gfortran.dg/goacc/asyncwait-5.f".

This I now have done myself, in trunk r269020 "[PR87924] OpenACC wait
clauses without async-arguments: remove XFAILs", attached.


Grüße
 Thomas
From 3eedad7ade57021a843447e8b6cba9619ac5b39e Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Tue, 19 Feb 2019 16:04:17 +0000
Subject: [PATCH] [PR87924] OpenACC wait clauses without async-arguments:
 remove XFAILs

... which the recent r269016 didn't do.

	gcc/testsuite/
	PR c/87924
	* c-c++-common/goacc/asyncwait-5.c: Remove XFAILs.
	* gfortran.dg/goacc/asyncwait-5.f: Likewise.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@269020 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/testsuite/ChangeLog                        |  6 ++++++
 gcc/testsuite/c-c++-common/goacc/asyncwait-5.c | 10 +++++-----
 gcc/testsuite/gfortran.dg/goacc/asyncwait-5.f  | 10 +++++-----
 3 files changed, 16 insertions(+), 10 deletions(-)

diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 5796a150c0d..1c91cec1a81 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,9 @@
+2019-02-19  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR c/87924
+	* c-c++-common/goacc/asyncwait-5.c: Remove XFAILs.
+	* gfortran.dg/goacc/asyncwait-5.f: Likewise.
+
 2019-02-19  H.J. Lu  <hongjiu.lu@intel.com>
 
 	PR target/89397
diff --git a/gcc/testsuite/c-c++-common/goacc/asyncwait-5.c b/gcc/testsuite/c-c++-common/goacc/asyncwait-5.c
index 80d4a8477b9..f4ba48a1833 100644
--- a/gcc/testsuite/c-c++-common/goacc/asyncwait-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/asyncwait-5.c
@@ -15,21 +15,21 @@ void f()
 
 #pragma acc parallel async (3) wait
   ;
-  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(-1\\) async\\(3\\)$" 1 "original" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(-1\\) async\\(3\\)$" 1 "original" } } */
 
 #pragma acc parallel async (4) wait (100) wait
   ;
-  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(-1\\) wait\\(100\\) async\\(4\\)$" 1 "original" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(-1\\) wait\\(100\\) async\\(4\\)$" 1 "original" } } */
 
 #pragma acc parallel async (5) wait wait (101)
   ;
-  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(101\\) wait\\(-1\\) async\\(5\\)$" 1 "original" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(101\\) wait\\(-1\\) async\\(5\\)$" 1 "original" } } */
 
 #pragma acc parallel async (6) wait wait (102, 103) wait wait
   ;
-  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(-1\\) wait\\(-1\\) wait\\(103\\) wait\\(102\\) wait\\(-1\\) async\\(6\\)$" 1 "original" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(-1\\) wait\\(-1\\) wait\\(103\\) wait\\(102\\) wait\\(-1\\) async\\(6\\)$" 1 "original" } } */
 
 #pragma acc parallel async (7) wait (104) wait wait (105, 106)
   ;
-  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(106\\) wait\\(105\\) wait\\(-1\\) wait\\(104\\) async\\(7\\)$" 1 "original" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel wait\\(106\\) wait\\(105\\) wait\\(-1\\) wait\\(104\\) async\\(7\\)$" 1 "original" } } */
 }
diff --git a/gcc/testsuite/gfortran.dg/goacc/asyncwait-5.f b/gcc/testsuite/gfortran.dg/goacc/asyncwait-5.f
index 7ad5813b8a0..89cd7923590 100644
--- a/gcc/testsuite/gfortran.dg/goacc/asyncwait-5.f
+++ b/gcc/testsuite/gfortran.dg/goacc/asyncwait-5.f
@@ -12,22 +12,22 @@
 
 !$ACC PARALLEL ASYNC (3) WAIT
 !$ACC END PARALLEL
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(3\\) wait\\(-1\\)$" 1 "original" { xfail *-*-* } } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(3\\) wait\\(-1\\)$" 1 "original" } }
 
 !$ACC PARALLEL ASYNC (4) WAIT (100) WAIT
 !$ACC END PARALLEL
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(4\\) wait\\(100\\) wait\\(-1\\)$" 1 "original" { xfail *-*-* } } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(4\\) wait\\(100\\) wait\\(-1\\)$" 1 "original" } }
 
 !$ACC PARALLEL ASYNC (5) WAIT WAIT (101)
 !$ACC END PARALLEL
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(5\\) wait\\(-1\\) wait\\(101\\)$" 1 "original" { xfail *-*-* } } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(5\\) wait\\(-1\\) wait\\(101\\)$" 1 "original" } }
 
 !$ACC PARALLEL ASYNC (6) WAIT WAIT (102, 103) WAIT WAIT
 !$ACC END PARALLEL
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(6\\) wait\\(-1\\) wait\\(102\\) wait\\(103\\) wait\\(-1\\) wait\\(-1\\)$" 1 "original" { xfail *-*-* } } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(6\\) wait\\(-1\\) wait\\(102\\) wait\\(103\\) wait\\(-1\\) wait\\(-1\\)$" 1 "original" } }
 
 !$ACC PARALLEL ASYNC (7) WAIT (104) WAIT WAIT (105, 106)
 !$ACC END PARALLEL
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(7\\) wait\\(104\\) wait\\(-1\\) wait\\(105\\) wait\\(106\\)$" 1 "original" { xfail *-*-* } } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel async\\(7\\) wait\\(104\\) wait\\(-1\\) wait\\(105\\) wait\\(106\\)$" 1 "original" } }
 
       END